 // Perhaps it's not very cache friendly, but it is
 // very simple to perform work-distribution for this variant
 
+#define __SET_MAP                               \
+  const int                                     \
+      x = blockIdx.x * blockDim.x + threadIdx.x \
+    , y = blockIdx.y * blockDim.y + threadIdx.y \
+    , xl = max_half_support - x                 \
+    , xr = max_half_support + x                 \
+    , yl = max_half_support - y                 \
+    , yr = max_half_support + y                 \
+    ;
+
+
 template <int max_half_support>
 __device__
 void ucs_common(
     cuDoubleComplex mesh[max_half_support * 2 + 1][max_half_support * 2 + 1]
   , double t2
   ){
-  const int
-      x = blockIdx.x * blockDim.x + threadIdx.x
-    , y = blockIdx.y * blockDim.y + threadIdx.y
-    ;
+  __SET_MAP
   double
       sc = double(max_half_support)
     , xs = double(x) / sc * t2

void cut_out(
    int supp
  , cuDoubleComplex * dst
  , const cuDoubleComplex src[max_half_support * oversample * 2 + 1][max_half_support * oversample * 2 + 1]
  ) {
  const int
      x = blockIdx.x * blockDim.x + threadIdx.x

void cut_out<256,8>(
    int supp
  , cuDoubleComplex * dst
  , const cuDoubleComplex src[4104][4104]
  );
#endif

template <
    int max_half_support
  , int oversample
  >
__device__
void extract_over(
    int overx
  , int overy
  , cuDoubleComplex dst[max_half_support * 2 + 1][max_half_support * 2 + 1]
  , const cuDoubleComplex src[(max_half_support * 2 + 1) * oversample][(max_half_support * 2 + 1) * oversample]
  ) {
  __SET_MAP
  const int
      sxl = xl * oversample + overx
    , sxr = xr * oversample + overx
    , syl = yl * oversample + overy
    , syr = yr * oversample + overy
    ;
  dst[xl][yl] = src[sxl][syl];
  dst[xl][yr] = src[sxl][syr];
  dst[xr][yl] = src[sxr][syl];
  dst[xr][yr] = src[sxr][syr];
  }

// test instantiation
template __device__
void extract_over<256,8>(
    int overx
  , int overy
  , cuDoubleComplex dst[513][513]
  , const cuDoubleComplex src[4104][4104]
  );
