From 29f64c8c9e081ddfe637a105c68766e93998e219 Mon Sep 17 00:00:00 2001 From: Zeng Jinle <32832641+sneaxiy@users.noreply.github.com> Date: Tue, 10 Dec 2019 05:18:12 -0600 Subject: [PATCH] refine some grad op makers, test=develop (#21629) --- .../fluid/op_use_default_grad_op_maker.spec | 4 -- paddle/fluid/operators/lstm_unit_op.cc | 29 +++++++++++++-- paddle/fluid/operators/lstm_unit_op.cu | 12 +++--- paddle/fluid/operators/lstm_unit_op.h | 3 -- paddle/fluid/operators/pool_with_index_op.cc | 37 ++++++++++++++----- .../operators/tensor_array_to_tensor_op.cc | 21 ++++++++++- 6 files changed, 76 insertions(+), 30 deletions(-) diff --git a/paddle/fluid/op_use_default_grad_op_maker.spec b/paddle/fluid/op_use_default_grad_op_maker.spec index 5defd153687..c6ffd5483ad 100644 --- a/paddle/fluid/op_use_default_grad_op_maker.spec +++ b/paddle/fluid/op_use_default_grad_op_maker.spec @@ -1,10 +1,7 @@ cos_sim fsp gru -lstm_unit match_matrix_tensor -max_pool2d_with_index -max_pool3d_with_index maxout pool2d pool3d @@ -15,6 +12,5 @@ reshape rnn_memory_helper sequence_softmax spp -tensor_array_to_tensor transpose unsqueeze diff --git a/paddle/fluid/operators/lstm_unit_op.cc b/paddle/fluid/operators/lstm_unit_op.cc index e3821f9dc20..4298cbc1458 100644 --- a/paddle/fluid/operators/lstm_unit_op.cc +++ b/paddle/fluid/operators/lstm_unit_op.cc @@ -13,6 +13,7 @@ See the License for the specific language governing permissions and limitations under the License. */ #include "paddle/fluid/operators/lstm_unit_op.h" +#include namespace paddle { namespace operators { @@ -94,14 +95,34 @@ class LstmUnitGradOp : public framework::OperatorWithKernel { } }; +template +class LstmUnitGradOpMaker : public framework::SingleGradOpMaker { + public: + using framework::SingleGradOpMaker::SingleGradOpMaker; + + protected: + std::unique_ptr Apply() const override { + std::unique_ptr op(new T()); + op->SetType("lstm_unit_grad"); + op->SetInput("X", this->Input("X")); + op->SetInput("C_prev", this->Input("C_prev")); + op->SetInput("C", this->Output("C")); + op->SetInput(framework::GradVarName("H"), this->OutputGrad("H")); + op->SetInput(framework::GradVarName("C"), this->OutputGrad("C")); + op->SetOutput(framework::GradVarName("X"), this->InputGrad("X")); + op->SetOutput(framework::GradVarName("C_prev"), this->InputGrad("C_prev")); + op->SetAttrMap(this->Attrs()); + return op; + } +}; + } // namespace operators } // namespace paddle namespace ops = paddle::operators; -REGISTER_OPERATOR( - lstm_unit, ops::LstmUnitOp, ops::LstmUnitOpMaker, - paddle::framework::DefaultGradOpMaker, - paddle::framework::DefaultGradOpMaker); +REGISTER_OPERATOR(lstm_unit, ops::LstmUnitOp, ops::LstmUnitOpMaker, + ops::LstmUnitGradOpMaker, + ops::LstmUnitGradOpMaker); REGISTER_OPERATOR(lstm_unit_grad, ops::LstmUnitGradOp); REGISTER_OP_CPU_KERNEL(lstm_unit, ops::LstmUnitKernel, diff --git a/paddle/fluid/operators/lstm_unit_op.cu b/paddle/fluid/operators/lstm_unit_op.cu index 87451cb1271..7d2279f16d3 100644 --- a/paddle/fluid/operators/lstm_unit_op.cu +++ b/paddle/fluid/operators/lstm_unit_op.cu @@ -62,9 +62,9 @@ __global__ void LSTMUnitKernel(const int nthreads, const int dim, template __global__ void LSTMUnitGradientKernel(const int nthreads, const int dim, const T* C_prev, const T* X, const T* C, - const T* H, const T* C_diff, - const T* H_diff, T* C_prev_diff, - T* X_diff, const T forget_bias) { + const T* C_diff, const T* H_diff, + T* C_prev_diff, T* X_diff, + const T forget_bias) { CUDA_1D_KERNEL_LOOP(index, nthreads) { const int n = index / dim; const int d = index % dim; @@ -146,7 +146,6 @@ class LstmUnitGradOpCUDAKernel : public framework::OpKernel { auto* X = x_tensor->data(); auto* C_prev = c_prev_tensor->data(); auto* C = c_tensor->data(); - auto* H = h_tensor->data(); auto* H_diff = hdiff_tensor->data(); auto* C_diff = cdiff_tensor->data(); @@ -163,9 +162,8 @@ class LstmUnitGradOpCUDAKernel : public framework::OpKernel { int n = N * D; int grid = (n + block - 1) / block; - LSTMUnitGradientKernel<<>>(n, D, C_prev, X, C, H, C_diff, - H_diff, C_prev_diff, X_diff, - forget_bias); + LSTMUnitGradientKernel<<>>( + n, D, C_prev, X, C, C_diff, H_diff, C_prev_diff, X_diff, forget_bias); } }; diff --git a/paddle/fluid/operators/lstm_unit_op.h b/paddle/fluid/operators/lstm_unit_op.h index 4ead9c22934..3fe7bda39b6 100644 --- a/paddle/fluid/operators/lstm_unit_op.h +++ b/paddle/fluid/operators/lstm_unit_op.h @@ -88,7 +88,6 @@ class LstmUnitGradKernel : public framework::OpKernel { auto x_tensor = ctx.Input("X"); auto c_prev_tensor = ctx.Input("C_prev"); auto c_tensor = ctx.Input("C"); - auto h_tensor = ctx.Input("H"); auto hdiff_tensor = ctx.Input(framework::GradVarName("H")); auto cdiff_tensor = ctx.Input(framework::GradVarName("C")); @@ -100,7 +99,6 @@ class LstmUnitGradKernel : public framework::OpKernel { auto* X = x_tensor->data(); auto* C_prev = c_prev_tensor->data(); auto* C = c_tensor->data(); - auto* H = h_tensor->data(); auto* H_diff = hdiff_tensor->data(); auto* C_diff = cdiff_tensor->data(); @@ -138,7 +136,6 @@ class LstmUnitGradKernel : public framework::OpKernel { C_prev += D; X += 4 * D; C += D; - H += D; C_diff += D; H_diff += D; X_diff += 4 * D; diff --git a/paddle/fluid/operators/pool_with_index_op.cc b/paddle/fluid/operators/pool_with_index_op.cc index f8148f946fc..7d8e1c49815 100644 --- a/paddle/fluid/operators/pool_with_index_op.cc +++ b/paddle/fluid/operators/pool_with_index_op.cc @@ -13,6 +13,7 @@ See the License for the specific language governing permissions and limitations under the License. */ #include "paddle/fluid/operators/pool_with_index_op.h" +#include namespace paddle { namespace operators { @@ -283,16 +284,33 @@ Example: } }; +template +class MaxPoolWithIndexGradOpMaker : public framework::SingleGradOpMaker { + public: + using framework::SingleGradOpMaker::SingleGradOpMaker; + + protected: + std::unique_ptr Apply() const override { + std::unique_ptr op(new T()); + op->SetType(this->ForwardOpType() + "_grad"); + op->SetAttrMap(this->Attrs()); + op->SetInput("X", this->Input("X")); + op->SetInput("Mask", this->Output("Mask")); + op->SetInput(framework::GradVarName("Out"), this->OutputGrad("Out")); + op->SetOutput(framework::GradVarName("X"), this->InputGrad("X")); + return op; + } +}; + } // namespace operators } // namespace paddle namespace ops = paddle::operators; -REGISTER_OPERATOR( - max_pool2d_with_index, ops::MaxPoolWithIndexOp, - ops::MaxPool2dWithIndexOpMaker, - paddle::framework::DefaultGradOpMaker, - paddle::framework::DefaultGradOpMaker); +REGISTER_OPERATOR(max_pool2d_with_index, ops::MaxPoolWithIndexOp, + ops::MaxPool2dWithIndexOpMaker, + ops::MaxPoolWithIndexGradOpMaker, + ops::MaxPoolWithIndexGradOpMaker); REGISTER_OPERATOR(max_pool2d_with_index_grad, ops::MaxPoolWithIndexOpGrad); REGISTER_OP_CPU_KERNEL( @@ -307,11 +325,10 @@ REGISTER_OP_CPU_KERNEL( ops::MaxPoolWithIndexGradKernel); -REGISTER_OPERATOR( - max_pool3d_with_index, ops::MaxPoolWithIndexOp, - ops::MaxPool3dWithIndexOpMaker, - paddle::framework::DefaultGradOpMaker, - paddle::framework::DefaultGradOpMaker); +REGISTER_OPERATOR(max_pool3d_with_index, ops::MaxPoolWithIndexOp, + ops::MaxPool3dWithIndexOpMaker, + ops::MaxPoolWithIndexGradOpMaker, + ops::MaxPoolWithIndexGradOpMaker); REGISTER_OPERATOR(max_pool3d_with_index_grad, ops::MaxPoolWithIndexOpGrad); REGISTER_OP_CPU_KERNEL( diff --git a/paddle/fluid/operators/tensor_array_to_tensor_op.cc b/paddle/fluid/operators/tensor_array_to_tensor_op.cc index e6f78f52d68..0c3cffa1a3e 100644 --- a/paddle/fluid/operators/tensor_array_to_tensor_op.cc +++ b/paddle/fluid/operators/tensor_array_to_tensor_op.cc @@ -273,6 +273,23 @@ class LoDTensorArray2TensorGradOp : public framework::OperatorBase { } }; +template +class TensorArrayToTensorGradOpMaker : public framework::SingleGradOpMaker { + public: + using framework::SingleGradOpMaker::SingleGradOpMaker; + + protected: + std::unique_ptr Apply() const override { + std::unique_ptr op(new T()); + op->SetType("tensor_array_to_tensor_grad"); + op->SetAttrMap(this->Attrs()); + op->SetInput("X", this->Input("X")); + op->SetInput(framework::GradVarName("Out"), this->OutputGrad("Out")); + op->SetOutput(framework::GradVarName("X"), this->InputGrad("X")); + return op; + } +}; + } // namespace operators } // namespace paddle USE_OP(concat); @@ -281,8 +298,8 @@ namespace ops = paddle::operators; REGISTER_OPERATOR( tensor_array_to_tensor, ops::LoDTensorArray2TensorOp, ops::LoDTensorArray2TensorOpMaker, ops::LoDTensorArray2TensorOpInferShape, - paddle::framework::DefaultGradOpMaker, - paddle::framework::DefaultGradOpMaker); + ops::TensorArrayToTensorGradOpMaker, + ops::TensorArrayToTensorGradOpMaker); REGISTER_OPERATOR(tensor_array_to_tensor_grad, ops::LoDTensorArray2TensorGradOp, ops::LoDTensorArray2TensorGradInferShape, ops::LoDTensorArray2TensorGradInferVarType); -- GitLab