Bias updates bug fix

Joseph Redmon

2015-01-27 153705226d8ca746478b69eeac9bc854766daa11

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
extern "C" {
#include "im2col.h"
#include "cuda.h"
}
 
__global__ void im2col_pad_kernel(float *im,  int offset,
     int channels,  int height,  int width,
     int ksize,  int stride, float *data_col)
{
    int c,h,w;
    int height_col = 1 + (height-1) / stride;
    int width_col = 1 + (width-1) / stride;
    int channels_col = channels * ksize * ksize;
 
    int pad = ksize/2;
 
    int id = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
    int col_size = height_col*width_col*channels_col;
    if (id >= col_size) return;
 
    int col_index = id;
    w = id % width_col;
    id /= width_col;
    h = id % height_col;
    id /= height_col;
    c = id % channels_col;
    id /= channels_col;
 
    int w_offset = c % ksize;
    int h_offset = (c / ksize) % ksize;
    int im_channel = c / ksize / ksize;
    int im_row = h_offset + h * stride - pad;
    int im_col = w_offset + w * stride - pad;
 
    int im_index = offset + im_col + width*(im_row + height*im_channel);
    float val = (im_row < 0 || im_col < 0 || im_row >= height || im_col >= width) ? 0 : im[im_index];
 
    data_col[col_index] = val;
}
 
__global__ void im2col_nopad_kernel(float *im,  int offset,
        int channels,  int height,  int width,
        int ksize,  int stride, float *data_col)
{
    int c,h,w;
    int height_col = (height - ksize) / stride + 1;
    int width_col = (width - ksize) / stride + 1;
    int channels_col = channels * ksize * ksize;
 
    int id = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
    int col_size = height_col*width_col*channels_col;
    if (id >= col_size) return;
 
    int col_index = id;
    w = id % width_col;
    id /= width_col;
    h = id % height_col;
    id /= height_col;
    c = id % channels_col;
    id /= channels_col;
 
    int w_offset = c % ksize;
    int h_offset = (c / ksize) % ksize;
    int im_channel = c / ksize / ksize;
    int im_row = h_offset + h * stride;
    int im_col = w_offset + w * stride;
 
    int im_index = offset + im_col + width*(im_row + height*im_channel);
    float val = (im_row < 0 || im_col < 0 || im_row >= height || im_col >= width) ? 0 : im[im_index];
 
    data_col[col_index] = val;
}
 
extern "C" void im2col_ongpu(float *im, int offset,
        int channels,  int height,  int width,
        int ksize,  int stride,  int pad, float *data_col)
{
 
    int height_col = (height - ksize) / stride + 1;
    int width_col = (width - ksize) / stride + 1;
    int channels_col = channels * ksize * ksize;
 
    if (pad){
        height_col = 1 + (height-1) / stride;
        width_col = 1 + (width-1) / stride;
    }
 
    size_t n = channels_col*height_col*width_col;
 
    if(pad)im2col_pad_kernel<<<cuda_gridsize(n),BLOCK>>>(im,  offset, channels, height, width, ksize, stride, data_col);
    else im2col_nopad_kernel<<<cuda_gridsize(n),BLOCK>>>(im,  offset, channels, height, width, ksize, stride, data_col);
    check_error(cudaPeekAtLastError());
}