Merge branch 'develop' into seqpool

doc/design/optimizer.md

+## Optimizer Design
+
+### The Problem
+
+A PaddlePaddle program, or a block, is a sequence of operators operating variables.  A training program needs to do three kinds of works:
+
+1. the forward pass, which computes intermediate results and the cost(s),
+1. the backward pass, which derives gradients from intermediate results and costs, and
+1. the optimization pass, which update model parameters to optimize the cost(s).
+
+These works rely on three kinds of operators:
+
+1. forward operators,
+1. gradient operators, and
+1. optimization operators.
+
+It's true that users should be able to create all these operators manually by calling some low-level API, but it would be much more convenient if they could only describe the forward pass and let PaddlePaddle create the backward and optimization operators automatically.
+
+In this design, we propose a high-level API that automatically derives the optimisation pass and operators from the forward pass.
+
+
+### High-level Python API to describe the training process
+
+1. User write code to describe the network:
+
+	```python
+	images = layer.data("images")
+	labels = layer.data("labels")
+	w1 = pd.var("w1")
+	b1 = pd.var("b1")
+	hidden = layer.fc(images, w=w1, b=b1)
+	cost = layer.mse(hidden, labels)
+	```
+
+	The above code snippet will create forward operators in [Block](https://github.com/PaddlePaddle/Paddle/blob/develop/doc/design/block.md).
+
+
+2. Users create a certain kind of Optimizer with some argument.
+
+	```python
+	optimizer = AdagradOptimizer(learing_rate=0.001)
+	```
+
+3. Users use the optimizer to `minimize` a certain `cost` through updating parameters in parameter_list.
+
+	```python
+	opt_op_list = optimizer.minimize(cost, parameter_list=[w1, b1])
+	```
+	The above code snippet will create gradient and optimization operators in Block. The return value of `minimize()` is list of optimization operators that will be run by session.
+
+4. Users use Session/Executor to run this opt_op_list as target to do training.
+
+	```python
+	sess.run(target= opt_op_list, ...)
+	```
+
+#### Optimizer Python interface:
+
+```python
+class Optimizer(object):
+    """Optimizer Base class.
+
+    """
+
+    def __init__(self):
+        pass
+
+    def create_backward_pass(self, loss, parameter_list=None):
+        """
+        create and add gradient Operators in BlockDesc to Compute gradients of `loss`
+        for parameters in parameter_list
+
+        Args:
+          loss: an variable generated by cost function.
+          parameter_list: parameters that need to compute gradient and update to optimize the lost.
+
+        Returns:
+          list of (parameters, gradients) pair.
+        """
+        return None
+
+    def create_optimization_pass(self, parameters_and_grads):
+        """Add optimization operators to update gradients to variables.
+
+        Args:
+          parameters_and_grads: a list of (variable, gradient) pair to update.
+
+        Returns:
+          optmization_op_list: a list of optimization operator that will update parameter using gradient.
+        """
+        return None
+
+    def minimize(self, loss, parameter_list):
+        """Add operations to minimize `loss` by updating `parameter_list`.
+
+        This method combines interface `create_backward_pass()` and
+        `create_optimization_pass()` into one.
+        """
+        params_grads = self.create_backward_pass(loss, parameter_list)
+        update_ops = self.create_optimization_pass(params_grads)
+        return update_ops
+
+```
+
+Users can inherit the Optimizer above to create their own Optimizer with some special logic, such as AdagradOptimizer.

doc/design/python_api.md

@@ -214,3 +214,7 @@ def fc_layer(input, size, ...):
     out.writer = op
     return out
 ```
+
+## Optimizer
+
+[Optimizer Design Doc](./optimizer.md)

doc/design/selected_rows.md

+# Design Doc: Selected Rows
+
+`SelectedRows` is a kind of sparse tensor data type, which is designed to support `embedding` operators. The gradient of embedding table is a sparse tensor. Only a few rows are non-zero values in that tensor. It is straightforward to represent the sparse tensor by the following sparse tensor data structure:
+
+```cpp
+class SelectedRows {
+ private:
+  vector<int> rows_;
+  Tensor value_;
+  int height_;
+};
+```
+
+The field `height_` shows the first dimension of `SelectedRows`. The `rows` are the indices of which rows of `SelectedRows` are non-zeros. The `value_` field is an N-dim tensor and shape is `[rows.size() /* NUM_ROWS */, ...]`, which supplies values for each row. The dimension of `SelectedRows` satisfies `[height_] + value_.shape[1:]`.
+
+Suppose that a SelectedRows-typed variable `x` has many rows, but only two of them have values -- row 73 is `[1, 2]` and row 84 is `[3, 4]`, the `SelectedRows` representation would be:
+
+```
+x = SelectedRow {
+  rows = [73, 84],
+  value = [[1, 2], [3,4]]
+}
+```
+
+
+## SelectedRows in Protobuf
+
+`SelectedRows` is a kind of `Variable`. `VarDesc` in protobuf should describe the `SelectedRows` information. Only the tensor dimension of a `SelectedRows` will be described in compile-time since the `rows_` and `value_` are related to training data. 
+So we use `TensorDesc` to unify `data_type` and `dims`. A LodTensorDesc contains a `TensorDesc` and `lod_level`. The description of `SelectedRows` is a Tensor description.
+
+```proto
+message TensorDesc {
+  required DataType data_type = 1;
+  repeated int64 dims = 2; // [UNK, 640, 480] is saved as [-1, 640, 480]
+}
+
+message LodTensorDesc {
+  required TensorDesc tensor = 1;
+  optional int lod_level = 2;
+}
+
+message VarDesc {
+  required string name = 1;
+  enum VarType { 
+    LOD_TENSOR = 0;
+    SELECTED_ROWS = 1;
+  }
+  required VarType type = 2;
+  optional LodTensorDesc lod_desc = 3;
+  optional TensorDesc selected_rows_desc = 4;
+  optional bool persistable = 5 [ default = false ];
+}
+```
+
+## InferShape for Selected Rows
+
+Just like `LoD` information, `InferShape` method will inference output tensor type as well. The operator should decide whether its output is a `SelectedRows` or `Dense` tensor.
+
+For example, the gradient operator of `TableLookup` will always generate `SelectedRows`. Its `InferShape` method should be like following
+
+```cpp
+void TableLookupGrad::InferShape(context) {
+  ...
+  context.SetDataType("Embedding.Grad", kSelectedRows);
+}
+```
+
+
+## Sparse Operators
+
+There are several operators should be written to support `SelectedRows`. They are:
+
+1. Operators which generates `SelectedRows` gradient. e.g. Gradient of `TableLookupOp`.
+2. Optimize operators which support `SelectedRows` gradient. e.g. `SGD` or `AdaGrad` for `SelectedRows`. However, there should be only one `SGD` operator. `OpWithKernel::Run` should select a suitable kernel for both `dense` tensor or `SelectedRows`.

paddle/framework/CMakeLists.txt

@@ -42,5 +42,12 @@ add_custom_command(TARGET framework_py_proto POST_BUILD
 cc_library(backward SRCS backward.cc DEPS net_op)
 cc_test(backward_test SRCS backward_test.cc DEPS backward recurrent_op device_context)
 
+cc_library(executor SRCS executor.cc DEPS op_registry device_context scope framework_proto backward ${GLOB_OP_LIB})
+if(WITH_GPU)
+    nv_test(executor_test SRCS executor_test.cc DEPS executor)
+else()
+    cc_test(executor_test SRCS executor_test.cc DEPS executor)
+endif()
+
 cc_library(tensor_array SRCS tensor_array.cc DEPS lod_tensor)
 cc_test(tensor_array_test SRCS tensor_array_test.cc DEPS tensor_array place)

paddle/framework/backward.cc

@@ -172,30 +172,14 @@ static std::unique_ptr<OperatorBase> BackwardRecursive(
                               std::to_string(i));
         net->ops_[op_offset]->Rename(name, dup_outputs.back());
       }
-      // collect all the offset to append `add` op for each alias
-      //
-      // 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}}}, {})});
-      }
+      // collect all the offset for each alias,
+      // insert a sum operator to add all aliases to output
+      insert_position.push_back(
+          {dup_op.back(), OpRegistry::CreateOp("sum", {{"X", dup_outputs}},
+                                               {{"Out", {name}}}, {})});
     }
 
-    // make sure the inserted `add` ops follow the BFS order.
+    // make sure the inserted `sum` ops follow the BFS order.
     insert_position.sort(
         [](const Pos& l, const Pos& r) { return l.first > r.first; });
 

paddle/framework/executor.cc

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+    http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License. */
+
+#include "paddle/framework/executor.h"
+
+#include <algorithm>
+#include <iostream>
+#include <memory>
+#include <set>
+#include <vector>
+
+#include "paddle/framework/lod_tensor.h"
+#include "paddle/framework/op_registry.h"
+#include "paddle/framework/scope.h"
+
+namespace paddle {
+namespace framework {
+
+const std::string kFeedOpType = "feed";
+const std::string kFetchOpType = "fetch";
+
+Executor::Executor(const std::vector<platform::Place>& places) {
+  PADDLE_ENFORCE_GT(places.size(), 0);
+  device_contexts_.resize(places.size());
+  for (size_t i = 0; i < places.size(); i++) {
+    if (platform::is_cpu_place(places[i])) {
+      device_contexts_[i] = new platform::CPUDeviceContext(
+          boost::get<platform::CPUPlace>(places[i]));
+    } else if (platform::is_gpu_place(places[i])) {
+#ifdef PADDLE_WITH_CUDA
+      device_contexts_[i] = new platform::CUDADeviceContext(
+          boost::get<platform::GPUPlace>(places[i]));
+#else
+      PADDLE_THROW(
+          "'GPUPlace' is not supported, Please re-compile with WITH_GPU "
+          "option");
+#endif
+    }
+  }
+}
+
+Executor::~Executor() {
+  for (auto& device_context : device_contexts_) {
+    delete device_context;
+  }
+}
+
+void Executor::Run(const ProgramDesc& pdesc, Scope* scope, int block_id) {
+  // TODO(tonyyang-svail):
+  //    - only runs on the first device (i.e. no interdevice communication)
+  //    - will change to use multiple blocks for RNN op and Cond Op
+  PADDLE_ENFORCE_GT(pdesc.blocks_size(), block_id);
+  auto& block = pdesc.blocks(block_id);
+  auto& device = device_contexts_[0];
+
+  // Instantiate all the vars in the global scope
+  for (auto& var : block.vars()) {
+    scope->NewVar(var.name());
+  }
+
+  Scope& local_scope = scope->NewScope();
+
+  std::vector<bool> should_run = Prune(pdesc, block_id);
+  PADDLE_ENFORCE_EQ(should_run.size(), static_cast<size_t>(block.ops_size()));
+  for (size_t i = 0; i < should_run.size(); ++i) {
+    if (should_run[i]) {
+      for (auto& var : block.ops(i).outputs()) {
+        for (auto& argu : var.arguments()) {
+          if (local_scope.FindVar(argu) == nullptr) {
+            local_scope.NewVar(argu);
+          }
+        }
+      }
+      auto op = paddle::framework::OpRegistry::CreateOp(block.ops(i));
+      op->Run(local_scope, *device);
+    }
+  }
+
+  // TODO(tonyyang-svail):
+  //  - Destroy local_scope
+}
+
+std::vector<bool> Prune(const ProgramDesc& pdesc, int block_id) {
+  // TODO(tonyyang-svail):
+  //    - will change to use multiple blocks for RNN op and Cond Op
+
+  auto& block = pdesc.blocks(block_id);
+  auto& ops = block.ops();
+
+  bool expect_feed = true;
+  for (auto& op_desc : ops) {
+    PADDLE_ENFORCE(op_desc.type() != kFeedOpType || expect_feed,
+                   "All FeedOps are at the beginning of the ProgramDesc");
+    expect_feed = (op_desc.type() == kFeedOpType);
+  }
+
+  bool expect_fetch = true;
+  for (auto op_iter = ops.rbegin(); op_iter != ops.rend(); ++op_iter) {
+    auto& op_desc = *op_iter;
+    PADDLE_ENFORCE(op_desc.type() != kFetchOpType || expect_fetch,
+                   "All FetchOps must at the end of the ProgramDesc");
+    expect_fetch = (op_desc.type() == kFetchOpType);
+  }
+
+  std::set<std::string> dependent_vars;
+  std::vector<bool> should_run;
+  for (auto op_iter = ops.rbegin(); op_iter != ops.rend(); ++op_iter) {
+    auto& op_desc = *op_iter;
+
+    bool found_dependent_vars = false;
+    for (auto& var : op_desc.outputs()) {
+      for (auto& argu : var.arguments()) {
+        if (dependent_vars.count(argu) != 0) {
+          found_dependent_vars = true;
+        }
+      }
+    }
+
+    if (op_desc.type() == kFetchOpType || found_dependent_vars) {
+      // erase its output to the dependency graph
+      for (auto& var : op_desc.outputs()) {
+        for (auto& argu : var.arguments()) {
+          dependent_vars.erase(argu);
+        }
+      }
+
+      // insert its input to the dependency graph
+      for (auto& var : op_desc.inputs()) {
+        for (auto& argu : var.arguments()) {
+          dependent_vars.insert(argu);
+        }
+      }
+
+      should_run.push_back(true);
+    } else {
+      should_run.push_back(false);
+    }
+  }
+
+  // TODO(tonyyang-svail):
+  //    - check this after integration of Init
+  // PADDLE_ENFORCE(dependent_vars.empty());
+
+  // since we are traversing the ProgramDesc in reverse order
+  // we reverse the should_run vector
+  std::reverse(should_run.begin(), should_run.end());
+
+  return should_run;
+}
+
+}  // namespace framework
+}  // namespace paddle

paddle/framework/executor.h

+/* 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/framework.pb.h"
+#include "paddle/framework/op_info.h"
+#include "paddle/framework/scope.h"
+#include "paddle/framework/tensor.h"
+
+namespace paddle {
+namespace framework {
+
+class Executor {
+ public:
+  explicit Executor(const std::vector<platform::Place>& places);
+  ~Executor();
+
+  /* @Brief
+   * Runtime evaluation of the given ProgramDesc under certain Scope
+   *
+   * @param
+   *  ProgramDesc
+   *  Scope
+   */
+  void Run(const ProgramDesc&, Scope*, int);
+
+ private:
+  std::vector<platform::DeviceContext*> device_contexts_;
+};
+
+/* @Brief
+ * Pruning the graph
+ *
+ * @param
+ *  ProgramDesc
+ *
+ * @return
+ *  vector<bool> Same size as ops. Indicates whether an op should be run.
+ */
+std::vector<bool> Prune(const ProgramDesc& pdesc, int block_id);
+
+}  // namespace framework
+}  // namespace paddle

paddle/framework/executor_test.cc

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+    http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License. */
+
+#include "paddle/framework/executor.h"
+
+#include <memory>
+#include <vector>
+
+#include "gtest/gtest.h"
+#include "paddle/framework/attribute.h"
+#include "paddle/framework/backward.h"
+#include "paddle/framework/block_desc.h"
+#include "paddle/framework/op_desc.h"
+#include "paddle/framework/op_registry.h"
+#include "paddle/framework/operator.h"
+
+USE_OP(elementwise_add);
+USE_OP(gaussian_random);
+USE_OP(feed);
+USE_OP(fetch);
+USE_OP(mul);
+USE_OP(sum);
+USE_OP(squared_l2_distance);
+USE_OP(fill_constant);
+USE_OP(sgd);
+
+using namespace paddle::platform;
+using namespace paddle::framework;
+
+void AddOp(const std::string& type, const VariableNameMap& inputs,
+           const VariableNameMap& outputs, AttributeMap attrs,
+           paddle::framework::BlockDescBind* block) {
+  // insert output
+  for (auto kv : outputs) {
+    for (auto v : kv.second) {
+      auto var = block->NewVar(v);
+      var->SetDataType(paddle::framework::DataType::FP32);
+    }
+  }
+
+  // insert op
+  auto op = block->AppendOp();
+  op->SetType(type);
+  for (auto& kv : inputs) {
+    op->SetInput(kv.first, kv.second);
+  }
+  for (auto& kv : outputs) {
+    op->SetOutput(kv.first, kv.second);
+  }
+  op->SetAttrMap(attrs);
+}
+
+// Tensors in feed value variable will only be in CPUPlace
+// So we can memcpy the data from vector<T> to feed_value
+template <typename T>
+void SetFeedVariable(const std::vector<std::vector<T>>& inputs,
+                     const std::vector<std::vector<int64_t>>& dims) {
+  Variable* g_feed_value = GetGlobalScope().FindVar("feed_value");
+  auto& feed_inputs =
+      *(g_feed_value->GetMutable<std::vector<paddle::framework::Tensor>>());
+  size_t size = inputs.size();
+  feed_inputs.resize(size);
+  for (size_t i = 0; i < size; i++) {
+    T* dst = feed_inputs[i].mutable_data<T>(make_ddim(dims[i]), CPUPlace());
+    memcpy(dst, inputs[i].data(), inputs[i].size() * sizeof(T));
+  }
+}
+
+// Tensors in fetch value variable will only be in CPUPlace
+// So we can memcpy the data from fetch_value to vector<T>
+template <typename T>
+std::vector<std::vector<T>> GetFetchVariable() {
+  Variable* g_fetch_value = GetGlobalScope().FindVar("fetch_value");
+  auto& fetch_outputs =
+      *(g_fetch_value->GetMutable<std::vector<paddle::framework::Tensor>>());
+
+  size_t size = fetch_outputs.size();
+  std::vector<std::vector<T>> result;
+  result.reserve(size);
+  for (size_t i = 0; i < size; i++) {
+    std::vector<T> tmp;
+    tmp.resize(fetch_outputs[i].numel());
+    memcpy(tmp.data(), fetch_outputs[i].data<T>(),
+           fetch_outputs[i].numel() * sizeof(T));
+    result.push_back(tmp);
+  }
+
+  return result;
+}
+
+class ExecutorTesterRandom : public ::testing::Test {
+ public:
+  virtual void SetUp() override {
+    int input_dim = 3, batch_size = 2, embed_dim = 5;
+
+    auto temp_init_root_block = init_pdesc_.add_blocks();
+    temp_init_root_block->set_idx(0);
+    temp_init_root_block->set_parent_idx(-1);
+    paddle::framework::ProgramDescBind& init_program =
+        paddle::framework::ProgramDescBind::Instance(&init_pdesc_);
+    paddle::framework::BlockDescBind* init_root_block = init_program.Block(0);
+
+    AddOp("gaussian_random", {}, {{"Out", {"w1"}}},
+          {{"dims", std::vector<int>{input_dim, embed_dim}}}, init_root_block);
+    AddOp("gaussian_random", {}, {{"Out", {"w2"}}},
+          {{"dims", std::vector<int>{embed_dim, input_dim}}}, init_root_block);
+    AddOp("fetch", {{"Input", {"w1"}}}, {}, {{"col", 0}}, init_root_block);
+    AddOp("fetch", {{"Input", {"w2"}}}, {}, {{"col", 1}}, init_root_block);
+
+    // flush
+    init_program.Proto();
+
+    // run block
+    auto temp_root_block = pdesc_.add_blocks();
+    temp_root_block->set_idx(0);
+    temp_root_block->set_parent_idx(-1);
+    paddle::framework::ProgramDescBind& program =
+        paddle::framework::ProgramDescBind::Instance(&pdesc_);
+    paddle::framework::BlockDescBind* root_block = program.Block(0);
+
+    // feed data
+    inputs_.push_back({1.0, 1.0, 1.0, 1.0, 1.0, 1.0});
+    dims_.push_back({batch_size, input_dim});
+    AddOp("feed", {}, {{"Out", {"a"}}},
+          {{"dims", std::vector<int>{batch_size, input_dim}}, {"col", 0}},
+          root_block);
+
+    // forward
+    AddOp("mul", {{"X", {"a"}}, {"Y", {"w1"}}}, {{"Out", {"b"}}}, {},
+          root_block);
+    AddOp("mul", {{"X", {"b"}}, {"Y", {"w2"}}}, {{"Out", {"a_out"}}}, {},
+          root_block);
+    AddOp("squared_l2_distance", {{"X", {"a"}}, {"Y", {"a_out"}}},
+          {{"Out", {"l2_distance"}}, {"sub_result", {"l2_distance_sub"}}}, {},
+          root_block);
+
+    // backward
+    AddOp("fill_constant", {}, {{"Out", {"l2_distance@GRAD"}}},
+          {{"shape", std::vector<int>{batch_size, 1}}, {"value", float(1.0)}},
+          root_block);
+    AppendBackward(program, {});
+
+    // update
+    AddOp("fill_constant", {}, {{"Out", {"learning_rate"}}},
+          {{"shape", std::vector<int>{1}}, {"value", float(0.001)}},
+          root_block);
+    AddOp("sgd", {{"Param", {"w1"}},
+                  {"LearningRate", {"learning_rate"}},
+                  {"Grad", {"w1@GRAD"}}},
+          {{"ParamOut", {"w1"}}}, {}, root_block);
+    AddOp("sgd", {{"Param", {"w2"}},
+                  {"LearningRate", {"learning_rate"}},
+                  {"Grad", {"w2@GRAD"}}},
+          {{"ParamOut", {"w2"}}}, {}, root_block);
+
+    AddOp("fetch", {{"Input", {"w1"}}}, {}, {{"col", 0}}, root_block);
+    AddOp("fetch", {{"Input", {"w2"}}}, {}, {{"col", 1}}, root_block);
+    AddOp("fetch", {{"Input", {"l2_distance"}}}, {}, {{"col", 0}}, root_block);
+
+    // flush
+    program.Proto();
+  }
+
+ protected:
+  ProgramDesc init_pdesc_;
+  ProgramDesc pdesc_;
+  std::vector<std::vector<float>> inputs_;
+  std::vector<std::vector<int64_t>> dims_;
+};
+
+class ExecutorTesterFeedAndFetch : public ::testing::Test {
+ public:
+  virtual void SetUp() override {
+    auto temp_root_block = pdesc_.add_blocks();
+    temp_root_block->set_idx(0);
+    temp_root_block->set_parent_idx(-1);
+
+    // wrap to BlockDescBind
+    paddle::framework::ProgramDescBind& program =
+        paddle::framework::ProgramDescBind::Instance(&pdesc_);
+    paddle::framework::BlockDescBind* root_block = program.Block(0);
+
+    std::vector<int> dim{6};
+
+    AddOp("feed", {}, {{"Out", {"a"}}}, {{"dims", dim}, {"col", 0}},
+          root_block);
+    AddOp("feed", {}, {{"Out", {"b"}}}, {{"dims", dim}, {"col", 1}},
+          root_block);
+    AddOp("fetch", {{"Input", {"a"}}}, {}, {{"col", 0}}, root_block);
+    AddOp("fetch", {{"Input", {"b"}}}, {}, {{"col", 1}}, root_block);
+
+    // flush
+    program.Proto();
+
+    std::vector<float> vec1 = {1.0, 2.0, 3.0, 4.0, 5.0, 6.0};
+    std::vector<float> vec2 = {4.0, 5.0, 6.0, 7.0, 8.0, 9.0};
+    inputs_.push_back(vec1);
+    inputs_.push_back(vec2);
+    dims_.push_back({static_cast<int64_t>(vec1.size())});
+    dims_.push_back({static_cast<int64_t>(vec2.size())});
+  }
+
+ protected:
+  ProgramDesc pdesc_;
+  std::vector<std::vector<float>> inputs_;
+  std::vector<std::vector<int64_t>> dims_;
+};
+
+#ifndef PADDLE_WITH_CUDA
+TEST_F(ExecutorTesterRandom, CPU) {
+  std::vector<Place> places;
+  CPUPlace cpu_place;
+  places.push_back(cpu_place);
+
+  // We have a global Scope and BuddyAllocator, and we must ensure
+  // global BuddyAllocator is initialized before global Scope. Thus,
+  // global Scope will deconstruct before BuddyAllocator. Otherwise,
+  // "pointer being freed was not allocated" error will appear.
+  paddle::memory::Used(cpu_place);
+
+  std::unique_ptr<Executor> executor(new Executor(places));
+
+  executor->Run(init_pdesc_, &GetGlobalScope(), 0);
+  SetFeedVariable<float>(inputs_, dims_);
+  executor->Run(pdesc_, &GetGlobalScope(), 0);
+  std::vector<std::vector<float>> result = GetFetchVariable<float>();
+}
+
+TEST_F(ExecutorTesterFeedAndFetch, CPU) {
+  std::vector<Place> places;
+  CPUPlace cpu_place;
+  places.push_back(cpu_place);
+
+  // We have a global Scope and BuddyAllocator, and we must ensure
+  // global BuddyAllocator is initialized before global Scope. Thus,
+  // global Scope will deconstruct before BuddyAllocator. Otherwise,
+  // "pointer being freed was not allocated" error will appear.
+  paddle::memory::Used(cpu_place);
+
+  std::unique_ptr<Executor> executor(new Executor(places));
+
+  for (int batch_id = 0; batch_id < 3; batch_id++) {
+    SetFeedVariable<float>(inputs_, dims_);
+    executor->Run(pdesc_, &GetGlobalScope(), 0);
+    std::vector<std::vector<float>> result = GetFetchVariable<float>();
+    PADDLE_ENFORCE_EQ(result.size(), inputs_.size());
+    for (size_t i = 0; i < result.size(); ++i) {
+      PADDLE_ENFORCE_EQ(result[i].size(), inputs_[i].size());
+      for (size_t j = 0; j < result[i].size(); ++j) {
+        PADDLE_ENFORCE_EQ(result[i][j], inputs_[i][j]);
+      }
+    }
+  }
+}
+#else
+TEST_F(ExecutorTesterRandom, GPU) {
+  std::vector<Place> places;
+  GPUPlace gpu_place(0);
+  places.push_back(gpu_place);
+
+  // We have a global Scope and BuddyAllocator, and we must ensure
+  // global BuddyAllocator is initialized before global Scope. Thus,
+  // global Scope will deconstruct before BuddyAllocator. Otherwise,
+  // "pointer being freed was not allocated" error will appear.
+  // If paddle is compiled with GPU, both CPU and GPU BuddyAllocator
+  // need to be used at first.
+  paddle::memory::Used(CPUPlace());
+  paddle::memory::Used(gpu_place);
+
+  std::unique_ptr<Executor> executor(new Executor(places));
+
+  executor->Run(init_pdesc_, &GetGlobalScope(), 0);
+  for (int batch_id = 0; batch_id < 3; batch_id++) {
+    SetFeedVariable<float>(inputs_, dims_);
+    executor->Run(pdesc_, &GetGlobalScope(), 0);
+  }
+}
+
+TEST_F(ExecutorTesterFeedAndFetch, GPU) {
+  std::vector<Place> places;
+  GPUPlace gpu_place(0);
+  places.push_back(gpu_place);
+  // We have a global Scope and BuddyAllocator, and we must ensure
+  // global BuddyAllocator is initialized before global Scope. Thus,
+  // global Scope will deconstruct before BuddyAllocator. Otherwise,
+  // "pointer being freed was not allocated" error will appear.
+  // If paddle is compiled with GPU, both CPU and GPU BuddyAllocator
+  // need to be used at first.
+  paddle::memory::Used(CPUPlace());
+  paddle::memory::Used(gpu_place);
+
+  std::unique_ptr<Executor> executor(new Executor(places));
+
+  for (int batch_id = 0; batch_id < 3; batch_id++) {
+    SetFeedVariable<float>(inputs_, dims_);
+    executor->Run(pdesc_, &GetGlobalScope(), 0);
+    std::vector<std::vector<float>> result = GetFetchVariable<float>();
+    PADDLE_ENFORCE_EQ(result.size(), inputs_.size());
+    for (size_t i = 0; i < result.size(); ++i) {
+      PADDLE_ENFORCE_EQ(result[i].size(), inputs_[i].size());
+      for (size_t j = 0; j < result[i].size(); ++j) {
+        PADDLE_ENFORCE_EQ(result[i][j], inputs_[i][j]);
+      }
+    }
+  }
+}
+#endif

paddle/framework/scope.cc

@@ -13,6 +13,9 @@ See the License for the specific language governing permissions and
 limitations under the License. */
 
 #include "paddle/framework/scope.h"
+
+#include <memory>  // for unique_ptr
+#include <mutex>   // for call_once
 #include "paddle/string/printf.h"
 
 namespace paddle {
@@ -62,5 +65,17 @@ void Scope::DropKids() {
   kids_.clear();
 }
 
+std::once_flag feed_variable_flag;
+
+framework::Scope& GetGlobalScope() {
+  static std::unique_ptr<framework::Scope> g_scope{nullptr};
+  std::call_once(feed_variable_flag, [&]() {
+    g_scope.reset(new framework::Scope());
+    g_scope->NewVar("feed_value");
+    g_scope->NewVar("fetch_value");
+  });
+  return *(g_scope.get());
+}
+
 }  // namespace framework
 }  // namespace paddle

paddle/framework/scope.h

@@ -73,5 +73,7 @@ class Scope {
   DISABLE_COPY_AND_ASSIGN(Scope);
 };
 
+framework::Scope& GetGlobalScope();
+
 }  // namespace framework
 }  // namespace paddle

paddle/operators/CMakeLists.txt

@@ -112,7 +112,9 @@ set(DEPS_OPS
     cond_op
     cross_entropy_op
     softmax_with_cross_entropy_op
-    sum_op)
+    sum_op
+    pool_op
+    pool_with_index_op)
 
 
 op_library(recurrent_op SRCS recurrent_op.cc rnn/recurrent_op_utils.cc
@@ -121,6 +123,8 @@ 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)
+op_library(pool_op DEPS pooling)
+op_library(pool_with_index_op DEPS pooling)
 
 list(REMOVE_ITEM GENERAL_OPS ${DEPS_OPS})
 foreach(src ${GENERAL_OPS})

paddle/operators/activation_op.cc

@@ -137,6 +137,24 @@ class TanhShrinkOpMaker : public framework::OpProtoAndCheckerMaker {
   }
 };
 
+template <typename AttrType>
+class HardShrinkOpMaker : public framework::OpProtoAndCheckerMaker {
+ public:
+  HardShrinkOpMaker(framework::OpProto *proto,
+                    framework::OpAttrChecker *op_checker)
+      : OpProtoAndCheckerMaker(proto, op_checker) {
+    AddInput("X", "Input of HardShrink operator");
+    AddOutput("Y", "Output of HardShrink operator");
+    AddComment(
+        "HardShrink activation operator, "
+        "hard_shrink(x) = x if x > lambda"
+        "hard_shrink(x) = x if x < -lambda"
+        "hard_shrink(x) = 0 otherwise");
+    AddAttr<AttrType>("threshold", "The value of threshold for HardShrink")
+        .SetDefault(static_cast<AttrType>(0.5));
+  }
+};
+
 class SqrtOpMaker : public framework::OpProtoAndCheckerMaker {
  public:
   SqrtOpMaker(framework::OpProto *proto, framework::OpAttrChecker *op_checker)
@@ -188,6 +206,17 @@ class SquareOpMaker : public framework::OpProtoAndCheckerMaker {
   }
 };
 
+class SoftplusOpMaker : public framework::OpProtoAndCheckerMaker {
+ public:
+  SoftplusOpMaker(framework::OpProto *proto,
+                  framework::OpAttrChecker *op_checker)
+      : OpProtoAndCheckerMaker(proto, op_checker) {
+    AddInput("X", "Input of Softplus operator");
+    AddOutput("Y", "Output of Softplus operator");
+    AddComment("Softplus activation operator, softplus(x) = log(1 + exp(x))");
+  }
+};
+
 class SoftsignOpMaker : public framework::OpProtoAndCheckerMaker {
  public:
   SoftsignOpMaker(framework::OpProto *proto,
@@ -333,6 +362,9 @@ REGISTER_OP(log, ops::ActivationOp, ops::LogOpMaker, log_grad,
 REGISTER_OP(square, ops::ActivationOp, ops::SquareOpMaker, square_grad,
             ops::ActivationOpGrad);
 
+REGISTER_OP(softplus, ops::ActivationOp, ops::SoftplusOpMaker, softplus_grad,
+            ops::ActivationOpGrad);
+
 REGISTER_OP(softsign, ops::ActivationOp, ops::SoftsignOpMaker, softsign_grad,
             ops::ActivationOpGrad);
 
@@ -357,6 +389,9 @@ REGISTER_OP(pow, ops::ActivationOp, ops::PowOpMaker<float>, pow_grad,
 REGISTER_OP(stanh, ops::ActivationOp, ops::STanhOpMaker<float>, stanh_grad,
             ops::ActivationOpGrad);
 
+REGISTER_OP(hard_shrink, ops::ActivationOp, ops::HardShrinkOpMaker<float>,
+            hard_shrink_grad, ops::ActivationOpGrad);
+
 #define REGISTER_ACTIVATION_CPU_KERNEL(act_type, functor, grad_functor)        \
   REGISTER_OP_CPU_KERNEL(                                                      \
       act_type,                                                                \

paddle/operators/activation_op.h

@@ -199,6 +199,39 @@ struct TanhShrinkGradFunctor : public BaseActivationFunctor<T> {
   }
 };
 
+// tanhshrink(x) = x - tanh(x)
+// where tanh(x) = (exp(x) - exp(-x)) / (exp(x) + exp(-x))
+template <typename T>
+struct HardShrinkFunctor : public BaseActivationFunctor<T> {
+  float threshold;
+
+  typename BaseActivationFunctor<T>::AttrPair GetAttrs() {
+    return {{"threshold", &threshold}};
+  }
+  template <typename Device, typename X, typename Y>
+  void operator()(Device d, X x, Y y) const {
+    auto temp1 = (x < (threshold * -1)).template cast<T>().eval();
+    auto temp2 = (x > threshold).template cast<T>().eval();
+    y.device(d) = x * (temp1 + temp2);
+  }
+};
+
+template <typename T>
+struct HardShrinkGradFunctor : public BaseActivationFunctor<T> {
+  float threshold;
+
+  typename BaseActivationFunctor<T>::AttrPair GetAttrs() {
+    return {{"threshold", &threshold}};
+  }
+
+  template <typename Device, typename X, typename Y, typename dY, typename dX>
+  void operator()(Device d, X x, Y y, dY dy, dX dx) const {
+    auto temp1 = (x < (threshold * -1)).template cast<T>().eval();
+    auto temp2 = (x > threshold).template cast<T>().eval();
+    dx.device(d) = dy * (temp1 + temp2).template cast<T>();
+  }
+};
+
 // softshrink(x) = x - lambda, if x > lambda; x + lambda, if x < lambda; 0
 // otherwise
 template <typename T>
@@ -351,8 +384,6 @@ template <typename T>
 struct Relu6Functor : public BaseActivationFunctor<T> {
   float threshold;
 
-  // NOTE: Explicit hides the `BaseActivationFunctor<T>::GetAttrs`
-  // not polymorphism for speed.
   typename BaseActivationFunctor<T>::AttrPair GetAttrs() {
     return {{"threshold", &threshold}};
   }
@@ -376,6 +407,33 @@ struct Relu6GradFunctor : public BaseActivationFunctor<T> {
   }
 };
 
+// softplus(x) = log(1 + exp(x))
+// When x is a very large positive number, exp(x) may explode to inf,
+// Using trick below for numerical stability
+// https://hips.seas.harvard.edu/blog/2013/01/09/computing-log-sum-exp/
+// Then: softplus(x) = max(x, 0) + log(exp(-max(x, 0)) + exp(x - max(x, 0)))
+template <typename T>
+struct SoftplusFunctor : public BaseActivationFunctor<T> {
+  template <typename Device, typename X, typename Y>
+  void operator()(Device d, X x, Y y) {
+    auto temp = x.cwiseMax(static_cast<T>(0));  // temp = max(x, 0)
+    y.device(d) = temp + (((-temp).exp() + (x - temp).exp()).log());
+  }
+};
+
+// d(softplus(x))/dx = exp(x) / (1 + exp(x))
+// For numerical stability:
+// d(softplus(x))/dx = exp(x - max(x, 0)) / (exp(-max(x, 0)) +
+// exp(x - max(x, 0)))
+template <typename T>
+struct SoftplusGradFunctor : public BaseActivationFunctor<T> {
+  template <typename Device, typename X, typename Y, typename dY, typename dX>
+  void operator()(Device d, X x, Y y, dY dy, dX dx) {
+    auto temp = x.cwiseMax(static_cast<T>(0));  // temp = max(x, 0)
+    dx.device(d) = dy * ((x - temp).exp() / ((-temp).exp() + (x - temp).exp()));
+  }
+};
+
 // softsign(x) = x / (1 + |x|)
 template <typename T>
 struct SoftsignFunctor : public BaseActivationFunctor<T> {
@@ -551,8 +609,10 @@ struct STanhGradFunctor : public BaseActivationFunctor<T> {
   __macro(soft_relu, SoftReluFunctor, SoftReluGradFunctor);       \
   __macro(pow, PowFunctor, PowGradFunctor);                       \
   __macro(stanh, STanhFunctor, STanhGradFunctor);                 \
+  __macro(softplus, SoftplusFunctor, SoftplusGradFunctor);        \
   __macro(softsign, SoftsignFunctor, SoftsignGradFunctor);        \
   __macro(relu6, Relu6Functor, Relu6GradFunctor);                 \
   __macro(leaky_relu, LeakyReluFunctor, LeakyReluGradFunctor);    \
   __macro(tanh_shrink, TanhShrinkFunctor, TanhShrinkGradFunctor); \
-  __macro(elu, ELUFunctor, ELUGradFunctor)
+  __macro(elu, ELUFunctor, ELUGradFunctor);                       \
+  __macro(hard_shrink, HardShrinkFunctor, HardShrinkGradFunctor)

paddle/operators/feed_op.cc

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+    http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License. */
+
+#include "paddle/operators/feed_op.h"
+
+namespace paddle {
+namespace operators {
+
+class FeedOp : public framework::OperatorWithKernel {
+ public:
+  using framework::OperatorWithKernel::OperatorWithKernel;
+
+ protected:
+  void InferShape(framework::InferShapeContext* ctx) const override {
+    PADDLE_ENFORCE(ctx->HasOutput("Out"), "Output should be not null.");
+    auto& shape = ctx->Attrs().Get<std::vector<int>>("dims");
+    std::vector<int64_t> shape_int64(shape.size(), 0);
+    std::transform(shape.begin(), shape.end(), shape_int64.begin(),
+                   [](int a) { return static_cast<int64_t>(a); });
+    ctx->SetOutputDim("Out", framework::make_ddim(shape_int64));
+    // TODO(qijun): need to handle LodTensor later
+  }
+
+  framework::DataType IndicateDataType(
+      const framework::ExecutionContext& ctx) const override {
+    return static_cast<framework::DataType>(Attr<int>("dataType"));
+  }
+};
+
+class FeedOpMaker : public framework::OpProtoAndCheckerMaker {
+ public:
+  FeedOpMaker(framework::OpProto* proto, framework::OpAttrChecker* op_checker)
+      : OpProtoAndCheckerMaker(proto, op_checker) {
+    AddAttr<int>("dataType", "output data type")
+        .SetDefault(framework::DataType::FP32);
+    AddAttr<int>("col", "The col in global feed variable").SetDefault(0);
+    AddAttr<std::vector<int>>("dims", "The dimension of feed tensor.");
+    AddOutput("Out", "The output of feed op.");
+    AddComment(R"DOC(Feed data from global feed variable)DOC");
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle
+
+namespace ops = paddle::operators;
+REGISTER_OP_WITHOUT_GRADIENT(feed, ops::FeedOp, ops::FeedOpMaker);
+REGISTER_OP_CPU_KERNEL(feed, ops::FeedKernel<float>);

paddle/operators/feed_op.cu

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License. */
+
+#include "paddle/operators/feed_op.h"
+
+namespace ops = paddle::operators;
+REGISTER_OP_GPU_KERNEL(feed, ops::FeedKernel<float>);

paddle/operators/feed_op.h

+/* 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 <typename T>
+class FeedKernel : public framework::OpKernel<T> {
+ public:
+  void Compute(const framework::ExecutionContext& ctx) const override {
+    framework::Tensor* out = ctx.Output<framework::Tensor>("Out");
+    out->mutable_data<T>(ctx.GetPlace());
+    framework::Variable* g_feed_variable =
+        framework::GetGlobalScope().FindVar("feed_value");
+    const auto& tensors =
+        g_feed_variable->Get<std::vector<framework::Tensor>>();
+    int col = ctx.template Attr<int>("col");
+    PADDLE_ENFORCE_GT(tensors.size(), static_cast<size_t>(col));
+    // TODO(qijun):
+    //   check tensors[col].dims() with attribute,
+    //   except the first dimenson.
+    out->CopyFrom<T>(tensors[col], ctx.GetPlace());
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle

paddle/operators/fetch_op.cc

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+    http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License. */
+
+#include "paddle/operators/fetch_op.h"
+
+namespace paddle {
+namespace operators {
+
+class FetchOp : public framework::OperatorWithKernel {
+ public:
+  using framework::OperatorWithKernel::OperatorWithKernel;
+
+ protected:
+  void InferShape(framework::InferShapeContext* ctx) const override {
+    PADDLE_ENFORCE(ctx->HasInput("Input"), "Input should be not null.");
+  }
+
+  framework::DataType IndicateDataType(
+      const framework::ExecutionContext& ctx) const override {
+    return static_cast<framework::DataType>(Attr<int>("dataType"));
+  }
+};
+
+class FetchOpMaker : public framework::OpProtoAndCheckerMaker {
+ public:
+  FetchOpMaker(framework::OpProto* proto, framework::OpAttrChecker* op_checker)
+      : OpProtoAndCheckerMaker(proto, op_checker) {
+    AddAttr<int>("dataType", "output data type")
+        .SetDefault(framework::DataType::FP32);
+    AddAttr<int>("col", "The col in global fetch variable").SetDefault(0);
+    AddInput("Input", "The output of fetch op.");
+    AddComment(R"DOC(Fetch data to global fetch variable)DOC");
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle
+
+namespace ops = paddle::operators;
+REGISTER_OP_WITHOUT_GRADIENT(fetch, ops::FetchOp, ops::FetchOpMaker);
+REGISTER_OP_CPU_KERNEL(fetch, ops::FetchKernel<float>);

paddle/operators/fetch_op.cu

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License. */
+
+#include "paddle/operators/fetch_op.h"
+
+namespace ops = paddle::operators;
+REGISTER_OP_GPU_KERNEL(fetch, ops::FetchKernel<float>);

paddle/operators/fetch_op.h

+/* 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 <typename T>
+class FetchKernel : public framework::OpKernel<T> {
+ public:
+  void Compute(const framework::ExecutionContext& ctx) const override {
+    const framework::Tensor* input = ctx.Input<framework::Tensor>("Input");
+    framework::Variable* g_fetch_variable =
+        framework::GetGlobalScope().FindVar("fetch_value");
+    auto* tensors =
+        g_fetch_variable->GetMutable<std::vector<framework::Tensor>>();
+    int col = ctx.template Attr<int>("col");
+    if (tensors->size() < static_cast<size_t>(col + 1)) {
+      tensors->resize(col + 1);
+    }
+    PADDLE_ENFORCE_GT(tensors->size(), static_cast<size_t>(col));
+    (*tensors)[col].Resize(input->dims());
+    (*tensors)[col].mutable_data<T>(platform::CPUPlace());
+    (*tensors)[col].CopyFrom<T>(*input, platform::CPUPlace());
+    // TODO(qijun): need to handle LodTensor later
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle

paddle/operators/math/CMakeLists.txt

 if(WITH_GPU)
-    nv_library(math_function SRCS math_function.cc math_function.cu im2col.cc im2col.cu pooling.cc pooling.cu DEPS cblas device_context operator)
+    nv_library(math_function SRCS math_function.cc math_function.cu im2col.cc im2col.cu DEPS cblas device_context operator)
     nv_test(math_function_test SRCS math_function_test.cc DEPS math_function tensor)
     nv_library(softmax SRCS softmax.cc softmax.cu DEPS operator)
     nv_library(cross_entropy SRCS cross_entropy.cc cross_entropy.cu DEPS operator)
+    nv_library(pooling SRCS pooling.cc pooling.cu DEPS device_context)
+    nv_library(vol2col SRCS vol2col.cc vol2col.cu DEPS device_context)
 else()
-    cc_library(math_function SRCS math_function.cc im2col.cc pooling.cc DEPS cblas device_context operator)
+    cc_library(math_function SRCS math_function.cc im2col.cc DEPS cblas device_context operator)
     cc_test(math_function_test SRCS math_function_test.cc DEPS math_function tensor)
     cc_library(softmax SRCS softmax.cc DEPS operator)
     cc_library(cross_entropy SRCS cross_entropy.cc DEPS operator)
+    cc_library(pooling SRCS pooling.cc DEPS device_context)
+    cc_library(vol2col SRCS vol2col.cc DEPS device_context)
 endif()
 
 cc_test(im2col_test SRCS im2col_test.cc DEPS math_function tensor)
+cc_test(vol2col_test SRCS vol2col_test.cc DEPS vol2col tensor)

paddle/operators/math/vol2col.cc

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+    http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License. */
+
+#include "paddle/operators/math/vol2col.h"
+
+namespace paddle {
+namespace operators {
+namespace math {
+
+/*
+ * vol = [input_channels, input_depth, input_height, input_width]
+ * col =
+ *   [input_channels, filter_depth, filter_height, filter_width,
+ *                    output_depth, output_height, output_width]
+ */
+template <class T>
+class Vol2ColFunctor<platform::CPUPlace, T> {
+ public:
+  void operator()(const platform::DeviceContext& context,
+                  const framework::Tensor& vol, framework::Tensor& col,
+                  int stride_depth, int stride_height, int stride_width,
+                  int padding_depth, int padding_height,
+                  int padding_width) const {
+    PADDLE_ENFORCE(vol.dims().size() == 4);
+    PADDLE_ENFORCE(col.dims().size() == 7);
+
+    int input_channels = vol.dims()[0];
+    int input_depth = vol.dims()[1];
+    int input_height = vol.dims()[2];
+    int input_width = vol.dims()[3];
+    int filter_depth = col.dims()[1];
+    int filter_height = col.dims()[2];
+    int filter_width = col.dims()[3];
+    int output_depth = col.dims()[4];
+    int output_height = col.dims()[5];
+    int output_width = col.dims()[6];
+    int channels_col =
+        input_channels * filter_depth * filter_height * filter_width;
+
+    const T* vol_data = vol.data<T>();
+    T* col_data = col.data<T>();
+
+    for (int c = 0; c < channels_col; ++c) {
+      int w_offset = c % filter_width;
+      int h_offset = (c / filter_width) % filter_height;
+      int d_offset = (c / filter_width / filter_height) % filter_depth;
+      int c_in = c / filter_width / filter_height / filter_depth;
+      for (int d = 0; d < output_depth; ++d) {
+        int d_pad = d * stride_depth - padding_depth + d_offset;
+        for (int h = 0; h < output_height; ++h) {
+          int h_pad = h * stride_height - padding_height + h_offset;
+          for (int w = 0; w < output_width; ++w) {
+            int w_pad = w * stride_width - padding_width + w_offset;
+
+            int col_idx =
+                ((c * output_depth + d) * output_height + h) * output_width + w;
+            if (h_pad < 0 || h_pad >= input_height || w_pad < 0 ||
+                w_pad >= input_width || d_pad < 0 || d_pad >= input_depth) {
+              col_data[col_idx] = static_cast<T>(0);
+            } else {
+              int vol_idx =
+                  ((c_in * input_depth + d_pad) * input_height + h_pad) *
+                      input_width +
+                  w_pad;
+              col_data[col_idx] = vol_data[vol_idx];
+            }
+          }
+        }
+      }
+    }
+  }
+};
+
+/*
+ * vol = [input_channels,input_depth, input_height, input_width]
+ * col =
+ *   [input_channels, filter_depth, filter_height, filter_width,
+ *                    output_depth, output_height, output_width]
+ */
+template <class T>
+class Col2VolFunctor<platform::CPUPlace, T> {
+ public:
+  void operator()(const platform::DeviceContext& context,
+                  framework::Tensor& vol, const framework::Tensor& col,
+                  int stride_depth, int stride_height, int stride_width,
+                  int padding_depth, int padding_height,
+                  int padding_width) const {
+    PADDLE_ENFORCE(vol.dims().size() == 4);
+    PADDLE_ENFORCE(col.dims().size() == 7);
+
+    int input_channels = vol.dims()[0];
+    int input_depth = vol.dims()[1];
+    int input_height = vol.dims()[2];
+    int input_width = vol.dims()[3];
+    int filter_depth = col.dims()[1];
+    int filter_height = col.dims()[2];
+    int filter_width = col.dims()[3];
+    int output_depth = col.dims()[4];
+    int output_height = col.dims()[5];
+    int output_width = col.dims()[6];
+    int channels_col =
+        input_channels * filter_depth * filter_height * filter_width;
+
+    T* vol_data = vol.data<T>();
+    const T* col_data = col.data<T>();
+
+    for (int c = 0; c < channels_col; ++c) {
+      int w_offset = c % filter_width;
+      int h_offset = (c / filter_width) % filter_height;
+      int d_offset = (c / filter_width / filter_height) % filter_depth;
+      int cIm = c / filter_width / filter_height / filter_depth;
+      for (int d = 0; d < output_depth; ++d) {
+        int d_pad = d * stride_depth - padding_depth + d_offset;
+        for (int h = 0; h < output_height; ++h) {
+          int h_pad = h * stride_height - padding_height + h_offset;
+          for (int w = 0; w < output_width; ++w) {
+            int w_pad = w * stride_width - padding_width + w_offset;
+
+            if (h_pad >= 0 && h_pad < input_height && w_pad >= 0 &&
+                w_pad < input_width && d_pad >= 0 && d_pad < input_depth) {
+              int vol_idx =
+                  ((cIm * input_depth + d_pad) * input_height + h_pad) *
+                      input_width +
+                  w_pad;
+              int col_idx =
+                  ((c * output_depth + d) * output_height + h) * output_width +
+                  w;
+              vol_data[vol_idx] += col_data[col_idx];
+            }
+          }
+        }
+      }
+    }
+  }
+};
+
+template class Vol2ColFunctor<platform::CPUPlace, float>;
+template class Vol2ColFunctor<platform::CPUPlace, double>;
+template class Col2VolFunctor<platform::CPUPlace, float>;
+template class Col2VolFunctor<platform::CPUPlace, double>;
+
+}  // namespace math
+}  // namespace operators
+}  // namespace paddle

paddle/operators/math/vol2col.cu

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+    http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License. */
+
+#include "paddle/operators/math/vol2col.h"
+#include "paddle/platform/cuda_helper.h"
+
+namespace paddle {
+namespace operators {
+namespace math {
+
+template <class T>
+__global__ void vol2col(int num_kernels, const T* data_vol, int depth,
+                        int height, int width, int filter_depth,
+                        int filter_height, int filter_width, int stride_depth,
+                        int stride_height, int stride_width, int padding_depth,
+                        int padding_height, int padding_width, int output_detph,
+                        int output_height, int output_width, T* data_col) {
+  for (int index = blockIdx.x * blockDim.x + threadIdx.x; index < num_kernels;
+       index += blockDim.x * gridDim.x) {
+    int w_out = index % output_width;
+    int h_out = (index / output_width) % output_height;
+    int d_out = (index / output_width / output_height) % output_detph;
+    int channel_in = index / output_width / output_height / output_detph;
+    int channel_out = channel_in * filter_depth * filter_height * filter_width;
+    int w_in = w_out * stride_width - padding_width;
+    int h_in = h_out * stride_height - padding_height;
+    int d_in = d_out * stride_depth - padding_depth;
+
+    data_col += ((channel_out * output_detph + d_out) * output_height + h_out) *
+                    output_width +
+                w_out;
+    data_vol += ((channel_in * depth + d_in) * height + h_in) * width + w_in;
+    for (int k = 0; k < filter_depth; ++k) {
+      for (int i = 0; i < filter_height; ++i) {
+        for (int j = 0; j < filter_width; ++j) {
+          int d = d_in + k;
+          int h = h_in + i;
+          int w = w_in + j;
+          *data_col = (d >= 0 && d < depth && h >= 0 && h < height && w >= 0 &&
+                       w < width)
+                          ? data_vol[(k * height + i) * width + j]
+                          : 0;
+          data_col += output_detph * output_height * output_width;
+        }
+      }
+    }
+  }
+}
+
+/*
+ * im = [input_channels,intpu_depth, input_height, input_width]
+ * col =
+ *   [input_channels, filter_depth, filter_height, filter_width,
+ *                    output_depth, output_height, output_width]
+ */
+template <class T>
+class Vol2ColFunctor<platform::GPUPlace, T> {
+ public:
+  void operator()(const platform::DeviceContext& context,
+                  const framework::Tensor& vol, framework::Tensor& col,
+                  int stride_depth, int stride_height, int stride_width,
+                  int padding_depth, int padding_height,
+                  int padding_width) const {
+    PADDLE_ENFORCE(vol.dims().size() == 4);
+    PADDLE_ENFORCE(col.dims().size() == 7);
+
+    int input_channels = vol.dims()[0];
+    int input_depth = vol.dims()[1];
+    int input_height = vol.dims()[2];
+    int input_width = vol.dims()[3];
+    int filter_depth = col.dims()[1];
+    int filter_height = col.dims()[2];
+    int filter_width = col.dims()[3];
+    int output_depth = col.dims()[4];
+    int output_height = col.dims()[5];
+    int output_width = col.dims()[6];
+
+    int num_outputs =
+        input_channels * output_depth * output_height * output_width;
+
+    const int threads = 1024;
+    const int blocks = (num_outputs + 1024 - 1) / 1024;
+    vol2col<T><<<blocks, threads, 0,
+                 reinterpret_cast<const platform::CUDADeviceContext&>(context)
+                     .stream()>>>(
+        num_outputs, vol.data<T>(), input_depth, input_height, input_width,
+        filter_depth, filter_height, filter_width, stride_depth, stride_height,
+        stride_width, padding_depth, padding_height, padding_width,
+        output_depth, output_height, output_width, col.data<T>());
+  }
+};
+
+template <class T>
+__global__ void col2vol(int num_kernels, const T* data_col, int depth,
+                        int height, int width, int filter_depth,
+                        int filter_height, int filter_width, int stride_depth,
+                        int stride_height, int stride_width, int padding_depth,
+                        int padding_height, int padding_width, int output_detph,
+                        int output_height, int output_width, T* data_vol) {
+  for (int index = blockIdx.x * blockDim.x + threadIdx.x; index < num_kernels;
+       index += blockDim.x * gridDim.x) {
+    T src_val = 0;
+    int w = index % width + padding_width;
+    int h = (index / width) % height + padding_height;
+    int d = (index / width / height) % depth + padding_depth;
+    int c = index / width / height / depth;
+    // compute the start and end of the output
+    int w_col_start =
+        (w < filter_width) ? 0 : (w - filter_width) / stride_width + 1;
+    int w_col_end = min(w / stride_width + 1, output_width);
+    int h_col_start =
+        (h < filter_height) ? 0 : (h - filter_height) / stride_height + 1;
+    int h_col_end = min(h / stride_height + 1, output_height);
+    int d_col_start =
+        (d < filter_depth) ? 0 : (d - filter_depth) / stride_depth + 1;
+    int d_col_end = min(d / stride_depth + 1, output_detph);
+
+    int offset = (c * filter_depth * filter_height * filter_width +
+                  d * filter_width * filter_height + h * filter_width + w) *
+                 output_detph * output_height * output_width;
+
+    int coeff_d_col =
+        (1 - stride_depth * filter_width * filter_height * output_detph) *
+        output_height * output_width;
+    int coeff_h_col =
+        (1 - stride_height * filter_width * output_detph * output_height) *
+        output_width;
+    int coeff_w_col =
+        (1 - stride_width * output_detph * output_height * output_width);
+
+    for (int d_col = d_col_start; d_col < d_col_end; ++d_col) {
+      for (int h_col = h_col_start; h_col < h_col_end; ++h_col) {
+        for (int w_col = w_col_start; w_col < w_col_end; ++w_col) {
+          src_val += data_col[offset + d_col * coeff_d_col +
+                              h_col * coeff_h_col + w_col * coeff_w_col];
+        }
+      }
+    }
+    data_vol[index] = src_val;
+  }
+}
+
+/*
+ * im = [input_channels, input_depth, input_height, input_width]
+ * col =
+ *   [input_channels, filter_depth, filter_height, filter_width,
+ *                    output_depth, output_height, output_width]
+ */
+template <class T>
+class Col2VolFunctor<platform::GPUPlace, T> {
+ public:
+  void operator()(const platform::DeviceContext& context,
+                  framework::Tensor& vol, const framework::Tensor& col,
+                  int stride_depth, int stride_height, int stride_width,
+                  int padding_depth, int padding_height,
+                  int padding_width) const {
+    PADDLE_ENFORCE(vol.dims().size() == 4);
+    PADDLE_ENFORCE(col.dims().size() == 7);
+
+    int input_channels = vol.dims()[0];
+    int input_depth = vol.dims()[1];
+    int input_height = vol.dims()[2];
+    int input_width = vol.dims()[3];
+    int filter_depth = col.dims()[1];
+    int filter_height = col.dims()[2];
+    int filter_width = col.dims()[3];
+    int output_depth = col.dims()[4];
+    int output_height = col.dims()[5];
+    int output_width = col.dims()[6];
+
+    int num_kernels = input_channels * input_depth * input_height * input_width;
+
+    const int threads = 1024;
+    const int blocks = (num_kernels + 1024 - 1) / 1024;
+
+    col2vol<T><<<blocks, threads, 0,
+                 reinterpret_cast<const platform::CUDADeviceContext&>(context)
+                     .stream()>>>(
+        num_kernels, col.data<T>(), input_depth, input_height, input_width,
+        filter_depth, filter_height, filter_width, stride_depth, stride_height,
+        stride_width, padding_depth, padding_height, padding_width,
+        output_depth, output_height, output_width, vol.data<T>());
+  }
+};
+
+template class Vol2ColFunctor<platform::GPUPlace, float>;
+template class Vol2ColFunctor<platform::GPUPlace, double>;
+template class Col2VolFunctor<platform::GPUPlace, float>;
+template class Col2VolFunctor<platform::GPUPlace, double>;
+
+}  // namespace math
+}  // namespace operators
+}  // namespace paddle

paddle/operators/math/vol2col.h

+/* 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/device_context.h"
+
+namespace paddle {
+namespace operators {
+namespace math {
+/*
+ * \brief Converts the feature data of four dimensions(CDHW) into a colData of
+ *        seven dimensions in the Vol2ColFunctor calculation,
+ *        And in the Col2VolFunctor calculation, it is reversed.
+ *
+ * \param volData   Vol data.
+ * \param volShape  The shape of volData,
+ *                 [input_channels, input_depth, input_height, input_width].
+ * \param colData  Column data.
+ * \param colShape The shape of colData.
+ *
+ * The shape of colData is:
+ * [input_channels, filter_depth, filter_height, filter_width, output_depth,
+ * output_height, output_width]
+ * So, it is easy to reshape into a convolution matrix for convolution
+ * calculation based on matrix multiplication.
+ * The shape of convolution matrix is [height, width], where the height is equal
+ * input_channels * filter_depth * filter_height * filter_width, and the width
+ * is equal output_depth * output_height * output_width.
+ *
+ * Reshape:
+ *     shape of colData           shape of convolution matrix
+ *     [input_channels,
+ *      filter_depth,
+ *      filter_height,
+ *      filter_width,      ======>      [height, width]
+ *      output_depth,
+ *      output_height,
+ *      output_width]
+ *
+ * \note The caller needs to ensure that volShape.inputChannels is equal to
+ *       colShape.inputChannels.
+ */
+template <typename Place, typename T>
+class Vol2ColFunctor {
+ public:
+  void operator()(const platform::DeviceContext& context,
+                  const framework::Tensor& vol, framework::Tensor& col,
+                  int stride_depth, int stride_height, int stride_width,
+                  int padding_depth, int padding_height,
+                  int padding_width) const;
+};
+
+template <typename Place, typename T>
+class Col2VolFunctor {
+ public:
+  void operator()(const platform::DeviceContext& context,
+                  framework::Tensor& vol, const framework::Tensor& col,
+                  int stride_depth, int stride_height, int stride_width,
+                  int padding_depth, int padding_height,
+                  int padding_width) const;
+};
+
+}  // namespace math
+}  // namespace operators
+}  // namespace paddle

paddle/operators/math/vol2col_test.cc

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+    http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License. */
+
+#include "paddle/operators/math/vol2col.h"
+#include <gtest/gtest.h>
+#include <iostream>
+
+template <typename Place>
+void testVol2col() {
+  paddle::framework::Tensor input;
+  paddle::framework::Tensor input_tmp;
+  paddle::framework::Tensor output;
+  paddle::framework::Tensor output_tmp;
+
+  auto* place = new Place();
+  paddle::platform::DeviceContext* context;
+  if (paddle::platform::is_cpu_place(*place)) {
+    context =
+        new paddle::platform::CPUDeviceContext(paddle::platform::CPUPlace());
+  } else {
+#ifdef PADDLE_WITH_CUDA
+    context =
+        new paddle::platform::CUDADeviceContext(paddle::platform::GPUPlace());
+#else
+    PADDLE_THROW("no GPU support");
+#endif  // PADDLE_WITH_CUDA
+  }
+
+  /**
+   * input = [[0, 1, 2,
+   *          3, 4, 5]
+   *          [6, 7, 8,
+   *          9, 10, 11]]
+   *
+   * output = [0, 1
+   *           1, 2
+   *           3, 4
+   *           4, 5
+   *           6, 7
+   *           7, 8
+   *           9, 10
+   *           10, 11]
+   *
+   * col2vol = [[0, 2, 2,
+   *             3, 8, 5]
+   *            [6, 14, 8,
+   *             9, 20, 11]]
+   *
+   */
+  int input_depth = 2;
+  int input_height = 2;
+  int input_width = 3;
+  int filter_size = 2;
+  int stride = 1;
+  int padding = 0;
+  int output_depth = (input_depth - filter_size + 2 * padding) / stride + 1;
+  int output_height = (input_height - filter_size + 2 * padding) / stride + 1;
+  int output_width = (input_width - filter_size + 2 * padding) / stride + 1;
+
+  // Vol2Col test
+  float* input_ptr =
+      input_tmp.mutable_data<float>({1, input_depth, input_height, input_width},
+                                    paddle::platform::CPUPlace());
+  float arr[12] = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11};
+  memcpy(input_ptr, arr, 12 * sizeof(float));
+
+  if (paddle::platform::is_cpu_place(*place)) {
+    input = input_tmp;
+  } else {
+    input.CopyFrom<float>(input_tmp, *place);
+  }
+  output.mutable_data<float>({1, filter_size, filter_size, filter_size,
+                              output_depth, output_height, output_width},
+                             *place);
+
+  paddle::operators::math::Vol2ColFunctor<Place, float> vol2col;
+  vol2col(*context, input, output, stride, stride, stride, padding, padding,
+          padding);
+
+  float vol_2_col[] = {0, 1, 1, 2, 3, 4, 4, 5, 6, 7, 7, 8, 9, 10, 10, 11};
+  float* out_cfo_ptr;
+  if (paddle::platform::is_cpu_place(*place)) {
+    out_cfo_ptr = output.data<float>();
+  } else {
+    output_tmp.CopyFrom<float>(output, paddle::platform::CPUPlace());
+    out_cfo_ptr = output_tmp.data<float>();
+  }
+
+  for (int i = 0; i < 16; ++i) {
+    EXPECT_EQ(out_cfo_ptr[i], vol_2_col[i]);
+  }
+
+  // Col2Vol test
+  float col_2_vol[] = {0, 2, 2, 3, 8, 5, 6, 14, 8, 9, 20, 11};
+  memset(input_ptr, 0, 12 * sizeof(float));
+  if (paddle::platform::is_cpu_place(*place)) {
+    input = input_tmp;
+  } else {
+    input.CopyFrom<float>(input_tmp, *place);
+  }
+
+  paddle::operators::math::Col2VolFunctor<Place, float> col2vol;
+  col2vol(*context, input, output, stride, stride, stride, padding, padding,
+          padding);
+
+  float* in_ptr;
+  if (paddle::platform::is_cpu_place(*place)) {
+    in_ptr = input.data<float>();
+  } else {
+    input_tmp.CopyFrom<float>(input, paddle::platform::CPUPlace());
+    in_ptr = input_tmp.data<float>();
+  }
+
+  for (int i = 0; i < 12; ++i) {
+    EXPECT_EQ(in_ptr[i], col_2_vol[i]);
+  }
+}
+
+TEST(math, vol2col) {
+  testVol2col<paddle::platform::CPUPlace>();
+#ifdef PADDLE_WITH_CUDA
+  testVol2col<paddle::platform::GPUPlace>();
+#endif  // PADDLE_WITH_CUDA
+}

paddle/operators/pool_op.cc

@@ -22,157 +22,181 @@ int OutputSizePool(int input_size, int filter_size, int padding, int stride) {
   return output_size;
 }
 
-class PoolOp : public framework::OperatorWithKernel {
- public:
-  using framework::OperatorWithKernel::OperatorWithKernel;
-
- protected:
-  void InferShape(framework::InferShapeContext *ctx) const override {
-    PADDLE_ENFORCE(ctx->HasInput("X"),
-                   "X(Input) of Pooling should not be null.");
-    PADDLE_ENFORCE(ctx->HasOutput("Out"),
-                   "Out(Output) of Pooling should not be null.");
-
-    auto in_x_dims = ctx->GetInputDim("X");
-
-    std::string pooling_type = ctx->Attrs().Get<std::string>("poolingType");
-    std::vector<int> ksize = ctx->Attrs().Get<std::vector<int>>("ksize");
-    std::vector<int> strides = ctx->Attrs().Get<std::vector<int>>("strides");
-    std::vector<int> paddings = ctx->Attrs().Get<std::vector<int>>("paddings");
-
-    PADDLE_ENFORCE(pooling_type == "max" || pooling_type == "avg",
-                   "pooling_type should be 'max' or 'avg'");
-    PADDLE_ENFORCE(in_x_dims.size() == 4 || in_x_dims.size() == 5,
-                   "Pooling intput should be 4-D or 5-D");
-
-    if (ctx->Attrs().Get<bool>("globalPooling")) {
-      ksize.resize(static_cast<size_t>(in_x_dims.size()) - 2);
-      for (size_t i = 0; i < ksize.size(); ++i)
-        ksize[i] = static_cast<int>(in_x_dims[i + 2]);
-    }
-
-    PADDLE_ENFORCE(in_x_dims.size() - ksize.size() == 2U,
-                   "Input size and Pooling size should be consistent.");
-    PADDLE_ENFORCE(ksize.size() == 2 || ksize.size() == 3,
-                   "Pooling size should be 2 elements. or 3 elements.");
-    PADDLE_ENFORCE_EQ(ksize.size(), strides.size(),
-                      "strides size and pooling size should be the same.");
-    PADDLE_ENFORCE_EQ(ksize.size(), paddings.size(),
-                      "paddings size and pooling size should be the same.");
-
-    std::vector<int64_t> output_shape({in_x_dims[0], in_x_dims[1]});
-    for (size_t i = 0; i < ksize.size(); ++i) {
-      output_shape.push_back(
-          OutputSizePool(in_x_dims[i + 2], ksize[i], paddings[i], strides[i]));
-    }
-    ctx->SetOutputDim("Out", framework::make_ddim(output_shape));
+void PoolOp::InferShape(framework::InferShapeContext *ctx) const {
+  PADDLE_ENFORCE(ctx->HasInput("X"), "X(Input) of Pooling should not be null.");
+  PADDLE_ENFORCE(ctx->HasOutput("Out"),
+                 "Out(Output) of Pooling should not be null.");
+
+  auto in_x_dims = ctx->GetInputDim("X");
+
+  std::string pooling_type = ctx->Attrs().Get<std::string>("poolingType");
+  std::vector<int> ksize = ctx->Attrs().Get<std::vector<int>>("ksize");
+  std::vector<int> strides = ctx->Attrs().Get<std::vector<int>>("strides");
+  std::vector<int> paddings = ctx->Attrs().Get<std::vector<int>>("paddings");
+
+  PADDLE_ENFORCE(in_x_dims.size() == 4 || in_x_dims.size() == 5,
+                 "Pooling intput should be 4-D or 5-D tensor.");
+
+  if (ctx->Attrs().Get<bool>("globalPooling")) {
+    ksize.resize(static_cast<size_t>(in_x_dims.size()) - 2);
+    for (size_t i = 0; i < ksize.size(); ++i)
+      ksize[i] = static_cast<int>(in_x_dims[i + 2]);
   }
-};
-
-class PoolOpGrad : public framework::OperatorWithKernel {
- public:
-  using framework::OperatorWithKernel::OperatorWithKernel;
-
- protected:
-  void InferShape(framework::InferShapeContext *ctx) const override {
-    PADDLE_ENFORCE(ctx->HasInput("X"),
-                   "X(Input) of Pooling should not be null.");
-    PADDLE_ENFORCE(ctx->HasOutput(framework::GradVarName("X")),
-                   "Input@Grad of Pooling should not be null.");
-    ctx->SetOutputDim(framework::GradVarName("X"), ctx->GetInputDim("X"));
+
+  PADDLE_ENFORCE(in_x_dims.size() - ksize.size() == 2U,
+                 "Input size and pooling size should be consistent.");
+  PADDLE_ENFORCE_EQ(ksize.size(), strides.size(),
+                    "Strides size and pooling size should be the same.");
+  PADDLE_ENFORCE_EQ(ksize.size(), paddings.size(),
+                    "Paddings size and pooling size should be the same.");
+
+  std::vector<int64_t> output_shape({in_x_dims[0], in_x_dims[1]});
+  for (size_t i = 0; i < ksize.size(); ++i) {
+    output_shape.push_back(
+        OutputSizePool(in_x_dims[i + 2], ksize[i], paddings[i], strides[i]));
   }
-};
-
-class Pool2dOpMaker : public framework::OpProtoAndCheckerMaker {
- public:
-  Pool2dOpMaker(framework::OpProto *proto, framework::OpAttrChecker *op_checker)
-      : OpProtoAndCheckerMaker(proto, op_checker) {
-    AddInput(
-        "X",
-        "The input tensor of pooling operator. "
-        "The format of input tensor is NCHW. Where N is batch size, C is the "
-        "number of channels, H and W is the height and width of feature.");
-    AddOutput("Out",
-              "The output tensor of pooling operator."
-              "The format of output tensor is also NCHW.");
-
-    AddAttr<std::string>("poolingType",
-                         "PoolingType of pooling operator."
-                         "Str constant equal to 'max' or 'avg'.")
-        .InEnum({"max", "avg"});
-    AddAttr<std::vector<int>>(
-        "ksize",
-        "Pooling size(depth, height, width) of pooling operator."
-        "If globalPooling = true, ksize is ignored and need not be "
-        "specified.");  // TODO(Add checker)
-    AddAttr<bool>(
-        "globalPooling",
-        "Whether to use the globalPooling."
-        "Bool constant equal to false or true."
-        "Default false."
-        "If globalPooling = true, ksize is ignored and need not be specified.")
-        .SetDefault(false);
-    AddAttr<std::vector<int>>("strides",
-                              "Strides(height, width) of pooling operator."
-                              "Default {1,1}")
-        .SetDefault({1, 1});  // TODO(Add checker)
-    AddAttr<std::vector<int>>("paddings",
-                              "Paddings(height, width) of pooling operator."
-                              "Default {0,0}.")
-        .SetDefault({0, 0});  // TODO(Add checker)
-    AddComment(R"DOC(
+  ctx->SetOutputDim("Out", framework::make_ddim(output_shape));
+}
+
+void PoolOpGrad::InferShape(framework::InferShapeContext *ctx) const {
+  PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) must not be null.");
+  PADDLE_ENFORCE(ctx->HasOutput(framework::GradVarName("X")),
+                 "Input(X@GRAD) should not be null.");
+  ctx->SetOutputDim(framework::GradVarName("X"), ctx->GetInputDim("X"));
+}
+
+Pool2dOpMaker::Pool2dOpMaker(framework::OpProto *proto,
+                             framework::OpAttrChecker *op_checker)
+    : OpProtoAndCheckerMaker(proto, op_checker) {
+  AddInput(
+      "X",
+      "(Tensor) The input tensor of pooling operator. "
+      "The format of input tensor is NCHW. Where N is batch size, C is the "
+      "number of channels, H and W is the height and width of feature.");
+  AddOutput("Out",
+            "(Tensor) The output tensor of pooling operator."
+            "The format of output tensor is also NCHW."
+            "Where N is batch size, C is "
+            "the number of channels, H and W is the height and "
+            "width of feature.");
+
+  AddAttr<std::string>("poolingType",
+                       "PoolingType of pooling operator."
+                       "Str constant equal to 'max' or 'avg'.")
+      .InEnum({"max", "avg"});
+
+  AddAttr<std::vector<int>>(
+      "ksize",
+      "The pooling window size(height, width) of pooling operator."
+      "If globalPooling = true, ksize is ignored and need not be "
+      "specified.");  // TODO(Chengduo): Add checker. (Currently,
+                      // TypedAttrChecker don't support vector type.)
+  AddAttr<bool>(
+      "globalPooling",
+      "Whether to use the globalPooling."
+      "Bool constant equal to false or true."
+      "Default false."
+      "If globalPooling = true, ksize is ignored and need not be specified.")
+      .SetDefault(false);
+  AddAttr<std::vector<int>>("strides",
+                            "The strides(height, width) of pooling window."
+                            "Default {1,1}.")
+      .SetDefault({1, 1});  // TODO(Chengduo): Add checker. (Currently,
+                            // TypedAttrChecker don't support vector type.)
+  AddAttr<std::vector<int>>("paddings",
+                            "The zero padding(height, width) size on both sides"
+                            "Default {0,0}.")
+      .SetDefault({0, 0});  // TODO(Chengduo): Add checker. (Currently,
+                            // TypedAttrChecker don't support vector type.)
+
+  AddComment(R"DOC(
 The pooling2d operation calculates the output based on
 the input, poolingType and ksize, strides, paddings parameters.
+Input(X) and output(Out) are in NCHW format. Where N is batch size, C is the
+number of channels, H and W is the height and width of feature.
+Parameters(ksize, strides, paddings) are two elements.
+These two elements represent height and width, respectively.
+The input(X) size and output(Out) size may be different.
+
+Example:
+  Input:
+       X shape: (N, C, H_in, W_in)
+  Output:
+       Out shape: (N, C, H_out, W_out)
+       Mask shape: (N, C, H_out, W_out)
+  where
+       H_out = (H_in - ksize[0] + 2 * paddings[0]) / strides[0] + 1;
+       W_out = (W_in - ksize[1] + 2 * paddings[1]) / strides[1] + 1;
 )DOC");
-  }
-};
-
-class Pool3dOpMaker : public framework::OpProtoAndCheckerMaker {
- public:
-  Pool3dOpMaker(framework::OpProto *proto, framework::OpAttrChecker *op_checker)
-      : OpProtoAndCheckerMaker(proto, op_checker) {
-    AddInput("X",
-             "The input tensor of pooling operator. "
-             "The format of input tensor is NCDHW. Where N is batch size, C is "
-             "the "
-             "number of channels, D, H and W is the depth, height and width of "
-             "feature.");
-    AddOutput("Out",
-              "The output tensor of pooling operator."
-              "The format of output tensor is also NCDHW.");
-
-    AddAttr<std::string>("poolingType",
-                         "PoolingType of pooling operator."
-                         "str constant equal to 'max' or 'avg'.")
-        .InEnum({"max", "avg"});
-    AddAttr<std::vector<int>>(
-        "ksize",
-        "Pooling size(depth, height, width) of pooling operator."
-        "If globalPooling = true, ksize is ignored and need not be "
-        "specified.");  // TODO(Add checker)
-    AddAttr<bool>(
-        "globalPooling",
-        "Whether to use the globalPooling."
-        "Bool constant equal to false or true."
-        "Default false."
-        "If globalPooling = true, ksize is ignored and need not be specified.")
-        .SetDefault(false);
-    AddAttr<std::vector<int>>(
-        "strides",
-        "Strides(depth, height, width) of pooling operator."
-        "Default {1,1,1}.")
-        .SetDefault({1, 1, 1});  // TODO(Add checker)
-    AddAttr<std::vector<int>>(
-        "paddings",
-        "Paddings(depth, height, width) of pooling operator."
-        "Default {0,0,0}.")
-        .SetDefault({0, 0, 0});  // TODO(Add checker)
-    AddComment(R"DOC(
+}
+
+Pool3dOpMaker::Pool3dOpMaker(framework::OpProto *proto,
+                             framework::OpAttrChecker *op_checker)
+    : OpProtoAndCheckerMaker(proto, op_checker) {
+  AddInput(
+      "X",
+      "(Tensor) The input tensor of pooling operator. "
+      "The format of input tensor is NCDHW. Where N is batch size, C is "
+      "the number of channels, D, H and W is the depth, height and width of "
+      "feature.");
+  AddOutput("Out",
+            "(Tensor) The output tensor of pooling operator."
+            "The format of output tensor is also NCDHW."
+            "Where N is batch size, C is "
+            "the number of channels, D, H and W is the depth, height and "
+            "width of feature.");
+
+  AddAttr<std::string>("poolingType",
+                       "PoolingType of pooling operator."
+                       "Str constant equal to 'max' or 'avg'.")
+      .InEnum({"max", "avg"});
+
+  AddAttr<std::vector<int>>(
+      "ksize",
+      "The pooling window size(depth, height, width) of pooling operator."
+      "If globalPooling = true, ksize is ignored and need not be "
+      "specified.");  // TODO(Chengduo): Add checker. (Currently,
+                      // TypedAttrChecker don't support vector type.)
+  AddAttr<bool>(
+      "globalPooling",
+      "Whether to use the globalPooling."
+      "Bool constant equal to false or true."
+      "Default false."
+      "If globalPooling = true, ksize is ignored and need not be specified.")
+      .SetDefault(false);
+  AddAttr<std::vector<int>>("strides",
+                            "Strides(depth, height, width) of pooling operator."
+                            "Default {1,1,1}.")
+      .SetDefault({1, 1, 1});  // TODO(Chengduo): Add checker. (Currently,
+                               // TypedAttrChecker don't support vector type.)
+  AddAttr<std::vector<int>>(
+      "paddings",
+      "Paddings(depth, height, width) of pooling operator."
+      "Default {0,0,0}.")
+      .SetDefault({0, 0, 0});  // TODO(Chengduo): Add checker. (Currently,
+                               // TypedAttrChecker don't support vector type.)
+
+  AddComment(R"DOC(
 The pooling3d operation calculates the output based on
 the input, poolingType and ksize, strides, paddings parameters.
+Input(X) and output(Out) are in NCDHW format. Where N is batch
+size, C is the number of channels, D, H and W is the depth, height and
+width of feature. Parameters(ksize, strides, paddings) are three elements.
+These three elements represent depth, height and width, respectively.
+The input(X) size and output(Out) size may be different.
+
+Example:
+  Input:
+       X shape: (N, C, D_in, H_in, W_in)
+  Output:
+       Out shape: (N, C, D_out, H_out, W_out)
+       Mask shape: (N, C, D_out, H_out, W_out)
+  where
+       D_out = (D_in - ksize[0] + 2 * paddings[0]) / strides[0] + 1;
+       H_out = (H_in - ksize[1] + 2 * paddings[1]) / strides[1] + 1;
+       W_out = (W_in - ksize[2] + 2 * paddings[2]) / strides[2] + 1;
 )DOC");
-  }
-};
+}
 }  // namespace operators
 }  // namespace paddle
 

paddle/operators/pool_op.h

@@ -24,6 +24,34 @@ namespace operators {
 
 using Tensor = framework::Tensor;
 
+class PoolOp : public framework::OperatorWithKernel {
+ public:
+  using framework::OperatorWithKernel::OperatorWithKernel;
+
+ protected:
+  void InferShape(framework::InferShapeContext* ctx) const override;
+};
+
+class PoolOpGrad : public framework::OperatorWithKernel {
+ public:
+  using framework::OperatorWithKernel::OperatorWithKernel;
+
+ protected:
+  void InferShape(framework::InferShapeContext* ctx) const override;
+};
+
+class Pool2dOpMaker : public framework::OpProtoAndCheckerMaker {
+ public:
+  Pool2dOpMaker(framework::OpProto* proto,
+                framework::OpAttrChecker* op_checker);
+};
+
+class Pool3dOpMaker : public framework::OpProtoAndCheckerMaker {
+ public:
+  Pool3dOpMaker(framework::OpProto* proto,
+                framework::OpAttrChecker* op_checker);
+};
+
 template <typename Place, typename T>
 class PoolKernel : public framework::OpKernel<T> {
  public:

paddle/operators/pool_with_index_op.cc

@@ -43,7 +43,7 @@ class MaxPoolWithIndexOp : public framework::OperatorWithKernel {
     std::vector<int> paddings = ctx->Attrs().Get<std::vector<int>>("paddings");
 
     PADDLE_ENFORCE(in_x_dims.size() == 4 || in_x_dims.size() == 5,
-                   "Pooling intput should be 4-D or 5-D");
+                   "Pooling intput should be 4-D or 5-D tensor.");
 
     if (ctx->Attrs().Get<bool>("globalPooling")) {
       ksize.resize(static_cast<size_t>(in_x_dims.size()) - 2);
@@ -52,7 +52,7 @@ class MaxPoolWithIndexOp : public framework::OperatorWithKernel {
     }
 
     PADDLE_ENFORCE(in_x_dims.size() - ksize.size() == 2U,
-                   "Intput size and pooling size should be consistent.");
+                   "Input size and pooling size should be consistent.");
     PADDLE_ENFORCE_EQ(ksize.size(), strides.size(),
                       "Strides size and pooling size should be the same.");
     PADDLE_ENFORCE_EQ(ksize.size(), paddings.size(),
@@ -74,6 +74,7 @@ class MaxPoolWithIndexOpGrad : public framework::OperatorWithKernel {
 
  protected:
   void InferShape(framework::InferShapeContext *ctx) const override {
+    PADDLE_ENFORCE(ctx->HasInput("Mask"), "Input(Mask) must not be null.");
     PADDLE_ENFORCE(ctx->HasInput("X"), "Input(X) must not be null.");
     PADDLE_ENFORCE(ctx->HasOutput(framework::GradVarName("X")),
                    "Input(X@GRAD) should not be null.");
@@ -88,17 +89,17 @@ class MaxPool2dWithIndexOpMaker : public framework::OpProtoAndCheckerMaker {
       : OpProtoAndCheckerMaker(proto, op_checker) {
     AddInput(
         "X",
-        "The input tensor of pooling operator. "
+        "(Tensor) The input tensor of pooling operator. "
         "The format of input tensor is NCHW. Where N is batch size, C is the "
         "number of channels, H and W is the height and width of image.");
     AddOutput("Out",
-              "The output tensor of pooling operator."
+              "(Tensor) The output tensor of pooling operator."
               "The format of output tensor is also NCHW."
               "Where N is batch size, C is "
               "the number of channels, H and W is the height and "
               "width of image.");
     AddOutput("Mask",
-              "The Mask tensor of pooling operator."
+              "(Tensor) The Mask tensor of pooling operator."
               "The format of output tensor is also NCHW."
               "Where N is batch size, C is the number of channels, H and W "
               "is the height and width of image."
@@ -106,7 +107,7 @@ class MaxPool2dWithIndexOpMaker : public framework::OpProtoAndCheckerMaker {
 
     AddAttr<std::vector<int>>(
         "ksize",
-        "The pooling size(height, width) of pooling operator."
+        "The pooling window size(height, width) of pooling operator."
         "If globalPooling = true, ksize is ignored and need not be "
         "specified.");  // TODO(Chengduo): Add checker. (Currently,
                         // TypedAttrChecker don't support vector type.)
@@ -118,13 +119,14 @@ class MaxPool2dWithIndexOpMaker : public framework::OpProtoAndCheckerMaker {
         "If globalPooling = true, ksize is ignored and need not be specified.")
         .SetDefault(false);
     AddAttr<std::vector<int>>("strides",
-                              "Strides(height, width) of pooling operator."
+                              "The strides(height, width) of pooling window."
                               "Default {1,1}.")
         .SetDefault({1, 1});  // TODO(Chengduo): Add checker. (Currently,
                               // TypedAttrChecker don't support vector type.)
-    AddAttr<std::vector<int>>("paddings",
-                              "Paddings(height, width) of pooling operator."
-                              "Default {0,0}.")
+    AddAttr<std::vector<int>>(
+        "paddings",
+        "The zero padding(height, width) size on both sides"
+        "Default {0,0}.")
         .SetDefault({0, 0});  // TODO(Chengduo): Add checker. (Currently,
                               // TypedAttrChecker don't support vector type.)
 
@@ -135,6 +137,17 @@ output(Out, Mask) are in NCHW format. Where N is batch size, C is the
 number of channels, H and W is the height and width of feature.
 Parameters(ksize, strides, paddings) are two elements.
 These two elements represent height and width, respectively.
+The input(X) size and output(Out, Mask) size may be different.
+
+Example:
+  Input:
+       X shape: (N, C, H_in, W_in)
+  Output:
+       Out shape: (N, C, H_out, W_out)
+       Mask shape: (N, C, H_out, W_out)
+  where
+       H_out = (H_in - ksize[0] + 2 * paddings[0]) / strides[0] + 1;
+       W_out = (W_in - ksize[1] + 2 * paddings[1]) / strides[1] + 1;
 )DOC");
   }
 };
@@ -146,18 +159,18 @@ class MaxPool3dWithIndexOpMaker : public framework::OpProtoAndCheckerMaker {
       : OpProtoAndCheckerMaker(proto, op_checker) {
     AddInput(
         "X",
-        "The input tensor of pooling operator. "
+        "(Tensor) The input tensor of pooling operator. "
         "The format of input tensor is NCDHW. Where N is batch size, C is "
         "the number of channels, D, H and W is the depth, height and width of "
         "image.");
     AddOutput("Out",
-              "The output tensor of pooling operator."
+              "(Tensor) The output tensor of pooling operator."
               "The format of output tensor is also NCDHW."
               "Where N is batch size, C is "
               "the number of channels, D, H and W is the depth, height and "
               "width of image.");
     AddOutput("Mask",
-              "The Mask tensor of pooling operator."
+              "(Tensor) The Mask tensor of pooling operator."
               "The format of output tensor is also NCDHW."
               "Where N is batch size, C is the number of channels, D, H and W "
               "is the depth, height and width of image."
@@ -165,7 +178,7 @@ class MaxPool3dWithIndexOpMaker : public framework::OpProtoAndCheckerMaker {
 
     AddAttr<std::vector<int>>(
         "ksize",
-        "The pooling size(depth, height, width) of pooling operator."
+        "The pooling window size(depth, height, width) of pooling operator."
         "If globalPooling = true, ksize is ignored and need not be "
         "specified.");  // TODO(Chengduo): Add checker. (Currently,
                         // TypedAttrChecker don't support vector type.)
@@ -196,6 +209,18 @@ Input(X) and output(Out, Mask) are in NCDHW format. Where N is batch
 size, C is the number of channels, D, H and W is the depth, height and
 width of feature. Parameters(ksize, strides, paddings) are three elements.
 These three elements represent depth, height and width, respectively.
+The input(X) size and output(Out, Mask) size may be different.
+
+Example:
+  Input:
+       X shape: (N, C, D_in, H_in, W_in)
+  Output:
+       Out shape: (N, C, D_out, H_out, W_out)
+       Mask shape: (N, C, D_out, H_out, W_out)
+  where
+       D_out = (D_in - ksize[0] + 2 * paddings[0]) / strides[0] + 1;
+       H_out = (H_in - ksize[1] + 2 * paddings[1]) / strides[1] + 1;
+       W_out = (W_in - ksize[2] + 2 * paddings[2]) / strides[2] + 1;
 )DOC");
   }
 };

paddle/platform/gpu_info.cc

@@ -43,6 +43,8 @@ int GetCurrentDeviceId() {
 }
 
 void SetDeviceId(int id) {
+  // TODO(qijun): find a better way to cache the cuda device count
+  PADDLE_ENFORCE_LT(id, GetCUDADeviceCount(), "id must less than GPU count");
   PADDLE_ENFORCE(cudaSetDevice(id),
                  "cudaSetDevice failed in paddle::platform::SetDeviceId");
 }

python/paddle/v2/framework/tests/test_activation_op.py

@@ -78,6 +78,26 @@ class TestTanhShrink(OpTest):
         self.check_grad(['X'], 'Y', max_relative_error=0.008)
 
 
+class TestHardShrink(OpTest):
+    def setUp(self):
+        self.op_type = "hard_shrink"
+        x = np.random.uniform(-1, 1, [4, 4]).astype("float32")
+        threshold = 0.5
+
+        self.inputs = {'X': x}
+        self.attrs = {'lambda': threshold}
+
+        t = np.copy(x)
+        t[(t >= -threshold) & (t <= threshold)] = 0
+        self.outputs = {'Y': t}
+
+    def test_check_output(self):
+        self.check_output()
+
+    def test_check_grad(self):
+        self.check_grad(['X'], 'Y', max_relative_error=0.005)
+
+
 class TestSoftShrink(OpTest):
     def setUp(self):
         self.op_type = "softshrink"
@@ -311,6 +331,21 @@ class TestSTanh(OpTest):
         self.check_grad(['X'], 'Y', max_relative_error=0.007)
 
 
+class TestSoftplus(OpTest):
+    def setUp(self):
+        self.op_type = "softplus"
+        self.inputs = {
+            'X': np.random.uniform(-1, 1, [11, 17]).astype("float32")
+        }
+        self.outputs = {'Y': np.log(1 + np.exp(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 TestSoftsign(OpTest):
     def setUp(self):
         self.op_type = "softsign"