   , int timesteps
   , int channels
   , bool is_half_gcf
+  >
+__device__ __inline__ void addBaselineToPsfi(
+    double scale
+  , double wstep
+  , int max_supp
+  , complexd psfi[grid_size][grid_size]
+  , const complexd * gcf[]
+  , const double3 uvw[timesteps * channels]
+  , const Double4c vis[timesteps * channels]
+  ) {
+  // FIXME: should share shared memory with grid computation
+  // Also that perhaps was a bug (we had each instantiation
+  // had it's own shared memory copy).
+  // FIXME: investigate this futher!
+  __shared__ Pregridded uvo_shared[timesteps * channels];
+  __shared__ int2 off_shared[timesteps * channels];
+  
+  loadUVWIntoSharedMem<
+      over
+    , grid_size
+    , timesteps
+    , channels
+    , is_half_gcf
+    >(scale
+    , wstep
+    , max_supp
+    , uvw
+    , uvo_shared
+    , off_shared
+    );
+  syncthreads();
+  psfiKernel_scatter_kernel<
+      grid_size
+    , timesteps
+    , channels
+    , is_half_gcf
+    >(max_supp
+    , gcf
+    , psfi
+    , uvo_shared
+    , off_shared
+  );
+}
+
+template <
+    int over
+  , int grid_size
+
+  , int timesteps
+  , int channels
+  , bool is_half_gcf
   , bool use_permutations
   >
 __device__ __inline__ void addBaselinesToGrid(

#include <math_functions.h>

template <
    int over
  , int grid_size


  , int channels
  , bool do_mirror
  >
__inline__ __device__ void loadUVWIntoSharedMem (
    double scale
  , double wstep
  , int max_supp // needed for u v translation
  , const double3 uvw[timesteps * channels]
  , Pregridded uvo_shared[timesteps * channels]
  , int2 off_shared[timesteps * channels]
  ) {
  for (int i = threadIdx.x; i < timesteps * channels; i += blockDim.x) {
    // uvo_shared[i] is passed by reference and updated!
    pregridPoint<grid_size, over, do_mirror>(scale, wstep, uvw[i], uvo_shared[i]);
    off_shared[i].x = uvo_shared[i].u % max_supp;
    off_shared[i].y = uvo_shared[i].v % max_supp;
  }
