Feature: Begun CPU support for Blip

spm · Mar 6, 2024 · 2ff7bb3 · 2ff7bb3
1 parent 4ab9237
commit 2ff7bb3
Show file tree

Hide file tree

Showing 5 changed files with 52 additions and 46 deletions.
diff --git a/README.md b/README.md
@@ -1,3 +1,3 @@
 # PushPull
 
-[![Build Status](https://github.com/JohnAshburner/PushPull.jl/actions/workflows/CI.yml/badge.svg?branch=main)](https://github.com/JohnAshburner/PushPull.jl/actions/workflows/CI.yml?query=branch%3Amain)
+[![Build Status](https://github.com/spm/PushPull.jl/actions/workflows/CI.yml/badge.svg?branch=main)](https://github.com/spm/PushPull.jl/actions/workflows/CI.yml?query=branch%3Amain)
diff --git a/artifacts/C/Makefile b/artifacts/C/Makefile
@@ -1,5 +1,6 @@
 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)
@@ -19,6 +20,8 @@ pushpull.$EXT : pushpull.c pushpull_dev.cu
 
 blip.ptx: blip.cu blip_dev.cu
 
+noise:
+	echo $(LIBDIR)
 
 
 #pushpull.$EXT : pushpull.c pushpull_dev.cu
@@ -41,6 +44,7 @@ blip.ptx: blip.cu blip_dev.cu
 #	$(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
@@ -50,6 +54,7 @@ blip.ptx: blip.cu blip_dev.cu
 #	gcc -O4 -shared -fPIC -Wl,-soname,$@ -o $@ -lm $<
 
 %.ptx : %.cu
+	echo $(PLATFORM)
 	$(NVCC) -ptx --ptxas-options --verbose -O3 $< 
 
 $(PTXDIR) :

diff --git a/artifacts/C/Makefile.var b/artifacts/C/Makefile.var
@@ -39,30 +39,28 @@ WARNFLAG     = -Wall -Wextra -Wpedantic
 USE_OPENMP  ?= 0
 SOSUF        = so
 
+
 ifndef PLATFORM
   PLATFORM   = $(shell $(UNAME))
 endif
 
 ##### Linux #####
 ifeq (Linux,$(PLATFORM))
-  MEXEXT     = mexa64
   MEXBIN    ?= mex
   #MEXOPTS   += CFLAGS='$$CFLAGS $(WARNFLAG)'
-  NVCC       = nvcc.exe -m64
+  NVCC       = nvcc -m64
   PTXDIR     = ../lib/ptxa64
   LIBDIR     = ../lib/liba64
   CC         = gcc
 endif
 
 ##### macOS #####
 ifeq (Darwin,$(PLATFORM))
-  MEXEXT     = mexmaci64
   MEXBIN    ?= mex
   # https://stackoverflow.com/questions/37362414/
   OMPFLAG    = -fopenmp=libiomp5
 endif
 ifeq (arm64,$(PLATFORM))
-  MEXEXT     = mexmaca64
   MEXBIN    ?= mex
   # https://stackoverflow.com/questions/37362414/
   OMPFLAG    = -fopenmp=libiomp5
@@ -71,7 +69,6 @@ endif
 
 ##### Windows #####
 ifeq (MINGW32,$(word 1,$(subst _, ,$(PLATFORM)))) # MSVC
-  MEXEXT     = mexw32
   PTXDIR     = ../lib/ptxw32
   LIBDIR     = ../lib/libw32
   MEXBIN    ?= cmd /c "mex.bat
@@ -86,7 +83,6 @@ ifeq (MINGW32,$(word 1,$(subst _, ,$(PLATFORM)))) # MSVC
   NVCC       = nvcc.exe -m32
 endif
 ifeq (MINGW64,$(word 1,$(subst _, ,$(PLATFORM)))) # MSVC
-  MEXEXT     = mexw64
   PTXDIR     = ../lib/ptxw64
   LIBDIR     = ../lib/libw64
   MEXBIN    ?= mex
@@ -97,7 +93,6 @@ ifeq (MINGW64,$(word 1,$(subst _, ,$(PLATFORM)))) # MSVC
   NVCC       = nvcc.exe -m64
 endif
 ifeq (MSYS,$(word 1,$(subst _, ,$(PLATFORM)))) # GCC
-  MEXEXT     = mexw64
   PTXDIR     = ../lib/ptxw64
   LIBDIR     = ../lib/libw64
   MEXBIN    ?= mex
@@ -106,46 +101,11 @@ ifeq (MSYS,$(word 1,$(subst _, ,$(PLATFORM)))) # GCC
   MOSUF      = obj
 endif
 
-#### Octave ####
-ifeq (octave,$(PLATFORM))
-  MEXEXT     = mex
-  MEXBIN    ?= mkoctfile
-  MEXOPTS    = --mex -DOCTAVE_MEX_FILE
-  #MEXOPTS   += $(WARNFLAG)
-  override PLATFORM = $(shell $(UNAME))
-  ifeq (MINGW64,$(word 1,$(subst _, ,$(PLATFORM))))
-    MEXOPTS += -DSPM_WIN32
-  endif
-  ifeq (MSYS,$(word 1,$(subst _, ,$(PLATFORM))))
-    MEXOPTS += -DSPM_WIN32
-  endif
-  OMPFLAG    =
-endif
-
-#### Otherwise ####
-ifndef MEXEXT
-  $(error Unknowm platform $(PLATFORM))
-endif
-
-ifeq (1,$(USE_OPENMP))
-  ifneq ($(OMPFLAG),)
-    MEXOPTS += CFLAGS='$$CFLAGS $(OMPFLAG)' LDFLAGS='$$LDFLAGS $(OMPFLAG)'
-  endif
-endif
-MEX          = $(MEXBIN) $(MEXOPTS)
-SUF          = $(MEXEXT)
-
-MATLABROOT   = $(realpath $(shell which $(firstword $(MEXBIN))))
-ifeq (mex,$(MEXEXT))
-  SPMEXE    ?= ../bin/spm-octave
-else
-  SPMEXE    ?= ../bin/spm-matlab
-endif
-
 define verb
 	@ echo "_____________________________________________________________"
 	@ echo ""
 	@ echo "        " $(1)
 	@ echo "_____________________________________________________________"
 	@ echo ""
 endef
+
diff --git a/artifacts/C/blip_dev.cu b/artifacts/C/blip_dev.cu
@@ -148,6 +148,47 @@ __device__ void blip_dev(USIZE_t i, USIZE_t j, USIZE_t k, const USIZE_t *d, floa
     u[ii] += t/(w0*1.000001);
 }
 
+__device__ void blip_nopad_dev(const USIZE_t *d, float *u,
+                                const float *g, const float *aa, const float *bb, const float *ab, const float *s)
+{
+    float t, v0 = s[0], v1 = s[1], v2 = s[2], w, w0, sv = v0+v1+v2, u0;
+    SSIZE_t o4, o5;
+
+    o4  = d[0];
+    o5  = 2*o4;
+    u0  = u[0];
+    t   = g[0] - (((u[-o5]-u0)* bb[-o4]        + (u[o5]-u0)* bb[o4])/(-4)
+                + ((u[-o4]-u0)*(ab[-o4]-ab[0]) + (u[o4]-u0)*(ab[0]-ab[o4]))/2
+                +          u0 *(aa[0]+(ab[-o4]-ab[o4])/2));
+    w0  = aa[0] + (bb[-o4] + bb[o4])/4;
+
+    w0 -= 2*(w = -4*v1*sv);
+    t  -= ((u[-o4]-u0)+(u[o4]-u0))*w;
+    w0 -= 2*(w = v1*v1);
+    t  -= ((u[-o5]-u0)+(u[o5]-u0))*w;
+
+    o5  = d[0]*d[1];
+    w0 -= 4*(w = 2*v0*v2);
+    t  -= ((u[-o5-1]-u0) + (u[-o5+1]-u0) + (u[o5-1]-u0) + (u[o5+1]-u0))*w;
+    w0 -= 4*(w = 2*v0*v1);
+    t  -= ((u[-o4-1]-u0) + (u[-o4+1]-u0) + (u[o4-1]-u0) + (u[o4+1]-u0))*w;
+    w0 -= 4*(w = 2*v1*v2);
+    t  -= ((u[-o4-o5]-u0) + (u[-o4+o5]-u0) + (u[o4-o5]-u0) + (u[o4+o5]-u0))*w;
+
+    w0 -= 2*(w = -4*v2*sv);
+    t  -= ((u[-o5]-u0) + (u[o5]-u0))*w;
+    o5  = 2*o5;
+    w0 -= 2*(w = v0*v0);
+    t  -= ((u[-2]-u0) + (u[2]-u0))*w;
+    w0 -= 2*(w = v2*v2);
+    t  -= ((u[-o5]-u0) + (u[o5]-u0))*w;
+    w0 -= 4*(w = 2*v0*v1);
+    t  -= ((u[-o4-1]-u0) + (u[-o4+1]-u0) + (u[o4-1]-u0) + (u[o4+1]-u0))*w;
+    w0 -= 2*(w = -4*v0*sv);
+    t  -= ((u[-1]-u0) + (u[1]-u0))*w;
+    u[0] += t/(w0*1.000001);
+}
+
 
 __device__ float hu_dev(USIZE_t i, USIZE_t j, USIZE_t k, const USIZE_t *d, const float *u,
                         const float *aa, const float *bb, const float *ab, const float *s)

diff --git a/src/locations.jl b/src/locations.jl
@@ -55,12 +55,12 @@ function basedir()
     if true
         ## Use the following for determining SHA values for Artifacts.toml
         #using Tar, Inflate, SHA
-        #filename = "pp.tar.gz";
+        #filename = "pp_lib.tar.gz";
         #println("sha256 = \"", bytes2hex(open(sha256, filename)), "\"");
         #println("    git-tree-sha1 = \"", Tar.tree_hash(IOBuffer(inflate_gzip(filename))), "\"");
 
         # This is the path to the Artifacts.toml we will manipulate
-        artifact_toml = joinpath(@__DIR__, "..", "Artifacts.toml")
+        artifact_toml = find_artifacts_toml(@__FILE__)
 
         # Query the `Artifacts.toml` file for the hash bound to the name "pp_lib"
         # (returns `nothing` if no such binding exists)