Star1486
3125
代码
代码
Issues321
Pull Requests704
Actions
项目讨论
Wiki
分析
Star1486
3125
已合并
fix: Modify the AIDD docs issue #9326
zwj223创建于 7月30日
fix: Modify the AIDD docs issue #9326
已合并
zwj223创建于 7月30日
从已删除 :master合入到cann/ops-transformermaster
共 22 个文件变更+102-98
MCHANGELOG.md+1-1
@@ -48,4 +48,4 @@ ops-transformer算子首个Beta版本8.5.0-beta.1 现已发布。本版本引入
48 48 
49- mc2通信域支持统一的torch.dist.group问题。([Issue47](https://gitcode.com/cann/ops-transformer/issues/47))49- mc2通信域支持统一的torch.dist.group问题。([Issue47](https://gitcode.com/cann/ops-transformer/issues/47))
50- add_example样例算子执行调用问题修复。([Issue174](https://gitcode.com/cann/ops-transformer/issues/174))50- add_example样例算子执行调用问题修复。([Issue174](https://gitcode.com/cann/ops-transformer/issues/174))
51-- 修复install_deps.sh脚本不支持openEuler系统问题。([Issue255](https://gitcode.com/cann/ops-transformer/issues/255))51+- 修复install_deps.sh脚本不支持openEuler系统的问题。([Issue255](https://gitcode.com/cann/ops-transformer/issues/255))
Mdocs/CONTRIBUTING_DOCS.md+1-1
@@ -84,7 +84,7 @@
84 - 图片：84 - 图片：
85 - 常用格式：推荐png格式，风格尽量与已有图片保持一致。85 - 常用格式：推荐png格式，风格尽量与已有图片保持一致。
86 - 分辨率与清晰度：需清晰且尺寸适中，避免模糊或过度压缩。86 - 分辨率与清晰度：需清晰且尺寸适中，避免模糊或过度压缩。
87- - 文件大小：单张图片不建议超过10M。87+ - 文件大小：单张图片不建议超过10MB。
88 - 版权：所有引用的图片、文献等资源，请确保合规性。88 - 版权：所有引用的图片、文献等资源，请确保合规性。
89 89 
90## 获取帮助90## 获取帮助
Mdocs/QUICKSTART.md+2-2
@@ -2,7 +2,7 @@
2 2 
3## 使用须知3## 使用须知
4 4 
5-本指南旨在帮助您快速上手CANN和`ops-transformer`算子仓的使用。为方便快速了解算子开发全流程，以**AddExample**算子为实践对象，算子源码位于`ops-transformer/examples/add_example`，操作流程如下：5+本指南旨在帮助开发者快速上手CANN和`ops-transformer`算子仓的使用。为方便快速了解算子开发全流程，以**AddExample**算子为实践对象，算子源码位于`ops-transformer/examples/add_example`，操作流程如下：
6 6 
71. **[前提条件](../README.md)**：参考项目README完成环境准备和源码下载，此处不再赘述。快速入门场景**推荐CANNLab或Docker部署**，操作简单。71. **[前提条件](../README.md)**：参考项目README完成环境准备和源码下载，此处不再赘述。快速入门场景**推荐CANNLab或Docker部署**，操作简单。
8 8 
@@ -269,4 +269,4 @@ int main() {
269 269 
270## 结语270## 结语
271 271 
272-体验完上述流程，您已基本完成算子开发过程，如果您想进一步贡献新算子或学习更多高阶开发、调试等技能，请访问本项目README学习[进阶教程](../README.md#学习教程)和[贡献指南](../README.md#相关信息)等。272+体验完上述流程，已基本完成算子开发过程，如果想进一步贡献新算子或学习更多高阶开发、调试等技能，请访问本项目README学习[进阶教程](../README.md#学习教程)和[贡献指南](../README.md#相关信息)等。
Mdocs/zh/ascend950_op_list.md+8-8
@@ -124,7 +124,7 @@ Ascend 950支持的算子分类和算子列表如下：
124 <td>✗</td>124 <td>✗</td>
125 <td>✓</td>125 <td>✓</td>
126 <td>AI Core</td>126 <td>AI Core</td>
127- <td>在Sparse Flash Attention的基础上支持了[Per-Token-Head-Tile-128量化]输入。</td>127+ <td>在Sparse Flash Attention的基础上支持了`Per-Token-Head-Tile-128量化`输入。</td>
128 </tr>128 </tr>
129 <tr>129 <tr>
130 <td>attention</td>130 <td>attention</td>
@@ -374,7 +374,7 @@ Ascend 950支持的算子分类和算子列表如下：
374 <td>✓</td>374 <td>✓</td>
375 <td>✓</td>375 <td>✓</td>
376 <td>AI Core</td>376 <td>AI Core</td>
377- <td>完成路由专家AlltoAllv、Permute、GroupedMatMul融合并实现与共享专家MatMul并行融合，先通信后计算。</td>377+ <td>完成路由专家AlltoAllV、Permute、GroupedMatMul融合并实现与共享专家MatMul并行融合，先通信后计算。</td>
378 </tr>378 </tr>
379 <tr>379 <tr>
380 <td>mc2</td>380 <td>mc2</td>
@@ -384,7 +384,7 @@ Ascend 950支持的算子分类和算子列表如下：
384 <td>✓</td>384 <td>✓</td>
385 <td>✓</td>385 <td>✓</td>
386 <td>AI Core</td>386 <td>AI Core</td>
387- <td>完成路由专家AlltoAllv、Permute、QuantGroupedMatMul融合并实现与共享专家MatMul并行融合。</td>387+ <td>完成路由专家AlltoAllV、Permute、QuantGroupedMatMul融合并实现与共享专家MatMul并行融合。</td>
388 </tr>388 </tr>
389 <tr>389 <tr>
390 <td>mc2</td>390 <td>mc2</td>
@@ -424,7 +424,7 @@ Ascend 950支持的算子分类和算子列表如下：
424 <td>✓</td>424 <td>✓</td>
425 <td>✓</td>425 <td>✓</td>
426 <td>AI Core</td>426 <td>AI Core</td>
427- <td>完成路由专家GroupedMatMul、Unpermute、AlltoAllv融合并实现与共享专家MatMul并行融合，先计算后通信。</td>427+ <td>完成路由专家GroupedMatMul、Unpermute、AlltoAllV融合并实现与共享专家MatMul并行融合，先计算后通信。</td>
428 </tr>428 </tr>
429 <tr>429 <tr>
430 <td>mc2</td>430 <td>mc2</td>
@@ -554,7 +554,7 @@ Ascend 950支持的算子分类和算子列表如下：
554 <td>✓</td>554 <td>✓</td>
555 <td>✓</td>555 <td>✓</td>
556 <td>AI Core</td>556 <td>AI Core</td>
557- <td>对Token数据进行量化（可选），当存在TP域通信时，先进行EP（Expert Parallelism）域的AllToAllV通信，再进行TP（Tensor Parallelism）域的AllGatherV通信；当不存在TP域通信时，进行EP（Expert Parallelism）域的AllToAllV通信。</td>557+ <td>对Token数据进行量化（可选），当存在TP域通信时，先进行EP（Expert Parallelism）域的AlltoAllV通信，再进行TP（Tensor Parallelism）域的AllGatherV通信；当不存在TP域通信时，进行EP（Expert Parallelism）域的AlltoAllV通信。</td>
558 </tr>558 </tr>
559 <tr>559 <tr>
560 <td>mc2</td>560 <td>mc2</td>
@@ -584,7 +584,7 @@ Ascend 950支持的算子分类和算子列表如下：
584 <td>✓</td>584 <td>✓</td>
585 <td>✓</td>585 <td>✓</td>
586 <td>AI Core</td>586 <td>AI Core</td>
587- <td>对Token数据进行量化（可选），当存在TP域通信时，先进行EP（Expert Parallelism）域的AllToAllV通信，再进行TP（Tensor Parallelism）域的AllGatherV通信；当不存在TP域通信时，进行EP（Expert Parallelism）域的AllToAllV通信。</td>587+ <td>对Token数据进行量化（可选），当存在TP域通信时，先进行EP（Expert Parallelism）域的AlltoAllV通信，再进行TP（Tensor Parallelism）域的AllGatherV通信；当不存在TP域通信时，进行EP（Expert Parallelism）域的AlltoAllV通信。</td>
588 </tr>588 </tr>
589 <tr>589 <tr>
590 <td>mc2</td>590 <td>mc2</td>
@@ -594,7 +594,7 @@ Ascend 950支持的算子分类和算子列表如下：
594 <td>✓</td>594 <td>✓</td>
595 <td>✓</td>595 <td>✓</td>
596 <td>AI Core</td>596 <td>AI Core</td>
597- <td>对Token数据进行量化（可选），当存在TP域通信时，先进行EP（Expert Parallelism）域的AllToAllV通信，再进行TP（Tensor Parallelism）域的AllGatherV通信；当不存在TP域通信时，进行EP（Expert Parallelism）域的AllToAllV通信。</td>597+ <td>对Token数据进行量化（可选），当存在TP域通信时，先进行EP（Expert Parallelism）域的AlltoAllV通信，再进行TP（Tensor Parallelism）域的AllGatherV通信；当不存在TP域通信时，进行EP（Expert Parallelism）域的AlltoAllV通信。</td>
598 </tr>598 </tr>
599 <tr>599 <tr>
600 <td>mc2</td>600 <td>mc2</td>
@@ -624,7 +624,7 @@ Ascend 950支持的算子分类和算子列表如下：
624 <td>✓</td>624 <td>✓</td>
625 <td>✓</td>625 <td>✓</td>
626 <td>AI Core</td>626 <td>AI Core</td>
627- <td>完成路由专家QuantGroupedMatMul、Unpermute、AlltoAllv融合并实现与共享专家MatMul并行融合。</td>627+ <td>完成路由专家QuantGroupedMatMul、Unpermute、AlltoAllV融合并实现与共享专家MatMul并行融合。</td>
628 </tr>628 </tr>
629 <tr>629 <tr>
630 <td>mc2</td>630 <td>mc2</td>
Mdocs/zh/context/aclnn_return_code.md+1-1
@@ -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 
Mdocs/zh/context/compile_and_run_sample.md+13-13
@@ -8,7 +8,7 @@
8 8 
9## 编译前准备9## 编译前准备
10 10 
11-本章以开发和运行环境合设场景为例，即带AI处理器的机器既作为开发环境又作为运行环境。该场景下，代码开发和代码运行在同一台机器上。这里以**FlashAttentionScore算子**为例，其他算子的调用逻辑、流程、编译脚本与FlashAttentionScore算子大致一样，请根据实际情况自行修改API调用脚本（\*.cpp）和编译脚本(CMakeLists)。11+本章以开发和运行环境合设场景为例，即带AI处理器的机器既作为开发环境又作为运行环境。该场景下，代码开发和代码运行在同一台机器上。这里以**FlashAttentionScore算子**为例，其他算子的调用逻辑、流程、编译脚本与FlashAttentionScore算子大致一样，请根据实际情况自行修改API调用脚本(\*.cpp)和编译脚本(CMakeLists)。
12 12 
13- **示例代码**13- **示例代码**
14 14 
@@ -124,21 +124,21 @@
124 124 
125 其中${INSTALL_DIR}为CANN软件安装后文件存储路径，请根据实际情况替换。125 其中${INSTALL_DIR}为CANN软件安装后文件存储路径，请根据实际情况替换。
126 3. 编译并运行。126 3. 编译并运行。
127- - 进入CMakeLists.txt所在目录，执行如下命令，新建build目录存放生成的编译文件。127+ - 进入CMakeLists.txt所在目录，执行如下命令，新建build目录存放生成的编译文件。
128 128 
129- ```bash129+ ```bash
130- mkdir -p build130+ mkdir -p build
131- ```131+ ```
132 132 
133- - 进入build目录，执行cmake命令编译，再执行make命令生成可执行文件。133+ - 进入build目录，执行cmake命令编译，再执行make命令生成可执行文件。
134 134 
135- ```bash135+ ```bash
136- cd build136+ cd build
137- cmake ../ -DCMAKE_CXX_COMPILER=g++ -DCMAKE_SKIP_RPATH=TRUE137+ cmake ../ -DCMAKE_CXX_COMPILER=g++ -DCMAKE_SKIP_RPATH=TRUE
138- make138+ make
139- ```139+ ```
140 140 
141- 编译成功后，会在build目录的bin文件夹下生成opapi\_test可执行文件。141+ 编译成功后，会在build目录的bin文件夹下生成opapi\_test可执行文件。
142 142 
143 - 进入bin目录，运行可执行文件opapi\_test。143 - 进入bin目录，运行可执行文件opapi\_test。
144 144 
@@ -179,5 +179,5 @@
179 179 
180 ```text180 ```text
181 aclnnFlashAttentionScoreGetWorkspaceSize failed. ERROR: 161001181 aclnnFlashAttentionScoreGetWorkspaceSize failed. ERROR: 161001
182- [ERROR msg][PID:xxxx] xxx(timesamp) AclNN_Parameter_Error(EZ1001): The query cannot be nullptr.182+ [ERROR msg][PID:xxxx] xxx(timestamp) AclNN_Parameter_Error(EZ1001): The query cannot be nullptr.
183 ```183 ```
Mdocs/zh/context/data_format.md+1-1
@@ -12,7 +12,7 @@
12 12 
13>**说明：**13>**说明：**
14>14>
15->- 两段式接口参数说明时，为简化描述，**将原始数据格式“ACL\_FORMAT\_XXXX_”简写为“_XXXX_”**。15+>- 两段式接口参数说明时，为简化描述，**将原始数据格式“ACL\_FORMAT\_XXXX_”简写为“\_XXXX_”**。
16>- 数据格式中各维度含义：N（Batch）表示批量大小、H（Height）表示特征图高度、W（Width）表示特征图宽度、C（Channels）表示特征图通道、D（Depth）表示特征图深度、L（Length）表示特征图长度。16>- 数据格式中各维度含义：N（Batch）表示批量大小、H（Height）表示特征图高度、W（Width）表示特征图宽度、C（Channels）表示特征图通道、D（Depth）表示特征图深度、L（Length）表示特征图长度。
17 17 
18## 常见数据格式18## 常见数据格式
Mdocs/zh/context/deduction_relationship.md+1-1
@@ -14,7 +14,7 @@ aclTensor支持的数据类型参见[数据类型](data_type.md)，其中部分
14>14>
15>- 为方便描述，表格中使用的数据类型是**简写形式**，代表的含义：ACL\_FLOAT\(f32\)、ACL\_FLOAT16\(f16\)、ACL\_DOUBLE\(f64\)、ACL\_BF16\(bf16\)、ACL\_INT8\(s8\)、ACL\_UINT8\(u8\)、ACL\_INT16\(s16\)、ACL\_UINT16\(u16\)、ACL\_INT32\(s32\)、ACL\_UINT32\(u32\)、ACL\_INT64\(s64\)、ACL\_UINT64\(u64\)、ACL\_BOOL\(bool\)、ACL\_COMPLEX32\(c32\)、ACL\_COMPLEX64\(c64\)、ACL\_COMPLEX128\(c128\)。15>- 为方便描述，表格中使用的数据类型是**简写形式**，代表的含义：ACL\_FLOAT\(f32\)、ACL\_FLOAT16\(f16\)、ACL\_DOUBLE\(f64\)、ACL\_BF16\(bf16\)、ACL\_INT8\(s8\)、ACL\_UINT8\(u8\)、ACL\_INT16\(s16\)、ACL\_UINT16\(u16\)、ACL\_INT32\(s32\)、ACL\_UINT32\(u32\)、ACL\_INT64\(s64\)、ACL\_UINT64\(u64\)、ACL\_BOOL\(bool\)、ACL\_COMPLEX32\(c32\)、ACL\_COMPLEX64\(c64\)、ACL\_COMPLEX128\(c128\)。
16>- 表格里表头和最左侧一列分别表示待推导的两个输入数据类型，表格中对应位置表示推导出的数据类型。16>- 表格里表头和最左侧一列分别表示待推导的两个输入数据类型，表格中对应位置表示推导出的数据类型。
17->- 表中叉号(×)表示这两种类型不能进行推导计算。17+>- 表中叉号（×）表示这两种类型不能进行推导计算。
18 18 
19**表1** 数据类型推导关系19**表1** 数据类型推导关系
20 20 
Mdocs/zh/context/two_phase_api.md+2-2
@@ -2,7 +2,7 @@
2 2 
3基于单算子API执行方式调用算子API时，通常分为“两段式”，样式形如：3基于单算子API执行方式调用算子API时，通常分为“两段式”，样式形如：
4 4 
5-```Cpp5+```cpp
6aclnnStatus aclxxXxxGetWorkspaceSize(const aclTensor *src, ..., aclTensor *out, ..., uint64_t *workspaceSize, aclOpExecutor **executor);6aclnnStatus aclxxXxxGetWorkspaceSize(const aclTensor *src, ..., aclTensor *out, ..., uint64_t *workspaceSize, aclOpExecutor **executor);
7aclnnStatus aclxxXxx(void *workspace, uint64_t workspaceSize, aclOpExecutor *executor, aclrtStream stream);7aclnnStatus aclxxXxx(void *workspace, uint64_t workspaceSize, aclOpExecutor *executor, aclrtStream stream);
8```8```
@@ -16,7 +16,7 @@ aclnnStatus aclxxXxx(void *workspace, uint64_t workspaceSize, aclOpExecutor *exe
16> - workspace是指除输入/输出外，算子在NPU上完成计算所需要的临时内存，workspaceSize表示临时内存的大小。16> - workspace是指除输入/输出外，算子在NPU上完成计算所需要的临时内存，workspaceSize表示临时内存的大小。
17> - 第二段接口aclxxXxx(...)不能重复调用，如下调用方式会出现异常：17> - 第二段接口aclxxXxx(...)不能重复调用，如下调用方式会出现异常：
18>18>
19-> ```Cpp19+> ```cpp
20> aclxxXxxGetWorkspaceSize(...)20> aclxxXxxGetWorkspaceSize(...)
21> aclxxXxx(...)21> aclxxXxx(...)
22> aclxxXxx(...)22> aclxxXxx(...)
Mdocs/zh/debug/cann_sim.md+10-10
@@ -30,7 +30,7 @@ CANN Simulator集成在CANN toolkit包里，参考[环境部署](../install/quic
30 30 
31# 快速开始31# 快速开始
32 32 
33-下面以[add_examples](../../../examples/add_example/)为例，对算子仿真进行详细说明33+下面以[add_example](../../../examples/add_example/)为例，对算子仿真进行详细说明
34 34 
35## 算子编译35## 算子编译
36 36 
@@ -104,19 +104,19 @@ 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 
1143. 命令完成后，会在默认路径或指定的“output”目录下生成以“cannsim_{timestamp}_${user_app}”命名的文件夹，结构示例如下：1143. 命令完成后，会在默认路径或指定的“output”目录下生成以“cannsim_{timestamp}_${user_app}”命名的文件夹，结构示例如下：
115 115 
116-```bash116+ ```bash
117-├─cannsim_{timestamp}_${user_app}117+ ├─cannsim_{timestamp}_${user_app}
118-├── cannsim.log118+ ├── cannsim.log
119-```119+ ```
120 120 
1214. 用户可以获取算子执行结果，并进行精度的对比，结果展示在cannsim.log，示例如下1214. 用户可以获取算子执行结果，并进行精度的对比，结果展示在cannsim.log，示例如下
122 122 
@@ -154,14 +154,14 @@ 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'
166 ```166 ```
167 167 
Mdocs/zh/debug/op_debug_prof.md+9-9
@@ -15,7 +15,7 @@
15 $HOME/ascend/log/debug/plog/plog-pid_*.log15 $HOME/ascend/log/debug/plog/plog-pid_*.log
16 ```16 ```
17 17 
18- 开启环境变量ASCEND_SLOG_PRINT_TO_STDOUT可以将log日志直接打屏显示(1:开启打屏，0：关闭打屏)，配置示例如下：18+ 开启环境变量ASCEND_SLOG_PRINT_TO_STDOUT可以将log日志直接打屏显示(1：开启打屏，0：关闭打屏)，配置示例如下：
19 19 
20 ```bash20 ```bash
21 export ASCEND_SLOG_PRINT_TO_STDOUT=121 export ASCEND_SLOG_PRINT_TO_STDOUT=1
@@ -27,15 +27,15 @@
27 27 
28 通过aclGetRecentErrMsg接口（参见[《Runtime运行时API》](https://hiascend.com/document/redirect/CannCommunityRuntimeApi)）获取aclnn接口调用过程中的异常信息，使用方法如下：28 通过aclGetRecentErrMsg接口（参见[《Runtime运行时API》](https://hiascend.com/document/redirect/CannCommunityRuntimeApi)）获取aclnn接口调用过程中的异常信息，使用方法如下：
29 29 
30- ```cpp30+ ```cpp
31- printf(aclGetRecentErrMsg());31+ printf("%s",aclGetRecentErrMsg());
32- ```32+ ```
33 33 
34 打印错误信息样例如下：34 打印错误信息样例如下：
35 35 
36- ```text36+ ```text
37- [PID:646612] 2026-01-24-11:53:44.671.727 AclNN_Parameter_Error(EZ1001): Expected a proper Tensor but got null for argument addmmTennsor.self.37+ [PID:646612] 2026-01-24-11:53:44.671.727 AclNN_Parameter_Error(EZ1001): Expected a proper Tensor but got null for argument addmmTensor.self.
38- ```38+ ```
39 39 
40### 2. Kernel调试40### 2. Kernel调试
41 41 
@@ -231,5 +231,5 @@
231 msprof op simulator --output=$PWD/pipeline_auto --kernel-name "AddExample" ./test_aclnn_add_example231 msprof op simulator --output=$PWD/pipeline_auto --kernel-name "AddExample" ./test_aclnn_add_example
232 ```232 ```
233 233 
234- 采集结果在本项目`$PWD/pipeline_auto/OPPROF_**`目录中。234+ 采集结果在本项目`$PWD/pipeline_auto/OPPROF_**`目录中。
235- 其中流水相关文件路径为`OPPROF**/simulator/visualize_data.bin`，可以借助[MindStudio Insight](https://www.hiascend.com/document/redirect/MindStudioInsight)工具中“基础操作 > 导入数据”章节查看如何导入流水数据。235+ 其中流水相关文件路径为`OPPROF_*/simulator/visualize_data.bin`，可以借助[MindStudio Insight](https://www.hiascend.com/document/redirect/MindStudioInsight)工具中“基础操作 > 导入数据”章节查看如何导入流水数据。
Mdocs/zh/develop/aicore_develop_guide.md+4-4
@@ -88,7 +88,7 @@ foreach(EXAMPLES_OP_NAME ${ASCEND_OP_NAME})
88 list(APPEND OP_DIR_LIST ${CMAKE_CURRENT_SOURCE_DIR}/examples/mc2/${EXAMPLES_OP_NAME})88 list(APPEND OP_DIR_LIST ${CMAKE_CURRENT_SOURCE_DIR}/examples/mc2/${EXAMPLES_OP_NAME})
89 # 在examples目录下新增算子分类时，参考mc2目录增加命令语句如下：89 # 在examples目录下新增算子分类时，参考mc2目录增加命令语句如下：
90 # elseif(IS_DIRECTORY ${EXAMPLES_${op_class}_DIR})90 # elseif(IS_DIRECTORY ${EXAMPLES_${op_class}_DIR})
91- # add_subdirectory(examples/${op_class}/${EXAMPLES_OP_NAME}")91+ # add_subdirectory(examples/${op_class}/${EXAMPLES_OP_NAME})
92 # list(APPEND OP_DIR_LIST ${CMAKE_CURRENT_SOURCE_DIR}/examples/${op_class}/${EXAMPLES_OP_NAME})92 # list(APPEND OP_DIR_LIST ${CMAKE_CURRENT_SOURCE_DIR}/examples/${op_class}/${EXAMPLES_OP_NAME})
93 endif()93 endif()
94endforeach()94endforeach()
@@ -286,7 +286,7 @@ Kernel入口文件，包含主函数和调度逻辑。
286 286 
287如需查看详细实现，请参考[add_example.cpp](../../../examples/add_example/op_kernel/add_example.cpp)。287如需查看详细实现，请参考[add_example.cpp](../../../examples/add_example/op_kernel/add_example.cpp)。
288 288 
289-```CPP289+```cpp
290// 1、核函数定义290// 1、核函数定义
291// schMode是一个模板参数，用于支持不同数据类型（如float和int32）的计算路径291// schMode是一个模板参数，用于支持不同数据类型（如float和int32）的计算路径
292// __global__ __aicore__表示该函数是个全局函数，可以在AI Core上执行292// __global__ __aicore__表示该函数是个全局函数，可以在AI Core上执行
@@ -315,7 +315,7 @@ __global__ __aicore__ void add_example(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR
315 315 
316如需查看详细实现，请参考[add_example.h](../../../examples/add_example/op_kernel/add_example.h)。316如需查看详细实现，请参考[add_example.h](../../../examples/add_example/op_kernel/add_example.h)。
317 317 
318-```C++318+```cpp
319// 2、定义Kernel类319// 2、定义Kernel类
320template <typename T>320template <typename T>
321class AddExample321class AddExample
@@ -462,7 +462,7 @@ __aicore__ inline void AddExample<T>::Process()
462 462 
463### UT验证463### UT验证
464 464 
465-主要交付件代码开发过程中，可通过UT验证方式进行快速验证，无需编译部署算子包。465+在主要交付件代码开发过程中，可通过UT验证方式进行快速验证，无需编译部署算子包。
466 466 
467UT目录结构如下，需用户手动创建：467UT目录结构如下，需用户手动创建：
468 468 
Mdocs/zh/develop/aicpu_develop_guide.md+5-5
@@ -177,9 +177,9 @@ uint32_t AddExampleCpuKernel::Compute(CpuKernelContext& ctx) {
177 case DT_FLOAT:177 case DT_FLOAT:
178 return AddCompute<float>(...);178 return AddCompute<float>(...);
179 case DT_INT32:179 case DT_INT32:
180- return AddCompute<int32>(...);180+ return AddCompute<int32_t>(...);
181 ....181 ....
182- default : return PARAM_INVALID;182+ default : return kParamInvalid;
183 }183 }
184}184}
185 185 
@@ -190,11 +190,11 @@ REGISTER_CPU_KERNEL(kAddExample, AddExampleCpuKernel);
190 190 
191## aclnn适配191## aclnn适配
192 192 
193-通常算子开发和编译完成后，会自动生成aclnn接口（一套基于C的API），无需做其他配置，可直接在应用程序中调用aclnn接口实现调用算子。193+通常算子开发和编译完成后，会自动生成aclnn接口（一套基于C的API），无需做其他配置，可直接在应用程序中调用aclnn接口实现算子调用。
194 194 
195## 编译部署195## 编译部署
196 196 
197-算子开发完成后，需对算子工程进行编译，生成自定义算子安装包\*\.run，具体操作如下：197+算子开发完成后，需对算子工程进行编译，生成自定义算子安装包 `*.run`，具体操作如下：
198 198 
1991. **准备工作。**1991. **准备工作。**
200 200 
@@ -234,7 +234,7 @@ REGISTER_CPU_KERNEL(kAddExample, AddExampleCpuKernel);
234 自定义算子包安装后在```${ASCEND_HOME_PATH}/opp/vendors/custom_transformer/scripts```目录会生成`uninstall.sh`，通过该脚本可卸载自定义算子包，命令如下：234 自定义算子包安装后在```${ASCEND_HOME_PATH}/opp/vendors/custom_transformer/scripts```目录会生成`uninstall.sh`，通过该脚本可卸载自定义算子包，命令如下：
235 235 
236 ```bash236 ```bash
237- bash ${ASCEND_HOME_PATH}/opp/vendors/custom_transformer/scripts/uninstall.sh237+ bash ${ASCEND_HOME_PATH}/opp/vendors/${vendor_name}_transformer/scripts/uninstall.sh
238 ```238 ```
239 239 
240## 算子验证240## 算子验证
Mdocs/zh/develop/cross_platform_migration_guide.md+11-7
@@ -99,8 +99,8 @@
99 <td>所有使用int4_t的算子需要切换到支持的数据类型（如int8），并更新量化解算逻辑</td>99 <td>所有使用int4_t的算子需要切换到支持的数据类型（如int8），并更新量化解算逻辑</td>
100 </tr>100 </tr>
101 <tr>101 <tr>
102- <td>不支持4：2稀疏矩阵计算</td>102+ <td>不支持4:2稀疏矩阵计算</td>
103- <td>原依赖4：2稀疏特性提速的kernel需要改为稠密或其他支持的稀疏策略，并更新性能预期说明</td>103+ <td>原依赖4:2稀疏特性提速的kernel需要改为稠密或其他支持的稀疏策略，并更新性能预期说明</td>
104 </tr>104 </tr>
105 <tr>105 <tr>
106 <td rowspan="1">存储单元</td>106 <td rowspan="1">存储单元</td>
@@ -168,7 +168,7 @@ TBuf<QuePosition::VECCALC> indexBuf_;
168for (int64_t j = 0; j < rows; j++) {168for (int64_t j = 0; j < rows; j++) {
169 INDICES_T index = GetIndex(yIdx, indiceEndIdx); // 标量读取索引169 INDICES_T index = GetIndex(yIdx, indiceEndIdx); // 标量读取索引
170 int64_t xIndex = index * tilingData_->innerSize;170 int64_t xIndex = index * tilingData_->innerSize;
171- DataCopyPad(xLocal[j * colsAlign], xGm[offset], dataCoptExtParams, dataCopyPadExtParams); // 批量连续数搬入171+ DataCopyPad(xLocal[j * colsAlign], xGm[offset], dataCopyExtParams, dataCopyPadExtParams); // 批量连续数搬入
172}172}
173inQueue_.EnQue<int8_t>(xLocal); // 入队等待输出173inQueue_.EnQue<int8_t>(xLocal); // 入队等待输出
174```174```
@@ -231,6 +231,8 @@ __simd_vf__ __aicore__ void GenIndexBuf(ubuf int32_t* helpAddr, int32_t colFacto
231 AscendC::MicroAPI::RegTensor<int32_t> v0;231 AscendC::MicroAPI::RegTensor<int32_t> v0;
232 AscendC::MicroAPI::RegTensor<int32_t> v1;232 AscendC::MicroAPI::RegTensor<int32_t> v1;
233 AscendC::MicroAPI::RegTensor<int32_t> vd1;233 AscendC::MicroAPI::RegTensor<int32_t> vd1;
234+ AscendC::MicroAPI::RegTensor<int32_t> vd2;
235+ AscendC::MicroAPI::RegTensor<int32_t> vd3;
234 236 
235 // 创建全量掩码237 // 创建全量掩码
236 AscendC::MicroAPI::MaskReg preg =238 AscendC::MicroAPI::MaskReg preg =
@@ -251,10 +253,12 @@ __simd_vf__ __aicore__ void GenIndexBuf(ubuf int32_t* helpAddr, int32_t colFacto
251 253 
252```cpp254```cpp
253// 动态掩码：处理尾部不完整数据255// 动态掩码：处理尾部不完整数据
254-__simd_vf__ __aicore__ void GatherProcess(ubuf int8_t* curYAddr, uint16_t repeatimes, uint16_t computeSize)256+__simd_vf__ __aicore__ void GatherProcess(ubuf int8_t* curXAddr, ubuf int8_t* curYAddr, uint16_t repeatTimes, uint16_t computeSize)
255{257{
256 MicroAPI::RegTensor<int8_t> vregTemp;258 MicroAPI::RegTensor<int8_t> vregTemp;
257 MicroAPI::MaskReg preg;259 MicroAPI::MaskReg preg;
260+ // sreg为剩余待处理元素（普通unit32_t标量计数）
261+ unit32_t sreg = static_cast<unit32_t>(repeatimes)*computeSize;
258 262 
259 for (uint16_t r = 0; r < repeatTimes; r++) {263 for (uint16_t r = 0; r < repeatTimes; r++) {
260 // 根据剩余元素数更新掩码264 // 根据剩余元素数更新掩码
@@ -343,7 +347,7 @@ Ascend 950新架构引入UB2L1 & L0C2UB间的直连通路，实现矩阵计算
343```cpp347```cpp
344// 1. 新增: 搬入接口增加UB2L1的Nd2Nz搬入，支持Src&Dst都是LocalTensor的形式348// 1. 新增: 搬入接口增加UB2L1的Nd2Nz搬入，支持Src&Dst都是LocalTensor的形式
345template <typename T>349template <typename T>
346-__aicore__ inline void DataCopy(const LocalTensor<T>& dst, const LocalTensor<T>& src, const Nd2NzParams& intriParams)；350+__aicore__ inline void DataCopy(const LocalTensor<T>& dst, const LocalTensor<T>& src, const Nd2NzParams& intriParams);
347 351 
348// 2. 新增: 搬出接口增加L0C2UB的搬出，支持直接从L0C搬出到UB,支持Src&Dst都是LocalTensor的形式352// 2. 新增: 搬出接口增加L0C2UB的搬出，支持直接从L0C搬出到UB,支持Src&Dst都是LocalTensor的形式
349template <typename T, typename U, const FixpipeConfig& config = CFG_ROW_MAJOR>353template <typename T, typename U, const FixpipeConfig& config = CFG_ROW_MAJOR>
@@ -356,9 +360,9 @@ struct FixpipeParamsC310 {
356 360 
357// 3. 能力增强: 核间同步接口新增模式3361// 3. 能力增强: 核间同步接口新增模式3
358template <uint8_t modeId, pipe_t pipe>362template <uint8_t modeId, pipe_t pipe>
359-__aicore__ inline void CrossCoreSetFlag(uint16_t flagId)363+__aicore__ inline void CrossCoreSetFlag(uint16_t flagId);
360template <uint8_t modeId = 0, pipe_t pipe = PIPE_S>364template <uint8_t modeId = 0, pipe_t pipe = PIPE_S>
361-__aicore__ inline void CrossCoreWaitFlag(uint16_t flagId)365+__aicore__ inline void CrossCoreWaitFlag(uint16_t flagId);
362 366 
363```367```
364 368 
Mdocs/zh/develop/graph_develop_guide.md+6-6
@@ -26,7 +26,7 @@ InferShape函数的作用是根据输入的shape推导输出的shape。
26 26 
27示例如下，`AddExample`算子完整代码请参考`examples/add_example/op_host`下[add_example_infershape.cpp](../../../examples/add_example/op_host/add_example_infershape.cpp)。27示例如下，`AddExample`算子完整代码请参考`examples/add_example/op_host`下[add_example_infershape.cpp](../../../examples/add_example/op_host/add_example_infershape.cpp)。
28 28 
29-```C++29+```cpp
30// AddExample算子逻辑是两个数相加，因此输出shape与输入shape一致30// AddExample算子逻辑是两个数相加，因此输出shape与输入shape一致
31static ge::graphStatus InferShapeAddExample(gert::InferShapeContext* context)31static ge::graphStatus InferShapeAddExample(gert::InferShapeContext* context)
32{32{
@@ -44,7 +44,7 @@ static ge::graphStatus InferShapeAddExample(gert::InferShapeContext* context)
44 int64_t dim = xShape->GetDim(i);44 int64_t dim = xShape->GetDim(i);
45 yShape->SetDim(i, dim);45 yShape->SetDim(i, dim);
46 }46 }
47- ....47+ ...
48}48}
49// inferShape注册49// inferShape注册
50IMPL_OP_INFERSHAPE(AddExample).InferShape(InferShapeAddExample);50IMPL_OP_INFERSHAPE(AddExample).InferShape(InferShapeAddExample);
@@ -56,16 +56,16 @@ InferDataType函数的作用是根据输入的DataType推导输出的DataType。
56 56 
57示例如下，`AddExample`算子完整代码请参考`examples/add_example/op_graph`下[add_example_graph_infer.cpp](../../../examples/add_example/op_graph/add_example_graph_infer.cpp)。57示例如下，`AddExample`算子完整代码请参考`examples/add_example/op_graph`下[add_example_graph_infer.cpp](../../../examples/add_example/op_graph/add_example_graph_infer.cpp)。
58 58 
59-```C++59+```cpp
60// AddExample算子逻辑是两个数相加，因此输出dataType与输入dataType一致60// AddExample算子逻辑是两个数相加，因此输出dataType与输入dataType一致
61static ge::graphStatus InferDataTypeAddExample(gert::InferDataTypeContext* context)61static ge::graphStatus InferDataTypeAddExample(gert::InferDataTypeContext* context)
62{62{
63- ....63+ ...
64 // 获取输入的dataType64 // 获取输入的dataType
65 ge::DataType sizeDtype = context->GetInputDataType(IDX_0);65 ge::DataType sizeDtype = context->GetInputDataType(IDX_0);
66 // 将输入dataType设置到输出66 // 将输入dataType设置到输出
67 context->SetOutputDataType(IDX_0, sizeDtype);67 context->SetOutputDataType(IDX_0, sizeDtype);
68- ....68+ ...
69}69}
70 70 
71// 注册InferDataType71// 注册InferDataType
@@ -109,7 +109,7 @@ IMPL_OP(AddExample).InferDataType(InferDataTypeAddExample);
109 109 
110示例代码如下，展示了如何注册`AddExample`算子：110示例代码如下，展示了如何注册`AddExample`算子：
111 111 
112-```CPP112+```cpp
113REG_OP(AddExample)113REG_OP(AddExample)
114 .INPUT(x1, TensorType({DT_FLOAT}))114 .INPUT(x1, TensorType({DT_FLOAT}))
115 .INPUT(x2, TensorType({DT_FLOAT}))115 .INPUT(x2, TensorType({DT_FLOAT}))
Mdocs/zh/install/compile.md+1-1
@@ -81,7 +81,7 @@
81 81 
82 自定义算子包安装路径为```${ASCEND_HOME_PATH}/opp/vendors```，\$\{ASCEND\_HOME\_PATH\}已通过环境变量配置，表示CANN toolkit包安装路径，一般为\$\{install\_path\}/cann。82 自定义算子包安装路径为```${ASCEND_HOME_PATH}/opp/vendors```，\$\{ASCEND\_HOME\_PATH\}已通过环境变量配置，表示CANN toolkit包安装路径，一般为\$\{install\_path\}/cann。
83 83 
84- 如果部署算子包时通过配置--install-path参数指定了算子包的安装目录，则在使用自定义算子前，需要执行```source ${install-path}/vendors/${vendor_name}/bin/set_env.bash```命令，set_env.bash脚本中将自定义算子包的安装路径追加到环境变量ASCEND_CUSTOM_OPP_PATH中，使自定义算子在当前环境中生效。84+ 如果部署算子包时通过配置--install-path参数指定了算子包的安装目录，则在使用自定义算子前，需要执行```source ${install_path}/vendors/${vendor_name}/bin/set_env.bash```命令，set_env.bash脚本中将自定义算子包的安装路径追加到环境变量ASCEND_CUSTOM_OPP_PATH中，使自定义算子在当前环境中生效。
85 85 
863. **（可选）删除自定义算子包**863. **（可选）删除自定义算子包**
87 87 
Mdocs/zh/install/dir_structure.md+1-1
@@ -46,7 +46,7 @@
46│ │ │ ├── ${op_name}_infershape.cpp # 可选，InferShape实现，根据算子形状推导输出shape，若未配置则输出shape与输入shape一样46│ │ │ ├── ${op_name}_infershape.cpp # 可选，InferShape实现，根据算子形状推导输出shape，若未配置则输出shape与输入shape一样
47│ │ │ ├── ${op_name}_tiling_${sub_case}.cpp # 可选，针对某些子场景下的Tiling优化，${sub_case}表示子场景，如${op_name}_tiling_arch35是针对arch35架构的优化，若无该文件表明该算子没有对应子场景的特定Tiling策略；注意：Tiling实现文件名须包含 `_tiling` 标识才会被编译系统识别并参与编译47│ │ │ ├── ${op_name}_tiling_${sub_case}.cpp # 可选，针对某些子场景下的Tiling优化，${sub_case}表示子场景，如${op_name}_tiling_arch35是针对arch35架构的优化，若无该文件表明该算子没有对应子场景的特定Tiling策略；注意：Tiling实现文件名须包含 `_tiling` 标识才会被编译系统识别并参与编译
48│ │ │ ├── ${op_name}_tiling_${sub_case}.h # 可选，${sub_case}子场景下Tiling实现用的头文件48│ │ │ ├── ${op_name}_tiling_${sub_case}.h # 可选，${sub_case}子场景下Tiling实现用的头文件
49-│ │ │ ├── ${op_name}_tiling.cpp # 可选，若无该文件表明对应场景下无Tiling实现(将张量划分为多个小块，区分数据类型进行并行计算)49+│ │ │ ├── ${op_name}_tiling.cpp # 可选，若无该文件表明对应场景下无Tiling实现（将张量划分为多个小块，区分数据类型进行并行计算）
50│ │ │ ├── ${op_name}_tiling.h # 可选，Tiling实现用的头文件50│ │ │ ├── ${op_name}_tiling.h # 可选，Tiling实现用的头文件
51│ │ │ └── CMakeLists.txt # Host侧cmakelist文件51│ │ │ └── CMakeLists.txt # Host侧cmakelist文件
52│ │ ├── op_api # 可选，算子aclnn实现文件目录，若未配置工程自动生成52│ │ ├── op_api # 可选，算子aclnn实现文件目录，若未配置工程自动生成
Mdocs/zh/install/quick_install.md+5-5
@@ -1,6 +1,6 @@
1# 环境部署1# 环境部署
2 2 
3-在学习各类[学习教程](../../../README.md#学习教程)操作之前，请您先参考下面步骤完成基础环境搭建和源码下载，确保已安装NPU驱动和CANN软件（`Ascend-cann-toolkit`和`Ascend-cann-ops`）等。3+在学习各类[学习教程](../../../README.md#学习教程)操作之前，请先参考下面步骤完成基础环境搭建和源码下载，确保已安装NPU驱动、固件和CANN软件（`Ascend-cann-toolkit`和`Ascend-cann-ops`）等。
4 4 
5## 环境安装5## 环境安装
6 6 
@@ -33,12 +33,12 @@
33 33 
34### 方式2：Docker部署34### 方式2：Docker部署
35 35 
36-对于有昇腾设备的开发者，若您想快速搭建昇腾环境，可使用Docker镜像部署。36+对于有昇腾设备的开发者，若想快速搭建昇腾环境，可使用Docker镜像部署。
37 37 
38> **说明**：38> **说明**：
39>39>
40-> - 镜像文件比较大，下载需要一定时间，请您耐心等待。关于docker命令的选项介绍可通过`docker --help`查询。40+> - 镜像文件比较大，下载需要一定时间，请耐心等待。关于docker命令的选项介绍可通过`docker --help`查询。
41-> - 环境默认安装最新版本CANN包，源码下载时注意与软件配套。41+> - 环境默认安装最新最新版本CANN包，源码下载时注意与软件配套。
42 42 
431.**安装驱动（运行态依赖）**431.**安装驱动（运行态依赖）**
44 44 
@@ -90,7 +90,7 @@ docker run --name cann_container --device /dev/davinci0 --device /dev/davinci_ma
90 90 
91### 方式3：手动安装91### 方式3：手动安装
92 92 
93-对于有昇腾设备的开发者，若您想手动搭建昇腾环境，请参考下述步骤。93+对于有昇腾设备的开发者，若想手动搭建昇腾环境，请参考下述步骤。
94 94 
95#### 安装软件95#### 安装软件
96 96 
Mdocs/zh/invocation/quick_op_invocation.md+5-5
@@ -36,7 +36,7 @@
36- 基于**自定义算子包**执行算子样例，命令如下：36- 基于**自定义算子包**执行算子样例，命令如下：
37 37 
38 ```bash38 ```bash
39- bash build.sh --run_example ${op} ${mode} ${pkg_mode} [--vendor_name=${vendor_name}] [--soc=${soc_version}]39+ bash build.sh --run_example ${op} ${mode} ${pkg_mode} [--example_name=${example_name}] [--vendor_name=${vendor_name}] [--soc=${soc_version}] [--simulator=${simulator}] [--experimental=${experimental}]
40 # 以FlashAttentionScore算子example执行为例40 # 以FlashAttentionScore算子example执行为例
41 # bash build.sh --run_example flash_attention_score eager cust --vendor_name=custom41 # bash build.sh --run_example flash_attention_score eager cust --vendor_name=custom
42 ```42 ```
@@ -125,7 +125,7 @@ mean result[1] is: 256.000000
125mean result[2] is: 256.000000125mean result[2] is: 256.000000
126mean result[3] is: 256.000000126mean result[3] is: 256.000000
127mean result[4] is: 256.000000127mean result[4] is: 256.000000
128-mean result[4] is: 256.000000128+mean result[5] is: 256.000000
129...129...
130mean result[65532] is: 256.000000130mean result[65532] is: 256.000000
131mean result[65533] is: 256.000000131mean result[65533] is: 256.000000
@@ -151,7 +151,7 @@ mean result[65535] is: 256.000000
151 151 
152#### 调用流程152#### 调用流程
153 153 
154-为方便调用算子，Host侧提供算子对应的C语言API（即以aclnn为前缀的API）实现算子调用，无需提供算子IR（Intermediate Representation）定义。aclnn API调用流程如下154+为方便调用算子，Host侧提供算子对应的C语言API（即以aclnn为前缀的API）实现算子调用，无需提供算子IR（Intermediate Representation）定义。aclnn API调用流程如下：
155 155 
156![原理图](../figures/aclnn调用.png)156![原理图](../figures/aclnn调用.png)
157 157 
@@ -257,7 +257,7 @@ int main()
257 257 
258- **调用自定义算子**：依赖自定义算子包258- **调用自定义算子**：依赖自定义算子包
259 259 
260- ```bash260+ ```cmake
261 cmake_minimum_required(VERSION 3.14)261 cmake_minimum_required(VERSION 3.14)
262 # 设置工程名262 # 设置工程名
263 project(ACLNN_EXAMPLE)263 project(ACLNN_EXAMPLE)
@@ -322,7 +322,7 @@ int main()
322 322 
323- **调用标准算子（内置算子）**：依赖ops-transformer整包323- **调用标准算子（内置算子）**：依赖ops-transformer整包
324 324 
325- ```bash325+ ```cmake
326 cmake_minimum_required(VERSION 3.14)326 cmake_minimum_required(VERSION 3.14)
327 # 设置工程名327 # 设置工程名
328 project(ACLNN_EXAMPLE)328 project(ACLNN_EXAMPLE)
Mdocs/zh/op_api_list.md+8-8
@@ -76,12 +76,12 @@
76|[aclnnFlashAttentionScoreGrad](../../attention/flash_attention_score_grad/docs/aclnnFlashAttentionScoreGrad.md)|训练场景下计算注意力的反向输出。|默认非确定性实现，支持配置开启| 默认非确定性实现，支持配置开启 |76|[aclnnFlashAttentionScoreGrad](../../attention/flash_attention_score_grad/docs/aclnnFlashAttentionScoreGrad.md)|训练场景下计算注意力的反向输出。|默认非确定性实现，支持配置开启| 默认非确定性实现，支持配置开启 |
77|[aclnnFlashAttentionScoreGradV2](../../attention/flash_attention_score_grad/docs/aclnnFlashAttentionScoreGradV2.md)|训练场景下计算注意力的反向输出，即[aclnnFlashAttentionScoreV2](../../attention/flash_attention_score/docs/aclnnFlashAttentionScoreV2.md)的反向计算。|默认非确定性实现，支持配置开启| 默认非确定性实现，支持配置开启 |77|[aclnnFlashAttentionScoreGradV2](../../attention/flash_attention_score_grad/docs/aclnnFlashAttentionScoreGradV2.md)|训练场景下计算注意力的反向输出，即[aclnnFlashAttentionScoreV2](../../attention/flash_attention_score/docs/aclnnFlashAttentionScoreV2.md)的反向计算。|默认非确定性实现，支持配置开启| 默认非确定性实现，支持配置开启 |
78|[aclnnFlashAttentionScoreGradV3](../../attention/flash_attention_score_grad/docs/aclnnFlashAttentionScoreGradV3.md)|训练场景下计算注意力的反向输出，即[aclnnFlashAttentionScoreV3](../../attention/flash_attention_score/docs/aclnnFlashAttentionScoreV3.md)的反向计算。|默认非确定性实现，支持配置开启| 默认非确定性实现，支持配置开启 |78|[aclnnFlashAttentionScoreGradV3](../../attention/flash_attention_score_grad/docs/aclnnFlashAttentionScoreGradV3.md)|训练场景下计算注意力的反向输出，即[aclnnFlashAttentionScoreV3](../../attention/flash_attention_score/docs/aclnnFlashAttentionScoreV3.md)的反向计算。|默认非确定性实现，支持配置开启| 默认非确定性实现，支持配置开启 |
79-|[aclnnFlashAttentionScoreGradV4](../../attention/flash_attention_score_grad/docs/aclnnFlashAttentionScoreGradV4.md)|训练场景下计算注意力的反向输出，即[FlashAttentionScoreV4](../../attention/flash_attention_score/docs/aclnnFlashAttentionScoreV4.md)的反向计算。该接口query、key、value参数支持多个长度相等或者长度不相等的sequence。|默认非确定性实现，支持配置开启| 默认非确定性实现，支持配置开启 |79+|[aclnnFlashAttentionScoreGradV4](../../attention/flash_attention_score_grad/docs/aclnnFlashAttentionScoreGradV4.md)|训练场景下计算注意力的反向输出，即[aclnnFlashAttentionScoreV4](../../attention/flash_attention_score/docs/aclnnFlashAttentionScoreV4.md)的反向计算。该接口query、key、value参数支持多个长度相等或者长度不相等的sequence。|默认非确定性实现，支持配置开启| 默认非确定性实现，支持配置开启 |
80|[aclnnFlashAttentionUnpaddingScoreGrad](../../attention/flash_attention_score_grad/docs/aclnnFlashAttentionUnpaddingScoreGrad.md)|训练场景下计算注意力的反向输出，即[aclnnFlashAttentionVarLenScore](../../attention/flash_attention_score/docs/aclnnFlashAttentionVarLenScore.md)的反向计算。|默认非确定性实现，支持配置开启| 默认非确定性实现，支持配置开启 |80|[aclnnFlashAttentionUnpaddingScoreGrad](../../attention/flash_attention_score_grad/docs/aclnnFlashAttentionUnpaddingScoreGrad.md)|训练场景下计算注意力的反向输出，即[aclnnFlashAttentionVarLenScore](../../attention/flash_attention_score/docs/aclnnFlashAttentionVarLenScore.md)的反向计算。|默认非确定性实现，支持配置开启| 默认非确定性实现，支持配置开启 |
81|[aclnnFlashAttentionUnpaddingScoreGradV2](../../attention/flash_attention_score_grad/docs/aclnnFlashAttentionUnpaddingScoreGradV2.md)|训练场景下计算注意力的反向输出，即[aclnnFlashAttentionVarLenScoreV2](../../attention/flash_attention_score/docs/aclnnFlashAttentionVarLenScoreV2.md)的反向计算。|默认非确定性实现，支持配置开启| 默认非确定性实现，支持配置开启 |81|[aclnnFlashAttentionUnpaddingScoreGradV2](../../attention/flash_attention_score_grad/docs/aclnnFlashAttentionUnpaddingScoreGradV2.md)|训练场景下计算注意力的反向输出，即[aclnnFlashAttentionVarLenScoreV2](../../attention/flash_attention_score/docs/aclnnFlashAttentionVarLenScoreV2.md)的反向计算。|默认非确定性实现，支持配置开启| 默认非确定性实现，支持配置开启 |
82|[aclnnFlashAttentionUnpaddingScoreGradV3](../../attention/flash_attention_score_grad/docs/aclnnFlashAttentionUnpaddingScoreGradV3.md)|训练场景下计算注意力的反向输出，即[aclnnFlashAttentionVarLenScoreV3](../../attention/flash_attention_score/docs/aclnnFlashAttentionVarLenScoreV3.md)的反向计算。该接口相较于[aclnnFlashAttentionUnpaddingScoreGradV2](../../attention/flash_attention_score_grad/docs/aclnnFlashAttentionUnpaddingScoreGradV2.md)接口，新增queryRope、keyRope、dqRope和dkRope参数。|默认非确定性实现，支持配置开启| 默认非确定性实现，支持配置开启 |82|[aclnnFlashAttentionUnpaddingScoreGradV3](../../attention/flash_attention_score_grad/docs/aclnnFlashAttentionUnpaddingScoreGradV3.md)|训练场景下计算注意力的反向输出，即[aclnnFlashAttentionVarLenScoreV3](../../attention/flash_attention_score/docs/aclnnFlashAttentionVarLenScoreV3.md)的反向计算。该接口相较于[aclnnFlashAttentionUnpaddingScoreGradV2](../../attention/flash_attention_score_grad/docs/aclnnFlashAttentionUnpaddingScoreGradV2.md)接口，新增queryRope、keyRope、dqRope和dkRope参数。|默认非确定性实现，支持配置开启| 默认非确定性实现，支持配置开启 |
83|[aclnnFlashAttentionUnpaddingScoreGradV4](../../attention/flash_attention_score_grad/docs/aclnnFlashAttentionUnpaddingScoreGradV4.md)|训练场景下计算注意力的反向输出。|默认非确定性实现，支持配置开启| 默认非确定性实现，支持配置开启 |83|[aclnnFlashAttentionUnpaddingScoreGradV4](../../attention/flash_attention_score_grad/docs/aclnnFlashAttentionUnpaddingScoreGradV4.md)|训练场景下计算注意力的反向输出。|默认非确定性实现，支持配置开启| 默认非确定性实现，支持配置开启 |
84-|[aclnnFlashAttentionUnpaddingScoreGradV5](../../attention/flash_attention_score_grad/docs/aclnnFlashAttentionUnpaddingScoreGradV5.md)|训练场景下，使用FlashAttention算法实现self-attention（自注意力）的计算。增加`sinkInOptional`可选输入。|默认非确定性实现，支持配置开启| 默认非确定性实现，支持配置开启 |84+|[aclnnFlashAttentionUnpaddingScoreGradV5](../../attention/flash_attention_score_grad/docs/aclnnFlashAttentionUnpaddingScoreGradV5.md)|训练场景下，使用FlashAttention算法实现self-attention（自注意力）的计算。增加`sinkInOptional`可选输入。本接口为反向输出。|默认非确定性实现，支持配置开启| 默认非确定性实现，支持配置开启 |
85|[aclnnFlashAttentionVarLenScore](../../attention/flash_attention_score/docs/aclnnFlashAttentionVarLenScore.md)|训练场景下，使用FlashAttention算法实现self-attention（自注意力）的计算。|默认确定性实现| 默认确定性实现 |85|[aclnnFlashAttentionVarLenScore](../../attention/flash_attention_score/docs/aclnnFlashAttentionVarLenScore.md)|训练场景下，使用FlashAttention算法实现self-attention（自注意力）的计算。|默认确定性实现| 默认确定性实现 |
86|[aclnnFlashAttentionVarLenScoreV2](../../attention/flash_attention_score/docs/aclnnFlashAttentionVarLenScoreV2.md)|训练场景下，使用FlashAttention算法实现self-attention（自注意力）的计算。|默认确定性实现| 默认确定性实现 |86|[aclnnFlashAttentionVarLenScoreV2](../../attention/flash_attention_score/docs/aclnnFlashAttentionVarLenScoreV2.md)|训练场景下，使用FlashAttention算法实现self-attention（自注意力）的计算。|默认确定性实现| 默认确定性实现 |
87|[aclnnFlashAttentionVarLenScoreV3](../../attention/flash_attention_score/docs/aclnnFlashAttentionVarLenScoreV3.md)|训练场景下，使用FlashAttention算法实现self-attention（自注意力）的计算。|默认确定性实现| 默认确定性实现 |87|[aclnnFlashAttentionVarLenScoreV3](../../attention/flash_attention_score/docs/aclnnFlashAttentionVarLenScoreV3.md)|训练场景下，使用FlashAttention算法实现self-attention（自注意力）的计算。|默认确定性实现| 默认确定性实现 |
@@ -102,9 +102,9 @@
102|[aclnnGroupedMatmulFinalizeRoutingV2](../../gmm/grouped_matmul_finalize_routing/docs/aclnnGroupedMatmulFinalizeRoutingV2.md)|GroupedMatmul和MoeFinalizeRouting的融合算子，GroupedMatmul计算后的输出按照索引做combine动作。|默认非确定性实现，支持配置开启| - |102|[aclnnGroupedMatmulFinalizeRoutingV2](../../gmm/grouped_matmul_finalize_routing/docs/aclnnGroupedMatmulFinalizeRoutingV2.md)|GroupedMatmul和MoeFinalizeRouting的融合算子，GroupedMatmul计算后的输出按照索引做combine动作。|默认非确定性实现，支持配置开启| - |
103|[aclnnGroupedMatmulFinalizeRoutingV3](../../gmm/grouped_matmul_finalize_routing/docs/aclnnGroupedMatmulFinalizeRoutingV3.md)|GroupedMatmul和MoeFinalizeRouting的融合算子，GroupedMatmul计算后的输出按照索引做combine动作。|默认非确定性实现，支持配置开启| 默认确定性实现 |103|[aclnnGroupedMatmulFinalizeRoutingV3](../../gmm/grouped_matmul_finalize_routing/docs/aclnnGroupedMatmulFinalizeRoutingV3.md)|GroupedMatmul和MoeFinalizeRouting的融合算子，GroupedMatmul计算后的输出按照索引做combine动作。|默认非确定性实现，支持配置开启| 默认确定性实现 |
104|[aclnnGroupedMatmulFinalizeRoutingWeightNz](../../gmm/grouped_matmul_finalize_routing/docs/aclnnGroupedMatmulFinalizeRoutingWeightNz.md)|GroupedMatMul和MoeFinalizeRouting的融合算子，GroupedMatmul计算后的输出按照索引做combine动作，支持输入Weight为AI处理器亲和数据排布格式(NZ)。|默认非确定性实现，支持配置开启| - |104|[aclnnGroupedMatmulFinalizeRoutingWeightNz](../../gmm/grouped_matmul_finalize_routing/docs/aclnnGroupedMatmulFinalizeRoutingWeightNz.md)|GroupedMatMul和MoeFinalizeRouting的融合算子，GroupedMatmul计算后的输出按照索引做combine动作，支持输入Weight为AI处理器亲和数据排布格式(NZ)。|默认非确定性实现，支持配置开启| - |
105-|[aclnnGroupedMatmulFinalizeRoutingWeightNzV2](../../gmm/grouped_matmul_finalize_routing/docs/aclnnGroupedMatmulFinalizeRoutingWeightNzV2.md)|GroupedMatmul和MoeFinalizeRouting的融合算子，GroupedMatmul计算后的输出按照索引做combine动作，支持w为AI处理器亲和数据排布格式(NZ)。|默认非确定性实现，支持配置开启| 默认确定性实现 |105+|[aclnnGroupedMatmulFinalizeRoutingWeightNzV2](../../gmm/grouped_matmul_finalize_routing/docs/aclnnGroupedMatmulFinalizeRoutingWeightNzV2.md)|GroupedMatmul和MoeFinalizeRouting的融合算子，GroupedMatmul计算后的输出按照索引做combine动作，支持输入Weight为AI处理器亲和数据排布格式(NZ)。|默认非确定性实现，支持配置开启| 默认确定性实现 |
106|[aclnnGroupedMatmulSwigluQuant](../../gmm/grouped_matmul_swiglu_quant/docs/aclnnGroupedMatmulSwigluQuant.md)|融合GroupedMatMul、Dequant、Swiglu和Quant。|默认确定性实现| - |106|[aclnnGroupedMatmulSwigluQuant](../../gmm/grouped_matmul_swiglu_quant/docs/aclnnGroupedMatmulSwigluQuant.md)|融合GroupedMatMul、Dequant、Swiglu和Quant。|默认确定性实现| - |
107-|[aclnnGroupedMatmulSwigluQuantV2](../../gmm/grouped_matmul_swiglu_quant_v2/docs/aclnnGroupedMatmulSwigluQuantV2.md)|融合GroupedMatmul 、dequant、swiglu和quant。|默认确定性实现| 默认确定性实现 |107+|[aclnnGroupedMatmulSwigluQuantV2](../../gmm/grouped_matmul_swiglu_quant_v2/docs/aclnnGroupedMatmulSwigluQuantV2.md)|融合GroupedMatmul 、Dequant、Swiglu和Quant。|默认确定性实现| 默认确定性实现 |
108|[aclnnGroupedMatmulSwigluQuantWeightNZ](../../gmm/grouped_matmul_swiglu_quant/docs/aclnnGroupedMatmulSwigluQuantWeightNZ.md)|融合GroupedMatMul、Dequant、Swiglu和Quant，输入权重Weight会被强制视为NZ格式。|默认确定性实现| - |108|[aclnnGroupedMatmulSwigluQuantWeightNZ](../../gmm/grouped_matmul_swiglu_quant/docs/aclnnGroupedMatmulSwigluQuantWeightNZ.md)|融合GroupedMatMul、Dequant、Swiglu和Quant，输入权重Weight会被强制视为NZ格式。|默认确定性实现| - |
109|[aclnnGroupedMatmulSwigluQuantWeightNzV2](../../gmm/grouped_matmul_swiglu_quant_v2/docs/aclnnGroupedMatmulSwigluQuantWeightNzV2.md)|融合GroupedMatMul、Dequant、Swiglu和Quant，输入权重Weight会被强制视为NZ格式。|默认确定性实现| 默认确定性实现 |109|[aclnnGroupedMatmulSwigluQuantWeightNzV2](../../gmm/grouped_matmul_swiglu_quant_v2/docs/aclnnGroupedMatmulSwigluQuantWeightNzV2.md)|融合GroupedMatMul、Dequant、Swiglu和Quant，输入权重Weight会被强制视为NZ格式。|默认确定性实现| 默认确定性实现 |
110|[aclnnGroupedMatmulWeightNz](../../gmm/grouped_matmul/docs/aclnnGroupedMatmulWeightNz.md)|实现分组矩阵乘计算，每组矩阵乘的维度大小可以不同，输入权重Weight会被强制视为NZ格式。|默认确定性实现| 默认确定性实现 |110|[aclnnGroupedMatmulWeightNz](../../gmm/grouped_matmul/docs/aclnnGroupedMatmulWeightNz.md)|实现分组矩阵乘计算，每组矩阵乘的维度大小可以不同，输入权重Weight会被强制视为NZ格式。|默认确定性实现| 默认确定性实现 |
@@ -117,7 +117,7 @@
117|[aclnnInterleaveRope](../../posembedding/interleave_rope/docs/aclnnInterleaveRope.md)|针对单输入x进行旋转位置编码。|- | 默认确定性实现 |117|[aclnnInterleaveRope](../../posembedding/interleave_rope/docs/aclnnInterleaveRope.md)|针对单输入x进行旋转位置编码。|- | 默认确定性实现 |
118|[aclnnKvCompressEpilog](../../attention/kv_compress_epilog/docs/aclnnKvCompressEpilog.md)|在KV Cache的Epilog阶段，对缓存进行原地逐组动态量化压缩更新，将bfloat16激活值压缩为FP8格式并按slotMapping散写。|- | 默认确定性实现 |118|[aclnnKvCompressEpilog](../../attention/kv_compress_epilog/docs/aclnnKvCompressEpilog.md)|在KV Cache的Epilog阶段，对缓存进行原地逐组动态量化压缩更新，将bfloat16激活值压缩为FP8格式并按slotMapping散写。|- | 默认确定性实现 |
119|[aclnnLightningIndexer](../../attention/lightning_indexer/docs/aclnnLightningIndexer.md)|稀疏attention前处理的计算，目的是选出关键的稀疏token位置。|默认确定性实现| 默认确定性实现 |119|[aclnnLightningIndexer](../../attention/lightning_indexer/docs/aclnnLightningIndexer.md)|稀疏attention前处理的计算，目的是选出关键的稀疏token位置。|默认确定性实现| 默认确定性实现 |
120-|[aclnnLightningIndexerGrad](../../attention/lightning_indexer_grad/docs/aclnnLightningIndexerGrad.md)|训练场景下，实现LightningIndexer反向，其中输入有Query, Key, Weights, Dy, Indices，反向主要利用正向计算的Indices从Key中提取TopK序列从而降低Matmul计算量。|默认非确定性实现，不支持配置开启|120+|[aclnnLightningIndexerGrad](../../attention/lightning_indexer_grad/docs/aclnnLightningIndexerGrad.md)|训练场景下，实现LightningIndexer反向，其中输入有Query, Key, Weights, Dy, Indices，反向主要利用正向计算的Indices从Key中提取TopK序列从而降低Matmul计算量。|默认非确定性实现，不支持配置开启| -|
121|[aclnnLightningIndexerV2](../../attention/lightning_indexer_v2/docs/aclnnLightningIndexerV2.md)|稀疏attention前处理的计算，目的是选出关键的稀疏token位置。支持KV压缩场景。|默认确定性实现| 默认确定性实现 |121|[aclnnLightningIndexerV2](../../attention/lightning_indexer_v2/docs/aclnnLightningIndexerV2.md)|稀疏attention前处理的计算，目的是选出关键的稀疏token位置。支持KV压缩场景。|默认确定性实现| 默认确定性实现 |
122|[aclnnLightningIndexerV2Metadata](../../attention/lightning_indexer_v2_metadata/docs/aclnnLightningIndexerV2Metadata.md)| aclnnLightningIndexerV2接口的前置接口，用于计算aclnnLightningIndexerV2的负载均衡。|默认确定性实现|默认确定性实现|122|[aclnnLightningIndexerV2Metadata](../../attention/lightning_indexer_v2_metadata/docs/aclnnLightningIndexerV2Metadata.md)| aclnnLightningIndexerV2接口的前置接口，用于计算aclnnLightningIndexerV2的负载均衡。|默认确定性实现|默认确定性实现|
123|[aclnnMaskedCausalConv1d](../../attention/masked_causal_conv1d/docs/aclnnMaskedCausalConv1d.md)|对hidden层的token之间进行带mask的因果一维分组卷积操作。| - | 默认确定性实现 |123|[aclnnMaskedCausalConv1d](../../attention/masked_causal_conv1d/docs/aclnnMaskedCausalConv1d.md)|对hidden层的token之间进行带mask的因果一维分组卷积操作。| - | 默认确定性实现 |
@@ -133,7 +133,7 @@
133|[aclnnQuantSparseFlashMla](../../attention/quant_sparse_flash_mla/docs/aclnnQuantSparseFlashMla.md)|支持全量化场景下SWA、CSA、HCA三类Attention计算场景。|-| 默认确定性实现 |133|[aclnnQuantSparseFlashMla](../../attention/quant_sparse_flash_mla/docs/aclnnQuantSparseFlashMla.md)|支持全量化场景下SWA、CSA、HCA三类Attention计算场景。|-| 默认确定性实现 |
134|[aclnnQuantSparseFlashMlaMetadata](../../attention/quant_sparse_flash_mla_metadata/docs/aclnnQuantSparseFlashMlaMetadata.md)| aclnnQuantSparseFlashMla接口的前置接口，用于计算aclnnQuantSparseFlashMla的负载均衡。| - | 默认确定性实现 |134|[aclnnQuantSparseFlashMlaMetadata](../../attention/quant_sparse_flash_mla_metadata/docs/aclnnQuantSparseFlashMlaMetadata.md)| aclnnQuantSparseFlashMla接口的前置接口，用于计算aclnnQuantSparseFlashMla的负载均衡。| - | 默认确定性实现 |
135|[aclnnMlaPreprocess](../../attention/mla_preprocess/docs/aclnnMlaPreprocess.md)|Multi-Head Latent Attention前处理的计算。|默认确定性实现| - |135|[aclnnMlaPreprocess](../../attention/mla_preprocess/docs/aclnnMlaPreprocess.md)|Multi-Head Latent Attention前处理的计算。|默认确定性实现| - |
136-|[aclnnMlaPreprocessV2](../../attention/mla_preprocess_v2/docs/aclnnMlaPreprocessV2.md)|推理场景，Multi-Head Latent Attention前处理的计算。主要计算过程如下：|默认确定性实现| - |136+|[aclnnMlaPreprocessV2](../../attention/mla_preprocess_v2/docs/aclnnMlaPreprocessV2.md)|推理场景，Multi-Head Latent Attention前处理的计算。|默认确定性实现| - |
137|[aclnnMlaProlog](../../attention/mla_prolog/docs/aclnnMlaProlog.md)|Multi-Head Latent Attention前处理的计算。|默认确定性实现| - |137|[aclnnMlaProlog](../../attention/mla_prolog/docs/aclnnMlaProlog.md)|Multi-Head Latent Attention前处理的计算。|默认确定性实现| - |
138|[aclnnMlaPrologV2WeightNz](../../attention/mla_prolog_v2/docs/aclnnMlaPrologV2WeightNz.md)|Multi-Head Latent Attention前处理的计算。|默认确定性实现| - |138|[aclnnMlaPrologV2WeightNz](../../attention/mla_prolog_v2/docs/aclnnMlaPrologV2WeightNz.md)|Multi-Head Latent Attention前处理的计算。|默认确定性实现| - |
139|[aclnnMlaPrologV3WeightNz](../../attention/mla_prolog_v3/docs/aclnnMlaPrologV3WeightNz.md)|Multi-Head Latent Attention前处理的计算。|默认非确定性实现，支持配置开启| 默认确定性实现 |139|[aclnnMlaPrologV3WeightNz](../../attention/mla_prolog_v3/docs/aclnnMlaPrologV3WeightNz.md)|Multi-Head Latent Attention前处理的计算。|默认非确定性实现，支持配置开启| 默认确定性实现 |
@@ -151,7 +151,7 @@
151|[aclnnMoeDistributeDispatchV3](../../mc2/moe_distribute_dispatch_v2/docs/aclnnMoeDistributeDispatchV3.md)|对token数据进行量化，当存在TP域通信时，先进行EP域的AllToAllV通信，再进行TP域的AllGatherV通信；当不存在TP域通信时，进行EP域的AllToAllV通信。|默认确定性实现| 默认确定性实现 |151|[aclnnMoeDistributeDispatchV3](../../mc2/moe_distribute_dispatch_v2/docs/aclnnMoeDistributeDispatchV3.md)|对token数据进行量化，当存在TP域通信时，先进行EP域的AllToAllV通信，再进行TP域的AllGatherV通信；当不存在TP域通信时，进行EP域的AllToAllV通信。|默认确定性实现| 默认确定性实现 |
152|[aclnnMoeDistributeDispatchV4](../../mc2/moe_distribute_dispatch_v2/docs/aclnnMoeDistributeDispatchV4.md)|对token数据进行量化，当存在TP域通信时，先进行EP域的AllToAllV通信，再进行TP域的AllGatherV通信；当不存在TP域通信时，进行EP域的AllToAllV通信。|默认确定性实现| 默认确定性实现 |152|[aclnnMoeDistributeDispatchV4](../../mc2/moe_distribute_dispatch_v2/docs/aclnnMoeDistributeDispatchV4.md)|对token数据进行量化，当存在TP域通信时，先进行EP域的AllToAllV通信，再进行TP域的AllGatherV通信；当不存在TP域通信时，进行EP域的AllToAllV通信。|默认确定性实现| 默认确定性实现 |
153|[aclnnMoeDistributeDispatchSetup](../../mc2/moe_distribute_dispatch_setup/docs/aclnnMoeDistributeDispatchSetup.md)|对Token数据进行量化（可选），根据token选择的topK专家在EP（Expert Parallelism）域的AllToAllV通信，只进行数据发送和通信状态发送，通信指令发出后算子即刻退出，无需等待通信完成。数据的接收和后处理由aclnnMoeDistributeDispatchTeardown接口完成。| - | 默认确定性实现 |153|[aclnnMoeDistributeDispatchSetup](../../mc2/moe_distribute_dispatch_setup/docs/aclnnMoeDistributeDispatchSetup.md)|对Token数据进行量化（可选），根据token选择的topK专家在EP（Expert Parallelism）域的AllToAllV通信，只进行数据发送和通信状态发送，通信指令发出后算子即刻退出，无需等待通信完成。数据的接收和后处理由aclnnMoeDistributeDispatchTeardown接口完成。| - | 默认确定性实现 |
154-|[aclnnMoeDistributeDispatchTeardown](../../mc2/moe_distribute_dispatch_teardown/docs/aclnnMoeDistributeDispatchTeardown.md)| 接收MOE层EP（Expert Parallelism）域的AllToAllV通信发过来的数据，数据发送端由aclnnMoeDistributeDispatchSetup完成，本接口内完成通信状态确认和数据整理。| - | 默认确定性实现 |154+|[aclnnMoeDistributeDispatchTeardown](../../mc2/moe_distribute_dispatch_teardown/docs/aclnnMoeDistributeDispatchTeardown.md)| 接收MoE层EP（Expert Parallelism）域的AllToAllV通信发过来的数据，数据发送端由aclnnMoeDistributeDispatchSetup完成，本接口内完成通信状态确认和数据整理。| - | 默认确定性实现 |
155|[aclnnMoeFinalizeRouting](../../moe/moe_finalize_routing/docs/aclnnMoeFinalizeRouting.md)|MoE计算中，最后处理合并MoE FFN的输出结果。|默认确定性实现| - |155|[aclnnMoeFinalizeRouting](../../moe/moe_finalize_routing/docs/aclnnMoeFinalizeRouting.md)|MoE计算中，最后处理合并MoE FFN的输出结果。|默认确定性实现| - |
156|[aclnnMoeFinalizeRoutingV2](../../moe/moe_finalize_routing_v2/docs/aclnnMoeFinalizeRoutingV2.md)|MoE计算中，最后处理合并MoE FFN的输出结果，支持配置dropPadMode。|默认确定性实现| 默认确定性实现 |156|[aclnnMoeFinalizeRoutingV2](../../moe/moe_finalize_routing_v2/docs/aclnnMoeFinalizeRoutingV2.md)|MoE计算中，最后处理合并MoE FFN的输出结果，支持配置dropPadMode。|默认确定性实现| 默认确定性实现 |
157|[aclnnMoeFinalizeRoutingV2Grad](../../moe/moe_finalize_routing_v2_grad/docs/aclnnMoeFinalizeRoutingV2Grad.md)|aclnnMoeFinalizeRoutingV2的反向传播。|默认确定性实现| 默认确定性实现 |157|[aclnnMoeFinalizeRoutingV2Grad](../../moe/moe_finalize_routing_v2_grad/docs/aclnnMoeFinalizeRoutingV2Grad.md)|aclnnMoeFinalizeRoutingV2的反向传播。|默认确定性实现| 默认确定性实现 |
@@ -178,7 +178,7 @@
178|[aclnnMoeTokenUnpermuteGrad](../../moe/moe_token_unpermute_grad/docs/aclnnMoeTokenUnpermuteGrad.md)|[aclnnMoeTokenUnpermute](../../moe/moe_token_unpermute/docs/aclnnMoeTokenUnpermute.md)的反向传播。|默认确定性实现| 默认确定性实现 |178|[aclnnMoeTokenUnpermuteGrad](../../moe/moe_token_unpermute_grad/docs/aclnnMoeTokenUnpermuteGrad.md)|[aclnnMoeTokenUnpermute](../../moe/moe_token_unpermute/docs/aclnnMoeTokenUnpermute.md)的反向传播。|默认确定性实现| 默认确定性实现 |
179|[aclnnMoeTokenUnpermuteWithEp](../../moe/moe_token_unpermute_with_ep/docs/aclnnMoeTokenUnpermuteWithEp.md)|根据sortedIndices存储的下标位置，去获取permutedTokens中的输入数据与probs相乘，并进行合并累加。|默认确定性实现| - |179|[aclnnMoeTokenUnpermuteWithEp](../../moe/moe_token_unpermute_with_ep/docs/aclnnMoeTokenUnpermuteWithEp.md)|根据sortedIndices存储的下标位置，去获取permutedTokens中的输入数据与probs相乘，并进行合并累加。|默认确定性实现| - |
180|[aclnnMoeTokenUnpermuteWithEpGrad](../../moe/moe_token_unpermute_with_ep_grad/docs/aclnnMoeTokenUnpermuteWithEpGrad.md)|[aclnnMoeTokenUnpermuteWithEp](../../moe/moe_token_unpermute_with_ep/docs/aclnnMoeTokenUnpermuteWithEp.md)的反向传播。|默认确定性实现| - |180|[aclnnMoeTokenUnpermuteWithEpGrad](../../moe/moe_token_unpermute_with_ep_grad/docs/aclnnMoeTokenUnpermuteWithEpGrad.md)|[aclnnMoeTokenUnpermuteWithEp](../../moe/moe_token_unpermute_with_ep/docs/aclnnMoeTokenUnpermuteWithEp.md)的反向传播。|默认确定性实现| - |
181-|[aclnnMoeTokenUnpermuteWithRoutingMap](../../moe/moe_token_unpermute_with_routing_map/docs/aclnnMoeTokenUnpermuteWithRoutingMap.md)|对经过aclnnMoeTokenpermuteWithRoutingMap处理的permutedTokens，累加回原unpermutedTokens。|默认确定性实现| 默认确定性实现 |181+|[aclnnMoeTokenUnpermuteWithRoutingMap](../../moe/moe_token_unpermute_with_routing_map/docs/aclnnMoeTokenUnpermuteWithRoutingMap.md)|对经过aclnnMoeTokenUpermuteWithRoutingMap处理的permutedTokens，累加回原unpermutedTokens。|默认确定性实现| 默认确定性实现 |
182|[aclnnMoeTokenUnpermuteWithRoutingMapGrad](../../moe/moe_token_unpermute_with_routing_map_grad/docs/aclnnMoeTokenUnpermuteWithRoutingMapGrad.md)|[aclnnMoeTokenUnpermuteWithRoutingMap](../../moe/moe_token_unpermute_with_routing_map/docs/aclnnMoeTokenUnpermuteWithRoutingMap.md)的反向传播。|默认非确定性实现，支持配置开启| 默认非确定性实现，支持配置开启 |182|[aclnnMoeTokenUnpermuteWithRoutingMapGrad](../../moe/moe_token_unpermute_with_routing_map_grad/docs/aclnnMoeTokenUnpermuteWithRoutingMapGrad.md)|[aclnnMoeTokenUnpermuteWithRoutingMap](../../moe/moe_token_unpermute_with_routing_map/docs/aclnnMoeTokenUnpermuteWithRoutingMap.md)的反向传播。|默认非确定性实现，支持配置开启| 默认非确定性实现，支持配置开启 |
183|[aclnnMoeUpdateExpert](../../mc2/moe_update_expert/docs/aclnnMoeUpdateExpert.md)|本API支持负载均衡和专家剪枝功能。经过映射后的专家表和Mask可传入MoE层进行数据分发和处理。|默认确定性实现| 默认确定性实现 |183|[aclnnMoeUpdateExpert](../../mc2/moe_update_expert/docs/aclnnMoeUpdateExpert.md)|本API支持负载均衡和专家剪枝功能。经过映射后的专家表和Mask可传入MoE层进行数据分发和处理。|默认确定性实现| 默认确定性实现 |
184|[aclnnMhcPre](../../mhc/mhc_pre/docs/aclnnMhcPre.md)|基于一系列计算得到MHC架构中hidden层的$H^{res}$和$H^{post}$投影矩阵以及Attention或MLP层的输入矩阵$h^{in}$。|默认确定性实现| - |184|[aclnnMhcPre](../../mhc/mhc_pre/docs/aclnnMhcPre.md)|基于一系列计算得到MHC架构中hidden层的$H^{res}$和$H^{post}$投影矩阵以及Attention或MLP层的输入矩阵$h^{in}$。|默认确定性实现| - |
Mdocs/zh/op_list.md+6-6
@@ -444,7 +444,7 @@
444 <td>✓</td>444 <td>✓</td>
445 <td>✗</td>445 <td>✗</td>
446 <td>AI Core</td>446 <td>AI Core</td>
447- <td>SparselightningIndexerGradKlLoss算子是LightningIndexer的反向算子，再额外融合了Loss计算功能输出。</td>447+ <td>SparselightningIndexerGradKlLoss算子是LightningIndexer的反向算子，并额外融合了Loss计算功能输出。</td>
448 </tr>448 </tr>
449 <tr>449 <tr>
450 <td>attention</td>450 <td>attention</td>
@@ -714,7 +714,7 @@
714 <td>✗</td>714 <td>✗</td>
715 <td>✓</td>715 <td>✓</td>
716 <td>AI Core</td>716 <td>AI Core</td>
717- <td>全量推理场景的FlashAttention算子。</td>717+ <td>面向Swin Transformer的注意力输出后处理融合算子，完成输出线性投影（MatMul+Bias）、窗口还原、反向循环移位以及残差加法。</td>
718 </tr>718 </tr>
719 <tr>719 <tr>
720 <td>ffn</td>720 <td>ffn</td>
@@ -1274,7 +1274,7 @@
1274 <td>✓</td>1274 <td>✓</td>
1275 <td>✗</td>1275 <td>✗</td>
1276 <td>AI Core</td>1276 <td>AI Core</td>
1277- <td>aclnnMoeTokenPermuteWithRoutingMap的反向传播。</td>1277+ <td>MoE的permute计算，将token和expert的标签作为routingMap传入，根据routingMap将tokens和可选probsOptional广播后排序。</td>
1278 </tr>1278 </tr>
1279 <tr>1279 <tr>
1280 <td>moe</td>1280 <td>moe</td>
@@ -1354,7 +1354,7 @@
1354 <td>✓</td>1354 <td>✓</td>
1355 <td>✓</td>1355 <td>✓</td>
1356 <td>AI Core</td>1356 <td>AI Core</td>
1357- <td>基于用Sinkhorn-Knopp迭代算法将超连接的混合矩阵投影到双随机矩阵流形，以此稳定深度网络信号传播、解决梯度消失 / 爆炸问题</td>1357+ <td>基于Sinkhorn-Knopp迭代算法将超连接的混合矩阵投影到双随机矩阵流形，以此稳定深度网络信号传播、解决梯度消失 / 爆炸问题</td>
1358 </tr>1358 </tr>
1359 <tr>1359 <tr>
1360 <td>mhc</td>1360 <td>mhc</td>
@@ -1384,7 +1384,7 @@
1384 <td>✓</td>1384 <td>✓</td>
1385 <td>✓</td>1385 <td>✓</td>
1386 <td>AI Core</td>1386 <td>AI Core</td>
1387- <td>mhc_post基于一系列计算对mHC架构中上一层输出进行Post Mapping，对上一层的输入进行Res Mapping，然后对二者进行残差连接，得到下一层的输入。该算子实现前述过程的反向</td>1387+ <td>mhc_post基于一系列计算对mHC架构中上一层输出进行Post Mapping，对上一层的输入进行Res Mapping，然后对二者进行残差连接，得到下一层的输入。该算子实现前述过程的反向计算。</td>
1388 </tr>1388 </tr>
1389 <tr>1389 <tr>
1390 <td>mhc</td>1390 <td>mhc</td>
@@ -1514,7 +1514,7 @@
1514 <td>✗</td>1514 <td>✗</td>
1515 <td>✓</td>1515 <td>✓</td>
1516 <td>AI Core</td>1516 <td>AI Core</td>
1517- <td>对输入张量的尾轴进行切分。</td>1517+ <td>对输入张量按照‘sizeSplits’对尾轴进行切分，划分为q、k、vOut，对q、k进行旋转位置编码，生成qOut与kOut，之后对kOut与vOut进行量化并按照‘indices’更新到kCacheRef和vCacheRef上。</td>
1518 </tr>1518 </tr>
1519 <tr>1519 <tr>
1520 <td>posembedding</td>1520 <td>posembedding</td>
Mdocs/zh/torch_api_list.md+1-1
@@ -15,7 +15,7 @@
15 ```python15 ```python
16 import torch16 import torch
17 import torch_npu17 import torch_npu
18- import cann-ops-transformer18+ import cann_ops_transformer
19 ```19 ```
20 20 
21- **V版本演进说明**21- **V版本演进说明**