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

    , gcf
    , uvw[bl]
    , vis[bl]
    );
}


#include <math_functions.h>

template <
    int timesteps
  , int channels
  >
__inline__ __device__ void loadVisIntoSharedMem (
    const Double4c vis[timesteps * channels]
  , Double4c vis_shared[timesteps * channels]
  ) {
  for (int i = threadIdx.x; i < timesteps * channels; i += blockDim.x) {
    vis_shared[i] = vis[i];
  }
}

template <
    int over
  , int grid_size

