Star1008
1936
代码
代码
Issues194
Pull Requests330
Actions
项目讨论
Wiki
分析
Star1008
1936
已合并
支持下一代实现MseLossGrad、max_pool_v3、max_pool_3d算子 #513
tianqiguang创建于 2025年12月26日
支持下一代实现MseLossGrad、max_pool_v3、max_pool_3d算子 #513
已合并
tianqiguang创建于 2025年12月26日
从tianqiguang/ops-nnmaster合入到cann/ops-nnmaster
共 107 个文件变更+38289-32
Mcommon/inc/tiling_base/tiling_util.h+9-1
@@ -20,11 +20,19 @@
20namespace Ops {20namespace Ops {
21namespace NN {21namespace NN {
22namespace OpTiling {22namespace OpTiling {
23+static const gert::Shape g_vec_1_shape = {1};
24+ 
23bool IsRegbaseSocVersion(const gert::TilingParseContext* context);25bool IsRegbaseSocVersion(const gert::TilingParseContext* context);
24 26 
25bool IsRegbaseSocVersion(const gert::TilingContext* context);27bool IsRegbaseSocVersion(const gert::TilingContext* context);
26 28 
27-const gert::Shape& EnsureNotScalar(const gert::Shape& inShape);29+inline const gert::Shape& EnsureNotScalar(const gert::Shape& inShape)
30+{
31+ if (inShape.IsScalar()) {
32+ return g_vec_1_shape;
33+ }
34+ return inShape;
35+}
28} // namespace OpTiling36} // namespace OpTiling
29} // namespace NN37} // namespace NN
30} // namespace Ops38} // namespace Ops
Aloss/mse_loss_grad/CMakeLists.txt+15-0
@@ -0,0 +1,15 @@
1+# This program is free software, you can redistribute it and/or modify.
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This file is a part of the CANN Open Software.
4+# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
7+# See LICENSE in the root of the software repository for the full text of the License.
8+#/
9+ 
10+# 设置算子定义时支持的芯片类型
11+set(SUPPORT_COMPUTE_UNIT "ascend910_95")
12+# 设置每种芯片类型对应的tiling文件目录，即采用op_host目录下哪个文件夹下的tiling文件编译
13+set(SUPPORT_TILING_DIR "arch35")
14+add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE mse_loss_grad ACLNNTYPE aclnn_exclude
15+ COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE DEPENDENCIES mse_loss_grad_v2)
Aloss/mse_loss_grad/README.md+90-0
@@ -0,0 +1,90 @@
1+# MseLossGrad
2+ 
3+## 产品支持情况
4+ 
5+|产品 | 是否支持 |
6+|:-------------------------|:----------:|
7+|<term>Ascend 950PR/Ascend 950DT</term> | √ |
8+ 
9+## 功能说明
10+ 
11+- 算子功能：均方误差函数[aclnnMseLoss](aclnnMseLoss.md)的反向传播。
12+ 
13+- 计算公式：
14+ 
15+ 当`reduction`为`mean`时：
16+ 
17+ $$
18+ MselossBackward(grad, x, y) = grad * (x - y) * 2 / x.numel()
19+ $$
20+ 
21+ 其中`x.numel()`表示`x`中的元素个数。如果`reduction`不是`mean`, 那么：
22+ 
23+ $$
24+ MselossBackward(grad, x, y) = grad * (x - y) * 2
25+ $$
26+ 
27+## 参数说明
28+ 
29+<table style="undefined;table-layout: fixed; width: 1576px"><colgroup>
30+ <col style="width: 170px">
31+ <col style="width: 170px">
32+ <col style="width: 310px">
33+ <col style="width: 212px">
34+ <col style="width: 100px">
35+ </colgroup>
36+ <thead>
37+ <tr>
38+ <th>参数名</th>
39+ <th>输入/输出/属性</th>
40+ <th>描述</th>
41+ <th>数据类型</th>
42+ <th>数据格式</th>
43+ </tr></thead>
44+ <tbody>
45+ <tr>
46+ <td>predict</td>
47+ <td>输入</td>
48+ <td>公式中的输入grad
49+ <td>BFLOAT16、FLOAT16、FLOAT</td>
50+ <td>ND</td>
51+ </tr>
52+ <tr>
53+ <td>label</td>
54+ <td>输入</td>
55+ <td>公式中的输入x。</td>
56+ <td>BFLOAT16、FLOAT16、FLOAT</td>
57+ <td>ND</td>
58+ </tr>
59+ <tr>
60+ <td>dout</td>
61+ <td>输入</td>
62+ <td>公式中的输入y。</td>
63+ <td>BFLOAT16、FLOAT16、FLOAT</td>
64+ <td>ND</td>
65+ </tr>
66+ <tr>
67+ <td>reduction</td>
68+ <td>输入</td>
69+ <td>公式中的输入reduction，指定损失函数的计算方式，支持 0('none') | 1('mean') | 2('sum')。'none' 表示不应用减少，'mean' 表示输出的总和将除以self中的元素数，'sum' 表示输出将被求和。</td>
70+ <td>INT64</td>
71+ <td>ND</td>
72+ </tr>
73+ <tr>
74+ <td>y</td>
75+ <td>输出</td>
76+ <td>公式中的输出MselossBackward。</td>
77+ <td>BFLOAT16、FLOAT16、FLOAT</td>
78+ <td>ND</td>
79+ </tr>
80+ </tbody></table>
81+ 
82+## 约束说明
83+ 
84+无
85+ 
86+## 调用说明
87+ 
88+| 调用方式 | 调用样例 | 说明 |
89+|--------------|------------------------------------------------------------------------|----------------------------------------------------------------|
90+| aclnn调用 | [test_aclnn_mse_loss_grad](./examples/test_aclnn_mse_loss_grad.cpp) | 通过[aclnnMseLossGrad](./docs/aclnnMseLossBackward.md)接口方式调用mse_loss_grad算子。 |
Aloss/mse_loss_grad/docs/aclnnMseLossBackward.md+247-0
@@ -0,0 +1,247 @@
1+# aclnnMseLossBackward
2+ 
3+## 产品支持情况
4+ 
5+| 产品 | 是否支持 |
6+| :----------------------------------------------------------- | :------: |
7+| <term>Ascend 950PR/Ascend 950DT</term> | √ |
8+ 
9+## 功能说明
10+ 
11+- 算子功能：均方误差函数[aclnnMseLoss](aclnnMseLoss.md)的反向传播。
12+ 
13+- 计算公式：
14+ 
15+ 当`reduction`为`mean`时：
16+ 
17+ $$
18+ MselossBackward(grad, x, y) = grad * (x - y) * 2 / x.numel()
19+ $$
20+ 
21+ 其中`x.numel()`表示`x`中的元素个数。如果`reduction`不是`mean`, 那么：
22+ 
23+ $$
24+ MselossBackward(grad, x, y) = grad * (x - y) * 2
25+ $$
26+ 
27+## 函数原型
28+ 
29+每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnMseLossBackwardGetWorkspaceSize”接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用“aclnnMseLossBackward”接口执行计算。
30+ 
31+ - `aclnnStatus aclnnMseLossBackwardGetWorkspaceSize(const aclTensor* gradOutput, const aclTensor* self, const aclTensor* target, int64_t reduction, aclTensor* out, uint64_t* workspaceSize, aclOpExecutor** executor)`
32+ - `aclnnStatus aclnnMseLossBackward(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream)`
33+ 
34+## aclnnMseLossBackwardGetWorkspaceSize
35+ 
36+- **参数说明：**
37+ 
38+ - gradOutput(aclTensor*, 计算输入)：公式中的输入`grad`，Device侧的aclTensor，gradOutput与self、target的数据类型一致，gradOutput与self、target的shape满足[broadcast关系](../../../docs/zh/context/broadcast关系.md)。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持ND，shape支持0到8维。
39+ - <term>Ascend 950PR/Ascend 950DT</term>：数据类型支持BFLOAT16、FLOAT16、FLOAT。
40+ - self(aclTensor*, 计算输入)：公式中的输入`x`，Device侧的aclTensor，gradOutput与self、target的数据类型一致，gradOutput与self、target的shape满足[broadcast关系](../../../docs/zh/context/broadcast关系.md)。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持ND，shape支持0到8维。
41+ - <term>Ascend 950PR/Ascend 950DT</term>：数据类型支持BFLOAT16、FLOAT16、FLOAT。
42+ - target(aclTensor*, 计算输入)：公式中的输入`y`，Device侧的aclTensor，gradOutput与self、target的数据类型一致，gradOutput与self、target的shape满足[broadcast关系](../../../docs/zh/context/broadcast关系.md)。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持ND，shape支持0到8维。
43+ - <term>Ascend 950PR/Ascend 950DT</term>：数据类型支持BFLOAT16、FLOAT16、FLOAT。
44+ 
45+ - reduction(int64_t, 计算输入)：公式中的参数`reduction`，指定损失函数的计算方式，支持 0('none') | 1('mean') | 2('sum')。
46+ 
47+ 'none' 表示不应用减少，'mean' 表示输出的总和将除以self中的元素数，'sum' 表示输出将被求和。
48+ 
49+ - out(aclTensor*, 计算输出)：公式中的输出`MselossBackward(grad, x, y)`，Device侧的aclTensor，out与gradOutput、self、target broadcast之后的tensor的shape一致。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)，[数据格式](../../../docs/zh/context/数据格式.md)支持ND。
50+ - <term>Ascend 950PR/Ascend 950DT</term>：数据类型支持BFLOAT16、FLOAT16、FLOAT。
51+ 
52+ - workspaceSize(uint64_t*, 出参)：返回需要在Device侧申请的workspace大小。
53+ 
54+ - executor(aclOpExecutor**, 出参)：返回op执行器，包含了算子计算流程。
55+ 
56+ 
57+- **返回值：**
58+ 
59+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
60+ 
61+ ```
62+ 第一段接口完成入参校验，出现如下场景时报错：
63+ 返回161001（ACLNN_ERR_PARAM_NULLPTR）：1. 传入的gradOutput、self、target或out是空指针时。
64+ 返回161002（ACLNN_ERR_PARAM_INVALID）：1. self的数据类型不在支持的范围之内。
65+ 2. gradOutput、target的数据类型和self不同。
66+ 3. gradOutput、self和target的shape无法做broadcast。
67+ 4. gradOutput、self和target做broadcast后的shape与out的shape不一致。
68+ 5. reduction值不在0~2范围之内。
69+ 6. gradOutput、self或target的shape超过8维。
70+ ```
71+ 
72+## aclnnMseLossBackward
73+ 
74+- **参数说明：**
75+ 
76+ - workspace(void*, 入参)：在Device侧申请的workspace内存地址。
77+ 
78+ - workspaceSize(uint64_t, 入参)：在Device侧申请的workspace大小，由第一段接口aclnnMseLossBackwardGetWorkspaceSize获取。
79+ 
80+ - executor(aclOpExecutor*, 入参)：op执行器，包含了算子计算流程。
81+ 
82+ - stream(aclrtStream, 入参)：指定执行任务的Stream。
83+ 
84+ 
85+- **返回值：**
86+ 
87+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
88+ 
89+## 约束说明
90+ 
91+无。
92+ 
93+## 调用示例
94+ 
95+示例代码如下，仅供参考，具体编译和执行过程请参考[编译与运行样例](../../../docs/zh/context/编译与运行样例.md)。
96+```Cpp
97+#include <iostream>
98+#include <vector>
99+#include "acl/acl.h"
100+#include "aclnnop/aclnn_mse_loss_backward.h"
101+ 
102+#define CHECK_RET(cond, return_expr) \
103+ do { \
104+ if (!(cond)) { \
105+ return_expr; \
106+ } \
107+ } while (0)
108+ 
109+#define LOG_PRINT(message, ...) \
110+ do { \
111+ printf(message, ##__VA_ARGS__); \
112+ } while (0)
113+ 
114+int64_t GetShapeSize(const std::vector<int64_t>& shape) {
115+ int64_t shapeSize = 1;
116+ for (auto i : shape) {
117+ shapeSize *= i;
118+ }
119+ return shapeSize;
120+}
121+ 
122+int Init(int32_t deviceId, aclrtStream* stream) {
123+ // 固定写法，资源初始化
124+ auto ret = aclInit(nullptr);
125+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
126+ ret = aclrtSetDevice(deviceId);
127+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
128+ ret = aclrtCreateStream(stream);
129+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
130+ return 0;
131+}
132+ 
133+template <typename T>
134+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
135+ aclDataType dataType, aclTensor** tensor) {
136+ auto size = GetShapeSize(shape) * sizeof(T);
137+ // 调用aclrtMalloc申请device侧内存
138+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
139+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
140+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
141+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
142+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
143+ 
144+ // 计算连续tensor的strides
145+ std::vector<int64_t> strides(shape.size(), 1);
146+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
147+ strides[i] = shape[i + 1] * strides[i + 1];
148+ }
149+ 
150+ // 调用aclCreateTensor接口创建aclTensor
151+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
152+ shape.data(), shape.size(), *deviceAddr);
153+ return 0;
154+}
155+ 
156+int main() {
157+ // 1. （固定写法）device/stream初始化，参考acl API手册
158+ // 根据自己的实际device填写deviceId
159+ int32_t deviceId = 0;
160+ aclrtStream stream;
161+ auto ret = Init(deviceId, &stream);
162+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
163+ 
164+ // 2. 构造输入与输出，需要根据API的接口自定义构造gradOutput
165+ std::vector<int64_t> gradOutputShape = {2, 2};
166+ std::vector<int64_t> selfShape = {2, 2};
167+ std::vector<int64_t> targetShape = {2, 2};
168+ std::vector<int64_t> outShape = {2, 2};
169+ void* gradOutputDeviceAddr = nullptr;
170+ void* selfDeviceAddr = nullptr;
171+ void* targetDeviceAddr = nullptr;
172+ void* outDeviceAddr = nullptr;
173+ aclTensor* gradOutput = nullptr;
174+ aclTensor* self = nullptr;
175+ aclTensor* target = nullptr;
176+ aclTensor* out = nullptr;
177+ std::vector<float> gradOutputHostData = {0, 1, 2, 3};
178+ std::vector<float> selfHostData = {0, 1, 2, 3};
179+ std::vector<float> targetHostData = {1, 1, 1, 1};
180+ std::vector<float> outHostData(4, 0);
181+ // 创建gradOutput aclTensor
182+ ret = CreateAclTensor(gradOutputHostData, gradOutputShape, &gradOutputDeviceAddr,
183+ aclDataType::ACL_FLOAT, &gradOutput);
184+ CHECK_RET(ret == ACL_SUCCESS, return ret);
185+ // 创建self aclTensor
186+ ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self);
187+ CHECK_RET(ret == ACL_SUCCESS, return ret);
188+ // 创建target aclTensor
189+ ret = CreateAclTensor(targetHostData, targetShape, &targetDeviceAddr, aclDataType::ACL_FLOAT, &target);
190+ CHECK_RET(ret == ACL_SUCCESS, return ret);
191+ // 创建out aclTensor
192+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
193+ CHECK_RET(ret == ACL_SUCCESS, return ret);
194+ // 创建reduction
195+ int64_t reduction = 1;
196+ 
197+ // 3. 调用CANN算子库API，需要修改为具体的Api名称
198+ uint64_t workspaceSize = 0;
199+ aclOpExecutor* executor;
200+ // 调用aclnnMseLossBackward第一段接口
201+ ret = aclnnMseLossBackwardGetWorkspaceSize(gradOutput, self, target, reduction, out, &workspaceSize, &executor);
202+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnMseLossBackwardGetWorkspaceSize failed. ERROR: %d\n", ret);
203+ return ret);
204+ // 根据第一段接口计算出的workspaceSize申请device内存
205+ void* workspaceAddr = nullptr;
206+ if (workspaceSize > 0) {
207+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
208+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
209+ }
210+ // 调用aclnnMseLossBackward第二段接口
211+ ret = aclnnMseLossBackward(workspaceAddr, workspaceSize, executor, stream);
212+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnMseLossBackward failed. ERROR: %d\n", ret); return ret);
213+ 
214+ // 4. （固定写法）同步等待任务执行结束
215+ ret = aclrtSynchronizeStream(stream);
216+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
217+ 
218+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
219+ auto size = GetShapeSize(outShape);
220+ std::vector<float> resultData(size, 0);
221+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr,
222+ size * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
223+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
224+ for (int64_t i = 0; i < size; i++) {
225+ LOG_PRINT("result[%ld] is: %f\n", i, resultData[i]);
226+ }
227+ 
228+ // 6. 释放aclTensor和aclScalar，需要根据具体API的接口定义修改
229+ aclDestroyTensor(gradOutput);
230+ aclDestroyTensor(self);
231+ aclDestroyTensor(target);
232+ aclDestroyTensor(out);
233+ 
234+ // 7. 释放device资源，需要根据具体API的接口定义修改
235+ aclrtFree(gradOutputDeviceAddr);
236+ aclrtFree(selfDeviceAddr);
237+ aclrtFree(targetDeviceAddr);
238+ aclrtFree(outDeviceAddr);
239+ if (workspaceSize > 0) {
240+ aclrtFree(workspaceAddr);
241+ }
242+ aclrtDestroyStream(stream);
243+ aclrtResetDevice(deviceId);
244+ aclFinalize();
245+ return 0;
246+}
247+```
Aloss/mse_loss_grad/examples/test_aclnn_mse_loss_backward.cpp+160-0
@@ -0,0 +1,160 @@
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 <iostream>
12+#include <vector>
13+#include "acl/acl.h"
14+#include "aclnnop/aclnn_mse_loss_backward.h"
15+ 
16+#define CHECK_RET(cond, return_expr) \
17+ do { \
18+ if (!(cond)) { \
19+ return_expr; \
20+ } \
21+ } while (0)
22+ 
23+#define LOG_PRINT(message, ...) \
24+ do { \
25+ printf(message, ##__VA_ARGS__); \
26+ } while (0)
27+ 
28+int64_t GetShapeSize(const std::vector<int64_t>& shape) {
29+ int64_t shapeSize = 1;
30+ for (auto i : shape) {
31+ shapeSize *= i;
32+ }
33+ return shapeSize;
34+}
35+ 
36+int Init(int32_t deviceId, aclrtStream* stream) {
37+ // 固定写法，资源初始化
38+ auto ret = aclInit(nullptr);
39+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
40+ ret = aclrtSetDevice(deviceId);
41+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
42+ ret = aclrtCreateStream(stream);
43+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
44+ return 0;
45+}
46+ 
47+template <typename T>
48+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
49+ aclDataType dataType, aclTensor** tensor) {
50+ auto size = GetShapeSize(shape) * sizeof(T);
51+ // 调用aclrtMalloc申请device侧内存
52+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
53+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
54+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
55+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
56+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
57+ 
58+ // 计算连续tensor的strides
59+ std::vector<int64_t> strides(shape.size(), 1);
60+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
61+ strides[i] = shape[i + 1] * strides[i + 1];
62+ }
63+ 
64+ // 调用aclCreateTensor接口创建aclTensor
65+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
66+ shape.data(), shape.size(), *deviceAddr);
67+ return 0;
68+}
69+ 
70+int main() {
71+ // 1. （固定写法）device/stream初始化，参考acl API手册
72+ // 根据自己的实际device填写deviceId
73+ int32_t deviceId = 0;
74+ aclrtStream stream;
75+ auto ret = Init(deviceId, &stream);
76+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
77+ 
78+ // 2. 构造输入与输出，需要根据API的接口自定义构造gradOutput
79+ std::vector<int64_t> gradOutputShape = {2, 2};
80+ std::vector<int64_t> selfShape = {2, 2};
81+ std::vector<int64_t> targetShape = {2, 2};
82+ std::vector<int64_t> outShape = {2, 2};
83+ void* gradOutputDeviceAddr = nullptr;
84+ void* selfDeviceAddr = nullptr;
85+ void* targetDeviceAddr = nullptr;
86+ void* outDeviceAddr = nullptr;
87+ aclTensor* gradOutput = nullptr;
88+ aclTensor* self = nullptr;
89+ aclTensor* target = nullptr;
90+ aclTensor* out = nullptr;
91+ std::vector<float> gradOutputHostData = {0, 1, 2, 3};
92+ std::vector<float> selfHostData = {0, 1, 2, 3};
93+ std::vector<float> targetHostData = {1, 1, 1, 1};
94+ std::vector<float> outHostData(4, 0);
95+ // 创建gradOutput aclTensor
96+ ret = CreateAclTensor(gradOutputHostData, gradOutputShape, &gradOutputDeviceAddr,
97+ aclDataType::ACL_FLOAT, &gradOutput);
98+ CHECK_RET(ret == ACL_SUCCESS, return ret);
99+ // 创建self aclTensor
100+ ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self);
101+ CHECK_RET(ret == ACL_SUCCESS, return ret);
102+ // 创建target aclTensor
103+ ret = CreateAclTensor(targetHostData, targetShape, &targetDeviceAddr, aclDataType::ACL_FLOAT, &target);
104+ CHECK_RET(ret == ACL_SUCCESS, return ret);
105+ // 创建out aclTensor
106+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
107+ CHECK_RET(ret == ACL_SUCCESS, return ret);
108+ // 创建reduction
109+ int64_t reduction = 1;
110+ 
111+ // 3. 调用CANN算子库API，需要修改为具体的Api名称
112+ uint64_t workspaceSize = 0;
113+ aclOpExecutor* executor;
114+ // 调用aclnnMseLossBackward第一段接口
115+ ret = aclnnMseLossBackwardGetWorkspaceSize(gradOutput, self, target, reduction, out, &workspaceSize, &executor);
116+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnMseLossBackwardGetWorkspaceSize failed. ERROR: %d\n", ret);
117+ return ret);
118+ // 根据第一段接口计算出的workspaceSize申请device内存
119+ void* workspaceAddr = nullptr;
120+ if (workspaceSize > 0) {
121+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
122+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
123+ }
124+ // 调用aclnnMseLossBackward第二段接口
125+ ret = aclnnMseLossBackward(workspaceAddr, workspaceSize, executor, stream);
126+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnMseLossBackward failed. ERROR: %d\n", ret); return ret);
127+ 
128+ // 4. （固定写法）同步等待任务执行结束
129+ ret = aclrtSynchronizeStream(stream);
130+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
131+ 
132+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
133+ auto size = GetShapeSize(outShape);
134+ std::vector<float> resultData(size, 0);
135+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr,
136+ size * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
137+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
138+ for (int64_t i = 0; i < size; i++) {
139+ LOG_PRINT("result[%ld] is: %f\n", i, resultData[i]);
140+ }
141+ 
142+ // 6. 释放aclTensor和aclScalar，需要根据具体API的接口定义修改
143+ aclDestroyTensor(gradOutput);
144+ aclDestroyTensor(self);
145+ aclDestroyTensor(target);
146+ aclDestroyTensor(out);
147+ 
148+ // 7. 释放device资源，需要根据具体API的接口定义修改
149+ aclrtFree(gradOutputDeviceAddr);
150+ aclrtFree(selfDeviceAddr);
151+ aclrtFree(targetDeviceAddr);
152+ aclrtFree(outDeviceAddr);
153+ if (workspaceSize > 0) {
154+ aclrtFree(workspaceAddr);
155+ }
156+ aclrtDestroyStream(stream);
157+ aclrtResetDevice(deviceId);
158+ aclFinalize();
159+ return 0;
160+}
Aloss/mse_loss_grad/op_graph/CMakelists.txt+11-0
@@ -0,0 +1,11 @@
1+# ----------------------------------------------------------------------------
2+# This program is free software, you can redistribute it and/or modify.
3+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
4+# This file is a part of the CANN Open Software.
5+# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6+# Please refer to the License for details. You may not use this file except in compliance with the License.
7+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------
10+ 
11+message(STATUS "=== DEBUG: start ops.loss.mse_loss_grad.graph_plugin.CMakeLists.txt")
Aloss/mse_loss_grad/op_graph/mse_loss_grad_proto.h+50-0
@@ -0,0 +1,50 @@
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+/*!
CANN-robot
CANN-robotCANN-robot2025年12月26日

代码结构与可维护性: 文件头注释中的文件名与实际文件名不一致。注释中描述的文件名是'nn_norm_ops.h'，但实际文件名是'mse_loss_grad_proto.h'。这种不一致会给代码维护带来困惑，特别是在大型项目中查找文件时。

问题类型: 代码结构与可维护性 文件路径: loss/mse_loss_grad/op_graph/mse_loss_grad_proto.h 行号: 11 问题代码:

/*!
 * \file nn_norm_ops.h
 * \brief
 */

