   int gcf_layer_offset;
 };
 
+/*
+  Suitable only for max_supp <= 32.
+  Must be launched as:
+  ...
+  dim3 blockDim(max_supp, max_supp);
+  gridKernel_scatter_small<<<baselines, blockDim ... >>>
+  ...
+ */
 template <
     int w_planes
   , int max_supp

}

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

  >
__global__
// grid must be initialized to 0s.
void gridKernel_scatter_small(
    const complexd gcf[w_planes][over][over][max_supp][max_supp]
  , Double4c grid[grid_size][grid_size]
  , const Pregridded uvw[baselines][timesteps * channels]
  , const Double4c vis[baselines][timesteps * channels]
  ) {
  int
      myU = threadIdx.x
    , myV = threadIdx.y
