Merge branch 'develop' of https://github.com/PaddlePaddle/Paddle into crop_op

Conflicts:
	paddle/pybind/pybind.cc

doc/design/ops/images/2_level_rnn.dot

+digraph G {
+
+  rnn [label="1-th level RNN" shape=box]
+
+  subgraph cluster0 {
+    label = "time step 0"
+
+    sent0 [label="sentence"]
+    sent1 [label="sentence"]
+
+    rnn1 [label="2-th level RNN" shape=box]
+
+    sent0 -> rnn1
+    sent1 -> rnn1
+  }
+
+  subgraph cluster1 {
+    label = "time step 1"
+
+    sent2 [label="sentence"]
+    sent3 [label="sentence"]
+
+    rnn2 [label="2-th level RNN" shape=box]
+
+    sent2 -> rnn2
+    sent3 -> rnn2
+  }
+
+  subgraph cluster2 {
+    label = "time step 2"
+
+    sent4 [label="sentence"]
+    sent5 [label="sentence"]
+
+    rnn3 [label="2-th level RNN" shape=box]
+
+    sent4 -> rnn3
+    sent5 -> rnn3
+  }
+
+
+  para0 [label="paragraph info 0"]
+  para1 [label="paragraph info 1"]
+  para2 [label="paragraph info 2"]
+
+  rnn1 -> para0
+  rnn2 -> para1
+  rnn3 -> para2
+
+  para0 -> rnn
+  para1 -> rnn
+  para2 -> rnn
+
+  chapter [label="chapter info"]
+  rnn -> chapter
+}

doc/design/ops/images/rnn.dot

+digraph G {
+  label = "simple RNN implementation" 
+
+  ranksep=2;
+
+  //graph [nodesep=1, ranksep=1];
+
+  node[nodesep=1]
+
+  subgraph cluster0 {
+    label = "global scope"
+    rankdir = TB
+    W
+    boot_memory
+    input
+    output
+  }
+
+  subgraph cluster1 {
+    label = "step-scope 0"
+    rankdir = TB
+    memory0[label="memory"]
+    prememory0[label="pre-memory"]
+    step_input0[label="step input"]
+    step_output0[label="step output"]
+  }
+
+  subgraph cluster2 {
+    label = "step-scope 1"
+    rankdir = TB
+    memory1[label="memory"]
+    prememory1[label="pre-memory"]
+    step_input1[label="step input"]
+    step_output1[label="step output"]
+  }
+
+  subgraph cluster3 {
+    label = "step-scope 2"
+    rankdir = TB
+    memory2[label="memory"]
+    prememory2[label="pre-memory"]
+    step_input2[label="step input"]
+    step_output2[label="step output"]
+  }
+
+  stepnet [shape=box]
+  stepnet0 [shape=box, style=dashed]
+  stepnet1 [shape=box, style=dashed]
+  stepnet2 [shape=box, style=dashed]
+
+
+  edge[color=blue]
+  boot_memory -> prememory0 [label="init" color="blue"]
+  memory0 -> prememory1  [label="copy/reference" color="blue"]
+  memory1 -> prememory2 [label="copy/reference" color="blue"]
+
+  edge[color=black]
+  W -> stepnet0[constraint=false, style=dashed]
+  W -> stepnet1[constraint=false, style=dashed]
+  W -> stepnet2[constraint=false, style=dashed]
+
+  memory0 -> stepnet0[style=dashed]
+  prememory0 -> stepnet0 -> step_output0[style=dashed]
+
+  memory1 -> stepnet1[style=dashed]
+  prememory1 -> stepnet1 -> step_output1[style=dashed]
+
+  memory2 -> stepnet2[style=dashed]
+  prememory2 -> stepnet2 -> step_output2[style=dashed]
+
+  input -> step_input0
+  input -> step_input1
+  input -> step_input2
+
+  step_input0 -> stepnet0 [style=dashed]
+  step_input1 -> stepnet1[style=dashed]
+  step_input2 -> stepnet2[style=dashed]
+
+  step_output0 -> output
+  step_output1 -> output
+  step_output2 -> output
+
+  stepnet0 -> stepnet[style=dashed]
+  stepnet1 -> stepnet[style=dashed]
+  stepnet2 -> stepnet[style=dashed]
+
+}

doc/design/ops/images/rnn_2level_data.dot

