已合并
fix: Modify the AIDD docs issue #8149
zwj223创建于 7月31日
fix: Modify the AIDD docs issue #8149
已合并
zwj223创建于 7月31日
22 个文件变更+162-163
@@ -6,7 +6,7 @@
6 6 
7发布日期:2026-03-307发布日期:2026-03-30
8 8 
9-本版本新增多项新增特性、问题修复及性能改进,支持最新的950硬件。9+本版本多项新增特性、问题修复及性能改进,支持最新的950硬件。
10我们诚挚欢迎社区反馈,以进一步提升ops-nn的稳定性和功能完备性。10我们诚挚欢迎社区反馈,以进一步提升ops-nn的稳定性和功能完备性。
11使用方式请参阅[官方文档](https://gitcode.com/cann/ops-nn/blob/master/README.md)。11使用方式请参阅[官方文档](https://gitcode.com/cann/ops-nn/blob/master/README.md)。
12 12 
@@ -77,7 +77,7 @@ ops-nn算子首个Beta版本v8.5.0-beta.1现已发布。
77 77 
78### 🔗 版本地址78### 🔗 版本地址
79 79 
80-[CANN 8.5.0-beta 1](https://ascend.devcloud.huaweicloud.com/cann/run/software/8.5.0-beta.1/)80+[CANN 8.5.0-beta.1](https://ascend.devcloud.huaweicloud.com/cann/run/software/8.5.0-beta.1/)
81 81 
82```text82```text
83版本目录说明如下:83版本目录说明如下:
@@ -31,8 +31,8 @@
31| 依赖 | 不涉及 | cmake/third_party/nlohmann_json.cmake | [https://gitcode.com/cann-src-third-party/json/releases/download/v3.11.3/include.zip](https://gitcode.com/cann-src-third-party/json/releases/download/v3.11.3/include.zip) | 从gitcode下载json源码,作用编译依赖 |31| 依赖 | 不涉及 | cmake/third_party/nlohmann_json.cmake | [https://gitcode.com/cann-src-third-party/json/releases/download/v3.11.3/include.zip](https://gitcode.com/cann-src-third-party/json/releases/download/v3.11.3/include.zip) | 从gitcode下载json源码,作用编译依赖 |
32| 依赖 | 不涉及 | cmake/third_party/gtest.cmake | [https://gitcode.com/cann-src-third-party/googletest/releases/download/v1.14.0/googletest-1.14.0.tar.gz](https://gitcode.com/cann-src-third-party/googletest/releases/download/v1.14.0/googletest-1.14.0.tar.gz) | 从gitcode下载googletest源码,作用编译依赖 |32| 依赖 | 不涉及 | cmake/third_party/gtest.cmake | [https://gitcode.com/cann-src-third-party/googletest/releases/download/v1.14.0/googletest-1.14.0.tar.gz](https://gitcode.com/cann-src-third-party/googletest/releases/download/v1.14.0/googletest-1.14.0.tar.gz) | 从gitcode下载googletest源码,作用编译依赖 |
33| 依赖 | 不涉及 | cmake/third_party/eigen.cmake | [https://gitcode.com/cann-src-third-party/eigen/releases/download/5.0.0-h0.trunk/eigen-5.0.0.tar.gz](https://gitcode.com/cann-src-third-party/eigen/releases/download/5.0.0-h0.trunk/eigen-5.0.0.tar.gz) | 从gitcode下载eigen源码,作用编译依赖 |33| 依赖 | 不涉及 | cmake/third_party/eigen.cmake | [https://gitcode.com/cann-src-third-party/eigen/releases/download/5.0.0-h0.trunk/eigen-5.0.0.tar.gz](https://gitcode.com/cann-src-third-party/eigen/releases/download/5.0.0-h0.trunk/eigen-5.0.0.tar.gz) | 从gitcode下载eigen源码,作用编译依赖 |
34-| 依赖 | 不涉及 | ops-nn/install_deps.sh | [https://apt.kitware.com/keys/kitware-archive-latest.asc](https://apt.kitware.com/keys/kitware-archive-latest.asc) | 从gitcode下载install_deps源码,作用编译依赖 |34+| 依赖 | 不涉及 | ops-nn/install_deps.sh | [https://apt.kitware.com/keys/kitware-archive-latest.asc](https://apt.kitware.com/keys/kitware-archive-latest.asc) | 从apt.kitware.com下载Kitware签名密钥,作用编译依赖 |
35-| 依赖 | 不涉及 | ops-nn/install_deps.sh | [https://apt.kitware.com/ubuntu/](https://apt.kitware.com/ubuntu/) | 从gitcode下载install_deps源码,作用编译依赖 |35+| 依赖 | 不涉及 | ops-nn/install_deps.sh | [https://apt.kitware.com/ubuntu/](https://apt.kitware.com/ubuntu/) | 从apt.kitware.com下载Kitware签名密钥,作用编译依赖 |
36| 依赖 | 不涉及 | cmake | [https://apt.kitware.com/keys/kitware-archive-latest.asc](https://apt.kitware.com/keys/kitware-archive-latest.asc) | 从kitware下载cmake软件,作用编译依赖 |36| 依赖 | 不涉及 | cmake | [https://apt.kitware.com/keys/kitware-archive-latest.asc](https://apt.kitware.com/keys/kitware-archive-latest.asc) | 从kitware下载cmake软件,作用编译依赖 |
37| 依赖 | 不涉及 | cmake | [https://apt.kitware.com/ubuntu/](https://apt.kitware.com/ubuntu/) | 从kitware下载cmake软件,作用编译依赖 |37| 依赖 | 不涉及 | cmake | [https://apt.kitware.com/ubuntu/](https://apt.kitware.com/ubuntu/) | 从kitware下载cmake软件,作用编译依赖 |
38 38 
@@ -60,6 +60,6 @@
60| 维护升级文件目录 | 770(rwxrwx---) |60| 维护升级文件目录 | 770(rwxrwx---) |
61| 业务数据文件 | 640(rw-r-----) |61| 业务数据文件 | 640(rw-r-----) |
62| 业务数据文件目录 | 750(rwxr-x---) |62| 业务数据文件目录 | 750(rwxr-x---) |
63-| 密钥组件、私钥、证书、密文文件目录 | 700(rwx-----) |63+| 密钥组件、私钥、证书、密文文件目录 | 700(rwx------) |
64| 密钥组件、私钥、证书、加密密文 | 600(rw-------) |64| 密钥组件、私钥、证书、加密密文 | 600(rw-------) |
65| 加解密接口、加解密脚本 | 500(r-x------) |65| 加解密接口、加解密脚本 | 500(r-x------) |
@@ -11,8 +11,8 @@
11|7 |[G.FMT.04-CPP每个变量单独一行进行声明或赋值](# G.FMT.04-CPP每个变量单独一行进行声明或赋值)|[G.FMT.04-CPP Each variable is declared or assigned on a separate line](# G.FMT.04-CPP每个变量单独一行进行声明或赋值)|check_multiple_var_dec:true|C++|版本级,门禁级|11|7 |[G.FMT.04-CPP每个变量单独一行进行声明或赋值](# G.FMT.04-CPP每个变量单独一行进行声明或赋值)|[G.FMT.04-CPP Each variable is declared or assigned on a separate line](# G.FMT.04-CPP每个变量单独一行进行声明或赋值)|check_multiple_var_dec:true|C++|版本级,门禁级|
12|8 |[G.FMT.05-CPP行宽不超过120个字符](# G.FMT.05-CPP行宽不超过120个字符)|[G.FMT.05-CPP Each line contains no more than 120 characters](# G.FMT.05-CPP行宽不超过120个字符)|column_width:120|C++|版本级,门禁级|12|8 |[G.FMT.05-CPP行宽不超过120个字符](# G.FMT.05-CPP行宽不超过120个字符)|[G.FMT.05-CPP Each line contains no more than 120 characters](# G.FMT.05-CPP行宽不超过120个字符)|column_width:120|C++|版本级,门禁级|
13|9 |[G.FMT.06-CPP换行时将操作符留在行末,新行缩进一层或进行同类对齐--函数调用参数换行](# G.FMT.06-CPP换行时将操作符留在行末,新行缩进一层或进行同类对齐--函数调用参数换行)|[G.FMT.06-CPP While line breaking, leave the operators at the end and indent or align the new line--Function call parameter wrapping](# G.FMT.06-CPP换行时将操作符留在行末,新行缩进一层或进行同类对齐--函数调用参数换行)|c_coding_standard_v5.0:true|C++|版本级,门禁级|13|9 |[G.FMT.06-CPP换行时将操作符留在行末,新行缩进一层或进行同类对齐--函数调用参数换行](# G.FMT.06-CPP换行时将操作符留在行末,新行缩进一层或进行同类对齐--函数调用参数换行)|[G.FMT.06-CPP While line breaking, leave the operators at the end and indent or align the new line--Function call parameter wrapping](# G.FMT.06-CPP换行时将操作符留在行末,新行缩进一层或进行同类对齐--函数调用参数换行)|c_coding_standard_v5.0:true|C++|版本级,门禁级|
14-|10 |[G.FMT.06-CPP换行时将操作符留在行末,新行缩进一层或进行同类对齐--函数声明参数换行](# G.FMT.06-CPP换行时将操作符留在行末,新行缩进一层或进行同类对齐--函数声明参数换行)|[G.FMT.06-CPP While line breaking, leave the operators at the end and indent or align the new line--Data initialization wrapping](# G.FMT.06-CPP换行时将操作符留在行末,新行缩进一层或进行同类对齐--函数声明参数换行)|-|C++|版本级,门禁级|14+|10 |[G.FMT.06-CPP换行时将操作符留在行末,新行缩进一层或进行同类对齐--函数声明参数换行](# G.FMT.06-CPP换行时将操作符留在行末,新行缩进一层或进行同类对齐--函数声明参数换行)|[G.FMT.06-CPP While line breaking, leave the operators at the end and indent or align the new line--Function declaration parameter wrapping](# G.FMT.06-CPP换行时将操作符留在行末,新行缩进一层或进行同类对齐--函数声明参数换行)|-|C++|版本级,门禁级|
15-|11 |[G.FMT.06-CPP换行时将操作符留在行末,新行缩进一层或进行同类对齐--数据初始化换行](# G.FMT.06-CPP换行时将操作符留在行末,新行缩进一层或进行同类对齐--数据初始化换行)|[G.FMT.06-CPP While line breaking, leave the operators at the end and indent or align the new line--Function declaration parameter wrapping](# G.FMT.06-CPP换行时将操作符留在行末,新行缩进一层或进行同类对齐--数据初始化换行)|-|C++|版本级,门禁级|15+|11 |[G.FMT.06-CPP换行时将操作符留在行末,新行缩进一层或进行同类对齐--数据初始化换行](# G.FMT.06-CPP换行时将操作符留在行末,新行缩进一层或进行同类对齐--数据初始化换行)|[G.FMT.06-CPP While line breaking, leave the operators at the end and indent or align the new line--Data initialization wrapping](# G.FMT.06-CPP换行时将操作符留在行末,新行缩进一层或进行同类对齐--数据初始化换行)|-|C++|版本级,门禁级|
16|12 |[G.FMT.07-CPP预处理的"#"统一放在行首,嵌套预处理语句时,"#"可以进行缩进](# G.FMT.07-CPP预处理的"#"统一放在行首,嵌套预处理语句时,"#"可以进行缩进)|[G.FMT.07-CPP The number sign (#) that starts a preprocessor directive should be placed at the beginning of a line and can be indented in nested](# G.FMT.07-CPP预处理的"#"统一放在行首,嵌套预处理语句时,"#"可以进行缩进)|-|C++|版本级,门禁级|16|12 |[G.FMT.07-CPP预处理的"#"统一放在行首,嵌套预处理语句时,"#"可以进行缩进](# G.FMT.07-CPP预处理的"#"统一放在行首,嵌套预处理语句时,"#"可以进行缩进)|[G.FMT.07-CPP The number sign (#) that starts a preprocessor directive should be placed at the beginning of a line and can be indented in nested](# G.FMT.07-CPP预处理的"#"统一放在行首,嵌套预处理语句时,"#"可以进行缩进)|-|C++|版本级,门禁级|
17|13 |[G.FMT.08-CPP遵循传统的类成员声明顺序](# G.FMT.08-CPP遵循传统的类成员声明顺序)|[G.FMT.08-CPP Use a conventional class member declaration order](# G.FMT.08-CPP遵循传统的类成员声明顺序)|-|C++|版本级,门禁级|17|13 |[G.FMT.08-CPP遵循传统的类成员声明顺序](# G.FMT.08-CPP遵循传统的类成员声明顺序)|[G.FMT.08-CPP Use a conventional class member declaration order](# G.FMT.08-CPP遵循传统的类成员声明顺序)|-|C++|版本级,门禁级|
18|14 |[G.FMT.09-CPP构造函数初始化列表放在同一行或按4空格缩进并排多行](# G.FMT.09-CPP构造函数初始化列表放在同一行或按4空格缩进并排多行)|[G.FMT.09-CPP A constructor initialization list is placed on the same line or on different lines indented by four spaces](# G.FMT.09-CPP构造函数初始化列表放在同一行或按4空格缩进并排多行)|-|C++|版本级,门禁级|18|14 |[G.FMT.09-CPP构造函数初始化列表放在同一行或按4空格缩进并排多行](# G.FMT.09-CPP构造函数初始化列表放在同一行或按4空格缩进并排多行)|[G.FMT.09-CPP A constructor initialization list is placed on the same line or on different lines indented by four spaces](# G.FMT.09-CPP构造函数初始化列表放在同一行或按4空格缩进并排多行)|-|C++|版本级,门禁级|
@@ -44,7 +44,7 @@
44使用//44使用//
45 45 
46```text46```text
47-// 单行注释 47+// 单行注释
48// 多行注释48// 多行注释
49// 第二行49// 第二行
50```50```
@@ -85,9 +85,8 @@ int Foo()
85在适当的时候,右边的注释上下对齐会更美观:85在适当的时候,右边的注释上下对齐会更美观:
86 86 
87```CPP87```CPP
88-const int A_CONST = 100; // 此处两行注释属于同类 88+const int A_CONST = 100; // 此处两行注释属于同类
89-const int ANOTHER_CONST = 200; // 可保持左侧对齐 89+const int ANOTHER_CONST = 200; // 可保持左侧对齐
90-```
91 90 
92## 【错误示例】91## 【错误示例】
93 92 
@@ -160,7 +159,7 @@ int y = 1; // 符合:注释位于代码上方
160版权许可内容解释如下:159版权许可内容解释如下:
161 160 
162- 2012-2018根据实际需要可以修改。161- 2012-2018根据实际需要可以修改。
163-- 2012是文件首次创建年份,而2018是最后文件修改年份。 162+- 2012是文件首次创建年份,而2018是最后文件修改年份。
164- 对文件有重大修改时,必须更新后面年份,如特性扩展,重大重构等。163- 对文件有重大修改时,必须更新后面年份,如特性扩展,重大重构等。
165- 可以只写一个创建年份,后续文件修改则不用更新版权声明。164- 可以只写一个创建年份,后续文件修改则不用更新版权声明。
166- 如:版权所有(c)华为技术有限公司2018165- 如:版权所有(c)华为技术有限公司2018
@@ -257,7 +256,7 @@ int Fun2();
257使用 /**/256使用 /**/
258 257 
259```CPP258```CPP
260-/* 单行函数头 */ 259+/* 单行函数头 */
261int Fun1();260int Fun1();
262 261 
263/*262/*
@@ -305,7 +304,7 @@ int Fun2();
305右大括号独占一行,除非后面跟着同一语句的剩余部分,如do语句中的while,或者if语句的else/else if,或者逗号、分号、小括号。304右大括号独占一行,除非后面跟着同一语句的剩余部分,如do语句中的while,或者if语句的else/else if,或者逗号、分号、小括号。
306 305 
307**Allman风格**306**Allman风格**
308-换行时,左大括另起并独占一行,保持与上一行相同缩进;307+换行时,左大括另起并独占一行,保持与上一行相同缩进;
309右大括号独占一行,除非后面跟着do语句中的while,或者逗号、分号。308右大括号独占一行,除非后面跟着do语句中的while,或者逗号、分号。
310 309 
311## 【修复建议】310## 【修复建议】
@@ -628,7 +627,7 @@ int FunctionName()
628 627 
629## 【描述】628## 【描述】
630 629 
631-类访问控制块的声明顺序默认依次是public:、protected:、 private:,缩进与class关键字对齐。630+类访问控制块的声明顺序默认依次是public:、protected:、private:,缩进与class关键字对齐。
632 631 
633## 【错误示例】632## 【错误示例】
634 633 
@@ -710,8 +709,8 @@ public: // 符合:和`class`对齐
710 709 
711## 【修复建议】710## 【修复建议】
712 711 
713-类访问控制块按顺序public:、protected:、 private:声明。712+类访问控制块按顺序public:、protected:、private:声明。
714-public:、protected:、 private:缩进和class关键字对齐。713+public:、protected:、private:缩进和class关键字对齐。
715在各个访问控制块中,建议将类似的声明放在一起,如果项目组没有制定声明顺序,可参考如下声明顺序:714在各个访问控制块中,建议将类似的声明放在一起,如果项目组没有制定声明顺序,可参考如下声明顺序:
716 715 
717- 类型(包括typedef,using和嵌套的结构体与类)716- 类型(包括typedef,using和嵌套的结构体与类)
@@ -786,7 +785,7 @@ SomeClass::SomeClass(int var)
786SomeClass::SomeClass(int var)785SomeClass::SomeClass(int var)
787 : someVar(var), // 缩进4个空格786 : someVar(var), // 缩进4个空格
788 someOtherVar(var + 1) // 与上一行的成员变量对齐787 someOtherVar(var + 1) // 与上一行的成员变量对齐
789-{ 788+{
790 DoSomething();789 DoSomething();
791}790}
792```791```
@@ -845,7 +844,7 @@ if (someConditions) {
845 ...844 ...
846} else if (...) { // 符合:else if与if在不同行845} else if (...) { // 符合:else if与if在不同行
847 ...846 ...
848-} else { // 符合:else与if在不同行 847+} else { // 符合:else与if在不同行
849 ...848 ...
850}849}
851```850```
@@ -864,11 +863,11 @@ case/default语句相对switch缩进一层,case中的语句相对case缩进一
864 863 
865```CPP864```CPP
866switch (var) {865switch (var) {
867-case CASE1: // 不符合:case未缩进 866+case CASE1: // 不符合:case未缩进
868 DoSomething1();867 DoSomething1();
869 break;868 break;
870...869...
871-default: // 不符合:default未缩进 870+default: // 不符合:default未缩进
872 break;871 break;
873}872}
874```873```
@@ -877,10 +876,10 @@ default: // 不符合:default未缩进
877 876 
878```CPP877```CPP
879switch (var) {878switch (var) {
880- case CASE1: // 符合:缩进 879+ case CASE1: // 符合:缩进
881- DoSomething1(); // 符合:缩进 880+ DoSomething1(); // 符合:缩进
882 break;881 break;
883- case CASE2: { // 符合:带大括号格式 882+ case CASE2: { // 符合:带大括号格式
884 DoSomething2();883 DoSomething2();
885 break;884 break;
886 }885 }
@@ -940,9 +939,10 @@ int & r2; // 不符合:两边都有空格
940```CPP939```CPP
941int &r; // 符合:"&"跟随变量名940int &r; // 符合:"&"跟随变量名
942struct Foo &CreateFoo(void); // 符合: "&"跟随函数名941struct Foo &CreateFoo(void); // 符合: "&"跟随函数名
943-场景2:选用的风格是跟随类型,但是实际开发时跟随变量名或函数名
944```942```
945 943 
944+场景2:选用的风格是跟随类型,但是实际开发时跟随变量名或函数名
945+ 
946修复示例:946修复示例:
947 947 
948```CPP948```CPP
@@ -1098,7 +1098,7 @@ const char* const VERSION = "V100"; // 符合:跟随类型(可选"*"两边
1098 1098 
1099```CPP1099```CPP
1100// 不符合:行尾有空格1100// 不符合:行尾有空格
1101-void Foo(int x); 1101+void Foo(int x);
1102```1102```
1103 1103 
1104场景2:小括号内两侧有空格1104场景2:小括号内两侧有空格
@@ -1181,7 +1181,7 @@ void Foo(int x, int y, int z);
1181```CPP1181```CPP
1182ret = DoSomething();1182ret = DoSomething();
1183 1183 
1184-if (ret != OK) { // 不符合:返回值判断应该紧跟函数调用 1184+if (ret != OK) { // 不符合:返回值判断应该紧跟函数调用
1185 return -1;1185 return -1;
1186}1186}
1187```1187```
@@ -1194,7 +1194,7 @@ int Foo()
1194 ...1194 ...
1195}1195}
1196 1196 
1197-int Bar() 1197+int Bar()
1198{1198{
1199 ...1199 ...
1200}1200}
@@ -27,7 +27,7 @@
27 27 
282. **文档修改**282. **文档修改**
29 29 
30- - 选择分支:请从master或其他Tag分支下载源码到本地。30+ - 选择分支:请从master分支或指定Tag下载源码到本地。
31 - 遵循格式:31 - 遵循格式:
32 - 本项目推荐使用**Markdown格式**32 - 本项目推荐使用**Markdown格式**
33 - 遵循项目既有的写作风格。33 - 遵循项目既有的写作风格。
@@ -42,7 +42,7 @@
42 42 
43 ```text43 ```text
44 简短说明(不超过50字符)44 简短说明(不超过50字符)
45- 45+ 
46 如有必要,在此处进行更详细描述。说明修改的原因和内容,而不是具体改了什么(代码本身会展示)。46 如有必要,在此处进行更详细描述。说明修改的原因和内容,而不是具体改了什么(代码本身会展示)。
47 关联的Issue: #12347 关联的Issue: #123
48 ```48 ```
@@ -76,15 +76,15 @@
76 - 确保所有代码示例可运行并经过测试。76 - 确保所有代码示例可运行并经过测试。
77 - 提供足够的上下文和解释。77 - 提供足够的上下文和解释。
78 - 注明代码运行所需的环境或前提条件。78 - 注明代码运行所需的环境或前提条件。
79- - 标点与格式: 79+ - 标点与格式:
80 - 中英文混排时,使用全角标点,标点符号必须符合中/英文语境。80 - 中英文混排时,使用全角标点,标点符号必须符合中/英文语境。
81- - 标题使用合适的层级(#、 ##、 ###)。81+ - 标题使用合适的层级(#、##、###)。
82- - 使用列表和表格来组织复杂信息。 82+ - 使用列表和表格来组织复杂信息。
83 - 链接:使用描述性链接文本,避免“点击这里”,确保链接资源真实可靠。83 - 链接:使用描述性链接文本,避免“点击这里”,确保链接资源真实可靠。
84 - 图片:84 - 图片:
85 - 常用格式:推荐png格式,风格尽量与已有图片保持一致。85 - 常用格式:推荐png格式,风格尽量与已有图片保持一致。
86 - 分辨率与清晰度:需清晰且尺寸适中,避免模糊或过度压缩。86 - 分辨率与清晰度:需清晰且尺寸适中,避免模糊或过度压缩。
87- - 文件大小:单张图片不建议超过10M 87+ - 文件大小:单张图片不建议超过10 MB
88 - 版权:所有引用的图片、文献等资源,请确保合规性。88 - 版权:所有引用的图片、文献等资源,请确保合规性。
89 89 
90## 获取帮助90## 获取帮助
@@ -119,7 +119,7 @@ add_example first input[7] is: 1.000000, second input[7] is: 1.000000, result[7]
119 119 
120### 1. 修改Kernel实现120### 1. 修改Kernel实现
121 121 
122-找到AddExample算子的核心kernel实现文件`ops-nn/examples/add_example/op_kernel/add_example.h`,尝试将算子中的Add操作改为Mul操作:122+找到AddExample算子的核心kernel实现文件`examples/add_example/op_kernel/add_example.h`,尝试将算子中的Add操作改为Mul操作:
123 123 
124```cpp124```cpp
125__aicore__ inline void AddExample<T>::Compute(int64_t currentNum)125__aicore__ inline void AddExample<T>::Compute(int64_t currentNum)
@@ -188,7 +188,7 @@ __aicore__ inline void AddExample<T>::Compute(int64_t currentNum)
188 188 
189 该接口支持打印Scalar类型数据,如整数、字符型、布尔型等,详细介绍请参见[《Ascend C API》](https://hiascend.com/document/redirect/CannCommunityAscendCApi)中“算子调测API > printf”。189 该接口支持打印Scalar类型数据,如整数、字符型、布尔型等,详细介绍请参见[《Ascend C API》](https://hiascend.com/document/redirect/CannCommunityAscendCApi)中“算子调测API > printf”。
190 190 
191- ```c++191+ ```cpp
192 blockLength_ = (tilingData->totalLength + AscendC::GetBlockNum() - 1) / AscendC::GetBlockNum();192 blockLength_ = (tilingData->totalLength + AscendC::GetBlockNum() - 1) / AscendC::GetBlockNum();
193 tileNum_ = tilingData->tileNum;193 tileNum_ = tilingData->tileNum;
194 tileLength_ = ((blockLength_ + tileNum_ - 1) / tileNum_ / BUFFER_NUM) ?194 tileLength_ = ((blockLength_ + tileNum_ - 1) / tileNum_ / BUFFER_NUM) ?
@@ -201,7 +201,7 @@ __aicore__ inline void AddExample<T>::Compute(int64_t currentNum)
201 201 
202 该接口支持Dump指定Tensor的内容,同时支持打印自定义附加信息,比如当前行号等,详细介绍请参见[《Ascend C API》](https://hiascend.com/document/redirect/CannCommunityAscendCApi)中“算子调测API > DumpTensor”。202 该接口支持Dump指定Tensor的内容,同时支持打印自定义附加信息,比如当前行号等,详细介绍请参见[《Ascend C API》](https://hiascend.com/document/redirect/CannCommunityAscendCApi)中“算子调测API > DumpTensor”。
203 203 
204- ```c++204+ ```cpp
205 AscendC::LocalTensor<T> zLocal = outputQueueZ.DeQue<T>();205 AscendC::LocalTensor<T> zLocal = outputQueueZ.DeQue<T>();
206 // 打印zLocal Tensor信息206 // 打印zLocal Tensor信息
207 DumpTensor(zLocal, 0, 128);207 DumpTensor(zLocal, 0, 128);
@@ -239,7 +239,7 @@ __aicore__ inline void AddExample<T>::Compute(int64_t currentNum)
239 239 
240**修改输入/输出数据**:修改输入、输出的shape信息,以及初始化数据,构造相应的输入、输出tensor。240**修改输入/输出数据**:修改输入、输出的shape信息,以及初始化数据,构造相应的输入、输出tensor。
241 241 
242-```c++242+```cpp
243int aclnnAddExampleTest(int32_t deviceId, aclrtStream& stream) {243int aclnnAddExampleTest(int32_t deviceId, aclrtStream& stream) {
244 // ... 初始化代码 ...244 // ... 初始化代码 ...
245 245 
@@ -1,7 +1,7 @@
1# aclnn返回码1# aclnn返回码
2 2 
3调用aclnn API时,常见的接口返回码如[表1](#table1)所示。3调用aclnn API时,常见的接口返回码如[表1](#table1)所示。
4-对于异常状态码值,可以通过aclGetRecentErrMsg接口(参见[《Runtime运行时API》](https://hiascend.com/document/redirect/CannCommunityRuntimeApi))获取异常信息,根据报错提示排查问题或联系技术支持。4+对于异常状态码值,可以通过aclGetRecentErrMsg接口(参见[《Runtime运行时API》](https://hiascend.com/document/redirect/CannCommunityRuntimeApi))获取异常信息,可根据报错提示排查问题或联系技术支持。
5 5 
6**表1** 返回状态码 <a id="table1"></a>6**表1** 返回状态码 <a id="table1"></a>
7 7 
@@ -7,11 +7,11 @@
7 7 
8## 编译前准备8## 编译前准备
9 9 
10-本章以开发和运行环境合设场景为例,即带AI处理器的机器既作为开发环境又作为运行环境。该场景下,代码开发和代码运行在同一台机器上。这里以**AddMatMul算子**为例,其他算子的调用逻辑、流程、编译脚本与AddMatMul算子大致一样,请根据实际情况自行修改API调用脚本(\*.cpp)和编译脚本(CMakeLists)。10+本章以开发和运行环境合设场景为例,即带AI处理器的机器既作为开发环境又作为运行环境。该场景下,代码开发和代码运行在同一台机器上。这里以**AddMatMul算子**为例,其他算子的调用逻辑、流程、编译脚本与AddMatMul算子大致一样,请根据实际情况自行修改API调用脚本(\*.cpp)和编译脚本(CMakeLists.txt)。
11 11 
12- **示例代码**12- **示例代码**
13 13 
14- 已知AddMatMul算子实现了张量加法运算,计算公式为:out = β self + α (mat1 @ mat2)。您可以从[aclnnAddmm&aclnnInplaceAddmm.md](../../../matmul/mat_mul_v3/docs/aclnnAddmm&aclnnInplaceAddmm.md)中“调用示例”获取示例代码,并将代码文件命名为“**test\_addmm.cpp**”。14+ 已知AddMatMul算子实现了张量加法运算,计算公式为:out = β self + α (mat1 @ mat2)。您可以从[aclnnAddmm&aclnnInplaceAddmm.md](../../../matmul/mat_mul_v3/docs/aclnnAddmm&aclnnInplaceAddmm.md)中“调用示例”获取示例代码,并将代码文件命名为“**test\_addmm.cpp**”。
15 15 
16- **CMakeLists文件**16- **CMakeLists文件**
17 17 
@@ -138,7 +138,7 @@
138 CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnAddmmGetWorkspaceSize failed. ERROR: %d\n[ERROR msg]%s", ret, aclGetRecentErrMsg()); return ret);138 CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnAddmmGetWorkspaceSize failed. ERROR: %d\n[ERROR msg]%s", ret, aclGetRecentErrMsg()); return ret);
139 ```139 ```
140 140 
141- 上述构造空指针问题获取报错信息示例如下:141+ 上述构造空指针问题获取报错信息示例如下
142 142 
143 ```bash143 ```bash
144 aclnnAddmmGetWorkspaceSize failed. ERROR: 161001144 aclnnAddmmGetWorkspaceSize failed. ERROR: 161001
@@ -8,7 +8,7 @@
8 8 
9## 使用说明9## 使用说明
10 10 
11-目前大部分算子API都是支持ND数据格式的。例如aclnnAdd接口,均标明支持的数据格式是ND(即多维Tensor,低维度优先连续排布的规则)。对于aclnnConvolution,其属于CNN类API,要求输入的aclTensor设置具有业务语义的格式,而不是ND格式。此类算子,在计算过程中需要知道Tensor中的业务语义才可以进行对应的计算。例如2D卷积中,需要知道Batch维度、Channel维度、Height维度、Width维度与Tensor维度的对应关系是什么。11+目前大部分算子API都是支持ND数据格式的。例如aclnnAdd接口,均标明支持的数据格式是ND(即多维Tensor,低维度优先连续排布的规则)。对于aclnnConvolution,其属于CNN类API,要求输入的aclTensor设置具有业务语义的格式,而不是ND格式。此类算子,在计算过程中需要知道Tensor中的业务语义才可以进行对应的计算。例如2D卷积中,需要知道Batch维度、Channel维度、Height维度、Width维度与Tensor维度的对应关系是什么。
12 12 
13>**说明:**13>**说明:**
14>14>
@@ -2,7 +2,7 @@
2 2 
3目前大部分算子API的输入aclTensor支持“**非连续的Tensor**”,即一个Tensor可以通过\(shape, strides, offset\)表示。3目前大部分算子API的输入aclTensor支持“**非连续的Tensor**”,即一个Tensor可以通过\(shape, strides, offset\)表示。
4 4 
5-说明:创建aclTensor可通过[《算子库接口》](https://hiascend.com/document/redirect/CannCommunityOplist)中公共接口>aclCreateTensor实现。5+说明:创建aclTensor可通过[《算子库接口》](https://hiascend.com/document/redirect/CannCommunityOplist)中"公共接口>aclCreateTensor"实现。
6 6 
7## 示例17## 示例1
8 8 
@@ -104,10 +104,10 @@ cannsim record [options] user_app
1042. 执行仿真命令,可参考以下使用示例1042. 执行仿真命令,可参考以下使用示例
105 105 
106 ```bash106 ```bash
107- 方式一:启用仿真,并将输出保存至 ./output目录,/path/to/app为算子程序107+ # 方式一:启用仿真,并将输出保存至 ./output目录,/path/to/app为算子程序
108 $ cannsim record /path/to/app -o ./output -s Ascend950108 $ cannsim record /path/to/app -o ./output -s Ascend950
109 109 
110- 方式二:启用仿真并生成报告,用于后续性能分析110+ # 方式二:启用仿真并生成报告,用于后续性能分析
111 $ cannsim record /path/to/app -o ./output -s Ascend950 --gen-report111 $ cannsim record /path/to/app -o ./output -s Ascend950 --gen-report
112 ```112 ```
113 113 
@@ -154,12 +154,12 @@ cannsim report [options]
154 154 
155## 使用示例155## 使用示例
156 156 
157-1. 参考仿真执行执行算子仿真,对比输出示例,确保对应的结果执行正确。157+1. 参考仿真执行执行算子仿真,对比输出示例,确保对应的结果执行正确。
1582. 执行仿真结果解析命令,可参考以下执行用例。1582. 执行仿真结果解析命令,可参考以下执行用例。
159 159 
160 ```bash160 ```bash
161 在当前目录下生成性能分析报告(默认仅分析核0)161 在当前目录下生成性能分析报告(默认仅分析核0)
162- cannsim report -e /path/to/cannsim_{timestamp}_${user_app} 162+ cannsim report -e /path/to/cannsim_{timestamp}_${user_app}
163 163 
164 在指定目录下生成核0、核1、核11、核12的性能分析报告164 在指定目录下生成核0、核1、核11、核12的性能分析报告
165 cannsim report -e /path/to/cannsim_{timestamp}_${user_app} -o /path/to/report -n '0-1, 11-12'165 cannsim report -e /path/to/cannsim_{timestamp}_${user_app} -o /path/to/report -n '0-1, 11-12'
@@ -210,11 +210,11 @@ cannsim --help
210```bash210```bash
211cannsim record --help211cannsim record --help
212```212```
213- 213+ 
214查询工具report子命令的帮助信息:214查询工具report子命令的帮助信息:
215 215 
216 ```bash216 ```bash
217- cannsim report --help 217+ cannsim report --help
218 ```218 ```
219 219 
220## 参数说明220## 参数说明
@@ -8,13 +8,13 @@
8 8 
9* **plog获取**9* **plog获取**
10 10 
11- 程序执行结束后,默认可在"$HOME/ascend/log"下查看,host日志文件存储路径如下:11+ 程序执行结束后,默认可在"$HOME/ascend/log/debug/plog"下查看,host日志文件存储路径如下:
12 12 
13 ```bash13 ```bash
14 $HOME/ascend/log/debug/plog/plog-pid_*.log14 $HOME/ascend/log/debug/plog/plog-pid_*.log
15 ```15 ```
16 16 
17- 开启环境变量ASCEND_SLOG_PRINT_TO_STDOUT可以将log日志直接打屏显示(1:开启打屏,0:关闭打屏),配置示例如下:17+ 开启环境变量ASCEND_SLOG_PRINT_TO_STDOUT可以将log日志直接打屏显示(1开启打屏,0:关闭打屏),配置示例如下:
18 18 
19 ```bash19 ```bash
20 export ASCEND_SLOG_PRINT_TO_STDOUT=120 export ASCEND_SLOG_PRINT_TO_STDOUT=1
@@ -26,8 +26,8 @@
26 26 
27 通过aclGetRecentErrMsg接口(参见[《Runtime运行时API》](https://hiascend.com/document/redirect/CannCommunityRuntimeApi))获取aclnn接口调用过程中的异常信息,使用方法如下:27 通过aclGetRecentErrMsg接口(参见[《Runtime运行时API》](https://hiascend.com/document/redirect/CannCommunityRuntimeApi))获取aclnn接口调用过程中的异常信息,使用方法如下:
28 28 
29- ```bash29+ ```cpp
30- printf(aclGetRecentErrMsg());30+ printf("%s",aclGetRecentErrMsg());
31 ```31 ```
32 32 
33 打印错误信息样例如下:33 打印错误信息样例如下:
@@ -81,7 +81,7 @@
81 81 
82 可通过如下宏打印算子执行过程中的日志信息,包括DEBUG、INFO、WARN、ERROR级别日志。82 可通过如下宏打印算子执行过程中的日志信息,包括DEBUG、INFO、WARN、ERROR级别日志。
83 83 
84- ```Cpp84+ ```c++
85 KERNEL_LOG_DEBUG(fmt, …) // fmt参数表示格式控制字符串85 KERNEL_LOG_DEBUG(fmt, …) // fmt参数表示格式控制字符串
86 KERNEL_LOG_INFO(fmt, …)86 KERNEL_LOG_INFO(fmt, …)
87 KERNEL_LOG_WARN(fmt, …)87 KERNEL_LOG_WARN(fmt, …)
@@ -256,8 +256,8 @@ graph LR
256Kernel一共需要两个交付件:```${op_name}.cpp``` ```${op_name}.h```256Kernel一共需要两个交付件:```${op_name}.cpp``` ```${op_name}.h```
257> 说明:257> 说明:
258>258>
259-> 1. `${op_name}.cpp`为kernel的入口函数只能放在`${op_name}/op_kernel`目录下;259+> 1. ```${op_name}.cpp```为kernel的入口函数只能放在`${op_name}/op_kernel`目录下;
260-> 2. `${op_name}.h`文件可以按照不同SoC或模板放在对应目录下,例如:`${op_name}/op_kernel/arch32`、`${op_name}/op_kernel/arch35`或`${op_name}/op_kernel/impl`等目录下;260+> 2. ```${op_name}.h```文件可以按照不同SoC或模板放在对应目录下,例如:`${op_name}/op_kernel/arch32`、`${op_name}/op_kernel/arch35`或`${op_name}/op_kernel/impl`等目录下;
261 261 
262**交付件1:${op_name}.cpp**262**交付件1:${op_name}.cpp**
263 263 
@@ -349,7 +349,7 @@ __aicore__ inline void AddExample<T>::Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, cons
349 blockLength_ = tilingData->totalLength / AscendC::GetBlockNum();349 blockLength_ = tilingData->totalLength / AscendC::GetBlockNum();
350 ...350 ...
351 // 3.2初始化GM地址351 // 3.2初始化GM地址
352- inputGMX.SetGlobalBuffer((__gm__ T*)x + blockLength_ * AscendC::GetBlockIdx(), blockLength_);352+ inputGMX_.SetGlobalBuffer((__gm__ T*)x + blockLength_ * AscendC::GetBlockIdx(), blockLength_);
353 ...353 ...
354 // 3.3初始化队列长度354 // 3.3初始化队列长度
355 pipe.InitBuffer(inputQueueX_, BUFFER_NUM, tileLength_ * sizeof(T));355 pipe.InitBuffer(inputQueueX_, BUFFER_NUM, tileLength_ * sizeof(T));
@@ -380,7 +380,7 @@ __aicore__ inline void AddExample<T>::Process()
380`scripts/kernel/binary_config`目录[ascendc_config.json](../../../scripts/kernel/binary_config/ascendc_config.json)中,注册算子的NPU型号和实现模式,示例如下,输入实际name和compute_units即可。380`scripts/kernel/binary_config`目录[ascendc_config.json](../../../scripts/kernel/binary_config/ascendc_config.json)中,注册算子的NPU型号和实现模式,示例如下,输入实际name和compute_units即可。
381 381 
382```json382```json
383-{"name":"AddExample", "compute_units": ["${soc_version}"], "auto_sync":true, "impl_mode" : "high_performance"},383+{"name":"AddExample", "compute_units": ["${soc_version}"], "auto_sync":true, "impl_mode" : "high_performance"}
384```384```
385 385 
386## 编译部署386## 编译部署
@@ -459,7 +459,7 @@ __aicore__ inline void AddExample<T>::Process()
459 459 
460### UT验证460### UT验证
461 461 
462-主要交付件代码开发过程中,可通过UT验证方式进行快速验证,无需编译部署算子包。462+主要交付件代码开发过程中,可通过UT验证方式进行快速验证,无需编译部署算子包。
463 463 
464UT目录结构如下,需用户手动创建:464UT目录结构如下,需用户手动创建:
465 465 
@@ -24,7 +24,7 @@
24 24 
256. [编译部署](#编译部署):通过工程编译脚本完成自定义算子的编译和安装。256. [编译部署](#编译部署):通过工程编译脚本完成自定义算子的编译和安装。
26 26 
27-7. [算子验证](#算子验证):通过常见算子调用方式,验证自定义算子功能。 27+7. [算子验证](#算子验证):通过常见算子调用方式,验证自定义算子功能。
28 28 
29## 工程创建29## 工程创建
30 30 
@@ -62,7 +62,7 @@ ${op_name} # 替换为实际算子名的小写下
62└── CMakeLists.txt # 算子cmakelist入口62└── CMakeLists.txt # 算子cmakelist入口
63```63```
64 64 
65- 若```${op_class}```为全新算子分类需额外在`CMakeLists`中添加```add_subdirectory(${op_class})```,否则无法正常编译。65+ 若`${op_class}`为全新算子分类需额外在`CMakeLists`中添加`add_subdirectory(${op_class})`,否则无法正常编译。
66 66 
67 ```bash67 ```bash
68 if(ENABLE_EXPERIMENTAL)68 if(ENABLE_EXPERIMENTAL)
@@ -78,7 +78,7 @@ ${op_name} # 替换为实际算子名的小写下
78 78 
79## 算子定义79## 算子定义
80 80 
81-算子定义需要完成两个交付件:`README.md` ```${op_name}.json```81+算子定义需要完成两个交付件:`README.md` `${op_name}.json`
82 82 
83**交付件1:README.md**83**交付件1:README.md**
84 84 
@@ -106,13 +106,13 @@ graph LR
106 106 
107### 代码实现107### 代码实现
108 108 
109-Kernel一共需要两个交付件:```${op_name}_aicpu.cpp``` ```${op_name}_aicpu.h```109+Kernel一共需要两个交付件:`${op_name}_aicpu.cpp` `${op_name}_aicpu.h`
110 110 
111**交付件1:${op_name}_aicpu.h**111**交付件1:${op_name}_aicpu.h**
112 112 
113算子类声明113算子类声明
114 114 
115-Kernel实现的第一步,需在头文件```op_kernel_aicpu/${op_name}_aicpu.h```进行算子类的声明,算子类需继承CpuKernel基类。115+Kernel实现的第一步,需在头文件`op_kernel_aicpu/${op_name}_aicpu.h`进行算子类的声明,算子类需继承CpuKernel基类。
116如需查看详细实现,请参考[add_example_aicpu.h](../../../examples/add_example_aicpu/op_kernel_aicpu/add_example_aicpu.h)。116如需查看详细实现,请参考[add_example_aicpu.h](../../../examples/add_example_aicpu/op_kernel_aicpu/add_example_aicpu.h)。
117 117 
118```CPP118```CPP
@@ -180,9 +180,9 @@ uint32_t AddExampleCpuKernel::Compute(CpuKernelContext& ctx) {
180 case DT_FLOAT:180 case DT_FLOAT:
181 return AddCompute<float>(...);181 return AddCompute<float>(...);
182 case DT_INT32:182 case DT_INT32:
183- return AddCompute<int32>(...);183+ return AddCompute<int32_t>(...);
184- ....184+ ...
185- default : return PARAM_INVALID;185+ default : return kParamInvalid;
186 }186 }
187}187}
188 188 
@@ -231,14 +231,14 @@ REGISTER_CPU_KERNEL(kAddExample, AddExampleCpuKernel);
231 ./build_out/cann-ops-nn-${vendor_name}_linux-${arch}.run231 ./build_out/cann-ops-nn-${vendor_name}_linux-${arch}.run
232 ```232 ```
233 233 
234- 自定义算子包安装在```${ASCEND_HOME_PATH}/opp/vendors```路径中,```${ASCEND_HOME_PATH}```表示CANN软件安装目录,可提前在环境变量中配置。234+ 自定义算子包安装在`${ASCEND_HOME_PATH}/opp/vendors`路径中,`${ASCEND_HOME_PATH}`表示CANN软件安装目录,可提前在环境变量中配置。
235- 235+ 
2364. **(可选)卸载自定义算子包。**2364. **(可选)卸载自定义算子包。**
237 237 
238- 自定义算子包安装后在```${ASCEND_HOME_PATH}/opp/vendors/custom_nn/scripts```目录会生成`uninstall.sh`,通过该脚本可卸载自定义算子包,命令如下:238+ 自定义算子包安装后在`${ASCEND_HOME_PATH}/opp/vendors/${vendor_name}_nn/scripts`目录会生成`uninstall.sh`,通过该脚本可卸载自定义算子包,命令如下:
239- 239+ 
240 ```bash240 ```bash
241- bash ${ASCEND_HOME_PATH}/opp/vendors/custom_nn/scripts/uninstall.sh241+ bash ${ASCEND_HOME_PATH}/opp/vendors/${vendor_name}_nn/scripts/uninstall.sh
242 ```242 ```
243 243 
244## 算子验证244## 算子验证
@@ -82,7 +82,7 @@
82 <td>可用ND2NZ/DN2NZ在MTE2阶段完成格式转换,减少中间buffer和格式转换开销;需关注步长、对齐与NZ形状映射</td>82 <td>可用ND2NZ/DN2NZ在MTE2阶段完成格式转换,减少中间buffer和格式转换开销;需关注步长、对齐与NZ形状映射</td>
83 </tr>83 </tr>
84 <tr>84 <tr>
85- <td>支持Cube-&gt;Vector高效内部数据通路:L1&lt;-&gt;UB、L0C-&gt;UB、FIXP-&gt;UB</td>85+ <td>支持Cube-&gt;Vector高效内部数据通路L1&lt;-&gt;UB、L0C-&gt;UB、FIXP-&gt;UB</td>
86 <td>可在UB侧做中间累加/激活/融合(如切K累加、后处理),减少GM往返;对应同步与管线切分需调整</td>86 <td>可在UB侧做中间累加/激活/融合(如切K累加、后处理),减少GM往返;对应同步与管线切分需调整</td>
87 </tr>87 </tr>
88 <tr>88 <tr>
@@ -251,7 +251,7 @@ __simd_vf__ __aicore__ void GenIndexBuf(ubuf int32_t* helpAddr, int32_t colFacto
251 251 
252```cpp252```cpp
253// 动态掩码:处理尾部不完整数据253// 动态掩码:处理尾部不完整数据
254-__simd_vf__ __aicore__ void GatherProcess(ubuf int8_t* curYAddr, uint16_t repeatimes, uint16_t computeSize)254+__simd_vf__ __aicore__ void GatherProcess(ubuf int8_t* curYAddr, uint16_t repeatTimes, uint16_t computeSize)
255{255{
256 MicroAPI::RegTensor<int8_t> vregTemp;256 MicroAPI::RegTensor<int8_t> vregTemp;
257 MicroAPI::MaskReg preg;257 MicroAPI::MaskReg preg;
@@ -343,7 +343,7 @@ Ascend 950新架构引入UB2L1 & L0C2UB间的直连通路,实现矩阵计算
343```cpp343```cpp
344// 1. 新增: 搬入接口增加UB2L1的Nd2Nz搬入,支持Src&Dst都是LocalTensor的形式344// 1. 新增: 搬入接口增加UB2L1的Nd2Nz搬入,支持Src&Dst都是LocalTensor的形式
345template <typename T>  345template <typename T>  
346-__aicore__ inline void DataCopy(const LocalTensor<T>& dst, const LocalTensor<T>& src, const Nd2NzParams& intriParams)346+__aicore__ inline void DataCopy(const LocalTensor<T>& dst, const LocalTensor<T>& src, const Nd2NzParams& intriParams);
347 347 
348// 2. 新增: 搬出接口增加L0C2UB的搬出,支持直接从L0C搬出到UB,支持Src&Dst都是LocalTensor的形式348// 2. 新增: 搬出接口增加L0C2UB的搬出,支持直接从L0C搬出到UB,支持Src&Dst都是LocalTensor的形式
349template <typename T, typename U, const FixpipeConfig& config = CFG_ROW_MAJOR>349template <typename T, typename U, const FixpipeConfig& config = CFG_ROW_MAJOR>
@@ -386,7 +386,7 @@ Ascend 950引入集合通信加速器CCU1.0,降低了访存需求,减少了
386以[MatmulAllReduce](https://gitcode.com/cann/ops-transformer/tree/master/mc2/matmul_all_reduce)算子迁移适配为例:386以[MatmulAllReduce](https://gitcode.com/cann/ops-transformer/tree/master/mc2/matmul_all_reduce)算子迁移适配为例:
387设置NnopbaseSetHcclServerType枚举值,A2为NNOPBASE_HCCL_SERVER_AICPU,950为NNOPBASE_HCCL_SERVER_TYPE_CCU。387设置NnopbaseSetHcclServerType枚举值,A2为NNOPBASE_HCCL_SERVER_AICPU,950为NNOPBASE_HCCL_SERVER_TYPE_CCU。
388 388 
389-```CPP389+```cpp
390// ...390// ...
391aclnnStatus aclnnMatmulAllReduce(391aclnnStatus aclnnMatmulAllReduce(
392 void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, const aclrtStream stream)392 void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, const aclrtStream stream)
@@ -18,7 +18,7 @@ ${op_name} # 替换为实际算子名的小写下
18 18 
19## Shape与DataType推导19## Shape与DataType推导
20 20 
21-图模式需要完成两个交付件```${op_name}_graph_infer.cpp``` ```${op_name}_infershape.cpp```21+图模式需要完成两个交付件`${op_name}_graph_infer.cpp` `${op_name}_infershape.cpp`
22 22 
23**交付件1:${op_name}_infershape.cpp**23**交付件1:${op_name}_infershape.cpp**
24 24 
@@ -107,7 +107,7 @@ IMPL_OP(AddExample).InferDataType(InferDataTypeAddExample);
107 107 
108示例代码如下,展示了如何注册`AddExample`算子:108示例代码如下,展示了如何注册`AddExample`算子:
109 109 
110-```CPP110+```cpp
111REG_OP(AddExample)111REG_OP(AddExample)
112 .INPUT(x1, TensorType({DT_FLOAT}))112 .INPUT(x1, TensorType({DT_FLOAT}))
113 .INPUT(x2, TensorType({DT_FLOAT}))113 .INPUT(x2, TensorType({DT_FLOAT}))
@@ -4,7 +4,7 @@
4 4 
5本文档介绍如何使用Ascend C和[PyTorch Extension](https://docs.pytorch.org/tutorials/extension.html)能力开发自定义NPU算子。通过PyTorch Extension机制,可以将Ascend C开发的算子封装成PyTorch可调用的扩展包,实现与PyTorch原生操作一致的使用体验。5本文档介绍如何使用Ascend C和[PyTorch Extension](https://docs.pytorch.org/tutorials/extension.html)能力开发自定义NPU算子。通过PyTorch Extension机制,可以将Ascend C开发的算子封装成PyTorch可调用的扩展包,实现与PyTorch原生操作一致的使用体验。
6 6 
7-**核心优势:**7+**核心优势**
8 8 
9- **单交付件:** 一个文件完成算子开发和PyTorch框架适配。9- **单交付件:** 一个文件完成算子开发和PyTorch框架适配。
10- **高效调用:** 使用`<<<>>>`语法启动核函数,流程简单高效。10- **高效调用:** 使用`<<<>>>`语法启动核函数,流程简单高效。
@@ -99,6 +99,7 @@ __global__ __aicore__ void add_kernel(GM_ADDR x, GM_ADDR y, GM_ADDR z, int64_t t
99 99 
100/**100/**
101 * 算子NPU调用实现101 * 算子NPU调用实现
102+ * 具体实现参考(./examples/fast_kernel_launch_example/csrc/add/ascend910b/add.cpp)文件
102 */103 */
103torch::Tensor add_npu(const torch::Tensor &x, const torch::Tensor &y)104torch::Tensor add_npu(const torch::Tensor &x, const torch::Tensor &y)
104{105{
@@ -155,7 +156,7 @@ add_sources()
155### 安装依赖156### 安装依赖
156 157 
157```bash158```bash
158-pip install build pyyaml 'numpy<2' pytest159+pip install build pyyaml "numpy<2" pytest
159```160```
160 161 
161### 编译命令162### 编译命令
@@ -163,7 +164,7 @@ pip install build pyyaml 'numpy<2' pytest
163`ops-nn`目录下执行以下命令进行编译打包:164`ops-nn`目录下执行以下命令进行编译打包:
164 165 
165```bash166```bash
166-bash build.sh --pkg --experimental --soc=${soc} --ops=${op1,op2}167+bash build.sh --pkg --experimental --soc=${soc} --ops=${ops}
167```168```
168 169 
169**参数说明:**170**参数说明:**
@@ -229,5 +230,5 @@ cpu_y = y.cpu()
229cpu_result = cpu_x + cpu_y230cpu_result = cpu_x + cpu_y
230 231 
231assert torch.allclose(cpu_result, result.cpu(), rtol=1e-6)232assert torch.allclose(cpu_result, result.cpu(), rtol=1e-6)
232-print("验证成功")233+print("验证成功!")
233```234```
@@ -40,7 +40,7 @@ bash build.sh --help
40| -O${n} | 可选 | 指定编译优化级别,支持O0/O1/O2/O3(如:-O3),${n}为优化级别标识。 |40| -O${n} | 可选 | 指定编译优化级别,支持O0/O1/O2/O3(如:-O3),${n}为优化级别标识。 |
41| -u | 可选 | 启用单元测试(UT)编译模式,编译所有UT目标。 |41| -u | 可选 | 启用单元测试(UT)编译模式,编译所有UT目标。 |
42| --help, -h | 可选 | 打印脚本使用帮助信息。 |42| --help, -h | 可选 | 打印脚本使用帮助信息。 |
43-| --ops | 可选 | 指定待编译的算子,如:mat_mul_v3mse_loss,多个算子用英文逗号“,”分隔,不可与--ophost、--opapi同时使用。 |43+| --ops | 可选 | 指定待编译的算子,如:mat_mul_v3,mse_loss,多个算子用英文逗号“,”分隔,不可与--ophost、--opapi同时使用。 |
44| --soc | 可选 | 指定NPU型号,每次编译只支持1个NPU型号。 |44| --soc | 可选 | 指定NPU型号,每次编译只支持1个NPU型号。 |
45| --jit | 可选 | 静态图场景下,编译`cann-${soc_name}-ops-nn_${cann_version}_linux-${arch}.run`整包时不需要编译算子二进制文件(图的运行态会在线编译),可以配置该选项,以提升编译速度。 |45| --jit | 可选 | 静态图场景下,编译`cann-${soc_name}-ops-nn_${cann_version}_linux-${arch}.run`整包时不需要编译算子二进制文件(图的运行态会在线编译),可以配置该选项,以提升编译速度。 |
46| --static | 可选 | 配置后,表示生成静态库文件,包含libcann_nn_static.a和aclnn接口头文件,搭配--pkg参数,生成静态库压缩包。|46| --static | 可选 | 配置后,表示生成静态库文件,包含libcann_nn_static.a和aclnn接口头文件,搭配--pkg参数,生成静态库压缩包。|
@@ -89,7 +89,7 @@
89├── docs # 项目相关文档目录89├── docs # 项目相关文档目录
90├── examples # 端到端算子开发和调用示例90├── examples # 端到端算子开发和调用示例
91│ ├── add_example # AI Core算子示例目录91│ ├── add_example # AI Core算子示例目录
92-│ │ ├── CMakeLists.txt # 算子编译配置文件 92+│ │ ├── CMakeLists.txt # 算子编译配置文件
93│ │ ├── examples # 算子使用示例目录93│ │ ├── examples # 算子使用示例目录
94│ │ ├── op_graph # 算子构图相关目录94│ │ ├── op_graph # 算子构图相关目录
95│ │ ├── op_host # 算子信息库、Tiling、InferShape相关实现目录95│ │ ├── op_host # 算子信息库、Tiling、InferShape相关实现目录
@@ -14,8 +14,8 @@
14| 安装方式 | 使用说明 | 使用场景 |14| 安装方式 | 使用说明 | 使用场景 |
15| ----- | ------ | ------ |15| ----- | ------ | ------ |
16| CANNLab | 一站式开发平台,提供在线直接运行的昇腾环境,无需手动安装。<br>当前可提供单机算力,**默认安装最新版本CANN包**。 | 适用于没有昇腾设备的开发者。|16| CANNLab | 一站式开发平台,提供在线直接运行的昇腾环境,无需手动安装。<br>当前可提供单机算力,**默认安装最新版本CANN包**。 | 适用于没有昇腾设备的开发者。|
17-| Docker | Docker镜像是一种高效部署方式,已预集成CANN包和必备依赖。<br>当前仅适用于Atlas A2系列产品,OS仅支持Ubuntu操作系统。**默认安装最新版本CANN包**。 |适用有昇腾设备,需要快速搭建环境的开发者。|17+| Docker | Docker镜像是一种高效部署方式,已预集成CANN包和必备依赖。<br>当前仅适用于Atlas A2系列产品,OS仅支持Ubuntu操作系统。**默认安装最新版本CANN包**。 |适用有昇腾设备,需要快速搭建环境的开发者。|
18-| 手动安装 | 手动安装CANN包和基础依赖,灵活性高。 |适用有昇腾设备,想体验手动安装CANN包或体验最新master分支能力的开发者。|18+| 手动安装 | 手动安装CANN包和基础依赖,灵活性高。 |适用有昇腾设备,想体验手动安装CANN包或体验最新master分支能力的开发者。|
19 19 
20### 方式1:CANNLab20### 方式1:CANNLab
21 21 
@@ -146,7 +146,7 @@
146 146 
1472. 安装依赖。1472. 安装依赖。
148 148 
149- 先通过项目根目录install\_deps.sh一键安装上述依赖,命令如下,若遇到不支持系统,请参考该文件自行适配。149+ 先通过项目根目录install\_deps.sh一键安装上述依赖,命令如下,若遇到不支持系统,请参考该文件自行适配。
150 150 
151 ```bash151 ```bash
152 bash install_deps.sh152 bash install_deps.sh
@@ -84,7 +84,7 @@
84 ```bash84 ```bash
85 # 静态库文件路径85 # 静态库文件路径
86 static_lib_path=""86 static_lib_path=""
87- 87+ 
88 # 环境变量生效88 # 环境变量生效
89 if [ -n "$ASCEND_INSTALL_PATH" ]; then89 if [ -n "$ASCEND_INSTALL_PATH" ]; then
90 _ASCEND_INSTALL_PATH=$ASCEND_INSTALL_PATH90 _ASCEND_INSTALL_PATH=$ASCEND_INSTALL_PATH
@@ -93,9 +93,9 @@
93 else93 else
94 _ASCEND_INSTALL_PATH="/usr/local/Ascend/cann"94 _ASCEND_INSTALL_PATH="/usr/local/Ascend/cann"
95 fi95 fi
96- 96+ 
97 source ${_ASCEND_INSTALL_PATH}/bin/setenv.bash97 source ${_ASCEND_INSTALL_PATH}/bin/setenv.bash
98- 98+ 
99 # 编译可执行文件99 # 编译可执行文件
100 g++ test_aclnn_transpose_batch_mat_mul.cpp \100 g++ test_aclnn_transpose_batch_mat_mul.cpp \
101 -I ${static_lib_path}/include \101 -I ${static_lib_path}/include \
@@ -108,12 +108,12 @@
108 -lpthread -lmmpa -lmetadef -lascendalog -lregister -lopp_registry -lops_base -lascendcl -ltiling_api -lplatform \108 -lpthread -lmmpa -lmetadef -lascendalog -lregister -lopp_registry -lops_base -lascendcl -ltiling_api -lplatform \
109 -ldl -lc_sec -lnnopbase -lruntime -lerror_manager -lunified_dlog \109 -ldl -lc_sec -lnnopbase -lruntime -lerror_manager -lunified_dlog \
110 -o test_aclnn_transpose_batch_mat_mul # 替换为实际算子可执行文件名110 -o test_aclnn_transpose_batch_mat_mul # 替换为实际算子可执行文件名
111- 111+ 
112 # 执行程序112 # 执行程序
113 ./test_aclnn_transpose_batch_mat_mul113 ./test_aclnn_transpose_batch_mat_mul
114 ```114 ```
115 115 
116- \$\{static_lib_path\}表示静态库统一放置路径;\$\{ASCEND_INSTALL_PATH\}已通过环境变量配置,表示CANN toolkit包安装路径;最终可执行文件名请替换为**实际算子可执行文件名**。 116+ \$\{static_lib_path\}表示静态库统一放置路径;\$\{ASCEND_INSTALL_PATH\}已通过环境变量配置,表示CANN toolkit包安装路径;最终可执行文件名请替换为**实际算子可执行文件名**。
117 117 
118 其中lcann_nn_static、lcann_math_static、lcann_legacy_static表示算子依赖的静态库文件,从静态库统一放置路径\$\{static_lib_path\}中获取;118 其中lcann_nn_static、lcann_math_static、lcann_legacy_static表示算子依赖的静态库文件,从静态库统一放置路径\$\{static_lib_path\}中获取;
119 lgraph、lmetadef等表示算子依赖的底层库文件,可在CANN toolkit包获取。119 lgraph、lmetadef等表示算子依赖的底层库文件,可在CANN toolkit包获取。
@@ -241,7 +241,7 @@
241 {241 {
242 int32_t deviceId = 0;242 int32_t deviceId = 0;
243 aclrtStream stream;243 aclrtStream stream;
244- 244+ 
245 auto ret = aclnnAddExampleTest(deviceId, stream);245 auto ret = aclnnAddExampleTest(deviceId, stream);
246 // 释放device资源以及acl去初始化246 // 释放device资源以及acl去初始化
247 aclrtDestroyStream(stream);247 aclrtDestroyStream(stream);
@@ -260,7 +260,7 @@
260 260 
261 - **调用自定义算子**:依赖自定义算子包261 - **调用自定义算子**:依赖自定义算子包
262 262 
263- ```bash263+ ```cmake
264 cmake_minimum_required(VERSION 3.14)264 cmake_minimum_required(VERSION 3.14)
265 # 设置工程名265 # 设置工程名
266 project(ACLNN_EXAMPLE)266 project(ACLNN_EXAMPLE)
@@ -269,18 +269,18 @@
269 add_compile_options(-std=c++11)269 add_compile_options(-std=c++11)
270 270 
271 # 设置编译输出目录为当前目录下的bin文件夹271 # 设置编译输出目录为当前目录下的bin文件夹
272- set(CMAKE_RUNTIME_OUTPUT_DIRECTORY "./bin") 272+ set(CMAKE_RUNTIME_OUTPUT_DIRECTORY "./bin")
273 273 
274 # 设置调试和发布模式的编译选项274 # 设置调试和发布模式的编译选项
275 set(CMAKE_CXX_FLAGS_DEBUG "-fPIC -O0 -g -Wall")275 set(CMAKE_CXX_FLAGS_DEBUG "-fPIC -O0 -g -Wall")
276 set(CMAKE_CXX_FLAGS_RELEASE "-fPIC -O2 -Wall")276 set(CMAKE_CXX_FLAGS_RELEASE "-fPIC -O2 -Wall")
277 277 
278 # 添加可执行文件(自定义:替换为实际调用算子的*.cpp文件)278 # 添加可执行文件(自定义:替换为实际调用算子的*.cpp文件)
279- add_executable(${test_aclnn_op_name} 279+ add_executable(${test_aclnn_op_name}
280- ${test_aclnn_op_name}.cpp) 280+ ${test_aclnn_op_name}.cpp)
281 281 
282 # ASCEND_PATH(如遇CANN包路径有误,请根据实际路径修改)282 # ASCEND_PATH(如遇CANN包路径有误,请根据实际路径修改)
283- if(NOT "$ENV{ASCEND_HOME_PATH}" STREQUAL "") 283+ if(NOT "$ENV{ASCEND_HOME_PATH}" STREQUAL "")
284 set(ASCEND_PATH $ENV{ASCEND_HOME_PATH})284 set(ASCEND_PATH $ENV{ASCEND_HOME_PATH})
285 else()285 else()
286 set(ASCEND_PATH "/usr/local/Ascend/cann")286 set(ASCEND_PATH "/usr/local/Ascend/cann")
@@ -296,7 +296,7 @@
296 endforeach()296 endforeach()
297 297 
298 if(NOT DEFINED TARGET_SUBDIR)298 if(NOT DEFINED TARGET_SUBDIR)
299- message(FATAL_ERROR "在路径${ASCEND_PATH}中未找到自定义算子包") 299+ message(FATAL_ERROR "在路径${ASCEND_PATH}中未找到自定义算子包")
300 endif()300 endif()
301 301 
302 # 设置头文件路径302 # 设置头文件路径
@@ -310,7 +310,7 @@
310 )310 )
311 311 
312 # 链接所需的动态库(自定义:替换为实际算子可执行文件)312 # 链接所需的动态库(自定义:替换为实际算子可执行文件)
313- target_link_libraries(${test_aclnn_op_name} PRIVATE 313+ target_link_libraries(${test_aclnn_op_name} PRIVATE
314 ${ASCEND_PATH}/lib64/libascendcl.so314 ${ASCEND_PATH}/lib64/libascendcl.so
315 ${ASCEND_PATH}/lib64/libnnopbase.so315 ${ASCEND_PATH}/lib64/libnnopbase.so
316 ${TARGET_SUBDIR}/op_api/lib/libcust_opapi.so # 链接自定义算子库文件316 ${TARGET_SUBDIR}/op_api/lib/libcust_opapi.so # 链接自定义算子库文件
@@ -319,53 +319,53 @@
319 "-Wl,-rpath,${TARGET_SUBDIR}/op_api/lib"319 "-Wl,-rpath,${TARGET_SUBDIR}/op_api/lib"
320 )320 )
321 321 
322- # 安装目标文件到bin目录(自定义:替换为实际算子可执行文件) 322+ # 安装目标文件到bin目录(自定义:替换为实际算子可执行文件)
323 install(TARGETS ${test_aclnn_op_name} DESTINATION ${CMAKE_RUNTIME_OUTPUT_DIRECTORY})323 install(TARGETS ${test_aclnn_op_name} DESTINATION ${CMAKE_RUNTIME_OUTPUT_DIRECTORY})
324 ```324 ```
325 325 
326 - **调用标准算子(内置算子)**:依赖op-nn整包326 - **调用标准算子(内置算子)**:依赖op-nn整包
327 327 
328- ```bash328+ ```cmake
329 cmake_minimum_required(VERSION 3.14)329 cmake_minimum_required(VERSION 3.14)
330 # 设置工程名330 # 设置工程名
331 project(ACLNN_EXAMPLE)331 project(ACLNN_EXAMPLE)
332- 332+ 
333 # 设置C++编译标准333 # 设置C++编译标准
334 add_compile_options(-std=c++11)334 add_compile_options(-std=c++11)
335- 335+ 
336 # 设置编译输出目录为当前目录下的bin文件夹336 # 设置编译输出目录为当前目录下的bin文件夹
337 set(CMAKE_RUNTIME_OUTPUT_DIRECTORY "./bin")337 set(CMAKE_RUNTIME_OUTPUT_DIRECTORY "./bin")
338- 338+ 
339 # 设置调试和发布模式的编译选项339 # 设置调试和发布模式的编译选项
340 set(CMAKE_CXX_FLAGS_DEBUG "-fPIC -O0 -g -Wall")340 set(CMAKE_CXX_FLAGS_DEBUG "-fPIC -O0 -g -Wall")
341 set(CMAKE_CXX_FLAGS_RELEASE "-fPIC -O2 -Wall")341 set(CMAKE_CXX_FLAGS_RELEASE "-fPIC -O2 -Wall")
342- 342+ 
343 # 添加可执行文件(自定义:替换为实际调用算子的*.cpp文件)343 # 添加可执行文件(自定义:替换为实际调用算子的*.cpp文件)
344 add_executable(${test_aclnn_op_name}344 add_executable(${test_aclnn_op_name}
345 ${test_aclnn_op_name}.cpp)345 ${test_aclnn_op_name}.cpp)
346- 346+ 
347 # ASCEND_PATH(如遇CANN包路径有误,请根据实际路径修改)347 # ASCEND_PATH(如遇CANN包路径有误,请根据实际路径修改)
348 if(NOT "$ENV{ASCEND_HOME_PATH}" STREQUAL "")348 if(NOT "$ENV{ASCEND_HOME_PATH}" STREQUAL "")
349 set(ASCEND_PATH $ENV{ASCEND_HOME_PATH})349 set(ASCEND_PATH $ENV{ASCEND_HOME_PATH})
350 else()350 else()
351 set(ASCEND_PATH "/usr/local/Ascend/cann")351 set(ASCEND_PATH "/usr/local/Ascend/cann")
352 endif()352 endif()
353- 353+ 
354 # 设置头文件路径354 # 设置头文件路径
355 set(INCLUDE_BASE_DIR "${ASCEND_PATH}/include")355 set(INCLUDE_BASE_DIR "${ASCEND_PATH}/include")
356 include_directories(356 include_directories(
357 ${INCLUDE_BASE_DIR}357 ${INCLUDE_BASE_DIR}
358 ${ASCEND_PATH}/include/aclnnop358 ${ASCEND_PATH}/include/aclnnop
359 )359 )
360- 360+ 
361 # 链接所需的动态库(自定义:替换为实际算子可执行文件)361 # 链接所需的动态库(自定义:替换为实际算子可执行文件)
362 target_link_libraries(${test_aclnn_op_name} PRIVATE362 target_link_libraries(${test_aclnn_op_name} PRIVATE
363 ${ASCEND_PATH}/lib64/libascendcl.so363 ${ASCEND_PATH}/lib64/libascendcl.so
364 ${ASCEND_PATH}/lib64/libnnopbase.so364 ${ASCEND_PATH}/lib64/libnnopbase.so
365 ${ASCEND_PATH}/lib64/libopapi_nn.so # 链接内置算子库文件365 ${ASCEND_PATH}/lib64/libopapi_nn.so # 链接内置算子库文件
366 )366 )
367- 367+ 
368- # 安装目标文件到bin目录(自定义:替换为实际算子可执行文件) 368+ # 安装目标文件到bin目录(自定义:替换为实际算子可执行文件)
369 install(TARGETS ${test_aclnn_op_name} DESTINATION ${CMAKE_RUNTIME_OUTPUT_DIRECTORY})369 install(TARGETS ${test_aclnn_op_name} DESTINATION ${CMAKE_RUNTIME_OUTPUT_DIRECTORY})
370 ```370 ```
371 371 
@@ -381,11 +381,11 @@
381 else381 else
382 _ASCEND_INSTALL_PATH="/usr/local/Ascend/cann"382 _ASCEND_INSTALL_PATH="/usr/local/Ascend/cann"
383 fi383 fi
384- 384+ 
385 source ${_ASCEND_INSTALL_PATH}/bin/setenv.bash385 source ${_ASCEND_INSTALL_PATH}/bin/setenv.bash
386- 386+ 
387 rm -rf build387 rm -rf build
388- mkdir -p build 388+ mkdir -p build
389 cd build389 cd build
390 cmake ../ -DCMAKE_CXX_COMPILER=g++ -DCMAKE_SKIP_RPATH=TRUE # 执行构建命令390 cmake ../ -DCMAKE_CXX_COMPILER=g++ -DCMAKE_SKIP_RPATH=TRUE # 执行构建命令
391 make391 make
@@ -431,41 +431,41 @@
431 int main() {431 int main() {
432 // 1. 创建图对象432 // 1. 创建图对象
433 Graph graph(graphName);433 Graph graph(graphName);
434- 434+ 
435 // 2. 图全局编译选项初始化435 // 2. 图全局编译选项初始化
436 Status ret = ge::GEInitialize(globalOptions);436 Status ret = ge::GEInitialize(globalOptions);
437- 437+ 
438 // 3. 创建AddExample算子实例438 // 3. 创建AddExample算子实例
439 auto add1 = op::AddExample("add1");439 auto add1 = op::AddExample("add1");
440- 440+ 
441 // 4. 定义图输入输出向量441 // 4. 定义图输入输出向量
442 std::vector<Operator> inputs{};442 std::vector<Operator> inputs{};
443 std::vector<Operator> outputs{};443 std::vector<Operator> outputs{};
444- 444+ 
445 // 5. 准备输入数据445 // 5. 准备输入数据
446 std::vector<int64_t> xShape = {32,4,4,4};446 std::vector<int64_t> xShape = {32,4,4,4};
447 // 宏展开方式处理变量赋值447 // 宏展开方式处理变量赋值
448 ADD_INPUT(1, x1, inDtype, xShape);448 ADD_INPUT(1, x1, inDtype, xShape);
449 ADD_INPUT(2, x2, inDtype, xShape);449 ADD_INPUT(2, x2, inDtype, xShape);
450 ADD_OUTPUT(1, y, inDtype, xShape);450 ADD_OUTPUT(1, y, inDtype, xShape);
451- 451+ 
452 outputs.push_back(add1);452 outputs.push_back(add1);
453- 453+ 
454 // 6. 设置图对象的输入算子和输出算子454 // 6. 设置图对象的输入算子和输出算子
455 graph.SetInputs(inputs).SetOutputs(outputs);455 graph.SetInputs(inputs).SetOutputs(outputs);
456- 456+ 
457 // 7. 创建session对象457 // 7. 创建session对象
458 ge::Session* session = new Session(buildOptions);458 ge::Session* session = new Session(buildOptions);
459- 459+ 
460 // 8. session添加图460 // 8. session添加图
461 ret = session->AddGraph(graphId, graph, graphOptions);461 ret = session->AddGraph(graphId, graph, graphOptions);
462- 462+ 
463 // 9. 运行图463 // 9. 运行图
464 ret = session->RunGraph(graphId, input, output);464 ret = session->RunGraph(graphId, input, output);
465- 465+ 
466 // 10. 释放资源466 // 10. 释放资源
467 GEFinalize();467 GEFinalize();
468- 468+ 
469 return 0;469 return 0;
470 }470 }
471 ```471 ```
@@ -474,12 +474,12 @@
474 474 
475`${test_geir_op_name}.cpp`同级目录下创建CMakeLists.txt文件,GE图引擎会根据配置好的环境变量自动加载已安装好的算子包(无论是自定义算子包或是标准内置算子包)库文件,无需特意区分。示例如下,仅供参考,请根据实际情况自行修改。475`${test_geir_op_name}.cpp`同级目录下创建CMakeLists.txt文件,GE图引擎会根据配置好的环境变量自动加载已安装好的算子包(无论是自定义算子包或是标准内置算子包)库文件,无需特意区分。示例如下,仅供参考,请根据实际情况自行修改。
476 476 
477- ```bash477+ ```cmake
478 cmake_minimum_required(VERSION 3.14)478 cmake_minimum_required(VERSION 3.14)
479- 479+ 
480 # 设置工程名480 # 设置工程名
481 project(GE_IR_EXAMPLE)481 project(GE_IR_EXAMPLE)
482- 482+ 
483 if(NOT "$ENV{ASCEND_OPP_PATH}" STREQUAL "")483 if(NOT "$ENV{ASCEND_OPP_PATH}" STREQUAL "")
484 get_filename_component(ASCEND_PATH $ENV{ASCEND_OPP_PATH} DIRECTORY)484 get_filename_component(ASCEND_PATH $ENV{ASCEND_OPP_PATH} DIRECTORY)
485 elseif(NOT "$ENV{ASCEND_HOME_PATH}" STREQUAL "")485 elseif(NOT "$ENV{ASCEND_HOME_PATH}" STREQUAL "")
@@ -487,31 +487,31 @@
487 else()487 else()
488 set(ASCEND_PATH "/usr/local/Ascend/cann")488 set(ASCEND_PATH "/usr/local/Ascend/cann")
489 endif()489 endif()
490- 490+ 
491 set(FWK_INCLUDE_DIR "${ASCEND_PATH}/compiler/include")491 set(FWK_INCLUDE_DIR "${ASCEND_PATH}/compiler/include")
492- 492+ 
493 message(STATUS "ASCEND_PATH: ${ASCEND_PATH}")493 message(STATUS "ASCEND_PATH: ${ASCEND_PATH}")
494- 494+ 
495 file(GLOB files CONFIGURE_DEPENDS495 file(GLOB files CONFIGURE_DEPENDS
496- ${test_geir_op_name}.cpp 496+ ${test_geir_op_name}.cpp
497 )497 )
498- 498+ 
499 # 添加可执行文件(请替换为实际算子可执行文件)499 # 添加可执行文件(请替换为实际算子可执行文件)
500- add_executable(${test_geir_op_name} ${files}) 500+ add_executable(${test_geir_op_name} ${files})
501- 501+ 
502 find_library(GRAPH_LIBRARY_DIR libgraph.so "${ASCEND_PATH}/compiler/lib64/stub")502 find_library(GRAPH_LIBRARY_DIR libgraph.so "${ASCEND_PATH}/compiler/lib64/stub")
503 find_library(GE_RUNNER_LIBRARY_DIR libge_runner.so "${ASCEND_PATH}/compiler/lib64/stub")503 find_library(GE_RUNNER_LIBRARY_DIR libge_runner.so "${ASCEND_PATH}/compiler/lib64/stub")
504 find_library(GRAPH_BASE_LIBRARY_DIR libgraph_base.so "${ASCEND_PATH}/compiler/lib64")504 find_library(GRAPH_BASE_LIBRARY_DIR libgraph_base.so "${ASCEND_PATH}/compiler/lib64")
505- 505+ 
506 # 链接所需的动态库506 # 链接所需的动态库
507- target_link_libraries(${test_geir_op_name} PRIVATE 507+ target_link_libraries(${test_geir_op_name} PRIVATE
508 ${GRAPH_LIBRARY_DIR}508 ${GRAPH_LIBRARY_DIR}
509 ${GE_RUNNER_LIBRARY_DIR}509 ${GE_RUNNER_LIBRARY_DIR}
510 ${GRAPH_BASE_LIBRARY_DIR}510 ${GRAPH_BASE_LIBRARY_DIR}
511 )511 )
512- 512+ 
513 # 设置头文件路径513 # 设置头文件路径
514- target_include_directories(${test_geir_op_name} PRIVATE 514+ target_include_directories(${test_geir_op_name} PRIVATE
515 ${FWK_INCLUDE_DIR}/graph/515 ${FWK_INCLUDE_DIR}/graph/
516 ${FWK_INCLUDE_DIR}/ge/516 ${FWK_INCLUDE_DIR}/ge/
517 ${ASCEND_PATH}/opp/built-in/op_proto/inc/517 ${ASCEND_PATH}/opp/built-in/op_proto/inc/
@@ -532,11 +532,11 @@
532 else532 else
533 _ASCEND_INSTALL_PATH="/usr/local/Ascend/cann"533 _ASCEND_INSTALL_PATH="/usr/local/Ascend/cann"
534 fi534 fi
535- 535+ 
536- source ${_ASCEND_INSTALL_PATH}/bin/setenv.bash 536+ source ${_ASCEND_INSTALL_PATH}/bin/setenv.bash
537- 537+ 
538- rm -rf build 538+ rm -rf build
539- mkdir -p build 539+ mkdir -p build
540 cd build540 cd build
541 cmake ../ -DCMAKE_CXX_COMPILER=g++ -DCMAKE_SKIP_RPATH=TRUE # 执行构建命令541 cmake ../ -DCMAKE_CXX_COMPILER=g++ -DCMAKE_SKIP_RPATH=TRUE # 执行构建命令
542 make542 make
@@ -8,8 +8,8 @@
8 8 
9 调用算子API时,需引用依赖的头文件和库文件,一般头文件默认在`${INSTALL_DIR}/include/aclnnop`,库文件默认在`${INSTALL_DIR}/lib64`,具体文件如下:9 调用算子API时,需引用依赖的头文件和库文件,一般头文件默认在`${INSTALL_DIR}/include/aclnnop`,库文件默认在`${INSTALL_DIR}/lib64`,具体文件如下:
10 10 
11- - 头文件:方式1 (推荐):引用算子仓总头文件aclnn\_ops\_\$\{ops\_project\}.h。方式2:引用单个算子API的头文件aclnn\_\*.h。11+ - 头文件:方式1(推荐):引用算子仓总头文件aclnn\_ops\_\$\{ops\_project\}.h。方式2:引用单个算子API的头文件aclnn\_\*.h。
12- - 库文件:引用算子仓对应的库文件libopapi_${ops_project}.so。注意,原所有算子仓总库文件libopapi.so后续会废弃,不推荐使用,也不支持与单个算子仓文件同时使用。12+ - 库文件:引用算子仓对应的库文件libopapi_${ops_project}.so。注意,原所有算子仓总库文件libopapi.so后续会废弃,不推荐使用,也不支持与单个算子仓文件同时使用。
13 13 
14 ${INSTALL_DIR}表示CANN安装后文件路径;\$\{ops\_project\}表示算子仓名(如math、nn、cv、transformer),请改为实际算子仓名。14 ${INSTALL_DIR}表示CANN安装后文件路径;\$\{ops\_project\}表示算子仓名(如math、nn、cv、transformer),请改为实际算子仓名。
15 15 
@@ -92,7 +92,7 @@
92| [aclnnBinaryCrossEntropyBackward](../../loss/binary_cross_entropy_grad/docs/aclnnBinaryCrossEntropyBackward.md) | 求二元交叉熵反向传播的梯度值。 | 默认确定性实现 | 默认确定性实现 |92| [aclnnBinaryCrossEntropyBackward](../../loss/binary_cross_entropy_grad/docs/aclnnBinaryCrossEntropyBackward.md) | 求二元交叉熵反向传播的梯度值。 | 默认确定性实现 | 默认确定性实现 |
93| [aclnnBinaryCrossEntropyWithLogits](../../loss/sigmoid_cross_entropy_with_logits_v2/docs/aclnnBinaryCrossEntropyWithLogits.md) | 计算输入logits与标签target之间的BCELoss损失。 | 默认确定性实现 | 默认确定性实现 |93| [aclnnBinaryCrossEntropyWithLogits](../../loss/sigmoid_cross_entropy_with_logits_v2/docs/aclnnBinaryCrossEntropyWithLogits.md) | 计算输入logits与标签target之间的BCELoss损失。 | 默认确定性实现 | 默认确定性实现 |
94| [aclnnBinaryCrossEntropyWithLogitsBackward](../../loss/sigmoid_cross_entropy_with_logits_grad_v2/docs/aclnnBinaryCrossEntropyWithLogitsBackward.md) | 将输入self执行logits计算,将得到的值与标签值target一起进行BCELoss的反向传播计算。 | 默认确定性实现 | 默认确定性实现 |94| [aclnnBinaryCrossEntropyWithLogitsBackward](../../loss/sigmoid_cross_entropy_with_logits_grad_v2/docs/aclnnBinaryCrossEntropyWithLogitsBackward.md) | 将输入self执行logits计算,将得到的值与标签值target一起进行BCELoss的反向传播计算。 | 默认确定性实现 | 默认确定性实现 |
95-| [aclnnBinaryCrossEntropyWithLogitsTargetBackward](../../activation/log_sigmoid/docs/aclnnBinaryCrossEntropyWithLogitsTargetBackward.md) | 将输入self执行logits计算,将得到的值与标签值target一起进行BECLoss关于target的反向传播计算。 | 默认确定性实现 | - |95+| [aclnnBinaryCrossEntropyWithLogitsTargetBackward](../../activation/log_sigmoid/docs/aclnnBinaryCrossEntropyWithLogitsTargetBackward.md) | 将输入self执行logits计算,将得到的值与标签值target一起进行BCELoss关于target的反向传播计算。 | 默认确定性实现 | - |
96| [aclnnBucketize](../../index/bucketize_v2/docs/aclnnBucketize.md) | 根据给定的边界数组(boundaries)确定输入张量中每个元素所属的区间索引。 | - | 默认确定性实现 |96| [aclnnBucketize](../../index/bucketize_v2/docs/aclnnBucketize.md) | 根据给定的边界数组(boundaries)确定输入张量中每个元素所属的区间索引。 | - | 默认确定性实现 |
97| [aclnnCelu&aclnnInplaceCelu](../../activation/celu_v2/docs/aclnnCelu&aclnnInplaceCelu.md) | aclnnCelu对输入张量self中的每个元素x调用连续可微指数线性单元激活函数CELU,并将得到的结果存入输出张量out中。 | 默认确定性实现 | 默认确定性实现 |97| [aclnnCelu&aclnnInplaceCelu](../../activation/celu_v2/docs/aclnnCelu&aclnnInplaceCelu.md) | aclnnCelu对输入张量self中的每个元素x调用连续可微指数线性单元激活函数CELU,并将得到的结果存入输出张量out中。 | 默认确定性实现 | 默认确定性实现 |
98| [aclnnChamferDistanceBackward](../../loss/chamfer_distance_grad/docs/aclnnChamferDistanceBackward.md) | ChamferDistance(倒角距离)的反向算子,根据正向的输入对输出的贡献及初始梯度求出输入对应的梯度。 | 默认非确定性实现,支持配置开启 | - |98| [aclnnChamferDistanceBackward](../../loss/chamfer_distance_grad/docs/aclnnChamferDistanceBackward.md) | ChamferDistance(倒角距离)的反向算子,根据正向的输入对输出的贡献及初始梯度求出输入对应的梯度。 | 默认非确定性实现,支持配置开启 | - |
@@ -302,7 +302,7 @@
302| [aclnnL1Loss](../../loss/lp_loss/docs/aclnnL1Loss.md) | 计算输入self和目标target中每个元素之间的平均绝对误差(Mean Absolute Error,简称MAE)。 | 默认确定性实现 | 默认确定性实现 |302| [aclnnL1Loss](../../loss/lp_loss/docs/aclnnL1Loss.md) | 计算输入self和目标target中每个元素之间的平均绝对误差(Mean Absolute Error,简称MAE)。 | 默认确定性实现 | 默认确定性实现 |
303| [aclnnL1LossBackward](../../loss/l1_loss_grad/docs/aclnnL1LossBackward.md) | 计算aclnnL1Loss的反向传播。reduction指定损失函数的计算方式。 | 默认确定性实现 | 默认确定性实现 |303| [aclnnL1LossBackward](../../loss/l1_loss_grad/docs/aclnnL1LossBackward.md) | 计算aclnnL1Loss的反向传播。reduction指定损失函数的计算方式。 | 默认确定性实现 | 默认确定性实现 |
304| [aclnnLayerNorm&aclnnLayerNormWithImplMode](../../norm/layer_norm_v4/docs/aclnnLayerNorm&aclnnLayerNormWithImplMode.md) | 对指定层进行均值为0、标准差为1的归一化计算。 | 默认确定性实现 | 默认确定性实现 |304| [aclnnLayerNorm&aclnnLayerNormWithImplMode](../../norm/layer_norm_v4/docs/aclnnLayerNorm&aclnnLayerNormWithImplMode.md) | 对指定层进行均值为0、标准差为1的归一化计算。 | 默认确定性实现 | 默认确定性实现 |
305-| [aclnnLayerNormBackward](../../norm/layer_norm_grad_v3/docs/aclnnLayerNormBackward.md) | [aclnnNorm](../../norm/lp_norm_v2/docs/aclnnNorm.md)的反向传播。用于计算输入张量的梯度,以便在反向传播过程中更新模型参数。 | 默认确定性实现 | 默认确定性实现 |305+| [aclnnLayerNormBackward](../../norm/layer_norm_grad_v3/docs/aclnnLayerNormBackward.md) | [aclnnLayerNorm](../../norm/layer_norm_v4/docs/aclnnLayerNorm&aclnnLayerNormWithImplMode.md)的反向传播。用于计算输入张量的梯度,以便在反向传播过程中更新模型参数。 | 默认确定性实现 | 默认确定性实现 |
306| [aclnnLayerNormQuant](../../norm/layer_norm_quant/docs/aclnnLayerNormQuant.md) | 算子将LayerNorm归一化输出和下游的量化算子融合起来,减少搬入搬出操作。 | 默认确定性实现 | 默认确定性实现 |306| [aclnnLayerNormQuant](../../norm/layer_norm_quant/docs/aclnnLayerNormQuant.md) | 算子将LayerNorm归一化输出和下游的量化算子融合起来,减少搬入搬出操作。 | 默认确定性实现 | 默认确定性实现 |
307| [aclnnLeakyRelu&aclnnInplaceLeakyRelu](../../activation/leaky_relu/docs/aclnnLeakyRelu&aclnnInplaceLeakyRelu.md) | 激活函数,用于解决Relu函数在输入小于0时输出为0的问题,避免神经元无法更新参数。 | 默认确定性实现 | 默认确定性实现 |307| [aclnnLeakyRelu&aclnnInplaceLeakyRelu](../../activation/leaky_relu/docs/aclnnLeakyRelu&aclnnInplaceLeakyRelu.md) | 激活函数,用于解决Relu函数在输入小于0时输出为0的问题,避免神经元无法更新参数。 | 默认确定性实现 | 默认确定性实现 |
308| [aclnnLeakyReluBackward](../../activation/leaky_relu_grad/docs/aclnnLeakyReluBackward.md) | LeakyRelu激活函数反向。 | 默认确定性实现 | 默认确定性实现 |308| [aclnnLeakyReluBackward](../../activation/leaky_relu_grad/docs/aclnnLeakyReluBackward.md) | LeakyRelu激活函数反向。 | 默认确定性实现 | 默认确定性实现 |
@@ -404,7 +404,6 @@
404| [aclnnSigmoidBackward](../../activation/sigmoid_grad/docs/aclnnSigmoidBackward.md) | 完成sigmoid的反向传播,根据sigmoid反向传播梯度与正向输出计算sigmoid的梯度输入。 | 默认确定性实现 | 默认确定性实现 |404| [aclnnSigmoidBackward](../../activation/sigmoid_grad/docs/aclnnSigmoidBackward.md) | 完成sigmoid的反向传播,根据sigmoid反向传播梯度与正向输出计算sigmoid的梯度输入。 | 默认确定性实现 | 默认确定性实现 |
405| [aclnnSilu](../../activation/swish/docs/aclnnSilu.md) | 该算子也被称为Swish函数。 | 默认确定性实现 | 默认确定性实现 |405| [aclnnSilu](../../activation/swish/docs/aclnnSilu.md) | 该算子也被称为Swish函数。 | 默认确定性实现 | 默认确定性实现 |
406| [aclnnSiluBackward](../../activation/silu_grad/docs/aclnnSiluBackward.md) | aclnnSilu的反向传播,根据silu反向传播梯度与正向输出计算silu的梯度输入。 | 默认确定性实现 | 默认确定性实现 |406| [aclnnSiluBackward](../../activation/silu_grad/docs/aclnnSiluBackward.md) | aclnnSilu的反向传播,根据silu反向传播梯度与正向输出计算silu的梯度输入。 | 默认确定性实现 | 默认确定性实现 |
407-| [aclnnSleep](../../control/sleep/docs/aclnnSleep.md) | 让NPU设备休眠指定的时钟周期数,通过忙等待实现精确的延时控制。 | - | 默认确定性实现 |
408| [aclnnSmoothL1Loss](../../loss/smooth_l1_loss_v2/docs/aclnnSmoothL1Loss.md) | 计算SmoothL1损失函数。 | 默认确定性实现 | 默认确定性实现 |407| [aclnnSmoothL1Loss](../../loss/smooth_l1_loss_v2/docs/aclnnSmoothL1Loss.md) | 计算SmoothL1损失函数。 | 默认确定性实现 | 默认确定性实现 |
409| [aclnnSmoothL1LossBackward](../../loss/smooth_l1_loss_grad_v2/docs/aclnnSmoothL1LossBackward.md) | 计算aclnnSmoothL1Loss api的反向传播。 | 默认确定性实现 | 默认确定性实现 |408| [aclnnSmoothL1LossBackward](../../loss/smooth_l1_loss_grad_v2/docs/aclnnSmoothL1LossBackward.md) | 计算aclnnSmoothL1Loss api的反向传播。 | 默认确定性实现 | 默认确定性实现 |
410| [aclnnSoftMarginLoss](../../loss/soft_margin_loss/docs/aclnnSoftMarginLoss.md) | 计算输入self和目标target的二分类逻辑损失函数。 | 默认确定性实现 | 默认确定性实现 |409| [aclnnSoftMarginLoss](../../loss/soft_margin_loss/docs/aclnnSoftMarginLoss.md) | 计算输入self和目标target的二分类逻辑损失函数。 | 默认确定性实现 | 默认确定性实现 |
@@ -443,7 +442,7 @@
443| [aclnnTransQuantParamV2](../../quant/trans_quant_param_v2/docs/aclnnTransQuantParamV2.md) | 完成量化计算参数scale数据类型的转换,将FLOAT32的数据类型转换为硬件需要的UINT64,INT64类型。 | 默认确定性实现 | 默认确定性实现 |442| [aclnnTransQuantParamV2](../../quant/trans_quant_param_v2/docs/aclnnTransQuantParamV2.md) | 完成量化计算参数scale数据类型的转换,将FLOAT32的数据类型转换为硬件需要的UINT64,INT64类型。 | 默认确定性实现 | 默认确定性实现 |
444| [aclnnTransQuantParamV3](../../quant/trans_quant_param_v2/docs/aclnnTransQuantParamV3.md) | 完成量化计算参数scale数据类型的转换,将Float32的数据类型转换为硬件需要的UINT64,INT64类型。 | 默认确定性实现 | 默认确定性实现 |443| [aclnnTransQuantParamV3](../../quant/trans_quant_param_v2/docs/aclnnTransQuantParamV3.md) | 完成量化计算参数scale数据类型的转换,将Float32的数据类型转换为硬件需要的UINT64,INT64类型。 | 默认确定性实现 | 默认确定性实现 |
445| [aclnnTopKTopPSample](../../index/top_k_top_p_sample/docs/aclnnTopKTopPSample.md) | 根据输入词频logits、topK/topP采样参数、随机采样权重分布q,进行topK-topP-sample采样计算,输出每个batch的最大词频logitsSelectIdx,以及topK-topP采样后的词频分布logitsTopKPSelect。 | 默认确定性实现 | - |444| [aclnnTopKTopPSample](../../index/top_k_top_p_sample/docs/aclnnTopKTopPSample.md) | 根据输入词频logits、topK/topP采样参数、随机采样权重分布q,进行topK-topP-sample采样计算,输出每个batch的最大词频logitsSelectIdx,以及topK-topP采样后的词频分布logitsTopKPSelect。 | 默认确定性实现 | - |
446-| [aclnnTopKTopPSampleV2](../../index/top_k_top_p_sample_v2/docs/aclnnTopKTopPSampleV2.md) | 根据输入词频logits、topK/topP/minP采样参数、随机采样权重分布q,进行topK-topP-minP-sample采样计算。当输入isNeedSampleResult为false时,输出每个batch的最大词频logitsSelectIdx,以及topK-topP-minP采样后的词频分布logitsTopKPSelect;当输入isNeedSampleResult为true时,输出topK-topP-minP采样后的中间计算结果logitsIdx和logitsSortMasked,其中logitsSortMasked为词频logits经过topK-topP-minP采样计算后的中间结果,logitsIdx为logitsSortMasked在logits中对应的索引。 | 默认确定性实现 | 默认确定性实现 |445+| [aclnnTopKTopPSampleV2](../../index/top_k_top_p_sample_v2/docs/aclnnTopKTopPSampleV2.md) | 根据输入词频logits、topK/topP/minP采样参数、随机采样权重分布q,进行topK-topP-minP-sample采样计算。当输入isNeedSampleResult为false时,输出每个batch的最大词频logitsSelectIdx,以及topK-topP-minP采样后的词频分布logitsTopKPSelect;当输入isNeedSampleResult为true时,输出topK-topP-minP采样后的中间计算结果logitsIdx和logitsSortMasked,其中logitsSortMasked为词频logits经过topK-topP-minP采样计算后的中间结果,logitsIdx为logitsSortMasked在logits中对应的索引。 | 默认确定性实现 | - |
447| [aclnnUnique](../../index/scatter_elements/docs/aclnnUnique.md) | 返回输入张量中的唯一元素。 | 默认确定性实现 | 默认确定性实现 |446| [aclnnUnique](../../index/scatter_elements/docs/aclnnUnique.md) | 返回输入张量中的唯一元素。 | 默认确定性实现 | 默认确定性实现 |
448| [aclnnUnique2](../../index/scatter_elements/docs/aclnnUnique2.md) | 对输入张量self进行去重,返回self中的唯一元素。unique功能的增强,新增返回值countsOut,表示valueOut中各元素在输入self中出现的次数,用returnCounts参数控制。 | 默认确定性实现 | 默认确定性实现 |447| [aclnnUnique2](../../index/scatter_elements/docs/aclnnUnique2.md) | 对输入张量self进行去重,返回self中的唯一元素。unique功能的增强,新增返回值countsOut,表示valueOut中各元素在输入self中出现的次数,用returnCounts参数控制。 | 默认确定性实现 | 默认确定性实现 |
449| [aclnnUniqueConsecutive](../../index/unique_consecutive/docs/aclnnUniqueConsecutive.md) | 去除每一个元素后的重复元素。当dim不为空时,去除对应维度上的每一个张量后的重复张量。 | 默认确定性实现 | 默认确定性实现 |448| [aclnnUniqueConsecutive](../../index/unique_consecutive/docs/aclnnUniqueConsecutive.md) | 去除每一个元素后的重复元素。当dim不为空时,去除对应维度上的每一个张量后的重复张量。 | 默认确定性实现 | 默认确定性实现 |
@@ -139,7 +139,7 @@
139 </tr>139 </tr>
140 <tr>140 <tr>
141 <td>activation</td>141 <td>activation</td>
142- <td><a href="../../activation/fast_gelu_v2/README.md">fast_gelu</a></td>142+ <td><a href="../../activation/fast_gelu_v2/README.md">fast_gelu_v2</a></td>
143 <td>✓</td>143 <td>✓</td>
144 <td>✓</td>144 <td>✓</td>
145 <td>✓</td>145 <td>✓</td>
@@ -2513,7 +2513,7 @@
2513 <td>✓</td>2513 <td>✓</td>
2514 <td>✓</td>2514 <td>✓</td>
2515 <td>✓</td>2515 <td>✓</td>
2516- <td>✓</td>02516+ <td>✓</td>
2517 <td>AI Core</td>2517 <td>AI Core</td>
2518 <td>该算子暂无Ascend C代码实现,欢迎开发者补充贡献,贡献方式参考<a href="../../CONTRIBUTING.md">贡献指南</a>。</td>2518 <td>该算子暂无Ascend C代码实现,欢迎开发者补充贡献,贡献方式参考<a href="../../CONTRIBUTING.md">贡献指南</a>。</td>
2519 </tr>2519 </tr>
@@ -3095,7 +3095,7 @@
3095 <td>✓</td>3095 <td>✓</td>
3096 <td>✓</td>3096 <td>✓</td>
3097 <td>AI Core</td>3097 <td>AI Core</td>
3098- <td>RmsNorm算子是大模型常用的归一化操作,相比LayerNorm算子,其去掉了减去均值的部分。DynamicQuant算子则是为输入张量进行对称动态量化的算子。AddRmsNormDynamicQuantV2算子将RmsNorm前的Add算子和RmsNorm归一化输出给到的1个或2个DynamicQuant算子融合起来,减少搬入搬出操作。AddRmsNormDynamicQuant算子相较于AddRmsNormDynamicQuantV2在RmsNorm计算过程中增加了偏置项betaOptional参数,即计算对应公式中的beta,以及新增输出配置项output_mask参数,用于配置是否输出对应位置的量化结果。</td>3098+ <td>RmsNorm算子是大模型常用的归一化操作,相比LayerNorm算子,其去掉了减去均值的部分。DynamicQuant算子则是为输入张量进行对称动态量化的算子。AddRmsNormDynamicQuant算子将RmsNorm前的Add算子和RmsNorm归一化输出给到的1个或2个DynamicQuant算子融合起来,减少搬入搬出操作。AddRmsNormDynamicQuant算子相较于AddRmsNormDynamicQuantV2在RmsNorm计算过程中增加了偏置项betaOptional参数,即计算对应公式中的beta,以及新增输出配置项output_mask参数,用于配置是否输出对应位置的量化结果。</td>
3099 </tr>3099 </tr>
3100 <tr>3100 <tr>
3101 <td>norm</td>3101 <td>norm</td>
@@ -3548,7 +3548,7 @@
3548 <td>AI Core</td>3548 <td>AI Core</td>
3549 <td>L2Normalize的反向计算。计算输入x的梯度dx。</td>3549 <td>L2Normalize的反向计算。计算输入x的梯度dx。</td>
3550 </tr>3550 </tr>
3551- <tr>3551+ <tr>
3552 <td>norm</td>3552 <td>norm</td>
3553 <td><a href="../../norm/rms_norm_grad_quant/README.md">rms_norm_grad_quant</a></td>3553 <td><a href="../../norm/rms_norm_grad_quant/README.md">rms_norm_grad_quant</a></td>
3554 <td>✓</td>3554 <td>✓</td>
@@ -3688,7 +3688,7 @@
3688 <td>AI Core</td>3688 <td>AI Core</td>
3689 <td>执行Adadelta优化器的单步参数更新,根据当前梯度、梯度平方累积accum和更新量平方累积accum_update,原地更新权重参数var以及accum、accum_update。</td>3689 <td>执行Adadelta优化器的单步参数更新,根据当前梯度、梯度平方累积accum和更新量平方累积accum_update,原地更新权重参数var以及accum、accum_update。</td>
3690 </tr>3690 </tr>
3691- <tr>3691+ <tr>
3692 <td>optim</td>3692 <td>optim</td>
3693 <td><a href="../../optim/inplace_apply_adadelta/README.md">inplace_apply_adadelta</a></td>3693 <td><a href="../../optim/inplace_apply_adadelta/README.md">inplace_apply_adadelta</a></td>
3694 <td>✓</td>3694 <td>✓</td>
@@ -3896,7 +3896,7 @@
3896 <td>✗</td>3896 <td>✗</td>
3897 <td>✓</td>3897 <td>✓</td>
3898 <td>AI Core</td>3898 <td>AI Core</td>
3899- <td>对模型中的一个参数(如权重),完成Adam优化算法的单步计算和更新,在adam_apply_one_with_decay基础上增加了赋值操作,确保更新后的值被正确地保存。</td>3899+ <td>对模型中的一个参数(如权重),完成Adam优化算法的单步计算和更新,在adam_apply_one_with_decay基础上增加了赋值操作,确保更新后的值被正确地保存。</td>
3900 </tr>3900 </tr>
3901 <tr>3901 <tr>
3902 <td>optim</td>3902 <td>optim</td>
@@ -3996,7 +3996,7 @@
3996 <td>✓</td>3996 <td>✓</td>
3997 <td>✗</td>3997 <td>✗</td>
3998 <td>AI Core</td>3998 <td>AI Core</td>
3999- <td>根据输入的output_size计算每次kernel的大小,对输入x进行3维最大池化操作,输出池化后的值y和索引indices。</td>3999+ <td>根据输入的output_size计算每次kernel的大小,对输入x进行2维最大池化操作,输出池化后的值y和索引indices。</td>
4000 </tr>4000 </tr>
4001 <tr>4001 <tr>
4002 <td>pooling</td>4002 <td>pooling</td>
@@ -4520,7 +4520,7 @@
4520 </tr>4520 </tr>
4521 <tr>4521 <tr>
4522 <td>quant</td>4522 <td>quant</td>
4523- <td><a href="../../quant/quant_max/README.md">group_max</a></td>4523+ <td><a href="../../quant/quant_max/README.md">quant_max</a></td>
4524 <td>✓</td>4524 <td>✓</td>
4525 <td>✓</td>4525 <td>✓</td>
4526 <td>✓</td>4526 <td>✓</td>
@@ -4530,7 +4530,7 @@
4530 </tr>4530 </tr>
4531 <tr>4531 <tr>
4532 <td>quant</td>4532 <td>quant</td>
4533- <td><a href="../../quant/grouped_quant_max/README.md">group_max</a></td>4533+ <td><a href="../../quant/grouped_quant_max/README.md">grouped_quant_max</a></td>
4534 <td>✓</td>4534 <td>✓</td>
4535 <td>✓</td>4535 <td>✓</td>
4536 <td>✓</td>4536 <td>✓</td>
@@ -4640,7 +4640,7 @@
4640 </tr>4640 </tr>
4641 <tr>4641 <tr>
4642 <td>rnn</td>4642 <td>rnn</td>
4643- <td><a href="../../rnn/bidirection_lstmv2/README.md">bidirection_lstm_v2</a></td>4643+ <td><a href="../../rnn/bidirection_lstmv2/README.md">bidirection_lstmv2</a></td>
4644 <td>✓</td>4644 <td>✓</td>
4645 <td>✓</td>4645 <td>✓</td>
4646 <td>✓</td>4646 <td>✓</td>
@@ -4949,8 +4949,7 @@
4949 <td>AI Core</td>4949 <td>AI Core</td>
4950 <td>稀疏版本的AdagradV2优化器算子,根据稀疏梯度和索引向量原地更新参数变量和累加器。</td>4950 <td>稀疏版本的AdagradV2优化器算子,根据稀疏梯度和索引向量原地更新参数变量和累加器。</td>
4951 </tr>4951 </tr>
4952- </tr>4952+ <tr>
4953- <tr>
4954 <td>optim</td>4953 <td>optim</td>
4955 <td><a href="../../optim/inplace_apply_power_sign/README.md">inplace_apply_power_sign</a></td>4954 <td><a href="../../optim/inplace_apply_power_sign/README.md">inplace_apply_power_sign</a></td>
4956 <td>✓</td>4955 <td>✓</td>