Star901
1818
代码
代码
Issues226
Pull Requests315
Actions
项目讨论
Wiki
分析
Star901
1818
已合并
Add AdaptiveMaxPool2d #1670
huohuo_wangyan创建于 2月9日
Add AdaptiveMaxPool2d #1670
已合并
huohuo_wangyan创建于 2月9日
从huohuo_wangyan/ops-nnAdaptiveMaxPool2d合入到cann/ops-nnmaster
共 24 个文件变更+3050-337
Mclassify_rule.yaml+2-0文件内容审核中，请稍后刷新重试
Apooling/adaptive_max_pool2d/CMakeLists.txt+11-0
@@ -0,0 +1,11 @@
1+# ----------------------------------------------------------------------------
2+# This program is free software, you can redistribute it and/or modify.
3+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
4+# This file is a part of the CANN Open Software.
5+# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6+# Please refer to the License for details. You may not use this file except in compliance with the License.
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.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------
10+ 
11+add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE adaptive_max_pool2d ACLNNTYPE aclnn_exclude DEPENDENCIES max_pool3d_with_argmax_v2 adaptive_max_pool3d)
Apooling/adaptive_max_pool2d/README.md+107-0
@@ -0,0 +1,107 @@
1+# AdaptiveMaxPool2d
2+ 
3+## 产品支持情况
4+ 
5+| 产品 | 是否支持 |
6+| ---- | :----:|
7+|Ascend 950PR/Ascend 950DT|√|
8+|Atlas A3 训练系列产品/Atlas A3 推理系列产品|x|
9+|Atlas A2 训练系列产品/Atlas A2 推理系列产品|x|
10+|Atlas 200I/500 A2推理产品|×|
11+|Atlas 推理系列产品|×|
12+|Atlas 训练系列产品|×|
13+|Kirin X90 处理器系列产品|x|
14+|Kirin 9030 处理器系列产品|x|
15+ 
16+## 功能说明
17+ 
18+- 算子功能：根据输入的output_size计算每次kernel的大小，对输入x进行3维最大池化操作，输出池化后的值y和索引indices。
19+ 
20+- 计算公式：
21+ y tensor中对于DHW轴上每个位置为$(l,m,n)$的元素来说，其计算公式为：
22+ $$
23+ H^{m}_{left} = floor((m*H)/H_o)
24+ $$
25+
26+ $$
27+ H^{m}_{right} = ceil(((m+1)*H)/H_o)
28+ $$
29+
30+ $$
31+ W^{n}_{left} = floor((n*W)/W_o)
32+ $$
33+
34+ $$
35+ W^{n}_{right} = ceil(((n+1)*W)/W_o)
36+ $$
37+
38+ $$
39+ y(N,C,m,n)=\underset {j\in [H^m_{left},H^m_{right}], k \in [W^n_{left},W^n_{right}] }{max} input(N,C,j,k)
40+ $$
41+
42+ $$
43+ indices(N,C,m,n)=\underset {j\in [H^m_{left},H^m_{right}], k \in [W^n_{left},W^n_{right}] }{argmax} input(N,C,j,k)
44+ $$
45+ 
46+## 参数说明
47+ 
48+<table style="undefined;table-layout: fixed; width: 1250px"><colgroup>
49+ <col style="width: 150px">
50+ <col style="width: 150px">
51+ <col style="width: 500px">
52+ <col style="width: 250px">
53+ <col style="width: 200px">
54+ </colgroup>
55+ <thead>
56+ <tr>
57+ <th>参数名</th>
58+ <th>输入/输出/属性</th>
59+ <th>描述</th>
60+ <th>数据类型</th>
61+ <th>数据格式</th>
62+ </tr></thead>
63+ <tbody>
64+ <tr>
65+ <td>x</td>
66+ <td>输入</td>
67+ <td>待进行AdaptiveMaxPool2d计算的入参。H轴W轴2个维度的乘积不能大于int32的最大表示。数据类型与出参`y`的保持一致。</td>
68+ <td>FLOAT、FLOAT16、BFLOAT16</td>
69+ <td>ND</td>
70+ </tr>
71+ <tr>
72+ <td>output_size</td>
73+ <td>属性</td>
74+ <td>表示输出结果在H，W维度上的空间大小。数据类型与入参`x`的保持一致。</td>
75+ <td>INT64</td>
76+ <td>-</td>
77+ </tr>
78+ <tr>
79+ <td>y</td>
80+ <td>输出</td>
81+ <td>待进行AdaptiveMaxPool2d计算的出参。shape与出参`indices`的保持一致，数据类型与入参`x`的保持一致。</td>
82+ <td>FLOAT、FLOAT16、BFLOAT16</td>
83+ <td>ND</td>
84+ </tr>
85+ <tr>
86+ <td>indices</td>
87+ <td>输出</td>
88+ <td>表示`y`元素在输入`x`中的索引位置。shape与出参`y`的保持一致。</td>
89+ <td>INT32</td>
90+ <td>ND</td>
91+ </tr>
92+ </tbody></table>
93+ 
94+ 
95+## 约束说明
96+Shape描述：
97+ - self.shape = (N, C, Hin, Win)
98+ - outputSize = [Hout, Wout]
99+ - outputOut.shape = (N, C, Hout, Wout)
100+ - indicesOut.shape = (N, C, Hout, Wout)
101+ 
102+ 
103+## 调用说明
104+ 
105+| 调用方式 | 样例代码 | 说明 |
106+| ---------------- | --------------------------- | --------------------------------------------------- |
107+| aclnn接口 | [test_aclnn_adaptive_max_pool2d.cpp](examples/test_aclnn_adaptive_max_pool2d.cpp) | 通过[aclnnAdaptiveMaxPool2d](docs/aclnnAdaptiveMaxPool2d.md)接口方式调用AdaptiveMaxPool2d算子。 |
Apooling/adaptive_max_pool2d/docs/aclnnAdaptiveMaxPool2d.md+396-0
@@ -0,0 +1,396 @@
1+`# aclnnAdaptiveMaxPool2d
2+ 
3+[📄 查看源码](https://gitcode.com/cann/ops-nn/tree/master/pooling/adaptive_max_pool3d)
4+ 
5+## 产品支持情况
6+ 
7+| 产品 | 是否支持 |
8+| :----------------------------------------------------------- | :------: |
9+| <term>Ascend 950PR/Ascend 950DT</term> | √ |
10+| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |
11+| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |
12+| <term>Atlas 200I/500 A2 推理产品</term> | × |
13+| <term>Atlas 推理系列产品</term> | × |
14+| <term>Atlas 训练系列产品</term> | √ |
15+ 
16+## 功能说明
17+ 
18+根据输入的outputSize计算每次kernel的大小，对输入self进行2维最大池化操作，输出池化后的值out和索引indices。aclnnAdaptiveMaxPool2d与aclnnMaxPool2d的区别在于，只需指定outputSize大小，并按outputSize的大小来划分pooling区域。
19+ 
20+## 函数原型
21+每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnAdaptiveMaxPool2dGetWorkspaceSize”接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用“aclnnAdaptiveMaxPool2d”接口执行计算。
22+ 
23+```Cpp
24+aclnnStatus aclnnAdaptiveMaxPool2dGetWorkspaceSize(
25+ const aclTensor *self,
26+ const aclIntArray *outputSize,
27+ aclTensor *outputOut,
28+ aclTensor *indicesOut,
29+ uint64_t *workspaceSize,
30+ aclOpExecutor **executor)
31+```
32+```Cpp
33+aclnnStatus aclnnAdaptiveMaxPool2d(
34+ void *workspace,
35+ uint64_t workspaceSize,
36+ aclOpExecutor *executor,
37+ aclrtStream stream)
38+```
39+ 
40+## aclnnAdaptiveMaxPool2dGetWorkspaceSize
41+ 
42+- **参数说明：**
43+ 
44+ <table style="undefined;table-layout: fixed; width: 1478px"><colgroup>
45+ <col style="width: 149px">
46+ <col style="width: 121px">
47+ <col style="width: 264px">
48+ <col style="width: 253px">
49+ <col style="width: 262px">
50+ <col style="width: 148px">
51+ <col style="width: 135px">
52+ <col style="width: 146px">
53+ </colgroup>
54+ <thead>
55+ <tr>
56+ <th>参数名</th>
57+ <th>输入/输出</th>
58+ <th>描述</th>
59+ <th>使用说明</th>
60+ <th>数据类型</th>
61+ <th>数据格式</th>
62+ <th>维度(shape)</th>
63+ <th>非连续Tensor</th>
64+ </tr></thead>
65+ <tbody>
66+ <tr>
67+ <td>self</td>
68+ <td>输入</td>
69+ <td>待计算张量。</td>
70+ <td>与outputOut的数据类型一致。</td>
71+ <td>FLOAT32、FLOAT16、BFLOAT16、DOUBLE</td>
72+ <td>NCHW、NCL、NHWC</td>
73+ <td>3-4</td>
74+ <td>√</td>
75+ </tr>
76+ <tr>
77+ <td>outputSize</td>
78+ <td>输入</td>
79+ <td>表示输出结果在H和W维度上的空间大小。</td>
80+ <td>-</td>
81+ <td>-</td>
82+ <td>-</td>
83+ <td>-</td>
84+ <td>-</td>
85+ </tr>
86+ <tr>
87+ <td>outputOut</td>
88+ <td>输出</td>
89+ <td>池化后的结果。</td>
90+ <td>与self的数据类型一致，shape与indicesOut一致。</td>
91+ <td>FLOAT32、FLOAT16、BFLOAT16、DOUBLE</td>
92+ <td>NCHW、NCL、NHWC</td>
93+ <td>3-4</td>
94+ <td>√</td>
95+ </tr>
96+ <tr>
97+ <td>indicesOut</td>
98+ <td>输出</td>
99+ <td>outputOut元素在输入self中的索引位置。</td>
100+ <td>shape与outputOut一致。</td>
101+ <td>INT64</td>
102+ <td>NCHW、NCL、NHWC</td>
103+ <td>3-4</td>
104+ <td>√</td>
105+ </tr>
106+ <tr>
107+ <td>workspaceSize</td>
108+ <td>输出</td>
109+ <td>返回需要在Device侧申请的workspace大小。</td>
110+ <td>-</td>
111+ <td>-</td>
112+ <td>-</td>
113+ <td>-</td>
114+ <td>-</td>
115+ </tr>
116+ <tr>
117+ <td>executor</td>
118+ <td>输出</td>
119+ <td>返回op执行器，包含了算子计算流程。</td>
120+ <td>-</td>
121+ <td>-</td>
122+ <td>-</td>
123+ <td>-</td>
124+ <td>-</td>
125+ </tr>
126+ </tbody></table>
127+ 
128+- **返回值：**
129+ 
130+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
131+ 
132+ 第一段接口完成入参校验，出现以下场景时报错：
133+ 
134+ <table style="undefined;table-layout: fixed; width: 1166px"><colgroup>
135+ <col style="width: 267px">
136+ <col style="width: 124px">
137+ <col style="width: 775px">
138+ </colgroup>
139+ <thead>
140+ <tr>
141+ <th>返回码</th>
142+ <th>错误码</th>
143+ <th>描述</th>
144+ </tr></thead>
145+ <tbody>
146+ <tr>
147+ <td>ACLNN_ERR_PARAM_NULLPTR</td>
148+ <td>161001</td>
149+ <td>传入的self、outputSize、outputOut或indicesOut是空指针。</td>
150+ </tr>
151+ <tr>
152+ <td rowspan="10">ACLNN_ERR_PARAM_INVALID</td>
153+ <td rowspan="10">161002</td>
154+ <td>self的数据类型不在支持的范围之内。</td>
155+ </tr>
156+ <tr>
157+ <td>self和outputOut的数据类型不一致。</td>
158+ </tr>
159+ <tr>
160+ <td>indicesOut的数据类型不为int64。</td>
161+ </tr>
162+ <tr>
163+ <td>self的shape不是3维或者4维。</td>
164+ </tr>
165+ <tr>
166+ <td>self在非第一维度上的size小于1。</td>
167+ </tr>
168+ <tr>
169+ <td>outputOut和indicesOut的shape不一致。</td>
170+ </tr>
171+ <tr>
172+ <td>outputSize的size大小不等于2。</td>
173+ </tr>
174+ <tr>
175+ <td>outputSize中元素值小于等于0。</td>
176+ </tr>
177+ <tr>
178+ <td>outputOut的shape与实际输出shape不匹配。</td>
179+ </tr>
180+ <tr>
181+ <td>self的format不是NCHW/NHWC/NCL这三种format。</td>
182+ </tr>
183+ <tr>
184+ <td>ACLNN_ERR_INNER_NULLPTR</td>
185+ <td>561103</td>
186+ <td>API内部校验错误，通常由于输入数据或属性的规格不在支持的范围之内导致。</td>
187+ </tr>
188+ </tbody>
189+ </table>
190+ 
191+## aclnnAdaptiveMaxPool2d
192+ 
193+- **参数说明：**
194+ <table style="undefined;table-layout: fixed; width: 1166px"><colgroup>
195+ <col style="width: 173px">
196+ <col style="width: 133px">
197+ <col style="width: 860px">
198+ </colgroup>
199+ <thead>
200+ <tr>
201+ <th>参数名</th>
202+ <th>输入/输出</th>
203+ <th>描述</th>
204+ </tr></thead>
205+ <tbody>
206+ <tr>
207+ <td>workspace</td>
208+ <td>输入</td>
209+ <td>在Device侧申请的workspace内存地址。</td>
210+ </tr>
211+ <tr>
212+ <td>workspaceSize</td>
213+ <td>输入</td>
214+ <td>在Device侧申请的workspace大小，由第一段接口aclnnAdaptiveMaxPool2dGetWorkspaceSize获取。</td>
215+ </tr>
216+ <tr>
217+ <td>executor</td>
218+ <td>输入</td>
219+ <td>op执行器，包含了算子计算流程。</td>
220+ </tr>
221+ <tr>
222+ <td>stream</td>
223+ <td>输入</td>
224+ <td>指定执行任务的Stream。</td>
225+ </tr>
226+ </tbody>
227+ </table>
228+- **返回值：**
229+ 
230+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
231+ 
232+## 约束说明
233+- 确定性计算：
234+ - aclnnAdaptiveMaxPool2d默认确定性实现。
235+ 
236+- Shape描述：
237+ - self.shape = (N, C, Hin, Win) 或者 (C, Hin, Win) 或 (N, Hin, Win, C)
238+ - outputSize = [Hout, Wout]
239+ - outputOut.shape = (N, C, Hout, Wout) 或者 (C, Hout, Wout) 或 (N, Hout, Wout, C)
240+ - indicesOut.shape = (N, C, Hout, Wout) 或者 (C, Hout, Wout) 或 (N, Hout, Wout, C)
241+ 
242+## 调用示例
243+示例代码如下，仅供参考，具体编译和执行过程请参考[编译与运行样例](../../../docs/zh/context/编译与运行样例.md)。
244+```Cpp
245+#include <iostream>
246+#include <vector>
247+#include "acl/acl.h"
248+#include "aclnnop/aclnn_adaptive_max_pool2d.h"
249+ 
250+#define CHECK_RET(cond, return_expr) \
251+ do { \
252+ if (!(cond)) { \
253+ return_expr; \
254+ } \
255+ } while (0)
256+ 
257+#define LOG_PRINT(message, ...) \
258+ do { \
259+ printf(message, ##__VA_ARGS__); \
260+ } while (0)
261+ 
262+int64_t GetShapeSize(const std::vector<int64_t>& shape) {
263+ int64_t shapeSize = 1;
264+ for (auto i : shape) {
265+ shapeSize *= i;
266+ }
267+ return shapeSize;
268+}
269+ 
270+int Init(int32_t deviceId, aclrtStream* stream) {
271+ // 固定写法，资源初始化
272+ auto ret = aclInit(nullptr);
273+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
274+ ret = aclrtSetDevice(deviceId);
275+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
276+ ret = aclrtCreateStream(stream);
277+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
278+ return 0;
279+}
280+ 
281+template <typename T>
282+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
283+ aclDataType dataType, aclTensor** tensor) {
284+ auto size = GetShapeSize(shape) * sizeof(T);
285+ // 调用aclrtMalloc申请device侧内存
286+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
287+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
288+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
289+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
290+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
291+ 
292+ // 计算连续tensor的strides
293+ std::vector<int64_t> strides(shape.size(), 1);
294+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
295+ strides[i] = shape[i + 1] * strides[i + 1];
296+ }
297+ 
298+ // 调用aclCreateTensor接口创建aclTensor
299+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_NCHW,
300+ shape.data(), shape.size(), *deviceAddr);
301+ return 0;
302+}
303+ 
304+int main() {
305+ // 1. （固定写法）device/stream初始化，参考acl API手册
306+ // 根据自己的实际device填写deviceId
307+ int32_t deviceId = 0;
308+ aclrtStream stream;
309+ auto ret = Init(deviceId, &stream);
310+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
311+ 
312+ // 2. 构造输入与输出，需要根据API的接口自定义构造
313+ std::vector<int64_t> selfShape = {1, 1, 4, 4};
314+ std::vector<int64_t> outShape = {1, 1, 2, 2};
315+ void* selfDeviceAddr = nullptr;
316+ void* outDeviceAddr = nullptr;
317+ void* indDeviceAddr = nullptr;
318+ aclTensor* self = nullptr;
319+ aclTensor* out = nullptr;
320+ aclTensor* indices = nullptr;
321+ std::vector<float> selfHostData = {0, 1, 2, 3, 4, 5, 6, 7,
322+ 8, 9, 10, 11, 12, 13, 14, 15};
323+ std::vector<float> outHostData = {0, 0, 0, 0.0};
324+ std::vector<int64_t> indicesHostData = {0, 0, 0, 0};
325+ // 创建self aclTensor
326+ ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self);
327+ CHECK_RET(ret == ACL_SUCCESS, return ret);
328+ 
329+ // 创建out aclTensor
330+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
331+ CHECK_RET(ret == ACL_SUCCESS, return ret);
332+ 
333+ // 创建indices aclTensor
334+ ret = CreateAclTensor(indicesHostData, outShape, &indDeviceAddr, aclDataType::ACL_INT64, &indices);
335+ CHECK_RET(ret == ACL_SUCCESS, return ret);
336+ 
337+ std::vector<int64_t> arraySize = {2, 2};
338+ const aclIntArray *outputSize = aclCreateIntArray(arraySize.data(), arraySize.size());
339+ CHECK_RET(outputSize != nullptr, return ACL_ERROR_INTERNAL_ERROR);
340+ 
341+ // 3. 调用CANN算子库API，需要修改为具体的Api名称
342+ uint64_t workspaceSize = 0;
343+ aclOpExecutor* executor;
344+ // 调用aclnnAdaptiveMaxPool2d第一段接口
345+ ret = aclnnAdaptiveMaxPool2dGetWorkspaceSize(self, outputSize, out, indices, &workspaceSize, &executor);
346+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnAdaptiveMaxPool2dGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
347+ // 根据第一段接口计算出的workspaceSize申请device内存
348+ void* workspaceAddr = nullptr;
349+ if (workspaceSize > 0) {
350+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
351+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
352+ }
353+ // 调用aclnnAdaptiveMaxPool2d第二段接口
354+ ret = aclnnAdaptiveMaxPool2d(workspaceAddr, workspaceSize, executor, stream);
355+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnAdaptiveMaxPool2d failed. ERROR: %d\n", ret); return ret);
356+ 
357+ // 4. （固定写法）同步等待任务执行结束
358+ ret = aclrtSynchronizeStream(stream);
359+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
360+ 
361+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
362+ auto size = GetShapeSize(outShape);
363+ std::vector<float> outData(size, 0);
364+ std::vector<int64_t> indicesData(size, 0);
365+ ret = aclrtMemcpy(outData.data(), outData.size() * sizeof(outData[0]), outDeviceAddr,
366+ size * sizeof(outData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
367+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
368+ ret = aclrtMemcpy(indicesData.data(), indicesData.size() * sizeof(indicesData[0]), indDeviceAddr,
369+ size * sizeof(indicesData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
370+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
371+ for (int64_t i = 0; i < size; i++) {
372+ LOG_PRINT("out[%ld] is: %f\n", i, outData[i]);
373+ }
374+ for (int64_t i = 0; i < size; i++) {
375+ LOG_PRINT("indices[%ld] is: %ld\n", i, indicesData[i]);
376+ }
377+ 
378+ // 6. 释放aclTensor和aclScalar，需要根据具体API的接口定义修改
379+ aclDestroyTensor(self);
380+ aclDestroyTensor(out);
381+ aclDestroyTensor(indices);
382+ aclDestroyIntArray(outputSize);
383+ 
384+ // 7. 释放device 资源
385+ aclrtFree(selfDeviceAddr);
386+ aclrtFree(outDeviceAddr);
387+ aclrtFree(indDeviceAddr);
388+ if (workspaceSize > 0) {
389+ aclrtFree(workspaceAddr);
390+ }
391+ aclrtDestroyStream(stream);
392+ aclrtResetDevice(deviceId);
393+ aclFinalize();
394+ return 0;
395+}
396+````
Rpooling/adaptive_max_pool3d/examples/test_aclnn_adaptive_max_pool2d.cpp→pooling/adaptive_max_pool2d/examples/test_aclnn_adaptive_max_pool2d.cpp+161-161
@@ -1,162 +1,162 @@
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 <iostream>11+#include <iostream>
12-#include <vector>12+#include <vector>
13-#include "acl/acl.h"13+#include "acl/acl.h"
14-#include "aclnnop/aclnn_adaptive_max_pool2d.h"14+#include "aclnnop/aclnn_adaptive_max_pool2d.h"
15- 15+ 
16-#define CHECK_RET(cond, return_expr) \16+#define CHECK_RET(cond, return_expr) \
17- do { \17+ do { \
18- if (!(cond)) { \18+ if (!(cond)) { \
19- return_expr; \19+ return_expr; \
20- } \20+ } \
21- } while (0)21+ } while (0)
22- 22+ 
23-#define LOG_PRINT(message, ...) \23+#define LOG_PRINT(message, ...) \
24- do { \24+ do { \
25- printf(message, ##__VA_ARGS__); \25+ printf(message, ##__VA_ARGS__); \
26- } while (0)26+ } while (0)
27- 27+ 
28-int64_t GetShapeSize(const std::vector<int64_t>& shape) {28+int64_t GetShapeSize(const std::vector<int64_t>& shape) {
29- int64_t shapeSize = 1;29+ int64_t shapeSize = 1;
30- for (auto i : shape) {30+ for (auto i : shape) {
31- shapeSize *= i;31+ shapeSize *= i;
32- }32+ }
33- return shapeSize;33+ return shapeSize;
34-}34+}
35- 35+ 
36-int Init(int32_t deviceId, aclrtStream* stream) {36+int Init(int32_t deviceId, aclrtStream* stream) {
37- // 固定写法，资源初始化37+ // 固定写法，资源初始化
38- auto ret = aclInit(nullptr);38+ auto ret = aclInit(nullptr);
39- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);39+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
40- ret = aclrtSetDevice(deviceId);40+ ret = aclrtSetDevice(deviceId);
41- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);41+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
42- ret = aclrtCreateStream(stream);42+ ret = aclrtCreateStream(stream);
43- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);43+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
44- return 0;44+ return 0;
45-}45+}
46- 46+ 
47-template <typename T>47+template <typename T>
48-int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,48+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
49- aclDataType dataType, aclTensor** tensor) {49+ aclDataType dataType, aclTensor** tensor) {
50- auto size = GetShapeSize(shape) * sizeof(T);50+ auto size = GetShapeSize(shape) * sizeof(T);
51- // 调用aclrtMalloc申请device侧内存51+ // 调用aclrtMalloc申请device侧内存
52- auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);52+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
53- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);53+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
54- // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上54+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
55- ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);55+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
56- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);56+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
57- 57+ 
58- // 计算连续tensor的strides58+ // 计算连续tensor的strides
59- std::vector<int64_t> strides(shape.size(), 1);59+ std::vector<int64_t> strides(shape.size(), 1);
60- for (int64_t i = shape.size() - 2; i >= 0; i--) {60+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
61- strides[i] = shape[i + 1] * strides[i + 1];61+ strides[i] = shape[i + 1] * strides[i + 1];
62- }62+ }
63- 63+ 
64- // 调用aclCreateTensor接口创建aclTensor64+ // 调用aclCreateTensor接口创建aclTensor
65- *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_NCHW,65+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_NCHW,
66- shape.data(), shape.size(), *deviceAddr);66+ shape.data(), shape.size(), *deviceAddr);
67- return 0;67+ return 0;
68-}68+}
69- 69+ 
70-int main() {70+int main() {
71- // 1. （固定写法）device/stream初始化，参考acl API手册71+ // 1. （固定写法）device/stream初始化，参考acl API手册
72- // 根据自己的实际device填写deviceId72+ // 根据自己的实际device填写deviceId
73- int32_t deviceId = 0;73+ int32_t deviceId = 0;
74- aclrtStream stream;74+ aclrtStream stream;
75- auto ret = Init(deviceId, &stream);75+ auto ret = Init(deviceId, &stream);
76- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);76+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
77- 77+ 
78- // 2. 构造输入与输出，需要根据API的接口自定义构造78+ // 2. 构造输入与输出，需要根据API的接口自定义构造
79- std::vector<int64_t> selfShape = {1, 2, 3, 4};79+ std::vector<int64_t> selfShape = {1, 2, 3, 4};
80- std::vector<int64_t> outShape = {1, 2, 2, 2};80+ std::vector<int64_t> outShape = {1, 2, 2, 2};
81- void* selfDeviceAddr = nullptr;81+ void* selfDeviceAddr = nullptr;
82- void* outDeviceAddr = nullptr;82+ void* outDeviceAddr = nullptr;
83- void* indDeviceAddr = nullptr;83+ void* indDeviceAddr = nullptr;
84- aclTensor* self = nullptr;84+ aclTensor* self = nullptr;
85- aclTensor* out = nullptr;85+ aclTensor* out = nullptr;
86- aclTensor* indices = nullptr;86+ aclTensor* indices = nullptr;
87- std::vector<float> selfHostData = {0, 1, 2, 3, 4, 5, 6, 7,87+ std::vector<float> selfHostData = {0, 1, 2, 3, 4, 5, 6, 7,
88- 8, 9, 10, 11, 12, 13, 14, 15,88+ 8, 9, 10, 11, 12, 13, 14, 15,
89- 16, 17, 18, 19, 20, 21, 22, 23};89+ 16, 17, 18, 19, 20, 21, 22, 23};
90- std::vector<float> outHostData = {0, 0, 0, 0.0};90+ std::vector<float> outHostData = {0, 0, 0, 0.0};
91- std::vector<int64_t> indicesHostData = {0, 0, 0, 0};91+ std::vector<int64_t> indicesHostData = {0, 0, 0, 0};
92- // 创建self aclTensor92+ // 创建self aclTensor
93- ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self);93+ ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self);
94- CHECK_RET(ret == ACL_SUCCESS, return ret);94+ CHECK_RET(ret == ACL_SUCCESS, return ret);
95- 95+ 
96- // 创建out aclTensor96+ // 创建out aclTensor
97- ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);97+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
98- CHECK_RET(ret == ACL_SUCCESS, return ret);98+ CHECK_RET(ret == ACL_SUCCESS, return ret);
99- 99+ 
100- // 创建indices aclTensor100+ // 创建indices aclTensor
101- ret = CreateAclTensor(indicesHostData, outShape, &indDeviceAddr, aclDataType::ACL_INT64, &indices);101+ ret = CreateAclTensor(indicesHostData, outShape, &indDeviceAddr, aclDataType::ACL_INT64, &indices);
102- CHECK_RET(ret == ACL_SUCCESS, return ret);102+ CHECK_RET(ret == ACL_SUCCESS, return ret);
103- 103+ 
104- std::vector<int64_t> arraySize = {2, 2};104+ std::vector<int64_t> arraySize = {2, 2};
105- const aclIntArray *outputSize = aclCreateIntArray(arraySize.data(), arraySize.size());105+ const aclIntArray *outputSize = aclCreateIntArray(arraySize.data(), arraySize.size());
106- CHECK_RET(outputSize != nullptr, return ACL_ERROR_INTERNAL_ERROR);106+ CHECK_RET(outputSize != nullptr, return ACL_ERROR_INTERNAL_ERROR);
107- 107+ 
108- // 3. 调用CANN算子库API，需要修改为具体的Api名称108+ // 3. 调用CANN算子库API，需要修改为具体的Api名称
109- uint64_t workspaceSize = 0;109+ uint64_t workspaceSize = 0;
110- aclOpExecutor* executor;110+ aclOpExecutor* executor;
111- // 调用aclnnAdaptiveMaxPool2d第一段接口111+ // 调用aclnnAdaptiveMaxPool2d第一段接口
112- ret = aclnnAdaptiveMaxPool2dGetWorkspaceSize(self, outputSize, out, indices, &workspaceSize, &executor);112+ ret = aclnnAdaptiveMaxPool2dGetWorkspaceSize(self, outputSize, out, indices, &workspaceSize, &executor);
113- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnAdaptiveMaxPool2dGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);113+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnAdaptiveMaxPool2dGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
114- // 根据第一段接口计算出的workspaceSize申请device内存114+ // 根据第一段接口计算出的workspaceSize申请device内存
115- void* workspaceAddr = nullptr;115+ void* workspaceAddr = nullptr;
116- if (workspaceSize > 0) {116+ if (workspaceSize > 0) {
117- ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);117+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
118- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);118+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
119- }119+ }
120- // 调用aclnnAdaptiveMaxPool2d第二段接口120+ // 调用aclnnAdaptiveMaxPool2d第二段接口
121- ret = aclnnAdaptiveMaxPool2d(workspaceAddr, workspaceSize, executor, stream);121+ ret = aclnnAdaptiveMaxPool2d(workspaceAddr, workspaceSize, executor, stream);
122- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnAdaptiveMaxPool2d failed. ERROR: %d\n", ret); return ret);122+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnAdaptiveMaxPool2d failed. ERROR: %d\n", ret); return ret);
123- 123+ 
124- // 4. （固定写法）同步等待任务执行结束124+ // 4. （固定写法）同步等待任务执行结束
125- ret = aclrtSynchronizeStream(stream);125+ ret = aclrtSynchronizeStream(stream);
126- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);126+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
127- 127+ 
128- // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改128+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
129- auto size = GetShapeSize(outShape);129+ auto size = GetShapeSize(outShape);
130- std::vector<float> outData(size, 0);130+ std::vector<float> outData(size, 0);
131- std::vector<int64_t> indicesData(size, 0);131+ std::vector<int64_t> indicesData(size, 0);
132- ret = aclrtMemcpy(outData.data(), outData.size() * sizeof(outData[0]), outDeviceAddr,132+ ret = aclrtMemcpy(outData.data(), outData.size() * sizeof(outData[0]), outDeviceAddr,
133- size * sizeof(outData[0]), ACL_MEMCPY_DEVICE_TO_HOST);133+ size * sizeof(outData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
134- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);134+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
135- ret = aclrtMemcpy(indicesData.data(), indicesData.size() * sizeof(indicesData[0]), indDeviceAddr,135+ ret = aclrtMemcpy(indicesData.data(), indicesData.size() * sizeof(indicesData[0]), indDeviceAddr,
136- size * sizeof(indicesData[0]), ACL_MEMCPY_DEVICE_TO_HOST);136+ size * sizeof(indicesData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
137- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);137+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
138- for (int64_t i = 0; i < size; i++) {138+ for (int64_t i = 0; i < size; i++) {
139- LOG_PRINT("out[%ld] is: %f\n", i, outData[i]);139+ LOG_PRINT("out[%ld] is: %f\n", i, outData[i]);
140- }140+ }
141- for (int64_t i = 0; i < size; i++) {141+ for (int64_t i = 0; i < size; i++) {
142- LOG_PRINT("indices[%ld] is: %ld\n", i, indicesData[i]);142+ LOG_PRINT("indices[%ld] is: %ld\n", i, indicesData[i]);
143- }143+ }
144- 144+ 
145- // 6. 释放aclTensor和aclScalar，需要根据具体API的接口定义修改145+ // 6. 释放aclTensor和aclScalar，需要根据具体API的接口定义修改
146- aclDestroyTensor(self);146+ aclDestroyTensor(self);
147- aclDestroyTensor(out);147+ aclDestroyTensor(out);
148- aclDestroyTensor(indices);148+ aclDestroyTensor(indices);
149- aclDestroyIntArray(outputSize);149+ aclDestroyIntArray(outputSize);
150- 150+ 
151- // 7. 释放device 资源151+ // 7. 释放device 资源
152- aclrtFree(selfDeviceAddr);152+ aclrtFree(selfDeviceAddr);
153- aclrtFree(outDeviceAddr);153+ aclrtFree(outDeviceAddr);
154- aclrtFree(indDeviceAddr);154+ aclrtFree(indDeviceAddr);
155- if (workspaceSize > 0) {155+ if (workspaceSize > 0) {
156- aclrtFree(workspaceAddr);156+ aclrtFree(workspaceAddr);
157- }157+ }
158- aclrtDestroyStream(stream);158+ aclrtDestroyStream(stream);
159- aclrtResetDevice(deviceId);159+ aclrtResetDevice(deviceId);
160- aclFinalize();160+ aclFinalize();
161- return 0;161+ return 0;
162}162}
Rpooling/adaptive_max_pool3d/op_api/aclnn_adaptive_max_pool2d.cpp→pooling/adaptive_max_pool2d/op_api/aclnn_adaptive_max_pool2d.cpp+379-172
@@ -9,18 +9,17 @@
9 */9 */
10 10 
11/*!11/*!
12- * \file aclnn_mv.cpp12+ * \file aclnn_adaptive_max_pool2d.cpp
13 * \brief13 * \brief
14 */14 */
15 15 
16#include "aclnn_kernels/contiguous.h"16#include "aclnn_kernels/contiguous.h"
17-#include "pooling/adaptive_avg_pool3d/op_api/shape_op.h"
18#include "aclnn_kernels/transdata.h"17#include "aclnn_kernels/transdata.h"
19#include "level0/squeeze.h"18#include "level0/squeeze.h"
20#include "level0/unsqueeze.h"19#include "level0/unsqueeze.h"
21#include "adaptive_max_pool2d.h"20#include "adaptive_max_pool2d.h"
22-#include "adaptive_max_pool3d.h"
23#include "level0/max_pool3d_with_argmax_v2.h"21#include "level0/max_pool3d_with_argmax_v2.h"
22+#include "../../adaptive_max_pool3d/op_api/adaptive_max_pool3d.h"
24#include "aclnn_adaptive_max_pool2d.h"23#include "aclnn_adaptive_max_pool2d.h"
25#include "aclnn/aclnn_base.h"24#include "aclnn/aclnn_base.h"
26#include "aclnn_kernels/reshape.h"25#include "aclnn_kernels/reshape.h"
@@ -71,8 +70,17 @@ static bool IsSelfDtypeDouble(const aclTensor* self)
71 70 
72static bool IsSocVersion910B()71static bool IsSocVersion910B()
73{72{
74- if (op::GetCurrentPlatformInfo().GetSocVersion() >= op::SocVersion::ASCEND910B &&73+ auto curArch = GetCurrentPlatformInfo().GetCurNpuArch();
75- op::GetCurrentPlatformInfo().GetSocVersion() <= op::SocVersion::ASCEND910E) {74+ if (curArch == NpuArch::DAV_2201) {
75+ return true;
76+ }
77+ return false;
78+}
79+ 
80+static bool IsSocVersion910D()
81+{
82+ auto curArch = GetCurrentPlatformInfo().GetCurNpuArch();
83+ if (Ops::NN::AclnnUtil::IsRegbase(curArch)) {
76 return true;84 return true;
77 }85 }
78 return false;86 return false;
@@ -80,8 +88,7 @@ static bool IsSocVersion910B()
80 88 
81static const std::initializer_list<DataType>& GetDtypeSupportList()89static const std::initializer_list<DataType>& GetDtypeSupportList()
82{90{
83- if (GetCurrentPlatformInfo().GetSocVersion() >= SocVersion::ASCEND910B &&91+ if (IsSocVersion910B() || IsSocVersion910D()) {
84- GetCurrentPlatformInfo().GetSocVersion() <= SocVersion::ASCEND910E) {
85 return DTYPE_SUPPORT_LIST_ASCEND910B;92 return DTYPE_SUPPORT_LIST_ASCEND910B;
86 } else {93 } else {
87 return DTYPE_SUPPORT_LIST;94 return DTYPE_SUPPORT_LIST;
@@ -187,12 +194,366 @@ static aclnnStatus CheckParams(
187 return ACLNN_SUCCESS;194 return ACLNN_SUCCESS;
188}195}
189 196 
197+static inline const aclTensor* View3Das5D(const aclTensor* input, aclOpExecutor* executor)
198+{
199+ // CHW -> unsqueeze -> reformat -> NCDHW
200+ // unsqueeze input into 4D
201+ const aclTensor* unsqueezedInput = l0op::UnsqueezeNd(input, 1, executor);
202+ CHECK_RET(unsqueezedInput != nullptr, nullptr);
203+ // unsqueeze input into 5D
204+ auto unsqueezedInput5D = l0op::UnsqueezeNd(unsqueezedInput, static_cast<int64_t>(0), executor);
205+ CHECK_RET(unsqueezedInput5D != nullptr, nullptr);
206+ // reformat to NCDHW
207+ auto reformatInput = l0op::ReFormat(unsqueezedInput5D, op::Format::FORMAT_NCDHW);
208+ CHECK_RET(reformatInput != nullptr, nullptr);
209+ 
210+ return reformatInput;
211+}
212+ 
213+static inline const aclTensor* View4Das5D(const aclTensor* input, aclOpExecutor* executor)
214+{
215+ // NCHW -> unsqueeze -> reformat -> NCDHW
216+ // unsqueeze input into 5D
217+ auto unsqueezedInput = l0op::UnsqueezeNd(input, 2, executor);
218+ CHECK_RET(unsqueezedInput != nullptr, nullptr);
219+ // reformat to NCDHW
220+ auto reformatInput = l0op::ReFormat(unsqueezedInput, op::Format::FORMAT_NCDHW);
221+ CHECK_RET(reformatInput != nullptr, nullptr);
222+ 
223+ return reformatInput;
224+}
225+ 
226+static inline const aclTensor* View5Das3D(const aclTensor* input, const op::Format& format, aclOpExecutor* executor)
227+{
228+ // NCDHW -> squeeze -> reformat -> CHW
229+ // squeeze out into 4D
230+ const aclTensor* squeezedInput = l0op::SqueezeNd(input, 2, executor);
231+ CHECK_RET(squeezedInput != nullptr, nullptr);
232+ // squeeze out into 3D
233+ auto squeezedInput3D = l0op::SqueezeNd(squeezedInput, static_cast<int64_t>(0), executor);
234+ CHECK_RET(squeezedInput != nullptr, nullptr);
235+ // reformat to NCL
236+ auto reformatInput = l0op::ReFormat(squeezedInput3D, format);
237+ CHECK_RET(reformatInput != nullptr, nullptr);
238+ 
239+ return reformatInput;
240+}
241+ 
242+static inline const aclTensor* View5Das4D(const aclTensor* input, const op::Format& format, aclOpExecutor* executor)
243+{
244+ // NCDHW -> squeeze -> reformat -> NCHW
245+ // squeeze out into 3D
246+ auto squeezedInput = l0op::SqueezeNd(input, 2, executor);
247+ CHECK_RET(squeezedInput != nullptr, nullptr);
248+ // reformat to NCHW
249+ auto reformatInput = l0op::ReFormat(squeezedInput, format);
250+ CHECK_RET(reformatInput != nullptr, nullptr);
251+ 
252+ return reformatInput;
253+}
254+ 
255+static const aclTensor* View3Das4D(const aclTensor* input, aclOpExecutor* executor)
256+{
257+ // NCL -> unsqueeze -> reformat -> NCHW
258+ // unsqueeze input into 4D
259+ const int64_t appendDim[] = {0};
260+ aclIntArray* dimUnsqueeze = executor->AllocIntArray(appendDim, 1);
261+ CHECK_RET(dimUnsqueeze != nullptr, nullptr);
262+ auto unsqueezedInput = l0op::UnsqueezeNd(input, dimUnsqueeze, executor);
263+ CHECK_RET(unsqueezedInput != nullptr, nullptr);
264+ // reformat to NCHW
265+ auto reformatInput = l0op::ReFormat(unsqueezedInput, op::Format::FORMAT_NCHW);
266+ CHECK_RET(reformatInput != nullptr, nullptr);
267+ 
268+ return reformatInput;
269+}
270+ 
271+static aclnnStatus ProcessAndCopyResults(
272+ const aclTensor* outResult,
273+ const aclTensor* indicesResultCast,
274+ aclTensor* outputOut,
275+ aclTensor* indicesOut,
276+ aclOpExecutor* executor)
277+{
278+ auto resShapeVector = op::ToShapeVector(outputOut->GetViewShape());
279+ aclIntArray* resShapeArray =
280+ executor->AllocIntArray(resShapeVector.data(), resShapeVector.size());
281+ CHECK_RET(resShapeArray != nullptr, ACLNN_ERR_INNER_NULLPTR);
282+ auto resTranReshapeOut = l0op::Reshape(outResult, resShapeArray, executor);
283+ CHECK_RET(resTranReshapeOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
284+ 
285+ resTranReshapeOut = l0op::ReFormat(resTranReshapeOut, outputOut->GetStorageFormat(), executor);
286+ CHECK_RET(resTranReshapeOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
287+ 
288+ auto viewCopyResultOut = l0op::ViewCopy(resTranReshapeOut, outputOut, executor);
289+ CHECK_RET(viewCopyResultOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
290+ 
291+ auto resTranReshapeIndices = l0op::Reshape(indicesResultCast, resShapeArray, executor);
292+ CHECK_RET(resTranReshapeIndices != nullptr, ACLNN_ERR_INNER_NULLPTR);
293+ 
294+ resTranReshapeIndices =
295+ l0op::ReFormat(resTranReshapeIndices, indicesOut->GetStorageFormat(), executor);
296+ CHECK_RET(resTranReshapeIndices != nullptr, ACLNN_ERR_INNER_NULLPTR);
297+ 
298+ auto viewCopyResultIndices = l0op::ViewCopy(resTranReshapeIndices, indicesOut, executor);
299+ CHECK_RET(viewCopyResultIndices != nullptr, ACLNN_ERR_INNER_NULLPTR);
300+
301+ return ACLNN_SUCCESS;
302+}
303+ 
304+// 转MaxPool3D
305+static aclnnStatus HandleMaxPool3DCase(
306+ const aclTensor* inputContiguousReshape,
307+ const std::vector<int64_t>& kernelSizeArr,
308+ aclTensor* outputOut,
309+ aclTensor* indicesOut,
310+ aclOpExecutor* executor)
311+{
312+ const aclIntArray* kernelSize = aclCreateIntArray(kernelSizeArr.data(), kernelSizeArr.size());
313+ const aclIntArray* stride = aclCreateIntArray(kernelSizeArr.data(), kernelSizeArr.size());
314+ std::vector<int64_t> paddingArr = {0, 0, 0};
315+ const aclIntArray* padding = aclCreateIntArray(paddingArr.data(), paddingArr.size());
316+ std::vector<int64_t> dilationArr = {1, 1, 1};
317+ const aclIntArray* dilation = aclCreateIntArray(dilationArr.data(), dilationArr.size());
318+ bool ceilMode = false;
319+ std::string dataFormat = "NCDHW";
320+ 
321+ inputContiguousReshape =
322+ l0op::ReFormat(inputContiguousReshape, op::Format::FORMAT_NCDHW, executor);
323+ CHECK_RET(inputContiguousReshape != nullptr, ACLNN_ERR_INNER_NULLPTR);
324+ 
325+ auto [outResult, indicesResult] = l0op::MaxPool3DWithArgmaxV2Ncdhw(
326+ inputContiguousReshape, kernelSize, stride, padding, dilation, ceilMode, dataFormat,
327+ executor);
328+ CHECK_RET(outResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
329+ CHECK_RET(indicesResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
330+ auto indicesResultCast = l0op::Cast(indicesResult, op::DataType::DT_INT64, executor);
331+ CHECK_RET(indicesResultCast != nullptr, ACLNN_ERR_INNER_NULLPTR);
332+ 
333+ auto ret = ProcessAndCopyResults(outResult, indicesResultCast, outputOut, indicesOut, executor);
334+ if (ret != ACLNN_SUCCESS) {
335+ return ret;
336+ }
337+ aclDestroyIntArray(kernelSize);
338+ aclDestroyIntArray(stride);
339+ aclDestroyIntArray(padding);
340+ aclDestroyIntArray(dilation);
341+
342+ return ACLNN_SUCCESS;
343+}
344+ 
345+// AdaptiveMaxPool2d的处理逻辑
346+static aclnnStatus ProcessAdaptiveMaxPoolResults(
347+ const aclTensor* outResult,
348+ const aclTensor* indicesResult,
349+ aclTensor* outputOut,
350+ aclTensor* indicesOut,
351+ aclOpExecutor* executor)
352+{
353+ auto indicesResultCast = l0op::Cast(indicesResult, op::DataType::DT_INT64, executor);
354+ CHECK_RET(indicesResultCast != nullptr, ACLNN_ERR_INNER_NULLPTR);
355+ 
356+ auto resShapeVector = op::ToShapeVector(outputOut->GetViewShape());
357+ aclIntArray* resShapeArray =
358+ executor->AllocIntArray(resShapeVector.data(), resShapeVector.size());
359+ CHECK_RET(resShapeArray != nullptr, ACLNN_ERR_INNER_NULLPTR);
360+ auto resTranReshapeOut = l0op::Reshape(outResult, resShapeArray, executor);
361+ CHECK_RET(resTranReshapeOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
362+ auto viewCopyResultOut = l0op::ViewCopy(resTranReshapeOut, outputOut, executor);
363+ CHECK_RET(viewCopyResultOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
364+ 
365+ auto resTranReshapeIndices = l0op::Reshape(indicesResultCast, resShapeArray, executor);
366+ CHECK_RET(resTranReshapeIndices != nullptr, ACLNN_ERR_INNER_NULLPTR);
367+ auto viewCopyResultIndices = l0op::ViewCopy(resTranReshapeIndices, indicesOut, executor);
368+ CHECK_RET(viewCopyResultIndices != nullptr, ACLNN_ERR_INNER_NULLPTR);
369+
370+ return ACLNN_SUCCESS;
371+}
372+ 
373+// NHWC->NCHW
374+static aclnnStatus ConvertNHWCtoNCHW(
375+ const aclTensor*& inputContiguous,
376+ const aclTensor* self,
377+ aclOpExecutor* executor)
378+{
379+ if (self->GetViewFormat() == op::Format::FORMAT_NHWC) {
380+ std::vector<int64_t> valuePerm{INDEX_DIM0, INDEX_DIM3, INDEX_DIM1, INDEX_DIM2};
381+ auto perm = executor->AllocIntArray(valuePerm.data(), NCHW_DIM_NUM);
382+ CHECK_RET(perm != nullptr, ACLNN_ERR_INNER_NULLPTR);
383+ inputContiguous = l0op::Transpose(inputContiguous, perm, executor);
384+ CHECK_RET(inputContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
385+ }
386+ return ACLNN_SUCCESS;
387+}
388+ 
389+static aclnnStatus Handle910DCase(
390+ const aclTensor* self, const aclIntArray* outputSize, aclTensor* outputOut, aclTensor* indicesOut,
391+ aclOpExecutor* executor)
392+{
393+ auto inputContiguous = l0op::Contiguous(self, executor);
394+ CHECK_RET(inputContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
395+ 
396+ // NHWC -> NCHW
397+ auto ret = ConvertNHWCtoNCHW(inputContiguous, self, executor);
398+ if (ret != ACLNN_SUCCESS) {
399+ return ret;
400+ }
401+ 
402+ // reshape NCHW/NCL -> NCDHW
403+ op::Shape inputContiguousShape = inputContiguous->GetViewShape();
404+ int64_t inputDimNum = static_cast<int64_t>(inputContiguousShape.GetDimNum());
405+ int64_t hiValue = inputDimNum == 3 ? inputContiguous->GetViewShape()[1] : inputContiguous->GetViewShape()[2];
406+ int64_t wiValue = inputDimNum == 3 ? inputContiguous->GetViewShape()[2] : inputContiguous->GetViewShape()[3];
407+ 
408+ int64_t hValueRemainder = hiValue % (*outputSize)[0];
409+ int64_t wValueRemainder = wiValue % (*outputSize)[1];
410+ bool ifTranMaxPool3D = hValueRemainder == 0 && wValueRemainder == 0;
411+ const aclTensor* inputContiguousReshape = inputContiguous;
412+ if (inputDimNum == 3 && !ifTranMaxPool3D) { // 转ada 2d
413+ inputContiguousReshape = View3Das4D(inputContiguous, executor);
414+ } else if (ifTranMaxPool3D) { // 转max3d
415+ inputContiguousReshape = inputDimNum == 3 ? View3Das5D(inputContiguous, executor) :
416+ View4Das5D(inputContiguous, executor);
417+ }
418+ if (ifTranMaxPool3D) {
419+ int64_t kernelDSize = 1;
420+ int64_t kernelHSize = inputContiguousReshape->GetViewShape()[INDEX_DIM3] / (*outputSize)[INDEX_DIM0];
421+ int64_t kernelWSize = inputContiguousReshape->GetViewShape()[INDEX_DIM4] / (*outputSize)[INDEX_DIM1];
422+ std::vector<int64_t> kernelSizeArr = {kernelDSize, kernelHSize, kernelWSize};
423+ ret = HandleMaxPool3DCase(inputContiguousReshape, kernelSizeArr, outputOut, indicesOut,
424+ executor);
425+ if (ret != ACLNN_SUCCESS) {
426+ return ret;
427+ }
428+ } else {
429+ auto [outResult, indicesResult] =
430+ l0op::AdaptiveMaxPool2d(inputContiguousReshape, outputSize, executor);
431+ CHECK_RET(outResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
432+ CHECK_RET(indicesResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
433+
434+ ret = ProcessAdaptiveMaxPoolResults(outResult, indicesResult, outputOut, indicesOut, executor);
435+ if (ret != ACLNN_SUCCESS) {
436+ return ret;
437+ }
438+ }
439+ return ACLNN_SUCCESS;
440+}
441+ 
442+static aclnnStatus Handle910BCase(
443+ const aclTensor* self, const aclIntArray* outputSize, aclTensor* outputOut, aclTensor* indicesOut,
444+ aclOpExecutor* executor)
445+{
446+ // 将2d参数转换为3d可以使用的参数
447+ int64_t newOutputSizeData[] = {1, (*outputSize)[0], (*outputSize)[1]};
448+ aclIntArray* newOutputSize = executor->AllocIntArray(newOutputSizeData, 3);
449+ 
450+ auto inputContiguous = l0op::Contiguous(self, executor);
451+ CHECK_RET(inputContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
452+ 
453+ // NHWC -> NCHW
454+ auto ret = ConvertNHWCtoNCHW(inputContiguous, self, executor);
455+ if (ret != ACLNN_SUCCESS) {
456+ return ret;
457+ }
458+ 
459+ // reshape NCHW/NCL -> NCDHW
460+ op::Shape inputContiguousShape = inputContiguous->GetViewShape();
461+ int64_t inputDimNum = static_cast<int64_t>(inputContiguousShape.GetDimNum());
462+ std::vector<int64_t> valueShape(NCDHW_DIM_NUM);
463+ valueShape[0] = inputDimNum == static_cast<int64_t>(NCHW_DIM_NUM) ? inputContiguousShape.GetDim(0) : 1;
464+ valueShape[1] = inputDimNum == static_cast<int64_t>(NCHW_DIM_NUM) ? inputContiguousShape.GetDim(1) :
465+ inputContiguousShape.GetDim(0);
466+ valueShape[DIM_D] = 1;
467+ for (int64_t i = inputDimNum - static_cast<int64_t>(OUTPUT_SIZE_NUM); i < inputDimNum; i++) {
468+ valueShape[NCDHW_DIM_NUM - inputDimNum + i] = inputContiguousShape.GetDim(i);
469+ }
470+ auto reshapeShape = executor->AllocIntArray(valueShape.data(), NCDHW_DIM_NUM);
471+ CHECK_RET(reshapeShape != nullptr, ACLNN_ERR_INNER_NULLPTR);
472+ auto inputContiguousReshape = l0op::Reshape(inputContiguous, reshapeShape, executor);
473+ CHECK_RET(inputContiguousReshape != nullptr, ACLNN_ERR_INNER_NULLPTR);
474+ 
475+ int64_t dValueRemainder = inputContiguousReshape->GetViewShape()[INDEX_DIM2] % (*newOutputSize)[INDEX_DIM0];
476+ int64_t hValueRemainder = inputContiguousReshape->GetViewShape()[INDEX_DIM3] % (*newOutputSize)[INDEX_DIM1];
477+ int64_t wValueRemainder = inputContiguousReshape->GetViewShape()[INDEX_DIM4] % (*newOutputSize)[INDEX_DIM2];
478+ if (dValueRemainder == 0 && hValueRemainder == 0 && wValueRemainder == 0) {
479+ int64_t kernelDSize = inputContiguousReshape->GetViewShape()[INDEX_DIM2] / (*newOutputSize)[INDEX_DIM0];
480+ int64_t kernelHSize = inputContiguousReshape->GetViewShape()[INDEX_DIM3] / (*newOutputSize)[INDEX_DIM1];
481+ int64_t kernelWSize = inputContiguousReshape->GetViewShape()[INDEX_DIM4] / (*newOutputSize)[INDEX_DIM2];
482+ std::vector<int64_t> kernelSizeArr = {kernelDSize, kernelHSize, kernelWSize};
483+
484+ ret = HandleMaxPool3DCase(inputContiguousReshape, kernelSizeArr, outputOut, indicesOut, executor);
485+ if (ret != ACLNN_SUCCESS) {
486+ return ret;
487+ }
488+ } else {
489+ auto [outResult, indicesResult] =
490+ l0op::AdaptiveMaxPool3d(inputContiguousReshape, newOutputSize, executor);
491+ CHECK_RET(outResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
492+ CHECK_RET(indicesResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
493+
494+ ret = ProcessAdaptiveMaxPoolResults(outResult, indicesResult, outputOut, indicesOut, executor);
495+ if (ret != ACLNN_SUCCESS) {
496+ return ret;
497+ }
498+ }
499+ return ACLNN_SUCCESS;
500+}
501+ 
502+static aclnnStatus HandleGeneralCase(
503+ const aclTensor* self, const aclIntArray* outputSize, aclTensor* outputOut, aclTensor* indicesOut,
504+ aclOpExecutor* executor)
505+{
506+ auto selfContiguous = l0op::Contiguous(self, executor);
507+ CHECK_RET(selfContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
508+ 
509+ const aclTensor* selfNewFormat = selfContiguous;
510+ if (self->GetViewShape().GetDimNum() != NCHW_DIM_NUM) {
511+ auto selfNd = l0op::UnsqueezeNd(selfContiguous, AXIS_DIM, executor);
512+ CHECK_RET(selfNd != nullptr, ACLNN_ERR_INNER_NULLPTR);
513+ 
514+ selfNewFormat = l0op::ReFormat(selfNd, static_cast<op::Format>(ACL_FORMAT_NCHW));
515+ CHECK_RET(selfNewFormat != nullptr, ACLNN_ERR_INNER_NULLPTR);
516+ }
517+ 
518+ CHECK_RET(CheckInputFormat(selfNewFormat), ACLNN_ERR_PARAM_INVALID);
519+ 
520+ auto result = l0op::AdaptiveMaxPool2d(selfNewFormat, outputSize, executor);
521+ const aclTensor* outputRst = std::get<0>(result);
522+ const aclTensor* indicesRst = std::get<1>(result);
523+ 
524+ CHECK_RET(outputRst != nullptr && indicesRst != nullptr, ACLNN_ERR_INNER_NULLPTR);
525+ 
526+ auto outputNewFormat = outputRst;
527+ auto indicesNewFormat = indicesRst;
528+ if (self->GetViewShape().GetDimNum() != NCHW_DIM_NUM) {
529+ auto outputNewShape = l0op::SqueezeNd(outputRst, AXIS_DIM, executor);
530+ CHECK_RET(outputNewShape != nullptr, ACLNN_ERR_INNER_NULLPTR);
531+ 
532+ auto indicesNewShape = l0op::SqueezeNd(indicesRst, AXIS_DIM, executor);
533+ CHECK_RET(indicesNewShape != nullptr, ACLNN_ERR_INNER_NULLPTR);
534+ 
535+ outputNewFormat = l0op::ReFormat(outputNewShape, outputOut->GetViewFormat());
536+ CHECK_RET(outputNewFormat != nullptr, ACLNN_ERR_INNER_NULLPTR);
537+ 
538+ indicesNewFormat = l0op::ReFormat(indicesNewShape, indicesOut->GetViewFormat());
539+ CHECK_RET(indicesNewFormat != nullptr, ACLNN_ERR_INNER_NULLPTR);
540+ }
541+ // check output shape
542+ CHECK_RET(CheckReduceOutShape(outputNewFormat, outputOut), ACLNN_ERR_PARAM_INVALID);
543+ CHECK_RET(CheckReduceOutShape(indicesNewFormat, indicesOut), ACLNN_ERR_PARAM_INVALID);
544+ auto viewCopyOutputResult = l0op::ViewCopy(outputNewFormat, outputOut, executor);
545+ CHECK_RET(viewCopyOutputResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
546+ 
547+ auto viewCopyIndicesResult = l0op::ViewCopy(indicesNewFormat, indicesOut, executor);
548+ CHECK_RET(viewCopyIndicesResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
549+ return ACLNN_SUCCESS;
550+}
551+ 
190aclnnStatus aclnnAdaptiveMaxPool2dGetWorkspaceSize(552aclnnStatus aclnnAdaptiveMaxPool2dGetWorkspaceSize(
191 const aclTensor* self, const aclIntArray* outputSize, aclTensor* outputOut, aclTensor* indicesOut,553 const aclTensor* self, const aclIntArray* outputSize, aclTensor* outputOut, aclTensor* indicesOut,
192 uint64_t* workspaceSize, aclOpExecutor** executor)554 uint64_t* workspaceSize, aclOpExecutor** executor)
193{555{
194 L2_DFX_PHASE_1(aclnnAdaptiveMaxPool2d, DFX_IN(self, outputSize), DFX_OUT(outputOut, indicesOut));556 L2_DFX_PHASE_1(aclnnAdaptiveMaxPool2d, DFX_IN(self, outputSize), DFX_OUT(outputOut, indicesOut));
195- 
196 auto uniqueExecutor = CREATE_EXECUTOR();557 auto uniqueExecutor = CREATE_EXECUTOR();
197 CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR);558 CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR);
198 559 
@@ -205,171 +566,17 @@ aclnnStatus aclnnAdaptiveMaxPool2dGetWorkspaceSize(
205 return ACLNN_SUCCESS;566 return ACLNN_SUCCESS;
206 }567 }
207 568 
208- if (IsSocVersion910B() && !(IsSelfDtypeDouble(self))) {569+ if (IsSocVersion910D() && !(IsSelfDtypeDouble(self))) {
209- // 将2d参数转换为3d可以使用的参数570+ CHECK_RET(
210- int64_t newOutputSizeData[] = {1, (*outputSize)[0], (*outputSize)[1]};571+ Handle910DCase(self, outputSize, outputOut, indicesOut, uniqueExecutor.get()) == ACLNN_SUCCESS,
211- aclIntArray* newOutputSize = uniqueExecutor.get()->AllocIntArray(newOutputSizeData, 3);572+ ACLNN_ERR_INNER_NULLPTR);
212- 573+ } else if (IsSocVersion910B() && !(IsSelfDtypeDouble(self))) {
213- auto inputContiguous = l0op::Contiguous(self, uniqueExecutor.get());574+ CHECK_RET(Handle910BCase(self, outputSize, outputOut, indicesOut, uniqueExecutor.get()) == ACLNN_SUCCESS,
214- CHECK_RET(inputContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);575+ ACLNN_ERR_INNER_NULLPTR);
215- 576+ } else {
216- // NHWC -> NCHW577+ CHECK_RET(HandleGeneralCase(self, outputSize, outputOut, indicesOut, uniqueExecutor.get()) == ACLNN_SUCCESS,
217- if (self->GetViewFormat() == op::Format::FORMAT_NHWC) {578+ ACLNN_ERR_INNER_NULLPTR);
218- std::vector<int64_t> valuePerm{INDEX_DIM0, INDEX_DIM3, INDEX_DIM1, INDEX_DIM2};
219- auto perm = uniqueExecutor.get()->AllocIntArray(valuePerm.data(), NCHW_DIM_NUM);
220- CHECK_RET(perm != nullptr, ACLNN_ERR_INNER_NULLPTR);
221- inputContiguous = l0op::Transpose(inputContiguous, perm, uniqueExecutor.get());
222- CHECK_RET(inputContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
223- }
224- 
225- // reshape NCHW/NCL -> NCDHW
226- op::Shape inputContiguousShape = inputContiguous->GetViewShape();
227- int64_t inputDimNum = static_cast<int64_t>(inputContiguousShape.GetDimNum());
228- std::vector<int64_t> valueShape(NCDHW_DIM_NUM);
229- valueShape[0] = inputDimNum == static_cast<int64_t>(NCHW_DIM_NUM) ? inputContiguousShape.GetDim(0) : 1;
230- valueShape[1] = inputDimNum == static_cast<int64_t>(NCHW_DIM_NUM) ? inputContiguousShape.GetDim(1) :
231- inputContiguousShape.GetDim(0);
232- valueShape[DIM_D] = 1;
233- for (int64_t i = inputDimNum - static_cast<int64_t>(OUTPUT_SIZE_NUM); i < inputDimNum; i++) {
234- valueShape[NCDHW_DIM_NUM - inputDimNum + i] = inputContiguousShape.GetDim(i);
235- }
236- auto reshapeShape = uniqueExecutor.get()->AllocIntArray(valueShape.data(), NCDHW_DIM_NUM);
237- CHECK_RET(reshapeShape != nullptr, ACLNN_ERR_INNER_NULLPTR);
238- auto inputContiguousReshape = l0op::Reshape(inputContiguous, reshapeShape, uniqueExecutor.get());
239- CHECK_RET(inputContiguousReshape != nullptr, ACLNN_ERR_INNER_NULLPTR);
240- 
241- int64_t dValueRemainder = inputContiguousReshape->GetViewShape()[INDEX_DIM2] % (*newOutputSize)[INDEX_DIM0];
242- int64_t hValueRemainder = inputContiguousReshape->GetViewShape()[INDEX_DIM3] % (*newOutputSize)[INDEX_DIM1];
243- int64_t wValueRemainder = inputContiguousReshape->GetViewShape()[INDEX_DIM4] % (*newOutputSize)[INDEX_DIM2];
244- if (dValueRemainder == 0 && hValueRemainder == 0 && wValueRemainder == 0) {
245- size_t kernelDSize = inputContiguousReshape->GetViewShape()[INDEX_DIM2] / (*newOutputSize)[INDEX_DIM0];
246- size_t kernelHSize = inputContiguousReshape->GetViewShape()[INDEX_DIM3] / (*newOutputSize)[INDEX_DIM1];
247- size_t kernelWSize = inputContiguousReshape->GetViewShape()[INDEX_DIM4] / (*newOutputSize)[INDEX_DIM2];
248- std::vector<int64_t> kernelSizeArr = {};
249- kernelSizeArr.push_back(kernelDSize);
250- kernelSizeArr.push_back(kernelHSize);
251- kernelSizeArr.push_back(kernelWSize);
252- const aclIntArray* kernelSize = aclCreateIntArray(kernelSizeArr.data(), kernelSizeArr.size());
253- const aclIntArray* stride = aclCreateIntArray(kernelSizeArr.data(), kernelSizeArr.size());
254- std::vector<int64_t> paddingArr = {0, 0, 0};
255- const aclIntArray* padding = aclCreateIntArray(paddingArr.data(), paddingArr.size());
256- std::vector<int64_t> dilationArr = {1, 1, 1};
257- const aclIntArray* dilation = aclCreateIntArray(dilationArr.data(), dilationArr.size());
258- bool ceilMode = false;
259- std::string dataFormat = "NCDHW";
260- 
261- inputContiguousReshape =
262- l0op::ReFormat(inputContiguousReshape, op::Format::FORMAT_NCDHW, uniqueExecutor.get());
263- CHECK_RET(inputContiguousReshape != nullptr, ACLNN_ERR_INNER_NULLPTR);
264- 
265- auto [outResult, indicesResult] = l0op::MaxPool3DWithArgmaxV2Ncdhw(
266- inputContiguousReshape, kernelSize, stride, padding, dilation, ceilMode, dataFormat,
267- uniqueExecutor.get());
268- CHECK_RET(outResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
269- CHECK_RET(indicesResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
270- auto indicesResultCast = l0op::Cast(indicesResult, op::DataType::DT_INT64, uniqueExecutor.get());
271- CHECK_RET(indicesResultCast != nullptr, ACLNN_ERR_INNER_NULLPTR);
272- 
273- auto resShapeVector = op::ToShapeVector(outputOut->GetViewShape());
274- aclIntArray* resShapeArray =
275- uniqueExecutor.get()->AllocIntArray(resShapeVector.data(), resShapeVector.size());
276- CHECK_RET(resShapeArray != nullptr, ACLNN_ERR_INNER_NULLPTR);
277- auto resTranReshapeOut = l0op::Reshape(outResult, resShapeArray, uniqueExecutor.get());
278- CHECK_RET(resTranReshapeOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
279- 
280- resTranReshapeOut = l0op::ReFormat(resTranReshapeOut, outputOut->GetStorageFormat(), uniqueExecutor.get());
281- CHECK_RET(resTranReshapeOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
282- 
283- auto viewCopyResultOut = l0op::ViewCopy(resTranReshapeOut, outputOut, uniqueExecutor.get());
284- CHECK_RET(viewCopyResultOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
285- 
286- auto resTranReshapeIndices = l0op::Reshape(indicesResultCast, resShapeArray, uniqueExecutor.get());
287- CHECK_RET(resTranReshapeIndices != nullptr, ACLNN_ERR_INNER_NULLPTR);
288- 
289- resTranReshapeIndices =
290- l0op::ReFormat(resTranReshapeIndices, indicesOut->GetStorageFormat(), uniqueExecutor.get());
291- CHECK_RET(resTranReshapeIndices != nullptr, ACLNN_ERR_INNER_NULLPTR);
292- 
293- auto viewCopyResultIndices = l0op::ViewCopy(resTranReshapeIndices, indicesOut, uniqueExecutor.get());
294- CHECK_RET(viewCopyResultIndices != nullptr, ACLNN_ERR_INNER_NULLPTR);
295- // 释放资源
296- aclDestroyIntArray(kernelSize);
297- aclDestroyIntArray(stride);
298- aclDestroyIntArray(padding);
299- aclDestroyIntArray(dilation);
300- } else {
301- auto [outResult, indicesResult] =
302- l0op::AdaptiveMaxPool3d(inputContiguousReshape, newOutputSize, uniqueExecutor.get());
303- CHECK_RET(outResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
304- CHECK_RET(indicesResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
305- auto indicesResultCast = l0op::Cast(indicesResult, op::DataType::DT_INT64, uniqueExecutor.get());
306- CHECK_RET(indicesResultCast != nullptr, ACLNN_ERR_INNER_NULLPTR);
307- 
308- auto resShapeVector = op::ToShapeVector(outputOut->GetViewShape());
309- aclIntArray* resShapeArray =
310- uniqueExecutor.get()->AllocIntArray(resShapeVector.data(), resShapeVector.size());
311- CHECK_RET(resShapeArray != nullptr, ACLNN_ERR_INNER_NULLPTR);
312- auto resTranReshapeOut = l0op::Reshape(outResult, resShapeArray, uniqueExecutor.get());
313- CHECK_RET(resTranReshapeOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
314- auto viewCopyResultOut = l0op::ViewCopy(resTranReshapeOut, outputOut, uniqueExecutor.get());
315- CHECK_RET(viewCopyResultOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
316- 
317- auto resTranReshapeIndices = l0op::Reshape(indicesResultCast, resShapeArray, uniqueExecutor.get());
318- CHECK_RET(resTranReshapeIndices != nullptr, ACLNN_ERR_INNER_NULLPTR);
319- auto viewCopyResultIndices = l0op::ViewCopy(resTranReshapeIndices, indicesOut, uniqueExecutor.get());
320- CHECK_RET(viewCopyResultIndices != nullptr, ACLNN_ERR_INNER_NULLPTR);
321- }
322- 
323- // 固定写法，获取计算过程中需要使用的workspace大小
324- *workspaceSize = uniqueExecutor->GetWorkspaceSize();
325- uniqueExecutor.ReleaseTo(executor);
326- return ACLNN_SUCCESS;
327 }579 }
328- auto selfContiguous = l0op::Contiguous(self, uniqueExecutor.get());
329- CHECK_RET(selfContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
330- 
331- const aclTensor* selfNewFormat = selfContiguous;
332- if (self->GetViewShape().GetDimNum() != NCHW_DIM_NUM) {
333- auto selfNd = l0op::UnsqueezeNd(selfContiguous, AXIS_DIM, uniqueExecutor.get());
334- CHECK_RET(selfNd != nullptr, ACLNN_ERR_INNER_NULLPTR);
335- 
336- selfNewFormat = l0op::ReFormat(selfNd, static_cast<op::Format>(ACL_FORMAT_NCHW));
337- CHECK_RET(selfNewFormat != nullptr, ACLNN_ERR_INNER_NULLPTR);
338- }
339- 
340- CHECK_RET(CheckInputFormat(selfNewFormat), ACLNN_ERR_PARAM_INVALID);
341- 
342- auto result = l0op::AdaptiveMaxPool2d(selfNewFormat, outputSize, uniqueExecutor.get());
343- const aclTensor* outputRst = std::get<0>(result);
344- const aclTensor* indicesRst = std::get<1>(result);
345- 
346- CHECK_RET(outputRst != nullptr && indicesRst != nullptr, ACLNN_ERR_INNER_NULLPTR);
347- 
348- auto outputNewFormat = outputRst;
349- auto indicesNewFormat = indicesRst;
350- if (self->GetViewShape().GetDimNum() != NCHW_DIM_NUM) {
351- auto outputNewShape = l0op::SqueezeNd(outputRst, AXIS_DIM, uniqueExecutor.get());
352- CHECK_RET(outputNewShape != nullptr, ACLNN_ERR_INNER_NULLPTR);
353- 
354- auto indicesNewShape = l0op::SqueezeNd(indicesRst, AXIS_DIM, uniqueExecutor.get());
355- CHECK_RET(indicesNewShape != nullptr, ACLNN_ERR_INNER_NULLPTR);
356- 
357- outputNewFormat = l0op::ReFormat(outputNewShape, outputOut->GetViewFormat());
358- CHECK_RET(outputNewFormat != nullptr, ACLNN_ERR_INNER_NULLPTR);
359- 
360- indicesNewFormat = l0op::ReFormat(indicesNewShape, indicesOut->GetViewFormat());
361- CHECK_RET(indicesNewFormat != nullptr, ACLNN_ERR_INNER_NULLPTR);
362- }
363- // check output shape
364- CHECK_RET(CheckReduceOutShape(outputNewFormat, outputOut), ACLNN_ERR_PARAM_INVALID);
365- CHECK_RET(CheckReduceOutShape(indicesNewFormat, indicesOut), ACLNN_ERR_PARAM_INVALID);
366- auto viewCopyOutputResult = l0op::ViewCopy(outputNewFormat, outputOut, uniqueExecutor.get());
367- CHECK_RET(viewCopyOutputResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
368- 
369- auto viewCopyIndicesResult = l0op::ViewCopy(indicesNewFormat, indicesOut, uniqueExecutor.get());
370- CHECK_RET(viewCopyIndicesResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
371- 
372- // 固定写法，获取计算过程中需要使用的workspace大小
373 *workspaceSize = uniqueExecutor->GetWorkspaceSize();580 *workspaceSize = uniqueExecutor->GetWorkspaceSize();
374 uniqueExecutor.ReleaseTo(executor);581 uniqueExecutor.ReleaseTo(executor);
375 return ACLNN_SUCCESS;582 return ACLNN_SUCCESS;
Rpooling/adaptive_max_pool3d/op_api/aclnn_adaptive_max_pool2d.h→pooling/adaptive_max_pool2d/op_api/aclnn_adaptive_max_pool2d.h+0-0
文件重命名但无更改。
Rpooling/adaptive_max_pool3d/op_api/adaptive_max_pool2d.cpp→pooling/adaptive_max_pool2d/op_api/adaptive_max_pool2d.cpp+41-4
@@ -15,6 +15,7 @@
15#include "opdev/op_dfx.h"15#include "opdev/op_dfx.h"
16#include "opdev/op_log.h"16#include "opdev/op_log.h"
17#include "opdev/shape_utils.h"17#include "opdev/shape_utils.h"
18+#include "op_api/aclnn_util.h"
18 19 
19using namespace op;20using namespace op;
20namespace l0op {21namespace l0op {
@@ -25,6 +26,17 @@ static constexpr size_t DIM_W = 1;
25static constexpr size_t NHWC_DIM_H = 3;26static constexpr size_t NHWC_DIM_H = 3;
26static constexpr size_t NHWC_DIM_W = 2;27static constexpr size_t NHWC_DIM_W = 2;
27 28 
29+static const std::initializer_list<op::DataType> AICORE_DTYPE_SUPPORT_LIST_950 = {
30+ op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16, op::DataType::DT_BF16};
31+ 
32+static bool IsAscend950iCoreSupport(const aclTensor* self)
33+{
34+ if (!Ops::NN::AclnnUtil::IsRegbase()) {
35+ return false;
36+ }
37+ return CheckType(self->GetDataType(), AICORE_DTYPE_SUPPORT_LIST_950);
38+}
39+ 
28std::tuple<aclTensor*, aclTensor*> AdapativeMaxPool2dAiCpu(40std::tuple<aclTensor*, aclTensor*> AdapativeMaxPool2dAiCpu(
29 const aclTensor* self, const aclIntArray* outputSize, aclTensor* outputOut, aclTensor* indicesOut,41 const aclTensor* self, const aclIntArray* outputSize, aclTensor* outputOut, aclTensor* indicesOut,
30 aclOpExecutor* executor)42 aclOpExecutor* executor)
@@ -41,6 +53,17 @@ std::tuple<aclTensor*, aclTensor*> AdapativeMaxPool2dAiCpu(
41 return std::tuple<aclTensor*, aclTensor*>(outputOut, indicesOut);53 return std::tuple<aclTensor*, aclTensor*>(outputOut, indicesOut);
42}54}
43 55 
56+std::tuple<aclTensor*, aclTensor*> AdapativeMaxPool2dAiCore(
57+ const aclTensor* self, const aclIntArray* outputSize, aclTensor* outputOut, aclTensor* indicesOut,
58+ aclOpExecutor* executor)
59+{
60+ L0_DFX(AdapativeMaxPool2dAiCore, self, outputSize, outputOut, indicesOut);
61+ ADD_TO_LAUNCHER_LIST_AICORE(
62+ AdaptiveMaxPool2d, OP_INPUT(self), OP_OUTPUT(outputOut, indicesOut), OP_ATTR(outputSize));
63+ return std::tuple<aclTensor*, aclTensor*>(outputOut, indicesOut);
64+}
65+ 
66+ 
44std::tuple<aclTensor*, aclTensor*> AdaptiveMaxPool2d(67std::tuple<aclTensor*, aclTensor*> AdaptiveMaxPool2d(
45 const aclTensor* self, const aclIntArray* outputSize, aclOpExecutor* executor)68 const aclTensor* self, const aclIntArray* outputSize, aclOpExecutor* executor)
46{69{
@@ -58,11 +81,25 @@ std::tuple<aclTensor*, aclTensor*> AdaptiveMaxPool2d(
58 }81 }
59 82 
60 auto outputOut = executor->AllocTensor(outShape, self->GetDataType(), self->GetStorageFormat());83 auto outputOut = executor->AllocTensor(outShape, self->GetDataType(), self->GetStorageFormat());
61- auto indicesOut = executor->AllocTensor(outShape, op::DataType::DT_INT64, self->GetStorageFormat());84+
62- if (outputOut == nullptr || indicesOut == nullptr) {85+ if (outputOut == nullptr) {
63- OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "outputOut or indicesOut is nullptr.");86+ OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "outputOut is nullptr.");
64 return std::tuple<aclTensor*, aclTensor*>(nullptr, nullptr);87 return std::tuple<aclTensor*, aclTensor*>(nullptr, nullptr);
65 }88 }
66- return AdapativeMaxPool2dAiCpu(self, outputSize, outputOut, indicesOut, executor);89+ if (IsAscend950iCoreSupport(self)) {
90+ auto indicesInt32Out = executor->AllocTensor(outShape, op::DataType::DT_INT32, self->GetStorageFormat());
91+ if (indicesInt32Out == nullptr) {
92+ OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "indicesInt32Out is nullptr.");
93+ return std::tuple<aclTensor*, aclTensor*>(nullptr, nullptr);
94+ }
95+ OP_LOGI("beigin adaptive maxpool2d.\n");
CANN-robot
CANN-robotCANN-robot2月9日

代码结构与可维护性: 日志信息中存在拼写错误：'beigin' 应为 'begin'。虽然不影响功能，但降低了代码的专业性和可读性。

问题类型: 代码结构与可维护性 文件路径: pooling/adaptive_max_pool2d/op_api/adaptive_max_pool2d.cpp 行号: 96 问题代码:

OP_LOGI("beigin adaptive maxpool2d.\n");

修改建议:

将日志信息更正为："begin adaptive maxpool2d."。

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

likedislike
96+ return AdapativeMaxPool2dAiCore(self, outputSize, outputOut, indicesInt32Out, executor);
97+ }
98+ auto indicesInt64Out = executor->AllocTensor(outShape, op::DataType::DT_INT64, self->GetStorageFormat());
99+ if (indicesInt64Out == nullptr) {
100+ OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "indicesInt64Out is nullptr.");
101+ return std::tuple<aclTensor*, aclTensor*>(nullptr, nullptr);
102+ }
103+ return AdapativeMaxPool2dAiCpu(self, outputSize, outputOut, indicesInt64Out, executor);
67}104}
68} // namespace l0op105} // namespace l0op
Rpooling/adaptive_max_pool3d/op_api/adaptive_max_pool2d.h→pooling/adaptive_max_pool2d/op_api/adaptive_max_pool2d.h+0-0
文件重命名但无更改。
Apooling/adaptive_max_pool2d/op_host/adaptive_max_pool2d_def.cpp+51-0
@@ -0,0 +1,51 @@
1+/**
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 of
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.
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.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file adaptive_max_pool2d_def.cpp
13+ * \brief
14+ */
15+ 
16+#include "register/op_def_registry.h"
17+constexpr int DTYPE_INT32 = 3;
18+namespace ops {
19+class AdaptiveMaxPool2d : public OpDef
20+{
21+public:
22+ explicit AdaptiveMaxPool2d(const char* name) : OpDef(name)
23+ {
24+ this->Input("x")
25+ .ParamType(REQUIRED)
26+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16})
27+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
28+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
29+ this->Output("y")
30+ .ParamType(REQUIRED)
31+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16})
32+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
33+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
34+ this->Output("indices")
35+ .ParamType(REQUIRED)
36+ .DataType({ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64})
37+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
38+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
39+ this->Attr("output_size").AttrType(REQUIRED).ListInt();
40+ this->Attr("indices_dtype").AttrType(OPTIONAL).Int(DTYPE_INT32);
41+ 
42+ OpAICoreConfig aicore_config;
43+ aicore_config.DynamicCompileStaticFlag(true)
44+ .DynamicRankSupportFlag(true)
45+ .DynamicShapeSupportFlag(true);
46+ this->AICore().AddConfig("ascend950");
47+ }
48+};
49+ 
50+OP_ADD(AdaptiveMaxPool2d);
51+} // namespace ops
Apooling/adaptive_max_pool2d/op_host/adaptive_max_pool2d_simt_tiling.h+81-0
@@ -0,0 +1,81 @@
1+/**
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 of
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.
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.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file adaptive_max_pool2d_simt_tiling.h
13+ * \brief simt imply for adaptive_max_pool2d
14+ */
15+ 
16+#ifndef AIR_CXX_RUNTIME_V2_OP_IMPL_ADAPTIVE_MAX_POOL2D_SIMT_TILING_H
17+#define AIR_CXX_RUNTIME_V2_OP_IMPL_ADAPTIVE_MAX_POOL2D_SIMT_TILING_H
18+ 
19+#include "adaptive_max_pool2d_tiling_base.h"
20+ 
21+namespace optiling
22+{
23+const int64_t NCHW_DIMS = 4;
24+const int64_t N_DIM_ = 0;
25+const int64_t C_DIM_ = 1;
26+const int64_t H_DIM_ = 2;
27+const int64_t W_DIM_ = 3;
28+const int64_t H_IDX_ = 0;
29+const int64_t W_IDX_ = 1;
30+const int64_t DOUB = 2;
31+const int64_t FIRPOS = 0;
32+const int64_t SECPOS = 1;
33+constexpr int64_t MAX_INT32 = 2147483647;
34+constexpr int64_t MAX_THREAD_NUM = 256;
35+ 
36+struct InputSIMTInfo {
37+ array<uint64_t, NCHW_DIMS> inputShape;
38+ array<uint64_t, NCHW_DIMS> outShape;
39+ int64_t kernelHMax;
40+ int64_t kernelWMax;
41+};
42+ 
43+class AdaMaxPool2dTilingSIMT : public AdaMaxPool2dBaseTiling
44+{
45+public:
46+ explicit AdaMaxPool2dTilingSIMT(gert::TilingContext* context) : AdaMaxPool2dBaseTiling(context)
47+ {
48+ }
49+ 
50+ ~AdaMaxPool2dTilingSIMT() override
51+ {
52+ }
53+ 
54+protected:
55+ bool IsCapable() override;
56+ ge::graphStatus GetPlatformInfo() override;
57+ // 获取INPUT/OUTPUT/ATTR信息
58+ ge::graphStatus GetShapeAttrsInfo() override;
59+ // 计算数据切分TilingData
60+ ge::graphStatus DoOpTiling() override;
61+ // 计算TilingKey
62+ uint64_t GetTilingKey() const;
63+ // 分配workspace
64+ ge::graphStatus GetWorkspaceSize() override;
65+ // 保存Tiling数据
66+ ge::graphStatus PostTiling() override;
67+ // tiling信息打屏
68+ void DumpTilingInfo() override;
69+ 
70+private:
71+ uint64_t GenerateTilingKey(uint64_t innerKey);
72+ ge::graphStatus CheckPlatformAndGetShapes();
73+ ge::graphStatus CheckDataTypeAndAttrs();
74+ AdaptiveMaxPool2dTilingData tiling;
75+ InputSIMTInfo inputData;
76+ uint64_t coreNum_ = 1;
77+ uint64_t ubSize_ = 0;
78+};
79+ 
80+} // namespace optiling
CANN-robotCANN-robot
CANN-robotCANN-robot2月9日

代码结构与可维护性: 文件末尾缺少换行符（根据diff标记显示'\ No newline at end of file'）。虽然这不是功能性问题，但违反了许多代码规范，可能导致某些工具处理时出现问题。

问题类型: 代码结构与可维护性 文件路径: pooling/adaptive_max_pool2d/op_host/adaptive_max_pool2d_simt_tiling.h 行号: 91 问题代码:

}  // namespace optiling
#endif  // MAX_POOL2D_WITH_ARGMAX_V2_SIMT_TILING_H

修改建议:

在文件末尾添加一个空行，确保以换行符结束。

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

likedislike
CANN-robotCANN-robot2月9日

文件结尾格式问题: 文件结尾缺少换行符，这不符合POSIX标准，可能导致某些工具（如cat、wc等）处理时出现问题。虽然现代编译器通常能处理这种情况，但为了代码规范性和可移植性，应该在文件末尾添加换行符。

问题类型: 文件结尾格式问题 文件路径: pooling/adaptive_max_pool2d/op_host/adaptive_max_pool2d_simt_tiling.h 行号: 91 问题代码:

}  // namespace optiling
#endif  // MAX_POOL2D_WITH_ARGMAX_V2_SIMT_TILING_H

修改建议:

在文件末尾添加一个空行，确保以换行符结束。这可以通过在最后一行后添加一个空行来实现。

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

likedislike
81+#endif // AIR_CXX_RUNTIME_V2_OP_IMPL_ADAPTIVE_MAX_POOL2D_SIMT_TILING_H
Apooling/adaptive_max_pool2d/op_host/adaptive_max_pool2d_tiling.cpp+158-0
@@ -0,0 +1,158 @@
1+/**
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 of
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.
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.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file adaptive_max_pool2d_tiling.cpp
13+ * \brief
14+ */
15+ 
16+#include "adaptive_max_pool2d_tiling_base.h"
17+ 
18+using Ops::NN::Optiling::TilingRegistry;
19+ 
20+namespace optiling {
21+constexpr uint64_t NCHW_DIM_N = 0;
22+constexpr uint64_t NCHW_DIM_C = 1;
23+constexpr uint64_t NCHW_DIM_H = 2;
24+constexpr uint64_t NCHW_DIM_W = 3;
25+constexpr uint64_t OUTPUTSIZE_DIM_MAX = 2;
26+constexpr uint64_t DIM_NUM_FOUR = 4;
27+static const gert::Shape g_vec_1_shape = {1};
28+ 
29+static const gert::Shape &EnsureNotScalar(const gert::Shape &inShape) {
30+ if (inShape.IsScalar()) {
31+ return g_vec_1_shape;
32+ }
33+ return inShape;
34+}
35+ 
36+bool AdaMaxPool2dBaseTiling::IsCapable()
37+{
38+ return true;
39+}
40+ 
41+ge::graphStatus AdaMaxPool2dBaseTiling::DoOpTiling()
42+{
43+ return ge::GRAPH_SUCCESS;
44+}
45+ 
46+ge::graphStatus AdaMaxPool2dBaseTiling::DoLibApiTiling()
47+{
48+ return ge::GRAPH_SUCCESS;
49+}
50+ 
51+uint64_t AdaMaxPool2dBaseTiling::GetTilingKey() const
52+{
53+ return 0;
54+}
55+ 
56+ge::graphStatus AdaMaxPool2dBaseTiling::GetPlatformInfo()
57+{
58+ auto compileInfo = context_->GetCompileInfo<AdaptiveMaxPool2dCompileInfo>();
59+ OP_CHECK_NULL_WITH_CONTEXT(context_, compileInfo);
60+ input_.coreNum = compileInfo->coreNum;
61+ input_.ubSizePlatForm = compileInfo->ubSizePlatForm;
62+ OP_CHECK_IF(input_.coreNum <= 0, OP_LOGE(context_, "GetPlatformInfo get corenum <= 0"), return ge::GRAPH_FAILED);
63+ sysWorkspaceSize_ = compileInfo->sysWorkspaceSize;
64+ return ge::GRAPH_SUCCESS;
65+}
66+ 
67+ge::graphStatus AdaMaxPool2dBaseTiling::GetShapeAttrsInfo()
68+{
69+ auto nodeName = context_->GetNodeName();
70+ OP_LOGD(nodeName, "GetShapeAttrsInfo begin.");
71+ 
72+ auto inputX = context_->GetInputShape(0);
73+ OP_CHECK_NULL_WITH_CONTEXT(context_, inputX);
74+ auto inputXDesc = context_->GetInputDesc(0);
75+ OP_CHECK_NULL_WITH_CONTEXT(context_, inputXDesc);
76+ auto xDtype = inputXDesc->GetDataType();
77+ OP_CHECK_IF(
78+ (xDtype != ge::DT_FLOAT && xDtype != ge::DT_FLOAT16 && xDtype != ge::DT_BF16),
79+ OP_LOGE(nodeName, "x datatype only support float, float16, bfloat16"), return ge::GRAPH_FAILED);
80+ input_.xDtype = xDtype;
81+ gert::Shape xShape = EnsureNotScalar(inputX->GetStorageShape());
82+ if (xShape.GetDimNum() == DIM_NUM_FOUR) {
83+ input_.N = xShape.GetDim(NCHW_DIM_N);
84+ input_.C = xShape.GetDim(NCHW_DIM_C);
85+ input_.Hi = xShape.GetDim(NCHW_DIM_H);
86+ input_.Wi = xShape.GetDim(NCHW_DIM_W);
87+ } else {
88+ OP_LOGE(nodeName, "xShape dim number should be 5");
89+ return ge::GRAPH_FAILED;
90+ }
91+ OP_CHECK_IF(
92+ input_.N < 1 || input_.C < 1 || input_.Hi < 1 || input_.Wi < 1,
93+ OP_LOGE(nodeName, "Invalid shape. Maybe empty tensor."), return ge::GRAPH_FAILED);
94+ OP_CHECK_IF(input_.Hi * input_.Wi > static_cast<int64_t>(std::numeric_limits<int32_t>::max()),
95+ OP_LOGE(nodeName, "no support for H*W of input greater than int32 max value"), return ge::GRAPH_FAILED);
96+ 
97+ auto attrPtr = context_->GetAttrs();
98+ OP_CHECK_NULL_WITH_CONTEXT(context_, attrPtr);
99+ auto outputSizePtr = attrPtr->GetAttrPointer<gert::ContinuousVector>(0);
100+ OP_CHECK_NULL_WITH_CONTEXT(context_, outputSizePtr);
101+ OP_CHECK_IF(
102+ outputSizePtr->GetSize() != OUTPUTSIZE_DIM_MAX, OP_LOGE(nodeName, "the size of outputsize only support 2"),
103+ return ge::GRAPH_FAILED);
104+ const int64_t* outputSize = static_cast<const int64_t*>(outputSizePtr->GetData());
105+ OP_CHECK_IF(
106+ outputSize[0] <= 0 || outputSize[1] <= 0,
107+ OP_LOGE(nodeName, "the value of outputsize should > 0"), return ge::GRAPH_FAILED);
108+ input_.Ho = outputSize[0];
109+ input_.Wo = outputSize[1];
110+ return ge::GRAPH_SUCCESS;
111+}
112+ 
113+ge::graphStatus AdaMaxPool2dBaseTiling::GetWorkspaceSize()
114+{
115+ size_t* workspaces = context_->GetWorkspaceSizes(1);
116+ workspaces[0] = sysWorkspaceSize_;
117+ return ge::GRAPH_SUCCESS;
118+}
119+ 
120+ge::graphStatus AdaMaxPool2dBaseTiling::PostTiling()
121+{
122+ return ge::GRAPH_SUCCESS;
123+}
124+ 
125+static ge::graphStatus Tiling4AdaptiveMaxPool2d(gert::TilingContext* context)
126+{
127+ return TilingRegistry::GetInstance().DoTilingImpl(context);
128+}
129+ 
130+static ge::graphStatus TilingPrepare4AdaptiveMaxPool2d(gert::TilingParseContext* context)
131+{
132+ OP_LOGD(context, "TilingPrepare4AdaptiveMaxPool2d enter.");
133+ 
134+ auto compileInfo = context->GetCompiledInfo<AdaptiveMaxPool2dCompileInfo>();
135+ OP_CHECK_NULL_WITH_CONTEXT(context, compileInfo);
136+ auto platformInfo = context->GetPlatformInfo();
137+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfo);
138+ 
139+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo);
140+ compileInfo->coreNum = ascendcPlatform.GetCoreNumAiv();
141+ uint64_t ubSizePlatForm;
142+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSizePlatForm);
143+ compileInfo->ubSizePlatForm = ubSizePlatForm;
144+ 
145+ size_t sysWorkspaceSize = ascendcPlatform.GetLibApiWorkSpaceSize();
146+ compileInfo->sysWorkspaceSize = sysWorkspaceSize;
147+ 
148+ OP_CHECK_IF((compileInfo->coreNum <= 0), OP_LOGE(context, "Failed to get corenum size"), return ge::GRAPH_FAILED);
149+ OP_CHECK_IF((compileInfo->ubSizePlatForm <= 0), OP_LOGE(context, "Failed to get ub size"), return ge::GRAPH_FAILED);
150+ OP_LOGD(context, "ub_size_platform is %lu", compileInfo->ubSizePlatForm);
151+ OP_LOGD(context, "TilingPrepare4AdaptiveMaxPool2d end");
152+ return ge::GRAPH_SUCCESS;
153+}
154+ 
155+IMPL_OP_OPTILING(AdaptiveMaxPool2d)
156+ .Tiling(Tiling4AdaptiveMaxPool2d)
157+ .TilingParse<AdaptiveMaxPool2dCompileInfo>(TilingPrepare4AdaptiveMaxPool2d);
158+} // namespace optiling
Apooling/adaptive_max_pool2d/op_host/adaptive_max_pool2d_tiling_base.h+136-0
@@ -0,0 +1,136 @@
1+/**
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 of
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.
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.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+ /*!
12+ * \file adaptive_max_pool2d_tiling_base.h
13+ * \brief tiling base imply for adaptive_max_pool2d
14+ */
15+ 
16+#ifndef AIR_CXX_RUNTIME_V2_OP_IMPL_ADAPTIVE_MAX_POOL2D_TILING_BASE_H_
17+#define AIR_CXX_RUNTIME_V2_OP_IMPL_ADAPTIVE_MAX_POOL2D_TILING_BASE_H_
18+ 
19+#include <array>
20+#include "log/log.h"
21+#include "register/op_impl_registry.h"
22+#include "register/tilingdata_base.h"
23+#include "tiling_base/tiling_base.h"
24+#include "tiling_base/tiling_templates_registry.h"
25+#include "util/math_util.h"
26+ 
27+using namespace std;
28+ 
29+namespace optiling {
30+using Ops::NN::Optiling::TilingBaseClass;
31+const int HW_DIMS = 2;
32+const uint32_t H_DIM = 0;
33+const uint32_t W_DIM = 1;
34+const uint32_t MAX_DIV = 2;
35+const uint32_t NCHW_CONV_ADDR_LIST_SIZE = 16;
36+const uint32_t MIN_TRANSPOSE_ROWS = 16;
37+const uint32_t INT64_FP32 = 2;
38+const uint32_t BINARY_SEARCH_COEFF = 2;
39+const uint32_t BLOCK_LEN_FP32 = 8;
40+const uint32_t BLOCK_LEN_FP16 = 16;
41+ 
42+struct AdaptiveMaxPool2dCompileInfo {
43+ uint64_t coreNum = 0;
44+ uint64_t ubSizePlatForm = 0;
45+ size_t sysWorkspaceSize = 0;
46+};
47+ 
48+struct InputInfo {
49+ uint64_t coreNum{0};
50+ uint64_t ubSizePlatForm{0};
51+ ge::DataType xDtype{ge::DT_FLOAT};
52+ ge::DataType indicesDtype{ge::DT_INT32};
53+ uint64_t N{0};
54+ uint64_t C{0};
55+ uint64_t Hi{0};
56+ uint64_t Wi{0};
57+ uint64_t Ho{0};
58+ uint64_t Wo{0};
59+};
60+ 
61+struct CalculateInfo {
62+ uint64_t useCoreNum{0};
63+ uint64_t totalIdx{0};
64+ uint64_t blockFactor{0};
65+ uint64_t blockTail{0};
66+ uint64_t ncFactor{0};
67+ uint64_t hoFactor{0};
68+ uint64_t woFactor{0};
69+ uint64_t ncOuter{0};
70+ uint64_t hoOuter{0};
71+ uint64_t woOuter{0};
72+ uint64_t ncTail{0};
73+ uint64_t hoTail{0};
74+ uint64_t woTail{0};
75+ uint64_t kernelHMax{0};
76+ uint64_t kernelWMax{0};
77+};
78+ 
79+BEGIN_TILING_DATA_DEF(AdaptiveMaxPool2dTilingData)
80+TILING_DATA_FIELD_DEF(int64_t, N);
81+TILING_DATA_FIELD_DEF(int64_t, C);
82+TILING_DATA_FIELD_DEF(int64_t, Hi);
83+TILING_DATA_FIELD_DEF(int64_t, Wi);
84+TILING_DATA_FIELD_DEF(int64_t, Ho);
85+TILING_DATA_FIELD_DEF(int64_t, Wo);
86+TILING_DATA_FIELD_DEF(int64_t, coreNums);
87+TILING_DATA_FIELD_DEF(int64_t, useCoreNum);
88+TILING_DATA_FIELD_DEF(int64_t, totalIdx);
89+TILING_DATA_FIELD_DEF(int64_t, blockFactor);
90+TILING_DATA_FIELD_DEF(int64_t, blockTail);
91+TILING_DATA_FIELD_DEF(int64_t, ncFactor);
92+TILING_DATA_FIELD_DEF(int64_t, hoFactor);
93+TILING_DATA_FIELD_DEF(int64_t, woFactor);
94+TILING_DATA_FIELD_DEF(int64_t, ncOuter);
95+TILING_DATA_FIELD_DEF(int64_t, hoOuter);
96+TILING_DATA_FIELD_DEF(int64_t, woOuter);
97+TILING_DATA_FIELD_DEF(int64_t, ncTail);
98+TILING_DATA_FIELD_DEF(int64_t, hoTail);
99+TILING_DATA_FIELD_DEF(int64_t, woTail);
100+ 
101+TILING_DATA_FIELD_DEF(int64_t, threadNums);
102+TILING_DATA_FIELD_DEF(int64_t, blockNums);
103+TILING_DATA_FIELD_DEF(int64_t, kMaxSizeH);
104+TILING_DATA_FIELD_DEF(int64_t, kMaxSizeW);
105+END_TILING_DATA_DEF;
106+ 
107+REGISTER_TILING_DATA_CLASS(AdaptiveMaxPool2d, AdaptiveMaxPool2dTilingData);
108+ 
109+ 
110+class AdaMaxPool2dBaseTiling : public TilingBaseClass {
111+public:
112+ explicit AdaMaxPool2dBaseTiling(gert::TilingContext* context) : TilingBaseClass(context)
113+ {}
114+ 
115+ ~AdaMaxPool2dBaseTiling() override
116+ {}
117+ 
118+protected:
119+ bool IsCapable() override;
120+ ge::graphStatus GetPlatformInfo() override;
121+ ge::graphStatus GetShapeAttrsInfo() override;
122+ ge::graphStatus DoOpTiling() override;
123+ ge::graphStatus DoLibApiTiling() override;
124+ uint64_t GetTilingKey() const override;
125+ ge::graphStatus GetWorkspaceSize() override;
126+ ge::graphStatus PostTiling() override;
127+ 
128+public:
129+ ge::DataType dtype = ge::DataType::DT_FLOAT;
130+ size_t sysWorkspaceSize_ = 0;
131+ InputInfo input_;
132+ CalculateInfo calInfo_;
133+};
134+} // namespace optiling
135+ 
136+#endif // AIR_CXX_RUNTIME_V2_OP_IMPL_ADAPTIVE_MAX_POOL2D_TILING_BASE_H_
Apooling/adaptive_max_pool2d/op_host/config/adaptive_max_pool2d_binary.json+317-0
@@ -0,0 +1,317 @@
1+{
2+ "op_type": "AdaptiveMaxPool2d",
3+ "op_list": [
4+ {
5+ "bin_filename": "AdaptiveMaxPool2d_be522d0caeae375632e98b1a501019f3",
6+ "inputs": [
7+ {
8+ "name": "x",
9+ "index": 0,
10+ "dtype": "float32",
11+ "format": "ND",
12+ "paramType": "required",
13+ "shape": [
14+ -2
15+ ],
16+ "format_match_mode": "FormatAgnostic"
17+ }
18+ ],
19+ "outputs": [
20+ {
21+ "name": "y",
22+ "index": 0,
23+ "dtype": "float32",
24+ "format": "ND",
25+ "paramType": "required",
26+ "shape": [
27+ -2
28+ ],
29+ "format_match_mode": "FormatAgnostic"
30+ },
31+ {
32+ "name": "indices",
33+ "index": 1,
34+ "dtype": "int32",
35+ "format": "ND",
36+ "paramType": "required",
37+ "shape": [
38+ -2
39+ ],
40+ "format_match_mode": "FormatAgnostic"
41+ }
42+ ],
43+ "attrs": [
44+ {
45+ "name": "output_size",
46+ "dtype": "list_int",
47+ "value": []
48+ },
49+ {
50+ "name": "indices_dtype",
51+ "dtype": "int",
52+ "value": 0
53+ }
54+ ]
55+ },
56+ {
57+ "bin_filename": "AdaptiveMaxPool2d_95c63351e3c3a70b21f41be389ff8ff4",
58+ "inputs": [
59+ {
60+ "name": "x",
61+ "index": 0,
62+ "dtype": "bfloat16",
63+ "format": "ND",
64+ "paramType": "required",
65+ "shape": [
66+ -2
67+ ],
68+ "format_match_mode": "FormatAgnostic"
69+ }
70+ ],
71+ "outputs": [
72+ {
73+ "name": "y",
74+ "index": 0,
75+ "dtype": "bfloat16",
76+ "format": "ND",
77+ "paramType": "required",
78+ "shape": [
79+ -2
80+ ],
81+ "format_match_mode": "FormatAgnostic"
82+ },
83+ {
84+ "name": "indices",
85+ "index": 1,
86+ "dtype": "int32",
87+ "format": "ND",
88+ "paramType": "required",
89+ "shape": [
90+ -2
91+ ],
92+ "format_match_mode": "FormatAgnostic"
93+ }
94+ ],
95+ "attrs": [
96+ {
97+ "name": "output_size",
98+ "dtype": "list_int",
99+ "value": []
100+ },
101+ {
102+ "name": "indices_dtype",
103+ "dtype": "int",
104+ "value": 0
105+ }
106+ ]
107+ },
108+ {
109+ "bin_filename": "AdaptiveMaxPool2d_d18392463d762fbcc48f41e994d30719",
110+ "inputs": [
111+ {
112+ "name": "x",
113+ "index": 0,
114+ "dtype": "float16",
115+ "format": "ND",
116+ "paramType": "required",
117+ "shape": [
118+ -2
119+ ],
120+ "format_match_mode": "FormatAgnostic"
121+ }
122+ ],
123+ "outputs": [
124+ {
125+ "name": "y",
126+ "index": 0,
127+ "dtype": "float16",
128+ "format": "ND",
129+ "paramType": "required",
130+ "shape": [
131+ -2
132+ ],
133+ "format_match_mode": "FormatAgnostic"
134+ },
135+ {
136+ "name": "indices",
137+ "index": 1,
138+ "dtype": "int64",
139+ "format": "ND",
140+ "paramType": "required",
141+ "shape": [
142+ -2
143+ ],
144+ "format_match_mode": "FormatAgnostic"
145+ }
146+ ],
147+ "attrs": [
148+ {
149+ "name": "output_size",
150+ "dtype": "list_int",
151+ "value": []
152+ },
153+ {
154+ "name": "indices_dtype",
155+ "dtype": "int",
156+ "value": 0
157+ }
158+ ]
159+ },
160+ {
161+ "bin_filename": "AdaptiveMaxPool2d_946b3699e34bac16354098363baf0f3c",
162+ "inputs": [
163+ {
164+ "name": "x",
165+ "index": 0,
166+ "dtype": "bfloat16",
167+ "format": "ND",
168+ "paramType": "required",
169+ "shape": [
170+ -2
171+ ],
172+ "format_match_mode": "FormatAgnostic"
173+ }
174+ ],
175+ "outputs": [
176+ {
177+ "name": "y",
178+ "index": 0,
179+ "dtype": "bfloat16",
180+ "format": "ND",
181+ "paramType": "required",
182+ "shape": [
183+ -2
184+ ],
185+ "format_match_mode": "FormatAgnostic"
186+ },
187+ {
188+ "name": "indices",
189+ "index": 1,
190+ "dtype": "int64",
191+ "format": "ND",
192+ "paramType": "required",
193+ "shape": [
194+ -2
195+ ],
196+ "format_match_mode": "FormatAgnostic"
197+ }
198+ ],
199+ "attrs": [
200+ {
201+ "name": "output_size",
202+ "dtype": "list_int",
203+ "value": []
204+ },
205+ {
206+ "name": "indices_dtype",
207+ "dtype": "int",
208+ "value": 0
209+ }
210+ ]
211+ },
212+ {
213+ "bin_filename": "AdaptiveMaxPool2d_3b64a0aab8fee8a04f831218e79592ba",
214+ "inputs": [
215+ {
216+ "name": "x",
217+ "index": 0,
218+ "dtype": "float32",
219+ "format": "ND",
220+ "paramType": "required",
221+ "shape": [
222+ -2
223+ ],
224+ "format_match_mode": "FormatAgnostic"
225+ }
226+ ],
227+ "outputs": [
228+ {
229+ "name": "y",
230+ "index": 0,
231+ "dtype": "float32",
232+ "format": "ND",
233+ "paramType": "required",
234+ "shape": [
235+ -2
236+ ],
237+ "format_match_mode": "FormatAgnostic"
238+ },
239+ {
240+ "name": "indices",
241+ "index": 1,
242+ "dtype": "int64",
243+ "format": "ND",
244+ "paramType": "required",
245+ "shape": [
246+ -2
247+ ],
248+ "format_match_mode": "FormatAgnostic"
249+ }
250+ ],
251+ "attrs": [
252+ {
253+ "name": "output_size",
254+ "dtype": "list_int",
255+ "value": []
256+ },
257+ {
258+ "name": "indices_dtype",
259+ "dtype": "int",
260+ "value": 0
261+ }
262+ ]
263+ },
264+ {
265+ "bin_filename": "AdaptiveMaxPool2d_4a1e4f3b6a197926a3518a0bcd1d53f9",
266+ "inputs": [
267+ {
268+ "name": "x",
269+ "index": 0,
270+ "dtype": "float16",
271+ "format": "ND",
272+ "paramType": "required",
273+ "shape": [
274+ -2
275+ ],
276+ "format_match_mode": "FormatAgnostic"
277+ }
278+ ],
279+ "outputs": [
280+ {
281+ "name": "y",
282+ "index": 0,
283+ "dtype": "float16",
284+ "format": "ND",
285+ "paramType": "required",
286+ "shape": [
287+ -2
288+ ],
289+ "format_match_mode": "FormatAgnostic"
290+ },
291+ {
292+ "name": "indices",
293+ "index": 1,
294+ "dtype": "int32",
295+ "format": "ND",
296+ "paramType": "required",
297+ "shape": [
298+ -2
299+ ],
300+ "format_match_mode": "FormatAgnostic"
301+ }
302+ ],
303+ "attrs": [
304+ {
305+ "name": "output_size",
306+ "dtype": "list_int",
307+ "value": []
308+ },
309+ {
310+ "name": "indices_dtype",
311+ "dtype": "int",
312+ "value": 0
313+ }
314+ ]
315+ }
316+ ]
317+}
Apooling/adaptive_max_pool2d/op_host/config/adaptive_max_pool2d_simplified_key.ini+13-0
@@ -0,0 +1,13 @@
1+; 该文件主要影响 opc 工具 编译二进制kernel时， --simplified_key_mode 选项中填写的值，格式如下所示：
2+; [某算子]
3+; default=xx
4+; ascendxx=xx
5+; 其中，default为默认mode，ascendxx为可选mode，如果不同芯片有差异化要求时，需要配置；
6+; 1)如果没有配置：非ascendC算子继续按空处理，即opc编译命令中不添加 --simplified_key_mode 选项，AscendC算子按照 simplified_key_mode=0 处理
7+; 2)如果仅有default配置：各个版本按default配置
8+; 3)如果仅有某些平台的配置，没有default配置：对应平台的按照配置的值传递，非对应平台的：非AscendC算子继续按空处理，AscendC算子按照 simplified_key_mode=0 处理
9+; 4)如果default配置和平台配置都有：对应平台的使用平台的配置，非对应的平台的以default值配置。
10+; 5)对于自定义simplified key的情况，需要在binary_simplified_key_mode.ini 文件中显式配置为None，不传入 --simplified_key_mode 选项，由opc工具和FE框架自行判断使用何种模式
11+; 6)是否是AscendC算子，由 ops/build-in/tbe/op_info_cfg/parser/ascendc_config.json 中配置的算子名字和对于的平台决定
12+[AdaptiveMaxPool2d]
13+default=0
Apooling/adaptive_max_pool2d/op_host/max_pool2d_with_argmax_v2_simt_tiling.cpp+226-0
@@ -0,0 +1,226 @@
1+/**
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 of
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.
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.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file max_pool2d_with_argmax_v2_simt_tiling.cpp
13+ * \brief
14+ */
15+ 
16+#include <cctype>
17+#include <algorithm>
18+#include "log/log.h"
19+#include "util/math_util.h"
20+#include "error_util.h"
21+#include "tiling_base/tiling_base.h"
22+#include "tiling_base/tiling_templates_registry.h"
23+#include "adaptive_max_pool2d_simt_tiling.h"
24+#include "op_common/op_host/util/platform_util.h"
25+#include "platform/platform_ascendc.h"
26+#include "register/op_def_registry.h"
27+#include "platform/platform_info.h"
28+#include "register/op_impl_registry.h"
29+ 
30+ 
31+using namespace ge;
32+ 
33+constexpr uint64_t CAL_KER_THRESHOLD = 10000;
34+constexpr int64_t N_IDX = 0;
35+constexpr int64_t C_IDX = 1;
36+constexpr int64_t H_IDX = 2;
37+constexpr int64_t W_IDX = 3;
38+ 
39+namespace optiling{
40+ 
41+static const gert::Shape g_vec_1_shape = {1};
42+ 
43+static const gert::Shape& EnsureNotScalar(const gert::Shape &inShape) {
44+ if (inShape.IsScalar()) {
45+ return g_vec_1_shape;
46+ }
47+ return inShape;
48+}
49+ 
50+bool AdaMaxPool2dTilingSIMT::IsCapable()
51+{
52+ return true;
53+}
54+ 
55+ge::graphStatus AdaMaxPool2dTilingSIMT::GetPlatformInfo()
56+{
57+ auto compileInfo = reinterpret_cast<const AdaptiveMaxPool2dCompileInfo*>(context_->GetCompileInfo());
58+ OP_CHECK_NULL_WITH_CONTEXT(context_, compileInfo);
59+ coreNum_ = compileInfo->coreNum;
60+ OP_CHECK_IF(coreNum_ <= 0, OP_LOGE(context_, "GetPlatformInfo get corenum <= 0"), return ge::GRAPH_FAILED);
61+ sysWorkspaceSize_ = compileInfo->sysWorkspaceSize;
62+ return ge::GRAPH_SUCCESS;
63+}
64+ 
65+ge::graphStatus AdaMaxPool2dTilingSIMT::CheckPlatformAndGetShapes() {
66+ auto platformInfo = context_->GetPlatformInfo();
67+ OP_CHECK_NULL_WITH_CONTEXT(context_, platformInfo);
68+ auto inputX = context_->GetInputShape(FIRPOS);
69+ OPS_CHECK_NULL_WITH_CONTEXT(context_, inputX);
70+ auto inputShape = EnsureNotScalar(inputX->GetStorageShape());
71+ auto outX = context_->GetOutputShape(FIRPOS);
72+ OPS_CHECK_NULL_WITH_CONTEXT(context_, outX);
73+ auto outShape = EnsureNotScalar(outX->GetStorageShape());
74+ auto indicesX = context_->GetOutputShape(SECPOS);
75+ OPS_CHECK_NULL_WITH_CONTEXT(context_, indicesX);
76+ if (inputShape.GetDimNum() != NCHW_DIMS) {
77+ VECTOR_INNER_ERR_REPORT_TILIING(context_->GetNodeName(),
78+ "AdaptiveMaxPool2d: input shape dim = %zu, should be equal 4",
79+ inputShape.GetDimNum());
80+ return ge::GRAPH_FAILED;
81+ }
82+ OP_CHECK_IF(
83+ inputShape.GetDim(N_IDX) < 1 || inputShape.GetDim(C_IDX) < 1 ||
84+ inputShape.GetDim(H_IDX) < 1 || inputShape.GetDim(W_IDX) < 1,
85+ OP_LOGE(context_->GetNodeName(), "Invalid shape. Maybe empty tensor."), return ge::GRAPH_FAILED);
86+ 
87+ inputData.inputShape =
88+ array<uint64_t, NCHW_DIMS>{uint64_t(inputShape.GetDim(N_IDX)), uint64_t(inputShape.GetDim(C_IDX)),
89+ uint64_t(inputShape.GetDim(H_IDX)), uint64_t(inputShape.GetDim(W_IDX))};
90+ inputData.outShape =
91+ array<uint64_t, NCHW_DIMS>{uint64_t(inputShape.GetDim(N_IDX)), uint64_t(inputShape.GetDim(C_IDX)),
92+ uint64_t(outShape.GetDim(H_IDX)), uint64_t(outShape.GetDim(W_IDX))};
93+ 
94+ return ge::GRAPH_SUCCESS;
95+}
96+ 
97+ge::graphStatus AdaMaxPool2dTilingSIMT::CheckDataTypeAndAttrs() {
98+ auto inputDesc = context_->GetInputDesc(0);
99+ OPS_CHECK_NULL_WITH_CONTEXT(context_, inputDesc);
100+ dtype = inputDesc->GetDataType();
101+ if (dtype != ge::DataType::DT_BF16 && dtype != ge::DataType::DT_FLOAT16 && dtype != ge::DataType::DT_FLOAT) {
102+ VECTOR_INNER_ERR_REPORT_TILIING(context_->GetNodeName(),
103+ "AdaptiveMaxPool2d: invalid dtype %s, should be BFloat16、Float16 or Float32",
104+ Ops::Base::ToString(dtype).c_str());
105+ return ge::GRAPH_FAILED;
106+ }
107+ 
108+ auto indicesX = context_->GetOutputShape(SECPOS);
109+ OPS_CHECK_NULL_WITH_CONTEXT(context_, indicesX);
110+ auto indicesShape = EnsureNotScalar(indicesX->GetStorageShape());
111+ if (indicesShape.GetDimNum() != NCHW_DIMS) {
112+ VECTOR_INNER_ERR_REPORT_TILIING(context_->GetNodeName(),
113+ "AdaptiveMaxPool2d: indices shape dim = %zu, should be 4",
114+ indicesShape.GetDimNum());
115+ return ge::GRAPH_FAILED;
116+ }
117+ 
118+ array<uint64_t, NCHW_DIMS> indicesArray{
119+ uint64_t(indicesShape.GetDim(N_IDX)),
120+ uint64_t(indicesShape.GetDim(C_IDX)),
121+ uint64_t(indicesShape.GetDim(H_IDX)),
122+ uint64_t(indicesShape.GetDim(W_IDX))
123+ };
124+ if (indicesArray != inputData.outShape) {
125+ VECTOR_INNER_ERR_REPORT_TILIING(context_->GetNodeName(),
126+ "AdaptiveMaxPool2d: indices shape and values shape is different");
127+ return ge::GRAPH_FAILED;
128+ }
129+ 
130+ auto attrPtr = context_->GetAttrs();
131+ OP_CHECK_NULL_WITH_CONTEXT(context_, attrPtr);
132+ auto outputSizePtr = attrPtr->GetAttrPointer<gert::ContinuousVector>(0);
133+ OP_CHECK_NULL_WITH_CONTEXT(context_, outputSizePtr);
134+ OP_CHECK_IF(
135+ outputSizePtr->GetSize() != DOUB,
136+ OP_LOGE(context_->GetNodeName(), "the size of outputsize only support 2"),
137+ return ge::GRAPH_FAILED);
138+ const int64_t* outputSize = static_cast<const int64_t*>(outputSizePtr->GetData());
139+ OP_CHECK_IF(
140+ outputSize[0] <= 0 || outputSize[1] <= 0,
141+ OP_LOGE(context_->GetNodeName(), "the value of outputsize should > 0"), return ge::GRAPH_FAILED);
142+ 
143+ return ge::GRAPH_SUCCESS;
144+}
145+ 
146+ge::graphStatus AdaMaxPool2dTilingSIMT::GetShapeAttrsInfo() {
147+ auto status = CheckPlatformAndGetShapes();
148+ if (status != ge::GRAPH_SUCCESS) return status;
149+ return CheckDataTypeAndAttrs();
150+}
151+ 
152+ge::graphStatus AdaMaxPool2dTilingSIMT::DoOpTiling()
153+{
154+ tiling.set_N(inputData.inputShape[N_DIM_]);
155+ tiling.set_C(inputData.inputShape[C_DIM_]);
156+ tiling.set_Hi(inputData.inputShape[H_DIM_]);
157+ tiling.set_Wi(inputData.inputShape[W_DIM_]);
158+ tiling.set_Ho(inputData.outShape[H_DIM_]);
159+ tiling.set_Wo(inputData.outShape[W_DIM_]);
160+ tiling.set_kMaxSizeH(inputData.kernelHMax);
161+ tiling.set_kMaxSizeW(inputData.kernelWMax);
162+ tiling.set_coreNums(0);
163+ tiling.set_useCoreNum(0);
164+ tiling.set_totalIdx(0);
165+ tiling.set_blockFactor(0);
166+ tiling.set_blockTail(0);
167+ tiling.set_ncFactor(0);
168+ tiling.set_hoFactor(0);
169+ tiling.set_woFactor(0);
170+ tiling.set_ncOuter(0);
171+ tiling.set_hoOuter(0);
172+ tiling.set_woOuter(0);
173+ tiling.set_ncTail(0);
174+ tiling.set_hoTail(0);
175+ tiling.set_woTail(0);
176+ int64_t outputDataCount = tiling.get_N() * tiling.get_C() * tiling.get_Hi() * tiling.get_Wi();
177+ int64_t threads = std::min(outputDataCount, MAX_THREAD_NUM);
178+ int64_t blockNum = Ops::Base::CeilDiv(outputDataCount, threads);
179+ blockNum = std::min(blockNum, static_cast<int64_t>(coreNum_));
180+ context_->SetBlockDim(blockNum);
181+ context_->SetTilingKey(GetTilingKey());
182+ tiling.set_threadNums(threads);
183+ tiling.set_blockNums(blockNum);
184+ return ge::GRAPH_SUCCESS;
185+}
186+ 
187+uint64_t AdaMaxPool2dTilingSIMT::GetTilingKey() const
188+{
189+ return 0;
190+}
191+ 
192+ge::graphStatus AdaMaxPool2dTilingSIMT::GetWorkspaceSize()
193+{
194+ size_t* currentWorkspace = context_->GetWorkspaceSizes(1);
195+ OP_CHECK_NULL_WITH_CONTEXT(context_, currentWorkspace);
196+ currentWorkspace[0] = sysWorkspaceSize_;
197+ return ge::GRAPH_SUCCESS;
198+}
199+ 
200+ge::graphStatus AdaMaxPool2dTilingSIMT::PostTiling()
201+{
202+ OP_CHECK_IF(context_->GetRawTilingData()->GetCapacity() < tiling.GetDataSize(),
203+ OP_LOGE(context_, "tiling data's[%zu] is larger than capacity[%zu].", tiling.GetDataSize(),
204+ context_->GetRawTilingData()->GetCapacity()),
205+ return ge::GRAPH_FAILED);
206+ tiling.SaveToBuffer(context_->GetRawTilingData()->GetData(), context_->GetRawTilingData()->GetCapacity());
207+ context_->GetRawTilingData()->SetDataSize(tiling.GetDataSize());
208+ return ge::GRAPH_SUCCESS;
209+}
210+ 
211+void AdaMaxPool2dTilingSIMT::DumpTilingInfo()
212+{
213+ std::string str;
214+ str += " threadNums:" + std::to_string(tiling.get_threadNums());
215+ str += " blockNums:" + std::to_string(tiling.get_blockNums());
216+ str += " nDim:" + std::to_string(tiling.get_N());
217+ str += " cDim:" + std::to_string(tiling.get_C());
218+ str += " hInDim:" + std::to_string(tiling.get_Hi());
219+ str += " wInDim:" + std::to_string(tiling.get_Wi());
220+ str += " hOutDim:" + std::to_string(tiling.get_Ho());
221+ str += " wOutDim:" + std::to_string(tiling.get_Wo());
222+ str += " kMaxSizeH:" + std::to_string(tiling.get_kMaxSizeH());
223+ str += " kMaxSizeW:" + std::to_string(tiling.get_kMaxSizeW());
224+}
225+REGISTER_TILING_TEMPLATE("AdaptiveMaxPool2d", AdaMaxPool2dTilingSIMT, 0);
226+} // namespace optiling
Apooling/adaptive_max_pool2d/op_kernel/adaptive_max_pool2d.cpp+33-0
@@ -0,0 +1,33 @@
1+/**
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 of
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.
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.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/* !
12+ * \file adaptive_max_pool2d.cpp
13+ * \brief adaptive_max_pool2d implied
14+ */
15+ 
16+#include <cstdint>
17+#include "kernel_operator.h"
18+#include "kernel_tiling/kernel_tiling.h"
19+#include "./arch35/adaptive_max_pool2d_simt.h"
20+using namespace AscendC;
21+ 
22+extern "C" __global__ __aicore__ void adaptive_max_pool2d(
23+ GM_ADDR x, GM_ADDR y, GM_ADDR indices, GM_ADDR workspace, GM_ADDR tiling)
24+{
25+ AscendC::TPipe pipeBase;
26+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
27+ GET_TILING_DATA(tilingData, tiling);
28+ if (TILING_KEY_IS(0)) {
29+ AdaptiveMaxPool2DWithSimt::AdaptiveMaxPool2DSimt<DTYPE_X, DTYPE_INDICES, int32_t, uint32_t> op(&pipeBase, &tilingData);
30+ op.Init(x, y, indices);
31+ op.Process();
32+ }
33+}
Apooling/adaptive_max_pool2d/op_kernel/arch35/adaptive_max_pool2d_simt.h+227-0
@@ -0,0 +1,227 @@
1+/**
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 of
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.
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.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file adaptive_max_pool2d_simt.h
13+ * \brief
14+ */
15+ 
16+#ifndef ADAPTIVE_MAX_POOL_2D_SIMT_H
17+#define ADAPTIVE_MAX_POOL_2D_SIMT_H
18+ 
19+#include "kernel_operator.h"
20+#include "../inc/load_store_utils.h"
21+#include "../inc/platform.h"
22+#include "../inc/kernel_utils.h"
23+ 
24+#ifdef __CCE_KT_TEST__
25+#define LAUNCH_BOUND(threads)
26+#endif
27+ 
28+namespace AdaptiveMaxPool2DWithSimt{
29+ using namespace AscendC;
30+ 
31+ constexpr uint32_t THREAD_DIM = 1024;
32+ constexpr size_t PARAM_NUM = 4;
33+ constexpr static uint32_t DIV_HW_IDX = 0;
34+ constexpr static uint32_t DIV_H_IDX = 2;
35+ constexpr static uint32_t DIV_W_IDX = 4;
36+ 
37+ template <typename DIV_T>
38+ __aicore__ __attribute__((always_inline)) inline static DIV_T startIndex(DIV_T outIdx, DIV_T magicOutLen, DIV_T shiftOutLen, DIV_T inLen)
39+ {
40+ DIV_T pStart = outIdx * inLen;
41+ return Simt::UintDiv<DIV_T>(pStart, magicOutLen, shiftOutLen);
42+ }
43+ 
44+ template <typename DIV_T>
45+ __aicore__ __attribute__((always_inline)) inline static DIV_T endIndex(DIV_T outIdx, DIV_T magicOutLen, DIV_T shiftOutLen, DIV_T inLen)
46+ {
47+ DIV_T pEnd = ((outIdx + 1) * inLen - 1);
48+ pEnd = Simt::UintDiv<DIV_T>(pEnd, magicOutLen, shiftOutLen);
49+ return pEnd + 1;
50+ }
51+ 
52+ template <typename VALUE_T, typename INDICES_T, typename FORMAT_T, typename DIV_T>
53+ __aicore__ __attribute__((always_inline)) inline static void executeFunc(FORMAT_T count, __gm__ VALUE_T* bottomData,
54+ FORMAT_T ncSize, FORMAT_T height, FORMAT_T width,
55+ FORMAT_T outputNc, FORMAT_T outputHeight, FORMAT_T outputWidth,
56+ __gm__ VALUE_T* valueData, __gm__ INDICES_T* indicesData,
57+ DIV_T magicHW, DIV_T shiftHW, DIV_T magicH, DIV_T shiftH,
58+ DIV_T magicW, DIV_T shiftW) {
59+ for (DIV_T index = Simt::GetBlockIdx() * Simt::GetThreadNum() + Simt::GetThreadIdx(); index < count;
60+ index += Simt::GetBlockNum() * Simt::GetThreadNum()) {
61+ // 计算NC索引 (index / (outputHeight * outputWidth))
62+ DIV_T ncId = Simt::UintDiv<DIV_T>(index, magicHW, shiftHW);
63+ DIV_T indexIdx = index - ncId * outputNc; // 在当前NC内的偏移
64+
65+ // 计算高度和宽度索引 (indexIdx = h * outputWidth + w)
66+ DIV_T hId = Simt::UintDiv<DIV_T>(indexIdx, magicW, shiftW);
67+ DIV_T wId = indexIdx - hId * outputWidth;
68+
69+ // 计算输入窗口的起始和结束位置
70+ FORMAT_T startInH = startIndex<DIV_T>(hId, magicH, shiftH, height);
71+ FORMAT_T endInH = endIndex<DIV_T>(hId, magicH, shiftH, height);
72+ FORMAT_T startInW = startIndex<DIV_T>(wId, magicW, shiftW, width);
73+ FORMAT_T endInW = endIndex<DIV_T>(wId, magicW, shiftW, width);
74+
75+ VALUE_T maxVal = AscendC::NumericLimits<VALUE_T>::NegativeInfinity();
76+ FORMAT_T maxIdx = startInH * width + startInW; // 2D索引: h * width + w
77+
78+ auto ncStartData = bottomData + ncId * ncSize;
79+ 
80+ // 2D池化: 只在高度和宽度上滑动
81+ for (FORMAT_T h = startInH; h < endInH; ++h) {
82+ for (FORMAT_T w = startInW; w < endInW; ++w) {
83+ FORMAT_T idxOffset = h * width + w; // 2D索引
84+ VALUE_T val = static_cast<VALUE_T>(ncStartData[idxOffset]);
85+ if ((static_cast<VALUE_T>(val) > maxVal) || Simt::IsNan(static_cast<float>(val))) {
86+ maxIdx = idxOffset;
87+ maxVal = val;
88+ }
89+ }
90+ }
91+ valueData[index] = static_cast<VALUE_T>(maxVal);
92+ indicesData[index] = static_cast<INDICES_T>(maxIdx);
93+ }
94+ }
95+ 
96+template <typename VALUE_T, typename INDICES_T, typename FORMAT_T, typename DIV_T>
97+class AdaptiveMaxPool2DSimt
98+{
99+public:
100+ __aicore__ inline AdaptiveMaxPool2DSimt(TPipe *pipe, const AdaptiveMaxPool2dTilingData* __restrict tilingData)
101+ : pipe_(pipe), tilingData_(tilingData)
102+ {
103+ }
104+
105+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR indices);
106+ __aicore__ inline void Process();
107+ 
108+private:
109+ TPipe *pipe_;
110+ AscendC::GlobalTensor<VALUE_T> x_;
111+ AscendC::GlobalTensor<VALUE_T> y_;
112+ AscendC::GlobalTensor<INDICES_T> indices_;
113+ const AdaptiveMaxPool2dTilingData* tilingData_;
114+ TBuf<TPosition::VECCALC> paramBuf_;
115+};
116+ 
117+template <typename VALUE_T, typename INDICES_T, typename FORMAT_T, typename DIV_T>
118+__aicore__ inline void AdaptiveMaxPool2DSimt<VALUE_T, INDICES_T, FORMAT_T, DIV_T>::Init(GM_ADDR x, GM_ADDR y, GM_ADDR indices)
119+{
120+ x_.SetGlobalBuffer((__gm__ VALUE_T*)(x));
121+ y_.SetGlobalBuffer((__gm__ VALUE_T*)(y));
122+ indices_.SetGlobalBuffer((__gm__ INDICES_T*)(indices));
123+ if constexpr (!(std::is_same<FORMAT_T, int32_t>::value && std::is_same<DIV_T, uint32_t>::value)){
124+ pipe_->InitBuffer(paramBuf_, PARAM_NUM * sizeof(DIV_T));
125+ }
126+}
127+ 
128+template <typename VALUE_T, typename INDICES_T, typename FORMAT_T, typename DIV_T>
129+__simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM) inline void AdaptiveMaxPool2DNchwUb(FORMAT_T count, __gm__ VALUE_T* bottomData,
130+ FORMAT_T ncSize, FORMAT_T height, FORMAT_T width,
131+ FORMAT_T outputNc, FORMAT_T outputHeight, FORMAT_T outputWidth,
132+ __gm__ VALUE_T* valueData, __gm__ INDICES_T* indicesData,
133+ __ubuf__ DIV_T* SimtParam)
134+{
135+ DIV_T magicHW = SimtParam[DIV_HW_IDX];
136+ DIV_T shiftHW = SimtParam[DIV_HW_IDX + 1];
137+ DIV_T magicH = SimtParam[DIV_H_IDX];
138+ DIV_T shiftH = SimtParam[DIV_H_IDX + 1];
139+ DIV_T magicW = SimtParam[DIV_W_IDX];
140+ DIV_T shiftW = SimtParam[DIV_W_IDX + 1];
141+ executeFunc<VALUE_T, INDICES_T, FORMAT_T, DIV_T>(count, bottomData, ncSize, height, width,
142+ outputNc, outputHeight, outputWidth,
143+ valueData, indicesData,
144+ magicHW, shiftHW, magicH, shiftH, magicW, shiftW);
145+}
146+ 
147+template <typename VALUE_T, typename INDICES_T, typename FORMAT_T, typename DIV_T>
148+__simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM) inline void AdaptiveMaxPool2DNchwFunc(FORMAT_T count, __gm__ VALUE_T* bottomData,
149+ FORMAT_T ncSize, FORMAT_T height, FORMAT_T width,
150+ FORMAT_T outputNc, FORMAT_T outputHeight, FORMAT_T outputWidth,
151+ __gm__ VALUE_T* valueData, __gm__ INDICES_T* indicesData,
152+ DIV_T magicHW, DIV_T shiftHW, DIV_T magicH, DIV_T shiftH,
153+ DIV_T magicW, DIV_T shiftW)
154+{
155+ executeFunc<VALUE_T, INDICES_T, FORMAT_T, DIV_T>(count, bottomData, ncSize, height, width,
156+ outputNc, outputHeight, outputWidth,
157+ valueData, indicesData,
158+ magicHW, shiftHW, magicH, shiftH, magicW, shiftW);
159+}
160+ 
161+template <typename VALUE_T, typename INDICES_T, typename FORMAT_T, typename DIV_T>
162+__aicore__ inline void AdaptiveMaxPool2DSimt<VALUE_T, INDICES_T, FORMAT_T, DIV_T>::Process()
163+{
164+ auto inputData = (__gm__ VALUE_T*)x_.GetPhyAddr();
165+ auto outputData = (__gm__ VALUE_T*)y_.GetPhyAddr();
166+ auto indicesData = (__gm__ INDICES_T*)indices_.GetPhyAddr();
167+
168+ DIV_T hw = tilingData_->Ho * tilingData_->Wo; // H * W
169+ FORMAT_T totalSize = hw * tilingData_->N * tilingData_->C; // N * C * H * W
170+ FORMAT_T ncSize = tilingData_->Hi * tilingData_->Wi; // 每个N*C的2D数据大小
171+
172+ DIV_T magicHW = 0;
173+ DIV_T shiftHW = 0;
174+ DIV_T magicH = 0;
175+ DIV_T shiftH = 0;
176+ DIV_T magicW = 0;
177+ DIV_T shiftW = 0;
178+
179+ GetUintDivMagicAndShift<DIV_T>(magicHW, shiftHW, hw);
180+ GetUintDivMagicAndShift<DIV_T>(magicH, shiftH, tilingData_->Ho);
181+ GetUintDivMagicAndShift<DIV_T>(magicW, shiftW, tilingData_->Wo);
182+
183+ if constexpr (std::is_same<FORMAT_T, int32_t>::value && std::is_same<DIV_T, uint32_t>::value){
184+ Simt::VF_CALL<AdaptiveMaxPool2DNchwFunc<VALUE_T, INDICES_T, FORMAT_T, DIV_T>>(
185+ Simt::Dim3(THREAD_DIM),
186+ static_cast<FORMAT_T>(totalSize),
187+ inputData,
188+ ncSize,
189+ static_cast<FORMAT_T>(tilingData_->Hi),
190+ static_cast<FORMAT_T>(tilingData_->Wi),
191+ static_cast<FORMAT_T>(hw),
192+ static_cast<FORMAT_T>(tilingData_->Ho),
193+ static_cast<FORMAT_T>(tilingData_->Wo),
194+ outputData,
195+ indicesData,
196+ magicHW, shiftHW, magicH, shiftH, magicW, shiftW
197+ );
198+ } else {
199+ LocalTensor<DIV_T> SimtParam = paramBuf_.Get<DIV_T>();
200+ SimtParam.SetValue(0, static_cast<DIV_T>(magicHW));
201+ SimtParam.SetValue(1, static_cast<DIV_T>(shiftHW));
202+ SimtParam.SetValue(DIV_H_IDX, static_cast<DIV_T>(magicH));
203+ SimtParam.SetValue(DIV_H_IDX + 1, static_cast<DIV_T>(shiftH));
204+ SimtParam.SetValue(DIV_W_IDX, static_cast<DIV_T>(magicW));
205+ SimtParam.SetValue(DIV_W_IDX + 1, static_cast<DIV_T>(shiftW));
206+ DataSyncBarrier<MemDsbT::UB>();
207+
208+ Simt::VF_CALL<AdaptiveMaxPool2DNchwUb<VALUE_T, INDICES_T, FORMAT_T, DIV_T>>(
209+ Simt::Dim3(THREAD_DIM),
210+ static_cast<FORMAT_T>(totalSize),
211+ inputData,
212+ ncSize,
213+ static_cast<FORMAT_T>(tilingData_->Hi),
214+ static_cast<FORMAT_T>(tilingData_->Wi),
215+ static_cast<FORMAT_T>(hw),
216+ static_cast<FORMAT_T>(tilingData_->Ho),
217+ static_cast<FORMAT_T>(tilingData_->Wo),
218+ outputData,
219+ indicesData,
220+ (__ubuf__ DIV_T*)SimtParam.GetPhyAddr()
221+ );
222+ }
223+}
224+ 
225+}
226+ 
227+#endif // ADAPTIVE_MAX_POOL_2D_SIMT_H
Apooling/adaptive_max_pool2d/test/CMakeLists.txt+18-0
@@ -0,0 +1,18 @@
1+# -----------------------------------------------------------------------------------------------------------
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 of
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.
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.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# -----------------------------------------------------------------------------------------------------------
10+ 
11+message(STATUS "=== Debug: start ops.pooling.adaptive_max_pool2d.tests.CMakeLists.txt ")
12+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
13+message(STATUS "=== Debug: CURRENT_DIRS =${CURRENT_DIRS} ")
14+foreach(SUB_DIR ${CURRENT_DIRS})
15+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
16+ add_subdirectory(${SUB_DIR})
17+ endif()
18+endforeach()
Apooling/adaptive_max_pool2d/test/op_api/CMakeLists.txt+14-0
@@ -0,0 +1,14 @@
1+# ----------------------------------------------------------------------------
2+# This program is free software, you can redistribute it and/or modify.
3+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
4+# This file is a part of the CANN Open Software.
5+# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6+# Please refer to the License for details. You may not use this file except in compliance with the License.
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.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------
10+ 
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()