diff --git a/cuda/shift.go b/cuda/shift.go index a9ad7e886..6f2b8186f 100644 --- a/cuda/shift.go +++ b/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()) diff --git a/cuda/shiftedgecarryx.cu b/cuda/shiftedgecarryx.cu new file mode 100644 index 000000000..2e952c796 --- /dev/null +++ b/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; + } +} diff --git a/cuda/shiftedgecarryy.cu b/cuda/shiftedgecarryy.cu new file mode 100644 index 000000000..1d52df68b --- /dev/null +++ b/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; + } +} diff --git a/engine/shift.go b/engine/shift.go index 20fea1433..b6338071b 100644 --- a/engine/shift.go +++ b/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 ? } } diff --git a/test/edgeCarryShift.mx3 b/test/edgeCarryShift.mx3 new file mode 100644 index 000000000..a2145fd18 --- /dev/null +++ b/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)