max pool layer is fixed

src/darknet.c

@@ -377,6 +377,7 @@ int main(int argc, char **argv)
 #else
     if(gpu_index >= 0){
         cuda_set_device(gpu_index);
+        check_error(cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync));
     }
 #endif
 

src/maxpool_layer.c

@@ -27,8 +27,8 @@ maxpool_layer make_maxpool_layer(int batch, int h, int w, int c, int size, int s
     l.w = w;
     l.c = c;
     l.pad = padding;
-    l.out_w = (w + 2 * padding - size) / stride + 1;
+    l.out_w = (w + padding - size) / stride + 1;
-    l.out_h = (h + 2 * 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 @@ void resize_maxpool_layer(maxpool_layer *l, int w, int h)
     l->w = w;
     l->inputs = h*w*l->c;
 
-    l->out_w = (w + 2 * l->pad - l->size) / l->stride + 1;
+    l->out_w = (w + l->pad - l->size) / l->stride + 1;
-    l->out_h = (h + 2 * 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 resize_maxpool_layer(maxpool_layer *l, int w, int h)
 void forward_maxpool_layer(const maxpool_layer l, network_state state)
 {
     int b,i,j,k,m,n;
-    int w_offset = -l.pad;
+    int w_offset = -l.pad / l.stride;
-    int h_offset = -l.pad;
+    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 @@ extern "C" {
 
 __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 h = (in_h + pad - size) / stride + 1;
-    int w = (in_w + 2 * 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 @@ __global__ void forward_maxpool_layer_kernel(int n, int in_h, int in_w, int in_c
     id /= c;
     int b = id;
 
-    int w_offset = -pad;
+    int w_offset = -pad / 2;
-    int h_offset = -pad;
+    int h_offset = -pad / 2;
 
     int out_index = j + w*(i + h*(k + c*b));
     float max = -INFINITY;
@@ -49,8 +49,8 @@ __global__ void forward_maxpool_layer_kernel(int n, int in_h, int in_w, int in_c
 
 __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 h = (in_h + pad - size) / stride + 1;
-    int w = (in_w + 2 * 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 @@ __global__ void backward_maxpool_layer_kernel(int n, int in_h, int in_w, int in_
     id /= in_c;
     int b = id;
 
-    int w_offset = -pad;
+    int w_offset = -pad / 2;
-    int h_offset = -pad;
+    int h_offset = -pad / 2;
 
     float d = 0;
     int l, m;

src/parser.c

@@ -457,7 +457,7 @@ maxpool_layer parse_maxpool(list *options, size_params params)
 {
     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;