data type extension of operators (PaddlePaddle#5039)

* op dtype extension * rename op_dtype_extension_contributing_guides * doc preview test * Update op_dtype_extension_contributing_guides_cn.rst
2017040264 · Jul 20, 2022 · 15496d0 · 15496d0
1 parent 22c9901
commit 15496d0
Show file tree

Hide file tree

Showing 4 changed files with 213 additions and 0 deletions.
diff --git a/docs/dev_guides/images/op_dtype_extension_contributing_guides_cn.png b/docs/dev_guides/images/op_dtype_extension_contributing_guides_cn.png
diff --git a/docs/dev_guides/index_cn.rst b/docs/dev_guides/index_cn.rst
@@ -12,6 +12,7 @@
 - `编译安装 <https://www.paddlepaddle.org.cn/documentation/docs/zh/install/compile/fromsource.html>`_ : 如何从源码编译安装Paddle。
 - `API开发指南 <./api_contributing_guides/api_contributing_guides_cn.html>`_ : API开发相关说明。
 - `算子性能优化贡献指南 <./op_optimization/op_optimization_contributing_guides_cn.html>`_ : 飞桨算子性能优化相关说明。
+- `算子数据类型扩展贡献指南 <./op_dtype_extension/op_dtype_extension_contributing_guides_cn.html>`_ : 飞桨算子数据类型扩展相关说明。
 - `曙光开发指南 <./sugon/index_cn.html>`_ : 曙光开发相关说明。
 - `自定义新硬件接入指南 <./custom_device_docs/index_cn.html>`_: 介绍如何通过自定义硬件功能为飞桨接入新硬件后端。
 - `文档贡献指南 <./docs_contributing_guides_cn.html>`_ : 飞桨文档贡献指南。
@@ -25,6 +26,7 @@
     git_guides/index_cn.rst
     api_contributing_guides/api_contributing_guides_cn.rst
     op_optimization/op_optimization_contributing_guides_cn.rst
+    op_dtype_extension/op_dtype_extension_contributing_guides_cn.rst
     sugon/index_cn.rst
     custom_device_docs/index_cn.rst
     docs_contributing_guides_cn.md

diff --git a/docs/dev_guides/op_dtype_extension/op_dtype_extension_acceptance_criteria_cn.md b/docs/dev_guides/op_dtype_extension/op_dtype_extension_acceptance_criteria_cn.md
@@ -0,0 +1,108 @@
+# 算子数据类型扩展 验收规范
+
+## 通过CI验证
+
+提交至 Paddle repo 的 Pull Request（简称 PR），涉及到的相关检测CI必须全部 Pass。用来验证对之前功能点的兼容和影响，保障新合入代码对历史代码不产生影响。
+
+新增代码必须要有相应的单测保障测试覆盖率达到准入要求（行覆盖率达到90%）。
+
+## 通过精度验证
+
+扩展数据类型后需要添加对应数据类型的单元测试，并通过算子的精度检查。单元测试需要注意以下规范：
+- [OP单测必须使用大尺寸输入](https://github.com/PaddlePaddle/Paddle/wiki/OP-test-input-shape-requirements)
+- [反向Op必须调用check_grad](https://github.com/PaddlePaddle/Paddle/wiki/Gradient-Check-Is-Required-for-Op-Test)
+- [单测精度中atol, rtol, eps, max_relative_error, 不允许自行放大阈值](https://github.com/PaddlePaddle/Paddle/wiki/OP-test-accuracy-requirements)
+
+## 通过性能验证
+
+深度学习框架通常支持多种数据类型的输入，其中低精度运算不仅能够减少显存占用，还可以加快计算的效率。在一些特定硬件上，使用半精度浮点数FP16的峰值计算能力最高可达单精度浮点数FP32的数倍，基于此原理实现的混合精度训练策略对模型也可以实现数倍加速。在完成数据类型的扩展后，可以使用飞桨的[OP Benchmark](https://github.com/PaddlePaddle/benchmark/tree/master/api)算子性能测试专业工具对算子性能进行测试对比，例如对于FP16数据类型，验收的基本要求是算子性能要明显优于使用FP32数据类型，同时我们也鼓励开发者针对FP16类型实现极致的加速。如下图所示，OP Benchmark 测试不同数据类型输入下的OP性能真实状态。
+
+- Conv2d算子，使用FP32数据类型：
+```
+===========================================================================
+-- paddle version             : 0.0.0
+-- paddle commit              : 5040e12e3ea3e640c14add6b9df70e9bfffb8271
+-- benchmark commit           : 26a577ca0c92a9eedb5723dd8c5a57994f967f0e
+-- benchmark last update time : Tue Jul 5 20:32:05 2022 +0800
+===========================================================================
+run command: nvprof --profile-from-start off /usr/bin/python /workspace/benchmark/api/dynamic_tests_v2/conv2d.py --task speed --framework paddle --testing_mode dynamic --json_file /workspace/benchmark/api/tests_v2/configs/conv2d.json --config_id 0 --profiler nvprof --backward False --use_gpu True --repeat 1000 --allow_adaptive_repeat False --log_level 0 --api_name paddle
+            Type  Time(%)      Time     Calls       Avg       Min       Max  Name
+ GPU activities:   56.86%  124.78ms      1000  124.78us  121.31us  128.77us  volta_sgemm_32x128_nn
+                   31.03%  68.097ms      1000  68.096us  66.207us  77.024us  void cudnn::winograd_nonfused::winogradForwardFilter4x4<float, float>(cudnn::winograd_nonfused::WinogradFilterParams<float, float>)
+                    6.27%  13.767ms      1000  13.766us  12.384us  21.216us  void cudnn::winograd_nonfused::winogradForwardData4x4<float, float>(cudnn::winograd_nonfused::WinogradDataParams<float, float>)
+                    5.84%  12.810ms      1000  12.810us  12.031us  18.463us  void cudnn::winograd_nonfused::winogradForwardOutput4x4<float, float>(cudnn::winograd_nonfused::WinogradOutputParams<float, float>)
+
+total gpu_time: 219.4513 ms
+
+W0706 08:37:25.901571 20400 gpu_context.cc:278] Please NOTE: device: 0, GPU Compute Capability: 7.0, Driver API Version: 11.4, Runtime API Version: 11.4
+W0706 08:37:25.901605 20400 gpu_context.cc:306] device: 0, cuDNN Version: 8.2.
+
+[paddle][conv2d] conv2d {
+  run_tf: True
+  run_torch: True
+  atol: 0.01
+  data_format: NCHW
+  dilation: [1, 1]
+  groups: 1
+  padding: [1, 1]
+  stride: 1
+  x_shape: [16, 512, 7, 7]
+  x_dtype: float32
+  weight_shape: [512, 512, 3, 3]
+  weight_dtype: float32
+}
+```
+
+- Conv2d算子，使用FP16数据类型：
+```
+===========================================================================
+-- paddle version             : 0.0.0
+-- paddle commit              : 5040e12e3ea3e640c14add6b9df70e9bfffb8271
+-- benchmark commit           : 26a577ca0c92a9eedb5723dd8c5a57994f967f0e
+-- benchmark last update time : Tue Jul 5 20:32:05 2022 +0800
+===========================================================================
+run command: nvprof --profile-from-start off /usr/bin/python /workspace/benchmark/api/dynamic_tests_v2/conv2d.py --task speed --framework paddle --testing_mode dynamic --json_file /workspace/benchmark/api/tests_v2/configs/conv2d.json --config_id 0 --profiler nvprof --backward False --use_gpu True --repeat 1000 --allow_adaptive_repeat False --log_level 0 --convert_to_fp16 True --api_name paddle
+            Type  Time(%)      Time     Calls       Avg       Min       Max  Name
+ GPU activities:   68.33%  98.867ms      1000  98.867us  98.207us  106.46us  void xmma_cudnn::gemm::kernel<xmma_cudnn::implicit_gemm::fprop_indexed::Kernel_traits<xmma_cudnn::Volta_hmma_fp32_traits, xmma_cudnn::Cta_tile<xmma_cudnn::Volta<int=0>, int=64, int=128, int=32, int=2, int=2, int=1, int=1>, xmma_cudnn::implicit_gemm::fprop_indexed::Gmem_tile_a_t<xmma_cudnn::Volta_hmma_fp32_traits, xmma_cudnn::Cta_tile<xmma_cudnn::Volta<int=0>, int=64, int=128, int=32, int=2, int=2, int=1, int=1>, xmma_cudnn::implicit_gemm::Input_related<int=0, int=0, int=0, bool=0>, int=16, bool=0, xmma_cudnn::implicit_gemm::fprop_indexed::Gmem_tile_base_a<xmma_cudnn::Volta_hmma_fp32_traits, xmma_cudnn::Cta_tile<xmma_cudnn::Volta<int=0>, int=64, int=128, int=32, int=2, int=2, int=1, int=1>, xmma_cudnn::implicit_gemm::Input_related<int=0, int=0, int=0, bool=0>, int=16, xmma_cudnn::Row, int=32, int=64>>, xmma_cudnn::implicit_gemm::fprop_indexed::Gmem_tile_c_t<xmma_cudnn::Volta_hmma_fp32_traits, xmma_cudnn::Cta_tile<xmma_cudnn::Volta<int=0>, int=64, int=128, int=32, int=2, int=2, int=1, int=1>, int=16, xmma_cudnn::Fragment_c<xmma_cudnn::Volta_hmma_fp32_traits, xmma_cudnn::Cta_tile<xmma_cudnn::Volta<int=0>, int=64, int=128, int=32, int=2, int=2, int=1, int=1>, bool=0>>, xmma_cudnn::implicit_gemm::Input_related<int=0, int=0, int=0, bool=0>, int=1>>(xmma_cudnn::Volta_hmma_fp32_traitsParams)
+                   26.13%  37.801ms      2000  18.900us  6.0160us  40.161us  void cudnn::ops::nchwToNhwcKernel<__half, __half, float, bool=0, bool=1, cudnnKernelDataType_t=0>(cudnn::ops::nchw2nhwc_params_t<float>, __half const *, __half*)
+                    4.30%  6.2263ms      1000  6.2260us  5.9200us  11.200us  void cudnn::ops::nhwcToNchwKernel<__half, __half, float, bool=1, bool=0, cudnnKernelDataType_t=0>(cudnn::ops::nhwc2nchw_params_t<float>, __half const *, __half*)
+                    1.24%  1.7880ms      1000  1.7880us  1.7270us  7.6480us  [CUDA memset]
+
+total gpu_time: 144.6905 ms
+
+W0706 08:37:25.901571 20400 gpu_context.cc:278] Please NOTE: device: 0, GPU Compute Capability: 7.0, Driver API Version: 11.4, Runtime API Version: 11.4
+W0706 08:37:25.901605 20400 gpu_context.cc:306] device: 0, cuDNN Version: 8.2.
+
+[paddle][conv2d] conv2d {
+  run_tf: True
+  run_torch: True
+  atol: 0.01
+  data_format: NCHW
+  dilation: [1, 1]
+  groups: 1
+  padding: [1, 1]
+  stride: 1
+  x_shape: [16, 512, 7, 7]
+  x_dtype: float16
+  weight_shape: [512, 512, 3, 3]
+  weight_dtype: float16
+}
+```
+
+## PR内容描述要求
+
+单元测试内容需要和开发代码放在同一个PR提交，后续修改也需要基于此PR。PR内容描述测试部分需要明确描述下列内容：
+
+1. 针对低精度数据类型的支持方法描述：概要说明计算精度是否对不同数据类型敏感，如Transpose算子的计算精度与数据类型无关
+
+2. 扩展数据类型后算子的性能状况：给出不同数据类型下算子性能，如FP32和FP16的性能对比
+
+3. PR性能优化方案概述：如果扩展数据类型时，还对算子进行了性能优化，则需要描述优化方案
+
+## 交流与改进
+
+PR的单测部分必须通过 Paddle 测试人员 review，确保完整覆盖了待测功能点后，会给予 approved。如果 review 过程中发现测试缺失和遗漏的测试点，会通过 GitHub 代码行 Comment 的和 Request Changes 的方式交流改进，待PR修改完毕后给予 approved。
+
+## 后续维护
+
+代码成功合入后，如果发现对框架造成了精度和性能下降影响，或者和部分功能存在严重冲突导致Bug，会对代码进行 Revert 并通过 ISSUE 告知相关的开发者，请提交 PR 修复问题，并重新合入。
diff --git a/docs/dev_guides/op_dtype_extension/op_dtype_extension_contributing_guides_cn.rst b/docs/dev_guides/op_dtype_extension/op_dtype_extension_contributing_guides_cn.rst
@@ -0,0 +1,103 @@
+#####################
+算子数据类型扩展 提交流程
+#####################
+
+飞桨作为一个开源项目，我们鼓励生态开发者为PaddlePaddle完善和优化各类算子，当你想为飞桨扩展数据类型的支持时，请遵守下述贡献流程，在 GitHub 上完成文档设计和代码设计并提交至相应的 GitHub 仓库。
+
+
+算子数据类型扩展贡献流程
+::::::::::::::::::::::
+
+.. image:: ../images/op_dtype_extension_contributing_guides_cn.png
+  :width: 900
+  :alt: op_dtype_extension_contributing_guides
+  :align: center
+
+
+流程介绍
+::::::::::::::::::::::
+
+**1. 任务认领**
+
+如果你想参与飞桨 OP 开源贡献，可以在 GitHub Paddle 项目上的 Issue 区域进行任务认领，飞桨官方会发布一些 OP 优化任务，你可以认领其中的算子数据类型扩展任务，并按照此贡献流程提交算子数据类型扩展的设计文档。
+
+**2. 签订贡献者许可协议（CLA）**
+
+对于你贡献的源代码，你将拥有合法的知识产权，为了保护你的权益，你需要先签署一份 `贡献者许可协议 <https://cla-assistant.io/PaddlePaddle/Paddle?pullRequest=39047>`_ 。
+
+**注意**：只有当你签署完CLA后，我们才能继续对你提交的设计方案和实现代码进行评审及合入。
+
+**3. 根据开发难点提交必要的设计文档**
+
+算子数据类型扩展是针对飞桨已有的算子，因此可以只在必要情况下提供设计文档：
+
+- 不需要提供设计文档：如果算子性能符合预期，比如算子使用FP16数据类型的计算性能明显优于使用FP32数据类型，代码实现中仅涉及了算子的功能增强，则不需要撰写设计文档。我们鼓励开发者在进行算子数据类型扩展时首先进行方案设计，你可以将自己的设计思路提交在PR中。
+- 需要提供设计文档：如果算子在扩展了数据类型后性能不符合预期，例如为算子扩展了FP16数据类型，但是性能却差于使用FP32数据类型计算，则需要针对FP16数据类型实现性能的优化，这种情况下需要添加算子性能优化设计文档，这也是为了促进社区开发者更容易的参与开源项目共建，开发者通过与飞桨专家和社区其他用户进行广泛的交流，完善设计方案和PR请求，在提交实现代码之前确保方案符合飞桨的设计理念，同时也便于后续的代码评审及合入工作。你需要将设计文档提交至 GitHub 开发者社区仓库，并根据本地开发指南提交PR。此过程请参考相应的开发规范，并提交以下内容：
+
+.. csv-table::
+    :header: "提交内容", "参考文档", "提交位置"
+    :widths: 10, 30, 30
+
+    "算子性能优化设计文档", "- `算子性能优化设计文档 模版 <https://github.com/PaddlePaddle/community/tree/master/rfcs/OPs-Perf/op_optimization_template.md>`_
+    - `算子性能优化设计文档 示例 <https://github.com/PaddlePaddle/community/blob/master/rfcs/OPs-Perf/op_optimization_example.md>`_ ", "`Github开发者社区仓库 <https://github.com/PaddlePaddle/community/tree/master/rfcs/OPs-Perf>`_"
+
+
+**4. 设计文档评审&公示**
+
+针对需要提供设计文档的算子，飞桨专家会对你提交的设计文档进行审核，同时此文档也将接受来自开发者社区的评估，所有开发者都可以在 PR 评论区进行广泛的交流。开发者根据飞桨专家和其他开发者的反馈意见进行讨论并做出修改，最终评审通过后会合入。
+
+如果你的设计方案比较复杂，会在社区中针对算子的设计文档发起评审会议。飞桨会提前在 PR 评论区公布会议时间、会议地址、参会人、议题等内容，请及时关注pr中最新动态，你也可以在评论区自行申请评审会。会议结束后，会在 PR 中发出会议结论。
+
+**5. 公布评审结果&合入文档**
+
+当设计文档评审&公示通过后，你的算子优化设计文档将会合入至 `飞桨开发者社区仓库 <https://github.com/PaddlePaddle/community>`_ ，并在开源社区中同步。
+
+**6. 提交算子优化代码**
+
+随后，你可以根据评审通过的设计内容进行代码开发。此过程请参考相应的开发规范，并提交以下内容：
+
+.. csv-table::
+    :header: "提交内容", "参考文档", "提交位置"
+    :widths: 10, 30,30
+
+    "1、算子数据类型扩展实现代码", "- `Paddle代码规范 <https://www.paddlepaddle.org.cn/documentation/docs/zh/dev_guides/style_guides_cn.html>`_
+    - `C++ OP开发指南 <../api_contributing_guides/new_cpp_op_cn.html>`_
+    ", "`Github飞桨训练框架仓库 <https://github.com/PaddlePaddle/Paddle>`_"
+    "2、API英文文档", "- `API文档书写规范 <https://github.com/PaddlePaddle/docs/wiki/%E9%A3%9E%E6%A1%A8API%E6%96%87%E6%A1%A3%E4%B9%A6%E5%86%99%E8%A7%84%E8%8C%83>`_
+    ", "`Github飞桨训练框架仓库 <https://github.com/PaddlePaddle/Paddle>`_"
+    "3、API中文文档", "- `API文档书写规范 <https://github.com/PaddlePaddle/docs/wiki/%E9%A3%9E%E6%A1%A8API%E6%96%87%E6%A1%A3%E4%B9%A6%E5%86%99%E8%A7%84%E8%8C%83>`_
+    - `中文文档贡献指南 <../docs_contributing_guides_cn.html>`_
+    ", "`Github飞桨文档仓库 <https://github.com/PaddlePaddle/docs>`_"
+    "4、算子不同数据类型性能对比", "- `OP Benchmark使用指南 <https://github.com/PaddlePaddle/benchmark/blob/master/api>`_
+    - `算子性能优化 优化方法 <../op_optimization/op_optimization_method_introduction_cn.html>`_
+    - `算子数据类型扩展 验收规范 <./op_dtype_extension_acceptance_criteria_cn.html>`_
+    ", "`Github飞桨训练框架仓库 <https://github.com/PaddlePaddle/Paddle>`_"
+
+
+当你完成以上代码设计后，需要将代码提交至 `Github飞桨训练框架仓库 <https://github.com/PaddlePaddle/Paddle>`_ 和 `Github飞桨中文文档仓库 <https://github.com/PaddlePaddle/docs>`_ ，并根据 `本地开发指南 <https://www.paddlepaddle.org.cn/documentation/docs/zh/develop/dev_guides/git_guides/local_dev_guide_cn.html>`_ 提交PR、准备接受社区的评审。
+
+**7. 实现代码评审&公示**
+
+飞桨官方会及时安排专家进行算子优化代码审核，代码也将接受来自开发者社区的评审，所有开发者可以在 PR 评论区进行交流。请你对飞桨专家和其他开发者的反馈意见进行讨论并做出修改，最终评审通过后会在开源社区中同步。
+
+如果你的代码实现逻辑比较复杂，会在社区中针对算子的设计文档发起评审会议。飞桨会提前在 PR 评论区公布会议时间、会议地址、参会人、议题等内容，请及时关注pr中最新动态，你也可以在评论区自行申请评审会。会议结束后，会在 PR 中发出会议结论。
+
+**8. 公布评审结果&合入代码**
+
+当算子优化代码评审&公示通过，官方会在开源社区中同步，你所实现的优化代码将会合入至 `Github飞桨训练框架仓库 <https://github.com/PaddlePaddle/Paddle>`_ 。
+
+**9. 通过集成测试、精度和性能验收**
+
+当你的代码合入 `Github飞桨训练框架仓库 <https://github.com/PaddlePaddle/Paddle>`_ 后，飞桨会对你的优化代码进行模型级别的集成测试、精度和性能的验收，并告知你测试结果。如果测试通过，恭喜你贡献流程已经全部完成。如果测试不通过，会通过 ISSUE 联系你进行代码修复，请及时关注 GitHub上的最新动态。
+
+**注意**：代码合入 develop 分支之后的第二天，你可以从官网下载 develop 安装包体验此功能。飞桨后续也会将此功能纳入下一个正式版的发版计划。
+
+**10. 贡献完成**
+
+感谢你的贡献！
+
+
+..  toctree::
+    :hidden:
+
+    op_dtype_extension_acceptance_criteria_cn.md