   , int channels
   >
 __inline__ __device__ void loadIntoSharedMem (
-    const double3 uvw[baselines][timesteps * channels]
+    const w_traits<w_planes>::uvw_type uvw[baselines][timesteps * channels]
   , const Double4c vis[baselines][timesteps * channels]
   , Pregridded uvo_shared[timesteps * channels]
   , Double4c vis_shared[timesteps * channels]
   ) {
   int bl = blockIdx.x;
   for (int i = threadIdx.x; i < timesteps * channels; i += blockDim.x) {
-    double3 coords = uvw[bl][i];
+    w_traits<w_planes>::uvw_type coords = uvw[bl][i];
     int
         u = __double2int_rz(coords.x)
       , v = __double2int_rz(coords.y)
       , over_u = __double2int_rz(over * (coords.x - u))
       , over_v = __double2int_rz(over * (coords.y - v))
-      , wplane = __double2int_rz(coords.z)
       ;
-    uvo_shared[i] = {short(u), short(v), (wplane * over + over_u) * over + over_v};
+    uvo_shared[i] = {short(u), short(v), (w_traits<w_planes>::get_wplane(coords.z) * over + over_u) * over + over_v};
     vis_shared[i] = vis[bl][i];
   }
 }

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


}

// Test instantiation
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 w_traits<1>::uvw_type uvw[(50*90)][20 * 1]
  , const Double4c vis[(50*90)][20 * 1]
  );
