   >
 __global__
 // grid must be initialized to 0s.
-void gridKernel_scatter(
+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]
   ) {
-  // We use 2-dimensional threadblocks (max_supp x max_supp)
-  // but the idea is entirely the same as Romein's
-  // FIXME: We are limited to 32x32 gcf layers thus. Add extra grid or block dimension
-  //   to account for this!!!
   int
       myU = threadIdx.x
     , myV = threadIdx.y

}

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

  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