修改建议:

将文件头注释中的文件名修改为实际文件名'mse_loss_grad_proto.h'，以保持一致性。

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

likedislike
12+ * \file nn_norm_ops.h
13+ * \brief
14+ */
15+#ifndef OPS_BUILT_IN_OP_PROTO_INC_NN_NORM_OPS_H_
CANN-robot
CANN-robotCANN-robot2025年12月26日

代码结构与可维护性: 头文件保护宏的名称与文件实际内容不匹配。宏定义为'OPS_BUILT_IN_OP_PROTO_INC_NN_NORM_OPS_H_'，暗示这是关于归一化操作的头文件，但实际内容定义的是MseLossGrad操作。这种命名不一致可能导致理解错误和潜在的宏冲突。

问题类型: 代码结构与可维护性 文件路径: loss/mse_loss_grad/op_graph/mse_loss_grad_proto.h 行号: 15 问题代码:

#ifndef OPS_BUILT_IN_OP_PROTO_INC_NN_NORM_OPS_H_
#define OPS_BUILT_IN_OP_PROTO_INC_NN_NORM_OPS_H_

修改建议:

将头文件保护宏修改为与文件内容相关的名称，例如'OPS_BUILT_IN_OP_PROTO_INC_MSE_LOSS_GRAD_H_'或与文件名保持一致。

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

likedislike
16+#define OPS_BUILT_IN_OP_PROTO_INC_NN_NORM_OPS_H_
17+ 
18+#include "graph/operator_reg.h"
19+#include "graph/types.h"
20+ 
21+namespace ge {
22+ 
23+/**
24+* @brief Computes gradients of mse loss.
25+ 
26+* @par Inputs:
27+* @li predict: An ND tensor of type float16, float32 or bfloat16.
28+* @li label: An ND tensor of type float16, float32 or bfloat16.
29+* @li dout: An ND tensor of type float16, float32 or bfloat16. \n
30+ 
31+* @par Attributes:
32+* reduction: An optional string.Defaults to "mean". \n
33+ 
34+* @par Outputs:
35+* y: An ND tensor tensor with the same shape and type as "predict". \n
36+ 
37+* @par Third-party framework compatibility
38+* Compatible with the Pytorch operator MseLossGrad.
39+*/
40+REG_OP(MseLossGrad)
41+ .INPUT(predict, TensorType({DT_FLOAT32, DT_FLOAT16, DT_BF16}))
42+ .INPUT(label, TensorType({DT_FLOAT32, DT_FLOAT16, DT_BF16}))
43+ .INPUT(dout, TensorType({DT_FLOAT32, DT_FLOAT16, DT_BF16}))
44+ .OUTPUT(y, TensorType({DT_FLOAT32, DT_FLOAT16, DT_BF16}))
45+ .ATTR(reduction, String, "mean")
46+ .OP_END_FACTORY_REG(MseLossGrad)
47+ 
48+ 
49+} // namespace ge
50+#endif // OPS_BUILT_IN_OP_PROTO_INC_NN_NORM_OPS_H_
CANN-robot
CANN-robotCANN-robot2025年12月26日

文件格式问题: 文件末尾缺少换行符。在 Git 等版本控制系统中，缺少末尾换行符可能导致警告或影响 diff 的显示。根据 POSIX 标准，文本文件的每一行（包括最后一行）都应以换行符结束。

问题类型: 文件格式问题 文件路径: loss/mse_loss_grad/op_graph/mse_loss_grad_proto.h 行号: 50 问题代码:

#endif  // OPS_BUILT_IN_OP_PROTO_INC_NN_NORM_OPS_H_

修改建议:

在文件末尾添加一个换行符。这通常可以通过文本编辑器或 IDE 的自动格式化功能完成。

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

