 __global__ void
 //__launch_bounds__(256, 6)
 // double2 is for 'in'
-degrid_kernel(complexd* out, double3* in, size_t npts, complexd* img, 
-                              size_t img_dim, complexd* gcf, int gcf_dim) {
+degrid_kernel(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) {

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

    } else {
       degrid_kernel<over>
         <<<timesteps_x_channels/32,dim3(32,8)>>>(
             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
