 #include <cmath>
 #include <cstdlib>
 #include <cstring>
-
-#include <device_functions.h>
 
 #include "caffeine/common.hpp"
 #include "caffeine/util/im2col.hpp"

    int w_col_end = min(w / stride + 1, width_col);
    int h_col_start = (h < ksize) ? 0 : (h - ksize) / stride + 1;
    int h_col_end = min(h / stride + 1, height_col);
    for (int h_col = h_col_start; h_col < h_col_end; ++h_col) {
      for (int w_col = w_col_start; w_col < w_col_end; ++w_col) {
        // the col location: [c * width * height + h_out, w_out]
        int c_col = c * ksize * ksize + (h - h_col * stride) * ksize + (w - w_col * stride); 
        data_im[index] += data_col[(c_col * height_col + h_col) * width_col + w_col];
      }
    }
  }

    const int height, const int width, const int ksize, const int stride,
    double* data_col);

template <typename Dtype>
__global__ void col2im_gpu_kernel(const int n, const Dtype* data_col,
  const int height, const int width, const int channels, const int ksize,
