   CHECK_SYNC("hl_matrix_rotate failed");
 }
 
-
-__global__ void keMatrixVol2Col(
-        int num_kernels, real*dataSrc, real* dataDst,
-        int depth, int height, int width,
-        int filterD, int filterH, int filterW,
-        int strideD, int strideH, int strideW,
-        int paddingD, int paddingH, int paddingW,
-        int depth_col, int height_col, int width_col){
-
-  for (int index = blockIdx.x * blockDim.x + threadIdx.x;
-       index < num_kernels;
-       index += blockDim.x * gridDim.x){
-
+__global__ void keMatrixVol2Col(int num_kernels,
+                                real* dataSrc,
+                                real* dataDst,
+                                int depth,
+                                int height,
+                                int width,
+                                int filterD,
+                                int filterH,
+                                int filterW,
+                                int strideD,
+                                int strideH,
+                                int strideW,
+                                int paddingD,
+                                int paddingH,
+                                int paddingW,
+                                int depth_col,
+                                int height_col,
+                                int width_col) {
+  for (int index = blockIdx.x * blockDim.x + threadIdx.x; index < num_kernels;
+       index += blockDim.x * gridDim.x) {
     int w_out = index % width_col;
-    int h_out = (index / width_col ) % height_col;
+    int h_out = (index / width_col) % height_col;
     int d_out = (index / width_col / height_col) % depth_col;
     int channel_in = index / width_col / height_col / depth_col;
     int channel_out = channel_in * filterD * filterH * filterW;

  const int threads = 512;
  const int blocks = DIVUP(num_kernels, threads);

  keMatrixVol2Col<<< blocks, threads >>>(
          num_kernels, dataSrc, dataDst,
                  depth, height, width,
                  filterD, filterH, filterW,
                  strideD, strideH, strideW,
                  paddingD, paddingH, paddingW,
                  depth_col, height_col, width_col);
  CHECK_SYNC("hl_matrix_vol2Col failed");
}

__global__ void keMatrixCol2Vol(
        int num_kernels, real*dataDst, real* dataSrc,
        int depth, int height, int width,
        int filterD, int filterH, int filterW,
        int strideD, int strideH, int strideW,
        int paddingD, int paddingH, int paddingW,
        int depth_col, int height_col, int width_col,
        real alpha, real beta){

  for (int index = blockIdx.x * blockDim.x + threadIdx.x;
       index < num_kernels;
       index += blockDim.x * gridDim.x) {

    real val = 0;
    int w = index % width + paddingW;
    int h = (index / width) % height + paddingH;
    int d = (index / width / height) % depth + paddingD;
    int c = index / (width * height * depth);
    // compute the start and end of the output
    int w_col_start = (w < filterW) ? 0 : (w - filterW) / strideW + 1;
    int w_col_end = min(w / strideW + 1, width_col);

}

void hl_matrix_vol2Col(real* dataSrc,
                       int channels,
                       int depth,
                       int height,
                       int width,
                       int filterD,
                       int filterH,
                       int filterW,
                       int strideD,
                       int strideH,
                       int strideW,
                       int paddingD,
                       int paddingH,
                       int paddingW,
                       real* dataDst) {
  int depth_col = (depth + 2 * paddingD - filterD) / strideD + 1;
  int height_col = (height + 2 * paddingH - filterH) / strideH + 1;
  int width_col = (width + 2 * paddingW - filterW) / strideW + 1;
