From 9d418102f4a44b2e2c437dec945189a646cbf3a4 Mon Sep 17 00:00:00 2001
From: Joseph Redmon <pjreddie@gmail.com>
Date: Sat, 21 Mar 2015 21:17:39 +0000
Subject: [PATCH] using caffe's im2col, it's so much better\!
---
src/col2im_kernels.cu | 144 ++++++++++++++++++++++++++++++++---------------
1 files changed, 98 insertions(+), 46 deletions(-)
diff --git a/src/col2im_kernels.cu b/src/col2im_kernels.cu
index 2fa2030..76a86e6 100644
--- a/src/col2im_kernels.cu
+++ b/src/col2im_kernels.cu
@@ -3,60 +3,112 @@
#include "cuda.h"
}
-__global__ void col2im_kernel(float *data_col,
- int channels, int height, int width,
- int ksize, int stride, int pad, float *data_im)
-{
+// 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
- 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;
+__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;
}
- data_im[index] = val;
+}
+
+void col2im_ongpu(float *im,
+ int channels, int height, int width,
+ int ksize, int stride, int pad, float *data_col){
+ // 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, 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)
+int channels, int height, int width,
+int ksize, int stride, int pad, float *data_im)
{
- size_t n = channels*height*width;
+size_t n = channels*height*width;
- col2im_kernel<<<cuda_gridsize(n), BLOCK>>>(data_col, channels, height, width, ksize, stride, pad, data_im);
- check_error(cudaPeekAtLastError());
+col2im_kernel<<<cuda_gridsize(n), BLOCK>>>(data_col, channels, height, width, ksize, stride, pad, data_im);
+check_error(cudaPeekAtLastError());
}
+ */
--
Gitblit v1.10.0