 
   }
 
+  // sync
   cudaThreadSynchronize();
 
   // clean up
   THCudaTensor_free(input);
   THCudaTensor_free(kernel);
 
+  // check potential errors
   cudaError errcode = cudaGetLastError();
   if(errcode != cudaSuccess)
     THError(cudaGetErrorString(errcode));

  int koffset = swapkernel ? kernel_w*kernel_h-1 : 0;

  // generate offsets according to block/thread ids
  int xx_start = threadIdx.x * patch_w;
  int xx_end = xx_start + patch_w; if (xx_end > output_w) xx_end = output_w;
  int yy_start = threadIdx.y * patch_h;
  int yy_end = yy_start + patch_h; if (yy_end > output_h) yy_end = output_h;
  int oo_start = blockIdx.x;
  int oo_end = oo_start+1;
  int ii_start = 0;

      // default 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+=yy_step) {
            for(xx = xx_start; xx < xx_end; xx+=xx_step) {
              // 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;
