Merge branch 'develop' into add_nest_sequence_select

.pre-commit-config.yaml

@@ -24,7 +24,7 @@
         description: Format files with ClangFormat.
         entry: clang-format -i
         language: system
-        files: \.(c|cc|cxx|cpp|h|hpp|hxx)$
+        files: \.(c|cc|cxx|cpp|cu|h|hpp|hxx|proto)$
 -   repo: https://github.com/PaddlePaddle/pre-commit-golang
     sha: 8337620115c25ff8333f1b1a493bd031049bd7c0
     hooks:

CMakeLists.txt

@@ -36,8 +36,8 @@ include(simd)
 ################################ Configurations #######################################
 option(WITH_GPU         "Compile PaddlePaddle with NVIDIA GPU"          ${CUDA_FOUND})
 option(WITH_AVX         "Compile PaddlePaddle with AVX intrinsics"      ${AVX_FOUND})
-option(WITH_MKLDNN      "Compile PaddlePaddle with mkl-dnn support."    OFF)
-option(WITH_MKLML       "Compile PaddlePaddle with mklml package."      OFF)
+option(WITH_MKLDNN      "Compile PaddlePaddle with mkl-dnn support."    ${AVX_FOUND})
+option(WITH_MKLML       "Compile PaddlePaddle with mklml package."      ${AVX_FOUND})
 option(WITH_DSO         "Compile PaddlePaddle with dynamic linked CUDA" ON)
 option(WITH_TESTING     "Compile PaddlePaddle with unit testing"        ON)
 option(WITH_SWIG_PY     "Compile PaddlePaddle with inference api"       ON)

Dockerfile

@@ -27,13 +27,16 @@ RUN apt-get update && \
     git python-pip python-dev openssh-server bison  \
     wget unzip unrar tar xz-utils bzip2 gzip coreutils ntp \
     curl sed grep graphviz libjpeg-dev zlib1g-dev  \
-    python-numpy python-matplotlib gcc-4.8 g++-4.8 \
+    python-matplotlib gcc-4.8 g++-4.8 \
     automake locales clang-format-3.8 swig doxygen cmake  \
     liblapack-dev liblapacke-dev libboost-dev \
     clang-3.8 llvm-3.8 libclang-3.8-dev \
     net-tools && \
     apt-get clean -y
 
+# paddle is using numpy.flip, which is introduced since 1.12.0
+RUN pip --no-cache-dir install 'numpy>=1.12.0'
+
 # Install Go and glide
 RUN wget -O go.tgz https://storage.googleapis.com/golang/go1.8.1.linux-amd64.tar.gz && \
     tar -C /usr/local -xzf go.tgz && \

cmake/configure.cmake

@@ -74,8 +74,6 @@ if(WITH_MKLDNN)
         set(OPENMP_FLAGS "-fopenmp")
         set(CMAKE_C_CREATE_SHARED_LIBRARY_FORBIDDEN_FLAGS ${OPENMP_FLAGS})
         set(CMAKE_CXX_CREATE_SHARED_LIBRARY_FORBIDDEN_FLAGS ${OPENMP_FLAGS})
-        set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} -L${MKLDNN_IOMP_DIR} -liomp5 -Wl,--as-needed")
-        set(CMAKE_EXE_LINKER_FLAGS "${CMAKE_EXE_LINKER_FLAGS} -L${MKLDNN_IOMP_DIR} -liomp5 -Wl,--as-needed")
         set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${OPENMP_FLAGS}")
         set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${OPENMP_FLAGS}")
     else()

cmake/cpplint.cmake

@@ -42,26 +42,21 @@ macro(add_style_check_target TARGET_NAME)
     if(WITH_STYLE_CHECK)
         set(SOURCES_LIST ${ARGN})
         list(REMOVE_DUPLICATES SOURCES_LIST)
-        list(SORT SOURCES_LIST)
-
         foreach(filename ${SOURCES_LIST})
-            set(LINT ON)
             foreach(pattern ${IGNORE_PATTERN})
                 if(filename MATCHES ${pattern})
-                    message(STATUS "DROP LINT ${filename}")
-                    set(LINT OFF)
+                    list(REMOVE_ITEM SOURCES_LIST ${filename})
                 endif()
             endforeach()
-            if(LINT MATCHES ON)
-                # cpplint code style
-                get_filename_component(base_filename ${filename} NAME)
-                set(CUR_GEN ${CMAKE_CURRENT_BINARY_DIR}/${base_filename}.cpplint)
-                add_custom_command(TARGET ${TARGET_NAME} PRE_BUILD
-                    COMMAND "${PYTHON_EXECUTABLE}" "${PROJ_ROOT}/paddle/scripts/cpplint.py"
-                            "--filter=${STYLE_FILTER}"
-                            "--write-success=${CUR_GEN}" ${filename}
-                    WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR})
-            endif()
         endforeach()
+
+        if(SOURCES_LIST)
+            add_custom_command(TARGET ${TARGET_NAME} POST_BUILD
+                COMMAND "${PYTHON_EXECUTABLE}" "${PROJ_ROOT}/paddle/scripts/cpplint.py"
+                        "--filter=${STYLE_FILTER}"
+                        ${SOURCES_LIST}
+                COMMENT "cpplint: Checking source code style"
+                WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR})        
+        endif()
     endif()
 endmacro()

cmake/external/any.cmake

@@ -7,7 +7,7 @@ INCLUDE_DIRECTORIES(${ANY_SOURCE_DIR}/src/extern_lib_any)
 ExternalProject_Add(
     extern_lib_any
     ${EXTERNAL_PROJECT_LOG_ARGS}
-    GIT_REPOSITORY  "https://github.com/thelink2012/any.git"
+    GIT_REPOSITORY  "https://github.com/PaddlePaddle/any.git"
     GIT_TAG         "8fef1e93710a0edf8d7658999e284a1142c4c020"
     PREFIX          ${ANY_SOURCE_DIR}
     UPDATE_COMMAND  ""

cmake/external/gflags.cmake

@@ -28,7 +28,14 @@ INCLUDE_DIRECTORIES(${GFLAGS_INCLUDE_DIR})
 ExternalProject_Add(
     extern_gflags
     ${EXTERNAL_PROJECT_LOG_ARGS}
-    GIT_REPOSITORY  "https://github.com/gflags/gflags.git"
+    # TODO(yiwang): The annoying warnings mentioned in
+    # https://github.com/PaddlePaddle/Paddle/issues/3277 are caused by
+    # gflags.  I fired a PR https://github.com/gflags/gflags/pull/230
+    # to fix it.  Before it gets accepted by the gflags team, we use
+    # my personal fork, which contains above fix, temporarily.  Let's
+    # change this back to the official Github repo once my PR is
+    # merged.
+    GIT_REPOSITORY  "https://github.com/wangkuiyi/gflags.git"
     PREFIX          ${GFLAGS_SOURCES_DIR}
     UPDATE_COMMAND  ""
     CMAKE_ARGS      -DCMAKE_CXX_COMPILER=${CMAKE_CXX_COMPILER}

cmake/external/openblas.cmake

@@ -69,8 +69,13 @@ ENDIF(NOT ${CBLAS_FOUND})
 MESSAGE(STATUS "BLAS library: ${CBLAS_LIBRARIES}")
 INCLUDE_DIRECTORIES(${CBLAS_INC_DIR})
 
-ADD_LIBRARY(cblas STATIC IMPORTED)
-SET_PROPERTY(TARGET cblas PROPERTY IMPORTED_LOCATION ${CBLAS_LIBRARIES})
+# FIXME(gangliao): generate cblas target to track all high performance
+# linear algebra libraries for cc_library(xxx SRCS xxx.c DEPS cblas)
+SET(dummyfile ${CMAKE_CURRENT_BINARY_DIR}/cblas_dummy.c)
+FILE(WRITE ${dummyfile} "const char * dummy = \"${dummyfile}\";")
+ADD_LIBRARY(cblas STATIC ${dummyfile})
+TARGET_LINK_LIBRARIES(cblas ${CBLAS_LIBRARIES})
+
 IF(NOT ${CBLAS_FOUND})
     ADD_DEPENDENCIES(cblas extern_openblas)
     LIST(APPEND external_project_dependencies cblas)

cmake/external/python.cmake

@@ -24,7 +24,6 @@ IF(WITH_PYTHON)
 ENDIF(WITH_PYTHON)
 
 SET(py_env "")
-SET(USE_VIRTUALENV_FOR_TEST 1)
 IF(PYTHONINTERP_FOUND)
     find_python_module(pip REQUIRED)
     find_python_module(numpy REQUIRED)

cmake/flags.cmake

@@ -115,7 +115,7 @@ set(COMMON_FLAGS
     -Wno-error=literal-suffix
     -Wno-error=sign-compare
     -Wno-error=unused-local-typedefs
-    -Wno-error=parentheses-equality # Warnings in Pybind11
+    -Wno-error=parentheses-equality # Warnings in pybind11
 )
 
 set(GPU_COMMON_FLAGS
@@ -195,6 +195,7 @@ endif()
 # Modern gpu architectures: Pascal
 if (CUDA_VERSION VERSION_GREATER "8.0" OR CUDA_VERSION VERSION_EQUAL "8.0")
       list(APPEND __arch_flags " -gencode arch=compute_60,code=sm_60")
+      list(APPEND CUDA_NVCC_FLAGS --expt-relaxed-constexpr)
 endif()
 
 # Custom gpu architecture

cmake/generic.cmake

@@ -403,3 +403,16 @@ function(py_proto_compile TARGET_NAME)
   protobuf_generate_python(py_srcs ${py_proto_compile_SRCS})
   add_custom_target(${TARGET_NAME} ALL DEPENDS ${py_srcs})
 endfunction()
+
+function(py_test TARGET_NAME)
+  if(WITH_TESTING)
+    set(options STATIC static SHARED shared)
+    set(oneValueArgs "")
+    set(multiValueArgs SRCS DEPS)
+    cmake_parse_arguments(py_test "${options}" "${oneValueArgs}" "${multiValueArgs}" ${ARGN})  
+    add_test(NAME ${TARGET_NAME}
+             COMMAND env PYTHONPATH=${PADDLE_PYTHON_PACKAGE_DIR}
+             python2 ${py_test_SRCS}
+             WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR})
+  endif()
+endfunction()

cmake/util.cmake

@@ -118,7 +118,6 @@ endfunction()
 macro(add_unittest_without_exec TARGET_NAME)
     add_executable(${TARGET_NAME} ${ARGN})
     link_paddle_test(${TARGET_NAME})
-    add_style_check_target(${TARGET_NAME} ${ARGN})
 endmacro()
 
 # add_unittest
@@ -150,9 +149,12 @@ endfunction()
 # Create a python unittest using run_python_tests.sh,
 # which takes care of making correct running environment
 function(add_python_test TEST_NAME)
-  add_test(NAME ${TEST_NAME}
-        COMMAND env PADDLE_PACKAGE_DIR=${PADDLE_PYTHON_PACKAGE_DIR}
-        bash ${PROJ_ROOT}/paddle/scripts/run_python_tests.sh
-        ${USE_VIRTUALENV_FOR_TEST} ${PYTHON_EXECUTABLE} ${ARGN}
-        WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR})
+    foreach(arg ${ARGN})
+        get_filename_component(py_fn ${arg} NAME_WE)
+        set(TRG_NAME ${TEST_NAME}_${py_fn})
+        add_test(NAME ${TRG_NAME}
+                COMMAND env PYTHONPATH=${PADDLE_PYTHON_PACKAGE_DIR}
+                python2 ${arg}
+                WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR})
+    endforeach()
 endfunction()

