Star921
1841
代码
代码
Issues231
Pull Requests361
Actions
项目讨论
Wiki
分析
Star921
1841
已合并
AddMaxPool3DGrad #799
huohuo_wangyan创建于 1月19日
AddMaxPool3DGrad #799
已合并
huohuo_wangyan创建于 1月19日
从huohuo_wangyan/ops-nnpooling合入到cann/ops-nnmaster
共 64 个文件变更+2635-1982
Mcmake/ut.cmake+2-0
@@ -335,6 +335,8 @@ function(AddOpTestCase opName supportedSocVersion otherCompileOptions)
335 COMMAND cp -r ${PROJECT_SOURCE_DIR}/conv/common/op_kernel/* ${CMAKE_BINARY_DIR}/tbe/ascendc/common335 COMMAND cp -r ${PROJECT_SOURCE_DIR}/conv/common/op_kernel/* ${CMAKE_BINARY_DIR}/tbe/ascendc/common
336 COMMAND ${CMAKE_COMMAND} -E make_directory ${CMAKE_BINARY_DIR}/tbe/ascendc/inc336 COMMAND ${CMAKE_COMMAND} -E make_directory ${CMAKE_BINARY_DIR}/tbe/ascendc/inc
337 COMMAND cp -r ${PROJECT_SOURCE_DIR}/common/inc/op_kernel/* ${CMAKE_BINARY_DIR}/tbe/ascendc/inc337 COMMAND cp -r ${PROJECT_SOURCE_DIR}/common/inc/op_kernel/* ${CMAKE_BINARY_DIR}/tbe/ascendc/inc
338+ COMMAND ${CMAKE_COMMAND} -E make_directory ${CMAKE_BINARY_DIR}/tbe/ascendc/pool_3d_common
339+ COMMAND cp -r ${PROJECT_SOURCE_DIR}/pooling/pool_3d_common/op_kernel/* ${CMAKE_BINARY_DIR}/tbe/ascendc/pool_3d_common
338 )340 )
339 ## kernel src copy341 ## kernel src copy
340 kernel_src_copy(342 kernel_src_copy(
Mpooling/adaptive_max_pool3d_grad/CMakeLists.txt+1-9
@@ -8,12 +8,4 @@
8# See LICENSE in the root of the software repository for the full text of the License.8# See LICENSE in the root of the software repository for the full text of the License.
9# ----------------------------------------------------------------------------9# ----------------------------------------------------------------------------
10 10 
11-file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)11+add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE adaptive_max_pool3d_grad ACLNNTYPE aclnn_exclude DEPENDENCIES max_pool3d_grad_with_argmax max_pool_grad_with_argmax_v3 pool_3d_common)
12-if(NOT ENABLE_TEST AND NOT BENCHMARK)
13- list(REMOVE_ITEM CURRENT_DIRS tests)
14-endif()
15-foreach(SUB_DIR ${CURRENT_DIRS})
16- if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
17- add_subdirectory(${SUB_DIR})
18- endif()
19-endforeach()
Mpooling/adaptive_max_pool3d_grad/README.md+5-0
@@ -4,8 +4,13 @@
4 4 
5| 产品 | 是否支持 |5| 产品 | 是否支持 |
6| ---- | :----:|6| ---- | :----:|
7+|昇腾910_95 AI处理器|×|
7|Atlas A3 训练系列产品/Atlas A3 推理系列产品|√|8|Atlas A3 训练系列产品/Atlas A3 推理系列产品|√|
8|Atlas A2 训练系列产品/Atlas 800I A2 推理产品/A200I A2 Box 异构组件|√|9|Atlas A2 训练系列产品/Atlas 800I A2 推理产品/A200I A2 Box 异构组件|√|
10+|Atlas 200I/500 A2推理产品|×|
11+|Atlas 推理系列产品|×|
12+|Atlas 训练系列产品|×|
13+|Atlas 200/300/500 推理产品|×|
9 14 
10## 功能说明15## 功能说明
11 16 
Mpooling/adaptive_max_pool3d_grad/docs/aclnnAdaptiveMaxPool2dBackward.md+182-33
@@ -1,11 +1,18 @@
1# aclnnAdaptiveMaxPool2dBackward1# aclnnAdaptiveMaxPool2dBackward
2 2 
3+[📄 查看源码](https://gitcode.com/cann/ops-nn/tree/master/pooling/adaptive_max_pool3d_grad)
4+ 
3## 产品支持情况5## 产品支持情况
4 6 
5| 产品 | 是否支持 |7| 产品 | 是否支持 |
6| :----------------------------------------------------------- | :------: |8| :----------------------------------------------------------- | :------: |
9+| <term>昇腾910_95 AI处理器</term> | × |
7| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |10| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |
8-| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |11+| <term>Atlas A2 训练系列产品/Atlas 800I A2 推理产品/A200I A2 Box 异构组件</term> | √ |
12+| <term>Atlas 200I/500 A2 推理产品</term> | × |
13+| <term>Atlas 推理系列产品 </term> | × |
14+| <term>Atlas 训练系列产品</term> | × |
15+| <term>Atlas 200/300/500 推理产品</term> | × |
9 16 
10## 功能说明17## 功能说明
11 18 
@@ -13,7 +20,6 @@
13 正向自适应最大池化的反向传播，将梯度回填到每个自适应窗口最大值的坐标处，相同坐标处累加。20 正向自适应最大池化的反向传播，将梯度回填到每个自适应窗口最大值的坐标处，相同坐标处累加。
14- 正向计算公式：21- 正向计算公式：
15 对于输入self维度$[N,C,H,W]$，N（Batch）表示批量大小、C（Channels）表示特征图通道、H（Height）表示特征图高度、W（Width）表示特征图宽度，outputSize值为$[H_o,W_o]$的场景，其输出output维度为$[N,C,H_o,W_o]$，索引indices维度为$[N,C,H_o,W_o]$，相应tensor中每个元素的计算公式如下：22 对于输入self维度$[N,C,H,W]$，N（Batch）表示批量大小、C（Channels）表示特征图通道、H（Height）表示特征图高度、W（Width）表示特征图宽度，outputSize值为$[H_o,W_o]$的场景，其输出output维度为$[N,C,H_o,W_o]$，索引indices维度为$[N,C,H_o,W_o]$，相应tensor中每个元素的计算公式如下：
16- 
17 $$23 $$
18 H_{left}^m = \lfloor(m*H)/H_o\rfloor \\24 H_{left}^m = \lfloor(m*H)/H_o\rfloor \\
19 H_{right}^m = \lceil(m*H)/H_o\rceil \\25 H_{right}^m = \lceil(m*H)/H_o\rceil \\
@@ -27,48 +33,191 @@
27 33 
28每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnAdaptiveMaxPool2dBackwardGetWorkspaceSize”接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用“aclnnAdaptiveMaxPool2dBackward”接口执行计算。34每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnAdaptiveMaxPool2dBackwardGetWorkspaceSize”接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用“aclnnAdaptiveMaxPool2dBackward”接口执行计算。
29 35 
30-- `aclnnStatus aclnnAdaptiveMaxPool2dBackwardGetWorkspaceSize(const aclTensor* gradOutput, const aclTensor* self, const aclTensor* indices, aclTensor* gradInput, uint64_t* workspaceSize, aclOpExecutor** executor)`36+```Cpp
31-- `aclnnStatus aclnnAdaptiveMaxPool2dBackward(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream)`37+aclnnStatus aclnnAdaptiveMaxPool2dBackwardGetWorkspaceSize(
38+ const aclTensor *gradOutput,
39+ const aclTensor *self,
40+ const aclTensor *indices,
41+ aclTensor *gradInput,
42+ uint64_t *workspaceSize,
43+ aclOpExecutor **executor)
44+```
45+```Cpp
46+aclnnStatus aclnnAdaptiveMaxPool2dBackward(
47+ void *workspace,
48+ uint64_t workspaceSize,
49+ aclOpExecutor *executor,
50+ aclrtStream stream)
51+```
32 52 
33## aclnnAdaptiveMaxPool2dBackwardGetWorkspaceSize53## aclnnAdaptiveMaxPool2dBackwardGetWorkspaceSize
34 54 
35- **参数说明：**55- **参数说明：**
36- 56+
37- - gradOutput（aclTensor \*， 计算输入）：梯度Tensor，Device侧aclTensor。数据类型支持FLOAT32、FLOAT16、BFLOAT16，和正向的输出shape一致。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持ND， 当输入是4维时，内部按照NCHW处理，当输入是3维时，在0维度处补1，内部按照NCHW处理。57+ <table style="undefined;table-layout: fixed; width: 1478px"><colgroup>
38- 58+ <col style="width: 149px">
39- - self（aclTensor \*， 计算输入）：正向的输入Tensor，Device侧aclTensor。数据类型支持FLOAT32、FLOAT16、BFLOAT16。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持ND， 当输入是4维时，内部按照NCHW处理，当输入是3维时，在0维度处补1，内部按照NCHW处理，与gradOutput一致。59+ <col style="width: 121px">
40- 60+ <col style="width: 264px">
41- - indices（aclTensor \*， 计算输入）：输入Tensor，是Device侧aclTensor。数据类型支持INT32，INT64。正向输入中最大元素的索引位置。[数据格式](../../../docs/zh/context/数据格式.md)与gradOutput保持一致。shape与gradOutput一致。61+ <col style="width: 253px">
42- 62+ <col style="width: 262px">
43- - gradInput（aclTensor \*，计算输出）：反向输出Tensor，是Device侧aclTensor。数据类型支持FLOAT32、FLOAT16、BFLOAT16，shape与self保持一致。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)与self保持一致。63+ <col style="width: 148px">
44- 64+ <col style="width: 135px">
45- - workspaceSize(uint64_t \*， 出参)：返回需要在Device侧申请的workspace大小。65+ <col style="width: 146px">
46- - executor（aclOpExecutor \*\*， 出参）：返回op执行器，包含了算子计算流程。66+ </colgroup>
67+ <thead>
68+ <tr>
69+ <th>参数名</th>
70+ <th>输入/输出</th>
71+ <th>描述</th>
72+ <th>使用说明</th>
73+ <th>数据类型</th>
74+ <th>数据格式</th>
75+ <th>维度(shape)</th>
76+ <th>非连续Tensor</th>
77+ </tr></thead>
78+ <tbody>
79+ <tr>
80+ <td>gradOutput</td>
81+ <td>输入</td>
82+ <td>当前节点的梯度。</td>
83+ <td>和正向的输出shape一致。</td>
84+ <td>FLOAT32、FLOAT16、BFLOAT16</td>
85+ <td>ND</td>
86+ <td>3-4</td>
87+ <td>√</td>
88+ </tr>
89+ <tr>
90+ <td>self</td>
91+ <td>输入</td>
92+ <td>正向的输入Tensor。</td>
93+ <td>数据类型与gradOutput一致。</td>
94+ <td>FLOAT32、FLOAT16、BFLOAT16</td>
95+ <td>ND</td>
96+ <td>3-4</td>
97+ <td>√</td>
98+ </tr>
99+ <tr>
100+ <td>indices</td>
101+ <td>输入</td>
102+ <td>正向输入中最大元素的索引位置。</td>
103+ <td>shape与gradOutput一致。</td>
104+ <td>INT32，INT64</td>
105+ <td>ND</td>
106+ <td>3-4</td>
107+ <td>√</td>
108+ </tr>
109+ <tr>
110+ <td>gradInput</td>
111+ <td>输出</td>
112+ <td>反向输出Tensor</td>
113+ <td>shape与self保持一致。</td>
114+ <td>FLOAT32、FLOAT16、BFLOAT16</td>
115+ <td>ND</td>
116+ <td>3-4</td>
117+ <td>√</td>
118+ </tr>
119+ <tr>
120+ <td>workspaceSize</td>
121+ <td>输出</td>
122+ <td>返回需要在Device侧申请的workspace大小。</td>
123+ <td>-</td>
124+ <td>-</td>
125+ <td>-</td>
126+ <td>-</td>
127+ <td>-</td>
128+ </tr>
129+ <tr>
130+ <td>executor</td>
131+ <td>输出</td>
132+ <td>返回op执行器，包含了算子计算流程。</td>
133+ <td>-</td>
134+ <td>-</td>
135+ <td>-</td>
136+ <td>-</td>
137+ <td>-</td>
138+ </tr>
139+ </tbody></table>
47- **返回值：**140- **返回值：**
48- 141+
49 aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。142 aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
50- 143+
51- ```
52 第一段接口完成入参校验，出现以下场景时报错：144 第一段接口完成入参校验，出现以下场景时报错：
53- 161001（ACLNN_ERR_PARAM_NULLPTR）：1. 传入的gradOutput、self或indices是空指针。145+ <table style="undefined;table-layout: fixed; width: 1166px"><colgroup>
54- 161002（ACLNN_ERR_PARAM_INVALID）：1. gradOutput、self、indices、gradInput的数据类型不在支持的范围内。146+ <col style="width: 267px">
55- 2. gradOutput、self、indices、gradInput的数据格式不在支持的范围内。147+ <col style="width: 124px">
56- 3. 输入输出的shape不是3维或者4维。148+ <col style="width: 775px">
57- 4. gradOutput与indices的shape不一致，self和gradInput的shape不一致。149+ </colgroup>
58- 5. height * width > max int32，超出了indices的表示范围。150+ <thead>
59- ```151+ <tr>
60- 152+ <th>返回码</th>
153+ <th>错误码</th>
154+ <th>描述</th>
155+ </tr></thead>
156+ <tbody>
157+ <tr>
158+ <td>ACLNN_ERR_PARAM_NULLPTR</td>
159+ <td>161001</td>
160+ <td>传入的gradOutput、self或indices是空指针。</td>
161+ </tr>
162+ <tr>
163+ <td rowspan="5">ACLNN_ERR_PARAM_INVALID</td>
164+ <td rowspan="5">161002</td>
165+ <td>gradOutput、self、indices、gradInput的数据类型不在支持的范围内。</td>
166+ </tr>
167+ <tr>
168+ <td>gradOutput、self、indices、gradInput的数据格式不在支持的范围内。</td>
169+ </tr>
170+ <tr>
171+ <td>输入输出的shape不是3维或者4维。</td>
172+ </tr>
173+ <tr>
174+ <td>gradOutput与indices的shape不一致，self和gradInput的shape不一致。</td>
175+ </tr>
176+ <tr>
177+ <td>height * width > max int32，超出了indices的表示范围。</td>
178+ </tr>
179+ </tbody>
180+ </table>
61## aclnnAdaptiveMaxPool2dBackward181## aclnnAdaptiveMaxPool2dBackward
62 182 
63- **参数说明：**183- **参数说明：**
184+ <table style="undefined;table-layout: fixed; width: 1166px"><colgroup>
185+ <col style="width: 173px">
186+ <col style="width: 133px">
187+ <col style="width: 860px">
188+ </colgroup>
189+ <thead>
190+ <tr>
191+ <th>参数名</th>
192+ <th>输入/输出</th>
193+ <th>描述</th>
194+ </tr></thead>
195+ <tbody>
196+ <tr>
197+ <td>workspace</td>
198+ <td>输入</td>
199+ <td>在Device侧申请的workspace内存地址。</td>
200+ </tr>
201+ <tr>
202+ <td>workspaceSize</td>
203+ <td>输入</td>
204+ <td>在Device侧申请的workspace大小，由第一段接口aclnnAdaptiveMaxPool2dBackwardGetWorkspaceSize获取。</td>
205+ </tr>
206+ <tr>
207+ <td>executor</td>
208+ <td>输入</td>
209+ <td>op执行器，包含了算子计算流程。</td>
210+ </tr>
211+ <tr>
212+ <td>stream</td>
213+ <td>输入</td>
214+ <td>指定执行任务的Stream。</td>
215+ </tr>
216+ </tbody>
217+ </table>
218+- **返回值：**
64 219 
65- - workspace（void \*， 入参）：在Device侧申请的workspace内存地址。220+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
66- - workspaceSize（uint64_t， 入参）：在Device侧申请的workspace大小，由第一段接口aclnnAdaptiveMaxPool2dBackwardGetWorkspaceSize获取。
67- - executor（aclOpExecutor \*， 入参）：op执行器，包含了算子计算流程。
68- - stream（aclrtStream， 入参）：指定执行任务的Stream。
69-- **返回值：**
70- 
71- aclnnStatus: 返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
72 221 
73## 约束说明222## 约束说明
74- 确定性计算：223- 确定性计算：
Mpooling/adaptive_max_pool3d_grad/docs/aclnnAdaptiveMaxPool3dBackward.md+197-48
@@ -1,79 +1,228 @@
1# aclnnAdaptiveMaxPool3dBackward1# aclnnAdaptiveMaxPool3dBackward
2 2 
3+[📄 查看源码](https://gitcode.com/cann/ops-nn/tree/master/pooling/adaptive_max_pool3d_grad)
4+ 
3## 产品支持情况5## 产品支持情况
4 6 
5| 产品 | 是否支持 |7| 产品 | 是否支持 |
6| :----------------------------------------------------------- | :------: |8| :----------------------------------------------------------- | :------: |
9+| <term>昇腾910_95 AI处理器</term> | × |
7| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |10| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |
8-| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |11+| <term>Atlas A2 训练系列产品/Atlas 800I A2 推理产品/A200I A2 Box 异构组件</term> | √ |
12+| <term>Atlas 200I/500 A2 推理产品</term> | × |
13+| <term>Atlas 推理系列产品 </term> | × |
14+| <term>Atlas 训练系列产品</term> | × |
15+| <term>Atlas 200/300/500 推理产品</term> | × |
9 16 
10## 功能说明17## 功能说明
11 18 
12- 算子功能：19- 算子功能：
13 正向自适应最大池化的反向传播，将梯度回填到每个自适应窗口最大值的坐标处，相同坐标处累加。20 正向自适应最大池化的反向传播，将梯度回填到每个自适应窗口最大值的坐标处，相同坐标处累加。
14- 正向计算公式：21- 正向计算公式：
15- > **说明**22+ N（Batch）表示批量大小、H（Height）表示特征图高度、W（Width）表示特征图宽度、C（Channels）表示特征图通道、D（Depth）表示特征图深度
16- > N（Batch）表示批量大小、H（Height）表示特征图高度、W（Width）表示特征图宽度、C（Channels）表示特征图通道、D（Depth）表示特征图深度23+
17- 24+ 对于输入self维度$[N,C,D,H,W]$、outputSize值为$[D_o,H_o,W_o]$的场景，其输出output维度为$[N,C,D_o,H_o,W_o]$，索引indices维度为$[N,C,D_o,H_o,W_o]$，相应tensor中每个元素的计算公式如下：
18- 对于输入self维度$[N,C,D,H,W]$、outputSize值为$[D_o,H_o,W_o]$的场景，其输出output维度为$[N,C,D_o,H_o,W_o]$，索引indices维度为$[N,C,D_o,H_o,W_o]$，相应tensor中每个元素的计算公式如下：25+
19- 26+ $$
20- $$27+ D_{left}^l = \lfloor(l*D)/D_o\rfloor \\
21- D_{left}^l = \lfloor(l*D)/D_o\rfloor \\28+ D_{right}^l = \lceil(l*D)/D_o\rceil \\
22- D_{right}^l = \lceil(l*D)/D_o\rceil \\29+ H_{left}^m = \lfloor(m*H)/H_o\rfloor \\
23- H_{left}^m = \lfloor(m*H)/H_o\rfloor \\30+ H_{right}^m = \lceil(m*H)/H_o\rceil \\
24- H_{right}^m = \lceil(m*H)/H_o\rceil \\31+ W_{left}^n = \lfloor(n*W)/W_o\rfloor \\
25- W_{left}^n = \lfloor(n*W)/W_o\rfloor \\32+ W_{right}^n = \lceil(n*W)/W_o\rceil \\
26- W_{right}^n = \lceil(n*W)/W_o\rceil \\33+ output(N,C,l,m,n) = \mathop{\max}\limits_{i \in [D_{left}^l,D_{right}^l],j\in[H_{left}^m,H_{right}^m],k\in[W_{left}^n,W_{right}^n]} input(N,C,i,j,k) \\
27- output(N,C,l,m,n) = \mathop{\max}\limits_{i \in [D_{left}^l,D_{right}^l],j\in[H_{left}^m,H_{right}^m],k\in[W_{left}^n,W_{right}^n]} input(N,C,i,j,k) \\34+ indices(N,C,l,m,n) = \mathop{\arg\max}\limits_{i \in [D_{left}^l,D_{right}^l],j\in[H_{left}^m,H_{right}^m],k\in[W_{left}^n,W_{right}^n]} input(N,C,i,j,k)
28- indices(N,C,l,m,n) = \mathop{\arg\max}\limits_{i \in [D_{left}^l,D_{right}^l],j\in[H_{left}^m,H_{right}^m],k\in[W_{left}^n,W_{right}^n]} input(N,C,i,j,k)35+ $$
29- $$
30 36 
31## 函数原型37## 函数原型
32 38 
33每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnAdaptiveMaxPool3dBackwardGetWorkspaceSize”接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用“aclnnAdaptiveMaxPool3dBackward”接口执行计算。39每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnAdaptiveMaxPool3dBackwardGetWorkspaceSize”接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用“aclnnAdaptiveMaxPool3dBackward”接口执行计算。
34 40 
35-- `aclnnStatus aclnnAdaptiveMaxPool3dBackwardGetWorkspaceSize(const aclTensor* gradOutput, const aclTensor* self, const aclTensor* indices, aclTensor* gradInput, uint64_t* workspaceSize, aclOpExecutor** executor)`41+```Cpp
36-- `aclnnStatus aclnnAdaptiveMaxPool3dBackward(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream)`42+aclnnStatus aclnnAdaptiveMaxPool3dBackwardGetWorkspaceSize(
37- 43+ const aclTensor *gradOutput,
44+ const aclTensor *self,
45+ const aclTensor *indices,
46+ aclTensor *gradInput,
47+ uint64_t *workspaceSize,
48+ aclOpExecutor **executor)
49+```
50+```Cpp
51+aclnnStatus aclnnAdaptiveMaxPool3dBackward(
52+ void *workspace,
53+ uint64_t workspaceSize,
54+ aclOpExecutor *executor,
55+ aclrtStream stream)
56+```
38## aclnnAdaptiveMaxPool3dBackwardGetWorkspaceSize57## aclnnAdaptiveMaxPool3dBackwardGetWorkspaceSize
39 58 
40- **参数说明：**59- **参数说明：**
41- 60+
42- - gradOutput（aclTensor \*， 计算输入）：梯度Tensor，Device侧aclTensor，数据类型支持FLOAT32、FLOAT16、BFLOAT16。和正向的输出shape一致。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持ND， 当输入是5维时，内部按照NCDHW处理，当输入是4维时，在0维度处补1，内部按照NCDHW处理。61+ <table style="undefined;table-layout: fixed; width: 1478px"><colgroup>
43- 62+ <col style="width: 149px">
44- - self（aclTensor \*， 计算输入）：正向的输入Tensor，Device侧aclTensor，数据类型支持FLOAT32、FLOAT16、BFLOAT16。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持ND， 当输入是5维时，内部按照NCDHW处理，当输入是4维时，在0维度处补1，内部按照NCDHW处理，与gradOutput一致。63+ <col style="width: 121px">
45- 64+ <col style="width: 264px">
46- - indices（aclTensor \*， 计算输入）：输入Tensor，是Device侧aclTensor，数据类型仅支持INT32。正向输入中最大元素的索引位置。[数据格式](../../../docs/zh/context/数据格式.md)与gradOutput保持一致。shape与gradOutput一致65+ <col style="width: 253px">
47- 66+ <col style="width: 262px">
48- - gradInput（aclTensor \*， 计算输出）：反向输出Tensor，是Device侧aclTensor，数据类型支持FLOAT32、FLOAT16、BFLOAT16。shape与self保持一致。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)与self保持一致。67+ <col style="width: 148px">
49- 68+ <col style="width: 135px">
50- - workspaceSize（uint64_t \*， 出参）：返回需要在Device侧申请的workspace大小。69+ <col style="width: 146px">
51- - executor（aclOpExecutor \*\*， 出参）：返回op执行器，包含了算子计算流程。70+ </colgroup>
71+ <thead>
72+ <tr>
73+ <th>参数名</th>
74+ <th>输入/输出</th>
75+ <th>描述</th>
76+ <th>使用说明</th>
77+ <th>数据类型</th>
78+ <th>数据格式</th>
79+ <th>维度(shape)</th>
80+ <th>非连续Tensor</th>
81+ </tr></thead>
82+ <tbody>
83+ <tr>
84+ <td>gradOutput</td>
85+ <td>输入</td>
86+ <td>当前节点的梯度。</td>
87+ <td>和正向的输出shape一致。</td>
88+ <td>FLOAT32、FLOAT16、BFLOAT16</td>
89+ <td>ND</td>
90+ <td>4-5</td>
91+ <td>√</td>
92+ </tr>
93+ <tr>
94+ <td>self</td>
95+ <td>输入</td>
96+ <td>正向的输入Tensor。</td>
97+ <td>数据类型与gradOutput一致。</td>
98+ <td>FLOAT32、FLOAT16、BFLOAT16</td>
99+ <td>ND</td>
100+ <td>4-5</td>
101+ <td>√</td>
102+ </tr>
103+ <tr>
104+ <td>indices</td>
105+ <td>输入</td>
106+ <td>正向输入中最大元素的索引位置。</td>
107+ <td>shape与gradOutput一致。</td>
108+ <td>INT32</td>
109+ <td>ND</td>
110+ <td>4-5</td>
111+ <td>√</td>
112+ </tr>
113+ <tr>
114+ <td>gradInput</td>
115+ <td>输出</td>
116+ <td>反向输出Tensor</td>
117+ <td>shape与self保持一致。</td>
118+ <td>FLOAT32、FLOAT16、BFLOAT16</td>
119+ <td>ND</td>
120+ <td>4-5</td>
121+ <td>√</td>
122+ </tr>
123+ <tr>
124+ <td>workspaceSize</td>
125+ <td>输出</td>
126+ <td>返回需要在Device侧申请的workspace大小。</td>
127+ <td>-</td>
128+ <td>-</td>
129+ <td>-</td>
130+ <td>-</td>
131+ <td>-</td>
132+ </tr>
133+ <tr>
134+ <td>executor</td>
135+ <td>输出</td>
136+ <td>返回op执行器，包含了算子计算流程。</td>
137+ <td>-</td>
138+ <td>-</td>
139+ <td>-</td>
140+ <td>-</td>
141+ <td>-</td>
142+ </tr>
143+ </tbody></table>
52- **返回值：**144- **返回值：**
53- 145+
54 aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。146 aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
55- 147+
56- ```148+ 第一段接口完成入参校验，出现以下场景时报错：
57- 第一段接口完成入参校验，出现以下场景时报错：149+ <table style="undefined;table-layout: fixed; width: 1166px"><colgroup>
58- 161001（ACLNN_ERR_PARAM_NULLPTR）：1. 传入的gradOutput、self或indices是空指针。150+ <col style="width: 267px">
59- 161002（ACLNN_ERR_PARAM_INVALID）：1. gradOutput、self、indices、gradInput的数据类型不在支持的范围内。151+ <col style="width: 124px">
60- 2. gradOutput、self、indices、gradInput的数据格式不在支持的范围内。152+ <col style="width: 775px">
61- 3. 输入输出的shape不是4维或者5维。153+ </colgroup>
62- 4. gradOutput与indices的shape不一致，self和gradInput的shape不一致。154+ <thead>
63- 5. depth * height * width > max int32，超出了indices的表示范围。155+ <tr>
64- ```156+ <th>返回码</th>
157+ <th>错误码</th>
158+ <th>描述</th>
159+ </tr></thead>
160+ <tbody>
161+ <tr>
162+ <td>ACLNN_ERR_PARAM_NULLPTR</td>
163+ <td>161001</td>
164+ <td>传入的gradOutput、self或indices是空指针。</td>
165+ </tr>
166+ <tr>
167+ <td rowspan="5">ACLNN_ERR_PARAM_INVALID</td>
168+ <td rowspan="5">161002</td>
169+ <td>gradOutput、self、indices、gradInput的数据类型不在支持的范围内。</td>
170+ </tr>
171+ <tr>
172+ <td>gradOutput、self、indices、gradInput的数据格式不在支持的范围内。</td>
173+ </tr>
174+ <tr>
175+ <td>输入输出的shape不是4维或者5维。</td>
176+ </tr>
177+ <tr>
178+ <td>gradOutput与indices的shape不一致，self和gradInput的shape不一致。</td>
179+ </tr>
180+ <tr>
181+ <td>depth * height * width > max int32，超出了indices的表示范围。</td>
182+ </tr>
183+ </tbody>
184+ </table>
65 185 
66## aclnnAdaptiveMaxPool3dBackward186## aclnnAdaptiveMaxPool3dBackward
67 187 
68- **参数说明：**188- **参数说明：**
189+ <table style="undefined;table-layout: fixed; width: 1166px"><colgroup>
190+ <col style="width: 173px">
191+ <col style="width: 133px">
192+ <col style="width: 860px">
193+ </colgroup>
194+ <thead>
195+ <tr>
196+ <th>参数名</th>
197+ <th>输入/输出</th>
198+ <th>描述</th>
199+ </tr></thead>
200+ <tbody>
201+ <tr>
202+ <td>workspace</td>
203+ <td>输入</td>
204+ <td>在Device侧申请的workspace内存地址。</td>
205+ </tr>
206+ <tr>
207+ <td>workspaceSize</td>
208+ <td>输入</td>
209+ <td>在Device侧申请的workspace大小，由第一段接口aclnnAdaptiveMaxPool3dBackwardGetWorkspaceSize获取。</td>
210+ </tr>
211+ <tr>
212+ <td>executor</td>
213+ <td>输入</td>
214+ <td>op执行器，包含了算子计算流程。</td>
215+ </tr>
216+ <tr>
217+ <td>stream</td>
218+ <td>输入</td>
219+ <td>指定执行任务的Stream。</td>
220+ </tr>
221+ </tbody>
222+ </table>
223+- **返回值：**
69 224 
70- - workspace（void \*， 入参）：在Device侧申请的workspace内存地址。225+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
71- - workspaceSize（uint64_t， 入参）：在Device侧申请的workspace大小，由第一段接口aclnnAdaptiveMaxPool3dBackwardGetWorkspaceSize获取。
72- - executor（aclOpExecutor \*， 入参）：op执行器，包含了算子计算流程。
73- - stream（aclrtStream， 入参）：指定执行任务的Stream。
74-- **返回值：**
75- 
76- aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
77 226 
78## 约束说明227## 约束说明
79- 确定性计算：228- 确定性计算：
Mpooling/adaptive_max_pool3d_grad/examples/test_aclnn_adaptive_max_pool3d_backward.cpp+4-4
@@ -1,10 +1,10 @@
1/**1/**
2+ * This program is free software, you can redistribute it and/or modify.
2 * Copyright (c) 2025 Huawei Technologies Co., Ltd.3 * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3- * This program is free software, you can redistribute it and/or modify it under the terms and conditions of4+ * This file is a part of the CANN Open Software.
4- * CANN Open Software License Agreement Version 2.0 (the "License").5+ * Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.6 * Please refer to the License for details. You may not use this file except in compliance with the License.
6- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,7+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
7- * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8 * See LICENSE in the root of the software repository for the full text of the License.8 * See LICENSE in the root of the software repository for the full text of the License.
9 */9 */
10 10 
Rpooling/adaptive_max_pool3d_grad/op_host/op_api/aclnn_adaptive_max_pool2d_backward.cpp→pooling/adaptive_max_pool3d_grad/op_api/aclnn_adaptive_max_pool2d_backward.cpp+0-0
文件重命名但无更改。
Rpooling/adaptive_max_pool3d_grad/op_host/op_api/aclnn_adaptive_max_pool2d_backward.h→pooling/adaptive_max_pool3d_grad/op_api/aclnn_adaptive_max_pool2d_backward.h+0-0
文件重命名但无更改。
Rpooling/adaptive_max_pool3d_grad/op_host/op_api/aclnn_adaptive_max_pool3d_backward.cpp→pooling/adaptive_max_pool3d_grad/op_api/aclnn_adaptive_max_pool3d_backward.cpp+0-0
文件重命名但无更改。
Rpooling/adaptive_max_pool3d_grad/op_host/op_api/aclnn_adaptive_max_pool3d_backward.h→pooling/adaptive_max_pool3d_grad/op_api/aclnn_adaptive_max_pool3d_backward.h+0-0
文件重命名但无更改。
Rpooling/adaptive_max_pool3d_grad/op_host/op_api/adaptive_max_pool3d_backward.cpp→pooling/adaptive_max_pool3d_grad/op_api/adaptive_max_pool3d_backward.cpp+0-0
文件重命名但无更改。
Rpooling/adaptive_max_pool3d_grad/op_host/op_api/adaptive_max_pool3d_backward.h→pooling/adaptive_max_pool3d_grad/op_api/adaptive_max_pool3d_backward.h+0-0
文件重命名但无更改。
Mpooling/adaptive_max_pool3d_grad/op_host/adaptive_max_pool3d_grad_normal_tiling.cpp+1-0
@@ -196,6 +196,7 @@ ge::graphStatus AdaptiveMaxPool3DGradNormalTiling::DoOpTiling()
196 bool res = SetNormalTilingParams();196 bool res = SetNormalTilingParams();
197 OP_CHECK_IF(!res, OP_LOGE(context_->GetNodeName(), "Normal cal tiling params failed."), return ge::GRAPH_FAILED);197 OP_CHECK_IF(!res, OP_LOGE(context_->GetNodeName(), "Normal cal tiling params failed."), return ge::GRAPH_FAILED);
198 maxPoolGradParams.tilingType = TILING_TYPE_NORMAL;198 maxPoolGradParams.tilingType = TILING_TYPE_NORMAL;
199+ context_->SetScheduleMode(1);
CANN-robot
CANN-robotCANN-robot1月19日

