diff --git a/paddle/fluid/op_use_default_grad_op_maker.spec b/paddle/fluid/op_use_default_grad_op_maker.spec index 5defd153687e34e75aa8b56bfdf4bf489aecf79b..c6ffd5483ad49e084611e3e8c3aabb6191765047 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 e3821f9dc203c08df6ff2e4cd69dd3664bb42597..4298cbc14582e9b4ca307eb87903f07eb7ad7921 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 87451cb1271cb68e6fc5e2f969e8028fac27d0c3..7d2279f16d35c0d39bbeb59fe4bf4450eb8dd13b 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 4ead9c22934dde6e42f9ede47cc1ddf502948fc4..3fe7bda39b68d6ab29a4131ccd59b1693bbdb9ea 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 f8148f946fcddfea23ab2f58b40216f5ff0f715f..7d8e1c498151cb502905fe8ba9e7cd9388c277cc 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 e6f78f52d682449051d31b82f1e9f50ce2680776..0c3cffa1a3e0a39a7d0ef95a3641a4cd72c453c9 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);