 #include "metrix.h"
 #include "OskarBinReader.h"
 
-template <int over>
-__global__ void
+template <int over, bool isbig> struct degridder {};
+
+template <int over> struct degridder<over, true>{
+
+static __device__ __inline__
 //__launch_bounds__(256, 6)
-// double2 is for 'in'
-degrid_kernel(complexd* out, const double3* in, size_t npts, const complexd* img, 
+// double2 is enough for 'in'
+void 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

  , int timesteps_x_channels
  , int img_dim
  ) {
  for(int bl = 0; bl < baselines; bl++) {
    // Always use (sorted) permutations vector
    bl = permutations[bl].bl;
    int gcf_dim = get_supp(permutations[bl].wp);
    // No mirrored GCF yet
    if (gcf_dim <= 16) {
       degrid_kernel_small_gcf<over>
         <<<timesteps_x_channels/32,dim3(32,32)>>>(
             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
           );
    } 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
          );
    }
  }
}

// Test instantiation
extern "C" void testtest(
    const BlWMap permutations[/* baselines */]
  , complexd* out_vis
  // Padded, points to the beginning

  }
};

template <
    int max_gcf_dim
  , int over
  , bool isBig
  >
__global__
void degridGPU(
    const BlWMap permutations[/* baselines */]
  , complexd * out_vis
  // Padded, points to the beginning
  // of physical data
  , const double3 * uvw
  // centered in 0 w-plane
  , const complexd * gcf[]
  , const complexd * img
  , size_t timesteps_x_channels
  , size_t img_dim
  , int blOff
  ) {
  // Temporarily use the second grid dimension as baseline dimension
  int bl = permutations[blockIdx.y + blOff].bl;
  int gcf_dim = get_supp(permutations[bl].wp);
  degridder<over, isBig>::degrid_kernel(
             out_vis + bl*timesteps_x_channels
           , 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
           );
}


// We assume permutations vector
// is sorted against abs(w_plane_index) compare function
template <
    int max_gcf_dim
  , int over
