Feature/fudgeShift

cuda/shift.go

@@ -15,6 +15,16 @@ func ShiftX(dst, src *data.Slice, shiftX int, clampL, clampR float32) {
 	k_shiftx_async(dst.DevPtr(0), src.DevPtr(0), N[X], N[Y], N[Z], shiftX, clampL, clampR, cfg)
 }
 
+// shift dst by shx cells (positive or negative) along X-axis.
+// new edge value is the current value at the border.
+func ShiftEdgeCarryX(dst, src *data.Slice, shiftX int) {
+	util.Argument(dst.NComp() == 1 && src.NComp() == 1)
+	util.Assert(dst.Len() == src.Len())
+	N := dst.Size()
+	cfg := make3DConf(N)
+	k_shiftedgecarryX_async(dst.DevPtr(0), src.DevPtr(0), N[X], N[Y], N[Z], shiftX, cfg)
+}
+
 func ShiftY(dst, src *data.Slice, shiftY int, clampL, clampR float32) {
 	util.Argument(dst.NComp() == 1 && src.NComp() == 1)
 	util.Assert(dst.Len() == src.Len())
@@ -23,6 +33,14 @@ func ShiftY(dst, src *data.Slice, shiftY int, clampL, clampR float32) {
 	k_shifty_async(dst.DevPtr(0), src.DevPtr(0), N[X], N[Y], N[Z], shiftY, clampL, clampR, cfg)
 }
 
+func ShiftEdgeCarryY(dst, src *data.Slice, shiftY int) {
+	util.Argument(dst.NComp() == 1 && src.NComp() == 1)
+	util.Assert(dst.Len() == src.Len())
+	N := dst.Size()
+	cfg := make3DConf(N)
+	k_shiftedgecarryY_async(dst.DevPtr(0), src.DevPtr(0), N[X], N[Y], N[Z], shiftY, cfg)
+}
+
 func ShiftZ(dst, src *data.Slice, shiftZ int, clampL, clampR float32) {
 	util.Argument(dst.NComp() == 1 && src.NComp() == 1)
 	util.Assert(dst.Len() == src.Len())

cuda/shiftedgecarryx.cu

@@ -0,0 +1,25 @@
+#include "stencil.h"
+
+// shift dst by shx cells (positive or negative) along X-axis.
+// new edge value is the current edge value.
+extern "C" __global__ void
+shiftedgecarryX(float* __restrict__  dst, float* __restrict__  src,
+    int Nx,  int Ny,  int Nz, int shx) {
+
+    int ix = blockIdx.x * blockDim.x + threadIdx.x;
+    int iy = blockIdx.y * blockDim.y + threadIdx.y;
+    int iz = blockIdx.z * blockDim.z + threadIdx.z;
+
+    if(ix < Nx && iy < Ny && iz < Nz) {
+        int ix2 = ix-shx;
+        float newval;
+        if (ix2 < 0) {
+            newval = src[idx(0, iy, iz)];
+        } else if (ix2 >= Nx) {
+            newval = src[idx(Nx-1, iy, iz)];
+        } else {
+            newval = src[idx(ix2, iy, iz)];
+        }
+        dst[idx(ix, iy, iz)] = newval;
+    }
+}

cuda/shiftedgecarryy.cu

@@ -0,0 +1,25 @@
+#include "stencil.h"
+
+// shift dst by shy cells (positive or negative) along Y-axis.
+// new edge value is the current edge value.
+extern "C" __global__ void
+shiftedgecarryY(float* __restrict__  dst, float* __restrict__  src,
+    int Nx,  int Ny,  int Nz, int shy) {
+
+    int ix = blockIdx.x * blockDim.x + threadIdx.x;
+    int iy = blockIdx.y * blockDim.y + threadIdx.y;
+    int iz = blockIdx.z * blockDim.z + threadIdx.z;
+
+    if(ix < Nx && iy < Ny && iz < Nz) {
+        int iy2 = iy-shy;
+        float newval;
+        if (iy2 < 0) {
+            newval = src[idx(ix, 0, iz)];
+        } else if (iy2 >= Ny) {
+            newval = src[idx(ix, Ny-1, iz)];
+        } else {
+            newval = src[idx(ix, iy2, iz)];
+        }
+        dst[idx(ix, iy, iz)] = newval;
+    }
+}

engine/shift.go

@@ -9,10 +9,12 @@ var (
 	TotalShift, TotalYShift                    float64                        // accumulated window shift (X and Y) in meter
 	ShiftMagL, ShiftMagR, ShiftMagU, ShiftMagD data.Vector                    // when shifting m, put these value at the left/right edge.
 	ShiftM, ShiftGeom, ShiftRegions            bool        = true, true, true // should shift act on magnetization, geometry, regions?
+	EdgeCarryShift                             bool        = false            // Use the values of M at the border for the new cells
 )
 
 func init() {
 	DeclFunc("Shift", Shift, "Shifts the simulation by +1/-1 cells along X")
+	DeclVar("EdgeCarryShift", &EdgeCarryShift, "Whether to use the current magnetization at the border for the cells inserted by Shift")
 	DeclVar("ShiftMagL", &ShiftMagL, "Upon shift, insert this magnetization from the left")
 	DeclVar("ShiftMagR", &ShiftMagR, "Upon shift, insert this magnetization from the right")
 	DeclVar("ShiftMagU", &ShiftMagU, "Upon shift, insert this magnetization from the top")
@@ -47,7 +49,11 @@ func shiftMag(m *data.Slice, dx int) {
 	defer cuda.Recycle(m2)
 	for c := 0; c < m.NComp(); c++ {
 		comp := m.Comp(c)
-		cuda.ShiftX(m2, comp, dx, float32(ShiftMagL[c]), float32(ShiftMagR[c]))
+		if EdgeCarryShift {
+			cuda.ShiftEdgeCarryX(m2, comp, dx)
+		} else {
+			cuda.ShiftX(m2, comp, dx, float32(ShiftMagL[c]), float32(ShiftMagR[c]))
+		}
 		data.Copy(comp, m2) // str0 ?
 	}
 }
@@ -72,7 +78,11 @@ func shiftMagY(m *data.Slice, dy int) {
 	defer cuda.Recycle(m2)
 	for c := 0; c < m.NComp(); c++ {
 		comp := m.Comp(c)
-		cuda.ShiftY(m2, comp, dy, float32(ShiftMagU[c]), float32(ShiftMagD[c]))
+		if EdgeCarryShift {
+			cuda.ShiftEdgeCarryY(m2, comp, dy)
+		} else {
+			cuda.ShiftY(m2, comp, dy, float32(ShiftMagU[c]), float32(ShiftMagD[c]))
+		}
 		data.Copy(comp, m2) // str0 ?
 	}
 }

test/edgeCarryShift.mx3

@@ -0,0 +1,16 @@
+SetGridSize(128, 64, 1)
+SetCellSize(4e-9, 4e-9, 4e-9)
+Msat = 1e6
+Aex = 10e-12
+alpha = 1
+m = RandomMag()
+steps(300)
+
+edgeCarryShift = true
+shift(32)
+ExpectV("edgeCarryShift = true; The added cells during shift are the same as the old border?", m.getCell(0, 20, 0), m.getCell(32, 20, 0), 1e-6)
+
+edgeCarryShift = false
+ShiftMagL = vector(0, 0, 1)
+shift(32)
+ExpectV("edgeCarryShift = false; The added cells during shift are ShiftMagL?", m.getCell(0, 20, 0), ShiftMagL, 1e-6)