 void gridKernel_scatter_small(
     const complexd gcf[w_planes][over][over][max_supp][max_supp]
   , Double4c grid[grid_size][grid_size]
-  , const double3 uvw[baselines][timesteps * channels]
+  , const w_traits<w_planes>::uvw_type uvw[baselines][timesteps * channels]
   , const Double4c vis[baselines][timesteps * channels]
   ) {
   __shared__ Pregridded uvo_shared[timesteps * channels];

void gridKernel_scatter(
    const complexd gcf[w_planes][over][over][max_supp][max_supp]
  , Double4c grid[grid_size][grid_size]
  , const double3 uvw[baselines][timesteps * channels]
  , const Double4c vis[baselines][timesteps * channels]
  ) {
  __shared__ Pregridded uvo_shared[timesteps * channels];

  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
