         {0,1,2,4}  // opposite 7
 };
 
+
+static __constant__ struct work work;
+
 template<
     unsigned BX,
     unsigned BY,

                r[idim]=s;
            }

            work.dst[idst]=tex3D(in,r[0],r[1],r[2]);
        }
    }
}

__global__
void
__launch_bounds__(BX*BY,1) // max threads, min blocks
resample_k(void) {
    const unsigned ox=threadIdx.x+blockIdx.x*BX;
    const unsigned oy=(threadIdx.y+blockIdx.y*BY)*WORK;
    const unsigned idst0=ox+oy*BX;
