已合并
fix: Modify the AIDD docs issue #4332
zwj223创建于 24 天前
fix: Modify the AIDD docs issue #4332
已合并
zwj223创建于 24 天前
18 个文件变更+93-91
MCHANGELOG.md+1-1
@@ -13,7 +13,7 @@ ops-math首个Beta版本v8.5.0-beta.1现已发布。
13 13 
14### 🔗 版本地址14### 🔗 版本地址
15 15 
16-[CANN 8.5.0-beta 1](https://ascend.devcloud.huaweicloud.com/cann/run/software/8.5.0-beta.1/)16+[CANN 8.5.0-beta.1](https://ascend.devcloud.huaweicloud.com/cann/run/software/8.5.0-beta.1/)
17 17 
18```text18```text
19版本目录说明如下:19版本目录说明如下:
Mdocs/CONTRIBUTING_DOCS.md+1-1
@@ -84,7 +84,7 @@
84 - 图片:84 - 图片:
85 - 常用格式:推荐png格式,风格尽量与已有图片保持一致。85 - 常用格式:推荐png格式,风格尽量与已有图片保持一致。
86 - 分辨率与清晰度:需清晰且尺寸适中,避免模糊或过度压缩。86 - 分辨率与清晰度:需清晰且尺寸适中,避免模糊或过度压缩。
87- - 文件大小:单张图片不建议超过10M87+ - 文件大小:单张图片不建议超过10MB
88 - 版权:所有引用的图片、文献等资源,请确保合规性。88 - 版权:所有引用的图片、文献等资源,请确保合规性。
89 89 
90## 获取帮助90## 获取帮助
Mdocs/zh/context/broadcast_relationship.md+2-2
@@ -14,12 +14,12 @@ broadcast(广播)描述了算子在运算期间如何处理不同形状的
14 14 
15 > 说明:15 > 说明:
16 > - 举例1:维度数(Number of Dimensions)是指张量(或数组)对应shape的维数,比如x.shape=(1,1,2,4),维度数是4 。16 > - 举例1:维度数(Number of Dimensions)是指张量(或数组)对应shape的维数,比如x.shape=(1,1,2,4),维度数是4 。
17- > - 举例2:比如计算a+b,其中a.shape=\(2, 2, 3\)、b.shape=\(2, 3\),那么数组b将被broadcast为b.shape=\(1, 2, 3\)。17+ > - 举例2:比如计算a+b,其中a.shape=(2, 2, 3)、b.shape=(2, 3),那么数组b将被broadcast为b.shape=(1, 2, 3)。
18 18 
19- 规则2:如果数组间维度数一致,且某个数组的某一维度为1,则该维度为1的数组将被拉伸以匹配另一个数组对应维度形状。19- 规则2:如果数组间维度数一致,且某个数组的某一维度为1,则该维度为1的数组将被拉伸以匹配另一个数组对应维度形状。
20 20 
21 > 说明:21 > 说明:
22- > 本场景下,只需保证在某一维度做broadcast即可。比如计算a+b,其中a.shape=\(1, 3\)、b.shape=\(3, 1\),那么两个数组会broadcast为a.shape=\(3, 3\)、b.shape=\(3, 3\)。22+ > 本场景下,只需保证在某一维度做broadcast即可。比如计算a+b,其中a.shape=(1, 3)、b.shape=(3, 1),那么两个数组会broadcast为a.shape=(3, 3)、b.shape=(3, 3)。
23 23 
24- 规则3:如果数组间维度数不一致,且均没有等于1的维度,则会报错。24- 规则3:如果数组间维度数不一致,且均没有等于1的维度,则会报错。
25 25 
Mdocs/zh/context/compile_and_run_sample.md+4-4
@@ -7,7 +7,7 @@
7 7 
8## 编译前准备8## 编译前准备
9 9 
10-本章以开发和运行环境合设场景为例,即带AI处理器的机器既作为开发环境又作为运行环境。该场景下,代码开发和代码运行在同一台机器上。这里以**Abs算子**为例,其他算子的调用逻辑、流程、编译脚本与Abs算子大致一样,请根据实际情况自行修改API调用脚本(\*.cpp)和编译脚本(CMakeLists)10+本章以开发和运行环境合设场景为例,即带AI处理器的机器既作为开发环境又作为运行环境。该场景下,代码开发和代码运行在同一台机器上。这里以**Abs算子**为例,其他算子的调用逻辑、流程、编译脚本与Abs算子大致一样,请根据实际情况自行修改API调用脚本(\*.cpp)和编译脚本CMakeLists
11 11 
12- **示例代码**12- **示例代码**
13 13 
@@ -17,7 +17,7 @@
17 17 
18 CMake文件示例如下,请根据实际情况修改:18 CMake文件示例如下,请根据实际情况修改:
19 19 
20- ```bash20+ ```cmake
21 # Copyright (c) Huawei Technologies Co., Ltd. 2019. All rights reserved.21 # Copyright (c) Huawei Technologies Co., Ltd. 2019. All rights reserved.
22 22 
23 # CMake lowest version requirement23 # CMake lowest version requirement
@@ -62,7 +62,7 @@
62 62 
63 对于集合通信和MatMul计算融合、并行的算子,统称为通算融合算子(简称MC2算子),包括AllGatherMatmul、AlltoAllAllGatherBatchMatMul、BatchMatMulReduceScatterAlltoAll、MatmulAllReduce、MatmulAllReduceAddRmsNorm、MatmulReduceScatter等。调用该类算子API时,一般会涉及多线程和HCCL(Huawei Collective Communication Library,集合通信库),因此CMake文件需要额外导入如下内容,否则无法成功编译。63 对于集合通信和MatMul计算融合、并行的算子,统称为通算融合算子(简称MC2算子),包括AllGatherMatmul、AlltoAllAllGatherBatchMatMul、BatchMatMulReduceScatterAlltoAll、MatmulAllReduce、MatmulAllReduceAddRmsNorm、MatmulReduceScatter等。调用该类算子API时,一般会涉及多线程和HCCL(Huawei Collective Communication Library,集合通信库),因此CMake文件需要额外导入如下内容,否则无法成功编译。
64 64 
65- ```text65+ ```cmake
66 # 设置链接的库文件路径66 # 设置链接的库文件路径
67 find_package(Threads REQUIRED)67 find_package(Threads REQUIRED)
68 target_link_libraries(opapi_test PRIVATE68 target_link_libraries(opapi_test PRIVATE
@@ -135,7 +135,7 @@
135 CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnAbsGetWorkspaceSize failed. ERROR: %d.\n[ERROR msg]%s", ret, aclGetRecentErrMsg()); return ret);135 CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnAbsGetWorkspaceSize failed. ERROR: %d.\n[ERROR msg]%s", ret, aclGetRecentErrMsg()); return ret);
136 ```136 ```
137 137 
138- 上述构造空指针问题获取报错信息示例如下:138+ 上述构造空指针问题获取报错信息示例如下
139 139 
140 ```bash140 ```bash
141 aclnnAbsGetWorkspaceSize failed. ERROR: 161001141 aclnnAbsGetWorkspaceSize failed. ERROR: 161001
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+2-2
@@ -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 
Mdocs/zh/debug/op_debug_prof.md+14-14
@@ -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
@@ -28,7 +28,7 @@
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 ```bash30 ```bash
31- printf(aclGetRecentErrMsg());31+ printf("%s",aclGetRecentErrMsg());
32 ```32 ```
33 33 
34 打印错误信息样例如下:34 打印错误信息样例如下:
@@ -82,11 +82,11 @@
82 82 
83 可通过如下宏打印算子执行过程中的日志信息,包括DEBUG、INFO、WARN、ERROR级别日志。83 可通过如下宏打印算子执行过程中的日志信息,包括DEBUG、INFO、WARN、ERROR级别日志。
84 84 
85- ```Cpp85+ ```c++
86- KERNEL_LOG_DEBUG(fmt, ) // fmt参数表示格式控制字符串86+ KERNEL_LOG_DEBUG(fmt, ...) // fmt参数表示格式控制字符串
87- KERNEL_LOG_INFO(fmt, )87+ KERNEL_LOG_INFO(fmt, ...)
88- KERNEL_LOG_WARN(fmt, )88+ KERNEL_LOG_WARN(fmt, ...)
89- KERNEL_LOG_ERROR(fmt, ) // 默认打印ERROR级别日志89+ KERNEL_LOG_ERROR(fmt, ...) // 默认打印ERROR级别日志
90 ```90 ```
91 91 
92 如需打印非ERROR级别日志,需提前配置环境变量`ASCEND_GLOBAL_LOG_LEVEL`,具体使用方法参见[《环境变量参考》](https://hiascend.com/document/redirect/CannCommunityEnvRef)。92 如需打印非ERROR级别日志,需提前配置环境变量`ASCEND_GLOBAL_LOG_LEVEL`,具体使用方法参见[《环境变量参考》](https://hiascend.com/document/redirect/CannCommunityEnvRef)。
@@ -106,7 +106,7 @@
106 106 
107 int64_t num_elements = input0->NumElements();107 int64_t num_elements = input0->NumElements();
108 // 打印输入元素个数108 // 打印输入元素个数
109- KERNEL_LOG_INFO("Num of elements is %ld", data_size);109+ KERNEL_LOG_INFO("Num of elements is %ld", num_elements);
110 ```110 ```
111 111 
112## 性能调优112## 性能调优
@@ -162,15 +162,15 @@
162 162 
163 执行仿真命令,生成仿真数据163 执行仿真命令,生成仿真数据
164 164 
165- ```165+ ```bash
166- cannsim record ./test_aclnn_add_example -s Ascend950 --gen-report166+ cannsim record ./test_aclnn_add_example -s Ascend950 --gen-report
167- ```167+ ```
168 168 
169 仿真结果在本项目`examples/add_example/examples/build/bin/cannsim_*/report/results/kernel_*/core_*`目录,流水相关文件为:169 仿真结果在本项目`examples/add_example/examples/build/bin/cannsim_*/report/results/kernel_*/core_*`目录,流水相关文件为:
170 170 
171- ```171+ ```text
172- trace_core0.json172+ trace_core0.json
173- ```173+ ```
174 174 
175 在Chrome浏览器中输入“chrome://tracing”地址,并将生成的指令流水图文件(trace_core0.json)拖到空白处打开,具体参数介绍参考CANN Simulator中[“仿真结果解析说明”](./cann_sim.md#仿真结果解析说明)章节。175 在Chrome浏览器中输入“chrome://tracing”地址,并将生成的指令流水图文件(trace_core0.json)拖到空白处打开,具体参数介绍参考CANN Simulator中[“仿真结果解析说明”](./cann_sim.md#仿真结果解析说明)章节。
176 176 
Mdocs/zh/develop/aicore_develop_guide.md+16-16
@@ -420,10 +420,10 @@ __aicore__ inline void AddExample<T>::Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, cons
420 blockLength_ = tilingData->totalLength / AscendC::GetBlockNum();420 blockLength_ = tilingData->totalLength / AscendC::GetBlockNum();
421 ...421 ...
422 // 3.2初始化GM地址422 // 3.2初始化GM地址
423- inputGMX.SetGlobalBuffer((__gm__ T*)x + blockLength_ * AscendC::GetBlockIdx(), blockLength_);423+ inputGMX_.SetGlobalBuffer((__gm__ T*)x + blockLength_ * AscendC::GetBlockIdx(), blockLength_);
424 ...424 ...
425 // 3.3初始化队列长度425 // 3.3初始化队列长度
426- pipe.InitBuffer(inputQueueX_, BUFFER_NUM, tileLength_ * sizeof(T));426+ pipe_.InitBuffer(inputQueueX_, BUFFER_NUM, tileLength_ * sizeof(T));
427 ...427 ...
428}428}
429 429 
@@ -456,7 +456,7 @@ __aicore__ inline void AddExample<T>::Process()
456`scripts/kernel/binary_config`目录[ascendc_config.json](../../../scripts/kernel/binary_config/ascendc_config.json)中,注册算子的NPU型号和实现模式,示例如下,输入实际name和compute_units即可。456`scripts/kernel/binary_config`目录[ascendc_config.json](../../../scripts/kernel/binary_config/ascendc_config.json)中,注册算子的NPU型号和实现模式,示例如下,输入实际name和compute_units即可。
457 457 
458```json458```json
459-{"name":"AddExample", "compute_units": ["${soc_version}"], "auto_sync":true, "impl_mode" : "high_performance"},459+{"name":"AddExample", "compute_units": ["${soc_version}"], "auto_sync":true, "impl_mode" : "high_performance"}
460```460```
461 461 
462## 编译部署462## 编译部署
@@ -567,7 +567,7 @@ UT编写指导如下,如需查看详细实现,请参考样例UT实现[test_a
567 567 
568测试类示例:568测试类示例:
569 569 
570-```CPP570+```cpp
571class ${OpName}InfershapeTest : public testing::Test {571class ${OpName}InfershapeTest : public testing::Test {
572protected:572protected:
573 static void SetUpTestCase()573 static void SetUpTestCase()
@@ -591,7 +591,7 @@ protected:
591 591 
592简化示例:592简化示例:
593 593 
594-```CPP594+```cpp
595TEST_F(${OpName}InfershapeTest, test_case_xxx)595TEST_F(${OpName}InfershapeTest, test_case_xxx)
596{596{
597 // 1.构造用例上下文597 // 1.构造用例上下文
@@ -632,7 +632,7 @@ UT编写指导如下,如需查看详细实现,请参考样例UT实现[test_a
632 632 
633测试类示例:633测试类示例:
634 634 
635-```CPP635+```cpp
636class ${OpName}TilingTest : public testing::Test {636class ${OpName}TilingTest : public testing::Test {
637protected:637protected:
638 static void SetUpTestCase()638 static void SetUpTestCase()
@@ -658,7 +658,7 @@ protected:
658 658 
659简化示例:659简化示例:
660 660 
661-```CPP661+```cpp
662TEST_F(${OpName}TilingTest, test_case_xxx)662TEST_F(${OpName}TilingTest, test_case_xxx)
663{663{
664 // 声明结构体并初始化一个结构体变量664 // 声明结构体并初始化一个结构体变量
@@ -712,7 +712,7 @@ UT编写指导如下,如需查看详细实现,请参考样例UT实现[test_a
712 712 
713测试类示例:713测试类示例:
714 714 
715-```CPP715+```cpp
716class ${OpName}KernelTest : public testing::Test {716class ${OpName}KernelTest : public testing::Test {
717protected:717protected:
718 static void SetUpTestCase()718 static void SetUpTestCase()
@@ -739,7 +739,7 @@ protected:
739 739 
740简化示例:740简化示例:
741 741 
742-```CPP742+```cpp
743extern "C" __global__ __aicore__ void ${op_name}(GM_ADDR x, GM_ADDR y, GM_ADDR z,743extern "C" __global__ __aicore__ void ${op_name}(GM_ADDR x, GM_ADDR y, GM_ADDR z,
744 GM_ADDR workspace, GM_ADDR tiling);744 GM_ADDR workspace, GM_ADDR tiling);
745 745 
@@ -779,7 +779,7 @@ TEST_F(${OpName}KernelTest, test_case_basic)
779- **手动构造**:适合字段少、逻辑简单。779- **手动构造**:适合字段少、逻辑简单。
780- **调用Tiling函数自动生成**:适合字段多、依赖属性/shape复杂。可复用`tests/ut/common/tiling_context_faker.h``tiling_case_executor.h`。示例:780- **调用Tiling函数自动生成**:适合字段多、依赖属性/shape复杂。可复用`tests/ut/common/tiling_context_faker.h``tiling_case_executor.h`。示例:
781 781 
782-```CPP782+```cpp
783gert::TilingContextPara para("OpName",783gert::TilingContextPara para("OpName",
784 {{{{2, 2, 2, 1}, {2, 2, 2, 1}}, ge::DT_FLOAT, ge::FORMAT_ND}},784 {{{{2, 2, 2, 1}, {2, 2, 2, 1}}, ge::DT_FLOAT, ge::FORMAT_ND}},
785 {{{{2, 1, 2, 2}, {2, 1, 2, 2}}, ge::DT_FLOAT, ge::FORMAT_ND}},785 {{{{2, 1, 2, 2}, {2, 1, 2, 2}}, ge::DT_FLOAT, ge::FORMAT_ND}},
@@ -884,7 +884,7 @@ export LD_LIBRARY_PATH=${ASCEND_HOME_PATH}/opp/vendors/${vendor_name}_math/op_ap
884 884 
885将原有${op_name}.cpp中算子信息库内容独立迁移至该文件,需要去掉SetInferShape和SetTiling内容。885将原有${op_name}.cpp中算子信息库内容独立迁移至该文件,需要去掉SetInferShape和SetTiling内容。
886 886 
887-```CPP887+```cpp
888// 原有${op_name}.cpp中算子信息库内容888// 原有${op_name}.cpp中算子信息库内容
889namespace ops {889namespace ops {
890class AddCustom : public OpDef {890class AddCustom : public OpDef {
@@ -940,7 +940,7 @@ OP_ADD(AddCustom);
940 940 
941图模式场景需要适配该文件,将原有${op_name}.cpp中shape推导部分独立迁至该文件,调用接口IMPL_OP_INFERSHAPE完成InferShape注册。941图模式场景需要适配该文件,将原有${op_name}.cpp中shape推导部分独立迁至该文件,调用接口IMPL_OP_INFERSHAPE完成InferShape注册。
942 942 
943-```CPP943+```cpp
944// 原有${op_name}.cpp中的InferShape944// 原有${op_name}.cpp中的InferShape
945namespace ge {945namespace ge {
946static graphStatus InferShape(gert::InferShapeContext *context)946static graphStatus InferShape(gert::InferShapeContext *context)
@@ -974,7 +974,7 @@ IMPL_OP_INFERSHAPE(AddCustom).InferShape(InferShape); // 在该文件中完成
974若是新增定义模板参数和模板参数组合,TilingFunc中需要同时配置模板参数tilingKey。974若是新增定义模板参数和模板参数组合,TilingFunc中需要同时配置模板参数tilingKey。
975可参考[add_example_tiling.cpp](../../../examples/add_example/op_host/add_example_tiling.cpp)。975可参考[add_example_tiling.cpp](../../../examples/add_example/op_host/add_example_tiling.cpp)。
976 976 
977-```CPP977+```cpp
978// 原有${op_name}.cpp中TilingFunc978// 原有${op_name}.cpp中TilingFunc
979namespace optiling {979namespace optiling {
980const uint32_t BLOCK_DIM = 8;980const uint32_t BLOCK_DIM = 8;
@@ -1032,7 +1032,7 @@ IMPL_OP_OPTILING(AddCustom).Tiling(TilingFunc); // 在该文件中完成Tiling
1032</div>1032</div>
1033图模式场景需要适配该文件,将原有${op_name}.cpp中类型推导独立迁移至该文件后,调用接口IMPL_OP完成InferDataType注册。1033图模式场景需要适配该文件,将原有${op_name}.cpp中类型推导独立迁移至该文件后,调用接口IMPL_OP完成InferDataType注册。
1034 1034 
1035-```CPP1035+```cpp
1036// 原有${op_name}.cpp中InferDataType1036// 原有${op_name}.cpp中InferDataType
1037namespace ge {1037namespace ge {
1038static graphStatus InferDataType(gert::InferDataTypeContext *context)1038static graphStatus InferDataType(gert::InferDataTypeContext *context)
@@ -1059,7 +1059,7 @@ IMPL_OP(AddCustom).InferDataType(InferDataType); // 在该文件中完成Infer
1059<p style="font-size:18px;"><b>op_kernel/{op_name}_tiling_data.h</b></p>1059<p style="font-size:18px;"><b>op_kernel/{op_name}_tiling_data.h</b></p>
1060</div>1060</div>
1061 1061 
1062-```CPP1062+```cpp
1063// 原有op_host/{op_name}_tiling.h中的宏定义TilingData结构体1063// 原有op_host/{op_name}_tiling.h中的宏定义TilingData结构体
1064namespace optiling {1064namespace optiling {
1065BEGIN_TILING_DATA_DEF(TilingData)1065BEGIN_TILING_DATA_DEF(TilingData)
@@ -1085,7 +1085,7 @@ struct TilingData {
1085<p style="font-size:18px;"><b>op_kernel/{op_name}.cpp</b></p>1085<p style="font-size:18px;"><b>op_kernel/{op_name}.cpp</b></p>
1086</div>1086</div>
1087 1087 
1088-```CPP1088+```cpp
1089// 原有op_kernel/{op_name}.cpp中的核函数实现1089// 原有op_kernel/{op_name}.cpp中的核函数实现
1090template<int D_T_X, int D_T_Y, int D_T_Z, int TILE_NUM, int IS_SPLIT>1090template<int D_T_X, int D_T_Y, int D_T_Z, int TILE_NUM, int IS_SPLIT>
1091 __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling)1091 __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling)
Mdocs/zh/develop/aicpu_develop_guide.md+3-3
@@ -115,7 +115,7 @@ Kernel一共需要两个交付件:```${op_name}_aicpu.cpp``` ```${op_name}_aic
115Kernel实现的第一步,需在头文件```op_kernel_aicpu/${op_name}_aicpu.h```进行算子类的声明,算子类需继承CpuKernel基类。115Kernel实现的第一步,需在头文件```op_kernel_aicpu/${op_name}_aicpu.h```进行算子类的声明,算子类需继承CpuKernel基类。
116如需查看详细实现,请参考[add_example_aicpu.h](../../../scripts/opgen/template/add_example_aicpu/op_kernel_aicpu/add_example_aicpu.h)。116如需查看详细实现,请参考[add_example_aicpu.h](../../../scripts/opgen/template/add_example_aicpu/op_kernel_aicpu/add_example_aicpu.h)。
117 117 
118-```CPP118+```cpp
119// 1、算子类声明119// 1、算子类声明
120// 包含AI CPU基础库头文件120// 包含AI CPU基础库头文件
121#include "cpu_kernel.h"121#include "cpu_kernel.h"
@@ -169,7 +169,7 @@ uint32_t AddExampleCpuKernel::Compute(CpuKernelContext& ctx) {
169 // 获取input tensor的数据地址,例如输入的数据类型是int32169 // 获取input tensor的数据地址,例如输入的数据类型是int32
170 auto input0_data = reinterpret_cast<int32_t*>(input0->GetData());170 auto input0_data = reinterpret_cast<int32_t*>(input0->GetData());
171 // 获取tensor的shape171 // 获取tensor的shape
172- auto input0_shape = input0->GetTensorShape();172+ auto tensor_shape = input0->GetTensorShape();
173 173 
174 // 获取output tensor的数据地址,例如输出的数据类型是int32174 // 获取output tensor的数据地址,例如输出的数据类型是int32
175 auto y = reinterpret_cast<int32_t*>(output->GetData());175 auto y = reinterpret_cast<int32_t*>(output->GetData());
@@ -180,7 +180,7 @@ 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 PARAM_INVALID;
186 }186 }
Mdocs/zh/develop/cross_platform_migration_guide.md+9-7
@@ -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>
@@ -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>不支持42稀疏矩阵计算</td>102+ <td>不支持4:2稀疏矩阵计算</td>
103- <td>原依赖42稀疏特性提速的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,7 +253,7 @@ __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* 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;
@@ -350,7 +352,7 @@ Ascend 950新架构引入UB2L1 & L0C2UB间的直连通路,实现矩阵计算
350```cpp352```cpp
351// 1.新增: 搬入接口增加UB2L1的Nd2Nz搬入,支持Src&Dst都是LocalTensor的形式353// 1.新增: 搬入接口增加UB2L1的Nd2Nz搬入,支持Src&Dst都是LocalTensor的形式
352template <typename T>  354template <typename T>  
353-__aicore__ inline void DataCopy(const LocalTensor<T>& dst, const LocalTensor<T>& src, const Nd2NzParams& intriParams)355+__aicore__ inline void DataCopy(const LocalTensor<T>& dst, const LocalTensor<T>& src, const Nd2NzParams& intriParams);
354 356 
355// 2.新增: 搬出接口增加L0C2UB的搬出,支持直接从L0C搬出到UB,支持Src&Dst都是LocalTensor的形式357// 2.新增: 搬出接口增加L0C2UB的搬出,支持直接从L0C搬出到UB,支持Src&Dst都是LocalTensor的形式
356template <typename T, typename U, const FixpipeConfig& config = CFG_ROW_MAJOR>358template <typename T, typename U, const FixpipeConfig& config = CFG_ROW_MAJOR>
@@ -393,7 +395,7 @@ Ascend 950引入集合通信加速器CCU1.0,降低了访存需求,减少了
393以[MatmulAllReduce](https://gitcode.com/cann/ops-transformer/tree/master/mc2/matmul_all_reduce)算子迁移适配为例:395以[MatmulAllReduce](https://gitcode.com/cann/ops-transformer/tree/master/mc2/matmul_all_reduce)算子迁移适配为例:
394设置NnopbaseSetHcclServerType枚举值,A2为NNOPBASE_HCCL_SERVER_AICPU,950为NNOPBASE_HCCL_SERVER_TYPE_CCU。396设置NnopbaseSetHcclServerType枚举值,A2为NNOPBASE_HCCL_SERVER_AICPU,950为NNOPBASE_HCCL_SERVER_TYPE_CCU。
395 397 
396-```CPP398+```cpp
397// ...399// ...
398aclnnStatus aclnnMatmulAllReduce(400aclnnStatus aclnnMatmulAllReduce(
399 void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, const aclrtStream stream)401 void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, const aclrtStream stream)
Mdocs/zh/develop/graph_develop_guide.md+3-3
@@ -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{
@@ -56,7 +56,7 @@ 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{
@@ -115,7 +115,7 @@ IMPL_OP(AddExample).InferDataType(InferDataTypeAddExample);
115 115 
116示例代码如下,展示了如何注册`AddExample`算子:116示例代码如下,展示了如何注册`AddExample`算子:
117 117 
118-```CPP118+```cpp
119REG_OP(AddExample)119REG_OP(AddExample)
120 .INPUT(x1, TensorType({DT_FLOAT}))120 .INPUT(x1, TensorType({DT_FLOAT}))
121 .INPUT(x2, TensorType({DT_FLOAT}))121 .INPUT(x2, TensorType({DT_FLOAT}))
Mdocs/zh/install/build.md+1-1
@@ -78,4 +78,4 @@ bash build.sh --help
78| --bisheng_flags | 可选 | 指定毕昇编译器编译参数,多个编译参数用英文逗号“,”分隔,不可与--mssanitizer、--oom、--dump_cce同时使用。 |78| --bisheng_flags | 可选 | 指定毕昇编译器编译参数,多个编译参数用英文逗号“,”分隔,不可与--mssanitizer、--oom、--dump_cce同时使用。 |
79| --kernel_template_input | 可选 | 指定编译kernel时的模板参数,多个模板参数用英文分号“;”分隔,与--ops同时使用且只能指定一个算子。 |79| --kernel_template_input | 可选 | 指定编译kernel时的模板参数,多个模板参数用英文分号“;”分隔,与--ops同时使用且只能指定一个算子。 |
80| --cann_3rd_lib_path | 可选 | 离线编译场景下第三方库存放的目录。 |80| --cann_3rd_lib_path | 可选 | 离线编译场景下第三方库存放的目录。 |
81-| --no_force | 可选 | 待编译的算子依赖其他算子时,不再编译其他算子二进制文件 |81+| --no_force | 可选 | 待编译的算子依赖其他算子时,不再编译其他算子二进制文件 |
Mdocs/zh/install/compile.md+2-2
@@ -13,7 +13,7 @@
13| json | 3.11.3 | [json-3.11.3.tar.gz](https://cann-3rd.obs.cn-north-4.myhuaweicloud.com/json/json-3.11.3.tar.gz) |13| json | 3.11.3 | [json-3.11.3.tar.gz](https://cann-3rd.obs.cn-north-4.myhuaweicloud.com/json/json-3.11.3.tar.gz) |
14| makeself | 2.5.0 | [makeself-release-2.5.0-patch1.tar.gz](https://gitcode.com/cann-src-third-party/makeself/releases/download/release-2.5.0-patch1.0/makeself-release-2.5.0-patch1.tar.gz) |14| makeself | 2.5.0 | [makeself-release-2.5.0-patch1.tar.gz](https://gitcode.com/cann-src-third-party/makeself/releases/download/release-2.5.0-patch1.0/makeself-release-2.5.0-patch1.tar.gz) |
15| eigen | 5.0.0 | [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) |15| eigen | 5.0.0 | [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) |
16-| protobuf | 25.1.0 | [protobuf-25.1.tar.gz](https://gitcode.com/cann-src-third-party/protobuf/releases/download/v25.1/protobuf-25.1.tar.gz) |16+| protobuf | 25.1 | [protobuf-25.1.tar.gz](https://gitcode.com/cann-src-third-party/protobuf/releases/download/v25.1/protobuf-25.1.tar.gz) |
17| abseil-cpp | 20230802.1 | [abseil-cpp-20230802.1.tar.gz](https://gitcode.com/cann-src-third-party/abseil-cpp/releases/download/20230802.1/abseil-cpp-20230802.1.tar.gz) |17| abseil-cpp | 20230802.1 | [abseil-cpp-20230802.1.tar.gz](https://gitcode.com/cann-src-third-party/abseil-cpp/releases/download/20230802.1/abseil-cpp-20230802.1.tar.gz) |
18| opbase(自CANN 9.0.0及以后版本需要下载) | master | [opbase](https://gitcode.com/cann/opbase) |18| opbase(自CANN 9.0.0及以后版本需要下载) | master | [opbase](https://gitcode.com/cann/opbase) |
19| cann-cmake | master-016 | [cmake-master-016.tar.gz](https://cann-3rd.obs.cn-north-4.myhuaweicloud.com/cmake/cmake-master-016.tar.gz) |19| cann-cmake | master-016 | [cmake-master-016.tar.gz](https://cann-3rd.obs.cn-north-4.myhuaweicloud.com/cmake/cmake-master-016.tar.gz) |
@@ -203,7 +203,7 @@
203 203 
2042. **查看产物位置**2042. **查看产物位置**
205 205 
206- 执行命令206+ 执行命令
207 207 
208 ```bash208 ```bash
209 spack location -i cann-ops-math209 spack location -i cann-ops-math
Mdocs/zh/install/spack_quick_install.md+8-8
@@ -45,7 +45,7 @@ spack config --scope system add "config:install_tree:root:/opt/spack"
45```45```
46 46 
47普通用户建议使用个人home下目录,避免被其他用户修改导致环境不稳定47普通用户建议使用个人home下目录,避免被其他用户修改导致环境不稳定
48-管理员建议使用全局路径,便于使用[Spack串联能力](https://spack.readthedocs.io/en/latest/chain.html)统一安装软件提供给其他用户使用:48+管理员建议使用全局路径,便于使用[Spack串联能力](https://spack.readthedocs.io/en/latest/chain.html)统一安装软件提供给其他用户使用
49 49 
50##### 步骤3:添加外部已安装工具到Spack环境中50##### 步骤3:添加外部已安装工具到Spack环境中
51 51 
@@ -55,7 +55,7 @@ spack compiler find # 配置本地已安装的编译器,如gcc
55 55 
56##### 步骤4:配置Spack官方仓库gitcode镜像源56##### 步骤4:配置Spack官方仓库gitcode镜像源
57 57 
58-Spack默认官方仓库地址为[github地址](https://github.com/spack/spack-packages.git),可以通过如下方式修改默认官方仓库,使用GitCode镜像源起到加速作用:58+Spack默认官方仓库地址为[Spack官方github地址](https://github.com/spack/spack-packages),可以通过如下方式修改默认官方仓库,使用GitCode镜像源起到加速作用:
59修改~/.spack目录下的repos.yaml,如不存在则新建该文件,写入以下内容59修改~/.spack目录下的repos.yaml,如不存在则新建该文件,写入以下内容
60 60 
61```yaml61```yaml
@@ -173,7 +173,7 @@ spack env activate cann-dev-env
173如果不想继续使用Spack,可使用`prepare_cann_env.sh`的clean参数卸载Spack并通过**重启终端**来清理环境变量173如果不想继续使用Spack,可使用`prepare_cann_env.sh`的clean参数卸载Spack并通过**重启终端**来清理环境变量
174 174 
175```bash175```bash
176-source prepare_cann_env.sh clean176+source spack/prepare_cann_env.sh clean
177```177```
178 178 
179## 五、ops-math的Spack构建变体179## 五、ops-math的Spack构建变体
@@ -197,7 +197,7 @@ spack find #列出所有软件包
197spack find -L #显示完整哈希197spack find -L #显示完整哈希
198spack find <package-name> #按包名过滤,不提供则为全部包198spack find <package-name> #按包名过滤,不提供则为全部包
199spack find --deps <package-name> #显示该包的依赖树199spack find --deps <package-name> #显示该包的依赖树
200-spack find --explicit <package-name> #显示手动安装的级包 缩写:-e200+spack find --explicit <package-name> #显示手动安装的级包 缩写:-e
201spack find -p <package-name> #显示完整安装路径201spack find -p <package-name> #显示完整安装路径
202spack find -lv <package-name> #显示完整哈希和变体信息202spack find -lv <package-name> #显示完整哈希和变体信息
203 203 
@@ -226,7 +226,7 @@ spack versions <package-name>
226 226 
227#清除Spack构建缓存:227#清除Spack构建缓存:
228spack clean228spack clean
229-spack clean -all #将旧源码与旧构建记录及缓存全部清除,慎用229+spack clean --all #将旧源码与旧构建记录及缓存全部清除,慎用
230 230 
231#查看当前软件包信息:查看该软件包所有支持的版本,Spack默认倾向于最新版本231#查看当前软件包信息:查看该软件包所有支持的版本,Spack默认倾向于最新版本
232spack info <package-name>232spack info <package-name>
@@ -303,7 +303,7 @@ $HOME/spack/var/spack/environments
303 303 
304### 1.执行prepare_cann_env.sh出现报错怎么办304### 1.执行prepare_cann_env.sh出现报错怎么办
305 305 
306-如果机器在以前安装过Spack可能会导致环境变量或配置文件残留,可在备份关键数据后执行```source prepare_cann_env.sh clean```并重启终端后重新执行306+如果机器在以前安装过Spack可能会导致环境变量或配置文件残留,可在备份关键数据后执行```source spack/prepare_cann_env.sh clean```并重启终端后重新执行
307 307 
308### 2.安装过程中出现ssl类报错怎么办308### 2.安装过程中出现ssl类报错怎么办
309 309 
@@ -311,8 +311,8 @@ $HOME/spack/var/spack/environments
311 311 
312### 3.不想使用Spack了如何卸载312### 3.不想使用Spack了如何卸载
313 313 
314-执行```source prepare_cann_env.sh clean```可自动卸载Spack,root用户还需要执行```rm -rf /opt/spack```卸载所有Spack软件包314+执行```source spack/prepare_cann_env.sh clean```可自动卸载Spack,root用户还需要执行```rm -rf /opt/spack```卸载所有Spack软件包
315 315 
316## 九、更多Spack操作请参考官方文档316## 九、更多Spack操作请参考官方文档
317 317 
318-<https://spack.readthedocs.io/en/latest/>318+[Spack 官方文档](https://spack.readthedocs.io/en/latest/)
Mdocs/zh/invocation/quick_op_invocation.md+10-10
@@ -43,12 +43,12 @@
43 # bash build.sh --experimental --run_example abs eager cust --vendor_name=custom43 # bash build.sh --experimental --run_example abs eager cust --vendor_name=custom
44 ```44 ```
45 45 
46- - \$\{op\}:表示待执行算子,算子名小写下划线形式,如abs。46+ - $\{op\}:表示待执行算子,算子名小写下划线形式,如abs。
47- - \$\{mode\}:表示调用方式,目前支持eager(aclnn调用)、graph(图模式调用)。47+ - $\{mode\}:表示调用方式,目前支持eager(aclnn调用)、graph(图模式调用)。
48- - \$\{pkg_mode\}:表示包模式,目前仅支持cust,即自定义算子包。48+ - $\{pkg_mode\}:表示包模式,目前仅支持cust,即自定义算子包。
49- - \$\{vendor\_name\}(可选):与构建的自定义算子包设置一致,默认名为custom。49+ - $\{vendor\_name\}(可选):与构建的自定义算子包设置一致,默认名为custom。
50- - \$\{soc_version\}(可选):表示NPU型号。50+ - $\{soc_version\}(可选):表示NPU型号。
51- - \$\{experimental\}(可选):表示执行用户保存在experimental贡献目录下的算子。51+ - $\{experimental\}(可选):表示执行用户保存在experimental贡献目录下的算子。
52 52 
53 说明:\$\{mode\}为graph时,不指定\$\{pkg_mode\}和\$\{vendor\_name\}53 说明:\$\{mode\}为graph时,不指定\$\{pkg_mode\}和\$\{vendor\_name\}
54 54 
@@ -170,7 +170,7 @@
170 170 
171 为方便理解,以`AddExample`算子为例,调用脚本如下,仅供参考,全量代码参见[test_aclnn_add_example.cpp](../../../examples/add_example/examples/test_aclnn_add_example.cpp)。171 为方便理解,以`AddExample`算子为例,调用脚本如下,仅供参考,全量代码参见[test_aclnn_add_example.cpp](../../../examples/add_example/examples/test_aclnn_add_example.cpp)。
172 172 
173- ```Cpp173+ ```cpp
174 int main()174 int main()
175 {175 {
176 // 1.调用acl进行device/stream初始化176 // 1.调用acl进行device/stream初始化
@@ -254,7 +254,7 @@
254 254 
255 - **调用自定义算子**:依赖自定义算子包255 - **调用自定义算子**:依赖自定义算子包
256 256 
257- ```bash257+ ```cmake
258 cmake_minimum_required(VERSION 3.14)258 cmake_minimum_required(VERSION 3.14)
259 # 设置工程名259 # 设置工程名
260 project(ACLNN_EXAMPLE)260 project(ACLNN_EXAMPLE)
@@ -423,7 +423,7 @@
423 423 
424 在目标算子`examples`目录下,新建调用脚本test\_geir\_\$\{op\_name\}.cpp,\$\{op\_name\}表示目标算子名。以`AddExample`算子为例,调用脚本如下,仅供参考,全量代码参见[test_geir_add_example.cpp](../../../examples/add_example/examples/test_geir_add_example.cpp)。424 在目标算子`examples`目录下,新建调用脚本test\_geir\_\$\{op\_name\}.cpp,\$\{op\_name\}表示目标算子名。以`AddExample`算子为例,调用脚本如下,仅供参考,全量代码参见[test_geir_add_example.cpp](../../../examples/add_example/examples/test_geir_add_example.cpp)。
425 425 
426- ```CPP426+ ```cpp
427 int main() {427 int main() {
428 // 1.创建图对象428 // 1.创建图对象
429 Graph graph(graphName);429 Graph graph(graphName);
@@ -470,7 +470,7 @@
470 470 
471 在test\_geir\_\$\{op\_name\}.cpp同级目录下创建CMakeLists.txt文件,以`AddExample`算子为例,示例如下,请根据实际情况自行修改。471 在test\_geir\_\$\{op\_name\}.cpp同级目录下创建CMakeLists.txt文件,以`AddExample`算子为例,示例如下,请根据实际情况自行修改。
472 472 
473- ```bash473+ ```cmake
474 cmake_minimum_required(VERSION 3.14)474 cmake_minimum_required(VERSION 3.14)
475 475 
476 # 设置工程名476 # 设置工程名
Mdocs/zh/op_api_list.md+4-4
@@ -78,7 +78,7 @@
78| [aclnnCast](../../math/cast/docs/aclnnCast.md) | 实现张量数据类型转换。 | 默认确定性实现| 默认确定性实现|78| [aclnnCast](../../math/cast/docs/aclnnCast.md) | 实现张量数据类型转换。 | 默认确定性实现| 默认确定性实现|
79| [aclnnCat](../../conversion/concat_d/docs/aclnnCat.md) | 将tensors中所有tensor按照维度dim进行级联。 | 默认确定性实现| 默认确定性实现 |79| [aclnnCat](../../conversion/concat_d/docs/aclnnCat.md) | 将tensors中所有tensor按照维度dim进行级联。 | 默认确定性实现| 默认确定性实现 |
80| [aclnnCdist](../../math/cdist/docs/aclnnCdist.md) | 计算两个向量集合中每个点之间的p范数距离。 | 默认确定性实现| - |80| [aclnnCdist](../../math/cdist/docs/aclnnCdist.md) | 计算两个向量集合中每个点之间的p范数距离。 | 默认确定性实现| - |
81-| [aclnnCdistBackward](../../math/cdist_grad/docs/aclnnCdistBackward.md) | 完成aclnnCdist的反向 | 默认确定性实现| 默认确定性实现 |81+| [aclnnCdistBackward](../../math/cdist_grad/docs/aclnnCdistBackward.md) | 完成aclnnCdist的反向 | 默认确定性实现| 默认确定性实现 |
82| [aclnnCeil&aclnnInplaceCeil](../../math/ceil/docs/aclnnCeil&aclnnInplaceCeil.md) | 返回输入tensor中每个元素向上取整的结果。 | 默认确定性实现| 默认确定性实现|82| [aclnnCeil&aclnnInplaceCeil](../../math/ceil/docs/aclnnCeil&aclnnInplaceCeil.md) | 返回输入tensor中每个元素向上取整的结果。 | 默认确定性实现| 默认确定性实现|
83| [aclnnChannelShuffle](../../conversion/transpose/docs/aclnnChannelShuffle.md) | 将(\*, C, H, W)张量的channels分成g个组,然后将每个通道组中的通道进行随机重排,最后将所有通道合并输出,同时保持最终输出张量的shape和输入张量保持一致。 | 默认确定性实现| 默认确定性实现|83| [aclnnChannelShuffle](../../conversion/transpose/docs/aclnnChannelShuffle.md) | 将(\*, C, H, W)张量的channels分成g个组,然后将每个通道组中的通道进行随机重排,最后将所有通道合并输出,同时保持最终输出张量的shape和输入张量保持一致。 | 默认确定性实现| 默认确定性实现|
84| [aclnnChunkCat](../../conversion/chunk_cat/docs/aclnnChunkCat.md) | 将tensors中所有tensor先按照维度dim切分为numChunks块,再按照dim后一维进行级联,最后转换为out的数据类型。 | 默认确定性实现| 默认确定性实现|84| [aclnnChunkCat](../../conversion/chunk_cat/docs/aclnnChunkCat.md) | 将tensors中所有tensor先按照维度dim切分为numChunks块,再按照dim后一维进行级联,最后转换为out的数据类型。 | 默认确定性实现| 默认确定性实现|
@@ -238,11 +238,11 @@
238| [aclnnPowTensorScalar&aclnnInplacePowTensorScalar](../../math/pow/docs/aclnnPowTensorScalar&aclnnInplacePowTensorScalar.md) | exponent每个元素作为input对应元素的幂完成计算。 | 默认确定性实现| 默认确定性实现 |238| [aclnnPowTensorScalar&aclnnInplacePowTensorScalar](../../math/pow/docs/aclnnPowTensorScalar&aclnnInplacePowTensorScalar.md) | exponent每个元素作为input对应元素的幂完成计算。 | 默认确定性实现| 默认确定性实现 |
239| [aclnnPowTensorTensor&aclnnInplacePowTensorTensor](../../math/pow/docs/aclnnPowTensorTensor&aclnnInplacePowTensorTensor.md) | exponent每个元素作为input对应元素的幂完成计算。 | 默认确定性实现| 默认确定性实现 |239| [aclnnPowTensorTensor&aclnnInplacePowTensorTensor](../../math/pow/docs/aclnnPowTensorTensor&aclnnInplacePowTensorTensor.md) | exponent每个元素作为input对应元素的幂完成计算。 | 默认确定性实现| 默认确定性实现 |
240| [aclnnPrecisionCompare](../../math/precision_compare/docs/aclnnPrecisionCompare.md) | 计算两个Tensor元素是否相同,进行精度比对,返回算子执行的状态码,不会触发管理面故障上报。 | 默认确定性实现| 默认确定性实现 |240| [aclnnPrecisionCompare](../../math/precision_compare/docs/aclnnPrecisionCompare.md) | 计算两个Tensor元素是否相同,进行精度比对,返回算子执行的状态码,不会触发管理面故障上报。 | 默认确定性实现| 默认确定性实现 |
241-| [aclnnProd](../../math/reduce_prod/docs/aclnnProd.md) | 回输入tensor中所有元素的乘积。 | 默认确定性实现| 默认确定性实现|241+| [aclnnProd](../../math/reduce_prod/docs/aclnnProd.md) | 回输入tensor中所有元素的乘积。 | 默认确定性实现| 默认确定性实现|
242| [aclnnProdDim](../../math/reduce_prod/docs/aclnnProdDim.md) | 返回输入tensor给定维度上每行的乘积。 | 默认确定性实现| 默认确定性实现|242| [aclnnProdDim](../../math/reduce_prod/docs/aclnnProdDim.md) | 返回输入tensor给定维度上每行的乘积。 | 默认确定性实现| 默认确定性实现|
243| [aclnnQr](../../math/qr/docs/aclnnQr.md) | 对输入Tensor进行正交分解。 | 默认确定性实现| - |243| [aclnnQr](../../math/qr/docs/aclnnQr.md) | 对输入Tensor进行正交分解。 | 默认确定性实现| - |
244-| [aclnnInplaceRandomWithoutFromTo](../../random/dsa_random_uniform/docs/aclnnInplaceRandomWithoutFromTo.md) | 给定seed/offset(int64),默认从 [0, DTYPE_MAX_VALUE] 的离散均匀分布中采样整数填充selfRef张量。 | -| 默认确定性实现 |244+| [aclnnInplaceRandomWithoutFromTo](../../random/dsa_random_uniform/docs/aclnnInplaceRandomWithoutFromTo.md) | 给定seed/offset(int64),默认从 [0, DTYPE_MAX_VALUE] 的离散均匀分布中采样整数填充selfRef张量。 | - | 默认确定性实现 |
245-| [aclnnInplaceRandomWithoutFromToTensor](../../random/dsa_random_uniform/docs/aclnnInplaceRandomWithoutFromToTensor.md) | 给定seed/offset(tensor),默认从 [0, DTYPE_MAX_VALUE] 的离散均匀分布中采样整数填充selfRef张量。 | -| 默认确定性实现 |245+| [aclnnInplaceRandomWithoutFromToTensor](../../random/dsa_random_uniform/docs/aclnnInplaceRandomWithoutFromToTensor.md) | 给定seed/offset(tensor),默认从 [0, DTYPE_MAX_VALUE] 的离散均匀分布中采样整数填充selfRef张量。 | - | 默认确定性实现 |
246| [aclnnRandperm](../../random/stateless_randperm/docs/aclnnRandperm.md) | 返回从0到n-1的整数随机排列。 | 默认确定性实现| 默认确定性实现 |246| [aclnnRandperm](../../random/stateless_randperm/docs/aclnnRandperm.md) | 返回从0到n-1的整数随机排列。 | 默认确定性实现| 默认确定性实现 |
247| [aclnnRange](../../math/range/docs/aclnnRange.md) | 从start起始到end结束按照step的间隔取值。 | 默认确定性实现| 默认确定性实现 |247| [aclnnRange](../../math/range/docs/aclnnRange.md) | 从start起始到end结束按照step的间隔取值。 | 默认确定性实现| 默认确定性实现 |
248| [aclnnReal](../../math/real/docs/aclnnReal.md) | 为输入张量的每一个元素取实数部分。 | 默认确定性实现| - |248| [aclnnReal](../../math/real/docs/aclnnReal.md) | 为输入张量的每一个元素取实数部分。 | 默认确定性实现| - |
Mdocs/zh/op_list.md+10-10
@@ -594,7 +594,7 @@
594 <td>×</td>594 <td>×</td>
595 <td>×</td>595 <td>×</td>
596 <td>AI Core</td>596 <td>AI Core</td>
597- <td>计算Tensor(A)每个元素和Tensor(B)中对应位置元素的按位或。</td>597+ <td>计算self的累积最小值,并返回最小值以及对应的索引。</td>
598 </tr>598 </tr>
599 <tr>599 <tr>
600 <td>math</td>600 <td>math</td>
@@ -1914,7 +1914,7 @@
1914 <td>√</td>1914 <td>√</td>
1915 <td>√</td>1915 <td>√</td>
1916 <td>AI Core/AI CPU</td>1916 <td>AI Core/AI CPU</td>
1917- <td>该算子对输入的每个元素进行正弦后输出。</td>1917+ <td>该算子对输入的每个元素进行双曲正弦操作后输出。</td>
1918 </tr>1918 </tr>
1919 <tr>1919 <tr>
1920 <td>math</td>1920 <td>math</td>
@@ -2104,7 +2104,7 @@
2104 <td>√</td>2104 <td>√</td>
2105 <td>√</td>2105 <td>√</td>
2106 <td>AI Core/AI CPU</td>2106 <td>AI Core/AI CPU</td>
2107- <td>该算子暂无Ascend C代码实现,欢迎开发者补充贡献,贡献方式参考<a href="../../CONTRIBUTING.md">贡献南</a>。</td>2107+ <td>对输入tensor沿着repeats中对每个维度定的复制次数进行复制。</td>
2108 </tr>2108 </tr>
2109 <tr>2109 <tr>
2110 <td>math</td>2110 <td>math</td>
@@ -2184,7 +2184,7 @@
2184 <td>√</td>2184 <td>√</td>
2185 <td>√</td>2185 <td>√</td>
2186 <td>AI Core</td>2186 <td>AI Core</td>
2187- <td>对输入Tensor截取整数部分。</td>2187+ <td>求解一个具有方形上或下三角形可逆矩阵A和多个右侧b的方程组。</td>
2188 </tr>2188 </tr>
2189 <tr>2189 <tr>
2190 <td>math</td>2190 <td>math</td>
@@ -2234,7 +2234,7 @@
2234 <td>×</td>2234 <td>×</td>
2235 <td>×</td>2235 <td>×</td>
2236 <td>AI Core</td>2236 <td>AI Core</td>
2237- <td>将输入的张量沿指定轴重复多次但不进行广播</td>2237+ <td>计算x1 * log(x2)逐元素运算,支持广播</td>
2238 </tr>2238 </tr>
2239 <tr>2239 <tr>
2240 <td>math</td>2240 <td>math</td>
@@ -2258,7 +2258,7 @@
2258 </tr>2258 </tr>
2259 <tr>2259 <tr>
2260 <td>math</td>2260 <td>math</td>
2261- <td><a href="../../math/zero_op/README.md">zeros_like</a></td>2261+ <td><a href="../../math/zero_op/README.md">zero_op</a></td>
2262 <td>√</td>2262 <td>√</td>
2263 <td>√</td>2263 <td>√</td>
2264 <td>√</td>2264 <td>√</td>
@@ -2594,7 +2594,7 @@
2594 <td>×</td>2594 <td>×</td>
2595 <td>√</td>2595 <td>√</td>
2596 <td>AI Core</td>2596 <td>AI Core</td>
2597- <td>允许用户通过定新的形状和步长来创建一个与原张量共享相同数据内存的张量视图。</td>2597+ <td>允许用户通过定新的形状和步长来创建一个与原张量共享相同数据内存的张量视图。</td>
2598 </tr>2598 </tr>
2599 <tr>2599 <tr>
2600 <td>conversion</td>2600 <td>conversion</td>
@@ -2974,7 +2974,7 @@
2974 <td>√</td>2974 <td>√</td>
2975 <td>√</td>2975 <td>√</td>
2976 <td>AI Core</td>2976 <td>AI Core</td>
2977- <td>完成张量置。</td>2977+ <td>将输入的self张量的最后二维(按shape从左向右数)沿对角线的左下部分参数diagonal可正可负,默认为零,正数表示主对角线向右上方向移动,负数表示主对角线向左下方向移动。</td>
2978 </tr>2978 </tr>
2979 <tr>2979 <tr>
2980 <td>conversion</td>2980 <td>conversion</td>
@@ -3214,7 +3214,7 @@
3214 <td>√</td>3214 <td>√</td>
3215 <td>√</td>3215 <td>√</td>
3216 <td>AI Core</td>3216 <td>AI Core</td>
3217- <td>根据随机种子和随机便宜量和保活系数生成随机掩码mask。</td>3217+ <td>根据随机种子和随机偏移量和保活系数生成随机掩码mask。</td>
3218 </tr>3218 </tr>
3219 <tr>3219 <tr>
3220 <td>random</td>3220 <td>random</td>
@@ -3244,7 +3244,7 @@
3244 <td>×</td>3244 <td>×</td>
3245 <td>√</td>3245 <td>√</td>
3246 <td>AI Core</td>3246 <td>AI Core</td>
3247- <td>根据输入x中值为true,选择出对应索引列表,其次根据seed, offset调用philox_random生成索引列表长度个随机数,利用随机数对索引列表做洗牌算法,得到随机交换后索引列表,返回count个。</td>3247+ <td>根据输入x中值为true,选择出对应索引列表,其次根据seed, offset调用philox_random生成索引列表长度个随机数,利用随机数对索引列表做洗牌算法,得到随机交换后索引列表,返回count个。</td>
3248 </tr>3248 </tr>
3249 <tr>3249 <tr>
3250 <td>random</td>3250 <td>random</td>