Skip to content

P
PaddleDetection

PaddlePaddle / PaddleDetection
大约 2 年 前同步成功

通知 708
Star 11112
Fork 2696
P
PaddleDetection
“9a3f50d8068a9904ca97a9860684bf6783b0d998”上不存在“...cluster/src/git@gitcode.net:BaiXuePrincess/Paddle.git”
提交 aa2f2fb4 编写于 作者: C chengduoZH
浏览文件
操作

remove conflict

上级 a6ef8758 0a8addf8
展开全部 隐藏空白更改
内联 并排
Showing with 4488 addition and 3062 deletion
RELEASE.cn.md
浏览文件 @ aa2f2fb4
# v0.11.0版本
## PaddlePaddle Fluid
- PaddlePaddle发布版本v0.11.0包含一个新的特性*PaddlePaddle Fluid*. Fluid 是设计用来让用户像Pytorch和Tensorflow Eager Execution一样执行程序。在这些系统中，不再有*模型*这个概念，应用也不再包含一个用于描述Operator图或者一系列层的符号描述，而是像通用程序那样描述训练或者预测的过程。而Fluid与PyTorch或Eager Execution的区别在于Fluid不依赖Python提供的控制流，例如 if-else-then或者for，而是提供了基于C++实现的控制流并暴露了对应的用with语法实现的Python接口。例如：
https://github.com/PaddlePaddle/Paddle/blob/3df78ed2a98d37f7ae6725894cc7514effd5664b/python/paddle/v2/fluid/tests/test_while_op.py#L36-L44
- 在v0.11.0版本中，我们提供了一个C++类`Executor`用于运行一个Fluid程序。Executor类似一个解释器。在未来的版本中，我们将提升和优化Executor成为一个调试器，就像GDB。并可能提供一些编译器，这个编译器会读取一个上文所描述的应用然后编译成一个等价的
源代码，这个源代码可以被nvcc编译成可以使用CUDA的二进制，或者被icc编译成可以充分利用Intel CPU的二进制。
## 新特点
* 发布 `PaddlePaddle Fluid`。
* 增加了用于模型预测的C-API。
* 用Fluid API实现了一个简单的GAN的例子。
* 增加了关于性能调优的文档。
* 为`paddle.v2.dataset`下载数据集提供了重试机制.
* C++中使用protobuf-lite替换protobuf减少了二进制的大小。
* 发布了新特性 [Elastic Deep Learning (EDL)](https://github.com/PaddlePaddle/cloud/tree/develop/doc/autoscale/experiment).
* 基于Bazel API利用cmake实现了一个的新的构建系统函数库。
* 当使用编译选项`WITH_MKL=ON`时自动下载和编译Intel® [MKLML](https://github.com/01org/mkl-dnn/releases/download/v0.11/mklml_lnx_2018.0.1.20171007.tgz) 函数库.
* [Intel® MKL-DNN on PaddlePaddle](https://github.com/PaddlePaddle/Paddle/tree/develop/doc/design/mkldnn):
- 完成了 11个 MKL-DNN 层: Convolution, Fully connectivity, Pooling, ReLU, Tanh, ELU, Softmax, BatchNorm, AddTo, Concat, LRN。
- 完成了 3个 MKL-DNN 网络: VGG-19, ResNet-50, GoogleNet
- 基于Intel Skylake 6148 CPU的[性能测试](https://github.com/PaddlePaddle/Paddle/blob/develop/benchmark/IntelOptimizedPaddle.md) : 相对于MKLML有2~3倍的训练加速。
* 增加 [softsign activation](http://www.paddlepaddle.org/docs/develop/documentation/zh/api/v2/config/activation.html#softsign)
* 增加 [dot product layer](http://www.paddlepaddle.org/docs/develop/documentation/zh/api/v2/config/layer.html#dot-prod)
* 增加 [L2 distance layer](http://www.paddlepaddle.org/docs/develop/documentation/zh/api/v2/config/layer.html#l2-distance)
* 增加 [sub-nested sequence layer](http://www.paddlepaddle.org/docs/develop/documentation/zh/api/v2/config/layer.html#sub-nested-seq)
* 增加 [kmax sequence score layer](http://www.paddlepaddle.org/docs/develop/documentation/zh/api/v2/config/layer.html#kmax-sequence-score)
* 增加 [sequence slice layer](http://www.paddlepaddle.org/docs/develop/documentation/zh/api/v2/config/layer.html#seq-slice)
* 增加 [row convolution layer](http://www.paddlepaddle.org/docs/develop/documentation/zh/api/v2/config/layer.html#row-conv)
* 增加移动端友好的网页
## 改进
* 使用一个Python`whl`包即可安装.
* [V2 API可以实现用户定制化评估](https://github.com/PaddlePaddle/models/tree/develop/ltr#训练过程中输出自定义评估指标)。
* 将 `PADDLE_ONLY_CPU` 改为 `PADDLE_WITH_GPU`, 因为我们会支持多种设备。
* 删除了有一些bug的BarrierStat。
* 清理和删除了paddle::Parameter中未使用的函数。
* 删除了ProtoDataProvider。
* Huber loss同时支持回归和分类。
* 为sequence pooling 层增加`stride`参数。
* v2 API自动使用cudnn batch normalization。
* 可以使用一个固定的参数名共享BN层的参数。
* 2D convolution operation支持variable-dimension input特性。
* 重构cmake中关于CUDA的部分并实现自动检测GPU架构的功能。
* 优化网页导航。
## 错误修复
* 修复ROI pooling的Bug. cc9a761
* 修复当label是dense vector是AUC变成0的问题. #5274
* 修复WarpCTC 层的Bug.
# v0.10.0版本 # v0.10.0版本
我们非常高兴发布了PaddlePaddle V0.10.0版，并开发了新的[Python API](http://research.baidu.com/paddlepaddles-new-api-simplifies-deep-learning-programs/)。 我们非常高兴发布了PaddlePaddle V0.10.0版，并开发了新的[Python API](http://research.baidu.com/paddlepaddles-new-api-simplifies-deep-learning-programs/)。
......
RELEASE.md
浏览文件 @ aa2f2fb4
# Release v0.11.0
## PaddlePaddle Fluid
- Release 0.11.0 includes a new feature *PaddlePaddle Fluid*. Fluid is
designed to allow users to program like PyTorch and TensorFlow Eager Execution.
In these systems, there is no longer the concept *model* and applications
do not include a symbolic description of a graph of operators nor a sequence
of layers. Instead, applications look exactly like a usual program that
describes a process of training or inference. The difference between
Fluid and PyTorch or Eager Execution is that Fluid doesn't rely on Python's
control-flow, `if-then-else` nor `for`. Instead, Fluid provides its
C++ implementations and their Python binding using the `with` statement. For an example
https://github.com/PaddlePaddle/Paddle/blob/3df78ed2a98d37f7ae6725894cc7514effd5664b/python/paddle/v2/fluid/tests/test_while_op.py#L36-L44
- In 0.11.0, we provides a C++ class `Executor` to run a Fluid program.
Executor works like an interpreter. In future version, we will improve
`Executor` into a debugger like GDB, and we might provide some compilers,
which, for example, takes an application like the above one, and outputs
an equivalent C++ source program, which can be compiled using
[`nvcc`](http://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html)
to generate binaries that use CUDA, or using
[`icc`](https://software.intel.com/en-us/c-compilers) to generate binaries
that make full use of Intel CPUs.
## New Features
* Release `PaddlePaddle Fluid`.
* Add C-API for model inference
* Use fluid API to create a simple GAN demo.
* Add develop guide about performance tunning.
* Add retry when download `paddle.v2.dataset`.
* Linking protobuf-lite not protobuf in C++. Reduce the binary size.
* Feature [Elastic Deep Learning (EDL)](https://github.com/PaddlePaddle/cloud/tree/develop/doc/autoscale/experiment) released.
* A new style cmake functions for Paddle. It is based on Bazel API.
* Automatically download and compile with Intel® [MKLML](https://github.com/01org/mkl-dnn/releases/download/v0.11/mklml_lnx_2018.0.1.20171007.tgz) library as CBLAS when build `WITH_MKL=ON`.
* [Intel® MKL-DNN on PaddlePaddle](https://github.com/PaddlePaddle/Paddle/tree/develop/doc/design/mkldnn):
- Complete 11 MKL-DNN layers: Convolution, Fully connectivity, Pooling, ReLU, Tanh, ELU, Softmax, BatchNorm, AddTo, Concat, LRN.
- Complete 3 MKL-DNN networks: VGG-19, ResNet-50, GoogleNet
- [Benchmark](https://github.com/PaddlePaddle/Paddle/blob/develop/benchmark/IntelOptimizedPaddle.md) on Intel Skylake 6148 CPU: 2~3x training speedup compared with MKLML.
* Add the [`softsign` activation](http://www.paddlepaddle.org/docs/develop/documentation/zh/api/v2/config/activation.html#softsign).
* Add the [dot product layer](http://www.paddlepaddle.org/docs/develop/documentation/zh/api/v2/config/layer.html#dot-prod).
* Add the [L2 distance layer](http://www.paddlepaddle.org/docs/develop/documentation/zh/api/v2/config/layer.html#l2-distance).
* Add the [sub-nested sequence layer](http://www.paddlepaddle.org/docs/develop/documentation/zh/api/v2/config/layer.html#sub-nested-seq).
* Add the [kmax sequence score layer](http://www.paddlepaddle.org/docs/develop/documentation/zh/api/v2/config/layer.html#kmax-sequence-score).
* Add the [sequence slice layer](http://www.paddlepaddle.org/docs/develop/documentation/zh/api/v2/config/layer.html#seq-slice).
* Add the [row convolution layer](http://www.paddlepaddle.org/docs/develop/documentation/zh/api/v2/config/layer.html#row-conv)
* Add mobile friendly webpages.
## Improvements
* Build and install using a single `whl` package.
* [Custom evaluating in V2 API](https://github.com/PaddlePaddle/models/tree/develop/ltr#训练过程中输出自定义评估指标).
* Change `PADDLE_ONLY_CPU` to `PADDLE_WITH_GPU`, since we will support many kinds of devices.
* Remove buggy BarrierStat.
* Clean and remove unused functions in paddle::Parameter.
* Remove ProtoDataProvider.
* Huber loss supports both regression and classification.
* Add the `stride` parameter for sequence pooling layers.
* Enable v2 API use cudnn batch normalization automatically.
* The BN layer's parameter can be shared by a fixed the parameter name.
* Support variable-dimension input feature for 2D convolution operation.
* Refine cmake about CUDA to automatically detect GPU architecture.
* Improved website navigation.
## Bug Fixes
* Fix bug in ROI pooling. cc9a761
* Fix AUC is zero when label is dense vector. #5274
* Fix bug in WarpCTC layer.
# Release v0.10.0 # Release v0.10.0
We are glad to release version 0.10.0. In this version, we are happy to release the new We are glad to release version 0.10.0. In this version, we are happy to release the new
......
benchmark/IntelOptimizedPaddle.md
浏览文件 @ aa2f2fb4
...@@ -19,6 +19,8 @@ On each machine, we will test and compare the performance of training on single ...@@ -19,6 +19,8 @@ On each machine, we will test and compare the performance of training on single
## Benchmark Model ## Benchmark Model
### Server ### Server
#### Training
Test on batch size 64, 128, 256 on Intel(R) Xeon(R) Gold 6148 CPU @ 2.40GHz 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 Input image size - 3 * 224 * 224, Time: images/second
...@@ -53,5 +55,33 @@ Input image size - 3 * 224 * 224, Time: images/second ...@@ -53,5 +55,33 @@ Input image size - 3 * 224 * 224, Time: images/second
<img src="figs/googlenet-cpu-train.png" width="500"> <img src="figs/googlenet-cpu-train.png" width="500">
#### Inference
Test on batch size 1, 2, 4, 8, 16 on Intel(R) Xeon(R) Gold 6148 CPU @ 2.40GHz
- VGG-19
| BatchSize | 1 | 2 | 4 | 8 | 16 |
|-----------|-------|-------|-------|-------|-------|
| OpenBLAS | 1.07 | 1.08 | 1.06 | 0.88 | 0.65 |
| MKLML | 5.58 | 9.80 | 15.15 | 21.21 | 28.67 |
| MKL-DNN | 75.07 | 88.64 | 82.58 | 92.29 | 96.75 |
- ResNet-50
| BatchSize | 1 | 2 | 4 | 8 | 16 |
|-----------|-------|--------|--------|--------|--------|
| OpenBLAS | 3.35 | 3.19 | 3.09 | 2.55 | 1.96 |
| MKLML | 6.33 | 12.02 | 22.88 | 40.53 | 63.09 |
| MKL-DNN | 107.83| 148.84 | 177.78 | 189.35 | 217.69 |
- GoogLeNet
| BatchSize | 1 | 2 | 4 | 8 | 16 |
|-----------|--------|--------|--------|--------|--------|
| OpenBLAS | 12.04 | 11.31 | 10.00 | 9.07 | 4.34 |
| MKLML | 22.74 | 41.56 | 81.22 | 133.47 | 210.53 |
| MKL-DNN | 175.10 | 272.92 | 450.70 | 512.00 | 600.94 |
### Laptop ### Laptop
TBD TBD
cmake/cblas.cmake
浏览文件 @ aa2f2fb4
...@@ -17,7 +17,7 @@ if(WITH_MKLML AND MKLML_INC_DIR AND MKLML_LIB) ...@@ -17,7 +17,7 @@ if(WITH_MKLML AND MKLML_INC_DIR AND MKLML_LIB)
set(CBLAS_INC_DIR ${MKLML_INC_DIR}) set(CBLAS_INC_DIR ${MKLML_INC_DIR})
set(CBLAS_LIBRARIES ${MKLML_LIB}) set(CBLAS_LIBRARIES ${MKLML_LIB})
add_definitions(-DPADDLE_USE_MKLML) add_definitions(-DPADDLE_WITH_MKLML)
add_definitions(-DLAPACK_FOUND) add_definitions(-DLAPACK_FOUND)
message(STATUS "Found cblas and lapack in MKLML " message(STATUS "Found cblas and lapack in MKLML "
......
cmake/external/mkldnn.cmake
浏览文件 @ aa2f2fb4
...@@ -67,5 +67,5 @@ ADD_LIBRARY(mkldnn SHARED IMPORTED GLOBAL) ...@@ -67,5 +67,5 @@ ADD_LIBRARY(mkldnn SHARED IMPORTED GLOBAL)
SET_PROPERTY(TARGET mkldnn PROPERTY IMPORTED_LOCATION ${MKLDNN_LIB}) SET_PROPERTY(TARGET mkldnn PROPERTY IMPORTED_LOCATION ${MKLDNN_LIB})
ADD_DEPENDENCIES(mkldnn ${MKLDNN_PROJECT}) ADD_DEPENDENCIES(mkldnn ${MKLDNN_PROJECT})
MESSAGE(STATUS "MKLDNN library: ${MKLDNN_LIB}") MESSAGE(STATUS "MKLDNN library: ${MKLDNN_LIB}")
add_definitions(-DPADDLE_USE_MKLDNN) add_definitions(-DPADDLE_WITH_MKLDNN)
LIST(APPEND external_project_dependencies mkldnn) LIST(APPEND external_project_dependencies mkldnn)
doc/howto/dev/contribute_to_paddle_cn.md
浏览文件 @ aa2f2fb4
...@@ -76,18 +76,18 @@ no changes added to commit (use "git add" and/or "git commit -a") ...@@ -76,18 +76,18 @@ no changes added to commit (use "git add" and/or "git commit -a")
## 构建和测试 ## 构建和测试
编译 PaddlePaddle 的源码以及生成文档需要多种开发工具。为了方便大家，我们的标准开发流程是把这些工具都装进一个Docker image，称为*开发镜像*，通常名字是 `paddle:dev`。然后所有用 `cmake && make` 的地方（比如IDE配置里）都用 `docker run paddle:dev`来代替。 编译 PaddlePaddle 的源码以及生成文档需要多种开发工具。为了方便大家，我们的标准开发流程是把这些工具都装进一个Docker image，称为*开发镜像*，通常名字是 `paddle:latest-dev` 或者 `paddle:[version tag]-dev` 如 `paddle:0.11.0-dev`。然后所有用 `cmake && make` 的地方（比如IDE配置里）都用 `docker run paddle:latest-dev`来代替。
如要build这个开发镜像，在源码目录树的根目录中运行： 如要build这个开发镜像，在源码目录树的根目录中运行：
```bash ```bash
➜ docker build -t paddle:dev . ➜ docker build -t paddle:latest-dev .
``` ```
随后可以用这个开发镜像开始build PaddlePaddle的源码。比如如果要build一个不依赖GPU，但是支持AVX指令集，并且包括unit tests的PaddlePaddle，可以： 随后可以用这个开发镜像开始build PaddlePaddle的源码。比如如果要build一个不依赖GPU，但是支持AVX指令集，并且包括unit tests的PaddlePaddle，可以：
```bash ```bash
➜ docker run -v $(pwd):/paddle -e "WITH_GPU=OFF" -e "WITH_AVX=ON" -e "WITH_TEST=ON" paddle:dev ➜ docker run -v $(pwd):/paddle -e "WITH_GPU=OFF" -e "WITH_AVX=ON" -e "WITH_TESTING=ON" paddle:latest-dev
``` ```
这个过程除了编译PaddlePaddle为 `./build/libpaddle.so`，并且输出一个 `./build/paddle.deb`文件之外，还会输出一个 `build/Dockerfile`。我们只需要运行下面命令把编译好的PaddlePaddle打包成一个*生产镜像*（`paddle:prod`）： 这个过程除了编译PaddlePaddle为 `./build/libpaddle.so`，并且输出一个 `./build/paddle.deb`文件之外，还会输出一个 `build/Dockerfile`。我们只需要运行下面命令把编译好的PaddlePaddle打包成一个*生产镜像*（`paddle:prod`）：
...@@ -99,7 +99,7 @@ no changes added to commit (use "git add" and/or "git commit -a") ...@@ -99,7 +99,7 @@ no changes added to commit (use "git add" and/or "git commit -a")
如果要运行所有的单元测试，可以用如下命令： 如果要运行所有的单元测试，可以用如下命令：
```bash ```bash
➜ docker run -it -v $(pwd):/paddle paddle:dev bash -c "cd /paddle/build && ctest" ➜ docker run -it -v $(pwd):/paddle paddle:latest-dev bash -c "cd /paddle/build && ctest"
``` ```
关于构建和测试的更多信息，请参见[这篇文档](https://github.com/PaddlePaddle/Paddle/blob/develop/doc/getstarted/build_and_install/docker_install_cn.rst)。 关于构建和测试的更多信息，请参见[这篇文档](https://github.com/PaddlePaddle/Paddle/blob/develop/doc/getstarted/build_and_install/docker_install_cn.rst)。
......
doc/howto/dev/new_op_cn.md
浏览文件 @ aa2f2fb4
# 如何写新的Operator # 如何写新的Operator
- [概念简介](#概念简介) - [概念简介](#概念简介)
- [实现C++类](#实现C++类) - [实现C++类](#实现c类)
- [定义ProtoMaker类](#定义ProtoMaker类) - [定义ProtoMaker类](#定义protomaker类)
- [定义Operator类](#定义Operator类) - [定义Operator类](#定义operator类)
- [定义OpKernel类](#定义OpKernel类) - [定义OpKernel类](#定义opkernel类)
- [注册Operator](#注册Operator) - [注册Operator](#注册operator)
- [编译](#编译) - [编译](#编译)
- [绑定Python](#绑定Python) - [绑定Python](#绑定python)
- [实现单元测试](#实现单元测试) - [实现单元测试](#实现单元测试)
- [前向Operator单测](#前向Operator单测) - [前向Operator单测](#前向operator单测)
- [反向Operator单测](#反向Operator单测) - [反向Operator单测](#反向operator单测)
- [编译和执行](#编译和执行) - [编译和执行](#编译和执行)
- [注意事项](#注意事项)
## 概念简介 ## 概念简介
...@@ -30,8 +31,8 @@ ...@@ -30,8 +31,8 @@
-------------- | :---------------------- -------------- | :----------------------
OpProtoMake定义 | `.cc`文件，Backward Op不需要定义OpProtoMake OpProtoMake定义 | `.cc`文件，Backward Op不需要定义OpProtoMake
Op定义 | `.cc`文件 Op定义 | `.cc`文件
Kernel实现 | CPU、GPU共享Kernel实现在`.h`文件中，否则，CPU 实现在`.cc`文件中，GPU 实现在`.cu`文件中。 Kernel实现 | CPU、CUDA共享Kernel实现在`.h`文件中，否则，CPU 实现在`.cc`文件中，CUDA 实现在`.cu`文件中。
注册Op | Op注册实现在`.cc`文件；Kernel注册CPU实现在`.cc`文件中，GPU实现在`.cu`文件中 注册Op | Op注册实现在`.cc`文件；Kernel注册CPU实现在`.cc`文件中，CUDA实现在`.cu`文件中
实现新的op都添加至目录[paddle/operators](https://github.com/PaddlePaddle/Paddle/tree/develop/paddle/operators)下，文件命名以`*_op.h`（如有） 、 `*_op.cc` 、`*_op.cu`（如有）结尾。**系统会根据文件名自动构建op和其对应的Python扩展。** 实现新的op都添加至目录[paddle/operators](https://github.com/PaddlePaddle/Paddle/tree/develop/paddle/operators)下，文件命名以`*_op.h`（如有） 、 `*_op.cc` 、`*_op.cu`（如有）结尾。**系统会根据文件名自动构建op和其对应的Python扩展。**
...@@ -43,7 +44,7 @@ Kernel实现 | CPU、GPU共享Kernel实现在`.h`文件中，否则，CPU ...@@ -43,7 +44,7 @@ Kernel实现 | CPU、GPU共享Kernel实现在`.h`文件中，否则，CPU
## 实现C++类 ## 实现C++类
### 1. 定义ProtoMaker类 ### 定义ProtoMaker类
矩阵乘法的公式：$Out = X * Y$, 可见该计算由两个输入，一个输出组成。 矩阵乘法的公式：$Out = X * Y$, 可见该计算由两个输入，一个输出组成。
...@@ -100,7 +101,7 @@ The equation is: Out = scale*X ...@@ -100,7 +101,7 @@ The equation is: Out = scale*X
- `AddAttr<AttrType>("scale", "...").SetDefault(1.0);` : 增加`scale`系数，作为参数属性，并且设置默认值为1.0。 - `AddAttr<AttrType>("scale", "...").SetDefault(1.0);` : 增加`scale`系数，作为参数属性，并且设置默认值为1.0。
### 2. 定义Operator类 ### 定义Operator类
下面的点实现了MulOp的定义： 下面的点实现了MulOp的定义：
...@@ -149,11 +150,11 @@ MulOp(const std::string &type, const framework::VariableNameMap &inputs, ...@@ -149,11 +150,11 @@ MulOp(const std::string &type, const framework::VariableNameMap &inputs,
通常`OpProtoMaker`和`Op`类的定义写在`.cc`文件中，和下面将要介绍的注册函数一起放在`.cc`中 通常`OpProtoMaker`和`Op`类的定义写在`.cc`文件中，和下面将要介绍的注册函数一起放在`.cc`中
### 3. 定义OpKernel类 ### 定义OpKernel类
`MulKernel`继承自`framework::OpKernel`，带有下面两个模板参数: `MulKernel`继承自`framework::OpKernel`，带有下面两个模板参数:
- `typename Place`: 表示设备类型，不同设备(CPU、GPU)共享同一个Kernel时，需加该模板参数，不共享则不加，一个不共享的例子是[`OnehotCrossEntropyOpKernel`](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/operators/cross_entropy_op.h#L43)。 - `typename DeviceContext`: 表示设备类型，不同设备(CPU、CUDA)共享同一个Kernel时，需加该模板参数，不共享则不加，一个不共享的例子是[`OnehotCrossEntropyOpKernel`](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/operators/cross_entropy_op.h#L43)。
- `typename T` : 表示数据类型，如`float`, `double`等。 - `typename T` : 表示数据类型，如`float`, `double`等。
...@@ -165,7 +166,7 @@ MulOp(const std::string &type, const framework::VariableNameMap &inputs, ...@@ -165,7 +166,7 @@ MulOp(const std::string &type, const framework::VariableNameMap &inputs,
下面是 `MulKernel` `Compute`的实现： 下面是 `MulKernel` `Compute`的实现：
```cpp ```cpp
template <typename Place, typename T> template <typename DeviceContext, typename T>
class MulKernel : public framework::OpKernel { class MulKernel : public framework::OpKernel {
public: public:
void Compute(const framework::ExecutionContext& context) const override { void Compute(const framework::ExecutionContext& context) const override {
...@@ -173,33 +174,32 @@ MulOp(const std::string &type, const framework::VariableNameMap &inputs, ...@@ -173,33 +174,32 @@ MulOp(const std::string &type, const framework::VariableNameMap &inputs,
auto* Y = context.Input<Tensor>("Y"); auto* Y = context.Input<Tensor>("Y");
auto* Z = context.Output<Tensor>("Out"); auto* Z = context.Output<Tensor>("Out");
Z->mutable_data<T>(context.GetPlace()); Z->mutable_data<T>(context.GetPlace());
auto* device_context = auto& device_context = context.template device_context<DeviceContext>();
const_cast<platform::DeviceContext*>(context.device_context_); math::matmul<DeviceContext, T>(*X, false, *Y, false, 1, Z, 0, device_context);
math::matmul<Place, T>(*X, false, *Y, false, 1, Z, 0, device_context);
} }
}; };
``` ```
需要注意：**不同设备(CPU、GPU)共享一个Op定义，是否则共享同一个`OpKernel`，取决于`Compute`调用的函数是否支持不同设备。** 需要注意：**不同设备(CPU、CUDA)共享一个Op定义，是否则共享同一个`OpKernel`，取决于`Compute`调用的函数是否支持不同设备。**
`MulOp`的CPU、GPU实现共享同一个`Kernel`。`OpKernel`不共享的例子可以参考：[`OnehotCrossEntropyOpKernel`](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/operators/cross_entropy_op.h#L43)。 `MulOp`的CPU、CUDA实现共享同一个`Kernel`。`OpKernel`不共享的例子可以参考：[`OnehotCrossEntropyOpKernel`](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/operators/cross_entropy_op.h#L43)。
为了使`OpKernel`的计算过程书写更加简单，并且CPU、GPU的代码可以复用，我们通常借助 Eigen unsupported Tensor模块来实现`Compute`接口。关于在PaddlePaddle中如何使用Eigen库，请参考[使用文档](https://github.com/PaddlePaddle/Paddle/blob/develop/doc/howto/dev/use_eigen_cn.md)。 为了使`OpKernel`的计算过程书写更加简单，并且CPU、CUDA的代码可以复用，我们通常借助 Eigen unsupported Tensor模块来实现`Compute`接口。关于在PaddlePaddle中如何使用Eigen库，请参考[使用文档](https://github.com/PaddlePaddle/Paddle/blob/develop/doc/howto/dev/use_eigen_cn.md)。
到此，前向Op实现完成。接下来，需要在`.cc`文件中注册该op和kernel。 到此，前向Op实现完成。接下来，需要在`.cc`文件中注册该op和kernel。
反向Op类的定义，反向OpKernel的定义与前向Op类似，这里不再赘述。**但需注意反向Op没有`ProtoMaker`**。 反向Op类的定义，反向OpKernel的定义与前向Op类似，这里不再赘述。**但需注意反向Op没有`ProtoMaker`**。
### 4. 注册Operator ### 注册Operator
- 在`.cc`文件中注册前向、反向Op类，注册CPU Kernel。 - 在`.cc`文件中注册前向、反向Op类，注册CPU Kernel。
```cpp ```cpp
namespace ops = paddle::operators; namespace ops = paddle::operators;
REGISTER_OP(mul, ops::MulOp, ops::MulOpMaker, mul_grad, ops::MulOpGrad); REGISTER_OP(mul, ops::MulOp, ops::MulOpMaker, mul_grad, ops::MulOpGrad);
REGISTER_OP_CPU_KERNEL(mul, ops::MulKernel<paddle::platform::CPUPlace, float>); REGISTER_OP_CPU_KERNEL(mul, ops::MulKernel<paddle::platform::CPUDeviceContext, float>);
REGISTER_OP_CPU_KERNEL(mul_grad, REGISTER_OP_CPU_KERNEL(mul_grad,
ops::MulGradKernel<paddle::platform::CPUPlace, float>); ops::MulGradKernel<paddle::platform::CPUDeviceContext, float>);
``` ```
在上面的代码中： 在上面的代码中：
...@@ -209,20 +209,20 @@ MulOp(const std::string &type, const framework::VariableNameMap &inputs, ...@@ -209,20 +209,20 @@ MulOp(const std::string &type, const framework::VariableNameMap &inputs,
- `REGISTER_OP_CPU_KERNEL` ：注册`ops::MulKernel`类，并特化模板参数为`paddle::platform::CPUPlace`和`float`类型，同理，注册`ops::MulGradKernel`类。 - `REGISTER_OP_CPU_KERNEL` ：注册`ops::MulKernel`类，并特化模板参数为`paddle::platform::CPUPlace`和`float`类型，同理，注册`ops::MulGradKernel`类。
- 在 `.cu`文件中注册GPU Kernel。 - 在 `.cu`文件中注册CUDA Kernel。
- 请注意，如果GPU Kernel的实现基于Eigen unsupported模块，那么在 `.cu`的开始请加上宏定义 `#define EIGEN_USE_GPU`，代码示例如下： - 请注意，如果CUDA Kernel的实现基于Eigen unsupported模块，那么在 `.cu`的开始请加上宏定义 `#define EIGEN_USE_GPU`，代码示例如下：
```cpp ```cpp
// if use Eigen unsupported module before include head files // if use Eigen unsupported module before include head files
// #define EIGEN_USE_GPU #define EIGEN_USE_GPU
namespace ops = paddle::operators; namespace ops = paddle::operators;
REGISTER_OP_GPU_KERNEL(mul, ops::MulKernel<paddle::platform::GPUPlace, float>); REGISTER_OP_CUDA_KERNEL(mul, ops::MulKernel<paddle::platform::CUDADeviceContext, float>);
REGISTER_OP_GPU_KERNEL(mul_grad, REGISTER_OP_CUDA_KERNEL(mul_grad,
ops::MulGradKernel<paddle::platform::GPUPlace, float>); ops::MulGradKernel<paddle::platform::CUDADeviceContext, float>);
``` ```
### 5. 编译 ### 编译
运行下面命令可以进行编译： 运行下面命令可以进行编译：
...@@ -236,71 +236,57 @@ make mul_op ...@@ -236,71 +236,57 @@ make mul_op
## 实现单元测试 ## 实现单元测试
单测包括对比前向Op不同设备(CPU、GPU)的实现、对比反向OP不同设备(CPU、GPU)的实现、反向Op的梯度测试。下面介绍介绍[`MulOp`的单元测试](https://github.com/PaddlePaddle/Paddle/blob/develop/python/paddle/v2/framework/tests/test_mul_op.py)。 单测包括对比前向Op不同设备(CPU、CUDA)的实现、对比反向OP不同设备(CPU、CUDA)的实现、反向Op的梯度测试。下面介绍介绍[`MulOp`的单元测试](https://github.com/PaddlePaddle/Paddle/blob/develop/python/paddle/v2/framework/tests/test_mul_op.py)。
### 前向Operator单元测试 ### 前向Operator单测
前向Op单元测试继承自`unittest.TestCase`，并定义元类`__metaclass__ = OpTestMeta`。各项更加具体的单元测试在`OpTestMeta`里完成。测试前向Operator，需要： Op单元测试继承自`OpTest`。各项更加具体的单元测试在`TestMulOp`里完成。测试Operator，需要：
1. 在`setUp`函数定义输入、输出，以及相关的属性参数。 1. 在`setUp`函数定义输入、输出，以及相关的属性参数。
2. 生成随机的输入数据。 2. 生成随机的输入数据。
3. 在Python脚本中实现与前向operator相同的计算逻辑，得到输出值，与operator前向计算的输出进行对比。 3. 在Python脚本中实现与前向operator相同的计算逻辑，得到输出值，与operator前向计算的输出进行对比。
4. 反向计算已经自动集成进测试框架，直接调用相应接口即可。
```python ```python
import unittest import unittest
import numpy as np import numpy as np
from gradient_checker import GradientChecker, create_op from op_test import OpTest
from op_test_util import OpTestMeta
class TestMulOp(unittest.TestCase):
__metaclass__ = OpTestMeta
class TestMulOp(OpTest):
def setUp(self): def setUp(self):
self.type = "mul" self.op_type = "mul"
self.inputs = { self.inputs = {
'X': np.random.random((32, 84)).astype("float32"), 'X': np.random.random((32, 84)).astype("float32"),
'Y': np.random.random((84, 100)).astype("float32") 'Y': np.random.random((84, 100)).astype("float32")
} }
self.outputs = {'Out': np.dot(self.inputs['X'], self.inputs['Y'])} self.outputs = {'Out': np.dot(self.inputs['X'], self.inputs['Y'])}
```
上面的代码首先导入依赖的包，下面是对`setUp`函数中操作的重要变量的详细解释：
- `self.type = "mul" ` : 定义类型，与operator注册时注册的类型一致。
- `self.inputs` : 定义输入，类型为`numpy.array`，并初始化。
- `self.outputs` : 定义输出，并在Python脚本中完成与operator同样的计算逻辑，返回Python端的计算结果。
def test_check_output(self):
self.check_output()
### 反向Operator单元测试 def test_check_grad_normal(self):
self.check_grad(['X', 'Y'], 'Out', max_relative_error=0.5)
反向Op单元测试继承自`GradientChecker`，而`GradientChecker`继承自`unittest.TestCase`，因此，**反向单元测试函数需要以`test_`开头**。 def test_check_grad_ingore_x(self):
self.check_grad(
['Y'], 'Out', max_relative_error=0.5, no_grad_set=set("X"))
```python def test_check_grad_ingore_y(self):
class TestMulGradOp(GradientChecker): self.check_grad(
def setUp(self): ['X'], 'Out', max_relative_error=0.5, no_grad_set=set('Y'))
self.op = create_op("mul") ```
self.inputs = {
'X': np.random.random((32, 84)).astype("float32"),
'Y': np.random.random((84, 100)).astype("float32")
}
def test_check_grad_normal(self):
# mul op will enlarge the relative error
self.check_grad(['X', 'Y'], 'Out', max_relative_error=0.5)
def test_check_grad_ingore_x(self): 上面的代码首先导入依赖的包，下面是对`setUp`函数中操作的重要变量的详细解释：
self.check_grad(
['Y'], 'Out', max_relative_error=0.5, no_grad_set=set("X"))
def test_check_grad_ingore_y(self): - `self.op_type = "mul" ` : 定义类型，与operator注册时注册的类型一致。
self.check_grad( - `self.inputs` : 定义输入，类型为`numpy.array`，并初始化。
['X'], 'Out', max_relative_error=0.5, no_grad_set=set('Y')) - `self.outputs` : 定义输出，并在Python脚本中完成与operator同样的计算逻辑，返回Python端的计算结果。
```
下面解释代码中一些关键的地方: ### 反向operator单测
- 调用`create_op("mul")`创建反向Op对应的前向Op。 而反向测试中：
- `test_check_grad_normal`中调用`check_grad`使用数值法检测梯度正确性和稳定性。 - `test_check_grad_normal`中调用`check_grad`使用数值法检测梯度正确性和稳定性。
- 第一个参数`["X", "Y"]` : 指定对输入变量`X`、`Y`做梯度检测。 - 第一个参数`["X", "Y"]` : 指定对输入变量`X`、`Y`做梯度检测。
- 第二个参数`"Out"` : 指定前向网络最终的输出目标变量`Out`。 - 第二个参数`"Out"` : 指定前向网络最终的输出目标变量`Out`。
...@@ -308,7 +294,7 @@ class TestMulGradOp(GradientChecker): ...@@ -308,7 +294,7 @@ class TestMulGradOp(GradientChecker):
- `test_check_grad_ingore_x`和`test_check_grad_ingore_y`分支用来测试只需要计算一个输入梯度的情况。 - `test_check_grad_ingore_x`和`test_check_grad_ingore_y`分支用来测试只需要计算一个输入梯度的情况。
### 编译和执行单元测试 ### 编译和执行
`python/paddle/v2/framework/tests` 目录下新增的 `test_*.py` 单元测试会被自动加入工程进行编译。 `python/paddle/v2/framework/tests` 目录下新增的 `test_*.py` 单元测试会被自动加入工程进行编译。
...@@ -328,5 +314,5 @@ ctest -R test_mul_op ...@@ -328,5 +314,5 @@ ctest -R test_mul_op
- 为每个Op创建单独的`*_op.h`（如有）、`*_op.cc`和`*_op.cu`（如有）。不允许一个文件中包含多个Op，这将会导致编译出错。 - 为每个Op创建单独的`*_op.h`（如有）、`*_op.cc`和`*_op.cu`（如有）。不允许一个文件中包含多个Op，这将会导致编译出错。
- 注册Op时的类型名，需要和该Op的名字一样。即不允许在`A_op.cc`里面，注册`REGISTER_OP(B, ...)`等，这将会导致单元测试出错。 - 注册Op时的类型名，需要和该Op的名字一样。即不允许在`A_op.cc`里面，注册`REGISTER_OP(B, ...)`等，这将会导致单元测试出错。
- 如果Op没有实现GPU Kernel，请不要创建空的`*_op.cu`，这将会导致单元测试出错。 - 如果Op没有实现CUDA Kernel，请不要创建空的`*_op.cu`，这将会导致单元测试出错。
- 如果多个Op依赖一些共用的函数，可以创建非`*_op.*`格式的文件来存放，如`gather.h`文件。 - 如果多个Op依赖一些共用的函数，可以创建非`*_op.*`格式的文件来存放，如`gather.h`文件。
doc/howto/dev/new_op_en.md
浏览文件 @ aa2f2fb4
# How to write a new operator # How to write a new operator
- [Background](#background) - [Background](#background)
- [Implementing C++ Types](#implementing-c++-types) - [Implementing C++ Types](#implementing-c-types)
- [Defining ProtoMaker](#defining-protoMaker) - [Defining ProtoMaker](#defining-protomaker)
- [Defining Operator](#defining-operator) - [Defining Operator](#defining-operator)
- [Registering Operator](#registering-operator) - [Registering Operator](#registering-operator)
- [Compilation](#compilation) - [Compilation](#compilation)
...@@ -28,8 +28,8 @@ An operator can be differentiated by whether in has kernel methods. An operator ...@@ -28,8 +28,8 @@ An operator can be differentiated by whether in has kernel methods. An operator
-------------- | :---------------------- -------------- | :----------------------
OpProtoMake definition | `.cc`files, Backward Op does not need an OpProtoMake interface. OpProtoMake definition | `.cc`files, Backward Op does not need an OpProtoMake interface.
Op definition | `.cc` files Op definition | `.cc` files
Kernel implementation | The kernel methods shared between CPU and GPU are defined in `.h` files. CPU-specific kernels live in `.cc` files, while GPU-specific kernels are implemented in `.cu`files. Kernel implementation | The kernel methods shared between CPU and CUDA are defined in `.h` files. CPU-specific kernels live in `.cc` files, while CUDA-specific kernels are implemented in `.cu`files.
Registering the Op | Ops are registered in `.cc` files; For Kernel registration, `.cc` files contain the CPU implementation, while `.cu` files contain the GPU implementation. Registering the Op | Ops are registered in `.cc` files; For Kernel registration, `.cc` files contain the CPU implementation, while `.cu` files contain the CUDA implementation.
New Operator implementations are added to the list [paddle/operators](https://github.com/PaddlePaddle/Paddle/tree/develop/paddle/operators), with file names in the format `*_op.h` (if applicable), `*_op.cc`, `*_op.cu` (if applicable).** The system will use the naming scheme to automatically build operators and their corresponding Python extensions. ** New Operator implementations are added to the list [paddle/operators](https://github.com/PaddlePaddle/Paddle/tree/develop/paddle/operators), with file names in the format `*_op.h` (if applicable), `*_op.cc`, `*_op.cu` (if applicable).** The system will use the naming scheme to automatically build operators and their corresponding Python extensions. **
...@@ -41,7 +41,7 @@ Let's take matrix multiplication operator, [MulOp](https://github.com/PaddlePadd ...@@ -41,7 +41,7 @@ Let's take matrix multiplication operator, [MulOp](https://github.com/PaddlePadd
## Implementing C++ Types ## Implementing C++ Types
### 1. Defining Class ProtoMaker ### Defining ProtoMaker
Matrix Multiplication can be written as $Out = X * Y$, meaning that the operation consists of two inputs and pne output. Matrix Multiplication can be written as $Out = X * Y$, meaning that the operation consists of two inputs and pne output.
...@@ -98,7 +98,7 @@ There are two changes in this example: ...@@ -98,7 +98,7 @@ There are two changes in this example:
- `AddAttr<AttrType>("scale", "...").SetDefault(1.0);` adds `scale`constant as an attribute, and sets the default value to 1.0. - `AddAttr<AttrType>("scale", "...").SetDefault(1.0);` adds `scale`constant as an attribute, and sets the default value to 1.0.
### 2. Defining Operator ### Defining Operator
The following code defines the interface for MulOp: The following code defines the interface for MulOp:
...@@ -147,11 +147,11 @@ MulOp(const std::string &type, const framework::VariableNameMap &inputs, ...@@ -147,11 +147,11 @@ MulOp(const std::string &type, const framework::VariableNameMap &inputs,
Usually `OpProtoMaker` and `Op`'s type definitions are written in `.cc` files, which also include the registration methods introduced later. Usually `OpProtoMaker` and `Op`'s type definitions are written in `.cc` files, which also include the registration methods introduced later.
### 3. Defining OpKernel ### Defining OpKernel
`MulKernel` inherits `framework::OpKernel`, which includes the following templates: `MulKernel` inherits `framework::OpKernel`, which includes the following templates:
- `typename Place` denotes device type. When different devices, namely the CPU and the GPU, share the same kernel, this template needs to be added. If they don't share kernels, this must not be added. An example of a non-sharing kernel is [`OnehotCrossEntropyOpKernel`](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/operators/cross_entropy_op.h#L43). - `typename DeviceContext` denotes device context type. When different devices, namely the CPUDeviceContext and the CUDADeviceContext, share the same kernel, this template needs to be added. If they don't share kernels, this must not be added. An example of a non-sharing kernel is [`OnehotCrossEntropyOpKernel`](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/operators/cross_entropy_op.h#L43).
- `typename T` denotes data type, such as `float` or `double`. - `typename T` denotes data type, such as `float` or `double`.
...@@ -163,7 +163,7 @@ Usually `OpProtoMaker` and `Op`'s type definitions are written in `.cc` files, w ...@@ -163,7 +163,7 @@ Usually `OpProtoMaker` and `Op`'s type definitions are written in `.cc` files, w
`MulKernel`'s implementation of `Compute` is as follows: `MulKernel`'s implementation of `Compute` is as follows:
```cpp ```cpp
template <typename Place, typename T> template <typename DeviceContext, typename T>
class MulKernel : public framework::OpKernel { class MulKernel : public framework::OpKernel {
public: public:
void Compute(const framework::ExecutionContext& context) const override { void Compute(const framework::ExecutionContext& context) const override {
...@@ -171,16 +171,15 @@ Usually `OpProtoMaker` and `Op`'s type definitions are written in `.cc` files, w ...@@ -171,16 +171,15 @@ Usually `OpProtoMaker` and `Op`'s type definitions are written in `.cc` files, w
auto* Y = context.Input<Tensor>("Y"); auto* Y = context.Input<Tensor>("Y");
auto* Z = context.Output<Tensor>("Out"); auto* Z = context.Output<Tensor>("Out");
Z->mutable_data<T>(context.GetPlace()); Z->mutable_data<T>(context.GetPlace());
auto* device_context = auto& device_context = context.template device_context<DeviceContext>();
const_cast<platform::DeviceContext*>(context.device_context_); math::matmul<DeviceContext, T>(*X, false, *Y, false, 1, Z, 0, device_context);
math::matmul<Place, T>(*X, false, *Y, false, 1, Z, 0, device_context);
} }
}; };
``` ```
Note that **different devices (CPU, GPU)share an Op definition; whether or not they share the same `OpKernel` depends on whether `Compute` calls functions that support both devices.** Note that **different devices (CPU, CUDA)share an Op definition; whether or not they share the same `OpKernel` depends on whether `Compute` calls functions that support both devices.**
`MulOp`'s CPU and GPU share the same `Kernel`. A non-sharing `OpKernel` example can be seen in [`OnehotCrossEntropyOpKernel`](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/operators/cross_entropy_op.h#L43). `MulOp`'s CPU and CUDA share the same `Kernel`. A non-sharing `OpKernel` example can be seen in [`OnehotCrossEntropyOpKernel`](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/operators/cross_entropy_op.h#L43).
To ease the writing of `OpKernel` compute, and for reusing code cross-device, [`Eigen-unsupported Tensor`](https://bitbucket.org/eigen/eigen/src/default/unsupported/Eigen/CXX11/src/Tensor/README.md?fileviewer=file-view-default) module is used to implement `Compute` interface. To learn about how the Eigen library is used in PaddlePaddle, please see [usage document](https://github.com/PaddlePaddle/Paddle/blob/develop/doc/howto/dev/use_eigen_cn.md). To ease the writing of `OpKernel` compute, and for reusing code cross-device, [`Eigen-unsupported Tensor`](https://bitbucket.org/eigen/eigen/src/default/unsupported/Eigen/CXX11/src/Tensor/README.md?fileviewer=file-view-default) module is used to implement `Compute` interface. To learn about how the Eigen library is used in PaddlePaddle, please see [usage document](https://github.com/PaddlePaddle/Paddle/blob/develop/doc/howto/dev/use_eigen_cn.md).
...@@ -189,16 +188,16 @@ This concludes the forward implementation of an operator. Next its operation and ...@@ -189,16 +188,16 @@ This concludes the forward implementation of an operator. Next its operation and
The definition of its corresponding backward operator, if applicable, is similar to that of an forward operator. **Note that a backward operator does not include a `ProtoMaker`**. The definition of its corresponding backward operator, if applicable, is similar to that of an forward operator. **Note that a backward operator does not include a `ProtoMaker`**.
### 4. Registering Operator ### Registering Operator
- In `.cc` files, register forward and backward operator classes and the CPU kernel. - In `.cc` files, register forward and backward operator classes and the CPU kernel.
```cpp ```cpp
namespace ops = paddle::operators; namespace ops = paddle::operators;
REGISTER_OP(mul, ops::MulOp, ops::MulOpMaker, mul_grad, ops::MulOpGrad); REGISTER_OP(mul, ops::MulOp, ops::MulOpMaker, mul_grad, ops::MulOpGrad);
REGISTER_OP_CPU_KERNEL(mul, ops::MulKernel<paddle::platform::CPUPlace, float>); REGISTER_OP_CPU_KERNEL(mul, ops::MulKernel<paddle::platform::CPUDeviceContext, float>);
REGISTER_OP_CPU_KERNEL(mul_grad, REGISTER_OP_CPU_KERNEL(mul_grad,
ops::MulGradKernel<paddle::platform::CPUPlace, float>); ops::MulGradKernel<paddle::platform::CPUDeviceContext, float>);
``` ```
In that code block, In that code block,
...@@ -208,20 +207,20 @@ The definition of its corresponding backward operator, if applicable, is similar ...@@ -208,20 +207,20 @@ The definition of its corresponding backward operator, if applicable, is similar
- `REGISTER_OP_CPU_KERNEL` registers `ops::MulKernel` class and specialized template types `paddle::platform::CPUPlace` and `float`, which also registers `ops::MulGradKernel`. - `REGISTER_OP_CPU_KERNEL` registers `ops::MulKernel` class and specialized template types `paddle::platform::CPUPlace` and `float`, which also registers `ops::MulGradKernel`.
- Registering GPU Kernel in `.cu` files - Registering CUDA Kernel in `.cu` files
- Note that if GPU Kernel is implemented using the `Eigen unsupported` module, then on top of `.cu`, a macro definition `#define EIGEN_USE_GPU` is needed, such as - Note that if CUDA Kernel is implemented using the `Eigen unsupported` module, then on top of `.cu`, a macro definition `#define EIGEN_USE_GPU` is needed, such as
```cpp ```cpp
// if use Eigen unsupported module before include head files // if use Eigen unsupported module before include head files
#define EIGEN_USE_GPU #define EIGEN_USE_GPU
namespace ops = paddle::operators; namespace ops = paddle::operators;
REGISTER_OP_GPU_KERNEL(mul, ops::MulKernel<paddle::platform::GPUPlace, float>); REGISTER_OP_CUDA_KERNEL(mul, ops::MulKernel<paddle::platform::CUDADeviceContext, float>);
REGISTER_OP_GPU_KERNEL(mul_grad, REGISTER_OP_CUDA_KERNEL(mul_grad,
ops::MulGradKernel<paddle::platform::GPUPlace, float>); ops::MulGradKernel<paddle::platform::CUDADeviceContext, float>);
``` ```
### 5. Compilation ### Compilation
Run the following commands to compile. Run the following commands to compile.
...@@ -253,62 +252,51 @@ A forward operator unit test inherits `unittest.TestCase` and defines metaclass ...@@ -253,62 +252,51 @@ A forward operator unit test inherits `unittest.TestCase` and defines metaclass
2. Generating random input data. 2. Generating random input data.
3. Implementing the same computation logic in a Python script: 3. Implementing the same computation logic in a Python script.
4. Call check gradient function to check the backward operator.
```python ```python
import unittest import unittest
import numpy as np import numpy as np
from gradient_checker import GradientChecker, create_op from op_test import OpTest
from op_test_util import OpTestMeta
class TestMulOp(unittest.TestCase):
__metaclass__ = OpTestMeta
class TestMulOp(OpTest):
def setUp(self): def setUp(self):
self.type = "mul" self.op_type = "mul"
self.inputs = { self.inputs = {
'X': np.random.random((32, 84)).astype("float32"), 'X': np.random.random((32, 84)).astype("float32"),
'Y': np.random.random((84, 100)).astype("float32") 'Y': np.random.random((84, 100)).astype("float32")
} }
self.outputs = {'Out': np.dot(self.inputs['X'], self.inputs['Y'])} self.outputs = {'Out': np.dot(self.inputs['X'], self.inputs['Y'])}
def test_check_output(self):
self.check_output()
def test_check_grad_normal(self):
self.check_grad(['X', 'Y'], 'Out', max_relative_error=0.5)
def test_check_grad_ingore_x(self):
self.check_grad(
['Y'], 'Out', max_relative_error=0.5, no_grad_set=set("X"))
def test_check_grad_ingore_y(self):
self.check_grad(
['X'], 'Out', max_relative_error=0.5, no_grad_set=set('Y'))
``` ```
Get its output, and compare it with the forward operator's own output. Get its output, and compare it with the forward operator's own output.
The code above first loads required packages. In addition, we have The code above first loads required packages. In addition, we have
- `self.type = "mul" ` defines the type that is identical to what the operator's registered type. - `self.op_type = "mul" ` defines the type that is identical to what the operator's registered type.
- `self.inputs` defines input, with type `numpy.array` and initializes it. - `self.inputs` defines input, with type `numpy.array` and initializes it.
- `self.outputs` defines output and completes the same operator computation in the Python script, and returns its result from the Python script. - `self.outputs` defines output and completes the same operator computation in the Python script, and returns its result from the Python script.
### Testing Backward Operators ### Testing Backward Operators
A backward operator unit test inherits `GradientChecker`, which inherits `unittest.TestCase`. As a result, **a backward operator unit test needs to be have the prefix `test_`**. Some key points in checking gradient above include:
```python
class TestMulGradOp(GradientChecker):
def setUp(self):
self.op = create_op("mul")
self.inputs = {
'X': np.random.random((32, 84)).astype("float32"),
'Y': np.random.random((84, 100)).astype("float32")
}
def test_check_grad_normal(self):
# mul op will enlarge the relative error
self.check_grad(['X', 'Y'], 'Out', max_relative_error=0.5)
def test_check_grad_ingore_x(self):
self.check_grad(
['Y'], 'Out', max_relative_error=0.5, no_grad_set=set("X"))
def test_check_grad_ingore_y(self):
self.check_grad(
['X'], 'Out', max_relative_error=0.5, no_grad_set=set('Y'))
```
Some key points in the code above include:
- `create_op("mul")` creates the backward operator's corresponding forward operator.
- `test_normal` calls `check_grad` to validate scaling tests' correctness and stability through numeric methods. - `test_normal` calls `check_grad` to validate scaling tests' correctness and stability through numeric methods.
- The first variable `["X", "Y"]` appoints `X` and `Y` to be scale tested. - The first variable `["X", "Y"]` appoints `X` and `Y` to be scale tested.
- The second variable `"Out"` points to the network's final output target `Out`. - The second variable `"Out"` points to the network's final output target `Out`.
...@@ -338,5 +326,5 @@ ctest -R test_mul_op ...@@ -338,5 +326,5 @@ ctest -R test_mul_op
- Every `*_op.h` (if applicable), `*_op.cc`, and `*_op.cu` (if applicable) must be created for a unique Op. Compiling will fail if multiple operators are included per file. - Every `*_op.h` (if applicable), `*_op.cc`, and `*_op.cu` (if applicable) must be created for a unique Op. Compiling will fail if multiple operators are included per file.
- The type with which an operator is registered needs to be identical to the Op's name. Registering `REGISTER_OP(B, ...)` in `A_op.cc` will cause unit testing failures. - The type with which an operator is registered needs to be identical to the Op's name. Registering `REGISTER_OP(B, ...)` in `A_op.cc` will cause unit testing failures.
- If the operator does not implement a GPU kernel, please refrain from creating an empty `*_op.cu` file, or else unit tests will fail. - If the operator does not implement a CUDA kernel, please refrain from creating an empty `*_op.cu` file, or else unit tests will fail.
- If multiple operators rely on some shared methods, a file NOT named `*_op.*` can be created to store them, such as `gather.h`. - If multiple operators rely on some shared methods, a file NOT named `*_op.*` can be created to store them, such as `gather.h`.
doc/howto/read_source.md 0 → 100644
浏览文件 @ aa2f2fb4
# PaddlePaddle Fluid Source Code Overview
Examples: https://github.com/PaddlePaddle/Paddle/tree/develop/python/paddle/v2/fluid/tests/book
Core: https://github.com/PaddlePaddle/Paddle/tree/develop/paddle/framework
Operator: https://github.com/PaddlePaddle/Paddle/tree/develop/paddle/operators
Optimizer: https://github.com/PaddlePaddle/Paddle/tree/develop/paddle/optimizer
Memory: https://github.com/PaddlePaddle/Paddle/tree/develop/paddle/memory
# Compile Time
The following **defines** the NN. The definition goes into this [protocol buffer](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/framework/framework.proto).
```python
x = fluid.layers.data(name='x', shape=[13], dtype='float32')
y = fluid.layers.data(name='y', shape=[1], dtype='float32')
y_predict = fluid.layers.fc(input=x, size=1, act=None)
cost = fluid.layers.square_error_cost(input=y_predict, label=y)
avg_cost = fluid.layers.mean(x=cost)
sgd_optimizer = fluid.optimizer.SGD(learning_rate=0.001)
sgd_optimizer.minimize(avg_cost)
```
- Variables: `x`, `y`, `y_predict`, `cost` and `avg_cost`. [Python](https://github.com/PaddlePaddle/Paddle/blob/develop/python/paddle/v2/fluid/framework.py#L93)
- Layers: `fluid.layers.data`, `fluid.layers.fc` and `fluid.layers.mean` are layers. [Python](https://github.com/PaddlePaddle/Paddle/blob/develop/python/paddle/v2/fluid/layers.py)
- Every Layer has one or more operators and variables/parameters
- All the operators are defined at [`paddle/operators/`](https://github.com/PaddlePaddle/Paddle/tree/develop/paddle/operators). Other worth-looking files:
- Base class: [`paddle/framework/operator.h`](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/framework/operator.h)
- Operator Registration: [`paddle/framework/op_registry.h`](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/framework/op_registry.h)
- Operator Lookup: [`paddle/framework/op_info.h`](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/framework/op_info.h)
- Optimizer: `fluid.optimizer.SGD`. It does the following
- Add backward operators. [[Python](https://github.com/PaddlePaddle/Paddle/blob/develop/python/paddle/v2/fluid/backward.py), [C++](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/framework/backward.cc)]
- Add optimizer operators. [[Python](https://github.com/PaddlePaddle/Paddle/blob/develop/python/paddle/v2/fluid/optimizer.py), [C++](https://github.com/PaddlePaddle/Paddle/tree/develop/paddle/optimizer)]
# Run Time
The following **evaluates** the NN. Instantiates all the variables, operators.
```python
place = fluid.CPUPlace()
feeder = fluid.DataFeeder(place=place, feed_list=[x, y])
exe = fluid.Executor(place)
# Allocate memory. Initialize Parameter.
exe.run(fluid.default_startup_program())
# Allocate memory. Do computation.
exe.run(fluid.default_main_program(),
feed=feeder.feed(data),
fetch_list=[avg_cost])
```
- Place: `place`. one of CPU, GPU or FPGA. [C++](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/platform/place.h)
- The device handle are at [paddle/platform/device_context.h](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/platform/device_context.h)
- Executor: `fluid.Executor(place)`. [[Python](https://github.com/PaddlePaddle/Paddle/blob/develop/python/paddle/v2/fluid/executor.py), [C++](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/framework/executor.cc)]
- Feeds the data: `feed=feeder.feed(data)`
- Evaluates all the operators
- Fetches the result: `fetch_list=[avg_cost]`
- Other worth looking files:
- Scope: [paddle/framework/scope.h](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/framework/scope.h). Where all the variables live
- Variable: [paddle/framework/variable.h](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/framework/variable.h). Where all the data (most likely tensors) live
- Tensor: [paddle/framework/tensor.h](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/framework/tensor.h). Where we allocate memory through [`paddle/memory/`](https://github.com/PaddlePaddle/Paddle/tree/develop/paddle/memory)
paddle/api/CMakeLists.txt
浏览文件 @ aa2f2fb4
...@@ -25,8 +25,18 @@ FILE(GLOB PY_PADDLE_PYTHON_FILES ${PADDLE_SOURCE_DIR}/paddle/py_paddle/*.py) ...@@ -25,8 +25,18 @@ FILE(GLOB PY_PADDLE_PYTHON_FILES ${PADDLE_SOURCE_DIR}/paddle/py_paddle/*.py)
SET_SOURCE_FILES_PROPERTIES(Paddle.i PROPERTIES CPLUSPLUS ON) SET_SOURCE_FILES_PROPERTIES(Paddle.i PROPERTIES CPLUSPLUS ON)
SET(SWIG_NEED_FLAGS
-ftls-model=global-dynamic
-Wno-parentheses-equality
-Wno-self-assign
-Wno-maybe-uninitialized
-Wno-missing-field-initializers)
FOREACH(flag ${SWIG_NEED_FLAGS})
safe_set_cxxflag(SWIG_CXX_FLAGS ${flag})
ENDFOREACH()
SET(CMAKE_SWIG_OUTDIR ${CMAKE_CURRENT_BINARY_DIR}) SET(CMAKE_SWIG_OUTDIR ${CMAKE_CURRENT_BINARY_DIR})
SET(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wno-parentheses-equality -Wno-missing-field-initializers -Wno-self-assign -ftls-model=global-dynamic") SET(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${SWIG_CXX_FLAGS}")
SET(SWIG_MODULE_swig_paddle_EXTRA_DEPS SET(SWIG_MODULE_swig_paddle_EXTRA_DEPS
paddle_parameter paddle_parameter
......
paddle/framework/backward.cc
浏览文件 @ aa2f2fb4
...@@ -190,8 +190,9 @@ static std::unique_ptr<OperatorBase> BackwardRecursive( ...@@ -190,8 +190,9 @@ static std::unique_ptr<OperatorBase> BackwardRecursive(
// collect all the offset for each alias, // collect all the offset for each alias,
// insert a sum operator to add all aliases to output // insert a sum operator to add all aliases to output
insert_position.push_back( insert_position.push_back(
{dup_op.back(), OpRegistry::CreateOp("sum", {{"X", dup_outputs}}, {dup_op.back(),
{{"Out", {name}}}, {})}); OpRegistry::CreateOp("sum", {{"X", dup_outputs}}, {{"Out", {name}}},
AttributeMap{})});
} }
// make sure the inserted `sum` ops follow the BFS order. // make sure the inserted `sum` ops follow the BFS order.
...@@ -216,7 +217,8 @@ static std::unique_ptr<OperatorBase> BackwardRecursive( ...@@ -216,7 +217,8 @@ static std::unique_ptr<OperatorBase> BackwardRecursive(
// If part of input gradient of that operator is not calculated, fill // If part of input gradient of that operator is not calculated, fill
// zero variables to that input gradient. // zero variables to that input gradient.
net->AppendOp(OpRegistry::CreateOp("fill_zeros_like", {{"X", {prefix}}}, net->AppendOp(OpRegistry::CreateOp("fill_zeros_like", {{"X", {prefix}}},
{{"Y", {grad_input}}}, {})); {{"Y", {grad_input}}},
AttributeMap{}));
} }
return false; return false;
}); });
...@@ -392,8 +394,9 @@ std::vector<std::unique_ptr<OpDescBind>> MakeOpGrad( ...@@ -392,8 +394,9 @@ std::vector<std::unique_ptr<OpDescBind>> MakeOpGrad(
0, in_name.size() - sizeof(kGradVarSuffix) / sizeof(char) + 1); 0, in_name.size() - sizeof(kGradVarSuffix) / sizeof(char) + 1);
std::string new_name = prefix + kZeroVarSuffix; std::string new_name = prefix + kZeroVarSuffix;
desc->Rename(in_name, new_name); desc->Rename(in_name, new_name);
std::unique_ptr<OpDescBind> fill_zeros_op(new OpDescBind( std::unique_ptr<OpDescBind> fill_zeros_op(
"fill_zeros_like", {{"X", {prefix}}}, {{"Y", {new_name}}}, {})); new OpDescBind("fill_zeros_like", {{"X", {prefix}}},
{{"Y", {new_name}}}, AttributeMap{}));
pending_fill_zeros_ops.push_back(std::move(fill_zeros_op)); pending_fill_zeros_ops.push_back(std::move(fill_zeros_op));
} }
} }
...@@ -483,8 +486,9 @@ std::vector<std::unique_ptr<OpDescBind>> MakeBlockBackward( ...@@ -483,8 +486,9 @@ std::vector<std::unique_ptr<OpDescBind>> MakeBlockBackward(
sum_op_inputs.emplace_back(new_name); sum_op_inputs.emplace_back(new_name);
next_g_name = sum_op_inputs.back(); next_g_name = sum_op_inputs.back();
} }
std::unique_ptr<OpDescBind> sum_op(new OpDescBind( std::unique_ptr<OpDescBind> sum_op(
"sum", {{"X", sum_op_inputs}}, {{"Out", {out_name}}}, {})); new OpDescBind("sum", {{"X", sum_op_inputs}}, {{"Out", {out_name}}},
AttributeMap{}));
pending_sum_ops.push_back({dup_op.back(), std::move(sum_op)}); pending_sum_ops.push_back({dup_op.back(), std::move(sum_op)});
} }
} }
......
paddle/framework/backward_test.cc
浏览文件 @ aa2f2fb4
...@@ -106,15 +106,15 @@ class FcOp : public operators::NetOp { ...@@ -106,15 +106,15 @@ class FcOp : public operators::NetOp {
FcOp(const std::string &type, const VariableNameMap &inputs, FcOp(const std::string &type, const VariableNameMap &inputs,
const VariableNameMap &outputs, const AttributeMap &attrs) const VariableNameMap &outputs, const AttributeMap &attrs)
: NetOp(type, inputs, outputs, attrs) { : NetOp(type, inputs, outputs, attrs) {
AppendOp(OpRegistry::CreateOp("mul", AppendOp(OpRegistry::CreateOp(
{{"X", {Input("X")}}, {"Y", {Input("W")}}}, "mul", {{"X", {Input("X")}}, {"Y", {Input("W")}}},
{{"Out", {Output("mul_result")}}}, {})); {{"Out", {Output("mul_result")}}}, AttributeMap{}));
auto input_b = Inputs("b"); auto input_b = Inputs("b");
std::string before_act = "mul_result"; std::string before_act = "mul_result";
if (input_b.size() != 0) { if (input_b.size() != 0) {
AppendOp(OpRegistry::CreateOp( AppendOp(OpRegistry::CreateOp(
"rowwise_add", {{"X", {Output("mul_result")}}, {"b", {input_b[0]}}}, "rowwise_add", {{"X", {Output("mul_result")}}, {"b", {input_b[0]}}},
{{"Out", {Output("add_result")}}}, {})); {{"Out", {Output("add_result")}}}, AttributeMap{}));
before_act = "add_result"; before_act = "add_result";
} else { } else {
auto out_varname = Output("add_result"); auto out_varname = Output("add_result");
...@@ -124,7 +124,7 @@ class FcOp : public operators::NetOp { ...@@ -124,7 +124,7 @@ class FcOp : public operators::NetOp {
} }
AppendOp(OpRegistry::CreateOp("sigmoid", {{"X", {Output(before_act)}}}, AppendOp(OpRegistry::CreateOp("sigmoid", {{"X", {Output(before_act)}}},
{{"Out", {Output("Out")}}}, {})); {{"Out", {Output("Out")}}}, AttributeMap{}));
CompleteAddOp(false); CompleteAddOp(false);
} }
}; };
...@@ -278,8 +278,9 @@ REGISTER_OPERATOR(scale, f::NoneOp); ...@@ -278,8 +278,9 @@ REGISTER_OPERATOR(scale, f::NoneOp);
REGISTER_OP_CPU_KERNEL(scale, f::NoneKernel<paddle::platform::CPUPlace, float>); REGISTER_OP_CPU_KERNEL(scale, f::NoneKernel<paddle::platform::CPUPlace, float>);
TEST(Backward, simple_op_not_need_grad) { TEST(Backward, simple_op_not_need_grad) {
auto fwd = f::OpRegistry::CreateOp( auto fwd =
"rowwise_add", {{"X", {"x"}}, {"b", {"b"}}}, {{"Out", {"out"}}}, {}); f::OpRegistry::CreateOp("rowwise_add", {{"X", {"x"}}, {"b", {"b"}}},
{{"Out", {"out"}}}, f::AttributeMap{});
ASSERT_NE(fwd, nullptr); ASSERT_NE(fwd, nullptr);
auto gop = f::Backward(*fwd, {"x"}); auto gop = f::Backward(*fwd, {"x"});
ASSERT_EQ(gop->Output(f::GradVarName("X")), f::kEmptyVarName); ASSERT_EQ(gop->Output(f::GradVarName("X")), f::kEmptyVarName);
...@@ -296,9 +297,10 @@ TEST(Backward, net_fc_backward_normal) { ...@@ -296,9 +297,10 @@ TEST(Backward, net_fc_backward_normal) {
{{"mul_result", {"mul_res"}}, {{"mul_result", {"mul_res"}},
{"add_result", {"add_re"}}, {"add_result", {"add_re"}},
{"Out", {"out"}}}, {"Out", {"out"}}},
{}); f::AttributeMap{});
ASSERT_NE(fwd, nullptr); ASSERT_NE(fwd, nullptr);
std::shared_ptr<f::OperatorBase> gop = f::Backward(*fwd, {}); std::shared_ptr<f::OperatorBase> gop =
f::Backward(*fwd, std::unordered_set<std::string>{});
ASSERT_TRUE(gop->IsNetOp()); ASSERT_TRUE(gop->IsNetOp());
auto net = static_cast<ops::NetOp *>(gop.get()); auto net = static_cast<ops::NetOp *>(gop.get());
...@@ -322,9 +324,10 @@ TEST(Backward, net_fc_backward_not_have_b) { ...@@ -322,9 +324,10 @@ TEST(Backward, net_fc_backward_not_have_b) {
{{"mul_result", {"mul_res"}}, {{"mul_result", {"mul_res"}},
{"add_result", {"add_res"}}, {"add_result", {"add_res"}},
{"Out", {"tmp"}}}, {"Out", {"tmp"}}},
{}); f::AttributeMap{});
ASSERT_NE(fwd, nullptr); ASSERT_NE(fwd, nullptr);
std::shared_ptr<f::OperatorBase> gop = f::Backward(*fwd, {}); std::shared_ptr<f::OperatorBase> gop =
f::Backward(*fwd, std::unordered_set<std::string>{});
ASSERT_TRUE(gop->IsNetOp()); ASSERT_TRUE(gop->IsNetOp());
auto net = static_cast<ops::NetOp *>(gop.get()); auto net = static_cast<ops::NetOp *>(gop.get());
...@@ -346,13 +349,13 @@ TEST(Backward, net_input_of_network_not_need_grad) { ...@@ -346,13 +349,13 @@ TEST(Backward, net_input_of_network_not_need_grad) {
{{"mul_result", {"mul_tmp_0"}}, {{"mul_result", {"mul_tmp_0"}},
{"add_result", {"add_tmp_0"}}, {"add_result", {"add_tmp_0"}},
{"Out", {"hidden0"}}}, {"Out", {"hidden0"}}},
{})); f::AttributeMap{}));
net.AppendOp(f::OpRegistry::CreateOp( net.AppendOp(f::OpRegistry::CreateOp(
"fc", {{"X", {"hidden0"}}, {"W", {"W2"}}, {"b", {"b2"}}}, "fc", {{"X", {"hidden0"}}, {"W", {"W2"}}, {"b", {"b2"}}},
{{"mul_result", {"mul_tmp_1"}}, {{"mul_result", {"mul_tmp_1"}},
{"add_result", {"add_tmp_1"}}, {"add_result", {"add_tmp_1"}},
{"Out", {"hidden1"}}}, {"Out", {"hidden1"}}},
{})); f::AttributeMap{}));
net.CompleteAddOp(); net.CompleteAddOp();
auto bwd = Backward(net, {"x"}); // x@GRAD is not need. auto bwd = Backward(net, {"x"}); // x@GRAD is not need.
ASSERT_TRUE(bwd->IsNetOp()); ASSERT_TRUE(bwd->IsNetOp());
...@@ -381,12 +384,13 @@ TEST(Backward, net_input_of_network_not_need_grad) { ...@@ -381,12 +384,13 @@ TEST(Backward, net_input_of_network_not_need_grad) {
TEST(Backward, net_shared_weight) { TEST(Backward, net_shared_weight) {
ops::NetOp net; ops::NetOp net;
net.AppendOp(f::OpRegistry::CreateOp("mul", {{"X", {"x"}}, {"Y", {"w"}}}, net.AppendOp(f::OpRegistry::CreateOp("mul", {{"X", {"x"}}, {"Y", {"w"}}},
{{"Out", {"out"}}}, {})); {{"Out", {"out"}}}, f::AttributeMap{}));
net.AppendOp(f::OpRegistry::CreateOp("mul", {{"X", {"out"}}, {"Y", {"w"}}}, net.AppendOp(f::OpRegistry::CreateOp("mul", {{"X", {"out"}}, {"Y", {"w"}}},
{{"Out", {"FinalOut"}}}, {})); {{"Out", {"FinalOut"}}},
f::AttributeMap{}));
net.CompleteAddOp(); net.CompleteAddOp();
auto bwd = f::Backward(net, {}); auto bwd = f::Backward(net, std::unordered_set<std::string>{});
ASSERT_TRUE(bwd->IsNetOp()); ASSERT_TRUE(bwd->IsNetOp());
auto bwd_net = static_cast<ops::NetOp *>(bwd.get()); auto bwd_net = static_cast<ops::NetOp *>(bwd.get());
ASSERT_EQ(3UL, bwd_net->ops_.size()); ASSERT_EQ(3UL, bwd_net->ops_.size());
...@@ -394,8 +398,9 @@ TEST(Backward, net_shared_weight) { ...@@ -394,8 +398,9 @@ TEST(Backward, net_shared_weight) {
} }
TEST(Backward, op_all_input_are_not_need) { TEST(Backward, op_all_input_are_not_need) {
auto fwd = f::OpRegistry::CreateOp( auto fwd =
"rowwise_add", {{"X", {"x"}}, {"b", {"b"}}}, {{"Out", {"out"}}}, {}); f::OpRegistry::CreateOp("rowwise_add", {{"X", {"x"}}, {"b", {"b"}}},
{{"Out", {"out"}}}, f::AttributeMap{});
auto backward = f::Backward(*fwd, {"x", "b"}); auto backward = f::Backward(*fwd, {"x", "b"});
ASSERT_TRUE(backward->IsNetOp()); ASSERT_TRUE(backward->IsNetOp());
auto net = static_cast<ops::NetOp *>(backward.get()); auto net = static_cast<ops::NetOp *>(backward.get());
...@@ -403,8 +408,9 @@ TEST(Backward, op_all_input_are_not_need) { ...@@ -403,8 +408,9 @@ TEST(Backward, op_all_input_are_not_need) {
} }
TEST(Backward, op_all_output_are_not_need) { TEST(Backward, op_all_output_are_not_need) {
auto fwd = f::OpRegistry::CreateOp( auto fwd =
"rowwise_add", {{"X", {"x"}}, {"b", {"b"}}}, {{"Out", {"out"}}}, {}); f::OpRegistry::CreateOp("rowwise_add", {{"X", {"x"}}, {"b", {"b"}}},
{{"Out", {"out"}}}, f::AttributeMap{});
auto backward = f::Backward(*fwd, {"out"}); auto backward = f::Backward(*fwd, {"out"});
ASSERT_TRUE(backward->IsNetOp()); ASSERT_TRUE(backward->IsNetOp());
auto net = static_cast<ops::NetOp *>(backward.get()); auto net = static_cast<ops::NetOp *>(backward.get());
...@@ -412,8 +418,9 @@ TEST(Backward, op_all_output_are_not_need) { ...@@ -412,8 +418,9 @@ TEST(Backward, op_all_output_are_not_need) {
} }
TEST(Backward, op_part_of_output_are_not_need) { TEST(Backward, op_part_of_output_are_not_need) {
auto fwd = f::OpRegistry::CreateOp("many_output_op", {{"x", {"X"}}}, auto fwd =
{{"y", {"Y"}}, {"z", {"Z"}}}, {}); f::OpRegistry::CreateOp("many_output_op", {{"x", {"X"}}},
{{"y", {"Y"}}, {"z", {"Z"}}}, f::AttributeMap{});
auto backward = f::Backward(*fwd, {"Z"}); auto backward = f::Backward(*fwd, {"Z"});
ASSERT_TRUE(backward->IsNetOp()); ASSERT_TRUE(backward->IsNetOp());
auto net = static_cast<ops::NetOp *>(backward.get()); auto net = static_cast<ops::NetOp *>(backward.get());
...@@ -437,7 +444,7 @@ TEST(Backward, op_part_of_output_are_not_need) { ...@@ -437,7 +444,7 @@ TEST(Backward, op_part_of_output_are_not_need) {
TEST(Backward, op_part_of_input_are_not_need) { TEST(Backward, op_part_of_input_are_not_need) {
auto fwd = f::OpRegistry::CreateOp("mul", {{"X", {"a"}}, {"Y", {"b"}}}, auto fwd = f::OpRegistry::CreateOp("mul", {{"X", {"a"}}, {"Y", {"b"}}},
{{"Out", {"out"}}}, {}); {{"Out", {"out"}}}, f::AttributeMap{});
auto backward = f::Backward(*fwd, {"a"}); auto backward = f::Backward(*fwd, {"a"});
auto &grad_mul = *backward; auto &grad_mul = *backward;
ASSERT_EQ(grad_mul.Type(), "mul_grad"); ASSERT_EQ(grad_mul.Type(), "mul_grad");
...@@ -458,19 +465,19 @@ TEST(Backward, linear_net_intermediate_variable_has_no_grad) { ...@@ -458,19 +465,19 @@ TEST(Backward, linear_net_intermediate_variable_has_no_grad) {
{{"mul_result", {"mul_out1"}}, {{"mul_result", {"mul_out1"}},
{"add_result", {"add_out1"}}, {"add_result", {"add_out1"}},
{"Out", {"out1"}}}, {"Out", {"out1"}}},
{})); f::AttributeMap{}));
net.AppendOp(f::OpRegistry::CreateOp( net.AppendOp(f::OpRegistry::CreateOp(
"fc", {{"X", {"out1"}}, {"W", {"w2"}}, {"b", {"b2"}}}, "fc", {{"X", {"out1"}}, {"W", {"w2"}}, {"b", {"b2"}}},
{{"mul_result", {"mul_out2"}}, {{"mul_result", {"mul_out2"}},
{"add_result", {"tmp_out2"}}, {"add_result", {"tmp_out2"}},
{"Out", {"out2"}}}, {"Out", {"out2"}}},
{})); f::AttributeMap{}));
net.AppendOp(f::OpRegistry::CreateOp( net.AppendOp(f::OpRegistry::CreateOp(
"fc", {{"X", {"out2"}}, {"W", {"w3"}}, {"b", {"b3"}}}, "fc", {{"X", {"out2"}}, {"W", {"w3"}}, {"b", {"b3"}}},
{{"mul_result", {"mul_out3"}}, {{"mul_result", {"mul_out3"}},
{"add_result", {"tmp_out3"}}, {"add_result", {"tmp_out3"}},
{"Out", {"out3"}}}, {"Out", {"out3"}}},
{})); f::AttributeMap{}));
net.CompleteAddOp(); net.CompleteAddOp();
auto backward = f::Backward(net, {"mul_out2", "tmp_out2", "out2"}); auto backward = f::Backward(net, {"mul_out2", "tmp_out2", "out2"});
...@@ -509,7 +516,8 @@ TEST(Backward, simple_single_op) { ...@@ -509,7 +516,8 @@ TEST(Backward, simple_single_op) {
auto target = f::VarDescBind("out"); auto target = f::VarDescBind("out");
target.SetShape({1}); target.SetShape({1});
auto var_to_grad = AppendBackward(program, target, {}); auto var_to_grad =
AppendBackward(program, target, std::unordered_set<std::string>{});
ASSERT_EQ(block->AllOps().size(), 3UL); ASSERT_EQ(block->AllOps().size(), 3UL);
f::OpDescBind *fill_op = block->AllOps()[1]; f::OpDescBind *fill_op = block->AllOps()[1];
...@@ -546,7 +554,7 @@ TEST(Backward, default_attribute) { ...@@ -546,7 +554,7 @@ TEST(Backward, default_attribute) {
auto target = f::VarDescBind("out"); auto target = f::VarDescBind("out");
target.SetShape({1}); target.SetShape({1});
AppendBackward(program, target, {}); AppendBackward(program, target, std::unordered_set<std::string>{});
ASSERT_EQ(block->AllOps().size(), 3UL); ASSERT_EQ(block->AllOps().size(), 3UL);
EXPECT_EQ(boost::get<int>(op->GetAttr("x_num_col_dims")), 1); EXPECT_EQ(boost::get<int>(op->GetAttr("x_num_col_dims")), 1);
...@@ -585,7 +593,8 @@ TEST(Backward, simple_mult_op) { ...@@ -585,7 +593,8 @@ TEST(Backward, simple_mult_op) {
auto target = f::VarDescBind("out3"); auto target = f::VarDescBind("out3");
target.SetShape({1}); target.SetShape({1});
size_t forward_len = block->AllOps().size(); size_t forward_len = block->AllOps().size();
auto var_to_grad = AppendBackward(program, target, {}); auto var_to_grad =
AppendBackward(program, target, std::unordered_set<std::string>{});
ASSERT_EQ(block->AllOps().size(), 6UL + 1); ASSERT_EQ(block->AllOps().size(), 6UL + 1);
f::OpDescBind *fill_op = block->AllOps()[forward_len]; f::OpDescBind *fill_op = block->AllOps()[forward_len];
...@@ -817,7 +826,8 @@ TEST(Backward, shared_var) { ...@@ -817,7 +826,8 @@ TEST(Backward, shared_var) {
auto target = f::VarDescBind("out3"); auto target = f::VarDescBind("out3");
target.SetShape({1}); target.SetShape({1});
size_t forward_len = block->AllOps().size(); size_t forward_len = block->AllOps().size();
auto var_to_grad = AppendBackward(program, target, {}); auto var_to_grad =
AppendBackward(program, target, std::unordered_set<std::string>{});
ASSERT_EQ(block->AllOps().size(), 8UL); ASSERT_EQ(block->AllOps().size(), 8UL);
f::OpDescBind *fill_op = block->AllOps()[forward_len]; f::OpDescBind *fill_op = block->AllOps()[forward_len];
......
paddle/framework/op_desc.cc
浏览文件 @ aa2f2fb4
...@@ -316,8 +316,8 @@ static void InitInferShapeFuncs() { ...@@ -316,8 +316,8 @@ static void InitInferShapeFuncs() {
for (auto &kern_pair : OperatorWithKernel::AllOpKernels()) { for (auto &kern_pair : OperatorWithKernel::AllOpKernels()) {
auto op_type = kern_pair.first; auto op_type = kern_pair.first;
auto &op_info = info_map.at(op_type); auto &op_info = info_map.at(op_type);
auto op = auto op = static_cast<OperatorWithKernel *>(op_info.Creator()(
static_cast<OperatorWithKernel *>(op_info.Creator()("", {}, {}, {})); "", VariableNameMap{}, VariableNameMap{}, AttributeMap{}));
if (op_info.infer_shape_) { // infer_shape has been registered. if (op_info.infer_shape_) { // infer_shape has been registered.
continue; continue;
} }
......
paddle/framework/op_registry.h
浏览文件 @ aa2f2fb4
...@@ -181,8 +181,8 @@ class OpKernelRegistrar : public Registrar { ...@@ -181,8 +181,8 @@ class OpKernelRegistrar : public Registrar {
return 0; \ return 0; \
} }
#define REGISTER_OP_GPU_KERNEL(op_type, ...) \ #define REGISTER_OP_CUDA_KERNEL(op_type, ...) \
REGISTER_OP_KERNEL(op_type, GPU, ::paddle::platform::GPUPlace, __VA_ARGS__) REGISTER_OP_KERNEL(op_type, CUDA, ::paddle::platform::GPUPlace, __VA_ARGS__)
#define REGISTER_OP_CPU_KERNEL(op_type, ...) \ #define REGISTER_OP_CPU_KERNEL(op_type, ...) \
REGISTER_OP_KERNEL(op_type, CPU, ::paddle::platform::CPUPlace, __VA_ARGS__) REGISTER_OP_KERNEL(op_type, CPU, ::paddle::platform::CPUPlace, __VA_ARGS__)
...@@ -217,7 +217,7 @@ class OpKernelRegistrar : public Registrar { ...@@ -217,7 +217,7 @@ class OpKernelRegistrar : public Registrar {
#else #else
#define USE_OP_KERNEL(op_type) \ #define USE_OP_KERNEL(op_type) \
USE_OP_DEVICE_KERNEL(op_type, CPU); \ USE_OP_DEVICE_KERNEL(op_type, CPU); \
USE_OP_DEVICE_KERNEL(op_type, GPU) USE_OP_DEVICE_KERNEL(op_type, CUDA)
#endif #endif
#define USE_NO_KERNEL_OP(op_type) USE_OP_ITSELF(op_type); #define USE_NO_KERNEL_OP(op_type) USE_OP_ITSELF(op_type);
...@@ -226,9 +226,9 @@ class OpKernelRegistrar : public Registrar { ...@@ -226,9 +226,9 @@ class OpKernelRegistrar : public Registrar {
USE_OP_ITSELF(op_type); \ USE_OP_ITSELF(op_type); \
USE_OP_DEVICE_KERNEL(op_type, CPU); USE_OP_DEVICE_KERNEL(op_type, CPU);
#define USE_GPU_ONLY_OP(op_type) \ #define USE_CUDA_ONLY_OP(op_type) \
USE_OP_ITSELF(op_type); \ USE_OP_ITSELF(op_type); \
USE_OP_DEVICE_KERNEL(op_type, GPU) USE_OP_DEVICE_KERNEL(op_type, CUDA)
#define USE_OP(op_type) \ #define USE_OP(op_type) \
USE_OP_ITSELF(op_type); \ USE_OP_ITSELF(op_type); \
......
paddle/framework/operator.cc
浏览文件 @ aa2f2fb4
...@@ -22,20 +22,6 @@ limitations under the License. */ ...@@ -22,20 +22,6 @@ limitations under the License. */
namespace paddle { namespace paddle {
namespace framework { namespace framework {
template <>
Eigen::DefaultDevice& ExecutionContext::GetEigenDevice<
platform::CPUPlace, Eigen::DefaultDevice>() const {
return *device_context_.GetEigenDevice<platform::CPUPlace>();
}
#ifdef PADDLE_WITH_CUDA
template <>
Eigen::GpuDevice&
ExecutionContext::GetEigenDevice<platform::GPUPlace, Eigen::GpuDevice>() const {
return *device_context_.GetEigenDevice<platform::GPUPlace>();
}
#endif
std::string OperatorBase::Input(const std::string& name) const { std::string OperatorBase::Input(const std::string& name) const {
auto& ins = Inputs(name); auto& ins = Inputs(name);
PADDLE_ENFORCE_LE(ins.size(), 1UL, PADDLE_ENFORCE_LE(ins.size(), 1UL,
...@@ -429,7 +415,7 @@ void OperatorWithKernel::Run(const Scope& scope, ...@@ -429,7 +415,7 @@ void OperatorWithKernel::Run(const Scope& scope,
} }
OpKernelType OperatorWithKernel::GetKernelType( OpKernelType OperatorWithKernel::GetKernelType(
const ExecutionContext& ctx) const { const ExecutionContext& ctx) const {
return OpKernelType(IndicateDataType(ctx), ctx.device_context()); return OpKernelType(IndicateDataType(ctx), ctx.GetPlace());
} }
DataType OperatorWithKernel::IndicateDataType( DataType OperatorWithKernel::IndicateDataType(
const ExecutionContext& ctx) const { const ExecutionContext& ctx) const {
......
paddle/framework/operator.h
浏览文件 @ aa2f2fb4
...@@ -276,17 +276,25 @@ class ExecutionContext { ...@@ -276,17 +276,25 @@ class ExecutionContext {
out_tensor->set_lod(in_tensor.lod()); out_tensor->set_lod(in_tensor.lod());
} }
template <typename PlaceType,
typename DeviceType = typename platform::EigenDeviceConverter<
PlaceType>::EigenDeviceType>
DeviceType& GetEigenDevice() const;
platform::Place GetPlace() const { return device_context_.GetPlace(); } platform::Place GetPlace() const { return device_context_.GetPlace(); }
template <typename DeviceContextType>
const DeviceContextType& device_context() const {
return *reinterpret_cast<const DeviceContextType*>(&device_context_);
}
const platform::DeviceContext& device_context() const { const platform::DeviceContext& device_context() const {
return device_context_; return device_context_;
} }
#ifdef PADDLE_WITH_CUDA
const inline platform::CUDADeviceContext& cuda_device_context() const {
PADDLE_ENFORCE(platform::is_gpu_place(device_context_.GetPlace()));
return *reinterpret_cast<const platform::CUDADeviceContext*>(
&device_context_);
}
#endif
//! Get actual name vector for this input. //! Get actual name vector for this input.
const std::vector<std::string>& Inputs(const std::string& name) const { const std::vector<std::string>& Inputs(const std::string& name) const {
return op_.Inputs(name); return op_.Inputs(name);
...@@ -297,14 +305,6 @@ class ExecutionContext { ...@@ -297,14 +305,6 @@ class ExecutionContext {
return op_.Outputs(name); return op_.Outputs(name);
} }
#ifdef PADDLE_WITH_CUDA
const inline platform::CUDADeviceContext& cuda_device_context() const {
PADDLE_ENFORCE(platform::is_gpu_place(device_context_.GetPlace()));
return *reinterpret_cast<const platform::CUDADeviceContext*>(
&device_context_);
}
#endif
private: private:
const OperatorBase& op_; const OperatorBase& op_;
const Scope& scope_; const Scope& scope_;
......
paddle/framework/operator_test.cc
浏览文件 @ aa2f2fb4
...@@ -115,7 +115,7 @@ class OpWithKernelTest : public OperatorWithKernel { ...@@ -115,7 +115,7 @@ class OpWithKernelTest : public OperatorWithKernel {
protected: protected:
void InferShape(framework::InferShapeContext* ctx) const override {} void InferShape(framework::InferShapeContext* ctx) const override {}
OpKernelType GetKernelType(const ExecutionContext& ctx) const override { OpKernelType GetKernelType(const ExecutionContext& ctx) const override {
return OpKernelType(DataType::FP32, ctx.device_context()); return OpKernelType(DataType::FP32, ctx.GetPlace());
} }
}; };
...@@ -261,7 +261,9 @@ class OperatorClone : public paddle::framework::OperatorBase { ...@@ -261,7 +261,9 @@ class OperatorClone : public paddle::framework::OperatorBase {
}; };
TEST(Operator, Clone) { TEST(Operator, Clone) {
OperatorClone a("ABC", {}, {}, {}); OperatorClone a("ABC", paddle::framework::VariableNameMap{},
paddle::framework::VariableNameMap{},
paddle::framework::AttributeMap{});
auto b = a.Clone(); auto b = a.Clone();
ASSERT_EQ(a.Type(), b->Type()); ASSERT_EQ(a.Type(), b->Type());
} }
paddle/framework/prune_test.cc
浏览文件 @ aa2f2fb4
...@@ -54,7 +54,8 @@ TEST(Prune, one_operator) { ...@@ -54,7 +54,8 @@ TEST(Prune, one_operator) {
f::ProgramDescBind program; f::ProgramDescBind program;
f::BlockDescBind *block = program.MutableBlock(0); f::BlockDescBind *block = program.MutableBlock(0);
AddOp("one_one", {{"input", {"a"}}}, {{"output", {"b"}}}, {}, block); AddOp("one_one", {{"input", {"a"}}}, {{"output", {"b"}}}, f::AttributeMap{},
block);
f::ProgramDesc *pdesc = program.Proto(); f::ProgramDesc *pdesc = program.Proto();
f::ProgramDesc pruned; f::ProgramDesc pruned;
...@@ -71,10 +72,14 @@ TEST(Prune, forward) { ...@@ -71,10 +72,14 @@ TEST(Prune, forward) {
f::ProgramDescBind program; f::ProgramDescBind program;
f::BlockDescBind *block = program.MutableBlock(0); f::BlockDescBind *block = program.MutableBlock(0);
AddOp("one_one", {{"input", {"a"}}}, {{"output", {"b"}}}, {}, block); AddOp("one_one", {{"input", {"a"}}}, {{"output", {"b"}}}, f::AttributeMap{},
AddOp("one_one", {{"input", {"b"}}}, {{"output", {"c"}}}, {}, block); block);
AddOp("one_one", {{"input", {"c"}}}, {{"output", {"d"}}}, {}, block); AddOp("one_one", {{"input", {"b"}}}, {{"output", {"c"}}}, f::AttributeMap{},
AddOp("one_one", {{"input", {"d"}}}, {{"output", {"e"}}}, {}, block); block);
AddOp("one_one", {{"input", {"c"}}}, {{"output", {"d"}}}, f::AttributeMap{},
block);
AddOp("one_one", {{"input", {"d"}}}, {{"output", {"e"}}}, f::AttributeMap{},
block);
f::ProgramDesc *pdesc = program.Proto(); f::ProgramDesc *pdesc = program.Proto();
...@@ -90,11 +95,14 @@ TEST(Prune, multi_input_op) { ...@@ -90,11 +95,14 @@ TEST(Prune, multi_input_op) {
f::ProgramDescBind program; f::ProgramDescBind program;
f::BlockDescBind *block = program.MutableBlock(0); f::BlockDescBind *block = program.MutableBlock(0);
AddOp("one_one", {{"input", {"a0"}}}, {{"output", {"b0"}}}, {}, block); AddOp("one_one", {{"input", {"a0"}}}, {{"output", {"b0"}}}, f::AttributeMap{},
AddOp("one_one", {{"input", {"a1"}}}, {{"output", {"b1"}}}, {}, block); block);
AddOp("one_one", {{"input", {"a2"}}}, {{"output", {"b2"}}}, {}, block); AddOp("one_one", {{"input", {"a1"}}}, {{"output", {"b1"}}}, f::AttributeMap{},
AddOp("three_one", {{"input", {"b0", "b1", "b2"}}}, {{"output", {"c"}}}, {},
block); block);
AddOp("one_one", {{"input", {"a2"}}}, {{"output", {"b2"}}}, f::AttributeMap{},
block);
AddOp("three_one", {{"input", {"b0", "b1", "b2"}}}, {{"output", {"c"}}},
f::AttributeMap{}, block);
f::ProgramDesc *pdesc = program.Proto(); f::ProgramDesc *pdesc = program.Proto();
pdesc->mutable_blocks(0)->mutable_ops(3)->set_is_target(true); pdesc->mutable_blocks(0)->mutable_ops(3)->set_is_target(true);
...@@ -108,9 +116,12 @@ TEST(Prune, multi_output_op) { ...@@ -108,9 +116,12 @@ TEST(Prune, multi_output_op) {
f::ProgramDescBind program; f::ProgramDescBind program;
f::BlockDescBind *block = program.MutableBlock(0); f::BlockDescBind *block = program.MutableBlock(0);
AddOp("one_two", {{"input", {"a"}}}, {{"output", {"b", "c"}}}, {}, block); AddOp("one_two", {{"input", {"a"}}}, {{"output", {"b", "c"}}},
AddOp("one_one", {{"input", {"b"}}}, {{"output", {"b1"}}}, {}, block); f::AttributeMap{}, block);
AddOp("one_one", {{"input", {"c"}}}, {{"output", {"c1"}}}, {}, block); AddOp("one_one", {{"input", {"b"}}}, {{"output", {"b1"}}}, f::AttributeMap{},
block);
AddOp("one_one", {{"input", {"c"}}}, {{"output", {"c1"}}}, f::AttributeMap{},
block);
f::ProgramDesc *pdesc = program.Proto(); f::ProgramDesc *pdesc = program.Proto();
pdesc->mutable_blocks(0)->mutable_ops(2)->set_is_target(true); pdesc->mutable_blocks(0)->mutable_ops(2)->set_is_target(true);
...@@ -124,9 +135,12 @@ TEST(Prune, multi_target) { ...@@ -124,9 +135,12 @@ TEST(Prune, multi_target) {
f::ProgramDescBind program; f::ProgramDescBind program;
f::BlockDescBind *block = program.MutableBlock(0); f::BlockDescBind *block = program.MutableBlock(0);
AddOp("one_two", {{"input", {"a"}}}, {{"output", {"b", "c"}}}, {}, block); AddOp("one_two", {{"input", {"a"}}}, {{"output", {"b", "c"}}},
AddOp("one_one", {{"input", {"b"}}}, {{"output", {"b1"}}}, {}, block); f::AttributeMap{}, block);
AddOp("one_one", {{"input", {"c"}}}, {{"output", {"c1"}}}, {}, block); AddOp("one_one", {{"input", {"b"}}}, {{"output", {"b1"}}}, f::AttributeMap{},
block);
AddOp("one_one", {{"input", {"c"}}}, {{"output", {"c1"}}}, f::AttributeMap{},
block);
f::ProgramDesc *pdesc = program.Proto(); f::ProgramDesc *pdesc = program.Proto();
pdesc->mutable_blocks(0)->mutable_ops(1)->set_is_target(true); pdesc->mutable_blocks(0)->mutable_ops(1)->set_is_target(true);
......
paddle/gserver/activations/ActivationFunction.cpp
浏览文件 @ aa2f2fb4
...@@ -24,7 +24,7 @@ limitations under the License. */ ...@@ -24,7 +24,7 @@ limitations under the License. */
#include "paddle/utils/ClassRegistrar.h" #include "paddle/utils/ClassRegistrar.h"
#include "paddle/utils/Logging.h" #include "paddle/utils/Logging.h"
#ifdef PADDLE_USE_MKLDNN #ifdef PADDLE_WITH_MKLDNN
#include "MKLDNNActivation.h" #include "MKLDNNActivation.h"
#endif #endif
...@@ -490,7 +490,7 @@ Error __must_check backward(Argument& act) { ...@@ -490,7 +490,7 @@ Error __must_check backward(Argument& act) {
END_DEFINE_ACTIVATION(log) END_DEFINE_ACTIVATION(log)
ActivationFunction* ActivationFunction::create(const std::string& type) { ActivationFunction* ActivationFunction::create(const std::string& type) {
#ifdef PADDLE_USE_MKLDNN #ifdef PADDLE_WITH_MKLDNN
if (!type.empty() && type.compare(0, 7, "mkldnn_") == 0) { if (!type.empty() && type.compare(0, 7, "mkldnn_") == 0) {
return MKLDNNActivation::create(type); return MKLDNNActivation::create(type);
} }
......
paddle/gserver/gradientmachines/NeuralNetwork.cpp
浏览文件 @ aa2f2fb4
...@@ -20,7 +20,7 @@ limitations under the License. */ ...@@ -20,7 +20,7 @@ limitations under the License. */
#include "paddle/utils/Logging.h" #include "paddle/utils/Logging.h"
#include "paddle/utils/Stat.h" #include "paddle/utils/Stat.h"
#ifdef PADDLE_USE_MKLDNN #ifdef PADDLE_WITH_MKLDNN
#include "paddle/gserver/layers/MKLDNNLayer.h" #include "paddle/gserver/layers/MKLDNNLayer.h"
#endif #endif
...@@ -307,7 +307,7 @@ void NeuralNetwork::backward(const UpdateCallback& callback) { ...@@ -307,7 +307,7 @@ void NeuralNetwork::backward(const UpdateCallback& callback) {
} }
void NeuralNetwork::finish() { void NeuralNetwork::finish() {
#ifdef PADDLE_USE_MKLDNN #ifdef PADDLE_WITH_MKLDNN
FOR_EACH_R(layer, layers_) { FOR_EACH_R(layer, layers_) {
MKLDNNLayerPtr dnnLayer = std::dynamic_pointer_cast<MKLDNNLayer>(*layer); MKLDNNLayerPtr dnnLayer = std::dynamic_pointer_cast<MKLDNNLayer>(*layer);
if (dnnLayer) { if (dnnLayer) {
......
paddle/math/Allocator.h
浏览文件 @ aa2f2fb4
...@@ -48,7 +48,7 @@ public: ...@@ -48,7 +48,7 @@ public:
*/ */
virtual void* alloc(size_t size) { virtual void* alloc(size_t size) {
void* ptr; void* ptr;
#ifdef PADDLE_USE_MKLDNN #ifdef PADDLE_WITH_MKLDNN
// refer to https://github.com/01org/mkl-dnn/blob/master/include/mkldnn.hpp // refer to https://github.com/01org/mkl-dnn/blob/master/include/mkldnn.hpp
// memory alignment // memory alignment
CHECK_EQ(posix_memalign(&ptr, 4096ul, size), 0); CHECK_EQ(posix_memalign(&ptr, 4096ul, size), 0);
......
paddle/math/MathFunctions.cpp
浏览文件 @ aa2f2fb4
...@@ -206,7 +206,7 @@ double dotProduct<double>(const int n, const double* x, const double* y) { ...@@ -206,7 +206,7 @@ double dotProduct<double>(const int n, const double* x, const double* y) {
} }
#endif #endif
#if defined(PADDLE_USE_MKLML) #if defined(PADDLE_WITH_MKLML)
template <> template <>
void vExp<float>(const int n, const float* a, float* r) { void vExp<float>(const int n, const float* a, float* r) {
......
paddle/math/MathFunctions.h
浏览文件 @ aa2f2fb4
...@@ -15,7 +15,7 @@ limitations under the License. */ ...@@ -15,7 +15,7 @@ limitations under the License. */
#ifndef MATHFUNCTIONS_H_ #ifndef MATHFUNCTIONS_H_
#define MATHFUNCTIONS_H_ #define MATHFUNCTIONS_H_
#ifdef PADDLE_USE_MKLML #ifdef PADDLE_WITH_MKLML
#include <mkl_cblas.h> #include <mkl_cblas.h>
#include <mkl_lapacke.h> #include <mkl_lapacke.h>
#include <mkl_vml_functions.h> #include <mkl_vml_functions.h>
......
paddle/math/Matrix.cpp
浏览文件 @ aa2f2fb4
...@@ -28,6 +28,7 @@ limitations under the License. */ ...@@ -28,6 +28,7 @@ limitations under the License. */
#include "hl_top_k.h" #include "hl_top_k.h"
#include "paddle/utils/Logging.h" #include "paddle/utils/Logging.h"
#include "NEONFunctions.h"
#include "paddle/function/GemmFunctor.h" #include "paddle/function/GemmFunctor.h"
#include "paddle/utils/ThreadLocal.h" #include "paddle/utils/ThreadLocal.h"
...@@ -4165,16 +4166,36 @@ void CpuMatrix::print(std::ostream& os) const { ...@@ -4165,16 +4166,36 @@ void CpuMatrix::print(std::ostream& os) const {
void CpuMatrix::paramReluForward(Matrix& data, Matrix& W) { void CpuMatrix::paramReluForward(Matrix& data, Matrix& W) {
real* input = data.getData(); real* input = data.getData();
real* w = W.getData(); real* w = W.getData();
real* output = data_;
size_t numElements = data.getWidth(); size_t numElements = data.getWidth();
size_t numSamples = data.getHeight(); size_t numSamples = data.getHeight();
size_t paraSize = W.getHeight() * W.getWidth(); size_t paraSize = W.getHeight() * W.getWidth();
CHECK(!(numElements % paraSize)); // this check from ParameterReluLayer::init CHECK(!(numElements % paraSize)); // this check from ParameterReluLayer::init
size_t partial_sum = numElements / paraSize; size_t partial_sum = numElements / paraSize;
if (paraSize == numElements) {
for (size_t n = 0; n < numSamples * numElements; ++n) {
output[n] = input[n] > 0 ? input[n] : input[n] * w[n % numElements];
}
return;
}
#if defined(__ARM_NEON__) || defined(__ARM_NEON)
for (size_t n = 0; n < numSamples; ++n) {
for (size_t i = 0; i < paraSize; i++) {
neon::prelu(
input + i * partial_sum, w[i], output + i * partial_sum, partial_sum);
}
input = input + numElements;
output = output + numElements;
}
#else
for (size_t n = 0, k = 0; n < numSamples; ++n) { for (size_t n = 0, k = 0; n < numSamples; ++n) {
for (size_t i = 0; i < numElements; ++i, ++k) { for (size_t i = 0; i < numElements; ++i, ++k) {
data_[k] = input[k] > 0 ? input[k] : input[k] * w[i / partial_sum]; output[k] = input[k] > 0 ? input[k] : input[k] * w[i / partial_sum];
} }
} }
#endif
} }
void CpuMatrix::paramReluBackwardW(Matrix& oGrad, Matrix& data) { void CpuMatrix::paramReluBackwardW(Matrix& oGrad, Matrix& data) {
......
paddle/math/NEONFunctions.cpp
浏览文件 @ aa2f2fb4
...@@ -49,6 +49,46 @@ void relu(const float* a, float* b, int len) { ...@@ -49,6 +49,46 @@ void relu(const float* a, float* b, int len) {
} }
} }
// b[i] = a[i] > 0.0f ? a[i] : a[i] * w
void prelu(const float* a, float w, float* b, int len) {
int offset = len % 16;
float32x4_t ma0, ma1, ma2, ma3;
float32x4_t zero = vdupq_n_f32(0.f);
float32x4_t vw = vdupq_n_f32(w);
for (int k = 0; k < len / 16; k++, a += 16, b += 16) {
ma0 = vld1q_f32(a);
ma1 = vld1q_f32(a + 4);
ma2 = vld1q_f32(a + 8);
ma3 = vld1q_f32(a + 12);
uint32x4_t flag0 = vcgtq_f32(ma0, zero);
uint32x4_t flag1 = vcgtq_f32(ma1, zero);
uint32x4_t flag2 = vcgtq_f32(ma2, zero);
uint32x4_t flag3 = vcgtq_f32(ma3, zero);
float32x4_t mul0 = vmulq_f32(ma0, vw);
float32x4_t mul1 = vmulq_f32(ma1, vw);
float32x4_t mul2 = vmulq_f32(ma2, vw);
float32x4_t mul3 = vmulq_f32(ma3, vw);
ma0 = vbslq_f32(flag0, ma0, mul0);
ma1 = vbslq_f32(flag1, ma1, mul1);
ma2 = vbslq_f32(flag2, ma2, mul2);
ma3 = vbslq_f32(flag3, ma3, mul3);
vst1q_f32(b, ma0);
vst1q_f32(b + 4, ma1);
vst1q_f32(b + 8, ma2);
vst1q_f32(b + 12, ma3);
}
for (int i = 0; i < offset; i++) {
b[i] = a[i] > 0.0f ? a[i] : a[i] * w;
}
}
} // namespace neon } // namespace neon
} // namespace paddle } // namespace paddle
......
paddle/math/NEONFunctions.h
浏览文件 @ aa2f2fb4
...@@ -18,6 +18,7 @@ namespace paddle { ...@@ -18,6 +18,7 @@ namespace paddle {
namespace neon { namespace neon {
void relu(const float* a, float* b, int len); void relu(const float* a, float* b, int len);
void prelu(const float* a, float w, float* b, int len);
} // namespace neon } // namespace neon
} // namespace paddle } // namespace paddle
paddle/math/float16.h
浏览文件 @ aa2f2fb4
...@@ -101,7 +101,7 @@ public: ...@@ -101,7 +101,7 @@ public:
half tmp = __float2half(val); half tmp = __float2half(val);
x = *reinterpret_cast<uint16_t*>(&tmp); x = *reinterpret_cast<uint16_t*>(&tmp);
#elif defined(PADDLE_NEON) #elif defined(PADDLE_WITH_NATIVE_FP16)
float32x4_t tmp = vld1q_dup_f32(&val); float32x4_t tmp = vld1q_dup_f32(&val);
float16_t res = vget_lane_f16(vcvt_f16_f32(tmp), 0); float16_t res = vget_lane_f16(vcvt_f16_f32(tmp), 0);
x = *reinterpret_cast<uint16_t*>(&res); x = *reinterpret_cast<uint16_t*>(&res);
...@@ -252,7 +252,7 @@ public: ...@@ -252,7 +252,7 @@ public:
half tmp = *reinterpret_cast<const half*>(this); half tmp = *reinterpret_cast<const half*>(this);
return __half2float(tmp); return __half2float(tmp);
#elif defined(PADDLE_NEON) #elif defined(PADDLE_WITH_NATIVE_FP16)
float16x4_t res = vld1_dup_f16(reinterpret_cast<const float16_t*>(this)); float16x4_t res = vld1_dup_f16(reinterpret_cast<const float16_t*>(this));
return vgetq_lane_f32(vcvt_f32_f16(res), 0); return vgetq_lane_f32(vcvt_f32_f16(res), 0);
......
paddle/math/tests/CMakeLists.txt
浏览文件 @ aa2f2fb4
...@@ -34,4 +34,4 @@ add_simple_unittest(test_FPException) ...@@ -34,4 +34,4 @@ add_simple_unittest(test_FPException)
add_simple_unittest(test_GpuProfiler) add_simple_unittest(test_GpuProfiler)
add_simple_unittest(test_BaseMatrix) add_simple_unittest(test_BaseMatrix)
add_simple_unittest(test_Matrix) add_simple_unittest(test_Matrix)
cc_test(test_float16 SRCS test_float16.cpp) add_simple_unittest(test_float16)
paddle/memory/detail/system_allocator.cc
浏览文件 @ aa2f2fb4
...@@ -43,7 +43,7 @@ void* CPUAllocator::Alloc(size_t& index, size_t size) { ...@@ -43,7 +43,7 @@ void* CPUAllocator::Alloc(size_t& index, size_t size) {
void* p; void* p;
#ifdef PADDLE_USE_MKLDNN #ifdef PADDLE_WITH_MKLDNN
// refer to https://github.com/01org/mkl-dnn/blob/master/include/mkldnn.hpp // refer to https://github.com/01org/mkl-dnn/blob/master/include/mkldnn.hpp
// memory alignment // memory alignment
PADDLE_ENFORCE_EQ(posix_memalign(&p, 4096ul, size), 0); PADDLE_ENFORCE_EQ(posix_memalign(&p, 4096ul, size), 0);
......
paddle/operators/CMakeLists.txt
浏览文件 @ aa2f2fb4
...@@ -138,7 +138,7 @@ function(op_library TARGET) ...@@ -138,7 +138,7 @@ function(op_library TARGET)
if ("${TARGET}" STREQUAL "nccl_op") if ("${TARGET}" STREQUAL "nccl_op")
set(pybind_flag 1) set(pybind_flag 1)
# It's enough to just adding one operator to pybind # It's enough to just adding one operator to pybind
file(APPEND ${pybind_file} "USE_GPU_ONLY_OP(ncclAllReduce);\n") file(APPEND ${pybind_file} "USE_CUDA_ONLY_OP(ncclAllReduce);\n")
endif() endif()
# reduce_op contains several operators # reduce_op contains several operators
......
paddle/operators/accuracy_op.cc
浏览文件 @ aa2f2fb4
...@@ -57,7 +57,7 @@ class AccuracyOp : public framework::OperatorWithKernel { ...@@ -57,7 +57,7 @@ class AccuracyOp : public framework::OperatorWithKernel {
const framework::ExecutionContext &ctx) const override { const framework::ExecutionContext &ctx) const override {
return framework::OpKernelType( return framework::OpKernelType(
framework::ToDataType(ctx.Input<Tensor>("Out")->type()), framework::ToDataType(ctx.Input<Tensor>("Out")->type()),
ctx.device_context()); ctx.GetPlace());
} }
}; };
......
paddle/operators/accuracy_op.cu
浏览文件 @ aa2f2fb4
...@@ -104,5 +104,6 @@ class AccuracyOpCUDAKernel : public framework::OpKernel<T> { ...@@ -104,5 +104,6 @@ class AccuracyOpCUDAKernel : public framework::OpKernel<T> {
// FIXME(typhoonzero): types of T is for inference data. // FIXME(typhoonzero): types of T is for inference data.
// label data is always int64 // label data is always int64
REGISTER_OP_GPU_KERNEL(accuracy, paddle::operators::AccuracyOpCUDAKernel<float>, REGISTER_OP_CUDA_KERNEL(accuracy,
paddle::operators::AccuracyOpCUDAKernel<double>); paddle::operators::AccuracyOpCUDAKernel<float>,
paddle::operators::AccuracyOpCUDAKernel<double>);
paddle/operators/accuracy_op.h
浏览文件 @ aa2f2fb4
...@@ -21,7 +21,7 @@ namespace operators { ...@@ -21,7 +21,7 @@ namespace operators {
using Tensor = framework::Tensor; using Tensor = framework::Tensor;
template <typename Place, typename T> template <typename DeviceContext, typename T>
class AccuracyKernel : public framework::OpKernel<T> { class AccuracyKernel : public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext& ctx) const override { void Compute(const framework::ExecutionContext& ctx) const override {
......
paddle/operators/activation_op.cc
浏览文件 @ aa2f2fb4
...@@ -611,16 +611,17 @@ REGISTER_OP(hard_sigmoid, ops::ActivationOp, ops::HardSigmoidOpMaker, ...@@ -611,16 +611,17 @@ REGISTER_OP(hard_sigmoid, ops::ActivationOp, ops::HardSigmoidOpMaker,
REGISTER_OP(swish, ops::ActivationOp, ops::SwishOpMaker, swish_grad, REGISTER_OP(swish, ops::ActivationOp, ops::SwishOpMaker, swish_grad,
ops::ActivationOpGrad); ops::ActivationOpGrad);
#define REGISTER_ACTIVATION_CPU_KERNEL(act_type, functor, grad_functor) \ #define REGISTER_ACTIVATION_CPU_KERNEL(act_type, functor, grad_functor) \
REGISTER_OP_CPU_KERNEL( \ REGISTER_OP_CPU_KERNEL( \
act_type, \ act_type, ops::ActivationKernel<paddle::platform::CPUDeviceContext, \
ops::ActivationKernel<paddle::platform::CPUPlace, ops::functor<float>>, \ ops::functor<float>>, \
ops::ActivationKernel<paddle::platform::CPUPlace, \ ops::ActivationKernel<paddle::platform::CPUDeviceContext, \
ops::functor<double>>); \ ops::functor<double>>); \
REGISTER_OP_CPU_KERNEL( \ REGISTER_OP_CPU_KERNEL( \
act_type##_grad, ops::ActivationGradKernel<paddle::platform::CPUPlace, \ act_type##_grad, \
ops::grad_functor<float>>, \ ops::ActivationGradKernel<paddle::platform::CPUDeviceContext, \
ops::ActivationGradKernel<paddle::platform::CPUPlace, \ ops::grad_functor<float>>, \
ops::ActivationGradKernel<paddle::platform::CPUDeviceContext, \
ops::grad_functor<double>>); ops::grad_functor<double>>);
FOR_EACH_KERNEL_FUNCTOR(REGISTER_ACTIVATION_CPU_KERNEL); FOR_EACH_KERNEL_FUNCTOR(REGISTER_ACTIVATION_CPU_KERNEL);
paddle/operators/activation_op.cu
浏览文件 @ aa2f2fb4
...@@ -17,16 +17,17 @@ ...@@ -17,16 +17,17 @@
namespace ops = paddle::operators; namespace ops = paddle::operators;
#define REGISTER_ACTIVATION_GPU_KERNEL(act_type, functor, grad_functor) \ #define REGISTER_ACTIVATION_CUDA_KERNEL(act_type, functor, grad_functor) \
REGISTER_OP_GPU_KERNEL( \ REGISTER_OP_CUDA_KERNEL( \
act_type, \ act_type, ops::ActivationKernel<paddle::platform::CUDADeviceContext, \
ops::ActivationKernel<paddle::platform::GPUPlace, ops::functor<float>>, \ ops::functor<float>>, \
ops::ActivationKernel<paddle::platform::GPUPlace, \ ops::ActivationKernel<paddle::platform::CUDADeviceContext, \
ops::functor<double>>); \ ops::functor<double>>); \
REGISTER_OP_GPU_KERNEL( \ REGISTER_OP_CUDA_KERNEL( \
act_type##_grad, ops::ActivationGradKernel<paddle::platform::GPUPlace, \ act_type##_grad, \
ops::grad_functor<float>>, \ ops::ActivationGradKernel<paddle::platform::CUDADeviceContext, \
ops::ActivationGradKernel<paddle::platform::GPUPlace, \ ops::grad_functor<float>>, \
ops::ActivationGradKernel<paddle::platform::CUDADeviceContext, \
ops::grad_functor<double>>); ops::grad_functor<double>>);
FOR_EACH_KERNEL_FUNCTOR(REGISTER_ACTIVATION_GPU_KERNEL); FOR_EACH_KERNEL_FUNCTOR(REGISTER_ACTIVATION_CUDA_KERNEL);
paddle/operators/activation_op.h
浏览文件 @ aa2f2fb4
...@@ -19,7 +19,7 @@ ...@@ -19,7 +19,7 @@
namespace paddle { namespace paddle {
namespace operators { namespace operators {
template <typename Place, typename Functor> template <typename DeviceContext, typename Functor>
class ActivationKernel class ActivationKernel
: public framework::OpKernel<typename Functor::ELEMENT_TYPE> { : public framework::OpKernel<typename Functor::ELEMENT_TYPE> {
public: public:
...@@ -32,18 +32,19 @@ class ActivationKernel ...@@ -32,18 +32,19 @@ class ActivationKernel
auto x = framework::EigenVector<T>::Flatten(*X); auto x = framework::EigenVector<T>::Flatten(*X);
auto y = framework::EigenVector<T>::Flatten(*Y); auto y = framework::EigenVector<T>::Flatten(*Y);
auto place = context.GetEigenDevice<Place>(); auto* place =
context.template device_context<DeviceContext>().eigen_device();
Functor functor; Functor functor;
auto attrs = functor.GetAttrs(); auto attrs = functor.GetAttrs();
for (auto& attr : attrs) { for (auto& attr : attrs) {
*attr.second = context.Attr<float>(attr.first); *attr.second = context.Attr<float>(attr.first);
} }
functor(place, x, y); functor(*place, x, y);
} }
}; };
template <typename Place, typename Functor> template <typename DeviceContext, typename Functor>
class ActivationGradKernel class ActivationGradKernel
: public framework::OpKernel<typename Functor::ELEMENT_TYPE> { : public framework::OpKernel<typename Functor::ELEMENT_TYPE> {
public: public:
...@@ -59,13 +60,14 @@ class ActivationGradKernel ...@@ -59,13 +60,14 @@ class ActivationGradKernel
auto x = framework::EigenVector<T>::Flatten(*X); auto x = framework::EigenVector<T>::Flatten(*X);
auto y = framework::EigenVector<T>::Flatten(*Y); auto y = framework::EigenVector<T>::Flatten(*Y);
auto dx = framework::EigenVector<T>::Flatten(*dX); auto dx = framework::EigenVector<T>::Flatten(*dX);
auto place = context.GetEigenDevice<Place>(); auto* place =
context.template device_context<DeviceContext>().eigen_device();
Functor functor; Functor functor;
auto attrs = functor.GetAttrs(); auto attrs = functor.GetAttrs();
for (auto& attr : attrs) { for (auto& attr : attrs) {
*attr.second = context.Attr<float>(attr.first); *attr.second = context.Attr<float>(attr.first);
} }
functor(place, x, y, dy, dx); functor(*place, x, y, dy, dx);
} }
}; };
......
paddle/operators/adadelta_op.cc
浏览文件 @ aa2f2fb4
...@@ -92,12 +92,12 @@ for gradient descent. ...@@ -92,12 +92,12 @@ for gradient descent.
Adadelta updates are as follows: Adadelta updates are as follows:
$$avgSquaredGradOut = \rho * avgSquaredGrad + (1 - \rho) * grad * grad \break $$
paramUpdate = - $\sqrt{((avgSquaredUpdate + \epsilon) / avg\_squared\_grad\_out = \rho * avg\_squared\_grad + (1 - \rho) * grad * grad \\
(avgSquaredGrad_out + \epsilon))}$ * grad \break param\_update = - \sqrt{\frac{avg\_squared\_update + \epsilon}{avg\_squared\_grad\_out + \epsilon}} * grad \\
avgSquaredUpdateOut = \rho * avgSquaredUpdate + (1 - \rho) * avg\_squared\_update\_out = \rho * avg\_squared\_update + (1 - \rho) * {param\_update}^2 \\
{(paramUpdate)}^2 \break param\_out = param + param\_update
paramOut = param + paramUpdate$$ $$
)DOC"); )DOC");
} }
...@@ -109,5 +109,5 @@ paramOut = param + paramUpdate$$ ...@@ -109,5 +109,5 @@ paramOut = param + paramUpdate$$
namespace ops = paddle::operators; namespace ops = paddle::operators;
REGISTER_OP_WITHOUT_GRADIENT(adadelta, ops::AdadeltaOp, ops::AdadeltaOpMaker); REGISTER_OP_WITHOUT_GRADIENT(adadelta, ops::AdadeltaOp, ops::AdadeltaOpMaker);
REGISTER_OP_CPU_KERNEL( REGISTER_OP_CPU_KERNEL(
adadelta, ops::AdadeltaOpKernel<paddle::platform::CPUPlace, float>, adadelta, ops::AdadeltaOpKernel<paddle::platform::CPUDeviceContext, float>,
ops::AdadeltaOpKernel<paddle::platform::CPUPlace, double>); ops::AdadeltaOpKernel<paddle::platform::CPUDeviceContext, double>);
paddle/operators/adadelta_op.cu
浏览文件 @ aa2f2fb4
...@@ -16,6 +16,6 @@ ...@@ -16,6 +16,6 @@
#include "paddle/operators/adadelta_op.h" #include "paddle/operators/adadelta_op.h"
namespace ops = paddle::operators; namespace ops = paddle::operators;
REGISTER_OP_GPU_KERNEL( REGISTER_OP_CUDA_KERNEL(
adadelta, ops::AdadeltaOpKernel<paddle::platform::GPUPlace, float>, adadelta, ops::AdadeltaOpKernel<paddle::platform::CUDADeviceContext, float>,
ops::AdadeltaOpKernel<paddle::platform::GPUPlace, double>); ops::AdadeltaOpKernel<paddle::platform::CUDADeviceContext, double>);
paddle/operators/adadelta_op.h
浏览文件 @ aa2f2fb4
...@@ -19,7 +19,7 @@ limitations under the License. */ ...@@ -19,7 +19,7 @@ limitations under the License. */
namespace paddle { namespace paddle {
namespace operators { namespace operators {
template <typename Place, typename T> template <typename DeviceContext, typename T>
class AdadeltaOpKernel : public framework::OpKernel<T> { class AdadeltaOpKernel : public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext& ctx) const override { void Compute(const framework::ExecutionContext& ctx) const override {
...@@ -51,7 +51,7 @@ class AdadeltaOpKernel : public framework::OpKernel<T> { ...@@ -51,7 +51,7 @@ class AdadeltaOpKernel : public framework::OpKernel<T> {
framework::EigenVector<T>::Flatten(*avg_squared_grad_out_tensor); framework::EigenVector<T>::Flatten(*avg_squared_grad_out_tensor);
auto avg_squared_update_out = auto avg_squared_update_out =
framework::EigenVector<T>::Flatten(*avg_squared_update_out_tensor); framework::EigenVector<T>::Flatten(*avg_squared_update_out_tensor);
auto place = ctx.GetEigenDevice<Place>(); auto& place = *ctx.template device_context<DeviceContext>().eigen_device();
avg_squared_grad_out.device(place) = avg_squared_grad_out.device(place) =
rho * avg_squared_grad + (1 - rho) * grad.square(); rho * avg_squared_grad + (1 - rho) * grad.square();
......
paddle/operators/adagrad_op.cc
浏览文件 @ aa2f2fb4
...@@ -80,8 +80,8 @@ Adaptive Gradient Algorithm (Adagrad). ...@@ -80,8 +80,8 @@ Adaptive Gradient Algorithm (Adagrad).
The update is done as follows: The update is done as follows:
$$momentOut = moment + grad * grad \break $$moment\_out = moment + grad * grad \\
paramOut = param - learningRate * grad / ($\sqrt{momentOut}$ + \epsilon) \break param\_out = param - \frac{learning\_rate * grad}{\sqrt{moment\_out} + \epsilon}
$$ $$
The original paper(http://www.jmlr.org/papers/volume12/duchi11a/duchi11a.pdf) The original paper(http://www.jmlr.org/papers/volume12/duchi11a/duchi11a.pdf)
...@@ -100,8 +100,8 @@ size_t FindPos(const std::vector<int64_t>& rows, int64_t value) { ...@@ -100,8 +100,8 @@ size_t FindPos(const std::vector<int64_t>& rows, int64_t value) {
} // namespace } // namespace
template <typename T> template <typename T>
struct SparseAdagradFunctor<platform::CPUPlace, T> { struct SparseAdagradFunctor<platform::CPUDeviceContext, T> {
void operator()(const platform::DeviceContext& context, void operator()(const platform::CPUDeviceContext& context,
const framework::SelectedRows& grad, const framework::SelectedRows& grad,
const framework::Tensor& learning_rate, T epsilon, const framework::Tensor& learning_rate, T epsilon,
framework::Tensor* moment, framework::Tensor* param) { framework::Tensor* moment, framework::Tensor* param) {
...@@ -120,7 +120,7 @@ struct SparseAdagradFunctor<platform::CPUPlace, T> { ...@@ -120,7 +120,7 @@ struct SparseAdagradFunctor<platform::CPUPlace, T> {
{static_cast<int64_t>(merge_rows.size()), grad_width}), {static_cast<int64_t>(merge_rows.size()), grad_width}),
context.GetPlace()); context.GetPlace());
math::SetConstant<platform::CPUPlace, T> constant_functor; math::SetConstant<platform::CPUDeviceContext, T> constant_functor;
constant_functor(context, grad_merge->mutable_value(), 0.0); constant_functor(context, grad_merge->mutable_value(), 0.0);
auto* grad_merge_data = grad_merge->mutable_value()->data<T>(); auto* grad_merge_data = grad_merge->mutable_value()->data<T>();
...@@ -144,9 +144,9 @@ struct SparseAdagradFunctor<platform::CPUPlace, T> { ...@@ -144,9 +144,9 @@ struct SparseAdagradFunctor<platform::CPUPlace, T> {
auto gs = auto gs =
framework::EigenVector<T>::Flatten(*(grad_square->mutable_value())); framework::EigenVector<T>::Flatten(*(grad_square->mutable_value()));
auto gm = framework::EigenVector<T>::Flatten(grad_merge->value()); auto gm = framework::EigenVector<T>::Flatten(grad_merge->value());
gs.device(*context.GetEigenDevice<platform::CPUPlace>()) = gm * gm; gs.device(*context.eigen_device()) = gm * gm;
math::SelectedRowsAddToTensor<platform::CPUPlace, T> functor; math::SelectedRowsAddToTensor<platform::CPUDeviceContext, T> functor;
functor(context, *grad_square, moment); functor(context, *grad_square, moment);
// 3. update parameter // 3. update parameter
...@@ -164,13 +164,13 @@ struct SparseAdagradFunctor<platform::CPUPlace, T> { ...@@ -164,13 +164,13 @@ struct SparseAdagradFunctor<platform::CPUPlace, T> {
} }
}; };
template struct SparseAdagradFunctor<platform::CPUPlace, float>; template struct SparseAdagradFunctor<platform::CPUDeviceContext, float>;
template struct SparseAdagradFunctor<platform::CPUPlace, double>; template struct SparseAdagradFunctor<platform::CPUDeviceContext, double>;
} // namespace operators } // namespace operators
} // namespace paddle } // namespace paddle
namespace ops = paddle::operators; namespace ops = paddle::operators;
REGISTER_OP_WITHOUT_GRADIENT(adagrad, ops::AdagradOp, ops::AdagradOpMaker); REGISTER_OP_WITHOUT_GRADIENT(adagrad, ops::AdagradOp, ops::AdagradOpMaker);
REGISTER_OP_CPU_KERNEL( REGISTER_OP_CPU_KERNEL(
adagrad, ops::AdagradOpKernel<paddle::platform::CPUPlace, float>, adagrad, ops::AdagradOpKernel<paddle::platform::CPUDeviceContext, float>,
ops::AdagradOpKernel<paddle::platform::CPUPlace, double>); ops::AdagradOpKernel<paddle::platform::CPUDeviceContext, double>);
paddle/operators/adagrad_op.cu
浏览文件 @ aa2f2fb4
...@@ -72,8 +72,8 @@ __global__ void SparseAdagradFunctorKernel(const T* grad, const int64_t* rows, ...@@ -72,8 +72,8 @@ __global__ void SparseAdagradFunctorKernel(const T* grad, const int64_t* rows,
} // namespace } // namespace
template <typename T> template <typename T>
struct SparseAdagradFunctor<platform::GPUPlace, T> { struct SparseAdagradFunctor<platform::CUDADeviceContext, T> {
void operator()(const platform::DeviceContext& context, void operator()(const platform::CUDADeviceContext& context,
const framework::SelectedRows& grad, const framework::SelectedRows& grad,
const framework::Tensor& learning_rate, T epsilon, const framework::Tensor& learning_rate, T epsilon,
framework::Tensor* moment, framework::Tensor* param) { framework::Tensor* moment, framework::Tensor* param) {
...@@ -92,7 +92,7 @@ struct SparseAdagradFunctor<platform::GPUPlace, T> { ...@@ -92,7 +92,7 @@ struct SparseAdagradFunctor<platform::GPUPlace, T> {
{static_cast<int64_t>(merge_rows.size()), grad_width}), {static_cast<int64_t>(merge_rows.size()), grad_width}),
context.GetPlace()); context.GetPlace());
math::SetConstant<platform::GPUPlace, T> constant_functor; math::SetConstant<platform::CUDADeviceContext, T> constant_functor;
constant_functor(context, grad_merge->mutable_value(), 0.0); constant_functor(context, grad_merge->mutable_value(), 0.0);
auto* grad_merge_data = grad_merge->mutable_value()->data<T>(); auto* grad_merge_data = grad_merge->mutable_value()->data<T>();
...@@ -119,9 +119,9 @@ struct SparseAdagradFunctor<platform::GPUPlace, T> { ...@@ -119,9 +119,9 @@ struct SparseAdagradFunctor<platform::GPUPlace, T> {
auto gs = auto gs =
framework::EigenVector<T>::Flatten(*(grad_square->mutable_value())); framework::EigenVector<T>::Flatten(*(grad_square->mutable_value()));
auto gm = framework::EigenVector<T>::Flatten(grad_merge->value()); auto gm = framework::EigenVector<T>::Flatten(grad_merge->value());
gs.device(*context.GetEigenDevice<platform::GPUPlace>()) = gm * gm; gs.device(*context.eigen_device()) = gm * gm;
math::SelectedRowsAddToTensor<platform::GPUPlace, T> functor; math::SelectedRowsAddToTensor<platform::CUDADeviceContext, T> functor;
functor(context, *grad_square, moment); functor(context, *grad_square, moment);
// 3. update parameter // 3. update parameter
...@@ -139,13 +139,13 @@ struct SparseAdagradFunctor<platform::GPUPlace, T> { ...@@ -139,13 +139,13 @@ struct SparseAdagradFunctor<platform::GPUPlace, T> {
} }
}; };
template struct SparseAdagradFunctor<platform::GPUPlace, float>; template struct SparseAdagradFunctor<platform::CUDADeviceContext, float>;
template struct SparseAdagradFunctor<platform::GPUPlace, double>; template struct SparseAdagradFunctor<platform::CUDADeviceContext, double>;
} // namespace operators } // namespace operators
} // namespace paddle } // namespace paddle
namespace ops = paddle::operators; namespace ops = paddle::operators;
REGISTER_OP_GPU_KERNEL( REGISTER_OP_CUDA_KERNEL(
adagrad, ops::AdagradOpKernel<paddle::platform::GPUPlace, float>, adagrad, ops::AdagradOpKernel<paddle::platform::CUDADeviceContext, float>,
ops::AdagradOpKernel<paddle::platform::GPUPlace, double>); ops::AdagradOpKernel<paddle::platform::CUDADeviceContext, double>);
paddle/operators/adagrad_op.h
浏览文件 @ aa2f2fb4
...@@ -19,15 +19,15 @@ limitations under the License. */ ...@@ -19,15 +19,15 @@ limitations under the License. */
namespace paddle { namespace paddle {
namespace operators { namespace operators {
template <typename Place, typename T> template <typename DeviceContext, typename T>
struct SparseAdagradFunctor { struct SparseAdagradFunctor {
void operator()(const platform::DeviceContext& context, void operator()(const DeviceContext& context,
const framework::SelectedRows& grad, const framework::SelectedRows& grad,
const framework::Tensor& learning_rate, T epsilon, const framework::Tensor& learning_rate, T epsilon,
framework::Tensor* moment, framework::Tensor* param); framework::Tensor* moment, framework::Tensor* param);
}; };
template <typename Place, typename T> template <typename DeviceContext, typename T>
class AdagradOpKernel : public framework::OpKernel<T> { class AdagradOpKernel : public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext& ctx) const override { void Compute(const framework::ExecutionContext& ctx) const override {
...@@ -52,11 +52,11 @@ class AdagradOpKernel : public framework::OpKernel<T> { ...@@ -52,11 +52,11 @@ class AdagradOpKernel : public framework::OpKernel<T> {
auto param_out = framework::EigenVector<T>::Flatten(*param_out_tensor); auto param_out = framework::EigenVector<T>::Flatten(*param_out_tensor);
auto moment_out = framework::EigenVector<T>::Flatten(*moment_out_tensor); auto moment_out = framework::EigenVector<T>::Flatten(*moment_out_tensor);
auto place = ctx.GetEigenDevice<Place>(); auto* place = ctx.template device_context<DeviceContext>().eigen_device();
moment_out.device(place) = moment + grad * grad; moment_out.device(*place) = moment + grad * grad;
Eigen::DSizes<int, 1> m_dsize(moment_out_tensor->numel()); Eigen::DSizes<int, 1> m_dsize(moment_out_tensor->numel());
param_out.device(place) = param_out.device(*place) =
param - lr.broadcast(m_dsize) * grad / (moment_out.sqrt() + epsilon); param - lr.broadcast(m_dsize) * grad / (moment_out.sqrt() + epsilon);
} else if (grad_var->IsType<framework::SelectedRows>()) { } else if (grad_var->IsType<framework::SelectedRows>()) {
auto* param_tensor = ctx.Input<framework::Tensor>("Param"); auto* param_tensor = ctx.Input<framework::Tensor>("Param");
...@@ -65,8 +65,9 @@ class AdagradOpKernel : public framework::OpKernel<T> { ...@@ -65,8 +65,9 @@ class AdagradOpKernel : public framework::OpKernel<T> {
auto* moment_tensor = ctx.Input<framework::Tensor>("Moment"); auto* moment_tensor = ctx.Input<framework::Tensor>("Moment");
PADDLE_ENFORCE_EQ(moment_tensor, moment_out_tensor); PADDLE_ENFORCE_EQ(moment_tensor, moment_out_tensor);
SparseAdagradFunctor<Place, T> functor; SparseAdagradFunctor<DeviceContext, T> functor;
functor(ctx.device_context(), *ctx.Input<framework::SelectedRows>("Grad"), functor(ctx.template device_context<DeviceContext>(),
*ctx.Input<framework::SelectedRows>("Grad"),
*ctx.Input<framework::Tensor>("LearningRate"), epsilon, *ctx.Input<framework::Tensor>("LearningRate"), epsilon,
moment_out_tensor, param_out_tensor); moment_out_tensor, param_out_tensor);
} else { } else {
......
paddle/operators/adam_op.cc
浏览文件 @ aa2f2fb4
...@@ -112,11 +112,13 @@ adaptive estimates of lower-order moments. ...@@ -112,11 +112,13 @@ adaptive estimates of lower-order moments.
Adam updates: Adam updates:
$$moment_1_{out} = \beta_1 * moment_1 + (1 - \beta_1) * grad \break $$
moment_2_{out} = \beta_2 * moment_2 + (1 - \beta_2) * grad * grad \break moment\_1\_out = \beta_1 * moment\_1 + (1 - \beta_1) * grad \\
learningRate = learningRate * moment\_2_\out = \beta_2 * moment\_2 + (1 - \beta_2) * grad * grad \\
$\sqrt{(1 - \beta_2_{pow})}$ / (1 - \beta_1_{pow}) \break learning\_rate = learning\_rate *
paramOut = param - learningRate * moment_1/ ($\sqrt{(moment_2)} + \epsilon)$$ \frac{\sqrt{1 - \beta_{2\_pow}}}{1 - \beta_{1\_pow}} \\
param\_out = param - learning\_rate * \frac{moment\_1}{\sqrt{moment\_2} + \epsilon}
$$
)DOC"); )DOC");
} }
...@@ -126,6 +128,6 @@ paramOut = param - learningRate * moment_1/ ($\sqrt{(moment_2)} + \epsilon)$$ ...@@ -126,6 +128,6 @@ paramOut = param - learningRate * moment_1/ ($\sqrt{(moment_2)} + \epsilon)$$
namespace ops = paddle::operators; namespace ops = paddle::operators;
REGISTER_OP_WITHOUT_GRADIENT(adam, ops::AdamOp, ops::AdamOpMaker); REGISTER_OP_WITHOUT_GRADIENT(adam, ops::AdamOp, ops::AdamOpMaker);
REGISTER_OP_CPU_KERNEL(adam, REGISTER_OP_CPU_KERNEL(
ops::AdamOpKernel<paddle::platform::CPUPlace, float>, adam, ops::AdamOpKernel<paddle::platform::CPUDeviceContext, float>,
ops::AdamOpKernel<paddle::platform::CPUPlace, double>); ops::AdamOpKernel<paddle::platform::CPUDeviceContext, double>);
paddle/operators/adam_op.cu
浏览文件 @ aa2f2fb4
...@@ -16,6 +16,6 @@ ...@@ -16,6 +16,6 @@
#include "paddle/operators/adam_op.h" #include "paddle/operators/adam_op.h"
namespace ops = paddle::operators; namespace ops = paddle::operators;
REGISTER_OP_GPU_KERNEL(adam, REGISTER_OP_CUDA_KERNEL(
ops::AdamOpKernel<paddle::platform::GPUPlace, float>, adam, ops::AdamOpKernel<paddle::platform::CUDADeviceContext, float>,
ops::AdamOpKernel<paddle::platform::GPUPlace, double>); ops::AdamOpKernel<paddle::platform::CUDADeviceContext, double>);
paddle/operators/adam_op.h
浏览文件 @ aa2f2fb4
...@@ -19,7 +19,7 @@ limitations under the License. */ ...@@ -19,7 +19,7 @@ limitations under the License. */
namespace paddle { namespace paddle {
namespace operators { namespace operators {
template <typename Place, typename T> template <typename DeviceContext, typename T>
class AdamOpKernel : public framework::OpKernel<T> { class AdamOpKernel : public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext& ctx) const override { void Compute(const framework::ExecutionContext& ctx) const override {
...@@ -52,17 +52,17 @@ class AdamOpKernel : public framework::OpKernel<T> { ...@@ -52,17 +52,17 @@ class AdamOpKernel : public framework::OpKernel<T> {
auto param_out = framework::EigenVector<T>::Flatten(*param_out_tensor); auto param_out = framework::EigenVector<T>::Flatten(*param_out_tensor);
auto moment1_out = framework::EigenVector<T>::Flatten(*moment1_out_tensor); auto moment1_out = framework::EigenVector<T>::Flatten(*moment1_out_tensor);
auto moment2_out = framework::EigenVector<T>::Flatten(*moment2_out_tensor); auto moment2_out = framework::EigenVector<T>::Flatten(*moment2_out_tensor);
auto place = ctx.GetEigenDevice<Place>(); auto* place = ctx.template device_context<DeviceContext>().eigen_device();
moment1_out.device(place) = beta1 * moment1 + (1 - beta1) * grad; moment1_out.device(*place) = beta1 * moment1 + (1 - beta1) * grad;
moment2_out.device(place) = beta2 * moment2 + (1 - beta2) * grad.square(); moment2_out.device(*place) = beta2 * moment2 + (1 - beta2) * grad.square();
// All of these are tensors of 1 element // All of these are tensors of 1 element
auto lr_t = lr * (1 - beta2_pow).sqrt() / (1 - beta1_pow); auto lr_t = lr * (1 - beta2_pow).sqrt() / (1 - beta1_pow);
// Eigen does not support automatic broadcast // Eigen does not support automatic broadcast
// Get dimensions of moment vector to broadcast lr_t // Get dimensions of moment vector to broadcast lr_t
Eigen::DSizes<int, 1> m_dsize(moment1_out_tensor->numel()); Eigen::DSizes<int, 1> m_dsize(moment1_out_tensor->numel());
param_out.device(place) = param_out.device(*place) =
param - param -
lr_t.broadcast(m_dsize) * lr_t.broadcast(m_dsize) *
(moment1_out / (moment2_out.sqrt() + epsilon)); (moment1_out / (moment2_out.sqrt() + epsilon));
......
paddle/operators/adamax_op.cc
浏览文件 @ aa2f2fb4
...@@ -108,10 +108,10 @@ Adam algorithm based on the infinity norm. ...@@ -108,10 +108,10 @@ Adam algorithm based on the infinity norm.
Adamax updates: Adamax updates:
$$ $$
momentOut = \beta_{1} * moment + (1 - \beta_{1}) * grad \\ moment\_out = \beta_1 * moment + (1 - \beta_1) * grad \\
infNormOut = max(\beta_{2} * infNorm + \epsilon, |grad|) \\ inf\_norm\_out = max(\beta_2 * inf\_norm + \epsilon, |grad|) \\
learningRate = \frac{learningRate}{1 - \beta_{1}^{Beta1Pow}} \\ learning\_rate = \frac{learning\_rate}{1 - \beta_{1\_pow}} \\
paramOut = param - learningRate * \frac{momentOut}{infNormOut} param\_out = param - learning\_rate * \frac{moment\_out}{inf\_norm\_out}
$$ $$
The original paper does not have an epsilon attribute. The original paper does not have an epsilon attribute.
...@@ -127,6 +127,6 @@ division by 0 error. ...@@ -127,6 +127,6 @@ division by 0 error.
namespace ops = paddle::operators; namespace ops = paddle::operators;
REGISTER_OP_WITHOUT_GRADIENT(adamax, ops::AdamaxOp, ops::AdamaxOpMaker); REGISTER_OP_WITHOUT_GRADIENT(adamax, ops::AdamaxOp, ops::AdamaxOpMaker);
REGISTER_OP_CPU_KERNEL(adamax, REGISTER_OP_CPU_KERNEL(
ops::AdamaxOpKernel<paddle::platform::CPUPlace, float>, adamax, ops::AdamaxOpKernel<paddle::platform::CPUDeviceContext, float>,
ops::AdamaxOpKernel<paddle::platform::CPUPlace, double>); ops::AdamaxOpKernel<paddle::platform::CPUDeviceContext, double>);
paddle/operators/adamax_op.cu
浏览文件 @ aa2f2fb4
...@@ -16,6 +16,6 @@ ...@@ -16,6 +16,6 @@
#include "paddle/operators/adamax_op.h" #include "paddle/operators/adamax_op.h"
namespace ops = paddle::operators; namespace ops = paddle::operators;
REGISTER_OP_GPU_KERNEL(adamax, REGISTER_OP_CUDA_KERNEL(
ops::AdamaxOpKernel<paddle::platform::GPUPlace, float>, adamax, ops::AdamaxOpKernel<paddle::platform::CUDADeviceContext, float>,
ops::AdamaxOpKernel<paddle::platform::GPUPlace, double>); ops::AdamaxOpKernel<paddle::platform::CUDADeviceContext, double>);
paddle/operators/adamax_op.h
浏览文件 @ aa2f2fb4
...@@ -19,7 +19,7 @@ limitations under the License. */ ...@@ -19,7 +19,7 @@ limitations under the License. */
namespace paddle { namespace paddle {
namespace operators { namespace operators {
template <typename Place, typename T> template <typename DeviceContext, typename T>
class AdamaxOpKernel : public framework::OpKernel<T> { class AdamaxOpKernel : public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext& ctx) const override { void Compute(const framework::ExecutionContext& ctx) const override {
...@@ -51,14 +51,14 @@ class AdamaxOpKernel : public framework::OpKernel<T> { ...@@ -51,14 +51,14 @@ class AdamaxOpKernel : public framework::OpKernel<T> {
auto moment_out = framework::EigenVector<T>::Flatten(*moment_out_tensor); auto moment_out = framework::EigenVector<T>::Flatten(*moment_out_tensor);
auto inf_norm_out = auto inf_norm_out =
framework::EigenVector<T>::Flatten(*inf_norm_out_tensor); framework::EigenVector<T>::Flatten(*inf_norm_out_tensor);
auto place = ctx.GetEigenDevice<Place>(); auto* place = ctx.template device_context<DeviceContext>().eigen_device();
moment_out.device(place) = beta1 * moment + (1 - beta1) * grad; moment_out.device(*place) = beta1 * moment + (1 - beta1) * grad;
inf_norm_out.device(place) = inf_norm_out.device(*place) =
grad.abs().cwiseMax((beta2 * inf_norm) + epsilon); grad.abs().cwiseMax((beta2 * inf_norm) + epsilon);
auto lr_t = lr / (1 - beta1_pow); auto lr_t = lr / (1 - beta1_pow);
Eigen::DSizes<int, 1> m_dsize(moment_out_tensor->numel()); Eigen::DSizes<int, 1> m_dsize(moment_out_tensor->numel());
param_out.device(place) = param_out.device(*place) =
param - lr_t.broadcast(m_dsize) * (moment_out / inf_norm_out); param - lr_t.broadcast(m_dsize) * (moment_out / inf_norm_out);
} }
}; };
......
paddle/operators/auc_op.h
浏览文件 @ aa2f2fb4
...@@ -25,7 +25,7 @@ template <typename T, int MajorType = Eigen::RowMajor, ...@@ -25,7 +25,7 @@ template <typename T, int MajorType = Eigen::RowMajor,
typename IndexType = Eigen::DenseIndex> typename IndexType = Eigen::DenseIndex>
using EigenVector = framework::EigenVector<T, MajorType, IndexType>; using EigenVector = framework::EigenVector<T, MajorType, IndexType>;
template <typename Place, typename T> template <typename DeviceContext, typename T>
class AucKernel : public framework::OpKernel<T> { class AucKernel : public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext& ctx) const override { void Compute(const framework::ExecutionContext& ctx) const override {
......
paddle/operators/batch_norm_op.cc
浏览文件 @ aa2f2fb4
...@@ -135,7 +135,8 @@ The required data format for this layer is one of the following: ...@@ -135,7 +135,8 @@ The required data format for this layer is one of the following:
}; };
template <typename T> template <typename T>
class BatchNormKernel<platform::CPUPlace, T> : public framework::OpKernel<T> { class BatchNormKernel<platform::CPUDeviceContext, T>
: public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext &ctx) const override { void Compute(const framework::ExecutionContext &ctx) const override {
const float epsilon = ctx.Attr<float>("epsilon"); const float epsilon = ctx.Attr<float>("epsilon");
...@@ -318,12 +319,12 @@ class BatchNormGradOp : public framework::OperatorWithKernel { ...@@ -318,12 +319,12 @@ class BatchNormGradOp : public framework::OperatorWithKernel {
PADDLE_THROW("can't find Y@GRAD"); PADDLE_THROW("can't find Y@GRAD");
} }
return framework::OpKernelType(framework::ToDataType(t->type()), return framework::OpKernelType(framework::ToDataType(t->type()),
ctx.device_context()); ctx.GetPlace());
} }
}; };
template <typename T> template <typename T>
class BatchNormGradKernel<platform::CPUPlace, T> class BatchNormGradKernel<platform::CPUDeviceContext, T>
: public framework::OpKernel<T> { : public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext &ctx) const override { void Compute(const framework::ExecutionContext &ctx) const override {
...@@ -436,8 +437,9 @@ class BatchNormGradKernel<platform::CPUPlace, T> ...@@ -436,8 +437,9 @@ class BatchNormGradKernel<platform::CPUPlace, T>
namespace ops = paddle::operators; namespace ops = paddle::operators;
REGISTER_OP(batch_norm, ops::BatchNormOp, ops::BatchNormOpMaker, REGISTER_OP(batch_norm, ops::BatchNormOp, ops::BatchNormOpMaker,
batch_norm_grad, ops::BatchNormGradOp); batch_norm_grad, ops::BatchNormGradOp);
REGISTER_OP_CPU_KERNEL(batch_norm, REGISTER_OP_CPU_KERNEL(
ops::BatchNormKernel<paddle::platform::CPUPlace, float>); batch_norm,
ops::BatchNormKernel<paddle::platform::CPUDeviceContext, float>);
REGISTER_OP_CPU_KERNEL( REGISTER_OP_CPU_KERNEL(
batch_norm_grad, batch_norm_grad,
ops::BatchNormGradKernel<paddle::platform::CPUPlace, float>); ops::BatchNormGradKernel<paddle::platform::CPUDeviceContext, float>);
paddle/operators/batch_norm_op.cu.cc
浏览文件 @ aa2f2fb4
...@@ -47,7 +47,8 @@ void ExtractNCWHD(const framework::DDim &dims, ...@@ -47,7 +47,8 @@ void ExtractNCWHD(const framework::DDim &dims,
} }
template <typename T> template <typename T>
class BatchNormKernel<platform::GPUPlace, T> : public framework::OpKernel<T> { class BatchNormKernel<platform::CUDADeviceContext, T>
: public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext &ctx) const override { void Compute(const framework::ExecutionContext &ctx) const override {
PADDLE_ENFORCE(platform::is_gpu_place(ctx.GetPlace()), PADDLE_ENFORCE(platform::is_gpu_place(ctx.GetPlace()),
...@@ -121,11 +122,12 @@ class BatchNormKernel<platform::GPUPlace, T> : public framework::OpKernel<T> { ...@@ -121,11 +122,12 @@ class BatchNormKernel<platform::GPUPlace, T> : public framework::OpKernel<T> {
saved_mean->mutable_data<T>(ctx.GetPlace()); saved_mean->mutable_data<T>(ctx.GetPlace());
saved_variance->mutable_data<T>(ctx.GetPlace()); saved_variance->mutable_data<T>(ctx.GetPlace());
math::SetConstant<platform::GPUPlace, T> functor; auto &dev_ctx = ctx.template device_context<platform::CUDADeviceContext>();
functor(ctx.device_context(), saved_mean, 0); math::SetConstant<platform::CUDADeviceContext, T> functor;
functor(ctx.device_context(), saved_variance, 0); functor(dev_ctx, saved_mean, 0);
functor(dev_ctx, saved_variance, 0);
auto handle = ctx.cuda_device_context().cudnn_handle(); auto handle = dev_ctx.cudnn_handle();
// Now, depending on whether we are running test or not, we have two paths. // Now, depending on whether we are running test or not, we have two paths.
if (is_test) { if (is_test) {
...@@ -171,7 +173,7 @@ class BatchNormKernel<platform::GPUPlace, T> : public framework::OpKernel<T> { ...@@ -171,7 +173,7 @@ class BatchNormKernel<platform::GPUPlace, T> : public framework::OpKernel<T> {
}; };
template <typename T> template <typename T>
class BatchNormGradKernel<platform::GPUPlace, T> class BatchNormGradKernel<platform::CUDADeviceContext, T>
: public framework::OpKernel<T> { : public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext &ctx) const override { void Compute(const framework::ExecutionContext &ctx) const override {
...@@ -244,11 +246,12 @@ class BatchNormGradKernel<platform::GPUPlace, T> ...@@ -244,11 +246,12 @@ class BatchNormGradKernel<platform::GPUPlace, T>
const void *saved_mean_data = saved_mean->template data<T>(); const void *saved_mean_data = saved_mean->template data<T>();
const void *saved_var_data = saved_var->template data<T>(); const void *saved_var_data = saved_var->template data<T>();
auto &dev_ctx = ctx.template device_context<platform::CUDADeviceContext>();
CUDNN_ENFORCE(platform::dynload::cudnnBatchNormalizationBackward( CUDNN_ENFORCE(platform::dynload::cudnnBatchNormalizationBackward(
ctx.cuda_device_context().cudnn_handle(), mode_, dev_ctx.cudnn_handle(), mode_, CudnnDataType<T>::kOne(),
CudnnDataType<T>::kOne(), CudnnDataType<T>::kZero(), CudnnDataType<T>::kZero(), CudnnDataType<T>::kOne(),
CudnnDataType<T>::kOne(), CudnnDataType<T>::kZero(), data_desc_, CudnnDataType<T>::kZero(), data_desc_, x->template data<T>(),
x->template data<T>(), data_desc_, d_y->template data<T>(), data_desc_, data_desc_, d_y->template data<T>(), data_desc_,
d_x->template mutable_data<T>(ctx.GetPlace()), bn_param_desc_, d_x->template mutable_data<T>(ctx.GetPlace()), bn_param_desc_,
scale->template data<T>(), scale->template data<T>(),
d_scale->template mutable_data<T>(ctx.GetPlace()), d_scale->template mutable_data<T>(ctx.GetPlace()),
...@@ -266,8 +269,9 @@ class BatchNormGradKernel<platform::GPUPlace, T> ...@@ -266,8 +269,9 @@ class BatchNormGradKernel<platform::GPUPlace, T>
} // namespace paddle } // namespace paddle
namespace ops = paddle::operators; namespace ops = paddle::operators;
REGISTER_OP_GPU_KERNEL(batch_norm, REGISTER_OP_CUDA_KERNEL(
ops::BatchNormKernel<paddle::platform::GPUPlace, float>); batch_norm,
REGISTER_OP_GPU_KERNEL( ops::BatchNormKernel<paddle::platform::CUDADeviceContext, float>);
REGISTER_OP_CUDA_KERNEL(
batch_norm_grad, batch_norm_grad,
ops::BatchNormGradKernel<paddle::platform::GPUPlace, float>); ops::BatchNormGradKernel<paddle::platform::CUDADeviceContext, float>);
paddle/operators/batch_norm_op.h
浏览文件 @ aa2f2fb4
...@@ -34,13 +34,13 @@ inline TensorFormat StringToTensorFormat(const std::string& str) { ...@@ -34,13 +34,13 @@ inline TensorFormat StringToTensorFormat(const std::string& str) {
} }
} }
template <typename Place, typename T> template <typename DeviceContext, typename T>
class BatchNormKernel : public framework::OpKernel<T> { class BatchNormKernel : public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext& ctx) const override; void Compute(const framework::ExecutionContext& ctx) const override;
}; };
template <typename Place, typename T> template <typename DeviceContext, typename T>
class BatchNormGradKernel : public framework::OpKernel<T> { class BatchNormGradKernel : public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext& ctx) const override; void Compute(const framework::ExecutionContext& ctx) const override;
......
paddle/operators/bilinear_tensor_product_op.cc
浏览文件 @ aa2f2fb4
...@@ -159,9 +159,12 @@ REGISTER_OP(bilinear_tensor_product, ops::BilinearTensorProductOp, ...@@ -159,9 +159,12 @@ REGISTER_OP(bilinear_tensor_product, ops::BilinearTensorProductOp,
ops::BilinearTensorProductOpGrad); ops::BilinearTensorProductOpGrad);
REGISTER_OP_CPU_KERNEL( REGISTER_OP_CPU_KERNEL(
bilinear_tensor_product, bilinear_tensor_product,
ops::BilinearTensorProductKernel<paddle::platform::CPUPlace, float>, ops::BilinearTensorProductKernel<paddle::platform::CPUDeviceContext, float>,
ops::BilinearTensorProductKernel<paddle::platform::CPUPlace, double>); ops::BilinearTensorProductKernel<paddle::platform::CPUDeviceContext,
double>);
REGISTER_OP_CPU_KERNEL( REGISTER_OP_CPU_KERNEL(
bilinear_tensor_product_grad, bilinear_tensor_product_grad,
ops::BilinearTensorProductGradKernel<paddle::platform::CPUPlace, float>, ops::BilinearTensorProductGradKernel<paddle::platform::CPUDeviceContext,
ops::BilinearTensorProductGradKernel<paddle::platform::CPUPlace, double>); float>,
ops::BilinearTensorProductGradKernel<paddle::platform::CPUDeviceContext,
double>);
paddle/operators/bilinear_tensor_product_op.cu
浏览文件 @ aa2f2fb4
...@@ -16,11 +16,15 @@ limitations under the License. */ ...@@ -16,11 +16,15 @@ limitations under the License. */
#include "paddle/operators/bilinear_tensor_product_op.h" #include "paddle/operators/bilinear_tensor_product_op.h"
namespace ops = paddle::operators; namespace ops = paddle::operators;
REGISTER_OP_GPU_KERNEL( REGISTER_OP_CUDA_KERNEL(
bilinear_tensor_product, bilinear_tensor_product,
ops::BilinearTensorProductKernel<paddle::platform::GPUPlace, float>, ops::BilinearTensorProductKernel<paddle::platform::CUDADeviceContext,
ops::BilinearTensorProductKernel<paddle::platform::GPUPlace, double>); float>,
REGISTER_OP_GPU_KERNEL( ops::BilinearTensorProductKernel<paddle::platform::CUDADeviceContext,
double>);
REGISTER_OP_CUDA_KERNEL(
bilinear_tensor_product_grad, bilinear_tensor_product_grad,
ops::BilinearTensorProductGradKernel<paddle::platform::GPUPlace, float>, ops::BilinearTensorProductGradKernel<paddle::platform::CUDADeviceContext,
ops::BilinearTensorProductGradKernel<paddle::platform::GPUPlace, double>); float>,
ops::BilinearTensorProductGradKernel<paddle::platform::CUDADeviceContext,
double>);
paddle/operators/bilinear_tensor_product_op.h
浏览文件 @ aa2f2fb4
...@@ -27,7 +27,7 @@ template <typename T, int MajorType = Eigen::RowMajor, ...@@ -27,7 +27,7 @@ template <typename T, int MajorType = Eigen::RowMajor,
typename IndexType = Eigen::DenseIndex> typename IndexType = Eigen::DenseIndex>
using EigenMatrix = framework::EigenMatrix<T, MajorType, IndexType>; using EigenMatrix = framework::EigenMatrix<T, MajorType, IndexType>;
template <typename Place, typename T> template <typename DeviceContext, typename T>
class BilinearTensorProductKernel : public framework::OpKernel<T> { class BilinearTensorProductKernel : public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext& ctx) const override { void Compute(const framework::ExecutionContext& ctx) const override {
...@@ -46,7 +46,8 @@ class BilinearTensorProductKernel : public framework::OpKernel<T> { ...@@ -46,7 +46,8 @@ class BilinearTensorProductKernel : public framework::OpKernel<T> {
int out_dim = weight_dims[0]; int out_dim = weight_dims[0];
auto x_dim = weight_dims[1]; auto x_dim = weight_dims[1];
auto y_dim = weight_dims[2]; auto y_dim = weight_dims[2];
auto place = ctx.GetEigenDevice<Place>(); auto& place = *ctx.template device_context<DeviceContext>().eigen_device();
auto& dev_ctx = ctx.template device_context<DeviceContext>();
// Create the intermediate variable to caculate the result of // Create the intermediate variable to caculate the result of
// Input(X) multiplied by Input(Weight_i), the formula is: // Input(X) multiplied by Input(Weight_i), the formula is:
...@@ -60,9 +61,9 @@ class BilinearTensorProductKernel : public framework::OpKernel<T> { ...@@ -60,9 +61,9 @@ class BilinearTensorProductKernel : public framework::OpKernel<T> {
auto output_col_vec = output_mat.chip(i, 1); auto output_col_vec = output_mat.chip(i, 1);
Tensor weight_mat = Tensor weight_mat =
weight->Slice(i, i + 1).Resize(framework::make_ddim({x_dim, y_dim})); weight->Slice(i, i + 1).Resize(framework::make_ddim({x_dim, y_dim}));
math::gemm<Place, T>(ctx.device_context(), CblasNoTrans, CblasNoTrans, math::gemm<DeviceContext, T>(dev_ctx, CblasNoTrans, CblasNoTrans,
batch_size, y_dim, x_dim, 1, x->data<T>(), batch_size, y_dim, x_dim, 1, x->data<T>(),
weight_mat.data<T>(), 0, left_mul.data<T>()); weight_mat.data<T>(), 0, left_mul.data<T>());
output_col_vec.device(place) = output_col_vec.device(place) =
(left_mul_mat * y_mat).sum(Eigen::DSizes<int, 1>(1)); (left_mul_mat * y_mat).sum(Eigen::DSizes<int, 1>(1));
} }
...@@ -74,7 +75,7 @@ class BilinearTensorProductKernel : public framework::OpKernel<T> { ...@@ -74,7 +75,7 @@ class BilinearTensorProductKernel : public framework::OpKernel<T> {
} }
}; };
template <typename Place, typename T> template <typename DeviceContext, typename T>
class BilinearTensorProductGradKernel : public framework::OpKernel<T> { class BilinearTensorProductGradKernel : public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext& ctx) const override { void Compute(const framework::ExecutionContext& ctx) const override {
...@@ -96,8 +97,8 @@ class BilinearTensorProductGradKernel : public framework::OpKernel<T> { ...@@ -96,8 +97,8 @@ class BilinearTensorProductGradKernel : public framework::OpKernel<T> {
auto x_mat = EigenMatrix<T>::From(*x); auto x_mat = EigenMatrix<T>::From(*x);
auto y_mat = EigenMatrix<T>::From(*y); auto y_mat = EigenMatrix<T>::From(*y);
auto d_out_mat = EigenMatrix<T>::From(*d_out); auto d_out_mat = EigenMatrix<T>::From(*d_out);
auto place = ctx.GetEigenDevice<Place>(); auto& place = *ctx.template device_context<DeviceContext>().eigen_device();
auto& dev_ctx = ctx.template device_context<DeviceContext>();
// Create the intermediate variable to caculate the Output(Y@Grad). // Create the intermediate variable to caculate the Output(Y@Grad).
Tensor x_scale; Tensor x_scale;
x_scale.mutable_data<T>(framework::make_ddim({batch_size, x_dim}), x_scale.mutable_data<T>(framework::make_ddim({batch_size, x_dim}),
...@@ -110,18 +111,18 @@ class BilinearTensorProductGradKernel : public framework::OpKernel<T> { ...@@ -110,18 +111,18 @@ class BilinearTensorProductGradKernel : public framework::OpKernel<T> {
ctx.GetPlace()); ctx.GetPlace());
auto y_scale_mat = EigenMatrix<T>::From(y_scale); auto y_scale_mat = EigenMatrix<T>::From(y_scale);
math::SetConstant<Place, T> set_zero; math::SetConstant<DeviceContext, T> set_zero;
// Set Output(X@Grad) be zero. // Set Output(X@Grad) be zero.
if (d_x) { if (d_x) {
d_x->mutable_data<T>(ctx.GetPlace()); d_x->mutable_data<T>(ctx.GetPlace());
set_zero(ctx.device_context(), d_x, static_cast<T>(0)); set_zero(dev_ctx, d_x, static_cast<T>(0));
} }
// Set Output(Y@Grad) be zero. // Set Output(Y@Grad) be zero.
if (d_y) { if (d_y) {
d_y->mutable_data<T>(ctx.GetPlace()); d_y->mutable_data<T>(ctx.GetPlace());
set_zero(ctx.device_context(), d_y, static_cast<T>(0)); set_zero(dev_ctx, d_y, static_cast<T>(0));
} }
// Caculate the Output(X@Grad) and Output(Y@Grad). // Caculate the Output(X@Grad) and Output(Y@Grad).
...@@ -137,18 +138,18 @@ class BilinearTensorProductGradKernel : public framework::OpKernel<T> { ...@@ -137,18 +138,18 @@ class BilinearTensorProductGradKernel : public framework::OpKernel<T> {
output_vec.reshape(Eigen::DSizes<int, 2>(batch_size, 1)) output_vec.reshape(Eigen::DSizes<int, 2>(batch_size, 1))
.broadcast(bcast_for_x) * .broadcast(bcast_for_x) *
y_mat; y_mat;
math::gemm<Place, T>(ctx.device_context(), CblasNoTrans, CblasTrans, math::gemm<DeviceContext, T>(
batch_size, x_dim, y_dim, 1, y_scale.data<T>(), dev_ctx, CblasNoTrans, CblasTrans, batch_size, x_dim, y_dim, 1,
weight_i.data<T>(), 1, d_x->data<T>()); y_scale.data<T>(), weight_i.data<T>(), 1, d_x->data<T>());
} }
if (d_y) { if (d_y) {
x_scale_mat.device(place) = x_scale_mat.device(place) =
output_vec.reshape(Eigen::DSizes<int, 2>(batch_size, 1)) output_vec.reshape(Eigen::DSizes<int, 2>(batch_size, 1))
.broadcast(bcast_for_y) * .broadcast(bcast_for_y) *
x_mat; x_mat;
math::gemm<Place, T>(ctx.device_context(), CblasNoTrans, CblasNoTrans, math::gemm<DeviceContext, T>(
batch_size, y_dim, x_dim, 1, x_scale.data<T>(), dev_ctx, CblasNoTrans, CblasNoTrans, batch_size, y_dim, x_dim, 1,
weight_i.data<T>(), 1, d_y->data<T>()); x_scale.data<T>(), weight_i.data<T>(), 1, d_y->data<T>());
} }
} }
} }
...@@ -165,9 +166,9 @@ class BilinearTensorProductGradKernel : public framework::OpKernel<T> { ...@@ -165,9 +166,9 @@ class BilinearTensorProductGradKernel : public framework::OpKernel<T> {
output_vec.reshape(Eigen::DSizes<int, 2>(batch_size, 1)) output_vec.reshape(Eigen::DSizes<int, 2>(batch_size, 1))
.broadcast(bcast_for_weight) * .broadcast(bcast_for_weight) *
x_mat; x_mat;
math::gemm<Place, T>(ctx.device_context(), CblasTrans, CblasNoTrans, math::gemm<DeviceContext, T>(dev_ctx, CblasTrans, CblasNoTrans, x_dim,
x_dim, y_dim, batch_size, 1, x_scale.data<T>(), y_dim, batch_size, 1, x_scale.data<T>(),
y->data<T>(), 0, d_weight_i.data<T>()); y->data<T>(), 0, d_weight_i.data<T>());
} }
} }
......
paddle/operators/cast_op.cc
浏览文件 @ aa2f2fb4
...@@ -68,10 +68,11 @@ class CastOpGradMaker : public framework::SingleGradOpDescMaker { ...@@ -68,10 +68,11 @@ class CastOpGradMaker : public framework::SingleGradOpDescMaker {
} // namespace paddle } // namespace paddle
namespace ops = paddle::operators; namespace ops = paddle::operators;
using CPU = paddle::platform::CPUPlace; using CPU = paddle::platform::CPUDeviceContext;
REGISTER_OP_WITH_KERNEL(cast, ops::CastOpGradMaker, ops::CastOpInferShape, REGISTER_OP_WITH_KERNEL(cast, ops::CastOpGradMaker, ops::CastOpInferShape,
ops::CastOpProtoMaker); ops::CastOpProtoMaker);
REGISTER_OP_CPU_KERNEL(cast, ops::CastOpKernel<CPU, float>, REGISTER_OP_CPU_KERNEL(cast, ops::CastOpKernel<CPU, float>,
ops::CastOpKernel<CPU, double>, ops::CastOpKernel<CPU, double>,
ops::CastOpKernel<CPU, int>, ops::CastOpKernel<CPU, int>,
ops::CastOpKernel<CPU, int64_t>); ops::CastOpKernel<CPU, int64_t>,
ops::CastOpKernel<CPU, bool>);
paddle/operators/cast_op.cu
浏览文件 @ aa2f2fb4
...@@ -16,7 +16,8 @@ ...@@ -16,7 +16,8 @@
template <typename T> template <typename T>
using CastOpKernel = using CastOpKernel =
paddle::operators::CastOpKernel<paddle::platform::GPUPlace, T>; paddle::operators::CastOpKernel<paddle::platform::CUDADeviceContext, T>;
REGISTER_OP_GPU_KERNEL(cast, CastOpKernel<float>, CastOpKernel<double>, REGISTER_OP_CUDA_KERNEL(cast, CastOpKernel<float>, CastOpKernel<double>,
CastOpKernel<int>, CastOpKernel<int64_t>); CastOpKernel<int>, CastOpKernel<int64_t>,
CastOpKernel<bool>);
paddle/operators/cast_op.h
浏览文件 @ aa2f2fb4
...@@ -27,13 +27,13 @@ struct CastOpTransformFunctor { ...@@ -27,13 +27,13 @@ struct CastOpTransformFunctor {
HOSTDEVICE OutT operator()(InT in) const { return static_cast<OutT>(in); } HOSTDEVICE OutT operator()(InT in) const { return static_cast<OutT>(in); }
}; };
template <typename Place, typename InT> template <typename DeviceContext, typename InT>
struct CastOpFunctor { struct CastOpFunctor {
const framework::Tensor* in_; const framework::Tensor* in_;
framework::Tensor* out_; framework::Tensor* out_;
const platform::DeviceContext& ctx_; const DeviceContext& ctx_;
CastOpFunctor(const framework::Tensor* in, framework::Tensor* out, CastOpFunctor(const framework::Tensor* in, framework::Tensor* out,
const platform::DeviceContext& ctx) const DeviceContext& ctx)
: in_(in), out_(out), ctx_(ctx) {} : in_(in), out_(out), ctx_(ctx) {}
template <typename OutT> template <typename OutT>
...@@ -42,13 +42,13 @@ struct CastOpFunctor { ...@@ -42,13 +42,13 @@ struct CastOpFunctor {
auto numel = in_->numel(); auto numel = in_->numel();
auto* in_end = in_begin + numel; auto* in_end = in_begin + numel;
auto* out_begin = out_->mutable_data<OutT>(ctx_.GetPlace()); auto* out_begin = out_->mutable_data<OutT>(ctx_.GetPlace());
platform::Transform<Place> trans; platform::Transform<DeviceContext> trans;
trans(ctx_, in_begin, in_end, out_begin, trans(ctx_, in_begin, in_end, out_begin,
CastOpTransformFunctor<InT, OutT>()); CastOpTransformFunctor<InT, OutT>());
} }
}; };
template <typename Place, typename InT> template <typename DeviceContext, typename InT>
class CastOpKernel : public framework::OpKernel<InT> { class CastOpKernel : public framework::OpKernel<InT> {
public: public:
void Compute(const framework::ExecutionContext& context) const override { void Compute(const framework::ExecutionContext& context) const override {
...@@ -56,7 +56,8 @@ class CastOpKernel : public framework::OpKernel<InT> { ...@@ -56,7 +56,8 @@ class CastOpKernel : public framework::OpKernel<InT> {
auto* out = context.Output<framework::Tensor>("Out"); auto* out = context.Output<framework::Tensor>("Out");
framework::VisitDataType( framework::VisitDataType(
static_cast<framework::DataType>(context.Attr<int>("out_dtype")), static_cast<framework::DataType>(context.Attr<int>("out_dtype")),
CastOpFunctor<Place, InT>(in, out, context.device_context())); CastOpFunctor<DeviceContext, InT>(
in, out, context.template device_context<DeviceContext>()));
} }
}; };
......
paddle/operators/chunk_eval_op.h
浏览文件 @ aa2f2fb4
...@@ -23,7 +23,7 @@ namespace operators { ...@@ -23,7 +23,7 @@ namespace operators {
using Tensor = framework::Tensor; using Tensor = framework::Tensor;
using LoDTensor = framework::LoDTensor; using LoDTensor = framework::LoDTensor;
template <typename Place, typename T> template <typename DeviceContext, typename T>
class ChunkEvalKernel : public framework::OpKernel<T> { class ChunkEvalKernel : public framework::OpKernel<T> {
public: public:
struct Segment { struct Segment {
......
paddle/operators/clip_by_norm_op.cc
浏览文件 @ aa2f2fb4
...@@ -71,4 +71,5 @@ namespace ops = paddle::operators; ...@@ -71,4 +71,5 @@ namespace ops = paddle::operators;
REGISTER_OP_WITHOUT_GRADIENT(clip_by_norm, ops::ClipByNormOp, REGISTER_OP_WITHOUT_GRADIENT(clip_by_norm, ops::ClipByNormOp,
ops::ClipByNormOpMaker); ops::ClipByNormOpMaker);
REGISTER_OP_CPU_KERNEL( REGISTER_OP_CPU_KERNEL(
clip_by_norm, ops::ClipByNormKernel<paddle::platform::CPUPlace, float>); clip_by_norm,
ops::ClipByNormKernel<paddle::platform::CPUDeviceContext, float>);
paddle/operators/clip_by_norm_op.cu
浏览文件 @ aa2f2fb4
...@@ -15,5 +15,6 @@ ...@@ -15,5 +15,6 @@
#include "paddle/operators/clip_by_norm_op.h" #include "paddle/operators/clip_by_norm_op.h"
namespace ops = paddle::operators; namespace ops = paddle::operators;
REGISTER_OP_GPU_KERNEL( REGISTER_OP_CUDA_KERNEL(
clip_by_norm, ops::ClipByNormKernel<paddle::platform::GPUPlace, float>); clip_by_norm,
ops::ClipByNormKernel<paddle::platform::CUDADeviceContext, float>);
paddle/operators/clip_by_norm_op.h
浏览文件 @ aa2f2fb4
...@@ -26,7 +26,7 @@ template <typename T, int MajorType = Eigen::RowMajor, ...@@ -26,7 +26,7 @@ template <typename T, int MajorType = Eigen::RowMajor,
typename IndexType = Eigen::DenseIndex> typename IndexType = Eigen::DenseIndex>
using EigenVector = framework::EigenVector<T, MajorType, IndexType>; using EigenVector = framework::EigenVector<T, MajorType, IndexType>;
template <typename Place, typename T> template <typename DeviceContext, typename T>
class ClipByNormKernel : public framework::OpKernel<T> { class ClipByNormKernel : public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext& context) const override { void Compute(const framework::ExecutionContext& context) const override {
...@@ -38,7 +38,8 @@ class ClipByNormKernel : public framework::OpKernel<T> { ...@@ -38,7 +38,8 @@ class ClipByNormKernel : public framework::OpKernel<T> {
auto x = EigenVector<T>::Flatten(*input); auto x = EigenVector<T>::Flatten(*input);
auto out = EigenVector<T>::Flatten(*output); auto out = EigenVector<T>::Flatten(*output);
auto x_norm = x.square().sum().sqrt(); auto x_norm = x.square().sum().sqrt();
auto place = context.GetEigenDevice<Place>(); auto& place =
*context.template device_context<DeviceContext>().eigen_device();
auto temp = (x_norm <= max_norm).template cast<T>().eval(); auto temp = (x_norm <= max_norm).template cast<T>().eval();
auto scaling = temp + (static_cast<T>(1) - temp) * max_norm / x_norm; auto scaling = temp + (static_cast<T>(1) - temp) * max_norm / x_norm;
......
paddle/operators/clip_op.cc
浏览文件 @ aa2f2fb4
...@@ -83,7 +83,7 @@ class ClipOpGrad : public framework::OperatorWithKernel { ...@@ -83,7 +83,7 @@ class ClipOpGrad : public framework::OperatorWithKernel {
namespace ops = paddle::operators; namespace ops = paddle::operators;
REGISTER_OP(clip, ops::ClipOp, ops::ClipOpMaker<float>, clip_grad, REGISTER_OP(clip, ops::ClipOp, ops::ClipOpMaker<float>, clip_grad,
ops::ClipOpGrad); ops::ClipOpGrad);
REGISTER_OP_CPU_KERNEL(clip, REGISTER_OP_CPU_KERNEL(
ops::ClipKernel<paddle::platform::CPUPlace, float>); clip, ops::ClipKernel<paddle::platform::CPUDeviceContext, float>);
REGISTER_OP_CPU_KERNEL(clip_grad, REGISTER_OP_CPU_KERNEL(
ops::ClipGradKernel<paddle::platform::CPUPlace, float>); clip_grad, ops::ClipGradKernel<paddle::platform::CPUDeviceContext, float>);
paddle/operators/clip_op.cu
浏览文件 @ aa2f2fb4
...@@ -15,7 +15,7 @@ ...@@ -15,7 +15,7 @@
#include "paddle/operators/clip_op.h" #include "paddle/operators/clip_op.h"
namespace ops = paddle::operators; namespace ops = paddle::operators;
REGISTER_OP_GPU_KERNEL(clip, REGISTER_OP_CUDA_KERNEL(
ops::ClipKernel<paddle::platform::GPUPlace, float>); clip, ops::ClipKernel<paddle::platform::CUDADeviceContext, float>);
REGISTER_OP_GPU_KERNEL(clip_grad, REGISTER_OP_CUDA_KERNEL(
ops::ClipGradKernel<paddle::platform::GPUPlace, float>); clip_grad, ops::ClipGradKernel<paddle::platform::CUDADeviceContext, float>);
paddle/operators/clip_op.h
浏览文件 @ aa2f2fb4
...@@ -55,7 +55,7 @@ class ClipGradFunctor { ...@@ -55,7 +55,7 @@ class ClipGradFunctor {
T max_; T max_;
}; };
template <typename Place, typename T> template <typename DeviceContext, typename T>
class ClipKernel : public framework::OpKernel<T> { class ClipKernel : public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext& context) const override { void Compute(const framework::ExecutionContext& context) const override {
...@@ -66,13 +66,13 @@ class ClipKernel : public framework::OpKernel<T> { ...@@ -66,13 +66,13 @@ class ClipKernel : public framework::OpKernel<T> {
T* out_data = out->mutable_data<T>(context.GetPlace()); T* out_data = out->mutable_data<T>(context.GetPlace());
const T* x_data = x->data<T>(); const T* x_data = x->data<T>();
int64_t numel = x->numel(); int64_t numel = x->numel();
Transform<Place> trans; Transform<DeviceContext> trans;
trans(context.device_context(), x_data, x_data + numel, out_data, trans(context.template device_context<DeviceContext>(), x_data,
ClipFunctor<T>(min, max)); x_data + numel, out_data, ClipFunctor<T>(min, max));
} }
}; };
template <typename Place, typename T> template <typename DeviceContext, typename T>
class ClipGradKernel : public framework::OpKernel<T> { class ClipGradKernel : public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext& context) const override { void Compute(const framework::ExecutionContext& context) const override {
...@@ -86,9 +86,9 @@ class ClipGradKernel : public framework::OpKernel<T> { ...@@ -86,9 +86,9 @@ class ClipGradKernel : public framework::OpKernel<T> {
auto* d_x_data = d_x->mutable_data<T>(context.GetPlace()); auto* d_x_data = d_x->mutable_data<T>(context.GetPlace());
const T* d_out_data = d_out->data<T>(); const T* d_out_data = d_out->data<T>();
const T* x_data = x->data<T>(); const T* x_data = x->data<T>();
Transform<Place> trans; Transform<DeviceContext> trans;
trans(context.device_context(), d_out_data, d_out_data + numel, x_data, trans(context.template device_context<DeviceContext>(), d_out_data,
d_x_data, ClipGradFunctor<T>(min, max)); d_out_data + numel, x_data, d_x_data, ClipGradFunctor<T>(min, max));
} }
} }
}; };
......
paddle/operators/compare_op.cu
浏览文件 @ aa2f2fb4
...@@ -14,10 +14,10 @@ ...@@ -14,10 +14,10 @@
#include "paddle/operators/compare_op.h" #include "paddle/operators/compare_op.h"
REGISTER_LOGICAL_KERNEL(less_than, GPU, paddle::operators::LessThanFunctor); REGISTER_LOGICAL_KERNEL(less_than, CUDA, paddle::operators::LessThanFunctor);
REGISTER_LOGICAL_KERNEL(less_equal, GPU, paddle::operators::LessEqualFunctor); REGISTER_LOGICAL_KERNEL(less_equal, CUDA, paddle::operators::LessEqualFunctor);
REGISTER_LOGICAL_KERNEL(greater_than, GPU, REGISTER_LOGICAL_KERNEL(greater_than, CUDA,
paddle::operators::GreaterThanFunctor); paddle::operators::GreaterThanFunctor);
REGISTER_LOGICAL_KERNEL(greater_equal, GPU, REGISTER_LOGICAL_KERNEL(greater_equal, CUDA,
paddle::operators::GreaterEqualFunctor); paddle::operators::GreaterEqualFunctor);
REGISTER_LOGICAL_KERNEL(equal, GPU, paddle::operators::EqualFunctor); REGISTER_LOGICAL_KERNEL(equal, CUDA, paddle::operators::EqualFunctor);
paddle/operators/compare_op.h
浏览文件 @ aa2f2fb4
...@@ -59,7 +59,7 @@ struct EqualFunctor { ...@@ -59,7 +59,7 @@ struct EqualFunctor {
} }
}; };
template <typename Place, typename Functor> template <typename DeviceContext, typename Functor>
class CompareOpKernel class CompareOpKernel
: public framework::OpKernel<typename Functor::ELEM_TYPE> { : public framework::OpKernel<typename Functor::ELEM_TYPE> {
public: public:
...@@ -69,24 +69,23 @@ class CompareOpKernel ...@@ -69,24 +69,23 @@ class CompareOpKernel
auto* y = context.Input<framework::Tensor>("Y"); auto* y = context.Input<framework::Tensor>("Y");
auto* out = context.Output<framework::Tensor>("Out"); auto* out = context.Output<framework::Tensor>("Out");
Functor binary_func; Functor binary_func;
platform::Transform<Place> trans; platform::Transform<DeviceContext> trans;
trans(context.device_context(), x->data<T>(), x->data<T>() + x->numel(), trans(context.template device_context<DeviceContext>(), x->data<T>(),
y->data<T>(), out->mutable_data<bool>(context.GetPlace()), x->data<T>() + x->numel(), y->data<T>(),
binary_func); out->mutable_data<bool>(context.GetPlace()), binary_func);
} }
}; };
} // namespace operators } // namespace operators
} // namespace paddle } // namespace paddle
#define REGISTER_LOGICAL_KERNEL(op_type, dev, functor) \ #define REGISTER_LOGICAL_KERNEL(op_type, dev, functor) \
REGISTER_OP_##dev##_KERNEL( \ REGISTER_OP_##dev##_KERNEL( \
op_type, \ op_type, ::paddle::operators::CompareOpKernel< \
::paddle::operators::CompareOpKernel<::paddle::platform::dev##Place, \ ::paddle::platform::dev##DeviceContext, functor<int>>, \
functor<int>>, \ ::paddle::operators::CompareOpKernel< \
::paddle::operators::CompareOpKernel<::paddle::platform::dev##Place, \ ::paddle::platform::dev##DeviceContext, functor<int64_t>>, \
functor<int64_t>>, \ ::paddle::operators::CompareOpKernel< \
::paddle::operators::CompareOpKernel<::paddle::platform::dev##Place, \ ::paddle::platform::dev##DeviceContext, functor<float>>, \
functor<float>>, \ ::paddle::operators::CompareOpKernel< \
::paddle::operators::CompareOpKernel<::paddle::platform::dev##Place, \ ::paddle::platform::dev##DeviceContext, functor<double>>);
functor<double>>);
paddle/operators/concat_op.cu.cc
浏览文件 @ aa2f2fb4
...@@ -14,7 +14,8 @@ limitations under the License. */ ...@@ -14,7 +14,8 @@ limitations under the License. */
#include "paddle/operators/concat_op.h" #include "paddle/operators/concat_op.h"
namespace ops = paddle::operators; namespace ops = paddle::operators;
REGISTER_OP_GPU_KERNEL(concat, REGISTER_OP_CUDA_KERNEL(
ops::ConcatKernel<paddle::platform::GPUPlace, float>); concat, ops::ConcatKernel<paddle::platform::CUDADeviceContext, float>);
REGISTER_OP_GPU_KERNEL( REGISTER_OP_CUDA_KERNEL(
concat_grad, ops::ConcatGradKernel<paddle::platform::GPUPlace, float>); concat_grad,
ops::ConcatGradKernel<paddle::platform::CUDADeviceContext, float>);
paddle/operators/concat_op.h
浏览文件 @ aa2f2fb4
...@@ -21,7 +21,7 @@ limitations under the License. */ ...@@ -21,7 +21,7 @@ limitations under the License. */
namespace paddle { namespace paddle {
namespace operators { namespace operators {
template <typename Place, typename T> template <typename DeviceContext, typename T>
class ConcatKernel : public framework::OpKernel<T> { class ConcatKernel : public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext& ctx) const override { void Compute(const framework::ExecutionContext& ctx) const override {
...@@ -43,7 +43,7 @@ class ConcatKernel : public framework::OpKernel<T> { ...@@ -43,7 +43,7 @@ class ConcatKernel : public framework::OpKernel<T> {
} }
}; };
template <typename Place, typename T> template <typename DeviceContext, typename T>
class ConcatGradKernel : public framework::OpKernel<T> { class ConcatGradKernel : public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext& ctx) const { void Compute(const framework::ExecutionContext& ctx) const {
......
paddle/operators/conditional_block_op.cc
浏览文件 @ aa2f2fb4
...@@ -142,9 +142,9 @@ class ConditionalBlockGradOp : public ConditionalOp { ...@@ -142,9 +142,9 @@ class ConditionalBlockGradOp : public ConditionalOp {
continue; continue;
} }
auto new_in_grad_name = cur_scope.Rename(in_grad_name); auto new_in_grad_name = cur_scope.Rename(in_grad_name);
auto assign = auto assign = framework::OpRegistry::CreateOp(
framework::OpRegistry::CreateOp("assign", {{"X", {new_in_grad_name}}}, "assign", {{"X", {new_in_grad_name}}}, {{"Out", {out_grad_name}}},
{{"Out", {out_grad_name}}}, {}); framework::AttributeMap{});
assign->Run(cur_scope, dev_ctx); assign->Run(cur_scope, dev_ctx);
cur_scope.Rename(new_in_grad_name, in_grad_name); cur_scope.Rename(new_in_grad_name, in_grad_name);
} }
......
paddle/operators/conv_cudnn_op.cc
浏览文件 @ aa2f2fb4
...@@ -57,18 +57,20 @@ REGISTER_OP(conv2d_cudnn, ops::ConvOp, ops::CudnnConv2DOpMaker, ...@@ -57,18 +57,20 @@ REGISTER_OP(conv2d_cudnn, ops::ConvOp, ops::CudnnConv2DOpMaker,
REGISTER_OP(conv3d_cudnn, ops::ConvOp, ops::CudnnConv3DOpMaker, REGISTER_OP(conv3d_cudnn, ops::ConvOp, ops::CudnnConv3DOpMaker,
conv3d_cudnn_grad, ops::ConvOpGrad); conv3d_cudnn_grad, ops::ConvOpGrad);
REGISTER_OP_CPU_KERNEL(conv2d_cudnn, REGISTER_OP_CPU_KERNEL(
ops::GemmConvKernel<paddle::platform::CPUPlace, float>, conv2d_cudnn,
ops::GemmConvKernel<paddle::platform::CPUPlace, double>); ops::GemmConvKernel<paddle::platform::CPUDeviceContext, float>,
ops::GemmConvKernel<paddle::platform::CPUDeviceContext, double>);
REGISTER_OP_CPU_KERNEL( REGISTER_OP_CPU_KERNEL(
conv2d_cudnn_grad, conv2d_cudnn_grad,
ops::GemmConvGradKernel<paddle::platform::CPUPlace, float>, ops::GemmConvGradKernel<paddle::platform::CPUDeviceContext, float>,
ops::GemmConvGradKernel<paddle::platform::CPUPlace, double>); ops::GemmConvGradKernel<paddle::platform::CPUDeviceContext, double>);
REGISTER_OP_CPU_KERNEL(conv3d_cudnn, REGISTER_OP_CPU_KERNEL(
ops::GemmConvKernel<paddle::platform::CPUPlace, float>, conv3d_cudnn,
ops::GemmConvKernel<paddle::platform::CPUPlace, double>); ops::GemmConvKernel<paddle::platform::CPUDeviceContext, float>,
ops::GemmConvKernel<paddle::platform::CPUDeviceContext, double>);
REGISTER_OP_CPU_KERNEL( REGISTER_OP_CPU_KERNEL(
conv3d_cudnn_grad, conv3d_cudnn_grad,
ops::GemmConvGradKernel<paddle::platform::CPUPlace, float>, ops::GemmConvGradKernel<paddle::platform::CPUDeviceContext, float>,
ops::GemmConvGradKernel<paddle::platform::CPUPlace, double>); ops::GemmConvGradKernel<paddle::platform::CPUDeviceContext, double>);
paddle/operators/conv_cudnn_op.cu.cc
浏览文件 @ aa2f2fb4
...@@ -118,7 +118,8 @@ class CudnnConvOpKernel : public framework::OpKernel<T> { ...@@ -118,7 +118,8 @@ class CudnnConvOpKernel : public framework::OpKernel<T> {
} }
// ------------------- cudnn conv algorithm --------------------- // ------------------- cudnn conv algorithm ---------------------
cudnnConvolutionFwdAlgo_t algo; cudnnConvolutionFwdAlgo_t algo;
auto handle = ctx.cuda_device_context().cudnn_handle(); auto& dev_ctx = ctx.template device_context<platform::CUDADeviceContext>();
auto handle = dev_ctx.cudnn_handle();
PADDLE_ENFORCE(platform::dynload::cudnnGetConvolutionForwardAlgorithm( PADDLE_ENFORCE(platform::dynload::cudnnGetConvolutionForwardAlgorithm(
handle, cudnn_input_desc, cudnn_filter_desc, cudnn_conv_desc, handle, cudnn_input_desc, cudnn_filter_desc, cudnn_conv_desc,
...@@ -238,7 +239,8 @@ class CudnnConvGradOpKernel : public framework::OpKernel<T> { ...@@ -238,7 +239,8 @@ class CudnnConvGradOpKernel : public framework::OpKernel<T> {
workspace_size_limit = user_workspace_size * 1024 * 1024; workspace_size_limit = user_workspace_size * 1024 * 1024;
} }
auto handle = ctx.cuda_device_context().cudnn_handle(); auto& dev_ctx = ctx.template device_context<platform::CUDADeviceContext>();
auto handle = dev_ctx.cudnn_handle();
if (input_grad) { if (input_grad) {
PADDLE_ENFORCE( PADDLE_ENFORCE(
platform::dynload::cudnnGetConvolutionBackwardDataAlgorithm( platform::dynload::cudnnGetConvolutionBackwardDataAlgorithm(
...@@ -313,16 +315,16 @@ class CudnnConvGradOpKernel : public framework::OpKernel<T> { ...@@ -313,16 +315,16 @@ class CudnnConvGradOpKernel : public framework::OpKernel<T> {
} // namespace operators } // namespace operators
} // namespace paddle } // namespace paddle
REGISTER_OP_GPU_KERNEL(conv2d_cudnn, REGISTER_OP_CUDA_KERNEL(conv2d_cudnn,
paddle::operators::CudnnConvOpKernel<float>, paddle::operators::CudnnConvOpKernel<float>,
paddle::operators::CudnnConvOpKernel<double>); paddle::operators::CudnnConvOpKernel<double>);
REGISTER_OP_GPU_KERNEL(conv2d_cudnn_grad, REGISTER_OP_CUDA_KERNEL(conv2d_cudnn_grad,
paddle::operators::CudnnConvGradOpKernel<float>, paddle::operators::CudnnConvGradOpKernel<float>,
paddle::operators::CudnnConvGradOpKernel<double>); paddle::operators::CudnnConvGradOpKernel<double>);
REGISTER_OP_GPU_KERNEL(conv3d_cudnn, REGISTER_OP_CUDA_KERNEL(conv3d_cudnn,
paddle::operators::CudnnConvOpKernel<float>, paddle::operators::CudnnConvOpKernel<float>,
paddle::operators::CudnnConvOpKernel<double>); paddle::operators::CudnnConvOpKernel<double>);
REGISTER_OP_GPU_KERNEL(conv3d_cudnn_grad, REGISTER_OP_CUDA_KERNEL(conv3d_cudnn_grad,
paddle::operators::CudnnConvGradOpKernel<float>, paddle::operators::CudnnConvGradOpKernel<float>,
paddle::operators::CudnnConvGradOpKernel<double>); paddle::operators::CudnnConvGradOpKernel<double>);
paddle/operators/conv_op.cc
浏览文件 @ aa2f2fb4
...@@ -235,16 +235,18 @@ namespace ops = paddle::operators; ...@@ -235,16 +235,18 @@ namespace ops = paddle::operators;
REGISTER_OP(conv3d, ops::ConvOp, ops::Conv3DOpMaker, conv3d_grad, REGISTER_OP(conv3d, ops::ConvOp, ops::Conv3DOpMaker, conv3d_grad,
ops::ConvOpGrad); ops::ConvOpGrad);
REGISTER_OP_CPU_KERNEL(conv2d,
ops::GemmConvKernel<paddle::platform::CPUPlace, float>,
ops::GemmConvKernel<paddle::platform::CPUPlace, double>);
REGISTER_OP_CPU_KERNEL( REGISTER_OP_CPU_KERNEL(
conv2d_grad, ops::GemmConvGradKernel<paddle::platform::CPUPlace, float>, conv2d, ops::GemmConvKernel<paddle::platform::CPUDeviceContext, float>,
ops::GemmConvGradKernel<paddle::platform::CPUPlace, double>); ops::GemmConvKernel<paddle::platform::CPUDeviceContext, double>);
REGISTER_OP_CPU_KERNEL(
conv2d_grad,
ops::GemmConvGradKernel<paddle::platform::CPUDeviceContext, float>,
ops::GemmConvGradKernel<paddle::platform::CPUDeviceContext, double>);
REGISTER_OP_CPU_KERNEL(conv3d,
ops::GemmConvKernel<paddle::platform::CPUPlace, float>,
ops::GemmConvKernel<paddle::platform::CPUPlace, double>);
REGISTER_OP_CPU_KERNEL( REGISTER_OP_CPU_KERNEL(
conv3d_grad, ops::GemmConvGradKernel<paddle::platform::CPUPlace, float>, conv3d, ops::GemmConvKernel<paddle::platform::CPUDeviceContext, float>,
ops::GemmConvGradKernel<paddle::platform::CPUPlace, double>); ops::GemmConvKernel<paddle::platform::CPUDeviceContext, double>);
REGISTER_OP_CPU_KERNEL(
conv3d_grad,
ops::GemmConvGradKernel<paddle::platform::CPUDeviceContext, float>,
ops::GemmConvGradKernel<paddle::platform::CPUDeviceContext, double>);
paddle/operators/conv_op.cu.cc
浏览文件 @ aa2f2fb4
...@@ -16,16 +16,18 @@ ...@@ -16,16 +16,18 @@
namespace ops = paddle::operators; namespace ops = paddle::operators;
REGISTER_OP_GPU_KERNEL(conv2d, REGISTER_OP_CUDA_KERNEL(
ops::GemmConvKernel<paddle::platform::GPUPlace, float>, conv2d, ops::GemmConvKernel<paddle::platform::CUDADeviceContext, float>,
ops::GemmConvKernel<paddle::platform::GPUPlace, double>); ops::GemmConvKernel<paddle::platform::CUDADeviceContext, double>);
REGISTER_OP_GPU_KERNEL( REGISTER_OP_CUDA_KERNEL(
conv2d_grad, ops::GemmConvGradKernel<paddle::platform::GPUPlace, float>, conv2d_grad,
ops::GemmConvGradKernel<paddle::platform::GPUPlace, double>); ops::GemmConvGradKernel<paddle::platform::CUDADeviceContext, float>,
ops::GemmConvGradKernel<paddle::platform::CUDADeviceContext, double>);
REGISTER_OP_GPU_KERNEL(conv3d, REGISTER_OP_CUDA_KERNEL(
ops::GemmConvKernel<paddle::platform::GPUPlace, float>, conv3d, ops::GemmConvKernel<paddle::platform::CUDADeviceContext, float>,
ops::GemmConvKernel<paddle::platform::GPUPlace, double>); ops::GemmConvKernel<paddle::platform::CUDADeviceContext, double>);
REGISTER_OP_GPU_KERNEL( REGISTER_OP_CUDA_KERNEL(
conv3d_grad, ops::GemmConvGradKernel<paddle::platform::GPUPlace, float>, conv3d_grad,
ops::GemmConvGradKernel<paddle::platform::GPUPlace, double>); ops::GemmConvGradKernel<paddle::platform::CUDADeviceContext, float>,
ops::GemmConvGradKernel<paddle::platform::CUDADeviceContext, double>);
paddle/operators/conv_op.h
浏览文件 @ aa2f2fb4
...@@ -72,7 +72,7 @@ class ConvOpGrad : public framework::OperatorWithKernel { ...@@ -72,7 +72,7 @@ class ConvOpGrad : public framework::OperatorWithKernel {
void InferShape(framework::InferShapeContext* ctx) const override; void InferShape(framework::InferShapeContext* ctx) const override;
}; };
template <typename Place, typename T> template <typename DeviceContext, typename T>
class GemmConvKernel : public framework::OpKernel<T> { class GemmConvKernel : public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext& context) const override { void Compute(const framework::ExecutionContext& context) const override {
...@@ -141,9 +141,10 @@ class GemmConvKernel : public framework::OpKernel<T> { ...@@ -141,9 +141,10 @@ class GemmConvKernel : public framework::OpKernel<T> {
int in_step = static_cast<int>(input->dims()[1]) / groups; int in_step = static_cast<int>(input->dims()[1]) / groups;
int out_step = static_cast<int>(output->dims()[1]) / groups; int out_step = static_cast<int>(output->dims()[1]) / groups;
math::Vol2ColFunctor<Place, T> vol2col; math::Vol2ColFunctor<DeviceContext, T> vol2col;
math::Im2ColFunctor<math::ColFormat::kCFO, Place, T> im2col; math::Im2ColFunctor<math::ColFormat::kCFO, DeviceContext, T> im2col;
auto& dev_ctx = context.template device_context<DeviceContext>();
for (int i = 0; i < batch_size; i++) { for (int i = 0; i < batch_size; i++) {
Tensor in_batch = input->Slice(i, i + 1).Resize(input_shape); Tensor in_batch = input->Slice(i, i + 1).Resize(input_shape);
Tensor out_batch = output->Slice(i, i + 1).Resize(output_matrix_shape); Tensor out_batch = output->Slice(i, i + 1).Resize(output_matrix_shape);
...@@ -157,27 +158,26 @@ class GemmConvKernel : public framework::OpKernel<T> { ...@@ -157,27 +158,26 @@ class GemmConvKernel : public framework::OpKernel<T> {
col_matrix.Resize(col_matrix_shape); col_matrix.Resize(col_matrix_shape);
} else if (data_dim == 2U) { } else if (data_dim == 2U) {
// im2col // im2col
im2col(context.device_context(), in_slice, dilations, strides, im2col(dev_ctx, in_slice, dilations, strides,
std::vector<int>{paddings[0], paddings[1], paddings[0], std::vector<int>{paddings[0], paddings[1], paddings[0],
paddings[1]}, paddings[1]},
&col); &col);
} else if (data_dim == 3U) { } else if (data_dim == 3U) {
// vol2col // vol2col
vol2col(context.device_context(), in_slice, dilations, strides, vol2col(dev_ctx, in_slice, dilations, strides, paddings, &col);
paddings, &col);
} }
// gemm // gemm
Tensor out_slice = out_batch.Slice(g * out_step, (g + 1) * out_step); Tensor out_slice = out_batch.Slice(g * out_step, (g + 1) * out_step);
Tensor filter_slice = filter.Slice(g * out_step, (g + 1) * out_step); Tensor filter_slice = filter.Slice(g * out_step, (g + 1) * out_step);
math::matmul<Place, T>(context.device_context(), filter_slice, false, math::matmul<DeviceContext, T>(dev_ctx, filter_slice, false, col_matrix,
col_matrix, false, T(1.0), &out_slice, T(0.0)); false, T(1.0), &out_slice, T(0.0));
} }
} }
} }
}; };
template <typename Place, typename T> template <typename DeviceContext, typename T>
class GemmConvGradKernel : public framework::OpKernel<T> { class GemmConvGradKernel : public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext& context) const override { void Compute(const framework::ExecutionContext& context) const override {
...@@ -256,17 +256,19 @@ class GemmConvGradKernel : public framework::OpKernel<T> { ...@@ -256,17 +256,19 @@ class GemmConvGradKernel : public framework::OpKernel<T> {
col_matrix.Resize(col_matrix_shape); col_matrix.Resize(col_matrix_shape);
} }
math::SetConstant<Place, T> set_zero; math::SetConstant<DeviceContext, T> set_zero;
auto& dev_ctx = context.template device_context<DeviceContext>();
if (input_grad) { if (input_grad) {
input_grad->mutable_data<T>(context.GetPlace()); input_grad->mutable_data<T>(context.GetPlace());
// if is_expand is false, the operation of set_zero is unnecessary, // if is_expand is false, the operation of set_zero is unnecessary,
// because math::matmul will reset input_grad. // because math::matmul will reset input_grad.
if (is_expand) { if (is_expand) {
set_zero(context.device_context(), input_grad, static_cast<T>(0)); set_zero(dev_ctx, input_grad, static_cast<T>(0));
} }
math::Col2VolFunctor<Place, T> col2vol; math::Col2VolFunctor<DeviceContext, T> col2vol;
math::Col2ImFunctor<math::ColFormat::kCFO, Place, T> col2im; math::Col2ImFunctor<math::ColFormat::kCFO, DeviceContext, T> col2im;
for (int i = 0; i < batch_size; i++) { for (int i = 0; i < batch_size; i++) {
Tensor out_grad_batch = Tensor out_grad_batch =
...@@ -285,18 +287,17 @@ class GemmConvGradKernel : public framework::OpKernel<T> { ...@@ -285,18 +287,17 @@ class GemmConvGradKernel : public framework::OpKernel<T> {
col_matrix.ShareDataWith(in_grad_slice); col_matrix.ShareDataWith(in_grad_slice);
col_matrix.Resize(col_matrix_shape); col_matrix.Resize(col_matrix_shape);
} }
math::matmul<Place, T>(context.device_context(), filter_slice, true, math::matmul<DeviceContext, T>(dev_ctx, filter_slice, true,
out_grad_slice, false, T(1.0), &col_matrix, out_grad_slice, false, T(1.0),
T(0.0)); &col_matrix, T(0.0));
if (is_expand && data_dim == 2U) { if (is_expand && data_dim == 2U) {
col2im(context.device_context(), col, dilations, strides, col2im(dev_ctx, col, dilations, strides,
std::vector<int>{paddings[0], paddings[1], paddings[0], std::vector<int>{paddings[0], paddings[1], paddings[0],
paddings[1]}, paddings[1]},
&in_grad_slice); &in_grad_slice);
} else if (is_expand && data_dim == 3U) { } else if (is_expand && data_dim == 3U) {
col2vol(context.device_context(), col, dilations, strides, paddings, col2vol(dev_ctx, col, dilations, strides, paddings, &in_grad_slice);
&in_grad_slice);
} }
} }
} }
...@@ -306,9 +307,9 @@ class GemmConvGradKernel : public framework::OpKernel<T> { ...@@ -306,9 +307,9 @@ class GemmConvGradKernel : public framework::OpKernel<T> {
filter_grad->mutable_data<T>(context.GetPlace()); filter_grad->mutable_data<T>(context.GetPlace());
Tensor filter_grad_ = *filter_grad; Tensor filter_grad_ = *filter_grad;
filter_grad_.Resize(filter_matrix_shape); filter_grad_.Resize(filter_matrix_shape);
set_zero(context.device_context(), filter_grad, static_cast<T>(0)); set_zero(dev_ctx, filter_grad, static_cast<T>(0));
math::Im2ColFunctor<math::ColFormat::kCFO, Place, T> im2col; math::Im2ColFunctor<math::ColFormat::kCFO, DeviceContext, T> im2col;
math::Vol2ColFunctor<Place, T> vol2col; math::Vol2ColFunctor<DeviceContext, T> vol2col;
for (int i = 0; i < batch_size; i++) { for (int i = 0; i < batch_size; i++) {
Tensor out_grad_batch = Tensor out_grad_batch =
output_grad->Slice(i, i + 1).Resize(output_matrix_shape); output_grad->Slice(i, i + 1).Resize(output_matrix_shape);
...@@ -324,21 +325,20 @@ class GemmConvGradKernel : public framework::OpKernel<T> { ...@@ -324,21 +325,20 @@ class GemmConvGradKernel : public framework::OpKernel<T> {
col_matrix.ShareDataWith(col); col_matrix.ShareDataWith(col);
col_matrix.Resize(col_matrix_shape); col_matrix.Resize(col_matrix_shape);
} else if (data_dim == 2U) { } else if (data_dim == 2U) {
im2col(context.device_context(), in_slice, dilations, strides, im2col(dev_ctx, in_slice, dilations, strides,
std::vector<int>{paddings[0], paddings[1], paddings[0], std::vector<int>{paddings[0], paddings[1], paddings[0],
paddings[1]}, paddings[1]},
&col); &col);
} else if (data_dim == 3U) { } else if (data_dim == 3U) {
vol2col(context.device_context(), in_slice, dilations, strides, vol2col(dev_ctx, in_slice, dilations, strides, paddings, &col);
paddings, &col);
} }
// gemm // gemm
Tensor filter_grad_slice = Tensor filter_grad_slice =
filter_grad_.Slice(g * out_step, (g + 1) * out_step); filter_grad_.Slice(g * out_step, (g + 1) * out_step);
math::matmul<Place, T>(context.device_context(), out_grad_slice, math::matmul<DeviceContext, T>(dev_ctx, out_grad_slice, false,
false, col_matrix, true, T(1.0), col_matrix, true, T(1.0),
&filter_grad_slice, T(1.0)); &filter_grad_slice, T(1.0));
} }
} }
} }
......
paddle/operators/conv_shift_op.cu
浏览文件 @ aa2f2fb4
...@@ -111,7 +111,8 @@ __global__ void ConvShiftDy(const T *x, const T *dout, int x_width, int y_width, ...@@ -111,7 +111,8 @@ __global__ void ConvShiftDy(const T *x, const T *dout, int x_width, int y_width,
} // namespace } // namespace
template <typename T> template <typename T>
class ConvShiftKernel<platform::GPUPlace, T> : public framework::OpKernel<T> { class ConvShiftKernel<platform::CUDADeviceContext, T>
: public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext &context) const override { void Compute(const framework::ExecutionContext &context) const override {
const Tensor *X = context.Input<Tensor>("X"); const Tensor *X = context.Input<Tensor>("X");
...@@ -132,7 +133,8 @@ class ConvShiftKernel<platform::GPUPlace, T> : public framework::OpKernel<T> { ...@@ -132,7 +133,8 @@ class ConvShiftKernel<platform::GPUPlace, T> : public framework::OpKernel<T> {
dim3 grid_dim(num_x_blocks, batch_size); dim3 grid_dim(num_x_blocks, batch_size);
auto stream = context.cuda_device_context().stream(); auto stream =
context.template device_context<platform::CUDADeviceContext>().stream();
ConvShiftForward<T><<<grid_dim, x_per_block, mem_per_block, stream>>>( ConvShiftForward<T><<<grid_dim, x_per_block, mem_per_block, stream>>>(
x_data, y_data, x_width, y_width, y_half_width, batch_size, out_data); x_data, y_data, x_width, y_width, y_half_width, batch_size, out_data);
...@@ -140,7 +142,7 @@ class ConvShiftKernel<platform::GPUPlace, T> : public framework::OpKernel<T> { ...@@ -140,7 +142,7 @@ class ConvShiftKernel<platform::GPUPlace, T> : public framework::OpKernel<T> {
}; };
template <typename T> template <typename T>
class ConvShiftGradKernel<platform::GPUPlace, T> class ConvShiftGradKernel<platform::CUDADeviceContext, T>
: public framework::OpKernel<T> { : public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext &context) const override { void Compute(const framework::ExecutionContext &context) const override {
...@@ -159,8 +161,9 @@ class ConvShiftGradKernel<platform::GPUPlace, T> ...@@ -159,8 +161,9 @@ class ConvShiftGradKernel<platform::GPUPlace, T>
int y_width = Y->dims()[1]; int y_width = Y->dims()[1];
int y_half_width = (y_width - 1) / 2; int y_half_width = (y_width - 1) / 2;
auto &device_ctx = context.cuda_device_context(); auto &device_ctx =
math::SetConstant<platform::GPUPlace, T> zero; context.template device_context<platform::CUDADeviceContext>();
math::SetConstant<platform::CUDADeviceContext, T> zero;
const int x_per_block = 256; const int x_per_block = 256;
int num_x_blocks = DivUp(x_width, x_per_block); int num_x_blocks = DivUp(x_width, x_per_block);
...@@ -186,8 +189,9 @@ class ConvShiftGradKernel<platform::GPUPlace, T> ...@@ -186,8 +189,9 @@ class ConvShiftGradKernel<platform::GPUPlace, T>
} // namespace paddle } // namespace paddle
namespace ops = paddle::operators; namespace ops = paddle::operators;
REGISTER_OP_GPU_KERNEL(conv_shift, REGISTER_OP_CUDA_KERNEL(
ops::ConvShiftKernel<paddle::platform::GPUPlace, float>); conv_shift,
REGISTER_OP_GPU_KERNEL( ops::ConvShiftKernel<paddle::platform::CUDADeviceContext, float>);
REGISTER_OP_CUDA_KERNEL(
conv_shift_grad, conv_shift_grad,
ops::ConvShiftGradKernel<paddle::platform::GPUPlace, float>); ops::ConvShiftGradKernel<paddle::platform::CUDADeviceContext, float>);
paddle/operators/conv_shift_op.h
浏览文件 @ aa2f2fb4
...@@ -18,13 +18,13 @@ ...@@ -18,13 +18,13 @@
namespace paddle { namespace paddle {
namespace operators { namespace operators {
template <typename Place, typename T> template <typename DeviceContext, typename T>
class ConvShiftKernel : public framework::OpKernel<T> { class ConvShiftKernel : public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext &context) const override; void Compute(const framework::ExecutionContext &context) const override;
}; };
template <typename Place, typename T> template <typename DeviceContext, typename T>
class ConvShiftGradKernel : public framework::OpKernel<T> { class ConvShiftGradKernel : public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext &context) const override; void Compute(const framework::ExecutionContext &context) const override;
......
paddle/operators/conv_transpose_cudnn_op.cc
浏览文件 @ aa2f2fb4
...@@ -61,12 +61,13 @@ REGISTER_OP(conv2d_transpose_cudnn, ops::ConvTransposeOp, ...@@ -61,12 +61,13 @@ REGISTER_OP(conv2d_transpose_cudnn, ops::ConvTransposeOp,
REGISTER_OP_CPU_KERNEL( REGISTER_OP_CPU_KERNEL(
conv2d_transpose_cudnn, conv2d_transpose_cudnn,
ops::GemmConvTransposeKernel<paddle::platform::CPUPlace, float>, ops::GemmConvTransposeKernel<paddle::platform::CPUDeviceContext, float>,
ops::GemmConvTransposeKernel<paddle::platform::CPUPlace, double>); ops::GemmConvTransposeKernel<paddle::platform::CPUDeviceContext, double>);
REGISTER_OP_CPU_KERNEL( REGISTER_OP_CPU_KERNEL(
conv2d_transpose_cudnn_grad, conv2d_transpose_cudnn_grad,
ops::GemmConvTransposeGradKernel<paddle::platform::CPUPlace, float>, ops::GemmConvTransposeGradKernel<paddle::platform::CPUDeviceContext, float>,
ops::GemmConvTransposeGradKernel<paddle::platform::CPUPlace, double>); ops::GemmConvTransposeGradKernel<paddle::platform::CPUDeviceContext,
double>);
REGISTER_OP(conv3d_transpose_cudnn, ops::ConvTransposeOp, REGISTER_OP(conv3d_transpose_cudnn, ops::ConvTransposeOp,
ops::CudnnConv3DTransposeOpMaker, conv3d_transpose_cudnn_grad, ops::CudnnConv3DTransposeOpMaker, conv3d_transpose_cudnn_grad,
...@@ -74,9 +75,10 @@ REGISTER_OP(conv3d_transpose_cudnn, ops::ConvTransposeOp, ...@@ -74,9 +75,10 @@ REGISTER_OP(conv3d_transpose_cudnn, ops::ConvTransposeOp,
REGISTER_OP_CPU_KERNEL( REGISTER_OP_CPU_KERNEL(
conv3d_transpose_cudnn, conv3d_transpose_cudnn,
ops::GemmConvTransposeKernel<paddle::platform::CPUPlace, float>, ops::GemmConvTransposeKernel<paddle::platform::CPUDeviceContext, float>,
ops::GemmConvTransposeKernel<paddle::platform::CPUPlace, double>); ops::GemmConvTransposeKernel<paddle::platform::CPUDeviceContext, double>);
REGISTER_OP_CPU_KERNEL( REGISTER_OP_CPU_KERNEL(
conv3d_transpose_cudnn_grad, conv3d_transpose_cudnn_grad,
ops::GemmConvTransposeGradKernel<paddle::platform::CPUPlace, float>, ops::GemmConvTransposeGradKernel<paddle::platform::CPUDeviceContext, float>,
ops::GemmConvTransposeGradKernel<paddle::platform::CPUPlace, double>); ops::GemmConvTransposeGradKernel<paddle::platform::CPUDeviceContext,
double>);
paddle/operators/conv_transpose_cudnn_op.cu.cc
浏览文件 @ aa2f2fb4
...@@ -83,7 +83,8 @@ class CudnnConvTransposeOpKernel : public framework::OpKernel<T> { ...@@ -83,7 +83,8 @@ class CudnnConvTransposeOpKernel : public framework::OpKernel<T> {
} }
// ------------------- cudnn conv algorithm --------------------- // ------------------- cudnn conv algorithm ---------------------
cudnnConvolutionBwdDataAlgo_t algo; cudnnConvolutionBwdDataAlgo_t algo;
auto handle = ctx.cuda_device_context().cudnn_handle(); auto& dev_ctx = ctx.template device_context<platform::CUDADeviceContext>();
auto handle = dev_ctx.cudnn_handle();
// Get the algorithm // Get the algorithm
PADDLE_ENFORCE(platform::dynload::cudnnGetConvolutionBackwardDataAlgorithm( PADDLE_ENFORCE(platform::dynload::cudnnGetConvolutionBackwardDataAlgorithm(
handle, cudnn_filter_desc, cudnn_input_desc, cudnn_conv_desc, handle, cudnn_filter_desc, cudnn_input_desc, cudnn_conv_desc,
...@@ -165,7 +166,8 @@ class CudnnConvTransposeGradOpKernel : public framework::OpKernel<T> { ...@@ -165,7 +166,8 @@ class CudnnConvTransposeGradOpKernel : public framework::OpKernel<T> {
workspace_size_limit = user_workspace_size * 1024 * 1024; workspace_size_limit = user_workspace_size * 1024 * 1024;
} }
auto handle = ctx.cuda_device_context().cudnn_handle(); auto& dev_ctx = ctx.template device_context<platform::CUDADeviceContext>();
auto handle = dev_ctx.cudnn_handle();
if (input_grad) { if (input_grad) {
// choose backward algorithm for data // choose backward algorithm for data
PADDLE_ENFORCE(platform::dynload::cudnnGetConvolutionForwardAlgorithm( PADDLE_ENFORCE(platform::dynload::cudnnGetConvolutionForwardAlgorithm(
...@@ -234,16 +236,16 @@ class CudnnConvTransposeGradOpKernel : public framework::OpKernel<T> { ...@@ -234,16 +236,16 @@ class CudnnConvTransposeGradOpKernel : public framework::OpKernel<T> {
namespace ops = paddle::operators; namespace ops = paddle::operators;
REGISTER_OP_GPU_KERNEL(conv2d_transpose_cudnn, REGISTER_OP_CUDA_KERNEL(conv2d_transpose_cudnn,
ops::CudnnConvTransposeOpKernel<float>, ops::CudnnConvTransposeOpKernel<float>,
ops::CudnnConvTransposeOpKernel<double>); ops::CudnnConvTransposeOpKernel<double>);
REGISTER_OP_GPU_KERNEL(conv2d_transpose_cudnn_grad, REGISTER_OP_CUDA_KERNEL(conv2d_transpose_cudnn_grad,
ops::CudnnConvTransposeGradOpKernel<float>, ops::CudnnConvTransposeGradOpKernel<float>,
ops::CudnnConvTransposeGradOpKernel<double>); ops::CudnnConvTransposeGradOpKernel<double>);
REGISTER_OP_GPU_KERNEL(conv3d_transpose_cudnn, REGISTER_OP_CUDA_KERNEL(conv3d_transpose_cudnn,
ops::CudnnConvTransposeOpKernel<float>, ops::CudnnConvTransposeOpKernel<float>,
ops::CudnnConvTransposeOpKernel<double>); ops::CudnnConvTransposeOpKernel<double>);
REGISTER_OP_GPU_KERNEL(conv3d_transpose_cudnn_grad, REGISTER_OP_CUDA_KERNEL(conv3d_transpose_cudnn_grad,
ops::CudnnConvTransposeGradOpKernel<float>, ops::CudnnConvTransposeGradOpKernel<float>,
ops::CudnnConvTransposeGradOpKernel<double>); ops::CudnnConvTransposeGradOpKernel<double>);
paddle/operators/conv_transpose_op.cc
浏览文件 @ aa2f2fb4
...@@ -197,21 +197,23 @@ REGISTER_OP(conv2d_transpose, ops::ConvTransposeOp, ops::Conv2DTransposeOpMaker, ...@@ -197,21 +197,23 @@ REGISTER_OP(conv2d_transpose, ops::ConvTransposeOp, ops::Conv2DTransposeOpMaker,
REGISTER_OP_CPU_KERNEL( REGISTER_OP_CPU_KERNEL(
conv2d_transpose, conv2d_transpose,
ops::GemmConvTransposeKernel<paddle::platform::CPUPlace, float>, ops::GemmConvTransposeKernel<paddle::platform::CPUDeviceContext, float>,
ops::GemmConvTransposeKernel<paddle::platform::CPUPlace, double>); ops::GemmConvTransposeKernel<paddle::platform::CPUDeviceContext, double>);
REGISTER_OP_CPU_KERNEL( REGISTER_OP_CPU_KERNEL(
conv2d_transpose_grad, conv2d_transpose_grad,
ops::GemmConvTransposeGradKernel<paddle::platform::CPUPlace, float>, ops::GemmConvTransposeGradKernel<paddle::platform::CPUDeviceContext, float>,
ops::GemmConvTransposeGradKernel<paddle::platform::CPUPlace, double>); ops::GemmConvTransposeGradKernel<paddle::platform::CPUDeviceContext,
double>);
REGISTER_OP(conv3d_transpose, ops::ConvTransposeOp, ops::Conv3DTransposeOpMaker, REGISTER_OP(conv3d_transpose, ops::ConvTransposeOp, ops::Conv3DTransposeOpMaker,
conv3d_transpose_grad, ops::ConvTransposeOpGrad); conv3d_transpose_grad, ops::ConvTransposeOpGrad);
REGISTER_OP_CPU_KERNEL( REGISTER_OP_CPU_KERNEL(
conv3d_transpose, conv3d_transpose,
ops::GemmConvTransposeKernel<paddle::platform::CPUPlace, float>, ops::GemmConvTransposeKernel<paddle::platform::CPUDeviceContext, float>,
ops::GemmConvTransposeKernel<paddle::platform::CPUPlace, double>); ops::GemmConvTransposeKernel<paddle::platform::CPUDeviceContext, double>);
REGISTER_OP_CPU_KERNEL( REGISTER_OP_CPU_KERNEL(
conv3d_transpose_grad, conv3d_transpose_grad,
ops::GemmConvTransposeGradKernel<paddle::platform::CPUPlace, float>, ops::GemmConvTransposeGradKernel<paddle::platform::CPUDeviceContext, float>,
ops::GemmConvTransposeGradKernel<paddle::platform::CPUPlace, double>); ops::GemmConvTransposeGradKernel<paddle::platform::CPUDeviceContext,
double>);
paddle/operators/conv_transpose_op.cu.cc
浏览文件 @ aa2f2fb4
...@@ -16,20 +16,24 @@ ...@@ -16,20 +16,24 @@
namespace ops = paddle::operators; namespace ops = paddle::operators;
REGISTER_OP_GPU_KERNEL( REGISTER_OP_CUDA_KERNEL(
conv2d_transpose, conv2d_transpose,
ops::GemmConvTransposeKernel<paddle::platform::GPUPlace, float>, ops::GemmConvTransposeKernel<paddle::platform::CUDADeviceContext, float>,
ops::GemmConvTransposeKernel<paddle::platform::GPUPlace, double>); ops::GemmConvTransposeKernel<paddle::platform::CUDADeviceContext, double>);
REGISTER_OP_GPU_KERNEL( REGISTER_OP_CUDA_KERNEL(
conv2d_transpose_grad, conv2d_transpose_grad,
ops::GemmConvTransposeGradKernel<paddle::platform::GPUPlace, float>, ops::GemmConvTransposeGradKernel<paddle::platform::CUDADeviceContext,
ops::GemmConvTransposeGradKernel<paddle::platform::GPUPlace, double>); float>,
ops::GemmConvTransposeGradKernel<paddle::platform::CUDADeviceContext,
double>);
REGISTER_OP_GPU_KERNEL( REGISTER_OP_CUDA_KERNEL(
conv3d_transpose, conv3d_transpose,
ops::GemmConvTransposeKernel<paddle::platform::GPUPlace, float>, ops::GemmConvTransposeKernel<paddle::platform::CUDADeviceContext, float>,
ops::GemmConvTransposeKernel<paddle::platform::GPUPlace, double>); ops::GemmConvTransposeKernel<paddle::platform::CUDADeviceContext, double>);
REGISTER_OP_GPU_KERNEL( REGISTER_OP_CUDA_KERNEL(
conv3d_transpose_grad, conv3d_transpose_grad,
ops::GemmConvTransposeGradKernel<paddle::platform::GPUPlace, float>, ops::GemmConvTransposeGradKernel<paddle::platform::CUDADeviceContext,
ops::GemmConvTransposeGradKernel<paddle::platform::GPUPlace, double>); float>,
ops::GemmConvTransposeGradKernel<paddle::platform::CUDADeviceContext,
double>);
paddle/operators/conv_transpose_op.h
浏览文件 @ aa2f2fb4
...@@ -52,7 +52,7 @@ class ConvTransposeOpGrad : public framework::OperatorWithKernel { ...@@ -52,7 +52,7 @@ class ConvTransposeOpGrad : public framework::OperatorWithKernel {
void InferShape(framework::InferShapeContext* ctx) const override; void InferShape(framework::InferShapeContext* ctx) const override;
}; };
template <typename Place, typename T> template <typename DeviceContext, typename T>
class GemmConvTransposeKernel : public framework::OpKernel<T> { class GemmConvTransposeKernel : public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext& context) const override { void Compute(const framework::ExecutionContext& context) const override {
...@@ -109,11 +109,12 @@ class GemmConvTransposeKernel : public framework::OpKernel<T> { ...@@ -109,11 +109,12 @@ class GemmConvTransposeKernel : public framework::OpKernel<T> {
filter.Resize(filter_matrix_shape); filter.Resize(filter_matrix_shape);
output->mutable_data<T>(context.GetPlace()); output->mutable_data<T>(context.GetPlace());
math::SetConstant<Place, T> set_zero; math::SetConstant<DeviceContext, T> set_zero;
set_zero(context.device_context(), output, static_cast<T>(0)); auto& dev_ctx = context.template device_context<DeviceContext>();
set_zero(dev_ctx, output, static_cast<T>(0));
math::Col2ImFunctor<math::ColFormat::kCFO, Place, T> col2im; math::Col2ImFunctor<math::ColFormat::kCFO, DeviceContext, T> col2im;
math::Col2VolFunctor<Place, T> col2vol; math::Col2VolFunctor<DeviceContext, T> col2vol;
std::vector<int> dilations({1, 1, 1}); std::vector<int> dilations({1, 1, 1});
// convolution transpose: gemm + col2im or col2vol (similar to conv-backward // convolution transpose: gemm + col2im or col2vol (similar to conv-backward
...@@ -127,29 +128,27 @@ class GemmConvTransposeKernel : public framework::OpKernel<T> { ...@@ -127,29 +128,27 @@ class GemmConvTransposeKernel : public framework::OpKernel<T> {
// col_matrix = filter * input_batch // col_matrix = filter * input_batch
// of shape (c * k_h * k_w, h * w) or (c * k_d * k_h * k_w, d * h * w) // of shape (c * k_h * k_w, h * w) or (c * k_d * k_h * k_w, d * h * w)
math::matmul<Place, T>(context.device_context(), filter, true, math::matmul<DeviceContext, T>(dev_ctx, filter, true, input_batch, false,
input_batch, false, static_cast<T>(1.0), static_cast<T>(1.0), &col_matrix,
&col_matrix, static_cast<T>(0.0)); static_cast<T>(0.0));
if (data_dim == 2U) { if (data_dim == 2U) {
// col2im: col_matrix -> dy // col2im: col_matrix -> dy
// from (c * k_h * k_w, h * w) to (c, o_h, o_w) // from (c * k_h * k_w, h * w) to (c, o_h, o_w)
col2im(context.device_context(), col, col2im(dev_ctx, col, std::vector<int>{dilations[0], dilations[1]},
std::vector<int>{dilations[0], dilations[1]}, strides, strides, std::vector<int>{paddings[0], paddings[1], paddings[0],
std::vector<int>{paddings[0], paddings[1], paddings[0], paddings[1]},
paddings[1]},
&output_batch); &output_batch);
} else if (data_dim == 3U) { } else if (data_dim == 3U) {
// col2vol: col_matrix -> dy // col2vol: col_matrix -> dy
// from (c * k_d * k_h * k_w, d * h * w) to (c, o_d, o_h, o_w) // from (c * k_d * k_h * k_w, d * h * w) to (c, o_d, o_h, o_w)
col2vol(context.device_context(), col, dilations, strides, paddings, col2vol(dev_ctx, col, dilations, strides, paddings, &output_batch);
&output_batch);
} }
} }
} }
}; };
template <typename Place, typename T> template <typename DeviceContext, typename T>
class GemmConvTransposeGradKernel : public framework::OpKernel<T> { class GemmConvTransposeGradKernel : public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext& context) const override { void Compute(const framework::ExecutionContext& context) const override {
...@@ -206,6 +205,7 @@ class GemmConvTransposeGradKernel : public framework::OpKernel<T> { ...@@ -206,6 +205,7 @@ class GemmConvTransposeGradKernel : public framework::OpKernel<T> {
// convolution transpose grad on input: // convolution transpose grad on input:
// im2col + gemm (similar to conv-forward) // im2col + gemm (similar to conv-forward)
// input need to compute gradient // input need to compute gradient
auto& dev_ctx = context.template device_context<DeviceContext>();
if (input_grad || filter_grad) { if (input_grad || filter_grad) {
Tensor col; Tensor col;
col.mutable_data<T>(col_shape, context.GetPlace()); col.mutable_data<T>(col_shape, context.GetPlace());
...@@ -217,10 +217,10 @@ class GemmConvTransposeGradKernel : public framework::OpKernel<T> { ...@@ -217,10 +217,10 @@ class GemmConvTransposeGradKernel : public framework::OpKernel<T> {
col_matrix.Resize(col_matrix_shape); col_matrix.Resize(col_matrix_shape);
Tensor filter_grad_; Tensor filter_grad_;
math::SetConstant<Place, T> set_zero; math::SetConstant<DeviceContext, T> set_zero;
math::Im2ColFunctor<math::ColFormat::kCFO, Place, T> im2col; math::Im2ColFunctor<math::ColFormat::kCFO, DeviceContext, T> im2col;
math::Vol2ColFunctor<Place, T> vol2col; math::Vol2ColFunctor<DeviceContext, T> vol2col;
std::vector<int> dilations({1, 1, 1}); std::vector<int> dilations({1, 1, 1});
if (input_grad) { if (input_grad) {
...@@ -229,7 +229,7 @@ class GemmConvTransposeGradKernel : public framework::OpKernel<T> { ...@@ -229,7 +229,7 @@ class GemmConvTransposeGradKernel : public framework::OpKernel<T> {
} }
if (filter_grad) { // filter size (m, c, k_h, k_w) if (filter_grad) { // filter size (m, c, k_h, k_w)
filter_grad->mutable_data<T>(context.GetPlace()); filter_grad->mutable_data<T>(context.GetPlace());
set_zero(context.device_context(), filter_grad, static_cast<T>(0)); set_zero(dev_ctx, filter_grad, static_cast<T>(0));
filter_grad_ = *filter_grad; filter_grad_ = *filter_grad;
filter_grad_.Resize(filter_matrix_shape); filter_grad_.Resize(filter_matrix_shape);
} }
...@@ -242,7 +242,7 @@ class GemmConvTransposeGradKernel : public framework::OpKernel<T> { ...@@ -242,7 +242,7 @@ class GemmConvTransposeGradKernel : public framework::OpKernel<T> {
if (data_dim == 2U) { if (data_dim == 2U) {
// im2col: dy -> col matrix // im2col: dy -> col matrix
// from (c, o_h, o_w) to (c * k_h * k_w, h * w) // from (c, o_h, o_w) to (c * k_h * k_w, h * w)
im2col(context.device_context(), output_grad_batch, im2col(dev_ctx, output_grad_batch,
std::vector<int>{dilations[0], dilations[1]}, strides, std::vector<int>{dilations[0], dilations[1]}, strides,
std::vector<int>{paddings[0], paddings[1], paddings[0], std::vector<int>{paddings[0], paddings[1], paddings[0],
paddings[1]}, paddings[1]},
...@@ -250,8 +250,8 @@ class GemmConvTransposeGradKernel : public framework::OpKernel<T> { ...@@ -250,8 +250,8 @@ class GemmConvTransposeGradKernel : public framework::OpKernel<T> {
} else if (data_dim == 3U) { } else if (data_dim == 3U) {
// vol2col: dy -> col_matrix // vol2col: dy -> col_matrix
// from (c, o_d, o_h, o_w) to (c * k_d * k_h * k_w, d * h * w) // from (c, o_d, o_h, o_w) to (c * k_d * k_h * k_w, d * h * w)
vol2col(context.device_context(), output_grad_batch, dilations, vol2col(dev_ctx, output_grad_batch, dilations, strides, paddings,
strides, paddings, &col); &col);
} }
if (input_grad) { if (input_grad) {
...@@ -263,9 +263,9 @@ class GemmConvTransposeGradKernel : public framework::OpKernel<T> { ...@@ -263,9 +263,9 @@ class GemmConvTransposeGradKernel : public framework::OpKernel<T> {
// or // or
// (m, c * k_d * k_h * k_w) * (c * k_d * k_h * k_w, d * h * w) -> (m, // (m, c * k_d * k_h * k_w) * (c * k_d * k_h * k_w, d * h * w) -> (m,
// d, h, w) // d, h, w)
math::matmul<Place, T>(context.device_context(), filter, false, math::matmul<DeviceContext, T>(
col_matrix, false, static_cast<T>(1.0), dev_ctx, filter, false, col_matrix, false, static_cast<T>(1.0),
&input_grad_batch, static_cast<T>(0.0)); &input_grad_batch, static_cast<T>(0.0));
} }
if (filter_grad) { if (filter_grad) {
// input batch // input batch
...@@ -275,9 +275,9 @@ class GemmConvTransposeGradKernel : public framework::OpKernel<T> { ...@@ -275,9 +275,9 @@ class GemmConvTransposeGradKernel : public framework::OpKernel<T> {
// or // or
// (m, d * h * w) * (d * h * w, c * k_d * k_h * k_w) -> (m, c * k_d * // (m, d * h * w) * (d * h * w, c * k_d * k_h * k_w) -> (m, c * k_d *
// k_h * k_w) // k_h * k_w)
math::matmul<Place, T>(context.device_context(), in_batch, false, math::matmul<DeviceContext, T>(dev_ctx, in_batch, false, col_matrix,
col_matrix, true, static_cast<T>(1.0), true, static_cast<T>(1.0),
&filter_grad_, static_cast<T>(1.0)); &filter_grad_, static_cast<T>(1.0));
} }
} }
} }
......
paddle/operators/cos_sim_op.cc
浏览文件 @ aa2f2fb4
...@@ -155,7 +155,8 @@ class CosSimOpGrad : public framework::OperatorWithKernel { ...@@ -155,7 +155,8 @@ class CosSimOpGrad : public framework::OperatorWithKernel {
namespace ops = paddle::operators; namespace ops = paddle::operators;
REGISTER_OP(cos_sim, ops::CosSimOp, ops::CosSimOpMaker, cos_sim_grad, REGISTER_OP(cos_sim, ops::CosSimOp, ops::CosSimOpMaker, cos_sim_grad,
ops::CosSimOpGrad); ops::CosSimOpGrad);
REGISTER_OP_CPU_KERNEL(cos_sim,
ops::CosSimKernel<paddle::platform::CPUPlace, float>);
REGISTER_OP_CPU_KERNEL( REGISTER_OP_CPU_KERNEL(
cos_sim_grad, ops::CosSimGradKernel<paddle::platform::CPUPlace, float>); cos_sim, ops::CosSimKernel<paddle::platform::CPUDeviceContext, float>);
REGISTER_OP_CPU_KERNEL(
cos_sim_grad,
ops::CosSimGradKernel<paddle::platform::CPUDeviceContext, float>);
paddle/operators/cos_sim_op.cu
浏览文件 @ aa2f2fb4
...@@ -16,7 +16,8 @@ ...@@ -16,7 +16,8 @@
#include "paddle/operators/cos_sim_op.h" #include "paddle/operators/cos_sim_op.h"
namespace ops = paddle::operators; namespace ops = paddle::operators;
REGISTER_OP_GPU_KERNEL(cos_sim, REGISTER_OP_CUDA_KERNEL(
ops::CosSimKernel<paddle::platform::GPUPlace, float>); cos_sim, ops::CosSimKernel<paddle::platform::CUDADeviceContext, float>);
REGISTER_OP_GPU_KERNEL( REGISTER_OP_CUDA_KERNEL(
cos_sim_grad, ops::CosSimGradKernel<paddle::platform::GPUPlace, float>); cos_sim_grad,
ops::CosSimGradKernel<paddle::platform::CUDADeviceContext, float>);
paddle/operators/cos_sim_op.h
浏览文件 @ aa2f2fb4
...@@ -27,7 +27,7 @@ template <typename T, int MajorType = Eigen::RowMajor, ...@@ -27,7 +27,7 @@ template <typename T, int MajorType = Eigen::RowMajor,
typename IndexType = Eigen::DenseIndex> typename IndexType = Eigen::DenseIndex>
using EigenVector = framework::EigenVector<T, MajorType, IndexType>; using EigenVector = framework::EigenVector<T, MajorType, IndexType>;
template <typename Place, typename T> template <typename DeviceContext, typename T>
class CosSimKernel : public framework::OpKernel<T> { class CosSimKernel : public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext& context) const override { void Compute(const framework::ExecutionContext& context) const override {
...@@ -51,7 +51,8 @@ class CosSimKernel : public framework::OpKernel<T> { ...@@ -51,7 +51,8 @@ class CosSimKernel : public framework::OpKernel<T> {
auto y_norm = EigenVector<T>::Flatten(*out_y_norm); auto y_norm = EigenVector<T>::Flatten(*out_y_norm);
// compute // compute
auto place = context.GetEigenDevice<Place>(); auto& place =
*context.template device_context<DeviceContext>().eigen_device();
auto row_along = Eigen::array<int, 1>({{1}}); auto row_along = Eigen::array<int, 1>({{1}});
x_norm.device(place) = x.square().sum(row_along).sqrt(); x_norm.device(place) = x.square().sum(row_along).sqrt();
y_norm.device(place) = y.square().sum(row_along).sqrt(); y_norm.device(place) = y.square().sum(row_along).sqrt();
...@@ -66,7 +67,7 @@ class CosSimKernel : public framework::OpKernel<T> { ...@@ -66,7 +67,7 @@ class CosSimKernel : public framework::OpKernel<T> {
} }
}; };
template <typename Place, typename T> template <typename DeviceContext, typename T>
class CosSimGradKernel : public framework::OpKernel<T> { class CosSimGradKernel : public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext& context) const override { void Compute(const framework::ExecutionContext& context) const override {
...@@ -96,7 +97,8 @@ class CosSimGradKernel : public framework::OpKernel<T> { ...@@ -96,7 +97,8 @@ class CosSimGradKernel : public framework::OpKernel<T> {
auto z_bcast = z.broadcast(bcast_cols); auto z_bcast = z.broadcast(bcast_cols);
auto dz_bcast = dz.broadcast(bcast_cols); auto dz_bcast = dz.broadcast(bcast_cols);
auto x_snorm_bcast = x_norm.square().eval().broadcast(bcast_cols); auto x_snorm_bcast = x_norm.square().eval().broadcast(bcast_cols);
auto place = context.GetEigenDevice<Place>(); auto& place =
*context.template device_context<DeviceContext>().eigen_device();
if (rows_x == rows_y) { if (rows_x == rows_y) {
auto y_snorm_bcast = y_norm.square().eval().broadcast(bcast_cols); auto y_snorm_bcast = y_norm.square().eval().broadcast(bcast_cols);
auto norm_prod_bcast = (x_norm * y_norm).eval().broadcast(bcast_cols); auto norm_prod_bcast = (x_norm * y_norm).eval().broadcast(bcast_cols);
......
paddle/operators/crf_decoding_op.cc
浏览文件 @ aa2f2fb4
...@@ -135,5 +135,6 @@ namespace ops = paddle::operators; ...@@ -135,5 +135,6 @@ namespace ops = paddle::operators;
REGISTER_OP_WITHOUT_GRADIENT(crf_decoding, ops::CRFDecodingOp, REGISTER_OP_WITHOUT_GRADIENT(crf_decoding, ops::CRFDecodingOp,
ops::CRFDecodingOpMaker); ops::CRFDecodingOpMaker);
REGISTER_OP_CPU_KERNEL( REGISTER_OP_CPU_KERNEL(
crf_decoding, ops::CRFDecodingOpKernel<paddle::platform::CPUPlace, float>, crf_decoding,
ops::CRFDecodingOpKernel<paddle::platform::CPUPlace, double>); ops::CRFDecodingOpKernel<paddle::platform::CPUDeviceContext, float>,
ops::CRFDecodingOpKernel<paddle::platform::CPUDeviceContext, double>);
paddle/operators/crf_decoding_op.h
浏览文件 @ aa2f2fb4
...@@ -24,7 +24,7 @@ using framework::LoDTensor; ...@@ -24,7 +24,7 @@ using framework::LoDTensor;
using framework::LoD; using framework::LoD;
using framework::Tensor; using framework::Tensor;
template <typename Place, typename T> template <typename DeviceContext, typename T>
class CRFDecodingOpKernel : public framework::OpKernel<T> { class CRFDecodingOpKernel : public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext& ctx) const override { void Compute(const framework::ExecutionContext& ctx) const override {
...@@ -44,8 +44,8 @@ class CRFDecodingOpKernel : public framework::OpKernel<T> { ...@@ -44,8 +44,8 @@ class CRFDecodingOpKernel : public framework::OpKernel<T> {
const size_t seq_num = lod[level].size() - 1; const size_t seq_num = lod[level].size() - 1;
int64_t* path = decoded_path->mutable_data<int64_t>(platform::CPUPlace()); int64_t* path = decoded_path->mutable_data<int64_t>(platform::CPUPlace());
math::SetConstant<platform::CPUPlace, int64_t>()(ctx.device_context(), math::SetConstant<DeviceContext, int64_t>()(
decoded_path, 0); ctx.template device_context<DeviceContext>(), decoded_path, 0);
for (size_t i = 0; i < seq_num; ++i) { for (size_t i = 0; i < seq_num; ++i) {
int start_pos = static_cast<int>(lod[level][i]); int start_pos = static_cast<int>(lod[level][i]);
int end_pos = static_cast<int>(lod[level][i + 1]); int end_pos = static_cast<int>(lod[level][i + 1]);
......
paddle/operators/crop_op.cc
浏览文件 @ aa2f2fb4
...@@ -133,5 +133,5 @@ class CropOpGrad : public framework::OperatorWithKernel { ...@@ -133,5 +133,5 @@ class CropOpGrad : public framework::OperatorWithKernel {
namespace ops = paddle::operators; namespace ops = paddle::operators;
REGISTER_OP(crop, ops::CropOp, ops::CropOpMaker, crop_grad, ops::CropOpGrad); REGISTER_OP(crop, ops::CropOp, ops::CropOpMaker, crop_grad, ops::CropOpGrad);
REGISTER_OP_CPU_KERNEL(crop, ops::CropKernel<float>); REGISTER_OP_CPU_KERNEL(crop, ops::CropKernel<float>);
REGISTER_OP_CPU_KERNEL(crop_grad, REGISTER_OP_CPU_KERNEL(
ops::CropGradKernel<paddle::platform::CPUPlace, float>); crop_grad, ops::CropGradKernel<paddle::platform::CPUDeviceContext, float>);
paddle/operators/crop_op.cu
浏览文件 @ aa2f2fb4
...@@ -16,6 +16,6 @@ ...@@ -16,6 +16,6 @@
#include "paddle/operators/crop_op.h" #include "paddle/operators/crop_op.h"
namespace ops = paddle::operators; namespace ops = paddle::operators;
REGISTER_OP_GPU_KERNEL(crop, ops::CropKernel<float>); REGISTER_OP_CUDA_KERNEL(crop, ops::CropKernel<float>);
REGISTER_OP_GPU_KERNEL(crop_grad, REGISTER_OP_CUDA_KERNEL(
ops::CropGradKernel<paddle::platform::GPUPlace, float>); crop_grad, ops::CropGradKernel<paddle::platform::CUDADeviceContext, float>);
paddle/operators/crop_op.h
浏览文件 @ aa2f2fb4
...@@ -49,7 +49,7 @@ class CropKernel : public framework::OpKernel<T> { ...@@ -49,7 +49,7 @@ class CropKernel : public framework::OpKernel<T> {
} }
}; };
template <typename Place, typename T, size_t D> template <typename DeviceContext, typename T, size_t D>
void CropGradFunction(const framework::ExecutionContext& context) { void CropGradFunction(const framework::ExecutionContext& context) {
auto* d_x = context.Output<Tensor>(framework::GradVarName("X")); auto* d_x = context.Output<Tensor>(framework::GradVarName("X"));
if (d_x != nullptr) { if (d_x != nullptr) {
...@@ -63,12 +63,13 @@ void CropGradFunction(const framework::ExecutionContext& context) { ...@@ -63,12 +63,13 @@ void CropGradFunction(const framework::ExecutionContext& context) {
} }
auto d_x_tensor = EigenTensor<T, D>::From(*d_x); auto d_x_tensor = EigenTensor<T, D>::From(*d_x);
auto d_out_tensor = EigenTensor<T, D>::From(*d_out); auto d_out_tensor = EigenTensor<T, D>::From(*d_out);
d_x_tensor.device(context.GetEigenDevice<Place>()) = d_x_tensor.device(
*context.template device_context<DeviceContext>().eigen_device()) =
d_out_tensor.pad(paddings, 0); d_out_tensor.pad(paddings, 0);
} }
} }
template <typename Place, typename T> template <typename DeviceContext, typename T>
class CropGradKernel : public framework::OpKernel<T> { class CropGradKernel : public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext& context) const override { void Compute(const framework::ExecutionContext& context) const override {
...@@ -76,22 +77,22 @@ class CropGradKernel : public framework::OpKernel<T> { ...@@ -76,22 +77,22 @@ class CropGradKernel : public framework::OpKernel<T> {
context.Input<Tensor>(framework::GradVarName("Out"))->dims().size(); context.Input<Tensor>(framework::GradVarName("Out"))->dims().size();
switch (rank) { switch (rank) {
case 1: case 1:
CropGradFunction<Place, T, 1>(context); CropGradFunction<DeviceContext, T, 1>(context);
break; break;
case 2: case 2:
CropGradFunction<Place, T, 2>(context); CropGradFunction<DeviceContext, T, 2>(context);
break; break;
case 3: case 3:
CropGradFunction<Place, T, 3>(context); CropGradFunction<DeviceContext, T, 3>(context);
break; break;
case 4: case 4:
CropGradFunction<Place, T, 4>(context); CropGradFunction<DeviceContext, T, 4>(context);
break; break;
case 5: case 5:
CropGradFunction<Place, T, 5>(context); CropGradFunction<DeviceContext, T, 5>(context);
break; break;
case 6: case 6:
CropGradFunction<Place, T, 6>(context); CropGradFunction<DeviceContext, T, 6>(context);
break; break;
default: default:
PADDLE_THROW( PADDLE_THROW(
......
paddle/operators/cross_entropy_op.cu
浏览文件 @ aa2f2fb4
...@@ -53,8 +53,9 @@ class CrossEntropyOpCUDAKernel : public framework::OpKernel<T> { ...@@ -53,8 +53,9 @@ class CrossEntropyOpCUDAKernel : public framework::OpKernel<T> {
Tensor* y = ctx.Output<Tensor>("Y"); Tensor* y = ctx.Output<Tensor>("Y");
y->mutable_data<T>(ctx.GetPlace()); y->mutable_data<T>(ctx.GetPlace());
math::CrossEntropyFunctor<platform::GPUPlace, T>()( math::CrossEntropyFunctor<platform::CUDADeviceContext, T>()(
ctx.device_context(), y, x, label, ctx.Attr<bool>("soft_label")); ctx.template device_context<platform::CUDADeviceContext>(), y, x, label,
ctx.Attr<bool>("soft_label"));
} }
}; };
...@@ -80,15 +81,17 @@ class CrossEntropyGradientOpCUDAKernel : public framework::OpKernel<T> { ...@@ -80,15 +81,17 @@ class CrossEntropyGradientOpCUDAKernel : public framework::OpKernel<T> {
int block = 512; int block = 512;
int grid = (batch_size * class_num + block - 1) / block; int grid = (batch_size * class_num + block - 1) / block;
auto stream = ctx.cuda_device_context().stream();
auto& dev_ctx = ctx.template device_context<platform::CUDADeviceContext>();
auto stream = dev_ctx.stream();
if (ctx.Attr<bool>("soft_label")) { if (ctx.Attr<bool>("soft_label")) {
auto* label_data = label->data<T>(); auto* label_data = label->data<T>();
SoftCrossEntropyGradientKernel<T><<<grid, block, 0, stream>>>( SoftCrossEntropyGradientKernel<T><<<grid, block, 0, stream>>>(
dx_data, dy_data, x_data, label_data, batch_size, class_num); dx_data, dy_data, x_data, label_data, batch_size, class_num);
} else { } else {
math::SetConstant<platform::GPUPlace, T> functor; math::SetConstant<platform::CUDADeviceContext, T> functor;
functor(ctx.device_context(), dx, 0); functor(dev_ctx, dx, 0);
auto* label_data = label->data<int64_t>(); auto* label_data = label->data<int64_t>();
grid = (batch_size + block - 1) / block; grid = (batch_size + block - 1) / block;
CrossEntropyGradientKernel<T><<<grid, block, 0, stream>>>( CrossEntropyGradientKernel<T><<<grid, block, 0, stream>>>(
...@@ -101,8 +104,8 @@ class CrossEntropyGradientOpCUDAKernel : public framework::OpKernel<T> { ...@@ -101,8 +104,8 @@ class CrossEntropyGradientOpCUDAKernel : public framework::OpKernel<T> {
} // namespace paddle } // namespace paddle
namespace ops = paddle::operators; namespace ops = paddle::operators;
REGISTER_OP_GPU_KERNEL(cross_entropy, ops::CrossEntropyOpCUDAKernel<float>, REGISTER_OP_CUDA_KERNEL(cross_entropy, ops::CrossEntropyOpCUDAKernel<float>,
ops::CrossEntropyOpCUDAKernel<double>); ops::CrossEntropyOpCUDAKernel<double>);
REGISTER_OP_GPU_KERNEL(cross_entropy_grad, REGISTER_OP_CUDA_KERNEL(cross_entropy_grad,
ops::CrossEntropyGradientOpCUDAKernel<float>, ops::CrossEntropyGradientOpCUDAKernel<float>,
ops::CrossEntropyGradientOpCUDAKernel<double>); ops::CrossEntropyGradientOpCUDAKernel<double>);
paddle/operators/cross_entropy_op.h
浏览文件 @ aa2f2fb4
...@@ -37,8 +37,9 @@ class CrossEntropyOpKernel : public framework::OpKernel<T> { ...@@ -37,8 +37,9 @@ class CrossEntropyOpKernel : public framework::OpKernel<T> {
Tensor* y = ctx.Output<Tensor>("Y"); Tensor* y = ctx.Output<Tensor>("Y");
y->mutable_data<T>(ctx.GetPlace()); y->mutable_data<T>(ctx.GetPlace());
math::CrossEntropyFunctor<platform::CPUPlace, T>()( math::CrossEntropyFunctor<platform::CPUDeviceContext, T>()(
ctx.device_context(), y, x, labels, ctx.Attr<bool>("soft_label")); ctx.template device_context<platform::CPUDeviceContext>(), y, x, labels,
ctx.Attr<bool>("soft_label"));
} }
}; };
...@@ -61,7 +62,8 @@ class CrossEntropyGradientOpKernel : public framework::OpKernel<T> { ...@@ -61,7 +62,8 @@ class CrossEntropyGradientOpKernel : public framework::OpKernel<T> {
auto lbl_mat = EigenMatrix<T>::From(*label); auto lbl_mat = EigenMatrix<T>::From(*label);
auto dx_mat = EigenMatrix<T>::From(*dx); auto dx_mat = EigenMatrix<T>::From(*dx);
dx_mat.device(ctx.GetEigenDevice<platform::CPUPlace>()) = dx_mat.device(*ctx.template device_context<platform::CPUDeviceContext>()
.eigen_device()) =
-(lbl_mat * -(lbl_mat *
dy_mat.broadcast(Eigen::DSizes<int64_t, 2>(1, class_num)) / x_mat); dy_mat.broadcast(Eigen::DSizes<int64_t, 2>(1, class_num)) / x_mat);
} else { } else {
...@@ -70,8 +72,8 @@ class CrossEntropyGradientOpKernel : public framework::OpKernel<T> { ...@@ -70,8 +72,8 @@ class CrossEntropyGradientOpKernel : public framework::OpKernel<T> {
const T* x_data = x->data<T>(); const T* x_data = x->data<T>();
const int64_t* label_data = label->data<int64_t>(); const int64_t* label_data = label->data<int64_t>();
math::SetConstant<platform::CPUPlace, T> functor; math::SetConstant<platform::CPUDeviceContext, T> functor;
functor(ctx.device_context(), dx, 0); functor(ctx.template device_context<platform::CPUDeviceContext>(), dx, 0);
for (int64_t i = 0; i < batch_size; ++i) { for (int64_t i = 0; i < batch_size; ++i) {
PADDLE_ASSERT(label_data[i] >= 0 || label_data[i] < class_num); PADDLE_ASSERT(label_data[i] >= 0 || label_data[i] < class_num);
......
paddle/operators/decayed_adagrad_op.cc
浏览文件 @ aa2f2fb4
...@@ -99,4 +99,4 @@ REGISTER_OP_WITHOUT_GRADIENT(decayed_adagrad, ops::DecayedAdagradOp, ...@@ -99,4 +99,4 @@ REGISTER_OP_WITHOUT_GRADIENT(decayed_adagrad, ops::DecayedAdagradOp,
ops::DecayedAdagradOpMaker); ops::DecayedAdagradOpMaker);
REGISTER_OP_CPU_KERNEL( REGISTER_OP_CPU_KERNEL(
decayed_adagrad, decayed_adagrad,
ops::DecayedAdagradOpKernel<paddle::platform::CPUPlace, float>); ops::DecayedAdagradOpKernel<paddle::platform::CPUDeviceContext, float>);
paddle/operators/decayed_adagrad_op.cu
浏览文件 @ aa2f2fb4
...@@ -16,6 +16,6 @@ ...@@ -16,6 +16,6 @@
#include "paddle/operators/decayed_adagrad_op.h" #include "paddle/operators/decayed_adagrad_op.h"
namespace ops = paddle::operators; namespace ops = paddle::operators;
REGISTER_OP_GPU_KERNEL( REGISTER_OP_CUDA_KERNEL(
decayed_adagrad, decayed_adagrad,
ops::DecayedAdagradOpKernel<paddle::platform::GPUPlace, float>); ops::DecayedAdagradOpKernel<paddle::platform::CUDADeviceContext, float>);
paddle/operators/decayed_adagrad_op.h
浏览文件 @ aa2f2fb4
...@@ -19,7 +19,7 @@ limitations under the License. */ ...@@ -19,7 +19,7 @@ limitations under the License. */
namespace paddle { namespace paddle {
namespace operators { namespace operators {
template <typename Place, typename T> template <typename DeviceContext, typename T>
class DecayedAdagradOpKernel : public framework::OpKernel<T> { class DecayedAdagradOpKernel : public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext& ctx) const override { void Compute(const framework::ExecutionContext& ctx) const override {
...@@ -43,7 +43,7 @@ class DecayedAdagradOpKernel : public framework::OpKernel<T> { ...@@ -43,7 +43,7 @@ class DecayedAdagradOpKernel : public framework::OpKernel<T> {
auto param_out = framework::EigenVector<T>::Flatten(*param_out_tensor); auto param_out = framework::EigenVector<T>::Flatten(*param_out_tensor);
auto moment_out = framework::EigenVector<T>::Flatten(*moment_out_tensor); auto moment_out = framework::EigenVector<T>::Flatten(*moment_out_tensor);
auto place = ctx.GetEigenDevice<Place>(); auto& place = *ctx.template device_context<DeviceContext>().eigen_device();
moment_out.device(place) = decay * moment + (1 - decay) * grad * grad; moment_out.device(place) = decay * moment + (1 - decay) * grad * grad;
Eigen::DSizes<int, 1> m_dsize(moment_out_tensor->numel()); Eigen::DSizes<int, 1> m_dsize(moment_out_tensor->numel());
......
paddle/operators/dropout_op.cc
浏览文件 @ aa2f2fb4
...@@ -100,6 +100,8 @@ namespace ops = paddle::operators; ...@@ -100,6 +100,8 @@ namespace ops = paddle::operators;
REGISTER_OP(dropout, ops::DropoutOp, ops::DropoutOpMaker<float>, dropout_grad, REGISTER_OP(dropout, ops::DropoutOp, ops::DropoutOpMaker<float>, dropout_grad,
ops::DropoutOpGrad<float>); ops::DropoutOpGrad<float>);
REGISTER_OP_CPU_KERNEL( REGISTER_OP_CPU_KERNEL(
dropout, ops::CPUDropoutKernel<paddle::platform::CPUPlace, float, float>); dropout,
ops::CPUDropoutKernel<paddle::platform::CPUDeviceContext, float, float>);
REGISTER_OP_CPU_KERNEL( REGISTER_OP_CPU_KERNEL(
dropout_grad, ops::DropoutGradKernel<paddle::platform::CPUPlace, float>); dropout_grad,
ops::DropoutGradKernel<paddle::platform::CPUDeviceContext, float>);
paddle/operators/dropout_op.cu
浏览文件 @ aa2f2fb4
...@@ -58,7 +58,7 @@ class GPUDropoutKernel : public framework::OpKernel<T> { ...@@ -58,7 +58,7 @@ class GPUDropoutKernel : public framework::OpKernel<T> {
auto X = EigenMatrix<T>::Reshape(*x, 1); auto X = EigenMatrix<T>::Reshape(*x, 1);
auto Y = EigenMatrix<T>::Reshape(*y, 1); auto Y = EigenMatrix<T>::Reshape(*y, 1);
auto place = context.GetEigenDevice<Place>(); auto& place = *context.template device_context<Place>().eigen_device();
if (!context.Attr<bool>("is_test")) { if (!context.Attr<bool>("is_test")) {
auto* mask = context.Output<Tensor>("Mask"); auto* mask = context.Output<Tensor>("Mask");
auto* mask_data = mask->mutable_data<T>(context.GetPlace()); auto* mask_data = mask->mutable_data<T>(context.GetPlace());
...@@ -80,7 +80,9 @@ class GPUDropoutKernel : public framework::OpKernel<T> { ...@@ -80,7 +80,9 @@ class GPUDropoutKernel : public framework::OpKernel<T> {
} // namespace paddle } // namespace paddle
namespace ops = paddle::operators; namespace ops = paddle::operators;
REGISTER_OP_GPU_KERNEL( REGISTER_OP_CUDA_KERNEL(
dropout, ops::GPUDropoutKernel<paddle::platform::GPUPlace, float, float>); dropout,
REGISTER_OP_GPU_KERNEL( ops::GPUDropoutKernel<paddle::platform::CUDADeviceContext, float, float>);
dropout_grad, ops::DropoutGradKernel<paddle::platform::GPUPlace, float>); REGISTER_OP_CUDA_KERNEL(
dropout_grad,
ops::DropoutGradKernel<paddle::platform::CUDADeviceContext, float>);
paddle/operators/dropout_op.h
浏览文件 @ aa2f2fb4
...@@ -25,7 +25,7 @@ template <typename T, int MajorType = Eigen::RowMajor, ...@@ -25,7 +25,7 @@ template <typename T, int MajorType = Eigen::RowMajor,
typename IndexType = Eigen::DenseIndex> typename IndexType = Eigen::DenseIndex>
using EigenMatrix = framework::EigenMatrix<T, MajorType, IndexType>; using EigenMatrix = framework::EigenMatrix<T, MajorType, IndexType>;
template <typename Place, typename T, typename AttrType> template <typename DeviceContext, typename T, typename AttrType>
class CPUDropoutKernel : public framework::OpKernel<T> { class CPUDropoutKernel : public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext& context) const override { void Compute(const framework::ExecutionContext& context) const override {
...@@ -55,13 +55,14 @@ class CPUDropoutKernel : public framework::OpKernel<T> { ...@@ -55,13 +55,14 @@ class CPUDropoutKernel : public framework::OpKernel<T> {
} else { } else {
auto X = EigenMatrix<T>::Reshape(*x, 1); auto X = EigenMatrix<T>::Reshape(*x, 1);
auto Y = EigenMatrix<T>::Reshape(*y, 1); auto Y = EigenMatrix<T>::Reshape(*y, 1);
auto place = context.GetEigenDevice<Place>(); auto& place =
*context.template device_context<DeviceContext>().eigen_device();
Y.device(place) = X * dropout_prob; Y.device(place) = X * dropout_prob;
} }
} }
}; };
template <typename Place, typename T> template <typename DeviceContext, typename T>
class DropoutGradKernel : public framework::OpKernel<T> { class DropoutGradKernel : public framework::OpKernel<T> {
public: public:
void Compute(const framework::ExecutionContext& context) const override { void Compute(const framework::ExecutionContext& context) const override {
...@@ -77,7 +78,8 @@ class DropoutGradKernel : public framework::OpKernel<T> { ...@@ -77,7 +78,8 @@ class DropoutGradKernel : public framework::OpKernel<T> {
auto dX = EigenMatrix<T>::Reshape(*grad_x, 1); auto dX = EigenMatrix<T>::Reshape(*grad_x, 1);
auto dY = EigenMatrix<T>::Reshape(*grad_y, 1); auto dY = EigenMatrix<T>::Reshape(*grad_y, 1);
auto place = context.GetEigenDevice<Place>(); auto& place =
*context.template device_context<DeviceContext>().eigen_device();
dX.device(place) = dY * M; dX.device(place) = dY * M;
} }
}; };
......
paddle/operators/elementwise_add_op.cc
浏览文件 @ aa2f2fb4
...@@ -34,13 +34,13 @@ REGISTER_OP(elementwise_add, ops::ElementwiseOp, ops::ElementwiseAddOpMaker, ...@@ -34,13 +34,13 @@ REGISTER_OP(elementwise_add, ops::ElementwiseOp, ops::ElementwiseAddOpMaker,
elementwise_add_grad, ops::ElementwiseOpGrad); elementwise_add_grad, ops::ElementwiseOpGrad);
REGISTER_OP_CPU_KERNEL( REGISTER_OP_CPU_KERNEL(
elementwise_add, elementwise_add,
ops::ElementwiseAddKernel<paddle::platform::CPUPlace, float>, ops::ElementwiseAddKernel<paddle::platform::CPUDeviceContext, float>,
ops::ElementwiseAddKernel<paddle::platform::CPUPlace, double>, ops::ElementwiseAddKernel<paddle::platform::CPUDeviceContext, double>,
ops::ElementwiseAddKernel<paddle::platform::CPUPlace, int>, ops::ElementwiseAddKernel<paddle::platform::CPUDeviceContext, int>,
ops::ElementwiseAddKernel<paddle::platform::CPUPlace, int64_t>); ops::ElementwiseAddKernel<paddle::platform::CPUDeviceContext, int64_t>);
REGISTER_OP_CPU_KERNEL( REGISTER_OP_CPU_KERNEL(
elementwise_add_grad, elementwise_add_grad,
ops::ElementwiseAddGradKernel<paddle::platform::CPUPlace, float>, ops::ElementwiseAddGradKernel<paddle::platform::CPUDeviceContext, float>,
ops::ElementwiseAddGradKernel<paddle::platform::CPUPlace, double>, ops::ElementwiseAddGradKernel<paddle::platform::CPUDeviceContext, double>,
ops::ElementwiseAddGradKernel<paddle::platform::CPUPlace, int>, ops::ElementwiseAddGradKernel<paddle::platform::CPUDeviceContext, int>,
ops::ElementwiseAddGradKernel<paddle::platform::CPUPlace, int64_t>); ops::ElementwiseAddGradKernel<paddle::platform::CPUDeviceContext, int64_t>);
paddle/operators/elementwise_add_op.cu
浏览文件 @ aa2f2fb4
...@@ -17,15 +17,16 @@ ...@@ -17,15 +17,16 @@
namespace ops = paddle::operators; namespace ops = paddle::operators;
REGISTER_OP_GPU_KERNEL( REGISTER_OP_CUDA_KERNEL(
elementwise_add, elementwise_add,
ops::ElementwiseAddKernel<paddle::platform::GPUPlace, float>, ops::ElementwiseAddKernel<paddle::platform::CUDADeviceContext, float>,
ops::ElementwiseAddKernel<paddle::platform::GPUPlace, double>, ops::ElementwiseAddKernel<paddle::platform::CUDADeviceContext, double>,
ops::ElementwiseAddKernel<paddle::platform::GPUPlace, int>, ops::ElementwiseAddKernel<paddle::platform::CUDADeviceContext, int>,
ops::ElementwiseAddKernel<paddle::platform::GPUPlace, int64_t>); ops::ElementwiseAddKernel<paddle::platform::CUDADeviceContext, int64_t>);
REGISTER_OP_GPU_KERNEL( REGISTER_OP_CUDA_KERNEL(
elementwise_add_grad, elementwise_add_grad,
ops::ElementwiseAddGradKernel<paddle::platform::GPUPlace, float>, ops::ElementwiseAddGradKernel<paddle::platform::CUDADeviceContext, float>,
ops::ElementwiseAddGradKernel<paddle::platform::GPUPlace, double>, ops::ElementwiseAddGradKernel<paddle::platform::CUDADeviceContext, double>,
ops::ElementwiseAddGradKernel<paddle::platform::GPUPlace, int>, ops::ElementwiseAddGradKernel<paddle::platform::CUDADeviceContext, int>,
ops::ElementwiseAddGradKernel<paddle::platform::GPUPlace, int64_t>); ops::ElementwiseAddGradKernel<paddle::platform::CUDADeviceContext,
int64_t>);
paddle/operators/expand_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/expand_op.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/expand_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/fill_op.cc 0 → 100644
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/ftrl_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/ftrl_op.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/ftrl_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/gather.cu.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/gather_op.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/gather_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/gru_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/gru_op.cu.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/gru_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/gru_unit_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/gru_unit_op.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/gru_unit_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/hinge_loss_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/hinge_loss_op.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/hinge_loss_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/huber_loss_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/huber_loss_op.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/huber_loss_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/l1_norm_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/l1_norm_op.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/l1_norm_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/lod_reset_op.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/lod_reset_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/log_loss_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/log_loss_op.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/log_loss_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/logical_op.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/logical_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/lookup_table_op.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/lrn_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/lrn_op.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/lrn_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/lstm_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/lstm_op.cu.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/lstm_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/lstm_unit_op.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/lstm_unit_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/math/gru_compute.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/math/im2col.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/math/im2col.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/math/im2col.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/math/matmul.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/math/maxouting.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/math/maxouting.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/math/maxouting.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/math/pooling.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/math/pooling.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/math/pooling.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/math/softmax.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/math/softmax.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/math/softmax.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/math/unpooling.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/math/unpooling.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/math/unpooling.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/math/vol2col.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/math/vol2col.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/math/vol2col.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/matmul_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/matmul_op.cu.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/matmul_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/maxout_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/maxout_op.cu.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/maxout_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/mean_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/mean_op.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/mean_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/minus_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/minus_op.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/minus_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/momentum_op.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/mul_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/mul_op.cu.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/mul_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/multiplex_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/multiplex_op.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/multiplex_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/nccl_op.cu.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/nccl_op_test.cu.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/nce_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/nce_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/net_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/net_op_test.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/pad_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/pad_op.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/pad_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/pool_cudnn_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/pool_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/pool_op.cu.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/pool_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/prelu_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/prelu_op.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/prelu_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/proximal_gd_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/proximal_gd_op.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/proximal_gd_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/rank_loss_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/rank_loss_op.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/rank_loss_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/recurrent_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/recv_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/reduce_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/reduce_op.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/reduce_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/reshape_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/reshape_op.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/reshape_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/rmsprop_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/rmsprop_op.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/rmsprop_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/roi_pool_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/roi_pool_op.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/roi_pool_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/row_conv_op.cc 0 → 100644
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/row_conv_op.cu 0 → 100644
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/row_conv_op.h 0 → 100644
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/scale_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/scale_op.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/scale_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/scatter_op.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/seq_expand_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/seq_expand_op.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/seq_expand_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/sequence_conv_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/sequence_pool_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/sgd_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/sgd_op.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/sgd_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/sign_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/sign_op.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/sign_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/softmax_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/softmax_op.cu.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/softmax_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/split_op.cu.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/split_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/sum_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/sum_op.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/sum_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/top_k_op.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/top_k_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/transpose_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/transpose_op.cu.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/transpose_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/unpool_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/unpool_op.cu.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/unpool_op.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/operators/while_op.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/platform/device_context.cc
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/platform/device_context.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/platform/transform.h
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/platform/transform_test.cu
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/scripts/docker/README.md
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/scripts/submit_local.sh.in
浏览文件 @ aa2f2fb4
此差异已折叠。
paddle/utils/Flags.cpp
浏览文件 @ aa2f2fb4
此差异已折叠。
python/paddle/v2/dataset/common.py
浏览文件 @ aa2f2fb4
此差异已折叠。
python/paddle/v2/fluid/layers.py
浏览文件 @ aa2f2fb4
此差异已折叠。
python/paddle/v2/fluid/nets.py
浏览文件 @ aa2f2fb4
此差异已折叠。
python/paddle/v2/fluid/optimizer.py
浏览文件 @ aa2f2fb4
此差异已折叠。
python/paddle/v2/fluid/registry.py 0 → 100644
浏览文件 @ aa2f2fb4
此差异已折叠。
python/paddle/v2/fluid/tests/test_fill_op.py 0 → 100644
浏览文件 @ aa2f2fb4
此差异已折叠。
python/paddle/v2/fluid/tests/test_registry.py 0 → 100644
浏览文件 @ aa2f2fb4
此差异已折叠。
python/paddle/v2/fluid/tests/test_row_conv_op.py 0 → 100644
浏览文件 @ aa2f2fb4
此差异已折叠。
Markdown is supported
0% 或 .
You are about to add 0 people to the discussion. Proceed with caution.
先完成此消息的编辑！
想要评论请 注册 或