   int koffset = swapkernel ? kernel_w*kernel_h-1 : 0;
 
   // generate offsets according to block/thread ids
+#ifdef CUDA_ORDER_FOR_COALESCING
+  int xx_start = threadIdx.x; // * patch_w;
+  int xx_end = output_w;
+  int xx_step = blockDim.x;
+#else
   int xx_start = threadIdx.x * patch_w;
   int xx_end = xx_start + patch_w; if (xx_end > output_w) xx_end = output_w;
+  int xx_step = 1;
+#endif
   int yy_start = threadIdx.y * patch_h;
   int yy_end = yy_start + patch_h; if (yy_end > output_h) yy_end = output_h;
+  int yy_step = 1;
   int oo_start = blockIdx.x;
   int oo_end = oo_start+1;
   int ii_start = 0;

    // convolution loop
    for(oo = oo_start; oo < oo_end; oo++) {
      for(ii = ii_start; ii < ii_end; ii++) {
        for(yy = yy_start; yy < yy_end; yy++) {
          for(xx = xx_start; xx < xx_end; xx++) {
            // Dot product in two dimensions... (between input image and the mask)
            float *input_p = input + ii*input_h*input_w + yy*stride_h*input_w + xx*stride_w;
            float *output_p = output + oo*output_h*output_w + yy*output_w + xx;

 */

#define CUDA_SHARED_MEM_SIZE (4*1024-32) // this is given by nVidia: max shared mem per block
#define CUDA_ORDER_FOR_COALESCING        // this reorders memory reads to enable coalescing

/*
 * Description:
