 }
 
 // Test instantiation
-template __global__ void gridKernel_scatter<32, 64, 2048, 8, 50*99, 20, 1>(
-    const complexd gcf[32][8][8][64][64]
+template __global__ void gridKernel_scatter<1, 64, 2048, 8, 50*99, 20, 1>(
+    const complexd gcf[1][8][8][64][64]
   , Double4c grid[2048][2048]
-  , const double3 uvw[(50*90)][20 * 1]
+  , const w_traits<1>::uvw_type uvw[(50*90)][20 * 1]
   , const Double4c vis[(50*90)][20 * 1]
   );

template __global__ void gridKernel_scatter_small<32, 64, 2048, 8, 50*99, 20, 1>(
    const complexd gcf[32][8][8][64][64]
  , Double4c grid[2048][2048]
  , const double3 uvw[(50*90)][20 * 1]
  , const Double4c vis[(50*90)][20 * 1]
  );


  int gcf_layer_offset;
};

template <int n> struct w_traits {
  typedef double3 uvw_type;
  static const int wplanes = n;
  static int get_wplane(const uvw_type & coords) {
    return __double2int_rz(coords.z);
  }
};

template <> struct w_traits<1> {
  typedef double2 uvw_type;
  static const int wplanes = 1;
  static int get_wplane(const uvw_type & coords) {
    return 0;
  }
};


template <
    int w_planes
  , int over
