AlexeyAB
2018-08-04 043289426b2d08d925fc1c980b0d2a01e2360e93
max pool layer is fixed
4 files modified
31 ■■■■ changed files
src/darknet.c 1 ●●●● patch | view | raw | blame | history
src/maxpool_layer.c 12 ●●●● patch | view | raw | blame | history
src/maxpool_layer_kernels.cu 16 ●●●● patch | view | raw | blame | history
src/parser.c 2 ●●● patch | view | raw | blame | history
src/darknet.c
@@ -377,6 +377,7 @@
#else
    if(gpu_index >= 0){
        cuda_set_device(gpu_index);
        check_error(cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync));
    }
#endif
src/maxpool_layer.c
@@ -27,8 +27,8 @@
    l.w = w;
    l.c = c;
    l.pad = padding;
    l.out_w = (w + 2 * padding - size) / stride + 1;
    l.out_h = (h + 2 * padding - size) / stride + 1;
    l.out_w = (w + padding - size) / stride + 1;
    l.out_h = (h + padding - size) / stride + 1;
    l.out_c = c;
    l.outputs = l.out_h * l.out_w * l.out_c;
    l.inputs = h*w*c;
@@ -58,8 +58,8 @@
    l->w = w;
    l->inputs = h*w*l->c;
    l->out_w = (w + 2 * l->pad - l->size) / l->stride + 1;
    l->out_h = (h + 2 * l->pad - l->size) / l->stride + 1;
    l->out_w = (w + l->pad - l->size) / l->stride + 1;
    l->out_h = (h + l->pad - l->size) / l->stride + 1;
    l->outputs = l->out_w * l->out_h * l->c;
    int output_size = l->outputs * l->batch;
@@ -80,8 +80,8 @@
void forward_maxpool_layer(const maxpool_layer l, network_state state)
{
    int b,i,j,k,m,n;
    int w_offset = -l.pad;
    int h_offset = -l.pad;
    int w_offset = -l.pad / l.stride;
    int h_offset = -l.pad / l.stride;
    int h = l.out_h;
    int w = l.out_w;
src/maxpool_layer_kernels.cu
@@ -9,8 +9,8 @@
__global__ void forward_maxpool_layer_kernel(int n, int in_h, int in_w, int in_c, int stride, int size, int pad, float *input, float *output, int *indexes)
{
    int h = (in_h + 2 * pad - size) / stride + 1;
    int w = (in_w + 2 * pad - size) / stride + 1;
    int h = (in_h + pad - size) / stride + 1;
    int w = (in_w + pad - size) / stride + 1;
    int c = in_c;
    int id = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
@@ -24,8 +24,8 @@
    id /= c;
    int b = id;
    int w_offset = -pad;
    int h_offset = -pad;
    int w_offset = -pad / 2;
    int h_offset = -pad / 2;
    int out_index = j + w*(i + h*(k + c*b));
    float max = -INFINITY;
@@ -49,8 +49,8 @@
__global__ void backward_maxpool_layer_kernel(int n, int in_h, int in_w, int in_c, int stride, int size, int pad, float *delta, float *prev_delta, int *indexes)
{
    int h = (in_h + 2 * pad - size) / stride + 1;
    int w = (in_w + 2 * pad - size) / stride + 1;
    int h = (in_h + pad - size) / stride + 1;
    int w = (in_w + pad - size) / stride + 1;
    int c = in_c;
    int area = (size-1)/stride;
@@ -66,8 +66,8 @@
    id /= in_c;
    int b = id;
    int w_offset = -pad;
    int h_offset = -pad;
    int w_offset = -pad / 2;
    int h_offset = -pad / 2;
    float d = 0;
    int l, m;
src/parser.c
@@ -457,7 +457,7 @@
{
    int stride = option_find_int(options, "stride",1);
    int size = option_find_int(options, "size",stride);
    int padding = option_find_int_quiet(options, "padding", (size-1)/2);
    int padding = option_find_int_quiet(options, "padding", size-1);
    int batch,h,w,c;
    h = params.h;