update data types and test cases

.travis.yml

@@ -30,6 +30,7 @@ addons:
       - automake
       - libtool
       - ccache
+  ssh_known_hosts: 52.76.173.135
 before_install:
   - if [[ "$JOB" == "check_style" ]]; then sudo ln -s /usr/bin/clang-format-3.8 /usr/bin/clang-format; fi
   # Paddle is using protobuf 3.1 currently. Protobuf 3.2 breaks the compatibility. So we specify the python
@@ -42,6 +43,14 @@ script:
   - |
     timeout 2580 paddle/scripts/travis/${JOB}.sh # 43min timeout
     RESULT=$?; if [ $RESULT -eq 0 ] || [ $RESULT -eq 142 ]; then true; else false; fi;
+  - |
+    if [[ "$JOB" != "build_doc" ]]; then exit 0; fi;
+    if [[ "$TRAVIS_PULL_REQUEST" != "false" ]]; then exit 0; fi;
+    if [[ "$TRAVIS_BRANCH" != "develop"  && ! "$TRAVIS_BRANCH" =~ ^v[[:digit:]]+\.[[:digit:]]+(\.[[:digit:]]+)?(-\S*)?$ ]]; then exit 0; fi;
+    export DEPLOY_DOCS_SH=https://raw.githubusercontent.com/PaddlePaddle/PaddlePaddle.org/master/scripts/deploy/deploy_docs.sh
+    export DOCS_DIR=`pwd`
+    cd ..
+    curl $DEPLOY_DOCS_SH | bash -s $CONTENT_DEC_PASSWD $TRAVIS_BRANCH $DOCS_DIR $DOCS_DIR/build/doc   
 notifications:
   email:
     on_success: change

benchmark/IntelOptimizedPaddle.md

+# Benchmark
+
+Machine:
+
+- Server
+ 	- Intel(R) Xeon(R) Gold 6148 CPU @ 2.40GHz, 2 Sockets, 20 Cores per socket
+- Laptop
+ 	- DELL XPS15-9560-R1745: i7-7700HQ 8G 256GSSD
+ 	- i5 MacBook Pro (Retina, 13-inch, Early 2015)
+- Desktop
+ 	- i7-6700k
+
+System: CentOS release 6.3 (Final), Docker 1.12.1.
+
+PaddlePaddle: paddlepaddle/paddle:latest (TODO: will rerun after 0.11.0)
+
+- MKL-DNN tag v0.10
+- MKLML 2018.0.20170720
+- OpenBLAS v0.2.20
+	 
+On each machine, we will test and compare the performance of training on single node using MKL-DNN / MKLML / OpenBLAS respectively.
+
+## Benchmark Model
+
+### Server
+Test on batch size 64, 128, 256 on Intel(R) Xeon(R) Gold 6148 CPU @ 2.40GHz
+
+Input image size - 3 * 224 * 224, Time: images/second
+
+- VGG-19
+
+| BatchSize    | 64    | 128  | 256     |
+|--------------|-------| -----| --------|
+| OpenBLAS     | 7.82  | 8.62  | 10.34  | 
+| MKLML        | 11.02 | 12.86 | 15.33  |
+| MKL-DNN      | 27.69 | 28.8 | 29.27  |
+
+
+chart on batch size 128
+TBD
+
+ - ResNet
+ - GoogLeNet
+
+### Laptop
+TBD
+### Desktop
+TBD

doc/design/parameter_average.md

