From f196ad0210aadb715c12cafea2798ca235d84940 Mon Sep 17 00:00:00 2001
From: Liu Yiqun <liuyiqun01@baidu.com>
Date: Tue, 5 Sep 2017 11:33:40 +0000
Subject: [PATCH 01/34] Port fully connected operator, the FCOp c++
 implementation and python unittest.

---
 paddle/operators/CMakeLists.txt               |   5 +-
 paddle/operators/fc_op.cc                     | 107 ++++++++++++++++++
 paddle/operators/scale_op.cc                  |   1 +
 paddle/pybind/CMakeLists.txt                  |   2 +-
 paddle/pybind/pybind.cc                       |   1 +
 .../paddle/v2/framework/tests/CMakeLists.txt  |   1 +
 .../paddle/v2/framework/tests/test_fc_op.py   |  30 +++++
 7 files changed, 145 insertions(+), 2 deletions(-)
 create mode 100644 paddle/operators/fc_op.cc
 create mode 100644 python/paddle/v2/framework/tests/test_fc_op.py

diff --git a/paddle/operators/CMakeLists.txt b/paddle/operators/CMakeLists.txt
index e5efcccb0e..2a8beda2c8 100644
--- a/paddle/operators/CMakeLists.txt
+++ b/paddle/operators/CMakeLists.txt
@@ -47,17 +47,20 @@ endfunction()
 add_subdirectory(math)
 
 list(REMOVE_ITEM GENERAL_OPS
+     fc_op
      net_op
      minus_op
      mul_op
      recurrent_op
      scale_op)
 
+op_library(fc_op SRCS fc_op.cc
+           DEPS mul_op rowwise_add_op scale_op softmax_op sigmoid_op)
 op_library(net_op SRCS net_op.cc)
 op_library(minus_op SRCS minus_op.cc minus_op.cu DEPS scale_op)
 op_library(mul_op SRCS mul_op.cc mul_op.cu DEPS math_function)
 op_library(recurrent_op SRCS recurrent_op.cc rnn/recurrent_op_utils.cc 
-  DEPS framework_proto tensor operator net_op)
+           DEPS framework_proto tensor operator net_op)
 op_library(scale_op SRCS scale_op.cc scale_op.cu DEPS net_op)
 
 foreach(src ${GENERAL_OPS})
