From d34aadb3f6c92ef0b7909cfaa7e006b1e9ff7561 Mon Sep 17 00:00:00 2001 From: "Yang Yang(Tony)" Date: Mon, 9 Oct 2017 14:17:13 -0700 Subject: [PATCH 01/36] Create executor.md --- doc/design/executor.md | 95 ++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 95 insertions(+) create mode 100644 doc/design/executor.md diff --git a/doc/design/executor.md b/doc/design/executor.md new file mode 100644 index 0000000000..bf7f05552c --- /dev/null +++ b/doc/design/executor.md @@ -0,0 +1,95 @@ +# Executor Desgin Doc + +## Overview + +`Executor` evaluates a `ProgramDesc`. Essentially, it instantializes Variables and Operators, then run all the operators + +```c++ +void Executor::Run(const ProgramDesc& pdesc, Scope* scope) { + auto& block = pdesc.blocks(0); + auto& device = device_contexts_[0]; + + // Instantiate all the vars in the global scope + for (auto& var : block.vars()) { + scope->NewVar(var.name()); + } + + // Decide which operator should be run + std::vector should_run = Preprocess(pdesc); + + // Run the block + Scope& local_scope = scope->NewScope(); + for (size_t i = 0; i < should_run.size(); ++i) { + if (should_run[i]) { + for (auto var : block.ops(i).outputs()) { + for (auto argu : var.arguments()) { + // Create variable in the local_scope + if (local_scope.FindVar(argu) == nullptr) { + local_scope.NewVar(argu); + } + } + } + auto op = paddle::framework::OpRegistry::CreateOp(block.ops(i)); + op->Run(local_scope, *device); + } + } +} +``` + +## Tasks + +As shown above, it is not hard to simply evaluate the graph. The real problem +is how do we actually construct the `ProgramDesc`. There are several different +situations that we need to consider. + +### 1. Init @tony @qijun + +##### Problem: + +Not sure which block to put init ops. Same concerns applys to `Load Model`. + +##### Solution: In seperate Blocks + +All `initop` and `parameter` goes to `block[0]`. Actual run starts from `block[1]`. + +When user writes `a = Parameter(Variable, init)`, a init op is inserted into +`block[0]`, and a `NOP` is inserted into `block[1]` to substitute init op. + +- Pro: + - Init Op can be run multiple times. + - Compatiable with current `Executor::Preprocessing` + - Still only one `ProgramDesc` +- Con: + - Let others know! + +### 2. IO + +#### 2.1 FeedOp and FetchOp + +Design Doc: https://github.com/PaddlePaddle/Paddle/pull/4599 + +FeedOp and FetchOp in distributed environment: +https://github.com/PaddlePaddle/Paddle/issues/4613 + +#### 2.2 ReaderOp and WriterOp + +### 3. Backward @jiayi + +Executor test case is a good place to test `backward` module, even though executor +is not necessarily depends on `backward`. Currently exposed issue: + +- Fill One: https://github.com/PaddlePaddle/Paddle/issues/4627 +- Attribute map: https://github.com/PaddlePaddle/Paddle/issues/4642 + +### 4. Optimizer @longfei + +Executor test case is a good place to test `optimizer `module, even though executor +is not necessarily depends on `optimizer `. + +### 5. RNN @chunwei + +To be discussed. + +- How to deal with multiple blocks +- How to deal with LoDTensor + From 3b879598c49b9caa0a2e1b2ad8bec103a15a63ef Mon Sep 17 00:00:00 2001 From: Yang Yang Date: Wed, 11 Oct 2017 01:23:34 +0000 Subject: [PATCH 02/36] update executor design doc --- doc/design/executor.md | 77 ++++++++++++++---------------------------- 1 file changed, 25 insertions(+), 52 deletions(-) diff --git a/doc/design/executor.md b/doc/design/executor.md index bf7f05552c..7376ecaef0 100644 --- a/doc/design/executor.md +++ b/doc/design/executor.md @@ -15,7 +15,7 @@ void Executor::Run(const ProgramDesc& pdesc, Scope* scope) { } // Decide which operator should be run - std::vector should_run = Preprocess(pdesc); + std::vector should_run = Prune(pdesc); // Run the block Scope& local_scope = scope->NewScope(); @@ -23,7 +23,7 @@ void Executor::Run(const ProgramDesc& pdesc, Scope* scope) { if (should_run[i]) { for (auto var : block.ops(i).outputs()) { for (auto argu : var.arguments()) { - // Create variable in the local_scope + // Create temp variable in the local_scope if (local_scope.FindVar(argu) == nullptr) { local_scope.NewVar(argu); } @@ -36,60 +36,33 @@ void Executor::Run(const ProgramDesc& pdesc, Scope* scope) { } ``` -## Tasks +## Challenge -As shown above, it is not hard to simply evaluate the graph. The real problem -is how do we actually construct the `ProgramDesc`. There are several different -situations that we need to consider. +It is not hard to simply evaluate a graph. However, it is hard to determine which op should be run. Consider the following different situations. -### 1. Init @tony @qijun - -##### Problem: - -Not sure which block to put init ops. Same concerns applys to `Load Model`. - -##### Solution: In seperate Blocks - -All `initop` and `parameter` goes to `block[0]`. Actual run starts from `block[1]`. - -When user writes `a = Parameter(Variable, init)`, a init op is inserted into -`block[0]`, and a `NOP` is inserted into `block[1]` to substitute init op. - -- Pro: - - Init Op can be run multiple times. - - Compatiable with current `Executor::Preprocessing` - - Still only one `ProgramDesc` -- Con: - - Let others know! - -### 2. IO - -#### 2.1 FeedOp and FetchOp - -Design Doc: https://github.com/PaddlePaddle/Paddle/pull/4599 - -FeedOp and FetchOp in distributed environment: -https://github.com/PaddlePaddle/Paddle/issues/4613 - -#### 2.2 ReaderOp and WriterOp - -### 3. Backward @jiayi - -Executor test case is a good place to test `backward` module, even though executor -is not necessarily depends on `backward`. Currently exposed issue: - -- Fill One: https://github.com/PaddlePaddle/Paddle/issues/4627 -- Attribute map: https://github.com/PaddlePaddle/Paddle/issues/4642 - -### 4. Optimizer @longfei +```python +# Case 1: run foward pass. +cost_np = executor.run(target=cost) +# Case 2: run backward passing. +opts_np, _ = executor.run(target=[cost, opt]) +# Case 3: run checkpointing +_ = executor.run(target=checkpoint) +``` -Executor test case is a good place to test `optimizer `module, even though executor -is not necessarily depends on `optimizer `. +We want to support the evaluation of both variables and operators. -### 5. RNN @chunwei +## Solution -To be discussed. +To support evaluation of operators, we add `is_target` field in the `OpDesc`. -- How to deal with multiple blocks -- How to deal with LoDTensor +```c++ +message OpDesc { + required string type = 3; + repeated Var inputs = 1; + repeated Var outputs = 2; + repeated Attr attrs = 4; + required bool is_target = 5 [ default = false ]; // true if the op is target +}; +``` +To support evaluation of variables, we add [fetch_op](https://github.com/PaddlePaddle/Paddle/pull/4599). For each variable in the `target`, we insert a `fetch_op` into the `ProgramDesc`. (Also, a user may want to overwrite a variable, so we also added [feed_op](https://github.com/PaddlePaddle/Paddle/pull/4599). ) From a6fbfed2c18b72ad08cc91180fbd8e090f223a61 Mon Sep 17 00:00:00 2001 From: "Yang Yang(Tony)" Date: Wed, 11 Oct 2017 10:37:28 -0700 Subject: [PATCH 03/36] Update executor.md --- doc/design/executor.md | 54 +++++++----------------------------------- 1 file changed, 9 insertions(+), 45 deletions(-) diff --git a/doc/design/executor.md b/doc/design/executor.md index 7376ecaef0..049ddb6a59 100644 --- a/doc/design/executor.md +++ b/doc/design/executor.md @@ -5,8 +5,8 @@ `Executor` evaluates a `ProgramDesc`. Essentially, it instantializes Variables and Operators, then run all the operators ```c++ -void Executor::Run(const ProgramDesc& pdesc, Scope* scope) { - auto& block = pdesc.blocks(0); +void Executor::Run(const ProgramDesc& pdesc, Scope* scope, int block_id) { + auto& block = pdesc.blocks(block_id); auto& device = device_contexts_[0]; // Instantiate all the vars in the global scope @@ -14,55 +14,19 @@ void Executor::Run(const ProgramDesc& pdesc, Scope* scope) { scope->NewVar(var.name()); } - // Decide which operator should be run - std::vector should_run = Prune(pdesc); - // Run the block Scope& local_scope = scope->NewScope(); for (size_t i = 0; i < should_run.size(); ++i) { - if (should_run[i]) { - for (auto var : block.ops(i).outputs()) { - for (auto argu : var.arguments()) { - // Create temp variable in the local_scope - if (local_scope.FindVar(argu) == nullptr) { - local_scope.NewVar(argu); - } + for (auto var : block.ops(i).outputs()) { + for (auto argu : var.arguments()) { + // Create temp variable in the local_scope + if (local_scope.FindVar(argu) == nullptr) { + local_scope.NewVar(argu); } } - auto op = paddle::framework::OpRegistry::CreateOp(block.ops(i)); - op->Run(local_scope, *device); } + auto op = paddle::framework::OpRegistry::CreateOp(block.ops(i)); + op->Run(local_scope, *device); } } ``` - -## Challenge - -It is not hard to simply evaluate a graph. However, it is hard to determine which op should be run. Consider the following different situations. - -```python -# Case 1: run foward pass. -cost_np = executor.run(target=cost) -# Case 2: run backward passing. -opts_np, _ = executor.run(target=[cost, opt]) -# Case 3: run checkpointing -_ = executor.run(target=checkpoint) -``` - -We want to support the evaluation of both variables and operators. - -## Solution - -To support evaluation of operators, we add `is_target` field in the `OpDesc`. - -```c++ -message OpDesc { - required string type = 3; - repeated Var inputs = 1; - repeated Var outputs = 2; - repeated Attr attrs = 4; - required bool is_target = 5 [ default = false ]; // true if the op is target -}; -``` - -To support evaluation of variables, we add [fetch_op](https://github.com/PaddlePaddle/Paddle/pull/4599). For each variable in the `target`, we insert a `fetch_op` into the `ProgramDesc`. (Also, a user may want to overwrite a variable, so we also added [feed_op](https://github.com/PaddlePaddle/Paddle/pull/4599). ) From 11c6dc6798d5e12f3ea7e0d4eea996e47ca8bb7d Mon Sep 17 00:00:00 2001 From: "Yang Yang(Tony)" Date: Wed, 11 Oct 2017 11:00:32 -0700 Subject: [PATCH 04/36] Update executor.md --- doc/design/executor.md | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/doc/design/executor.md b/doc/design/executor.md index 049ddb6a59..30b76b3385 100644 --- a/doc/design/executor.md +++ b/doc/design/executor.md @@ -16,16 +16,16 @@ void Executor::Run(const ProgramDesc& pdesc, Scope* scope, int block_id) { // Run the block Scope& local_scope = scope->NewScope(); - for (size_t i = 0; i < should_run.size(); ++i) { - for (auto var : block.ops(i).outputs()) { - for (auto argu : var.arguments()) { + for (auto& op_desc : block.ops()) { + for (auto& var : op_desc) { + for (auto& argu : var.arguments()) { // Create temp variable in the local_scope if (local_scope.FindVar(argu) == nullptr) { local_scope.NewVar(argu); } } } - auto op = paddle::framework::OpRegistry::CreateOp(block.ops(i)); + auto op = paddle::framework::OpRegistry::CreateOp(op_desc); op->Run(local_scope, *device); } } From 9d3d82f9799b64a410da66c17c0fbd3cff6a25da Mon Sep 17 00:00:00 2001 From: "Yang Yang(Tony)" Date: Wed, 11 Oct 2017 11:04:29 -0700 Subject: [PATCH 05/36] Update executor.md --- doc/design/executor.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/doc/design/executor.md b/doc/design/executor.md index 30b76b3385..777f451b11 100644 --- a/doc/design/executor.md +++ b/doc/design/executor.md @@ -17,7 +17,7 @@ void Executor::Run(const ProgramDesc& pdesc, Scope* scope, int block_id) { // Run the block Scope& local_scope = scope->NewScope(); for (auto& op_desc : block.ops()) { - for (auto& var : op_desc) { + for (auto& var : op_desc.outputs()) { for (auto& argu : var.arguments()) { // Create temp variable in the local_scope if (local_scope.FindVar(argu) == nullptr) { From 1c9e461520a3e7dc9c8753b29a2c39dda9ce1600 Mon Sep 17 00:00:00 2001 From: "Yang Yang(Tony)" Date: Wed, 11 Oct 2017 11:48:39 -0700 Subject: [PATCH 06/36] add detailed overview --- doc/design/executor.md | 22 ++++++++++++++++++++-- 1 file changed, 20 insertions(+), 2 deletions(-) diff --git a/doc/design/executor.md b/doc/design/executor.md index 777f451b11..ab2d6c3558 100644 --- a/doc/design/executor.md +++ b/doc/design/executor.md @@ -1,8 +1,26 @@ -# Executor Desgin Doc +# Executor Design Doc + +## Motivation + +We use executor to do the runtime evaluation of a `ProgramDesc`. ## Overview -`Executor` evaluates a `ProgramDesc`. Essentially, it instantializes Variables and Operators, then run all the operators +An executor takes a `ProgramDesc`, a `block_id` and a `Scope`. The `ProgramDesc` is a list of blocks and each block contains the protobuf definition of all the parameters and operators. The `block_id` specifies the entrance block. And the `Scope` is the container of all the variable instance, which is persistent throughout different runs. + +### What does executor do? + +It evaluates all the operators in the `block_id`th block of a `ProgramDesc`. + +### What does executor NOT do? + +It does not do runtime optimization, meaning intelligently parse the dependency of each op a choose which one to be run and in which order they should be run. + +It does not do graph partitioning, meaning dividing the `ProgramDesc` into several small pieces and executing them on different devices. + +## Implementation + +`Executor` evaluates a `ProgramDesc`. Essentially, it instantiates Variables and Operators, then run all the operators ```c++ void Executor::Run(const ProgramDesc& pdesc, Scope* scope, int block_id) { From 1b1cb44f13242f2e315b6f648679cf936eb999a2 Mon Sep 17 00:00:00 2001 From: Yu Yang Date: Fri, 13 Oct 2017 14:05:28 -0700 Subject: [PATCH 07/36] Complete infer_var_type --- paddle/framework/CMakeLists.txt | 3 + paddle/framework/details/op_registry.h | 19 +++- paddle/framework/op_desc.cc | 14 +++ paddle/framework/op_desc.h | 2 + paddle/framework/op_info.h | 2 +- paddle/framework/type_defs.h | 9 ++ paddle/framework/var_type_inference.h | 29 ++++++ paddle/framework/var_type_inference_test.cc | 103 ++++++++++++++++++++ 8 files changed, 178 insertions(+), 3 deletions(-) create mode 100644 paddle/framework/var_type_inference.h create mode 100644 paddle/framework/var_type_inference_test.cc diff --git a/paddle/framework/CMakeLists.txt b/paddle/framework/CMakeLists.txt index 14947b6f21..2c61ae40a5 100644 --- a/paddle/framework/CMakeLists.txt +++ b/paddle/framework/CMakeLists.txt @@ -53,3 +53,6 @@ endif() cc_library(tensor_array SRCS tensor_array.cc DEPS lod_tensor) cc_test(tensor_array_test SRCS tensor_array_test.cc DEPS tensor_array place) + +cc_test(var_type_inference_test SRCS var_type_inference_test.cc DEPS op_registry + proto_desc) diff --git a/paddle/framework/details/op_registry.h b/paddle/framework/details/op_registry.h index ca8584b78a..71ee143018 100644 --- a/paddle/framework/details/op_registry.h +++ b/paddle/framework/details/op_registry.h @@ -18,6 +18,7 @@ #include "paddle/framework/op_info.h" #include "paddle/framework/op_proto_maker.h" #include "paddle/framework/operator.h" +#include "paddle/framework/var_type_inference.h" namespace paddle { namespace framework { @@ -26,7 +27,8 @@ namespace details { enum OpInfoFillType { kOperator = 0, kOpProtoAndCheckerMaker = 1, - kGradOpDescMaker = 2 + kGradOpDescMaker = 2, + kVarTypeInference = 3 }; template @@ -38,7 +40,9 @@ struct OpInfoFillTypeID { ? kOpProtoAndCheckerMaker : (std::is_base_of::value ? kGradOpDescMaker - : static_cast(-1))); + : (std::is_base_of::value + ? kVarTypeInference + : static_cast(-1)))); } }; @@ -105,6 +109,17 @@ struct OpInfoFiller { }; } }; + +template +struct OpInfoFiller { + void operator()(const char* op_type, OpInfo* info) const { + info->infer_var_type_ = [](const OpDescBind& fwd_op, BlockDescBind* block) { + T inference; + inference(fwd_op, block); + }; + } +}; + } // namespace details } // namespace framework diff --git a/paddle/framework/op_desc.cc b/paddle/framework/op_desc.cc index a5d515bbca..09a544fb9e 100644 --- a/paddle/framework/op_desc.cc +++ b/paddle/framework/op_desc.cc @@ -236,5 +236,19 @@ void OpDescBind::InferShape(const BlockDescBind &block) const { it->second(&ctx); } +void OpDescBind::InferVarType(BlockDescBind *block) const { + auto &info = OpInfoMap::Instance().Get(this->Type()); + if (info.infer_var_type_) { + info.infer_var_type_(*this, block); + } else { + // all output type is LoDTensor by default + for (auto &out_pair : this->outputs_) { + for (auto &out_var_name : out_pair.second) { + block->Var(out_var_name)->SetType(VarDesc::LOD_TENSOR); + } + } + } +} + } // namespace framework } // namespace paddle diff --git a/paddle/framework/op_desc.h b/paddle/framework/op_desc.h index 90155fadea..d05ee0875d 100644 --- a/paddle/framework/op_desc.h +++ b/paddle/framework/op_desc.h @@ -104,6 +104,8 @@ class OpDescBind { void InferShape(const BlockDescBind &block) const; + void InferVarType(BlockDescBind *block) const; + private: template static std::vector MapKeys(const MapType &map) { diff --git a/paddle/framework/op_info.h b/paddle/framework/op_info.h index c504f69e30..e926180780 100644 --- a/paddle/framework/op_info.h +++ b/paddle/framework/op_info.h @@ -19,7 +19,6 @@ #include #include "paddle/framework/attribute.h" -#include "paddle/framework/op_desc.h" #include "paddle/framework/type_defs.h" #include "paddle/platform/macros.h" @@ -31,6 +30,7 @@ struct OpInfo { GradOpMakerFN grad_op_maker_; OpProto* proto_{nullptr}; OpAttrChecker* checker_{nullptr}; + InferVarTypeFN infer_var_type_; bool HasOpProtoAndChecker() const { return proto_ != nullptr && checker_ != nullptr; diff --git a/paddle/framework/type_defs.h b/paddle/framework/type_defs.h index 7e1b79c97b..a4e8253bfd 100644 --- a/paddle/framework/type_defs.h +++ b/paddle/framework/type_defs.h @@ -16,12 +16,18 @@ #include #include #include +#include +#include +#include +#include #include "paddle/platform/variant.h" namespace paddle { namespace framework { class OperatorBase; class OpDescBind; +class BlockDescBind; +class BlockDesc; using VariableNameMap = std::map>; // The order should be as same as framework.proto @@ -39,5 +45,8 @@ using OpCreator = std::function>( const OpDescBind&, const std::unordered_set& /*no_grad_set*/)>; +using InferVarTypeFN = std::function; + } // namespace framework } // namespace paddle diff --git a/paddle/framework/var_type_inference.h b/paddle/framework/var_type_inference.h new file mode 100644 index 0000000000..32abbeb334 --- /dev/null +++ b/paddle/framework/var_type_inference.h @@ -0,0 +1,29 @@ +/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + + Licensed under the Apache License, Version 2.0 (the "License"); + you may not use this file except in compliance with the License. + You may obtain a copy of the License at + + http://www.apache.org/licenses/LICENSE-2.0 + + Unless required by applicable law or agreed to in writing, software + distributed under the License is distributed on an "AS IS" BASIS, + WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + See the License for the specific language governing permissions and + limitations under the License. */ + +#pragma once +#include "paddle/framework/type_defs.h" + +namespace paddle { +namespace framework { + +class VarTypeInference { + public: + virtual ~VarTypeInference() {} + virtual void operator()(const OpDescBind& op_desc, + BlockDescBind* block) const = 0; +}; + +} // namespace framework +} // namespace paddle diff --git a/paddle/framework/var_type_inference_test.cc b/paddle/framework/var_type_inference_test.cc new file mode 100644 index 0000000000..e3f4893f1a --- /dev/null +++ b/paddle/framework/var_type_inference_test.cc @@ -0,0 +1,103 @@ +/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + + Licensed under the Apache License, Version 2.0 (the "License"); + you may not use this file except in compliance with the License. + You may obtain a copy of the License at + + http://www.apache.org/licenses/LICENSE-2.0 + + Unless required by applicable law or agreed to in writing, software + distributed under the License is distributed on an "AS IS" BASIS, + WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + See the License for the specific language governing permissions and + limitations under the License. */ + +#include "paddle/framework/var_type_inference.h" +#include "gtest/gtest.h" +#include "paddle/framework/op_registry.h" +#include "paddle/framework/operator.h" +#include "paddle/framework/program_desc.h" + +namespace paddle { +namespace framework { + +class SumOpMaker : public OpProtoAndCheckerMaker { + public: + SumOpMaker(OpProto *proto, OpAttrChecker *op_checker) + : OpProtoAndCheckerMaker(proto, op_checker) { + AddInput("X", "").AsDuplicable(); + AddOutput("Out", ""); + AddComment(""); + } +}; + +class SumOpVarTypeInference : public VarTypeInference { + public: + void operator()(const OpDescBind &op_desc, + BlockDescBind *block) const override { + auto default_var_type = VarDesc::LOD_TENSOR; + for (auto &in_var_name : op_desc.Input("X")) { + auto in_var_type = block->Var(in_var_name)->GetType(); + if (in_var_type != default_var_type) { + default_var_type = in_var_type; + break; + } + } + auto out_var_name = op_desc.Output("Out").front(); + block->Var(out_var_name)->SetType(default_var_type); + } +}; +} // namespace framework +} // namespace paddle + +REGISTER_OPERATOR(sum, paddle::framework::NOP, paddle::framework::SumOpMaker, + paddle::framework::SumOpVarTypeInference); +REGISTER_OPERATOR(sum_without_infer_var_type, paddle::framework::NOP, + paddle::framework::SumOpMaker); + +namespace paddle { +namespace framework { + +TEST(InferVarType, sum_op) { + auto &prog = ProgramDescBind::Instance(&GetProgramDesc()); + auto *op = prog.Block(0)->AppendOp(); + op->SetType("sum"); + op->SetInput("X", {"test_a", "test_b", "test_c"}); + op->SetOutput("Out", {"test_out"}); + + prog.Block(0)->NewVar("test_a")->SetType(VarDesc_VarType_LOD_TENSOR); + prog.Block(0)->NewVar("test_b")->SetType(VarDesc_VarType_LOD_TENSOR); + prog.Block(0)->NewVar("test_c")->SetType(VarDesc_VarType_LOD_TENSOR); + prog.Block(0)->NewVar("test_out"); + + op->InferVarType(prog.Block(0)); + + ASSERT_EQ(VarDesc_VarType_LOD_TENSOR, + prog.Block(0)->Var("test_out")->GetType()); + + prog.Block(0)->Var("test_b")->SetType(VarDesc_VarType_SELECTED_ROWS); + op->InferVarType(prog.Block(0)); + ASSERT_EQ(VarDesc_VarType_SELECTED_ROWS, + prog.Block(0)->Var("test_out")->GetType()); +} + +TEST(InferVarType, sum_op_without_infer_var_type) { + auto &prog = ProgramDescBind::Instance(&GetProgramDesc()); + auto *op = prog.Block(0)->AppendOp(); + op->SetType("sum_without_infer_var_type"); + op->SetInput("X", {"test2_a", "test2_b", "test2_c"}); + op->SetOutput("Out", {"test2_out"}); + + prog.Block(0)->NewVar("test2_a")->SetType(VarDesc_VarType_LOD_TENSOR); + prog.Block(0)->NewVar("test2_b")->SetType(VarDesc_VarType_SELECTED_ROWS); + prog.Block(0)->NewVar("test2_c")->SetType(VarDesc_VarType_LOD_TENSOR); + prog.Block(0)->NewVar("test2_out"); + + op->InferVarType(prog.Block(0)); + + ASSERT_EQ(VarDesc_VarType_LOD_TENSOR, + prog.Block(0)->Var("test2_out")->GetType()); +} + +} // namespace framework +} // namespace paddle \ No newline at end of file From 5be10872f95c18434bce79c0111717efa1994029 Mon Sep 17 00:00:00 2001 From: qijun Date: Fri, 13 Oct 2017 14:59:20 -0700 Subject: [PATCH 08/36] add selected_rows add cpu functor --- paddle/operators/cross_entropy_op.cu | 3 +- paddle/operators/cross_entropy_op.h | 3 +- paddle/operators/math/math_function.cc | 60 +++++++++++++++++++++ paddle/operators/math/math_function.h | 22 ++++++-- paddle/operators/math/math_function_test.cc | 60 +++++++++++++++++++-- 5 files changed, 137 insertions(+), 11 deletions(-) diff --git a/paddle/operators/cross_entropy_op.cu b/paddle/operators/cross_entropy_op.cu index 5e2024e0ea..07b0388b60 100644 --- a/paddle/operators/cross_entropy_op.cu +++ b/paddle/operators/cross_entropy_op.cu @@ -91,7 +91,8 @@ class CrossEntropyGradientOpCUDAKernel : public framework::OpKernel { .stream()>>>(dx_data, dy_data, x_data, label_data, batch_size, class_num); } else { - math::SetConstant(ctx.device_context(), dx, 0); + math::SetConstant functor; + functor(ctx.device_context(), dx, 0); auto* label_data = label->data(); grid = (batch_size + block - 1) / block; CrossEntropyGradientKernel<<< diff --git a/paddle/operators/cross_entropy_op.h b/paddle/operators/cross_entropy_op.h index d2d321aa7e..19c276d23f 100644 --- a/paddle/operators/cross_entropy_op.h +++ b/paddle/operators/cross_entropy_op.h @@ -70,7 +70,8 @@ class CrossEntropyGradientOpKernel : public framework::OpKernel { const T* x_data = x->data(); const int* label_data = label->data(); - math::SetConstant(ctx.device_context(), dx, 0); + math::SetConstant functor; + functor(ctx.device_context(), dx, 0); for (int i = 0; i < batch_size; ++i) { PADDLE_ASSERT(label_data[i] >= 0 || label_data[i] < class_num); diff --git a/paddle/operators/math/math_function.cc b/paddle/operators/math/math_function.cc index ba653afa2c..75a705b346 100644 --- a/paddle/operators/math/math_function.cc +++ b/paddle/operators/math/math_function.cc @@ -13,6 +13,7 @@ See the License for the specific language governing permissions and limitations under the License. */ #include "paddle/operators/math/math_function.h" +#include namespace paddle { namespace operators { @@ -130,6 +131,65 @@ void matmul( matrix_b.data(), beta, matrix_out->data()); } +template struct SetConstant; + +namespace detail { +size_t FindPos(const std::vector& rows, int64_t value) { + for (size_t i = 0; i < rows.size(); i++) { + if (rows[i] == value) { + return i; + } + } + return 0; +} +} // namespace detail + +template +struct SelectedRowsAdd { + void operator()(const platform::DeviceContext& context, + const framework::SelectedRows& input1, + const framework::SelectedRows& input2, + framework::SelectedRows* output) { + auto in1_height = input1.height(); + PADDLE_ENFORCE_EQ(in1_height, input2.height()); + PADDLE_ENFORCE_EQ(in1_height, output->height()); + + auto& in1_rows = input1.rows(); + auto& in2_rows = input2.rows(); + auto& out_rows = output->rows(); + + auto* out_value = output->mutable_value(); + auto& in1_value = input1.value(); + auto& in2_value = input2.value(); + + auto in1_row_numel = in1_value.numel() / in1_rows.size(); + PADDLE_ENFORCE_EQ(in1_row_numel, in2_value.numel() / in2_rows.size()); + PADDLE_ENFORCE_EQ(in1_row_numel, out_value->numel() / out_rows.size()); + + SetConstant functor; + functor(context, out_value, 0.0); + auto* out_data = out_value->data(); + + auto* in1_data = in1_value.data(); + for (size_t i = 0; i < in1_rows.size(); i++) { + auto row = detail::FindPos(out_rows, in1_rows[i]); + for (size_t j = 0; j < in1_row_numel; j++) { + out_data[row * in1_row_numel + j] += in1_data[i * in1_row_numel + j]; + } + } + + auto* in2_data = in2_value.data(); + for (size_t i = 0; i < in2_rows.size(); i++) { + auto row = detail::FindPos(out_rows, in2_rows[i]); + for (size_t j = 0; j < in1_row_numel; j++) { + out_data[row * in1_row_numel + j] += in2_data[i * in1_row_numel + j]; + } + } + } +}; + +template struct SelectedRowsAdd; + } // namespace math } // namespace operators } // namespace paddle diff --git a/paddle/operators/math/math_function.h b/paddle/operators/math/math_function.h index 473eff4d19..f298f34bab 100644 --- a/paddle/operators/math/math_function.h +++ b/paddle/operators/math/math_function.h @@ -53,6 +53,7 @@ int LAPACKE_dgetri(int matrix_layout, int n, double* a, int lda, #include #include "paddle/framework/eigen.h" +#include "paddle/framework/selected_rows.h" #include "paddle/framework/tensor.h" #include "paddle/platform/device_context.h" #include "paddle/platform/enforce.h" @@ -86,11 +87,22 @@ void matmul(const platform::DeviceContext& context, framework::Tensor* matrix_out, T beta); template -void SetConstant(const platform::DeviceContext& context, - framework::Tensor* tensor, T num) { - auto t = framework::EigenVector::Flatten(*tensor); - t.device(*context.GetEigenDevice()) = t.constant(static_cast(num)); -} +struct SetConstant { + void operator()(const platform::DeviceContext& context, + framework::Tensor* tensor, T num) { + auto t = framework::EigenVector::Flatten(*tensor); + t.device(*context.GetEigenDevice()) = + t.constant(static_cast(num)); + } +}; + +template +struct SelectedRowsAdd { + void operator()(const platform::DeviceContext& context, + const framework::SelectedRows& input1, + const framework::SelectedRows& input2, + framework::SelectedRows* output); +}; } // namespace math } // namespace operators diff --git a/paddle/operators/math/math_function_test.cc b/paddle/operators/math/math_function_test.cc index c87d200c3a..43760bc601 100644 --- a/paddle/operators/math/math_function_test.cc +++ b/paddle/operators/math/math_function_test.cc @@ -1,4 +1,5 @@ #include "paddle/operators/math/math_function.h" +#include "glog/logging.h" #include "gtest/gtest.h" #ifdef PADDLE_WITH_CUDA @@ -253,18 +254,69 @@ TEST(math_function, zero) { auto* cpu_place = new paddle::platform::CPUPlace(); float* t = tensor.mutable_data({2, 2}, *cpu_place); paddle::platform::CPUDeviceContext context(*cpu_place); - paddle::operators::math::SetConstant( - context, &tensor, 0); + paddle::operators::math::SetConstant + functor; + functor(context, &tensor, 0); EXPECT_EQ(t[0], 0); EXPECT_EQ(t[1], 0); EXPECT_EQ(t[2], 0); EXPECT_EQ(t[3], 0); - paddle::operators::math::SetConstant( - context, &tensor, 1); + functor(context, &tensor, 1); EXPECT_EQ(t[0], 1); EXPECT_EQ(t[1], 1); EXPECT_EQ(t[2], 1); EXPECT_EQ(t[3], 1); } + +TEST(math_function, selected_rows_add) { + using namespace paddle::framework; + using namespace paddle::platform; + using namespace paddle::operators::math; + + CPUPlace cpu_place; + CPUDeviceContext ctx(cpu_place); + SetConstant functor; + int64_t height = 10; + int64_t row_numel = 10; + + std::vector rows1{0, 4, 7}; + std::unique_ptr selected_rows1{new SelectedRows(rows1, height)}; + auto* in1_value = selected_rows1->mutable_value(); + in1_value->mutable_data( + make_ddim({static_cast(rows1.size()), row_numel}), cpu_place); + functor(ctx, in1_value, 2.0); + + std::vector rows2{0, 5, 7, 9}; + std::unique_ptr selected_rows2{new SelectedRows(rows2, height)}; + auto* in2_value = selected_rows2->mutable_value(); + in2_value->mutable_data( + make_ddim({static_cast(rows2.size()), row_numel}), cpu_place); + functor(ctx, in2_value, 1.0); + + std::unique_ptr output{new SelectedRows()}; + output->set_height(height); + std::vector out_rows = {0, 4, 5, 7, 9}; + output->set_rows(out_rows); + + auto* out_value = output->mutable_value(); + out_value->mutable_data(make_ddim({5, 10}), cpu_place); + + SelectedRowsAdd add_functor; + add_functor(ctx, *selected_rows1, *selected_rows2, output.get()); + + auto* data = output->value().data(); + // out_rows[0] = 0 + EXPECT_EQ(data[0 * row_numel + 0], 3.0); + EXPECT_EQ(data[0 * row_numel + 8], 3.0); + // out_rows[1] = 4 + EXPECT_EQ(data[1 * row_numel + 1], 2.0); + // out_rows[2] = 5 + EXPECT_EQ(data[2 * row_numel + 6], 1.0); + // out_rows[3] = 7 + EXPECT_EQ(data[3 * row_numel + 3], 3.0); + EXPECT_EQ(data[3 * row_numel + 8], 3.0); + // out_rows[4] = 9 + EXPECT_EQ(data[4 * row_numel + 4], 1.0); +} From 7b1834330537594b2dc9267d9249f9d84fef391b Mon Sep 17 00:00:00 2001 From: qijun Date: Fri, 13 Oct 2017 15:02:52 -0700 Subject: [PATCH 09/36] remove unused header file --- paddle/operators/math/math_function.cc | 1 - 1 file changed, 1 deletion(-) diff --git a/paddle/operators/math/math_function.cc b/paddle/operators/math/math_function.cc index 75a705b346..306612b65f 100644 --- a/paddle/operators/math/math_function.cc +++ b/paddle/operators/math/math_function.cc @@ -13,7 +13,6 @@ See the License for the specific language governing permissions and limitations under the License. */ #include "paddle/operators/math/math_function.h" -#include namespace paddle { namespace operators { From 931572e2109d25c825d1d55b17358020c4d4c317 Mon Sep 17 00:00:00 2001 From: qijun Date: Fri, 13 Oct 2017 16:25:22 -0700 Subject: [PATCH 10/36] SelectedRowsAddTensor method --- paddle/framework/selected_rows.h | 3 + paddle/operators/math/math_function.cc | 68 ++++++++++++++---- paddle/operators/math/math_function.h | 9 +++ paddle/operators/math/math_function_test.cc | 79 +++++++++++++++------ 4 files changed, 124 insertions(+), 35 deletions(-) diff --git a/paddle/framework/selected_rows.h b/paddle/framework/selected_rows.h index f9f563051e..ddc6dec194 100644 --- a/paddle/framework/selected_rows.h +++ b/paddle/framework/selected_rows.h @@ -45,6 +45,9 @@ class SelectedRows { } private: + // Notice: rows can be duplicate. We can have {0, 4, 7, 0, 5, 7, 9} here. + // SelectedRows are simplely concated when adding together. Until a + // SelectedRows add a Tensor, will the duplicate rows be handled. std::vector rows_; std::unique_ptr value_{nullptr}; int64_t height_; diff --git a/paddle/operators/math/math_function.cc b/paddle/operators/math/math_function.cc index 306612b65f..ed49a0a549 100644 --- a/paddle/operators/math/math_function.cc +++ b/paddle/operators/math/math_function.cc @@ -13,6 +13,8 @@ See the License for the specific language governing permissions and limitations under the License. */ #include "paddle/operators/math/math_function.h" +#include "paddle/framework/eigen.h" +#include "paddle/memory/memcpy.h" namespace paddle { namespace operators { @@ -151,11 +153,17 @@ struct SelectedRowsAdd { framework::SelectedRows* output) { auto in1_height = input1.height(); PADDLE_ENFORCE_EQ(in1_height, input2.height()); - PADDLE_ENFORCE_EQ(in1_height, output->height()); + output->set_height(in1_height); auto& in1_rows = input1.rows(); auto& in2_rows = input2.rows(); - auto& out_rows = output->rows(); + std::vector out_rows; + out_rows.reserve(in1_rows.size() + in2_rows.size()); + + // concat rows + out_rows.insert(out_rows.end(), in1_rows.begin(), in1_rows.end()); + out_rows.insert(out_rows.end(), in2_rows.begin(), in2_rows.end()); + output->set_rows(out_rows); auto* out_value = output->mutable_value(); auto& in1_value = input1.value(); @@ -165,29 +173,59 @@ struct SelectedRowsAdd { PADDLE_ENFORCE_EQ(in1_row_numel, in2_value.numel() / in2_rows.size()); PADDLE_ENFORCE_EQ(in1_row_numel, out_value->numel() / out_rows.size()); - SetConstant functor; - functor(context, out_value, 0.0); auto* out_data = out_value->data(); auto* in1_data = in1_value.data(); - for (size_t i = 0; i < in1_rows.size(); i++) { - auto row = detail::FindPos(out_rows, in1_rows[i]); - for (size_t j = 0; j < in1_row_numel; j++) { - out_data[row * in1_row_numel + j] += in1_data[i * in1_row_numel + j]; - } - } + memory::Copy(platform::CPUPlace(), out_data, platform::CPUPlace(), in1_data, + in1_value.numel() * sizeof(T)); auto* in2_data = in2_value.data(); - for (size_t i = 0; i < in2_rows.size(); i++) { - auto row = detail::FindPos(out_rows, in2_rows[i]); - for (size_t j = 0; j < in1_row_numel; j++) { - out_data[row * in1_row_numel + j] += in2_data[i * in1_row_numel + j]; + memory::Copy(platform::CPUPlace(), out_data + in1_value.numel(), + platform::CPUPlace(), in2_data, in2_value.numel() * sizeof(T)); + } +}; + +template struct SelectedRowsAdd; + +template +struct SelectedRowsAddTensor { + void operator()(const platform::DeviceContext& context, + const framework::SelectedRows& input1, + const framework::Tensor& input2, framework::Tensor* output) { + auto in1_height = input1.height(); + auto in2_dims = input2.dims(); + auto out_dims = output->dims(); + PADDLE_ENFORCE_EQ(in1_height, in2_dims[0]); + PADDLE_ENFORCE_EQ(in1_height, out_dims[0]); + + auto& in1_value = input1.value(); + auto& in1_rows = input1.rows(); + + int64_t in1_row_numel = in1_value.numel() / in1_rows.size(); + PADDLE_ENFORCE_EQ(in1_row_numel, input2.numel() / in1_height); + PADDLE_ENFORCE_EQ(in1_row_numel, output->numel() / in1_height); + + SetConstant functor; + functor(context, output, 0.0); + + auto* in1_data = in1_value.data(); + auto* out_data = output->data(); + + for (size_t i = 0; i < in1_rows.size(); i++) { + for (int64_t j = 0; j < in1_row_numel; j++) { + out_data[in1_rows[i] * in1_row_numel + j] += + in1_data[i * in1_row_numel + j]; } } + + auto out_eigen = framework::EigenVector::Flatten(*output); + auto in2_eigen = framework::EigenVector::Flatten(input2); + out_eigen.device(*context.GetEigenDevice()) = + out_eigen + in2_eigen; } }; -template struct SelectedRowsAdd; +template struct SelectedRowsAddTensor; } // namespace math } // namespace operators diff --git a/paddle/operators/math/math_function.h b/paddle/operators/math/math_function.h index f298f34bab..0d0d4cdd73 100644 --- a/paddle/operators/math/math_function.h +++ b/paddle/operators/math/math_function.h @@ -96,6 +96,8 @@ struct SetConstant { } }; +// SelectedRows + SelectedRows will simplely concat value and rows. +// The real computation happens in dealing with LoDTensor. template struct SelectedRowsAdd { void operator()(const platform::DeviceContext& context, @@ -104,6 +106,13 @@ struct SelectedRowsAdd { framework::SelectedRows* output); }; +template +struct SelectedRowsAddTensor { + void operator()(const platform::DeviceContext& context, + const framework::SelectedRows& input1, + const framework::Tensor& input2, framework::Tensor* output); +}; + } // namespace math } // namespace operators } // namespace paddle diff --git a/paddle/operators/math/math_function_test.cc b/paddle/operators/math/math_function_test.cc index 43760bc601..e3186171d1 100644 --- a/paddle/operators/math/math_function_test.cc +++ b/paddle/operators/math/math_function_test.cc @@ -286,37 +286,76 @@ TEST(math_function, selected_rows_add) { auto* in1_value = selected_rows1->mutable_value(); in1_value->mutable_data( make_ddim({static_cast(rows1.size()), row_numel}), cpu_place); - functor(ctx, in1_value, 2.0); + functor(ctx, in1_value, 1.0); std::vector rows2{0, 5, 7, 9}; std::unique_ptr selected_rows2{new SelectedRows(rows2, height)}; auto* in2_value = selected_rows2->mutable_value(); in2_value->mutable_data( make_ddim({static_cast(rows2.size()), row_numel}), cpu_place); - functor(ctx, in2_value, 1.0); + functor(ctx, in2_value, 2.0); std::unique_ptr output{new SelectedRows()}; - output->set_height(height); - std::vector out_rows = {0, 4, 5, 7, 9}; - output->set_rows(out_rows); - auto* out_value = output->mutable_value(); - out_value->mutable_data(make_ddim({5, 10}), cpu_place); + + // simplely concat two SelectedRows + out_value->mutable_data(make_ddim({7, 10}), cpu_place); SelectedRowsAdd add_functor; add_functor(ctx, *selected_rows1, *selected_rows2, output.get()); - auto* data = output->value().data(); - // out_rows[0] = 0 - EXPECT_EQ(data[0 * row_numel + 0], 3.0); - EXPECT_EQ(data[0 * row_numel + 8], 3.0); - // out_rows[1] = 4 - EXPECT_EQ(data[1 * row_numel + 1], 2.0); - // out_rows[2] = 5 - EXPECT_EQ(data[2 * row_numel + 6], 1.0); - // out_rows[3] = 7 - EXPECT_EQ(data[3 * row_numel + 3], 3.0); - EXPECT_EQ(data[3 * row_numel + 8], 3.0); - // out_rows[4] = 9 - EXPECT_EQ(data[4 * row_numel + 4], 1.0); + auto out_height = output->height(); + EXPECT_EQ(out_height, height); + + auto& out_rows = output->rows(); + + // input1 rows + EXPECT_EQ(out_rows[0], 0); + EXPECT_EQ(out_rows[1], 4); + EXPECT_EQ(out_rows[2], 7); + // input2 rows + EXPECT_EQ(out_rows[3], 0); + EXPECT_EQ(out_rows[4], 5); + EXPECT_EQ(out_rows[5], 7); + EXPECT_EQ(out_rows[6], 9); + + auto* out_data = output->value().data(); + // input1 value + EXPECT_EQ(out_data[0 * row_numel + 0], 1.0); + EXPECT_EQ(out_data[0 * row_numel + 8], 1.0); + EXPECT_EQ(out_data[1 * row_numel + 1], 1.0); + EXPECT_EQ(out_data[2 * row_numel + 6], 1.0); + // input2 value + EXPECT_EQ(out_data[3 * row_numel + 3], 2.0); + EXPECT_EQ(out_data[3 * row_numel + 8], 2.0); + EXPECT_EQ(out_data[4 * row_numel + 4], 2.0); + EXPECT_EQ(out_data[5 * row_numel + 7], 2.0); + EXPECT_EQ(out_data[6 * row_numel + 9], 2.0); + + std::unique_ptr tensor1{new Tensor()}; + tensor1->mutable_data(make_ddim({height, row_numel}), cpu_place); + SetConstant constant_functor; + constant_functor(ctx, tensor1.get(), 3.0); + + std::unique_ptr tensor2{new Tensor()}; + tensor2->mutable_data(make_ddim({height, row_numel}), cpu_place); + + SelectedRowsAddTensor add_tensor_functor; + add_tensor_functor(ctx, *output, *tensor1, tensor2.get()); + + auto* tensor2_data = tensor2->data(); + // row0: 1.0 + 2.0 + 3.0 + EXPECT_EQ(tensor2_data[0 * row_numel + 0], 6.0); + // row1: 3.0 + EXPECT_EQ(tensor2_data[1 * row_numel + 1], 3.0); + // row4 : 1.0 + 3.0 + EXPECT_EQ(tensor2_data[4 * row_numel + 6], 4.0); + // row5: 2.0 + 3.0 + EXPECT_EQ(tensor2_data[5 * row_numel + 7], 5.0); + // row6: 3.0 + EXPECT_EQ(tensor2_data[6 * row_numel + 1], 3.0); + // row7: 1.0 + 2.0 + 3.0 + EXPECT_EQ(tensor2_data[7 * row_numel + 3], 6.0); + // row9: 2.0 + 3.0 + EXPECT_EQ(tensor2_data[9 * row_numel + 6], 5.0); } From 4741266d6fcac37c678e7815a2b93994fcddaec7 Mon Sep 17 00:00:00 2001 From: qijun Date: Fri, 13 Oct 2017 16:27:39 -0700 Subject: [PATCH 11/36] remove unused method --- paddle/operators/math/math_function.cc | 11 ----------- paddle/operators/math/math_function_test.cc | 1 - 2 files changed, 12 deletions(-) diff --git a/paddle/operators/math/math_function.cc b/paddle/operators/math/math_function.cc index ed49a0a549..ddb904aa46 100644 --- a/paddle/operators/math/math_function.cc +++ b/paddle/operators/math/math_function.cc @@ -134,17 +134,6 @@ void matmul( template struct SetConstant; -namespace detail { -size_t FindPos(const std::vector& rows, int64_t value) { - for (size_t i = 0; i < rows.size(); i++) { - if (rows[i] == value) { - return i; - } - } - return 0; -} -} // namespace detail - template struct SelectedRowsAdd { void operator()(const platform::DeviceContext& context, diff --git a/paddle/operators/math/math_function_test.cc b/paddle/operators/math/math_function_test.cc index e3186171d1..fe0d1981fa 100644 --- a/paddle/operators/math/math_function_test.cc +++ b/paddle/operators/math/math_function_test.cc @@ -1,5 +1,4 @@ #include "paddle/operators/math/math_function.h" -#include "glog/logging.h" #include "gtest/gtest.h" #ifdef PADDLE_WITH_CUDA From 29819ba7646d5a44d927347d49ad6d6ab36039c9 Mon Sep 17 00:00:00 2001 From: Yu Yang Date: Sat, 14 Oct 2017 14:29:44 -0700 Subject: [PATCH 12/36] Fix unittest --- paddle/framework/var_type_inference_test.cc | 37 +++++++++++---------- 1 file changed, 19 insertions(+), 18 deletions(-) diff --git a/paddle/framework/var_type_inference_test.cc b/paddle/framework/var_type_inference_test.cc index e3f4893f1a..97b8c64748 100644 --- a/paddle/framework/var_type_inference_test.cc +++ b/paddle/framework/var_type_inference_test.cc @@ -35,14 +35,17 @@ class SumOpVarTypeInference : public VarTypeInference { public: void operator()(const OpDescBind &op_desc, BlockDescBind *block) const override { - auto default_var_type = VarDesc::LOD_TENSOR; - for (auto &in_var_name : op_desc.Input("X")) { - auto in_var_type = block->Var(in_var_name)->GetType(); - if (in_var_type != default_var_type) { - default_var_type = in_var_type; - break; - } + auto &inputs = op_desc.Input("X"); + auto default_var_type = VarDesc::SELECTED_ROWS; + + bool any_input_is_lod_tensor = std::any_of( + inputs.begin(), inputs.end(), [block](const std::string &name) { + return block->Var(name)->GetType() == VarDesc::LOD_TENSOR; + }); + if (any_input_is_lod_tensor) { + default_var_type = VarDesc::LOD_TENSOR; } + auto out_var_name = op_desc.Output("Out").front(); block->Var(out_var_name)->SetType(default_var_type); } @@ -65,20 +68,18 @@ TEST(InferVarType, sum_op) { op->SetInput("X", {"test_a", "test_b", "test_c"}); op->SetOutput("Out", {"test_out"}); - prog.Block(0)->NewVar("test_a")->SetType(VarDesc_VarType_LOD_TENSOR); - prog.Block(0)->NewVar("test_b")->SetType(VarDesc_VarType_LOD_TENSOR); - prog.Block(0)->NewVar("test_c")->SetType(VarDesc_VarType_LOD_TENSOR); + prog.Block(0)->NewVar("test_a")->SetType(VarDesc::SELECTED_ROWS); + prog.Block(0)->NewVar("test_b")->SetType(VarDesc::SELECTED_ROWS); + prog.Block(0)->NewVar("test_c")->SetType(VarDesc::SELECTED_ROWS); prog.Block(0)->NewVar("test_out"); op->InferVarType(prog.Block(0)); - ASSERT_EQ(VarDesc_VarType_LOD_TENSOR, - prog.Block(0)->Var("test_out")->GetType()); + ASSERT_EQ(VarDesc::SELECTED_ROWS, prog.Block(0)->Var("test_out")->GetType()); - prog.Block(0)->Var("test_b")->SetType(VarDesc_VarType_SELECTED_ROWS); + prog.Block(0)->Var("test_b")->SetType(VarDesc::LOD_TENSOR); op->InferVarType(prog.Block(0)); - ASSERT_EQ(VarDesc_VarType_SELECTED_ROWS, - prog.Block(0)->Var("test_out")->GetType()); + ASSERT_EQ(VarDesc::LOD_TENSOR, prog.Block(0)->Var("test_out")->GetType()); } TEST(InferVarType, sum_op_without_infer_var_type) { @@ -88,9 +89,9 @@ TEST(InferVarType, sum_op_without_infer_var_type) { op->SetInput("X", {"test2_a", "test2_b", "test2_c"}); op->SetOutput("Out", {"test2_out"}); - prog.Block(0)->NewVar("test2_a")->SetType(VarDesc_VarType_LOD_TENSOR); - prog.Block(0)->NewVar("test2_b")->SetType(VarDesc_VarType_SELECTED_ROWS); - prog.Block(0)->NewVar("test2_c")->SetType(VarDesc_VarType_LOD_TENSOR); + prog.Block(0)->NewVar("test2_a")->SetType(VarDesc::SELECTED_ROWS); + prog.Block(0)->NewVar("test2_b")->SetType(VarDesc::SELECTED_ROWS); + prog.Block(0)->NewVar("test2_c")->SetType(VarDesc::SELECTED_ROWS); prog.Block(0)->NewVar("test2_out"); op->InferVarType(prog.Block(0)); From f59a7c1d36d2e930c48e118ad90f35a541e00223 Mon Sep 17 00:00:00 2001 From: qijun Date: Sat, 14 Oct 2017 14:57:07 -0700 Subject: [PATCH 13/36] add gpu functor for SelectedRows --- paddle/framework/lod_tensor.h | 3 - paddle/framework/selected_rows.h | 7 +- paddle/framework/selected_rows_test.cc | 2 +- paddle/framework/type_defs.h | 6 +- paddle/operators/math/CMakeLists.txt | 2 +- paddle/operators/math/math_function.cc | 17 +- paddle/operators/math/math_function.cu | 107 ++++++++ paddle/operators/math/math_function_test.cc | 179 ------------- paddle/operators/math/math_function_test.cu | 277 ++++++++++++++++++++ 9 files changed, 405 insertions(+), 195 deletions(-) create mode 100644 paddle/operators/math/math_function_test.cu diff --git a/paddle/framework/lod_tensor.h b/paddle/framework/lod_tensor.h index 4db36ee766..ee040a9144 100644 --- a/paddle/framework/lod_tensor.h +++ b/paddle/framework/lod_tensor.h @@ -1,11 +1,8 @@ /* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. - Licensed under the Apache License, Version 2.0 (the "License"); you may not use this file except in compliance with the License. You may obtain a copy of the License at - http://www.apache.org/licenses/LICENSE-2.0 - Unless required by applicable law or agreed to in writing, software distributed under the License is distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. diff --git a/paddle/framework/selected_rows.h b/paddle/framework/selected_rows.h index ddc6dec194..cd90781371 100644 --- a/paddle/framework/selected_rows.h +++ b/paddle/framework/selected_rows.h @@ -10,6 +10,7 @@ See the License for the specific language governing permissions and limitations under the License. */ #pragma once +#include "paddle/framework/lod_tensor.h" #include "paddle/framework/tensor.h" namespace paddle { @@ -34,9 +35,9 @@ class SelectedRows { void set_height(int64_t height) { height_ = height; } - const std::vector& rows() const { return rows_; } + const Vector& rows() const { return rows_; } - void set_rows(const std::vector& rows) { rows_ = rows; } + void set_rows(const Vector& rows) { rows_ = rows; } DDim GetCompleteDims() const { std::vector dims = vectorize(value_->dims()); @@ -48,7 +49,7 @@ class SelectedRows { // Notice: rows can be duplicate. We can have {0, 4, 7, 0, 5, 7, 9} here. // SelectedRows are simplely concated when adding together. Until a // SelectedRows add a Tensor, will the duplicate rows be handled. - std::vector rows_; + Vector rows_; std::unique_ptr value_{nullptr}; int64_t height_; }; diff --git a/paddle/framework/selected_rows_test.cc b/paddle/framework/selected_rows_test.cc index 4ee13a65d7..055b867600 100644 --- a/paddle/framework/selected_rows_test.cc +++ b/paddle/framework/selected_rows_test.cc @@ -18,7 +18,7 @@ namespace framework { class SelectedRowsTester : public ::testing::Test { public: virtual void SetUp() override { - std::vector rows{0, 4, 7}; + Vector rows{0, 4, 7}; int64_t height = 10; int64_t row_numel = 100; selected_rows_.reset(new SelectedRows(rows, height)); diff --git a/paddle/framework/type_defs.h b/paddle/framework/type_defs.h index 7e1b79c97b..0c0a72de31 100644 --- a/paddle/framework/type_defs.h +++ b/paddle/framework/type_defs.h @@ -1,11 +1,8 @@ /* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. - Licensed under the Apache License, Version 2.0 (the "License"); you may not use this file except in compliance with the License. You may obtain a copy of the License at - http://www.apache.org/licenses/LICENSE-2.0 - Unless required by applicable law or agreed to in writing, software distributed under the License is distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. @@ -37,7 +34,8 @@ using OpCreator = std::function; using GradOpMakerFN = std::function>( - const OpDescBind&, const std::unordered_set& /*no_grad_set*/)>; + const OpDescBind&, const std::unordered_set& /*no_grad_set*/, + std::unordered_map* /*grad_to_var*/)>; } // namespace framework } // namespace paddle diff --git a/paddle/operators/math/CMakeLists.txt b/paddle/operators/math/CMakeLists.txt index 1a2f623ce7..a7f275bae5 100644 --- a/paddle/operators/math/CMakeLists.txt +++ b/paddle/operators/math/CMakeLists.txt @@ -1,6 +1,6 @@ if(WITH_GPU) nv_library(math_function SRCS math_function.cc math_function.cu im2col.cc im2col.cu DEPS cblas device_context operator) - nv_test(math_function_test SRCS math_function_test.cc DEPS math_function tensor) + nv_test(math_function_gpu_test SRCS math_function_test.cu DEPS math_function tensor) nv_library(softmax SRCS softmax.cc softmax.cu DEPS operator) nv_library(cross_entropy SRCS cross_entropy.cc cross_entropy.cu DEPS operator) nv_library(pooling SRCS pooling.cc pooling.cu DEPS device_context) diff --git a/paddle/operators/math/math_function.cc b/paddle/operators/math/math_function.cc index ddb904aa46..a1faafb7c4 100644 --- a/paddle/operators/math/math_function.cc +++ b/paddle/operators/math/math_function.cc @@ -162,15 +162,24 @@ struct SelectedRowsAdd { PADDLE_ENFORCE_EQ(in1_row_numel, in2_value.numel() / in2_rows.size()); PADDLE_ENFORCE_EQ(in1_row_numel, out_value->numel() / out_rows.size()); - auto* out_data = out_value->data(); + auto in1_place = input1.place(); + PADDLE_ENFORCE(platform::is_cpu_place(in1_place)); + auto in2_place = input2.place(); + PADDLE_ENFORCE(platform::is_cpu_place(in2_place)); + auto out_place = context.GetPlace(); + PADDLE_ENFORCE(platform::is_cpu_place(out_place)); + auto* out_data = out_value->data(); auto* in1_data = in1_value.data(); - memory::Copy(platform::CPUPlace(), out_data, platform::CPUPlace(), in1_data, + memory::Copy(boost::get(out_place), out_data, + boost::get(in1_place), in1_data, in1_value.numel() * sizeof(T)); auto* in2_data = in2_value.data(); - memory::Copy(platform::CPUPlace(), out_data + in1_value.numel(), - platform::CPUPlace(), in2_data, in2_value.numel() * sizeof(T)); + memory::Copy(boost::get(out_place), + out_data + in1_value.numel(), + boost::get(in2_place), in2_data, + in2_value.numel() * sizeof(T)); } }; diff --git a/paddle/operators/math/math_function.cu b/paddle/operators/math/math_function.cu index 649f1f352c..26bf0ec2f1 100644 --- a/paddle/operators/math/math_function.cu +++ b/paddle/operators/math/math_function.cu @@ -155,6 +155,113 @@ void matmul( matrix_b.data(), beta, matrix_out->data()); } +template +struct SelectedRowsAdd { + void operator()(const platform::DeviceContext& context, + const framework::SelectedRows& input1, + const framework::SelectedRows& input2, + framework::SelectedRows* output) { + auto in1_height = input1.height(); + PADDLE_ENFORCE_EQ(in1_height, input2.height()); + output->set_height(in1_height); + + auto& in1_rows = input1.rows(); + auto& in2_rows = input2.rows(); + std::vector out_rows; + out_rows.reserve(in1_rows.size() + in2_rows.size()); + + // concat rows + out_rows.insert(out_rows.end(), in1_rows.begin(), in1_rows.end()); + out_rows.insert(out_rows.end(), in2_rows.begin(), in2_rows.end()); + output->set_rows(out_rows); + + auto* out_value = output->mutable_value(); + auto& in1_value = input1.value(); + auto& in2_value = input2.value(); + + auto in1_row_numel = in1_value.numel() / in1_rows.size(); + PADDLE_ENFORCE_EQ(in1_row_numel, in2_value.numel() / in2_rows.size()); + PADDLE_ENFORCE_EQ(in1_row_numel, out_value->numel() / out_rows.size()); + + auto* out_data = out_value->data(); + auto* in1_data = in1_value.data(); + + auto in1_place = input1.place(); + PADDLE_ENFORCE(platform::is_gpu_place(in1_place)); + auto in2_place = input2.place(); + PADDLE_ENFORCE(platform::is_gpu_place(in2_place)); + auto out_place = context.GetPlace(); + PADDLE_ENFORCE(platform::is_gpu_place(out_place)) + + memory::Copy( + boost::get(out_place), out_data, + boost::get(in1_place), in1_data, + in1_value.numel() * sizeof(T), + reinterpret_cast(context).stream()); + + auto* in2_data = in2_value.data(); + memory::Copy( + boost::get(out_place), out_data + in1_value.numel(), + boost::get(in2_place), in2_data, + in2_value.numel() * sizeof(T), + reinterpret_cast(context).stream()); + } +}; + +template struct SelectedRowsAdd; + +namespace { +template +__global__ void SelectedRowsAddTensorKernel(T* selected_rows, int64_t* rows, + T* tensor_in, T* tensor_out, + const int64_t row_numel) { + const ty = blockIdx.y; + int tid = threadIdx.x; + + selected_rows += ty * row_numel; + tensor_in += rows[ty] * row_numel; + tensor_out += rows[ty] * row_numel; + + for (int index = tid; index < row_numel; index += block_size) { + tensor_out[index] = tensor_in[index] + selected_rows[index]; + } +} +} + +template +struct SelectedRowsAddTensor { + void operator()(const platform::DeviceContext& context, + const framework::SelectedRows& input1, + const framework::Tensor& input2, framework::Tensor* output) { + auto in1_height = input1.height(); + auto in2_dims = input2.dims(); + auto out_dims = output->dims(); + PADDLE_ENFORCE_EQ(in1_height, in2_dims[0]); + PADDLE_ENFORCE_EQ(in1_height, out_dims[0]); + + auto& in1_value = input1.value(); + auto& in1_rows = input1.rows(); + + int64_t in1_row_numel = in1_value.numel() / in1_rows.size(); + PADDLE_ENFORCE_EQ(in1_row_numel, input2.numel() / in1_height); + PADDLE_ENFORCE_EQ(in1_row_numel, output->numel() / in1_height); + + auto* in1_data = in1_value.data(); + auto* in2_data = input2.data(); + auto* out_data = output->data(); + + const int block_size = 256; + dim3 threads(block_size, 1); + dim3 grid(1, in1_height); + SelectedRowsAddTensorKernel<<< + grid, threads, 0, + reinterpret_cast(ctx).stream()>>>( + in1_data, in1_rows.data(), in2_data, out_data, in1_row_numel); + } +}; + +template struct SelectedRowsAddTensor; + } // namespace math } // namespace operators } // namespace paddle diff --git a/paddle/operators/math/math_function_test.cc b/paddle/operators/math/math_function_test.cc index fe0d1981fa..33c561f6c6 100644 --- a/paddle/operators/math/math_function_test.cc +++ b/paddle/operators/math/math_function_test.cc @@ -1,185 +1,6 @@ #include "paddle/operators/math/math_function.h" #include "gtest/gtest.h" -#ifdef PADDLE_WITH_CUDA -TEST(math_function, notrans_mul_trans) { - paddle::framework::Tensor input1; - paddle::framework::Tensor input1_gpu; - paddle::framework::Tensor input2_gpu; - paddle::framework::Tensor out_gpu; - paddle::framework::Tensor out; - - auto* cpu_place = new paddle::platform::CPUPlace(); - float* input1_ptr = input1.mutable_data({2, 3}, *cpu_place); - float arr[6] = {0, 1, 2, 3, 4, 5}; - memcpy(input1_ptr, arr, 6 * sizeof(float)); - - auto* gpu_place = new paddle::platform::GPUPlace(0); - paddle::platform::CUDADeviceContext context(*gpu_place); - - input1_gpu.CopyFrom(input1, *gpu_place, context); - input2_gpu.CopyFrom(input1, *gpu_place, context); - - out_gpu.mutable_data({2, 2}, *gpu_place); - - paddle::operators::math::matmul( - context, input1_gpu, false, input2_gpu, true, 1, &out_gpu, 0); - - out.CopyFrom(out_gpu, *cpu_place, context); - - float* out_ptr = out.data(); - context.Wait(); - EXPECT_EQ(out_ptr[0], 5); - EXPECT_EQ(out_ptr[1], 14); - EXPECT_EQ(out_ptr[2], 14); - EXPECT_EQ(out_ptr[3], 50); - delete gpu_place; -} - -TEST(math_function, trans_mul_notrans) { - paddle::framework::Tensor input1; - paddle::framework::Tensor input1_gpu; - paddle::framework::Tensor input2_gpu; - paddle::framework::Tensor out_gpu; - paddle::framework::Tensor out; - - auto* cpu_place = new paddle::platform::CPUPlace(); - float* input1_ptr = input1.mutable_data({2, 3}, *cpu_place); - float arr[6] = {0, 1, 2, 3, 4, 5}; - memcpy(input1_ptr, arr, 6 * sizeof(float)); - - auto* gpu_place = new paddle::platform::GPUPlace(0); - paddle::platform::CUDADeviceContext context(*gpu_place); - - input1_gpu.CopyFrom(input1, *gpu_place, context); - input2_gpu.CopyFrom(input1, *gpu_place, context); - - out_gpu.mutable_data({3, 3}, *gpu_place); - - paddle::operators::math::matmul( - context, input1_gpu, true, input2_gpu, false, 1, &out_gpu, 0); - - out.CopyFrom(out_gpu, *cpu_place, context); - - float* out_ptr = out.data(); - context.Wait(); - EXPECT_EQ(out_ptr[0], 9); - EXPECT_EQ(out_ptr[1], 12); - EXPECT_EQ(out_ptr[2], 15); - EXPECT_EQ(out_ptr[3], 12); - EXPECT_EQ(out_ptr[4], 17); - EXPECT_EQ(out_ptr[5], 22); - EXPECT_EQ(out_ptr[6], 15); - EXPECT_EQ(out_ptr[7], 22); - EXPECT_EQ(out_ptr[8], 29); - delete gpu_place; -} - -TEST(math_function, gemm_notrans_cublas) { - paddle::framework::Tensor input1; - paddle::framework::Tensor input2; - paddle::framework::Tensor input3; - paddle::framework::Tensor input1_gpu; - paddle::framework::Tensor input2_gpu; - paddle::framework::Tensor input3_gpu; - - int m = 2; - int n = 3; - int k = 3; - auto* cpu_place = new paddle::platform::CPUPlace(); - float* input1_ptr = input1.mutable_data({2, 3}, *cpu_place); - float arr1[6] = {0, 1, 2, 3, 4, 5}; - memcpy(input1_ptr, arr1, 6 * sizeof(float)); - float* input2_ptr = input2.mutable_data({3, 4}, *cpu_place); - float arr2[12] = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11}; - memcpy(input2_ptr, arr2, 12 * sizeof(float)); - float* input3_ptr = input3.mutable_data({2, 4}, *cpu_place); - float arr3[8] = {0, 1, 2, 3, 4, 5, 6, 7}; - memcpy(input3_ptr, arr3, 8 * sizeof(float)); - - auto* gpu_place = new paddle::platform::GPUPlace(0); - paddle::platform::CUDADeviceContext context(*gpu_place); - - input1_gpu.CopyFrom(input1, *gpu_place, context); - input2_gpu.CopyFrom(input2, *gpu_place, context); - input3_gpu.CopyFrom(input3, *gpu_place, context); - float* a = input1_gpu.data(); - float* b = input2_gpu.data(); - float* c = input3_gpu.mutable_data(*gpu_place); - - paddle::operators::math::gemm( - context, false, false, m, n, k, 1, a, 3, b + 1, 4, 1, c + 1, 4); - - input3.CopyFrom(input3_gpu, *cpu_place, context); - - // numpy code: - // a = np.arange(6).reshape(2, 3) - // b = np.arange(12).reshape(3, 4)[:, 1:] - // c = np.arange(8).reshape(2, 4)[:, 1:] - // out = np.arange(8).reshape(2, 4) - // out[:, 1:] = np.dot(a, b) + c - context.Wait(); - EXPECT_EQ(input3_ptr[0], 0); - EXPECT_EQ(input3_ptr[1], 24); - EXPECT_EQ(input3_ptr[2], 28); - EXPECT_EQ(input3_ptr[3], 32); - EXPECT_EQ(input3_ptr[4], 4); - EXPECT_EQ(input3_ptr[5], 73); - EXPECT_EQ(input3_ptr[6], 86); - EXPECT_EQ(input3_ptr[7], 99); - delete gpu_place; -} - -TEST(math_function, gemm_trans_cublas) { - paddle::framework::Tensor input1; - paddle::framework::Tensor input2; - paddle::framework::Tensor input3; - paddle::framework::Tensor input1_gpu; - paddle::framework::Tensor input2_gpu; - paddle::framework::Tensor input3_gpu; - - int m = 2; - int n = 3; - int k = 3; - auto* cpu_place = new paddle::platform::CPUPlace(); - float* input1_ptr = input1.mutable_data({2, 3}, *cpu_place); - float arr1[6] = {0, 1, 2, 3, 4, 5}; - memcpy(input1_ptr, arr1, 6 * sizeof(float)); - float* input2_ptr = input2.mutable_data({4, 3}, *cpu_place); - float arr2[12] = {0, 4, 8, 1, 5, 9, 2, 6, 10, 3, 7, 11}; - memcpy(input2_ptr, arr2, 12 * sizeof(float)); - float* input3_ptr = input3.mutable_data({2, 4}, *cpu_place); - float arr3[8] = {0, 1, 2, 3, 4, 5, 6, 7}; - memcpy(input3_ptr, arr3, 8 * sizeof(float)); - - auto* gpu_place = new paddle::platform::GPUPlace(0); - paddle::platform::CUDADeviceContext context(*gpu_place); - - input1_gpu.CopyFrom(input1, *gpu_place, context); - input2_gpu.CopyFrom(input2, *gpu_place, context); - input3_gpu.CopyFrom(input3, *gpu_place, context); - float* a = input1_gpu.data(); - float* b = input2_gpu.data(); - float* c = input3_gpu.mutable_data(*gpu_place); - - paddle::operators::math::gemm( - context, false, true, m, n, k, 1, a, 3, b + 3, 3, 1, c + 1, 4); - - input3.CopyFrom(input3_gpu, *cpu_place, context); - context.Wait(); - - EXPECT_EQ(input3_ptr[0], 0); - EXPECT_EQ(input3_ptr[1], 24); - EXPECT_EQ(input3_ptr[2], 28); - EXPECT_EQ(input3_ptr[3], 32); - EXPECT_EQ(input3_ptr[4], 4); - EXPECT_EQ(input3_ptr[5], 73); - EXPECT_EQ(input3_ptr[6], 86); - EXPECT_EQ(input3_ptr[7], 99); - delete gpu_place; -} -#endif - TEST(math_function, gemm_notrans_cblas) { paddle::framework::Tensor input1; paddle::framework::Tensor input2; diff --git a/paddle/operators/math/math_function_test.cu b/paddle/operators/math/math_function_test.cu new file mode 100644 index 0000000000..e691078bb6 --- /dev/null +++ b/paddle/operators/math/math_function_test.cu @@ -0,0 +1,277 @@ +#include "gtest/gtest.h" +#include "paddle/operators/math/math_function.h" + +TEST(math_function, notrans_mul_trans) { + paddle::framework::Tensor input1; + paddle::framework::Tensor input1_gpu; + paddle::framework::Tensor input2_gpu; + paddle::framework::Tensor out_gpu; + paddle::framework::Tensor out; + + auto* cpu_place = new paddle::platform::CPUPlace(); + float* input1_ptr = input1.mutable_data({2, 3}, *cpu_place); + float arr[6] = {0, 1, 2, 3, 4, 5}; + memcpy(input1_ptr, arr, 6 * sizeof(float)); + + auto* gpu_place = new paddle::platform::GPUPlace(0); + paddle::platform::CUDADeviceContext context(*gpu_place); + + input1_gpu.CopyFrom(input1, *gpu_place, context); + input2_gpu.CopyFrom(input1, *gpu_place, context); + + out_gpu.mutable_data({2, 2}, *gpu_place); + + paddle::operators::math::matmul( + context, input1_gpu, false, input2_gpu, true, 1, &out_gpu, 0); + + out.CopyFrom(out_gpu, *cpu_place, context); + + float* out_ptr = out.data(); + context.Wait(); + EXPECT_EQ(out_ptr[0], 5); + EXPECT_EQ(out_ptr[1], 14); + EXPECT_EQ(out_ptr[2], 14); + EXPECT_EQ(out_ptr[3], 50); + delete gpu_place; +} + +TEST(math_function, trans_mul_notrans) { + paddle::framework::Tensor input1; + paddle::framework::Tensor input1_gpu; + paddle::framework::Tensor input2_gpu; + paddle::framework::Tensor out_gpu; + paddle::framework::Tensor out; + + auto* cpu_place = new paddle::platform::CPUPlace(); + float* input1_ptr = input1.mutable_data({2, 3}, *cpu_place); + float arr[6] = {0, 1, 2, 3, 4, 5}; + memcpy(input1_ptr, arr, 6 * sizeof(float)); + + auto* gpu_place = new paddle::platform::GPUPlace(0); + paddle::platform::CUDADeviceContext context(*gpu_place); + + input1_gpu.CopyFrom(input1, *gpu_place, context); + input2_gpu.CopyFrom(input1, *gpu_place, context); + + out_gpu.mutable_data({3, 3}, *gpu_place); + + paddle::operators::math::matmul( + context, input1_gpu, true, input2_gpu, false, 1, &out_gpu, 0); + + out.CopyFrom(out_gpu, *cpu_place, context); + + float* out_ptr = out.data(); + context.Wait(); + EXPECT_EQ(out_ptr[0], 9); + EXPECT_EQ(out_ptr[1], 12); + EXPECT_EQ(out_ptr[2], 15); + EXPECT_EQ(out_ptr[3], 12); + EXPECT_EQ(out_ptr[4], 17); + EXPECT_EQ(out_ptr[5], 22); + EXPECT_EQ(out_ptr[6], 15); + EXPECT_EQ(out_ptr[7], 22); + EXPECT_EQ(out_ptr[8], 29); + delete gpu_place; +} + +TEST(math_function, gemm_notrans_cublas) { + paddle::framework::Tensor input1; + paddle::framework::Tensor input2; + paddle::framework::Tensor input3; + paddle::framework::Tensor input1_gpu; + paddle::framework::Tensor input2_gpu; + paddle::framework::Tensor input3_gpu; + + int m = 2; + int n = 3; + int k = 3; + auto* cpu_place = new paddle::platform::CPUPlace(); + float* input1_ptr = input1.mutable_data({2, 3}, *cpu_place); + float arr1[6] = {0, 1, 2, 3, 4, 5}; + memcpy(input1_ptr, arr1, 6 * sizeof(float)); + float* input2_ptr = input2.mutable_data({3, 4}, *cpu_place); + float arr2[12] = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11}; + memcpy(input2_ptr, arr2, 12 * sizeof(float)); + float* input3_ptr = input3.mutable_data({2, 4}, *cpu_place); + float arr3[8] = {0, 1, 2, 3, 4, 5, 6, 7}; + memcpy(input3_ptr, arr3, 8 * sizeof(float)); + + auto* gpu_place = new paddle::platform::GPUPlace(0); + paddle::platform::CUDADeviceContext context(*gpu_place); + + input1_gpu.CopyFrom(input1, *gpu_place, context); + input2_gpu.CopyFrom(input2, *gpu_place, context); + input3_gpu.CopyFrom(input3, *gpu_place, context); + float* a = input1_gpu.data(); + float* b = input2_gpu.data(); + float* c = input3_gpu.mutable_data(*gpu_place); + + paddle::operators::math::gemm( + context, false, false, m, n, k, 1, a, 3, b + 1, 4, 1, c + 1, 4); + + input3.CopyFrom(input3_gpu, *cpu_place, context); + + // numpy code: + // a = np.arange(6).reshape(2, 3) + // b = np.arange(12).reshape(3, 4)[:, 1:] + // c = np.arange(8).reshape(2, 4)[:, 1:] + // out = np.arange(8).reshape(2, 4) + // out[:, 1:] = np.dot(a, b) + c + context.Wait(); + EXPECT_EQ(input3_ptr[0], 0); + EXPECT_EQ(input3_ptr[1], 24); + EXPECT_EQ(input3_ptr[2], 28); + EXPECT_EQ(input3_ptr[3], 32); + EXPECT_EQ(input3_ptr[4], 4); + EXPECT_EQ(input3_ptr[5], 73); + EXPECT_EQ(input3_ptr[6], 86); + EXPECT_EQ(input3_ptr[7], 99); + delete gpu_place; +} + +TEST(math_function, gemm_trans_cublas) { + paddle::framework::Tensor input1; + paddle::framework::Tensor input2; + paddle::framework::Tensor input3; + paddle::framework::Tensor input1_gpu; + paddle::framework::Tensor input2_gpu; + paddle::framework::Tensor input3_gpu; + + int m = 2; + int n = 3; + int k = 3; + auto* cpu_place = new paddle::platform::CPUPlace(); + float* input1_ptr = input1.mutable_data({2, 3}, *cpu_place); + float arr1[6] = {0, 1, 2, 3, 4, 5}; + memcpy(input1_ptr, arr1, 6 * sizeof(float)); + float* input2_ptr = input2.mutable_data({4, 3}, *cpu_place); + float arr2[12] = {0, 4, 8, 1, 5, 9, 2, 6, 10, 3, 7, 11}; + memcpy(input2_ptr, arr2, 12 * sizeof(float)); + float* input3_ptr = input3.mutable_data({2, 4}, *cpu_place); + float arr3[8] = {0, 1, 2, 3, 4, 5, 6, 7}; + memcpy(input3_ptr, arr3, 8 * sizeof(float)); + + auto* gpu_place = new paddle::platform::GPUPlace(0); + paddle::platform::CUDADeviceContext context(*gpu_place); + + input1_gpu.CopyFrom(input1, *gpu_place, context); + input2_gpu.CopyFrom(input2, *gpu_place, context); + input3_gpu.CopyFrom(input3, *gpu_place, context); + float* a = input1_gpu.data(); + float* b = input2_gpu.data(); + float* c = input3_gpu.mutable_data(*gpu_place); + + paddle::operators::math::gemm( + context, false, true, m, n, k, 1, a, 3, b + 3, 3, 1, c + 1, 4); + + input3.CopyFrom(input3_gpu, *cpu_place, context); + context.Wait(); + + EXPECT_EQ(input3_ptr[0], 0); + EXPECT_EQ(input3_ptr[1], 24); + EXPECT_EQ(input3_ptr[2], 28); + EXPECT_EQ(input3_ptr[3], 32); + EXPECT_EQ(input3_ptr[4], 4); + EXPECT_EQ(input3_ptr[5], 73); + EXPECT_EQ(input3_ptr[6], 86); + EXPECT_EQ(input3_ptr[7], 99); + delete gpu_place; +} + +TEST(math_function, selected_rows_add) { + using namespace paddle::framework; + using namespace paddle::platform; + using namespace paddle::operators::math; + + CPUPlace gpu_place(0); + CUDADeviceContext ctx(gpu_place); + SetConstant functor; + int64_t height = 10; + int64_t row_numel = 10; + + Vector rows1{0, 4, 7}; + std::unique_ptr selected_rows1{new SelectedRows(rows1, height)}; + auto* in1_value = selected_rows1->mutable_value(); + in1_value->mutable_data( + make_ddim({static_cast(rows1.size()), row_numel}), gpu_place); + functor(ctx, in1_value, 1.0); + + Vector rows2{0, 5, 7, 9}; + std::unique_ptr selected_rows2{new SelectedRows(rows2, height)}; + auto* in2_value = selected_rows2->mutable_value(); + in2_value->mutable_data( + make_ddim({static_cast(rows2.size()), row_numel}), gpu_place); + functor(ctx, in2_value, 2.0); + + std::unique_ptr output{new SelectedRows()}; + auto* out_value = output->mutable_value(); + + // simplely concat two SelectedRows + out_value->mutable_data(make_ddim({7, 10}), gpu_place); + + SelectedRowsAdd add_functor; + add_functor(ctx, *selected_rows1, *selected_rows2, output.get()); + + auto out_height = output->height(); + EXPECT_EQ(out_height, height); + + auto& out_rows = output->rows(); + + // input1 rows + EXPECT_EQ(out_rows[0], 0); + EXPECT_EQ(out_rows[1], 4); + EXPECT_EQ(out_rows[2], 7); + // input2 rows + EXPECT_EQ(out_rows[3], 0); + EXPECT_EQ(out_rows[4], 5); + EXPECT_EQ(out_rows[5], 7); + EXPECT_EQ(out_rows[6], 9); + + Tensor out_cpu; + out_cpu.CopyFrom(*out_value, platform::CPUPlace(), ctx); + ctx.Wait(); + + auto* out_cpu_data = out_cpu.data(); + // input1 value + EXPECT_EQ(out_cpu_data[0 * row_numel + 0], 1.0); + EXPECT_EQ(out_cpu_data[0 * row_numel + 8], 1.0); + EXPECT_EQ(out_cpu_data[1 * row_numel + 1], 1.0); + EXPECT_EQ(out_cpu_data[2 * row_numel + 6], 1.0); + // input2 value + EXPECT_EQ(out_cpu_data[3 * row_numel + 3], 2.0); + EXPECT_EQ(out_cpu_data[3 * row_numel + 8], 2.0); + EXPECT_EQ(out_cpu_data[4 * row_numel + 4], 2.0); + EXPECT_EQ(out_cpu_data[5 * row_numel + 7], 2.0); + EXPECT_EQ(out_cpu_data[6 * row_numel + 9], 2.0); + + std::unique_ptr tensor1{new Tensor()}; + tensor1->mutable_data(make_ddim({height, row_numel}), gpu_place); + SetConstant constant_functor; + constant_functor(ctx, tensor1.get(), 3.0); + + std::unique_ptr tensor2{new Tensor()}; + tensor2->mutable_data(make_ddim({height, row_numel}), gpu_place); + + SelectedRowsAddTensor add_tensor_functor; + add_tensor_functor(ctx, *output, *tensor1, tensor2.get()); + + Tensor tensor2_cpu; + tensor2_cpu.CopyFrom(*tensor2, platform::CPUPlace(), ctx); + ctx.Wait(); + + auto* tensor2_cpu_data = tensor2_cpu->data(); + // row0: 1.0 + 2.0 + 3.0 + EXPECT_EQ(tensor2_cpu_data[0 * row_numel + 0], 6.0); + // row1: 3.0 + EXPECT_EQ(tensor2_cpu_data[1 * row_numel + 1], 3.0); + // row4 : 1.0 + 3.0 + EXPECT_EQ(tensor2_cpu_data[4 * row_numel + 6], 4.0); + // row5: 2.0 + 3.0 + EXPECT_EQ(tensor2_cpu_data[5 * row_numel + 7], 5.0); + // row6: 3.0 + EXPECT_EQ(tensor2_cpu_data[6 * row_numel + 1], 3.0); + // row7: 1.0 + 2.0 + 3.0 + EXPECT_EQ(tensor2_cpu_data[7 * row_numel + 3], 6.0); + // row9: 2.0 + 3.0 + EXPECT_EQ(tensor2_cpu_data[9 * row_numel + 6], 5.0); +} From 7ef568e8935768e7ba9ebd4b82bc42b1fcf6ade0 Mon Sep 17 00:00:00 2001 From: qijun Date: Sat, 14 Oct 2017 23:13:04 +0000 Subject: [PATCH 14/36] fix gpu unittest error --- paddle/operators/math/CMakeLists.txt | 4 +-- paddle/operators/math/math_function.cu | 40 ++++++++++++++------- paddle/operators/math/math_function_test.cu | 13 +++---- 3 files changed, 36 insertions(+), 21 deletions(-) diff --git a/paddle/operators/math/CMakeLists.txt b/paddle/operators/math/CMakeLists.txt index a7f275bae5..77a3603eb6 100644 --- a/paddle/operators/math/CMakeLists.txt +++ b/paddle/operators/math/CMakeLists.txt @@ -1,18 +1,18 @@ if(WITH_GPU) nv_library(math_function SRCS math_function.cc math_function.cu im2col.cc im2col.cu DEPS cblas device_context operator) - nv_test(math_function_gpu_test SRCS math_function_test.cu DEPS math_function tensor) + nv_test(math_function_gpu_test SRCS math_function_test.cu DEPS math_function tensor selected_rows) nv_library(softmax SRCS softmax.cc softmax.cu DEPS operator) nv_library(cross_entropy SRCS cross_entropy.cc cross_entropy.cu DEPS operator) nv_library(pooling SRCS pooling.cc pooling.cu DEPS device_context) nv_library(vol2col SRCS vol2col.cc vol2col.cu DEPS device_context) else() cc_library(math_function SRCS math_function.cc im2col.cc DEPS cblas device_context operator) - cc_test(math_function_test SRCS math_function_test.cc DEPS math_function tensor) cc_library(softmax SRCS softmax.cc DEPS operator) cc_library(cross_entropy SRCS cross_entropy.cc DEPS operator) cc_library(pooling SRCS pooling.cc DEPS device_context) cc_library(vol2col SRCS vol2col.cc DEPS device_context) endif() +cc_test(math_function_test SRCS math_function_test.cc DEPS math_function tensor selected_rows) cc_test(im2col_test SRCS im2col_test.cc DEPS math_function tensor) cc_test(vol2col_test SRCS vol2col_test.cc DEPS vol2col tensor) diff --git a/paddle/operators/math/math_function.cu b/paddle/operators/math/math_function.cu index 26bf0ec2f1..d31b223b2c 100644 --- a/paddle/operators/math/math_function.cu +++ b/paddle/operators/math/math_function.cu @@ -13,6 +13,7 @@ See the License for the specific language governing permissions and limitations under the License. */ #include "paddle/operators/math/math_function.h" +#include "paddle/platform/cuda_helper.h" namespace paddle { namespace operators { @@ -191,7 +192,7 @@ struct SelectedRowsAdd { auto in2_place = input2.place(); PADDLE_ENFORCE(platform::is_gpu_place(in2_place)); auto out_place = context.GetPlace(); - PADDLE_ENFORCE(platform::is_gpu_place(out_place)) + PADDLE_ENFORCE(platform::is_gpu_place(out_place)); memory::Copy( boost::get(out_place), out_data, @@ -211,22 +212,26 @@ struct SelectedRowsAdd { template struct SelectedRowsAdd; namespace { -template -__global__ void SelectedRowsAddTensorKernel(T* selected_rows, int64_t* rows, - T* tensor_in, T* tensor_out, - const int64_t row_numel) { - const ty = blockIdx.y; +template +__global__ void SelectedRowsAddTensorKernel(const T* selected_rows, + const int64_t* rows, + T* tensor_out, + int64_t row_numel, + int block_size) { + const int ty = blockIdx.y; int tid = threadIdx.x; selected_rows += ty * row_numel; - tensor_in += rows[ty] * row_numel; tensor_out += rows[ty] * row_numel; for (int index = tid; index < row_numel; index += block_size) { - tensor_out[index] = tensor_in[index] + selected_rows[index]; + // Since index in rows of SelectedRows can be duplicate, we can not use + // tensor_out[index] += selected_rows[index]; Instead, we have to use + // AtomicAdd to avoid concurrent write error. + paddle::platform::CudaAtomicAdd(&tensor_out[index], selected_rows[index]); } } -} +} // namespace template struct SelectedRowsAddTensor { @@ -250,13 +255,22 @@ struct SelectedRowsAddTensor { auto* in2_data = input2.data(); auto* out_data = output->data(); - const int block_size = 256; + SetConstant functor; + functor(context, output, 0.0); + + int block_size = 256; dim3 threads(block_size, 1); dim3 grid(1, in1_height); - SelectedRowsAddTensorKernel<<< + SelectedRowsAddTensorKernel<<< grid, threads, 0, - reinterpret_cast(ctx).stream()>>>( - in1_data, in1_rows.data(), in2_data, out_data, in1_row_numel); + reinterpret_cast(context).stream() + >>>(in1_data, in1_rows.data(), + out_data, in1_row_numel, block_size); + + auto out_eigen = framework::EigenVector::Flatten(*output); + auto in2_eigen = framework::EigenVector::Flatten(input2); + out_eigen.device(*context.GetEigenDevice()) = + out_eigen + in2_eigen; } }; diff --git a/paddle/operators/math/math_function_test.cu b/paddle/operators/math/math_function_test.cu index e691078bb6..1acc5f66a6 100644 --- a/paddle/operators/math/math_function_test.cu +++ b/paddle/operators/math/math_function_test.cu @@ -183,20 +183,21 @@ TEST(math_function, selected_rows_add) { using namespace paddle::platform; using namespace paddle::operators::math; - CPUPlace gpu_place(0); + GPUPlace gpu_place(0); + CPUPlace cpu_place; CUDADeviceContext ctx(gpu_place); SetConstant functor; int64_t height = 10; int64_t row_numel = 10; - Vector rows1{0, 4, 7}; + std::vector rows1{0, 4, 7}; std::unique_ptr selected_rows1{new SelectedRows(rows1, height)}; auto* in1_value = selected_rows1->mutable_value(); in1_value->mutable_data( make_ddim({static_cast(rows1.size()), row_numel}), gpu_place); functor(ctx, in1_value, 1.0); - Vector rows2{0, 5, 7, 9}; + std::vector rows2{0, 5, 7, 9}; std::unique_ptr selected_rows2{new SelectedRows(rows2, height)}; auto* in2_value = selected_rows2->mutable_value(); in2_value->mutable_data( @@ -228,7 +229,7 @@ TEST(math_function, selected_rows_add) { EXPECT_EQ(out_rows[6], 9); Tensor out_cpu; - out_cpu.CopyFrom(*out_value, platform::CPUPlace(), ctx); + out_cpu.CopyFrom(*out_value, cpu_place, ctx); ctx.Wait(); auto* out_cpu_data = out_cpu.data(); @@ -256,10 +257,10 @@ TEST(math_function, selected_rows_add) { add_tensor_functor(ctx, *output, *tensor1, tensor2.get()); Tensor tensor2_cpu; - tensor2_cpu.CopyFrom(*tensor2, platform::CPUPlace(), ctx); + tensor2_cpu.CopyFrom(*tensor2, cpu_place, ctx); ctx.Wait(); - auto* tensor2_cpu_data = tensor2_cpu->data(); + auto* tensor2_cpu_data = tensor2_cpu.data(); // row0: 1.0 + 2.0 + 3.0 EXPECT_EQ(tensor2_cpu_data[0 * row_numel + 0], 6.0); // row1: 3.0 From df2d1769fd530fa3a57b92e50819d341768a7e80 Mon Sep 17 00:00:00 2001 From: qijun Date: Sat, 14 Oct 2017 16:21:26 -0700 Subject: [PATCH 15/36] fix code style --- paddle/framework/lod_tensor.h | 3 +++ paddle/framework/type_defs.h | 3 +++ paddle/operators/math/math_function.cu | 16 +++++++--------- 3 files changed, 13 insertions(+), 9 deletions(-) diff --git a/paddle/framework/lod_tensor.h b/paddle/framework/lod_tensor.h index ee040a9144..4db36ee766 100644 --- a/paddle/framework/lod_tensor.h +++ b/paddle/framework/lod_tensor.h @@ -1,8 +1,11 @@ /* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + Licensed under the Apache License, Version 2.0 (the "License"); you may not use this file except in compliance with the License. You may obtain a copy of the License at + http://www.apache.org/licenses/LICENSE-2.0 + Unless required by applicable law or agreed to in writing, software distributed under the License is distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. diff --git a/paddle/framework/type_defs.h b/paddle/framework/type_defs.h index 0c0a72de31..0d1564a751 100644 --- a/paddle/framework/type_defs.h +++ b/paddle/framework/type_defs.h @@ -1,8 +1,11 @@ /* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + Licensed under the Apache License, Version 2.0 (the "License"); you may not use this file except in compliance with the License. You may obtain a copy of the License at + http://www.apache.org/licenses/LICENSE-2.0 + Unless required by applicable law or agreed to in writing, software distributed under the License is distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. diff --git a/paddle/operators/math/math_function.cu b/paddle/operators/math/math_function.cu index d31b223b2c..fc16d1b0a7 100644 --- a/paddle/operators/math/math_function.cu +++ b/paddle/operators/math/math_function.cu @@ -214,10 +214,8 @@ template struct SelectedRowsAdd; namespace { template __global__ void SelectedRowsAddTensorKernel(const T* selected_rows, - const int64_t* rows, - T* tensor_out, - int64_t row_numel, - int block_size) { + const int64_t* rows, T* tensor_out, + int64_t row_numel, int block_size) { const int ty = blockIdx.y; int tid = threadIdx.x; @@ -261,11 +259,11 @@ struct SelectedRowsAddTensor { int block_size = 256; dim3 threads(block_size, 1); dim3 grid(1, in1_height); - SelectedRowsAddTensorKernel<<< - grid, threads, 0, - reinterpret_cast(context).stream() - >>>(in1_data, in1_rows.data(), - out_data, in1_row_numel, block_size); + SelectedRowsAddTensorKernel< + T><<(context) + .stream()>>>(in1_data, in1_rows.data(), out_data, + in1_row_numel, block_size); auto out_eigen = framework::EigenVector::Flatten(*output); auto in2_eigen = framework::EigenVector::Flatten(input2); From 89758adb83cea198ee6d2d31bc8e1d9bad5e827e Mon Sep 17 00:00:00 2001 From: qijun Date: Sat, 14 Oct 2017 16:21:26 -0700 Subject: [PATCH 16/36] fix code style --- paddle/framework/lod_tensor.h | 3 +++ paddle/framework/selected_rows_test.cc | 2 +- paddle/framework/type_defs.h | 3 +++ paddle/operators/math/math_function.cu | 16 +++++++--------- 4 files changed, 14 insertions(+), 10 deletions(-) diff --git a/paddle/framework/lod_tensor.h b/paddle/framework/lod_tensor.h index ee040a9144..4db36ee766 100644 --- a/paddle/framework/lod_tensor.h +++ b/paddle/framework/lod_tensor.h @@ -1,8 +1,11 @@ /* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + Licensed under the Apache License, Version 2.0 (the "License"); you may not use this file except in compliance with the License. You may obtain a copy of the License at + http://www.apache.org/licenses/LICENSE-2.0 + Unless required by applicable law or agreed to in writing, software distributed under the License is distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. diff --git a/paddle/framework/selected_rows_test.cc b/paddle/framework/selected_rows_test.cc index 055b867600..4ee13a65d7 100644 --- a/paddle/framework/selected_rows_test.cc +++ b/paddle/framework/selected_rows_test.cc @@ -18,7 +18,7 @@ namespace framework { class SelectedRowsTester : public ::testing::Test { public: virtual void SetUp() override { - Vector rows{0, 4, 7}; + std::vector rows{0, 4, 7}; int64_t height = 10; int64_t row_numel = 100; selected_rows_.reset(new SelectedRows(rows, height)); diff --git a/paddle/framework/type_defs.h b/paddle/framework/type_defs.h index 0c0a72de31..0d1564a751 100644 --- a/paddle/framework/type_defs.h +++ b/paddle/framework/type_defs.h @@ -1,8 +1,11 @@ /* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + Licensed under the Apache License, Version 2.0 (the "License"); you may not use this file except in compliance with the License. You may obtain a copy of the License at + http://www.apache.org/licenses/LICENSE-2.0 + Unless required by applicable law or agreed to in writing, software distributed under the License is distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. diff --git a/paddle/operators/math/math_function.cu b/paddle/operators/math/math_function.cu index d31b223b2c..fc16d1b0a7 100644 --- a/paddle/operators/math/math_function.cu +++ b/paddle/operators/math/math_function.cu @@ -214,10 +214,8 @@ template struct SelectedRowsAdd; namespace { template __global__ void SelectedRowsAddTensorKernel(const T* selected_rows, - const int64_t* rows, - T* tensor_out, - int64_t row_numel, - int block_size) { + const int64_t* rows, T* tensor_out, + int64_t row_numel, int block_size) { const int ty = blockIdx.y; int tid = threadIdx.x; @@ -261,11 +259,11 @@ struct SelectedRowsAddTensor { int block_size = 256; dim3 threads(block_size, 1); dim3 grid(1, in1_height); - SelectedRowsAddTensorKernel<<< - grid, threads, 0, - reinterpret_cast(context).stream() - >>>(in1_data, in1_rows.data(), - out_data, in1_row_numel, block_size); + SelectedRowsAddTensorKernel< + T><<(context) + .stream()>>>(in1_data, in1_rows.data(), out_data, + in1_row_numel, block_size); auto out_eigen = framework::EigenVector::Flatten(*output); auto in2_eigen = framework::EigenVector::Flatten(input2); From 7c0facd1958faa2e48d5764b3274459124e51c53 Mon Sep 17 00:00:00 2001 From: qijun Date: Sun, 15 Oct 2017 22:01:51 -0700 Subject: [PATCH 17/36] init --- paddle/framework/scope.h | 62 +++++++++++++++++++ paddle/pybind/pybind.cc | 6 ++ .../framework/tests/test_feed_fetch_method.py | 50 +++++++++++++++ 3 files changed, 118 insertions(+) create mode 100644 python/paddle/v2/framework/tests/test_feed_fetch_method.py diff --git a/paddle/framework/scope.h b/paddle/framework/scope.h index a7fce3514b..7376245b53 100644 --- a/paddle/framework/scope.h +++ b/paddle/framework/scope.h @@ -18,8 +18,10 @@ limitations under the License. */ #include #include +#include "paddle/framework/lod_tensor.h" #include "paddle/framework/variable.h" #include "paddle/platform/macros.h" +#include "paddle/platform/place.h" namespace paddle { namespace framework { @@ -75,5 +77,65 @@ class Scope { framework::Scope& GetGlobalScope(); +// template +// void SetFeedVariable(const std::vector& input, const Lod& lod, +// const std::vector& dims, +// const std::string& var_name, size_t index) { +// Variable* g_feed_value = GetGlobalScope().Var("var_name"); +// // feed variable holds vector +// auto& feed_inputs = +// *(g_feed_value->GetMutable< +// std::vector>()); +// if (index >= feed_inputs.size()) { +// feed_inputs.resize(index); +// } +// // copy tensor +// T* dst = feed_inputs[index].mutable_data(make_ddim(dims), +// platform::CPUPlace()); +// memcpy(dst, inputs[i].data(), inputs[i].size() * sizeof(T)); +// // copy lod +// feed_inputs[index].set_lod(lod); +// } + +template +void SetFeedVariable(const LoDTensor& input, const std::string& var_name, + size_t index) { + std::cout << "into SetFeedVariable" << std::endl; + std::cout << var_name << std::endl; + std::cout << index << std::endl; + Variable* g_feed_value = GetGlobalScope().Var(var_name); + auto& feed_inputs = + *(g_feed_value->GetMutable>()); + if (index >= feed_inputs.size()) { + feed_inputs.resize(index + 1); + } + // shared data with input tensor + feed_inputs[index].ShareDataWith(input); + // set lod + feed_inputs[index].set_lod(input.lod()); +} + +// template +// std::vector GetFetchVariable(const std::string& var_name, size_t index) { +// Variable* g_fetch_value = GetGlobalScope().Var(var_name); +// auto& fetch_outputs = +// *(g_fetch_value->GetMutable< +// std::vector>()); +// std::vector result; +// result.resize(fetch_outputs[index].numel()); +// memcpy(result.data(), fetch_outputs[i].data(), +// fetch_outputs[i].numel() * sizeof(T)); +// } + +template +LoDTensor& GetFetchVariable(const std::string& var_name, size_t index) { + Variable* g_fetch_value = GetGlobalScope().Var(var_name); + auto& fetch_outputs = + *(g_fetch_value->GetMutable>()); + std::cout << "into GetFetchVariable" << std::endl; + PADDLE_ENFORCE_LT(index, fetch_outputs.size()); + return fetch_outputs[index]; +} + } // namespace framework } // namespace paddle diff --git a/paddle/pybind/pybind.cc b/paddle/pybind/pybind.cc index b143cb9f59..a4d14c9303 100644 --- a/paddle/pybind/pybind.cc +++ b/paddle/pybind/pybind.cc @@ -394,6 +394,12 @@ All parameter, weight, gradient are variables in Paddle. m.def("unique_integer", UniqueIntegerGenerator); m.def("is_compile_gpu", IsCompileGPU); + m.def("set_feed_variable", SetFeedVariable); + // m.def("set_feed_variable", SetFeedVariable); + // m.def("set_feed_variable", SetFeedVariable); + m.def("get_fetch_variable", GetFetchVariable); + // m.def("get_fetch_variable", GetFetchVariable); + // m.def("get_fetch_variable", GetFetchVariable); BindProgramDesc(m); BindBlockDesc(m); diff --git a/python/paddle/v2/framework/tests/test_feed_fetch_method.py b/python/paddle/v2/framework/tests/test_feed_fetch_method.py new file mode 100644 index 0000000000..cd8eeee68f --- /dev/null +++ b/python/paddle/v2/framework/tests/test_feed_fetch_method.py @@ -0,0 +1,50 @@ +import paddle.v2.framework.core as core +import unittest +import numpy as np + +# class TestFeedFetch(unittest.TestCase): +# def test_feed_fetch(self): +# place = core.CPUPlace() +# input_tensor = core.LoDTensor([[0, 2, 4]]) +# input_tensor.set_dims([4, 4, 6]) +# input_tensor.alloc_int(place) +# input_array = np.array(input_tensor) +# input_array[0, 0, 0] = 3 +# input_array[3, 3, 5] = 10 +# input_tensor.set(input_array, place) + +# core.set_feed_variable(input_tensor, "feed", 0) + +# output_tensor = core.get_fetch_variable("feed", 0) +# print type(output_tensor) + +# output_lod = output_tensor.lod() +# print type(output_lod) +# print output_lod[0] +# print output_lod[0][0] +# print output_lod[0][1] +# print output_lod[0][2] +# # self.assertEqual(0, output_lod[0][0]) +# # self.assertEqual(0, output_lod[0][0]) +# # self.assertEqual(2, output_lod[0][1]) +# # self.assertEqual(4, output_lod[0][2]) + +# # output_array = np.array(output_tensor) +# # self.assertEqual(3, output_array[0, 0, 0]) +# # self.assertEqual(10, output_array[3, 3, 5]); + + +class TestFeedFetch(unittest.TestCase): + def test_feed_fetch(self): + place = core.CPUPlace() + input_tensor = core.LoDTensor([[0, 2, 4]]) + input_tensor.set_dims([4, 4, 6]) + input_tensor.alloc_float(place) + input_array = np.array(input_tensor) + input_array[0, 0, 0] = 3 + input_array[3, 3, 5] = 10 + input_tensor.set(input_array, place) + + +if __name__ == "__main__": + unittest.main() From ab5dc9fe185215760ee351e36d772ad245e1eb3d Mon Sep 17 00:00:00 2001 From: qijun Date: Sun, 15 Oct 2017 22:45:52 -0700 Subject: [PATCH 18/36] remove SelectedRows functors to selected_rows_functor.h --- paddle/operators/math/CMakeLists.txt | 8 +- paddle/operators/math/math_function.cc | 93 ------------ paddle/operators/math/math_function.cu | 119 +-------------- paddle/operators/math/math_function.h | 18 --- paddle/operators/math/math_function_test.cc | 90 ----------- paddle/operators/math/math_function_test.cu | 99 ------------ .../operators/math/selected_rows_functor.cc | 114 ++++++++++++++ .../operators/math/selected_rows_functor.cu | 142 ++++++++++++++++++ paddle/operators/math/selected_rows_functor.h | 41 +++++ .../math/selected_rows_functor_test.cc | 106 +++++++++++++ .../math/selected_rows_functor_test.cu | 115 ++++++++++++++ 11 files changed, 525 insertions(+), 420 deletions(-) create mode 100644 paddle/operators/math/selected_rows_functor.cc create mode 100644 paddle/operators/math/selected_rows_functor.cu create mode 100644 paddle/operators/math/selected_rows_functor.h create mode 100644 paddle/operators/math/selected_rows_functor_test.cc create mode 100644 paddle/operators/math/selected_rows_functor_test.cu diff --git a/paddle/operators/math/CMakeLists.txt b/paddle/operators/math/CMakeLists.txt index 77a3603eb6..72ce858504 100644 --- a/paddle/operators/math/CMakeLists.txt +++ b/paddle/operators/math/CMakeLists.txt @@ -1,18 +1,22 @@ if(WITH_GPU) nv_library(math_function SRCS math_function.cc math_function.cu im2col.cc im2col.cu DEPS cblas device_context operator) - nv_test(math_function_gpu_test SRCS math_function_test.cu DEPS math_function tensor selected_rows) + nv_test(math_function_gpu_test SRCS math_function_test.cu DEPS math_function tensor) + nv_library(selected_rows_functor SRCS selected_rows_functor.cc selected_rows_functor.cu DEPS selected_rows math_function) + nv_test(selected_rows_functor_gpu_test SRCS selected_rows_functor_test.cu DEPS selected_rows_functor) nv_library(softmax SRCS softmax.cc softmax.cu DEPS operator) nv_library(cross_entropy SRCS cross_entropy.cc cross_entropy.cu DEPS operator) nv_library(pooling SRCS pooling.cc pooling.cu DEPS device_context) nv_library(vol2col SRCS vol2col.cc vol2col.cu DEPS device_context) else() cc_library(math_function SRCS math_function.cc im2col.cc DEPS cblas device_context operator) + cc_library(selected_rows_functor SRCS selected_rows_functor.cc DEPS selected_rows math_function) cc_library(softmax SRCS softmax.cc DEPS operator) cc_library(cross_entropy SRCS cross_entropy.cc DEPS operator) cc_library(pooling SRCS pooling.cc DEPS device_context) cc_library(vol2col SRCS vol2col.cc DEPS device_context) endif() -cc_test(math_function_test SRCS math_function_test.cc DEPS math_function tensor selected_rows) +cc_test(math_function_test SRCS math_function_test.cc DEPS math_function tensor) +cc_test(selected_rows_functor_test SRCS selected_rows_functor_test.cc DEPS selected_rows_functor) cc_test(im2col_test SRCS im2col_test.cc DEPS math_function tensor) cc_test(vol2col_test SRCS vol2col_test.cc DEPS vol2col tensor) diff --git a/paddle/operators/math/math_function.cc b/paddle/operators/math/math_function.cc index a1faafb7c4..77a1e22b41 100644 --- a/paddle/operators/math/math_function.cc +++ b/paddle/operators/math/math_function.cc @@ -13,8 +13,6 @@ See the License for the specific language governing permissions and limitations under the License. */ #include "paddle/operators/math/math_function.h" -#include "paddle/framework/eigen.h" -#include "paddle/memory/memcpy.h" namespace paddle { namespace operators { @@ -134,97 +132,6 @@ void matmul( template struct SetConstant; -template -struct SelectedRowsAdd { - void operator()(const platform::DeviceContext& context, - const framework::SelectedRows& input1, - const framework::SelectedRows& input2, - framework::SelectedRows* output) { - auto in1_height = input1.height(); - PADDLE_ENFORCE_EQ(in1_height, input2.height()); - output->set_height(in1_height); - - auto& in1_rows = input1.rows(); - auto& in2_rows = input2.rows(); - std::vector out_rows; - out_rows.reserve(in1_rows.size() + in2_rows.size()); - - // concat rows - out_rows.insert(out_rows.end(), in1_rows.begin(), in1_rows.end()); - out_rows.insert(out_rows.end(), in2_rows.begin(), in2_rows.end()); - output->set_rows(out_rows); - - auto* out_value = output->mutable_value(); - auto& in1_value = input1.value(); - auto& in2_value = input2.value(); - - auto in1_row_numel = in1_value.numel() / in1_rows.size(); - PADDLE_ENFORCE_EQ(in1_row_numel, in2_value.numel() / in2_rows.size()); - PADDLE_ENFORCE_EQ(in1_row_numel, out_value->numel() / out_rows.size()); - - auto in1_place = input1.place(); - PADDLE_ENFORCE(platform::is_cpu_place(in1_place)); - auto in2_place = input2.place(); - PADDLE_ENFORCE(platform::is_cpu_place(in2_place)); - auto out_place = context.GetPlace(); - PADDLE_ENFORCE(platform::is_cpu_place(out_place)); - - auto* out_data = out_value->data(); - auto* in1_data = in1_value.data(); - memory::Copy(boost::get(out_place), out_data, - boost::get(in1_place), in1_data, - in1_value.numel() * sizeof(T)); - - auto* in2_data = in2_value.data(); - memory::Copy(boost::get(out_place), - out_data + in1_value.numel(), - boost::get(in2_place), in2_data, - in2_value.numel() * sizeof(T)); - } -}; - -template struct SelectedRowsAdd; - -template -struct SelectedRowsAddTensor { - void operator()(const platform::DeviceContext& context, - const framework::SelectedRows& input1, - const framework::Tensor& input2, framework::Tensor* output) { - auto in1_height = input1.height(); - auto in2_dims = input2.dims(); - auto out_dims = output->dims(); - PADDLE_ENFORCE_EQ(in1_height, in2_dims[0]); - PADDLE_ENFORCE_EQ(in1_height, out_dims[0]); - - auto& in1_value = input1.value(); - auto& in1_rows = input1.rows(); - - int64_t in1_row_numel = in1_value.numel() / in1_rows.size(); - PADDLE_ENFORCE_EQ(in1_row_numel, input2.numel() / in1_height); - PADDLE_ENFORCE_EQ(in1_row_numel, output->numel() / in1_height); - - SetConstant functor; - functor(context, output, 0.0); - - auto* in1_data = in1_value.data(); - auto* out_data = output->data(); - - for (size_t i = 0; i < in1_rows.size(); i++) { - for (int64_t j = 0; j < in1_row_numel; j++) { - out_data[in1_rows[i] * in1_row_numel + j] += - in1_data[i * in1_row_numel + j]; - } - } - - auto out_eigen = framework::EigenVector::Flatten(*output); - auto in2_eigen = framework::EigenVector::Flatten(input2); - out_eigen.device(*context.GetEigenDevice()) = - out_eigen + in2_eigen; - } -}; - -template struct SelectedRowsAddTensor; - } // namespace math } // namespace operators } // namespace paddle diff --git a/paddle/operators/math/math_function.cu b/paddle/operators/math/math_function.cu index fc16d1b0a7..7fbc03acf2 100644 --- a/paddle/operators/math/math_function.cu +++ b/paddle/operators/math/math_function.cu @@ -13,7 +13,6 @@ See the License for the specific language governing permissions and limitations under the License. */ #include "paddle/operators/math/math_function.h" -#include "paddle/platform/cuda_helper.h" namespace paddle { namespace operators { @@ -156,123 +155,7 @@ void matmul( matrix_b.data(), beta, matrix_out->data()); } -template -struct SelectedRowsAdd { - void operator()(const platform::DeviceContext& context, - const framework::SelectedRows& input1, - const framework::SelectedRows& input2, - framework::SelectedRows* output) { - auto in1_height = input1.height(); - PADDLE_ENFORCE_EQ(in1_height, input2.height()); - output->set_height(in1_height); - - auto& in1_rows = input1.rows(); - auto& in2_rows = input2.rows(); - std::vector out_rows; - out_rows.reserve(in1_rows.size() + in2_rows.size()); - - // concat rows - out_rows.insert(out_rows.end(), in1_rows.begin(), in1_rows.end()); - out_rows.insert(out_rows.end(), in2_rows.begin(), in2_rows.end()); - output->set_rows(out_rows); - - auto* out_value = output->mutable_value(); - auto& in1_value = input1.value(); - auto& in2_value = input2.value(); - - auto in1_row_numel = in1_value.numel() / in1_rows.size(); - PADDLE_ENFORCE_EQ(in1_row_numel, in2_value.numel() / in2_rows.size()); - PADDLE_ENFORCE_EQ(in1_row_numel, out_value->numel() / out_rows.size()); - - auto* out_data = out_value->data(); - auto* in1_data = in1_value.data(); - - auto in1_place = input1.place(); - PADDLE_ENFORCE(platform::is_gpu_place(in1_place)); - auto in2_place = input2.place(); - PADDLE_ENFORCE(platform::is_gpu_place(in2_place)); - auto out_place = context.GetPlace(); - PADDLE_ENFORCE(platform::is_gpu_place(out_place)); - - memory::Copy( - boost::get(out_place), out_data, - boost::get(in1_place), in1_data, - in1_value.numel() * sizeof(T), - reinterpret_cast(context).stream()); - - auto* in2_data = in2_value.data(); - memory::Copy( - boost::get(out_place), out_data + in1_value.numel(), - boost::get(in2_place), in2_data, - in2_value.numel() * sizeof(T), - reinterpret_cast(context).stream()); - } -}; - -template struct SelectedRowsAdd; - -namespace { -template -__global__ void SelectedRowsAddTensorKernel(const T* selected_rows, - const int64_t* rows, T* tensor_out, - int64_t row_numel, int block_size) { - const int ty = blockIdx.y; - int tid = threadIdx.x; - - selected_rows += ty * row_numel; - tensor_out += rows[ty] * row_numel; - - for (int index = tid; index < row_numel; index += block_size) { - // Since index in rows of SelectedRows can be duplicate, we can not use - // tensor_out[index] += selected_rows[index]; Instead, we have to use - // AtomicAdd to avoid concurrent write error. - paddle::platform::CudaAtomicAdd(&tensor_out[index], selected_rows[index]); - } -} -} // namespace - -template -struct SelectedRowsAddTensor { - void operator()(const platform::DeviceContext& context, - const framework::SelectedRows& input1, - const framework::Tensor& input2, framework::Tensor* output) { - auto in1_height = input1.height(); - auto in2_dims = input2.dims(); - auto out_dims = output->dims(); - PADDLE_ENFORCE_EQ(in1_height, in2_dims[0]); - PADDLE_ENFORCE_EQ(in1_height, out_dims[0]); - - auto& in1_value = input1.value(); - auto& in1_rows = input1.rows(); - - int64_t in1_row_numel = in1_value.numel() / in1_rows.size(); - PADDLE_ENFORCE_EQ(in1_row_numel, input2.numel() / in1_height); - PADDLE_ENFORCE_EQ(in1_row_numel, output->numel() / in1_height); - - auto* in1_data = in1_value.data(); - auto* in2_data = input2.data(); - auto* out_data = output->data(); - - SetConstant functor; - functor(context, output, 0.0); - - int block_size = 256; - dim3 threads(block_size, 1); - dim3 grid(1, in1_height); - SelectedRowsAddTensorKernel< - T><<(context) - .stream()>>>(in1_data, in1_rows.data(), out_data, - in1_row_numel, block_size); - - auto out_eigen = framework::EigenVector::Flatten(*output); - auto in2_eigen = framework::EigenVector::Flatten(input2); - out_eigen.device(*context.GetEigenDevice()) = - out_eigen + in2_eigen; - } -}; - -template struct SelectedRowsAddTensor; +template struct SetConstant; } // namespace math } // namespace operators diff --git a/paddle/operators/math/math_function.h b/paddle/operators/math/math_function.h index 0d0d4cdd73..6f92d83aab 100644 --- a/paddle/operators/math/math_function.h +++ b/paddle/operators/math/math_function.h @@ -53,7 +53,6 @@ int LAPACKE_dgetri(int matrix_layout, int n, double* a, int lda, #include #include "paddle/framework/eigen.h" -#include "paddle/framework/selected_rows.h" #include "paddle/framework/tensor.h" #include "paddle/platform/device_context.h" #include "paddle/platform/enforce.h" @@ -96,23 +95,6 @@ struct SetConstant { } }; -// SelectedRows + SelectedRows will simplely concat value and rows. -// The real computation happens in dealing with LoDTensor. -template -struct SelectedRowsAdd { - void operator()(const platform::DeviceContext& context, - const framework::SelectedRows& input1, - const framework::SelectedRows& input2, - framework::SelectedRows* output); -}; - -template -struct SelectedRowsAddTensor { - void operator()(const platform::DeviceContext& context, - const framework::SelectedRows& input1, - const framework::Tensor& input2, framework::Tensor* output); -}; - } // namespace math } // namespace operators } // namespace paddle diff --git a/paddle/operators/math/math_function_test.cc b/paddle/operators/math/math_function_test.cc index 33c561f6c6..3b9f92e7ae 100644 --- a/paddle/operators/math/math_function_test.cc +++ b/paddle/operators/math/math_function_test.cc @@ -89,93 +89,3 @@ TEST(math_function, zero) { EXPECT_EQ(t[2], 1); EXPECT_EQ(t[3], 1); } - -TEST(math_function, selected_rows_add) { - using namespace paddle::framework; - using namespace paddle::platform; - using namespace paddle::operators::math; - - CPUPlace cpu_place; - CPUDeviceContext ctx(cpu_place); - SetConstant functor; - int64_t height = 10; - int64_t row_numel = 10; - - std::vector rows1{0, 4, 7}; - std::unique_ptr selected_rows1{new SelectedRows(rows1, height)}; - auto* in1_value = selected_rows1->mutable_value(); - in1_value->mutable_data( - make_ddim({static_cast(rows1.size()), row_numel}), cpu_place); - functor(ctx, in1_value, 1.0); - - std::vector rows2{0, 5, 7, 9}; - std::unique_ptr selected_rows2{new SelectedRows(rows2, height)}; - auto* in2_value = selected_rows2->mutable_value(); - in2_value->mutable_data( - make_ddim({static_cast(rows2.size()), row_numel}), cpu_place); - functor(ctx, in2_value, 2.0); - - std::unique_ptr output{new SelectedRows()}; - auto* out_value = output->mutable_value(); - - // simplely concat two SelectedRows - out_value->mutable_data(make_ddim({7, 10}), cpu_place); - - SelectedRowsAdd add_functor; - add_functor(ctx, *selected_rows1, *selected_rows2, output.get()); - - auto out_height = output->height(); - EXPECT_EQ(out_height, height); - - auto& out_rows = output->rows(); - - // input1 rows - EXPECT_EQ(out_rows[0], 0); - EXPECT_EQ(out_rows[1], 4); - EXPECT_EQ(out_rows[2], 7); - // input2 rows - EXPECT_EQ(out_rows[3], 0); - EXPECT_EQ(out_rows[4], 5); - EXPECT_EQ(out_rows[5], 7); - EXPECT_EQ(out_rows[6], 9); - - auto* out_data = output->value().data(); - // input1 value - EXPECT_EQ(out_data[0 * row_numel + 0], 1.0); - EXPECT_EQ(out_data[0 * row_numel + 8], 1.0); - EXPECT_EQ(out_data[1 * row_numel + 1], 1.0); - EXPECT_EQ(out_data[2 * row_numel + 6], 1.0); - // input2 value - EXPECT_EQ(out_data[3 * row_numel + 3], 2.0); - EXPECT_EQ(out_data[3 * row_numel + 8], 2.0); - EXPECT_EQ(out_data[4 * row_numel + 4], 2.0); - EXPECT_EQ(out_data[5 * row_numel + 7], 2.0); - EXPECT_EQ(out_data[6 * row_numel + 9], 2.0); - - std::unique_ptr tensor1{new Tensor()}; - tensor1->mutable_data(make_ddim({height, row_numel}), cpu_place); - SetConstant constant_functor; - constant_functor(ctx, tensor1.get(), 3.0); - - std::unique_ptr tensor2{new Tensor()}; - tensor2->mutable_data(make_ddim({height, row_numel}), cpu_place); - - SelectedRowsAddTensor add_tensor_functor; - add_tensor_functor(ctx, *output, *tensor1, tensor2.get()); - - auto* tensor2_data = tensor2->data(); - // row0: 1.0 + 2.0 + 3.0 - EXPECT_EQ(tensor2_data[0 * row_numel + 0], 6.0); - // row1: 3.0 - EXPECT_EQ(tensor2_data[1 * row_numel + 1], 3.0); - // row4 : 1.0 + 3.0 - EXPECT_EQ(tensor2_data[4 * row_numel + 6], 4.0); - // row5: 2.0 + 3.0 - EXPECT_EQ(tensor2_data[5 * row_numel + 7], 5.0); - // row6: 3.0 - EXPECT_EQ(tensor2_data[6 * row_numel + 1], 3.0); - // row7: 1.0 + 2.0 + 3.0 - EXPECT_EQ(tensor2_data[7 * row_numel + 3], 6.0); - // row9: 2.0 + 3.0 - EXPECT_EQ(tensor2_data[9 * row_numel + 6], 5.0); -} diff --git a/paddle/operators/math/math_function_test.cu b/paddle/operators/math/math_function_test.cu index 1acc5f66a6..14359d835b 100644 --- a/paddle/operators/math/math_function_test.cu +++ b/paddle/operators/math/math_function_test.cu @@ -177,102 +177,3 @@ TEST(math_function, gemm_trans_cublas) { EXPECT_EQ(input3_ptr[7], 99); delete gpu_place; } - -TEST(math_function, selected_rows_add) { - using namespace paddle::framework; - using namespace paddle::platform; - using namespace paddle::operators::math; - - GPUPlace gpu_place(0); - CPUPlace cpu_place; - CUDADeviceContext ctx(gpu_place); - SetConstant functor; - int64_t height = 10; - int64_t row_numel = 10; - - std::vector rows1{0, 4, 7}; - std::unique_ptr selected_rows1{new SelectedRows(rows1, height)}; - auto* in1_value = selected_rows1->mutable_value(); - in1_value->mutable_data( - make_ddim({static_cast(rows1.size()), row_numel}), gpu_place); - functor(ctx, in1_value, 1.0); - - std::vector rows2{0, 5, 7, 9}; - std::unique_ptr selected_rows2{new SelectedRows(rows2, height)}; - auto* in2_value = selected_rows2->mutable_value(); - in2_value->mutable_data( - make_ddim({static_cast(rows2.size()), row_numel}), gpu_place); - functor(ctx, in2_value, 2.0); - - std::unique_ptr output{new SelectedRows()}; - auto* out_value = output->mutable_value(); - - // simplely concat two SelectedRows - out_value->mutable_data(make_ddim({7, 10}), gpu_place); - - SelectedRowsAdd add_functor; - add_functor(ctx, *selected_rows1, *selected_rows2, output.get()); - - auto out_height = output->height(); - EXPECT_EQ(out_height, height); - - auto& out_rows = output->rows(); - - // input1 rows - EXPECT_EQ(out_rows[0], 0); - EXPECT_EQ(out_rows[1], 4); - EXPECT_EQ(out_rows[2], 7); - // input2 rows - EXPECT_EQ(out_rows[3], 0); - EXPECT_EQ(out_rows[4], 5); - EXPECT_EQ(out_rows[5], 7); - EXPECT_EQ(out_rows[6], 9); - - Tensor out_cpu; - out_cpu.CopyFrom(*out_value, cpu_place, ctx); - ctx.Wait(); - - auto* out_cpu_data = out_cpu.data(); - // input1 value - EXPECT_EQ(out_cpu_data[0 * row_numel + 0], 1.0); - EXPECT_EQ(out_cpu_data[0 * row_numel + 8], 1.0); - EXPECT_EQ(out_cpu_data[1 * row_numel + 1], 1.0); - EXPECT_EQ(out_cpu_data[2 * row_numel + 6], 1.0); - // input2 value - EXPECT_EQ(out_cpu_data[3 * row_numel + 3], 2.0); - EXPECT_EQ(out_cpu_data[3 * row_numel + 8], 2.0); - EXPECT_EQ(out_cpu_data[4 * row_numel + 4], 2.0); - EXPECT_EQ(out_cpu_data[5 * row_numel + 7], 2.0); - EXPECT_EQ(out_cpu_data[6 * row_numel + 9], 2.0); - - std::unique_ptr tensor1{new Tensor()}; - tensor1->mutable_data(make_ddim({height, row_numel}), gpu_place); - SetConstant constant_functor; - constant_functor(ctx, tensor1.get(), 3.0); - - std::unique_ptr tensor2{new Tensor()}; - tensor2->mutable_data(make_ddim({height, row_numel}), gpu_place); - - SelectedRowsAddTensor add_tensor_functor; - add_tensor_functor(ctx, *output, *tensor1, tensor2.get()); - - Tensor tensor2_cpu; - tensor2_cpu.CopyFrom(*tensor2, cpu_place, ctx); - ctx.Wait(); - - auto* tensor2_cpu_data = tensor2_cpu.data(); - // row0: 1.0 + 2.0 + 3.0 - EXPECT_EQ(tensor2_cpu_data[0 * row_numel + 0], 6.0); - // row1: 3.0 - EXPECT_EQ(tensor2_cpu_data[1 * row_numel + 1], 3.0); - // row4 : 1.0 + 3.0 - EXPECT_EQ(tensor2_cpu_data[4 * row_numel + 6], 4.0); - // row5: 2.0 + 3.0 - EXPECT_EQ(tensor2_cpu_data[5 * row_numel + 7], 5.0); - // row6: 3.0 - EXPECT_EQ(tensor2_cpu_data[6 * row_numel + 1], 3.0); - // row7: 1.0 + 2.0 + 3.0 - EXPECT_EQ(tensor2_cpu_data[7 * row_numel + 3], 6.0); - // row9: 2.0 + 3.0 - EXPECT_EQ(tensor2_cpu_data[9 * row_numel + 6], 5.0); -} diff --git a/paddle/operators/math/selected_rows_functor.cc b/paddle/operators/math/selected_rows_functor.cc new file mode 100644 index 0000000000..f2305ea169 --- /dev/null +++ b/paddle/operators/math/selected_rows_functor.cc @@ -0,0 +1,114 @@ +/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +Licensed under the Apache License, Version 2.0 (the "License"); +you may not use this file except in compliance with the License. +You may obtain a copy of the License at + + http://www.apache.org/licenses/LICENSE-2.0 + +Unless required by applicable law or agreed to in writing, software +distributed under the License is distributed on an "AS IS" BASIS, +WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +See the License for the specific language governing permissions and +limitations under the License. */ + +#include "paddle/operators/math/selected_rows_functor.h" +#include "paddle/operators/math/math_function.h" + +namespace paddle { +namespace operators { +namespace math { +template +struct SelectedRowsAdd { + void operator()(const platform::DeviceContext& context, + const framework::SelectedRows& input1, + const framework::SelectedRows& input2, + framework::SelectedRows* output) { + auto in1_height = input1.height(); + PADDLE_ENFORCE_EQ(in1_height, input2.height()); + output->set_height(in1_height); + + auto& in1_rows = input1.rows(); + auto& in2_rows = input2.rows(); + std::vector out_rows; + out_rows.reserve(in1_rows.size() + in2_rows.size()); + + // concat rows + out_rows.insert(out_rows.end(), in1_rows.begin(), in1_rows.end()); + out_rows.insert(out_rows.end(), in2_rows.begin(), in2_rows.end()); + output->set_rows(out_rows); + + auto* out_value = output->mutable_value(); + auto& in1_value = input1.value(); + auto& in2_value = input2.value(); + + auto in1_row_numel = in1_value.numel() / in1_rows.size(); + PADDLE_ENFORCE_EQ(in1_row_numel, in2_value.numel() / in2_rows.size()); + PADDLE_ENFORCE_EQ(in1_row_numel, out_value->numel() / out_rows.size()); + + auto in1_place = input1.place(); + PADDLE_ENFORCE(platform::is_cpu_place(in1_place)); + auto in2_place = input2.place(); + PADDLE_ENFORCE(platform::is_cpu_place(in2_place)); + auto out_place = context.GetPlace(); + PADDLE_ENFORCE(platform::is_cpu_place(out_place)); + + auto* out_data = out_value->data(); + auto* in1_data = in1_value.data(); + memory::Copy(boost::get(out_place), out_data, + boost::get(in1_place), in1_data, + in1_value.numel() * sizeof(T)); + + auto* in2_data = in2_value.data(); + memory::Copy(boost::get(out_place), + out_data + in1_value.numel(), + boost::get(in2_place), in2_data, + in2_value.numel() * sizeof(T)); + } +}; + +template struct SelectedRowsAdd; + +template +struct SelectedRowsAddTensor { + void operator()(const platform::DeviceContext& context, + const framework::SelectedRows& input1, + const framework::Tensor& input2, framework::Tensor* output) { + auto in1_height = input1.height(); + auto in2_dims = input2.dims(); + auto out_dims = output->dims(); + PADDLE_ENFORCE_EQ(in1_height, in2_dims[0]); + PADDLE_ENFORCE_EQ(in1_height, out_dims[0]); + + auto& in1_value = input1.value(); + auto& in1_rows = input1.rows(); + + int64_t in1_row_numel = in1_value.numel() / in1_rows.size(); + PADDLE_ENFORCE_EQ(in1_row_numel, input2.numel() / in1_height); + PADDLE_ENFORCE_EQ(in1_row_numel, output->numel() / in1_height); + + SetConstant functor; + functor(context, output, 0.0); + + auto* in1_data = in1_value.data(); + auto* out_data = output->data(); + + for (size_t i = 0; i < in1_rows.size(); i++) { + for (int64_t j = 0; j < in1_row_numel; j++) { + out_data[in1_rows[i] * in1_row_numel + j] += + in1_data[i * in1_row_numel + j]; + } + } + + auto out_eigen = framework::EigenVector::Flatten(*output); + auto in2_eigen = framework::EigenVector::Flatten(input2); + out_eigen.device(*context.GetEigenDevice()) = + out_eigen + in2_eigen; + } +}; + +template struct SelectedRowsAddTensor; + +} // namespace math +} // namespace operators +} // namespace paddle diff --git a/paddle/operators/math/selected_rows_functor.cu b/paddle/operators/math/selected_rows_functor.cu new file mode 100644 index 0000000000..a406bef39a --- /dev/null +++ b/paddle/operators/math/selected_rows_functor.cu @@ -0,0 +1,142 @@ +/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +Licensed under the Apache License, Version 2.0 (the "License"); +you may not use this file except in compliance with the License. +You may obtain a copy of the License at + + http://www.apache.org/licenses/LICENSE-2.0 + +Unless required by applicable law or agreed to in writing, software +distributed under the License is distributed on an "AS IS" BASIS, +WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +See the License for the specific language governing permissions and +limitations under the License. */ + +#include "paddle/operators/math/math_function.h" +#include "paddle/operators/math/selected_rows_functor.h" +#include "paddle/platform/cuda_helper.h" + +namespace paddle { +namespace operators { +namespace math { +template +struct SelectedRowsAdd { + void operator()(const platform::DeviceContext& context, + const framework::SelectedRows& input1, + const framework::SelectedRows& input2, + framework::SelectedRows* output) { + auto in1_height = input1.height(); + PADDLE_ENFORCE_EQ(in1_height, input2.height()); + output->set_height(in1_height); + + auto& in1_rows = input1.rows(); + auto& in2_rows = input2.rows(); + std::vector out_rows; + out_rows.reserve(in1_rows.size() + in2_rows.size()); + + // concat rows + out_rows.insert(out_rows.end(), in1_rows.begin(), in1_rows.end()); + out_rows.insert(out_rows.end(), in2_rows.begin(), in2_rows.end()); + output->set_rows(out_rows); + + auto* out_value = output->mutable_value(); + auto& in1_value = input1.value(); + auto& in2_value = input2.value(); + + auto in1_row_numel = in1_value.numel() / in1_rows.size(); + PADDLE_ENFORCE_EQ(in1_row_numel, in2_value.numel() / in2_rows.size()); + PADDLE_ENFORCE_EQ(in1_row_numel, out_value->numel() / out_rows.size()); + + auto* out_data = out_value->data(); + auto* in1_data = in1_value.data(); + + auto in1_place = input1.place(); + PADDLE_ENFORCE(platform::is_gpu_place(in1_place)); + auto in2_place = input2.place(); + PADDLE_ENFORCE(platform::is_gpu_place(in2_place)); + auto out_place = context.GetPlace(); + PADDLE_ENFORCE(platform::is_gpu_place(out_place)); + + memory::Copy( + boost::get(out_place), out_data, + boost::get(in1_place), in1_data, + in1_value.numel() * sizeof(T), + reinterpret_cast(context).stream()); + + auto* in2_data = in2_value.data(); + memory::Copy( + boost::get(out_place), out_data + in1_value.numel(), + boost::get(in2_place), in2_data, + in2_value.numel() * sizeof(T), + reinterpret_cast(context).stream()); + } +}; + +template struct SelectedRowsAdd; + +namespace { +template +__global__ void SelectedRowsAddTensorKernel(const T* selected_rows, + const int64_t* rows, T* tensor_out, + int64_t row_numel, int block_size) { + const int ty = blockIdx.y; + int tid = threadIdx.x; + + selected_rows += ty * row_numel; + tensor_out += rows[ty] * row_numel; + + for (int index = tid; index < row_numel; index += block_size) { + // Since index in rows of SelectedRows can be duplicate, we can not use + // tensor_out[index] += selected_rows[index]; Instead, we have to use + // AtomicAdd to avoid concurrent write error. + paddle::platform::CudaAtomicAdd(&tensor_out[index], selected_rows[index]); + } +} +} // namespace + +template +struct SelectedRowsAddTensor { + void operator()(const platform::DeviceContext& context, + const framework::SelectedRows& input1, + const framework::Tensor& input2, framework::Tensor* output) { + auto in1_height = input1.height(); + auto in2_dims = input2.dims(); + auto out_dims = output->dims(); + PADDLE_ENFORCE_EQ(in1_height, in2_dims[0]); + PADDLE_ENFORCE_EQ(in1_height, out_dims[0]); + + auto& in1_value = input1.value(); + auto& in1_rows = input1.rows(); + + int64_t in1_row_numel = in1_value.numel() / in1_rows.size(); + PADDLE_ENFORCE_EQ(in1_row_numel, input2.numel() / in1_height); + PADDLE_ENFORCE_EQ(in1_row_numel, output->numel() / in1_height); + + auto* in1_data = in1_value.data(); + auto* in2_data = input2.data(); + auto* out_data = output->data(); + + SetConstant functor; + functor(context, output, 0.0); + + int block_size = 256; + dim3 threads(block_size, 1); + dim3 grid(1, in1_height); + SelectedRowsAddTensorKernel< + T><<(context) + .stream()>>>(in1_data, in1_rows.data(), out_data, + in1_row_numel, block_size); + + auto out_eigen = framework::EigenVector::Flatten(*output); + auto in2_eigen = framework::EigenVector::Flatten(input2); + out_eigen.device(*context.GetEigenDevice()) = + out_eigen + in2_eigen; + } +}; + +template struct SelectedRowsAddTensor; + +} // namespace math +} // namespace operators +} // namespace paddle diff --git a/paddle/operators/math/selected_rows_functor.h b/paddle/operators/math/selected_rows_functor.h new file mode 100644 index 0000000000..53ab240ca6 --- /dev/null +++ b/paddle/operators/math/selected_rows_functor.h @@ -0,0 +1,41 @@ +/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +Licensed under the Apache License, Version 2.0 (the "License"); +you may not use this file except in compliance with the License. +You may obtain a copy of the License at + + http://www.apache.org/licenses/LICENSE-2.0 + +Unless required by applicable law or agreed to in writing, software +distributed under the License is distributed on an "AS IS" BASIS, +WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +See the License for the specific language governing permissions and +limitations under the License. */ +#pragma once +#include "paddle/framework/selected_rows.h" +#include "paddle/platform/device_context.h" + +namespace paddle { +namespace operators { +namespace math { + +// SelectedRows + SelectedRows will simplely concat value and rows. +// The real computation happens in dealing with LoDTensor. +template +struct SelectedRowsAdd { + void operator()(const platform::DeviceContext& context, + const framework::SelectedRows& input1, + const framework::SelectedRows& input2, + framework::SelectedRows* output); +}; + +template +struct SelectedRowsAddTensor { + void operator()(const platform::DeviceContext& context, + const framework::SelectedRows& input1, + const framework::Tensor& input2, framework::Tensor* output); +}; + +} // namespace math +} // namespace operators +} // namespace paddle diff --git a/paddle/operators/math/selected_rows_functor_test.cc b/paddle/operators/math/selected_rows_functor_test.cc new file mode 100644 index 0000000000..4f7760cb71 --- /dev/null +++ b/paddle/operators/math/selected_rows_functor_test.cc @@ -0,0 +1,106 @@ +/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +Licensed under the Apache License, Version 2.0 (the "License"); +you may not use this file except in compliance with the License. +You may obtain a copy of the License at + + http://www.apache.org/licenses/LICENSE-2.0 + +Unless required by applicable law or agreed to in writing, software +distributed under the License is distributed on an "AS IS" BASIS, +WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +See the License for the specific language governing permissions and +limitations under the License. */ + +#include "paddle/operators/math/selected_rows_functor.h" +#include "gtest/gtest.h" +#include "paddle/operators/math/math_function.h" + +TEST(selected_rows_functor, cpu_add) { + using namespace paddle::framework; + using namespace paddle::platform; + using namespace paddle::operators::math; + + CPUPlace cpu_place; + CPUDeviceContext ctx(cpu_place); + SetConstant functor; + int64_t height = 10; + int64_t row_numel = 10; + + std::vector rows1{0, 4, 7}; + std::unique_ptr selected_rows1{new SelectedRows(rows1, height)}; + auto* in1_value = selected_rows1->mutable_value(); + in1_value->mutable_data( + make_ddim({static_cast(rows1.size()), row_numel}), cpu_place); + functor(ctx, in1_value, 1.0); + + std::vector rows2{0, 5, 7, 9}; + std::unique_ptr selected_rows2{new SelectedRows(rows2, height)}; + auto* in2_value = selected_rows2->mutable_value(); + in2_value->mutable_data( + make_ddim({static_cast(rows2.size()), row_numel}), cpu_place); + functor(ctx, in2_value, 2.0); + + std::unique_ptr output{new SelectedRows()}; + auto* out_value = output->mutable_value(); + + // simplely concat two SelectedRows + out_value->mutable_data(make_ddim({7, 10}), cpu_place); + + SelectedRowsAdd add_functor; + add_functor(ctx, *selected_rows1, *selected_rows2, output.get()); + + auto out_height = output->height(); + EXPECT_EQ(out_height, height); + + auto& out_rows = output->rows(); + + // input1 rows + EXPECT_EQ(out_rows[0], 0); + EXPECT_EQ(out_rows[1], 4); + EXPECT_EQ(out_rows[2], 7); + // input2 rows + EXPECT_EQ(out_rows[3], 0); + EXPECT_EQ(out_rows[4], 5); + EXPECT_EQ(out_rows[5], 7); + EXPECT_EQ(out_rows[6], 9); + + auto* out_data = output->value().data(); + // input1 value + EXPECT_EQ(out_data[0 * row_numel + 0], 1.0); + EXPECT_EQ(out_data[0 * row_numel + 8], 1.0); + EXPECT_EQ(out_data[1 * row_numel + 1], 1.0); + EXPECT_EQ(out_data[2 * row_numel + 6], 1.0); + // input2 value + EXPECT_EQ(out_data[3 * row_numel + 3], 2.0); + EXPECT_EQ(out_data[3 * row_numel + 8], 2.0); + EXPECT_EQ(out_data[4 * row_numel + 4], 2.0); + EXPECT_EQ(out_data[5 * row_numel + 7], 2.0); + EXPECT_EQ(out_data[6 * row_numel + 9], 2.0); + + std::unique_ptr tensor1{new Tensor()}; + tensor1->mutable_data(make_ddim({height, row_numel}), cpu_place); + functor(ctx, tensor1.get(), 3.0); + + std::unique_ptr tensor2{new Tensor()}; + tensor2->mutable_data(make_ddim({height, row_numel}), cpu_place); + + SelectedRowsAddTensor add_tensor_functor; + add_tensor_functor(ctx, *output, *tensor1, tensor2.get()); + + auto* tensor2_data = tensor2->data(); + // row0: 1.0 + 2.0 + 3.0 + EXPECT_EQ(tensor2_data[0 * row_numel + 0], 6.0); + // row1: 3.0 + EXPECT_EQ(tensor2_data[1 * row_numel + 1], 3.0); + // row4 : 1.0 + 3.0 + EXPECT_EQ(tensor2_data[4 * row_numel + 6], 4.0); + // row5: 2.0 + 3.0 + EXPECT_EQ(tensor2_data[5 * row_numel + 7], 5.0); + // row6: 3.0 + EXPECT_EQ(tensor2_data[6 * row_numel + 1], 3.0); + // row7: 1.0 + 2.0 + 3.0 + EXPECT_EQ(tensor2_data[7 * row_numel + 3], 6.0); + // row9: 2.0 + 3.0 + EXPECT_EQ(tensor2_data[9 * row_numel + 6], 5.0); +} diff --git a/paddle/operators/math/selected_rows_functor_test.cu b/paddle/operators/math/selected_rows_functor_test.cu new file mode 100644 index 0000000000..8a9f25b982 --- /dev/null +++ b/paddle/operators/math/selected_rows_functor_test.cu @@ -0,0 +1,115 @@ +/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +Licensed under the Apache License, Version 2.0 (the "License"); +you may not use this file except in compliance with the License. +You may obtain a copy of the License at + + http://www.apache.org/licenses/LICENSE-2.0 + +Unless required by applicable law or agreed to in writing, software +distributed under the License is distributed on an "AS IS" BASIS, +WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +See the License for the specific language governing permissions and +limitations under the License. */ + +#include "gtest/gtest.h" +#include "paddle/operators/math/math_function.h" +#include "paddle/operators/math/selected_rows_functor.h" + +TEST(selected_rows_functor, gpu_add) { + using namespace paddle::framework; + using namespace paddle::platform; + using namespace paddle::operators::math; + + GPUPlace gpu_place(0); + CPUPlace cpu_place; + CUDADeviceContext ctx(gpu_place); + SetConstant functor; + int64_t height = 10; + int64_t row_numel = 10; + + std::vector rows1{0, 4, 7}; + std::unique_ptr selected_rows1{new SelectedRows(rows1, height)}; + auto* in1_value = selected_rows1->mutable_value(); + in1_value->mutable_data( + make_ddim({static_cast(rows1.size()), row_numel}), gpu_place); + functor(ctx, in1_value, 1.0); + + std::vector rows2{0, 5, 7, 9}; + std::unique_ptr selected_rows2{new SelectedRows(rows2, height)}; + auto* in2_value = selected_rows2->mutable_value(); + in2_value->mutable_data( + make_ddim({static_cast(rows2.size()), row_numel}), gpu_place); + functor(ctx, in2_value, 2.0); + + std::unique_ptr output{new SelectedRows()}; + auto* out_value = output->mutable_value(); + + // simplely concat two SelectedRows + out_value->mutable_data(make_ddim({7, 10}), gpu_place); + + SelectedRowsAdd add_functor; + add_functor(ctx, *selected_rows1, *selected_rows2, output.get()); + + auto out_height = output->height(); + EXPECT_EQ(out_height, height); + + auto& out_rows = output->rows(); + + // input1 rows + EXPECT_EQ(out_rows[0], 0); + EXPECT_EQ(out_rows[1], 4); + EXPECT_EQ(out_rows[2], 7); + // input2 rows + EXPECT_EQ(out_rows[3], 0); + EXPECT_EQ(out_rows[4], 5); + EXPECT_EQ(out_rows[5], 7); + EXPECT_EQ(out_rows[6], 9); + + Tensor out_cpu; + out_cpu.CopyFrom(*out_value, cpu_place, ctx); + ctx.Wait(); + + auto* out_cpu_data = out_cpu.data(); + // input1 value + EXPECT_EQ(out_cpu_data[0 * row_numel + 0], 1.0); + EXPECT_EQ(out_cpu_data[0 * row_numel + 8], 1.0); + EXPECT_EQ(out_cpu_data[1 * row_numel + 1], 1.0); + EXPECT_EQ(out_cpu_data[2 * row_numel + 6], 1.0); + // input2 value + EXPECT_EQ(out_cpu_data[3 * row_numel + 3], 2.0); + EXPECT_EQ(out_cpu_data[3 * row_numel + 8], 2.0); + EXPECT_EQ(out_cpu_data[4 * row_numel + 4], 2.0); + EXPECT_EQ(out_cpu_data[5 * row_numel + 7], 2.0); + EXPECT_EQ(out_cpu_data[6 * row_numel + 9], 2.0); + + std::unique_ptr tensor1{new Tensor()}; + tensor1->mutable_data(make_ddim({height, row_numel}), gpu_place); + functor(ctx, tensor1.get(), 3.0); + + std::unique_ptr tensor2{new Tensor()}; + tensor2->mutable_data(make_ddim({height, row_numel}), gpu_place); + + SelectedRowsAddTensor add_tensor_functor; + add_tensor_functor(ctx, *output, *tensor1, tensor2.get()); + + Tensor tensor2_cpu; + tensor2_cpu.CopyFrom(*tensor2, cpu_place, ctx); + ctx.Wait(); + + auto* tensor2_cpu_data = tensor2_cpu.data(); + // row0: 1.0 + 2.0 + 3.0 + EXPECT_EQ(tensor2_cpu_data[0 * row_numel + 0], 6.0); + // row1: 3.0 + EXPECT_EQ(tensor2_cpu_data[1 * row_numel + 1], 3.0); + // row4 : 1.0 + 3.0 + EXPECT_EQ(tensor2_cpu_data[4 * row_numel + 6], 4.0); + // row5: 2.0 + 3.0 + EXPECT_EQ(tensor2_cpu_data[5 * row_numel + 7], 5.0); + // row6: 3.0 + EXPECT_EQ(tensor2_cpu_data[6 * row_numel + 1], 3.0); + // row7: 1.0 + 2.0 + 3.0 + EXPECT_EQ(tensor2_cpu_data[7 * row_numel + 3], 6.0); + // row9: 2.0 + 3.0 + EXPECT_EQ(tensor2_cpu_data[9 * row_numel + 6], 5.0); +} From 0d7b1fd4feed8d5249c6904f10126c7d58319503 Mon Sep 17 00:00:00 2001 From: qijun Date: Mon, 16 Oct 2017 08:28:03 +0000 Subject: [PATCH 19/36] fix bug in SelectedRowsAddTensorKernel --- paddle/operators/math/selected_rows_functor.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/paddle/operators/math/selected_rows_functor.cu b/paddle/operators/math/selected_rows_functor.cu index a406bef39a..ea149ebbc1 100644 --- a/paddle/operators/math/selected_rows_functor.cu +++ b/paddle/operators/math/selected_rows_functor.cu @@ -89,7 +89,7 @@ __global__ void SelectedRowsAddTensorKernel(const T* selected_rows, // Since index in rows of SelectedRows can be duplicate, we can not use // tensor_out[index] += selected_rows[index]; Instead, we have to use // AtomicAdd to avoid concurrent write error. - paddle::platform::CudaAtomicAdd(&tensor_out[index], selected_rows[index]); + paddle::platform::CudaAtomicAdd(tensor_out + index, selected_rows[index]); } } } // namespace @@ -121,7 +121,7 @@ struct SelectedRowsAddTensor { int block_size = 256; dim3 threads(block_size, 1); - dim3 grid(1, in1_height); + dim3 grid(1, in1_rows.size()); SelectedRowsAddTensorKernel< T><<(context) From 4df6cf4d16ad271101bd37de7f84fb054f1d788a Mon Sep 17 00:00:00 2001 From: Yu Yang Date: Mon, 16 Oct 2017 09:59:58 -0700 Subject: [PATCH 20/36] Rewrite feed/fetch op (#4815) * Feed/Fetch op just plain operator, not a OpWithKernel * Do not register OpInfoMaker since Feed/Fetch will never be configured by users * Feed/Fetch op has empty gradient * Feed/Fetch op do not hard code `feed_variable`, `fetch_variable` as its input and output, make it as a plain Operator input/output --- paddle/framework/executor_test.cc | 55 ++++++++++------- paddle/framework/feed_fetch_type.h | 24 ++++++++ paddle/framework/grad_op_desc_maker.h | 8 +++ paddle/framework/op_desc.cc | 9 ++- paddle/framework/tensor.h | 16 +++++ paddle/operators/feed_op.cc | 86 +++++++++++++-------------- paddle/operators/feed_op.cu | 18 ------ paddle/operators/feed_op.h | 42 ------------- paddle/operators/fetch_op.cc | 78 ++++++++++++++---------- paddle/operators/fetch_op.cu | 18 ------ paddle/operators/fetch_op.h | 45 -------------- 11 files changed, 173 insertions(+), 226 deletions(-) create mode 100644 paddle/framework/feed_fetch_type.h delete mode 100644 paddle/operators/feed_op.cu delete mode 100644 paddle/operators/feed_op.h delete mode 100644 paddle/operators/fetch_op.cu delete mode 100644 paddle/operators/fetch_op.h diff --git a/paddle/framework/executor_test.cc b/paddle/framework/executor_test.cc index fcd2e47cff..e08d31e361 100644 --- a/paddle/framework/executor_test.cc +++ b/paddle/framework/executor_test.cc @@ -28,8 +28,8 @@ limitations under the License. */ USE_OP(elementwise_add); USE_OP(gaussian_random); -USE_OP(feed); -USE_OP(fetch); +USE_NO_KERNEL_OP(feed); +USE_NO_KERNEL_OP(fetch); USE_OP(mul); USE_OP(sum); USE_OP(squared_l2_distance); @@ -37,6 +37,9 @@ USE_OP(fill_constant); USE_OP(mean); USE_OP(sgd); +constexpr auto kFeedValueName = "feed_value"; +constexpr auto kFetchValueName = "fetch_value"; + using namespace paddle::platform; using namespace paddle::framework; @@ -77,9 +80,9 @@ void AddOp(const std::string& type, const VariableNameMap& inputs, template void SetFeedVariable(const std::vector>& inputs, const std::vector>& dims) { - Variable* g_feed_value = GetGlobalScope().FindVar("feed_value"); + Variable* g_feed_value = GetGlobalScope().FindVar(kFeedValueName); auto& feed_inputs = - *(g_feed_value->GetMutable>()); + *(g_feed_value->GetMutable>()); size_t size = inputs.size(); feed_inputs.resize(size); for (size_t i = 0; i < size; i++) { @@ -92,9 +95,9 @@ void SetFeedVariable(const std::vector>& inputs, // So we can memcpy the data from fetch_value to vector template std::vector> GetFetchVariable() { - Variable* g_fetch_value = GetGlobalScope().FindVar("fetch_value"); + Variable* g_fetch_value = GetGlobalScope().FindVar(kFetchValueName); auto& fetch_outputs = - *(g_fetch_value->GetMutable>()); + *(g_fetch_value->GetMutable>()); size_t size = fetch_outputs.size(); std::vector> result; @@ -126,8 +129,10 @@ class ExecutorTesterRandom : public ::testing::Test { {{"dims", std::vector{input_dim, embed_dim}}}, init_root_block); AddOp("gaussian_random", {}, {{"Out", {"w2"}}}, {{"dims", std::vector{embed_dim, input_dim}}}, init_root_block); - AddOp("fetch", {{"Input", {"w1"}}}, {}, {{"col", 0}}, init_root_block); - AddOp("fetch", {{"Input", {"w2"}}}, {}, {{"col", 1}}, init_root_block); + AddOp("fetch", {{"Input", {"w1"}}}, {{"Out", {kFetchValueName}}}, + {{"col", 0}}, init_root_block); + AddOp("fetch", {{"Input", {"w2"}}}, {{"Out", {kFetchValueName}}}, + {{"col", 1}}, init_root_block); // flush init_program.Proto(); @@ -143,7 +148,7 @@ class ExecutorTesterRandom : public ::testing::Test { // feed data inputs_.push_back({1.0, 1.0, 1.0, 1.0, 1.0, 1.0}); dims_.push_back({batch_size, input_dim}); - AddOp("feed", {}, {{"Out", {"a"}}}, + AddOp("feed", {{"Input", {kFeedValueName}}}, {{"Out", {"a"}}}, {{"dims", std::vector{batch_size, input_dim}}, {"col", 0}}, root_block); @@ -175,9 +180,12 @@ class ExecutorTesterRandom : public ::testing::Test { {"Grad", {"w2@GRAD"}}}, {{"ParamOut", {"w2"}}}, {}, root_block); - AddOp("fetch", {{"Input", {"w1"}}}, {}, {{"col", 0}}, root_block); - AddOp("fetch", {{"Input", {"w2"}}}, {}, {{"col", 1}}, root_block); - AddOp("fetch", {{"Input", {"l2_distance"}}}, {}, {{"col", 0}}, root_block); + AddOp("fetch", {{"Input", {"w1"}}}, {{"Out", {kFetchValueName}}}, + {{"col", 0}}, root_block); + AddOp("fetch", {{"Input", {"w2"}}}, {{"Out", {kFetchValueName}}}, + {{"col", 1}}, root_block); + AddOp("fetch", {{"Input", {"l2_distance"}}}, {{"Out", {kFetchValueName}}}, + {{"col", 0}}, root_block); // flush program.Proto(); @@ -204,12 +212,14 @@ class ExecutorTesterFeedAndFetch : public ::testing::Test { std::vector dim{6}; - AddOp("feed", {}, {{"Out", {"a"}}}, {{"dims", dim}, {"col", 0}}, - root_block); - AddOp("feed", {}, {{"Out", {"b"}}}, {{"dims", dim}, {"col", 1}}, - root_block); - AddOp("fetch", {{"Input", {"a"}}}, {}, {{"col", 0}}, root_block); - AddOp("fetch", {{"Input", {"b"}}}, {}, {{"col", 1}}, root_block); + AddOp("feed", {{"Input", {kFeedValueName}}}, {{"Out", {"a"}}}, + {{"dims", dim}, {"col", 0}}, root_block); + AddOp("feed", {{"Input", {kFeedValueName}}}, {{"Out", {"b"}}}, + {{"dims", dim}, {"col", 1}}, root_block); + AddOp("fetch", {{"Input", {"a"}}}, {{"Out", {kFetchValueName}}}, + {{"col", 0}}, root_block); + AddOp("fetch", {{"Input", {"b"}}}, {{"Out", {kFetchValueName}}}, + {{"col", 1}}, root_block); // flush program.Proto(); @@ -241,7 +251,6 @@ TEST_F(ExecutorTesterRandom, CPU) { paddle::memory::Used(cpu_place); std::unique_ptr executor(new Executor(places)); - executor->Run(init_pdesc_, &GetGlobalScope(), 0); SetFeedVariable(inputs_, dims_); executor->Run(pdesc_, &GetGlobalScope(), 0); @@ -251,7 +260,7 @@ TEST_F(ExecutorTesterRandom, CPU) { TEST_F(ExecutorTesterFeedAndFetch, CPU) { std::vector places; CPUPlace cpu_place; - places.push_back(cpu_place); + places.emplace_back(cpu_place); // We have a global Scope and BuddyAllocator, and we must ensure // global BuddyAllocator is initialized before global Scope. Thus, @@ -265,11 +274,11 @@ TEST_F(ExecutorTesterFeedAndFetch, CPU) { SetFeedVariable(inputs_, dims_); executor->Run(pdesc_, &GetGlobalScope(), 0); std::vector> result = GetFetchVariable(); - PADDLE_ENFORCE_EQ(result.size(), inputs_.size()); + ASSERT_EQ(result.size(), inputs_.size()); for (size_t i = 0; i < result.size(); ++i) { - PADDLE_ENFORCE_EQ(result[i].size(), inputs_[i].size()); + ASSERT_EQ(result[i].size(), inputs_[i].size()); for (size_t j = 0; j < result[i].size(); ++j) { - PADDLE_ENFORCE_EQ(result[i][j], inputs_[i][j]); + ASSERT_EQ(result[i][j], inputs_[i][j]); } } } diff --git a/paddle/framework/feed_fetch_type.h b/paddle/framework/feed_fetch_type.h new file mode 100644 index 0000000000..bc4ae440fc --- /dev/null +++ b/paddle/framework/feed_fetch_type.h @@ -0,0 +1,24 @@ +/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + + Licensed under the Apache License, Version 2.0 (the "License"); + you may not use this file except in compliance with the License. + You may obtain a copy of the License at + + http://www.apache.org/licenses/LICENSE-2.0 + + Unless required by applicable law or agreed to in writing, software + distributed under the License is distributed on an "AS IS" BASIS, + WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + See the License for the specific language governing permissions and + limitations under the License. */ + +#pragma once +#include +#include "paddle/framework/lod_tensor.h" + +namespace paddle { +namespace framework { +using FeedFetchType = LoDTensor; +using FeedFetchList = std::vector; +} // namespace framework +} // namespace paddle diff --git a/paddle/framework/grad_op_desc_maker.h b/paddle/framework/grad_op_desc_maker.h index 1219e04875..94944c79b6 100644 --- a/paddle/framework/grad_op_desc_maker.h +++ b/paddle/framework/grad_op_desc_maker.h @@ -149,5 +149,13 @@ class DefaultGradOpDescMaker : public SingleGradOpDescMaker { } }; +class EmptyGradOpMaker : public GradOpDescMakerBase { + public: + using GradOpDescMakerBase::GradOpDescMakerBase; + std::vector> operator()() const override { + return {}; + } +}; + } // namespace framework } // namespace paddle diff --git a/paddle/framework/op_desc.cc b/paddle/framework/op_desc.cc index ef207dc54e..7f7cebb026 100644 --- a/paddle/framework/op_desc.cc +++ b/paddle/framework/op_desc.cc @@ -220,9 +220,12 @@ static InferShapeFuncMap &InferShapeFuncs() { void OpDescBind::CheckAttrs() { PADDLE_ENFORCE(!Type().empty(), "CheckAttr() can not be called before type is setted."); - const auto *checker = OpInfoMap::Instance().Get(Type()).Checker(); - PADDLE_ENFORCE_NOT_NULL(checker, "Operator \"%s\" has no registered checker.", - Type()); + auto *checker = OpInfoMap::Instance().Get(Type()).Checker(); + if (checker == nullptr) { + // checker is not configured. That operator could be generated by Paddle, + // not by users. + return; + } checker->Check(attrs_); } diff --git a/paddle/framework/tensor.h b/paddle/framework/tensor.h index 3304d857ae..bc430852de 100644 --- a/paddle/framework/tensor.h +++ b/paddle/framework/tensor.h @@ -100,6 +100,22 @@ class Tensor { inline void CopyFrom(const Tensor& src, const platform::Place& dst_place, const platform::DeviceContext& ctx); + // FIXME(yuyang18): CopyFrom should without template T, use the replace + // `CopyFrom` with `CopyFromTensor` + inline void CopyFromTensor(const Tensor& src, + const platform::Place& dst_place, + const platform::DeviceContext& ctx) { + // NOLINTNEXTLINES_8 cpplint.py will recognize below lines as functions. + // That is a bug of cpplint.py. Just ignore lint these lines. + if (src.type() == std::type_index(typeid(double))) { + CopyFrom(src, dst_place, ctx); + } else if (src.type() == std::type_index(typeid(float))) { + CopyFrom(src, dst_place, ctx); + } else if (src.type() == std::type_index(typeid(int))) { + CopyFrom(src, dst_place, ctx); + } + } + /** * @brief Copy the content of an external vector to a tensor. * diff --git a/paddle/operators/feed_op.cc b/paddle/operators/feed_op.cc index fa325bb282..d742bbe51b 100644 --- a/paddle/operators/feed_op.cc +++ b/paddle/operators/feed_op.cc @@ -1,59 +1,57 @@ /* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. -Licensed under the Apache License, Version 2.0 (the "License"); -you may not use this file except in compliance with the License. -You may obtain a copy of the License at + Licensed under the Apache License, Version 2.0 (the "License"); + you may not use this file except in compliance with the License. + You may obtain a copy of the License at - http://www.apache.org/licenses/LICENSE-2.0 + http://www.apache.org/licenses/LICENSE-2.0 -Unless required by applicable law or agreed to in writing, software -distributed under the License is distributed on an "AS IS" BASIS, -WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -See the License for the specific language governing permissions and -limitations under the License. */ + Unless required by applicable law or agreed to in writing, software + distributed under the License is distributed on an "AS IS" BASIS, + WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + See the License for the specific language governing permissions and + limitations under the License. */ -#include "paddle/operators/feed_op.h" +#include "paddle/framework/feed_fetch_type.h" +#include "paddle/framework/op_registry.h" +#include "paddle/framework/operator.h" namespace paddle { namespace operators { - -class FeedOp : public framework::OperatorWithKernel { - public: - using framework::OperatorWithKernel::OperatorWithKernel; - - protected: - void InferShape(framework::InferShapeContext* ctx) const override { - PADDLE_ENFORCE(ctx->HasOutput("Out"), "Output should be not null."); - auto& shape = ctx->Attrs().Get>("dims"); - std::vector shape_int64(shape.size(), 0); - std::transform(shape.begin(), shape.end(), shape_int64.begin(), - [](int a) { return static_cast(a); }); - ctx->SetOutputDim("Out", framework::make_ddim(shape_int64)); - // TODO(qijun): need to handle LodTensor later - } - - framework::DataType IndicateDataType( - const framework::ExecutionContext& ctx) const override { - return static_cast(Attr("dataType")); - } -}; - -class FeedOpMaker : public framework::OpProtoAndCheckerMaker { +class FeedOp : public framework::OperatorBase { public: - FeedOpMaker(framework::OpProto* proto, framework::OpAttrChecker* op_checker) - : OpProtoAndCheckerMaker(proto, op_checker) { - AddAttr("dataType", "output data type") - .SetDefault(framework::DataType::FP32); - AddAttr("col", "The col in global feed variable").SetDefault(0); - AddAttr>("dims", "The dimension of feed tensor."); - AddOutput("Out", "The output of feed op."); - AddComment(R"DOC(Feed data from global feed variable)DOC"); + FeedOp(const std::string &type, const framework::VariableNameMap &inputs, + const framework::VariableNameMap &outputs, + const framework::AttributeMap &attrs) + : OperatorBase(type, inputs, outputs, attrs) {} + void Run(const framework::Scope &scope, + const platform::DeviceContext &dev_ctx) const override { + auto feed_var_name = Input("Input"); + auto *feed_var = scope.FindVar(feed_var_name); + PADDLE_ENFORCE(feed_var != nullptr, + "Cannot find feed_var in scope, feed_var_name is %s", + feed_var_name); + + auto out_name = this->Output("Out"); + auto *out_var = scope.FindVar(out_name); + PADDLE_ENFORCE(out_var != nullptr, + "Cannot find out_var in scope, out_var_name is %s", + out_name); + + auto col = Attr("col"); + + auto &feed_list = feed_var->Get(); + auto &feed_item = feed_list.at(static_cast(col)); + auto *out_item = out_var->GetMutable(); + out_item->CopyFromTensor(feed_item, dev_ctx.GetPlace(), dev_ctx); + out_item->set_lod(feed_item.lod()); } }; } // namespace operators } // namespace paddle -namespace ops = paddle::operators; -REGISTER_OP_WITHOUT_GRADIENT(feed, ops::FeedOp, ops::FeedOpMaker); -REGISTER_OP_CPU_KERNEL(feed, ops::FeedKernel); +// We do not need to register OpInfoMaker, +// since feed operator will not be used by end users directly +REGISTER_OPERATOR(feed, paddle::operators::FeedOp, + paddle::framework::EmptyGradOpMaker); diff --git a/paddle/operators/feed_op.cu b/paddle/operators/feed_op.cu deleted file mode 100644 index 7b6a2ac91e..0000000000 --- a/paddle/operators/feed_op.cu +++ /dev/null @@ -1,18 +0,0 @@ -/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. - -Licensed under the Apache License, Version 2.0 (the "License"); -you may not use this file except in compliance with the License. -You may obtain a copy of the License at - -http://www.apache.org/licenses/LICENSE-2.0 - -Unless required by applicable law or agreed to in writing, software -distributed under the License is distributed on an "AS IS" BASIS, -WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -See the License for the specific language governing permissions and -limitations under the License. */ - -#include "paddle/operators/feed_op.h" - -namespace ops = paddle::operators; -REGISTER_OP_GPU_KERNEL(feed, ops::FeedKernel); diff --git a/paddle/operators/feed_op.h b/paddle/operators/feed_op.h deleted file mode 100644 index e756cd1842..0000000000 --- a/paddle/operators/feed_op.h +++ /dev/null @@ -1,42 +0,0 @@ -/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. - -Licensed under the Apache License, Version 2.0 (the "License"); -you may not use this file except in compliance with the License. -You may obtain a copy of the License at - - http://www.apache.org/licenses/LICENSE-2.0 - -Unless required by applicable law or agreed to in writing, software -distributed under the License is distributed on an "AS IS" BASIS, -WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -See the License for the specific language governing permissions and -limitations under the License. */ - -#pragma once -#include "paddle/framework/eigen.h" -#include "paddle/framework/op_registry.h" - -namespace paddle { -namespace operators { - -template -class FeedKernel : public framework::OpKernel { - public: - void Compute(const framework::ExecutionContext& ctx) const override { - framework::Tensor* out = ctx.Output("Out"); - out->mutable_data(ctx.GetPlace()); - framework::Variable* g_feed_variable = - framework::GetGlobalScope().FindVar("feed_value"); - const auto& tensors = - g_feed_variable->Get>(); - int col = ctx.template Attr("col"); - PADDLE_ENFORCE_GT(tensors.size(), static_cast(col)); - // TODO(qijun): - // check tensors[col].dims() with attribute, - // except the first dimenson. - out->CopyFrom(tensors[col], ctx.GetPlace(), ctx.device_context()); - } -}; - -} // namespace operators -} // namespace paddle diff --git a/paddle/operators/fetch_op.cc b/paddle/operators/fetch_op.cc index 90737c8c55..55d6ac0939 100644 --- a/paddle/operators/fetch_op.cc +++ b/paddle/operators/fetch_op.cc @@ -1,52 +1,64 @@ /* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. -Licensed under the Apache License, Version 2.0 (the "License"); -you may not use this file except in compliance with the License. -You may obtain a copy of the License at + Licensed under the Apache License, Version 2.0 (the "License"); + you may not use this file except in compliance with the License. + You may obtain a copy of the License at - http://www.apache.org/licenses/LICENSE-2.0 + http://www.apache.org/licenses/LICENSE-2.0 -Unless required by applicable law or agreed to in writing, software -distributed under the License is distributed on an "AS IS" BASIS, -WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -See the License for the specific language governing permissions and -limitations under the License. */ + Unless required by applicable law or agreed to in writing, software + distributed under the License is distributed on an "AS IS" BASIS, + WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + See the License for the specific language governing permissions and + limitations under the License. */ -#include "paddle/operators/fetch_op.h" +#include "paddle/framework/feed_fetch_type.h" +#include "paddle/framework/op_registry.h" namespace paddle { namespace operators { -class FetchOp : public framework::OperatorWithKernel { +class FetchOp : public framework::OperatorBase { public: - using framework::OperatorWithKernel::OperatorWithKernel; + FetchOp(const std::string &type, const framework::VariableNameMap &inputs, + const framework::VariableNameMap &outputs, + const framework::AttributeMap &attrs) + : OperatorBase(type, inputs, outputs, attrs) {} - protected: - void InferShape(framework::InferShapeContext* ctx) const override { - PADDLE_ENFORCE(ctx->HasInput("Input"), "Input should be not null."); - } + void Run(const framework::Scope &scope, + const platform::DeviceContext &dev_ctx) const override { + auto fetch_var_name = Input("Input"); + auto *fetch_var = scope.FindVar(fetch_var_name); + PADDLE_ENFORCE(fetch_var != nullptr, + "Cannot find fetch variable in scope, fetch_var_name is %s", + fetch_var_name); - framework::DataType IndicateDataType( - const framework::ExecutionContext& ctx) const override { - return static_cast(Attr("dataType")); - } -}; + auto out_name = this->Output("Out"); + auto *out_var = scope.FindVar(out_name); + PADDLE_ENFORCE(out_var != nullptr, + "Cannot find out_var in scope, out_var_name is %s", + out_name); -class FetchOpMaker : public framework::OpProtoAndCheckerMaker { - public: - FetchOpMaker(framework::OpProto* proto, framework::OpAttrChecker* op_checker) - : OpProtoAndCheckerMaker(proto, op_checker) { - AddAttr("dataType", "output data type") - .SetDefault(framework::DataType::FP32); - AddAttr("col", "The col in global fetch variable").SetDefault(0); - AddInput("Input", "The output of fetch op."); - AddComment(R"DOC(Fetch data to global fetch variable)DOC"); + auto col = static_cast(Attr("col")); + + auto *fetch_list = out_var->GetMutable(); + auto &src_item = fetch_var->Get(); + + if (col >= fetch_list->size()) { + fetch_list->resize(col + 1); + } + auto &dst_item = fetch_list->at(col); + + // FIXME(yuyang18): Should we assume the fetch operator always generate + // CPU outputs? + dst_item.CopyFromTensor(src_item, platform::CPUPlace(), dev_ctx); } }; } // namespace operators } // namespace paddle -namespace ops = paddle::operators; -REGISTER_OP_WITHOUT_GRADIENT(fetch, ops::FetchOp, ops::FetchOpMaker); -REGISTER_OP_CPU_KERNEL(fetch, ops::FetchKernel); +// We do not need to register OpInfoMaker, +// since fetch operator will not be used by end users directly +REGISTER_OPERATOR(fetch, paddle::operators::FetchOp, + paddle::framework::EmptyGradOpMaker); diff --git a/paddle/operators/fetch_op.cu b/paddle/operators/fetch_op.cu deleted file mode 100644 index ca39d24c79..0000000000 --- a/paddle/operators/fetch_op.cu +++ /dev/null @@ -1,18 +0,0 @@ -/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. - -Licensed under the Apache License, Version 2.0 (the "License"); -you may not use this file except in compliance with the License. -You may obtain a copy of the License at - -http://www.apache.org/licenses/LICENSE-2.0 - -Unless required by applicable law or agreed to in writing, software -distributed under the License is distributed on an "AS IS" BASIS, -WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -See the License for the specific language governing permissions and -limitations under the License. */ - -#include "paddle/operators/fetch_op.h" - -namespace ops = paddle::operators; -REGISTER_OP_GPU_KERNEL(fetch, ops::FetchKernel); diff --git a/paddle/operators/fetch_op.h b/paddle/operators/fetch_op.h deleted file mode 100644 index b2a6e95875..0000000000 --- a/paddle/operators/fetch_op.h +++ /dev/null @@ -1,45 +0,0 @@ -/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. - -Licensed under the Apache License, Version 2.0 (the "License"); -you may not use this file except in compliance with the License. -You may obtain a copy of the License at - - http://www.apache.org/licenses/LICENSE-2.0 - -Unless required by applicable law or agreed to in writing, software -distributed under the License is distributed on an "AS IS" BASIS, -WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -See the License for the specific language governing permissions and -limitations under the License. */ - -#pragma once -#include "paddle/framework/eigen.h" -#include "paddle/framework/op_registry.h" - -namespace paddle { -namespace operators { - -template -class FetchKernel : public framework::OpKernel { - public: - void Compute(const framework::ExecutionContext& ctx) const override { - const framework::Tensor* input = ctx.Input("Input"); - framework::Variable* g_fetch_variable = - framework::GetGlobalScope().FindVar("fetch_value"); - auto* tensors = - g_fetch_variable->GetMutable>(); - int col = ctx.template Attr("col"); - if (tensors->size() < static_cast(col + 1)) { - tensors->resize(col + 1); - } - PADDLE_ENFORCE_GT(tensors->size(), static_cast(col)); - (*tensors)[col].Resize(input->dims()); - (*tensors)[col].mutable_data(platform::CPUPlace()); - (*tensors)[col].CopyFrom(*input, platform::CPUPlace(), - ctx.device_context()); - // TODO(qijun): need to handle LodTensor later - } -}; - -} // namespace operators -} // namespace paddle From 5993497cb746ccd5e6d99b1605e5c062484b63dd Mon Sep 17 00:00:00 2001 From: qijun Date: Mon, 16 Oct 2017 10:04:51 -0700 Subject: [PATCH 21/36] fix merge conflict --- paddle/operators/sequence_pool_op.h | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/paddle/operators/sequence_pool_op.h b/paddle/operators/sequence_pool_op.h index ce68204d41..a5569d1aac 100644 --- a/paddle/operators/sequence_pool_op.h +++ b/paddle/operators/sequence_pool_op.h @@ -111,7 +111,8 @@ class SequencePoolGradKernel : public framework::OpKernel { in_g->mutable_data(context.GetPlace()); if (strategy == LAST || strategy == FIRST) { // set X@Grad be zero at first when strategy is LAST/FIRST - math::SetConstant(context.device_context(), in_g, 0); + math::SetConstant functor; + functor(context.device_context(), in_g, 0); } auto place = context.GetEigenDevice(); for (int i = 0; i < static_cast(lod.size()) - 1; ++i) { From c87e060c187b29cfcf1272ddbaf741915ffcdee4 Mon Sep 17 00:00:00 2001 From: qijun Date: Mon, 16 Oct 2017 10:40:42 -0700 Subject: [PATCH 22/36] export feed/fetch method to Python --- paddle/framework/feed_fetch_method.h | 50 ++++++++++++++++ paddle/framework/scope.h | 60 ------------------- paddle/pybind/pybind.cc | 16 +++-- .../framework/tests/test_feed_fetch_method.py | 50 +++++----------- 4 files changed, 72 insertions(+), 104 deletions(-) create mode 100644 paddle/framework/feed_fetch_method.h diff --git a/paddle/framework/feed_fetch_method.h b/paddle/framework/feed_fetch_method.h new file mode 100644 index 0000000000..be96dc3267 --- /dev/null +++ b/paddle/framework/feed_fetch_method.h @@ -0,0 +1,50 @@ +/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. + +Licensed under the Apache License, Version 2.0 (the "License"); +you may not use this file except in compliance with the License. +You may obtain a copy of the License at + + http://www.apache.org/licenses/LICENSE-2.0 + +Unless required by applicable law or agreed to in writing, software +distributed under the License is distributed on an "AS IS" BASIS, +WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +See the License for the specific language governing permissions and +limitations under the License. */ + +#pragma once +#include "paddle/framework/scope.h" +#include "paddle/framework/variable.h" + +namespace paddle { +namespace framework { + +template +void SetFeedVariable(const LoDTensor& input, const std::string& var_name, + size_t index) { + // If var_name Variable is not found in GlobalScope, a new variable will + // be created. + Variable* g_feed_value = GetGlobalScope().Var(var_name); + auto& feed_inputs = + *(g_feed_value->GetMutable>()); + if (index >= feed_inputs.size()) { + feed_inputs.resize(index + 1); + } + // shared data with input tensor + feed_inputs[index].ShareDataWith(input); + // set lod + feed_inputs[index].set_lod(input.lod()); +} + +LoDTensor& GetFetchVariable(const std::string& var_name, size_t index) { + // If var_name Variable is not found in GlobalScope, a new variable will + // be created. + Variable* g_fetch_value = GetGlobalScope().Var(var_name); + auto& fetch_outputs = + *(g_fetch_value->GetMutable>()); + PADDLE_ENFORCE_LT(index, fetch_outputs.size()); + return fetch_outputs[index]; +} + +} // namespace framework +} // namespace paddle diff --git a/paddle/framework/scope.h b/paddle/framework/scope.h index 7376245b53..78eca28009 100644 --- a/paddle/framework/scope.h +++ b/paddle/framework/scope.h @@ -77,65 +77,5 @@ class Scope { framework::Scope& GetGlobalScope(); -// template -// void SetFeedVariable(const std::vector& input, const Lod& lod, -// const std::vector& dims, -// const std::string& var_name, size_t index) { -// Variable* g_feed_value = GetGlobalScope().Var("var_name"); -// // feed variable holds vector -// auto& feed_inputs = -// *(g_feed_value->GetMutable< -// std::vector>()); -// if (index >= feed_inputs.size()) { -// feed_inputs.resize(index); -// } -// // copy tensor -// T* dst = feed_inputs[index].mutable_data(make_ddim(dims), -// platform::CPUPlace()); -// memcpy(dst, inputs[i].data(), inputs[i].size() * sizeof(T)); -// // copy lod -// feed_inputs[index].set_lod(lod); -// } - -template -void SetFeedVariable(const LoDTensor& input, const std::string& var_name, - size_t index) { - std::cout << "into SetFeedVariable" << std::endl; - std::cout << var_name << std::endl; - std::cout << index << std::endl; - Variable* g_feed_value = GetGlobalScope().Var(var_name); - auto& feed_inputs = - *(g_feed_value->GetMutable>()); - if (index >= feed_inputs.size()) { - feed_inputs.resize(index + 1); - } - // shared data with input tensor - feed_inputs[index].ShareDataWith(input); - // set lod - feed_inputs[index].set_lod(input.lod()); -} - -// template -// std::vector GetFetchVariable(const std::string& var_name, size_t index) { -// Variable* g_fetch_value = GetGlobalScope().Var(var_name); -// auto& fetch_outputs = -// *(g_fetch_value->GetMutable< -// std::vector>()); -// std::vector result; -// result.resize(fetch_outputs[index].numel()); -// memcpy(result.data(), fetch_outputs[i].data(), -// fetch_outputs[i].numel() * sizeof(T)); -// } - -template -LoDTensor& GetFetchVariable(const std::string& var_name, size_t index) { - Variable* g_fetch_value = GetGlobalScope().Var(var_name); - auto& fetch_outputs = - *(g_fetch_value->GetMutable>()); - std::cout << "into GetFetchVariable" << std::endl; - PADDLE_ENFORCE_LT(index, fetch_outputs.size()); - return fetch_outputs[index]; -} - } // namespace framework } // namespace paddle diff --git a/paddle/pybind/pybind.cc b/paddle/pybind/pybind.cc index 1c8d2cfd61..983f19b30d 100644 --- a/paddle/pybind/pybind.cc +++ b/paddle/pybind/pybind.cc @@ -12,10 +12,10 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the License for the specific language governing permissions and limitations under the License. */ -#include "paddle/pybind/protobuf.h" - +#include "paddle/pybind/pybind.h" #include "paddle/framework/backward.h" #include "paddle/framework/executor.h" +#include "paddle/framework/feed_fetch_method.h" #include "paddle/framework/lod_tensor.h" #include "paddle/framework/tensor_array.h" #include "paddle/operators/cond_op.h" @@ -25,7 +25,7 @@ limitations under the License. */ #include "paddle/platform/enforce.h" #include "paddle/platform/place.h" #include "paddle/pybind/exception.h" -#include "paddle/pybind/pybind.h" +#include "paddle/pybind/protobuf.h" #include "paddle/pybind/tensor_py.h" #include "paddle/string/to_string.h" @@ -403,12 +403,10 @@ All parameter, weight, gradient are variables in Paddle. m.def("unique_integer", UniqueIntegerGenerator); m.def("is_compile_gpu", IsCompileGPU); - m.def("set_feed_variable", SetFeedVariable); - // m.def("set_feed_variable", SetFeedVariable); - // m.def("set_feed_variable", SetFeedVariable); - m.def("get_fetch_variable", GetFetchVariable); - // m.def("get_fetch_variable", GetFetchVariable); - // m.def("get_fetch_variable", GetFetchVariable); + m.def("set_feed_variable_float", framework::SetFeedVariable); + m.def("set_feed_variable_double", framework::SetFeedVariable); + m.def("set_feed_variable_int", framework::SetFeedVariable); + m.def("get_fetch_variable", framework::GetFetchVariable); BindProgramDesc(m); BindBlockDesc(m); diff --git a/python/paddle/v2/framework/tests/test_feed_fetch_method.py b/python/paddle/v2/framework/tests/test_feed_fetch_method.py index cd8eeee68f..47eedddcb6 100644 --- a/python/paddle/v2/framework/tests/test_feed_fetch_method.py +++ b/python/paddle/v2/framework/tests/test_feed_fetch_method.py @@ -2,49 +2,29 @@ import paddle.v2.framework.core as core import unittest import numpy as np -# class TestFeedFetch(unittest.TestCase): -# def test_feed_fetch(self): -# place = core.CPUPlace() -# input_tensor = core.LoDTensor([[0, 2, 4]]) -# input_tensor.set_dims([4, 4, 6]) -# input_tensor.alloc_int(place) -# input_array = np.array(input_tensor) -# input_array[0, 0, 0] = 3 -# input_array[3, 3, 5] = 10 -# input_tensor.set(input_array, place) - -# core.set_feed_variable(input_tensor, "feed", 0) - -# output_tensor = core.get_fetch_variable("feed", 0) -# print type(output_tensor) - -# output_lod = output_tensor.lod() -# print type(output_lod) -# print output_lod[0] -# print output_lod[0][0] -# print output_lod[0][1] -# print output_lod[0][2] -# # self.assertEqual(0, output_lod[0][0]) -# # self.assertEqual(0, output_lod[0][0]) -# # self.assertEqual(2, output_lod[0][1]) -# # self.assertEqual(4, output_lod[0][2]) - -# # output_array = np.array(output_tensor) -# # self.assertEqual(3, output_array[0, 0, 0]) -# # self.assertEqual(10, output_array[3, 3, 5]); - class TestFeedFetch(unittest.TestCase): def test_feed_fetch(self): place = core.CPUPlace() - input_tensor = core.LoDTensor([[0, 2, 4]]) - input_tensor.set_dims([4, 4, 6]) - input_tensor.alloc_float(place) - input_array = np.array(input_tensor) + input_array = np.ones((4, 4, 6)).astype("float32") input_array[0, 0, 0] = 3 input_array[3, 3, 5] = 10 + input_tensor = core.LoDTensor([[0, 2, 4]]) input_tensor.set(input_array, place) + core.set_feed_variable_float(input_tensor, "feed", 0) + + output_tensor = core.get_fetch_variable("feed", 0) + + output_lod = output_tensor.lod() + self.assertEqual(0, output_lod[0][0]) + self.assertEqual(2, output_lod[0][1]) + self.assertEqual(4, output_lod[0][2]) + + output_array = np.array(output_tensor) + self.assertEqual(3, output_array[0, 0, 0]) + self.assertEqual(10, output_array[3, 3, 5]) + if __name__ == "__main__": unittest.main() From 701c90f6f74f2907baaa8d580e64b51b2db3fba0 Mon Sep 17 00:00:00 2001 From: qijun Date: Mon, 16 Oct 2017 10:47:08 -0700 Subject: [PATCH 23/36] remove unused header file --- paddle/framework/scope.h | 2 -- 1 file changed, 2 deletions(-) diff --git a/paddle/framework/scope.h b/paddle/framework/scope.h index 78eca28009..a7fce3514b 100644 --- a/paddle/framework/scope.h +++ b/paddle/framework/scope.h @@ -18,10 +18,8 @@ limitations under the License. */ #include #include -#include "paddle/framework/lod_tensor.h" #include "paddle/framework/variable.h" #include "paddle/platform/macros.h" -#include "paddle/platform/place.h" namespace paddle { namespace framework { From 5fe3f8f630cf4459a30a399d63f27db7b474ca8e Mon Sep 17 00:00:00 2001 From: qijun Date: Mon, 16 Oct 2017 10:50:20 -0700 Subject: [PATCH 24/36] refine code --- paddle/framework/feed_fetch_method.h | 6 +++--- paddle/pybind/pybind.cc | 5 +++-- 2 files changed, 6 insertions(+), 5 deletions(-) diff --git a/paddle/framework/feed_fetch_method.h b/paddle/framework/feed_fetch_method.h index be96dc3267..826d180bfc 100644 --- a/paddle/framework/feed_fetch_method.h +++ b/paddle/framework/feed_fetch_method.h @@ -37,9 +37,9 @@ void SetFeedVariable(const LoDTensor& input, const std::string& var_name, } LoDTensor& GetFetchVariable(const std::string& var_name, size_t index) { - // If var_name Variable is not found in GlobalScope, a new variable will - // be created. - Variable* g_fetch_value = GetGlobalScope().Var(var_name); + // Since we want to fetch LodTensor from a variable, the variable must + // be created alreadly. + Variable* g_fetch_value = GetGlobalScope().FindVar(var_name); auto& fetch_outputs = *(g_fetch_value->GetMutable>()); PADDLE_ENFORCE_LT(index, fetch_outputs.size()); diff --git a/paddle/pybind/pybind.cc b/paddle/pybind/pybind.cc index 983f19b30d..008a9441db 100644 --- a/paddle/pybind/pybind.cc +++ b/paddle/pybind/pybind.cc @@ -12,7 +12,8 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the License for the specific language governing permissions and limitations under the License. */ -#include "paddle/pybind/pybind.h" +#include "paddle/pybind/protobuf.h" + #include "paddle/framework/backward.h" #include "paddle/framework/executor.h" #include "paddle/framework/feed_fetch_method.h" @@ -25,7 +26,7 @@ limitations under the License. */ #include "paddle/platform/enforce.h" #include "paddle/platform/place.h" #include "paddle/pybind/exception.h" -#include "paddle/pybind/protobuf.h" +#include "paddle/pybind/pybind.h" #include "paddle/pybind/tensor_py.h" #include "paddle/string/to_string.h" From 3e427441910bbccddf0222d4bd05ecfdea9302e7 Mon Sep 17 00:00:00 2001 From: "Yang Yang(Tony)" Date: Mon, 16 Oct 2017 11:15:10 -0700 Subject: [PATCH 25/36] update based on review --- doc/design/executor.md | 29 +---------------------------- 1 file changed, 1 insertion(+), 28 deletions(-) diff --git a/doc/design/executor.md b/doc/design/executor.md index ab2d6c3558..b5fb6c5c3c 100644 --- a/doc/design/executor.md +++ b/doc/design/executor.md @@ -20,31 +20,4 @@ It does not do graph partitioning, meaning dividing the `ProgramDesc` into sever ## Implementation -`Executor` evaluates a `ProgramDesc`. Essentially, it instantiates Variables and Operators, then run all the operators - -```c++ -void Executor::Run(const ProgramDesc& pdesc, Scope* scope, int block_id) { - auto& block = pdesc.blocks(block_id); - auto& device = device_contexts_[0]; - - // Instantiate all the vars in the global scope - for (auto& var : block.vars()) { - scope->NewVar(var.name()); - } - - // Run the block - Scope& local_scope = scope->NewScope(); - for (auto& op_desc : block.ops()) { - for (auto& var : op_desc.outputs()) { - for (auto& argu : var.arguments()) { - // Create temp variable in the local_scope - if (local_scope.FindVar(argu) == nullptr) { - local_scope.NewVar(argu); - } - } - } - auto op = paddle::framework::OpRegistry::CreateOp(op_desc); - op->Run(local_scope, *device); - } -} -``` +`Executor` evaluates a `ProgramDesc`. Essentially, it instantiates Variables and Operators, then run all the operators in sequence. [[code]](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/framework/executor.cc) From 219f46ae16fa3f3e74f28bcbf1b7f815b9b5ac92 Mon Sep 17 00:00:00 2001 From: qijun Date: Mon, 16 Oct 2017 11:56:15 -0700 Subject: [PATCH 26/36] export SelectedRows to Python --- paddle/pybind/pybind.cc | 17 +++++++++ .../v2/framework/tests/test_selected_rows.py | 37 +++++++++++++++++++ 2 files changed, 54 insertions(+) create mode 100644 python/paddle/v2/framework/tests/test_selected_rows.py diff --git a/paddle/pybind/pybind.cc b/paddle/pybind/pybind.cc index afc80b25b1..23e76011c9 100644 --- a/paddle/pybind/pybind.cc +++ b/paddle/pybind/pybind.cc @@ -17,6 +17,7 @@ limitations under the License. */ #include "paddle/framework/backward.h" #include "paddle/framework/executor.h" #include "paddle/framework/lod_tensor.h" +#include "paddle/framework/selected_rows.h" #include "paddle/framework/tensor_array.h" #include "paddle/operators/cond_op.h" #include "paddle/operators/dynamic_recurrent_op.h" @@ -138,6 +139,22 @@ PYBIND11_PLUGIN(core) { #endif }); + py::class_(m, "SelectedRows") + .def("__init__", + [](SelectedRows &instance) { new (&instance) SelectedRows(); }) + .def("__init__", + [](SelectedRows &instance, const std::vector rows, + const int64_t &height) { + new (&instance) SelectedRows(rows, height); + }) + .def("get_tensor", + [](SelectedRows &self) { return self.mutable_value(); }, + py::return_value_policy::reference) + .def("set_height", &SelectedRows::set_height) + .def("height", &SelectedRows::height) + .def("set_rows", &SelectedRows::set_rows) + .def("rows", &SelectedRows::rows, py::return_value_policy::reference); + py::class_(m, "Variable", R"DOC(Variable Class. All parameter, weight, gradient are variables in Paddle. diff --git a/python/paddle/v2/framework/tests/test_selected_rows.py b/python/paddle/v2/framework/tests/test_selected_rows.py new file mode 100644 index 0000000000..661e818179 --- /dev/null +++ b/python/paddle/v2/framework/tests/test_selected_rows.py @@ -0,0 +1,37 @@ +import paddle.v2.framework.core as core +import unittest +import numpy as np + + +class TestSelectedRows(unittest.TestCase): + def test_selected_rows(self): + place = core.CPUPlace() + height = 10 + rows = [0, 4, 7] + row_numel = 10 + selcted_rows = core.SelectedRows(rows, row_numel) + np_array = np.ones((len(rows), height)).astype("float32") + np_array[0, 0] = 2.0 + np_array[2, 8] = 4.0 + tensor = selcted_rows.get_tensor() + tensor.set(np_array, place) + + # compare rows + self.assertEqual(0, selcted_rows.rows()[0]) + self.assertEqual(4, selcted_rows.rows()[1]) + self.assertEqual(7, selcted_rows.rows()[2]) + + # compare height + self.assertEqual(10, selcted_rows.height()) + + # compare tensor + self.assertAlmostEqual(2.0, + selcted_rows.get_tensor().get_float_element(0)) + self.assertAlmostEqual(1.0, + selcted_rows.get_tensor().get_float_element(1)) + self.assertAlmostEqual( + 4.0, selcted_rows.get_tensor().get_float_element(2 * row_numel + 8)) + + +if __name__ == "__main__": + unittest.main() From eb27c735cb95accf5abdf97fb35b8e7c39519856 Mon Sep 17 00:00:00 2001 From: qijun Date: Mon, 16 Oct 2017 13:35:28 -0700 Subject: [PATCH 27/36] remove unused C++ class OpRegistrar --- paddle/framework/op_registry.h | 10 ---------- 1 file changed, 10 deletions(-) diff --git a/paddle/framework/op_registry.h b/paddle/framework/op_registry.h index 226e8ddcd4..0bda87dfa1 100644 --- a/paddle/framework/op_registry.h +++ b/paddle/framework/op_registry.h @@ -82,16 +82,6 @@ class OpRegistry { static std::unique_ptr CreateOp(const OpDescBind& op_desc); }; -template -class OpRegistrar : public Registrar { - public: - explicit OpRegistrar(const char* op_type) { OpRegistrar(op_type, ""); } - OpRegistrar(const char* op_type, const char* grad_op_type) { - OpRegistry::RegisterOp(op_type, - grad_op_type); - } -}; - template struct OpKernelRegistrarFunctor; From 186d1655650c8e07615609c5f299ffc15512ad74 Mon Sep 17 00:00:00 2001 From: Yu Yang Date: Mon, 16 Oct 2017 14:53:05 -0700 Subject: [PATCH 28/36] AttributeChecker Better error log and speicalize bool (#4840) * AttributeChecker Better error log and speicalize bool Since lots of types can be cast to bool * add FIXME comment --- paddle/framework/attribute.h | 56 ++++++++++++++++++++++++++++++++++-- 1 file changed, 54 insertions(+), 2 deletions(-) diff --git a/paddle/framework/attribute.h b/paddle/framework/attribute.h index d13530e340..8a7a949346 100644 --- a/paddle/framework/attribute.h +++ b/paddle/framework/attribute.h @@ -120,6 +120,57 @@ class EnumInContainer { std::unordered_set container_; }; +template +struct ExtractAttribute { + explicit ExtractAttribute(const std::string& attr_name) + : attr_name_(attr_name) {} + + T* operator()(Attribute& attr) const { + T* attr_value = nullptr; + try { + attr_value = &boost::get(attr); + } catch (boost::bad_get& bad_get) { + PADDLE_THROW("Cannot get attribute %s by type %s, its type is %s", + attr_name_, typeid(T).name(), attr.type().name()); + } + return attr_value; + } + + const std::string& attr_name_; +}; + +// special handle bool +// FIXME(yuyang18): Currently we cast bool into int in python binding. It is +// hard to change the logic there. In another way, we should correct handle +// if the user set `some_flag=1`. +// +// FIX ME anytime if there is a better solution. +template <> +struct ExtractAttribute { + explicit ExtractAttribute(const std::string& attr_name) + : attr_name_(attr_name) {} + + bool* operator()(Attribute& attr) const { + if (attr.type() == typeid(int)) { // NOLINT + int val = boost::get(attr); + attr = static_cast(val); + } else if (attr.type() == typeid(float)) { // NOLINT + float val = boost::get(attr); + attr = static_cast(val); + } + bool* attr_value = nullptr; + try { + attr_value = &boost::get(attr); + } catch (boost::bad_get& bad_get) { + PADDLE_THROW("Cannot get attribute %s by type bool, its type is %s", + attr_name_, attr.type().name()); + } + return attr_value; + } + + const std::string& attr_name_; +}; + // check whether a certain attribute fit its limits // an attribute can have more than one limits template @@ -171,9 +222,10 @@ class TypedAttrChecker { attr_map[attr_name_] = val; } Attribute& attr = attr_map.at(attr_name_); - T& attr_value = boost::get(attr); + ExtractAttribute extract_attr(attr_name_); + T* attr_value = extract_attr(attr); for (const auto& checker : value_checkers_) { - checker(attr_value); + checker(*attr_value); } } From 790b9ce4c5532d5e10e2911581b9358554e5de29 Mon Sep 17 00:00:00 2001 From: Yu Yang Date: Mon, 16 Oct 2017 15:01:05 -0700 Subject: [PATCH 29/36] Update design doc for Python Layer (#4698) * Update design doc for Python Layer * Update document --- doc/design/python_api.md | 114 ++++++++++++++++++++++++++++++--------- 1 file changed, 89 insertions(+), 25 deletions(-) diff --git a/doc/design/python_api.md b/doc/design/python_api.md index 56ae1d925a..cb5fdc765b 100644 --- a/doc/design/python_api.md +++ b/doc/design/python_api.md @@ -179,40 +179,104 @@ init_attr={ `optimize_op_attrs` is not in the `VarDesc` message, but kept in the Python instance, as it will be used in the Python space when creating the optimize operator's `OpDesc`, and will be in the `OpDesc` message. -## Layer Functions +## Layer Function -A layer is a Python function that creates some operators and variables. Layers simplify the work of application programmers. +A layer is a Python function that creates some operators and variables. Layers simplify the work of application programmers. -### Data Layer +Layer functions take `Variable` and configuration parameters as its input and return the output variable(s). + +For example, `FullyConnected` take one or more variable as its input. The input could be input data or another layer's output. There are many configuration options for a `FullyConnected` layer, such as layer size, activation, parameter names, initialization strategies of parameters, and so on. The `FullyConnected` layer will return an output variable. + + +### Necessity for reusing code between layer functions + +There are a lot of code that can be reused. Such as + +* Give the default value of configuration. e.g., default initialize strategy for parameters is uniform random with `min = -1.0`, `max = 1.0`. and default initialize strategy for bias is to fill zero. +* Append the activation operator. +* Create a temporary variable. +* Create parameter. +* Generate a unique name. +* Add a bias. +* ... + +A mechanism to reuse code between layer functions is necessary. It will be around [150 lines of code](https://github.com/PaddlePaddle/Paddle/pull/4724/files#diff-823b27e07e93914ada859232ae23f846R12) if we write a `FullyConnected` layer without any helper functions. + + + +### Comparision between global functions and helper class + +The `FullyConnected` layer will be as follow when we provide global functions: ```python -def data_layer(name, type, column_name): - block = the_current_program.glolal_block() - var = block.create_global_var( - name=name, - shape=[None] + type.dims(), - dtype=type.dtype) - block.prepend_operator(block, - type="Feed", - inputs = None, - outputs = [var], - {column_name: column_name}) - return var +def fc_layer(input, size, param_attr=None, bias_attr=None, act=None, name=None): + if name is None: + name = unique_name("fc") + input = multiple_input(input) + param_attr = default_param_attr(param_attr) + param_attr = multiple_param_attr(param_attr, len(input)) + + # mul + mul_results = [] + for ipt, attr in zip(input, param_attr): + shape = ipt.shape[1:] + [size] + w = g_program.global_block().create_parameter(shape, ipt.dtype, name, attr) + tmp = create_tmp_var(name) + g_program.current_block().append_op("mul", {ipt, w}, {tmp}) + mul_results.append(tmp) + + # add sum + ... + # add bias + ... + # add activation + ... + return out ``` -The input to the feed operator is a special variable in the global scope, which is the output of [Python readers](https://github.com/PaddlePaddle/Paddle/blob/develop/doc/design/reader/README.md). +We can provide many helpers functions for layer developers. However, there are several disadvantages for global helper functions: + +1. We need a namespace for these methods, then layer developers can quickly figure out what method they can use. +2. Global functions will force layer developers to pass its parameter time by time. + +So we provide a helper class, `LayerHelper`, to share code between layer functions. The `FullyConnected` Layer will be as follow. + +```python +def fc_layer(input, size, param_attr=None, bias_attr=None, act=None, name=None): + helper = LayerHelper(locals()) # pass all parameter to LayerHelper + + mul_results = [] + for ipt, param in helper.iter_multiple_input_and_param(): + w = helper.create_parameter(shape=ipt.shape[1:] + [size], dtype = ipt.dtype) + tmp = helper.create_tmp_variable() + helper.append_op('mul', {ipt, w}, {tmp}) + mul_results.append(tmp) + + pre_bias = helper.add_sum(mul_results) + pre_activation = helper.add_bias(pre_bias) + return helper.add_activation(pre_activation) +``` + +We not only use the fewer lines of code to write `fc_layer` but also make the code clearer to understand. At the same time, layer developers can figure out what function they can invoke by typing `helper.` in a python editor. + + +### Implementation of layer helper -### FC Layer +We just keep all parameters of a layer function as a dictionary in layer helper as a private data member. Every method of layer helper will look up the dictionary after it is invoked. In that way, we can implement a layer helper for all layer functions even some layer does not contain some operator. For example, The `activation` is used by the FullyConnected layer or convolution layers, but a cross-entropy layer does not use it. The example code of `add_activation` are: ```python -def fc_layer(input, size, ...): - block = program.current_block() - w = block.create_parameter(...) - b = block.create_parameter(...) - out = block.create_var() - op = block.append_operator("FC", X=input, W=w, b=b, out=out) - out.writer = op - return out +class LayerHelper(object): + def __init__(self, **kwargs): # kwargs is short for `keyword arguments` + self.kwargs = kwargs + + def add_activation(self, input_var): + act = self.kwargs.get("act", None) # default value is None + if act is None: # do nothing if no act + return input_var + + tmp = self.create_tmp_var(self) + self.append_op(type=act, input=input_var, output=tmp) + return tmp ``` ## Optimizer From 75d0c7901530248afbcad7fc79c12241c1e4f00d Mon Sep 17 00:00:00 2001 From: Yu Yang Date: Mon, 16 Oct 2017 15:01:39 -0700 Subject: [PATCH 30/36] Change Name convention of operator attributes (#4807) * Change dataType to data_type Follow PEP8 * Change name_convention to fit PEP8 --- paddle/framework/backward.cc | 2 +- paddle/operators/cross_entropy_op.cc | 22 ++++++++-------- paddle/operators/cross_entropy_op.cu | 4 +-- paddle/operators/cross_entropy_op.h | 4 +-- paddle/operators/fill_constant_op.cc | 4 +-- paddle/operators/name_convention.md | 2 +- paddle/operators/pool_op.cc | 26 +++++++++---------- paddle/operators/pool_op.h | 8 +++--- paddle/operators/pool_with_index_op.cc | 18 ++++++------- paddle/operators/pool_with_index_op.h | 4 +-- .../softmax_with_cross_entropy_op.cc | 14 +++++----- .../softmax_with_cross_entropy_op.cu | 4 +-- .../operators/softmax_with_cross_entropy_op.h | 4 +-- .../framework/tests/test_cross_entropy_op.py | 4 +-- .../v2/framework/tests/test_pool2d_op.py | 4 +-- .../v2/framework/tests/test_pool3d_op.py | 4 +-- .../v2/framework/tests/test_pool_max_op.py | 2 +- .../test_softmax_with_cross_entropy_op.py | 2 +- 18 files changed, 66 insertions(+), 66 deletions(-) diff --git a/paddle/framework/backward.cc b/paddle/framework/backward.cc index e3d7dacd7f..c78e056071 100644 --- a/paddle/framework/backward.cc +++ b/paddle/framework/backward.cc @@ -433,7 +433,7 @@ ParamGradInfoMap AppendBackward( new OpDescBind("fill_constant", {}, {{"Out", {fill_one_op_out}}}, {{"shape", std::vector{1}}, {"value", static_cast(1.0)}, - {"dataType", framework::DataType::FP32}})); + {"data_type", framework::DataType::FP32}})); all_ops.push_back(std::move(fill_one_op)); size_t forward_op_num = all_ops.size(); size_t forward_block_num = program_desc.Size(); diff --git a/paddle/operators/cross_entropy_op.cc b/paddle/operators/cross_entropy_op.cc index 708e80e96a..923ae5be6c 100644 --- a/paddle/operators/cross_entropy_op.cc +++ b/paddle/operators/cross_entropy_op.cc @@ -34,13 +34,13 @@ class CrossEntropyOp : public framework::OperatorWithKernel { PADDLE_ENFORCE_EQ(x_dims[0], label_dims[0], "The 1st dimension of Input(X) and Input(Label) should " "be equal."); - if (ctx->Attrs().Get("softLabel")) { + if (ctx->Attrs().Get("soft_label")) { PADDLE_ENFORCE_EQ(x_dims[1], label_dims[1], - "If Attr(softLabel) == true, the 2nd dimension of " + "If Attr(soft_label) == true, the 2nd dimension of " "Input(X) and Input(Label) should be equal."); } else { PADDLE_ENFORCE_EQ(label_dims[1], 1, - "If Attr(softLabel) == false, the 2nd dimension of " + "If Attr(soft_label) == false, the 2nd dimension of " "Input(Label) should be 1."); } @@ -82,13 +82,13 @@ class CrossEntropyGradientOp : public framework::OperatorWithKernel { "be equal."); PADDLE_ENFORCE_EQ(dy_dims[1], 1, "The 2nd dimension of Input(Y@Grad) should be 1."); - if (ctx->Attrs().Get("softLabel")) { + if (ctx->Attrs().Get("soft_label")) { PADDLE_ENFORCE_EQ(x_dims[1], label_dims[1], - "When Attr(softLabel) == true, the 2nd dimension of " + "When Attr(soft_label) == true, the 2nd dimension of " "Input(X) and Input(Label) should be equal."); } else { PADDLE_ENFORCE_EQ(label_dims[1], 1, - "When Attr(softLabel) == false, the 2nd dimension of " + "When Attr(soft_label) == false, the 2nd dimension of " "Input(Label) should be 1."); } ctx->SetOutputDim(framework::GradVarName("X"), x_dims); @@ -115,15 +115,15 @@ class CrossEntropyOpMaker : public framework::OpProtoAndCheckerMaker { "Label", "(Tensor, default Tensor), the ground truth which is " "a 2-D tensor. " - "When softLabel is set to false, `Label` is a Tensor with shape " + "When soft_label is set to false, `Label` is a Tensor with shape " "[N x 1]. " - "When softLabel is set to true, `Label` is a Tensor " + "When soft_label is set to true, `Label` is a Tensor " "with shape [N x K]."); AddOutput("Y", "(Tensor, default Tensor), a 2-D tensor " "with shape [N x 1]. The cross entropy loss."); AddAttr( - "softLabel", + "soft_label", "(bool, default false), a flag to indicate whether to interpretate " "the given labels as soft labels.") .SetDefault(false); @@ -133,12 +133,12 @@ CrossEntropy Operator. It supports both standard cross-entropy and soft-label cross-entropy loss computation. 1) One-hot cross-entropy: - softLabel = false, Label[i, 0] indicates the class index for sample i: + soft_label = false, Label[i, 0] indicates the class index for sample i: Y[i] = -log(X[i, Label[i]]) 2) Soft-label cross-entropy: - softLabel = true, Label[i, j] indicates the soft label of class j + soft_label = true, Label[i, j] indicates the soft label of class j for sample i: Y[i] = \sum_j{-Label[i, j] * log(X[i, j])} diff --git a/paddle/operators/cross_entropy_op.cu b/paddle/operators/cross_entropy_op.cu index 07b0388b60..c492dddb09 100644 --- a/paddle/operators/cross_entropy_op.cu +++ b/paddle/operators/cross_entropy_op.cu @@ -56,7 +56,7 @@ class CrossEntropyOpCUDAKernel : public framework::OpKernel { y->mutable_data(ctx.GetPlace()); math::CrossEntropyFunctor()( - ctx.device_context(), y, x, label, ctx.Attr("softLabel")); + ctx.device_context(), y, x, label, ctx.Attr("soft_label")); } }; @@ -83,7 +83,7 @@ class CrossEntropyGradientOpCUDAKernel : public framework::OpKernel { int block = 512; int grid = (batch_size * class_num + block - 1) / block; - if (ctx.Attr("softLabel")) { + if (ctx.Attr("soft_label")) { auto* label_data = label->data(); SoftCrossEntropyGradientKernel<<< grid, block, 0, reinterpret_cast( diff --git a/paddle/operators/cross_entropy_op.h b/paddle/operators/cross_entropy_op.h index 19c276d23f..42f282103b 100644 --- a/paddle/operators/cross_entropy_op.h +++ b/paddle/operators/cross_entropy_op.h @@ -38,7 +38,7 @@ class CrossEntropyOpKernel : public framework::OpKernel { y->mutable_data(ctx.GetPlace()); math::CrossEntropyFunctor()( - ctx.device_context(), y, x, labels, ctx.Attr("softLabel")); + ctx.device_context(), y, x, labels, ctx.Attr("soft_label")); } }; @@ -55,7 +55,7 @@ class CrossEntropyGradientOpKernel : public framework::OpKernel { T* dx_data = dx->mutable_data(ctx.GetPlace()); int class_num = x->dims()[1]; - if (ctx.Attr("softLabel")) { + if (ctx.Attr("soft_label")) { auto x_mat = EigenMatrix::From(*x); auto dy_mat = EigenMatrix::From(*dy); auto lbl_mat = EigenMatrix::From(*label); diff --git a/paddle/operators/fill_constant_op.cc b/paddle/operators/fill_constant_op.cc index 65d03d5fa4..a56832e202 100644 --- a/paddle/operators/fill_constant_op.cc +++ b/paddle/operators/fill_constant_op.cc @@ -35,7 +35,7 @@ class FillConstantOp : public framework::OperatorWithKernel { framework::DataType IndicateDataType( const framework::ExecutionContext &ctx) const override { - return static_cast(ctx.Attr("dataType")); + return static_cast(ctx.Attr("data_type")); } }; @@ -44,7 +44,7 @@ class FillConstantOpMaker : public framework::OpProtoAndCheckerMaker { FillConstantOpMaker(framework::OpProto *proto, framework::OpAttrChecker *op_checker) : framework::OpProtoAndCheckerMaker(proto, op_checker) { - AddAttr("dataType", + AddAttr("data_type", "(int, default 5 (FP32)) " "Output data type") .SetDefault(framework::DataType::FP32); diff --git a/paddle/operators/name_convention.md b/paddle/operators/name_convention.md index 379385dc5d..5a21690795 100644 --- a/paddle/operators/name_convention.md +++ b/paddle/operators/name_convention.md @@ -11,7 +11,7 @@ When defining an operator in Paddle, a corresponding [OpProtoMaker](https://gith - If an operator's Input/Output are tensors in math, not match to any meaningful words, input name should starts from `X`. e.g. `X`, `Y`, and output name should starts from `Out`. e.g. `Out`. This rule intends making operators which have few inputs/outputs unified. - Attribute. - - Attribute name follows the **camelCase**. e.g. `x`, `y`, `axis`, `rowwiseMatrix`. Also, attribute name prefers to meaningful English words. + - Attribute name follows the **snake_case**. e.g. `x`, `y`, `axis`, `rowwise_matrix`. Also, attribute name prefers to meaningful English words. - Comments. - Input/Output/Attr comment follow the format of **(type,default value) usage**, corresponding to which type it can be and how it will be used in the operator. e.g. Attribute in Accumulator`"gamma" `,`(float, default 1.0) Accumulation multiplier`. diff --git a/paddle/operators/pool_op.cc b/paddle/operators/pool_op.cc index c6d9aae133..a326839c0f 100644 --- a/paddle/operators/pool_op.cc +++ b/paddle/operators/pool_op.cc @@ -29,7 +29,7 @@ void PoolOp::InferShape(framework::InferShapeContext *ctx) const { auto in_x_dims = ctx->GetInputDim("X"); - std::string pooling_type = ctx->Attrs().Get("poolingType"); + std::string pooling_type = ctx->Attrs().Get("pooling_type"); std::vector ksize = ctx->Attrs().Get>("ksize"); std::vector strides = ctx->Attrs().Get>("strides"); std::vector paddings = ctx->Attrs().Get>("paddings"); @@ -37,7 +37,7 @@ void PoolOp::InferShape(framework::InferShapeContext *ctx) const { PADDLE_ENFORCE(in_x_dims.size() == 4 || in_x_dims.size() == 5, "Pooling intput should be 4-D or 5-D tensor."); - if (ctx->Attrs().Get("globalPooling")) { + if (ctx->Attrs().Get("global_pooling")) { ksize.resize(static_cast(in_x_dims.size()) - 2); for (size_t i = 0; i < ksize.size(); ++i) ksize[i] = static_cast(in_x_dims[i + 2]); @@ -80,23 +80,23 @@ Pool2dOpMaker::Pool2dOpMaker(framework::OpProto *proto, "the number of channels, H and W is the height and " "width of feature."); - AddAttr("poolingType", - "PoolingType of pooling operator." + AddAttr("pooling_type", + "Pooling_type of pooling operator." "Str constant equal to 'max' or 'avg'.") .InEnum({"max", "avg"}); AddAttr>( "ksize", "The pooling window size(height, width) of pooling operator." - "If globalPooling = true, ksize is ignored and need not be " + "If global_pooling = true, ksize is ignored and need not be " "specified."); // TODO(Chengduo): Add checker. (Currently, // TypedAttrChecker don't support vector type.) AddAttr( - "globalPooling", - "Whether to use the globalPooling." + "global_pooling", + "Whether to use the global_pooling." "Bool constant equal to false or true." "Default false." - "If globalPooling = true, ksize is ignored and need not be specified.") + "If global_pooling = true, ksize is ignored and need not be specified.") .SetDefault(false); AddAttr>("strides", "The strides(height, width) of pooling window." @@ -146,7 +146,7 @@ Pool3dOpMaker::Pool3dOpMaker(framework::OpProto *proto, "the number of channels, D, H and W is the depth, height and " "width of feature."); - AddAttr("poolingType", + AddAttr("pooling_type", "PoolingType of pooling operator." "Str constant equal to 'max' or 'avg'.") .InEnum({"max", "avg"}); @@ -154,15 +154,15 @@ Pool3dOpMaker::Pool3dOpMaker(framework::OpProto *proto, AddAttr>( "ksize", "The pooling window size(depth, height, width) of pooling operator." - "If globalPooling = true, ksize is ignored and need not be " + "If global_pooling = true, ksize is ignored and need not be " "specified."); // TODO(Chengduo): Add checker. (Currently, // TypedAttrChecker don't support vector type.) AddAttr( - "globalPooling", - "Whether to use the globalPooling." + "global_pooling", + "Whether to use the global_pooling." "Bool constant equal to false or true." "Default false." - "If globalPooling = true, ksize is ignored and need not be specified.") + "If global_pooling = true, ksize is ignored and need not be specified.") .SetDefault(false); AddAttr>("strides", "Strides(depth, height, width) of pooling operator." diff --git a/paddle/operators/pool_op.h b/paddle/operators/pool_op.h index e5016d573d..f6169e05b3 100644 --- a/paddle/operators/pool_op.h +++ b/paddle/operators/pool_op.h @@ -59,11 +59,11 @@ class PoolKernel : public framework::OpKernel { const Tensor* in_x = context.Input("X"); Tensor* out = context.Output("Out"); - std::string pooling_type = context.Attr("poolingType"); + std::string pooling_type = context.Attr("pooling_type"); std::vector ksize = context.Attr>("ksize"); std::vector strides = context.Attr>("strides"); std::vector paddings = context.Attr>("paddings"); - if (context.Attr("globalPooling")) { + if (context.Attr("global_pooling")) { for (size_t i = 0; i < ksize.size(); ++i) { ksize[i] = static_cast(in_x->dims()[i + 2]); } @@ -119,12 +119,12 @@ class PoolGradKernel : public framework::OpKernel { context.Input(framework::GradVarName("Out")); Tensor* in_x_grad = context.Output(framework::GradVarName("X")); - std::string pooling_type = context.Attr("poolingType"); + std::string pooling_type = context.Attr("pooling_type"); std::vector ksize = context.Attr>("ksize"); std::vector strides = context.Attr>("strides"); std::vector paddings = context.Attr>("paddings"); - if (context.Attr("globalPooling")) { + if (context.Attr("global_pooling")) { for (size_t i = 0; i < ksize.size(); ++i) ksize[i] = static_cast(in_x->dims()[i + 2]); } diff --git a/paddle/operators/pool_with_index_op.cc b/paddle/operators/pool_with_index_op.cc index 005ee88693..8eee20c2a9 100644 --- a/paddle/operators/pool_with_index_op.cc +++ b/paddle/operators/pool_with_index_op.cc @@ -45,7 +45,7 @@ class MaxPoolWithIndexOp : public framework::OperatorWithKernel { PADDLE_ENFORCE(in_x_dims.size() == 4 || in_x_dims.size() == 5, "Pooling intput should be 4-D or 5-D tensor."); - if (ctx->Attrs().Get("globalPooling")) { + if (ctx->Attrs().Get("global_pooling")) { ksize.resize(static_cast(in_x_dims.size()) - 2); for (size_t i = 0; i < ksize.size(); ++i) ksize[i] = static_cast(in_x_dims[i + 2]); @@ -108,15 +108,15 @@ class MaxPool2dWithIndexOpMaker : public framework::OpProtoAndCheckerMaker { AddAttr>( "ksize", "The pooling window size(height, width) of pooling operator." - "If globalPooling = true, ksize is ignored and need not be " + "If global_pooling = true, ksize is ignored and need not be " "specified."); // TODO(Chengduo): Add checker. (Currently, // TypedAttrChecker don't support vector type.) AddAttr( - "globalPooling", - "Whether to use the globalPooling." + "global_pooling", + "Whether to use the global_pooling." "Bool constant equal to false or true." "Default false." - "If globalPooling = true, ksize is ignored and need not be specified.") + "If global_pooling = true, ksize is ignored and need not be specified.") .SetDefault(false); AddAttr>("strides", "The strides(height, width) of pooling window." @@ -179,15 +179,15 @@ class MaxPool3dWithIndexOpMaker : public framework::OpProtoAndCheckerMaker { AddAttr>( "ksize", "The pooling window size(depth, height, width) of pooling operator." - "If globalPooling = true, ksize is ignored and need not be " + "If global_pooling = true, ksize is ignored and need not be " "specified."); // TODO(Chengduo): Add checker. (Currently, // TypedAttrChecker don't support vector type.) AddAttr( - "globalPooling", - "Whether to use the globalPooling." + "global_pooling", + "Whether to use the global_pooling." "Bool constant equal to false or true." "Default false." - "If globalPooling = true, ksize is ignored and need not be specified.") + "If global_pooling = true, ksize is ignored and need not be specified.") .SetDefault(false); AddAttr>( "strides", diff --git a/paddle/operators/pool_with_index_op.h b/paddle/operators/pool_with_index_op.h index 01b961ca82..455c453efc 100644 --- a/paddle/operators/pool_with_index_op.h +++ b/paddle/operators/pool_with_index_op.h @@ -35,7 +35,7 @@ class MaxPoolWithIndexKernel : public framework::OpKernel { std::vector ksize = context.Attr>("ksize"); std::vector strides = context.Attr>("strides"); std::vector paddings = context.Attr>("paddings"); - if (context.Attr("globalPooling")) { + if (context.Attr("global_pooling")) { for (size_t i = 0; i < ksize.size(); ++i) { ksize[i] = static_cast(in_x->dims()[i + 2]); } @@ -70,7 +70,7 @@ class MaxPoolWithIndexGradKernel : public framework::OpKernel { std::vector ksize = context.Attr>("ksize"); std::vector strides = context.Attr>("strides"); std::vector paddings = context.Attr>("paddings"); - if (context.Attr("globalPooling")) { + if (context.Attr("global_pooling")) { for (size_t i = 0; i < ksize.size(); ++i) { ksize[i] = static_cast(in_x_grad->dims()[i + 2]); } diff --git a/paddle/operators/softmax_with_cross_entropy_op.cc b/paddle/operators/softmax_with_cross_entropy_op.cc index 5431a1657c..9121609f10 100644 --- a/paddle/operators/softmax_with_cross_entropy_op.cc +++ b/paddle/operators/softmax_with_cross_entropy_op.cc @@ -46,7 +46,7 @@ class SoftmaxWithCrossEntropyOpMaker "(Tensor, default: Tensor), A 2-D tensor. The cross " "entropy loss with shape [N x 1]."); AddAttr( - "softLabel", + "soft_label", "(bool, default: false), A flag to indicate whether to interpretate " "the given labels as soft labels.") .SetDefault(false); @@ -100,13 +100,13 @@ class SoftmaxWithCrossEntropyOp : public framework::OperatorWithKernel { PADDLE_ENFORCE_EQ(labels_dims.size(), 2UL, "The labels should be a 2-D tensor."); - if (ctx->Attrs().Get("softLabel")) { + if (ctx->Attrs().Get("soft_label")) { PADDLE_ENFORCE_EQ(logits_dims[1], labels_dims[1], - "If Attr(softLabel) == true, the 2nd dimension of " + "If Attr(soft_label) == true, the 2nd dimension of " "Input(X) and Input(Label) should be equal."); } else { PADDLE_ENFORCE_EQ(labels_dims[1], 1UL, - "If Attr(softLabel) == false, the 2nd dimension of " + "If Attr(soft_label) == false, the 2nd dimension of " "Input(Label) should be 1."); } @@ -142,13 +142,13 @@ class SoftmaxWithCrossEntropyOpGrad : public framework::OperatorWithKernel { PADDLE_ENFORCE_EQ(labels_dims.size(), 2UL, "The labels should be a 2-D tensor."); - if (ctx->Attrs().Get("softLabel")) { + if (ctx->Attrs().Get("soft_label")) { PADDLE_ENFORCE_EQ(softmax_dims[1], labels_dims[1], - "When Attr(softLabel) == true, the 2nd dimension of " + "When Attr(soft_label) == true, the 2nd dimension of " "Input(X) and Input(Label) should be equal."); } else { PADDLE_ENFORCE_EQ(labels_dims[1], 1UL, - "When Attr(softLabel) == false, the 2nd dimension of " + "When Attr(soft_label) == false, the 2nd dimension of " "Input(Label) should be 1."); } diff --git a/paddle/operators/softmax_with_cross_entropy_op.cu b/paddle/operators/softmax_with_cross_entropy_op.cu index 2bc53ecf87..d03a1a7658 100644 --- a/paddle/operators/softmax_with_cross_entropy_op.cu +++ b/paddle/operators/softmax_with_cross_entropy_op.cu @@ -70,7 +70,7 @@ class SoftmaxWithCrossEntropyCUDAKernel : public framework::OpKernel { logits, softmax); math::CrossEntropyFunctor()( context.device_context(), loss, softmax, labels, - context.Attr("softLabel")); + context.Attr("soft_label")); } }; @@ -93,7 +93,7 @@ class SoftmaxWithCrossEntropyGradCUDAKernel : public framework::OpKernel { int block = 512; int grid = (batch_size * class_num + block - 1) / block; - if (context.Attr("softLabel")) { + if (context.Attr("soft_label")) { const T* label_data = labels->data(); SoftCrossEntropyGradientKernel<<< grid, block, 0, reinterpret_cast( diff --git a/paddle/operators/softmax_with_cross_entropy_op.h b/paddle/operators/softmax_with_cross_entropy_op.h index cffd422f18..66d7bc1569 100644 --- a/paddle/operators/softmax_with_cross_entropy_op.h +++ b/paddle/operators/softmax_with_cross_entropy_op.h @@ -44,7 +44,7 @@ class SoftmaxWithCrossEntropyKernel : public framework::OpKernel { logits, softmax); math::CrossEntropyFunctor()( context.device_context(), loss, softmax, labels, - context.Attr("softLabel")); + context.Attr("soft_label")); } }; @@ -60,7 +60,7 @@ class SoftmaxWithCrossEntropyGradKernel : public framework::OpKernel { logit_grad->ShareDataWith(*context.Input("Softmax")); const int class_num = logit_grad->dims()[1]; - if (context.Attr("softLabel")) { + if (context.Attr("soft_label")) { auto out_grad_mat = EigenMatrix::From(*out_grad); auto logit_grad_mat = EigenMatrix::From(*logit_grad); auto lbl_mat = EigenMatrix::From(*labels); diff --git a/python/paddle/v2/framework/tests/test_cross_entropy_op.py b/python/paddle/v2/framework/tests/test_cross_entropy_op.py index 4ea14da7fd..919b6c3f67 100644 --- a/python/paddle/v2/framework/tests/test_cross_entropy_op.py +++ b/python/paddle/v2/framework/tests/test_cross_entropy_op.py @@ -49,7 +49,7 @@ class TestCrossEntropyOp2(OpTest): self.inputs = {"X": X, "Label": label} self.outputs = {"Y": cross_entropy} - self.attrs = {"softLabel": True} + self.attrs = {"soft_label": True} def test_check_output(self): self.check_output() @@ -82,7 +82,7 @@ class TestCrossEntropyOp3(OpTest): self.inputs = {"X": X, "Label": label.astype(np.float32)} self.outputs = {"Y": cross_entropy} - self.attrs = {"softLabel": True} + self.attrs = {"soft_label": True} def test_check_output(self): self.check_output() diff --git a/python/paddle/v2/framework/tests/test_pool2d_op.py b/python/paddle/v2/framework/tests/test_pool2d_op.py index 2941fda81b..3fcd8941d4 100644 --- a/python/paddle/v2/framework/tests/test_pool2d_op.py +++ b/python/paddle/v2/framework/tests/test_pool2d_op.py @@ -56,8 +56,8 @@ class TestPool2d_Op(OpTest): 'strides': self.strides, 'paddings': self.paddings, 'ksize': self.ksize, - 'poolingType': self.pool_type, - 'globalPooling': self.global_pool, + 'pooling_type': self.pool_type, + 'global_pooling': self.global_pool, } self.outputs = {'Out': output} diff --git a/python/paddle/v2/framework/tests/test_pool3d_op.py b/python/paddle/v2/framework/tests/test_pool3d_op.py index 8792b492e3..f4e938041f 100644 --- a/python/paddle/v2/framework/tests/test_pool3d_op.py +++ b/python/paddle/v2/framework/tests/test_pool3d_op.py @@ -64,8 +64,8 @@ class TestPool3d_Op(OpTest): 'strides': self.strides, 'paddings': self.paddings, 'ksize': self.ksize, - 'poolingType': self.pool_type, - 'globalPooling': self.global_pool, + 'pooling_type': self.pool_type, + 'global_pooling': self.global_pool, } self.outputs = {'Out': output} diff --git a/python/paddle/v2/framework/tests/test_pool_max_op.py b/python/paddle/v2/framework/tests/test_pool_max_op.py index f0f8aa6089..b78f9bba05 100644 --- a/python/paddle/v2/framework/tests/test_pool_max_op.py +++ b/python/paddle/v2/framework/tests/test_pool_max_op.py @@ -86,7 +86,7 @@ class TestMaxPoolWithIndex_Op(OpTest): 'strides': self.strides, 'paddings': self.paddings, 'ksize': self.ksize, - 'globalPooling': self.global_pool, + 'global_pooling': self.global_pool, } self.inputs = {'X': input} diff --git a/python/paddle/v2/framework/tests/test_softmax_with_cross_entropy_op.py b/python/paddle/v2/framework/tests/test_softmax_with_cross_entropy_op.py index 377d07fb59..05ba954c0b 100644 --- a/python/paddle/v2/framework/tests/test_softmax_with_cross_entropy_op.py +++ b/python/paddle/v2/framework/tests/test_softmax_with_cross_entropy_op.py @@ -57,7 +57,7 @@ class TestSoftmaxWithCrossEntropyOp2(OpTest): self.inputs = {"Logits": logits, "Label": labels} self.outputs = {"Softmax": softmax, "Loss": cross_entropy} - self.attrs = {"softLabel": True} + self.attrs = {"soft_label": True} def test_check_output(self): self.check_output() From c65bdd95656fa4abee93d02429b8694787667f0e Mon Sep 17 00:00:00 2001 From: qijun Date: Mon, 16 Oct 2017 15:49:44 -0700 Subject: [PATCH 31/36] fix SelectedRows rows() method gpu runtime error --- paddle/pybind/pybind.cc | 12 +++++++++++- 1 file changed, 11 insertions(+), 1 deletion(-) diff --git a/paddle/pybind/pybind.cc b/paddle/pybind/pybind.cc index 23e76011c9..f313eda964 100644 --- a/paddle/pybind/pybind.cc +++ b/paddle/pybind/pybind.cc @@ -153,7 +153,17 @@ PYBIND11_PLUGIN(core) { .def("set_height", &SelectedRows::set_height) .def("height", &SelectedRows::height) .def("set_rows", &SelectedRows::set_rows) - .def("rows", &SelectedRows::rows, py::return_value_policy::reference); + .def("rows", [](SelectedRows &self) { +#ifndef PADDLE_WITH_CUDA + return self.rows(); +#else + auto rows = self.rows(); + std::vector new_rows; + new_rows.reserve(rows.size()); + std::copy(rows.begin(), rows.end(), std::back_inserter(new_rows)); + return new_rows; +#endif + }); py::class_(m, "Variable", R"DOC(Variable Class. From 8e52b34a0c2f14175724571b89814277cbf28d59 Mon Sep 17 00:00:00 2001 From: Yu Yang Date: Mon, 16 Oct 2017 16:08:01 -0700 Subject: [PATCH 32/36] Implement FC layer with helper (#4726) * Implement FC layer with helper * Update LayerHelper * Add debug string for Python ProtoBuf and Rename `Sync` to `Flush` * Add check of ProtoBuf initialization * Layer wrapper for FC * Fix unittest * Fix CI * Add code generator * AttributeChecker Better error log and speicalize bool Since lots of types can be cast to bool * Complete mlp, fit_a_line --- python/paddle/v2/framework/framework.py | 14 +- python/paddle/v2/framework/layer_helper.py | 160 ++++++++++++++++++ python/paddle/v2/framework/layers.py | 143 ++++++++++++++++ .../paddle/v2/framework/tests/test_layers.py | 43 +++++ .../v2/framework/tests/test_operator_desc.py | 2 +- 5 files changed, 354 insertions(+), 8 deletions(-) create mode 100644 python/paddle/v2/framework/layer_helper.py create mode 100644 python/paddle/v2/framework/layers.py create mode 100644 python/paddle/v2/framework/tests/test_layers.py diff --git a/python/paddle/v2/framework/framework.py b/python/paddle/v2/framework/framework.py index d649e69d58..a17f988bf4 100644 --- a/python/paddle/v2/framework/framework.py +++ b/python/paddle/v2/framework/framework.py @@ -153,7 +153,8 @@ class OpProtoHolder(object): self.op_proto_map[proto.type] = proto def get_op_proto(self, type): - assert type in self.op_proto_map, "Operator \"%s\" has not been registered." % type + if type not in self.op_proto_map: + raise ValueError("Operator \"%s\" has not been registered." % type) return self.op_proto_map[type] @@ -374,10 +375,10 @@ class Program(object): cls._instance = cls() return cls._instance - def __init__(self): - assert not hasattr(self.__class__, - '_instance'), 'Do not call constructor directly!' - self.desc = core.ProgramDesc.instance() + def __init__(self, desc=None): + if desc is None: + desc = core.ProgramDesc.instance() + self.desc = desc self.blocks = [Block(self, 0)] self.current_block_idx = 0 @@ -428,7 +429,6 @@ class Parameter(Variable): if each < 0: raise ValueError("Parameter shape should not be related with " "batch-size") - Variable.__init__(self, block, shape=shape, dtype=dtype, **kwargs) self.trainable = kwargs.get('trainable', True) self.init_attr = kwargs.get('initialize_attr', { @@ -441,7 +441,7 @@ class Parameter(Variable): self._append_initialize_ops_() def _append_initialize_ops_(self): - attr = copy.deepcopy(self.init_attr) + attr = self.init_attr op_type = attr.pop('type', None) block = self.block assert isinstance(block, Block) diff --git a/python/paddle/v2/framework/layer_helper.py b/python/paddle/v2/framework/layer_helper.py new file mode 100644 index 0000000000..26d3e04310 --- /dev/null +++ b/python/paddle/v2/framework/layer_helper.py @@ -0,0 +1,160 @@ +from paddle.v2.framework.framework import Variable, OpProtoHolder, g_program +import paddle.v2.framework.core as core +import copy +import itertools + + +def unique_name(prefix): + uid = core.unique_integer() # unique during whole process. + return "_".join([prefix, str(uid)]) + + +class LayerHelper(object): + def __init__(self, layer_type, **kwargs): + self.kwargs = kwargs + self.layer_type = layer_type + name = self.kwargs.get('name', None) + if name is None: + self.kwargs['name'] = unique_name(self.layer_type) + + @property + def name(self): + return self.kwargs['name'] + + @property + def program(self): + prog = self.kwargs.get('program', None) + if prog is None: + return g_program + else: + return prog + + def append_op(self, *args, **kwargs): + return self.program.current_block().append_op(*args, **kwargs) + + def multiple_input(self, input_param_name='input'): + inputs = self.kwargs.get(input_param_name, []) + type_error = TypeError( + "Input of {0} layer should be Variable or sequence of Variable". + format(self.layer_type)) + if isinstance(inputs, Variable): + inputs = [inputs] + elif not isinstance(inputs, list) and not isinstance(inputs, tuple): + raise type_error + else: + for each in inputs: + if not isinstance(each, Variable): + raise type_error + return inputs + + def input(self, input_param_name='input'): + inputs = self.multiple_input(input_param_name) + if len(inputs) != 1: + raise "{0} layer only takes one input".format(self.layer_type) + return inputs[0] + + @property + def param_attr(self): + default = { + 'name': None, + 'init_attr': { + 'type': 'uniform_random', + 'min': -1.0, + 'max': 1.0 + } + } + actual = self.kwargs.get('param_attr', None) + return actual if actual is not None else default + + def bias_attr(self, size, dtype): + bias_attr = self.kwargs.get('bias_attr', False) + if bias_attr is None or bias_attr: + bias_attr = { + 'name': None, + 'init_attr': { + 'type': 'fill_constant', + 'value': 0.0, + 'shape': [size], + 'dataType': dtype + } + } + return bias_attr + + def multiple_param_attr(self, length): + param_attr = self.param_attr + if isinstance(param_attr, dict): + param_attr = [param_attr] + + if len(param_attr) != 1 and len(param_attr) != length: + raise ValueError("parameter number mismatch") + elif len(param_attr) == 1 and length != 1: + tmp = [None] * length + for i in xrange(length): + tmp[i] = copy.deepcopy(param_attr[0]) + param_attr = tmp + return param_attr + + def iter_inputs_and_params(self, input_param_name='input'): + inputs = self.multiple_input(input_param_name) + param_attrs = self.multiple_param_attr(len(inputs)) + for ipt, param_attr in itertools.izip(inputs, param_attrs): + yield ipt, param_attr + + def input_dtype(self, input_param_name='input'): + inputs = self.multiple_input(input_param_name) + dtype = None + for each in inputs: + if dtype is None: + dtype = each.data_type + elif dtype != each.data_type: + raise ValueError("Data Type mismatch") + return dtype + + def create_parameter(self, attr, shape, dtype, suffix='w'): + if attr['name'] is None: + attr['name'] = unique_name(".".join([self.name, suffix])) + return self.program.global_block().create_parameter( + name=attr['name'], + dtype=dtype, + shape=shape, + initialize_attr=attr['init_attr']) + + def create_tmp_variable(self, dtype): + return self.program.current_block().create_var( + name=unique_name(".".join([self.name, 'tmp'])), dtype=dtype) + + def create_global_variable(self, *args, **kwargs): + return self.program.global_block().create_var(*args, **kwargs) + + def append_bias_op(self, input_var): + bias_attr = self.bias_attr( + self.kwargs['size'], dtype=input_var.data_type) + if not bias_attr: + return input_var + b = self.create_parameter( + attr=bias_attr, + shape=[self.kwargs['size']], + dtype=input_var.data_type, + suffix='b') + tmp = self.create_tmp_variable(dtype=input_var.data_type) + self.append_op( + type='elementwise_add', + inputs={'X': [input_var], + 'Y': [b]}, + outputs={'Out': [tmp]}) + return tmp + + def append_activation(self, input_var): + act = self.kwargs.get('act', None) + if act is None: + return input_var + if isinstance(act, basestring): + act = {'type': act} + tmp = self.create_tmp_variable(dtype=input_var.data_type) + act_type = act.pop('type') + self.append_op( + type=act_type, + inputs={"X": [input_var]}, + outputs={"Y": [tmp]}, + attrs=act) + return tmp diff --git a/python/paddle/v2/framework/layers.py b/python/paddle/v2/framework/layers.py new file mode 100644 index 0000000000..44b587b116 --- /dev/null +++ b/python/paddle/v2/framework/layers.py @@ -0,0 +1,143 @@ +from paddle.v2.framework.layer_helper import LayerHelper +import paddle.v2.framework.core as core +from paddle.v2.framework.framework import OpProtoHolder, Variable +import re + +__all__ = ['fc_layer', 'data_layer', 'cross_entropy'] + + +def fc_layer(input, + size, + param_attr=None, + bias_attr=True, + name=None, + act=None, + num_flatten_dims=1, + program=None): + # create helper + helper = LayerHelper('fc', **locals()) + + dtype = helper.input_dtype() + + # mul + mul_results = [] + for input_var, param_attr in helper.iter_inputs_and_params(): + input_shape = input_var.shape + param_shape = list(input_shape[num_flatten_dims:]) + [size] + w = helper.create_parameter( + attr=param_attr, shape=param_shape, dtype=dtype) + tmp = helper.create_tmp_variable(dtype) + helper.append_op( + type="mul", + inputs={ + "X": input_var, + "Y": w, + }, + outputs={"Out": tmp}, + attrs={'x_num_col_dims': num_flatten_dims}) + mul_results.append(tmp) + + # sum + if len(mul_results) == 1: + pre_bias = mul_results[0] + else: + pre_bias = helper.create_tmp_variable(dtype) + helper.append_op( + type="sum", inputs={"X": mul_results}, outputs={"Out": pre_bias}) + # add bias + pre_activation = helper.append_bias_op(pre_bias) + # add activation + return helper.append_activation(pre_activation) + + +def data_layer(name, + shape, + data_type='float32', + type=core.VarDesc.VarType.LOD_TENSOR, + program=None): + helper = LayerHelper('data', **locals()) + shape = [-1] + shape # append batch size as -1 + return helper.create_global_variable( + name=name, shape=shape, dtype=data_type, type=type) + + +def _convert_(name): + s1 = re.sub('(.)([A-Z][a-z]+)', r'\1_\2', name) + return re.sub('([a-z0-9])([A-Z])', r'\1_\2', s1).lower() + + +def _create_op_func_(op_type): + op_proto = OpProtoHolder.instance().get_op_proto(op_type) + if len(op_proto.outputs) != 1: + raise ValueError( + "Only one output operator can be automatically generated") + + if op_proto.outputs[0].duplicable: + raise ValueError( + "Only not duplicable op can be automatically generated") + + o_name = op_proto.outputs[0].name + + def func(**kwargs): + helper = LayerHelper(op_type, **kwargs) + inputs = dict() + dtype = None + for ipt in op_proto.inputs: + name = _convert_(ipt.name) + val = kwargs.pop(name, []) + if not isinstance(val, list) and not isinstance(val, tuple): + val = [val] + for each in val: + if not isinstance(each, Variable): + raise ValueError("input of {0} must be variable".format( + op_type)) + + if dtype is None: + dtype = each.data_type + elif dtype != each.data_type: + raise ValueError( + "operator {0} must input same dtype".format(op_type)) + inputs[ipt.name] = val + + out = helper.create_tmp_variable(dtype=dtype) + helper.append_op( + type=op_type, inputs=inputs, outputs={o_name: [out]}, attrs=kwargs) + return out + + func.__name__ = op_type + globals()[op_type] = func + global __all__ + __all__.append(op_type) + + +_create_op_func_('mean') + + +def cross_entropy(input, label, **kwargs): + helper = LayerHelper('cross_entropy', **kwargs) + out = helper.create_tmp_variable(dtype=input.data_type) + helper.append_op( + type='cross_entropy', + inputs={'X': [input], + 'Label': [label]}, + outputs={'Y': [out]}, + attrs=kwargs) + return out + + +def square_error_cost(input, label, **kwargs): + helper = LayerHelper('square_error_cost', **kwargs) + minus_out = helper.create_tmp_variable(dtype=input.data_type) + helper.append_op( + type='elementwise_sub', + inputs={'X': [input], + 'Y': [label]}, + outputs={'Out': [minus_out]}) + + square_out = helper.create_tmp_variable(dtype=input.data_type) + helper.append_op( + type='pow', + inputs={'X': [minus_out]}, + outputs={'Y': [square_out]}, + attrs={'factor': 2.0}) + return square_out diff --git a/python/paddle/v2/framework/tests/test_layers.py b/python/paddle/v2/framework/tests/test_layers.py new file mode 100644 index 0000000000..1ef2591cca --- /dev/null +++ b/python/paddle/v2/framework/tests/test_layers.py @@ -0,0 +1,43 @@ +from paddle.v2.framework.layers import fc_layer, data_layer, cross_entropy, mean, square_error_cost +from paddle.v2.framework.framework import Program, g_program +import paddle.v2.framework.core as core +import unittest + + +class TestBook(unittest.TestCase): + def test_fit_a_line(self): + pd = core.ProgramDesc.__create_program_desc__() + program = Program(desc=pd) + x = data_layer( + name='x', shape=[13], data_type='float32', program=program) + y_predict = fc_layer(input=x, size=1, act=None, program=program) + + y = data_layer( + name='y', shape=[1], data_type='float32', program=program) + cost = square_error_cost(input=y_predict, label=y, program=program) + + avg_cost = mean(x=cost, program=program) + self.assertIsNotNone(avg_cost) + print str(program) + + def test_recognize_digits_mlp(self): + pd = core.ProgramDesc.__create_program_desc__() + program = Program(desc=pd) + + # Change g_program, so the rest layers use `g_program` + images = data_layer( + name='pixel', shape=[784], data_type='float32', program=program) + label = data_layer( + name='label', shape=[1], data_type='int32', program=program) + hidden1 = fc_layer(input=images, size=128, act='relu', program=program) + hidden2 = fc_layer(input=hidden1, size=64, act='relu', program=program) + predict = fc_layer( + input=hidden2, size=10, act='softmax', program=program) + cost = cross_entropy(input=predict, label=label, program=program) + avg_cost = mean(x=cost, program=program) + self.assertIsNotNone(avg_cost) + print str(program) + + +if __name__ == '__main__': + unittest.main() diff --git a/python/paddle/v2/framework/tests/test_operator_desc.py b/python/paddle/v2/framework/tests/test_operator_desc.py index dfe39c98f7..af4e980b8e 100644 --- a/python/paddle/v2/framework/tests/test_operator_desc.py +++ b/python/paddle/v2/framework/tests/test_operator_desc.py @@ -16,7 +16,7 @@ class TestOperator(unittest.TestCase): try: block.append_op(type="no_such_op") self.assertFail() - except AssertionError as a_err: + except ValueError as a_err: self.assertEqual(a_err.message, "Operator \"no_such_op\" has not been registered.") From b10cd435547290d6417b00688c190bc98169d1df Mon Sep 17 00:00:00 2001 From: Qiao Longfei Date: Mon, 16 Oct 2017 16:35:55 -0700 Subject: [PATCH 33/36] rm cpp executor_test, rewrite in python later (#4849) * rm cpp executor_test, rewrite in python later * remove executor_test code in CMakeList.txt --- paddle/framework/CMakeLists.txt | 7 - paddle/framework/executor_test.cc | 348 ------------------------------ 2 files changed, 355 deletions(-) delete mode 100644 paddle/framework/executor_test.cc diff --git a/paddle/framework/CMakeLists.txt b/paddle/framework/CMakeLists.txt index c8d9dac21d..6585fb7b5d 100644 --- a/paddle/framework/CMakeLists.txt +++ b/paddle/framework/CMakeLists.txt @@ -43,13 +43,6 @@ cc_library(backward SRCS backward.cc DEPS net_op) cc_test(backward_test SRCS backward_test.cc DEPS backward recurrent_op device_context) cc_library(executor SRCS executor.cc DEPS op_registry device_context scope framework_proto backward) -set(EXECUTOR_TEST_OP elementwise_add_op gaussian_random_op feed_op fetch_op - mul_op sum_op squared_l2_distance_op fill_constant_op sgd_op mean_op) -if(WITH_GPU) - nv_test(executor_test SRCS executor_test.cc DEPS executor ${EXECUTOR_TEST_OP}) -else() - cc_test(executor_test SRCS executor_test.cc DEPS executor ${EXECUTOR_TEST_OP}) -endif() cc_library(tensor_array SRCS tensor_array.cc DEPS lod_tensor) cc_test(tensor_array_test SRCS tensor_array_test.cc DEPS tensor_array place) diff --git a/paddle/framework/executor_test.cc b/paddle/framework/executor_test.cc deleted file mode 100644 index e08d31e361..0000000000 --- a/paddle/framework/executor_test.cc +++ /dev/null @@ -1,348 +0,0 @@ -/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve. - -Licensed under the Apache License, Version 2.0 (the "License"); -you may not use this file except in compliance with the License. -You may obtain a copy of the License at - - http://www.apache.org/licenses/LICENSE-2.0 - -Unless required by applicable law or agreed to in writing, software -distributed under the License is distributed on an "AS IS" BASIS, -WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -See the License for the specific language governing permissions and -limitations under the License. */ - -#include "paddle/framework/executor.h" - -#include -#include - -#include "gflags/gflags.h" -#include "gtest/gtest.h" -#include "paddle/framework/attribute.h" -#include "paddle/framework/backward.h" -#include "paddle/framework/block_desc.h" -#include "paddle/framework/op_desc.h" -#include "paddle/framework/op_registry.h" -#include "paddle/framework/operator.h" - -USE_OP(elementwise_add); -USE_OP(gaussian_random); -USE_NO_KERNEL_OP(feed); -USE_NO_KERNEL_OP(fetch); -USE_OP(mul); -USE_OP(sum); -USE_OP(squared_l2_distance); -USE_OP(fill_constant); -USE_OP(mean); -USE_OP(sgd); - -constexpr auto kFeedValueName = "feed_value"; -constexpr auto kFetchValueName = "fetch_value"; - -using namespace paddle::platform; -using namespace paddle::framework; - -void AddOp(const std::string& type, const VariableNameMap& inputs, - const VariableNameMap& outputs, AttributeMap attrs, - paddle::framework::BlockDescBind* block) { - // insert output - for (auto kv : outputs) { - for (auto v : kv.second) { - // <<<<<<< HEAD - // auto var = block->Var(v); - // var->SetType(VarDesc::LOD_TENSOR); - // var->SetDataType(paddle::framework::DataType::FP32); - // ======= - if (!block->HasVar(v)) { - auto var = block->Var(v); - var->SetDataType(paddle::framework::DataType::FP32); - } - // >>>>>>> origin/develop - } - } - - // insert op - auto op = block->AppendOp(); - op->SetType(type); - for (auto& kv : inputs) { - op->SetInput(kv.first, kv.second); - } - for (auto& kv : outputs) { - op->SetOutput(kv.first, kv.second); - } - op->SetAttrMap(attrs); - op->CheckAttrs(); -} - -// Tensors in feed value variable will only be in CPUPlace -// So we can memcpy the data from vector to feed_value -template -void SetFeedVariable(const std::vector>& inputs, - const std::vector>& dims) { - Variable* g_feed_value = GetGlobalScope().FindVar(kFeedValueName); - auto& feed_inputs = - *(g_feed_value->GetMutable>()); - size_t size = inputs.size(); - feed_inputs.resize(size); - for (size_t i = 0; i < size; i++) { - T* dst = feed_inputs[i].mutable_data(make_ddim(dims[i]), CPUPlace()); - memcpy(dst, inputs[i].data(), inputs[i].size() * sizeof(T)); - } -} - -// Tensors in fetch value variable will only be in CPUPlace -// So we can memcpy the data from fetch_value to vector -template -std::vector> GetFetchVariable() { - Variable* g_fetch_value = GetGlobalScope().FindVar(kFetchValueName); - auto& fetch_outputs = - *(g_fetch_value->GetMutable>()); - - size_t size = fetch_outputs.size(); - std::vector> result; - result.reserve(size); - for (size_t i = 0; i < size; i++) { - std::vector tmp; - tmp.resize(fetch_outputs[i].numel()); - memcpy(tmp.data(), fetch_outputs[i].data(), - fetch_outputs[i].numel() * sizeof(T)); - result.push_back(tmp); - } - - return result; -} - -class ExecutorTesterRandom : public ::testing::Test { - public: - virtual void SetUp() override { - int input_dim = 3, batch_size = 2, embed_dim = 5; - - auto temp_init_root_block = init_pdesc_.add_blocks(); - temp_init_root_block->set_idx(0); - temp_init_root_block->set_parent_idx(-1); - paddle::framework::ProgramDescBind& init_program = - paddle::framework::ProgramDescBind::Instance(&init_pdesc_); - paddle::framework::BlockDescBind* init_root_block = init_program.Block(0); - - AddOp("gaussian_random", {}, {{"Out", {"w1"}}}, - {{"dims", std::vector{input_dim, embed_dim}}}, init_root_block); - AddOp("gaussian_random", {}, {{"Out", {"w2"}}}, - {{"dims", std::vector{embed_dim, input_dim}}}, init_root_block); - AddOp("fetch", {{"Input", {"w1"}}}, {{"Out", {kFetchValueName}}}, - {{"col", 0}}, init_root_block); - AddOp("fetch", {{"Input", {"w2"}}}, {{"Out", {kFetchValueName}}}, - {{"col", 1}}, init_root_block); - - // flush - init_program.Proto(); - - // run block - auto temp_root_block = pdesc_.add_blocks(); - temp_root_block->set_idx(0); - temp_root_block->set_parent_idx(-1); - paddle::framework::ProgramDescBind& program = - paddle::framework::ProgramDescBind::Instance(&pdesc_); - paddle::framework::BlockDescBind* root_block = program.Block(0); - - // feed data - inputs_.push_back({1.0, 1.0, 1.0, 1.0, 1.0, 1.0}); - dims_.push_back({batch_size, input_dim}); - AddOp("feed", {{"Input", {kFeedValueName}}}, {{"Out", {"a"}}}, - {{"dims", std::vector{batch_size, input_dim}}, {"col", 0}}, - root_block); - - // forward - AddOp("mul", {{"X", {"a"}}, {"Y", {"w1"}}}, {{"Out", {"b"}}}, {}, - root_block); - AddOp("mul", {{"X", {"b"}}, {"Y", {"w2"}}}, {{"Out", {"a_out"}}}, {}, - root_block); - AddOp("squared_l2_distance", {{"X", {"a"}}, {"Y", {"a_out"}}}, - {{"Out", {"l2_distance"}}, {"sub_result", {"l2_distance_sub"}}}, {}, - root_block); - AddOp("mean", {{"X", {"l2_distance"}}}, {{"Out", {"mean_out"}}}, {}, - root_block); - - // backward - auto target = VarDescBind("mean_out"); - AppendBackward(program, target, {}); - - // update - AddOp("fill_constant", {}, {{"Out", {"learning_rate"}}}, - {{"shape", std::vector{1}}, {"value", float(0.001)}}, - root_block); - AddOp("sgd", {{"Param", {"w1"}}, - {"LearningRate", {"learning_rate"}}, - {"Grad", {"w1@GRAD"}}}, - {{"ParamOut", {"w1"}}}, {}, root_block); - AddOp("sgd", {{"Param", {"w2"}}, - {"LearningRate", {"learning_rate"}}, - {"Grad", {"w2@GRAD"}}}, - {{"ParamOut", {"w2"}}}, {}, root_block); - - AddOp("fetch", {{"Input", {"w1"}}}, {{"Out", {kFetchValueName}}}, - {{"col", 0}}, root_block); - AddOp("fetch", {{"Input", {"w2"}}}, {{"Out", {kFetchValueName}}}, - {{"col", 1}}, root_block); - AddOp("fetch", {{"Input", {"l2_distance"}}}, {{"Out", {kFetchValueName}}}, - {{"col", 0}}, root_block); - - // flush - program.Proto(); - } - - protected: - ProgramDesc init_pdesc_; - ProgramDesc pdesc_; - std::vector> inputs_; - std::vector> dims_; -}; - -class ExecutorTesterFeedAndFetch : public ::testing::Test { - public: - virtual void SetUp() override { - auto temp_root_block = pdesc_.add_blocks(); - temp_root_block->set_idx(0); - temp_root_block->set_parent_idx(-1); - - // wrap to BlockDescBind - paddle::framework::ProgramDescBind& program = - paddle::framework::ProgramDescBind::Instance(&pdesc_); - paddle::framework::BlockDescBind* root_block = program.Block(0); - - std::vector dim{6}; - - AddOp("feed", {{"Input", {kFeedValueName}}}, {{"Out", {"a"}}}, - {{"dims", dim}, {"col", 0}}, root_block); - AddOp("feed", {{"Input", {kFeedValueName}}}, {{"Out", {"b"}}}, - {{"dims", dim}, {"col", 1}}, root_block); - AddOp("fetch", {{"Input", {"a"}}}, {{"Out", {kFetchValueName}}}, - {{"col", 0}}, root_block); - AddOp("fetch", {{"Input", {"b"}}}, {{"Out", {kFetchValueName}}}, - {{"col", 1}}, root_block); - - // flush - program.Proto(); - - std::vector vec1 = {1.0, 2.0, 3.0, 4.0, 5.0, 6.0}; - std::vector vec2 = {4.0, 5.0, 6.0, 7.0, 8.0, 9.0}; - inputs_.push_back(vec1); - inputs_.push_back(vec2); - dims_.push_back({static_cast(vec1.size())}); - dims_.push_back({static_cast(vec2.size())}); - } - - protected: - ProgramDesc pdesc_; - std::vector> inputs_; - std::vector> dims_; -}; - -#ifndef PADDLE_WITH_CUDA -TEST_F(ExecutorTesterRandom, CPU) { - std::vector places; - CPUPlace cpu_place; - places.push_back(cpu_place); - - // We have a global Scope and BuddyAllocator, and we must ensure - // global BuddyAllocator is initialized before global Scope. Thus, - // global Scope will deconstruct before BuddyAllocator. Otherwise, - // "pointer being freed was not allocated" error will appear. - paddle::memory::Used(cpu_place); - - std::unique_ptr executor(new Executor(places)); - executor->Run(init_pdesc_, &GetGlobalScope(), 0); - SetFeedVariable(inputs_, dims_); - executor->Run(pdesc_, &GetGlobalScope(), 0); - std::vector> result = GetFetchVariable(); -} - -TEST_F(ExecutorTesterFeedAndFetch, CPU) { - std::vector places; - CPUPlace cpu_place; - places.emplace_back(cpu_place); - - // We have a global Scope and BuddyAllocator, and we must ensure - // global BuddyAllocator is initialized before global Scope. Thus, - // global Scope will deconstruct before BuddyAllocator. Otherwise, - // "pointer being freed was not allocated" error will appear. - paddle::memory::Used(cpu_place); - - std::unique_ptr executor(new Executor(places)); - - for (int batch_id = 0; batch_id < 3; batch_id++) { - SetFeedVariable(inputs_, dims_); - executor->Run(pdesc_, &GetGlobalScope(), 0); - std::vector> result = GetFetchVariable(); - ASSERT_EQ(result.size(), inputs_.size()); - for (size_t i = 0; i < result.size(); ++i) { - ASSERT_EQ(result[i].size(), inputs_[i].size()); - for (size_t j = 0; j < result[i].size(); ++j) { - ASSERT_EQ(result[i][j], inputs_[i][j]); - } - } - } -} -#else -TEST_F(ExecutorTesterRandom, GPU) { - std::vector places; - GPUPlace gpu_place(0); - places.push_back(gpu_place); - - // We have a global Scope and BuddyAllocator, and we must ensure - // global BuddyAllocator is initialized before global Scope. Thus, - // global Scope will deconstruct before BuddyAllocator. Otherwise, - // "pointer being freed was not allocated" error will appear. - // If paddle is compiled with GPU, both CPU and GPU BuddyAllocator - // need to be used at first. - paddle::memory::Used(CPUPlace()); - paddle::memory::Used(gpu_place); - - std::unique_ptr executor(new Executor(places)); - - executor->Run(init_pdesc_, &GetGlobalScope(), 0); - for (int batch_id = 0; batch_id < 3; batch_id++) { - SetFeedVariable(inputs_, dims_); - executor->Run(pdesc_, &GetGlobalScope(), 0); - } -} - -TEST_F(ExecutorTesterFeedAndFetch, GPU) { - std::vector places; - GPUPlace gpu_place(0); - places.push_back(gpu_place); - // We have a global Scope and BuddyAllocator, and we must ensure - // global BuddyAllocator is initialized before global Scope. Thus, - // global Scope will deconstruct before BuddyAllocator. Otherwise, - // "pointer being freed was not allocated" error will appear. - // If paddle is compiled with GPU, both CPU and GPU BuddyAllocator - // need to be used at first. - paddle::memory::Used(CPUPlace()); - paddle::memory::Used(gpu_place); - - std::unique_ptr executor(new Executor(places)); - - for (int batch_id = 0; batch_id < 3; batch_id++) { - SetFeedVariable(inputs_, dims_); - executor->Run(pdesc_, &GetGlobalScope(), 0); - std::vector> result = GetFetchVariable(); - PADDLE_ENFORCE_EQ(result.size(), inputs_.size()); - for (size_t i = 0; i < result.size(); ++i) { - PADDLE_ENFORCE_EQ(result[i].size(), inputs_[i].size()); - for (size_t j = 0; j < result[i].size(); ++j) { - PADDLE_ENFORCE_EQ(result[i][j], inputs_[i][j]); - } - } - } -} - -DECLARE_double(fraction_of_gpu_memory_to_use); - -int main(int argc, char** argv) { - testing::InitGoogleTest(&argc, argv); - // Use less GPU memory for unittest. - FLAGS_fraction_of_gpu_memory_to_use = 0.25; - return RUN_ALL_TESTS(); -} - -#endif From fafc54d6c08b06a8b86217118eace48a9a5a122f Mon Sep 17 00:00:00 2001 From: Qiao Longfei Date: Mon, 16 Oct 2017 19:00:22 -0700 Subject: [PATCH 34/36] fix var_type_inference_test (#4854) --- paddle/framework/var_type_inference_test.cc | 16 ++++++++-------- 1 file changed, 8 insertions(+), 8 deletions(-) diff --git a/paddle/framework/var_type_inference_test.cc b/paddle/framework/var_type_inference_test.cc index 97b8c64748..87399208e9 100644 --- a/paddle/framework/var_type_inference_test.cc +++ b/paddle/framework/var_type_inference_test.cc @@ -68,10 +68,10 @@ TEST(InferVarType, sum_op) { op->SetInput("X", {"test_a", "test_b", "test_c"}); op->SetOutput("Out", {"test_out"}); - prog.Block(0)->NewVar("test_a")->SetType(VarDesc::SELECTED_ROWS); - prog.Block(0)->NewVar("test_b")->SetType(VarDesc::SELECTED_ROWS); - prog.Block(0)->NewVar("test_c")->SetType(VarDesc::SELECTED_ROWS); - prog.Block(0)->NewVar("test_out"); + prog.Block(0)->Var("test_a")->SetType(VarDesc::SELECTED_ROWS); + prog.Block(0)->Var("test_b")->SetType(VarDesc::SELECTED_ROWS); + prog.Block(0)->Var("test_c")->SetType(VarDesc::SELECTED_ROWS); + prog.Block(0)->Var("test_out"); op->InferVarType(prog.Block(0)); @@ -89,10 +89,10 @@ TEST(InferVarType, sum_op_without_infer_var_type) { op->SetInput("X", {"test2_a", "test2_b", "test2_c"}); op->SetOutput("Out", {"test2_out"}); - prog.Block(0)->NewVar("test2_a")->SetType(VarDesc::SELECTED_ROWS); - prog.Block(0)->NewVar("test2_b")->SetType(VarDesc::SELECTED_ROWS); - prog.Block(0)->NewVar("test2_c")->SetType(VarDesc::SELECTED_ROWS); - prog.Block(0)->NewVar("test2_out"); + prog.Block(0)->Var("test2_a")->SetType(VarDesc::SELECTED_ROWS); + prog.Block(0)->Var("test2_b")->SetType(VarDesc::SELECTED_ROWS); + prog.Block(0)->Var("test2_c")->SetType(VarDesc::SELECTED_ROWS); + prog.Block(0)->Var("test2_out"); op->InferVarType(prog.Block(0)); From 73a8b78a720e2998634fc597efb78cc8fe53b2d7 Mon Sep 17 00:00:00 2001 From: Yu Yang Date: Mon, 16 Oct 2017 19:11:42 -0700 Subject: [PATCH 35/36] Correct OpWithKernel's infershape (#4847) They are public now --- paddle/operators/accuracy_op.cc | 1 - paddle/operators/activation_op.cc | 2 -- paddle/operators/adadelta_op.cc | 1 - paddle/operators/adagrad_op.cc | 1 - paddle/operators/adam_op.cc | 1 - paddle/operators/adamax_op.cc | 1 - paddle/operators/clip_op.cc | 2 -- paddle/operators/concat_op.cc | 2 -- paddle/operators/conv2d_op.h | 2 -- paddle/operators/conv_shift_op.cc | 2 -- paddle/operators/cos_sim_op.cc | 2 -- paddle/operators/crop_op.cc | 2 -- paddle/operators/cross_entropy_op.cc | 4 ++-- paddle/operators/decayed_adagrad_op.cc | 1 - paddle/operators/dropout_op.cc | 2 -- paddle/operators/elementwise_op.h | 2 -- paddle/operators/fill_constant_op.cc | 2 +- paddle/operators/fill_zeros_like_op.cc | 1 - paddle/operators/gather_op.cc | 4 ++-- paddle/operators/gaussian_random_op.cc | 2 +- paddle/operators/gru_unit_op.cc | 2 -- paddle/operators/lookup_table_op.cc | 4 ++-- paddle/operators/lstm_unit_op.cc | 2 -- paddle/operators/margin_rank_loss_op.cc | 2 -- paddle/operators/mean_op.cc | 2 -- paddle/operators/minus_op.cc | 1 - paddle/operators/modified_huber_loss_op.cc | 2 -- paddle/operators/mul_op.cc | 2 -- paddle/operators/multiplex_op.cc | 4 ++-- paddle/operators/pad_op.cc | 2 -- paddle/operators/pool_op.h | 2 -- paddle/operators/pool_with_index_op.cc | 2 -- paddle/operators/prelu_op.cc | 2 -- paddle/operators/rank_loss_op.cc | 2 -- paddle/operators/reduce_op.cc | 2 -- paddle/operators/reshape_op.cc | 2 -- paddle/operators/rmsprop_op.cc | 1 - paddle/operators/scale_op.cc | 2 -- paddle/operators/scatter_op.cc | 4 ++-- paddle/operators/sequence_concat_op.cc | 2 -- paddle/operators/sequence_pool_op.cc | 2 -- paddle/operators/sequence_softmax_op.cc | 2 -- paddle/operators/sgd_op.cc | 1 - paddle/operators/sigmoid_cross_entropy_with_logits_op.cc | 2 -- paddle/operators/smooth_l1_loss_op.cc | 2 -- paddle/operators/softmax_op.cc | 2 -- paddle/operators/softmax_with_cross_entropy_op.cc | 4 ++-- paddle/operators/split_op.cc | 1 - paddle/operators/squared_l2_distance_op.cc | 2 -- paddle/operators/sum_op.cc | 1 - paddle/operators/top_k_op.cc | 1 - paddle/operators/transpose_op.cc | 2 -- paddle/operators/uniform_random_op.cc | 2 +- 53 files changed, 15 insertions(+), 90 deletions(-) diff --git a/paddle/operators/accuracy_op.cc b/paddle/operators/accuracy_op.cc index c5fb113e0f..037bb49abc 100644 --- a/paddle/operators/accuracy_op.cc +++ b/paddle/operators/accuracy_op.cc @@ -21,7 +21,6 @@ class AccuracyOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext *ctx) const override { PADDLE_ENFORCE(ctx->HasInput("Inference"), "Input(Inference) of AccuracyOp should not be null."); diff --git a/paddle/operators/activation_op.cc b/paddle/operators/activation_op.cc index 84c3775b4f..ee4f9b0ef2 100644 --- a/paddle/operators/activation_op.cc +++ b/paddle/operators/activation_op.cc @@ -21,7 +21,6 @@ class ActivationOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext *ctx) const override { ctx->SetOutputDim("Y", ctx->GetInputDim("X")); ctx->ShareLoD("X", /*->*/ "Y"); @@ -32,7 +31,6 @@ class ActivationOpGrad : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext *ctx) const override { ctx->SetOutputDim(framework::GradVarName("X"), ctx->GetInputDim("Y")); } diff --git a/paddle/operators/adadelta_op.cc b/paddle/operators/adadelta_op.cc index cf1bca1658..24e419b532 100644 --- a/paddle/operators/adadelta_op.cc +++ b/paddle/operators/adadelta_op.cc @@ -21,7 +21,6 @@ class AdadeltaOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext *ctx) const override { PADDLE_ENFORCE(ctx->HasInput("Param"), "Input(Param) of AdadeltaOp should not be null."); diff --git a/paddle/operators/adagrad_op.cc b/paddle/operators/adagrad_op.cc index a17747efb7..bc081f87dc 100644 --- a/paddle/operators/adagrad_op.cc +++ b/paddle/operators/adagrad_op.cc @@ -21,7 +21,6 @@ class AdagradOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext *ctx) const override { PADDLE_ENFORCE(ctx->HasInput("Param"), "Input(Param) of AdagradOp should not be null."); diff --git a/paddle/operators/adam_op.cc b/paddle/operators/adam_op.cc index 293b37b775..e3db70ea12 100644 --- a/paddle/operators/adam_op.cc +++ b/paddle/operators/adam_op.cc @@ -21,7 +21,6 @@ class AdamOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext *ctx) const override { PADDLE_ENFORCE(ctx->HasInput("Param"), "Input(Param) of AdamOp should not be null."); diff --git a/paddle/operators/adamax_op.cc b/paddle/operators/adamax_op.cc index 5cf727742c..e848333ef8 100644 --- a/paddle/operators/adamax_op.cc +++ b/paddle/operators/adamax_op.cc @@ -21,7 +21,6 @@ class AdamaxOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext *ctx) const override { PADDLE_ENFORCE(ctx->HasInput("Param"), "Input(Param) of AdamaxOp should not be null."); diff --git a/paddle/operators/clip_op.cc b/paddle/operators/clip_op.cc index 3e9b0d82ba..2d029394dd 100644 --- a/paddle/operators/clip_op.cc +++ b/paddle/operators/clip_op.cc @@ -21,7 +21,6 @@ class ClipOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) of ClipOp should not be null."); @@ -60,7 +59,6 @@ class ClipOpGrad : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) should not be null"); PADDLE_ENFORCE(ctx->HasInput(framework::GradVarName("Out")), diff --git a/paddle/operators/concat_op.cc b/paddle/operators/concat_op.cc index 235c4449ac..e11e51b458 100644 --- a/paddle/operators/concat_op.cc +++ b/paddle/operators/concat_op.cc @@ -23,7 +23,6 @@ class ConcatOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext *ctx) const override { PADDLE_ENFORCE_GE(ctx->Inputs("X").size(), 1UL, "Inputs(X) of ConcatOp should be empty.") @@ -82,7 +81,6 @@ class ConcatOpGrad : public framework::OperatorWithKernel { const framework::AttributeMap &attrs) : OperatorWithKernel(type, inputs, outputs, attrs) {} - protected: void InferShape(framework::InferShapeContext *ctx) const override { ctx->SetOutputsDim(framework::GradVarName("X"), ctx->GetInputsDim("X")); } diff --git a/paddle/operators/conv2d_op.h b/paddle/operators/conv2d_op.h index 7ebdbe81cb..bd1734879e 100644 --- a/paddle/operators/conv2d_op.h +++ b/paddle/operators/conv2d_op.h @@ -44,7 +44,6 @@ class Conv2DOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override; }; @@ -52,7 +51,6 @@ class Conv2DOpGrad : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override; }; diff --git a/paddle/operators/conv_shift_op.cc b/paddle/operators/conv_shift_op.cc index e1e321ed5f..6156a2d6af 100644 --- a/paddle/operators/conv_shift_op.cc +++ b/paddle/operators/conv_shift_op.cc @@ -27,7 +27,6 @@ class ConvShiftOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext *ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) should be not null."); PADDLE_ENFORCE(ctx->HasInput("Y"), "Input(Y) should be not null."); @@ -54,7 +53,6 @@ class ConvShiftGradOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext *ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) should be not null."); PADDLE_ENFORCE(ctx->HasInput("Y"), "Input(Y) should be not null."); diff --git a/paddle/operators/cos_sim_op.cc b/paddle/operators/cos_sim_op.cc index 2b4c4b9c45..55f69fb03a 100644 --- a/paddle/operators/cos_sim_op.cc +++ b/paddle/operators/cos_sim_op.cc @@ -23,7 +23,6 @@ class CosSimOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { // notnull check PADDLE_ENFORCE(ctx->HasInput("X"), @@ -97,7 +96,6 @@ class CosSimOpGrad : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { // notnull check PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) must not be null."); diff --git a/paddle/operators/crop_op.cc b/paddle/operators/crop_op.cc index a1424993cc..a994d91676 100644 --- a/paddle/operators/crop_op.cc +++ b/paddle/operators/crop_op.cc @@ -24,7 +24,6 @@ class CropOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) of CropOp should not be null."); @@ -114,7 +113,6 @@ class CropOpGrad : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) should not be null"); PADDLE_ENFORCE(ctx->HasInput(framework::GradVarName("Out")), diff --git a/paddle/operators/cross_entropy_op.cc b/paddle/operators/cross_entropy_op.cc index 923ae5be6c..a865991db3 100644 --- a/paddle/operators/cross_entropy_op.cc +++ b/paddle/operators/cross_entropy_op.cc @@ -21,7 +21,6 @@ class CrossEntropyOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) should be not null."); PADDLE_ENFORCE(ctx->HasInput("Label"), "Input(Label) should be not null."); @@ -48,6 +47,7 @@ class CrossEntropyOp : public framework::OperatorWithKernel { ctx->ShareLoD("X", /*->*/ "Y"); } + protected: // CrossEntropy's data type just determined by "X" framework::DataType IndicateDataType( const framework::ExecutionContext& ctx) const override { @@ -59,7 +59,6 @@ class CrossEntropyGradientOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) should be not null."); PADDLE_ENFORCE(ctx->HasInput("Label"), "Input(Label) should be not null."); @@ -94,6 +93,7 @@ class CrossEntropyGradientOp : public framework::OperatorWithKernel { ctx->SetOutputDim(framework::GradVarName("X"), x_dims); } + protected: // CrossEntropy's data type just determined by "X" framework::DataType IndicateDataType( const framework::ExecutionContext& ctx) const override { diff --git a/paddle/operators/decayed_adagrad_op.cc b/paddle/operators/decayed_adagrad_op.cc index 7f583f18c8..17b394aa07 100644 --- a/paddle/operators/decayed_adagrad_op.cc +++ b/paddle/operators/decayed_adagrad_op.cc @@ -21,7 +21,6 @@ class DecayedAdagradOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext *ctx) const override { PADDLE_ENFORCE(ctx->HasInput("Param"), "Input(Param) of DecayedAdagradOp should not be null."); diff --git a/paddle/operators/dropout_op.cc b/paddle/operators/dropout_op.cc index 708ccfa0bf..29858c9083 100644 --- a/paddle/operators/dropout_op.cc +++ b/paddle/operators/dropout_op.cc @@ -23,7 +23,6 @@ class DropoutOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) must not be null."); PADDLE_ENFORCE_GE(ctx->Attrs().Get("dropout_prob"), 0); @@ -69,7 +68,6 @@ class DropoutOpGrad : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE_EQ(ctx->Attrs().Get("is_training"), 1, "GradOp is only callable when is_training is true"); diff --git a/paddle/operators/elementwise_op.h b/paddle/operators/elementwise_op.h index 66f1910a47..fce4b24a22 100644 --- a/paddle/operators/elementwise_op.h +++ b/paddle/operators/elementwise_op.h @@ -23,7 +23,6 @@ class ElementwiseOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: using Tensor = framework::Tensor; void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), @@ -105,7 +104,6 @@ class ElementwiseOpGrad : public framework::OperatorWithKernel { using framework::OperatorWithKernel::OperatorWithKernel; using Tensor = framework::Tensor; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) should not be null"); PADDLE_ENFORCE(ctx->HasInput("Y"), "Input(Y) should not be null"); diff --git a/paddle/operators/fill_constant_op.cc b/paddle/operators/fill_constant_op.cc index a56832e202..0438d4d085 100644 --- a/paddle/operators/fill_constant_op.cc +++ b/paddle/operators/fill_constant_op.cc @@ -21,7 +21,6 @@ class FillConstantOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext *ctx) const override { PADDLE_ENFORCE(ctx->HasOutput("Out"), "Output(Out) of FillConstantOp should not be null."); @@ -33,6 +32,7 @@ class FillConstantOp : public framework::OperatorWithKernel { ctx->SetOutputDim("Out", dims); } + protected: framework::DataType IndicateDataType( const framework::ExecutionContext &ctx) const override { return static_cast(ctx.Attr("data_type")); diff --git a/paddle/operators/fill_zeros_like_op.cc b/paddle/operators/fill_zeros_like_op.cc index 4c70b9a36b..ed529ac40a 100644 --- a/paddle/operators/fill_zeros_like_op.cc +++ b/paddle/operators/fill_zeros_like_op.cc @@ -21,7 +21,6 @@ class FillZerosLikeOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext *ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) of FillZerosLikeOp should not be null."); diff --git a/paddle/operators/gather_op.cc b/paddle/operators/gather_op.cc index fb99c6c016..f6c7f472da 100644 --- a/paddle/operators/gather_op.cc +++ b/paddle/operators/gather_op.cc @@ -22,7 +22,6 @@ class GatherOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) of GatherOp should not be null."); @@ -40,6 +39,7 @@ class GatherOp : public framework::OperatorWithKernel { ctx->SetOutputDim("Out", output_dims); } + protected: framework::DataType IndicateDataType( const framework::ExecutionContext& ctx) const override { return framework::ToDataType(ctx.Input("X")->type()); @@ -50,11 +50,11 @@ class GatherGradOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { ctx->SetOutputDim(framework::GradVarName("X"), ctx->GetInputDim("X")); } + protected: framework::DataType IndicateDataType( const framework::ExecutionContext& ctx) const override { return framework::ToDataType(ctx.Input("X")->type()); diff --git a/paddle/operators/gaussian_random_op.cc b/paddle/operators/gaussian_random_op.cc index ca7fb38505..f59f497d9f 100644 --- a/paddle/operators/gaussian_random_op.cc +++ b/paddle/operators/gaussian_random_op.cc @@ -42,7 +42,6 @@ class GaussianRandomOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasOutput("Out"), "Output(Out) of GaussianRandomOp should not be null."); @@ -57,6 +56,7 @@ class GaussianRandomOp : public framework::OperatorWithKernel { ctx->SetOutputDim("Out", framework::make_ddim(temp)); } + protected: framework::DataType IndicateDataType( const framework::ExecutionContext& ctx) const override { return static_cast(Attr("data_type")); diff --git a/paddle/operators/gru_unit_op.cc b/paddle/operators/gru_unit_op.cc index 24f84597cd..72dd841c85 100644 --- a/paddle/operators/gru_unit_op.cc +++ b/paddle/operators/gru_unit_op.cc @@ -23,7 +23,6 @@ class GRUUnitOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput("Input"), "Input(%s) of GRUUnitOp should not be null.", "Input"); @@ -131,7 +130,6 @@ class GRUUnitGradOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput("Input"), "Input(%s) of GRUUnitGradOp should not be null.", "Input"); diff --git a/paddle/operators/lookup_table_op.cc b/paddle/operators/lookup_table_op.cc index 3f8d4ab857..b88cd14d78 100644 --- a/paddle/operators/lookup_table_op.cc +++ b/paddle/operators/lookup_table_op.cc @@ -21,7 +21,6 @@ class LookupTableOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput("W"), "Input(W) of LookupTableOp should not be null."); @@ -37,6 +36,7 @@ class LookupTableOp : public framework::OperatorWithKernel { ctx->ShareLoD("Ids", /*->*/ "Out"); } + protected: framework::DataType IndicateDataType( const framework::ExecutionContext& ctx) const override { return framework::ToDataType(ctx.Input("W")->type()); @@ -69,12 +69,12 @@ class LookupTableOpGrad : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { auto table_dims = ctx->GetInputDim("W"); ctx->SetOutputDim(framework::GradVarName("W"), table_dims); } + protected: framework::DataType IndicateDataType( const framework::ExecutionContext& ctx) const override { return framework::ToDataType(ctx.Input("W")->type()); diff --git a/paddle/operators/lstm_unit_op.cc b/paddle/operators/lstm_unit_op.cc index 13a45ec246..5d63017208 100644 --- a/paddle/operators/lstm_unit_op.cc +++ b/paddle/operators/lstm_unit_op.cc @@ -21,7 +21,6 @@ class LstmUnitOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) of LSTM should not be null."); PADDLE_ENFORCE(ctx->HasInput("C_prev"), @@ -76,7 +75,6 @@ class LstmUnitGradOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput(framework::GradVarName("C")), "Input(C@GRAD) should not be null"); diff --git a/paddle/operators/margin_rank_loss_op.cc b/paddle/operators/margin_rank_loss_op.cc index 5be61dfec3..638a99addc 100644 --- a/paddle/operators/margin_rank_loss_op.cc +++ b/paddle/operators/margin_rank_loss_op.cc @@ -21,7 +21,6 @@ class MarginRankLossOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext *ctx) const override { // input check PADDLE_ENFORCE(ctx->HasInput("Label"), "Input(Label) shouldn't be null."); @@ -94,7 +93,6 @@ class MarginRankLossGradOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext *ctx) const override { PADDLE_ENFORCE(ctx->HasInput("Label"), "Input(Label) shouldn't be null."); PADDLE_ENFORCE(ctx->HasInput("X1"), "Input(X1) shouldn't be null."); diff --git a/paddle/operators/mean_op.cc b/paddle/operators/mean_op.cc index 441543049f..9556fdf731 100644 --- a/paddle/operators/mean_op.cc +++ b/paddle/operators/mean_op.cc @@ -21,7 +21,6 @@ class MeanOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) of MeanOp should not be null."); @@ -46,7 +45,6 @@ class MeanGradOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { ctx->SetOutputDim(framework::GradVarName("X"), ctx->GetInputDim("X")); } diff --git a/paddle/operators/minus_op.cc b/paddle/operators/minus_op.cc index d7fd2f901b..f7943e99ac 100644 --- a/paddle/operators/minus_op.cc +++ b/paddle/operators/minus_op.cc @@ -25,7 +25,6 @@ class MinusOp : public framework::OperatorWithKernel { const framework::AttributeMap &attrs) : OperatorWithKernel(type, inputs, outputs, attrs) {} - protected: void InferShape(framework::InferShapeContext *ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) of MinusOp should not be null."); diff --git a/paddle/operators/modified_huber_loss_op.cc b/paddle/operators/modified_huber_loss_op.cc index 6522327fdc..7b9e952895 100644 --- a/paddle/operators/modified_huber_loss_op.cc +++ b/paddle/operators/modified_huber_loss_op.cc @@ -21,7 +21,6 @@ class ModifiedHuberLossOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "X must be initialized."); PADDLE_ENFORCE(ctx->HasInput("Y"), "Y must be initialized."); @@ -73,7 +72,6 @@ class ModifiedHuberLossGradOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "X must be initialized."); PADDLE_ENFORCE(ctx->HasInput("Y"), "Y must be initialized."); diff --git a/paddle/operators/mul_op.cc b/paddle/operators/mul_op.cc index ec0683d887..943f81e949 100644 --- a/paddle/operators/mul_op.cc +++ b/paddle/operators/mul_op.cc @@ -23,7 +23,6 @@ class MulOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) of MulOp should not be null."); PADDLE_ENFORCE(ctx->HasInput("Y"), "Input(Y) of MulOp should not be null."); @@ -96,7 +95,6 @@ class MulOpGrad : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) should not be null"); PADDLE_ENFORCE(ctx->HasInput("Y"), "Input(Y) should not be null"); diff --git a/paddle/operators/multiplex_op.cc b/paddle/operators/multiplex_op.cc index 051051b051..4d86769026 100644 --- a/paddle/operators/multiplex_op.cc +++ b/paddle/operators/multiplex_op.cc @@ -23,7 +23,6 @@ class MultiplexOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput("Ids"), "Input(Ids) shouldn't be null."); PADDLE_ENFORCE(!ctx->Inputs("X").empty(), @@ -51,6 +50,7 @@ class MultiplexOp : public framework::OperatorWithKernel { ctx->SetOutputDim("Out", in_dim); } + protected: framework::DataType IndicateDataType( const framework::ExecutionContext& ctx) const override { return framework::ToDataType(ctx.MultiInput("X")[0]->type()); @@ -89,7 +89,6 @@ class MultiplexGradOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(!ctx->Inputs("X").empty(), "Input(X) should not be null."); PADDLE_ENFORCE(!ctx->Outputs(framework::GradVarName("X")).empty(), @@ -105,6 +104,7 @@ class MultiplexGradOp : public framework::OperatorWithKernel { ctx->SetOutputsDim(framework::GradVarName("X"), d_ins); } + protected: framework::DataType IndicateDataType( const framework::ExecutionContext& ctx) const override { return framework::ToDataType(ctx.MultiInput("X")[0]->type()); diff --git a/paddle/operators/pad_op.cc b/paddle/operators/pad_op.cc index 2f26ada85e..73a0b8baff 100644 --- a/paddle/operators/pad_op.cc +++ b/paddle/operators/pad_op.cc @@ -23,7 +23,6 @@ class PadOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) of PadOp should not be null."); PADDLE_ENFORCE(ctx->HasOutput("Out"), @@ -97,7 +96,6 @@ class PadOpGrad : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) should not be null"); PADDLE_ENFORCE(ctx->HasInput(framework::GradVarName("Out")), diff --git a/paddle/operators/pool_op.h b/paddle/operators/pool_op.h index f6169e05b3..ada9565019 100644 --- a/paddle/operators/pool_op.h +++ b/paddle/operators/pool_op.h @@ -28,7 +28,6 @@ class PoolOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override; }; @@ -36,7 +35,6 @@ class PoolOpGrad : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override; }; diff --git a/paddle/operators/pool_with_index_op.cc b/paddle/operators/pool_with_index_op.cc index 8eee20c2a9..29d0322a27 100644 --- a/paddle/operators/pool_with_index_op.cc +++ b/paddle/operators/pool_with_index_op.cc @@ -27,7 +27,6 @@ class MaxPoolWithIndexOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext *ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "X(Input) of Pooling should not be null."); @@ -72,7 +71,6 @@ class MaxPoolWithIndexOpGrad : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext *ctx) const override { PADDLE_ENFORCE(ctx->HasInput("Mask"), "Input(Mask) must not be null."); PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) must not be null."); diff --git a/paddle/operators/prelu_op.cc b/paddle/operators/prelu_op.cc index 166fe26824..eef2e34eaa 100644 --- a/paddle/operators/prelu_op.cc +++ b/paddle/operators/prelu_op.cc @@ -25,7 +25,6 @@ class PReluOp : public framework::OperatorWithKernel { const framework::AttributeMap &attrs) : OperatorWithKernel(type, inputs, outputs, attrs) {} - protected: void InferShape(framework::InferShapeContext *ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) should not be null"); PADDLE_ENFORCE(ctx->HasInput("Alpha"), "Input(Alpha) should not be null"); @@ -62,7 +61,6 @@ class PReluGradOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext *ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) must not be null."); PADDLE_ENFORCE(ctx->HasInput(framework::GradVarName("Out")), diff --git a/paddle/operators/rank_loss_op.cc b/paddle/operators/rank_loss_op.cc index e0abbc4db1..17ef2b1d01 100644 --- a/paddle/operators/rank_loss_op.cc +++ b/paddle/operators/rank_loss_op.cc @@ -24,7 +24,6 @@ class RankLossOp : public framework::OperatorWithKernel { const framework::AttributeMap &attrs) : OperatorWithKernel(type, inputs, outputs, attrs) {} - protected: void InferShape(framework::InferShapeContext *ctx) const override { // input check PADDLE_ENFORCE(ctx->HasInput("Label"), "Input(Label) shouldn't be null"); @@ -89,7 +88,6 @@ class RankLossGradOp : public framework::OperatorWithKernel { const framework::AttributeMap &attrs) : OperatorWithKernel(type, inputs, outputs, attrs) {} - protected: void InferShape(framework::InferShapeContext *ctx) const override { PADDLE_ENFORCE(ctx->HasInput("Label"), "Input(Label) shouldn't be null."); PADDLE_ENFORCE(ctx->HasInput("Left"), "Input(Left) shouldn't be null."); diff --git a/paddle/operators/reduce_op.cc b/paddle/operators/reduce_op.cc index 005f88b57c..5e878353ce 100644 --- a/paddle/operators/reduce_op.cc +++ b/paddle/operators/reduce_op.cc @@ -23,7 +23,6 @@ class ReduceOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext *ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) of ReduceOp should not be null."); @@ -57,7 +56,6 @@ class ReduceGradOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext *ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) should not be null."); PADDLE_ENFORCE(ctx->HasInput(framework::GradVarName("Out")), diff --git a/paddle/operators/reshape_op.cc b/paddle/operators/reshape_op.cc index 3cd54930a0..a8eb8d45ee 100644 --- a/paddle/operators/reshape_op.cc +++ b/paddle/operators/reshape_op.cc @@ -25,7 +25,6 @@ class ReshapeOp : public framework::OperatorWithKernel { const framework::AttributeMap &attrs) : OperatorWithKernel(type, inputs, outputs, attrs) {} - protected: void InferShape(framework::InferShapeContext *ctx) const override { // input check PADDLE_ENFORCE(ctx->HasInput("X"), @@ -93,7 +92,6 @@ class ReshapeGradOp : public framework::OperatorWithKernel { const framework::AttributeMap &attrs) : OperatorWithKernel(type, inputs, outputs, attrs) {} - protected: void InferShape(framework::InferShapeContext *ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) shouldn't be null."); PADDLE_ENFORCE(ctx->HasInput(framework::GradVarName("Out")), diff --git a/paddle/operators/rmsprop_op.cc b/paddle/operators/rmsprop_op.cc index ada6f2bc3c..fd5567a365 100644 --- a/paddle/operators/rmsprop_op.cc +++ b/paddle/operators/rmsprop_op.cc @@ -21,7 +21,6 @@ class RmspropOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext *ctx) const override { PADDLE_ENFORCE(ctx->HasInput("Param"), "Input(Param) of RmspropOp should not be null."); diff --git a/paddle/operators/scale_op.cc b/paddle/operators/scale_op.cc index ac297da6b7..7f1a21bea7 100644 --- a/paddle/operators/scale_op.cc +++ b/paddle/operators/scale_op.cc @@ -25,7 +25,6 @@ class ScaleOp : public framework::OperatorWithKernel { const framework::AttributeMap &attrs) : OperatorWithKernel(type, inputs, outputs, attrs) {} - protected: void InferShape(framework::InferShapeContext *ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) of ScaleOp should not be null."); @@ -56,7 +55,6 @@ class ScaleGradMaker : public framework::SingleGradOpDescMaker { public: using framework::SingleGradOpDescMaker::SingleGradOpDescMaker; - protected: std::unique_ptr Apply() const override { auto *grad_op = new framework::OpDescBind(); grad_op->SetType("scale"); diff --git a/paddle/operators/scatter_op.cc b/paddle/operators/scatter_op.cc index fbea01a8db..62e6c70b45 100644 --- a/paddle/operators/scatter_op.cc +++ b/paddle/operators/scatter_op.cc @@ -22,7 +22,6 @@ class ScatterOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput("Ref"), "Input(Ref) of ScatterOp should not be null."); @@ -49,6 +48,7 @@ class ScatterOp : public framework::OperatorWithKernel { ctx->SetOutputDim("Out", ref_dims); } + protected: framework::DataType IndicateDataType( const framework::ExecutionContext& ctx) const override { return framework::ToDataType(ctx.Input("Ref")->type()); @@ -59,13 +59,13 @@ class ScatterGradOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { ctx->SetOutputDim(framework::GradVarName("Updates"), ctx->GetInputDim("Updates")); ctx->SetOutputDim(framework::GradVarName("Ref"), ctx->GetInputDim("Ref")); } + protected: framework::DataType IndicateDataType( const framework::ExecutionContext& ctx) const override { return framework::ToDataType(ctx.Input("Ref")->type()); diff --git a/paddle/operators/sequence_concat_op.cc b/paddle/operators/sequence_concat_op.cc index 287fb1942e..1fce96cdfe 100644 --- a/paddle/operators/sequence_concat_op.cc +++ b/paddle/operators/sequence_concat_op.cc @@ -21,7 +21,6 @@ class SequenceConcatOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInputs("X"), "Inputs(X) of SequenceConcatOp should not be null."); @@ -105,7 +104,6 @@ class SequenceConcatGradOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput(framework::GradVarName("Out")), "The gradient of Out should not be null."); diff --git a/paddle/operators/sequence_pool_op.cc b/paddle/operators/sequence_pool_op.cc index 8dc4a59ba8..e3f5d509a8 100644 --- a/paddle/operators/sequence_pool_op.cc +++ b/paddle/operators/sequence_pool_op.cc @@ -21,7 +21,6 @@ class SequencePoolOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) of SequencePoolOp should not be null."); @@ -72,7 +71,6 @@ class SequencePoolGradOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput(framework::GradVarName("Out")), "Gradient of Out should not be null."); diff --git a/paddle/operators/sequence_softmax_op.cc b/paddle/operators/sequence_softmax_op.cc index ea217ba459..c891ab1fdc 100644 --- a/paddle/operators/sequence_softmax_op.cc +++ b/paddle/operators/sequence_softmax_op.cc @@ -21,7 +21,6 @@ class SequenceSoftmaxOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) of SequenceSoftmaxOp should not be null."); @@ -66,7 +65,6 @@ class SequenceSoftmaxGradOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput("Out"), "Input(Out) of SequenceSoftmaxGradOp should not be null."); diff --git a/paddle/operators/sgd_op.cc b/paddle/operators/sgd_op.cc index 2a6a162a02..0f78eeab9b 100644 --- a/paddle/operators/sgd_op.cc +++ b/paddle/operators/sgd_op.cc @@ -21,7 +21,6 @@ class SGDOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext *ctx) const override { PADDLE_ENFORCE(ctx->HasInput("Param"), "Input(Param) of SGDOp should not be null."); diff --git a/paddle/operators/sigmoid_cross_entropy_with_logits_op.cc b/paddle/operators/sigmoid_cross_entropy_with_logits_op.cc index b6653e1cc7..e781c8db20 100644 --- a/paddle/operators/sigmoid_cross_entropy_with_logits_op.cc +++ b/paddle/operators/sigmoid_cross_entropy_with_logits_op.cc @@ -23,7 +23,6 @@ class SigmoidCrossEntropyWithLogitsOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) should be not null."); PADDLE_ENFORCE(ctx->HasInput("Labels"), @@ -52,7 +51,6 @@ class SigmoidCrossEntropyWithLogitsGradOp public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) should be not null."); PADDLE_ENFORCE(ctx->HasInput("Labels"), diff --git a/paddle/operators/smooth_l1_loss_op.cc b/paddle/operators/smooth_l1_loss_op.cc index 91391dc945..a4f0f37764 100644 --- a/paddle/operators/smooth_l1_loss_op.cc +++ b/paddle/operators/smooth_l1_loss_op.cc @@ -21,7 +21,6 @@ class SmoothL1LossOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "X must be initialized."); PADDLE_ENFORCE(ctx->HasInput("Y"), "Y must be initialized."); @@ -93,7 +92,6 @@ class SmoothL1LossGradOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { auto in_dims = ctx->GetInputDim("X"); auto out_dims = ctx->GetInputDim(framework::GradVarName("Out")); diff --git a/paddle/operators/softmax_op.cc b/paddle/operators/softmax_op.cc index 4c131ed44d..00fd0b32a9 100644 --- a/paddle/operators/softmax_op.cc +++ b/paddle/operators/softmax_op.cc @@ -21,7 +21,6 @@ class SoftmaxOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) of SoftmaxOp should not be null."); @@ -68,7 +67,6 @@ class SoftmaxOpGrad : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput("Y"), "Input(Y) should be not null."); PADDLE_ENFORCE(ctx->HasInput(framework::GradVarName("Y")), diff --git a/paddle/operators/softmax_with_cross_entropy_op.cc b/paddle/operators/softmax_with_cross_entropy_op.cc index 9121609f10..942fbb42df 100644 --- a/paddle/operators/softmax_with_cross_entropy_op.cc +++ b/paddle/operators/softmax_with_cross_entropy_op.cc @@ -82,7 +82,6 @@ class SoftmaxWithCrossEntropyOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput("Logits"), "Input(Logits) should be not null."); @@ -117,6 +116,7 @@ class SoftmaxWithCrossEntropyOp : public framework::OperatorWithKernel { ctx->ShareLoD("Logits", /*->*/ "Loss"); } + protected: framework::DataType IndicateDataType( const framework::ExecutionContext& ctx) const override { return framework::ToDataType(ctx.Input("Logits")->type()); @@ -127,7 +127,6 @@ class SoftmaxWithCrossEntropyOpGrad : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput(framework::GradVarName("Loss")), "Input(Loss@Grad) should not be null."); @@ -156,6 +155,7 @@ class SoftmaxWithCrossEntropyOpGrad : public framework::OperatorWithKernel { ctx->GetInputDim("Softmax")); } + protected: framework::DataType IndicateDataType( const framework::ExecutionContext& ctx) const override { return framework::ToDataType( diff --git a/paddle/operators/split_op.cc b/paddle/operators/split_op.cc index d5dd4df2a2..4a6c50f797 100644 --- a/paddle/operators/split_op.cc +++ b/paddle/operators/split_op.cc @@ -23,7 +23,6 @@ class SplitOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext *ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) of SplitOp should not be null."); diff --git a/paddle/operators/squared_l2_distance_op.cc b/paddle/operators/squared_l2_distance_op.cc index cce4e527c3..e360c19b47 100644 --- a/paddle/operators/squared_l2_distance_op.cc +++ b/paddle/operators/squared_l2_distance_op.cc @@ -21,7 +21,6 @@ class SquaredL2DistanceOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) of SquaredL2DistanceOp should not be null."); @@ -85,7 +84,6 @@ class SquaredL2DistanceGradOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput(framework::GradVarName("Out")), "Gradient of Out should not be null"); diff --git a/paddle/operators/sum_op.cc b/paddle/operators/sum_op.cc index 573487b835..5214a8413e 100644 --- a/paddle/operators/sum_op.cc +++ b/paddle/operators/sum_op.cc @@ -21,7 +21,6 @@ class SumOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInputs("X"), "Inputs(X) should not be null"); auto x_dims = ctx->GetInputsDim("X"); diff --git a/paddle/operators/top_k_op.cc b/paddle/operators/top_k_op.cc index c954819912..d5c2c91a5f 100644 --- a/paddle/operators/top_k_op.cc +++ b/paddle/operators/top_k_op.cc @@ -21,7 +21,6 @@ class TopkOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext *ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) of TopkOp should not be null."); diff --git a/paddle/operators/transpose_op.cc b/paddle/operators/transpose_op.cc index 1101bbe3ef..d785e57c83 100644 --- a/paddle/operators/transpose_op.cc +++ b/paddle/operators/transpose_op.cc @@ -23,7 +23,6 @@ class TransposeOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) should not be null"); PADDLE_ENFORCE(ctx->HasOutput("Out"), "Output(Out) should not be null"); @@ -92,7 +91,6 @@ class TransposeOpGrad : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) should not be null"); PADDLE_ENFORCE(ctx->HasInput(framework::GradVarName("Out")), diff --git a/paddle/operators/uniform_random_op.cc b/paddle/operators/uniform_random_op.cc index 75928f1ec8..612bdd70db 100644 --- a/paddle/operators/uniform_random_op.cc +++ b/paddle/operators/uniform_random_op.cc @@ -46,7 +46,6 @@ class UniformRandomOp : public framework::OperatorWithKernel { public: using framework::OperatorWithKernel::OperatorWithKernel; - protected: void InferShape(framework::InferShapeContext* ctx) const override { PADDLE_ENFORCE(ctx->HasOutput("Out"), "Output(Out) of UniformRandomOp should not be null."); @@ -63,6 +62,7 @@ class UniformRandomOp : public framework::OperatorWithKernel { ctx->SetOutputDim("Out", framework::make_ddim(temp)); } + protected: framework::DataType IndicateDataType( const framework::ExecutionContext& ctx) const override { return static_cast(Attr("data_type")); From f43b1a90d8d0dccb54b4243f33df87660f7679d6 Mon Sep 17 00:00:00 2001 From: Yu Yang Date: Mon, 16 Oct 2017 19:32:10 -0700 Subject: [PATCH 36/36] Design Doc: infer_var_type (#4795) --- doc/design/infer_var_type.md | 78 ++++++++++++++++++++++++++++++++++++ 1 file changed, 78 insertions(+) create mode 100644 doc/design/infer_var_type.md diff --git a/doc/design/infer_var_type.md b/doc/design/infer_var_type.md new file mode 100644 index 0000000000..d9d5397bec --- /dev/null +++ b/doc/design/infer_var_type.md @@ -0,0 +1,78 @@ +# Design Doc: InferVarType + +## The Problem Posed + +The variable in our design can hold variant types. Such as `LoDTensor` and `SelectedRows`. An operator should be able to inference the variable types of its output. + +For example, a `lookup table` operator takes two `LoDTensor`; one is a float tensor as the embedding table, the other is an int tensor as word ID. The gradient operator of `lookup table` will generate a `SelectedRows` as its output. A `sum` operator can take both `LoDTensor` and `SelectedRows` as its inputs and will generate a `LoDTensor` if any of its inputs is `LoDTensor`, otherwise, the `sum` operator will generate `SelectedRows` as its output. + +The variable type will be constant at runtime. Every variable's type can either be set by the user (input data and parameter) or be inferred by the operator in compile time. + +## Proposed Solution + +The `InferVarType` is a compile-time function which is registered to each operator. The inferface of that function is: + + +```c++ +using InferVarTypeFN = std::function< + void (const OpDescBind& /*op_desc*/, BlockDescBind* /*block*/)>; +``` + +It takes an operator description as its input and will write the output variable type and store them in block description. + +The `InferVarTypeFN` will be registered in `OpInfo`, to replace `infer_var_type_` field. The `OpInfo` should be + +```cpp +struct OpInfo { + InferVarTypeFN infer_var_type_; + ... +}; +``` + +The default `InferVarType` will set output type as `LoDTensor`. It can be done by `GetInferVarType()`. + +```cpp +void DefaultInferVarType(const OpDescBind& op_desc, BlockDescBind* block) { + // set the output type of variable as `LoDTensor`. + // ... +} + +struct OpInfo { + InferVarTypeFN infer_var_type_; + InferVarTypeFN GetInferVarType() const { + if (infer_var_type_) { + return infer_var_type_; + } else { + return DefaultInferVarType; + } + } +}; +``` + +## Register InferVarType + +We provide a thin base class for registering an `InferVarTypeFN`. To use a base class will ease the implementation of registry since we can detect the registry entry is an `InferVarTypeFN` or not. + +```cpp +class VarTypeInferer { +public: + virtual void operator()(const OpDescBind& op_desc, BlockDescBind* block) const = 0; +} +``` + +Operator developers can write the specialize `VarTypeInferer` as follow. + +```cpp +class SpecialVarTypeInferer : public VarTypeInferer { +public: + virtual void operator()(const OpDescBind& op_desc, BlockDescBind* block) const { + // .. own logic + } +} +``` + +Then user can register the `InferVarType` just like `GradOpDescMaker` and `OpInfoMaker`. + +``` +REGISTER_OPERATOR(some_op, OpType, SpecialVarTypeInferer, ...); +```