From a723e1c62a27aeb39aaf7fcdeb3beb4e89fba32d Mon Sep 17 00:00:00 2001
From: Alexey <AlexeyAB@users.noreply.github.com>
Date: Wed, 15 Aug 2018 20:52:09 +0000
Subject: [PATCH] Merge pull request #766 from HotChick91/AlexeyAB-mask
---
src/maxpool_layer_kernels.cu | 18 +++++++++---------
1 files changed, 9 insertions(+), 9 deletions(-)
diff --git a/src/maxpool_layer_kernels.cu b/src/maxpool_layer_kernels.cu
index fc54f52..78b7f39 100644
--- a/src/maxpool_layer_kernels.cu
+++ b/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 + 1)/stride + 1;
- int w = (in_w + 2*pad - size + 1)/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 + 1)/stride + 1;
- int w = (in_w + 2*pad - size + 1)/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;
@@ -92,7 +92,7 @@
size_t n = h*w*c*layer.batch;
- forward_maxpool_layer_kernel<<<cuda_gridsize(n), BLOCK>>>(n, layer.h, layer.w, layer.c, layer.stride, layer.size, layer.pad, state.input, layer.output_gpu, layer.indexes_gpu);
+ forward_maxpool_layer_kernel<<<cuda_gridsize(n), BLOCK, 0, get_cuda_stream()>>>(n, layer.h, layer.w, layer.c, layer.stride, layer.size, layer.pad, state.input, layer.output_gpu, layer.indexes_gpu);
check_error(cudaPeekAtLastError());
}
--
Gitblit v1.10.0