diff --git a/paddle/operators/fc_op.cc b/paddle/operators/fc_op.cc
new file mode 100644
index 0000000000..ebf8908db7
--- /dev/null
+++ b/paddle/operators/fc_op.cc
@@ -0,0 +1,107 @@
+/* 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/op_registry.h"
+#include "paddle/operators/net_op.h"
+
+namespace paddle {
+namespace operators {
+
+class FCOp : public NetOp {
+ public:
+  FCOp(const std::string &type, const framework::VariableNameMap &inputs,
+       const framework::VariableNameMap &outputs,
+       const framework::AttributeMap &attrs)
+      : NetOp(type, inputs, outputs, attrs) {
+    AppendOp(framework::OpRegistry::CreateOp(
+        "mul", {{"X", {Input("X")}}, {"Y", {Input("W")}}},
+        {{"Out", {Output("mul_out")}}}, {}));
+    auto b = Input("b");
+    if (b != framework::kEmptyVarName) {
+      AppendOp(framework::OpRegistry::CreateOp(
+          "rowwise_add", {{"X", {Output("mul_out")}}, {"b", {Input("b")}}},
+          {{"Out", {Output("mul_out")}}}, {}));
+    }
+
+    auto activation = GetAttr<std::string>("activation");
+    AppendOp(framework::OpRegistry::CreateOp(
+        activation, {{"X", {Output("mul_out")}}}, {{"Y", {Output("Y")}}}, {}));
+    CompleteAddOp(false);
+  }
+};
+
+class FCOpMaker : public framework::OpProtoAndCheckerMaker {
+ public:
+  FCOpMaker(framework::OpProto *proto, framework::OpAttrChecker *op_checker)
+      : OpProtoAndCheckerMaker(proto, op_checker) {
+    AddInput("X", "The 2D input matrix of FC operator.");
+    AddInput("W", "The 2D weight matrix of FC operator.");
+    AddInput("b", "The 1D bias vector of FC operator");
+
+    AddOutput("Y", "The activated output matrix of FC operator");
+    AddOutput("mul_out", "The non-actived output of FC operator, X * W + b")
+        .AsIntermediate();
+    AddAttr<std::string>("activation", "The activation type of FC operator.")
+        .SetDefault("identity")
+        .InEnum({"identity", "sigmoid", "softmax"});
+
+    AddComment(R"DOC(
+Fully Connected Operator, known as Fully Connected Layer or Inner Product Layer
+in Convolutional Neural Networks. Neurons in a fully connected layer have
+full connections to all activations in the previous layer.
+It computes an inner product of a set of
+learned weights with a matrix multiplication followed by a bias offset
+(optionally).
+
+Equation:
+  Y = Act(sum_n{X_i * W_i} + b)
+
+where X_i is a 2D matrix of size (M x K), usually M is the minibatch size and
+K is the number of features. W_i is also a 2D matrix of size (K x N),
+where N means the number of neurons in the fully connected layer.
+b is a 1D vector of size N. Thus, the output Y is a 2D matrix of size (M x N).
+Activation type can be set to `identity` (default), `sigmoid` or `softmax`.
+
+  The config api is `paddle.v2.layer.fc`.
+)DOC");
+  }
+};
+
+class FCGradOp : public NetOp {
+ public:
+  FCGradOp(const std::string &type, const framework::VariableNameMap &inputs,
+           const framework::VariableNameMap &outputs,
+           const framework::AttributeMap &attrs)
+      : NetOp(type, inputs, outputs, attrs) {
+    auto y_grad = Input(framework::GradVarName("Y"));
+    auto mul_out_grad = Input(framework::GradVarName("mul_out"));
+    auto x_grad = Output(framework::GradVarName("X"));
+    auto w_grad = Output(framework::GradVarName("W"));
+    auto b_grad = Output(framework::GradVarName("b"));
+
+    CompleteAddOp(false);
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle
+
+USE_OP(mul);
+USE_OP(rowwise_add);
+USE_NO_KERNEL_OP(identity);
+USE_OP(sigmoid);
+USE_OP(softmax);
+
+namespace ops = paddle::operators;
+REGISTER_OP(fc, ops::FCOp, ops::FCOpMaker, fc_grad, ops::FCGradOp);
diff --git a/paddle/operators/scale_op.cc b/paddle/operators/scale_op.cc
index 8e96a74c94..ffc2f02b0b 100644
--- a/paddle/operators/scale_op.cc
+++ b/paddle/operators/scale_op.cc
@@ -89,6 +89,7 @@ class IdentityOp : public NetOp {
     AppendOp(framework::OpRegistry::CreateOp(
         "scale", {{"X", {Input("X")}}}, {{"Out", {Output("Out")}}},
         {{"scale", static_cast<AttrType>(1)}}));
+    CompleteAddOp(false);
   }
 };
 
diff --git a/paddle/pybind/CMakeLists.txt b/paddle/pybind/CMakeLists.txt
index 0003005070..4f05406c7f 100644
--- a/paddle/pybind/CMakeLists.txt
+++ b/paddle/pybind/CMakeLists.txt
@@ -1,5 +1,5 @@
 if(WITH_PYTHON)
-cc_library(paddle_pybind SHARED
+  cc_library(paddle_pybind SHARED
     SRCS pybind.cc
     DEPS pybind python backward
     ${GLOB_OP_LIB})
diff --git a/paddle/pybind/pybind.cc b/paddle/pybind/pybind.cc
index 6896422617..ff6bae8f85 100644
--- a/paddle/pybind/pybind.cc
+++ b/paddle/pybind/pybind.cc
@@ -45,6 +45,7 @@ USE_OP(uniform_random);
 USE_OP(lookup_table);
 USE_OP(scale);
 USE_NO_KERNEL_OP(identity);
+USE_NO_KERNEL_OP(fc);
 USE_OP(minus);
 USE_CPU_ONLY_OP(gather);
 USE_CPU_ONLY_OP(scatter);
diff --git a/python/paddle/v2/framework/tests/CMakeLists.txt b/python/paddle/v2/framework/tests/CMakeLists.txt
index 661ebd8964..807ca2961e 100644
--- a/python/paddle/v2/framework/tests/CMakeLists.txt
+++ b/python/paddle/v2/framework/tests/CMakeLists.txt
@@ -16,6 +16,7 @@ py_test(test_cross_entropy_op SRCS test_cross_entropy_op.py)
 py_test(test_gather_op SRCS test_gather_op.py)
 py_test(test_scatter_op SRCS test_scatter_op.py)
 py_test(test_fill_zeros_like_op SRCS test_fill_zeros_like_op.py)
+py_test(test_fc_op SRCS test_fc_op.py)
 
 py_test(gradient_checker SRCS gradient_checker.py)
 
diff --git a/python/paddle/v2/framework/tests/test_fc_op.py b/python/paddle/v2/framework/tests/test_fc_op.py
new file mode 100644
index 0000000000..bc469a5f47
--- /dev/null
+++ b/python/paddle/v2/framework/tests/test_fc_op.py
@@ -0,0 +1,30 @@
+import unittest
+import numpy as np
+from gradient_checker import GradientChecker, create_op
+from op_test_util import OpTestMeta
+
+
+class TestFCOp(unittest.TestCase):
+    __metaclass__ = OpTestMeta
+
+    def setUp(self):
+        self.type = "fc"
+        self.inputs = {
+            "X": np.random.random((32, 784)).astype("float32"),
+            "W": np.random.random((784, 1000)).astype("float32"),
+            "b": np.random.random(1000).astype("float32")
+        }
+        self.attrs = {"activation": "sigmoid"}
+        mul_out = np.dot(self.inputs["X"], self.inputs["W"])
+        add_out = np.add(mul_out, self.inputs["b"])
+        sigmoid_out = 1 / (1 + np.exp(-add_out))
+        self.outputs = {"mul_out": add_out, "Y": sigmoid_out}
+
+
+class TestFCGradOp(GradientChecker):
+    def test_normal(self):
+        print "nothing"
+
+
+if __name__ == '__main__':
+    unittest.main()

From 734a9eeaa464510de1374c196e40325efd2f8edb Mon Sep 17 00:00:00 2001
From: Liu Yiqun <liuyiqun01@baidu.com>
Date: Thu, 7 Sep 2017 09:29:23 +0000
Subject: [PATCH 02/34] Correct the definition of Operator in TestFCGradOp, and
 rename the output name of identity to Y.

---
 paddle/operators/fc_op.cc                     | 24 +++++++++----------
 paddle/operators/identity_op.cc               |  4 ++--
 paddle/operators/minus_op.cc                  |  2 +-
 .../paddle/v2/framework/tests/CMakeLists.txt  |  1 +
 .../v2/framework/tests/gradient_checker.py    |  4 ----
 .../paddle/v2/framework/tests/test_fc_op.py   | 17 +++++++++----
 .../v2/framework/tests/test_minus_op.py       |  4 ++--
 .../tests/test_scale_and_identity_op.py       |  4 ++--
 8 files changed, 32 insertions(+), 28 deletions(-)

diff --git a/paddle/operators/fc_op.cc b/paddle/operators/fc_op.cc
index 60bf6e9dae..40b5128bff 100644
--- a/paddle/operators/fc_op.cc
+++ b/paddle/operators/fc_op.cc
@@ -24,30 +24,30 @@ class FCOp : public NetOp {
        const framework::VariableNameMap &outputs,
        const framework::AttributeMap &attrs)
       : NetOp(type, inputs, outputs, attrs) {
+    // mul_out = X * W
     AppendOp(framework::OpRegistry::CreateOp(
         "mul", {{"X", {Input("X")}}, {"Y", {Input("W")}}},
         {{"Out", {Output("mul_out")}}}, {}));
+
+    std::string add_out_name = "mul_out";
     auto b = Input("b");
     if (b != framework::kEmptyVarName) {
+      // add_out = mul_out + b
       AppendOp(framework::OpRegistry::CreateOp(
           "rowwise_add", {{"X", {Output("mul_out")}}, {"b", {Input("b")}}},
           {{"Out", {Output("add_out")}}}, {}));
+      add_out_name = "add_out";
     } else {
-      AppendOp(framework::OpRegistry::CreateOp(
-          "identity", {{"X", {Output("mul_out")}}},
-          {{"Out", {Output("add_out")}}}, {}));
+      auto add_out = Output("add_out");
+      if (add_out != framework::kEmptyVarName) {
+        this->Rename(add_out, framework::kEmptyVarName);
+      }
     }
 
     auto activation = GetAttr<std::string>("activation");
-    if (activation == "identity") {
-      AppendOp(framework::OpRegistry::CreateOp(activation,
-                                               {{"X", {Output("add_out")}}},
-                                               {{"Out", {Output("Out")}}}, {}));
-    } else {
-      AppendOp(framework::OpRegistry::CreateOp(activation,
-                                               {{"X", {Output("add_out")}}},
-                                               {{"Y", {Output("Out")}}}, {}));
-    }
+    AppendOp(framework::OpRegistry::CreateOp(activation,
+                                             {{"X", {Output(add_out_name)}}},
+                                             {{"Y", {Output("Out")}}}, {}));
     CompleteAddOp(false);
   }
 };
diff --git a/paddle/operators/identity_op.cc b/paddle/operators/identity_op.cc
index b67240fb9f..b9f0a450fd 100644
--- a/paddle/operators/identity_op.cc
+++ b/paddle/operators/identity_op.cc
@@ -27,7 +27,7 @@ class IdentityOpMaker : public framework::OpProtoAndCheckerMaker {
                   framework::OpAttrChecker *op_checker)
       : OpProtoAndCheckerMaker(proto, op_checker) {
     AddInput("X", "input tensor of identity op");
-    AddOutput("Out", "output tensor of identity op");
+    AddOutput("Y", "output tensor of identity op");
     AddComment("identity operator. Just a alias of scale op which scale = 1.0");
   }
 };
@@ -40,7 +40,7 @@ class IdentityOp : public NetOp {
              const framework::AttributeMap &attrs)
       : NetOp(type, inputs, outputs, attrs) {
     AppendOp(framework::OpRegistry::CreateOp(
-        "scale", {{"X", {Input("X")}}}, {{"Out", {Output("Out")}}},
+        "scale", {{"X", {Input("X")}}}, {{"Out", {Output("Y")}}},
         {{"scale", static_cast<AttrType>(1)}}));
     CompleteAddOp(false);
   }
diff --git a/paddle/operators/minus_op.cc b/paddle/operators/minus_op.cc
index 069fb5e1ab..1d7276a194 100644
--- a/paddle/operators/minus_op.cc
+++ b/paddle/operators/minus_op.cc
@@ -65,7 +65,7 @@ class MinusGradOp : public NetOp {
 
     // x_grad = out_grad
     AppendOp(framework::OpRegistry::CreateOp("identity", {{"X", {out_grad}}},
-                                             {{"Out", {x_grad}}}, {}));
+                                             {{"Y", {x_grad}}}, {}));
 
     framework::AttributeMap scale_attr;
     scale_attr["scale"] = static_cast<AttrType>(-1);
diff --git a/python/paddle/v2/framework/tests/CMakeLists.txt b/python/paddle/v2/framework/tests/CMakeLists.txt
index f9c787a446..60ee996e4a 100644
--- a/python/paddle/v2/framework/tests/CMakeLists.txt
+++ b/python/paddle/v2/framework/tests/CMakeLists.txt
@@ -18,6 +18,7 @@ py_test(test_gather_op SRCS test_gather_op.py)
 py_test(test_scatter_op SRCS test_scatter_op.py)
 py_test(test_fill_zeros_like_op SRCS test_fill_zeros_like_op.py)
 py_test(test_fc_op SRCS test_fc_op.py)
+py_test(test_minus_op SRCS test_minus_op.py)
 
 py_test(gradient_checker SRCS gradient_checker.py)
 
diff --git a/python/paddle/v2/framework/tests/gradient_checker.py b/python/paddle/v2/framework/tests/gradient_checker.py
index fdb06b7988..0607275a4e 100644
--- a/python/paddle/v2/framework/tests/gradient_checker.py
+++ b/python/paddle/v2/framework/tests/gradient_checker.py
@@ -277,10 +277,6 @@ class GradientChecker(unittest.TestCase):
         if no_grad_set is None:
             no_grad_set = set()
 
-        no_tmp_out = forward_op.no_intermediate_outputs()
-        if len(no_tmp_out) != 1:
-            raise ValueError("non temp out_names should be 1")
-
         inputs = forward_op.inputs()
         in_names = [item for k in inputs for item in inputs[k]]
         for no_grad in no_grad_set:
diff --git a/python/paddle/v2/framework/tests/test_fc_op.py b/python/paddle/v2/framework/tests/test_fc_op.py
index 140442db96..76b68ad614 100644
--- a/python/paddle/v2/framework/tests/test_fc_op.py
+++ b/python/paddle/v2/framework/tests/test_fc_op.py
@@ -29,13 +29,20 @@ class TestFCOp(unittest.TestCase):
 class TestFCGradOp(GradientChecker):
     def test_normal(self):
         self.inputs = {
-            "X": np.random.random((4, 4)).astype("float32"),
-            "W": np.random.random((4, 4)).astype("float32"),
-            "b": np.random.random(4).astype("float32")
+            "X": np.random.random((32, 256)).astype("float32"),
+            "W": np.random.random((256, 100)).astype("float32"),
+            "b": np.random.random(100).astype("float32")
         }
         op = Operator(
-            "fc", X="X", W="W", b="b", Out="Out", activation="sigmoid")
-        #self.check_grad(op, self.inputs, ["X", "W", "b"], "Out")
+            "fc",
+            X="X",
+            W="W",
+            b="b",
+            Out="Out",
+            mul_out="mul_out",
+            add_out="add_out",
+            activation="sigmoid")
+        self.check_grad(op, self.inputs, ["X", "W", "b"], "Out")
 
 
 if __name__ == '__main__':
diff --git a/python/paddle/v2/framework/tests/test_minus_op.py b/python/paddle/v2/framework/tests/test_minus_op.py
index 5abdd4a69b..aa05d87baa 100644
--- a/python/paddle/v2/framework/tests/test_minus_op.py
+++ b/python/paddle/v2/framework/tests/test_minus_op.py
@@ -4,7 +4,7 @@ from gradient_checker import GradientChecker, create_op
 from op_test_util import OpTestMeta
 
 
-class MinusOpTest(unittest.TestCase):
+class TestMinusOp(unittest.TestCase):
     __metaclass__ = OpTestMeta
 
     def setUp(self):
@@ -16,7 +16,7 @@ class MinusOpTest(unittest.TestCase):
         self.outputs = {'Out': (self.inputs['X'] - self.inputs['Y'])}
 
 
-class MinusGradTest(GradientChecker):
+class TestMinusGrad(GradientChecker):
     def test_left(self):
         op = create_op("minus")
         inputs = {
diff --git a/python/paddle/v2/framework/tests/test_scale_and_identity_op.py b/python/paddle/v2/framework/tests/test_scale_and_identity_op.py
index 69b301c376..4c1d484991 100644
--- a/python/paddle/v2/framework/tests/test_scale_and_identity_op.py
+++ b/python/paddle/v2/framework/tests/test_scale_and_identity_op.py
@@ -11,14 +11,14 @@ class IdentityTest(unittest.TestCase):
     def setUp(self):
         self.type = "identity"
         self.inputs = {'X': np.random.random((32, 784)).astype("float32")}
-        self.outputs = {'Out': self.inputs['X']}
+        self.outputs = {'Y': self.inputs['X']}
 
 
 class IdentityGradOpTest(GradientChecker):
     def test_normal(self):
         op = create_op("identity")
         inputs = {"X": np.random.random((10, 10)).astype("float32")}
-        self.check_grad(op, inputs, set("X"), "Out")
+        self.check_grad(op, inputs, set("X"), "Y")
 
 
 class ScaleTest(unittest.TestCase):

From d874fca46bb77be244ed38d3edcc029a446cfafa Mon Sep 17 00:00:00 2001
From: Liu Yiqun <liuyiqun01@baidu.com>
Date: Mon, 11 Sep 2017 04:29:38 +0000
Subject: [PATCH 03/34] Support multiple inputs in FCOp.

---
 paddle/operators/fc_op.cc                     | 55 +++++++++++------
 .../paddle/v2/framework/tests/test_fc_op.py   | 61 +++++++++----------
 2 files changed, 67 insertions(+), 49 deletions(-)

diff --git a/paddle/operators/fc_op.cc b/paddle/operators/fc_op.cc
index 40b5128bff..ec76d6c659 100644
--- a/paddle/operators/fc_op.cc
+++ b/paddle/operators/fc_op.cc
@@ -24,30 +24,49 @@ class FCOp : public NetOp {
        const framework::VariableNameMap &outputs,
        const framework::AttributeMap &attrs)
       : NetOp(type, inputs, outputs, attrs) {
-    // mul_out = X * W
-    AppendOp(framework::OpRegistry::CreateOp(
-        "mul", {{"X", {Input("X")}}, {"Y", {Input("W")}}},
-        {{"Out", {Output("mul_out")}}}, {}));
+    auto x = Inputs("X");
+    auto w = Inputs("W");
+    PADDLE_ENFORCE_EQ(
+        x.size(), w.size(),
+        "The size of inputs X(%d) should be the same as that of weights W(%d).",
+        x.size(), w.size());
+
+    int n = x.size();
+    PADDLE_ENFORCE_GE(n, 1,
+                      "The size of inputs X(%d) should be no less than 1.", n);
+
+    // mul_out = X[0] * W[0] + ... + X[n-1] * W[n-1]
+    AppendOp(
+        framework::OpRegistry::CreateOp("mul", {{"X", {x[0]}}, {"W", {w[0]}}},
+                                        {{"Out", {Output("mul_out")}}}, {}));
+
+    for (int i = 1; i < n; i++) {
+      // mul_out = mul_out + X[i] * W[i]
+      AppendOp(
+          framework::OpRegistry::CreateOp("mul", {{"X", {x[i]}}, {"Y", {w[i]}}},
+                                          {{"Out", {Output("add_out")}}}, {}));
+      AppendOp(framework::OpRegistry::CreateOp(
+          "add", {{"X", {Output("mul_out")}}, {"Y", {Output("add_out")}}},
+          {{"Out", {Output("mul_out")}}}, {}));
+    }
 
-    std::string add_out_name = "mul_out";
     auto b = Input("b");
+    std::string add_out = "mul_out";
     if (b != framework::kEmptyVarName) {
       // add_out = mul_out + b
       AppendOp(framework::OpRegistry::CreateOp(
           "rowwise_add", {{"X", {Output("mul_out")}}, {"b", {Input("b")}}},
           {{"Out", {Output("add_out")}}}, {}));
-      add_out_name = "add_out";
+      add_out = "add_out";
     } else {
-      auto add_out = Output("add_out");
-      if (add_out != framework::kEmptyVarName) {
-        this->Rename(add_out, framework::kEmptyVarName);
+      if (Output("add_out") != framework::kEmptyVarName) {
+        this->Rename(Output("add_out"), framework::kEmptyVarName);
       }
     }
 
-    auto activation = GetAttr<std::string>("activation");
-    AppendOp(framework::OpRegistry::CreateOp(activation,
-                                             {{"X", {Output(add_out_name)}}},
-                                             {{"Y", {Output("Out")}}}, {}));
+    auto activation = Attr<std::string>("activation");
+    AppendOp(framework::OpRegistry::CreateOp(
+        activation, {{"X", {Output(add_out)}}}, {{"Y", {Output("Y")}}}, {}));
     CompleteAddOp(false);
   }
 };
@@ -56,11 +75,11 @@ class FCOpMaker : public framework::OpProtoAndCheckerMaker {
  public:
   FCOpMaker(framework::OpProto *proto, framework::OpAttrChecker *op_checker)
       : OpProtoAndCheckerMaker(proto, op_checker) {
-    AddInput("X", "The 2D input matrix of FC operator.");
-    AddInput("W", "The 2D weight matrix of FC operator.");
-    AddInput("b", "The 1D bias vector of FC operator");
+    AddInput("X", "The 2-D input matrix of FC operator.").AsDuplicable();
+    AddInput("W", "The 2-D weight matrix of FC operator.").AsDuplicable();
+    AddInput("b", "The 1-D bias vector of FC operator");
 
-    AddOutput("Out", "The activated output matrix of FC operator");
+    AddOutput("Y", "The activated output matrix of FC operator");
     AddOutput("mul_out", "The non-actived output of FC operator, X * W")
         .AsIntermediate();
     AddOutput("add_out", "The non-actived output of FC operator, X * W + b")
@@ -78,7 +97,7 @@ learned weights with a matrix multiplication followed by a bias offset
 (optionally).
 
 Equation:
-  Out = Act(sum_n{X_i * W_i} + b)
+  Y = Act(sum_n{X_i * W_i} + b)
 
 where X_i is a 2D matrix of size (M x K), usually M is the minibatch size and
 K is the number of features. W_i is also a 2D matrix of size (K x N),
diff --git a/python/paddle/v2/framework/tests/test_fc_op.py b/python/paddle/v2/framework/tests/test_fc_op.py
index 76b68ad614..72d750111c 100644
--- a/python/paddle/v2/framework/tests/test_fc_op.py
+++ b/python/paddle/v2/framework/tests/test_fc_op.py
@@ -1,48 +1,47 @@
 import unittest
 import numpy as np
-from gradient_checker import GradientChecker, create_op
-from op_test_util import OpTestMeta
+from op_test import OpTest
+import paddle.v2.framework.core as core
 from paddle.v2.framework.op import Operator
 
 
-class TestFCOp(unittest.TestCase):
-    __metaclass__ = OpTestMeta
-
+class TestFCOp(OpTest):
     def setUp(self):
-        self.type = "fc"
+        print "Run"
+        self.op_type = "fc"
+        x0 = np.random.random((32, 256)).astype("float32")
+        x1 = np.random.random((32, 256)).astype("float32")
+        w0 = np.random.random((256, 100)).astype("float32")
+        w1 = np.random.random((256, 100)).astype("float32")
+        b = np.random.random(100).astype("float32")
         self.inputs = {
-            "X": np.random.random((32, 784)).astype("float32"),
-            "W": np.random.random((784, 1000)).astype("float32"),
-            "b": np.random.random(1000).astype("float32")
+            "X": {
+                "X0": x0,
+                "X1": x1
+            },
+            "W": {
+                "W0": w0,
+                "W1": w1
+            },
+            "b": b
         }
-        self.attrs = {"activation": "sigmoid"}
-        mul_out = np.dot(self.inputs["X"], self.inputs["W"])
-        add_out = np.add(mul_out, self.inputs["b"])
-        sigmoid_out = 1 / (1 + np.exp(-add_out))
+        #self.attrs = {"activation": "sigmoid"}
+        mul_out = np.dot(x0, w0) + np.dot(x1, w1)
+        add_out = np.add(mul_out, b)
+        #sigmoid_out = 1 / (1 + np.exp(-add_out))
+        sigmoid_out = add_out
         self.outputs = {
             "mul_out": mul_out,
             "add_out": add_out,
-            "Out": sigmoid_out
+            "Y": sigmoid_out
         }
 
+    def test_check_output(self):
+        self.check_output(core.CPUPlace())
+        self.check_output(core.GPUPlace(0))
 
-class TestFCGradOp(GradientChecker):
-    def test_normal(self):
-        self.inputs = {
-            "X": np.random.random((32, 256)).astype("float32"),
-            "W": np.random.random((256, 100)).astype("float32"),
-            "b": np.random.random(100).astype("float32")
-        }
-        op = Operator(
-            "fc",
-            X="X",
-            W="W",
-            b="b",
-            Out="Out",
-            mul_out="mul_out",
-            add_out="add_out",
-            activation="sigmoid")
-        self.check_grad(op, self.inputs, ["X", "W", "b"], "Out")
+    #def test_check_grad(self):
+    #    self.check_grad(["X0", "X1", "W0", "W1", "b"], "Y")
 
 
 if __name__ == '__main__':

From 4223ff8c27a8681096ad73659a6ea03441cf3831 Mon Sep 17 00:00:00 2001
From: Liu Yiqun <liuyiqun01@baidu.com>
Date: Mon, 11 Sep 2017 07:00:04 +0000
Subject: [PATCH 04/34] Correct the key name of "mul" op in FCOp, and add some
 annotations for debug.

---
 paddle/operators/fc_op.cc                      | 4 +++-
 python/paddle/v2/framework/tests/op_test.py    | 8 ++++----
 python/paddle/v2/framework/tests/test_fc_op.py | 5 +----
 3 files changed, 8 insertions(+), 9 deletions(-)

diff --git a/paddle/operators/fc_op.cc b/paddle/operators/fc_op.cc
index ec76d6c659..6e6a09bc3f 100644
--- a/paddle/operators/fc_op.cc
+++ b/paddle/operators/fc_op.cc
@@ -37,7 +37,7 @@ class FCOp : public NetOp {
 
     // mul_out = X[0] * W[0] + ... + X[n-1] * W[n-1]
     AppendOp(
-        framework::OpRegistry::CreateOp("mul", {{"X", {x[0]}}, {"W", {w[0]}}},
+        framework::OpRegistry::CreateOp("mul", {{"X", {x[0]}}, {"Y", {w[0]}}},
                                         {{"Out", {Output("mul_out")}}}, {}));
 
     for (int i = 1; i < n; i++) {
@@ -68,6 +68,8 @@ class FCOp : public NetOp {
     AppendOp(framework::OpRegistry::CreateOp(
         activation, {{"X", {Output(add_out)}}}, {{"Y", {Output("Y")}}}, {}));
     CompleteAddOp(false);
+
+    std::cout << DebugString() << std::endl;
   }
 };
 
diff --git a/python/paddle/v2/framework/tests/op_test.py b/python/paddle/v2/framework/tests/op_test.py
index 3a6a5dca4c..b524f88551 100644
--- a/python/paddle/v2/framework/tests/op_test.py
+++ b/python/paddle/v2/framework/tests/op_test.py
@@ -36,8 +36,8 @@ def create_op(scope, op_type, inputs, outputs, attrs=None):
                 var = scope.new_var(out_name)
                 kwargs[out_name].append(out_name)
 
-    for attr_name in Operator.get_op_attr_names(op_type):
-        kwargs[attr_name] = attrs[attr_name]
+    #for attr_name in Operator.get_op_attr_names(op_type):
+    #    kwargs[attr_name] = attrs[attr_name]
     return Operator(op_type, **kwargs)
 
 
@@ -186,14 +186,14 @@ class OpTest(unittest.TestCase):
                     self.assertTrue(
                         np.allclose(
                             actual, expect, atol=1e-05),
-                        "output name: " + out_name + "has diff")
+                        "output name: " + out_name + " has diff")
             else:
                 actual = np.array(self.scope.find_var(out_name).get_tensor())
                 expect = self.outputs[out_name]
                 self.assertTrue(
                     np.allclose(
                         actual, expect, atol=1e-05),
-                    "output name: " + out_name + "has diff")
+                    "output name: " + out_name + " has diff")
 
     def check_output(self):
         places = [core.CPUPlace()]
diff --git a/python/paddle/v2/framework/tests/test_fc_op.py b/python/paddle/v2/framework/tests/test_fc_op.py
index 72d750111c..959bd7e405 100644
--- a/python/paddle/v2/framework/tests/test_fc_op.py
+++ b/python/paddle/v2/framework/tests/test_fc_op.py
@@ -1,8 +1,6 @@
 import unittest
 import numpy as np
 from op_test import OpTest
-import paddle.v2.framework.core as core
-from paddle.v2.framework.op import Operator
 
 
 class TestFCOp(OpTest):
@@ -37,8 +35,7 @@ class TestFCOp(OpTest):
         }
 
     def test_check_output(self):
-        self.check_output(core.CPUPlace())
-        self.check_output(core.GPUPlace(0))
+        self.check_output()
 
     #def test_check_grad(self):
     #    self.check_grad(["X0", "X1", "W0", "W1", "b"], "Y")

From 4f2ee63c4466e018b9a1196281c1a059094698f3 Mon Sep 17 00:00:00 2001
From: Liu Yiqun <liuyiqun01@baidu.com>
Date: Mon, 11 Sep 2017 11:10:25 +0000
Subject: [PATCH 05/34] Get rid of the calling of inplace op in FCOp.

---
 paddle/operators/fc_op.cc                     | 57 +++++++++-------
 .../paddle/v2/framework/tests/test_fc_op.py   | 65 ++++++++++++++-----
 2 files changed, 84 insertions(+), 38 deletions(-)

diff --git a/paddle/operators/fc_op.cc b/paddle/operators/fc_op.cc
index 6e6a09bc3f..1c6c045427 100644
--- a/paddle/operators/fc_op.cc
+++ b/paddle/operators/fc_op.cc
@@ -26,38 +26,43 @@ class FCOp : public NetOp {
       : NetOp(type, inputs, outputs, attrs) {
     auto x = Inputs("X");
     auto w = Inputs("W");
+    auto mul_out = Outputs("mul_out");
     PADDLE_ENFORCE_EQ(
         x.size(), w.size(),
         "The size of inputs X(%d) should be the same as that of weights W(%d).",
         x.size(), w.size());
+    PADDLE_ENFORCE_EQ(mul_out.size(), x.size(),
+                      "The size of intermediate mul_out(%d) should be the same "
+                      "as that of inputs X(%d).",
+                      mul_out.size(), x.size());
 
     int n = x.size();
     PADDLE_ENFORCE_GE(n, 1,
                       "The size of inputs X(%d) should be no less than 1.", n);
 
-    // mul_out = X[0] * W[0] + ... + X[n-1] * W[n-1]
-    AppendOp(
-        framework::OpRegistry::CreateOp("mul", {{"X", {x[0]}}, {"Y", {w[0]}}},
-                                        {{"Out", {Output("mul_out")}}}, {}));
+    // mul_out[i] = X[i] * W[i]
+    for (int i = 0; i < n; i++) {
+      AppendOp(framework::OpRegistry::CreateOp(
+          "mul", {{"X", {x[i]}}, {"Y", {w[i]}}}, {{"Out", {mul_out[i]}}}, {}));
+    }
 
-    for (int i = 1; i < n; i++) {
-      // mul_out = mul_out + X[i] * W[i]
-      AppendOp(
-          framework::OpRegistry::CreateOp("mul", {{"X", {x[i]}}, {"Y", {w[i]}}},
-                                          {{"Out", {Output("add_out")}}}, {}));
+    // sum_out = X[0] * W[0] + ... + X[n-1] * W[n-1]
+    if (n > 1) {
+      AppendOp(framework::OpRegistry::CreateOp(
+          "sum", {{"X", {mul_out}}}, {{"Out", {Output("sum_out")}}}, {}));
+    } else {
       AppendOp(framework::OpRegistry::CreateOp(
-          "add", {{"X", {Output("mul_out")}}, {"Y", {Output("add_out")}}},
-          {{"Out", {Output("mul_out")}}}, {}));
+          "identity", {{"X", {mul_out[0]}}}, {{"Y", {Output("sum_out")}}}, {}));
     }
 
+    // add_out = sum_out + b
     auto b = Input("b");
-    std::string add_out = "mul_out";
+    std::string add_out = "sum_out";
     if (b != framework::kEmptyVarName) {
-      // add_out = mul_out + b
-      AppendOp(framework::OpRegistry::CreateOp(
-          "rowwise_add", {{"X", {Output("mul_out")}}, {"b", {Input("b")}}},
-          {{"Out", {Output("add_out")}}}, {}));
       add_out = "add_out";
+      AppendOp(framework::OpRegistry::CreateOp(
+          "rowwise_add", {{"X", {Output("sum_out")}}, {"b", {Input("b")}}},
+          {{"Out", {Output(add_out)}}}, {}));
     } else {
       if (Output("add_out") != framework::kEmptyVarName) {
         this->Rename(Output("add_out"), framework::kEmptyVarName);
@@ -68,8 +73,6 @@ class FCOp : public NetOp {
     AppendOp(framework::OpRegistry::CreateOp(
         activation, {{"X", {Output(add_out)}}}, {{"Y", {Output("Y")}}}, {}));
     CompleteAddOp(false);
-
-    std::cout << DebugString() << std::endl;
   }
 };
 
@@ -77,14 +80,24 @@ class FCOpMaker : public framework::OpProtoAndCheckerMaker {
  public:
   FCOpMaker(framework::OpProto *proto, framework::OpAttrChecker *op_checker)
       : OpProtoAndCheckerMaker(proto, op_checker) {
-    AddInput("X", "The 2-D input matrix of FC operator.").AsDuplicable();
-    AddInput("W", "The 2-D weight matrix of FC operator.").AsDuplicable();
+    AddInput("X", "The inputs of FC operator, a ordered vector of 2-D matrix.")
+        .AsDuplicable();
+    AddInput("W", "The weights of FC operator, a ordered vector of 2-D matrix.")
+        .AsDuplicable();
     AddInput("b", "The 1-D bias vector of FC operator");
 
     AddOutput("Y", "The activated output matrix of FC operator");
-    AddOutput("mul_out", "The non-actived output of FC operator, X * W")
+    AddOutput("mul_out",
+              "The intermediate outputs of FC operator, "
+              "saving the product of X[i] * W[i]")
+        .AsIntermediate()
+        .AsDuplicable();
+    AddOutput("sum_out",
+              "The intermediate output of FC operator, "
+              "saving the sum of products, sum(X[i] * W[i])")
         .AsIntermediate();
-    AddOutput("add_out", "The non-actived output of FC operator, X * W + b")
+    AddOutput("add_out",
+              "The non-actived output of FC operator, saving X * W + b")
         .AsIntermediate();
     AddAttr<std::string>("activation", "The activation type of FC operator.")
         .SetDefault("identity")
diff --git a/python/paddle/v2/framework/tests/test_fc_op.py b/python/paddle/v2/framework/tests/test_fc_op.py
index 959bd7e405..4355191223 100644
--- a/python/paddle/v2/framework/tests/test_fc_op.py
+++ b/python/paddle/v2/framework/tests/test_fc_op.py
@@ -3,33 +3,65 @@ import numpy as np
 from op_test import OpTest
 
 
-class TestFCOp(OpTest):
+class TestFCOp1(OpTest):
     def setUp(self):
-        print "Run"
         self.op_type = "fc"
-        x0 = np.random.random((32, 256)).astype("float32")
-        x1 = np.random.random((32, 256)).astype("float32")
-        w0 = np.random.random((256, 100)).astype("float32")
-        w1 = np.random.random((256, 100)).astype("float32")
-        b = np.random.random(100).astype("float32")
+        x1 = np.random.random((16, 32)).astype("float32")
+        w1 = np.random.random((32, 10)).astype("float32")
+        b = np.random.random(10).astype("float32")
+        self.inputs = {"X": {"X1": x1}, "W": {"W1": w1}, "b": b}
+        mul_out1 = np.dot(x1, w1)
+        sum_out = mul_out1
+        add_out = sum_out + b
+        identity_out = add_out
+        self.outputs = {
+            "mul_out": {
+                "mul_out1": mul_out1,
+            },
+            "sum_out": sum_out,
+            "add_out": add_out,
+            "Y": identity_out
+        }
+
+    def test_check_output(self):
+        self.check_output()
+
+    def test_check_grad(self):
+        self.check_grad(["X1", "W1", "b"], "Y", max_relative_error=0.05)
+
+
+class TestFCOp2(OpTest):
+    def setUp(self):
+        self.op_type = "fc"
+        x1 = np.random.random((16, 32)).astype("float32")
+        x2 = np.random.random((16, 32)).astype("float32")
+        w1 = np.random.random((32, 10)).astype("float32")
+        w2 = np.random.random((32, 10)).astype("float32")
+        b = np.random.random(10).astype("float32")
         self.inputs = {
             "X": {
-                "X0": x0,
-                "X1": x1
+                "X1": x1,
+                "X2": x2
             },
             "W": {
-                "W0": w0,
-                "W1": w1
+                "W1": w1,
+                "W2": w2
             },
             "b": b
         }
         #self.attrs = {"activation": "sigmoid"}
-        mul_out = np.dot(x0, w0) + np.dot(x1, w1)
-        add_out = np.add(mul_out, b)
+        mul_out1 = np.dot(x1, w1)
+        mul_out2 = np.dot(x2, w2)
+        sum_out = mul_out1 + mul_out2
+        add_out = np.add(sum_out, b)
         #sigmoid_out = 1 / (1 + np.exp(-add_out))
         sigmoid_out = add_out
         self.outputs = {
-            "mul_out": mul_out,
+            "mul_out": {
+                "mul_out0": mul_out1,
+                "mul_out1": mul_out2
+            },
+            "sum_out": sum_out,
             "add_out": add_out,
             "Y": sigmoid_out
         }
@@ -37,8 +69,9 @@ class TestFCOp(OpTest):
     def test_check_output(self):
         self.check_output()
 
-    #def test_check_grad(self):
-    #    self.check_grad(["X0", "X1", "W0", "W1", "b"], "Y")
+    def test_check_grad(self):
+        self.check_grad(
+            ["X1", "X2", "W1", "W2", "b"], "Y", max_relative_error=0.05)
 
 
 if __name__ == '__main__':

From 297b3d0f0f03461bf067d6196863454f685c3567 Mon Sep 17 00:00:00 2001
From: zhangchao41 <zhangchao41@baidu.com>
Date: Tue, 12 Sep 2017 11:57:14 +0800
Subject: [PATCH 06/34] add test

---
 test | 0
 1 file changed, 0 insertions(+), 0 deletions(-)
 create mode 100644 test

diff --git a/test b/test
new file mode 100644
index 0000000000..e69de29bb2

From b90461b9d978913d2210ef1f02d4e62197a1458c Mon Sep 17 00:00:00 2001
From: zhangchao41 <zhangchao41@baidu.com>
Date: Tue, 12 Sep 2017 13:40:43 +0800
Subject: [PATCH 07/34] fix the typo of the param description in
 sequence_conv_pool

---
 python/paddle/trainer_config_helpers/networks.py | 6 +++---
 test                                             | 0
 2 files changed, 3 insertions(+), 3 deletions(-)
 delete mode 100644 test

diff --git a/python/paddle/trainer_config_helpers/networks.py b/python/paddle/trainer_config_helpers/networks.py
index 34be203ee2..2f604ee45a 100644
--- a/python/paddle/trainer_config_helpers/networks.py
+++ b/python/paddle/trainer_config_helpers/networks.py
@@ -68,7 +68,7 @@ def sequence_conv_pool(input,
     :type context_len: int
     :param hidden_size: FC Layer size.
     :type hidden_size: int
-    :param context_start: context projection length. See
+    :param context_start: context start position. See
                           context_projection's context_start.
     :type context_start: int or None
     :param pool_type: pooling layer type. See pooling_layer's document.
@@ -1406,7 +1406,7 @@ def inputs(layers, *args):
     if len(args) != 0:
         layers.extend(args)
 
-    Inputs(*[l.name for l in layers])
+    Inputs(* [l.name for l in layers])
 
 
 def outputs(layers, *args):
@@ -1456,7 +1456,7 @@ def outputs(layers, *args):
     assert len(layers) > 0
 
     if HasInputsSet():  # input already set
-        Outputs(*[l.name for l in layers])
+        Outputs(* [l.name for l in layers])
         return  # just return outputs.
 
     if len(layers) != 1:
diff --git a/test b/test
deleted file mode 100644
index e69de29bb2..0000000000

From af2eb94909a0b938b23b0959b42a1b8c36236778 Mon Sep 17 00:00:00 2001
From: Liu Yiqun <liuyiqun01@baidu.com>
Date: Wed, 13 Sep 2017 12:48:22 +0000
Subject: [PATCH 08/34] Support inputs and weights of multi-dimensions and
 refine the output names.

---
 paddle/operators/fc_op.cc                     | 53 ++++++++++++-------
 .../paddle/v2/framework/tests/test_fc_op.py   | 49 ++++++++++-------
 2 files changed, 64 insertions(+), 38 deletions(-)

diff --git a/paddle/operators/fc_op.cc b/paddle/operators/fc_op.cc
index 1c6c045427..3e6cd8f76a 100644
--- a/paddle/operators/fc_op.cc
+++ b/paddle/operators/fc_op.cc
@@ -26,7 +26,7 @@ class FCOp : public NetOp {
       : NetOp(type, inputs, outputs, attrs) {
     auto x = Inputs("X");
     auto w = Inputs("W");
-    auto mul_out = Outputs("mul_out");
+    auto mul_out = Outputs("MulOut");
     PADDLE_ENFORCE_EQ(
         x.size(), w.size(),
         "The size of inputs X(%d) should be the same as that of weights W(%d).",
@@ -36,36 +36,51 @@ class FCOp : public NetOp {
                       "as that of inputs X(%d).",
                       mul_out.size(), x.size());
 
-    int n = x.size();
-    PADDLE_ENFORCE_GE(n, 1,
+    size_t n = x.size();
+    PADDLE_ENFORCE_GE(n, static_cast<size_t>(1),
                       "The size of inputs X(%d) should be no less than 1.", n);
 
+    auto x_num_col_dims = Attr<std::vector<int>>("xNumColDims");
+    auto w_num_col_dims = Attr<std::vector<int>>("wNumColDims");
+    PADDLE_ENFORCE_EQ(x_num_col_dims.size(), n,
+                      "The size of attribute xNumColDims(%d) should be the "
+                      "same as that of inputs X(%d).",
+                      x_num_col_dims.size(), n);
+    PADDLE_ENFORCE_EQ(w_num_col_dims.size(), n,
+                      "The size of attribute wNumColDims(%d) should be the "
+                      "same as that of inputs X(%d).",
+                      w_num_col_dims.size(), n)
+
     // mul_out[i] = X[i] * W[i]
-    for (int i = 0; i < n; i++) {
-      AppendOp(framework::OpRegistry::CreateOp(
-          "mul", {{"X", {x[i]}}, {"Y", {w[i]}}}, {{"Out", {mul_out[i]}}}, {}));
+    for (size_t i = 0; i < n; i++) {
+      framework::AttributeMap mul_attr;
+      mul_attr["x_num_col_dims"] = static_cast<int>(x_num_col_dims[i]);
+      mul_attr["y_num_col_dims"] = static_cast<int>(w_num_col_dims[i]);
+      AppendOp(
+          framework::OpRegistry::CreateOp("mul", {{"X", {x[i]}}, {"Y", {w[i]}}},
+                                          {{"Out", {mul_out[i]}}}, mul_attr));
     }
 
     // sum_out = X[0] * W[0] + ... + X[n-1] * W[n-1]
     if (n > 1) {
       AppendOp(framework::OpRegistry::CreateOp(
-          "sum", {{"X", {mul_out}}}, {{"Out", {Output("sum_out")}}}, {}));
+          "sum", {{"X", {mul_out}}}, {{"Out", {Output("SumOut")}}}, {}));
     } else {
       AppendOp(framework::OpRegistry::CreateOp(
-          "identity", {{"X", {mul_out[0]}}}, {{"Y", {Output("sum_out")}}}, {}));
+          "identity", {{"X", {mul_out[0]}}}, {{"Y", {Output("SumOut")}}}, {}));
     }
 
     // add_out = sum_out + b
-    auto b = Input("b");
-    std::string add_out = "sum_out";
+    auto b = Input("B");
+    std::string add_out = "SumOut";
     if (b != framework::kEmptyVarName) {
-      add_out = "add_out";
+      add_out = "AddOut";
       AppendOp(framework::OpRegistry::CreateOp(
-          "rowwise_add", {{"X", {Output("sum_out")}}, {"b", {Input("b")}}},
+          "rowwise_add", {{"X", {Output("SumOut")}}, {"b", {Input("B")}}},
           {{"Out", {Output(add_out)}}}, {}));
     } else {
-      if (Output("add_out") != framework::kEmptyVarName) {
-        this->Rename(Output("add_out"), framework::kEmptyVarName);
+      if (Output("AddOut") != framework::kEmptyVarName) {
+        this->Rename(Output("AddOut"), framework::kEmptyVarName);
       }
     }
 
@@ -84,24 +99,26 @@ class FCOpMaker : public framework::OpProtoAndCheckerMaker {
         .AsDuplicable();
     AddInput("W", "The weights of FC operator, a ordered vector of 2-D matrix.")
         .AsDuplicable();
-    AddInput("b", "The 1-D bias vector of FC operator");
+    AddInput("B", "The 1-D bias vector of FC operator");
 
     AddOutput("Y", "The activated output matrix of FC operator");
-    AddOutput("mul_out",
+    AddOutput("MulOut",
               "The intermediate outputs of FC operator, "
               "saving the product of X[i] * W[i]")
         .AsIntermediate()
         .AsDuplicable();
-    AddOutput("sum_out",
+    AddOutput("SumOut",
               "The intermediate output of FC operator, "
               "saving the sum of products, sum(X[i] * W[i])")
         .AsIntermediate();
-    AddOutput("add_out",
+    AddOutput("AddOut",
               "The non-actived output of FC operator, saving X * W + b")
         .AsIntermediate();
     AddAttr<std::string>("activation", "The activation type of FC operator.")
         .SetDefault("identity")
         .InEnum({"identity", "sigmoid", "softmax"});
+    AddAttr<std::vector<int>>("xNumColDims", "");
+    AddAttr<std::vector<int>>("wNumColDims", "");
 
     AddComment(R"DOC(
 Fully Connected Operator, known as Fully Connected Layer or Inner Product Layer
diff --git a/python/paddle/v2/framework/tests/test_fc_op.py b/python/paddle/v2/framework/tests/test_fc_op.py
index 00c4870997..39906c8b33 100644
--- a/python/paddle/v2/framework/tests/test_fc_op.py
+++ b/python/paddle/v2/framework/tests/test_fc_op.py
@@ -5,52 +5,61 @@ from op_test import OpTest
 
 class TestFCOp1(OpTest):
     def setUp(self):
-        self.op_type = "fc"
         x0 = np.random.random((16, 32)).astype("float32")
         w0 = np.random.random((32, 10)).astype("float32")
         b = np.random.random(10).astype("float32")
-        self.inputs = {"X": [("X0", x0)], "W": [("W0", w0)], "b": b}
+
         mul_out0 = np.dot(x0, w0)
         sum_out = mul_out0
         add_out = sum_out + b
         identity_out = add_out
+
+        self.op_type = "fc"
+        self.inputs = {"X": [("X0", x0)], "W": [("W0", w0)], "B": b}
         self.outputs = {
-            "mul_out": [("mul_out0", mul_out0)],
-            "sum_out": sum_out,
-            "add_out": add_out,
+            "MulOut": [("MulOut0", mul_out0)],
+            "SumOut": sum_out,
+            "AddOut": add_out,
             "Y": identity_out
         }
+        self.attrs = {"xNumColDims": [1], "wNumColDims": [1]}
 
     def test_check_output(self):
         self.check_output()
 
     def test_check_grad(self):
-        self.check_grad(["X0", "W0", "b"], "Y", max_relative_error=0.01)
+        self.check_grad(["X0", "W0", "B"], "Y", max_relative_error=0.01)
 
 
 class TestFCOp2(OpTest):
     def setUp(self):
-        self.op_type = "fc"
-        x0 = np.random.random((16, 32)).astype("float32")
+        x0 = np.random.random((16, 4, 8)).astype("float32")
         x1 = np.random.random((16, 32)).astype("float32")
         w0 = np.random.random((32, 10)).astype("float32")
-        w1 = np.random.random((32, 10)).astype("float32")
+        w1 = np.random.random((4, 8, 10)).astype("float32")
         b = np.random.random(10).astype("float32")
+
+        mul_out0 = np.dot(x0.reshape(16, 4 * 8), w0)
+        mul_out1 = np.dot(x1, w1.reshape(4 * 8, 10))
+        sum_out = mul_out0 + mul_out1
+        add_out = np.add(sum_out, b)
+        sigmoid_out = 1 / (1 + np.exp(-add_out))
+
+        self.op_type = "fc"
         self.inputs = {
             "X": [("X0", x0), ("X1", x1)],
             "W": [("W0", w0), ("W1", w1)],
-            "b": b
+            "B": b
+        }
+        self.attrs = {
+            "xNumColDims": [1, 1],
+            "wNumColDims": [1, 2],
+            "activation": "sigmoid"
         }
-        self.attrs = {"activation": "sigmoid"}
-        mul_out0 = np.dot(x0, w0)
-        mul_out1 = np.dot(x1, w1)
-        sum_out = mul_out0 + mul_out1
-        add_out = np.add(sum_out, b)
-        sigmoid_out = 1 / (1 + np.exp(-add_out))
         self.outputs = {
-            "mul_out": [("mul_out0", mul_out0), ("mul_out1", mul_out1)],
-            "sum_out": sum_out,
-            "add_out": add_out,
+            "MulOut": [("MulOut0", mul_out0), ("MulOut1", mul_out1)],
+            "SumOut": sum_out,
+            "AddOut": add_out,
             "Y": sigmoid_out
         }
 
@@ -59,7 +68,7 @@ class TestFCOp2(OpTest):
 
     def test_check_grad(self):
         self.check_grad(
-            ["X0", "X1", "W0", "W1", "b"], "Y", max_relative_error=0.01)
+            ["X0", "X1", "W0", "W1", "B"], "Y", max_relative_error=0.01)
 
 
 if __name__ == '__main__':

From 0b21b854ecded15161d5281f1d44eb0867bfac92 Mon Sep 17 00:00:00 2001
From: Liu Yiqun <liuyiqun01@baidu.com>
Date: Thu, 14 Sep 2017 03:07:48 +0000
Subject: [PATCH 09/34] Make the weights of FCOp a fixed 2-D matrix and refine
 some comments in FCOp.

---
 paddle/operators/fc_op.cc                     | 70 ++++++++++++-------
 .../paddle/v2/framework/tests/test_fc_op.py   | 14 ++--
 2 files changed, 49 insertions(+), 35 deletions(-)

diff --git a/paddle/operators/fc_op.cc b/paddle/operators/fc_op.cc
index 3e6cd8f76a..be0fca3c71 100644
--- a/paddle/operators/fc_op.cc
+++ b/paddle/operators/fc_op.cc
@@ -41,21 +41,16 @@ class FCOp : public NetOp {
                       "The size of inputs X(%d) should be no less than 1.", n);
 
     auto x_num_col_dims = Attr<std::vector<int>>("xNumColDims");
-    auto w_num_col_dims = Attr<std::vector<int>>("wNumColDims");
     PADDLE_ENFORCE_EQ(x_num_col_dims.size(), n,
                       "The size of attribute xNumColDims(%d) should be the "
                       "same as that of inputs X(%d).",
                       x_num_col_dims.size(), n);
-    PADDLE_ENFORCE_EQ(w_num_col_dims.size(), n,
-                      "The size of attribute wNumColDims(%d) should be the "
-                      "same as that of inputs X(%d).",
-                      w_num_col_dims.size(), n)
 
     // mul_out[i] = X[i] * W[i]
     for (size_t i = 0; i < n; i++) {
       framework::AttributeMap mul_attr;
       mul_attr["x_num_col_dims"] = static_cast<int>(x_num_col_dims[i]);
-      mul_attr["y_num_col_dims"] = static_cast<int>(w_num_col_dims[i]);
+      mul_attr["y_num_col_dims"] = static_cast<int>(1);
       AppendOp(
           framework::OpRegistry::CreateOp("mul", {{"X", {x[i]}}, {"Y", {w[i]}}},
                                           {{"Out", {mul_out[i]}}}, mul_attr));
@@ -95,30 +90,54 @@ class FCOpMaker : public framework::OpProtoAndCheckerMaker {
  public:
   FCOpMaker(framework::OpProto *proto, framework::OpAttrChecker *op_checker)
       : OpProtoAndCheckerMaker(proto, op_checker) {
-    AddInput("X", "The inputs of FC operator, a ordered vector of 2-D matrix.")
+    AddInput("X",
+             "(A vector of Tensors) each input Tensor can be of arbitrary "
+             "dimension, and will be reshaped to a 2-D matrix of size "
+             "(minibatch, number_of_input_features) according to attribute "
+             "xNumColDims.")
         .AsDuplicable();
-    AddInput("W", "The weights of FC operator, a ordered vector of 2-D matrix.")
+    AddInput("W",
+             "(A vector of Tensors) the weights of FC operator, a "
+             "vector of 2-D matrix of size "
+             "(number_of_input_features, number_of_neurons).")
         .AsDuplicable();
-    AddInput("B", "The 1-D bias vector of FC operator");
+    AddInput("B",
+             "(Tensor) the bias of FC operator, a 1-D vector of size "
+             "number_of_neurons.");
 
-    AddOutput("Y", "The activated output matrix of FC operator");
+    AddOutput("Y",
+              "(Tensor) the activated output matrix of FC operator, a 2-D "
+              "matrix of size (minibatch, number_of_neurons).");
     AddOutput("MulOut",
-              "The intermediate outputs of FC operator, "
-              "saving the product of X[i] * W[i]")
+              "(A vector of Tensors) the intermediate outputs of FC operator, "
+              "each Tensor saving the product of X_i * W_i.")
         .AsIntermediate()
         .AsDuplicable();
-    AddOutput("SumOut",
-              "The intermediate output of FC operator, "
-              "saving the sum of products, sum(X[i] * W[i])")
+    AddOutput(
+        "SumOut",
+        "(Tensor) the intermediate output of FC operator, "
+        "saving the sum of the products of X and W, that is sum{X_i * W_i}.")
         .AsIntermediate();
     AddOutput("AddOut",
-              "The non-actived output of FC operator, saving X * W + b")
+              "(Tensor) the non-actived output of FC operator, "
+              "saving sum{X_i * W_i} + B.")
         .AsIntermediate();
-    AddAttr<std::string>("activation", "The activation type of FC operator.")
+    AddAttr<std::string>(
+        "activation",
+        "(string, default identity) the activation type of FC operator.")
         .SetDefault("identity")
         .InEnum({"identity", "sigmoid", "softmax"});
-    AddAttr<std::vector<int>>("xNumColDims", "");
-    AddAttr<std::vector<int>>("wNumColDims", "");
+    AddAttr<std::vector<int>>(
+        "xNumColDims",
+        "(std::vector<int>) The inputs Tensors of FC operator can be of "
+        "more than 2 dimensions. In that case, each input Tensor `X_i` will be "
+        "reshaped to a 2-D matrix. The matrix's first dimension "
+        "(the length of column) will be the product of `X_i`'s last "
+        "`xNumColDims_i` dimensions, that is "
+        "`X_i.dims[0] x ... x X_i.dims[xNumColDims_i - 1]`. "
+        "The matrix's second dimension (the length of row) will be the product "
+        "of `X_i`'s first `rank - xNumColDims_i` dimensions, that is "
+        "`X_i.dims[xNumColDims_i] x ... x X_i.dims[rank - 1]`)");
 
     AddComment(R"DOC(
 Fully Connected Operator, known as Fully Connected Layer or Inner Product Layer
@@ -129,15 +148,14 @@ learned weights with a matrix multiplication followed by a bias offset
 (optionally).
 
 Equation:
-  Y = Act(sum_n{X_i * W_i} + b)
+  Y = Act(sum_n{X_i * W_i} + B)
 
-where X_i is a 2D matrix of size (M x K), usually M is the minibatch size and
-K is the number of features. W_i is also a 2D matrix of size (K x N),
-where N means the number of neurons in the fully connected layer.
-b is a 1D vector of size N. Thus, the output Y is a 2D matrix of size (M x N).
+where X_i is Tensor that will be reshaped to a 2-D matrix of size (M x K),
+usually M is the minibatch size and K is the number of input features.
+W_i is a 2-D matrix of size (K x N), where N means the number of neurons
+in the fully connected layer. B is a 1-D vector of size N.
+Thus, the output Y is a 2-D matrix of size (M x N).
 Activation type can be set to `identity` (default), `sigmoid` or `softmax`.
-
-  The config api is `paddle.v2.layer.fc`.
 )DOC");
   }
 };
diff --git a/python/paddle/v2/framework/tests/test_fc_op.py b/python/paddle/v2/framework/tests/test_fc_op.py
index 39906c8b33..f646fad337 100644
--- a/python/paddle/v2/framework/tests/test_fc_op.py
+++ b/python/paddle/v2/framework/tests/test_fc_op.py
@@ -22,7 +22,7 @@ class TestFCOp1(OpTest):
             "AddOut": add_out,
             "Y": identity_out
         }
-        self.attrs = {"xNumColDims": [1], "wNumColDims": [1]}
+        self.attrs = {"xNumColDims": [1]}
 
     def test_check_output(self):
         self.check_output()
@@ -34,13 +34,13 @@ class TestFCOp1(OpTest):
 class TestFCOp2(OpTest):
     def setUp(self):
         x0 = np.random.random((16, 4, 8)).astype("float32")
-        x1 = np.random.random((16, 32)).astype("float32")
+        x1 = np.random.random((4, 4, 32)).astype("float32")
         w0 = np.random.random((32, 10)).astype("float32")
-        w1 = np.random.random((4, 8, 10)).astype("float32")
+        w1 = np.random.random((32, 10)).astype("float32")
         b = np.random.random(10).astype("float32")
 
         mul_out0 = np.dot(x0.reshape(16, 4 * 8), w0)
-        mul_out1 = np.dot(x1, w1.reshape(4 * 8, 10))
+        mul_out1 = np.dot(x1.reshape(4 * 4, 32), w1)
         sum_out = mul_out0 + mul_out1
         add_out = np.add(sum_out, b)
         sigmoid_out = 1 / (1 + np.exp(-add_out))
@@ -51,11 +51,7 @@ class TestFCOp2(OpTest):
             "W": [("W0", w0), ("W1", w1)],
             "B": b
         }
-        self.attrs = {
-            "xNumColDims": [1, 1],
-            "wNumColDims": [1, 2],
-            "activation": "sigmoid"
-        }
+        self.attrs = {"xNumColDims": [1, 2], "activation": "sigmoid"}
         self.outputs = {
             "MulOut": [("MulOut0", mul_out0), ("MulOut1", mul_out1)],
             "SumOut": sum_out,

From fe2ab2ee7fc9c0e139edb3274e8e4e6d7889a6ab Mon Sep 17 00:00:00 2001
From: Liu Yiqun <liuyiqun01@baidu.com>
Date: Thu, 14 Sep 2017 05:49:58 +0000
Subject: [PATCH 10/34] Set the default value of xNumColDims and rename the
 output to "Out" in FCOp.

---
 paddle/operators/fc_op.cc                     | 28 +++++++++++++------
 .../paddle/v2/framework/tests/test_fc_op.py   |  9 +++---
 2 files changed, 23 insertions(+), 14 deletions(-)

diff --git a/paddle/operators/fc_op.cc b/paddle/operators/fc_op.cc
index be0fca3c71..14a7fa8467 100644
--- a/paddle/operators/fc_op.cc
+++ b/paddle/operators/fc_op.cc
@@ -41,10 +41,19 @@ class FCOp : public NetOp {
                       "The size of inputs X(%d) should be no less than 1.", n);
 
     auto x_num_col_dims = Attr<std::vector<int>>("xNumColDims");
-    PADDLE_ENFORCE_EQ(x_num_col_dims.size(), n,
-                      "The size of attribute xNumColDims(%d) should be the "
-                      "same as that of inputs X(%d).",
-                      x_num_col_dims.size(), n);
+
+    // Set all values or set no values (use the default value)
+    if (!x_num_col_dims.empty()) {
+      PADDLE_ENFORCE_EQ(x_num_col_dims.size(), n,
+                        "The size of attribute xNumColDims(%d) should be the "
+                        "same as that of inputs X(%d).",
+                        x_num_col_dims.size(), n);
+    } else {
+      x_num_col_dims.resize(n);
+      for (size_t i = 0; i < n; i++) {
+        x_num_col_dims[i] = 1;
+      }
+    }
 
     // mul_out[i] = X[i] * W[i]
     for (size_t i = 0; i < n; i++) {
@@ -81,7 +90,7 @@ class FCOp : public NetOp {
 
     auto activation = Attr<std::string>("activation");
     AppendOp(framework::OpRegistry::CreateOp(
-        activation, {{"X", {Output(add_out)}}}, {{"Y", {Output("Y")}}}, {}));
+        activation, {{"X", {Output(add_out)}}}, {{"Y", {Output("Out")}}}, {}));
     CompleteAddOp(false);
   }
 };
@@ -105,7 +114,7 @@ class FCOpMaker : public framework::OpProtoAndCheckerMaker {
              "(Tensor) the bias of FC operator, a 1-D vector of size "
              "number_of_neurons.");
 
-    AddOutput("Y",
+    AddOutput("Out",
               "(Tensor) the activated output matrix of FC operator, a 2-D "
               "matrix of size (minibatch, number_of_neurons).");
     AddOutput("MulOut",
@@ -137,7 +146,8 @@ class FCOpMaker : public framework::OpProtoAndCheckerMaker {
         "`X_i.dims[0] x ... x X_i.dims[xNumColDims_i - 1]`. "
         "The matrix's second dimension (the length of row) will be the product "
         "of `X_i`'s first `rank - xNumColDims_i` dimensions, that is "
-        "`X_i.dims[xNumColDims_i] x ... x X_i.dims[rank - 1]`)");
+        "`X_i.dims[xNumColDims_i] x ... x X_i.dims[rank - 1]`)")
+        .SetDefault(std::vector<int>{});
 
     AddComment(R"DOC(
 Fully Connected Operator, known as Fully Connected Layer or Inner Product Layer
@@ -148,13 +158,13 @@ learned weights with a matrix multiplication followed by a bias offset
 (optionally).
 
 Equation:
-  Y = Act(sum_n{X_i * W_i} + B)
+  Out = Act(sum_n{X_i * W_i} + B)
 
 where X_i is Tensor that will be reshaped to a 2-D matrix of size (M x K),
 usually M is the minibatch size and K is the number of input features.
 W_i is a 2-D matrix of size (K x N), where N means the number of neurons
 in the fully connected layer. B is a 1-D vector of size N.
-Thus, the output Y is a 2-D matrix of size (M x N).
+Thus, the output Out is a 2-D matrix of size (M x N).
 Activation type can be set to `identity` (default), `sigmoid` or `softmax`.
 )DOC");
   }
diff --git a/python/paddle/v2/framework/tests/test_fc_op.py b/python/paddle/v2/framework/tests/test_fc_op.py
index f646fad337..ed8d869a40 100644
--- a/python/paddle/v2/framework/tests/test_fc_op.py
+++ b/python/paddle/v2/framework/tests/test_fc_op.py
@@ -20,15 +20,14 @@ class TestFCOp1(OpTest):
             "MulOut": [("MulOut0", mul_out0)],
             "SumOut": sum_out,
             "AddOut": add_out,
-            "Y": identity_out
+            "Out": identity_out
         }
-        self.attrs = {"xNumColDims": [1]}
 
     def test_check_output(self):
         self.check_output()
 
     def test_check_grad(self):
-        self.check_grad(["X0", "W0", "B"], "Y", max_relative_error=0.01)
+        self.check_grad(["X0", "W0", "B"], "Out", max_relative_error=0.01)
 
 
 class TestFCOp2(OpTest):
@@ -56,7 +55,7 @@ class TestFCOp2(OpTest):
             "MulOut": [("MulOut0", mul_out0), ("MulOut1", mul_out1)],
             "SumOut": sum_out,
             "AddOut": add_out,
-            "Y": sigmoid_out
+            "Out": sigmoid_out
         }
 
     def test_check_output(self):
@@ -64,7 +63,7 @@ class TestFCOp2(OpTest):
 
     def test_check_grad(self):
         self.check_grad(
-            ["X0", "X1", "W0", "W1", "B"], "Y", max_relative_error=0.01)
+            ["X0", "X1", "W0", "W1", "B"], "Out", max_relative_error=0.01)
 
 
 if __name__ == '__main__':

From 989e8358b3803c5c15ae4ec0a3fa93fd7b915302 Mon Sep 17 00:00:00 2001
From: Liu Yiqun <liuyiqun01@baidu.com>
Date: Thu, 14 Sep 2017 06:50:08 +0000
Subject: [PATCH 11/34] Reuse the output of mul when there is only one input in
 FCOp.

---
 paddle/operators/fc_op.cc                     | 23 +++++++++++--------
 python/paddle/v2/framework/tests/op_test.py   | 14 ++++++-----
 .../paddle/v2/framework/tests/test_fc_op.py   | 16 ++++---------
 3 files changed, 25 insertions(+), 28 deletions(-)

diff --git a/paddle/operators/fc_op.cc b/paddle/operators/fc_op.cc
index 14a7fa8467..5549a836c9 100644
--- a/paddle/operators/fc_op.cc
+++ b/paddle/operators/fc_op.cc
@@ -66,22 +66,25 @@ class FCOp : public NetOp {
     }
 
     // sum_out = X[0] * W[0] + ... + X[n-1] * W[n-1]
+    auto sum_out = mul_out[0];
     if (n > 1) {
-      AppendOp(framework::OpRegistry::CreateOp(
-          "sum", {{"X", {mul_out}}}, {{"Out", {Output("SumOut")}}}, {}));
+      sum_out = Output("SumOut");
+      AppendOp(framework::OpRegistry::CreateOp("sum", {{"X", {mul_out}}},
+                                               {{"Out", {sum_out}}}, {}));
     } else {
-      AppendOp(framework::OpRegistry::CreateOp(
-          "identity", {{"X", {mul_out[0]}}}, {{"Y", {Output("SumOut")}}}, {}));
+      if (Output("SumOut") != framework::kEmptyVarName) {
+        this->Rename(Output("SumOut"), framework::kEmptyVarName);
+      }
     }
 
     // add_out = sum_out + b
     auto b = Input("B");
-    std::string add_out = "SumOut";
+    auto add_out = sum_out;
     if (b != framework::kEmptyVarName) {
-      add_out = "AddOut";
+      add_out = Output("AddOut");
       AppendOp(framework::OpRegistry::CreateOp(
-          "rowwise_add", {{"X", {Output("SumOut")}}, {"b", {Input("B")}}},
-          {{"Out", {Output(add_out)}}}, {}));
+          "rowwise_add", {{"X", {sum_out}}, {"b", {Input("B")}}},
+          {{"Out", {add_out}}}, {}));
     } else {
       if (Output("AddOut") != framework::kEmptyVarName) {
         this->Rename(Output("AddOut"), framework::kEmptyVarName);
@@ -89,8 +92,8 @@ class FCOp : public NetOp {
     }
 
     auto activation = Attr<std::string>("activation");
-    AppendOp(framework::OpRegistry::CreateOp(
-        activation, {{"X", {Output(add_out)}}}, {{"Y", {Output("Out")}}}, {}));
+    AppendOp(framework::OpRegistry::CreateOp(activation, {{"X", {add_out}}},
+                                             {{"Y", {Output("Out")}}}, {}));
     CompleteAddOp(false);
   }
 };
diff --git a/python/paddle/v2/framework/tests/op_test.py b/python/paddle/v2/framework/tests/op_test.py
index c6e4c59881..41690961b5 100644
--- a/python/paddle/v2/framework/tests/op_test.py
+++ b/python/paddle/v2/framework/tests/op_test.py
@@ -193,12 +193,14 @@ class OpTest(unittest.TestCase):
                             actual, expect, atol=1e-05),
                         "output name: " + out_name + " has diff")
             else:
-                actual = np.array(self.scope.find_var(out_name).get_tensor())
-                expect = self.outputs[out_name]
-                self.assertTrue(
-                    np.allclose(
-                        actual, expect, atol=1e-05),
-                    "output name: " + out_name + " has diff")
+                var = self.scope.find_var(out_name)
+                if var is not None:
+                    actual = np.array(var.get_tensor())
+                    expect = self.outputs[out_name]
+                    self.assertTrue(
+                        np.allclose(
+                            actual, expect, atol=1e-05),
+                        "output name: " + out_name + " has diff")
 
     def check_output(self):
         places = [core.CPUPlace()]
diff --git a/python/paddle/v2/framework/tests/test_fc_op.py b/python/paddle/v2/framework/tests/test_fc_op.py
index ed8d869a40..9f56fe5049 100644
--- a/python/paddle/v2/framework/tests/test_fc_op.py
+++ b/python/paddle/v2/framework/tests/test_fc_op.py
@@ -7,27 +7,19 @@ class TestFCOp1(OpTest):
     def setUp(self):
         x0 = np.random.random((16, 32)).astype("float32")
         w0 = np.random.random((32, 10)).astype("float32")
-        b = np.random.random(10).astype("float32")
 
         mul_out0 = np.dot(x0, w0)
-        sum_out = mul_out0
-        add_out = sum_out + b
-        identity_out = add_out
+        identity_out = mul_out0
 
         self.op_type = "fc"
-        self.inputs = {"X": [("X0", x0)], "W": [("W0", w0)], "B": b}
-        self.outputs = {
-            "MulOut": [("MulOut0", mul_out0)],
-            "SumOut": sum_out,
-            "AddOut": add_out,
-            "Out": identity_out
-        }
+        self.inputs = {"X": [("X0", x0)], "W": [("W0", w0)]}
+        self.outputs = {"MulOut": [("MulOut0", mul_out0)], "Out": identity_out}
 
     def test_check_output(self):
         self.check_output()
 
     def test_check_grad(self):
-        self.check_grad(["X0", "W0", "B"], "Out", max_relative_error=0.01)
+        self.check_grad(["X0", "W0"], "Out", max_relative_error=0.01)
 
 
 class TestFCOp2(OpTest):

From 9e7da6799bd52891685c9b0f982749cc6d8dc5d0 Mon Sep 17 00:00:00 2001
From: hedaoyuan <hedaoyuan@github.com>
Date: Fri, 15 Sep 2017 10:55:09 +0800
Subject: [PATCH 12/34] Add a neon::relu.

---
 paddle/math/BaseMatrix.cu     |  8 +++++
 paddle/math/NEONFunctions.cpp | 55 +++++++++++++++++++++++++++++++++++
 paddle/math/NEONFunctions.h   | 23 +++++++++++++++
 3 files changed, 86 insertions(+)
 create mode 100644 paddle/math/NEONFunctions.cpp
 create mode 100644 paddle/math/NEONFunctions.h

diff --git a/paddle/math/BaseMatrix.cu b/paddle/math/BaseMatrix.cu
index 5435808fb7..53dd538360 100644
--- a/paddle/math/BaseMatrix.cu
+++ b/paddle/math/BaseMatrix.cu
@@ -17,6 +17,7 @@ limitations under the License. */
 #include <cmath>
 #include "BaseMatrix.h"
 #include "MathFunctions.h"
