已合并
modified md files(for readability improvement) #977
gitee-duhuiping创建于 6月12日
modified md files(for readability improvement) #977
已合并
gitee-duhuiping创建于 6月12日
65 个文件变更+487-489
@@ -13,7 +13,7 @@
13 13 
14## 类型标签14## 类型标签
15<!-- [x] 表示选中 -->15<!-- [x] 表示选中 -->
16-- [ ] 🐛 Bug 修复16+- [ ] 🐛 Bug修复
17- [ ] ✨ 新特性17- [ ] ✨ 新特性
18- [ ] ⚡ 性能优化18- [ ] ⚡ 性能优化
19- [ ] ♻️ 重构19- [ ] ♻️ 重构
@@ -6,9 +6,9 @@
6 6 
7发布日期:2025-12-307发布日期:2025-12-30
8 8 
9-ops-cv 算子首个 Beta 版本 v8.5.0-beta.1 现已发布。9+ops-cv算子首个Beta版本v8.5.0-beta.1现已发布。
10本版本引入了多项新增特性、问题修复及性能改进,目前仍处于测试阶段。10本版本引入了多项新增特性、问题修复及性能改进,目前仍处于测试阶段。
11-我们诚挚欢迎社区反馈,以进一步提升 ops-cv 的稳定性和功能完备性。11+我们诚挚欢迎社区反馈,以进一步提升ops-cv的稳定性和功能完备性。
12使用方式请参阅[官方文档](https://gitcode.com/cann/ops-cv/blob/master/README.md)。12使用方式请参阅[官方文档](https://gitcode.com/cann/ops-cv/blob/master/README.md)。
13 13 
14### 🔗 版本地址14### 🔗 版本地址
@@ -10,7 +10,7 @@
10 10 
11## 🚀概述11## 🚀概述
12 12 
13-ops-cv是[CANN](https://hiascend.com/software/cann) (Compute Architecture for Neural Networks)算子库中提供图像处理、目标检测等能力的高阶算子库,包括image类、objdetect类算子,覆盖常见的图像处理操作,子库在架构图中的位置如下。13+ops-cv是[CANN](https://hiascend.com/software/cann)(Compute Architecture for Neural Networks)算子库中提供图像处理、目标检测等能力的高阶算子库,包括image类、objdetect类算子,覆盖常见的图像处理操作,子库在架构图中的位置如下。
14 14 
15<img src="docs/zh/figures/architecture.png" alt="架构图" width="700px" height="326px">15<img src="docs/zh/figures/architecture.png" alt="架构图" width="700px" height="326px">
16 16 
@@ -27,7 +27,7 @@ ops-cv是[CANN](https://hiascend.com/software/cann) (Compute Architecture for
27 27 
28环境准备好后,下载与CANN版本配套的分支源码,命令如下,\$\{tag\_version\}替换为分支标签名。28环境准备好后,下载与CANN版本配套的分支源码,命令如下,\$\{tag\_version\}替换为分支标签名。
29 29 
30- > 说明:若环境中已存在配套分支源码,**可跳过本步骤**,例如CANNLab默认已提供最新商发版CANN对应的源码 30+ > 说明:若环境中已存在配套分支源码,**可跳过本步骤**,例如CANNLab默认已提供最新商发版CANN对应的源码。
31 31 
32```bash32```bash
33# 通用命令:git clone -b ${tag_version} https://gitcode.com/cann/ops-cv.git33# 通用命令:git clone -b ${tag_version} https://gitcode.com/cann/ops-cv.git
@@ -23,7 +23,7 @@
231. **准备工作**231. **准备工作**
24 24 
25 - 确定任务:如有文档问题可新建Issues,建议标签类别`[Documentation|文档反馈]`,并提供详细描述。基于已有Issues列表,确定待解决的文档issue。25 - 确定任务:如有文档问题可新建Issues,建议标签类别`[Documentation|文档反馈]`,并提供详细描述。基于已有Issues列表,确定待解决的文档issue。
26- - 认领任务:在对应的Issue下评论`/assign @yourself` ,表明您将处理它,避免重复劳动。26+ - 认领任务:在对应的Issue下评论`/assign @yourself`,表明您将处理它,避免重复劳动。
27 27 
282. **文档修改**282. **文档修改**
29 29 
@@ -44,7 +44,7 @@
44 简短说明(不超过50字符)44 简短说明(不超过50字符)
45 45
46 如有必要,在此处进行更详细描述。说明修改的原因和内容,而不是具体改了什么(代码本身会展示)。46 如有必要,在此处进行更详细描述。说明修改的原因和内容,而不是具体改了什么(代码本身会展示)。
47- 关联的 Issue: #12347+ 关联的Issue: #123
48 ```48 ```
49 49 
504. **发起Pull Request**504. **发起Pull Request**
@@ -97,7 +97,7 @@ bash build.sh --run_example add_example eager cust --vendor_name=custom
97 97 
98预期输出:打印算子`AddExample`的加法计算结果,表明算子已成功部署并正确执行。98预期输出:打印算子`AddExample`的加法计算结果,表明算子已成功部署并正确执行。
99 99 
100-```100+```bash
101add_example first input[0] is: 1.000000, second input[0] is: 1.000000, result[0] is: 2.000000101add_example first input[0] is: 1.000000, second input[0] is: 1.000000, result[0] is: 2.000000
102add_example first input[1] is: 1.000000, second input[1] is: 1.000000, result[1] is: 2.000000102add_example first input[1] is: 1.000000, second input[1] is: 1.000000, result[1] is: 2.000000
103add_example first input[2] is: 1.000000, second input[2] is: 1.000000, result[2] is: 2.000000103add_example first input[2] is: 1.000000, second input[2] is: 1.000000, result[2] is: 2.000000
@@ -158,7 +158,7 @@ __aicore__ inline void AddExample<T>::Compute(int64_t currentNum)
158 158 
1594. **成功标志**:输出结果变成乘法结果。1594. **成功标志**:输出结果变成乘法结果。
160 160 
161- ```161+ ```bash
162 add_example first input[0] is: 1.000000, second input[0] is: 1.000000, result[0] is: 1.000000162 add_example first input[0] is: 1.000000, second input[0] is: 1.000000, result[0] is: 1.000000
163 add_example first input[1] is: 1.000000, second input[1] is: 1.000000, result[1] is: 1.000000163 add_example first input[1] is: 1.000000, second input[1] is: 1.000000, result[1] is: 1.000000
164 add_example first input[2] is: 1.000000, second input[2] is: 1.000000, result[2] is: 1.000000164 add_example first input[2] is: 1.000000, second input[2] is: 1.000000, result[2] is: 1.000000
@@ -238,7 +238,7 @@ __aicore__ inline void AddExample<T>::Compute(int64_t currentNum)
238 238 
239```c++239```c++
240int main() {240int main() {
241- // ... 初始化代码 ...241+ // ... 初始化代码...
242 242
243 // === ① 修改selfX的输入 ===243 // === ① 修改selfX的输入 ===
244 // 修改前:shape = {32, 4, 4, 4}, 数值全为1244 // 修改前:shape = {32, 4, 4, 4}, 数值全为1
@@ -252,7 +252,7 @@ int main() {
252 }252 }
253 // === ② 参考selfX,同理修改selfY和out,并确保hostData长度与shape元素数一致 ===253 // === ② 参考selfX,同理修改selfY和out,并确保hostData长度与shape元素数一致 ===
254 254
255- // ... 后续执行代码 ...255+ // ... 后续执行代码...
256}256}
257```257```
258 258 
@@ -4,7 +4,7 @@
4 4 
5Docs目录结构说明如下:5Docs目录结构说明如下:
6 6 
7-```7+```text
8├── zh8├── zh
9│ ├── context # 公共文档,如术语、基础概念等9│ ├── context # 公共文档,如术语、基础概念等
10│ ├── debug # 算子调试指导文档10│ ├── debug # 算子调试指导文档
@@ -31,6 +31,7 @@ Docs目录结构说明如下:
31```31```
32 32 
33## 进阶教程33## 进阶教程
34+ 
34### 指南类文档35### 指南类文档
35 36 
36| 文档 | 说明 |37| 文档 | 说明 |
@@ -49,6 +50,7 @@ Docs目录结构说明如下:
49| [aclnn列表](zh/op_api_list.md) | 介绍项目包含的所有算子aclnn API清单。为方便用户在Host侧调用算子,提供了C语言API,即aclnn前缀的API。 |50| [aclnn列表](zh/op_api_list.md) | 介绍项目包含的所有算子aclnn API清单。为方便用户在Host侧调用算子,提供了C语言API,即aclnn前缀的API。 |
50 51 
51### 工具类文档52### 工具类文档
53+ 
52| 文档 | 说明 |54| 文档 | 说明 |
53| ----------------------- | ---------------------- |55| ----------------------- | ---------------------- |
54| [Simulator仿真工具](zh/debug/cann_sim.md) | 面向算子开发场景的SoC级仿真工具,用于分析运行在AI仿真器上AI任务在各阶段精度和性能数据。 |56| [Simulator仿真工具](zh/debug/cann_sim.md) | 面向算子开发场景的SoC级仿真工具,用于分析运行在AI仿真器上AI任务在各阶段精度和性能数据。 |
@@ -12,7 +12,7 @@
12>- 表格里表头表示待推导的输入Tensor数据类型,最左侧一列分别表示待推导的输入Scalar数据类型,表格中对应位置表示推导出的数据类型。12>- 表格里表头表示待推导的输入Tensor数据类型,最左侧一列分别表示待推导的输入Scalar数据类型,表格中对应位置表示推导出的数据类型。
13>- 表中叉号(×)表示这两种类型不能进行推导计算。13>- 表中叉号(×)表示这两种类型不能进行推导计算。
14 14 
15-**表 1** 数据类型推导关系表15+**表1** 数据类型推导关系表
16 16 
17| 数据类型 | f32 | f16 | f64 | bf16 | s8 | u8 | s16 | u16 | s32 | u32 | s64 | u64 | bool | c32 | c64 | c128 |17| 数据类型 | f32 | f16 | f64 | bf16 | s8 | u8 | s16 | u16 | s32 | u32 | s64 | u64 | bool | c32 | c64 | c128 |
18| :------: | :--: | :--: | :--: | :--: | :--: | :--: | :--: | :--: | :--: | :--: | :--: | :--: | :--: | :--: | :--: | :--: |18| :------: | :--: | :--: | :--: | :--: | :--: | :--: | :--: | :--: | :--: | :--: | :--: | :--: | :--: | :--: | :--: | :--: |
@@ -3,7 +3,7 @@
3调用aclnn API时,常见的接口返回码如[表1](#zh-cn_topic_0000001563019104_table8155243135018)所示。3调用aclnn API时,常见的接口返回码如[表1](#zh-cn_topic_0000001563019104_table8155243135018)所示。
4对于异常状态码值,可以通过aclGetRecentErrMsg接口(参见[《acl API(C)》](https://hiascend.com/document/redirect/CannCommunityCppApi))获取异常信息,您可以根据报错提示排查问题或者联系技术支持。4对于异常状态码值,可以通过aclGetRecentErrMsg接口(参见[《acl API(C)》](https://hiascend.com/document/redirect/CannCommunityCppApi))获取异常信息,您可以根据报错提示排查问题或者联系技术支持。
5 5 
6-**表 1** 返回状态码6+**表1** 返回状态码
7 7 
8<a name="zh-cn_topic_0000001563019104_table8155243135018"></a>8<a name="zh-cn_topic_0000001563019104_table8155243135018"></a>
9<table><thead align="left"><tr id="zh-cn_topic_0000001563019104_row111561243135019"><th class="cellrowborder" valign="top" width="30.543054305430545%" id="mcps1.2.4.1.1"><p id="zh-cn_topic_0000001563019104_p6676115185014"><a name="zh-cn_topic_0000001563019104_p6676115185014"></a><a name="zh-cn_topic_0000001563019104_p6676115185014"></a>状态码名称</p>9<table><thead align="left"><tr id="zh-cn_topic_0000001563019104_row111561243135019"><th class="cellrowborder" valign="top" width="30.543054305430545%" id="mcps1.2.4.1.1"><p id="zh-cn_topic_0000001563019104_p6676115185014"><a name="zh-cn_topic_0000001563019104_p6676115185014"></a><a name="zh-cn_topic_0000001563019104_p6676115185014"></a>状态码名称</p>
@@ -55,7 +55,7 @@
55 55 
56更多关于ACLNN_ERR_INNER_XXX类状态码的说明如[表2](#zh-cn_topic_0000001563019104_table3354143205413)所示。56更多关于ACLNN_ERR_INNER_XXX类状态码的说明如[表2](#zh-cn_topic_0000001563019104_table3354143205413)所示。
57 57 
58-**表 2** 异常状态码58+**表2** 异常状态码
59 59 
60<a name="zh-cn_topic_0000001563019104_table3354143205413"></a>60<a name="zh-cn_topic_0000001563019104_table3354143205413"></a>
61<table><thead align="left"><tr id="zh-cn_topic_0000001563019104_row15354124318546"><th class="cellrowborder" valign="top" width="30.183018301830185%" id="mcps1.2.4.1.1"><p id="zh-cn_topic_0000001563019104_p5354164365416"><a name="zh-cn_topic_0000001563019104_p5354164365416"></a><a name="zh-cn_topic_0000001563019104_p5354164365416"></a>状态码名称</p>61<table><thead align="left"><tr id="zh-cn_topic_0000001563019104_row15354124318546"><th class="cellrowborder" valign="top" width="30.183018301830185%" id="mcps1.2.4.1.1"><p id="zh-cn_topic_0000001563019104_p5354164365416"><a name="zh-cn_topic_0000001563019104_p5354164365416"></a><a name="zh-cn_topic_0000001563019104_p5354164365416"></a>状态码名称</p>
@@ -12,7 +12,7 @@ aclTensor支持的数据类型参见[数据类型](数据类型.md),其中部
12>- 表格里表头和最左侧一列分别表示待推导的两个输入数据类型,表格中对应位置表示推导出的数据类型。12>- 表格里表头和最左侧一列分别表示待推导的两个输入数据类型,表格中对应位置表示推导出的数据类型。
13>- 表中叉号(×)表示这两种类型不能进行推导计算。13>- 表中叉号(×)表示这两种类型不能进行推导计算。
14 14 
15-**表 1** 数据类型推导关系15+**表1** 数据类型推导关系
16 16 
17| 数据类型 | f32 | f16 | f64 | bf16 | s8 | u8 | s16 | u16 | s32 | u32 | s64 | u64 | bool | c32 | c64 | c128 |17| 数据类型 | f32 | f16 | f64 | bf16 | s8 | u8 | s16 | u16 | s32 | u32 | s64 | u64 | bool | c32 | c64 | c128 |
18| :------: | :--: | :--: | :--: | :--: | :--: | :--: | :--: | :--: | :--: | :--: | :--: | :--: | :--: | :--: | :--: | :--: |18| :------: | :--: | :--: | :--: | :--: | :--: | :--: | :--: | :--: | :--: | :--: | :--: | :--: | :--: | :--: | :--: | :--: |
@@ -3,7 +3,7 @@
3通过**aclCreateTensor**接口创建aclTensor时,支持的全量数据类型参见[《acl API(C)》](https://hiascend.com/document/redirect/CannCommunityCppApi)中“数据类型及其操作接口>aclDataType”。3通过**aclCreateTensor**接口创建aclTensor时,支持的全量数据类型参见[《acl API(C)》](https://hiascend.com/document/redirect/CannCommunityCppApi)中“数据类型及其操作接口>aclDataType”。
4两段式接口参数说明时,为了方便描述,支持的数据类型将采用如下简写形式。4两段式接口参数说明时,为了方便描述,支持的数据类型将采用如下简写形式。
5 5 
6-**表 1** 数据类型简写表6+**表1** 数据类型简写表
7 7 
8| 原始数据类型 | 简写形式(不区分大小写) |8| 原始数据类型 | 简写形式(不区分大小写) |
9| :---------------: | :----------------------: |9| :---------------: | :----------------------: |
@@ -18,7 +18,7 @@
18> a<sub>8,0</sub> , a<sub>8,1</sub> , a<sub>8,2</sub> , a<sub>8,3</sub> , a<sub>8,4</sub> , a<sub>8,5</sub> , a<sub>8,6</sub> , a<sub>8,7</sub> , a<sub>8,8</sub> , a<sub>8,9</sub> 18> a<sub>8,0</sub> , a<sub>8,1</sub> , a<sub>8,2</sub> , a<sub>8,3</sub> , a<sub>8,4</sub> , a<sub>8,5</sub> , a<sub>8,6</sub> , a<sub>8,7</sub> , a<sub>8,8</sub> , a<sub>8,9</sub>
19> a<sub>9,0</sub> , a<sub>9,1</sub> , a<sub>9,2</sub> , a<sub>9,3</sub> , a<sub>9,4</sub> , a<sub>9,5</sub> , a<sub>9,6</sub> , a<sub>9,7</sub> , a<sub>9,8</sub> , a<sub>9,9</sub> 19> a<sub>9,0</sub> , a<sub>9,1</sub> , a<sub>9,2</sub> , a<sub>9,3</sub> , a<sub>9,4</sub> , a<sub>9,5</sub> , a<sub>9,6</sub> , a<sub>9,7</sub> , a<sub>9,8</sub> , a<sub>9,9</sub>
20 20 
21-即该Tensor排布如上图的深色位置。这个完整的Tensor在内存排布上是不连续的,strides描述Tensor维度上相邻两个元素的间隔,如果在维度1上的stride为1, 该维度是连续的;如果在维度0上的stride为10,那么相邻的元素间隔10个元素,即非连续。offset表示这个Tensor的首元素相对addr的偏移。21+即该Tensor排布如上图的深色位置。这个完整的Tensor在内存排布上是不连续的,strides描述Tensor维度上相邻两个元素的间隔,如果在维度1上的stride为1,该维度是连续的;如果在维度0上的stride为10,那么相邻的元素间隔10个元素,即非连续。offset表示这个Tensor的首元素相对addr的偏移。
22 22 
23## 示例223## 示例2
24 24 
@@ -35,4 +35,4 @@
35> a<sub>8,0</sub> , a<sub>8,1</sub> , **a<sub>8,2</sub>** , a<sub>8,3</sub> , **a<sub>8,4</sub>** , a<sub>8,5</sub> , **a<sub>8,6</sub>** , a<sub>8,7</sub> , a<sub>8,8</sub> , a<sub>8,9</sub> 35> a<sub>8,0</sub> , a<sub>8,1</sub> , **a<sub>8,2</sub>** , a<sub>8,3</sub> , **a<sub>8,4</sub>** , a<sub>8,5</sub> , **a<sub>8,6</sub>** , a<sub>8,7</sub> , a<sub>8,8</sub> , a<sub>8,9</sub>
36> a<sub>9,0</sub> , a<sub>9,1</sub> , a<sub>9,2</sub> , a<sub>9,3</sub> , a<sub>9,4</sub> , a<sub>9,5</sub> , a<sub>9,6</sub> , a<sub>9,7</sub> , a<sub>9,8</sub> , a<sub>9,9</sub> 36> a<sub>9,0</sub> , a<sub>9,1</sub> , a<sub>9,2</sub> , a<sub>9,3</sub> , a<sub>9,4</sub> , a<sub>9,5</sub> , a<sub>9,6</sub> , a<sub>9,7</sub> , a<sub>9,8</sub> , a<sub>9,9</sub>
37 37 
38-即该Tensor排布如上图的深色位置。这个完整的Tensor在内存排布上是不连续的,strides描述Tensor维度上相邻两个元素的间隔,如果在维度1上的stride为2, 该维度上间隔1个元素;如果在维度0上的stride为20,那么相邻的元素间隔20个元素,即非连续。offset表示这个Tensor的首元素相对addr的偏移。38+即该Tensor排布如上图的深色位置。这个完整的Tensor在内存排布上是不连续的,strides描述Tensor维度上相邻两个元素的间隔,如果在维度1上的stride为2,该维度上间隔1个元素;如果在维度0上的stride为20,那么相邻的元素间隔20个元素,即非连续。offset表示这个Tensor的首元素相对addr的偏移。
@@ -4,7 +4,7 @@ CANN Simulator是一款面向算子开发场景的SoC级芯片仿真工具,用
4 4 
5# 主要功能5# 主要功能
6 6 
7-该工具与板上运行保持二进制兼容(同一 kernel可同时在仿真和AI处理器执行),主要用途如下:7+该工具与板上运行保持二进制兼容(同一kernel可同时在仿真和AI处理器执行),主要用途如下:
8 8 
9* 精度仿真:输出bit级精度结果,协助用户完成算子的精度验证。9* 精度仿真:输出bit级精度结果,协助用户完成算子的精度验证。
10* 性能仿真:输出指令流水图,协助用户定位算子性能瓶颈问题。10* 性能仿真:输出指令流水图,协助用户定位算子性能瓶颈问题。
@@ -43,7 +43,7 @@ bash build.sh --pkg --soc=Ascend950 --vendor_name=custom --ops=add_example
43./build_out/cann-ops-cv-${vendor_name}_linux-${arch}.run43./build_out/cann-ops-cv-${vendor_name}_linux-${arch}.run
44```44```
45 45 
46-* 参考[aclnn调用](../invocation/quick_op_invocation.md#aclnn调用)完成test_aclnn_add_example.cpp的编译,编出可执行文件test_aclnn_add_example46+* 参考[aclnn调用](../invocation/quick_op_invocation.md#aclnn-api)完成test_aclnn_add_example.cpp的编译,编出可执行文件test_aclnn_add_example
47 47 
48## 执行仿真命令48## 执行仿真命令
49 49 
@@ -51,8 +51,7 @@ bash build.sh --pkg --soc=Ascend950 --vendor_name=custom --ops=add_example
51cannsim record ./test_aclnn_add_example -s Ascend950 --gen-report51cannsim record ./test_aclnn_add_example -s Ascend950 --gen-report
52```52```
53 53 
54-仿真工具执行日志文件在examples/add_example/examples/build/bin/cannsim_*目录,执行日志文件为 cannsim.log。54+仿真工具执行日志文件在examples/add_example/examples/build/bin/cannsim_*目录,执行日志文件为cannsim.log。
55- 
56 55 
57从仿真工具日志文件可以看到示例中的打印信息:56从仿真工具日志文件可以看到示例中的打印信息:
58 57 
@@ -88,16 +87,16 @@ cannsim record [options] user_app
88 87 
89## 参数说明88## 参数说明
90 89 
91-表1 仿真执行参数说明90+**表1仿真执行参数说明**
92 91 
93|参数|可选/必选|说明|92|参数|可选/必选|说明|
94| --- | --- | --- |93| --- | --- | --- |
95-|-s --soc-version | 必选 | 指定模拟目标芯片版本(如:Ascend950)。|94+|-s或--soc-version | 必选 | 指定模拟目标芯片版本(如:Ascend950)。|
96-|-o --output | 可选 | 生成文件所在路径,可配置为绝对路径或者相对路径,并且执行工具的用户需要具有读写权限。如果未指定路径,则默认在当前目录下保存数据。|95+|-o或--output | 可选 | 生成文件所在路径,可配置为绝对路径或者相对路径,并且执行工具的用户需要具有读写权限。如果未指定路径,则默认在当前目录下保存数据。|
97-|-g --gen-report | 可选 | 启用仿真完成后是否进行自动解析,并生成分析报告。默认不自动解析。|96+|-g或--gen-report | 可选 | 启用仿真完成后是否进行自动解析,并生成分析报告。默认不自动解析。|
98-|-u --user-option | 可选 | 用户自定义算子参数,以命令行选项形式传递给算子程序。|97+|-u或--user-option | 可选 | 用户自定义算子参数,以命令行选项形式传递给算子程序。|
99-|-n --core-id | 可选 | 仿真期间启用日志的AI Core,格式同 report -n:'all'、'0-2,12-14'、'5'。默认全开;配合 -g 且未指定时回退到 core 0。|98+|-n或--core-id | 可选 | 仿真期间启用日志的AI Core,格式同report -n:'all'、'0-2,12-14'、'5'。默认全开;配合-g且未指定时回退到core 0。|
100-|user_app|必选|待运行的算子程序或命令(如 ./app, python train.py, bash run.sh)。|99+|user_app|必选|待运行的算子程序或命令(如./app, python train.py, bash run.sh)。|
101 100 
102## 使用示例101## 使用示例
103 102 
@@ -105,7 +104,7 @@ cannsim record [options] user_app
1052. 执行仿真命令,可参考以下使用示例1042. 执行仿真命令,可参考以下使用示例
106 105 
107 ```bash106 ```bash
108- 方式一: 启用仿真,并将输出保存至 ./output 目录,/path/to/app 为算子程序107+ 方式一:启用仿真,并将输出保存至./output目录,/path/to/app为算子程序
109 $ cannsim record /path/to/app -o ./output -s Ascend950108 $ cannsim record /path/to/app -o ./output -s Ascend950
110 109 
111 方式二:启用仿真并生成报告,用于后续性能分析110 方式二:启用仿真并生成报告,用于后续性能分析
@@ -114,10 +113,10 @@ cannsim record [options] user_app
114 113 
1153. 命令完成后,会在默认路径或指定的“output”目录下生成以“cannsim_{timestamp}_${user_app}”命名的文件夹,结构示例如下:1143. 命令完成后,会在默认路径或指定的“output”目录下生成以“cannsim_{timestamp}_${user_app}”命名的文件夹,结构示例如下:
116 115 
117-```bash116+ ```bash
118-├─cannsim_{timestamp}_${user_app}117+ ├─cannsim_{timestamp}_${user_app}
119-├── cannsim.log118+ ├── cannsim.log
120-```119+ ```
121 120 
1224. 用户可以获取算子执行结果,并进行精度的对比,结果展示在cannsim.log,示例如下1214. 用户可以获取算子执行结果,并进行精度的对比,结果展示在cannsim.log,示例如下
123 122 
@@ -144,14 +143,14 @@ cannsim report [options]
144 143 
145## 参数说明144## 参数说明
146 145 
147-1 仿真结果解析参数说明146+**2仿真结果解析参数说明**
148 147 
149|参数 | 可选/必选 | 说明|148|参数 | 可选/必选 | 说明|
150| --- | --- | --- |149| --- | --- | --- |
151-|-e --export | 必选 | 仿真执行结果目录,指定到cannsim_{timestamp}_${user_app}层,可配置为绝对路径或者相对路径,且执行用户需具有读写权限。|150+|-e或--export | 必选 | 仿真执行结果目录,指定到cannsim_{timestamp}_${user_app}层,可配置为绝对路径或者相对路径,且执行用户需具有读写权限。|
152-|-o --output | 可选 | 指令流水图输出目录,可配置为绝对路径或者相对路径,且执行用户需具有读写权限。若未指定路径,默认与 export 目录相同。|151+|-o或--output | 可选 | 指令流水图输出目录,可配置为绝对路径或者相对路径,且执行用户需具有读写权限。若未指定路径,默认与export目录相同。|
153-|-n --core-id | 可选 | 指定生成指令流水的核ID,支持格式:'all'、'0-2,12-14'、'5'。不指定默认生成0核的指令流水。|152+|-n或--core-id | 可选 | 指定生成指令流水的核ID,支持格式:'all'、'0-2,12-14'、'5'。不指定默认生成0核的指令流水。|
154-|-f --object-file | 可选 | 设备对象文件路径,用于辅助生成报告。|153+|-f或--object-file | 可选 | 设备对象文件路径,用于辅助生成报告。|
155 154 
156## 使用示例155## 使用示例
157 156 
@@ -178,7 +177,7 @@ cannsim report [options]
178 在Chrome浏览器中输入“chrome://tracing”地址,并将生成的指令流水图文件(trace.json)拖到空白处打开,键盘上输入快捷键(W:放大,S:缩小,A:左移,D:右移)可进行查看。177 在Chrome浏览器中输入“chrome://tracing”地址,并将生成的指令流水图文件(trace.json)拖到空白处打开,键盘上输入快捷键(W:放大,S:缩小,A:左移,D:右移)可进行查看。
179 ![指令流水图](../figures/指令流水图.png)178 ![指令流水图](../figures/指令流水图.png)
180 179 
181-2 关键字段说明180+ **3关键字段说明**
182 181 
183 |字段名|字段含义|182 |字段名|字段含义|
184 | --- | --- |183 | --- | --- |
@@ -206,13 +205,13 @@ cannsim report [options]
206cannsim --help205cannsim --help
207```206```
208 207 
209-查询工具 record 子命令的帮助信息:208+查询工具record子命令的帮助信息:
210 209 
211```bash210```bash
212cannsim record --help211cannsim record --help
213```212```
214 213
215-查询工具 report 子命令的帮助信息:214+查询工具report子命令的帮助信息:
216 215 
217 ```bash216 ```bash
218 cannsim report --help 217 cannsim report --help
@@ -8,13 +8,13 @@
8 8 
9* **plog获取**9* **plog获取**
10 10 
11- 程序执行结束后,默认可在"$HOME/ascend/log"下查看,host日志文件存储路径如下:11+ 程序执行结束后,默认可在$HOME/ascend/log下查看,host日志文件存储路径如下:
12 12 
13 ```sh13 ```sh
14 $HOME/ascend/log/debug/plog/plog-pid_*.log14 $HOME/ascend/log/debug/plog/plog-pid_*.log
15 ```15 ```
16 16 
17- 开启环境变量ASCEND_SLOG_PRINT_TO_STDOUT可以将log日志直接打屏显示(1:开启打屏,0:关闭打屏),配置示例如下:17+ 开启环境变量ASCEND_SLOG_PRINT_TO_STDOUT可以将log日志直接打屏显示1开启打屏,0:关闭打屏,配置示例如下:
18 18 
19 ```sh19 ```sh
20 export ASCEND_SLOG_PRINT_TO_STDOUT=120 export ASCEND_SLOG_PRINT_TO_STDOUT=1
@@ -23,7 +23,7 @@
23 日志相关介绍参见[《日志参考》](https://hiascend.com/document/redirect/CannCommunitylogref),环境变量介绍参见[《环境变量参考》](https://hiascend.com/document/redirect/CannCommunityEnvRef)。23 日志相关介绍参见[《日志参考》](https://hiascend.com/document/redirect/CannCommunitylogref),环境变量介绍参见[《环境变量参考》](https://hiascend.com/document/redirect/CannCommunityEnvRef)。
24 24 
25* **aclnn异常错误信息获取**25* **aclnn异常错误信息获取**
26- 26+ 
27 通过aclGetRecentErrMsg接口(参见[《acl API(C)》](https://hiascend.com/document/redirect/CannCommunityCppApi))获取aclnn接口调用过程中的异常信息,使用方法如下:27 通过aclGetRecentErrMsg接口(参见[《acl API(C)》](https://hiascend.com/document/redirect/CannCommunityCppApi))获取aclnn接口调用过程中的异常信息,使用方法如下:
28 28 
29 ```sh29 ```sh
@@ -161,17 +161,17 @@
161 161 
162 执行仿真命令,生成仿真数据162 执行仿真命令,生成仿真数据
163 163 
164- ```164+ ```bash
165 cannsim record ./test_aclnn_add_example -s Ascend950 --gen-report165 cannsim record ./test_aclnn_add_example -s Ascend950 --gen-report
166 ```166 ```
167 167 
168 仿真结果在本项目`examples/add_example/examples/build/bin/cannsim_*`目录,流水相关文件为:168 仿真结果在本项目`examples/add_example/examples/build/bin/cannsim_*`目录,流水相关文件为:
169 169 
170- ```170+ ```bash
171 trace_core0.json171 trace_core0.json
172 ```172 ```
173 173 
174- 在Chrome浏览器中输入“chrome://tracing”地址,并将生成的指令流水图文件(trace_core0.json)拖到空白处打开,具体参数介绍参考CANN Simulator中[“仿真结果解析”](./cann_sim.md#仿真结果解析)章节。174+ 在Chrome浏览器中输入“chrome://tracing”地址,并将生成的指令流水图文件(trace_core0.json)拖到空白处打开,具体参数介绍参考CANN Simulator中[“仿真结果解析说明”](./cann_sim.md#仿真结果解析说明)章节。
175 175 
176 * **针对Atlas A2/A3系列产品,可使用[msProf](https://www.hiascend.com/document/redirect/CannCommunityToolMsprof)工具,执行仿真命令,生成仿真数据**176 * **针对Atlas A2/A3系列产品,可使用[msProf](https://www.hiascend.com/document/redirect/CannCommunityToolMsprof)工具,执行仿真命令,生成仿真数据**
177 177 
@@ -144,10 +144,10 @@ flowchart TB
144| 步骤 | 主要文件 | 关键API | 输出产物 |144| 步骤 | 主要文件 | 关键API | 输出产物 |
145|------|----------|---------|----------|145|------|----------|---------|----------|
146| 1. 算子设计 | - | - | 算子规格文档 |146| 1. 算子设计 | - | - | 算子规格文档 |
147-| 2. 算子定义 | `op_host/{op}_def.cpp` | `OpDef`, `Input()`, `Output()`, `OP_ADD()` | 算子原型注册 |147+| 2. 算子定义 | `op_host/{op}_def.cpp` | `OpDef``Input()``Output()``OP_ADD()` | 算子原型注册 |
148-| 3. Tiling实现 | `op_host/{op}_tiling.cpp` | `TilingFunc`, `GetTilingData<>()`, `SetBlockDim()` | Tiling参数 |148+| 3. Tiling实现 | `op_host/{op}_tiling.cpp` | `TilingFunc``GetTilingData<>()` `SetBlockDim()` | Tiling参数 |
149-| 4. Kernel实现 | `op_kernel/{op}.cpp` | `__global__ __aicore__`, `GET_TILING_DATA` | 核函数二进制 |149+| 4. Kernel实现 | `op_kernel/{op}.cpp` | `__global__ __aicore__``GET_TILING_DATA` | 核函数二进制 |
150-| 5. 图模式适配 | `op_graph/{op}_proto.h` | `REG_OP()`, `INPUT()`, `OUTPUT()` | IR定义 |150+| 5. 图模式适配 | `op_graph/{op}_proto.h` | `REG_OP()``INPUT()``OUTPUT()` | IR定义 |
151| 6. aclnn适配 | `CMakeLists.txt` | `ACLNNTYPE aclnn` | API动态库 |151| 6. aclnn适配 | `CMakeLists.txt` | `ACLNNTYPE aclnn` | API动态库 |
152 152 
153---153---
@@ -157,7 +157,7 @@ flowchart TB
157### 算子工程目录结构157### 算子工程目录结构
158 158 
159```tex159```tex
160-{op_name}/ # 算子根目录(如 add/)160+{op_name}/ # 算子根目录(如add/)
161├── CMakeLists.txt # 构建配置(必需)161├── CMakeLists.txt # 构建配置(必需)
162├── README.md # 算子说明162├── README.md # 算子说明
163163
@@ -350,7 +350,7 @@ StatelessRandom → stateless_random
350> | **IR定义** | 只需**枚举**各输入输出支持的类型 | 无需标明输入输出的对应关系 |350> | **IR定义** | 只需**枚举**各输入输出支持的类型 | 无需标明输入输出的对应关系 |
351>351>
352> ```c++352> ```c++
353-> // 算子原型注册 (op_host/{op}_def.cpp) - 需要排列组合353+> // 算子原型注册(op_host/{op}_def.cpp) - 需要排列组合
354> this->Input("x")354> this->Input("x")
355> .DataType({ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_FLOAT16}) // 第1,2,3个分别对应输出的第1,2,3个355> .DataType({ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_FLOAT16}) // 第1,2,3个分别对应输出的第1,2,3个
356> .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});356> .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
@@ -358,7 +358,7 @@ StatelessRandom → stateless_random
358> .DataType({ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_FLOAT}) // float32->float32, float16->float16, float16->float32358> .DataType({ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_FLOAT}) // float32->float32, float16->float16, float16->float32
359> .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});359> .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
360>360>
361-> // IR定义 (op_graph/{op}_proto.h) - 只需枚举361+> // IR定义(op_graph/{op}_proto.h) - 只需枚举
362> .INPUT(x, TensorType({DT_FLOAT, DT_FLOAT16}))362> .INPUT(x, TensorType({DT_FLOAT, DT_FLOAT16}))
363> .OUTPUT(y, TensorType({DT_FLOAT, DT_FLOAT16}))363> .OUTPUT(y, TensorType({DT_FLOAT, DT_FLOAT16}))
364> ```364> ```
@@ -372,7 +372,7 @@ this->Input("x")
372 .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});372 .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
373```373```
374 374 
375-**表1 输入输出参数说明**375+**表1输入输出参数说明**
376 376 
377| 原型定义 | 参数 | 具体描述 |377| 原型定义 | 参数 | 具体描述 |
378| ------------ | --------- | ------------------------------------------------------------ |378| ------------ | --------- | ------------------------------------------------------------ |
@@ -461,7 +461,7 @@ this->Attr("isKeepDim")
461| 原型定义 | 注册方式 | 具体描述 |461| 原型定义 | 注册方式 | 具体描述 |
462| -------- | ----------------- | ------------------------------------------------------------ |462| -------- | ----------------- | ------------------------------------------------------------ |
463| Attr | AttrType | 设置算子属性类型,取值为:OPTIONAL(可选)、REQUIRED(必选)。 |463| Attr | AttrType | 设置算子属性类型,取值为:OPTIONAL(可选)、REQUIRED(必选)。 |
464-| | Bool/Float/Int... | 设置算子属性数据类型为Bool/Float/Int...。 |464+| | Bool/Float/Int... | 设置算子属性数据类型为Bool/Float/Int...。 |
465 465 
466### AI处理器上相关实现信息466### AI处理器上相关实现信息
467 467 
@@ -588,9 +588,9 @@ Tiling实现完成后,获取到的Tiling切分算法相关参数,会传递
588 588 
589```mermaid589```mermaid
590flowchart TD590flowchart TD
591- A["算子的输入、输出以及属性信息<br/>通过 TilingContext* context 获取"]591+ A["算子的输入、输出以及属性信息<br/>通过TilingContext* context获取"]
592 B["Tiling"]592 B["Tiling"]
593- C["TilingData、BlockDim、TilingKey、workspace size...<br/>设置到 TilingContext* context 中"]593+ C["TilingData、BlockDim、TilingKey、workspace size...<br/>设置到TilingContext* context中"]
594 594 
595 A --> B595 A --> B
596 B --> C596 B --> C
@@ -710,7 +710,7 @@ workspaceSizes[0] = sysWorkspaceSize + usrWorkspaceSize;
710 710 
711### Tiling结构体定义方式对比711### Tiling结构体定义方式对比
712 712 
713-| 特性 | 标准C++语法 | 宏定义方式(BEGIN_TILING_DATA_DEF) |713+| 特性 | 标准C++语法 | 宏定义方式BEGIN_TILING_DATA_DEF |
714|------|-------------|-----------------------------------|714|------|-------------|-----------------------------------|
715| 支持bool类型 | ✅ | ❌ |715| 支持bool类型 | ✅ | ❌ |
716| 支持数组/列表初始化 | ✅ | ❌ |716| 支持数组/列表初始化 | ✅ | ❌ |
@@ -1266,34 +1266,34 @@ ExtendCfgInfo("opFile.value", "{op_name_snake}_apt");
1266 1266 
1267## 常见问题1267## 常见问题
1268 1268 
1269-### Q1: Tiling结构体定义放在哪个目录?1269+### Q1Tiling结构体定义放在哪个目录?
1270 1270 
1271-**A:** Tiling结构体头文件应放置在`op_kernel/`目录下,因为只有该目录下的文件会被打包进算子包。如果放在其他目录,可能导致在线编译失败。1271+**A** Tiling结构体头文件应放置在`op_kernel/`目录下,因为只有该目录下的文件会被打包进算子包。如果放在其他目录,可能导致在线编译失败。
1272 1272 
1273-### Q2: 核函数参数顺序可以调整吗?1273+### Q2核函数参数顺序可以调整吗?
1274 1274 
1275-**A:** 不可以。核函数参数必须按照 **输入 → 输出 → workspace → tiling** 的固定顺序排布。1275+**A** 不可以。核函数参数必须按照 **输入 → 输出 → workspace → tiling** 的固定顺序排布。
1276 1276 
1277-### Q3: 如何选择BlockDim的值?1277+### Q3如何选择BlockDim的值?
1278 1278 
1279-**A:**1279+**A**
1280 1280 
1281- 耦合模式:使用`GetCoreNumAiv()``GetCoreNumAic()`获取核数1281- 耦合模式:使用`GetCoreNumAiv()``GetCoreNumAic()`获取核数
1282- 分离模式Vector算子:设置为Vector核数1282- 分离模式Vector算子:设置为Vector核数
1283- 分离模式Cube算子:设置为Cube核数1283- 分离模式Cube算子:设置为Cube核数
1284- 一般建议设置为物理核数以充分利用硬件资源1284- 一般建议设置为物理核数以充分利用硬件资源
1285 1285 
1286-### Q4: TilingData获取后需要初始化吗?1286+### Q4TilingData获取后需要初始化吗?
1287 1287 
1288-**A:** 是的。通过`GetTilingData<T>()`获取的Tiling结构体不包含初值,必须显式赋值所有需要使用的成员变量。1288+**A** 是的。通过`GetTilingData<T>()`获取的Tiling结构体不包含初值,必须显式赋值所有需要使用的成员变量。
1289 1289 
1290-### Q5: 算子输入输出同名怎么处理?1290+### Q5算子输入输出同名怎么处理?
1291 1291 
1292-**A:** 输出参数会增加`ref`后缀。例如输入`x`,输出也是`x`时,核函数参数为`x`和`x_ref`。1292+**A** 输出参数会增加`ref`后缀。例如输入`x`,输出也是`x`时,核函数参数为`x`和`x_ref`。
1293 1293 
1294-### Q6: 不同芯片架构的代码如何隔离?1294+### Q6不同芯片架构的代码如何隔离?
1295 1295 
1296-**A:**1296+**A**
1297 1297 
12981.`op_host/``op_kernel/`下分别创建`arch32/``arch35/`目录12981.`op_host/``op_kernel/`下分别创建`arch32/``arch35/`目录
12992.`{op}_def.cpp`中通过`ExtendCfgInfo("opFile.value", ...)`配置不同入口12992.`{op}_def.cpp`中通过`ExtendCfgInfo("opFile.value", ...)`配置不同入口
@@ -6,7 +6,7 @@
6 6 
7算子根据运行的硬件单元不同,可分为AI Core算子和AI CPU算子(少数)。AI Core算子使用Ascend C语言开发,运行在AI Core硬件单元;AI CPU算子使用C++语言开发,运行在AI CPU硬件单元。7算子根据运行的硬件单元不同,可分为AI Core算子和AI CPU算子(少数)。AI Core算子使用Ascend C语言开发,运行在AI Core硬件单元;AI CPU算子使用C++语言开发,运行在AI CPU硬件单元。
8 8 
9-本文旨在介绍如何基于标准工程开发AI Core算子,如果您想贡献AI CPU算子,请参考[AI CPU算子开发指南](./aicpu_develop_guide.md)。9+本文旨在介绍如何基于标准工程开发AI Core算子,如果您想贡献AI CPU算子,请参考[AI CPU算子开发指南](./aicpu_develop_guide.md)。
10 10 
11算子开发前,请先了解如下信息:11算子开发前,请先了解如下信息:
12 12 
@@ -56,7 +56,7 @@ Create the initial directory for ${op_name} under ${op_class} success
56 56 
57创建完成后,目录结构如下所示:57创建完成后,目录结构如下所示:
58 58 
59-```59+```tex
60${op_name} # 替换为实际算子名的小写下划线形式60${op_name} # 替换为实际算子名的小写下划线形式
61├── examples # 算子调用示例61├── examples # 算子调用示例
62│ ├── test_aclnn_${op_name}.cpp # 算子aclnn调用示例62│ ├── test_aclnn_${op_name}.cpp # 算子aclnn调用示例
@@ -77,7 +77,7 @@ ${op_name} # 替换为实际算子名的小写下
77 77 
78`${op_class}`为全新算子分类需额外在`CMakeLists.txt`中添加`add_subdirectory(${op_class})`,否则无法正常编译。78`${op_class}`为全新算子分类需额外在`CMakeLists.txt`中添加`add_subdirectory(${op_class})`,否则无法正常编译。
79 79 
80-```80+```bash
81if(ENABLE_EXPERIMENTAL)81if(ENABLE_EXPERIMENTAL)
82 # genop新增experimental算子分类82 # genop新增experimental算子分类
83 # add_subdirectory(${op_class})83 # add_subdirectory(${op_class})
@@ -93,8 +93,6 @@ endif()
93 93 
94算子定义需要完成两个交付件:`README.md````${op_name}_def.cpp```94算子定义需要完成两个交付件:`README.md````${op_name}_def.cpp```
95 95 
96-> 💡 **进阶内容**:关于算子原型定义的详细说明,包括输入/输出/属性定义、AI处理器配置、多硬件平台差异化注册等,请参考[《AI Core算子开发进阶指南 - 算子原型定义》](./aicore_develop_advanced_guide.md#算子原型定义)。
97- 
98**交付件1:README.md**96**交付件1:README.md**
99 97 
100开发算子前需要先确定目标算子的功能和计算逻辑。98开发算子前需要先确定目标算子的功能和计算逻辑。
@@ -109,8 +107,6 @@ endif()
109 107 
110## Tiling实现108## Tiling实现
111 109 
112-> 💡 **进阶内容**:关于Host侧Tiling实现的详细说明,包括基本流程、Tiling结构体定义、Tiling模板编程等,请参考[《AI Core算子开发进阶指南 - Host侧Tiling实现》](./aicore_develop_advanced_guide.md#host侧tiling实现)。
113- 
114### Tiling简介110### Tiling简介
115 111 
116因NPU中AI Core内部存储空间有限,无法一次性将整个张量数据加载到计算单元中处理,因此需要将输入张量切分为多个小块(Tile),逐块进行计算,这一过程称为Tiling。112因NPU中AI Core内部存储空间有限,无法一次性将整个张量数据加载到计算单元中处理,因此需要将输入张量切分为多个小块(Tile),逐块进行计算,这一过程称为Tiling。
@@ -120,7 +116,9 @@ endif()
120### 代码实现116### 代码实现
121 117 
122Tiling一共需要三个交付件:```${op_name}_tiling.cpp``` ```${op_name}_tiling_key.h``` ```${op_name}_tiling_data.h```118Tiling一共需要三个交付件:```${op_name}_tiling.cpp``` ```${op_name}_tiling_key.h``` ```${op_name}_tiling_data.h```
119+ 
123> 说明:120> 说明:
121+>
124> 1. `${op_name}_tiling.cpp`放在`${op_name}/op_host`目录下;122> 1. `${op_name}_tiling.cpp`放在`${op_name}/op_host`目录下;
125> 2. `${op_name}_tiling_key.h`和`${op_name}_tiling_data.h`放在`${op_name}/op_kernel`目录下;123> 2. `${op_name}_tiling_key.h`和`${op_name}_tiling_data.h`放在`${op_name}/op_kernel`目录下;
126> 3. 如果`${op_name}_tiling.cpp`中需要引用`${op_name}_tiling_data.h`,请使用相对路径的方式,例如:`#include "../op_kernel/${op_name}_tiling_data.h"`。124> 3. 如果`${op_name}_tiling.cpp`中需要引用`${op_name}_tiling_data.h`,请使用相对路径的方式,例如:`#include "../op_kernel/${op_name}_tiling_data.h"`。
@@ -132,7 +130,7 @@ Tiling主要切分逻辑。
132如需查看详细实现,请参考[add_example_tiling.cpp](../../../examples/add_example/op_host/add_example_tiling.cpp)。130如需查看详细实现,请参考[add_example_tiling.cpp](../../../examples/add_example/op_host/add_example_tiling.cpp)。
133 131 
134> **样例中函数空实现说明:**132> **样例中函数空实现说明:**
135- 133+>
136> 1. **TilingParse**:图模式标准交付件,保留函数定义以满足框架调用规范,无实际逻辑时可置空。134> 1. **TilingParse**:图模式标准交付件,保留函数定义以满足框架调用规范,无实际逻辑时可置空。
137> 2. **CompileInfo**:图模式标准交付件,保留函数定义以满足框架调用规范,无实际逻辑时可置空。135> 2. **CompileInfo**:图模式标准交付件,保留函数定义以满足框架调用规范,无实际逻辑时可置空。
138 136 
@@ -203,6 +201,7 @@ static ge::graphStatus TilingFunc(gert::TilingContext* context){
203// 3.Tiling注册入口201// 3.Tiling注册入口
204IMPL_OP_OPTILING(${op_name}).Tiling(TilingFunc).TilingParse<CompileInfo>(TilingParse);202IMPL_OP_OPTILING(${op_name}).Tiling(TilingFunc).TilingParse<CompileInfo>(TilingParse);
205```203```
204+ 
206**交付件2:${op_name}_tiling_key.h**205**交付件2:${op_name}_tiling_key.h**
207 206 
208TilingKey是一个算子内为了区分不同的实现而将kernel代码进行区分的方法,kernel侧可以通过TilingKey来选择不同的算法逻辑。207TilingKey是一个算子内为了区分不同的实现而将kernel代码进行区分的方法,kernel侧可以通过TilingKey来选择不同的算法逻辑。
@@ -237,8 +236,6 @@ struct ${op_name}TilingData {
237 236 
238## Kernel实现237## Kernel实现
239 238 
240-> 💡 **进阶内容**:关于Kernel侧算子实现的详细说明,包括核函数定义、GET_TILING_DATA获取Tiling参数、核函数内推导输入数据类型和格式等,请参考[《AI Core算子开发进阶指南 - Kernel侧算子实现》](./aicore_develop_advanced_guide.md#kernel侧算子实现)。
241- 
242### Kernel简介239### Kernel简介
243 240 
244Kernel是算子在NPU执行的核心部分,负责张量数据的加载、计算和存储,是算子功能实现的最终载体。Kernel的实现需要与Tiling策略紧密配合,根据Tiling提供的`TilingData``TilingKey`信息进行内存分配和计算调度。241Kernel是算子在NPU执行的核心部分,负责张量数据的加载、计算和存储,是算子功能实现的最终载体。Kernel的实现需要与Tiling策略紧密配合,根据Tiling提供的`TilingData``TilingKey`信息进行内存分配和计算调度。
@@ -247,10 +244,10 @@ Kernel实现包括如下步骤,整个流程通过`Process`函数串联,实
247 244 
248```mermaid245```mermaid
249graph LR246graph LR
250- H([核函数定义]) -->A([定义Kernel类])247+ H([核函数定义]) -->A([定义Kernel类])
251- A -->B([初始化函数<br>Init])248+ A -->B([初始化函数<br>Init])
252 B -->C([主处理函数<br>Process])249 B -->C([主处理函数<br>Process])
253- subgraph C [主处理函数 Process]250+ subgraph C [主处理函数Process]
254 D([数据搬入<br>CopyIn]) -->E([计算<br>Compute]) -->F([数据搬出<br>CopyOut])251 D([数据搬入<br>CopyIn]) -->E([计算<br>Compute]) -->F([数据搬出<br>CopyOut])
255 end252 end
256 F -->G([Kernel执行完成])253 F -->G([Kernel执行完成])
@@ -260,6 +257,7 @@ graph LR
260 257 
261Kernel一共需要两个交付件:```${op_name}.cpp``` ```${op_name}.h```258Kernel一共需要两个交付件:```${op_name}.cpp``` ```${op_name}.h```
262> 说明:259> 说明:
260+>
263> 1. `${op_name}.cpp`为kernel的入口函数只能放在`${op_name}/op_kernel`目录下;261> 1. `${op_name}.cpp`为kernel的入口函数只能放在`${op_name}/op_kernel`目录下;
264> 2. `${op_name}.h`文件可以按照不同SoC或模板放在对应目录下,例如:`${op_name}/op_kernel/arch32`、`${op_name}/op_kernel/arch35`或`${op_name}/op_kernel/impl`等目录下;262> 2. `${op_name}.h`文件可以按照不同SoC或模板放在对应目录下,例如:`${op_name}/op_kernel/arch32`、`${op_name}/op_kernel/arch35`或`${op_name}/op_kernel/impl`等目录下;
265 263 
@@ -326,7 +324,7 @@ private:
326 TQue<QuePosition::VECIN, BUFFER_NUM> inputQueueX_;324 TQue<QuePosition::VECIN, BUFFER_NUM> inputQueueX_;
327 // 输入队列Y,从GM拷贝到LM,BUFFER_NUM表示buffer数量,开启double buff达到流水并行,为2325 // 输入队列Y,从GM拷贝到LM,BUFFER_NUM表示buffer数量,开启double buff达到流水并行,为2
328 TQue<QuePosition::VECIN, BUFFER_NUM> inputQueueY_;326 TQue<QuePosition::VECIN, BUFFER_NUM> inputQueueY_;
329- // 输出队列Z,从LM拷贝到GM,BUFFER_NUM表示 buffer数量,这里开启double buff达到流水并行,为2327+ // 输出队列Z,从LM拷贝到GM,BUFFER_NUM表示buffer数量,这里开启double buff达到流水并行,为2
330 TQue<QuePosition::VECOUT, BUFFER_NUM> outputQueueZ_;328 TQue<QuePosition::VECOUT, BUFFER_NUM> outputQueueZ_;
331 329 
332 // 输入X的GM地址330 // 输入X的GM地址
@@ -377,16 +375,12 @@ __aicore__ inline void AddExample<T>::Process()
377 375 
378## 图模式适配376## 图模式适配
379 377 
380-> 💡 **进阶内容**:关于GE图模式原型定义的详细说明,包括REG_OP接口、TensorType类等,请参考[《AI Core算子开发进阶指南 - GE图模式原型定义》](./aicore_develop_advanced_guide.md#ge图模式原型定义)。
381- 
382图模式一共需要三个交付件:```${op_name}_graph_infer.cpp``` ```${op_name}_infershape.cpp``` ```${op_name}_proto.h```378图模式一共需要三个交付件:```${op_name}_graph_infer.cpp``` ```${op_name}_infershape.cpp``` ```${op_name}_proto.h```
383详细说明见图模式适配指南[graph_develop_guide.md](./graph_develop_guide.md)。379详细说明见图模式适配指南[graph_develop_guide.md](./graph_develop_guide.md)。
384 380 
385## aclnn适配381## aclnn适配
386 382 
387-> 💡关于Aclnn接口的详细说明包括自动生成配置方式、动态库路径等请参考[《AI Core算子开发进阶指南 - Aclnn指导》](./aicore_develop_advanced_guide.md#aclnn指导)383+通常算子开发和编译完成后自动生成aclnn接口(一套基于C的API)可直接在应用程序中调用aclnn接口实现调用算子。
388- 
389-通常算子开发和编译完成后,会自动生成aclnn接口(一套基于C 的API),可直接在应用程序中调用aclnn接口实现调用算子。
390 384 
391为实现该调用方式,需提前生成算子对应的二进制包,增加二进制编译json文件,以`AddExample`算子为例:385为实现该调用方式,需提前生成算子对应的二进制包,增加二进制编译json文件,以`AddExample`算子为例:
392 386 
@@ -405,8 +399,8 @@ __aicore__ inline void AddExample<T>::Process()
405 参考[工程创建](#工程创建)完成基础环境搭建,同时检查算子开发交付件是否完备,是否在对应算子分类目录下。399 参考[工程创建](#工程创建)完成基础环境搭建,同时检查算子开发交付件是否完备,是否在对应算子分类目录下。
406 400 
4072. **配置环境变量。**4012. **配置环境变量。**
408- 402+ 
409- 根据实际场景,选择合适的命令。403+ 根据实际场景,选择合适的命令。
410 404 
411 ```bash405 ```bash
412 # 默认路径安装,以root用户为例(非root用户,将/usr/local替换为${HOME})406 # 默认路径安装,以root用户为例(非root用户,将/usr/local替换为${HOME})
@@ -466,9 +460,9 @@ __aicore__ inline void AddExample<T>::Process()
466 460 
467算子开发过程中,可通过如下方式进行验证:461算子开发过程中,可通过如下方式进行验证:
468 462 
469-1. [UT验证](#UT验证): 验证交付件代码能否正常运行。UT验证无需NPU环境。463+1. [UT验证](#ut验证)验证交付件代码能否正常运行。UT验证无需NPU环境。
470 464 
471-2. [aclnn调用验证](#aclnn调用验证): 验证算子在NPU环境上的功能。aclnn调用验证需要NPU环境。465+2. [aclnn调用验证](#aclnn调用验证)验证算子在NPU环境上的功能。aclnn调用验证需要NPU环境。
472 466 
473### UT验证467### UT验证
474 468 
@@ -498,7 +492,7 @@ UT编写指导如下,如需查看详细实现,请参考样例UT实现[test_a
498 492 
499**1. 组织结构与命名建议**493**1. 组织结构与命名建议**
500 494 
501-- **头文件**:统一包含`iostream`, `gtest/gtest.h`、`infershape_context_faker.h`、`infershape_case_executor.h`。495+- **头文件**:统一包含`iostream``gtest/gtest.h`、`infershape_context_faker.h`、`infershape_case_executor.h`。
502- **测试类**:继承`testing::Test`,实现`SetUpTestCase/TearDownTestCase`统一做数据准备与清理。496- **测试类**:继承`testing::Test`,实现`SetUpTestCase/TearDownTestCase`统一做数据准备与清理。
503- **命名**:测试类建议`${OpName}InfershapeTest`,用例名建议`test_case_xxx`,可读性更高。497- **命名**:测试类建议`${OpName}InfershapeTest`,用例名建议`test_case_xxx`,可读性更高。
504 498 
@@ -562,7 +556,7 @@ UT编写指导如下,如需查看详细实现,请参考样例UT实现[test_a
562 556 
563**1. 组织结构与命名建议**557**1. 组织结构与命名建议**
564 558 
565-- **头文件**:统一包含`iostream`, `gtest/gtest.h`、`tiling_context_faker.h`、`tiling_case_executor.h`。559+- **头文件**:统一包含`iostream``gtest/gtest.h`、`tiling_context_faker.h`、`tiling_case_executor.h`。
566 - 若tiling头文件中已经定义CompileInfo结构体,则也需引入。560 - 若tiling头文件中已经定义CompileInfo结构体,则也需引入。
567- **测试类**:继承`testing::Test`,实现`SetUpTestCase/TearDownTestCase`统一做数据准备与清理。561- **测试类**:继承`testing::Test`,实现`SetUpTestCase/TearDownTestCase`统一做数据准备与清理。
568- **命名**:测试类建议`${OpName}TilingTest`,用例名建议`test_case_xxx`,可读性更高。562- **命名**:测试类建议`${OpName}TilingTest`,用例名建议`test_case_xxx`,可读性更高。
@@ -751,11 +745,7 @@ export LD_LIBRARY_PATH=${ASCEND_HOME_PATH}/opp/vendors/${vendor_name}_cv/op_api/
751 745 
752## 附录746## 附录
753 747 
754-自定义算子如需运行图模式,不需要aclnn适配,详细内容请参考[图模式开发指南](./graph_develop_guide.md)。748+自定义算子如需运行图模式,不需要aclnn适配,详细内容请参考[图模式开发指南](./graph_develop_guide.md)。
755- 
756-> 💡 **进阶内容**
757->
758-> - 关于多芯片代际隔离的详细说明,包括芯片架构映射、隔离位置清单、Kernel入口配置等,请参考[《AI Core算子开发进阶指南 - 代际隔离说明》](./aicore_develop_advanced_guide.md#代际隔离说明)。
759 749 
760### 算子工程迁移750### 算子工程迁移
761 751 
@@ -1078,7 +1068,7 @@ template<int D_T_X, int D_T_Y, int D_T_Z, int TILE_NUM, int IS_SPLIT>
1078 1068 
1079保留原有op\_kernel/tiling\_key\_{op\_name}.h中算子的模板参数定义,若不存在op\_kernel/tiling\_key\_{op\_name}.h,请参考[add_example_tiling_key.h](../../../examples/add_example/op_kernel/add_example_tiling_key.h)新增定义模板参数和模板参数组合。1069保留原有op\_kernel/tiling\_key\_{op\_name}.h中算子的模板参数定义,若不存在op\_kernel/tiling\_key\_{op\_name}.h,请参考[add_example_tiling_key.h](../../../examples/add_example/op_kernel/add_example_tiling_key.h)新增定义模板参数和模板参数组合。
1080 1070 
1081-### 算子跨平台迁移1071+### 算子跨平台迁移
1082 1072 
1083完成算子代码开发后,如需实现多平台间(如Atlas A2/A3等)的算子代码迁移,需考虑硬件结构差异引发的软件实现变更。1073完成算子代码开发后,如需实现多平台间(如Atlas A2/A3等)的算子代码迁移,需考虑硬件结构差异引发的软件实现变更。
1084 1074 
@@ -51,7 +51,7 @@ Create the AI CPU initial directory for ${op_name} under ${op_class} success
51 51 
52创建完成后,目录结构如下所示:52创建完成后,目录结构如下所示:
53 53 
54-```54+```tex
55${op_name} # 替换为实际算子名的小写下划线形式55${op_name} # 替换为实际算子名的小写下划线形式
56├── examples # 算子调用示例56├── examples # 算子调用示例
57│ └── test_aclnn_${op_name}.cpp # 算子aclnn调用示例57│ └── test_aclnn_${op_name}.cpp # 算子aclnn调用示例
@@ -68,7 +68,7 @@ ${op_name} # 替换为实际算子名的小写下
68 68 
69```${op_class}```为全新算子分类需额外在`CMakeLists`中添加`add_subdirectory(${op_class})`,否则无法正常编译。69```${op_class}```为全新算子分类需额外在`CMakeLists`中添加`add_subdirectory(${op_class})`,否则无法正常编译。
70 70 
71-```71+```Cpp
72if(ENABLE_EXPERIMENTAL)72if(ENABLE_EXPERIMENTAL)
73 # genop新增experimental算子分类73 # genop新增experimental算子分类
74 # add_subdirectory(${op_class})74 # add_subdirectory(${op_class})
@@ -197,7 +197,7 @@ REGISTER_CPU_KERNEL(kAddExample, AddExampleCpuKernel);
197 197 
198## aclnn适配198## aclnn适配
199 199 
200-通常算子开发和编译完成后,会自动生成aclnn接口(一套基于C 的API),无需做其他配置,可直接在应用程序中调用aclnn接口实现调用算子。200+通常算子开发和编译完成后,会自动生成aclnn接口(一套基于C的API),无需做其他配置,可直接在应用程序中调用aclnn接口实现调用算子。
201 201 
202## 编译部署202## 编译部署
203 203 
@@ -234,12 +234,13 @@ REGISTER_CPU_KERNEL(kAddExample, AddExampleCpuKernel);
234 # 安装run包234 # 安装run包
235 ./build_out/cann-ops-cv-${vendor_name}_linux-${arch}.run235 ./build_out/cann-ops-cv-${vendor_name}_linux-${arch}.run
236 ```236 ```
237+ 
237 自定义算子包安装在```${ASCEND_HOME_PATH}/opp/vendors```路径中,```${ASCEND_HOME_PATH}```表示CANN软件安装目录,可提前在环境变量中配置。238 自定义算子包安装在```${ASCEND_HOME_PATH}/opp/vendors```路径中,```${ASCEND_HOME_PATH}```表示CANN软件安装目录,可提前在环境变量中配置。
238- 239+ 
2394. **(可选)卸载自定义算子包。**2404. **(可选)卸载自定义算子包。**
240 241 
241 自定义算子包安装后在```${ASCEND_HOME_PATH}/opp/vendors/custom_cv/scripts```目录会生成`uninstall.sh`,通过该脚本可卸载自定义算子包,命令如下:242 自定义算子包安装后在```${ASCEND_HOME_PATH}/opp/vendors/custom_cv/scripts```目录会生成`uninstall.sh`,通过该脚本可卸载自定义算子包,命令如下:
242- 243+ 
243 ```bash244 ```bash
244 bash ${ASCEND_HOME_PATH}/opp/vendors/custom_cv/scripts/uninstall.sh245 bash ${ASCEND_HOME_PATH}/opp/vendors/custom_cv/scripts/uninstall.sh
245 ```246 ```
@@ -50,12 +50,12 @@
50 50
51 <tr>51 <tr>
52 <td rowspan="2">Memory</td>52 <td rowspan="2">Memory</td>
53- <td>Memory 容量(GB)</td>53+ <td>Memory容量(GB)</td>
54 <td>64</td>54 <td>64</td>
55 <td>128</td>55 <td>128</td>
56 </tr>56 </tr>
57 <tr>57 <tr>
58- <td>Memory 带宽</td>58+ <td>Memory带宽</td>
59 <td>1.6TB/s</td>59 <td>1.6TB/s</td>
60 <td>1.6TB/s</td>60 <td>1.6TB/s</td>
61 </tr>61 </tr>
@@ -83,7 +83,7 @@
83 <td>可用ND2NZ/DN2NZ在MTE2阶段完成格式转换,减少中间buffer和格式转换开销;需关注步长、对齐与NZ形状映射</td>83 <td>可用ND2NZ/DN2NZ在MTE2阶段完成格式转换,减少中间buffer和格式转换开销;需关注步长、对齐与NZ形状映射</td>
84 </tr>84 </tr>
85 <tr>85 <tr>
86- <td>支持Cube-&gt;Vector高效内部 数据通路:L1&lt;-&gt;UB、L0C-&gt;UB、FIXP-&gt;UB</td>86+ <td>支持Cube-&gt;Vector高效内部数据通路:L1&lt;-&gt;UB、L0C-&gt;UB、FIXP-&gt;UB</td>
87 <td>可在UB侧做中间累加/激活/融合(如切K累加、后处理),减少GM往返;对应同步与管线切分需调整</td>87 <td>可在UB侧做中间累加/激活/融合(如切K累加、后处理),减少GM往返;对应同步与管线切分需调整</td>
88 </tr>88 </tr>
89 <tr>89 <tr>
@@ -124,7 +124,7 @@
1241. 确认算子涉及的计算单元(Cube/Vector)和对应单元支持的数据类型是否在平台间存在差异。1241. 确认算子涉及的计算单元(Cube/Vector)和对应单元支持的数据类型是否在平台间存在差异。
1252. 确认涉及的数据搬运单元(ND-&gt;NZ、GM&lt;-&gt;Lx、集合通信等)是否在平台间存在差异。1252. 确认涉及的数据搬运单元(ND-&gt;NZ、GM&lt;-&gt;Lx、集合通信等)是否在平台间存在差异。
1263. 按硬件能力变更点逐项对照修改(Vector架构、Cube支持数据类型、L1/L0/UB大小、CCU通信等)。1263. 按硬件能力变更点逐项对照修改(Vector架构、Cube支持数据类型、L1/L0/UB大小、CCU通信等)。
127-4. 参考算子迁移样例调整/补齐 Atlas A2/Ascend 950 分支逻辑。127+4. 参考算子迁移样例调整/补齐Atlas A2/Ascend 950分支逻辑。
128 128 
129## 四、算子迁移样例129## 四、算子迁移样例
130 130 
@@ -337,7 +337,7 @@ Ascend 950新架构引入UB2L1 & L0C2UB间的直连通路,实现矩阵计算
337 337 
338**矩阵搬出**338**矩阵搬出**
339 339 
340-启用L0C至UB(L0C2UB)直连通路,通过DataCopy接口,支持融合算子的矩阵计算结果直接搬入UB进行后续向量计算。 340+启用L0C至UB(L0C2UB)直连通路,通过DataCopy接口,支持融合算子的矩阵计算结果直接搬入UB进行后续向量计算。
341 341 
342对于切K或多阶段融合场景,可将"L0C搬回GM再读回UB"改为"L0C直达UB累加/后处理",降低GM往返带宽压力和时延。迁移时建议把中间结果归并、激活/量化前处理放到UB侧完成,并显式梳理MTE1/MTE2/MTE3与计算单元的事件同步顺序,确保跨单元流水连续,避免由于新增通路引入数据可见性或同步时序问题。关键接口定义可参考:342对于切K或多阶段融合场景,可将"L0C搬回GM再读回UB"改为"L0C直达UB累加/后处理",降低GM往返带宽压力和时延。迁移时建议把中间结果归并、激活/量化前处理放到UB侧完成,并显式梳理MTE1/MTE2/MTE3与计算单元的事件同步顺序,确保跨单元流水连续,避免由于新增通路引入数据可见性或同步时序问题。关键接口定义可参考:
343 343 
@@ -162,7 +162,7 @@
162 162 
163 \$\{static\_lib\_path\}表示静态库解压路径。解压后目录结构如下:163 \$\{static\_lib\_path\}表示静态库解压路径。解压后目录结构如下:
164 164 
165- ```165+ ```bash
166 ├── cann-${soc_name}-ops-cv-static_${cann_version}_linux-${arch}166 ├── cann-${soc_name}-ops-cv-static_${cann_version}_linux-${arch}
167 │ ├── lib64167 │ ├── lib64
168 │ │ ├── libcann_cv_static.a # 静态库文件168 │ │ ├── libcann_cv_static.a # 静态库文件
@@ -178,7 +178,7 @@
178 178 
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
@@ -232,7 +232,7 @@
232 232 
233 未联网和联网场景下编译得到算子包结果一样,默认存放于项目根目录build_out目录下,并且安装和卸载的操作命令也一样,具体参见[联网编译](#联网编译)。233 未联网和联网场景下编译得到算子包结果一样,默认存放于项目根目录build_out目录下,并且安装和卸载的操作命令也一样,具体参见[联网编译](#联网编译)。
234 234 
235-## 本地验证 235+## 本地验证
236 236 
237源码包部署后,可通过项目根目录build.sh执行UT用例,验证项目功能是否正常。237源码包部署后,可通过项目根目录build.sh执行UT用例,验证项目功能是否正常。
238 238 
@@ -121,7 +121,7 @@
121 ./test_aclnn_grid_sample3_d121 ./test_aclnn_grid_sample3_d
122 ```122 ```
123 123 
124- \${static_lib_path}表示静态库统一放置路径;\${ASCEND_INSTALL_PATH}已通过环境变量配置,表示CANN toolkit包安装路径; 最终可执行文件名请替换为**实际算子可执行文件名**。124+ \${static_lib_path}表示静态库统一放置路径;\${ASCEND_INSTALL_PATH}已通过环境变量配置,表示CANN toolkit包安装路径;最终可执行文件名请替换为**实际算子可执行文件名**。
125 125 
126 其中lcann_cv_static、lcann_legacy_static表示算子依赖的静态库文件,从静态库统一放置路径\${static_lib_path}中获取;126 其中lcann_cv_static、lcann_legacy_static表示算子依赖的静态库文件,从静态库统一放置路径\${static_lib_path}中获取;
127 lgraph、lmetadef等表示算子依赖的底层库文件,可在CANN toolkit包获取。127 lgraph、lmetadef等表示算子依赖的底层库文件,可在CANN toolkit包获取。
@@ -8,7 +8,7 @@
8 8 
9 调用算子API时,需引用依赖的头文件和库文件,一般头文件默认在`${INSTALL_DIR}/include/aclnnop`,库文件默认在`${INSTALL_DIR}/lib64`,具体文件如下:9 调用算子API时,需引用依赖的头文件和库文件,一般头文件默认在`${INSTALL_DIR}/include/aclnnop`,库文件默认在`${INSTALL_DIR}/lib64`,具体文件如下:
10 10 
11- - 头文件:方式1 (推荐):引用算子仓总头文件aclnn\_ops\_\$\{ops\_project\}.h。方式2:引用单个算子API的头文件aclnn\_\*.h。11+ - 头文件:方式1(推荐):引用算子仓总头文件aclnn\_ops\_\$\{ops\_project\}.h。方式2:引用单个算子API的头文件aclnn\_\*.h。
12 - 库文件:引用算子仓对应的库文件libopapi_${ops_project}.so。注意,原所有算子仓总库文件libopapi.so后续会废弃,不推荐使用,也不支持与单个算子仓库文件同时使用。12 - 库文件:引用算子仓对应的库文件libopapi_${ops_project}.so。注意,原所有算子仓总库文件libopapi.so后续会废弃,不推荐使用,也不支持与单个算子仓库文件同时使用。
13 13 
14 ${INSTALL_DIR}表示CANN安装后文件路径;\$\{ops\_project\}表示算子仓名(如math、nn、cv、transformer),请改为实际算子仓名。14 ${INSTALL_DIR}表示CANN安装后文件路径;\$\{ops\_project\}表示算子仓名(如math、nn、cv、transformer),请改为实际算子仓名。
@@ -73,6 +73,6 @@
73| [aclnnRoiAlign](../../objdetect/roi_align/docs/aclnnRoiAlign.md) | RoIAlign是一种池化层,用于非均匀输入尺寸的特征图,并输出固定尺寸的特征图。 |默认确定性实现|- |73| [aclnnRoiAlign](../../objdetect/roi_align/docs/aclnnRoiAlign.md) | RoIAlign是一种池化层,用于非均匀输入尺寸的特征图,并输出固定尺寸的特征图。 |默认确定性实现|- |
74| [aclnnRoiAlignV2](../../objdetect/roi_align/docs/aclnnRoiAlignV2.md) | RoIAlign是一种池化层,用于非均匀输入尺寸的特征图,并输出固定尺寸的特征图。 |默认确定性实现|- |74| [aclnnRoiAlignV2](../../objdetect/roi_align/docs/aclnnRoiAlignV2.md) | RoIAlign是一种池化层,用于非均匀输入尺寸的特征图,并输出固定尺寸的特征图。 |默认确定性实现|- |
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) | 从批处理输入张量中提取滑动局部块,将滑动局部块数组合并为一个大张量。 |默认确定性实现|默认确定性实现|
@@ -5,7 +5,7 @@
5>- **算子目录**:目录名为算子名小写下划线形式,每个目录承载该算子所有交付件,包括代码实现、examples、文档等,目录介绍参见[项目目录](./install/dir_structure.md)。5>- **算子目录**:目录名为算子名小写下划线形式,每个目录承载该算子所有交付件,包括代码实现、examples、文档等,目录介绍参见[项目目录](./install/dir_structure.md)。
6>- **算子执行硬件单元**:大部分算子运行在AI Core,少部分算子运行在AI CPU。默认情况下,项目中提到的算子一般指AI Core算子。关于AI Core和AI CPU详细介绍参见[《Ascend C算子开发》](https://hiascend.com/document/redirect/CannCommunityOpdevAscendC)中“概念原理和术语 > 硬件架构与数据处理原理”。6>- **算子执行硬件单元**:大部分算子运行在AI Core,少部分算子运行在AI CPU。默认情况下,项目中提到的算子一般指AI Core算子。关于AI Core和AI CPU详细介绍参见[《Ascend C算子开发》](https://hiascend.com/document/redirect/CannCommunityOpdevAscendC)中“概念原理和术语 > 硬件架构与数据处理原理”。
7>- **算子接口列表**:为方便调用算子,CANN提供一套C API执行算子,一般以aclnn为前缀,全量接口参见[aclnn列表](op_api_list.md)。7>- **算子接口列表**:为方便调用算子,CANN提供一套C API执行算子,一般以aclnn为前缀,全量接口参见[aclnn列表](op_api_list.md)。
8-> - **V版本演进说明**:部分算子存在多个V版本,使用时选择最高V版本即可(高版本算子已兼容低版本算子的所有能力)。8+>- **V版本演进说明**:部分算子存在多个V版本,使用时选择最高V版本即可(高版本算子已兼容低版本算子的所有能力)。
9 9 
10项目提供的所有算子分类和算子列表如下:10项目提供的所有算子分类和算子列表如下:
11 11 
@@ -44,7 +44,7 @@
44 <td>&cross;</td>44 <td>&cross;</td>
45 <td>&check;</td>45 <td>&check;</td>
46 <td>AI CPU</td>46 <td>AI CPU</td>
47- <td>对 RGB 图像的饱和度进行调整。</td>47+ <td>对RGB图像的饱和度进行调整。</td>
48 </tr>48 </tr>
49 <tr>49 <tr>
50 <td>image</td>50 <td>image</td>
@@ -514,7 +514,7 @@
514 <td>&check;</td>514 <td>&check;</td>
515 <td>&cross;</td>515 <td>&cross;</td>
516 <td>AI Core</td>516 <td>AI Core</td>
517- <td>对输入特征图按 ROI(感兴趣区域)进行池化,在每个 ROI 内按空间划分格子,对每个格子做最大池化,并输出池化结果及最大值在通道内的一维索引。</td>517+ <td>对输入特征图按ROI(感兴趣区域)进行池化,在每个ROI内按空间划分格子,对每个格子做最大池化,并输出池化结果及最大值在通道内的一维索引。</td>
518 </tr>518 </tr>
519 <tr>519 <tr>
520 <td>objdetect</td>520 <td>objdetect</td>
@@ -547,4 +547,4 @@
547 <td>完成带掩码非极大值抑制计算。</td>547 <td>完成带掩码非极大值抑制计算。</td>
548 </tr>548 </tr>
549</tbody>549</tbody>
550-</table>550+</table>
@@ -4,7 +4,7 @@
4 4 
5本文档演示如何使用Ascend C和[PyTorch Extension](https://docs.pytorch.org/tutorials/extension.html)能力开发自定义NPU算子。5本文档演示如何使用Ascend C和[PyTorch Extension](https://docs.pytorch.org/tutorials/extension.html)能力开发自定义NPU算子。
6 6 
7-**核心优势:**7+**核心优势**
8 8 
9- **单交付件:** 一个文件完成算子开发和PyTorch框架适配。9- **单交付件:** 一个文件完成算子开发和PyTorch框架适配。
10 10 
@@ -22,14 +22,14 @@
22 22 
231. 进入`examples/fast_kernel_launch_example`目录。231. 进入`examples/fast_kernel_launch_example`目录。
24 24 
25-2. 安装依赖 | Install Dependencies:25+2. 安装依赖 | Install Dependencies
26- 26+ 
27 ```sh27 ```sh
28 python3 -m pip install -r requirements.txt28 python3 -m pip install -r requirements.txt
29 ```29 ```
30 30 
31-3. 构建Wheel包 | Build the Wheel:31+3. 构建Wheel包 | Build the Wheel
32- 32+ 
33 ```sh33 ```sh
34 # -n: non-isolated build (uses existing environment)34 # -n: non-isolated build (uses existing environment)
35 python3 -m build --wheel -n35 python3 -m build --wheel -n
@@ -38,7 +38,7 @@
38 构建完成后,产物在当前目录的`dist`文件夹下,产物名`ascend_ops-1.0.0-${python_version}-abi3-${arch}.whl`,38 构建完成后,产物在当前目录的`dist`文件夹下,产物名`ascend_ops-1.0.0-${python_version}-abi3-${arch}.whl`,
39 `${python_version}`表示当前环境中的python版本(python3.8.3为cp38),`${arch}`表示CPU架构。39 `${python_version}`表示当前环境中的python版本(python3.8.3为cp38),`${arch}`表示CPU架构。
40 40 
41-4. 安装Wheel包 | Install Package:41+4. 安装Wheel包 | Install Package
42 42
43 ```sh43 ```sh
44 python3 -m pip install dist/*.whl --force-reinstall --no-deps44 python3 -m pip install dist/*.whl --force-reinstall --no-deps
@@ -89,10 +89,10 @@ print("Verification successful!")
89 89 
90 这里`dav-2201`为ascend910b芯片对应的编译参数,获取方法参考[NpuArch说明和使用指导](https://gitcode.com/cann/ops-math/wiki/NpuArch%E8%AF%B4%E6%98%8E%E5%92%8C%E4%BD%BF%E7%94%A8%E6%8C%87%E5%AF%BC.md)。90 这里`dav-2201`为ascend910b芯片对应的编译参数,获取方法参考[NpuArch说明和使用指导](https://gitcode.com/cann/ops-math/wiki/NpuArch%E8%AF%B4%E6%98%8E%E5%92%8C%E4%BD%BF%E7%94%A8%E6%8C%87%E5%AF%BC.md)。
91 91 
92-3. 在soc目录下新建一个`add.cpp`(建议使用算子名为文件名)。这个文件包含了开发一个AI Core算子所需要的全部模块。92+3. 在soc目录下新建一个`add.cpp`建议使用算子名为文件名。这个文件包含了开发一个AI Core算子所需要的全部模块。
93 - 算子Schema注册93 - 算子Schema注册
94 - 算子Meta Function实现 & 注册94 - 算子Meta Function实现 & 注册
95- - 算子Kernel实现 (Ascend C)95+ - 算子Kernel实现Ascend C
96 - 算子NPU调用实现 & 注册96 - 算子NPU调用实现 & 注册
97 97 
98 ```cpp98 ```cpp
@@ -160,7 +160,7 @@ print("Verification successful!")
160 */160 */
161 torch::Tensor add_npu(const torch::Tensor &x, const torch::Tensor &y)161 torch::Tensor add_npu(const torch::Tensor &x, const torch::Tensor &y)
162 {162 {
163- // OptionalDeviceGuard 确保后续操作在正确的设备上下文执行163+ // OptionalDeviceGuard确保后续操作在正确的设备上下文执行
164 // 它会记录当前设备状态,执行完作用域代码后自动恢复164 // 它会记录当前设备状态,执行完作用域代码后自动恢复
165 const c10::OptionalDeviceGuard guard(x.device());165 const c10::OptionalDeviceGuard guard(x.device());
166 auto z = add_meta(x, y);166 auto z = add_meta(x, y);
@@ -9,7 +9,7 @@
9 9 
10## 功能说明10## 功能说明
11 11 
12-- 算子功能:根据人脸区域累积图、累积 mask、最大 mask 与背景图计算融合结果。12+- 算子功能:根据人脸区域累积图、累积mask、最大mask与背景图计算融合结果。
13- 计算公式:13- 计算公式:
14 14 
15 ```text15 ```text
@@ -47,14 +47,14 @@
47 <tr>47 <tr>
48 <td>acc_mask</td>48 <td>acc_mask</td>
49 <td>输入</td>49 <td>输入</td>
50- <td>人脸区域累积 mask。</td>50+ <td>人脸区域累积mask。</td>
51 <td>FLOAT32</td>51 <td>FLOAT32</td>
52 <td>ND</td>52 <td>ND</td>
53 </tr>53 </tr>
54 <tr>54 <tr>
55 <td>max_mask</td>55 <td>max_mask</td>
56 <td>输入</td>56 <td>输入</td>
57- <td>融合权重 mask,取值通常位于 [0, 1]。</td>57+ <td>融合权重mask,取值通常位于 [0, 1]。</td>
58 <td>FLOAT32</td>58 <td>FLOAT32</td>
59 <td>ND</td>59 <td>ND</td>
60 </tr>60 </tr>
@@ -68,7 +68,7 @@
68 <tr>68 <tr>
69 <td>epsilon</td>69 <td>epsilon</td>
70 <td>属性</td>70 <td>属性</td>
71- <td>除法数值稳定项,默认值为 1e-12。</td>71+ <td>除法数值稳定项,默认值为1e-12。</td>
72 <td>FLOAT32</td>72 <td>FLOAT32</td>
73 <td>-</td>73 <td>-</td>
74 </tr>74 </tr>
@@ -84,13 +84,13 @@
84 84 
85## 约束说明85## 约束说明
86 86 
87-- `acc_face`、`acc_mask`、`max_mask`、`bg_img` 和 `fused_img` 的 shape 需保持一致。87+- `acc_face`、`acc_mask`、`max_mask`、`bg_img` 和 `fused_img` 的shape需保持一致。
88-- `acc_face`、`acc_mask`、`max_mask` 和 `fused_img` 目前只支持 FLOAT32。88+- `acc_face`、`acc_mask`、`max_mask` 和 `fused_img` 目前只支持FLOAT32。
89-- `bg_img` 支持 FLOAT32 UINT8。89+- `bg_img` 支持FLOAT32和UINT8。
90-- 当前仅支持 ND 格式。90+- 当前仅支持ND格式。
91 91 
92## 调用说明92## 调用说明
93 93 
94| 调用方式 | 调用样例 | 说明 |94| 调用方式 | 调用样例 | 说明 |
95| -------- | -------- | ---- |95| -------- | -------- | ---- |
96-| aclnn 调用 | examples/test_aclnn_blend_face_bg_part_two.cpp | aclnn 调用示例。 |96+| aclnn调用 | examples/test_aclnn_blend_face_bg_part_two.cpp | aclnn调用示例。 |
@@ -8,8 +8,8 @@
8 8 
9## 功能说明9## 功能说明
10 10 
11-- 算子功能:对候选框执行非极大值抑制(NMS),输出候选框保留掩码,支持 IOU 阈值、置信度阈值过滤。11+- 算子功能:对候选框执行非极大值抑制(NMS),输出候选框保留掩码,支持IOU阈值、置信度阈值过滤。
12-- 输入坐标格式为 <x1,y1,x2,y2>(左下角和右上角坐标),在代码中以一维数组形式存储([N,4] 维度展开为长度 N*4 的一维数据)。12+- 输入坐标格式为 <x1,y1,x2,y2>(左下角和右上角坐标),在代码中以一维数组形式存储([N,4] 维度展开为长度N*4的一维数据)。
13- 计算公式:13- 计算公式:
14 - 输入节点:14 - 输入节点:
15 - x (shape[N,4], FLOAT32) - 候选框坐标(x1,y1,x2,y2)15 - x (shape[N,4], FLOAT32) - 候选框坐标(x1,y1,x2,y2)
@@ -18,12 +18,12 @@
18 - scores_threshold (shape[1], FLOAT32) - 置信度过滤阈值18 - scores_threshold (shape[1], FLOAT32) - 置信度过滤阈值
19 19 
20 - 计算节点:20 - 计算节点:
21- - Step1: 置信度降序排序索引 idx_sorted ( 该算子默认已降序排序,需用户自行保证)21+ - Step1: 置信度降序排序索引idx_sorted该算子默认已降序排序,需用户自行保证
22- - Step2: 初始化临时掩码标记所有候选框为保留状态(值为 1),输出掩码初始化为 0;22+ - Step2: 初始化临时掩码标记所有候选框为保留状态(值为1),输出掩码初始化为0;
23- - Step3: 遍历候选框,跳过已标记为过滤的框,过滤置信度低于scores_threshold的框(临时掩码置 0);23+ - Step3: 遍历候选框,跳过已标记为过滤的框,过滤置信度低于scores_threshold的框(临时掩码置0);
24- - Step4: 对保留的候选框标记输出掩码为 1,计算其与后续所有候选框的 IOU(IOU = 交集面积 /(当前框面积 + 对比框面积 - 交集面积));24+ - Step4: 对保留的候选框标记输出掩码为1,计算其与后续所有候选框的IOU($IOU = 交集面积 /(当前框面积 + 对比框面积 - 交集面积)$);
25- - Step5:将 IOU 高于iou_threshold的后续框标记为过滤(临时掩码置 0);25+ - Step5:将IOU高于iou_threshold的后续框标记为过滤(临时掩码置0);
26- - 重复步骤 Step3-Step5 直至所有框处理完成。26+ - 重复步骤Step3-Step5直至所有框处理完成。
27 27 
28 - 输出节点:28 - 输出节点:
29 - z (shape[N], UINT8) - 候选框保留掩码29 - z (shape[N], UINT8) - 候选框保留掩码
@@ -1,29 +1,30 @@
1# RoiAlignV21# RoiAlignV2
2 2 
3## 产品支持情况3## 产品支持情况
4+ 
4| 产品 | 是否支持 |5| 产品 | 是否支持 |
5| ---- | :----:|6| ---- | :----:|
6|Atlas A2 训练系列产品/Atlas A2 推理系列产品|√|7|Atlas A2 训练系列产品/Atlas A2 推理系列产品|√|
7 8 
8## 功能说明9## 功能说明
9 10 
10-- 算子功能:对输入特征图执行 ROI Align 操作,对每个感兴趣区域(ROI)进行双线性插值采样,输出固定大小的特征图。11+- 算子功能:对输入特征图执行ROI Align操作,对每个感兴趣区域(ROI)进行双线性插值采样,输出固定大小的特征图。
11-- 输入 ROI 坐标格式为 <batch_index, x1, y1, x2, y2>(左上角和右下角坐标),在代码中以一维数组形式存储([numRois, 5] 维度展开为长度 numRois*5 的一维数据)。12+- 输入ROI坐标格式为 <batch_index, x1, y1, x2, y2>(左上角和右下角坐标),在代码中以一维数组形式存储([numRois, 5] 维度展开为长度numRois*5的一维数据)。
12- 计算公式:13- 计算公式:
13 - 输入节点:14 - 输入节点:
14 - features (shape[N, C, H, W], FLOAT32) - 输入特征图15 - features (shape[N, C, H, W], FLOAT32) - 输入特征图
15 - rois (shape[numRois, 5], FLOAT32) - 感兴趣区域坐标(batch_index, x1, y1, x2, y2)16 - rois (shape[numRois, 5], FLOAT32) - 感兴趣区域坐标(batch_index, x1, y1, x2, y2)
16 17 
17 - 计算节点:18 - 计算节点:
18- - Step1: 将 ROI 坐标乘以 spatial_scale 进行缩放,并转换为 (x, y, w, h) 格式;19+ - Step1: 将ROI坐标乘以spatial_scale进行缩放,并转换为(x, y, w, h)格式;
19- - Step2: 根据 pooled_height pooled_width ROI 区域划分为均匀的 bin,计算每个 bin 的宽高 (bin_w, bin_h);20+ - Step2: 根据pooled_height和pooled_width将ROI区域划分为均匀的bin,计算每个bin的宽高(bin_w, bin_h);
20- - Step3: 根据 sampling_ratio 确定每个 bin 内的采样网格大小 (grid_h, grid_w),若 sampling_ratio > 0 则固定为该值,否则自适应计算 (ceil(roi_h / pooled_height), ceil(roi_w / pooled_width));21+ - Step3: 根据sampling_ratio确定每个bin内的采样网格大小(grid_h, grid_w),若sampling_ratio > 0则固定为该值,否则自适应计算(ceil(roi_h / pooled_height), ceil(roi_w / pooled_width));
21- - Step4: 对每个 bin 内的每个采样点,计算其在特征图上的坐标,通过双线性插值获取特征值;22+ - Step4: 对每个bin内的每个采样点,计算其在特征图上的坐标,通过双线性插值获取特征值;
22- - Step5: 对每个 bin 内所有采样点的特征值取平均,作为该位置的输出值;23+ - Step5: 对每个bin内所有采样点的特征值取平均,作为该位置的输出值;
23- - 重复上述步骤直至所有 ROI 的所有通道处理完成。24+ - 重复上述步骤直至所有ROI的所有通道处理完成。
24 25 
25 - 输出节点:26 - 输出节点:
26- - output (shape[numRois, C, pooled_height, pooled_width], FLOAT32) - 对齐后的 ROI 特征图27+ - output (shape[numRois, C, pooled_height, pooled_width], FLOAT32) - 对齐后的ROI特征图
27 28 
28## 参数说明29## 参数说明
29 30 
@@ -53,14 +54,14 @@
53 <tr>54 <tr>
54 <td>rois</td>55 <td>rois</td>
55 <td>输入</td>56 <td>输入</td>
56- <td>感兴趣区域坐标,shape [numRois, 5],每行为 (batch_index, x1, y1, x2, y2)。</td>57+ <td>感兴趣区域坐标,shape [numRois, 5],每行为(batch_index, x1, y1, x2, y2)。</td>
57 <td>FLOAT32</td>58 <td>FLOAT32</td>
58 <td>ND</td>59 <td>ND</td>
59 </tr>60 </tr>
60 <tr>61 <tr>
61 <td>output</td>62 <td>output</td>
62 <td>输出</td>63 <td>输出</td>
63- <td>对齐后的 ROI 特征图,shape [numRois, C, pooled_height, pooled_width]。</td>64+ <td>对齐后的ROI特征图,shape [numRois, C, pooled_height, pooled_width]。</td>
64 <td>FLOAT32</td>65 <td>FLOAT32</td>
65 <td>ND</td>66 <td>ND</td>
66 </tr>67 </tr>
@@ -81,14 +82,14 @@
81 <tr>82 <tr>
82 <td>spatial_scale</td>83 <td>spatial_scale</td>
83 <td>属性(可选)</td>84 <td>属性(可选)</td>
84- <td>空间缩放因子,用于将 ROI 坐标映射到特征图尺度。</td>85+ <td>空间缩放因子,用于将ROI坐标映射到特征图尺度。</td>
85 <td>FLOAT</td>86 <td>FLOAT</td>
86 <td>-</td>87 <td>-</td>
87 </tr>88 </tr>
88 <tr>89 <tr>
89 <td>sampling_ratio</td>90 <td>sampling_ratio</td>
90 <td>属性(可选)</td>91 <td>属性(可选)</td>
91- <td>每个 bin 的采样点数。大于 0 时固定为该值,否则自适应计算。</td>92+ <td>每个bin的采样点数。大于0时固定为该值,否则自适应计算。</td>
92 <td>INT32</td>93 <td>INT32</td>
93 <td>-</td>94 <td>-</td>
94 </tr>95 </tr>
@@ -96,11 +97,11 @@
96 97 
97## 约束说明98## 约束说明
98 99 
99-- 目前只支持 float32 输入100+- 目前只支持float32输入
100-- 目前只支持 ascend910b101+- 目前只支持ascend910b
101 102 
102## 调用说明103## 调用说明
103 104 
104| 调用方式 | 调用样例 | 说明 |105| 调用方式 | 调用样例 | 说明 |
105|--------------|------------------------------------------------------------------------|--------------------------------------------------------------|106|--------------|------------------------------------------------------------------------|--------------------------------------------------------------|
106-| aclnn调用 | [test_aclnn_roi_align_v2](./examples/test_aclnn_roi_align_v2.cpp) | 通过 aclnnRoiAlignV2 接口方式调用 RoiAlignV2 算子。 |107+| aclnn调用 | [test_aclnn_roi_align_v2](./examples/test_aclnn_roi_align_v2.cpp) | 通过aclnnRoiAlignV2接口方式调用RoiAlignV2算子。 |
@@ -13,20 +13,20 @@
13 13 
14## 功能说明14## 功能说明
15 15 
16-- 算子功能:对 RGB 图像的饱和度进行调整。16+- 算子功能:对RGB图像的饱和度进行调整。
17 17 
18## 参数说明18## 参数说明
19 19 
20| 参数名 | 输入/输出/属性 | 描述 | 数据类型 | 数据格式 |20| 参数名 | 输入/输出/属性 | 描述 | 数据类型 | 数据格式 |
21|-----|----------|----------------------------------------------------------------------------------------------------------------------------------------------------|----------------|------|21|-----|----------|----------------------------------------------------------------------------------------------------------------------------------------------------|----------------|------|
22-| images | 输入 | 输入 Tensor,至少为三维,最后一维大小必须为 3,按 RGB 三通道图像解释。 | FLOAT16、FLOAT | ND |22+| images | 输入 | 输入Tensor,至少为三维,最后一维大小必须为3,按RGB三通道图像解释。 | FLOAT16、FLOAT | ND |
23-| scale | 输入 | 包含 1 个元素的一维 Tensor,其唯一元素表示饱和度缩放因子。 | FLOAT | ND |23+| scale | 输入 | 包含1个元素的一维Tensor,其唯一元素表示饱和度缩放因子。 | FLOAT | ND |
24-| y | 输出 | 输出 Tensor,shape 和数据类型与 images 相同。 | FLOAT16、FLOAT | ND |24+| y | 输出 | 输出Tensor,shape和数据类型与images相同。 | FLOAT16、FLOAT | ND |
25 25 
26## 约束说明26## 约束说明
27 27 
28-- images 至少为三维,且最后一维必须为 3。28+- images至少为三维,且最后一维必须为3。
29-- scale 必须为仅包含 1 个元素的 FLOAT Tensor。29+- scale必须为仅包含1个元素的FLOAT Tensor。
30-- y 的数据类型和数据大小必须与 images 一致。30+- y的数据类型和数据大小必须与images一致。
31 31 
32-## 调用说明32+## 调用说明
@@ -23,7 +23,7 @@
23 23 
24 2. 根据传入的图片格式和转换参数,对图片进行色域转换,计算公式如下:24 2. 根据传入的图片格式和转换参数,对图片进行色域转换,计算公式如下:
25 25 
26- - YUV RGB:26+ - YUV转RGB:
27 27 
28 ```text28 ```text
29 | R | | matrix_r0c0 matrix_r0c1 matrix_r0c2 | | Y - input_bias_0 |29 | R | | matrix_r0c0 matrix_r0c1 matrix_r0c2 | | Y - input_bias_0 |
@@ -31,7 +31,7 @@
31 | B | | matrix_r2c0 matrix_r2c1 matrix_r2c2 | | V - input_bias_2 |31 | B | | matrix_r2c0 matrix_r2c1 matrix_r2c2 | | V - input_bias_2 |
32 ```32 ```
33 33 
34- - RGB YUV:34+ - RGB转YUV:
35 35 
36 ```text36 ```text
37 | Y | | matrix_r0c0 matrix_r0c1 matrix_r0c2 | | R | | output_bias_0 |37 | Y | | matrix_r0c0 matrix_r0c1 matrix_r0c2 | | R | | output_bias_0 |
@@ -96,7 +96,7 @@
96 96 
97## 约束说明97## 约束说明
98 98 
99-- aipp_config_path配置文件说明:支持动态AIPP和静态AIPP配置,其中静态AIPP可配置图片裁剪参数配置(Crop)、色域转换参数配置(CSC)、输出类型转换参数配置(DTC);动态AIPP可配置输入处理标识(related_input_rank)参数和输入图像最大尺寸(max_src_image_size)参数。99+- aipp_config_path配置文件说明:支持动态AIPP和静态AIPP配置,其中静态AIPP可配置图片裁剪参数配置(Crop)、色域转换参数配置(CSC)、输出类型转换参数配置DTC;动态AIPP可配置输入处理标识related_input_rank参数和输入图像最大尺寸max_src_image_size参数。
100 100 
101- aipp_config_path文件配置示例如下,具体参数说明参见下表。101- aipp_config_path文件配置示例如下,具体参数说明参见下表。
102 102 
@@ -237,7 +237,7 @@
237 <td>INT32</td>237 <td>INT32</td>
238 </tr>238 </tr>
239 <tr>239 <tr>
240- <td rowspan="6">CSC 参数设置</td>240+ <td rowspan="6">CSC参数设置</td>
241 <td>csc_switch</td>241 <td>csc_switch</td>
242 <td>可选属性</td>242 <td>可选属性</td>
243 <td>色域转换开关,静态AIPP配置,true表示开启色域转换,false表示关闭。</td>243 <td>色域转换开关,静态AIPP配置,true表示开启色域转换,false表示关闭。</td>
@@ -20,7 +20,7 @@
20 20
21 考虑一个形状为 $(N,C,*)$的批处理input张量,其中$N$是批处理维度,$C$是通道维度,而$*$表示任意空间维度。21 考虑一个形状为 $(N,C,*)$的批处理input张量,其中$N$是批处理维度,$C$是通道维度,而$*$表示任意空间维度。
22 22 
23- 此操作将input空间维度内的每个滑动kernel_size大小的块展平为形状是$(N,C×\prod(kernel\_size),L)$ 的 3-D output张量的列(即最后一维)。23+ 此操作将input空间维度内的每个滑动kernel_size大小的块展平为形状是$(N,C×\prod(kernel\_size),L)$ 的3-D output张量的列(即最后一维)。
24 24 
25 其中:25 其中:
26 - $C×\prod(kernel\_size)$ 是每个块内的值的数量(一个块有$\prod(kernel\_size)$ 个空间位置,每个空间位置都包含一个$C$ 通道向量),而$L$是这些块的总数:26 - $C×\prod(kernel\_size)$ 是每个块内的值的数量(一个块有$\prod(kernel\_size)$ 个空间位置,每个空间位置都包含一个$C$ 通道向量),而$L$是这些块的总数:
@@ -246,7 +246,7 @@ aclnnStatus aclnnIm2colBackward(
246 <tr>246 <tr>
247 <td>workspace</td>247 <td>workspace</td>
248 <td>输入</td>248 <td>输入</td>
249- <td>在 Device 侧申请的 workspace 内存地址。</td>249+ <td>在Device侧申请的workspace内存地址。</td>
250 </tr>250 </tr>
251 <tr>251 <tr>
252 <td>workspaceSize</td>252 <td>workspaceSize</td>
@@ -256,12 +256,12 @@ aclnnStatus aclnnIm2colBackward(
256 <tr>256 <tr>
257 <td>executor</td>257 <td>executor</td>
258 <td>输入</td>258 <td>输入</td>
259- <td>op 执行器,包含了算子计算流程。</td>259+ <td>op执行器,包含了算子计算流程。</td>
260 </tr>260 </tr>
261 <tr>261 <tr>
262 <td>stream</td>262 <td>stream</td>
263 <td>输入</td>263 <td>输入</td>
264- <td>指定执行任务的 Stream。</td>264+ <td>指定执行任务的Stream。</td>
265 </tr>265 </tr>
266 </tbody>266 </tbody>
267 </table>267 </table>
@@ -22,44 +22,44 @@
22 22 
23 1. 计算缩放比例$height\_scale$和$width\_scale$:23 1. 计算缩放比例$height\_scale$和$width\_scale$:
24 24 
25- $$25+ $$
26- height\_scale = \begin{cases} (y_2 - y_1) \times (image\_height - 1) / (crop\_height - 1), & crop\_height > 1 \\ 0, & crop\_height = 1 \end{cases}26+ height\_scale = \begin{cases} (y_2 - y_1) \times (image\_height - 1) / (crop\_height - 1), & crop\_height > 1 \\ 0, & crop\_height = 1 \end{cases}
27- $$27+ $$
28 28 
29- $$29+ $$
30- width\_scale = \begin{cases} (x_2 - x_1) \times (image\_width - 1) / (crop\_width - 1), & crop\_width > 1 \\ 0, & crop\_width = 1 \end{cases}30+ width\_scale = \begin{cases} (x_2 - x_1) \times (image\_width - 1) / (crop\_width - 1), & crop\_width > 1 \\ 0, & crop\_width = 1 \end{cases}
31- $$31+ $$
32 32
33 2. 将输出坐标$(p_y, p_x)$映射到输入图像坐标$(in\_y, in\_x)$:33 2. 将输出坐标$(p_y, p_x)$映射到输入图像坐标$(in\_y, in\_x)$:
34 34 
35- $$35+ $$
36- in\_y = \begin{cases} y_1 \times (image\_height - 1) + p_y \times height\_scale, & crop\_height > 1 \\ 0.5 \times (y_1 + y_2) \times (image\_height - 1), & crop\_height = 1 \end{cases}36+ in\_y = \begin{cases} y_1 \times (image\_height - 1) + p_y \times height\_scale, & crop\_height > 1 \\ 0.5 \times (y_1 + y_2) \times (image\_height - 1), & crop\_height = 1 \end{cases}
37- $$37+ $$
38 38
39- $$39+ $$
40- in\_x = \begin{cases} x_1 \times (image\_width - 1) + p_x \times width\_scale, & crop\_width > 1 \\ 0.5 \times (x_1 + x_2) \times (image\_width - 1), & crop\_width = 1 \end{cases}40+ in\_x = \begin{cases} x_1 \times (image\_width - 1) + p_x \times width\_scale, & crop\_width > 1 \\ 0.5 \times (x_1 + x_2) \times (image\_width - 1), & crop\_width = 1 \end{cases}
41- $$41+ $$
42 42
43- 若$in\_y < 0$或$in\_y > image\_height - 1$或$in\_x < 0$ 或 $in\_x > image\_width - 1$,则:43+ 若$in\_y < 0$或$in\_y > image\_height - 1$或$in\_x < 0$ 或 $in\_x > image\_width - 1$,则:
44 44
45- $$45+ $$
46- y(i, p_y, p_x, d) = extrapolation\_value46+ y(i, p_y, p_x, d) = extrapolation\_value
47- $$47+ $$
48+ 
49+ 否则,进行双线性插值。令$top = \lfloor in\_y \rfloor$,$bottom = \lceil in\_y \rceil$,$left = \lfloor in\_x \rfloor$,$right = \lceil in\_x \rceil$,$y\_ratio = in\_y - top$,$x\_ratio = in\_x - left$,则:
48 50
49- 否则,进行双线性插值。令$top = \lfloor in\_y \rfloor$,$bottom = \lceil in\_y \rceil$,$left = \lfloor in\_x \rfloor$,$right = \lceil in\_x \rceil$,$y\_ratio = in\_y - top$,$x\_ratio = in\_x - left$,则:51+ $$
50- 52+ \begin{aligned}
51- $$53+ y(i, p_y, p_x, d) =\; & (1 - y\_ratio)(1 - x\_ratio) \cdot x(b, top, left, d) \\
52- \begin{aligned}54+ +\; & (1 - y\_ratio) \cdot x\_ratio \cdot x(b, top, right, d) \\
53- y(i, p_y, p_x, d) =\; & (1 - y\_ratio)(1 - x\_ratio) \cdot x(b, top, left, d) \\55+ +\; & y\_ratio \cdot (1 - x\_ratio) \cdot x(b, bottom, left, d) \\
54- +\; & (1 - y\_ratio) \cdot x\_ratio \cdot x(b, top, right, d) \\56+ +\; & y\_ratio \cdot x\_ratio \cdot x(b, bottom, right, d)
55- +\; & y\_ratio \cdot (1 - x\_ratio) \cdot x(b, bottom, left, d) \\57+ \end{aligned}
56- +\; & y\_ratio \cdot x\_ratio \cdot x(b, bottom, right, d)58+ $$
57- \end{aligned}
58- $$
59 59 
60 - method = nearest时:60 - method = nearest时:
61 61 
62- $height\_scale$、$width\_scale$、$in\_y$、$in\_x$ 的计算方式与 bilinear 方法相同。62+ $height\_scale$、$width\_scale$、$in\_y$、$in\_x$ 的计算方式与bilinear方法相同。
63 63 
64 若$in\_y$或$in\_x$越界(条件同bilinear),则:64 若$in\_y$或$in\_x$越界(条件同bilinear),则:
65 65 
@@ -79,44 +79,44 @@
79 79 
80 1. 首先将归一化框坐标映射为输入图像上的像素坐标,计算裁剪窗口。令$y_{1o} = \lfloor y_1 \times image\_height \rfloor$,$x_{1o} = \lfloor x_1 \times image\_width \rfloor$,$y_{2o} = \lfloor y_2 \times image\_height \rfloor$,$x_{2o} = \lfloor x_2 \times image\_width \rfloor$,则裁剪窗口的高$h$和宽$w$为:80 1. 首先将归一化框坐标映射为输入图像上的像素坐标,计算裁剪窗口。令$y_{1o} = \lfloor y_1 \times image\_height \rfloor$,$x_{1o} = \lfloor x_1 \times image\_width \rfloor$,$y_{2o} = \lfloor y_2 \times image\_height \rfloor$,$x_{2o} = \lfloor x_2 \times image\_width \rfloor$,则裁剪窗口的高$h$和宽$w$为:
81 81 
82- $$82+ $$
83- h = \max(y_{2o} - y_{1o} + 1,\; 1), \quad w = \max(x_{2o} - x_{1o} + 1,\; 1)83+ h = \max(y_{2o} - y_{1o} + 1,\; 1), \quad w = \max(x_{2o} - x_{1o} + 1,\; 1)
84- $$84+ $$
85 85 
86 2. 对于输出位置$(p_y, p_x)$,计算其在裁剪窗口中的浮点索引$r_y$和$r_x$:86 2. 对于输出位置$(p_y, p_x)$,计算其在裁剪窗口中的浮点索引$r_y$和$r_x$:
87 87 
88- $$88+ $$
89- r_y = (p_y + 0.5) \times h / crop\_height - 0.589+ r_y = (p_y + 0.5) \times h / crop\_height - 0.5
90- $$90+ $$
91 91 
92- $$92+ $$
93- r_x = (p_x + 0.5) \times w / crop\_width - 0.593+ r_x = (p_x + 0.5) \times w / crop\_width - 0.5
94- $$94+ $$
95 95 
96 3. 计算插值下标和权重。令$clamp(v, lo, hi) = \max(\min(v, hi), lo)$,则:96 3. 计算插值下标和权重。令$clamp(v, lo, hi) = \max(\min(v, hi), lo)$,则:
97 97 
98- $$98+ $$
99- lower\_y = clamp(\lfloor r_y \rfloor, 0, h - 1), \quad upper\_y = clamp(\lceil r_y \rceil, 0, h - 1)99+ lower\_y = clamp(\lfloor r_y \rfloor, 0, h - 1), \quad upper\_y = clamp(\lceil r_y \rceil, 0, h - 1)
100- $$100+ $$
101 101 
102- $$102+ $$
103- lower\_x = clamp(\lfloor r_x \rfloor, 0, w - 1), \quad upper\_x = clamp(\lceil r_x \rceil, 0, w - 1)103+ lower\_x = clamp(\lfloor r_x \rfloor, 0, w - 1), \quad upper\_x = clamp(\lceil r_x \rceil, 0, w - 1)
104- $$104+ $$
105 105 
106- $$106+ $$
107- y\_ratio = r_y - lower\_y, \quad x\_ratio = r_x - lower\_x107+ y\_ratio = r_y - lower\_y, \quad x\_ratio = r_x - lower\_x
108- $$108+ $$
109 109 
110 4. 双线性插值:110 4. 双线性插值:
111 111 
112- $$112+ $$
113- \begin{aligned}113+ \begin{aligned}
114- y(i, p_y, p_x, d) =\; & (1 - y\_ratio)(1 - x\_ratio) \cdot x(b, y_{1o} + lower\_y, x_{1o} + lower\_x, d) \\114+ y(i, p_y, p_x, d) =\; & (1 - y\_ratio)(1 - x\_ratio) \cdot x(b, y_{1o} + lower\_y, x_{1o} + lower\_x, d) \\
115- +\; & (1 - y\_ratio) \cdot x\_ratio \cdot x(b, y_{1o} + lower\_y, x_{1o} + upper\_x, d) \\115+ +\; & (1 - y\_ratio) \cdot x\_ratio \cdot x(b, y_{1o} + lower\_y, x_{1o} + upper\_x, d) \\
116- +\; & y\_ratio \cdot (1 - x\_ratio) \cdot x(b, y_{1o} + upper\_y, x_{1o} + lower\_x, d) \\116+ +\; & y\_ratio \cdot (1 - x\_ratio) \cdot x(b, y_{1o} + upper\_y, x_{1o} + lower\_x, d) \\
117- +\; & y\_ratio \cdot x\_ratio \cdot x(b, y_{1o} + upper\_y, x_{1o} + upper\_x, d)117+ +\; & y\_ratio \cdot x\_ratio \cdot x(b, y_{1o} + upper\_y, x_{1o} + upper\_x, d)
118- \end{aligned}118+ \end{aligned}
119- $$119+ $$
120 120 
121## 参数说明121## 参数说明
122 122 
@@ -57,22 +57,22 @@
57 - 对input采样时,会根据interpolationMode进行不同处理:57 - 对input采样时,会根据interpolationMode进行不同处理:
58 58 
59 - interpolationMode=0,表示取(x, y)周围四个坐标的加权平均值。59 - interpolationMode=0,表示取(x, y)周围四个坐标的加权平均值。
60- 60+ 
61 $$61 $$
62 output(N, C, H_{out}, W_{out}) = \sum_{i=0}^{2}\sum_{j=0}^{2}{w(i, j)} * {f(x', y')}62 output(N, C, H_{out}, W_{out}) = \sum_{i=0}^{2}\sum_{j=0}^{2}{w(i, j)} * {f(x', y')}
63 $$63 $$
64- 64+ 
65 其中:65 其中:
66 - $f(x', y')$是原图像在$(x', y')$的像素值。66 - $f(x', y')$是原图像在$(x', y')$的像素值。
67 - $w(i, j)$是双线性插值周边4个点的权重,计算公式为:67 - $w(i, j)$是双线性插值周边4个点的权重,计算公式为:
68- 68+ 
69 $$69 $$
70 w(i) = \begin{cases}70 w(i) = \begin{cases}
71 1 - |x'_i - x_i| & |x'_i - x_i| < 1 \\71 1 - |x'_i - x_i| & |x'_i - x_i| < 1 \\
72 0 & otherwise72 0 & otherwise
73 \end{cases}73 \end{cases}
74 $$74 $$
75- 75+ 
76 $$76 $$
77 w(j) = \begin{cases}77 w(j) = \begin{cases}
78 1 - |y'_j - y_j| & |y'_j - y_j| < 1 \\78 1 - |y'_j - y_j| & |y'_j - y_j| < 1 \\
@@ -368,7 +368,7 @@ aclnnStatus aclnnGridSampler2D(
368## 约束说明368## 约束说明
369 369 
370- 参数`input``grid``out`的数据格式只支持(N, C, H, W),当输入其他数据格式时,默认按照(N, C, H, W)格式处理。370- 参数`input``grid``out`的数据格式只支持(N, C, H, W),当输入其他数据格式时,默认按照(N, C, H, W)格式处理。
371-- 输入`input`的(H轴的大小 * W轴的大小) < INT32的最大值。371+- 输入`input`的(H轴的大小 * W轴的大小)< INT32的最大值。
372- 当grid的输入值*图片(长或宽)大于24位的二进制数(16777216)时,采样点可能存在误差,精度可能产生偏差。372- 当grid的输入值*图片(长或宽)大于24位的二进制数(16777216)时,采样点可能存在误差,精度可能产生偏差。
373- 如果grid含有大量超过[-1, 1]范围的数据,使用zeros或者border的填充策略时,计算结果中的值会大量重复。373- 如果grid含有大量超过[-1, 1]范围的数据,使用zeros或者border的填充策略时,计算结果中的值会大量重复。
374- <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term><term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>374- <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term><term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>
@@ -29,33 +29,33 @@
29 其中input、grid、output中的N是一致的,input和output中的C是一致的,grid和output中的$D_{out}$、$H_{out}$、$W_{out}$是一致的,grid最后一维大小为3,表示input像素位置信息为(x, y, z),会将x、y、z的取值范围归一化到[-1, 1]之间。29 其中input、grid、output中的N是一致的,input和output中的C是一致的,grid和output中的$D_{out}$、$H_{out}$、$W_{out}$是一致的,grid最后一维大小为3,表示input像素位置信息为(x, y, z),会将x、y、z的取值范围归一化到[-1, 1]之间。
30 30 
31 - 反归一化的计算公式:31 - 反归一化的计算公式:
32- - alignCorners=true,表示特征值位于像素中心。32+ - alignCorners=true,表示特征值位于像素中心。
33 33 
34- $$34+ $$
35- x' = (grid\_x + 1) / 2 * (D_{in} - 1)35+ x' = (grid\_x + 1) / 2 * (D_{in} - 1)
36- $$36+ $$
37 37 
38- $$38+ $$
39- y' = (grid\_y +1) / 2 * (H_{in} - 1)39+ y' = (grid\_y +1) / 2 * (H_{in} - 1)
40- $$40+ $$
41 41 
42- $$42+ $$
43- z' = (grid\_z +1) / 2 * (W_{in} - 1)43+ z' = (grid\_z +1) / 2 * (W_{in} - 1)
44- $$44+ $$
45 45 
46- - alignCorners=false,表示特征值位于像素的角点。46+ - alignCorners=false,表示特征值位于像素的角点。
47 47 
48- $$48+ $$
49- x' = ((grid\_x +1) * D_{in} - 1) / 249+ x' = ((grid\_x +1) * D_{in} - 1) / 2
50- $$50+ $$
51 51 
52- $$52+ $$
53- y' = ((grid\_y +1) * H_{in} - 1) / 253+ y' = ((grid\_y +1) * H_{in} - 1) / 2
54- $$54+ $$
55 55 
56- $$56+ $$
57- z' = ((grid\_z +1) * W_{in} - 1) / 257+ z' = ((grid\_z +1) * W_{in} - 1) / 2
58- $$58+ $$
59 59 
60 - 对于超出范围的坐标,会根据paddingMode进行不同处理:60 - 对于超出范围的坐标,会根据paddingMode进行不同处理:
61 - paddingMode=0,表示对越界位置用0填充。61 - paddingMode=0,表示对越界位置用0填充。
@@ -52,7 +52,7 @@
52 2. padding_mode对梯度乘子的影响:52 2. padding_mode对梯度乘子的影响:
53 - padding_mode="zeros",`gix_mult`不变53 - padding_mode="zeros",`gix_mult`不变
54 - padding_mode="border",$gix\_mult = gix\_mult × grad\_clip$(坐标在边界外时grad_clip=0,否则=1)54 - padding_mode="border",$gix\_mult = gix\_mult × grad\_clip$(坐标在边界外时grad_clip=0,否则=1)
55- - padding_mode="reflection",$gix\_mult = gix\_mult × grad\_refl × grad\_clip$(grad_refl是反射坐标变换函数对输入坐标的导数,表示反射后输出坐标随输入坐标变化的方向和速率。取值为-1,0,155+ - padding_mode="reflection",$gix\_mult = gix\_mult × grad\_refl × grad\_clip$(grad_refl是反射坐标变换函数对输入坐标的导数,表示反射后输出坐标随输入坐标变化的方向和速率。取值为-1,0,1)
56 56 
57 3. 各插值模式的梯度公式:57 3. 各插值模式的梯度公式:
58 - Bilinear(双线性插值)58 - Bilinear(双线性插值)
@@ -61,10 +61,10 @@
61 61 
62 | 角点 | 坐标$(i_p, j_p)$ | 权重$w_p$ |62 | 角点 | 坐标$(i_p, j_p)$ | 权重$w_p$ |
63 |:------:|:------:|:----------:|63 |:------:|:------:|:----------:|
64- | nw (西北) | $(iy_{nw}, ix_{nw})$ | $(ix_{se} - ix) × (iy_{se} - iy)$ |64+ | nw西北 | $(iy_{nw}, ix_{nw})$ | $(ix_{se} - ix) × (iy_{se} - iy)$ |
65- | ne (东北) | $(iy_{ne}, ix_{ne})$ | $(ix - ix_{sw}) × (iy_{sw} - iy)$ |65+ | ne东北 | $(iy_{ne}, ix_{ne})$ | $(ix - ix_{sw}) × (iy_{sw} - iy)$ |
66- | sw (西南) | $(iy_{sw}, ix_{sw})$ | $(ix_{ne} - ix) × (iy - iy_{ne})$ |66+ | sw西南 | $(iy_{sw}, ix_{sw})$ | $(ix_{ne} - ix) × (iy - iy_{ne})$ |
67- | se (东南) | $(iy_{se}, ix_{se})$ | $(ix - ix_{nw}) × (iy - iy_{nw})$ |67+ | se东南 | $(iy_{se}, ix_{se})$ | $(ix - ix_{nw}) × (iy - iy_{nw})$ |
68 68 
69 其中:69 其中:
70 70 
@@ -79,15 +79,15 @@
79 iy_{se} = floor(iy) + 179 iy_{se} = floor(iy) + 1
80 $$80 $$
81 81 
82- - dx(input 梯度):将上游梯度按权重散射到input对应位置82+ - dx(input梯度):将上游梯度按权重散射到input对应位置
83- 83+ 
84 $$84 $$
85 dx(N, C, i_p, j_p) \mathrel{+}= w_p \cdot grad(N, C, H_{out}, W_{out})85 dx(N, C, i_p, j_p) \mathrel{+}= w_p \cdot grad(N, C, H_{out}, W_{out})
86 $$86 $$
87 87 
88 即对每个输出像素(h, w),将其梯度乘以双线性权重,累加到input的四个相邻像素位置(越界位置不累加)。88 即对每个输出像素(h, w),将其梯度乘以双线性权重,累加到input的四个相邻像素位置(越界位置不累加)。
89- - dgrid(grid 梯度):对(ix, iy)的偏导89+ - dgrid(grid梯度):对(ix, iy)的偏导
90- 90+ 
91 $$91 $$
92 gix = \sum_{c} \left[ -V_{nw} \cdot (iy_{se} - iy) + V_{ne} \cdot (iy_{sw} - iy) - V_{sw} \cdot (iy - iy_{ne}) + V_{se} \cdot (iy - iy_{nw}) \right] \cdot grad(N, C, H_{out}, W_{out})92 gix = \sum_{c} \left[ -V_{nw} \cdot (iy_{se} - iy) + V_{ne} \cdot (iy_{sw} - iy) - V_{sw} \cdot (iy - iy_{ne}) + V_{se} \cdot (iy - iy_{nw}) \right] \cdot grad(N, C, H_{out}, W_{out})
93 $$93 $$
@@ -107,26 +107,26 @@
107 dgrid(N, H_{out}, W_{out}, 1) = giy\_mult \cdot giy107 dgrid(N, H_{out}, W_{out}, 1) = giy\_mult \cdot giy
108 $$108 $$
109 109 
110- - Nearest(最邻近插值)110+ - Nearest(最邻近插值)
111- - dx:将上游梯度直接累加到最近邻位置111+ - dx:将上游梯度直接累加到最近邻位置
112 112 
113 $$113 $$
114 dx(N, C, \text{round}(iy), \text{round}(ix)) \mathrel{+}= grad(N, C, H_{out}, W_{out})114 dx(N, C, \text{round}(iy), \text{round}(ix)) \mathrel{+}= grad(N, C, H_{out}, W_{out})
115 $$115 $$
116 116 
117- - dgrid:最邻近插值对坐标不可导,因此 **dgrid = 0**。117+ - dgrid:最邻近插值对坐标不可导,因此 **dgrid = 0**。
118 118 
119- - Bicubic(双三次插值)119+ - Bicubic(双三次插值)
120- - dx:120+ - dx:
121 121 
122 $$122 $$
123 dx(N, C, iy', ix') \mathrel{+}= grad(N, C, H_{out}, W_{out}) \cdot x\_coeffs[i] \cdot y\_coeffs[j]123 dx(N, C, iy', ix') \mathrel{+}= grad(N, C, H_{out}, W_{out}) \cdot x\_coeffs[i] \cdot y\_coeffs[j]
124 $$124 $$
125 125 
126 其中:126 其中:
127- 127+ 
128 $(ix', iy') = (ix_{nw}-1+i, iy_{nw}-1+j)$,$i,j \in \{0,1,2,3\}$,越界位置根据padding_mode处理。128 $(ix', iy') = (ix_{nw}-1+i, iy_{nw}-1+j)$,$i,j \in \{0,1,2,3\}$,越界位置根据padding_mode处理。
129- 129+ 
130 $$130 $$
131 A = -0.75 \\131 A = -0.75 \\
132 x_0 = x + 1.0 \\132 x_0 = x + 1.0 \\
@@ -168,7 +168,7 @@
168 y\_coeffs[3] = ((A * y_3 - 5* A) * y_3 + 8 * A) * y_3 - 4 * A168 y\_coeffs[3] = ((A * y_3 - 5* A) * y_3 + 8 * A) * y_3 - 4 * A
169 $$169 $$
170 170 
171- - dgrid:171+ - dgrid:
172 172 
173 $$173 $$
174 gix = -\sum_{C}\sum_{i=0}^{3}\sum_{j=0}^{3} V_{ij} \cdot x\_coeffs\_grad[i] \cdot y\_coeffs[j] \cdot grad(N, C, H_{out}, W_{out})174 gix = -\sum_{C}\sum_{i=0}^{3}\sum_{j=0}^{3} V_{ij} \cdot x\_coeffs\_grad[i] \cdot y\_coeffs[j] \cdot grad(N, C, H_{out}, W_{out})
@@ -55,7 +55,7 @@
55 2. paddingMode对梯度乘子的影响:55 2. paddingMode对梯度乘子的影响:
56 - paddingMode="zeros",`gix_mult`不变56 - paddingMode="zeros",`gix_mult`不变
57 - paddingMode="border",$gix\_mult = gix\_mult × grad\_clip$(坐标在边界外时grad_clip=0,否则=1)57 - paddingMode="border",$gix\_mult = gix\_mult × grad\_clip$(坐标在边界外时grad_clip=0,否则=1)
58- - paddingMode="reflection",$gix\_mult = gix\_mult × grad\_refl × grad\_clip$(grad_refl是反射坐标变换函数对输入坐标的导数,表示反射后输出坐标随输入坐标变化的方向和速率。取值为-1,0,158+ - paddingMode="reflection",$gix\_mult = gix\_mult × grad\_refl × grad\_clip$(grad_refl是反射坐标变换函数对输入坐标的导数,表示反射后输出坐标随输入坐标变化的方向和速率。取值为-1,0,1)
59 59 
60 3. 各插值模式的梯度公式:60 3. 各插值模式的梯度公式:
61 - Bilinear(双线性插值)61 - Bilinear(双线性插值)
@@ -82,15 +82,15 @@
82 iy_{se} = floor(iy) + 182 iy_{se} = floor(iy) + 1
83 $$83 $$
84 84 
85- - dx(input 梯度):将上游梯度按权重散射到input对应位置85+ - dx(input梯度):将上游梯度按权重散射到input对应位置
86- 86+ 
87 $$87 $$
88 dx(N, C, i_p, j_p) \mathrel{+}= w_p \cdot grad(N, C, H_{out}, W_{out})88 dx(N, C, i_p, j_p) \mathrel{+}= w_p \cdot grad(N, C, H_{out}, W_{out})
89 $$89 $$
90 90 
91 即对每个输出像素(h, w),将其梯度乘以双线性权重,累加到input的四个相邻像素位置(越界位置不累加)。91 即对每个输出像素(h, w),将其梯度乘以双线性权重,累加到input的四个相邻像素位置(越界位置不累加)。
92- - dgrid(grid 梯度):对(ix, iy)的偏导92+ - dgrid(grid梯度):对(ix, iy)的偏导
93- 93+ 
94 $$94 $$
95 gix = \sum_{c} \left[ -V_{nw} \cdot (iy_{se} - iy) + V_{ne} \cdot (iy_{sw} - iy) - V_{sw} \cdot (iy - iy_{ne}) + V_{se} \cdot (iy - iy_{nw}) \right] \cdot grad(N, C, H_{out}, W_{out})95 gix = \sum_{c} \left[ -V_{nw} \cdot (iy_{se} - iy) + V_{ne} \cdot (iy_{sw} - iy) - V_{sw} \cdot (iy - iy_{ne}) + V_{se} \cdot (iy - iy_{nw}) \right] \cdot grad(N, C, H_{out}, W_{out})
96 $$96 $$
@@ -100,6 +100,7 @@
100 $$100 $$
101 101 
102 其中 $V_p = input(N, C, i_p, j_p)$(仅当角点在边界内时参与计算)。102 其中 $V_p = input(N, C, i_p, j_p)$(仅当角点在边界内时参与计算)。
103+ 
103 - 最终:104 - 最终:
104 105 
105 $$106 $$
@@ -110,26 +111,27 @@
110 dgrid(N, H_{out}, W_{out}, 1) = giy\_mult \cdot giy111 dgrid(N, H_{out}, W_{out}, 1) = giy\_mult \cdot giy
111 $$112 $$
112 113 
113- - Nearest(最邻近插值)114+ - Nearest(最邻近插值)
114- - dx:将上游梯度直接累加到最近邻位置
115 115 
116- $$116+ - dx:将上游梯度直接累加到最近邻位置
117- dx(N, C, \text{round}(iy), \text{round}(ix)) \mathrel{+}= grad(N, C, H_{out}, W_{out})
118- $$
119 117 
120- - dgrid:最邻近插值对坐标不可导,因此 **dgrid = 0**。118+ $$
119+ dx(N, C, \text{round}(iy), \text{round}(ix)) \mathrel{+}= grad(N, C, H_{out}, W_{out})
120+ $$
121 121 
122- - Bicubic(双三次插值122+ - dgrid:最邻近插值对坐标不可导,因此 **dgrid = 0**。
123- - dx:123+ 
124+ - Bicubic(双三次插值)
125+ - dx:
124 126 
125 $$127 $$
126 dx(N, C, iy', ix') \mathrel{+}= grad(N, C, H_{out}, W_{out}) \cdot x\_coeffs[i] \cdot y\_coeffs[j]128 dx(N, C, iy', ix') \mathrel{+}= grad(N, C, H_{out}, W_{out}) \cdot x\_coeffs[i] \cdot y\_coeffs[j]
127 $$129 $$
128 130 
129 其中:131 其中:
130- 132+ 
131 $(ix', iy') = (ix_{nw}-1+i, iy_{nw}-1+j)$,$i,j \in \{0,1,2,3\}$,越界位置根据paddingMode处理。133 $(ix', iy') = (ix_{nw}-1+i, iy_{nw}-1+j)$,$i,j \in \{0,1,2,3\}$,越界位置根据paddingMode处理。
132- 134+ 
133 $$135 $$
134 A = -0.75 \\136 A = -0.75 \\
135 x_0 = x + 1.0 \\137 x_0 = x + 1.0 \\
@@ -171,7 +173,7 @@
171 y\_coeffs[3] = ((A * y_3 - 5* A) * y_3 + 8 * A) * y_3 - 4 * A173 y\_coeffs[3] = ((A * y_3 - 5* A) * y_3 + 8 * A) * y_3 - 4 * A
172 $$174 $$
173 175 
174- - dgrid:176+ - dgrid:
175 177 
176 $$178 $$
177 gix = -\sum_{C}\sum_{i=0}^{3}\sum_{j=0}^{3} V_{ij} \cdot x\_coeffs\_grad[i] \cdot y\_coeffs[j] \cdot grad(N, C, H_{out}, W_{out})179 gix = -\sum_{C}\sum_{i=0}^{3}\sum_{j=0}^{3} V_{ij} \cdot x\_coeffs\_grad[i] \cdot y\_coeffs[j] \cdot grad(N, C, H_{out}, W_{out})
@@ -55,7 +55,7 @@
55 2. padding_mode对梯度乘子的影响:55 2. padding_mode对梯度乘子的影响:
56 - padding_mode="zeros",`gix_mult`不变56 - padding_mode="zeros",`gix_mult`不变
57 - padding_mode="border",$gix\_mult = gix\_mult × grad\_clip$(坐标在边界外时grad_clip=0,否则=1)57 - padding_mode="border",$gix\_mult = gix\_mult × grad\_clip$(坐标在边界外时grad_clip=0,否则=1)
58- - padding_mode="reflection",$gix\_mult = gix\_mult × grad\_refl × grad\_clip$(grad_refl是反射坐标变换函数对输入坐标的导数,表示反射后输出坐标随输入坐标变化的方向和速率。取值为-1,0,158+ - padding_mode="reflection",$gix\_mult = gix\_mult × grad\_refl × grad\_clip$(grad_refl是反射坐标变换函数对输入坐标的导数,表示反射后输出坐标随输入坐标变化的方向和速率。取值为-1,0,1)
59 59 
60 3. 各插值模式的梯度公式:60 3. 各插值模式的梯度公式:
61 - Bilinear(三线性插值,Trilinear)61 - Bilinear(三线性插值,Trilinear)
@@ -64,14 +64,14 @@
64 64 
65 | 角点 | 坐标$(i_p, j_p, k_p)$ | 权重$w_p$ |65 | 角点 | 坐标$(i_p, j_p, k_p)$ | 权重$w_p$ |
66 |:------:|:------:|:----------:|66 |:------:|:------:|:----------:|
67- | tnw (顶-北-西) | $(⌊iz⌋, ⌊iy⌋, ⌊ix⌋)$ | $(ix_{bse} - ix) × (iy_{bse} - iy) × (iz_{bse} - iz)$ |67+ | tnw顶-北-西 | $(⌊iz⌋, ⌊iy⌋, ⌊ix⌋)$ | $(ix_{bse} - ix) × (iy_{bse} - iy) × (iz_{bse} - iz)$ |
68- | tne (顶-北-东) | $(⌊iz⌋, ⌊iy⌋, ⌊ix⌋+1)$ | $(ix - ix_{bsw}) × (iy_{bsw} - iy) × (iz_{bsw} - iz)$ |68+ | tne顶-北-东 | $(⌊iz⌋, ⌊iy⌋, ⌊ix⌋+1)$ | $(ix - ix_{bsw}) × (iy_{bsw} - iy) × (iz_{bsw} - iz)$ |
69- | tsw (顶-南-西) | $(⌊iz⌋, ⌊iy⌋+1, ⌊ix⌋)$ | $(ix_{bne} - ix) × (iy - iy_{bne}) × (iz_{bne} - iz)$ |69+ | tsw顶-南-西 | $(⌊iz⌋, ⌊iy⌋+1, ⌊ix⌋)$ | $(ix_{bne} - ix) × (iy - iy_{bne}) × (iz_{bne} - iz)$ |
70- | tse (顶-南-东) | $(⌊iz⌋, ⌊iy⌋+1, ⌊ix⌋+1)$ | $(ix - ix_{bnw}) × (iy - iy_{bnw}) × (iz_{bnw} - iz)$ |70+ | tse顶-南-东 | $(⌊iz⌋, ⌊iy⌋+1, ⌊ix⌋+1)$ | $(ix - ix_{bnw}) × (iy - iy_{bnw}) × (iz_{bnw} - iz)$ |
71- | bnw (底-北-西) | $(⌊iz⌋+1, ⌊iy⌋, ⌊ix⌋)$ | $(ix_{tse} - ix) × (iy_{tse} - iy) × (iz - iz_{tse})$|71+ | bnw底-北-西 | $(⌊iz⌋+1, ⌊iy⌋, ⌊ix⌋)$ | $(ix_{tse} - ix) × (iy_{tse} - iy) × (iz - iz_{tse})$|
72- | bne (底-北-东) | $(⌊iz⌋+1, ⌊iy⌋, ⌊ix⌋+1)$ | $(ix - ix_{tsw}) × (iy_{tsw} - iy) × (iz - iz_{tsw})$ |72+ | bne底-北-东 | $(⌊iz⌋+1, ⌊iy⌋, ⌊ix⌋+1)$ | $(ix - ix_{tsw}) × (iy_{tsw} - iy) × (iz - iz_{tsw})$ |
73- | bsw (底-南-西) | $(⌊iz⌋+1, ⌊iy⌋+1, ⌊ix⌋)$ | $(ix_{tne} - ix) × (iy - iy_{tne}) × (iz - iz_{tne})$ |73+ | bsw底-南-西 | $(⌊iz⌋+1, ⌊iy⌋+1, ⌊ix⌋)$ | $(ix_{tne} - ix) × (iy - iy_{tne}) × (iz - iz_{tne})$ |
74- | bse (底-南-东) | $(⌊iz⌋+1, ⌊iy⌋+1, ⌊ix⌋+1)$ | $(ix - ix_{tnw}) × (iy - iy_{tnw}) × (iz - iz_{tnw})$ |74+ | bse底-南-东 | $(⌊iz⌋+1, ⌊iy⌋+1, ⌊ix⌋+1)$ | $(ix - ix_{tnw}) × (iy - iy_{tnw}) × (iz - iz_{tnw})$ |
75 75
76 其中:76 其中:
77 77 
@@ -123,15 +123,15 @@
123 iz_{bse} = iz_{tnw} + 1 \\123 iz_{bse} = iz_{tnw} + 1 \\
124 $$124 $$
125 125 
126- - dx(input 梯度):将上游梯度按三线性权重散射到input对应位置126+ - dx(input梯度):将上游梯度按三线性权重散射到input对应位置
127 127
128 $$128 $$
129 dx(N, C, i_p, j_p, k_p) \mathrel{+}= w_p \cdot grad(N, C, D_{out}, H_{out}, W_{out})129 dx(N, C, i_p, j_p, k_p) \mathrel{+}= w_p \cdot grad(N, C, D_{out}, H_{out}, W_{out})
130 $$130 $$
131 131 
132 即对每个输出像素(d, h, w),将其梯度乘以三线性权重,累加到input的8个相邻体素位置(越界位置不累加)。132 即对每个输出像素(d, h, w),将其梯度乘以三线性权重,累加到input的8个相邻体素位置(越界位置不累加)。
133- - dgrid(grid 梯度):对(ix, iy, iz)的偏导133+ - dgrid(grid梯度):对(ix, iy, iz)的偏导
134- 134+ 
135 $$135 $$
136 gix = \sum_{c} \left[ -V_{tnw} \cdot (iy_{bse}-iy)(iz_{bse}-iz) + V_{tne} \cdot (iy_{bsw}-iy)(iz_{bsw}-iz) - V_{tsw} \cdot (iy-iy_{bne})(iz_{bne}-iz) + V_{tse} \cdot (iy-iy_{bnw})(iz_{bnw}-iz) - V_{bnw} \cdot (iy_{tse}-iy)(iz-iz_{tse}) + V_{bne} \cdot (iy_{tsw}-iy)(iz-iz_{tsw}) - V_{bsw} \cdot (iy-iy_{tne})(iz-iz_{tne}) + V_{bse} \cdot (iy-iy_{tnw})(iz-iz_{tnw}) \right] \cdot grad(N, C, D_{out}, H_{out}, W_{out})136 gix = \sum_{c} \left[ -V_{tnw} \cdot (iy_{bse}-iy)(iz_{bse}-iz) + V_{tne} \cdot (iy_{bsw}-iy)(iz_{bsw}-iz) - V_{tsw} \cdot (iy-iy_{bne})(iz_{bne}-iz) + V_{tse} \cdot (iy-iy_{bnw})(iz_{bnw}-iz) - V_{bnw} \cdot (iy_{tse}-iy)(iz-iz_{tse}) + V_{bne} \cdot (iy_{tsw}-iy)(iz-iz_{tsw}) - V_{bsw} \cdot (iy-iy_{tne})(iz-iz_{tne}) + V_{bse} \cdot (iy-iy_{tnw})(iz-iz_{tnw}) \right] \cdot grad(N, C, D_{out}, H_{out}, W_{out})
137 $$137 $$
@@ -159,14 +159,14 @@
159 dgrid(N, D_{out}, H_{out}, W_{out}, 2) = giy\_mult \cdot giz159 dgrid(N, D_{out}, H_{out}, W_{out}, 2) = giy\_mult \cdot giz
160 $$160 $$
161 161 
162- - Nearest(最邻近插值)162+ - Nearest(最邻近插值)
163- - dx:将上游梯度直接累加到最近邻位置163+ - dx:将上游梯度直接累加到最近邻位置
164 164 
165 $$165 $$
166 dx(N, C, \text{round}(iz), \text{round}(iy), \text{round}(ix)) \mathrel{+}= grad(N, C, D_{out}, H_{out}, W_{out})166 dx(N, C, \text{round}(iz), \text{round}(iy), \text{round}(ix)) \mathrel{+}= grad(N, C, D_{out}, H_{out}, W_{out})
167 $$167 $$
168 168 
169- - dgrid:最邻近插值对坐标不可导,因此 **dgrid = 0**。169+ - dgrid:最邻近插值对坐标不可导,因此 **dgrid = 0**。
170 170 
171## 参数说明171## 参数说明
172 172 
@@ -57,7 +57,7 @@
57 2. padding_mode对梯度乘子的影响:57 2. padding_mode对梯度乘子的影响:
58 - paddingMode="zeros",`gix_mult`不变58 - paddingMode="zeros",`gix_mult`不变
59 - paddingMode="border",$gix\_mult = gix\_mult × grad\_clip$(坐标在边界外时grad_clip=0,否则=1)59 - paddingMode="border",$gix\_mult = gix\_mult × grad\_clip$(坐标在边界外时grad_clip=0,否则=1)
60- - paddingMode="reflection",$gix\_mult = gix\_mult × grad\_refl × grad\_clip$(grad_refl是反射坐标变换函数对输入坐标的导数,表示反射后输出坐标随输入坐标变化的方向和速率。取值为-1,0,160+ - paddingMode="reflection",$gix\_mult = gix\_mult × grad\_refl × grad\_clip$(grad_refl是反射坐标变换函数对输入坐标的导数,表示反射后输出坐标随输入坐标变化的方向和速率。取值为-1,0,1)
61 61 
62 3. 各插值模式的梯度公式:62 3. 各插值模式的梯度公式:
63 - Bilinear63 - Bilinear
@@ -66,14 +66,14 @@
66 66 
67 | 角点 | 坐标$(i_p, j_p, k_p)$ | 权重$w_p$ |67 | 角点 | 坐标$(i_p, j_p, k_p)$ | 权重$w_p$ |
68 |:------:|:------:|:----------:|68 |:------:|:------:|:----------:|
69- | tnw (顶-北-西) | $(⌊iz⌋, ⌊iy⌋, ⌊ix⌋)$ | $(ix_{bse} - ix) × (iy_{bse} - iy) × (iz_{bse} - iz)$ |69+ | tnw顶-北-西 | $(⌊iz⌋, ⌊iy⌋, ⌊ix⌋)$ | $(ix_{bse} - ix) × (iy_{bse} - iy) × (iz_{bse} - iz)$ |
70- | tne (顶-北-东) | $(⌊iz⌋, ⌊iy⌋, ⌊ix⌋+1)$ | $(ix - ix_{bsw}) × (iy_{bsw} - iy) × (iz_{bsw} - iz)$ |70+ | tne顶-北-东 | $(⌊iz⌋, ⌊iy⌋, ⌊ix⌋+1)$ | $(ix - ix_{bsw}) × (iy_{bsw} - iy) × (iz_{bsw} - iz)$ |
71- | tsw (顶-南-西) | $(⌊iz⌋, ⌊iy⌋+1, ⌊ix⌋)$ | $(ix_{bne} - ix) × (iy - iy_{bne}) × (iz_{bne} - iz)$ |71+ | tsw顶-南-西 | $(⌊iz⌋, ⌊iy⌋+1, ⌊ix⌋)$ | $(ix_{bne} - ix) × (iy - iy_{bne}) × (iz_{bne} - iz)$ |
72- | tse (顶-南-东) | $(⌊iz⌋, ⌊iy⌋+1, ⌊ix⌋+1)$ | $(ix - ix_{bnw}) × (iy - iy_{bnw}) × (iz_{bnw} - iz)$ |72+ | tse顶-南-东 | $(⌊iz⌋, ⌊iy⌋+1, ⌊ix⌋+1)$ | $(ix - ix_{bnw}) × (iy - iy_{bnw}) × (iz_{bnw} - iz)$ |
73- | bnw (底-北-西) | $(⌊iz⌋+1, ⌊iy⌋, ⌊ix⌋)$ | $(ix_{tse} - ix) × (iy_{tse} - iy) × (iz - iz_{tse})$|73+ | bnw底-北-西 | $(⌊iz⌋+1, ⌊iy⌋, ⌊ix⌋)$ | $(ix_{tse} - ix) × (iy_{tse} - iy) × (iz - iz_{tse})$|
74- | bne (底-北-东) | $(⌊iz⌋+1, ⌊iy⌋, ⌊ix⌋+1)$ | $(ix - ix_{tsw}) × (iy_{tsw} - iy) × (iz - iz_{tsw})$ |74+ | bne底-北-东 | $(⌊iz⌋+1, ⌊iy⌋, ⌊ix⌋+1)$ | $(ix - ix_{tsw}) × (iy_{tsw} - iy) × (iz - iz_{tsw})$ |
75- | bsw (底-南-西) | $(⌊iz⌋+1, ⌊iy⌋+1, ⌊ix⌋)$ | $(ix_{tne} - ix) × (iy - iy_{tne}) × (iz - iz_{tne})$ |75+ | bsw底-南-西 | $(⌊iz⌋+1, ⌊iy⌋+1, ⌊ix⌋)$ | $(ix_{tne} - ix) × (iy - iy_{tne}) × (iz - iz_{tne})$ |
76- | bse (底-南-东) | $(⌊iz⌋+1, ⌊iy⌋+1, ⌊ix⌋+1)$ | $(ix - ix_{tnw}) × (iy - iy_{tnw}) × (iz - iz_{tnw})$ |76+ | bse底-南-东 | $(⌊iz⌋+1, ⌊iy⌋+1, ⌊ix⌋+1)$ | $(ix - ix_{tnw}) × (iy - iy_{tnw}) × (iz - iz_{tnw})$ |
77 77
78 其中:78 其中:
79 79 
@@ -125,15 +125,15 @@
125 iz_{bse} = iz_{tnw} + 1 \\125 iz_{bse} = iz_{tnw} + 1 \\
126 $$126 $$
127 127 
128- - dx(input 梯度):将上游梯度按三线性权重散射到input对应位置128+ - dx(input梯度):将上游梯度按三线性权重散射到input对应位置
129- 129+ 
130 $$130 $$
131 dx(N, C, i_p, j_p, k_p) \mathrel{+}= w_p \cdot grad(N, C, D_{out}, H_{out}, W_{out})131 dx(N, C, i_p, j_p, k_p) \mathrel{+}= w_p \cdot grad(N, C, D_{out}, H_{out}, W_{out})
132 $$132 $$
133 133 
134 即对每个输出像素(d, h, w),将其梯度乘以三线性权重,累加到input的8个相邻体素位置(越界位置不累加)。134 即对每个输出像素(d, h, w),将其梯度乘以三线性权重,累加到input的8个相邻体素位置(越界位置不累加)。
135- - dgrid(grid 梯度):对(ix, iy, iz)的偏导135+ - dgrid(grid梯度):对(ix, iy, iz)的偏导
136- 136+ 
137 $$137 $$
138 gix = \sum_{c} \left[ -V_{tnw} \cdot (iy_{bse}-iy)(iz_{bse}-iz) + V_{tne} \cdot (iy_{bsw}-iy)(iz_{bsw}-iz) - V_{tsw} \cdot (iy-iy_{bne})(iz_{bne}-iz) + V_{tse} \cdot (iy-iy_{bnw})(iz_{bnw}-iz) - V_{bnw} \cdot (iy_{tse}-iy)(iz-iz_{tse}) + V_{bne} \cdot (iy_{tsw}-iy)(iz-iz_{tsw}) - V_{bsw} \cdot (iy-iy_{tne})(iz-iz_{tne}) + V_{bse} \cdot (iy-iy_{tnw})(iz-iz_{tnw}) \right] \cdot grad(N, C, D_{out}, H_{out}, W_{out})138 gix = \sum_{c} \left[ -V_{tnw} \cdot (iy_{bse}-iy)(iz_{bse}-iz) + V_{tne} \cdot (iy_{bsw}-iy)(iz_{bsw}-iz) - V_{tsw} \cdot (iy-iy_{bne})(iz_{bne}-iz) + V_{tse} \cdot (iy-iy_{bnw})(iz_{bnw}-iz) - V_{bnw} \cdot (iy_{tse}-iy)(iz-iz_{tse}) + V_{bne} \cdot (iy_{tsw}-iy)(iz-iz_{tsw}) - V_{bsw} \cdot (iy-iy_{tne})(iz-iz_{tne}) + V_{bse} \cdot (iy-iy_{tnw})(iz-iz_{tnw}) \right] \cdot grad(N, C, D_{out}, H_{out}, W_{out})
139 $$139 $$
@@ -161,14 +161,14 @@
161 dgrid(N, D_{out}, H_{out}, W_{out}, 2) = giy\_mult \cdot giz161 dgrid(N, D_{out}, H_{out}, W_{out}, 2) = giy\_mult \cdot giz
162 $$162 $$
163 163 
164- - Nearest(最邻近插值)164+ - Nearest(最邻近插值)
165- - dx:将上游梯度直接累加到最近邻位置165+ - dx:将上游梯度直接累加到最近邻位置
166 166 
167 $$167 $$
168 dx(N, C, \text{round}(iz), \text{round}(iy), \text{round}(ix)) \mathrel{+}= grad(N, C, D_{out}, H_{out}, W_{out})168 dx(N, C, \text{round}(iz), \text{round}(iy), \text{round}(ix)) \mathrel{+}= grad(N, C, D_{out}, H_{out}, W_{out})
169 $$169 $$
170 170 
171- - dgrid:最邻近插值对坐标不可导,因此 **dgrid = 0**。171+ - dgrid:最邻近插值对坐标不可导,因此 **dgrid = 0**。
172 172 
173## 函数原型173## 函数原型
174 174 
@@ -13,19 +13,19 @@
13 13 
14## 功能说明14## 功能说明
15 15 
16-- 算子功能:该算子根据张量 offsets 对输入张量 images 进行偏移变换,生成warp_images张量。16+- 算子功能:该算子根据张量offsets对输入张量images进行偏移变换,生成warp_images张量。
17 17 
18## 参数说明18## 参数说明
19 19 
20| 参数名 | 输入/输出/属性 | 描述 | 数据类型 | 数据格式 |20| 参数名 | 输入/输出/属性 | 描述 | 数据类型 | 数据格式 |
21|---------------------|----------|----------------------------------------------------------------------------|---------------------|------|21|---------------------|----------|----------------------------------------------------------------------------|---------------------|------|
22-| images | 输入 | 四维 Tensor,其 shape 为(batch, image_height, image_width, 3)。 | UINT8、FLOAT16、FLOAT | ND |22+| images | 输入 | 四维Tensor,其shape为(batch, image_height, image_width, 3)。 | UINT8、FLOAT16、FLOAT | ND |
23-| offsets | 输入 | 四维 Tensor,其 shape 为(batch, 4, new_height, new_width)。 | FLOAT、INT32 | ND |23+| offsets | 输入 | 四维Tensor,其shape为(batch, 4, new_height, new_width)。 | FLOAT、INT32 | ND |
24-| warp_images | 输出 | 五维 Tensor,其 shape 为(batch, 4, new_height, new_width, 3),数据类型与输入 images 相同。 | UINT8、FLOAT16、FLOAT | ND |24+| warp_images | 输出 | 五维Tensor,其shape为(batch, 4, new_height, new_width, 3),数据类型与输入images相同。 | UINT8、FLOAT16、FLOAT | ND |
25 25 
26## 约束说明26## 约束说明
27 27 
28-- 当 offsets 的数据类型为 INT32 时,images 必须为 FLOAT16。28+- 当offsets的数据类型为INT32时,images必须为FLOAT16。
29 29 
30## 调用说明30## 调用说明
31 31 
@@ -11,7 +11,6 @@
11| <term>Atlas 推理系列产品</term> | × |11| <term>Atlas 推理系列产品</term> | × |
12| <term>Atlas 训练系列产品</term> | × |12| <term>Atlas 训练系列产品</term> | × |
13 13 
14- 
15## 功能说明14## 功能说明
16 15 
17- 算子功能:对边界框进行非极大值抑制(NMS)处理,输出经过NMS后的选中的框、索引以及掩码。用于目标检测后处理,去除重复检测的框。16- 算子功能:对边界框进行非极大值抑制(NMS)处理,输出经过NMS后的选中的框、索引以及掩码。用于目标检测后处理,去除重复检测的框。
@@ -20,19 +19,17 @@
20 19 
21| 参数名 | 输入/输出/属性 | 描述 | 数据类型 | 数据格式 |20| 参数名 | 输入/输出/属性 | 描述 | 数据类型 | 数据格式 |
22|:-------------------------|:----------:|:-----|:-----|:----:|21|:-------------------------|:----------:|:-----|:-----|:----:|
23-| box_scores | 输入 | 二维 Tensor,其 shape (num_boxes, 5),其中 5 表示 [y1, x1, y2, x2, score],num_boxes不超过39936。 | FLOAT16、FLOAT、BF16 | ND |22+| box_scores | 输入 | 二维Tensor,其shape为(num_boxes, 5),其中5表示[y1, x1, y2, x2, score],num_boxes不超过39936。 | FLOAT16、FLOAT、BF16 | ND |
24-| iou_threshold | 属性 | 浮点数,表示用于判断候选框是否在 IoU(交并比)上重叠过多的阈值。默认值为 0.5。 | FLOAT | - |23+| iou_threshold | 属性 | 浮点数,表示用于判断候选框是否在IoU(交并比)上重叠过多的阈值。默认值为0.5。 | FLOAT | - |
25-| selected_boxes | 输出 | 二维 Tensor,其 shape (num_boxes, 5),和原输入的box_scores一样。 | FLOAT16、FLOAT、BF16 | ND |24+| selected_boxes | 输出 | 二维Tensor,其shape为(num_boxes, 5),和原输入的box_scores一样。 | FLOAT16、FLOAT、BF16 | ND |
26-| selected_idx | 输出 | 一维 Tensor,其 shape (num_boxes),表示0到 num_boxes - 1 的序列数。 | INT32 | ND |25+| selected_idx | 输出 | 一维Tensor,其shape为(num_boxes),表示0到num_boxes - 1的序列数。 | INT32 | ND |
27-| selected_mask | 输出 | 一维 Tensor,其 shape (num_boxes),表示目标框的掩码情况。 | UINT8 | ND |26+| selected_mask | 输出 | 一维Tensor,其shape为(num_boxes),表示目标框的掩码情况。 | UINT8 | ND |
28- 
29 27 
30## 约束说明28## 约束说明
31 29 
32-- 输入的 box_scores 最后一维必须为 5,依次表示 [y1, x1, y2, x2, score]。30+- 输入的box_scores最后一维必须为5,依次表示 [y1, x1, y2, x2, score]。
33-- 输出 selected_boxes、selected_idx、selected_mask batch 维度与输入 box_scores batch 维度保持一致。31+- 输出selected_boxes、selected_idx、selected_mask的batch维度与输入box_scores的batch维度保持一致。
34-- 在 FLOAT16 BF16 场景下,算子进行排序和计算对比标杆可能会引入计算误差。32+- 在FLOAT16或BF16场景下,算子进行排序和计算对比标杆可能会引入计算误差。
35- 
36 33 
37## 调用说明34## 调用说明
38 35 
@@ -19,17 +19,17 @@
19 19 
20| 参数名 | 输入/输出/属性 | 描述 | 数据类型 | 数据格式 |20| 参数名 | 输入/输出/属性 | 描述 | 数据类型 | 数据格式 |
21|-------------------|----------|-----------------------------------------------------------------------------------|---------------|------|21|-------------------|----------|-----------------------------------------------------------------------------------|---------------|------|
22-| boxes | 输入 | 二维 Tensor,其 shape (num_boxes, 4)。输入格式为 (x1, y1, x2, y2),并且要求 x1 < x2 y1 < y2。 | FLOAT16、FLOAT | ND |22+| boxes | 输入 | 二维Tensor,其shape为(num_boxes, 4)。输入格式为(x1, y1, x2, y2),并且要求x1 < x2且y1 < y2。 | FLOAT16、FLOAT | ND |
23-| scores | 输入 | 一维 Tensor,其 shape (num_boxes)。表示与每个候选框对应的分数(与 boxes 的每一行一一对应)。 | FLOAT16、FLOAT | ND |23+| scores | 输入 | 一维Tensor,其shape为(num_boxes)。表示与每个候选框对应的分数(与boxes的每一行一一对应)。 | FLOAT16、FLOAT | ND |
24-| max_output_size | 输入 | 标量整数 Tensor,表示非极大值抑制(Non-Maximum Suppression)最多选择的候选框数量。 | INT32 | ND |24+| max_output_size | 输入 | 标量整数Tensor,表示非极大值抑制(Non-Maximum Suppression)最多选择的候选框数量。 | INT32 | ND |
25-| iou_threshold | 输入 | 标量浮点 Tensor,表示用于判断候选框是否在 IoU(交并比)上重叠过多的阈值。 | FLOAT16、FLOAT | ND |25+| iou_threshold | 输入 | 标量浮点Tensor,表示用于判断候选框是否在IoU(交并比)上重叠过多的阈值。 | FLOAT16、FLOAT | ND |
26-| score_threshold | 输入 | 标量浮点 Tensor,表示用于根据分数移除候选框的阈值。 | FLOAT16、FLOAT | ND |26+| score_threshold | 输入 | 标量浮点Tensor,表示用于根据分数移除候选框的阈值。 | FLOAT16、FLOAT | ND |
27-| offset | 可选属性 | • 可选整数。<br>• 默认值为 0。 | INT | - |27+| offset | 可选属性 | • 可选整数。<br>• 默认值为0。 | INT | - |
28-| selected_indices | 输出 | 一维整数 Tensor,shape (M),表示从输入 boxes Tensor 中选中的索引,其中 M <= max_output_size。 | INT32 | ND |28+| selected_indices | 输出 | 一维整数Tensor,shape为(M),表示从输入boxes Tensor中选中的索引,其中M <= max_output_size。 | INT32 | ND |
29 29 
30## 约束说明30## 约束说明
31 31 
32-- 输入的 boxes scores 必须是 float 类型。32+- 输入的boxes和scores必须是float类型。
33 33 
34## 调用说明34## 调用说明
35 35 
@@ -16,44 +16,45 @@
16- 算子功能:实现光栅化计算。根据给定的三维空间中的点和面,获取屏幕中每个像素点的最小深度及其对应的面片索引,并计算该面片的重心坐标透视矫正插值。16- 算子功能:实现光栅化计算。根据给定的三维空间中的点和面,获取屏幕中每个像素点的最小深度及其对应的面片索引,并计算该面片的重心坐标透视矫正插值。
17 17 
18- 计算公式:18- 计算公式:
19+ 
19 $findices$记录每个像素点最小深度对应的面索引,$barycentric$记录每个顶点相对于$findices$中记录的面的重心坐标透视矫正插值。20 $findices$记录每个像素点最小深度对应的面索引,$barycentric$记录每个顶点相对于$findices$中记录的面的重心坐标透视矫正插值。
20 计算过程中使用的Z-Buffer记录每个像素点$(x, y)$的最小深度$z_{\min}(x, y)$以及该深度对应的三角形面片索引$\text{face\_idx}(x, y)$。21 计算过程中使用的Z-Buffer记录每个像素点$(x, y)$的最小深度$z_{\min}(x, y)$以及该深度对应的三角形面片索引$\text{face\_idx}(x, y)$。
21 22
22 计算过程如下:23 计算过程如下:
23 对空间中的每个三角形面片$f$:24 对空间中的每个三角形面片$f$:
24 25
25- 1. 将$f$的三个顶点坐标$v_0$,$v_1$,$v_2$转换为屏幕坐标$v_{s0}$,$v_{s1}$,$v_{s2}$26+ 1. 将$f$的三个顶点坐标$(v_0, v_1, v_2)$转换为屏幕坐标$(v_{s0}, v_{s1}, v_{s2})$
26- 2. 根据$v_{s0}$,$v_{s1}$,$v_{s2}$计算包围$f$的矩形范围27+ 2. 根据$(v_{s0}, v_{s1}, v_{s2})$计算包围$f$的矩形范围
27 3. 对矩形内每个像素点$v_i = (x_i, y_i)$,执行以下操作:28 3. 对矩形内每个像素点$v_i = (x_i, y_i)$,执行以下操作:
28- 29+ 
29 a. 计算像素中心坐标$v_c$ 30 a. 计算像素中心坐标$v_c$
30 b. 计算$v_c$相对于三角形$f$的重心坐标$(\alpha, \beta, \gamma)$ 31 b. 计算$v_c$相对于三角形$f$的重心坐标$(\alpha, \beta, \gamma)$
31- c. 根据$(\alpha, \beta, \gamma)$判断$v_c$是否在三角形内部。若$v_c$不在三角形内部,则处理矩形内下个像素点,否则执行下述步骤 32+ c. 根据$(\alpha, \beta, \gamma)$判断$v_c$是否在三角形内部。若$v_c$不在三角形内部,则处理矩形内下个像素点,否则执行下述步骤
32- d. 使用$(\alpha, \beta, \gamma)$和$v_{s0}$,$v_{s1}$,$v_{s2}$得到当前像素的深度值depth 33+ d. 使用$(\alpha, \beta, \gamma)$和$(v_{s0}, v_{s1}, v_{s2})$得到当前像素的深度值depth
33 e. 若启用了深度先验;否则,直接执行下一步“Z-Buffer更新”34 e. 若启用了深度先验;否则,直接执行下一步“Z-Buffer更新”
34- 35+ 
35 - 使用深度先验图计算深度阈值depth_thres36 - 使用深度先验图计算深度阈值depth_thres
36 - 如果depth < depth_thres,处理矩形内下个像素点,否则执行下述步骤37 - 如果depth < depth_thres,处理矩形内下个像素点,否则执行下述步骤
37- 38+ 
38 f. Z-Buffer更新:39 f. Z-Buffer更新:
39 40
40 - 若$depth < z_{\min}(x_i, y_i)$:41 - 若$depth < z_{\min}(x_i, y_i)$:
41 42
42- $$43+ $$
43- \quad z_{\min}(x_i, y_i) \gets \text{depth} \\44+ \quad z_{\min}(x_i, y_i) \gets \text{depth} \\
44- \quad \text{face\_idx}(x_i, y_i) \gets f45+ \quad \text{face\_idx}(x_i, y_i) \gets f
45- $$46+ $$
46 47
47 - 若$depth = z_{\min}(x_i, y_i)$:48 - 若$depth = z_{\min}(x_i, y_i)$:
48 49
49- $$50+ $$
50- \quad \text{face\_idx}(x_i, y_i) \gets \min(\text{face\_idx}(x_i, y_i),\ f)51+ \quad \text{face\_idx}(x_i, y_i) \gets \min(\text{face\_idx}(x_i, y_i),\ f)
51- $$52+ $$
52 53
53 按上述步骤对空间中所有的三角形面片进行处理后,对大小为$height * width$的屏幕上每个像素点$v_i = (x_i, y_i)$:54 按上述步骤对空间中所有的三角形面片进行处理后,对大小为$height * width$的屏幕上每个像素点$v_i = (x_i, y_i)$:
54 55
55 1. 取Z-Buffer中$v_i$对应的面片索引$f_{idx}$,$findices (x_i, y_i) \gets f_{idx}$56 1. 取Z-Buffer中$v_i$对应的面片索引$f_{idx}$,$findices (x_i, y_i) \gets f_{idx}$
56- 2. 将$f$的三个顶点坐标$v_0$,$v_1$,$v_2$转换为屏幕坐标$v_{s0}$,$v_{s1}$,$v_{s2}$57+ 2. 将$f$的三个顶点坐标$(v_0, v_1, v_2)$转换为屏幕坐标$(v_{s0}, v_{s1}, v_{s2})$
57 3. 计算$v_i$的中心点坐标$v_c$58 3. 计算$v_i$的中心点坐标$v_c$
58 4. 计算$v_c$相对于三角形$f$的重心坐标$(\alpha, \beta, \gamma)$59 4. 计算$v_c$相对于三角形$f$的重心坐标$(\alpha, \beta, \gamma)$
59 5. 使用$(\alpha, \beta, \gamma)$计算透视矫正插值$(\tilde{\alpha}, \tilde{\beta}, \tilde{\gamma})$60 5. 使用$(\alpha, \beta, \gamma)$计算透视矫正插值$(\tilde{\alpha}, \tilde{\beta}, \tilde{\gamma})$
@@ -69,10 +70,10 @@
69 z_s = z / w * 0.49999 + 0.570 z_s = z / w * 0.49999 + 0.5
70 $$71 $$
71 72
72- - 点$v$相对于三角形$(v_0, v_1, v_2)$的重心坐标$(\alpha, \beta, \gamma)$73+ - 点$v$相对于三角形 $(v_0, v_1, v_2)$的重心坐标$(\alpha, \beta, \gamma)$
73 74
74 1. 分别计算计算三角形$(v_0, v_1, v_2)$、$(v_0, v, v_2)$和$(v_0, v_1, v)$的有向面积$area$、$beta\_tri$和$gamma\_tri$75 1. 分别计算计算三角形$(v_0, v_1, v_2)$、$(v_0, v, v_2)$和$(v_0, v_1, v)$的有向面积$area$、$beta\_tri$和$gamma\_tri$
75- 2. 若$area$为0,则$\alpha = \beta = \gamma = -1$, 否则76+ 2. 若$area$为0,则$\alpha = \beta = \gamma = -1$,否则
76 77
77 $$78 $$
78 \beta = beta\_tri / area\\79 \beta = beta\_tri / area\\
@@ -80,13 +81,13 @@
80 \alpha = 1 - \beta - \gamma81 \alpha = 1 - \beta - \gamma
81 $$82 $$
82 83 
83- - 由顶点$v_0 = (x_0, y_0, z_0)$,$v_1 = (x_1, y_1, z_1)$和$v_2 = (x_2, y_2, z_2)$组成的三角形的有向面积84+ - 由顶点$v_0 = (x_0, y_0, z_0)$$v_1 = (x_1, y_1, z_1)$和$v_2 = (x_2, y_2, z_2)$组成的三角形的有向面积
84 85
85 $$86 $$
86 area = (x_2 - x_0) * (y_1 - y_0) - (x_1 - x_0) * (y_2 - y_0)87 area = (x_2 - x_0) * (y_1 - y_0) - (x_1 - x_0) * (y_2 - y_0)
87 $$88 $$
88 89
89- - 结合重心坐标$(\alpha, \beta, \gamma)$和三角形屏幕坐标$v_0 = (x_0, y_0, z_0)$, $v_1 = (x_1, y_1, z_1)$和$v_2 = (x_2, y_2, z_2)$计算像素点$v = (x, y)$的深度$depth$90+ - 结合重心坐标$(\alpha, \beta, \gamma)$和三角形屏幕坐标$(v_0 = (x_0, y_0, z_0), v_1 = (x_1, y_1, z_1), v_2 = (x_2, y_2, z_2))$计算像素点$v = (x, y)$ 的深度$depth$
90 91
91 $$92 $$
92 depth = \alpha * z_0 + \beta * z_1 + \gamma * z_293 depth = \alpha * z_0 + \beta * z_1 + \gamma * z_2
@@ -100,8 +101,8 @@
100 101
101 - 根据重心坐标$(\alpha, \beta, \gamma)$判断顶点是否在三角形内102 - 根据重心坐标$(\alpha, \beta, \gamma)$判断顶点是否在三角形内
102 如果$\alpha >= 0$且$\beta >= 0$且$\gamma >= 0$则点在三角形内(包括在三角形边上),否则点不在三角形内。103 如果$\alpha >= 0$且$\beta >= 0$且$\gamma >= 0$则点在三角形内(包括在三角形边上),否则点不在三角形内。
103- - 结合重心坐标$(\lambda_0, \lambda_1, \lambda_2)$以及三角形的三个顶点坐标$v_0 = (x_0, y_0, z_0, w_0)$,$v_1 = (x_1, y_1, z_1, w_1)$和$v_2 = (x_2, y_2, z_2, w_2)$计算透视矫正插值$(\lambda_0^{corrected}, \lambda_1^{corrected}, \lambda_2^{corrected})$104+ - 结合重心坐标$(\lambda_0, \lambda_1, \lambda_2)$以及三角形的三个顶点坐标$v_0 = (x_0, y_0, z_0, w_0)$$v_1 = (x_1, y_1, z_1, w_1)$和$v_2 = (x_2, y_2, z_2, w_2)$计算透视矫正插值$(\lambda_0^{corrected}, \lambda_1^{corrected}, \lambda_2^{corrected})$
104- 105+ 
105 $$106 $$
106 \lambda_i^{corrected} = \frac{\lambda_i / w_i} { \sum (\lambda_j / w_j)}107 \lambda_i^{corrected} = \frac{\lambda_i / w_i} { \sum (\lambda_j / w_j)}
107 $$108 $$
@@ -18,25 +18,26 @@
18- 接口功能:实现光栅化计算。根据给定的三维空间中的点和面,获取屏幕中每个像素点的最小深度及其对应的面片索引,并计算该面片的重心坐标透视矫正插值。18- 接口功能:实现光栅化计算。根据给定的三维空间中的点和面,获取屏幕中每个像素点的最小深度及其对应的面片索引,并计算该面片的重心坐标透视矫正插值。
19 19 
20- 计算公式:20- 计算公式:
21+ 
21 $findices$记录每个像素点最小深度对应的面索引,$barycentric$记录每个顶点相对于$findices$中记录的面的重心坐标透视矫正插值。22 $findices$记录每个像素点最小深度对应的面索引,$barycentric$记录每个顶点相对于$findices$中记录的面的重心坐标透视矫正插值。
22 计算过程中使用的Z-Buffer记录每个像素点$(x, y)$的最小深度$z_{\min}(x, y)$以及该深度对应的三角形面片索引$\text{face\_idx}(x, y)$。23 计算过程中使用的Z-Buffer记录每个像素点$(x, y)$的最小深度$z_{\min}(x, y)$以及该深度对应的三角形面片索引$\text{face\_idx}(x, y)$。
23 24
24 计算过程如下:25 计算过程如下:
25 对空间中的每个三角形面片$f$:26 对空间中的每个三角形面片$f$:
26 27
27- 1. 将$f$的三个顶点坐标$v_0$, $v_1$, $v_2$转换为屏幕坐标$v_{s0}$,$v_{s1}$,$v_{s2}$28+ 1. 将$f$的三个顶点坐标$(v_0, v_1, v_2)$转换为屏幕坐标$(v_{s0}, v_{s1}, v_{s2})$
28- 2. 根据$v_{s0}$,$v_{s1}$,$v_{s2}$计算包围$f$的矩形范围29+ 2. 根据$(v_{s0}, v_{s1}, v_{s2})$计算包围$f$的矩形范围
29 3. 对矩形内每个像素点$v_i = (x_i, y_i)$,执行以下操作:30 3. 对矩形内每个像素点$v_i = (x_i, y_i)$,执行以下操作:
30- 31+ 
31 a. 计算像素中心坐标$v_c$ 32 a. 计算像素中心坐标$v_c$
32 b. 计算$v_c$相对于三角形$f$的重心坐标$(\alpha, \beta, \gamma)$ 33 b. 计算$v_c$相对于三角形$f$的重心坐标$(\alpha, \beta, \gamma)$
33- c. 根据$(\alpha, \beta, \gamma)$判断$v_c$是否在三角形内部。若$v_c$不在三角形内部,则处理矩形内下个像素点,否则执行下述步骤 34+ c. 根据$(\alpha, \beta, \gamma)$判断$v_c$是否在三角形内部。若$v_c$不在三角形内部,则处理矩形内下个像素点,否则执行下述步骤
34- d. 使用$(\alpha, \beta, \gamma)$和$v_{s0}$,$v_{s1}$,$v_{s2}$得到当前像素的深度值depth 35+ d. 使用$(\alpha, \beta, \gamma)$和$(v_{s0}, v_{s1}, v_{s2})$得到当前像素的深度值depth
35 e. 若启用了深度先验;否则,直接执行下一步“Z-Buffer更新”36 e. 若启用了深度先验;否则,直接执行下一步“Z-Buffer更新”
36- 37+ 
37 - 使用深度先验图计算深度阈值depth_thres38 - 使用深度先验图计算深度阈值depth_thres
38 - 如果depth < depth_thres,处理矩形内下个像素点,否则执行下述步骤39 - 如果depth < depth_thres,处理矩形内下个像素点,否则执行下述步骤
39- 40+ 
40 f. Z-Buffer更新:41 f. Z-Buffer更新:
41 42
42 - 若$depth < z_{\min}(x_i, y_i)$:43 - 若$depth < z_{\min}(x_i, y_i)$:
@@ -55,7 +56,7 @@
55 按上述步骤对空间中所有的三角形面片进行处理后,对大小为$height * width$的屏幕上每个像素点$v_i = (x_i, y_i)$:56 按上述步骤对空间中所有的三角形面片进行处理后,对大小为$height * width$的屏幕上每个像素点$v_i = (x_i, y_i)$:
56 57
57 1. 取Z-Buffer中$v_i$对应的面片索引$f_{idx}$,$findices (x_i, y_i) \gets f_{idx}$58 1. 取Z-Buffer中$v_i$对应的面片索引$f_{idx}$,$findices (x_i, y_i) \gets f_{idx}$
58- 2. 将$f$的三个顶点坐标$v_0$,$v_1$,$v_2$转换为屏幕坐标$v_{s0}$,$v_{s1}$,$v_{s2}$59+ 2. 将$f$的三个顶点坐标$(v_0, v_1, v_2)$转换为屏幕坐标$(v_{s0}, v_{s1}, v_{s2})$
59 3. 计算$v_i$的中心点坐标$v_c$60 3. 计算$v_i$的中心点坐标$v_c$
60 4. 计算$v_c$相对于三角形$f$的重心坐标$(\alpha, \beta, \gamma)$61 4. 计算$v_c$相对于三角形$f$的重心坐标$(\alpha, \beta, \gamma)$
61 5. 使用$(\alpha, \beta, \gamma)$计算透视矫正插值$(\tilde{\alpha}, \tilde{\beta}, \tilde{\gamma})$62 5. 使用$(\alpha, \beta, \gamma)$计算透视矫正插值$(\tilde{\alpha}, \tilde{\beta}, \tilde{\gamma})$
@@ -74,7 +75,7 @@
74 - 点$v$相对于三角形 $(v_0, v_1, v_2)$的重心坐标$(\alpha, \beta, \gamma)$75 - 点$v$相对于三角形 $(v_0, v_1, v_2)$的重心坐标$(\alpha, \beta, \gamma)$
75 76
76 1. 分别计算计算三角形$(v_0, v_1, v_2)$、$(v_0, v, v_2)$和$(v_0, v_1, v)$的有向面积$area$、$beta\_tri$和$gamma\_tri$77 1. 分别计算计算三角形$(v_0, v_1, v_2)$、$(v_0, v, v_2)$和$(v_0, v_1, v)$的有向面积$area$、$beta\_tri$和$gamma\_tri$
77- 2. 若$area$为0,则$\alpha = \beta = \gamma = -1$, 否则78+ 2. 若$area$为0,则$\alpha = \beta = \gamma = -1$,否则
78 79
79 $$80 $$
80 \beta = beta\_tri / area\\81 \beta = beta\_tri / area\\
@@ -82,13 +83,13 @@
82 \alpha = 1 - \beta - \gamma83 \alpha = 1 - \beta - \gamma
83 $$84 $$
84 85 
85- - 由顶点$v_0 = (x_0, y_0, z_0)$, $v_1 = (x_1, y_1, z_1)$和$v_2 = (x_2, y_2, z_2)$组成的三角形的有向面积86+ - 由顶点$v_0 = (x_0, y_0, z_0)$$v_1 = (x_1, y_1, z_1)$和$v_2 = (x_2, y_2, z_2)$组成的三角形的有向面积
86 87
87 $$88 $$
88 area = (x_2 - x_0) * (y_1 - y_0) - (x_1 - x_0) * (y_2 - y_0)89 area = (x_2 - x_0) * (y_1 - y_0) - (x_1 - x_0) * (y_2 - y_0)
89 $$90 $$
90 91
91- - 结合重心坐标$(\alpha, \beta, \gamma)$和三角形屏幕坐标$v_0 = (x_0, y_0, z_0)$, $v_1 = (x_1, y_1, z_1)$和$v_2 = (x_2, y_2, z_2)$计算像素点$v = (x, y)$ 的深度$depth$92+ - 结合重心坐标$(\alpha, \beta, \gamma)$和三角形屏幕坐标$(v_0 = (x_0, y_0, z_0), v_1 = (x_1, y_1, z_1), v_2 = (x_2, y_2, z_2))$计算像素点$v = (x, y)$ 的深度$depth$
92 93
93 $$94 $$
94 depth = \alpha * z_0 + \beta * z_1 + \gamma * z_295 depth = \alpha * z_0 + \beta * z_1 + \gamma * z_2
@@ -102,8 +103,8 @@
102 103
103 - 根据重心坐标$(\alpha, \beta, \gamma)$判断顶点是否在三角形内104 - 根据重心坐标$(\alpha, \beta, \gamma)$判断顶点是否在三角形内
104 如果$\alpha >= 0$且$\beta >= 0$且$\gamma >= 0$则点在三角形内(包括在三角形边上),否则点不在三角形内。105 如果$\alpha >= 0$且$\beta >= 0$且$\gamma >= 0$则点在三角形内(包括在三角形边上),否则点不在三角形内。
105- - 结合重心坐标$(\lambda_0, \lambda_1, \lambda_2)$以及三角形的三个顶点坐标$v_0 = (x_0, y_0, z_0, w_0)$, $v_1 = (x_1, y_1, z_1, w_1)$和$v_2 = (x_2, y_2, z_2, w_2)$计算透视矫正插值$(\lambda_0^{corrected}, \lambda_1^{corrected}, \lambda_2^{corrected})$106+ - 结合重心坐标$(\lambda_0, \lambda_1, \lambda_2)$以及三角形的三个顶点坐标$v_0 = (x_0, y_0, z_0, w_0)$$v_1 = (x_1, y_1, z_1, w_1)$和$v_2 = (x_2, y_2, z_2, w_2)$计算透视矫正插值$(\lambda_0^{corrected}, \lambda_1^{corrected}, \lambda_2^{corrected})$
106- 107+ 
107 $$108 $$
108 \lambda_i^{corrected} = \frac{\lambda_i / w_i} { \sum (\lambda_j / w_j)}109 \lambda_i^{corrected} = \frac{\lambda_i / w_i} { \sum (\lambda_j / w_j)}
109 $$110 $$
@@ -175,7 +176,7 @@ aclnnStatus aclnnRasterizer(
175 <td>f(aclTensor*)</td>176 <td>f(aclTensor*)</td>
176 <td>输入</td>177 <td>输入</td>
177 <td>表示空间中的面的输入张量。</td>178 <td>表示空间中的面的输入张量。</td>
178- <td><ul><li>不支持空Tensor。</li><li>shape为(numFaces, 3),其中 numFaces表示空间中面的数量,为正整数。每个面是一个三角形,三角形每个顶点表示为顶点在v中的索引,因此f中元素取值应当是对v中元素的合法索引,即取值范围为[0, numVertices-1]。由调用者保证f中元素合法。</li></ul></td>179+ <td><ul><li>不支持空Tensor。</li><li>shape为(numFaces, 3),其中numFaces表示空间中面的数量,为正整数。每个面是一个三角形,三角形每个顶点表示为顶点在v中的索引,因此f中元素取值应当是对v中元素的合法索引,即取值范围为[0, numVertices-1]。由调用者保证f中元素合法。</li></ul></td>
179 <td>INT32</td>180 <td>INT32</td>
180 <td>ND</td>181 <td>ND</td>
181 <td>2</td>182 <td>2</td>
@@ -11,7 +11,7 @@
11## 功能说明11## 功能说明
12 12 
13- 算子功能:使用双三次插值调整图像大小到指定的大小。13- 算子功能:使用双三次插值调整图像大小到指定的大小。
14-- 计算公式: 14+- 计算公式:
15 15 
16 周边16个点的像素位置:16 周边16个点的像素位置:
17 17 
@@ -31,25 +31,25 @@
31 31
32 对应的梯度累加公式如下:32 对应的梯度累加公式如下:
33 33 
34- 左上点$Q_{11}$:34+ 左上点$Q_{11}$
35 35
36 $$36 $$
37 y(N, C, x_1, y_1) += grads(N, C, h', w') \cdot (1 - d_h) \cdot (1 - d_w)37 y(N, C, x_1, y_1) += grads(N, C, h', w') \cdot (1 - d_h) \cdot (1 - d_w)
38 $$38 $$
39 39 
40- 右上点$Q_{12}$:40+ 右上点$Q_{12}$
41 41
42 $$42 $$
43 y(N, C, x_1, y_2) += grads(N, C, h', w') \cdot (1 - d_h) \cdot d_w43 y(N, C, x_1, y_2) += grads(N, C, h', w') \cdot (1 - d_h) \cdot d_w
44 $$44 $$
45 45 
46- 左下点$Q_{21}$:46+ 左下点$Q_{21}$
47 47
48 $$48 $$
49 y(N, C, x_2, y_1) += grads(N, C, h', w') \cdot d_h \cdot (1 - d_w)49 y(N, C, x_2, y_1) += grads(N, C, h', w') \cdot d_h \cdot (1 - d_w)
50 $$50 $$
51 51 
52- 右下点$Q_{22}$:52+ 右下点$Q_{22}$
53 53
54 $$54 $$
55 y(N, C, x_2, y_2) += grads(N, C, h', w') \cdot d_h \cdot d_w55 y(N, C, x_2, y_2) += grads(N, C, h', w') \cdot d_h \cdot d_w
@@ -223,7 +223,7 @@ aclnnStatus aclnnUpsampleBilinear2dBackward(
223 </table>223 </table>
224 224 
225 - <term>Atlas 推理系列产品</term><term>Atlas 训练系列产品</term><term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term><term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>225 - <term>Atlas 推理系列产品</term><term>Atlas 训练系列产品</term><term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term><term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>
226- - 参数`gradOut`、`out`的数据类型不支持BFLOAT16.226+ - 参数`gradOut`、`out`的数据类型不支持BFLOAT16
227 - 参数`out`的数据类型与`gradOut`的数据类型保持一致。227 - 参数`out`的数据类型与`gradOut`的数据类型保持一致。
228 - <term>Ascend 950PR/Ascend 950DT</term>228 - <term>Ascend 950PR/Ascend 950DT</term>
229 -`gradOut`的数据类型不是FLOAT时,`out`的数据类型与`gradOut`的数据类型保持一致。229 -`gradOut`的数据类型不是FLOAT时,`out`的数据类型与`gradOut`的数据类型保持一致。
@@ -11,7 +11,7 @@
11## 功能说明11## 功能说明
12 12 
13- 算子功能:使用单线性插值调整图像大小到指定的大小。13- 算子功能:使用单线性插值调整图像大小到指定的大小。
14-- 计算公式: 14+- 计算公式:
15 15 
16 对一维数据使用周围两个点进行加权插值16 对一维数据使用周围两个点进行加权插值
17 $$17 $$
@@ -93,4 +93,4 @@
93| 调用方式 | 样例代码 | 说明 |93| 调用方式 | 样例代码 | 说明 |
94| :---- | :---- | :---- |94| :---- | :---- | :---- |
95| aclnn接口 | [test_aclnn_upsample_nearest2d_grad](../upsample_nearest2d_grad/examples/test_aclnn_upsample_nearest2d_grad.cpp) | 通过[aclnnUpsampleNearest2dBackward](../upsample_nearest2d_grad/docs/aclnnUpsampleNearest2dBackward.md)接口方式调用ResizeNearestNeighborV2Grad算子。 |95| aclnn接口 | [test_aclnn_upsample_nearest2d_grad](../upsample_nearest2d_grad/examples/test_aclnn_upsample_nearest2d_grad.cpp) | 通过[aclnnUpsampleNearest2dBackward](../upsample_nearest2d_grad/docs/aclnnUpsampleNearest2dBackward.md)接口方式调用ResizeNearestNeighborV2Grad算子。 |
96-| 图模式 | - | 通过[算子IR](op_graph/resize_nearest_neighbor_v2_grad_proto.h)构图方式调用ResizeNearestNeighborV2Grad算子。 |96+| 图模式 | - | 通过[算子IR](op_graph/resize_nearest_neighbor_v2_grad_proto.h)构图方式调用ResizeNearestNeighborV2Grad算子。 |
@@ -19,24 +19,24 @@
19 19 
20| 参数名 | 输入/输出/属性 | 描述 | 数据类型 | 数据格式 |20| 参数名 | 输入/输出/属性 | 描述 | 数据类型 | 数据格式 |
21|-----|----------|--------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------|---------------------------------------------------------|------|21|-----|----------|--------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------|---------------------------------------------------------|------|
22-| images | 输入 | 输入为四维 Tensor,shape (batch, input_height, input_width, channels)。 | INT8、UINT8、INT16、UINT16、INT32、INT64、FLOAT16、FLOAT、DOUBLE | NHWC |22+| images | 输入 | 输入为四维Tensor,shape为(batch, input_height, input_width, channels)。 | INT8、UINT8、INT16、UINT16、INT32、INT64、FLOAT16、FLOAT、DOUBLE | NHWC |
23-| size | 输入 | 包含 2 个元素的一维 Tensor,依次表示输出高度和输出宽度 [output_height, output_width],两个值都必须为正。 | INT32 | ND |23+| size | 输入 | 包含2个元素的一维Tensor,依次表示输出高度和输出宽度 [output_height, output_width],两个值都必须为正。 | INT32 | ND |
24-| scale | 输入 | 包含 2 个元素的一维 Tensor,依次表示行、列方向的缩放因子 [row_scale, col_scale],两个值都必须大于 0。 | FLOAT | ND |24+| scale | 输入 | 包含2个元素的一维Tensor,依次表示行、列方向的缩放因子 [row_scale, col_scale],两个值都必须大于0。 | FLOAT | ND |
25-| translation | 输入 | 包含 2 个元素的一维 Tensor,依次表示行、列方向的平移量 [row_translation, col_translation]。 | FLOAT | ND |25+| translation | 输入 | 包含2个元素的一维Tensor,依次表示行、列方向的平移量 [row_translation, col_translation]。 | FLOAT | ND |
26-| kernel_type | 可选属性 | • 指定重采样核类型。支持 "lanczos1","lanczos3","lanczos5","gaussian","box","triangle","keyscubic","mitchellcubic"。<br>• 默认值为 "lanczos3"。 | STRING | - |26+| kernel_type | 可选属性 | • 指定重采样核类型。支持"lanczos1","lanczos3","lanczos5","gaussian","box","triangle","keyscubic","mitchellcubic"。<br>• 默认值为"lanczos3"。 | STRING | - |
27-| antialias | 可选属性 | • 指定是否在缩小场景下按更大的有效核半径生成采样权重,以降低混叠。<br>• 默认值为 true。 | BOOL | - |27+| antialias | 可选属性 | • 指定是否在缩小场景下按更大的有效核半径生成采样权重,以降低混叠。<br>• 默认值为true。 | BOOL | - |
28-| y | 输出 | 返回四维 FLOAT Tensor,shape (batch, output_height, output_width, channels)。 | FLOAT | NHWC |28+| y | 输出 | 返回四维FLOAT Tensor,shape为(batch, output_height, output_width, channels)。 | FLOAT | NHWC |
29 29 
30## 约束说明30## 约束说明
31 31 
32-- images 必须为四维 Tensor,且 input_height、input_width、channels 必须大于 0。32+- images必须为四维Tensor,且input_height、input_width、channels必须大于0。
33-- size 必须为包含 2 个元素的一维 INT32 Tensor,且 output_height、output_width 必须为正。33+- size必须为包含2个元素的一维INT32 Tensor,且output_height、output_width必须为正。
34-- scale 按源码实现读取前 2 FLOAT 元素,分别表示 row_scale col_scale,且两者都必须大于 0。34+- scale按源码实现读取前2个FLOAT元素,分别表示row_scale和col_scale,且两者都必须大于0。
35-- translation 按源码实现读取前 2 FLOAT 元素,分别表示 row_translation col_translation。35+- translation按源码实现读取前2个FLOAT元素,分别表示row_translation和col_translation。
36-- y 的数据类型固定为 FLOAT,shape images batch、channels 以及 size 指定的输出高宽共同决定。36+- y的数据类型固定为FLOAT,shape由images的batch、channels以及size指定的输出高宽共同决定。
37 37 
38## 调用说明38## 调用说明
39 39 
40| 调用方式 | 调用样例 | 说明 |40| 调用方式 | 调用样例 | 说明 |
41|--------------|------------------------------------------------------------------------|----------------------------------------------------------------|41|--------------|------------------------------------------------------------------------|----------------------------------------------------------------|
42-| 图模式调用 | [test_geir_scale_and_translate](./examples/test_geir_scale_and_translate.cpp) | 通过[算子IR](./op_graph/scale_and_translate_proto.h)构图方式调用ScaleAndTranslate算子。 |42+| 图模式调用 | [test_geir_scale_and_translate](./examples/test_geir_scale_and_translate.cpp) | 通过[算子IR](./op_graph/scale_and_translate_proto.h)构图方式调用ScaleAndTranslate算子。 |
@@ -13,7 +13,7 @@
13 13 
14## 功能说明14## 功能说明
15 15 
16-算子功能:Spatial Transformer Network (STN) 算子用于对输入张量进行仿射变换。该算子通过变换矩阵 theta 对输入图像 x 进行空间变换,输出变换后的图像 y。16+算子功能:Spatial Transformer Network (STN) 算子用于对输入张量进行仿射变换。该算子通过变换矩阵theta对输入图像x进行空间变换,输出变换后的图像y。
17 17 
18## 参数说明18## 参数说明
19 19 
@@ -60,28 +60,28 @@
60 <tr>60 <tr>
61 <td>output_size</td>61 <td>output_size</td>
62 <td>属性</td>62 <td>属性</td>
63- <td>指定输出的高度和宽度,包含 2 个整数。默认为 [-1, -1],表示使用输入尺寸。</td>63+ <td>指定输出的高度和宽度,包含2个整数。默认为 [-1, -1],表示使用输入尺寸。</td>
64 <td>ListInt</td>64 <td>ListInt</td>
65 <td>-</td>65 <td>-</td>
66 </tr>66 </tr>
67 <tr>67 <tr>
68 <td>default_theta</td>68 <td>default_theta</td>
69 <td>属性</td>69 <td>属性</td>
70- <td>默认的仿射变换参数,当 use_default_theta true 时使用。默认为空列表。</td>70+ <td>默认的仿射变换参数,当use_default_theta为true时使用。默认为空列表。</td>
71 <td>ListFloat</td>71 <td>ListFloat</td>
72 <td>-</td>72 <td>-</td>
73 </tr>73 </tr>
74 <tr>74 <tr>
75 <td>align_corners</td>75 <td>align_corners</td>
76 <td>属性</td>76 <td>属性</td>
77- <td>如果为 true,则输入和输出张量的 4 个角像素中心对齐,保留角像素的值。默认为 false。</td>77+ <td>如果为true,则输入和输出张量的4个角像素中心对齐,保留角像素的值。默认为false。</td>
78 <td>Bool</td>78 <td>Bool</td>
79 <td>-</td>79 <td>-</td>
80 </tr>80 </tr>
81 <tr>81 <tr>
82 <td>use_default_theta</td>82 <td>use_default_theta</td>
83 <td>属性</td>83 <td>属性</td>
84- <td>指定哪些 theta 参数从 default_theta 使用。1 表示使用默认值,0 表示使用输入 theta。默认为空列表。</td>84+ <td>指定哪些theta参数从default_theta使用。1表示使用默认值,0表示使用输入theta。默认为空列表。</td>
85 <td>ListInt</td>85 <td>ListInt</td>
86 <td>-</td>86 <td>-</td>
87 </tr>87 </tr>
@@ -89,11 +89,11 @@
89 89 
90## 约束说明90## 约束说明
91 91 
92-- 输入张量 x 的格式必须为 NCHW NC1HWC0。92+- 输入张量x的格式必须为NCHW或NC1HWC0。
93-- 输出张量 y 的格式必须与输入张量 x 的格式一致。93+- 输出张量y的格式必须与输入张量x的格式一致。
94-- 变换矩阵 theta 的形状必须为 [batch, 2, 3] 或 [2, 3]。94+- 变换矩阵theta的形状必须为 [batch, 2, 3] 或 [2, 3]。
95-- 当 use_default_theta 为空列表时,使用输入 theta 进行变换。95+- 当use_default_theta为空列表时,使用输入theta进行变换。
96-- 当 use_default_theta 不为空时,对应位置为 1 的参数使用 default_theta 中的值,为 0 的参数使用输入 theta 中的值。96+- 当use_default_theta不为空时,对应位置为1的参数使用default_theta中的值,为0的参数使用输入theta中的值。
97 97 
98## 调用说明98## 调用说明
99 99 
@@ -15,7 +15,7 @@
15 15 
16## 功能说明16## 功能说明
17 17 
18-- 算子功能:对由多个输入通道组成的输入信号应用双三次抗锯齿算法进行上采样。如果输入Tensor x的shape为(N, C, H, W) ,则输出Tensor out的shape为(N, C, outputSize[0], outputSize[1])。18+- 算子功能:对由多个输入通道组成的输入信号应用双三次抗锯齿算法进行上采样。如果输入Tensor x的shape为(N, C, H, W),则输出Tensor out的shape为(N, C, outputSize[0], outputSize[1])。
19- 计算公式:对于一个二维插值点$(N, C, h, w)$,插值$out(N, C, h, w)$可以表示为:19- 计算公式:对于一个二维插值点$(N, C, h, w)$,插值$out(N, C, h, w)$可以表示为:
20 20
21 $$21 $$
@@ -15,7 +15,7 @@
15 15 
16## 功能说明16## 功能说明
17 17 
18-- 接口功能:对由多个输入通道组成的输入信号应用双三次抗锯齿算法进行上采样。如果输入Tensor x的shape为(N, C, H, W) ,则输出Tensor out的shape为(N, C, outputSize[0], outputSize[1])。18+- 接口功能:对由多个输入通道组成的输入信号应用双三次抗锯齿算法进行上采样。如果输入Tensor x的shape为(N, C, H, W),则输出Tensor out的shape为(N, C, outputSize[0], outputSize[1])。
19 19 
20- 计算公式:对于一个二维插值点$(N, C, h, w)$,插值$out(N, C, h, w)$可以表示为:20- 计算公式:对于一个二维插值点$(N, C, h, w)$,插值$out(N, C, h, w)$可以表示为:
21 21
@@ -16,7 +16,7 @@
16## 功能说明16## 功能说明
17 17 
18- 算子功能:对由多个输入通道组成的输入信号应用2D双线性抗锯齿采样。18- 算子功能:对由多个输入通道组成的输入信号应用2D双线性抗锯齿采样。
19-- 计算公式:对于一个二维插值点$(N, C, H, W)$, 插值$I(N, C, H, W)$可以表示为:19+- 计算公式:对于一个二维插值点$(N, C, H, W)$插值$I(N, C, H, W)$可以表示为:
20 20
21 $$21 $$
22 {I(N, C, H, W)} = \sum_{i=0}^{kW}\sum_{j=0}^{kH}{w(i) * w(j)} * {f(h_i, w_j)}/\sum_{i=0}^{kW}w(i)/\sum_{j=0}^{kH}w(j)22 {I(N, C, H, W)} = \sum_{i=0}^{kW}\sum_{j=0}^{kH}{w(i) * w(j)} * {f(h_i, w_j)}/\sum_{i=0}^{kW}w(i)/\sum_{j=0}^{kH}w(j)
@@ -17,7 +17,7 @@
17 17 
18- 接口功能:对由多个输入通道组成的输入信号应用2D双线性抗锯齿采样。18- 接口功能:对由多个输入通道组成的输入信号应用2D双线性抗锯齿采样。
19 19 
20-- 计算公式:对于一个二维插值点$(N, C, H, W)$, 插值$I(N, C, H, W)$可以表示为:20+- 计算公式:对于一个二维插值点$(N, C, H, W)$插值$I(N, C, H, W)$可以表示为:
21 21
22 $$22 $$
23 {I(N, C, H, W)} = \sum_{i=0}^{kW}\sum_{j=0}^{kH}{w(i) * w(j)} * {f(h_i, w_j)}/\sum_{i=0}^{kW}w(i)/\sum_{j=0}^{kH}w(j)23 {I(N, C, H, W)} = \sum_{i=0}^{kW}\sum_{j=0}^{kH}{w(i) * w(j)} * {f(h_i, w_j)}/\sum_{i=0}^{kW}w(i)/\sum_{j=0}^{kH}w(j)
@@ -195,7 +195,7 @@ aclnnStatus aclnnUpsampleBilinear2dAA(
195 </tbody>195 </tbody>
196 </table>196 </table>
197 197 
198-* **返回值**:198+- **返回值**:
199 199 
200 aclnnStatus:返回状态码,具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。200 aclnnStatus:返回状态码,具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
201 201
@@ -15,7 +15,7 @@
15 15 
16- 算子功能:[UpsampleBilinear2dAA](../upsample_bilinear2d_aa/README.md)的反向传播。16- 算子功能:[UpsampleBilinear2dAA](../upsample_bilinear2d_aa/README.md)的反向传播。
17 17 
18-- 计算公式:对于一个二维插值点$(N, C, H, W)$, 插值$I(N, C, H, W)$可以表示为:18+- 计算公式:对于一个二维插值点$(N, C, H, W)$插值$I(N, C, H, W)$可以表示为:
19 19
20 $$20 $$
21 {I(N, C, H, W)} = \sum_{i=0}^{kW}\sum_{j=0}^{kH}{w(i) * w(j)} * {f(h_i, w_j)}/\sum_{i=0}^{kW}w(i)/\sum_{j=0}^{kH}w(j)21 {I(N, C, H, W)} = \sum_{i=0}^{kW}\sum_{j=0}^{kH}{w(i) * w(j)} * {f(h_i, w_j)}/\sum_{i=0}^{kW}w(i)/\sum_{j=0}^{kH}w(j)
@@ -228,7 +228,7 @@ aclnnStatus aclnnUpsampleBilinear2dBackwardV2(
228 228 
229 - <term>Atlas 训练系列产品</term><term>Atlas 推理系列产品</term>229 - <term>Atlas 训练系列产品</term><term>Atlas 推理系列产品</term>
230 230
231- 参数`gradOut`、`out`的数据类型不支持BFLOAT16.231+ 参数`gradOut`、`out`的数据类型不支持BFLOAT16
232 232 
233- **返回值**233- **返回值**
234 234 
@@ -95,7 +95,7 @@
95 <tr>95 <tr>
96 <td>output_size</td>96 <td>output_size</td>
97 <td>可选属性</td>97 <td>可选属性</td>
98- <td><ul><li>表示输入`grad_output`在D、H和W维度上的空间大小。size为3,且各元素均大于零。 包含3个元素:[depth, height, width]。只能指定'scales'和'output_size'中的一个。必须满足:grad_output_tensor_size[2] == floor(input_size[2] * scales[0]) == output_size[0];grad_output_tensor_size[3] == floor(input_size[3] * scales[1]) == output_size[1];grad_output_tensor_size[4] == floor(input_size[4] * scales[2]) == output_size[2]。</li><li>默认为空。</li></ul></td>98+ <td><ul><li>表示输入`grad_output`在D、H和W维度上的空间大小。size为3,且各元素均大于零。包含3个元素:[depth, height, width]。只能指定'scales'和'output_size'中的一个。必须满足:grad_output_tensor_size[2] == floor(input_size[2] * scales[0]) == output_size[0];grad_output_tensor_size[3] == floor(input_size[3] * scales[1]) == output_size[1];grad_output_tensor_size[4] == floor(input_size[4] * scales[2]) == output_size[2]。</li><li>默认为空。</li></ul></td>
99 <td>LISTINT</td>99 <td>LISTINT</td>
100 <td>-</td>100 <td>-</td>
101 </tr>101 </tr>
@@ -73,7 +73,7 @@
73 73 
74## 约束说明74## 约束说明
75 75 
76-* 确定性计算:aclnnBlendImagesCustom默认确定性实现。76+
77 77 
78## 调用说明78## 调用说明
79 79 
@@ -225,7 +225,7 @@ aclnnStatus aclnnCIoU(
225 225 
226- **参数说明**226- **参数说明**
227 227 
228- <table style="undefined;table-layout: fixed; width: 1155px"><colgroup>228+ <table style="undefined;table-layout: fixed; width: 1155px"><colgroup>
229 <col style="width: 319px">229 <col style="width: 319px">
230 <col style="width: 144px">230 <col style="width: 144px">
231 <col style="width: 671px">231 <col style="width: 671px">
@@ -21,7 +21,7 @@
21 21 
22 $$22 $$
23 IOU = \frac {Area_3} {Area_1 + Area_2 - Area_3} \\23 IOU = \frac {Area_3} {Area_1 + Area_2 - Area_3} \\
24- IOF = \frac {Area_3} {Area_2} 24+ IOF = \frac {Area_3} {Area_2}
25 $$25 $$
26 26 
27 其中,Area_1为bBox的面积,Area_2为gtBox的面积,Area_3为两者重叠部分面积,x和y的定义见参数说明。27 其中,Area_1为bBox的面积,Area_2为gtBox的面积,Area_3为两者重叠部分面积,x和y的定义见参数说明。
@@ -102,11 +102,11 @@
102 </tr>102 </tr>
103 </tbody></table>103 </tbody></table>
104 104 
105-- Kirin X90/Kirin 9030 处理器系列产品: 不支持BFLOAT16。105+- Kirin X90/Kirin 9030 处理器系列产品不支持BFLOAT16。
106 106 
107## 约束说明107## 约束说明
108 108 
109-* 输入shape限制:输入shape为(N, 4)的二维张量,第二维的四个值(X1, Y1, X2, Y2)需满足X1 < X2, Y1 < Y2109+- 输入shape限制:输入shape为(N, 4)的二维张量,第二维的四个值(X1, Y1, X2, Y2)需满足X1 < X2, Y1 < Y2
110 110 
111## 调用说明111## 调用说明
112 112 
@@ -170,6 +170,7 @@ aclnnStatus aclnnIou(
170 aclnnStatus:返回状态码,具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。170 aclnnStatus:返回状态码,具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
171 171 
172 第一段接口完成入参校验,出现以下场景时报错:172 第一段接口完成入参校验,出现以下场景时报错:
173+
173 <table style="undefined;table-layout: fixed; width: 1148px"><colgroup>174 <table style="undefined;table-layout: fixed; width: 1148px"><colgroup>
174 <col style="width: 290px">175 <col style="width: 290px">
175 <col style="width: 134px">176 <col style="width: 134px">
@@ -17,7 +17,11 @@
17 17 
18- 接口功能:完成张量rgb和张量alpha的透明度乘法计算18- 接口功能:完成张量rgb和张量alpha的透明度乘法计算
19 19 
20-- 计算公式:out = rgb * ((broadcast)alpha/255)20+- 计算公式:
21+ 
22+ $$
23+ out = rgb * ((broadcast)alpha/255)
24+ $$
21 25 
22- 示例:26- 示例:
23 假设rgb是一张三通道彩色图片,alpha是其对应的透明度(单通道),使用该算子后可以将该图片生成带透明度的三通道图片。27 假设rgb是一张三通道彩色图片,alpha是其对应的透明度(单通道),使用该算子后可以将该图片生成带透明度的三通道图片。
@@ -15,25 +15,24 @@
15 15 
16## 功能说明16## 功能说明
17 17 
18-- 接口功能:对每个类别的检测框,先按置信度分数降序排序,依次选取分数最高的框作为选中框,并抑制与选中框 IoU 超过 `iouThreshold` 的候选框,同时移除分数不超过 `scoreThreshold` 的输出框。18+- 接口功能:对每个类别的检测框,先按置信度分数降序排序,依次选取分数最高的框作为选中框,并抑制与选中框IoU超过 `iouThreshold` 的候选框,同时移除分数不超过 `scoreThreshold` 的输出框。
19 19 
20- 计算公式:20- 计算公式:
21- - IoU 计算公式:对于两个边界框 A B,IoU 定义为交集面积与并集面积之比:21+ - IoU计算公式:对于两个边界框A和B,IoU定义为交集面积与并集面积之比:
22 22 
23 $$\text{IoU} = \frac{\text{Area}(A \cap B)}{\text{Area}(A \cup B)} = \frac{\text{Area}(A \cap B)}{\text{Area}(A) + \text{Area}(B) - \text{Area}(A \cap B)}$$23 $$\text{IoU} = \frac{\text{Area}(A \cap B)}{\text{Area}(A \cup B)} = \frac{\text{Area}(A \cap B)}{\text{Area}(A) + \text{Area}(B) - \text{Area}(A \cap B)}$$
24 24 
25 - 阈值比较规则25 - 阈值比较规则
26 26 
27- IoU 抑制采用**严格大于**比较,等价于以下代数变换:27+ IoU抑制采用**严格大于**比较,等价于以下代数变换:
28 28 
29 $$\text{IoU} > \text{iouThreshold} \quad \Leftrightarrow \quad \text{Area}(A \cap B) > \bigl(\text{Area}(A) + \text{Area}(B)\bigr) \times \frac{\text{iouThreshold}}{1 + \text{iouThreshold}}$$29 $$\text{IoU} > \text{iouThreshold} \quad \Leftrightarrow \quad \text{Area}(A \cap B) > \bigl(\text{Area}(A) + \text{Area}(B)\bigr) \times \frac{\text{iouThreshold}}{1 + \text{iouThreshold}}$$
30 30 
31 即当交集面积超过右式阈值时,候选框被抑制。31 即当交集面积超过右式阈值时,候选框被抑制。
32 32 
33- - **iouThreshold**(`aclFloatArray*`):取值范围 `[0, 1]`,控制 NMS 抑制强度。值越大,保留的框越多。33+ - **iouThreshold**(`aclFloatArray*`):取值范围 `[0, 1]`,控制NMS抑制强度。值越大,保留的框越多。
34 - **scoreThreshold**`aclFloatArray*`):取值范围 `[0, 1]`,最终输出时过滤低分框,分数 ≤ `scoreThreshold` 的框填充 `-1` 表示无效。34 - **scoreThreshold**`aclFloatArray*`):取值范围 `[0, 1]`,最终输出时过滤低分框,分数 ≤ `scoreThreshold` 的框填充 `-1` 表示无效。
35 35 
36- 
37## 函数原型36## 函数原型
38 37 
39每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md),必须先调用“aclnnNonMaxSuppressionGetWorkspaceSize”接口获取计算所需workspace大小以及包含了算子计算流程的执行器,再调用“aclnnNonMaxSuppression”接口执行计算。38每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md),必须先调用“aclnnNonMaxSuppressionGetWorkspaceSize”接口获取计算所需workspace大小以及包含了算子计算流程的执行器,再调用“aclnnNonMaxSuppression”接口执行计算。
@@ -210,7 +209,7 @@ aclnnStatus aclnnNonMaxSuppression(
210 <td>boxes、scores和maxOutputBoxesPerClass的数据类型不在支持的范围内。</td>209 <td>boxes、scores和maxOutputBoxesPerClass的数据类型不在支持的范围内。</td>
211 </tr>210 </tr>
212 <tr>211 <tr>
213- <td>boxes、scores和 selectedIndices 的数据格式不在支持的范围内。</td>212+ <td>boxes、scores和selectedIndices的数据格式不在支持的范围内。</td>
214 </tr>213 </tr>
215 <tr>214 <tr>
216 <td>boxes、scores需为3维。</td>215 <td>boxes、scores需为3维。</td>
@@ -102,8 +102,8 @@
102 102 
103## 约束说明103## 约束说明
104 104 
105-* 输入shape必须为[N,C,W,H]的Tensor,其中N、C、W和H均为INT32类型正整数。C的取值在(0, 1024]之间。105+- 输入shape必须为[N,C,W,H]的Tensor,其中N、C、W和H均为INT32类型正整数。C的取值在(0, 1024]之间。
106-* rois的shape必须为[n, 6]的Tensor。n的取值范围为[1, 8192]。对于每个ROI而言,其组成为[batch_idx, center_x, center_y, w, h, angle],其中batch_idx取值范围在[0, N)之间,会进行强制类型转换(float->int),center_x、center_y、w、h为FLOAT32类型正浮点数,center_x与w的取值范围为[0, W),center_y与h的取值范围为[0, H),angle的取值为FLOAT32类型浮点数,取值范围为[0, π)。106+- rois的shape必须为[n, 6]的Tensor。n的取值范围为[1, 8192]。对于每个ROI而言,其组成为[batch_idx, center_x, center_y, w, h, angle],其中batch_idx取值范围在[0, N)之间,会进行强制类型转换(float->int),center_x、center_y、w、h为FLOAT32类型正浮点数,center_x与w的取值范围为[0, W),center_y与h的取值范围为[0, H),angle的取值为FLOAT32类型浮点数,取值范围为[0, π)。
107 107 
108## 调用说明108## 调用说明
109 109 
@@ -135,7 +135,7 @@
135* gradOutput、argmax的shape[2]等于pooledH和shape[3]等于pooledW135* gradOutput、argmax的shape[2]等于pooledH和shape[3]等于pooledW
136* rois的值大于等于0136* rois的值大于等于0
137* pooledH、pooledW大于0。137* pooledH、pooledW大于0。
138-* rois[:, 1] 小于 rois[:, 2] 且 rois[:, 3] 小于 rois[:, 4]138+* rois[:, 1] 小于rois[:, 2] 且rois[:, 3] 小于rois[:, 4]
139* rois.shape[0]、gradOutput.shape[0]小于等于1024139* rois.shape[0]、gradOutput.shape[0]小于等于1024
140 140 
141## 调用说明141## 调用说明
@@ -223,7 +223,7 @@ aclnnStatus aclnnRoiPoolingGradWithArgMax(
223 <td>pooledH、pooledW大于0。</td>223 <td>pooledH、pooledW大于0。</td>
224 </tr>224 </tr>
225 <tr>225 <tr>
226- <td>rois[:, 1] 小于 rois[:, 2] 且 rois[:, 3] 小于 rois[:, 4]。</td>226+ <td>rois[:, 1] 小于rois[:, 2] 且rois[:, 3] 小于rois[:, 4]。</td>
227 </tr>227 </tr>
228 <tr>228 <tr>
229 <td>rois.shape[0]、gradOutput.shape[0]小于等于1024。</td>229 <td>rois.shape[0]、gradOutput.shape[0]小于等于1024。</td>
@@ -288,7 +288,7 @@ aclnnStatus aclnnRoiPoolingGradWithArgMax(
288- gradOutput、argmax的shape[2]等于pooledH和shape[3]等于pooledW。288- gradOutput、argmax的shape[2]等于pooledH和shape[3]等于pooledW。
289- rois的值大于等于0。289- rois的值大于等于0。
290- pooledH、pooledW大于0。290- pooledH、pooledW大于0。
291-- rois[:, 1] 小于 rois[:, 2] 且 rois[:, 3] 小于 rois[:, 4]。291+- rois[:, 1] 小于rois[:, 2] 且rois[:, 3] 小于rois[:, 4]。
292- rois.shape[0]、gradOutput.shape[0]小于等于1024。292- rois.shape[0]、gradOutput.shape[0]小于等于1024。
293- 确定性计算:293- 确定性计算:
294 - <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品、Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>:aclnnRoiPoolingGradWithArgMax默认确定性实现。294 - <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品、Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>:aclnnRoiPoolingGradWithArgMax默认确定性实现。
@@ -13,13 +13,13 @@
13 13 
14## 功能说明14## 功能说明
15 15 
16-- 算子功能:对输入特征图按 ROI(感兴趣区域)进行池化,在每个 ROI 内按空间划分为 `pooled_h × pooled_w` 个格子,对每个格子做最大池化,并输出池化结果及最大值在通道内的一维索引(argmax)。16+- 算子功能:对输入特征图按ROI(感兴趣区域)进行池化,在每个ROI内按空间划分为 `pooled_h × pooled_w` 个格子,对每个格子做最大池化,并输出池化结果及最大值在通道内的一维索引(argmax)。
17 17 
18- 计算公式:18- 计算公式:
19 19 
20- 输入特征图 $x$ 的 shape 为 $(N, C, H, W)$,ROI 张量 $\text{rois}$ 的 shape 为 $(\text{num\_rois}, 5)$,每行表示 $(b_n, x_1, y_1, x_2, y_2)$。标量参数为 $s_h$、$s_w$(spatial_scale)以及 $\text{pooled\_h}$、$\text{pooled\_w}$。下标 $n$ 表示 ROI 索引,$c$ 表示通道,$(\text{ph}, \text{pw})$ 表示池化格点。20+ 输入特征图 $x$ 的shape为 $(N, C, H, W)$,ROI张量 $\text{rois}$ 的shape为 $(\text{num\_rois}, 5)$,每行表示 $(b_n, x_1, y_1, x_2, y_2)$。标量参数为 $s_h$、$s_w$(spatial_scale)以及 $\text{pooled\_h}$、$\text{pooled\_w}$。下标 $n$ 表示ROI索引,$c$ 表示通道,$(\text{ph}, \text{pw})$ 表示池化格点。
21 21 
22- - **ROI 映射到特征图**:将 ROI 坐标乘以 spatial_scale 得到特征图上的浮点区间:22+ - **ROI映射到特征图**:将ROI坐标乘以spatial_scale得到特征图上的浮点区间:
23 23 
24 $$24 $$
25 \tilde{x}_1 = x_1 s_w,\quad \tilde{y}_1 = y_1 s_h,\quad \tilde{x}_2 = (x_2+1)s_w,\quad \tilde{y}_2 = (y_2+1)s_h25 \tilde{x}_1 = x_1 s_w,\quad \tilde{y}_1 = y_1 s_h,\quad \tilde{x}_2 = (x_2+1)s_w,\quad \tilde{y}_2 = (y_2+1)s_h
@@ -29,9 +29,9 @@
29 W_{\text{roi}} = \tilde{x}_2 - \tilde{x}_1,\qquad H_{\text{roi}} = \tilde{y}_2 - \tilde{y}_129 W_{\text{roi}} = \tilde{x}_2 - \tilde{x}_1,\qquad H_{\text{roi}} = \tilde{y}_2 - \tilde{y}_1
30 $$30 $$
31 31 
32- 若 $W_{\text{roi}} \le 0$ 或 $H_{\text{roi}} \le 0$,该 ROI 的 $y$ 全为 0,$\text{argmax}$ 全为 -1。32+ 若 $W_{\text{roi}} \le 0$ 或 $H_{\text{roi}} \le 0$,该ROI的 $y$ 全为0,$\text{argmax}$ 全为 -1。
33 33 
34- - **Bin 步长与区间**:每个池化格 (ph, pw) 对应 ROI 内一个 bin,步长与浮点区间为:34+ - **Bin步长与区间**:每个池化格(ph, pw)对应ROI内一个bin,步长与浮点区间为:
35 35 
36 $$36 $$
37 \Delta w = \frac{W_{\text{roi}}}{\text{pooled\_w}},\qquad \Delta h = \frac{H_{\text{roi}}}{\text{pooled\_h}}37 \Delta w = \frac{W_{\text{roi}}}{\text{pooled\_w}},\qquad \Delta h = \frac{H_{\text{roi}}}{\text{pooled\_h}}
@@ -55,23 +55,23 @@
55 h_1 = \text{clip}(\lfloor\tilde{h}_1\rfloor,\, 0,\, H),\quad h_2 = \text{clip}(\lceil\tilde{h}_2\rceil,\, 0,\, H)55 h_1 = \text{clip}(\lfloor\tilde{h}_1\rfloor,\, 0,\, H),\quad h_2 = \text{clip}(\lceil\tilde{h}_2\rceil,\, 0,\, H)
56 $$56 $$
57 57 
58- 其中 $\text{clip}(a,l,u) = \min(\max(a,l), u)$。若 $w_2 \le w_1$ 或 $h_2 \le h_1$,该 bin 为空:$y=0$,$\text{argmax}=-1$。58+ 其中 $\text{clip}(a,l,u) = \min(\max(a,l), u)$。若 $w_2 \le w_1$ 或 $h_2 \le h_1$,该bin为空:$y=0$,$\text{argmax}=-1$。
59 59 
60- - **池化输出与 Argmax**:记 $b = \text{rois}[n,0]$,bin 区域 $R = \{(h,w) : h_1 \le h < h_2,\, w_1 \le w < w_2\}$,则60+ - **池化输出与Argmax**:记 $b = \text{rois}[n,0]$,bin区域 $R = \{(h,w) : h_1 \le h < h_2,\, w_1 \le w < w_2\}$,则
61 61 
62 $$62 $$
63 y[n,c,\text{ph},\text{pw}] = \max_{(h,w) \in R} x[b,c,h,w]63 y[n,c,\text{ph},\text{pw}] = \max_{(h,w) \in R} x[b,c,h,w]
64 $$64 $$
65 65 
66- (空 $R$ 时为 066+ (空 $R$ 时为0)
67 67 
68 $$68 $$
69 \text{argmax}[n,c,\text{ph},\text{pw}] = h^* W + w^*69 \text{argmax}[n,c,\text{ph},\text{pw}] = h^* W + w^*
70 $$70 $$
71 71 
72- $(h^*, w^*)$ 为 bin 内最大值位置(多解取第一个);空 $R$ 为 -1。72+ $(h^*, w^*)$ 为bin内最大值位置(多解取第一个);空 $R$ 为 -1。
73 73 
74- - **输出 Shape**:74+ - **输出Shape**:
75 75 
76 | 输出 | Shape | 数据类型 |76 | 输出 | Shape | 数据类型 |
77 |------|--------|----------|77 |------|--------|----------|
@@ -170,10 +170,10 @@
170 170 
171## 约束说明171## 约束说明
172 172 
173-* x、rois、y、argmax 的数据类型或格式在支持的范围之内。173+* x、rois、y、argmax的数据类型或格式在支持的范围之内。
174-* x shape 4 维(NCHW)。174+* x的shape是4维(NCHW)。
175-* rois shape 第二维是 5。175+* rois的shape第二维是5。
176-* pooled_h、pooled_w、spatial_scale_h、spatial_scale_w 大于 0。176+* pooled_h、pooled_w、spatial_scale_h、spatial_scale_w大于0。
177* x、argmax、rois的shape[0]相等。177* x、argmax、rois的shape[0]相等。
178* rois.shape[0]、x.shape[0]小于等于1024。178* rois.shape[0]、x.shape[0]小于等于1024。
179* x.shape[1]等于pool_channel。179* x.shape[1]等于pool_channel。
@@ -15,13 +15,13 @@
15 15 
16## 功能说明16## 功能说明
17 17 
18-- 接口功能:对输入特征图按 ROI(感兴趣区域)进行池化,在每个 ROI 内按空间划分为 `pooled_h × pooled_w` 个格子,对每个格子做最大池化,并输出池化结果及最大值在通道内的一维索引(argmax)。18+- 接口功能:对输入特征图按ROI(感兴趣区域)进行池化,在每个ROI内按空间划分为 `pooled_h × pooled_w` 个格子,对每个格子做最大池化,并输出池化结果及最大值在通道内的一维索引(argmax)。
19 19 
20- 计算公式:20- 计算公式:
21 21 
22- 输入特征图 $x$ 的 shape 为 $(N, C, H, W)$,ROI 张量 $\text{rois}$ 的 shape 为 $(\text{num\_rois}, 5)$,每行表示 $(b_n, x_1, y_1, x_2, y_2)$。标量参数为 $s_h$、$s_w$(spatial_scale)以及 $\text{pooled\_h}$、$\text{pooled\_w}$。下标 $n$ 表示 ROI 索引,$c$ 表示通道,$(\text{ph}, \text{pw})$ 表示池化格点。22+ 输入特征图 $x$ 的shape为 $(N, C, H, W)$,ROI张量 $\text{rois}$ 的shape为 $(\text{num\_rois}, 5)$,每行表示 $(b_n, x_1, y_1, x_2, y_2)$。标量参数为 $s_h$、$s_w$(spatial_scale)以及 $\text{pooled\_h}$、$\text{pooled\_w}$。下标 $n$ 表示ROI索引,$c$ 表示通道,$(\text{ph}, \text{pw})$ 表示池化格点。
23 23 
24- - **ROI 映射到特征图**:将 ROI 坐标乘以 spatial_scale 得到特征图上的浮点区间:24+ - **ROI映射到特征图**:将ROI坐标乘以spatial_scale得到特征图上的浮点区间:
25 25 
26 $$26 $$
27 \tilde{x}_1 = x_1 s_w,\quad \tilde{y}_1 = y_1 s_h,\quad \tilde{x}_2 = (x_2+1)s_w,\quad \tilde{y}_2 = (y_2+1)s_h27 \tilde{x}_1 = x_1 s_w,\quad \tilde{y}_1 = y_1 s_h,\quad \tilde{x}_2 = (x_2+1)s_w,\quad \tilde{y}_2 = (y_2+1)s_h
@@ -31,9 +31,9 @@
31 W_{\text{roi}} = \tilde{x}_2 - \tilde{x}_1,\qquad H_{\text{roi}} = \tilde{y}_2 - \tilde{y}_131 W_{\text{roi}} = \tilde{x}_2 - \tilde{x}_1,\qquad H_{\text{roi}} = \tilde{y}_2 - \tilde{y}_1
32 $$32 $$
33 33 
34- 若 $W_{\text{roi}} \le 0$ 或 $H_{\text{roi}} \le 0$,该 ROI 的 $y$ 全为 0,$\text{argmax}$ 全为 -1。34+ 若 $W_{\text{roi}} \le 0$ 或 $H_{\text{roi}} \le 0$,该ROI的 $y$ 全为0,$\text{argmax}$ 全为 -1。
35 35 
36- - **Bin 步长与区间**:每个池化格 (ph, pw) 对应 ROI 内一个 bin,步长与浮点区间为:36+ - **Bin步长与区间**:每个池化格(ph, pw)对应ROI内一个bin,步长与浮点区间为:
37 37 
38 $$38 $$
39 \Delta w = \frac{W_{\text{roi}}}{\text{pooled\_w}},\qquad \Delta h = \frac{H_{\text{roi}}}{\text{pooled\_h}}39 \Delta w = \frac{W_{\text{roi}}}{\text{pooled\_w}},\qquad \Delta h = \frac{H_{\text{roi}}}{\text{pooled\_h}}
@@ -57,23 +57,23 @@
57 h_1 = \text{clip}(\lfloor\tilde{h}_1\rfloor,\, 0,\, H),\quad h_2 = \text{clip}(\lceil\tilde{h}_2\rceil,\, 0,\, H)57 h_1 = \text{clip}(\lfloor\tilde{h}_1\rfloor,\, 0,\, H),\quad h_2 = \text{clip}(\lceil\tilde{h}_2\rceil,\, 0,\, H)
58 $$58 $$
59 59 
60- 其中 $\text{clip}(a,l,u) = \min(\max(a,l), u)$。若 $w_2 \le w_1$ 或 $h_2 \le h_1$,该 bin 为空:$y=0$,$\text{argmax}=-1$。60+ 其中 $\text{clip}(a,l,u) = \min(\max(a,l), u)$。若 $w_2 \le w_1$ 或 $h_2 \le h_1$,该bin为空:$y=0$,$\text{argmax}=-1$。
61 61 
62- - **池化输出与 Argmax**:记 $b = \text{rois}[n,0]$,bin 区域 $R = \{(h,w) : h_1 \le h < h_2,\, w_1 \le w < w_2\}$,则62+ - **池化输出与Argmax**:记 $b = \text{rois}[n,0]$,bin区域 $R = \{(h,w) : h_1 \le h < h_2,\, w_1 \le w < w_2\}$,则
63 63 
64 $$64 $$
65 y[n,c,\text{ph},\text{pw}] = \max_{(h,w) \in R} x[b,c,h,w]65 y[n,c,\text{ph},\text{pw}] = \max_{(h,w) \in R} x[b,c,h,w]
66 $$66 $$
67 67 
68- (空 $R$ 时为 068+ (空$R$时为0)
69 69 
70 $$70 $$
71 \text{argmax}[n,c,\text{ph},\text{pw}] = h^* W + w^*71 \text{argmax}[n,c,\text{ph},\text{pw}] = h^* W + w^*
72 $$72 $$
73 73 
74- $(h^*, w^*)$ 为 bin 内最大值位置(多解取第一个);空 $R$ 为 -1。74+ $(h^*, w^*)$ 为bin内最大值位置(多解取第一个);空 $R$ 为 -1。
75 75 
76- - **输出 Shape**:76+ - **输出Shape**:
77 77 
78 | 输出 | Shape | 数据类型 |78 | 输出 | Shape | 数据类型 |
79 |------|--------|----------|79 |------|--------|----------|
@@ -82,7 +82,7 @@
82 82 
83## 函数原型83## 函数原型
84 84 
85-每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md),必须先调用“aclnnRoiPoolingWithArgMaxGetWorkspaceSize”接口获取计算所需 workspace 大小以及包含了算子计算流程的执行器,再调用“aclnnRoiPoolingWithArgMax”接口执行计算。85+每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md),必须先调用“aclnnRoiPoolingWithArgMaxGetWorkspaceSize”接口获取计算所需workspace大小以及包含了算子计算流程的执行器,再调用“aclnnRoiPoolingWithArgMax”接口执行计算。
86 86 
87```Cpp87```Cpp
88aclnnStatus aclnnRoiPoolingWithArgMaxGetWorkspaceSize(88aclnnStatus aclnnRoiPoolingWithArgMaxGetWorkspaceSize(
@@ -135,8 +135,8 @@ aclnnStatus aclnnRoiPoolingWithArgMax(
135 <tr>135 <tr>
136 <td>x(aclTensor*)</td>136 <td>x(aclTensor*)</td>
137 <td>输入</td>137 <td>输入</td>
138- <td>输入特征图,格式为 NCHW,(N, C, H, W)。</td>138+ <td>输入特征图,格式为NCHW,(N, C, H, W)。</td>
139- <td><ul><li>不支持空 Tensor。</li><li>输入维度必须为 4 维。</li><li>N需要为16的倍数。</li><li>N小于等于1024。</li></ul></td>139+ <td><ul><li>不支持空Tensor。</li><li>输入维度必须为4维。</li><li>N需要为16的倍数。</li><li>N小于等于1024。</li></ul></td>
140 <td>FLOAT32、FLOAT16</td>140 <td>FLOAT32、FLOAT16</td>
141 <td>ND</td>141 <td>ND</td>
142 <td>4</td>142 <td>4</td>
@@ -145,8 +145,8 @@ aclnnStatus aclnnRoiPoolingWithArgMax(
145 <tr>145 <tr>
146 <td>rois(aclTensor*)</td>146 <td>rois(aclTensor*)</td>
147 <td>输入</td>147 <td>输入</td>
148- <td>ROI 框,每行 5 个元素:batch_idx, x1, y1, x2, y2。</td>148+ <td>ROI框,每行5个元素:batch_idx, x1, y1, x2, y2。</td>
149- <td><ul><li>shape 为(num_rois,5),不支持空 Tensor。</li><li>第0维小于等于1024。</li><li>x1, y1, x2, y2大于等于0.0。</li></ul></td>149+ <td><ul><li>shape为(num_rois,5),不支持空Tensor。</li><li>第0维小于等于1024。</li><li>x1, y1, x2, y2大于等于0.0。</li></ul></td>
150 <td>FLOAT32、FLOAT16</td>150 <td>FLOAT32、FLOAT16</td>
151 <td>ND</td>151 <td>ND</td>
152 <td>2</td>152 <td>2</td>
@@ -156,7 +156,7 @@ aclnnStatus aclnnRoiPoolingWithArgMax(
156 <td>pooled_h(int64_t)</td>156 <td>pooled_h(int64_t)</td>
157 <td>输入</td>157 <td>输入</td>
158 <td>池化输出高度。</td>158 <td>池化输出高度。</td>
159- <td>必须大于 0。</td>159+ <td>必须大于0。</td>
160 <td>-</td>160 <td>-</td>
161 <td>-</td>161 <td>-</td>
162 <td>-</td>162 <td>-</td>
@@ -166,7 +166,7 @@ aclnnStatus aclnnRoiPoolingWithArgMax(
166 <td>pooled_w(int64_t)</td>166 <td>pooled_w(int64_t)</td>
167 <td>输入</td>167 <td>输入</td>
168 <td>池化输出宽度。</td>168 <td>池化输出宽度。</td>
169- <td>必须大于 0。</td>169+ <td>必须大于0。</td>
170 <td>-</td>170 <td>-</td>
171 <td>-</td>171 <td>-</td>
172 <td>-</td>172 <td>-</td>
@@ -175,8 +175,8 @@ aclnnStatus aclnnRoiPoolingWithArgMax(
175 <tr>175 <tr>
176 <td>spatial_scale_h(float)</td>176 <td>spatial_scale_h(float)</td>
177 <td>输入</td>177 <td>输入</td>
178- <td>ROI 坐标映射到特征图时在高度方向的缩放比例。</td>178+ <td>ROI坐标映射到特征图时在高度方向的缩放比例。</td>
179- <td>必须大于 0。</td>179+ <td>必须大于0。</td>
180 <td>-</td>180 <td>-</td>
181 <td>-</td>181 <td>-</td>
182 <td>-</td>182 <td>-</td>
@@ -185,8 +185,8 @@ aclnnStatus aclnnRoiPoolingWithArgMax(
185 <tr>185 <tr>
186 <td>spatial_scale_w(float)</td>186 <td>spatial_scale_w(float)</td>
187 <td>输入</td>187 <td>输入</td>
188- <td>ROI 坐标映射到特征图时在宽度方向的缩放比例。</td>188+ <td>ROI坐标映射到特征图时在宽度方向的缩放比例。</td>
189- <td>必须大于 0。</td>189+ <td>必须大于0。</td>
190 <td>-</td>190 <td>-</td>
191 <td>-</td>191 <td>-</td>
192 <td>-</td>192 <td>-</td>
@@ -195,8 +195,8 @@ aclnnStatus aclnnRoiPoolingWithArgMax(
195 <tr>195 <tr>
196 <td>y(aclTensor*)</td>196 <td>y(aclTensor*)</td>
197 <td>输出</td>197 <td>输出</td>
198- <td>池化结果,shape 为(num_rois,C,pooled_h,pooled_w)。</td>198+ <td>池化结果,shape为(num_rois,C,pooled_h,pooled_w)。</td>
199- <td><ul><li>不支持空 Tensor。</li><li>数据类型与 x 一致。</li></ul></td>199+ <td><ul><li>不支持空Tensor。</li><li>数据类型与x一致。</li></ul></td>
200 <td>FLOAT32、FLOAT16</td>200 <td>FLOAT32、FLOAT16</td>
201 <td>ND</td>201 <td>ND</td>
202 <td>4</td>202 <td>4</td>
@@ -206,7 +206,7 @@ aclnnStatus aclnnRoiPoolingWithArgMax(
206 <td>argmax(aclTensor*)</td>206 <td>argmax(aclTensor*)</td>
207 <td>输出</td>207 <td>输出</td>
208 <td>每个池化格点最大值在通道内的线性偏移索引。</td>208 <td>每个池化格点最大值在通道内的线性偏移索引。</td>
209- <td><ul><li>不支持空 Tensor。</li><li>shape y 一致。</li></ul></td>209+ <td><ul><li>不支持空Tensor。</li><li>shape与y一致。</li></ul></td>
210 <td>INT32</td>210 <td>INT32</td>
211 <td>ND</td>211 <td>ND</td>
212 <td>4</td>212 <td>4</td>
@@ -215,7 +215,7 @@ aclnnStatus aclnnRoiPoolingWithArgMax(
215 <tr>215 <tr>
216 <td>workspaceSize(uint64_t*)</td>216 <td>workspaceSize(uint64_t*)</td>
217 <td>输出</td>217 <td>输出</td>
218- <td>返回需要在 Device 侧申请的 workspace 大小。</td>218+ <td>返回需要在Device侧申请的workspace大小。</td>
219 <td>-</td>219 <td>-</td>
220 <td>-</td>220 <td>-</td>
221 <td>-</td>221 <td>-</td>
@@ -225,7 +225,7 @@ aclnnStatus aclnnRoiPoolingWithArgMax(
225 <tr>225 <tr>
226 <td>executor(aclOpExecutor**)</td>226 <td>executor(aclOpExecutor**)</td>
227 <td>输出</td>227 <td>输出</td>
228- <td>返回 op 执行器,包含了算子计算流程。</td>228+ <td>返回op执行器,包含了算子计算流程。</td>
229 <td>-</td>229 <td>-</td>
230 <td>-</td>230 <td>-</td>
231 <td>-</td>231 <td>-</td>
@@ -264,13 +264,13 @@ aclnnStatus aclnnRoiPoolingWithArgMax(
264 <td rowspan="5">161002</td>264 <td rowspan="5">161002</td>
265 </tr>265 </tr>
266 <tr>266 <tr>
267- <td>x、rois、y、argmax 的数据类型或格式不在支持范围内。</td>267+ <td>x、rois、y、argmax的数据类型或格式不在支持范围内。</td>
268 </tr>268 </tr>
269- <tr><td>x shape 不是 4 维(NCHW)。</td>269+ <tr><td>x的shape不是4维(NCHW)。</td>
270 </tr>270 </tr>
271- <tr><td>rois shape 第二维不是 5。</td>271+ <tr><td>rois的shape第二维不是5。</td>
272 </tr>272 </tr>
273- <tr><td>pooled_h、pooled_w、spatial_scale_h、spatial_scale_w 不大于 0。</td>273+ <tr><td>pooled_h、pooled_w、spatial_scale_h、spatial_scale_w不大于0。</td>
274 </tr>274 </tr>
275 </tbody></table>275 </tbody></table>
276 276 
@@ -293,22 +293,22 @@ aclnnStatus aclnnRoiPoolingWithArgMax(
293 <tr>293 <tr>
294 <td>workspace</td>294 <td>workspace</td>
295 <td>输入</td>295 <td>输入</td>
296- <td>在 Device 侧申请的 workspace 内存地址。</td>296+ <td>在Device侧申请的workspace内存地址。</td>
297 </tr>297 </tr>
298 <tr>298 <tr>
299 <td>workspaceSize</td>299 <td>workspaceSize</td>
300 <td>输入</td>300 <td>输入</td>
301- <td>在 Device 侧申请的 workspace 大小,由第一段接口 aclnnRoiPoolingWithArgMaxGetWorkspaceSize 获取。</td>301+ <td>在Device侧申请的workspace大小,由第一段接口aclnnRoiPoolingWithArgMaxGetWorkspaceSize获取。</td>
302 </tr>302 </tr>
303 <tr>303 <tr>
304 <td>executor</td>304 <td>executor</td>
305 <td>输入</td>305 <td>输入</td>
306- <td>op 执行器,包含了算子计算流程。</td>306+ <td>op执行器,包含了算子计算流程。</td>
307 </tr>307 </tr>
308 <tr>308 <tr>
309 <td>stream</td>309 <td>stream</td>
310 <td>输入</td>310 <td>输入</td>
311- <td>指定执行任务的 Stream。</td>311+ <td>指定执行任务的Stream。</td>
312 </tr>312 </tr>
313 </tbody>313 </tbody>
314 </table>314 </table>
@@ -320,12 +320,12 @@ aclnnStatus aclnnRoiPoolingWithArgMax(
320## 约束说明320## 约束说明
321 321 
322- 确定性计算:322- 确定性计算:
323- - aclnnRoiPoolingWithArgMax 默认确定性实现。323+ - aclnnRoiPoolingWithArgMax默认确定性实现。
324- x、argmax、rois的shape[0]相等324- x、argmax、rois的shape[0]相等
325 325 
326## 调用示例326## 调用示例
327 327 
328-示例代码如下,仅供参考,具体编译和执行过程请参考[编译与运行样例](../../../docs/zh/context/编译与运行样例.md)。实际调用时需先通过 opgen 生成 `aclnnop/aclnn_roi_pooling_with_arg_max.h`,若生成的头文件或接口签名不同,请以生成接口为准。328+示例代码如下,仅供参考,具体编译和执行过程请参考[编译与运行样例](../../../docs/zh/context/编译与运行样例.md)。实际调用时需先通过opgen生成 `aclnnop/aclnn_roi_pooling_with_arg_max.h`,若生成的头文件或接口签名不同,请以生成接口为准。
329 329 
330```Cpp330```Cpp
331#include <iostream>331#include <iostream>