魔数与常量: 代码中直接使用了魔数1作为SetScheduleMode的参数，缺乏明确的含义说明。在代码检视的上下文中，虽然只关注第199行的改动，但结合前后代码分析，SetScheduleMode(1)中的参数1是一个硬编码的魔法数字，没有定义为有意义的常量。这降低了代码的可读性和可维护性，其他开发者难以理解参数1的具体含义。

问题类型: 魔数与常量 文件路径: pooling/adaptive_max_pool3d_grad/op_host/adaptive_max_pool3d_grad_normal_tiling.cpp 行号: 199 问题代码:

    context_->SetScheduleMode(1);

修改建议:

将魔法数字1定义为有意义的常量，例如：constexpr int SCHEDULE_MODE_NORMAL = 1; 并在调用时使用该常量：context_->SetScheduleMode(SCHEDULE_MODE_NORMAL);。同时建议添加注释说明该调度模式的具体含义。

此评论由代码审查工具自动生成

likedislike
199 SetOtherTilingParams();200 SetOtherTilingParams();
200 SetBaseTilingData();201 SetBaseTilingData();
201 SetNormalTilingData();202 SetNormalTilingData();
Mpooling/adaptive_max_pool3d_grad/op_host/adaptive_max_pool3d_grad_scatter_tiling.cpp+21-88
@@ -14,6 +14,7 @@
14 */14 */
15#include "adaptive_max_pool3d_grad_tiling.h"15#include "adaptive_max_pool3d_grad_tiling.h"
16#include <iostream>16#include <iostream>
17+#include "../../pool_3d_common/op_host/arch32/max_pool3d_grad_tiling_common.h"
17 18 
18namespace optiling {19namespace optiling {
19ge::graphStatus AdaptiveMaxPool3DGradScatterTiling::GetShapeAttrsInfo()20ge::graphStatus AdaptiveMaxPool3DGradScatterTiling::GetShapeAttrsInfo()
@@ -32,106 +33,36 @@ bool AdaptiveMaxPool3DGradScatterTiling::IsCapable()
32 33 
33bool AdaptiveMaxPool3DGradScatterTiling::SetScatterTilingParams()34bool AdaptiveMaxPool3DGradScatterTiling::SetScatterTilingParams()
34{35{
35- const uint64_t doDim = maxPoolGradParams.doDim;36+ return CalculateScatterTilingParams(
36- const uint64_t hoDim = maxPoolGradParams.hoDim;37+ maxPoolGradParams,
37- const uint64_t woDim = maxPoolGradParams.woDim;38+ maxPoolGradParams.doDim,
38- const uint64_t xDtypeSize = maxPoolGradParams.xDtypeSize;39+ maxPoolGradParams.hoDim,
39- const uint64_t indexDtypeSize = maxPoolGradParams.indexDtypeSize;40+ maxPoolGradParams.woDim,
40- 41+ maxPoolGradParams.xDtypeSize,
41- uint64_t ncPreCore = Ops::Base::CeilDiv(maxPoolGradParams.ncDim, maxPoolGradParams.totalCoreNum);42+ maxPoolGradParams.indexDtypeSize,
42- maxPoolGradParams.usedCoreNum = Ops::Base::CeilDiv(maxPoolGradParams.ncDim, ncPreCore);43+ MAX_BLOCK_COUNT);
43- 
44- // Scatter main tiling cal
45- // 1. All Tensor full size, cut nc between cores, without cut in one core
46- uint64_t noCutSize = ncPreCore * doDim * hoDim * woDim * (xDtypeSize + indexDtypeSize);
47- if (noCutSize <= maxPoolGradParams.maxUbSize) {
48- maxPoolGradParams.baseNc = ncPreCore;
49- maxPoolGradParams.baseDo = doDim;
50- maxPoolGradParams.baseHo = hoDim;
51- maxPoolGradParams.baseWo = woDim;
52- maxPoolGradParams.ubCutAxis = TILING_UB_NO_CUT;
53- return true;
54- }
55- 
56- // 2. Cut nc
57- uint64_t perNcSize = 1UL * doDim * hoDim * woDim * (xDtypeSize + indexDtypeSize);
58- if (perNcSize <= maxPoolGradParams.maxUbSize) {
59- uint64_t baseNc = maxPoolGradParams.maxUbSize / perNcSize;
60- if (baseNc > MAX_BLOCK_COUNT) { // Use NC for blockCount.
61- baseNc = MAX_BLOCK_COUNT;
62- }
63- maxPoolGradParams.baseNc = baseNc;
64- maxPoolGradParams.baseDo = doDim;
65- maxPoolGradParams.baseHo = hoDim;
66- maxPoolGradParams.baseWo = woDim;
67- maxPoolGradParams.ubCutAxis = TILING_UB_CUT_NC;
68- return true;
69- }
70- maxPoolGradParams.baseNc = 1UL;
71- 
72- // 3. Cut do
73- uint64_t perDoSize = 1UL * 1UL * hoDim * woDim * (xDtypeSize + indexDtypeSize);
74- if (perDoSize <= maxPoolGradParams.maxUbSize) {
75- maxPoolGradParams.baseDo = maxPoolGradParams.maxUbSize / perDoSize;
76- maxPoolGradParams.baseHo = hoDim;
77- maxPoolGradParams.baseWo = woDim;
78- maxPoolGradParams.ubCutAxis = TILING_UB_CUT_DO;
79- return true;
80- }
81- maxPoolGradParams.baseDo = 1UL;
82- 
83- // 4. Cut ho
84- uint64_t perHoSize = 1UL * 1UL * 1UL * woDim * (xDtypeSize + indexDtypeSize);
85- if (perHoSize <= maxPoolGradParams.maxUbSize) {
86- maxPoolGradParams.baseHo = maxPoolGradParams.maxUbSize / perHoSize;
87- maxPoolGradParams.baseWo = woDim;
88- maxPoolGradParams.ubCutAxis = TILING_UB_CUT_HO;
89- return true;
90- }
91- maxPoolGradParams.baseHo = 1UL;
92- 
93- // 5. Cut wo
94- uint64_t perWoSize = 1UL * 1UL * 1UL * 1UL * (xDtypeSize + indexDtypeSize);
95- if (perWoSize <= maxPoolGradParams.maxUbSize) {
96- maxPoolGradParams.baseWo = maxPoolGradParams.maxUbSize / perWoSize;
97- maxPoolGradParams.ubCutAxis = TILING_UB_CUT_WO;
98- return true;
99- }
100- maxPoolGradParams.baseWo = 1UL;
101- 
102- return false;
103}44}
104 45 
105void AdaptiveMaxPool3DGradScatterTiling::SetOtherTilingParams()46void AdaptiveMaxPool3DGradScatterTiling::SetOtherTilingParams()
106{47{
107 SetCntTailTilingParams();48 SetCntTailTilingParams();
108- maxPoolGradParams.ncRound = Ops::Base::CeilDiv(maxPoolGradParams.ncCnt, maxPoolGradParams.usedCoreNum);49+
109- maxPoolGradParams.preCoreNum = maxPoolGradParams.ncCnt % maxPoolGradParams.usedCoreNum;50+ CalculateRoundParams(
110- maxPoolGradParams.ncRoundTail =51+ maxPoolGradParams,
111- maxPoolGradParams.preCoreNum == 0UL ? maxPoolGradParams.ncRound : maxPoolGradParams.ncRound - 1UL;52+ maxPoolGradParams.isOverLap,
112- maxPoolGradParams.totalRound =53+ maxPoolGradParams.diDim,
113- maxPoolGradParams.ncRound * maxPoolGradParams.doCnt * maxPoolGradParams.hoCnt * maxPoolGradParams.woCnt;54+ maxPoolGradParams.hiDim,
114- if (maxPoolGradParams.xDtypeSize != DTYPE_LEN_B32 && maxPoolGradParams.isOverLap) {55+ maxPoolGradParams.wiDim);
115- maxPoolGradParams.workspaceSize = maxPoolGradParams.ncDim * maxPoolGradParams.diDim * maxPoolGradParams.hiDim *
116- maxPoolGradParams.wiDim * sizeof(float);
117- } else {
118- maxPoolGradParams.workspaceSize = 0UL;
119- }
120}56}
121 57 
122void AdaptiveMaxPool3DGradScatterTiling::SetScatterTilingData()58void AdaptiveMaxPool3DGradScatterTiling::SetScatterTilingData()
123{59{
124- tilingData.set_ncRound(maxPoolGradParams.ncRound);60+ SetScatterTilingDataCommon(tilingData, maxPoolGradParams);
125- tilingData.set_ncRoundTail(maxPoolGradParams.ncRoundTail);
126- tilingData.set_totalRound(maxPoolGradParams.totalRound);
127- tilingData.set_preCoreNum(maxPoolGradParams.preCoreNum);
128}61}
129 62 
130void AdaptiveMaxPool3DGradScatterTiling::PrintScatterTilingData()63void AdaptiveMaxPool3DGradScatterTiling::PrintScatterTilingData()
131{64{
132- OP_LOGI(65+ PrintScatterTilingDataCommon(context_->GetNodeName(), tilingData);
133- context_->GetNodeName(), "TilingData ncRound: %lu, ncRoundTail: %lu, totalRound: %lu.",
134- tilingData.get_ncRound(), tilingData.get_ncRoundTail(), tilingData.get_totalRound());
135}66}
136 67 
137ge::graphStatus AdaptiveMaxPool3DGradScatterTiling::DoOpTiling()68ge::graphStatus AdaptiveMaxPool3DGradScatterTiling::DoOpTiling()
@@ -144,8 +75,10 @@ ge::graphStatus AdaptiveMaxPool3DGradScatterTiling::DoOpTiling()
144 SetScatterTilingData();75 SetScatterTilingData();
145 PrintTilingData();76 PrintTilingData();
146 PrintScatterTilingData();77 PrintScatterTilingData();
78+
147 return ge::GRAPH_SUCCESS;79 return ge::GRAPH_SUCCESS;
148}80}
149 81 
82+// 注册
150REGISTER_TILING_TEMPLATE("AdaptiveMaxPool3DGrad", AdaptiveMaxPool3DGradScatterTiling, 2);83REGISTER_TILING_TEMPLATE("AdaptiveMaxPool3DGrad", AdaptiveMaxPool3DGradScatterTiling, 2);
151-} // namespace optiling84+} // namespace optiling
Mpooling/adaptive_max_pool3d_grad/op_host/adaptive_max_pool3d_grad_tiling.h+3-40
@@ -11,12 +11,6 @@
11/*!11/*!
12 * \file adaptive_max_pool3d_grad_tiling.h12 * \file adaptive_max_pool3d_grad_tiling.h
13 * \brief13 * \brief
14- * ATTENTION: MAKE SURE 'BEGIN_TILING_DATA_DEF' STAY IN THE SAME LINE (35) USING BLANK LINES.
15- *
16- *
17- *
18- *
19- *
20 */14 */
21 15 
22#ifndef OPS_BUILD_IN_OP_TILING_RUNTIME_ADAPTIVE_MAX_POOL3D_GRAD_H16#ifndef OPS_BUILD_IN_OP_TILING_RUNTIME_ADAPTIVE_MAX_POOL3D_GRAD_H
@@ -28,6 +22,7 @@
28#include "tiling_base/tiling_base.h"22#include "tiling_base/tiling_base.h"
29#include "tiling_base/tiling_templates_registry.h"23#include "tiling_base/tiling_templates_registry.h"
30#include "util/math_util.h"24#include "util/math_util.h"
25+#include "../../pool_3d_common/op_host/arch32/max_pool3d_grad_tiling_constants.h"
31 26 
32namespace optiling {27namespace optiling {
33using Ops::NN::Optiling::TilingBaseClass;28using Ops::NN::Optiling::TilingBaseClass;
@@ -72,40 +67,7 @@ END_TILING_DATA_DEF;
72 67 
73REGISTER_TILING_DATA_CLASS(AdaptiveMaxPool3DGrad, AdaptiveMaxPool3DGradTilingData)68REGISTER_TILING_DATA_CLASS(AdaptiveMaxPool3DGrad, AdaptiveMaxPool3DGradTilingData)
74 69 
75-// Index const
76-constexpr uint32_t X_INDEX = 0;
77-constexpr uint32_t GRAD_INDEX = 1;
78-constexpr uint32_t ARGMAX_INDEX = 2;
79-constexpr size_t KSIZE_ATTR_INDEX = 0U;
80-constexpr size_t STRIDES_ATTR_INDEX = 1U;
81-constexpr size_t PADS_ATTR_INDEX = 2U;
82-constexpr size_t DILATION_ATTR_INDEX = 3U;
83-constexpr size_t CEIL_MODE_ATTR_INDEX = 4U;
84-// Params const
85-constexpr size_t NC_DIM_NUM = 2;
86-constexpr size_t NCDHW_DIM_NUM = 5;
87-constexpr uint32_t DTYPE_LEN_B8 = 1;
88-constexpr uint32_t DTYPE_LEN_B16 = 2;
89-constexpr uint32_t DTYPE_LEN_B32 = 4;
90-constexpr uint32_t BLOCK_SIZE = 32;
91constexpr uint32_t MAX_BLOCK_COUNT = 65535;70constexpr uint32_t MAX_BLOCK_COUNT = 65535;
92-constexpr uint32_t NUM_PER_REP_B16 = 128;
93-constexpr uint32_t NUM_PER_REP_B32 = 64;
94-constexpr uint32_t SELECT_RESERVED_UB_SIZE = 8192;
95-constexpr uint64_t MAX_INT32 = 2147483647;
96-// Tiling const
97-constexpr uint32_t TILING_OVERLAP = 100;
98-constexpr uint32_t TILING_UB_NO_CUT = 0;
99-constexpr uint32_t TILING_UB_CUT_NC = 10;
100-constexpr uint32_t TILING_UB_CUT_DO = 20;
101-constexpr uint32_t TILING_UB_CUT_HO = 30;
102-constexpr uint32_t TILING_UB_CUT_WO = 40;
103-constexpr uint32_t TILING_UB_CUT_KD = 50;
104-constexpr uint32_t TILING_UB_CUT_KH = 60;
105-constexpr uint32_t TILING_UB_CUT_KW = 70;
106-constexpr uint32_t TILING_TYPE_NORMAL = 0;
107-constexpr uint32_t TILING_TYPE_CUTK = 1;
108-constexpr uint32_t TILING_TYPE_SCATTER = 2;
109 71 
110/*72/*
111 * @brief: get m and n greatest common divisor73 * @brief: get m and n greatest common divisor
@@ -266,5 +228,6 @@ private:
266 void SetScatterTilingData();228 void SetScatterTilingData();
267 void PrintScatterTilingData();229 void PrintScatterTilingData();
268};230};
231+ 
269} // namespace optiling232} // namespace optiling
270-#endif // OPS_BUILD_IN_OP_TILING_RUNTIME_ADAPTIVE_MAX_POOL3D_GRAD_H233+#endif // OPS_BUILD_IN_OP_TILING_RUNTIME_ADAPTIVE_MAX_POOL3D_GRAD_H
Mpooling/adaptive_max_pool3d_grad/op_host/adaptive_max_pool3d_grad_tiling_base.cpp+0-1
@@ -224,7 +224,6 @@ ge::graphStatus AdaptiveMaxPool3DGradTilingBase::PostTiling()
224 context_->SetBlockDim(maxPoolGradParams.usedCoreNum);224 context_->SetBlockDim(maxPoolGradParams.usedCoreNum);
225 tilingData.SaveToBuffer(context_->GetRawTilingData()->GetData(), context_->GetRawTilingData()->GetCapacity());225 tilingData.SaveToBuffer(context_->GetRawTilingData()->GetData(), context_->GetRawTilingData()->GetCapacity());
226 context_->GetRawTilingData()->SetDataSize(tilingData.GetDataSize());226 context_->GetRawTilingData()->SetDataSize(tilingData.GetDataSize());
227- 
228 size_t usrWorkspaceSize = maxPoolGradParams.workspaceSize;227 size_t usrWorkspaceSize = maxPoolGradParams.workspaceSize;
229 size_t sysWorkSpaceSize = 16 * 1024 * 1024;228 size_t sysWorkSpaceSize = 16 * 1024 * 1024;
230 size_t* currentWorkspace = context_->GetWorkspaceSizes(1);229 size_t* currentWorkspace = context_->GetWorkspaceSizes(1);
Mpooling/adaptive_max_pool3d_grad/op_kernel/adaptive_max_pool3d_grad.cpp+1-9
@@ -28,15 +28,7 @@ using namespace AdaptiveMaxPool3DGrad;
28extern "C" __global__ __aicore__ void adaptive_max_pool3d_grad(28extern "C" __global__ __aicore__ void adaptive_max_pool3d_grad(
29 GM_ADDR x, GM_ADDR grad, GM_ADDR argmax, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)29 GM_ADDR x, GM_ADDR grad, GM_ADDR argmax, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
30{30{
31- if (workspace == nullptr) {31+ if (workspace == nullptr || GetUserWorkspace(workspace) == nullptr || g_coreType == AIC) {
32- return;
33- }
34- 
35- GM_ADDR userWS = GetUserWorkspace(workspace);
36- if (userWS == nullptr) {
37- return;
38- }
39- if (g_coreType == AIC) {
40 return;32 return;
41 }33 }
42 TPipe pipe;34 TPipe pipe;
Mpooling/adaptive_max_pool3d_grad/op_kernel/adaptive_max_pool3d_grad_common.h+11-179
@@ -15,48 +15,20 @@
15 15 
16#ifndef ADAPTIVE_MAX_POOL3D_GRAD_COMMON_H16#ifndef ADAPTIVE_MAX_POOL3D_GRAD_COMMON_H
17#define ADAPTIVE_MAX_POOL3D_GRAD_COMMON_H17#define ADAPTIVE_MAX_POOL3D_GRAD_COMMON_H
18+ 
19+#include "../pool_3d_common/arch32/max_pool3d_grad_common.h"
18#include "kernel_tiling/kernel_tiling.h"20#include "kernel_tiling/kernel_tiling.h"
19 21 
20namespace AdaptiveMaxPool3DGradComm {22namespace AdaptiveMaxPool3DGradComm {
21using namespace AscendC;23using namespace AscendC;
22- 24+using namespace MaxPool3DGradCommon;
23-constexpr uint64_t TRANS_ADDR_LEN = 16;25+struct BlockParams : public BlockParamsCommon {
24-constexpr uint64_t UINT8_BITS = 8;
25-constexpr uint64_t BLOCK_SIZE = 32;
26-constexpr uint64_t BLOCK_NUM_16 = BLOCK_SIZE / sizeof(half);
27-constexpr uint64_t BLOCK_NUM_32 = BLOCK_SIZE / sizeof(float);
28-constexpr uint64_t UNIT_BLOCK_LEN = BLOCK_SIZE / sizeof(float);
29-constexpr float ZERO = 0.0f;
30- 
31-struct BlockParams {
32- uint64_t ncCntIndex = 0;
33- uint64_t doCntIndex = 0;
34- uint64_t hoCntIndex = 0;
35- uint64_t woCntIndex = 0;
36- uint64_t ncShape = 0;
37- uint64_t doShape = 0;
38- uint64_t hoShape = 0;
39- uint64_t woShape = 0;
40- uint64_t diShape = 0;
41- uint64_t hiShape = 0;
42- uint64_t wiShape = 0;
43 float coeffD = 0.0;26 float coeffD = 0.0;
44 float coeffH = 0.0;27 float coeffH = 0.0;
45 float hwDims = 0.0;28 float hwDims = 0.0;
46 uint64_t maxKwAlign8 = 0;29 uint64_t maxKwAlign8 = 0;
47 uint64_t maxKwAlign16 = 0;30 uint64_t maxKwAlign16 = 0;
48 uint64_t maxKwAlignDtype = 0;31 uint64_t maxKwAlignDtype = 0;
49- uint64_t dihiwiAlign = 0;
50- uint64_t diValid = 0;
51- uint64_t hiValid = 0;
52- uint64_t wiValid = 0;
53- uint64_t dohowoShape = 0;
54- uint64_t dohowoAlign8 = 0;
55- uint64_t dohowoAlign16 = 0;
56- uint64_t offsetX = 0;
57- uint64_t offsetGrad = 0;
58- uint64_t offsetArgmax = 0;
59- uint64_t offsetY = 0;
60 uint64_t startD = 0;32 uint64_t startD = 0;
61 uint64_t startH = 0;33 uint64_t startH = 0;
62 uint64_t startW = 0;34 uint64_t startW = 0;
@@ -64,162 +36,22 @@ struct BlockParams {
64 uint64_t deltaH = 0;36 uint64_t deltaH = 0;
65 uint64_t deltaW = 0;37 uint64_t deltaW = 0;
66 uint64_t kernelShape = 0;38 uint64_t kernelShape = 0;
67- uint64_t baseNcOffset = 0;
68- uint64_t ShapeSum = 0;
69};39};
70 40 
71-struct TilingParams {41+struct TilingParams : public TilingParamsCommon {
72- uint64_t ncDim;42+ uint64_t baseNcTail = 0;
73- uint64_t diDim;43+ uint64_t maxKd = 0;
74- uint64_t hiDim;44+ uint64_t maxKh = 0;
75- uint64_t wiDim;45+ uint64_t maxKw = 0;
76- uint64_t doDim;46+ uint64_t maxKdhwLen = 0;
77- uint64_t hoDim;47+ uint64_t addMode = 0;
78- uint64_t woDim;
79- uint64_t singleCoreNc;
80- uint64_t singleCoreDo;
81- uint64_t singleCoreHo;
82- uint64_t singleCoreWo;
83- uint64_t baseNc;
84- uint64_t baseDo;
85- uint64_t baseHo;
86- uint64_t baseWo;
87- uint64_t ncCnt;
88- uint64_t dCnt;
89- uint64_t hCnt;
90- uint64_t wCnt;
91- uint64_t baseNcTail;
92- uint64_t ncTail;
93- uint64_t doTail;
94- uint64_t hoTail;
95- uint64_t woTail;
96- uint64_t totalCnt;
97- uint64_t maxKd;
98- uint64_t maxKh;
99- uint64_t maxKw;
100- uint64_t maxKdhwLen;
101- uint64_t needInitOutput;
102- uint64_t usedCoreNum;
103- uint64_t preCoreNum;
104- uint64_t round;
105- uint64_t realRound;
106- uint64_t addMode;
107- uint64_t ncIndex;
108- uint64_t ncCntRound;
109- uint64_t ncRealRound;
110- uint64_t diHiWiLen;
111- uint64_t ubSize;
112- uint64_t initLen;
113- uint64_t initOffset;
114};48};
115 49 
116-template <typename Tp, Tp v>
117-struct integral_constant {
118- static constexpr Tp value = v;
119-};
120-using true_type = integral_constant<bool, true>;
121-using false_type = integral_constant<bool, false>;
122-template <typename, typename>
123-struct is_same : public false_type {
124-};
125-template <typename Tp>
126-struct is_same<Tp, Tp> : public true_type {
127-};
128- 
129-__aicore__ inline uint64_t CeilDiv(uint64_t x, uint64_t y)
130-{
131- return y == 0 ? x : (x + y - 1) / y;
132-}
133- 
134__aicore__ inline uint64_t FloorDiv(uint64_t x, uint64_t y)50__aicore__ inline uint64_t FloorDiv(uint64_t x, uint64_t y)
135{51{
136 return y == 0 ? x : (uint64_t)(x / y);52 return y == 0 ? x : (uint64_t)(x / y);
137}53}
138 54 
139-// only support float/int32_t
140-// [row, col] -> [col, row]: row:align16, col:align8
141-template <typename T>
142-__aicore__ inline void TransposeBase16M8(LocalTensor<T>& dstUb, LocalTensor<T>& srcUb, uint64_t rowNum, uint64_t colNum)
143-{
144- uint64_t srcAddrList[TRANS_ADDR_LEN];
145- uint64_t dstAddrList[TRANS_ADDR_LEN];
146- 
147- for (uint64_t r = 0; r < rowNum / TRANS_ADDR_LEN; r++) {
148- for (uint64_t i = 0; i < TRANS_ADDR_LEN; i++) {
149- srcAddrList[i] = (uint64_t)(srcUb[r * TRANS_ADDR_LEN * colNum + i * colNum].GetPhyAddr());
150- dstAddrList[i] = (uint64_t)(dstUb[r * TRANS_ADDR_LEN + i / 2 * rowNum + i % 2 * BLOCK_NUM_32].GetPhyAddr());
151- }
152- struct TransDataTo5HDParams transDataParams;
153- transDataParams.repeatTimes = colNum / BLOCK_NUM_32;
154- if (transDataParams.repeatTimes == 1) {
155- transDataParams.srcRepStride = 0;
156- transDataParams.dstRepStride = 0;
157- } else {
158- transDataParams.srcRepStride = 1;
159- transDataParams.dstRepStride = rowNum;
160- }
161- 
162- TransDataTo5HD<float>(dstAddrList, srcAddrList, transDataParams);
163- }
164-}
165- 
166-// only support float/int32_t
167-// [row, col] -> [col, row]: row:align8, col:align16
168-template <typename T>
169-__aicore__ inline void TransposeBase8M16(LocalTensor<T>& dstUb, LocalTensor<T>& srcUb, uint64_t rowNum, uint64_t colNum)
170-{
171- uint64_t srcAddrList[TRANS_ADDR_LEN];
172- uint64_t dstAddrList[TRANS_ADDR_LEN];
173- 
174- for (uint64_t r = 0; r < colNum / TRANS_ADDR_LEN; r++) {
175- for (uint64_t i = 0; i < TRANS_ADDR_LEN; i++) {
176- srcAddrList[i] =
177- (uint64_t)(srcUb[r * TRANS_ADDR_LEN + i % BLOCK_NUM_32 * colNum + i / BLOCK_NUM_32 * BLOCK_NUM_32]
178- .GetPhyAddr());
179- dstAddrList[i] =
180- (uint64_t)(dstUb[r * TRANS_ADDR_LEN * rowNum + (i % 2 * BLOCK_NUM_32 + i / 2) * rowNum].GetPhyAddr());
181- }
182- struct TransDataTo5HDParams transDataParams;
183- transDataParams.repeatTimes = rowNum / BLOCK_NUM_32;
184- if (transDataParams.repeatTimes == 1) {
185- transDataParams.srcRepStride = 0;
186- transDataParams.dstRepStride = 0;
187- } else {
188- transDataParams.srcRepStride = colNum;
189- transDataParams.dstRepStride = 1;
190- }
191- 
192- TransDataTo5HD<float>(dstAddrList, srcAddrList, transDataParams);
193- }
194-}
195- 
196-// only support float16/bfloat16
197-// [row, col] -> [col, row]: row:align16, col:align16
198-template <typename T>
199-__aicore__ inline void TransposeBase16M16(
200- LocalTensor<T>& dstUb, LocalTensor<T>& srcUb, uint64_t rowNum, uint64_t colNum)
201-{
202- uint64_t srcAddrList[TRANS_ADDR_LEN];
203- uint64_t dstAddrList[TRANS_ADDR_LEN];
204- 
205- for (uint64_t r = 0; r < rowNum / TRANS_ADDR_LEN; r++) {
206- for (uint64_t i = 0; i < TRANS_ADDR_LEN; i++) {
207- srcAddrList[i] = (uint64_t)(srcUb[r * TRANS_ADDR_LEN * colNum + i * colNum].GetPhyAddr());
208- dstAddrList[i] = (uint64_t)(dstUb[r * TRANS_ADDR_LEN + i * rowNum].GetPhyAddr());
209- }
210- struct TransDataTo5HDParams transDataParams;
211- transDataParams.repeatTimes = colNum / BLOCK_NUM_16;
212- if (transDataParams.repeatTimes == 1) {
213- transDataParams.srcRepStride = 0;
214- transDataParams.dstRepStride = 0;
215- } else {
216- transDataParams.srcRepStride = 1;
217- transDataParams.dstRepStride = rowNum;
218- }
219- TransDataTo5HD<half>(dstAddrList, srcAddrList, transDataParams);
220- }
221-}
222- 
223} // namespace AdaptiveMaxPool3DGradComm55} // namespace AdaptiveMaxPool3DGradComm
224 56 
225#endif // ADAPTIVE_MAX_POOL3D_GRAD_COMMON_H57#endif // ADAPTIVE_MAX_POOL3D_GRAD_COMMON_H
Mpooling/adaptive_max_pool3d_grad/op_kernel/adaptive_max_pool3d_grad_normal.h+18-14
@@ -59,9 +59,9 @@ public:
59 params_.baseHo = tiling->baseHo;59 params_.baseHo = tiling->baseHo;
60 params_.baseWo = tiling->baseWo;60 params_.baseWo = tiling->baseWo;
61 params_.ncCnt = tiling->ncCnt;61 params_.ncCnt = tiling->ncCnt;
62- params_.dCnt = tiling->doCnt;62+ params_.doCnt = tiling->doCnt;
63- params_.hCnt = tiling->hoCnt;63+ params_.hoCnt = tiling->hoCnt;
64- params_.wCnt = tiling->woCnt;64+ params_.woCnt = tiling->woCnt;
65 params_.totalCnt = tiling->totalCnt;65 params_.totalCnt = tiling->totalCnt;
66 params_.baseNcTail = tiling->ncTail;66 params_.baseNcTail = tiling->ncTail;
67 params_.doTail = tiling->doTail;67 params_.doTail = tiling->doTail;
@@ -78,10 +78,14 @@ public:
78 78 
79 __aicore__ inline void InitInputsOutputs(GM_ADDR x, GM_ADDR grad, GM_ADDR argmax, GM_ADDR y, GM_ADDR usrWorkspace)79 __aicore__ inline void InitInputsOutputs(GM_ADDR x, GM_ADDR grad, GM_ADDR argmax, GM_ADDR y, GM_ADDR usrWorkspace)
80 {80 {
81- gradGm.SetGlobalBuffer((__gm__ TGrad*)grad);81+ gradGm.SetGlobalBuffer((__gm__ TGrad*)grad, params_.ncDim * params_.doDim * params_.hoDim * params_.woDim);
82- argmaxGm.SetGlobalBuffer((__gm__ TArgmax*)argmax);82+ argmaxGm.SetGlobalBuffer((__gm__ TArgmax*)argmax, params_.ncDim * params_.doDim * params_.hoDim * params_.woDim);
83- yGm.SetGlobalBuffer((__gm__ TY*)y);83+ yGm.SetGlobalBuffer((__gm__ TY*)y, params_.ncDim * params_.diHiWiLen);
84- workspaceGm.SetGlobalBuffer((__gm__ float*)usrWorkspace);84+ if constexpr (!is_same<TY, float>::value && IsOverlap) {
85+ workspaceGm.SetGlobalBuffer((__gm__ float*)usrWorkspace, params_.ncDim * params_.diHiWiLen);
86+ } else {
87+ workspaceGm.SetGlobalBuffer((__gm__ float*)usrWorkspace);
88+ }
85 if (GetBlockIdx() == 0) {89 if (GetBlockIdx() == 0) {
86 if constexpr (is_same<TY, float>::value) {90 if constexpr (is_same<TY, float>::value) {
87 InitGlobalMemory(yGm, params_.ncDim * params_.diHiWiLen, 0.0f);91 InitGlobalMemory(yGm, params_.ncDim * params_.diHiWiLen, 0.0f);
@@ -143,14 +147,14 @@ public:
143 147 
144 for (uint64_t totalIndex = 0; totalIndex < params_.totalCnt; totalIndex++) {148 for (uint64_t totalIndex = 0; totalIndex < params_.totalCnt; totalIndex++) {
145 if (GetBlockIdx() == totalIndex % GetBlockNum()) {149 if (GetBlockIdx() == totalIndex % GetBlockNum()) {
146- core_.ncCntIndex = totalIndex / (params_.dCnt * params_.hCnt * params_.wCnt);150+ core_.ncCntIndex = totalIndex / (params_.doCnt * params_.hoCnt * params_.woCnt);
147- core_.doCntIndex = totalIndex / (params_.hCnt * params_.wCnt) % params_.dCnt;151+ core_.doCntIndex = totalIndex / (params_.hoCnt * params_.woCnt) % params_.doCnt;
148- core_.hoCntIndex = totalIndex / params_.wCnt % params_.hCnt;152+ core_.hoCntIndex = totalIndex / params_.woCnt % params_.hoCnt;
149- core_.woCntIndex = totalIndex % params_.wCnt;153+ core_.woCntIndex = totalIndex % params_.woCnt;
150 core_.ncShape = core_.ncCntIndex == params_.ncCnt - 1 ? params_.baseNcTail : params_.singleCoreNc;154 core_.ncShape = core_.ncCntIndex == params_.ncCnt - 1 ? params_.baseNcTail : params_.singleCoreNc;
151- core_.doShape = core_.doCntIndex == params_.dCnt - 1 ? params_.doTail : params_.singleCoreDo;155+ core_.doShape = core_.doCntIndex == params_.doCnt - 1 ? params_.doTail : params_.singleCoreDo;
152- core_.hoShape = core_.hoCntIndex == params_.hCnt - 1 ? params_.hoTail : params_.singleCoreHo;156+ core_.hoShape = core_.hoCntIndex == params_.hoCnt - 1 ? params_.hoTail : params_.singleCoreHo;
153- core_.woShape = core_.woCntIndex == params_.wCnt - 1 ? params_.woTail : params_.singleCoreWo;157+ core_.woShape = core_.woCntIndex == params_.woCnt - 1 ? params_.woTail : params_.singleCoreWo;
154 SubProcess();158 SubProcess();
155 }159 }
156 }160 }
Mpooling/adaptive_max_pool3d_grad/op_kernel/adaptive_max_pool3d_grad_scatter.h+15-102
@@ -15,119 +15,32 @@
15 15 
16#ifndef ADAPTIVE_MAX_POOL3D_GRAD_SCATTER_H16#ifndef ADAPTIVE_MAX_POOL3D_GRAD_SCATTER_H
17#define ADAPTIVE_MAX_POOL3D_GRAD_SCATTER_H17#define ADAPTIVE_MAX_POOL3D_GRAD_SCATTER_H
18-#include "kernel_tiling/kernel_tiling.h"
19-#include "kernel_operator.h"
20-#include "adaptive_max_pool3d_grad_common.h"
21#include "adaptive_max_pool3d_grad_scatter_base.h"18#include "adaptive_max_pool3d_grad_scatter_base.h"
19+#include "adaptive_max_pool3d_grad_common.h"
20+#include "../pool_3d_common/arch32/max_pool3d_grad_scatter_unified.h"
22 21 
23namespace AdaptiveMaxPool3DGrad {22namespace AdaptiveMaxPool3DGrad {
24using namespace AscendC;23using namespace AscendC;
25using namespace AdaptiveMaxPool3DGradComm;24using namespace AdaptiveMaxPool3DGradComm;
26 25 
27template <typename TX, typename TGrad, typename TArgmax, typename TY>26template <typename TX, typename TGrad, typename TArgmax, typename TY>
28-class AdaptiveMaxPool3DGradScatter : public AdaptiveMaxPool3DGradScatterBase<TX, TGrad, TArgmax, TY>27+class AdaptiveMaxPool3DGradScatter :
28+ public MaxPool3DGradCommon::MaxPool3DGradScatterUnified<
29+ TX, TGrad, TArgmax, TY,
30+ AdaptiveMaxPool3DGradTilingData,
31+ TilingParams,
32+ BlockParams,
33+ MaxPool3DGradScatterInternal::MaxPool3DGradScatterBaseTemplate>
29{34{
30public:35public:
31 __aicore__ inline AdaptiveMaxPool3DGradScatter(TPipe* pipe)36 __aicore__ inline AdaptiveMaxPool3DGradScatter(TPipe* pipe)
32- : AdaptiveMaxPool3DGradScatterBase<TX, TGrad, TArgmax, TY>(pipe)37+ : MaxPool3DGradCommon::MaxPool3DGradScatterUnified<
38+ TX, TGrad, TArgmax, TY,
39+ AdaptiveMaxPool3DGradTilingData,
40+ TilingParams,
41+ BlockParams,
42+ MaxPool3DGradScatterInternal::MaxPool3DGradScatterBaseTemplate>(pipe)
33 {}43 {}
34- __aicore__ inline void Init(
35- GM_ADDR x, GM_ADDR grad, GM_ADDR argmax, GM_ADDR y, GM_ADDR usrWorkspace,
36- const AdaptiveMaxPool3DGradTilingData* __restrict__ tiling)
37- {
38- // load tiling data
39- this->InitParams(tiling);
40- // set global buffer
41- this->InitInputsOutputs(x, grad, argmax, y, usrWorkspace);
42- // init global memory
43- InitYGMGlobalMemory(x, grad, argmax, y, usrWorkspace);
44- // ub buffer init
45- this->InitUbBuffer();
46- }
47- 
48- __aicore__ inline void InitYGMGlobalMemory(GM_ADDR x, GM_ADDR grad, GM_ADDR argmax, GM_ADDR y, GM_ADDR usrWorkspace)
49- {
50- InitGlobalMemory(this->yGm, this->params_.initLen, static_cast<TY>(0));
51- event_t eventMTE3ToS = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE3_S));
52- SetFlag<HardEvent::MTE3_S>(eventMTE3ToS);
53- WaitFlag<HardEvent::MTE3_S>(eventMTE3ToS);
54- }
55- 
56- __aicore__ inline void CalcOutOffset()
57- {
58- LocalTensor<TArgmax> argmaxUb = this->argmaxQue.template DeQue<TArgmax>(); // need free in the end
59- LocalTensor<TGrad> gradUb = this->gradQue.template DeQue<TGrad>(); // need free in the end
60- 
61- event_t eventMTE2ToS = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE2_S));
62- SetFlag<HardEvent::MTE2_S>(eventMTE2ToS);
63- WaitFlag<HardEvent::MTE2_S>(eventMTE2ToS);
64- 
65- uint64_t basedhwLen = this->block_.doShape * this->block_.woShape * this->block_.hoShape;
66- for (uint64_t ncIdx = 0; ncIdx < this->block_.ncShape; ncIdx++) {
67- uint64_t ncOffset = this->block_.baseNcOffset * this->params_.diHiWiLen;
68- for (uint64_t dhwIdx = 0; dhwIdx < basedhwLen; dhwIdx++) {
69- uint64_t ubOffset = ncIdx * basedhwLen + dhwIdx;
70- uint64_t outOffset = ncOffset + (uint64_t)argmaxUb.GetValue(ubOffset);
71- TGrad gradValue = gradUb.GetValue(ubOffset);
72- this->yGm.SetValue(outOffset, gradValue);
73- DataCacheCleanAndInvalid<TGrad, CacheLine::SINGLE_CACHE_LINE>(this->yGm[outOffset]);
74- }
75- this->block_.ShapeSum += basedhwLen;
76- if (this->block_.ShapeSum == this->params_.doDim * this->params_.hoDim * this->params_.woDim) {
77- this->block_.baseNcOffset += 1;
78- this->block_.ShapeSum = 0;
79- }
80- }
81- this->gradQue.FreeTensor(gradUb);
82- this->argmaxQue.FreeTensor(argmaxUb);
83- }
84- 
85- __aicore__ inline void CalcBlock()
86- {
87- this->CopyInGrad();
88- this->CopyInArgmax();
89- CalcOutOffset();
90- PipeBarrier<PIPE_ALL>();
91- }
92- 
93- __aicore__ inline void Process()
94- {
95- uint64_t ncIndex = this->params_.ncIndex;
96- for (uint64_t i = 0; i < this->params_.ncRealRound; i++) {
97- if (ncIndex < this->params_.ncCnt) {
98- this->block_.ncCntIndex = ncIndex;
99- this->block_.ncShape =
100- this->block_.ncCntIndex >= (this->params_.ncCnt - 1) ? this->params_.ncTail : this->params_.baseNc;
101- for (uint64_t j = 0; j < this->params_.dCnt; j++) {
102- this->block_.doCntIndex = j;
103- this->block_.doShape = this->block_.doCntIndex >= (this->params_.dCnt - 1) ? this->params_.doTail :
104- this->params_.baseDo;
105- for (uint64_t k = 0; k < this->params_.hCnt; k++) {
106- this->block_.hoCntIndex = k;
107- this->block_.hoShape = this->block_.hoCntIndex >= (this->params_.hCnt - 1) ?
108- this->params_.hoTail :
109- this->params_.baseHo;
110- for (uint64_t l = 0; l < this->params_.wCnt; l++) {
111- this->block_.woCntIndex = l;
112- this->block_.woShape = this->block_.woCntIndex >= (this->params_.wCnt - 1) ?
113- this->params_.woTail :
114- this->params_.baseWo;
115- this->block_.offsetGrad =
116- this->block_.ncCntIndex * this->params_.baseNc * this->params_.doDim *
117- this->params_.hoDim * this->params_.woDim +
118- this->block_.doCntIndex * this->params_.baseDo * this->params_.hoDim *
119- this->params_.woDim +
120- this->block_.hoCntIndex * this->params_.baseHo * this->params_.woDim +
121- this->block_.woCntIndex * this->params_.baseWo;
122- this->block_.offsetArgmax = this->block_.offsetGrad;
123- CalcBlock();
124- }
125- }
126- }
127- ncIndex += 1; // 当前ncCntIndex
128- }
129- }
130- }
131};44};
132} // namespace AdaptiveMaxPool3DGrad45} // namespace AdaptiveMaxPool3DGrad
133#endif // ADAPTIVE_MAX_POOL3D_GRAD_SCATTER_H46#endif // ADAPTIVE_MAX_POOL3D_GRAD_SCATTER_H
Mpooling/adaptive_max_pool3d_grad/op_kernel/adaptive_max_pool3d_grad_scatter_base.h+8-210
@@ -16,222 +16,20 @@
16#ifndef ADAPTIVE_MAX_POOL3D_GRAD_SCATTER_BASE_H16#ifndef ADAPTIVE_MAX_POOL3D_GRAD_SCATTER_BASE_H
17#define ADAPTIVE_MAX_POOL3D_GRAD_SCATTER_BASE_H17#define ADAPTIVE_MAX_POOL3D_GRAD_SCATTER_BASE_H
18 18 
19-#include "kernel_tiling/kernel_tiling.h"
20-#include "kernel_operator.h"
21#include "adaptive_max_pool3d_grad_common.h"19#include "adaptive_max_pool3d_grad_common.h"
20+#include "../pool_3d_common/arch32/max_pool3d_grad_scatter_base_template.h"
22 21 
23namespace AdaptiveMaxPool3DGrad {22namespace AdaptiveMaxPool3DGrad {
24using namespace AscendC;23using namespace AscendC;
25using namespace AdaptiveMaxPool3DGradComm;24using namespace AdaptiveMaxPool3DGradComm;
26 25 
27-const uint64_t REPEAT_BASE_BLOCK_LEN = 32;
28- 
29template <typename TX, typename TGrad, typename TArgmax, typename TY>26template <typename TX, typename TGrad, typename TArgmax, typename TY>
30-class AdaptiveMaxPool3DGradScatterBase27+using AdaptiveMaxPool3DGradScatterBase =
31-{28+ MaxPool3DGradScatterInternal::MaxPool3DGradScatterBaseTemplate<
32-public:29+ TX, TGrad, TArgmax, TY,
33- __aicore__ inline AdaptiveMaxPool3DGradScatterBase(TPipe* pipe)30+ AdaptiveMaxPool3DGradTilingData,
34- {31+ TilingParams,
35- pipe_ = pipe;32+ BlockParams>;
36- }
37 33 
38- __aicore__ inline void InitParams(const AdaptiveMaxPool3DGradTilingData* __restrict__ tiling)
39- {
40- params_.ncDim = tiling->ncDim; // 池化正向输出结果NC维度
41- params_.doDim = tiling->doDim; // 池化正向输出结果d维度
42- params_.hoDim = tiling->hoDim; // 池化正向输出结果h维度
43- params_.woDim = tiling->woDim; // 池化正向输出结果w维度
44- params_.diDim = tiling->diDim; // 池化正向输入d维度
45- params_.hiDim = tiling->hiDim; // 池化正向输入h维度
46- params_.wiDim = tiling->wiDim; // 池化正向输入w维度
47- 
48- params_.baseNc = tiling->baseNc; // 每次CalcOutOffset矩阵的NC维度
49- params_.baseDo = tiling->baseDo; // 每次CalcOutOffset矩阵的d维度
50- params_.baseHo = tiling->baseHo; // 每次CalcOutOffset矩阵的h维度
51- params_.baseWo = tiling->baseWo; // 每次CalcOutOffset矩阵的w维度
52- params_.ncTail = tiling->ncTail;
53- params_.doTail = tiling->doTail;
54- params_.hoTail = tiling->hoTail;
55- params_.woTail = tiling->woTail;
56- params_.ncCnt = tiling->ncCnt; // nc方向base矩阵个数
57- params_.dCnt = tiling->doCnt; // h方向base矩阵个数
58- params_.hCnt = tiling->hoCnt; // h方向base矩阵个数
59- params_.wCnt = tiling->woCnt; // w方向base矩阵个数
60- params_.usedCoreNum = tiling->usedCoreNum;
61- 
62- params_.totalCnt = tiling->totalCnt; // 需要处理base矩阵个数
63- params_.ncCntRound = tiling->ncRound; // 多核切nc，先分nc,向上取整,
64- params_.preCoreNum = tiling->preCoreNum; // 每个核均分完后剩余nce由前preCoreNum个核进行填充
65- params_.diHiWiLen = params_.diDim * params_.hiDim * params_.wiDim;
66- params_.ncRealRound = 0;
67- params_.ubSize = tiling->totalUBSize;
68- uint64_t blockId = GetBlockIdx();
69- if (params_.preCoreNum == 0 || blockId < params_.preCoreNum) { // 前preCoreNum个核
70- params_.ncIndex =
71- blockId * params_.ncCntRound; // 由于轮数为向上取整，所以当前核填充数的起始位置为 填充数*核数
72- params_.ncRealRound = params_.ncCntRound;
73- } else {
74- params_.ncIndex =
75- params_.preCoreNum * (params_.ncCntRound) + (blockId - params_.preCoreNum) * tiling->ncRoundTail;
76- params_.ncRealRound = tiling->ncRoundTail;
77- }
78- params_.realRound = params_.ncRealRound * params_.dCnt * params_.hCnt *
79- params_.wCnt; // 处理的base矩阵个数，包括base矩阵和tail矩阵
80- }
81- 
82- __aicore__ inline void InitInputsOutputs(GM_ADDR x, GM_ADDR grad, GM_ADDR argmax, GM_ADDR y, GM_ADDR usrWorkspace)
83- {
84- gradGm.SetGlobalBuffer((__gm__ TGrad*)grad);
85- argmaxGm.SetGlobalBuffer((__gm__ TArgmax*)argmax);
86- 
87- // 起始地址
88- uint64_t initOffset = params_.ncIndex * params_.baseNc * params_.diHiWiLen;
89- uint64_t initLen = 0;
90- uint64_t ncIndex = params_.ncIndex;
91- for (uint64_t j = 0; j < params_.ncRealRound; j++) {
92- block_.ncShape = ncIndex >= (params_.ncCnt - 1) ? params_.ncTail : params_.baseNc;
93- initLen += block_.ncShape * params_.diHiWiLen;
94- ncIndex += 1; // 当前ncCntIndex
95- }
96- params_.initLen = initLen;
97- params_.initOffset = initOffset;
98- 
99- yGm.SetGlobalBuffer((__gm__ TY*)y + initOffset, initLen);
100- workspaceGm.SetGlobalBuffer((__gm__ float*)usrWorkspace + initOffset, initLen);
101- }
102- 
103- __aicore__ inline void InitUbBuffer()
104- {
105- pipe_->InitBuffer(
106- gradQue, 1, params_.baseNc * params_.baseDo * params_.baseHo * params_.baseWo * sizeof(TGrad));
107- pipe_->InitBuffer(
108- argmaxQue, 1, params_.baseNc * params_.baseDo * params_.baseHo * params_.baseWo * sizeof(TArgmax));
109- }
110- 
111- __aicore__ inline void CopyInGrad()
112- {
113- LocalTensor<TGrad> gradUb = gradQue.AllocTensor<TGrad>();
114- if (params_.baseWo == params_.woDim) {
115- // repeat = VL, blockLen = baseHo * baseWo
116- uint32_t baseblockLen = block_.doShape * block_.hoShape * block_.woShape * sizeof(TGrad);
117- if (baseblockLen > REPEAT_BASE_BLOCK_LEN && baseblockLen % REPEAT_BASE_BLOCK_LEN == 0) {
118- DataCopyExtParams copyParamsGrad;
119- copyParamsGrad.blockCount = block_.ncShape;
120- copyParamsGrad.blockLen = baseblockLen;
121- copyParamsGrad.srcStride = 0;
122- copyParamsGrad.dstStride = 0;
123- DataCopyPadExtParams<TGrad> padGrad{false, 0, 0, 0};
124- 
125- DataCopyPad(gradUb, gradGm[block_.offsetGrad], copyParamsGrad, padGrad);
126- } else {
127- DataCopyExtParams copyParamsGrad;
128- uint32_t totalBlockLen = block_.ncShape * baseblockLen;
129- uint16_t baseBlockCount = totalBlockLen / REPEAT_BASE_BLOCK_LEN;
130- copyParamsGrad.blockCount = baseBlockCount;
131- copyParamsGrad.blockLen = REPEAT_BASE_BLOCK_LEN;
132- copyParamsGrad.srcStride = 0;
133- copyParamsGrad.dstStride = 0;
134- DataCopyPadExtParams<TGrad> padGrad{false, 0, 0, 0};
135- DataCopyPad(gradUb, gradGm[block_.offsetGrad], copyParamsGrad, padGrad);
136- 
137- uint32_t tailBlockLen = totalBlockLen % REPEAT_BASE_BLOCK_LEN;
138- if (tailBlockLen != 0) {
139- copyParamsGrad.blockCount = 1;
140- copyParamsGrad.blockLen = tailBlockLen;
141- padGrad.isPad = true;
142- DataCopyPad(
143- gradUb[baseBlockCount * (REPEAT_BASE_BLOCK_LEN / sizeof(TGrad))],
144- gradGm[block_.offsetGrad + baseBlockCount * (REPEAT_BASE_BLOCK_LEN / sizeof(TGrad))],
145- copyParamsGrad, padGrad);
146- }
147- }
148- } else {
149- // VL is for loop, repeat = baseHo, blockLen = baseWo
150- for (uint64_t loopidx = 0; loopidx < block_.ncShape; loopidx++) {
151- DataCopyExtParams copyParamsGrad;
152- copyParamsGrad.blockCount = params_.baseHo;
153- copyParamsGrad.blockLen = params_.baseWo * sizeof(TGrad);
154- copyParamsGrad.srcStride = (params_.woDim - block_.woShape) * sizeof(TGrad);
155- copyParamsGrad.dstStride = 0;
156- DataCopyPadExtParams<TGrad> padGrad{false, 0, 0, 0};
157- DataCopyPad(
158- gradUb[loopidx * block_.hoShape * block_.woShape],
159- gradGm[block_.offsetGrad + loopidx * params_.hoDim * params_.woDim], copyParamsGrad, padGrad);
160- }
161- }
162- gradQue.EnQue(gradUb);
163- }
164- 
165- __aicore__ inline void CopyInArgmax()
166- {
167- LocalTensor<TArgmax> argmaxUb = argmaxQue.AllocTensor<TArgmax>();
168- if (params_.baseWo == params_.woDim) {
169- uint32_t baseblockLen = block_.doShape * block_.hoShape * block_.woShape * sizeof(TArgmax);
170- if (baseblockLen > REPEAT_BASE_BLOCK_LEN && baseblockLen % REPEAT_BASE_BLOCK_LEN == 0) {
171- // repeat = VL, blockLen = baseHo * baseWo
172- DataCopyExtParams copyParamsArgmax;
173- copyParamsArgmax.blockCount = block_.ncShape;
174- copyParamsArgmax.blockLen = baseblockLen;
175- copyParamsArgmax.srcStride = 0;
176- copyParamsArgmax.dstStride = 0;
177- DataCopyPadExtParams<TArgmax> padArgmax{false, 0, 0, 0};
178- DataCopyPad(argmaxUb, argmaxGm[block_.offsetArgmax], copyParamsArgmax, padArgmax);
179- } else {
180- DataCopyExtParams copyParamsArgmax;
181- uint32_t totalBlockLen = block_.ncShape * baseblockLen;
182- uint16_t baseBlockCount = totalBlockLen / REPEAT_BASE_BLOCK_LEN;
183- copyParamsArgmax.blockCount = baseBlockCount;
184- copyParamsArgmax.blockLen = REPEAT_BASE_BLOCK_LEN;
185- copyParamsArgmax.srcStride = 0;
186- copyParamsArgmax.dstStride = 0;
187- DataCopyPadExtParams<TArgmax> padArgmax{false, 0, 0, 0};
188- DataCopyPad(argmaxUb, argmaxGm[block_.offsetArgmax], copyParamsArgmax, padArgmax);
189- 
190- uint16_t tailBlockLen = totalBlockLen % REPEAT_BASE_BLOCK_LEN;
191- if (tailBlockLen != 0) {
192- copyParamsArgmax.blockCount = 1;
193- copyParamsArgmax.blockLen = tailBlockLen;
194- copyParamsArgmax.srcStride = 0;
195- copyParamsArgmax.dstStride = 0;
196- padArgmax.isPad = true;
197- DataCopyPad(
198- argmaxUb[baseBlockCount * (REPEAT_BASE_BLOCK_LEN / sizeof(TArgmax))],
199- argmaxGm[block_.offsetArgmax + baseBlockCount * (REPEAT_BASE_BLOCK_LEN / sizeof(TArgmax))],
200- copyParamsArgmax, padArgmax);
201- }
202- }
203- } else {
204- // VL is for loop, repeat = baseHo, blockLen = baseWo
205- for (uint64_t loopidx = 0; loopidx < block_.ncShape; loopidx++) {
206- DataCopyExtParams copyParamsArgmax;
207- copyParamsArgmax.blockCount = params_.baseHo;
208- copyParamsArgmax.blockLen = params_.baseWo * sizeof(TArgmax);
209- copyParamsArgmax.srcStride = (params_.woDim - block_.woShape) * sizeof(TArgmax);
210- copyParamsArgmax.dstStride = 0;
211- DataCopyPadExtParams<TArgmax> padGrad{false, 0, 0, 0};
212- DataCopyPad(
213- argmaxUb[loopidx * block_.hoShape * block_.woShape],
214- argmaxGm[block_.offsetArgmax + loopidx * params_.hoDim * params_.woDim], copyParamsArgmax, padGrad);
215- }
216- }
217- argmaxQue.EnQue(argmaxUb);
218- }
219- 
220-public:
221- TilingParams params_;
222- BlockParams block_;
223- TPipe* pipe_ = nullptr;
224- 
225- GlobalTensor<TGrad> gradGm;
226- GlobalTensor<TArgmax> argmaxGm;
227- 
228- GlobalTensor<TY> yGm;
229- GlobalTensor<float> workspaceGm;
230- 
231- TQue<QuePosition::VECIN, 1> gradQue;
232- TQue<QuePosition::VECIN, 1> argmaxQue;
233- TQue<QuePosition::VECIN, 1> wsQue;
234- TQue<QuePosition::VECOUT, 1> yQue;
235-};
236} // namespace AdaptiveMaxPool3DGrad34} // namespace AdaptiveMaxPool3DGrad
237-#endif // ADAPTIVE_MAX_POOL3D_GRAD_SCATTER_BASE_H35+#endif // ADAPTIVE_MAX_POOL3D_GRAD_SCATTER_BASE_H
Mpooling/adaptive_max_pool3d_grad/op_kernel/adaptive_max_pool3d_grad_scatter_overlap.h+15-187
@@ -14,204 +14,32 @@
14 14 
15#ifndef ADAPTIVE_MAX_POOL3D_GRAD_SCATTER_OVERLAP_H15#ifndef ADAPTIVE_MAX_POOL3D_GRAD_SCATTER_OVERLAP_H
16#define ADAPTIVE_MAX_POOL3D_GRAD_SCATTER_OVERLAP_H16#define ADAPTIVE_MAX_POOL3D_GRAD_SCATTER_OVERLAP_H
17-#include "kernel_tiling/kernel_tiling.h"
18-#include "kernel_operator.h"
19-#include "adaptive_max_pool3d_grad_common.h"
20#include "adaptive_max_pool3d_grad_scatter_base.h"17#include "adaptive_max_pool3d_grad_scatter_base.h"
18+#include "adaptive_max_pool3d_grad_common.h"
19+#include "../pool_3d_common/arch32/max_pool3d_grad_scatter_overlap_unified.h"
CANN-robot
CANN-robotCANN-robot1月19日

代码结构与可维护性: 头文件包含路径使用相对路径 '../pool_3d_common/arch32/max_pool3d_grad_scatter_overlap_unified.h'，这种依赖关系不够清晰，且当文件目录结构发生变化时容易导致编译错误。相对路径增加了模块间的耦合度。

问题类型: 代码结构与可维护性 文件路径: pooling/adaptive_max_pool3d_grad/op_kernel/adaptive_max_pool3d_grad_scatter_overlap.h 行号: 19 问题代码:

#include "../pool_3d_common/arch32/max_pool3d_grad_scatter_overlap_unified.h"

修改建议:

1. 改为使用绝对路径或项目统一的包含路径策略。
2. 如果这是项目内的标准包含方式，建议添加注释说明这种包含关系的必要性。
3. 考虑将公共头文件放在统一的include目录下，使用 <> 包含。

此评论由代码审查工具自动生成

likedislike
21 20 
22namespace AdaptiveMaxPool3DGrad {21namespace AdaptiveMaxPool3DGrad {
23using namespace AscendC;22using namespace AscendC;
24using namespace AdaptiveMaxPool3DGradComm;23using namespace AdaptiveMaxPool3DGradComm;
25 24 
26template <typename TX, typename TGrad, typename TArgmax, typename TY>25template <typename TX, typename TGrad, typename TArgmax, typename TY>
27-class AdaptiveMaxPool3DGradScatterOverlap : public AdaptiveMaxPool3DGradScatterBase<TX, TGrad, TArgmax, TY>26+class AdaptiveMaxPool3DGradScatterOverlap :
27+ public MaxPool3DGradCommon::MaxPool3DGradScatterOverlapUnified<
28+ TX, TGrad, TArgmax, TY,
29+ AdaptiveMaxPool3DGradTilingData,
30+ TilingParams,
31+ BlockParams,
32+ MaxPool3DGradScatterInternal::MaxPool3DGradScatterBaseTemplate>
28{33{
29public:34public:
30 __aicore__ inline AdaptiveMaxPool3DGradScatterOverlap(TPipe* pipe)35 __aicore__ inline AdaptiveMaxPool3DGradScatterOverlap(TPipe* pipe)
31- : AdaptiveMaxPool3DGradScatterBase<TX, TGrad, TArgmax, TY>(pipe)36+ : MaxPool3DGradCommon::MaxPool3DGradScatterOverlapUnified<
37+ TX, TGrad, TArgmax, TY,
38+ AdaptiveMaxPool3DGradTilingData,
39+ TilingParams,
40+ BlockParams,
41+ MaxPool3DGradScatterInternal::MaxPool3DGradScatterBaseTemplate>(pipe)
32 {}42 {}
33- __aicore__ inline void Init(
34- GM_ADDR x, GM_ADDR grad, GM_ADDR argmax, GM_ADDR y, GM_ADDR usrWorkspace,
35- const AdaptiveMaxPool3DGradTilingData* __restrict__ tiling)
36- {
37- // load tiling data
38- this->InitParams(tiling);
39- // set global buffer
40- this->InitInputsOutputs(x, grad, argmax, y, usrWorkspace);
41- // init global memory
42- InitOutGlobalMemory(x, grad, argmax, y, usrWorkspace);
43- // ub buffer init
44- this->InitUbBuffer();
45- }
46- 
47- __aicore__ inline void InitOutGlobalMemory(GM_ADDR x, GM_ADDR grad, GM_ADDR argmax, GM_ADDR y, GM_ADDR usrWorkspace)
48- {
49- if constexpr (!is_same<TY, float>::value) {
50- InitGlobalMemory(this->workspaceGm, this->params_.initLen, 0.0f);
51- } else {
52- InitGlobalMemory(this->yGm, this->params_.initLen, static_cast<TY>(0));
53- }
54- event_t eventMTE3ToS = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE3_S));
55- SetFlag<HardEvent::MTE3_S>(eventMTE3ToS);
56- WaitFlag<HardEvent::MTE3_S>(eventMTE3ToS);
57- }
58- 
59- __aicore__ inline void CalcOutOffset()
60- {
61- LocalTensor<TArgmax> argmaxUb = this->argmaxQue.template DeQue<TArgmax>(); // need free in the end
62- LocalTensor<TGrad> gradUb = this->gradQue.template DeQue<TGrad>(); // need free in the end
63- 
64- event_t eventMTE2ToS = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE2_S));
65- SetFlag<HardEvent::MTE2_S>(eventMTE2ToS);
66- WaitFlag<HardEvent::MTE2_S>(eventMTE2ToS);
67- 
68- uint64_t basedhwLen = this->block_.doShape * this->block_.hoShape * this->block_.woShape;
69- for (uint64_t ncIdx = 0; ncIdx < this->block_.ncShape; ncIdx++) {
70- uint64_t ncOffset = this->block_.baseNcOffset * this->params_.diHiWiLen;
71- for (uint64_t dhwIdx = 0; dhwIdx < basedhwLen; dhwIdx++) {
72- uint64_t ubOffset = ncIdx * basedhwLen + dhwIdx;
73- uint64_t gmOffset = ncOffset + (uint64_t)argmaxUb.GetValue(ubOffset);
74- float gradValueFloat;
75- DataCacheCleanAndInvalid<float, CacheLine::ENTIRE_DATA_CACHE>(this->workspaceGm);
76- if constexpr (is_same<TY, float>::value) {
77- gradValueFloat = (this->yGm.GetValue(gmOffset) + gradUb.GetValue(ubOffset));
78- this->yGm.SetValue(gmOffset, gradValueFloat);
79- } else {
80- if constexpr (IsSameType<TGrad, bfloat16_t>::value) {
81- float ubValueFloat32 = ToFloat(gradUb.GetValue(ubOffset));
82- float gmValueFloat32 = this->workspaceGm.GetValue(gmOffset);
83- gradValueFloat = gmValueFloat32 + ubValueFloat32;
84- } else {
85- gradValueFloat = (this->workspaceGm.GetValue(gmOffset) + (float)gradUb.GetValue(ubOffset));
86- }
87- this->workspaceGm.SetValue(gmOffset, gradValueFloat);
88- }
89- DataCacheCleanAndInvalid<float, CacheLine::ENTIRE_DATA_CACHE>(this->workspaceGm);
90- }
91- this->block_.ShapeSum += basedhwLen;
92- if (this->block_.ShapeSum == this->params_.doDim * this->params_.hoDim * this->params_.woDim) {
93- this->block_.baseNcOffset += 1;
94- this->block_.ShapeSum = 0;
95- }
96- }
97- 
98- this->gradQue.FreeTensor(gradUb);
99- this->argmaxQue.FreeTensor(argmaxUb);
100- }
101- 
102- __aicore__ inline void CalcBlock()
103- {
104- this->CopyInGrad();
105- this->CopyInArgmax();
106- CalcOutOffset();
107- PipeBarrier<PIPE_ALL>();
108- }
109- 
110- __aicore__ inline void InitCastUbBuffer()
111- {
112- this->pipe_->Reset();
113- uint64_t maxCalcNum = this->params_.ubSize / (sizeof(half) + sizeof(float));
114- this->pipe_->InitBuffer(this->wsQue, 1, maxCalcNum * sizeof(float));
115- this->pipe_->InitBuffer(this->yQue, 1, maxCalcNum * sizeof(half));
116- }
117- 
118- __aicore__ inline void ProcessCast()
119- {
120- uint64_t maxCalcNum = this->params_.ubSize / (sizeof(half) + sizeof(float));
121- uint64_t totalLoops = CeilDiv(this->params_.initLen, maxCalcNum);
122- uint64_t calcTail = this->params_.initLen - (totalLoops - 1) * maxCalcNum;
123- for (uint64_t loopIndex = 0; loopIndex < totalLoops; loopIndex++) {
124- uint64_t calcNum = (loopIndex == totalLoops - 1) ? calcTail : maxCalcNum;
125- CopyInWorkspace(loopIndex * maxCalcNum, calcNum);
126- ComputeCast(calcNum);
127- CopyOutCast(loopIndex * maxCalcNum, calcNum);
128- }
129- }
130- 
131- __aicore__ inline void CopyInWorkspace(uint64_t gmOffset, uint64_t calcNum)
132- {
133- LocalTensor<float> fp32Ub = this->wsQue.template AllocTensor<float>();
134- 
135- DataCopyExtParams copyParamsWs;
136- copyParamsWs.blockCount = 1;
137- copyParamsWs.blockLen = calcNum * sizeof(float);
138- copyParamsWs.srcStride = 0;
139- copyParamsWs.dstStride = 0;
140- DataCopyPadExtParams<float> padWs{false, 0, 0, 0};
141- 
142- DataCopyPad(fp32Ub, this->workspaceGm[gmOffset], copyParamsWs, padWs);
143- this->wsQue.EnQue(fp32Ub);
144- }
145- 
146- __aicore__ inline void ComputeCast(uint64_t calcNum)
147- {
148- LocalTensor<float> fp32Ub = this->wsQue.template DeQue<float>();
149- LocalTensor<TY> b16Ub = this->yQue.template AllocTensor<TY>();
150- if constexpr (is_same<TY, half>::value) {
151- Cast(b16Ub, fp32Ub, RoundMode::CAST_NONE, calcNum); // 也可以只cast valid
152- } else if constexpr (is_same<TY, bfloat16_t>::value) {
153- Cast(b16Ub, fp32Ub, RoundMode::CAST_RINT, calcNum);
154- }
155- this->wsQue.template FreeTensor(fp32Ub);
156- this->yQue.template EnQue(b16Ub);
157- }
158- 
159- __aicore__ inline void CopyOutCast(uint64_t gmOffset, uint64_t calcNum)
160- {
161- LocalTensor<TY> yUb = this->yQue.template DeQue<TY>();
162- DataCopyExtParams copyParamsY;
163- copyParamsY.blockCount = 1;
164- copyParamsY.blockLen = calcNum * sizeof(TY);
165- copyParamsY.srcStride = 0;
166- copyParamsY.dstStride = 0;
167- DataCopyPad(this->yGm[gmOffset], yUb, copyParamsY);
168- this->yQue.template FreeTensor(yUb);
169- }
170- 
171- __aicore__ inline void Process()
172- {
173- uint64_t ncIndex = this->params_.ncIndex;
174- for (uint64_t i = 0; i < this->params_.ncRealRound; i++) {
175- if (ncIndex < this->params_.ncCnt) {
176- this->block_.ncCntIndex = ncIndex;
177- this->block_.ncShape =
178- this->block_.ncCntIndex >= (this->params_.ncCnt - 1) ? this->params_.ncTail : this->params_.baseNc;
179- for (uint64_t j = 0; j < this->params_.dCnt; j++) {
180- this->block_.doCntIndex = j;
181- this->block_.doShape = this->block_.doCntIndex >= (this->params_.dCnt - 1) ? this->params_.doTail :
182- this->params_.baseDo;
183- for (uint64_t k = 0; k < this->params_.hCnt; k++) {
184- this->block_.hoCntIndex = k;
185- this->block_.hoShape = this->block_.hoCntIndex >= (this->params_.hCnt - 1) ?
186- this->params_.hoTail :
187- this->params_.baseHo;
188- for (uint64_t w = 0; w < this->params_.wCnt; w++) {
189- this->block_.woCntIndex = w;
190- this->block_.woShape = this->block_.woCntIndex >= (this->params_.wCnt - 1) ?
191- this->params_.woTail :
192- this->params_.baseWo;
193- this->block_.offsetGrad =
194- this->block_.ncCntIndex * this->params_.baseNc * this->params_.doDim *
195- this->params_.hoDim * this->params_.woDim +
196- this->block_.doCntIndex * this->params_.baseDo * this->params_.hoDim *
197- this->params_.woDim +
198- this->block_.hoCntIndex * this->params_.baseHo * this->params_.woDim +
199- this->block_.woCntIndex * this->params_.baseWo;
200- this->block_.offsetArgmax = this->block_.offsetGrad;
201- CalcBlock();
202- }
203- }
204- }
205- ncIndex += 1; // 当前ncCntIndex
206- }
207- }
208- if constexpr (!is_same<TY, float>::value) {
209- PipeBarrier<PIPE_ALL>();
210- DataCacheCleanAndInvalid<float, CacheLine::ENTIRE_DATA_CACHE>(this->workspaceGm);
211- InitCastUbBuffer();
212- ProcessCast();
213- }
214- }
215};43};
216} // namespace AdaptiveMaxPool3DGrad44} // namespace AdaptiveMaxPool3DGrad
217#endif // ADAPTIVE_MAX_POOL3D_GRAD_SCATTER_OVERLAP_H45#endif // ADAPTIVE_MAX_POOL3D_GRAD_SCATTER_OVERLAP_H
Rpooling/adaptive_max_pool3d_grad/op_host/CMakeLists.txt→pooling/adaptive_max_pool3d_grad/tests/ut/op_api/CMakeLists.txt+8-6
@@ -1,12 +1,14 @@
1# ----------------------------------------------------------------------------1# ----------------------------------------------------------------------------
2+# This program is free software, you can redistribute it and/or modify.
2# Copyright (c) 2025 Huawei Technologies Co., Ltd.3# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of 4+# This file is a part of the CANN Open Software.
4-# CANN Open Software License Agreement Version 2.0 (the "License").5+# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
5# Please refer to the License for details. You may not use this file except in compliance with the License.6# Please refer to the License for details. You may not use this file except in compliance with the License.
6-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, 7+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
7-# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8# See LICENSE in the root of the software repository for the full text of the License.8# See LICENSE in the root of the software repository for the full text of the License.
9# ----------------------------------------------------------------------------9# ----------------------------------------------------------------------------
10-message(STATUS "=== Debug: start ops.pooling.adaptive_max_pool3d_grad.op_host.CMakeLists.txt ")
11 10 
12-add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE adaptive_max_pool3d_grad ACLNNTYPE aclnn_exclude)11+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+if(UT_TEST_ALL OR OP_API_UT)
13+ add_modules_ut_sources(HOSTNAME ${OP_API_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
14+endif()
Rpooling/adaptive_max_pool3d_grad/tests/ut/op_host/test_aclnn_AdaptiveMaxPool2dBackward.cpp→pooling/adaptive_max_pool3d_grad/tests/ut/op_api/test_aclnn_adaptive_max_pool2d_backward.cpp+5-5
@@ -1,10 +1,10 @@
1/**1/**
2+ * This program is free software, you can redistribute it and/or modify.
2 * Copyright (c) 2025 Huawei Technologies Co., Ltd.3 * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3- * This program is free software, you can redistribute it and/or modify it under the terms and conditions of4+ * This file is a part of the CANN Open Software.
4- * CANN Open Software License Agreement Version 2.0 (the "License").5+ * Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.6 * Please refer to the License for details. You may not use this file except in compliance with the License.
6- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,7+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
7- * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8 * See LICENSE in the root of the software repository for the full text of the License.8 * See LICENSE in the root of the software repository for the full text of the License.
9 */9 */
10#include <vector>10#include <vector>
@@ -12,7 +12,7 @@
12#include "gtest/gtest.h"12#include "gtest/gtest.h"
13 13 
14#include "opdev/op_log.h"14#include "opdev/op_log.h"
15-#include "../../../op_host/op_api/aclnn_adaptive_max_pool2d_backward.h"15+#include "../../../op_api/aclnn_adaptive_max_pool2d_backward.h"
16 16 
17#include "op_api_ut_common/tensor_desc.h"17#include "op_api_ut_common/tensor_desc.h"
18#include "op_api_ut_common/scalar_desc.h"18#include "op_api_ut_common/scalar_desc.h"
Rpooling/adaptive_max_pool3d_grad/tests/ut/op_host/test_aclnn_AdaptiveMaxPool3dBackward.cpp→pooling/adaptive_max_pool3d_grad/tests/ut/op_api/test_aclnn_adaptive_max_pool3d_backward.cpp+5-5
@@ -1,10 +1,10 @@
1/**1/**
2+ * This program is free software, you can redistribute it and/or modify.
2 * Copyright (c) 2025 Huawei Technologies Co., Ltd.3 * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3- * This program is free software, you can redistribute it and/or modify it under the terms and conditions of4+ * This file is a part of the CANN Open Software.
4- * CANN Open Software License Agreement Version 2.0 (the "License").5+ * Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.6 * Please refer to the License for details. You may not use this file except in compliance with the License.
6- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,7+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
7- * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8 * See LICENSE in the root of the software repository for the full text of the License.8 * See LICENSE in the root of the software repository for the full text of the License.
9 */9 */
10#include <vector>10#include <vector>
@@ -12,7 +12,7 @@
12#include "gtest/gtest.h"12#include "gtest/gtest.h"
13 13 
14#include "opdev/op_log.h"14#include "opdev/op_log.h"
15-#include "../../../op_host/op_api/aclnn_adaptive_max_pool3d_backward.h"15+#include "../../../op_api/aclnn_adaptive_max_pool3d_backward.h"
16 16 
17#include "op_api_ut_common/tensor_desc.h"17#include "op_api_ut_common/tensor_desc.h"
18#include "op_api_ut_common/scalar_desc.h"18#include "op_api_ut_common/scalar_desc.h"
Mpooling/adaptive_max_pool3d_grad/tests/ut/op_kernel/adaptive_max_pool3d_grad_tiling_def.h+3-3
@@ -67,8 +67,8 @@ inline void InitTilingData(uint8_t* tiling, AdaptiveMaxPool3DGradTilingData* con
67 AdaptiveMaxPool3DGradTilingData tiling_data; \67 AdaptiveMaxPool3DGradTilingData tiling_data; \
68 InitTilingData(tiling_arg, &tiling_data)68 InitTilingData(tiling_arg, &tiling_data)
69 69 
70-#define DTYPE_X half70+#define DTYPE_X float
71-#define DTYPE_GRAD half71+#define DTYPE_GRAD float
72#define DTYPE_ARGMAX int32_t72#define DTYPE_ARGMAX int32_t
73-#define DTYPE_Y half73+#define DTYPE_Y float
74#endif // _ADA_MAX_POOL3d_GRAD_TILING_DEF_H_74#endif // _ADA_MAX_POOL3d_GRAD_TILING_DEF_H_
Mpooling/adaptive_max_pool3d_grad/tests/ut/op_kernel/test_adaptive_max_pool3d_grad.cpp+33-6
@@ -41,7 +41,7 @@ struct AdaptiveMaxPool3DGradTestParam {
41 41 
42 int64_t blockDim = 0;42 int64_t blockDim = 0;
43 43 
44- size_t dataTypeSize;44+ size_t dataTypeSize = 4;
45 45 
46 uint32_t tilingKey;46 uint32_t tilingKey;
47 AdaptiveMaxPool3DGradTilingData tiling;47 AdaptiveMaxPool3DGradTilingData tiling;
@@ -84,14 +84,13 @@ TEST_P(AdaptiveMaxPool3DGradTest, test_case_adaptive_max_pool3d_grad)
84 int64_t gradByteSize = outputShapeSize * param.dataTypeSize;84 int64_t gradByteSize = outputShapeSize * param.dataTypeSize;
85 int64_t argmaxByteSize = outputShapeSize * sizeof(int32_t);85 int64_t argmaxByteSize = outputShapeSize * sizeof(int32_t);
86 int64_t dxByteSize = inputShapeSize * param.dataTypeSize;86 int64_t dxByteSize = inputShapeSize * param.dataTypeSize;
87- 
88 int64_t workspaceSize = 0;87 int64_t workspaceSize = 0;
89 int64_t tilingDataSize = sizeof(AdaptiveMaxPool3DGradTilingData);88 int64_t tilingDataSize = sizeof(AdaptiveMaxPool3DGradTilingData);
90 89 
91 uint8_t* x = (uint8_t*)AscendC::GmAlloc(xByteSize);90 uint8_t* x = (uint8_t*)AscendC::GmAlloc(xByteSize);
92 uint8_t* grad = (uint8_t*)AscendC::GmAlloc(gradByteSize);91 uint8_t* grad = (uint8_t*)AscendC::GmAlloc(gradByteSize);
93 uint8_t* argmax = (uint8_t*)AscendC::GmAlloc(argmaxByteSize);92 uint8_t* argmax = (uint8_t*)AscendC::GmAlloc(argmaxByteSize);
94- uint8_t* dy = (uint8_t*)AscendC::GmAlloc(dxByteSize);93+ uint8_t* dx = (uint8_t*)AscendC::GmAlloc(dxByteSize);
95 uint8_t* workspace = (uint8_t*)AscendC::GmAlloc(workspaceSize);94 uint8_t* workspace = (uint8_t*)AscendC::GmAlloc(workspaceSize);
96 uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(tilingDataSize);95 uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(tilingDataSize);
97 96 
@@ -137,7 +136,7 @@ TEST_P(AdaptiveMaxPool3DGradTest, test_case_adaptive_max_pool3d_grad)
137 136 
138 ICPU_SET_TILING_KEY(tilingKey);137 ICPU_SET_TILING_KEY(tilingKey);
139 AscendC::SetKernelMode(KernelMode::AIV_MODE);138 AscendC::SetKernelMode(KernelMode::AIV_MODE);
140- ICPU_RUN_KF(adaptive_max_pool3d_grad, blockDim, x, grad, argmax, dy, workspace, (uint8_t*)(tilingDatafromBin));139+ ICPU_RUN_KF(adaptive_max_pool3d_grad, blockDim, x, grad, argmax, dx, workspace, (uint8_t*)(tilingDatafromBin));
141 140 
142 AscendC::GmFree(x);141 AscendC::GmFree(x);
143 AscendC::GmFree(grad);142 AscendC::GmFree(grad);
@@ -156,11 +155,39 @@ static AdaptiveMaxPool3DGradTestParam cases[] = {
156 1,155 1,
157 1,156 1,
158 1,157 1,
159- sizeof(float),
160 40,158 40,
159+ sizeof(float),
161 2,160 2,
162 {3200, 1, 64, 64, 1, 1, 1, 1, 64, 64, 80, 1, 1, 1, 3200, 1, 1,161 {3200, 1, 64, 64, 1, 1, 1, 1, 64, 64, 80, 1, 1, 1, 3200, 1, 1,
163- 1, 80, 1, 1, 1, 40, 1, 1, 1, 40, 40, 196352, 1, 0, 1, 0}}162+ 1, 80, 1, 1, 1, 40, 1, 1, 1, 40, 40, 196352, 1, 0, 1, 0}},
163+ {"test_case_adaptive_max_pool3d_grad_scatter_overlap",
164+ 1,
165+ 1,
166+ 70,
167+ 70,
168+ 70,
169+ 3,
170+ 3,
171+ 3,
172+ 1,
173+ sizeof(float),
174+ 102,
175+ {1, 70, 70, 70, 3, 3, 3, 24, 24, 24, 1, 3, 3, 3, 1, 1, 1, 1, 1, 3, 3,
176+ 3, 1, 1, 1, 1, 1, 1, 196352, 1, 1, 1, 0}},
177+ {"test_case_adaptive_max_pool3d_grad_normal_overlap",
178+ 1,
179+ 1,
180+ 7,
181+ 7,
182+ 7,
183+ 3,
184+ 3,
185+ 3,
186+ 40,
187+ sizeof(float),
188+ 100,
189+ {1, 7, 7, 7, 3, 3, 3, 3, 3, 3, 1, 1, 1, 3, 64, 3, 3, 3, 1, 3, 3,
190+ 3, 1, 1, 1, 1, 1, 1, 196352, 0, 0, 0, 0}}
164};191};
165 192 
166INSTANTIATE_TEST_CASE_P(AdaptiveMaxPool3DGrad, AdaptiveMaxPool3DGradTest, testing::ValuesIn(cases));193INSTANTIATE_TEST_CASE_P(AdaptiveMaxPool3DGrad, AdaptiveMaxPool3DGradTest, testing::ValuesIn(cases));
Mpooling/max_pool3d_grad_with_argmax/CMakeLists.txt+1-9
@@ -8,12 +8,4 @@
8# See LICENSE in the root of the software repository for the full text of the License.8# See LICENSE in the root of the software repository for the full text of the License.
9# ----------------------------------------------------------------------------9# ----------------------------------------------------------------------------
10 10 
11-file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)11+add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE max_pool3d_grad_with_argmax ACLNNTYPE aclnn_exclude DEPENDENCIES max_pool_grad_with_argmax_v3 pool_3d_common)
12-if(NOT ENABLE_TEST AND NOT BENCHMARK)
13- list(REMOVE_ITEM CURRENT_DIRS tests)
14-endif()
15-foreach(SUB_DIR ${CURRENT_DIRS})
16- if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
17- add_subdirectory(${SUB_DIR})
18- endif()
19-endforeach()
Mpooling/max_pool3d_grad_with_argmax/README.md+5-0
@@ -4,8 +4,13 @@
4 4 
5| 产品 | 是否支持 |5| 产品 | 是否支持 |
6| ---- | :----:|6| ---- | :----:|
7+|昇腾910_95 AI处理器|×|
7|Atlas A3 训练系列产品/Atlas A3 推理系列产品|√|8|Atlas A3 训练系列产品/Atlas A3 推理系列产品|√|
8|Atlas A2 训练系列产品/Atlas 800I A2 推理产品/A200I A2 Box 异构组件|√|9|Atlas A2 训练系列产品/Atlas 800I A2 推理产品/A200I A2 Box 异构组件|√|
10+|Atlas 200I/500 A2推理产品|×|
11+|Atlas 推理系列产品|×|
12+|Atlas 训练系列产品|×|
13+|Atlas 200/300/500 推理产品|×|
9 14 
10## 功能说明15## 功能说明
11 16 
Mpooling/max_pool3d_grad_with_argmax/docs/aclnnMaxPool2dWithIndicesBackward.md+254-54
@@ -1,101 +1,301 @@
1# aclnnMaxPool2dWithIndicesBackward1# aclnnMaxPool2dWithIndicesBackward
2 2 
3+[📄 查看源码](https://gitcode.com/cann/ops-nn/tree/master/pooling/max_pool3d_grad_with_argmax)
4+ 
3## 产品支持情况5## 产品支持情况
4 6 
5| 产品 | 是否支持 |7| 产品 | 是否支持 |
6| :----------------------------------------------------------- | :------: |8| :----------------------------------------------------------- | :------: |
9+| <term>昇腾910_95 AI处理器</term> | √ |
7| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |10| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |
8-| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |11+| <term>Atlas A2 训练系列产品/Atlas 800I A2 推理产品/A200I A2 Box 异构组件</term> | √ |
12+| <term>Atlas 200I/500 A2 推理产品</term> | × |
13+| <term>Atlas 推理系列产品 </term> | × |
14+| <term>Atlas 训练系列产品</term> | × |
15+| <term>Atlas 200/300/500 推理产品</term> | × |
9 16 
10## 功能说明17## 功能说明
11 18 
12-- 算子功能：
13正向最大池化[aclnnMaxPool2dWithIndices](../../max_pool3d_with_argmax_v2/docs/aclnnMaxPool2dWithIndices.md)的反向传播。19正向最大池化[aclnnMaxPool2dWithIndices](../../max_pool3d_with_argmax_v2/docs/aclnnMaxPool2dWithIndices.md)的反向传播。
14- 输入tensor的推导公式：20- 输入tensor的推导公式：
15 - 当ceilMode=False时，indices tensor的shape中H和W维度推导公式：21 - 当ceilMode=False时，indices tensor的shape中H和W维度推导公式：
16- 
17 $$22 $$
18 [H_{out}, W_{out}]=[\lfloor{\frac{H_{in}+ padding\_size_{Htop} + padding\_size_{Hbottom} - {dilation\_size \times(k_h - 1) - 1}}{s_h}}\rfloor + 1,\lfloor{\frac{W_{in}+ padding\_size_{Wleft} + padding\_size_{Wright} - {dilation\_size \times(k_w - 1) - 1}}{s_w}}\rfloor + 1]23 [H_{out}, W_{out}]=[\lfloor{\frac{H_{in}+ padding\_size_{Htop} + padding\_size_{Hbottom} - {dilation\_size \times(k_h - 1) - 1}}{s_h}}\rfloor + 1,\lfloor{\frac{W_{in}+ padding\_size_{Wleft} + padding\_size_{Wright} - {dilation\_size \times(k_w - 1) - 1}}{s_w}}\rfloor + 1]
19 $$24 $$
20 25 
21 - 当ceilMode=True时，out tensor的shape中H和W维度推导公式：26 - 当ceilMode=True时，out tensor的shape中H和W维度推导公式：
22- 
23 $$27 $$
24 [H_{out}, W_{out}]=[\lceil{\frac{H_{in}+ padding\_size_{Htop} + padding\_size_{Hbottom} - {dilation\_size \times(k_h - 1) - 1}}{s_h}}\rceil + 1,\lceil{\frac{W_{in}+ padding\_size_{Wleft} + padding\_size_{Wright} - {dilation\_size \times(k_w - 1) - 1}}{s_w}}\rceil + 1]28 [H_{out}, W_{out}]=[\lceil{\frac{H_{in}+ padding\_size_{Htop} + padding\_size_{Hbottom} - {dilation\_size \times(k_h - 1) - 1}}{s_h}}\rceil + 1,\lceil{\frac{W_{in}+ padding\_size_{Wleft} + padding\_size_{Wright} - {dilation\_size \times(k_w - 1) - 1}}{s_w}}\rceil + 1]
25 $$29 $$
26 30 
27 - 滑窗左上角起始位处在下或右侧pad填充位上或者界外（无法取到有效值）时，舍弃该滑窗结果，在上述推导公式基础上对应空间轴shape需减去1：31 - 滑窗左上角起始位处在下或右侧pad填充位上或者界外（无法取到有效值）时，舍弃该滑窗结果，在上述推导公式基础上对应空间轴shape需减去1：
28- 
29 $$32 $$
30 \begin{cases}33 \begin{cases}
31 H_{out}=H_{out} - 1& \text{if } (H_{out}-1)*s_h>=H_{in}+padding\_size_{Htop} \\34 H_{out}=H_{out} - 1& \text{if } (H_{out}-1)*s_h>=H_{in}+padding\_size_{Htop} \\
32 W_{out}=W_{out} - 1& \text{if } (W_{out}-1)*s_w>=W_{in}+padding\_size_{Wleft} \\35 W_{out}=W_{out} - 1& \text{if } (W_{out}-1)*s_w>=W_{in}+padding\_size_{Wleft} \\
33 \end{cases}\\36 \end{cases}\\
34 $$37 $$
35- 
36## 函数原型38## 函数原型
37每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnMaxPool2dWithIndicesBackwardGetWorkspaceSize”接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用“aclnnMaxPool2dWithIndicesBackward”接口执行计算。39每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnMaxPool2dWithIndicesBackwardGetWorkspaceSize”接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用“aclnnMaxPool2dWithIndicesBackward”接口执行计算。
38 40 
39-- `aclnnStatus aclnnMaxPool2dWithIndicesBackwardGetWorkspaceSize(const aclTensor *gradOutput, const aclTensor *self, const aclTensor *indices, const aclIntArray *kernelSize, const aclIntArray *stride, const aclIntArray *padding, const aclIntArray *dilation, bool ceilMode, aclTensor *gradInput, uint64_t *workspaceSize, aclOpExecutor **executor)`41+```Cpp
40-- `aclnnStatus aclnnMaxPool2dWithIndicesBackward(void *workspace, uint64_t workspaceSize, aclOpExecutor *executor, aclrtStream stream)`42+aclnnStatus aclnnMaxPool2dWithIndicesBackwardGetWorkspaceSize(
43+ const aclTensor *gradOutput,
44+ const aclTensor *self,
45+ const aclTensor *indices,
46+ const aclIntArray *kernelSize,
47+ const aclIntArray *stride,
48+ const aclIntArray *padding,
49+ const aclIntArray *dilation,
50+ bool ceilMode,
51+ aclTensor *gradInput,
52+ uint64_t *workspaceSize,
53+ aclOpExecutor **executor)
54+```
55+```Cpp
56+aclnnStatus aclnnMaxPool2dWithIndicesBackward(
57+ void *workspace,
58+ uint64_t workspaceSize,
59+ aclOpExecutor *executor,
60+ aclrtStream stream)
61+```
41 62 
42## aclnnMaxPool2dWithIndicesBackwardGetWorkspaceSize63## aclnnMaxPool2dWithIndicesBackwardGetWorkspaceSize
43 64 
44- **参数说明：**65- **参数说明：**
45- * gradOutput(const aclTensor \*, 计算输入): 反向传播过程中上一步输出的梯度，Device侧aclTensor。和正向的输出shape一致。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)。数据格式与self保持一致。66+ <table style="undefined;table-layout: fixed; width: 1478px"><colgroup>
46- - <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term>、<term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：数据类型支持FLOAT。[数据格式](../../../docs/zh/context/数据格式.md)支持NCHW和CHW。67+ <col style="width: 149px">
47- * self(const aclTensor \*, 计算输入): 正向的输入数据，Device侧aclTensor。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)。68+ <col style="width: 121px">
48- - <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term>、<term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：数据类型支持FLOAT。[数据格式](../../../docs/zh/context/数据格式.md)支持NCHW和CHW。69+ <col style="width: 264px">
49- * indices(aclTensor \*, 计算输入): 正向输出的索引，是Device侧aclTensor，shape和输入gradOutput一致。正向输出中最大元素的索引位置，数据格式与self保持一致。70+ <col style="width: 253px">
50- - <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term>、<term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：数据类型仅支持INT32。[数据格式](../../../docs/zh/context/数据格式.md)支持NCHW和CHW。71+ <col style="width: 262px">
51- * kernelSize(const aclIntArray \*, 计算输入): 池化操作中使用的滑动窗口大小，Host侧的aclIntArray，长度仅支持1、2。72+ <col style="width: 148px">
52- - 当kernelSize中元素个数为1时，窗口大小为（kernelSize[0], kernelSize[0]）。73+ <col style="width: 135px">
53- - 当kernelSize中元素个数为2时，窗口大小为（kernelSize[0], kernelSize[1]）。74+ <col style="width: 146px">
54- * stride(aclIntArray*, 计算输入): 窗口移动的步长，Host侧的aclIntArray，长度仅支持0、1、2。stride的长度为0时，stride的数值等于kernelSize的值。75+ </colgroup>
55- - 当stride中元素个数为0时，步幅长度与kernelSize一致。76+ <thead>
56- - 当stride中元素个数为1时，步幅长度为（stride[0], stride[0]）。77+ <tr>
57- - 当stride中元素个数为2时，步幅长度为（stride[0], stride[1]）。78+ <th>参数名</th>
58- * padding(const aclIntArray \*, 计算输入): 输入数据的填充，表示输入每个维度上的填充量，影响池化窗口覆盖整个输入张量的行为，Host侧的aclIntArray，长度仅支持1、2。79+ <th>输入/输出</th>
59- - 当padding中元素个数为1时，对H与W轴的头部与尾部分别填充长度为padding[0]的`-Inf`。80+ <th>描述</th>
60- - 当padding中元素个数为2时，对H轴的头部与尾部填充长度为padding[0]的`-Inf`,对W轴的头部与尾部填充长度为padding[1]的`-Inf`。81+ <th>使用说明</th>
61- * dilation(const aclIntArray \*, 计算输入): 池化操作的扩张因子，扩张操作增加了池化窗口中元素间的距离，Host侧的aclIntArray。82+ <th>数据类型</th>
62- - <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term>、<term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：仅支持dilation为（1，1）。83+ <th>数据格式</th>
63- * ceilMode(const bool \*, 计算输入): 控制是否开启池化操作的输出大小为向上取整模式，Host侧的bool。为True时表示计算输出形状时，采用向上取整的方法；为False即向下取整。84+ <th>维度(shape)</th>
64- * gradInput(aclTensor \*, 计算输出): 反向传播输出的梯度，是Device侧aclTensor。shape与self保持一致。数据格式与self保持一致。85+ <th>非连续Tensor</th>
65- - <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term>、<term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：数据类型仅支持FLOAT。[数据格式](../../../docs/zh/context/数据格式.md)支持NCHW和CHW。86+ </tr></thead>
66- * workspaceSize(uint64_t \*, 出参): 返回需要在Device侧申请的workspace大小。87+ <tbody>
67- * executor(aclOpExecutor \*\*, 出参): 返回op执行器，包含了算子计算流程。88+ <tr>
89+ <td>gradOutput</td>
90+ <td>输入</td>
91+ <td>反向传播过程中上一步输出的梯度。</td>
92+ <td>和正向的输出shape一致，数据格式和self一致。</td>
93+ <td>FLOAT、FLOAT16、BFLOAT16</td>
94+ <td>ND、NCHW、NHWC。</td>
95+ <td>3-4</td>
96+ <td>√</td>
97+ </tr>
98+ <tr>
99+ <td>self</td>
100+ <td>输入</td>
101+ <td>正向的输入数据。</td>
102+ <td>-</td>
103+ <td>FLOAT、FLOAT16、BFLOAT16</td>
104+ <td>ND、NCHW、NHWC</td>
105+ <td>3-4</td>
106+ <td>√</td>
107+ </tr>
108+ <tr>
109+ <td>indices</td>
110+ <td>输入</td>
111+ <td>正向输出的索引。</td>
112+ <td>shape和输入gradOutput一致，正向输出中最大元素的索引位置，数据格式与self保持一致。</td>
113+ <td>INT32、INT64</td>
114+ <td>ND、NCHW、NHWC</td>
115+ <td>3-4</td>
116+ <td>√</td>
117+ </tr>
118+ <tr>
119+ <td>kernelSize</td>
120+ <td>输入</td>
121+ <td>池化操作中使用的滑动窗口大小。</td>
122+ <td>长度仅支持1、2。</td>
123+ <td>INT64</td>
124+ <td>-</td>
125+ <td>-</td>
126+ <td>-</td>
127+ </tr>
128+ <tr>
129+ <td>stride</td>
130+ <td>输入</td>
131+ <td>窗口移动的步长。</td>
132+ <td>长度仅支持0、1、2。stride的长度为0时，stride的数值等于kernelSize的值。</td>
133+ <td>INT64</td>
134+ <td>-</td>
135+ <td>-</td>
136+ <td>-</td>
137+ </tr>
138+ <tr>
139+ <td>padding</td>
140+ <td>输入</td>
141+ <td>输入数据的填充，表示输入每个维度上的填充量，影响池化窗口覆盖整个输入张量的行为。</td>
142+ <td>长度仅支持1、2。</td>
143+ <td>INT64</td>
144+ <td>-</td>
145+ <td>-</td>
146+ <td>-</td>
147+ </tr>
148+ <tr>
149+ <td>dilation</td>
150+ <td>输入</td>
151+ <td>池化操作的扩张因子，扩张操作增加了池化窗口中元素间的距离。</td>
152+ <td>-</td>
153+ <td>INT64</td>
154+ <td>-</td>
155+ <td>-</td>
156+ <td>-</td>
157+ </tr>
158+ <tr>
159+ <td>ceilMode</td>
160+ <td>输入</td>
161+ <td>计算输出形状时取整的方法。</td>
162+ <td>为True时表示计算输出形状时用向上取整的方法，为False时则表示向下取整。</td>
163+ <td>BOOL</td>
164+ <td>-</td>
165+ <td>-</td>
166+ <td>-</td>
167+ </tr>
168+ <tr>
169+ <td>gradInput</td>
170+ <td>输出</td>
171+ <td>反向传播输出的梯度。</td>
172+ <td>shape、数据格式与self保持一致。</td>
173+ <td>BFLOAT16、FLOAT16、FLOAT32</td>
174+ <td>ND、NCHW、NHWC。</td>
175+ <td>3-4</td>
176+ <td>√</td>
177+ </tr>
178+ <tr>
179+ <td>workspaceSize</td>
180+ <td>输出</td>
181+ <td>返回需要在Device侧申请的workspace大小。</td>
182+ <td>-</td>
183+ <td>-</td>
184+ <td>-</td>
185+ <td>-</td>
186+ <td>-</td>
187+ </tr>
188+ <tr>
189+ <td>executor</td>
190+ <td>输出</td>
191+ <td>返回op执行器，包含了算子计算流程。</td>
192+ <td>-</td>
193+ <td>-</td>
194+ <td>-</td>
195+ <td>-</td>
196+ <td>-</td>
197+ </tr>
198+ </tbody></table>
199+ <term>昇腾910_95 AI处理器</term>：gradOutput和self数据类型支持FLOAT、FLOAT16、BFLOAT16，数据格式支持ND、NCHW、NHWC。
200+ indices数据类型支持INT32、INT64。
201+ 支持dilation中的元素值大于0，支持1维或者2维输入。
202+ 
203+ <term>Atlas A2 训练系列产品/Atlas 800I A2 推理产品/A200I A2 Box 异构组件</term>、<term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：gradOutput和self数据类型支持FLOAT。数据格式支持NCHW和CHW。indices数据类型支持INT32。仅支持dilation为（1，1）。
68 204 
69- **返回值：**205- **返回值：**
70- 
71 aclnnStatus: 返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。206 aclnnStatus: 返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
72- 
73-```
74 第一段接口完成入参校验，出现以下场景时报错：207 第一段接口完成入参校验，出现以下场景时报错：
75- 161001(ACLNN_ERR_PARAM_NULLPTR)：1. 传入的self、indices是空指针。208+ <table style="undefined;table-layout: fixed; width: 1166px"><colgroup>
76- 161002(ACLNN_ERR_PARAM_INVALID)：1. gradOutput、self、indices、gradInput的数据类型不在支持的范围内。209+ <col style="width: 267px">
77- 2. gradOutput、self、indices、gradInput的数据格式不在支持的范围内。210+ <col style="width: 124px">
78- 3. gradOutput与indices的shape不一致，self和gradInput的shape不一致。211+ <col style="width: 775px">
79- 4. kernelSize的长度不等于1或者2。212+ </colgroup>
80- 5. kernelSize中的数值中存在小于等于0的数值。213+ <thead>
81- 6. stride的长度不等于0，1或2。214+ <tr>
82- 8. stride的数值中存在小于等于0的值。215+ <th>返回码</th>
83- 9. padding的元素个数不等于1或2.216+ <th>错误码</th>
84- 10. padding的数值中存在小于0或者大于kernelSize/2的值。217+ <th>描述</th>
85- 11. dilation的元素数值不符合入参要求。218+ </tr></thead>
86-```219+ <tbody>
87- 220+ <tr>
221+ <td>ACLNN_ERR_PARAM_NULLPTR</td>
222+ <td>161001</td>
223+ <td>传入的self、indices是空指针。</td>
224+ </tr>
225+ <tr>
226+ <td rowspan="11">ACLNN_ERR_PARAM_INVALID</td>
227+ <td rowspan="11">161002</td>
228+ <td>gradOutput、self、indices、gradInput的数据类型不在支持的范围内。</td>
229+ </tr>
230+ <tr>
231+ <td>gradOutput、self、indices、gradInput的数据格式不在支持的范围内。</td>
232+ </tr>
233+ <tr>
234+ <td>gradOutput与indices的shape不一致，self和gradInput的shape不一致。</td>
235+ </tr>
236+ <tr>
237+ <td>kernelSize的长度不等于1或者2。</td>
238+ </tr>
239+ <tr>
240+ <td>kernelSize中的数值中存在小于等于0的数值。</td>
241+ </tr>
242+ <tr>
243+ <td>stride的长度不等于0，1或2。</td>
244+ </tr>
245+ <tr>
246+ <td>stride的数值中存在小于等于0的值。</td>
247+ </tr>
248+ <tr>
249+ <td>padding的元素个数不等于1或2</td>
250+ </tr>
251+ <tr>
252+ <td>padding的数值中存在小于0或者大于kernelSize</td>
253+ </tr>
254+ <tr>
255+ <td>dilation的元素数值不符合入参要求。</td>
256+ </tr>
257+ </tbody>
258+ </table>
88## aclnnMaxPool2dWithIndicesBackward259## aclnnMaxPool2dWithIndicesBackward
89 260 
90- **参数说明：**261- **参数说明：**
91- * workspace(void \*, 入参): 在Device侧申请的workspace内存地址。262+ <table style="undefined;table-layout: fixed; width: 1166px"><colgroup>
92- * workspaceSize(uint64_t, 入参): 在Device侧申请的workspace大小，由第一段接口aclnnMaxPool2dWithIndicesBackwardGetWorkspaceSize获取。263+ <col style="width: 173px">
93- * executor(aclOpExecutor \*, 入参): op执行器，包含了算子计算流程。264+ <col style="width: 133px">
94- * stream(aclrtStream, 入参): 指定执行任务的Stream。265+ <col style="width: 860px">
266+ </colgroup>
267+ <thead>
268+ <tr>
269+ <th>参数名</th>
270+ <th>输入/输出</th>
271+ <th>描述</th>
272+ </tr></thead>
273+ <tbody>
274+ <tr>
275+ <td>workspace</td>
276+ <td>输入</td>
277+ <td>在Device侧申请的workspace内存地址。</td>
278+ </tr>
279+ <tr>
280+ <td>workspaceSize</td>
281+ <td>输入</td>
282+ <td>在Device侧申请的workspace大小，由第一段接口aclnnMaxPool2dWithIndicesBackwardGetWorkspaceSize获取。</td>
283+ </tr>
284+ <tr>
285+ <td>executor</td>
286+ <td>输入</td>
287+ <td>op执行器，包含了算子计算流程。</td>
288+ </tr>
289+ <tr>
290+ <td>stream</td>
291+ <td>输入</td>
292+ <td>指定执行任务的Stream。</td>
293+ </tr>
294+ </tbody>
295+ </table>
296+- **返回值：**
95 297 
96-- **返回值：**298+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
97- 
98- aclnnStatus: 返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
99 299 
100## 约束说明300## 约束说明
101- 确定性计算：301- 确定性计算：
Mpooling/max_pool3d_grad_with_argmax/docs/aclnnMaxPool2dWithMaskBackward.md+251-47
@@ -1,76 +1,278 @@
1# aclnnMaxPool2dWithMaskBackward1# aclnnMaxPool2dWithMaskBackward
2 2 
3+[📄 查看源码](https://gitcode.com/cann/ops-nn/tree/master/pooling/max_pool3d_grad_with_argmax)
4+ 
3## 产品支持情况5## 产品支持情况
4 6 
5| 产品 | 是否支持 |7| 产品 | 是否支持 |
6| :----------------------------------------------------------- | :------: |8| :----------------------------------------------------------- | :------: |
9+| <term>昇腾910_95 AI处理器</term> | × |
7| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |10| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |
8-| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |11+| <term>Atlas A2 训练系列产品/Atlas 800I A2 推理产品/A200I A2 Box 异构组件</term> | √ |
12+| <term>Atlas 200I/500 A2 推理产品</term> | × |
13+| <term>Atlas 推理系列产品 </term> | × |
14+| <term>Atlas 训练系列产品</term> | √ |
15+| <term>Atlas 200/300/500 推理产品</term> | × |
9 16 
10## 功能说明17## 功能说明
11- 
12-- 算子功能：
13正向最大池化[aclnnMaxPool2dWithMask](../../max_pool3d_with_argmax_v2/docs/aclnnMaxPool2dWithMask.md)的反向传播。18正向最大池化[aclnnMaxPool2dWithMask](../../max_pool3d_with_argmax_v2/docs/aclnnMaxPool2dWithMask.md)的反向传播。
14 19 
15## 函数原型20## 函数原型
16每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnMaxPool2dWithMaskBackwardGetWorkspaceSize”接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用“aclnnMaxPool2dWithMaskBackward”接口执行计算。21每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnMaxPool2dWithMaskBackwardGetWorkspaceSize”接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用“aclnnMaxPool2dWithMaskBackward”接口执行计算。
17 22 
18-- `aclnnStatus aclnnMaxPool2dWithMaskBackwardGetWorkspaceSize(const aclTensor *gradOutput, const aclTensor *self, const aclTensor *indices, const aclIntArray *kernelSize, const aclIntArray *stride, const aclIntArray *padding, const aclIntArray *dilation, bool ceilMode, aclTensor *gradInput, uint64_t *workspaceSize, aclOpExecutor **executor)`23+```Cpp
19-- `aclnnStatus aclnnMaxPool2dWithMaskBackward(void *workspace, uint64_t workspaceSize, aclOpExecutor *executor, aclrtStream stream)`24+aclnnStatus aclnnMaxPool2dWithMaskBackwardGetWorkspaceSize(
20- 25+ const aclTensor *gradOutput,
26+ const aclTensor *self,
27+ const aclTensor *indices,
28+ const aclIntArray *kernelSize,
29+ const aclIntArray *stride,
30+ const aclIntArray *padding,
31+ const aclIntArray *dilation,
32+ bool ceilMode,
33+ aclTensor *gradInput,
34+ uint64_t *workspaceSize,
35+ aclOpExecutor **executor)
36+```
37+```Cpp
38+aclnnStatus aclnnMaxPool2dWithMaskBackward(
39+ void *workspace,
40+ uint64_t workspaceSize,
41+ aclOpExecutor *executor,
42+ aclrtStream stream)
43+```
21## aclnnMaxPool2dWithMaskBackwardGetWorkspaceSize44## aclnnMaxPool2dWithMaskBackwardGetWorkspaceSize
22 45 
23- **参数说明：**46- **参数说明：**
24- * gradOutput(aclTensor*, 计算输入): 反向传播过程中上一步输出的梯度，Device侧aclTensor。和正向的输出shape一致。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持NCHW。47+ <table style="undefined;table-layout: fixed; width: 1478px"><colgroup>
25- * <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term>、<term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：数据类型支持FLOAT32、FLOAT16、BFLOAT16。48+ <col style="width: 149px">
26- * self(aclTensor*, 计算输入): 正向的输入数据，Device侧aclTensor。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持NCHW，与gradOutput一致。49+ <col style="width: 121px">
27- * <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term>、<term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：数据类型支持FLOAT32、FLOAT16、BFLOAT16。50+ <col style="width: 264px">
28- * indices(aclTensor \*, 计算输入): 正向输出的索引，是Device侧aclTensor。最大值在求mask的kernel位置的bit值组成的Tensor。数据类型仅支持INT8。[数据格式](../../../docs/zh/context/数据格式.md)支持NCHW，与self保持一致。51+ <col style="width: 253px">
29- * kernelSize(aclIntArray*, 计算输入): 池化操作中使用的滑动窗口大小，Host侧的aclIntArray，长度仅支持1、2。52+ <col style="width: 262px">
30- - 当kernelSize中元素个数为1时，窗口大小为（kernelSize[0], kernelSize[0]）。53+ <col style="width: 148px">
31- - 当kernelSize中元素个数为2时，窗口大小为（kernelSize[0], kernelSize[1]）。54+ <col style="width: 135px">
32- * stride(aclIntArray*, 计算输入): 窗口移动的步长，Host侧的aclIntArray，长度仅支持0、1、2。stride的长度为0时，stride的数值等于kernelSize的值。55+ <col style="width: 146px">
33- - 当stride中元素个数为0时，步幅长度与kernelSize一致。56+ </colgroup>
34- - 当stride中元素个数为1时，步幅长度为（stride[0], stride[0]）。57+ <thead>
35- - 当stride中元素个数为2时，步幅长度为（stride[0], stride[1]）。58+ <tr>
36- * padding(aclIntArray*, 计算输入): 每一条边补充的层数，补充的位置填写“负无穷”，Host侧的aclIntArray，长度仅支持1、2。59+ <th>参数名</th>
37- * dilation(aclIntArray*, 计算输入): 控制窗口中元素的步幅，Host侧的aclIntArray，长度仅支持1、2，值仅支持1。60+ <th>输入/输出</th>
38- * ceilMode(const bool, 计算输入): 控制是否开启池化操作的输出大小为向上取整模式，Host侧的bool。为True时表示计算输出形状时用向上取整的方法；为False时即向下取整。61+ <th>描述</th>
39- * gradInput(aclTensor \*, 计算输出): 反向传播输出的梯度，是Device侧aclTensor。shape与self保持一致。[数据格式](../../../docs/zh/context/数据格式.md)支持NCHW，与self保持一致。62+ <th>使用说明</th>
40- * <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term>、<term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：数据类型支持FLOAT32、FLOAT16、BFLOAT16。63+ <th>数据类型</th>
41- * workspaceSize(uint64_t \*, 出参): 返回需要在Device侧申请的workspace大小。64+ <th>数据格式</th>
42- * executor(aclOpExecutor \*\*, 出参): 返回op执行器，包含了算子计算流程。65+ <th>维度(shape)</th>
66+ <th>非连续Tensor</th>
67+ </tr></thead>
68+ <tbody>
69+ <tr>
70+ <td>gradOutput</td>
71+ <td>输入</td>
72+ <td>反向传播过程中上一步输出的梯度。</td>
73+ <td>和正向的输出shape一致，数据格式和self一致。</td>
74+ <td>FLOAT32、FLOAT16、BFLOAT16</td>
75+ <td>NCHW</td>
76+ <td>4</td>
77+ <td>√</td>
78+ </tr>
79+ <tr>
80+ <td>self</td>
81+ <td>输入</td>
82+ <td>正向的输入数据。</td>
83+ <td>-</td>
84+ <td>FLOAT、FLOAT16、BFLOAT16</td>
85+ <td>NCHW</td>
86+ <td>4</td>
87+ <td>√</td>
88+ </tr>
89+ <tr>
90+ <td>indices</td>
91+ <td>输入</td>
92+ <td>正向输出的索引。</td>
93+ <td>最大值在求mask的kernel位置的bit值组成的Tensor。</td>
94+ <td>INT8</td>
95+ <td>NCHW</td>
96+ <td>4</td>
97+ <td>√</td>
98+ </tr>
99+ <tr>
100+ <td>kernelSize</td>
101+ <td>输入</td>
102+ <td>池化操作中使用的滑动窗口大小。</td>
103+ <td>长度仅支持1、2。</td>
104+ <td>INT64</td>
105+ <td>-</td>
106+ <td>-</td>
107+ <td>-</td>
108+ </tr>
109+ <tr>
110+ <td>stride</td>
111+ <td>输入</td>
112+ <td>窗口移动的步长。</td>
113+ <td>长度仅支持0、1、2。stride的长度为0时，stride的数值等于kernelSize的值。</td>
114+ <td>INT64</td>
115+ <td>-</td>
116+ <td>-</td>
117+ <td>-</td>
118+ </tr>
119+ <tr>
120+ <td>padding</td>
121+ <td>输入</td>
122+ <td>输入数据的填充，表示输入每个维度上的填充量，影响池化窗口覆盖整个输入张量的行为。</td>
123+ <td>长度仅支持1、2。</td>
124+ <td>INT64</td>
125+ <td>-</td>
126+ <td>-</td>
127+ <td>-</td>
128+ </tr>
129+ <tr>
130+ <td>dilation</td>
131+ <td>输入</td>
132+ <td>控制窗口中元素的步幅。</td>
133+ <td>长度仅支持1、2，值仅支持1。</td>
134+ <td>INT64</td>
135+ <td>-</td>
136+ <td>-</td>
137+ <td>-</td>
138+ </tr>
139+ <tr>
140+ <td>ceilMode</td>
141+ <td>输入</td>
142+ <td>计算输出形状时取整的方法。</td>
143+ <td>为True时表示计算输出形状时用向上取整的方法，为False时则表示向下取整。</td>
144+ <td>BOOL</td>
145+ <td>-</td>
146+ <td>-</td>
147+ <td>-</td>
148+ </tr>
149+ <tr>
150+ <td>gradInput</td>
151+ <td>输出</td>
152+ <td>反向传播输出的梯度。</td>
153+ <td>shape和数据格式与self保持一致。</td>
154+ <td>BFLOAT16、FLOAT16、FLOAT32</td>
155+ <td>NCHW</td>
156+ <td>4</td>
157+ <td>√</td>
158+ </tr>
159+ <tr>
160+ <td>workspaceSize</td>
161+ <td>输出</td>
162+ <td>返回需要在Device侧申请的workspace大小。</td>
163+ <td>-</td>
164+ <td>-</td>
165+ <td>-</td>
166+ <td>-</td>
167+ <td>-</td>
168+ </tr>
169+ <tr>
170+ <td>executor</td>
171+ <td>输出</td>
172+ <td>返回op执行器，包含了算子计算流程。</td>
173+ <td>-</td>
174+ <td>-</td>
175+ <td>-</td>
176+ <td>-</td>
177+ <td>-</td>
178+ </tr>
179+ </tbody></table>
43 180 
44- **返回值：**181- **返回值：**
45 182 
46 aclnnStatus: 返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。183 aclnnStatus: 返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
47- 
48- ```
49 第一段接口完成入参校验，出现以下场景时报错：184 第一段接口完成入参校验，出现以下场景时报错：
50- 161001(ACLNN_ERR_PARAM_NULLPTR)：1. 传入的self、indices是空指针。185+ <table style="undefined;table-layout: fixed; width: 1166px"><colgroup>
51- 161002(ACLNN_ERR_PARAM_INVALID)：1. gradOutput、self、indices、gradInput的数据类型不在支持的范围内。186+ <col style="width: 267px">
52- 2. gradOutput、self、indices、gradInput的数据格式不在支持的范围内。187+ <col style="width: 124px">
53- 3. gradOutput与indices的shape不一致，self和gradInput的shape不一致。188+ <col style="width: 775px">
54- 4. kernelSize的长度不等于1或者2。189+ </colgroup>
55- 5. kernelSize中的数值中存在小于等于0的数值。190+ <thead>
56- 6. stride的长度不等于0，1或2。191+ <tr>
57- 8. stride的数值中存在小于等于0的值。192+ <th>返回码</th>
58- 9. padding的长度不等于1或2.193+ <th>错误码</th>
59- 10. padding的数值中存在小于0或者大于kernelSize/2的值。194+ <th>描述</th>
60- 11. dilation的数值不等于1。195+ </tr></thead>
61- ```196+ <tbody>
62- 197+ <tr>
198+ <td>ACLNN_ERR_PARAM_NULLPTR</td>
199+ <td>161001</td>
200+ <td>传入的self、indices是空指针。</td>
201+ </tr>
202+ <tr>
203+ <td rowspan="11">ACLNN_ERR_PARAM_INVALID</td>
204+ <td rowspan="11">161002</td>
205+ <td>gradOutput、self、indices、gradInput的数据类型不在支持的范围内。</td>
206+ </tr>
207+ <tr>
208+ <td>gradOutput、self、indices、gradInput的数据格式不在支持的范围内。</td>
209+ </tr>
210+ <tr>
211+ <td>gradOutput与indices的shape不一致，self和gradInput的shape不一致。</td>
212+ </tr>
213+ <tr>
214+ <td>kernelSize的长度不等于1或者2。</td>
215+ </tr>
216+ <tr>
217+ <td>kernelSize中的数值中存在小于等于0的数值。</td>
218+ </tr>
219+ <tr>
220+ <td>stride的长度不等于0，1或2。</td>
221+ </tr>
222+ <tr>
223+ <td>stride的数值中存在小于等于0的值。</td>
224+ </tr>
225+ <tr>
226+ <td>padding的长度不等于1或2。</td>
227+ </tr>
228+ <tr>
229+ <td>padding的数值中存在小于0或者大于kernelSize</td>
230+ </tr>
231+ <tr>
232+ <td>dilation的数值不等于1。</td>
233+ </tr>
234+ </tbody>
235+ </table>
63## aclnnMaxPool2dWithMaskBackward236## aclnnMaxPool2dWithMaskBackward
64 237 
65- **参数说明：**238- **参数说明：**
66- * workspace(void \*, 入参): 在Device侧申请的workspace内存地址。239+ <table style="undefined;table-layout: fixed; width: 1166px"><colgroup>
67- * workspaceSize(uint64_t, 入参): 在Device侧申请的workspace大小，由第一段接口aclnnMaxPool2dWithMaskBackwardGetWorkspaceSize获取。240+ <col style="width: 173px">
68- * executor(aclOpExecutor \*, 入参): op执行器，包含了算子计算流程。241+ <col style="width: 133px">
69- * stream(aclrtStream, 入参): 指定执行任务的Stream。242+ <col style="width: 860px">
243+ </colgroup>
244+ <thead>
245+ <tr>
246+ <th>参数名</th>
247+ <th>输入/输出</th>
248+ <th>描述</th>
249+ </tr></thead>
250+ <tbody>
251+ <tr>
252+ <td>workspace</td>
253+ <td>输入</td>
254+ <td>在Device侧申请的workspace内存地址。</td>
255+ </tr>
256+ <tr>
257+ <td>workspaceSize</td>
258+ <td>输入</td>
259+ <td>在Device侧申请的workspace大小，由第一段接口aclnnMaxPool2dWithMaskBackwardGetWorkspaceSize获取。</td>
260+ </tr>
261+ <tr>
262+ <td>executor</td>
263+ <td>输入</td>
264+ <td>op执行器，包含了算子计算流程。</td>
265+ </tr>
266+ <tr>
267+ <td>stream</td>
268+ <td>输入</td>
269+ <td>指定执行任务的Stream。</td>
270+ </tr>
271+ </tbody>
272+ </table>
273+- **返回值：**
70 274 
71-- **返回值：**275+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
72- 
73- aclnnStatus: 返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
74 276 
75## 约束说明277## 约束说明
76- 确定性计算：278- 确定性计算：
@@ -78,6 +280,8 @@
78 280 
79- 输入数据暂不支持NaN、-Inf。281- 输入数据暂不支持NaN、-Inf。
80 282 
283+- <term>Atlas 训练系列产品</term>：当输入数据是FLOAT类型时，会转换为FLOAT16类型进行计算，存在一定程度的精度损失。
284+ 
81## 调用示例285## 调用示例
82示例代码如下，仅供参考，具体编译和执行过程请参考[编译与运行样例](../../../docs/zh/context/编译与运行样例.md)。286示例代码如下，仅供参考，具体编译和执行过程请参考[编译与运行样例](../../../docs/zh/context/编译与运行样例.md)。
83```Cpp287```Cpp
Mpooling/max_pool3d_grad_with_argmax/docs/aclnnMaxPool3dWithArgmaxBackward.md+256-43
@@ -1,74 +1,287 @@
1# aclnnMaxPool3dWithArgmaxBackward1# aclnnMaxPool3dWithArgmaxBackward
2 2 
3+[📄 查看源码](https://gitcode.com/cann/ops-nn/tree/master/pooling/max_pool3d_grad_with_argmax)
4+ 
3## 产品支持情况5## 产品支持情况
4 6 
5| 产品 | 是否支持 |7| 产品 | 是否支持 |
6| :----------------------------------------------------------- | :------: |8| :----------------------------------------------------------- | :------: |
9+| <term>昇腾910_95 AI处理器</term> | × |
7| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |10| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |
8-| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |11+| <term>Atlas A2 训练系列产品/Atlas 800I A2 推理产品/A200I A2 Box 异构组件</term> | √ |
12+| <term>Atlas 200I/500 A2 推理产品</term> | × |
13+| <term>Atlas 推理系列产品 </term> | × |
14+| <term>Atlas 训练系列产品</term> | × |
15+| <term>Atlas 200/300/500 推理产品</term> | × |
9 16 
10## 功能说明17## 功能说明
11- 
12-- 算子功能：
13正向最大池化[aclnnMaxPool3dWithArgmax](../../max_pool3d_with_argmax_v2/docs/aclnnMaxPool3dWithArgmax.md)的反向传播，将梯度回填到每个窗口最大值的坐标处，相同坐标处累加。18正向最大池化[aclnnMaxPool3dWithArgmax](../../max_pool3d_with_argmax_v2/docs/aclnnMaxPool3dWithArgmax.md)的反向传播，将梯度回填到每个窗口最大值的坐标处，相同坐标处累加。
14 19 
15## 函数原型20## 函数原型
16每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnMaxPool3dWithArgmaxBackwardGetWorkspaceSize”接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用“aclnnMaxPool3dWithArgmaxBackward”接口执行计算。21每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnMaxPool3dWithArgmaxBackwardGetWorkspaceSize”接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用“aclnnMaxPool3dWithArgmaxBackward”接口执行计算。
17 22 
18-- `aclnnStatus aclnnMaxPool3dWithArgmaxBackwardGetWorkspaceSize(const aclTensor *gradOutput, const aclTensor *self, const aclTensor *indices, const aclIntArray *kernelSize, const aclIntArray *stride, const aclIntArray *padding, const aclIntArray *dilation, bool ceilMode, aclTensor *gradInput, uint64_t *workspaceSize, aclOpExecutor **executor)`
19-- `aclnnStatus aclnnMaxPool3dWithArgmaxBackward(void *workspace, uint64_t workspaceSize, aclOpExecutor *executor, aclrtStream stream)`
20 23 
24+```Cpp
25+aclnnStatus aclnnMaxPool3dWithArgmaxBackwardGetWorkspaceSize(
26+ const aclTensor *gradOutput,
27+ const aclTensor *self,
28+ const aclTensor *indices,
29+ const aclIntArray *kernelSize,
30+ const aclIntArray *stride,
31+ const aclIntArray *padding,
32+ const aclIntArray *dilation,
33+ bool ceilMode,
34+ aclTensor *gradInput,
35+ uint64_t *workspaceSize,
36+ aclOpExecutor **executor)
37+```
38+```Cpp
39+aclnnStatus aclnnMaxPool3dWithArgmaxBackward(
40+ void *workspace,
41+ uint64_t workspaceSize,
42+ aclOpExecutor *executor,
43+ aclrtStream stream)
44+```
21## aclnnMaxPool3dWithArgmaxBackwardGetWorkspaceSize45## aclnnMaxPool3dWithArgmaxBackwardGetWorkspaceSize
22 46 
23- **参数说明：**47- **参数说明：**
24- * gradOutput(aclTensor*, 计算输入): 梯度Tensor，Device侧aclTensor。和正向的输出shape一致。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持ND, 当输入是5维时，内部按照NCDHW处理，当输入是4维时，在0维度处补1，内部按照NCDHW处理。48+ <table style="undefined;table-layout: fixed; width: 1478px"><colgroup>
25- * <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term>、<term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：数据类型支持FLOAT32、FLOAT16、BFLOAT16。49+ <col style="width: 149px">
26- * self(aclTensor*, 计算输入): 正向的输入Tensor，Device侧aclTensor。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持ND, 当输入是5维时，内部按照NCDHW处理，当输入是4维时，在0维度处补1，内部按照NCDHW处理，与gradOutput一致。50+ <col style="width: 121px">
27- * <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term>、<term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：数据类型支持FLOAT32、FLOAT16、BFLOAT16。51+ <col style="width: 264px">
28- * indices(aclTensor \*, 计算输入): 输入Tensor，是Device侧aclTensor。正向输入中最大元素的索引位置。[数据格式](../../../docs/zh/context/数据格式.md)支持NCDHW，与self保持一致。shape与gradOutput一致。52+ <col style="width: 253px">
29- * <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term>、<term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：数据类型仅支持INT3253+ <col style="width: 262px">
30- * kernelSize(aclIntArray*, 计算输入): 表示最大池化的窗口大小。Host侧的aclIntArray，表示池化窗口的大小，INT64类型数组，长度为1 ($kD = kH = kW$) 或3 ($kD, kH, kW$)。54+ <col style="width: 148px">
31- * stride(aclIntArray*, 计算输入): Host侧的aclIntArray，表示池化操作的步长，INT64类型的数组，长度为0（$sD = kD, sH = kH, sW = kW$）或者1（$sD = sH = sW$）或3（$sD, sH, sW$）。55+ <col style="width: 135px">
32- * padding(aclIntArray*, 计算输入): Host侧的aclIntArray，表示在输入的D、H、W方向上padding补0的层数，INT64类型数组，长度为1（$padD = padH = padW$）或3（$padD, padH, padW$）。56+ <col style="width: 146px">
33- * dilation(aclIntArray*, 计算输入): Host侧的aclIntArray，表示控制窗口中元素的步幅，INT64类型数组，长度为1（$dD = dH = dW$）或3（$dD, dH, dW$），值仅支持1。57+ </colgroup>
34- * ceilMode(bool, 计算输入): 表示正向平均池化过程中推导的输出的shape是否向上取整。数据类型支持BOOL。58+ <thead>
35- * gradInput(aclTensor \*, 计算输出): 反向输出Tensor，是Device侧aclTensor。shape与self保持一致。[数据格式](../../../docs/zh/context/数据格式.md)支持NCDHW，与self保持一致。59+ <tr>
36- * <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term>、<term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：数据类型支持FLOAT32、FLOAT16、BFLOAT16。60+ <th>参数名</th>
37- * workspaceSize(uint64_t \*, 出参): 返回需要在Device侧申请的workspace大小。61+ <th>输入/输出</th>
38- * executor(aclOpExecutor \*\*, 出参): 返回op执行器，包含了算子计算流程。62+ <th>描述</th>
39- 63+ <th>使用说明</th>
64+ <th>数据类型</th>
65+ <th>数据格式</th>
66+ <th>维度(shape)</th>
67+ <th>非连续Tensor</th>
68+ </tr></thead>
69+ <tbody>
70+ <tr>
71+ <td>gradOutput</td>
72+ <td>输入</td>
73+ <td>反向传播过程中上一步输出的梯度。</td>
74+ <td>和正向的输出shape一致，数据格式和self一致。</td>
75+ <td>FLOAT32、FLOAT16、BFLOAT16</td>
76+ <td>ND</td>
77+ <td>4-5</td>
78+ <td>√</td>
79+ </tr>
80+ <tr>
81+ <td>self</td>
82+ <td>输入</td>
83+ <td>正向的输入数据。</td>
84+ <td>-</td>
85+ <td>FLOAT、FLOAT16、BFLOAT16</td>
86+ <td>ND</td>
87+ <td>4-5</td>
88+ <td>√</td>
89+ </tr>
90+ <tr>
91+ <td>indices</td>
92+ <td>输入</td>
93+ <td>正向输入中最大元素的索引位置。</td>
94+ <td>-</td>
95+ <td>INT32</td>
96+ <td>ND</td>
97+ <td>4-5</td>
98+ <td>√</td>
99+ </tr>
100+ <tr>
101+ <td>kernelSize</td>
102+ <td>输入</td>
103+ <td>池化操作中使用的滑动窗口大小。</td>
104+ <td>长度为1 (kD = kH = kW) 或3 (kD, kH, kW)。</td>
105+ <td>INT64</td>
106+ <td>-</td>
107+ <td>-</td>
108+ <td>-</td>
109+ </tr>
110+ <tr>
111+ <td>stride</td>
112+ <td>输入</td>
113+ <td>窗口移动的步长。</td>
114+ <td>长度为0（sD = kD, sH = kH, sW = kW）或者1（sD = sH = sW）或3（sD, sH, sW）。</td>
115+ <td>INT64</td>
116+ <td>-</td>
117+ <td>-</td>
118+ <td>-</td>
119+ </tr>
120+ <tr>
121+ <td>padding</td>
122+ <td>输入</td>
123+ <td>在输入的D、H、W方向上padding补0的层数。</td>
124+ <td>长度为1（padD = padH = padW）或3（padD, padH, padW）。</td>
125+ <td>INT64</td>
126+ <td>-</td>
127+ <td>-</td>
128+ <td>-</td>
129+ </tr>
130+ <tr>
131+ <td>dilation</td>
132+ <td>输入</td>
133+ <td>控制窗口中元素的步幅。</td>
134+ <td>长度为1或3，值仅支持1。</td>
135+ <td>INT64</td>
136+ <td>-</td>
137+ <td>-</td>
138+ <td>-</td>
139+ </tr>
140+ <tr>
141+ <td>ceilMode</td>
142+ <td>输入</td>
143+ <td>计算输出形状时取整的方法。</td>
144+ <td>为True时表示计算输出形状时用向上取整的方法，为False时则表示向下取整。</td>
145+ <td>BOOL</td>
146+ <td>-</td>
147+ <td>-</td>
148+ <td>-</td>
149+ </tr>
150+ <tr>
151+ <td>gradInput</td>
152+ <td>输出</td>
153+ <td>反向传播输出的梯度。</td>
154+ <td>shape和数据格式与self保持一致。</td>
155+ <td>BFLOAT16、FLOAT16、FLOAT32</td>
156+ <td>ND</td>
157+ <td>4-5</td>
158+ <td>√</td>
159+ </tr>
160+ <tr>
161+ <td>workspaceSize</td>
162+ <td>输出</td>
163+ <td>返回需要在Device侧申请的workspace大小。</td>
164+ <td>-</td>
165+ <td>-</td>
166+ <td>-</td>
167+ <td>-</td>
168+ <td>-</td>
169+ </tr>
170+ <tr>
171+ <td>executor</td>
172+ <td>输出</td>
173+ <td>返回op执行器，包含了算子计算流程。</td>
174+ <td>-</td>
175+ <td>-</td>
176+ <td>-</td>
177+ <td>-</td>
178+ <td>-</td>
179+ </tr>
180+ </tbody></table>
181+
40- **返回值：**182- **返回值：**
41 183 
42 aclnnStatus: 返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。184 aclnnStatus: 返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
43 185 
44- ```
45 第一段接口完成入参校验，出现以下场景时报错：186 第一段接口完成入参校验，出现以下场景时报错：
46- 161001(ACLNN_ERR_PARAM_NULLPTR)：1. 传入的self、indices是空指针。187+ <table style="undefined;table-layout: fixed; width: 1166px"><colgroup>
47- 161002(ACLNN_ERR_PARAM_INVALID)：1. gradOutput、self、indices、gradInput的数据类型不在支持的范围内。188+ <col style="width: 267px">
48- 2. gradOutput、self、indices、gradInput的数据格式不在支持的范围内。189+ <col style="width: 124px">
49- 3. gradOutput与indices的shape不一致，self和gradInput的shape不一致。190+ <col style="width: 775px">
50- 4. kernelSize的长度不等于1或者3。191+ </colgroup>
51- 5. kernelSize中存在小于等于0的数值。192+ <thead>
52- 6. stride的长度不等于0，1或3。193+ <tr>
53- 7. stride中存在小于等于0的数值。194+ <th>返回码</th>
54- 8. padding的长度不等于1或3.195+ <th>错误码</th>
55- 9. padding中存在小于0或者大于kernelSize/2的数值。196+ <th>描述</th>
56- 10. dilation的长度不等于1或者3。197+ </tr></thead>
57- 11. 平台不支持198+ <tbody>
58- 12. depth * height * width > max int32，超出了indices的表示范围。199+ <tr>
59- ```200+ <td>ACLNN_ERR_PARAM_NULLPTR</td>
201+ <td>161001</td>
202+ <td>传入的self、indices是空指针。</td>
203+ </tr>
204+ <tr>
205+ <td rowspan="12">ACLNN_ERR_PARAM_INVALID</td>
206+ <td rowspan="12">161002</td>
207+ <td>gradOutput、self、indices、gradInput的数据类型不在支持的范围内。</td>
208+ </tr>
209+ <tr>
210+ <td>gradOutput、self、indices、gradInput的数据格式不在支持的范围内。</td>
211+ </tr>
212+ <tr>
213+ <td>gradOutput与indices的shape不一致，self和gradInput的shape不一致。</td>
214+ </tr>
215+ <tr>
216+ <td>kernelSize的长度不等于1或者3。</td>
217+ </tr>
218+ <tr>
219+ <td>kernelSize中存在小于等于0的数值。</td>
220+ </tr>
221+ <tr>
222+ <td>stride的长度不等于0，1或3。</td>
223+ </tr>
224+ <tr>
225+ <td>stride中存在小于等于0的数值。</td>
226+ </tr>
227+ <tr>
228+ <td>padding的长度不等于1或3。</td>
229+ </tr>
230+ <tr>
231+ <td>padding中存在小于0或者大于kernelSize/2的数值。</td>
232+ </tr>
233+ <tr>
234+ <td>dilation的长度不等于1或者3。</td>
235+ </tr>
236+ <tr>
237+ <td>平台不支持。</td>
238+ </tr>
239+ <tr>
240+ <td>depth * height * width > max int32，超出了indices的表示范围。</td>
241+ </tr>
242+ </tbody>
243+ </table>
60 244 
61## aclnnMaxPool3dWithArgmaxBackward245## aclnnMaxPool3dWithArgmaxBackward
62 246 
63- **参数说明：**247- **参数说明：**
64- * workspace(void \*, 入参): 在Device侧申请的workspace内存地址。248+ <table style="undefined;table-layout: fixed; width: 1166px"><colgroup>
65- * workspaceSize(uint64_t, 入参): 在Device侧申请的workspace大小，由第一段接口aclnnMaxPool3dWithArgmaxBackwardGetWorkspaceSize获取。249+ <col style="width: 173px">
66- * executor(aclOpExecutor \*, 入参): op执行器，包含了算子计算流程。250+ <col style="width: 133px">
67- * stream(aclrtStream, 入参): 指定执行任务的Stream。251+ <col style="width: 860px">
252+ </colgroup>
253+ <thead>
254+ <tr>
255+ <th>参数名</th>
256+ <th>输入/输出</th>
257+ <th>描述</th>
258+ </tr></thead>
259+ <tbody>
260+ <tr>
261+ <td>workspace</td>
262+ <td>输入</td>
263+ <td>在Device侧申请的workspace内存地址。</td>
264+ </tr>
265+ <tr>
266+ <td>workspaceSize</td>
267+ <td>输入</td>
268+ <td>在Device侧申请的workspace大小，由第一段接口aclnnMaxPool3dWithArgmaxBackwardGetWorkspaceSize获取。</td>
269+ </tr>
270+ <tr>
271+ <td>executor</td>
272+ <td>输入</td>
273+ <td>op执行器，包含了算子计算流程。</td>
274+ </tr>
275+ <tr>
276+ <td>stream</td>
277+ <td>输入</td>
278+ <td>指定执行任务的Stream。</td>
279+ </tr>
280+ </tbody>
281+ </table>
282+- **返回值：**
68 283 
69-- **返回值：**284+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
70- 
71- aclnnStatus: 返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
72 285 
73## 约束说明286## 约束说明
74- 确定性计算：287- 确定性计算：
Mpooling/max_pool3d_grad_with_argmax/examples/test_aclnn_max_pool2d_with_mask_backward_exec2d.cpp+3-3
@@ -1,12 +1,12 @@
1/**1/**
2 * Copyright (c) 2025 Huawei Technologies Co., Ltd.2 * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3- * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
6- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7 * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.7 * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8 * See LICENSE in the root of the software repository for the full text of the License.8 * See LICENSE in the root of the software repository for the full text of the License.
9- */9+*/
10#include <iostream>10#include <iostream>
11#include <vector>11#include <vector>
12#include "acl/acl.h"12#include "acl/acl.h"
Rpooling/max_pool3d_grad_with_argmax/op_host/op_api/aclnn_max_pool2d_with_indices_backward.cpp→pooling/max_pool3d_grad_with_argmax/op_api/aclnn_max_pool2d_with_indices_backward.cpp+4-4
@@ -1,16 +1,16 @@
1/**1/**
2 * Copyright (c) 2025 Huawei Technologies Co., Ltd.2 * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3- * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
6- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7 * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.7 * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8 * See LICENSE in the root of the software repository for the full text of the License.8 * See LICENSE in the root of the software repository for the full text of the License.
9- */9+*/
10 10 
11#include "aclnn_max_pool2d_with_indices_backward.h"11#include "aclnn_max_pool2d_with_indices_backward.h"
12#include "max_pool_grad_with_argmax_v1.h"12#include "max_pool_grad_with_argmax_v1.h"
13-#include "../../../max_pool_grad_with_argmax_v3/op_api/max_pool_grad_with_argmax_v3.h"13+#include "pooling/max_pool_grad_with_argmax_v3/op_api/max_pool_grad_with_argmax_v3.h"
14#include "max_pool3d_grad_with_argmax.h"14#include "max_pool3d_grad_with_argmax.h"
15#include "aclnn_kernels/contiguous.h"15#include "aclnn_kernels/contiguous.h"
16#include "level0/unsqueeze.h"16#include "level0/unsqueeze.h"
Rpooling/max_pool3d_grad_with_argmax/op_host/op_api/aclnn_max_pool2d_with_indices_backward.h→pooling/max_pool3d_grad_with_argmax/op_api/aclnn_max_pool2d_with_indices_backward.h+0-0
文件重命名但无更改。
Rpooling/max_pool3d_grad_with_argmax/op_host/op_api/aclnn_max_pool3d_with_argmax_backward.cpp→pooling/max_pool3d_grad_with_argmax/op_api/aclnn_max_pool3d_with_argmax_backward.cpp+0-0
文件重命名但无更改。
Rpooling/max_pool3d_grad_with_argmax/op_host/op_api/aclnn_max_pool3d_with_argmax_backward.h→pooling/max_pool3d_grad_with_argmax/op_api/aclnn_max_pool3d_with_argmax_backward.h+0-0
文件重命名但无更改。
Rpooling/max_pool3d_grad_with_argmax/op_host/op_api/max_pool3d_grad_with_argmax.cpp→pooling/max_pool3d_grad_with_argmax/op_api/max_pool3d_grad_with_argmax.cpp+0-0
文件重命名但无更改。
Rpooling/max_pool3d_grad_with_argmax/op_host/op_api/max_pool3d_grad_with_argmax.h→pooling/max_pool3d_grad_with_argmax/op_api/max_pool3d_grad_with_argmax.h+0-0
文件重命名但无更改。
Rpooling/max_pool3d_grad_with_argmax/op_host/op_api/max_pool_grad_with_argmax_v1.cpp→pooling/max_pool3d_grad_with_argmax/op_api/max_pool_grad_with_argmax_v1.cpp+3-3
@@ -1,12 +1,12 @@
1/**1/**
2 * Copyright (c) 2025 Huawei Technologies Co., Ltd.2 * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3- * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
6- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7 * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.7 * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8 * See LICENSE in the root of the software repository for the full text of the License.8 * See LICENSE in the root of the software repository for the full text of the License.
9- */9+*/
10 10 
11#include "max_pool_grad_with_argmax_v1.h"11#include "max_pool_grad_with_argmax_v1.h"
12#include "opdev/data_type_utils.h"12#include "opdev/data_type_utils.h"
Rpooling/max_pool3d_grad_with_argmax/op_host/op_api/max_pool_grad_with_argmax_v1.h→pooling/max_pool3d_grad_with_argmax/op_api/max_pool_grad_with_argmax_v1.h+3-3
@@ -1,12 +1,12 @@
1/**1/**
2 * Copyright (c) 2025 Huawei Technologies Co., Ltd.2 * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3- * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 * CANN Open Software License Agreement Version 2.0 (the "License").4 * CANN Open Software License Agreement Version 2.0 (the "License").
5 * Please refer to the License for details. You may not use this file except in compliance with the License.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
6- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7 * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.7 * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8 * See LICENSE in the root of the software repository for the full text of the License.8 * See LICENSE in the root of the software repository for the full text of the License.
9- */9+*/
10#ifndef OP_API_INC_LEVEL0_MAX_POOL_GRAD_WITH_ARGMAX_V1_H_10#ifndef OP_API_INC_LEVEL0_MAX_POOL_GRAD_WITH_ARGMAX_V1_H_
11#define OP_API_INC_LEVEL0_MAX_POOL_GRAD_WITH_ARGMAX_V1_H_11#define OP_API_INC_LEVEL0_MAX_POOL_GRAD_WITH_ARGMAX_V1_H_
12 12 
Mpooling/max_pool3d_grad_with_argmax/op_host/max_pool3d_grad_with_argmax_cutk_tiling.cpp+4-0
@@ -273,6 +273,7 @@ void MaxPool3DGradWithArgmaxCutKTiling::SetOtherTilingParams()
273 } else {273 } else {
274 maxPoolGradParams.workspaceSize = 0UL;274 maxPoolGradParams.workspaceSize = 0UL;
275 }275 }
276+ maxPoolGradParams.isNeedWorkspace = maxPoolGradParams.xDtypeSize != DTYPE_LEN_B32 && maxPoolGradParams.isOverLap;
276}277}
277 278 
278ge::graphStatus MaxPool3DGradWithArgmaxCutKTiling::DoOpTiling()279ge::graphStatus MaxPool3DGradWithArgmaxCutKTiling::DoOpTiling()
@@ -283,6 +284,9 @@ ge::graphStatus MaxPool3DGradWithArgmaxCutKTiling::DoOpTiling()
283 SetOtherTilingParams();284 SetOtherTilingParams();
284 SetBaseTilingData();285 SetBaseTilingData();
285 PrintTilingData();286 PrintTilingData();
287+ if (maxPoolGradParams.isInitOutput || maxPoolGradParams.isNeedWorkspace) {
288+ context_->SetScheduleMode(1);
289+ }
286 return ge::GRAPH_SUCCESS;290 return ge::GRAPH_SUCCESS;
287}291}
288 292 
Mpooling/max_pool3d_grad_with_argmax/op_host/max_pool3d_grad_with_argmax_normal_tiling.cpp+5-1
@@ -294,7 +294,8 @@ void MaxPool3DGradWithArgmaxNormalTiling::SetOtherTilingParams()
294 maxPoolGradParams.totalCnt =294 maxPoolGradParams.totalCnt =
295 maxPoolGradParams.ncCnt * maxPoolGradParams.doCnt * maxPoolGradParams.hoCnt * maxPoolGradParams.woCnt;295 maxPoolGradParams.ncCnt * maxPoolGradParams.doCnt * maxPoolGradParams.hoCnt * maxPoolGradParams.woCnt;
296 maxPoolGradParams.usedCoreNum = std::min(maxPoolGradParams.totalCnt, maxPoolGradParams.totalCoreNum);296 maxPoolGradParams.usedCoreNum = std::min(maxPoolGradParams.totalCnt, maxPoolGradParams.totalCoreNum);
297- if (maxPoolGradParams.xDtypeSize != DTYPE_LEN_B32 && maxPoolGradParams.isOverLap) {297+ maxPoolGradParams.isNeedWorkspace = maxPoolGradParams.xDtypeSize != DTYPE_LEN_B32 && maxPoolGradParams.isOverLap;
298+ if (maxPoolGradParams.isNeedWorkspace) {
298 maxPoolGradParams.workspaceSize = maxPoolGradParams.ncDim * maxPoolGradParams.diDim * maxPoolGradParams.hiDim *299 maxPoolGradParams.workspaceSize = maxPoolGradParams.ncDim * maxPoolGradParams.diDim * maxPoolGradParams.hiDim *
299 maxPoolGradParams.wiDim * sizeof(float);300 maxPoolGradParams.wiDim * sizeof(float);
300 } else {301 } else {
@@ -329,6 +330,9 @@ ge::graphStatus MaxPool3DGradWithArgmaxNormalTiling::DoOpTiling()
329 SetBaseTilingData();330 SetBaseTilingData();
330 SetNormalTilingData();331 SetNormalTilingData();
331 PrintTilingData();332 PrintTilingData();
333+ if (maxPoolGradParams.isInitOutput || maxPoolGradParams.isNeedWorkspace) {
334+ context_->SetScheduleMode(1);
335+ }
332 PrintNormalTilingData();336 PrintNormalTilingData();
333 return ge::GRAPH_SUCCESS;337 return ge::GRAPH_SUCCESS;
334}338}
Mpooling/max_pool3d_grad_with_argmax/op_host/max_pool3d_grad_with_argmax_scatter_tiling.cpp+24-90
@@ -14,6 +14,7 @@
14 */14 */
15#include <iostream>15#include <iostream>
16#include "max_pool3d_grad_with_argmax_tiling.h"16#include "max_pool3d_grad_with_argmax_tiling.h"
17+#include "../../pool_3d_common/op_host/arch32/max_pool3d_grad_tiling_common.h"
17 18 
18namespace optiling {19namespace optiling {
19ge::graphStatus MaxPool3DGradWithArgmaxScatterTiling::GetShapeAttrsInfo()20ge::graphStatus MaxPool3DGradWithArgmaxScatterTiling::GetShapeAttrsInfo()
@@ -32,121 +33,54 @@ bool MaxPool3DGradWithArgmaxScatterTiling::IsCapable()
32 33 
33bool MaxPool3DGradWithArgmaxScatterTiling::SetScatterTilingParams()34bool MaxPool3DGradWithArgmaxScatterTiling::SetScatterTilingParams()
34{35{
35- const uint64_t doDim = maxPoolGradParams.doDim;36+ return CalculateScatterTilingParams(
36- const uint64_t hoDim = maxPoolGradParams.hoDim;37+ maxPoolGradParams,
37- const uint64_t woDim = maxPoolGradParams.woDim;38+ maxPoolGradParams.doDim,
38- const uint64_t xDtypeSize = maxPoolGradParams.xDtypeSize;39+ maxPoolGradParams.hoDim,
39- const uint64_t indexDtypeSize = maxPoolGradParams.indexDtypeSize;40+ maxPoolGradParams.woDim,
40- 41+ maxPoolGradParams.xDtypeSize,
41- uint64_t ncPreCore = Ops::Base::CeilDiv(maxPoolGradParams.ncDim, maxPoolGradParams.totalCoreNum);42+ maxPoolGradParams.indexDtypeSize,
42- maxPoolGradParams.usedCoreNum = Ops::Base::CeilDiv(maxPoolGradParams.ncDim, ncPreCore);43+ MAX_BLOCK_COUNT,
43- 44+ BLOCK_SIZE);
44- // Scatter main tiling cal
45- // 1. All Tensor full size, cut nc between cores, without cut in one core
46- uint64_t noCutSize = ncPreCore * doDim * hoDim * woDim * (xDtypeSize + indexDtypeSize);
47- if (noCutSize <= (maxPoolGradParams.maxUbSize - BLOCK_SIZE)) {
48- maxPoolGradParams.baseNc = ncPreCore;
49- maxPoolGradParams.baseDo = doDim;
50- maxPoolGradParams.baseHo = hoDim;
51- maxPoolGradParams.baseWo = woDim;
52- maxPoolGradParams.ubCutAxis = TILING_UB_NO_CUT;
53- return true;
54- }
55- 
56- // 2. Cut nc
57- uint64_t perNcSize = 1UL * doDim * hoDim * woDim * (xDtypeSize + indexDtypeSize);
58- if (perNcSize <= (maxPoolGradParams.maxUbSize - BLOCK_SIZE)) {
59- uint64_t baseNc = (maxPoolGradParams.maxUbSize - BLOCK_SIZE) / perNcSize;
60- if (baseNc > MAX_BLOCK_COUNT) { // Use NC for blockCount.
61- baseNc = MAX_BLOCK_COUNT;
62- }
63- maxPoolGradParams.baseNc = baseNc;
64- maxPoolGradParams.baseDo = doDim;
65- maxPoolGradParams.baseHo = hoDim;
66- maxPoolGradParams.baseWo = woDim;
67- maxPoolGradParams.ubCutAxis = TILING_UB_CUT_NC;
68- return true;
69- }
70- maxPoolGradParams.baseNc = 1UL;
71- 
72- // 3. Cut do
73- uint64_t perDoSize = 1UL * 1UL * hoDim * woDim * (xDtypeSize + indexDtypeSize);
74- if (perDoSize <= (maxPoolGradParams.maxUbSize - BLOCK_SIZE)) {
75- maxPoolGradParams.baseDo = (maxPoolGradParams.maxUbSize - BLOCK_SIZE) / perDoSize;
76- maxPoolGradParams.baseHo = hoDim;
77- maxPoolGradParams.baseWo = woDim;
78- maxPoolGradParams.ubCutAxis = TILING_UB_CUT_DO;
79- return true;
80- }
81- maxPoolGradParams.baseDo = 1UL;
82- 
83- // 4. Cut ho
84- uint64_t perHoSize = 1UL * 1UL * 1UL * woDim * (xDtypeSize + indexDtypeSize);
85- if (perHoSize <= (maxPoolGradParams.maxUbSize - BLOCK_SIZE)) {
86- maxPoolGradParams.baseHo = (maxPoolGradParams.maxUbSize - BLOCK_SIZE) / perHoSize;
87- maxPoolGradParams.baseWo = woDim;
88- maxPoolGradParams.ubCutAxis = TILING_UB_CUT_HO;
89- return true;
90- }
91- maxPoolGradParams.baseHo = 1UL;
92- 
93- // 5. Cut wo
94- uint64_t perWoSize = 1UL * 1UL * 1UL * 1UL * (xDtypeSize + indexDtypeSize);
95- if (perWoSize <= (maxPoolGradParams.maxUbSize - BLOCK_SIZE)) {
96- maxPoolGradParams.baseWo = (maxPoolGradParams.maxUbSize - BLOCK_SIZE) / perWoSize;
97- maxPoolGradParams.ubCutAxis = TILING_UB_CUT_WO;
98- return true;
99- }
100- maxPoolGradParams.baseWo = 1UL;
101- 
102- return false;
103}45}
104 46 
105void MaxPool3DGradWithArgmaxScatterTiling::SetOtherTilingParams()47void MaxPool3DGradWithArgmaxScatterTiling::SetOtherTilingParams()
106{48{
107 SetCntTailTilingParams();49 SetCntTailTilingParams();
108- maxPoolGradParams.ncRound = Ops::Base::CeilDiv(maxPoolGradParams.ncCnt, maxPoolGradParams.usedCoreNum);50+
109- maxPoolGradParams.preCoreNum = maxPoolGradParams.ncCnt % maxPoolGradParams.usedCoreNum;51+ CalculateRoundParams(
110- maxPoolGradParams.ncRoundTail =52+ maxPoolGradParams,
111- maxPoolGradParams.preCoreNum == 0UL ? maxPoolGradParams.ncRound : maxPoolGradParams.ncRound - 1UL;53+ maxPoolGradParams.isOverLap,
112- maxPoolGradParams.totalRound =54+ maxPoolGradParams.diDim,
113- maxPoolGradParams.ncRound * maxPoolGradParams.doCnt * maxPoolGradParams.hoCnt * maxPoolGradParams.woCnt;55+ maxPoolGradParams.hiDim,
114- if (maxPoolGradParams.xDtypeSize != DTYPE_LEN_B32 && maxPoolGradParams.isOverLap) {56+ maxPoolGradParams.wiDim);
115- maxPoolGradParams.workspaceSize = maxPoolGradParams.ncDim * maxPoolGradParams.diDim * maxPoolGradParams.hiDim *
116- maxPoolGradParams.wiDim * sizeof(float);
117- } else {
118- maxPoolGradParams.workspaceSize = 0UL;
119- }
120}57}
121 58 
122void MaxPool3DGradWithArgmaxScatterTiling::SetScatterTilingData()59void MaxPool3DGradWithArgmaxScatterTiling::SetScatterTilingData()
123{60{
124- tilingData.set_ncRound(maxPoolGradParams.ncRound);61+ SetScatterTilingDataCommon(tilingData, maxPoolGradParams);
125- tilingData.set_ncRoundTail(maxPoolGradParams.ncRoundTail);
126- tilingData.set_totalRound(maxPoolGradParams.totalRound);
127- tilingData.set_preCoreNum(maxPoolGradParams.preCoreNum);
128}62}
129 63 
130void MaxPool3DGradWithArgmaxScatterTiling::PrintScatterTilingData()64void MaxPool3DGradWithArgmaxScatterTiling::PrintScatterTilingData()
131{65{
132- OP_LOGI(66+ PrintScatterTilingDataCommon(context_->GetNodeName(), tilingData);
133- context_->GetNodeName(), "TilingData ncRound: %lu, ncRoundTail: %lu, totalRound: %lu.",
134- tilingData.get_ncRound(), tilingData.get_ncRoundTail(), tilingData.get_totalRound());
135}67}
136 68 
137ge::graphStatus MaxPool3DGradWithArgmaxScatterTiling::DoOpTiling()69ge::graphStatus MaxPool3DGradWithArgmaxScatterTiling::DoOpTiling()
138{70{
139 bool res = SetScatterTilingParams();71 bool res = SetScatterTilingParams();
140- OP_CHECK_IF(72+ OP_CHECK_IF(!res, OP_LOGE(context_->GetNodeName(), "Scatter cal tiling params failed."), return ge::GRAPH_FAILED);
141- !res, OP_LOGE(context_->GetNodeName(), "Scatter cal tiling params failed."), return ge::GRAPH_FAILED);
142 maxPoolGradParams.tilingType = TILING_TYPE_SCATTER;73 maxPoolGradParams.tilingType = TILING_TYPE_SCATTER;
143 SetOtherTilingParams();74 SetOtherTilingParams();
144 SetBaseTilingData();75 SetBaseTilingData();
145 SetScatterTilingData();76 SetScatterTilingData();
146 PrintTilingData();77 PrintTilingData();
147 PrintScatterTilingData();78 PrintScatterTilingData();
79+
148 return ge::GRAPH_SUCCESS;80 return ge::GRAPH_SUCCESS;
149}81}
150 82 
151-REGISTER_TILING_TEMPLATE("MaxPool3DGradWithArgmax", MaxPool3DGradWithArgmaxScatterTiling, 6);83+REGISTER_TILING_TEMPLATE("MaxPool3DGradWithArgmax",
84+ MaxPool3DGradWithArgmaxScatterTiling, 6);
85+ 
152} // namespace optiling86} // namespace optiling
Mpooling/max_pool3d_grad_with_argmax/op_host/max_pool3d_grad_with_argmax_tiling.h+3-41
@@ -11,14 +11,6 @@
11/*!11/*!
12 * \file max_pool3d_grad_with_argmax_tiling.h12 * \file max_pool3d_grad_with_argmax_tiling.h
13 * \brief13 * \brief
14- * ATTENTION: MAKE SURE 'BEGIN_TILING_DATA_DEF' STAY IN THE SAME LINE (45) USING BLANK LINES.
15- *
16- *
17- *
18- *
19- *
20- *
21- *
22 */14 */
23#ifndef OPS_BUILD_IN_OP_TILING_RUNTIME_MAX_POOL3D_GRAD_WITH_ARGMAX_H15#ifndef OPS_BUILD_IN_OP_TILING_RUNTIME_MAX_POOL3D_GRAD_WITH_ARGMAX_H
24#define OPS_BUILD_IN_OP_TILING_RUNTIME_MAX_POOL3D_GRAD_WITH_ARGMAX_H16#define OPS_BUILD_IN_OP_TILING_RUNTIME_MAX_POOL3D_GRAD_WITH_ARGMAX_H
@@ -29,6 +21,7 @@
29#include "tiling_base/tiling_base.h"21#include "tiling_base/tiling_base.h"
30#include "tiling_base/tiling_templates_registry.h"22#include "tiling_base/tiling_templates_registry.h"
31#include "util/math_util.h"23#include "util/math_util.h"
24+#include "../../pool_3d_common/op_host/arch32/max_pool3d_grad_tiling_constants.h"
32 25 
33namespace optiling {26namespace optiling {
34using Ops::NN::Optiling::TilingBaseClass;27using Ops::NN::Optiling::TilingBaseClass;
@@ -116,43 +109,10 @@ struct UBBufferSize {
116 uint64_t valSize;109 uint64_t valSize;
117};110};
118 111 
119-// Index const
120-constexpr uint32_t X_INDEX = 0;
121-constexpr uint32_t GRAD_INDEX = 1;
122-constexpr uint32_t ARGMAX_INDEX = 2;
123constexpr uint32_t Y_INDEX = 0;112constexpr uint32_t Y_INDEX = 0;
124-constexpr size_t KSIZE_ATTR_INDEX = 0U;
125-constexpr size_t STRIDES_ATTR_INDEX = 1U;
126-constexpr size_t PADS_ATTR_INDEX = 2U;
127-constexpr size_t DILATION_ATTR_INDEX = 3U;
128-constexpr size_t CEIL_MODE_ATTR_INDEX = 4U;
129-// Params const
130constexpr uint64_t NUM_TWO = 2;113constexpr uint64_t NUM_TWO = 2;
131-constexpr size_t NC_DIM_NUM = 2;
132constexpr size_t DHW_DIM_NUM = 3;114constexpr size_t DHW_DIM_NUM = 3;
133-constexpr size_t NCDHW_DIM_NUM = 5;
134-constexpr uint32_t DTYPE_LEN_B8 = 1;
135-constexpr uint32_t DTYPE_LEN_B16 = 2;
136-constexpr uint32_t DTYPE_LEN_B32 = 4;
137-constexpr uint32_t BLOCK_SIZE = 32;
138constexpr uint32_t MAX_BLOCK_COUNT = 4095;115constexpr uint32_t MAX_BLOCK_COUNT = 4095;
139-constexpr uint64_t MAX_INT32 = 2147483647;
140-constexpr uint32_t NUM_PER_REP_B16 = 128;
141-constexpr uint32_t NUM_PER_REP_B32 = 64;
142-constexpr uint32_t SELECT_RESERVED_UB_SIZE = 8192;
143-// Tiling const
144-constexpr uint32_t TILING_OVERLAP = 100;
145-constexpr uint32_t TILING_UB_NO_CUT = 0;
146-constexpr uint32_t TILING_UB_CUT_NC = 10;
147-constexpr uint32_t TILING_UB_CUT_DO = 20;
148-constexpr uint32_t TILING_UB_CUT_HO = 30;
149-constexpr uint32_t TILING_UB_CUT_WO = 40;
150-constexpr uint32_t TILING_UB_CUT_KD = 50;
151-constexpr uint32_t TILING_UB_CUT_KH = 60;
152-constexpr uint32_t TILING_UB_CUT_KW = 70;
153-constexpr uint32_t TILING_TYPE_NORMAL = 0;
154-constexpr uint32_t TILING_TYPE_CUTK = 1;
155-constexpr uint32_t TILING_TYPE_SCATTER = 2;
156 116 
157struct Tiling4MaxPool3DGradWithArgmaxCompileInfo {117struct Tiling4MaxPool3DGradWithArgmaxCompileInfo {
158 platform_ascendc::SocVersion curSocVersion = platform_ascendc::SocVersion::ASCEND910B;118 platform_ascendc::SocVersion curSocVersion = platform_ascendc::SocVersion::ASCEND910B;
@@ -228,6 +188,8 @@ struct MaxPoolGradWithArgmaxTilingParams {
228 uint32_t ubCutAxis{0};188 uint32_t ubCutAxis{0};
229 bool ceilMode{false};189 bool ceilMode{false};
230 bool isOverLap{false};190 bool isOverLap{false};
191+ bool isInitOutput{false};
192+ bool isNeedWorkspace{false};
231};193};
232 194 
233class MaxPool3DGradWithArgmaxTilingBase : public TilingBaseClass {195class MaxPool3DGradWithArgmaxTilingBase : public TilingBaseClass {
Mpooling/max_pool3d_grad_with_argmax/op_host/max_pool3d_grad_with_argmax_tiling_base.cpp+19-6
@@ -260,15 +260,15 @@ ge::graphStatus MaxPool3DGradWithArgmaxTilingBase::CheckInputValid()
260 int64_t doExpected, hoExpected, woExpected;260 int64_t doExpected, hoExpected, woExpected;
261 if (maxPoolGradParams.ceilMode) {261 if (maxPoolGradParams.ceilMode) {
262 doExpected =262 doExpected =
263- Ops::Base::CeilDiv((maxPoolGradParams.diDim + NUM_TWO * pDTop - dilationD * (kd - 1) - 1), sd) + 1;263+ Ops::Base::CeilDiv((maxPoolGradParams.diDim + NUM_TWO * pDTop + sd - dilationD * (kd - 1) - 1), sd);
264 hoExpected =264 hoExpected =
265- Ops::Base::CeilDiv((maxPoolGradParams.hiDim + NUM_TWO * pHTop - dilationH * (kh - 1) - 1), sh) + 1;265+ Ops::Base::CeilDiv((maxPoolGradParams.hiDim + NUM_TWO * pHTop + sh - dilationH * (kh - 1) - 1), sh);
266 woExpected =266 woExpected =
267- Ops::Base::CeilDiv((maxPoolGradParams.wiDim + NUM_TWO * pWTop - dilationW * (kw - 1) - 1), sw) + 1;267+ Ops::Base::CeilDiv((maxPoolGradParams.wiDim + NUM_TWO * pWTop + sw - dilationW * (kw - 1) - 1), sw);
268 } else {268 } else {
269- doExpected = (maxPoolGradParams.diDim + NUM_TWO * pDTop - dilationD * (kd - 1) - 1) / sd + 1;269+ doExpected = (maxPoolGradParams.diDim + NUM_TWO * pDTop + sd - dilationD * (kd - 1) - 1) / sd;
270- hoExpected = (maxPoolGradParams.hiDim + NUM_TWO * pHTop - dilationH * (kh - 1) - 1) / sh + 1;270+ hoExpected = (maxPoolGradParams.hiDim + NUM_TWO * pHTop + sh - dilationH * (kh - 1) - 1) / sh;
271- woExpected = (maxPoolGradParams.wiDim + NUM_TWO * pWTop - dilationW * (kw - 1) - 1) / sw + 1;271+ woExpected = (maxPoolGradParams.wiDim + NUM_TWO * pWTop + sw - dilationW * (kw - 1) - 1) / sw;
272 }272 }
273 doExpected = ((doExpected - 1) * sd >= maxPoolGradParams.diDim + pDTop) ? doExpected - 1 : doExpected;273 doExpected = ((doExpected - 1) * sd >= maxPoolGradParams.diDim + pDTop) ? doExpected - 1 : doExpected;
274 hoExpected = ((hoExpected - 1) * sh >= maxPoolGradParams.hiDim + pHTop) ? hoExpected - 1 : hoExpected;274 hoExpected = ((hoExpected - 1) * sh >= maxPoolGradParams.hiDim + pHTop) ? hoExpected - 1 : hoExpected;
@@ -430,6 +430,19 @@ void MaxPool3DGradWithArgmaxTilingBase::SetBaseTilingData()
430 430 
431void MaxPool3DGradWithArgmaxTilingBase::PrintTilingData()431void MaxPool3DGradWithArgmaxTilingBase::PrintTilingData()
432{432{
433+ maxPoolGradParams.isInitOutput =
434+ (maxPoolGradParams.doDim * maxPoolGradParams.kd < maxPoolGradParams.diDim + maxPoolGradParams.pDTop +
435+ maxPoolGradParams.pDBottom) ||
436+ (maxPoolGradParams.hoDim * maxPoolGradParams.kh < maxPoolGradParams.hiDim + maxPoolGradParams.pHTop +
437+ maxPoolGradParams.pHBottom) ||
438+ (maxPoolGradParams.woDim * maxPoolGradParams.kw < maxPoolGradParams.wiDim + maxPoolGradParams.pWTop +
439+ maxPoolGradParams.pWBottom) ||
440+ (maxPoolGradParams.doDim - 1) * maxPoolGradParams.sd + maxPoolGradParams.kd < maxPoolGradParams.diDim +
441+ maxPoolGradParams.pDTop ||
442+ (maxPoolGradParams.hoDim - 1) * maxPoolGradParams.sh + maxPoolGradParams.kh < maxPoolGradParams.hiDim +
443+ maxPoolGradParams.pHTop ||
444+ (maxPoolGradParams.woDim - 1) * maxPoolGradParams.sw + maxPoolGradParams.kw < maxPoolGradParams.wiDim +
445+ maxPoolGradParams.pWTop || maxPoolGradParams.isOverLap;
433 OP_LOGI(446 OP_LOGI(
434 context_->GetNodeName(),447 context_->GetNodeName(),
435 "TilingData nc: %lu, di: %lu, hi: %lu, wi: %lu do: %lu, ho: %lu, wo: %lu, "448 "TilingData nc: %lu, di: %lu, hi: %lu, wi: %lu do: %lu, ho: %lu, wo: %lu, "
Mpooling/max_pool3d_grad_with_argmax/op_kernel/max_pool3d_grad_with_argmax.cpp+1-9
@@ -41,15 +41,7 @@ using namespace MaxPool3DGradWithArgmax;
41extern "C" __global__ __aicore__ void max_pool3d_grad_with_argmax(41extern "C" __global__ __aicore__ void max_pool3d_grad_with_argmax(
42 GM_ADDR x, GM_ADDR grad, GM_ADDR argmax, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)42 GM_ADDR x, GM_ADDR grad, GM_ADDR argmax, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
43{43{
44- if (workspace == nullptr) {44+ if (workspace == nullptr || GetUserWorkspace(workspace) == nullptr || g_coreType == AIC) {
45- return;
46- }
47- 
48- GM_ADDR userWS = GetUserWorkspace(workspace);
49- if (userWS == nullptr) {
50- return;
51- }
52- if (g_coreType == AIC) {
53 return;45 return;
54 }46 }
55 47 
Mpooling/max_pool3d_grad_with_argmax/op_kernel/max_pool3d_grad_with_argmax_base.h+1-1
@@ -8,7 +8,7 @@
8 * See LICENSE in the root of the software repository for the full text of the License.8 * See LICENSE in the root of the software repository for the full text of the License.
9 */9 */
10 10 
11-/*!11+/* !
12 * \file max_pool3d_grad_with_argmax_base.h12 * \file max_pool3d_grad_with_argmax_base.h
13 * \brief13 * \brief
14 */14 */
Mpooling/max_pool3d_grad_with_argmax/op_kernel/max_pool3d_grad_with_argmax_common.h+5-174
@@ -14,57 +14,29 @@
14 */14 */
15#ifndef MAX_POOL_GRAD3D_WITH_ARGMAX_COMMON15#ifndef MAX_POOL_GRAD3D_WITH_ARGMAX_COMMON
16#define MAX_POOL_GRAD3D_WITH_ARGMAX_COMMON16#define MAX_POOL_GRAD3D_WITH_ARGMAX_COMMON
17-#include "kernel_tiling/kernel_tiling.h"17+ 
18+#include "../pool_3d_common/arch32/max_pool3d_grad_common.h"
18 19 
19namespace MaxPool3DGradWithArgmaxComm {20namespace MaxPool3DGradWithArgmaxComm {
20using namespace AscendC;21using namespace AscendC;
22+using namespace MaxPool3DGradCommon;
21 23 
22constexpr uint64_t VL_FP32 = 64;24constexpr uint64_t VL_FP32 = 64;
23constexpr uint64_t VL_FP16 = 128;25constexpr uint64_t VL_FP16 = 128;
24constexpr uint64_t MAX_LIST_NUM = 4;26constexpr uint64_t MAX_LIST_NUM = 4;
25-constexpr uint64_t TRANS_ADDR_LEN = 16;
26-constexpr uint64_t BLOCK_SIZE = 32;
27-constexpr uint64_t BLOCK_NUM_16 = BLOCK_SIZE / sizeof(half);
28-constexpr uint64_t BLOCK_NUM_32 = BLOCK_SIZE / sizeof(float);
29constexpr uint64_t LARGE_KERNEL = 0;27constexpr uint64_t LARGE_KERNEL = 0;
30constexpr uint64_t LARGE_HO = 1;28constexpr uint64_t LARGE_HO = 1;
31constexpr uint64_t LARGE_WO = 2;29constexpr uint64_t LARGE_WO = 2;
32constexpr uint64_t UINT16_BITS = 16;30constexpr uint64_t UINT16_BITS = 16;
33-constexpr uint64_t UINT8_BITS = 8;
34-constexpr float ZERO = 0.0f;
35constexpr uint64_t B32_VECTOR_MASK = 64;31constexpr uint64_t B32_VECTOR_MASK = 64;
36constexpr uint64_t B16_VECTOR_MASK = 128;32constexpr uint64_t B16_VECTOR_MASK = 128;
37const uint64_t NUM_TWO = 2;33const uint64_t NUM_TWO = 2;
38const uint64_t BITSIZE = 16;34const uint64_t BITSIZE = 16;
39- 
40constexpr uint32_t FLOAT_BLOCK_ELEM = 8;35constexpr uint32_t FLOAT_BLOCK_ELEM = 8;
41constexpr uint32_t MAX_REP_NUM = 255;36constexpr uint32_t MAX_REP_NUM = 255;
42-constexpr uint32_t UNIT_BLOCK_LEN = 32;
43 37 
44-struct BlockParams {38+struct BlockParams : public BlockParamsCommon {
45- uint64_t ncCntIndex = 0;
46- uint64_t doCntIndex = 0;
47- uint64_t hoCntIndex = 0;
48- uint64_t woCntIndex = 0;
49- uint64_t ncShape = 0;
50- uint64_t doShape = 0;
51- uint64_t hoShape = 0;
52- uint64_t woShape = 0;
53- uint64_t diShape = 0;
54- uint64_t hiShape = 0;
55- uint64_t wiShape = 0;
56 uint64_t wiShapeAlign = 0;39 uint64_t wiShapeAlign = 0;
57- uint64_t diValid = 0;
58- uint64_t hiValid = 0;
59- uint64_t wiValid = 0;
60- uint64_t dohowoShape = 0;
61- uint64_t dihiwiAlign = 0;
62- uint64_t dohowoAlign8 = 0;
63- uint64_t dohowoAlign16 = 0;
64- uint64_t offsetX = 0;
65- uint64_t offsetGrad = 0;
66- uint64_t offsetArgmax = 0;
67- uint64_t offsetY = 0;
68 uint64_t offsetYD = 0;40 uint64_t offsetYD = 0;
69 uint64_t offsetYH = 0;41 uint64_t offsetYH = 0;
70 uint64_t offsetYW = 0;42 uint64_t offsetYW = 0;
@@ -74,18 +46,9 @@ struct BlockParams {
74 uint64_t padDBottom = 0;46 uint64_t padDBottom = 0;
75 uint64_t padHBottom = 0;47 uint64_t padHBottom = 0;
76 uint64_t padWBottom = 0;48 uint64_t padWBottom = 0;
77- uint64_t baseNcOffset = 0;
78- uint64_t ShapeSum = 0;
79};49};
80 50 
81-struct TilingParams {51+struct TilingParams : public TilingParamsCommon {
82- uint64_t ncDim;
83- uint64_t diDim;
84- uint64_t hiDim;
85- uint64_t wiDim;
86- uint64_t doDim;
87- uint64_t hoDim;
88- uint64_t woDim;
89 uint64_t kd;52 uint64_t kd;
90 uint64_t kh;53 uint64_t kh;
91 uint64_t kw;54 uint64_t kw;
@@ -98,14 +61,6 @@ struct TilingParams {
98 uint64_t padHBottom;61 uint64_t padHBottom;
99 uint64_t padWTop;62 uint64_t padWTop;
100 uint64_t padWBottom;63 uint64_t padWBottom;
101- uint64_t singleCoreNc;
102- uint64_t singleCoreDo;
103- uint64_t singleCoreHo;
104- uint64_t singleCoreWo;
105- uint64_t baseNc;
106- uint64_t baseDo;
107- uint64_t baseHo;
108- uint64_t baseWo;
109 uint64_t baseDoHoWo;64 uint64_t baseDoHoWo;
110 uint64_t baseDoHoWoAlign8;65 uint64_t baseDoHoWoAlign8;
111 uint64_t baseDoHoWoAlign16;66 uint64_t baseDoHoWoAlign16;
@@ -117,134 +72,10 @@ struct TilingParams {
117 uint64_t baseDiHiWiAlign8;72 uint64_t baseDiHiWiAlign8;
118 uint64_t baseDiHiWiAlign16;73 uint64_t baseDiHiWiAlign16;
119 uint64_t baseDiHiWiAlign;74 uint64_t baseDiHiWiAlign;
120- uint64_t ncCnt;
121- uint64_t doCnt;
122- uint64_t hoCnt;
123- uint64_t woCnt;
124- uint64_t ncTail;
125- uint64_t doTail;
126- uint64_t hoTail;
127- uint64_t woTail;
128- uint64_t totalCnt;
129- uint64_t needInitOutput;
130 uint64_t padGmOffset;75 uint64_t padGmOffset;
131- uint64_t ubSize;
132- uint64_t usedCoreNum;
133- uint64_t preCoreNum;
134- uint64_t round;
135- uint64_t realRound;
136 uint64_t outputDataSize;76 uint64_t outputDataSize;
137- uint64_t ncIndex;
138- uint64_t ncCntRound;
139- uint64_t ncRealRound;
140- uint64_t diHiWiLen;
141- uint64_t initLen;
142- uint64_t initOffset;
143};77};
144 78 
145-__aicore__ inline uint64_t CeilDiv(uint64_t x, uint64_t y)
146-{
147- return y == 0 ? x : (x + y - 1) / y;
148-}
149- 
150-template <typename Tp, Tp v>
151-struct integral_constant {
152- static constexpr Tp value = v;
153-};
154-using true_type = integral_constant<bool, true>;
155-using false_type = integral_constant<bool, false>;
156-template <typename, typename>
157-struct is_same : public false_type {};
158-template <typename Tp>
159-struct is_same<Tp, Tp> : public true_type {};
160- 
161-// only support float/int32_t
162-// [row, col] -> [col, row]: row:align16, col:align8
163-template <typename T>
164-__aicore__ inline void TransposeBase16M8(
165- const LocalTensor<T>& dstUb, const LocalTensor<T>& srcUb, uint64_t rowNum, uint64_t colNum)
166-{
167- uint64_t srcAddrList[TRANS_ADDR_LEN];
168- uint64_t dstAddrList[TRANS_ADDR_LEN];
169- 
170- for (uint64_t r = 0; r < rowNum / TRANS_ADDR_LEN; r++) {
171- for (uint64_t i = 0; i < TRANS_ADDR_LEN; i++) {
172- srcAddrList[i] = (uint64_t)(srcUb[r * TRANS_ADDR_LEN * colNum + i * colNum].GetPhyAddr());
173- dstAddrList[i] = (uint64_t)(dstUb[r * TRANS_ADDR_LEN + i / 2 * rowNum + i % 2 * BLOCK_NUM_32].GetPhyAddr());
174- }
175- struct TransDataTo5HDParams transDataParams;
176- transDataParams.repeatTimes = colNum / BLOCK_NUM_32;
177- if (transDataParams.repeatTimes == 1) {
178- transDataParams.srcRepStride = 0;
179- transDataParams.dstRepStride = 0;
180- } else {
181- transDataParams.srcRepStride = 1;
182- transDataParams.dstRepStride = rowNum;
183- }
184- 
185- TransDataTo5HD<T>(dstAddrList, srcAddrList, transDataParams);
186- }
187-}
188- 
189-// only support float/int32_t
190-// [row, col] -> [col, row]: row:align8, col:align16
191-template <typename T>
192-__aicore__ inline void TransposeBase8M16(
193- const LocalTensor<T>& dstUb, const LocalTensor<T>& srcUb, uint64_t rowNum, uint64_t colNum)
194-{
195- uint64_t srcAddrList[TRANS_ADDR_LEN];
196- uint64_t dstAddrList[TRANS_ADDR_LEN];
197- 
198- for (uint64_t r = 0; r < colNum / TRANS_ADDR_LEN; r++) {
199- for (uint64_t i = 0; i < TRANS_ADDR_LEN; i++) {
200- srcAddrList[i] =
201- (uint64_t)(srcUb[r * TRANS_ADDR_LEN + i % BLOCK_NUM_32 * colNum + i / BLOCK_NUM_32 * BLOCK_NUM_32]
202- .GetPhyAddr());
203- dstAddrList[i] =
204- (uint64_t)(dstUb[r * TRANS_ADDR_LEN * rowNum + (i % 2 * BLOCK_NUM_32 + i / 2) * rowNum].GetPhyAddr());
205- }
206- struct TransDataTo5HDParams transDataParams;
207- transDataParams.repeatTimes = rowNum / BLOCK_NUM_32;
208- if (transDataParams.repeatTimes == 1) {
209- transDataParams.srcRepStride = 0;
210- transDataParams.dstRepStride = 0;
211- } else {
212- transDataParams.srcRepStride = colNum;
213- transDataParams.dstRepStride = 1;
214- }
215- 
216- TransDataTo5HD<T>(dstAddrList, srcAddrList, transDataParams);
217- }
218-}
219- 
220-// only support float16/bfloat16
221-// [row, col] -> [col, row]: row:align16, col:align16
222-template <typename T>
223-__aicore__ inline void TransposeBase16M16(
224- const LocalTensor<T>& dstUb, const LocalTensor<T>& srcUb, uint64_t rowNum, uint64_t colNum)
225-{
226- uint64_t srcAddrList[TRANS_ADDR_LEN];
227- uint64_t dstAddrList[TRANS_ADDR_LEN];
228- 
229- for (uint64_t r = 0; r < rowNum / TRANS_ADDR_LEN; r++) {
230- for (uint64_t i = 0; i < TRANS_ADDR_LEN; i++) {
231- srcAddrList[i] = (uint64_t)(srcUb[r * TRANS_ADDR_LEN * colNum + i * colNum].GetPhyAddr());
232- dstAddrList[i] = (uint64_t)(dstUb[r * TRANS_ADDR_LEN + i * rowNum].GetPhyAddr());
233- }
234- struct TransDataTo5HDParams transDataParams;
235- transDataParams.repeatTimes = colNum / BLOCK_NUM_16;
236- if (transDataParams.repeatTimes == 1) {
237- transDataParams.srcRepStride = 0;
238- transDataParams.dstRepStride = 0;
239- } else {
240- transDataParams.srcRepStride = 1;
241- transDataParams.dstRepStride = rowNum;
242- }
243- // TransData5HD does not support bfloat16, only use half
244- TransDataTo5HD<half>(dstAddrList, srcAddrList, transDataParams);
245- }
246-}
247- 
248// only support float/int32_t79// only support float/int32_t
249// [row, col] -> [col, row]: row:align16,max:64, col:align880// [row, col] -> [col, row]: row:align16,max:64, col:align8
250__aicore__ inline void TransposeAddrBase16M8(81__aicore__ inline void TransposeAddrBase16M8(
Mpooling/max_pool3d_grad_with_argmax/op_kernel/max_pool3d_grad_with_argmax_scatter.h+16-116
@@ -8,6 +8,7 @@
8 * See LICENSE in the root of the software repository for the full text of the License.8 * See LICENSE in the root of the software repository for the full text of the License.
9 */9 */
10 10 
11+ 
11/* !12/* !
12 * \file max_pool3d_grad_with_argmax_scatter.h13 * \file max_pool3d_grad_with_argmax_scatter.h
13 * \brief14 * \brief
@@ -18,131 +19,30 @@
18#include "kernel_operator.h"19#include "kernel_operator.h"
19#include "max_pool3d_grad_with_argmax_common.h"20#include "max_pool3d_grad_with_argmax_common.h"
20#include "max_pool3d_grad_with_argmax_scatter_base.h"21#include "max_pool3d_grad_with_argmax_scatter_base.h"
22+#include "../pool_3d_common/arch32/max_pool3d_grad_scatter_unified.h"
21 23 
22namespace MaxPool3DGradWithArgmax {24namespace MaxPool3DGradWithArgmax {
23using namespace AscendC;25using namespace AscendC;
24using namespace MaxPool3DGradWithArgmaxComm;26using namespace MaxPool3DGradWithArgmaxComm;
25 27 
26template <typename TX, typename TGrad, typename TArgmax, typename TY>28template <typename TX, typename TGrad, typename TArgmax, typename TY>
27-class MaxPoolGradWithArgScatter : public MaxPoolGradWithArgScatterBase<TX, TGrad, TArgmax, TY> {29+class MaxPoolGradWithArgScatter :
30+ public MaxPool3DGradCommon::MaxPool3DGradScatterUnified<
31+ TX, TGrad, TArgmax, TY,
32+ MaxPool3DGradWithArgmaxTilingData,
33+ TilingParams,
34+ BlockParams,
35+ MaxPool3DGradScatterInternal::MaxPool3DGradScatterBaseTemplate>
36+{
28public:37public:
29 __aicore__ inline MaxPoolGradWithArgScatter(TPipe* pipe)38 __aicore__ inline MaxPoolGradWithArgScatter(TPipe* pipe)
30- : MaxPoolGradWithArgScatterBase<TX, TGrad, TArgmax, TY>(pipe)39+ : MaxPool3DGradCommon::MaxPool3DGradScatterUnified<
40+ TX, TGrad, TArgmax, TY,
41+ MaxPool3DGradWithArgmaxTilingData,
42+ TilingParams,
43+ BlockParams,
44+ MaxPool3DGradScatterInternal::MaxPool3DGradScatterBaseTemplate>(pipe)
31 {}45 {}
32- __aicore__ inline void Init(
33- GM_ADDR x, GM_ADDR grad, GM_ADDR argmax, GM_ADDR y, GM_ADDR usrWorkspace,
34- const MaxPool3DGradWithArgmaxTilingData* __restrict__ tiling)
35- {
36- // load tiling data
37- this->InitParams(tiling);
38- // set global buffer
39- this->InitInputsOutputs(x, grad, argmax, y, usrWorkspace);
40- // init global memory
41- InitYGMGlobalMemory(x, grad, argmax, y, usrWorkspace);
42- // ub buffer init
43- this->InitUbBuffer();
44- this->pipe_->InitBuffer(this->yQue, 1, BLOCK_SIZE);
45- }
46- 
47- __aicore__ inline void InitYGMGlobalMemory(GM_ADDR x, GM_ADDR grad, GM_ADDR argmax, GM_ADDR y, GM_ADDR usrWorkspace)
48- {
49- InitGlobalMemory(this->yGm, this->params_.initLen, static_cast<TY>(0));
50- event_t eventMTE3ToS = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE3_S));
51- SetFlag<HardEvent::MTE3_S>(eventMTE3ToS);
52- WaitFlag<HardEvent::MTE3_S>(eventMTE3ToS);
53- }
54- 
55- __aicore__ inline void CalcOutOffset()
56- {
57- LocalTensor<TArgmax> argmaxUb = this->argmaxQue.template DeQue<TArgmax>(); // need free in the end
58- LocalTensor<TGrad> gradUb = this->gradQue.template DeQue<TGrad>(); // need free in the end
59- LocalTensor<TGrad> yUb = this->yQue.template AllocTensor<TGrad>(); // need free in the end
60- 
61- event_t eventMTE2ToS = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE2_S));
62- SetFlag<HardEvent::MTE2_S>(eventMTE2ToS);
63- WaitFlag<HardEvent::MTE2_S>(eventMTE2ToS);
64- 
65- uint64_t basedhwLen = this->block_.doShape * this->block_.hoShape * this->block_.woShape;
66- for (uint64_t ncIdx = 0; ncIdx < this->block_.ncShape; ncIdx++) {
67- uint64_t ncOffset = this->block_.baseNcOffset * this->params_.diHiWiLen;
68- for (uint64_t dhwIdx = 0; dhwIdx < basedhwLen; dhwIdx++) {
69- uint64_t ubOffset = ncIdx * basedhwLen + dhwIdx;
70- uint64_t outOffset = ncOffset + (uint64_t)argmaxUb.GetValue(ubOffset);
71- 
72- TGrad gradValue = gradUb.GetValue(ubOffset);
73- yUb.SetValue(0, gradValue);
74- event_t eventSToMTE3 = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::S_MTE3));
75- SetFlag<HardEvent::S_MTE3>(eventSToMTE3);
76- WaitFlag<HardEvent::S_MTE3>(eventSToMTE3);
77- DataCopyPad(this->yGm[outOffset], yUb, {1, static_cast<uint16_t>(sizeof(TGrad)), 0, 0});
78- event_t eventMTE3ToS = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE3_S));
79- SetFlag<HardEvent::MTE3_S>(eventMTE3ToS);
80- WaitFlag<HardEvent::MTE3_S>(eventMTE3ToS);
81- }
82- this->block_.ShapeSum += basedhwLen;
83- if (this->block_.ShapeSum == this->params_.doDim * this->params_.hoDim * this->params_.woDim) {
84- this->block_.baseNcOffset += 1;
85- this->block_.ShapeSum = 0;
86- }
87- }
88- 
89- this->gradQue.FreeTensor(gradUb);
90- this->argmaxQue.FreeTensor(argmaxUb);
91- this->yQue.FreeTensor(yUb);
92- }
93- 
94- __aicore__ inline void CalcBlock()
95- {
96- this->CopyInGrad();
97- this->CopyInArgmax();
98- CalcOutOffset();
99- PipeBarrier<PIPE_ALL>();
100- }
101- 
102- __aicore__ inline void Process()
103- {
104- uint64_t totalCnt = this->params_.totalCnt; // 需要处理base矩阵个数
105- uint64_t round = this->params_.round; // 每个核需要处理多少个base矩阵，向上取整
106- uint64_t preCoreNum = this->params_.preCoreNum; // 每个核均分完后剩余base由前preCoreNum个核进行填充
107- uint64_t index = this->params_.ncIndex; // 当前核填充数的起始index
108- uint64_t realRound = this->params_.realRound; // 填充轮数
109- 
110- uint64_t ncIndex = this->params_.ncIndex;
111- for (uint64_t i = 0; i < this->params_.ncRealRound; i++) {
112- if (ncIndex < this->params_.ncCnt) {
113- this->block_.ncCntIndex = ncIndex;
114- this->block_.ncShape = this->block_.ncCntIndex >= (this->params_.ncCnt - 1UL) ? this->params_.ncTail :
115- this->params_.baseNc;
116- for (uint64_t j = 0; j < this->params_.doCnt; j++) {
117- this->block_.doCntIndex = j;
118- this->block_.doShape = this->block_.doCntIndex >= (this->params_.doCnt - 1) ? this->params_.doTail :
119- this->params_.baseDo;
120- for (uint64_t k = 0; k < this->params_.hoCnt; k++) {
121- this->block_.hoCntIndex = k;
122- this->block_.hoShape = this->block_.hoCntIndex >= (this->params_.hoCnt - 1) ?
123- this->params_.hoTail :
124- this->params_.baseHo;
125- for (uint64_t l = 0; l < this->params_.woCnt; l++) {
126- this->block_.woCntIndex = l;
127- this->block_.woShape = this->block_.woCntIndex >= (this->params_.woCnt - 1) ?
128- this->params_.woTail :
129- this->params_.baseWo;
130- this->block_.offsetGrad =
131- this->block_.ncCntIndex * this->params_.baseNc * this->params_.doDim *
132- this->params_.hoDim * this->params_.woDim +
133- this->block_.doCntIndex * this->params_.baseDo * this->params_.hoDim *
134- this->params_.woDim +
135- this->block_.hoCntIndex * this->params_.baseHo * this->params_.woDim +
136- this->block_.woCntIndex * this->params_.baseWo;
137- this->block_.offsetArgmax = this->block_.offsetGrad;
138- CalcBlock();
139- }
140- }
141- }
142- ncIndex += 1; // 当前ncCntIndex
143- }
144- }
145- }
146};46};
147} // namespace MaxPool3DGradWithArgmax47} // namespace MaxPool3DGradWithArgmax
148#endif // MAX_POOL_GRAD3D_WITH_ARGMAX_SCATTER_H48#endif // MAX_POOL_GRAD3D_WITH_ARGMAX_SCATTER_H
Rpooling/max_pool3d_grad_with_argmax/op_host/CMakeLists.txt→pooling/max_pool3d_grad_with_argmax/tests/ut/op_api/CMakeLists.txt+9-6
Rpooling/max_pool3d_grad_with_argmax/tests/ut/op_host/test_aclnn_max_pool2d_with_indices_backward.cpp→pooling/max_pool3d_grad_with_argmax/tests/ut/op_api/test_aclnn_max_pool2d_with_indices_backward.cpp+5-5
Rpooling/max_pool3d_grad_with_argmax/tests/ut/op_host/test_aclnn_max_pool2d_with_mask_backward.cpp→pooling/max_pool3d_grad_with_argmax/tests/ut/op_api/test_aclnn_max_pool2d_with_mask_backward.cpp+5-5
Rpooling/max_pool3d_grad_with_argmax/tests/ut/op_host/test_aclnn_max_pool3d_with_argmax_backward.cpp→pooling/max_pool3d_grad_with_argmax/tests/ut/op_api/test_aclnn_max_pool3d_with_argmax_backward.cpp+5-5