doc/design/mkldnn/README.MD

+# Intel® MKL-DNN on PaddlePaddle: Design Doc
+
+我们计划将Intel深度神经网络数学库(**MKL-DNN**\[[1](#references)\])集成到PaddlePaddle，充分展现英特尔平台的优势，有效提升PaddlePaddle在英特尔架构上的性能。
+
+我们短期内的基本目标是：
+
+- 完成常用layer的MKL-DNN实现。
+- 完成常见深度神经网络VGG，GoogLeNet 和 ResNet的MKL-DNN实现。
+
+
+## Contents
+
+- [Overview](#overview)
+- [Actions](#actions)
+ 	- [CMake](#cmake)
+	- [Layers](#layers)
+	- [Activations](#activations)
+	- [Unit Tests](#unit-tests)
+	- [Protobuf Messages](#protobuf-messages)
+	- [Python API](#python-api)
+	- [Demos](#demos)
+	- [Benchmarking](#benchmarking)
+	- [Others](#others)
+- [Design Concerns](#design-concerns)
+
+## Overview
+
+我们会把MKL-DNN作为第三方库集成进PaddlePaddle，整体框架图
+<div align="center">
+<img src="image/overview.png" width=350><br/>
+Figure 1. PaddlePaddle on IA.
+</div>
+
+## Actions
+我们把集成方案大致分为了如下几个方面。
+
+### CMake
+我们会在`CMakeLists.txt`中会添加`WITH_MKLDNN`的选项，当设置这个值为`ON`的时候会启用编译MKL-DNN功能。同时会自动开启OpenMP用于提高MKL-DNN的性能。
+
+同时，我们会引入`WITH_MKLML`选项，用于选择是否使用MKL-DNN自带的MKLML安装包。这个安装包可以独立于MKL-DNN使用，但是建议在开启MKL-DNN的同时也打开MKLML的开关，这样才能发挥最好的性能。
+
+所以，我们会在`cmake/external`目录新建`mkldnn.cmake`和`mklml.cmake`文件，它们会在编译PaddlePaddle的时候下载对应的软件包，并放到PaddlePaddle的third party目录中。
+
+**备注**：当`WITH_MKLML=ON`的时候，会优先使用这个包作为PaddlePaddle的CBLAS和LAPACK库，所以会稍微改动`cmake/cblas.cmake`中的逻辑。
+
+### Layers
+所有MKL-DNN相关的C++ layers，都会按照PaddlePaddle的目录结构存放在
+`paddle/gserver/layers`中，并且文件名都会一以*Mkldnn*开头。
+
+所有MKL-DNN的layers都会继承于一个叫做`MkldnnLayer`的父类，该父类继承于PaddlePaddle的基类`Layer`。
+
+### Activations
+由于在PaddlePaddle中，激活函数是独立于layer概念的，所以会在`paddle/gserver/activations`目录下添加一个`MkldnnActivation.h`文件定义一些用于MKL-DNN的接口，实现方法还是会在`ActivationFunction.cpp`文件。
+
+### Unit Tests
+会在`paddle/gserver/test`目录下添加`test_Mkldnn.cpp`和`MkldnnTester.*`用于MKL-DNN的测试。
+
+Activation的测试，计划在PaddlePaddle原有的测试文件上直接添加新的测试type。
+
+### Protobuf Messages
+根据具体layer的需求可能会在`proto/ModelConfig.proto`里面添加必要的选项。
+
+### Python API
+目前只考虑**v1 API**。
+
+计划在`python/paddle/trainer/config_parser.py`里面添加`use_mkldnn`这个选择，方便用户选择使用MKL-DNN的layers。
+
+具体实现方式比如：
+
+```python
+use_mkldnn = bool(int(g_command_config_args.get("use_mkldnn", 0)))
+if use_mkldnn
+    self.layer_type = mkldnn_*
+```
+
+所有MKL-DNN的layer type会以*mkldnn_*开头，以示区分。 
+
+并且可能在`python/paddle/trainer_config_helper`目录下的`activations.py `和`layers.py`里面添加必要的MKL-DNN的接口。
+
+### Demos
+
+会在`v1_api_demo`目录下添加一个`mkldnn`的文件夹，里面放入一些用于MKL-DNN测试的demo脚本。
+
+### Benchmarking
+会考虑添加部分逻辑在`benchmark/paddle/image/run.sh`，添加使用MKL-DNN的测试。
+
+### Others
+1. 如果在使用MKL-DNN的情况下，会把CPU的Buffer对齐为64。
+2. 深入PaddlePaddle，寻找有没有其他可以优化的可能，进一步优化。比如可能会用OpenMP改进SGD的更新性能。
+
+## Design Concerns
+
+为了更好的符合PaddlePaddle的代码风格\[[2](#references)\]，同时又尽可能少的牺牲MKL-DNN的性能\[[3](#references)\]。
+
+我们总结出一些特别需要注意的点：
+
+1. 使用**deviceId_**。为了尽可能少的在父类Layer中添加变量或者函数，我们决定使用已有的`deviceId_`变量来区分layer的属性，定义`-2`为`MkldnnLayer`特有的设备ID。
+2. 重写父类Layer的**init**函数，修改`deviceId_`为`-2`，代表这个layer是用于跑在MKL-DNN的环境下。
+3. 创建`MkldnnMatrix`，用于管理MKL-DNN会用到的相关memory函数、接口以及会用的到格式信息。
+4. 创建`MkldnnBase`，定义一些除了layer和memory相关的类和函数。包括MKL-DNN会用到`MkldnnStream`和`CpuEngine`，和未来可能还会用到`FPGAEngine`等。
+5. 在**Argument**里添加两个`MkldnnMatrixPtr`，取名为`mkldnnValue`和`mkldnnGrad`，用于存放`MkldnnLayer`会用到的memory buffer。 并且添加函数cvt(会修改为一个更加合适的函数名)，用于处理"CPU device"和"MKL-DNN device"之间memory的相互转化。
+6. 在父类`Layer`中的`getOutput`函数中添加一段逻辑，用于判断`deviceId`，并针对device在MKL-DNN和CPU之间不统一的情况，做一个前期转换。 也就是调用`Argument`的cvt函数把output统一到需要的device上。
+7. 在原来的`FLAGS`中添加一个`use_mkldnn`的flag，用于选择是否使用MKL-DNN的相关功能。
+
+## References
+
+1. [Intel Math Kernel Library for Deep Neural Networks (Intel MKL-DNN)](https://github.com/01org/mkl-dnn "Intel MKL-DNN")
+2. [原来的方案](https://github.com/PaddlePaddle/Paddle/pull/3096)会引入**nextLayer**的信息。但是在PaddlePaddle中，无论是重构前的layer还是重构后的op，都不会想要知道next layer/op的信息。
+3. MKL-DNN的高性能格式与PaddlePaddle原有的`NCHW`不同(PaddlePaddle中的CUDNN部分使用的也是`NCHW`，所以不存在这个问题)，所以需要引入一个转换方法，并且只需要在必要的时候转换这种格式，才能更好的发挥MKL-DNN的性能。
+

paddle/.set_python_path.sh

@@ -21,22 +21,15 @@
 # 
 # It same as PYTHONPATH=${YOUR_PYTHON_PATH}:$PYTHONPATH {exec...}
 #
-
-if ! python -c "import paddle" >/dev/null 2>/dev/null; then
-  PYPATH=""
-  set -x
-  while getopts "d:" opt; do
-    case $opt in
-      d)
-        PYPATH=$OPTARG
-        ;;
-    esac
-  done
-  shift $(($OPTIND - 1))
-  export PYTHONPATH=$PYPATH:$PYTHONPATH
-  $@
-else
-  echo "paddle package is already in your PYTHONPATH. But unittest need a clean environment."
-  echo "Please uninstall paddle package before start unittest. Try to 'pip uninstall paddle'"
-  exit 1
-fi
+PYPATH=""
+set -x
+while getopts "d:" opt; do
+  case $opt in
+    d)
+      PYPATH=$OPTARG
+      ;;
+  esac
+done
+shift $(($OPTIND - 1))
+export PYTHONPATH=$PYPATH:$PYTHONPATH
+$@

paddle/api/test/CMakeLists.txt

-add_python_test(test_swig_api
-    testArguments.py testGradientMachine.py testMatrix.py testVector.py testTrain.py testTrainer.py)
+py_test(testTrain SRCS testTrain.py)
+py_test(testMatrix SRCS testMatrix.py)
+py_test(testVector SRCS testVector.py)
+py_test(testTrainer SRCS testTrainer.py)
+py_test(testArguments SRCS testArguments.py)
+py_test(testGradientMachine SRCS testGradientMachine.py)

paddle/cuda/src/hl_batch_transpose.cu

@@ -12,17 +12,15 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 See the License for the specific language governing permissions and
 limitations under the License. */
 
-#include "hl_batch_transpose.h"
 #include "hl_base.h"
+#include "hl_batch_transpose.h"
 
 const int TILE_DIM = 64;
 const int BLOCK_ROWS = 16;
 
 // No bank-conflict transpose for a batch of data.
-__global__ void batchTransposeNoBankConflicts(real* odata,
-                                              const real* idata,
-                                              int numSamples, int width,
-                                              int height) {
+__global__ void batchTransposeNoBankConflicts(
+    real* odata, const real* idata, int numSamples, int width, int height) {
   __shared__ float tile[TILE_DIM][TILE_DIM + 1];
 
   const int x = blockIdx.x * TILE_DIM + threadIdx.x;
@@ -50,12 +48,12 @@ __global__ void batchTransposeNoBankConflicts(real* odata,
           newX] = tile[threadIdx.x][j];
 }
 
-void batchTranspose(const real* input, real* output, int width, int height,
-                    int batchSize) {
+void batchTranspose(
+    const real* input, real* output, int width, int height, int batchSize) {
   dim3 dimBlock(TILE_DIM, BLOCK_ROWS, 1);
   dim3 dimGrid(DIVUP(width, TILE_DIM), DIVUP(height, TILE_DIM), batchSize);
-  batchTransposeNoBankConflicts<<<dimGrid, dimBlock, 0, STREAM_DEFAULT>>>
-      (output, input, batchSize, width, height);
+  batchTransposeNoBankConflicts<<<dimGrid, dimBlock, 0, STREAM_DEFAULT>>>(
+      output, input, batchSize, width, height);
 
   CHECK_SYNC("batchTranspose failed!");
 }

paddle/cuda/src/hl_cuda_aggregate.cu

@@ -12,27 +12,23 @@ 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 "hl_aggregate.h"
 #include "hl_base.h"
 #include "hl_cuda.h"
 #include "hl_cuda.ph"
-#include "hl_aggregate.h"
-#include "hl_thread.ph"
 #include "hl_matrix_base.cuh"
+#include "hl_thread.ph"
 #include "paddle/utils/Logging.h"
 
 /**
  * @brief   matrix row operator.
  */
-template<class Agg, int blockSize>
-__global__ void KeMatrixRowOp(Agg agg,
-                              real *E,
-                              real *Sum,
-                              int dimN) {
+template <class Agg, int blockSize>
+__global__ void KeMatrixRowOp(Agg agg, real *E, real *Sum, int dimN) {
   __shared__ real sum_s[blockSize];
-  int cnt = (dimN + blockSize -1) / blockSize;
-  int rowId = blockIdx.x + blockIdx.y*gridDim.x;
-  int index = rowId*dimN;
+  int cnt = (dimN + blockSize - 1) / blockSize;
+  int rowId = blockIdx.x + blockIdx.y * gridDim.x;
+  int index = rowId * dimN;
   int tid = threadIdx.x;
   int lmt = tid;
 
@@ -44,7 +40,7 @@ __global__ void KeMatrixRowOp(Agg agg,
   sum_s[tid] = tmp;
   __syncthreads();
 
-  for (int stride = blockSize/2; stride > 0; stride = stride/2) {
+  for (int stride = blockSize / 2; stride > 0; stride = stride / 2) {
     if (tid < stride) {
       sum_s[tid] = agg(sum_s[tid], sum_s[tid + stride]);
     }
@@ -58,29 +54,21 @@ __global__ void KeMatrixRowOp(Agg agg,
 }
 
 template <class Agg>
-void hl_matrix_row_op(Agg agg,
-                      real *A_d,
-                      real *C_d,
-                      int dimM,
-                      int dimN) {
+void hl_matrix_row_op(Agg agg, real *A_d, real *C_d, int dimM, int dimN) {
   int blocksX = dimM;
   int blocksY = 1;
   dim3 threads(128, 1);
   dim3 grid(blocksX, blocksY);
 
-  KeMatrixRowOp<Agg, 128><<< grid, threads, 0, STREAM_DEFAULT >>>
-           (agg, A_d, C_d, dimN);
+  KeMatrixRowOp<Agg, 128><<<grid, threads, 0, STREAM_DEFAULT>>>(
+      agg, A_d, C_d, dimN);
 }
 
 void hl_matrix_row_sum(real *A_d, real *C_d, int dimM, int dimN) {
   CHECK_NOTNULL(A_d);
   CHECK_NOTNULL(C_d);
 
-  hl_matrix_row_op(aggregate::sum(),
-                   A_d,
-                   C_d,
-                   dimM,
-                   dimN);
+  hl_matrix_row_op(aggregate::sum(), A_d, C_d, dimM, dimN);
   CHECK_SYNC("hl_matrix_row_sum failed");
 }
 
@@ -88,11 +76,7 @@ void hl_matrix_row_max(real *A_d, real *C_d, int dimM, int dimN) {
   CHECK_NOTNULL(A_d);
   CHECK_NOTNULL(C_d);
 
-  hl_matrix_row_op(aggregate::max(),
-                   A_d,
-                   C_d,
-                   dimM,
-                   dimN);
+  hl_matrix_row_op(aggregate::max(), A_d, C_d, dimM, dimN);
   CHECK_SYNC("hl_matrix_row_max failed");
 }
 
@@ -100,23 +84,16 @@ void hl_matrix_row_min(real *A_d, real *C_d, int dimM, int dimN) {
   CHECK_NOTNULL(A_d);
   CHECK_NOTNULL(C_d);
 
-  hl_matrix_row_op(aggregate::min(),
-                   A_d,
-                   C_d,
-                   dimM,
-                   dimN);
+  hl_matrix_row_op(aggregate::min(), A_d, C_d, dimM, dimN);
   CHECK_SYNC("hl_matrix_row_min failed");
 }
 
 /**
  * @brief   matrix column operator.
  */
-template<class Agg>
-__global__ void KeMatrixColumnOp(Agg agg,
-                                 real *E,
-                                 real *Sum,
-                                 int dimM,
-                                 int dimN) {
+template <class Agg>
+__global__ void KeMatrixColumnOp(
+    Agg agg, real *E, real *Sum, int dimM, int dimN) {
   int rowIdx = blockIdx.x * blockDim.x + threadIdx.x;
   real tmp = agg.init();
   if (rowIdx < dimN) {
@@ -127,15 +104,12 @@ __global__ void KeMatrixColumnOp(Agg agg,
   }
 }
 
-template<class Agg, int blockDimX, int blockDimY>
-__global__ void KeMatrixColumnOp_S(Agg agg,
-                                   real *E,
-                                   real *Sum,
-                                   int dimM,
-                                   int dimN) {
-    __shared__ real _sum[blockDimX*blockDimY];
-    int rowIdx = blockIdx.x * blockDim.x + threadIdx.x;
-    int index = threadIdx.y;
+template <class Agg, int blockDimX, int blockDimY>
+__global__ void KeMatrixColumnOp_S(
+    Agg agg, real *E, real *Sum, int dimM, int dimN) {
+  __shared__ real _sum[blockDimX * blockDimY];
+  int rowIdx = blockIdx.x * blockDim.x + threadIdx.x;
+  int index = threadIdx.y;
 
   real tmp = agg.init();
   if (rowIdx < dimN) {
@@ -144,14 +118,14 @@ __global__ void KeMatrixColumnOp_S(Agg agg,
       index += blockDimY;
     }
   }
-  _sum[threadIdx.x + threadIdx.y*blockDimX] = tmp;
+  _sum[threadIdx.x + threadIdx.y * blockDimX] = tmp;
   __syncthreads();
 
   if (rowIdx < dimN) {
-    if (threadIdx.y ==0) {
+    if (threadIdx.y == 0) {
       real tmp = agg.init();
-      for (int i=0; i < blockDimY; i++) {
-        tmp = agg(tmp, _sum[threadIdx.x + i*blockDimX]);
+      for (int i = 0; i < blockDimY; i++) {
+        tmp = agg(tmp, _sum[threadIdx.x + i * blockDimX]);
       }
       Sum[rowIdx] = tmp;
     }
@@ -159,25 +133,21 @@ __global__ void KeMatrixColumnOp_S(Agg agg,
 }
 
 template <class Agg>
-void hl_matrix_column_op(Agg agg,
-                         real *A_d,
-                         real *C_d,
-                         int dimM,
-                         int dimN) {
+void hl_matrix_column_op(Agg agg, real *A_d, real *C_d, int dimM, int dimN) {
   if (dimN >= 8192) {
-    int blocksX = (dimN + 128 -1) / 128;
+    int blocksX = (dimN + 128 - 1) / 128;
     int blocksY = 1;
     dim3 threads(128, 1);
     dim3 grid(blocksX, blocksY);
-    KeMatrixColumnOp<Agg><<< grid, threads, 0, STREAM_DEFAULT >>>
-             (agg, A_d, C_d, dimM, dimN);
+    KeMatrixColumnOp<Agg><<<grid, threads, 0, STREAM_DEFAULT>>>(
+        agg, A_d, C_d, dimM, dimN);
   } else {
-    int blocksX = (dimN + 32 -1) / 32;
+    int blocksX = (dimN + 32 - 1) / 32;
     int blocksY = 1;
     dim3 threads(32, 32);
     dim3 grid(blocksX, blocksY);
-    KeMatrixColumnOp_S<Agg, 32, 32><<< grid, threads, 0, STREAM_DEFAULT>>>
-             (agg, A_d, C_d, dimM, dimN);
+    KeMatrixColumnOp_S<Agg, 32, 32><<<grid, threads, 0, STREAM_DEFAULT>>>(
+        agg, A_d, C_d, dimM, dimN);
   }
 
   return;
@@ -187,11 +157,7 @@ void hl_matrix_column_sum(real *A_d, real *C_d, int dimM, int dimN) {
   CHECK_NOTNULL(A_d);
   CHECK_NOTNULL(C_d);
 
-  hl_matrix_column_op(aggregate::sum(),
-                      A_d,
-                      C_d,
-                      dimM,
-                      dimN);
+  hl_matrix_column_op(aggregate::sum(), A_d, C_d, dimM, dimN);
 
   CHECK_SYNC("hl_matrix_column_sum failed");
 }
@@ -200,11 +166,7 @@ void hl_matrix_column_max(real *A_d, real *C_d, int dimM, int dimN) {
   CHECK_NOTNULL(A_d);
   CHECK_NOTNULL(C_d);
 
-  hl_matrix_column_op(aggregate::max(),
-                      A_d,
-                      C_d,
-                      dimM,
-                      dimN);
+  hl_matrix_column_op(aggregate::max(), A_d, C_d, dimM, dimN);
 
   CHECK_SYNC("hl_matrix_column_max failed");
 }
@@ -213,11 +175,7 @@ void hl_matrix_column_min(real *A_d, real *C_d, int dimM, int dimN) {
   CHECK_NOTNULL(A_d);
   CHECK_NOTNULL(C_d);
 
-  hl_matrix_column_op(aggregate::min(),
-                      A_d,
-                      C_d,
-                      dimM,
-                      dimN);
+  hl_matrix_column_op(aggregate::min(), A_d, C_d, dimM, dimN);
 
   CHECK_SYNC("hl_matrix_column_min failed");
 }
@@ -226,16 +184,16 @@ template <int blockSize>
 __global__ void KeVectorSum(real *E, real *Sum, int dimM) {
   __shared__ double sum_s[blockSize];
   int tid = threadIdx.x;
-  int index = blockIdx.y*blockDim.x+threadIdx.x;
+  int index = blockIdx.y * blockDim.x + threadIdx.x;
 
   sum_s[tid] = 0.0f;
   while (index < dimM) {
     sum_s[tid] += E[index];
-    index += blockDim.x*gridDim.y;
+    index += blockDim.x * gridDim.y;
   }
   __syncthreads();
 
-  for (int stride = blockSize/2; stride > 0; stride = stride/2) {
+  for (int stride = blockSize / 2; stride > 0; stride = stride / 2) {
     if (tid < stride) {
       sum_s[tid] += sum_s[tid + stride];
     }
@@ -259,38 +217,39 @@ void hl_vector_sum(real *A_d, real *C_h, int dimM) {
   dim3 threads(blockSize, 1);
   dim3 grid(blocksX, blocksY);
 
-  struct _hl_event_st hl_event_st  = {.cu_event = t_resource.event};
+  struct _hl_event_st hl_event_st = {.cu_event = t_resource.event};
   hl_event_t hl_event = &hl_event_st;
-  while (!hl_cuda_event_is_ready(hl_event)) {}
+  while (!hl_cuda_event_is_ready(hl_event)) {
+  }
 
-  KeVectorSum<128><<< grid, threads, 0, STREAM_DEFAULT >>>
-           (A_d, t_resource.gpu_mem, dimM);
-  KeVectorSum<128><<< 1, threads, 0, STREAM_DEFAULT >>>
-           (t_resource.gpu_mem, t_resource.cpu_mem, 128);
+  KeVectorSum<128><<<grid, threads, 0, STREAM_DEFAULT>>>(
+      A_d, t_resource.gpu_mem, dimM);
+  KeVectorSum<128><<<1, threads, 0, STREAM_DEFAULT>>>(
+      t_resource.gpu_mem, t_resource.cpu_mem, 128);
 
   hl_memcpy_async(C_h, t_resource.cpu_mem, sizeof(real), HPPL_STREAM_DEFAULT);
   hl_stream_record_event(HPPL_STREAM_DEFAULT, hl_event);
 
   hl_stream_synchronize(HPPL_STREAM_DEFAULT);
   cudaError_t err = (cudaError_t)hl_get_device_last_error();
-  CHECK_EQ(cudaSuccess, err)
-    << "CUDA error: " << hl_get_device_error_string((size_t)err);
+  CHECK_EQ(cudaSuccess, err) << "CUDA error: "
+                             << hl_get_device_error_string((size_t)err);
 }
 
 template <int blockSize>
 __global__ void KeVectorAbsSum(real *E, real *Sum, int dimM) {
   __shared__ double sum_s[blockSize];
   int tid = threadIdx.x;
-  int index = blockIdx.y*blockDim.x+threadIdx.x;
+  int index = blockIdx.y * blockDim.x + threadIdx.x;
 
   sum_s[tid] = 0.0f;
   while (index < dimM) {
     sum_s[tid] += abs(E[index]);
-    index += blockDim.x*gridDim.y;
+    index += blockDim.x * gridDim.y;
   }
   __syncthreads();
 
-  for (int stride = blockSize/2; stride > 0; stride = stride/2) {
+  for (int stride = blockSize / 2; stride > 0; stride = stride / 2) {
     if (tid < stride) {
       sum_s[tid] += sum_s[tid + stride];
     }
@@ -314,20 +273,21 @@ void hl_vector_abs_sum(real *A_d, real *C_h, int dimM) {
   dim3 threads(blockSize, 1);
   dim3 grid(blocksX, blocksY);
 
-  struct _hl_event_st hl_event_st  = {.cu_event = t_resource.event};
+  struct _hl_event_st hl_event_st = {.cu_event = t_resource.event};
   hl_event_t hl_event = &hl_event_st;
-  while (!hl_cuda_event_is_ready(hl_event)) {}
+  while (!hl_cuda_event_is_ready(hl_event)) {
+  }
 
-  KeVectorAbsSum<128><<< grid, threads, 0, STREAM_DEFAULT >>>
-           (A_d, t_resource.gpu_mem, dimM);
-  KeVectorAbsSum<128><<< 1, threads, 0, STREAM_DEFAULT >>>
-           (t_resource.gpu_mem, t_resource.cpu_mem, 128);
+  KeVectorAbsSum<128><<<grid, threads, 0, STREAM_DEFAULT>>>(
+      A_d, t_resource.gpu_mem, dimM);
+  KeVectorAbsSum<128><<<1, threads, 0, STREAM_DEFAULT>>>(
+      t_resource.gpu_mem, t_resource.cpu_mem, 128);
 
   hl_memcpy_async(C_h, t_resource.cpu_mem, sizeof(real), HPPL_STREAM_DEFAULT);
   hl_stream_record_event(HPPL_STREAM_DEFAULT, hl_event);
 
   hl_stream_synchronize(HPPL_STREAM_DEFAULT);
   cudaError_t err = (cudaError_t)hl_get_device_last_error();
-  CHECK_EQ(cudaSuccess, err)
-    << "CUDA error: " << hl_get_device_error_string((size_t)err);
+  CHECK_EQ(cudaSuccess, err) << "CUDA error: "
+                             << hl_get_device_error_string((size_t)err);
 }

paddle/cuda/src/hl_cuda_sequence.cu

@@ -16,36 +16,36 @@ limitations under the License. */
 #include "hl_device_functions.cuh"
 #include "paddle/utils/Logging.h"
 
-__global__ void KeMaxSequenceForward(real *input,
-                                     const int *sequence,
+__global__ void KeMaxSequenceForward(real* input,
+                                     const int* sequence,
                                      real* output,
-                                     int *index,
+                                     int* index,
                                      int numSequences,
                                      int dim) {
   int dimIdx = threadIdx.x;
   int sequenceId = blockIdx.x;
   if (sequenceId >= numSequences) return;
   int start = sequence[sequenceId];
-  int end = sequence[sequenceId+1];
+  int end = sequence[sequenceId + 1];
 
   for (int i = dimIdx; i < dim; i += blockDim.x) {
     real tmp = -HL_FLOAT_MAX;
     int tmpId = -1;
     for (int insId = start; insId < end; insId++) {
-      if (tmp < input[insId*dim + i]) {
-        tmp = input[insId*dim + i];
+      if (tmp < input[insId * dim + i]) {
+        tmp = input[insId * dim + i];
         tmpId = insId;
       }
     }
-    output[sequenceId*dim + i] = tmp;
-    index[sequenceId*dim + i] = tmpId;
+    output[sequenceId * dim + i] = tmp;
+    index[sequenceId * dim + i] = tmpId;
   }
 }
 
 void hl_max_sequence_forward(real* input,
                              const int* sequence,
                              real* output,
-                             int *index,
+                             int* index,
                              int numSequences,
                              int dim) {
   CHECK_NOTNULL(input);
@@ -55,29 +55,23 @@ void hl_max_sequence_forward(real* input,
 
   dim3 threads(256, 1);
   dim3 grid(numSequences, 1);
-  KeMaxSequenceForward<<< grid, threads, 0, STREAM_DEFAULT >>>
-      (input, sequence, output, index, numSequences, dim);
+  KeMaxSequenceForward<<<grid, threads, 0, STREAM_DEFAULT>>>(
+      input, sequence, output, index, numSequences, dim);
   CHECK_SYNC("hl_max_sequence_forward failed");
 }
 
-__global__ void KeMaxSequenceBackward(real *outputGrad,
-                                      int *index,
-                                      real* inputGrad,
-                                      int numSequences,
-                                      int dim) {
+__global__ void KeMaxSequenceBackward(
+    real* outputGrad, int* index, real* inputGrad, int numSequences, int dim) {
   int idx = threadIdx.x + blockIdx.x * blockDim.x;
   int colIdx = idx % dim;
-  if (idx < numSequences*dim) {
+  if (idx < numSequences * dim) {
     int insId = index[idx];
     inputGrad[insId * dim + colIdx] += outputGrad[idx];
   }
 }
 
-void hl_max_sequence_backward(real* outputGrad,
-                              int *index,
-                              real* inputGrad,
-                              int numSequences,
-                              int dim) {
+void hl_max_sequence_backward(
+    real* outputGrad, int* index, real* inputGrad, int numSequences, int dim) {
   CHECK_NOTNULL(outputGrad);
   CHECK_NOTNULL(index);
   CHECK_NOTNULL(inputGrad);
@@ -85,12 +79,12 @@ void hl_max_sequence_backward(real* outputGrad,
   unsigned int blocks = (numSequences * dim + 128 - 1) / 128;
   dim3 threads(128, 1);
   dim3 grid(blocks, 1);
-  KeMaxSequenceBackward<<< grid, threads, 0, STREAM_DEFAULT >>>
-      (outputGrad, index, inputGrad, numSequences, dim);
+  KeMaxSequenceBackward<<<grid, threads, 0, STREAM_DEFAULT>>>(
+      outputGrad, index, inputGrad, numSequences, dim);
   CHECK_SYNC("hl_max_sequence_backward failed");
 }
 
-template<int blockDimX, int blockDimY, int gridDimX, bool AddRow>
+template <int blockDimX, int blockDimY, int gridDimX, bool AddRow>
 __global__ void KeMatrixAddRows(real* output,
                                 real* table,
                                 int* ids,
@@ -104,8 +98,8 @@ __global__ void KeMatrixAddRows(real* output,
   while (sampleId < numSamples) {
     int tableId = ids[sampleId];
     if ((0 <= tableId) && (tableId < tableSize)) {
-      real *outputData = output + sampleId * dim;
-      real *tableData = table + tableId * dim;
+      real* outputData = output + sampleId * dim;
+      real* tableData = table + tableId * dim;
       for (int i = idx; i < dim; i += blockDimX) {
         if (AddRow == 0) {
           outputData[i] += tableData[i];
@@ -114,24 +108,27 @@ __global__ void KeMatrixAddRows(real* output,
         }
       }
     }
-    sampleId += blockDimY*gridDimX;
+    sampleId += blockDimY * gridDimX;
   }
 }
 
-template<int blockDimX, int blockDimY, int gridDimX, bool seq2batch, bool isAdd>
-__global__
-void KeSequence2Batch(real *batch,
-                      real *sequence,
-                      const int *batchIndex,
-                      int seqWidth,
-                      int batchCount) {
+template <int blockDimX,
+          int blockDimY,
+          int gridDimX,
+          bool seq2batch,
+          bool isAdd>
+__global__ void KeSequence2Batch(real* batch,
+                                 real* sequence,
+                                 const int* batchIndex,
+                                 int seqWidth,
+                                 int batchCount) {
   int idx = threadIdx.x;
   int idy = threadIdx.y;
   int id = blockIdx.x + idy * gridDimX;
   while (id < batchCount) {
     int seqId = batchIndex[id];
-    real* batchData = batch + id*seqWidth;
-    real* seqData = sequence + seqId*seqWidth;
+    real* batchData = batch + id * seqWidth;
+    real* seqData = sequence + seqId * seqWidth;
     for (int i = idx; i < seqWidth; i += blockDimX) {
       if (seq2batch) {
         if (isAdd) {
@@ -147,13 +144,13 @@ void KeSequence2Batch(real *batch,
         }
       }
     }
-    id += blockDimY*gridDimX;
+    id += blockDimY * gridDimX;
   }
 }
 
-void hl_sequence2batch_copy(real *batch,
-                            real *sequence,
-                            const int *batchIndex,
+void hl_sequence2batch_copy(real* batch,
+                            real* sequence,
+                            const int* batchIndex,
                             int seqWidth,
                             int batchCount,
                             bool seq2batch) {
@@ -164,18 +161,18 @@ void hl_sequence2batch_copy(real *batch,
   dim3 threads(128, 8);
   dim3 grid(8, 1);
   if (seq2batch) {
-    KeSequence2Batch<128, 8, 8, 1, 0><<< grid, threads, 0, STREAM_DEFAULT >>>
-      (batch, sequence, batchIndex, seqWidth, batchCount);
+    KeSequence2Batch<128, 8, 8, 1, 0><<<grid, threads, 0, STREAM_DEFAULT>>>(
+        batch, sequence, batchIndex, seqWidth, batchCount);
   } else {
-    KeSequence2Batch<128, 8, 8, 0, 0><<< grid, threads, 0, STREAM_DEFAULT >>>
-      (batch, sequence, batchIndex, seqWidth, batchCount);
+    KeSequence2Batch<128, 8, 8, 0, 0><<<grid, threads, 0, STREAM_DEFAULT>>>(
+        batch, sequence, batchIndex, seqWidth, batchCount);
   }
   CHECK_SYNC("hl_sequence2batch_copy failed");
 }
 
-void hl_sequence2batch_add(real *batch,
-                           real *sequence,
-                           int *batchIndex,
+void hl_sequence2batch_add(real* batch,
+                           real* sequence,
+                           int* batchIndex,
                            int seqWidth,
                            int batchCount,
                            bool seq2batch) {
@@ -186,23 +183,22 @@ void hl_sequence2batch_add(real *batch,
   dim3 threads(128, 8);
   dim3 grid(8, 1);
   if (seq2batch) {
-    KeSequence2Batch<128, 8, 8, 1, 1><<< grid, threads, 0, STREAM_DEFAULT >>>
-      (batch, sequence, batchIndex, seqWidth, batchCount);
+    KeSequence2Batch<128, 8, 8, 1, 1><<<grid, threads, 0, STREAM_DEFAULT>>>(
+        batch, sequence, batchIndex, seqWidth, batchCount);
   } else {
-    KeSequence2Batch<128, 8, 8, 0, 1><<< grid, threads, 0, STREAM_DEFAULT >>>
-      (batch, sequence, batchIndex, seqWidth, batchCount);
+    KeSequence2Batch<128, 8, 8, 0, 1><<<grid, threads, 0, STREAM_DEFAULT>>>(
+        batch, sequence, batchIndex, seqWidth, batchCount);
   }
   CHECK_SYNC("hl_sequence2batch_add failed");
 }
 
-template<bool normByTimes, bool seq2batch>
-__global__
-void KeSequence2BatchPadding(real* batch,
-                             real* sequence,
-                             const int* sequenceStartPositions,
-                             const size_t sequenceWidth,
-                             const size_t maxSequenceLength,
-                             const size_t numSequences) {
+template <bool normByTimes, bool seq2batch>
+__global__ void KeSequence2BatchPadding(real* batch,
+                                        real* sequence,
+                                        const int* sequenceStartPositions,
+                                        const size_t sequenceWidth,
+                                        const size_t maxSequenceLength,
+                                        const size_t numSequences) {
   int batchIdx = blockIdx.y;
   int sequenceStart = sequenceStartPositions[batchIdx];
   int sequenceLength = sequenceStartPositions[batchIdx + 1] - sequenceStart;
@@ -276,37 +272,49 @@ void hl_sequence2batch_copy_padding(real* batch,
   if (seq2batch) {
     /* sequence -> batch */
     if (normByTimes) {
-      KeSequence2BatchPadding<1, 1><<< grid, threads, 0, STREAM_DEFAULT >>>(
-              batch, sequence, sequenceStartPositions,
-              sequenceWidth, maxSequenceLength, numSequences);
+      KeSequence2BatchPadding<1, 1><<<grid, threads, 0, STREAM_DEFAULT>>>(
+          batch,
+          sequence,
+          sequenceStartPositions,
+          sequenceWidth,
+          maxSequenceLength,
+          numSequences);
     } else {
-      KeSequence2BatchPadding<0, 1><<< grid, threads, 0, STREAM_DEFAULT >>>(
-              batch, sequence, sequenceStartPositions,
-              sequenceWidth, maxSequenceLength, numSequences);
+      KeSequence2BatchPadding<0, 1><<<grid, threads, 0, STREAM_DEFAULT>>>(
+          batch,
+          sequence,
+          sequenceStartPositions,
+          sequenceWidth,
+          maxSequenceLength,
+          numSequences);
     }
   } else {
     /* batch -> sequence */
     if (normByTimes) {
-      KeSequence2BatchPadding<1, 0><<< grid, threads, 0, STREAM_DEFAULT >>>(
-              batch, sequence, sequenceStartPositions,
-              sequenceWidth, maxSequenceLength, numSequences);
+      KeSequence2BatchPadding<1, 0><<<grid, threads, 0, STREAM_DEFAULT>>>(
+          batch,
+          sequence,
+          sequenceStartPositions,
+          sequenceWidth,
+          maxSequenceLength,
+          numSequences);
     } else {
-      KeSequence2BatchPadding<0, 0><<< grid, threads, 0, STREAM_DEFAULT >>>(
-              batch, sequence, sequenceStartPositions,
-              sequenceWidth, maxSequenceLength, numSequences);
+      KeSequence2BatchPadding<0, 0><<<grid, threads, 0, STREAM_DEFAULT>>>(
+          batch,
+          sequence,
+          sequenceStartPositions,
+          sequenceWidth,
+          maxSequenceLength,
+          numSequences);
     }
   }
 
   CHECK_SYNC("hl_sequence2batch_copy_padding failed");
 }
 
-__device__ inline float my_rsqrt(float x) {
-  return rsqrtf(x);
-}
+__device__ inline float my_rsqrt(float x) { return rsqrtf(x); }
 
-__device__ inline double my_rsqrt(double x) {
-  return rsqrt(x);
-}
+__device__ inline double my_rsqrt(double x) { return rsqrt(x); }
 
 __global__ void KeSequenceAvgForward(real* dst,
                                      real* src,
@@ -327,8 +335,8 @@ __global__ void KeSequenceAvgForward(real* dst,
     for (int i = start; i < end; i++) {
       sum += src[i * width + col];
     }
-    sum = mode == 1 ? sum :
-        (mode == 0 ? sum / seqLength : sum * my_rsqrt((real)seqLength));
+    sum = mode == 1 ? sum : (mode == 0 ? sum / seqLength
+                                       : sum * my_rsqrt((real)seqLength));
     dst[gid] += sum;
   }
 }
@@ -347,10 +355,10 @@ void hl_sequence_avg_forward(real* dst,
   int grid = DIVUP(width * height, 512);
 
   CHECK(mode == 0 || mode == 1 || mode == 2)
-    << "mode error in hl_sequence_avg_forward!";
+      << "mode error in hl_sequence_avg_forward!";
 
-  KeSequenceAvgForward<<< grid, block, 0, STREAM_DEFAULT >>>
-           (dst, src, starts, height, width, mode);
+  KeSequenceAvgForward<<<grid, block, 0, STREAM_DEFAULT>>>(
+      dst, src, starts, height, width, mode);
   CHECK_SYNC("hl_sequence_avg_forward failed");
 }
 
@@ -370,8 +378,8 @@ __global__ void KeSequenceAvgBackward(real* dst,
     int seqLength = end - start;
     if (seqLength == 0) return;
     real grad = src[gid];
-    grad = mode == 1 ? grad :
-        (mode == 0 ? grad / seqLength : grad * my_rsqrt((real)seqLength));
+    grad = mode == 1 ? grad : (mode == 0 ? grad / seqLength
+                                         : grad * my_rsqrt((real)seqLength));
     for (int i = start; i < end; i++) {
       dst[i * width + col] += grad;
     }
@@ -392,9 +400,9 @@ void hl_sequence_avg_backward(real* dst,
   int grid = DIVUP(width * height, 512);
 
   CHECK(mode == 0 || mode == 1 || mode == 2)
-    << "mode error in hl_sequence_avg_backward!";
+      << "mode error in hl_sequence_avg_backward!";
 
-  KeSequenceAvgBackward<<< grid, block, 0, STREAM_DEFAULT >>>
-           (dst, src, starts, height, width, mode);
+  KeSequenceAvgBackward<<<grid, block, 0, STREAM_DEFAULT>>>(
+      dst, src, starts, height, width, mode);
   CHECK_SYNC("hl_sequence_avg_backward failed");
 }

paddle/cuda/src/hl_perturbation_util.cu

@@ -12,13 +12,12 @@ 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 <cmath>
 #include <stdlib.h>
-#include "hl_cuda.h"
-#include "hl_time.h"
+#include <cmath>
 #include "hl_base.h"
+#include "hl_cuda.h"
 #include "hl_perturbation_util.cuh"
+#include "hl_time.h"
 
 #define _USE_MATH_DEFINES
 
@@ -30,10 +29,16 @@ limitations under the License. */
  * centerX, centerY: translation.
  * sourceX, sourceY: output coordinates in the original image.
  */
-__device__ void getTranformCoord(int x, int y, real theta, real scale,
-                                 real tgtCenter, real imgCenter,
-                                 real centerR, real centerC,
-                                 int* sourceX, int* sourceY) {
+__device__ void getTranformCoord(int x,
+                                 int y,
+                                 real theta,
+                                 real scale,
+                                 real tgtCenter,
+                                 real imgCenter,
+                                 real centerR,
+                                 real centerC,
+                                 int* sourceX,
+                                 int* sourceY) {
   real H[4] = {cosf(-theta), -sinf(-theta), sinf(-theta), cosf(-theta)};
 
   // compute coornidates in the rotated and scaled image
@@ -57,11 +62,17 @@ __device__ void getTranformCoord(int x, int y, real theta, real scale,
  * created by Wei Xu (genome), converted by Jiang Wang
  */
 
-__global__ void kSamplingPatches(const real* imgs, real* targets,
-                                 int imgSize, int tgtSize, const int channels,
-                                 int samplingRate, const real* thetas,
-                                 const real* scales, const int* centerRs,
-                                 const int* centerCs, const real padValue,
+__global__ void kSamplingPatches(const real* imgs,
+                                 real* targets,
+                                 int imgSize,
+                                 int tgtSize,
+                                 const int channels,
+                                 int samplingRate,
+                                 const real* thetas,
+                                 const real* scales,
+                                 const int* centerRs,
+                                 const int* centerCs,
+                                 const real padValue,
                                  const int numImages) {
   const int caseIdx = blockIdx.x * 4 + threadIdx.x;
   const int pxIdx = blockIdx.y * 128 + threadIdx.y;
@@ -80,8 +91,15 @@ __global__ void kSamplingPatches(const real* imgs, real* targets,
     const int pxY = pxIdx / tgtSize;
 
     int srcPxX, srcPxY;
-    getTranformCoord(pxX, pxY, thetas[imgIdx], scales[imgIdx], tgtCenter,
-                     imgCenter, centerCs[caseIdx], centerRs[caseIdx], &srcPxX,
+    getTranformCoord(pxX,
+                     pxY,
+                     thetas[imgIdx],
+                     scales[imgIdx],
+                     tgtCenter,
+                     imgCenter,
+                     centerCs[caseIdx],
+                     centerRs[caseIdx],
+                     &srcPxX,
                      &srcPxY);
 
     imgs += (imgIdx * imgPixels + srcPxY * imgSize + srcPxX) * channels;
@@ -100,10 +118,15 @@ __global__ void kSamplingPatches(const real* imgs, real* targets,
  *
  * created by Wei Xu
  */
-void hl_generate_disturb_params(real*& gpuAngle, real*& gpuScaleRatio,
-                                int*& gpuCenterR, int*& gpuCenterC,
-                                int numImages, int imgSize, real rotateAngle,
-                                real scaleRatio, int samplingRate,
+void hl_generate_disturb_params(real*& gpuAngle,
+                                real*& gpuScaleRatio,
+                                int*& gpuCenterR,
+                                int*& gpuCenterC,
+                                int numImages,
+                                int imgSize,
+                                real rotateAngle,
+                                real scaleRatio,
+                                int samplingRate,
                                 bool isTrain) {
   // The number of output samples.
   int numPatches = numImages * samplingRate;
@@ -123,7 +146,8 @@ void hl_generate_disturb_params(real*& gpuAngle, real*& gpuScaleRatio,
     for (int i = 0; i < numImages; i++) {
       r_angle[i] =
           (rotateAngle * M_PI / 180.0) * (rand() / (RAND_MAX + 1.0)  // NOLINT
-                                          - 0.5);
+                                          -
+                                          0.5);
       s_ratio[i] =
           1 + (rand() / (RAND_MAX + 1.0) - 0.5) * scaleRatio;  // NOLINT
     }
@@ -140,8 +164,10 @@ void hl_generate_disturb_params(real*& gpuAngle, real*& gpuScaleRatio,
         int pxY =
             (int)(real(imgSize - 1) * rand() / (RAND_MAX + 1.0));  // NOLINT
 
-        const real H[4] = {cos(-r_angle[i]), -sin(-r_angle[i]),
-                           sin(-r_angle[i]), cos(-r_angle[i])};
+        const real H[4] = {cos(-r_angle[i]),
+                           -sin(-r_angle[i]),
+                           sin(-r_angle[i]),
+                           cos(-r_angle[i])};
         real x = pxX - imgCenter;
         real y = pxY - imgCenter;
         real xx = H[0] * x + H[1] * y;
@@ -185,9 +211,12 @@ void hl_generate_disturb_params(real*& gpuAngle, real*& gpuScaleRatio,
   delete[] center_c;
 }
 
-void hl_conv_random_disturb_with_params(const real* images, int imgSize,
-                                        int tgtSize, int channels,
-                                        int numImages, int samplingRate,
+void hl_conv_random_disturb_with_params(const real* images,
+                                        int imgSize,
+                                        int tgtSize,
+                                        int channels,
+                                        int numImages,
+                                        int samplingRate,
                                         const real* gpuRotationAngle,
                                         const real* gpuScaleRatio,
                                         const int* gpuCenterR,
@@ -202,29 +231,59 @@ void hl_conv_random_disturb_with_params(const real* images, int imgSize,
   dim3 threadsPerBlock(4, 128);
   dim3 numBlocks(DIVUP(numPatches, 4), DIVUP(targetSize, 128));
 
-  kSamplingPatches <<<numBlocks, threadsPerBlock>>>
-      (images, target, imgSize, tgtSize, channels, samplingRate,
-      gpuRotationAngle, gpuScaleRatio, gpuCenterR, gpuCenterC,
-      paddingValue, numImages);
+  kSamplingPatches<<<numBlocks, threadsPerBlock>>>(images,
+                                                   target,
+                                                   imgSize,
+                                                   tgtSize,
+                                                   channels,
+                                                   samplingRate,
+                                                   gpuRotationAngle,
+                                                   gpuScaleRatio,
+                                                   gpuCenterR,
+                                                   gpuCenterC,
+                                                   paddingValue,
+                                                   numImages);
 
   hl_device_synchronize();
 }
 
-void hl_conv_random_disturb(const real* images, int imgSize,
-                            int tgtSize, int channels, int numImages,
-                            real scaleRatio, real rotateAngle,
-                            int samplingRate, real* gpu_r_angle,
-                            real* gpu_s_ratio, int* gpu_center_r,
-                            int* gpu_center_c, int paddingValue,
-                            bool isTrain, real* targets) {
+void hl_conv_random_disturb(const real* images,
+                            int imgSize,
+                            int tgtSize,
+                            int channels,
+                            int numImages,
+                            real scaleRatio,
+                            real rotateAngle,
+                            int samplingRate,
+                            real* gpu_r_angle,
+                            real* gpu_s_ratio,
+                            int* gpu_center_r,
+                            int* gpu_center_c,
+                            int paddingValue,
+                            bool isTrain,
+                            real* targets) {
   // generate the random disturbance sequence and the sampling locations
-  hl_generate_disturb_params(gpu_r_angle, gpu_s_ratio, gpu_center_r,
-                  gpu_center_c, numImages, imgSize, rotateAngle,
-                  scaleRatio, samplingRate, isTrain);
-
-  hl_conv_random_disturb_with_params(
-                  images, imgSize, tgtSize, channels, numImages,
-                  samplingRate, gpu_r_angle, gpu_s_ratio,
-                  gpu_center_r, gpu_center_r, paddingValue,
-                  targets);
+  hl_generate_disturb_params(gpu_r_angle,
+                             gpu_s_ratio,
+                             gpu_center_r,
+                             gpu_center_c,
+                             numImages,
+                             imgSize,
+                             rotateAngle,
+                             scaleRatio,
+                             samplingRate,
+                             isTrain);
+
+  hl_conv_random_disturb_with_params(images,
+                                     imgSize,
+                                     tgtSize,
+                                     channels,
+                                     numImages,
+                                     samplingRate,
+                                     gpu_r_angle,
+                                     gpu_s_ratio,
+                                     gpu_center_r,
+                                     gpu_center_r,
+                                     paddingValue,
+                                     targets);
 }

paddle/cuda/src/hl_table_apply.cu

@@ -12,15 +12,16 @@ 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 "hl_base.h"
-#include "hl_device_functions.cuh"
 #include "hl_cuda.h"
+#include "hl_device_functions.cuh"
 #include "paddle/utils/Logging.h"
 
-template<int blockDimX, int blockDimY, int gridDimX, bool AddRow>
-__global__ void KeMatrixAddRows(real* output, int ldo,
-                                real* table, int ldt,
+template <int blockDimX, int blockDimY, int gridDimX, bool AddRow>
+__global__ void KeMatrixAddRows(real* output,
+                                int ldo,
+                                real* table,
+                                int ldt,
                                 int* ids,
                                 int numSamples,
                                 int tableSize,
@@ -31,8 +32,8 @@ __global__ void KeMatrixAddRows(real* output, int ldo,
   while (idy < numSamples) {
     int tableId = ids[idy];
     if ((0 <= tableId) && (tableId < tableSize)) {
-      real *out = output + idy * ldo;
-      real *tab = table + tableId * ldt;
+      real* out = output + idy * ldo;
+      real* tab = table + tableId * ldt;
       for (int i = idx; i < dim; i += blockDimX) {
         if (AddRow) {
           paddle::paddleAtomicAdd(&tab[i], out[i]);
@@ -45,8 +46,10 @@ __global__ void KeMatrixAddRows(real* output, int ldo,
   }
 }
 
-void hl_matrix_select_rows(real* output, int ldo,
-                           real* table, int ldt,
+void hl_matrix_select_rows(real* output,
+                           int ldo,
+                           real* table,
+                           int ldt,
                            int* ids,
                            int numSamples,
                            int tableSize,
@@ -57,14 +60,16 @@ void hl_matrix_select_rows(real* output, int ldo,
 
   dim3 threads(128, 8);
   dim3 grid(8, 1);
-  KeMatrixAddRows<128, 8, 8, 0><<< grid, threads, 0, STREAM_DEFAULT >>>
-    (output, ldo, table, ldt, ids, numSamples, tableSize, dim);
+  KeMatrixAddRows<128, 8, 8, 0><<<grid, threads, 0, STREAM_DEFAULT>>>(
+      output, ldo, table, ldt, ids, numSamples, tableSize, dim);
 
   CHECK_SYNC("hl_matrix_select_rows failed");
 }
 
-void hl_matrix_add_to_rows(real* table, int ldt,
-                           real* input, int ldi,
+void hl_matrix_add_to_rows(real* table,
+                           int ldt,
+                           real* input,
+                           int ldi,
                            int* ids,
                            int numSamples,
                            int tableSize,
@@ -75,16 +80,15 @@ void hl_matrix_add_to_rows(real* table, int ldt,
 
   dim3 threads(128, 8);
   dim3 grid(8, 1);
-  KeMatrixAddRows<128, 8, 8, 1><<< grid, threads, 0, STREAM_DEFAULT >>>
-    (input, ldi, table, ldt, ids, numSamples, tableSize, dim);
+  KeMatrixAddRows<128, 8, 8, 1><<<grid, threads, 0, STREAM_DEFAULT>>>(
+      input, ldi, table, ldt, ids, numSamples, tableSize, dim);
 
   CHECK_SYNC("hl_matrix_add_to_rows failed");
 }
 
-template<class T, int blockDimX, int gridDimX>
-__global__ void KeVectorSelect(T* dst, int sized,
-                               const T* src, int sizes,
-                               const int* ids, int sizei) {
+template <class T, int blockDimX, int gridDimX>
+__global__ void KeVectorSelect(
+    T* dst, int sized, const T* src, int sizes, const int* ids, int sizei) {
   int idx = threadIdx.x + blockDimX * blockIdx.x;
   while (idx < sizei) {
     int index = ids[idx];
@@ -95,9 +99,8 @@ __global__ void KeVectorSelect(T* dst, int sized,
 }
 
 template <class T>
-void hl_vector_select_from(T* dst, int sized,
-                           const T* src, int sizes,
-                           const int* ids, int sizei) {
+void hl_vector_select_from(
+    T* dst, int sized, const T* src, int sizes, const int* ids, int sizei) {
   CHECK_NOTNULL(dst);
   CHECK_NOTNULL(src);
   CHECK_NOTNULL(ids);
@@ -105,18 +108,17 @@ void hl_vector_select_from(T* dst, int sized,
 
   dim3 threads(512, 1);
   dim3 grid(8, 1);
-  KeVectorSelect<T, 512, 8><<< grid, threads, 0, STREAM_DEFAULT >>>
-    (dst, sized, src, sizes, ids, sizei);
+  KeVectorSelect<T, 512, 8><<<grid, threads, 0, STREAM_DEFAULT>>>(
+      dst, sized, src, sizes, ids, sizei);
 
   CHECK_SYNC("hl_vector_select_from failed");
 }
 
-template
-void hl_vector_select_from(real* dst, int sized,
-                           const real* src, int sizes,
-                           const int* ids, int sizei);
-template
-void hl_vector_select_from(int* dst, int sized,
-                           const int* src, int sizes,
-                           const int* ids, int sizei);
-
+template void hl_vector_select_from(real* dst,
+                                    int sized,
+                                    const real* src,
+                                    int sizes,
+                                    const int* ids,
+                                    int sizei);
+template void hl_vector_select_from(
+    int* dst, int sized, const int* src, int sizes, const int* ids, int sizei);

paddle/framework/CMakeLists.txt

@@ -12,13 +12,15 @@ cc_test(variable_test SRCS variable_test.cc)
 cc_library(scope SRCS scope.cc)
 cc_test(scope_test SRCS scope_test.cc DEPS scope)
 
-proto_library(attr_type SRCS attr_type.proto)
-proto_library(op_proto SRCS op_proto.proto DEPS attr_type)
-proto_library(op_desc SRCS op_desc.proto DEPS attr_type)
+proto_library(attribute_proto SRCS attribute.proto)
+proto_library(op_proto SRCS op_proto.proto DEPS attribute_proto)
+proto_library(op_desc SRCS op_desc.proto DEPS attribute_proto)
 cc_test(op_proto_test SRCS op_proto_test.cc DEPS op_proto protobuf)
 cc_test(op_desc_test SRCS op_desc_test.cc DEPS op_desc protobuf)
 
-cc_library(operator SRCS operator.cc DEPS op_desc device_context tensor scope)
+cc_library(attribute SRCS attribute.cc DEPS op_desc op_proto)
+
+cc_library(operator SRCS operator.cc DEPS op_desc device_context tensor scope attribute)
 cc_test(operator_test SRCS operator_test.cc DEPS operator op_registry)
 
 cc_library(grad_op_builder SRCS grad_op_builder.cc DEPS op_proto operator)
@@ -26,7 +28,7 @@ cc_library(op_registry SRCS op_registry.cc DEPS op_desc grad_op_builder)
 cc_test(op_registry_test SRCS op_registry_test.cc DEPS op_registry)
 cc_test(grad_op_builder_test SRCS grad_op_builder_test.cc DEPS grad_op_builder op_registry add_op)
 
-py_proto_compile(framework_py_proto SRCS attr_type.proto op_proto.proto op_desc.proto)
+py_proto_compile(framework_py_proto SRCS attribute.proto op_proto.proto op_desc.proto)
 # Generate an empty __init__.py to make framework_py_proto as a valid python module.
 add_custom_target(framework_py_proto_init ALL COMMAND ${CMAKE_COMMAND} -E touch __init__.py)
 add_dependencies(framework_py_proto framework_py_proto_init)

paddle/framework/attribute.cc

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+    http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License. */
+
+#include "paddle/framework/attribute.h"
+
+#include <vector>
+
+namespace paddle {
+namespace framework {
+
+template <>
+AttrType AttrTypeID<int>() {
+  return INT;
+}
+template <>
+AttrType AttrTypeID<float>() {
+  return FLOAT;
+}
+template <>
+AttrType AttrTypeID<std::string>() {
+  return STRING;
+}
+template <>
+AttrType AttrTypeID<std::vector<int>>() {
+  return INTS;
+}
+template <>
+AttrType AttrTypeID<std::vector<float>>() {
+  return FLOATS;
+}
+template <>
+AttrType AttrTypeID<std::vector<std::string>>() {
+  return STRINGS;
+}
+
+Attribute GetAttrValue(const AttrDesc& attr_desc) {
+  switch (attr_desc.type()) {
+    case paddle::framework::AttrType::INT: {
+      return attr_desc.i();
+    }
+    case paddle::framework::AttrType::FLOAT: {
+      return attr_desc.f();
+    }
+    case paddle::framework::AttrType::STRING: {
+      return attr_desc.s();
+    }
+    case paddle::framework::AttrType::INTS: {
+      std::vector<int> val(attr_desc.ints_size());
+      for (int i = 0; i < attr_desc.ints_size(); ++i) {
+        val[i] = attr_desc.ints(i);
+      }
+      return val;
+    }
+    case paddle::framework::AttrType::FLOATS: {
+      std::vector<float> val(attr_desc.floats_size());
+      for (int i = 0; i < attr_desc.floats_size(); ++i) {
+        val[i] = attr_desc.floats(i);
+      }
+      return val;
+    }
+    case paddle::framework::AttrType::STRINGS: {
+      std::vector<std::string> val(attr_desc.strings_size());
+      for (int i = 0; i < attr_desc.strings_size(); ++i) {
+        val[i] = attr_desc.strings(i);
+      }
+      return val;
+    }
+  }
+  PADDLE_ENFORCE(false, "Unknown OpDesc::AttrDesc::type !");
+  return boost::blank();
+}
+
+}  // namespace framework
+}  // namespace paddle

paddle/framework/attr_checker.h → paddle/framework/attribute.h

+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+    http://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License. */
+
 #pragma once
 
 #include <boost/variant.hpp>
@@ -6,6 +20,9 @@
 #include <unordered_map>
 #include <unordered_set>
 #include <vector>
+
+#include "paddle/framework/attribute.pb.h"
+#include "paddle/framework/op_desc.pb.h"
 #include "paddle/platform/enforce.h"
 
 namespace paddle {
@@ -14,13 +31,19 @@ namespace framework {
 typedef boost::variant<boost::blank, int, float, std::string, std::vector<int>,
                        std::vector<float>, std::vector<std::string>>
     Attribute;
+
 typedef std::unordered_map<std::string, Attribute> AttributeMap;
 
+template <typename T>
+AttrType AttrTypeID();
+
+Attribute GetAttrValue(const AttrDesc& attr_desc);
+
 // check whether a value(attribute) fit a certain limit
 template <typename T>
 class LargerThanChecker {
  public:
-  LargerThanChecker(T lower_bound) : lower_bound_(lower_bound) {}
+  explicit LargerThanChecker(T lower_bound) : lower_bound_(lower_bound) {}
   void operator()(T& value) const {
     PADDLE_ENFORCE(value > lower_bound_, "larger_than check fail");
   }
@@ -35,7 +58,8 @@ class LargerThanChecker {
 template <typename T>
 class DefaultValueSetter {
  public:
-  DefaultValueSetter(T default_value) : default_value_(default_value) {}
+  explicit DefaultValueSetter(T default_value)
+      : default_value_(default_value) {}
   void operator()(T& value) const { value = default_value_; }
 
  private:
@@ -78,7 +102,8 @@ class TypedAttrChecker {
   typedef std::function<void(T&)> ValueChecker;
 
  public:
-  TypedAttrChecker(const std::string& attr_name) : attr_name_(attr_name) {}
+  explicit TypedAttrChecker(const std::string& attr_name)
+      : attr_name_(attr_name) {}
 
   TypedAttrChecker& InEnum(const std::unordered_set<T>& range) {
     value_checkers_.push_back(EnumInContainer<T>(range));

paddle/framework/attr_type.proto → paddle/framework/attribute.proto

@@ -12,17 +12,17 @@ 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. */
 
-syntax="proto2";
+syntax = "proto2";
 package paddle.framework;
 
 // Attribute Type for paddle's Op.
 // Op contains many attributes. Each type of attributes could be different.
 // The AttrType will be shared between AttrDesc and AttrProto.
 enum AttrType {
-    INT = 0;
-    FLOAT = 1;
-    STRING = 2;
-    INTS = 3;
-    FLOATS = 4;
-    STRINGS = 5;
+  INT = 0;
+  FLOAT = 1;
+  STRING = 2;
+  INTS = 3;
+  FLOATS = 4;
+  STRINGS = 5;
 }
\ No newline at end of file

paddle/framework/backward.cc

@@ -59,19 +59,17 @@ std::shared_ptr<OperatorBase> BackwardRecursive(
   //  If all input gradients of forwarding operator do not need to calculate,
   //  just return an NOP. Not return null ptr because NOP does not take
   //  too much time for calculation, but it is useful for simplifying logic.
-  if (AllInSet(forwardOp.inputs_, OperatorBase::GRAD_VAR_SUFFIX(),
-               no_grad_names)) {
+  if (AllInSet(forwardOp.inputs_, kGradVarSuffix, no_grad_names)) {
     return NOP();
   }
 
   //  All output gradients of forwarding operator do not need to calculate.
   //  Then all input gradients cannot be computed at all, and we put them into
   //  `no_grad_names` set. Return an NOP.
-  if (AllInSet(forwardOp.outputs_, OperatorBase::GRAD_VAR_SUFFIX(),
-               no_grad_names)) {
+  if (AllInSet(forwardOp.outputs_, kGradVarSuffix, no_grad_names)) {
     for (auto& name : forwardOp.inputs_) {
       // Mark all input is not need
-      no_grad_names.insert(name + OperatorBase::GRAD_VAR_SUFFIX());
+      no_grad_names.insert(name + kGradVarSuffix);
     }
     return NOP();
   }
@@ -134,9 +132,9 @@ std::shared_ptr<OperatorBase> BackwardRecursive(
     std::shared_ptr<OperatorBase> grad_op = OpRegistry::CreateGradOp(forwardOp);
     for (std::string& grad_input : grad_op->inputs_) {
       if (no_grad_names.count(grad_input)) {
-        std::string prefix = grad_input.substr(
-            0, grad_input.size() - OperatorBase::GRAD_VAR_SUFFIX().size());
-        grad_input = prefix + OperatorBase::ZERO_VAR_SUFFIX();
+        std::string prefix =
+            grad_input.substr(0, grad_input.size() - kGradVarSuffix.size());
+        grad_input = prefix + kZeroVarSuffix;
 
         // If part of input gradient of that operator is not calculated, fill
         // zero variables to that input gradient.
@@ -147,7 +145,7 @@ std::shared_ptr<OperatorBase> BackwardRecursive(
 
     for (std::string& grad_output : grad_op->outputs_) {
       if (no_grad_names.count(grad_output)) {
-        grad_output = OperatorBase::EMPTY_VAR_NAME();
+        grad_output = kEmptyVarName;
       }
     }
 
@@ -168,14 +166,14 @@ std::shared_ptr<OperatorBase> Backward(
   std::unordered_set<std::string> no_grad_names;
   no_grad_names.reserve(no_grad_vars.size());
 
-  no_grad_names.insert(OperatorBase::EMPTY_VAR_NAME() +
-                       OperatorBase::GRAD_VAR_SUFFIX());
+  no_grad_names.insert(kEmptyVarName + kGradVarSuffix);
 
   for (auto& name : no_grad_vars) {
-    no_grad_names.insert(name + OperatorBase::GRAD_VAR_SUFFIX());
+    no_grad_names.insert(name + kGradVarSuffix);
   }
   size_t uid = 0;
   return BackwardRecursive(forwardOp, no_grad_names, uid);
 }
+
 }  // namespace framework
 }  // namespace paddle

paddle/framework/grad_op_builder.cc

@@ -56,8 +56,7 @@ static void TransOpArg(const OperatorBase* src_op, OperatorBase* dst_op,
 
   for (const auto& arg : src_arg_list) {
     std::string src_name = arg.name();
-    std::string dst_name =
-        is_grad ? src_name + OperatorBase::GRAD_VAR_SUFFIX() : src_name;
+    std::string dst_name = is_grad ? src_name + kGradVarSuffix : src_name;
     (*dst_op->in_out_idxs_)[dst_name] = idx++;
     int src_arg_idx = src_op->in_out_idxs_->at(src_name);
     int src_begin =
@@ -65,10 +64,9 @@ static void TransOpArg(const OperatorBase* src_op, OperatorBase* dst_op,
     int src_end = src_format == nullptr ? src_arg_idx + 1
                                         : src_format->at(src_arg_idx + 1);
     for (int i = src_begin; i < src_end; ++i) {
-      std::string s = is_grad ? src_inout[i] + OperatorBase::GRAD_VAR_SUFFIX()
-                              : arg.ignore_gradient()
-                                    ? OperatorBase::EMPTY_VAR_NAME()
-                                    : src_inout[i];
+      std::string s =
+          is_grad ? src_inout[i] + kGradVarSuffix
+                  : (arg.ignore_gradient() ? kEmptyVarName : src_inout[i]);
       dst_inout.emplace_back(s);
     }
     if (dst_format != nullptr) {

paddle/framework/pybind.cc

@@ -163,8 +163,8 @@ All parameter, weight, gradient are variables in Paddle.
   m.def_submodule(
        "var_names",
        "The module will return special predefined variable name in Paddle")
-      .def("empty", OperatorBase::EMPTY_VAR_NAME)
-      .def("temp", OperatorBase::TMP_VAR_NAME);
+      .def("empty", []() { return kEmptyVarName; })
+      .def("temp", []() { return kTempVarName; });
   // clang-format off
   py::class_<paddle::platform::DeviceContext>(m, "DeviceContext")
       .def_static("create",