Star886
1788
代码
代码
Issues236
Pull Requests322
Actions
项目讨论
Wiki
分析
Star886
1788
已合并
AscendC实现KlDivLossGrad算子贡献 #5358
hqnzw创建于 5月28日
AscendC实现KlDivLossGrad算子贡献 #5358
已合并
hqnzw创建于 5月28日
从hqnzw/ops-nn_7534master合入到cann/ops-nnmaster
共 24 个文件变更+2594-0
Aexperimental/loss/kl_div_loss_grad/CMakeLists.txt+19-0
@@ -0,0 +1,19 @@
1+ # ----------------------------------------------------------------------------
2+ # Copyright (c) 2026 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(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+if(NOT ENABLE_TEST AND NOT BENCHMARK)
13+ list(REMOVE_ITEM CURRENT_DIRS tests)
14+endif()
15+foreach(SUB_DIR ${CURRENT_DIRS})
16+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
17+ add_subdirectory(${SUB_DIR})
18+ endif()
19+endforeach()
Aexperimental/loss/kl_div_loss_grad/README.md+36-0
@@ -0,0 +1,36 @@
1+# KlDivLossGrad
2+ 
3+## 产品支持情况
4+ 
5+|产品 | 是否支持 |
6+|:-------------------------|:----------:|
7+| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |
8+ 
9+## 功能说明
10+ 
11+算子功能：进行[aclnnKlDiv](https://gitcode.com/cann/ops-math/blob/master/math/kl_div_v2/docs/aclnnKlDiv.md) api的结果的反向计算。
12+ 
13+## 参数说明
14+ 
15+- gradOutput(aclTensor*, 计算输入)：Device侧的aclTensor，[数据格式](../../../docs/zh/context/数据格式.md)支持ND。shape需要与self满足[broadcast关系](../../../docs/zh/context/broadcast关系.md)。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)。数据类型支持FLOAT、FLOAT16、BFLOAT16。
16+- self(aclTensor*, 计算输入)：Device侧的aclTensor。[数据格式](../../../docs/zh/context/数据格式.md)支持ND。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)。数据类型支持FLOAT、FLOAT16、BFLOAT16。
17+- target(aclTensor*, 计算输入)：Device侧的aclTensor，[数据格式](../../../docs/zh/context/数据格式.md)支持ND。shape需要与self满足[broadcast关系](../../../docs/zh/context/broadcast关系.md)。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)。数据类型支持FLOAT、FLOAT16、BFLOAT16。
18+- reduction(int64_t, 计算输入)：Host侧的int64_t，[数据格式](../../../docs/zh/context/数据格式.md)支持ND。指定要应用到输出的缩减。支持0（‘none’）| 1（‘mean’）| 2（‘sum’）|3（‘batchmean’）。‘none’表示不应用减少，‘mean’表示输出的总和将除以输出中的元素数，‘sum’表示输出将被求和，‘batchmean’表示输出的总和将除以batch的个数。
19+- logTarget(bool, 计算输入)：Host侧的BOOL类型，是否对target进行log空间转换。
20+- out(aclTensor*, 计算输出)：Device侧的aclTensor。支持[非连续的Tensor](../../../docs/zh/context/非连续的Tensor.md)。数据类型支持FLOAT、FLOAT16、BFLOAT16。
21+ 
22+## 约束说明
23+ 
24+无。
25+ 
26+## 调用说明
27+ 
28+| 调用方式 | 样例代码 | 说明 |
29+| ---------------- | --------------------------- | --------------------------------------------------- |
30+| aclnn接口 | [test_aclnn_kl_div_loss_grad.cpp](examples/test_aclnn_kl_div_loss_grad.cpp) | 通过[aclnnKlDivBackward](docs/aclnnKlDivBackward.md)接口方式调用KlDivLossGrad算子。 |
31+ 
32+## 贡献说明
33+ 
34+| 贡献者 | 贡献方 | 贡献算子 | 贡献时间 | 贡献内容 |
35+| ---- | ---- | ---- | ---- | ---- |
36+| 韩智惟 | 北京交通大学-赵宏智老师团队 | KlDivLossGrad | 2026/05/28 | KlDivLossGrad算子适配开源仓 |
Aexperimental/loss/kl_div_loss_grad/docs/aclnnKlDivBackward.md+394-0
@@ -0,0 +1,394 @@
1+# aclnnKlDivBackward
2+ 
3+## 产品支持情况
4+ 
5+|产品 | 是否支持 |
6+|:-------------------------|:----------:|
7+| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |
8+ 
9+## 功能说明
10+ 
11+进行[aclnnKlDiv](https://gitcode.com/cann/ops-math/blob/master/math/kl_div_v2/docs/aclnnKlDiv.md) api的结果的反向计算。
12+ 
13+## 函数原型
14+ 
15+每个算子分为[两段式接口](../../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnKlDivBackwardGetWorkspaceSize”接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用“aclnnKlDivBackward”接口执行计算。
16+ 
17+```Cpp
18+aclnnStatus aclnnKlDivBackwardGetWorkspaceSize(
19+ const aclTensor* gradOutput,
20+ const aclTensor* self,
21+ const aclTensor* target,
22+ int64_t reduction,
23+ bool logTarget,
24+ aclTensor* out,
25+ uint64_t* workspaceSize,
26+ aclOpExecutor** executor)
27+```
28+ 
29+```Cpp
30+aclnnStatus aclnnKlDivBackward(
31+ void* workspace,
32+ uint64_t workspaceSize,
33+ aclOpExecutor* executor,
34+ aclrtStream stream)
35+```
36+ 
37+## aclnnKlDivBackwardGetWorkspaceSize
38+ 
39+- **参数说明：**
40+ 
41+ <table class="tg" style="undefined;table-layout: fixed; width: 1435px"><colgroup>
42+ <col style="width: 205px">
43+ <col style="width: 120px">
44+ <col style="width: 300px">
45+ <col style="width: 300px">
46+ <col style="width: 130px">
47+ <col style="width: 115px">
48+ <col style="width: 120px">
49+ <col style="width: 145px">
50+ </colgroup>
51+ <thead>
52+ <tr>
53+ <th class="tg-0pky">参数名</th>
54+ <th class="tg-0pky">输入/输出</th>
55+ <th class="tg-0pky">描述</th>
56+ <th class="tg-0pky">使用说明</th>
57+ <th class="tg-0pky">数据类型</th>
58+ <th class="tg-0pky">数据格式</th>
59+ <th class="tg-0pky">维度(shape)</th>
60+ <th class="tg-0pky">非连续Tensor</th>
61+ </tr></thead>
62+ <tbody>
63+ <tr>
64+ <td class="tg-0pky">gradOutput（aclTensor*）</td>
65+ <td class="tg-0pky">输入</td>
66+ <td class="tg-0pky">梯度反向输入。</td>
67+ <td class="tg-0pky">shape需要与self满足<a href="../../../../docs/zh/context/broadcast关系.md" target="_blank">broadcast</a>关系。</td>
68+ <td class="tg-0pky">FLOAT、FLOAT16、BFLOAT16</td>
69+ <td class="tg-0pky">ND</td>
70+ <td class="tg-0pky">1-8</td>
71+ <td class="tg-0pky">√</td>
72+ </tr>
73+ <tr>
74+ <td class="tg-0pky">self（aclTensor*）</td>
75+ <td class="tg-0pky">输入</td>
76+ <td class="tg-0pky">输入张量。</td>
77+ <td class="tg-0pky">-</td>
78+ <td class="tg-0pky">FLOAT、FLOAT16、BFLOAT16</td>
79+ <td class="tg-0pky">ND</td>
80+ <td class="tg-0pky">1-8</td>
81+ <td class="tg-0pky">√</td>
82+ </tr>
83+ <tr>
84+ <td class="tg-0pky">target（aclTensor*）</td>
85+ <td class="tg-0pky">输入</td>
86+ <td class="tg-0pky">真实的标签。</td>
87+ <td class="tg-0pky">shape需要与self满足<a href="../../../../docs/zh/context/broadcast关系.md" target="_blank">broadcast</a>关系。</td>
88+ <td class="tg-0pky">FLOAT、FLOAT16、BFLOAT16</td>
89+ <td class="tg-0pky">ND</td>
90+ <td class="tg-0pky">1-8</td>
91+ <td class="tg-0pky">√</td>
92+ </tr>
93+ <tr>
94+ <td class="tg-0pky">reduction（int64_t）</td>
95+ <td class="tg-0pky">输入</td>
96+ <td class="tg-0pky">指定要应用到输出的缩减。</td>
97+ <td class="tg-0pky">支持0(none)|1(mean)|2(sum)|3(batchmean)。<br>'none'表示不应用缩减。<br>'mean'表示输出的总和将除以输出中的元素数。<br>'sum'表示输出将被求和。<br>'batchmean'表示输出的总和将除以batch的个数。</td>
98+ <td class="tg-0pky">INT64</td>
99+ <td class="tg-0pky">-</td>
100+ <td class="tg-0pky">-</td>
101+ <td class="tg-0pky">-</td>
102+ </tr>
103+ <tr>
104+ <td class="tg-0pky">logTarget（bool）</td>
105+ <td class="tg-0pky">输入</td>
106+ <td class="tg-0pky">是否对target进行log空间转换。</td>
107+ <td class="tg-0pky">-</td>
108+ <td class="tg-0pky">BOOL</td>
109+ <td class="tg-0pky">-</td>
110+ <td class="tg-0pky">-</td>
111+ <td class="tg-0pky">√</td>
112+ </tr>
113+ <tr>
114+ <td class="tg-0pky">out（aclTensor*）</td>
115+ <td class="tg-0pky">输出</td>
116+ <td class="tg-0pky">输出的损失。</td>
117+ <td class="tg-0pky">-</td>
118+ <td class="tg-0pky">与self保持一致</td>
119+ <td class="tg-0pky">ND</td>
120+ <td class="tg-0pky">1-8</td>
121+ <td class="tg-0pky">√</td>
122+ </tr>
123+ <tr>
124+ <td class="tg-0pky">workspaceSize（uint64_t*）</td>
125+ <td class="tg-0pky">输出</td>
126+ <td class="tg-0pky">返回需要在Device侧申请的workspace大小。</td>
127+ <td class="tg-0pky">-</td>
128+ <td class="tg-0pky">-</td>
129+ <td class="tg-0pky">-</td>
130+ <td class="tg-0pky">-</td>
131+ <td class="tg-0pky">-</td>
132+ </tr>
133+ <tr>
134+ <td class="tg-0pky">executor（aclOpExecutor**）</td>
135+ <td class="tg-0pky">输出</td>
136+ <td class="tg-0pky">返回op执行器，包含了算子计算流程。</td>
137+ <td class="tg-0pky">-</td>
138+ <td class="tg-0pky">-</td>
139+ <td class="tg-0pky">-</td>
140+ <td class="tg-0pky">-</td>
141+ <td class="tg-0pky">-</td>
142+ </tr>
143+ </tbody></table>
144+ 
145+- **返回值：**
146+ 
147+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../../docs/zh/context/aclnn返回码.md)。
148+ 
149+ 第一段接口完成入参校验，出现以下场景时报错：
150+ 
151+ <table class="tg" style="undefined;table-layout: fixed; width: 991px"><colgroup>
152+ <col style="width: 269px">
153+ <col style="width: 90px">
154+ <col style="width: 632px">
155+ </colgroup>
156+ <thead>
157+ <tr>
158+ <th class="tg-0pky">返回值</th>
159+ <th class="tg-0pky">错误码</th>
160+ <th class="tg-0pky">描述</th>
161+ </tr></thead>
162+ <tbody>
163+ <tr>
164+ <td class="tg-0pky">ACLNN_ERR_PARAM_NULLPTR</td>
165+ <td class="tg-0pky">161001</td>
166+ <td class="tg-0pky">传入的gradOutput、self、target和out是空指针。</td>
167+ </tr>
168+ <tr>
169+ <td class="tg-0pky" rowspan="6">ACLNN_ERR_PARAM_INVALID</td>
170+ <td class="tg-0pky" rowspan="6">161002</td>
171+ <td class="tg-0pky">gradOutput、self、target和out的数据类型不在支持的范围内时。</td>
172+ </tr>
173+ <tr>
174+ <td class="tg-0pky">self、out的数据类型不一致。</td>
175+ </tr>
176+ <tr>
177+ <td class="tg-0pky">gradOutput的shape不能向self或者target做<a href="../../../../docs/zh/context/broadcast关系.md" target="_blank">broadcast</a>。</td>
178+ </tr>
179+ <tr>
180+ <td class="tg-0pky">target的shape和self的shape不满足<a href="../../../../docs/zh/context/broadcast关系.md" target="_blank">broadcast</a>关系。</td>
181+ </tr>
182+ <tr>
183+ <td class="tg-0lax">self的shape与out的shape不相同。</td>
184+ </tr>
185+ <tr>
186+ <td class="tg-0lax">gradOutput、self、target或者out维度大于8。</td>
187+ </tr>
188+ </tbody>
189+ </table>
190+ 
191+## aclnnKlDivBackward
192+ 
193+- **参数说明：**
194+ 
195+ <table style="undefined;table-layout: fixed; width: 1244px"><colgroup>
196+ <col style="width: 200px">
197+ <col style="width: 162px">
198+ <col style="width: 882px">
199+ </colgroup>
200+ <thead>
201+ <tr>
202+ <th>参数名</th>
203+ <th>输入/输出</th>
204+ <th>描述</th>
205+ </tr></thead>
206+ <tbody>
207+ <tr>
208+ <td>workspace</td>
209+ <td>输入</td>
210+ <td>在Device侧申请的workspace内存地址。</td>
211+ </tr>
212+ <tr>
213+ <td>workspaceSize</td>
214+ <td>输入</td>
215+ <td>在Device侧申请的workspace大小，由第一段接口aclnnKlDivBackwardGetWorkspaceSize获取。</td>
216+ </tr>
217+ <tr>
218+ <td>executor</td>
219+ <td>输入</td>
220+ <td>op执行器，包含了算子计算流程。</td>
221+ </tr>
222+ <tr>
223+ <td>stream</td>
224+ <td>输入</td>
225+ <td>指定执行任务的Stream。</td>
226+ </tr>
227+ </tbody>
228+ </table>
229+ 
230+- **返回值：**
231+ 
232+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../../docs/zh/context/aclnn返回码.md)。
233+ 
234+## 约束说明
235+ 
236+- 确定性计算：
237+ - aclnnKlDivBackward默认确定性实现。
238+ 
239+## 调用示例
240+ 
241+示例代码如下，仅供参考，具体编译和执行过程请参考[编译与运行样例](../../../../docs/zh/context/编译与运行样例.md)。
242+ 
243+```Cpp
244+#include <iostream>
245+#include <vector>
246+#include "acl/acl.h"
247+#include "aclnnop/aclnn_kl_div_backward.h"
248+ 
249+#define CHECK_RET(cond, return_expr) \
250+ do { \
251+ if (!(cond)) { \
252+ return_expr; \
253+ } \
254+ } while (0)
255+ 
256+#define LOG_PRINT(message, ...) \
257+ do { \
258+ printf(message, ##__VA_ARGS__); \
259+ } while (0)
260+ 
261+int64_t GetShapeSize(const std::vector<int64_t>& shape) {
262+ int64_t shapeSize = 1;
263+ for (auto i : shape) {
264+ shapeSize *= i;
265+ }
266+ return shapeSize;
267+}
268+ 
269+int Init(int32_t deviceId, aclrtStream* stream) {
270+ // 固定写法，资源初始化
271+ auto ret = aclInit(nullptr);
272+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
273+ ret = aclrtSetDevice(deviceId);
274+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
275+ ret = aclrtCreateStream(stream);
276+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
277+ return 0;
278+}
279+ 
280+template <typename T>
281+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
282+ aclDataType dataType, aclTensor** tensor) {
283+ auto size = GetShapeSize(shape) * sizeof(T);
284+ // 调用aclrtMalloc申请device侧内存
285+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
286+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
287+ 
288+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
289+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
290+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
291+ 
292+ // 计算连续tensor的strides
293+ std::vector<int64_t> strides(shape.size(), 1);
294+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
295+ strides[i] = shape[i + 1] * strides[i + 1];
296+ }
297+ 
298+ // 调用aclCreateTensor接口创建aclTensor
299+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
300+ shape.data(), shape.size(), *deviceAddr);
301+ return 0;
302+}
303+ 
304+int main() {
305+ // 1. （固定写法）device/stream初始化，参考acl API手册
306+ // 根据自己的实际device填写deviceId
307+ int32_t deviceId = 0;
308+ aclrtStream stream;
309+ auto ret = Init(deviceId, &stream);
310+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
311+ 
312+ // 2. 构造输入与输出，需要根据API的接口自定义构造
313+ std::vector<int64_t> gradOutputShape = {2, 2};
314+ std::vector<int64_t> selfShape = {2, 2};
315+ std::vector<int64_t> targetShape = {2, 2};
316+ std::vector<int64_t> outShape = {2, 2};
317+ void* gradOutputDeviceAddr = nullptr;
318+ void* selfDeviceAddr = nullptr;
319+ void* targetDeviceAddr = nullptr;
320+ void* outDeviceAddr = nullptr;
321+ aclTensor* gradOutput = nullptr;
322+ aclTensor* self = nullptr;
323+ aclTensor* target = nullptr;
324+ aclTensor* out = nullptr;
325+ std::vector<float> gradOutputHostData = {2, 3, 5, 8};
326+ std::vector<float> selfHostData = {2, 3, 5, 8};
327+ std::vector<float> targetHostData = {2, 3, 5, 8};
328+ std::vector<float> outHostData = {2, 3, 5, 8};
329+ int64_t reduction = 0;
330+ bool logTarget = false;
331+ // 创建gradOutput aclTensor
332+ ret = CreateAclTensor(gradOutputHostData, gradOutputShape, &gradOutputDeviceAddr, aclDataType::ACL_FLOAT, &gradOutput);
333+ CHECK_RET(ret == ACL_SUCCESS, return ret);
334+ // 创建self aclTensor
335+ ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self);
336+ CHECK_RET(ret == ACL_SUCCESS, return ret);
337+ // 创建target aclTensor
338+ ret = CreateAclTensor(targetHostData, targetShape, &targetDeviceAddr, aclDataType::ACL_FLOAT, &target);
339+ CHECK_RET(ret == ACL_SUCCESS, return ret);
340+ // 创建out aclTensor
341+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
342+ CHECK_RET(ret == ACL_SUCCESS, return ret);
343+ 
344+ // 3. 调用CANN算子库API
345+ uint64_t workspaceSize = 0;
346+ aclOpExecutor* executor;
347+ // 调用aclnnKlDivBackward第一段接口
348+ ret = aclnnKlDivBackwardGetWorkspaceSize(gradOutput, self, target, reduction, logTarget, out, &workspaceSize, &executor);
349+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnKlDivBackwardGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
350+ // 根据第一段接口计算出的workspaceSize申请device内存
351+ void* workspaceAddr = nullptr;
352+ if (workspaceSize > 0) {
353+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
354+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
355+ }
356+ // 调用aclnnKlDivBackward第二段接口
357+ ret = aclnnKlDivBackward(workspaceAddr, workspaceSize, executor, stream);
358+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnKlDivBackward failed. ERROR: %d\n", ret); return ret);
359+ 
360+ // 4. （固定写法）同步等待任务执行结束
361+ ret = aclrtSynchronizeStream(stream);
362+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
363+ 
364+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
365+ auto size = GetShapeSize(outShape);
366+ std::vector<float> resultData(size, 0);
367+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]),
368+ outDeviceAddr, size * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
369+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy resultData from device to host failed. ERROR: %d\n", ret);
370+ return ret);
371+ for (int64_t i = 0; i < size; i++) {
372+ LOG_PRINT("resultData[%ld] is: %f\n", i, resultData[i]);
373+ }
374+ 
375+ // 6. 释放aclTensor，需要根据具体API的接口定义修改
376+ aclDestroyTensor(gradOutput);
377+ aclDestroyTensor(self);
378+ aclDestroyTensor(target);
379+ aclDestroyTensor(out);
380+ 
381+ // 7. 释放device资源，需要根据具体API的接口定义修改
382+ aclrtFree(gradOutputDeviceAddr);
383+ aclrtFree(selfDeviceAddr);
384+ aclrtFree(targetDeviceAddr);
385+ aclrtFree(outDeviceAddr);
386+ if (workspaceSize > 0) {
387+ aclrtFree(workspaceAddr);
388+ }
389+ aclrtDestroyStream(stream);
390+ aclrtResetDevice(deviceId);
391+ aclFinalize();
392+ return 0;
393+}
394+```
Aexperimental/loss/kl_div_loss_grad/examples/test_aclnn_kl_div_loss_grad.cpp+169-0
@@ -0,0 +1,169 @@
1+/**
2+ * Copyright (c) 2026 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 "aclnn_kl_div_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+{
30+ int64_t shapeSize = 1;
31+ for (auto i : shape) {
32+ shapeSize *= i;
33+ }
34+ return shapeSize;
35+}
36+ 
37+int Init(int32_t deviceId, aclrtStream* stream)
38+{
39+ // 固定写法，资源初始化
40+ auto ret = aclInit(nullptr);
41+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
42+ ret = aclrtSetDevice(deviceId);
43+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
44+ ret = aclrtCreateStream(stream);
45+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
46+ return 0;
47+}
48+ 
49+template <typename T>
50+int CreateAclTensor(
51+ const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr, aclDataType dataType,
52+ aclTensor** tensor)
53+{
54+ auto size = GetShapeSize(shape) * sizeof(T);
55+ // 调用aclrtMalloc申请device侧内存
56+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
57+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
58+ 
59+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
60+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
61+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
62+ 
63+ // 计算连续tensor的strides
64+ std::vector<int64_t> strides(shape.size(), 1);
65+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
66+ strides[i] = shape[i + 1] * strides[i + 1];
67+ }
68+ 
69+ // 调用aclCreateTensor接口创建aclTensor
70+ *tensor = aclCreateTensor(
71+ shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND, shape.data(), shape.size(),
72+ *deviceAddr);
73+ return 0;
74+}
75+ 
76+int main()
77+{
78+ // 1. （固定写法）device/stream初始化，参考acl API手册
79+ // 根据自己的实际device填写deviceId
80+ int32_t deviceId = 0;
81+ aclrtStream stream;
82+ auto ret = Init(deviceId, &stream);
83+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
84+ 
85+ // 2. 构造输入与输出，需要根据API的接口自定义构造
86+ std::vector<int64_t> gradOutputShape = {2, 2};
87+ std::vector<int64_t> selfShape = {2, 2};
88+ std::vector<int64_t> targetShape = {2, 2};
89+ std::vector<int64_t> outShape = {2, 2};
90+ void* gradOutputDeviceAddr = nullptr;
91+ void* selfDeviceAddr = nullptr;
92+ void* targetDeviceAddr = nullptr;
93+ void* outDeviceAddr = nullptr;
94+ aclTensor* gradOutput = nullptr;
95+ aclTensor* self = nullptr;
96+ aclTensor* target = nullptr;
97+ aclTensor* out = nullptr;
98+ std::vector<float> gradOutputHostData = {2, 3, 5, 8};
99+ std::vector<float> selfHostData = {2, 3, 5, 8};
100+ std::vector<float> targetHostData = {2, 3, 5, 8};
101+ std::vector<float> outHostData = {2, 3, 5, 8};
102+ int64_t reduction = 0;
103+ bool logTarget = false;
104+ // 创建gradOutput aclTensor
105+ ret = CreateAclTensor(
106+ gradOutputHostData, gradOutputShape, &gradOutputDeviceAddr, aclDataType::ACL_FLOAT, &gradOutput);
107+ CHECK_RET(ret == ACL_SUCCESS, return ret);
108+ // 创建self aclTensor
109+ ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self);
110+ CHECK_RET(ret == ACL_SUCCESS, return ret);
111+ // 创建target aclTensor
112+ ret = CreateAclTensor(targetHostData, targetShape, &targetDeviceAddr, aclDataType::ACL_FLOAT, &target);
113+ CHECK_RET(ret == ACL_SUCCESS, return ret);
114+ // 创建out aclTensor
115+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
116+ CHECK_RET(ret == ACL_SUCCESS, return ret);
117+ 
118+ // 3. 调用CANN算子库API
119+ uint64_t workspaceSize = 0;
120+ aclOpExecutor* executor;
121+ // 调用aclnnKlDivBackward第一段接口
122+ ret = aclnnKlDivBackwardGetWorkspaceSize(
123+ gradOutput, self, target, reduction, logTarget, out, &workspaceSize, &executor);
124+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnKlDivBackwardGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
125+ // 根据第一段接口计算出的workspaceSize申请device内存
126+ void* workspaceAddr = nullptr;
127+ if (workspaceSize > 0) {
128+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
129+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
130+ }
131+ // 调用aclnnKlDivBackward第二段接口
132+ ret = aclnnKlDivBackward(workspaceAddr, workspaceSize, executor, stream);
133+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnKlDivBackward failed. ERROR: %d\n", ret); return ret);
134+ 
135+ // 4. （固定写法）同步等待任务执行结束
136+ ret = aclrtSynchronizeStream(stream);
137+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
138+ 
139+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
140+ auto size = GetShapeSize(outShape);
141+ std::vector<float> resultData(size, 0);
142+ ret = aclrtMemcpy(
143+ resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr, size * sizeof(resultData[0]),
144+ ACL_MEMCPY_DEVICE_TO_HOST);
145+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy resultData from device to host failed. ERROR: %d\n", ret);
146+ return ret);
147+ for (int64_t i = 0; i < size; i++) {
148+ LOG_PRINT("resultData[%ld] is: %f\n", i, resultData[i]);
149+ }
150+ 
151+ // 6. 释放aclTensor，需要根据具体API的接口定义修改
152+ aclDestroyTensor(gradOutput);
153+ aclDestroyTensor(self);
154+ aclDestroyTensor(target);
155+ aclDestroyTensor(out);
156+ 
157+ // 7. 释放device资源，需要根据具体API的接口定义修改
158+ aclrtFree(gradOutputDeviceAddr);
159+ aclrtFree(selfDeviceAddr);
160+ aclrtFree(targetDeviceAddr);
161+ aclrtFree(outDeviceAddr);
162+ if (workspaceSize > 0) {
163+ aclrtFree(workspaceAddr);
164+ }
165+ aclrtDestroyStream(stream);
166+ aclrtResetDevice(deviceId);
167+ aclFinalize();
168+ return 0;
169+}
Aexperimental/loss/kl_div_loss_grad/op_host/CMakeLists.txt+11-0
@@ -0,0 +1,11 @@
1+ # ----------------------------------------------------------------------------
2+ # Copyright (c) 2026 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+add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE kl_div_loss_grad ACLNNTYPE aclnn_exclude)
Aexperimental/loss/kl_div_loss_grad/op_host/kl_div_loss_grad_def.cpp+35-0
@@ -0,0 +1,35 @@
1+/**
F
Ffulltower5月29日
  1. License 年份错误：所有 13 个源文件使用 Copyright (c) 2025，应为 2026。涉及文件：kl_div_loss_grad_def.cpp、kl_div_loss_grad_infershape.cpp、kl_div_loss_grad_tiling.cpp、kl_div_loss_grad.cpp（kernel）、kl_div_loss_grad.h、kl_div_loss_grad_tiling_data.h、kl_div_loss_grad_tiling_key.h、aclnn_kl_div_backward.cpp、aclnn_kl_div_backward.h、kl_div_loss_grad.cpp（op_api）、kl_div_loss_grad.h、test_aclnn_kl_div_loss_grad.cpp、CMakeLists.txt
likedislike
2+ * Copyright (c) 2026 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 kl_div_loss_grad.cpp
13+ * \brief
14+ */
15+#include "register/op_def_registry.h"
16+ 
17+namespace ops {
18+static const std::vector<ge::DataType> dataType = {ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT};
19+static const std::vector<ge::Format> dataFormat = {ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND};
20+class KlDivLossGrad : public OpDef {
21+public:
22+ explicit KlDivLossGrad(const char* name) : OpDef(name)
23+ {
24+ this->Input("grad").ParamType(REQUIRED).DataType(dataType).Format(dataFormat);
25+ this->Input("input").ParamType(REQUIRED).DataType(dataType).Format(dataFormat);
26+ this->Input("target").ParamType(REQUIRED).DataType(dataType).Format(dataFormat);
27+ this->Output("y").ParamType(REQUIRED).DataType(dataType).Format(dataFormat);
28+ this->Attr("reduction").AttrType(OPTIONAL).String("mean");
29+ this->Attr("log_target").AttrType(OPTIONAL).Bool(false);
30+ 
31+ this->AICore().AddConfig("ascend910b");
32+ }
33+};
34+OP_ADD(KlDivLossGrad);
35+} // namespace ops
Aexperimental/loss/kl_div_loss_grad/op_host/kl_div_loss_grad_infershape.cpp+47-0
@@ -0,0 +1,47 @@
1+/**
2+ * Copyright (c) 2026 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 kl_div_loss_grad_infer.cpp
13+ * \brief
14+ */
15+#include "register/op_impl_registry.h"
16+#include "log/log.h"
17+ 
18+using namespace ge;
19+ 
20+namespace ops {
21+ge::graphStatus CopyShapeInput2OutputWithForKlDivLossGrad(
22+ gert::InferShapeContext* context, int64_t input_idx, int64_t output_idx)
23+{
24+ auto in_shape = context->GetInputShape(input_idx);
25+ OP_CHECK_NULL_WITH_CONTEXT(context, in_shape);
26+ auto out_shape = context->GetOutputShape(output_idx);
27+ OP_CHECK_NULL_WITH_CONTEXT(context, out_shape);
28+ *out_shape = *in_shape;
29+ return ge::GRAPH_SUCCESS;
30+}
31+ 
32+static ge::graphStatus InferShape4InIdx1AndOutIdx0(gert::InferShapeContext* context)
33+{
34+ constexpr size_t input_index = 1;
35+ constexpr size_t output_index = 0;
36+ return CopyShapeInput2OutputWithForKlDivLossGrad(context, input_index, output_index);
37+}
38+ 
39+graphStatus InferDtypeForKlDivLossGrad(gert::InferDataTypeContext* context)
40+{
41+ auto xDtype = context->GetInputDataType(0);
42+ context->SetOutputDataType(0, xDtype);
43+ return GRAPH_SUCCESS;
44+}
45+ 
46+IMPL_OP_INFERSHAPE(KlDivLossGrad).InferShape(InferShape4InIdx1AndOutIdx0).InferDataType(InferDtypeForKlDivLossGrad);
47+} // namespace ops
Aexperimental/loss/kl_div_loss_grad/op_host/kl_div_loss_grad_tiling.cpp+282-0
@@ -0,0 +1,282 @@
1+/**
2+ * Copyright (c) 2026 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 kl_div_loss_grad_tiling.cpp
13+ * \brief
14+ */
15+ 
16+#include "log/log.h"
17+#include "util/math_util.h"
18+#include "util/platform_util.h"
19+#include "op_host/tiling_util.h"
20+#include "op_host/tiling_templates_registry.h"
21+#include "../op_kernel/kl_div_loss_grad_tiling_data.h"
22+#include "../op_kernel/kl_div_loss_grad_tiling_key.h"
23+ 
24+namespace optiling {
25+ 
26+struct KlDivLossGradCompileInfo {};
27+ 
28+enum Reduction {
29+ NONE,
30+ MEAN,
F
Ffulltower5月29日
  1. BLOCK_SIZE 硬编码为 32，不同芯片型号下可能存在兼容性问题。建议改为 uint32_t BLOCK_SIZE = Ops::Base::GetUbBlockSize(context); 动态获取。
likedislike
31+ SUM,
32+ BATCHMEAN,
33+};
34+ 
35+static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t& ubSize, int64_t& coreNum)
36+{
37+ // 获取平台信息指针
38+ fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo();
39+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr);
40+ 
41+ // 创建AscendC平台对象
42+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
43+ 
44+ // 获取AI Core数量
45+ coreNum = ascendcPlatform.GetCoreNumAiv();
46+ OP_CHECK_IF(coreNum == 0, OP_LOGE(context, "coreNum is 0"), return ge::GRAPH_FAILED);
47+ 
48+ // 获取UB(统一缓冲区)大小
49+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
50+ OP_CHECK_IF(ubSize == 0, OP_LOGE(context, "ubSize is 0"), return ge::GRAPH_FAILED);
51+ 
52+ return ge::GRAPH_SUCCESS;
53+}
54+ 
55+static ge::graphStatus GetShapeInfo(
56+ gert::TilingContext* context, int64_t& inputNum, int64_t& gradNum, int64_t& batchSize)
57+{
58+ // 获取输入grad的shape信息
59+ auto inputGrad = context->GetInputShape(0);
60+ OP_CHECK_NULL_WITH_CONTEXT(context, inputGrad);
61+ 
62+ // 如果输入shape是标量，转换为{1}，否则保持原shape不变
63+ auto inputShapeGrad = Ops::NN::OpTiling::EnsureNotScalar(inputGrad->GetStorageShape());
64+ gradNum = inputShapeGrad.GetShapeSize();
65+ 
66+ // 获取输入input的shape信息
67+ auto inputInput = context->GetInputShape(1);
68+ OP_CHECK_NULL_WITH_CONTEXT(context, inputInput);
69+ 
70+ // 如果输入shape是标量，转换为{1}，否则保持原shape不变
71+ auto inputShapeInput = Ops::NN::OpTiling::EnsureNotScalar(inputInput->GetStorageShape());
72+ 
73+ // 获取shape的大小与batch大小(用于后续batchmean计算)
74+ inputNum = inputShapeInput.GetShapeSize();
75+ batchSize = inputShapeInput.GetDim(0);
76+ 
77+ return ge::GRAPH_SUCCESS;
78+}
79+ 
80+static ge::graphStatus GetAndValidateDtype(gert::TilingContext* context, ge::DataType& inputType)
81+{
82+ const std::set<ge::DataType> supportedDtype = {ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16};
83+ auto inputDesc = context->GetInputDesc(0);
84+ OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc);
85+ inputType = inputDesc->GetDataType();
86+ 
87+ OP_CHECK_IF(
88+ supportedDtype.count(inputType) == 0,
89+ OP_LOGE(context, "KlDivLossGrad: invalid dtype, only support float, fp16 and bf16"), return ge::GRAPH_FAILED);
90+ 
91+ return ge::GRAPH_SUCCESS;
92+}
93+ 
94+static ge::graphStatus ParseAttrs(gert::TilingContext* context, Reduction& reduction, bool& logTarget)
95+{
96+ const gert::RuntimeAttrs* attrs = context->GetAttrs();
97+ OP_CHECK_NULL_WITH_CONTEXT(context, attrs);
98+ 
99+ const char* reductionPtr = attrs->GetStr(0);
100+ OP_CHECK_NULL_WITH_CONTEXT(context, reductionPtr);
101+ 
102+ const bool* logTargetPtr = attrs->GetBool(1);
103+ OP_CHECK_NULL_WITH_CONTEXT(context, logTargetPtr);
104+ logTarget = *logTargetPtr;
105+ 
106+ // 解析reduction属性
107+ if (std::strcmp(reductionPtr, "none") == 0) {
108+ reduction = NONE;
109+ } else if (std::strcmp(reductionPtr, "mean") == 0) {
110+ reduction = MEAN;
111+ } else if (std::strcmp(reductionPtr, "sum") == 0) {
112+ reduction = SUM;
113+ } else if (std::strcmp(reductionPtr, "batchmean") == 0) {
114+ reduction = BATCHMEAN;
115+ } else {
116+ OP_LOGE(context, "KlDivLossGrad: invalid reduction type: %s", reductionPtr);
117+ return ge::GRAPH_FAILED;
118+ }
119+ 
120+ return ge::GRAPH_SUCCESS;
121+}
122+ 
123+ge::graphStatus GetWorkspaceSize(gert::TilingContext* context)
124+{
125+ // 获取工作空间大小数组
126+ size_t* currentWorkspace = context->GetWorkspaceSizes(1);
127+ OP_CHECK_NULL_WITH_CONTEXT(context, currentWorkspace);
128+ currentWorkspace[0] = 0;
129+ return ge::GRAPH_SUCCESS;
130+}
131+ 
132+static ge::graphStatus ComputeTilingParams(
133+ gert::TilingContext* context, uint64_t ubSize, int64_t coreNum, int64_t inputNum, int64_t gradNum,
134+ ge::DataType inputType, Reduction reduction, int64_t batchSize, int64_t& bigCoreDataNum, int64_t& smallCoreDataNum,
135+ int64_t& tileDataNum, int64_t& bigCoreNum, float& coff, int64_t& usedCoreNum)
136+{
137+ int64_t ubBlockSize = Ops::Base::GetUbBlockSize(context);
138+ int64_t inputTypeLength = (inputType == ge::DT_FLOAT) ? 4 : 2;
139+ int64_t blockElemNum = ubBlockSize / inputTypeLength;
140+ 
141+ int64_t ubTileNum;
142+ if (gradNum != 1) {
143+ ubTileNum = (inputTypeLength == 2) ? 12 : 6;
144+ } else {
145+ ubTileNum = (inputTypeLength == 2) ? 8 : 4;
146+ }
147+ int64_t ubBlockNum = Ops::Base::FloorDiv(ubSize, static_cast<uint64_t>(ubBlockSize));
148+ int64_t tileBlockNum = Ops::Base::FloorDiv(ubBlockNum, ubTileNum);
149+ OP_CHECK_IF(tileBlockNum == 0, OP_LOGE(context, "tileBlockNum is 0"), return ge::GRAPH_FAILED);
150+ 
151+ tileDataNum = tileBlockNum * blockElemNum;
152+ 
153+ int64_t inputBlockNum = Ops::Base::CeilDiv(inputNum, blockElemNum);
154+ 
155+ // 核切分：优先使用更多的核并行计算
156+ int64_t maxCoreNum = Ops::Base::CeilDiv(inputBlockNum, tileBlockNum);
157+ usedCoreNum = std::max(std::min(coreNum, maxCoreNum), static_cast<int64_t>(1));
158+ 
159+ int64_t smallCoreBlockNum = Ops::Base::FloorDiv(inputBlockNum, usedCoreNum);
160+ bigCoreNum = inputBlockNum % usedCoreNum;
161+ 
162+ smallCoreDataNum = smallCoreBlockNum * blockElemNum;
163+ bigCoreDataNum = smallCoreDataNum + blockElemNum;
164+ 
165+ // 计算系数
166+ coff = 1.0f;
167+ if (reduction == BATCHMEAN) {
168+ OP_CHECK_IF(batchSize == 0, OP_LOGE(context, "batchSize is 0"), return ge::GRAPH_FAILED);
169+ coff = 1.0f / batchSize;
170+ } else if (reduction == MEAN) {
171+ OP_CHECK_IF(inputNum == 0, OP_LOGE(context, "inputNum is 0"), return ge::GRAPH_FAILED);
172+ coff = 1.0f / inputNum;
173+ }
174+ 
175+ return ge::GRAPH_SUCCESS;
176+}
177+ 
178+static ge::graphStatus SetTilingInfo(
179+ gert::TilingContext* context, int64_t bigCoreDataNum, int64_t smallCoreDataNum, int64_t tileDataNum,
180+ int64_t bigCoreNum, float coff, int64_t usedCoreNum, bool logTarget, bool isScalarGrad)
181+{
182+ KlDivLossGradTilingData* tiling = context->GetTilingData<KlDivLossGradTilingData>();
183+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
184+ 
185+ // 初始化tiling数据为0
186+ OP_CHECK_IF(
187+ memset_s(tiling, sizeof(KlDivLossGradTilingData), 0, sizeof(KlDivLossGradTilingData)) != EOK,
188+ OP_LOGE(context, "set tiling data error"), return ge::GRAPH_FAILED);
189+ 
190+ tiling->bigCoreDataNum = bigCoreDataNum;
191+ tiling->smallCoreDataNum = smallCoreDataNum;
192+ tiling->tileDataNum = tileDataNum;
193+ tiling->bigCoreNum = bigCoreNum;
194+ tiling->coff = coff;
195+ 
196+ // 设置使用的AI Core数量
197+ context->SetBlockDim(usedCoreNum);
198+ 
199+ // 根据属性设置tiling key
200+ uint64_t tilingKey = GET_TPL_TILING_KEY(logTarget, isScalarGrad);
201+ context->SetTilingKey(tilingKey);
202+ 
203+ return ge::GRAPH_SUCCESS;
204+}
205+ 
206+static ge::graphStatus KlDivLossGradTilingComputeAndSet(
207+ gert::TilingContext* context, uint64_t ubSize, int64_t coreNum, int64_t inputNum, int64_t gradNum,
208+ ge::DataType inputType, Reduction reduction, int64_t batchSize, bool logTarget)
209+{
210+ // 计算Tiling切分信息
211+ int64_t bigCoreDataNum;
212+ int64_t smallCoreDataNum;
213+ int64_t tileDataNum;
214+ int64_t bigCoreNum;
215+ float coff;
216+ int64_t usedCoreNum;
217+ OP_CHECK_IF(
218+ ComputeTilingParams(
219+ context, ubSize, coreNum, inputNum, gradNum, inputType, reduction, batchSize, bigCoreDataNum,
220+ smallCoreDataNum, tileDataNum, bigCoreNum, coff, usedCoreNum) != ge::GRAPH_SUCCESS,
221+ OP_LOGE(context, "ComputeTilingParams error"), return ge::GRAPH_FAILED);
222+ 
223+ // 设置tiling信息
224+ OP_CHECK_IF(
225+ SetTilingInfo(
226+ context, bigCoreDataNum, smallCoreDataNum, tileDataNum, bigCoreNum, coff, usedCoreNum, logTarget,
227+ (gradNum == 1)) != ge::GRAPH_SUCCESS,
228+ OP_LOGE(context, "SetTilingInfo error"), return ge::GRAPH_FAILED);
229+ 
230+ return ge::GRAPH_SUCCESS;
231+}
232+ 
233+static ge::graphStatus KlDivLossGradTilingFunc(gert::TilingContext* context)
234+{
235+ // 1、获取平台运行时信息
236+ uint64_t ubSize;
237+ int64_t coreNum;
238+ OP_CHECK_IF(
239+ GetPlatformInfo(context, ubSize, coreNum) != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetPlatformInfo error"),
240+ return ge::GRAPH_FAILED);
241+ 
242+ // 2、获取shape信息
243+ int64_t inputNum;
244+ int64_t gradNum;
245+ int64_t batchSize;
246+ OP_CHECK_IF(
247+ GetShapeInfo(context, inputNum, gradNum, batchSize) != ge::GRAPH_SUCCESS,
248+ OP_LOGE(context, "GetShapeInfo error"), return ge::GRAPH_FAILED);
249+ 
250+ // 3、获取并校验数据类型
251+ ge::DataType inputType;
252+ OP_CHECK_IF(
253+ GetAndValidateDtype(context, inputType) != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetAndValidateDtype error"),
254+ return ge::GRAPH_FAILED);
255+ 
256+ // 4、解析属性
257+ Reduction reduction;
258+ bool logTarget;
259+ OP_CHECK_IF(
260+ ParseAttrs(context, reduction, logTarget) != ge::GRAPH_SUCCESS, OP_LOGE(context, "ParseAttrs error"),
261+ return ge::GRAPH_FAILED);
262+ 
263+ // 5、获取WorkspaceSize信息
264+ OP_CHECK_IF(
265+ GetWorkspaceSize(context) != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetWorkspaceSize error"),
266+ return ge::GRAPH_FAILED);
267+ 
268+ // 6、计算Tiling切分信息并设置
269+ return KlDivLossGradTilingComputeAndSet(
270+ context, ubSize, coreNum, inputNum, gradNum, inputType, reduction, batchSize, logTarget);
271+}
272+ 
273+static ge::graphStatus TilingParseForKlDivLossGrad([[maybe_unused]] gert::TilingParseContext* context)
274+{
275+ return ge::GRAPH_SUCCESS;
276+}
277+ 
278+// tiling注册入口.
279+IMPL_OP_OPTILING(KlDivLossGrad)
280+ .Tiling(KlDivLossGradTilingFunc)
281+ .TilingParse<KlDivLossGradCompileInfo>(TilingParseForKlDivLossGrad);
282+} // namespace optiling
Aexperimental/loss/kl_div_loss_grad/op_host/op_api/aclnn_kl_div_backward.cpp+266-0
@@ -0,0 +1,266 @@
1+/**
2+ * Copyright (c) 2026 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 "aclnn_kl_div_backward.h"
12+#include "aclnn_kernels/cast.h"
13+#include "aclnn_kernels/contiguous.h"
14+#include "kl_div_loss_grad.h"
15+#include "level0/reduce_sum_op.h"
16+#include "level0/broadcast_to.h"
17+#include "aclnn_kernels/reshape.h"
18+#include "aclnn_kernels/common/op_error_check.h"
19+#include "opdev/common_types.h"
20+#include "opdev/data_type_utils.h"
21+#include "opdev/format_utils.h"
22+#include "opdev/op_dfx.h"
23+#include "opdev/op_executor.h"
24+#include "opdev/op_log.h"
25+#include "opdev/shape_utils.h"
26+#include "opdev/tensor_view_utils.h"
27+#include "loss/common/level2_base_loss.h"
28+#include "op_api/aclnn_util.h"
29+ 
30+using namespace op;
31+#ifdef __cplusplus
32+extern "C" {
33+#endif
34+ 
35+enum Reduction { None = 0, Mean = 1, Sum = 2, Batchmean = 3, End };
36+ 
37+static const char* REDUCTION_NONE = "none";
38+static const char* REDUCTION_MEAN = "mean";
39+static const char* REDUCTION_SUM = "sum";
40+static const char* REDUCTION_BATCHMEAN = "batchmean";
41+ 
42+static const std::initializer_list<op::DataType> ASCEND910B_DTYPE_SUPPORT_LIST = {
43+ op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16, op::DataType::DT_BF16};
44+ 
45+static const inline std::initializer_list<DataType>& GetSupportDtypeList()
46+{
47+ static const std::initializer_list<DataType> emptyDtypes = {};
48+ if (GetCurrentPlatformInfo().GetCurNpuArch() == NpuArch::DAV_2201) {
F
Ffulltower6月4日

GetSupportDtypeList() 对非 DAV_2201 平台返回空 emptyDtypes，导致后续 CheckDtypeValid 中的 OP_CHECK_DTYPE_NOT_SUPPORT 会因列表为空而报错，但错误信息不明确（仅提示 dtype not support，未说明是平台不支持）。建议对不支持平台显式打印 OP_LOGE 说明原因，或使用 IsRegbase() 判断兼容后续平台。

likedislike
H
hqnzw
6月4日 评论：
49+ return ASCEND910B_DTYPE_SUPPORT_LIST;
50+ }
51+ return emptyDtypes;
52+}
53+ 
54+static bool CheckDtypeValid(
55+ const aclTensor* gradOutput, const aclTensor* self, const aclTensor* target, const aclTensor* out)
56+{
57+ const auto& supportList = GetSupportDtypeList();
58+ if (supportList.size() == 0) {
59+ OP_LOGE(
60+ ACLNN_ERR_PARAM_INVALID, "support for npuArch %u is not implemented",
61+ static_cast<uint32_t>(GetCurrentPlatformInfo().GetCurNpuArch()));
62+ return false;
63+ }
64+ 
65+ OP_CHECK_RESULT_DTYPE_CAST_FAILED(self->GetDataType(), out->GetDataType(), return false);
66+ // 检查gradOutput的数据类型是否在支持列表内
67+ OP_CHECK_DTYPE_NOT_SUPPORT(gradOutput, supportList, return false);
68+ // 检查self的数据类型是否在支持列表内
69+ OP_CHECK_DTYPE_NOT_SUPPORT(self, supportList, return false);
70+ // 检查target的数据类型是否在支持列表内
71+ OP_CHECK_DTYPE_NOT_SUPPORT(target, supportList, return false);
72+ // 检查out的数据类型是否在支持列表内
73+ OP_CHECK_DTYPE_NOT_SUPPORT(out, supportList, return false);
74+ 
75+ return true;
76+}
77+ 
78+constexpr size_t MAX_DIM_LEN = 8;
79+ 
80+static bool CheckShape(
81+ const aclTensor* gradOutput, const aclTensor* self, const aclTensor* target, const aclTensor* out)
82+{
83+ OP_CHECK_MAX_DIM(gradOutput, MAX_DIM_LEN, return false);
84+ OP_CHECK_MAX_DIM(self, MAX_DIM_LEN, return false);
85+ OP_CHECK_MAX_DIM(target, MAX_DIM_LEN, return false);
86+ OP_CHECK_MAX_DIM(out, MAX_DIM_LEN, return false);
87+ 
88+ op::Shape broadcastShape;
89+ op::Shape broadcastGradShape;
90+ OP_CHECK_BROADCAST_AND_INFER_SHAPE(self, target, broadcastShape, return false);
91+ if (!BroadcastInferShape(gradOutput->GetViewShape(), broadcastShape, broadcastGradShape) ||
92+ broadcastShape != broadcastGradShape) {
93+ OP_LOGE(
94+ ACLNN_ERR_PARAM_INVALID, "Except shape of gradOutput must broadcast to %s, but current is %s.",
95+ op::ToString(broadcastShape).GetString(), op::ToString(gradOutput->GetViewShape()).GetString());
96+ return false;
97+ }
98+ OP_CHECK_SHAPE_NOT_EQUAL(self, out, return false);
99+ 
100+ return true;
101+}
102+ 
103+static aclnnStatus CheckParams(
104+ const aclTensor* gradOutput, const aclTensor* self, const aclTensor* target, aclTensor* out)
105+{
106+ // 1. 检查参数是否为空指针
107+ CHECK_RET(CheckNotNull4Tensor(gradOutput, self, target, out), ACLNN_ERR_PARAM_NULLPTR);
F
Ffulltower5月29日
  1. 注释序号错误且描述有误："4. 检查输出输出shape" 应为 "3. 检查输入输出shape"（序号跳跃，"输出输出" 应为 "输入输出"）。
likedislike
108+ 
109+ // 2. 检查输入的数据类型是否在API支持的数据类型范围之内，需要根据api定义校验
110+ CHECK_RET(CheckDtypeValid(gradOutput, self, target, out), ACLNN_ERR_PARAM_INVALID);
111+ 
112+ // 3. 检查输入输出shape
113+ CHECK_RET(CheckShape(gradOutput, self, target, out), ACLNN_ERR_PARAM_INVALID);
114+ 
115+ return ACLNN_SUCCESS;
116+}
117+ 
118+static const char* GetReductionStr(int64_t reduction)
119+{
120+ if (reduction == None) {
121+ return REDUCTION_NONE;
122+ } else if (reduction == Mean) {
123+ return REDUCTION_MEAN;
124+ } else if (reduction == Sum) {
125+ return REDUCTION_SUM;
126+ } else if (reduction == Batchmean) {
127+ return REDUCTION_BATCHMEAN;
128+ } else {
129+ return REDUCTION_NONE;
130+ }
131+}
132+ 
133+static const aclTensor* BroadcastTensor(const aclTensor* self, const op::Shape broadcastShape, aclOpExecutor* executor)
134+{
135+ // 如果self的shape与broadcast的不一致，进行BroadcastTo
136+ if (self->GetViewShape() != broadcastShape) {
137+ auto broadcastShapeIntArray = GetBroadcastShapeLossBackward(broadcastShape, executor);
138+ if (broadcastShapeIntArray != nullptr) {
139+ return l0op::BroadcastTo(self, broadcastShapeIntArray, executor);
140+ }
141+ }
142+ return self;
143+}
144+ 
145+static const aclTensor* ReduceSumTensor(const aclTensor* grad, const op::Shape outShape, aclOpExecutor* executor)
146+{
147+ // 如果grad的shape与outShape不一致，进行ReduceSum
148+ if (grad->GetViewShape() != outShape) {
149+ size_t outDimNum = outShape.GetDimNum();
150+ size_t gradDimNum = grad->GetViewShape().GetDimNum();
151+ size_t startDim = gradDimNum - outDimNum;
152+ size_t dimIdx = startDim;
153+ std::vector<int64_t> appendDim;
154+ for (size_t i = 0; i < startDim; ++i) {
155+ appendDim.push_back(i);
156+ }
157+ for (size_t j = startDim; j < gradDimNum; ++j) {
158+ if (outShape[j - startDim] != (grad->GetViewShape())[j]) {
159+ appendDim.push_back(j);
160+ dimIdx++;
161+ }
162+ }
163+ auto axes = executor->AllocIntArray(appendDim.data(), dimIdx);
164+ auto out = l0op::ReduceSumOp(grad, axes, true, executor);
165+ CHECK_RET(out != nullptr, nullptr);
166+ auto outShapeIntArray = GetBroadcastShapeLossBackward(outShape, executor);
167+ return l0op::Reshape(out, outShapeIntArray, executor);
168+ }
169+ return grad;
170+}
171+ 
172+aclnnStatus aclnnKlDivBackwardGetWorkspaceSize(
173+ const aclTensor* gradOutput, const aclTensor* self, const aclTensor* target, int64_t reduction, bool logTarget,
174+ aclTensor* out, uint64_t* workspaceSize, aclOpExecutor** executor)
175+{
176+ OP_CHECK_COMM_INPUT(workspaceSize, executor);
177+ 
178+ L2_DFX_PHASE_1(aclnnKlDivBackward, DFX_IN(gradOutput, self, target, reduction, logTarget), DFX_OUT(out));
179+ // 固定写法，创建OpExecutor
180+ auto uniqueExecutor = CREATE_EXECUTOR();
181+ CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR);
182+ 
183+ // 固定写法，参数检查
184+ auto ret = CheckParams(gradOutput, self, target, out);
185+ CHECK_RET(ret == ACLNN_SUCCESS, ret);
186+ 
187+ if (self->IsEmpty()) {
188+ // 根据实际支持情况补充
189+ *workspaceSize = 0;
190+ uniqueExecutor.ReleaseTo(executor);
191+ return ACLNN_SUCCESS;
192+ }
193+ 
194+ // 固定写法，将输入gradOutput转换成连续的tensor
195+ auto gradOutputContiguous = l0op::Contiguous(gradOutput, uniqueExecutor.get());
196+ CHECK_RET(gradOutputContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
197+ 
198+ // 固定写法，将输入self转换成连续的tensor
199+ auto selfContiguous = l0op::Contiguous(self, uniqueExecutor.get());
200+ CHECK_RET(selfContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
201+ 
202+ // 固定写法，将输入target转换成连续的tensor
203+ auto targetContiguous = l0op::Contiguous(target, uniqueExecutor.get());
204+ CHECK_RET(targetContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
205+ 
206+ auto gradOutputCasted = gradOutputContiguous;
207+ auto selfCasted = selfContiguous;
208+ auto targetCasted = targetContiguous;
209+ if (!(gradOutput->GetDataType() == op::DataType::DT_FLOAT16 && self->GetDataType() == op::DataType::DT_FLOAT16 &&
210+ target->GetDataType() == op::DataType::DT_FLOAT16)) {
211+ auto promoteType = op::DataType::DT_FLOAT;
212+ 
213+ // 将输入gradoutput的数据类型转换成隐式数据类型，根据具体算子语义按需调用
214+ gradOutputCasted = l0op::Cast(gradOutputContiguous, promoteType, uniqueExecutor.get());
215+ CHECK_RET(gradOutputCasted != nullptr, ACLNN_ERR_INNER_NULLPTR);
216+ 
217+ // 将输入self的数据类型转换成隐式数据类型，根据具体算子语义按需调用
218+ selfCasted = l0op::Cast(selfContiguous, promoteType, uniqueExecutor.get());
219+ CHECK_RET(selfCasted != nullptr, ACLNN_ERR_INNER_NULLPTR);
220+ 
221+ // 将输入target的数据类型转换成隐式数据类型，根据具体算子语义按需调用
222+ targetCasted = l0op::Cast(targetContiguous, promoteType, uniqueExecutor.get());
223+ CHECK_RET(targetCasted != nullptr, ACLNN_ERR_INNER_NULLPTR);
224+ }
225+ 
226+ op::Shape broadcastShape;
227+ BroadcastInferShape(target->GetViewShape(), self->GetViewShape(), broadcastShape);
228+ 
229+ // 判断self是否需要进行broadcast
230+ auto selfBroadcast = BroadcastTensor(selfCasted, broadcastShape, uniqueExecutor.get());
231+ CHECK_RET(selfBroadcast != nullptr, ACLNN_ERR_INNER_NULLPTR);
232+ 
233+ // 进行计算
234+ auto grad = l0op::KlDivLossGrad(
235+ gradOutputCasted, selfBroadcast, targetCasted, GetReductionStr(reduction), logTarget, uniqueExecutor.get());
236+ CHECK_RET(grad != nullptr, ACLNN_ERR_INNER_NULLPTR);
237+ 
238+ // 根据grad的shape是否与out的shape相同，判断是否需要reduce
239+ auto gradReduce = ReduceSumTensor(grad, out->GetViewShape(), uniqueExecutor.get());
240+ CHECK_RET(gradReduce != nullptr, ACLNN_ERR_INNER_NULLPTR);
241+ 
242+ // 固定写法，将计算结果转换成输出out的数据类型
243+ auto castOut = l0op::Cast(gradReduce, out->GetDataType(), uniqueExecutor.get());
244+ CHECK_RET(castOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
245+ 
246+ // 固定写法，将计算结果拷贝到输出out上，out可能是非连续的tensor
247+ auto viewCopyResult = l0op::ViewCopy(castOut, out, uniqueExecutor.get());
248+ CHECK_RET(viewCopyResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
249+ 
250+ // 固定写法，获取计算过程中需要使用的workspace大小
251+ *workspaceSize = uniqueExecutor->GetWorkspaceSize();
252+ // 需要把 uniqueExecutor持有executor转移给executor
253+ uniqueExecutor.ReleaseTo(executor);
254+ return ACLNN_SUCCESS;
255+}
256+ 
257+aclnnStatus aclnnKlDivBackward(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream)
258+{
259+ L2_DFX_PHASE_2(aclnnKlDivBackward);
260+ // 固定写法，调用框架能力，完成计算
261+ return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);
262+}
263+ 
264+#ifdef __cplusplus
265+}
266+#endif
Aexperimental/loss/kl_div_loss_grad/op_host/op_api/aclnn_kl_div_backward.h+39-0
@@ -0,0 +1,39 @@
1+/**
2+ * Copyright (c) 2026 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 OP_API_INC_KL_DIV_BACKWARD_H_
12+#define OP_API_INC_KL_DIV_BACKWARD_H_
13+ 
14+#include "aclnn/aclnn_base.h"
15+#include "aclnn_util.h"
16+ 
17+#ifdef __cplusplus
18+extern "C" {
19+#endif
20+ 
21+/**
22+ * @brief aclnnKlDivBackward的第一段接口，根据具体的计算流程，计算workspace大小。
23+ * @domain aclnn_ops_train
24+ */
25+ACLNN_API aclnnStatus aclnnKlDivBackwardGetWorkspaceSize(
26+ const aclTensor* gradOutput, const aclTensor* self, const aclTensor* target, int64_t reduction, bool logTarget,
27+ aclTensor* out, uint64_t* workspaceSize, aclOpExecutor** executor);
28+ 
29+/**
30+ * @brief aclnnKlDivBackward的第二段接口，用于执行计算。
31+ */
32+ACLNN_API aclnnStatus
33+aclnnKlDivBackward(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);
34+ 
35+#ifdef __cplusplus
36+}
37+#endif
38+ 
39+#endif // OP_API_INC_KL_DIV_BACKWARD_H_
Aexperimental/loss/kl_div_loss_grad/op_host/op_api/kl_div_loss_grad.cpp+66-0
@@ -0,0 +1,66 @@
1+/**
2+ * Copyright (c) 2026 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 kl_div_loss_grad.cpp
13+ * \brief
14+ */
15+ 
16+#include "kl_div_loss_grad.h"
17+#include "opdev/make_op_executor.h"
18+#include "opdev/aicpu/aicpu_task.h"
19+#include "opdev/op_def.h"
20+#include "opdev/op_dfx.h"
21+#include "opdev/op_executor.h"
22+#include "opdev/op_log.h"
23+#include "opdev/shape_utils.h"
24+ 
25+using namespace op;
26+ 
27+namespace l0op {
28+OP_TYPE_REGISTER(KlDivLossGrad);
29+ 
30+static const std::initializer_list<op::DataType> AICORE_DTYPE_SUPPORT_LIST = {
31+ op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16, op::DataType::DT_BF16};
32+// 根据芯片类型、dtype判断算子是否支持走aicore
33+static bool IsAiCoreSupport(const aclTensor* self)
34+{
35+ // KlDivLossGrad只需要判断dtype
36+ return CheckType(self->GetDataType(), AICORE_DTYPE_SUPPORT_LIST);
37+}
38+ 
39+// AICORE算子kernel
40+static const aclTensor* KlDivLossGradAiCore(
41+ const aclTensor* gradOutput, const aclTensor* self, const aclTensor* target, const char* reduction, bool logTarget,
42+ aclTensor* out, aclOpExecutor* executor)
43+{
44+ L0_DFX(KlDivLossGradAiCore, gradOutput, self, target, reduction, logTarget);
45+ auto ret = ADD_TO_LAUNCHER_LIST_AICORE(
46+ KlDivLossGrad, OP_INPUT(gradOutput, self, target), OP_OUTPUT(out), OP_ATTR(reduction, logTarget));
47+ OP_CHECK(
48+ ret == ACLNN_SUCCESS,
49+ OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "KlDivLossGradAiCore ADD_TO_LAUNCHER_LIST_AICORE failed."), return nullptr);
50+ 
51+ return out;
52+}
53+ 
54+const aclTensor* KlDivLossGrad(
55+ const aclTensor* gradOutput, const aclTensor* self, const aclTensor* target, const char* reduction, bool logTarget,
56+ aclOpExecutor* executor)
57+{
58+ auto out = executor->AllocTensor(self->GetViewShape(), self->GetDataType(), self->GetStorageFormat());
59+ if (IsAiCoreSupport(self)) {
60+ // 只走aicore
61+ return KlDivLossGradAiCore(gradOutput, self, target, reduction, logTarget, out, executor);
62+ }
63+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "KlDivLossGrad not supported: dtype=%d.", static_cast<int>(self->GetDataType()));
64+ return nullptr;
65+}
66+} // namespace l0op
Aexperimental/loss/kl_div_loss_grad/op_host/op_api/kl_div_loss_grad.h+26-0
@@ -0,0 +1,26 @@
1+/**
2+ * Copyright (c) 2026 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 kl_div_loss_grad.h
13+ * \brief
14+ */
15+#ifndef PTA_NPU_OP_API_INC_LEVEL0_OP_KL_DIV_LOSS_GRAD_OP_H_
16+#define PTA_NPU_OP_API_INC_LEVEL0_OP_KL_DIV_LOSS_GRAD_OP_H_
17+ 
18+#include "opdev/op_executor.h"
19+ 
20+namespace l0op {
21+const aclTensor* KlDivLossGrad(
22+ const aclTensor* gradOutput, const aclTensor* self, const aclTensor* target, const char* reduction, bool logTarget,
23+ aclOpExecutor* executor);
24+}
25+ 
26+#endif // PTA_NPU_OP_API_INC_LEVEL0_OP_KL_DIV_LOSS_GRAD_OP_H_
Aexperimental/loss/kl_div_loss_grad/op_kernel/kl_div_loss_grad.cpp+59-0
@@ -0,0 +1,59 @@
1+/**
2+ * Copyright (c) 2026 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 kl_div_loss_grad.cpp
13+ * \brief
14+ */
15+ 
16+#include "kl_div_loss_grad.h"
17+ 
18+using namespace NsKlDivLossGrad;
19+ 
20+template <bool logTarget, bool broadcast>
21+struct PolicySelector;
22+ 
23+template <>
24+struct PolicySelector<false, false> {
25+ template <class T>
26+ using type = ComputeImpl<T>;
27+};
28+ 
29+template <>
30+struct PolicySelector<true, false> {
31+ template <class T>
32+ using type = ComputeImplLog<T>;
33+};
34+ 
35+template <>
36+struct PolicySelector<false, true> {
37+ template <class T>
38+ using type = ComputeImplBroadCast<T>;
39+};
40+ 
41+template <>
42+struct PolicySelector<true, true> {
43+ template <class T>
44+ using type = ComputeImplBroadCastLog<T>;
45+};
46+ 
47+template <bool logTarget, bool broadcast>
48+__global__ __aicore__ void kl_div_loss_grad(
49+ GM_ADDR grad, GM_ADDR input, GM_ADDR target, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
50+{
51+ REGISTER_TILING_DEFAULT(KlDivLossGradTilingData);
52+ GET_TILING_DATA_WITH_STRUCT(KlDivLossGradTilingData, tilingData, tiling);
53+ KernelKlDivLossGrad<DTYPE_GRAD, PolicySelector<logTarget, broadcast>::template type> op;
54+ AscendC::TPipe pipe;
55+ op.Init(
56+ grad, input, target, y, tilingData.bigCoreDataNum, tilingData.smallCoreDataNum, tilingData.tileDataNum,
57+ tilingData.bigCoreNum, tilingData.coff, &pipe);
58+ op.Process();
59+}
Aexperimental/loss/kl_div_loss_grad/op_kernel/kl_div_loss_grad.h+214-0
@@ -0,0 +1,214 @@
1+/**
2+ * Copyright (c) 2026 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 kl_div_loss_grad.h
13+ * \brief
14+ */
15+#ifndef __KL_DIV_LOSS_GRAD_H__
16+#define __KL_DIV_LOSS_GRAD_H__
17+ 
18+#include "kernel_operator.h"
19+#include "kernel_tiling/kernel_tiling.h"
20+#include "kl_div_loss_grad_tiling_data.h"
21+#include "kl_div_loss_grad_tiling_key.h"
22+ 
23+namespace NsKlDivLossGrad {
24+template <class T, template <class> class Policy>
25+class KernelKlDivLossGrad {
26+ using PolicyType = Policy<float>;
27+ 
28+public:
29+ __aicore__ inline KernelKlDivLossGrad() {}
30+ __aicore__ inline void Init(
31+ GM_ADDR grad, GM_ADDR input, GM_ADDR target, GM_ADDR y, uint32_t bigCoreDataNum, uint32_t smallCoreDataNum,
32+ uint32_t tileDataNum, uint32_t bigCoreNum, float coeff, AscendC::TPipe* pipe)
33+ {
34+ uint32_t globalBufferIndex = bigCoreDataNum * AscendC::GetBlockIdx();
35+ if (AscendC::GetBlockIdx() < bigCoreNum) {
36+ this->coreDataNum = bigCoreDataNum;
37+ } else {
38+ this->coreDataNum = smallCoreDataNum;
39+ globalBufferIndex -= (bigCoreDataNum - smallCoreDataNum) * (AscendC::GetBlockIdx() - bigCoreNum);
40+ }
41+ if constexpr (!PolicyType::broadcast) {
42+ gradGm.SetGlobalBuffer((__gm__ T*)grad + globalBufferIndex, this->coreDataNum);
43+ } else {
44+ gradGm.SetGlobalBuffer((__gm__ T*)grad, 1);
45+ }
46+ targetGm.SetGlobalBuffer((__gm__ T*)target + globalBufferIndex, this->coreDataNum);
47+ yGm.SetGlobalBuffer((__gm__ T*)y + globalBufferIndex, this->coreDataNum);
48+ 
49+ if constexpr (!PolicyType::broadcast) {
50+ pipe->InitBuffer(inQueGrad, 2, tileDataNum * sizeof(T));
51+ if constexpr (!std::is_same_v<T, float>) {
52+ pipe->InitBuffer(gradBuf, tileDataNum * sizeof(float));
53+ }
54+ }
55+ pipe->InitBuffer(inQueTarget, 2, tileDataNum * sizeof(T));
56+ if constexpr (!std::is_same_v<T, float>) {
57+ pipe->InitBuffer(targetBuf, tileDataNum * sizeof(float));
58+ }
59+ pipe->InitBuffer(outQueY, 2, tileDataNum * sizeof(T));
60+ if constexpr (!std::is_same_v<T, float>) {
61+ pipe->InitBuffer(yBuf, tileDataNum * sizeof(float));
62+ }
63+ 
64+ this->tileDataNum = tileDataNum;
65+ if constexpr (PolicyType::broadcast) {
66+ if constexpr (std::is_same_v<T, bfloat16_t>) {
67+ policy.grad = AscendC::ToFloat(gradGm.GetValue(0));
68+ } else {
69+ policy.grad = static_cast<float>(gradGm.GetValue(0));
70+ }
71+ }
72+ policy.negCoeff = -coeff;
73+ }
74+ __aicore__ inline void Process()
75+ {
76+ uint64_t coreDataNum = this->coreDataNum;
77+ uint64_t tileDataNum = this->tileDataNum;
78+ for (uint64_t offset = 0; offset < coreDataNum; offset += tileDataNum) {
79+ uint32_t processDataNum = AscendC::Std::min(tileDataNum, coreDataNum - offset);
80+ CopyIn(offset, processDataNum);
81+ Compute(processDataNum);
82+ CopyOut(offset, processDataNum);
83+ }
84+ }
85+ 
86+private:
87+ __aicore__ inline void CopyIn(uint32_t offset, uint32_t processDataNum)
88+ {
89+ AscendC::LocalTensor<T> targetLocal = inQueTarget.AllocTensor<T>();
90+ AscendC::DataCopy(targetLocal, targetGm[offset], processDataNum);
91+ inQueTarget.EnQue(targetLocal);
92+ if constexpr (!PolicyType::broadcast) {
93+ AscendC::LocalTensor<T> gradLocal = inQueGrad.AllocTensor<T>();
94+ AscendC::DataCopy(gradLocal, gradGm[offset], processDataNum);
95+ inQueGrad.EnQue(gradLocal);
96+ }
97+ }
98+ __aicore__ inline void Compute(uint32_t processDataNum)
99+ {
100+ AscendC::LocalTensor<T> yLocal = outQueY.AllocTensor<T>();
101+ AscendC::LocalTensor<T> targetLocal = inQueTarget.DeQue<T>();
102+ if constexpr (!std::is_same_v<T, float>) {
103+ AscendC::LocalTensor<float> targetFp32 = targetBuf.Get<float>();
104+ AscendC::LocalTensor<float> yFp32 = yBuf.Get<float>();
105+ AscendC::Cast(targetFp32, targetLocal, AscendC::RoundMode::CAST_NONE, processDataNum);
106+ if constexpr (!PolicyType::broadcast) {
107+ AscendC::LocalTensor<T> gradLocal = inQueGrad.DeQue<T>();
108+ AscendC::LocalTensor<float> gradFp32 = gradBuf.Get<float>();
109+ AscendC::Cast(gradFp32, gradLocal, AscendC::RoundMode::CAST_NONE, processDataNum);
110+ inQueGrad.FreeTensor(gradLocal);
111+ policy.compute(gradFp32, targetFp32, yFp32, processDataNum);
112+ } else {
113+ policy.compute(targetFp32, yFp32, processDataNum);
114+ }
115+ AscendC::Cast(yLocal, yFp32, AscendC::RoundMode::CAST_RINT, processDataNum);
116+ } else {
117+ if constexpr (!PolicyType::broadcast) {
118+ AscendC::LocalTensor<T> gradLocal = inQueGrad.DeQue<T>();
119+ policy.compute(gradLocal, targetLocal, yLocal, processDataNum);
120+ inQueGrad.FreeTensor(gradLocal);
121+ } else {
122+ policy.compute(targetLocal, yLocal, processDataNum);
123+ }
124+ }
125+ outQueY.EnQue(yLocal);
126+ inQueTarget.FreeTensor(targetLocal);
127+ }
128+ __aicore__ inline void CopyOut(uint32_t offset, uint32_t processDataNum)
129+ {
130+ AscendC::LocalTensor<T> yLocal = outQueY.DeQue<T>();
131+ AscendC::DataCopy(yGm[offset], yLocal, processDataNum);
132+ outQueY.FreeTensor(yLocal);
133+ }
134+ 
135+private:
136+ AscendC::TQue<AscendC::TPosition::VECIN, 1> inQueGrad, inQueTarget;
137+ AscendC::TQue<AscendC::TPosition::VECOUT, 1> outQueY;
138+ AscendC::TBuf<AscendC::TPosition::VECCALC> gradBuf, targetBuf, yBuf;
139+ AscendC::GlobalTensor<T> gradGm, targetGm;
140+ AscendC::GlobalTensor<T> yGm;
141+ 
142+ uint32_t coreDataNum;
143+ uint32_t tileDataNum;
144+ 
145+ PolicyType policy;
146+};
147+ 
148+template <class T>
149+class ComputeImpl {
150+public:
151+ static constexpr bool broadcast = false;
152+ 
153+ __aicore__ inline void compute(
154+ const AscendC::LocalTensor<T>& grad, const AscendC::LocalTensor<T>& target, const AscendC::LocalTensor<T>& y,
155+ uint32_t processDataNum) const
156+ {
157+ AscendC::Mul(y, grad, target, processDataNum);
158+ AscendC::Muls(y, y, negCoeff, processDataNum);
159+ }
160+ 
161+ T negCoeff;
162+};
163+ 
164+template <class T>
165+class ComputeImplLog {
166+public:
167+ static constexpr bool broadcast = false;
168+ 
169+ __aicore__ inline void compute(
170+ const AscendC::LocalTensor<T>& grad, const AscendC::LocalTensor<T>& target, const AscendC::LocalTensor<T>& y,
171+ uint32_t processDataNum) const
172+ {
173+ AscendC::Exp(y, target, processDataNum);
174+ AscendC::Mul(target, y, grad, processDataNum);
175+ AscendC::Muls(y, target, negCoeff, processDataNum);
176+ }
177+ 
178+ T negCoeff;
179+};
180+ 
181+template <class T>
182+class ComputeImplBroadCast {
183+public:
184+ static constexpr bool broadcast = true;
185+ 
186+ __aicore__ inline void compute(
187+ const AscendC::LocalTensor<T>& target, const AscendC::LocalTensor<T>& y, uint32_t processDataNum) const
188+ {
189+ AscendC::Muls(y, target, grad, processDataNum);
190+ AscendC::Muls(y, y, negCoeff, processDataNum);
191+ }
192+ 
193+ T grad;
194+ T negCoeff;
195+};
196+ 
197+template <class T>
198+class ComputeImplBroadCastLog {
199+public:
200+ static constexpr bool broadcast = true;
201+ 
202+ __aicore__ inline void compute(
203+ const AscendC::LocalTensor<T>& target, const AscendC::LocalTensor<T>& y, uint32_t processDataNum) const
204+ {
205+ AscendC::Exp(y, target, processDataNum);
206+ AscendC::Muls(target, y, grad, processDataNum);
207+ AscendC::Muls(y, target, negCoeff, processDataNum);
208+ }
209+ 
210+ T grad;
211+ T negCoeff;
212+};
213+} // namespace NsKlDivLossGrad
214+#endif // KL_DIV_LOSS_GRAD_H
Aexperimental/loss/kl_div_loss_grad/op_kernel/kl_div_loss_grad_tiling_data.h+27-0
@@ -0,0 +1,27 @@
1+/**
2+ * Copyright (c) 2026 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 kl_div_loss_grad_tiling_data.h
13+ * \brief tiling data struct
14+ */
15+ 
16+#ifndef __KL_DIV_LOSS_GRAD_TILLING_DATA_H__
17+#define __KL_DIV_LOSS_GRAD_TILLING_DATA_H__
18+ 
19+struct KlDivLossGradTilingData {
20+ uint32_t bigCoreDataNum;
21+ uint32_t smallCoreDataNum;
22+ uint32_t tileDataNum;
23+ uint32_t bigCoreNum;
24+ float coff;
25+};
26+ 
27+#endif
Aexperimental/loss/kl_div_loss_grad/op_kernel/kl_div_loss_grad_tiling_key.h+27-0
@@ -0,0 +1,27 @@
1+/**
2+ * Copyright (c) 2026 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 kl_div_loss_grad_tiling_key.h
13+ * \brief kl_div_loss_grad tiling key declare
14+ */
15+ 
16+#ifndef __KL_DIV_LOSS_GRAD_TILING_KEY_H__
17+#define __KL_DIV_LOSS_GRAD_TILING_KEY_H__
18+ 
19+#include "ascendc/host_api/tiling/template_argument.h"
20+ 
21+/* 模板参数 */
22+ASCENDC_TPL_ARGS_DECL(klDivLossGrad, ASCENDC_TPL_BOOL_DECL(logTarget, 0, 1), ASCENDC_TPL_BOOL_DECL(broadcast, 0, 1), );
23+ 
24+/* 模板参数组合 */
25+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_BOOL_SEL(logTarget, 0, 1), ASCENDC_TPL_BOOL_SEL(broadcast, 0, 1)), );
26+ 
27+#endif
Aexperimental/loss/kl_div_loss_grad/tests/CMakeLists.txt+18-0
@@ -0,0 +1,18 @@
1+ # ----------------------------------------------------------------------------
2+ # Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ # This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ # CANN Open Software License Agreement Version 2.0 (the "License").
5+ # Please refer to the License for details. You may not use this file except in compliance with the License.
6+ # THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ # INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ # See LICENSE in the root of the software repository for the full text of the License.
9+ # ----------------------------------------------------------------------------
10+ 
11+message(STATUS "=== Debug: start ops.loss.kl_div_loss_grad.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()
Aexperimental/loss/kl_div_loss_grad/tests/ut/CMakeLists.txt+18-0
@@ -0,0 +1,18 @@
1+ # ----------------------------------------------------------------------------
2+ # Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ # This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ # CANN Open Software License Agreement Version 2.0 (the "License").
5+ # Please refer to the License for details. You may not use this file except in compliance with the License.
6+ # THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ # INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ # See LICENSE in the root of the software repository for the full text of the License.
9+ # ----------------------------------------------------------------------------
10+ 
11+message(STATUS "=== Debug: start ops.loss.kl_div_loss_grad.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()
Aexperimental/loss/kl_div_loss_grad/tests/ut/op_api/CMakeLists.txt+12-0
@@ -0,0 +1,12 @@
1+# ----------------------------------------------------------------------------
2+ # Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ # This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ # CANN Open Software License Agreement Version 2.0 (the "License").
5+ # Please refer to the License for details. You may not use this file except in compliance with the License.
6+ # THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ # INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ # See LICENSE in the root of the software repository for the full text of the License.
9+ # ----------------------------------------------------------------------------
10+ 
11+message(STATUS "=== Debug: target_sources add test_kl_div_loss_grad")
12+add_modules_llt_sources(HOSTNAME ${OPTEST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
Aexperimental/loss/kl_div_loss_grad/tests/ut/op_api/test_aclnn_kl_div_backward.cpp+493-0
@@ -0,0 +1,493 @@
1+/**
2+ * Copyright (c) 2026 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 <array>
12+#include <vector>
13+#include "gtest/gtest.h"
14+#include "../../../../op_host/op_api/aclnn_kl_div_backward.h"
15+#include "op_api_ut_common/op_api_ut.h"
16+#include "op_api_ut_common/scalar_desc.h"
17+#include "op_api_ut_common/tensor_desc.h"
18+ 
19+#include <unistd.h>
20+ 
21+using namespace op;
22+using namespace std;
23+ 
24+enum Reduction { None, Mean, Sum, BatchMean, END };
25+ 
26+class kl_div_loss_backward_test : public testing::Test {
27+protected:
28+ static void SetUpTestCase() { std::cout << "kl_div_loss_backward_test SetUp" << std::endl; }
29+ 
30+ static void TearDownTestCase() { std::cout << "kl_div_loss_backward_test TearDown" << std::endl; }
31+};
32+ 
33+TEST_F(kl_div_loss_backward_test, aclnnKlDivBackward_01_float_none_broadcast)
34+{
35+ auto gradOutputDesc = TensorDesc({3, 1}, ACL_FLOAT, ACL_FORMAT_ND);
36+ auto selfDesc = TensorDesc({3, 5}, ACL_FLOAT, ACL_FORMAT_ND);
37+ auto targetDesc = TensorDesc({3, 5}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(0, 4);
38+ int64_t reduction = Reduction::None;
39+ bool logTarget = false;
40+ 
41+ auto outDesc = TensorDesc({3, 5}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.0001, 0.0001);
42+ 
43+ auto ut = OP_API_UT(
44+ aclnnKlDivBackward, INPUT(gradOutputDesc, selfDesc, targetDesc, reduction, logTarget), OUTPUT(outDesc));
45+ 
46+ uint64_t workspace_size = 0;
47+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
48+ EXPECT_EQ(aclRet, ACL_SUCCESS);
49+ 
50+ ut.TestPrecision();
51+}
52+ 
53+TEST_F(kl_div_loss_backward_test, aclnnKlDivBackward_02_float_nchw_mean)
54+{
55+ auto gradOutputDesc = TensorDesc({3, 5, 2, 4}, ACL_FLOAT, ACL_FORMAT_NCHW);
56+ auto selfDesc = TensorDesc({3, 5, 2, 4}, ACL_FLOAT, ACL_FORMAT_NCHW).ValueRange(-1, 1);
57+ auto targetDesc = TensorDesc({3, 5, 2, 4}, ACL_FLOAT, ACL_FORMAT_NCHW).ValueRange(-1, 1);
58+ int64_t reduction = Reduction::Mean;
59+ bool logTarget = false;
60+ 
61+ auto outDesc = TensorDesc({3, 5, 2, 4}, ACL_FLOAT, ACL_FORMAT_NCHW).Precision(0.0001, 0.0001);
62+ 
63+ auto ut = OP_API_UT(
64+ aclnnKlDivBackward, INPUT(gradOutputDesc, selfDesc, targetDesc, reduction, logTarget), OUTPUT(outDesc));
65+ 
66+ uint64_t workspace_size = 0;
67+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
68+ EXPECT_EQ(aclRet, ACL_SUCCESS);
69+ 
70+ ut.TestPrecision();
71+}
72+ 
73+TEST_F(kl_div_loss_backward_test, aclnnKlDivBackward_03_float16_ncdhw_sum)
74+{
75+ auto gradOutputDesc = TensorDesc({1, 2, 3, 4, 5}, ACL_FLOAT16, ACL_FORMAT_NCDHW);
76+ auto selfDesc = TensorDesc({1, 2, 3, 4, 5}, ACL_FLOAT16, ACL_FORMAT_NCDHW).ValueRange(-1, 1);
77+ auto targetDesc = TensorDesc({1, 2, 3, 4, 5}, ACL_FLOAT16, ACL_FORMAT_NCDHW).ValueRange(-1, 1);
78+ int64_t reduction = Reduction::Sum;
79+ bool logTarget = false;
80+ 
81+ auto outDesc = TensorDesc({1, 2, 3, 4, 5}, ACL_FLOAT16, ACL_FORMAT_NCDHW).Precision(0.0001, 0.0001);
82+ 
83+ auto ut = OP_API_UT(
84+ aclnnKlDivBackward, INPUT(gradOutputDesc, selfDesc, targetDesc, reduction, logTarget), OUTPUT(outDesc));
85+ 
86+ uint64_t workspace_size = 0;
87+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
88+ EXPECT_EQ(aclRet, ACL_SUCCESS);
89+ 
90+ ut.TestPrecision();
91+}
92+ 
93+TEST_F(kl_div_loss_backward_test, aclnnKlDivBackward_04_float_hwcn_batchmean)
94+{
95+ auto gradOutputDesc = TensorDesc({3, 5, 4, 6}, ACL_FLOAT, ACL_FORMAT_HWCN);
96+ auto selfDesc = TensorDesc({3, 5, 4, 6}, ACL_FLOAT, ACL_FORMAT_HWCN).ValueRange(-1, 1);
97+ auto targetDesc = TensorDesc({3, 5, 4, 6}, ACL_FLOAT, ACL_FORMAT_HWCN).ValueRange(-1, 1);
98+ int64_t reduction = Reduction::BatchMean;
99+ bool logTarget = false;
100+ 
101+ auto outDesc = TensorDesc({3, 5, 4, 6}, ACL_FLOAT, ACL_FORMAT_HWCN).Precision(0.0001, 0.0001);
102+ 
103+ auto ut = OP_API_UT(
104+ aclnnKlDivBackward, INPUT(gradOutputDesc, selfDesc, targetDesc, reduction, logTarget), OUTPUT(outDesc));
105+ 
106+ uint64_t workspace_size = 0;
107+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
108+ EXPECT_EQ(aclRet, ACL_SUCCESS);
109+ 
110+ ut.TestPrecision();
111+}
112+ 
113+TEST_F(kl_div_loss_backward_test, aclnnKlDivBackward_05_float_none_logtarget_true)
114+{
115+ auto gradOutputDesc = TensorDesc({3, 5, 4, 6}, ACL_FLOAT, ACL_FORMAT_ND);
116+ auto selfDesc = TensorDesc({3, 5, 4, 6}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
117+ auto targetDesc = TensorDesc({3, 5, 4, 6}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
118+ int64_t reduction = Reduction::None;
119+ bool logTarget = true;
120+ 
121+ auto outDesc = TensorDesc({3, 5, 4, 6}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.0001, 0.0001);
122+ 
123+ auto ut = OP_API_UT(
124+ aclnnKlDivBackward, INPUT(gradOutputDesc, selfDesc, targetDesc, reduction, logTarget), OUTPUT(outDesc));
125+ 
126+ uint64_t workspace_size = 0;
127+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
128+ EXPECT_EQ(aclRet, ACL_SUCCESS);
129+ 
130+ ut.TestPrecision();
131+}
132+ 
133+TEST_F(kl_div_loss_backward_test, aclnnKlDivBackward_06_float16_to_float_logtarget_true)
134+{
135+ auto gradOutputDesc = TensorDesc({3, 5, 4, 6}, ACL_FLOAT16, ACL_FORMAT_ND);
136+ auto selfDesc = TensorDesc({3, 5, 4, 6}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
137+ auto targetDesc = TensorDesc({3, 5, 4, 6}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-1, 1);
138+ int64_t reduction = Reduction::None;
139+ bool logTarget = true;
140+ 
141+ auto outDesc = TensorDesc({3, 5, 4, 6}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.0001, 0.0001);
142+ 
143+ auto ut = OP_API_UT(
144+ aclnnKlDivBackward, INPUT(gradOutputDesc, selfDesc, targetDesc, reduction, logTarget), OUTPUT(outDesc));
145+ 
146+ uint64_t workspace_size = 0;
147+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
148+ EXPECT_EQ(aclRet, ACL_SUCCESS);
149+ 
150+ ut.TestPrecision();
151+}
152+ 
153+TEST_F(kl_div_loss_backward_test, aclnnKlDivBackward_07_float_nhwc_logtarget_true)
154+{
155+ auto gradOutputDesc = TensorDesc({3, 1, 2, 5}, ACL_FLOAT, ACL_FORMAT_NHWC).ValueRange(-1, 1);
156+ auto selfDesc = TensorDesc({3, 1, 2, 5}, ACL_FLOAT, ACL_FORMAT_NHWC).ValueRange(-1, 1);
157+ auto targetDesc = TensorDesc({3, 1, 2, 5}, ACL_FLOAT, ACL_FORMAT_NHWC).ValueRange(-1, 1);
158+ int64_t reduction = Reduction::None;
159+ bool logTarget = true;
160+ 
161+ auto outDesc = TensorDesc({3, 1, 2, 5}, ACL_FLOAT, ACL_FORMAT_NHWC).Precision(0.0001, 0.0001);
162+ 
163+ auto ut = OP_API_UT(
164+ aclnnKlDivBackward, INPUT(gradOutputDesc, selfDesc, targetDesc, reduction, logTarget), OUTPUT(outDesc));
165+ 
166+ uint64_t workspace_size = 0;
167+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
168+ EXPECT_EQ(aclRet, ACL_SUCCESS);
169+ 
170+ ut.TestPrecision();
171+}
172+ 
173+TEST_F(kl_div_loss_backward_test, aclnnKlDivBackward_08_float_ndhwc_logtarget_true)
174+{
175+ auto gradOutputDesc = TensorDesc({3, 1, 2, 5, 4}, ACL_FLOAT, ACL_FORMAT_NDHWC).ValueRange(-1, 1);
176+ auto selfDesc = TensorDesc({3, 1, 2, 5, 4}, ACL_FLOAT, ACL_FORMAT_NDHWC).ValueRange(-1, 1);
177+ auto targetDesc = TensorDesc({3, 1, 2, 5, 4}, ACL_FLOAT, ACL_FORMAT_NDHWC).ValueRange(-1, 1);
178+ int64_t reduction = Reduction::None;
179+ bool logTarget = true;
180+ 
181+ auto outDesc = TensorDesc({3, 1, 2, 5, 4}, ACL_FLOAT, ACL_FORMAT_NDHWC).Precision(0.0001, 0.0001);
182+ 
183+ auto ut = OP_API_UT(
184+ aclnnKlDivBackward, INPUT(gradOutputDesc, selfDesc, targetDesc, reduction, logTarget), OUTPUT(outDesc));
185+ 
186+ uint64_t workspace_size = 0;
187+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
188+ EXPECT_EQ(aclRet, ACL_SUCCESS);
189+ 
190+ ut.TestPrecision();
191+}
192+ 
193+TEST_F(kl_div_loss_backward_test, aclnnKlDivBackward_09_empty_tensor)
194+{
195+ auto gradOutputDesc = TensorDesc({3, 1, 0, 5}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
196+ auto selfDesc = TensorDesc({3, 1, 0, 5}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
197+ auto targetDesc = TensorDesc({3, 1, 0, 5}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
198+ int64_t reduction = Reduction::None;
199+ bool logTarget = true;
200+ 
201+ auto outDesc = TensorDesc({3, 1, 0, 5}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.0001, 0.0001);
202+ 
203+ auto ut = OP_API_UT(
204+ aclnnKlDivBackward, INPUT(gradOutputDesc, selfDesc, targetDesc, reduction, logTarget), OUTPUT(outDesc));
205+ 
206+ uint64_t workspace_size = 0;
207+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
208+ EXPECT_EQ(aclRet, ACL_SUCCESS);
209+ 
210+ ut.TestPrecision();
211+}
212+ 
213+TEST_F(kl_div_loss_backward_test, aclnnKlDivBackward_10_onedim_tensor_logtarget_true)
214+{
215+ auto gradOutputDesc = TensorDesc({3}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
216+ auto selfDesc = TensorDesc({3}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
217+ auto targetDesc = TensorDesc({3}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
218+ int64_t reduction = Reduction::None;
219+ bool logTarget = true;
220+ 
221+ auto outDesc = TensorDesc({3}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.0001, 0.0001);
222+ 
223+ auto ut = OP_API_UT(
224+ aclnnKlDivBackward, INPUT(gradOutputDesc, selfDesc, targetDesc, reduction, logTarget), OUTPUT(outDesc));
225+ 
226+ uint64_t workspace_size = 0;
227+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
228+ EXPECT_EQ(aclRet, ACL_SUCCESS);
229+ 
230+ ut.TestPrecision();
231+}
232+ 
233+TEST_F(kl_div_loss_backward_test, aclnnKlDivBackward_11_threedim_tensor_broadcast)
234+{
235+ auto gradOutputDesc = TensorDesc({3, 4, 1}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
236+ auto selfDesc = TensorDesc({3, 4, 6}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
237+ auto targetDesc = TensorDesc({3, 4, 6}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
238+ int64_t reduction = Reduction::None;
239+ bool logTarget = true;
240+ 
241+ auto outDesc = TensorDesc({3, 4, 6}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.0001, 0.0001);
242+ 
243+ auto ut = OP_API_UT(
244+ aclnnKlDivBackward, INPUT(gradOutputDesc, selfDesc, targetDesc, reduction, logTarget), OUTPUT(outDesc));
245+ 
246+ uint64_t workspace_size = 0;
247+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
248+ EXPECT_EQ(aclRet, ACL_SUCCESS);
249+ 
250+ ut.TestPrecision();
251+}
252+ 
253+TEST_F(kl_div_loss_backward_test, aclnnKlDivBackward_12_fivedim_tensor_logtarget_true)
254+{
255+ auto gradOutputDesc = TensorDesc({3, 4, 6, 2, 3}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
256+ auto selfDesc = TensorDesc({3, 4, 6, 2, 3}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
257+ auto targetDesc = TensorDesc({3, 4, 6, 2, 3}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
258+ int64_t reduction = Reduction::None;
259+ bool logTarget = true;
260+ 
261+ auto outDesc = TensorDesc({3, 4, 6, 2, 3}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.0001, 0.0001);
262+ 
263+ auto ut = OP_API_UT(
264+ aclnnKlDivBackward, INPUT(gradOutputDesc, selfDesc, targetDesc, reduction, logTarget), OUTPUT(outDesc));
265+ 
266+ uint64_t workspace_size = 0;
267+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
268+ EXPECT_EQ(aclRet, ACL_SUCCESS);
269+ 
270+ ut.TestPrecision();
271+}
272+ 
273+TEST_F(kl_div_loss_backward_test, aclnnKlDivBackward_13_dtype_promote_float16_to_float)
274+{
275+ auto gradOutputDesc = TensorDesc({3, 4, 1, 2, 3}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-1, 1);
276+ auto selfDesc = TensorDesc({3, 4, 1, 2, 3}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
277+ auto targetDesc = TensorDesc({3, 4, 1, 2, 3}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-1, 1);
278+ int64_t reduction = Reduction::None;
279+ bool logTarget = true;
280+ 
281+ auto outDesc = TensorDesc({3, 4, 1, 2, 3}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.0001, 0.0001);
282+ 
283+ auto ut = OP_API_UT(
284+ aclnnKlDivBackward, INPUT(gradOutputDesc, selfDesc, targetDesc, reduction, logTarget), OUTPUT(outDesc));
285+ 
286+ uint64_t workspace_size = 0;
287+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
288+ EXPECT_EQ(aclRet, ACL_SUCCESS);
289+ 
290+ ut.TestPrecision();
291+}
292+ 
293+TEST_F(kl_div_loss_backward_test, aclnnKlDivBackward_14_reduction_error)
294+{
295+ auto gradOutputDesc = TensorDesc({3, 5, 4, 6}, ACL_FLOAT, ACL_FORMAT_ND);
296+ auto selfDesc = TensorDesc({3, 5, 4, 6}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
297+ auto targetDesc = TensorDesc({3, 5, 4, 6}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
298+ int64_t reduction = Reduction::END;
299+ bool logTarget = false;
300+ 
301+ auto outDesc = TensorDesc({3, 5, 4, 6}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.0001, 0.0001);
302+ 
303+ auto ut = OP_API_UT(
304+ aclnnKlDivBackward, INPUT(gradOutputDesc, selfDesc, targetDesc, reduction, logTarget), OUTPUT(outDesc));
305+ 
306+ uint64_t workspace_size = 0;
307+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
308+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
309+}
310+ 
311+TEST_F(kl_div_loss_backward_test, aclnnKlDivBackward_15_input_out_nullptr)
312+{
313+ auto tensorDesc = TensorDesc({10, 3, 5, 24}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
314+ int64_t reduction = Reduction::None;
315+ bool logTarget = true;
316+ 
317+ auto ut_grad = OP_API_UT(
318+ aclnnKlDivBackward, INPUT((aclTensor*)nullptr, tensorDesc, tensorDesc, reduction, logTarget),
319+ OUTPUT(tensorDesc));
320+ uint64_t workspace_size = 0;
321+ aclnnStatus aclRet = ut_grad.TestGetWorkspaceSize(&workspace_size);
322+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_NULLPTR);
323+ 
324+ auto ut_self = OP_API_UT(
325+ aclnnKlDivBackward, INPUT(tensorDesc, (aclTensor*)nullptr, tensorDesc, reduction, logTarget),
326+ OUTPUT(tensorDesc));
327+ aclRet = ut_self.TestGetWorkspaceSize(&workspace_size);
328+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_NULLPTR);
329+ 
330+ auto ut_tar = OP_API_UT(
331+ aclnnKlDivBackward, INPUT(tensorDesc, tensorDesc, (aclTensor*)nullptr, reduction, logTarget),
332+ OUTPUT(tensorDesc));
333+ aclRet = ut_tar.TestGetWorkspaceSize(&workspace_size);
334+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_NULLPTR);
335+ 
336+ auto ut_o = OP_API_UT(
337+ aclnnKlDivBackward, INPUT(tensorDesc, tensorDesc, tensorDesc, reduction, logTarget),
338+ OUTPUT((aclTensor*)nullptr));
339+ aclRet = ut_o.TestGetWorkspaceSize(&workspace_size);
340+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_NULLPTR);
341+}
342+ 
343+TEST_F(kl_div_loss_backward_test, aclnnKlDivBackward_16_input_error_shape_len)
344+{
345+ auto tensorDesc9 = TensorDesc({10, 24, 3, 5, 10, 22, 42, 30, 24}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
346+ auto tensorDesc4 = TensorDesc({3, 1, 0, 5}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
347+ int64_t reduction = Reduction::None;
348+ bool logTarget = true;
349+ 
350+ auto ut_grad = OP_API_UT(
351+ aclnnKlDivBackward, INPUT(tensorDesc9, tensorDesc4, tensorDesc4, reduction, logTarget), OUTPUT(tensorDesc4));
352+ uint64_t workspace_size = 0;
353+ aclnnStatus aclRet = ut_grad.TestGetWorkspaceSize(&workspace_size);
354+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
355+ 
356+ auto ut_self = OP_API_UT(
357+ aclnnKlDivBackward, INPUT(tensorDesc4, tensorDesc9, tensorDesc4, reduction, logTarget), OUTPUT(tensorDesc4));
358+ aclRet = ut_self.TestGetWorkspaceSize(&workspace_size);
359+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
360+ 
361+ auto ut_tar = OP_API_UT(
362+ aclnnKlDivBackward, INPUT(tensorDesc4, tensorDesc4, tensorDesc9, reduction, logTarget), OUTPUT(tensorDesc4));
363+ aclRet = ut_tar.TestGetWorkspaceSize(&workspace_size);
364+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
365+}
366+ 
367+TEST_F(kl_div_loss_backward_test, aclnnKlDivBackward_17_error_input_dtype)
368+{
369+ auto gradOutputDesc = TensorDesc({3, 4, 1, 1}, ACL_INT32, ACL_FORMAT_ND).ValueRange(-1, 1);
370+ auto selfDesc = TensorDesc({3, 4, 6, 3}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
371+ auto targetDesc = TensorDesc({3, 4, 6, 1}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-1, 1);
372+ int64_t reduction = Reduction::None;
373+ bool logTarget = true;
374+ 
375+ auto outDesc = TensorDesc({3, 4, 6, 3}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.0001, 0.0001);
376+ 
377+ auto ut = OP_API_UT(
378+ aclnnKlDivBackward, INPUT(gradOutputDesc, selfDesc, targetDesc, reduction, logTarget), OUTPUT(outDesc));
379+ uint64_t workspace_size = 0;
380+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
381+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
382+}
383+ 
384+TEST_F(kl_div_loss_backward_test, aclnnKlDivBackward_18_target_not_broadcast)
385+{
386+ auto gradOutputDesc = TensorDesc({3, 4, 1, 1}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-1, 1);
387+ auto selfDesc = TensorDesc({3, 4, 6, 3}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
388+ auto targetDesc = TensorDesc({3, 4, 6, 4}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-1, 1);
389+ int64_t reduction = Reduction::None;
390+ bool logTarget = true;
391+ 
392+ auto outDesc = TensorDesc({3, 4, 6, 3}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.0001, 0.0001);
393+ 
394+ auto ut = OP_API_UT(
395+ aclnnKlDivBackward, INPUT(gradOutputDesc, selfDesc, targetDesc, reduction, logTarget), OUTPUT(outDesc));
396+ uint64_t workspace_size = 0;
397+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
398+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
399+}
400+ 
401+TEST_F(kl_div_loss_backward_test, aclnnKlDivBackward_19_out_shape_not_equal_self)
402+{
403+ auto gradOutputDesc = TensorDesc({3, 4, 1, 1}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-1, 1);
404+ auto selfDesc = TensorDesc({3, 4, 6, 3}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
405+ auto targetDesc = TensorDesc({3, 4, 6, 1}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-1, 1);
406+ int64_t reduction = Reduction::None;
407+ bool logTarget = true;
408+ 
409+ auto outDesc = TensorDesc({3, 4, 6, 5}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.0001, 0.0001);
410+ 
411+ auto ut = OP_API_UT(
412+ aclnnKlDivBackward, INPUT(gradOutputDesc, selfDesc, targetDesc, reduction, logTarget), OUTPUT(outDesc));
413+ uint64_t workspace_size = 0;
414+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
415+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
416+}
417+ 
418+TEST_F(kl_div_loss_backward_test, aclnnKlDivBackward_20_self_not_equal_broadcast_shape)
419+{
420+ auto gradOutputDesc = TensorDesc({3, 4, 1, 3}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-1, 1);
421+ auto selfDesc = TensorDesc({3, 4, 6, 1}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
422+ auto targetDesc = TensorDesc({3, 4, 6, 1}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-1, 1);
423+ int64_t reduction = Reduction::None;
424+ bool logTarget = true;
425+ 
426+ auto outDesc = TensorDesc({3, 4, 6, 1}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.0001, 0.0001);
427+ 
428+ auto ut = OP_API_UT(
429+ aclnnKlDivBackward, INPUT(gradOutputDesc, selfDesc, targetDesc, reduction, logTarget), OUTPUT(outDesc));
430+ uint64_t workspace_size = 0;
431+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
432+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
433+}
434+ 
435+TEST_F(kl_div_loss_backward_test, aclnnKlDivBackward_21_broadcast_grad_1d)
436+{
437+ auto gradOutputDesc = TensorDesc({3, 1}, ACL_FLOAT, ACL_FORMAT_ND);
438+ auto selfDesc = TensorDesc({3, 1}, ACL_FLOAT, ACL_FORMAT_ND);
439+ auto targetDesc = TensorDesc({3, 5}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(0, 4);
440+ int64_t reduction = Reduction::None;
441+ bool logTarget = false;
442+ 
443+ auto outDesc = TensorDesc({3, 1}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.0001, 0.0001);
444+ 
445+ auto ut = OP_API_UT(
446+ aclnnKlDivBackward, INPUT(gradOutputDesc, selfDesc, targetDesc, reduction, logTarget), OUTPUT(outDesc));
447+ 
448+ uint64_t workspace_size = 0;
449+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
450+ EXPECT_EQ(aclRet, ACL_SUCCESS);
451+ 
452+ ut.TestPrecision();
453+}
454+ 
455+TEST_F(kl_div_loss_backward_test, aclnnKlDivBackward_22_broadcast_target_1d)
456+{
457+ auto gradOutputDesc = TensorDesc({3, 1}, ACL_FLOAT, ACL_FORMAT_ND);
458+ auto selfDesc = TensorDesc({3, 5}, ACL_FLOAT, ACL_FORMAT_ND);
459+ auto targetDesc = TensorDesc({3, 1}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(0, 4);
460+ int64_t reduction = Reduction::None;
461+ bool logTarget = false;
462+ 
463+ auto outDesc = TensorDesc({3, 5}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.0001, 0.0001);
464+ 
465+ auto ut = OP_API_UT(
466+ aclnnKlDivBackward, INPUT(gradOutputDesc, selfDesc, targetDesc, reduction, logTarget), OUTPUT(outDesc));
467+ 
468+ uint64_t workspace_size = 0;
469+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
470+ EXPECT_EQ(aclRet, ACL_SUCCESS);
471+ 
472+ ut.TestPrecision();
473+}
474+ 
475+TEST_F(kl_div_loss_backward_test, aclnnKlDivBackward_23_broadcast_mixed_rank)
476+{
477+ auto gradOutputDesc = TensorDesc({3, 1, 1}, ACL_FLOAT, ACL_FORMAT_ND);
478+ auto selfDesc = TensorDesc({3, 5, 1}, ACL_FLOAT, ACL_FORMAT_ND);
479+ auto targetDesc = TensorDesc({6, 3, 1, 7}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(0, 4);
480+ int64_t reduction = Reduction::None;
481+ bool logTarget = false;
482+ 
483+ auto outDesc = TensorDesc({3, 5, 1}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.0001, 0.0001);
484+ 
485+ auto ut = OP_API_UT(
486+ aclnnKlDivBackward, INPUT(gradOutputDesc, selfDesc, targetDesc, reduction, logTarget), OUTPUT(outDesc));
487+ 
488+ uint64_t workspace_size = 0;
489+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
490+ EXPECT_EQ(aclRet, ACL_SUCCESS);
491+ 
492+ ut.TestPrecision();
493+}
Aexperimental/loss/kl_div_loss_grad/tests/ut/op_host/CMakeLists.txt+15-0
@@ -0,0 +1,15 @@
1+# ----------------------------------------------------------------------------
2+ # Copyright (c) 2026 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(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()
Aexperimental/loss/kl_div_loss_grad/tests/ut/op_host/test_kl_div_loss_grad_tiling.cpp+222-0
@@ -0,0 +1,222 @@
1+/**
2+ * Copyright (c) 2026 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 <gtest/gtest.h>
14+#include "log/log.h"
15+#include "kernel_run_context_facker.h"
16+#include "exe_graph/runtime/storage_format.h"
17+#include "exe_graph/runtime/storage_shape.h"
18+#include "test_cube_util.h"
19+#include "register/op_impl_registry.h"
20+#include "ut_op_util.h"
21+#include "ut_op_common.h"
22+#include "platform/platform_infos_def.h"
23+ 
24+using namespace ut_util;
25+using namespace std;
26+using namespace ge;
27+ 
28+class KlDivLossGradTiling : public testing::Test {
29+protected:
30+ static void SetUpTestCase() { cout << "KlDivLossGradTiling SetUp" << endl; }
31+ 
32+ static void TearDownTestCase() { cout << "KlDivLossGradTiling TearDown " << endl; }
33+};
34+ 
35+TEST_F(KlDivLossGradTiling, kl_div_loss_grad_float32_success)
36+{
37+ gert::StorageShape x1Shape = {{1, 2, 8, 16}, {1, 2, 8, 16}};
38+ gert::StorageShape x2Shape = {{1, 2, 8, 16}, {1, 2, 8, 16}};
39+ gert::StorageShape x3Shape = {{1, 2, 8, 16}, {1, 2, 8, 16}};
40+ gert::StorageShape yShape = {{1, 2, 8, 16}, {1, 2, 8, 16}};
41+ string compileInfoString = R"({
42+ "hardware_info": {"BT_SIZE": 0, "load3d_constraints": "1",
43+ "Intrinsic_fix_pipe_l0c2out": false, "Intrinsic_data_move_l12ub": true, "Intrinsic_data_move_l0c2ub": true, "Intrinsic_data_move_out2l1_nd2nz": false,
44+ "UB_SIZE": 196608, "L2_SIZE": 33554432, "L1_SIZE": 524288,
45+ "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072,
46+ "CORE_NUM": 48}
47+ })";
48+ map<string, string> socInfos;
49+ map<string, string> aicoreSpec;
50+ map<string, string> intrinsics;
51+ GetPlatFormInfos(compileInfoString.c_str(), socInfos, aicoreSpec, intrinsics);
52+ 
53+ fe::PlatFormInfos platformInfo;
54+ platformInfo.Init();
55+ 
56+ struct KlDivLossGradTilingCompileInfo {};
57+ KlDivLossGradTilingCompileInfo compileInfo;
58+ 
59+ std::string opType("KlDivLossGrad");
60+ auto tilingFunc = gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str())->tiling;
61+ 
62+ auto kernelHolder =
63+ gert::KernelRunContextFaker()
64+ .KernelIONum(3, 1)
65+ .Inputs({const_cast<char*>(compileInfoString.c_str()), reinterpret_cast<void*>(&platformInfo)})
66+ .Outputs({&compileInfo})
67+ .Build();
68+ kernelHolder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", socInfos);
69+ kernelHolder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicoreSpec);
70+ kernelHolder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
71+ kernelHolder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes(
72+ "AICoreintrinsicDtypeMap", intrinsics);
73+ 
74+ auto param = gert::TilingData::CreateCap(4096);
75+ auto workspaceSizeHoler = gert::ContinuousVector::Create<size_t>(4096);
76+ auto wsSize = reinterpret_cast<gert::ContinuousVector*>(workspaceSizeHoler.get());
77+ ASSERT_NE(param, nullptr);
78+ auto holder = gert::TilingContextFaker()
79+ .SetOpType("KlDivLossGrad")
80+ .NodeIoNum(3, 1)
81+ .IrInstanceNum({1, 1})
82+ .InputShapes({&x1Shape, &x2Shape, &x3Shape})
83+ .OutputShapes({&yShape})
84+ .CompileInfo(&compileInfo)
85+ .PlatformInfo(reinterpret_cast<char*>(&platformInfo))
86+ .NodeInputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
87+ .NodeInputTd(1, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
88+ .NodeInputTd(2, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
89+ .NodeOutputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
90+ .TilingData(param.get())
91+ .Workspace(wsSize)
92+ .Build();
93+ gert::TilingContext* tilingContext = holder.GetContext<gert::TilingContext>();
94+ ASSERT_NE(tilingContext, nullptr);
95+ EXPECT_EQ(tilingFunc(tilingContext), ge::GRAPH_SUCCESS);
96+}
97+ 
98+TEST_F(KlDivLossGradTiling, kl_div_loss_grad_float16_success)
99+{
100+ gert::StorageShape x1Shape = {{1, 2, 8, 16}, {1, 2, 8, 16}};
101+ gert::StorageShape x2Shape = {{1, 2, 8, 16}, {1, 2, 8, 16}};
102+ gert::StorageShape x3Shape = {{1, 2, 8, 16}, {1, 2, 8, 16}};
103+ gert::StorageShape yShape = {{1, 2, 8, 16}, {1, 2, 8, 16}};
104+ string compileInfoString = R"({
105+ "hardware_info": {"BT_SIZE": 0, "load3d_constraints": "1",
106+ "Intrinsic_fix_pipe_l0c2out": false, "Intrinsic_data_move_l12ub": true, "Intrinsic_data_move_l0c2ub": true, "Intrinsic_data_move_out2l1_nd2nz": false,
107+ "UB_SIZE": 196608, "L2_SIZE": 33554432, "L1_SIZE": 524288,
108+ "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072,
109+ "CORE_NUM": 48}
110+ })";
111+ map<string, string> socInfos;
112+ map<string, string> aicoreSpec;
113+ map<string, string> intrinsics;
114+ GetPlatFormInfos(compileInfoString.c_str(), socInfos, aicoreSpec, intrinsics);
115+ 
116+ fe::PlatFormInfos platformInfo;
117+ platformInfo.Init();
118+ 
119+ struct KlDivLossGradTilingCompileInfo {};
120+ KlDivLossGradTilingCompileInfo compileInfo;
121+ 
122+ std::string opType("KlDivLossGrad");
123+ auto tilingFunc = gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str())->tiling;
124+ 
125+ auto kernelHolder =
126+ gert::KernelRunContextFaker()
127+ .KernelIONum(3, 1)
128+ .Inputs({const_cast<char*>(compileInfoString.c_str()), reinterpret_cast<void*>(&platformInfo)})
129+ .Outputs({&compileInfo})
130+ .Build();
131+ kernelHolder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", socInfos);
132+ kernelHolder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicoreSpec);
133+ kernelHolder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
134+ kernelHolder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes(
135+ "AICoreintrinsicDtypeMap", intrinsics);
136+ 
137+ auto param = gert::TilingData::CreateCap(4096);
138+ auto workspaceSizeHoler = gert::ContinuousVector::Create<size_t>(4096);
139+ auto wsSize = reinterpret_cast<gert::ContinuousVector*>(workspaceSizeHoler.get());
140+ ASSERT_NE(param, nullptr);
141+ auto holder = gert::TilingContextFaker()
142+ .SetOpType("KlDivLossGrad")
143+ .NodeIoNum(3, 1)
144+ .IrInstanceNum({1, 1})
145+ .InputShapes({&x1Shape, &x2Shape, &x3Shape})
146+ .OutputShapes({&yShape})
147+ .CompileInfo(&compileInfo)
148+ .PlatformInfo(reinterpret_cast<char*>(&platformInfo))
149+ .NodeInputTd(0, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
150+ .NodeInputTd(1, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
151+ .NodeInputTd(2, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
152+ .NodeOutputTd(0, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
153+ .TilingData(param.get())
154+ .Workspace(wsSize)
155+ .Build();
156+ gert::TilingContext* tilingContext = holder.GetContext<gert::TilingContext>();
157+ ASSERT_NE(tilingContext, nullptr);
158+ EXPECT_EQ(tilingFunc(tilingContext), ge::GRAPH_SUCCESS);
159+}
160+ 
161+TEST_F(KlDivLossGradTiling, kl_div_loss_grad_bf16_success)
162+{
163+ gert::StorageShape x1Shape = {{1, 2, 8, 16}, {1, 2, 8, 16}};
164+ gert::StorageShape x2Shape = {{1, 2, 8, 16}, {1, 2, 8, 16}};
165+ gert::StorageShape x3Shape = {{1, 2, 8, 16}, {1, 2, 8, 16}};
166+ gert::StorageShape yShape = {{1, 2, 8, 16}, {1, 2, 8, 16}};
167+ string compileInfoString = R"({
168+ "hardware_info": {"BT_SIZE": 0, "load3d_constraints": "1",
169+ "Intrinsic_fix_pipe_l0c2out": false, "Intrinsic_data_move_l12ub": true, "Intrinsic_data_move_l0c2ub": true, "Intrinsic_data_move_out2l1_nd2nz": false,
170+ "UB_SIZE": 196608, "L2_SIZE": 33554432, "L1_SIZE": 524288,
171+ "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072,
172+ "CORE_NUM": 48}
173+ })";
174+ map<string, string> socInfos;
175+ map<string, string> aicoreSpec;
176+ map<string, string> intrinsics;
177+ GetPlatFormInfos(compileInfoString.c_str(), socInfos, aicoreSpec, intrinsics);
178+ 
179+ fe::PlatFormInfos platformInfo;
180+ platformInfo.Init();
181+ 
182+ struct KlDivLossGradTilingCompileInfo {};
183+ KlDivLossGradTilingCompileInfo compileInfo;
184+ 
185+ std::string opType("KlDivLossGrad");
186+ auto tilingFunc = gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str())->tiling;
187+ 
188+ auto kernelHolder =
189+ gert::KernelRunContextFaker()
190+ .KernelIONum(3, 1)
191+ .Inputs({const_cast<char*>(compileInfoString.c_str()), reinterpret_cast<void*>(&platformInfo)})
192+ .Outputs({&compileInfo})
193+ .Build();
194+ kernelHolder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", socInfos);
195+ kernelHolder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicoreSpec);
196+ kernelHolder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
197+ kernelHolder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes(
198+ "AICoreintrinsicDtypeMap", intrinsics);
199+ 
200+ auto param = gert::TilingData::CreateCap(4096);
201+ auto workspaceSizeHoler = gert::ContinuousVector::Create<size_t>(4096);
202+ auto wsSize = reinterpret_cast<gert::ContinuousVector*>(workspaceSizeHoler.get());
203+ ASSERT_NE(param, nullptr);
204+ auto holder = gert::TilingContextFaker()
205+ .SetOpType("KlDivLossGrad")
206+ .NodeIoNum(3, 1)
207+ .IrInstanceNum({1, 1})
208+ .InputShapes({&x1Shape, &x2Shape, &x3Shape})
209+ .OutputShapes({&yShape})
210+ .CompileInfo(&compileInfo)
211+ .PlatformInfo(reinterpret_cast<char*>(&platformInfo))
212+ .NodeInputTd(0, ge::DT_BF16, ge::FORMAT_ND, ge::FORMAT_ND)
213+ .NodeInputTd(1, ge::DT_BF16, ge::FORMAT_ND, ge::FORMAT_ND)
214+ .NodeInputTd(2, ge::DT_BF16, ge::FORMAT_ND, ge::FORMAT_ND)
215+ .NodeOutputTd(0, ge::DT_BF16, ge::FORMAT_ND, ge::FORMAT_ND)
216+ .TilingData(param.get())
217+ .Workspace(wsSize)
218+ .Build();
219+ gert::TilingContext* tilingContext = holder.GetContext<gert::TilingContext>();
220+ ASSERT_NE(tilingContext, nullptr);
221+ EXPECT_EQ(tilingFunc(tilingContext), ge::GRAPH_SUCCESS);
222+}
Aexperimental/loss/kl_div_loss_grad/tests/ut/op_kernel/CMakeLists.txt+30-0
@@ -0,0 +1,30 @@
1+# ----------------------------------------------------------------------------
2+ # Copyright (c) 2026 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+if ((UT_TEST_ALL OR OP_KERNEL_UT) AND NOT UT_DONE)
13+ # 需要将Tiling依赖的文件添加到CMakeLists.txt中
14+ # set(elewise_common_tiling_files
15+ # ${CANN_ROOT}/ops/built-in/op_tiling/runtime/elewise_tiling.cc
16+ # )
17+ # 算子自己的tiling文件路径
18+ set(kl_div_loss_grad_tiling_files
19+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/kl_div_loss_grad_tiling.cpp
20+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/kl_div_loss_grad_infershape.cpp
21+ # ${elewise_common_tiling_files}
22+ )
23+ # 使用AddOpTestCase
24+ # param1：算子名称，以kernel方式命名
25+ # param2：soc版本，多个以分号分隔，例如："ascend950pr_9599;AscendB1"
26+ # param3：自定义编译选项，一般填写测试的一种典型数据类型组合，不需要则传入空字符串，例如："-DDTYPE_X=float"，多个使用空格分隔，例如："-DDTYPE_X=float -DDTYPE_Y=float"
27+ # param4：该算子依赖的所有tiling源码文件
28+ # AddOpTestCase(kl_div_loss_grad "ascend910b" "${kl_div_loss_grad_tiling_files}")
29+ AddOpTestCase(kl_div_loss_grad "ascend910B1" "-DDTYPE_GRAD=float" "${kl_div_loss_grad_tiling_files}")
30+endif()
Aexperimental/loss/kl_div_loss_grad/tests/ut/op_kernel/test_kl_div_loss_grad.cpp+69-0
@@ -0,0 +1,69 @@
1+/**
2+ * Copyright (c) 2026 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 <array>
12+#include <vector>
13+#include "gtest/gtest.h"
14+ 
15+#ifdef __CCE_KT_TEST__
16+#include "tikicpulib.h"
17+#include "data_utils.h"
18+#include "string.h"
19+#include <iostream>
20+#include <string>
21+#endif
22+#include "../../../op_kernel/kl_div_loss_grad.h"
23+#include <cstdint>
24+ 
25+using namespace std;
26+ 
27+class kl_div_loss_grad_test : public testing::Test {
28+protected:
29+ static void SetUpTestCase() { cout << "kl_div_loss_grad_test SetUp\n" << endl; }
30+ static void TearDownTestCase() { cout << "kl_div_loss_grad_test TearDown\n" << endl; }
31+};
32+ 
33+TEST_F(kl_div_loss_grad_test, test_case_0)
34+{
35+ size_t gradByteSize = 32 * 4 * 4 * 4 * sizeof(float);
36+ size_t inputByteSize = 32 * 4 * 4 * 4 * sizeof(float);
37+ size_t targetByteSize = 32 * 4 * 4 * 4 * sizeof(float);
38+ size_t yByteSize = 32 * 4 * 4 * 4 * sizeof(float);
39+ size_t tilingDataSize = sizeof(KlDivLossGradTilingData);
40+ uint32_t blockDim = 1;
41+ 
42+ uint8_t* grad = (uint8_t*)AscendC::GmAlloc(gradByteSize);
43+ uint8_t* input = (uint8_t*)AscendC::GmAlloc(inputByteSize);
44+ uint8_t* target = (uint8_t*)AscendC::GmAlloc(targetByteSize);
45+ uint8_t* y = (uint8_t*)AscendC::GmAlloc(yByteSize);
46+ 
47+ uint8_t* workspace = (uint8_t*)AscendC::GmAlloc(1024 * 1024 * 16);
48+ uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(tilingDataSize);
49+ 
50+ auto* tilingDataFromBin = reinterpret_cast<KlDivLossGradTilingData*>(tiling);
51+ tilingDataFromBin->bigCoreDataNum = 2048;
52+ tilingDataFromBin->smallCoreDataNum = 2048;
53+ tilingDataFromBin->tileDataNum = 2048;
54+ tilingDataFromBin->bigCoreNum = 1;
55+ tilingDataFromBin->coff = 1.0f;
56+ 
57+ auto klDivLossGradKernel = kl_div_loss_grad<false, false>;
58+ 
59+ ICPU_SET_TILING_KEY(0);
60+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
61+ ICPU_RUN_KF(klDivLossGradKernel, blockDim, grad, input, target, y, workspace, (uint8_t*)(tilingDataFromBin));
62+ 
63+ AscendC::GmFree(grad);
64+ AscendC::GmFree(input);
65+ AscendC::GmFree(target);
66+ AscendC::GmFree(y);
67+ AscendC::GmFree(workspace);
68+ AscendC::GmFree(tiling);
69+}