+#include "NEONFunctions.h"
 #include "SIMDFunctions.h"
 #include "hl_matrix_apply.cuh"
 #include "hl_matrix_base.cuh"
@@ -666,6 +667,13 @@ void BaseMatrixT<T>::relu(BaseMatrixT& b) {
   applyBinary(binary::Relu<T>(), b);
 }
 
+#if defined(__ARM_NEON__) || defined(__ARM_NEON)
+template <>
+void BaseMatrixT<float>::relu(BaseMatrixT& b) {
+  neon::relu(data_, b.data_, height_ * width_);
+}
+#endif
+
 DEFINE_MATRIX_BINARY_OP(ReluDerivative, a *= (b > 0.0f ? 1.0f : 0.0f));
 template <class T>
 void BaseMatrixT<T>::reluDerivative(BaseMatrixT& b) {
diff --git a/paddle/math/NEONFunctions.cpp b/paddle/math/NEONFunctions.cpp
new file mode 100644
index 0000000000..3bf47901f1
--- /dev/null
+++ b/paddle/math/NEONFunctions.cpp
@@ -0,0 +1,55 @@
+/* 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. */
+
+#if defined(__ARM_NEON__) || defined(__ARM_NEON)
+
+#include "NEONFunctions.h"
+#include <arm_neon.h>
+
+namespace paddle {
+namespace neon {
+
+// b[i] = a[i] > 0.0f ? a[i] : 0.0f
+void relu(const float* a, float* b, int len) {
+  int offset = len % 16;
+  float32x4_t ma0, ma1, ma2, ma3;
+  float32x4_t mb0, mb1, mb2, mb3;
+
+  float32x4_t zero = vdupq_n_f32(0.f);
+  for (int k = 0; k < len / 16; k++, a += 16, b += 16) {
+    ma0 = vld1q_f32(a);
+    ma1 = vld1q_f32(a + 4);
+    ma2 = vld1q_f32(a + 8);
+    ma3 = vld1q_f32(a + 12);
+
+    mb0 = vmaxq_f32(ma0, zero);
+    mb1 = vmaxq_f32(ma1, zero);
+    mb2 = vmaxq_f32(ma2, zero);
+    mb3 = vmaxq_f32(ma3, zero);
+
+    vst1q_f32(b, mb0);
+    vst1q_f32(b + 4, mb1);
+    vst1q_f32(b + 8, mb2);
+    vst1q_f32(b + 12, mb3);
+  }
+
+  for (int i = 0; i < offset; i++) {
+    b[i] = a[i] > 0.0f ? a[i] : 0.0f;
+  }
+}
+
+}  // namespace neon
+}  // namespace paddle
+
+#endif
diff --git a/paddle/math/NEONFunctions.h b/paddle/math/NEONFunctions.h
new file mode 100644
index 0000000000..69085e3335
--- /dev/null
+++ b/paddle/math/NEONFunctions.h
@@ -0,0 +1,23 @@
+/* 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
+
+namespace paddle {
+namespace neon {
+
+void relu(const float* a, float* b, int len);
+
+}  // namespace neon
+}  // namespace paddle

From 2a110af3097ed04cf492fbe565602f8ccb71c109 Mon Sep 17 00:00:00 2001
From: hedaoyuan <hedaoyuan@github.com>
Date: Fri, 15 Sep 2017 16:11:13 +0800
Subject: [PATCH 13/34] Use Matrix::addBias.

---
 paddle/gserver/layers/ExpandConvLayer.cpp | 19 +++++++++++++------
 1 file changed, 13 insertions(+), 6 deletions(-)

diff --git a/paddle/gserver/layers/ExpandConvLayer.cpp b/paddle/gserver/layers/ExpandConvLayer.cpp
index 20de475fc3..7c1190a414 100644
--- a/paddle/gserver/layers/ExpandConvLayer.cpp
+++ b/paddle/gserver/layers/ExpandConvLayer.cpp
@@ -155,11 +155,12 @@ void ExpandConvLayer::forward(PassType passType) {
 
   /* add the bias-vector */
   if (biases_.get()) {
-    if (sharedBiases_) {
-      addSharedBias();
-    } else {
-      addUnsharedBias();
-    }
+    MatrixPtr bias = Matrix::create(biases_->getW()->getData(),
+                                    1,
+                                    biases_->getW()->getElementCnt(),
+                                    false,
+                                    useGpu_);
+    output_.value->addBias(*bias, 1.0, sharedBiases_);
   }
 
   /* activation */
@@ -171,7 +172,13 @@ void ExpandConvLayer::backward(const UpdateCallback &callback) {
 
   MatrixPtr outGrad = getOutputGrad();
   if (biases_ && biases_->getWGrad()) {
-    bpropBiases(outGrad);
+    // bpropBiases(outGrad);
+    MatrixPtr bias = Matrix::create(biases_->getWGrad()->getData(),
+                                    1,
+                                    biases_->getWGrad()->getElementCnt(),
+                                    false,
+                                    useGpu_);
+    bias->collectBias(*getOutputGrad(), 1, sharedBiases_);
     /* Increasing the number of gradient */
     biases_->getParameterPtr()->incUpdate(callback);
   }

From eef1ccbf08605f6fb784a472540ba9c1cc959a67 Mon Sep 17 00:00:00 2001
From: Liu Yiqun <liuyiqun01@baidu.com>
Date: Fri, 15 Sep 2017 08:18:28 +0000
Subject: [PATCH 14/34] Add the check of inputs and outputs in all operators.

---
 paddle/operators/accuracy_op.cc               | 11 +++++++---
 paddle/operators/add_op.cc                    |  7 +++++++
 paddle/operators/concat_op.cc                 |  3 +++
 paddle/operators/cond_op.cc                   | 21 ++++++++++++++-----
 paddle/operators/cos_sim_op.cc                | 12 +++++++++--
 paddle/operators/elementwise_mul_op.cc        | 10 +++++++--
 paddle/operators/fill_zeros_like_op.cc        |  7 +++++++
 paddle/operators/gather_op.cc                 |  7 +++++++
 paddle/operators/gaussian_random_op.cc        |  8 +++++--
 paddle/operators/identity_op.cc               |  5 +++++
 paddle/operators/lookup_table_op.cc           | 15 +++++++++----
 paddle/operators/mean_op.cc                   |  4 +++-
 paddle/operators/minus_op.cc                  |  7 +++++++
 paddle/operators/mul_op.cc                    |  7 +++++++
 paddle/operators/onehot_cross_entropy_op.cc   | 10 +++++++++
 paddle/operators/pad_op.cc                    |  5 +++++
 paddle/operators/reshape_op.cc                |  6 +++++-
 paddle/operators/rowwise_add_op.cc            |  7 +++++++
 paddle/operators/scale_op.cc                  |  5 +++++
 paddle/operators/scatter_op.cc                |  9 ++++++++
 paddle/operators/sequence_avg_pool_op.cc      |  9 +++++---
 paddle/operators/sgd_op.cc                    |  7 +++++++
 paddle/operators/sigmoid_op.cc                |  5 +++++
 paddle/operators/softmax_op.cc                |  5 +++++
 paddle/operators/squared_l2_distance_op.cc    | 18 ++++++++++------
 paddle/operators/sum_op.cc                    |  5 +++++
 paddle/operators/top_k_op.cc                  |  7 ++++++-
 paddle/operators/uniform_random_op.cc         |  4 ++++
 .../{test_add_two_op.py => test_add_op.py}    |  0
 .../tests/test_gaussian_random_op.py          |  2 +-
 .../v2/framework/tests/test_identity_op.py    | 20 ++++++++++++++++++
 ...ookup_table.py => test_lookup_table_op.py} |  0
 .../v2/framework/tests/test_minus_op.py       |  2 +-
 ..._op.py => test_onehot_cross_entropy_op.py} |  2 +-
 ...le_and_identity_op.py => test_scale_op.py} | 15 +------------
 .../paddle/v2/framework/tests/test_sgd_op.py  |  2 +-
 .../v2/framework/tests/test_sigmoid_op.py     |  2 +-
 .../v2/framework/tests/test_top_k_op.py       |  6 ++++++
 .../framework/tests/test_uniform_random_op.py |  2 +-
 39 files changed, 229 insertions(+), 50 deletions(-)
 rename python/paddle/v2/framework/tests/{test_add_two_op.py => test_add_op.py} (100%)
 create mode 100644 python/paddle/v2/framework/tests/test_identity_op.py
 rename python/paddle/v2/framework/tests/{test_lookup_table.py => test_lookup_table_op.py} (100%)
 rename python/paddle/v2/framework/tests/{test_cross_entropy_op.py => test_onehot_cross_entropy_op.py} (95%)
 rename python/paddle/v2/framework/tests/{test_scale_and_identity_op.py => test_scale_op.py} (56%)

diff --git a/paddle/operators/accuracy_op.cc b/paddle/operators/accuracy_op.cc
index 4a6c6381b0..0c813748b2 100644
--- a/paddle/operators/accuracy_op.cc
+++ b/paddle/operators/accuracy_op.cc
@@ -23,10 +23,15 @@ class AccuracyOp : public framework::OperatorWithKernel {
 
  protected:
   void InferShape(const framework::InferShapeContext &ctx) const override {
-    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("Inference"),
-                            "Input of Inference must be initialized.");
+    PADDLE_ENFORCE_NOT_NULL(
+        ctx.InputVar("Inference"),
+        "Input(Inference) of AccuracyOp should not be null.");
     PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("Label"),
-                            "Input of Inference must be initialized.");
+                            "Input(Label) of AccuracyOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(
+        ctx.OutputVar("Accuracy"),
+        "Output(Accuracy) of AccuracyOp should not be null.");
+
     auto *inference = ctx.Input<framework::Tensor>("Inference");
     auto *label = ctx.Input<framework::Tensor>("Label");
 