+# Averaging Parameter in PaddlePaddle
+
+## Why Averaging
+In a large scale machine learning setup where the size of the training data is huge, it could take us a large number of iterations over the training data before we can achieve the optimal values of parameters of our model. Looking at the problem setup, it is desirable if we can obtain the optimal values of parameters by going through the data in as few passes as we can.
+
+Polyak and Juditsky (1992) showed that the test performance of simple average of parameters obtained by Stochastic Gradient Descent (SGD) is as good as that of parameter values that are obtained by training the model over and over again, over the training dataset.
+
+Hence, to accelerate the speed of Stochastic Gradient Descent, Averaged Stochastic Gradient Descent (ASGD) was proposed in Polyak and Juditsky (1992). For ASGD, the running average of parameters obtained by SGD, is used as the estimator for <img src="./images/theta_star.gif"/><br/> . The averaging is done as follows:
+
+<img src="./images/asgd.gif" align="center"/><br/>
+
+We propose averaging for any optimizer similar to how ASGD performs it, as mentioned above.
+
+### How to perform Parameter Averaging in PaddlePaddle
+
+Parameter Averaging in PaddlePaddle works in the following way during training :
+1. It will take in an instance of a normal optimizer as an input, e.g. RMSPropOptimizer
+2. The optimizer itself is responsible for updating the parameters.
+3. The ParameterAverageOptimizer maintains a separate copy of the parameters for itself:
+    1. In concept, the values of this copy are the average of the values of the parameters in the most recent N batches.
+    2. However, saving all the N instances of the parameters in memory is not feasible.
+    3. Therefore, an approximation algorithm is used.
+
+Hence, overall we have have two copies of the parameters: one for the optimizer itself, and one for the ParameterAverageOptimizer. The former should be used in back propagation, while the latter should be used during testing and should be saved.
+
+During the testing/ saving the model phase, we perform the following steps:
+1. Perform the delayed operations.
+2. Save current values of the parameters to a temporary variable.
+3. Replace the values of the parameters with the averaged values.
+4. Perform testing and/or save the parameters.
+5. Restore the values of the parameters once done.
+
+### How to implement Averaging of Parameter in PaddlePaddle
+
+We can add the ParameterAverageOptimizer op to the graph through Python API. Using this approach, we manually add this op to the graph and direct the output of the optimizer op to this op during training.
+
+	**Advantages**:
+    - Allows for greater flexibility to the users of PaddlePaddle. Using this approach, the users can plug different optimizers into ParameterAverageOptimizer by passing in the optimizer to the op.
+    - Makes it easy for the users to customize and extend the framework.
+
+	**Disadvantages**:
+    - Implementation requires re-writing the averaging methodology in Python.  
+
+### Low-Level implementation
+
+In the new design, we propose to create a new operation for averaging parameter updates (ParameterAverageOptimizer). For now, we can add an op that takes in the following as input:
+- the optimizer
+- the window_size to keep the updates
+
+The ParameterAverageOptimizer op can be like any other operator with its own CPU/GPU implementation either using Eigen or separate CPU and GPU kernels. As the initial implementation, we can implement the kernel using Eigen following the abstraction pattern implemented for [Operators](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/operators/rmsprop_op.h). We also want to support the case when the Trainer/Optimizer runs on the GPU while ParameterAverageOptimizer runs on a CPU.
+
+The idea of building an op for averaging is in sync with the refactored PaddlePaddle philosophy of using operators to represent any computation unit. The way the op will be added to the computation graph will be decided by the [layer functions](https://github.com/PaddlePaddle/Paddle/blob/develop/doc/design/python_api.md#layer-function) in Python API.
+
+### Python API implementation for ParameterAverageOptimizer
+
+Based on Polyak and Juditsky (1992), we can generalize the averaging of updates to any optimizer. The input to the op would be the following:
+- Any optimizer (RMSProp , AdaGrad etc.)
+- A window size. The op keeps accumulating updated parameter values over a window of N batches and takes an average. Move the averaged value to a buffer when window is full to avoid loss of precision.
+
+Using the ParameterAverageOptimizer op, any user can add the operation to their computation graphs. However, this will require a lot of lines of code and we should design Python APIs that support averaging. As per the PaddlePaddle [Python API design](https://github.com/PaddlePaddle/Paddle/blob/develop/doc/design/python_api.md), the layer functions are responsible for creating operators, operator parameters and variables. Since ParameterAverageOptimizer will be an operator, it makes sense to create it in the layer functions.
+We will have a wrapper written in Python that will support the functionality and implement the actual core computation in C++ core as we have done for other [Optimizers](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/operators/rmsprop_op.cc)
+
+#### Creation of the ParameterAverageOptimizer operator
+There are two ways for creating the ParameterAverageOptimizer op:
+1. We create the op immediately while building the computation graph.
+2. We add the op in a lazy manner, just before the backward pass, similar to the way the optimization ops are added.
+
+The proposal is to add the op immediately while building the computation graph.
+
+#### High-level API
+
+In PaddlePaddle Python API, users will primarily rely on [layer functions](https://github.com/PaddlePaddle/Paddle/blob/develop/doc/design/python_api.md#layer-function) to create neural network layers. Hence, we also need to provide parameter average functionality in layer functions.

doc/howto/cross_compiling/cross_compiling_for_android.md

+# Build PaddlePaddle for Android
+
+There are two approaches to build PaddlePaddle for Android: using Docker and on Linux without Docker. 
+
+## Cross-Compiling Using Docker
+
+Docker-based cross-compiling is the recommended approach because Docker runs on all major operating systems, including Linux, Mac OS X, and Windows.
+
+### Build the Docker Image
+
+The following steps pack all the tools that we need to build PaddlePaddle into a Docker image.
+
+```bash
+$ git clone https://github.com/PaddlePaddle/Paddle.git
+$ cd Paddle
+$ docker build -t paddle:dev-android . -f Dockerfile.android
+```
+
+### Build the Inference Library
+
+We can run the Docker image we just created to build the inference library of PaddlePaddle for Android using the command below:
+
+```bash
+$ docker run -it --rm -v $PWD:/paddle -e "ANDROID_ABI=armeabi-v7a" -e "ANDROID_API=21" paddle:dev-android
+```
+
+The Docker image accepts two arguments `ANDROID_ABI` and `ANDROID_API`:
+
+| Argument        | Optional Values         | Default |
+|-----------------|-------------------------|---------|
+|`ANDROID_ABI`    |`armeabi-v7a, arm64-v8a` | `armeabi-v7a` |
+|`ANDROID_API`    |`>= 21` | `21` |
+
+The ARM-64 architecture (`arm64-v8a`) requires at least level 21 of Android API.
+
+The default entry-point of the Docker image, [`paddle/scripts/docker/build_android.sh`](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/scripts/docker/build_android.sh) generates the [Android cross-compiling standalone toolchain](https://developer.android.com/ndk/guides/standalone_toolchain.html) based on the argument: `ANDROID_ABI` or `ANDROID_API`.  For information about other configuration arguments, please continue reading.
+
+The above command generates and outputs the inference library in `$PWD/install_android` and puts third-party libraries in `$PWD/install_android/third_party`.
+
+## Cross-Compiling on Linux
+
+The Linux-base approach to cross-compile is to run steps in `Dockerfile.android` manually on a Linux x64 computer.
+
+### Setup the Environment
+
+To build for Android's, we need [Android NDK](
+https://developer.android.com/ndk/downloads/index.html):
+
+```bash
+wget -q https://dl.google.com/android/repository/android-ndk-r14b-linux-x86_64.zip
+unzip -q android-ndk-r14b-linux-x86_64.zip
+```
+
+Android NDK includes everything we need to build the [*standalone toolchain*](https://developer.android.com/ndk/guides/standalone_toolchain.html), which in then used to build PaddlePaddle for Android.  (We plan to remove the intermediate stage of building the standalone toolchain in the near future.)
+
+- To build the standalone toolchain for `armeabi-v7a` and Android API level 21:
+
+  ```bash
+  your/path/to/android-ndk-r14b-linux-x86_64/build/tools/make-standalone-toolchain.sh \
+          --arch=arm --platform=android-21 --install-dir=your/path/to/arm_standalone_toolchain
+  ```
+  
+  The generated standalone toolchain will be in `your/path/to/arm_standalone_toolchain`.
+
+- To build the standalone toolchain for `arm64-v8a` and Android API level 21:
+
+  ```bash
+  your/path/to/android-ndk-r14b-linux-x86_64/build/tools/make-standalone-toolchain.sh \
+          --arch=arm64 --platform=android-21 --install-dir=your/path/to/arm64_standalone_toolchain
+  ```
+
+  The generated standalone toolchain will be in `your/path/to/arm64_standalone_toolchain`.
+
+**Please be aware that the minimum level of Android API required by PaddlePaddle is 21.**
+
+### Cross-Compiling Arguments
+
+CMake supports [choosing the toolchain](https://cmake.org/cmake/help/v3.0/manual/cmake-toolchains.7.html#cross-compiling).  PaddlePaddle provides [`android.cmake`](https://github.com/PaddlePaddle/Paddle/blob/develop/cmake/cross_compiling/android.cmake), which configures the Android cross-compiling toolchain for CMake.  `android.cmake` is not required for CMake >= 3.7, which support Android cross-compiling. PaddlePaddle detects the CMake version, for those newer than 3.7, it uses [the official version](https://cmake.org/cmake/help/v3.7/manual/cmake-toolchains.7.html#cross-compiling).
+
+Some other CMake arguments you need to know:
+
+- `CMAKE_SYSTEM_NAME` must be `Android`.  This tells PaddlePaddle's CMake system to cross-compile third-party dependencies.  This also changes some other CMake arguments like `WITH_GPU=OFF`, `WITH_AVX=OFF`, `WITH_PYTHON=OFF`, and `WITH_RDMA=OFF`.
+- `WITH_C_API` must be `ON`, to build the C-based inference library for Android.
+- `WITH_SWIG_PY` must be `OFF` because the Android platform doesn't support SWIG-based API.
+
+Some Android-specific arguments:
+
+- `ANDROID_STANDALONE_TOOLCHAIN`: the absolute path of the Android standalone toolchain, or the path relative to the CMake build directory.  PaddlePaddle's CMake extensions would derive the cross-compiler, sysroot and Android API level from this argument.
+- `ANDROID_TOOLCHAIN`: could be `gcc` or `clang`.  The default value is `clang`.
+  - For CMake >= 3.7, it should anyway be `clang`.  For older versions, it could be `gcc`.
+  - Android's official `clang` requires `glibc` >= 2.15.
+- `ANDROID_ABI`: could be `armeabi-v7a` or `arm64-v8a`.  The default value is `armeabi-v7a`.
+- `ANDROID_NATIVE_API_LEVEL`: could be derived from the value of `ANDROID_STANDALONE_TOOLCHAIN`.
+- `ANROID_ARM_MODE`:
+  - could be `ON` or `OFF`, and defaults to `ON`, when `ANDROID_ABI=armeabi-v7a`;
+  - no need to specify when `ANDROID_ABI=arm64-v8a`.
+- `ANDROID_ARM_NEON`: indicates if to use NEON instructions.
+  - could be `ON` or `OFF`, and defaults to `ON`, when `ANDROID_ABI=armeabi-v7a`;
+  - no need to specify when `ANDROID_ABI=arm64-v8a`.
+
+Other useful arguments:
+
+- `USE_EIGEN_FOR_BLAS`: indicates if using Eigen.  Could be `ON` or `OFF`, defaults to `OFF`.
+- `HOST_C/CXX_COMPILER`: specifies the host compiler, which is used to build the host-specific protoc and target-specific OpenBLAS.  It defaults to the value of the environment variable `CC`, or `cc`.
+
+Some frequent configurations for your reference:
+
+```bash
+cmake -DCMAKE_SYSTEM_NAME=Android \
+      -DANDROID_STANDALONE_TOOLCHAIN=your/path/to/arm_standalone_toolchain \
+      -DANDROID_ABI=armeabi-v7a \
+      -DANDROID_ARM_NEON=ON \
+      -DANDROID_ARM_MODE=ON \
+      -DUSE_EIGEN_FOR_BLAS=ON \
+      -DCMAKE_INSTALL_PREFIX=your/path/to/install \
+      -DWITH_C_API=ON \
+      -DWITH_SWIG_PY=OFF \
+      ..
+```
+
+```
+cmake -DCMAKE_SYSTEM_NAME=Android \
+      -DANDROID_STANDALONE_TOOLCHAIN=your/path/to/arm64_standalone_toolchain \
+      -DANDROID_ABI=arm64-v8a \
+      -DUSE_EIGEN_FOR_BLAS=OFF \
+      -DCMAKE_INSTALL_PREFIX=your/path/to/install \
+      -DWITH_C_API=ON \
+      -DWITH_SWIG_PY=OFF \
+      ..
+```
+
+
+There are some other arguments you might want to configure.
+
+- `CMAKE_BUILD_TYPE=MinSizeRel` minimizes the size of library.
+- `CMAKE_BUILD_TYPE-Release` optimizes the runtime performance.
+
+Our own tip for performance optimization to use clang and Eigen or OpenBLAS:
+- `CMAKE_BUILD_TYPE=Release`
+- `ANDROID_TOOLCHAIN=clang`
+- `USE_EIGEN_BLAS=ON` for `armeabi-v7a`, or `USE_EIGEN_FOR_BLAS=OFF` for `arm64-v8a`.
+
+### Build and Install
+
+After running `cmake`, we can run `make; make install` to build and install.
+
+Before building, you might want to remove the `third_party` and `build` directories including pre-built libraries for other architectures.
+
+After building，in the directory `CMAKE_INSTALL_PREFIX`, you will find three sub-directories:
+
+- `include`: the header file of the inference library,
+- `lib`: the inference library built for various Android ABIs,
+- `third_party`: dependent third-party libraries built for Android.

doc/howto/cross_compiling/cross_compiling_for_android_cn.md

 # 构建Android平台上的PaddlePaddle库
 
 用户可通过如下两种方式，交叉编译Android平台上适用的PaddlePaddle库：
-- 基于Docker容器的编译方式
+- 基于Docker容器的编译方式
 - 基于Linux交叉编译环境的编译方式
 
 ## 基于Docker容器的编译方式
@@ -26,14 +26,14 @@ Android的Docker开发镜像向用户提供两个可配置的参数：
 |`ANDROID_API`    |`>= 21` | `21` |
 
 - 编译`armeabi-v7a`，`Android API 21`的PaddlePaddle库
-```bash
-$ docker run -it --rm -v $PWD:/paddle -e "ANDROID_ABI=armeabi-v7a" -e "ANDROID_API=21" username/paddle-android:dev
-```
+  ```bash
+  $ docker run -it --rm -v $PWD:/paddle -e "ANDROID_ABI=armeabi-v7a" -e "ANDROID_API=21" username/paddle-android:dev
+  ```
 
-- 编译`arm64-v8a`，`Android API 21`的PaddlePaddle库
-```bash
-$ docker run -it --rm -v $PWD:/paddle -e "ANDROID_ABI=arm64-v8a" -e "ANDROID_API=21" username/paddle-android:dev
-```
+- 编译`arm64-v8a`，`Android API 21`的PaddlePaddle库
+  ```bash
+  $ docker run -it --rm -v $PWD:/paddle -e "ANDROID_ABI=arm64-v8a" -e "ANDROID_API=21" username/paddle-android:dev
+  ```
 
 执行上述`docker run`命令时，容器默认执行[paddle/scripts/docker/build_android.sh](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/scripts/docker/build_android.sh)脚本。该脚本中记录了交叉编译Android版PaddlePaddle库常用的CMake配置，并且会根据`ANDROID_ABI`和`ANDROID_API`自动构建独立工具链、进行编译和安装。由于arm64架构要求Android API不小于21。因此当`ANDROID_ABI=arm64-v8a`，`ANDROID_API<21`时，Docker容器中将默认使用`Android API 21`的编译工具链。用户可以参考下文**配置交叉编译参数**章节，根据个人的需求修改定制Docker容器所执行的脚本。编译安装结束之后，PaddlePaddle的C-API库将被安装到`$PWD/install_android`目录，所依赖的第三方库同时也被安装到`$PWD/install_android/third_party`目录。
 
@@ -82,16 +82,16 @@ CMake系统对交叉编译提供了支持[cmake-toolchains](https://cmake.org/cm
 Android平台可选配置参数：
 
 - `ANDROID_STANDALONE_TOOLCHAIN`，独立工具链所在的绝对路径，或者相对于构建目录的相对路径。PaddlePaddle的CMake系统将根据该值自动推导和设置需要使用的交叉编译器、sysroot、以及Android API级别；否则，用户需要在cmake时手动设置这些值。无默认值。
-- `ANDROID_TOOLCHAIN`，目标工具链。可设置`gcc/clang`，默认值为`clang`。
-	- CMake 3.7以上，将会始终使用`clang`工具链；CMake 3.7以下，可设置`ANDROID_TOOLCHAIN=gcc`以使用`gcc`工具链。
+- `ANDROID_TOOLCHAIN`，目标工具链。可设置`gcc/clang`，默认值为`clang`。
+	- CMake 3.7以上，将会始终使用`clang`工具链；CMake 3.7以下，可设置`ANDROID_TOOLCHAIN=gcc`以使用`gcc`工具链。
 	- Android官方提供的`clang`编译器要求系统支持`GLIBC 2.15`以上。
 - `ANDROID_ABI`，目标架构ABI。目前支持`armeabi-v7a`和`arm64-v8a`，默认值为`armeabi-v7a`。
 - `ANDROID_NATIVE_API_LEVEL`，工具链的Android API级别。若没有显式设置，PaddlePaddle将根据`ANDROID_STANDALONE_TOOLCHAIN`的值自动推导得到。
-- `ANROID_ARM_MODE`，是否使用ARM模式。
-	- `ANDROID_ABI=armeabi-v7a`时，可设置`ON/OFF`，默认值为`ON`；
+- `ANROID_ARM_MODE`，是否使用ARM模式。
+	- `ANDROID_ABI=armeabi-v7a`时，可设置`ON/OFF`，默认值为`ON`；
 	- `ANDROID_ABI=arm64-v8a`时，不需要设置。
-- `ANDROID_ARM_NEON`，是否使用NEON指令。
-	- `ANDROID_ABI=armeabi-v7a`时，可设置`ON/OFF`，默认值为`ON`；
+- `ANDROID_ARM_NEON`，是否使用NEON指令。
+	- `ANDROID_ABI=armeabi-v7a`时，可设置`ON/OFF`，默认值为`ON`；
 	- `ANDROID_ABI=arm64-v8a`时，不需要设置。
 
 其他配置参数：
@@ -119,7 +119,7 @@ cmake -DCMAKE_SYSTEM_NAME=Android \
       -DANDROID_STANDALONE_TOOLCHAIN=your/path/to/arm64_standalone_toolchain \
       -DANDROID_ABI=arm64-v8a \
       -DUSE_EIGEN_FOR_BLAS=OFF \
-      -DCMAKE_INSTALL_PREFIX=your/path/to/install \  
+      -DCMAKE_INSTALL_PREFIX=your/path/to/install \
       -DWITH_C_API=ON \
       -DWITH_SWIG_PY=OFF \
       ..
@@ -128,8 +128,8 @@ cmake -DCMAKE_SYSTEM_NAME=Android \
 用户还可根据自己的需求设置其他编译参数。比如希望最小化生成的库的大小，可以设置`CMAKE_BUILD_TYPE`为`MinSizeRel`；若希望最快的执行速度，则可设置`CMAKE_BUILD_TYPE`为`Release`。亦可以通过手动设置`CMAKE_C/CXX_FLAGS_MINSIZEREL/RELEASE`来影响PaddlePaddle的编译过程。
 
 **性能TIPS**，为了达到最快的计算速度，在CMake参数配置上，有以下建议：
-- 设置`CMAKE_BUILD_TYPE`为`Release`
-- 使用`clang`编译工具链
+- 设置`CMAKE_BUILD_TYPE`为`Release`
+- 使用`clang`编译工具链
 - `armeabi-v7a`时，设置`USE_EIGEN_BLAS=ON`，使用Eigen进行矩阵计算；`arm64-v8a`时，设置`USE_EIGEN_FOR_BLAS=OFF`，使用OpenBLAS进行矩阵计算
 
 ### 编译和安装

doc/howto/cross_compiling/cross_compiling_for_ios_cn.md

+# 构建iOS平台上的PaddlePaddle库
+交叉编译iOS平台上适用的PaddlePaddle库，需要在MacOS系统上进行。本文的将介绍在MacOS上，从源码交叉编译iOS平台上适用的PaddlePaddle库。
+
+## 准备交叉编译环境
+Apple官方为iOS开发提供了完整的交叉编译工具和集成开发环境，用户从App Store下载安装Xcode即可。也可自行前往官网下载，[Xcode](https://developer.apple.com/cn/xcode/)。安装完成之后，可在命令行执行`xcodebuild -version`，判断是否安装成功。
+
+```bash
+$ xcodebuild -version
+Xcode 9.0
+Build version 9A235
+```
+
+## 配置交叉编译参数
+
+PaddlePaddle为交叉编译提供了工具链配置文档[cmake/cross_compiling/ios.cmake](https://github.com/PaddlePaddle/Paddle/blob/develop/cmake/cross_compiling/ios.cmake)，以提供一些默认的编译器和编译参数配置。
+
+交叉编译iOS版本的PaddlePaddle库时，有一些必须配置的参数：
+
+- `CMAKE_SYSTEM_NAME`，CMake编译的目标平台，必须设置为`iOS`。在设置`CMAKE_SYSTEM_NAME=iOS`后，PaddlePaddle的CMake系统会自动编译所有的第三方依赖库，并且强制设置一些PaddlePaddle参数的值（`WITH_C_API=ON`、`WITH_GPU=OFF`、`WITH_AVX=OFF`、`WITH_PYTHON=OFF`、`WITH_RDMA=OFF`）。
+- `WITH_C_API`，是否编译C-API预测库，必须设置为ON。在iOS平台上只支持使用C-API来预测。
+- `WITH_SWIG_PY`，必须设置为ON。在iOS平台上不支持通过swig调用来训练或者预测。
+
+iOS平台可选配置参数：
+
+- `IOS_PLATFORM`，可设置为`OS/SIMULATOR`，默认值为`OS`。
+  - `OS`，构建目标为`arm`架构的iPhone或者iPad等物理设备。
+  - `SIMULATOR`，构建目标为`x86`架构的模拟器平台。
+- `IOS_ARCH`，目标架构。针对不同的`IOS_PLATFORM`，可设置的目标架构如下表所示：
+
+   | IOS_PLATFORM | IOS_ARCH             |
+   |--------------|----------------------|
+   |   OS         | armv7, armv7s, arm64 (默认) |
+   | SIMULATOR    | i386, x86_64 (默认)         |   
+
+- `IOS_DEPLOYMENT_TARGET`，最小的iOS部署版本，默认值为`7.0`。
+- `IOS_ENABLE_BITCODE`，是否使能[Bitcode](https://developer.apple.com/library/content/documentation/IDEs/Conceptual/AppDistributionGuide/AppThinning/AppThinning.html#//apple_ref/doc/uid/TP40012582-CH35-SW3)，可设置`ON/OFF`，默认值为`ON`。
+- `IOS_USE_VECLIB_FOR_BLAS`，是否使用[vecLib](https://developer.apple.com/documentation/accelerate/veclib)框架进行BLAS矩阵计算，可设置`ON/OFF`，默认值为`OFF`。
+- `IOS_DEVELOPMENT_ROOT`，`Developer`目录，可显式指定为`/path/to/platform/Developer`。若未显式指定，PaddlePaddle将会根据`IOS_PLATFORM`自动选择`Xcode`对应`platform`的`Developer`目录。
+- `IOS_SDK_ROOT`，所使用`SDK`的根目录，可显式指定为`/path/to/platform/Developer/SDKs/SDK`。若未显式指定，PaddlePaddle将会自动选择`IOS_DEVELOPMENT_ROOT`目录下最新的`SDK`版本。
+
+其他配置参数：
+
+- `USE_EIGEN_FOR_BLAS`，是否使用Eigen库进行矩阵计算，在`IOS_USE_VECLIB_FOR_BLAS=OFF`时有效。可设置`ON/OFF`，默认值为`OFF`。
+- `HOST_C/CXX_COMPILER`，宿主机的C/C++编译器。默认值为环境变量`CC/CXX`的值；若环境变量`CC/CXX`未设置，则使用`cc/c++`编译器。
+
+常用的cmake配置如下：
+
+```bash
+cmake -DCMAKE_SYSTEM_NAME=iOS \
+      -DIOS_PLATFORM=OS \
+      -DIOS_ARCH="arm64" \
+      -DIOS_ENABLE_BITCODE=ON \
+      -DIOS_USE_VECLIB_FOR_BLAS=ON \
+      -DCMAKE_INSTALL_PREFIX=your/path/to/install \
+      -DWITH_C_API=ON \
+      -DWITH_TESTING=OFF \
+      -DWITH_SWIG_PY=OFF \
+      ..
+```
+
+```bash
+cmake -DCMAKE_SYSTEM_NAME=iOS \
+      -DIOS_PLATFORM=SIMULATOR \
+      -DIOS_ARCH="x86_64" \
+      -DIOS_USE_VECLIB_FOR_BLAS=ON \
+      -DCMAKE_INSTALL_PREFIX=your/path/to/install \
+      -DWITH_C_API=ON \
+      -DWITH_TESTING=OFF \
+      -DWITH_SWIG_PY=OFF \
+      ..
+```
+
+用户还可根据自己的需求设置其他编译参数。比如希望最小化生成库的大小，可以设置`CMAKE_BUILD_TYPE`为`MinSizeRel`；若希望得到最快的执行速度，则可设置`CMAKE_BUILD_TYPE`为`Release`。亦可以通过手动设置`CMAKE_C/CXX_FLAGS`来影响PaddlePaddle的编译过程。
+
+**性能TIPS**，为了达到最快的计算速度，在CMake参数配置上，有以下建议：
+
+- 设置`CMAKE_BUILD_TYPE`为`Release`
+- 设置`IOS_USE_VECLIB_FOR_BLAS=ON`，调用`vecLib`框架提供的BLAS函数进行矩阵计算。
+
+## 编译和安装
+
+CMake配置完成后，执行以下命令，PaddlePaddle将自动下载和编译所有第三方依赖库、编译和安装PaddlePaddle预测库。
+
+```
+$ make
+$ make install
+```
+
+注意：如果你曾在源码目录下编译过其他平台的PaddlePaddle库，请先使用`rm -rf`命令删除`third_party`目录和`build`目录，以确保所有的第三方依赖库和PaddlePaddle代码都是针对新的CMake配置重新编译的。
+
+执行完安装命令后，`your/path/to/install`目录中会包含以下内容：
+
+- `include`目录，其中包含所有C-API的头文件
+- `lib`目录，其中包含PaddlePaddle的C-API静态库
+- `third_party`目录，其中包含所依赖的所有第三方库
+
+注意，不同架构的PaddlePaddle库建议安装到不同的目录下，然后使用`lipo`工具将多个静态库合并成一个支持多个架构的fat库。
+
+自此，PaddlePaddle库已经安装完成，用户可将合成的fat库用于深度学习相关的iOS App中，调用方法见C-API文档。

doc/howto/cross_compiling/cross_compiling_for_raspberry_cn.md

@@ -59,4 +59,4 @@ make install
 
 注意：如果你曾经在源码目录下编译过其他平台的PaddlePaddle库，请先使用`rm -rf`命令删除`third_party`目录和`build`目录，以确保所有的第三方依赖库和PaddlePaddle代码都是针对新的CMake配置重新编译的。
 
-执行完安装命令后，，`your/path/to/install`目录中会包含`include`和`lib`目录，其中`include`中包含C-API的头文件，`lib`中包含一个Raspberry Pi版本的库。
+执行完安装命令后，`your/path/to/install`目录中会包含`include`和`lib`目录，其中`include`中包含C-API的头文件，`lib`中包含一个Raspberry Pi版本的库。

doc/howto/cross_compiling/cross_compiling_for_raspberry_en.md

@@ -44,7 +44,7 @@ cmake -DCMAKE_SYSTEM_NAME=RPi \
       ..
 ```
 
-To build the inference library, please set the argument WITH_API to ON: `WITH_C_API=ON`.
+To build the inference library, please set the argument WITH\_C\_API to ON: `WITH_C_API=ON`.
 
 You can add more arguments. For example, to minimize the size of the generated inference library, you may use `CMAKE_BUILD_TYPE=MinSizeRel`. For performance optimization, you may use `CMAKE_BUILD_TYPE=Release`.
 

doc/howto/usage/cluster/cluster_train_cn.md

@@ -19,7 +19,7 @@
      * [启动集群作业](#启动集群作业-1)
   * [在Kubernetes集群中提交训练作业](#在kubernetes集群中提交训练作业)
 
-# 概述
+## 概述
 本文将介绍如何使用PaddlePaddle在不同的集群框架下完成分布式训练。分布式训练架构如下图所示：
 
 <img src="https://user-images.githubusercontent.com/13348433/31772175-5f419eca-b511-11e7-9db7-5231fe3d9ccb.png" width="500">
@@ -32,7 +32,7 @@
 
 在使用同步SGD训练神经网络时，PaddlePaddle使用同步屏障（barrier），使梯度的提交和参数的更新按照顺序方式执行。在异步SGD中，则并不会等待所有trainer提交梯度才更新参数，这样极大地提高了计算的并行性：参数服务器之间不相互依赖，并行地接收梯度和更新参数，参数服务器也不会等待计算节点全部都提交梯度之后才开始下一步，计算节点之间也不会相互依赖，并行地执行模型的训练。可以看出，虽然异步SGD方式会提高参数更新并行度, 但是并不能保证参数同步更新，在任意时间某一台参数服务器上保存的参数可能比另一台要更新，与同步SGD相比，梯度会有噪声。
 
-# 环境准备
+## 环境准备
 
 1. 准备您的计算集群。计算集群通常由一组（几台到几千台规模）的Linux服务器组成。服务器之间可以通过局域网（LAN）联通，每台服务器具有集群中唯一的IP地址（或者可被DNS解析的主机名）。集群中的每台计算机通常被成为一个“节点”。
 1. 我们需要在集群的所有节点上安装 PaddlePaddle。 如果要启用GPU，还需要在节点上安装对应的GPU驱动以及CUDA。PaddlePaddle的安装可以参考[build_and_install](https://github.com/PaddlePaddle/Paddle/tree/develop/doc/getstarted/build_and_install)的多种安装方式。我们推荐使用[Docker](https://github.com/PaddlePaddle/Paddle/blob/develop/doc/getstarted/build_and_install/docker_install_cn.rst)安装方式来快速安装PaddlePaddle。
@@ -51,8 +51,8 @@ PaddlePaddle 0.10.0, compiled with
 
 下面以`doc/howto/usage/cluster/src/word2vec`中的代码作为实例，介绍使用PaddlePaddle v2 API完成分布式训练。
 
-# 启动参数说明
-## 启动参数服务器
+## 启动参数说明
+### 启动参数服务器
 执行以下的命令启动一个参数服务器并等待和计算节点的数据交互
 ```bash
 $ paddle pserver --port=7164 --ports_num=1 --ports_num_for_sparse=1 --num_gradient_servers=1
@@ -70,7 +70,7 @@ $ stdbuf -oL /usr/bin/nohup paddle pserver --port=7164 --ports_num=1 --ports_num
 | ports_num_for_sparse  | 必选 | 1 | 用于稀疏类型参数通信的端口个数  |
 | num_gradient_servers  | 必选 | 1 | 当前训练任务pserver总数 |
 
-## 启动计算节点
+### 启动计算节点
 执行以下命令启动使用python编写的trainer程序（文件名为任意文件名，如train.py）
 ```bash
 $ python train.py
@@ -117,7 +117,7 @@ paddle.init(
 | pservers  | 必选 | 127.0.0.1 | 当前训练任务启动的pserver的IP列表，多个IP使用“,”隔开 |
 
 
-## 准备数据集
+### 准备数据集
 
 参考样例数据准备脚本[prepare.py](https://github.com/PaddlePaddle/Paddle/tree/develop/doc/howto/usage/cluster/src/word2vec/prepare.py)，准备训练数据和验证数据集，我们使用paddle.dataset.imikolov数据集，并根据分布式训练并发数（trainer节点个数），在`prepare.py`开头部分指定`SPLIT_COUNT`将数据切分成多份。
 
@@ -149,7 +149,7 @@ test.txt-00002
 
 对于不同的训练任务，训练数据格式和训练程序的`reader()`会大不相同，所以开发者需要根据自己训练任务的实际场景完成训练数据的分割和`reader()`的编写。
 
-## 准备训练程序
+### 准备训练程序
 
 我们会对每个训练任务都会在每个节点上创建一个工作空间（workspace），其中包含了用户的训练程序、程序依赖、挂载或下载的训练数据分片。
 
@@ -184,7 +184,7 @@ test.txt-00002
 - `train_data_dir`：包含训练数据的目录，可以是从分布式存储挂载过来的，也可以是在任务启动前下载到本地的。
 - `test_data_dir`：包含测试数据集的目录。
 
-# 使用分布式计算平台或工具
+## 使用分布式计算平台或工具
 
 PaddlePaddle可以使用多种分布式计算平台构建分布式计算任务，包括：
 - [Kubernetes](http://kubernetes.io) Google开源的容器集群的调度框架，支持大规模集群生产环境的完整集群方案。
@@ -195,12 +195,12 @@ PaddlePaddle可以使用多种分布式计算平台构建分布式计算任务
 
 在使用分布式计算平台进行训练时，任务被调度在集群中时，分布式计算平台通常会通过API或者环境变量提供任务运行需要的参数，比如节点的ID、IP和任务节点个数等。
 
-## 使用Fabric启动集群作业
+### 使用Fabric启动集群作业
 
-### 准备一个Linux集群
+#### 准备一个Linux集群
 可以在`paddle/scripts/cluster_train_v2/fabric/docker_cluster`目录下，执行`kubectl -f ssh_servers.yaml`启动一个测试集群，并使用`kubectl get po -o wide`获得这些节点的IP地址。
 
-### 启动集群作业
+#### 启动集群作业
 
 `paddle.py` 提供了自动化脚本来启动不同节点中的所有 PaddlePaddle 集群进程。默认情况下，所有命令行选项可以设置为 `paddle.py` 命令选项并且 `paddle.py` 将透明、自动地将这些选项应用到 PaddlePaddle 底层进程。
 
@@ -216,10 +216,10 @@ sh run.sh
 
 集群作业将会在几秒后启动。
 
-### 终止集群作业
+#### 终止集群作业
 `paddle.py`能获取`Ctrl + C` SIGINT 信号来自动终止它启动的所有进程。只需中断 `paddle.py` 任务来终止集群作业。如果程序崩溃你也可以手动终止。
 
-### 检查集群训练结果
+#### 检查集群训练结果
 详细信息请检查 $workspace/log 里的日志，每一个节点都有相同的日志结构。
 
 `paddle_trainer.INFO`
@@ -234,13 +234,13 @@ sh run.sh
 `train.log`
 提供训练过程的 stderr 和 stdout。训练失败时可以检查错误日志。
 
-### 检查模型输出
+#### 检查模型输出
 运行完成后，模型文件将被写入节点 0 的 `output` 目录中。
 工作空间中的 `nodefile` 表示当前集群作业的节点 ID。
 
-## 在OpenMPI集群中提交训练作业
+### 在OpenMPI集群中提交训练作业
 
-### 准备OpenMPI集群
+#### 准备OpenMPI集群
 
 执行下面的命令以启动3个节点的OpenMPI集群和一个"head"节点：
 
@@ -252,7 +252,7 @@ kubectl create -f mpi-nodes.yaml
 
 然后可以从head节点ssh无密码登录到OpenMPI的每个节点上。
 
-### 启动集群作业
+#### 启动集群作业
 
 您可以按照下面的步骤在OpenMPI集群中提交paddle训练任务：
 
@@ -280,6 +280,6 @@ scp train.txt-00002 test.txt-00002 [node3IP]:/home/tutorial
 mpirun -hostfile machines -n 3  /home/tutorial/start_mpi_train.sh
 ```
 
-## 在Kubernetes集群中提交训练作业
+### 在Kubernetes集群中提交训练作业
 
 此部分的使用方法可以参考[here](../k8s/k8s_distributed_cn.md)。

doc/howto/usage/cluster/cluster_train_en.md

@@ -19,7 +19,7 @@
       * [Launching Cluster Job](#launching-cluster-job-1)
    * [Cluster Training Using Kubernetes](#cluster-training-using-kubernetes)
 
-# Introduction
+## Introduction
 
 In this article, we'll explain how to run distributed training jobs with PaddlePaddle on different types of clusters. The diagram below shows the main architecture of a distributed trainning job:
 
@@ -33,7 +33,7 @@ PaddlePaddle can support both synchronize stochastic gradient descent (SGD) and
 
 When training with synchronize SGD, PaddlePaddle uses an internal "synchronize barrier" which makes gradients update and parameter download in strict order. On the other hand, asynchronous SGD won't wait for all trainers to finish upload at a single step, this will increase the parallelism of distributed training: parameter servers do not depend on each other, they'll do parameter optimization concurrently. Parameter servers will not wait for trainers, so trainers will also do their work concurrently. But asynchronous SGD will introduce more randomness and noises in the gradient.
 
-# Preparations
+## Preparations
 1. Prepare your computer cluster. It's normally a bunch of Linux servers connected by LAN. Each server will be assigned a unique IP address. The computers in the cluster can be called "nodes".
 2. Install PaddlePaddle on every node. If you are going to take advantage of GPU cards, you'll also need to install proper driver and CUDA libraries. To install PaddlePaddle please read [this build and install](https://github.com/PaddlePaddle/Paddle/tree/develop/doc/getstarted/build_and_install) document. We strongly recommend using [Docker installation](https://github.com/PaddlePaddle/Paddle/blob/develop/doc/getstarted/build_and_install/docker_install_en.rst).
 
@@ -52,9 +52,9 @@ PaddlePaddle 0.10.0rc, compiled with
 
 We'll take `doc/howto/usage/cluster/src/word2vec` as an example to introduce distributed training using PaddlePaddle v2 API.
 
-# Command-line arguments
+## Command-line arguments
 
-## Starting parameter server
+### Starting parameter server
 
 Type the below command to start a parameter server which will wait for trainers to connect:
 
@@ -74,7 +74,7 @@ $ stdbuf -oL /usr/bin/nohup paddle pserver --port=7164 --ports_num=1 --ports_num
 | ports_num_for_sparse  | required | 1 | number of ports which serves sparse parameter update  |
 | num_gradient_servers  | required | 1 | total number of gradient servers |
 
-## Starting trainer
+### Starting trainer
 Type the command below to start the trainer(name the file whatever you want, like "train.py")
 
 ```bash
@@ -122,7 +122,7 @@ paddle.init(
 | trainer_id  | required | 0 | ID for every trainer, start from 0 |
 | pservers  | required | 127.0.0.1 | list of IPs of parameter servers, separated by "," |
 
-## Prepare Training Dataset
+### Prepare Training Dataset
 
 Here's some example code [prepare.py](https://github.com/PaddlePaddle/Paddle/tree/develop/doc/howto/usage/cluster/src/word2vec/prepare.py), it will download public `imikolov` dataset and split it into multiple files according to job parallelism(trainers count). Modify `SPLIT_COUNT` at the begining of `prepare.py` to change the count of output files.
 
@@ -155,7 +155,7 @@ When job started, every trainer needs to get it's own part of data. In some dist
 
 Different training jobs may have different data format and `reader()` function, developers may need to write different data prepare scripts and `reader()` functions for their job.
 
-## Prepare Training program
+### Prepare Training program
 
 We'll create a *workspace* directory on each node, storing your training program, dependencies, mounted or downloaded dataset directory.
 
@@ -191,7 +191,7 @@ Your workspace may looks like:
 - `train_data_dir`: containing training data. Mount from storage service or copy trainning data to here.
 - `test_data_dir`: containing testing data.
 
-# Use cluster platforms or cluster management tools
+## Use cluster platforms or cluster management tools
 
 PaddlePaddle supports running jobs on several platforms including:
 - [Kubernetes](http://kubernetes.io) open-source system for automating deployment, scaling, and management of containerized applications from Google.
@@ -202,13 +202,13 @@ We'll introduce cluster job management on these platforms. The examples can be f
 
 These cluster platforms provide API or environment variables for training processes, when the job is dispatched to different nodes. Like node ID, IP or total number of nodes etc.
 
-## Cluster Training Using Fabric
+### Cluster Training Using Fabric
 
-### Prepare a Linux cluster
+#### Prepare a Linux cluster
 
 Run `kubectl -f ssh_servers.yaml` under the directory:  `paddle/scripts/cluster_train_v2/fabric/docker_cluster` will launch a demo cluster. Run `kubectl get po -o wide` to get IP addresses of these nodes.
 
-### Launching Cluster Job
+#### Launching Cluster Job
 `paddle.py` provides automatical scripts to start all PaddlePaddle cluster processes in different nodes. By default, all command line options can be set as `paddle.py` command options and `paddle.py` will transparently and automatically set these options to PaddlePaddle lower level processes.
 
 `paddle.py`provides two distinguished command option for easy job launching.
@@ -224,10 +224,10 @@ sh run.sh
 
 The cluster Job will start in several seconds.
 
-### Kill Cluster Job
+#### Kill Cluster Job
 `paddle.py` can capture `Ctrl + C` SIGINT signal to automatically kill all processes launched by it. So just stop `paddle.py` to kill cluster job. You should manually kill the job if the program crashed.
 
-### Check Cluster Training Result
+#### Check Cluster Training Result
 Check log in $workspace/log for details, each node owns same log structure.
 
 `paddle_trainer.INFO`
@@ -242,13 +242,13 @@ It provides stderr and stdout of parameter server process. Check error log if tr
 `train.log`
 It provides stderr and stdout of trainer process. Check error log if training crashes.
 
-### Check Model Output
+#### Check Model Output
 After one pass finished, model files will be written in `output` directory in node 0.
 `nodefile` in workspace indicates the node id of current cluster job.
 
-## Cluster Training Using OpenMPI
+### Cluster Training Using OpenMPI
 
-### Prepare an OpenMPI cluster
+#### Prepare an OpenMPI cluster
 
 Run the following command to start a 3-node MPI cluster and one "head" node.
 
@@ -260,7 +260,7 @@ kubectl create -f mpi-nodes.yaml
 
 Then you can log in to every OpenMPI node using ssh without input any passwords.
 
-### Launching Cluster Job
+#### Launching Cluster Job
 
 Follow the steps to launch a PaddlePaddle training job in OpenMPI cluster:\
 
@@ -288,6 +288,6 @@ scp train.txt-00002 test.txt-00002 [node3IP]:/home/tutorial
 mpirun -hostfile machines -n 3  /home/tutorial/start_mpi_train.sh
 ```
 
-## Cluster Training Using Kubernetes
+### Cluster Training Using Kubernetes
 
 The details can be found [here](../k8s/k8s_cn.md)

paddle/cuda/include/hl_matrix.h

@@ -300,4 +300,12 @@ extern void hl_matrix_col2Vol(real* dataDst,
                               real alpha,
                               real beta);
 
+/**
+ * @brief  Matrix col2Vol: Convert col matrix into 3D volume
+ * @param[out]  out     output int vector.
+ * @param[in]   vec     input float vector.
+ * @param[in]   size    size of the vector.
+ */
+extern void hl_vector_cast2int(int* out, real* vec, int size);
+
 #endif /* HL_MATRIX_H_ */

paddle/cuda/include/stub/hl_matrix_stub.h

@@ -133,4 +133,6 @@ inline void hl_matrix_col2Vol(real* dataDst,
                               real alpha,
                               real beta) {}
 
+inline void hl_vector_cast2int(int* out, real* vec, int size) {}
+
 #endif  // HL_MATRIX_STUB_H_

paddle/cuda/src/hl_cuda_matrix.cu

@@ -793,3 +793,14 @@ void hl_matrix_col2Vol(real* dataDst,
 
   CHECK_SYNC("hl_matrix_col2Vol failed");
 }
+
+__global__ void keVectorCast2Int(int* out, real* vec, int size) {
+  for (int i = threadIdx.x; i < (size); i += blockDim.x) {
+    out[i] = int(vec[i]);
+  }
+}
+
+void hl_vector_cast2int(int* out, real* vec, int size) {
+  keVectorCast2Int<<<1, 512, 0, STREAM_DEFAULT>>>(out, vec, size);
+  CHECK_SYNC("hl_vector_cast2int failed");
+}

paddle/framework/CMakeLists.txt

@@ -20,7 +20,8 @@ cc_test(scope_test SRCS scope_test.cc DEPS scope)
 
 
 cc_library(attribute SRCS attribute.cc DEPS framework_proto)
-cc_test(program_desc_test SRCS program_desc_test.cc DEPS proto_desc)
+cc_test(program_desc_test SRCS program_desc_test.cc DEPS proto_desc
+device_context)
 cc_library(op_proto_maker SRCS op_proto_maker.cc DEPS framework_proto attribute)
 cc_test(op_proto_maker_test SRCS op_proto_maker_test.cc DEPS op_proto_maker)
 cc_library(op_info SRCS op_info.cc DEPS attribute framework_proto)

paddle/framework/attribute.cc

@@ -19,7 +19,7 @@ limitations under the License. */
 namespace paddle {
 namespace framework {
 
-Attribute GetAttrValue(const OpDesc::Attr& attr_desc, ProgramDesc* program) {
+Attribute GetAttrValue(const OpDesc::Attr& attr_desc) {
   switch (attr_desc.type()) {
     case framework::AttrType::BOOLEAN: {
       return attr_desc.b();
@@ -61,13 +61,9 @@ Attribute GetAttrValue(const OpDesc::Attr& attr_desc, ProgramDesc* program) {
       }
       return val;
     }
-    case framework::AttrType::BLOCK: {
-      PADDLE_ENFORCE(program != nullptr,
-                     "Need to specify ProgramDesc when get a block attr");
-      return program->mutable_blocks(attr_desc.block_idx());
-    }
+    default:
+      PADDLE_THROW("Unsupport attr type %d", attr_desc.type());
   }
-  PADDLE_ENFORCE(false, "Unknown OpDesc::AttrDesc::type !");
   return boost::blank();
 }
 

paddle/framework/attribute.h

@@ -32,7 +32,7 @@ inline AttrType AttrTypeID() {
   return static_cast<AttrType>(tmp.which() - 1);
 }
 
-Attribute GetAttrValue(const OpDesc::Attr& attr_desc, ProgramDesc* desc);
+Attribute GetAttrValue(const OpDesc::Attr& attr_desc);
 
 class AttrReader {
  public:

paddle/framework/backward.cc

@@ -18,12 +18,12 @@
 #include <deque>
 #include <list>
 #include <memory>
+#include <unordered_set>
 
 #include "paddle/framework/block_desc.h"
 #include "paddle/framework/op_registry.h"
 #include "paddle/operators/dynamic_recurrent_op.h"
 #include "paddle/operators/net_op.h"
-#include "paddle/operators/recurrent_op.h"
 
 namespace paddle {
 namespace framework {
@@ -37,7 +37,7 @@ static inline std::unique_ptr<OperatorBase> CreateGradOp(
   op_desc.SetType(op.Type());
   op_desc.SetAttrMap(op.Attrs());
   auto& info = OpInfoMap::Instance().Get(op.Type());
-  auto grad_descs = info.GradOpMaker()(op_desc, no_grad_set, grad_to_var);
+  auto grad_descs = info.GradOpMaker()(op_desc, no_grad_set, grad_to_var, {});
   std::vector<std::unique_ptr<OperatorBase>> grad_ops;
   grad_ops.reserve(grad_descs.size());
   std::transform(grad_descs.begin(), grad_descs.end(),
@@ -219,19 +219,7 @@ static std::unique_ptr<OperatorBase> BackwardRecursive(
                    });
 
     // process recurrent gradient op as a special operator.
-    if (forwardOp.Type() == "recurrent") {
-      // NOTE clean up cycle call somewhere (RNN's stepnet constains itself),
-      // or this will result in infinite loop.
-      const auto& rnnop =
-          *static_cast<const operators::RecurrentOp*>(&forwardOp);
-      auto rnn_grad_op =
-          static_cast<operators::RecurrentGradientOp*>(grad_op.get());
-      const auto& stepnet_op =
-          *static_cast<const OperatorBase*>(&rnnop.stepnet());
-      // create stepnet's gradient op
-      rnn_grad_op->set_stepnet(
-          BackwardRecursive(stepnet_op, no_grad_names, grad_to_var, uniq_id));
-    } else if (forwardOp.Type() == "dynamic_recurrent") {
+    if (forwardOp.Type() == "dynamic_recurrent") {
       // NOTE clean up cycle call somewhere (RNN's stepnet constains itself),
       // or this will result in infinite loop.
       const auto& rnnop =
@@ -285,6 +273,15 @@ static bool AllGradInSet(const std::vector<std::string>& names,
   return true;
 }
 
+static std::string FwdName(const std::string& grad_name) {
+  auto pos = grad_name.find("@GRAD");
+  if (pos == std::string::npos) {
+    return "";
+  } else {
+    return grad_name.substr(0, pos);
+  }
+}
+
 static void CreateGradVarInBlock(
     size_t grad_op_start_index,
     const std::unordered_map<std::string, std::string>& param_name_map,
@@ -294,6 +291,7 @@ static void CreateGradVarInBlock(
   for (size_t op_index = grad_op_start_index; op_index < ops.size();
        ++op_index) {
     bool need_infer_shape = false;
+    std::unordered_set<std::string> new_vars;
     ForEachVarName(ops[op_index]->Outputs(),
                    [&](const std::string& grad_var_name) {
                      if (block_desc->HasVar(grad_var_name)) {
@@ -301,8 +299,7 @@ static void CreateGradVarInBlock(
                      }
                      need_infer_shape = true;
                      auto var = block_desc->Var(grad_var_name);
-                     // FIXME(qiao) infer the datatype
-                     var->SetDataType(framework::DataType::FP32);
+                     new_vars.insert(var->Name());
                      auto it = param_name_map.find(grad_var_name);
                      if (it == param_name_map.end()) {
                        return false;
@@ -316,6 +313,21 @@ static void CreateGradVarInBlock(
                    });
     if (need_infer_shape) {
       ops[op_index]->InferVarType(block_desc);
+      for (auto& arg : ops[op_index]->OutputArgumentNames()) {
+        if (new_vars.find(arg) == new_vars.end()) {
+          continue;
+        }
+        auto pname = FwdName(arg);
+        auto* param = block_desc->FindVarRecursive(pname);
+        auto* grad = block_desc->FindVar(arg);
+        if (param == nullptr) {
+          LOG(WARNING) << "Cannot find forward variable of " << arg
+                       << ". Set its gradient to FP32";
+          grad->SetDataType(DataType::FP32);
+        } else {
+          grad->SetDataType(param->GetDataType());
+        }
+      }
       ops[op_index]->InferShape(*block_desc);
     }
   }
@@ -323,7 +335,9 @@ static void CreateGradVarInBlock(
 
 std::vector<std::unique_ptr<OpDescBind>> MakeOpGrad(
     const OpDescBind* op_desc, std::unordered_set<std::string>* no_grad_vars,
-    std::unordered_map<std::string, std::string>* grad_to_var) {
+    std::unordered_map<std::string, std::string>* grad_to_var,
+    const std::vector<BlockDescBind*>& grad_block =
+        std::vector<BlockDescBind*>()) {
   std::vector<std::unique_ptr<OpDescBind>> grad_op_descs;
   // All input gradients of forwarding operator do not need to calculate.
   const std::vector<std::string>& inputs = op_desc->InputArgumentNames();
@@ -339,9 +353,10 @@ std::vector<std::unique_ptr<OpDescBind>> MakeOpGrad(
     return grad_op_descs;  // empty vector
   }
 
-  grad_op_descs = OpInfoMap::Instance()
-                      .Get(op_desc->Type())
-                      .GradOpMaker()(*op_desc, *no_grad_vars, grad_to_var);
+  grad_op_descs =
+      OpInfoMap::Instance()
+          .Get(op_desc->Type())
+          .GradOpMaker()(*op_desc, *no_grad_vars, grad_to_var, grad_block);
 
   std::list<std::unique_ptr<OpDescBind>> pending_fill_zeros_ops;
   for (auto& desc : grad_op_descs) {
@@ -368,28 +383,27 @@ std::vector<std::unique_ptr<OpDescBind>> MakeBlockBackward(
     ProgramDescBind& program_desc, int block_idx,
     std::unordered_set<std::string>* no_grad_vars,
     std::unordered_map<std::string, std::string>* grad_to_var) {
-  BlockDescBind* cur_block = program_desc.Block(block_idx);
+  BlockDescBind* cur_block = program_desc.MutableBlock(block_idx);
   std::vector<OpDescBind*> op_descs = cur_block->AllOps();
   std::unordered_map<std::string, std::vector<size_t>> dup_out_ops;
   size_t grad_desc_idx = 0;
   std::vector<std::unique_ptr<OpDescBind>> backward_descs;
 
   for (auto it = op_descs.rbegin(); it != op_descs.rend(); ++it) {
-    std::vector<std::unique_ptr<OpDescBind>> op_grads =
-        MakeOpGrad(*it, no_grad_vars, grad_to_var);
+    std::vector<std::unique_ptr<OpDescBind>> op_grads;
 
     if ((*it)->Type() == "recurrent") {
-      PADDLE_ENFORCE_EQ(
-          op_grads.size(), static_cast<size_t>(1),
-          "rnn_op's gradient process should contain only one op.");
       int step_block_idx = (*it)->GetBlockAttr("step_block");
       auto backward_block_op_descs = MakeBlockBackward(
           program_desc, step_block_idx, no_grad_vars, grad_to_var);
-      BlockDescBind* backward_block = program_desc.AppendBlock(*cur_block);
+      BlockDescBind* backward_block =
+          program_desc.AppendBlock(*program_desc.MutableBlock(step_block_idx));
       for (auto& ptr : backward_block_op_descs) {
         backward_block->AppendAllocatedOp(std::move(ptr));
       }
-      op_grads[0]->SetBlockAttr("step_block", *backward_block);
+      op_grads = MakeOpGrad(*it, no_grad_vars, grad_to_var, {backward_block});
+    } else {
+      op_grads = MakeOpGrad(*it, no_grad_vars, grad_to_var);
     }
 
     for (const auto& desc : op_grads) {
@@ -443,7 +457,7 @@ ParamGradInfoMap AppendBackward(
   }
 
   const int root_block_idx = 0;
-  auto root_block = program_desc.Block(root_block_idx);
+  auto root_block = program_desc.MutableBlock(root_block_idx);
 
   // insert fill one op for target
   // TODO(qiao) add some check to the target.
@@ -492,7 +506,7 @@ ParamGradInfoMap AppendBackward(
   CreateGradVarInBlock(forward_op_num, grad_to_var, root_block, &retv);
   for (size_t block_index = forward_block_num;
        block_index < program_desc.Size(); ++block_index) {
-    CreateGradVarInBlock(0, grad_to_var, program_desc.Block(block_index),
+    CreateGradVarInBlock(0, grad_to_var, program_desc.MutableBlock(block_index),
                          &retv);
   }
   return retv;

paddle/framework/backward_test.cc

@@ -499,7 +499,7 @@ TEST(Backward, linear_net_intermediate_variable_has_no_grad) {
 
 TEST(Backward, simple_single_op) {
   f::ProgramDescBind program;
-  f::BlockDescBind *block = program.Block(0);
+  f::BlockDescBind *block = program.MutableBlock(0);
 
   f::OpDescBind *op = block->AppendOp();
   op->SetType("rowwise_add");
@@ -535,7 +535,7 @@ TEST(Backward, simple_single_op) {
 
 TEST(Backward, default_attribute) {
   f::ProgramDescBind program;
-  f::BlockDescBind *block = program.Block(0);
+  f::BlockDescBind *block = program.MutableBlock(0);
   f::OpDescBind *op = block->AppendOp();
   op->SetType("mul");
   op->SetInput("X", {"x"});
@@ -561,7 +561,7 @@ TEST(Backward, default_attribute) {
 
 TEST(Backward, simple_mult_op) {
   f::ProgramDescBind program;
-  f::BlockDescBind *block = program.Block(0);
+  f::BlockDescBind *block = program.MutableBlock(0);
   f::OpDescBind *op1 = block->AppendOp();
   op1->SetType("rowwise_add");
   op1->SetInput("X", {"x1"});
@@ -644,7 +644,7 @@ TEST(Backward, simple_mult_op) {
 
 TEST(Backward, intermedia_var_no_grad) {
   f::ProgramDescBind program;
-  f::BlockDescBind *block = program.Block(0);
+  f::BlockDescBind *block = program.MutableBlock(0);
   f::OpDescBind *op1 = block->AppendOp();
   op1->SetType("rowwise_add");
   op1->SetInput("X", {"x1"});
@@ -714,7 +714,7 @@ TEST(Backward, intermedia_var_no_grad) {
 
 TEST(Backward, var_no_grad) {
   f::ProgramDescBind program;
-  f::BlockDescBind *block = program.Block(0);
+  f::BlockDescBind *block = program.MutableBlock(0);
   f::OpDescBind *op1 = block->AppendOp();
   op1->SetType("mult_in_out");
   op1->SetInput("X", {"x1"});
@@ -790,7 +790,7 @@ TEST(Backward, var_no_grad) {
 
 TEST(Backward, shared_var) {
   f::ProgramDescBind program;
-  f::BlockDescBind *block = program.Block(0);
+  f::BlockDescBind *block = program.MutableBlock(0);
   f::OpDescBind *op1 = block->AppendOp();
   op1->SetType("rowwise_add");
   op1->SetInput("X", {"x1"});
@@ -880,7 +880,7 @@ TEST(Backward, shared_var) {
 
 TEST(Backward, half_backward) {
   f::ProgramDescBind program;
-  f::BlockDescBind *block = program.Block(0);
+  f::BlockDescBind *block = program.MutableBlock(0);
   auto *op1 = block->AppendOp();
   op1->SetType("minus");
   op1->SetInput("X", {"a"});

paddle/framework/block_desc.cc

@@ -113,7 +113,7 @@ BlockDescBind *BlockDescBind::ParentBlock() const {
   if (this->desc_->parent_idx() == kNoneBlockIndex) {
     return nullptr;
   }
-  return prog_->Block(static_cast<size_t>(this->desc_->parent_idx()));
+  return prog_->MutableBlock(static_cast<size_t>(this->desc_->parent_idx()));
 }
 
 BlockDesc *BlockDescBind::Proto() {

paddle/framework/block_desc.h

@@ -88,6 +88,8 @@ class BlockDescBind {
 
   BlockDesc *Proto();
 
+  ProgramDescBind *Program() { return this->prog_; }
+
  private:
   void ClearPBOps();
   void ClearPBVars();

paddle/framework/details/op_registry.h

@@ -108,8 +108,9 @@ struct OpInfoFiller<T, kGradOpDescMaker> {
     info->grad_op_maker_ = [](
         const OpDescBind& fwd_op,
         const std::unordered_set<std::string>& no_grad_set,
-        std::unordered_map<std::string, std::string>* grad_to_var) {
-      T maker(fwd_op, no_grad_set, grad_to_var);
+        std::unordered_map<std::string, std::string>* grad_to_var,
+        const std::vector<BlockDescBind*>& grad_block) {
+      T maker(fwd_op, no_grad_set, grad_to_var, grad_block);
       return maker();
     };
   }

paddle/framework/executor.cc

@@ -31,7 +31,7 @@ namespace framework {
 const std::string kFeedOpType = "feed";
 const std::string kFetchOpType = "fetch";
 
-Executor::Executor(const std::vector<platform::Place>& places) {
+Executor::Executor(const std::vector<platform::Place>& places) : own_(true) {
   PADDLE_ENFORCE_GT(places.size(), 0);
   device_contexts_.resize(places.size());
   for (size_t i = 0; i < places.size(); i++) {
@@ -52,8 +52,10 @@ Executor::Executor(const std::vector<platform::Place>& places) {
 }
 
 Executor::~Executor() {
-  for (auto& device_context : device_contexts_) {
-    delete device_context;
+  if (own_) {
+    for (auto& device_context : device_contexts_) {
+      delete device_context;
+    }
   }
 }
 
@@ -66,45 +68,61 @@ static void CreateTensor(Variable* var, VarDesc::VarType var_type) {
     var->GetMutable<FeedFetchList>();
   } else if (var_type == VarDesc::FETCH_LIST) {
     var->GetMutable<FeedFetchList>();
+  } else if (var_type == VarDesc::STEP_SCOPES) {
+    var->GetMutable<std::vector<framework::Scope>>();
   } else {
     PADDLE_THROW(
-        "Variable type must be "
-        "LoDTensor/SelectedRows/FEED_MINIBATCH/FETCH_LIST.");
+        "Variable type %d is not in "
+        "[LoDTensor, SelectedRows, FEED_MINIBATCH, FETCH_LIST]",
+        var_type);
   }
 }
 
-void Executor::Run(const ProgramDesc& pdesc, Scope* scope, int block_id) {
+void Executor::Run(const ProgramDescBind& pdesc, Scope* scope, int block_id,
+                   bool create_local_scope) {
   // TODO(tonyyang-svail):
   //    - only runs on the first device (i.e. no interdevice communication)
   //    - will change to use multiple blocks for RNN op and Cond Op
-  PADDLE_ENFORCE_GT(pdesc.blocks_size(), block_id);
-  auto& block = pdesc.blocks(block_id);
+  PADDLE_ENFORCE_LT(static_cast<size_t>(block_id), pdesc.Size());
+  auto& block = pdesc.Block(block_id);
   auto& device = device_contexts_[0];
 
-  Scope& local_scope = scope->NewScope();
-
-  for (auto& var : block.vars()) {
-    if (var.persistable()) {
-      auto* ptr = scope->Var(var.name());
-      CreateTensor(ptr, var.type());
-      VLOG(3) << "Create Variable " << var.name()
-              << " global, which pointer is " << ptr;
-    } else {
-      auto* ptr = local_scope.Var(var.name());
-      CreateTensor(ptr, var.type());
-      VLOG(3) << "Create Variable " << var.name()
-              << " locally, which pointer is " << ptr;
+  Scope* local_scope = scope;
+  if (create_local_scope) {
+    local_scope = &scope->NewScope();
+    for (auto& var : block.AllVars()) {
+      if (var->Persistable()) {
+        auto* ptr = scope->Var(var->Name());
+        CreateTensor(ptr, var->GetType());
+        VLOG(3) << "Create Variable " << var->Name()
+                << " global, which pointer is " << ptr;
+      } else {
+        auto* ptr = local_scope->Var(var->Name());
+        CreateTensor(ptr, var->GetType());
+        VLOG(3) << "Create Variable " << var->Name()
+                << " locally, which pointer is " << ptr;
+      }
+    }
+  } else {
+    for (auto& var : block.AllVars()) {
+      auto* ptr = local_scope->Var(var->Name());
+      CreateTensor(ptr, var->GetType());
+      VLOG(3) << "Create variable " << var->Name() << ", which pointer is "
+              << ptr;
     }
   }
 
-  for (auto& op_desc : block.ops()) {
-    auto op = paddle::framework::OpRegistry::CreateOp(
-        op_desc, const_cast<ProgramDesc*>(&pdesc));
-    op->Run(local_scope, *device);
+  for (auto& op_desc : block.AllOps()) {
+    auto op = paddle::framework::OpRegistry::CreateOp(*op_desc);
+    op->Run(*local_scope, *device);
+  }
+  if (create_local_scope) {
+    scope->DeleteScope(local_scope);
   }
-
-  scope->DeleteScope(&local_scope);
 }
 
+Executor::Executor(const platform::DeviceContext& device)
+    : device_contexts_({&device}), own_(false) {}
+
 }  // namespace framework
 }  // namespace paddle

paddle/framework/executor.h

@@ -14,8 +14,8 @@ limitations under the License. */
 
 #pragma once
 
-#include "paddle/framework/framework.pb.h"
 #include "paddle/framework/op_info.h"
+#include "paddle/framework/program_desc.h"
 #include "paddle/framework/scope.h"
 #include "paddle/framework/tensor.h"
 
@@ -25,6 +25,7 @@ namespace framework {
 class Executor {
  public:
   explicit Executor(const std::vector<platform::Place>& places);
+  explicit Executor(const platform::DeviceContext& devices);
   ~Executor();
 
   /* @Brief
@@ -34,10 +35,11 @@ class Executor {
    *  ProgramDesc
    *  Scope
    */
-  void Run(const ProgramDesc&, Scope*, int);
+  void Run(const ProgramDescBind&, Scope*, int, bool create_local_scope = true);
 
  private:
-  std::vector<platform::DeviceContext*> device_contexts_;
+  std::vector<const platform::DeviceContext*> device_contexts_;
+  bool own_;
 };
 
 }  // namespace framework

paddle/framework/grad_op_desc_maker.h

@@ -15,6 +15,7 @@
 #pragma once
 #include <string>
 #include <unordered_set>
+#include <vector>
 #include "paddle/framework/op_desc.h"
 #include "paddle/framework/operator.h"
 
@@ -26,8 +27,13 @@ class GradOpDescMakerBase {
   explicit GradOpDescMakerBase(
       const OpDescBind& fwd_op,
       const std::unordered_set<std::string>& no_grad_set,
-      std::unordered_map<std::string, std::string>* grad_to_var)
-      : fwd_op_(fwd_op), no_grad_set_(no_grad_set), grad_to_var_(grad_to_var) {}
+      std::unordered_map<std::string, std::string>* grad_to_var,
+      const std::vector<BlockDescBind*>& grad_block =
+          std::vector<BlockDescBind*>())
+      : fwd_op_(fwd_op),
+        no_grad_set_(no_grad_set),
+        grad_to_var_(grad_to_var),
+        grad_block_(grad_block) {}
 
   virtual ~GradOpDescMakerBase() = default;
   virtual std::vector<std::unique_ptr<OpDescBind>> operator()() const = 0;
@@ -102,6 +108,9 @@ class GradOpDescMakerBase {
   const OpDescBind& fwd_op_;
   const std::unordered_set<std::string>& no_grad_set_;
   std::unordered_map<std::string, std::string>* grad_to_var_;
+
+ protected:
+  std::vector<BlockDescBind*> grad_block_;
 };
 
 class SingleGradOpDescMaker : public GradOpDescMakerBase {

paddle/framework/lod_tensor_test.cu

@@ -36,8 +36,8 @@ TEST(LoDTensor, LoDInGPU) {
   lod_tensor.mutable_data<float>(place);
 
   lod_tensor.set_lod(src_lod);
-  CHECK_EQ(lod_tensor.lod_element(0, 2).first, 4UL);
-  CHECK_EQ(lod_tensor.lod_element(0, 4).first, 8UL);
+  EXPECT_EQ(lod_tensor.lod_element(0, 2).first, 4UL);
+  EXPECT_EQ(lod_tensor.lod_element(0, 4).first, 8UL);
 
   auto lod = lod_tensor.lod();
 
@@ -45,6 +45,6 @@ TEST(LoDTensor, LoDInGPU) {
   cudaDeviceSynchronize();
 
   for (size_t i = 0; i < src_lod[0].size(); ++i) {
-    CHECK_EQ(lod[0].data()[i], src_lod[0].data()[i] * 2);
+    EXPECT_EQ(lod[0].data()[i], src_lod[0].data()[i] * 2);
   }
-}
+}
\ No newline at end of file

paddle/framework/op_desc.cc

@@ -52,6 +52,22 @@ class CompileTimeInferShapeContext : public InferShapeContext {
   const std::vector<std::string> &Outputs(
       const std::string &name) const override;
 
+  void ShareLoD(const std::string &in, const std::string &out, size_t i = 0,
+                size_t j = 0) const override {
+    PADDLE_ENFORCE_LT(i, Inputs(in).size());
+    PADDLE_ENFORCE_LT(j, Outputs(out).size());
+    auto *in_var = block_.FindVarRecursive(Inputs(in)[i]);
+    auto *out_var = block_.FindVarRecursive(Outputs(out)[j]);
+    if (in_var->GetType() != VarDesc::LOD_TENSOR) {
+      VLOG(3) << "input " << in << "is not LodTensor";
+      return;
+    }
+    PADDLE_ENFORCE_EQ(in_var->GetType(), VarDesc::LOD_TENSOR,
+                      "The %d-th output of Output(%s) must be LoDTensor.", j,
+                      out);
+    in_var->SetLoDLevel(out_var->GetLodLevel());
+  }
+
  private:
   DDim GetDim(const std::string &name) const override;
 
@@ -98,7 +114,12 @@ OpDescBind::OpDescBind(const OpDesc &desc, ProgramDescBind *prog)
   // restore attrs_
   for (const OpDesc::Attr &attr : desc_.attrs()) {
     std::string attr_name = attr.name();
-    attrs_[attr_name] = GetAttrValue(attr, prog->Proto());
+    if (attr.type() != AttrType::BLOCK) {
+      attrs_[attr_name] = GetAttrValue(attr);
+    } else {
+      auto bid = attr.block_idx();
+      attrs_[attr_name] = prog->MutableBlock(bid);
+    }
   }
 }
 
@@ -172,8 +193,7 @@ void OpDescBind::SetAttr(const std::string &name, const Attribute &v) {
 }
 
 void OpDescBind::SetBlockAttr(const std::string &name, BlockDescBind &block) {
-  BlockDesc *desc = block.Proto();
-  this->attrs_[name] = desc;
+  this->attrs_[name] = &block;
   need_update_ = true;
 }
 
@@ -192,7 +212,7 @@ Attribute OpDescBind::GetAttr(const std::string &name) const {
 int OpDescBind::GetBlockAttr(const std::string &name) const {
   auto it = attrs_.find(name);
   PADDLE_ENFORCE(it != attrs_.end(), "Attribute %s is not found", name);
-  return boost::get<BlockDesc *>(it->second)->idx();
+  return boost::get<BlockDescBind *>(it->second)->ID();
 }
 
 const std::unordered_map<std::string, Attribute> &OpDescBind::GetAttrMap()
@@ -307,6 +327,19 @@ void OpDescBind::InferShape(const BlockDescBind &block) const {
   PADDLE_ENFORCE(static_cast<bool>(infer_shape),
                  "%s's infer_shape has not been registered", this->Type());
   CompileTimeInferShapeContext ctx(*this, block);
+  if (VLOG_IS_ON(10)) {
+    std::ostringstream sout;
+    auto inames = this->InputArgumentNames();
+    sout << " From [";
+    std::copy(inames.begin(), inames.end(),
+              std::ostream_iterator<std::string>(sout, ", "));
+    sout << "] to [";
+    auto onames = this->OutputArgumentNames();
+    std::copy(onames.begin(), onames.end(),
+              std::ostream_iterator<std::string>(sout, ", "));
+    sout << "]";
+    VLOG(10) << sout.str();
+  }
   infer_shape(&ctx);
 }
 

paddle/framework/op_registry.cc

@@ -43,13 +43,15 @@ static VariableNameMap ConvertOpDescVarsToVarNameMap(
   return ret_val;
 }
 
-std::unique_ptr<OperatorBase> OpRegistry::CreateOp(const OpDesc& op_desc,
-                                                   ProgramDesc* program) {
+std::unique_ptr<OperatorBase> OpRegistry::CreateOp(const OpDesc& op_desc) {
+  VLOG(1) << "CreateOp directly from OpDesc is deprecated. It should only be"
+             "used in unit tests. Use CreateOp(const OpDescBind& op_desc) "
+             "instead.";
   VariableNameMap inputs = ConvertOpDescVarsToVarNameMap(op_desc.inputs());
   VariableNameMap outputs = ConvertOpDescVarsToVarNameMap(op_desc.outputs());
   AttributeMap attrs;
   for (auto& attr : op_desc.attrs()) {
-    attrs[attr.name()] = GetAttrValue(attr, program);
+    attrs[attr.name()] = GetAttrValue(attr);
   }
 
   return CreateOp(op_desc.type(), inputs, outputs, attrs);

paddle/framework/op_registry.h

@@ -77,8 +77,7 @@ class OpRegistry {
                                                 const VariableNameMap& outputs,
                                                 AttributeMap attrs);
 
-  static std::unique_ptr<OperatorBase> CreateOp(const OpDesc& op_desc,
-                                                ProgramDesc* program);
+  static std::unique_ptr<OperatorBase> CreateOp(const OpDesc& op_desc);
 
   static std::unique_ptr<OperatorBase> CreateOp(const OpDescBind& op_desc);
 };

paddle/framework/op_registry_test.cc

@@ -74,7 +74,7 @@ TEST(OpRegistry, CreateOp) {
   attr->set_type(paddle::framework::AttrType::FLOAT);
   attr->set_f(scale);
 
-  auto op = paddle::framework::OpRegistry::CreateOp(op_desc, nullptr);
+  auto op = paddle::framework::OpRegistry::CreateOp(op_desc);
   paddle::framework::Scope scope;
   paddle::platform::CPUDeviceContext dev_ctx;
   op->Run(scope, dev_ctx);
@@ -95,7 +95,7 @@ TEST(OpRegistry, IllegalAttr) {
 
   bool caught = false;
   try {
-    paddle::framework::OpRegistry::CreateOp(op_desc, nullptr);
+    paddle::framework::OpRegistry::CreateOp(op_desc);
   } catch (paddle::platform::EnforceNotMet err) {
     caught = true;
     std::string msg = "larger_than check fail";
@@ -115,7 +115,7 @@ TEST(OpRegistry, DefaultValue) {
 
   ASSERT_TRUE(op_desc.IsInitialized());
 
-  auto op = paddle::framework::OpRegistry::CreateOp(op_desc, nullptr);
+  auto op = paddle::framework::OpRegistry::CreateOp(op_desc);
   paddle::framework::Scope scope;
   paddle::platform::CPUDeviceContext dev_ctx;
   op->Run(scope, dev_ctx);
@@ -131,7 +131,7 @@ TEST(OpRegistry, CustomChecker) {
   // attr 'test_attr' is not set
   bool caught = false;
   try {
-    paddle::framework::OpRegistry::CreateOp(op_desc, nullptr);
+    paddle::framework::OpRegistry::CreateOp(op_desc);
   } catch (paddle::platform::EnforceNotMet err) {
     caught = true;
     std::string msg = "Attribute 'test_attr' is required!";
@@ -149,7 +149,7 @@ TEST(OpRegistry, CustomChecker) {
   attr->set_i(3);
   caught = false;
   try {
-    paddle::framework::OpRegistry::CreateOp(op_desc, nullptr);
+    paddle::framework::OpRegistry::CreateOp(op_desc);
   } catch (paddle::platform::EnforceNotMet err) {
     caught = true;
     std::string msg = "'test_attr' must be even!";
@@ -166,7 +166,7 @@ TEST(OpRegistry, CustomChecker) {
   attr->set_name("test_attr");
   attr->set_type(paddle::framework::AttrType::INT);
   attr->set_i(4);
-  auto op = paddle::framework::OpRegistry::CreateOp(op_desc, nullptr);
+  auto op = paddle::framework::OpRegistry::CreateOp(op_desc);
   paddle::platform::CPUDeviceContext dev_ctx;
   paddle::framework::Scope scope;
   op->Run(scope, dev_ctx);

paddle/framework/operator.cc

@@ -37,32 +37,32 @@ ExecutionContext::GetEigenDevice<platform::GPUPlace, Eigen::GpuDevice>() const {
 std::string OperatorBase::Input(const std::string& name) const {
   auto& ins = Inputs(name);
   PADDLE_ENFORCE_LE(ins.size(), 1UL,
-                    "Op %s input %s should contain only one variable", type_,
-                    name);
+                    "Operator %s's input %s should contain only one variable.",
+                    type_, name);
   return ins.empty() ? kEmptyVarName : ins[0];
 }
 
 const std::vector<std::string>& OperatorBase::Inputs(
     const std::string& name) const {
   auto it = inputs_.find(name);
-  PADDLE_ENFORCE(it != inputs_.end(), "Op %s do not have input %s", type_,
-                 name);
+  PADDLE_ENFORCE(it != inputs_.end(), "Operator %s does not have the input %s.",
+                 type_, name);
   return it->second;
 }
 
 std::string OperatorBase::Output(const std::string& name) const {
   auto& outs = Outputs(name);
   PADDLE_ENFORCE_LE(outs.size(), 1UL,
-                    "Op %s output %s should contain only one variable", type_,
-                    name);
+                    "Operator %s's output %s should contain only one variable.",
+                    type_, name);
   return outs.empty() ? kEmptyVarName : outs[0];
 }
 
 const std::vector<std::string>& OperatorBase::Outputs(
     const std::string& name) const {
   auto it = outputs_.find(name);
-  PADDLE_ENFORCE(it != outputs_.end(), "Op %s does not have output called %s",
-                 type_, name);
+  PADDLE_ENFORCE(it != outputs_.end(),
+                 "Operator %s does not have an output called %s.", type_, name);
   return it->second;
 }
 
@@ -126,7 +126,7 @@ OperatorBase::OperatorBase(const std::string& type,
 
 std::vector<std::string> OperatorBase::InputVars() const {
   std::vector<std::string> ret_val;
-  for (auto& o : outputs_) {
+  for (auto& o : inputs_) {
     ret_val.reserve(ret_val.size() + o.second.size());
     ret_val.insert(ret_val.end(), o.second.begin(), o.second.end());
   }
@@ -351,6 +351,20 @@ class RuntimeInferShapeContext : public InferShapeContext {
     return op_.Outputs(name);
   }
 
+  void ShareLoD(const std::string& in, const std::string& out, size_t i = 0,
+                size_t j = 0) const override {
+    PADDLE_ENFORCE_LT(i, Inputs(in).size());
+    PADDLE_ENFORCE_LT(j, Outputs(out).size());
+    Variable* in_var = scope_.FindVar(Inputs(in)[i]);
+    Variable* out_var = scope_.FindVar(Outputs(out)[j]);
+    if (!in_var->IsType<LoDTensor>()) return;
+    PADDLE_ENFORCE(out_var->IsType<LoDTensor>(),
+                   "The %d-th output of Output(%s) must be LoDTensor.", j, out);
+    auto in_tensor = in_var->Get<LoDTensor>();
+    auto* out_tensor = out_var->GetMutable<LoDTensor>();
+    out_tensor->set_lod(in_tensor.lod());
+  }
+
  private:
   DDim GetDim(const std::string& name) const override {
     Variable* var = scope_.FindVar(name);
@@ -380,7 +394,19 @@ class RuntimeInferShapeContext : public InferShapeContext {
 
 void OperatorWithKernel::Run(const Scope& scope,
                              const platform::DeviceContext& dev_ctx) const {
-  VLOG(3) << "Running operator " << this->Type();
+  if (VLOG_IS_ON(1)) {
+    auto inputs = this->InputVars();
+    auto outputs = this->OutputVars(true);
+    std::ostringstream sout;
+    sout << "Run operator " << this->Type() << " From [";
+    std::ostream_iterator<std::string> out_it(sout, ",");
+    std::copy(inputs.begin(), inputs.end(), out_it);
+    sout << "] to [";
+    std::copy(outputs.begin(), outputs.end(), out_it);
+    sout << "]";
+    VLOG(1) << sout.str();
+  }
+
   RuntimeInferShapeContext infer_shape_ctx(*this, scope);
   this->InferShape(&infer_shape_ctx);
 

paddle/framework/operator.h

@@ -427,7 +427,8 @@ class OperatorWithKernel : public OperatorBase {
             int tmp = static_cast<int>(ToDataType(t->type()));
             VLOG(3) << "Input " << ipt_name << " with data_type " << tmp;
             PADDLE_ENFORCE(tmp == data_type || data_type == -1,
-                           "DataType of Paddle Op %s must be same.", Type());
+                           "DataType of Paddle Op %s must be the same.",
+                           Type());
             data_type = tmp;
           }
         }

paddle/framework/operator_test.cc

@@ -83,7 +83,7 @@ TEST(OperatorBase, all) {
   paddle::platform::CPUDeviceContext device_context;
   paddle::framework::Scope scope;
 
-  auto op = paddle::framework::OpRegistry::CreateOp(op_desc, nullptr);
+  auto op = paddle::framework::OpRegistry::CreateOp(op_desc);
   scope.Var("OUT1");
   ASSERT_EQ(paddle::framework::op_run_num, 0);
   op->Run(scope, device_context);
@@ -208,7 +208,7 @@ TEST(OpKernel, all) {
   paddle::platform::CPUDeviceContext cpu_device_context;
   paddle::framework::Scope scope;
 
-  auto op = paddle::framework::OpRegistry::CreateOp(op_desc, nullptr);
+  auto op = paddle::framework::OpRegistry::CreateOp(op_desc);
   ASSERT_EQ(paddle::framework::cpu_kernel_run_num, 0);
   op->Run(scope, cpu_device_context);
   ASSERT_EQ(paddle::framework::cpu_kernel_run_num, 1);
@@ -244,7 +244,7 @@ TEST(OpKernel, multi_inputs) {
   scope.Var("y0")->GetMutable<LoDTensor>();
   scope.Var("y1")->GetMutable<LoDTensor>();
 
-  auto op = paddle::framework::OpRegistry::CreateOp(op_desc, nullptr);
+  auto op = paddle::framework::OpRegistry::CreateOp(op_desc);
   op->Run(scope, cpu_device_context);
 }
 

paddle/framework/program_desc.h

@@ -37,7 +37,9 @@ class ProgramDescBind {
 
   BlockDescBind *AppendBlock(const BlockDescBind &parent);
 
-  BlockDescBind *Block(size_t idx) { return blocks_[idx].get(); }
+  BlockDescBind *MutableBlock(size_t idx) { return blocks_[idx].get(); }
+
+  const BlockDescBind &Block(size_t idx) const { return *blocks_[idx]; }
 
   size_t Size() const { return blocks_.size(); }
 

paddle/framework/program_desc_test.cc

@@ -20,7 +20,7 @@ namespace paddle {
 namespace framework {
 TEST(ProgramDesc, copy_ctor) {
   ProgramDescBind program;
-  auto* global_block = program.Block(0);
+  auto* global_block = program.MutableBlock(0);
   auto* x = global_block->Var("X");
   x->SetType(VarDesc_VarType_LOD_TENSOR);
   x->SetLoDLevel(0);
@@ -44,7 +44,7 @@ TEST(ProgramDesc, copy_ctor) {
 
   ProgramDescBind program_copy(program);
 
-  auto* global_block_copy = program_copy.Block(0);
+  auto* global_block_copy = program_copy.MutableBlock(0);
   ASSERT_NE(global_block, global_block_copy);
 
   auto assert_same_var = [&](const std::string& name, VarDescBind* var_before) {
@@ -82,7 +82,7 @@ TEST(ProgramDesc, copy_ctor) {
 
 TEST(ProgramDescBind, serialize_and_deserialize) {
   ProgramDescBind program_origin;
-  auto* global_block = program_origin.Block(0);
+  auto* global_block = program_origin.MutableBlock(0);
   auto* x = global_block->Var("X");
   x->SetType(VarDesc_VarType_LOD_TENSOR);
   x->SetLoDLevel(0);
@@ -108,7 +108,7 @@ TEST(ProgramDescBind, serialize_and_deserialize) {
   program_origin.Proto()->SerializeToString(&binary_str);
 
   ProgramDescBind program_restored(binary_str);
-  auto* global_block_restored = program_restored.Block(0);
+  auto* global_block_restored = program_restored.MutableBlock(0);
   ASSERT_NE(global_block, global_block_restored);
 
   auto assert_same_var = [&](const std::string& name, VarDescBind* var_before) {

paddle/framework/prune_test.cc

@@ -52,7 +52,7 @@ void AddOp(const std::string &type, const f::VariableNameMap &inputs,
 
 TEST(Prune, one_operator) {
   f::ProgramDescBind program;
-  f::BlockDescBind *block = program.Block(0);
+  f::BlockDescBind *block = program.MutableBlock(0);
 
   AddOp("one_one", {{"input", {"a"}}}, {{"output", {"b"}}}, {}, block);
 
@@ -69,7 +69,7 @@ TEST(Prune, one_operator) {
 
 TEST(Prune, forward) {
   f::ProgramDescBind program;
-  f::BlockDescBind *block = program.Block(0);
+  f::BlockDescBind *block = program.MutableBlock(0);
 
   AddOp("one_one", {{"input", {"a"}}}, {{"output", {"b"}}}, {}, block);
   AddOp("one_one", {{"input", {"b"}}}, {{"output", {"c"}}}, {}, block);
@@ -88,7 +88,7 @@ TEST(Prune, forward) {
 
 TEST(Prune, multi_input_op) {
   f::ProgramDescBind program;
-  f::BlockDescBind *block = program.Block(0);
+  f::BlockDescBind *block = program.MutableBlock(0);
 
   AddOp("one_one", {{"input", {"a0"}}}, {{"output", {"b0"}}}, {}, block);
   AddOp("one_one", {{"input", {"a1"}}}, {{"output", {"b1"}}}, {}, block);
@@ -106,7 +106,7 @@ TEST(Prune, multi_input_op) {
 
 TEST(Prune, multi_output_op) {
   f::ProgramDescBind program;
-  f::BlockDescBind *block = program.Block(0);
+  f::BlockDescBind *block = program.MutableBlock(0);
 
   AddOp("one_two", {{"input", {"a"}}}, {{"output", {"b", "c"}}}, {}, block);
   AddOp("one_one", {{"input", {"b"}}}, {{"output", {"b1"}}}, {}, block);
@@ -122,7 +122,7 @@ TEST(Prune, multi_output_op) {
 
 TEST(Prune, multi_target) {
   f::ProgramDescBind program;
-  f::BlockDescBind *block = program.Block(0);
+  f::BlockDescBind *block = program.MutableBlock(0);
 
   AddOp("one_two", {{"input", {"a"}}}, {{"output", {"b", "c"}}}, {}, block);
   AddOp("one_one", {{"input", {"b"}}}, {{"output", {"b1"}}}, {}, block);

paddle/framework/scope.cc

@@ -47,8 +47,12 @@ Variable* Scope::Var(const std::string& name) {
   return v;
 }
 
-Variable* Scope::Var() {
-  return Var(string::Sprintf("%p.%d", this, vars_.size()));
+Variable* Scope::Var(std::string* name) {
+  auto var_name = string::Sprintf("%p.%d", this, vars_.size());
+  if (name != nullptr) {
+    *name = var_name;
+  }
+  return Var(var_name);
 }
 
 Variable* Scope::FindVar(const std::string& name) const {

paddle/framework/scope.h

@@ -49,7 +49,7 @@ class Scope {
   Variable* Var(const std::string& name);
 
   /// Create a variable with a scope-unique name.
-  Variable* Var();
+  Variable* Var(std::string* name = nullptr);
 
   /// Find a variable in the scope or any of its ancestors.  Returns
   /// nullptr if cannot find.

paddle/framework/shape_inference.cc

@@ -28,9 +28,6 @@ void InferShapeContext::SetOutputsDim(
   SetDims(names, dims);
 }
 
-void InferShapeContext::ShareLoD(const std::string &in, const std::string &out,
-                                 size_t i, size_t j) const {}
-
 std::vector<framework::DDim> InferShapeContext::GetDims(
     const std::vector<std::string> &names) const {
   std::vector<framework::DDim> ret;

paddle/framework/shape_inference.h

@@ -43,9 +43,8 @@ class InferShapeContext {
   virtual const std::vector<std::string> &Outputs(
       const std::string &name) const = 0;
 
-  // TODO(qiao) implement this function
-  void ShareLoD(const std::string &in, const std::string &out, size_t i = 0,
-                size_t j = 0) const;
+  virtual void ShareLoD(const std::string &in, const std::string &out,
+                        size_t i = 0, size_t j = 0) const = 0;
 
  protected:
   virtual framework::DDim GetDim(const std::string &name) const = 0;

paddle/framework/tensor.h

@@ -118,12 +118,14 @@ class Tensor {
                              const platform::DeviceContext& ctx);
 
   /**
-   * @brief   Return the slice of the tensor.
+   * @brief  Return a sub-tensor of the given tensor.
    *
-   * @param[in] begin_idx   The begin index of the slice.
-   * @param[in] end_idx     The end index of the slice.
+   * @param[in] begin_idx   The index of the start row(inclusive) to slice.
+   *                        The index number begins from 0.
+   * @param[in] end_idx     The index of the end row(exclusive) to slice.
+   *                        The index number begins from 0.
    */
-  inline Tensor Slice(const int& begin_idx, const int& end_idx) const;
+  inline Tensor Slice(int begin_idx, int end_idx) const;
 
   platform::Place place() const {
     PADDLE_ENFORCE_NOT_NULL(

paddle/framework/tensor_impl.h

@@ -112,9 +112,10 @@ inline void* Tensor::mutable_data(platform::Place place, std::type_index type) {
   if (holder_ != nullptr) {
     holder_->set_type(type);
   }
-  PADDLE_ENFORCE_GT(numel(), 0,
-                    "Tensor's numel must be larger than zero to call "
-                    "Tensor::mutable_data. Call Tensor::set_dim first.");
+  PADDLE_ENFORCE_GT(
+      numel(), 0,
+      "When calling this method, the Tensor's numel must be larger than zero. "
+      "Please check Tensor::Resize has been called first.");
   int64_t size = numel() * SizeOfType(type);
   /* some versions of boost::variant don't have operator!= */
   if (holder_ == nullptr || !(holder_->place() == place) ||
@@ -227,12 +228,14 @@ inline void Tensor::CopyFromVector(const std::vector<T>& src,
 #endif
 }
 
-inline Tensor Tensor::Slice(const int& begin_idx, const int& end_idx) const {
+inline Tensor Tensor::Slice(int begin_idx, int end_idx) const {
   check_memory_size();
-  PADDLE_ENFORCE_GE(begin_idx, 0, "Slice begin index is less than zero.");
-  PADDLE_ENFORCE_LE(end_idx, dims_[0], "Slice end index is out of bound.");
-  PADDLE_ENFORCE_LT(begin_idx, end_idx,
-                    "Begin index must be less than end index.");
+  PADDLE_ENFORCE_GE(begin_idx, 0,
+                    "The start row index must be greater than 0.");
+  PADDLE_ENFORCE_LE(end_idx, dims_[0], "The end row index is out of bound.");
+  PADDLE_ENFORCE_LT(
+      begin_idx, end_idx,
+      "The start row index must be lesser than the end row index.");
 
   if (dims_[0] == 1) {
     return *this;

paddle/framework/type_defs.h

@@ -29,6 +29,7 @@ class OpDescBind;
 class BlockDescBind;
 class BlockDesc;
 class InferShapeContext;
+class BlockDescBind;
 
 using VariableNameMap = std::map<std::string, std::vector<std::string>>;
 
@@ -36,7 +37,7 @@ using VariableNameMap = std::map<std::string, std::vector<std::string>>;
 using Attribute =
     boost::variant<boost::blank, int, float, std::string, std::vector<int>,
                    std::vector<float>, std::vector<std::string>, bool,
-                   std::vector<bool>, BlockDesc*>;
+                   std::vector<bool>, BlockDescBind*>;
 
 using AttributeMap = std::unordered_map<std::string, Attribute>;
 
@@ -46,7 +47,8 @@ using OpCreator = std::function<OperatorBase*(
 
 using GradOpMakerFN = std::function<std::vector<std::unique_ptr<OpDescBind>>(
     const OpDescBind&, const std::unordered_set<std::string>& /*no_grad_set*/,
-    std::unordered_map<std::string, std::string>* /*grad_to_var*/)>;
+    std::unordered_map<std::string, std::string>* /*grad_to_var*/,
+    const std::vector<BlockDescBind*>& grad_block)>;
 
 using InferVarTypeFN = std::function<void(const OpDescBind& /*op_desc*/,
                                           BlockDescBind* /*block*/)>;

paddle/framework/var_type_inference_test.cc

@@ -63,41 +63,43 @@ namespace framework {
 
 TEST(InferVarType, sum_op) {
   ProgramDescBind prog;
-  auto *op = prog.Block(0)->AppendOp();
+  auto *op = prog.MutableBlock(0)->AppendOp();
   op->SetType("sum");
   op->SetInput("X", {"test_a", "test_b", "test_c"});
   op->SetOutput("Out", {"test_out"});
 
-  prog.Block(0)->Var("test_a")->SetType(VarDesc::SELECTED_ROWS);
-  prog.Block(0)->Var("test_b")->SetType(VarDesc::SELECTED_ROWS);
-  prog.Block(0)->Var("test_c")->SetType(VarDesc::SELECTED_ROWS);
-  prog.Block(0)->Var("test_out");
+  prog.MutableBlock(0)->Var("test_a")->SetType(VarDesc::SELECTED_ROWS);
+  prog.MutableBlock(0)->Var("test_b")->SetType(VarDesc::SELECTED_ROWS);
+  prog.MutableBlock(0)->Var("test_c")->SetType(VarDesc::SELECTED_ROWS);
+  prog.MutableBlock(0)->Var("test_out");
 
-  op->InferVarType(prog.Block(0));
+  op->InferVarType(prog.MutableBlock(0));
 
-  ASSERT_EQ(VarDesc::SELECTED_ROWS, prog.Block(0)->Var("test_out")->GetType());
+  ASSERT_EQ(VarDesc::SELECTED_ROWS,
+            prog.MutableBlock(0)->Var("test_out")->GetType());
 
-  prog.Block(0)->Var("test_b")->SetType(VarDesc::LOD_TENSOR);
-  op->InferVarType(prog.Block(0));
-  ASSERT_EQ(VarDesc::LOD_TENSOR, prog.Block(0)->Var("test_out")->GetType());
+  prog.MutableBlock(0)->Var("test_b")->SetType(VarDesc::LOD_TENSOR);
+  op->InferVarType(prog.MutableBlock(0));
+  ASSERT_EQ(VarDesc::LOD_TENSOR,
+            prog.MutableBlock(0)->Var("test_out")->GetType());
 }
 
 TEST(InferVarType, sum_op_without_infer_var_type) {
   ProgramDescBind prog;
-  auto *op = prog.Block(0)->AppendOp();
+  auto *op = prog.MutableBlock(0)->AppendOp();
   op->SetType("sum_without_infer_var_type");
   op->SetInput("X", {"test2_a", "test2_b", "test2_c"});
   op->SetOutput("Out", {"test2_out"});
 
-  prog.Block(0)->Var("test2_a")->SetType(VarDesc::SELECTED_ROWS);
-  prog.Block(0)->Var("test2_b")->SetType(VarDesc::SELECTED_ROWS);
-  prog.Block(0)->Var("test2_c")->SetType(VarDesc::SELECTED_ROWS);
-  prog.Block(0)->Var("test2_out");
+  prog.MutableBlock(0)->Var("test2_a")->SetType(VarDesc::SELECTED_ROWS);
+  prog.MutableBlock(0)->Var("test2_b")->SetType(VarDesc::SELECTED_ROWS);
+  prog.MutableBlock(0)->Var("test2_c")->SetType(VarDesc::SELECTED_ROWS);
+  prog.MutableBlock(0)->Var("test2_out");
 
-  op->InferVarType(prog.Block(0));
+  op->InferVarType(prog.MutableBlock(0));
 
   ASSERT_EQ(VarDesc_VarType_LOD_TENSOR,
-            prog.Block(0)->Var("test2_out")->GetType());
+            prog.MutableBlock(0)->Var("test2_out")->GetType());
 }
 
 }  // namespace framework

paddle/gserver/evaluators/Evaluator.cpp

@@ -395,14 +395,24 @@ real AucEvaluator::evalImp(std::vector<Argument>& arguments) {
   CHECK_LE(arguments.size(), (size_t)3);
   MatrixPtr output = arguments[0].value;
   IVectorPtr label = arguments[1].ids;
+  MatrixPtr labelval = arguments[1].value;
   bool supportWeight = (3 == arguments.size()) ? true : false;
   MatrixPtr weight = supportWeight ? arguments[2].value : nullptr;
-  if (nullptr == output || nullptr == label ||
-      (supportWeight && nullptr == weight)) {
+
+  if (nullptr == output || (supportWeight && nullptr == weight)) {
     return 0;
   }
   size_t insNum = output->getHeight();
   size_t outputDim = output->getWidth();
+  // Copy label from value to a vector.
+  if (nullptr == label && nullptr != labelval) {
+    // label width is 1
+    CHECK_EQ(1U, labelval->getWidth());
+    VectorPtr vec =
+        Vector::create(labelval->getData(), insNum, output->useGpu());
+    label = vec->castToInt();
+  }
+
   CHECK_EQ(insNum, label->getSize());
   if (supportWeight) {
     CHECK_EQ(insNum, weight->getHeight());
@@ -443,6 +453,7 @@ real AucEvaluator::evalImp(std::vector<Argument>& arguments) {
   int* labelD = label->getData();
   real* weightD = supportWeight ? weight->getData() : nullptr;
   size_t pos = realColumnIdx_;
+
   for (size_t i = 0; i < insNum; ++i) {
     real value = outputD[pos];
     uint32_t binIdx = static_cast<uint32_t>(value * kBinNum_);

paddle/gserver/layers/CRFLayer.cpp

@@ -101,8 +101,10 @@ void CRFLayer::backward(const UpdateCallback& callback) {
                               : real(1.0f);
     instanceWeight *= coeff_;
 
-    MatrixPtr grad = output.grad->subRowMatrix(starts[i], starts[i + 1]);
-    grad->add(*crfs_[i].getXGrad(), real(1.0f), instanceWeight);
+    if (output.grad) {
+      MatrixPtr grad = output.grad->subRowMatrix(starts[i], starts[i + 1]);
+      grad->add(*crfs_[i].getXGrad(), real(1.0f), instanceWeight);
+    }
     if (needWGrad) {
       weight_->getWGrad()->add(
           *crfs_[i].getWGrad(), real(1.0f), instanceWeight);

paddle/gserver/layers/LinearChainCRF.cpp

@@ -102,7 +102,6 @@ real LinearChainCRF::forward(real* x, int* s, int length) {
 }
 
 void LinearChainCRF::backward(real* x, int* s, int length, bool needWGrad) {
-  MatrixPtr matX = Matrix::create(x, length, numClasses_);
   Matrix::resizeOrCreate(matGrad_, length, numClasses_);
   Matrix::resizeOrCreate(beta_, length, numClasses_);
   real* b = b_->getData();

paddle/gserver/layers/SequenceReshapeLayer.cpp

@@ -70,11 +70,23 @@ void SequenceReshapeLayer::forward(PassType passType) {
   size_t outDim = getSize();
 
   size_t numSequences = input.getNumSequences();
-  auto startPositions = input.sequenceStartPositions->getVector(false);
-  const int* starts = startPositions->getData();
 
-  CHECK_EQ(starts[numSequences], input.getBatchSize());
-  CHECK_EQ(numSequences, startPositions->getSize() - 1);
+  // by default, we assume each instance as a sequence
+  IVectorPtr seqStarts;
+  IVector::resizeOrCreate(seqStarts, input.getBatchSize() + 1, false);
+  int* startsData = seqStarts->getData();
+  for (int i = 0; i < input.getBatchSize() + 1; i++) {
+    startsData[i] = i;
+  }
+  const int* starts = startsData;
+
+  // if there is sequence, then use start positions
+  if (input.sequenceStartPositions) {
+    auto startPositions = input.sequenceStartPositions->getVector(false);
+    starts = startPositions->getData();
+    CHECK_EQ(starts[numSequences], input.getBatchSize());
+    CHECK_EQ(numSequences, startPositions->getSize() - 1);
+  }
 
   for (size_t seqID = 0; seqID < numSequences; seqID++) {
     size_t inNumIns = starts[seqID + 1] - starts[seqID];

paddle/gserver/tests/MKLDNNTester.cpp

@@ -273,31 +273,37 @@ void MKLDNNTester::printVector(const VectorPtr& v) {
   VLOG(MKLDNN_ALL) << std::endl << ostr.str();
 }
 
-double MKLDNNTester::getDelta(const real* d1,
-                              const real* d2,
+double MKLDNNTester::getDelta(const real* refer,
+                              const real* value,
                               size_t len,
                               const float failRate,
                               const float thres) {
   double delta = 0, sum = 0;
   int failCnt = 0;
   const double eps = 1e-5;
-  double maxOut = 0;
+  double maxRatio = 0;
   for (size_t i = 0; i < len; ++i) {
-    double ref = fabs(d2[i]);
-    double diff = fabs(d1[i] - d2[i]);
+    double ref = fabs(refer[i]);
+    double val = fabs(value[i]);
+    double diff = fabs(refer[i] - value[i]);
     delta += diff;
     sum += ref;
-    if (ref > eps && fabs(d1[i]) > eps && diff / ref > thres) {
-      maxOut = std::max(maxOut, diff / ref);
+    if (ref < eps && val < eps) {  // both values are very small
+      continue;
+    }
+    double ratio = diff / ref;
+    if (ratio > thres) {
+      maxRatio = std::max(maxRatio, ratio);
       failCnt++;
     }
   }
-  EXPECT_TRUE(std::isnormal(sum));
   EXPECT_FALSE(std::isinf(sum));
+  EXPECT_FALSE(std::isnan(sum));
   EXPECT_FALSE(std::isnan(delta));
   VLOG(MKLDNN_ALL) << "reference avg data: " << sum / len
                    << ", delta: " << delta / sum << ", failCnt:" << failCnt;
-  return (failCnt / (float)len) > failRate ? maxOut : delta / sum;
+  double res = sum > eps ? delta / sum : eps;
+  return (failCnt / (float)len) > failRate ? maxRatio : res;
 }
 
 double MKLDNNTester::compareMatrix(const MatrixPtr& m1, const MatrixPtr& m2) {
@@ -515,12 +521,16 @@ void MKLDNNTester::getOutResult(const std::string& configPath,
     gradientMachine->forward(in.inArgs[i], &outArgs, PASS_TRAIN);
     // save forward result
     for (size_t k = 0; k < outArgs.size(); k++) {
-      MatrixPtr value = Matrix::create(outArgs[k].value->getHeight(),
-                                       outArgs[k].value->getWidth(),
-                                       false,
-                                       false);
-      value->copyFrom(*outArgs[k].value);
-      out.outValues.push_back(value);
+      const MatrixPtr& src = outArgs[k].value;
+      MatrixPtr dst =
+          Matrix::create(src->getHeight(), src->getWidth(), false, false);
+      if (typeid(*src) == typeid(MKLDNNMatrix)) {
+        MKLDNNMatrixPtr dnnSrc = std::dynamic_pointer_cast<MKLDNNMatrix>(src);
+        dnnSrc->copyTo(*dst);
+      } else {
+        dst->copyFrom(*src);
+      }
+      out.outValues.push_back(dst);
     }
 
     // random backward input
@@ -543,19 +553,19 @@ void MKLDNNTester::getOutResult(const std::string& configPath,
 void MKLDNNTester::compareResult(DataOut& ref, DataOut& dnn, float eps) {
   CHECK_EQ(ref.outValues.size(), dnn.outValues.size());
   CHECK_EQ(ref.paraValues.size(), dnn.paraValues.size());
-  VLOG(MKLDNN_TESTS) << "compare value size: " << ref.outValues.size();
   for (size_t i = 0; i < ref.outValues.size(); i++) {
+    VLOG(MKLDNN_TESTS) << "compare value index: " << i;
     EXPECT_LE(fabs(compareMatrix(ref.outValues[i], dnn.outValues[i])), eps);
   }
-  VLOG(MKLDNN_TESTS) << "compare param size: " << ref.outValues.size();
   for (size_t i = 0; i < ref.paraValues.size(); i++) {
+    VLOG(MKLDNN_TESTS) << "compare param index: " << i;
     EXPECT_LE(fabs(compareVector(ref.paraValues[i], dnn.paraValues[i])), eps);
   }
 }
 
-void MKLDNNTester::runBranchesTest(const std::string& configPath,
-                                   size_t iter,
-                                   float eps) {
+void MKLDNNTester::runNetTest(const std::string& configPath,
+                              size_t iter,
+                              float eps) {
   DataIn in;
   initArgument(in, configPath, iter);
   DataOut outCpu, outDnn;

paddle/gserver/tests/MKLDNNTester.h

@@ -85,17 +85,17 @@ public:
            bool printDetails = false,
            size_t iter = 3,
            float epsilon = 1e-4);
-  static void runBranchesTest(const std::string& configPath,
-                              size_t iter = 3,
-                              float eps = 1e-4);
+  static void runNetTest(const std::string& configPath,
+                         size_t iter = 2,
+                         float eps = 1e-4);
   static void initArgument(DataIn& data,
                            const std::string& configPath,
-                           size_t iter = 3);
+                           size_t iter = 2);
   static void getOutResult(const std::string& configPath,
                            DataIn& in,
                            DataOut& out,
                            bool use_mkldnn,
-                           size_t iter = 3);
+                           size_t iter = 2);
 
 private:
   void reset(const TestConfig& dnn, const TestConfig& ref, size_t batchSize);
@@ -128,13 +128,13 @@ private:
 
   /**
    * Get delta percent
-   * if many(>failRate) wrong(abs(dnn-ref)/abs(ref)>thres) points return the
-   * max(diff/ref)
-   * else return sum(abs(a-b)) / sum(abs(b))
+   * if many(>failRate) wrong(abs(val-ref)/abs(ref) > thres) points
+   * return the max(diff/ref)
+   * else return sum(abs(diff)) / sum(abs(ref))
    * The return value should be smaller than eps when passing.
    */
-  static double getDelta(const real* d1,
-                         const real* d2,
+  static double getDelta(const real* refer,
+                         const real* value,
                          size_t len,
                          const float failRate = 1e-3,
                          const float thres = 0.1);

paddle/trainer/tests/sample_trainer_config_branch_net.conf → paddle/gserver/tests/mkldnn_branch_net.conf

@@ -14,36 +14,82 @@
 
 from paddle.trainer_config_helpers import *
 
-################################### Data Configuration ###################################
-TrainData(ProtoData(files = "trainer/tests/mnist.list"))
-################################### Algorithm Configuration ###################################
-settings(batch_size = 128,
-         learning_method = MomentumOptimizer(momentum=0.5, sparse=False))
-################################### Network Configuration ###################################
-data = data_layer(name ="input", size=784)
+settings(batch_size=16)
+channels = get_config_arg("channels", int, 2)
+
+def two_conv(input, group_name):
+  out1 = img_conv_layer(input=input,
+              name=group_name+'_conv1_',
+              filter_size=1,
+              num_filters=channels,
+              padding=0,
+              shared_biases=True,
+              act=ReluActivation())
+
+  out2 = img_conv_layer(input=input,
+              name=group_name+'_conv2_',
+              filter_size=3,
+              num_filters=channels,
+              padding=1,
+              shared_biases=True,
+              act=ReluActivation())
+  return out1, out2
+
+def two_conv_bn(input, group_name):
+  out1, out2 = two_conv(input, group_name)
+  out1 = batch_norm_layer(input=out1,
+              name=group_name+'_bn1_',
+              use_global_stats=False,
+              act=ReluActivation())
+
+  out2 = batch_norm_layer(input=out2,
+              name=group_name+'_bn2_',
+              use_global_stats=False,
+              act=ReluActivation())
+  return out1, out2
+
+def two_conv_pool(input, group_name):
+  out1, out2 = two_conv(input, group_name)
+  out1 = img_pool_layer(input=out1,
+              name=group_name+'_pool1_',
+              pool_size=3,
+              stride=2,
+              padding=0,
+              pool_type=MaxPooling())
+
+  out2 = img_pool_layer(input=out2,
+              name=group_name+'_pool2_',
+              pool_size=5,
+              stride=2,
+              padding=1,
+              pool_type=MaxPooling())
+  return out1, out2
+
+def two_fc(input, group_name):
+  out1 = fc_layer(input=input,
+            name=group_name+'_fc1_',
+            size=channels,
+            bias_attr=False,
+            act=LinearActivation())
 
-tmp = img_conv_layer(input=data,
-            num_channels=1,
-            filter_size=3,
-            num_filters=32,
-            padding=1,
-            shared_biases=True,
-            act=ReluActivation())
+  out2 = fc_layer(input=input,
+            name=group_name+'_fc2_',
+            size=channels,
+            bias_attr=False,
+            act=LinearActivation())
+  return out1, out2
 
-a1 = img_conv_layer(input=tmp,
-            filter_size=1,
-            num_filters=32,
-            padding=0,
-            shared_biases=True,
-            act=ReluActivation())
+data = data_layer(name ="input", size=channels*16*16)
 
-a2 = img_conv_layer(input=tmp,
+tmp = img_conv_layer(input=data,
+            num_channels=channels,
             filter_size=3,
-            num_filters=32,
+            num_filters=channels,
             padding=1,
             shared_biases=True,
             act=ReluActivation())
 
+a1, a2 = two_conv(tmp, 'conv_branch')
 tmp = addto_layer(input=[a1, a2],
             act=ReluActivation(),
             bias_attr=False)
@@ -54,36 +100,11 @@ tmp = img_pool_layer(input=tmp,
             padding=1,
             pool_type=AvgPooling())
 
-b1 = img_conv_layer(input=tmp,
-            filter_size=3,
-            num_filters=32,
-            padding=1,
-            shared_biases=True,
-            act=ReluActivation())
-
-b1 = img_pool_layer(input=b1,
-            pool_size=3,
-            stride=2,
-            padding=0,
-            pool_type=MaxPooling())
-
-b2 = img_conv_layer(input=tmp,
-            filter_size=3,
-            num_filters=64,
-            padding=1,
-            shared_biases=True,
-            act=ReluActivation())
-
-b2 = img_pool_layer(input=b2,
-            pool_size=5,
-            stride=2,
-            padding=1,
-            pool_type=MaxPooling())
-
+b1, b2 = two_conv_pool(tmp, 'pool_branch')
 tmp = concat_layer(input=[b1, b2])
 
 tmp = img_pool_layer(input=tmp,
-            num_channels=96,
+            num_channels=channels*2,
             pool_size=3,
             stride=2,
             padding=1,
@@ -91,8 +112,9 @@ tmp = img_pool_layer(input=tmp,
 
 tmp = img_conv_layer(input=tmp,
             filter_size=3,
-            num_filters=32,
+            num_filters=channels,
             padding=1,
+            stride=2,
             shared_biases=True,
             act=LinearActivation(),
             bias_attr=False)
@@ -101,33 +123,20 @@ tmp = batch_norm_layer(input=tmp,
             use_global_stats=False,
             act=ReluActivation())
 
-c1 = img_conv_layer(input=tmp,
-            filter_size=1,
-            num_filters=32,
-            padding=0,
-            shared_biases=True,
-            act=ReluActivation())
-
-c2 = img_conv_layer(input=tmp,
-            filter_size=3,
-            num_filters=32,
-            padding=1,
-            shared_biases=True,
-            act=ReluActivation())
-
+c1, c2 = two_conv_bn(tmp, 'bn_branch')
 tmp = addto_layer(input=[c1, c2],
             act=ReluActivation(),
             bias_attr=False)
 
-tmp = fc_layer(input=tmp, size=64,
-            bias_attr=False,
-            act=TanhActivation())
+tmp = fc_layer(input=tmp, size=channels,
+            bias_attr=True,
+            act=ReluActivation())
 
-output = fc_layer(input=tmp, size=10,
+d1, d2 = two_fc(tmp, 'fc_branch')
+tmp = addto_layer(input=[d1, d2])
+
+out = fc_layer(input=tmp, size=10,
             bias_attr=True,
             act=SoftmaxActivation())
 
-lbl = data_layer(name ="label", size=10)
-
-cost = classification_cost(input=output, label=lbl)
-outputs(cost)
+outputs(out)

paddle/gserver/tests/mkldnn_branches_fc.conf

-# Copyright (c) 2017 PaddlePaddle Authors. All Rights Reserved
-#
-# 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.
-
-from paddle.trainer_config_helpers import *
-
-settings(batch_size=16)
-channels = get_config_arg("channels", int, 2)
-
-def two_fc(input, group_name):
-  out1 = fc_layer(input=input,
-            name=group_name+'_fc1',
-            size=channels,
-            bias_attr=False,
-            act=LinearActivation())
-
-  out2 = fc_layer(input=input,
-            name=group_name+'_fc2',
-            size=channels,
-            bias_attr=False,
-            act=LinearActivation())
-  return out1, out2
-
-data = data_layer(name ="input", size=channels*16*16)
-
-conv = img_conv_layer(input=data,
-            num_channels=channels,
-            filter_size=3,
-            num_filters=channels,
-            padding=1,
-            shared_biases=True,
-            act=LinearActivation())
-
-pool = img_pool_layer(input=conv,
-            pool_size=3,
-            stride=2,
-            padding=1,
-            pool_type=AvgPooling())
-
-a1, a2 = two_fc(input=pool, group_name='a')
-
-concat = concat_layer(input=[a1, a2])
-
-b1, b2 = two_fc(input=pool, group_name='b')
-
-addto = addto_layer(input=[b1, b2])
-
-outputs([concat, addto])

paddle/gserver/tests/mkldnn_branches_pool.conf

-# Copyright (c) 2017 PaddlePaddle Authors. All Rights Reserved
-#
-# 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.
-
-from paddle.trainer_config_helpers import *
-
-settings(batch_size=16)
-channels = get_config_arg("channels", int, 2)
-
-def two_pool(input, group_name):
-  out1 = img_pool_layer(input=input,
-            name=group_name+'_pool1',
-            pool_size=3,
-            stride=2,
-            padding=0,
-            pool_type=MaxPooling())
-
-  out2 = img_pool_layer(input=input,
-            name=group_name+'_pool2',
-            pool_size=5,
-            stride=2,
-            padding=1,
-            pool_type=MaxPooling())
-  return out1, out2
-
-data = data_layer(name ="input", size=channels*16*16)
-
-conv = img_conv_layer(input=data,
-            num_channels=channels,
-            filter_size=3,
-            num_filters=channels,
-            padding=1,
-            shared_biases=True,
-            act=LinearActivation())
-
-pool = img_pool_layer(input=conv,
-            pool_size=3,
-            stride=1,
-            padding=1,
-            pool_type=AvgPooling())
-
-a1, a2 = two_pool(input=pool, group_name='a')
-
-concat = concat_layer(input=[a1, a2])
-
-b1, b2 = two_pool(input=pool, group_name='b')
-
-addto = addto_layer(input=[b1, b2])
-
-outputs([concat, addto])

paddle/gserver/tests/mkldnn_branches_conv.conf → paddle/gserver/tests/mkldnn_simple_net.conf

@@ -17,40 +17,48 @@ from paddle.trainer_config_helpers import *
 settings(batch_size=16)
 channels = get_config_arg("channels", int, 2)
 
-def two_conv(input, group_name):
-  out1 = img_conv_layer(input=input,
-            name=group_name+'_conv1',
-            filter_size=1,
-            num_filters=channels,
-            padding=0,
-            shared_biases=True,
-            act=ReluActivation())
+data = data_layer(name ="input", size=channels*16*16)
 
-  out2 = img_conv_layer(input=input,
-            name=group_name+'_conv2',
+tmp = img_conv_layer(input=data,
+            num_channels=channels,
             filter_size=3,
             num_filters=channels,
             padding=1,
             shared_biases=True,
             act=ReluActivation())
-  return out1, out2
 
-data = data_layer(name ="input", size=channels*16*16)
+tmp = img_pool_layer(input=tmp,
+            pool_size=3,
+            stride=1,
+            padding=0,
+            pool_type=AvgPooling())
 
-conv = img_conv_layer(input=data,
-            num_channels=channels,
+tmp = img_conv_layer(input=tmp,
             filter_size=3,
             num_filters=channels,
             padding=1,
             shared_biases=True,
-            act=ReluActivation())
+            act=LinearActivation(),
+            bias_attr=False)
 
-a1, a2 = two_conv(input=conv, group_name='a')
+tmp = batch_norm_layer(input=tmp,
+            use_global_stats=False,
+            act=ReluActivation())
 
-concat = concat_layer(input=[a1, a2])
+tmp = img_pool_layer(input=tmp,
+            pool_size=3,
+            stride=2,
+            padding=1,
+            pool_type=MaxPooling())
 
-b1, b2 = two_conv(input=conv, group_name='b')
+tmp = fc_layer(input=tmp,
+            size=channels,
+            bias_attr=False,
+            act=ReluActivation())
 
-addto = addto_layer(input=[b1, b2])
+out = fc_layer(input=tmp,
+            size=10,
+            bias_attr=True,
+            act=SoftmaxActivation())
 
-outputs([concat, addto])
+outputs(out)

paddle/gserver/tests/test_MKLDNN.cpp

@@ -234,8 +234,7 @@ static void getMKLDNNBatchNormConfig(TestConfig& cfg,
   cfg.inputDefs.push_back({INPUT_DATA, "layer_2_moving_var", 1, size_t(pm.ic)});
   cfg.inputDefs.back().isStatic = true;
   LayerInputConfig* input = cfg.layerConfig.add_inputs();
-  // TODO(TJ): uncomment me when refine and support comparing all zeroes vector
-  // cfg.layerConfig.set_active_type("relu");
+  cfg.layerConfig.set_active_type("relu");
   cfg.layerConfig.add_inputs();
   cfg.layerConfig.add_inputs();
   ImageConfig* img_conf = input->mutable_image_conf();
@@ -309,15 +308,15 @@ TEST(MKLDNNActivation, Activations) {
 }
 
 DECLARE_string(config_args);
-TEST(MKLDNNLayer, branches) {
-  std::vector<std::string> cases = {"conv", "pool", "fc"};
+TEST(MKLDNNNet, net) {
+  std::vector<std::string> cases = {"simple", "branch"};
   for (auto name : cases) {
-    std::string config = "./gserver/tests/mkldnn_branches_" + name + ".conf";
+    std::string config = "./gserver/tests/mkldnn_" + name + "_net.conf";
     for (auto channels : {2, 32}) {
       std::ostringstream oss;
       oss << "channels=" << channels;
       FLAGS_config_args = oss.str();
-      MKLDNNTester::runBranchesTest(config);
+      MKLDNNTester::runNetTest(config);
     }
   }
 }

paddle/math/MKLDNNMatrix.h

@@ -102,6 +102,11 @@ public:
     m_->copyFrom(src);
   }
 
+  void copyTo(Matrix& dst) {
+    // TODO(TJ): reorder data if this format is not nchw or x
+    dst.copyFrom(*m_);
+  }
+
 public:
   /**
    * Reorder this MKLDNNMatrix from other format.

paddle/math/Vector.cpp

@@ -18,6 +18,7 @@ limitations under the License. */
 #include <memory>
 #include "Matrix.h"
 #include "hl_gpu.h"
+#include "hl_matrix.h"
 #include "hl_table_apply.h"
 #include "paddle/utils/Flags.h"
 #include "paddle/utils/Logging.h"
@@ -99,6 +100,19 @@ MatrixPtr VectorT<int>::toOneHotSparseMatrix(size_t idRange, bool useGpu) {
   return mat;
 }
 
+template <>
+std::shared_ptr<VectorT<int>> VectorT<real>::castToInt() {
+  std::shared_ptr<VectorT<int>> ret = IVector::create(this->getSize(), useGpu_);
+  if (useGpu_) {
+    hl_vector_cast2int(ret->getData(), this->getData(), this->getSize());
+  } else {
+    for (size_t i = 0; i < getSize(); ++i) {
+      ret->getData()[i] = int(this->getData()[i]);
+    }
+  }
+  return ret;
+}
+
 template <class T>
 GpuVectorT<T>::GpuVectorT(size_t size)
     : VectorT<T>(size,

paddle/math/Vector.h

@@ -162,6 +162,13 @@ public:
    */
   std::shared_ptr<Matrix> toOneHotSparseMatrix(size_t idRange, bool useGpu);
 
+  /**
+   * @brief cast vector of "real" elements to "int" elements.
+   *
+   * @note: float -> int must be casted, or you'll get wrong data.
+   */
+  std::shared_ptr<VectorT<int>> castToInt();
+
   /**
    * This function will crash if the size of src and dest is different.
    */

paddle/memory/detail/buddy_allocator.cc

@@ -27,11 +27,11 @@ BuddyAllocator::BuddyAllocator(SystemAllocator* system_allocator,
       system_allocator_(std::move(system_allocator)) {}
 
 BuddyAllocator::~BuddyAllocator() {
-  VLOG(3) << "BuddyAllocator Disconstructor makes sure that all of these "
-             "have actually been freed";
+  VLOG(10) << "BuddyAllocator Disconstructor makes sure that all of these "
+              "have actually been freed";
   while (!pool_.empty()) {
     auto block = static_cast<MemoryBlock*>(std::get<2>(*pool_.begin()));
-    VLOG(3) << "Free from block (" << block << ", " << max_chunk_size_ << ")";
+    VLOG(10) << "Free from block (" << block << ", " << max_chunk_size_ << ")";
 
     system_allocator_->Free(block, max_chunk_size_, block->index(cache_));
     cache_.invalidate(block);
@@ -51,11 +51,12 @@ void* BuddyAllocator::Alloc(size_t unaligned_size) {
   // acquire the allocator lock
   std::lock_guard<std::mutex> lock(mutex_);
 
-  VLOG(3) << "Allocate " << unaligned_size << " bytes from chunk size " << size;
+  VLOG(10) << "Allocate " << unaligned_size << " bytes from chunk size "
+           << size;
 
   // if the allocation is huge, send directly to the system allocator
   if (size > max_chunk_size_) {
-    VLOG(3) << "Allocate from system allocator.";
+    VLOG(10) << "Allocate from system allocator.";
     return SystemAlloc(size);
   }
 
@@ -70,9 +71,9 @@ void* BuddyAllocator::Alloc(size_t unaligned_size) {
       return nullptr;
     }
   } else {
-    VLOG(3) << "Allocation from existing memory block " << std::get<2>(*it)
-            << " at address "
-            << reinterpret_cast<MemoryBlock*>(std::get<2>(*it))->data();
+    VLOG(10) << "Allocation from existing memory block " << std::get<2>(*it)
+             << " at address "
+             << reinterpret_cast<MemoryBlock*>(std::get<2>(*it))->data();
   }
 
   total_used_ += size;
@@ -89,10 +90,10 @@ void BuddyAllocator::Free(void* p) {
   // Acquire the allocator lock
   std::lock_guard<std::mutex> lock(mutex_);
 
-  VLOG(3) << "Free from address " << block;
+  VLOG(10) << "Free from address " << block;
 
   if (block->type(cache_) == MemoryBlock::HUGE_CHUNK) {
-    VLOG(3) << "Free directly from system allocator";
+    VLOG(10) << "Free directly from system allocator";
     system_allocator_->Free(block, block->total_size(cache_),
                             block->index(cache_));
 
@@ -109,8 +110,8 @@ void BuddyAllocator::Free(void* p) {
 
   // Trying to merge the right buddy
   if (block->has_right_buddy(cache_)) {
-    VLOG(3) << "Merging this block " << block << " with its right buddy "
-            << block->right_buddy(cache_);
+    VLOG(10) << "Merging this block " << block << " with its right buddy "
+             << block->right_buddy(cache_);
 
     auto right_buddy = block->right_buddy(cache_);
 
@@ -127,8 +128,8 @@ void BuddyAllocator::Free(void* p) {
 
   // Trying to merge the left buddy
   if (block->has_left_buddy(cache_)) {
-    VLOG(3) << "Merging this block " << block << " with its left buddy "
-            << block->left_buddy(cache_);
+    VLOG(10) << "Merging this block " << block << " with its left buddy "
+             << block->left_buddy(cache_);
 
     auto left_buddy = block->left_buddy(cache_);
 
@@ -144,8 +145,8 @@ void BuddyAllocator::Free(void* p) {
   }
 
   // Dumping this block into pool
-  VLOG(3) << "Inserting free block (" << block << ", "
-          << block->total_size(cache_) << ")";
+  VLOG(10) << "Inserting free block (" << block << ", "
+           << block->total_size(cache_) << ")";
   pool_.insert(
       IndexSizeAddress(block->index(cache_), block->total_size(cache_), block));
 
@@ -164,7 +165,7 @@ void* BuddyAllocator::SystemAlloc(size_t size) {
   size_t index = 0;
   void* p = system_allocator_->Alloc(index, size);
 
-  VLOG(3) << "Allocated " << p << " from system allocator.";
+  VLOG(10) << "Allocated " << p << " from system allocator.";
 
   if (p == nullptr) return nullptr;
 
@@ -190,8 +191,8 @@ BuddyAllocator::PoolSet::iterator BuddyAllocator::RefillPool() {
 
   if (p == nullptr) return pool_.end();
 
-  VLOG(3) << "Creating and inserting new block " << p
-          << " from system allocator";
+  VLOG(10) << "Creating and inserting new block " << p
+           << " from system allocator";
 
   static_cast<MemoryBlock*>(p)->init(cache_, MemoryBlock::FREE_CHUNK, index,
                                      max_chunk_size_, nullptr, nullptr);
@@ -235,19 +236,19 @@ void* BuddyAllocator::SplitToAlloc(BuddyAllocator::PoolSet::iterator it,
   auto block = static_cast<MemoryBlock*>(std::get<2>(*it));
   pool_.erase(it);
 
-  VLOG(3) << "Split block (" << block << ", " << block->total_size(cache_)
-          << ") into";
+  VLOG(10) << "Split block (" << block << ", " << block->total_size(cache_)
+           << ") into";
   block->split(cache_, size);
 
-  VLOG(3) << "Left block (" << block << ", " << block->total_size(cache_)
-          << ")";
+  VLOG(10) << "Left block (" << block << ", " << block->total_size(cache_)
+           << ")";
   block->set_type(cache_, MemoryBlock::ARENA_CHUNK);
 
   // the rest of memory if exist
   if (block->has_right_buddy(cache_)) {
     if (block->right_buddy(cache_)->type(cache_) == MemoryBlock::FREE_CHUNK) {
-      VLOG(3) << "Insert right block (" << block->right_buddy(cache_) << ", "
-              << block->right_buddy(cache_)->total_size(cache_) << ")";
+      VLOG(10) << "Insert right block (" << block->right_buddy(cache_) << ", "
+               << block->right_buddy(cache_)->total_size(cache_) << ")";
 
       pool_.insert(
           IndexSizeAddress(block->right_buddy(cache_)->index(cache_),
@@ -274,7 +275,7 @@ void BuddyAllocator::CleanIdleFallBackAlloc() {
       return;
     }
 
-    VLOG(3) << "Return block " << block << " to fallback allocator.";
+    VLOG(10) << "Return block " << block << " to fallback allocator.";
 
     system_allocator_->Free(block, max_chunk_size_, block->index(cache_));
     cache_.invalidate(block);
@@ -310,7 +311,7 @@ void BuddyAllocator::CleanIdleNormalAlloc() {
 
     MemoryBlock* block = static_cast<MemoryBlock*>(std::get<2>(*pool));
 
-    VLOG(3) << "Return block " << block << " to base allocator.";
+    VLOG(10) << "Return block " << block << " to base allocator.";
 
     system_allocator_->Free(block, max_chunk_size_, block->index(cache_));
     cache_.invalidate(block);

paddle/memory/detail/meta_cache.cc

@@ -30,7 +30,7 @@ Metadata MetadataCache::load(const MemoryBlock* block) {
     return existing_metadata->second;
   } else {
     auto* meta = reinterpret_cast<const Metadata*>(block);
-    VLOG(3) << "Load MetaData type=" << meta->type;
+    VLOG(10) << "Load MetaData type=" << meta->type;
     PADDLE_ASSERT(meta->check_guards());
     return *reinterpret_cast<const Metadata*>(block);
   }

paddle/memory/detail/system_allocator.cc

@@ -41,7 +41,16 @@ void* CPUAllocator::Alloc(size_t& index, size_t size) {
 
   index = 0;  // unlock memory
 
-  void* p = malloc(size);
+  void* p;
+
+#ifdef PADDLE_USE_MKLDNN
+  // refer to https://github.com/01org/mkl-dnn/blob/master/include/mkldnn.hpp
+  // memory alignment
+  PADDLE_ENFORCE_EQ(posix_memalign(&p, 4096ul, size), 0);
+#else
+  PADDLE_ENFORCE_EQ(posix_memalign(&p, 32ul, size), 0);
+#endif
+  PADDLE_ENFORCE(p, "Fail to allocate CPU memory: size = %d .", size);
 
   if (p != nullptr) {
     if (FLAGS_use_pinned_memory) {

paddle/memory/memory.cc

@@ -39,15 +39,15 @@ BuddyAllocator* GetCPUBuddyAllocator() {
 
 template <>
 void* Alloc<platform::CPUPlace>(platform::CPUPlace place, size_t size) {
-  VLOG(3) << "Allocate " << size << " bytes on " << platform::Place(place);
+  VLOG(10) << "Allocate " << size << " bytes on " << platform::Place(place);
   void* p = GetCPUBuddyAllocator()->Alloc(size);
-  VLOG(3) << "  pointer=" << p;
+  VLOG(10) << "  pointer=" << p;
   return p;
 }
 
 template <>
 void Free<platform::CPUPlace>(platform::CPUPlace place, void* p) {
-  VLOG(3) << "Free pointer=" << p << " on " << platform::Place(place);
+  VLOG(10) << "Free pointer=" << p << " on " << platform::Place(place);
   GetCPUBuddyAllocator()->Free(p);
 }
 
@@ -69,11 +69,12 @@ BuddyAllocator* GetGPUBuddyAllocator(int gpu_id) {
                                    platform::GpuMinChunkSize(),
                                    platform::GpuMaxChunkSize());
     }
-    VLOG(3) << "\n\nNOTE: each GPU device use "
-            << FLAGS_fraction_of_gpu_memory_to_use * 100 << "% of GPU memory.\n"
-            << "You can set environment variable '"
-            << platform::kEnvFractionGpuMemoryToUse
-            << "' to change the fraction of GPU usage.\n\n";
+    VLOG(10) << "\n\nNOTE: each GPU device use "
+             << FLAGS_fraction_of_gpu_memory_to_use * 100
+             << "% of GPU memory.\n"
+             << "You can set environment variable '"
+             << platform::kEnvFractionGpuMemoryToUse
+             << "' to change the fraction of GPU usage.\n\n";
   }
   platform::SetDeviceId(gpu_id);
   return as[gpu_id];

paddle/operators/CMakeLists.txt

@@ -131,9 +131,10 @@ add_subdirectory(math)
 add_subdirectory(nccl)
 
 set(DEPS_OPS
-    recurrent_op
     cond_op
     cross_entropy_op
+    recurrent_op
+    dynamic_recurrent_op
     softmax_with_cross_entropy_op
     sum_op
     pool_op
@@ -142,9 +143,6 @@ set(DEPS_OPS
     sequence_conv_op
     lstm_op)
 
-
-op_library(recurrent_op SRCS recurrent_op.cc rnn/recurrent_op_utils.cc
-  DEPS framework_proto tensor net_op)
 op_library(cond_op SRCS cond_op.cc DEPS framework_proto tensor operator net_op)
 op_library(cross_entropy_op DEPS cross_entropy)
 op_library(softmax_with_cross_entropy_op DEPS cross_entropy softmax)
@@ -156,7 +154,9 @@ op_library(nccl_op DEPS nccl_common)
 endif()
 op_library(sequence_conv_op DEPS context_project)
 op_library(lstm_op DEPS sequence2batch lstm_compute)
-
+op_library(dynamic_recurrent_op SRCS dynamic_recurrent_op.cc rnn/recurrent_op_utils.cc
+        DEPS net_op tensor_array)
+op_library(recurrent_op SRCS recurrent_op.cc DEPS executor)
 list(REMOVE_ITEM GENERAL_OPS ${DEPS_OPS})
 foreach(src ${GENERAL_OPS})
     op_library(${src})
@@ -168,8 +168,9 @@ cc_test(gather_test SRCS gather_test.cc DEPS tensor)
 cc_test(net_op_test SRCS net_op_test.cc DEPS net_op)
 cc_test(scatter_test SRCS scatter_test.cc DEPS tensor)
 cc_test(strided_memcpy_test SRCS strided_memcpy_test.cc DEPS tensor paddle_memory)
-cc_test(dynamic_recurrent_op_test SRCS dynamic_recurrent_op_test.cc DEPS dynamic_recurrent_op recurrent_op tensor_array)
-
+cc_test(dynamic_recurrent_op_test SRCS dynamic_recurrent_op_test.cc
+        rnn/recurrent_op_utils.cc
+        DEPS dynamic_recurrent_op)
 if(WITH_GPU)
   nv_test(nccl_op_test SRCS nccl_op_test.cu DEPS nccl_op gpu_info device_context)
 endif()

paddle/operators/activation_op.cc

@@ -43,7 +43,12 @@ class SigmoidOpMaker : public framework::OpProtoAndCheckerMaker {
       : OpProtoAndCheckerMaker(proto, op_checker) {
     AddInput("X", "Input of Sigmoid operator");
     AddOutput("Y", "Output of Sigmoid operator");
-    AddComment("Sigmoid activation operator, sigmoid = 1 / (1 + exp(-x))");
+    AddComment(R"DOC(
+Sigmoid activation operator.
+
+$y = 1 / (1 + e^{-x})$
+
+)DOC");
   }
 };
 
@@ -54,8 +59,12 @@ class LogSigmoidOpMaker : public framework::OpProtoAndCheckerMaker {
       : OpProtoAndCheckerMaker(proto, op_checker) {
     AddInput("X", "Input of LogSigmoid operator");
     AddOutput("Y", "Output of LogSigmoid operator");
-    AddComment(
-        "Logsigmoid activation operator, logsigmoid = log (1 / (1 + exp(-x)))");
+    AddComment(R"DOC(
+Logsigmoid activation operator.
+
+$y = \log(1 / (1 + e^{-x}))$
+
+)DOC");
   }
 };
 
@@ -65,7 +74,12 @@ class ExpOpMaker : public framework::OpProtoAndCheckerMaker {
       : OpProtoAndCheckerMaker(proto, op_checker) {
     AddInput("X", "Input of Exp operator");
     AddOutput("Y", "Output of Exp operator");
-    AddComment("Exp activation operator, exp(x) = e^x");
+    AddComment(R"DOC(
+Exp activation operator.
+
+$y = e^x$
+
+)DOC");
   }
 };
 
@@ -75,7 +89,12 @@ class ReluOpMaker : public framework::OpProtoAndCheckerMaker {
       : OpProtoAndCheckerMaker(proto, op_checker) {
     AddInput("X", "Input of Relu operator");
     AddOutput("Y", "Output of Relu operator");
-    AddComment("Relu activation operator, relu(x) = max(x, 0)");
+    AddComment(R"DOC(
+Relu activation operator.
+
+$y = \max(x, 0)$
+
+)DOC");
   }
 };
 
@@ -87,11 +106,14 @@ class LeakyReluOpMaker : public framework::OpProtoAndCheckerMaker {
       : OpProtoAndCheckerMaker(proto, op_checker) {
     AddInput("X", "Input of LeakyRelu operator");
     AddOutput("Y", "Output of LeakyRelu operator");
-    AddComment(
-        "LeakyRelu activation operator, "
-        "leaky_relu = max(x, alpha * x)");
     AddAttr<AttrType>("alpha", "The small negative slope")
         .SetDefault(static_cast<AttrType>(0.02f));
+    AddComment(R"DOC(
+LeakyRelu activation operator.
+
+$y = \max(x, \alpha * x)$
+
+)DOC");
   }
 };
 
@@ -103,12 +125,20 @@ class SoftShrinkOpMaker : public framework::OpProtoAndCheckerMaker {
       : OpProtoAndCheckerMaker(proto, op_checker) {
     AddInput("X", "Input of Softshrink operator");
     AddOutput("Y", "Output of Softshrink operator");
-    AddComment(
-        "Softshrink activation operator, "
-        "softshrink = x - lambda, if x > lambda;"
-        " x + lambda, if x < lambda; 0 otherwise");
     AddAttr<AttrType>("lambda", "non-negative offset")
         .SetDefault(static_cast<AttrType>(0.5f));
+    AddComment(R"DOC(
+Softshrink activation operator.
+
+$$
+y = \begin{cases} 
+    x - \lambda, \text{if } x > \lambda \\
+    x + \lambda, \text{if } x < -\lambda \\
+    0,  \text{otherwise}
+    \end{cases}
+$$
+
+)DOC");
   }
 };
 
@@ -118,9 +148,12 @@ class TanhOpMaker : public framework::OpProtoAndCheckerMaker {
       : OpProtoAndCheckerMaker(proto, op_checker) {
     AddInput("X", "Input of Tanh operator");
     AddOutput("Y", "Output of Tanh operator");
-    AddComment(
-        "Tanh activation operator, tanh = (exp(x) - exp(-x)) / (exp(x) + "
-        "exp(-x))");
+    AddComment(R"DOC(
+Tanh activation operator.
+
+$$y = \frac{e^{x} - e^{-x}}{e^{x} + e^{-x}}$$
+
+)DOC");
   }
 };
 
@@ -131,7 +164,12 @@ class TanhShrinkOpMaker : public framework::OpProtoAndCheckerMaker {
       : OpProtoAndCheckerMaker(proto, op_checker) {
     AddInput("X", "Input of TanhShrink operator");
     AddOutput("Y", "Output of TanhShrink operator");
-    AddComment("TanhShrink activation operator, tanhshrink(x) = x - tanh(x)");
+    AddComment(R"DOC(
+TanhShrink activation operator.
+
+$$y = x - \frac{e^{x} - e^{-x}}{e^{x} + e^{-x}}$$
+
+)DOC");
   }
 };
 
@@ -143,13 +181,20 @@ class HardShrinkOpMaker : public framework::OpProtoAndCheckerMaker {
       : OpProtoAndCheckerMaker(proto, op_checker) {
     AddInput("X", "Input of HardShrink operator");
     AddOutput("Y", "Output of HardShrink operator");
-    AddComment(
-        "HardShrink activation operator, "
-        "hard_shrink(x) = x if x > lambda"
-        "hard_shrink(x) = x if x < -lambda"
-        "hard_shrink(x) = 0 otherwise");
     AddAttr<AttrType>("threshold", "The value of threshold for HardShrink")
         .SetDefault(static_cast<AttrType>(0.5));
+    AddComment(R"DOC(
+HardShrink activation operator.
+
+$$
+y = \begin{cases} 
+    x, \text{if } x > \lambda \\
+    x, \text{if } x < -\lambda \\
+    0,  \text{otherwise}
+    \end{cases}
+$$
+
+)DOC");
   }
 };
 
@@ -159,7 +204,12 @@ class SqrtOpMaker : public framework::OpProtoAndCheckerMaker {
       : OpProtoAndCheckerMaker(proto, op_checker) {
     AddInput("X", "Input of Sqrt operator");
     AddOutput("Y", "Output of Sqrt operator");
-    AddComment("Sqrt activation operator, sqrt(x) = x^(1/2)");
+    AddComment(R"DOC(
+Sqrt activation operator.
+
+$y = \sqrt{x}$
+
+)DOC");
   }
 };
 
@@ -169,7 +219,12 @@ class AbsOpMaker : public framework::OpProtoAndCheckerMaker {
       : OpProtoAndCheckerMaker(proto, op_checker) {
     AddInput("X", "Input of Abs operator");
     AddOutput("Y", "Output of Abs operator");
-    AddComment("Abs activation operator, abs(x) = |x|");
+    AddComment(R"DOC(
+Abs activation operator.
+
+$y = |x|$
+
+)DOC");
   }
 };
 
@@ -180,7 +235,12 @@ class ReciprocalOpMaker : public framework::OpProtoAndCheckerMaker {
       : OpProtoAndCheckerMaker(proto, op_checker) {
     AddInput("X", "Input of Reciprocal operator");
     AddOutput("Y", "Output of Reciprocal operator");
-    AddComment("Reciprocal activation operator, reciprocal(x) = 1 / x");
+    AddComment(R"DOC(
+Reciprocal activation operator.
+
+$$y = \frac{1}{x}$$
+
+)DOC");
   }
 };
 
@@ -190,7 +250,14 @@ class LogOpMaker : public framework::OpProtoAndCheckerMaker {
       : OpProtoAndCheckerMaker(proto, op_checker) {
     AddInput("X", "Input of Log operator");
     AddOutput("Y", "Output of Log operator");
-    AddComment("Log activation operator, log(x) = natural logarithm of x");
+    AddComment(R"DOC(
+Log activation operator.
+
+$y = \ln(x)$
+
+Natural logarithm of x.
+
+)DOC");
   }
 };
 
@@ -200,7 +267,12 @@ class SquareOpMaker : public framework::OpProtoAndCheckerMaker {
       : OpProtoAndCheckerMaker(proto, op_checker) {
     AddInput("X", "Input of Square operator");
     AddOutput("Y", "Output of Square operator");
-    AddComment("Square activation operator, square(x) = x^2");
+    AddComment(R"DOC(
+Square activation operator.
+
+$y = x^2$
+
+)DOC");
   }
 };
 
@@ -211,7 +283,12 @@ class SoftplusOpMaker : public framework::OpProtoAndCheckerMaker {
       : OpProtoAndCheckerMaker(proto, op_checker) {
     AddInput("X", "Input of Softplus operator");
     AddOutput("Y", "Output of Softplus operator");
-    AddComment("Softplus activation operator, softplus(x) = log(1 + exp(x))");
+    AddComment(R"DOC(
+Softplus activation operator.
+
+$y = \ln(1 + e^{x})$
+
+)DOC");
   }
 };
 
@@ -222,7 +299,12 @@ class SoftsignOpMaker : public framework::OpProtoAndCheckerMaker {
       : OpProtoAndCheckerMaker(proto, op_checker) {
     AddInput("X", "Input of Softsign operator");
     AddOutput("Y", "Output of Softsign operator");
-    AddComment("Softsign activation operator, softsign(x) = x / (1 + |x|)");
+    AddComment(R"DOC(
+Softsign activation operator.
+
+$$y = \frac{x}{1 + |x|}$$
+
+)DOC");
   }
 };
 
@@ -233,11 +315,16 @@ class BReluOpMaker : public framework::OpProtoAndCheckerMaker {
       : OpProtoAndCheckerMaker(proto, op_checker) {
     AddInput("X", "Input of BRelu operator");
     AddOutput("Y", "Output of BRelu operator");
-    AddComment("BRelu activation operator, brelu = max(min(x, t_min), t_max)");
     AddAttr<AttrType>("t_min", "The min marginal value of BRelu")
         .SetDefault(static_cast<AttrType>(0));
     AddAttr<AttrType>("t_max", "The max marginal value of BRelu")
         .SetDefault(static_cast<AttrType>(24));
+    AddComment(R"DOC(
+BRelu activation operator.
+
+$y = \max(\min(x, t_{min}), t_{max})$
+
+)DOC");
   }
 };
 
@@ -249,11 +336,14 @@ class SoftReluOpMaker : public framework::OpProtoAndCheckerMaker {
       : OpProtoAndCheckerMaker(proto, op_checker) {
     AddInput("X", "Input of SoftRelu operator");
     AddOutput("Y", "Output of SoftRelu operator");
-    AddComment(
-        "SoftRelu activation operator, soft_relu = log(1 + exp(max(min(x, "
-        "threshold), threshold)))");
     AddAttr<AttrType>("threshold", "The threshold value of SoftRelu")
         .SetDefault(static_cast<AttrType>(40));
+    AddComment(R"DOC(
+SoftRelu activation operator.
+
+$y = \ln(1 + \exp(\max(\min(x, threshold), threshold))$
+
+)DOC");
   }
 };
 
@@ -262,19 +352,19 @@ class ELUOpMaker : public framework::OpProtoAndCheckerMaker {
  public:
   ELUOpMaker(framework::OpProto *proto, framework::OpAttrChecker *op_checker)
       : OpProtoAndCheckerMaker(proto, op_checker) {
-    AddInput("X",
-             "(Tensor) The input of ELU operator, it shouldn't be empty. Input "
-             "is flattened and treated as a 1D array.");
-    AddOutput("Y",
-              "(Tensor) The output of ELU operator. It has the same shape as "
-              "the input.");
-    AddAttr<AttrType>(
-        "alpha", "(float, default 1.0) Alpha value in the elu formulation.")
-        .SetDefault(static_cast<AttrType>(1.));
+    AddInput("X", "Input of ELU operator");
+    AddOutput("Y", "Output of ELU operator");
+    AddAttr<AttrType>("alpha", "The alpha value of ELU")
+        .SetDefault(static_cast<AttrType>(1.0f));
     AddComment(R"DOC(
-        ELU activation operator. It applies this element-wise computation on
-        the input: f(x) = max(0, x) + min(0, alpha * (exp(x) - 1)).
-        Check .. _Link: https://arxiv.org/abs/1511.07289 for more details.)DOC");
+ELU activation operator.
+
+Applies the following element-wise computation on the input according to
+https://arxiv.org/abs/1511.07289.
+
+$y = \max(0, x) + \min(0, \alpha * (e^x - 1))$
+
+)DOC");
   }
 };
 
@@ -285,9 +375,14 @@ class Relu6OpMaker : public framework::OpProtoAndCheckerMaker {
       : OpProtoAndCheckerMaker(proto, op_checker) {
     AddInput("X", "Input of Relu6 operator");
     AddOutput("Y", "Output of Relu6 operator");
-    AddComment("Relu6 activation operator, relu6 = min(max(0, x), 6)");
     AddAttr<AttrType>("threshold", "The threshold value of Relu6")
         .SetDefault(static_cast<AttrType>(6));
+    AddComment(R"DOC(
+Relu6 activation operator.
+
+$y = \min(\max(0, x), 6)$
+
+)DOC");
   }
 };
 
@@ -298,9 +393,14 @@ class PowOpMaker : public framework::OpProtoAndCheckerMaker {
       : OpProtoAndCheckerMaker(proto, op_checker) {
     AddInput("X", "Input of Pow operator");
     AddOutput("Y", "Output of Pow operator");
-    AddComment("Pow activation operator, pow(x, factor) = x^factor");
     AddAttr<AttrType>("factor", "The exponential factor of Pow")
         .SetDefault(static_cast<AttrType>(1));
+    AddComment(R"DOC(
+Pow activation operator.
+
+$y = x^{factor}$
+
+)DOC");
   }
 };
 
@@ -311,11 +411,16 @@ class STanhOpMaker : public framework::OpProtoAndCheckerMaker {
       : OpProtoAndCheckerMaker(proto, op_checker) {
     AddInput("X", "Input of STanh operator");
     AddOutput("Y", "Output of STanh operator");
-    AddComment("STanh activation operator, stanh = b * tanh(a * x)");
     AddAttr<AttrType>("scale_a", "The scale parameter of a for the input")
         .SetDefault(static_cast<AttrType>(2 / 3));
     AddAttr<AttrType>("scale_b", "The scale parameter of b for the input")
         .SetDefault(static_cast<AttrType>(1.7159));
+    AddComment(R"DOC(
+STanh activation operator.
+
+$$y = b * \frac{e^{a * x} - e^{-a * x}}{e^{a * x} + e^{-a * x}}$$
+
+)DOC");
   }
 };
 
@@ -327,12 +432,19 @@ class ThresholdedReluOpMaker : public framework::OpProtoAndCheckerMaker {
       : OpProtoAndCheckerMaker(proto, op_checker) {
     AddInput("X", "Input of ThresholdedRelu operator");
     AddOutput("Y", "Output of ThresholdedRelu operator");
-    AddComment(
-        "ThresholdedRelu activation operator, "
-        "thresholded_relu = x for x > threshold, "
-        "thresholded_relu = 0 otherwise.");
     AddAttr<AttrType>("threshold", "The threshold location of activation")
         .SetDefault(static_cast<AttrType>(1.0));
+    AddComment(R"DOC(
+ThresholdedRelu activation operator.
+
+$$
+y = \begin{cases} 
+    x, \text{if } x > threshold \\
+    0,  \text{otherwise}
+    \end{cases}
+$$
+
+)DOC");
   }
 };
 
@@ -344,27 +456,23 @@ class HardSigmoidOpMaker : public framework::OpProtoAndCheckerMaker {
       : OpProtoAndCheckerMaker(proto, op_checker) {
     AddInput("X", "Input of HardSigmoid operator");
     AddOutput("Y", "Output of HardSigmoid operator");
+    AddAttr<AttrType>("slope", "Slope for linear approximation of sigmoid")
+        .SetDefault(static_cast<AttrType>(0.2));
+    AddAttr<AttrType>("offset", "Offset for linear approximation of sigmoid")
+        .SetDefault(static_cast<AttrType>(0.5));
     AddComment(R"DOC(
-Hard Sigmoid activation operator.
+HardSigmoid activation operator.
 
-Segment-wise linear approximation of sigmoid[1].
-This is much faster than sigmoid.
+Segment-wise linear approximation of sigmoid(https://arxiv.org/abs/1603.00391), 
+which is much faster than sigmoid.
 
-hard_sigmoid = max(0, min(1, slope * x + shift))
+$y = \max(0, \min(1, slope * x + shift))$
 
 The slope should be positive. The offset can be either positive or negative.
-The default slope and shift are set from [1].
+The default slope and shift are set according to the above reference.
 It is recommended to use the defaults for this activation.
 
-References:
-  [1] Noisy Activation Functions
-      (https://arxiv.org/abs/1603.00391)
-
-    )DOC");
-    AddAttr<AttrType>("slope", "Slope for linear approximation of sigmoid")
-        .SetDefault(static_cast<AttrType>(0.2));
-    AddAttr<AttrType>("offset", "Offset for linear approximation of sigmoid")
-        .SetDefault(static_cast<AttrType>(0.5));
+)DOC");
   }
 };
 

paddle/operators/activation_op.h

@@ -232,7 +232,7 @@ struct HardShrinkGradFunctor : public BaseActivationFunctor<T> {
   }
 };
 
-// softshrink(x) = x - lambda, if x > lambda; x + lambda, if x < lambda; 0
+// softshrink(x) = x - lambda, if x > lambda; x + lambda, if x < -lambda; 0
 // otherwise
 template <typename T>
 struct SoftShrinkFunctor : public BaseActivationFunctor<T> {

paddle/operators/conv2d_transpose_cudnn_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/conv2d_transpose_op.h"
+
+namespace paddle {
+namespace operators {
+
+class CudnnConv2DTransposeOpMaker : public Conv2DTransposeOpMaker {
+ public:
+  CudnnConv2DTransposeOpMaker(framework::OpProto* proto,
+                              framework::OpAttrChecker* op_checker)
+      : Conv2DTransposeOpMaker(proto, op_checker) {
+    AddAttr<std::vector<int>>("dilations", "dilations of convolution operator.")
+        .SetDefault(std::vector<int>{1, 1});
+    AddAttr<int>("workspace_size_MB",
+                 "workspace size for cudnn, in MB, "
+                 "workspace is a section of GPU memory which will be "
+                 "allocated/freed each time the operator runs, larger "
+                 "workspace size can increase performance but also requires "
+                 "better hardward. This size should be carefully setted.")
+        .SetDefault(4096);
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle
+
+namespace ops = paddle::operators;
+REGISTER_OP(conv2d_transpose_cudnn, ops::Conv2DTransposeOp,
+            ops::CudnnConv2DTransposeOpMaker, conv2d_transpose_cudnn_grad,
+            ops::Conv2DTransposeOpGrad);
+
+REGISTER_OP_CPU_KERNEL(
+    conv2d_transpose_cudnn,
+    ops::GemmConv2DTransposeKernel<paddle::platform::CPUPlace, float>);
+REGISTER_OP_CPU_KERNEL(
+    conv2d_transpose_cudnn_grad,
+    ops::GemmConv2DTransposeGradKernel<paddle::platform::CPUPlace, float>);

paddle/operators/conv2d_transpose_cudnn_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/framework/eigen.h"
+#include "paddle/framework/op_registry.h"
+#include "paddle/memory/memory.h"
+#include "paddle/operators/conv2d_transpose_op.h"
+#include "paddle/platform/assert.h"
+#include "paddle/platform/cudnn_helper.h"
+
+namespace paddle {
+namespace operators {
+
+using Tensor = framework::Tensor;
+using ScopedTensorDescriptor = platform::ScopedTensorDescriptor;
+using ScopedFilterDescriptor = platform::ScopedFilterDescriptor;
+using ScopedConvolutionDescriptor = platform::ScopedConvolutionDescriptor;
+using DataLayout = platform::DataLayout;
+using CUDADeviceContext = platform::CUDADeviceContext;
+
+static constexpr size_t kConvCudnnWorkspaceLimitBytes = 1024 * 1024 * 1024;
+
+template <typename T>
+class CudnnConvTransposeOpKernel : public framework::OpKernel<T> {
+ public:
+  void Compute(const framework::ExecutionContext& ctx) const override {
+    PADDLE_ENFORCE(platform::is_gpu_place(ctx.GetPlace()),
+                   "It must use GPUPlace.");
+    auto* input = ctx.Input<Tensor>("Input");
+    auto* filter = ctx.Input<Tensor>("Filter");
+    auto* output = ctx.Output<Tensor>("Output");
+
+    std::vector<int> strides = ctx.Attr<std::vector<int>>("strides");
+    std::vector<int> paddings = ctx.Attr<std::vector<int>>("paddings");
+    // cudnn v5 does not support dilations
+    std::vector<int> dilations = ctx.Attr<std::vector<int>>("dilations");
+    int user_workspace_size = ctx.Attr<int>("workspace_size_MB");
+
+    const T* input_data = input->data<T>();
+    const T* filter_data = filter->data<T>();
+    T* output_data = output->mutable_data<T>(ctx.GetPlace());
+    // ------------------- cudnn descriptors ---------------------
+    ScopedTensorDescriptor input_desc;
+    ScopedTensorDescriptor output_desc;
+    ScopedFilterDescriptor filter_desc;
+    ScopedConvolutionDescriptor conv_desc;
+    DataLayout layout = DataLayout::kNCHW;
+
+    // N, M, H, W
+    cudnnTensorDescriptor_t cudnn_input_desc = input_desc.descriptor<T>(
+        layout, framework::vectorize2int(input->dims()));
+    // N, C, O_h, O_w
+    cudnnTensorDescriptor_t cudnn_output_desc = output_desc.descriptor<T>(
+        layout, framework::vectorize2int(output->dims()));
+    // M, C, K_h, K_w
+    cudnnFilterDescriptor_t cudnn_filter_desc = filter_desc.descriptor<T>(
+        layout, framework::vectorize2int(filter->dims()));
+    cudnnConvolutionDescriptor_t cudnn_conv_desc =
+        conv_desc.descriptor<T>(paddings, strides, dilations);
+
+    // ------------------- cudnn conv workspace ---------------------
+    void* cudnn_workspace = nullptr;
+    size_t workspace_size_in_bytes;  // final workspace to allocate.
+    size_t workspace_size_limit = kConvCudnnWorkspaceLimitBytes;
+    if (user_workspace_size > 0) {
+      workspace_size_limit = user_workspace_size * 1024 * 1024;
+    }
+    // ------------------- cudnn conv algorithm ---------------------
+    cudnnConvolutionBwdDataAlgo_t algo;
+    auto handle = ctx.cuda_device_context().cudnn_handle();
+    // Get the algorithm
+    PADDLE_ENFORCE(platform::dynload::cudnnGetConvolutionBackwardDataAlgorithm(
+        handle, cudnn_filter_desc, cudnn_input_desc, cudnn_conv_desc,
+        // dxDesc: Handle to the previously initialized output tensor
+        // descriptor.
+        cudnn_output_desc, CUDNN_CONVOLUTION_BWD_DATA_SPECIFY_WORKSPACE_LIMIT,
+        workspace_size_limit, &algo));
+
+    // get workspace size able to allocate
+    PADDLE_ENFORCE(
+        platform::dynload::cudnnGetConvolutionBackwardDataWorkspaceSize(
+            handle, cudnn_filter_desc, cudnn_input_desc, cudnn_conv_desc,
+            cudnn_output_desc, algo, &workspace_size_in_bytes));
+
+    // Allocate on GPU memory
+    platform::GPUPlace gpu = boost::get<platform::GPUPlace>(ctx.GetPlace());
+    cudnn_workspace = paddle::memory::Alloc(gpu, workspace_size_in_bytes);
+
+    // ------------------- cudnn conv transpose forward ---------------------
+    T alpha = 1.0f, beta = 0.0f;
+    PADDLE_ENFORCE(platform::dynload::cudnnConvolutionBackwardData(
+        handle, &alpha, cudnn_filter_desc, filter_data, cudnn_input_desc,
+        input_data, cudnn_conv_desc, algo, cudnn_workspace,
+        workspace_size_in_bytes, &beta, cudnn_output_desc, output_data));
+
+    // Release the cudnn workspace
+    paddle::memory::Free(gpu, cudnn_workspace);
+  }
+};
+
+template <typename T>
+class CudnnConvTransposeGradOpKernel : public framework::OpKernel<T> {
+ public:
+  void Compute(const framework::ExecutionContext& ctx) const override {
+    PADDLE_ENFORCE(platform::is_gpu_place(ctx.GetPlace()),
+                   "It must use GPUPlace.");
+    auto input = ctx.Input<Tensor>("Input");
+    auto filter = ctx.Input<Tensor>("Filter");
+    auto output_grad = ctx.Input<Tensor>(framework::GradVarName("Output"));
+    auto input_grad = ctx.Output<Tensor>(framework::GradVarName("Input"));
+    auto filter_grad = ctx.Output<Tensor>(framework::GradVarName("Filter"));
+    const T* input_data = input->data<T>();
+    const T* output_grad_data = output_grad->data<T>();
+    const T* filter_data = filter->data<T>();
+
+    std::vector<int> strides = ctx.Attr<std::vector<int>>("strides");
+    std::vector<int> paddings = ctx.Attr<std::vector<int>>("paddings");
+    // cudnn v5 does not support dilations
+    std::vector<int> dilations = ctx.Attr<std::vector<int>>("dilations");
+    int user_workspace_size = ctx.Attr<int>("workspace_size_MB");
+
+    // ------------------- cudnn descriptors ---------------------
+    ScopedTensorDescriptor input_desc;
+    ScopedTensorDescriptor output_desc;
+    ScopedFilterDescriptor filter_desc;
+    ScopedConvolutionDescriptor conv_desc;
+    DataLayout layout = DataLayout::kNCHW;
+
+    // Input: (N, M, H, W)
+    cudnnTensorDescriptor_t cudnn_input_desc = input_desc.descriptor<T>(
+        layout, framework::vectorize2int(input->dims()));
+    // Output: (N, C, O_H, O_W)
+    cudnnTensorDescriptor_t cudnn_output_desc = output_desc.descriptor<T>(
+        layout, framework::vectorize2int(output_grad->dims()));
+    // Filter (M, C, K_H, K_W)
+    cudnnFilterDescriptor_t cudnn_filter_desc = filter_desc.descriptor<T>(
+        layout, framework::vectorize2int(filter->dims()));
+
+    cudnnConvolutionDescriptor_t cudnn_conv_desc =
+        conv_desc.descriptor<T>(paddings, strides, dilations);
+
+    // ------------------- cudnn backward algorithm ---------------------
+    cudnnConvolutionFwdAlgo_t data_algo;
+    cudnnConvolutionBwdFilterAlgo_t filter_algo;
+    size_t bwd_filter_ws_size, fwd_ws_size;
+    size_t workspace_size_in_bytes = 0;
+    size_t workspace_size_limit = kConvCudnnWorkspaceLimitBytes;
+    if (user_workspace_size > 0) {
+      workspace_size_limit = user_workspace_size * 1024 * 1024;
+    }
+
+    auto handle = ctx.cuda_device_context().cudnn_handle();
+    if (input_grad) {
+      // choose backward algorithm for data
+      PADDLE_ENFORCE(platform::dynload::cudnnGetConvolutionForwardAlgorithm(
+          handle, cudnn_output_desc, cudnn_filter_desc, cudnn_conv_desc,
+          cudnn_input_desc, CUDNN_CONVOLUTION_FWD_SPECIFY_WORKSPACE_LIMIT,
+          workspace_size_limit, &data_algo));
+      PADDLE_ENFORCE(platform::dynload::cudnnGetConvolutionForwardWorkspaceSize(
+          handle, cudnn_output_desc, cudnn_filter_desc, cudnn_conv_desc,
+          cudnn_input_desc, data_algo, &fwd_ws_size));
+      workspace_size_in_bytes = std::max(workspace_size_in_bytes, fwd_ws_size);
+    }
+
+    if (filter_grad) {
+      // choose backward algorithm for filter
+      PADDLE_ENFORCE(
+          platform::dynload::cudnnGetConvolutionBackwardFilterAlgorithm(
+              handle, cudnn_output_desc, cudnn_input_desc, cudnn_conv_desc,
+              cudnn_filter_desc,
+              CUDNN_CONVOLUTION_BWD_FILTER_SPECIFY_WORKSPACE_LIMIT,
+              workspace_size_limit, &filter_algo));
+
+      // get workspace for backwards filter algorithm
+      PADDLE_ENFORCE(
+          platform::dynload::cudnnGetConvolutionBackwardFilterWorkspaceSize(
+              handle, cudnn_output_desc, cudnn_input_desc, cudnn_conv_desc,
+              cudnn_filter_desc, filter_algo, &bwd_filter_ws_size));
+      workspace_size_in_bytes =
+          std::max(workspace_size_in_bytes, bwd_filter_ws_size);
+    }
+
+    // ------------------- cudnn conv workspace ---------------------
+    // Already on GPU
+    void* cudnn_workspace = nullptr;
+    platform::GPUPlace gpu = boost::get<platform::GPUPlace>(ctx.GetPlace());
+    cudnn_workspace = paddle::memory::Alloc(gpu, workspace_size_in_bytes);
+    // ------------------- cudnn conv backward data ---------------------
+    // FIXME(typhoonzero): template type T may not be the same as cudnn call.
+    T alpha = 1.0f, beta = 0.0f;
+    if (input_grad) {
+      T* input_grad_data = input_grad->mutable_data<T>(ctx.GetPlace());
+      auto t = framework::EigenVector<T>::Flatten(*input_grad);
+      t.device(ctx.GetEigenDevice<platform::GPUPlace>()) =
+          t.constant(static_cast<T>(0));
+
+      PADDLE_ENFORCE(platform::dynload::cudnnConvolutionForward(
+          handle, &alpha, cudnn_output_desc, output_grad_data,
+          cudnn_filter_desc, filter_data, cudnn_conv_desc, data_algo,
+          cudnn_workspace, workspace_size_in_bytes, &beta, cudnn_input_desc,
+          input_grad_data));
+    }
+
+    // ------------------- cudnn conv backward filter ---------------------
+    if (filter_grad) {
+      T* filter_grad_data = filter_grad->mutable_data<T>(ctx.GetPlace());
+      auto t = framework::EigenVector<T>::Flatten(*filter_grad);
+      t.device(ctx.GetEigenDevice<platform::GPUPlace>()) =
+          t.constant(static_cast<T>(0));
+      // Gradient with respect to the filter
+      PADDLE_ENFORCE(platform::dynload::cudnnConvolutionBackwardFilter(
+          handle, &alpha, cudnn_output_desc, output_grad_data, cudnn_input_desc,
+          input_data, cudnn_conv_desc, filter_algo, cudnn_workspace,
+          workspace_size_in_bytes, &beta, cudnn_filter_desc, filter_grad_data));
+    }
+    // Release the cudnn workspace
+    paddle::memory::Free(gpu, cudnn_workspace);
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle
+
+namespace ops = paddle::operators;
+
+REGISTER_OP_GPU_KERNEL(conv2d_transpose_cudnn,
+                       ops::CudnnConvTransposeOpKernel<float>);
+REGISTER_OP_GPU_KERNEL(conv2d_transpose_cudnn_grad,
+                       ops::CudnnConvTransposeGradOpKernel<float>);

paddle/operators/conv2dtranspose_op.cc → paddle/operators/conv2d_transpose_op.cc

@@ -12,7 +12,7 @@
    See the License for the specific language governing permissions and
    limitations under the License. */
 
-#include "paddle/operators/conv2dtranspose_op.h"
+#include "paddle/operators/conv2d_transpose_op.h"
 
 namespace paddle {
 namespace operators {
@@ -95,13 +95,13 @@ void Conv2DTransposeOpGrad::InferShape(
 }  // namespace paddle
 
 namespace ops = paddle::operators;
-REGISTER_OP(conv2dtranspose, ops::Conv2DTransposeOp,
-            ops::Conv2DTransposeOpMaker, conv2dtranspose_grad,
+REGISTER_OP(conv2d_transpose, ops::Conv2DTransposeOp,
+            ops::Conv2DTransposeOpMaker, conv2d_transpose_grad,
             ops::Conv2DTransposeOpGrad);
 
 REGISTER_OP_CPU_KERNEL(
-    conv2dtranspose,
+    conv2d_transpose,
     ops::GemmConv2DTransposeKernel<paddle::platform::CPUPlace, float>);
 REGISTER_OP_CPU_KERNEL(
-    conv2dtranspose_grad,
+    conv2d_transpose_grad,
     ops::GemmConv2DTransposeGradKernel<paddle::platform::CPUPlace, float>);

paddle/operators/conv2dtranspose_op.cu → paddle/operators/conv2d_transpose_op.cu

@@ -12,13 +12,13 @@
    See the License for the specific language governing permissions and
    limitations under the License. */
 
-#include "paddle/operators/conv2dtranspose_op.h"
+#include "paddle/operators/conv2d_transpose_op.h"
 
 namespace ops = paddle::operators;
 
 REGISTER_OP_GPU_KERNEL(
-    conv2dtranspose,
+    conv2d_transpose,
     ops::GemmConv2DTransposeKernel<paddle::platform::GPUPlace, float>);
 REGISTER_OP_GPU_KERNEL(
-    conv2dtranspose_grad,
+    conv2d_transpose_grad,
     ops::GemmConv2DTransposeGradKernel<paddle::platform::GPUPlace, float>);

paddle/operators/conv2dtranspose_op.h → paddle/operators/conv2d_transpose_op.h

@@ -62,7 +62,7 @@ class GemmConv2DTransposeKernel : public framework::OpKernel<T> {
     std::vector<int> strides = context.Attr<std::vector<int>>("strides");
 
     // TODO(Zhuoyuan): Paddings can be added in future.
-    // groups will alway be disabled in conv2dtranspose.
+    // groups will alway be disabled in conv2d_transpose.
 
     const int batch_size = input->dims()[0];
     const int m = input->dims()[1];

paddle/operators/cross_entropy_op.cc

@@ -28,8 +28,9 @@ class CrossEntropyOp : public framework::OperatorWithKernel {
 
     auto x_dims = ctx->GetInputDim("X");
     auto label_dims = ctx->GetInputDim("Label");
-    PADDLE_ENFORCE_EQ(x_dims.size(), 2, "Input(X)'s rank should be 2.");
-    PADDLE_ENFORCE_EQ(label_dims.size(), 2, "Input(Label)'s rank should be 2.");
+    PADDLE_ENFORCE_EQ(x_dims.size(), 2UL, "Input(X)'s rank should be 2.");
+    PADDLE_ENFORCE_EQ(label_dims.size(), 2UL,
+                      "Input(Label)'s rank should be 2.");
     PADDLE_ENFORCE_EQ(x_dims[0], label_dims[0],
                       "The 1st dimension of Input(X) and Input(Label) should "
                       "be equal.");
@@ -38,8 +39,8 @@ class CrossEntropyOp : public framework::OperatorWithKernel {
                         "If Attr(soft_label) == true, the 2nd dimension of "
                         "Input(X) and Input(Label) should be equal.");
     } else {
-      PADDLE_ENFORCE_EQ(label_dims[1], 1,
-                        "If Attr(soft_label) == false, the 2nd dimension of "
+      PADDLE_ENFORCE_EQ(label_dims[1], 1UL,
+                        "If Attr(softLabel) == false, the 2nd dimension of "
                         "Input(Label) should be 1.");
     }
 
@@ -48,7 +49,8 @@ class CrossEntropyOp : public framework::OperatorWithKernel {
   }
 
  protected:
-  // CrossEntropy's data type just determined by "X"
+  // Explicitly set that data type of the output of the cross_entropy operator
+  // is determined by its input "X".
   framework::DataType IndicateDataType(
       const framework::ExecutionContext& ctx) const override {
     return framework::ToDataType(ctx.Input<Tensor>("X")->type());

paddle/operators/dynamic_recurrent_op_test.cc

@@ -51,7 +51,7 @@ class RNNAlgorithmTestHelper : public ::testing::Test {
     CreateGlobalVariables();
 
     auto op_desc = CreateOpDesc();
-    op = paddle::framework::OpRegistry::CreateOp(op_desc, nullptr);
+    op = paddle::framework::OpRegistry::CreateOp(op_desc);
     dop = &(dynamic_cast<DynamicRecurrentOp*>(op.get())->rnn);
     InitCacheManually();
     InitStepNet();

paddle/operators/fill_constant_batch_size_like_op.cc

@@ -36,7 +36,12 @@ class FillConstantBatchSizeLikeOp : public framework::OperatorWithKernel {
                    [](int a) { return static_cast<int64_t>(a); });
     auto dims = framework::make_ddim(shape_int64);
 
-    dims[0] = ctx->GetInputDim("Input")[0];
+    int dim_idx = ctx->Attrs().Get<int>("dim_idx");
+    PADDLE_ENFORCE_GE(dim_idx, 0);
+    PADDLE_ENFORCE_GT(static_cast<int>(shape.size()), dim_idx);
+    PADDLE_ENFORCE_GT(ctx->GetInputDim("Input").size(), dim_idx);
+
+    dims[dim_idx] = ctx->GetInputDim("Input")[dim_idx];
     ctx->SetOutputDim("Out", dims);
   }
 
@@ -57,15 +62,18 @@ class FillConstantBatchSizeLikeOpMaker
                  "(int, default 5 (FP32)) "
                  "Output data type")
         .SetDefault(framework::DataType::FP32);
-    AddAttr<std::vector<int>>("shape", "(vector<int>) The shape of the output");
-    AddAttr<float>("value", "(float, default 0) The value to be filled")
-        .SetDefault(0.0f);
     AddInput("Input",
              "(Tensor) Tensor "
-             "whose first dimension is used to specify the batch_size");
+             "whose dim_idx th dimension is used to specify the batch_size");
     AddOutput("Out",
               "(Tensor) Tensor of specified shape will be filled "
               "with the specified value");
+    AddAttr<std::vector<int>>("shape", "(vector<int>) The shape of the output");
+    AddAttr<int>("dim_idx",
+                 "(int, default 0) the index of batch size dimension")
+        .SetDefault(0);
+    AddAttr<float>("value", "(float, default 0) The value to be filled")
+        .SetDefault(0.0f);
     AddComment(R"DOC(Fill up a variable with specified constant value.)DOC");
   }
 };

paddle/operators/gaussian_random_op.cc

@@ -45,14 +45,14 @@ class GaussianRandomOp : public framework::OperatorWithKernel {
   void InferShape(framework::InferShapeContext* ctx) const override {
     PADDLE_ENFORCE(ctx->HasOutput("Out"),
                    "Output(Out) of GaussianRandomOp should not be null.");
-    auto dims = ctx->Attrs().Get<std::vector<int>>("dims");
+    auto shape = ctx->Attrs().Get<std::vector<int>>("shape");
     std::vector<int64_t> temp;
-    temp.reserve(dims.size());
-    for (auto dim : dims) {
+    temp.reserve(shape.size());
+    for (auto dim : shape) {
       temp.push_back(static_cast<int64_t>(dim));
     }
-    PADDLE_ENFORCE(dims.size() > 0UL,
-                   "dims can be one int or array. dims must be set.");
+    PADDLE_ENFORCE(shape.size() > 0UL,
+                   "shape can be one int or array. shape must be set.");
     ctx->SetOutputDim("Out", framework::make_ddim(temp));
   }
 
@@ -74,7 +74,7 @@ GaussianRandom operator.
 Use to initialize tensor with gaussian random generator.
 )DOC");
 
-    AddAttr<std::vector<int>>("dims", "The dimension of random tensor.");
+    AddAttr<std::vector<int>>("shape", "The dimension of random tensor.");
     AddAttr<float>("mean", "mean of random tensor.").SetDefault(.0f);
     AddAttr<float>("std", "std of random tensor.").SetDefault(1.0f);
     AddAttr<int>("seed",

paddle/operators/linear_chain_crf_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/linear_chain_crf_op.h"
+
+namespace paddle {
+namespace operators {
+
+class LinearChainCRFOpMaker : public framework::OpProtoAndCheckerMaker {
+ public:
+  LinearChainCRFOpMaker(framework::OpProto* proto,
+                        framework::OpAttrChecker* op_checker)
+      : OpProtoAndCheckerMaker(proto, op_checker) {
+    AddInput(
+        "Emission",
+        "(LoDTensor, default: LoDTensor<float>). "
+        "The unscaled emission weight matrix for the linear chain CRF. "
+        "This input is a LoDTensor with shape [N x D] where N is the size of "
+        "the mini-batch and D is the total tag number.");
+    AddInput(
+        "Transition",
+        "(Tensor, default: Tensor<float>). A Tensor with shape [(D + 2) x D]. "
+        "The learnable parameter for the linear_chain_crf operator. "
+        "See more details in the operator's comments.");
+    AddInput(
+        "Label",
+        "(LoDTensor, default: LoDTensor<int>). The ground truth which is a 2-D "
+        "LoDTensor with shape [N x 1], where N is the total element number in "
+        "a mini-batch.");
+    AddOutput(
+        "Alpha",
+        "Tensor, default: Tensor<float>. The forward vectors for the entire "
+        "batch. A two dimensional tensor with shape [N x D], "
+        "denoted as \f$\alpha\f$. \f$\alpha$\f is a memo table used to "
+        "calculate the normalization factor in CRF. \f$\alpha[k, v]$\f stores "
+        "the unnormalized probabilites of all possible unfinished sequences of "
+        "tags that end at position \f$k$\f with tag \f$v$\f. For each \f$k$\f, "
+        "\f$\alpha[k, v]$\f is a vector of length \f$D$\f with a component for "
+        "each tag value \f$v$\f. This vector is called a forward vecotr and "
+        "will also be used in backward computations.")
+        .AsIntermediate();
+    AddOutput("EmissionExps",
+              "The exponentials of Input(Emission). This is an intermediate "
+              "computational result in forward computation, and will be reused "
+              "in backward computation.")
+        .AsIntermediate();
+    AddOutput("TransitionExps",
+              "The exponentials of Input(Transition). This is an intermediate "
+              "computational result in forward computation, and will be reused "
+              "in backward computation.")
+        .AsIntermediate();
+    AddOutput(
+        "LogLikelihood",
+        "(Tensor, default: Tensor<float>). The logarithm of the conditional "
+        "likelihood of each training sample in a mini-batch. This is a 2-D "
+        "tensor with shape [S x 1], where S is the sequence number in a "
+        "mini-batch. Note: S is equal to the sequence number in a mini-batch. "
+        "The output is no longer a LoDTensor.");
+    AddComment(R"DOC(
+Conditional Random Field defines an undirected probabilistic graph with nodes
+denoting random variables and edges denoting dependencies between these
+variables. CRF learns the conditional probability \f$P(Y|X)\f$, where
+\f$X = (x_1, x_2, ... , x_n)\f$ are structured inputs and
+\f$Y = (y_1, y_2, ... , y_n)\f$ are labels for the inputs.
+
+Linear chain CRF is a special case of CRF that is useful for sequence labeling
+task. Sequence labeling tasks do not assume a lot of conditional
+independences among inputs. The only constraint they impose is that the input
+and output must be linear sequences. Thus, the graph of such a CRF is a simple
+chain or a line, which results in the linear chain CRF.
+
+This operator implements the Forward-Backward algorithm for the linear chain
+CRF. Please see http://www.cs.columbia.edu/~mcollins/fb.pdf and
+http://cseweb.ucsd.edu/~elkan/250Bwinter2012/loglinearCRFs.pdf for reference.
+
+Equation:
+
+- Denote Input(Emission) to this operator as \f$x\f$ here.
+- The first D values of Input(Transition) to this operator are for starting
+weights, denoted as \f$a\f$ here.
+- The next D values of Input(Transition) of this operator are for ending
+weights, denoted as \f$b\f$ here.
+- The remaning values of Input(Transition) are for transition weights,
+denoted as \f$w\f$ here.
+- Denote Input(Label) as \f$s\f$ here.
+
+The probability of a sequence \f$s\f$ of length \f$L\f$ is defined as:
+\f$P(s) = (1/Z) exp(a_{s_1} + b_{s_L}
+                 + \sum_{l=1}^L x_{s_l}
+                 + \sum_{l=2}^L w_{s_{l-1},s_l})\f$
+where \f$Z\f$ is a normalization value so that the sum of \f$P(s)\f$ over
+all possible sequences is \f$1\f$, and \f$x\f$ is the emission feature weight
+to the linear chain CRF.
+
+Finaly, the linear chain CRF operator outputs the logarithm of the conditional
+likelihood of each training sample in a mini-batch.
+
+NOTE:
+1. The feature function for a CRF is made up of the emission features and the
+transition features. The emission feature weights are NOT computed in
+this operator. They MUST be computed first before this operator is called.
+
+2. Because this operator performs global normalization over all possible
+sequences internally, it expects UNSCALED emission feature weights.
+Please do not call this op with the emission feature being output of any
+nonlinear activation.
+
+3. The 2nd dimension of Input(Emission) MUST be equal to the tag number.
+
+)DOC");
+  }
+};
+
+class LinearChainCRFOp : public framework::OperatorWithKernel {
+ public:
+  using framework::OperatorWithKernel::OperatorWithKernel;
+
+  void InferShape(framework::InferShapeContext* ctx) const override {
+    PADDLE_ENFORCE(ctx->HasInput("Emission"),
+                   "Input(Emission) should be not null.");
+    PADDLE_ENFORCE(ctx->HasInput("Transition"),
+                   "Input(Transition) should be not null.");
+    PADDLE_ENFORCE(ctx->HasInput("Label"), "Input(Label) should be not null.");
+
+    PADDLE_ENFORCE(ctx->HasOutput("Alpha"),
+                   "Output(Alpha) should be not null.");
+    PADDLE_ENFORCE(ctx->HasOutput("EmissionExps"),
+                   "Output(EmissionExps) should be not null.");
+    PADDLE_ENFORCE(ctx->HasOutput("TransitionExps"),
+                   "Output(TransitionExps) should be not null.");
+    PADDLE_ENFORCE(ctx->HasOutput("LogLikelihood"),
+                   "Output(LogLikelihood) should be not null.");
+
+    auto emission_dims = ctx->GetInputDim("Emission");
+    PADDLE_ENFORCE_EQ(emission_dims.size(), 2UL,
+                      "The Input(Emission) should be a 2-D tensor.");
+    PADDLE_ENFORCE(emission_dims[0], "An empty mini-batch is not allowed.");
+
+    auto transition_dims = ctx->GetInputDim("Transition");
+    PADDLE_ENFORCE_EQ(transition_dims.size(), 2UL,
+                      "The Input(Transition) should be a 2-D tensor.");
+    PADDLE_ENFORCE_EQ(
+        transition_dims[0] - 2, transition_dims[1],
+        "An invalid dimension for the Input(Transition), which should "
+        "be a 2-D tensor with shape [(D + 2) x D].");
+    PADDLE_ENFORCE_EQ(
+        emission_dims[1], transition_dims[1],
+        "The 2nd dimension of the Input(Emission) and the Input(Transition) "
+        "should be equal to the tag number.");
+
+    auto label_dims = ctx->GetInputDim("Label");
+    PADDLE_ENFORCE(label_dims.size() == 2UL && label_dims[1] == 1UL,
+                   "The Input(Label) should be a 2-D tensor with the 2nd "
+                   "dimensions fixed to 1.");
+    PADDLE_ENFORCE_EQ(
+        emission_dims[0], label_dims[0],
+        "The height of Input(Emission) and the height of Input(Label) "
+        "should be the same.");
+
+    ctx->SetOutputDim("Alpha", emission_dims);
+    ctx->SetOutputDim("EmissionExps", emission_dims);
+    ctx->SetOutputDim("TransitionExps", transition_dims);
+    // TODO(caoying) This is tricky. The 1st dimension of Output(LogLikelihood)
+    // is the sequence number in a mini-batch. The dimension set here should be
+    // resized to its correct size in the function Compute. Fix this once we can
+    // get LoD information in the InferShape interface.
+    ctx->SetOutputDim("LogLikelihood", {emission_dims[0], 1});
+  }
+
+ protected:
+  // Explicitly set that the data type of output of the linear_chain_crf
+  // operator is determined by its input "Emission".
+  framework::DataType IndicateDataType(
+      const framework::ExecutionContext& ctx) const override {
+    return framework::ToDataType(ctx.Input<LoDTensor>("Emission")->type());
+  }
+};
+
+class LinearChainCRFGradOp : public framework::OperatorWithKernel {
+ public:
+  using framework::OperatorWithKernel::OperatorWithKernel;
+
+  void InferShape(framework::InferShapeContext* ctx) const override {
+    PADDLE_ENFORCE(ctx->HasInput("EmissionExps"),
+                   "Input(EmissionExps) should be not null.");
+    PADDLE_ENFORCE(ctx->HasInput("TransitionExps"),
+                   "Input(TransitionExps) should be not null.");
+    PADDLE_ENFORCE(ctx->HasInput(framework::GradVarName("LogLikelihood")),
+                   "Input(LogLikelihood@GRAD) shoudl be not null.");
+
+    auto emission_exps_dims = ctx->GetInputDim("EmissionExps");
+    PADDLE_ENFORCE_EQ(emission_exps_dims.size(), 2UL,
+                      "The Input(EmissionExps) should be a 2-D tensor.");
+    PADDLE_ENFORCE(emission_exps_dims[0],
+                   "An empty mini-batch is not allowed.");
+
+    auto transition_exps_dims = ctx->GetInputDim("TransitionExps");
+    PADDLE_ENFORCE_EQ(transition_exps_dims.size(), 2UL,
+                      "The Input(TransitionExps) should be a 2-D tensor.");
+    PADDLE_ENFORCE_EQ(
+        transition_exps_dims[0] - 2, transition_exps_dims[1],
+        "An invalid dimension for the Input(TransitionExps), which should "
+        "be a 2-D tensor with shape [(D + 2) x D].");
+    PADDLE_ENFORCE_EQ(
+        emission_exps_dims[1], transition_exps_dims[1],
+        "The 2nd dimension of the Input(EmissionExps) and the "
+        "Input(TransitionExps) should be equal to the tag number.");
+
+    auto label_dims = ctx->GetInputDim("Label");
+    PADDLE_ENFORCE(label_dims.size() == 2UL && label_dims[1] == 1UL,
+                   "The Input(Label) should be a 2-D tensor with the 2nd "
+                   "dimensions fixed to 1.");
+    PADDLE_ENFORCE_EQ(
+        emission_exps_dims[0], label_dims[0],
+        "The height of Input(EmissionExps) and the height of Input(Label) "
+        "should be the same.");
+
+    if (ctx->HasOutput(framework::GradVarName("Emission"))) {
+      ctx->SetOutputDim(framework::GradVarName("Emission"), emission_exps_dims);
+    }
+    if (ctx->HasOutput(framework::GradVarName("Transition"))) {
+      ctx->SetOutputDim(framework::GradVarName("Transition"),
+                        transition_exps_dims);
+    }
+  }
+
+ protected:
+  // Explicitly set that the data type of output of the linear_chain_crf_grad
+  // operator is determined by its input: gradients of LogLikelihood.
+  framework::DataType IndicateDataType(
+      const framework::ExecutionContext& ctx) const override {
+    return framework::ToDataType(
+        ctx.Input<LoDTensor>(framework::GradVarName("LogLikelihood"))->type());
+  }
+};
+
+}  // namespace operators
+}  // namespace paddle
+
+namespace ops = paddle::operators;
+REGISTER_OP(linear_chain_crf, ops::LinearChainCRFOp, ops::LinearChainCRFOpMaker,
+            linear_chain_crf_grad, ops::LinearChainCRFGradOp);
+REGISTER_OP_CPU_KERNEL(
+    linear_chain_crf,
+    ops::LinearChainCRFOpKernel<paddle::platform::CPUPlace, float>,
+    ops::LinearChainCRFOpKernel<paddle::platform::CPUPlace, double>);
+REGISTER_OP_CPU_KERNEL(
+    linear_chain_crf_grad,
+    ops::LinearChainCRFGradOpKernel<paddle::platform::CPUPlace, float>,
+    ops::LinearChainCRFGradOpKernel<paddle::platform::CPUPlace, double>);

paddle/operators/math/detail/hl_avx_functions.h → paddle/operators/linear_chain_crf_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
+   Licensed under the Apache License, Version 2.0 (the "License");
+   you may not use this file except in compliance with the License.
+   You may obtain a copy of the License at
 
-    http://www.apache.org/licenses/LICENSE-2.0
+http://www.apache.org/licenses/LICENSE-2.0
 
 Unless required by applicable law or agreed to in writing, software
 distributed under the License is distributed on an "AS IS" BASIS,
@@ -12,21 +12,15 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 See the License for the specific language governing permissions and
 limitations under the License. */
 
-#ifndef HL_AVX_FUNCTIONS_H_
-#define HL_AVX_FUNCTIONS_H_
+#include "paddle/operators/linear_chain_crf_op.h"
 
-#include <immintrin.h>
+namespace ops = paddle::operators;
 
-namespace hppl {
-__m256 relu(const __m256 a);
-__m256 sigmoid(const __m256 a);
-__m256 tanh(const __m256 a);
-__m256 linear(const __m256 a);
-
-__m256 relu(const __m256 a, const __m256 b);
-__m256 sigmoid(const __m256 a, const __m256 b);
-__m256 tanh(const __m256 a, const __m256 b);
-__m256 linear(const __m256 a, const __m256 b);
-}  // namespace hppl
-
-#endif  // HL_AVX_FUNCTIONS_H_
+REGISTER_OP_GPU_KERNEL(
+    linear_chain_crf,
+    ops::LinearChainCRFOpKernel<paddle::platform::GPUPlace, float>,
+    ops::LinearChainCRFOpKernel<paddle::platform::GPUPlace, double>);
+REGISTER_OP_GPU_KERNEL(
+    linear_chain_crf_grad,
+    ops::LinearChainCRFGradOpKernel<paddle::platform::GPUPlace, float>,
+    ops::LinearChainCRFGradOpKernel<paddle::platform::GPUPlace, double>);

paddle/operators/lookup_table_op.cc

@@ -43,7 +43,7 @@ class LookupTableOp : public framework::OperatorWithKernel {
  protected:
   framework::DataType IndicateDataType(
       const framework::ExecutionContext& ctx) const override {
-    return framework::ToDataType(ctx.Input<Tensor>("W")->type());
+    return framework::ToDataType(ctx.Input<LoDTensor>("W")->type());
   }
 };
 
@@ -93,7 +93,7 @@ class LookupTableOpGrad : public framework::OperatorWithKernel {
  protected:
   framework::DataType IndicateDataType(
       const framework::ExecutionContext& ctx) const override {
-    return framework::ToDataType(ctx.Input<Tensor>("W")->type());
+    return framework::ToDataType(ctx.Input<LoDTensor>("W")->type());
   }
 };
 

paddle/operators/lookup_table_op.cu

@@ -61,16 +61,16 @@ template <typename T>
 class LookupTableCUDAKernel : public framework::OpKernel<T> {
  public:
   void Compute(const framework::ExecutionContext& context) const override {
-    auto table_t = context.Input<Tensor>("W");
-    auto ids_t = context.Input<Tensor>("Ids");
-    auto output_t = context.Output<Tensor>("Out");
+    auto* table_t = context.Input<LoDTensor>("W");
+    auto* ids_t = context.Input<LoDTensor>("Ids");
+    auto* output_t = context.Output<LoDTensor>("Out");
 
     size_t N = table_t->dims()[0];
     size_t D = table_t->dims()[1];
     size_t K = ids_t->numel();
-    auto ids = ids_t->data<int64_t>();
-    auto table = table_t->data<T>();
-    auto output = output_t->mutable_data<T>(context.GetPlace());
+    auto* ids = ids_t->data<int64_t>();
+    auto* table = table_t->data<T>();
+    auto* output = output_t->mutable_data<T>(context.GetPlace());
 
     dim3 threads(128, 8);
     dim3 grids(8, 1);
@@ -87,9 +87,9 @@ class LookupTableGradCUDAKernel : public framework::OpKernel<T> {
   void Compute(const framework::ExecutionContext& context) const override {
     bool is_sparse = context.Attr<bool>("is_sparse");
     if (is_sparse) {
-      auto* ids = context.Input<Tensor>("Ids");
-      auto* table = context.Input<Tensor>("W");
-      auto* d_output = context.Input<Tensor>(framework::GradVarName("Out"));
+      auto* ids = context.Input<LoDTensor>("Ids");
+      auto* table = context.Input<LoDTensor>("W");
+      auto* d_output = context.Input<LoDTensor>(framework::GradVarName("Out"));
       auto* d_table = context.Output<SelectedRows>(framework::GradVarName("W"));
 
       auto* ids_data = ids->data<int64_t>();
@@ -116,12 +116,12 @@ class LookupTableGradCUDAKernel : public framework::OpKernel<T> {
       auto* d_output_data = d_output->data<T>();
       PADDLE_ENFORCE_EQ(d_table_value->dims(), d_output->dims());
       memory::Copy(gpu_place, d_table_data, gpu_place, d_output_data,
-                   d_output->numel(), stream);
+                   d_output->numel() * sizeof(T), stream);
 
     } else {
-      auto ids_t = context.Input<Tensor>("Ids");
-      auto d_output_t = context.Input<Tensor>(framework::GradVarName("Out"));
-      auto d_table_t = context.Output<Tensor>(framework::GradVarName("W"));
+      auto ids_t = context.Input<LoDTensor>("Ids");
+      auto d_output_t = context.Input<LoDTensor>(framework::GradVarName("Out"));
+      auto d_table_t = context.Output<LoDTensor>(framework::GradVarName("W"));
 
       int N = d_table_t->dims()[0];
       int D = d_table_t->dims()[1];

paddle/operators/lookup_table_op.h

@@ -19,22 +19,22 @@
 namespace paddle {
 namespace operators {
 
-using Tensor = framework::Tensor;
+using LoDTensor = framework::LoDTensor;
 using SelectedRows = framework::SelectedRows;
 
 template <typename T>
 class LookupTableKernel : public framework::OpKernel<T> {
  public:
   void Compute(const framework::ExecutionContext& context) const override {
-    auto table_t = context.Input<Tensor>("W");      // float tensor
-    auto ids_t = context.Input<Tensor>("Ids");      // int tensor
-    auto output_t = context.Output<Tensor>("Out");  // float tensor
+    auto* table_t = context.Input<LoDTensor>("W");      // float tensor
+    auto* ids_t = context.Input<LoDTensor>("Ids");      // int tensor
+    auto* output_t = context.Output<LoDTensor>("Out");  // float tensor
 
     int N = table_t->dims()[0];
     int D = table_t->dims()[1];
-    auto ids = ids_t->data<int64_t>();
-    auto table = table_t->data<T>();
-    auto output = output_t->mutable_data<T>(context.GetPlace());
+    auto* ids = ids_t->data<int64_t>();
+    auto* table = table_t->data<T>();
+    auto* output = output_t->mutable_data<T>(context.GetPlace());
     for (int64_t i = 0; i < ids_t->numel(); ++i) {
       PADDLE_ENFORCE_LT(ids[i], N);
       PADDLE_ENFORCE_GE(ids[i], 0);
@@ -49,9 +49,9 @@ class LookupTableGradKernel : public framework::OpKernel<T> {
   void Compute(const framework::ExecutionContext& context) const override {
     bool is_sparse = context.Attr<bool>("is_sparse");
     if (is_sparse) {
-      auto* ids = context.Input<Tensor>("Ids");
-      auto* table = context.Input<Tensor>("W");
-      auto* d_output = context.Input<Tensor>(framework::GradVarName("Out"));
+      auto* ids = context.Input<LoDTensor>("Ids");
+      auto* table = context.Input<LoDTensor>("W");
+      auto* d_output = context.Input<LoDTensor>(framework::GradVarName("Out"));
       auto* d_table = context.Output<SelectedRows>(framework::GradVarName("W"));
 
       auto* ids_data = ids->data<int64_t>();
@@ -76,10 +76,10 @@ class LookupTableGradKernel : public framework::OpKernel<T> {
       PADDLE_ENFORCE_EQ(d_table_value->dims(), d_output->dims());
       memcpy(d_table_data, d_output_data, sizeof(T) * d_output->numel());
     } else {
-      auto* ids = context.Input<Tensor>("Ids");
-      auto* d_output = context.Input<Tensor>(framework::GradVarName("Out"));
-      auto* d_table = context.Output<Tensor>(framework::GradVarName("W"));
-      auto* table = context.Input<Tensor>("W");
+      auto* ids = context.Input<LoDTensor>("Ids");
+      auto* d_output = context.Input<LoDTensor>(framework::GradVarName("Out"));
+      auto* d_table = context.Output<LoDTensor>(framework::GradVarName("W"));
+      auto* table = context.Input<LoDTensor>("W");
 
       auto* ids_data = ids->data<int64_t>();
       auto ids_dim = ids->dims();
@@ -90,11 +90,13 @@ class LookupTableGradKernel : public framework::OpKernel<T> {
       auto* d_output_data = d_output->data<T>();
       auto* d_table_data = d_table->mutable_data<T>(context.GetPlace());
 
+      memset(d_table_data, 0, d_table->numel() * sizeof(T));
+
       for (int64_t i = 0; i < ids->numel(); ++i) {
         PADDLE_ENFORCE_LT(ids_data[i], N);
         PADDLE_ENFORCE_GE(ids_data[i], 0);
         for (int j = 0; j < D; ++j) {
-          d_table_data[ids_data[i] * D + j] = d_output_data[i * D + j];
+          d_table_data[ids_data[i] * D + j] += d_output_data[i * D + j];
         }
       }
     }

paddle/operators/lstm_op.cc

@@ -21,7 +21,6 @@ class LSTMOp : public framework::OperatorWithKernel {
  public:
   using framework::OperatorWithKernel::OperatorWithKernel;
 
- protected:
   void InferShape(framework::InferShapeContext* ctx) const override {
     PADDLE_ENFORCE(ctx->HasInput("Input"),
                    "Input(Input) of LSTM should not be null.");
@@ -29,9 +28,13 @@ class LSTMOp : public framework::OperatorWithKernel {
                    "Output(Hidden) of LSTM should not be null.");
     PADDLE_ENFORCE(ctx->HasOutput("Cell"),
                    "Output(Cell) of LSTM should not be null.");
+    PADDLE_ENFORCE(ctx->HasOutput("BatchGate"),
+                   "Output(BatchGate) of LSTM should not be null.");
+    PADDLE_ENFORCE(ctx->HasOutput("BatchCellPreAct"),
+                   "Output(BatchGate) of LSTM should not be null.");
 
-    auto x_dims = ctx->GetInputDim("Input");
-    PADDLE_ENFORCE_EQ(x_dims.size(), 2, "Input(X)'s rank must be 2.");
+    auto in_dims = ctx->GetInputDim("Input");
+    PADDLE_ENFORCE_EQ(in_dims.size(), 2, "Input(X)'s rank must be 2.");
 
     if (ctx->HasInput("H0")) {
       PADDLE_ENFORCE(ctx->HasInput("C0"),
@@ -44,7 +47,7 @@ class LSTMOp : public framework::OperatorWithKernel {
                      "should be the same.");
     }
 
-    int frame_size = x_dims[1] / 4;
+    int frame_size = in_dims[1] / 4;
     auto w_dims = ctx->GetInputDim("Weight");
     PADDLE_ENFORCE_EQ(w_dims.size(), 2,
                       "The rank of Input(Weight) should be 2.");
@@ -71,12 +74,21 @@ class LSTMOp : public framework::OperatorWithKernel {
                         "4 * %d if disable peepholes connection",
                         frame_size);
     }
-    ctx->SetOutputDim("Hidden", {x_dims[0], frame_size});
-    ctx->SetOutputDim("Cell", {x_dims[0], frame_size});
-    ctx->SetOutputDim("BatchGate", x_dims);
+    framework::DDim out_dims({in_dims[0], frame_size});
+    ctx->SetOutputDim("Hidden", out_dims);
+    ctx->SetOutputDim("Cell", out_dims);
+    ctx->SetOutputDim("BatchGate", in_dims);
+    ctx->SetOutputDim("BatchCellPreAct", out_dims);
     ctx->ShareLoD("Input", "Hidden");
     ctx->ShareLoD("Input", "Cell");
   }
+
+ protected:
+  framework::DataType IndicateDataType(
+      const framework::ExecutionContext& ctx) const override {
+    return framework::ToDataType(
+        ctx.Input<framework::LoDTensor>("Input")->type());
+  }
 };
 
 class LSTMOpMaker : public framework::OpProtoAndCheckerMaker {
@@ -86,16 +98,18 @@ class LSTMOpMaker : public framework::OpProtoAndCheckerMaker {
     AddInput("Input",
              "(LoDTensor) the first input is a LodTensor, which support "
              "variable-time length input sequence. The underlying tensor in "
-             "this LoDTensor is a matrix with shape (T X 4D), where, T is the "
+             "this LoDTensor is a matrix with shape (T X 4D), where T is the "
              "total time steps in this mini-batch, D is the hidden size.");
     AddInput("H0",
              "(Tensor, optional) the initial hidden state is an optional "
              "input. This is a tensor with shape (N x D), where N is the "
-             "batch size, D is the hidden size.");
+             "batch size, D is the hidden size.")
+        .AsDispensable();
     AddInput("C0",
              "(Tensor, optional) the initial cell state is an optional "
              "input. This is a tensor with shape (N x D), where N is the "
-             "batch size. `H0` and `C0` can be NULL but only at the same time");
+             "batch size. `H0` and `C0` can be NULL but only at the same time")
+        .AsDispensable();
     AddInput("Weight",
              "(Tensor) the learnable hidden-hidden weights."
              " - The shape is (D x 4D), where D is the hidden size. "
@@ -109,22 +123,27 @@ class LSTMOpMaker : public framework::OpProtoAndCheckerMaker {
              " - Bias = {b_c, b_i, b_f, b_o}."
              "2. `usePeepholes = True` "
              " - The shape is (1 x 7D). "
-             " - Bias = {b_c, b_i, b_f, b_o, W_ic, W_fc, W_oc}.");
+             " - Bias = {b_c, b_i, b_f, b_o, W_ic, W_fc, W_oc}.")
+        .AsDispensable();
+    AddOutput("Hidden",
+              "(LoDTensor) the hidden state of LSTM operator. "
+              "The shape is (T x D), and lod is the same with the `Input`.");
+    AddOutput("Cell",
+              "(LoDTensor) the cell state of LSTM operator. "
+              "The shape is (T x D), and lod is the same with the `Input`.");
     AddOutput("BatchGate",
               "(LoDTensor) This LoDTensor contains input gate, forget gate "
               "and output gate after the nonlinear computation. This "
               "LoDTensor has the same shape with the reorganized input, which "
-              "was also be called batch input. The LoD size is 2. The first "
+              "is also be called batch input. The LoD size is 2. The first "
               "LoD is the batch offsets and the second LoD contains the "
               "indexes, which denote the position of reorganized sequence "
               "in the raw input.")
         .AsIntermediate();
-    AddOutput("Hidden",
-              "(LoDTensor) the hidden state lod tensor of LSTM operator. "
-              "The shape and lod is the same with the `Input`.");
-    AddOutput("Cell",
-              "(LoDTensor) the cell state lod tensor of LSTM operator. "
-              "The shape and lod is the same with the `Input`.");
+    AddOutput("BatchCellPreAct",
+              "(LoDTensor) This LoDTensor is got in the forward and used "
+              "in the backward.")
+        .AsIntermediate();
     AddAttr<bool>("usePeepholes",
                   "(bool, defalut: True) "
                   "whether to enable diagonal/peephole connections.")
@@ -202,15 +221,37 @@ class LSTMGradOp : public framework::OperatorWithKernel {
  public:
   using framework::OperatorWithKernel::OperatorWithKernel;
 
- protected:
   void InferShape(framework::InferShapeContext* ctx) const override {
-    PADDLE_ENFORCE(ctx->HasInput(framework::GradVarName("Hidden")),
-                   "Input(Hidden@GRAD) should not be null");
-    PADDLE_ENFORCE(ctx->HasInput(framework::GradVarName("Cell")),
-                   "Input(Cell@GRAD) should not be null");
-    ctx->SetOutputDim(framework::GradVarName("Weight"),
-                      ctx->GetInputDim("Weight"));
-    ctx->SetOutputDim(framework::GradVarName("Bias"), ctx->GetInputDim("Bias"));
+    PADDLE_ENFORCE(ctx->HasInput("Input"),
+                   "Input(Input) of LSTM should not be null.");
+    PADDLE_ENFORCE(ctx->HasInput("Hidden"),
+                   "Input(Hidden) of LSTM should not be null.");
+    PADDLE_ENFORCE(ctx->HasInput("Cell"),
+                   "Input(Cell) of LSTM should not be null.");
+
+    PADDLE_ENFORCE(ctx->HasInput("BatchGate"),
+                   "Input(BatchGate) of LSTM should not be null.");
+    PADDLE_ENFORCE(ctx->HasInput("BatchCellPreAct"),
+                   "Input(BatchGate) of LSTM should not be null.");
+
+    auto in_g_name = framework::GradVarName("Input");
+    if (ctx->HasOutput(in_g_name))
+      ctx->SetOutputDim(in_g_name, ctx->GetInputDim("Input"));
+
+    auto w_g_name = framework::GradVarName("Weight");
+    if (ctx->HasOutput(w_g_name))
+      ctx->SetOutputDim(w_g_name, ctx->GetInputDim("Weight"));
+
+    auto b_g_name = framework::GradVarName("Bias");
+    if (ctx->HasOutput(b_g_name))
+      ctx->SetOutputDim(b_g_name, ctx->GetInputDim("Bias"));
+  }
+
+ protected:
+  framework::DataType IndicateDataType(
+      const framework::ExecutionContext& ctx) const override {
+    return framework::ToDataType(
+        ctx.Input<framework::LoDTensor>("Input")->type());
   }
 };
 

paddle/operators/lstm_op.h

@@ -21,8 +21,9 @@ limitations under the License. */
 namespace paddle {
 namespace operators {
 
-using framework::LoDTensor;
-using framework::Tensor;
+using LoDTensor = framework::LoDTensor;
+using Tensor = framework::Tensor;
+
 template <typename T, int MajorType = Eigen::RowMajor,
           typename IndexType = Eigen::DenseIndex>
 using EigenMatrix = framework::EigenMatrix<T, MajorType, IndexType>;
@@ -31,15 +32,15 @@ template <typename Place, typename T>
 class LSTMKernel : public framework::OpKernel<T> {
  public:
   void Compute(const framework::ExecutionContext& ctx) const override {
-    auto* input = ctx.Input<framework::LoDTensor>("Input");
-    auto* weight = ctx.Input<framework::Tensor>("Weight");
-    auto* bias = ctx.Input<framework::Tensor>("Bias");
+    auto* input = ctx.Input<LoDTensor>("Input");
+    auto* weight = ctx.Input<Tensor>("Weight");
+    auto* bias = ctx.Input<Tensor>("Bias");
 
-    auto* batch_gate = ctx.Output<framework::LoDTensor>("BatchGate");
+    auto* batch_gate = ctx.Output<LoDTensor>("BatchGate");
     batch_gate->mutable_data<T>(ctx.GetPlace());
-    auto* hidden_out = ctx.Output<framework::LoDTensor>("Hidden");
+    auto* hidden_out = ctx.Output<LoDTensor>("Hidden");
     hidden_out->mutable_data<T>(ctx.GetPlace());
-    auto* cell_out = ctx.Output<framework::LoDTensor>("Cell");
+    auto* cell_out = ctx.Output<LoDTensor>("Cell");
     cell_out->mutable_data<T>(ctx.GetPlace());
 
     // Now the function ShareLoD in InferShape is not implemented.
@@ -49,7 +50,8 @@ class LSTMKernel : public framework::OpKernel<T> {
 
     bool is_reverse = ctx.Attr<bool>("isReverse");
     math::LoDTensor2BatchFunctor<Place, T> to_batch;
-    to_batch(ctx.device_context(), *input, *batch_gate, is_reverse);
+    auto& device_ctx = ctx.device_context();
+    to_batch(device_ctx, *input, *batch_gate, true, is_reverse);
 
     auto in_dims = input->dims();
     int frame_size = static_cast<int>(in_dims[1] / 4);
@@ -69,17 +71,26 @@ class LSTMKernel : public framework::OpKernel<T> {
     }
 
     math::LstmMetaValue<T> lstm_value;
-    T* bias_data = const_cast<T*>(bias->data<T>());
-    // the code style in LstmMetaValue will be updated later.
-    lstm_value.checkIg = bias_data + 4 * frame_size;
-    lstm_value.checkFg = lstm_value.checkIg + frame_size;
-    lstm_value.checkOg = lstm_value.checkFg + frame_size;
+    if (bias) {
+      T* bias_data = const_cast<T*>(bias->data<T>());
+      // the code style in LstmMetaValue will be updated later.
+
+      lstm_value.checkIg = bias_data + 4 * frame_size;
+      lstm_value.checkFg = lstm_value.checkIg + frame_size;
+      lstm_value.checkOg = lstm_value.checkFg + frame_size;
+    } else {
+      lstm_value.checkIg = nullptr;
+      lstm_value.checkFg = nullptr;
+      lstm_value.checkOg = nullptr;
+    }
     lstm_value.prevStateValue = nullptr;
 
-    framework::LoDTensor batch_out, batch_cell, batch_cell_pre_act;
-    batch_out.mutable_data<T>(dims, ctx.GetPlace());
+    // Use the local variable as here.
+    LoDTensor batch_hidden, batch_cell;
+    auto* batch_cell_pre_act = ctx.Output<LoDTensor>("BatchCellPreAct");
+    batch_hidden.mutable_data<T>(dims, ctx.GetPlace());
     batch_cell.mutable_data<T>(dims, ctx.GetPlace());
-    batch_cell_pre_act.mutable_data<T>(dims, ctx.GetPlace());
+    batch_cell_pre_act->mutable_data<T>(dims, ctx.GetPlace());
 
     auto batch_starts = batch_gate->lod()[0];
     size_t num_batch = batch_starts.size() - 1;
@@ -92,18 +103,18 @@ class LSTMKernel : public framework::OpKernel<T> {
       int bend = static_cast<int>(batch_starts[n + 1]);
 
       Tensor gate_t = batch_gate->Slice(bstart, bend);
-      Tensor out_t = batch_out.Slice(bstart, bend);
+      Tensor out_t = batch_hidden.Slice(bstart, bend);
       Tensor cell_t = batch_cell.Slice(bstart, bend);
-      Tensor cell_pre_act_t = batch_cell_pre_act.Slice(bstart, bend);
+      Tensor cell_pre_act_t = batch_cell_pre_act->Slice(bstart, bend);
 
       int cur_batch_size = bend - bstart;
 
       if (n != 0) {
         int pre_h_start = static_cast<int>(batch_starts[n - 1]);
         int pre_h_end = pre_h_start + cur_batch_size;
-        auto pre_hidden_t = batch_out.Slice(pre_h_start, pre_h_end);
-        math::matmul<Place, T>(ctx.device_context(), pre_hidden_t, false,
-                               *weight, false, static_cast<T>(1.0), &gate_t,
+        auto pre_hidden_t = batch_hidden.Slice(pre_h_start, pre_h_end);
+        math::matmul<Place, T>(device_ctx, pre_hidden_t, false, *weight, false,
+                               static_cast<T>(1.0), &gate_t,
                                static_cast<T>(1.0));
       }
       // else if : FIXME support the initial hidden and cell
@@ -112,27 +123,186 @@ class LSTMKernel : public framework::OpKernel<T> {
       lstm_value.outputValue = out_t.data<T>();
       lstm_value.stateValue = cell_t.data<T>();
       lstm_value.stateActiveValue = cell_pre_act_t.data<T>();
-      math::LstmUnitFunctor<Place, T>::compute(ctx.device_context(), lstm_value,
+      math::LstmUnitFunctor<Place, T>::compute(device_ctx, lstm_value,
                                                frame_size, cur_batch_size,
                                                gate_act, cell_act, cand_act);
       lstm_value.prevStateValue = lstm_value.stateValue;
     }
 
     math::Batch2LoDTensorFunctor<Place, T> to_seq;
-    batch_out.set_lod(batch_gate->lod());
+    batch_hidden.set_lod(batch_gate->lod());
     // restore the output hidden in LoDTensor from the batch hidden
-    to_seq(ctx.device_context(), batch_out, *hidden_out);
+    to_seq(device_ctx, batch_hidden, *hidden_out);
 
     batch_cell.set_lod(batch_gate->lod());
     // restore the output cell state in LoDTensor from the batch cell
-    to_seq(ctx.device_context(), batch_cell, *cell_out);
+    to_seq(device_ctx, batch_cell, *cell_out);
   }
 };
 
 template <typename Place, typename T>
 class LSTMGradKernel : public framework::OpKernel<T> {
  public:
-  void Compute(const framework::ExecutionContext& ctx) const override {}
+  void Compute(const framework::ExecutionContext& ctx) const override {
+    auto* input = ctx.Input<LoDTensor>("Input");
+    auto* weight = ctx.Input<Tensor>("Weight");
+    auto* bias = ctx.Input<Tensor>("Bias");
+
+    auto* hidden_out = ctx.Input<LoDTensor>("Hidden");
+    auto* cell_out = ctx.Input<LoDTensor>("Cell");
+
+    auto* batch_gate = ctx.Input<LoDTensor>("BatchGate");
+    auto* batch_cell_pre_act = ctx.Input<LoDTensor>("BatchCellPreAct");
+
+    auto* hidden_g = ctx.Input<LoDTensor>(framework::GradVarName("Hidden"));
+
+    auto* in_g = ctx.Output<LoDTensor>(framework::GradVarName("Input"));
+    auto* weight_g = ctx.Output<Tensor>(framework::GradVarName("Weight"));
+    auto* bias_g = ctx.Output<Tensor>(framework::GradVarName("Bias"));
+
+    auto& device_ctx = ctx.device_context();
+    math::SetConstant<Place, T> zero;
+    if (weight_g) {
+      weight_g->mutable_data<T>(ctx.GetPlace());
+      zero(device_ctx, weight_g, static_cast<T>(0.0));
+    }
+
+    auto in_dims = input->dims();
+    auto out_dims = hidden_g->dims();
+    int frame_size = static_cast<int>(in_dims[1] / 4);
+    PADDLE_ENFORCE_EQ(frame_size, out_dims[1]);
+
+    math::LstmMetaValue<T> lstm_value;
+    if (bias) {
+      T* bias_data = const_cast<T*>(bias->data<T>());
+      lstm_value.checkIg = bias_data + 4 * frame_size;
+      lstm_value.checkFg = lstm_value.checkIg + frame_size;
+      lstm_value.checkOg = lstm_value.checkFg + frame_size;
+    } else {
+      lstm_value.checkIg = nullptr;
+      lstm_value.checkFg = nullptr;
+      lstm_value.checkOg = nullptr;
+    }
+
+    math::LstmMetaGrad<T> lstm_grad;
+    if (bias && bias_g) {
+      T* bias_g_data = const_cast<T*>(bias_g->mutable_data<T>(ctx.GetPlace()));
+      zero(device_ctx, bias_g, static_cast<T>(0.0));
+      lstm_grad.checkIgGrad = bias_g_data + 4 * frame_size;
+      lstm_grad.checkFgGrad = lstm_grad.checkIgGrad + frame_size;
+      lstm_grad.checkOgGrad = lstm_grad.checkFgGrad + frame_size;
+    } else {
+      lstm_grad.checkIgGrad = nullptr;
+      lstm_grad.checkFgGrad = nullptr;
+      lstm_grad.checkOgGrad = nullptr;
+    }
+
+    math::LoDTensor2BatchFunctor<Place, T> to_batch;
+
+    // use the local variable as here.
+    LoDTensor batch_hidden;
+    batch_hidden.mutable_data<T>(out_dims, ctx.GetPlace());
+    batch_hidden.set_lod(batch_gate->lod());
+    to_batch(device_ctx, *hidden_out, batch_hidden, false);
+
+    LoDTensor batch_hidden_g;
+    batch_hidden_g.mutable_data<T>(out_dims, ctx.GetPlace());
+    batch_hidden_g.set_lod(batch_gate->lod());
+    to_batch(device_ctx, *hidden_g, batch_hidden_g, false);
+
+    LoDTensor batch_cell;
+    batch_cell.mutable_data<T>(out_dims, ctx.GetPlace());
+    batch_cell.set_lod(batch_gate->lod());
+    to_batch(device_ctx, *cell_out, batch_cell, false);
+
+    LoDTensor batch_cell_g;
+    batch_cell_g.mutable_data<T>(out_dims, ctx.GetPlace());
+    batch_cell_g.set_lod(batch_gate->lod());
+    // TODO(qingqing) support the case output cell has gradient.
+    // to_batch(device_ctx, *cell_g, batch_cell_g, false);
+    zero(device_ctx, &batch_cell_g, static_cast<T>(0.0));
+
+    LoDTensor batch_gate_g;
+    batch_gate_g.mutable_data<T>(batch_gate->dims(), ctx.GetPlace());
+    batch_gate_g.set_lod(batch_gate->lod());
+
+    auto gate_act = ctx.Attr<std::string>("gateActivation");
+    auto cell_act = ctx.Attr<std::string>("cellActivation");
+    auto cand_act = ctx.Attr<std::string>("candidateActivation");
+
+    auto batch_starts = batch_gate->lod()[0];
+    size_t num_batch = batch_starts.size() - 1;
+    for (int n = static_cast<int>(num_batch) - 1; n >= 0; n--) {
+      int bstart = static_cast<int>(batch_starts[n]);
+      int bend = static_cast<int>(batch_starts[n + 1]);
+
+      Tensor gate = batch_gate->Slice(bstart, bend);
+      Tensor cell = batch_cell.Slice(bstart, bend);
+      Tensor cell_pre_act = batch_cell_pre_act->Slice(bstart, bend);
+      lstm_value.gateValue = gate.data<T>();
+      lstm_value.stateValue = cell.data<T>();
+      lstm_value.stateActiveValue = cell_pre_act.data<T>();
+
+      Tensor out_g = batch_hidden_g.Slice(bstart, bend);
+      Tensor gate_g = batch_gate_g.Slice(bstart, bend);
+      Tensor cell_g = batch_cell_g.Slice(bstart, bend);
+      lstm_grad.stateGrad = cell_g.data<T>();
+      lstm_grad.gateGrad = gate_g.data<T>();
+      lstm_grad.outputGrad = out_g.data<T>();
+
+      if (n) {
+        int bstart_pre = static_cast<int>(batch_starts[n - 1]);
+        Tensor cell_pre = batch_cell.Slice(bstart_pre, bstart);
+        Tensor cell_pre_g = batch_cell_g.Slice(bstart_pre, bstart);
+        lstm_value.prevStateValue = cell_pre.data<T>();
+        lstm_grad.prevStateGrad = cell_pre_g.data<T>();
+      } else {
+        lstm_value.prevStateValue = nullptr;
+        lstm_grad.prevStateGrad = nullptr;
+      }
+
+      int cur_batch_size = bend - bstart;
+      math::LstmUnitGradFunctor<Place, T>::compute(
+          device_ctx, lstm_value, lstm_grad, frame_size, cur_batch_size,
+          gate_act, cell_act, cand_act);
+
+      if (n != 0) {
+        int pre_h_start = static_cast<int>(batch_starts[n - 1]);
+        int pre_h_end = pre_h_start + cur_batch_size;
+        auto pre_hidden_g = batch_hidden_g.Slice(pre_h_start, pre_h_end);
+        math::matmul<Place, T>(device_ctx, gate_g, false, *weight, true,
+                               static_cast<T>(1.0), &pre_hidden_g,
+                               static_cast<T>(1.0));
+        if (weight_g) {
+          /* backward weight */
+          auto pre_hidden = batch_hidden.Slice(pre_h_start, pre_h_end);
+          math::matmul<Place, T>(device_ctx, pre_hidden, true, gate_g, false,
+                                 static_cast<T>(1.0), weight_g,
+                                 static_cast<T>(1.0));
+        }
+      }
+    }
+
+    math::Batch2LoDTensorFunctor<Place, T> to_seq;
+    if (in_g) {
+      /* backward data */
+      in_g->mutable_data<T>(ctx.GetPlace());
+      to_seq(device_ctx, batch_gate_g, *in_g);
+    }
+    if (bias && bias_g) {
+      /* backward bias */
+      int m = static_cast<int>(batch_gate_g.dims()[0]);
+      int n = static_cast<int>(batch_gate_g.dims()[1]);
+
+      Tensor ones;
+      ones.mutable_data<T>({m}, ctx.GetPlace());
+      math::SetConstant<Place, T> set;
+      set(device_ctx, &ones, static_cast<T>(1.0));
+
+      math::gemv<Place, T>(device_ctx, true, m, n, 1., batch_gate_g.data<T>(),
+                           ones.data<T>(), 0., bias_g->data<T>());
+    }
+  }
 };
 
 }  // namespace operators

paddle/operators/lstm_unit_op.cu

@@ -12,6 +12,10 @@
    See the License for the specific language governing permissions and
    limitations under the License. */
 
+/* Acknowledgement: the following code is strongly inspired by
+https://github.com/caffe2/caffe2/blob/master/caffe2/operators/lstm_unit_op_gpu.cu
+*/
+
 #include "paddle/framework/op_registry.h"
 #include "paddle/operators/cross_entropy_op.h"
 #include "paddle/platform/assert.h"

paddle/operators/lstm_unit_op.h

@@ -12,6 +12,10 @@
    See the License for the specific language governing permissions and
    limitations under the License. */
 
+/* Acknowledgement: the following code is strongly inspired by
+https://github.com/caffe2/caffe2/blob/master/caffe2/operators/lstm_unit_op.h
+*/
+
 #pragma once
 #include "glog/logging.h"
 #include "paddle/framework/op_registry.h"

paddle/operators/math/detail/CMakeLists.txt

 if(WITH_AVX)
-    cc_library(activation_functions SRCS hl_cpu_functions.cc hl_avx_functions.cc)
-else()
-    cc_library(activation_functions SRCS hl_cpu_functions.cc)
+    cc_library(activation_functions SRCS avx_functions.cc)
 endif()

paddle/operators/math/math_function.cc

@@ -211,6 +211,26 @@ void batched_gemm<platform::CPUPlace, double>(
 }
 #endif
 
+template <>
+void gemv<platform::CPUPlace, float>(const platform::DeviceContext& context,
+                                     const bool trans_a, const int M,
+                                     const int N, const float alpha,
+                                     const float* A, const float* B,
+                                     const float beta, float* C) {
+  CBLAS_TRANSPOSE transA = (trans_a == false) ? CblasNoTrans : CblasTrans;
+  cblas_sgemv(CblasRowMajor, transA, M, N, alpha, A, N, B, 1, beta, C, 1);
+}
+
+template <>
+void gemv<platform::CPUPlace, double>(const platform::DeviceContext& context,
+                                      const bool trans_a, const int M,
+                                      const int N, const double alpha,
+                                      const double* A, const double* B,
+                                      const double beta, double* C) {
+  CBLAS_TRANSPOSE transA = (trans_a == false) ? CblasNoTrans : CblasTrans;
+  cblas_dgemv(CblasRowMajor, transA, M, N, alpha, A, N, B, 1, beta, C, 1);
+}
+
 template struct SetConstant<platform::CPUPlace, float>;
 
 }  // namespace math