extern "C" { #include "col2im.h" #include "cuda.h" } // src: https://github.com/BVLC/caffe/blob/master/src/caffe/util/im2col.cu // You may also want to read: https://github.com/BVLC/caffe/blob/master/LICENSE __global__ void col2im_gpu_kernel(const int n, const float* data_col, const int height, const int width, const int ksize, const int pad, const int stride, const int height_col, const int width_col, float *data_im) { int index = blockIdx.x*blockDim.x+threadIdx.x; for(; index < n; index += blockDim.x*gridDim.x){ float val = 0; int w = index % width + pad; int h = (index / width) % height + pad; int c = index / (width * height); // compute the start and end of the output int w_col_start = (w < ksize) ? 0 : (w - ksize) / stride + 1; 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); // equivalent implementation int offset = (c * ksize * ksize + h * ksize + w) * height_col * width_col; int coeff_h_col = (1 - stride * ksize * height_col) * width_col; int coeff_w_col = (1 - stride * height_col * width_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) { val += data_col[offset + h_col * coeff_h_col + w_col * coeff_w_col]; } } data_im[index] = val; } } void col2im_ongpu(float *data_col, int channels, int height, int width, int ksize, int stride, int pad, float *data_im){ // We are going to launch channels * height_col * width_col kernels, each // kernel responsible for copying a single-channel grid. pad = pad ? ksize/2 : 0; int height_col = (height + 2 * pad - ksize) / stride + 1; int width_col = (width + 2 * pad - ksize) / stride + 1; int num_kernels = channels * height * width; col2im_gpu_kernel<<<(num_kernels+BLOCK-1)/BLOCK, BLOCK>>>( num_kernels, data_col, height, width, ksize, pad, stride, height_col, width_col, data_im); } /* __global__ void col2im_kernel(float *data_col, int channels, int height, int width, int ksize, int stride, int pad, float *data_im) { int height_col = (height - ksize) / stride + 1; int width_col = (width - ksize) / stride + 1; if (pad){ height_col = 1 + (height-1) / stride; width_col = 1 + (width-1) / stride; pad = ksize/2; } int id = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x; if(id >= channels*height*width) return; int index = id; int w = id%width + pad; id /= width; int h = id%height + pad; id /= height; int c = id%channels; int w_start = (w-ksize+stride)/stride; int w_end = w/stride + 1; int h_start = (h-ksize+stride)/stride; int h_end = h/stride + 1; // int rows = channels * ksize * ksize; // int cols = height_col*width_col; int col_offset = (c*ksize*ksize + h * ksize + w)*height_col*width_col; int h_coeff = (1-stride*ksize*height_col)*width_col; int w_coeff = 1-stride*height_col*width_col; float val = 0; int h_col, w_col; for(h_col = h_start; h_col < h_end; ++h_col){ for(w_col = w_start; w_col < w_end; ++w_col){ int col_index = col_offset +h_col*h_coeff + w_col*w_coeff; float part = (w_col < 0 || h_col < 0 || h_col >= height_col || w_col >= width_col) ? 0 : data_col[col_index]; val += part; } } data_im[index] = val; } extern "C" void col2im_ongpu(float *data_col, int channels, int height, int width, int ksize, int stride, int pad, float *data_im) { size_t n = channels*height*width; col2im_kernel<<>>(data_col, channels, height, width, ksize, stride, pad, data_im); check_error(cudaPeekAtLastError()); } */