diff --git a/paddle/operators/add_op.cc b/paddle/operators/add_op.cc
index b43c09d4f0..e83c1efeaf 100644
--- a/paddle/operators/add_op.cc
+++ b/paddle/operators/add_op.cc
@@ -23,6 +23,13 @@ class AddOp : public framework::OperatorWithKernel {
 
  protected:
   void InferShape(const framework::InferShapeContext &ctx) const override {
+    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("X"),
+                            "Input(X) of AddOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("Y"),
+                            "Input(Y) of AddOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(ctx.OutputVar("Out"),
+                            "Output(Out) of AddOp should not be null.");
+
     PADDLE_ENFORCE_EQ(ctx.Input<Tensor>("X")->dims(),
                       ctx.Input<Tensor>("Y")->dims(),
                       "Two input of Add Op's dimension must be same.");
diff --git a/paddle/operators/concat_op.cc b/paddle/operators/concat_op.cc
index 72fd179354..223bb0ffe6 100644
--- a/paddle/operators/concat_op.cc
+++ b/paddle/operators/concat_op.cc
@@ -25,6 +25,9 @@ class ConcatOp : public framework::OperatorWithKernel {
 
  protected:
   void InferShape(const framework::InferShapeContext &ctx) const override {
+    PADDLE_ENFORCE_NOT_NULL(ctx.OutputVar("Out"),
+                            "Output(Out) of ConcatOp should not be null.");
+
     auto ins = ctx.MultiInput<framework::Tensor>("X");
     auto *out = ctx.Output<framework::LoDTensor>("Out");
     size_t axis = static_cast<size_t>(ctx.Attr<int>("axis"));
diff --git a/paddle/operators/cond_op.cc b/paddle/operators/cond_op.cc
index b2e1ca395d..8262a7a5c8 100644
--- a/paddle/operators/cond_op.cc
+++ b/paddle/operators/cond_op.cc
@@ -33,7 +33,8 @@ using DDim = framework::DDim;
 
 void CondOp::CreateScope(const Scope& scope) const {
   auto sub_scopes_var = scope.FindVar("SubScopes");
-  PADDLE_ENFORCE(sub_scopes_var != nullptr, "");
+  PADDLE_ENFORCE_NOT_NULL(sub_scopes_var,
+                          "Output(SubScopes) of CondOp should not be null.");
   auto sub_scopes = sub_scopes_var->GetMutable<std::vector<Scope*>>();
   auto& sub_scope = scope.NewScope();
   sub_scopes->push_back(&sub_scope);
@@ -41,7 +42,8 @@ void CondOp::CreateScope(const Scope& scope) const {
 
 void CondOp::CreateIndexTensor(const Scope& scope) const {
   auto index_tensors_var = scope.FindVar("IndexTensors");
-  PADDLE_ENFORCE(index_tensors_var != nullptr, "");
+  PADDLE_ENFORCE_NOT_NULL(index_tensors_var,
+                          "Output(IndexTensors) of CondOp should not be null.");
   auto& index_tensors =
       *index_tensors_var->GetMutable<std::vector<LoDTensor>>();
   index_tensors.push_back(LoDTensor());
@@ -49,7 +51,8 @@ void CondOp::CreateIndexTensor(const Scope& scope) const {
 
 void CondOp::InferShape(const Scope& scope) const {
   auto sub_scopes_var = scope.FindVar("SubScopes");
-  PADDLE_ENFORCE_NOT_NULL(sub_scopes_var);
+  PADDLE_ENFORCE_NOT_NULL(sub_scopes_var,
+                          "Output(SubScopes) of CondOp should not be null.");
   auto& sub_scopes = *sub_scopes_var->GetMutable<std::vector<Scope*>>();
 
   for (int i = 0; i < 2; ++i) {
@@ -63,7 +66,8 @@ void CondOp::InferShape(const Scope& scope) const {
     // branch
     CreateIndexTensor(scope);
 
-    PADDLE_ENFORCE(!Inputs("Xs").empty(), "Inputs can't be empty");
+    PADDLE_ENFORCE(!Inputs("Xs").empty(),
+                   "Inputs(Xs) of CondOp can't be empty.");
     for (auto& input : Inputs("Xs")) {
       // Create a new tensor in sub-scope for input-type tensor
       Variable* v = sub_scopes[i]->NewVar(input);
@@ -108,13 +112,18 @@ void CondOp::InferShape(const Scope& scope) const {
 void CondOp::Run(const Scope& scope,
                  const platform::DeviceContext& dev_ctx) const {
   auto* sub_scopes_var = scope.FindVar("SubScopes");
+  PADDLE_ENFORCE_NOT_NULL(sub_scopes_var,
+                          "Output(SubScopes) of CondOp should not be null.");
   auto sub_scopes = sub_scopes_var->Get<std::vector<Scope*>>();
   auto* index_tensors_var = scope.FindVar("IndexTensors");
+  PADDLE_ENFORCE_NOT_NULL(index_tensors_var,
+                          "Output(IndexTensors) of CondOp should not be null.");
   auto index_tensors = index_tensors_var->Get<std::vector<LoDTensor>>();
 
   std::string cond_name = Input("Cond");
   Variable* cond_var = scope.FindVar(cond_name);
-  PADDLE_ENFORCE_NOT_NULL(cond_var);
+  PADDLE_ENFORCE_NOT_NULL(cond_var,
+                          "Input(Cond) of CondOp should not be null.");
   const LoDTensor* cond = cond_var->GetMutable<LoDTensor>();
 
   // Step 1: get the true/false index at runtime
@@ -171,6 +180,8 @@ void CondOp::Run(const Scope& scope,
   }
 
   // Step 4: merge output results
+  PADDLE_ENFORCE(!Outputs("Outs").empty(),
+                 "Outputs(Outs) of CondOp can't be empty.");
   for (int i = 0; i < 2; ++i) {
     // i= 0/i for True and False branches respectively
     for (auto& output : Outputs("Outs")) {
diff --git a/paddle/operators/cos_sim_op.cc b/paddle/operators/cos_sim_op.cc
index 253b17d8a1..72c4464936 100644
--- a/paddle/operators/cos_sim_op.cc
+++ b/paddle/operators/cos_sim_op.cc
@@ -26,8 +26,16 @@ class CosSimOp : public framework::OperatorWithKernel {
  protected:
   void InferShape(const framework::InferShapeContext &ctx) const override {
     // notnull check
-    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("X"), "Input(X) must not be null.");
-    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("Y"), "Input(Y) must not be null.");
+    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("X"),
+                            "Input(X) of CosSimOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("Y"),
+                            "Input(Y) of CosSimOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(ctx.OutputVar("Out"),
+                            "Output(Out) of CosSimOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(ctx.OutputVar("XNorm"),
+                            "Output(XNorm) of CosSimOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(ctx.OutputVar("YNorm"),
+                            "Output(YNorm) of CosSimOp should not be null.");
 
     // shape check
     auto x_dims = ctx.Input<Tensor>("X")->dims();
diff --git a/paddle/operators/elementwise_mul_op.cc b/paddle/operators/elementwise_mul_op.cc
index e37c582adb..ee6e975b44 100644
--- a/paddle/operators/elementwise_mul_op.cc
+++ b/paddle/operators/elementwise_mul_op.cc
@@ -25,8 +25,14 @@ class ElementWiseMulOp : public framework::OperatorWithKernel {
 
  protected:
   void InferShape(const framework::InferShapeContext &ctx) const override {
-    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("X"), "Input(X) should not be null");
-    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("Y"), "Input(Y) should not be null");
+    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("X"),
+                            "Input(X) of ElementWiseMulOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("Y"),
+                            "Input(Y) of ElementWiseMulOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(
+        ctx.OutputVar("Out"),
+        "Output(Out) of ElementWiseMulOp should not be null.");
+
     auto x_dim = ctx.Input<Tensor>("X")->dims();
     auto y_dim = ctx.Input<Tensor>("Y")->dims();
     PADDLE_ENFORCE_GE(x_dim.size(), y_dim.size(),
diff --git a/paddle/operators/fill_zeros_like_op.cc b/paddle/operators/fill_zeros_like_op.cc
index 0c9734892a..ba7857cc65 100644
--- a/paddle/operators/fill_zeros_like_op.cc
+++ b/paddle/operators/fill_zeros_like_op.cc
@@ -23,6 +23,13 @@ class FillZerosLikeOp : public framework::OperatorWithKernel {
 
  protected:
   void InferShape(const framework::InferShapeContext &ctx) const override {
+    PADDLE_ENFORCE_NOT_NULL(
+        ctx.InputVar("Src"),
+        "Input(Src) of FillZerosLikeOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(
+        ctx.OutputVar("Dst"),
+        "Output(Dst) of FillZerosLikeOp should not be null.");
+
     ctx.Output<framework::LoDTensor>("Dst")->Resize(
         ctx.Input<framework::Tensor>("Src")->dims());
   }
diff --git a/paddle/operators/gather_op.cc b/paddle/operators/gather_op.cc
index 8883d6d5fe..d445b61c16 100644
--- a/paddle/operators/gather_op.cc
+++ b/paddle/operators/gather_op.cc
@@ -24,6 +24,13 @@ class GatherOp : public framework::OperatorWithKernel {
 
  protected:
   void InferShape(const framework::InferShapeContext &ctx) const override {
+    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("X"),
+                            "Input(X) of GatherOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("Index"),
+                            "Input(Index) of GatherOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(ctx.OutputVar("Out"),
+                            "Output(Out) of GatherOp should not be null.");
+
     int batch_size = ctx.Input<Tensor>("Index")->dims()[0];
     PADDLE_ENFORCE_GE(batch_size, 0, "Batch size must be >0");
     framework::DDim output_dims(ctx.Input<Tensor>("X")->dims());
diff --git a/paddle/operators/gaussian_random_op.cc b/paddle/operators/gaussian_random_op.cc
index 25b0776a37..c0e161bbc0 100644
--- a/paddle/operators/gaussian_random_op.cc
+++ b/paddle/operators/gaussian_random_op.cc
@@ -43,8 +43,12 @@ class GaussianRandomOp : public framework::OperatorWithKernel {
   using framework::OperatorWithKernel::OperatorWithKernel;
 
  protected:
-  void InferShape(const framework::InferShapeContext& context) const override {
-    auto* tensor = context.Output<framework::LoDTensor>("Out");
+  void InferShape(const framework::InferShapeContext& ctx) const override {
+    PADDLE_ENFORCE_NOT_NULL(
+        ctx.OutputVar("Out"),
+        "Output(Out) of GaussianRandomOp should not be null.");
+
+    auto* tensor = ctx.Output<framework::LoDTensor>("Out");
     auto dims = Attr<std::vector<int>>("dims");
     std::vector<int64_t> temp;
     temp.reserve(dims.size());
diff --git a/paddle/operators/identity_op.cc b/paddle/operators/identity_op.cc
index 7d9d4fa519..b67ca5f6f8 100644
--- a/paddle/operators/identity_op.cc
+++ b/paddle/operators/identity_op.cc
@@ -42,6 +42,11 @@ class IdentityOp : public NetOp {
              const framework::VariableNameMap &outputs,
              const framework::AttributeMap &attrs)
       : NetOp(type, inputs, outputs, attrs) {
+    PADDLE_ENFORCE_NE(Input("X"), framework::kEmptyVarName,
+                      "Input(X) of IdentityOp should not be null.");
+    PADDLE_ENFORCE_NE(Output("Out"), framework::kEmptyVarName,
+                      "Output(Out) of IdentityOp should not be null.");
+
     AppendOp(framework::OpRegistry::CreateOp(
         "scale", {{"X", {Input("X")}}}, {{"Out", {Output("Out")}}},
         {{"scale", static_cast<AttrType>(1)}}));
diff --git a/paddle/operators/lookup_table_op.cc b/paddle/operators/lookup_table_op.cc
index b3d15f1ec9..07f6dfabca 100644
--- a/paddle/operators/lookup_table_op.cc
+++ b/paddle/operators/lookup_table_op.cc
@@ -22,10 +22,17 @@ class LookupTableOp : public framework::OperatorWithKernel {
   using framework::OperatorWithKernel::OperatorWithKernel;
 
  protected:
-  void InferShape(const framework::InferShapeContext &context) const override {
-    auto table_t = context.Input<Tensor>("W");
-    auto ids_t = context.Input<Tensor>("Ids");
-    auto output_t = context.Output<framework::LoDTensor>("Out");
+  void InferShape(const framework::InferShapeContext &ctx) const override {
+    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("W"),
+                            "Input(W) of LookupTableOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("Ids"),
+                            "Input(Ids) of LookupTableOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(ctx.OutputVar("Out"),
+                            "Output(Out) of LookupTableOp should not be null.");
+
+    auto table_t = ctx.Input<Tensor>("W");
+    auto ids_t = ctx.Input<Tensor>("Ids");
+    auto output_t = ctx.Output<framework::LoDTensor>("Out");
 
     output_t->Resize({ids_t->dims()[0], table_t->dims()[1]});
   }
diff --git a/paddle/operators/mean_op.cc b/paddle/operators/mean_op.cc
index 3e523d31b6..7d7eeb59a2 100644
--- a/paddle/operators/mean_op.cc
+++ b/paddle/operators/mean_op.cc
@@ -24,7 +24,9 @@ class MeanOp : public framework::OperatorWithKernel {
  protected:
   void InferShape(const framework::InferShapeContext &ctx) const override {
     PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("X"),
-                            "Input of MeanOp must be initialized.");
+                            "Input(X) of MeanOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(ctx.OutputVar("Out"),
+                            "Output(Out) of MeanOp should not be null.");
     ctx.Output<framework::LoDTensor>("Out")->Resize({1});
   }
 };
diff --git a/paddle/operators/minus_op.cc b/paddle/operators/minus_op.cc
index 8a583f24ed..61fe49ce32 100644
--- a/paddle/operators/minus_op.cc
+++ b/paddle/operators/minus_op.cc
@@ -27,6 +27,13 @@ class MinusOp : public framework::OperatorWithKernel {
 
  protected:
   void InferShape(const framework::InferShapeContext &ctx) const override {
+    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("X"),
+                            "Input(X) of MinusOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("Y"),
+                            "Input(Y) of MinusOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(ctx.OutputVar("Out"),
+                            "Output(Out) of MinusOp should not be null.");
+
     auto *left_tensor = ctx.Input<framework::Tensor>("X");
     auto *right_tensor = ctx.Input<framework::Tensor>("Y");
 
diff --git a/paddle/operators/mul_op.cc b/paddle/operators/mul_op.cc
index 015e13de9a..b6d320b415 100644
--- a/paddle/operators/mul_op.cc
+++ b/paddle/operators/mul_op.cc
@@ -26,6 +26,13 @@ class MulOp : public framework::OperatorWithKernel {
 
  protected:
   void InferShape(const framework::InferShapeContext &ctx) const override {
+    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("X"),
+                            "Input(X) of MulOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("Y"),
+                            "Input(Y) of MulOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(ctx.OutputVar("Out"),
+                            "Output(Out) of MulOp should not be null.");
+
     auto x_dims = ctx.Input<Tensor>("X")->dims();
     auto y_dims = ctx.Input<Tensor>("Y")->dims();
     int x_num_col_dims = Attr<int>("x_num_col_dims");
diff --git a/paddle/operators/onehot_cross_entropy_op.cc b/paddle/operators/onehot_cross_entropy_op.cc
index a9baada1cd..f38be3549f 100644
--- a/paddle/operators/onehot_cross_entropy_op.cc
+++ b/paddle/operators/onehot_cross_entropy_op.cc
@@ -23,6 +23,16 @@ class OnehotCrossEntropyOp : public framework::OperatorWithKernel {
 
  protected:
   void InferShape(const framework::InferShapeContext &ctx) const override {
+    PADDLE_ENFORCE_NOT_NULL(
+        ctx.InputVar("X"),
+        "Input(X) of OnehotCrossEntropyOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(
+        ctx.InputVar("label"),
+        "Input(label) of OnehotCrossEntropyOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(
+        ctx.OutputVar("Y"),
+        "Output(Y) of OnehotCrossEntropyOp should not be null.");
+
     auto *X = ctx.Input<Tensor>("X");
     auto *label = ctx.Input<Tensor>("label");
 
diff --git a/paddle/operators/pad_op.cc b/paddle/operators/pad_op.cc
index 6cf7bd6f35..a0b1c6b631 100644
--- a/paddle/operators/pad_op.cc
+++ b/paddle/operators/pad_op.cc
@@ -25,6 +25,11 @@ class PadOp : public framework::OperatorWithKernel {
 
  protected:
   void InferShape(const framework::InferShapeContext &ctx) const override {
+    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("X"),
+                            "Input(X) of PadOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(ctx.OutputVar("Out"),
+                            "Output(Out) of PadOp should not be null.");
+
     auto x_dim = ctx.Input<Tensor>("X")->dims();
     auto paddings = Attr<std::vector<int>>("paddings");
     PADDLE_ENFORCE_EQ(x_dim.size() * 2, int64_t(paddings.size()),
diff --git a/paddle/operators/reshape_op.cc b/paddle/operators/reshape_op.cc
index d281702092..0d05e34414 100644
--- a/paddle/operators/reshape_op.cc
+++ b/paddle/operators/reshape_op.cc
@@ -28,7 +28,11 @@ class ReshapeOp : public framework::OperatorWithKernel {
  protected:
   void InferShape(const framework::InferShapeContext &ctx) const override {
     // input check
-    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("X"), "Input(X) shouldn't be null");
+    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("X"),
+                            "Input(X) of ReshapeOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(ctx.OutputVar("Out"),
+                            "Output(Out) of ReshapeOp should not be null.");
+
     auto shape = ctx.Attr<std::vector<int>>("shape");
     PADDLE_ENFORCE(shape.size() > 0, "Attr(shape) shouldn't be empty.");
     for (auto dim : shape) {
diff --git a/paddle/operators/rowwise_add_op.cc b/paddle/operators/rowwise_add_op.cc
index c6101685a3..2a3fd3be94 100644
--- a/paddle/operators/rowwise_add_op.cc
+++ b/paddle/operators/rowwise_add_op.cc
@@ -25,6 +25,13 @@ class RowwiseAddOp : public framework::OperatorWithKernel {
 
  protected:
   void InferShape(const framework::InferShapeContext &ctx) const override {
+    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("X"),
+                            "Input(X) of RowwiseAddOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("b"),
+                            "Input(b) of RowwiseAddOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(ctx.OutputVar("Out"),
+                            "Output(Out) of RowwiseAddOp should not be null.");
+
     auto x_dims = ctx.Input<Tensor>("X")->dims();
     auto b_dims = ctx.Input<Tensor>("b")->dims();
     PADDLE_ENFORCE_GT(
diff --git a/paddle/operators/scale_op.cc b/paddle/operators/scale_op.cc
index 35e6b70ba9..d1f42e8662 100644
--- a/paddle/operators/scale_op.cc
+++ b/paddle/operators/scale_op.cc
@@ -27,6 +27,11 @@ class ScaleOp : public framework::OperatorWithKernel {
 
  protected:
   void InferShape(const framework::InferShapeContext &ctx) const override {
+    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("X"),
+                            "Input(X) of ScaleOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(ctx.OutputVar("Out"),
+                            "Output(Out) of ScaleOp should not be null.");
+
     auto *in = ctx.Input<framework::Tensor>("X");
     auto *out = ctx.Output<framework::LoDTensor>("Out");
     out->Resize(in->dims());
diff --git a/paddle/operators/scatter_op.cc b/paddle/operators/scatter_op.cc
index 0f7510983e..8820262732 100644
--- a/paddle/operators/scatter_op.cc
+++ b/paddle/operators/scatter_op.cc
@@ -24,6 +24,15 @@ class ScatterOp : public framework::OperatorWithKernel {
 
  protected:
   void InferShape(const framework::InferShapeContext &ctx) const override {
+    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("Ref"),
+                            "Input(Ref) of ScatterOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("Index"),
+                            "Input(Index) of ScatterOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("Updates"),
+                            "Input(Updates) of ScatterOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(ctx.OutputVar("Out"),
+                            "Output(Out) of ScatterOp should not be null.");
+
     PADDLE_ENFORCE_EQ(ctx.Input<Tensor>("Index")->dims().size(), 1,
                       "Update Index should be 1-D.");
     PADDLE_ENFORCE_EQ(ctx.Input<Tensor>("Ref")->dims().size(),
diff --git a/paddle/operators/sequence_avg_pool_op.cc b/paddle/operators/sequence_avg_pool_op.cc
index c15a5833de..eb3e37655b 100644
--- a/paddle/operators/sequence_avg_pool_op.cc
+++ b/paddle/operators/sequence_avg_pool_op.cc
@@ -23,9 +23,12 @@ class SequenceAvgPoolOp : public framework::OperatorWithKernel {
 
  protected:
   void InferShape(const framework::InferShapeContext& ctx) const override {
-    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("X"),
-                            "Input of SequenceAvgPoolOp"
-                            "must be initialized.");
+    PADDLE_ENFORCE_NOT_NULL(
+        ctx.InputVar("X"), "Input(X) of SequenceAvgPoolOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(
+        ctx.OutputVar("Out"),
+        "Output(Out) of SequenceAvgPoolOp should not be null.");
+
     auto* x = ctx.Input<framework::LoDTensor>("X");
     auto dims = x->dims();
     auto lod = x->lod();
diff --git a/paddle/operators/sgd_op.cc b/paddle/operators/sgd_op.cc
index 7997bf6907..1232e64c7f 100644
--- a/paddle/operators/sgd_op.cc
+++ b/paddle/operators/sgd_op.cc
@@ -23,6 +23,13 @@ class SGDOp : public framework::OperatorWithKernel {
 
  protected:
   void InferShape(const framework::InferShapeContext &ctx) const override {
+    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("param"),
+                            "Input(param) of SGDOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("grad"),
+                            "Input(grad) of SGDOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(ctx.OutputVar("param_out"),
+                            "Output(param_out) of SGDOp should not be null.");
+
     PADDLE_ENFORCE_EQ(ctx.Input<Tensor>("param")->dims(),
                       ctx.Input<Tensor>("grad")->dims(),
                       "Two input of SGD Op's dimension must be same.");
diff --git a/paddle/operators/sigmoid_op.cc b/paddle/operators/sigmoid_op.cc
index de6a1ba773..992b19965e 100644
--- a/paddle/operators/sigmoid_op.cc
+++ b/paddle/operators/sigmoid_op.cc
@@ -23,6 +23,11 @@ class SigmoidOp : public framework::OperatorWithKernel {
 
  protected:
   void InferShape(const framework::InferShapeContext &ctx) const override {
+    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("X"),
+                            "Input(X) of SigmoidOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(ctx.OutputVar("Y"),
+                            "Output(Y) of SigmoidOp should not be null.");
+
     ctx.Output<framework::LoDTensor>("Y")->Resize(
         ctx.Input<Tensor>("X")->dims());
   }
diff --git a/paddle/operators/softmax_op.cc b/paddle/operators/softmax_op.cc
index 239d3d141e..c67eb028c8 100644
--- a/paddle/operators/softmax_op.cc
+++ b/paddle/operators/softmax_op.cc
@@ -23,6 +23,11 @@ class SoftmaxOp : public framework::OperatorWithKernel {
 
  protected:
   void InferShape(const framework::InferShapeContext &ctx) const override {
+    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("X"),
+                            "Input(X) of SoftmaxOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(ctx.OutputVar("Y"),
+                            "Output(Y) of SoftmaxOp should not be null.");
+
     PADDLE_ENFORCE(ctx.Input<Tensor>("X")->dims().size() == 2UL,
                    "The input of softmax op must be a matrix.");
     ctx.Output<framework::LoDTensor>("Y")->Resize(
diff --git a/paddle/operators/squared_l2_distance_op.cc b/paddle/operators/squared_l2_distance_op.cc
index ebe5bd352e..39f4305877 100644
--- a/paddle/operators/squared_l2_distance_op.cc
+++ b/paddle/operators/squared_l2_distance_op.cc
@@ -23,12 +23,18 @@ class SquaredL2DistanceOp : public framework::OperatorWithKernel {
 
  protected:
   void InferShape(const framework::InferShapeContext& ctx) const override {
-    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("X"),
-                            "Input of SquaredL2DistanceOp "
-                            "must be initialized.");
-    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("Y"),
-                            "Target of SquaredL2DistanceOp "
-                            "must be initialized.");
+    PADDLE_ENFORCE_NOT_NULL(
+        ctx.InputVar("X"),
+        "Input(X) of SquaredL2DistanceOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(
+        ctx.InputVar("Y"),
+        "Input(Y) of SquaredL2DistanceOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(
+        ctx.OutputVar("sub_result"),
+        "Output(sub_result) of SquaredL2DistanceOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(
+        ctx.OutputVar("Out"),
+        "Output(Out) of SquaredL2DistanceOp should not be null.");
 
     auto* x = ctx.Input<Tensor>("X");
     auto x_dims = x->dims();
diff --git a/paddle/operators/sum_op.cc b/paddle/operators/sum_op.cc
index 7170e7256c..41e05c27f9 100644
--- a/paddle/operators/sum_op.cc
+++ b/paddle/operators/sum_op.cc
@@ -22,6 +22,11 @@ class SumOp : public framework::OperatorWithKernel {
 
  protected:
   void InferShape(const framework::InferShapeContext &ctx) const override {
+    PADDLE_ENFORCE(!ctx.MultiInputVar("X").empty(),
+                   "Input(X) of SumOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(ctx.OutputVar("Out"),
+                            "Output(Out) of SumOp should not be null.");
+
     auto ins = ctx.MultiInput<framework::Tensor>("X");
     auto *out = ctx.Output<framework::LoDTensor>("Out");
     int N = ins.size();
diff --git a/paddle/operators/top_k_op.cc b/paddle/operators/top_k_op.cc
index ff0e77a344..169b815fef 100644
--- a/paddle/operators/top_k_op.cc
+++ b/paddle/operators/top_k_op.cc
@@ -24,7 +24,12 @@ class TopkOp : public framework::OperatorWithKernel {
  protected:
   void InferShape(const framework::InferShapeContext &ctx) const override {
     PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("X"),
-                            "Input of TopkOP must be initialized.");
+                            "Input(X) of TopkOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(ctx.OutputVar("Out"),
+                            "Output(Out) of TopkOp should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(ctx.OutputVar("Indices"),
+                            "Output(Indices) of TopkOp should not be null.");
+
     auto *input = ctx.Input<framework::Tensor>("X");
     const int k = static_cast<int>(ctx.Attr<int>("k"));
 
diff --git a/paddle/operators/uniform_random_op.cc b/paddle/operators/uniform_random_op.cc
index ed79736936..184bcbc29c 100644
--- a/paddle/operators/uniform_random_op.cc
+++ b/paddle/operators/uniform_random_op.cc
@@ -48,6 +48,10 @@ class UniformRandomOp : public framework::OperatorWithKernel {
 
  protected:
   void InferShape(const framework::InferShapeContext& ctx) const override {
+    PADDLE_ENFORCE_NOT_NULL(
+        ctx.OutputVar("Out"),
+        "Output(Out) of UniformRandomOp should not be null.");
+
     PADDLE_ENFORCE(Attr<float>("min") < Attr<float>("max"),
                    "uniform_random's min must less then max");
     auto* tensor = ctx.Output<framework::LoDTensor>("Out");
diff --git a/python/paddle/v2/framework/tests/test_add_two_op.py b/python/paddle/v2/framework/tests/test_add_op.py
similarity index 100%
rename from python/paddle/v2/framework/tests/test_add_two_op.py
rename to python/paddle/v2/framework/tests/test_add_op.py
diff --git a/python/paddle/v2/framework/tests/test_gaussian_random_op.py b/python/paddle/v2/framework/tests/test_gaussian_random_op.py
index 1f9e4db783..1888ee28f9 100644
--- a/python/paddle/v2/framework/tests/test_gaussian_random_op.py
+++ b/python/paddle/v2/framework/tests/test_gaussian_random_op.py
@@ -4,7 +4,7 @@ from paddle.v2.framework.op import Operator
 import numpy
 
 
-class GaussianRandomTest(unittest.TestCase):
+class TestGaussianRandomOp(unittest.TestCase):
     def test_cpu(self):
         self.gaussian_random_test(place=core.CPUPlace())
 
diff --git a/python/paddle/v2/framework/tests/test_identity_op.py b/python/paddle/v2/framework/tests/test_identity_op.py
new file mode 100644
index 0000000000..2e95e7c786
--- /dev/null
+++ b/python/paddle/v2/framework/tests/test_identity_op.py
@@ -0,0 +1,20 @@
+import unittest
+import numpy as np
+from op_test import OpTest
+
+
+class TestIdentityOp(OpTest):
+    def setUp(self):
+        self.op_type = "identity"
+        self.inputs = {'X': np.random.random((10, 10)).astype("float32")}
+        self.outputs = {'Out': self.inputs['X']}
+
+    def test_check_output(self):
+        self.check_output()
+
+    def test_check_grad(self):
+        self.check_grad(['X'], 'Out')
+
+
+if __name__ == "__main__":
+    unittest.main()
diff --git a/python/paddle/v2/framework/tests/test_lookup_table.py b/python/paddle/v2/framework/tests/test_lookup_table_op.py
similarity index 100%
rename from python/paddle/v2/framework/tests/test_lookup_table.py
rename to python/paddle/v2/framework/tests/test_lookup_table_op.py
diff --git a/python/paddle/v2/framework/tests/test_minus_op.py b/python/paddle/v2/framework/tests/test_minus_op.py
index dea797a1fe..c56d7cb548 100644
--- a/python/paddle/v2/framework/tests/test_minus_op.py
+++ b/python/paddle/v2/framework/tests/test_minus_op.py
@@ -3,7 +3,7 @@ import numpy as np
 from op_test import OpTest
 
 
-class MinusOpTest(OpTest):
+class TestMinusOp(OpTest):
     def setUp(self):
         self.op_type = "minus"
         self.inputs = {
diff --git a/python/paddle/v2/framework/tests/test_cross_entropy_op.py b/python/paddle/v2/framework/tests/test_onehot_cross_entropy_op.py
similarity index 95%
rename from python/paddle/v2/framework/tests/test_cross_entropy_op.py
rename to python/paddle/v2/framework/tests/test_onehot_cross_entropy_op.py
index 253e7b8a24..fd3cbdb803 100644
--- a/python/paddle/v2/framework/tests/test_cross_entropy_op.py
+++ b/python/paddle/v2/framework/tests/test_onehot_cross_entropy_op.py
@@ -3,7 +3,7 @@ import numpy
 from op_test import OpTest
 
 
-class TestCrossEntropy(OpTest):
+class TestOnehotCrossEntropyOp(OpTest):
     def setUp(self):
         self.op_type = "onehot_cross_entropy"
         batch_size = 30
diff --git a/python/paddle/v2/framework/tests/test_scale_and_identity_op.py b/python/paddle/v2/framework/tests/test_scale_op.py
similarity index 56%
rename from python/paddle/v2/framework/tests/test_scale_and_identity_op.py
rename to python/paddle/v2/framework/tests/test_scale_op.py
index 05d76d4282..2ea1e18547 100644
--- a/python/paddle/v2/framework/tests/test_scale_and_identity_op.py
+++ b/python/paddle/v2/framework/tests/test_scale_op.py
@@ -3,20 +3,7 @@ import numpy as np
 from op_test import OpTest
 
 
-class IdentityTest(OpTest):
-    def setUp(self):
-        self.op_type = "identity"
-        self.inputs = {'X': np.random.random((10, 10)).astype("float32")}
-        self.outputs = {'Out': self.inputs['X']}
-
-    def test_check_output(self):
-        self.check_output()
-
-    def test_check_grad(self):
-        self.check_grad(['X'], 'Out')
-
-
-class ScaleTest(OpTest):
+class TestScaleOp(OpTest):
     def setUp(self):
         self.op_type = "scale"
         self.inputs = {'X': np.random.random((10, 10)).astype("float32")}
diff --git a/python/paddle/v2/framework/tests/test_sgd_op.py b/python/paddle/v2/framework/tests/test_sgd_op.py
index 557cf15ace..64e54d1500 100644
--- a/python/paddle/v2/framework/tests/test_sgd_op.py
+++ b/python/paddle/v2/framework/tests/test_sgd_op.py
@@ -3,7 +3,7 @@ import numpy as np
 from op_test import OpTest
 
 
-class TestSGD(OpTest):
+class TestSGDOp(OpTest):
     def setUp(self):
         self.op_type = "sgd"
         w = np.random.random((102, 105)).astype("float32")
diff --git a/python/paddle/v2/framework/tests/test_sigmoid_op.py b/python/paddle/v2/framework/tests/test_sigmoid_op.py
index 2316e49eff..d65d887db4 100644
--- a/python/paddle/v2/framework/tests/test_sigmoid_op.py
+++ b/python/paddle/v2/framework/tests/test_sigmoid_op.py
@@ -3,7 +3,7 @@ import numpy as np
 from op_test import OpTest
 
 
-class TestSigmoid(OpTest):
+class TestSigmoidOp(OpTest):
     def setUp(self):
         self.op_type = "sigmoid"
         self.inputs = {
diff --git a/python/paddle/v2/framework/tests/test_top_k_op.py b/python/paddle/v2/framework/tests/test_top_k_op.py
index cab799256d..694f37d612 100644
--- a/python/paddle/v2/framework/tests/test_top_k_op.py
+++ b/python/paddle/v2/framework/tests/test_top_k_op.py
@@ -21,6 +21,9 @@ class TestTopkOp(OpTest):
 
         self.outputs = {'Out': output, 'Indices': indices}
 
+    def test_check_output(self):
+        self.check_output()
+
 
 class TestTopkOp3d(OpTest):
     def setUp(self):
@@ -42,6 +45,9 @@ class TestTopkOp3d(OpTest):
 
         self.outputs = {'Out': output, 'Indices': indices}
 
+    def test_check_output(self):
+        self.check_output()
+
 
 if __name__ == "__main__":
     unittest.main()
diff --git a/python/paddle/v2/framework/tests/test_uniform_random_op.py b/python/paddle/v2/framework/tests/test_uniform_random_op.py
index 76a5e36e56..9e8898fb59 100644
--- a/python/paddle/v2/framework/tests/test_uniform_random_op.py
+++ b/python/paddle/v2/framework/tests/test_uniform_random_op.py
@@ -4,7 +4,7 @@ import paddle.v2.framework.core as core
 import numpy
 
 
-class UniformRandomTest(unittest.TestCase):
+class TestUniformRandomOp(unittest.TestCase):
     def test_uniform_random_cpu(self):
         self.uniform_random_test(place=core.CPUPlace())
 

From b3a50d53c90ebf3174867abf66ff24c3dd82724b Mon Sep 17 00:00:00 2001
From: hedaoyuan <hedaoyuan@github.com>
Date: Fri, 15 Sep 2017 17:03:14 +0800
Subject: [PATCH 15/34] Remove useless code.

---
 paddle/gserver/layers/ExpandConvBaseLayer.cpp | 63 -------------------
 paddle/gserver/layers/ExpandConvBaseLayer.h   | 13 ----
 2 files changed, 76 deletions(-)

diff --git a/paddle/gserver/layers/ExpandConvBaseLayer.cpp b/paddle/gserver/layers/ExpandConvBaseLayer.cpp
index 2b7bef0a75..5cf7de41f9 100644
--- a/paddle/gserver/layers/ExpandConvBaseLayer.cpp
+++ b/paddle/gserver/layers/ExpandConvBaseLayer.cpp
@@ -58,67 +58,4 @@ size_t ExpandConvBaseLayer::getOutputSize() {
   return layerSize;
 }
 
-void ExpandConvBaseLayer::addSharedBias() {
-  size_t mapW = getOutputSize() / numFilters_;
-  size_t mapH = getOutputValue()->getElementCnt() / mapW;
-  MatrixPtr out =
-      Matrix::create(getOutputValue()->getData(), mapH, mapW, false, useGpu_);
-
-  Matrix::resizeOrCreate(transOutValue_, mapW, mapH, false, useGpu_);
-
-  out->transpose(transOutValue_, false);  // false means no memory allocation
-  transOutValue_->reshape(transOutValue_->getElementCnt() / numFilters_,
-                          numFilters_);
-
-  MatrixPtr bias = Matrix::create(biases_->getW()->getData(),
-                                  1,
-                                  biases_->getW()->getElementCnt(),
-                                  false,
-                                  useGpu_);
-  transOutValue_->addBias(*bias, 1.0f);
-
-  transOutValue_->reshape(mapW, mapH);
-  transOutValue_->transpose(out, false);  // false means no memory allocation
-
-  out->clear();
-  bias->clear();
-}
-
-void ExpandConvBaseLayer::addUnsharedBias() {
-  MatrixPtr outValue = getOutputValue();
-  MatrixPtr bias = Matrix::create(biases_->getW()->getData(),
-                                  1,
-                                  biases_->getW()->getElementCnt(),
-                                  false,
-                                  useGpu_);
-  outValue->addBias(*bias, 1.0f);
-}
-
-void ExpandConvBaseLayer::bpropSharedBias(MatrixPtr biases, MatrixPtr v) {
-  size_t mapW = getOutputSize() / numFilters_;
-  size_t mapH = v->getElementCnt() / mapW;
-  MatrixPtr vTmp = Matrix::create(v->getData(), mapH, mapW, false, useGpu_);
-
-  Matrix::resizeOrCreate(transOutValue_, mapW, mapH, false, useGpu_);
-
-  vTmp->transpose(transOutValue_, false);  // false means no memory allocation
-  transOutValue_->reshape(transOutValue_->getElementCnt() / numFilters_,
-                          numFilters_);
-  biases->collectBias(*transOutValue_, 1.0f);
-}
-
-void ExpandConvBaseLayer::bpropBiases(MatrixPtr v) {
-  MatrixPtr biases = Matrix::create(biases_->getWGrad()->getData(),
-                                    1,
-                                    biases_->getWGrad()->getElementCnt(),
-                                    false,
-                                    useGpu_);
-  if (sharedBiases_) {
-    bpropSharedBias(biases, v);
-  } else {
-    biases->collectBias(*v, 1.0f);
-  }
-  biases->clear();
-}
-
 }  // namespace paddle
diff --git a/paddle/gserver/layers/ExpandConvBaseLayer.h b/paddle/gserver/layers/ExpandConvBaseLayer.h
index 01c699d234..74b3296119 100644
--- a/paddle/gserver/layers/ExpandConvBaseLayer.h
+++ b/paddle/gserver/layers/ExpandConvBaseLayer.h
@@ -39,19 +39,6 @@ public:
             const ParameterMap& parameterMap) override;
 
   size_t getOutputSize();
-
-  /**
-   * Add shared bias.
-   */
-  void addSharedBias();
-
-  /**
-   * Add unshared bias.
-   */
-  void addUnsharedBias();
-
-  void bpropSharedBias(MatrixPtr biases, MatrixPtr v);
-  void bpropBiases(MatrixPtr v);
 };
 
 }  // namespace paddle

From 3616ef77390b250e543cc7ba177860905f1ec975 Mon Sep 17 00:00:00 2001
From: hedaoyuan <hedaoyuan@github.com>
Date: Fri, 15 Sep 2017 17:52:26 +0800
Subject: [PATCH 16/34] Remove ExpandConvBaseLayer.

---
 paddle/gserver/layers/ExpandConvLayer.cpp | 36 ++++++++++++++++++++++-
 paddle/gserver/layers/ExpandConvLayer.h   |  7 +++--
 2 files changed, 39 insertions(+), 4 deletions(-)

diff --git a/paddle/gserver/layers/ExpandConvLayer.cpp b/paddle/gserver/layers/ExpandConvLayer.cpp
index 7c1190a414..abc36aaef7 100644
--- a/paddle/gserver/layers/ExpandConvLayer.cpp
+++ b/paddle/gserver/layers/ExpandConvLayer.cpp
@@ -36,7 +36,35 @@ inline bool isDepthwiseConv(int channels, int groups) {
 bool ExpandConvLayer::init(const LayerMap &layerMap,
                            const ParameterMap &parameterMap) {
   /* Initialize the basic convolutional parent class */
-  ExpandConvBaseLayer::init(layerMap, parameterMap);
+  ConvBaseLayer::init(layerMap, parameterMap);
+
+  int index = 0;
+  for (auto &inputConfig : config_.inputs()) {
+    const ConvConfig &conf = inputConfig.conv_conf();
+    /* Consistent caffe mode for multiple input */
+    caffeMode_ = conf.caffe_mode();
+
+    // create a new weight
+    size_t height, width;
+    height = filterPixels_[index] * filterChannels_[index];
+    width = (!isDeconv_) ? numFilters_ : channels_[index];
+    CHECK_EQ(parameters_[index]->getSize(), width * height);
+    Weight *w = new Weight(height, width, parameters_[index]);
+    weights_.emplace_back(w);
+    index++;
+  }
+  if (biasParameter_.get()) {
+    if (sharedBiases_) {
+      CHECK_EQ((size_t)numFilters_, biasParameter_->getSize());
+      biases_ =
+          std::unique_ptr<Weight>(new Weight(numFilters_, 1, biasParameter_));
+    } else {
+      biases_ =
+          std::unique_ptr<Weight>(new Weight(getSize(), 1, biasParameter_));
+    }
+  }
+
+  getOutputSize();
 
   size_t numInputs = config_.inputs_size();
   inputShape_.resize(numInputs);
@@ -108,6 +136,12 @@ bool ExpandConvLayer::init(const LayerMap &layerMap,
   return true;
 }
 
+size_t ExpandConvLayer::getOutputSize() {
+  CHECK_NE(inputLayers_.size(), 0UL);
+  size_t layerSize = ConvBaseLayer::calOutputSize();
+  return layerSize;
+}
+
 // i is the index of input layers
 #define BACKWARD_INPUT(i, inputs, outputs) \
   backward_[2 * i]->calc(inputs, outputs)
diff --git a/paddle/gserver/layers/ExpandConvLayer.h b/paddle/gserver/layers/ExpandConvLayer.h
index a1f943d152..698c37fb3f 100644
--- a/paddle/gserver/layers/ExpandConvLayer.h
+++ b/paddle/gserver/layers/ExpandConvLayer.h
@@ -28,10 +28,9 @@ namespace paddle {
  * The config file api is img_conv_layer.
  */
 
-class ExpandConvLayer : public ExpandConvBaseLayer {
+class ExpandConvLayer : public ConvBaseLayer {
 public:
-  explicit ExpandConvLayer(const LayerConfig& config)
-      : ExpandConvBaseLayer(config) {}
+  explicit ExpandConvLayer(const LayerConfig& config) : ConvBaseLayer(config) {}
 
   ~ExpandConvLayer() {}
 
@@ -41,6 +40,8 @@ public:
   void forward(PassType passType) override;
   void backward(const UpdateCallback& callback) override;
 
+  size_t getOutputSize();
+
 protected:
   std::vector<TensorShape> inputShape_;
   std::vector<TensorShape> filterShape_;

From 9e74b898776065fcb2504b70b1f5fc3ea137fc8c Mon Sep 17 00:00:00 2001
From: peterzhang2029 <zhangchao41@baidu.com>
Date: Fri, 15 Sep 2017 17:55:40 +0800
Subject: [PATCH 17/34] update notation in networks.py roughly

---
 .../paddle/trainer_config_helpers/networks.py | 338 +++++++++---------
 1 file changed, 171 insertions(+), 167 deletions(-)

diff --git a/python/paddle/trainer_config_helpers/networks.py b/python/paddle/trainer_config_helpers/networks.py
index 2f604ee45a..04bb9ce1d3 100644
--- a/python/paddle/trainer_config_helpers/networks.py
+++ b/python/paddle/trainer_config_helpers/networks.py
@@ -11,10 +11,7 @@
 # WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 # See the License for the specific language governing permissions and
 # limitations under the License.
-"""
 
-"""
-# from activations import *
 from activations import LinearActivation, ReluActivation, SoftmaxActivation, \
     IdentityActivation, TanhActivation, SequenceSoftmaxActivation
 from attrs import ExtraAttr
@@ -55,13 +52,13 @@ def sequence_conv_pool(input,
                        context_attr=None,
                        pool_attr=None):
     """
-    Text convolution pooling layers helper.
+    Text convolution pooling group.
 
     Text input => Context Projection => FC Layer => Pooling => Output.
 
-    :param name: name of output layer(pooling layer name)
+    :param name: group name.
     :type name: basestring
-    :param input: name of input layer
+    :param input: input layer.
     :type input: LayerOutput
     :param context_len: context projection length. See
                         context_projection's document.
@@ -70,34 +67,34 @@ def sequence_conv_pool(input,
     :type hidden_size: int
     :param context_start: context start position. See
                           context_projection's context_start.
-    :type context_start: int or None
+    :type context_start: int|None
     :param pool_type: pooling layer type. See pooling_layer's document.
-    :type pool_type: BasePoolingType.
+    :type pool_type: BasePoolingType
     :param context_proj_layer_name: context projection layer name.
                                     None if user don't care.
     :type context_proj_layer_name: basestring
-    :param context_proj_param_attr: context projection parameter attribute.
-                                    None if user don't care.
-    :type context_proj_param_attr: ParameterAttribute or None.
+    :param context_proj_param_attr: padding parameter attribute of context projection layer.
+                                    If false, it means padding always be zero.
+    :type context_proj_param_attr: ParameterAttribute|None
     :param fc_layer_name: fc layer name. None if user don't care.
     :type fc_layer_name: basestring
     :param fc_param_attr: fc layer parameter attribute. None if user don't care.
-    :type fc_param_attr: ParameterAttribute or None
+    :type fc_param_attr: ParameterAttribute|None
     :param fc_bias_attr: fc bias parameter attribute. False if no bias,
                          None if user don't care.
-    :type fc_bias_attr: ParameterAttribute or None
-    :param fc_act: fc layer activation type. None means tanh
+    :type fc_bias_attr: ParameterAttribute|False|None
+    :param fc_act: fc layer activation type. None means tanh.
     :type fc_act: BaseActivation
-    :param pool_bias_attr: pooling layer bias attr. None if don't care.
-                           False if no bias.
-    :type pool_bias_attr: ParameterAttribute or None.
+    :param pool_bias_attr: pooling layer bias attr. False if no bias.
+                           None if user don't care.
+    :type pool_bias_attr: ParameterAttribute|False|None
     :param fc_attr: fc layer extra attribute.
     :type fc_attr: ExtraLayerAttribute
     :param context_attr: context projection layer extra attribute.
     :type context_attr: ExtraLayerAttribute
     :param pool_attr: pooling layer extra attribute.
     :type pool_attr: ExtraLayerAttribute
-    :return: output layer name.
+    :return: layer's output.
     :rtype: LayerOutput
     """
     # Set Default Value to param
@@ -163,45 +160,45 @@ def simple_img_conv_pool(input,
     """
     Simple image convolution and pooling group.
 
-    Input => conv => pooling
+    Img input => Conv => Pooling => Output.
 
-    :param name: group name
+    :param name: group name.
     :type name: basestring
-    :param input: input layer name.
+    :param input: input layer.
     :type input: LayerOutput
-    :param filter_size: see img_conv_layer for details
+    :param filter_size: see img_conv_layer for details.
     :type filter_size: int
-    :param num_filters: see img_conv_layer for details
+    :param num_filters: see img_conv_layer for details.
     :type num_filters: int
-    :param pool_size: see img_pool_layer for details
+    :param pool_size: see img_pool_layer for details.
     :type pool_size: int
-    :param pool_type: see img_pool_layer for details
+    :param pool_type: see img_pool_layer for details.
     :type pool_type: BasePoolingType
-    :param act: see img_conv_layer for details
+    :param act: see img_conv_layer for details.
     :type act: BaseActivation
-    :param groups: see img_conv_layer for details
+    :param groups: see img_conv_layer for details.
     :type groups: int
-    :param conv_stride: see img_conv_layer for details
+    :param conv_stride: see img_conv_layer for details.
     :type conv_stride: int
-    :param conv_padding: see img_conv_layer for details
+    :param conv_padding: see img_conv_layer for details.
     :type conv_padding: int
-    :param bias_attr: see img_conv_layer for details
+    :param bias_attr: see img_conv_layer for details.
     :type bias_attr: ParameterAttribute
-    :param num_channel: see img_conv_layer for details
+    :param num_channel: see img_conv_layer for details.
     :type num_channel: int
-    :param param_attr: see img_conv_layer for details
+    :param param_attr: see img_conv_layer for details.
     :type param_attr: ParameterAttribute
-    :param shared_bias: see img_conv_layer for details
+    :param shared_bias: see img_conv_layer for details.
     :type shared_bias: bool
-    :param conv_layer_attr: see img_conv_layer for details
+    :param conv_layer_attr: see img_conv_layer for details.
     :type conv_layer_attr: ExtraLayerAttribute
-    :param pool_stride: see img_pool_layer for details
+    :param pool_stride: see img_pool_layer for details.
     :type pool_stride: int
-    :param pool_padding: see img_pool_layer for details
+    :param pool_padding: see img_pool_layer for details.
     :type pool_padding: int
-    :param pool_layer_attr: see img_pool_layer for details
+    :param pool_layer_attr: see img_pool_layer for details.
     :type pool_layer_attr: ExtraLayerAttribute
-    :return: Layer's output
+    :return: layer's output
     :rtype: LayerOutput
     """
     _conv_ = img_conv_layer(
@@ -252,48 +249,52 @@ def img_conv_bn_pool(input,
                      pool_layer_attr=None):
     """
     Convolution, batch normalization, pooling group.
+    
+    Img input => Conv => BN => Pooling => Output.
 
-    :param name: group name
+    :param name: group name.
     :type name: basestring
-    :param input: layer's input
-    :type input: LayerOutput
-    :param filter_size: see img_conv_layer's document
+    :param input: input layer.
+    :type input: LayerOutput 
+    :param filter_size: see img_conv_layer for details.
     :type filter_size: int
-    :param num_filters: see img_conv_layer's document
+    :param num_filters: see img_conv_layer for details.
     :type num_filters: int
-    :param pool_size: see img_pool_layer's document.
+    :param pool_size: see img_pool_layer for details.
     :type pool_size: int
-    :param pool_type: see img_pool_layer's document.
+    :param pool_type: see img_pool_layer for details.
     :type pool_type: BasePoolingType
-    :param act: see batch_norm_layer's document.
+    :param act: see batch_norm_layer for details.
     :type act: BaseActivation
-    :param groups: see img_conv_layer's document
+    :param groups: see img_conv_layer for details.
     :type groups: int
-    :param conv_stride: see img_conv_layer's document.
+    :param conv_stride: see img_conv_layer for details.
     :type conv_stride: int
-    :param conv_padding: see img_conv_layer's document.
+    :param conv_padding: see img_conv_layer for details.
     :type conv_padding: int
-    :param conv_bias_attr: see img_conv_layer's document.
+    :param conv_bias_attr: see img_conv_layer for details.
     :type conv_bias_attr: ParameterAttribute
-    :param num_channel: see img_conv_layer's document.
+    :param num_channel: see img_conv_layer for details.
     :type num_channel: int
-    :param conv_param_attr: see img_conv_layer's document.
+    :param conv_param_attr: see img_conv_layer for details.
     :type conv_param_attr: ParameterAttribute
-    :param shared_bias: see img_conv_layer's document.
+    :param shared_bias: see img_conv_layer for details.
     :type shared_bias: bool
-    :param conv_layer_attr: see img_conv_layer's document.
+    :param conv_layer_attr: see img_conv_layer for details.
     :type conv_layer_attr: ExtraLayerOutput
-    :param bn_param_attr: see batch_norm_layer's document.
-    :type bn_param_attr: ParameterAttribute.
-    :param bn_bias_attr: see batch_norm_layer's document.
-    :param bn_layer_attr: ParameterAttribute.
-    :param pool_stride: see img_pool_layer's document.
+    :param bn_param_attr: see batch_norm_layer for details.
+    :type bn_param_attr: ParameterAttribute
+    :param bn_bias_attr: see batch_norm_layer for details.
+    :type bn_bias_attr: ParameterAttribute
+    :param bn_layer_attr: see batch_norm_layer for details.
+    :type bn_layer_attr: ExtraLayerAttribute
+    :param pool_stride: see img_pool_layer for details.
     :type pool_stride: int
-    :param pool_padding: see img_pool_layer's document.
+    :param pool_padding: see img_pool_layer for details.
     :type pool_padding: int
-    :param pool_layer_attr: see img_pool_layer's document.
+    :param pool_layer_attr: see img_pool_layer for details.
     :type pool_layer_attr: ExtraLayerAttribute
-    :return: Layer groups output
+    :return: layer's output
     :rtype: LayerOutput
     """
     __conv__ = img_conv_layer(
@@ -348,10 +349,10 @@ def img_conv_group(input,
     :param conv_batchnorm_drop_rate: if conv_with_batchnorm[i] is true,
         conv_batchnorm_drop_rate[i] represents the drop rate of each batch norm.
     :type conv_batchnorm_drop_rate: list
-    :param input: layer's input.
+    :param input: input layer.
     :type input: LayerOutput
-    :param conv_num_filter: output channels num.
-    :type conv_num_filter: int
+    :param conv_num_filter: list of output channels num.
+    :type conv_num_filter: list|tuple
     :param pool_size: pooling filter size.
     :type pool_size: int
     :param num_channels: input channels num.
@@ -362,18 +363,18 @@ def img_conv_group(input,
     :type conv_filter_size: int
     :param conv_act: activation funciton after convolution.
     :type conv_act: BaseActivation
-    :param conv_with_batchnorm: conv_with_batchnorm[i] represents
-        if there is a batch normalization after each convolution.
+    :param conv_with_batchnorm: if conv_with_batchnorm[i] is true,
+        there is a batch normalization operation after each convolution.
     :type conv_with_batchnorm: list
     :param pool_stride: pooling stride size.
     :type pool_stride: int
     :param pool_type: pooling type.
     :type pool_type: BasePoolingType
-    :param param_attr: Convolution param attribute.
-        None means default attribute.
+    :param param_attr: param attribute of convolution layer,
+                       None means default attribute.
     :type param_attr: ParameterAttribute
-    :return: Layer's output
-    :type: LayerOutput
+    :return: layer's output
+    :rtype: LayerOutput
     """
     tmp = input
 
@@ -466,12 +467,14 @@ def vgg_16_network(input_image, num_channels, num_classes=1000):
     """
     Same model from https://gist.github.com/ksimonyan/211839e770f7b538e2d8
 
-    :param num_classes:
-    :param input_image:
+    :param num_classes: number of class.
+    :type num_classes: int
+    :param input_image: input layer.
     :type input_image: LayerOutput
-    :param num_channels:
+    :param num_channels: input channels num.
     :type num_channels: int
-    :return:
+    :return: layer's output
+    :rtype: LayerOutput
     """
 
     tmp = img_conv_group(
@@ -560,8 +563,8 @@ def simple_lstm(input,
     """
     Simple LSTM Cell.
 
-    It just combine a mixed layer with fully_matrix_projection and a lstmemory
-    layer. The simple lstm cell was implemented as follow equations.
+    It just combines a mixed layer with fully_matrix_projection and a lstmemory
+    layer. The simple lstm cell was implemented with follow equations.
 
     ..  math::
 
@@ -575,37 +578,37 @@ def simple_lstm(input,
 
         h_t & = o_t tanh(c_t)
 
-    Please refer **Generating Sequences With Recurrent Neural Networks** if you
-    want to know what lstm is. Link_ is here.
+    Please refer to **Generating Sequences With Recurrent Neural Networks** for more
+    details about lstm. Link_ is here.
 
     .. _Link: http://arxiv.org/abs/1308.0850
 
     :param name: lstm layer name.
     :type name: basestring
-    :param input: input layer name.
+    :param input: layer's input.
     :type input: LayerOutput
     :param size: lstm layer size.
     :type size: int
-    :param reverse: whether to process the input data in a reverse order
+    :param reverse: process the input in a reverse order or not.
     :type reverse: bool
-    :param mat_param_attr: mixed layer's matrix projection parameter attribute.
+    :param mat_param_attr: parameter attribute of matrix projection in mixed layer.
     :type mat_param_attr: ParameterAttribute
     :param bias_param_attr: bias parameter attribute. False means no bias, None
                             means default bias.
     :type bias_param_attr: ParameterAttribute|False
-    :param inner_param_attr: lstm cell parameter attribute.
+    :param inner_param_attr: parameter attribute of lstm cell.
     :type inner_param_attr: ParameterAttribute
-    :param act: lstm final activiation type
+    :param act: last activiation type of lstm.
     :type act: BaseActivation
-    :param gate_act: lstm gate activiation type
+    :param gate_act: gate activiation type of lstm.
     :type gate_act: BaseActivation
-    :param state_act: lstm state activiation type.
+    :param state_act: state activiation type of lstm.
     :type state_act: BaseActivation
-    :param mixed_layer_attr: mixed layer's extra attribute.
+    :param mixed_layer_attr: extra attribute of mixed layer.
     :type mixed_layer_attr: ExtraLayerAttribute
-    :param lstm_cell_attr: lstm layer's extra attribute.
+    :param lstm_cell_attr: extra attribute of lstm.
     :type lstm_cell_attr: ExtraLayerAttribute
-    :return: lstm layer name.
+    :return: layer's output.
     :rtype: LayerOutput
     """
     fc_name = 'lstm_transform_%s' % name
@@ -643,9 +646,9 @@ def lstmemory_unit(input,
                    lstm_bias_attr=None,
                    lstm_layer_attr=None):
     """
-    Define calculations that a LSTM unit performs during a single time step.
-    This function itself is not a recurrent layer, so it can not be
-    directly used to process sequence inputs. This function is always used in
+    lstmemory_unit defines the caculation process of a LSTM unit during a 
+    single time step. This function is not a recurrent layer, so it can not be
+    directly used to process sequence input. This function is always used in
     recurrent_group (see layers.py for more details) to implement attention
     mechanism.
 
@@ -676,7 +679,7 @@ def lstmemory_unit(input,
                                    state_act=TanhActivation())
 
 
-    :param input: input layer name.
+    :param input: input layer.
     :type input: LayerOutput
     :param out_memory: output of previous time step
     :type out_memory: LayerOutput | None
@@ -684,15 +687,15 @@ def lstmemory_unit(input,
     :type name: basestring
     :param size: lstmemory unit size.
     :type size: int
-    :param param_attr: Parameter config, None if use default.
+    :param param_attr: parameter attribute, None means default attribute.
     :type param_attr: ParameterAttribute
-    :param act: lstm final activiation type
+    :param act: last activiation type of lstm.
     :type act: BaseActivation
-    :param gate_act: lstm gate activiation type
+    :param gate_act: gate activiation type of lstm.
     :type gate_act: BaseActivation
-    :param state_act: lstm state activiation type.
+    :param state_act: state activiation type of lstm.
     :type state_act: BaseActivation
-    :param input_proj_bias_attr: bias attribute for input-to-hidden projection.
+    :param input_proj_bias_attr: bias attribute for input to hidden projection.
                 False means no bias, None means default bias.
     :type input_proj_bias_attr: ParameterAttribute|False|None
     :param input_proj_layer_attr: extra layer attribute for input to hidden
@@ -700,8 +703,8 @@ def lstmemory_unit(input,
     :type input_proj_layer_attr: ExtraLayerAttribute
     :param lstm_bias_attr: bias parameter attribute of lstm layer.
                 False means no bias, None means default bias.
-    :type lstm_bias_attr: ParameterAttribute|False
-    :param lstm_layer_attr: lstm layer's extra attribute.
+    :type lstm_bias_attr: ParameterAttribute|False|None
+    :param lstm_layer_attr: extra attribute of lstm layer.
     :type lstm_layer_attr: ExtraLayerAttribute
     :return: lstmemory unit name.
     :rtype: LayerOutput
@@ -758,9 +761,9 @@ def lstmemory_group(input,
     lstm_group is a recurrent_group version of Long Short Term Memory. It
     does exactly the same calculation as the lstmemory layer (see lstmemory in
     layers.py for the maths) does. A promising benefit is that LSTM memory
-    cell states, or hidden states in every time step are accessible to the
+    cell states(or hidden states) in every time step are accessible to the
     user. This is especially useful in attention model. If you do not need to
-    access the internal states of the lstm, but merely use its outputs,
+    access the internal states of the lstm and merely use its outputs,
     it is recommended to use the lstmemory, which is relatively faster than
     lstmemory_group.
 
@@ -781,28 +784,28 @@ def lstmemory_group(input,
                                     gate_act=SigmoidActivation(),
                                     state_act=TanhActivation())
 
-    :param input: input layer name.
+    :param input: input layer.
     :type input: LayerOutput
     :param size: lstmemory group size.
     :type size: int
-    :param name: name of the lstmemory group.
+    :param name: name of lstmemory group.
     :type name: basestring
-    :param out_memory: output of previous time step
+    :param out_memory: output of previous time step.
     :type out_memory: LayerOutput | None
-    :param reverse: is lstm reversed
+    :param reverse: process the input in a reverse order or not.
     :type reverse: bool
-    :param param_attr: Parameter config, None if use default.
+    :param param_attr: parameter attribute, None means default attribute.
     :type param_attr: ParameterAttribute
-    :param act: lstm final activiation type
+    :param act: last activiation type of lstm.
     :type act: BaseActivation
-    :param gate_act: lstm gate activiation type
+    :param gate_act: gate activiation type of lstm.
     :type gate_act: BaseActivation
-    :param state_act: lstm state activiation type.
+    :param state_act: state activiation type of lstm.
     :type state_act: BaseActivation
     :param lstm_bias_attr: bias parameter attribute of lstm layer.
                            False means no bias, None means default bias.
-    :type lstm_bias_attr: ParameterAttribute|False
-    :param input_proj_bias_attr: bias attribute for input-to-hidden projection.
+    :type lstm_bias_attr: ParameterAttribute|False|None
+    :param input_proj_bias_attr: bias attribute for input to hidden projection.
                 False means no bias, None means default bias.
     :type input_proj_bias_attr: ParameterAttribute|False|None
     :param input_proj_layer_attr: extra layer attribute for input to hidden
@@ -848,15 +851,15 @@ def gru_unit(input,
              gru_layer_attr=None,
              naive=False):
     """
-    Define calculations that a gated recurrent unit performs in a single time
-    step. This function itself is not a recurrent layer, so it can not be
-    directly used to process sequence inputs. This function is always used in
+    gru_unit defines the calculation process of a gated recurrent unit during a single 
+    time step. This function is not a recurrent layer, so it can not be
+    directly used to process sequence input. This function is always used in
     the recurrent_group (see layers.py for more details) to implement attention
     mechanism.
 
     Please see grumemory in layers.py for the details about the maths.
 
-    :param input: input layer name.
+    :param input: input layer.
     :type input: LayerOutput
     :param memory_boot: the initialization state of the LSTM cell.
     :type memory_boot: LayerOutput | None
@@ -864,12 +867,12 @@ def gru_unit(input,
     :type name: basestring
     :param size: hidden size of the gru.
     :type size: int
-    :param act: type of the activation
+    :param act: activation type of gru
     :type act: BaseActivation
-    :param gate_act: type of the gate activation
+    :param gate_act: gate activation type or gru
     :type gate_act: BaseActivation
-    :param gru_layer_attr: Extra parameter attribute of the gru layer.
-    :type gru_layer_attr: ParameterAttribute|False
+    :param gru_layer_attr: Extra attribute of the gru layer.
+    :type gru_layer_attr: ExtraLayerAttribute
     :return: the gru output layer.
     :rtype: LayerOutput
     """
@@ -915,7 +918,7 @@ def gru_group(input,
     does exactly the same calculation as the grumemory layer does. A promising
     benefit is that gru hidden states are accessible to the user. This is
     especially useful in attention model. If you do not need to access
-    any internal state, but merely use the outputs of a GRU, it is recommended
+    any internal state and merely use the outputs of a GRU, it is recommended
     to use the grumemory, which is relatively faster.
 
     Please see grumemory in layers.py for more detail about the maths.
@@ -924,12 +927,12 @@ def gru_group(input,
 
     ..  code-block:: python
 
-        gru = gur_group(input=[layer1],
+        gru = gru_group(input=[layer1],
                         size=256,
                         act=TanhActivation(),
                         gate_act=SigmoidActivation())
 
-    :param input: input layer name.
+    :param input: input layer.
     :type input: LayerOutput
     :param memory_boot: the initialization state of the LSTM cell.
     :type memory_boot: LayerOutput | None
@@ -937,16 +940,17 @@ def gru_group(input,
     :type name: basestring
     :param size: hidden size of the gru.
     :type size: int
-    :param reverse: whether to process the input data in a reverse order
+    :param reverse: process the input in a reverse order or not.
     :type reverse: bool
-    :param act: type of the activiation
+    :param act: activiation type of gru
     :type act: BaseActivation
-    :param gate_act: type of the gate activiation
+    :param gate_act: gate activiation type of gru
     :type gate_act: BaseActivation
-    :param gru_bias_attr: bias. False means no bias, None means default bias.
-    :type gru_bias_attr: ParameterAttribute|False
-    :param gru_layer_attr: Extra parameter attribute of the gru layer.
-    :type gru_layer_attr: ParameterAttribute|False
+    :param gru_bias_attr: bias parameter attribute of gru layer,
+                          False means no bias, None means default bias.
+    :type gru_bias_attr: ParameterAttribute|False|None
+    :param gru_layer_attr: Extra attribute of the gru layer.
+    :type gru_layer_attr: ExtraLayerAttribute
     :return: the gru group.
     :rtype: LayerOutput
     """
@@ -986,11 +990,11 @@ def simple_gru(input,
                gru_layer_attr=None,
                naive=False):
     """
-    You maybe see gru_step_layer, grumemory in layers.py, gru_unit, gru_group,
+    You may see gru_step_layer, grumemory in layers.py, gru_unit, gru_group,
     simple_gru in network.py. The reason why there are so many interfaces is
     that we have two ways to implement recurrent neural network. One way is to
     use one complete layer to implement rnn (including simple rnn, gru and lstm)
-    with multiple time steps, such as recurrent_layer, lstmemory, grumemory. But,
+    with multiple time steps, such as recurrent_layer, lstmemory, grumemory. But 
     the multiplication operation :math:`W x_t` is not computed in these layers.
     See details in their interfaces in layers.py.
     The other implementation is to use an recurrent group which can ensemble a
@@ -1018,22 +1022,23 @@ def simple_gru(input,
 
         gru = simple_gru(input=[layer1], size=256)
 
-    :param input: input layer name.
+    :param input: input layer.
     :type input: LayerOutput
     :param name: name of the gru group.
     :type name: basestring
     :param size: hidden size of the gru.
     :type size: int
-    :param reverse: whether to process the input data in a reverse order
+    :param reverse: process the input in a reverse order or not.
     :type reverse: bool
-    :param act: type of the activiation
+    :param act: activiation type of gru
     :type act: BaseActivation
-    :param gate_act: type of the gate activiation
+    :param gate_act: gate activiation type of gru
     :type gate_act: BaseActivation
-    :param gru_bias_attr: bias. False means no bias, None means default bias.
-    :type gru_bias_attr: ParameterAttribute|False
-    :param gru_layer_attr: Extra parameter attribute of the gru layer.
-    :type gru_layer_attr: ParameterAttribute|False
+    :param gru_bias_attr: bias parameter attribute of gru layer,
+                          False means no bias, None means default bias.
+    :type gru_bias_attr: ParameterAttribute|False|None
+    :param gru_layer_attr: Extra attribute of the gru layer.
+    :type gru_layer_attr: ExtraLayerAttribute
     :return: the gru group.
     :rtype: LayerOutput
     """
@@ -1071,8 +1076,8 @@ def simple_gru2(input,
                 mixed_layer_attr=None,
                 gru_cell_attr=None):
     """
-    simple_gru2 is the same with simple_gru, but using grumemory instead
-    Please see grumemory in layers.py for more detail about the maths.
+    simple_gru2 is the same with simple_gru, but using grumemory instead.
+    Please refer to grumemory in layers.py for more detail about the math.
     simple_gru2 is faster than simple_gru.
 
     The example usage is:
@@ -1081,22 +1086,23 @@ def simple_gru2(input,
 
         gru = simple_gru2(input=[layer1], size=256)
 
-    :param input: input layer name.
+    :param input: input layer.
     :type input: LayerOutput
     :param name: name of the gru group.
     :type name: basestring
     :param size: hidden size of the gru.
     :type size: int
-    :param reverse: whether to process the input data in a reverse order
+    :param reverse: process the input in a reverse order or not.
     :type reverse: bool
-    :param act: type of the activiation
+    :param act: activiation type of gru
     :type act: BaseActivation
-    :param gate_act: type of the gate activiation
+    :param gate_act: gate activiation type of gru
     :type gate_act: BaseActivation
-    :param gru_bias_attr: bias. False means no bias, None means default bias.
-    :type gru_bias_attr: ParameterAttribute|False
-    :param gru_layer_attr: Extra parameter attribute of the gru layer.
-    :type gru_layer_attr: ParameterAttribute|False
+    :param gru_bias_attr: bias parameter attribute of gru layer, 
+                          False means no bias, None means default bias.
+    :type gru_bias_attr: ParameterAttribute|False|None
+    :param gru_layer_attr: Extra attribute of the gru layer.
+    :type gru_layer_attr: ExtraLayerAttribute
     :return: the gru group.
     :rtype: LayerOutput
     """
@@ -1145,7 +1151,7 @@ def bidirectional_gru(input,
                       concat_act=None):
     """
     A bidirectional_gru is a recurrent unit that iterates over the input
-    sequence both in forward and bardward orders, and then concatenate two
+    sequence both in forward and backward orders, and then concatenate two
     outputs to form a final output. However, concatenation of two outputs
     is not the only way to form the final output, you can also, for example,
     just add them together.
@@ -1162,11 +1168,10 @@ def bidirectional_gru(input,
     :type input: LayerOutput
     :param size: gru layer size.
     :type size: int
-    :param return_seq: If set False, outputs of the last time step are
-                       concatenated and returned.
-                       If set True, the entire output sequences that are
-                       processed in forward and backward directions are
+    :param return_seq: If set False, the last time step of output are
                        concatenated and returned.
+                       If set True, the entire output sequences in forward 
+                       and backward directions are concatenated and returned.
     :type return_seq: bool
     :return: LayerOutput object.
     :rtype: LayerOutput
@@ -1230,8 +1235,8 @@ def bidirectional_lstm(input,
                        concat_act=None):
     """
     A bidirectional_lstm is a recurrent unit that iterates over the input
-    sequence both in forward and bardward orders, and then concatenate two
-    outputs form a final output. However, concatenation of two outputs
+    sequence both in forward and backward orders, and then concatenate two
+    outputs to form a final output. However, concatenation of two outputs
     is not the only way to form the final output, you can also, for example,
     just add them together.
 
@@ -1252,13 +1257,12 @@ def bidirectional_lstm(input,
     :type input: LayerOutput
     :param size: lstm layer size.
     :type size: int
-    :param return_seq: If set False, outputs of the last time step are
-                       concatenated and returned.
-                       If set True, the entire output sequences that are
-                       processed in forward and backward directions are
+    :param return_seq: If set False, the last time step of output are
                        concatenated and returned.
+                       If set True, the entire output sequences in forward 
+                       and backward directions are concatenated and returned.
     :type return_seq: bool
-    :return: LayerOutput object accroding to the return_seq.
+    :return: LayerOutput object.
     :rtype: LayerOutput
     """
     args = locals()
@@ -1303,7 +1307,7 @@ def simple_attention(encoded_sequence,
                      weight_act=None,
                      name=None):
     """
-    Calculate and then return a context vector by attention machanism.
+    Calculate and return a context vector with attention mechanism.
     Size of the context vector equals to size of the encoded_sequence.
 
     ..  math::
@@ -1336,10 +1340,10 @@ def simple_attention(encoded_sequence,
     :param name: name of the attention model.
     :type name: basestring
     :param softmax_param_attr: parameter attribute of sequence softmax
-                               that is used to produce attention weight
+                               that is used to produce attention weight.
     :type softmax_param_attr: ParameterAttribute
-    :param weight_act: activation of the attention model
-    :type weight_act: Activation
+    :param weight_act: activation of the attention model.
+    :type weight_act: BaseActivation
     :param encoded_sequence: output of the encoder
     :type encoded_sequence: LayerOutput
     :param encoded_proj: attention weight is computed by a feed forward neural
@@ -1411,7 +1415,7 @@ def inputs(layers, *args):
 
 def outputs(layers, *args):
     """
-    Declare the outputs of network. If user have not defined the inputs of
+    Declare the outputs of network. If user has not defined the inputs of
     network, this method will calculate the input order by dfs travel.
 
     :param layers: Output layers.

From 88b31422b3385f7aa1d61946660bf966f2173a43 Mon Sep 17 00:00:00 2001
From: hedaoyuan <hedaoyuan@github.com>
Date: Fri, 15 Sep 2017 17:59:02 +0800
Subject: [PATCH 18/34] Delete ExpandConvBaseLayer.cpp and
 ExpandConvBaseLayer.h

---
 paddle/gserver/layers/ExpandConvBaseLayer.cpp | 61 -------------------
 paddle/gserver/layers/ExpandConvBaseLayer.h   | 44 -------------
 2 files changed, 105 deletions(-)
 delete mode 100644 paddle/gserver/layers/ExpandConvBaseLayer.cpp
 delete mode 100644 paddle/gserver/layers/ExpandConvBaseLayer.h

diff --git a/paddle/gserver/layers/ExpandConvBaseLayer.cpp b/paddle/gserver/layers/ExpandConvBaseLayer.cpp
deleted file mode 100644
index 5cf7de41f9..0000000000
--- a/paddle/gserver/layers/ExpandConvBaseLayer.cpp
+++ /dev/null
@@ -1,61 +0,0 @@
-/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
-
-Licensed under the Apache License, Version 2.0 (the "License");
-you may not use this file except in compliance with the License.
-You may obtain a copy of the License at
-
-    http://www.apache.org/licenses/LICENSE-2.0
-
-Unless required by applicable law or agreed to in writing, software
-distributed under the License is distributed on an "AS IS" BASIS,
-WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-See the License for the specific language governing permissions and
-limitations under the License. */
-
-#include "ExpandConvBaseLayer.h"
-
-#include "paddle/utils/Logging.h"
-namespace paddle {
-
-bool ExpandConvBaseLayer::init(const LayerMap &layerMap,
-                               const ParameterMap &parameterMap) {
-  /* Initialize the basic convolutional parent class */
-  ConvBaseLayer::init(layerMap, parameterMap);
-
-  int index = 0;
-  for (auto &inputConfig : config_.inputs()) {
-    const ConvConfig &conf = inputConfig.conv_conf();
-    /* Consistent caffe mode for multiple input */
-    caffeMode_ = conf.caffe_mode();
-
-    // create a new weight
-    size_t height, width;
-    height = filterPixels_[index] * filterChannels_[index];
-    width = (!isDeconv_) ? numFilters_ : channels_[index];
-    CHECK_EQ(parameters_[index]->getSize(), width * height);
-    Weight *w = new Weight(height, width, parameters_[index]);
-    weights_.emplace_back(w);
-    index++;
-  }
-  if (biasParameter_.get()) {
-    if (sharedBiases_) {
-      CHECK_EQ((size_t)numFilters_, biasParameter_->getSize());
-      biases_ =
-          std::unique_ptr<Weight>(new Weight(numFilters_, 1, biasParameter_));
-    } else {
-      biases_ =
-          std::unique_ptr<Weight>(new Weight(getSize(), 1, biasParameter_));
-    }
-  }
-  getOutputSize();
-
-  return true;
-}
-
-size_t ExpandConvBaseLayer::getOutputSize() {
-  CHECK_NE(inputLayers_.size(), 0UL);
-  size_t layerSize = ConvBaseLayer::calOutputSize();
-  return layerSize;
-}
-
-}  // namespace paddle
diff --git a/paddle/gserver/layers/ExpandConvBaseLayer.h b/paddle/gserver/layers/ExpandConvBaseLayer.h
deleted file mode 100644
index 74b3296119..0000000000
--- a/paddle/gserver/layers/ExpandConvBaseLayer.h
+++ /dev/null
@@ -1,44 +0,0 @@
-/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
-
-Licensed under the Apache License, Version 2.0 (the "License");
-you may not use this file except in compliance with the License.
-You may obtain a copy of the License at
-
-    http://www.apache.org/licenses/LICENSE-2.0
-
-Unless required by applicable law or agreed to in writing, software
-distributed under the License is distributed on an "AS IS" BASIS,
-WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-See the License for the specific language governing permissions and
-limitations under the License. */
-
-#pragma once
-
-#include <vector>
-#include "ConvBaseLayer.h"
-#include "paddle/math/Matrix.h"
-
-namespace paddle {
-
-/**
- * @brief A subclass of ConvBaseLayer that is a superclass of both
- * ExpandConvLayer and ExpandConvTransLayer
- */
-class ExpandConvBaseLayer : public ConvBaseLayer {
-protected:
-  /// The transpose of output, which is an auxiliary matrix.
-  MatrixPtr transOutValue_;
-
-public:
-  explicit ExpandConvBaseLayer(const LayerConfig& config)
-      : ConvBaseLayer(config) {}
-
-  ~ExpandConvBaseLayer() {}
-
-  bool init(const LayerMap& layerMap,
-            const ParameterMap& parameterMap) override;
-
-  size_t getOutputSize();
-};
-
-}  // namespace paddle

From dae249b1cbb57e1e148788df5fb6b048404a4b00 Mon Sep 17 00:00:00 2001
From: Liu Yiqun <liuyiqun01@baidu.com>
Date: Fri, 15 Sep 2017 10:01:11 +0000
Subject: [PATCH 19/34] Delete USE_OP statements and add more ENFORCE
 statements to check the inputs and outputs in FCOp.

---
 paddle/operators/fc_op.cc       | 23 +++++++++++++++++------
 paddle/operators/identity_op.cc |  4 ++--
 2 files changed, 19 insertions(+), 8 deletions(-)

diff --git a/paddle/operators/fc_op.cc b/paddle/operators/fc_op.cc
index 5549a836c9..e5d0f3c372 100644
--- a/paddle/operators/fc_op.cc
+++ b/paddle/operators/fc_op.cc
@@ -24,6 +24,15 @@ class FCOp : public NetOp {
        const framework::VariableNameMap &outputs,
        const framework::AttributeMap &attrs)
       : NetOp(type, inputs, outputs, attrs) {
+    PADDLE_ENFORCE(!Inputs("X").empty(),
+                   "Inputs(X) of FCOp should not be null.");
+    PADDLE_ENFORCE(!Inputs("W").empty(),
+                   "Inputs(W) of FCOp should not be null.");
+    PADDLE_ENFORCE(!Outputs("MulOut").empty(),
+                   "Outputs(MulOut) of FCOp should not be null.");
+    PADDLE_ENFORCE_NE(Output("Out"), framework::kEmptyVarName,
+                      "Output(Out) of FCOp should not be null.");
+
     auto x = Inputs("X");
     auto w = Inputs("W");
     auto mul_out = Outputs("MulOut");
@@ -68,6 +77,10 @@ class FCOp : public NetOp {
     // sum_out = X[0] * W[0] + ... + X[n-1] * W[n-1]
     auto sum_out = mul_out[0];
     if (n > 1) {
+      PADDLE_ENFORCE_NE(Output("SumOut"), framework::kEmptyVarName,
+                        "Output(SumOut) of FCOp should not be null when the "
+                        "size of Inputs(X) > 1.");
+
       sum_out = Output("SumOut");
       AppendOp(framework::OpRegistry::CreateOp("sum", {{"X", {mul_out}}},
                                                {{"Out", {sum_out}}}, {}));
@@ -81,6 +94,10 @@ class FCOp : public NetOp {
     auto b = Input("B");
     auto add_out = sum_out;
     if (b != framework::kEmptyVarName) {
+      PADDLE_ENFORCE_NE(
+          Output("AddOut"), framework::kEmptyVarName,
+          "Output(AddOut) of FCOp should not be null when Input(B) is set.");
+
       add_out = Output("AddOut");
       AppendOp(framework::OpRegistry::CreateOp(
           "rowwise_add", {{"X", {sum_out}}, {"b", {Input("B")}}},
@@ -176,11 +193,5 @@ Activation type can be set to `identity` (default), `sigmoid` or `softmax`.
 }  // namespace operators
 }  // namespace paddle
 
-USE_OP(mul);
-USE_OP(rowwise_add);
-USE_NO_KERNEL_OP(identity);
-USE_OP(sigmoid);
-USE_OP(softmax);
-
 namespace ops = paddle::operators;
 REGISTER_OP_WITHOUT_GRADIENT(fc, ops::FCOp, ops::FCOpMaker);
diff --git a/paddle/operators/identity_op.cc b/paddle/operators/identity_op.cc
index 5ab8c0fadc..2cc632205e 100644
--- a/paddle/operators/identity_op.cc
+++ b/paddle/operators/identity_op.cc
@@ -44,8 +44,8 @@ class IdentityOp : public NetOp {
       : NetOp(type, inputs, outputs, attrs) {
     PADDLE_ENFORCE_NE(Input("X"), framework::kEmptyVarName,
                       "Input(X) of IdentityOp should not be null.");
-    PADDLE_ENFORCE_NE(Output("Out"), framework::kEmptyVarName,
-                      "Output(Out) of IdentityOp should not be null.");
+    PADDLE_ENFORCE_NE(Output("Y"), framework::kEmptyVarName,
+                      "Output(Y) of IdentityOp should not be null.");
 
     AppendOp(framework::OpRegistry::CreateOp(
         "scale", {{"X", {Input("X")}}}, {{"Out", {Output("Y")}}},

From e4c340e4e7dc7ec898c1ce0f2514c9579cba460d Mon Sep 17 00:00:00 2001
From: hedaoyuan <hedaoyuan@github.com>
Date: Fri, 15 Sep 2017 18:01:33 +0800
Subject: [PATCH 20/34] Bug fix.

---
 paddle/gserver/layers/ExpandConvLayer.h | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/paddle/gserver/layers/ExpandConvLayer.h b/paddle/gserver/layers/ExpandConvLayer.h
index 698c37fb3f..a0873de192 100644
--- a/paddle/gserver/layers/ExpandConvLayer.h
+++ b/paddle/gserver/layers/ExpandConvLayer.h
@@ -15,7 +15,7 @@ limitations under the License. */
 #pragma once
 
 #include <vector>
-#include "ExpandConvBaseLayer.h"
+#include "ConvBaseLayer.h"
 #include "paddle/math/Matrix.h"
 
 namespace paddle {

From d6a0280eb99565b39cf6ba079676f465faccd3b6 Mon Sep 17 00:00:00 2001
From: dangqingqing <dangqingqing@baidu.com>
Date: Fri, 15 Sep 2017 18:21:42 +0800
Subject: [PATCH 21/34] Enhance unit testing framework for operator with
 LoDTensor.

---
 paddle/operators/sequence_avg_pool_op.cc    |  4 +++-
 paddle/operators/sequence_avg_pool_op.h     | 13 ++++++++-----
 python/paddle/v2/framework/tests/op_test.py | 15 +++++++++++----
 3 files changed, 22 insertions(+), 10 deletions(-)

diff --git a/paddle/operators/sequence_avg_pool_op.cc b/paddle/operators/sequence_avg_pool_op.cc
index c15a5833de..00b5a6c697 100644
--- a/paddle/operators/sequence_avg_pool_op.cc
+++ b/paddle/operators/sequence_avg_pool_op.cc
@@ -60,7 +60,9 @@ class SequenceAvgPoolGradOp : public framework::OperatorWithKernel {
  protected:
   void InferShape(const framework::InferShapeContext& ctx) const override {
     PADDLE_ENFORCE_NOT_NULL(ctx.InputVar(framework::GradVarName("Out")),
-                            "Gradient of Out should not be null");
+                            "Gradient of Out should not be null.");
+    PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("X"),
+                            "The input X should not be null.");
     auto og_dims =
         ctx.Input<framework::LoDTensor>(framework::GradVarName("Out"))->dims();
     auto x_dims = ctx.Input<framework::LoDTensor>("X")->dims();
diff --git a/paddle/operators/sequence_avg_pool_op.h b/paddle/operators/sequence_avg_pool_op.h
index 6e343b87e2..ebe0956344 100644
--- a/paddle/operators/sequence_avg_pool_op.h
+++ b/paddle/operators/sequence_avg_pool_op.h
@@ -21,6 +21,9 @@ namespace operators {
 
 using Tensor = framework::Tensor;
 using LoDTensor = framework::LoDTensor;
+template <typename T, int MajorType = Eigen::RowMajor,
+          typename IndexType = Eigen::DenseIndex>
+using EigenVector = framework::EigenVector<T, MajorType, IndexType>;
 template <typename T, int MajorType = Eigen::RowMajor,
           typename IndexType = Eigen::DenseIndex>
 using EigenMatrix = framework::EigenMatrix<T, MajorType, IndexType>;
@@ -43,8 +46,8 @@ class SequenceAvgPoolKernel : public framework::OpKernel {
                                  static_cast<int>(lod[0][i + 1]));
       Tensor out_t = out->Slice<T>(i, i + 1);
       int64_t h = static_cast<int64_t>(lod[0][i + 1] - lod[0][i]);
-      auto in_e = EigenMatrix<T>::From(in_t, {h, w});
-      auto out_e = EigenMatrix<T>::From(out_t, {h, w});
+      auto in_e = EigenMatrix<T>::From(in_t, framework::make_ddim({h, w}));
+      auto out_e = EigenVector<T>::Flatten(out_t);
       out_e.device(place) = in_e.mean(Eigen::array<int, 1>({{0}}));
     }
   }
@@ -54,9 +57,9 @@ template <typename Place, typename T>
 class SequenceAvgPoolGradKernel : public framework::OpKernel {
  public:
   void Compute(const framework::ExecutionContext& context) const override {
-    auto* in = context.Output<LoDTensor>("X");
-    auto* in_g = context.Output<LoDTensor>(framework::GradVarName("X"));
+    auto* in = context.Input<LoDTensor>("X");
     auto* out_g = context.Input<LoDTensor>(framework::GradVarName("Out"));
+    auto* in_g = context.Output<LoDTensor>(framework::GradVarName("X"));
 
     auto dims = in->dims();
     auto lod = in->lod();
@@ -71,7 +74,7 @@ class SequenceAvgPoolGradKernel : public framework::OpKernel {
       int64_t h = static_cast<int64_t>(lod[0][i + 1] - lod[0][i]);
       auto in_g_e = EigenMatrix<T>::From(in_g_t, {h, w});
       auto out_g_e = EigenMatrix<T>::From(out_g_t, {1, w});
-      Eigen::DSizes<int, 2> bcast(h, w);
+      Eigen::DSizes<int, 2> bcast(h, 1);
       in_g_e.device(place) = (out_g_e / static_cast<T>(h)).broadcast(bcast);
     }
   }
diff --git a/python/paddle/v2/framework/tests/op_test.py b/python/paddle/v2/framework/tests/op_test.py
index 9936fd76ba..a0533efacd 100644
--- a/python/paddle/v2/framework/tests/op_test.py
+++ b/python/paddle/v2/framework/tests/op_test.py
@@ -47,17 +47,24 @@ def set_input(scope, op, inputs, place):
         if in_name in inputs:
             if in_dup:
                 sub_in = inputs[in_name]
-                for sub_in_name, sub_in_array in sub_in:
+                for sub_in_name, sub_in_val in sub_in:
                     var = scope.find_var(sub_in_name)
                     tensor = var.get_tensor()
+                    sub_in_array = sub_in_val[0] \
+                        if isinstance(sub_in_val, tuple) else sub_in_val
                     tensor.set_dims(sub_in_array.shape)
                     tensor.set(sub_in_array, place)
+                    if isinstance(sub_in_val, tuple):
+                        tensor.set_lod(sub_in_val[1])
             else:
                 var = scope.find_var(in_name)
                 tensor = var.get_tensor()
-                arr = inputs[in_name]
-                tensor.set_dims(arr.shape)
-                tensor.set(arr, place)
+                in_val = inputs[in_name]
+                in_array = in_val[0] if isinstance(in_val, tuple) else in_val
+                tensor.set_dims(in_array.shape)
+                tensor.set(in_array, place)
+                if isinstance(in_val, tuple):
+                    tensor.set_lod(in_val[1])
 
 
 def set_output_grad(scope, op, outputs, place):

From 05680d0d5a88adaae83d144558eeee5af1ce0c02 Mon Sep 17 00:00:00 2001
From: dangqingqing <dangqingqing@baidu.com>
Date: Fri, 15 Sep 2017 18:31:40 +0800
Subject: [PATCH 22/34] Add unit testing for sequence average pool operator.

---
 .../v2/framework/tests/test_seq_pool.py       | 51 +++++++++++++++++++
 1 file changed, 51 insertions(+)
 create mode 100644 python/paddle/v2/framework/tests/test_seq_pool.py

diff --git a/python/paddle/v2/framework/tests/test_seq_pool.py b/python/paddle/v2/framework/tests/test_seq_pool.py
new file mode 100644
index 0000000000..cf864936af
--- /dev/null
+++ b/python/paddle/v2/framework/tests/test_seq_pool.py
@@ -0,0 +1,51 @@
+import unittest
+import numpy as np
+from op_test import OpTest
+
+
+class TestSeqAvgPool1D(OpTest):
+    def setUp(self):
+        self.op_type = 'sequence_avg_pool'
+        # one level, batch size is 4
+        x = np.random.uniform(0.1, 1, [11, 23]).astype('float32')
+        lod = [[0, 4, 5, 8, 11]]
+
+        out = np.zeros((4, 23)).astype('float32')
+        for i in range(4):
+            sub_x = x[lod[0][i]:lod[0][i + 1], :]
+            out[i] = sub_x.mean(axis=0)
+
+        self.inputs = {'X': (x, lod)}
+        self.outputs = {'Out': out}
+
+    def test_check_output(self):
+        self.check_output()
+
+    def test_check_grad(self):
+        self.check_grad(["X"], "Out")
+
+
+class TestSeqAvgPool2D(OpTest):
+    def setUp(self):
+        self.op_type = 'sequence_avg_pool'
+        # one level, batch size is 4
+        x = np.random.uniform(0.1, 1, [13, 3, 17]).astype('float32')
+        lod = [[0, 4, 5, 8, 13]]
+
+        out = np.zeros((4, 3, 17)).astype('float32')
+        for i in range(4):
+            sub_x = np.reshape(x[lod[0][i]:lod[0][i + 1], :], (-1, 3 * 17))
+            out[i] = np.reshape(sub_x.mean(axis=0), (3, 17))
+
+        self.inputs = {'X': (x, lod)}
+        self.outputs = {'Out': out}
+
+    def test_check_output(self):
+        self.check_output()
+
+    def test_check_grad(self):
+        self.check_grad(["X"], "Out")
+
+
+if __name__ == '__main__':
+    unittest.main()

From b1bca06667c1e7a05c6db5ad13799c5380225091 Mon Sep 17 00:00:00 2001
From: hedaoyuan <hedaoyuan@github.com>
Date: Fri, 15 Sep 2017 19:10:24 +0800
Subject: [PATCH 23/34] Refine the ExpandConvLayer.

---
 paddle/gserver/layers/ExpandConvLayer.cpp | 22 ++++++----------------
 1 file changed, 6 insertions(+), 16 deletions(-)

diff --git a/paddle/gserver/layers/ExpandConvLayer.cpp b/paddle/gserver/layers/ExpandConvLayer.cpp
index abc36aaef7..48dfcb49a4 100644
--- a/paddle/gserver/layers/ExpandConvLayer.cpp
+++ b/paddle/gserver/layers/ExpandConvLayer.cpp
@@ -53,14 +53,15 @@ bool ExpandConvLayer::init(const LayerMap &layerMap,
     weights_.emplace_back(w);
     index++;
   }
+
   if (biasParameter_.get()) {
     if (sharedBiases_) {
       CHECK_EQ((size_t)numFilters_, biasParameter_->getSize());
-      biases_ =
-          std::unique_ptr<Weight>(new Weight(numFilters_, 1, biasParameter_));
+      biases_ = std::unique_ptr<Weight>(
+          new Weight(1, numFilters_, biasParameter_, 0));
     } else {
       biases_ =
-          std::unique_ptr<Weight>(new Weight(getSize(), 1, biasParameter_));
+          std::unique_ptr<Weight>(new Weight(1, getSize(), biasParameter_, 0));
     }
   }
 
@@ -189,12 +190,7 @@ void ExpandConvLayer::forward(PassType passType) {
 
   /* add the bias-vector */
   if (biases_.get()) {
-    MatrixPtr bias = Matrix::create(biases_->getW()->getData(),
-                                    1,
-                                    biases_->getW()->getElementCnt(),
-                                    false,
-                                    useGpu_);
-    output_.value->addBias(*bias, 1.0, sharedBiases_);
+    output_.value->addBias(*biases_->getW(), 1.0, sharedBiases_);
   }
 
   /* activation */
@@ -206,13 +202,7 @@ void ExpandConvLayer::backward(const UpdateCallback &callback) {
 
   MatrixPtr outGrad = getOutputGrad();
   if (biases_ && biases_->getWGrad()) {
-    // bpropBiases(outGrad);
-    MatrixPtr bias = Matrix::create(biases_->getWGrad()->getData(),
-                                    1,
-                                    biases_->getWGrad()->getElementCnt(),
-                                    false,
-                                    useGpu_);
-    bias->collectBias(*getOutputGrad(), 1, sharedBiases_);
+    biases_->getWGrad()->collectBias(*getOutputGrad(), 1, sharedBiases_);
     /* Increasing the number of gradient */
     biases_->getParameterPtr()->incUpdate(callback);
   }

From 7fe03f7b3a72778daceb7c48ab712b37a9a5b0a7 Mon Sep 17 00:00:00 2001
From: hedaoyuan <hedaoyuan@github.com>
Date: Fri, 15 Sep 2017 23:29:33 +0800
Subject: [PATCH 24/34] Bug fix for add<float32x4_t>.

---
 paddle/cuda/include/hl_tensor_ops.h | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/paddle/cuda/include/hl_tensor_ops.h b/paddle/cuda/include/hl_tensor_ops.h
index 93d38b7d22..b2bf334dab 100644
--- a/paddle/cuda/include/hl_tensor_ops.h
+++ b/paddle/cuda/include/hl_tensor_ops.h
@@ -461,7 +461,7 @@ class add<float32x4_t> {
 public:
   INLINE float32x4_t operator()(const float32x4_t a,
                                 const float32x4_t b) const {
-    return vmulq_f32(a, b);
+    return vaddq_f32(a, b);
   }
 };
 

From 17e1aa8d436bcd57fd707e518bc05648f64982f0 Mon Sep 17 00:00:00 2001
From: Yu Yang <yuyang18@baidu.com>
Date: Fri, 15 Sep 2017 18:54:00 -0700
Subject: [PATCH 25/34] Do not invoke GPU method when use_gpu=false

---
 .../gserver/gradientmachines/RecurrentGradientMachine.cpp  | 7 ++++++-
 1 file changed, 6 insertions(+), 1 deletion(-)

diff --git a/paddle/gserver/gradientmachines/RecurrentGradientMachine.cpp b/paddle/gserver/gradientmachines/RecurrentGradientMachine.cpp
index 9f29b97466..b71431b907 100644
--- a/paddle/gserver/gradientmachines/RecurrentGradientMachine.cpp
+++ b/paddle/gserver/gradientmachines/RecurrentGradientMachine.cpp
@@ -18,6 +18,7 @@ limitations under the License. */
 #include <cmath>
 #include <functional>
 #include <limits>
+#include <memory>
 #include "NeuralNetwork.h"
 #include "paddle/gserver/layers/AgentLayer.h"
 #include "paddle/utils/Flags.h"
@@ -429,7 +430,11 @@ void RecurrentGradientMachine::reorganizeInput(PassType passType) {
   }
 
   {
-    AsyncGpuBlock asyncGpuBlock;
+    std::unique_ptr<AsyncGpuBlock> asyncBlock;
+
+    if (useGpu_) {
+      asyncBlock.reset(new AsyncGpuBlock());
+    }
 
     // inFrameLine select rows in real layer one time
     for (size_t i = 0; i < inFrameLines_.size(); i++) {

From fc3b129b08795aecc2697ea4304855b44a5bb207 Mon Sep 17 00:00:00 2001
From: peterzhang2029 <zhangchao41@baidu.com>
Date: Sun, 17 Sep 2017 11:14:15 +0800
Subject: [PATCH 26/34] delete the unused comments

---
 python/paddle/trainer_config_helpers/networks.py | 1 +
 1 file changed, 1 insertion(+)

diff --git a/python/paddle/trainer_config_helpers/networks.py b/python/paddle/trainer_config_helpers/networks.py
index 6a9bfbd5bc..93e8ac173e 100644
--- a/python/paddle/trainer_config_helpers/networks.py
+++ b/python/paddle/trainer_config_helpers/networks.py
@@ -12,6 +12,7 @@
 # See the License for the specific language governing permissions and
 # limitations under the License.
 
+
 from activations import LinearActivation, ReluActivation, SoftmaxActivation, \
     IdentityActivation, TanhActivation, SequenceSoftmaxActivation
 from attrs import ExtraAttr

From b012013a5887d7d7b387e1369de94460039292a3 Mon Sep 17 00:00:00 2001
From: Yu Yang <yuyang18@baidu.com>
Date: Sat, 16 Sep 2017 20:24:35 -0700
Subject: [PATCH 27/34] Revert "Do not invoke GPU method when use_gpu=false"

---
 .../gserver/gradientmachines/RecurrentGradientMachine.cpp  | 7 +------
 1 file changed, 1 insertion(+), 6 deletions(-)

diff --git a/paddle/gserver/gradientmachines/RecurrentGradientMachine.cpp b/paddle/gserver/gradientmachines/RecurrentGradientMachine.cpp
index b71431b907..9f29b97466 100644
--- a/paddle/gserver/gradientmachines/RecurrentGradientMachine.cpp
+++ b/paddle/gserver/gradientmachines/RecurrentGradientMachine.cpp
@@ -18,7 +18,6 @@ limitations under the License. */
 #include <cmath>
 #include <functional>
 #include <limits>
-#include <memory>
 #include "NeuralNetwork.h"
 #include "paddle/gserver/layers/AgentLayer.h"
 #include "paddle/utils/Flags.h"
@@ -430,11 +429,7 @@ void RecurrentGradientMachine::reorganizeInput(PassType passType) {
   }
 
   {
-    std::unique_ptr<AsyncGpuBlock> asyncBlock;
-
-    if (useGpu_) {
-      asyncBlock.reset(new AsyncGpuBlock());
-    }
+    AsyncGpuBlock asyncGpuBlock;
 
     // inFrameLine select rows in real layer one time
     for (size_t i = 0; i < inFrameLines_.size(); i++) {

From 59c48f9831e4a189ea042cd1ebb9be39e98b6d03 Mon Sep 17 00:00:00 2001
From: Yan Chunwei <yanchunwei@outlook.com>
Date: Sun, 17 Sep 2017 02:09:54 -0400
Subject: [PATCH 28/34] block design (#3708)

* add block
---
 doc/design/block.md | 338 ++++++++++++++++++++++++++++++++++++++++++++
 1 file changed, 338 insertions(+)
 create mode 100644 doc/design/block.md

diff --git a/doc/design/block.md b/doc/design/block.md
new file mode 100644
index 0000000000..be88001220
--- /dev/null
+++ b/doc/design/block.md
@@ -0,0 +1,338 @@
+# Design Doc: Block and Scope
+
+## The Representation of Computation
+
+Both deep learning systems and programming languages help users describe computation procedures.  These systems use various representations of computation:
+
+- Caffe, Torch, and Paddle: sequences of layers.
+- TensorFlow, Caffe2, Mxnet: graphs of operators.
+- PaddlePaddle: nested blocks, like C++ and Java programs.
+
+## Block in Programming Languages and Deep Learning
+
+In programming languages, a block is a pair of curly braces that includes local variables definitions and a sequence of instructions, or operators.
+
+Blocks work with control flow structures like `if`, `else`, and `for`, which have equivalents in deep learning:
+
+| programming languages | PaddlePaddle          |
+|-----------------------|-----------------------|
+| for, while loop       | RNN, WhileOp          |
+| if, if-else, switch   | IfElseOp, SwitchOp    |
+| sequential execution  | a sequence of layers  |
+
+A key difference is that a C++ program describes a one pass computation, whereas a deep learning program describes both the forward and backward passes.
+
+## Stack Frames and the Scope Hierarchy
+
+The existence of the backward makes the execution of a block of traditional programs and PaddlePaddle different to each other:
+
+| programming languages | PaddlePaddle                  |
+|-----------------------|-------------------------------|
+| stack                 | scope hierarchy               |
+| stack frame           | scope                         |
+| push at entering block| push at entering block        |
+| pop at leaving block  | destroy at minibatch completes|
+
+1. In traditional programs:
+
+   - When the execution enters the left curly brace of a block, the runtime pushes a frame into the stack, where it realizes local variables.
+   - After the execution leaves the right curly brace, the runtime pops the frame.
+   - The maximum number of frames in the stack is the maximum depth of nested blocks.
+
+1. In PaddlePaddle
+
+   - When the execution enters a block, PaddlePaddle adds a new scope, where it realizes variables.
+   - PaddlePaddle doesn't pop a scope after the execution of the block because variables therein are to be used by the backward pass.  So it has a stack forest known as a *scope hierarchy*.
+   - The height of the highest tree is the maximum depth of nested blocks.
+   - After the process of a minibatch, PaddlePaddle destroys the scope hierarchy.
+
+## Use Blocks in C++ and PaddlePaddle Programs
+
+Let us consolidate the discussion by presenting some examples.
+
+### Blocks with `if-else` and `IfElseOp`
+
+The following C++ programs shows how blocks are used with the `if-else` structure:
+
+```c++
+int x = 10;
+int y = 20;
+int out;
+bool cond = false;
+if (cond) {
+  int z = x + y;
+  out = softmax(z);
+} else {
+  int z = fc(x);
+  out = z;
+}
+```
+
+An equivalent PaddlePaddle program from the design doc of the [IfElseOp operator](./if_else_op.md) is as follows:
+
+```python
+import paddle as pd
+
+x = var(10)
+y = var(20)
+cond = var(false)
+ie = pd.create_ifelseop(inputs=[x], output_num=1)
+with ie.true_block():
+    x = ie.inputs(true, 0)
+    z = operator.add(x, y)
+    ie.set_output(true, 0, operator.softmax(z))
+with ie.false_block():
+    x = ie.inputs(false, 0)
+    z = layer.fc(x)
+    ie.set_output(true, 0, operator.softmax(z))
+out = b(cond)
+```
+
+In both examples, the left branch computes `softmax(x+y)` and the right branch computes `fc(x)`.
+
+A difference is that variables in the C++ program contain scalar values, whereas those in the PaddlePaddle programs are mini-batches of instances.  The `ie.input(true, 0)` invocation returns instances in the 0-th input, `x`, that corresponds to true values in `cond` as the local variable `x`, where `ie.input(false, 0)` returns instances corresponding to false values.
+
+### Blocks with `for` and `RNNOp`
+
+The following RNN model from the [RNN design doc](./rnn.md)
+
+```python
+x = sequence([10, 20, 30])
+m = var(0)
+W = tensor()
+U = tensor()
+
+rnn = create_rnn(inputs=[input])
+with rnn.stepnet() as net:
+  x = net.set_inputs(0)
+  h = net.add_memory(init=m)
+  fc_out = pd.matmul(W, x)
+  hidden_out = pd.matmul(U, h.pre(n=1))
+  sum = pd.add_two(fc_out, hidden_out)
+  act = pd.sigmoid(sum)
+  h.update(act)                       # update memory with act
+  net.set_outputs(0, act, hidden_out) # two outputs
+
+o1, o2 = rnn()
+print o1, o2
+```
+
+has its equivalent C++ program as follows
+
+```c++
+int* x = {10, 20, 30};
+int m = 0;
+int W = some_value();
+int U = some_other_value();
+
+int mem[sizeof(x) / sizeof(x[0]) + 1];
+int o1[sizeof(x) / sizeof(x[0]) + 1];
+int o2[sizeof(x) / sizeof(x[0]) + 1];
+for (int i = 1; i <= sizeof(x)/sizeof(x[0]); ++i) {
+  int x = x[i-1];
+  if (i == 1) mem[0] = m;
+  int fc_out = W * x;
+  int hidden_out = Y * mem[i-1];
+  int sum = fc_out + hidden_out;
+  int act = sigmoid(sum);
+  mem[i] = act;
+  o1[i] = act;
+  o2[i] = hidden_out;
+}
+
+print_array(o1);
+print_array(o2);
+```
+
+
+## Compilation and Execution
+
+Like TensorFlow programs, a PaddlePaddle program is written in Python.  The first part describes a neural network as a protobuf message, and the rest part executes the message for training or inference.
+
+The generation of this protobuf message is like what a compiler generates a binary executable file.  The execution of the message that the OS executes the binary file.
+
+## The "Binary Executable File Format"
+
+The definition of the protobuf message is as follows:
+
+```protobuf
+message BlockDesc {
+  repeated VarDesc vars = 1;
+  repeated OpDesc ops = 2;
+}
+```
+
+The step net in above RNN example would look like
+
+```
+BlockDesc {
+  vars = {
+    VarDesc {...} // x
+    VarDesc {...} // h
+    VarDesc {...} // fc_out
+    VarDesc {...} // hidden_out
+    VarDesc {...} // sum
+    VarDesc {...} // act
+  }
+  ops = {
+    OpDesc {...} // matmul
+    OpDesc {...} // add_two
+    OpDesc {...} // sigmoid
+  }
+};
+```
+
+Also, the RNN operator in above example is serialized into a protobuf message of type `OpDesc` and would look like:
+
+```
+OpDesc {
+  inputs = {0} // the index of x
+  outputs = {5, 3} // indices of act and hidden_out
+  attrs {
+    "memories" : {1} // the index of h
+    "step_net" : <above step net>
+  }
+};
+```
+
+This `OpDesc` value is in the `ops` field of the `BlockDesc` value representing the global block.
+
+
+## The Compilation of Blocks
+
+During the generation of the Protobuf message, the Block should store VarDesc (the Protobuf message which describes Variable) and OpDesc (the Protobuf message which describes Operator).
+
+VarDesc in a block should have its name scope to avoid local variables affect parent block's name scope.
+Child block's name scopes should inherit the parent's so that OpDesc in child block can reference a VarDesc that stored in parent block. For example
+
+```python
+a = pd.Varaible(shape=[20, 20])
+b = pd.fc(a, params=["fc.w", "fc.b"])
+
+rnn = pd.create_rnn()
+with rnn.stepnet() as net:
+    x = net.set_inputs(a)
+    # reuse fc's parameter
+    fc_without_b = pd.get_variable("fc.w")
+    net.set_outputs(fc_without_b)
+
+out = rnn()
+```
+the method `pd.get_variable` can help retrieve a Variable by a name, a Variable may store in a parent block, but might be retrieved in a child block, so block should have a variable scope that supports inheritance.
+
+In compiler design, the symbol table is a data structure created and maintained by compilers to store information about the occurrence of various entities such as variable names, function names, classes, etc.
+
+To store the definition of variables and operators, we define a C++ class `SymbolTable`, like the one used in compilers.
+
+`SymbolTable` can do the following stuff:
+
+- store the definitions (some names and attributes) of variables and operators,
+- to verify if a variable was declared,
+- to make it possible to implement type checking (offer Protobuf message pointers to `InferShape` handlers).
+
+
+```c++
+// Information in SymbolTable is enough to trace the dependency graph. So maybe
+// the Eval() interface takes a SymbolTable is enough.
+class SymbolTable {
+ public:
+  SymbolTable(SymbolTable* parent) : parent_(parent) {}
+
+  OpDesc* NewOp(const string& name="");
+
+  // TODO determine whether name is generated by python or C++
+  // currently assume that a unique name will be generated by C++ if the
+  // argument name left default.
+  VarDesc* NewVar(const string& name="");
+
+  // find a VarDesc by name, if recursive true, find parent's SymbolTable
+  // recursively.
+  // this interface is introduced to support InferShape, find protobuf messages
+  // of variables and operators, pass pointers into InferShape.
+  // operator
+  //
+  // NOTE maybe some C++ classes such as VarDescBuilder and OpDescBuilder should
+  // be proposed and embedded into pybind to enable python operate on C++ pointers.
+  VarDesc* FindVar(const string& name, bool recursive=true);
+
+  OpDesc* FindOp(const string& name);
+
+  BlockDesc Compile() const;
+
+ private:
+  SymbolTable* parent_;
+
+  map<string, OpDesc> ops_;
+  map<string, VarDesc> vars_;
+};
+```
+
+After all the description of variables and operators is added into SymbolTable,
+the block has enough information to run.
+
+The `Block` class takes a `BlockDesc` as input, and provide `Run` and `InferShape` functions.
+
+
+```c++
+namespace {
+
+class Block : OperatorBase {
+public:
+  Block(const BlockDesc& desc) desc_(desc) {}
+
+  void InferShape(const framework::Scope& scope) const override {
+    if (!symbols_ready_) {
+      CreateVariables(scope);
+      CreateOperators();
+    }
+    // should run InferShape first.
+    for (auto& op : runtime_table_.ops()) {
+      op->InferShape(scope);
+    }
+  }
+
+  void Run(const framework::Scope& scope,
+           const platform::DeviceContext& dev_ctx) const override {
+    PADDLE_ENFORCE(symbols_ready_, "operators and variables should be created first.");
+    for (auto& op : runtime_table_.ops()) {
+      op->Run(scope, dev_ctx);
+    }
+  }
+
+  void CreateVariables(const framework::Scope& scope);
+  void CreateOperators();
+
+  // some other necessary interfaces of NetOp are list below
+  // ...
+
+private:
+  BlockDesc desc_;
+  bool symbols_ready_{false};
+};
+```
+
+## The Execution of Blocks
+
+Block inherits from OperatorBase, which has a Run method.
+Block's Run method will run its operators sequentially.
+
+There is another important interface called `Eval`, which take some arguments called targets, and generate a minimal graph which takes targets as the end points and creates a new Block,
+after `Run`, `Eval` will get the latest value and return the targets.
+
+The definition of Eval is as follows:
+
+```c++
+// clean a block description by targets using the corresponding dependency graph.
+// return a new BlockDesc with minimal number of operators.
+// NOTE not return a Block but the block's description so that this can be distributed
+// to a cluster.
+BlockDesc Prune(const BlockDesc& desc, vector<string> targets);
+
+void Block::Eval(const vector<string>& targets,
+                 const framework::Scope& scope,
+                 const platform::DeviceContext& dev_ctx) {
+  BlockDesc min_desc = Prune(desc_, targets);
+  Block min_block(min_desc);
+  min_block.Run(scope, dev_ctx);
+}
+```

From 8580dce308118c5903224c9ac40213c9105ad179 Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?=E6=AD=A6=E6=AF=85?= <typhoonzero1986@gmail.com>
Date: Mon, 18 Sep 2017 15:49:34 +0800
Subject: [PATCH 29/34] Refine accuracy_op CUDA kernel (#4097)

* refind accuracy_op

* follow comments

* follow comments
---
 paddle/operators/accuracy_op.cu               | 38 ++++++++++++-------
 paddle/platform/cuda_helper.h                 |  5 +++
 .../v2/framework/tests/test_accuracy_op.py    |  9 +++--
 3 files changed, 35 insertions(+), 17 deletions(-)

diff --git a/paddle/operators/accuracy_op.cu b/paddle/operators/accuracy_op.cu
index 4e6d1ef965..0a6a0fd15c 100644
--- a/paddle/operators/accuracy_op.cu
+++ b/paddle/operators/accuracy_op.cu
@@ -12,26 +12,38 @@ 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 <thrust/execution_policy.h>
+#include <thrust/reduce.h>
 #include "paddle/operators/accuracy_op.h"
+#include "paddle/platform/cuda_helper.h"
 
 namespace paddle {
 namespace operators {
+using platform::PADDLE_CUDA_NUM_THREADS;
 
-__global__ void AccuracySingleKernel(const int N, const int D, const int top_k,
-                                     const int* Xdata, const int* labelData,
-                                     float* accuracy) {
-  int correct = 0;
-  for (int row = 0; row < N; row++) {
-    const int label = labelData[row];
-    for (int col = 0; col < D; col++) {
-      const int pred = Xdata[row * D + col];
-      if (pred == label) {
-        ++correct;
+template <int BlockSize>
+__global__ void AccuracyCudaKernel(const int N, const int D, const int* Xdata,
+                                   const int* labeldata, float* accuracy) {
+  int count = 0;
+  __shared__ int total[BlockSize];
+
+  // support only 1 block
+  for (int i = threadIdx.x; i < (N); i += BlockSize) {
+    for (int j = 0; j < D; ++j) {
+      if (Xdata[i * D + j] == labeldata[i]) {
+        ++count;
         break;
       }
     }
   }
-  *accuracy = static_cast<float>(correct) / static_cast<float>(N);
+  total[threadIdx.x] = count;
+  __syncthreads();
+
+  // reduce the count with init value 0, and output accuracy.
+  int result = thrust::reduce(thrust::device, total, total + BlockSize, 0);
+  if (threadIdx.x == 0) {
+    *accuracy = static_cast<float>(result) / static_cast<float>(N);
+  }
 }
 
 template <typename T>
@@ -57,8 +69,8 @@ class AccuracyOpCUDAKernel : public framework::OpKernel {
       return;
     }
 
-    AccuracySingleKernel<<<1, 1>>>(num_samples, infer_width, 1, inference_data,
-                                   label_data, accuracy_data);
+    AccuracyCudaKernel<PADDLE_CUDA_NUM_THREADS><<<1, PADDLE_CUDA_NUM_THREADS>>>(
+        num_samples, infer_width, inference_data, label_data, accuracy_data);
   }
 };
 
diff --git a/paddle/platform/cuda_helper.h b/paddle/platform/cuda_helper.h
index 6feec0d7f8..a7d99cde10 100644
--- a/paddle/platform/cuda_helper.h
+++ b/paddle/platform/cuda_helper.h
@@ -24,6 +24,11 @@ namespace platform {
 #define USE_CUDA_ATOMIC(op, T) \
   CUDA_ATOMIC_WRAPPER(op, T) { return atomic##op(address, val); }
 
+// Default thread count per block(or block size).
+// TODO(typhoonzero): need to benchmark against setting this value
+//                    to 1024.
+constexpr int PADDLE_CUDA_NUM_THREADS = 512;
+
 // For atomicAdd.
 USE_CUDA_ATOMIC(Add, float);
 
diff --git a/python/paddle/v2/framework/tests/test_accuracy_op.py b/python/paddle/v2/framework/tests/test_accuracy_op.py
index 43d60eb90d..b6f3a35d6f 100644
--- a/python/paddle/v2/framework/tests/test_accuracy_op.py
+++ b/python/paddle/v2/framework/tests/test_accuracy_op.py
@@ -6,16 +6,17 @@ from op_test import OpTest
 class TestAccuracyOp(OpTest):
     def setUp(self):
         self.op_type = "accuracy"
-        infer = np.random.randint(0, 2, (32, 1)).astype("int")
-        label = np.random.randint(0, 2, (32, )).astype("int")
+        n = 8192
+        infer = np.random.randint(0, 2, (n, 1)).astype("int")
+        label = np.random.randint(0, 2, (n, )).astype("int")
         self.inputs = {'Inference': infer, "Label": label}
         num_correct = 0
-        for rowid in xrange(32):
+        for rowid in xrange(n):
             for ele in infer[rowid]:
                 if ele == label[rowid]:
                     num_correct += 1
                     break
-        self.outputs = {'Accuracy': [num_correct / 32.0]}
+        self.outputs = {'Accuracy': [num_correct / float(n)]}
 
     def test_check_output(self):
         self.check_output()

From 8067a42befef3eec593546c8fba11d3b802f2a62 Mon Sep 17 00:00:00 2001
From: Luo Tao <luotao02@baidu.com>
Date: Mon, 18 Sep 2017 15:31:49 +0800
Subject: [PATCH 30/34] refine avg-pooling, which is exclusive. refine related
 code.

---
 paddle/cuda/src/hl_cuda_cnn.cu |  31 ++--
 paddle/math/Matrix.cpp         | 283 +++++++++++++++------------------
 2 files changed, 142 insertions(+), 172 deletions(-)

diff --git a/paddle/cuda/src/hl_cuda_cnn.cu b/paddle/cuda/src/hl_cuda_cnn.cu
index 9ba3d14261..596dd6d1bf 100644
--- a/paddle/cuda/src/hl_cuda_cnn.cu
+++ b/paddle/cuda/src/hl_cuda_cnn.cu
@@ -211,13 +211,11 @@ __global__ void KeAvgPoolForward(const int nthreads,
 
     int hstart = ph * strideH - padH;
     int wstart = pw * strideW - padW;
-    int hend = min(hstart + sizeY, height + padH);
-    int wend = min(wstart + sizeX, width + padW);
-    int pool_size = (hend - hstart) * (wend - wstart);
+    int hend = min(hstart + sizeY, height);
+    int wend = min(wstart + sizeX, width);
     hstart = max(hstart, 0);
     wstart = max(wstart, 0);
-    hend = min(hend, height);
-    wend = min(wend, width);
+    int pool_size = (hend - hstart) * (wend - wstart);
 
     real aveval = 0;
     inputData += (frameNum * channels + c) * height * width;
@@ -299,11 +297,11 @@ __global__ void KeAvgPoolBackward(const int nthreads,
     outGrad += (frameNum * outStride + offsetC * pooledH * pooledW);
 
     for (int ph = phstart; ph < phend; ++ph) {
+      int hstart = ph * strideH - padH;
+      int hend = min(hstart + sizeY, height + padH);
       for (int pw = pwstart; pw < pwend; ++pw) {
         // figure out the pooling size
-        int hstart = ph * strideH - padH;
         int wstart = pw * strideW - padW;
-        int hend = min(hstart + sizeY, height + padH);
         int wend = min(wstart + sizeX, width + padW);
         int poolsize = (hend - hstart) * (wend - wstart);
         gradient += outGrad[ph * pooledW + pw] / poolsize;
@@ -600,16 +598,13 @@ __global__ void KeAvgPool3DForward(const int nthreads,
     int dstart = pd * strideD - padD;
     int hstart = ph * strideH - padH;
     int wstart = pw * strideW - padW;
-    int dend = min(dstart + sizeZ, depth + padD);
-    int hend = min(hstart + sizeY, height + padH);
-    int wend = min(wstart + sizeX, width + padW);
-    int pool_size = (dend - dstart) * (hend - hstart) * (wend - wstart);
+    int dend = min(dstart + sizeZ, depth);
+    int hend = min(hstart + sizeY, height);
+    int wend = min(wstart + sizeX, width);
     dstart = max(dstart, 0);
     hstart = max(hstart, 0);
     wstart = max(wstart, 0);
-    dend = min(dend, depth);
-    hend = min(hend, height);
-    wend = min(wend, width);
+    int pool_size = (dend - dstart) * (hend - hstart) * (wend - wstart);
 
     real aveval = 0;
     inputData += (frameNum * channels + c) * depth * height * width;
@@ -712,14 +707,14 @@ __global__ void KeAvgPool3DBackward(const int nthreads,
     outGrad += (frameNum * channels + offsetC) * pooledD * pooledH * pooledW;
 
     for (int pd = pdstart; pd < pdend; ++pd) {
+      int dstart = pd * strideD - padD;
+      int dend = min(dstart + sizeZ, depth + padD);
       for (int ph = phstart; ph < phend; ++ph) {
+        int hstart = ph * strideH - padH;
+        int hend = min(hstart + sizeY, height + padH);
         for (int pw = pwstart; pw < pwend; ++pw) {
           // figure out the pooling size
-          int dstart = pd * strideD - padD;
-          int hstart = ph * strideH - padH;
           int wstart = pw * strideW - padW;
-          int dend = min(dstart + sizeZ, depth + padD);
-          int hend = min(hstart + sizeY, height + padH);
           int wend = min(wstart + sizeX, width + padW);
           int poolsize = (dend - dstart) * (hend - hstart) * (wend - wstart);
           gradient += outGrad[(pd * pooledH + ph) * pooledW + pw] / poolsize;
diff --git a/paddle/math/Matrix.cpp b/paddle/math/Matrix.cpp
index 4a2132c8d1..83c06a04c6 100644
--- a/paddle/math/Matrix.cpp
+++ b/paddle/math/Matrix.cpp
@@ -1033,17 +1033,15 @@ void GpuMatrix::maxPoolForward(Matrix& inputMat,
 
   real* inputData = inputMat.getData();
   size_t frameNum = inputMat.getHeight();
-  size_t width = imgSizeW;
-  size_t height = imgSizeH;
-  CHECK(height * width * channels == inputMat.getWidth());
+  CHECK(imgSizeH * imgSizeW * channels == inputMat.getWidth());
   CHECK(height_ == inputMat.getHeight());
   CHECK(width_ == outputH * outputW * channels);
 
   hl_maxpool_forward(frameNum,
                      inputData,
                      channels,
-                     height,
-                     width,
+                     imgSizeH,
+                     imgSizeW,
                      outputH,
                      outputW,
                      sizeX,
@@ -1080,11 +1078,9 @@ void GpuMatrix::maxPoolBackward(Matrix& inputMat,
   real* outDiff = outGrad.getData();
   size_t frameNum = inputMat.getHeight();
   size_t channels = outV.getWidth() / outputH / outputW;
-  size_t width = imgSizeW;
-  size_t height = imgSizeH;
-  CHECK(height * width * channels == inputMat.getWidth());
+  CHECK(imgSizeH * imgSizeW * channels == inputMat.getWidth());
   CHECK(height_ == inputMat.getHeight());
-  CHECK(width_ == width * height * channels);
+  CHECK(width_ == imgSizeW * imgSizeH * channels);
   CHECK(outGrad.getHeight() == outV.getHeight() &&
         outGrad.getWidth() == outV.getWidth());
 
@@ -1093,8 +1089,8 @@ void GpuMatrix::maxPoolBackward(Matrix& inputMat,
                       outData,
                       outDiff,
                       channels,
-                      height,
-                      width,
+                      imgSizeH,
+                      imgSizeW,
                       outputH,
                       outputW,
                       sizeX,
@@ -1125,17 +1121,15 @@ void GpuMatrix::avgPoolForward(Matrix& inputMat,
 
   real* inputData = inputMat.getData();
   size_t frameNum = inputMat.getHeight();
-  size_t height = imgSizeH;
-  size_t width = imgSizeW;
-  CHECK(height * width * channels == inputMat.getWidth());
+  CHECK(imgSizeH * imgSizeW * channels == inputMat.getWidth());
   CHECK(height_ == inputMat.getHeight());
   CHECK(width_ == outputH * outputW * channels);
 
   hl_avgpool_forward(frameNum,
                      inputData,
                      channels,
-                     height,
-                     width,
+                     imgSizeH,
+                     imgSizeW,
                      outputH,
                      outputW,
                      sizeX,
@@ -1166,17 +1160,15 @@ void GpuMatrix::avgPoolBackward(Matrix& outGrad,
   real* outDiff = outGrad.getData();
   size_t frameNum = outGrad.getHeight();
   size_t channels = outGrad.getWidth() / outputH / outputW;
-  size_t height = imgSizeH;
-  size_t width = imgSizeW;
-  CHECK(height * width * channels == width_);
+  CHECK(imgSizeH * imgSizeW * channels == width_);
   CHECK(height_ == outGrad.getHeight());
   CHECK(outGrad.getWidth() == outputH * outputW * channels);
 
   hl_avgpool_backward(frameNum,
                       outDiff,
                       channels,
-                      height,
-                      width,
+                      imgSizeH,
+                      imgSizeW,
                       outputH,
                       outputW,
                       sizeX,
@@ -1214,19 +1206,16 @@ void GpuMatrix::maxPool3DForward(Matrix& inputMat,
   real* inputData = inputMat.getData();
   real* maxPoolIdxData = maxPoolIdx.getData();
   size_t num = inputMat.getHeight();
-  size_t width = imgSizeW;
-  size_t height = imgSizeH;
-  size_t depth = imgSizeD;
-  CHECK(depth * height * width * channels == inputMat.getWidth());
+  CHECK(imgSizeD * imgSizeH * imgSizeW * channels == inputMat.getWidth());
   CHECK(height_ == inputMat.getHeight());
   CHECK(width_ == outputD * outputH * outputW * channels);
 
   hl_maxpool3D_forward(num,
                        inputData,
                        channels,
-                       depth,
-                       height,
-                       width,
+                       imgSizeD,
+                       imgSizeH,
+                       imgSizeW,
                        outputD,
                        outputH,
                        outputW,
@@ -1269,20 +1258,17 @@ void GpuMatrix::maxPool3DBackward(Matrix& outGrad,
   real* maxPoolIdxData = maxPoolIdx.getData();
   size_t frameNum = getHeight();
   size_t channels = outGrad.getWidth() / outputD / outputH / outputW;
-  size_t width = imgSizeW;
-  size_t height = imgSizeH;
-  size_t depth = imgSizeD;
-  CHECK(depth * height * width * channels == getWidth());
-  CHECK(width_ == depth * width * height * channels);
+  CHECK(imgSizeD * imgSizeH * imgSizeW * channels == getWidth());
+  CHECK(width_ == imgSizeD * imgSizeH * imgSizeW * channels);
   CHECK(outGrad.getHeight() == maxPoolIdx.getHeight() &&
         outGrad.getWidth() == maxPoolIdx.getWidth());
 
   hl_maxpool3D_backward(frameNum,
                         outDiff,
                         channels,
-                        depth,
-                        height,
-                        width,
+                        imgSizeD,
+                        imgSizeH,
+                        imgSizeW,
                         outputD,
                         outputH,
                         outputW,
@@ -1323,19 +1309,16 @@ void GpuMatrix::avgPool3DForward(Matrix& inputMat,
 
   real* inputData = inputMat.getData();
   size_t frameNum = inputMat.getHeight();
-  size_t height = imgSizeH;
-  size_t width = imgSizeW;
-  size_t depth = imgSizeD;
-  CHECK(depth * height * width * channels == inputMat.getWidth());
+  CHECK(imgSizeD * imgSizeH * imgSizeW * channels == inputMat.getWidth());
   CHECK(height_ == inputMat.getHeight());
   CHECK(width_ == outputD * outputH * outputW * channels);
 
   hl_avgpool3D_forward(frameNum,
                        inputData,
                        channels,
-                       depth,
-                       height,
-                       width,
+                       imgSizeD,
+                       imgSizeH,
+                       imgSizeW,
                        outputD,
                        outputH,
                        outputW,
@@ -1375,19 +1358,16 @@ void GpuMatrix::avgPool3DBackward(Matrix& outGrad,
   real* outDiff = outGrad.getData();
   size_t frameNum = outGrad.getHeight();
   size_t channels = outGrad.getWidth() / outputD / outputH / outputW;
-  size_t height = imgSizeH;
-  size_t width = imgSizeW;
-  size_t depth = imgSizeD;
-  CHECK(depth * height * width * channels == width_);
+  CHECK(imgSizeD * imgSizeH * imgSizeW * channels == width_);
   CHECK(height_ == outGrad.getHeight());
   CHECK(outGrad.getWidth() == outputD * outputH * outputW * channels);
 
   hl_avgpool3D_backward(frameNum,
                         outDiff,
                         channels,
-                        depth,
-                        height,
-                        width,
+                        imgSizeD,
+                        imgSizeH,
+                        imgSizeW,
                         outputD,
                         outputH,
                         outputW,
@@ -1999,11 +1979,11 @@ void CpuMatrix::maxPoolForward(Matrix& inputMat,
   real* inputData = inputMat.getData();
   real* outData = data_;
   size_t num = inputMat.getHeight();
-  size_t inWidth = imgSizeW;
-  size_t inHeight = imgSizeH;
-  CHECK(inHeight * inWidth == inputMat.getWidth() / channels);
+  size_t inLength = imgSizeH * imgSizeW;
+  size_t outLength = outputH * outputW;
+  CHECK(inLength == inputMat.getWidth() / channels);
   CHECK_EQ(num, this->getHeight());
-  CHECK_EQ(channels * outputH * outputW, this->getWidth());
+  CHECK_EQ(channels * outLength, this->getWidth());
   size_t outStride = getStride();
 
   /* initialize the data_ */
@@ -2020,24 +2000,24 @@ void CpuMatrix::maxPoolForward(Matrix& inputMat,
     }
     for (size_t c = 0; c < channels; ++c) {  // channel by channel
       for (size_t ph = 0; ph < outputH; ++ph) {
+        int hstart = ph * strideH - paddingH;
+        int hend = std::min(hstart + sizeY, imgSizeH);
+        hstart = std::max(hstart, 0);
         for (size_t pw = 0; pw < outputW; ++pw) {
-          int hstart = ph * strideH - paddingH;
           int wstart = pw * strideW - paddingW;
-          int hend = std::min(hstart + sizeY, inHeight);
-          int wend = std::min(wstart + sizeX, inWidth);
-          hstart = std::max(hstart, 0);
+          int wend = std::min(wstart + sizeX, imgSizeW);
           wstart = std::max(wstart, 0);
           for (int h = hstart; h < hend; ++h) {
             for (int w = wstart; w < wend; ++w) {
-              outData[ph * outputW + pw] = std::max(outData[ph * outputW + pw],
-                                                    inputData[h * inWidth + w]);
+              outData[ph * outputW + pw] = std::max(
+                  outData[ph * outputW + pw], inputData[h * imgSizeW + w]);
             }
           }
         }
       }
       // compute offset
-      inputData += inHeight * inWidth;
-      outData += outputH * outputW;
+      inputData += inLength;
+      outData += outLength;
     }
   }
 }
@@ -2058,8 +2038,10 @@ void CpuMatrix::maxPoolBackward(Matrix& image,
                                 size_t paddingH,
                                 size_t paddingW) {
   size_t num = image.getHeight();
-  size_t channels = size_t(width_ / imgSizeH / imgSizeW);
-  CHECK(image.getWidth() == imgSizeH * imgSizeW * channels);
+  size_t inLength = imgSizeH * imgSizeW;
+  size_t outLength = outputH * outputW;
+  size_t channels = size_t(width_ / inLength);
+  CHECK(image.getWidth() == inLength * channels);
   CHECK(image.getHeight() == height_ && image.getWidth() == width_);
   CHECK(outV.getHeight() == outGrad.getHeight() &&
         outV.getWidth() == outGrad.getWidth());
@@ -2080,12 +2062,12 @@ void CpuMatrix::maxPoolBackward(Matrix& image,
     }
     for (size_t c = 0; c < channels; ++c) {
       for (size_t ph = 0; ph < outputH; ++ph) {
+        int hstart = ph * strideH - paddingH;
+        int hend = std::min(hstart + sizeY, imgSizeH);
+        hstart = std::max(hstart, 0);
         for (size_t pw = 0; pw < outputW; ++pw) {
-          int hstart = ph * strideH - paddingH;
           int wstart = pw * strideW - paddingW;
-          int hend = std::min(hstart + sizeY, imgSizeH);
           int wend = std::min(wstart + sizeX, imgSizeW);
-          hstart = std::max(hstart, 0);
           wstart = std::max(wstart, 0);
           for (int h = hstart; h < hend; ++h) {
             for (int w = wstart; w < wend; ++w) {
@@ -2098,10 +2080,10 @@ void CpuMatrix::maxPoolBackward(Matrix& image,
         }
       }
       // offset
-      inData += imgSizeH * imgSizeW;
-      tgtGrad += imgSizeH * imgSizeW;
-      otData += outputH * outputW;
-      otGrad += outputH * outputW;
+      inData += inLength;
+      tgtGrad += inLength;
+      otData += outLength;
+      otGrad += outLength;
     }
   }
 }
@@ -2120,10 +2102,10 @@ void CpuMatrix::avgPoolForward(Matrix& input,
                                size_t paddingW) {
   // The main loop
   size_t num = input.getHeight();
-  size_t inHeight = imgSizeH;
-  size_t inWidth = imgSizeW;
-  CHECK(inHeight * inWidth * channels == input.getWidth());
-  CHECK(outputH * outputW * channels * num == height_ * width_);
+  size_t inLength = imgSizeH * imgSizeW;
+  size_t outLength = outputH * outputW;
+  CHECK(inLength * channels == input.getWidth());
+  CHECK(outLength * channels * num == height_ * width_);
   real* tgtData = data_;
   real* inData = input.getData();
 
@@ -2133,30 +2115,27 @@ void CpuMatrix::avgPoolForward(Matrix& input,
     }
     for (size_t c = 0; c < channels; ++c) {
       for (size_t ph = 0; ph < outputH; ++ph) {
+        int hstart = ph * strideH - paddingH;
+        int hend = std::min(hstart + sizeY, imgSizeH);
+        hstart = std::max(hstart, 0);
         for (size_t pw = 0; pw < outputW; ++pw) {
-          int hstart = ph * strideH - paddingH;
           int wstart = pw * strideW - paddingW;
-          int hend = std::min(hstart + sizeY, inHeight + paddingH);
-          int wend = std::min(wstart + sizeX, inWidth + paddingW);
-          int poolSize = (hend - hstart) * (wend - wstart);
-          hstart = std::max(hstart, 0);
+          int wend = std::min(wstart + sizeX, imgSizeW);
           wstart = std::max(wstart, 0);
-          hend = std::min(hend, static_cast<int>(inHeight));
-          wend = std::min(wend, static_cast<int>(inWidth));
-
-          CHECK(poolSize);
           tgtData[ph * outputW + pw] = 0;  // clear
           for (int h = hstart; h < hend; ++h) {
             for (int w = wstart; w < wend; ++w) {
-              tgtData[ph * outputW + pw] += inData[h * inWidth + w];
+              tgtData[ph * outputW + pw] += inData[h * imgSizeW + w];
             }
           }
+          int poolSize = (hend - hstart) * (wend - wstart);
+          CHECK(poolSize);
           tgtData[ph * outputW + pw] /= poolSize;
         }
       }
       // compute offset
-      inData += inHeight * inWidth;
-      tgtData += outputH * outputW;
+      inData += inLength;
+      tgtData += outLength;
     }
   }
 }
@@ -2176,7 +2155,9 @@ void CpuMatrix::avgPoolBackward(Matrix& input,
                                 size_t paddingW) {
   size_t num = input.getHeight();
   size_t channels = input.getWidth() / outputH / outputW;
-  CHECK(imgSizeH * imgSizeW * channels == getWidth());
+  size_t inLength = imgSizeH * imgSizeW;
+  size_t outLength = outputH * outputW;
+  CHECK(inLength * channels == getWidth());
   real* inData = input.getData();
   real* outData = getData();
 
@@ -2186,16 +2167,14 @@ void CpuMatrix::avgPoolBackward(Matrix& input,
     }
     for (size_t c = 0; c < channels; ++c) {
       for (size_t ph = 0; ph < outputH; ++ph) {
+        int hstart = ph * strideH - paddingH;
+        int hend = std::min(hstart + sizeY, imgSizeH);
+        hstart = std::max(hstart, 0);
         for (size_t pw = 0; pw < outputW; ++pw) {
-          int hstart = ph * strideH - paddingH;
           int wstart = pw * strideW - paddingW;
-          int hend = std::min(hstart + sizeY, imgSizeH + paddingH);
-          int wend = std::min(wstart + sizeX, imgSizeW + paddingW);
-          int poolSize = (hend - hstart) * (wend - wstart);
-          hstart = std::max(hstart, 0);
+          int wend = std::min(wstart + sizeX, imgSizeW);
           wstart = std::max(wstart, 0);
-          hend = std::min(hend, static_cast<int>(imgSizeH));
-          wend = std::min(wend, static_cast<int>(imgSizeW));
+          int poolSize = (hend - hstart) * (wend - wstart);
           CHECK(poolSize);
 
           for (int h = hstart; h < hend; ++h) {
@@ -2206,8 +2185,8 @@ void CpuMatrix::avgPoolBackward(Matrix& input,
         }
       }
       // offset
-      outData += imgSizeH * imgSizeW;
-      inData += outputH * outputW;
+      outData += inLength;
+      inData += outLength;
     }
   }
 }
@@ -2234,12 +2213,11 @@ void CpuMatrix::maxPool3DForward(Matrix& inputMat,
   real* outData = getData();
   real* maxPoolIdxData = maxPoolIdx.getData();
   size_t num = inputMat.getHeight();
-  size_t inWidth = imgSizeW;
-  size_t inHeight = imgSizeH;
-  size_t inDepth = imgSizeD;
-  CHECK(inHeight * inWidth * inDepth == inputMat.getWidth() / channels);
+  size_t inLength = imgSizeH * imgSizeW * imgSizeD;
+  size_t outLength = outputH * outputW * outputD;
+  CHECK(inLength == inputMat.getWidth() / channels);
   CHECK_EQ(num, this->getHeight());
-  CHECK_EQ(channels * outputH * outputW * outputD, this->getWidth());
+  CHECK_EQ(channels * outLength, this->getWidth());
   size_t outStride = getStride();
 
   /* initialize the data_ */
@@ -2258,16 +2236,16 @@ void CpuMatrix::maxPool3DForward(Matrix& inputMat,
     }
     for (size_t c = 0; c < channels; ++c) {  // channel by channel
       for (size_t pd = 0; pd < outputD; ++pd) {
+        int dstart = pd * strideD - paddingD;
+        int dend = std::min(dstart + sizeZ, imgSizeD);
+        dstart = std::max(dstart, 0);
         for (size_t ph = 0; ph < outputH; ++ph) {
+          int hstart = ph * strideH - paddingH;
+          int hend = std::min(hstart + sizeY, imgSizeH);
+          hstart = std::max(hstart, 0);
           for (size_t pw = 0; pw < outputW; ++pw) {
-            int dstart = pd * strideD - paddingD;
-            int hstart = ph * strideH - paddingH;
             int wstart = pw * strideW - paddingW;
-            int dend = std::min(dstart + sizeZ, inDepth);
-            int hend = std::min(hstart + sizeY, inHeight);
-            int wend = std::min(wstart + sizeX, inWidth);
-            dstart = std::max(dstart, 0);
-            hstart = std::max(hstart, 0);
+            int wend = std::min(wstart + sizeX, imgSizeW);
             wstart = std::max(wstart, 0);
             int maxIdx = -1;
             real maxOutData = outData[(pd * outputH + ph) * outputW + pw];
@@ -2275,9 +2253,9 @@ void CpuMatrix::maxPool3DForward(Matrix& inputMat,
               for (int h = hstart; h < hend; ++h) {
                 for (int w = wstart; w < wend; ++w) {
                   if (maxOutData <
-                      inputData[(d * inHeight + h) * inWidth + w]) {
-                    maxOutData = inputData[(d * inHeight + h) * inWidth + w];
-                    maxIdx = (d * inHeight + h) * inWidth + w;
+                      inputData[(d * imgSizeH + h) * imgSizeW + w]) {
+                    maxOutData = inputData[(d * imgSizeH + h) * imgSizeW + w];
+                    maxIdx = (d * imgSizeH + h) * imgSizeW + w;
                   }
                 }
               }
@@ -2288,9 +2266,9 @@ void CpuMatrix::maxPool3DForward(Matrix& inputMat,
         }
       }
       // compute offset
-      inputData += inDepth * inHeight * inWidth;
-      outData += outputD * outputH * outputW;
-      maxPoolIdxData += outputD * outputH * outputW;
+      inputData += inLength;
+      outData += outLength;
+      maxPoolIdxData += outLength;
     }
   }
 }
@@ -2315,7 +2293,9 @@ void CpuMatrix::maxPool3DBackward(Matrix& outGrad,
                                   real scaleTargets,
                                   real scaleOutput) {
   size_t num = getHeight();
-  size_t channels = size_t(width_ / imgSizeD / imgSizeH / imgSizeW);
+  size_t inLength = imgSizeH * imgSizeW * imgSizeD;
+  size_t outLength = outputH * outputW * outputD;
+  size_t channels = size_t(width_ / inLength);
   CHECK(maxPoolIdx.getHeight() == outGrad.getHeight() &&
         maxPoolIdx.getWidth() == outGrad.getWidth());
 
@@ -2341,9 +2321,9 @@ void CpuMatrix::maxPool3DBackward(Matrix& outGrad,
         }
       }
       // offset
-      tgtGrad += imgSizeD * imgSizeH * imgSizeW;
-      otGrad += outputD * outputH * outputW;
-      maxPoolIdxData += outputD * outputH * outputW;
+      tgtGrad += inLength;
+      otGrad += outLength;
+      maxPoolIdxData += outLength;
     }
   }
 }
@@ -2367,11 +2347,10 @@ void CpuMatrix::avgPool3DForward(Matrix& input,
                                  size_t paddingW) {
   // The main loop
   size_t num = input.getHeight();
-  size_t inDepth = imgSizeD;
-  size_t inHeight = imgSizeH;
-  size_t inWidth = imgSizeW;
-  CHECK(inDepth * inHeight * inWidth * channels == input.getWidth());
-  CHECK(outputD * outputH * outputW * channels * num == height_ * width_);
+  size_t inLength = imgSizeH * imgSizeW * imgSizeD;
+  size_t outLength = outputH * outputW * outputD;
+  CHECK(inLength * channels == input.getWidth());
+  CHECK(outLength * channels * num == height_ * width_);
   real* tgtData = getData();
   real* inData = input.getData();
 
@@ -2381,39 +2360,36 @@ void CpuMatrix::avgPool3DForward(Matrix& input,
     }
     for (size_t c = 0; c < channels; ++c) {
       for (size_t pd = 0; pd < outputD; ++pd) {
+        int dstart = pd * strideD - paddingD;
+        int dend = std::min(dstart + sizeZ, imgSizeD);
+        dstart = std::max(dstart, 0);
         for (size_t ph = 0; ph < outputH; ++ph) {
+          int hstart = ph * strideH - paddingH;
+          int hend = std::min(hstart + sizeY, imgSizeH);
+          hstart = std::max(hstart, 0);
           for (size_t pw = 0; pw < outputW; ++pw) {
-            int dstart = pd * strideD - paddingD;
-            int hstart = ph * strideH - paddingH;
             int wstart = pw * strideW - paddingW;
-            int dend = std::min(dstart + sizeZ, inDepth + paddingD);
-            int hend = std::min(hstart + sizeY, inHeight + paddingH);
-            int wend = std::min(wstart + sizeX, inWidth + paddingW);
-            int poolSize = (dend - dstart) * (hend - hstart) * (wend - wstart);
-            dstart = std::max(dstart, 0);
-            hstart = std::max(hstart, 0);
+            int wend = std::min(wstart + sizeX, imgSizeW);
             wstart = std::max(wstart, 0);
-            dend = std::min(dend, static_cast<int>(inDepth));
-            hend = std::min(hend, static_cast<int>(inHeight));
-            wend = std::min(wend, static_cast<int>(inWidth));
 
-            CHECK(poolSize);
             tgtData[(pd * outputH + ph) * outputW + pw] = 0;  // clear
             for (int d = dstart; d < dend; ++d) {
               for (int h = hstart; h < hend; ++h) {
                 for (int w = wstart; w < wend; ++w) {
                   tgtData[(pd * outputH + ph) * outputW + pw] +=
-                      inData[(d * inHeight + h) * inWidth + w];
+                      inData[(d * imgSizeH + h) * imgSizeW + w];
                 }
               }
             }
+            int poolSize = (dend - dstart) * (hend - hstart) * (wend - wstart);
+            CHECK(poolSize);
             tgtData[(pd * outputH + ph) * outputW + pw] /= poolSize;
           }
         }
       }
       // compute offset
-      inData += inDepth * inHeight * inWidth;
-      tgtData += outputD * outputH * outputW;
+      inData += inLength;
+      tgtData += outLength;
     }
   }
 }
@@ -2437,8 +2413,10 @@ void CpuMatrix::avgPool3DBackward(Matrix& input,
                                   real scaleTargets,
                                   real scaleOutput) {
   size_t num = input.getHeight();
-  size_t channels = input.getWidth() / outputD / outputH / outputW;
-  CHECK(imgSizeD * imgSizeH * imgSizeW * channels == getWidth());
+  size_t inLength = imgSizeH * imgSizeW * imgSizeD;
+  size_t outLength = outputH * outputW * outputD;
+  size_t channels = input.getWidth() / outLength;
+  CHECK(inLength * channels == getWidth());
   real* inData = input.getData();
   real* outData = getData();
 
@@ -2448,21 +2426,18 @@ void CpuMatrix::avgPool3DBackward(Matrix& input,
     }
     for (size_t c = 0; c < channels; ++c) {
       for (size_t pd = 0; pd < outputD; ++pd) {
+        int dstart = pd * strideD - paddingD;
+        int dend = std::min(dstart + sizeZ, imgSizeD);
+        dstart = std::max(dstart, 0);
         for (size_t ph = 0; ph < outputH; ++ph) {
+          int hstart = ph * strideH - paddingH;
+          int hend = std::min(hstart + sizeY, imgSizeH);
+          hstart = std::max(hstart, 0);
           for (size_t pw = 0; pw < outputW; ++pw) {
-            int dstart = pd * strideD - paddingD;
-            int hstart = ph * strideH - paddingH;
             int wstart = pw * strideW - paddingW;
-            int dend = std::min(dstart + sizeZ, imgSizeD + paddingD);
-            int hend = std::min(hstart + sizeY, imgSizeH + paddingH);
-            int wend = std::min(wstart + sizeX, imgSizeW + paddingW);
-            int poolSize = (dend - dstart) * (hend - hstart) * (wend - wstart);
-            dstart = std::max(dstart, 0);
-            hstart = std::max(hstart, 0);
+            int wend = std::min(wstart + sizeX, imgSizeW);
             wstart = std::max(wstart, 0);
-            dend = std::min(dend, static_cast<int>(imgSizeD));
-            hend = std::min(hend, static_cast<int>(imgSizeH));
-            wend = std::min(wend, static_cast<int>(imgSizeW));
+            int poolSize = (dend - dstart) * (hend - hstart) * (wend - wstart);
             CHECK(poolSize);
             for (int d = dstart; d < dend; ++d) {
               for (int h = hstart; h < hend; ++h) {
@@ -2476,8 +2451,8 @@ void CpuMatrix::avgPool3DBackward(Matrix& input,
         }
       }
       // offset
-      outData += imgSizeD * imgSizeH * imgSizeW;
-      inData += outputD * outputH * outputW;
+      outData += inLength;
+      inData += outLength;
     }
   }
 }

From 56b1b7014294f6e4c82817a81b2b939b794d1069 Mon Sep 17 00:00:00 2001
From: Yancey <yancey1989@gmail.com>
Date: Mon, 18 Sep 2017 17:56:03 +0800
Subject: [PATCH 31/34] Split operator with CPU kernel (#4046)

Split Op CPU Kernel
---
 paddle/operators/split_op.cc                  | 118 ++++++++++++++++++
 paddle/operators/split_op.h                   |  62 +++++++++
 python/paddle/v2/framework/tests/op_test.py   |   7 +-
 .../v2/framework/tests/test_split_op.py       |  26 ++++
 4 files changed, 211 insertions(+), 2 deletions(-)
 create mode 100644 paddle/operators/split_op.cc
 create mode 100644 paddle/operators/split_op.h
 create mode 100644 python/paddle/v2/framework/tests/test_split_op.py

diff --git a/paddle/operators/split_op.cc b/paddle/operators/split_op.cc
new file mode 100644
index 0000000000..61296f5c81
--- /dev/null
+++ b/paddle/operators/split_op.cc
@@ -0,0 +1,118 @@
+/* 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/split_op.h"
+#include "paddle/operators/net_op.h"
+
+namespace paddle {
+namespace operators {
+using framework::Tensor;
+
+class SplitOp : public framework::OperatorWithKernel {
+ public:
+  using framework::OperatorWithKernel::OperatorWithKernel;
+
+ protected:
+  void InferShape(const framework::InferShapeContext &ctx) const override {
+    // infershape
+    auto *in = ctx.Input<framework::Tensor>("X");
+    auto outs = ctx.MultiOutput<framework::LoDTensor>("Out");
+    size_t axis = static_cast<size_t>(ctx.Attr<int>("axis"));
+    size_t num = static_cast<size_t>(ctx.Attr<int>("num"));
+    std::vector<int> sections =
+        static_cast<std::vector<int>>(ctx.Attr<std::vector<int>>("sections"));
+    const size_t n = outs.size();
+
+    if (num > 0) {
+      int64_t in_axis_dim = in->dims()[axis];
+      PADDLE_ENFORCE_EQ(in_axis_dim % num, 0,
+                        "tensor split does not result"
+                        " in an equal division");
+      size_t out_axis_dim = in_axis_dim / num;
+      for (size_t i = 0; i < n; ++i) {
+        auto dim = in->dims();
+        dim[axis] = out_axis_dim;
+        outs[i]->Resize(dim);
+      }
+    } else if (sections.size() > 0) {
+      PADDLE_ENFORCE_EQ(sections.size(), n,
+                        "tensor split sections size"
+                        "should be equal to output size.");
+      for (size_t i = 0; i < n; ++i) {
+        auto dim = in->dims();
+        dim[axis] = sections[i];
+        outs[i]->Resize(dim);
+      }
+    } else {
+      PADDLE_ENFORCE_NOT_NULL(nullptr, "split operator should",
+                              " specify indices or sections.");
+    }
+  }
+};
+
+class SplitOpMaker : public framework::OpProtoAndCheckerMaker {
+ public:
+  SplitOpMaker(framework::OpProto *proto, framework::OpAttrChecker *op_checker)
+      : OpProtoAndCheckerMaker(proto, op_checker) {
+    AddInput("X", "the input tensor of split operator.");
+    AddOutput("Out", "the output tensors of split operator.").AsDuplicable();
+    AddComment(R"DOC(
+      Split the input tensor into multiple sub-tensors.
+      Example:
+        Input = [[1,2],
+                 [3,4],
+                 [5,6]]
+        sections = [2,1]
+        axis = 0
+        Output[0] = [[1,2],
+                     [3,4]]
+        Output[1] = [[5,6]]
+
+    )DOC");
+    AddAttr<std::vector<int>>("sections",
+                              "the length for each"
+                              "output along with the specify axis.")
+        .SetDefault(std::vector<int>{});
+    AddAttr<int>("num",
+                 "number of the sub-tensors, it must evenly divide "
+                 "Input.dims()[axis]")
+        .SetDefault(0);
+    AddAttr<int>("axis", "The axis which the input will be splited on.")
+        .SetDefault(0);
+  }
+};
+
+class SplitOpGrad : public NetOp {
+ public:
+  SplitOpGrad(const std::string &type, const framework::VariableNameMap &inputs,
+              const framework::VariableNameMap &outputs,
+              const framework::AttributeMap &attrs)
+      : NetOp(type, inputs, outputs, attrs) {
+    auto out_grad = Inputs(framework::GradVarName("Out"));
+    auto x_grad = Output(framework::GradVarName("X"));
+    AppendOp(framework::OpRegistry::CreateOp("concat", {{"X", out_grad}},
+                                             {{"Out", {x_grad}}}, attrs));
+    CompleteAddOp(false);
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle
+
+namespace ops = paddle::operators;
+USE_CPU_ONLY_OP(concat);
+REGISTER_OP(split, ops::SplitOp, ops::SplitOpMaker, split_grad,
+            ops::SplitOpGrad);
+REGISTER_OP_CPU_KERNEL(split,
+                       ops::SplitKernel<paddle::platform::CPUPlace, float>);
diff --git a/paddle/operators/split_op.h b/paddle/operators/split_op.h
new file mode 100644
index 0000000000..860690ee89
--- /dev/null
+++ b/paddle/operators/split_op.h
@@ -0,0 +1,62 @@
+/* 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 <vector>
+#include "paddle/framework/op_registry.h"
+
+namespace paddle {
+namespace operators {
+
+template <typename Place, typename T>
+class SplitKernel : public framework::OpKernel {
+ public:
+  void Compute(const framework::ExecutionContext& ctx) const override {
+    auto* in = ctx.Input<framework::Tensor>("X");
+    auto outs = ctx.MultiOutput<framework::Tensor>("Out");
+    int64_t axis = static_cast<int64_t>(ctx.Attr<int>("axis"));
+    size_t before = 1, after = 1;
+    const size_t n = outs.size();
+    size_t input_axis_dim = in->dims()[axis];
+
+    for (int64_t i = 0; i < in->dims().size(); ++i) {
+      if (i == axis) {
+        continue;
+      }
+      if (i < axis) {
+        before *= in->dims()[i];
+      } else {
+        after *= in->dims()[i];
+      }
+    }
+    size_t input_offset = 0;
+    for (size_t i = 0; i < n; i++) {
+      auto& out = outs[i];
+      size_t axis_dim = out->dims()[axis];
+      for (size_t j = 0; j < before; j++) {
+        size_t len = axis_dim * after * sizeof(T);
+        T* dest =
+            out->mutable_data<T>(platform::CPUPlace()) + axis_dim * after * j;
+        const T* src =
+            in->data<T>() + input_offset + input_axis_dim * after * j;
+        memcpy(dest, src, len);
+      }
+      input_offset += axis_dim * after;
+    }
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle
diff --git a/python/paddle/v2/framework/tests/op_test.py b/python/paddle/v2/framework/tests/op_test.py
index 75b689982a..6bbea22c5f 100644
--- a/python/paddle/v2/framework/tests/op_test.py
+++ b/python/paddle/v2/framework/tests/op_test.py
@@ -194,10 +194,13 @@ class OpTest(unittest.TestCase):
         for out_name, out_dup in Operator.get_op_outputs(self.op.type()):
             if out_dup:
                 sub_out = self.outputs[out_name]
-                for sub_out_name, sub_out_array in sub_out:
+                if not isinstance(sub_out, list):
+                    raise AssertionError("sub_out type %s is not list",
+                                         type(sub_out))
+
+                for sub_out_name, expect in sub_out:
                     actual = np.array(
                         self.scope.find_var(sub_out_name).get_tensor())
-                    expect = sub_out_array
                     self.assertTrue(
                         np.allclose(
                             actual, expect, atol=1e-05),
diff --git a/python/paddle/v2/framework/tests/test_split_op.py b/python/paddle/v2/framework/tests/test_split_op.py
new file mode 100644
index 0000000000..b4420db9d7
--- /dev/null
+++ b/python/paddle/v2/framework/tests/test_split_op.py
@@ -0,0 +1,26 @@
+import unittest
+import numpy as np
+from op_test import OpTest
+
+
+class TestSplitOp(OpTest):
+    def setUp(self):
+        self.op_type = "split"
+        axis = 0
+        num = 2
+        x = np.random.random((4, 2)).astype('float32')
+        out = np.split(x, num, axis)
+        self.inputs = {'X': x}
+        self.attrs = {'axis': axis, 'num': num}
+        self.outputs = {'Out': [('out%d' % i, out[i]) \
+            for i in xrange(len(out))]}
+
+    def test_check_output(self):
+        self.check_output()
+
+    def test_check_grad(self):
+        self.check_grad(['X'], ['out0', 'out1'])
+
+
+if __name__ == '__main__':
+    unittest.main()

From 73da97207e065d57e3abd33f9d9f6ca5c5840fbd Mon Sep 17 00:00:00 2001
From: Luo Tao <luotao02@baidu.com>
Date: Mon, 18 Sep 2017 17:40:14 +0800
Subject: [PATCH 32/34] fix AvgPoolBackward and using CaffeMode in
 test_AvgPoolFwdBwd

---
 paddle/cuda/src/hl_cuda_cnn.cu           | 15 ++++++++++-----
 paddle/math/tests/test_matrixCompare.cpp | 10 ++++------
 2 files changed, 14 insertions(+), 11 deletions(-)

diff --git a/paddle/cuda/src/hl_cuda_cnn.cu b/paddle/cuda/src/hl_cuda_cnn.cu
index 596dd6d1bf..58674febdc 100644
--- a/paddle/cuda/src/hl_cuda_cnn.cu
+++ b/paddle/cuda/src/hl_cuda_cnn.cu
@@ -298,11 +298,13 @@ __global__ void KeAvgPoolBackward(const int nthreads,
 
     for (int ph = phstart; ph < phend; ++ph) {
       int hstart = ph * strideH - padH;
-      int hend = min(hstart + sizeY, height + padH);
+      int hend = min(hstart + sizeY, height);
+      hstart = max(hstart, 0);
       for (int pw = pwstart; pw < pwend; ++pw) {
         // figure out the pooling size
         int wstart = pw * strideW - padW;
-        int wend = min(wstart + sizeX, width + padW);
+        int wend = min(wstart + sizeX, width);
+        wstart = max(wstart, 0);
         int poolsize = (hend - hstart) * (wend - wstart);
         gradient += outGrad[ph * pooledW + pw] / poolsize;
       }
@@ -708,14 +710,17 @@ __global__ void KeAvgPool3DBackward(const int nthreads,
 
     for (int pd = pdstart; pd < pdend; ++pd) {
       int dstart = pd * strideD - padD;
-      int dend = min(dstart + sizeZ, depth + padD);
+      int dend = min(dstart + sizeZ, depth);
+      dstart = max(dstart, 0);
       for (int ph = phstart; ph < phend; ++ph) {
         int hstart = ph * strideH - padH;
-        int hend = min(hstart + sizeY, height + padH);
+        int hend = min(hstart + sizeY, height);
+        hstart = max(hstart, 0);
         for (int pw = pwstart; pw < pwend; ++pw) {
           // figure out the pooling size
           int wstart = pw * strideW - padW;
-          int wend = min(wstart + sizeX, width + padW);
+          int wend = min(wstart + sizeX, width);
+          wstart = max(wstart, 0);
           int poolsize = (dend - dstart) * (hend - hstart) * (wend - wstart);
           gradient += outGrad[(pd * pooledH + ph) * pooledW + pw] / poolsize;
         }
diff --git a/paddle/math/tests/test_matrixCompare.cpp b/paddle/math/tests/test_matrixCompare.cpp
index 103f06acc5..061fb22e3f 100644
--- a/paddle/math/tests/test_matrixCompare.cpp
+++ b/paddle/math/tests/test_matrixCompare.cpp
@@ -825,9 +825,8 @@ void testMaxPoolFwdBwd(int numSamples,
                        int strideW,
                        int padH,
                        int padW) {
-  int outH = 0, outW = 0;
-  outH = (imgSizeH - ksizeH + 2 * padH + strideH - 1) / strideH + 1;
-  outW = (imgSizeW - ksizeW + 2 * padW + strideW - 1) / strideW + 1;
+  int outH = outputSize(imgSizeH, ksizeH, padH, strideH, true);
+  int outW = outputSize(imgSizeW, ksizeW, padW, strideW, true);
 
   int inWidth = imgSizeH * imgSizeW * channels;
   MatrixPtr input = CpuMatrix::create(numSamples, inWidth, false, false);
@@ -927,9 +926,8 @@ void testAvgPoolFwdBwd(int numSamples,
                        int strideW,
                        int padH,
                        int padW) {
-  int outH = 0, outW = 0;
-  outH = (imgSizeH - ksizeH + 2 * padH + strideH - 1) / strideH + 1;
-  outW = (imgSizeW - ksizeW + 2 * padW + strideW - 1) / strideW + 1;
+  int outH = outputSize(imgSizeH, ksizeH, padH, strideH, true);
+  int outW = outputSize(imgSizeW, ksizeW, padW, strideW, true);
 
   int inWidth = imgSizeH * imgSizeW * channels;
   MatrixPtr input = CpuMatrix::create(numSamples, inWidth, false, false);

From cc7a7557a27d6cb335c044bc6abcfdc834c109e1 Mon Sep 17 00:00:00 2001
From: Luo Tao <luotao02@baidu.com>
Date: Mon, 18 Sep 2017 20:09:01 +0800
Subject: [PATCH 33/34] set CudnnPooling don't include padded values in default

---
 paddle/cuda/include/hl_cuda_cudnn.h      | 6 +++---
 paddle/cuda/src/hl_cuda_cudnn.cc         | 6 +++---
 paddle/gserver/layers/CudnnPoolLayer.cpp | 4 ++--
 3 files changed, 8 insertions(+), 8 deletions(-)

diff --git a/paddle/cuda/include/hl_cuda_cudnn.h b/paddle/cuda/include/hl_cuda_cudnn.h
index 3f68c62de6..b44b071bd1 100644
--- a/paddle/cuda/include/hl_cuda_cudnn.h
+++ b/paddle/cuda/include/hl_cuda_cudnn.h
@@ -22,10 +22,10 @@ limitations under the License. */
  */
 typedef enum {
   HL_POOLING_MAX = 0,
-  // average includes padded values
-  HL_POOLING_AVERAGE = 1,
   // average does not include padded values
-  HL_POOLING_AVERAGE_EXCLUDE_PADDING = 2,
+  HL_POOLING_AVERAGE = 1,
+  // average includes padded values
+  HL_POOLING_AVERAGE_INCLUDE_PADDING = 2,
   HL_POOLING_END
 } hl_pooling_mode_t;
 
diff --git a/paddle/cuda/src/hl_cuda_cudnn.cc b/paddle/cuda/src/hl_cuda_cudnn.cc
index f38ef69255..b8caf48f9c 100644
--- a/paddle/cuda/src/hl_cuda_cudnn.cc
+++ b/paddle/cuda/src/hl_cuda_cudnn.cc
@@ -432,11 +432,11 @@ void hl_create_pooling_descriptor(hl_pooling_descriptor* pooling_desc,
       cudnn_mode = CUDNN_POOLING_MAX;
       break;
     case HL_POOLING_AVERAGE:
-      cudnn_mode = CUDNN_POOLING_AVERAGE_COUNT_INCLUDE_PADDING;
-      break;
-    case HL_POOLING_AVERAGE_EXCLUDE_PADDING:
       cudnn_mode = CUDNN_POOLING_AVERAGE_COUNT_EXCLUDE_PADDING;
       break;
+    case HL_POOLING_AVERAGE_INCLUDE_PADDING:
+      cudnn_mode = CUDNN_POOLING_AVERAGE_COUNT_INCLUDE_PADDING;
+      break;
     default:
       LOG(FATAL) << "parameter mode error";
   }
diff --git a/paddle/gserver/layers/CudnnPoolLayer.cpp b/paddle/gserver/layers/CudnnPoolLayer.cpp
index 4adb2d4709..810a1af2d0 100644
--- a/paddle/gserver/layers/CudnnPoolLayer.cpp
+++ b/paddle/gserver/layers/CudnnPoolLayer.cpp
@@ -29,9 +29,9 @@ bool CudnnPoolLayer::typeCheck(const std::string &poolType,
     if (mode) {
       *mode = HL_POOLING_AVERAGE;
     }
-  } else if (poolType == "cudnn-avg-excl-pad-pool") {
+  } else if (poolType == "cudnn-avg-incl-pad-pool") {
     if (mode) {
-      *mode = HL_POOLING_AVERAGE_EXCLUDE_PADDING;
+      *mode = HL_POOLING_AVERAGE_INCLUDE_PADDING;
     }
   } else {
     return false;

From 780e26826df8f81b47c9209a17f8bff2d932e5e8 Mon Sep 17 00:00:00 2001
From: Luo Tao <luotao02@baidu.com>
Date: Mon, 18 Sep 2017 21:24:17 +0800
Subject: [PATCH 34/34] remove unused CHECK in matrix.cpp

---
 paddle/math/Matrix.cpp | 2 --
 1 file changed, 2 deletions(-)

diff --git a/paddle/math/Matrix.cpp b/paddle/math/Matrix.cpp
index 83c06a04c6..0023b4d0f5 100644
--- a/paddle/math/Matrix.cpp
+++ b/paddle/math/Matrix.cpp
@@ -1080,7 +1080,6 @@ void GpuMatrix::maxPoolBackward(Matrix& inputMat,
   size_t channels = outV.getWidth() / outputH / outputW;
   CHECK(imgSizeH * imgSizeW * channels == inputMat.getWidth());
   CHECK(height_ == inputMat.getHeight());
-  CHECK(width_ == imgSizeW * imgSizeH * channels);
   CHECK(outGrad.getHeight() == outV.getHeight() &&
         outGrad.getWidth() == outV.getWidth());
 
@@ -1259,7 +1258,6 @@ void GpuMatrix::maxPool3DBackward(Matrix& outGrad,
   size_t frameNum = getHeight();
   size_t channels = outGrad.getWidth() / outputD / outputH / outputW;
   CHECK(imgSizeD * imgSizeH * imgSizeW * channels == getWidth());
-  CHECK(width_ == imgSizeD * imgSizeH * imgSizeW * channels);
   CHECK(outGrad.getHeight() == maxPoolIdx.getHeight() &&
         outGrad.getWidth() == maxPoolIdx.getWidth());