From f6e69d74124995fd0dac30b6d7d3717ef2a9241a Mon Sep 17 00:00:00 2001 From: chengduoZH Date: Mon, 25 Sep 2017 23:10:03 +0800 Subject: [PATCH] fix maxpool backward functor --- paddle/operators/math/pool_test_maxPool2d.cc | 22 ++++++++++++-------- paddle/operators/math/pool_test_maxPool3d.cc | 21 +++++++++++-------- paddle/operators/math/pooling.cc | 4 ++-- paddle/operators/pool_op.h | 12 ++++------- 4 files changed, 31 insertions(+), 28 deletions(-) diff --git a/paddle/operators/math/pool_test_maxPool2d.cc b/paddle/operators/math/pool_test_maxPool2d.cc index 8ddf1d79c..63b6b5657 100644 --- a/paddle/operators/math/pool_test_maxPool2d.cc +++ b/paddle/operators/math/pool_test_maxPool2d.cc @@ -23,16 +23,17 @@ #ifndef PADDLE_ONLY_CPU -template -void testPool2d(paddle::platform::DeviceContext& context, PooType pool_process, - paddle::framework::Tensor& input, +template +void testPool2d(paddle::platform::DeviceContext& context, PoolType pool_process, + PoolGradType poolGrad_process, paddle::framework::Tensor& input, paddle::framework::Tensor& input_grad, paddle::framework::Tensor& output, paddle::framework::Tensor& output_grad, std::vector& ksize, std::vector& strides, std::vector& paddings) { paddle::operators::math::Pool2dForwardFunctor + PoolType, float> pool2d_forward; + pool2d_forward(context, input, output, ksize, strides, paddings, pool_process); @@ -44,10 +45,11 @@ void testPool2d(paddle::platform::DeviceContext& context, PooType pool_process, start = clock(); for (int i = 0; i < times; ++i) { paddle::operators::math::Pool2dBackwardFunctor + PoolGradType, float> pool2d_backward; pool2d_backward(context, input, input_grad, output, output_grad, ksize, - strides, paddings, pool_process); + strides, paddings, poolGrad_process); + PADDLE_ENFORCE(cudaStreamSynchronize(0), "cudaStreamSynchronize failed in pool2d_backward CopyFrom"); } @@ -136,10 +138,12 @@ void test2dPool() { paddle::platform::DeviceContext* context = new paddle::platform::CUDADeviceContext(paddle::platform::GPUPlace()); paddle::operators::math::pool::maxPool pool_process; + paddle::operators::math::pool::maxPoolGrad poolGrad_process; - testPool2d>( - *context, pool_process, input, input_grad, output, output_grad, ksize, - strides, paddings); + testPool2d, + paddle::operators::math::pool::maxPoolGrad>( + *context, pool_process, poolGrad_process, input, input_grad, output, + output_grad, ksize, strides, paddings); } int main() { diff --git a/paddle/operators/math/pool_test_maxPool3d.cc b/paddle/operators/math/pool_test_maxPool3d.cc index 000b006a1..ce572164c 100644 --- a/paddle/operators/math/pool_test_maxPool3d.cc +++ b/paddle/operators/math/pool_test_maxPool3d.cc @@ -23,15 +23,15 @@ #ifndef PADDLE_ONLY_CPU -template -void testPool3d(paddle::platform::DeviceContext& context, PooType pool_process, - paddle::framework::Tensor& input, +template +void testPool3d(paddle::platform::DeviceContext& context, PoolType pool_process, + PoolGradType poolGrad_process, paddle::framework::Tensor& input, paddle::framework::Tensor& input_grad, paddle::framework::Tensor& output, paddle::framework::Tensor& output_grad, std::vector& ksize, std::vector& strides, std::vector& paddings) { paddle::operators::math::Pool3dForwardFunctor + PoolType, float> pool3d_forward; pool3d_forward(context, input, output, ksize, strides, paddings, pool_process); @@ -44,10 +44,10 @@ void testPool3d(paddle::platform::DeviceContext& context, PooType pool_process, start = clock(); for (int i = 0; i < times; ++i) { paddle::operators::math::Pool3dBackwardFunctor + PoolGradType, float> pool3d_backward; pool3d_backward(context, input, input_grad, output, output_grad, ksize, - strides, paddings, pool_process); + strides, paddings, poolGrad_process); PADDLE_ENFORCE(cudaStreamSynchronize(0), "cudaStreamSynchronize failed in pool3d_backward CopyFrom"); } @@ -145,9 +145,12 @@ void test3dPool() { new paddle::platform::CUDADeviceContext(paddle::platform::GPUPlace()); paddle::operators::math::pool::maxPool pool_process; - testPool3d>( - *context, pool_process, input, input_grad, output, output_grad, ksize, - strides, paddings); + paddle::operators::math::pool::maxPoolGrad poolGrad_process; + + testPool3d, + paddle::operators::math::pool::maxPoolGrad>( + *context, pool_process, poolGrad_process, input, input_grad, output, + output_grad, ksize, strides, paddings); } int main() { test3dPool(); } diff --git a/paddle/operators/math/pooling.cc b/paddle/operators/math/pooling.cc index 6539b8067..cb1f9d728 100644 --- a/paddle/operators/math/pooling.cc +++ b/paddle/operators/math/pooling.cc @@ -196,7 +196,7 @@ class MaxPool2dBackwardFunctor { }; template class MaxPool2dBackwardFunctor; -template class MaxPool2dBackwardFunctor; +// template class MaxPool2dBackwardFunctor; template class Pool2dForwardFunctor< platform::CPUPlace, paddle::operators::math::pool::maxPool, float>; @@ -443,7 +443,7 @@ class MaxPool3dBackwardFunctor { }; template class MaxPool3dBackwardFunctor; -template class MaxPool3dBackwardFunctor; +// template class MaxPool3dBackwardFunctor; template class Pool3dForwardFunctor< platform::CPUPlace, paddle::operators::math::pool::maxPool, float>; diff --git a/paddle/operators/pool_op.h b/paddle/operators/pool_op.h index 6f1d05272..763103d88 100644 --- a/paddle/operators/pool_op.h +++ b/paddle/operators/pool_op.h @@ -111,12 +111,10 @@ class PoolGradKernel : public framework::OpKernel { switch (ksize.size()) { case 2: { if (pooling_type == "max") { - paddle::operators::math::Pool2dBackwardFunctor< - Place, paddle::operators::math::pool::maxPoolGrad, T> + paddle::operators::math::MaxPool2dBackwardFunctor pool2d_backward; - paddle::operators::math::pool::maxPoolGrad pool_process; pool2d_backward(context.device_context(), *in_X, *in_X_grad, *out, - *out_grad, ksize, strides, paddings, pool_process); + *out_grad, ksize, strides, paddings); } else if (pooling_type == "avg") { paddle::operators::math::Pool2dBackwardFunctor< Place, paddle::operators::math::pool::avgPoolGrad, T> @@ -128,12 +126,10 @@ class PoolGradKernel : public framework::OpKernel { } break; case 3: { if (pooling_type == "max") { - paddle::operators::math::Pool3dBackwardFunctor< - Place, paddle::operators::math::pool::maxPoolGrad, T> + paddle::operators::math::MaxPool3dBackwardFunctor pool3d_backward; - paddle::operators::math::pool::maxPoolGrad pool_process; pool3d_backward(context.device_context(), *in_X, *in_X_grad, *out, - *out_grad, ksize, strides, paddings, pool_process); + *out_grad, ksize, strides, paddings); } else if (pooling_type == "avg") { paddle::operators::math::Pool3dBackwardFunctor< Place, paddle::operators::math::pool::avgPoolGrad, T> -- GitLab