已合并
fix: Modify the AIDD docs issue #1195
zwj223创建于 7月29日
fix: Modify the AIDD docs issue #1195
已合并
共 19 个文件变更+104-106
| @@ -13,7 +13,7 @@ ops-cv算子首个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 | ```text | 18 | ```text |
| 19 | 版本目录说明如下： | 19 | 版本目录说明如下： |
| @@ -40,9 +40,9 @@ | |||
| 40 | 40 | ||
| 41 | - 撰写清晰的提交信息： | 41 | - 撰写清晰的提交信息： |
| 42 | 42 | ||
| 43 | - ```bash | 43 | + ```text |
| 44 | 简短说明（不超过50字符） | 44 | 简短说明（不超过50字符） |
| 45 | - | 45 | + |
| 46 | 如有必要，在此处进行更详细描述。说明修改的原因和内容，而不是具体改了什么（代码本身会展示）。 | 46 | 如有必要，在此处进行更详细描述。说明修改的原因和内容，而不是具体改了什么（代码本身会展示）。 |
| 47 | 关联的Issue: #123 | 47 | 关联的Issue: #123 |
| 48 | ``` | 48 | ``` |
| @@ -84,7 +84,7 @@ | |||
| 84 | - 图片： | 84 | - 图片： |
| 85 | - 常用格式：推荐png格式，风格尽量与已有图片保持一致。 | 85 | - 常用格式：推荐png格式，风格尽量与已有图片保持一致。 |
| 86 | - 分辨率与清晰度：需清晰且尺寸适中，避免模糊或过度压缩。 | 86 | - 分辨率与清晰度：需清晰且尺寸适中，避免模糊或过度压缩。 |
| 87 | - - 文件大小：单张图片不建议超过10M。 | 87 | + - 文件大小：单张图片不建议超过10MB。 |
| 88 | - 版权：所有引用的图片、文献等资源，请确保合规性。 | 88 | - 版权：所有引用的图片、文献等资源，请确保合规性。 |
| 89 | 89 | ||
| 90 | ## 获取帮助 | 90 | ## 获取帮助 |
| @@ -2,7 +2,7 @@ | |||
| 2 | 2 | ||
| 3 | ## 使用须知 | 3 | ## 使用须知 |
| 4 | 4 | ||
| 5 | -本指南旨在帮助您快速上手CANN和`ops-cv`算子仓的使用。为方便您快速了解算子开发全流程，以**AddExample**算子为实践对象，算子源码位于`ops-cv/examples/add_example`，操作流程如下： | 5 | +本指南旨在帮助开发者快速上手CANN和`ops-cv`算子仓的使用。为方便快速了解算子开发全流程，以**AddExample**算子为实践对象，算子源码位于`ops-cv/examples/add_example`，操作流程如下： |
| 6 | 6 | ||
| 7 | 1. **[前提条件](../README.md)**：参考项目README完成环境准备和源码下载，此处不再赘述。快速入门场景**推荐CANNLab或Docker部署**，操作简单。 | 7 | 1. **[前提条件](../README.md)**：参考项目README完成环境准备和源码下载，此处不再赘述。快速入门场景**推荐CANNLab或Docker部署**，操作简单。 |
| 8 | 8 | ||
| @@ -178,9 +178,9 @@ __aicore__ inline void AddExample<T>::Compute(int64_t currentNum) | |||
| 178 | 178 | ||
| 179 | - **printf** | 179 | - **printf** |
| 180 | 180 | ||
| 181 | - 该接口支持打印Scalar类型数据，如整数、字符型、布尔型等，详细介绍请参见[《Ascend C API》](https://hiascend.com/document/redirect/CannCommunityAscendCApi)中"算子调测API > printf"。 | 181 | + 该接口支持打印Scalar类型数据，如整数、字符型、布尔型等，详细介绍请参见[《Ascend C API》](https://hiascend.com/document/redirect/CannCommunityAscendCApi)中“算子调测API > printf”。 |
| 182 | 182 | ||
| 183 | - ```c++ | 183 | + ```cpp |
| 184 | blockLength_ = (remainderLength > tilingData->blockFactor) ? tilingData->blockFactor : remainderLength; | 184 | blockLength_ = (remainderLength > tilingData->blockFactor) ? tilingData->blockFactor : remainderLength; |
| 185 | ubLength_ = tilingData->ubFactor; | 185 | ubLength_ = tilingData->ubFactor; |
| 186 | // 打印当前核计算Block长度 | 186 | // 打印当前核计算Block长度 |
| @@ -191,7 +191,7 @@ __aicore__ inline void AddExample<T>::Compute(int64_t currentNum) | |||
| 191 | 191 | ||
| 192 | 该接口支持Dump指定Tensor的内容，同时支持打印自定义附加信息，比如当前行号等，详细介绍请参见[《Ascend C API》](https://hiascend.com/document/redirect/CannCommunityAscendCApi)中“算子调测API > DumpTensor”。 | 192 | 该接口支持Dump指定Tensor的内容，同时支持打印自定义附加信息，比如当前行号等，详细介绍请参见[《Ascend C API》](https://hiascend.com/document/redirect/CannCommunityAscendCApi)中“算子调测API > DumpTensor”。 |
| 193 | 193 | ||
| 194 | - ```c++ | 194 | + ```cpp |
| 195 | AscendC::LocalTensor<T> xLocal = inputQueueX.DeQue<T>(); | 195 | AscendC::LocalTensor<T> xLocal = inputQueueX.DeQue<T>(); |
| 196 | AscendC::LocalTensor<T> yLocal = inputQueueY.DeQue<T>(); | 196 | AscendC::LocalTensor<T> yLocal = inputQueueY.DeQue<T>(); |
| 197 | AscendC::LocalTensor<T> zLocal = outputQueueZ.AllocTensor<T>(); | 197 | AscendC::LocalTensor<T> zLocal = outputQueueZ.AllocTensor<T>(); |
| @@ -232,7 +232,7 @@ __aicore__ inline void AddExample<T>::Compute(int64_t currentNum) | |||
| 232 | 232 | ||
| 233 | **修改输入/输出数据**：修改输入、输出的shape信息，以及初始化数据，构造相应的输入、输出tensor。 | 233 | **修改输入/输出数据**：修改输入、输出的shape信息，以及初始化数据，构造相应的输入、输出tensor。 |
| 234 | 234 | ||
| 235 | -```c++ | 235 | +```cpp |
| 236 | int main() { | 236 | int main() { |
| 237 | // ... 初始化代码... | 237 | // ... 初始化代码... |
| 238 | 238 | ||
| @@ -266,4 +266,4 @@ int main() { | |||
| 266 | 266 | ||
| 267 | ## 结语 | 267 | ## 结语 |
| 268 | 268 | ||
| 269 | -体验完上述流程，您已基本完成算子开发过程，如果您想进一步贡献新算子或学习更多高阶开发、调试等技能，请访问本项目README学习[进阶教程](../README.md#学习教程)和[贡献指南](../README.md#相关信息)等。 | 269 | +体验完上述流程，已基本完成算子开发过程，如果想要进一步贡献新算子或学习更多高阶开发、调试等技能，请访问本项目README学习[进阶教程](../README.md#学习教程)和[贡献指南](../README.md#相关信息)等。 |
| @@ -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 | ||
| @@ -11,11 +11,11 @@ broadcast（广播）描述了算子在运算期间如何处理不同形状的 | |||
| 11 | 一般进行广播计算时，需要理解以下规则： | 11 | 一般进行广播计算时，需要理解以下规则： |
| 12 | 12 | ||
| 13 | - 规则1：如果数组间维度数不一致，所有数组向最长形状的数组看齐，形状不足的部分在**左侧**填充1，直至维度数相同。 | 13 | - 规则1：如果数组间维度数不一致，所有数组向最长形状的数组看齐，形状不足的部分在**左侧**填充1，直至维度数相同。 |
| 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 | > 说明： |
| @@ -25,7 +25,7 @@ broadcast（广播）描述了算子在运算期间如何处理不同形状的 | |||
| 25 | 25 | ||
| 26 | 基于上述规则，广播过程一般先按**规则1**进行扩维，再按**规则2**进行形状拉伸，具体例子如下： | 26 | 基于上述规则，广播过程一般先按**规则1**进行扩维，再按**规则2**进行形状拉伸，具体例子如下： |
| 27 | 27 | ||
| 28 | -```tex | 28 | +```text |
| 29 | 假设a.shape=(2,2,3)，取值形如: | 29 | 假设a.shape=(2,2,3)，取值形如: |
| 30 | [[[1 2 3],[4 5 6]], | 30 | [[[1 2 3],[4 5 6]], |
| 31 | [[1 2 3],[4 5 6]]] | 31 | [[1 2 3],[4 5 6]]] |
| @@ -19,27 +19,27 @@ | |||
| 19 | 19 | ||
| 20 | ```CMake | 20 | ```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 requirement | 23 | # CMake lowest version requirement |
| 24 | cmake_minimum_required(VERSION 3.14) | 24 | cmake_minimum_required(VERSION 3.14) |
| 25 | - | 25 | + |
| 26 | # 设置工程名 | 26 | # 设置工程名 |
| 27 | project(ACLNN_EXAMPLE) | 27 | project(ACLNN_EXAMPLE) |
| 28 | - | 28 | + |
| 29 | # Compile options | 29 | # Compile options |
| 30 | add_compile_options(-std=c++11) | 30 | add_compile_options(-std=c++11) |
| 31 | - | 31 | + |
| 32 | # 设置编译选项 | 32 | # 设置编译选项 |
| 33 | - set(CMAKE_RUNTIME_OUTPUT_DIRECTORY "./bin") | 33 | + set(CMAKE_RUNTIME_OUTPUT_DIRECTORY "./bin") |
| 34 | set(CMAKE_CXX_FLAGS_DEBUG "-fPIC -O0 -g -Wall") | 34 | set(CMAKE_CXX_FLAGS_DEBUG "-fPIC -O0 -g -Wall") |
| 35 | set(CMAKE_CXX_FLAGS_RELEASE "-fPIC -O2 -Wall") | 35 | set(CMAKE_CXX_FLAGS_RELEASE "-fPIC -O2 -Wall") |
| 36 | - | 36 | + |
| 37 | # 设置可执行文件名（如opapi_test），并指定待运行算子文件*.cpp所在目录 | 37 | # 设置可执行文件名（如opapi_test），并指定待运行算子文件*.cpp所在目录 |
| 38 | add_executable(opapi_test | 38 | add_executable(opapi_test |
| 39 | test_grid_sampler2_d.cpp) | 39 | test_grid_sampler2_d.cpp) |
| 40 | - | 40 | + |
| 41 | # 设置ASCEND_PATH（CANN软件包目录，请根据实际路径修改）和INCLUDE_BASE_DIR（头文件目录） | 41 | # 设置ASCEND_PATH（CANN软件包目录，请根据实际路径修改）和INCLUDE_BASE_DIR（头文件目录） |
| 42 | - if(NOT "$ENV{ASCEND_CUSTOM_PATH}" STREQUAL "") | 42 | + if(NOT "$ENV{ASCEND_CUSTOM_PATH}" STREQUAL "") |
| 43 | set(ASCEND_PATH $ENV{ASCEND_CUSTOM_PATH}) | 43 | set(ASCEND_PATH $ENV{ASCEND_CUSTOM_PATH}) |
| 44 | else() | 44 | else() |
| 45 | set(ASCEND_PATH "/usr/local/Ascend/cann") | 45 | set(ASCEND_PATH "/usr/local/Ascend/cann") |
| @@ -49,20 +49,20 @@ | |||
| 49 | ${INCLUDE_BASE_DIR} | 49 | ${INCLUDE_BASE_DIR} |
| 50 | ${INCLUDE_BASE_DIR}/aclnn | 50 | ${INCLUDE_BASE_DIR}/aclnn |
| 51 | ) | 51 | ) |
| 52 | - | 52 | + |
| 53 | # 设置链接的库文件路径 | 53 | # 设置链接的库文件路径 |
| 54 | target_link_libraries(opapi_test PRIVATE | 54 | target_link_libraries(opapi_test PRIVATE |
| 55 | ${ASCEND_PATH}/lib64/libacl_rt.so | 55 | ${ASCEND_PATH}/lib64/libacl_rt.so |
| 56 | ${ASCEND_PATH}/lib64/libnnopbase.so | 56 | ${ASCEND_PATH}/lib64/libnnopbase.so |
| 57 | ${ASCEND_PATH}/lib64/libopapi_math.so | 57 | ${ASCEND_PATH}/lib64/libopapi_math.so |
| 58 | ${ASCEND_PATH}/lib64/libopapi_cv.so) | 58 | ${ASCEND_PATH}/lib64/libopapi_cv.so) |
| 59 | - | 59 | + |
| 60 | # 可执行文件在CMakeLists文件所在目录的bin目录下 | 60 | # 可执行文件在CMakeLists文件所在目录的bin目录下 |
| 61 | install(TARGETS opapi_test DESTINATION ${CMAKE_RUNTIME_OUTPUT_DIRECTORY}) | 61 | install(TARGETS opapi_test DESTINATION ${CMAKE_RUNTIME_OUTPUT_DIRECTORY}) |
| 62 | ``` | 62 | ``` |
| 63 | 63 | ||
| 64 | 对于集合通信和MatMul计算融合、并行的算子，统称为通算融合算子（简称MC2算子），包括AllGatherMatmul、AlltoAllAllGatherBatchMatMul、BatchMatMulReduceScatterAlltoAll、MatmulAllReduce、MatmulAllReduceAddRmsNorm、MatmulReduceScatter等。调用该类算子API时，一般会涉及多线程和HCCL（Huawei Collective Communication Library，集合通信库），因此CMake文件需要额外导入如下内容，否则无法成功编译。 | 64 | 对于集合通信和MatMul计算融合、并行的算子，统称为通算融合算子（简称MC2算子），包括AllGatherMatmul、AlltoAllAllGatherBatchMatMul、BatchMatMulReduceScatterAlltoAll、MatmulAllReduce、MatmulAllReduceAddRmsNorm、MatmulReduceScatter等。调用该类算子API时，一般会涉及多线程和HCCL（Huawei Collective Communication Library，集合通信库），因此CMake文件需要额外导入如下内容，否则无法成功编译。 |
| 65 | - | 65 | + |
| 66 | ```CMake | 66 | ```CMake |
| 67 | # 设置链接的库文件路径 | 67 | # 设置链接的库文件路径 |
| 68 | find_package(Threads REQUIRED) | 68 | find_package(Threads REQUIRED) |
| @@ -76,11 +76,11 @@ | |||
| 76 | ``` | 76 | ``` |
| 77 | 77 | ||
| 78 | 其中“find_package(Threads REQUIRED)”是CMake用于查找线程库的命令，可自动链接线程库依赖的头文件或间接依赖的库文件。 | 78 | 其中“find_package(Threads REQUIRED)”是CMake用于查找线程库的命令，可自动链接线程库依赖的头文件或间接依赖的库文件。 |
| 79 | - | 79 | + |
| 80 | ## 编译与运行 | 80 | ## 编译与运行 |
| 81 | 81 | ||
| 82 | 1. 提前准备好算子的调用代码（\*.cpp）和编译脚本（CMakeLists.txt）。 | 82 | 1. 提前准备好算子的调用代码（\*.cpp）和编译脚本（CMakeLists.txt）。 |
| 83 | - 2. 配置环境变量。 | 83 | + 2. 配置环境变量。 |
| 84 | 84 | ||
| 85 | 安装CANN软件后，使用CANN运行用户登录环境，执行如下命令生效环境变量。 | 85 | 安装CANN软件后，使用CANN运行用户登录环境，执行如下命令生效环境变量。 |
| 86 | 86 | ||
| @@ -88,14 +88,14 @@ | |||
| 88 | source ${INSTALL_DIR}/set_env.sh | 88 | source ${INSTALL_DIR}/set_env.sh |
| 89 | ``` | 89 | ``` |
| 90 | 90 | ||
| 91 | - 其中${INSTALL_DIR}为CANN软件安装后文件存储路径，请根据实际情况替换。 | 91 | + 其中${INSTALL_DIR}为CANN软件安装后文件存储路径，请根据实际情况替换。 |
| 92 | 3. 编译并运行。 | 92 | 3. 编译并运行。 |
| 93 | - 进入CMakeLists.txt所在目录，执行如下命令，新建build目录存放生成的编译文件。 | 93 | - 进入CMakeLists.txt所在目录，执行如下命令，新建build目录存放生成的编译文件。 |
| 94 | 94 | ||
| 95 | ```sh | 95 | ```sh |
| 96 | - mkdir -p build | 96 | + mkdir -p build |
| 97 | ``` | 97 | ``` |
| 98 | - | 98 | + |
| 99 | - 进入build目录，执行cmake命令编译，再执行make命令生成可执行文件。 | 99 | - 进入build目录，执行cmake命令编译，再执行make命令生成可执行文件。 |
| 100 | 100 | ||
| 101 | ```sh | 101 | ```sh |
| @@ -130,14 +130,14 @@ | |||
| 130 | 若执行结果报错，未出现预期结果，可以使用aclGetRecentErrMsg接口获取报错具体信息。 | 130 | 若执行结果报错，未出现预期结果，可以使用aclGetRecentErrMsg接口获取报错具体信息。 |
| 131 | 调用aclnnGridSampler2DGetWorkspaceSize报错获取异常信息示例如下： | 131 | 调用aclnnGridSampler2DGetWorkspaceSize报错获取异常信息示例如下： |
| 132 | 132 | ||
| 133 | - ```sh | 133 | + ```cpp |
| 134 | // input is nullptr | 134 | // input is nullptr |
| 135 | ret = aclnnGridSampler2DGetWorkspaceSize( | 135 | ret = aclnnGridSampler2DGetWorkspaceSize( |
| 136 | input, grid, interpolationMode, paddingMode, alignCorners, out, & workspaceSize, &executor); | 136 | input, grid, interpolationMode, paddingMode, alignCorners, out, & workspaceSize, &executor); |
| 137 | CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnGridSampler2DGetWorkspaceSize failed. ERROR: %d.\n[ERROR msg]%s", ret, aclGetRecentErrMsg()); return ret); | 137 | CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnGridSampler2DGetWorkspaceSize failed. ERROR: %d.\n[ERROR msg]%s", ret, aclGetRecentErrMsg()); return ret); |
| 138 | ``` | 138 | ``` |
| 139 | 139 | ||
| 140 | - 上述构造空指针问题获取报错信息示例如下: | 140 | + 上述构造空指针问题获取报错信息示例如下： |
| 141 | 141 | ||
| 142 | ```sh | 142 | ```sh |
| 143 | aclnnGridSampler2DGetWorkspaceSize failed. ERROR: 161001 | 143 | aclnnGridSampler2DGetWorkspaceSize failed. ERROR: 161001 |
| @@ -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 | ||
| @@ -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 | ||
| @@ -38,12 +38,12 @@ CANN Simulator集成在CANN toolkit包里，参考[环境部署](../install/quic | |||
| 38 | 38 | ||
| 39 | ```bash | 39 | ```bash |
| 40 | # 说明：进入项目根目录，执行如下编译命令，命令仅供参考，详细可以查看算子调用的说明。 | 40 | # 说明：进入项目根目录，执行如下编译命令，命令仅供参考，详细可以查看算子调用的说明。 |
| 41 | -bash build.sh --pkg --soc=ascend950 --vendor_name=custom --ops=add_example | 41 | +bash build.sh --pkg --soc=Ascend950 --vendor_name=custom --ops=add_example |
| 42 | # 安装自定义算子包 | 42 | # 安装自定义算子包 |
| 43 | ./build_out/cann-ops-cv-${vendor_name}_linux-${arch}.run | 43 | ./build_out/cann-ops-cv-${vendor_name}_linux-${arch}.run |
| 44 | ``` | 44 | ``` |
| 45 | 45 | ||
| 46 | -* 参考[aclnn调用](../invocation/quick_op_invocation.md#aclnn-api)完成test_aclnn_add_example.cpp的编译，编出可执行文件test_aclnn_add_example | 46 | +* 参考[aclnn调用](../invocation/quick_op_invocation.md#aclnn-api)完成test_aclnn_add_example.cpp的编译，编译出可执行文件test_aclnn_add_example |
| 47 | 47 | ||
| 48 | ## 执行仿真命令 | 48 | ## 执行仿真命令 |
| 49 | 49 | ||
| @@ -154,7 +154,7 @@ cannsim report [options] | |||
| 154 | 154 | ||
| 155 | ## 使用示例 | 155 | ## 使用示例 |
| 156 | 156 | ||
| 157 | -1. 参考仿真执行执行算子仿真，对比输出示例，确保对应的结果执行正确。 | 157 | +1. 参考仿真执行算子仿真，对比输出示例，确保对应的结果执行正确。 |
| 158 | 2. 执行仿真结果解析命令，可参考以下执行用例。 | 158 | 2. 执行仿真结果解析命令，可参考以下执行用例。 |
| 159 | 159 | ||
| 160 | ```bash | 160 | ```bash |
| @@ -39,7 +39,7 @@ | |||
| 39 | 39 | ||
| 40 | 目录创建是算子开发的重要步骤，为后续代码编写、编译构建和调试提供统一的目录结构和文件组织方式。 | 40 | 目录创建是算子开发的重要步骤，为后续代码编写、编译构建和调试提供统一的目录结构和文件组织方式。 |
| 41 | 41 | ||
| 42 | -本项目`build.sh`，支持快速创建算子目录。进入项目根目录，执行以下命令： | 42 | +本项目提供`build.sh`脚本，支持快速创建算子目录。进入项目根目录，执行以下命令： |
| 43 | 43 | ||
| 44 | ```bash | 44 | ```bash |
| 45 | # 创建指定算子目录，如bash build.sh --genop=examples/example_ops | 45 | # 创建指定算子目录，如bash build.sh --genop=examples/example_ops |
| @@ -56,7 +56,7 @@ Create the initial directory for ${op_name} under ${op_class} success | |||
| 56 | 56 | ||
| 57 | 创建完成后，目录结构如下所示： | 57 | 创建完成后，目录结构如下所示： |
| 58 | 58 | ||
| 59 | -```tex | 59 | +```text |
| 60 | ${op_name} # 替换为实际算子名的小写下划线形式 | 60 | ${op_name} # 替换为实际算子名的小写下划线形式 |
| 61 | ├── examples # 算子调用示例 | 61 | ├── examples # 算子调用示例 |
| 62 | │ ├── test_aclnn_${op_name}.cpp # 算子aclnn调用示例 | 62 | │ ├── test_aclnn_${op_name}.cpp # 算子aclnn调用示例 |
| @@ -134,7 +134,7 @@ Tiling主要切分逻辑。 | |||
| 134 | > 1. **TilingParse**：图模式标准交付件，保留函数定义以满足框架调用规范，无实际逻辑时可置空。 | 134 | > 1. **TilingParse**：图模式标准交付件，保留函数定义以满足框架调用规范，无实际逻辑时可置空。 |
| 135 | > 2. **CompileInfo**：图模式标准交付件，保留函数定义以满足框架调用规范，无实际逻辑时可置空。 | 135 | > 2. **CompileInfo**：图模式标准交付件，保留函数定义以满足框架调用规范，无实际逻辑时可置空。 |
| 136 | 136 | ||
| 137 | -```CPP | 137 | +```cpp |
| 138 | // ${op_name}_tiling.cpp | 138 | // ${op_name}_tiling.cpp |
| 139 | // 1.Tiling需要获取运行环境信息，包括可用核数、UB(Unified Buffer)大小，并将获取到的信息传递给CompileInfo, 自动生成aclnn不调用该函数，直接返回ge::GRAPH_SUCCESS即可。 | 139 | // 1.Tiling需要获取运行环境信息，包括可用核数、UB(Unified Buffer)大小，并将获取到的信息传递给CompileInfo, 自动生成aclnn不调用该函数，直接返回ge::GRAPH_SUCCESS即可。 |
| 140 | static ge::graphStatus TilingParse(gert::TilingParseContext* context) | 140 | static ge::graphStatus TilingParse(gert::TilingParseContext* context) |
| @@ -164,7 +164,7 @@ static ge::graphStatus TilingFunc(gert::TilingContext* context){ | |||
| 164 | OP_CHECK_IF( | 164 | OP_CHECK_IF( |
| 165 | GetPlatformInfo(context, ubSize, coreNum) != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetPlatformInfo error"), | 165 | GetPlatformInfo(context, ubSize, coreNum) != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetPlatformInfo error"), |
| 166 | return ge::GRAPH_FAILED); | 166 | return ge::GRAPH_FAILED); |
| 167 | - | 167 | + |
| 168 | // 2.2获取输入信息 | 168 | // 2.2获取输入信息 |
| 169 | // 获取输入张量shape信息 | 169 | // 获取输入张量shape信息 |
| 170 | auto inputX = context->GetInputShape(0); | 170 | auto inputX = context->GetInputShape(0); |
| @@ -210,7 +210,7 @@ TilingKey是一个算子内为了区分不同的实现而将kernel代码进行 | |||
| 210 | 210 | ||
| 211 | > **说明：** 如需实现复杂参数组合完成分支选择（涉及多TilingKey场景），请参考[《Ascend C算子开发接口》](https://hiascend.com/document/redirect/CannCommunityAscendCApi)中“Utils API > Tiling模版编程 > 模版参数含义”。 | 211 | > **说明：** 如需实现复杂参数组合完成分支选择（涉及多TilingKey场景），请参考[《Ascend C算子开发接口》](https://hiascend.com/document/redirect/CannCommunityAscendCApi)中“Utils API > Tiling模版编程 > 模版参数含义”。 |
| 212 | 212 | ||
| 213 | -```CPP | 213 | +```cpp |
| 214 | // ${op_name}_tiling_key.h | 214 | // ${op_name}_tiling_key.h |
| 215 | ASCENDC_TPL_ARGS_DECL( | 215 | ASCENDC_TPL_ARGS_DECL( |
| 216 | ${op_name}, | 216 | ${op_name}, |
| @@ -226,7 +226,7 @@ ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL( | |||
| 226 | 226 | ||
| 227 | 如需查看详细实现，请参考[add_example_tiling_data.h](../../../examples/add_example/op_kernel/add_example_tiling_data.h)。 | 227 | 如需查看详细实现，请参考[add_example_tiling_data.h](../../../examples/add_example/op_kernel/add_example_tiling_data.h)。 |
| 228 | 228 | ||
| 229 | -```CPP | 229 | +```cpp |
| 230 | // ${op_name}_tiling_data.h | 230 | // ${op_name}_tiling_data.h |
| 231 | struct ${op_name}TilingData { | 231 | struct ${op_name}TilingData { |
| 232 | int64_t totalLength; | 232 | int64_t totalLength; |
| @@ -267,7 +267,7 @@ Kernel入口文件，包含主函数和调度逻辑。 | |||
| 267 | 267 | ||
| 268 | 如需查看详细实现，请参考[add_example.cpp](../../../examples/add_example/op_kernel/add_example.cpp)。 | 268 | 如需查看详细实现，请参考[add_example.cpp](../../../examples/add_example/op_kernel/add_example.cpp)。 |
| 269 | 269 | ||
| 270 | -```CPP | 270 | +```cpp |
| 271 | // 1、核函数定义 | 271 | // 1、核函数定义 |
| 272 | // schMode是一个模板参数，用于支持不同数据类型（如float和int32）的计算路径 | 272 | // schMode是一个模板参数，用于支持不同数据类型（如float和int32）的计算路径 |
| 273 | // __global__ __aicore__表示该函数是个全局函数，可以在AI Core上执行 | 273 | // __global__ __aicore__表示该函数是个全局函数，可以在AI Core上执行 |
| @@ -296,14 +296,14 @@ __global__ __aicore__ void add_example(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR | |||
| 296 | 296 | ||
| 297 | 如需查看详细实现，请参考[add_example.h](../../../examples/add_example/op_kernel/add_example.h)。 | 297 | 如需查看详细实现，请参考[add_example.h](../../../examples/add_example/op_kernel/add_example.h)。 |
| 298 | 298 | ||
| 299 | -```C++ | 299 | +```cpp |
| 300 | // 2、定义Kernel类 | 300 | // 2、定义Kernel类 |
| 301 | template <typename T> | 301 | template <typename T> |
| 302 | class AddExample | 302 | class AddExample |
| 303 | { | 303 | { |
| 304 | public: | 304 | public: |
| 305 | // 默认构造函数，__aicore__表示该函数在AI Core上运行 | 305 | // 默认构造函数，__aicore__表示该函数在AI Core上运行 |
| 306 | - __aicore__ inline AddExample(){}; | 306 | + __aicore__ inline AddExample(){}; |
| 307 | // 初始化函数，用于设置输入输出地址和Tiling切分信息计算 | 307 | // 初始化函数，用于设置输入输出地址和Tiling切分信息计算 |
| 308 | __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, const AddExampleTilingData* tilingData); | 308 | __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, const AddExampleTilingData* tilingData); |
| 309 | // 主处理函数，执行数据拷贝和计算 | 309 | // 主处理函数，执行数据拷贝和计算 |
| @@ -333,7 +333,7 @@ private: | |||
| 333 | GlobalTensor<T> inputGMY_; | 333 | GlobalTensor<T> inputGMY_; |
| 334 | // 输出Z的GM地址 | 334 | // 输出Z的GM地址 |
| 335 | GlobalTensor<T> outputGMZ_; | 335 | GlobalTensor<T> outputGMZ_; |
| 336 | - | 336 | + |
| 337 | // 总数据长度 | 337 | // 总数据长度 |
| 338 | int64_t blockLength_ = 0; | 338 | int64_t blockLength_ = 0; |
| 339 | // 每个block被划分多少块 | 339 | // 每个block被划分多少块 |
| @@ -351,10 +351,10 @@ __aicore__ inline void AddExample<T>::Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, cons | |||
| 351 | blockLength_ = tilingData->totalLength / AscendC::GetBlockNum(); | 351 | blockLength_ = tilingData->totalLength / AscendC::GetBlockNum(); |
| 352 | ... | 352 | ... |
| 353 | // 3.2 初始化GM地址 | 353 | // 3.2 初始化GM地址 |
| 354 | - inputGMX.SetGlobalBuffer((__gm__ T*)x + blockLength_ * AscendC::GetBlockIdx(), blockLength_); | 354 | + inputGMX_.SetGlobalBuffer((__gm__ T*)x + blockLength_ * AscendC::GetBlockIdx(), blockLength_); |
| 355 | ... | 355 | ... |
| 356 | // 3.3 初始化队列长度 | 356 | // 3.3 初始化队列长度 |
| 357 | - pipe.InitBuffer(inputQueueX_, BUFFER_NUM, tileLength_ * sizeof(T)); | 357 | + pipe_.InitBuffer(inputQueueX_, BUFFER_NUM, tileLength_ * sizeof(T)); |
| 358 | ... | 358 | ... |
| 359 | } | 359 | } |
| 360 | 360 | ||
| @@ -392,7 +392,7 @@ __aicore__ inline void AddExample<T>::Process() | |||
| 392 | 392 | ||
| 393 | ## 编译部署 | 393 | ## 编译部署 |
| 394 | 394 | ||
| 395 | -算子开发完成后，需对算子工程进行编译，生成自定义算子安装包\*\.run，详细的编译操作如下： | 395 | +算子开发完成后，需对算子工程进行编译，生成自定义算子安装包*.run，详细的编译操作如下： |
| 396 | 396 | ||
| 397 | 1. **准备工作。** | 397 | 1. **准备工作。** |
| 398 | 398 | ||
| @@ -420,7 +420,7 @@ __aicore__ inline void AddExample<T>::Process() | |||
| 420 | ```bash | 420 | ```bash |
| 421 | # 编译指定算子，如bash build.sh --pkg --ops=add_example -j16 | 421 | # 编译指定算子，如bash build.sh --pkg --ops=add_example -j16 |
| 422 | bash build.sh --pkg --soc=${soc_version} --vendor_name=${vendor_name} --ops=${op_list} [-j${n}] | 422 | bash build.sh --pkg --soc=${soc_version} --vendor_name=${vendor_name} --ops=${op_list} [-j${n}] |
| 423 | - | 423 | + |
| 424 | # 编译experimental目录下指定算子 | 424 | # 编译experimental目录下指定算子 |
| 425 | bash build.sh --pkg --soc=${soc_version} --vendor_name=${vendor_name} --ops=${op_list} [--experimental] [-j${n}] | 425 | bash build.sh --pkg --soc=${soc_version} --vendor_name=${vendor_name} --ops=${op_list} [--experimental] [-j${n}] |
| 426 | ``` | 426 | ``` |
| @@ -440,7 +440,7 @@ __aicore__ inline void AddExample<T>::Process() | |||
| 440 | 4. **安装自定义算子包。** | 440 | 4. **安装自定义算子包。** |
| 441 | 441 | ||
| 442 | 执行以下命令进行安装： | 442 | 执行以下命令进行安装： |
| 443 | - | 443 | + |
| 444 | ```bash | 444 | ```bash |
| 445 | # 安装run包 | 445 | # 安装run包 |
| 446 | ./build_out/cann-ops-cv-${vendor_name}_linux-${arch}.run | 446 | ./build_out/cann-ops-cv-${vendor_name}_linux-${arch}.run |
| @@ -498,7 +498,7 @@ UT编写指导如下，如需查看详细实现，请参考样例UT实现[test_a | |||
| 498 | 498 | ||
| 499 | 测试类示例： | 499 | 测试类示例： |
| 500 | 500 | ||
| 501 | -```CPP | 501 | +```cpp |
| 502 | class ${OpName}InfershapeTest : public testing::Test { | 502 | class ${OpName}InfershapeTest : public testing::Test { |
| 503 | protected: | 503 | protected: |
| 504 | static void SetUpTestCase() | 504 | static void SetUpTestCase() |
| @@ -522,7 +522,7 @@ protected: | |||
| 522 | 522 | ||
| 523 | 简化示例： | 523 | 简化示例： |
| 524 | 524 | ||
| 525 | -```CPP | 525 | +```cpp |
| 526 | TEST_F(${OpName}InfershapeTest, test_case_xxx) | 526 | TEST_F(${OpName}InfershapeTest, test_case_xxx) |
| 527 | { | 527 | { |
| 528 | // 1. 构造用例上下文 | 528 | // 1. 构造用例上下文 |
| @@ -563,7 +563,7 @@ UT编写指导如下，如需查看详细实现，请参考样例UT实现[test_a | |||
| 563 | 563 | ||
| 564 | 测试类示例： | 564 | 测试类示例： |
| 565 | 565 | ||
| 566 | -```CPP | 566 | +```cpp |
| 567 | class ${OpName}TilingTest : public testing::Test { | 567 | class ${OpName}TilingTest : public testing::Test { |
| 568 | protected: | 568 | protected: |
| 569 | static void SetUpTestCase() | 569 | static void SetUpTestCase() |
| @@ -589,7 +589,7 @@ protected: | |||
| 589 | 589 | ||
| 590 | 简化示例： | 590 | 简化示例： |
| 591 | 591 | ||
| 592 | -```CPP | 592 | +```cpp |
| 593 | TEST_F(${OpName}TilingTest, test_case_xxx) | 593 | TEST_F(${OpName}TilingTest, test_case_xxx) |
| 594 | { | 594 | { |
| 595 | // 声明结构体并初始化一个结构体变量 | 595 | // 声明结构体并初始化一个结构体变量 |
| @@ -643,7 +643,7 @@ UT编写指导如下，如需查看详细实现，请参考样例UT实现[test_a | |||
| 643 | 643 | ||
| 644 | 测试类示例： | 644 | 测试类示例： |
| 645 | 645 | ||
| 646 | -```CPP | 646 | +```cpp |
| 647 | class ${OpName}KernelTest : public testing::Test { | 647 | class ${OpName}KernelTest : public testing::Test { |
| 648 | protected: | 648 | protected: |
| 649 | static void SetUpTestCase() | 649 | static void SetUpTestCase() |
| @@ -670,7 +670,7 @@ protected: | |||
| 670 | 670 | ||
| 671 | 简化示例： | 671 | 简化示例： |
| 672 | 672 | ||
| 673 | -```CPP | 673 | +```cpp |
| 674 | extern "C" __global__ __aicore__ void ${op_name}(GM_ADDR x, GM_ADDR y, GM_ADDR z, | 674 | extern "C" __global__ __aicore__ void ${op_name}(GM_ADDR x, GM_ADDR y, GM_ADDR z, |
| 675 | GM_ADDR workspace, GM_ADDR tiling); | 675 | GM_ADDR workspace, GM_ADDR tiling); |
| 676 | 676 | ||
| @@ -710,7 +710,7 @@ TEST_F(${OpName}KernelTest, test_case_basic) | |||
| 710 | - **手动构造**：适合字段少、逻辑简单。 | 710 | - **手动构造**：适合字段少、逻辑简单。 |
| 711 | - **调用Tiling函数自动生成**：适合字段多、依赖属性/shape复杂。可复用`tests/ut/common/tiling_context_faker.h`与`tiling_case_executor.h`。示例： | 711 | - **调用Tiling函数自动生成**：适合字段多、依赖属性/shape复杂。可复用`tests/ut/common/tiling_context_faker.h`与`tiling_case_executor.h`。示例： |
| 712 | 712 | ||
| 713 | -```CPP | 713 | +```cpp |
| 714 | gert::TilingContextPara para("OpName", | 714 | gert::TilingContextPara para("OpName", |
| 715 | {{{{2, 2, 2, 1}, {2, 2, 2, 1}}, ge::DT_FLOAT, ge::FORMAT_ND}}, | 715 | {{{{2, 2, 2, 1}, {2, 2, 2, 1}}, ge::DT_FLOAT, ge::FORMAT_ND}}, |
| 716 | {{{{2, 1, 2, 2}, {2, 1, 2, 2}}, ge::DT_FLOAT, ge::FORMAT_ND}}, | 716 | {{{{2, 1, 2, 2}, {2, 1, 2, 2}}, ge::DT_FLOAT, ge::FORMAT_ND}}, |
| @@ -797,7 +797,7 @@ export LD_LIBRARY_PATH=${ASCEND_HOME_PATH}/opp/vendors/${vendor_name}_cv/op_api/ | |||
| 797 | <td>op_kernel/{op_name}.cpp</td> | 797 | <td>op_kernel/{op_name}.cpp</td> |
| 798 | <td>将原有op_host/{op_name}.cpp中kernel实现的核函数实现迁移至cpp文件，同时： | 798 | <td>将原有op_host/{op_name}.cpp中kernel实现的核函数实现迁移至cpp文件，同时： |
| 799 | <br>. 新增REGISTER_TILING_DEFAULT调用注册Tiling结构体，使用GET_TILING_DATA_WITH_STRUCT获取TilingData | 799 | <br>. 新增REGISTER_TILING_DEFAULT调用注册Tiling结构体，使用GET_TILING_DATA_WITH_STRUCT获取TilingData |
| 800 | - <br>. 添加tiling模板，支持模板参数的传入，根据模板参数的分支判断，选择不同的kernel侧是实现 | 800 | + <br>. 添加tiling模板，支持模板参数的传入，根据模板参数的分支判断，选择不同的kernel侧的实现 |
| 801 | </td> | 801 | </td> |
| 802 | <td><a href="#op_kernel/{op_name}.cpp">op_kernel/{op_name}.cpp</a></td> | 802 | <td><a href="#op_kernel/{op_name}.cpp">op_kernel/{op_name}.cpp</a></td> |
| 803 | </tr> | 803 | </tr> |
| @@ -815,7 +815,7 @@ export LD_LIBRARY_PATH=${ASCEND_HOME_PATH}/opp/vendors/${vendor_name}_cv/op_api/ | |||
| 815 | 815 | ||
| 816 | 将原有${op_name}.cpp中算子信息库内容独立迁移至该文件，需要去掉SetInferShape和SetTiling内容。 | 816 | 将原有${op_name}.cpp中算子信息库内容独立迁移至该文件，需要去掉SetInferShape和SetTiling内容。 |
| 817 | 817 | ||
| 818 | -```CPP | 818 | +```cpp |
| 819 | // 原有${op_name}.cpp中算子信息库内容 | 819 | // 原有${op_name}.cpp中算子信息库内容 |
| 820 | namespace ops { | 820 | namespace ops { |
| 821 | class AddCustom : public OpDef { | 821 | class AddCustom : public OpDef { |
| @@ -871,7 +871,7 @@ OP_ADD(AddCustom); | |||
| 871 | 871 | ||
| 872 | 图模式场景需要适配该文件，将原有${op_name}.cpp中shape推导部分独立迁至该文件，调用接口IMPL_OP_INFERSHAPE完成InferShape注册。 | 872 | 图模式场景需要适配该文件，将原有${op_name}.cpp中shape推导部分独立迁至该文件，调用接口IMPL_OP_INFERSHAPE完成InferShape注册。 |
| 873 | 873 | ||
| 874 | -```CPP | 874 | +```cpp |
| 875 | // 原有${op_name}.cpp中的InferShape | 875 | // 原有${op_name}.cpp中的InferShape |
| 876 | namespace ge { | 876 | namespace ge { |
| 877 | static graphStatus InferShape(gert::InferShapeContext *context) | 877 | static graphStatus InferShape(gert::InferShapeContext *context) |
| @@ -905,7 +905,7 @@ IMPL_OP_INFERSHAPE(AddCustom).InferShape(InferShape); // 在该文件中完成 | |||
| 905 | 若是新增定义模板参数和模板参数组合，TilingFunc中需要同时配置模板参数tilingKey。 | 905 | 若是新增定义模板参数和模板参数组合，TilingFunc中需要同时配置模板参数tilingKey。 |
| 906 | 可参考[add_example_tiling.cpp](../../../examples/add_example/op_host/add_example_tiling.cpp)。 | 906 | 可参考[add_example_tiling.cpp](../../../examples/add_example/op_host/add_example_tiling.cpp)。 |
| 907 | 907 | ||
| 908 | -```CPP | 908 | +```cpp |
| 909 | // 原有${op_name}.cpp中TilingFunc | 909 | // 原有${op_name}.cpp中TilingFunc |
| 910 | namespace optiling { | 910 | namespace optiling { |
| 911 | const uint32_t BLOCK_DIM = 8; | 911 | const uint32_t BLOCK_DIM = 8; |
| @@ -963,7 +963,7 @@ IMPL_OP_OPTILING(AddCustom).Tiling(TilingFunc); // 在该文件中完成Tiling | |||
| 963 | </div> | 963 | </div> |
| 964 | 图模式场景需要适配该文件，将原有${op_name}.cpp中类型推导独立迁移至该文件后，调用接口IMPL_OP完成InferDataType注册。 | 964 | 图模式场景需要适配该文件，将原有${op_name}.cpp中类型推导独立迁移至该文件后，调用接口IMPL_OP完成InferDataType注册。 |
| 965 | 965 | ||
| 966 | -```CPP | 966 | +```cpp |
| 967 | // 原有${op_name}.cpp中InferDataType | 967 | // 原有${op_name}.cpp中InferDataType |
| 968 | namespace ge { | 968 | namespace ge { |
| 969 | static graphStatus InferDataType(gert::InferDataTypeContext *context) | 969 | static graphStatus InferDataType(gert::InferDataTypeContext *context) |
| @@ -990,7 +990,7 @@ IMPL_OP(AddCustom).InferDataType(InferDataType); // 在该文件中完成Infer | |||
| 990 | <p style="font-size:18px;"><b>op_kernel/{op_name}_tiling_data.h</b></p> | 990 | <p style="font-size:18px;"><b>op_kernel/{op_name}_tiling_data.h</b></p> |
| 991 | </div> | 991 | </div> |
| 992 | 992 | ||
| 993 | -```CPP | 993 | +```cpp |
| 994 | // 原有op_host/{op_name}_tiling.h中的宏定义TilingData结构体 | 994 | // 原有op_host/{op_name}_tiling.h中的宏定义TilingData结构体 |
| 995 | namespace optiling { | 995 | namespace optiling { |
| 996 | BEGIN_TILING_DATA_DEF(TilingData) | 996 | BEGIN_TILING_DATA_DEF(TilingData) |
| @@ -1016,7 +1016,7 @@ struct TilingData { | |||
| 1016 | <p style="font-size:18px;"><b>op_kernel/{op_name}.cpp</b></p> | 1016 | <p style="font-size:18px;"><b>op_kernel/{op_name}.cpp</b></p> |
| 1017 | </div> | 1017 | </div> |
| 1018 | 1018 | ||
| 1019 | -```CPP | 1019 | +```cpp |
| 1020 | // 原有op_kernel/{op_name}.cpp中的核函数实现 | 1020 | // 原有op_kernel/{op_name}.cpp中的核函数实现 |
| 1021 | template<int D_T_X, int D_T_Y, int D_T_Z, int TILE_NUM, int IS_SPLIT> | 1021 | template<int D_T_X, int D_T_Y, int D_T_Z, int TILE_NUM, int IS_SPLIT> |
| 1022 | __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling) | 1022 | __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling) |
| @@ -184,9 +184,9 @@ uint32_t AddExampleCpuKernel::Compute(CpuKernelContext& ctx) { | |||
| 184 | case DT_FLOAT: | 184 | case DT_FLOAT: |
| 185 | return AddCompute<float>(...); | 185 | return AddCompute<float>(...); |
| 186 | case DT_INT32: | 186 | case DT_INT32: |
| 187 | - return AddCompute<int32>(...); | 187 | + return AddCompute<int32_t>(...); |
| 188 | .... | 188 | .... |
| 189 | - default : return PARAM_INVALID; | 189 | + default : return kParamInvalid; |
| 190 | } | 190 | } |
| 191 | } | 191 | } |
| 192 | 192 | ||
| @@ -201,7 +201,7 @@ REGISTER_CPU_KERNEL(kAddExample, AddExampleCpuKernel); | |||
| 201 | 201 | ||
| 202 | ## 编译部署 | 202 | ## 编译部署 |
| 203 | 203 | ||
| 204 | -算子开发完成后，需对算子工程进行编译，生成自定义算子安装包\*\.run，具体操作如下： | 204 | +算子开发完成后，需对算子工程进行编译，生成自定义算子安装包*.run，具体操作如下： |
| 205 | 205 | ||
| 206 | 1. **准备工作。** | 206 | 1. **准备工作。** |
| 207 | 207 | ||
| @@ -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> |
| @@ -159,7 +159,7 @@ gather_v2算子根据合轴后的尾轴为单位进行gather，因此模板选 | |||
| 159 | 159 | ||
| 160 | SIMD实现采用传统的向量化编程模型，需显式管理UB缓冲区和流水队列： | 160 | SIMD实现采用传统的向量化编程模型，需显式管理UB缓冲区和流水队列： |
| 161 | 161 | ||
| 162 | -```Cpp | 162 | +```cpp |
| 163 | // SIMD: 使用队列机制管理数据缓冲 | 163 | // SIMD: 使用队列机制管理数据缓冲 |
| 164 | TQueBind<QuePosition::VECIN, QuePosition::VECOUT, BUFFER_NUM> inQueue_; | 164 | TQueBind<QuePosition::VECIN, QuePosition::VECOUT, BUFFER_NUM> inQueue_; |
| 165 | TBuf<QuePosition::VECCALC> indexBuf_; | 165 | TBuf<QuePosition::VECCALC> indexBuf_; |
| @@ -168,14 +168,14 @@ TBuf<QuePosition::VECCALC> indexBuf_; | |||
| 168 | for (int64_t j = 0; j < rows; j++) { | 168 | for (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 | } |
| 173 | inQueue_.EnQue<int8_t>(xLocal); // 入队等待输出 | 173 | inQueue_.EnQue<int8_t>(xLocal); // 入队等待输出 |
| 174 | ``` | 174 | ``` |
| 175 | 175 | ||
| 176 | SIMT采用线程级并行模型，每个线程独立处理元素： | 176 | SIMT采用线程级并行模型，每个线程独立处理元素： |
| 177 | 177 | ||
| 178 | -```Cpp | 178 | +```cpp |
| 179 | // SIMT: 使用线程级并行，无需显式buffer管理 | 179 | // SIMT: 使用线程级并行，无需显式buffer管理 |
| 180 | __simt_vf__ LAUNCH_BOUND(2048) void GatherSimt(...) { | 180 | __simt_vf__ LAUNCH_BOUND(2048) void GatherSimt(...) { |
| 181 | for (INDEX_SIZE_T index = Simt::GetThreadIdx(); | 181 | for (INDEX_SIZE_T index = Simt::GetThreadIdx(); |
| @@ -340,7 +340,7 @@ Ascend 950新架构引入UB2L1 & L0C2UB间的直连通路，实现矩阵计算 | |||
| 340 | 340 | ||
| 341 | 对于切K或多阶段融合场景，可将"L0C搬回GM再读回UB"改为"L0C直达UB累加/后处理"，降低GM往返带宽压力和时延。迁移时建议把中间结果归并、激活/量化前处理放到UB侧完成，并显式梳理MTE1/MTE2/MTE3与计算单元的事件同步顺序，确保跨单元流水连续，避免由于新增通路引入数据可见性或同步时序问题。关键接口定义可参考： | 341 | 对于切K或多阶段融合场景，可将"L0C搬回GM再读回UB"改为"L0C直达UB累加/后处理"，降低GM往返带宽压力和时延。迁移时建议把中间结果归并、激活/量化前处理放到UB侧完成，并显式梳理MTE1/MTE2/MTE3与计算单元的事件同步顺序，确保跨单元流水连续，避免由于新增通路引入数据可见性或同步时序问题。关键接口定义可参考： |
| 342 | 342 | ||
| 343 | -```Cpp | 343 | +```cpp |
| 344 | // 1. 新增: 搬入接口增加UB2L1的Nd2Nz搬入，支持Src&Dst都是LocalTensor的形式 | 344 | // 1. 新增: 搬入接口增加UB2L1的Nd2Nz搬入，支持Src&Dst都是LocalTensor的形式 |
| 345 | template <typename T> | 345 | template <typename T> |
| 346 | __aicore__ inline void DataCopy(const LocalTensor<T>& dst, const LocalTensor<T>& src, const Nd2NzParams& intriParams); | 346 | __aicore__ inline void DataCopy(const LocalTensor<T>& dst, const LocalTensor<T>& src, const Nd2NzParams& intriParams); |
| @@ -4,7 +4,7 @@ | |||
| 4 | 4 | ||
| 5 | 本文档介绍自定义算子的图模式适配方法，整体流程与算子开发指南（[AI Core算子开发指南](aicore_develop_guide.md)/[AI CPU算子开发指南](aicpu_develop_guide.md)）一致。值得注意的是，**不需要进行aclnn适配**，只需做如下交付件适配，即可实现图模式调用算子。 | 5 | 本文档介绍自定义算子的图模式适配方法，整体流程与算子开发指南（[AI Core算子开发指南](aicore_develop_guide.md)/[AI CPU算子开发指南](aicpu_develop_guide.md)）一致。值得注意的是，**不需要进行aclnn适配**，只需做如下交付件适配，即可实现图模式调用算子。 |
| 6 | 6 | ||
| 7 | -```Cpp | 7 | +```cpp |
| 8 | ${op_name} # 替换为实际算子名的小写下划线形式 | 8 | ${op_name} # 替换为实际算子名的小写下划线形式 |
| 9 | ├── op_host # Host侧实现 | 9 | ├── op_host # Host侧实现 |
| 10 | │ └── ${op_name}_infershape.cpp # InferShape实现，实现算子形状推导，在运行时推导输出shape | 10 | │ └── ${op_name}_infershape.cpp # InferShape实现，实现算子形状推导，在运行时推导输出shape |
| @@ -26,7 +26,7 @@ InferShape函数的作用是根据输入的shape推导输出的shape。 | |||
| 26 | 26 | ||
| 27 | 示例如下，`AddExample`算子完整代码请参考`examples/add_example_aicpu/op_host`下[add_example_infershape.cpp](../../../examples/add_example_aicpu/op_host/add_example_infershape.cpp)。 | 27 | 示例如下，`AddExample`算子完整代码请参考`examples/add_example_aicpu/op_host`下[add_example_infershape.cpp](../../../examples/add_example_aicpu/op_host/add_example_infershape.cpp)。 |
| 28 | 28 | ||
| 29 | -```C++ | 29 | +```cpp |
| 30 | // AddExample算子逻辑是两个数相加，因此输出shape与输入shape一致 | 30 | // AddExample算子逻辑是两个数相加，因此输出shape与输入shape一致 |
| 31 | static ge::graphStatus InferShapeAddExample(gert::InferShapeContext* context) | 31 | static ge::graphStatus InferShapeAddExample(gert::InferShapeContext* context) |
| 32 | { | 32 | { |
| @@ -56,7 +56,7 @@ InferDataType函数的作用是根据输入的DataType推导输出的DataType。 | |||
| 56 | 56 | ||
| 57 | 示例如下，`AddExample`算子完整代码请参考`examples/add_example_aicpu/op_graph`下[add_example_graph_infer.cpp](../../../examples/add_example_aicpu/op_graph/add_example_graph_infer.cpp)。 | 57 | 示例如下，`AddExample`算子完整代码请参考`examples/add_example_aicpu/op_graph`下[add_example_graph_infer.cpp](../../../examples/add_example_aicpu/op_graph/add_example_graph_infer.cpp)。 |
| 58 | 58 | ||
| 59 | -```C++ | 59 | +```cpp |
| 60 | // AddExample算子逻辑是两个数相加，因此输出dataType与输入dataType一致 | 60 | // AddExample算子逻辑是两个数相加，因此输出dataType与输入dataType一致 |
| 61 | static ge::graphStatus InferDataTypeAddExample(gert::InferDataTypeContext* context) | 61 | static ge::graphStatus InferDataTypeAddExample(gert::InferDataTypeContext* context) |
| 62 | { | 62 | { |
| @@ -39,7 +39,7 @@ bash build.sh --help | |||
| 39 | | -v | 可选 | 查看CMake编译配置信息。 | | 39 | | -v | 可选 | 查看CMake编译配置信息。 | |
| 40 | | -O${n} | 可选 | 指定编译优化级别，支持O0/O1/O2/O3（如：-O3），${n}为优化级别标识。 | | 40 | | -O${n} | 可选 | 指定编译优化级别，支持O0/O1/O2/O3（如：-O3），${n}为优化级别标识。 | |
| 41 | | -u | 可选 | 启用单元测试（UT）编译模式，编译所有UT目标。 | | 41 | | -u | 可选 | 启用单元测试（UT）编译模式，编译所有UT目标。 | |
| 42 | -| --help, -h | 可选 | 打印脚本使用帮助信息。 | | 42 | +| --help,-h | 可选 | 打印脚本使用帮助信息。 | |
| 43 | | --ops | 可选 | 指定待编译的算子，如：grid_sample,iou_v2，多个算子用英文逗号“,”分隔，不可与--ophost、--opapi、--opgraph同时使用。 | | 43 | | --ops | 可选 | 指定待编译的算子，如：grid_sample,iou_v2，多个算子用英文逗号“,”分隔，不可与--ophost、--opapi、--opgraph同时使用。 | |
| 44 | | --soc | 可选 | 指定NPU型号，每次编译只支持1个NPU型号。 | | 44 | | --soc | 可选 | 指定NPU型号，每次编译只支持1个NPU型号。 | |
| 45 | | --jit | 可选 | 静态图场景下，编译`cann-${soc_name}-ops-cv_${cann_version}_linux-${arch}.run`整包时不需要编译算子二进制文件（图的运行态会在线编译），可以配置该选项，以提升编译速度。 | | 45 | | --jit | 可选 | 静态图场景下，编译`cann-${soc_name}-ops-cv_${cann_version}_linux-${arch}.run`整包时不需要编译算子二进制文件（图的运行态会在线编译），可以配置该选项，以提升编译速度。 | |
| @@ -18,7 +18,7 @@ | |||
| 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-044 | [cmake-master-044.tar.gz](https://cann-3rd.obs.cn-north-4.myhuaweicloud.com/cmake/cmake-master-044.tar.gz) | | 19 | | cann-cmake | master-044 | [cmake-master-044.tar.gz](https://cann-3rd.obs.cn-north-4.myhuaweicloud.com/cmake/cmake-master-044.tar.gz) | |
| 20 | 20 | ||
| 21 | -若您的编译环境可以访问网络，请参考[联网编译](#联网编译)，编译脚本会自动联网下载第三方软件。否则，请参考[未联网编译](#未联网编译)手动下载第三方软件。 | 21 | +若编译环境可以访问网络，请参考[联网编译](#联网编译)，编译脚本会自动联网下载第三方软件。否则，请参考[未联网编译](#未联网编译)手动下载第三方软件。 |
| 22 | 22 | ||
| 23 | 准备好开源第三方软件后，可采用如下编译方式，请按需选择： | 23 | 准备好开源第三方软件后，可采用如下编译方式，请按需选择： |
| 24 | 24 | ||
| @@ -32,7 +32,7 @@ | |||
| 32 | 32 | ||
| 33 | - **ops-cv静态库**： | 33 | - **ops-cv静态库**： |
| 34 | 34 | ||
| 35 | - > 说明：若您需要**基于本项目进行二次发布**并且对**软件包大小有要求**时，建议采用静态库编译，该库可以链接您的应用开发程序，仅保留业务所需的算子，从而实现软件最小化部署。 | 35 | + > 说明：如果需要**基于本项目进行二次发布**并且对**软件包大小有要求**时，建议采用静态库编译，该库可以链接应用开发程序，仅保留业务所需的算子，从而实现软件最小化部署。 |
| 36 | 36 | ||
| 37 | 表示整个项目编译为一个静态库文件，包含libcann_cv_static.a和aclnn接口头文件。该包仅支持aclnn调用AI Core算子。 | 37 | 表示整个项目编译为一个静态库文件，包含libcann_cv_static.a和aclnn接口头文件。该包仅支持aclnn调用AI Core算子。 |
| 38 | 38 | ||
| @@ -110,7 +110,7 @@ | |||
| 110 | ```bash | 110 | ```bash |
| 111 | Self-extractable archive "cann-${soc_name}-ops-cv_${cann_version}_linux-${arch}.run" successfully created. | 111 | Self-extractable archive "cann-${soc_name}-ops-cv_${cann_version}_linux-${arch}.run" successfully created. |
| 112 | ``` | 112 | ``` |
| 113 | - | 113 | + |
| 114 | \$\{soc\_name\}表示NPU型号名称，即\$\{soc\_version\}删除“ascend”后剩余的内容。编译成功后，run包存放于build_out目录下。 | 114 | \$\{soc\_name\}表示NPU型号名称，即\$\{soc\_version\}删除“ascend”后剩余的内容。编译成功后，run包存放于build_out目录下。 |
| 115 | 115 | ||
| 116 | 2. **安装ops-cv包** | 116 | 2. **安装ops-cv包** |
| @@ -179,9 +179,9 @@ | |||
| 179 | 在联网环境中提前下载第三方软件，目前有如下方式，请按需选择： | 179 | 在联网环境中提前下载第三方软件，目前有如下方式，请按需选择： |
| 180 | 180 | ||
| 181 | - 方式1：根据[安装第三方依赖](#安装第三方依赖)提供的表格手动下载，若从其他地址下载，请确保版本号一致。 | 181 | - 方式1：根据[安装第三方依赖](#安装第三方依赖)提供的表格手动下载，若从其他地址下载，请确保版本号一致。 |
| 182 | - | 182 | + |
| 183 | - 方式2：通过[third_lib_download.py](../../../scripts/tools/third_lib_download.py)脚本一键下载，该脚本在本项目`scripts/tools/`目录，下载该脚本并执行如下命令： | 183 | - 方式2：通过[third_lib_download.py](../../../scripts/tools/third_lib_download.py)脚本一键下载，该脚本在本项目`scripts/tools/`目录，下载该脚本并执行如下命令： |
| 184 | - | 184 | + |
| 185 | ```bash | 185 | ```bash |
| 186 | python ${scripts_dir}/third_lib_download.py | 186 | python ${scripts_dir}/third_lib_download.py |
| 187 | ``` | 187 | ``` |
| @@ -193,36 +193,34 @@ | |||
| 193 | 将下载好的第三方软件上传至离线环境，可存放在`third_party`目录或自定义目录下。**推荐前者，其编译命令与联网编译场景下的命令一致。** | 193 | 将下载好的第三方软件上传至离线环境，可存放在`third_party`目录或自定义目录下。**推荐前者，其编译命令与联网编译场景下的命令一致。** |
| 194 | 194 | ||
| 195 | - **third\_party目录**（推荐） | 195 | - **third\_party目录**（推荐） |
| 196 | - | 196 | + |
| 197 | - 请在本项目根目录创建`third_party`目录（若有则无需创建），将第三方软件拷贝到该指定目录。此时编译命令与联网编译命令一致，具体参考[联网编译](#联网编译)。 | ||
| 198 | - | ||
| 199 | - **自定义目录** | 197 | - **自定义目录** |
| 200 | - | 198 | + |
| 201 | 在离线环境的任意位置新建`${cann_3rd_lib_path}`目录，将第三方软件拷贝到该目录，请确保该目录有权限访问。 | 199 | 在离线环境的任意位置新建`${cann_3rd_lib_path}`目录，将第三方软件拷贝到该目录，请确保该目录有权限访问。 |
| 202 | 200 | ||
| 203 | ```bash | 201 | ```bash |
| 204 | mkdir -p ${cann_3rd_lib_path} | 202 | mkdir -p ${cann_3rd_lib_path} |
| 205 | ``` | 203 | ``` |
| 206 | - | 204 | + |
| 207 | 此时编译命令需在联网编译命令基础上额外增加`--cann_3rd_lib_path=${cann_3rd_lib_path}`用于指定第三方软件所在路径。假设存放路径为`/path/cann_3rd_lib_path`，不同编译方式对应的命令如下： | 205 | 此时编译命令需在联网编译命令基础上额外增加`--cann_3rd_lib_path=${cann_3rd_lib_path}`用于指定第三方软件所在路径。假设存放路径为`/path/cann_3rd_lib_path`，不同编译方式对应的命令如下： |
| 208 | - | 206 | + |
| 209 | - 自定义算子包 | 207 | - 自定义算子包 |
| 210 | - | 208 | + |
| 211 | ```bash | 209 | ```bash |
| 212 | bash build.sh --pkg --soc=${soc_version} [--vendor_name=${vendor_name}] [--ops=${op_list}] --cann_3rd_lib_path=${cann_3rd_lib_path} | 210 | bash build.sh --pkg --soc=${soc_version} [--vendor_name=${vendor_name}] [--ops=${op_list}] --cann_3rd_lib_path=${cann_3rd_lib_path} |
| 213 | # 以GridSample算子编译为例 | 211 | # 以GridSample算子编译为例 |
| 214 | # bash build.sh --pkg --soc=ascend910b --ops=grid_sample -j16 --cann_3rd_lib_path=/path/cann_3rd_lib_path | 212 | # bash build.sh --pkg --soc=ascend910b --ops=grid_sample -j16 --cann_3rd_lib_path=/path/cann_3rd_lib_path |
| 215 | ``` | 213 | ``` |
| 216 | - | 214 | + |
| 217 | - ops-cv整包 | 215 | - ops-cv整包 |
| 218 | - | 216 | + |
| 219 | ```bash | 217 | ```bash |
| 220 | bash build.sh --pkg --soc=${soc_version} --cann_3rd_lib_path=${cann_3rd_lib_path} | 218 | bash build.sh --pkg --soc=${soc_version} --cann_3rd_lib_path=${cann_3rd_lib_path} |
| 221 | # bash build.sh --pkg --soc=ascend910b --cann_3rd_lib_path=/path/cann_3rd_lib_path | 219 | # bash build.sh --pkg --soc=ascend910b --cann_3rd_lib_path=/path/cann_3rd_lib_path |
| 222 | ``` | 220 | ``` |
| 223 | - | 221 | + |
| 224 | - ops-cv静态库 | 222 | - ops-cv静态库 |
| 225 | - | 223 | + |
| 226 | ```bash | 224 | ```bash |
| 227 | bash build.sh --pkg --static --soc=${soc_version} --cann_3rd_lib_path=${cann_3rd_lib_path} | 225 | bash build.sh --pkg --static --soc=${soc_version} --cann_3rd_lib_path=${cann_3rd_lib_path} |
| 228 | # bash build.sh --pkg --static --soc=ascend910b --cann_3rd_lib_path=/path/cann_3rd_lib_path | 226 | # bash build.sh --pkg --static --soc=ascend910b --cann_3rd_lib_path=/path/cann_3rd_lib_path |
| @@ -9,7 +9,7 @@ | |||
| 9 | 9 | ||
| 10 | 项目全量目录层级介绍如下： | 10 | 项目全量目录层级介绍如下： |
| 11 | 11 | ||
| 12 | -```Cpp | 12 | +```cpp |
| 13 | ├── cmake # 项目工程编译目录 | 13 | ├── cmake # 项目工程编译目录 |
| 14 | │ ├── aclnn_ops_cv.h.in # aclnn汇总头文件模板 | 14 | │ ├── aclnn_ops_cv.h.in # aclnn汇总头文件模板 |
| 15 | │ └── ... | 15 | │ └── ... |
| @@ -95,7 +95,7 @@ | |||
| 95 | ├── docs # 项目相关文档目录 | 95 | ├── docs # 项目相关文档目录 |
| 96 | ├── examples # 端到端算子开发和调用示例 | 96 | ├── examples # 端到端算子开发和调用示例 |
| 97 | │ ├── add_example # AI Core算子示例目录 | 97 | │ ├── add_example # AI Core算子示例目录 |
| 98 | -│ │ ├── CMakeLists.txt # 算子编译配置文件 | 98 | +│ │ ├── CMakeLists.txt # 算子编译配置文件 |
| 99 | │ │ ├── examples # 算子使用示例目录 | 99 | │ │ ├── examples # 算子使用示例目录 |
| 100 | │ │ ├── op_graph # 算子构图相关目录 | 100 | │ │ ├── op_graph # 算子构图相关目录 |
| 101 | │ │ ├── op_host # 算子信息库、Tiling、InferShape相关实现目录 | 101 | │ │ ├── op_host # 算子信息库、Tiling、InferShape相关实现目录 |
| @@ -117,7 +117,7 @@ | |||
| 117 | - \$\{soc\_name\}：表示NPU型号名称。 | 117 | - \$\{soc\_name\}：表示NPU型号名称。 |
| 118 | - \$\{install\_path\}：表示指定安装路径，ops包需与toolkit包安装在相同路径，root用户默认安装在`/usr/local/Ascend`目录。 | 118 | - \$\{install\_path\}：表示指定安装路径，ops包需与toolkit包安装在相同路径，root用户默认安装在`/usr/local/Ascend`目录。 |
| 119 | 119 | ||
| 120 | - - **场景2：体验已发布版本能力或基于已发布版本进行开发** | 120 | +- **场景2：体验已发布版本能力或基于已发布版本进行开发** |
| 121 | 121 | ||
| 122 | 请访问[CANN官网下载中心](https://www.hiascend.com/cann/download)，选择发布版本（仅支持CANN 8.5.0及后续版本），并根据产品型号和环境架构下载对应包，最后参考网页提供的命令完成安装。 | 122 | 请访问[CANN官网下载中心](https://www.hiascend.com/cann/download)，选择发布版本（仅支持CANN 8.5.0及后续版本），并根据产品型号和环境架构下载对应包，最后参考网页提供的命令完成安装。 |
| 123 | 123 | ||
| @@ -490,7 +490,7 @@ resultData[3] is: 4.000000 | |||
| 490 | } | 490 | } |
| 491 | ``` | 491 | ``` |
| 492 | 492 | ||
| 493 | -3. 创建CMakelist文件。 | 493 | +3. 创建CMakelists文件。 |
| 494 | 494 | ||
| 495 | 在`${test_geir_op_name}.cpp`同级目录下创建CMakeLists.txt文件，GE图引擎会根据配置好的环境变量自动加载已安装好的算子包（无论是自定义算子包或是标准内置算子包）库文件，无需特意区分。示例如下，仅供参考，请根据实际情况自行修改。 | 495 | 在`${test_geir_op_name}.cpp`同级目录下创建CMakeLists.txt文件，GE图引擎会根据配置好的环境变量自动加载已安装好的算子包（无论是自定义算子包或是标准内置算子包）库文件，无需特意区分。示例如下，仅供参考，请根据实际情况自行修改。 |
| 496 | 496 | ||
| @@ -24,7 +24,7 @@ | |||
| 24 | > - 配置说明：因CANN或NPU型号不同等原因，可能无法保证同一个算子多次运行结果一致。在相同条件下（平台、设备、版本号和其他随机性参数等），部分算子接口可通过`aclrtCtxSetSysParamOpt`（参见[《Runtime运行时API》](https://hiascend.com/document/redirect/CannCommunityRuntimeApi)）开启确定性算法，使多次运行结果一致。 | 24 | > - 配置说明：因CANN或NPU型号不同等原因，可能无法保证同一个算子多次运行结果一致。在相同条件下（平台、设备、版本号和其他随机性参数等），部分算子接口可通过`aclrtCtxSetSysParamOpt`（参见[《Runtime运行时API》](https://hiascend.com/document/redirect/CannCommunityRuntimeApi)）开启确定性算法，使多次运行结果一致。 |
| 25 | > - 性能说明：同一个算子采用确定性计算通常比非确定性慢，因此模型单次运行性能可能会下降。但在实验、调试调测等需要保证多次运行结果相同来定位问题的场景，确定性计算可以提升效率。 | 25 | > - 性能说明：同一个算子采用确定性计算通常比非确定性慢，因此模型单次运行性能可能会下降。但在实验、调试调测等需要保证多次运行结果相同来定位问题的场景，确定性计算可以提升效率。 |
| 26 | > - 线程说明：同一线程中只能设置一次确定性状态，多次设置以最后一次有效设置为准。有效设置是指设置确定性状态后，真正执行了一次算子任务下发。如果仅设置，没有算子下发，只能是确定性变量开启但未下发给算子，因此不执行算子。 | 26 | > - 线程说明：同一线程中只能设置一次确定性状态，多次设置以最后一次有效设置为准。有效设置是指设置确定性状态后，真正执行了一次算子任务下发。如果仅设置，没有算子下发，只能是确定性变量开启但未下发给算子，因此不执行算子。 |
| 27 | -> 解决方案：暂不推荐一个线程多次设置确定性。该问题在二进制开启和关闭情况下均存在，在后续版本中会解决该问题。 | 27 | +> 解决方案：暂不推荐一个线程多次设置确定性状态。该问题在二进制开启和关闭情况下均存在，在后续版本中会解决该问题。 |
| 28 | 28 | ||
| 29 | 算子接口列表如下： | 29 | 算子接口列表如下： |
| 30 | 30 | ||
| @@ -75,4 +75,4 @@ | |||
| 75 | | [aclnnRoiAlignV2Backward](../../objdetect/roi_align_grad/docs/aclnnRoiAlignV2Backward.md) |[aclnnRoiAlignV2](../../objdetect/roi_align/docs/aclnnRoiAlignV2.md)的反向传播。 |默认非确定性实现，支持配置开启|默认确定性实现| | 75 | | [aclnnRoiAlignV2Backward](../../objdetect/roi_align_grad/docs/aclnnRoiAlignV2Backward.md) |[aclnnRoiAlignV2](../../objdetect/roi_align/docs/aclnnRoiAlignV2.md)的反向传播。 |默认非确定性实现，支持配置开启|默认确定性实现| |
| 76 | | [aclnnRoiPoolingWithArgMax](../../objdetect/roi_pooling_with_arg_max/docs/aclnnRoiPoolingWithArgMax.md) | 对输入特征图按ROI（感兴趣区域）进行池化，在每个ROI内按空间划分为pooled_h × pooled_w个格子，对每个格子做最大池化，并输出池化结果及最大值在通道内的一维索引（argmax）。| - |默认确定性实现| | 76 | | [aclnnRoiPoolingWithArgMax](../../objdetect/roi_pooling_with_arg_max/docs/aclnnRoiPoolingWithArgMax.md) | 对输入特征图按ROI（感兴趣区域）进行池化，在每个ROI内按空间划分为pooled_h × pooled_w个格子，对每个格子做最大池化，并输出池化结果及最大值在通道内的一维索引（argmax）。| - |默认确定性实现| |
| 77 | | [aclnnRoiPoolingGradWithArgMax](../../objdetect/roi_pooling_grad_with_arg_max/docs/aclnnRoiPoolingGradWithArgMax.md) | [aclnnRoiPoolingWithArgMax](../../objdetect/roi_pooling_with_arg_max/docs/aclnnRoiPoolingWithArgMax.md)的反向传播。 | - |默认非确定性实现，不支持配置开启| | 77 | | [aclnnRoiPoolingGradWithArgMax](../../objdetect/roi_pooling_grad_with_arg_max/docs/aclnnRoiPoolingGradWithArgMax.md) | [aclnnRoiPoolingWithArgMax](../../objdetect/roi_pooling_with_arg_max/docs/aclnnRoiPoolingWithArgMax.md)的反向传播。 | - |默认非确定性实现，不支持配置开启| |
| 78 | -| [aclnnIm2colBackward](../../image/col2im/docs/aclnnIm2colBackward.md) | 从批处理输入张量中提取滑动局部块，将滑动局部块数组合并为一个大张量。 |默认确定性实现|默认确定性实现| | 78 | +| [aclnnIm2colBackward](../../image/col2im/docs/aclnnIm2colBackward.md) | 将梯度张量还原为原始输入张量的梯度分布，完成反向传播计算。 |默认确定性实现|默认确定性实现| |
| @@ -74,7 +74,7 @@ | |||
| 74 | <td>✗</td> | 74 | <td>✗</td> |
| 75 | <td>✓</td> | 75 | <td>✓</td> |
| 76 | <td>AI CPU</td> | 76 | <td>AI CPU</td> |
| 77 | - <td>从输入图像中提取多个裁剪区域,并将它们统一调整为指定大小，支持双线性插值和最近邻插值。</td> | 77 | + <td>从输入图像中提取多个裁剪区域，并将它们统一调整为指定大小，支持双线性插值和最近邻插值。</td> |
| 78 | </tr> | 78 | </tr> |
| 79 | <tr> | 79 | <tr> |
| 80 | <td>image</td> | 80 | <td>image</td> |
| @@ -334,7 +334,7 @@ | |||
| 334 | <td>✓</td> | 334 | <td>✓</td> |
| 335 | <td>✗</td> | 335 | <td>✗</td> |
| 336 | <td>AI Core</td> | 336 | <td>AI Core</td> |
| 337 | - <td>对由多个输入通道组成的输入信号应用2D双三次上采样。如果输入Tensorx的shape为(N,C,H,W)，则输出Tensorout的shape为(N,C,outputSize[0],outputSize[1])。</td> | 337 | + <td>对由多个输入通道组成的输入信号应用2D双三次上采样。如果输入Tensor x的shape为(N,C,H,W)，则输出Tensor out的shape为(N,C,outputSize[0],outputSize[1])。</td> |
| 338 | </tr> | 338 | </tr> |
| 339 | <tr> | 339 | <tr> |
| 340 | <td>image</td> | 340 | <td>image</td> |
| @@ -344,7 +344,7 @@ | |||
| 344 | <td>✓</td> | 344 | <td>✓</td> |
| 345 | <td>✗</td> | 345 | <td>✗</td> |
| 346 | <td>AI Core</td> | 346 | <td>AI Core</td> |
| 347 | - <td>对由多个输入通道组成的输入信号应用双三次抗锯齿算法进行上采样。如果输入Tensorx的shape为(N,C,H,W)，则输出Tensorout的shape为(N,C,outputSize[0],outputSize[1])。</td> | 347 | + <td>对由多个输入通道组成的输入信号应用双三次抗锯齿算法进行上采样。如果输入Tensor x的shape为(N,C,H,W)，则输出Tensor out的shape为(N,C,outputSize[0],outputSize[1])。</td> |
| 348 | </tr> | 348 | </tr> |
| 349 | <tr> | 349 | <tr> |
| 350 | <td>image</td> | 350 | <td>image</td> |
| @@ -434,7 +434,7 @@ | |||
| 434 | <td>✓</td> | 434 | <td>✓</td> |
| 435 | <td>✗</td> | 435 | <td>✗</td> |
| 436 | <td>AI Core</td> | 436 | <td>AI Core</td> |
| 437 | - <td>[UpsampleNearest](../../image/upsample_nearest/README.md)在exact_mode为true时的反向传播。</td> | 437 | + <td><a href="../../image/upsample_nearest/README.md">UpsampleNearest</a>在exact_mode为true时的反向传播。</td> |
| 438 | </tr> | 438 | </tr> |
| 439 | <tr> | 439 | <tr> |
| 440 | <td>image</td> | 440 | <td>image</td> |