     , grid
     , uvo_shared
     , vis_shared
+    , off_shared
   );
 }
 

  , const Double4c vis[timesteps * channels]
  , Pregridded uvo_shared[timesteps * channels]
  , Double4c vis_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, w_planes, do_mirror>(scale, wstep, uvw[i], uvo_shared[i]);
    vis_shared[i] = vis[i];
  }
}


  , const complexd * gcf[]
  , Double4c grid[grid_size][grid_size]
  , const Pregridded uvo_shared[timesteps * channels]
  , const Double4c vis_shared[timesteps * channels]
  , const int2 off_shared[timesteps * channels]
  ) {
  for (int i = threadIdx.x; i < max_supp * max_supp; i += blockDim.x) {
    int
