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 new file mode 100644 index 00000000000..8a8d450466c Binary files /dev/null and b/docs/dev_guides/images/op_dtype_extension_contributing_guides_cn.png differ diff --git a/docs/dev_guides/index_cn.rst b/docs/dev_guides/index_cn.rst index 0e7e0938e14..40213432afb 100644 --- a/docs/dev_guides/index_cn.rst +++ b/docs/dev_guides/index_cn.rst @@ -12,6 +12,7 @@ - `编译安装 `_ : 如何从源码编译安装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 new file mode 100644 index 00000000000..4acad719101 --- /dev/null +++ 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(cudnn::winograd_nonfused::WinogradFilterParams) + 6.27% 13.767ms 1000 13.766us 12.384us 21.216us void cudnn::winograd_nonfused::winogradForwardData4x4(cudnn::winograd_nonfused::WinogradDataParams) + 5.84% 12.810ms 1000 12.810us 12.031us 18.463us void cudnn::winograd_nonfused::winogradForwardOutput4x4(cudnn::winograd_nonfused::WinogradOutputParams) + +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, int=64, int=128, int=32, int=2, int=2, int=1, int=1>, xmma_cudnn::implicit_gemm::fprop_indexed::Gmem_tile_a_t, int=64, int=128, int=32, int=2, int=2, int=1, int=1>, xmma_cudnn::implicit_gemm::Input_related, int=16, bool=0, xmma_cudnn::implicit_gemm::fprop_indexed::Gmem_tile_base_a, int=64, int=128, int=32, int=2, int=2, int=1, int=1>, xmma_cudnn::implicit_gemm::Input_related, int=16, xmma_cudnn::Row, int=32, int=64>>, xmma_cudnn::implicit_gemm::fprop_indexed::Gmem_tile_c_t, int=64, int=128, int=32, int=2, int=2, int=1, int=1>, int=16, xmma_cudnn::Fragment_c, int=64, int=128, int=32, int=2, int=2, int=1, int=1>, bool=0>>, xmma_cudnn::implicit_gemm::Input_related, 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, __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, __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 new file mode 100644 index 00000000000..b5ad1fdca7d --- /dev/null +++ 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）** + +对于你贡献的源代码，你将拥有合法的知识产权，为了保护你的权益，你需要先签署一份 `贡献者许可协议 `_ 。 + +**注意**：只有当你签署完CLA后，我们才能继续对你提交的设计方案和实现代码进行评审及合入。 + +**3. 根据开发难点提交必要的设计文档** + +算子数据类型扩展是针对飞桨已有的算子，因此可以只在必要情况下提供设计文档： + +- 不需要提供设计文档：如果算子性能符合预期，比如算子使用FP16数据类型的计算性能明显优于使用FP32数据类型，代码实现中仅涉及了算子的功能增强，则不需要撰写设计文档。我们鼓励开发者在进行算子数据类型扩展时首先进行方案设计，你可以将自己的设计思路提交在PR中。 +- 需要提供设计文档：如果算子在扩展了数据类型后性能不符合预期，例如为算子扩展了FP16数据类型，但是性能却差于使用FP32数据类型计算，则需要针对FP16数据类型实现性能的优化，这种情况下需要添加算子性能优化设计文档，这也是为了促进社区开发者更容易的参与开源项目共建，开发者通过与飞桨专家和社区其他用户进行广泛的交流，完善设计方案和PR请求，在提交实现代码之前确保方案符合飞桨的设计理念，同时也便于后续的代码评审及合入工作。你需要将设计文档提交至 GitHub 开发者社区仓库，并根据本地开发指南提交PR。此过程请参考相应的开发规范，并提交以下内容： + +.. csv-table:: + :header: "提交内容", "参考文档", "提交位置" + :widths: 10, 30, 30 + + "算子性能优化设计文档", "- `算子性能优化设计文档 模版 `_ + - `算子性能优化设计文档 示例 `_ ", "`Github开发者社区仓库 `_" + + +**4. 设计文档评审&公示** + +针对需要提供设计文档的算子，飞桨专家会对你提交的设计文档进行审核，同时此文档也将接受来自开发者社区的评估，所有开发者都可以在 PR 评论区进行广泛的交流。开发者根据飞桨专家和其他开发者的反馈意见进行讨论并做出修改，最终评审通过后会合入。 + +如果你的设计方案比较复杂，会在社区中针对算子的设计文档发起评审会议。飞桨会提前在 PR 评论区公布会议时间、会议地址、参会人、议题等内容，请及时关注pr中最新动态，你也可以在评论区自行申请评审会。会议结束后，会在 PR 中发出会议结论。 + +**5. 公布评审结果&合入文档** + +当设计文档评审&公示通过后，你的算子优化设计文档将会合入至 `飞桨开发者社区仓库 `_ ，并在开源社区中同步。 + +**6. 提交算子优化代码** + +随后，你可以根据评审通过的设计内容进行代码开发。此过程请参考相应的开发规范，并提交以下内容： + +.. csv-table:: + :header: "提交内容", "参考文档", "提交位置" + :widths: 10, 30,30 + + "1、算子数据类型扩展实现代码", "- `Paddle代码规范 `_ + - `C++ OP开发指南 <../api_contributing_guides/new_cpp_op_cn.html>`_ + ", "`Github飞桨训练框架仓库 `_" + "2、API英文文档", "- `API文档书写规范 `_ + ", "`Github飞桨训练框架仓库 `_" + "3、API中文文档", "- `API文档书写规范 `_ + - `中文文档贡献指南 <../docs_contributing_guides_cn.html>`_ + ", "`Github飞桨文档仓库 `_" + "4、算子不同数据类型性能对比", "- `OP Benchmark使用指南 `_ + - `算子性能优化 优化方法 <../op_optimization/op_optimization_method_introduction_cn.html>`_ + - `算子数据类型扩展 验收规范 <./op_dtype_extension_acceptance_criteria_cn.html>`_ + ", "`Github飞桨训练框架仓库 `_" + + +当你完成以上代码设计后，需要将代码提交至 `Github飞桨训练框架仓库 `_ 和 `Github飞桨中文文档仓库 `_ ，并根据 `本地开发指南 `_ 提交PR、准备接受社区的评审。 + +**7. 实现代码评审&公示** + +飞桨官方会及时安排专家进行算子优化代码审核，代码也将接受来自开发者社区的评审，所有开发者可以在 PR 评论区进行交流。请你对飞桨专家和其他开发者的反馈意见进行讨论并做出修改，最终评审通过后会在开源社区中同步。 + +如果你的代码实现逻辑比较复杂，会在社区中针对算子的设计文档发起评审会议。飞桨会提前在 PR 评论区公布会议时间、会议地址、参会人、议题等内容，请及时关注pr中最新动态，你也可以在评论区自行申请评审会。会议结束后，会在 PR 中发出会议结论。 + +**8. 公布评审结果&合入代码** + +当算子优化代码评审&公示通过，官方会在开源社区中同步，你所实现的优化代码将会合入至 `Github飞桨训练框架仓库 `_ 。 + +**9. 通过集成测试、精度和性能验收** + +当你的代码合入 `Github飞桨训练框架仓库 `_ 后，飞桨会对你的优化代码进行模型级别的集成测试、精度和性能的验收，并告知你测试结果。如果测试通过，恭喜你贡献流程已经全部完成。如果测试不通过，会通过 ISSUE 联系你进行代码修复，请及时关注 GitHub上的最新动态。 + +**注意**：代码合入 develop 分支之后的第二天，你可以从官网下载 develop 安装包体验此功能。飞桨后续也会将此功能纳入下一个正式版的发版计划。 + +**10. 贡献完成** + +感谢你的贡献！ + + +.. toctree:: + :hidden: + + op_dtype_extension_acceptance_criteria_cn.md