     if (gcf_dim <= 16) {
        degrid_kernel_small_gcf<over>
          <<<timesteps_x_channels/32,dim3(32,32)>>>(
-             out
+             out_vis
            , uvw+bl*timesteps_x_channels, timesteps_x_channels
            , img+degridPadder<max_gcf_dim>::bare_offset(img_dim), img_dim
            , gcf[permutations[bl].wp], gcf_dim

      complexd tmp;
      tmp.x = sum_r;
      tmp.y = sum_i;
      if (threadIdx.x == 0) {
         out[n+q] = tmp;
      }
   }

template <int over>
__global__ void
//__launch_bounds__(256, 6)
degrid_kernel_small_gcf(complexd* out, const double3* in, size_t npts, const complexd* img, 
                              size_t img_dim, const complexd* gcf, int gcf_dim) {
   
   //TODO remove hard-coded 32
   for (int n = 32*blockIdx.x; n<npts; n+= 32*gridDim.x) {