+digraph G {
+  chapter [label="chapter"]
+
+  subgraph cluster0 {
+    label = "paragraph 0"
+
+    top_rnn0[label="top rnn step 0" shape=box]
+
+    p0 [label="paragraph 0"]
+    p1 [label="paragraph 1"]
+  }
+
+  subgraph cluster1{
+    label = "paragraph 1"
+
+    top_rnn1[label="top rnn step 1" shape=box]
+
+    p2 [label="paragraph 0"]
+    p3 [label="paragraph 1"]
+  }
+
+  subgraph cluster_p0 {
+    label = "sentence 0"
+
+    low_rnn0 [label="low rnn step 0" shape=box]
+    s00 [label="sentence 0"]
+    s01 [label="sentence 1"]
+
+    low_rnn0 -> s00
+    low_rnn0 -> s01
+  }
+
+  subgraph cluster_p1 {
+    label = "sentence 1"
+    low_rnn1 [label="low rnn step 1" shape=box]
+    s10 [label="sentence 0"]
+    s11 [label="sentence 1"]
+    low_rnn1 -> s10
+    low_rnn1 -> s11
+  }
+
+  subgraph cluster_p2 {
+    label = "sentence 1"
+    low_rnn2 [label="low rnn step 0" shape=box]
+    s20 [label="sentence 0"]
+    s21 [label="sentence 1"]
+    low_rnn2 -> s20
+    low_rnn2 -> s21
+  }
+
+  subgraph cluster_p3 {
+    label = "sentence 1"
+    low_rnn3 [label="low rnn step 1" shape=box]
+    s30 [label="sentence 0"]
+    s31 [label="sentence 1"]
+    low_rnn3 -> s30
+    low_rnn3 -> s31
+  }
+
+
+  chapter -> top_rnn0
+  chapter -> top_rnn1
+
+  top_rnn0 -> p0
+  top_rnn0 -> p1
+  top_rnn1 -> p2
+  top_rnn1 -> p3
+
+
+  p0 -> low_rnn0
+  p1 -> low_rnn1
+  p2 -> low_rnn2
+  p3 -> low_rnn3
+
+}

doc/design/ops/rnn.md

