diff --git a/paddle/operators/math/pool_test_maxPool2d.cc b/paddle/operators/math/pool_test_maxPool2d.cc index 8ddf1d79c7a9213064dd4049fd2ded68dece983e..63b6b5657bea457d3d5aa2a4727842d8f9632401 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 000b006a1281d99d3427740bc7b0f9ed3da5d1dd..ce572164c4f800c712e04fd2bb359ac793c042dd 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 6539b80678431d3edf5d9cdf8da5a1e27a3b85e0..cb1f9d72857592bd8cb6d4532170d03a11e45343 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 6f1d05272cd1e66f8858d19454fe7dc08ce7b45d..763103d884bb223a0ec6767299a22ca3259c0cc8 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>