New feature: Progress towards running Blip on CPU

spm · Mar 7, 2024 · 97583c2 · 97583c2
1 parent c5b1d5e
commit 97583c2
Show file tree

Hide file tree

Showing 6 changed files with 113 additions and 15 deletions.
diff --git a/Artifacts.toml b/Artifacts.toml
@@ -1,19 +1,19 @@
 [[pp_lib]]
 arch = "x86_64"
 os = "linux"
-git-tree-sha1 = "368835caf6eb5017e11c2181aec8c46452d44441"
+git-tree-sha1 = "3d5b880c83373c8d9d91e3ee3a57b49d67d15b73"
 lazy = true
     [[pp_lib.download]]
-    sha256 = "d0c5124ba78beb35706dc44e8747a88731e5f9e91da70dbd2837300b26dcebfd"
+    sha256 = "4e642e478f9f2ae485a287ebf01f3ba0af135545a17a567dea5bff5930a697cc"
     #url = "file:///home/john/work/julia/pp_lib.tar.gz"
     url = "https://raw.githubusercontent.com/spm/PushPull.jl/main/artifacts/pp_lib.tar.gz"
 
 [[pp_lib]]
 arch = "x86_64"
 os = "windows"
-git-tree-sha1 = "7c43d07a531ec2432ca5673f96cb037d7e16b6c9"
+git-tree-sha1 = "3d5b880c83373c8d9d91e3ee3a57b49d67d15b73"
 lazy = true
     [[pp_lib.download]]
-    sha256 = "d0c5124ba78beb35706dc44e8747a88731e5f9e91da70dbd2837300b26dcebfd"
+    sha256 = "4e642e478f9f2ae485a287ebf01f3ba0af135545a17a567dea5bff5930a697cc"
     url = "https://raw.githubusercontent.com/spm/PushPull.jl/main/artifacts/pp_lib.tar.gz"
 
diff --git a/artifacts/C/Makefile b/artifacts/C/Makefile
@@ -3,15 +3,17 @@ include Makefile.var
 
 #PTXFILES = pushpull.ptx lbessi.ptx operator.ptx shootfun.ptx TVdenoise3d.ptx TVdenoise2d.ptx TVdenoise2d_old.ptx
 PTXFILES = pushpull.ptx lbessi.ptx sparse_operator.ptx TVdenoise3d.ptx TVdenoise2d.ptx blip.ptx
-LIBFILES = pushpull.$(SOSUF) TVdenoise3d.$(SOSUF) sparse_operator.$(SOSUF)
-
+LIBFILES = pushpull.$(SOSUF) TVdenoise3d.$(SOSUF) sparse_operator.$(SOSUF) blip.$(SOSUF)
+PPLIB    = ../pp_lib.tar.gz
 #CC       = x86_64-w64-mingw32-gcc-win32
 
 all: $(PTXFILES) $(LIBFILES)
 
 install: all $(PTXDIR) $(LIBDIR)
 	$(COPY) $(PTXFILES) $(PTXDIR)
 	$(COPY) $(LIBFILES) $(LIBDIR)
+	tar -C .. -cvf - lib | gzip -v > $(PPLIB)
+	julia get_sha.jl $(PPLIB)
 
 clean:
 	$(DEL) $(PTXFILES) $(LIBFILES)
@@ -20,9 +22,6 @@ pushpull.$EXT : pushpull.c pushpull_dev.cu
 
 blip.ptx: blip.cu blip_dev.cu
 
-noise:
-	echo $(LIBDIR)
-
 
 #pushpull.$EXT : pushpull.c pushpull_dev.cu
 #	$(CC) -Wall -fPIC -O3 -c pushpull.c
@@ -44,7 +43,6 @@ noise:
 #	$(CC) -shared -Wl,-soname,TVdenoise3.so -o TVdenoise3.so TVdenoise3.o
 
 %.so : %.c
-	echo $(PLATFORM)
 	$(CC) -O4 -shared -fPIC -Wl,-soname,$@ -o $@ -lm $<
 
 %.dll : %.c
@@ -54,7 +52,6 @@ noise:
 #	gcc -O4 -shared -fPIC -Wl,-soname,$@ -o $@ -lm $<
 
 %.ptx : %.cu
-	echo $(PLATFORM)
 	$(NVCC) -ptx --ptxas-options --verbose -O3 $< 
 
 $(PTXDIR) :

