diff --git a/cmake/configure.cmake b/cmake/configure.cmake index 51c3b918cc4ef4cf6c8052ccc14028a872309fcf..c1c93e17fd82ea048ba27b127b1527d9a8c9da41 100644 --- a/cmake/configure.cmake +++ b/cmake/configure.cmake @@ -49,11 +49,12 @@ if(NOT WITH_GOLANG) endif(NOT WITH_GOLANG) if(NOT WITH_GPU) - add_definitions(-DPADDLE_ONLY_CPU) add_definitions(-DHPPL_STUB_FUNC) list(APPEND CMAKE_CXX_SOURCE_FILE_EXTENSIONS cu) else() + add_definitions(-DPADDLE_WITH_CUDA) + FIND_PACKAGE(CUDA REQUIRED) if(${CUDA_VERSION_MAJOR} VERSION_LESS 7) diff --git a/doc/design/refactor/session.md b/doc/design/refactor/session.md new file mode 100644 index 0000000000000000000000000000000000000000..1d9a26683c14f54e3b5fe41675cd03b5620646b8 --- /dev/null +++ b/doc/design/refactor/session.md @@ -0,0 +1,180 @@ +# Design Doc: Session + +## Abstract + +The *session* object encapsulates the environment in which the +computation graph is executed. + +We will have the *local* session and *remote* session, they offer the +same [interface](#interface). The local session encapsulates the local +runtime environment and the remote session encapsulates the cluster +runtime environment. + +The local runtime environment contains: + +1. computation devices (i.e., CPU, GPU) handles, and +1. the [scope](../scope.md) which holds all variables. + +The remote runtime environment contains: + +1. computation devices (i.e., CPU and GPU on node 0, 1) in a cluster, + and +1. the distributed [scope](../scope.md) in a cluster which holds all + variables. + +The user can create a remote session on Paddle Cloud and evaluate the +computation graph with it. In this way, the user can control the +remote computation resource in a cluster from his local computer. + + +## Background + +The current design has an implicit global session in which +`paddle.eval()` is executed. The pain point is: + +Since the user is not able to explicitly switch between runtime +environments, the user cannot run a topology in two independent +environments. + +For example, in reinforcement learning, the user may want to have a +stale model for inference and a fresh model for training, and only +replace the stale model with the fresh model periodically. + +Furthermore, we have no concept that encapsulates a remote environment +that executes a computation graph. + +We need the session object to address above issues. + + +## Session + +A session is an object that owns the runtime environment. All +computations are executed through `session.eval()`. + + +### Interface + +```python +eval( + targets, + feed_dict=None, +) +``` + +Evaluates the target Operations or Variables in `targets`. + +- *targets*: the evaluation targets. Can be a single Operation or + Variable, or a list with the Operations or Variables as + elements. The value returned by `eval()` has the same shape as the + `target` argument. + + The PaddlePaddle program is represented by + the [ProgramDesc](../design/program.md), `eval()` will infer the + ProgramDesc from the given targets and run the PaddlePaddle + program. Please + see + [this graph](./distributed_architecture.md#local-training-architecture) for + the detailed illustration for the local session + and + [this graph](./distributed_architecture.md#distributed-training-architecture) for + the detailed illustration for the remote session. + +- *feed_dict*: a dictionary that contains the tensors which override + the edges of the computation graph. + + feed_dict not only can provide the input data, it can override any + OP's input as well: + + ```python + a = pd.constant(2.0, name="a") + b = pd.variable(name="b") + c = pd.mul(a,b) + sess.eval(targets=c, feed_dict={"b":3.0}) # returns 6.0 + ``` + +```python +close() +``` + +Closes the session and releases the scope that the session owns. + + +### Create a Local Session + +```python +session( + devices=None +) +``` + +Creates a new session. One session owns one global scope, so creating +multiple sessions will create different scopes. + +- *devices*: a single `string` or a list of `string` of device names, + the corresponding devices will be the computation devices for + `eval()`. If not specified, all available devices (e.g., all GPUs) + will be used. The user doesn't need to specify the CPU device since + it will be always used. Multiple sessions can use the same device. + + +#### Example + +```Python +a = paddle.constant(1.0) +b = paddle.constant(2.0) +c = a + b +sess = paddle.session(devices=["gpu:0", "gpu:1", "fpga:0"]) +sess.eval(c) +sess.close() +``` + +### Create a Remote Session + +```python +create_cloud_job( + name, + num_trainer, + mem_per_trainer, + gpu_per_trainer, + cpu_per_trainer, + num_ps, + mem_per_ps, + cpu_per_ps, +) +``` + +Creates a Paddle Cloud job. Fails if the job name exists. + +```python +get_cloud_job( + name +) +``` + +Gets a Paddle Cloud job. + +```python +remote_session( + job +) +``` + +- *job*: the Paddle Cloud job. + +#### Example + +```Python +reader = paddle.reader.recordio("/pfs/home/peter/mnist-train-*") # data stored on Paddle Cloud +image = reader.column(0) +label = reader.column(1) +fc1 = paddle.op.fc(image, size=256, act="sigmoid") +fc2 = paddle.op.fc(fc1, size=10, act="softmax") +cost = paddle.op.cross_entropy(fc2, label) +opt = paddle.optimizer.sgd(cost) + +job = paddle.create_cloud_job("test", 3, "1G", 1, 1, 2, "1G", 1) +sess = paddle.remote_ession(job) +for i in range(1000): + sess.eval(opt) +sess.close() +``` diff --git a/doc/design/register_grad_op.md b/doc/design/register_grad_op.md index 12b04fb2713d8bb862c08ddb538ddf860bb2983c..3cf8a59446d244bb3a388b87b14273d9096c839a 100644 --- a/doc/design/register_grad_op.md +++ b/doc/design/register_grad_op.md @@ -33,22 +33,45 @@ The mapping relationship between an operator and its gradient operators is a fun ```cpp // (OpDesc) --> vector -using GradOpDescMaker = std::function(const OpDesc&)>; +std::function(const OpDescBind&)>; ``` -The function take a `OpDesc` of the forward operator and return one or many gradient operator descriptions. +The function takes an `OpDescBind` of the forward operator and returns one or many gradient operator descriptions. `OpDescBind` is a C++ wrapper for protobuf message `OpDesc` to manipulate `OpDesc` fast. The `GradOpDescMaker` will be registered in `OpInfo`, to replace `grad_op_type_` field. The `OpInfo` should be ```cpp struct OpInfo { - GradOpDescMaker grad_op_maker_; + std::function>(const OpDescBind&)> grad_op_maker_; ... }; ``` The `grad_op_maker_ ` is `nullptr` if the operator does not have associated gradient operators. +We propose a base class called `GradOpDescMakerBase` to let operator developers generate `Gradient Operators` easily. The public interface of that class is + +```cpp +class GradOpDescMakerBase { +public: + GradOpDescMakerBase(const OpDescBind& ); + virtual std::vector> operator()()const = 0; +}; +``` + +We can convert `GradOpDescMakerBase` to `std::function>(const OpDescBind&)>` by + +```cpp +using GradOpMaker = ...; +std::function(const OpDescBind&)> func; +func = [] (const OpDescBind& fwd_op) { + GradOpMaker maker(fwd_op); + return maker(); +}; +``` + +We can write many helper functions since the `GradOpDescMakerBase` is a class now. The basic helper functions get the variables of `Input`, `Output`, `InputGradient` and `OutputGradient` in the forwarding operator. + We should chagne register macros at the same time. In the current solution, there is no difference between forwarding operators and backward operators. So `REGISTER_OP` just register one operator. If the `REGISTER_OPERATOR ` contains `OpProtoAndCheckerMaker` and `GradOpDescMaker`, we just list them in the same macro. It can be done by a macro contains `__VA_ARGS__`. The user interface should be diff --git a/paddle/api/Util.cpp b/paddle/api/Util.cpp index d369df5d4e04b4a8d822db0e72a8051150868ce6..11bd05c09d1ecbbcec6b6596c16416c26635a072 100644 --- a/paddle/api/Util.cpp +++ b/paddle/api/Util.cpp @@ -47,7 +47,7 @@ bool isUsingGpu() { return FLAGS_use_gpu; } void setUseGpu(bool useGpu) { FLAGS_use_gpu = useGpu; } bool isGpuVersion() { -#ifdef PADDLE_ONLY_CPU +#ifndef PADDLE_WITH_CUDA return false; #else return true; diff --git a/paddle/capi/Matrix.cpp b/paddle/capi/Matrix.cpp index d898ebe2612d749ca261d35139d1cd45bd355eef..4547afaf1dc9af8bc7909a684db766fdd7b159c0 100644 --- a/paddle/capi/Matrix.cpp +++ b/paddle/capi/Matrix.cpp @@ -46,7 +46,7 @@ paddle_error paddle_matrix_set_row(paddle_matrix mat, if (rowID >= ptr->mat->getHeight()) return kPD_OUT_OF_RANGE; paddle::real* buf = ptr->mat->getRowBuf(rowID); size_t width = ptr->mat->getWidth(); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA hl_memcpy(buf, rowArray, sizeof(paddle::real) * width); #else std::copy(rowArray, rowArray + width, buf); diff --git a/paddle/framework/backward.cc b/paddle/framework/backward.cc index 0ec18de5b8a0e7cebdb91c30d2b45596b02dfa51..c0188c0e55a6db04a86d80c216fd162ad5259593 100644 --- a/paddle/framework/backward.cc +++ b/paddle/framework/backward.cc @@ -141,9 +141,26 @@ static std::unique_ptr BackwardRecursive( net->ops_[op_offset]->Rename(name, dup_outputs.back()); } // collect all the offset to append `add` op for each alias - insert_position.push_back( - {dup_op.back(), OpRegistry::CreateOp("add", {{"X", {dup_outputs}}}, - {{"Out", {name}}}, {})}); + // + // one variable is shared between multiple operators. + // insert add operator one by one, then add it to output + for (size_t output_idx = 0; output_idx < dup_outputs.size() - 1; + ++output_idx) { + auto insert_add_x = dup_outputs[output_idx]; + auto insert_add_y = dup_outputs[output_idx + 1]; + auto insert_add_out = name + "@SHARED@" + std::to_string(output_idx); + // first add op inserted + if (output_idx == dup_outputs.size() - 2) { + insert_add_out = name; + } + if (output_idx != 0) { + insert_add_y = name + "@SHARED@" + std::to_string(output_idx - 1); + } + insert_position.push_back( + {dup_op.back(), + OpRegistry::CreateOp("sum", {{"X", {insert_add_x, insert_add_y}}}, + {{"Out", {insert_add_out}}}, {})}); + } } // make sure the inserted `add` ops follow the BFS order. @@ -182,7 +199,8 @@ static std::unique_ptr BackwardRecursive( // process recurrent gradient op as a special operator. if (forwardOp.Type() == "recurrent") { - // NOTE clean up cycle call somewhere (RNN's stepnet constains itself), or + // NOTE clean up cycle call somewhere (RNN's stepnet constains itself), + // or // this will result in infinite loop. const auto& rnnop = *static_cast(&forwardOp); diff --git a/paddle/framework/backward_test.cc b/paddle/framework/backward_test.cc index 6932f5b989a3e21ebc44ec4fec9f5223f2547d7a..a36e7bde8c61c6a8bfc6eefd46ce5e09907476ef 100644 --- a/paddle/framework/backward_test.cc +++ b/paddle/framework/backward_test.cc @@ -133,15 +133,18 @@ class FillZeroOpMaker : public OpProtoAndCheckerMaker { } }; -class AddOpMaker : public OpProtoAndCheckerMaker { +class SumOpMaker : public framework::OpProtoAndCheckerMaker { public: - AddOpMaker(OpProto *proto, OpAttrChecker *op_checker) + SumOpMaker(framework::OpProto *proto, framework::OpAttrChecker *op_checker) : OpProtoAndCheckerMaker(proto, op_checker) { - AddInput("X", "x").AsDuplicable(); - AddOutput("Out", "out"); + AddInput("X", "the input tensors of sum operator.") + .AsDuplicable() + .NotInGradient(); + AddOutput("Out", "the output tensor of sum operator.").NotInGradient(); AddComment(""); } }; + } // namespace framework } // namespace paddle @@ -154,7 +157,7 @@ REGISTER_OP(mul, f::NOP, f::MulOpMaker, mul_grad, f::NOP); REGISTER_OP(sigmoid, f::NOP, f::SigmoidOpMaker, sigmoid_grad, f::NOP); REGISTER_OP_WITHOUT_GRADIENT(nograd, f::NOP, f::NoGradOpMaker); REGISTER_OP_WITHOUT_GRADIENT(fill_zeros_like, f::NOP, f::FillZeroOpMaker); -REGISTER_OP(add, f::NOP, f::AddOpMaker, add_grad, f::NOP); +REGISTER_OP(sum, f::NOP, f::SumOpMaker, sum_grad, f::NOP); REGISTER_OP_WITHOUT_GRADIENT(fc, f::FcOp, f::FcOpMaker); REGISTER_OP(many_output_op, f::NOP, f::ManyOutputOpMaker, many_output_op_grad, f::NOP); @@ -283,7 +286,7 @@ TEST(Backward, net_shared_weight) { ASSERT_TRUE(bwd->IsNetOp()); auto bwd_net = static_cast(bwd.get()); ASSERT_EQ(3UL, bwd_net->ops_.size()); - ASSERT_EQ("add", bwd_net->ops_[2]->Type()); + ASSERT_EQ("sum", bwd_net->ops_[2]->Type()); } TEST(Backward, op_register_grad_not_for_network) { diff --git a/paddle/framework/block_desc.h b/paddle/framework/block_desc.h index 1a1135bab44cd27bb7d784c3b486188aa40635e4..59513ede33ebb41acbeb0e1acab66be5947a9a3c 100644 --- a/paddle/framework/block_desc.h +++ b/paddle/framework/block_desc.h @@ -19,6 +19,7 @@ limitations under the License. */ #include #include "paddle/framework/op_desc.h" #include "paddle/framework/var_desc.h" +#include "paddle/platform/macros.h" namespace paddle { namespace framework { @@ -34,9 +35,6 @@ class BlockDescBind { BlockDescBind(ProgramDescBind *prog, BlockDesc *desc) : prog_(prog), desc_(desc), need_update_(false) {} - BlockDescBind(const BlockDescBind &o) = delete; - BlockDescBind &operator=(const BlockDescBind &o) = delete; - int32_t ID() const { return desc_->idx(); } int32_t Parent() const { return desc_->parent_idx(); } @@ -66,6 +64,8 @@ class BlockDescBind { std::deque> ops_; std::unordered_map> vars_; + + DISABLE_COPY_AND_ASSIGN(BlockDescBind); }; } // namespace framework } // namespace paddle diff --git a/paddle/framework/details/op_registry.h b/paddle/framework/details/op_registry.h index d2516ccc1eec21682276c2fddf049984453404d4..daa474e8c5a223589018720da29a5c3363b5934d 100644 --- a/paddle/framework/details/op_registry.h +++ b/paddle/framework/details/op_registry.h @@ -14,6 +14,7 @@ #pragma once +#include "paddle/framework/grad_op_desc_maker.h" #include "paddle/framework/op_info.h" #include "paddle/framework/op_proto_maker.h" #include "paddle/framework/operator.h" @@ -96,7 +97,10 @@ struct OpInfoFiller { template struct OpInfoFiller { void operator()(const char* op_type, OpInfo* info) const { - info->grad_op_maker_ = new T(); + info->grad_op_maker_ = [](const OpDescBind& fwd_op) { + T maker(fwd_op); + return maker(); + }; } }; } // namespace details diff --git a/paddle/framework/grad_op_builder_test.cc b/paddle/framework/grad_op_builder_test.cc index 55c5fa420e2441260999ba15cf86484f9fd6e890..793780ea44b80a0d750d35cbea601bfbfd5ccda4 100644 --- a/paddle/framework/grad_op_builder_test.cc +++ b/paddle/framework/grad_op_builder_test.cc @@ -39,28 +39,6 @@ class IOIgnoredOpMaker : public OpProtoAndCheckerMaker { namespace f = paddle::framework; -TEST(GradOpBuilder, AddTwo) { - std::shared_ptr add_op(f::OpRegistry::CreateOp( - "sum", {{"X", {"x", "y"}}}, {{"Out", {"out"}}}, {})); - std::shared_ptr grad_add_op = - f::OpRegistry::CreateGradOp(*add_op); - - EXPECT_EQ(grad_add_op->Inputs().size(), 1UL); - EXPECT_EQ(grad_add_op->Outputs().size(), 1UL); - EXPECT_EQ(grad_add_op->Input(f::GradVarName("Out")), f::GradVarName("out")); - auto &outputs = grad_add_op->Outputs(f::GradVarName("X")); - EXPECT_EQ(2UL, outputs.size()); - auto in_output = [&outputs](const std::string &name) { - for (auto &output_name : outputs) { - if (output_name == name) return true; - } - return false; - }; - - EXPECT_TRUE(in_output(f::GradVarName("x"))); - EXPECT_TRUE(in_output(f::GradVarName("y"))); -} - REGISTER_OP(mult_io, f::NOP, f::MutiInOutOpMaker, mult_io_grad, f::NOP); REGISTER_OP(io_ignored, f::NOP, f::IOIgnoredOpMaker, io_ignored_grad, f::NOP); @@ -205,4 +183,4 @@ TEST(GradOpDescBuilder, IOIgnoredInGradient) { {f::GradVarName("in3_1"), f::GradVarName("in3_2")})); delete forw_op; delete grad_op; -} \ No newline at end of file +} diff --git a/paddle/framework/grad_op_desc_maker.h b/paddle/framework/grad_op_desc_maker.h new file mode 100644 index 0000000000000000000000000000000000000000..e9ae6e22060850fe229998d3b651d08a5ca2033a --- /dev/null +++ b/paddle/framework/grad_op_desc_maker.h @@ -0,0 +1,124 @@ +/* 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/op_desc.h" +#include "paddle/framework/operator.h" + +namespace paddle { +namespace framework { + +class GradOpDescMakerBase { + public: + explicit GradOpDescMakerBase(const OpDescBind& fwd_op) : fwd_op_(fwd_op) {} + + virtual ~GradOpDescMakerBase() = default; + virtual std::vector> operator()() const = 0; + + protected: + static std::vector ToGradNames( + const std::vector& var_names) { + std::vector ret_val; + ret_val.reserve(var_names.size()); + std::transform(var_names.begin(), var_names.end(), + std::back_inserter(ret_val), GradVarName); + return ret_val; + } + + std::vector InputGrad(const std::string& name) const { + return ToGradNames(fwd_op_.Input(name)); + } + + std::vector OutputGrad(const std::string& name) const { + return ToGradNames(fwd_op_.Output(name)); + } + + std::vector InputNames() const { + return this->fwd_op_.InputNames(); + } + + std::vector OutputNames() const { + return this->fwd_op_.OutputNames(); + } + + std::vector Input(const std::string& name) const { + return fwd_op_.Input(name); + } + + std::vector Output(const std::string& name) const { + return fwd_op_.Output(name); + } + + const std::unordered_map& Attrs() const { + return fwd_op_.GetAttrMap(); + } + + const Attribute& GetAttr(const std::string& name) const { + auto& map = fwd_op_.GetAttrMap(); + auto it = map.find(name); + PADDLE_ENFORCE(it != map.end(), "Cannot find attribute %s", name); + return it->second; + } + + std::string ForwardOpType() const { return this->fwd_op_.Type(); } + + private: + const OpDescBind& fwd_op_; +}; + +class SingleGradOpDescMaker : public GradOpDescMakerBase { + public: + using GradOpDescMakerBase::GradOpDescMakerBase; + + std::vector> operator()() const { + std::vector> retv; + retv.emplace_back(this->Apply()); + return retv; + } + + protected: + virtual std::unique_ptr Apply() const = 0; +}; + +class DefaultGradOpDescMaker : public SingleGradOpDescMaker { + public: + using SingleGradOpDescMaker::SingleGradOpDescMaker; + + protected: + virtual std::unique_ptr Apply() const { + auto* grad = new OpDescBind(); + grad->SetType(this->GradOpType()); + + for (auto& input_param : this->InputNames()) { + grad->SetInput(input_param, this->Input(input_param)); + grad->SetOutput(GradVarName(input_param), this->InputGrad(input_param)); + } + + for (auto& output_param : this->OutputNames()) { + grad->SetInput(output_param, this->Output(output_param)); + grad->SetInput(GradVarName(output_param), this->OutputGrad(output_param)); + } + + grad->SetAttrMap(this->Attrs()); + + return std::unique_ptr(grad); + } + + virtual std::string GradOpType() const { + return this->ForwardOpType() + "_grad"; + } +}; + +} // namespace framework +} // namespace paddle diff --git a/paddle/framework/lod_tensor.h b/paddle/framework/lod_tensor.h index 49786a4a6635f1b39356dbf9633c4e7da443f04e..4db36ee76609ac6360fe2fc7b4a366e0284d1016 100644 --- a/paddle/framework/lod_tensor.h +++ b/paddle/framework/lod_tensor.h @@ -15,7 +15,7 @@ #pragma once #include -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA #include #include #include @@ -29,7 +29,7 @@ namespace paddle { namespace framework { -#ifdef PADDLE_ONLY_CPU +#ifndef PADDLE_WITH_CUDA template using Vector = std::vector; #else diff --git a/paddle/framework/lod_tensor.md b/paddle/framework/lod_tensor.md index 07bbdf9416c432052b3222757a61ac4bfd70fe14..d147f1c4257eec14664301edab8d1fe2f128d2b0 100644 --- a/paddle/framework/lod_tensor.md +++ b/paddle/framework/lod_tensor.md @@ -1,147 +1,175 @@ # Design Doc: LoD (Level-of-Detail) Tensor -PaddlePaddle's RNN doesn't require that all instances have the same length. To do so, we introduce an extension to Tensor, namely, LoD Tensor. +Like other deep learning systems, PaddlePaddle supports training models from sequence data. Also, like other systems, PaddlePaddle represent a mini-batch of sequences as a Tensor. What is different is that PaddlePaddle doesn't require all sequences in a mini-batch to be of the same length. Thus no need for padding zeros. -## Challenge of Variable-length Inputs +| | TensorFlow | PaddlePaddle | +|-----------------------|------------|--------------| +| RNN | Support | Support | +| recursive RNN | Support | Support | +| padding zeros | Must | No need | +| blob data type | Tensor | LoDTensor | -People usually represent a mini-batch by a Tensor. For example, a mini-batch of 10 images, each of size 32x32, is a 10x32x32 Tensor. So a transformation, T, of all images can be a matrix multiplication of the 10xOx32-dimensional tensor T and the 10x32x32 Tensor. +PaddlePaddle achieves this flexibility by passing through a new data type, *LoD Tensor*, which is a Tensor attached with segmentation index known as *LoD*, between operators. The LoD index doesn't only segment a tensor, but also recursively segments sub-sequences. This document presents the design of LoD and LoDTensor. -Another example is that each mini-batch contains 32 sentences, where each word is a D-dimensional one-hot vector. If all sentences have the same length L, we can represent this mini-batch by a 32xLxD tensor. However, in most cases, sentences have variable lengths, and we will need an index data structure to record these variable lengths. -## LoD as a Solution +## The Challenge: Variable-length Sequences -### Mini-Batch of variable-length sentences +Most deep learning systems represent a mini-batch as a Tensor. For example, a mini-batch of 10 images, each of size 32x32, is a 10x32x32 Tensor. Another example is that each mini-batch contains N sentences, where each word is a D-dimensional one-hot vector. Suppose that all sentences have the same length L, we can represent this mini-batch by a NxLxD tensor. -Let's imagine a mini-batch of 3 variable lengths sentences, containing 3, 1, and 2 words respectively. We can represent it by a (3+1+2)xD tensor plus some index information: +Both examples show that the elements of sequences are usually of the same size. In the first example, all images are 32x32, and in the second one, all words are D-dimensional vectors. It doesn't make sense to allow variable-sized images, as that would require transformations like convolution to handle variable-sized Tensors. + +The real challenge is that in most cases, sentences have variable lengths, and we will need an index data structure to segment the tensor into sequences. Also, sequences might consist of sub-sequences. + + +## A Solution: The LoD Index + +To understand our solution, it is best to look at some examples. + +### A Mini-Batch of Sentences + +Let's imagine a mini-batch of 3 variable lengths sentences composed of 3, 1, and 2 words, respectively. We can represent the mini-batch by a (3+1+2)xD tensor plus some index information: ``` - 3 3 1 2 ||| | || ``` -Each `|` represents a D-dimensional word vectors. The number 3 on top indicate 3 sentences, and numbers 3, 1, and 2 on the second level represent the number of words in each sentence. +where each `|` represents a D-dimensional word vector. The numbers, 3, 1, and 2, form a 1-level LoD. + +### Recursive Sequences + +Let check another example of a 2-level LoD Tensor. Consider a mini-batch of three articles with 3, 1, and 2 sentences, and each sentence consists of a variable number of words: + +``` +3 1 2 +3 2 4 1 2 3 +||| || |||| | || ||| +``` -### Mini-Batch of variable-length videos +### A Mini-Batch of Videos -This approach generalizes to the case where elements are not words, but higher dimensional objects, like images. Suppose that a mini-batch contains videos of the same frame size 640x480. If a mini-batch contains 3 videos of 3, 1, and 2 frames respectively. The underlying tensor is of size (3+1+2)x640x480. The index information illustrates as: +LoD tensors generalize to the case where elements are higher dimensional objects, like images. Suppose that a mini-batch contains videos of the same frame size 640x480. Here is a mini-batch of 3 videos with 3, 1, and 2 frames, respectively. ``` - 3 3 1 2 口口口 口 口口 ``` -where each `口` represents an image. +The underlying tensor is of size (3+1+2)x640x480, and each `口` represents a 640x480 image. -### Mini-Batch of fixed-size images +### A Mini-Batch of Images -Let's get back to a typical example, image classification, where each mini-batch has M fixed-sized images. The LoD Tensor representation is +In traditional cases like a mini-batch with N fixed-sized images, the LoD Tensor representation is as ``` - M 1 1 1 1 1 口口口口 ... 口 ``` -The many 1's on the second level seem duplicated. For this particular case of 2 levels and the second level always have length 1, we can ignore the LoD index. - -### Design and summarization +In this case, we don't lose any information by ignoring the many 1's in the index and simply considering this LoD Tensor as a usual Tensor: -In summary, as long as that the essential elements (words or images) have the same size, we can represent mini-batches by a LoD Tensor: +``` +口口口口 ... 口 +``` -- The underlying tensor has size LxD1xD2x..., where D1xD2... is the size of the essential elements, and -- The first dimension size L has an additonal property -- a LoD index as a nested vector: +### Model Parameters - ```c++ - typedef std::vector> LoD; - ``` +A model parameter is just a usual Tensor, which, just like the above example, is a **0-level LoD Tensor**. -- The LoD index is not necessary when there are only two levels and all elements of the second level have length 1. -## Slicing of LoD Tensor +## The LoD Tensor -Consider that we have a network with three levels of RNN: the top level one handles articles, the second level one handles sentences, and the basic level one handles words. This network requires that mini-batches represented by 3 level LoD Tensor, for example, +Let us revisit above example of the 2-level LoD Tensor ``` - 3 3 1 2 3 2 4 1 2 3 ||| || |||| | || ||| ``` -To allow each level of RNN to handle its input, we define **the slicing of a LoD Tensor is defined as getting the j-th sequence on level i, or the -slice** +It is indeed a tree, where leaves are elementary sequences identified by **branches**. + +For example, the third sentence in above example is identified by branch <0,2>, where 0 indicates the first article with length 3, and 2 indicates the third sentence in this article with length 4. + +### The LoD Index -For example, the <2,1>-slice of above slice is +We can save the LoD index in the above example ``` -2 -|| +3 1 2 +3 2 4 1 2 3 ``` -and the <1,2>-slice of above example is +in a not-full 2D matrix: +```c++ +typedef std::vector > LoD; ``` -2 -2 3 -|| ||| -``` -Let's go on slicing this slice. Its <1,1>-slice is +where + +- `LoD.size()` is the number of levels, or the maximum length of branches, +- `LoD[i][j]` is the length of the j-th segment at the i-th level. + +## The Offset Representation + +To quickly access elementary sequences, we adopt an offset representation -- instead of saving the lengths, we save the beginning and ending elements of sequences. + +In the above example, we accumulate the length of elementary sequences: ``` -1 -1 -| +3 2 4 1 2 3 ``` -### The Slicing Algorithm +into offsets -The algorithm, with over-simplified data structure, is defined as +``` +0 3 5 9 10 12 15 + = = = = = = + 3 2+3 4+5 1+9 2+10 3+12 +``` -```c++ -typedef std::vector> LoD; +so we know that the first sentence is from word 0 to word 3, and the second sentence from work 3 to word 5. -struct LoDTensor { - LoD lod_; - float* tensor_; -}; +Similarly, the lengths in the top level LoD -LoDTensor Slice(const LoDTensor& lodt, int level, int sequence); +``` +3 1 2 ``` -Let us revisit the example above +are transformed into offsets of elements/words as follows: ``` - 3 -3 1 2 -3 2 4 1 2 3 -||| || |||| | || ||| +0 9 10 15 + = = = + 3+2+4 1+9 2+3+10 ``` -Suppose that we want to retrieve the <1,2>-slice +so we can tell that the first article is from word 0 to word 9, and the second article is from word 9 to word 10. + +The complete offset representation is as follows: ``` -2 -2 3 -|| ||| +0 9 10 15 +0 3 5 9 10 12 15 + ||| || |||| | || ||| ``` -we will need to find out the starting position of this slice by summing over all leaf nodes in `LoD` to the left of the slice, i.e., 3 + 2 + 4 + 1 = 10. +## Slicing of LoD Tensors + +When we use the above 2-level LoD Tensor as the input to a nested-RNN, we need to retrieve certain sequences. Here we define the sequence identified by branch as the **-slice**. -To avoid the traversal of the LoD tree at slicing time, we can do it at the construction time -- instead of saving the lengths of the next level in the LoD tree, we can save the starting offset of the next level. For example, above LoD Tensor can be transformed into +For example, the <2>-slice of above example is ``` - 0 -0 9 10 -0 3 5 9 10 12 -||| || |||| | || ||| +10 15 +10 12 15 + || ||| ``` -We don't really need the 0 on top, so the LoD Tensor could be +and the <2,0>-slice of above slice is ``` -0 9 10 -0 3 5 9 10 12 -||| || |||| | || ||| +10 12 + || ``` diff --git a/paddle/framework/op_desc.cc b/paddle/framework/op_desc.cc index 33a064890cc62a4e3f04c79bc000951e76fd477c..852f0f1eb8ea09b336d656afea4f5b146c8d45de 100644 --- a/paddle/framework/op_desc.cc +++ b/paddle/framework/op_desc.cc @@ -31,15 +31,6 @@ const std::vector &OpDescBind::Input( return it->second; } -std::vector OpDescBind::InputNames() const { - std::vector retv; - retv.reserve(this->inputs_.size()); - for (auto &ipt : this->inputs_) { - retv.push_back(ipt.first); - } - return retv; -} - void OpDescBind::SetInput(const std::string ¶m_name, const std::vector &args) { need_update_ = true; @@ -54,15 +45,6 @@ const std::vector &OpDescBind::Output( return it->second; } -std::vector OpDescBind::OutputNames() const { - std::vector retv; - retv.reserve(this->outputs_.size()); - for (auto &ipt : this->outputs_) { - retv.push_back(ipt.first); - } - return retv; -} - void OpDescBind::SetOutput(const std::string ¶m_name, const std::vector &args) { need_update_ = true; diff --git a/paddle/framework/op_desc.h b/paddle/framework/op_desc.h index 0af416971556c812db238caeb7959a0a488c6e1b..397393f796ed11312578859495529f4aedd17e78 100644 --- a/paddle/framework/op_desc.h +++ b/paddle/framework/op_desc.h @@ -35,15 +35,11 @@ class OpDescBind { const std::vector &Input(const std::string &name) const; - std::vector InputNames() const; - void SetInput(const std::string ¶m_name, const std::vector &args); const std::vector &Output(const std::string &name) const; - std::vector OutputNames() const; - void SetOutput(const std::string ¶m_name, const std::vector &args); @@ -61,9 +57,6 @@ class OpDescBind { void SetBlockAttr(const std::string &name, BlockDescBind &block); - // Only be used in C++ - void SetAttrMap(const AttributeMap &attr_map); - Attribute GetAttr(const std::string &name) const; int GetBlockAttr(const std::string &name) const; @@ -71,7 +64,23 @@ class OpDescBind { // Only be used in C++ const AttributeMap &GetAttrMap() const; + // Only be used in C++ + void SetAttrMap(const AttributeMap &attr_map); + + std::vector InputNames() const { return MapKeys(inputs_); } + std::vector OutputNames() const { return MapKeys(outputs_); } + private: + template + static std::vector MapKeys(const MapType &map) { + std::vector ret_val; + ret_val.reserve(map.size()); + std::transform( + map.begin(), map.end(), std::back_inserter(ret_val), + [](const typename MapType::value_type &pair) { return pair.first; }); + return ret_val; + } + void Sync(); OpDesc op_desc_; diff --git a/paddle/framework/op_info.h b/paddle/framework/op_info.h index 470336d367566b342af6a7dac9c9cdf465101ea2..8b7882485ff82860b884d3a8bc3d5543e2426898 100644 --- a/paddle/framework/op_info.h +++ b/paddle/framework/op_info.h @@ -20,20 +20,15 @@ #include "paddle/framework/attribute.h" #include "paddle/framework/op_desc.h" #include "paddle/framework/type_defs.h" +#include "paddle/platform/macros.h" namespace paddle { namespace framework { -class GradOpDescMakerBase { - public: - virtual ~GradOpDescMakerBase() = default; - virtual std::vector operator()(const OpDescBind&) const = 0; -}; - struct OpInfo { OpCreator creator_; std::string grad_op_type_; - GradOpDescMakerBase* grad_op_maker_{nullptr}; + GradOpMakerFN grad_op_maker_; OpProto* proto_{nullptr}; OpAttrChecker* checker_{nullptr}; @@ -67,11 +62,6 @@ class OpInfoMap { public: static OpInfoMap& Instance(); - OpInfoMap(const OpInfoMap& o) = delete; - OpInfoMap(OpInfoMap&& o) = delete; - OpInfoMap& operator=(const OpInfoMap& o) = delete; - OpInfoMap& operator=(OpInfoMap&& o) = delete; - bool Has(const std::string& op_type) const { return map_.find(op_type) != map_.end(); } @@ -107,6 +97,8 @@ class OpInfoMap { private: OpInfoMap() = default; std::unordered_map map_; + + DISABLE_COPY_AND_ASSIGN(OpInfoMap); }; } // namespace framework diff --git a/paddle/framework/op_proto_maker_test.cc b/paddle/framework/op_proto_maker_test.cc index b01e30f75371ca4aa63dae86ddfb966b1d4c7830..988a14cf4de8fdf052ca7e8c41bff0c05ba2daaa 100644 --- a/paddle/framework/op_proto_maker_test.cc +++ b/paddle/framework/op_proto_maker_test.cc @@ -48,4 +48,4 @@ TEST(ProtoMaker, DuplicatedInOut) { paddle::framework::OpAttrChecker op_checker; auto proto_maker = TestInOutProtoMaker(&op_proto, &op_checker); ASSERT_THROW(proto_maker.Validate(), paddle::platform::EnforceNotMet); -} \ No newline at end of file +} diff --git a/paddle/framework/op_registry.h b/paddle/framework/op_registry.h index 4ee2c7d27561c3855059c42e1604f353c65bfa41..958cf581f53b2bb0252b655267f657c8d9ab371c 100644 --- a/paddle/framework/op_registry.h +++ b/paddle/framework/op_registry.h @@ -211,7 +211,7 @@ class OpKernelRegistrar : public Registrar { // TODO(fengjiayi): The following macros // seems ugly, do we have better method? -#ifdef PADDLE_ONLY_CPU +#ifndef PADDLE_WITH_CUDA #define USE_OP_KERNEL(op_type) USE_OP_DEVICE_KERNEL(op_type, CPU) #else #define USE_OP_KERNEL(op_type) \ diff --git a/paddle/framework/op_registry_test.cc b/paddle/framework/op_registry_test.cc index f89f40b444f893d898595db50e245824a19284f6..b860fe6cac773d1e85adecc43f5dfec42b6c7661 100644 --- a/paddle/framework/op_registry_test.cc +++ b/paddle/framework/op_registry_test.cc @@ -183,4 +183,4 @@ class CosineOpComplete : public paddle::framework::CosineOp { TEST(OperatorRegistrar, Test) { using namespace paddle::framework; OperatorRegistrar reg("cos"); -} \ No newline at end of file +} diff --git a/paddle/framework/operator.cc b/paddle/framework/operator.cc index 1012a30b0a6b28e79b92007244bb6e0139cf39f9..2ca838f838ad0b9211a59bf9247c48d283484d50 100644 --- a/paddle/framework/operator.cc +++ b/paddle/framework/operator.cc @@ -25,7 +25,7 @@ Eigen::DefaultDevice& ExecutionContext::GetEigenDevice< return *device_context_.GetEigenDevice(); } -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA template <> Eigen::GpuDevice& ExecutionContext::GetEigenDevice() const { diff --git a/paddle/framework/program_desc.h b/paddle/framework/program_desc.h index 06ffcd4b15078f62ea8b7a3714e73de799530785..9b34a06aeff94e6fa855f6f287a73889e2a4faee 100644 --- a/paddle/framework/program_desc.h +++ b/paddle/framework/program_desc.h @@ -16,6 +16,7 @@ limitations under the License. */ #include #include "paddle/framework/framework.pb.h" +#include "paddle/platform/macros.h" namespace paddle { namespace framework { @@ -26,9 +27,6 @@ class ProgramDescBind { public: static ProgramDescBind &Instance(ProgramDesc *prog); - ProgramDescBind(const ProgramDescBind &o) = delete; - ProgramDescBind &operator=(const ProgramDescBind &o) = delete; - BlockDescBind *AppendBlock(const BlockDescBind &parent); BlockDescBind *Block(size_t idx) { return blocks_[idx].get(); } @@ -46,6 +44,8 @@ class ProgramDescBind { ProgramDesc *prog_; std::vector> blocks_; + + DISABLE_COPY_AND_ASSIGN(ProgramDescBind); }; } // namespace framework } // namespace paddle diff --git a/paddle/framework/scope.h b/paddle/framework/scope.h index c93b03e48130afe9568089b6a7586c4185d1d5b4..7047f0d55e9844aec19892631fe4b5b387bf89ca 100644 --- a/paddle/framework/scope.h +++ b/paddle/framework/scope.h @@ -19,6 +19,7 @@ limitations under the License. */ #include #include "paddle/framework/variable.h" +#include "paddle/platform/macros.h" namespace paddle { namespace framework { @@ -38,11 +39,6 @@ class Scope { Scope() {} ~Scope(); - // Disable Copy, Assign, Move. - Scope(const Scope& other) = delete; - Scope& operator=(const Scope& other) = delete; - Scope(Scope&& other) = delete; - /// Create a sub-scope. Returns a reference other than a pointer so /// to prevent from manual deletion. /// Mark it to const because that new kid scope cannot change parent scope. @@ -73,6 +69,8 @@ class Scope { std::unordered_map vars_; mutable std::list kids_; Scope const* parent_{nullptr}; + + DISABLE_COPY_AND_ASSIGN(Scope); }; } // namespace framework diff --git a/paddle/framework/tensor_array.h b/paddle/framework/tensor_array.h index e76f33d2c0818aa6091d6d5edb1cdfd1e9122476..22ae6a966f90c47fe8b4bbaaf5eb227c39d84173 100644 --- a/paddle/framework/tensor_array.h +++ b/paddle/framework/tensor_array.h @@ -47,13 +47,6 @@ class TensorArray { // max number of values allowed to store. const size_t MAX_SIZE{100000}; - /* - * Inputs: - * - value_shared: share memory between tensors. - */ - explicit TensorArray(bool values_shared = true) - : values_shared_(values_shared) {} - /* * Read the value at location `index` in the `TensorArray`. */ @@ -111,7 +104,6 @@ class TensorArray { private: mutable std::vector values_; - bool values_shared_; }; // class TensorArray } // namespace framework diff --git a/paddle/framework/tensor_impl.h b/paddle/framework/tensor_impl.h index a5405f9c31543b5733f9db923c2a6f8b968cfc2d..379eac94f985c9fa10b6c773065551575f57f033 100644 --- a/paddle/framework/tensor_impl.h +++ b/paddle/framework/tensor_impl.h @@ -65,7 +65,7 @@ inline T* Tensor::mutable_data(platform::Place place) { holder_.reset(new PlaceholderImpl( boost::get(place), size)); } else if (platform::is_gpu_place(place)) { -#ifdef PADDLE_ONLY_CPU +#ifndef PADDLE_WITH_CUDA PADDLE_THROW("'GPUPlace' is not supported in CPU only device."); } #else @@ -103,7 +103,7 @@ inline void Tensor::CopyFrom(const Tensor& src, memory::Copy(boost::get(dst_place), dst_ptr, boost::get(src_place), src_ptr, size); } -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA else if (platform::is_gpu_place(src_place) && platform::is_cpu_place(dst_place)) { memory::Copy(boost::get(dst_place), dst_ptr, diff --git a/paddle/framework/tensor_test.cc b/paddle/framework/tensor_test.cc index e2ec738de35c90c6a06c9a46b062d4cce55f5eda..58cf0fc3cb6cf0bad693118ca57d71fb21c55a40 100644 --- a/paddle/framework/tensor_test.cc +++ b/paddle/framework/tensor_test.cc @@ -74,7 +74,7 @@ TEST(Tensor, MutableData) { EXPECT_EQ(p1, p2); } -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA { Tensor src_tensor; float* p1 = nullptr; @@ -126,7 +126,7 @@ TEST(Tensor, ShareDataWith) { ASSERT_EQ(src_tensor.data(), dst_tensor.data()); } -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA { Tensor src_tensor; Tensor dst_tensor; @@ -163,7 +163,7 @@ TEST(Tensor, Slice) { EXPECT_EQ(src_data_address + 3 * 4 * 1 * sizeof(int), slice_data_address); } -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA { Tensor src_tensor; src_tensor.mutable_data(make_ddim({6, 9}), GPUPlace()); @@ -218,7 +218,7 @@ TEST(Tensor, CopyFrom) { EXPECT_EQ(dst_ptr[i], slice_ptr[i]); } } -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA { Tensor src_tensor; Tensor gpu_tensor; diff --git a/paddle/framework/type_defs.h b/paddle/framework/type_defs.h index dec5066f1e649a27c9f47fc5f93faa1498b37de7..a5b94722136eca7ac0a22eb7ef113532330428ab 100644 --- a/paddle/framework/type_defs.h +++ b/paddle/framework/type_defs.h @@ -20,6 +20,7 @@ namespace paddle { namespace framework { class OperatorBase; +class OpDescBind; using VariableNameMap = std::map>; // The order should be as same as framework.proto @@ -34,5 +35,8 @@ using OpCreator = std::function; +using GradOpMakerFN = + std::function>(const OpDescBind&)>; + } // namespace framework } // namespace paddle diff --git a/paddle/function/BlockExpandOp.cpp b/paddle/function/BlockExpandOp.cpp index a89b6bba45843d81264819cad6ba053f28314f6b..bd0fe119ce46df9c333258c9c1ad7b5b2bdc544f 100644 --- a/paddle/function/BlockExpandOp.cpp +++ b/paddle/function/BlockExpandOp.cpp @@ -194,7 +194,7 @@ public: REGISTER_TYPED_FUNC(BlockExpand, CPU, BlockExpandForward); REGISTER_TYPED_FUNC(BlockExpandGrad, CPU, BlockExpandBackward); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA REGISTER_TYPED_FUNC(BlockExpand, GPU, BlockExpandForward); REGISTER_TYPED_FUNC(BlockExpandGrad, GPU, BlockExpandBackward); #endif diff --git a/paddle/function/ContextProjectionOp.cpp b/paddle/function/ContextProjectionOp.cpp index b87750b74247bd0eb822340bc5a85d41b86ecec2..23916c0f4b6319004ca0793bc9305b8a1dd0ae89 100644 --- a/paddle/function/ContextProjectionOp.cpp +++ b/paddle/function/ContextProjectionOp.cpp @@ -395,7 +395,7 @@ REGISTER_TYPED_FUNC(ContextProjectionForward, REGISTER_TYPED_FUNC(ContextProjectionBackward, CPU, ContextProjectionBackwardFunc); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA REGISTER_TYPED_FUNC(ContextProjectionForward, GPU, ContextProjectionForwardFunc); diff --git a/paddle/function/CosSimOp.cpp b/paddle/function/CosSimOp.cpp index 7ece7b2dfedaf460741c97b5a700eb632d85cabc..2e5c281f37d8ffb1062121b5dc5b4f790ab52089 100644 --- a/paddle/function/CosSimOp.cpp +++ b/paddle/function/CosSimOp.cpp @@ -233,7 +233,7 @@ private: REGISTER_TYPED_FUNC(CosSimForward, CPU, CosSimForwardFunc); REGISTER_TYPED_FUNC(CosSimBackward, CPU, CosSimBackwardFunc); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA REGISTER_TYPED_FUNC(CosSimForward, GPU, CosSimForwardFunc); REGISTER_TYPED_FUNC(CosSimBackward, GPU, CosSimBackwardFunc); #endif diff --git a/paddle/function/CropOp.cpp b/paddle/function/CropOp.cpp index f12ee43e3d72f9ac776eaff93914228850694dd2..46f98f12c1f150fdf3ed53a7a96e5cf0020e14a4 100644 --- a/paddle/function/CropOp.cpp +++ b/paddle/function/CropOp.cpp @@ -169,7 +169,7 @@ private: REGISTER_TYPED_FUNC(Crop, CPU, CropFunc); REGISTER_TYPED_FUNC(CropGrad, CPU, CropGradFunc); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA REGISTER_TYPED_FUNC(Crop, GPU, CropFunc); REGISTER_TYPED_FUNC(CropGrad, GPU, CropGradFunc); #endif diff --git a/paddle/function/CrossMapNormalOp.cpp b/paddle/function/CrossMapNormalOp.cpp index ef878bfbba961bdd3d5212e19fb83bb1e285e47f..9e88669d37bd50179dcc0464e8c1cd6e2fed74db 100644 --- a/paddle/function/CrossMapNormalOp.cpp +++ b/paddle/function/CrossMapNormalOp.cpp @@ -336,7 +336,7 @@ private: REGISTER_TYPED_FUNC(CrossMapNormal, CPU, CrossMapNormalFunc); REGISTER_TYPED_FUNC(CrossMapNormalGrad, CPU, CrossMapNormalGradFunc); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA REGISTER_TYPED_FUNC(CrossMapNormal, GPU, CrossMapNormalFunc); REGISTER_TYPED_FUNC(CrossMapNormalGrad, GPU, CrossMapNormalGradFunc); #endif diff --git a/paddle/function/DepthwiseConvOp.cpp b/paddle/function/DepthwiseConvOp.cpp index 2f3112fe657cd381891dc53c7179e7520911e8c9..9863e3ae1d5fcb1eece5267fd4f2a6b593b799df 100644 --- a/paddle/function/DepthwiseConvOp.cpp +++ b/paddle/function/DepthwiseConvOp.cpp @@ -292,7 +292,7 @@ REGISTER_TYPED_FUNC(DepthwiseConvGradInput, REGISTER_TYPED_FUNC(DepthwiseConvGradFilter, CPU, DepthwiseConvGradFilterFunction); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA REGISTER_TYPED_FUNC(DepthwiseConv, GPU, DepthwiseConvFunction); REGISTER_TYPED_FUNC(DepthwiseConvGradInput, GPU, diff --git a/paddle/function/DepthwiseConvOpTest.cpp b/paddle/function/DepthwiseConvOpTest.cpp index d8e8c889d5c23bf9b2b5fd0b0393395883188fd8..b1a90da7db2b647dd384e3772820294140e5ec9d 100644 --- a/paddle/function/DepthwiseConvOpTest.cpp +++ b/paddle/function/DepthwiseConvOpTest.cpp @@ -17,7 +17,7 @@ limitations under the License. */ namespace paddle { -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA TEST(DepthwiseConv, Forward) { DepthwiseConvolution( "GemmConv-CPU", "DepthwiseConv-GPU", forward); diff --git a/paddle/function/GemmConvOp.cpp b/paddle/function/GemmConvOp.cpp index f8cf4ebea8d724f0291b981647622b63e3d84495..bdb56ddac38b91d756fc6f31282f29c0489fd660 100644 --- a/paddle/function/GemmConvOp.cpp +++ b/paddle/function/GemmConvOp.cpp @@ -340,7 +340,7 @@ public: REGISTER_TYPED_FUNC(GemmConv, CPU, GemmConvFunction); REGISTER_TYPED_FUNC(GemmConvGradInput, CPU, GemmConvGradInputFunction); REGISTER_TYPED_FUNC(GemmConvGradFilter, CPU, GemmConvGradFilterFunction); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA REGISTER_TYPED_FUNC(GemmConv, GPU, GemmConvFunction); REGISTER_TYPED_FUNC(GemmConvGradInput, GPU, GemmConvGradInputFunction); REGISTER_TYPED_FUNC(GemmConvGradFilter, GPU, GemmConvGradFilterFunction); diff --git a/paddle/function/GemmConvOpTest.cpp b/paddle/function/GemmConvOpTest.cpp index 5283d79a5a53d979ae4e134f7e46b7ee106e9c44..b5b5e1f35b79e422b14f7495bc321533cc1d618a 100644 --- a/paddle/function/GemmConvOpTest.cpp +++ b/paddle/function/GemmConvOpTest.cpp @@ -24,7 +24,7 @@ TEST(GemmConv, NaiveConv) { "NaiveConv-CPU", "GemmConv-CPU", forward); } -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA TEST(GemmConv, Forward) { Convolution( "GemmConv-CPU", "GemmConv-GPU", forward); diff --git a/paddle/function/Im2ColTest.cpp b/paddle/function/Im2ColTest.cpp index acc88a553abe7ac58b629aba9b850df58cee7f81..a0a01a5fc7fc055dce6ddb3ee51c7ab18f8a4ca7 100644 --- a/paddle/function/Im2ColTest.cpp +++ b/paddle/function/Im2ColTest.cpp @@ -116,7 +116,7 @@ void TestIm2ColFunctor() { TEST(Im2ColFunctor, CPU) { TestIm2ColFunctor(); } -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA TEST(Im2ColFunctor, GPU) { TestIm2ColFunctor(); } diff --git a/paddle/function/MulOp.cpp b/paddle/function/MulOp.cpp index 25e41edad54bec0f76a3de4799fab14241407272..704a8c41325ef86067a3bd8ed6d772b77df147c5 100644 --- a/paddle/function/MulOp.cpp +++ b/paddle/function/MulOp.cpp @@ -341,7 +341,7 @@ private: }; REGISTER_TYPED_FUNC(MulOp, CPU, MulFunc); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA REGISTER_TYPED_FUNC(MulOp, GPU, MulFunc); #endif } // namespace paddle diff --git a/paddle/function/PadOp.cpp b/paddle/function/PadOp.cpp index adba7c92ece505eecc74edce6b393cf27fa10ccc..eed2f2e3089b6b6167ef7c5a7acb7ecaa08945e1 100644 --- a/paddle/function/PadOp.cpp +++ b/paddle/function/PadOp.cpp @@ -207,7 +207,7 @@ private: REGISTER_TYPED_FUNC(Pad, CPU, PadFunc); REGISTER_TYPED_FUNC(PadGrad, CPU, PadGradFunc); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA REGISTER_TYPED_FUNC(Pad, GPU, PadFunc); REGISTER_TYPED_FUNC(PadGrad, GPU, PadGradFunc); #endif diff --git a/paddle/function/RowConvOp.cpp b/paddle/function/RowConvOp.cpp index b6501e8f4db7fd33891cd80e07a6f36dd0b34532..7c802d66273c6f7aa56b2f460e3dff4401967517 100644 --- a/paddle/function/RowConvOp.cpp +++ b/paddle/function/RowConvOp.cpp @@ -217,7 +217,7 @@ public: REGISTER_TYPED_FUNC(RowConv, CPU, RowConvFunc); REGISTER_TYPED_FUNC(RowConvGrad, CPU, RowConvGradFunc); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA REGISTER_TYPED_FUNC(RowConv, GPU, RowConvFunc); REGISTER_TYPED_FUNC(RowConvGrad, GPU, RowConvGradFunc); #endif diff --git a/paddle/function/SwitchOp.cpp b/paddle/function/SwitchOp.cpp index 01e252a8dc0cd5fa1e964efa01d04cf282b3dfe7..597723a2dded6a6a116e05b7d4c942cd633e2c99 100644 --- a/paddle/function/SwitchOp.cpp +++ b/paddle/function/SwitchOp.cpp @@ -132,7 +132,7 @@ public: REGISTER_TYPED_FUNC(NCHW2NHWC, CPU, NCHW2NHWCFunc); REGISTER_TYPED_FUNC(NHWC2NCHW, CPU, NHWC2NCHWFunc); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA REGISTER_TYPED_FUNC(NCHW2NHWC, GPU, NCHW2NHWCFunc); REGISTER_TYPED_FUNC(NHWC2NCHW, GPU, NHWC2NCHWFunc); #endif diff --git a/paddle/gserver/layers/BatchNormBaseLayer.cpp b/paddle/gserver/layers/BatchNormBaseLayer.cpp index f7a80e23e1bd49549bec57b360587adc6b423794..bc7d1c83a48aefeb4bc6d3baa32b78aba712e58d 100644 --- a/paddle/gserver/layers/BatchNormBaseLayer.cpp +++ b/paddle/gserver/layers/BatchNormBaseLayer.cpp @@ -16,7 +16,7 @@ limitations under the License. */ #include "BatchNormalizationLayer.h" #include "Layer.h" #include "paddle/utils/Stat.h" -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA #include "CudnnBatchNormLayer.h" #endif diff --git a/paddle/gserver/layers/BatchNormalizationLayer.cpp b/paddle/gserver/layers/BatchNormalizationLayer.cpp index 412762d38475422be98ffeb87ffcfb028c3e035f..dacff25e5927daf9c991577a71be86b160228317 100644 --- a/paddle/gserver/layers/BatchNormalizationLayer.cpp +++ b/paddle/gserver/layers/BatchNormalizationLayer.cpp @@ -13,7 +13,7 @@ See the License for the specific language governing permissions and limitations under the License. */ #include "paddle/utils/Stat.h" -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA #include "hl_batch_transpose.h" #endif #include "BatchNormalizationLayer.h" @@ -90,7 +90,7 @@ void BatchNormalizationLayer::expandMat(const MatrixPtr& in, MatrixPtr& out) { size_t batchSize = in->getHeight(); CHECK_EQ(out->getHeight(), batchSize * imgPixels_); if (useGpu_) { -#ifdef PADDLE_ONLY_CPU +#ifndef PADDLE_WITH_CUDA LOG(FATAL) << "paddle is compiled only for cpu"; #else batchTranspose( @@ -127,7 +127,7 @@ void BatchNormalizationLayer::shrinkMat(const MatrixPtr& in, MatrixPtr& out) { } CHECK_EQ(in->getHeight(), static_cast(batchSize * imgPixels_)); if (useGpu_) { -#ifdef PADDLE_ONLY_CPU +#ifndef PADDLE_WITH_CUDA LOG(FATAL) << "paddle is compiled only for cpu"; #else batchTranspose( diff --git a/paddle/gserver/layers/PoolLayer.cpp b/paddle/gserver/layers/PoolLayer.cpp index 96d5c54accc047b685502a178de2d290f3158731..7b932d5a76e9c4fe7cbe5882bbc19eb3de4b503a 100644 --- a/paddle/gserver/layers/PoolLayer.cpp +++ b/paddle/gserver/layers/PoolLayer.cpp @@ -15,7 +15,7 @@ limitations under the License. */ #include "PoolLayer.h" #include "PoolProjectionLayer.h" #include "paddle/utils/Logging.h" -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA #include "CudnnPoolLayer.h" #endif namespace paddle { @@ -53,7 +53,7 @@ Layer* PoolLayer::create(const LayerConfig& config) { const std::string& pool = config.inputs(0).pool_conf().pool_type(); if (pool == "max-projection" || pool == "avg-projection") { return new PoolProjectionLayer(config); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA } else if (CudnnPoolLayer::typeCheck(pool)) { return new CudnnPoolLayer(config); #endif diff --git a/paddle/gserver/tests/LayerGradUtil.cpp b/paddle/gserver/tests/LayerGradUtil.cpp index a38880e14cdfcef05461dae567d198e5400c6bb1..cd957c7c0bca4c6089cc07e8f4226b8260190f07 100644 --- a/paddle/gserver/tests/LayerGradUtil.cpp +++ b/paddle/gserver/tests/LayerGradUtil.cpp @@ -674,7 +674,7 @@ void testLayerGradKernel(TestConfig testConf, bool useGpu, bool useWeight, float epsilon) { -#ifdef PADDLE_ONLY_CPU +#ifndef PADDLE_WITH_CUDA if (useGpu) return; #endif FLAGS_use_gpu = useGpu; diff --git a/paddle/gserver/tests/test_BatchNorm.cpp b/paddle/gserver/tests/test_BatchNorm.cpp index 659eefa31bdb1f2433d03a59d5bf4782c71bdecf..050fde9d0af54e6d3fd46e19903b926019ddd229 100644 --- a/paddle/gserver/tests/test_BatchNorm.cpp +++ b/paddle/gserver/tests/test_BatchNorm.cpp @@ -119,7 +119,7 @@ TEST(Layer, batchNorm) { CHECK_EQ(static_cast(convLayer->getOutputValue()->getWidth()), 576); } -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA void batchNormInference(int n, int c, int h, int w) { MatrixPtr input = std::make_shared(n, c * h * w); MatrixPtr cudnnOut = std::make_shared(n, c * h * w); diff --git a/paddle/gserver/tests/test_ConvUnify.cpp b/paddle/gserver/tests/test_ConvUnify.cpp index e7325e0cc3b7195b5fec77c878e3e087cfc643e0..ffcc47e2a84356807b6591ba11e670d6d3f336ee 100644 --- a/paddle/gserver/tests/test_ConvUnify.cpp +++ b/paddle/gserver/tests/test_ConvUnify.cpp @@ -117,7 +117,7 @@ MatrixPtr doOneConvTest(size_t imgSize, } TEST(Layer, convParaUnified) { -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA MatrixPtr input, resultCpu, resultGpu; /// TEST1 for conv /// diff --git a/paddle/gserver/tests/test_DetectionOutput.cpp b/paddle/gserver/tests/test_DetectionOutput.cpp index af43dc51fad35c834635b543b1a016f6d717de1e..dc39c97a87f8b346dc9cc09d6158b1b4069bcf2d 100644 --- a/paddle/gserver/tests/test_DetectionOutput.cpp +++ b/paddle/gserver/tests/test_DetectionOutput.cpp @@ -150,7 +150,7 @@ TEST(Layer, detectionOutputLayerFwd) { useGpu, result2); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA // GPU case 1. useGpu = true; inputLoc = Matrix::create(1, 16, false, useGpu); diff --git a/paddle/gserver/tests/test_Evaluator.cpp b/paddle/gserver/tests/test_Evaluator.cpp index 93996392d221d531f65caf465decbffdbc2d0384..62a131171fa5ae973cb3069151a582aaeac9ee0e 100644 --- a/paddle/gserver/tests/test_Evaluator.cpp +++ b/paddle/gserver/tests/test_Evaluator.cpp @@ -51,7 +51,7 @@ void testEvaluator(TestConfig testConf, string testEvaluatorName, size_t batchSize, bool useGpu) { -#ifdef PADDLE_ONLY_CPU +#ifndef PADDLE_WITH_CUDA if (useGpu) return; #endif FLAGS_use_gpu = useGpu; diff --git a/paddle/gserver/tests/test_KmaxSeqScore.cpp b/paddle/gserver/tests/test_KmaxSeqScore.cpp index 308abe6816428bc0f98ec32e892622fa4a23b1ae..6386259882f8c70da59e22574c60fadda636aa42 100644 --- a/paddle/gserver/tests/test_KmaxSeqScore.cpp +++ b/paddle/gserver/tests/test_KmaxSeqScore.cpp @@ -97,7 +97,7 @@ TEST(Layer, kmaxSeqScoreLayer) { Matrix::create(subSeqStartPosition.back(), 1, false, false); std::vector mode = {false}; -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA mode.push_back(true); #endif diff --git a/paddle/gserver/tests/test_LayerGrad.cpp b/paddle/gserver/tests/test_LayerGrad.cpp index 090bde7b203652e3ffb1662b8f5b8937885d2608..90a3352898863a66819e23dd9ecba375ae104a60 100644 --- a/paddle/gserver/tests/test_LayerGrad.cpp +++ b/paddle/gserver/tests/test_LayerGrad.cpp @@ -12,7 +12,7 @@ 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. */ -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA #include #endif #include @@ -258,7 +258,7 @@ void testProjectionConv(size_t groups, bool isDeconv) { true); } -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA TEST(Projection, conv) { /// test ConvProjection testProjectionConv(1, false); @@ -422,7 +422,7 @@ TEST(Layer, depthwiseConvLayer) { // 'depthwise_conv' is a sepecial case of 'exconv' whose // groups size equals to the input channels size. testDepthwiseConvLayer("exconv", /* useGpu= */ false); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA testDepthwiseConvLayer("exconv", /* useGpu= */ true); #endif } @@ -480,7 +480,7 @@ void testConvLayer(const string& type, bool trans, bool useGpu) { TEST(Layer, convLayer) { testConvLayer("exconv", /* trans= */ false, /* useGpu= */ false); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA testConvLayer("exconv", /* trans= */ false, /* useGpu= */ true); testConvLayer("cudnn_conv", /* trans= */ false, /* useGpu= */ true); #endif @@ -525,7 +525,7 @@ TEST(Layer, convTransLayer) { for (auto useGpu : {false, true}) { testConvTransLayer("exconvt", /* trans= */ false, /* useGpu= */ useGpu); } -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA testConvTransLayer("cudnn_convt", /* trans= */ false, /* useGpu= */ true); #endif } @@ -638,7 +638,7 @@ TEST(Layer, SelectiveFullyConnectedLayer) { /* trans= */ false, /* useGup= */ false, false); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA testLayerGrad(config, "selective_fc", 100, @@ -1210,7 +1210,7 @@ void testPoolLayer(const string& poolType, bool trans, bool useGpu) { testLayerGrad(config, "pool", 100, trans, useGpu); } -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA void testPoolLayer2(const string& poolType, bool trans, bool useGpu) { TestConfig config; config.inputDefs.push_back({INPUT_DATA, "layer_0", 3200, 0}); @@ -1236,7 +1236,7 @@ TEST(Layer, PoolLayer) { testPoolLayer("avg-projection", /* trans= */ false, /* useGpu= */ false); testPoolLayer("max-projection", /* trans= */ false, /* useGpu= */ false); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA testPoolLayer("avg-projection", /* trans= */ false, /* useGpu= */ true); testPoolLayer("max-projection", /* trans= */ false, /* useGpu= */ true); testPoolLayer("cudnn-max-pool", /* trans= */ false, /* useGpu= */ true); @@ -1309,7 +1309,7 @@ void testPool3DLayer(const string& poolType, bool trans, bool useGpu) { TEST(Layer, Pool3DLayer) { testPool3DLayer("avg", /* trans= */ false, /* useGpu= */ false); testPool3DLayer("max", /* trans= */ false, /* useGpu= */ false); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA testPool3DLayer("avg", /* trans= */ false, /* useGpu= */ true); testPool3DLayer("max", /* trans= */ false, /* useGpu= */ true); #endif @@ -1695,7 +1695,7 @@ void testBatchNormLayer(const string& type, bool trans, bool useGpu) { TEST(Layer, BatchNormalizationLayer) { testBatchNormLayer("batch_norm", false, false); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA testBatchNormLayer("batch_norm", false, true); if (hl_get_cudnn_lib_version() >= int(4000)) { testBatchNormLayer("cudnn_batch_norm", false, true); @@ -1744,7 +1744,7 @@ void testBatchNorm3DLayer(const string& type, bool trans, bool useGpu) { TEST(Layer, testBatchNorm3DLayer) { testBatchNorm3DLayer("batch_norm", false, false); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA testBatchNorm3DLayer("batch_norm", false, true); if (hl_get_cudnn_lib_version() >= int(4000)) { testBatchNorm3DLayer("cudnn_batch_norm", false, true); @@ -2262,7 +2262,7 @@ void test3DConvLayer(const string& type, bool trans, bool useGpu) { TEST(Layer, test3DConvLayer) { test3DConvLayer("conv3d", /* trans= */ false, /* useGpu= */ false); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA test3DConvLayer("conv3d", /* trans= */ false, /* useGpu= */ true); #endif } @@ -2339,7 +2339,7 @@ void test3DDeConvLayer(const string& type, bool trans, bool useGpu) { TEST(Layer, test3DDeConvLayer) { test3DDeConvLayer("deconv3d", /* trans= */ false, /* useGpu= */ false); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA test3DDeConvLayer("deconv3d", /* trans= */ false, /* useGpu= */ true); #endif } diff --git a/paddle/gserver/tests/test_NetworkCompare.cpp b/paddle/gserver/tests/test_NetworkCompare.cpp index d36f72360f8ebd2033fb3e8c0e1b30911abba362..2b92211936aad1a034369bda0830bed3438cf401 100644 --- a/paddle/gserver/tests/test_NetworkCompare.cpp +++ b/paddle/gserver/tests/test_NetworkCompare.cpp @@ -243,7 +243,7 @@ TEST(Compare, concat_slice) { compareNetwork(config_file_a, config_file_b); } -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA TEST(Compare, img_pool) { std::string config_file_a = "./gserver/tests/img_pool_a.conf"; std::string config_file_b = "./gserver/tests/img_pool_b.conf"; diff --git a/paddle/gserver/tests/test_PriorBox.cpp b/paddle/gserver/tests/test_PriorBox.cpp index ae0e3bc3d24c54eb84c7b5f5053e629607ef4310..8dc5568784295b5a2e7d4decd178d612432a1a18 100644 --- a/paddle/gserver/tests/test_PriorBox.cpp +++ b/paddle/gserver/tests/test_PriorBox.cpp @@ -151,7 +151,7 @@ TEST(Layer, priorBoxLayerFwd) { useGpu, result); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA // reset the input parameters variance[1] = 0.1; variance[3] = 0.2; diff --git a/paddle/gserver/tests/test_ProtoDataProvider.cpp b/paddle/gserver/tests/test_ProtoDataProvider.cpp index e11bf402c27898b8fdbd3fceeb8aeff8906352db..af6472619d1840e82787974d265d601b4a406c09 100644 --- a/paddle/gserver/tests/test_ProtoDataProvider.cpp +++ b/paddle/gserver/tests/test_ProtoDataProvider.cpp @@ -485,7 +485,7 @@ TEST(ProtoDataProvider, test) { // Currently in async mode, useGpu is not supported continue; } -#ifdef PADDLE_ONLY_CPU +#ifndef PADDLE_WITH_CUDA if (useGpu) { continue; } @@ -525,7 +525,7 @@ TEST(ProtoDataProvider, constant_slots) { for (int numConstantSlots : {1, 2}) { for (int useGpu : numTwoArray) { for (int dataCompression : numTwoArray) { -#ifdef PADDLE_ONLY_CPU +#ifndef PADDLE_WITH_CUDA if (useGpu) { continue; } @@ -708,7 +708,7 @@ TEST(ProtoSequenceDataProvider, test) { // Currently in async mode, useGpu is not supported continue; } -#ifdef PADDLE_ONLY_CPU +#ifndef PADDLE_WITH_CUDA if (useGpu) { continue; } diff --git a/paddle/gserver/tests/test_PyDataProvider.cpp b/paddle/gserver/tests/test_PyDataProvider.cpp index db883543c306c1938eb9da188ce20ed768018efb..fe54799259d86064c4fcaec0e53707247981a1b4 100644 --- a/paddle/gserver/tests/test_PyDataProvider.cpp +++ b/paddle/gserver/tests/test_PyDataProvider.cpp @@ -37,7 +37,7 @@ TEST(PyDataProvider, py_fill_slots) { config.clear_files(); std::string dataFile = "gserver/tests/pyDataProvider/pyDataProviderList"; config.set_files(dataFile); -#ifdef PADDLE_ONLY_CPU +#ifndef PADDLE_WITH_CUDA bool useGpu = false; #else bool useGpu = true; @@ -71,7 +71,7 @@ TEST(PyDataProvider, py_fill_nest_slots) { std::string dataFile = "gserver/tests/pyDataProvider/pyDataProviderList"; config.set_files(dataFile); EXPECT_EQ(config.IsInitialized(), true); -#ifdef PADDLE_ONLY_CPU +#ifndef PADDLE_WITH_CUDA bool useGpu = false; #else bool useGpu = true; diff --git a/paddle/gserver/tests/test_SelectiveFCLayer.cpp b/paddle/gserver/tests/test_SelectiveFCLayer.cpp index ab23d00a2cb6077147f5b89664a8e2437b4cd63b..4c87fe1bba1eff3c081754b3b255a0d7d1b3dfbe 100644 --- a/paddle/gserver/tests/test_SelectiveFCLayer.cpp +++ b/paddle/gserver/tests/test_SelectiveFCLayer.cpp @@ -321,7 +321,7 @@ TEST(Layer, SelectiveFcLayer_train_dense_mul) { "filelist=gserver/tests/SelectiveFcTest/dense_mul_list"; for (auto useGpu : {false, true}) { -#ifdef PADDLE_ONLY_CPU +#ifndef PADDLE_WITH_CUDA if (useGpu) { break; } @@ -388,7 +388,7 @@ void testSelectiveFcLayerTrainSparseMul(const LayerConfig& config, outMatSelfc->getWidth(), outMatSelfc->getElementCnt())); cpuOutMatSelfc->copyFrom(*outMatSelfc, HPPL_STREAM_DEFAULT); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA if (useGpu) { hl_stream_synchronize(HPPL_STREAM_DEFAULT); } @@ -418,7 +418,7 @@ void testSelectiveFcLayerTrainSparseMul(const LayerConfig& config, MatrixPtr cpuOutMatFc( new CpuMatrix(outMatFc->getHeight(), outMatFc->getWidth())); cpuOutMatFc->copyFrom(*outMatFc, HPPL_STREAM_DEFAULT); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA if (useGpu) { hl_stream_synchronize(HPPL_STREAM_DEFAULT); } @@ -443,7 +443,7 @@ TEST(Layer, SelectiveFcLayer_train_sparse_mul) { selLayerConfig.set_size(fcLayerWidth); testSelectiveFcLayerTrainSparseMul(selLayerConfig, false); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA testSelectiveFcLayerTrainSparseMul(selLayerConfig, true); #endif } diff --git a/paddle/gserver/tests/test_SeqSliceLayerGrad.cpp b/paddle/gserver/tests/test_SeqSliceLayerGrad.cpp index e1d4ae16176433b898ba88dd60550e44b4fe37be..3366002ca113a3fd1c8a8c597e619a50f1e29931 100644 --- a/paddle/gserver/tests/test_SeqSliceLayerGrad.cpp +++ b/paddle/gserver/tests/test_SeqSliceLayerGrad.cpp @@ -195,7 +195,7 @@ TEST(Layer, SeqSliceLayer) { vector> ends; std::vector mode = {false}; -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA mode.push_back(true); #endif genSeqInfo(seqStartPos, subSeqStartPos); diff --git a/paddle/gserver/tests/test_WarpCTCLayer.cpp b/paddle/gserver/tests/test_WarpCTCLayer.cpp index 55427e2f12fd7b77c6eea1f65b3229e6fd29d71d..da829460061d38f363317e33daeb65cfa705bb8e 100644 --- a/paddle/gserver/tests/test_WarpCTCLayer.cpp +++ b/paddle/gserver/tests/test_WarpCTCLayer.cpp @@ -199,7 +199,7 @@ TEST(Layer, WarpCTCLayer) { for (auto batchSize : {1, 10, 32}) { for (auto normByTimes : {false, true}) { for (auto useGpu : {false, true}) { -#ifdef PADDLE_ONLY_CPU +#ifndef PADDLE_WITH_CUDA if (useGpu) continue; #endif LOG(INFO) << "layerSize=" << layerSize << " batchSize=" << batchSize diff --git a/paddle/math/Matrix.cpp b/paddle/math/Matrix.cpp index 0023b4d0f5da500f380ecb836b7c54e050b13d67..c3e34d5309d9ca8a32d7b0a8043e668cdb5be54b 100644 --- a/paddle/math/Matrix.cpp +++ b/paddle/math/Matrix.cpp @@ -670,7 +670,7 @@ void GpuMatrix::leftMul(Matrix& a, real scaleAB, real scaleT) { } void GpuMatrix::selectRows(Matrix& table, IVector& ids) { -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA CHECK(dynamic_cast(&table)); CHECK(table.useGpu()); CHECK(ids.useGpu()); @@ -694,7 +694,7 @@ void GpuMatrix::selectRows(Matrix& table, IVector& ids) { } void GpuMatrix::addToRows(Matrix& table, IVector& ids) { -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA CHECK(dynamic_cast(&table)); CHECK(table.useGpu()); CHECK(ids.useGpu()); @@ -741,7 +741,7 @@ void GpuMatrix::rowMax(Matrix& max) { } void GpuMatrix::rowMax(IVector& maxIds, Matrix& maxVal) { -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA CHECK(maxIds.useGpu() && maxVal.useGpu()) << "Matrix type are not equal"; size_t numSamples = getHeight(); size_t beam = maxVal.getWidth(); diff --git a/paddle/math/SparseMatrix.cpp b/paddle/math/SparseMatrix.cpp index 6370c77386688a334fa0de8b4e2b272882e9e2b0..284b68d590ba655395c0186d8ea86d6855c6fc50 100644 --- a/paddle/math/SparseMatrix.cpp +++ b/paddle/math/SparseMatrix.cpp @@ -836,7 +836,7 @@ void GpuSparseMatrix::zeroMem() { } void GpuSparseMatrix::rowMax(IVector& maxIds, Matrix& maxVal) { -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA CHECK(maxIds.useGpu() && maxVal.useGpu()) << "Matrix type are not equal"; size_t numSamples = getHeight(); size_t beam = maxVal.getWidth(); diff --git a/paddle/math/Vector.cpp b/paddle/math/Vector.cpp index eb87ee9bb7936d27c0c32a1a4b35ff49871c0a10..ff72672e3ab77212b309fcfea835839a916fa632 100644 --- a/paddle/math/Vector.cpp +++ b/paddle/math/Vector.cpp @@ -172,7 +172,7 @@ void GpuVectorT::isEqualTo(const VectorT& b, const T& value) { template void GpuVectorT::selectFrom(const VectorT& src, const VectorT& ids) { -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA hl_vector_select_from(this->getData(), this->getSize(), src.getData(), @@ -850,7 +850,7 @@ CpuGpuVectorT::CpuGpuVectorT(CpuGpuVectorT& src, size_t size) : sync_(nullptr) { CHECK_LE(offset + size, static_cast(src.getSize())); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA SyncedFlag* flag = src.getSync(); if (*flag == DATA_AT_CPU) { src.copyToGpu(); // will set synchronous data between CPU and GPU @@ -861,7 +861,7 @@ CpuGpuVectorT::CpuGpuVectorT(CpuGpuVectorT& src, auto cMemHandle = (src.getVector(false))->getMemoryHandle(); cpuVectorT_ = std::make_shared>( size, std::dynamic_pointer_cast(cMemHandle), offset); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA auto gMemHandle = (src.getVector(true))->getMemoryHandle(); gpuVectorT_ = std::make_shared>( size, std::dynamic_pointer_cast(gMemHandle), offset); diff --git a/paddle/math/tests/test_Allocator.cpp b/paddle/math/tests/test_Allocator.cpp index 1ca70ea84c867b83013625eaee141f5b75fad4ae..1fecf659e5080c7d25f5f76b92b15f75eaab6ce3 100644 --- a/paddle/math/tests/test_Allocator.cpp +++ b/paddle/math/tests/test_Allocator.cpp @@ -68,7 +68,7 @@ void testPoolAllocator() { TEST(Allocator, Pool) { testPoolAllocator(); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA testPoolAllocator(); #endif } @@ -92,7 +92,7 @@ TEST(MemoryHandle, Cpu) { EXPECT_EQ(ptr1, ptr2); } -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA TEST(MemoryHandle, Gpu) { int numGpu = hl_get_device_count(); diff --git a/paddle/math/tests/test_BaseMatrix.cpp b/paddle/math/tests/test_BaseMatrix.cpp index 22ce39701fca7b650fc03794cb0701e0987d2dae..1766257860b0b13e9f0ce898438e7c2d644f545e 100644 --- a/paddle/math/tests/test_BaseMatrix.cpp +++ b/paddle/math/tests/test_BaseMatrix.cpp @@ -12,7 +12,7 @@ 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. */ -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA /** * This test file use autotest::AutoCompare and cmpWithoutArg to compares the * implementation of CPU and GPU member function in diff --git a/paddle/math/tests/test_CpuGpuVector.cpp b/paddle/math/tests/test_CpuGpuVector.cpp index 58bc43a38ba9465a832fcd0652e6309c403577e3..c72f89c8244b1209e490b09387c2ee6352426ce1 100644 --- a/paddle/math/tests/test_CpuGpuVector.cpp +++ b/paddle/math/tests/test_CpuGpuVector.cpp @@ -12,7 +12,7 @@ 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. */ -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA #include #include "paddle/math/Vector.h" diff --git a/paddle/math/tests/test_ExecViaCpu.cpp b/paddle/math/tests/test_ExecViaCpu.cpp index 04c856453d2ec4ad764e37ae430e3e30ac0dea0b..25e0ba11ded96dd78aedc3c297507d0555d80d74 100644 --- a/paddle/math/tests/test_ExecViaCpu.cpp +++ b/paddle/math/tests/test_ExecViaCpu.cpp @@ -94,7 +94,7 @@ void testWrapper(F&& f) { } } -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA TEST(ExecViaCpu, test1) { testWrapper(f); testWrapper(&f); diff --git a/paddle/math/tests/test_GpuProfiler.cpp b/paddle/math/tests/test_GpuProfiler.cpp index e6b5dba446b5a0022ade76b188895c4e0e2a22b4..9402bd3ec48fbed381ef1f676e8b179cabd4cb9f 100644 --- a/paddle/math/tests/test_GpuProfiler.cpp +++ b/paddle/math/tests/test_GpuProfiler.cpp @@ -12,7 +12,7 @@ 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. */ -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA #include #include "paddle/math/Matrix.h" diff --git a/paddle/math/tests/test_Matrix.cpp b/paddle/math/tests/test_Matrix.cpp index 1c21da5b76e95603258a5006d0c57b00126e65b9..2f99fa3581e14b91acc0b294856619f4ae2b3483 100644 --- a/paddle/math/tests/test_Matrix.cpp +++ b/paddle/math/tests/test_Matrix.cpp @@ -12,7 +12,7 @@ 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. */ -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA /** * This test file use autotest::AutoCompare and cmpWithArg to compares the * implementation of CPU and GPU member function in Matrix.cpp. diff --git a/paddle/math/tests/test_SparseMatrix.cpp b/paddle/math/tests/test_SparseMatrix.cpp index c0572dfdbf738a4dfad04811b3a3e1b65487ff6d..8abbe8d82e02b7d1738fe7e6d0c8d494166e7892 100644 --- a/paddle/math/tests/test_SparseMatrix.cpp +++ b/paddle/math/tests/test_SparseMatrix.cpp @@ -47,7 +47,7 @@ struct MatrixPara { SparseFormat format; }; -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA void test_sparse_matrix_mul(MatrixPara paraA, MatrixPara paraB, MatrixPara paraC) { @@ -452,7 +452,7 @@ TEST(Matrix, SparseMatrixCSRFormatTrimFrom) { matB->trimFrom(*mat); checkSMatrixEqual2(matA, matB); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA GpuSparseMatrixPtr matC = std::make_shared( height, trimedWidth, height, FLOAT_VALUE, SPARSE_CSR, true); matC->trimFrom(*mat); @@ -546,7 +546,7 @@ TEST(Matrix, SparseMatrixCSCFormatTrimFrom) { matB->trimFrom(*mat); checkSMatrixEqual2(matA, matB); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA GpuSparseMatrixPtr matC = std::make_shared( height, trimedWidth, height, FLOAT_VALUE, SPARSE_CSC, true); matC->trimFrom(*mat); diff --git a/paddle/math/tests/test_Tensor.cu b/paddle/math/tests/test_Tensor.cu index 31b693afa8bd50f77a8efb67769e6215dd755bd3..d03698dee25fdd6dd49f2a3fdb5c605333440f49 100644 --- a/paddle/math/tests/test_Tensor.cu +++ b/paddle/math/tests/test_Tensor.cu @@ -270,7 +270,7 @@ TEST(Unary, BaseOp) { TestUnaryVectorT testCpuIVector( testUnaryBaseOpInt); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_GPU TestUnaryMatrix testGpuMatrix(testUnaryBaseOp); TestUnaryVectorT testGpuVector(testUnaryBaseOp); TestUnaryVectorT testGpuIVector( @@ -317,7 +317,7 @@ void testUnayrMathOp(Tensor& A1, Tensor& A2) { TEST(Unary, MathOp) { TestUnaryMatrix testCpu(testUnayrMathOp); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_GPU TestUnaryMatrix testGpu(testUnayrMathOp); #endif } @@ -374,7 +374,7 @@ void testUnayrCompareOp(Tensor& A1, Tensor& A2) { TEST(Unary, CompareOp) { TestUnaryMatrix testCpu(testUnayrCompareOp); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_GPU TestUnaryMatrix testGpu(testUnayrCompareOp); #endif } @@ -536,7 +536,7 @@ void testBinaryBaseOp(Tensor& A1, Tensor& A2, Tensor& B) { TEST(Binary, BaseOp) { TestBinaryMatrix testCpu(testBinaryBaseOp); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_GPU TestBinaryMatrix testGpu(testBinaryBaseOp); #endif } @@ -710,7 +710,7 @@ void testBinaryMathOp(Tensor& A1, Tensor& A2, Tensor& B) { TEST(Binary, MathOp) { TestBinaryMatrix testCpu(testBinaryMathOp); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_GPU TestBinaryMatrix testGpu(testBinaryMathOp); #endif } @@ -810,7 +810,7 @@ void testBinaryCompareOp(Tensor& A1, Tensor& A2, Tensor& B) { TEST(Binary, CompareOp) { TestBinaryMatrix testCpu(testBinaryCompareOp); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_GPU TestBinaryMatrix testGpu(testBinaryCompareOp); #endif } @@ -955,7 +955,7 @@ void testTernaryBaseOp(Tensor& A1, Tensor& A2, Tensor& B, Tensor& C) { TEST(Ternary, BaseOp) { TestTernaryMatrix testCpu(testTernaryBaseOp); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_GPU TestTernaryMatrix testGpu(testTernaryBaseOp); #endif } @@ -1058,7 +1058,7 @@ void testTernaryCompareOp(Tensor& A1, Tensor& A2, Tensor& B, Tensor& C) { TEST(Ternary, CompareOp) { TestTernaryMatrix testCpu(testTernaryCompareOp); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_GPU TestTernaryMatrix testGpu(testTernaryCompareOp); #endif } @@ -1086,7 +1086,7 @@ void testQuaternaryAdd( TEST(Quaternary, BaseOp) { TestQuaternaryMatrix testCpu(testQuaternaryAdd); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_GPU TestQuaternaryMatrix testGpu(testQuaternaryAdd); #endif } @@ -1156,7 +1156,7 @@ void testQuaternaryCompareOp( TEST(Quaternary, CompareOp) { TestQuaternaryMatrix testCpu(testQuaternaryCompareOp); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_GPU TestQuaternaryMatrix testGpu(testQuaternaryCompareOp); #endif } diff --git a/paddle/math/tests/test_TrainingAlgorithm.cpp b/paddle/math/tests/test_TrainingAlgorithm.cpp index 4a88844b43ef40af988d2b391d2bef4568dea9b7..5ae0aa036f6bfc1e5bd4e955277c4efff8c739ce 100644 --- a/paddle/math/tests/test_TrainingAlgorithm.cpp +++ b/paddle/math/tests/test_TrainingAlgorithm.cpp @@ -91,7 +91,7 @@ int VectorCheckErr(const VectorPtr& vector1, const VectorPtr& vector2) { typedef std::function testMatrixFunc; void testCase(testMatrixFunc matrixFunc) { -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA for (auto useGpu : {false, true}) { #else for (auto useGpu : {false}) { diff --git a/paddle/math/tests/test_batchTranspose.cpp b/paddle/math/tests/test_batchTranspose.cpp index 4eb9837909ffaaf0f483ab65ece7a0b29fd49319..b70a61976402fd0a7cfee8382fd926fcf28486d5 100644 --- a/paddle/math/tests/test_batchTranspose.cpp +++ b/paddle/math/tests/test_batchTranspose.cpp @@ -17,7 +17,7 @@ limitations under the License. */ using namespace paddle; // NOLINT -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA TEST(MatrixBatchTransTest, test_batch_matrix_transpose) { const int nx = 100; const int ny = 50; diff --git a/paddle/math/tests/test_lazyAssign.cu b/paddle/math/tests/test_lazyAssign.cu index 92afab4ff7f5ff4acc219c5ac783733340c5726a..04f23cff55db45c39049538545430bc5996cce5d 100644 --- a/paddle/math/tests/test_lazyAssign.cu +++ b/paddle/math/tests/test_lazyAssign.cu @@ -72,7 +72,7 @@ void testLazyAssign(int height, int width) { TEST(lazyAssign, CPU) { testMatrixCase(testLazyAssign); } -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_GPU TEST(lazyAssign, GPU) { testMatrixCase(testLazyAssign); } #endif @@ -142,6 +142,6 @@ void testSgdUpdate(int height, int width) { TEST(sgdUpdate, CPU) { testMatrixCase(testSgdUpdate); } -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_GPU TEST(sgdUpdate, GPU) { testMatrixCase(testSgdUpdate); } #endif diff --git a/paddle/math/tests/test_matrixCompare.cpp b/paddle/math/tests/test_matrixCompare.cpp index 061fb22e3fd744d9d9895fd1008089e4a6ce6a0f..7e5a1db44a5302e3b4e5d2768755824666e880ba 100644 --- a/paddle/math/tests/test_matrixCompare.cpp +++ b/paddle/math/tests/test_matrixCompare.cpp @@ -12,7 +12,7 @@ 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. */ -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA /// This unittest checks GpuMatrix/CpuMatrix get same result, so disable when /// only cpu version. diff --git a/paddle/math/tests/test_perturbation.cpp b/paddle/math/tests/test_perturbation.cpp index 60ebae015381a3901c14d0cd4c1225e54ac5726f..c7c07c817a08d78ddcbf8218e8c4a9d22f4990bc 100644 --- a/paddle/math/tests/test_perturbation.cpp +++ b/paddle/math/tests/test_perturbation.cpp @@ -12,7 +12,7 @@ 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. */ -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA #include #include diff --git a/paddle/math/tests/test_sparseMatrixCompare.cpp b/paddle/math/tests/test_sparseMatrixCompare.cpp index a9185a4b24b13ca0287b0f67375c4599e8b9ac78..2b2a391b9d04a9f7fa4986a6b6dd5cd8e5385f1f 100644 --- a/paddle/math/tests/test_sparseMatrixCompare.cpp +++ b/paddle/math/tests/test_sparseMatrixCompare.cpp @@ -12,7 +12,7 @@ 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. */ -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA /// This unittest checks GpuSparseMatrix/CpuSparseMatrix get same result, // so disable when /// only cpu version. diff --git a/paddle/memory/detail/buddy_allocator.cc b/paddle/memory/detail/buddy_allocator.cc index bb44970109c05d239e6b92d90b2079b752fa0104..fdc5ed19dc2973e744676c3b795c8ab86da58590 100644 --- a/paddle/memory/detail/buddy_allocator.cc +++ b/paddle/memory/detail/buddy_allocator.cc @@ -175,7 +175,7 @@ void* BuddyAllocator::SystemAlloc(size_t size) { } BuddyAllocator::PoolSet::iterator BuddyAllocator::RefillPool() { -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA if (system_allocator_->UseGpu()) { if ((total_used_ + total_free_) == 0) { // Compute the maximum allocation size for the first allocation. diff --git a/paddle/memory/detail/system_allocator.cc b/paddle/memory/detail/system_allocator.cc index a270bd59581520859d43cddd2fc0cfa72080f46d..6c9a46dd09c15347fca1a30971e7e732d887bc8e 100644 --- a/paddle/memory/detail/system_allocator.cc +++ b/paddle/memory/detail/system_allocator.cc @@ -62,7 +62,7 @@ void CPUAllocator::Free(void* p, size_t size, size_t index) { bool CPUAllocator::UseGpu() const { return false; } -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA void* GPUAllocator::Alloc(size_t& index, size_t size) { // CUDA documentation doesn't explain if cudaMalloc returns nullptr diff --git a/paddle/memory/detail/system_allocator.h b/paddle/memory/detail/system_allocator.h index 82ba322e057575c460b1d51d719c9b0fa459273e..ee9b012f91a9647839cf465c4074082f2d3509a6 100644 --- a/paddle/memory/detail/system_allocator.h +++ b/paddle/memory/detail/system_allocator.h @@ -40,7 +40,7 @@ class CPUAllocator : public SystemAllocator { virtual bool UseGpu() const; }; -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA class GPUAllocator : public SystemAllocator { public: virtual void* Alloc(size_t& index, size_t size); diff --git a/paddle/memory/detail/system_allocator_test.cc b/paddle/memory/detail/system_allocator_test.cc index ba44e06ddb68e92e4086a8006b868557b0c89b50..cd563844e7fa23241bb0bb56d1365ef34826c4a8 100644 --- a/paddle/memory/detail/system_allocator_test.cc +++ b/paddle/memory/detail/system_allocator_test.cc @@ -56,7 +56,7 @@ TEST(CPUAllocator, LockMem) { TestAllocator(a, 0); } -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA TEST(GPUAllocator, Alloc) { paddle::memory::detail::GPUAllocator a; TestAllocator(a, 2048); diff --git a/paddle/memory/memcpy.cc b/paddle/memory/memcpy.cc index c96a697a7e022684688b31c05da43e52812100d8..790420a8ab41b1a61ee35dc086c8b95fa1a02019 100644 --- a/paddle/memory/memcpy.cc +++ b/paddle/memory/memcpy.cc @@ -26,7 +26,7 @@ void Copy(platform::CPUPlace, void* dst, std::memcpy(dst, src, num); } -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA template <> void Copy(platform::CPUPlace dst_place, void* dst, diff --git a/paddle/memory/memcpy.h b/paddle/memory/memcpy.h index 2b9c0eada6e8406fc81baec7f331a8dd5b8b0ec1..0bccee58c3a22379c75523467e0c717b98b08bcf 100644 --- a/paddle/memory/memcpy.h +++ b/paddle/memory/memcpy.h @@ -33,7 +33,7 @@ namespace memory { template void Copy(DstPlace, void* dst, SrcPlace, const void* src, size_t num); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA /** * \brief Copy memory from one place to another place. diff --git a/paddle/memory/memory.cc b/paddle/memory/memory.cc index 29bc26f9d3bca0e30896657431f9a9bb1dac0d1d..30ce8a82e16ed26a41b009ce5d52dd1a2a1b7c21 100644 --- a/paddle/memory/memory.cc +++ b/paddle/memory/memory.cc @@ -62,7 +62,7 @@ size_t Used(platform::CPUPlace place) { return GetCPUBuddyAllocator()->Used(); } -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA BuddyAllocator* GetGPUBuddyAllocator(int gpu_id) { using BuddyAllocVec = std::vector; @@ -77,7 +77,7 @@ BuddyAllocator* GetGPUBuddyAllocator(int gpu_id) { // GPU buddy allocator initialization std::call_once(gpu_allocator_flag, [&]() { - int gpu_num = platform::GetDeviceCount(); + int gpu_num = platform::GetCUDADeviceCount(); allocators.reserve(gpu_num); for (int gpu = 0; gpu < gpu_num; gpu++) { platform::SetDeviceId(gpu); diff --git a/paddle/memory/memory_test.cc b/paddle/memory/memory_test.cc index 53cc63a098d0802479e3a371717adb7596c249ed..0d402038a06f4ad93fd15946fc44aaeac58ada40 100644 --- a/paddle/memory/memory_test.cc +++ b/paddle/memory/memory_test.cc @@ -80,7 +80,7 @@ TEST(BuddyAllocator, CPUMultAlloc) { } } -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA size_t align(size_t size, paddle::platform::GPUPlace place) { size += sizeof(paddle::memory::detail::Metadata); diff --git a/paddle/operators/CMakeLists.txt b/paddle/operators/CMakeLists.txt index 43eb4de2c1d4ff61f295c05af9dfd5f26c2ba365..0fa1fca2bcd3117e1e9a6a54c343b2d0d8c3822b 100644 --- a/paddle/operators/CMakeLists.txt +++ b/paddle/operators/CMakeLists.txt @@ -103,12 +103,16 @@ set(DEPS_OPS recurrent_op cond_op cross_entropy_op - softmax_with_cross_entropy_op) + softmax_with_cross_entropy_op + sum_op) + + op_library(recurrent_op SRCS recurrent_op.cc rnn/recurrent_op_utils.cc DEPS framework_proto tensor net_op) op_library(cond_op SRCS cond_op.cc DEPS framework_proto tensor operator net_op) op_library(cross_entropy_op DEPS cross_entropy) op_library(softmax_with_cross_entropy_op DEPS cross_entropy softmax) +op_library(sum_op DEPS net_op) list(REMOVE_ITEM GENERAL_OPS ${DEPS_OPS}) foreach(src ${GENERAL_OPS}) diff --git a/paddle/operators/activation_op.cc b/paddle/operators/activation_op.cc index 7ae4d2f6b6c0b0f30c06adc34c811bfe34b59fa6..5f2ecc2673a5bf87d07e99faf60d585fb02621d0 100644 --- a/paddle/operators/activation_op.cc +++ b/paddle/operators/activation_op.cc @@ -69,6 +69,22 @@ class ReluOpMaker : public framework::OpProtoAndCheckerMaker { } }; +template +class LeakyReluOpMaker : public framework::OpProtoAndCheckerMaker { + public: + LeakyReluOpMaker(framework::OpProto *proto, + framework::OpAttrChecker *op_checker) + : OpProtoAndCheckerMaker(proto, op_checker) { + AddInput("X", "Input of LeakyRelu operator"); + AddOutput("Y", "Output of LeakyRelu operator"); + AddComment( + "LeakyRelu activation operator, " + "leaky_relu = max(x, alpha * x)"); + AddAttr("alpha", "The small negative slope") + .SetDefault(static_cast(0.02f)); + } +}; + class TanhOpMaker : public framework::OpProtoAndCheckerMaker { public: TanhOpMaker(framework::OpProto *proto, framework::OpAttrChecker *op_checker) @@ -240,6 +256,9 @@ REGISTER_OP(softsign, ops::ActivationOp, ops::SoftsignOpMaker, softsign_grad, REGISTER_OP(brelu, ops::ActivationOp, ops::BReluOpMaker, brelu_grad, ops::ActivationOpGrad); +REGISTER_OP(leaky_relu, ops::ActivationOp, ops::LeakyReluOpMaker, + leaky_relu_grad, ops::ActivationOpGrad); + REGISTER_OP(soft_relu, ops::ActivationOp, ops::SoftReluOpMaker, soft_relu_grad, ops::ActivationOpGrad); diff --git a/paddle/operators/activation_op.h b/paddle/operators/activation_op.h index ff35c2d97e856ab76581c74512a0b451ea6fe60c..dae66cc77d9103a6a2b13e69b3f014f8de313209 100644 --- a/paddle/operators/activation_op.h +++ b/paddle/operators/activation_op.h @@ -309,6 +309,33 @@ struct SoftReluGradFunctor : public BaseActivationFunctor { } }; +template +struct LeakyReluFunctor : public BaseActivationFunctor { + float alpha; + typename BaseActivationFunctor::AttrPair GetAttrs() { + return {{"alpha", &alpha}}; + } + + template + void operator()(Device d, X x, Y y) const { + y.device(d) = x.cwiseMax(alpha * x); + } +}; + +template +struct LeakyReluGradFunctor : public BaseActivationFunctor { + float alpha; + typename BaseActivationFunctor::AttrPair GetAttrs() { + return {{"alpha", &alpha}}; + } + template + void operator()(Device d, X x, Y y, dY dy, dX dx) const { + auto temp1 = alpha * (x < static_cast(0)).template cast().eval(); + auto temp2 = (x >= static_cast(0)).template cast().eval(); + dx.device(d) = dy * (temp1 + temp2).template cast(); + } +}; + template struct PowFunctor : public BaseActivationFunctor { float factor; @@ -379,4 +406,5 @@ struct STanhGradFunctor : public BaseActivationFunctor { __macro(soft_relu, SoftReluFunctor, SoftReluGradFunctor); \ __macro(pow, PowFunctor, PowGradFunctor); \ __macro(stanh, STanhFunctor, STanhGradFunctor); \ - __macro(softsign, SoftsignFunctor, SoftsignGradFunctor) + __macro(softsign, SoftsignFunctor, SoftsignGradFunctor); \ + __macro(leaky_relu, LeakyReluFunctor, LeakyReluGradFunctor) diff --git a/paddle/operators/adadelta_op.cc b/paddle/operators/adadelta_op.cc new file mode 100644 index 0000000000000000000000000000000000000000..bd8c93b4a193c169a4472ff20efca779ddc5c804 --- /dev/null +++ b/paddle/operators/adadelta_op.cc @@ -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 "paddle/operators/adadelta_op.h" + +namespace paddle { +namespace operators { + +class AdadeltaOp : public framework::OperatorWithKernel { + public: + using framework::OperatorWithKernel::OperatorWithKernel; + + protected: + void InferShape(framework::InferShapeContextBase *ctx) const override { + PADDLE_ENFORCE(ctx->HasInput("Param"), + "Input(Param) of AdadeltaOp should not be null."); + PADDLE_ENFORCE(ctx->HasInput("Grad"), + "Input(Grad) of AdadeltaOp should not be null."); + PADDLE_ENFORCE(ctx->HasInput("AvgSquaredGrad"), + "Input(AvgSquaredGrad) of AdadeltaOp should not be null."); + PADDLE_ENFORCE(ctx->HasInput("AvgSquaredUpdate"), + "Input(AvgSquaredUpdate) of AdadeltaOp should not be null."); + + PADDLE_ENFORCE(ctx->HasOutput("ParamOut"), + "Output(ParamOut) of AdadeltaOp should not be null."); + PADDLE_ENFORCE( + ctx->HasOutput("AvgSquaredGradOut"), + "Output(AvgSquaredGradOut) of AdadeltaOp should not be null."); + PADDLE_ENFORCE( + ctx->HasOutput("AvgSquaredUpdateOut"), + "Output(AvgSquaredUpdateOut) of AdadeltaOp should not be null."); + + auto param_dim = ctx->GetInputDim("Param"); + PADDLE_ENFORCE_EQ( + param_dim, ctx->GetInputDim("Grad"), + "param and grad input of AdadeltaOp should have same dimension"); + PADDLE_ENFORCE_EQ(param_dim, ctx->GetInputDim("AvgSquaredGrad"), + "Param and AvgSquaredGrad input of AdadeltaOp " + "should have same dimension"); + PADDLE_ENFORCE_EQ(param_dim, ctx->GetInputDim("AvgSquaredUpdate"), + "Param and AvgSquaredUpdate input of AdadeltaOp " + "should have same dimension"); + + ctx->SetOutputDim("ParamOut", param_dim); + ctx->SetOutputDim("AvgSquaredGradOut", param_dim); + ctx->SetOutputDim("AvgSquaredUpdateOut", param_dim); + } +}; + +class AdadeltaOpMaker : public framework::OpProtoAndCheckerMaker { + public: + AdadeltaOpMaker(framework::OpProto *proto, + framework::OpAttrChecker *op_checker) + : OpProtoAndCheckerMaker(proto, op_checker) { + AddInput("Param", "(Tensor) Input parameter"); + AddInput("Grad", "(Tensor) Input gradient"); + AddInput("AvgSquaredGrad", + "(Tensor) Input expectation of squared gradient"); + AddInput("AvgSquaredUpdate", + "(Tensor) Input expectation of squared parameter updates"); + + AddOutput("ParamOut", "(Tensor) Output parameter"); + AddOutput("AvgSquaredGradOut", + "(Tensor) Output expectation of squared gradient"); + AddOutput("AvgSquaredUpdateOut", + "(Tensor) Output expectation of squared parameter updates"); + + AddAttr("rho", + "(float, default 0.95) Exponential decay rate " + "for squared gradients.") + .SetDefault(0.95f); + AddAttr("epsilon", + "(float, default 1.0e-6) Constant for " + "numerical stability") + .SetDefault(1.0e-6f); + AddComment(R"DOC( +Adadelta Updates Operator. + +This implements the Adadelta optimizer[1]. Adadelta is a per-dimension +adaptive learning rate method for gradient descent. + +Adadelta updates: + +avg_squared_grad_out = rho * avg_squared_grad + (1 - rho) * grad * grad +param_update = - sqrt((avg_squared_update + epsilon) / + (avg_squared_grad_out + epsilon)) * grad +avg_squared_update_out = rho * avg_squared_update + (1 - rho) * param_update**2 +param_out = param + param_update + +References: + [1] ADADELTA: An Adaptive Learning Rate Method + https://arxiv.org/abs/1212.5701 + +)DOC"); + } +}; + +} // namespace operators +} // namespace paddle + +namespace ops = paddle::operators; +REGISTER_OP_WITHOUT_GRADIENT(adadelta, ops::AdadeltaOp, ops::AdadeltaOpMaker); +REGISTER_OP_CPU_KERNEL( + adadelta, ops::AdadeltaOpKernel); diff --git a/paddle/operators/adadelta_op.cu b/paddle/operators/adadelta_op.cu new file mode 100644 index 0000000000000000000000000000000000000000..3af1c8c8e9861138a33b3156818f704c3b20363f --- /dev/null +++ b/paddle/operators/adadelta_op.cu @@ -0,0 +1,20 @@ +/* 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. */ + +#define EIGEN_USE_GPU +#include "paddle/operators/adadelta_op.h" + +namespace ops = paddle::operators; +REGISTER_OP_GPU_KERNEL( + adadelta, ops::AdadeltaOpKernel); diff --git a/paddle/operators/adadelta_op.h b/paddle/operators/adadelta_op.h new file mode 100644 index 0000000000000000000000000000000000000000..d29e15c43583bd447fbacb548a326f303f7d1463 --- /dev/null +++ b/paddle/operators/adadelta_op.h @@ -0,0 +1,69 @@ +/* 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 AdadeltaOpKernel : public framework::OpKernel { + public: + void Compute(const framework::ExecutionContext& ctx) const override { + auto param_out_tensor = ctx.Output("ParamOut"); + auto avg_squared_grad_out_tensor = + ctx.Output("AvgSquaredGradOut"); + auto avg_squared_update_out_tensor = + ctx.Output("AvgSquaredUpdateOut"); + + param_out_tensor->mutable_data(ctx.GetPlace()); + avg_squared_grad_out_tensor->mutable_data(ctx.GetPlace()); + avg_squared_update_out_tensor->mutable_data(ctx.GetPlace()); + + float rho = ctx.Attr("rho"); + float epsilon = ctx.Attr("epsilon"); + + auto param = framework::EigenVector::Flatten( + *ctx.Input("Param")); + auto grad = framework::EigenVector::Flatten( + *ctx.Input("Grad")); + // Squared gradient accumulator + auto avg_squared_grad = framework::EigenVector::Flatten( + *ctx.Input("AvgSquaredGrad")); + // Squared updates accumulator + auto avg_squared_update = framework::EigenVector::Flatten( + *ctx.Input("AvgSquaredUpdate")); + auto param_out = framework::EigenVector::Flatten(*param_out_tensor); + auto avg_squared_grad_out = + framework::EigenVector::Flatten(*avg_squared_grad_out_tensor); + auto avg_squared_update_out = + framework::EigenVector::Flatten(*avg_squared_update_out_tensor); + auto place = ctx.GetEigenDevice(); + + avg_squared_grad_out.device(place) = + rho * avg_squared_grad + (1 - rho) * grad.square(); + auto update = + -((avg_squared_update + epsilon) / (avg_squared_grad_out + epsilon)) + .sqrt() * + grad; + avg_squared_update_out.device(place) = + rho * avg_squared_update + (1 - rho) * update.square(); + param_out.device(place) = param + update; + } +}; + +} // namespace operators +} // namespace paddle diff --git a/paddle/operators/cond_op.cc b/paddle/operators/cond_op.cc index db20b69f3fab1a5a5bfcd445536437dc4b428559..2737104a205cbc1e18ce4a3a45592a416d38a874 100644 --- a/paddle/operators/cond_op.cc +++ b/paddle/operators/cond_op.cc @@ -126,8 +126,7 @@ void CondOp::PrepareDataForSubnet( dim[0] = index_tensors[i].dims()[0]; tensor_child->mutable_data(dim, platform::CPUPlace()); - Gather(dev_ctx.GetPlace(), tensor_parent, &index_tensors[i], - tensor_child); + CPUGather(dev_ctx, *tensor_parent, index_tensors[i], tensor_child); } } @@ -188,7 +187,7 @@ void CondOp::MergeDataFromSubnet(const framework::Scope& scope, Variable* var_child = sub_scopes[i]->FindVar(output); PADDLE_ENFORCE_NOT_NULL(var_child); auto* tensor_child = &var_child->Get(); - ScatterUpdate(dev_ctx.GetPlace(), tensor_child, &index_tensors[i], + ScatterAssign(dev_ctx, *tensor_child, index_tensors[i], tensor_parent); } } diff --git a/paddle/operators/detail/strided_memcpy.h b/paddle/operators/detail/strided_memcpy.h index b165224b37fb091c094a823179256c3dd40a37c9..068c82f399316a1587d7322d8dab75823656800e 100644 --- a/paddle/operators/detail/strided_memcpy.h +++ b/paddle/operators/detail/strided_memcpy.h @@ -34,7 +34,7 @@ struct StridedMemcpyFunctor { auto& cpu_place = boost::get(place); memory::Copy(cpu_place, dst, cpu_place, src, sizeof(T) * dst_dim.head); } else { -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA auto& gpu_place = boost::get(place); auto& cuda_ctx = reinterpret_cast(dev_ctx); diff --git a/paddle/operators/gather.cu.h b/paddle/operators/gather.cu.h new file mode 100644 index 0000000000000000000000000000000000000000..8d04ecd284226c7b4c6cdd5531915fee2d94ce61 --- /dev/null +++ b/paddle/operators/gather.cu.h @@ -0,0 +1,79 @@ +/* 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/tensor.h" +#include "paddle/platform/place.h" + +namespace paddle { +namespace operators { + +using framework::Tensor; +using platform::Place; + +#define CUDA_1D_KERNEL_LOOP(i, n) \ + for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < (n); \ + i += blockDim.x * gridDim.x) + +template +__global__ void GatherCUDAKernel(const T* params, const int* indices, T* output, + size_t index_size, size_t slice_size) { + CUDA_1D_KERNEL_LOOP(i, index_size * slice_size) { + int indices_i = i / slice_size; + int slice_i = i - indices_i * slice_size; // offset inside the slice + int gather_i = indices[indices_i]; + int params_i = gather_i * slice_size + slice_i; + *(output + i) = *(params + params_i); + } +} + +/** + * A thin wrapper on gpu tensor + * Return a new tensor from source tensor, gathered according to index + * input[src]: type-T source Tensor + * input[index]: type-int index Tensor (1-D) + * return: output tensor + */ +template +void GPUGather(const platform::DeviceContext& ctx, const Tensor& src, + const Tensor& index, Tensor* output) { + // PADDLE_ENFORCE(platform::is_gpu_place(place)); + // check index of shape 1-D + PADDLE_ENFORCE(index.dims().size() == 1); + int index_size = index.dims()[0]; + + auto src_dims = src.dims(); + framework::DDim output_dims(src_dims); + output_dims[0] = index_size; + + // slice size + int slice_size = 1; + for (int i = 1; i < src_dims.size(); ++i) slice_size *= src_dims[i]; + + const T* p_src = src.data(); + const int* p_index = index.data(); + T* p_output = output->data(); + + int block = 512; + int n = slice_size * index_size; + int grid = (n + block - 1) / block; + + GatherCUDAKernel<<< + grid, block, 0, + reinterpret_cast(ctx).stream()>>>( + p_src, p_index, p_output, index_size, slice_size); +} + +} // namespace operators +} // namespace paddle diff --git a/paddle/operators/gather.h b/paddle/operators/gather.h index 92fb51ec17709bc6f8abb2f516a9240fb5dc3a77..052db49cb3c2594eca8b9a5e3716689480089703 100644 --- a/paddle/operators/gather.h +++ b/paddle/operators/gather.h @@ -24,49 +24,40 @@ limitations under the License. */ namespace paddle { namespace operators { -// Implementation of CPU copy -template -void CPUGather(const T* src, const int* indices, const int slice_size, - const int index_size, T* output) { - const size_t slice_bytes = slice_size * sizeof(T); - - for (int i = 0; i < index_size; ++i) { - int index_ = indices[i]; - memcpy(output + i * slice_size, src + index_ * slice_size, slice_bytes); - } -} - -// Implementation of GPU copy: -template -void GPUGather(const T* src, const int* index, const int slice_size, - const int index_size, T* output); +using framework::Tensor; /** + * A thin wrapper for gathering on cpu tensor * Return a new tensor from source tensor, gathered according to index * input[src]: type-T source Tensor * input[index]: type-int index Tensor (1-D) * return: output tensor */ template -void Gather(const platform::Place& place, const paddle::framework::Tensor* src, - const paddle::framework::Tensor* index, - paddle::framework::Tensor* output) { +void CPUGather(const platform::DeviceContext& ctx, const Tensor& src, + const Tensor& index, Tensor* output) { + PADDLE_ENFORCE(platform::is_cpu_place(ctx.GetPlace())); // check index of shape 1-D - PADDLE_ENFORCE(index->dims().size() == 1); - int index_size = index->dims()[0]; + PADDLE_ENFORCE(index.dims().size() == 1); + int index_size = index.dims()[0]; - auto src_dims = src->dims(); + auto src_dims = src.dims(); framework::DDim output_dims(src_dims); output_dims[0] = index_size; + const T* p_src = src.data(); + const int* p_index = index.data(); + T* p_output = output->data(); + // slice size int slice_size = 1; for (int i = 1; i < src_dims.size(); ++i) slice_size *= src_dims[i]; - // Gathering - if (platform::is_cpu_place(place)) { - CPUGather(src->data(), index->data(), slice_size, index_size, - output->data()); + const size_t slice_bytes = slice_size * sizeof(T); + + for (int i = 0; i < index_size; ++i) { + int index_ = p_index[i]; + memcpy(p_output + i * slice_size, p_src + index_ * slice_size, slice_bytes); } } diff --git a/paddle/operators/gather_op.cc b/paddle/operators/gather_op.cc index da22bd0c52c27d7decd10e2e2b34fa38d0620da8..fe305337cbebd7c679ae1b8ee8aa2740472ee109 100644 --- a/paddle/operators/gather_op.cc +++ b/paddle/operators/gather_op.cc @@ -31,6 +31,8 @@ class GatherOp : public framework::OperatorWithKernel { PADDLE_ENFORCE(ctx->HasOutput("Out"), "Output(Out) of GatherOp should not be null."); + auto index_dims = ctx->GetInputDim("Index"); + PADDLE_ENFORCE(index_dims.size() == 1); int batch_size = ctx->GetInputDim("Index")[0]; PADDLE_ENFORCE_GE(batch_size, 0, "Batch size must be >0"); framework::DDim output_dims(ctx->GetInputDim("X")); @@ -79,8 +81,5 @@ Out = X[Index] namespace ops = paddle::operators; REGISTER_OP(gather, ops::GatherOp, ops::GatherOpMaker, gather_grad, ops::GatherGradOp); -REGISTER_OP_CPU_KERNEL(gather, - ops::GatherOpKernel); -REGISTER_OP_CPU_KERNEL( - gather_grad, - ops::GatherGradientOpKernel); +REGISTER_OP_CPU_KERNEL(gather, ops::GatherOpKernel); +REGISTER_OP_CPU_KERNEL(gather_grad, ops::GatherGradientOpKernel); diff --git a/paddle/operators/gather_op.cu b/paddle/operators/gather_op.cu new file mode 100644 index 0000000000000000000000000000000000000000..92219d6a433e6db0bb9886ed8670cbafaa843ff8 --- /dev/null +++ b/paddle/operators/gather_op.cu @@ -0,0 +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 + + 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 "gather.cu.h" +#include "paddle/framework/eigen.h" +#include "paddle/operators/gather_op.h" +#include "scatter.cu.h" + +namespace paddle { +namespace operators { + +template +class GatherOpCUDAKernel : public framework::OpKernel { + public: + void Compute(const framework::ExecutionContext &ctx) const override { + PADDLE_ENFORCE(platform::is_gpu_place(ctx.GetPlace()), + "This kernel only runs on GPU device."); + auto *x = ctx.Input("X"); + auto *index = ctx.Input("Index"); + auto *output = ctx.Output("Out"); + + output->mutable_data(ctx.GetPlace()); + + GPUGather(ctx.device_context(), *x, *index, output); + } +}; + +template +class GatherGradOpCUDAKernel : public framework::OpKernel { + public: + void Compute(const framework::ExecutionContext &ctx) const override { + PADDLE_ENFORCE(platform::is_gpu_place(ctx.GetPlace()), + "This kernel only runs on GPU device."); + auto *Index = ctx.Input("Index"); + auto *dX = ctx.Output(framework::GradVarName("X")); + auto *dO = ctx.Input(framework::GradVarName("Out")); + auto *x = ctx.Input("X"); + + dX->mutable_data(ctx.GetPlace()); + auto dxt = framework::EigenVector::Flatten(*dX); + auto place = ctx.GetEigenDevice(); + dxt.device(place) = dxt.constant(static_cast(0)); + + GPUScatterAssign(ctx.device_context(), *dO, *Index, dX); + } +}; + +} // namespace operators +} // namespace paddle + +namespace ops = paddle::operators; +REGISTER_OP_GPU_KERNEL(gather, ops::GatherOpCUDAKernel); +REGISTER_OP_GPU_KERNEL(gather_grad, ops::GatherGradOpCUDAKernel); diff --git a/paddle/operators/gather_op.h b/paddle/operators/gather_op.h index 073e566e8f6962d62cc1b738672843421dcb4ee5..8276ed0d3d8b676aafab45fae70942e78b72b8e6 100644 --- a/paddle/operators/gather_op.h +++ b/paddle/operators/gather_op.h @@ -23,29 +23,40 @@ namespace operators { using Tensor = framework::Tensor; -template +template class GatherOpKernel : public framework::OpKernel { public: void Compute(const framework::ExecutionContext &ctx) const override { - auto *X = ctx.Input("X"); - auto *Index = ctx.Input("Index"); - auto *Y = ctx.Output("Out"); + PADDLE_ENFORCE(platform::is_cpu_place(ctx.GetPlace()), + "This kernel only runs on CPU."); + + auto *x = ctx.Input("X"); + auto *index = ctx.Input("Index"); + auto *output = ctx.Output("Out"); + + output->mutable_data(ctx.GetPlace()); - Y->mutable_data(ctx.GetPlace()); - Gather(ctx.GetPlace(), X, Index, Y); + CPUGather(ctx.device_context(), *x, *index, output); } }; -template +template class GatherGradientOpKernel : public framework::OpKernel { public: void Compute(const framework::ExecutionContext &ctx) const override { + PADDLE_ENFORCE(platform::is_cpu_place(ctx.GetPlace()), + "This kernel only runs on CPU."); + auto *Index = ctx.Input("Index"); auto *dX = ctx.Output(framework::GradVarName("X")); auto *dO = ctx.Input(framework::GradVarName("Out")); dX->mutable_data(ctx.GetPlace()); - ScatterUpdate(ctx.GetPlace(), dO, Index, dX); + auto dxt = framework::EigenVector::Flatten(*dX); + auto place = ctx.GetEigenDevice(); + dxt.device(place) = dxt.constant(static_cast(0)); + + ScatterAssign(ctx.device_context(), *dO, *Index, dX); } }; diff --git a/paddle/operators/gather_test.cc b/paddle/operators/gather_test.cc index 0ae1e99452973feb6d085dd6ef51e2afca988f59..cbd86b87961ee24aa889e208de5ac38e03a33135 100644 --- a/paddle/operators/gather_test.cc +++ b/paddle/operators/gather_test.cc @@ -41,7 +41,9 @@ TEST(Gather, GatherData) { int* p_output = output->mutable_data(make_ddim({2, 4}), CPUPlace()); - Gather(CPUPlace(), src, index, output); + auto* cpu_place = new paddle::platform::CPUPlace(); + paddle::platform::CPUDeviceContext ctx(*cpu_place); + CPUGather(ctx, *src, *index, output); for (int i = 0; i < 4; ++i) EXPECT_EQ(p_output[i], i + 4); for (int i = 4; i < 8; ++i) EXPECT_EQ(p_output[i], i - 4); diff --git a/paddle/operators/math/im2col_test.cc b/paddle/operators/math/im2col_test.cc index f0b8c885918afe7f80edc465c6d9be7c11ac066f..40bdbfe73351a609a4ab9fdc27ac5ff6710df2a2 100644 --- a/paddle/operators/math/im2col_test.cc +++ b/paddle/operators/math/im2col_test.cc @@ -71,7 +71,7 @@ void testIm2col() { context = new paddle::platform::CPUDeviceContext(paddle::platform::CPUPlace()); } else { -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA context = new paddle::platform::CUDADeviceContext(paddle::platform::GPUPlace()); #else @@ -116,7 +116,7 @@ void testIm2col() { TEST(math, im2col) { testIm2col(); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA testIm2col(); #endif } diff --git a/paddle/operators/math/math_function_test.cc b/paddle/operators/math/math_function_test.cc index 22468a0c4a4b0aca343fe766c8c9d63393a338eb..9945ba101d719848aa0c06fa65629d59f167c083 100644 --- a/paddle/operators/math/math_function_test.cc +++ b/paddle/operators/math/math_function_test.cc @@ -1,7 +1,7 @@ #include "paddle/operators/math/math_function.h" #include "gtest/gtest.h" -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA TEST(math_function, notrans_mul_trans) { paddle::framework::Tensor input1; paddle::framework::Tensor input1_gpu; diff --git a/paddle/operators/recurrent_op.cc b/paddle/operators/recurrent_op.cc index 80de229c333f645fb3098b97fa076c6b77bb7ca9..04c4c24951f5db572486ded5edfc26948a821682 100644 --- a/paddle/operators/recurrent_op.cc +++ b/paddle/operators/recurrent_op.cc @@ -30,36 +30,39 @@ using LoDTensor = framework::LoDTensor; void RecurrentAlgorithm::Run(const Scope& scope, const platform::DeviceContext& dev_ctx) const { - auto step_scopes = GetStepScopes(scope); - rnn::SegmentInputs(step_scopes, arg_->inlinks, seq_len_, - false /*infer_shape_mode*/); - InitMemories(step_scopes[0], false /*infer_shape_mode*/); + auto* input0 = scope.FindVar(arg_->inlinks[0]); + PADDLE_ENFORCE_NOT_NULL(input0); + size_t seq_len = input0->GetMutable()->dims()[0]; + PADDLE_ENFORCE_GT(seq_len, 0); - for (size_t step_id = 0; step_id < seq_len_; step_id++) { - // create output alias variables + CreateScopes(scope, seq_len); + auto& step_scopes = GetStepScopes(scope); + rnn::SegmentInputs(step_scopes, arg_->inlinks, seq_len); + InitMemories(step_scopes[0]); + + for (size_t step_id = 0; step_id < seq_len; step_id++) { if (step_id > 0) { - rnn::LinkMemories(step_scopes, arg_->memories, step_id, -1, - false /*infer_shape_mode*/); + rnn::LinkMemories(step_scopes, arg_->memories, step_id, -1); } (*stepnet_)->Run(*step_scopes[step_id], dev_ctx); } - rnn::ConcatOutputs(step_scopes, arg_->outlinks, seq_len_, - false /*infer_shape_mode*/); + rnn::ConcatOutputs(step_scopes, arg_->outlinks, seq_len); } -void RecurrentAlgorithm::CreateScopes(const Scope& scope) const { +void RecurrentAlgorithm::CreateScopes(const Scope& scope, + size_t seq_len) const { // TODO(superjom) Only two scopes are needed for inference, this case will be // supported later. - auto step_scopes_var = scope.FindVar(arg_->step_scopes); + auto* step_scopes_var = scope.FindVar(arg_->step_scopes); PADDLE_ENFORCE(step_scopes_var != nullptr, ""); - auto step_scopes = step_scopes_var->GetMutable>(); + auto* step_scopes = step_scopes_var->GetMutable>(); // Now all variables in scope must be created outside of op. PADDLE_ENFORCE_NOT_NULL(stepnet_); PADDLE_ENFORCE(!(*stepnet_)->Outputs().empty(), "stepnet_ op has no outputs"); - if (seq_len_ > step_scopes->size()) { - for (size_t i = step_scopes->size(); i < seq_len_; ++i) { + if (seq_len > step_scopes->size()) { + for (size_t i = step_scopes->size(); i < seq_len; ++i) { auto& step_scope = scope.NewScope(); // create step net's temp inputs @@ -82,8 +85,7 @@ void RecurrentAlgorithm::CreateScopes(const Scope& scope) const { } } -void RecurrentAlgorithm::InitMemories(Scope* step_scope, - bool infer_shape_mode) const { +void RecurrentAlgorithm::InitMemories(Scope* step_scope) const { for (auto& attr : arg_->memories) { auto* pre_mem = step_scope->NewVar(attr.pre_var)->GetMutable(); PADDLE_ENFORCE(step_scope->FindVar(attr.boot_var) != nullptr, @@ -91,12 +93,9 @@ void RecurrentAlgorithm::InitMemories(Scope* step_scope, attr.boot_var); auto* boot_mem = step_scope->FindVar(attr.boot_var)->GetMutable(); - if (infer_shape_mode) { - pre_mem->Resize(boot_mem->dims()); - PADDLE_ENFORCE_EQ(pre_mem->dims().size(), 2); - } else { - pre_mem->ShareDataWith(*boot_mem); - } + pre_mem->Resize(boot_mem->dims()); + PADDLE_ENFORCE_EQ(pre_mem->dims().size(), 2); + pre_mem->ShareDataWith(*boot_mem); } } @@ -146,23 +145,23 @@ class RecurrentAlgorithmProtoAndCheckerMaker void RecurrentGradientAlgorithm::Run( const Scope& scope, const platform::DeviceContext& dev_ctx) const { - auto step_scopes = GetStepScopes(scope); - rnn::SegmentInputs(step_scopes, arg_->inlinks, seq_len_, - false /*infer_shape_mode*/); - for (int step_id = seq_len_ - 1; step_id >= 0; --step_id) { - if (static_cast(step_id) != seq_len_ - 1) { - rnn::LinkMemories(step_scopes, arg_->memories, step_id, 1, - false /*infer_shape_mode*/); + auto* input0 = scope.FindVar(arg_->inlinks[0]); + PADDLE_ENFORCE_NOT_NULL(input0); + size_t seq_len = input0->GetMutable()->dims()[0]; + auto& step_scopes = GetStepScopes(scope); + rnn::SegmentInputs(step_scopes, arg_->inlinks, seq_len); + for (int step_id = seq_len - 1; step_id >= 0; --step_id) { + if (step_id != seq_len - 1) { + rnn::LinkMemories(step_scopes, arg_->memories, step_id, 1); } (*stepnet_)->Run(*step_scopes[step_id], dev_ctx); } - LinkBootMemoryGradients(step_scopes[0], false); - rnn::ConcatOutputs(step_scopes, arg_->outlinks, seq_len_, - false /*infer_shape_mode*/); + rnn::ConcatOutputs(step_scopes, arg_->outlinks, seq_len); + LinkBootMemoryGradients(step_scopes[0]); } void RecurrentGradientAlgorithm::LinkBootMemoryGradients( - Scope* step_scope, bool infer_shape_mode) const { + Scope* step_scope) const { for (auto& attr : arg_->memories) { PADDLE_ENFORCE(step_scope->FindVar(attr.var) != nullptr, "memory variable [%s] does not exists", attr.var); @@ -171,11 +170,8 @@ void RecurrentGradientAlgorithm::LinkBootMemoryGradients( auto* mem_grad = step_scope->NewVar(attr.var)->GetMutable(); auto* boot_mem_grad = step_scope->NewVar(attr.boot_var)->GetMutable(); - if (infer_shape_mode) { - boot_mem_grad->Resize(mem_grad->dims()); - } else { - boot_mem_grad->ShareDataWith(*mem_grad); - } + boot_mem_grad->Resize(mem_grad->dims()); + boot_mem_grad->ShareDataWith(*mem_grad); } } diff --git a/paddle/operators/recurrent_op.h b/paddle/operators/recurrent_op.h index c6b9a5533eece9057449b5c875ddcb3cefe716f0..253d7e3284360ceaddce9ef5f8f9a3ea4793d740 100644 --- a/paddle/operators/recurrent_op.h +++ b/paddle/operators/recurrent_op.h @@ -48,7 +48,7 @@ class RecurrentAlgorithm { * NOTE the scopes are reused in both the forward and backward, so just * create once and expand its size if more steps need. */ - void CreateScopes(const framework::Scope& scope) const; + void CreateScopes(const framework::Scope& scope, size_t seq_len) const; const std::vector& GetStepScopes( const framework::Scope& scope) const { @@ -56,12 +56,11 @@ class RecurrentAlgorithm { ->GetMutable>(); } - void InitMemories(framework::Scope* step_scopes, bool infer_shape_mode) const; + void InitMemories(framework::Scope* step_scopes) const; private: std::unique_ptr* stepnet_; rnn::Argument* arg_; - mutable size_t seq_len_; }; class RecurrentGradientAlgorithm { @@ -86,8 +85,7 @@ class RecurrentGradientAlgorithm { void Run(const framework::Scope& scope, const platform::DeviceContext& dev_ctx) const; - void LinkBootMemoryGradients(framework::Scope* step_scopes, - bool infer_shape_mode) const; + void LinkBootMemoryGradients(framework::Scope* step_scopes) const; protected: inline const std::vector& GetStepScopes( @@ -98,7 +96,6 @@ class RecurrentGradientAlgorithm { private: rnn::Argument* arg_; - mutable size_t seq_len_; std::unique_ptr* stepnet_; }; @@ -123,6 +120,7 @@ class RecurrentOp : public framework::OperatorBase { void set_stepnet(std::unique_ptr net) { stepnet_ = std::move(net); } + const OperatorBase& stepnet() const { return *stepnet_; } static const rnn::ArgumentName kArgName; diff --git a/paddle/operators/rnn/recurrent_op_utils.cc b/paddle/operators/rnn/recurrent_op_utils.cc index a767009d2366e20d2ebd35f562b8df7d408f2d4e..ef317a71f12c6de974bd8715bb08122b761fae37 100644 --- a/paddle/operators/rnn/recurrent_op_utils.cc +++ b/paddle/operators/rnn/recurrent_op_utils.cc @@ -25,7 +25,7 @@ using LoDTensor = framework::LoDTensor; void SegmentInputs(const std::vector& step_scopes, const std::vector& inlinks, - const size_t seq_len, bool infer_shape_mode) { + const size_t seq_len) { PADDLE_ENFORCE(!inlinks.empty(), "no in links are provided."); for (size_t i = 0; i < inlinks.size(); ++i) { // global inputs @@ -41,11 +41,9 @@ void SegmentInputs(const std::vector& step_scopes, for (size_t j = 0; j < seq_len; j++) { Tensor* step_input = step_scopes[j]->NewVar(inlinks[i])->GetMutable(); - if (!infer_shape_mode) { - // The input of operators of each step is Tensor here. - // Maybe need to modify Slice function. - *step_input = input->Slice(j, j + 1); - } + // The input of operators of each step is Tensor here. + // Maybe need to modify Slice function. + *step_input = input->Slice(j, j + 1); step_input->Resize(step_dims); } } @@ -53,39 +51,35 @@ void SegmentInputs(const std::vector& step_scopes, void ConcatOutputs(const std::vector& step_scopes, const std::vector& outlinks, - const size_t seq_len, bool infer_shape_mode) { + const size_t seq_len) { for (size_t i = 0; i < outlinks.size(); i++) { - auto output_var = step_scopes[0]->parent().FindVar(outlinks[i]); + auto* output_var = step_scopes[0]->parent().FindVar(outlinks[i]); PADDLE_ENFORCE_NOT_NULL(output_var, "output link [%s] is not in scope.", outlinks[i]); LoDTensor* output = output_var->GetMutable(); - if (infer_shape_mode) { - auto step_scope_var = step_scopes[0]->FindVar(outlinks[i]); - PADDLE_ENFORCE_NOT_NULL(step_scope_var, "%s not in scope", outlinks[i]); - f::DDim step_dims = - step_scope_var->template GetMutable()->dims(); - std::vector dims_vec = vectorize(step_dims); - dims_vec.insert(dims_vec.begin(), seq_len); - output->Resize(f::make_ddim(dims_vec)); - } else { - output->mutable_data(platform::CPUPlace()); - for (size_t j = 0; j < seq_len; j++) { - LoDTensor* step_output = - step_scopes[j]->FindVar(outlinks[i])->GetMutable(); - // TODO(luotao02) data type and platform::DeviceContext() should set - // correctly - (output->Slice(j, j + 1)) - .CopyFrom(*step_output, platform::CPUPlace()); - } + auto* step_scope_var = step_scopes[0]->FindVar(outlinks[i]); + PADDLE_ENFORCE_NOT_NULL(step_scope_var, "%s not in scope", outlinks[i]); + f::DDim step_dims = + step_scope_var->template GetMutable()->dims(); + std::vector dims_vec = vectorize(step_dims); + dims_vec.insert(dims_vec.begin(), seq_len); + output->Resize(f::make_ddim(dims_vec)); + output->mutable_data(platform::CPUPlace()); + for (size_t j = 0; j < seq_len; j++) { + LoDTensor* step_output = + step_scopes[j]->FindVar(outlinks[i])->GetMutable(); + // TODO(luotao02) data type and platform::DeviceContext() should set + // correctly + (output->Slice(j, j + 1)) + .CopyFrom(*step_output, platform::CPUPlace()); } } } void LinkMemories(const std::vector& scopes, const std::vector& memories, - const size_t step_id, const int offset, - bool infer_shape_mode) { + const size_t step_id, const int offset) { PADDLE_ENFORCE_LT(step_id, scopes.size(), "step [%d] is out of range of step scopes' size [%d]", step_id, scopes.size()); @@ -95,16 +89,13 @@ void LinkMemories(const std::vector& scopes, step_id + offset, scopes.size(), "offset [%d] is out of range, it must be less than (%d - %d)", offset, scopes.size(), step_id); - auto scope = scopes[step_id]; - auto linked_scope = scopes[step_id + offset]; + auto* scope = scopes[step_id]; + auto* linked_scope = scopes[step_id + offset]; for (auto& attr : memories) { - auto mem = scope->FindVar(attr.pre_var)->GetMutable(); - auto linked_mem = linked_scope->FindVar(attr.var)->GetMutable(); - if (infer_shape_mode) { - mem->Resize(linked_mem->dims()); - } else { - mem->ShareDataWith(*linked_mem); - } + auto* mem = scope->FindVar(attr.pre_var)->GetMutable(); + auto* linked_mem = linked_scope->FindVar(attr.var)->GetMutable(); + mem->Resize(linked_mem->dims()); + mem->ShareDataWith(*linked_mem); } } @@ -115,11 +106,11 @@ void InitArgument(const ArgumentName& name, Argument* arg, arg->inlinks = op.Inputs(name.inlinks); arg->outlinks = op.Outputs(name.outlinks); - auto boot_memories = + auto& boot_memories = is_grad ? op.Outputs(name.boot_memories) : op.Inputs(name.boot_memories); // attributes - auto memories = op.Attr>(name.memories); - auto pre_memories = op.Attr>(name.pre_memories); + auto& memories = op.Attr>(name.memories); + auto& pre_memories = op.Attr>(name.pre_memories); PADDLE_ENFORCE(memories.size() == boot_memories.size(), "the size of memories, boot_memories don't match:%d,%d", diff --git a/paddle/operators/rnn/recurrent_op_utils.h b/paddle/operators/rnn/recurrent_op_utils.h index 9c777f1e9067a3e2ceb9d23f7bf7d3c73343c91f..fd17b9b88915cf458ff2836b5c5d8f84cd9b65b5 100644 --- a/paddle/operators/rnn/recurrent_op_utils.h +++ b/paddle/operators/rnn/recurrent_op_utils.h @@ -64,18 +64,18 @@ struct ArgumentName { */ void SegmentInputs(const std::vector& step_scopes, const std::vector& inlinks, - const size_t seq_len, bool infer_shape_mode); + const size_t seq_len); /** * Process outputs of step nets and merge to variables. */ void ConcatOutputs(const std::vector& step_scopes, const std::vector& outlinks, - const size_t seq_len, bool infer_shape_mode); + const size_t seq_len); void LinkMemories(const std::vector& step_scopes, const std::vector& memories, const size_t step_id, - const int offset, bool infer_shape_mode); + const int offset); void InitArgument(const ArgumentName& name, Argument* arg, const framework::OperatorBase& op, bool is_grad = false); diff --git a/paddle/operators/scatter.cu.h b/paddle/operators/scatter.cu.h new file mode 100644 index 0000000000000000000000000000000000000000..d95436be4f25b9df4aaef57ddb249ecf944f0666 --- /dev/null +++ b/paddle/operators/scatter.cu.h @@ -0,0 +1,80 @@ +/* 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/tensor.h" +#include "paddle/platform/place.h" + +namespace paddle { +namespace operators { + +using Tensor = framework::Tensor; + +#define CUDA_1D_KERNEL_LOOP(i, n) \ + for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < (n); \ + i += blockDim.x * gridDim.x) + +template +__global__ void ScatterCUDAKernel(const T* params, const int* indices, + T* output, size_t index_size, + size_t slice_size) { + CUDA_1D_KERNEL_LOOP(i, index_size * slice_size) { + int indices_i = i / slice_size; + int slice_i = i - indices_i * slice_size; // offset inside the slice + int scatter_i = indices[indices_i]; + int out_i = scatter_i * slice_size + slice_i; + *(output + out_i) = *(params + i); + } +} + +/** + * A thin wrapper on gpu tensor + * Return a new updated tensor from source tensor, scatter-assigned according to + * index + * input[src]: type-T source Tensor + * input[index]: type-int index Tensor (1-D) + * return: output tensor + */ +template +void GPUScatterAssign(const platform::DeviceContext& ctx, const Tensor& src, + const Tensor& index, Tensor* output) { + // PADDLE_ENFORCE(platform::is_gpu_place(place)); + // check index of shape 1-D + PADDLE_ENFORCE(index.dims().size() == 1); + int index_size = index.dims()[0]; + + auto src_dims = src.dims(); + framework::DDim output_dims(src_dims); + output_dims[0] = index_size; + + // slice size + int slice_size = 1; + for (int i = 1; i < src_dims.size(); ++i) slice_size *= src_dims[i]; + + const T* p_src = src.data(); + const int* p_index = index.data(); + T* p_output = output->data(); + + int block = 512; + int n = slice_size * index_size; + int grid = (n + block - 1) / block; + + ScatterCUDAKernel<<< + grid, block, 0, + reinterpret_cast(ctx).stream()>>>( + p_src, p_index, p_output, index_size, slice_size); +} + +} // namespace operators +} // namespace paddle diff --git a/paddle/operators/scatter.h b/paddle/operators/scatter.h index fec1046a9748637ac41aa35143a8c2cf5528913f..c1fb844ebd2ff7ca7dbdb8e8ac3c1fff4c0c6607 100644 --- a/paddle/operators/scatter.h +++ b/paddle/operators/scatter.h @@ -24,63 +24,42 @@ namespace paddle { namespace operators { using Tensor = framework::Tensor; -template -using EigenVector = framework::EigenVector; - -// Implementation of CPU copy -template -void CPUScatterUpdate(const paddle::framework::Tensor* src, const int* index, - const size_t index_size, - paddle::framework::Tensor* output) { - paddle::framework::DDim output_dims = output->dims(); - - for (size_t i = 0; i < index_size; ++i) { - int index_ = index[i]; - - paddle::framework::Tensor src_ = *src; - paddle::framework::Tensor output_ = *output; - if (index_size > 1) src_ = src->Slice(i, i + 1); - if (output_dims[0] > 1) output_ = output->Slice(index_, index_ + 1); - - auto X = EigenVector::Flatten(src_); - auto Y = EigenVector::Flatten(output_); - - Y = X + Y; - } -} - -// Implementation of GPU scatter: -template -void GPUScatterUpdate(const T* src, const int* index, const int slice_size, - const int index_size, T* output); /** * Return a updated tensor from source tensor, scattered according to index: - * dst[i] += src[index[i]] + * dst[i] = src[index[i]] * input[src]: type-T source Tensor * input[index]: type-int index Tensor (1-D) * return: output tensor */ template -void ScatterUpdate(const platform::Place& place, - const paddle::framework::Tensor* src, - const paddle::framework::Tensor* index, - paddle::framework::Tensor* output) { +void ScatterAssign(const platform::DeviceContext& ctx, const Tensor& src, + const Tensor& index, Tensor* output) { + PADDLE_ENFORCE(platform::is_cpu_place(ctx.GetPlace())); // check index of shape 1-D - PADDLE_ENFORCE(index->dims().size() == 1); - int index_size = index->dims()[0]; + PADDLE_ENFORCE(index.dims().size() == 1); + int index_size = index.dims()[0]; - auto src_dims = src->dims(); + auto src_dims = src.dims(); auto dst_dims = output->dims(); + const T* p_src = src.data(); + const int* p_index = index.data(); + T* p_output = output->data(); + // check src shape and dst shape should match for (int i = 1; i < src_dims.size(); i++) PADDLE_ENFORCE(src_dims[i] == dst_dims[i]); - if (platform::is_cpu_place(place)) { - CPUScatterUpdate(src, index->data(), index_size, output); - } else { + // slice size + size_t slice_size = 1; + for (int i = 1; i < src_dims.size(); ++i) slice_size *= src_dims[i]; + + const size_t slice_bytes = slice_size * sizeof(T); + + for (int i = 0; i < index_size; ++i) { + int index_ = p_index[i]; + memcpy(p_output + index_ * slice_size, p_src + i * slice_size, slice_bytes); } } diff --git a/paddle/operators/scatter_op.cc b/paddle/operators/scatter_op.cc index cadd8841b6ab3a3674054240265eb6d4b474db1e..d15ba151539987c133ac57e102df53551483c6dd 100644 --- a/paddle/operators/scatter_op.cc +++ b/paddle/operators/scatter_op.cc @@ -97,8 +97,5 @@ Out[Index] = Ref[Index] + Updates namespace ops = paddle::operators; REGISTER_OP(scatter, ops::ScatterOp, ops::ScatterOpMaker, scatter_grad, ops::ScatterGradOp); -REGISTER_OP_CPU_KERNEL(scatter, - ops::ScatterOpKernel); -REGISTER_OP_CPU_KERNEL( - scatter_grad, - ops::ScatterGradientOpKernel); +REGISTER_OP_CPU_KERNEL(scatter, ops::ScatterOpKernel); +REGISTER_OP_CPU_KERNEL(scatter_grad, ops::ScatterGradientOpKernel); diff --git a/paddle/operators/scatter_op.cu b/paddle/operators/scatter_op.cu new file mode 100644 index 0000000000000000000000000000000000000000..06f4d759447b6dcd28b50576dfc246fc466d9336 --- /dev/null +++ b/paddle/operators/scatter_op.cu @@ -0,0 +1,63 @@ +/* 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 "gather.cu.h" +#include "paddle/operators/gather_op.h" +#include "scatter.cu.h" + +namespace paddle { +namespace operators { + +template +class ScatterOpCUDAKernel : public framework::OpKernel { + public: + void Compute(const framework::ExecutionContext &ctx) const override { + PADDLE_ENFORCE(platform::is_gpu_place(ctx.GetPlace()), + "This kernel only runs on GPU device."); + auto *Ref = ctx.Input("Ref"); + auto *Index = ctx.Input("Index"); + auto *Updates = ctx.Input("Updates"); + auto *Out = ctx.Output("Out"); + + Out->ShareDataWith(*Ref); + + GPUScatterAssign(ctx.device_context(), *Updates, *Index, Out); + } +}; + +template +class ScatterGradOpCUDAKernel : public framework::OpKernel { + public: + void Compute(const framework::ExecutionContext &ctx) const override { + PADDLE_ENFORCE(platform::is_gpu_place(ctx.GetPlace()), + "This kernel only runs on GPU device."); + auto *dRef = ctx.Output(framework::GradVarName("Ref")); + auto *dUpdates = ctx.Output(framework::GradVarName("Updates")); + auto *Index = ctx.Input("Index"); + auto *dOut = ctx.Input(framework::GradVarName("Out")); + + // In place gradient: dRef = dO + dRef->ShareDataWith(*dOut); + dUpdates->mutable_data(ctx.GetPlace()); + // Gradient by Gather: dUpdates = dO[Index] + GPUGather(ctx.device_context(), *dOut, *Index, dUpdates); + } +}; + +} // namespace operators +} // namespace paddle + +namespace ops = paddle::operators; +REGISTER_OP_GPU_KERNEL(scatter, ops::ScatterOpCUDAKernel); +REGISTER_OP_GPU_KERNEL(scatter_grad, ops::ScatterGradOpCUDAKernel); diff --git a/paddle/operators/scatter_op.h b/paddle/operators/scatter_op.h index a8eb54399a932913de208e1ddc90a6ff0dfaa452..6101219006414e4865f676e3ca5d2a88949ad17a 100644 --- a/paddle/operators/scatter_op.h +++ b/paddle/operators/scatter_op.h @@ -23,10 +23,12 @@ namespace operators { using Tensor = framework::Tensor; -template +template class ScatterOpKernel : public framework::OpKernel { public: void Compute(const framework::ExecutionContext &ctx) const override { + PADDLE_ENFORCE(platform::is_cpu_place(ctx.GetPlace()), + "This kernel only runs on CPU."); auto *Ref = ctx.Input("Ref"); auto *Index = ctx.Input("Index"); auto *Updates = ctx.Input("Updates"); @@ -35,14 +37,16 @@ class ScatterOpKernel : public framework::OpKernel { // In place output: Out = Ref, Out[Index] += Updates Out->ShareDataWith(*Ref); // Apply ScatterUpdate: Out[index] += Updates[:] - ScatterUpdate(ctx.GetPlace(), Updates, Index, Out); + ScatterAssign(ctx.device_context(), *Updates, *Index, Out); } }; -template +template class ScatterGradientOpKernel : public framework::OpKernel { public: void Compute(const framework::ExecutionContext &ctx) const override { + PADDLE_ENFORCE(platform::is_cpu_place(ctx.GetPlace()), + "This kernel only runs on CPU."); auto *dRef = ctx.Output(framework::GradVarName("Ref")); auto *dUpdates = ctx.Output(framework::GradVarName("Updates")); auto *Index = ctx.Input("Index"); @@ -52,7 +56,7 @@ class ScatterGradientOpKernel : public framework::OpKernel { dRef->ShareDataWith(*dOut); dUpdates->mutable_data(ctx.GetPlace()); // Gradient by Gather: dUpdates += dO[Index] - Gather(ctx.GetPlace(), dOut, Index, dUpdates); + CPUGather(ctx.device_context(), *dOut, *Index, dUpdates); } }; diff --git a/paddle/operators/scatter_test.cc b/paddle/operators/scatter_test.cc index 26fdaff1460a297fa638181641991f732533fe52..00dbdacbfef7af826790472acc6caa285c259e0e 100644 --- a/paddle/operators/scatter_test.cc +++ b/paddle/operators/scatter_test.cc @@ -40,7 +40,9 @@ TEST(scatter, ScatterUpdate) { float* p_output = output->mutable_data(make_ddim({4, 4}), CPUPlace()); - ScatterUpdate(CPUPlace(), src, index, output); + auto* cpu_place = new paddle::platform::CPUPlace(); + paddle::platform::CPUDeviceContext ctx(*cpu_place); + ScatterAssign(ctx, *src, *index, output); for (size_t i = 0; i < 4; ++i) EXPECT_EQ(p_output[i], float(0)); for (size_t i = 0; i < 4; ++i) EXPECT_EQ(output->data()[i], float(0)); diff --git a/paddle/operators/sgd_op.cc b/paddle/operators/sgd_op.cc index 3bce95535cf10c0df95b503c6e362b3f0ba2e723..31d491f130e36289f7dd3dc18710d35d5ba6ab34 100644 --- a/paddle/operators/sgd_op.cc +++ b/paddle/operators/sgd_op.cc @@ -23,17 +23,22 @@ class SGDOp : public framework::OperatorWithKernel { protected: void InferShape(framework::InferShapeContextBase *ctx) const override { - PADDLE_ENFORCE(ctx->HasInput("param"), - "Input(param) of SGDOp should not be null."); - PADDLE_ENFORCE(ctx->HasInput("grad"), - "Input(grad) of SGDOp should not be null."); - PADDLE_ENFORCE(ctx->HasOutput("param_out"), - "Output(param_out) of SGDOp should not be null."); - - auto param_dim = ctx->GetInputDim("param"); - PADDLE_ENFORCE_EQ(param_dim, ctx->GetInputDim("grad"), + PADDLE_ENFORCE(ctx->HasInput("Param"), + "Input(Param) of SGDOp should not be null."); + PADDLE_ENFORCE(ctx->HasInput("Grad"), + "Input(Grad) of SGDOp should not be null."); + PADDLE_ENFORCE(ctx->HasInput("LearningRate"), + "Input(LearningRate) of SGDOp should not be null."); + PADDLE_ENFORCE(ctx->HasOutput("ParamOut"), + "Output(ParamOut) of SGDOp should not be null."); + + auto lr_dims = ctx->GetInputDim("LearningRate"); + PADDLE_ENFORCE_EQ(framework::product(lr_dims), 1, + "Learning rate should have 1 element"); + auto param_dim = ctx->GetInputDim("Param"); + PADDLE_ENFORCE_EQ(param_dim, ctx->GetInputDim("Grad"), "Two input of SGD Op's dimension must be same."); - ctx->SetOutputDim("param_out", param_dim); + ctx->SetOutputDim("ParamOut", param_dim); } }; @@ -41,10 +46,10 @@ class SGDOpMaker : public framework::OpProtoAndCheckerMaker { public: SGDOpMaker(framework::OpProto *proto, framework::OpAttrChecker *op_checker) : OpProtoAndCheckerMaker(proto, op_checker) { - AddInput("param", "input parameter"); - AddInput("grad", "input gradient"); - AddOutput("param_out", "output parameter"); - AddAttr("learning_rate", "learning rate of sgd"); + AddInput("Param", "Input parameter"); + AddInput("LearningRate", "Learning rate of SGD"); + AddInput("Grad", "Input gradient"); + AddOutput("ParamOut", "output parameter"); AddComment(R"DOC( Simplest sgd algorithm. diff --git a/paddle/operators/sgd_op.h b/paddle/operators/sgd_op.h index a3fe3308942f98e2c28376b589b6fc930e6878a1..26f4012f258771794c736dbfad4af174b017f410 100644 --- a/paddle/operators/sgd_op.h +++ b/paddle/operators/sgd_op.h @@ -19,28 +19,25 @@ limitations under the License. */ namespace paddle { namespace operators { -using Tensor = framework::Tensor; -template -using EigenVector = framework::EigenVector; - template class SGDOpKernel : public framework::OpKernel { public: void Compute(const framework::ExecutionContext& ctx) const override { - auto param = ctx.Input("param"); - auto grad = ctx.Input("grad"); - auto param_out = ctx.Output("param_out"); - float lr = ctx.Attr("learning_rate"); + auto param = ctx.Input("Param"); + auto grad = ctx.Input("Grad"); + auto param_out = ctx.Output("ParamOut"); + auto learning_rate = ctx.Input("LearningRate"); param_out->mutable_data(ctx.GetPlace()); - auto p = EigenVector::Flatten(*param); - auto g = EigenVector::Flatten(*grad); - auto o = EigenVector::Flatten(*param_out); + auto p = framework::EigenVector::Flatten(*param); + auto g = framework::EigenVector::Flatten(*grad); + auto o = framework::EigenVector::Flatten(*param_out); + auto lr = framework::EigenVector::Flatten(*learning_rate); auto place = ctx.GetEigenDevice(); - o.device(place) = p - lr * g; + Eigen::DSizes grad_dsize(grad->numel()); + o.device(place) = p - lr.broadcast(grad_dsize) * g; } }; diff --git a/paddle/operators/strided_memcpy_test.cc b/paddle/operators/strided_memcpy_test.cc index 05882a88738cfc9cc23480efe0afe504008377ca..68f064eaee5851333ddf9767b7138da83a28503d 100644 --- a/paddle/operators/strided_memcpy_test.cc +++ b/paddle/operators/strided_memcpy_test.cc @@ -72,7 +72,7 @@ TEST(StridedMemcpy, CPUConcat) { } } -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA TEST(StridedMemcpy, GPUCrop) { // clang-format off int src[] = { @@ -157,4 +157,4 @@ TEST(StridedMemcpy, GPUConcat) { #endif } // namespace operators -} // namespace paddle \ No newline at end of file +} // namespace paddle diff --git a/paddle/operators/sum_op.cc b/paddle/operators/sum_op.cc index 5d76313aeb96c0c8204f64aee1057f753ec85d6b..7c422b477083fba4661aeb427422abb623b172bb 100644 --- a/paddle/operators/sum_op.cc +++ b/paddle/operators/sum_op.cc @@ -11,6 +11,7 @@ limitations under the License. */ #include "paddle/operators/sum_op.h" #include +#include "paddle/operators/net_op.h" namespace paddle { namespace operators { @@ -22,14 +23,15 @@ class SumOp : public framework::OperatorWithKernel { protected: void InferShape(framework::InferShapeContextBase* ctx) const override { + PADDLE_ENFORCE(ctx->HasInputs("X"), "Inputs(X) should not be null"); auto x_dims = ctx->GetInputsDim("X"); - PADDLE_ENFORCE(!x_dims.empty(), "Input(X) of SumOp should not be null."); PADDLE_ENFORCE(ctx->HasOutput("Out"), "Output(Out) of SumOp should not be null."); - auto in_dim = x_dims[0]; size_t N = x_dims.size(); PADDLE_ENFORCE_GT(N, 1, "Input tensors count should > 1."); + + auto in_dim = x_dims[0]; for (size_t i = 1; i < N; i++) { auto dim = x_dims[i]; PADDLE_ENFORCE(in_dim == dim, "Input tensors must have same shape"); @@ -56,21 +58,23 @@ or not. But the output only shares the LoD with the first input. } }; -class SumGradOp : public framework::OperatorWithKernel { +class SumGradOp : public NetOp { public: - using framework::OperatorWithKernel::OperatorWithKernel; + SumGradOp(const std::string& type, const framework::VariableNameMap& inputs, + const framework::VariableNameMap& outputs, + const framework::AttributeMap& attrs) + : NetOp(type, inputs, outputs, attrs) { + auto& x_grad_names = Outputs(framework::GradVarName("X")); + auto out_grad_name = this->Input(framework::GradVarName("Out")); - protected: - void InferShape(framework::InferShapeContextBase* ctx) const override { - auto out_grad_dims = ctx->GetInputDim(framework::GradVarName("Out")); - auto x_grad_names = ctx->Outputs(framework::GradVarName("X")); - size_t x_length = x_grad_names.size(); - std::vector x_grad_dims; - x_grad_dims.reserve(x_length); - for (size_t i = 0; i < x_length; ++i) { - x_grad_dims.push_back(out_grad_dims); + framework::AttributeMap grad_attrs; + grad_attrs["scale"] = 1.0f; + for (auto& x_grad_name : x_grad_names) { + AppendOp(framework::OpRegistry::CreateOp( + "scale", {{"X", {out_grad_name}}}, {{"Out", {x_grad_name}}}, + grad_attrs)); } - ctx->SetOutputsDim(framework::GradVarName("X"), x_grad_dims); + CompleteAddOp(false); } }; @@ -80,5 +84,3 @@ class SumGradOp : public framework::OperatorWithKernel { namespace ops = paddle::operators; REGISTER_OP(sum, ops::SumOp, ops::SumOpMaker, sum_grad, ops::SumGradOp); REGISTER_OP_CPU_KERNEL(sum, ops::SumKernel); -REGISTER_OP_CPU_KERNEL(sum_grad, - ops::SumGradKernel); diff --git a/paddle/operators/sum_op.cu b/paddle/operators/sum_op.cu index a465cf3659ba7c51338abadfc62962fb6755a39d..b1896d3cd87f47bd2573287ee37b1b72ae9ec6e8 100644 --- a/paddle/operators/sum_op.cu +++ b/paddle/operators/sum_op.cu @@ -14,5 +14,3 @@ limitations under the License. */ namespace ops = paddle::operators; REGISTER_OP_GPU_KERNEL(sum, ops::SumKernel); -REGISTER_OP_GPU_KERNEL(sum_grad, - ops::SumGradKernel); diff --git a/paddle/operators/sum_op.h b/paddle/operators/sum_op.h index 7e8fbb9e41c694df9169ea583ce47c33d3bcf2bb..91e5da8b40d452db8715990cdbe2731b3aea44b9 100644 --- a/paddle/operators/sum_op.h +++ b/paddle/operators/sum_op.h @@ -42,24 +42,5 @@ class SumKernel : public framework::OpKernel { } }; -template -class SumGradKernel : public framework::OpKernel { - public: - void Compute(const framework::ExecutionContext& context) const override { - auto* input = context.Input(framework::GradVarName("Out")); - auto outs = context.MultiOutput(framework::GradVarName("X")); - for (auto out : outs) { - out->mutable_data(context.GetPlace()); - } - - auto place = context.GetEigenDevice(); - auto in = EigenVector::Flatten(*input); - for (auto out : outs) { - auto result = EigenVector::Flatten(*out); - result.device(place) = in; - } - } -}; - } // namespace operators } // namespace paddle diff --git a/paddle/platform/device_context.cc b/paddle/platform/device_context.cc index 36af1ac677f6bb3e5b6392ff0de678afe7e47950..a9b6b799036a4f2ba93ef52398131db4fcb599f5 100644 --- a/paddle/platform/device_context.cc +++ b/paddle/platform/device_context.cc @@ -35,7 +35,7 @@ Eigen::DefaultDevice* CPUDeviceContext::eigen_device() const { Place CPUDeviceContext::GetPlace() const { return CPUPlace(); } -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA template <> Eigen::GpuDevice* diff --git a/paddle/platform/device_context.h b/paddle/platform/device_context.h index d805d2ab085f76e119edf1c6f2acb9715883d755..ef5f19214d9ccb23b9c946bee28cb764122bd7cd 100644 --- a/paddle/platform/device_context.h +++ b/paddle/platform/device_context.h @@ -14,7 +14,7 @@ limitations under the License. */ #include "paddle/platform/enforce.h" #include "paddle/platform/place.h" -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA #include "paddle/platform/dynload/cublas.h" #include "paddle/platform/dynload/cudnn.h" #include "paddle/platform/gpu_info.h" @@ -61,7 +61,7 @@ class CPUDeviceContext : public DeviceContext { std::unique_ptr eigen_device_; }; -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA template <> struct EigenDeviceConverter { using EigenDeviceType = Eigen::GpuDevice; diff --git a/paddle/platform/device_context_test.cc b/paddle/platform/device_context_test.cc index f4b00c57dee5196e535816d8985fd7e831c4c226..8bf5174c4a5579f6f5602dd38e5a87ed3ef444a7 100644 --- a/paddle/platform/device_context_test.cc +++ b/paddle/platform/device_context_test.cc @@ -20,7 +20,7 @@ TEST(Device, Init) { using paddle::platform::CUDADeviceContext; using paddle::platform::GPUPlace; - int count = paddle::platform::GetDeviceCount(); + int count = paddle::platform::GetCUDADeviceCount(); for (int i = 0; i < count; i++) { DeviceContext* device_context = new CUDADeviceContext(GPUPlace(i)); Eigen::GpuDevice* gpu_device = @@ -34,7 +34,7 @@ TEST(Device, CUDADeviceContext) { using paddle::platform::CUDADeviceContext; using paddle::platform::GPUPlace; - int count = paddle::platform::GetDeviceCount(); + int count = paddle::platform::GetCUDADeviceCount(); for (int i = 0; i < count; i++) { CUDADeviceContext* device_context = new CUDADeviceContext(GPUPlace(i)); Eigen::GpuDevice* gpu_device = device_context->eigen_device(); diff --git a/paddle/platform/enforce.h b/paddle/platform/enforce.h index 52bd23039b82fc0290edb640f5d5851ba9abee61..15d8446cd8dceb2fdc03536e1f7bbcde73403a23 100644 --- a/paddle/platform/enforce.h +++ b/paddle/platform/enforce.h @@ -29,7 +29,7 @@ limitations under the License. */ #include // for __cxa_demangle #endif -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA #include "paddle/platform/dynload/cublas.h" #include "paddle/platform/dynload/cudnn.h" @@ -113,7 +113,7 @@ inline typename std::enable_if::type throw_on_error( } } -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA template inline typename std::enable_if::type throw_on_error( diff --git a/paddle/platform/enforce_test.cc b/paddle/platform/enforce_test.cc index 80bdee3d9dfbe38ef707a6ba60cdb7f7b99714de..8206a055eabf4abf584962e921610d5029e2f571 100644 --- a/paddle/platform/enforce_test.cc +++ b/paddle/platform/enforce_test.cc @@ -213,4 +213,4 @@ TEST(ENFORCE_USER_DEFINED_CLASS, EQ) { TEST(ENFORCE_USER_DEFINED_CLASS, NE) { Dims a{{1, 2, 3, 4}}, b{{5, 6, 7, 8}}; ASSERT_THROW(PADDLE_ENFORCE_EQ(a, b), paddle::platform::EnforceNotMet); -} \ No newline at end of file +} diff --git a/paddle/platform/gpu_info.cc b/paddle/platform/gpu_info.cc index be381a4e26cf0eb41f5b3de88bd03ad8901683cc..70ad611d5dd61937e6bf7d980e34b5c9023977b2 100644 --- a/paddle/platform/gpu_info.cc +++ b/paddle/platform/gpu_info.cc @@ -26,11 +26,11 @@ DEFINE_double(fraction_of_gpu_memory_to_use, 0.95, namespace paddle { namespace platform { -int GetDeviceCount() { +int GetCUDADeviceCount() { int count; PADDLE_ENFORCE( cudaGetDeviceCount(&count), - "cudaGetDeviceCount failed in paddle::platform::GetDeviceCount"); + "cudaGetDeviceCount failed in paddle::platform::GetCUDADeviceCount"); return count; } diff --git a/paddle/platform/gpu_info.h b/paddle/platform/gpu_info.h index f0c825bd9b0bc41396b8fdb95f0b4337cbe3db02..fb33db07bd54d37dec2e5d687ecefb01cc330e44 100644 --- a/paddle/platform/gpu_info.h +++ b/paddle/platform/gpu_info.h @@ -14,7 +14,7 @@ limitations under the License. */ #pragma once -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA #include #include @@ -28,7 +28,7 @@ const std::string kEnvFractionGpuMemoryToUse = "PADDLE_FRACTION_GPU_MEMORY_TO_USE"; //! Get the total number of GPU devices in system. -int GetDeviceCount(); +int GetCUDADeviceCount(); //! Get the current GPU device id in system. int GetCurrentDeviceId(); diff --git a/paddle/platform/macros.h b/paddle/platform/macros.h index 4a04a38c0c6f905639004dea2f4416ecc57c8620..feae7bdd77e3a0d02f33fb33991648408f542d0e 100644 --- a/paddle/platform/macros.h +++ b/paddle/platform/macros.h @@ -16,8 +16,10 @@ limitations under the License. */ // Disable the copy and assignment operator for a class. #ifndef DISABLE_COPY_AND_ASSIGN -#define DISABLE_COPY_AND_ASSIGN(classname) \ - private: \ - classname(const classname&) = delete; \ - classname& operator=(const classname&) = delete +#define DISABLE_COPY_AND_ASSIGN(classname) \ + private: \ + classname(const classname&) = delete; \ + classname(const classname&&) = delete; \ + classname& operator=(const classname&) = delete; \ + classname& operator=(const classname&&) = delete #endif diff --git a/paddle/platform/variant.h b/paddle/platform/variant.h index 16ee00efe7a9b0406f8459e19a55e1e1b9ca7419..619897ca19eb2e6f4dbfd9160edf8c4bc58c89a9 100644 --- a/paddle/platform/variant.h +++ b/paddle/platform/variant.h @@ -16,7 +16,7 @@ #include -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA // Because boost's variadic templates has bug on nvcc, boost will disable // variadic template support when GPU enabled on nvcc. diff --git a/paddle/pserver/test/SocketTest.cpp b/paddle/pserver/test/SocketTest.cpp index 6f6c9e596cfb7a2547d5b6c5de69381eb9c29132..b43461d61bab21747e85090bbf7af21a87a670c6 100644 --- a/paddle/pserver/test/SocketTest.cpp +++ b/paddle/pserver/test/SocketTest.cpp @@ -215,7 +215,7 @@ int main(int argc, char** argv) { uint64_t dataSize = FLAGS_dim * sizeof(real); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA GpuVector gpuParam(FLAGS_dim); GpuVector gpuGrad(FLAGS_dim); #else diff --git a/paddle/pserver/test/test_ProtoServer.cpp b/paddle/pserver/test/test_ProtoServer.cpp index 04236fda2fb62b928b5c06ff38acfd3eb7217b08..ad8ffed9c1c8e4bdef27689ab21950db6b5cf0a2 100644 --- a/paddle/pserver/test/test_ProtoServer.cpp +++ b/paddle/pserver/test/test_ProtoServer.cpp @@ -99,7 +99,7 @@ TEST(ProtoServer, regular) { } TEST(ProtoServer, extended) { -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA ProtoClient* client; if (FLAGS_rdma_tcp == "rdma") client = new ProtoClient(FLAGS_server_addr, FLAGS_port, F_RDMA); diff --git a/paddle/pybind/pybind.cc b/paddle/pybind/pybind.cc index f4121e9d71824296770f86c1e94c096f767dec0a..cff54b174134879a3779c7738cfc3b43a074f8d7 100644 --- a/paddle/pybind/pybind.cc +++ b/paddle/pybind/pybind.cc @@ -34,7 +34,7 @@ static size_t UniqueIntegerGenerator() { } bool IsCompileGPU() { -#ifdef PADDLE_ONLY_CPU +#ifndef PADDLE_WITH_CUDA return false; #else return true; @@ -78,7 +78,7 @@ PYBIND11_PLUGIN(core) { .def("set", PyCPUTensorSetFromArray) .def("set", PyCPUTensorSetFromArray) .def("set", PyCPUTensorSetFromArray) -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA .def("set", PyCUDATensorSetFromArray) .def("set", PyCUDATensorSetFromArray) .def("set", PyCUDATensorSetFromArray) @@ -96,7 +96,7 @@ PYBIND11_PLUGIN(core) { .def( "__init__", [](LoDTensor &instance, const std::vector> &lod) { -#ifdef PADDLE_ONLY_CPU +#ifndef PADDLE_WITH_CUDA new (&instance) LoDTensor(lod); #else LoD new_lod; @@ -107,7 +107,7 @@ PYBIND11_PLUGIN(core) { }) .def("set_lod", [](LoDTensor &self, const std::vector> &lod) { -#ifdef PADDLE_ONLY_CPU +#ifndef PADDLE_WITH_CUDA self.set_lod(lod); #else LoD new_lod; @@ -117,7 +117,7 @@ PYBIND11_PLUGIN(core) { #endif }) .def("lod", [](LoDTensor &self) -> std::vector> { -#ifdef PADDLE_ONLY_CPU +#ifndef PADDLE_WITH_CUDA return self.lod(); #else auto lod = self.lod(); @@ -143,6 +143,13 @@ All parameter, weight, gradient are variables in Paddle. .def("set_int", [](Variable &var, int val) -> void { *var.GetMutable() = val; }) .def("get_int", [](const Variable &var) -> int { return var.Get(); }) + .def("is_float", [](const Variable &var) { return var.IsType(); }) + .def("set_float", + [](Variable &var, float val) -> void { + *var.GetMutable() = val; + }) + .def("get_float", + [](const Variable &var) -> float { return var.Get(); }) .def("get_tensor", [](Variable &self) -> LoDTensor * { return self.GetMutable(); @@ -196,7 +203,7 @@ All parameter, weight, gradient are variables in Paddle. .def_static("create", [](paddle::platform::GPUPlace& place) -> paddle::platform::DeviceContext* { -#ifdef PADDLE_ONLY_CPU +#ifndef PADDLE_WITH_CUDA PADDLE_THROW("GPUPlace is not supported in CPU device."); #else return new paddle::platform::CUDADeviceContext(place); diff --git a/paddle/pybind/tensor_py.h b/paddle/pybind/tensor_py.h index 3e3e6bc0312974fab50e17d428c7dea9ca547d1e..9e73f79cbdd545db558bd8641bc52e4bf3b0664f 100644 --- a/paddle/pybind/tensor_py.h +++ b/paddle/pybind/tensor_py.h @@ -106,7 +106,7 @@ void PyCPUTensorSetFromArray( std::memcpy(dst, array.data(), sizeof(T) * array.size()); } -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA template void PyCUDATensorSetFromArray( framework::Tensor &self, diff --git a/paddle/string/to_string_test.cc b/paddle/string/to_string_test.cc index 542c771a98ec8ae187cd4f821ed1ee4373427041..971484dd0c073762e99f3926576eb21b96197769 100644 --- a/paddle/string/to_string_test.cc +++ b/paddle/string/to_string_test.cc @@ -36,4 +36,4 @@ TEST(to_string, user_defined) { using namespace paddle::string; UserDefinedClass instance; ASSERT_EQ(kOutputString, to_string(instance)); -} \ No newline at end of file +} diff --git a/paddle/trainer/MergeModel.cpp b/paddle/trainer/MergeModel.cpp index 91d89b61a32259b8bbe70fda2579f87ec6b9af00..6c52eaf4494bb247324b29981d94d7e97e0f212a 100644 --- a/paddle/trainer/MergeModel.cpp +++ b/paddle/trainer/MergeModel.cpp @@ -29,7 +29,7 @@ int main(int argc, char** argv) { initMain(argc, argv); initPython(argc, argv); string confFile = TrainerConfigHelper::getConfigNameFromPath(FLAGS_model_dir); -#ifdef PADDLE_ONLY_CPU +#ifndef PADDLE_WITH_CUDA FLAGS_use_gpu = false; #endif auto config = std::make_shared(confFile); diff --git a/paddle/trainer/tests/test_Compare.cpp b/paddle/trainer/tests/test_Compare.cpp index e855a8fe2e09aa0f16a73f3e7bcc2f32921092f8..f3a964acb69be059a43470f7b68910a3b6cecaab 100644 --- a/paddle/trainer/tests/test_Compare.cpp +++ b/paddle/trainer/tests/test_Compare.cpp @@ -146,7 +146,7 @@ void compareGradient(comData& comDataCpu, comData& comDataGpu) { } int main(int argc, char** argv) { -#ifdef PADDLE_ONLY_CPU +#ifndef PADDLE_WITH_CUDA exit(0); #endif paddle::initMain(argc, argv); diff --git a/paddle/trainer/tests/test_CompareSparse.cpp b/paddle/trainer/tests/test_CompareSparse.cpp index 813275518e411d6e963e23df634541f771096e0f..5f1834bd730375fc10762fc19788d0c693f8e752 100644 --- a/paddle/trainer/tests/test_CompareSparse.cpp +++ b/paddle/trainer/tests/test_CompareSparse.cpp @@ -174,7 +174,7 @@ TEST(compareSparse, multiGradientMachine) { FLAGS_local = local; FLAGS_ports_num_for_sparse = 5; for (bool useGpu : {false, true}) { -#ifdef PADDLE_ONLY_CPU +#ifndef PADDLE_WITH_CUDA if (useGpu) continue; #endif FLAGS_parallel_nn = useGpu; @@ -198,7 +198,7 @@ TEST(compareSparse, NeuralNetwork) { FLAGS_local = local; FLAGS_ports_num_for_sparse = 5; for (bool useGpu : {false, true}) { -#ifdef PADDLE_ONLY_CPU +#ifndef PADDLE_WITH_CUDA if (useGpu) continue; #endif FLAGS_parallel_nn = useGpu; diff --git a/paddle/trainer/tests/test_Trainer.cpp b/paddle/trainer/tests/test_Trainer.cpp index 264bc46ebcd0aa17fd605e537fcb2c316ef31162..425b3d10a38086463784ba2a18db1293efe96e92 100644 --- a/paddle/trainer/tests/test_Trainer.cpp +++ b/paddle/trainer/tests/test_Trainer.cpp @@ -51,7 +51,7 @@ void checkGradientTest(const string& configFile, TEST(checkGradient, cpu) { checkGradientTest(configFile1, false, false); } -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA TEST(checkGradient, gpu) { checkGradientTest(configFile1, true, false); } TEST(checkGradient, multiGpu) { @@ -97,7 +97,7 @@ TEST(checkGradient, hsigmoid) { checkGradientTest(configFile2, false, false); } TEST(checkGradient, chunk) { checkGradientTest(configFile3, false, false); -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA checkGradientTest(configFile3, true, true); #endif } diff --git a/paddle/trainer/tests/test_TrainerOnePass.cpp b/paddle/trainer/tests/test_TrainerOnePass.cpp index 00ba61377aeff17d82e03f7560c0d71b3570d14f..b2a93d4d5eea37ad716b59427f2aa4409d2f537d 100644 --- a/paddle/trainer/tests/test_TrainerOnePass.cpp +++ b/paddle/trainer/tests/test_TrainerOnePass.cpp @@ -79,7 +79,7 @@ void trainerOnePassTest(const string& configFile, // 1. test trainer (cpu, gpu). TEST(trainerOnePass, cpu) { trainerOnePassTest(configFile1, false, false); } -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA TEST(trainerOnePass, gpu) { trainerOnePassTest(configFile1, true, false); } TEST(trainerOnePass, gpu2) { trainerOnePassTest(configFile1, true, false, 2); } @@ -94,7 +94,7 @@ TEST(trainerOnePass, parallel) { #endif // 2. test average_window. -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA TEST(average_window, gpu) { trainerOnePassTest(configFile1, true, false, 4, 0.01); } @@ -266,7 +266,7 @@ TEST(checkRemoteUpdater, cpuTrainerOldUpdater) { checkRemoteParameterUpdaterTest(configFile1, false, false, 1, true); } -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA TEST(checkRemoteUpdater, gpuTrainer) { checkRemoteParameterUpdaterTest(configFile1, true, false); } diff --git a/paddle/trainer/tests/test_recurrent_machine_generation.cpp b/paddle/trainer/tests/test_recurrent_machine_generation.cpp index 1322e77178a4f5674f41943f886a17be8337bd75..a8fbe31c2b1e228107dfc19483444409bfcbf788 100644 --- a/paddle/trainer/tests/test_recurrent_machine_generation.cpp +++ b/paddle/trainer/tests/test_recurrent_machine_generation.cpp @@ -113,7 +113,7 @@ void testGeneration(const string& configFile, #ifndef PADDLE_TYPE_DOUBLE TEST(RecurrentGradientMachine, test_generation) { -#ifdef PADDLE_ONLY_CPU +#ifndef PADDLE_WITH_CUDA const auto useGpuConfs = {false}; #else const auto useGpuConfs = {true, false}; diff --git a/paddle/utils/Flags.cpp b/paddle/utils/Flags.cpp index ab1c181c62cdbee8cc5f804ec9aaf63ac5464ad6..8f100f02e90bcbc7fdcf6f053aec6f95cfb09c1a 100644 --- a/paddle/utils/Flags.cpp +++ b/paddle/utils/Flags.cpp @@ -14,7 +14,7 @@ limitations under the License. */ #include "Flags.h" -#ifdef PADDLE_ONLY_CPU +#ifndef PADDLE_WITH_CUDA DEFINE_bool(use_gpu, false, "Only support CPU training"); #else DEFINE_bool(use_gpu, true, "Whether to use GPU for training"); diff --git a/paddle/utils/Util.h b/paddle/utils/Util.h index 22ce2534d3468ded36221810aa61c15b37f13f3d..9579881ea3b92abab0189631184bab515afb67a3 100644 --- a/paddle/utils/Util.h +++ b/paddle/utils/Util.h @@ -218,7 +218,7 @@ protected: * *d2* is peer device to enable direct access to by the d1 device. */ inline void enablePeerAccess(int d1, int d2) { -#ifndef PADDLE_ONLY_CPU +#ifdef PADDLE_WITH_CUDA if (hl_device_can_access_peer(d1, d2)) { SetDevice dev(d1); hl_device_enable_peer_access(d2); diff --git a/paddle/utils/Version.h b/paddle/utils/Version.h index f53d6420bbbdf66f8f355af95c6b11c30a3bfab9..004d62451cddfee8fbd687938086e04ecb2332a9 100644 --- a/paddle/utils/Version.h +++ b/paddle/utils/Version.h @@ -48,7 +48,7 @@ void printVersion(std::ostream& os); * @return return true if paddle compiled with GPU */ constexpr bool isWithGpu() { -#ifdef PADDLE_ONLY_CPU +#ifndef PADDLE_WITH_CUDA return false; #else return true; diff --git a/python/paddle/v2/framework/tests/op_test.py b/python/paddle/v2/framework/tests/op_test.py index 75df2eeddfe67269d4709887c7cfdb8fab108bd8..81067f38bbf64ac1ab4ccf02aa43b0a38b7d48ad 100644 --- a/python/paddle/v2/framework/tests/op_test.py +++ b/python/paddle/v2/framework/tests/op_test.py @@ -46,12 +46,17 @@ def create_op(scope, op_type, inputs, outputs, attrs): def set_input(scope, op, inputs, place): def __set_input__(var_name, var): - tensor = scope.find_var(var_name).get_tensor() - if isinstance(var, tuple): - tensor.set_lod(var[1]) - var = var[0] - tensor.set_dims(var.shape) - tensor.set(var, place) + if isinstance(var, tuple) or isinstance(var, np.ndarray): + tensor = scope.find_var(var_name).get_tensor() + if isinstance(var, tuple): + tensor.set_lod(var[1]) + var = var[0] + tensor.set_dims(var.shape) + tensor.set(var, place) + elif isinstance(var, float): + scope.find_var(var_name).set_float(var) + elif isinstance(var, int): + scope.find_var(var_name).set_int(var) for in_name, in_dup in Operator.get_op_inputs(op.type()): if in_name in inputs: diff --git a/python/paddle/v2/framework/tests/test_activation_op.py b/python/paddle/v2/framework/tests/test_activation_op.py index c44eb849063592fbda417ec1516d195dd4358612..f232996a55da86399140ea2aad27428926c1a055 100644 --- a/python/paddle/v2/framework/tests/test_activation_op.py +++ b/python/paddle/v2/framework/tests/test_activation_op.py @@ -122,6 +122,23 @@ class TestBRelu(OpTest): self.check_grad(['X'], 'Y', max_relative_error=0.02) +class TestLeakyRelu(OpTest): + def setUp(self): + self.op_type = "leaky_relu" + alpha = 0.02 + self.attrs = {'alpha': alpha} + self.inputs = {'X': np.random.uniform(-3, 3, [4, 4]).astype("float32")} + self.outputs = { + 'Y': np.maximum(self.inputs['X'], alpha * self.inputs['X']) + } + + def test_check_output(self): + self.check_output() + + def test_check_grad(self): + self.check_grad(['X'], 'Y', max_relative_error=0.007) + + class TestSoftRelu(OpTest): def setUp(self): self.op_type = "soft_relu" diff --git a/python/paddle/v2/framework/tests/test_adadelta_op.py b/python/paddle/v2/framework/tests/test_adadelta_op.py new file mode 100644 index 0000000000000000000000000000000000000000..7105593a98aee9885ba16e3ee0649a6024033ee7 --- /dev/null +++ b/python/paddle/v2/framework/tests/test_adadelta_op.py @@ -0,0 +1,96 @@ +import unittest +import numpy as np +from op_test import OpTest + + +class TestAdadeltaOp1(OpTest): + def setUp(self): + self.op_type = "adadelta" + param = np.random.uniform(-1, 1, (102, 105)).astype("float32") + grad = np.random.uniform(-1, 1, (102, 105)).astype("float32") + # The squared gradient is positive + avg_squared_grad = np.random.random((102, 105)).astype("float32") + # The squared update is positive + avg_squared_update = np.random.random((102, 105)).astype("float32") + + rho = 0.95 + epsilon = 1e-6 + + self.inputs = { + 'Param': param, + 'Grad': grad, + 'AvgSquaredGrad': avg_squared_grad, + 'AvgSquaredUpdate': avg_squared_update + } + + self.attrs = {'rho': rho, 'epsilon': epsilon} + + avg_squared_grad_out = rho * avg_squared_grad + \ + (1 - rho) * np.square(grad) + update = -np.multiply( + np.sqrt( + np.divide(avg_squared_update + epsilon, avg_squared_grad_out + + epsilon)), grad) + + avg_squared_update_out = rho * avg_squared_update + \ + (1 - rho) * np.square(update) + + param_out = param + update + + self.outputs = { + 'ParamOut': param_out, + 'AvgSquaredGradOut': avg_squared_grad_out, + 'AvgSquaredUpdateOut': avg_squared_update_out + } + + def test_check_output(self): + self.check_output() + + +class TestAdadeltaOp2(OpTest): + '''Test Adadelta op with default attribute values + ''' + + def setUp(self): + self.op_type = "adadelta" + param = np.random.uniform(-1, 1, (102, 105)).astype("float32") + grad = np.random.uniform(-1, 1, (102, 105)).astype("float32") + # The squared gradient is positive + avg_squared_grad = np.random.random((102, 105)).astype("float32") + # The squared update is positive + avg_squared_update = np.random.random((102, 105)).astype("float32") + + rho = 0.95 + epsilon = 1e-6 + + self.inputs = { + 'Param': param, + 'Grad': grad, + 'AvgSquaredGrad': avg_squared_grad, + 'AvgSquaredUpdate': avg_squared_update + } + + avg_squared_grad_out = rho * avg_squared_grad + \ + (1 - rho) * np.square(grad) + update = -np.multiply( + np.sqrt( + np.divide(avg_squared_update + epsilon, avg_squared_grad_out + + epsilon)), grad) + + avg_squared_update_out = rho * avg_squared_update + \ + (1 - rho) * np.square(update) + + param_out = param + update + + self.outputs = { + 'ParamOut': param_out, + 'AvgSquaredGradOut': avg_squared_grad_out, + 'AvgSquaredUpdateOut': avg_squared_update_out + } + + def test_check_output(self): + self.check_output() + + +if __name__ == "__main__": + unittest.main() diff --git a/python/paddle/v2/framework/tests/test_recurrent_op.py b/python/paddle/v2/framework/tests/test_recurrent_op.py index 92161ae5dd93d34d898a2027435cc5e55611bcd0..1f114432c09f29fab6cd56de00dff341785ae0e4 100644 --- a/python/paddle/v2/framework/tests/test_recurrent_op.py +++ b/python/paddle/v2/framework/tests/test_recurrent_op.py @@ -16,14 +16,17 @@ class PySimpleRNN(object): ''' def __init__(self, input_dim=30, batch_size=50, weight_dim=15, sent_len=11): - self.x = np.random.normal(size=(sent_len, batch_size, input_dim)) - self.W = np.random.normal(size=(input_dim, input_dim)) - self.U = np.random.normal(size=(input_dim, input_dim)) - self.h_boot = np.random.normal(size=(batch_size, input_dim)) + self.x = np.random.normal(size=(sent_len, batch_size, + input_dim)).astype("float32") + self.W = np.random.normal(size=(input_dim, input_dim)).astype("float32") + self.U = np.random.normal(size=(input_dim, input_dim)).astype("float32") + self.h_boot = np.random.normal(size=(batch_size, + input_dim)).astype("float32") # memories self.mems = [ - np.zeros(shape=(batch_size, input_dim)) for i in range(sent_len) + np.zeros(shape=(batch_size, input_dim)).astype("float32") + for i in range(sent_len) ] def forward(self): @@ -36,7 +39,7 @@ class PySimpleRNN(object): return [self.x[i] for i in range(self.x.shape[0])] def concat_outputs(self): - return np.array(self.mems) + return np.array(self.mems).astype("float32") def step(self, step_id, x): ''' @@ -47,8 +50,8 @@ class PySimpleRNN(object): pre_mem = self.mems[step_id - 1] else: pre_mem = self.h_boot - xW = np.matmul(x, self.W) - hU = np.matmul(pre_mem, self.U) + xW = np.matmul(x, self.W).astype("float32") + hU = np.matmul(pre_mem, self.U).astype("float32") sum = xW + hU self.mems[step_id] = py_sigmoid(sum) @@ -102,7 +105,8 @@ class RecurrentOpTest(unittest.TestCase): self.create_step_net() ctx = core.DeviceContext.create(core.CPUPlace()) self.rnnop.run(self.scope, ctx) - return np.array(self.scope.find_var("h@mem").get_tensor()) + return np.array(self.scope.find_var("h@mem").get_tensor()).astype( + "float32") def create_global_variables(self): # create inlink @@ -142,7 +146,7 @@ class RecurrentOpTest(unittest.TestCase): stepnet = core.Net.create() x_fc_op = Operator("mul", X="x", Y="W", Out="Wx") h_fc_op = Operator("mul", X="h@pre", Y="U", Out="Uh") - sum_op = Operator("add", X="Wx", Y="Uh", Out="sum") + sum_op = Operator("sum", X=["Wx", "Uh"], Out="sum") sig_op = Operator("sigmoid", X="sum", Y="h@mem") for op in [x_fc_op, h_fc_op, sum_op, sig_op]: @@ -179,7 +183,7 @@ class RecurrentGradientOpTest(unittest.TestCase): stepnet = core.Net.create() x_fc_op = Operator("mul", X="x@alias", Y="W", Out="Wx") h_fc_op = Operator("mul", X="h@pre", Y="U", Out="Uh") - sum_op = Operator("add", X="Wx", Y="Uh", Out="sum") + sum_op = Operator("sum", X=["Wx", "Uh"], Out="sum") sig_op = Operator("sigmoid", X="sum", Y="h@alias") for op in [x_fc_op, h_fc_op, sum_op, sig_op]: @@ -197,7 +201,4 @@ class RecurrentGradientOpTest(unittest.TestCase): if __name__ == '__main__': - exit( - 0 - ) # FIXME(yuyang18): InferShape has been removed, this unittest may error unittest.main() diff --git a/python/paddle/v2/framework/tests/test_scatter_op.py b/python/paddle/v2/framework/tests/test_scatter_op.py index 33c73c52631a09ea0fefdeb9467991ae9c04321c..1032269d5dfb02e3518b9ef2820d5d0dcc8a51a0 100644 --- a/python/paddle/v2/framework/tests/test_scatter_op.py +++ b/python/paddle/v2/framework/tests/test_scatter_op.py @@ -10,7 +10,7 @@ class TestScatterOp(OpTest): index_np = np.array([1, 2]).astype("int32") updates_np = np.random.random((2, 3)).astype("float32") output_np = np.copy(ref_np) - output_np[index_np] += updates_np + output_np[index_np] = updates_np self.inputs = {'Ref': ref_np, 'Index': index_np, 'Updates': updates_np} self.outputs = {'Out': output_np} @@ -18,7 +18,7 @@ class TestScatterOp(OpTest): self.check_output() def test_check_grad(self): - self.check_grad(['Updates', 'Ref'], 'Out', in_place=True) + self.check_grad(['Updates'], 'Out', in_place=True) if __name__ == "__main__": diff --git a/python/paddle/v2/framework/tests/test_sgd_op.py b/python/paddle/v2/framework/tests/test_sgd_op.py index 64e54d1500c1bc134cc1efe33d41a16dbc08f2d4..2dd881e5e107249277a91bd8e3a72567269e1cd4 100644 --- a/python/paddle/v2/framework/tests/test_sgd_op.py +++ b/python/paddle/v2/framework/tests/test_sgd_op.py @@ -8,11 +8,10 @@ class TestSGDOp(OpTest): self.op_type = "sgd" w = np.random.random((102, 105)).astype("float32") g = np.random.random((102, 105)).astype("float32") - lr = 0.1 + lr = np.array([0.1]).astype("float32") - self.inputs = {'param': w, 'grad': g} - self.attrs = {'learning_rate': lr} - self.outputs = {'param_out': w - lr * g} + self.inputs = {'Param': w, 'Grad': g, 'LearningRate': lr} + self.outputs = {'ParamOut': w - lr * g} def test_check_output(self): self.check_output()