Star899
1810
代码
代码
Issues248
Pull Requests312
Actions
项目讨论
Wiki
分析
Star899
1810
已合并
支持MaxPool3DGradWithArgmax算子 #1193
小王！创建于 1月29日
支持MaxPool3DGradWithArgmax算子 #1193
已合并
小王！创建于 1月29日
从已删除 :MaxPool3DGradWithArgmax合入到cann/ops-nnmaster
共 25 个文件变更+6055-687
Mpooling/max_pool3d_grad_with_argmax/docs/aclnnMaxPool3dWithArgmaxBackward.md+8-14
@@ -6,13 +6,12 @@
6 6 
7| 产品 | 是否支持 |7| 产品 | 是否支持 |
8| :----------------------------------------------------------- | :------: |8| :----------------------------------------------------------- | :------: |
9-| <term>昇腾910_95 AI处理器</term> | × |9+| <term>Ascend 950PR/Ascend 950DT</term> | √ |
10| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |10| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |
11-| <term>Atlas A2 训练系列产品/Atlas 800I A2 推理产品/A200I A2 Box 异构组件</term> | √ |11+| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |
12| <term>Atlas 200I/500 A2 推理产品</term> | × |12| <term>Atlas 200I/500 A2 推理产品</term> | × |
13-| <term>Atlas 推理系列产品 </term> | × |13+| <term>Atlas 推理系列产品</term> | × |
14| <term>Atlas 训练系列产品</term> | × |14| <term>Atlas 训练系列产品</term> | × |
15-| <term>Atlas 200/300/500 推理产品</term> | × |
16 15 
17## 功能说明16## 功能说明
18正向最大池化[aclnnMaxPool3dWithArgmax](../../max_pool3d_with_argmax_v2/docs/aclnnMaxPool3dWithArgmax.md)的反向传播，将梯度回填到每个窗口最大值的坐标处，相同坐标处累加。17正向最大池化[aclnnMaxPool3dWithArgmax](../../max_pool3d_with_argmax_v2/docs/aclnnMaxPool3dWithArgmax.md)的反向传播，将梯度回填到每个窗口最大值的坐标处，相同坐标处累加。
@@ -20,7 +19,6 @@
20## 函数原型19## 函数原型
21每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnMaxPool3dWithArgmaxBackwardGetWorkspaceSize”接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用“aclnnMaxPool3dWithArgmaxBackward”接口执行计算。20每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnMaxPool3dWithArgmaxBackwardGetWorkspaceSize”接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用“aclnnMaxPool3dWithArgmaxBackward”接口执行计算。
22 21 
23- 
24```Cpp22```Cpp
25aclnnStatus aclnnMaxPool3dWithArgmaxBackwardGetWorkspaceSize(23aclnnStatus aclnnMaxPool3dWithArgmaxBackwardGetWorkspaceSize(
26 const aclTensor *gradOutput,24 const aclTensor *gradOutput,
@@ -35,6 +33,7 @@ aclnnStatus aclnnMaxPool3dWithArgmaxBackwardGetWorkspaceSize(
35 uint64_t *workspaceSize,33 uint64_t *workspaceSize,
36 aclOpExecutor **executor)34 aclOpExecutor **executor)
37```35```
36+ 
38```Cpp37```Cpp
39aclnnStatus aclnnMaxPool3dWithArgmaxBackward(38aclnnStatus aclnnMaxPool3dWithArgmaxBackward(
40 void *workspace,39 void *workspace,
@@ -82,7 +81,7 @@ aclnnStatus aclnnMaxPool3dWithArgmaxBackward(
82 <td>输入</td>81 <td>输入</td>
83 <td>正向的输入数据。</td>82 <td>正向的输入数据。</td>
84 <td>-</td>83 <td>-</td>
85- <td>FLOAT、FLOAT16、BFLOAT16</td>84+ <td>FLOAT32、FLOAT16、BFLOAT16</td>
86 <td>ND</td>85 <td>ND</td>
87 <td>4-5</td>86 <td>4-5</td>
88 <td>√</td>87 <td>√</td>
@@ -92,7 +91,7 @@ aclnnStatus aclnnMaxPool3dWithArgmaxBackward(
92 <td>输入</td>91 <td>输入</td>
93 <td>正向输入中最大元素的索引位置。</td>92 <td>正向输入中最大元素的索引位置。</td>
94 <td>-</td>93 <td>-</td>
95- <td>INT32</td>94+ <td>INT32、INT64</td>
96 <td>ND</td>95 <td>ND</td>
97 <td>4-5</td>96 <td>4-5</td>
98 <td>√</td>97 <td>√</td>
@@ -178,6 +177,8 @@ aclnnStatus aclnnMaxPool3dWithArgmaxBackward(
178 <td>-</td>177 <td>-</td>
179 </tr>178 </tr>
180 </tbody></table>179 </tbody></table>
180+ 
181+ - <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>、<term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term>：`indices` 数据类型不支持INT64。depth * height * width 不支持大于 max int32。
181 182
182- **返回值：**183- **返回值：**
183 184 
@@ -236,9 +237,6 @@ aclnnStatus aclnnMaxPool3dWithArgmaxBackward(
236 <tr>237 <tr>
237 <td>平台不支持。</td>238 <td>平台不支持。</td>
238 </tr>239 </tr>
239- <tr>
240- <td>depth * height * width > max int32，超出了indices的表示范围。</td>
241- </tr>
242 </tbody>240 </tbody>
243 </table>241 </table>
244 242 
@@ -287,10 +285,6 @@ aclnnStatus aclnnMaxPool3dWithArgmaxBackward(
287- 确定性计算：285- 确定性计算：
288 - aclnnMaxPool3dWithArgmaxBackward默认非确定性实现，支持通过aclrtCtxSetSysParamOpt开启确定性。286 - aclnnMaxPool3dWithArgmaxBackward默认非确定性实现，支持通过aclrtCtxSetSysParamOpt开启确定性。
289 287 
290-- 功能维度：
291- - 数据类型支持：indices支持INT32。
292- - 数据格式支持：ND。
293- 
294- 未支持类型说明：288- 未支持类型说明：
295 - DOUBLE：指令不支持DOUBLE。289 - DOUBLE：指令不支持DOUBLE。
296 - 是否支持空tensor：不支持空进空出。290 - 是否支持空tensor：不支持空进空出。
Mpooling/max_pool3d_grad_with_argmax/op_api/aclnn_max_pool3d_with_argmax_backward.cpp+22-13
@@ -33,6 +33,8 @@ static const std::initializer_list<DataType> NULL_DTYPE_SUPPORT_LIST = {};
33static const std::initializer_list<DataType> GRAD_DTYPE_SUPPORT_LIST = {33static const std::initializer_list<DataType> GRAD_DTYPE_SUPPORT_LIST = {
34 DataType::DT_BF16, DataType::DT_FLOAT16, DataType::DT_FLOAT};34 DataType::DT_BF16, DataType::DT_FLOAT16, DataType::DT_FLOAT};
35static const std::initializer_list<op::DataType> INDICES_DTYPE_SUPPORT_LIST = {op::DataType::DT_INT32};35static const std::initializer_list<op::DataType> INDICES_DTYPE_SUPPORT_LIST = {op::DataType::DT_INT32};
36+static const std::initializer_list<op::DataType> INDICES_DTYPE_SUPPORT_LIST_950 = {
37+ op::DataType::DT_INT32, op::DataType::DT_INT64};
36 38 
37static const size_t CDHW_DIMS = 4;39static const size_t CDHW_DIMS = 4;
38static const size_t NCDHW_DIMS = 5;40static const size_t NCDHW_DIMS = 5;
@@ -64,6 +66,7 @@ static const std::initializer_list<op::DataType> GetDtypeSupportListBySocVersion
64 auto socVersion = GetCurrentPlatformInfo().GetSocVersion();66 auto socVersion = GetCurrentPlatformInfo().GetSocVersion();
65 switch (socVersion) {67 switch (socVersion) {
66 case SocVersion::ASCEND910B:68 case SocVersion::ASCEND910B:
69+ case SocVersion::ASCEND950:
67 case SocVersion::ASCEND910_93: {70 case SocVersion::ASCEND910_93: {
68 return GRAD_DTYPE_SUPPORT_LIST;71 return GRAD_DTYPE_SUPPORT_LIST;
69 }72 }
@@ -81,7 +84,11 @@ static bool CheckDtypeValid(
81{84{
82 auto dtypeSupportList = GetDtypeSupportListBySocVersion();85 auto dtypeSupportList = GetDtypeSupportListBySocVersion();
83 OP_CHECK_DTYPE_NOT_SUPPORT(self, dtypeSupportList, return false);86 OP_CHECK_DTYPE_NOT_SUPPORT(self, dtypeSupportList, return false);
84- OP_CHECK_DTYPE_NOT_SUPPORT(indices, INDICES_DTYPE_SUPPORT_LIST, return false);87+ if (GetCurrentPlatformInfo().GetSocVersion() == SocVersion::ASCEND950) {
88+ OP_CHECK_DTYPE_NOT_SUPPORT(indices, INDICES_DTYPE_SUPPORT_LIST_950, return false);
89+ } else {
90+ OP_CHECK_DTYPE_NOT_SUPPORT(indices, INDICES_DTYPE_SUPPORT_LIST, return false);
91+ }
85 OP_CHECK_DTYPE_NOT_SAME(self, gradOutput, return false);92 OP_CHECK_DTYPE_NOT_SAME(self, gradOutput, return false);
86 OP_CHECK_DTYPE_NOT_SAME(self, gradInput, return false);93 OP_CHECK_DTYPE_NOT_SAME(self, gradInput, return false);
87 return true;94 return true;
@@ -228,12 +235,16 @@ static bool CheckSelfShapeSupport(const aclTensor* self)
228 const auto& selfDimD = selfShape.GetDim(selfDimNum + D_DIM);235 const auto& selfDimD = selfShape.GetDim(selfDimNum + D_DIM);
229 236 
230 const int64_t selfSize = selfDimW * selfDimH * selfDimD;237 const int64_t selfSize = selfDimW * selfDimH * selfDimD;
231- OP_CHECK(238+ 
232- (selfSize <= MAX_INT32),239+ if (GetCurrentPlatformInfo().GetSocVersion() != SocVersion::ASCEND950) {
233- OP_LOGE(240+ OP_CHECK(
234- ACLNN_ERR_PARAM_INVALID, "The size of self should be less than or equal to 2^32 - 1, but got selfSize:%ld",241+ (selfSize <= MAX_INT32),
235- selfSize),242+ OP_LOGE(
236- return false);243+ ACLNN_ERR_PARAM_INVALID,
244+ "The size of self should be less than or equal to 2^32 - 1, but got selfSize:%ld", selfSize),
245+ return false);
246+ }
247+ 
237 return true;248 return true;
238}249}
239 250 
@@ -330,14 +341,12 @@ aclnnStatus aclnnMaxPool3dWithArgmaxBackwardGetWorkspaceSize(
330 l0op::ReFormat(selfContiguous, op::Format::FORMAT_NCDHW, uniqueExecutor.get());341 l0op::ReFormat(selfContiguous, op::Format::FORMAT_NCDHW, uniqueExecutor.get());
331 CHECK_RET(selfUnsqueezed != nullptr, ACLNN_ERR_INNER_NULLPTR);342 CHECK_RET(selfUnsqueezed != nullptr, ACLNN_ERR_INNER_NULLPTR);
332 343 
333- auto gradOutputUnsqueezed =344+ auto gradOutputUnsqueezed = isSelf4D ? View4Das5D(gradOutputContiguous, uniqueExecutor.get()) :
334- isSelf4D ? View4Das5D(gradOutputContiguous, uniqueExecutor.get()) :345+ l0op::ReFormat(gradOutputContiguous, op::Format::FORMAT_NCDHW, uniqueExecutor.get());
335- l0op::ReFormat(gradOutputContiguous, op::Format::FORMAT_NCDHW, uniqueExecutor.get());
336 CHECK_RET(gradOutputUnsqueezed != nullptr, ACLNN_ERR_INNER_NULLPTR);346 CHECK_RET(gradOutputUnsqueezed != nullptr, ACLNN_ERR_INNER_NULLPTR);
337 347 
338- auto indicesUnsqueezed = isSelf4D ?348+ auto indicesUnsqueezed = isSelf4D ? View4Das5D(indicesContiguous, uniqueExecutor.get()) :
339- View4Das5D(indicesContiguous, uniqueExecutor.get()) :349+ l0op::ReFormat(indicesContiguous, op::Format::FORMAT_NCDHW, uniqueExecutor.get());
340- l0op::ReFormat(indicesContiguous, op::Format::FORMAT_NCDHW, uniqueExecutor.get());
341 CHECK_RET(indicesUnsqueezed != nullptr, ACLNN_ERR_INNER_NULLPTR);350 CHECK_RET(indicesUnsqueezed != nullptr, ACLNN_ERR_INNER_NULLPTR);
342 351 
343 auto gradInputResult = l0op::MaxPool3DGradWithArgmax(352 auto gradInputResult = l0op::MaxPool3DGradWithArgmax(
Mpooling/max_pool3d_grad_with_argmax/op_api/max_pool3d_grad_with_argmax.cpp+10-3
@@ -35,6 +35,7 @@ static const inline std::initializer_list<op::DataType> GetDtypeSupportListBySoc
35 auto socVersion = GetCurrentPlatformInfo().GetSocVersion();35 auto socVersion = GetCurrentPlatformInfo().GetSocVersion();
36 switch (socVersion) {36 switch (socVersion) {
37 case SocVersion::ASCEND910_93:37 case SocVersion::ASCEND910_93:
38+ case SocVersion::ASCEND950:
38 case SocVersion::ASCEND910B: {39 case SocVersion::ASCEND910B: {
39 return GRAD_DTYPE_SUPPORT_LIST;40 return GRAD_DTYPE_SUPPORT_LIST;
40 }41 }
@@ -68,9 +69,15 @@ const inline aclTensor* MaxPool3DGradWithArgmaxAiCore(
68 L0_DFX(69 L0_DFX(
69 MaxPool3DGradWithArgmaxAiCore, gradOutput, self, indices, kernelSize, stride, padding, dilation, ceilMode,70 MaxPool3DGradWithArgmaxAiCore, gradOutput, self, indices, kernelSize, stride, padding, dilation, ceilMode,
70 gradInput);71 gradInput);
71- ADD_TO_LAUNCHER_LIST_AICORE(72+ if (GetCurrentPlatformInfo().GetSocVersion() == SocVersion::ASCEND950) {
72- MaxPool3DGradWithArgmax, OP_INPUT(self, gradOutput, indices), OP_OUTPUT(gradInput),73+ ADD_TO_LAUNCHER_LIST_AICORE(
73- OP_ATTR(kernelSize, stride, padding, dilation, ceilMode));74+ MaxPool3DGradWithArgmax, OP_INPUT(self, gradOutput, indices), OP_OUTPUT(gradInput),
75+ OP_ATTR(kernelSize, stride, padding, dilation, ceilMode, "NCDHW"));
76+ } else {
77+ ADD_TO_LAUNCHER_LIST_AICORE(
78+ MaxPool3DGradWithArgmax, OP_INPUT(self, gradOutput, indices), OP_OUTPUT(gradInput),
79+ OP_ATTR(kernelSize, stride, padding, dilation, ceilMode));
80+ }
74 return gradInput;81 return gradInput;
75}82}
76 83 
Apooling/max_pool3d_grad_with_argmax/op_graph/max_pool3d_grad_with_argmax_proto.h+85-0
@@ -0,0 +1,85 @@
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_pool3d_grad_with_argmax_proto.h
13+ * \brief
14+ */
15+
16+#ifndef OPS_BUILT_IN_OP_PROTO_INC_MAX_POOL3D_GRAD_WITH_ARGMAX_PROTO_H_
17+#define OPS_BUILT_IN_OP_PROTO_INC_MAX_POOL3D_GRAD_WITH_ARGMAX_PROTO_H_
18+ 
19+#include "graph/operator_reg.h"
20+#include "graph/operator.h"
21+ 
22+namespace ge {
23+ 
24+/**
25+* @brief Performs the backpropagation of MaxPool3DGradWithArgmax.
26+ 
27+* @par Inputs:
28+* Three inputs, including:
29+* @li x: An 5D Tensor. Supported type:float16, bfloat16, float32
30+* Must set the format, supported format list ["NCDHW, NDHWC"].
31+* @li grad: An 5D Tensor. Supported type:float16, bfloat16, float32
32+* Must set the format, supported format list ["NCDHW, NDHWC"].
33+* @li argmax: An 5D Tensor. Supported type:int32, int64
34+* Must set the format, supported format list ["NCDHW, NDHWC"]. \n
35+ 
36+* @par Attributes:
37+* @li ksize: A required list of int8, int16, int32, or int64 values,
38+* specifying the size of the window for each dimension (D/H/W) of the input tensor.
39+* @li strides: A required list of int8, int16, int32, or int64 values,
40+* specifying the strides of the sliding window for each dimension (D/H/W) of the input tensor.
41+* @li pads: A required list of int8, int16, int32, or int64 values,
42+* specifying the pads of the sliding window for each dimension of the input tensor.
43+* @li dilation: An optional list of int8, int16, int32, or int64 values,
44+* specifying the dilation of the sliding window for each dimension of the input tensor, default value [1, 1, 1].
45+* @li ceil_mode: An optional Boolean value. When true, will use ceil instead of floor
46+* in the formula to compute the output shape. Defaults to false.
47+* @li data_format: An optional string, supported values: ["NCDHW", "NDHWC"],
48+* default value: ["NCDHW"]. \n
49+ 
50+* @par Outputs:
51+* y: A Tensor. Has the same dtype, shape and format as input "x".
52+ 
53+* @attention Constraints:
54+* @li "ksize" is a list that has length 1 or 3(one value for each of D/H/W),
55+* every element in the list must be a numeric greater than 0.
56+* @li "strides" is a list that has length 0 or 1 or 3(one value for each of D/H/W),
CANN-robot
CANN-robotCANN-robot1月29日

逻辑运算与副作用: 注释中关于 strides 属性的约束描述存在逻辑矛盾。第59行注释说 strides 列表长度可以是0，并解释为“use default ksize for each of D/H/W”。然而，strides（步长）和 ksize（窗口大小）是两个不同的概念，用 ksize 作为 strides 的默认值在语义上不合理，可能导致运行时错误。

问题类型: 逻辑运算与副作用 文件路径: pooling/max_pool3d_grad_with_argmax/op_graph/max_pool3d_grad_with_argmax_proto.h 行号: 59 问题代码:

* @li "strides" is a list that has length 0 or 1 or 3(one value for each of D/H/W),
* length 0 means use default ksize for each of D/H/W,

修改建议:

修正注释：如果框架确实支持长度为0的 `strides`，应明确其实际默认行为（例如，步长默认为1）。建议修改为：`length 0 means use default stride (usually 1) for each dimension。` 并确保底层实现与此一致。

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

likedislike
小王！
小王！
1月31日 评论：
57+* length 0 means use default ksize for each of D/H/W,
58+* every element in the list must be a numeric greater than 0.
59+* @li "pads" is a list that has length 1 or 3(one value for each of D/H/W),
CANN-robot
CANN-robotCANN-robot1月29日

逻辑运算与副作用: 注释中关于 pads 属性的约束描述不完整。第62行要求 pads 的每个元素必须大于0，但填充（padding）可以为0（表示无填充）。要求大于0会错误地排除有效的无填充情况。

问题类型: 逻辑运算与副作用 文件路径: pooling/max_pool3d_grad_with_argmax/op_graph/max_pool3d_grad_with_argmax_proto.h 行号: 62 问题代码:

* @li "pads" is a list that has length 1 or 3(one value for each of D/H/W),
* every element in the list must be a numeric greater than 0.

修改建议:

将约束修改为：`every element in the list must be a numeric greater than or equal to 0。` 以允许零填充。

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

likedislike
小王！
小王！
1月31日 评论：
60+* every element in the list must be a numeric greater than or equal to 0.
61+* additionally, two strict size limits apply: \n
62+* 1. Each pads value (pD/pH/pW for D/H/W dimensions) must be less than or equal to half of the corresponding "ksize" value (kD/kH/kW / 2). \n
63+* 2. Each pads value (pD/pH/pW) must be less than or equal to ((corresponding ksize - 1) * corresponding dilation + 1) / 2
64+* (i.e., pD ≤ ((kD - 1) * dD + 1) / 2, pH ≤ ((kH - 1) * dH + 1) / 2, pW ≤ ((kW - 1) * dW + 1) / 2).
65+* @li "dilation" is a list that has length 1 or 3(one value for each of D/H/W),
66+* every element in the list must be a numeric greater than 0.
67+* @li "ceil_mode": This operator currently supports ceil_mode = false or true.
68+ 
69+* @par Third-party framework compatibility
70+* Compatible with the Torch operator MaxPool3DGradWithArgmax.
71+*/
72+REG_OP(MaxPool3DGradWithArgmax)
73+ .INPUT(x, TensorType({DT_FLOAT16, DT_FLOAT32, DT_BF16}))
74+ .INPUT(grad, TensorType({DT_FLOAT16, DT_FLOAT32, DT_BF16}))
75+ .INPUT(argmax, TensorType({DT_INT32, DT_INT64}))
76+ .OUTPUT(y, TensorType({DT_FLOAT16, DT_FLOAT32, DT_BF16}))
77+ .REQUIRED_ATTR(ksize, ListInt)
78+ .REQUIRED_ATTR(strides, ListInt)
79+ .REQUIRED_ATTR(pads, ListInt)
80+ .ATTR(dilation, ListInt, {1, 1, 1})
81+ .ATTR(ceil_mode, Bool, false)
82+ .ATTR(data_format, String, "NCDHW")
83+ .OP_END_FACTORY_REG(MaxPool3DGradWithArgmax)
84+} // namespace ge
85+#endif // OPS_BUILT_IN_OP_PROTO_INC_NN_POOLING_OPS_H
Mpooling/max_pool3d_grad_with_argmax/op_host/config/ascend910_93/max_pool3d_grad_with_argmax_binary.json+265-250
@@ -1,251 +1,266 @@
1-{1+{
2- "op_type": "MaxPool3DGradWithArgmax",2+ "op_type": "MaxPool3DGradWithArgmax",
3- "op_list": [3+ "op_list": [
4- {4+ {
5- "bin_filename": "MaxPool3DGradWithArgmax_fp32",5+ "bin_filename": "MaxPool3DGradWithArgmax_fp32",
6- "inputs": [6+ "inputs": [
7- {7+ {
8- "name": "x",8+ "name": "x",
9- "index": 0,9+ "index": 0,
10- "dtype": "float32",10+ "dtype": "float32",
11- "format": "NCDHW",11+ "format": "NCDHW",
12- "paramType": "required",12+ "paramType": "required",
13- "shape": [13+ "shape": [
14- -214+ -2
15- ],15+ ],
16- "format_match_mode": "FormatFixed"16+ "format_match_mode": "FormatFixed"
17- },17+ },
18- {18+ {
19- "name": "grad",19+ "name": "grad",
20- "index": 1,20+ "index": 1,
21- "dtype": "float32",21+ "dtype": "float32",
22- "format": "NCDHW",22+ "format": "NCDHW",
23- "paramType": "required",23+ "paramType": "required",
24- "shape": [24+ "shape": [
25- -225+ -2
26- ],26+ ],
27- "format_match_mode": "FormatFixed"27+ "format_match_mode": "FormatFixed"
28- },28+ },
29- {29+ {
30- "name": "argmax",30+ "name": "argmax",
31- "index": 2,31+ "index": 2,
32- "dtype": "int32",32+ "dtype": "int32",
33- "format": "NCDHW",33+ "format": "NCDHW",
34- "paramType": "required",34+ "paramType": "required",
35- "shape": [35+ "shape": [
36- -236+ -2
37- ],37+ ],
38- "format_match_mode": "FormatFixed"38+ "format_match_mode": "FormatFixed"
39- }39+ }
40- ],40+ ],
41- "outputs": [41+ "outputs": [
42- {42+ {
43- "name": "y",43+ "name": "y",
44- "index": 0,44+ "index": 0,
45- "dtype": "float32",45+ "dtype": "float32",
46- "format": "NCDHW",46+ "format": "NCDHW",
47- "paramType": "required",47+ "paramType": "required",
48- "shape": [48+ "shape": [
49- -249+ -2
50- ],50+ ],
51- "format_match_mode": "FormatFixed"51+ "format_match_mode": "FormatFixed"
52- }52+ }
53- ],53+ ],
54- "attrs": [54+ "attrs": [
55- {55+ {
56- "name": "ksize",56+ "name": "ksize",
57- "dtype": "list_int",57+ "dtype": "list_int",
58- "value": null58+ "value": null
59- },59+ },
60- {60+ {
61- "name": "strides",61+ "name": "strides",
62- "dtype": "list_int",62+ "dtype": "list_int",
63- "value": null63+ "value": null
64- },64+ },
65- {65+ {
66- "name": "pads",66+ "name": "pads",
67- "dtype": "list_int",67+ "dtype": "list_int",
68- "value": null68+ "value": null
69- },69+ },
70- {70+ {
71- "name": "dilation",71+ "name": "dilation",
72- "dtype": "list_int",72+ "dtype": "list_int",
73- "value": [73+ "value": [
74- 1,74+ 1,
75- 1,75+ 1,
76- 176+ 1
77- ]77+ ]
78- },78+ },
79- {79+ {
80- "name": "ceil_mode",80+ "name": "ceil_mode",
81- "dtype": "bool",81+ "dtype": "bool",
82- "value": null82+ "value": null
83- }83+ },
84- ]84+ {
85- },85+ "name": "data_format",
86- {86+ "dtype": "string",
87- "bin_filename": "MaxPool3DGradWithArgmax_fp16",87+ "value": "NCDHW"
88- "inputs": [88+ }
89- {89+ ]
90- "name": "x",90+ },
91- "index": 0,91+ {
92- "dtype": "float16",92+ "bin_filename": "MaxPool3DGradWithArgmax_fp16",
93- "format": "NCDHW",93+ "inputs": [
94- "paramType": "required",94+ {
95- "shape": [95+ "name": "x",
96- -296+ "index": 0,
97- ],97+ "dtype": "float16",
98- "format_match_mode": "FormatFixed"98+ "format": "NCDHW",
99- },99+ "paramType": "required",
100- {100+ "shape": [
101- "name": "grad",101+ -2
102- "index": 1,102+ ],
103- "dtype": "float16",103+ "format_match_mode": "FormatFixed"
104- "format": "NCDHW",104+ },
105- "paramType": "required",105+ {
106- "shape": [106+ "name": "grad",
107- -2107+ "index": 1,
108- ],108+ "dtype": "float16",
109- "format_match_mode": "FormatFixed"109+ "format": "NCDHW",
110- },110+ "paramType": "required",
111- {111+ "shape": [
112- "name": "argmax",112+ -2
113- "index": 2,113+ ],
114- "dtype": "int32",114+ "format_match_mode": "FormatFixed"
115- "format": "NCDHW",115+ },
116- "paramType": "required",116+ {
117- "shape": [117+ "name": "argmax",
118- -2118+ "index": 2,
119- ],119+ "dtype": "int32",
120- "format_match_mode": "FormatFixed"120+ "format": "NCDHW",
121- }121+ "paramType": "required",
122- ],122+ "shape": [
123- "outputs": [123+ -2
124- {124+ ],
125- "name": "y",125+ "format_match_mode": "FormatFixed"
126- "index": 0,126+ }
127- "dtype": "float16",127+ ],
128- "format": "NCDHW",128+ "outputs": [
129- "paramType": "required",129+ {
130- "shape": [130+ "name": "y",
131- -2131+ "index": 0,
132- ],132+ "dtype": "float16",
133- "format_match_mode": "FormatFixed"133+ "format": "NCDHW",
134- }134+ "paramType": "required",
135- ],135+ "shape": [
136- "attrs": [136+ -2
137- {137+ ],
138- "name": "ksize",138+ "format_match_mode": "FormatFixed"
139- "dtype": "list_int",139+ }
140- "value": null140+ ],
141- },141+ "attrs": [
142- {142+ {
143- "name": "strides",143+ "name": "ksize",
144- "dtype": "list_int",144+ "dtype": "list_int",
145- "value": null145+ "value": null
146- },146+ },
147- {147+ {
148- "name": "pads",148+ "name": "strides",
149- "dtype": "list_int",149+ "dtype": "list_int",
150- "value": null150+ "value": null
151- },151+ },
152- {152+ {
153- "name": "dilation",153+ "name": "pads",
154- "dtype": "list_int",154+ "dtype": "list_int",
155- "value": [155+ "value": null
156- 1,156+ },
157- 1,157+ {
158- 1158+ "name": "dilation",
159- ]159+ "dtype": "list_int",
160- },160+ "value": [
161- {161+ 1,
162- "name": "ceil_mode",162+ 1,
163- "dtype": "bool",163+ 1
164- "value": null164+ ]
165- }165+ },
166- ]166+ {
167- },167+ "name": "ceil_mode",
168- {168+ "dtype": "bool",
169- "bin_filename": "MaxPool3DGradWithArgmax_bf16",169+ "value": null
170- "inputs": [170+ },
171- {171+ {
172- "name": "x",172+ "name": "data_format",
173- "index": 0,173+ "dtype": "string",
174- "dtype": "bfloat16",174+ "value": "NCDHW"
175- "format": "NCDHW",175+ }
176- "paramType": "required",176+ ]
177- "shape": [177+ },
178- -2178+ {
179- ],179+ "bin_filename": "MaxPool3DGradWithArgmax_bf16",
180- "format_match_mode": "FormatFixed"180+ "inputs": [
181- },181+ {
182- {182+ "name": "x",
183- "name": "grad",183+ "index": 0,
184- "index": 1,184+ "dtype": "bfloat16",
185- "dtype": "bfloat16",185+ "format": "NCDHW",
186- "format": "NCDHW",186+ "paramType": "required",
187- "paramType": "required",187+ "shape": [
188- "shape": [188+ -2
189- -2189+ ],
190- ],190+ "format_match_mode": "FormatFixed"
191- "format_match_mode": "FormatFixed"191+ },
192- },192+ {
193- {193+ "name": "grad",
194- "name": "argmax",194+ "index": 1,
195- "index": 2,195+ "dtype": "bfloat16",
196- "dtype": "int32",196+ "format": "NCDHW",
197- "format": "NCDHW",197+ "paramType": "required",
198- "paramType": "required",198+ "shape": [
199- "shape": [199+ -2
200- -2200+ ],
201- ],201+ "format_match_mode": "FormatFixed"
202- "format_match_mode": "FormatFixed"202+ },
203- }203+ {
204- ],204+ "name": "argmax",
205- "outputs": [205+ "index": 2,
206- {206+ "dtype": "int32",
207- "name": "y",207+ "format": "NCDHW",
208- "index": 0,208+ "paramType": "required",
209- "dtype": "bfloat16",209+ "shape": [
210- "format": "NCDHW",210+ -2
211- "paramType": "required",211+ ],
212- "shape": [212+ "format_match_mode": "FormatFixed"
213- -2213+ }
214- ],214+ ],
215- "format_match_mode": "FormatFixed"215+ "outputs": [
216- }216+ {
217- ],217+ "name": "y",
218- "attrs": [218+ "index": 0,
219- {219+ "dtype": "bfloat16",
220- "name": "ksize",220+ "format": "NCDHW",
221- "dtype": "list_int",221+ "paramType": "required",
222- "value": null222+ "shape": [
223- },223+ -2
224- {224+ ],
225- "name": "strides",225+ "format_match_mode": "FormatFixed"
226- "dtype": "list_int",226+ }
227- "value": null227+ ],
228- },228+ "attrs": [
229- {229+ {
230- "name": "pads",230+ "name": "ksize",
231- "dtype": "list_int",231+ "dtype": "list_int",
232- "value": null232+ "value": null
233- },233+ },
234- {234+ {
235- "name": "dilation",235+ "name": "strides",
236- "dtype": "list_int",236+ "dtype": "list_int",
237- "value": [237+ "value": null
238- 1,238+ },
239- 1,239+ {
240- 1240+ "name": "pads",
241- ]241+ "dtype": "list_int",
242- },242+ "value": null
243- {243+ },
244- "name": "ceil_mode",244+ {
245- "dtype": "bool",245+ "name": "dilation",
246- "value": null246+ "dtype": "list_int",
247- }247+ "value": [
248- ]248+ 1,
249- }249+ 1,
250- ]250+ 1
251+ ]
252+ },
253+ {
254+ "name": "ceil_mode",
255+ "dtype": "bool",
256+ "value": null
257+ },
258+ {
259+ "name": "data_format",
260+ "dtype": "string",
261+ "value": "NCDHW"
262+ }
263+ ]
264+ }
265+ ]
251}266}
Mpooling/max_pool3d_grad_with_argmax/op_host/config/ascend910b/max_pool3d_grad_with_argmax_binary.json+265-250
@@ -1,251 +1,266 @@
1-{1+{
2- "op_type": "MaxPool3DGradWithArgmax",2+ "op_type": "MaxPool3DGradWithArgmax",
3- "op_list": [3+ "op_list": [
4- {4+ {
5- "bin_filename": "MaxPool3DGradWithArgmax_fp32",5+ "bin_filename": "MaxPool3DGradWithArgmax_fp32",
6- "inputs": [6+ "inputs": [
7- {7+ {
8- "name": "x",8+ "name": "x",
9- "index": 0,9+ "index": 0,
10- "dtype": "float32",10+ "dtype": "float32",
11- "format": "NCDHW",11+ "format": "NCDHW",
12- "paramType": "required",12+ "paramType": "required",
13- "shape": [13+ "shape": [
14- -214+ -2
15- ],15+ ],
16- "format_match_mode": "FormatFixed"16+ "format_match_mode": "FormatFixed"
17- },17+ },
18- {18+ {
19- "name": "grad",19+ "name": "grad",
20- "index": 1,20+ "index": 1,
21- "dtype": "float32",21+ "dtype": "float32",
22- "format": "NCDHW",22+ "format": "NCDHW",
23- "paramType": "required",23+ "paramType": "required",
24- "shape": [24+ "shape": [
25- -225+ -2
26- ],26+ ],
27- "format_match_mode": "FormatFixed"27+ "format_match_mode": "FormatFixed"
28- },28+ },
29- {29+ {
30- "name": "argmax",30+ "name": "argmax",
31- "index": 2,31+ "index": 2,
32- "dtype": "int32",32+ "dtype": "int32",
33- "format": "NCDHW",33+ "format": "NCDHW",
34- "paramType": "required",34+ "paramType": "required",
35- "shape": [35+ "shape": [
36- -236+ -2
37- ],37+ ],
38- "format_match_mode": "FormatFixed"38+ "format_match_mode": "FormatFixed"
39- }39+ }
40- ],40+ ],
41- "outputs": [41+ "outputs": [
42- {42+ {
43- "name": "y",43+ "name": "y",
44- "index": 0,44+ "index": 0,
45- "dtype": "float32",45+ "dtype": "float32",
46- "format": "NCDHW",46+ "format": "NCDHW",
47- "paramType": "required",47+ "paramType": "required",
48- "shape": [48+ "shape": [
49- -249+ -2
50- ],50+ ],
51- "format_match_mode": "FormatFixed"51+ "format_match_mode": "FormatFixed"
52- }52+ }
53- ],53+ ],
54- "attrs": [54+ "attrs": [
55- {55+ {
56- "name": "ksize",56+ "name": "ksize",
57- "dtype": "list_int",57+ "dtype": "list_int",
58- "value": null58+ "value": null
59- },59+ },
60- {60+ {
61- "name": "strides",61+ "name": "strides",
62- "dtype": "list_int",62+ "dtype": "list_int",
63- "value": null63+ "value": null
64- },64+ },
65- {65+ {
66- "name": "pads",66+ "name": "pads",
67- "dtype": "list_int",67+ "dtype": "list_int",
68- "value": null68+ "value": null
69- },69+ },
70- {70+ {
71- "name": "dilation",71+ "name": "dilation",
72- "dtype": "list_int",72+ "dtype": "list_int",
73- "value": [73+ "value": [
74- 1,74+ 1,
75- 1,75+ 1,
76- 176+ 1
77- ]77+ ]
78- },78+ },
79- {79+ {
80- "name": "ceil_mode",80+ "name": "ceil_mode",
81- "dtype": "bool",81+ "dtype": "bool",
82- "value": null82+ "value": null
83- }83+ },
84- ]84+ {
85- },85+ "name": "data_format",
86- {86+ "dtype": "string",
87- "bin_filename": "MaxPool3DGradWithArgmax_fp16",87+ "value": "NCDHW"
88- "inputs": [88+ }
89- {89+ ]
90- "name": "x",90+ },
91- "index": 0,91+ {
92- "dtype": "float16",92+ "bin_filename": "MaxPool3DGradWithArgmax_fp16",
93- "format": "NCDHW",93+ "inputs": [
94- "paramType": "required",94+ {
95- "shape": [95+ "name": "x",
96- -296+ "index": 0,
97- ],97+ "dtype": "float16",
98- "format_match_mode": "FormatFixed"98+ "format": "NCDHW",
99- },99+ "paramType": "required",
100- {100+ "shape": [
101- "name": "grad",101+ -2
102- "index": 1,102+ ],
103- "dtype": "float16",103+ "format_match_mode": "FormatFixed"
104- "format": "NCDHW",104+ },
105- "paramType": "required",105+ {
106- "shape": [106+ "name": "grad",
107- -2107+ "index": 1,
108- ],108+ "dtype": "float16",
109- "format_match_mode": "FormatFixed"109+ "format": "NCDHW",
110- },110+ "paramType": "required",
111- {111+ "shape": [
112- "name": "argmax",112+ -2
113- "index": 2,113+ ],
114- "dtype": "int32",114+ "format_match_mode": "FormatFixed"
115- "format": "NCDHW",115+ },
116- "paramType": "required",116+ {
117- "shape": [117+ "name": "argmax",
118- -2118+ "index": 2,
119- ],119+ "dtype": "int32",
120- "format_match_mode": "FormatFixed"120+ "format": "NCDHW",
121- }121+ "paramType": "required",
122- ],122+ "shape": [
123- "outputs": [123+ -2
124- {124+ ],
125- "name": "y",125+ "format_match_mode": "FormatFixed"
126- "index": 0,126+ }
127- "dtype": "float16",127+ ],
128- "format": "NCDHW",128+ "outputs": [
129- "paramType": "required",129+ {
130- "shape": [130+ "name": "y",
131- -2131+ "index": 0,
132- ],132+ "dtype": "float16",
133- "format_match_mode": "FormatFixed"133+ "format": "NCDHW",
134- }134+ "paramType": "required",
135- ],135+ "shape": [
136- "attrs": [136+ -2
137- {137+ ],
138- "name": "ksize",138+ "format_match_mode": "FormatFixed"
139- "dtype": "list_int",139+ }
140- "value": null140+ ],
141- },141+ "attrs": [
142- {142+ {
143- "name": "strides",143+ "name": "ksize",
144- "dtype": "list_int",144+ "dtype": "list_int",
145- "value": null145+ "value": null
146- },146+ },
147- {147+ {
148- "name": "pads",148+ "name": "strides",
149- "dtype": "list_int",149+ "dtype": "list_int",
150- "value": null150+ "value": null
151- },151+ },
152- {152+ {
153- "name": "dilation",153+ "name": "pads",
154- "dtype": "list_int",154+ "dtype": "list_int",
155- "value": [155+ "value": null
156- 1,156+ },
157- 1,157+ {
158- 1158+ "name": "dilation",
159- ]159+ "dtype": "list_int",
160- },160+ "value": [
161- {161+ 1,
162- "name": "ceil_mode",162+ 1,
163- "dtype": "bool",163+ 1
164- "value": null164+ ]
165- }165+ },
166- ]166+ {
167- },167+ "name": "ceil_mode",
168- {168+ "dtype": "bool",
169- "bin_filename": "MaxPool3DGradWithArgmax_bf16",169+ "value": null
170- "inputs": [170+ },
171- {171+ {
172- "name": "x",172+ "name": "data_format",
173- "index": 0,173+ "dtype": "string",
174- "dtype": "bfloat16",174+ "value": "NCDHW"
175- "format": "NCDHW",175+ }
176- "paramType": "required",176+ ]
177- "shape": [177+ },
178- -2178+ {
179- ],179+ "bin_filename": "MaxPool3DGradWithArgmax_bf16",
180- "format_match_mode": "FormatFixed"180+ "inputs": [
181- },181+ {
182- {182+ "name": "x",
183- "name": "grad",183+ "index": 0,
184- "index": 1,184+ "dtype": "bfloat16",
185- "dtype": "bfloat16",185+ "format": "NCDHW",
186- "format": "NCDHW",186+ "paramType": "required",
187- "paramType": "required",187+ "shape": [
188- "shape": [188+ -2
189- -2189+ ],
190- ],190+ "format_match_mode": "FormatFixed"
191- "format_match_mode": "FormatFixed"191+ },
192- },192+ {
193- {193+ "name": "grad",
194- "name": "argmax",194+ "index": 1,
195- "index": 2,195+ "dtype": "bfloat16",
196- "dtype": "int32",196+ "format": "NCDHW",
197- "format": "NCDHW",197+ "paramType": "required",
198- "paramType": "required",198+ "shape": [
199- "shape": [199+ -2
200- -2200+ ],
201- ],201+ "format_match_mode": "FormatFixed"
202- "format_match_mode": "FormatFixed"202+ },
203- }203+ {
204- ],204+ "name": "argmax",
205- "outputs": [205+ "index": 2,
206- {206+ "dtype": "int32",
207- "name": "y",207+ "format": "NCDHW",
208- "index": 0,208+ "paramType": "required",
209- "dtype": "bfloat16",209+ "shape": [
210- "format": "NCDHW",210+ -2
211- "paramType": "required",211+ ],
212- "shape": [212+ "format_match_mode": "FormatFixed"
213- -2213+ }
214- ],214+ ],
215- "format_match_mode": "FormatFixed"215+ "outputs": [
216- }216+ {
217- ],217+ "name": "y",
218- "attrs": [218+ "index": 0,
219- {219+ "dtype": "bfloat16",
220- "name": "ksize",220+ "format": "NCDHW",
221- "dtype": "list_int",221+ "paramType": "required",
222- "value": null222+ "shape": [
223- },223+ -2
224- {224+ ],
225- "name": "strides",225+ "format_match_mode": "FormatFixed"
226- "dtype": "list_int",226+ }
227- "value": null227+ ],
228- },228+ "attrs": [
229- {229+ {
230- "name": "pads",230+ "name": "ksize",
231- "dtype": "list_int",231+ "dtype": "list_int",
232- "value": null232+ "value": null
233- },233+ },
234- {234+ {
235- "name": "dilation",235+ "name": "strides",
236- "dtype": "list_int",236+ "dtype": "list_int",
237- "value": [237+ "value": null
238- 1,238+ },
239- 1,239+ {
240- 1240+ "name": "pads",
241- ]241+ "dtype": "list_int",
242- },242+ "value": null
243- {243+ },
244- "name": "ceil_mode",244+ {
245- "dtype": "bool",245+ "name": "dilation",
246- "value": null246+ "dtype": "list_int",
247- }247+ "value": [
248- ]248+ 1,
249- }249+ 1,
250- ]250+ 1
251+ ]
252+ },
253+ {
254+ "name": "ceil_mode",
255+ "dtype": "bool",
256+ "value": null
257+ },
258+ {
259+ "name": "data_format",
260+ "dtype": "string",
261+ "value": "NCDHW"
262+ }
263+ ]
264+ }
265+ ]
251}266}
Apooling/max_pool3d_grad_with_argmax/op_host/config/ascend950/max_pool3d_grad_with_argmax_binary.json+527-0
@@ -0,0 +1,527 @@
1+{
2+ "op_type": "MaxPool3DGradWithArgmax",
3+ "op_list": [
4+ {
5+ "bin_filename": "MaxPool3DGradWithArgmax_fp32_int32",
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+ "name": "grad",
20+ "index": 1,
21+ "dtype": "float32",
22+ "format": "ND",
23+ "paramType": "required",
24+ "shape": [
25+ -2
26+ ],
27+ "format_match_mode": "FormatAgnostic"
28+ },
29+ {
30+ "name": "argmax",
31+ "index": 2,
32+ "dtype": "int32",
33+ "format": "ND",
34+ "paramType": "required",
35+ "shape": [
36+ -2
37+ ],
38+ "format_match_mode": "FormatAgnostic"
39+ }
40+ ],
41+ "outputs": [
42+ {
43+ "name": "y",
44+ "index": 0,
45+ "dtype": "float32",
46+ "format": "ND",
47+ "paramType": "required",
48+ "shape": [
49+ -2
50+ ],
51+ "format_match_mode": "FormatAgnostic"
52+ }
53+ ],
54+ "attrs": [
55+ {
56+ "name": "ksize",
57+ "dtype": "list_int",
58+ "value": null
59+ },
60+ {
61+ "name": "strides",
62+ "dtype": "list_int",
63+ "value": null
64+ },
65+ {
66+ "name": "pads",
67+ "dtype": "list_int",
68+ "value": null
69+ },
70+ {
71+ "name": "dilation",
72+ "dtype": "list_int",
73+ "value": [
74+ 1,
75+ 1,
76+ 1
77+ ]
78+ },
79+ {
80+ "name": "ceil_mode",
81+ "dtype": "bool",
82+ "value": null
83+ },
84+ {
85+ "name": "data_format",
86+ "dtype": "string",
87+ "value": "NCDHW"
88+ }
89+ ]
90+ },
91+ {
92+ "bin_filename": "MaxPool3DGradWithArgmax_fp16_int32",
93+ "inputs": [
94+ {
95+ "name": "x",
96+ "index": 0,
97+ "dtype": "float16",
98+ "format": "ND",
99+ "paramType": "required",
100+ "shape": [
101+ -2
102+ ],
103+ "format_match_mode": "FormatAgnostic"
104+ },
105+ {
106+ "name": "grad",
107+ "index": 1,
108+ "dtype": "float16",
109+ "format": "ND",
110+ "paramType": "required",
111+ "shape": [
112+ -2
113+ ],
114+ "format_match_mode": "FormatAgnostic"
115+ },
116+ {
117+ "name": "argmax",
118+ "index": 2,
119+ "dtype": "int32",
120+ "format": "ND",
121+ "paramType": "required",
122+ "shape": [
123+ -2
124+ ],
125+ "format_match_mode": "FormatAgnostic"
126+ }
127+ ],
128+ "outputs": [
129+ {
130+ "name": "y",
131+ "index": 0,
132+ "dtype": "float16",
133+ "format": "ND",
134+ "paramType": "required",
135+ "shape": [
136+ -2
137+ ],
138+ "format_match_mode": "FormatAgnostic"
139+ }
140+ ],
141+ "attrs": [
142+ {
143+ "name": "ksize",
144+ "dtype": "list_int",
145+ "value": null
146+ },
147+ {
148+ "name": "strides",
149+ "dtype": "list_int",
150+ "value": null
151+ },
152+ {
153+ "name": "pads",
154+ "dtype": "list_int",
155+ "value": null
156+ },
157+ {
158+ "name": "dilation",
159+ "dtype": "list_int",
160+ "value": [
161+ 1,
162+ 1,
163+ 1
164+ ]
165+ },
166+ {
167+ "name": "ceil_mode",
168+ "dtype": "bool",
169+ "value": null
170+ },
171+ {
172+ "name": "data_format",
173+ "dtype": "string",
174+ "value": "NCDHW"
175+ }
176+ ]
177+ },
178+ {
179+ "bin_filename": "MaxPool3DGradWithArgmax_bf16_int32",
180+ "inputs": [
181+ {
182+ "name": "x",
183+ "index": 0,
184+ "dtype": "bfloat16",
185+ "format": "ND",
186+ "paramType": "required",
187+ "shape": [
188+ -2
189+ ],
190+ "format_match_mode": "FormatAgnostic"
191+ },
192+ {
193+ "name": "grad",
194+ "index": 1,
195+ "dtype": "bfloat16",
196+ "format": "ND",
197+ "paramType": "required",
198+ "shape": [
199+ -2
200+ ],
201+ "format_match_mode": "FormatAgnostic"
202+ },
203+ {
204+ "name": "argmax",
205+ "index": 2,
206+ "dtype": "int32",
207+ "format": "ND",
208+ "paramType": "required",
209+ "shape": [
210+ -2
211+ ],
212+ "format_match_mode": "FormatAgnostic"
213+ }
214+ ],
215+ "outputs": [
216+ {
217+ "name": "y",
218+ "index": 0,
219+ "dtype": "bfloat16",
220+ "format": "ND",
221+ "paramType": "required",
222+ "shape": [
223+ -2
224+ ],
225+ "format_match_mode": "FormatAgnostic"
226+ }
227+ ],
228+ "attrs": [
229+ {
230+ "name": "ksize",
231+ "dtype": "list_int",
232+ "value": null
233+ },
234+ {
235+ "name": "strides",
236+ "dtype": "list_int",
237+ "value": null
238+ },
239+ {
240+ "name": "pads",
241+ "dtype": "list_int",
242+ "value": null
243+ },
244+ {
245+ "name": "dilation",
246+ "dtype": "list_int",
247+ "value": [
248+ 1,
249+ 1,
250+ 1
251+ ]
252+ },
253+ {
254+ "name": "ceil_mode",
255+ "dtype": "bool",
256+ "value": null
257+ },
258+ {
259+ "name": "data_format",
260+ "dtype": "string",
261+ "value": "NCDHW"
262+ }
263+ ]
264+ },
265+ {
266+ "bin_filename": "MaxPool3DGradWithArgmax_fp32_int64",
267+ "inputs": [
268+ {
269+ "name": "x",
270+ "index": 0,
271+ "dtype": "float32",
272+ "format": "ND",
273+ "paramType": "required",
274+ "shape": [
275+ -2
276+ ],
277+ "format_match_mode": "FormatAgnostic"
278+ },
279+ {
280+ "name": "grad",
281+ "index": 1,
282+ "dtype": "float32",
283+ "format": "ND",
284+ "paramType": "required",
285+ "shape": [
286+ -2
287+ ],
288+ "format_match_mode": "FormatAgnostic"
289+ },
290+ {
291+ "name": "argmax",
292+ "index": 2,
293+ "dtype": "int64",
294+ "format": "ND",
295+ "paramType": "required",
296+ "shape": [
297+ -2
298+ ],
299+ "format_match_mode": "FormatAgnostic"
300+ }
301+ ],
302+ "outputs": [
303+ {
304+ "name": "y",
305+ "index": 0,
306+ "dtype": "float32",
307+ "format": "ND",
308+ "paramType": "required",
309+ "shape": [
310+ -2
311+ ],
312+ "format_match_mode": "FormatAgnostic"
313+ }
314+ ],
315+ "attrs": [
316+ {
317+ "name": "ksize",
318+ "dtype": "list_int",
319+ "value": null
320+ },
321+ {
322+ "name": "strides",
323+ "dtype": "list_int",
324+ "value": null
325+ },
326+ {
327+ "name": "pads",
328+ "dtype": "list_int",
329+ "value": null
330+ },
331+ {
332+ "name": "dilation",
333+ "dtype": "list_int",
334+ "value": [
335+ 1,
336+ 1,
337+ 1
338+ ]
339+ },
340+ {
341+ "name": "ceil_mode",
342+ "dtype": "bool",
343+ "value": null
344+ },
345+ {
346+ "name": "data_format",
347+ "dtype": "string",
348+ "value": "NCDHW"
349+ }
350+ ]
351+ },
352+ {
353+ "bin_filename": "MaxPool3DGradWithArgmax_fp16_int64",
354+ "inputs": [
355+ {
356+ "name": "x",
357+ "index": 0,
358+ "dtype": "float16",
359+ "format": "ND",
360+ "paramType": "required",
361+ "shape": [
362+ -2
363+ ],
364+ "format_match_mode": "FormatAgnostic"
365+ },
366+ {
367+ "name": "grad",
368+ "index": 1,
369+ "dtype": "float16",
370+ "format": "ND",
371+ "paramType": "required",
372+ "shape": [
373+ -2
374+ ],
375+ "format_match_mode": "FormatAgnostic"
376+ },
377+ {
378+ "name": "argmax",
379+ "index": 2,
380+ "dtype": "int64",
381+ "format": "ND",
382+ "paramType": "required",
383+ "shape": [
384+ -2
385+ ],
386+ "format_match_mode": "FormatAgnostic"
387+ }
388+ ],
389+ "outputs": [
390+ {
391+ "name": "y",
392+ "index": 0,
393+ "dtype": "float16",
394+ "format": "ND",
395+ "paramType": "required",
396+ "shape": [
397+ -2
398+ ],
399+ "format_match_mode": "FormatAgnostic"
400+ }
401+ ],
402+ "attrs": [
403+ {
404+ "name": "ksize",
405+ "dtype": "list_int",
406+ "value": null
407+ },
408+ {
409+ "name": "strides",
410+ "dtype": "list_int",
411+ "value": null
412+ },
413+ {
414+ "name": "pads",
415+ "dtype": "list_int",
416+ "value": null
417+ },
418+ {
419+ "name": "dilation",
420+ "dtype": "list_int",
421+ "value": [
422+ 1,
423+ 1,
424+ 1
425+ ]
426+ },
427+ {
428+ "name": "ceil_mode",
429+ "dtype": "bool",
430+ "value": null
431+ },
432+ {
433+ "name": "data_format",
434+ "dtype": "string",
435+ "value": "NCDHW"
436+ }
437+ ]
438+ },
439+ {
440+ "bin_filename": "MaxPool3DGradWithArgmax_bf16_int64",
441+ "inputs": [
442+ {
443+ "name": "x",
444+ "index": 0,
445+ "dtype": "bfloat16",
446+ "format": "ND",
447+ "paramType": "required",
448+ "shape": [
449+ -2
450+ ],
451+ "format_match_mode": "FormatAgnostic"
452+ },
453+ {
454+ "name": "grad",
455+ "index": 1,
456+ "dtype": "bfloat16",
457+ "format": "ND",
458+ "paramType": "required",
459+ "shape": [
460+ -2
461+ ],
462+ "format_match_mode": "FormatAgnostic"
463+ },
464+ {
465+ "name": "argmax",
466+ "index": 2,
467+ "dtype": "int64",
468+ "format": "ND",
469+ "paramType": "required",
470+ "shape": [
471+ -2
472+ ],
473+ "format_match_mode": "FormatAgnostic"
474+ }
475+ ],
476+ "outputs": [
477+ {
478+ "name": "y",
479+ "index": 0,
480+ "dtype": "bfloat16",
481+ "format": "ND",
482+ "paramType": "required",
483+ "shape": [
484+ -2
485+ ],
486+ "format_match_mode": "FormatAgnostic"
487+ }
488+ ],
489+ "attrs": [
490+ {
491+ "name": "ksize",
492+ "dtype": "list_int",
493+ "value": null
494+ },
495+ {
496+ "name": "strides",
497+ "dtype": "list_int",
498+ "value": null
499+ },
500+ {
501+ "name": "pads",
502+ "dtype": "list_int",
503+ "value": null
504+ },
505+ {
506+ "name": "dilation",
507+ "dtype": "list_int",
508+ "value": [
509+ 1,
510+ 1,
511+ 1
512+ ]
513+ },
514+ {
515+ "name": "ceil_mode",
516+ "dtype": "bool",
517+ "value": null
518+ },
519+ {
520+ "name": "data_format",
521+ "dtype": "string",
522+ "value": "NCDHW"
523+ }
524+ ]
525+ }
526+ ]
527+}
Apooling/max_pool3d_grad_with_argmax/op_host/config/ascend950/max_pool3d_grad_with_argmax_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+[MaxPool3DGradWithArgmax]
13+default=0
Mpooling/max_pool3d_grad_with_argmax/op_host/max_pool3d_grad_with_argmax_cutk_tiling.cpp+1-1
@@ -290,5 +290,5 @@ ge::graphStatus MaxPool3DGradWithArgmaxCutKTiling::DoOpTiling()
290 return ge::GRAPH_SUCCESS;290 return ge::GRAPH_SUCCESS;
291}291}
292 292 
293-REGISTER_TILING_TEMPLATE("MaxPool3DGradWithArgmax", MaxPool3DGradWithArgmaxCutKTiling, 0);293+REGISTER_TILING_TEMPLATE("MaxPool3DGradWithArgmax", MaxPool3DGradWithArgmaxCutKTiling, 10);
294} // namespace optiling294} // namespace optiling
Mpooling/max_pool3d_grad_with_argmax/op_host/max_pool3d_grad_with_argmax_def.cpp+102-65
@@ -1,66 +1,103 @@
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-/*!11+/*!
12- * \file max_pool3d_grad_with_argmax.cpp12+ * \file max_pool3d_grad_with_argmax.cpp
13- * \brief13+ * \brief
14- */14+ */
15- 15+ 
16-#include "register/op_def_registry.h"16+#include "register/op_def_registry.h"
17- 17+ 
18-namespace ops {18+namespace ops {
19-class MaxPool3DGradWithArgmax : public OpDef {19+constexpr int32_t DTYPE_INT32 = 3;
20-public:20+class MaxPool3DGradWithArgmax : public OpDef {
21- explicit MaxPool3DGradWithArgmax(const char* name) : OpDef(name)21+public:
22- {22+ explicit MaxPool3DGradWithArgmax(const char* name) : OpDef(name)
23- this->Input("x")23+ {
24- .ParamType(REQUIRED)24+ this->Input("x")
25- .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})25+ .ParamType(REQUIRED)
26- .Format({ge::FORMAT_NCDHW, ge::FORMAT_NCDHW, ge::FORMAT_NCDHW})26+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})
27- .UnknownShapeFormat({ge::FORMAT_NCDHW, ge::FORMAT_NCDHW, ge::FORMAT_NCDHW})27+ .Format({ge::FORMAT_NCDHW, ge::FORMAT_NCDHW, ge::FORMAT_NCDHW})
28- .AutoContiguous();28+ .UnknownShapeFormat({ge::FORMAT_NCDHW, ge::FORMAT_NCDHW, ge::FORMAT_NCDHW})
29- this->Input("grad")29+ .AutoContiguous();
30- .ParamType(REQUIRED)30+ this->Input("grad")
31- .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})31+ .ParamType(REQUIRED)
32- .Format({ge::FORMAT_NCDHW, ge::FORMAT_NCDHW, ge::FORMAT_NCDHW})32+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})
33- .UnknownShapeFormat({ge::FORMAT_NCDHW, ge::FORMAT_NCDHW, ge::FORMAT_NCDHW})33+ .Format({ge::FORMAT_NCDHW, ge::FORMAT_NCDHW, ge::FORMAT_NCDHW})
34- .AutoContiguous();34+ .UnknownShapeFormat({ge::FORMAT_NCDHW, ge::FORMAT_NCDHW, ge::FORMAT_NCDHW})
35- this->Input("argmax")35+ .AutoContiguous();
36- .ParamType(REQUIRED)36+ this->Input("argmax")
37- .DataType({ge::DT_INT32, ge::DT_INT32, ge::DT_INT32})37+ .ParamType(REQUIRED)
38- .Format({ge::FORMAT_NCDHW, ge::FORMAT_NCDHW, ge::FORMAT_NCDHW})38+ .DataType({ge::DT_INT32, ge::DT_INT32, ge::DT_INT32})
39- .UnknownShapeFormat({ge::FORMAT_NCDHW, ge::FORMAT_NCDHW, ge::FORMAT_NCDHW})39+ .Format({ge::FORMAT_NCDHW, ge::FORMAT_NCDHW, ge::FORMAT_NCDHW})
40- .AutoContiguous();40+ .UnknownShapeFormat({ge::FORMAT_NCDHW, ge::FORMAT_NCDHW, ge::FORMAT_NCDHW})
41- this->Output("y")41+ .AutoContiguous();
42- .ParamType(REQUIRED)42+ this->Output("y")
43- .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})43+ .ParamType(REQUIRED)
44- .Format({ge::FORMAT_NCDHW, ge::FORMAT_NCDHW, ge::FORMAT_NCDHW})44+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})
45- .UnknownShapeFormat({ge::FORMAT_NCDHW, ge::FORMAT_NCDHW, ge::FORMAT_NCDHW});45+ .Format({ge::FORMAT_NCDHW, ge::FORMAT_NCDHW, ge::FORMAT_NCDHW})
46- this->Attr("ksize").AttrType(REQUIRED).ListInt();46+ .UnknownShapeFormat({ge::FORMAT_NCDHW, ge::FORMAT_NCDHW, ge::FORMAT_NCDHW});
47- this->Attr("strides").AttrType(REQUIRED).ListInt();47+ this->Attr("ksize").AttrType(REQUIRED).ListInt();
48- this->Attr("pads").AttrType(REQUIRED).ListInt();48+ this->Attr("strides").AttrType(REQUIRED).ListInt();
49- this->Attr("dilation").AttrType(OPTIONAL).ListInt({1, 1, 1});49+ this->Attr("pads").AttrType(REQUIRED).ListInt();
50- this->Attr("ceil_mode").AttrType(OPTIONAL).Bool(false);50+ this->Attr("dilation").AttrType(OPTIONAL).ListInt({1, 1, 1});
51- 51+ this->Attr("ceil_mode").AttrType(OPTIONAL).Bool(false);
52- OpAICoreConfig aicore_config;52+ 
53- aicore_config.DynamicCompileStaticFlag(true)53+ OpAICoreConfig aicore_config;
54- .DynamicRankSupportFlag(true)54+ aicore_config.DynamicCompileStaticFlag(true)
55- .DynamicShapeSupportFlag(true)55+ .DynamicRankSupportFlag(true)
56- .ExtendCfgInfo("opFile.value", "max_pool3d_grad_with_argmax")56+ .DynamicShapeSupportFlag(true)
57- .ExtendCfgInfo("opInterface.value", "max_pool3d_grad_with_argmax")57+ .ExtendCfgInfo("opFile.value", "max_pool3d_grad_with_argmax")
58- .ExtendCfgInfo("aclnnSupport.value", "support_aclnn");58+ .ExtendCfgInfo("opInterface.value", "max_pool3d_grad_with_argmax")
59- 59+ .ExtendCfgInfo("aclnnSupport.value", "support_aclnn");
60- this->AICore().AddConfig("ascend910b", aicore_config);60+ 
61- this->AICore().AddConfig("ascend910_93", aicore_config);61+ this->AICore().AddConfig("ascend910b", aicore_config);
62- }62+ this->AICore().AddConfig("ascend910_93", aicore_config);
63-};63+ 
64- 64+ OpAICoreConfig aicoreConfig950;
65-OP_ADD(MaxPool3DGradWithArgmax);65+ aicoreConfig950.Input("x")
66+ .ParamType(REQUIRED)
67+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})
68+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
69+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
70+ .AutoContiguous();
71+ aicoreConfig950.Input("grad")
72+ .ParamType(REQUIRED)
73+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})
74+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
75+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
76+ .AutoContiguous();
77+ aicoreConfig950.Input("argmax")
78+ .ParamType(REQUIRED)
79+ .DataType({ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64})
80+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
81+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
82+ .AutoContiguous();
83+ aicoreConfig950.Output("y")
84+ .ParamType(REQUIRED)
85+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})
86+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
87+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
88+ this->Attr("data_format").AttrType(OPTIONAL).String("NCDHW");
89+ 
90+ aicoreConfig950.DynamicCompileStaticFlag(true)
91+ .DynamicFormatFlag(false)
92+ .DynamicRankSupportFlag(true)
93+ .DynamicShapeSupportFlag(true)
94+ .NeedCheckSupportFlag(false)
95+ .PrecisionReduceFlag(true)
96+ .ExtendCfgInfo("opFile.value", "max_pool3d_grad_with_argmax")
97+ .ExtendCfgInfo("opInterface.value", "max_pool3d_grad_with_argmax");
98+ this->AICore().AddConfig("ascend950", aicoreConfig950);
99+ }
100+};
101+ 
102+OP_ADD(MaxPool3DGradWithArgmax);
66} // namespace ops103} // namespace ops
Apooling/max_pool3d_grad_with_argmax/op_host/max_pool3d_grad_with_argmax_infershape_arch35.cpp+300-0
@@ -0,0 +1,300 @@
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_pool3d_grad_with_argmax_infershape_arch35.cpp
13+ * \brief
14+ */
15+#include <string>
16+#include "error_util.h"
17+#include "exe_graph/runtime/infer_shape_context.h"
18+#include "graph/utils/type_utils.h"
19+#include "log/log.h"
20+#include "register/op_impl_registry.h"
21+#include "util/shape_util.h"
22+ 
23+using namespace ge;
24+namespace ops {
25+static constexpr size_t ATTR_INDEX_KSIZE = 0;
26+static constexpr size_t ATTR_INDEX_STRIDES = 1;
27+static constexpr size_t ATTR_INDEX_PADS = 2;
28+static constexpr size_t ATTR_INDEX_DILATION = 3;
29+static constexpr size_t ATTR_INDEX_CEIL_MODE = 4;
30+static constexpr size_t ATTR_INDEX_DATA_FORMAT = 5;
31+static constexpr size_t INDEX_OUT_MAX = 0;
32+static constexpr size_t ATTR_LIST_SHAPE_SIZE = 3;
33+static constexpr size_t INDEX_ZERO = 0;
34+static constexpr size_t INDEX_ONE = 1;
35+static constexpr size_t INDEX_TWO = 2;
36+static constexpr size_t CDHW_DIM_NUM = 4;
37+static constexpr size_t NUMBER_TWO = 2;
38+static constexpr size_t PARAM_NUM = 4;
39+static constexpr size_t SHAPE_D_DIM = 2;
40+static constexpr size_t SHAPE_H_DIM = 3;
41+static constexpr size_t SHAPE_W_DIM = 4;
42+static constexpr int64_t UNKNOWN_DIM_VALUE_ = -1LL;
43+ 
44+static int64_t DivRtn(int64_t x, int64_t y)
45+{
46+ if (y == 0) {
47+ OP_LOGE("MaxPool3DGradWithArgmax", "strides value cannot be zero.");
48+ return GRAPH_FAILED;
49+ }
50+ if (x < 0) {
51+ OP_LOGE("MaxPool3DGradWithArgmax", "x value cannot small than zero.");
52+ return GRAPH_FAILED;
53+ }
54+ int64_t q = x / y;
55+ return q;
56+}
57+ 
58+static void UpdateMaxShape(
59+ const int64_t (&param)[PARAM_NUM], bool ceil_mode, const int64_t& dim_size, int64_t& out_max_shape)
60+{
61+ int64_t ksize = param[ATTR_INDEX_KSIZE];
62+ int64_t strides = param[ATTR_INDEX_STRIDES];
63+ int64_t pad = param[ATTR_INDEX_PADS];
64+ int64_t dilation = param[PARAM_NUM - 1];
65+ int64_t exact_size = dim_size + 2 * pad - dilation * (ksize - 1) - 1 + (ceil_mode ? (strides - 1) : 0);
66+ out_max_shape = DivRtn(exact_size, strides) + 1;
67+ if (ceil_mode) {
68+ if ((out_max_shape - 1) * strides >= dim_size + pad) {
69+ out_max_shape = out_max_shape - 1;
70+ }
71+ }
72+}
73+ 
74+static int64_t CalcOutDim(int64_t input_dim, int64_t k, int64_t s, int64_t p, int64_t d, bool ceil_mode)
75+{
76+ int64_t param[PARAM_NUM] = {k, s, p, d};
77+ int64_t out_dim = 0;
78+ UpdateMaxShape(param, ceil_mode, input_dim, out_dim);
79+ return out_dim;
80+}
81+ 
82+inline ge::graphStatus SetAllUnknownDim(const int64_t rank, gert::Shape* output_shape)
83+{
84+ OP_CHECK_IF(
85+ output_shape == nullptr, OP_LOGD("SetAllUnknownDim", "the output_shape is nullptr, return unsuccess"),
86+ return ge::GRAPH_FAILED);
87+ output_shape->SetDimNum(rank);
88+ for (int64_t i = 0; i < rank; ++i) {
89+ output_shape->SetDim(i, UNKNOWN_DIM_VALUE_);
90+ }
91+ OP_LOGD("SetAllUnknownDim", "set all dim = -1, output = %s", Ops::Base::ToString(*output_shape).c_str());
92+ return ge::GRAPH_SUCCESS;
93+}
94+ 
95+ge::graphStatus InferShapeForMaxPool3DGradWithArgmax(gert::InferShapeContext* context)
96+{
97+ if (context == nullptr) {
98+ return ge::GRAPH_FAILED;
99+ }
100+ 
101+ OP_LOGD(context->GetNodeName(), "runtime2.0 MaxPool3DGradWithArgmax infershape running");
102+ auto inputXDesc = context->GetInputDesc(0);
103+ OP_CHECK_NULL_WITH_CONTEXT(context, inputXDesc);
104+ auto inputXFormat = inputXDesc->GetOriginFormat();
105+ OP_CHECK_IF(
106+ inputXFormat != FORMAT_ND && inputXFormat != FORMAT_NCDHW && inputXFormat != FORMAT_NDHWC &&
107+ inputXFormat != FORMAT_NCHW,
108+ OP_LOGE(context->GetNodeName(), "format only supports ND, NCDHW, NDHWC"), return GRAPH_FAILED);
109+ 
110+ size_t input_d_dim = SHAPE_D_DIM;
111+ size_t input_h_dim = SHAPE_H_DIM;
112+ size_t input_w_dim = SHAPE_W_DIM;
113+ 
114+ auto attrs = context->GetAttrs();
115+ OP_CHECK_NULL_WITH_CONTEXT(context, attrs);
116+ 
117+ auto ksize = attrs->GetAttrPointer<gert::ContinuousVector>(ATTR_INDEX_KSIZE);
118+ OP_CHECK_NULL_WITH_CONTEXT(context, ksize);
119+ OP_CHECK_IF(
120+ ksize->GetSize() != 1 && ksize->GetSize() != ATTR_LIST_SHAPE_SIZE,
121+ OP_LOGE(context->GetNodeName(), "Length of ksize %lu must be equal 1 or 3!", ksize->GetSize()),
122+ return GRAPH_FAILED);
123+ 
124+ auto ksize_data = reinterpret_cast<const int64_t*>(ksize->GetData());
125+ for (uint32_t i = 0; i < static_cast<uint32_t>(ksize->GetSize()); i++) {
126+ OP_CHECK_IF(
127+ (ksize_data[i] <= 0),
128+ OP_LOGE(
129+ context->GetNodeName(), "Attr value invalid, ksize_data[%u] is %ld, should bigger than 0.", i,
130+ ksize_data[i]),
131+ return ge::GRAPH_FAILED);
132+ }
133+ 
134+ int64_t kD = ksize_data[INDEX_ZERO];
135+ int64_t kH = ksize->GetSize() == 1 ? kD : ksize_data[INDEX_ONE];
136+ int64_t kW = ksize->GetSize() == 1 ? kD : ksize_data[INDEX_TWO];
137+ 
138+ auto strides = attrs->GetAttrPointer<gert::ContinuousVector>(ATTR_INDEX_STRIDES);
139+ OP_CHECK_NULL_WITH_CONTEXT(context, strides);
140+ OP_CHECK_IF(
141+ strides->GetSize() != 0 && strides->GetSize() != 1 && strides->GetSize() != ATTR_LIST_SHAPE_SIZE,
142+ OP_LOGE(context->GetNodeName(), "Length of strides %lu must be equal 0 or 1 or 3!", strides->GetSize()),
143+ return GRAPH_FAILED);
144+ 
145+ auto strides_data = reinterpret_cast<const int64_t*>(strides->GetData());
146+ for (uint32_t i = 0; i < static_cast<uint32_t>(strides->GetSize()); i++) {
147+ OP_CHECK_IF(
148+ (strides_data[i] <= 0),
149+ OP_LOGE(
150+ context->GetNodeName(), "Attr value invalid, strides_data[%u] is %ld, should bigger than 0.", i,
151+ strides_data[i]),
152+ return ge::GRAPH_FAILED);
153+ }
154+ 
155+ int64_t sD = kD;
156+ int64_t sH = kH;
157+ int64_t sW = kW;
158+ if (strides->GetSize() > 0) {
159+ sD = strides_data[INDEX_ZERO];
160+ sH = (strides->GetSize() > 1) ? strides_data[INDEX_ONE] : sD;
161+ sW = (strides->GetSize() > 1) ? strides_data[INDEX_TWO] : sD;
162+ }
163+ 
164+ auto pads = attrs->GetAttrPointer<gert::ContinuousVector>(ATTR_INDEX_PADS);
165+ OP_CHECK_NULL_WITH_CONTEXT(context, pads);
166+ OP_CHECK_IF(
167+ pads->GetSize() != 1 && pads->GetSize() != ATTR_LIST_SHAPE_SIZE,
168+ OP_LOGE(context->GetNodeName(), "Length of pads %lu must be equal 1 or 3!", pads->GetSize()),
169+ return GRAPH_FAILED);
170+ 
171+ auto pads_data = reinterpret_cast<const int64_t*>(pads->GetData());
172+ for (uint32_t i = 0; i < static_cast<uint32_t>(pads->GetSize()); i++) {
173+ OP_CHECK_IF(
174+ (pads_data[i] < 0),
175+ OP_LOGE(
176+ context->GetNodeName(), "Attr value invalid, pads_data[%u] is %ld, should bigger or equal 0.", i,
177+ pads_data[i]),
178+ return ge::GRAPH_FAILED);
179+ }
180+ 
181+ int64_t pD = pads_data[INDEX_ZERO];
182+ int64_t pH = pads->GetSize() == 1 ? pD : pads_data[INDEX_ONE];
183+ int64_t pW = pads->GetSize() == 1 ? pD : pads_data[INDEX_TWO];
184+ 
185+ auto dilation = attrs->GetAttrPointer<gert::ContinuousVector>(ATTR_INDEX_DILATION);
186+ OP_CHECK_NULL_WITH_CONTEXT(context, dilation);
187+ OP_CHECK_IF(
188+ dilation->GetSize() != 1 && dilation->GetSize() != ATTR_LIST_SHAPE_SIZE,
189+ OP_LOGE(context->GetNodeName(), "Length of dilation %lu must be equal 1 or 3!", dilation->GetSize()),
190+ return GRAPH_FAILED);
191+ 
192+ auto dilation_data = reinterpret_cast<const int64_t*>(dilation->GetData());
193+ for (uint32_t i = 0; i < static_cast<uint32_t>(dilation->GetSize()); i++) {
194+ OP_CHECK_IF(
195+ (dilation_data[i] <= 0),
196+ OP_LOGE(
197+ context->GetNodeName(), "Attr value invalid, dilation_data[%u] is %ld, should bigger than 0.", i,
198+ dilation_data[i]),
199+ return ge::GRAPH_FAILED);
200+ }
201+ 
202+ int64_t dD = dilation_data[INDEX_ZERO];
203+ int64_t dH = dilation->GetSize() == 1 ? dD : dilation_data[INDEX_ONE];
204+ int64_t dW = dilation->GetSize() == 1 ? dD : dilation_data[INDEX_TWO];
205+ 
206+ auto ceil_mode = attrs->GetAttrPointer<bool>(ATTR_INDEX_CEIL_MODE);
207+ OP_CHECK_NULL_WITH_CONTEXT(context, ceil_mode);
208+ 
209+ const char* data_format = attrs->GetAttrPointer<char>(ATTR_INDEX_DATA_FORMAT);
210+ OP_CHECK_NULL_WITH_CONTEXT(context, data_format);
211+ 
212+ OP_CHECK_IF(
213+ (pD > (kD / 2)) || (pH > (kH / 2)) || (pW > (kW / 2)),
214+ OP_LOGE(context->GetNodeName(), "Attr size invalid, padSize should smaller than kernelSize div 2"),
215+ return ge::GRAPH_FAILED);
216+ 
217+ OP_CHECK_IF(
218+ (pD > ((kD - 1) * dD + 1) / 2) || (pH > ((kH - 1) * dH + 1) / 2) || (pW > ((kW - 1) * dW + 1) / 2),
219+ OP_LOGE(
220+ context->GetNodeName(),
221+ "Attr size invalid, padSize should smaller than ((kernelSize - 1) * dilation + 1) / 2."),
222+ return ge::GRAPH_FAILED);
223+ 
224+ const gert::Shape* xShape = context->GetInputShape(0);
225+ OP_CHECK_NULL_WITH_CONTEXT(context, xShape);
226+ const gert::Shape* gradShape = context->GetInputShape(1);
227+ OP_CHECK_NULL_WITH_CONTEXT(context, gradShape);
228+ const gert::Shape* argmaxShape = context->GetInputShape(2);
229+ OP_CHECK_NULL_WITH_CONTEXT(context, argmaxShape);
230+ gert::Shape* yShape = context->GetOutputShape(0);
231+ OP_CHECK_NULL_WITH_CONTEXT(context, yShape);
232+ size_t xDimNum = xShape->GetDimNum();
233+ 
234+ std::string data_format_str = data_format;
CANN-robot
CANN-robotCANN-robot1月29日

逻辑错误: data_format_str从char*构造，但data_format可能不是以null结尾的字符串。attrs->GetAttrPointer()返回的是指向单个char的指针，而不是字符串。

问题类型: 逻辑错误 文件路径: pooling/max_pool3d_grad_with_argmax/op_host/max_pool3d_grad_with_argmax_infershape_arch35.cpp 行号: 234 问题代码:

    std::string data_format_str = data_format;

修改建议:

检查data_format的获取方式。如果它确实是字符串，应该使用const char*类型。否则需要修改为获取字符串属性。

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

likedislike
小王！
小王！
1月31日 评论：
235+ if (data_format_str == "NDHWC" || xDimNum == CDHW_DIM_NUM) {
236+ input_d_dim = input_d_dim - 1;
237+ input_h_dim = input_h_dim - 1;
238+ input_w_dim = input_w_dim - 1;
239+ }
240+ 
241+ if (Ops::Base::IsUnknownShape(*xShape) || Ops::Base::IsUnknownShape(*gradShape) ||
242+ Ops::Base::IsUnknownShape(*argmaxShape)) {
243+ SetAllUnknownDim(xDimNum, yShape);
244+ OP_LOGD(context->GetNodeName(), "runtime2.0 MaxPool3DGradWithArgmax infershape handle unknown shape.");
245+ return ge::GRAPH_SUCCESS;
246+ }
247+ 
248+ if (Ops::Base::IsUnknownRank(*xShape)) {
249+ Ops::Base::SetUnknownRank(*yShape);
250+ OP_LOGD(context->GetNodeName(), "runtime2.0 MaxPool3DGradWithArgmax infershape handle unknown rank.");
251+ return ge::GRAPH_SUCCESS;
252+ }
253+ yShape->SetDimNum(xDimNum);
254+ *yShape = *xShape;
255+ 
256+ // Check gradShape and argmaxShape input dim invaild
257+ int64_t doExpected = CalcOutDim(xShape->GetDim(input_d_dim), kD, sD, pD, dD, *ceil_mode);
258+ int64_t hoExpected = CalcOutDim(xShape->GetDim(input_h_dim), kH, sH, pH, dH, *ceil_mode);
259+ int64_t woExpected = CalcOutDim(xShape->GetDim(input_w_dim), kW, sW, pW, dW, *ceil_mode);
260+ 
261+ OP_CHECK_IF(
262+ (!Ops::Base::IsUnknownRank(*gradShape) && !Ops::Base::IsUnknownShape(*gradShape)) &&
263+ ((doExpected <= 0) || (static_cast<uint64_t>(doExpected) != gradShape->GetDim(input_d_dim)) ||
CANN-robot
CANN-robotCANN-robot1月29日

类型转换问题: 将int64_t的doExpected转换为uint64_t与gradShape->GetDim()比较，但GetDim()可能返回int64_t。直接比较可能因符号问题导致错误。且当doExpected为负数时，转换为uint64_t会产生大正数。

问题类型: 类型转换问题 文件路径: pooling/max_pool3d_grad_with_argmax/op_host/max_pool3d_grad_with_argmax_infershape_arch35.cpp 行号: 263 问题代码:

            ((doExpected <= 0) || (static_cast<uint64_t>(doExpected) != gradShape->GetDim(input_d_dim)) ||

修改建议:

直接使用int64_t比较：doExpected != static_cast<int64_t>(gradShape->GetDim(input_d_dim))。同时确保GetDim()返回类型与doExpected匹配。

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

likedislike
小王！
小王！
1月31日 评论：
264+ (hoExpected <= 0) || (static_cast<uint64_t>(hoExpected) != gradShape->GetDim(input_h_dim)) ||
265+ (woExpected <= 0) || (static_cast<uint64_t>(woExpected) != gradShape->GetDim(input_w_dim))),
266+ OP_LOGE(
267+ context->GetNodeName(), "GradShape D / H / W size invalid, expected d: %ld, h: %ld, w: %ld.", doExpected,
268+ hoExpected, woExpected),
269+ return ge::GRAPH_FAILED);
270+ 
271+ OP_CHECK_IF(
272+ (!Ops::Base::IsUnknownRank(*argmaxShape) && !Ops::Base::IsUnknownShape(*argmaxShape)) &&
273+ ((doExpected <= 0) || (static_cast<uint64_t>(doExpected) != argmaxShape->GetDim(input_d_dim)) ||
274+ (hoExpected <= 0) || (static_cast<uint64_t>(hoExpected) != argmaxShape->GetDim(input_h_dim)) ||
275+ (woExpected <= 0) || (static_cast<uint64_t>(woExpected) != argmaxShape->GetDim(input_w_dim))),
276+ OP_LOGE(
277+ context->GetNodeName(), "ArgmaxShape D / H / W size invalid, expected d: %ld, h: %ld, w: %ld.", doExpected,
278+ hoExpected, woExpected),
279+ return ge::GRAPH_FAILED);
280+ 
281+ OP_LOGD(context->GetNodeName(), "runtime2.0 MaxPool3DGradWithArgmax infershape run success.");
282+ return ge::GRAPH_SUCCESS;
283+}
284+ 
285+static ge::graphStatus InferDataTypeForMaxPool3DGradWithArgmax(gert::InferDataTypeContext* context)
286+{
287+ if (context == nullptr) {
288+ return ge::GRAPH_FAILED;
289+ }
290+ 
291+ const ge::DataType xDtype = context->GetInputDataType(0);
292+ context->SetOutputDataType(INDEX_OUT_MAX, xDtype);
293+ 
294+ return ge::GRAPH_SUCCESS;
295+}
296+ 
297+IMPL_OP_INFERSHAPE(MaxPool3DGradWithArgmax)
298+ .InferShape(InferShapeForMaxPool3DGradWithArgmax)
299+ .InferDataType(InferDataTypeForMaxPool3DGradWithArgmax);
300+} // namespace ops
Mpooling/max_pool3d_grad_with_argmax/op_host/max_pool3d_grad_with_argmax_normal_tiling.cpp+1-1
@@ -337,5 +337,5 @@ ge::graphStatus MaxPool3DGradWithArgmaxNormalTiling::DoOpTiling()
337 return ge::GRAPH_SUCCESS;337 return ge::GRAPH_SUCCESS;
338}338}
339 339 
340-REGISTER_TILING_TEMPLATE("MaxPool3DGradWithArgmax", MaxPool3DGradWithArgmaxNormalTiling, 2);340+REGISTER_TILING_TEMPLATE("MaxPool3DGradWithArgmax", MaxPool3DGradWithArgmaxNormalTiling, 12);
341} // namespace optiling341} // namespace optiling
Mpooling/max_pool3d_grad_with_argmax/op_host/max_pool3d_grad_with_argmax_scatter_tiling.cpp+1-1
@@ -81,6 +81,6 @@ ge::graphStatus MaxPool3DGradWithArgmaxScatterTiling::DoOpTiling()
81}81}
82 82 
83REGISTER_TILING_TEMPLATE("MaxPool3DGradWithArgmax", 83REGISTER_TILING_TEMPLATE("MaxPool3DGradWithArgmax",
84- MaxPool3DGradWithArgmaxScatterTiling, 6);84+ MaxPool3DGradWithArgmaxScatterTiling, 16);
85 85 
86} // namespace optiling86} // namespace optiling
Apooling/max_pool3d_grad_with_argmax/op_host/max_pool3d_grad_with_argmax_simd_tiling.cpp+490-0
@@ -0,0 +1,490 @@
1+ 
2+/**
3+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
4+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
5+ * 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,
8+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9+ * See LICENSE in the root of the software repository for the full text of the License.
10+ */
11+ 
12+/*!
13+ * \file max_pool3d_grad_with_argmax_simd_tiling.cpp
14+ * \brief
15+ */
16+#include "platform/platform_info.h"
17+#include "tiling_base/tiling_templates_registry.h"
18+#include "max_pool3d_grad_with_argmax_tiling_arch35.h"
19+ 
20+namespace optiling {
21+static constexpr int64_t FLOAT16_SIZE = 2;
22+static constexpr int64_t FLOAT32_SIZE = 4;
23+static constexpr int64_t INT32_SIZE = 4;
24+static constexpr int64_t INT64_SIZE = 8;
25+static constexpr int64_t UB_RESVERVED_SIZE = 2048;
26+static constexpr int64_t T3_INT64 = 10;
27+static constexpr int64_t DOUBLE_BUFFER = 2;
28+static constexpr int64_t THRESHOLD= 2;
29+ 
30+void MaxPool3DGradWithArgmaxNCDHWTiling::InitializationVars()
31+{
32+ baseData.vRegSize = Ops::Base::GetVRegSize(context_);
33+ baseData.ubBlockSize = Ops::Base::GetUbBlockSize(context_);
34+ baseData.inputBytes = inputData.inputDtype == ge::DT_FLOAT ? FLOAT32_SIZE : FLOAT16_SIZE;
35+ baseData.indexBytes = inputData.indexDtype == ge::DT_INT32 ? INT32_SIZE : INT64_SIZE;
36+ baseData.availableUb = ubSize_ - UB_RESVERVED_SIZE;
37+ baseData.totalCoreNum = coreNum_;
38+ baseData.coreUsedForBestPerformance = baseData.totalCoreNum;
39+ 
40+ int64_t oneBlockNumT1 = baseData.ubBlockSize / baseData.inputBytes;
41+ int64_t oneBlockNumT2 = baseData.ubBlockSize / baseData.indexBytes;
42+ 
43+ baseData.maxDataNumInOneBlock = std::max(oneBlockNumT1, oneBlockNumT2);
44+ 
45+ baseData.proDataNumInOneBeatT2 = baseData.vRegSize / baseData.ubBlockSize * oneBlockNumT2;
46+ baseData.inputNCSize = inputData.nX * inputData.cX;
47+ 
48+ baseData.isPad = 0;
49+ if (inputData.hPad != 0 || inputData.wPad != 0 || inputData.dPad != 0) {
50+ baseData.isPad = 1;
51+ }
52+ baseData.dProBatchSize = 1;
53+ if (inputData.dKernel > inputData.dStride) {
54+ baseData.dProBatchSize = Ops::Base::CeilDiv(inputData.dKernel, inputData.dStride);
55+ }
56+ 
57+ baseData.hProBatchSize = 1;
58+ if (inputData.hKernel > inputData.hStride) {
59+ baseData.hProBatchSize = Ops::Base::CeilDiv(inputData.hKernel, inputData.hStride);
60+ }
61+ 
62+ baseData.wProBatchSize = 1;
63+ if (inputData.wKernel > inputData.wStride) {
64+ baseData.wProBatchSize = Ops::Base::CeilDiv(inputData.wKernel, inputData.wStride);
65+ }
66+ 
67+ baseData.isOverlap = 0;
68+ if (baseData.wProBatchSize != 1 || baseData.hProBatchSize != 1 || baseData.dProBatchSize != 1) {
69+ baseData.isOverlap = 1;
70+ }
71+}
72+ 
73+bool MaxPool3DGradWithArgmaxNCDHWTiling::IsCapable()
74+{
75+ InitializationVars();
76+ if(inputData.inputFormat != ge::Format::FORMAT_NCDHW)
77+ {
78+ return false;
79+ }
80+ if (baseData.hProBatchSize >= inputData.hGrad && baseData.wProBatchSize >= inputData.wGrad && baseData.wProBatchSize >= inputData.wGrad) {
81+ return false;
82+ }
83+ if(inputData.dDilation != 1 || inputData.hDilation != 1 || inputData.wDilation != 1) {
84+ return false;
85+ }
86+ // ub is not enough
87+ splitData.highAxisInner = 1;
88+ splitData.dOutputInner = 1;
89+ splitData.hOutputInner = 1;
90+ splitData.wOutputInner = std::min(inputData.wX, baseData.proDataNumInOneBeatT2);
91+ DoBufferCalculate();
92+ return splitData.totalBufferSize <= baseData.availableUb;
93+}
94+ 
95+uint64_t MaxPool3DGradWithArgmaxNCDHWTiling::GetTilingKey() const
96+{
97+ int64_t outDataCount = inputData.nX * inputData.cX * inputData.dX * inputData.hX * inputData.wX;
98+ uint32_t idxDtype = outDataCount <= static_cast<int64_t>(MAX_INT32) ? TPL_INT32 : TPL_INT64;
99+ uint32_t isChannelLast = 0;
100+ uint32_t isSimt = 0;
101+ return GET_TPL_TILING_KEY(idxDtype, isSimt, isChannelLast, splitData.isCheckRange);
102+}
103+ 
104+ 
105+void MaxPool3DGradWithArgmaxNCDHWTiling::DoBufferCalculate()
106+{
107+ // The calculation only involves inner.
108+ int64_t dInputInner = Ops::Base::CeilDiv(splitData.dOutputInner + inputData.dKernel - 1, inputData.dStride);
109+ int64_t hInputInner = Ops::Base::CeilDiv(splitData.hOutputInner + inputData.hKernel - 1, inputData.hStride);
110+ int64_t wInputInner = Ops::Base::CeilDiv(splitData.wOutputInner + inputData.wKernel - 1, inputData.wStride);
111+ int64_t wInputInnerAligned = Ops::Base::CeilAlign(wInputInner, baseData.maxDataNumInOneBlock);
112+ int64_t wOutputInnerAligned = Ops::Base::CeilAlign(splitData.wOutputInner, baseData.maxDataNumInOneBlock);
113+ 
114+ int64_t inputPlaneSizeDHW = dInputInner * hInputInner * wInputInnerAligned;
115+ int64_t outputPlaneSizeDHW = splitData.dOutputInner * splitData.hOutputInner * wOutputInnerAligned;
116+ 
117+ splitData.gradBufferSize = splitData.highAxisInner * inputPlaneSizeDHW * baseData.inputBytes;
118+ splitData.argmaxBufferSize = splitData.highAxisInner * inputPlaneSizeDHW * baseData.indexBytes;
119+ splitData.outputBufferSize = splitData.highAxisInner * outputPlaneSizeDHW * FLOAT32_SIZE;
120+ 
121+ int64_t tmpTotalBufferSize = splitData.outputBufferSize + splitData.gradBufferSize + splitData.argmaxBufferSize;
122+ splitData.totalBufferSize = tmpTotalBufferSize * DOUBLE_BUFFER;
123+}
124+ 
125+bool MaxPool3DGradWithArgmaxNCDHWTiling::IsMeetTargetCoreNum() const
126+{
127+ int64_t tmpWOutputOuter = Ops::Base::CeilDiv(inputData.wX, splitData.wOutputInner);
128+ int64_t tmpHOutputOuter = Ops::Base::CeilDiv(inputData.hX, splitData.hOutputInner);
129+ int64_t tmpDOutputOuter = Ops::Base::CeilDiv(inputData.dX, splitData.dOutputInner);
130+ int64_t tmpHighAxisOutputOuter = Ops::Base::CeilDiv(baseData.inputNCSize, splitData.highAxisInner);
131+ 
132+ return tmpDOutputOuter * tmpWOutputOuter * tmpHOutputOuter * tmpHighAxisOutputOuter >= baseData.coreUsedForBestPerformance;
133+}
134+ 
135+bool MaxPool3DGradWithArgmaxNCDHWTiling::IsMeetUBSize()
136+{
137+ DoBufferCalculate();
138+ return splitData.totalBufferSize <= baseData.availableUb;
139+}
140+ 
141+bool MaxPool3DGradWithArgmaxNCDHWTiling::TrySplitNC()
142+{
143+ splitData.wOutputInner = inputData.wX;
144+ splitData.hOutputInner = inputData.hX;
145+ splitData.dOutputInner = inputData.dX;
146+ splitData.highAxisInner = Ops::Base::CeilDiv(baseData.inputNCSize, baseData.coreUsedForBestPerformance);
147+ if (IsMeetUBSize() && IsMeetTargetCoreNum()) {
148+ return true;
149+ }
150+ 
151+ splitData.highAxisInner = 1;
152+ if (IsMeetUBSize() && IsMeetTargetCoreNum()) {
153+ int64_t left = 1;
154+ int64_t right = baseData.inputNCSize;
155+ int64_t bestSplit = 1;
156+ while (left <= right) {
157+ int64_t mid = left + (right - left) / 2;
158+ splitData.highAxisInner = mid;
159+ 
160+ if (IsMeetUBSize() && IsMeetTargetCoreNum()) {
161+ bestSplit = mid;
162+ left = mid + 1;
163+ } else {
164+ right = mid - 1;
165+ }
166+ }
167+ 
168+ splitData.highAxisInner = bestSplit;
169+ return true;
170+ } else {
171+ return false;
172+ }
173+}
174+ 
175+bool MaxPool3DGradWithArgmaxNCDHWTiling::TrySplitAlignD()
176+{
177+ splitData.highAxisInner = 1;
178+ splitData.hOutputInner = inputData.hX;
179+ splitData.wOutputInner = inputData.wX;
180+ int64_t halfInput = inputData.dX / 2;
181+ splitData.dOutputInner = inputData.dStride;
182+ if (IsMeetUBSize() && IsMeetTargetCoreNum()) {
183+ int64_t left = 1;
184+ int64_t right = Ops::Base::CeilDiv(halfInput, inputData.dStride);
185+ int64_t bestSplit = 1;
186+ 
187+ while (left <= right) {
188+ int64_t mid = left + (right - left) / 2;
189+ splitData.dOutputInner = mid * inputData.dStride;
190+ 
191+ if (IsMeetUBSize() && IsMeetTargetCoreNum()) {
192+ bestSplit = mid;
193+ left = mid + 1;
194+ } else {
195+ right = mid - 1;
196+ }
197+ }
198+ 
199+ splitData.dOutputInner = bestSplit * inputData.dStride;
200+ return true;
201+ } else {
202+ return false;
203+ }
204+}
205+ 
206+bool MaxPool3DGradWithArgmaxNCDHWTiling::TrySplitAlignH()
207+{
208+ splitData.highAxisInner = 1;
209+ splitData.dOutputInner = inputData.dX;
210+ splitData.wOutputInner = inputData.wX;
211+ 
212+ splitData.hOutputInner = inputData.hStride;
213+ int64_t halfInput = inputData.hX / 2;
214+ if (IsMeetUBSize() && IsMeetTargetCoreNum()) {
215+ int64_t left = 1;
216+ int64_t right = Ops::Base::CeilDiv(halfInput, inputData.hStride);
217+ int64_t bestSplit = 1;
218+ 
219+ while (left <= right) {
220+ int64_t mid = left + (right - left) / 2;
221+ splitData.hOutputInner = mid * inputData.hStride;
222+ 
223+ if (IsMeetUBSize() && IsMeetTargetCoreNum()) {
224+ bestSplit = mid;
225+ left = mid + 1;
226+ } else {
227+ right = mid - 1;
228+ }
229+ }
230+ 
231+ splitData.hOutputInner = bestSplit * inputData.hStride;
232+ return true;
233+ } else {
234+ return false;
235+ }
236+}
237+ 
238+bool MaxPool3DGradWithArgmaxNCDHWTiling::TrySplitAlignW()
239+{
240+ splitData.highAxisInner = 1;
241+ splitData.hOutputInner = inputData.hStride;
242+ splitData.dOutputInner = inputData.dStride;
243+ splitData.wOutputInner = inputData.wStride;
244+ int64_t halfInput = inputData.wX / 2;
245+ if (IsMeetUBSize() && IsMeetTargetCoreNum()) {
246+ int64_t left = 1;
247+ int64_t right = Ops::Base::CeilDiv(halfInput, inputData.wStride);
248+ int64_t bestSplit = 1;
249+ 
250+ while (left <= right) {
251+ int64_t mid = left + (right - left) / 2;
252+ splitData.wOutputInner = mid * inputData.wStride;
253+ 
254+ if (IsMeetUBSize() && IsMeetTargetCoreNum()) {
255+ bestSplit = mid;
256+ left = mid + 1;
257+ } else {
258+ right = mid - 1;
259+ }
260+ }
261+ 
262+ splitData.wOutputInner = bestSplit * inputData.wStride;
263+ return true;
264+ } else {
265+ return false;
266+ }
267+}
268+ 
269+void MaxPool3DGradWithArgmaxNCDHWTiling::SplitUnalignDHW()
270+{
271+ splitData.highAxisInner = 1;
272+ if (baseData.isPad == 0 && baseData.isOverlap == 0) {
273+ splitData.hOutputInner = inputData.hStride;
274+ splitData.wOutputInner = inputData.wStride;
275+ splitData.dOutputInner = inputData.dStride;
276+ } else {
277+ splitData.hOutputInner = inputData.hX;
278+ splitData.wOutputInner = inputData.wX;
279+ splitData.dOutputInner = inputData.dX;
280+ }
281+ 
282+ splitData.wOutputOuter = Ops::Base::CeilDiv(inputData.wX, splitData.wOutputInner);
283+ splitData.hOutputOuter = Ops::Base::CeilDiv(inputData.hX, splitData.hOutputInner);
284+ splitData.dOutputOuter = Ops::Base::CeilDiv(inputData.dX, splitData.dOutputInner);
285+ 
286+ while (splitData.hOutputInner != 1 || splitData.dOutputInner != 1 || splitData.wOutputInner > baseData.proDataNumInOneBeatT2) {
287+ if (!IsMeetTargetCoreNum() || !IsMeetUBSize()) {
288+ DynamicAdjustmentDWH();
289+ } else {
290+ return;
291+ }
292+ }
293+ 
294+ splitData.wOutputInner = std::min(inputData.wX, baseData.proDataNumInOneBeatT2);
295+ return;
296+}
297+ 
298+void MaxPool3DGradWithArgmaxNCDHWTiling::DynamicAdjustmentDWH()
299+{
300+ if (splitData.dOutputInner != 1) {
301+ splitData.dOutputOuter++;
302+ splitData.dOutputInner = Ops::Base::CeilDiv(inputData.dX, splitData.dOutputOuter);
303+ return;
304+ }
305+ if (splitData.hOutputInner != 1) {
306+ splitData.hOutputOuter++;
307+ splitData.hOutputInner = Ops::Base::CeilDiv(inputData.hX, splitData.hOutputOuter);
308+ return;
309+ }
310+ splitData.wOutputOuter++;
311+ splitData.wOutputInner = Ops::Base::CeilDiv(inputData.wX, splitData.wOutputOuter);
312+}
313+ 
314+void MaxPool3DGradWithArgmaxNCDHWTiling::SearchBestTiling()
315+{
316+ splitData.isCheckRange = 0;
317+ if (TrySplitNC()) {
318+ return;
319+ }
320+ if (baseData.isPad == 0 && baseData.isOverlap == 0) {
321+ if (TrySplitAlignD()) {
322+ return;
323+ }
324+ 
325+ if (TrySplitAlignH()) {
326+ return;
327+ }
328+ 
329+ if (TrySplitAlignW()) {
330+ return;
331+ }
332+ }
333+ splitData.isCheckRange = 1;
334+ SplitUnalignDHW();
335+ return;
336+}
337+ 
338+void MaxPool3DGradWithArgmaxNCDHWTiling::DoUBTiling()
339+{
340+ SearchBestTiling();
341+ DoBufferCalculate();
342+ splitData.wOutputOuter = Ops::Base::CeilDiv(inputData.wX, splitData.wOutputInner);
343+ int64_t tempWOutputTail = inputData.wX % splitData.wOutputInner;
344+ splitData.wOutputTail = tempWOutputTail == 0 ? splitData.wOutputInner : tempWOutputTail;
345+ 
346+ splitData.hOutputOuter = Ops::Base::CeilDiv(inputData.hX, splitData.hOutputInner);
347+ int64_t tempHOutputTail = inputData.hX % splitData.hOutputInner;
348+ splitData.hOutputTail = tempHOutputTail == 0 ? splitData.hOutputInner : tempHOutputTail;
349+ 
350+ splitData.dOutputOuter = Ops::Base::CeilDiv(inputData.dX, splitData.dOutputInner);
351+ int64_t tempDOutputTail = inputData.dX % splitData.dOutputInner;
352+ splitData.dOutputTail = tempDOutputTail == 0 ? splitData.dOutputInner : tempDOutputTail;
353+ 
354+ splitData.highAxisOuter = Ops::Base::CeilDiv(baseData.inputNCSize, splitData.highAxisInner);
355+ int64_t tempHighAxisTail = baseData.inputNCSize % splitData.highAxisInner;
356+ splitData.highAxisTail = tempHighAxisTail == 0 ? splitData.highAxisInner : tempHighAxisTail;
357+}
358+ 
359+void MaxPool3DGradWithArgmaxNCDHWTiling::DoBlockTiling()
360+{
361+ splitData.totalBaseBlockNum = splitData.highAxisOuter * splitData.hOutputOuter * splitData.wOutputOuter * splitData.dOutputOuter;
362+ splitData.normalCoreProcessNum = Ops::Base::CeilDiv(splitData.totalBaseBlockNum, baseData.totalCoreNum);
363+ splitData.usedCoreNum = Ops::Base::CeilDiv(splitData.totalBaseBlockNum, splitData.normalCoreProcessNum);
364+ splitData.tailCoreProcessNum =
365+ splitData.totalBaseBlockNum - splitData.normalCoreProcessNum * (splitData.usedCoreNum - 1);
366+}
367+ 
368+void MaxPool3DGradWithArgmaxNCDHWTiling::PrintBaseData() const
369+{
370+ OP_LOGD("MaxPool3DGradWithArgmaxNCDHW", "[MaxPool3DGradWithArgmaxNCDHW] PrintBaseData start running");
371+ 
372+ std::ostringstream info;
373+ info << "baseData.vRegSize: " << baseData.vRegSize << std::endl;
374+ info << "baseData.ubBlockSize: " << baseData.ubBlockSize << std::endl;
375+ info << "baseData.inputBytes: " << baseData.inputBytes << std::endl;
376+ info << "baseData.indexBytes: " << baseData.indexBytes << std::endl;
377+ info << "baseData.availableUb: " << baseData.availableUb << std::endl;
378+ info << "baseData.maxDataNumInOneBlock: " << baseData.maxDataNumInOneBlock << std::endl;
379+ info << "baseData.proDataNumInOneBeatT2: " << baseData.proDataNumInOneBeatT2 << std::endl;
380+ info << "baseData.totalCoreNum: " << baseData.totalCoreNum << std::endl;
381+ info << "baseData.coreUsedForBestPerformance: " << baseData.coreUsedForBestPerformance << std::endl;
382+ info << "baseData.isPad: " << baseData.isPad << std::endl;
383+ info << "baseData.isOverlap: " << baseData.isOverlap << std::endl;
384+ info << "baseData.hProBatchSize: " << baseData.hProBatchSize << std::endl;
385+ info << "baseData.wProBatchSize: " << baseData.wProBatchSize << std::endl;
386+ info << "baseData.dProBatchSize: " << baseData.dProBatchSize << std::endl;
387+ info << "baseData.inputNCSize: " << baseData.inputNCSize << std::endl;
388+ 
389+ OP_LOGI("MaxPool3DGradWithArgmaxNCDHW", "%s", info.str().c_str());
390+}
391+ 
392+void MaxPool3DGradWithArgmaxNCDHWTiling::PrintSplitData() const
393+{
394+ OP_LOGD("MaxPool3DGradWithArgmaxNCDHW", "[MaxPool3DGradWithArgmaxNCDHW] PrintSplitData start running");
395+ 
396+ std::ostringstream info;
397+ info << "splitData.isCheckRange: " << splitData.isCheckRange << std::endl;
398+ 
399+ info << "splitData.highAxisInner: " << splitData.highAxisInner << std::endl;
400+ info << "splitData.highAxisTail: " << splitData.highAxisTail << std::endl;
401+ info << "splitData.highAxisOuter: " << splitData.highAxisOuter << std::endl;
402+ 
403+ info << "splitData.hOutputInner: " << splitData.hOutputInner << std::endl;
404+ info << "splitData.hOutputTail: " << splitData.hOutputTail << std::endl;
405+ info << "splitData.hOutputOuter: " << splitData.hOutputOuter << std::endl;
406+ 
407+ info << "splitData.wOutputInner: " << splitData.wOutputInner << std::endl;
408+ info << "splitData.wOutputTail: " << splitData.wOutputTail << std::endl;
409+ info << "splitData.wOutputOuter: " << splitData.wOutputOuter << std::endl;
410+ 
411+ info << "splitData.dOutputInner: " << splitData.dOutputInner << std::endl;
412+ info << "splitData.dOutputTail: " << splitData.dOutputTail << std::endl;
413+ info << "splitData.dOutputOuter: " << splitData.dOutputOuter << std::endl;
414+ 
415+ info << "splitData.normalCoreProcessNum: " << splitData.normalCoreProcessNum << std::endl;
416+ info << "splitData.tailCoreProcessNum: " << splitData.tailCoreProcessNum << std::endl;
417+ info << "splitData.usedCoreNum: " << splitData.usedCoreNum << std::endl;
418+ info << "splitData.totalBaseBlockNum: " << splitData.totalBaseBlockNum << std::endl;
419+ 
420+ info << "splitData.outputBufferSize: " << splitData.outputBufferSize << std::endl;
421+ info << "splitData.gradBufferSize: " << splitData.gradBufferSize << std::endl;
422+ info << "splitData.argmaxBufferSize: " << splitData.argmaxBufferSize << std::endl;
423+ info << "splitData.totalBufferSize: " << splitData.totalBufferSize << std::endl;
424+ 
425+ OP_LOGI("MaxPool3DGradWithArgmaxNCDHW", "%s", info.str().c_str());
426+}
427+ 
428+void MaxPool3DGradWithArgmaxNCDHWTiling::SetTilingData()
429+{
430+ tilingData->dArgmax=inputData.dGrad;
431+ tilingData->hArgmax=inputData.hGrad;
432+ tilingData->wArgmax=inputData.wGrad;
433+ tilingData->dOutput=inputData.dX;
434+ tilingData->hOutput=inputData.hX;
435+ tilingData->wOutput=inputData.wX;
436+ tilingData->dKernel=inputData.dKernel;
437+ tilingData->hKernel=inputData.hKernel;
438+ tilingData->wKernel=inputData.wKernel;
439+ tilingData->dStride=inputData.dStride;
440+ tilingData->hStride=inputData.hStride;
441+ tilingData->wStride=inputData.wStride;
442+ tilingData->padD=inputData.dPad;
443+ tilingData->padH=inputData.hPad;
444+ tilingData->padW=inputData.wPad;
445+ tilingData->dilationD=inputData.dDilation;
446+ tilingData->dilationH=inputData.hDilation;
447+ tilingData->dilationW=inputData.wDilation;
448+ tilingData->highAxisInner=splitData.highAxisInner;
449+ tilingData->highAxisTail=splitData.highAxisTail;
450+ tilingData->highAxisOuter=splitData.highAxisOuter;
451+ tilingData->dOutputInner=splitData.dOutputInner;
452+ tilingData->dOutputTail=splitData.dOutputTail;
453+ tilingData->dOutputOuter=splitData.dOutputOuter;
454+ tilingData->hOutputInner=splitData.hOutputInner;
455+ tilingData->hOutputTail=splitData.hOutputTail;
456+ tilingData->hOutputOuter=splitData.hOutputOuter;
457+ tilingData->wOutputInner=splitData.wOutputInner;
458+ tilingData->wOutputTail=splitData.wOutputTail;
459+ tilingData->wOutputOuter=splitData.wOutputOuter;
460+ tilingData->normalCoreProcessNum=splitData.normalCoreProcessNum;
461+ tilingData->tailCoreProcessNum=splitData.tailCoreProcessNum;
462+ tilingData->usedCoreNum=splitData.usedCoreNum;
463+ tilingData->outputBufferSize=splitData.outputBufferSize;
464+ tilingData->gradBufferSize=splitData.gradBufferSize;
465+ tilingData->argmaxBufferSize=splitData.argmaxBufferSize;
466+ tilingData->dProBatchSize=baseData.dProBatchSize;
467+ tilingData->hProBatchSize=baseData.hProBatchSize;
468+ tilingData->wProBatchSize=baseData.wProBatchSize;
469+}
470+ 
471+ge::graphStatus MaxPool3DGradWithArgmaxNCDHWTiling::DoOpTiling()
472+{
473+ DoUBTiling();
474+ DoBlockTiling();
475+ SetTilingData();
476+ PrintBaseData();
477+ PrintSplitData();
478+ return ge::GRAPH_SUCCESS;
479+}
480+ 
481+ge::graphStatus MaxPool3DGradWithArgmaxNCDHWTiling::PostTiling()
482+{
483+ context_->SetTilingKey(GetTilingKey());
484+ context_->SetBlockDim(tilingData->usedCoreNum);
485+ return ge::GRAPH_SUCCESS;
486+}
487+ 
488+REGISTER_TILING_TEMPLATE("MaxPool3DGradWithArgmax", MaxPool3DGradWithArgmaxNCDHWTiling, 0);
489+ 
490+} // namespace optiling
Apooling/max_pool3d_grad_with_argmax/op_host/max_pool3d_grad_with_argmax_simt_tiling_arch35.cpp+78-0
@@ -0,0 +1,78 @@
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_pool3d_grad_with_argmax_simt_tiling_arch35.cpp
13+ * \brief
14+ */
15+#include <iostream>
16+#include "max_pool3d_grad_with_argmax_tiling_arch35.h"
17+ 
18+using namespace MaxPool3DGradWithArgmaxOp;
19+ 
20+namespace optiling {
21+ 
22+static constexpr uint64_t DCACHE_SIZE = 128 * 1024UL;
23+static constexpr int64_t MAX_THREAD_NUM = 1024;
24+ 
25+bool MaxPool3DGradWithArgmaxTilingSimt::IsCapable()
26+{
27+ return true;
28+}
29+ 
30+ge::graphStatus MaxPool3DGradWithArgmaxTilingSimt::DoOpTiling()
31+{
32+ OP_LOGD(context_->GetNodeName(), "Enter MaxPool3DGradWithArgmaxTilingSimt DoOpTiling.");
33+ tilingData_->nDim = inputData.nX;
34+ tilingData_->cDim = inputData.cX;
35+ tilingData_->dInDim = inputData.dX;
36+ tilingData_->hInDim = inputData.hX;
37+ tilingData_->wInDim = inputData.wX;
38+ tilingData_->dOutDim = inputData.dGrad;
39+ tilingData_->hOutDim = inputData.hGrad;
40+ tilingData_->wOutDim = inputData.wGrad;
41+ tilingData_->kSizeD = inputData.dKernel;
42+ tilingData_->kSizeH = inputData.hKernel;
43+ tilingData_->kSizeW = inputData.wKernel;
44+ tilingData_->strideD = inputData.dStride;
45+ tilingData_->strideH = inputData.hStride;
46+ tilingData_->strideW = inputData.wStride;
47+ tilingData_->padD = inputData.dPad;
48+ tilingData_->padH = inputData.hPad;
49+ tilingData_->padW = inputData.wPad;
50+ tilingData_->dilationD = inputData.dDilation;
51+ tilingData_->dilationH = inputData.hDilation;
52+ tilingData_->dilationW = inputData.wDilation;
53+ return ge::GRAPH_SUCCESS;
54+}
55+ 
56+ge::graphStatus MaxPool3DGradWithArgmaxTilingSimt::PostTiling()
57+{
58+ int64_t outDataCount = inputData.nX * inputData.cX * inputData.dX * inputData.hX * inputData.wX;
59+ int64_t threads = std::min(outDataCount, MAX_THREAD_NUM);
60+ int64_t blockNum = Ops::Base::CeilDiv(outDataCount, threads);
61+ blockNum = std::min(blockNum, static_cast<int64_t>(coreNum_));
62+ context_->SetBlockDim(blockNum);
63+ context_->SetLocalMemorySize(ubSize_ - DCACHE_SIZE);
64+ return ge::GRAPH_SUCCESS;
65+}
66+ 
67+uint64_t MaxPool3DGradWithArgmaxTilingSimt::GetTilingKey() const
68+{
69+ int64_t outDataCount = inputData.nX * inputData.cX * inputData.dX * inputData.hX * inputData.wX;
70+ uint32_t idxDtype = outDataCount <= static_cast<int64_t>(MAX_INT32) ? TPL_INT32 : TPL_INT64;
71+ uint32_t isChannelLast = (inputData.inputFormat == ge::Format::FORMAT_NDHWC) ? 1 : 0;
72+ uint32_t isSimt = 1;
73+ uint32_t isCheckRange = 0;
74+ return GET_TPL_TILING_KEY(idxDtype, isSimt, isChannelLast, isCheckRange);
75+}
76+ 
77+REGISTER_TILING_TEMPLATE("MaxPool3DGradWithArgmax", MaxPool3DGradWithArgmaxTilingSimt, 5);
78+} // namespace optiling
Apooling/max_pool3d_grad_with_argmax/op_host/max_pool3d_grad_with_argmax_tiling_arch35.h+206-0
@@ -0,0 +1,206 @@
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_pool3d_grad_with_argmax_tiling_arch35.h
13+ * \brief
14+ */
15+#ifndef OPS_BUILD_IN_OP_TILING_RUNTIME_MAX_POOL3D_GRAD_WITH_ARGMAX_ARCH35_H
16+#define OPS_BUILD_IN_OP_TILING_RUNTIME_MAX_POOL3D_GRAD_WITH_ARGMAX_ARCH35_H
17+ 
18+#include "log/log.h"
19+#include "register/op_impl_registry.h"
20+#include "register/tilingdata_base.h"
21+#include "tiling_base/tiling_base.h"
22+#include "tiling_base/tiling_templates_registry.h"
23+#include "util/math_util.h"
24+#include "max_pool3d_grad_with_argmax_tiling.h"
25+#include "../op_kernel/arch35/max_pool3d_grad_with_argmax_struct.h"
26+#include "../../pool_3d_common/op_host/arch32/max_pool3d_grad_tiling_constants.h"
27+#include "register/op_def_registry.h"
28+#include "tiling/tiling_api.h"
29+#include "op_common/op_host/util/platform_util.h"
30+ 
31+using namespace MaxPool3DGradWithArgmaxOp;
32+ 
33+namespace optiling {
34+using Ops::NN::Optiling::TilingBaseClass;
35+ 
36+struct MaxPool3DGradWithArgmaxInputInfo {
37+ int64_t dPad{0};
38+ int64_t hPad{0};
39+ int64_t wPad{0};
40+ int64_t dStride{1};
41+ int64_t hStride{1};
42+ int64_t wStride{1};
43+ int64_t dKernel{1};
44+ int64_t hKernel{1};
45+ int64_t wKernel{1};
46+ int64_t dDilation{1};
47+ int64_t hDilation{1};
48+ int64_t wDilation{1};
49+ int64_t nX{1};
50+ int64_t cX{1};
51+ int64_t dX{1};
52+ int64_t hX{1};
53+ int64_t wX{1};
54+ int64_t nGrad{1};
55+ int64_t cGrad{1};
56+ int64_t dGrad{1};
57+ int64_t hGrad{1};
58+ int64_t wGrad{1};
59+ bool ceilMode{false};
60+ int64_t gradShapeSize{0};
61+ ge::DataType inputDtype{ge::DataType::DT_FLOAT};
62+ ge::DataType indexDtype{ge::DataType::DT_INT32};
63+ ge::Format inputFormat{ge::Format::FORMAT_NCDHW};
64+ int64_t isInt32Meet{1};
65+};
66+ 
67+struct MaxPool3DGradWithArgmaxNCDHWBaseInfo {
68+ int64_t vRegSize{0};
69+ int64_t ubBlockSize{0};
70+ int64_t inputBytes{0};
71+ int64_t indexBytes{0};
72+ int64_t availableUb{0};
73+ int64_t totalCoreNum{0};
74+ int64_t coreUsedForBestPerformance{0};
75+ int64_t dProBatchSize{0};
76+ int64_t hProBatchSize{0};
77+ int64_t wProBatchSize{0};
78+ int64_t inputNCSize{0};
79+ int64_t maxDataNumInOneBlock{0};
80+ int64_t proDataNumInOneBeatT2{0};
81+ int64_t isPad{0};
82+ int64_t isOverlap{0};
83+};
84+ 
85+struct MaxPool3DGradWithArgmaxNCDHWSplitInfo {
86+ // DoUBTiling
87+ int64_t isCheckRange{0};
88+ 
89+ int64_t highAxisInner{0};
90+ int64_t highAxisTail{0};
91+ int64_t highAxisOuter{0};
92+ 
93+ int64_t dOutputInner{0};
94+ int64_t dOutputTail{0};
95+ int64_t dOutputOuter{0};
96+ 
97+ int64_t hOutputInner{0};
98+ int64_t hOutputTail{0};
99+ int64_t hOutputOuter{0};
100+ 
101+ int64_t wOutputInner{0};
102+ int64_t wOutputTail{0};
103+ int64_t wOutputOuter{0};
104+ 
105+ // DoBlockTiling
106+ int64_t normalCoreProcessNum{0};
107+ int64_t tailCoreProcessNum{0};
108+ int64_t usedCoreNum{0};
109+ int64_t totalBaseBlockNum{0};
110+ 
111+ // DoBufferCalculate
112+ int64_t outputBufferSize{0};
113+ int64_t gradBufferSize{0};
114+ int64_t argmaxBufferSize{0};
115+ int64_t totalBufferSize{0};
116+};
117+ 
118+class MaxPool3DGradWithArgmaxTilingBaseV35 : public TilingBaseClass {
119+public:
120+ explicit MaxPool3DGradWithArgmaxTilingBaseV35(gert::TilingContext* context) : TilingBaseClass(context)
121+ {}
122+ ~MaxPool3DGradWithArgmaxTilingBaseV35() override
123+ {}
124+ 
125+ const std::string nodeName = "MaxPool3DGradWithArgmax";
126+ MaxPool3DGradWithArgmaxTilingDataV35* tilingData_ = context_->GetTilingData<MaxPool3DGradWithArgmaxTilingDataV35>();
127+ MaxPool3DGradWithArgmaxInputInfo inputData;
128+ int64_t coreNum_{0};
129+ int64_t ubSize_{0};
130+ 
131+ bool CheckInputShape();
132+ ge::graphStatus CheckInputDtype();
133+ ge::graphStatus CheckAttrShape();
134+ ge::graphStatus CheckInputValid();
135+ ge::graphStatus SetInputParams();
136+ ge::graphStatus SetAttrParams();
137+ void SetCntTailTilingParams();
138+ void SetOtherInputParams();
139+ 
140+protected:
141+ ge::graphStatus GetShapeAttrsInfo() override;
142+ ge::graphStatus GetPlatformInfo() override;
143+ bool IsCapable() override;
144+ ge::graphStatus DoOpTiling() override;
145+ ge::graphStatus DoLibApiTiling() override;
146+ ge::graphStatus GetWorkspaceSize() override;
147+ ge::graphStatus PostTiling() override;
148+ uint64_t GetTilingKey() const override;
149+};
150+ 
151+class MaxPool3DGradWithArgmaxTilingSimt : public MaxPool3DGradWithArgmaxTilingBaseV35 {
152+public:
153+ explicit MaxPool3DGradWithArgmaxTilingSimt(gert::TilingContext* context)
154+ : MaxPool3DGradWithArgmaxTilingBaseV35(context)
155+ {}
156+ ~MaxPool3DGradWithArgmaxTilingSimt() override
157+ {}
158+ 
159+protected:
160+ bool IsCapable() override;
161+ ge::graphStatus DoOpTiling() override;
162+ ge::graphStatus PostTiling() override;
163+ uint64_t GetTilingKey() const override;
164+};
165+ 
166+class MaxPool3DGradWithArgmaxNCDHWTiling : public MaxPool3DGradWithArgmaxTilingBaseV35 {
167+public:
168+ explicit MaxPool3DGradWithArgmaxNCDHWTiling(gert::TilingContext* context)
169+ : MaxPool3DGradWithArgmaxTilingBaseV35(context)
170+ {}
171+ 
172+ ~MaxPool3DGradWithArgmaxNCDHWTiling() override
173+ {}
174+ MaxPool3DGradWithArgmaxOp::MaxPool3DGradWithArgmaxNCDHWTilingData* tilingData =
175+ context_->GetTilingData<MaxPool3DGradWithArgmaxOp::MaxPool3DGradWithArgmaxNCDHWTilingData>();
176+ 
177+private:
178+ void DoUBTiling();
179+ void InitializationVars();
180+ bool TrySplitNC();
181+ bool TrySplitAlignD();
182+ bool TrySplitAlignH();
183+ bool TrySplitAlignW();
184+ void SplitUnalignDHW();
185+ bool IsMeetTargetCoreNum() const;
186+ bool IsMeetUBSize();
187+ void SearchBestTiling();
188+ void DynamicAdjustmentDWH();
189+ void SetTilingData();
190+ uint64_t GetTilingKey() const override;
191+ void PrintBaseData() const;
192+ void PrintSplitData() const;
193+ void DoBlockTiling();
194+ void DoBufferCalculate();
195+ bool IsCapable() override;
196+ ge::graphStatus DoOpTiling() override;
197+ ge::graphStatus PostTiling() override;
198+ 
199+ MaxPool3DGradWithArgmaxNCDHWBaseInfo baseData;
200+ MaxPool3DGradWithArgmaxNCDHWSplitInfo splitData;
201+};
202+ 
203+ 
204+} // namespace optiling
205+ 
206+#endif
Apooling/max_pool3d_grad_with_argmax/op_host/max_pool3d_grad_with_argmax_tiling_base_arch35.cpp+402-0
@@ -0,0 +1,402 @@
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_pool3d_grad_with_argmax_tiling_base_arch35.cpp
13+ * \brief
14+ */
15+ 
16+#include "error_util.h"
17+#include "max_pool3d_grad_with_argmax_tiling_arch35.h"
18+ 
19+namespace optiling {
20+ 
21+constexpr size_t CDHW_DIM_NUM = 4U;
22+constexpr size_t DATA_FORMAT_ATTR_INDEX = 5U;
23+constexpr size_t C_DIM_OFFSET = 4; // pos = dim - offset
24+constexpr size_t D_DIM_OFFSET = 3;
25+constexpr size_t H_DIM_OFFSET = 2;
26+constexpr size_t W_DIM_OFFSET = 1;
27+constexpr size_t D_ATTR_INDEX = 0;
28+constexpr size_t H_ATTR_INDEX = 1;
29+constexpr size_t W_ATTR_INDEX = 2;
30+constexpr int64_t WS_SYS_SIZE = 16 * 1024 * 1024;
31+static const gert::Shape g_vec_1_shape = {1};
32+ 
33+static const gert::Shape &EnsureNotScalar(const gert::Shape &inShape) {
34+ if (inShape.IsScalar()) {
35+ return g_vec_1_shape;
36+ }
37+ return inShape;
38+}
39+ 
40+static inline bool IsGreaterThanInt32Max(const MaxPool3DGradWithArgmaxInputInfo& inputData)
41+{
42+ if (inputData.indexDtype == ge::DataType::DT_INT32) {
43+ return false;
44+ }
45+ 
46+ int64_t cubeSize = inputData.dX * inputData.hX * inputData.wX;
47+ return cubeSize > static_cast<int64_t>(INT32_MAX);
48+}
49+ 
50+bool MaxPool3DGradWithArgmaxTilingBaseV35::CheckInputShape()
51+{
52+ const gert::StorageShape* xShape = context_->GetInputShape(X_INDEX);
53+ const gert::StorageShape* gradShape = context_->GetInputShape(GRAD_INDEX);
54+ const gert::StorageShape* argmaxShape = context_->GetInputShape(ARGMAX_INDEX);
55+ size_t xDimNum = EnsureNotScalar(xShape->GetStorageShape()).GetDimNum();
56+ size_t gradDimNum = EnsureNotScalar(gradShape->GetStorageShape()).GetDimNum();
57+ size_t argmaxDimNum = EnsureNotScalar(argmaxShape->GetStorageShape()).GetDimNum();
58+ auto attrs = context_->GetAttrs();
59+ OP_CHECK_NULL_WITH_CONTEXT(context_, attrs);
60+ const char* data_format = attrs->GetAttrPointer<char>(DATA_FORMAT_ATTR_INDEX);
61+ OPS_CHECK_NULL_WITH_CONTEXT(context_, data_format);
62+ std::string data_formatStr = data_format;
63+ 
64+ // data_format should be NCDHW or NDHWC
65+ OP_CHECK_IF(!(data_formatStr == "NCDHW" || data_formatStr == "NDHWC"),
66+ OP_LOGE(context_->GetNodeName(), "ATTR data_format is %s ,expect [NDHWC] or [NCDHW].", data_format),
67+ return false);
68+ 
69+ // xDimNum should be 5 or 4
70+ OP_CHECK_IF(((xDimNum != NCDHW_DIM_NUM) || (gradDimNum != NCDHW_DIM_NUM) || (argmaxDimNum != NCDHW_DIM_NUM)) &&
71+ ((xDimNum != CDHW_DIM_NUM) || (gradDimNum != CDHW_DIM_NUM) || (argmaxDimNum != CDHW_DIM_NUM)),
72+ OP_LOGE(context_->GetNodeName(),
73+ "Input dim num should equal = %lu or %lu, actual is xDim: %lu, gradDim: %lu, argmaxDim: %lu.",
74+ NCDHW_DIM_NUM, CDHW_DIM_NUM, xDimNum, gradDimNum, argmaxDimNum),
75+ return false);
76+ for (uint32_t i = 0; i < xDimNum; i++) {
77+ OP_CHECK_IF(xShape->GetStorageShape().GetDim(i) == 0,
78+ OP_LOGE(context_->GetNodeName(), "Input x shape can not be 0."), return false);
79+ }
80+ 
81+ // gradShape&argmaxShape's shape should be equal
82+ for (size_t i = 0; i < xDimNum; i++) {
83+ uint64_t gradDimValue = gradShape->GetStorageShape().GetDim(i);
84+ uint64_t argmaxDimValue = argmaxShape->GetStorageShape().GetDim(i);
85+ OP_CHECK_IF(gradDimValue != argmaxDimValue,
86+ OP_LOGE(context_->GetNodeName(),
87+ "Input dim check invalid, grad[%lu] is %lu, argmax[%lu] is %lu, not equal.", i,
88+ gradDimValue, i, argmaxDimValue),
89+ return false);
90+ }
91+ 
92+ // Input NCDim should be equal
93+ uint32_t cPosIdx = (data_formatStr == "NDHWC") ? xDimNum - 1 : xDimNum - 4;
94+ uint64_t xNDim = (xDimNum == CDHW_DIM_NUM) ? 1 : xShape->GetStorageShape().GetDim(0);
95+ uint64_t gradNDim = (gradDimNum == CDHW_DIM_NUM) ? 1 : gradShape->GetStorageShape().GetDim(0);
96+ uint64_t xCDim = xShape->GetStorageShape().GetDim(cPosIdx);
97+ uint64_t gradCDim = gradShape->GetStorageShape().GetDim(cPosIdx);
98+ OP_CHECK_IF((xNDim != gradNDim) || (xCDim != gradCDim),
99+ OP_LOGE(context_->GetNodeName(), "Input N,C dim check invalid, grad(%lu,%lu), x(%lu,%lu), not equal.",
100+ gradNDim, gradCDim, xNDim, xCDim),
101+ return false);
102+ 
103+ return true;
104+}
105+ 
106+ge::graphStatus MaxPool3DGradWithArgmaxTilingBaseV35::CheckInputDtype()
107+{
108+ OP_CHECK_NULL_WITH_CONTEXT(context_, context_->GetInputDesc(X_INDEX));
109+ OP_CHECK_NULL_WITH_CONTEXT(context_, context_->GetInputDesc(GRAD_INDEX));
110+ OP_CHECK_NULL_WITH_CONTEXT(context_, context_->GetInputDesc(ARGMAX_INDEX));
111+ OP_CHECK_NULL_WITH_CONTEXT(context_, context_->GetOutputDesc(Y_INDEX));
112+ auto xDataType = context_->GetInputDesc(X_INDEX)->GetDataType();
113+ auto gradDataType = context_->GetInputDesc(GRAD_INDEX)->GetDataType();
114+ auto argmaxDataType = context_->GetInputDesc(ARGMAX_INDEX)->GetDataType();
115+ auto yOutDataType = context_->GetOutputDesc(Y_INDEX)->GetDataType();
116+ 
117+ OP_CHECK_IF(xDataType != gradDataType,
118+ OP_LOGE(context_->GetNodeName(), "Data type invalid, x data type not equal grad data type."),
119+ return ge::GRAPH_FAILED);
120+ OP_CHECK_IF(xDataType != yOutDataType,
121+ OP_LOGE(context_->GetNodeName(), "Data type invalid, x data type not equal y data type."),
122+ return ge::GRAPH_FAILED);
123+ OP_CHECK_IF((xDataType != ge::DT_FLOAT) && (xDataType != ge::DT_FLOAT16) && (xDataType != ge::DT_BF16),
124+ OP_LOGE(context_->GetNodeName(), "Data type invalid, x data type not fp32/fp16/bf16."),
125+ return ge::GRAPH_FAILED);
126+ OP_CHECK_IF((argmaxDataType != ge::DT_INT32) && (argmaxDataType != ge::DT_INT64),
127+ OP_LOGE(context_->GetNodeName(), "Data type invalid, argmax data type not equal int32/int64."),
128+ return ge::GRAPH_FAILED);
129+ return ge::GRAPH_SUCCESS;
130+}
131+ 
132+ge::graphStatus MaxPool3DGradWithArgmaxTilingBaseV35::CheckAttrShape()
133+{
134+ auto attrs = context_->GetAttrs();
135+ OP_CHECK_NULL_WITH_CONTEXT(context_, attrs);
136+ int32_t kSizeDimNum = attrs->GetListInt(KSIZE_ATTR_INDEX)->GetSize();
137+ int32_t stridesDimNum = attrs->GetListInt(STRIDES_ATTR_INDEX)->GetSize();
138+ int32_t padsDimNum = attrs->GetListInt(PADS_ATTR_INDEX)->GetSize();
139+ int32_t dilationsDimNum = attrs->GetListInt(DILATION_ATTR_INDEX)->GetSize();
140+ 
141+ // Check attr dim num
142+ OP_CHECK_IF((kSizeDimNum != DHW_DIM_NUM) && (kSizeDimNum != 1),
143+ OP_LOGE(context_->GetNodeName(), "Attr kSize dim num invalid, dim num should equal 3 or 1."),
144+ return ge::GRAPH_FAILED);
145+ OP_CHECK_IF((stridesDimNum != DHW_DIM_NUM) && (stridesDimNum != 1) && (stridesDimNum != 0),
146+ OP_LOGE(context_->GetNodeName(), "Attr strides dim num invalid, dim num should equal 3 or 1 or 0."),
147+ return ge::GRAPH_FAILED);
148+ OP_CHECK_IF((padsDimNum != DHW_DIM_NUM) && (padsDimNum != 1),
149+ OP_LOGE(context_->GetNodeName(), "Attr pads dim num invalid, dim num should equal 3 or 1."),
150+ return ge::GRAPH_FAILED);
151+ OP_CHECK_IF((dilationsDimNum != DHW_DIM_NUM) && (dilationsDimNum != 1),
152+ OP_LOGE(context_->GetNodeName(), "Attr dilations dim num invalid, dim num should equal 3 or 1."),
153+ return ge::GRAPH_FAILED);
154+ 
155+ // Check attr value bigger than 0
156+ auto kSizeVector = attrs->GetListInt(KSIZE_ATTR_INDEX)->GetData();
157+ auto stridesVector = attrs->GetListInt(STRIDES_ATTR_INDEX)->GetData();
158+ auto padsVector = attrs->GetListInt(PADS_ATTR_INDEX)->GetData();
159+ auto dilationsVector = attrs->GetListInt(DILATION_ATTR_INDEX)->GetData();
160+ for (uint32_t i = 0; i < (uint32_t)kSizeDimNum; i++) {
161+ OP_CHECK_IF((kSizeVector[i] <= 0),
162+ OP_LOGE(context_->GetNodeName(), "Attr value invalid, kSize[%u] is %ld, should bigger than 0.",
163+ i, kSizeVector[i]),
164+ return ge::GRAPH_FAILED);
165+ }
166+ for (uint32_t i = 0; i < (uint32_t)stridesDimNum; i++) {
167+ OP_CHECK_IF((stridesVector[i] <= 0),
168+ OP_LOGE(context_->GetNodeName(), "Attr value invalid, strides[%u] is %ld, should bigger than 0.",
169+ i, stridesVector[i]),
170+ return ge::GRAPH_FAILED);
171+ }
172+ for (uint32_t i = 0; i < (uint32_t)padsDimNum; i++) {
173+ OP_CHECK_IF((padsVector[i] < 0),
174+ OP_LOGE(context_->GetNodeName(), "Attr value invalid, pads[%u] is %ld, should bigger or equal 0.",
175+ i, padsVector[i]),
176+ return ge::GRAPH_FAILED);
177+ }
178+ for (uint32_t i = 0; i < (uint32_t)dilationsDimNum; i++) {
179+ OP_CHECK_IF((dilationsVector[i] <= 0),
180+ OP_LOGE(context_->GetNodeName(), "Attr value invalid, dilations[%u] is %ld, should bigger than 0.",
181+ i, dilationsVector[i]),
182+ return ge::GRAPH_FAILED);
183+ }
184+ return ge::GRAPH_SUCCESS;
185+}
186+ 
187+ge::graphStatus MaxPool3DGradWithArgmaxTilingBaseV35::SetInputParams()
188+{
189+ const gert::Shape xShape = context_->GetInputShape(X_INDEX)->GetStorageShape();
190+ const gert::Shape gradShape = context_->GetInputShape(GRAD_INDEX)->GetStorageShape();
191+ size_t xDimNum = xShape.GetDimNum();
192+ auto attrs = context_->GetAttrs();
193+ OPS_CHECK_NULL_WITH_CONTEXT(context_, attrs);
194+ const char* data_format = attrs->GetAttrPointer<char>(DATA_FORMAT_ATTR_INDEX);
195+ std::string data_formatStr = data_format;
196+ 
197+ uint32_t cPosIdx = xDimNum - C_DIM_OFFSET;
198+ uint32_t dPosIdx = xDimNum - D_DIM_OFFSET;
199+ uint32_t hPosIdx = xDimNum - H_DIM_OFFSET;
200+ uint32_t wPosIdx = xDimNum - W_DIM_OFFSET;
201+
202+ inputData.inputFormat = ge::Format::FORMAT_NCDHW;
203+ 
204+ if (data_formatStr == "NDHWC") {
205+ inputData.inputFormat = ge::Format::FORMAT_NDHWC;
206+ dPosIdx = dPosIdx - 1;
207+ hPosIdx = hPosIdx - 1;
208+ wPosIdx = wPosIdx - 1;
209+ cPosIdx = xDimNum - 1;
210+ }
211+ 
212+ inputData.nX = (xDimNum == CDHW_DIM_NUM) ? 1 : xShape.GetDim(0);
213+ inputData.cX = xShape.GetDim(cPosIdx);
214+ inputData.dX = xShape.GetDim(dPosIdx);
215+ inputData.hX = xShape.GetDim(hPosIdx);
216+ inputData.wX = xShape.GetDim(wPosIdx);
217+ inputData.nGrad = (xDimNum == CDHW_DIM_NUM) ? 1 : gradShape.GetDim(0);
218+ inputData.cGrad = gradShape.GetDim(cPosIdx);
219+ inputData.dGrad = gradShape.GetDim(dPosIdx);
220+ inputData.hGrad = gradShape.GetDim(hPosIdx);
221+ inputData.wGrad = gradShape.GetDim(wPosIdx);
222+ inputData.gradShapeSize = gradShape.GetShapeSize();
223+ return ge::GRAPH_SUCCESS;
224+}
225+ 
226+ge::graphStatus MaxPool3DGradWithArgmaxTilingBaseV35::SetAttrParams()
227+{
228+ auto attrs = context_->GetAttrs();
229+ OP_CHECK_NULL_WITH_CONTEXT(context_, attrs);
230+ int32_t kSizeDimNum = attrs->GetListInt(KSIZE_ATTR_INDEX)->GetSize();
231+ int32_t stridesDimNum = attrs->GetListInt(STRIDES_ATTR_INDEX)->GetSize();
232+ int32_t padsDimNum = attrs->GetListInt(PADS_ATTR_INDEX)->GetSize();
233+ int32_t dilationsDimNum = attrs->GetListInt(DILATION_ATTR_INDEX)->GetSize();
234+ auto kSizeVector = attrs->GetListInt(KSIZE_ATTR_INDEX)->GetData();
235+ auto stridesVector = attrs->GetListInt(STRIDES_ATTR_INDEX)->GetData();
236+ auto padsVector = attrs->GetListInt(PADS_ATTR_INDEX)->GetData();
237+ auto dilationsVector = attrs->GetListInt(DILATION_ATTR_INDEX)->GetData();
238+ bool ceilMode = *attrs->GetBool(CEIL_MODE_ATTR_INDEX);
239+ inputData.ceilMode = ceilMode;
240+ inputData.dKernel = kSizeVector[D_ATTR_INDEX];
241+ inputData.hKernel = (kSizeDimNum == 1) ? inputData.dKernel : kSizeVector[H_ATTR_INDEX];
242+ inputData.wKernel = (kSizeDimNum == 1) ? inputData.dKernel : kSizeVector[W_ATTR_INDEX];
243+ if (stridesDimNum == 0) {
244+ inputData.dStride = inputData.dKernel;
245+ inputData.hStride = inputData.hKernel;
246+ inputData.wStride = inputData.wKernel;
247+ } else {
248+ inputData.dStride = stridesVector[D_ATTR_INDEX];
249+ inputData.hStride = (stridesDimNum == 1) ? inputData.dStride : stridesVector[H_ATTR_INDEX];
250+ inputData.wStride = (stridesDimNum == 1) ? inputData.dStride : stridesVector[W_ATTR_INDEX];
251+ }
252+ inputData.dPad = padsVector[D_ATTR_INDEX];
253+ inputData.hPad = (padsDimNum == 1) ? inputData.dPad : padsVector[H_ATTR_INDEX];
254+ inputData.wPad = (padsDimNum == 1) ? inputData.dPad : padsVector[W_ATTR_INDEX];
255+ inputData.dDilation = dilationsVector[D_ATTR_INDEX];
256+ inputData.hDilation = (dilationsDimNum == 1) ? inputData.dDilation : dilationsVector[H_ATTR_INDEX];
257+ inputData.wDilation = (dilationsDimNum == 1) ? inputData.dDilation : dilationsVector[W_ATTR_INDEX];
258+ return ge::GRAPH_SUCCESS;
259+}
260+ 
261+ge::graphStatus MaxPool3DGradWithArgmaxTilingBaseV35::CheckInputValid()
262+{
263+ const uint64_t kd = inputData.dKernel;
264+ const uint64_t kh = inputData.hKernel;
265+ const uint64_t kw = inputData.wKernel;
266+ const uint64_t sd = inputData.dStride;
267+ const uint64_t sh = inputData.hStride;
268+ const uint64_t sw = inputData.wStride;
269+ const uint64_t pDTop = inputData.dPad;
270+ const uint64_t pHTop = inputData.hPad;
271+ const uint64_t pWTop = inputData.wPad;
272+ const uint64_t dilationD = inputData.dDilation;
273+ const uint64_t dilationH = inputData.hDilation;
274+ const uint64_t dilationW = inputData.wDilation;
275+ 
276+ // check 1
277+ OP_CHECK_IF((pDTop > (kd / 2)) || (pHTop > (kh / 2)) || (pWTop > (kw / 2)),
278+ OP_LOGE(context_->GetNodeName(), "Attr size invalid, padSize should smaller than kernelSize div 2"),
279+ return ge::GRAPH_FAILED);
280+ // check 2
281+ OP_CHECK_IF((pDTop > ((kd - 1) * dilationD + 1) / 2) || (pHTop > ((kh - 1) * dilationH + 1) / 2) ||
282+ (pWTop > ((kw - 1) * dilationW + 1) / 2),
283+ OP_LOGE(context_->GetNodeName(),
284+ "Attr size invalid, padSize should smaller than ((kernelSize - 1) * dilation + 1) / 2."),
285+ return ge::GRAPH_FAILED);
286+ // check 3
287+ // Check outerDim invaild
288+ int64_t doExpected, hoExpected, woExpected;
289+ if (inputData.ceilMode) {
290+ doExpected = Ops::Base::CeilDiv((inputData.dX + NUM_TWO * pDTop + sd - dilationD * (kd - 1) - 1), sd);
291+ hoExpected = Ops::Base::CeilDiv((inputData.hX + NUM_TWO * pHTop + sh - dilationH * (kh - 1) - 1), sh);
292+ woExpected = Ops::Base::CeilDiv((inputData.wX + NUM_TWO * pWTop + sw - dilationW * (kw - 1) - 1), sw);
293+ } else {
294+ doExpected = (inputData.dX + NUM_TWO * pDTop + sd - dilationD * (kd - 1) - 1) / sd;
295+ hoExpected = (inputData.hX + NUM_TWO * pHTop + sh - dilationH * (kh - 1) - 1) / sh;
296+ woExpected = (inputData.wX + NUM_TWO * pWTop + sw - dilationW * (kw - 1) - 1) / sw;
297+ }
298+ doExpected = ((doExpected - 1) * sd >= inputData.dX + pDTop) ? doExpected - 1 : doExpected;
299+ hoExpected = ((hoExpected - 1) * sh >= inputData.hX + pHTop) ? hoExpected - 1 : hoExpected;
300+ woExpected = ((woExpected - 1) * sw >= inputData.wX + pWTop) ? woExpected - 1 : woExpected;
301+ OP_CHECK_IF(
302+ (doExpected <= 0) || (doExpected != inputData.dGrad) || (hoExpected <= 0) || (hoExpected != inputData.hGrad) ||
303+ (woExpected <= 0) || (woExpected != inputData.wGrad),
304+ OP_LOGE(context_->GetNodeName(), "OuterDim size invalid, doExpected: %ld, hoExpected: %ld, woExpected: %ld.",
305+ doExpected, hoExpected, woExpected),
306+ return ge::GRAPH_FAILED);
307+ 
308+ return ge::GRAPH_SUCCESS;
309+}
310+ 
311+void MaxPool3DGradWithArgmaxTilingBaseV35::SetOtherInputParams()
312+{
313+ inputData.inputDtype = context_->GetInputDesc(X_INDEX)->GetDataType();
314+ inputData.indexDtype = context_->GetInputDesc(ARGMAX_INDEX)->GetDataType();
315+ inputData.isInt32Meet = IsGreaterThanInt32Max(inputData) ? 0 : 1;
316+}
317+ 
318+ge::graphStatus MaxPool3DGradWithArgmaxTilingBaseV35::GetShapeAttrsInfo()
319+{
320+ auto platformInfo = context_->GetPlatformInfo();
321+ OP_CHECK_NULL_WITH_CONTEXT(context_, platformInfo);
322+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo);
323+ auto socVersion = ascendcPlatform.GetSocVersion();
324+ if (socVersion != platform_ascendc::SocVersion::ASCEND950) {
325+ // Skip the current template
326+ return ge::GRAPH_PARAM_INVALID;
327+ }
328+ 
329+ OP_LOGD(context_->GetNodeName(), "Enter MaxPool3DGradWithArgmaxTilingBaseV35 GetShapeAttrsInfo.");
330+ OP_CHECK_IF(ge::GRAPH_SUCCESS != CheckInputDtype(), OP_LOGE(context_->GetNodeName(), "The input dtype is invalid."),
331+ return ge::GRAPH_FAILED);
332+ OP_CHECK_IF(!CheckInputShape(), OP_LOGE(context_->GetNodeName(), "The input relationship is invalid."),
333+ return ge::GRAPH_FAILED);
334+ OP_CHECK_IF(ge::GRAPH_SUCCESS != CheckAttrShape(), OP_LOGE(context_->GetNodeName(), "The attr shape is invalid."),
335+ return ge::GRAPH_FAILED);
336+ OP_CHECK_IF(ge::GRAPH_SUCCESS != SetInputParams(), OP_LOGE(context_->GetNodeName(), "Set input shape failed."),
337+ return ge::GRAPH_FAILED);
338+ OP_CHECK_IF(ge::GRAPH_SUCCESS != SetAttrParams(), OP_LOGE(context_->GetNodeName(), "Set attr shape failed."),
339+ return ge::GRAPH_FAILED);
340+ OP_CHECK_IF(ge::GRAPH_SUCCESS != CheckInputValid(), OP_LOGE(context_->GetNodeName(), "The input shape is invalid."),
341+ return ge::GRAPH_FAILED);
342+ SetOtherInputParams();
343+ return ge::GRAPH_SUCCESS;
344+}
345+ 
346+bool MaxPool3DGradWithArgmaxTilingBaseV35::IsCapable()
347+{
348+ return false;
349+}
350+ 
351+ge::graphStatus MaxPool3DGradWithArgmaxTilingBaseV35::DoOpTiling()
352+{
353+ return ge::GRAPH_SUCCESS;
354+}
355+ 
356+ge::graphStatus MaxPool3DGradWithArgmaxTilingBaseV35::DoLibApiTiling()
357+{
358+ return ge::GRAPH_SUCCESS;
359+}
360+ 
361+ge::graphStatus MaxPool3DGradWithArgmaxTilingBaseV35::GetPlatformInfo()
362+{
363+ auto platformPtr = context_->GetPlatformInfo();
364+ if (platformPtr == nullptr) {
365+ auto compileInfoPtr =
366+ reinterpret_cast<const Tiling4MaxPool3DGradWithArgmaxCompileInfo*>(context_->GetCompileInfo());
367+ OP_CHECK_IF(compileInfoPtr == nullptr, OP_LOGE(context_->GetNodeName(), "compile info is null"),
368+ return ge::GRAPH_FAILED);
369+ coreNum_ = compileInfoPtr->totalCoreNum;
370+ ubSize_ = compileInfoPtr->maxUbSize;
371+ } else {
372+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformPtr);
373+ coreNum_ = ascendcPlatform.GetCoreNumAiv();
374+ 
375+ uint64_t ubSizePlatform;
376+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSizePlatform);
377+ ubSize_ = static_cast<int64_t>(ubSizePlatform);
378+ }
379+ 
380+ OP_CHECK_IF(coreNum_ == 0, OP_LOGE(context_->GetNodeName(), "coreNum is 0"), return ge::GRAPH_FAILED);
381+ return ge::GRAPH_SUCCESS;
382+}
383+ 
384+ge::graphStatus MaxPool3DGradWithArgmaxTilingBaseV35::GetWorkspaceSize()
385+{
386+ auto sys_workspace = WS_SYS_SIZE;
387+ size_t* currentWorkspace = context_->GetWorkspaceSizes(1);
388+ OP_CHECK_NULL_WITH_CONTEXT(context_, currentWorkspace);
389+ currentWorkspace[0] = sys_workspace;
390+ return ge::GRAPH_SUCCESS;
391+}
392+ 
393+ge::graphStatus MaxPool3DGradWithArgmaxTilingBaseV35::PostTiling()
394+{
395+ return ge::GRAPH_SUCCESS;
396+}
397+ 
398+uint64_t MaxPool3DGradWithArgmaxTilingBaseV35::GetTilingKey() const
399+{
400+ return 0;
401+}
402+} // namespace optiling
Apooling/max_pool3d_grad_with_argmax/op_kernel/arch35/max_pool3d_grad_with_argmax_simd.h+1000-0文件内容审核中，请稍后刷新重试
Apooling/max_pool3d_grad_with_argmax/op_kernel/arch35/max_pool3d_grad_with_argmax_simt.h+289-0文件内容审核中，请稍后刷新重试
Apooling/max_pool3d_grad_with_argmax/op_kernel/arch35/max_pool3d_grad_with_argmax_struct.h+118-0
@@ -0,0 +1,118 @@
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_pool3d_grad_with_argmax_struct.h
13+ * \brief max_pool3d_grad_with_argmax_struct
14+ */
15+#ifndef MAX_POOL3D_GRAD_WITH_ARGMAX_STRUCT_H_
16+#define MAX_POOL3D_GRAD_WITH_ARGMAX_STRUCT_H_
17+ 
18+#include <cstdint>
19+#include "kernel_tiling/kernel_tiling.h"
20+#include "ascendc/host_api/tiling/template_argument.h"
21+ 
22+namespace MaxPool3DGradWithArgmaxOp {
23+ 
24+#define TPL_INT32 1
25+#define TPL_INT64 2
26+ 
27+ASCENDC_TPL_ARGS_DECL(MaxPool3DGradWithArgmax,
28+ ASCENDC_TPL_DTYPE_DECL(INDEX_DTYPE, TPL_INT32, TPL_INT64),
29+ ASCENDC_TPL_BOOL_DECL(IS_SIMT, 0, 1),
30+ ASCENDC_TPL_BOOL_DECL(IS_CHANNEL_LAST, 0, 1),
31+ ASCENDC_TPL_BOOL_DECL(IS_CHECK_RANGE, 0, 1)
32+);
33+ 
34+ASCENDC_TPL_SEL(
35+ ASCENDC_TPL_ARGS_SEL(
36+ ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_AIV_ONLY),
37+ ASCENDC_TPL_DTYPE_SEL(INDEX_DTYPE, TPL_INT32),
38+ ASCENDC_TPL_BOOL_SEL(IS_SIMT, 0, 1),
39+ ASCENDC_TPL_BOOL_SEL(IS_CHANNEL_LAST, 0, 1),
40+ ASCENDC_TPL_BOOL_SEL(IS_CHECK_RANGE, 0, 1)
41+ ),
42+ 
43+ ASCENDC_TPL_ARGS_SEL(
44+ ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_AIV_ONLY),
45+ ASCENDC_TPL_DTYPE_SEL(INDEX_DTYPE, TPL_INT64),
46+ ASCENDC_TPL_BOOL_SEL(IS_SIMT, 0, 1),
47+ ASCENDC_TPL_BOOL_SEL(IS_CHANNEL_LAST, 0, 1),
48+ ASCENDC_TPL_BOOL_SEL(IS_CHECK_RANGE, 0, 1)
49+ )
50+);
51+ 
52+struct MaxPool3DGradWithArgmaxTilingDataV35 {
53+ int64_t nDim = 0;
54+ int64_t cDim = 0;
55+ int64_t dInDim = 0;
56+ int64_t hInDim = 0;
57+ int64_t wInDim = 0;
58+ int64_t dOutDim = 0;
59+ int64_t hOutDim = 0;
60+ int64_t wOutDim = 0;
61+ int64_t kSizeD = 0;
62+ int64_t kSizeH = 0;
63+ int64_t kSizeW = 0;
64+ int64_t strideD = 0;
65+ int64_t strideH = 0;
66+ int64_t strideW = 0;
67+ int64_t padD = 0;
68+ int64_t padH = 0;
69+ int64_t padW = 0;
70+ int64_t dilationD = 0;
71+ int64_t dilationH = 0;
72+ int64_t dilationW = 0;
73+};
74+ 
75+struct MaxPool3DGradWithArgmaxNCDHWTilingData {
76+ int64_t dArgmax = 0;
77+ int64_t hArgmax = 0;
78+ int64_t wArgmax = 0;
79+ int64_t dOutput = 0;
80+ int64_t hOutput = 0;
81+ int64_t wOutput = 0;
82+ int64_t dKernel = 0;
83+ int64_t hKernel = 0;
84+ int64_t wKernel = 0;
85+ int64_t dStride = 0;
86+ int64_t hStride = 0;
87+ int64_t wStride = 0;
88+ int64_t padD = 0;
89+ int64_t padH = 0;
90+ int64_t padW = 0;
91+ int64_t dilationD = 0;
92+ int64_t dilationH = 0;
93+ int64_t dilationW = 0;
94+ int64_t highAxisInner = 0;
95+ int64_t highAxisTail = 0;
96+ int64_t highAxisOuter = 0;
97+ int64_t dOutputInner = 0;
98+ int64_t dOutputTail = 0;
99+ int64_t dOutputOuter = 0;
100+ int64_t hOutputInner = 0;
101+ int64_t hOutputTail = 0;
102+ int64_t hOutputOuter = 0;
103+ int64_t wOutputInner = 0;
104+ int64_t wOutputTail = 0;
105+ int64_t wOutputOuter = 0;
106+ int64_t normalCoreProcessNum = 0;
107+ int64_t tailCoreProcessNum = 0;
108+ int64_t usedCoreNum = 0;
109+ int64_t outputBufferSize = 0;
110+ int64_t gradBufferSize = 0;
111+ int64_t argmaxBufferSize = 0;
112+ int64_t dProBatchSize = 0;
113+ int64_t hProBatchSize = 0;
114+ int64_t wProBatchSize = 0;
115+};
116+ 
117+} // namespace MaxPool3DGradWithArgmaxOp
118+#endif // MAX_POOL3D_GRAD_WITH_ARGMAX_STRUCT_H_
Mpooling/max_pool3d_grad_with_argmax/op_kernel/max_pool3d_grad_with_argmax.cpp+142-88
@@ -1,88 +1,142 @@
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-/* !11+/* !
12- * \file max_pool3d_grad_with_argmax.cpp12+ * \file max_pool3d_grad_with_argmax.cpp
13- * \brief13+ * \brief
14- */14+ */
15- 15+#if defined(__CCE_AICORE__) && __CCE_AICORE__ == 310
16-#include "max_pool3d_grad_with_argmax_normal.h"16+#include "arch35/max_pool3d_grad_with_argmax_simt.h"
17-#include "max_pool3d_grad_with_argmax_scatter.h"17+#include "arch35/max_pool3d_grad_with_argmax_simd_impl.h"
18-#include "max_pool3d_grad_with_argmax_scatter_overlap.h"18+#else
19-#include "max_pool3d_grad_with_argmax_cutk_d.h"19+#include "max_pool3d_grad_with_argmax_normal.h"
20-#include "max_pool3d_grad_with_argmax_cutk_dh.h"20+#include "max_pool3d_grad_with_argmax_scatter.h"
21-#include "max_pool3d_grad_with_argmax_cutk_dhw.h"21+#include "max_pool3d_grad_with_argmax_scatter_overlap.h"
22- 22+#include "max_pool3d_grad_with_argmax_cutk_d.h"
23-using namespace MaxPool3DGradWithArgmax;23+#include "max_pool3d_grad_with_argmax_cutk_dh.h"
24- 24+#include "max_pool3d_grad_with_argmax_cutk_dhw.h"
25-#define GENERAL_OP_IMPL(templateClass, ...) \25+ 
26- do { \26+using namespace MaxPool3DGradWithArgmax;
27- GET_TILING_DATA(tilingData, tiling); \27+ 
28- templateClass<__VA_ARGS__> op(&pipe); \28+#define GENERAL_OP_IMPL(templateClass, ...) \
29- op.Init(x, grad, argmax, y, workspace, &tilingData); \29+ do { \
30- op.Process(); \30+ GET_TILING_DATA(tilingData, tiling); \
31- } while (0)31+ templateClass<__VA_ARGS__> op(&pipe); \
32- 32+ op.Init(x, grad, argmax, y, workspace, &tilingData); \
33-#define GENERAL_OP_IMPL_CUTNC(templateClass, ...) \33+ op.Process(); \
34- do { \34+ } while (0)
35- GET_TILING_DATA(tilingData, tiling); \35+ 
36- templateClass<__VA_ARGS__> op(&pipe); \36+#define GENERAL_OP_IMPL_CUTNC(templateClass, ...) \
37- op.Init(x, grad, argmax, y, workspace, &tilingData); \37+ do { \
38- op.ProcessCutNc(); \38+ GET_TILING_DATA(tilingData, tiling); \
39- } while (0)39+ templateClass<__VA_ARGS__> op(&pipe); \
40- 40+ op.Init(x, grad, argmax, y, workspace, &tilingData); \
41-extern "C" __global__ __aicore__ void max_pool3d_grad_with_argmax(41+ op.ProcessCutNc(); \
42- GM_ADDR x, GM_ADDR grad, GM_ADDR argmax, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)42+ } while (0)
43-{43+#endif
44- if (workspace == nullptr || GetUserWorkspace(workspace) == nullptr || g_coreType == AIC) {44+ 
45- return;45+#if defined(__CCE_AICORE__) && __CCE_AICORE__ == 310
46- }46+using namespace MaxPool3DGradWithArgmaxOp;
47- 47+template <uint64_t INDEX_DTYPE = TPL_INT32, uint64_t IS_SIMT = 0, uint64_t IS_CHANNEL_LAST = 0, uint64_t IS_CHECK_RANGE = 0>
48- TPipe pipe;48+__global__ __aicore__ void max_pool3d_grad_with_argmax(
49- // The percentile determines if overlap occurs49+ GM_ADDR x, GM_ADDR grad, GM_ADDR argmax, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
50- if (TILING_KEY_IS(0)) { // Normal Kernel50+{
51- GENERAL_OP_IMPL(MaxPool3DGradWithArgmaxNormal, DTYPE_X, DTYPE_X, int32_t, DTYPE_X, false);51+ if (workspace == nullptr || GetUserWorkspace(workspace) == nullptr || g_coreType == AIC) {
52- } else if (TILING_KEY_IS(100)) {52+ return;
53- GENERAL_OP_IMPL(MaxPool3DGradWithArgmaxNormal, DTYPE_X, DTYPE_X, int32_t, DTYPE_X, true);53+ }
54- } else if (TILING_KEY_IS(2)) { // Scatter Kernel54+ TPipe pipe;
55- GENERAL_OP_IMPL(MaxPoolGradWithArgScatter, DTYPE_X, DTYPE_X, int32_t, DTYPE_X);55+ if constexpr (INDEX_DTYPE == TPL_INT32 && IS_SIMT == 1) {
56- } else if (TILING_KEY_IS(102)) {56+ REGISTER_TILING_DEFAULT(MaxPool3DGradWithArgmaxTilingDataV35);
57- GENERAL_OP_IMPL(MaxPoolGradWithArgScatterOverlap, DTYPE_X, DTYPE_GRAD, int32_t, DTYPE_Y);57+ GET_TILING_DATA_WITH_STRUCT(MaxPool3DGradWithArgmaxTilingDataV35, tilingData, tiling);
58- } else if (TILING_KEY_IS(1)) { // CutK Kernel, no cut58+ MaxPool3DGradWithArgmaxSimt<DTYPE_X, DTYPE_ARGMAX, int32_t, IS_CHANNEL_LAST> op(&pipe, &tilingData);
59- GENERAL_OP_IMPL(MaxPool3DGradWithArgmaxCutKD, DTYPE_X, DTYPE_GRAD, int32_t, DTYPE_Y, false);59+ op.Init(x, grad, argmax, y);
60- } else if (TILING_KEY_IS(21)) { // CutK Kernel, cut do60+ op.Process();
61- GENERAL_OP_IMPL(MaxPool3DGradWithArgmaxCutKD, DTYPE_X, DTYPE_GRAD, int32_t, DTYPE_Y, false);61+ } else if constexpr (INDEX_DTYPE == TPL_INT64 && IS_SIMT == 1) {
62- } else if (TILING_KEY_IS(31)) { // CutK Kernel, cut do, kd, ho62+ REGISTER_TILING_DEFAULT(MaxPool3DGradWithArgmaxTilingDataV35);
63- GENERAL_OP_IMPL(MaxPool3DGradWithArgmaxCutKD, DTYPE_X, DTYPE_GRAD, int32_t, DTYPE_Y, false);63+ GET_TILING_DATA_WITH_STRUCT(MaxPool3DGradWithArgmaxTilingDataV35, tilingData, tiling);
64- } else if (TILING_KEY_IS(41)) { // CutK Kernel, cut do, kd, ho, kh, wo64+ MaxPool3DGradWithArgmaxSimt<DTYPE_X, DTYPE_ARGMAX, int64_t, IS_CHANNEL_LAST> op(&pipe, &tilingData);
65- GENERAL_OP_IMPL(MaxPool3DGradWithArgmaxCutKDH, DTYPE_X, DTYPE_GRAD, int32_t, DTYPE_Y, false);65+ op.Init(x, grad, argmax, y);
66- } else if (TILING_KEY_IS(51)) { // CutK Kernel, cut do, kd66+ op.Process();
67- GENERAL_OP_IMPL(MaxPool3DGradWithArgmaxCutKD, DTYPE_X, DTYPE_GRAD, int32_t, DTYPE_Y, false);67+ } else if constexpr (INDEX_DTYPE == TPL_INT32 && IS_SIMT == 0 && IS_CHECK_RANGE == 0) {
68- } else if (TILING_KEY_IS(61)) { // CutK Kernel, cut do, kd, ho, kh68+ REGISTER_TILING_FOR_TILINGKEY("INDEX_DTYPE == TPL_INT32 && IS_SIMT == 0 && IS_CHECK_RANGE == 0", MaxPool3DGradWithArgmaxNCDHWTilingData);
69- GENERAL_OP_IMPL(MaxPool3DGradWithArgmaxCutKDH, DTYPE_X, DTYPE_X, int32_t, DTYPE_X, false);69+ GET_TILING_DATA_WITH_STRUCT(MaxPool3DGradWithArgmaxNCDHWTilingData, tilingData, tiling);
70- } else if (TILING_KEY_IS(71)) { // CutK Kernel, cut do, kd, ho, kh, wo, kw70+ MaxPool3DGradWithArgmaxNCDHWNameSpace::MaxPool3DGradWithArgmaxNCDHWKernel<DTYPE_X, DTYPE_ARGMAX, int32_t, false> op;
71- GENERAL_OP_IMPL(MaxPool3DGradWithArgmaxCutKDHW, DTYPE_X, DTYPE_GRAD, int32_t, DTYPE_Y, false);71+ op.Init(x, grad, argmax, y, pipe, tilingData);
72- } else if (TILING_KEY_IS(101)) {72+ op.Process();
73- GENERAL_OP_IMPL_CUTNC(MaxPool3DGradWithArgmaxCutKD, DTYPE_X, DTYPE_GRAD, int32_t, DTYPE_Y, true);73+ } else if constexpr (INDEX_DTYPE == TPL_INT32 && IS_SIMT == 0 && IS_CHECK_RANGE == 1) {
74- } else if (TILING_KEY_IS(121)) {74+ REGISTER_TILING_FOR_TILINGKEY("INDEX_DTYPE == TPL_INT32 && IS_SIMT == 0 && IS_CHECK_RANGE == 1", MaxPool3DGradWithArgmaxNCDHWTilingData);
75- GENERAL_OP_IMPL_CUTNC(MaxPool3DGradWithArgmaxCutKD, DTYPE_X, DTYPE_GRAD, int32_t, DTYPE_Y, true);75+ GET_TILING_DATA_WITH_STRUCT(MaxPool3DGradWithArgmaxNCDHWTilingData, tilingData, tiling);
76- } else if (TILING_KEY_IS(131)) {76+ MaxPool3DGradWithArgmaxNCDHWNameSpace::MaxPool3DGradWithArgmaxNCDHWKernel<DTYPE_X, DTYPE_ARGMAX, int32_t, true> op;
77- GENERAL_OP_IMPL_CUTNC(MaxPool3DGradWithArgmaxCutKD, DTYPE_X, DTYPE_GRAD, int32_t, DTYPE_Y, true);77+ op.Init(x, grad, argmax, y, pipe, tilingData);
78- } else if (TILING_KEY_IS(141)) {78+ op.Process();
79- GENERAL_OP_IMPL_CUTNC(MaxPool3DGradWithArgmaxCutKDH, DTYPE_X, DTYPE_GRAD, int32_t, DTYPE_Y, true);79+ } else if constexpr (INDEX_DTYPE == TPL_INT64 && IS_SIMT == 0 && IS_CHECK_RANGE == 0) {
80- } else if (TILING_KEY_IS(151)) {80+ REGISTER_TILING_FOR_TILINGKEY("INDEX_DTYPE == TPL_INT64 && IS_SIMT == 0 && IS_CHECK_RANGE == 0", MaxPool3DGradWithArgmaxNCDHWTilingData);
81- GENERAL_OP_IMPL_CUTNC(MaxPool3DGradWithArgmaxCutKD, DTYPE_X, DTYPE_X, int32_t, DTYPE_X, true);81+ GET_TILING_DATA_WITH_STRUCT(MaxPool3DGradWithArgmaxNCDHWTilingData, tilingData, tiling);
82- } else if (TILING_KEY_IS(161)) {82+ MaxPool3DGradWithArgmaxNCDHWNameSpace::MaxPool3DGradWithArgmaxNCDHWKernel<DTYPE_X, DTYPE_ARGMAX, int64_t, false> op;
83- GENERAL_OP_IMPL_CUTNC(MaxPool3DGradWithArgmaxCutKDH, DTYPE_X, DTYPE_X, int32_t, DTYPE_X, true);83+ op.Init(x, grad, argmax, y, pipe, tilingData);
84- } else if (TILING_KEY_IS(171)) {84+ op.Process();
85- GENERAL_OP_IMPL_CUTNC(MaxPool3DGradWithArgmaxCutKDHW, DTYPE_X, DTYPE_X, int32_t, DTYPE_X, true);85+ } else if constexpr (INDEX_DTYPE == TPL_INT64 && IS_SIMT == 0 && IS_CHECK_RANGE == 1) {
86- }86+ REGISTER_TILING_FOR_TILINGKEY("INDEX_DTYPE == TPL_INT64 && IS_SIMT == 0 && IS_CHECK_RANGE == 1", MaxPool3DGradWithArgmaxNCDHWTilingData);
87- return;87+ GET_TILING_DATA_WITH_STRUCT(MaxPool3DGradWithArgmaxNCDHWTilingData, tilingData, tiling);
88-}88+ MaxPool3DGradWithArgmaxNCDHWNameSpace::MaxPool3DGradWithArgmaxNCDHWKernel<DTYPE_X, DTYPE_ARGMAX, int64_t, true> op;
89+ op.Init(x, grad, argmax, y, pipe, tilingData);
90+ op.Process();
91+ }
92+}
93+#else
94+extern "C" __global__ __aicore__ void max_pool3d_grad_with_argmax(
95+ GM_ADDR x, GM_ADDR grad, GM_ADDR argmax, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
96+{
97+ if (workspace == nullptr || GetUserWorkspace(workspace) == nullptr || g_coreType == AIC) {
98+ return;
99+ }
100+ 
101+ TPipe pipe;
102+ // The percentile determines if overlap occurs
103+ if (TILING_KEY_IS(0)) { // Normal Kernel
104+ GENERAL_OP_IMPL(MaxPool3DGradWithArgmaxNormal, DTYPE_X, DTYPE_X, int32_t, DTYPE_X, false);
105+ } else if (TILING_KEY_IS(100)) {
106+ GENERAL_OP_IMPL(MaxPool3DGradWithArgmaxNormal, DTYPE_X, DTYPE_X, int32_t, DTYPE_X, true);
107+ } else if (TILING_KEY_IS(2)) { // Scatter Kernel
108+ GENERAL_OP_IMPL(MaxPoolGradWithArgScatter, DTYPE_X, DTYPE_X, int32_t, DTYPE_X);
109+ } else if (TILING_KEY_IS(102)) {
110+ GENERAL_OP_IMPL(MaxPoolGradWithArgScatterOverlap, DTYPE_X, DTYPE_GRAD, int32_t, DTYPE_Y);
111+ } else if (TILING_KEY_IS(1)) { // CutK Kernel, no cut
112+ GENERAL_OP_IMPL(MaxPool3DGradWithArgmaxCutKD, DTYPE_X, DTYPE_GRAD, int32_t, DTYPE_Y, false);
113+ } else if (TILING_KEY_IS(21)) { // CutK Kernel, cut do
114+ GENERAL_OP_IMPL(MaxPool3DGradWithArgmaxCutKD, DTYPE_X, DTYPE_GRAD, int32_t, DTYPE_Y, false);
115+ } else if (TILING_KEY_IS(31)) { // CutK Kernel, cut do, kd, ho
116+ GENERAL_OP_IMPL(MaxPool3DGradWithArgmaxCutKD, DTYPE_X, DTYPE_GRAD, int32_t, DTYPE_Y, false);
117+ } else if (TILING_KEY_IS(41)) { // CutK Kernel, cut do, kd, ho, kh, wo
118+ GENERAL_OP_IMPL(MaxPool3DGradWithArgmaxCutKDH, DTYPE_X, DTYPE_GRAD, int32_t, DTYPE_Y, false);
119+ } else if (TILING_KEY_IS(51)) { // CutK Kernel, cut do, kd
120+ GENERAL_OP_IMPL(MaxPool3DGradWithArgmaxCutKD, DTYPE_X, DTYPE_GRAD, int32_t, DTYPE_Y, false);
121+ } else if (TILING_KEY_IS(61)) { // CutK Kernel, cut do, kd, ho, kh
122+ GENERAL_OP_IMPL(MaxPool3DGradWithArgmaxCutKDH, DTYPE_X, DTYPE_X, int32_t, DTYPE_X, false);
123+ } else if (TILING_KEY_IS(71)) { // CutK Kernel, cut do, kd, ho, kh, wo, kw
124+ GENERAL_OP_IMPL(MaxPool3DGradWithArgmaxCutKDHW, DTYPE_X, DTYPE_GRAD, int32_t, DTYPE_Y, false);
125+ } else if (TILING_KEY_IS(101)) {
126+ GENERAL_OP_IMPL_CUTNC(MaxPool3DGradWithArgmaxCutKD, DTYPE_X, DTYPE_GRAD, int32_t, DTYPE_Y, true);
127+ } else if (TILING_KEY_IS(121)) {
128+ GENERAL_OP_IMPL_CUTNC(MaxPool3DGradWithArgmaxCutKD, DTYPE_X, DTYPE_GRAD, int32_t, DTYPE_Y, true);
129+ } else if (TILING_KEY_IS(131)) {
130+ GENERAL_OP_IMPL_CUTNC(MaxPool3DGradWithArgmaxCutKD, DTYPE_X, DTYPE_GRAD, int32_t, DTYPE_Y, true);
131+ } else if (TILING_KEY_IS(141)) {
132+ GENERAL_OP_IMPL_CUTNC(MaxPool3DGradWithArgmaxCutKDH, DTYPE_X, DTYPE_GRAD, int32_t, DTYPE_Y, true);
133+ } else if (TILING_KEY_IS(151)) {
134+ GENERAL_OP_IMPL_CUTNC(MaxPool3DGradWithArgmaxCutKD, DTYPE_X, DTYPE_X, int32_t, DTYPE_X, true);
135+ } else if (TILING_KEY_IS(161)) {
136+ GENERAL_OP_IMPL_CUTNC(MaxPool3DGradWithArgmaxCutKDH, DTYPE_X, DTYPE_X, int32_t, DTYPE_X, true);
137+ } else if (TILING_KEY_IS(171)) {
138+ GENERAL_OP_IMPL_CUTNC(MaxPool3DGradWithArgmaxCutKDHW, DTYPE_X, DTYPE_X, int32_t, DTYPE_X, true);
139+ }
140+ return;
141+}
142+#endif
Apooling/max_pool3d_grad_with_argmax/tests/ut/op_host/test_max_pool3d_grad_with_argmax_infershape.cpp+95-0
@@ -0,0 +1,95 @@
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+#include <gtest/gtest.h>
12+#include <iostream>
13+#include "../../../op_graph/max_pool3d_grad_with_argmax_proto.h"
14+#include "infershape_test_util.h"
15+#include "ut_op_common.h"
16+#include "error_util.h"
17+#include "log/log.h"
18+ 
19+using namespace ge;
20+using namespace op;
21+ 
22+class MaxPool3DGradWithArgmaxInferShapeTest : public testing::Test
23+{
24+protected:
25+ static void SetUpTestCase()
26+ {
27+ std::cout << "MaxPool3DGradWithArgmax InferShape Test SetUp" << std::endl;
28+ }
29+ 
30+ static void TearDownTestCase()
31+ {
32+ std::cout << "MaxPool3DGradWithArgmax InferShape Test TearDown" << std::endl;
33+ }
34+};
35+ 
36+TEST_F(MaxPool3DGradWithArgmaxInferShapeTest, max_pool3d_grad_with_argmax_infershape_test1)
37+{
38+ auto inferShapeFunc = gert::OpImplRegistry::GetInstance().GetOpImpl("MaxPool3DGradWithArgmax")->infer_shape;
39+ 
40+ gert::StorageShape xShape = {{3698, 2, 2, 3, 2}, {3698, 2, 2, 3, 2}};
41+ gert::StorageShape gradShape = {{3698, 2, 2, 1, 2}, {3698, 2, 2, 1, 2}};
42+ gert::StorageShape yShape = {{3698, 2, 2, 3, 2}, {3698, 2, 2, 3, 2}};
43+ gert::StorageShape indicesShape = {{3698, 2, 2, 1, 2}, {3698, 2, 2, 1, 2}};
44+ auto holder = gert::InferShapeContextFaker()
45+ .NodeIoNum(3, 1)
46+ .IrInstanceNum({1, 1, 1})
47+ .NodeInputTd(0, ge::DT_FLOAT, ge::Format::FORMAT_NCDHW, ge::Format::FORMAT_RESERVED)
48+ .NodeInputTd(1, ge::DT_FLOAT, ge::Format::FORMAT_NCDHW, ge::Format::FORMAT_RESERVED)
49+ .NodeInputTd(2, ge::DT_INT32, ge::Format::FORMAT_NCDHW, ge::Format::FORMAT_RESERVED)
50+ .NodeOutputTd(0, ge::DT_FLOAT, ge::Format::FORMAT_NCDHW, ge::Format::FORMAT_RESERVED)
51+ .NodeAttrs({{"ksize", Ops::NN::AnyValue::CreateFrom<std::vector<int64_t>>({1, 7, 1})},
52+ {"strides", Ops::NN::AnyValue::CreateFrom<std::vector<int64_t>>({})},
53+ {"pads", Ops::NN::AnyValue::CreateFrom<std::vector<int64_t>>({0, 1, 0})},
54+ {"dilation", Ops::NN::AnyValue::CreateFrom<std::vector<int64_t>>({2474, 1, 1})},
55+ {"ceil_mode", Ops::NN::AnyValue::CreateFrom<bool>(true)},
56+ {"data_format", Ops::NN::AnyValue::CreateFrom<std::string>("NCDHW")}})
57+ .InputShapes({&xShape, &gradShape, &indicesShape})
58+ .OutputShapes({&yShape})
59+ .Build();
60+ 
61+ ASSERT_EQ(inferShapeFunc(holder.GetContext<gert::InferShapeContext>()), ge::GRAPH_SUCCESS);
62+ gert::Shape* output = holder.GetContext<gert::InferShapeContext>()->GetOutputShape(0);
63+ ASSERT_EQ(Shape2String(*output), "[3698, 2, 2, 3, 2]");
64+}
65+ 
66+TEST_F(MaxPool3DGradWithArgmaxInferShapeTest, max_pool3d_grad_with_argmax_infershape_test2)
67+{
68+ auto inferShapeFunc = gert::OpImplRegistry::GetInstance().GetOpImpl("MaxPool3DGradWithArgmax")->infer_shape;
69+ 
70+ gert::StorageShape xShape = {{8, 5, 6, 8, 8}, {8, 5, 6, 8, 8}};
71+ gert::StorageShape gradShape = {{8, 1, 1, 2, 8}, {8, 1, 1, 2, 8}};
72+ gert::StorageShape yShape = {{8, 5, 6, 8, 8}, {8, 5, 6, 8, 8}};
73+ gert::StorageShape indicesShape = {{8, 1, 1, 2, 8}, {8, 1, 1, 2, 8}};
74+ 
75+ auto holder = gert::InferShapeContextFaker()
76+ .NodeIoNum(3, 1)
77+ .IrInstanceNum({1, 1, 1})
78+ .NodeInputTd(0, ge::DT_FLOAT, ge::Format::FORMAT_NDHWC, ge::Format::FORMAT_RESERVED)
79+ .NodeInputTd(1, ge::DT_FLOAT, ge::Format::FORMAT_NDHWC, ge::Format::FORMAT_RESERVED)
80+ .NodeInputTd(2, ge::DT_INT64, ge::Format::FORMAT_NDHWC, ge::Format::FORMAT_RESERVED)
81+ .NodeOutputTd(0, ge::DT_FLOAT, ge::Format::FORMAT_NDHWC, ge::Format::FORMAT_RESERVED)
82+ .NodeAttrs({{"ksize", Ops::NN::AnyValue::CreateFrom<std::vector<int64_t>>({6, 7, 7})},
83+ {"strides", Ops::NN::AnyValue::CreateFrom<std::vector<int64_t>>({7404, 6, 6})},
84+ {"pads", Ops::NN::AnyValue::CreateFrom<std::vector<int64_t>>({0})},
85+ {"dilation", Ops::NN::AnyValue::CreateFrom<std::vector<int64_t>>({1})},
86+ {"ceil_mode", Ops::NN::AnyValue::CreateFrom<bool>(true)},
87+ {"data_format", Ops::NN::AnyValue::CreateFrom<std::string>("NDHWC")}})
88+ .InputShapes({&xShape, &gradShape, &indicesShape})
89+ .OutputShapes({&yShape})
90+ .Build();
91+ 
92+ ASSERT_EQ(inferShapeFunc(holder.GetContext<gert::InferShapeContext>()), ge::GRAPH_SUCCESS);
93+ gert::Shape* output = holder.GetContext<gert::InferShapeContext>()->GetOutputShape(0);
94+ ASSERT_EQ(Shape2String(*output), "[8, 5, 6, 8, 8]");
95+}
Mpooling/max_pool3d_grad_with_argmax/tests/ut/op_host/test_max_pool3d_grad_with_argmax_tiling.cpp+219-0
@@ -552,3 +552,222 @@ TEST_F(MaxPool3dGradWithArgmaxTiling, max_pool3d_grad_with_argmax_tilingkey_2_ca
552 {"ceil_mode", Ops::NN::AnyValue::CreateFrom<bool>(false)}};552 {"ceil_mode", Ops::NN::AnyValue::CreateFrom<bool>(false)}};
553 TestMaxPool3dGradWithArgmaxTiling(xShape, gradShape, argmaxShape, dxShape, attrList, ge::DT_FLOAT, 2);553 TestMaxPool3dGradWithArgmaxTiling(xShape, gradShape, argmaxShape, dxShape, attrList, ge::DT_FLOAT, 2);
554}554}
555+ 
556+static void ExecuteTestCase(gert::StorageShape xShape, gert::StorageShape gradShape,gert::StorageShape argmaxShape,
557+ gert::StorageShape yShape, std::vector<int64_t> ksize, std::vector<int64_t> strides,
558+ std::vector<int64_t> pads, std::vector<int64_t> dilation, ge::DataType dtype,
559+ int64_t index_dtype, bool ceil_mode, std::string data_format, uint64_t except_tilingkey,
560+ std::string expect)
561+{
562+ dlog_setlevel(0, 0, 0);
563+ 
564+ string compile_info_string = R"({
565+ "hardware_info": {"BT_SIZE": 0, "load3d_constraints": "1",
566+ "Intrinsic_fix_pipe_l0c2out": false,
567+ "Intrinsic_data_move_l12ub": true,
568+ "Intrinsic_data_move_l0c2ub": true,
569+ "Intrinsic_data_move_out2l1_nd2nz": false,
570+ "UB_SIZE": 245760, "L2_SIZE": 33554432, "L1_SIZE": 524288,
571+ "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072,
572+ "CORE_NUM": 64}
573+ })";
574+ map<string, string> soc_infos;
575+ map<string, string> aicore_spec;
576+ map<string, string> intrinsics;
577+ GetPlatFormInfos(compile_info_string.c_str(), soc_infos, aicore_spec, intrinsics);
578+ std::map<std::string, std::string> soc_version_infos = {{"Short_SoC_version", "Ascend950"}};
579+ 
580+ // platform info
581+ fe::PlatFormInfos platform_info;
582+ platform_info.Init();
583+ // compile info
584+ optiling::Tiling4MaxPool3DGradWithArgmaxCompileInfo compile_info;
585+ 
586+ std::string op_type("MaxPool3DGradWithArgmax");
587+ ASSERT_NE(gert::OpImplRegistry::GetInstance().GetOpImpl(op_type.c_str()), nullptr);
588+ auto tiling_func = gert::OpImplRegistry::GetInstance().GetOpImpl(op_type.c_str())->tiling;
589+ auto tiling_parse_func = gert::OpImplRegistry::GetInstance().GetOpImpl(op_type.c_str())->tiling_parse;
590+ 
591+ // tilingParseFunc simulate
592+ auto kernel_holder =
593+ gert::KernelRunContextFaker()
594+ .KernelIONum(2, 1)
595+ .Inputs({const_cast<char*>(compile_info_string.c_str()), reinterpret_cast<void*>(&platform_info)})
596+ .Outputs({&compile_info})
597+ .Build();
598+ 
599+ ASSERT_TRUE(kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->Init());
600+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", soc_infos);
601+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicore_spec);
602+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
603+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap",
604+ intrinsics);
605+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("version",
606+ soc_version_infos);
607+ 
608+ ASSERT_EQ(tiling_parse_func(kernel_holder.GetContext<gert::KernelContext>()), ge::GRAPH_SUCCESS);
609+ 
610+ // tilingFunc simulate
611+ auto param = gert::TilingData::CreateCap(4096);
612+ auto workspace_size_holer = gert::ContinuousVector::Create<size_t>(4096);
613+ auto ws_size = reinterpret_cast<gert::ContinuousVector*>(workspace_size_holer.get());
614+ ge::DataType indicesDtype = (index_dtype == 3) ? ge::DT_INT32 : ge::DT_INT64;
615+ ASSERT_NE(param, nullptr);
616+ auto holder = gert::TilingContextFaker()
617+ .SetOpType(op_type)
618+ .NodeIoNum(3, 1)
619+ .IrInstanceNum({1})
620+ .InputShapes({&xShape, &gradShape, &argmaxShape})
621+ .OutputShapes({&yShape})
622+ .CompileInfo(&compile_info)
623+ .PlatformInfo(reinterpret_cast<char*>(&platform_info))
624+ .NodeInputTd(0, dtype, ge::FORMAT_ND, ge::FORMAT_ND)
625+ .NodeInputTd(1, dtype, ge::FORMAT_ND, ge::FORMAT_ND)
626+ .NodeInputTd(2, indicesDtype, ge::FORMAT_ND, ge::FORMAT_ND)
627+ .NodeOutputTd(0, dtype, ge::FORMAT_ND, ge::FORMAT_ND)
628+ .NodeAttrs({{"ksize", Ops::NN::AnyValue::CreateFrom<std::vector<int64_t>>(ksize)},
629+ {"strides", Ops::NN::AnyValue::CreateFrom<std::vector<int64_t>>(strides)},
630+ {"pads", Ops::NN::AnyValue::CreateFrom<std::vector<int64_t>>(pads)},
631+ {"dilation", Ops::NN::AnyValue::CreateFrom<std::vector<int64_t>>(dilation)},
632+ {"ceil_mode", Ops::NN::AnyValue::CreateFrom<bool>(ceil_mode)},
633+ {"data_format", Ops::NN::AnyValue::CreateFrom<std::string>(data_format)},
634+ {"dtype", Ops::NN::AnyValue::CreateFrom<int64_t>(index_dtype)}})
635+ .TilingData(param.get())
636+ .Workspace(ws_size)
637+ .Build();
638+ 
639+ gert::TilingContext* tiling_context = holder.GetContext<gert::TilingContext>();
640+ ASSERT_NE(tiling_context->GetPlatformInfo(), nullptr);
641+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", soc_infos);
642+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicore_spec);
643+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
644+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics);
645+
646+ // workspaces nullptr return failed
647+ EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_SUCCESS);
648+ auto tiling_key = tiling_context->GetTilingKey();
649+ ASSERT_EQ(tiling_key, except_tilingkey);
650+ auto tilingData = tiling_context->GetRawTilingData();
651+ ASSERT_NE(tilingData, nullptr);
652+ dlog_setlevel(0, 3, 0);
653+}
654+ 
655+TEST_F(MaxPool3dGradWithArgmaxTiling, MaxPool3DGradWithArgmax_tiling_simt_test_0001_NCDHW)
656+{
657+ gert::StorageShape xShape = {{2,4,5,6,3}, {2,4,5,6,3}};
658+ gert::StorageShape gradShape = {{2,2,2,3,3}, {2,2,2,3,3}};
659+ gert::StorageShape argmaxShape = {{2,2,2,3,3}, {2,2,2,3,3}};
660+ gert::StorageShape yShape = {{2,4,5,6,3}, {2,4,5,6,3}};
661+ std::vector<int64_t> ksize = {2,2,2};
662+ std::vector<int64_t> strides = {2,2,2};
663+ std::vector<int64_t> pads = {0,0,0};
664+ std::vector<int64_t> dilation = {1,1,1};
665+ ge::DataType dtype = ge::DT_FLOAT;
666+ int64_t index_dtype = 3;
667+ bool ceil_mode = false;
668+ std::string data_format = "NDHWC";
669+ uint64_t except_tilingkey = 769;
670+ std::string expect = " ";
671+ ExecuteTestCase(xShape, gradShape, argmaxShape, yShape, ksize, strides, pads, dilation, dtype, index_dtype, ceil_mode,
672+ data_format, except_tilingkey, expect);
673+}
674+ 
675+TEST_F(MaxPool3dGradWithArgmaxTiling, MaxPool3DGradWithArgmax_tiling_simd_test_0001_NCDHW)
676+{
677+ gert::StorageShape xShape = {{1,2,8,10,514}, {1,2,8,10,514}};
678+ gert::StorageShape gradShape = {{1,2,4,5,257}, {1,2,4,5,257}};
679+ gert::StorageShape argmaxShape = {{1,2,4,5,257}, {1,2,4,5,257}};
680+ gert::StorageShape yShape = {{1,2,8,10,514}, {1,2,8,10,514}};
681+ std::vector<int64_t> ksize = {2,2,2};
682+ std::vector<int64_t> strides = {2,2,2};
683+ std::vector<int64_t> pads = {0,0,0};
684+ std::vector<int64_t> dilation = {1,1,1};
685+ ge::DataType dtype = ge::DT_FLOAT;
686+ int64_t index_dtype = 3;
687+ bool ceil_mode = false;
688+ std::string data_format = "NCDHW";
689+ uint64_t except_tilingkey = 1;
690+ std::string expect = " ";
691+ ExecuteTestCase(xShape, gradShape, argmaxShape, yShape, ksize, strides, pads, dilation, dtype, index_dtype, ceil_mode,
692+ data_format, except_tilingkey, expect);
693+}
694+ 
695+TEST_F(MaxPool3dGradWithArgmaxTiling, MaxPool3DGradWithArgmax_tiling_simd_test_0002_NCDHW)
696+{
697+ gert::StorageShape xShape = {{2,2,10,11,1024}, {2,2,10,11,1024}};
698+ gert::StorageShape gradShape = {{2,2,11,6,342}, {2,2,11,6,342}};
699+ gert::StorageShape argmaxShape = {{2,2,11,6,342}, {2,2,11,6,342}};
700+ gert::StorageShape yShape = {{2,2,10,11,1024}, {2,2,10,11,1024}};
701+ std::vector<int64_t> ksize = {2,4,5};
702+ std::vector<int64_t> strides = {1,2,3};
703+ std::vector<int64_t> pads = {1,2,2};
704+ std::vector<int64_t> dilation = {1,1,1};
705+ ge::DataType dtype = ge::DT_FLOAT;
706+ int64_t index_dtype = 3;
707+ bool ceil_mode = false;
708+ std::string data_format = "NCDHW";
709+ uint64_t except_tilingkey = 1025;
710+ std::string expect = " ";
711+ ExecuteTestCase(xShape, gradShape, argmaxShape, yShape, ksize, strides, pads, dilation, dtype, index_dtype, ceil_mode,
712+ data_format, except_tilingkey, expect);
713+}
714+ 
715+TEST_F(MaxPool3dGradWithArgmaxTiling, MaxPool3DGradWithArgmax_tiling_simd_test_0003_NCDHW)
716+{
717+ gert::StorageShape xShape = {{38,235,4,4,4}, {38,235,4,4,4}};
718+ gert::StorageShape gradShape = {{38,235,3,3,3}, {38,235,3,3,3}};
719+ gert::StorageShape argmaxShape = {{38,235,3,3,3}, {38,235,3,3,3}};
720+ gert::StorageShape yShape = {{38,235,4,4,4}, {38,235,4,4,4}};
721+ std::vector<int64_t> ksize = {2,2,2};
722+ std::vector<int64_t> strides = {2,2,2};
723+ std::vector<int64_t> pads = {1,1,1};
724+ std::vector<int64_t> dilation = {1,1,1};
725+ ge::DataType dtype = ge::DT_FLOAT;
726+ int64_t index_dtype = 3;
727+ bool ceil_mode = false;
728+ std::string data_format = "NCDHW";
729+ uint64_t except_tilingkey = 1;
730+ std::string expect = " ";
731+ ExecuteTestCase(xShape, gradShape, argmaxShape, yShape, ksize, strides, pads, dilation, dtype, index_dtype, ceil_mode,
732+ data_format, except_tilingkey, expect);
733+}
734+ 
735+TEST_F(MaxPool3dGradWithArgmaxTiling, MaxPool3DGradWithArgmax_tiling_simd_test_0004_NCDHW)
736+{
737+ gert::StorageShape xShape = {{2,2,9,10,960}, {2,2,9,10,960}};
738+ gert::StorageShape gradShape = {{2,2,9,5,240}, {2,2,9,5,240}};
739+ gert::StorageShape argmaxShape = {{2,2,9,5,240}, {2,2,9,5,240}};
740+ gert::StorageShape yShape = {{2,2,9,10,960}, {2,2,9,10,960}};
741+ std::vector<int64_t> ksize = {3,2,7};
742+ std::vector<int64_t> strides = {1,2,4};
743+ std::vector<int64_t> pads = {1,0,2};
744+ std::vector<int64_t> dilation = {1,1,1};
745+ ge::DataType dtype = ge::DT_FLOAT;
746+ int64_t index_dtype = 3;
747+ bool ceil_mode = false;
748+ std::string data_format = "NCDHW";
749+ uint64_t except_tilingkey = 1025;
750+ std::string expect = " ";
751+ ExecuteTestCase(xShape, gradShape, argmaxShape, yShape, ksize, strides, pads, dilation, dtype, index_dtype, ceil_mode,
752+ data_format, except_tilingkey, expect);
753+}
754+ 
755+TEST_F(MaxPool3dGradWithArgmaxTiling, MaxPool3DGradWithArgmax_tiling_simd_test_0005_NCDHW)
756+{
757+ gert::StorageShape xShape = {{1,8,12,14,1152}, {1,8,12,14,1152}};
758+ gert::StorageShape gradShape = {{1,8,11,7,287}, {1,8,11,7,287}};
759+ gert::StorageShape argmaxShape = {{1,8,11,7,287}, {1,8,11,7,287}};
760+ gert::StorageShape yShape = {{1,8,12,14,1152}, {1,8,12,14,1152}};
761+ std::vector<int64_t> ksize = {4,3,7};
762+ std::vector<int64_t> strides = {1,2,4};
763+ std::vector<int64_t> pads = {1,1,1};
764+ std::vector<int64_t> dilation = {1,1,1};
765+ ge::DataType dtype = ge::DT_FLOAT;
766+ int64_t index_dtype = 3;
767+ bool ceil_mode = false;
768+ std::string data_format = "NCDHW";
769+ uint64_t except_tilingkey = 1025;
770+ std::string expect = " ";
771+ ExecuteTestCase(xShape, gradShape, argmaxShape, yShape, ksize, strides, pads, dilation, dtype, index_dtype, ceil_mode,
772+ data_format, except_tilingkey, expect);
773+}
Mscripts/kernel/binary_config/ascendc_config.json+1-1
@@ -389,7 +389,7 @@
389 {"name": "LayerNormGradV3", "compute_units": ["ascend910b", "ascend910_93", "ascend950"], "auto_sync": false, "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},389 {"name": "LayerNormGradV3", "compute_units": ["ascend910b", "ascend910_93", "ascend950"], "auto_sync": false, "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},
390 {"name": "LayerNormGrad", "compute_units": ["ascend950"], "auto_sync": false, "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},390 {"name": "LayerNormGrad", "compute_units": ["ascend950"], "auto_sync": false, "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},
391 {"name":"MaxPool3DWithArgmaxV2", "compute_units": ["ascend910b", "ascend910_93", "ascend950", "kirinx90"], "auto_sync" : false, "impl_mode" : ""},391 {"name":"MaxPool3DWithArgmaxV2", "compute_units": ["ascend910b", "ascend910_93", "ascend950", "kirinx90"], "auto_sync" : false, "impl_mode" : ""},
392- {"name":"MaxPool3DGradWithArgmax", "compute_units": ["ascend910b", "ascend910_93"], "auto_sync" : false},392+ {"name":"MaxPool3DGradWithArgmax", "compute_units": ["ascend910b", "ascend910_93", "ascend950"], "auto_sync" : false},
393 {"name":"AdaptiveMaxPool3DGrad", "compute_units": ["ascend910b", "ascend910_93", "ascend950"], "auto_sync" : false},393 {"name":"AdaptiveMaxPool3DGrad", "compute_units": ["ascend910b", "ascend910_93", "ascend950"], "auto_sync" : false},
394 {"name":"LinSpace", "compute_units": ["ascend910_93", "ascend910b", "ascend910", "ascend310p", "ascend950"], "auto_sync" : true},394 {"name":"LinSpace", "compute_units": ["ascend910_93", "ascend910b", "ascend910", "ascend310p", "ascend950"], "auto_sync" : true},
395 {"name":"ApplyCamePart3", "compute_units": ["ascend910b", "ascend910_93"], "auto_sync" : true},395 {"name":"ApplyCamePart3", "compute_units": ["ascend910b", "ascend910_93"], "auto_sync" : true},