diff --git a/artifacts/C/blip.c b/artifacts/C/blip.c
@@ -0,0 +1,90 @@
+#include "cuheader.h"
+#include "blip_dev.cu"
+
+#define min(a,b) ((a)<(b) ? (a) : (b))
+#define max(a,b) ((a)>(b) ? (a) : (b))
+
+
+void blip_pad(const USIZE_t i_start, const USIZE_t j_start, const USIZE_t k_start,
+              const USIZE_t i_stop,  const USIZE_t j_stop,  const USIZE_t k_stop,
+              float *u, const USIZE_t *d, const float *g,
+              const float *aa, const float *bb, const float *ab,
+              const float *s)
+{
+    USIZE_t i0, j0, k0;
+    for(k0=k_start; k0<k_stop; k0++)
+        for(j0=j_start; j0<j_stop; j0++)
+            for(i0=i_start; i0<i_stop; i0++)
+            {
+                USIZE_t i, j, k;
+                for(k=k0; k<k_stop; k+=5)
+                    for(j=j0; j<j_stop; j+=5)
+                        for(i=i0; i<i_stop; i+=5)
+                            blip_dev(i, j, k, d, u, g, aa, bb, ab, s);
+            }
+}
+
+
+void blip_nopad(const USIZE_t i_start, const USIZE_t j_start, const USIZE_t k_start,
+                const USIZE_t i_stop,  const USIZE_t j_stop,  const USIZE_t k_stop,
+                float *u, const USIZE_t *d, const float *g,
+                const float *aa, const float *bb, const float *ab,
+                const float *s)
+{
+    USIZE_t i0, j0, k0;
+    for(k0=k_start; k0<k_stop; k0++)
+        for(j0=j_start; j0<j_stop; j0++)
+            for(i0=i_start; i0<i_stop; i0++)
+            {
+                USIZE_t i, j, k;
+                for(k=k0; k<k_stop; k+=5)
+                {
+                    USIZE_t ok = k*d[1];
+                    for(j=j0; j<j_stop; j+=5)
+                    {
+                        USIZE_t oj = d[0]*(j + ok); 
+                        for(i=i0; i<i_stop; i+=5)
+                        {
+                            USIZE_t oi = oj + i;
+                            blip_nopad_dev(d, u+oi, g+oi, aa+oi, bb+oi, ab+oi, s);
+                        }
+                    }
+                }
+            }
+}
+
+
+void blip(float *u, const USIZE_t *d, const float *g,
+          const float *aa, const float *bb, const float *ab,
+          const float *s)
+{
+    USIZE_t i0, i1, j0, j1, k0, k1;
+    i0 = min(3,  d[0]);
+    j0 = min(3,  d[1]);
+    k0 = min(3,  d[2]);
+    i1 = max(i0, d[0]-3);
+    j1 = max(j0, d[1]-3);
+    k1 = max(k0, d[2]-3);
+    blip_pad(0 ,  0,  0, d[0], d[1],   k0, u, d, g, aa, bb, ab, s);
+    blip_pad(0,   0, k0, d[0],   j0,   k1, u, d, g, aa, bb, ab, s);
+    blip_pad(0,  j0, k0,   i0,   j1,   k1, u, d, g, aa, bb, ab, s);
+    blip_nopad(i0, j0, k0, i1,   j1,   k1, u, d, g, aa, bb, ab, s);
+    blip_pad(i1, j0, k0, d[0],   j1,   k1, u, d, g, aa, bb, ab, s);
+    blip_pad(0,  j1, k0, d[0], d[1],   k1, u, d, g, aa, bb, ab, s);
+    blip_pad(0,   0, k1, d[0], d[1], d[2], u, d, g, aa, bb, ab, s);
+}
+
+
+void hu(float *r, const USIZE_t *d, const float *u,
+                   const float *aa, const float *bb, const float *ab, const float *s)
+{
+    USIZE_t i, j, k;
+    for(k=0; k<d[2]; k++)
+        for(j=0; j<d[1]; j++)
+            for(i=0; i<d[0]; i++)
+            {
+                USIZE_t ijk = i+d[0]*(j+d[1]*k);
+                r[ijk]      = hu_dev(i, j, k, d, u, aa, bb, ab, s);
+            }
+}
+
diff --git a/artifacts/C/blip.cu b/artifacts/C/blip.cu
@@ -1,10 +1,6 @@
 #define CUDA
 #include "cuheader.h"
 #include "blip_dev.cu"
-/*
-    MAXD - maximum filter size
-    MAXN - maximum number of gradient fields
-*/
 
 __global__ void blip(float *u, const USIZE_t *d, const float *g,
                      const float *aa, const float *bb, const float *ab,

diff --git a/artifacts/C/get_sha.jl b/artifacts/C/get_sha.jl
@@ -0,0 +1,15 @@
+using Tar, Inflate, SHA
+filename = "../pp_lib.tar.gz"
+if length(ARGS)>=1
+    filename = ARGS[1]
+end
+sha1    = Tar.tree_hash(IOBuffer(inflate_gzip(filename)))
+sha2 = bytes2hex(open(sha256, filename))
+println("Edit the Artifacts.toml using the following:\n")
+println("git-tree-sha1 = \"", sha1, "\"");
+println("    sha256 = \"", sha2, "\"");
+
+println("\nor run the following (in Unix):")
+artif = "../../Artifacts.toml"
+println("\ncp ", artif, " Artifacts.prev ; sed < Artifacts.prev '/git-tree-sha1/s/\"[0-9,a-c].*\"/\"", sha1,"\"/' | sed '/sha256/s/\"[0-9,a-c].*\"/\"",sha2,"\"/' > ", artif, "\n")
+
diff --git a/artifacts/pp_lib.tar.gz b/artifacts/pp_lib.tar.gz