   mesh[xr][yr] = r2;
 }
 
-// test instantiation
-template __device__
-void ucs_common<256>(cuDoubleComplex mesh[513][513], double t2);
-
-
+extern "C" __global__ void r2(cuDoubleComplex mesh[513][513], double t2) {
+  ucs_common<256>(mesh, t2);
+}
+
+
+#if 0
 template <int max_half_support>
 __device__
 void calc_inplace(

  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]
  );

//
__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__
void normalize(
    double norm
  , cuDoubleComplex v[(max_half_support * 2 + 1) * (max_half_support * 2 + 1)]
  ) {

  dst[xr][yr] = res;
}

extern "C" __global__ void wkernff(
    cuDoubleComplex dst[513][513]
  , const cuDoubleComplex src[513][513]
  , double w
  ){
  calc<256>(dst, src, w);
}


template <
    int max_half_support
  , int oversample
  >
__device__ __inline__
void copy_ucs_2_over(
    cuDoubleComplex dst[(max_half_support * 2 + 1) * oversample][(max_half_support * 2 + 1) * oversample]
  , const cuDoubleComplex src[max_half_support * 2 + 1][max_half_support * 2 + 1]
  ){
  const int dst_center = (max_half_support * 2 + 1) * oversample / 2;
  __SET_MAP
  dst[dst_center - x][dst_center - y] = src[xl][yl];
