From 028696bf15efeca3acb3db8c42a96f7b9e0f55ff Mon Sep 17 00:00:00 2001
From: iovodov <b@ovdv.ru>
Date: Thu, 03 May 2018 13:33:46 +0000
Subject: [PATCH] Output improvements for detector results: When printing detector results, output was done in random order, obfuscating results for interpreting. Now: 1. Text output includes coordinates of rects in (left,right,top,bottom in pixels) along with label and score 2. Text output is sorted by rect lefts to simplify finding appropriate rects on image 3. If several class probs are > thresh for some detection, the most probable is written first and coordinates for others are not repeated 4. Rects are imprinted in image in order by their best class prob, so most probable rects are always on top and not overlayed by less probable ones 5. Most probable label for rect is always written first Also: 6. Message about low GPU memory include required amount
---
src/col2im_kernels.cu | 98 +++++++++++++++++++++++-------------------------
1 files changed, 47 insertions(+), 51 deletions(-)
diff --git a/src/col2im_kernels.cu b/src/col2im_kernels.cu
index 2fa2030..aed2df9 100644
--- a/src/col2im_kernels.cu
+++ b/src/col2im_kernels.cu
@@ -1,62 +1,58 @@
+#include "cuda_runtime.h"
+#include "curand.h"
+#include "cublas_v2.h"
+
extern "C" {
#include "col2im.h"
#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;
}
-
-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<<<cuda_gridsize(n), BLOCK>>>(data_col, channels, height, width, ksize, stride, pad, data_im);
- check_error(cudaPeekAtLastError());
+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.
+ 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);
}
+
--
Gitblit v1.10.0