 
 
 template <int max_half_support>
-__device__
+__device__ __inline__
 void ucs_common(
     cuDoubleComplex mesh[max_half_support * 2 + 1][max_half_support * 2 + 1]
   , double t2
   ){
   __SET_MAP
   double
-      sc = double(max_half_support)
-    , xs = double(x) / sc * t2
-    , ys = double(y) / sc * t2
+      t2_div_sc = t2 / double(max_half_support)
+    , xs = double(x) * t2_div_sc
+    , ys = double(y) * t2_div_sc
     ;
   cuDoubleComplex r2 = make_cuDoubleComplex(xs * xs + ys * ys, 0.0);
 

  dst[dst_center + x][dst_center + y] = src[xr][yr];
}

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


template <
    int max_half_support
  , int oversample
  >
__device__
void extract_over(
    int overx
  , int overy

  dst[xr][yr] = src[sxr][syr];
  }


extern "C" __global__ void wextract0(
    int overx
  , int overy
  , cuDoubleComplex dst[513][513]
  , const cuDoubleComplex src[4104][4104]
  ){
  extract_over<256,8>(overx, overy, dst, src);
}


//
__device__ static __inline__ cuDoubleComplex cuMulComplexByDouble(cuDoubleComplex v,
                                                             double y){
  return make_cuDoubleComplex ( v.x * y
                              , v.y * y
                              );
}

// The work-distribution scheme here is very different from those above (and below).
template <
    int max_half_support
  >
__device__ __inline__
void normalize_kernel(
    double norm
  , cuDoubleComplex v[(max_half_support * 2 + 1) * (max_half_support * 2 + 1)]
  ) {