+# RNNOp design
+
+This document is about an RNN operator which requires that instances in a mini-batch have the same length.  We will have a more flexible RNN operator.
+
+## RNN Algorithm Implementation
+
+<p aligh="center">
+<img src="./images/rnn.jpg"/>
+</p>
+
+The above diagram shows an RNN unrolled into a full network.
+
+There are several important concepts:
+
+- *step-net*: the sub-graph to run at each step,
+- *memory*, $h_t$, the state of the current step,
+- *ex-memory*, $h_{t-1}$, the state of the previous step,
+- *initial memory value*, the ex-memory of the first step.
+
+### Step-scope
+
+There could be local variables defined in step-nets.  PaddlePaddle runtime realizes these variables in *step-scopes* -- scopes created for each step.
+
+<p aligh="center">
+<img src="./images/rnn.png"/><br/>
+Figure 2 the RNN's data flow
+</p>
+
+Please be aware that all steps run the same step-net.  Each step
+
+1. creates the step-scope,
+2. realizes local variables, including step-outputs, in the step-scope, and
+3. runs the step-net, which could use these variables.
+
+The RNN operator will compose its output from step outputs in step scopes.
+
+### Memory and Ex-memory
+
+Let's give more details about memory and ex-memory via a simply example:
+
+$$
+h_t = U h_{t-1} + W x_t
+$$,
+
+where $h_t$ and $h_{t-1}$ are the memory and ex-memory of step $t$'s respectively.
+
+In the implementation, we can make an ex-memory variable either "refers to" the memory variable of the previous step,
+or copy the value of the previous memory value to the current ex-memory variable.
+
+### Usage in Python
+
+For more information on Block, please refer to the [design doc](https://github.com/PaddlePaddle/Paddle/blob/develop/doc/design/block.md).
+
+We can define an RNN's step-net using Block:
+
+```python
+import paddle as pd
+
+X = some_op() # x is some operator's output, and is a LoDTensor
+a = some_op()
+
+# declare parameters
+W = pd.Variable(shape=[20, 30])
+U = pd.Variable(shape=[20, 30])
+
+rnn = pd.create_rnn_op(output_num=1)
+with rnn.stepnet():
+    x = rnn.add_input(X)
+    # declare a memory (rnn's step)
+    h = rnn.add_memory(init=a)
+    # h.pre_state() means previous memory of rnn
+    new_state = pd.add_two( pd.matmul(W, x) + pd.matmul(U, h.pre_state()))
+    # update current memory
+    h.update(new_state)
+    # indicate that h variables in all step scopes should be merged
+    rnn.add_outputs(h)
+
+out = rnn()
+```
+
+Python API functions in above example:
+
+- `rnn.add_input` indicates the parameter is a variable that will be segmented into step-inputs.
+- `rnn.add_memory` creates a variable used as the memory.
+- `rnn.add_outputs` mark the variables that will be concatenated across steps into the RNN output.
+
+### Nested RNN and LoDTensor
+
+An RNN whose step-net includes other RNN operators is known as an *nested RNN*.
+
+For example, we could have a 2-level RNN, where the top level corresponds to paragraphs, and the lower level corresponds to sentences.
+
+The following figure illustrates the feeding of text into the lower level, one sentence each step, and the feeding of step outputs to the top level. The final top level output is about the whole text.
+
+<p aligh="center">
+<img src="./images/2_level_rnn.png"/>
+</p>
+
+```python
+import paddle as pd
+
+W = pd.Variable(shape=[20, 30])
+U = pd.Variable(shape=[20, 30])
+
+W0 = pd.Variable(shape=[20, 30])
+U0 = pd.Variable(shape=[20, 30])
+
+# a is output of some op
+a = some_op()
+
+# chapter_data is a set of 128-dim word vectors
+# the first level of LoD is sentence
+# the second level of LoD is chapter
+chapter_data = pd.Variable(shape=[None, 128], type=pd.lod_tensor, level=2)
+
+def lower_level_rnn(paragraph):
+    '''
+    x: the input
+    '''
+    rnn = pd.create_rnn_op(output_num=1)
+    with rnn.stepnet():
+        sentence = rnn.add_input(paragraph, level=0)
+        h = rnn.add_memory(shape=[20, 30])
+        h.update(
+            pd.matmul(W, sentence) + pd.matmul(U, h.pre_state()))
+        # get the last state as sentence's info
+        rnn.add_outputs(h)
+    return rnn
+
+top_level_rnn = pd.create_rnn_op(output_num=1)
+with top_level_rnn.stepnet():
+    paragraph_data = rnn.add_input(chapter_data, level=1)
+    low_rnn = lower_level_rnn(paragraph_data)
+    paragraph_out = low_rnn()
+
+    h = rnn.add_memory(init=a)
+    h.update(
+        pd.matmul(W0, paragraph_data) + pd.matmul(U0, h.pre_state()))
+    top_level_rnn.add_outputs(h)
+
+# just output the last step
+chapter_out = top_level_rnn(output_all_steps=False)
+```
+
+in above example, the construction of the `top_level_rnn` calls  `lower_level_rnn`.  The input is a LoD Tensor. The top level RNN segments input text data into paragraphs, and the lower level RNN segments each paragraph into sentences.
+
+By default, the `RNNOp` will concatenate the outputs from all the time steps,
+if the `output_all_steps` set to False, it will only output the final time step.
+
+
+<p align="center">
+<img src="images/rnn_2level_data.png"/>
+</p>

paddle/operators/accuracy_op.cc

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+    http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License. */
+
+#include "paddle/operators/accuracy_op.h"
+
+namespace paddle {
+namespace operators {
+
+class AccuracyOp : public framework::OperatorWithKernel {
+ public:
+  using framework::OperatorWithKernel::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("Label"),
+                            "Input of Inference must be initialized.");
+    auto *inference = ctx.Input<framework::Tensor>("Inference");
+    auto *label = ctx.Input<framework::Tensor>("Label");
+
+    PADDLE_ENFORCE_EQ(label->dims().size(), 1, "label must be a vector");
+    PADDLE_ENFORCE_EQ(inference->dims()[0], label->dims()[0],
+                      "inference size must be the same as label size");
+
+    ctx.Output<Tensor>("Accuracy")->Resize({1});
+  }
+};
+
+class AccuracyOpMaker : public framework::OpProtoAndCheckerMaker {
+ public:
+  AccuracyOpMaker(framework::OpProto *proto,
+                  framework::OpAttrChecker *op_checker)
+      : OpProtoAndCheckerMaker(proto, op_checker) {
+    // TODO(typhoonzero): support both inference value and indices.
+    AddInput("Inference", "topk(indices) the network output");
+    AddInput("Label", "Label of the training data");
+    // TODO(typhoonzero): AddInput("Weight", ...
+    AddOutput("Accuracy", "The accuracy of current batch");
+
+    AddComment(
+        R"DOC(Accuracy. It will print accuracy rate for classification.
+The accuracy is:
+..  math::
+accuracy = \\frac{NumOfCorrectPredicts}{NumOfAllSamples})DOC");
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle
+
+namespace ops = paddle::operators;
+REGISTER_OP_WITHOUT_GRADIENT(accuracy, ops::AccuracyOp, ops::AccuracyOpMaker);
+REGISTER_OP_CPU_KERNEL(accuracy,
+                       ops::AccuracyKernel<paddle::platform::CPUPlace, float>);

paddle/operators/accuracy_op.cu

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+    http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License. */
+
+#include "paddle/operators/accuracy_op.h"
+
+namespace paddle {
+namespace operators {
+
+__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;
+        break;
+      }
+    }
+  }
+  *accuracy = static_cast<float>(correct) / static_cast<float>(N);
+}
+
+template <typename T>
+class AccuracyOpCUDAKernel : public framework::OpKernel {
+ public:
+  void Compute(const framework::ExecutionContext& ctx) const override {
+    PADDLE_ENFORCE(platform::is_gpu_place(ctx.GetPlace()),
+                   "It must use GPUPlace.");
+    auto* inference = ctx.Input<Tensor>("Inference");
+    auto* label = ctx.Input<Tensor>("Label");
+    auto* accuracy = ctx.Output<Tensor>("Accuracy");
+    // FIXME(typhoonzero): only support indices currently
+    // if add support for output values, how to detect the data type?
+    const int* inference_data = inference->data<int>();
+    const int* label_data = label->data<int>();
+    float* accuracy_data = accuracy->mutable_data<float>(ctx.GetPlace());
+
+    size_t num_samples = inference->dims()[0];
+    size_t infer_width = inference->dims()[1];
+    cudaMemset((void**)&accuracy_data, 0, sizeof(float));
+
+    if (num_samples == 0) {
+      return;
+    }
+
+    AccuracySingleKernel<<<1, 1>>>(num_samples, infer_width, 1, inference_data,
+                                   label_data, accuracy_data);
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle
+
+REGISTER_OP_GPU_KERNEL(accuracy,
+                       paddle::operators::AccuracyOpCUDAKernel<float>);

paddle/operators/accuracy_op.h

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+    http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License. */
+
+#pragma once
+#include <algorithm>
+#include "paddle/framework/eigen.h"
+#include "paddle/framework/op_registry.h"
+
+namespace paddle {
+namespace operators {
+
+using Tensor = framework::Tensor;
+
+template <typename T, int MajorType = Eigen::RowMajor,
+          typename IndexType = Eigen::DenseIndex>
+using EigenMatrix = framework::EigenMatrix<T, MajorType, IndexType>;
+
+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 EigenScalar = framework::EigenScalar<T, MajorType, IndexType>;
+
+template <typename Place, typename T>
+class AccuracyKernel : public framework::OpKernel {
+ public:
+  void Compute(const framework::ExecutionContext& ctx) const override {
+    auto* inference = ctx.Input<Tensor>("Inference");
+    auto* label = ctx.Input<Tensor>("Label");
+    auto* accuracy = ctx.Output<Tensor>("Accuracy");
+
+    float* accuracy_data = accuracy->mutable_data<float>(ctx.GetPlace());
+
+    const T* inference_data = inference->data<T>();
+    const T* label_data = label->data<T>();
+
+    size_t num_samples = inference->dims()[0];
+    size_t class_dim = inference->dims()[1];
+    *accuracy_data = 0.0f;
+
+    if (num_samples == 0) {
+      return;
+    }
+
+    int num_correct = 0;
+    // assume inference is already the topk of the output
+    for (size_t i = 0; i < num_samples; ++i) {
+      PADDLE_ENFORCE_GE(label_data[i], 0, "label must >= 0");
+      for (size_t j = 0; j < class_dim; ++j) {
+        if (inference_data[i * class_dim + j] == label_data[i]) {
+          ++num_correct;
+          break;
+        }
+      }
+    }
+
+    // FIXME(typhoonzero): we don't accumulate the accuracy for now.
+    *accuracy_data =
+        static_cast<float>(num_correct) / static_cast<float>(num_samples);
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle

paddle/operators/crop_op.cc

@@ -98,7 +98,8 @@ Out = [[1, 2],
     AddAttr<std::vector<int>>("shape",
                               "A list<int> describing the shape of output."
                               "The size of shape list should be as same as "
-                              "dimension size of  input X.");
+                              "dimension size of  input X.")
+        .SetDefault(std::vector<int>());
   }
 };
 
@@ -113,9 +114,10 @@ class CropOpGrad : public framework::OperatorWithKernel {
                             "Input(Out@GRAD) should not be null");
     auto x_dims = ctx.Input<Tensor>("X")->dims();
     auto *x_grad = ctx.Output<Tensor>(framework::GradVarName("X"));
-
+    if (x_grad != nullptr) {
       x_grad->Resize(x_dims);
     }
+  }
 };
 
 int64_t transIndex(std::vector<int64_t> out_shape, std::vector<int64_t> x_shape,

paddle/operators/crop_op.cu

@@ -13,12 +13,9 @@
    limitations under the License. */
 
 #define EIGEN_USE_GPU
+#include <stdio.h>
 #include "paddle/operators/crop_op.h"
 
-#define CUDA_1D_KERNEL_LOOP(i, n)                            \
-  for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < n; \
-       i += blockDim.x * gridDim.x)
-
 namespace paddle {
 namespace operators {
 
@@ -28,39 +25,52 @@ template <typename T, int D>
 __global__ void CropKernel(const int N, const int64_t* out_shape,
                            const int64_t* x_shape, const int* crop_rules,
                            const T* x_data, T* out_data) {
-  CUDA_1D_KERNEL_LOOP(index, N) {
-    // int64_t dim_size = out_shape.size();
   int64_t pos[D];
-
+  int tmp;
+  int64_t x_index;
+  for (int out_index = blockIdx.x * blockDim.x + threadIdx.x; out_index < N;
+       out_index += blockDim.x * gridDim.x) {
+    tmp = out_index;
     for (int64_t i = D - 1; i >= 0; --i) {
-      pos[i] = (index % out_shape[i]) + crop_rules[i * 2];
-      index = index / out_shape[i];
+      pos[i] = (tmp % out_shape[i]) + crop_rules[i * 2];
+      tmp = tmp / out_shape[i];
     }
 
-    int64_t result = pos[0];
+    x_index = pos[0];
     for (size_t i = 1; i < D; ++i) {
-      result = result * x_shape[i] + pos[i];
+      x_index = x_index * x_shape[i] + pos[i];
     }
-
-    out_data[index] = x_data[result];
+    out_data[out_index] = x_data[x_index];
   }
 }
 
 template <typename T, int D>
 void CropCUDAFunctoin(const framework::ExecutionContext& context) {
+  PADDLE_ENFORCE(platform::is_gpu_place(context.GetPlace()),
+                 "It must use GPUPlace.");
   auto* x = context.Input<Tensor>("X");
   auto* out = context.Output<Tensor>("Out");
   auto x_data = x->data<T>();
-  T* out_data = out->mutable_data<T>(paddle::platform::CPUPlace());
+  T* out_data = out->mutable_data<T>(paddle::platform::GPUPlace());
   auto x_dims = x->dims();
   auto out_dims = out->dims();
   int64_t out_count = framework::product(out_dims);
-  int64_t* x_shape = &(framework::vectorize(x_dims))[0];
-  int64_t* out_shape = &(framework::vectorize(out_dims))[0];
-
+  int64_t x_shape[D];
+  int64_t out_shape[D];
+  for (int i = 0; i < D; ++i) {
+    x_shape[i] = x_dims[i];
+    out_shape[i] = out_dims[i];
+  }
+  int64_t* x_shape_gpu;
+  int64_t* out_shape_gpu;
+  cudaMalloc((void**)&x_shape_gpu, sizeof(int64_t) * D);
+  cudaMemcpy(x_shape_gpu, x_shape, sizeof(int64_t) * D, cudaMemcpyHostToDevice);
+  cudaMalloc((void**)&out_shape_gpu, sizeof(int64_t) * D);
+  cudaMemcpy(out_shape_gpu, out_shape, sizeof(int64_t) * D,
+             cudaMemcpyHostToDevice);
   auto offsets = context.op().Attr<std::vector<int>>("offsets");
   PADDLE_ENFORCE_EQ(
-      x_dims.size(), offsets.size(),
+      D, offsets.size(),
       "Offsets size should be equal to dimension size of input tensor.");
 
   int crop_rules[D * 2];
@@ -69,13 +79,20 @@ void CropCUDAFunctoin(const framework::ExecutionContext& context) {
     crop_rules[i * 2 + 1] = x_dims[i] - out_dims[i] - offsets[i];
   }
 
+  int* crop_rules_gpu;
+  cudaMalloc((void**)&crop_rules_gpu, sizeof(int) * D * 2);
+  cudaMemcpy(crop_rules_gpu, crop_rules, sizeof(int) * D * 2,
+             cudaMemcpyHostToDevice);
+
   int n = out_dims[0];
   int d = out_dims[1];
   int block = 512;
   int grid = (n * d + block - 1) / block;
-
-  CropKernel<T, D><<<grid, block>>>(out_count, out_shape, x_shape, crop_rules,
-                                    x_data, out_data);
+  CropKernel<T, D><<<grid, block>>>(out_count, out_shape_gpu, x_shape_gpu,
+                                    crop_rules_gpu, x_data, out_data);
+  cudaFree(crop_rules_gpu);
+  cudaFree(x_shape_gpu);
+  cudaFree(out_shape_gpu);
 }
 
 template <typename T>

paddle/operators/crop_op.h

@@ -30,22 +30,21 @@ template <typename Place, typename T, size_t D>
 void CropGradFunction(const framework::ExecutionContext& context) {
   auto* d_out = context.Input<Tensor>(framework::GradVarName("Out"));
   auto* d_x = context.Output<Tensor>(framework::GradVarName("X"));
+  if (d_x != nullptr) {
     d_x->mutable_data<T>(context.GetPlace());
     auto d_x_dims = d_x->dims();
     auto d_out_dims = d_out->dims();
-
     auto offsets = context.op().Attr<std::vector<int>>("offsets");
-
     Eigen::array<std::pair<int, int>, D> paddings;
     for (int i = 0; i < d_out_dims.size(); ++i) {
       paddings[i].first = offsets[i];
       paddings[i].second = d_x_dims[i] - d_out_dims[i] - offsets[i];
     }
-
     auto d_x_tensor = EigenTensor<T, D>::From(*d_x);
     auto d_out_tensor = EigenTensor<T, D>::From(*d_out);
     auto place = context.GetEigenDevice<Place>();
     d_x_tensor.device(place) = d_out_tensor.pad(paddings, 0);
+  }
 }
 
 template <typename Place, typename T>

paddle/operators/name_convention.md

@@ -38,9 +38,11 @@ public:
   AccumulateOpMaker(framework::OpProto *proto,
                             framework::OpAttrChecker *op_checker)
     : OpProtoAndCheckerMaker(proto, op_checker) {
-    AddInput("X", "(Tensor) The input tensor that has to be accumulated to the output tensor. If the output size is not the same as input size, the output tensor is first reshaped and initialized to zero, and only then, accumulation is done.");
+    AddInput("X", "(Tensor) The input tensor that has to be accumulated to the output tensor. 
+    If the output size is not the same as input size, 
+    the output tensor is first reshaped and initialized to zero, and only then, accumulation is done.");
     AddOutput("Out", "(Tensor) Accumulated output tensor");
-    AddAttr<float>("gamma", "(float, default 1.0) Accumulation multiplier");
+    AddAttr<float>("gamma", "(float, default 1.0) Accumulation multiplier").SetDefault(1.0f);
     AddComment(R"DOC(
 Accumulate operator accumulates the input tensor to the output tensor. If the
 output tensor already has the right size, we add to it; otherwise, we first
@@ -51,7 +53,7 @@ Accumulation is done as shown:
 
 Out = 1*X + gamma*Out
 
-where X is the input tensor, Y is the output tensor and gamma is the multiplier
+where X is the input tensor, Out is the output tensor and gamma is the multiplier
 argument.
 )DOC");
   }

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

+import unittest
+import numpy as np
+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")
+        self.inputs = {'Inference': infer, "Label": label}
+        num_correct = 0
+        for rowid in xrange(32):
+            for ele in infer[rowid]:
+                if ele == label[rowid]:
+                    num_correct += 1
+                    break
+        self.outputs = {'Accuracy': [num_correct / 32.0]}
+
+    def test_check_output(self):
+        self.check_output()
+
+
+if __name__ == '__main__':
+    unittest.main()

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

 import unittest
 import numpy as np
-from paddle.v2.framework.op import Operator
-from gradient_checker import GradientChecker
-from op_test_util import OpTestMeta
+from op_test import OpTest
 
 
 def crop(data, offsets, crop_shape):
@@ -26,57 +24,68 @@ def crop(data, offsets, crop_shape):
     return np.array(result).reshape(crop_shape)
 
 
-class TCropOp(OpTest):
+class TestCropOp(OpTest):
     def setUp(self):
-        self.initTestCase()
-        self.type = "crop"
-        self.inputs = {'X': np.random.random(self.shape).astype("float32"), }
+        self.op_type = "crop"
+        self.crop_by_input = False
         self.attrs = {}
+        self.initTestCase()
         self.attrs['offsets'] = self.offsets
+        if self.crop_by_input:
+            self.inputs = {
+                'X': np.random.random(self.x_shape).astype("float32"),
+                'Y': np.random.random(self.crop_shape).astype("float32")
+            }
+        else:
             self.attrs['shape'] = self.crop_shape
+            self.inputs = {
+                'X': np.random.random(self.x_shape).astype("float32"),
+            }
         self.outputs = {
             'Out': crop(self.inputs['X'], self.offsets, self.crop_shape)
         }
-        print "input=%s" % self.inputs['X']
 
     def initTestCase(self):
-        self.shape = (8, 8, 8)
-        self.crop_shape = [2, 2, 2]
-        self.offsets = [0, 0, 0]
-
-
-#class TCase1(TCropOp):
-#    def initTestCase(self):
-#        self.shape = (16, 16, 16)
-#        self.crop_shape = [2, 2, 3]
-#        self.offsets = [1, 5, 3]
-
-#class TCropGradOp(GradientChecker):
-
-#    def initTestCase(self):
-#        self.shape = (4, 4)
-#        self.crop_shape = [2, 2]
-#        self.offsets = [0, 0]
-
-#    def setUp(self):
-#        self.initTestCase()
-#        self.op = Operator(
-#            type="crop", X="X", Out="Out", offsets=self.offsets, shape=self.crop_shape)
-#        self.inputs = {'X': np.random.random(self.shape).astype("float32"), }
-#
-#    def test_normal(self):
-#        self.check_grad(
-#            self.op, self.inputs, set(["X"]), "Out", max_relative_error=0.5)
-
-#def test_cpu_gpu_compare(self):
-#    self.compare_grad(self.op, self.inputs)
-
-#class TestGradCase1(TestCropGradOp):
-
-#    def initTestCase(self):
-#        self.shape = (16, 16)
-#        self.crop_shape = [8, 8]
-#        self.offsets = [1, 1]
+        self.x_shape = (8, 8)
+        self.crop_shape = [2, 2]
+        self.offsets = [1, 2]
+
+    def test_check_output(self):
+        self.check_output()
+
+    def test_check_grad_normal(self):
+        self.check_grad(['X'], 'Out', max_relative_error=0.006)
+
+
+class TestCase1(TestCropOp):
+    def initTestCase(self):
+        self.x_shape = (16, 16, 16)
+        self.crop_shape = [2, 2, 3]
+        self.offsets = [1, 5, 3]
+
+
+class TestCase2(TestCropOp):
+    def initTestCase(self):
+        self.x_shape = (4, 4)
+        self.crop_shape = [4, 4]
+        self.offsets = [0, 0]
+
+
+class TestCase3(TestCropOp):
+    def initTestCase(self):
+        self.x_shape = (16, 16, 16)
+        self.crop_shape = [2, 2, 3]
+        self.offsets = [1, 5, 3]
+        self.crop_by_input = True
+
+
+class TestCase4(TestCropOp):
+    def initTestCase(self):
+        self.x_shape = (4, 4)
+        self.crop_shape = [4, 4]
+        self.offsets = [0, 0]
+        self.crop_by_input = True
+
 
 if __name__ == '__main__':
     unittest.main()