From d0b9326a352ed2fbc3ae66fdef40b4533a2f211d Mon Sep 17 00:00:00 2001
From: Joseph Redmon <pjreddie@gmail.com>
Date: Tue, 11 Aug 2015 06:22:27 +0000
Subject: [PATCH] Hacks to get nightmare to not break gridsizing
---
src/maxpool_layer_kernels.cu | 10 +++++-----
1 files changed, 5 insertions(+), 5 deletions(-)
diff --git a/src/maxpool_layer_kernels.cu b/src/maxpool_layer_kernels.cu
index a5c8209..8f69f90 100644
--- a/src/maxpool_layer_kernels.cu
+++ b/src/maxpool_layer_kernels.cu
@@ -77,10 +77,10 @@
d += (valid && indexes[out_index] == index) ? delta[out_index] : 0;
}
}
- prev_delta[index] = d;
+ prev_delta[index] += d;
}
-extern "C" void forward_maxpool_layer_gpu(maxpool_layer layer, float *input)
+extern "C" void forward_maxpool_layer_gpu(maxpool_layer layer, network_state state)
{
int h = (layer.h-1)/layer.stride + 1;
int w = (layer.w-1)/layer.stride + 1;
@@ -88,15 +88,15 @@
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, input, layer.output_gpu, layer.indexes_gpu);
+ forward_maxpool_layer_kernel<<<cuda_gridsize(n), BLOCK>>>(n, layer.h, layer.w, layer.c, layer.stride, layer.size, state.input, layer.output_gpu, layer.indexes_gpu);
check_error(cudaPeekAtLastError());
}
-extern "C" void backward_maxpool_layer_gpu(maxpool_layer layer, float * delta)
+extern "C" void backward_maxpool_layer_gpu(maxpool_layer layer, network_state state)
{
size_t n = layer.h*layer.w*layer.c*layer.batch;
- backward_maxpool_layer_kernel<<<cuda_gridsize(n), BLOCK>>>(n, layer.h, layer.w, layer.c, layer.stride, layer.size, layer.delta_gpu, delta, layer.indexes_gpu);
+ backward_maxpool_layer_kernel<<<cuda_gridsize(n), BLOCK>>>(n, layer.h, layer.w, layer.c, layer.stride, layer.size, layer.delta_gpu, state.delta, layer.indexes_gpu);
check_error(cudaPeekAtLastError());
}
--
Gitblit v1.10.0