 }
 
 template <
+    int grid_size
+
+  , int timesteps
+  , int channels
+  , bool is_half_gcf
+  >
+__inline__ __device__
+// grid must be initialized to 0s.
+void psfiKernel_scatter_kernel(
+    int max_supp
+  , const complexd * gcf[]
+  , complexd psfi[grid_size][grid_size]
+  , const Pregridded uvo_shared[timesteps * channels]
+  , const int2 off_shared[timesteps * channels]
+  ) {
+  for (int i = threadIdx.x; i < max_supp * max_supp; i += blockDim.x) {
+    int
+        myU = i % max_supp
+      , myV = i / max_supp
+      ;
+    psfiKernel_scatter_kernel_small<
+      grid_size
+
+    , timesteps
+    , channels
+    , is_half_gcf
+    > (max_supp, gcf, psfi, uvo_shared, off_shared, myU, myV);
+  }
+}
+
+template <
     int over
   , int grid_size
 

#include <math_functions.h>

template <
    int over
  , int grid_size


    , wstep
    , max_supp
    , uvw
    , uvo_shared
    , off_shared
    );
  loadVisIntoSharedMem<
      timesteps
    , channels
    >(vis
    , vis_shared
    );
  syncthreads();
  gridKernel_scatter_kernel<
