     long inputHeight  = input->size[dimh];
     long outputWidth  = (inputWidth - kW) / dW + 1;
     long outputHeight = (inputHeight - kH) / dH + 1;
-    
+
     luaL_argcheck(L, kW == kH, 1, "filters must be square (kW == kH)");
     luaL_argcheck(L, dW == dH, 1, "stride must be square (dW == dH)");
-    
+
     if (input->nDimension == 3) {
         // Resize output
         THCudaTensor_resize3d(output, nOutputPlane, outputHeight, outputWidth);
-        
+
         // Resize temporary columns
         THCudaTensor_resize2d(columns, outputHeight*outputWidth, nOutputPlane*kW*kH);
 

        // Resize temporary columns
        THCudaTensor_resize2d(columns, nInputPlane*kW*kH, outputHeight*outputWidth);

	/* add bias */
	{
	  /* 
	     batchSize/4 blocks
	     32 warps per block, 
	     4 batches per block, 
	     8 warps per batch-slice 
	     Each warp handles 1 batch's nOutputPlane/8 
	  */
	  dim3 blocks(batchSize/4); /* 128/4 = 32 */
	  dim3 threads(1024); 
	  fillBiasBatch <<<blocks,threads>>> (THCudaTensor_data(output), THCudaTensor_data(bias),
					      batchSize, nOutputPlane, outputHeight, outputWidth);
	}

        // Helper    
        THCudaTensor *output_n = THCudaTensor_new();

    } else {
        // Batch size + input planes
        long batchSize = input->size[0];
        luaL_argcheck(L, batchSize % 4 == 0, 1, "batch size should be a multiple of 4");
        luaL_argcheck(L, nOutputPlane % 8 == 0, 1, "nOutputPlane should be a multiple of 8");

        // Resize output
        THCudaTensor_resize4d(output, batchSize, nOutputPlane, outputHeight, outputWidth);
