diff --git a/.travis.yml b/.travis.yml
index d0e2696f100e55f320e410afd6a3038db647f76f..c51e02eb79a9e53a2b8d1d663e8f0c3e0d8c3a61 100644
--- a/.travis.yml
+++ b/.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
diff --git a/benchmark/IntelOptimizedPaddle.md b/benchmark/IntelOptimizedPaddle.md
new file mode 100644
index 0000000000000000000000000000000000000000..040f5ffa41968cbf93a817faa1db86c18956341e
--- /dev/null
+++ b/benchmark/IntelOptimizedPaddle.md
@@ -0,0 +1,48 @@
+# 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
diff --git a/doc/design/images/asgd.gif b/doc/design/images/asgd.gif
new file mode 100644
index 0000000000000000000000000000000000000000..4a0da7bf6df9326a2aab1638b77c5455c18b8c4e
Binary files /dev/null and b/doc/design/images/asgd.gif differ
diff --git a/doc/design/images/theta_star.gif b/doc/design/images/theta_star.gif
new file mode 100644
index 0000000000000000000000000000000000000000..dd24d33e124396be3fc410c9b12f33148f64efe2
Binary files /dev/null and b/doc/design/images/theta_star.gif differ
diff --git a/doc/design/parameter_average.md b/doc/design/parameter_average.md
new file mode 100644
index 0000000000000000000000000000000000000000..2c4edee9fe31d502ea62b9fe5c8757c0a4c5e79f
--- /dev/null
+++ b/doc/design/parameter_average.md
@@ -0,0 +1,72 @@
+# 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 
. The averaging is done as follows:
+
+
+
+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.
diff --git a/doc/faq/parameter/index_cn.rst b/doc/faq/parameter/index_cn.rst
index c721b623183cc7d8d17e2c9fb1635ea07b8970cc..6fa0c64413be1616a435640b0347904a49873349 100644
--- a/doc/faq/parameter/index_cn.rst
+++ b/doc/faq/parameter/index_cn.rst
@@ -75,7 +75,7 @@ PaddlePaddle目前支持8种learning_rate_schedule,这8种learning_rate_schedu
optimizer = paddle.optimizer.Adam(
learning_rate=1e-3,
- learning_rate_schedule="manual",
+ learning_rate_schedule="pass_manual",
learning_rate_args="1:1.0,2:0.9,3:0.8",)
在该示例中,当已训练pass数小于等于1时,学习率为 :code:`1e-3 * 1.0`;当已训练pass数大于1小于等于2时,学习率为 :code:`1e-3 * 0.9`;当已训练pass数大于2时,学习率为 :code:`1e-3 * 0.8`。
diff --git a/doc/howto/cross_compiling/cross_compiling_for_android.md b/doc/howto/cross_compiling/cross_compiling_for_android.md
new file mode 100644
index 0000000000000000000000000000000000000000..161863e5c0a2c002af7d7611dad53c2c19148722
--- /dev/null
+++ b/doc/howto/cross_compiling/cross_compiling_for_android.md
@@ -0,0 +1,153 @@
+# 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.
diff --git a/doc/howto/cross_compiling/cross_compiling_for_android_cn.md b/doc/howto/cross_compiling/cross_compiling_for_android_cn.md
index 1fc58c37cc9151d5e4d99b939e30c29aa99e04f1..58e4dd9c3fe43f963d00152aa4f456fadbb12bf3 100644
--- a/doc/howto/cross_compiling/cross_compiling_for_android_cn.md
+++ b/doc/howto/cross_compiling/cross_compiling_for_android_cn.md
@@ -1,7 +1,7 @@
# 构建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进行矩阵计算
### 编译和安装
diff --git a/paddle/cuda/include/hl_matrix.h b/paddle/cuda/include/hl_matrix.h
index c7f25109972195fb56b9e96c4b68d952363e6338..7daca18761b80eac0f876b21377a6ccc6a853485 100644
--- a/paddle/cuda/include/hl_matrix.h
+++ b/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_ */
diff --git a/paddle/cuda/include/stub/hl_matrix_stub.h b/paddle/cuda/include/stub/hl_matrix_stub.h
index 6ac332945c8f09fef23f35680ba5bb1d9ba9f4fd..46e77e140768dd80fd327dd4eb3b0f62a3370950 100644
--- a/paddle/cuda/include/stub/hl_matrix_stub.h
+++ b/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_
diff --git a/paddle/cuda/src/hl_cuda_matrix.cu b/paddle/cuda/src/hl_cuda_matrix.cu
index b41a3a1e06db7b2566acef19ce430645f79d486d..607efb4f6b0aa0d22a2789397b8743f7a5271d5b 100644
--- a/paddle/cuda/src/hl_cuda_matrix.cu
+++ b/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");
+}
diff --git a/paddle/framework/CMakeLists.txt b/paddle/framework/CMakeLists.txt
index f4fef055daf39e9be0645deaafdad4132fc7e35f..2be21e825ae1b028eefe820e4e152a0666d67f10 100644
--- a/paddle/framework/CMakeLists.txt
+++ b/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)
diff --git a/paddle/framework/backward.cc b/paddle/framework/backward.cc
index dbd5a14f9f3b681f0b77b9bd507b34edfaa78766..ed94540c268e5ed990c1d92859c6a2093c052868 100644
--- a/paddle/framework/backward.cc
+++ b/paddle/framework/backward.cc
@@ -24,7 +24,6 @@
#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 {
@@ -38,7 +37,7 @@ static inline std::unique_ptr 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> grad_ops;
grad_ops.reserve(grad_descs.size());
std::transform(grad_descs.begin(), grad_descs.end(),
@@ -220,19 +219,7 @@ static std::unique_ptr BackwardRecursive(
});
// process recurrent gradient op as a special operator.
- if (forwardOp.Type() == "recurrent") {
- // NOTE clean up cycle call somewhere (RNN's stepnet constains itself),
- // or this will result in infinite loop.
- const auto& rnnop =
- *static_cast(&forwardOp);
- auto rnn_grad_op =
- static_cast(grad_op.get());
- const auto& stepnet_op =
- *static_cast(&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 =
@@ -331,7 +318,7 @@ static void CreateGradVarInBlock(
continue;
}
auto pname = FwdName(arg);
- auto* param = block_desc->FindVar(pname);
+ auto* param = block_desc->FindVarRecursive(pname);
auto* grad = block_desc->FindVar(arg);
if (param == nullptr) {
LOG(WARNING) << "Cannot find forward variable of " << arg
@@ -348,7 +335,9 @@ static void CreateGradVarInBlock(
std::vector> MakeOpGrad(
const OpDescBind* op_desc, std::unordered_set* no_grad_vars,
- std::unordered_map* grad_to_var) {
+ std::unordered_map* grad_to_var,
+ const std::vector& grad_block =
+ std::vector()) {
std::vector> grad_op_descs;
// All input gradients of forwarding operator do not need to calculate.
const std::vector& inputs = op_desc->InputArgumentNames();
@@ -364,9 +353,10 @@ std::vector> 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> pending_fill_zeros_ops;
for (auto& desc : grad_op_descs) {
@@ -400,21 +390,20 @@ std::vector> MakeBlockBackward(
std::vector> backward_descs;
for (auto it = op_descs.rbegin(); it != op_descs.rend(); ++it) {
- std::vector> op_grads =
- MakeOpGrad(*it, no_grad_vars, grad_to_var);
+ std::vector> op_grads;
if ((*it)->Type() == "recurrent") {
- PADDLE_ENFORCE_EQ(
- op_grads.size(), static_cast(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) {
diff --git a/paddle/framework/block_desc.h b/paddle/framework/block_desc.h
index 72f77a88a24434fd7d2ed685ac850c88888d6808..26adf6a20ff09483b84f479db08efcf402135053 100644
--- a/paddle/framework/block_desc.h
+++ b/paddle/framework/block_desc.h
@@ -88,6 +88,8 @@ class BlockDescBind {
BlockDesc *Proto();
+ ProgramDescBind *Program() { return this->prog_; }
+
private:
void ClearPBOps();
void ClearPBVars();
diff --git a/paddle/framework/details/op_registry.h b/paddle/framework/details/op_registry.h
index b731840ef2a4b2d5d82b019d28ad6517fa4b7607..f91e0e03410c95f84a65f02beed38b7bbfdcaa86 100644
--- a/paddle/framework/details/op_registry.h
+++ b/paddle/framework/details/op_registry.h
@@ -108,8 +108,9 @@ struct OpInfoFiller {
info->grad_op_maker_ = [](
const OpDescBind& fwd_op,
const std::unordered_set& no_grad_set,
- std::unordered_map* grad_to_var) {
- T maker(fwd_op, no_grad_set, grad_to_var);
+ std::unordered_map* grad_to_var,
+ const std::vector& grad_block) {
+ T maker(fwd_op, no_grad_set, grad_to_var, grad_block);
return maker();
};
}
diff --git a/paddle/framework/executor.cc b/paddle/framework/executor.cc
index 9bf2311dc835c701c9311880b8adba486a7d446c..52fefe4ea30899880cd386587340d691ee97547b 100644
--- a/paddle/framework/executor.cc
+++ b/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& places) {
+Executor::Executor(const std::vector& 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& places) {
}
Executor::~Executor() {
- for (auto& device_context : device_contexts_) {
- delete device_context;
+ if (own_) {
+ for (auto& device_context : device_contexts_) {
+ delete device_context;
+ }
}
}
@@ -66,44 +68,61 @@ static void CreateTensor(Variable* var, VarDesc::VarType var_type) {
var->GetMutable();
} else if (var_type == VarDesc::FETCH_LIST) {
var->GetMutable();
+ } else if (var_type == VarDesc::STEP_SCOPES) {
+ var->GetMutable>();
} 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 ProgramDescBind& 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_LT(block_id, pdesc.Size());
+ PADDLE_ENFORCE_LT(static_cast(block_id), pdesc.Size());
auto& block = pdesc.Block(block_id);
auto& device = device_contexts_[0];
- 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());
+ 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()
- << " locally, which pointer is " << ptr;
+ VLOG(3) << "Create variable " << var->Name() << ", which pointer is "
+ << ptr;
}
}
for (auto& op_desc : block.AllOps()) {
auto op = paddle::framework::OpRegistry::CreateOp(*op_desc);
- op->Run(local_scope, *device);
+ 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
diff --git a/paddle/framework/executor.h b/paddle/framework/executor.h
index c78bfe8f9f07f1324515f0baaca4a94cc0fe844e..b745f4f6474ef688774f4c833a3958942e9aa8cb 100644
--- a/paddle/framework/executor.h
+++ b/paddle/framework/executor.h
@@ -25,6 +25,7 @@ namespace framework {
class Executor {
public:
explicit Executor(const std::vector& places);
+ explicit Executor(const platform::DeviceContext& devices);
~Executor();
/* @Brief
@@ -34,10 +35,11 @@ class Executor {
* ProgramDesc
* Scope
*/
- void Run(const ProgramDescBind&, Scope*, int);
+ void Run(const ProgramDescBind&, Scope*, int, bool create_local_scope = true);
private:
- std::vector device_contexts_;
+ std::vector device_contexts_;
+ bool own_;
};
} // namespace framework
diff --git a/paddle/framework/grad_op_desc_maker.h b/paddle/framework/grad_op_desc_maker.h
index 94944c79b64d38e799df436de874cabc3661e30a..998186e33915a11f2864eb5387d19ed1bfbab51c 100644
--- a/paddle/framework/grad_op_desc_maker.h
+++ b/paddle/framework/grad_op_desc_maker.h
@@ -15,6 +15,7 @@
#pragma once
#include
#include
+#include
#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& no_grad_set,
- std::unordered_map* grad_to_var)
- : fwd_op_(fwd_op), no_grad_set_(no_grad_set), grad_to_var_(grad_to_var) {}
+ std::unordered_map* grad_to_var,
+ const std::vector& grad_block =
+ std::vector())
+ : 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> operator()() const = 0;
@@ -102,6 +108,9 @@ class GradOpDescMakerBase {
const OpDescBind& fwd_op_;
const std::unordered_set& no_grad_set_;
std::unordered_map* grad_to_var_;
+
+ protected:
+ std::vector grad_block_;
};
class SingleGradOpDescMaker : public GradOpDescMakerBase {
diff --git a/paddle/framework/op_desc.cc b/paddle/framework/op_desc.cc
index 0779137639e6cd9f6ecf3bbbc24d081cae3de9c0..c96166f35d1425218a4a74f50dc5ed542d677b68 100644
--- a/paddle/framework/op_desc.cc
+++ b/paddle/framework/op_desc.cc
@@ -327,6 +327,19 @@ void OpDescBind::InferShape(const BlockDescBind &block) const {
PADDLE_ENFORCE(static_cast(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(sout, ", "));
+ sout << "] to [";
+ auto onames = this->OutputArgumentNames();
+ std::copy(onames.begin(), onames.end(),
+ std::ostream_iterator(sout, ", "));
+ sout << "]";
+ VLOG(10) << sout.str();
+ }
infer_shape(&ctx);
}
diff --git a/paddle/framework/operator.cc b/paddle/framework/operator.cc
index aa46829fdde82b58a649108bf708901299cd8153..9295d36c2b2e66130ad273ebd3a40de739efeea7 100644
--- a/paddle/framework/operator.cc
+++ b/paddle/framework/operator.cc
@@ -37,32 +37,32 @@ ExecutionContext::GetEigenDevice() 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& 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& 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 OperatorBase::InputVars() const {
std::vector 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());
}
@@ -394,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 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);
diff --git a/paddle/framework/operator.h b/paddle/framework/operator.h
index 93885fa3028e072bc0bd021ea9287087678f3621..b8a7040ed024fc7b19980beef3d8b367dfdd7f50 100644
--- a/paddle/framework/operator.h
+++ b/paddle/framework/operator.h
@@ -427,7 +427,8 @@ class OperatorWithKernel : public OperatorBase {
int tmp = static_cast(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;
}
}
diff --git a/paddle/framework/scope.cc b/paddle/framework/scope.cc
index 14cc530448379eb6d4bf0435f607494aa01ef5b5..fb2c69105627f663ddcce07d31526c9e4278e863 100644
--- a/paddle/framework/scope.cc
+++ b/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 {
diff --git a/paddle/framework/scope.h b/paddle/framework/scope.h
index ac334da5ef0c8ad563b6be5413df33f5d0bdbcf8..fb660949394149ebf2c6172a0ac3f4c7594f4286 100644
--- a/paddle/framework/scope.h
+++ b/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.
diff --git a/paddle/framework/tensor.h b/paddle/framework/tensor.h
index 7b9a5b75e1087a1cc3b6c6c7a6e4dc185c32dd42..28d0fcf94ec31c82476e093f93ccee222a0c9d9a 100644
--- a/paddle/framework/tensor.h
+++ b/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(
diff --git a/paddle/framework/tensor_impl.h b/paddle/framework/tensor_impl.h
index 29ac683f48fcde4dd3b5ad7f04b5d1d7434706ba..d78a2c4c21149ef3c800991b9a144ea198f1bdcf 100644
--- a/paddle/framework/tensor_impl.h
+++ b/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& 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;
diff --git a/paddle/framework/type_defs.h b/paddle/framework/type_defs.h
index afeeb1914ac30188b93c3b9da30bb5ceaf74416e..baeb98c9bd49ec65da5931bcbe33ab788f86f3e8 100644
--- a/paddle/framework/type_defs.h
+++ b/paddle/framework/type_defs.h
@@ -29,6 +29,7 @@ class OpDescBind;
class BlockDescBind;
class BlockDesc;
class InferShapeContext;
+class BlockDescBind;
using VariableNameMap = std::map>;
@@ -46,7 +47,8 @@ using OpCreator = std::function>(
const OpDescBind&, const std::unordered_set& /*no_grad_set*/,
- std::unordered_map* /*grad_to_var*/)>;
+ std::unordered_map* /*grad_to_var*/,
+ const std::vector& grad_block)>;
using InferVarTypeFN = std::function;
diff --git a/paddle/gserver/evaluators/Evaluator.cpp b/paddle/gserver/evaluators/Evaluator.cpp
index 9db6d252d97bfeee3fe376bcda431fe94c65a678..8e66b1f0db5d8a365a5aa9b98d2fb3f867458411 100644
--- a/paddle/gserver/evaluators/Evaluator.cpp
+++ b/paddle/gserver/evaluators/Evaluator.cpp
@@ -395,14 +395,24 @@ real AucEvaluator::evalImp(std::vector& 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& 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(value * kBinNum_);
diff --git a/paddle/gserver/layers/CRFLayer.cpp b/paddle/gserver/layers/CRFLayer.cpp
index 0b544420097e9150f8489731b6379dea633e992c..867303b4fa0d490297ab152fc2ad266e92e29baf 100644
--- a/paddle/gserver/layers/CRFLayer.cpp
+++ b/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);
diff --git a/paddle/gserver/layers/LinearChainCRF.cpp b/paddle/gserver/layers/LinearChainCRF.cpp
index dc3dc156792bdf32c3b948a292597d0e9eca5d8b..abaa1802b763a49f748214dbd4dec1d2bac53b59 100644
--- a/paddle/gserver/layers/LinearChainCRF.cpp
+++ b/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();
diff --git a/paddle/gserver/tests/MKLDNNTester.cpp b/paddle/gserver/tests/MKLDNNTester.cpp
index 73b7e8857f35d194e71b2b5b341f89b77fd1f8b0..7670cb88fb67dec0ab1d170458d102da166dc7b6 100644
--- a/paddle/gserver/tests/MKLDNNTester.cpp
+++ b/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(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;
diff --git a/paddle/gserver/tests/MKLDNNTester.h b/paddle/gserver/tests/MKLDNNTester.h
index 19d8848f74f2ee4a809e42164a0eb180abd2a4e1..ca55a45bc77b4e171619ab788d7c7dfeefcd036a 100644
--- a/paddle/gserver/tests/MKLDNNTester.h
+++ b/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);
diff --git a/paddle/gserver/tests/mkldnn_branch_net.conf b/paddle/gserver/tests/mkldnn_branch_net.conf
new file mode 100644
index 0000000000000000000000000000000000000000..8d5146abb0ebd7f5d6c512457f3cb5c84eac20f5
--- /dev/null
+++ b/paddle/gserver/tests/mkldnn_branch_net.conf
@@ -0,0 +1,142 @@
+# 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_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())
+
+ 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)
+
+tmp = img_conv_layer(input=data,
+ num_channels=channels,
+ filter_size=3,
+ 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)
+
+tmp = img_pool_layer(input=tmp,
+ pool_size=3,
+ stride=2,
+ padding=1,
+ pool_type=AvgPooling())
+
+b1, b2 = two_conv_pool(tmp, 'pool_branch')
+tmp = concat_layer(input=[b1, b2])
+
+tmp = img_pool_layer(input=tmp,
+ num_channels=channels*2,
+ pool_size=3,
+ stride=2,
+ padding=1,
+ pool_type=MaxPooling())
+
+tmp = img_conv_layer(input=tmp,
+ filter_size=3,
+ num_filters=channels,
+ padding=1,
+ stride=2,
+ shared_biases=True,
+ act=LinearActivation(),
+ bias_attr=False)
+
+tmp = batch_norm_layer(input=tmp,
+ use_global_stats=False,
+ 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=channels,
+ bias_attr=True,
+ act=ReluActivation())
+
+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())
+
+outputs(out)
diff --git a/paddle/gserver/tests/mkldnn_branches_fc.conf b/paddle/gserver/tests/mkldnn_branches_fc.conf
deleted file mode 100644
index fb85425c2b63c7604d636e2b0c5d20d91fb5de1b..0000000000000000000000000000000000000000
--- a/paddle/gserver/tests/mkldnn_branches_fc.conf
+++ /dev/null
@@ -1,58 +0,0 @@
-# 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])
diff --git a/paddle/gserver/tests/mkldnn_branches_pool.conf b/paddle/gserver/tests/mkldnn_branches_pool.conf
deleted file mode 100644
index ca17c74752ab0777a69f818d9f43275a6140cb4c..0000000000000000000000000000000000000000
--- a/paddle/gserver/tests/mkldnn_branches_pool.conf
+++ /dev/null
@@ -1,60 +0,0 @@
-# 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])
diff --git a/paddle/gserver/tests/mkldnn_branches_conv.conf b/paddle/gserver/tests/mkldnn_simple_net.conf
similarity index 64%
rename from paddle/gserver/tests/mkldnn_branches_conv.conf
rename to paddle/gserver/tests/mkldnn_simple_net.conf
index 2628509db43e6a5f69a4f5ea956bffdc2837e32a..8bbe91e56d0ba6da06475ad16f3162ee1103ee02 100644
--- a/paddle/gserver/tests/mkldnn_branches_conv.conf
+++ b/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)
diff --git a/paddle/gserver/tests/test_MKLDNN.cpp b/paddle/gserver/tests/test_MKLDNN.cpp
index 85d4f437c2664135a7975c6ed3270d8f1ddbeaf4..d60b0f04a1613acc3711e711cfe18ced5f0f924d 100644
--- a/paddle/gserver/tests/test_MKLDNN.cpp
+++ b/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 cases = {"conv", "pool", "fc"};
+TEST(MKLDNNNet, net) {
+ std::vector 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);
}
}
}
diff --git a/paddle/math/MKLDNNMatrix.h b/paddle/math/MKLDNNMatrix.h
index 5f5b819017b83579ce58522198b3f13311297d42..54cfefe23b3dc70fd12fd2ca8886c941047b59f7 100644
--- a/paddle/math/MKLDNNMatrix.h
+++ b/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.
diff --git a/paddle/math/Vector.cpp b/paddle/math/Vector.cpp
index ff72672e3ab77212b309fcfea835839a916fa632..346008439c35a2bcbcd2e9dfd36d689e01d7495f 100644
--- a/paddle/math/Vector.cpp
+++ b/paddle/math/Vector.cpp
@@ -18,6 +18,7 @@ limitations under the License. */
#include
#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::toOneHotSparseMatrix(size_t idRange, bool useGpu) {
return mat;
}
+template <>
+std::shared_ptr> VectorT::castToInt() {
+ std::shared_ptr> 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
GpuVectorT::GpuVectorT(size_t size)
: VectorT(size,
diff --git a/paddle/math/Vector.h b/paddle/math/Vector.h
index 80b9775fccf10c57bb48145ef56165ec7c86d8b8..f965a5809209da313c78a545c44e7aa39e95ac65 100644
--- a/paddle/math/Vector.h
+++ b/paddle/math/Vector.h
@@ -162,6 +162,13 @@ public:
*/
std::shared_ptr 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> castToInt();
+
/**
* This function will crash if the size of src and dest is different.
*/
diff --git a/paddle/operators/CMakeLists.txt b/paddle/operators/CMakeLists.txt
index bbc782bb294ab6f433efbae28fdb947a1d94edff..a84ca8e3b821693baa6469462ef466447fae235a 100644
--- a/paddle/operators/CMakeLists.txt
+++ b/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
@@ -143,9 +144,6 @@ set(DEPS_OPS
lstm_op
gru_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)
@@ -158,7 +156,9 @@ endif()
op_library(sequence_conv_op DEPS context_project)
op_library(lstm_op DEPS sequence2batch lstm_compute)
op_library(gru_op DEPS sequence2batch gru_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})
@@ -170,8 +170,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()
diff --git a/paddle/operators/activation_op.cc b/paddle/operators/activation_op.cc
index 90f1535fcd387c34ea39d84d9c2ec78fcbc3c764..483f9888973edc9db6317723c136778d40cc7878 100644
--- a/paddle/operators/activation_op.cc
+++ b/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("alpha", "The small negative slope")
.SetDefault(static_cast(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("lambda", "non-negative offset")
.SetDefault(static_cast(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("threshold", "The value of threshold for HardShrink")
.SetDefault(static_cast(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("t_min", "The min marginal value of BRelu")
.SetDefault(static_cast(0));
AddAttr("t_max", "The max marginal value of BRelu")
.SetDefault(static_cast(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("threshold", "The threshold value of SoftRelu")
.SetDefault(static_cast(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(
- "alpha", "(float, default 1.0) Alpha value in the elu formulation.")
- .SetDefault(static_cast(1.));
+ AddInput("X", "Input of ELU operator");
+ AddOutput("Y", "Output of ELU operator");
+ AddAttr("alpha", "The alpha value of ELU")
+ .SetDefault(static_cast(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("threshold", "The threshold value of Relu6")
.SetDefault(static_cast(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("factor", "The exponential factor of Pow")
.SetDefault(static_cast(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("scale_a", "The scale parameter of a for the input")
.SetDefault(static_cast(2 / 3));
AddAttr("scale_b", "The scale parameter of b for the input")
.SetDefault(static_cast(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("threshold", "The threshold location of activation")
.SetDefault(static_cast(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("slope", "Slope for linear approximation of sigmoid")
+ .SetDefault(static_cast(0.2));
+ AddAttr("offset", "Offset for linear approximation of sigmoid")
+ .SetDefault(static_cast(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("slope", "Slope for linear approximation of sigmoid")
- .SetDefault(static_cast(0.2));
- AddAttr("offset", "Offset for linear approximation of sigmoid")
- .SetDefault(static_cast(0.5));
+)DOC");
}
};
diff --git a/paddle/operators/activation_op.h b/paddle/operators/activation_op.h
index ddd966e26c9abad0d83f8b5c6e3e7d9ad65158a8..ceb4b4e40b67473f42e67e3f02f8e012e1b1eb50 100644
--- a/paddle/operators/activation_op.h
+++ b/paddle/operators/activation_op.h
@@ -232,7 +232,7 @@ struct HardShrinkGradFunctor : public BaseActivationFunctor {
}
};
-// 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
struct SoftShrinkFunctor : public BaseActivationFunctor {
diff --git a/paddle/operators/conv2d_transpose_cudnn_op.cc b/paddle/operators/conv2d_transpose_cudnn_op.cc
new file mode 100644
index 0000000000000000000000000000000000000000..8ce94e0f04f14e1eae7e7d01280601cc72dea8c4
--- /dev/null
+++ b/paddle/operators/conv2d_transpose_cudnn_op.cc
@@ -0,0 +1,50 @@
+/* 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>("dilations", "dilations of convolution operator.")
+ .SetDefault(std::vector{1, 1});
+ AddAttr("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);
+REGISTER_OP_CPU_KERNEL(
+ conv2d_transpose_cudnn_grad,
+ ops::GemmConv2DTransposeGradKernel);
diff --git a/paddle/operators/conv2d_transpose_cudnn_op.cu b/paddle/operators/conv2d_transpose_cudnn_op.cu
new file mode 100644
index 0000000000000000000000000000000000000000..61fcfb3bd8fa57f2c45fbf3a980dbe41041cff18
--- /dev/null
+++ b/paddle/operators/conv2d_transpose_cudnn_op.cu
@@ -0,0 +1,240 @@
+/* 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
+class CudnnConvTransposeOpKernel : public framework::OpKernel {
+ public:
+ void Compute(const framework::ExecutionContext& ctx) const override {
+ PADDLE_ENFORCE(platform::is_gpu_place(ctx.GetPlace()),
+ "It must use GPUPlace.");
+ auto* input = ctx.Input("Input");
+ auto* filter = ctx.Input("Filter");
+ auto* output = ctx.Output("Output");
+
+ std::vector strides = ctx.Attr>("strides");
+ std::vector paddings = ctx.Attr>("paddings");
+ // cudnn v5 does not support dilations
+ std::vector dilations = ctx.Attr>("dilations");
+ int user_workspace_size = ctx.Attr("workspace_size_MB");
+
+ const T* input_data = input->data();
+ const T* filter_data = filter->data();
+ T* output_data = output->mutable_data(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(
+ layout, framework::vectorize2int(input->dims()));
+ // N, C, O_h, O_w
+ cudnnTensorDescriptor_t cudnn_output_desc = output_desc.descriptor(
+ layout, framework::vectorize2int(output->dims()));
+ // M, C, K_h, K_w
+ cudnnFilterDescriptor_t cudnn_filter_desc = filter_desc.descriptor(
+ layout, framework::vectorize2int(filter->dims()));
+ cudnnConvolutionDescriptor_t cudnn_conv_desc =
+ conv_desc.descriptor(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(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
+class CudnnConvTransposeGradOpKernel : public framework::OpKernel {
+ public:
+ void Compute(const framework::ExecutionContext& ctx) const override {
+ PADDLE_ENFORCE(platform::is_gpu_place(ctx.GetPlace()),
+ "It must use GPUPlace.");
+ auto input = ctx.Input("Input");
+ auto filter = ctx.Input("Filter");
+ auto output_grad = ctx.Input(framework::GradVarName("Output"));
+ auto input_grad = ctx.Output(framework::GradVarName("Input"));
+ auto filter_grad = ctx.Output(framework::GradVarName("Filter"));
+ const T* input_data = input->data();
+ const T* output_grad_data = output_grad->data();
+ const T* filter_data = filter->data();
+
+ std::vector strides = ctx.Attr>("strides");
+ std::vector paddings = ctx.Attr>("paddings");
+ // cudnn v5 does not support dilations
+ std::vector dilations = ctx.Attr>("dilations");
+ int user_workspace_size = ctx.Attr("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(
+ layout, framework::vectorize2int(input->dims()));
+ // Output: (N, C, O_H, O_W)
+ cudnnTensorDescriptor_t cudnn_output_desc = output_desc.descriptor(
+ layout, framework::vectorize2int(output_grad->dims()));
+ // Filter (M, C, K_H, K_W)
+ cudnnFilterDescriptor_t cudnn_filter_desc = filter_desc.descriptor(
+ layout, framework::vectorize2int(filter->dims()));
+
+ cudnnConvolutionDescriptor_t cudnn_conv_desc =
+ conv_desc.descriptor(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(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(ctx.GetPlace());
+ auto t = framework::EigenVector::Flatten(*input_grad);
+ t.device(ctx.GetEigenDevice()) =
+ t.constant(static_cast(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(ctx.GetPlace());
+ auto t = framework::EigenVector::Flatten(*filter_grad);
+ t.device(ctx.GetEigenDevice()) =
+ t.constant(static_cast(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);
+REGISTER_OP_GPU_KERNEL(conv2d_transpose_cudnn_grad,
+ ops::CudnnConvTransposeGradOpKernel);
diff --git a/paddle/operators/conv2dtranspose_op.cc b/paddle/operators/conv2d_transpose_op.cc
similarity index 95%
rename from paddle/operators/conv2dtranspose_op.cc
rename to paddle/operators/conv2d_transpose_op.cc
index c1b231906e2f172b6f9cee55f850d1a5ec6c3221..348527728bdd4ed60676d6e6e44c4e761b803096 100644
--- a/paddle/operators/conv2dtranspose_op.cc
+++ b/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);
REGISTER_OP_CPU_KERNEL(
- conv2dtranspose_grad,
+ conv2d_transpose_grad,
ops::GemmConv2DTransposeGradKernel);
diff --git a/paddle/operators/conv2dtranspose_op.cu b/paddle/operators/conv2d_transpose_op.cu
similarity index 89%
rename from paddle/operators/conv2dtranspose_op.cu
rename to paddle/operators/conv2d_transpose_op.cu
index 761bc1959e69be94f43571728e6b92a322558b99..931ac9eed294c4fe7c726d8cc2c4d9a39ec12828 100644
--- a/paddle/operators/conv2dtranspose_op.cu
+++ b/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);
REGISTER_OP_GPU_KERNEL(
- conv2dtranspose_grad,
+ conv2d_transpose_grad,
ops::GemmConv2DTransposeGradKernel);
diff --git a/paddle/operators/conv2dtranspose_op.h b/paddle/operators/conv2d_transpose_op.h
similarity index 99%
rename from paddle/operators/conv2dtranspose_op.h
rename to paddle/operators/conv2d_transpose_op.h
index 8c70b3dcec1e26ab3d8a42d88040764c643b5ae6..cab7788227690621a0e5b744197b86c515bbef72 100644
--- a/paddle/operators/conv2dtranspose_op.h
+++ b/paddle/operators/conv2d_transpose_op.h
@@ -62,7 +62,7 @@ class GemmConv2DTransposeKernel : public framework::OpKernel {
std::vector strides = context.Attr>("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];
diff --git a/paddle/operators/cross_entropy_op.cc b/paddle/operators/cross_entropy_op.cc
index d94b96200c2a5cd112b17e45aa6cd4a63bdd04d0..39df19da677a7dee7d0989d491f8d5511f73a9c7 100644
--- a/paddle/operators/cross_entropy_op.cc
+++ b/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("X")->type());
diff --git a/paddle/operators/fill_constant_batch_size_like_op.cc b/paddle/operators/fill_constant_batch_size_like_op.cc
index 58c9f1cd2c79c150aaed7753641f6ad6120dd0f5..0244adb42392c707d755e95c7abdebd826c219b4 100644
--- a/paddle/operators/fill_constant_batch_size_like_op.cc
+++ b/paddle/operators/fill_constant_batch_size_like_op.cc
@@ -36,7 +36,12 @@ class FillConstantBatchSizeLikeOp : public framework::OperatorWithKernel {
[](int a) { return static_cast(a); });
auto dims = framework::make_ddim(shape_int64);
- dims[0] = ctx->GetInputDim("Input")[0];
+ int dim_idx = ctx->Attrs().Get("dim_idx");
+ PADDLE_ENFORCE_GE(dim_idx, 0);
+ PADDLE_ENFORCE_GT(static_cast(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>("shape", "(vector) The shape of the output");
- AddAttr("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>("shape", "(vector) The shape of the output");
+ AddAttr("dim_idx",
+ "(int, default 0) the index of batch size dimension")
+ .SetDefault(0);
+ AddAttr("value", "(float, default 0) The value to be filled")
+ .SetDefault(0.0f);
AddComment(R"DOC(Fill up a variable with specified constant value.)DOC");
}
};
diff --git a/paddle/operators/linear_chain_crf_op.cc b/paddle/operators/linear_chain_crf_op.cc
new file mode 100644
index 0000000000000000000000000000000000000000..605dbba5af1bb8b0d718833be6af45fdaeac70ac
--- /dev/null
+++ b/paddle/operators/linear_chain_crf_op.cc
@@ -0,0 +1,261 @@
+/* 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). "
+ "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). 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). 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. 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). 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("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(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,
+ ops::LinearChainCRFOpKernel);
+REGISTER_OP_CPU_KERNEL(
+ linear_chain_crf_grad,
+ ops::LinearChainCRFGradOpKernel,
+ ops::LinearChainCRFGradOpKernel);
diff --git a/paddle/operators/linear_chain_crf_op.cu b/paddle/operators/linear_chain_crf_op.cu
new file mode 100644
index 0000000000000000000000000000000000000000..6fc8995f4c2ce05f89ffb58129695113f89159fa
--- /dev/null
+++ b/paddle/operators/linear_chain_crf_op.cu
@@ -0,0 +1,26 @@
+/* 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 ops = paddle::operators;
+
+REGISTER_OP_GPU_KERNEL(
+ linear_chain_crf,
+ ops::LinearChainCRFOpKernel,
+ ops::LinearChainCRFOpKernel);
+REGISTER_OP_GPU_KERNEL(
+ linear_chain_crf_grad,
+ ops::LinearChainCRFGradOpKernel,
+ ops::LinearChainCRFGradOpKernel);
diff --git a/paddle/operators/linear_chain_crf_op.h b/paddle/operators/linear_chain_crf_op.h
new file mode 100644
index 0000000000000000000000000000000000000000..56fb0c9102bee6e2fefd1180ef20237891573f70
--- /dev/null
+++ b/paddle/operators/linear_chain_crf_op.h
@@ -0,0 +1,543 @@
+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+ http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License. */
+
+#pragma once
+#include "paddle/framework/eigen.h"
+#include "paddle/framework/op_registry.h"
+#include "paddle/operators/math/math_function.h"
+
+namespace paddle {
+namespace operators {
+
+template
+static inline T NormalizeL1(T* x, size_t len) {
+ T sum = 0.;
+ for (size_t i = 0; i < len; ++i) sum += x[i];
+ // (This comment is from the old LinearChainCRFLayer.)
+ // Right now, we just bet that sum won't be zero. If this really happens, we
+ // will figure out what should be done then.
+ PADDLE_ENFORCE(sum,
+ "The unnormalized probabilities of all possible unfinished "
+ "sequences must be greater than 0.");
+ T s = 1. / sum;
+ for (size_t i = 0; i < len; ++i) x[i] *= s;
+ return sum;
+}
+
+template
+struct ScalarMul {
+ explicit ScalarMul(const T& scalar) : scalar(scalar) {}
+ T operator()(const T& val) const { return val * scalar; }
+
+ T scalar;
+};
+
+using framework::LoDTensor;
+using framework::LoD;
+using framework::Tensor;
+template
+using EigenMatrix = framework::EigenMatrix;
+
+template
+class LinearChainCRFOpKernel : public framework::OpKernel {
+ public:
+ void Compute(const framework::ExecutionContext& ctx) const override {
+ // TODO(caoying) The checks related to LoD information should be
+ // moved into InferShape once after the InferShape is refactored.
+ PADDLE_ENFORCE_EQ(ctx.Input("Emission")->NumLevels(), 1UL,
+ "The Input(Emission) should be a sequence.");
+ PADDLE_ENFORCE_EQ(ctx.Input("Label")->NumLevels(), 1UL,
+ "The Input(Label) should be a sequence.");
+ auto in_lod = ctx.Input("Label")->lod();
+ PADDLE_ENFORCE(in_lod.size(), "Input(Label) must be a sequence.");
+ const size_t level = 0;
+ const size_t seq_num = in_lod[level].size() - 1;
+
+ // These local variables hold the inputs and outputs, garanteeing them on
+ // CPU memory, to provide a consistent reference.
+ // TODO(caoying) Fix this by moving all these local variables into the
+ // class's data members once we can profile the whole training process.
+ LoDTensor* emission_weights = nullptr;
+ LoDTensor emission_weight_tensor;
+ Tensor* transition_weights = nullptr;
+ Tensor transition_weight_tensor;
+ LoDTensor* label = nullptr;
+ LoDTensor label_tensor;
+
+ Tensor* emission_exps = nullptr;
+ Tensor emission_exps_tensor;
+ Tensor* transition_exps = nullptr;
+ Tensor transition_exps_tensor;
+ Tensor* alpha = nullptr;
+ Tensor alpha_tensor;
+ Tensor* ll = nullptr;
+ Tensor ll_tensor;
+
+ if (platform::is_gpu_place(ctx.GetPlace())) {
+ emission_weights = &emission_weight_tensor;
+ transition_weights = &transition_weight_tensor;
+ label = &label_tensor;
+
+ CopyInputsToCpuMemory(
+ ctx.device_context(), *ctx.Input("Emission"),
+ *ctx.Input("Transition"), *ctx.Input("Label"),
+ emission_weights, transition_weights, label);
+
+ emission_exps = &emission_exps_tensor;
+ emission_exps->Resize(emission_weights->dims());
+
+ transition_exps = &transition_exps_tensor;
+ transition_exps->Resize(transition_weights->dims());
+
+ alpha = &alpha_tensor;
+ alpha->Resize(ctx.Output("Alpha")->dims());
+
+ ll = &ll_tensor;
+ } else {
+ emission_weights =
+ const_cast(ctx.Input("Emission"));
+ transition_weights = const_cast(ctx.Input("Transition"));
+ label = const_cast(ctx.Input("Label"));
+
+ emission_exps = ctx.Output("EmissionExps");
+ transition_exps = ctx.Output("TransitionExps");
+ alpha = ctx.Output("Alpha");
+ ll = ctx.Output("LogLikelihood");
+ }
+
+ // Because the computation codes only runs on CPU, here the memory for all
+ // the outputs is FIXED to be allocated on the CPU memory.
+ emission_exps->mutable_data(platform::CPUPlace());
+ transition_exps->mutable_data(platform::CPUPlace());
+ alpha->mutable_data(platform::CPUPlace());
+
+ // Resize the output tensor to its correct dimension.
+ ll->Resize({static_cast(seq_num), 1});
+ ll->mutable_data(platform::CPUPlace());
+
+ // Now, all the inputs and outputs should be on the CPU memory.
+ auto emission_dims = emission_weights->dims();
+ const size_t batch_size = emission_dims[0];
+ const size_t tag_num = emission_dims[1];
+
+ Tensor emission_row_max;
+ emission_row_max.mutable_data(
+ framework::make_ddim({static_cast(batch_size), 1}),
+ platform::CPUPlace());
+
+ auto place = ctx.GetEigenDevice();
+ auto x = EigenMatrix::From(*emission_weights);
+ auto x_row_max = EigenMatrix::From(emission_row_max);
+ x_row_max.device(place) =
+ x.maximum(Eigen::DSizes(1))
+ .reshape(Eigen::DSizes(int(batch_size), 1));
+
+ auto x_exps = EigenMatrix::From(*emission_exps);
+ x_exps.device(place) =
+ (x - x_row_max.broadcast(Eigen::DSizes(1, tag_num))).exp();
+
+ auto w = EigenMatrix::From(*transition_weights);
+ auto w_exps = EigenMatrix::From(*transition_exps);
+ w_exps.device(place) = w.exp();
+
+ T* log_likelihood = ll->data();
+ for (size_t i = 0; i < seq_num; ++i) {
+ int start_pos = static_cast(in_lod[level][i]);
+ int end_pos = static_cast(in_lod[level][i + 1]);
+ if (end_pos == start_pos) {
+ // If an empty input sequence is given, pad 0 for its cost.
+ log_likelihood[i] = 0.;
+ continue;
+ }
+
+ const Tensor one_seq = emission_weights->Slice(start_pos, end_pos);
+ Tensor one_seq_row_max = emission_row_max.Slice(start_pos, end_pos);
+ Tensor one_seq_exps = emission_exps->Slice(start_pos, end_pos);
+ const Tensor one_seq_label = label->Slice(start_pos, end_pos);
+ Tensor one_seq_alpha = alpha->Slice(start_pos, end_pos);
+
+ log_likelihood[i] = ForwardOneSequence(
+ one_seq, one_seq_row_max, one_seq_exps, *transition_weights,
+ *transition_exps, one_seq_label, &one_seq_alpha);
+ }
+
+ if (platform::is_gpu_place(ctx.GetPlace())) {
+ CopyOutputsToGpuMemory(
+ ctx.device_context(), *emission_exps, *transition_exps, *alpha, *ll,
+ ctx.Output("EmissionExps"),
+ ctx.Output("TransitionExps"), ctx.Output("Alpha"),
+ ctx.Output