likedislike
Aloss/mse_loss_grad/op_host/arch35/mse_loss_grad_tiling_arch35.cpp+283-0
@@ -0,0 +1,283 @@
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 mse_loss_grad_tiling.cc
13+ * \brief mse_loss_grad_tiling
14+ */
15+ 
16+#include <graph/utils/type_utils.h>
17+#include "mse_loss_grad_tiling_arch35.h"
18+#include "log/log.h"
19+#include "atvoss/broadcast/broadcast_tiling.h"
20+#include "loss/mse_loss_grad/op_kernel/arch35/mse_loss_grad_tiling_key.h"
21+#include "loss/mse_loss_grad/op_kernel/arch35/mse_loss_grad_dag.h"
22+ 
23+using namespace AscendC;
24+using namespace ge;
25+using namespace Ops::Base;
26+ 
27+namespace optiling {
28+ 
29+constexpr static uint64_t COMMON_TILING_PRIORITY = 0;
30+constexpr static int32_t INPUT_PREDICT_IDX = 0;
31+constexpr static int32_t INPUT_LABEL_IDX = 1;
32+constexpr static int32_t INPUT_DOUT_IDX = 2;
33+constexpr static int32_t OUTPUT_IDX = 0;
34+static const std::map<std::string, uint32_t> STR_2_INT = {{"none", 0}, {"sum", 1}, {"mean", 2}};
35+ 
36+inline const gert::Shape &EnsureNotScalar(const gert::Shape &in_shape) {
37+ if (in_shape.IsScalar()) {
38+ return g_vec_1_shape;
39+ }
40+ return in_shape;
41+}
42+ 
43+ge::graphStatus MseLossGradTilingClass::GetShapeAttrsInfo()
44+{
45+ auto predictDesc = context_->GetInputDesc(INPUT_PREDICT_IDX);
46+ OP_CHECK_NULL_WITH_CONTEXT(context_, predictDesc);
47+ auto labelDesc = context_->GetInputDesc(INPUT_LABEL_IDX);
48+ OP_CHECK_NULL_WITH_CONTEXT(context_, labelDesc);
49+ auto doutDesc = context_->GetInputDesc(INPUT_DOUT_IDX);
50+ OP_CHECK_NULL_WITH_CONTEXT(context_, doutDesc);
51+ auto outputDesc = context_->GetOutputDesc(OUTPUT_IDX);
52+ OP_CHECK_NULL_WITH_CONTEXT(context_, outputDesc);
53+ 
54+ ge::DataType predictDType = predictDesc->GetDataType();
55+ ge::DataType labelDType = labelDesc->GetDataType();
56+ ge::DataType doutDType = doutDesc->GetDataType();
57+ ge::DataType outputDType = outputDesc->GetDataType();
58+ OP_CHECK_IF((predictDType != labelDType),
59+ OP_LOGE(context_->GetNodeName(), "dtype of predict[%s] and dtype of label[%s] not same",
60+ ge::TypeUtils::DataTypeToSerialString(predictDType).c_str(),
61+ ge::TypeUtils::DataTypeToSerialString(labelDType).c_str()),
62+ return ge::GRAPH_FAILED);
63+ 
64+ OP_CHECK_IF((predictDType != doutDType),
65+ OP_LOGE(context_->GetNodeName(), "dtype of predict[%s] and dtype of dout[%s] not same",
66+ ge::TypeUtils::DataTypeToSerialString(predictDType).c_str(),
67+ ge::TypeUtils::DataTypeToSerialString(doutDType).c_str()),
68+ return ge::GRAPH_FAILED);
69+ 
70+ OP_CHECK_IF((predictDType != outputDType),
71+ OP_LOGE(context_->GetNodeName(), "dtype of predict[%s] and dtype of y[%s] not same",
72+ ge::TypeUtils::DataTypeToSerialString(predictDType).c_str(),
73+ ge::TypeUtils::DataTypeToSerialString(outputDType).c_str()),
74+ return ge::GRAPH_FAILED);
75+ this->inputDtype = predictDType;
76+ return ge::GRAPH_SUCCESS;
77+}
78+ 
79+bool MseLossGradTilingClass::IsCapable()
80+{
81+ return true;
82+}
83+ 
84+ge::graphStatus MseLossGradTilingClass::CheckDoutIsScalar()
85+{
86+ auto inputDoutShape = context_->GetInputShape(INPUT_DOUT_IDX);
87+ OP_CHECK_NULL_WITH_CONTEXT(context_, inputDoutShape);
88+ auto storageShape = inputDoutShape->GetStorageShape();
89+ if (storageShape.IsScalar() || storageShape.GetShapeSize() == 1) {
CANN-robot
CANN-robotCANN-robot2025年12月26日

变量初始化: 第86-93行的CheckDoutIsScalar函数中，当storageShape不是标量且GetShapeSize() != 1时，doutIsScalar成员变量不会被赋值。在DoOpTiling函数中直接使用this->doutIsScalar进行比较，如果doutIsScalar未初始化，将使用未定义的值。

问题类型: 变量初始化 文件路径: loss/mse_loss_grad/op_host/arch35/mse_loss_grad_tiling_arch35.cpp 行号: 89 问题代码:

    if (storageShape.IsScalar() || storageShape.GetShapeSize() == 1) {
        this->doutIsScalar = static_cast<uint32_t>(ATTR_IS_TRUE);
    }

修改建议:

在函数开头或类构造函数中初始化doutIsScalar为默认值（ATTR_IS_FALSE）：
    this->doutIsScalar = static_cast<uint32_t>(ATTR_IS_FALSE);
    if (storageShape.IsScalar() || storageShape.GetShapeSize() == 1) {
        this->doutIsScalar = static_cast<uint32_t>(ATTR_IS_TRUE);
    }

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

likedislike
90+ this->doutIsScalar = static_cast<uint32_t>(ATTR_IS_TRUE);
91+ }
92+ OP_LOGD(context_->GetNodeName(), "[TilingData] : doutIsScalar is %u", this->doutIsScalar);
93+ return ge::GRAPH_SUCCESS;
94+}
95+ 
96+ge::graphStatus MseLossGradTilingClass::CalcReduceMeanCof()
97+{
98+ auto attrs = context_->GetAttrs();
99+ OP_CHECK_NULL_WITH_CONTEXT(context_, attrs);
100+ 
101+ this->reducationStr = attrs->GetAttrPointer<char>(0);
102+ auto iter = STR_2_INT.find(this->reducationStr);
103+ OP_CHECK_IF((iter == STR_2_INT.end()),
104+ OP_LOGE(context_->GetNodeName(), "reduction is not in [none, mean, sum]"),
105+ return ge::GRAPH_FAILED);
106+ this->reduceMeanCof = 2.0f;
107+ int64_t dimVal = 1;
108+ if (strcmp(this->reducationStr, "mean") == 0) {
109+ auto inputStorageShape = context_->GetInputShape(INPUT_PREDICT_IDX);
110+ OP_CHECK_NULL_WITH_CONTEXT(context_, inputStorageShape);
111+ const gert::Shape& inputShape = EnsureNotScalar(inputStorageShape->GetStorageShape());
112+ const size_t dimLen = inputShape.GetDimNum();
113+ for (uint32_t i = 0; i < dimLen; i++) {
114+ if (inputShape.GetDim(i) != 0) {
115+ dimVal = dimVal * inputShape.GetDim(i);
116+ } else {
117+ OP_LOGE(context_->GetNodeName(), "the shape[%u] of output is 0, do not supported", i);
118+ return ge::GRAPH_FAILED;
119+ }
120+ }
121+ this->reduceMeanCof = static_cast<float>(this->reduceMeanCof / static_cast<double>(dimVal));
122+ }
123+ OP_LOGD(context_->GetNodeName(), "[TilingData] : reduceMeanCof = %f", this->reduceMeanCof);
124+ return ge::GRAPH_SUCCESS;
125+}
126+ 
127+ge::graphStatus MseLossGradTilingClass::DoScalarDagOpTiling()
128+{
129+ if (this->inputDtype == ge::DT_FLOAT16) {
130+ BroadcastBaseTiling<MseLossGradOp::MseLossGradScalarDag<half, float>::OpDag> brcBaseTiling(context_);
131+ brcBaseTiling.SetScalar(this->reduceMeanCof);
132+ OP_CHECK_IF((brcBaseTiling.DoTiling() == ge::GRAPH_FAILED),
133+ OP_LOGE(context_->GetNodeName(), "Do tiling failed. Please check the detailed log."),
134+ return ge::GRAPH_FAILED);
135+ this->tilingKey = GET_TPL_TILING_KEY(brcBaseTiling.GetSchMode(), this->doutIsScalar);
136+ } else if (this->inputDtype == ge::DT_BF16) {
137+ BroadcastBaseTiling<MseLossGradOp::MseLossGradScalarDag<bfloat16_t, float>::OpDag> brcBaseTiling(context_);
138+ brcBaseTiling.SetScalar(this->reduceMeanCof);
139+ OP_CHECK_IF((brcBaseTiling.DoTiling() == ge::GRAPH_FAILED),
140+ OP_LOGE(context_->GetNodeName(), "Do tiling failed. Please check the detailed log."),
141+ return ge::GRAPH_FAILED);
142+ this->tilingKey = GET_TPL_TILING_KEY(brcBaseTiling.GetSchMode(), this->doutIsScalar);
143+ } else if (this->inputDtype == ge::DT_FLOAT) {
144+ BroadcastBaseTiling<MseLossGradOp::MseLossGradScalarDag<float, float>::OpDag> brcBaseTiling(context_);
145+ brcBaseTiling.SetScalar(this->reduceMeanCof);
146+ OP_CHECK_IF((brcBaseTiling.DoTiling() == ge::GRAPH_FAILED),
147+ OP_LOGE(context_->GetNodeName(), "Do tiling failed. Please check the detailed log."),
148+ return ge::GRAPH_FAILED);
149+ this->tilingKey = GET_TPL_TILING_KEY(brcBaseTiling.GetSchMode(), this->doutIsScalar);
150+ } else {
151+ OP_LOGE(context_->GetNodeName(), "input dtype is only support float16, bfloat16, float32, while got %s!",
152+ ge::TypeUtils::DataTypeToSerialString(inputDtype).c_str());
153+ return ge::GRAPH_FAILED;
154+ }
155+ return ge::GRAPH_SUCCESS;
156+}
157+ 
158+ge::graphStatus MseLossGradTilingClass::DoTensorDagOpTiling()
159+{
160+ if (this->inputDtype == ge::DT_FLOAT16) {
161+ BroadcastBaseTiling<MseLossGradOp::MseLossGradDag<half, float>::OpDag> brcBaseTiling(context_);
162+ brcBaseTiling.SetScalar(this->reduceMeanCof);
163+ OP_CHECK_IF((brcBaseTiling.DoTiling() == ge::GRAPH_FAILED),
164+ OP_LOGE(context_->GetNodeName(), "Do tiling failed. Please check the detailed log."),
165+ return ge::GRAPH_FAILED);
166+ this->tilingKey = GET_TPL_TILING_KEY(brcBaseTiling.GetSchMode(), this->doutIsScalar);
167+ } else if (this->inputDtype == ge::DT_BF16) {
168+ BroadcastBaseTiling<MseLossGradOp::MseLossGradDag<bfloat16_t, float>::OpDag> brcBaseTiling(context_);
169+ brcBaseTiling.SetScalar(this->reduceMeanCof);
170+ OP_CHECK_IF((brcBaseTiling.DoTiling() == ge::GRAPH_FAILED),
171+ OP_LOGE(context_->GetNodeName(), "Do tiling failed. Please check the detailed log."),
172+ return ge::GRAPH_FAILED);
173+ this->tilingKey = GET_TPL_TILING_KEY(brcBaseTiling.GetSchMode(), this->doutIsScalar);
174+ } else if (this->inputDtype == ge::DT_FLOAT) {
175+ BroadcastBaseTiling<MseLossGradOp::MseLossGradDag<float, float>::OpDag> brcBaseTiling(context_);
176+ brcBaseTiling.SetScalar(this->reduceMeanCof);
177+ OP_CHECK_IF((brcBaseTiling.DoTiling() == ge::GRAPH_FAILED),
178+ OP_LOGE(context_->GetNodeName(), "Do tiling failed. Please check the detailed log."),
179+ return ge::GRAPH_FAILED);
180+ this->tilingKey = GET_TPL_TILING_KEY(brcBaseTiling.GetSchMode(), this->doutIsScalar);
181+ } else {
182+ OP_LOGE(context_->GetNodeName(), "input tensor dtype is only support float16, bfloat16, float32, while got %s!",
183+ ge::TypeUtils::DataTypeToSerialString(inputDtype).c_str());
184+ return ge::GRAPH_FAILED;
185+ }
186+ 
187+ return ge::GRAPH_SUCCESS;
188+}
189+ 
190+ge::graphStatus MseLossGradTilingClass::DoOpTiling()
191+{
192+ OP_CHECK_IF(CheckDoutIsScalar() == ge::GRAPH_FAILED,
193+ OP_LOGE(context_->GetNodeName(), "check dout is scalar failed"),
194+ return ge::GRAPH_FAILED);
195+ OP_CHECK_IF(CalcReduceMeanCof() == ge::GRAPH_FAILED,
196+ OP_LOGE(context_->GetNodeName(), "get reduceMeanCof failed"), return ge::GRAPH_FAILED);
197+ if (this->doutIsScalar == static_cast<uint32_t>(ATTR_IS_TRUE)) {
CANN-robotCANN-robot
CANN-robotCANN-robot2025年12月26日

代码结构与可维护性: 第197-205行的条件分支中，this->doutIsScalar与ATTR_IS_TRUE比较，但ATTR_IS_TRUE未在代码片段中定义。这种魔法数字/常量应该定义为命名常量，提高代码可读性。

问题类型: 代码结构与可维护性 文件路径: loss/mse_loss_grad/op_host/arch35/mse_loss_grad_tiling_arch35.cpp 行号: 197 问题代码:

    if (this->doutIsScalar == static_cast<uint32_t>(ATTR_IS_TRUE)) {

修改建议:

在文件头部定义常量：
constexpr static uint32_t ATTR_IS_TRUE = 1;
constexpr static uint32_t ATTR_IS_FALSE = 0;

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

likedislike
CANN-robotCANN-robot2025年12月26日

未初始化变量使用: 在DoOpTiling函数中，使用this->doutIsScalar进行比较，但该变量在CheckDoutIsScalar中只有在条件满足时才被赋值。如果storageShape既不是标量也不是大小为1，doutIsScalar将保持未初始化状态（假设类成员变量未在构造函数中初始化）。

问题类型: 未初始化变量使用 文件路径: loss/mse_loss_grad/op_host/arch35/mse_loss_grad_tiling_arch35.cpp 行号: 197 问题代码:

if (this->doutIsScalar == static_cast<uint32_t>(ATTR_IS_TRUE)) {

修改建议:

在CheckDoutIsScalar函数中确保doutIsScalar有明确的初始值：
this->doutIsScalar = static_cast<uint32_t>(ATTR_IS_FALSE); // 假设ATTR_IS_FALSE存在
if (storageShape.IsScalar() || storageShape.GetShapeSize() == 1) {
    this->doutIsScalar = static_cast<uint32_t>(ATTR_IS_TRUE);
}

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

likedislike
198+ OP_CHECK_IF(DoScalarDagOpTiling() == ge::GRAPH_FAILED,
199+ OP_LOGE(context_->GetNodeName(), "do tiling failed, when dout is scalar"),
200+ return ge::GRAPH_FAILED);
201+ } else {
202+ OP_CHECK_IF(DoTensorDagOpTiling() == ge::GRAPH_FAILED,
203+ OP_LOGE(context_->GetNodeName(), "do tiling failed, when dout is tensor"),
204+ return ge::GRAPH_FAILED);
205+ }
206+ return ge::GRAPH_SUCCESS;
207+}
208+ 
209+ge::graphStatus MseLossGradTilingClass::DoLibApiTiling()
210+{
211+ return ge::GRAPH_SUCCESS;
212+}
213+ 
214+uint64_t MseLossGradTilingClass::GetTilingKey() const
215+{
216+ return tilingKey;
217+}
218+ 
219+ge::graphStatus MseLossGradTilingClass::GetWorkspaceSize()
220+{
221+ return ge::GRAPH_SUCCESS;
222+}
223+ 
224+ge::graphStatus MseLossGradTilingClass::PostTiling()
225+{
226+ return ge::GRAPH_SUCCESS;
227+}
228+ 
229+ge::graphStatus MseLossGradTilingClass::GetPlatformInfo()
230+{
231+ return ge::GRAPH_SUCCESS;
232+}
233+ 
234+ge::graphStatus TilingForMseLossGrad(gert::TilingContext* context)
235+{
236+ OP_LOGD("MseLossGradTiling", "Enter TilingForMseLossGrad");
237+ if (context == nullptr) {
238+ OP_LOGE("MseLossGradTiling", "Tiling context is nullptr");
239+ return ge::GRAPH_FAILED;
240+ }
241+ 
242+ OP_LOGD(context, "Enter ascendc MseLossGradTiling");
243+ return Ops::NN::Optiling::TilingRegistry::GetInstance().DoTilingImpl(context);
244+}
245+ 
246+ge::graphStatus MseLossGradTilingPrepareAscendC(gert::TilingParseContext* context)
247+{
248+ fe::PlatFormInfos* platformInfo = context->GetPlatformInfo();
249+ auto compileInfo = context->GetCompiledInfo<MseLossGradCompileInfo>();
250+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfo);
251+ 
252+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo);
253+ compileInfo->coreNum = ascendcPlatform.GetCoreNumAiv();
254+ OP_CHECK_IF((compileInfo->coreNum <= 0),
255+ OP_LOGE(context->GetNodeName(), "Get core num failed, core num: %lu", compileInfo->coreNum),
256+ return ge::GRAPH_FAILED);
257+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, compileInfo->ubSize);
258+ 
259+ OP_CHECK_IF((compileInfo->ubSize <= 0),
260+ OP_LOGE(context->GetNodeName(), "Get ub size failed, ub size: %lu", compileInfo->ubSize),
261+ return ge::GRAPH_FAILED);
262+ return ge::GRAPH_SUCCESS;
263+}
264+ 
265+ge::graphStatus TilingPrepareForMseLossGrad(gert::TilingParseContext* context)
266+{
267+ OP_LOGD("TilingPrepareForMseLossGrad", "Enter TilingPrepareForMseLossGrad");
268+ if (context == nullptr) {
269+ OP_LOGE("TilingPrepareForMseLossGrad", "Tiling context is nullptr");
270+ return ge::GRAPH_FAILED;
271+ }
272+ 
273+ OP_LOGD(context, "Enter MseLossGradTilingPrepareAscendC");
274+ return MseLossGradTilingPrepareAscendC(context);
275+}
276+ 
277+IMPL_OP_OPTILING(MseLossGrad)
278+ .Tiling(TilingForMseLossGrad)
279+ .TilingParse<MseLossGradCompileInfo>(TilingPrepareForMseLossGrad);
280+ 
281+REGISTER_OPS_TILING_TEMPLATE(MseLossGrad, MseLossGradTilingClass, COMMON_TILING_PRIORITY);
282+ 
283+} // namespace optiling
Aloss/mse_loss_grad/op_host/arch35/mse_loss_grad_tiling_arch35.h+59-0
@@ -0,0 +1,59 @@
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 mse_loss_grad_tiling.h
13+ * \brief mse_loss_grad_tiling
14+ */
15+ 
16+#ifndef OPS_BUILD_IN_OP_TILING_RUNTIME_MSE_LOSS_GRAD_TILING_H
17+#define OPS_BUILD_IN_OP_TILING_RUNTIME_MSE_LOSS_GRAD_TILING_H
18+ 
19+#include "register/op_impl_registry.h"
20+#include "tiling_base/tiling_base.h"
21+#include "tiling_base/tiling_templates_registry.h"
22+ 
23+namespace optiling {
24+ 
25+class MseLossGradTilingClass : public Ops::NN::Optiling::TilingBaseClass {
26+public:
27+ explicit MseLossGradTilingClass(gert::TilingContext* context) : Ops::NN::Optiling::TilingBaseClass(context)
28+ {}
29+ 
30+protected:
31+ bool IsCapable() override;
32+ ge::graphStatus GetPlatformInfo() override;
33+ ge::graphStatus GetShapeAttrsInfo() override;
34+ ge::graphStatus DoOpTiling() override;
35+ ge::graphStatus DoLibApiTiling() override;
36+ uint64_t GetTilingKey() const override;
37+ ge::graphStatus GetWorkspaceSize() override;
38+ ge::graphStatus PostTiling() override;
39+ ge::graphStatus CalcReduceMeanCof();
40+ ge::graphStatus CheckDoutIsScalar();
41+ ge::graphStatus DoScalarDagOpTiling();
42+ ge::graphStatus DoTensorDagOpTiling();
43+ 
44+private:
45+ uint64_t tilingKey = 0;
46+ const char* reducationStr = "";
CANN-robot
CANN-robotCANN-robot2025年12月26日

拼写错误: 变量名 reducationStr 疑似拼写错误，正确的应该是 reductionStr（减少一个'd'）。拼写错误可能导致代码可读性下降，也可能在后续代码中使用时造成混淆。

问题类型: 拼写错误 文件路径: loss/mse_loss_grad/op_host/arch35/mse_loss_grad_tiling_arch35.h 行号: 46 问题代码:

const char* reducationStr = "";

修改建议:

将变量名更正为正确的拼写：`const char* reductionStr = "";`。同时检查代码库中所有使用该变量的地方，确保一致修改。

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

likedislike
47+ ge::DataType inputDtype;
48+ float reduceMeanCof = 2.0f;
49+ uint32_t doutIsScalar = 0;
50+};
51+ 
52+struct MseLossGradCompileInfo {
53+ uint64_t coreNum = 0;
54+ uint64_t ubSize = 0;
55+};
56+ 
57+} // namespace optiling
58+ 
59+#endif // OPS_BUILD_IN_OP_TILING_RUNTIME_MSE_LOSS_GRAD_TILING_H
Aloss/mse_loss_grad/op_host/config/ascend910_95/mse_loss_grad_binary.json+491-0
@@ -0,0 +1,491 @@
1+{
2+ "op_type": "MseLossGrad",
3+ "op_list": [
4+ {
5+ "bin_filename": "MseLossGrad_9985a68d4a894b5b70a5154886ddf7a3",
6+ "inputs": [
7+ {
8+ "name": "predict",
9+ "index": 0,
10+ "dtype": "bfloat16",
11+ "format": "ND",
12+ "paramType": "required",
13+ "shape": [
14+ -2
15+ ]
16+ },
17+ {
18+ "name": "label",
19+ "index": 1,
20+ "dtype": "bfloat16",
21+ "format": "ND",
22+ "paramType": "required",
23+ "shape": [
24+ -2
25+ ]
26+ },
27+ {
28+ "name": "dout",
29+ "index": 2,
30+ "dtype": "bfloat16",
31+ "format": "ND",
32+ "paramType": "required",
33+ "shape": [
34+ -2
35+ ]
36+ }
37+ ],
38+ "outputs": [
39+ {
40+ "name": "y",
41+ "index": 0,
42+ "dtype": "bfloat16",
43+ "format": "ND",
44+ "paramType": "required",
45+ "shape": [
46+ -2
47+ ]
48+ }
49+ ],
50+ "attrs": [
51+ {
52+ "name": "reduction",
53+ "dtype": "string",
54+ "value": "mean"
55+ }
56+ ]
57+ },
58+ {
59+ "bin_filename": "MseLossGrad_859d76b7bfb4ed6d9521ed0fec544b9f",
60+ "inputs": [
61+ {
62+ "name": "predict",
63+ "index": 0,
64+ "dtype": "bfloat16",
65+ "format": "ND",
66+ "paramType": "required",
67+ "shape": [
68+ -2
69+ ]
70+ },
71+ {
72+ "name": "label",
73+ "index": 1,
74+ "dtype": "bfloat16",
75+ "format": "ND",
76+ "paramType": "required",
77+ "shape": [
78+ -2
79+ ]
80+ },
81+ {
82+ "name": "dout",
83+ "index": 2,
84+ "dtype": "bfloat16",
85+ "format": "ND",
86+ "paramType": "required",
87+ "shape": [
88+ -2
89+ ]
90+ }
91+ ],
92+ "outputs": [
93+ {
94+ "name": "y",
95+ "index": 0,
96+ "dtype": "bfloat16",
97+ "format": "ND",
98+ "paramType": "required",
99+ "shape": [
100+ -2
101+ ]
102+ }
103+ ],
104+ "attrs": [
105+ {
106+ "name": "reduction",
107+ "dtype": "string",
108+ "value": "none"
109+ }
110+ ]
111+ },
112+ {
113+ "bin_filename": "MseLossGrad_4ed4454eae143978153c8f14d809bc3c",
114+ "inputs": [
115+ {
116+ "name": "predict",
117+ "index": 0,
118+ "dtype": "bfloat16",
119+ "format": "ND",
120+ "paramType": "required",
121+ "shape": [
122+ -2
123+ ]
124+ },
125+ {
126+ "name": "label",
127+ "index": 1,
128+ "dtype": "bfloat16",
129+ "format": "ND",
130+ "paramType": "required",
131+ "shape": [
132+ -2
133+ ]
134+ },
135+ {
136+ "name": "dout",
137+ "index": 2,
138+ "dtype": "bfloat16",
139+ "format": "ND",
140+ "paramType": "required",
141+ "shape": [
142+ -2
143+ ]
144+ }
145+ ],
146+ "outputs": [
147+ {
148+ "name": "y",
149+ "index": 0,
150+ "dtype": "bfloat16",
151+ "format": "ND",
152+ "paramType": "required",
153+ "shape": [
154+ -2
155+ ]
156+ }
157+ ],
158+ "attrs": [
159+ {
160+ "name": "reduction",
161+ "dtype": "string",
162+ "value": "sum"
163+ }
164+ ]
165+ },
166+ {
167+ "bin_filename": "MseLossGrad_34bb81d9f78c504c0dfd2fd2fa5d690a",
168+ "inputs": [
169+ {
170+ "name": "predict",
171+ "index": 0,
172+ "dtype": "float16",
173+ "format": "ND",
174+ "paramType": "required",
175+ "shape": [
176+ -2
177+ ]
178+ },
179+ {
180+ "name": "label",
181+ "index": 1,
182+ "dtype": "float16",
183+ "format": "ND",
184+ "paramType": "required",
185+ "shape": [
186+ -2
187+ ]
188+ },
189+ {
190+ "name": "dout",
191+ "index": 2,
192+ "dtype": "float16",
193+ "format": "ND",
194+ "paramType": "required",
195+ "shape": [
196+ -2
197+ ]
198+ }
199+ ],
200+ "outputs": [
201+ {
202+ "name": "y",
203+ "index": 0,
204+ "dtype": "float16",
205+ "format": "ND",
206+ "paramType": "required",
207+ "shape": [
208+ -2
209+ ]
210+ }
211+ ],
212+ "attrs": [
213+ {
214+ "name": "reduction",
215+ "dtype": "string",
216+ "value": "mean"
217+ }
218+ ]
219+ },
220+ {
221+ "bin_filename": "MseLossGrad_f7921eccb5136ca755254840ef1b8968",
222+ "inputs": [
223+ {
224+ "name": "predict",
225+ "index": 0,
226+ "dtype": "float16",
227+ "format": "ND",
228+ "paramType": "required",
229+ "shape": [
230+ -2
231+ ]
232+ },
233+ {
234+ "name": "label",
235+ "index": 1,
236+ "dtype": "float16",
237+ "format": "ND",
238+ "paramType": "required",
239+ "shape": [
240+ -2
241+ ]
242+ },
243+ {
244+ "name": "dout",
245+ "index": 2,
246+ "dtype": "float16",
247+ "format": "ND",
248+ "paramType": "required",
249+ "shape": [
250+ -2
251+ ]
252+ }
253+ ],
254+ "outputs": [
255+ {
256+ "name": "y",
257+ "index": 0,
258+ "dtype": "float16",
259+ "format": "ND",
260+ "paramType": "required",
261+ "shape": [
262+ -2
263+ ]
264+ }
265+ ],
266+ "attrs": [
267+ {
268+ "name": "reduction",
269+ "dtype": "string",
270+ "value": "none"
271+ }
272+ ]
273+ },
274+ {
275+ "bin_filename": "MseLossGrad_ac5523d5903675b7a32625a2aecdea65",
276+ "inputs": [
277+ {
278+ "name": "predict",
279+ "index": 0,
280+ "dtype": "float16",
281+ "format": "ND",
282+ "paramType": "required",
283+ "shape": [
284+ -2
285+ ]
286+ },
287+ {
288+ "name": "label",
289+ "index": 1,
290+ "dtype": "float16",
291+ "format": "ND",
292+ "paramType": "required",
293+ "shape": [
294+ -2
295+ ]
296+ },
297+ {
298+ "name": "dout",
299+ "index": 2,
300+ "dtype": "float16",
301+ "format": "ND",
302+ "paramType": "required",
303+ "shape": [
304+ -2
305+ ]
306+ }
307+ ],
308+ "outputs": [
309+ {
310+ "name": "y",
311+ "index": 0,
312+ "dtype": "float16",
313+ "format": "ND",
314+ "paramType": "required",
315+ "shape": [
316+ -2
317+ ]
318+ }
319+ ],
320+ "attrs": [
321+ {
322+ "name": "reduction",
323+ "dtype": "string",
324+ "value": "sum"
325+ }
326+ ]
327+ },
328+ {
329+ "bin_filename": "MseLossGrad_dd2c8ff8d413b8282514a161e2ae82b4",
330+ "inputs": [
331+ {
332+ "name": "predict",
333+ "index": 0,
334+ "dtype": "float32",
335+ "format": "ND",
336+ "paramType": "required",
337+ "shape": [
338+ -2
339+ ]
340+ },
341+ {
342+ "name": "label",
343+ "index": 1,
344+ "dtype": "float32",
345+ "format": "ND",
346+ "paramType": "required",
347+ "shape": [
348+ -2
349+ ]
350+ },
351+ {
352+ "name": "dout",
353+ "index": 2,
354+ "dtype": "float32",
355+ "format": "ND",
356+ "paramType": "required",
357+ "shape": [
358+ -2
359+ ]
360+ }
361+ ],
362+ "outputs": [
363+ {
364+ "name": "y",
365+ "index": 0,
366+ "dtype": "float32",
367+ "format": "ND",
368+ "paramType": "required",
369+ "shape": [
370+ -2
371+ ]
372+ }
373+ ],
374+ "attrs": [
375+ {
376+ "name": "reduction",
377+ "dtype": "string",
378+ "value": "mean"
379+ }
380+ ]
381+ },
382+ {
383+ "bin_filename": "MseLossGrad_be5c7ffea0cb6c370ee3b3f33a85c587",
384+ "inputs": [
385+ {
386+ "name": "predict",
387+ "index": 0,
388+ "dtype": "float32",
389+ "format": "ND",
390+ "paramType": "required",
391+ "shape": [
392+ -2
393+ ]
394+ },
395+ {
396+ "name": "label",
397+ "index": 1,
398+ "dtype": "float32",
399+ "format": "ND",
400+ "paramType": "required",
401+ "shape": [
402+ -2
403+ ]
404+ },
405+ {
406+ "name": "dout",
407+ "index": 2,
408+ "dtype": "float32",
409+ "format": "ND",
410+ "paramType": "required",
411+ "shape": [
412+ -2
413+ ]
414+ }
415+ ],
416+ "outputs": [
417+ {
418+ "name": "y",
419+ "index": 0,
420+ "dtype": "float32",
421+ "format": "ND",
422+ "paramType": "required",
423+ "shape": [
424+ -2
425+ ]
426+ }
427+ ],
428+ "attrs": [
429+ {
430+ "name": "reduction",
431+ "dtype": "string",
432+ "value": "none"
433+ }
434+ ]
435+ },
436+ {
437+ "bin_filename": "MseLossGrad_abbb884f8c684789ab1377b0e54868b9",
438+ "inputs": [
439+ {
440+ "name": "predict",
441+ "index": 0,
442+ "dtype": "float32",
443+ "format": "ND",
444+ "paramType": "required",
445+ "shape": [
446+ -2
447+ ]
448+ },
449+ {
450+ "name": "label",
451+ "index": 1,
452+ "dtype": "float32",
453+ "format": "ND",
454+ "paramType": "required",
455+ "shape": [
456+ -2
457+ ]
458+ },
459+ {
460+ "name": "dout",
461+ "index": 2,
462+ "dtype": "float32",
463+ "format": "ND",
464+ "paramType": "required",
465+ "shape": [
466+ -2
467+ ]
468+ }
469+ ],
470+ "outputs": [
471+ {
472+ "name": "y",
473+ "index": 0,
474+ "dtype": "float32",
475+ "format": "ND",
476+ "paramType": "required",
477+ "shape": [
478+ -2
479+ ]
480+ }
481+ ],
482+ "attrs": [
483+ {
484+ "name": "reduction",
485+ "dtype": "string",
486+ "value": "sum"
487+ }
488+ ]
489+ }
490+ ]
491+}
Aloss/mse_loss_grad/op_host/config/ascend910_95/mse_loss_grad_simplified_key.ini+2-0
@@ -0,0 +1,2 @@
1+[MseLossGrad]
2+default=0
Aloss/mse_loss_grad/op_host/mse_loss_grad_def.cpp+61-0
@@ -0,0 +1,61 @@
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 mse_loss_grad.cpp
13+ * \brief mse_loss_grad def
14+ */
15+ 
16+ #include <cstdint>
17+ #include "register/op_def_registry.h"
18+
19+ namespace ops{
20+ static const std::vector<ge::DataType> dataType = {
21+ ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT
22+ };
23+
24+ static const std::vector<ge::Format> dataFormat = {
25+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND
26+ };
27+
28+ class MseLossGrad : public OpDef {
29+ public:
30+ explicit MseLossGrad(const char* name) : OpDef(name)
31+ {
32+ this->Input("predict")
33+ .ParamType(REQUIRED)
34+ .DataType(dataType)
35+ .Format(dataFormat);
36+ this->Input("label")
37+ .ParamType(REQUIRED)
38+ .DataType(dataType)
39+ .Format(dataFormat);
40+ this->Input("dout")
41+ .ParamType(REQUIRED)
42+ .DataType(dataType)
43+ .Format(dataFormat);
44+ this->Output("y")
45+ .ParamType(REQUIRED)
46+ .DataType(dataType)
47+ .Format(dataFormat);
48+ this->Attr("reduction").AttrType(OPTIONAL).String("mean");
49+
50+ OpAICoreConfig aicoreConfig;
51+ aicoreConfig.DynamicCompileStaticFlag(true)
52+ .DynamicRankSupportFlag(true)
53+ .DynamicShapeSupportFlag(true)
54+ .PrecisionReduceFlag(false)
55+ .ExtendCfgInfo("opFile.value", "mse_loss_grad_apt");
56+ this->AICore().AddConfig("ascend910_95", aicoreConfig);
57+ }
58+ };
59+
60+ OP_ADD(MseLossGrad);
61+ } // namespace ops
Aloss/mse_loss_grad/op_host/mse_loss_grad_infershape.cpp+51-0
@@ -0,0 +1,51 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file mse_loss_grad_infer.cc
13+ * \brief mse_loss_grad_infer
14+ */
15+ 
16+#include "register/op_impl_registry.h"
17+#include "infershape_broadcast_util.h"
18+#include "log/log.h"
19+ 
20+using namespace ge;
21+using namespace Ops::Base;
22+ 
23+namespace ops {
24+ 
25+ge::graphStatus Infershape4MseLossGrad(gert::InferShapeContext* context) {
26+ const size_t inputCount = 3;
27+ std::vector<const gert::Shape*> to_broadcast_shapes(inputCount);
28+ for (size_t i = 0; i < inputCount; i++) {
29+ auto in_shape = context->GetInputShape(i);
30+ OP_CHECK_NULL_WITH_CONTEXT(context, in_shape);
31+ to_broadcast_shapes[i] = in_shape;
32+ }
33+ auto out_shape = context->GetOutputShape(0);
34+ OP_CHECK_NULL_WITH_CONTEXT(context, out_shape);
35+ 
36+ OP_CHECK_IF(!BroadcastShape(to_broadcast_shapes, out_shape),
37+ OP_LOGE(context->GetNodeName(), "BroadcastShape failed!"), return ge::GRAPH_FAILED);
38+ 
39+ return ge::GRAPH_SUCCESS;
40+}
41+ 
42+ge::graphStatus InferDataType4MseLossGrad(gert::InferDataTypeContext* context) {
43+ OP_LOGD(context->GetNodeName(), "InferDataType4MseLossGrad enter");
44+ auto input_x_dtype = context->GetInputDataType(0);
45+ context->SetOutputDataType(0, input_x_dtype);
46+ OP_LOGD(context->GetNodeName(), "InferDataType4MseLossGrad end");
47+ return GRAPH_SUCCESS;
48+}
49+ 
50+IMPL_OP_INFERSHAPE(MseLossGrad).InferShape(Infershape4MseLossGrad).InferDataType(InferDataType4MseLossGrad);
51+} // namespace ops
Aloss/mse_loss_grad/op_kernel/arch35/mse_loss_grad_dag.h+67-0
@@ -0,0 +1,67 @@
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 mse_loss_grad_dag.h
13+ * \brief mse_loss_grad_dag
14+ */
15+ 
16+ #ifndef MSE_LOSS_GRAD_DAG_H
17+ #define MSE_LOSS_GRAD_DAG_H
18+ #include "atvoss/util/dag.h"
19+ #include "atvoss/util/vec.h"
20+ #include "atvoss/util/placeholder.h"
21+
22+ namespace MseLossGradOp{
23+ using namespace AscendC;
24+ using namespace Ops::Base;
25+
26+ template <typename T, typename U>
27+ struct MseLossGradDag {
28+ using OpInputPredict = Bind<Vec::CopyInBrc<T>, Placeholder::In0<T>>;
29+ using OpInputLabel = Bind<Vec::CopyInBrc<T>, Placeholder::In1<T>>;
30+ using OpInputDout = Bind<Vec::CopyInBrc<T>, Placeholder::In2<T>>;
31+
32+ using OpPredictCast = Bind<Vec::Cast<U, T, 0>, OpInputPredict>;
33+ using OpLabelCast = Bind<Vec::Cast<U, T, 0>, OpInputLabel>;
34+ using OpDoutCast = Bind<Vec::Cast<U, T, 0>, OpInputDout>;
35+ using OpSubRes = Bind<Vec::Sub<U>, OpPredictCast, OpLabelCast>;
36+ using OpNormGrad = Bind<Vec::Muls<U>, OpSubRes, Placeholder::Var<U, 0>>;
37+ using OpOutput = Bind<Vec::Mul<U>, OpNormGrad, OpDoutCast>;
38+ using OpOutputCast = Bind<Vec::Cast<T, U, 1>, OpOutput>;
39+
40+ using OpCopyOut = Bind<Vec::CopyOut<T>, Placeholder::Out0<T>, OpOutputCast>;
41+ using Outputs = Elems<OpCopyOut>;
42+ using MemCfg = MemOptCfg<MemLevel::LEVEL_2>;
43+ using OpDag = DAGSch<Outputs, void, MemCfg>;
44+ };
45+
46+ template <typename T, typename U>
47+ struct MseLossGradScalarDag {
48+ using OpInputPredict = Bind<Vec::CopyInBrc<T>, Placeholder::In0<T>>;
49+ using OpInputLabel = Bind<Vec::CopyInBrc<T>, Placeholder::In1<T>>;
50+ using OpInputDout = Bind<Vec::Duplicate<T>, Placeholder::In2<T, Placeholder::ScalarAttr<true>>>;
51+
52+ using OpPredictCast = Bind<Vec::Cast<U, T, 0>, OpInputPredict>;
53+ using OpLabelCast = Bind<Vec::Cast<U, T, 0>, OpInputLabel>;
54+ using OpDoutCast = Bind<Vec::Cast<U, T, 0>, OpInputDout>;
55+
56+ using OpSubRes = Bind<Vec::Sub<U>, OpPredictCast, OpLabelCast>;
57+ using OpNormGrad = Bind<Vec::Muls<U>, OpSubRes, Placeholder::Var<U, 0>>;
58+ using OpOutput = Bind<Vec::Mul<U>, OpNormGrad, OpDoutCast>;
59+ using OpOutputCast = Bind<Vec::Cast<T, U, 1>, OpOutput>;
60+ using OpCopyOut = Bind<Vec::CopyOut<T>, Placeholder::Out0<T>, OpOutputCast>;
61+
62+ using Outputs = Elems<OpCopyOut>;
63+ using MemCfg = MemOptCfg<MemLevel::LEVEL_2>;
64+ using OpDag = DAGSch<Outputs, void, MemCfg>;
65+ };
66+ }
67+ #endif // MSE_LOSS_GRAD_DAG_H
Aloss/mse_loss_grad/op_kernel/arch35/mse_loss_grad_tiling_key.h+34-0
@@ -0,0 +1,34 @@
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 mse_loss_grad_tiling_key.h
13+ * \brief mse_loss_grad_tiling_key
14+ */
15+ 
16+ #ifndef MSE_LOSS_GRAD_STRUCT_H
17+ #define MSE_LOSS_GRAD_STRUCT_H
18+
19+ #include "atvoss/broadcast/broadcast_base_struct.h"
20+
21+ #define ATTR_BIT_WIDTH 1
22+ #define ATTR_IS_TRUE 1
23+ // 算子自定义的tiling key字段
24+ ASCENDC_TPL_ARGS_DECL(MseLossGrad,
25+ BRC_TEMP_SCH_MODE_KEY_DECL(schMode),
26+ ASCENDC_TPL_UINT_DECL(doutIsScalar, ATTR_BIT_WIDTH, ASCENDC_TPL_UI_LIST, 0, ATTR_IS_TRUE)
27+ );
28+
29+ ASCENDC_TPL_SEL(
30+ ASCENDC_TPL_ARGS_SEL(BRC_TEMP_SCH_MODE_KEY_SEL(schMode),
31+ ASCENDC_TPL_UINT_SEL(doutIsScalar, ASCENDC_TPL_UI_LIST, 0, ATTR_IS_TRUE))
32+ );
33+
34+ #endif // MSE_LOSS_GRAD_STRUCT_H
Aloss/mse_loss_grad/op_kernel/mse_loss_grad_apt.cpp+35-0
@@ -0,0 +1,35 @@
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 mse_loss_grad.cpp
13+ * \brief mse_loss_grad
14+ */
15+ 
16+ #include "kernel_operator.h"
17+ #include "arch35/mse_loss_grad_dag.h"
18+ #include "arch35/mse_loss_grad_tiling_key.h"
19+ #include "atvoss/broadcast/broadcast_sch.h"
20+
21+ using namespace AscendC;
22+
23+ template <uint64_t schMode, uint32_t doutIsScalar>
24+ __global__ __aicore__ void mse_loss_grad(GM_ADDR predict, GM_ADDR label, GM_ADDR dout, GM_ADDR y, GM_ADDR workspace,
25+ GM_ADDR tiling) {
26+ if constexpr (doutIsScalar == static_cast<uint32_t>(ATTR_IS_TRUE)) {
27+ using OpDag = MseLossGradOp::MseLossGradScalarDag<DTYPE_PREDICT, float>::OpDag;
28+ Ops::Base::BroadcastSch<schMode, OpDag> sch(tiling);
29+ sch.Process(predict, label, dout, y);
30+ } else {
31+ using OpDag = MseLossGradOp::MseLossGradDag<DTYPE_PREDICT, float>::OpDag;
32+ Ops::Base::BroadcastSch<schMode, OpDag> sch(tiling);
33+ sch.Process(predict, label, dout, y);
34+ }
35+ }
Aloss/mse_loss_grad/tests/CMakeLists.txt+18-0
@@ -0,0 +1,18 @@
1+#
2+# This program is free software, you can redistribute it and/or modify.
3+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
4+# This file is a part of the CANN Open Software.
5+# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6+# Please refer to the License for details. You may not use this file except in compliance with the License.
7+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+#/
10+ 
11+message(STATUS "=== Debug: start ops.math.is_finite.tests.CMakeLists.txt ")
12+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
13+message(STATUS "=== Debug: CURRENT_DIRS =${CURRENT_DIRS} ")
14+foreach(SUB_DIR ${CURRENT_DIRS})
15+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
16+ add_subdirectory(${SUB_DIR})
17+ endif()
18+endforeach()
Aloss/mse_loss_grad/tests/ut/CMakeLists.txt+18-0
@@ -0,0 +1,18 @@
1+#
2+# This program is free software, you can redistribute it and/or modify.
3+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
4+# This file is a part of the CANN Open Software.
5+# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6+# Please refer to the License for details. You may not use this file except in compliance with the License.
7+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+#/
10+ 
11+# 每个目录下需要生成的可执行文件，具体参考：ops/built-in/test/CMakeLists.txt： 50~124
12+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
13+message(STATUS "=== Debug: CURRENT_DIRS =${CURRENT_DIRS} ")
14+foreach(SUB_DIR ${CURRENT_DIRS})
15+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
16+ add_subdirectory(${SUB_DIR})
17+ endif()
18+endforeach()
Aloss/mse_loss_grad/tests/ut/op_host/CMakeLists.txt+15-0
@@ -0,0 +1,15 @@
1+#
2+# This program is free software, you can redistribute it and/or modify.
3+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
4+# This file is a part of the CANN Open Software.
5+# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6+# Please refer to the License for details. You may not use this file except in compliance with the License.
7+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+#/
10+ 
11+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+if(UT_TEST_ALL OR OP_HOST_UT)
13+ add_modules_ut_sources(HOSTNAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
14+ add_modules_ut_sources(HOSTNAME ${OP_INFERSHAPE_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
15+endif()
Aloss/mse_loss_grad/tests/ut/op_host/test_mse_loss_grad_infershape.cpp+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+ * \file test_mse_loss_grad_proto.cpp
12+ * \brief
13+ */
14+ 
15+#include <iostream>
16+#include <gtest/gtest.h>
17+#include "register/op_impl_registry.h"
18+#include "kernel_run_context_facker.h"
19+#include "infershape_test_util.h"
20+#include "../../../op_graph/mse_loss_grad_proto.h"
21+#include "exe_graph/runtime/storage_format.h"
22+#include "exe_graph/runtime/storage_shape.h"
23+#include "log/log.h"
24+#include "ut_op_common.h"
25+#include "platform/platform_info.h"
26+#include "../../../../../tests/ut/common/any_value.h"
27+ 
28+class MseLossGradTest : public testing::Test {
29+protected:
30+ static void SetUpTestCase() {
31+ std::cout << "MseLossGrad Proto Test SetUp" << std::endl;
32+ }
33+ 
34+ static void TearDownTestCase() {
35+ std::cout << "MseLossGrad Proto Test TearDown" << std::endl;
36+ }
37+};
38+ 
39+TEST_F(MseLossGradTest, mse_loss_grad_infer_shape_test1) {
40+ ge::op::MseLossGrad op;
41+ 
42+ ge::DataType dtype = ge::DT_FLOAT;
43+ ge::Format format = ge::FORMAT_ND;
44+
45+ auto input_tensor = create_desc_with_ori({182,4}, dtype, format, {182,4}, format);
46+
47+ op.UpdateInputDesc("predict", input_tensor);
48+ op.UpdateInputDesc("label", input_tensor);
49+ op.UpdateInputDesc("dout", input_tensor);
50+ 
51+ op.SetAttr("reduction", "mean");
52+ Runtime2TestParam param{{"reduction"}};
53+ EXPECT_EQ(InferShapeTest(op, param), ge::GRAPH_SUCCESS);
54+ 
55+ auto output_desc = op.GetOutputDescByName("y");
56+ 
57+ EXPECT_EQ(output_desc.GetShape().GetDimNum(), 2);
58+}
59+ 
60+TEST_F(MseLossGradTest, mse_loss_grad_infer_data_type)
61+{
62+ ASSERT_NE(gert::OpImplRegistry::GetInstance().GetOpImpl("MseLossGrad"), nullptr);
63+ auto data_type_func = gert::OpImplRegistry::GetInstance().GetOpImpl("MseLossGrad")->infer_datatype;
64+ ASSERT_NE(data_type_func, nullptr);
65+ 
66+ ge::DataType input_x = ge::DT_FLOAT;
67+ ge::DataType y_datatype = ge::DT_FLOAT;
68+ auto context_holder = gert::InferDataTypeContextFaker()
69+ .IrInputNum(3)
70+ .NodeIoNum(3, 1)
71+ .NodeInputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
72+ .NodeInputTd(1, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
73+ .NodeInputTd(2, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
74+ .NodeOutputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
75+ .NodeAttrs({{"reduction", Ops::NN::AnyValue::CreateFrom<std::string>("mean")}})
76+ .InputDataTypes({&input_x, &input_x, &input_x})
77+ .OutputDataTypes({&y_datatype})
78+ .Build();
79+ auto context = context_holder.GetContext<gert::InferDataTypeContext>();
80+ EXPECT_EQ(data_type_func(context), ge::GRAPH_SUCCESS);
81+ ASSERT_NE(context, nullptr);
82+ 
83+ EXPECT_EQ(context->GetInputDataType(0), input_x);
84+ EXPECT_EQ(context->GetOutputDataType(0), y_datatype);
85+}
Aloss/mse_loss_grad/tests/ut/op_host/test_mse_loss_grad_tiling.cpp+399-0
@@ -0,0 +1,399 @@
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 <iostream>
12+#include <fstream>
13+#include <vector>
14+#include <gtest/gtest.h>
15+ 
16+#include "log/log.h"
17+#include "kernel_run_context_facker.h"
18+#include "test_cube_util.h"
19+#include "exe_graph/runtime/storage_format.h"
20+#include "exe_graph/runtime/storage_shape.h"
21+#include "platform/platform_infos_def.h"
22+#include "ut_op_util.h"
23+#include "../../../op_host/arch35/mse_loss_grad_tiling_arch35.h"
24+ 
25+using namespace ut_util;
26+using namespace std;
27+using namespace ge;
28+ 
29+class MseLossGradTiling : public testing::Test {
30+ protected:
31+ static void SetUpTestCase() {
32+ std::cout << "MseLossGradTiling SetUp" << std::endl;
33+ }
34+ 
35+ static void TearDownTestCase() {
36+ std::cout << "MseLossGradTiling TearDown" << std::endl;
37+ }
38+};
39+ 
40+// TEST_F(MseLossGradTiling, mse_loss_grad_testcase_001)
41+// {
42+// gert::StorageShape input_shape = {{182,4}, {182,4}};
43+// gert::StorageShape output_shape = {{182,4}, {182,4}};
44+ 
45+// std::map<std::string, std::string> soc_infos;
46+// std::map<std::string, std::string> aicore_spec;
47+// std::map<std::string, std::string> intrinsics;
48+// std::map<std::string, std::string> soc_version_infos = {{"Short_SoC_version", "Ascend910_95"}};
49+// std::string compile_info_string = R"({
50+// "hardware_info": {
51+// "BT_SIZE": 0, "load3d_constraints": "1",
52+// "Intrinsic_fix_pipe_l0c2out": false, "Intrinsic_data_move_l12ub": true,
53+// "Intrinsic_data_move_l0c2ub": true, "Intrinsic_data_move_out2l1_nd2nz": false,
54+// "UB_SIZE": 245760, "L2_SIZE": 33554432, "L1_SIZE": 524288,
55+// "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072, "CORE_NUM": 64
56+// }
57+// })";
58+// std::string op_type("MseLossGrad");
59+ 
60+// GetPlatFormInfos(compile_info_string.c_str(), soc_infos, aicore_spec, intrinsics);
61+ 
62+// fe::PlatFormInfos platform_info;
63+// platform_info.Init();
64+ 
65+// struct MseLossGradCompileInfo {
66+// uint64_t coreNum = 0;
67+// uint64_t ubSize = 0;
68+// };
69+// MseLossGradCompileInfo compile_info;
70+ 
71+// auto tiling_func = gert::OpImplRegistry::GetInstance().GetOpImpl("MseLossGrad")->tiling;
72+// auto tiling_parse_func = gert::OpImplRegistry::GetInstance().GetOpImpl("MseLossGrad")->tiling_parse;
73+// auto gen_simplifiedkey_func = gert::OpImplRegistry::GetInstance().GetOpImpl("MseLossGrad")->gen_simplifiedkey;
74+ 
75+// auto kernel_holder =
76+// gert::KernelRunContextFaker()
77+// .KernelIONum(3, 1)
78+// .Inputs({const_cast<char*>(compile_info_string.c_str()), reinterpret_cast<void*>(&platform_info)})
79+// .Outputs({&compile_info})
80+// .Build();
81+ 
82+// ASSERT_TRUE(kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->Init());
83+// kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", soc_infos);
84+// kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicore_spec);
85+// kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
86+// kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap",
87+// intrinsics);
88+// kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("version",
89+// soc_version_infos);
90+// ASSERT_EQ(tiling_parse_func(kernel_holder.GetContext<gert::KernelContext>()), ge::GRAPH_SUCCESS);
91+ 
92+// auto param = gert::TilingData::CreateCap(4096);
93+// auto workspace_size_holder = gert::ContinuousVector::Create<size_t>(4096);
94+// auto ws_size = reinterpret_cast<gert::ContinuousVector*>(workspace_size_holder.get());
95+// ASSERT_NE(param, nullptr);
96+ 
97+// auto holder = gert::TilingContextFaker()
98+// .SetOpType(op_type)
99+// .NodeIoNum(3, 1)
100+// .IrInstanceNum({1,1,1})
101+// .InputShapes({&input_shape, &input_shape, &input_shape})
102+// .OutputShapes({&output_shape})
103+// .CompileInfo(&compile_info)
104+// .PlatformInfo(reinterpret_cast<char*>(&platform_info))
105+// .NodeInputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
106+// .NodeInputTd(1, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
107+// .NodeInputTd(2, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
108+// .NodeOutputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
109+// .NodeAttrs({{"reduction", Ops::NN::AnyValue::CreateFrom<std::string>("mean")}})
110+// .TilingData(param.get())
111+// .Workspace(ws_size)
112+// .Build();
113+ 
114+ 
115+// gert::TilingContext* tiling_context = holder.GetContext<gert::TilingContext>();
116+// ASSERT_NE(tiling_context->GetPlatformInfo(), nullptr);
117+ 
118+// tiling_context->GetPlatformInfo()->SetPlatformRes("SoCInfo", soc_infos);
119+// tiling_context->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicore_spec);
120+// tiling_context->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
121+// tiling_context->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics);
122+ 
123+// EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_SUCCESS);
124+// auto tiling_key = tiling_context->GetTilingKey();
125+// ASSERT_EQ(tiling_key, 7);
126+// auto block_dim = tiling_context->GetBlockDim();
127+// ASSERT_EQ(block_dim, 4);
128+// }
129+ 
130+// TEST_F(MseLossGradTiling, mse_loss_grad_testcase_002)
131+// {
132+// gert::StorageShape input_shape = {{182,4}, {182,4}};
133+// gert::StorageShape output_shape = {{182,4}, {182,4}};
134+ 
135+// std::map<std::string, std::string> soc_infos;
136+// std::map<std::string, std::string> aicore_spec;
137+// std::map<std::string, std::string> intrinsics;
138+// std::map<std::string, std::string> soc_version_infos = {{"Short_SoC_version", "Ascend910_95"}};
139+// std::string compile_info_string = R"({
140+// "hardware_info": {
141+// "BT_SIZE": 0, "load3d_constraints": "1",
142+// "Intrinsic_fix_pipe_l0c2out": false, "Intrinsic_data_move_l12ub": true,
143+// "Intrinsic_data_move_l0c2ub": true, "Intrinsic_data_move_out2l1_nd2nz": false,
144+// "UB_SIZE": 245760, "L2_SIZE": 33554432, "L1_SIZE": 524288,
145+// "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072, "CORE_NUM": 64
146+// }
147+// })";
148+// std::string op_type("MseLossGrad");
149+ 
150+// GetPlatFormInfos(compile_info_string.c_str(), soc_infos, aicore_spec, intrinsics);
151+ 
152+// fe::PlatFormInfos platform_info;
153+// platform_info.Init();
154+ 
155+// struct MseLossGradCompileInfo {
156+// uint64_t coreNum = 0;
157+// uint64_t ubSize = 0;
158+// };
159+// MseLossGradCompileInfo compile_info;
160+ 
161+// auto tiling_func = gert::OpImplRegistry::GetInstance().GetOpImpl("MseLossGrad")->tiling;
162+// auto tiling_parse_func = gert::OpImplRegistry::GetInstance().GetOpImpl("MseLossGrad")->tiling_parse;
163+// auto gen_simplifiedkey_func = gert::OpImplRegistry::GetInstance().GetOpImpl("MseLossGrad")->gen_simplifiedkey;
164+ 
165+// auto kernel_holder =
166+// gert::KernelRunContextFaker()
167+// .KernelIONum(3, 1)
168+// .Inputs({const_cast<char*>(compile_info_string.c_str()), reinterpret_cast<void*>(&platform_info)})
169+// .Outputs({&compile_info})
170+// .Build();
171+ 
172+// ASSERT_TRUE(kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->Init());
173+// kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", soc_infos);
174+// kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicore_spec);
175+// kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
176+// kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap",
177+// intrinsics);
178+// kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("version",
179+// soc_version_infos);
180+// ASSERT_EQ(tiling_parse_func(kernel_holder.GetContext<gert::KernelContext>()), ge::GRAPH_SUCCESS);
181+ 
182+// auto param = gert::TilingData::CreateCap(4096);
183+// auto workspace_size_holder = gert::ContinuousVector::Create<size_t>(4096);
184+// auto ws_size = reinterpret_cast<gert::ContinuousVector*>(workspace_size_holder.get());
185+// ASSERT_NE(param, nullptr);
186+ 
187+// auto holder = gert::TilingContextFaker()
188+// .SetOpType(op_type)
189+// .NodeIoNum(3, 1)
190+// .IrInstanceNum({1,1,1})
191+// .InputShapes({&input_shape, &input_shape, &input_shape})
192+// .OutputShapes({&output_shape})
193+// .CompileInfo(&compile_info)
194+// .PlatformInfo(reinterpret_cast<char*>(&platform_info))
195+// .NodeInputTd(0, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
196+// .NodeInputTd(1, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
197+// .NodeInputTd(2, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
198+// .NodeOutputTd(0, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
199+// .NodeAttrs({{"reduction", Ops::NN::AnyValue::CreateFrom<std::string>("mean")}})
200+// .TilingData(param.get())
201+// .Workspace(ws_size)
202+// .Build();
203+ 
204+ 
205+// gert::TilingContext* tiling_context = holder.GetContext<gert::TilingContext>();
206+// ASSERT_NE(tiling_context->GetPlatformInfo(), nullptr);
207+ 
208+// tiling_context->GetPlatformInfo()->SetPlatformRes("SoCInfo", soc_infos);
209+// tiling_context->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicore_spec);
210+// tiling_context->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
211+// tiling_context->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics);
212+ 
213+// EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_SUCCESS);
214+// auto tiling_key = tiling_context->GetTilingKey();
215+// ASSERT_EQ(tiling_key, 7);
216+// auto block_dim = tiling_context->GetBlockDim();
217+// ASSERT_EQ(block_dim, 5);
218+// }
219+ 
220+// TEST_F(MseLossGradTiling, mse_loss_grad_testcase_003)
221+// {
222+// gert::StorageShape input_shape = {{182,4}, {182,4}};
223+// gert::StorageShape output_shape = {{182,4}, {182,4}};
224+ 
225+// std::map<std::string, std::string> soc_infos;
226+// std::map<std::string, std::string> aicore_spec;
227+// std::map<std::string, std::string> intrinsics;
228+// std::map<std::string, std::string> soc_version_infos = {{"Short_SoC_version", "Ascend910_95"}};
229+// std::string compile_info_string = R"({
230+// "hardware_info": {
231+// "BT_SIZE": 0, "load3d_constraints": "1",
232+// "Intrinsic_fix_pipe_l0c2out": false, "Intrinsic_data_move_l12ub": true,
233+// "Intrinsic_data_move_l0c2ub": true, "Intrinsic_data_move_out2l1_nd2nz": false,
234+// "UB_SIZE": 245760, "L2_SIZE": 33554432, "L1_SIZE": 524288,
235+// "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072, "CORE_NUM": 64
236+// }
237+// })";
238+// std::string op_type("MseLossGrad");
239+ 
240+// GetPlatFormInfos(compile_info_string.c_str(), soc_infos, aicore_spec, intrinsics);
241+ 
242+// fe::PlatFormInfos platform_info;
243+// platform_info.Init();
244+ 
245+// struct MseLossGradCompileInfo {
246+// uint64_t coreNum = 0;
247+// uint64_t ubSize = 0;
248+// };
249+// MseLossGradCompileInfo compile_info;
250+ 
251+// auto tiling_func = gert::OpImplRegistry::GetInstance().GetOpImpl("MseLossGrad")->tiling;
252+// auto tiling_parse_func = gert::OpImplRegistry::GetInstance().GetOpImpl("MseLossGrad")->tiling_parse;
253+// auto gen_simplifiedkey_func = gert::OpImplRegistry::GetInstance().GetOpImpl("MseLossGrad")->gen_simplifiedkey;
254+ 
255+// auto kernel_holder =
256+// gert::KernelRunContextFaker()
257+// .KernelIONum(3, 1)
258+// .Inputs({const_cast<char*>(compile_info_string.c_str()), reinterpret_cast<void*>(&platform_info)})
259+// .Outputs({&compile_info})
260+// .Build();
261+ 
262+// ASSERT_TRUE(kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->Init());
263+// kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", soc_infos);
264+// kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicore_spec);
265+// kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
266+// kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap",
267+// intrinsics);
268+// kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("version",
269+// soc_version_infos);
270+// ASSERT_EQ(tiling_parse_func(kernel_holder.GetContext<gert::KernelContext>()), ge::GRAPH_SUCCESS);
271+ 
272+// auto param = gert::TilingData::CreateCap(4096);
273+// auto workspace_size_holder = gert::ContinuousVector::Create<size_t>(4096);
274+// auto ws_size = reinterpret_cast<gert::ContinuousVector*>(workspace_size_holder.get());
275+// ASSERT_NE(param, nullptr);
276+ 
277+// auto holder = gert::TilingContextFaker()
278+// .SetOpType(op_type)
279+// .NodeIoNum(3, 1)
280+// .IrInstanceNum({1,1,1})
281+// .InputShapes({&input_shape, &input_shape, &input_shape})
282+// .OutputShapes({&output_shape})
283+// .CompileInfo(&compile_info)
284+// .PlatformInfo(reinterpret_cast<char*>(&platform_info))
285+// .NodeInputTd(0, ge::DT_BF16, ge::FORMAT_ND, ge::FORMAT_ND)
286+// .NodeInputTd(1, ge::DT_BF16, ge::FORMAT_ND, ge::FORMAT_ND)
287+// .NodeInputTd(2, ge::DT_BF16, ge::FORMAT_ND, ge::FORMAT_ND)
288+// .NodeOutputTd(0, ge::DT_BF16, ge::FORMAT_ND, ge::FORMAT_ND)
289+// .NodeAttrs({{"reduction", Ops::NN::AnyValue::CreateFrom<std::string>("mean")}})
290+// .TilingData(param.get())
291+// .Workspace(ws_size)
292+// .Build();
293+ 
294+ 
295+// gert::TilingContext* tiling_context = holder.GetContext<gert::TilingContext>();
296+// ASSERT_NE(tiling_context->GetPlatformInfo(), nullptr);
297+ 
298+// tiling_context->GetPlatformInfo()->SetPlatformRes("SoCInfo", soc_infos);
299+// tiling_context->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicore_spec);
300+// tiling_context->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
301+// tiling_context->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics);
302+ 
303+// EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_SUCCESS);
304+// auto tiling_key = tiling_context->GetTilingKey();
305+// ASSERT_EQ(tiling_key, 7);
306+// auto block_dim = tiling_context->GetBlockDim();
307+// ASSERT_EQ(block_dim, 5);
308+// }
309+ 
310+// TEST_F(MseLossGradTiling, mse_loss_grad_testcase_004)
311+// {
312+// gert::StorageShape input_shape = {{182,4}, {182,4}};
313+// gert::StorageShape dout_shape = {{1}, {1}};
314+// gert::StorageShape output_shape = {{182,4}, {182,4}};
315+ 
316+// std::map<std::string, std::string> soc_infos;
317+// std::map<std::string, std::string> aicore_spec;
318+// std::map<std::string, std::string> intrinsics;
319+// std::map<std::string, std::string> soc_version_infos = {{"Short_SoC_version", "Ascend910_95"}};
320+// std::string compile_info_string = R"({
321+// "hardware_info": {
322+// "BT_SIZE": 0, "load3d_constraints": "1",
323+// "Intrinsic_fix_pipe_l0c2out": false, "Intrinsic_data_move_l12ub": true,
324+// "Intrinsic_data_move_l0c2ub": true, "Intrinsic_data_move_out2l1_nd2nz": false,
325+// "UB_SIZE": 245760, "L2_SIZE": 33554432, "L1_SIZE": 524288,
326+// "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072, "CORE_NUM": 64
327+// }
328+// })";
329+// std::string op_type("MseLossGrad");
330+ 
331+// GetPlatFormInfos(compile_info_string.c_str(), soc_infos, aicore_spec, intrinsics);
332+ 
333+// fe::PlatFormInfos platform_info;
334+// platform_info.Init();
335+ 
336+// struct MseLossGradCompileInfo {
337+// uint64_t coreNum = 0;
338+// uint64_t ubSize = 0;
339+// };
340+// MseLossGradCompileInfo compile_info;
341+ 
342+// auto tiling_func = gert::OpImplRegistry::GetInstance().GetOpImpl("MseLossGrad")->tiling;
343+// auto tiling_parse_func = gert::OpImplRegistry::GetInstance().GetOpImpl("MseLossGrad")->tiling_parse;
344+// auto gen_simplifiedkey_func = gert::OpImplRegistry::GetInstance().GetOpImpl("MseLossGrad")->gen_simplifiedkey;
345+ 
346+// auto kernel_holder =
347+// gert::KernelRunContextFaker()
348+// .KernelIONum(3, 1)
349+// .Inputs({const_cast<char*>(compile_info_string.c_str()), reinterpret_cast<void*>(&platform_info)})
350+// .Outputs({&compile_info})
351+// .Build();
352+ 
353+// ASSERT_TRUE(kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->Init());
354+// kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", soc_infos);
355+// kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicore_spec);
356+// kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
357+// kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap",
358+// intrinsics);
359+// kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("version",
360+// soc_version_infos);
361+// ASSERT_EQ(tiling_parse_func(kernel_holder.GetContext<gert::KernelContext>()), ge::GRAPH_SUCCESS);
362+ 
363+// auto param = gert::TilingData::CreateCap(4096);
364+// auto workspace_size_holder = gert::ContinuousVector::Create<size_t>(4096);
365+// auto ws_size = reinterpret_cast<gert::ContinuousVector*>(workspace_size_holder.get());
366+// ASSERT_NE(param, nullptr);
367+ 
368+// auto holder = gert::TilingContextFaker()
369+// .SetOpType(op_type)
370+// .NodeIoNum(3, 1)
371+// .IrInstanceNum({1,1,1})
372+// .InputShapes({&input_shape, &input_shape, &dout_shape})
373+// .OutputShapes({&output_shape})
374+// .CompileInfo(&compile_info)
375+// .PlatformInfo(reinterpret_cast<char*>(&platform_info))
376+// .NodeInputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
377+// .NodeInputTd(1, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
378+// .NodeInputTd(2, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
379+// .NodeOutputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
380+// .NodeAttrs({{"reduction", Ops::NN::AnyValue::CreateFrom<std::string>("mean")}})
381+// .TilingData(param.get())
382+// .Workspace(ws_size)
383+// .Build();
384+ 
385+ 
386+// gert::TilingContext* tiling_context = holder.GetContext<gert::TilingContext>();
387+// ASSERT_NE(tiling_context->GetPlatformInfo(), nullptr);
388+ 
389+// tiling_context->GetPlatformInfo()->SetPlatformRes("SoCInfo", soc_infos);
390+// tiling_context->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicore_spec);
391+// tiling_context->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
392+// tiling_context->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics);
393+ 
394+// EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_SUCCESS);
395+// auto tiling_key = tiling_context->GetTilingKey();
396+// ASSERT_EQ(tiling_key, 65543);
397+// auto block_dim = tiling_context->GetBlockDim();
398+// ASSERT_EQ(block_dim, 3);
399+// }
Rpooling/max_pool_v3/op_host/CMakeLists.txt→pooling/max_pool3_d/CMakeLists.txt+6-4
@@ -4,9 +4,11 @@
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-message(STATUS "=== Debug: start ops.pooling.max_pool_v3.op_host.CMakeLists.txt ")10+# 设置算子定义时支持的芯片类型
11- 11+set(SUPPORT_COMPUTE_UNIT "ascend910_95")
12-add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE max_pool_v3 ACLNNTYPE aclnn_exclude)12+# 设置每种芯片类型对应的tiling文件目录，即采用op_host目录下哪个文件夹下的tiling文件编译
13+set(SUPPORT_TILING_DIR "arch35")
14+add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} OPTYPE max_pool3_d ACLNNTYPE aclnn_exclude DISABLE_IN_OPP TRUE DEPENDENCIES pool_3d)
Apooling/max_pool3_d/op_graph/max_pool3_d_proto.h+106-0
@@ -0,0 +1,106 @@
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+#ifndef OPS_BUILT_IN_OP_PROTO_INC_NN_POOLING_OPS_H_
12+#define OPS_BUILT_IN_OP_PROTO_INC_NN_POOLING_OPS_H_
13+ 
14+#include "graph/operator_reg.h"
15+#include "graph/operator.h"
16+ 
17+namespace ge {
18+/**
19+* @brief Performs max pooling on the input . \n
20+ 
21+* @par Inputs:
22+* One input:
23+* x: A 5D tensor. Supported type:float16, float32. Additional support for bfloat16 in Ascend 910_95 AI Processor.
24+* The double type is reserved but currently unsupported. Support format: NHDWC, NCDHW. Additional support for ND in Ascend 910_95 AI Processor.
25+ 
26+* @par Attributes:
27+* @li ksize: A required list of int8, int16, int32, or int64 values,
28+* specifying the size of the window for each dimension of the input tensor.
29+* @li strides: A required list of int8, int16, int32, or int64 values,
30+* specifying the stride of the sliding window for each dimension of
31+* the input tensor. No default value.
32+* @li padding: An required string, specifying the type of
33+* the padding algorithm to use. It support "SAME", "VALID" or "CALCULATED".
34+* @li pads: An optional list of int8, int16, int32, or int64 values.
35+* It must have 6 elements, specifying the front, backend, top, bottom, left, and right padding for input tensor.
36+* The pads should be greater than or equal to 0 and smaller than the corresponding kernel size.
37+* It only takes effect when padding is "CALCULATED". Default value is {0,0,0,0,0,0}.
38+* @li dilation: Dilation of kernel. default value is {1,1,1,1,1}.
39+* @li ceil_mode: Use the floor or ceil function to calculate output depth, height and width.
40+* It support 0(floor) or 1(ceil). Default value is 0.
41+* It can be set to 1 only when padding is "CALCULATED".
42+* @li data_format: An optional string, specify the data format of the input and
43+* output data. It support "NDHWC"(default), "NCDHW"". \n
44+ 
45+* @par Outputs:
46+* y: A 5D tensor. Has the same type and format as input "x" . \n
47+ 
48+* @attention Constraints:
49+* @li "ksize" is a list that has length 1, 3, or 5. The ksize of the H and W dimensions should be greater than 0.
50+* The ksize of the N and C dimensions should be 1. e.g. For "data_format" is "NCDHW", ksize[0] = 1 and ksize[1] = 1.
51+* For "data_format" is "NDHWC", ksize[0] = 1 and ksize[4] = 1. \n
52+* For Atlas Training Series Product, Atlas A2 Training Series Product/Atlas 800I A2 Inference Product,
53+* Atlas A3 Training Series Product: The produce of the ksize in D, H and W dimensions
54+* should be less than or equal to 255. e.g. For "data_format" is "NCDHW", ksize[2] * ksize[3] * ksize[4] <= 255. \n
55+* @li "strides" is a list that has length 1, 3, or 5. The stride of the N and C dimensions should be 1. \n
56+* For Atlas Training Series Product, Atlas A2 Training Series Product/Atlas 800I A2 Inference Product,
57+* Atlas A3 Training Series Product: The stride of the D, H and W dimensions should be greater than 0 and
58+* smaller than 64. \n
59+* The stride of the D, H and W dimensions should be greater than 0.
60+* @li "data_format" only support "NCDHW" and "NDHWC". \n
61+* @li The ouput "y" shape at the N and C dimensions should be equal with input "x" shape at same dimensions. The output
62+* shape at the D, H and W dimensions is calculated by below formula: \n
63+* @code{.c}
64+ when "padding" is "SAME":
65+ out_depth = (in_depth + stride_d - 1) / stride_d
66+ out_height = (in_height + stride_h - 1) / stride_h
67+ out_width = (in_width + stride_w - 1) / stride_w
68+ when "padding" is "VALID":
69+ out_depth = (in_depth + stride_d - ((ksize_d - 1) * dilation_d + 1)) / stride_d
70+ out_height = (in_height + stride_h - ((ksize_h - 1) * dilation_h + 1)) / stride_h
71+ out_width = (in_width + stride_w - ((ksize_w - 1) * dilation_w + 1)) / stride_w
72+ when "padding" is "CALCULATED":
73+ if "ceil_mode" is 0:
74+ out_depth = (in_depth + pad_front + pad_backend - ((ksize_d - 1) * dilation_d + 1)) / stride_d + 1
75+ out_height = (in_height + pad_top + pad_bottom - ((ksize_h - 1) * dilation_h + 1)) / stride_h + 1
76+ out_width = (in_width + pad_left + pad_right - ((ksize_w - 1) * dilation_w + 1)) / stride_w + 1
77+ else :
78+ out_depth = (in_depth + pad_front + pad_backend - ((ksize_d - 1) * dilation_d + 1) + stride_d - 1) / stride_d + 1
79+ out_height = (in_height + pad_top + pad_bottom - ((ksize_h - 1) * dilation_h + 1) + stride_h - 1) / stride_h + 1
80+ out_width = (in_width + in_width + pad_right - ((ksize_w - 1) * dilation_w + 1) + stride_w - 1) / stride_w + 1
81+ if (out_depth - 1) * stride_d >= in_depth + pad_front :
82+ out_depth = out_depth - 1
83+ if (out_height - 1) * stride_h >= in_height + pad_top :
84+ out_height = out_height - 1
85+ if (out_width - 1) * stride_w >= in_width + in_width :
86+ out_width = out_width - 1
87+ It not support out_height < 0 or out_width < 0.
88+* @endcode
89+* @par Third-party framework compatibility
90+* Compatible with the TensorFlow operator MaxPool3D.
91+*/
92+REG_OP(MaxPool3D)
93+ .INPUT(x, TensorType({DT_FLOAT16, DT_FLOAT32, DT_DOUBLE, DT_BF16}))
94+ .OUTPUT(y, TensorType({DT_FLOAT16, DT_FLOAT32, DT_DOUBLE, DT_BF16}))
95+ .REQUIRED_ATTR(ksize, ListInt)
96+ .REQUIRED_ATTR(strides, ListInt)
97+ .REQUIRED_ATTR(padding, String)
98+ .ATTR(pads, ListInt, {0,0,0,0,0,0})
99+ .ATTR(dilation, ListInt, {1, 1, 1, 1, 1})
100+ .ATTR(ceil_mode, Int, 0)
101+ .ATTR(data_format, String, "NDHWC")
102+ .OP_END_FACTORY_REG(MaxPool3D)
103+ 
104+} // namespace ge
105+ 
106+#endif
Apooling/max_pool3_d/op_host/arch35/max_pool_3d_tiling.cpp+47-0
@@ -0,0 +1,47 @@
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_pool_3d_tiling.cpp
13+ * \brief
14+ */
15+ 
16+#include "pooling/pool_3d/op_host/arch35/pool_tiling_templates_registry.h"
17+#include "pooling/pool_3d/op_host/arch35/max_pool_3d_tiling_common.h"
18+#include "log/log.h"
19+ 
20+using namespace AscendC;
21+using optiling::PoolTilingRegistry;
22+namespace optiling {
23+ge::graphStatus Tiling4MaxPool3D(gert::TilingContext* context)
24+{
25+ OP_LOGD(context->GetNodeName(), "Tiling for max_pool_3d is running.");
26+ auto compileInfoPtr = context->GetCompileInfo<MaxPool3DCompileInfo>();
27+ OP_CHECK_NULL_WITH_CONTEXT(context, compileInfoPtr);
28+ return PoolTilingRegistry::GetInstance().DoTilingImpl(context);
29+}
30+ 
31+ge::graphStatus TilingPrepare4MaxPool3D(gert::TilingParseContext* context)
32+{
33+ OP_CHECK_IF(nullptr == context, OP_LOGE("MaxPool3D", "Context is null"), return ge::GRAPH_FAILED);
34+ auto compileInfoPtr = context->GetCompiledInfo<MaxPool3DCompileInfo>();
35+ OP_CHECK_NULL_WITH_CONTEXT(context, compileInfoPtr);
36+ 
37+ fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo();
38+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr);
39+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
40+ compileInfoPtr->coreNum = ascendcPlatform.GetCoreNumAiv();
41+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, compileInfoPtr->ubSize);
42+ return ge::GRAPH_SUCCESS;
43+}
44+ 
45+IMPL_OP_OPTILING(MaxPool3D).Tiling(Tiling4MaxPool3D).TilingParse<MaxPool3DCompileInfo>(TilingPrepare4MaxPool3D);
46+ 
47+} // namespace optiling
Apooling/max_pool3_d/op_host/config/ascend910_95/max_pool3_d_binary.json+221-0
@@ -0,0 +1,221 @@
1+{
2+ "op_type": "MaxPool3D",
3+ "op_list": [
4+ {
5+ "bin_filename": "MaxPool3D_612bebee9e1c9b3a7c2f9b32fd8b1820",
6+ "inputs": [
7+ {
8+ "name": "x",
9+ "index": 0,
10+ "dtype": "float16",
11+ "format": "ND",
12+ "paramType": "required",
13+ "shape": [
14+ -2
15+ ],
16+ "format_match_mode": "FormatAgnostic"
17+ }
18+ ],
19+ "outputs": [
20+ {
21+ "name": "y",
22+ "index": 0,
23+ "dtype": "float16",
24+ "format": "ND",
25+ "paramType": "required",
26+ "shape": [
27+ -2
28+ ],
29+ "format_match_mode": "FormatAgnostic"
30+ }
31+ ],
32+ "attrs": [
33+ {
34+ "name": "ksize",
35+ "dtype": "list_int",
36+ "value": null
37+ },
38+ {
39+ "name": "strides",
40+ "dtype": "list_int",
41+ "value": null
42+ },
43+ {
44+ "name": "padding",
45+ "dtype": "string",
46+ "value": "CALCULATED"
47+ },
48+ {
49+ "name": "pads",
50+ "dtype": "list_int",
51+ "value": null
52+ },
53+ {
54+ "name": "dilation",
55+ "dtype": "list_int",
56+ "value": [
57+ 1,
58+ 1,
59+ 1,
60+ 1,
61+ 1
62+ ]
63+ },
64+ {
65+ "name": "ceil_mode",
66+ "dtype": "int",
67+ "value": 0
68+ },
69+ {
70+ "name": "data_format",
71+ "dtype": "string",
72+ "value": "NCDHW"
73+ }
74+ ]
75+ },
76+ {
77+ "bin_filename": "MaxPool3D_5381b87740abaf5ce279a9d1323c9354",
78+ "inputs": [
79+ {
80+ "name": "x",
81+ "index": 0,
82+ "dtype": "float32",
83+ "format": "ND",
84+ "paramType": "required",
85+ "shape": [
86+ -2
87+ ],
88+ "format_match_mode": "FormatAgnostic"
89+ }
90+ ],
91+ "outputs": [
92+ {
93+ "name": "y",
94+ "index": 0,
95+ "dtype": "float32",
96+ "format": "ND",
97+ "paramType": "required",
98+ "shape": [
99+ -2
100+ ],
101+ "format_match_mode": "FormatAgnostic"
102+ }
103+ ],
104+ "attrs": [
105+ {
106+ "name": "ksize",
107+ "dtype": "list_int",
108+ "value": null
109+ },
110+ {
111+ "name": "strides",
112+ "dtype": "list_int",
113+ "value": null
114+ },
115+ {
116+ "name": "padding",
117+ "dtype": "string",
118+ "value": "CALCULATED"
119+ },
120+ {
121+ "name": "pads",
122+ "dtype": "list_int",
123+ "value": null
124+ },
125+ {
126+ "name": "dilation",
127+ "dtype": "list_int",
128+ "value": [
129+ 1,
130+ 1,
131+ 1,
132+ 1,
133+ 1
134+ ]
135+ },
136+ {
137+ "name": "ceil_mode",
138+ "dtype": "int",
139+ "value": 0
140+ },
141+ {
142+ "name": "data_format",
143+ "dtype": "string",
144+ "value": "NDHWC"
145+ }
146+ ]
147+ },
148+ {
149+ "bin_filename": "MaxPool3D_c9eaa18d4decb5251914c2d6442198c4",
150+ "inputs": [
151+ {
152+ "name": "x",
153+ "index": 0,
154+ "dtype": "bfloat16",
155+ "format": "ND",
156+ "paramType": "required",
157+ "shape": [
158+ -2
159+ ],
160+ "format_match_mode": "FormatAgnostic"
161+ }
162+ ],
163+ "outputs": [
164+ {
165+ "name": "y",
166+ "index": 0,
167+ "dtype": "bfloat16",
168+ "format": "ND",
169+ "paramType": "required",
170+ "shape": [
171+ -2
172+ ],
173+ "format_match_mode": "FormatAgnostic"
174+ }
175+ ],
176+ "attrs": [
177+ {
178+ "name": "ksize",
179+ "dtype": "list_int",
180+ "value": null
181+ },
182+ {
183+ "name": "strides",
184+ "dtype": "list_int",
185+ "value": null
186+ },
187+ {
188+ "name": "padding",
189+ "dtype": "string",
190+ "value": "CALCULATED"
191+ },
192+ {
193+ "name": "pads",
194+ "dtype": "list_int",
195+ "value": null
196+ },
197+ {
198+ "name": "dilation",
199+ "dtype": "list_int",
200+ "value": [
201+ 1,
202+ 1,
203+ 1,
204+ 1,
205+ 1
206+ ]
207+ },
208+ {
209+ "name": "ceil_mode",
210+ "dtype": "int",
211+ "value": 1
212+ },
213+ {
214+ "name": "data_format",
215+ "dtype": "string",
216+ "value": "NCDHW"
217+ }
218+ ]
219+ }
220+ ]
221+}
Apooling/max_pool3_d/op_host/config/ascend910_95/max_pool3_d_simplified_key.ini+2-0
@@ -0,0 +1,2 @@
1+[MaxPool3D]
2+default=0
Apooling/max_pool3_d/op_host/max_pool3_d_def.cpp+54-0
@@ -0,0 +1,54 @@
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 <cstdint>
12+#include "register/op_def_registry.h"
13+ 
14+namespace ops {
15+class MaxPool3D : public OpDef {
16+public:
17+ const std::vector<ge::DataType> maxPool3DXDataType = {ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16};
18+ const std::vector<ge::Format> maxPool3DXFormat = {ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND};
19+ explicit MaxPool3D(const char* name) : OpDef(name)
20+ {
21+ this->Input("x")
22+ .ParamType(REQUIRED)
23+ .DataType(maxPool3DXDataType)
24+ .Format(maxPool3DXFormat)
25+ .UnknownShapeFormat(maxPool3DXFormat)
26+ .AutoContiguous();
27+ this->Output("y")
28+ .ParamType(REQUIRED)
29+ .DataType(maxPool3DXDataType)
30+ .Format(maxPool3DXFormat)
31+ .UnknownShapeFormat(maxPool3DXFormat)
32+ .AutoContiguous();
33+ this->Attr("ksize").AttrType(REQUIRED).ListInt();
34+ this->Attr("strides").AttrType(REQUIRED).ListInt();
35+ this->Attr("padding").AttrType(REQUIRED).String();
36+ this->Attr("pads").AttrType(OPTIONAL).ListInt({0, 0, 0, 0, 0, 0});
37+ this->Attr("dilation").AttrType(OPTIONAL).ListInt({1, 1, 1, 1, 1});
38+ this->Attr("ceil_mode").AttrType(OPTIONAL).Int(0);
39+ this->Attr("data_format").AttrType(OPTIONAL).String("NDHWC");
40+ OpAICoreConfig aiCoreConfig;
41+ aiCoreConfig.DynamicCompileStaticFlag(true)
42+ .DynamicFormatFlag(false)
43+ .DynamicRankSupportFlag(true)
44+ .DynamicShapeSupportFlag(true)
45+ .NeedCheckSupportFlag(false)
46+ .PrecisionReduceFlag(true)
47+ .ExtendCfgInfo("opFile.value", "max_pool3_d_apt");
48+ this->AICore().AddConfig("ascend910_95", aiCoreConfig);
49+ this->AICore().AddConfig("mc62cm12a", aiCoreConfig);
50+ }
51+};
52+ 
53+OP_ADD(MaxPool3D);
54+} // namespace ops
Apooling/max_pool3_d/op_host/max_pool3_d_infershape.cpp+277-0
@@ -0,0 +1,277 @@
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_pool_3d_infershape.cpp
13+ * \brief
14+ */
15+#include "register/op_impl_registry.h"
16+#include "log/log.h"
17+#include "util/shape_util.h"
18+ 
19+using namespace ge;
20+namespace ops {
21+constexpr size_t INDEX_KSIZE = 0;
22+constexpr size_t INDEX_STRIDES = 1;
23+constexpr size_t INDEX_PADDING = 2;
24+constexpr size_t INDEX_PADS = 3;
25+constexpr size_t INDEX_DILATION = 4;
26+constexpr size_t INDEX_CEIL_MODE = 5;
27+constexpr size_t INDEX_DATA_FORMAT = 6;
28+constexpr size_t SHAPE_SIZE = 5;
29+constexpr size_t DIM_SIZE6 = 6;
30+constexpr size_t PAD_SIZE = 6;
31+constexpr size_t PAD_FRONT = 0;
32+constexpr size_t PAD_BACK = 1;
33+constexpr size_t PAD_TOP = 2;
34+constexpr size_t PAD_BOTTOM = 3;
35+constexpr size_t PAD_LEFT = 4;
36+constexpr size_t PAD_RIGHT = 5;
37+ 
38+typedef ge::graphStatus (*InferShapePaddingFunc)(
39+ gert::InferShapeContext*, size_t, size_t, size_t, const gert::RuntimeAttrs*);
40+ 
41+static std::string Int64ToString(const int64_t* data, size_t size)
42+{
43+ std::string r = "[";
44+ for (size_t i = 0; i < size; i++) {
45+ r = r + std::to_string(data[i]) + " ";
46+ }
47+ r = r + "]";
48+ return r;
49+}
50+ 
51+static int64_t SameUpdateDim(const int64_t ksize, const int64_t strides, int64_t dim_size)
52+{
53+ return (strides == 0) ? (dim_size - ksize + 1) : ((dim_size - ksize + strides) / strides);
54+}
55+ 
56+static void CalculateUpdateDim(
57+ const int64_t ksize, const int64_t stride, const int64_t dilation, const int32_t ceil_mode, int64_t& dim_size)
58+{
59+ if (ceil_mode) {
60+ dim_size = (stride == 0) ? (dim_size - ksize + 1) :
61+ (dim_size - dilation * (ksize - 1) - 1 + stride + stride - 1) / stride;
62+ } else {
63+ dim_size = (stride == 0) ? (dim_size - ksize + 1) : (dim_size - dilation * (ksize - 1) - 1 + stride) / stride;
64+ }
65+}
66+ 
67+static ge::graphStatus InferShapePaddingCalculated(
68+ gert::InferShapeContext* context, size_t d_dim, size_t h_dim, size_t w_dim, const gert::RuntimeAttrs* attrs)
69+{
70+ auto ksize = attrs->GetAttrPointer<gert::ContinuousVector>(INDEX_KSIZE);
71+ OP_CHECK_NULL_WITH_CONTEXT(context, ksize);
72+ OP_CHECK_IF(
73+ ksize->GetSize() != SHAPE_SIZE,
74+ OP_LOGE(context->GetNodeName(), "Length of ksize %zu must be 5!", ksize->GetSize()), return GRAPH_FAILED);
75+ auto ksize_data = reinterpret_cast<const int64_t*>(ksize->GetData());
76+ auto pads = attrs->GetAttrPointer<gert::ContinuousVector>(INDEX_PADS);
77+ OP_CHECK_NULL_WITH_CONTEXT(context, pads);
78+ OP_CHECK_IF(
79+ pads->GetSize() != PAD_SIZE, OP_LOGE(context->GetNodeName(), "Length of pads %zu must be 6!", pads->GetSize()),
80+ return GRAPH_FAILED);
81+ auto ceil_mode = attrs->GetAttrPointer<int32_t>(INDEX_CEIL_MODE);
82+ OP_CHECK_NULL_WITH_CONTEXT(context, ceil_mode);
83+ auto strides = attrs->GetAttrPointer<gert::ContinuousVector>(INDEX_STRIDES);
84+ OP_CHECK_NULL_WITH_CONTEXT(context, strides);
85+ OP_CHECK_IF(
86+ strides->GetSize() != SHAPE_SIZE,
87+ OP_LOGE(context->GetNodeName(), "Length of strides %zu must be 5!", strides->GetSize()), return GRAPH_FAILED);
88+ auto strides_data = reinterpret_cast<const int64_t*>(strides->GetData());
89+ OP_CHECK_IF(
90+ strides_data[d_dim] <= 0 || strides_data[h_dim] <= 0 || strides_data[w_dim] <= 0,
91+ OP_LOGE(
92+ context->GetNodeName(), "%s h %lld and w %lld must be greater than 0.",
93+ Int64ToString(strides_data, strides->GetSize()).c_str(), strides_data[d_dim], strides_data[h_dim],
94+ strides_data[w_dim]),
95+ return GRAPH_FAILED);
96+ auto pads_data = reinterpret_cast<const int64_t*>(pads->GetData());
97+ 
98+ auto dilations = attrs->GetAttrPointer<gert::ContinuousVector>(INDEX_DILATION);
99+ OP_CHECK_NULL_WITH_CONTEXT(context, dilations);
100+ OP_CHECK_IF(
101+ ((dilations->GetSize() != SHAPE_SIZE) && (dilations->GetSize() != DIM_SIZE6)),
102+ OP_LOGE(context->GetNodeName(), "Length of dilation %zu must be 5!", dilations->GetSize()),
103+ return GRAPH_FAILED);
104+ auto dilations_data = reinterpret_cast<const int64_t*>(dilations->GetData());
105+ if (dilations->GetSize() == DIM_SIZE6) {
106+ OP_LOGW("InferShapePaddingCalculated", "The size of dilationList: 6 is deprecated, should be 5");
107+ }
108+ 
109+ auto in_shape = context->GetInputShape(0);
110+ OP_CHECK_NULL_WITH_CONTEXT(context, in_shape);
111+ auto out_shape = context->GetOutputShape(0);
112+ OP_CHECK_NULL_WITH_CONTEXT(context, out_shape);
113+ 
114+ *out_shape = *in_shape;
115+ int64_t dim_size = in_shape->GetDim(d_dim);
116+ int64_t out_dim_size = dim_size + pads_data[PAD_FRONT] + pads_data[PAD_BACK];
117+ CalculateUpdateDim(ksize_data[d_dim], strides_data[d_dim], dilations_data[d_dim], *ceil_mode, out_dim_size);
118+ if ((out_dim_size - 1) * strides_data[d_dim] >= dim_size + pads_data[PAD_FRONT]) {
119+ out_dim_size = out_dim_size - 1;
120+ }
121+ out_shape->SetDim(d_dim, out_dim_size);
122+ 
123+ dim_size = in_shape->GetDim(h_dim);
124+ out_dim_size = dim_size + pads_data[PAD_TOP] + pads_data[PAD_BOTTOM];
125+ CalculateUpdateDim(ksize_data[h_dim], strides_data[h_dim], dilations_data[h_dim], *ceil_mode, out_dim_size);
126+ if ((out_dim_size - 1) * strides_data[h_dim] >= dim_size + pads_data[PAD_TOP]) {
127+ out_dim_size = out_dim_size - 1;
128+ }
129+ out_shape->SetDim(h_dim, out_dim_size);
130+ 
131+ dim_size = in_shape->GetDim(w_dim);
132+ out_dim_size = dim_size + pads_data[PAD_LEFT] + pads_data[PAD_RIGHT];
133+ CalculateUpdateDim(ksize_data[w_dim], strides_data[w_dim], dilations_data[w_dim], *ceil_mode, out_dim_size);
134+ if ((out_dim_size - 1) * strides_data[w_dim] >= dim_size + pads_data[PAD_LEFT]) {
135+ out_dim_size = out_dim_size - 1;
136+ }
137+ out_shape->SetDim(w_dim, out_dim_size);
138+ 
139+ return ge::GRAPH_SUCCESS;
140+}
141+ 
142+static ge::graphStatus InferShapePaddingValid(
143+ gert::InferShapeContext* context, size_t d_dim, size_t h_dim, size_t w_dim, const gert::RuntimeAttrs* attrs)
144+{
145+ auto ksize = attrs->GetAttrPointer<gert::ContinuousVector>(INDEX_KSIZE);
146+ OP_CHECK_NULL_WITH_CONTEXT(context, ksize);
147+ OP_CHECK_IF(
148+ ksize->GetSize() != SHAPE_SIZE,
149+ OP_LOGE(context->GetNodeName(), "Length of ksize %zu must be 5!", ksize->GetSize()), return GRAPH_FAILED);
150+ auto ksize_data = reinterpret_cast<const int64_t*>(ksize->GetData());
151+ auto strides = attrs->GetAttrPointer<gert::ContinuousVector>(INDEX_STRIDES);
152+ OP_CHECK_NULL_WITH_CONTEXT(context, strides);
153+ OP_CHECK_IF(
154+ strides->GetSize() != SHAPE_SIZE,
155+ OP_LOGE(context->GetNodeName(), "Length of strides %zu must be 5!", strides->GetSize()), return GRAPH_FAILED);
156+ auto strides_data = reinterpret_cast<const int64_t*>(strides->GetData());
157+ OP_CHECK_IF(
158+ strides_data[d_dim] <= 0 || strides_data[h_dim] <= 0 || strides_data[w_dim] <= 0,
159+ OP_LOGE(
160+ context->GetNodeName(), "%s h %lld and w %lld must be greater than 0.",
161+ Int64ToString(strides_data, strides->GetSize()).c_str(), strides_data[d_dim], strides_data[h_dim],
162+ strides_data[w_dim]),
163+ return GRAPH_FAILED);
164+ 
165+ auto in_shape = context->GetInputShape(0);
166+ OP_CHECK_NULL_WITH_CONTEXT(context, in_shape);
167+ auto out_shape = context->GetOutputShape(0);
168+ OP_CHECK_NULL_WITH_CONTEXT(context, out_shape);
169+ 
170+ *out_shape = *in_shape;
171+ 
172+ int64_t dim_size = in_shape->GetDim(d_dim);
173+ out_shape->SetDim(d_dim, SameUpdateDim(ksize_data[d_dim], strides_data[d_dim], dim_size));
174+ dim_size = in_shape->GetDim(h_dim);
175+ out_shape->SetDim(h_dim, SameUpdateDim(ksize_data[h_dim], strides_data[h_dim], dim_size));
176+ dim_size = in_shape->GetDim(w_dim);
177+ out_shape->SetDim(w_dim, SameUpdateDim(ksize_data[w_dim], strides_data[w_dim], dim_size));
178+ 
179+ return ge::GRAPH_SUCCESS;
180+}
181+ 
182+static ge::graphStatus InferShapePaddingSame(
183+ gert::InferShapeContext* context, size_t d_dim, size_t h_dim, size_t w_dim, const gert::RuntimeAttrs* attrs)
184+{
185+ auto strides = attrs->GetAttrPointer<gert::ContinuousVector>(INDEX_STRIDES);
186+ OP_CHECK_NULL_WITH_CONTEXT(context, strides);
187+ OP_CHECK_IF(
188+ strides->GetSize() != SHAPE_SIZE,
189+ OP_LOGE(context->GetNodeName(), "Length of strides %zu must be 5!", strides->GetSize()), return GRAPH_FAILED);
190+ auto strides_data = reinterpret_cast<const int64_t*>(strides->GetData());
191+ OP_CHECK_IF(
192+ strides_data[d_dim] <= 0 || strides_data[h_dim] <= 0 || strides_data[w_dim] <= 0,
193+ OP_LOGE(
194+ context->GetNodeName(), "%s h %lld and w %lld must be greater than 0.",
195+ Int64ToString(strides_data, strides->GetSize()).c_str(), strides_data[d_dim], strides_data[h_dim],
196+ strides_data[w_dim]),
197+ return GRAPH_FAILED);
198+ 
199+ auto in_shape = context->GetInputShape(0);
200+ OP_CHECK_NULL_WITH_CONTEXT(context, in_shape);
201+ auto out_shape = context->GetOutputShape(0);
202+ OP_CHECK_NULL_WITH_CONTEXT(context, out_shape);
203+ 
204+ *out_shape = *in_shape;
205+ 
206+ int64_t dim_size = in_shape->GetDim(d_dim);
207+ out_shape->SetDim(d_dim, SameUpdateDim(1, strides_data[d_dim], dim_size));
208+ dim_size = in_shape->GetDim(h_dim);
209+ out_shape->SetDim(h_dim, SameUpdateDim(1, strides_data[h_dim], dim_size));
210+ dim_size = in_shape->GetDim(w_dim);
211+ out_shape->SetDim(w_dim, SameUpdateDim(1, strides_data[w_dim], dim_size));
212+ 
213+ return ge::GRAPH_SUCCESS;
214+}
215+ 
216+static const std::vector<std::pair<std::string, InferShapePaddingFunc>> kFuncMap = {
217+ {"CALCULATED", InferShapePaddingCalculated},
218+ {"SAME", InferShapePaddingSame},
219+ {"VALID", InferShapePaddingValid},
220+};
221+ 
222+static ge::graphStatus InferShape4MaxPool3D(gert::InferShapeContext* context)
223+{
224+ const gert::Shape* xShape = context->GetInputShape(0);
225+ OP_CHECK_NULL_WITH_CONTEXT(context, xShape);
226+ 
227+ gert::Shape* yShape = context->GetOutputShape(0);
228+ OP_CHECK_NULL_WITH_CONTEXT(context, yShape);
229+ 
230+ if (Ops::Base::IsUnknownRank(*xShape)) {
231+ Ops::Base::SetUnknownRank(*yShape);
232+ return ge::GRAPH_SUCCESS;
233+ }
234+ 
235+ if (Ops::Base::IsUnknownShape(*xShape)) {
236+ Ops::Base::SetUnknownShape(xShape->GetDimNum(), *yShape);
237+ return ge::GRAPH_SUCCESS;
238+ }
239+ 
240+ auto src_td = context->GetInputDesc(0);
241+ OP_CHECK_NULL_WITH_CONTEXT(context, src_td);
242+ auto input_format = src_td->GetOriginFormat();
243+ size_t d_dim = input_format == FORMAT_NDHWC ? 1 : 2;
244+ size_t h_dim = input_format == FORMAT_NDHWC ? 2 : 3;
245+ size_t w_dim = input_format == FORMAT_NDHWC ? 3 : 4;
246+ 
247+ auto attrs = context->GetAttrs();
248+ OP_CHECK_NULL_WITH_CONTEXT(context, attrs);
249+ auto padding_mode = attrs->GetAttrPointer<char>(INDEX_PADDING);
250+ OP_CHECK_NULL_WITH_CONTEXT(context, padding_mode);
251+ auto it = std::find_if(
252+ kFuncMap.begin(), kFuncMap.end(),
253+ [&padding_mode](const std::pair<std::string, InferShapePaddingFunc>& item) -> bool {
254+ return item.first == padding_mode;
255+ });
256+ OP_CHECK_IF(
257+ it == kFuncMap.end(),
258+ OP_LOGE(context->GetNodeName(), "padding_mode %s must in (CALCULATED, VALID, SAME).", padding_mode),
259+ return GRAPH_FAILED);
260+ 
261+ // when padding_mode in (CALCULATED, VALID, SAME)
262+ return it->second(context, d_dim, h_dim, w_dim, attrs);
263+}
264+ 
265+static ge::graphStatus InferDataTypeForMaxPool3D(gert::InferDataTypeContext* context)
266+{
267+ if (context == nullptr) {
268+ return GRAPH_FAILED;
269+ }
270+ 
271+ const ge::DataType xDtype = context->GetInputDataType(0);
272+ context->SetOutputDataType(0, xDtype);
273+ return GRAPH_SUCCESS;
274+}
275+ 
276+IMPL_OP_INFERSHAPE(MaxPool3D).InferShape(InferShape4MaxPool3D).InferDataType(InferDataTypeForMaxPool3D);
277+} // namespace ops
Rpooling/max_pool_v3/op_host/op_api/aclnn_max_pool.cpp→pooling/max_pool_v3/op_api/aclnn_max_pool.cpp+0-0
Rpooling/max_pool_v3/op_host/op_api/aclnn_max_pool.h→pooling/max_pool_v3/op_api/aclnn_max_pool.h+0-0
Rpooling/max_pool_v3/op_host/op_api/max_pool_v3.cpp→pooling/max_pool_v3/op_api/max_pool_v3.cpp+0-0
Rpooling/max_pool_v3/op_host/op_api/max_pool_v3.h→pooling/max_pool_v3/op_api/max_pool_v3.h+0-0
Rpooling/max_pool_v3/tests/ut/op_host/test_aclnn_max_pool.cpp→pooling/max_pool_v3/tests/ut/op_api/test_aclnn_max_pool.cpp+5-5