Star889
1799
代码
代码
Issues259
Pull Requests295
Actions
项目讨论
Wiki
分析
Star889
1799
已合并
ThresholdGradV2D 算子贡献 #3222
松柏创建于 3月26日
ThresholdGradV2D 算子贡献 #3222
已合并
松柏创建于 3月26日
从ilovescrapy/thresholdgradv2dmaster合入到cann/ops-nnmaster
共 24 个文件变更+2070-0
Aexperimental/activation/threshold_grad_v2_d/CMakeLists.txt+20-0
@@ -0,0 +1,20 @@
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(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
13+if(NOT ENABLE_TEST)
14+ list(REMOVE_ITEM CURRENT_DIRS tests)
15+endif()
16+foreach(SUB_DIR ${CURRENT_DIRS})
17+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
18+ add_subdirectory(${SUB_DIR})
19+ endif()
20+endforeach()
Aexperimental/activation/threshold_grad_v2_d/README.md+89-0
@@ -0,0 +1,89 @@
1+# ThresholdGradV2D
2+ 
3+## 产品支持情况
4+ 
5+ 
6+| 产品 | 是否支持 |
7+| ------------------------------------------------------------------ | :------: |
8+| Atlas A2 训练系列产品/Atlas 800I A2 推理产品/A200I A2 Box 异构组件 | √ |
9+ 
10+## 功能说明
11+ 
12+- 算子功能：求thresholdv2d函数梯度。
13+- 计算公式：
14+ 
15+$$
16+ y =
17+ \begin{cases}
18+ input\_ gradient, & input\_ feature > threshold \\
19+ 0, & input\_ feature <= threshold
20+ \end{cases}
21+$$
22+ 
23+## 参数说明
24+ 
25+<table style="undefined;table-layout: fixed; width: 820px"><colgroup>
26+ <col style="width: 100px">
27+ <col style="width: 150px">
28+ <col style="width: 190px">
29+ <col style="width: 260px">
30+ <col style="width: 120px">
31+ </colgroup>
32+ <thead>
33+ <tr>
34+ <th>参数名</th>
35+ <th>输入/输出/属性</th>
36+ <th>描述</th>
37+ <th>数据类型</th>
38+ <th>数据格式</th>
39+ </tr></thead>
40+ <tbody>
41+ <tr>
42+ <td>input_gradient</td>
43+ <td>输入</td>
44+ <td>反向传播梯度</td>
45+ <td>FLOAT、FLOAT16、BFLOAT16、INT32、INT8、UINT8</td>
46+ <td>ND</td>
47+ </tr>
48+ <tr>
49+ <td>input_feature</td>
50+ <td>输入</td>
51+ <td>与"input_gradient"具有相同的类型、格式和形状。</td>
52+ <td>FLOAT、FLOAT16、BFLOAT16、INT32、INT8、UINT8</td>
53+ <td>ND</td>
54+ </tr>
55+ <tr>
56+ <td>threshold</td>
57+ <td>属性</td>
58+ <td></td>
59+ <td>FLOAT</td>
60+ <td></td>
61+ </tr>
62+ <tr>
63+ <td>output_backprops</td>
64+ <td>输出</td>
65+ <td>公式中的输出张量</td>
66+ <td>FLOAT、FLOAT16、BFLOAT16、INT32、INT8、UINT8</td>
67+ <td>ND</td>
68+ </tr>
69+ </tbody></table>
70+ 
71+ 
72+ 
73+## 调用说明
74+ 
75+| 调用方式 | 样例代码 | 说明 |
76+| ---------------- | --------------------------- | --------------------------------------------------- |
77+| Aclnn模式 | [test_aclnn_threshold_backward.cpp](./examples/test_aclnn_threshold_backward.cpp) | 通过Aclnn接口调用ThresholdGradV2D算子。 |
78+ 
79+ 
80+## 约束说明
81+ 
82+无
83+ 
84+## 贡献说明
85+ 
86+ 
87+| 贡献者 | 贡献方 | 贡献算子 | 贡献时间 | 贡献内容 |
88+| ----------- | ---------- | --------- | ---------- | ----------------------- |
89+| ilovescrapy | 个人开发者 | ThresholdGradV2D | 2026/3/23 | ThresholdGradV2D算子适配开源仓 |
Aexperimental/activation/threshold_grad_v2_d/docs/aclnnThresholdBackward.md+366-0
@@ -0,0 +1,366 @@
1+# aclnnThresholdBackward
2+ 
3+## 产品支持情况
4+ 
5+|产品 | 是否支持 |
6+|:-------------------------|:----------:|
7+| <term>Atlas A2 训练系列产品/Atlas 800I A2 推理产品/A200I A2 Box 异构组件</term> | √ |
8+ 
9+## 功能说明
10+ 
11+- 接口功能：完成[aclnnThreshold](../../../../activation/threshold/docs/aclnnThreshold&aclnnInplaceThreshold.md)的反向。
12+- 计算公式：
13+ 
14+ $$
15+ output =
16+ \begin{cases}
17+ gradOutput(i) & \text{if } self(i) > threshold \\
18+ 0 & \text{otherwise}
19+ \end{cases}
20+ $$
21+ 
22+## 函数原型
23+ 
24+每个算子分为[两段式接口](../../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnThresholdBackwardGetWorkspaceSize”接口获取入参并根据计算流程计算所需workspace大小，再调用“aclnnThresholdBackward”接口执行计算。
25+ 
26+```Cpp
27+aclnnStatus aclnnThresholdBackwardGetWorkspaceSize(
28+ const aclTensor *gradOutput,
29+ const aclTensor *self,
30+ const aclScalar *threshold,
31+ aclTensor *out,
32+ uint64_t *workspaceSize,
33+ aclOpExecutor **executor)
34+```
35+ 
36+```Cpp
37+aclnnStatus aclnnThresholdBackward(
38+ void *workspace,
39+ uint64_t workspaceSize,
40+ aclOpExecutor *executor,
41+ const aclrtStream stream)
42+```
43+ 
44+## aclnnThresholdBackwardGetWorkspaceSize
45+ 
46+- **参数说明：**
47+ 
48+ <table style="undefined;table-layout: fixed; width: 1330px"><colgroup>
49+ <col style="width: 171px">
50+ <col style="width: 115px">
51+ <col style="width: 150px">
52+ <col style="width: 230px">
53+ <col style="width: 177px">
54+ <col style="width: 104px">
55+ <col style="width: 238px">
56+ <col style="width: 145px">
57+ </colgroup>
58+ <thead>
59+ <tr>
60+ <th>参数名</th>
61+ <th>输入/输出</th>
62+ <th>描述</th>
63+ <th>使用说明</th>
64+ <th>数据类型</th>
65+ <th>数据格式</th>
66+ <th>维度(shape)</th>
67+ <th>非连续Tensor</th>
68+ </tr></thead>
69+ <tbody>
70+ <tr>
71+ <td>gradOutput</td>
72+ <td>输入</td>
73+ <td>公式中的gradOutput。</td>
74+ <td><ul><li>支持空Tensor。</li><li>dtype需要与self保持一致。</li><li>shape需要与self满足<a href="../../../../docs/zh/context/broadcast关系.md" target="_blank">broadcast关系</a>。</li></ul></td>
75+ <td>FLOAT、BFLOAT16、FLOAT16、INT32、INT8、UINT8、INT64</td>
76+ <td>ND</td>
77+ <td>0-8</td>
78+ <td>√</td>
79+ </tr>
80+ <tr>
81+ <td>self</td>
82+ <td>输入</td>
83+ <td>公式中的threshold。</td>
84+ <td>数据类型与gradOutput的数据类型满足数据类型推导规则（参见<a href="../../../../docs/zh/context/互推导关系.md" target="_blank">互推导关系</a>）。</td>
85+ <td>FLOAT、BFLOAT16、FLOAT16、INT32、INT8、UINT8、INT64</td>
86+ <td>ND</td>
87+ <td>0-8</td>
88+ <td>√</td>
89+ </tr>
90+ <tr>
91+ <td>threshold</td>
92+ <td>输入</td>
93+ <td>公式中的self。</td>
94+ <td><ul><li>支持空Tensor。</li><li>dtype需要与gradOutput保持一致。</li><li>shape需要与gradOutput满足<a href="../../../../docs/zh/context/broadcast关系.md" target="_blank">broadcast关系。</a></li></ul></td>
95+ <td>FLOAT、BFLOAT16、FLOAT16、INT32、INT8、UINT8、INT64</td>
96+ <td>ND</td>
97+ <td>0-8</td>
98+ <td>√</td>
99+ </tr>
100+ <tr>
101+ <tr>
102+ <td>out</td>
103+ <td>输出</td>
104+ <td>公式中的out。</td>
105+ <td><ul><li>dtype需要与self相同。</li><li>shape需要与self和gradOutput进行broadcast之后的shape相等。</li></ul></td>
106+ <td>FLOAT、BFLOAT16、FLOAT16、INT32、INT8、UINT8、INT64</td>
107+ <td>ND</td>
108+ <td>0-8</td>
109+ <td>√</td>
110+ </tr>
111+ <tr>
112+ <td>workspaceSize</td>
113+ <td>输出</td>
114+ <td>返回需要在Device侧申请的workspace大小。</td>
115+ <td>-</td>
116+ <td>-</td>
117+ <td>-</td>
118+ <td>-</td>
119+ <td>-</td>
120+ </tr>
121+ <tr>
122+ <td>executor</td>
123+ <td>输出</td>
124+ <td>返回op执行器，包含了算子计算流程。</td>
125+ <td>-</td>
126+ <td>-</td>
127+ <td>-</td>
128+ <td>-</td>
129+ <td>-</td>
130+ </tr>
131+ </tbody>
132+ </table>
133+
134+ - <term>Atlas A2 训练系列产品/Atlas 800I A2 推理产品/A200I A2 Box 异构组件</term>、<term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：数据类型支持FLOAT、BFLOAT16、FLOAT16、INT32、INT8、UINT8。
135+- **返回值：**
136+ 
137+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../../docs/zh/context/aclnn返回码.md)。
138+ 第一段接口会完成入参校验，出现以下场景时报错：
139+ <table style="undefined;table-layout: fixed;width: 979px"><colgroup>
140+ <col style="width: 272px">
141+ <col style="width: 103px">
142+ <col style="width: 604px">
143+ </colgroup>
144+ <thead>
145+ <tr>
146+ <th>返回码</th>
147+ <th>错误码</th>
148+ <th>描述</th>
149+ </tr>
150+ </thead>
151+ <tbody>
152+ <tr>
153+ <td>ACLNN_ERR_PARAM_NULLPTR</td>
154+ <td>161001</td>
155+ <td>传入的gradOutput或self是空指针。</td>
156+ </tr>
157+ <tr>
158+ <td rowspan="8">ACLNN_ERR_PARAM_INVALID</td>
159+ <td rowspan="8">161002</td>
160+ <td>gradOutput或self的数据类型不在支持的范围之内。</td>
161+ </tr>
162+ <tr>
163+ <td>gradOutput或self的shape超过8维。</td>
164+ </tr>
165+ <tr>
166+ <td>gradOutput、out与self数据类型不一致。</td>
167+ </tr>
168+ </tbody></table>
169+ 
170+## aclnnThresholdBackward
171+ 
172+- **参数说明：**
173+ 
174+ <table style="undefined;table-layout: fixed; width: 953px"><colgroup>
175+ <col style="width: 173px">
176+ <col style="width: 112px">
177+ <col style="width: 668px">
178+ </colgroup>
179+ <thead>
180+ <tr>
181+ <th>参数名</th>
182+ <th>输入/输出</th>
183+ <th>描述</th>
184+ </tr></thead>
185+ <tbody>
186+ <tr>
187+ <td>workspace</td>
188+ <td>输入</td>
189+ <td>在Device侧申请的workspace内存地址。</td>
190+ </tr>
191+ <tr>
192+ <td>workspaceSize</td>
193+ <td>输入</td>
194+ <td>在Device侧申请的workspace大小，由第一段接口aclnnThresholdBackwardGetWorkspaceSize获取。</td>
195+ </tr>
196+ <tr>
197+ <td>executor</td>
198+ <td>输入</td>
199+ <td>op执行器，包含了算子计算流程。</td>
200+ </tr>
201+ <tr>
202+ <td>stream</td>
203+ <td>输入</td>
204+ <td>指定执行任务的Stream。</td>
205+ </tr>
206+ </tbody>
207+ </table>
208+ 
209+- **返回值：**
210+ 
211+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../../docs/zh/context/aclnn返回码.md)。
212+ 
213+## 约束说明
214+ 
215+- 确定性计算：
216+ - aclnnThresholdBackward默认确定性实现。
217+ 
218+## 调用示例
219+ 
220+示例代码如下，仅供参考，具体编译和执行过程请参考[编译与运行样例](../../../../docs/zh/context/编译与运行样例.md)。
221+ 
222+```Cpp
223+#include <iostream>
224+#include <vector>
225+#include "acl/acl.h"
226+#include "aclnnop/aclnn_threshold_backward.h"
227+ 
228+#define CHECK_RET(cond, return_expr) \
229+ do { \
230+ if (!(cond)) { \
231+ return_expr; \
232+ } \
233+ } while (0)
234+ 
235+#define LOG_PRINT(message, ...) \
236+ do { \
237+ printf(message, ##__VA_ARGS__); \
238+ } while (0)
239+ 
240+int64_t GetShapeSize(const std::vector<int64_t>& shape) {
241+ int64_t shapeSize = 1;
242+ for (auto i : shape) {
243+ shapeSize *= i;
244+ }
245+ return shapeSize;
246+}
247+ 
248+int Init(int32_t deviceId, aclrtStream* stream) {
249+ // 固定写法，资源初始化
250+ auto ret = aclInit(nullptr);
251+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
252+ ret = aclrtSetDevice(deviceId);
253+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
254+ ret = aclrtCreateStream(stream);
255+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
256+ return 0;
257+}
258+ 
259+template <typename T>
260+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
261+ aclDataType dataType, aclTensor** tensor) {
262+ auto size = GetShapeSize(shape) * sizeof(T);
263+ // 调用aclrtMalloc申请device侧内存
264+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
265+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
266+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
267+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
268+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
269+ 
270+ // 计算连续tensor的strides
271+ std::vector<int64_t> strides(shape.size(), 1);
272+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
273+ strides[i] = shape[i + 1] * strides[i + 1];
274+ }
275+ 
276+ // 调用aclCreateTensor接口创建aclTensor
277+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
278+ shape.data(), shape.size(), *deviceAddr);
279+ return 0;
280+}
281+ 
282+int main() {
283+ // 1. （固定写法）device/stream初始化，参考acl API手册
284+ // 根据自己的实际device填写deviceId
285+ int32_t deviceId = 0;
286+ aclrtStream stream;
287+ auto ret = Init(deviceId, &stream);
288+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
289+ 
290+ // 2. 构造输入与输出，需要根据API的接口自定义构造
291+ std::vector<int64_t> selfShape = {2, 2};
292+ std::vector<int64_t> gradOutputShape = {2, 2};
293+ std::vector<int64_t> outShape = {2, 2};
294+ void* selfDeviceAddr = nullptr;
295+ void* gradOutputDeviceAddr = nullptr;
296+ void* outDeviceAddr = nullptr;
297+ aclTensor* self = nullptr;
298+ aclTensor* gradOutput = nullptr;
299+ aclScalar* threshold = nullptr;
300+ aclTensor* out = nullptr;
301+ std::vector<float> selfHostData = {0.2, 1.2, 2.2, 3.2};
302+ std::vector<float> gradOutputHostData = {4.5, 4.4, 4.3, 4.2};
303+ std::vector<float> outHostData = {0.0, 0.0, 0.0, 0.0};
304+ float thresholdValue = 1.0f;
305+ // 创建self aclTensor
306+ ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self);
307+ CHECK_RET(ret == ACL_SUCCESS, return ret);
308+ // 创建gradOutput aclTensor
309+ ret = CreateAclTensor(gradOutputHostData, gradOutputShape, &gradOutputDeviceAddr, aclDataType::ACL_FLOAT, &gradOutput);
310+ CHECK_RET(ret == ACL_SUCCESS, return ret);
311+ // 创建threshold aclScalar
312+ threshold = aclCreateScalar(&thresholdValue, aclDataType::ACL_FLOAT);
313+ CHECK_RET(threshold != nullptr, return ret);
314+ // 创建out aclTensor
315+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
316+ CHECK_RET(ret == ACL_SUCCESS, return ret);
317+ 
318+ // 3. 调用CANN算子库API，需要修改为具体的API名称
319+ uint64_t workspaceSize = 0;
320+ aclOpExecutor* executor;
321+ // 调用aclnnThresholdBackward第一段接口
322+ ret = aclnnThresholdBackwardGetWorkspaceSize(gradOutput, self, threshold, out, &workspaceSize, &executor);
323+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnThresholdBackwardGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
324+ // 根据第一段接口计算出的workspaceSize申请device内存
325+ void* workspaceAddr = nullptr;
326+ if (workspaceSize > 0) {
327+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
328+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
329+ }
330+ // 调用aclnnThresholdBackward第二段接口
331+ ret = aclnnThresholdBackward(workspaceAddr, workspaceSize, executor, stream);
332+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnThresholdBackward failed. ERROR: %d\n", ret); return ret);
333+ 
334+ // 4. （固定写法）同步等待任务执行结束
335+ ret = aclrtSynchronizeStream(stream);
336+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
337+ 
338+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
339+ auto size = GetShapeSize(outShape);
340+ std::vector<float> resultData(size, 0);
341+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr,
342+ size * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
343+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
344+ for (int64_t i = 0; i < size; i++) {
345+ LOG_PRINT("result[%ld] is: %f\n", i, resultData[i]);
346+ }
347+ 
348+ // 6. 释放aclTensor和aclScalar，需要根据具体API的接口定义修改
349+ aclDestroyTensor(self);
350+ aclDestroyTensor(gradOutput);
351+ aclDestroyScalar(threshold);
352+ aclDestroyTensor(out);
353+ 
354+ // 7. 释放device资源，需要根据具体API的接口定义修改
355+ aclrtFree(selfDeviceAddr);
356+ aclrtFree(gradOutputDeviceAddr);
357+ aclrtFree(outDeviceAddr);
358+ if (workspaceSize > 0) {
359+ aclrtFree(workspaceAddr);
360+ }
361+ aclrtDestroyStream(stream);
362+ aclrtResetDevice(deviceId);
363+ aclFinalize();
364+ return 0;
365+}
366+```
Aexperimental/activation/threshold_grad_v2_d/examples/test_aclnn_threshold_backward.cpp+152-0
@@ -0,0 +1,152 @@
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+#include <iostream>
11+#include <vector>
12+#include "acl/acl.h"
13+#include "aclnnop/aclnn_threshold_backward.h"
14+ 
15+#define CHECK_RET(cond, return_expr) \
16+ do { \
17+ if (!(cond)) { \
18+ return_expr; \
19+ } \
20+ } while (0)
21+ 
22+#define LOG_PRINT(message, ...) \
23+ do { \
24+ printf(message, ##__VA_ARGS__); \
25+ } while (0)
26+ 
27+int64_t GetShapeSize(const std::vector<int64_t>& shape) {
28+ int64_t shapeSize = 1;
29+ for (auto i : shape) {
30+ shapeSize *= i;
31+ }
32+ return shapeSize;
33+}
34+ 
35+int Init(int32_t deviceId, aclrtStream* stream) {
36+ // 固定写法，资源初始化
37+ auto ret = aclInit(nullptr);
38+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
39+ ret = aclrtSetDevice(deviceId);
40+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
41+ ret = aclrtCreateStream(stream);
42+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
43+ return 0;
44+}
45+ 
46+template <typename T>
47+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
48+ aclDataType dataType, aclTensor** tensor) {
49+ auto size = GetShapeSize(shape) * sizeof(T);
50+ // 调用aclrtMalloc申请device侧内存
51+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
52+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
53+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
54+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
55+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
56+ 
57+ // 计算连续tensor的strides
58+ std::vector<int64_t> strides(shape.size(), 1);
59+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
60+ strides[i] = shape[i + 1] * strides[i + 1];
61+ }
62+ 
63+ // 调用aclCreateTensor接口创建aclTensor
64+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
65+ shape.data(), shape.size(), *deviceAddr);
66+ return 0;
67+}
68+ 
69+int main() {
70+ // 1. （固定写法）device/stream初始化，参考acl API手册
71+ // 根据自己的实际device填写deviceId
72+ int32_t deviceId = 0;
73+ aclrtStream stream;
74+ auto ret = Init(deviceId, &stream);
75+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
76+ 
77+ // 2. 构造输入与输出，需要根据API的接口自定义构造
78+ std::vector<int64_t> selfShape = {2, 2};
79+ std::vector<int64_t> gradOutputShape = {2, 2};
80+ std::vector<int64_t> outShape = {2, 2};
81+ void* selfDeviceAddr = nullptr;
82+ void* gradOutputDeviceAddr = nullptr;
83+ void* outDeviceAddr = nullptr;
84+ aclTensor* self = nullptr;
85+ aclTensor* gradOutput = nullptr;
86+ aclScalar* threshold = nullptr;
87+ aclTensor* out = nullptr;
88+ std::vector<float> selfHostData = {0.2, 1.2, 2.2, 3.2};
89+ std::vector<float> gradOutputHostData = {4.5, 4.4, 4.3, 4.2};
90+ std::vector<float> outHostData = {0.0, 0.0, 0.0, 0.0};
91+ float thresholdValue = 1.0f;
92+ // 创建self aclTensor
93+ ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self);
94+ CHECK_RET(ret == ACL_SUCCESS, return ret);
95+ // 创建gradOutput aclTensor
96+ ret = CreateAclTensor(gradOutputHostData, gradOutputShape, &gradOutputDeviceAddr, aclDataType::ACL_FLOAT, &gradOutput);
97+ CHECK_RET(ret == ACL_SUCCESS, return ret);
98+ // 创建threshold aclScalar
99+ threshold = aclCreateScalar(&thresholdValue, aclDataType::ACL_FLOAT);
100+ CHECK_RET(threshold != nullptr, return ret);
101+ // 创建out aclTensor
102+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
103+ CHECK_RET(ret == ACL_SUCCESS, return ret);
104+ 
105+ // 3. 调用CANN算子库API，需要修改为具体的API名称
106+ uint64_t workspaceSize = 0;
107+ aclOpExecutor* executor;
108+ // 调用aclnnThresholdBackward第一段接口
109+ ret = aclnnThresholdBackwardGetWorkspaceSize(gradOutput, self, threshold, out, &workspaceSize, &executor);
110+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnThresholdBackwardGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
111+ // 根据第一段接口计算出的workspaceSize申请device内存
112+ void* workspaceAddr = nullptr;
113+ if (workspaceSize > 0) {
114+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
115+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
116+ }
117+ // 调用aclnnThresholdBackward第二段接口
118+ ret = aclnnThresholdBackward(workspaceAddr, workspaceSize, executor, stream);
119+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnThresholdBackward failed. ERROR: %d\n", ret); return ret);
120+ 
121+ // 4. （固定写法）同步等待任务执行结束
122+ ret = aclrtSynchronizeStream(stream);
123+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
124+ 
125+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
126+ auto size = GetShapeSize(outShape);
127+ std::vector<float> resultData(size, 0);
128+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr,
129+ size * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
130+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
131+ for (int64_t i = 0; i < size; i++) {
132+ LOG_PRINT("result[%ld] is: %f\n", i, resultData[i]);
133+ }
134+ 
135+ // 6. 释放aclTensor和aclScalar，需要根据具体API的接口定义修改
136+ aclDestroyTensor(self);
137+ aclDestroyTensor(gradOutput);
138+ aclDestroyScalar(threshold);
139+ aclDestroyTensor(out);
140+ 
141+ // 7. 释放device资源，需要根据具体API的接口定义修改
142+ aclrtFree(selfDeviceAddr);
143+ aclrtFree(gradOutputDeviceAddr);
144+ aclrtFree(outDeviceAddr);
145+ if (workspaceSize > 0) {
146+ aclrtFree(workspaceAddr);
147+ }
148+ aclrtDestroyStream(stream);
149+ aclrtResetDevice(deviceId);
150+ aclFinalize();
151+ return 0;
152+}
Aexperimental/activation/threshold_grad_v2_d/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 threshold_grad_v2_d ACLNNTYPE aclnn_exclude)
Aexperimental/activation/threshold_grad_v2_d/op_host/op_api/aclnn_threshold_backward.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+#include "aclnn_threshold_backward.h"
11+#include "../../../../../activation/relu_grad/op_host/op_api/relu_grad.h"
12+#include "threshold_grad.h"
13+#include "aclnn_kernels/contiguous.h"
14+#include "aclnn_kernels/common/op_error_check.h"
15+#include "aclnn/aclnn_base.h"
16+#include "opdev/common_types.h"
17+#include "opdev/data_type_utils.h"
18+#include "opdev/shape_utils.h"
19+#include "opdev/format_utils.h"
20+#include "opdev/op_dfx.h"
21+#include "opdev/op_executor.h"
22+#include "opdev/op_log.h"
23+#include "opdev/tensor_view_utils.h"
24+#include "opdev/make_op_executor.h"
25+#include "op_api/op_api_def.h"
26+ 
27+using namespace op;
28+#ifdef __cplusplus
29+extern "C" {
30+#endif
31+ 
32+namespace {
33+float thresholdVal_ = 0.0;
34+// 根据API定义，需要列出所能支持的所有dtype
35+static const std::initializer_list<op::DataType> ASCEND910_DTYPE_DTYPE_SUPPORT_LIST = {
36+ op::DataType::DT_FLOAT, op::DataType::DT_INT32, op::DataType::DT_FLOAT16,
37+ op::DataType::DT_INT8, op::DataType::DT_UINT8};
38+ 
39+static const std::initializer_list<op::DataType> ASCEND910B_DTYPE_DTYPE_SUPPORT_LIST = {
40+ op::DataType::DT_FLOAT, op::DataType::DT_INT32, op::DataType::DT_FLOAT16,
41+ op::DataType::DT_INT8, op::DataType::DT_UINT8, op::DataType::DT_BF16};
42+ 
43+static const std::initializer_list<op::DataType> REGBASE_DTYPE_DTYPE_SUPPORT_LIST = {
44+ op::DataType::DT_FLOAT, op::DataType::DT_INT32, op::DataType::DT_FLOAT16,
45+ op::DataType::DT_INT8, op::DataType::DT_UINT8, op::DataType::DT_BF16, op::DataType::DT_INT64};
46+ 
47+static bool IsFloatEqual(float a, float b) {
48+ return std::abs(a - b) <= std::numeric_limits<float>::epsilon();
49+}
50+ 
51+static bool CheckPtrValid(const aclTensor *gradOutput, const aclTensor *self) {
52+ OP_CHECK_NULL(gradOutput, return false);
53+ OP_CHECK_NULL(self, return false);
54+ return true;
55+}
56+ 
57+static const std::initializer_list<DataType>& GetDtypeSupportList() {
58+ auto curArch = GetCurrentPlatformInfo().GetCurNpuArch();
59+ if (curArch == NpuArch::DAV_2201 || Ops::NN::AclnnUtil::IsRegbase(curArch)) {
60+ if (Ops::NN::AclnnUtil::IsRegbase(curArch) && IsFloatEqual(thresholdVal_, 0.0)) {
61+ // relugrad 支持int64
62+ return REGBASE_DTYPE_DTYPE_SUPPORT_LIST;
63+ } else {
64+ return ASCEND910B_DTYPE_DTYPE_SUPPORT_LIST;
65+ }
66+ } else {
67+ return ASCEND910_DTYPE_DTYPE_SUPPORT_LIST;
68+ }
69+}
70+ 
71+static bool CheckDtypeValid(const aclTensor *gradOutput, const aclTensor *self, const aclTensor *out) {
72+ const auto& supportList = GetDtypeSupportList();
73+ // 检查gradOutput和self数据类型是否在ThresholdBackward算子的支持列表内
74+ OP_CHECK_DTYPE_NOT_SUPPORT(gradOutput, supportList, return false);
75+ OP_CHECK_DTYPE_NOT_SUPPORT(self, supportList, return false);
76+ OP_CHECK_DTYPE_NOT_MATCH(gradOutput, self->GetDataType(), return false);
77+ OP_CHECK_DTYPE_NOT_MATCH(out, self->GetDataType(), return false);
78+ return true;
79+}
80+ 
81+static bool CheckShape(const aclTensor *gradOutput, const aclTensor *self) {
82+ OP_CHECK_MAX_DIM(gradOutput, MAX_SUPPORT_DIMS_NUMS, return false);
83+ OP_CHECK_MAX_DIM(self, MAX_SUPPORT_DIMS_NUMS, return false);
84+ return true;
85+}
86+ 
87+static aclnnStatus CheckParams(const aclTensor *gradOutput, const aclTensor *self, const aclTensor *out) {
88+ // 错误码等DFX方案细化后刷新，错误日志在check接口内打印
89+ // 1. 检查参数是否为空指针
90+ CHECK_RET(CheckPtrValid(gradOutput, self), ACLNN_ERR_PARAM_NULLPTR);
91+ 
92+ // 2. 检查输入的数据类型是否在API支持的数据类型范围之内，需要根据api定义校验
93+ CHECK_RET(CheckDtypeValid(gradOutput, self, out), ACLNN_ERR_PARAM_INVALID);
94+ 
95+ // 3.输入维度校验
96+ CHECK_RET(CheckShape(gradOutput, self), ACLNN_ERR_PARAM_INVALID);
97+ 
98+ return ACLNN_SUCCESS;
99+}
100+}
101+ 
102+aclnnStatus aclnnThresholdBackwardGetWorkspaceSize(const aclTensor *gradOutput, const aclTensor *self,
103+ const aclScalar *threshold, aclTensor *out,
104+ uint64_t *workspaceSize, aclOpExecutor **executor) {
105+ L2_DFX_PHASE_1(aclnnThresholdBackward, DFX_IN(gradOutput, self, threshold), DFX_OUT(out));
106+ // 固定写法，创建OpExecutor
107+ auto uniqueExecutor = CREATE_EXECUTOR();
108+ CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR);
109+ // 固定写法，参数检查
110+ CHECK_RET(threshold != nullptr, ACLNN_ERR_PARAM_NULLPTR);
111+ thresholdVal_ = threshold->ToFloat();
112+ CHECK_RET(out != nullptr, ACLNN_ERR_PARAM_NULLPTR);
113+ CHECK_RET(workspaceSize != nullptr, ACLNN_ERR_PARAM_NULLPTR);
114+ auto ret = CheckParams(gradOutput, self, out);
115+ CHECK_RET(ret == ACLNN_SUCCESS, ret);
116+ // 校验输入shape是否可broadcast
117+ op::Shape broadcastShape;
118+ OP_CHECK_BROADCAST_AND_INFER_SHAPE(self, gradOutput, broadcastShape, return ACLNN_ERR_PARAM_INVALID);
119+ // 校验输出shape是否与推导出的broadcastShape相符
120+ if (broadcastShape != out->GetViewShape()) {
121+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Shape of out should be %s, but current is %s.",
122+ op::ToString(broadcastShape).GetString(), op::ToString(out->GetViewShape()).GetString());
123+ return ACLNN_ERR_PARAM_INVALID;
124+ }
125+ 
126+ // 算子的空tensor在kernel中支持，对标竞品根据算子实际情况补充
127+ if (self->IsEmpty() || gradOutput->IsEmpty()) {
128+ *workspaceSize = 0;
129+ uniqueExecutor.ReleaseTo(executor);
130+ return ACLNN_SUCCESS;
131+ }
132+ 
133+ // 固定写法，将输入gradOutput转换成连续的tensor
134+ auto gradOutputContiguous = l0op::Contiguous(gradOutput, uniqueExecutor.get());
135+ CHECK_RET(gradOutputContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
136+ 
137+ // 固定写法，将输入self转换成连续的tensor
138+ auto selfContiguous = l0op::Contiguous(self, uniqueExecutor.get());
139+ CHECK_RET(selfContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
140+ 
141+ // 调用 ReluGrad or ThresholdGradV2D 算子kernel
142+ const aclTensor* opOut;
143+ if (IsFloatEqual(thresholdVal_, 0.0)) {
144+ opOut = l0op::ReluGrad(gradOutputContiguous, selfContiguous, uniqueExecutor.get());
145+ } else {
146+ opOut = l0op::ThresholdGradV2D(gradOutputContiguous, selfContiguous, thresholdVal_, uniqueExecutor.get());
147+ }
148+ 
149+ CHECK_RET(opOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
150+ // 固定写法，将计算结果拷贝到输出out上，out可能是非连续的tensor
151+ auto viewCopyResult = l0op::ViewCopy(opOut, out, uniqueExecutor.get());
152+ CHECK_RET(viewCopyResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
153+ 
154+ // 固定写法，获取计算过程中需要使用的workspace大小
155+ *workspaceSize = uniqueExecutor->GetWorkspaceSize();
156+ uniqueExecutor.ReleaseTo(executor);
157+ return ACLNN_SUCCESS;
158+}
159+ 
160+aclnnStatus aclnnThresholdBackward(void *workspace, uint64_t workspaceSize, aclOpExecutor *executor,
161+ const aclrtStream stream) {
162+ L2_DFX_PHASE_2(aclnnThresholdBackward);
163+ // 固定写法，调用框架能力，完成计算
164+ return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);
165+}
166+ 
167+#ifdef __cplusplus
168+}
169+#endif
Aexperimental/activation/threshold_grad_v2_d/op_host/op_api/aclnn_threshold_backward.h+71-0
@@ -0,0 +1,71 @@
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+#ifndef OP_API_INC_THRESHOLD_BACKWARG_H_
11+#define OP_API_INC_THRESHOLD_BACKWARG_H_
12+ 
13+#include "aclnn/aclnn_base.h"
14+#include "aclnn_util.h"
15+ 
16+#ifdef __cplusplus
17+extern "C" {
18+#endif
19+ 
20+/**
21+ * @brief aclnnThresholdBackward的第一段接口，根据具体的计算流程，计算workspace大小。
22+ * @domain aclnn_ops_train
23+ *
24+ * 算子功能：完成threshold_forward对应反向
25+ * 计算公式：
26+ * res(i) = gradOutput(i) if self(i) > threshold else 0
27+ *
28+ * 实现说明：
29+ * api计算的基本路径：
30+ * ```mermaid
31+ * graph LR
32+ * A[(self)] --> B([l0op::Contiguous])
33+ * B --> C([l0op::ThresholdGradV2D or l0op::ReluGrad])
34+ * D[(grad_output)] --> E([l0op::Contiguous])
35+ * E --> C
36+ * F[(threshold)] --> C
37+ * C--> G([l0op::ViewCopy])
38+ * G --> H[(out)]
39+ * ```
40+ *
41+ * @param [in] gradOutput: npu device侧的aclTensor，数据类型支持FLOAT、BFLOAT16、FLOAT16、INT32、INT8、UINT8，shape需要与self一致。
42+ * 支持非连续的Tensor，数据格式支持ND，且数据格式需要与self一致。
43+ * @param [in] self: npu device侧的aclTensor，数据类型支持FLOAT、BFLOAT16、FLOAT16、INT32、INT8、UINT8。
44+ * 支持非连续的Tensor，数据格式支持ND。
45+ * @param [in] threshold: host侧的aclScalar，数据类型需要可转换成self与other推导后的数据类型。
46+ * @param [in] out: npu device侧的aclTensor，数据类型支持FLOAT、BFLOAT16、FLOAT16、INT32、INT8、UINT8，shape需要与self一致。
47+ * 支持非连续的Tensor，数据格式支持ND，且数据格式需要与self一致。
48+ * @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。
49+ * @param [out] executor: 返回op执行器，包含算子计算流程。
50+ * @return aclnnStatus: 返回状态码。
51+ */
52+ACLNN_API aclnnStatus aclnnThresholdBackwardGetWorkspaceSize(const aclTensor *gradOutput, const aclTensor *self,
53+ const aclScalar *threshold, aclTensor *out,
54+ uint64_t *workspaceSize, aclOpExecutor **executor);
55+/**
56+ * @brief aclnnAdd的第二段接口，用于执行计算。
57+ *
58+ * @param [in] workspace: 在npu device侧申请的workspace内存起址。
59+ * @param [in] workspaceSize: 在npu device侧申请的workspace大小，由第一段接口aclnnAddGetWorkspaceSize获取。
60+ * @param [in] stream: acl stream流。
61+ * @param [in] executor: op执行器，包含了算子计算流程。
62+ * @return aclnnStatus: 返回状态码。
63+ */
64+ACLNN_API aclnnStatus aclnnThresholdBackward(void *workspace, uint64_t workspaceSize, aclOpExecutor *executor,
65+ const aclrtStream stream);
66+ 
67+#ifdef __cplusplus
68+}
69+#endif
70+ 
71+#endif // OP_API_INC_THRESHOLD_BACKWARG_H_
Aexperimental/activation/threshold_grad_v2_d/op_host/op_api/threshold_grad.cpp+41-0
@@ -0,0 +1,41 @@
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 "threshold_grad.h"
12+#include "aclnn_kernels/common/op_error_check.h"
13+#include "opdev/op_log.h"
14+#include "opdev/op_executor.h"
15+#include "opdev/make_op_executor.h"
16+#include "opdev/shape_utils.h"
17+#include "opdev/shape_utils.h"
18+#include "opdev/op_def.h"
19+#include "opdev/op_dfx.h"
20+ 
21+using namespace op;
22+ 
23+namespace l0op {
24+OP_TYPE_REGISTER(ThresholdGradV2D);
25+ 
26+const aclTensor *ThresholdGradV2D(const aclTensor *gradOutput, const aclTensor *self, float threshold,
27+ aclOpExecutor *executor) {
28+ L0_DFX(ThresholdGradV2D, gradOutput, self, threshold);
29+ // 根据推导出的输出shape申请输出tensor
30+ Shape broadcastShape;
31+ OP_CHECK_BROADCAST_AND_INFER_SHAPE(self, gradOutput, broadcastShape, return nullptr);
32+ // 第一个参数是输出shape，第二个参数是输出的dtype
33+ auto out = executor->AllocTensor(broadcastShape, self->GetDataType());
34+ 
35+ ADD_TO_LAUNCHER_LIST_AICORE(ThresholdGradV2D,
36+ OP_INPUT(gradOutput, self),
37+ OP_OUTPUT(out),
38+ OP_ATTR(threshold));
39+ return out;
40+}
41+} // namespace l0op
Aexperimental/activation/threshold_grad_v2_d/op_host/op_api/threshold_grad.h+22-0
@@ -0,0 +1,22 @@
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+#ifndef PTA_NPU_OP_API_INC_LEVEL0_OP_THRESHOLD_GRAD_H_
11+#define PTA_NPU_OP_API_INC_LEVEL0_OP_THRESHOLD_GRAD_H_
12+ 
13+#include "opdev/op_executor.h"
14+ 
15+namespace l0op {
16+const aclTensor *ThresholdGradV2D(const aclTensor *gradOutput,
17+ const aclTensor *self,
18+ float threshold,
19+ aclOpExecutor *executor);
20+}
21+ 
22+#endif // PTA_NPU_OP_API_INC_LEVEL0_OP_THRESHOLD_GRAD_H_
Aexperimental/activation/threshold_grad_v2_d/op_host/threshold_grad_v2_d_def.cpp+50-0
@@ -0,0 +1,50 @@
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 threshold_grad_v2_d.cpp
13+ * \brief
14+ */
15+#include "register/op_def_registry.h"
16+ 
17+namespace ops {
18+class ThresholdGradV2D : public OpDef {
19+public:
20+ explicit ThresholdGradV2D(const char* name) : OpDef(name)
21+ {
22+ this->Input("input_gradient")
23+ .ParamType(REQUIRED)
24+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16, ge::DT_INT32, ge::DT_INT8, ge::DT_UINT8})
25+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
26+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
27+ this->Input("input_feature")
28+ .ParamType(REQUIRED)
29+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16, ge::DT_INT32, ge::DT_INT8, ge::DT_UINT8})
30+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
31+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
32+ this->Output("output_backprops")
33+ .ParamType(REQUIRED)
34+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16, ge::DT_INT32, ge::DT_INT8, ge::DT_UINT8})
35+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
36+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
37+ this->Attr("threshold").AttrType(REQUIRED).Float(1.0);
38+ OpAICoreConfig aicoreConfig;
39+ aicoreConfig.DynamicCompileStaticFlag(true)
40+ .DynamicFormatFlag(false)
41+ .DynamicRankSupportFlag(true)
42+ .DynamicShapeSupportFlag(true)
43+ .NeedCheckSupportFlag(false)
44+ .PrecisionReduceFlag(true)
45+ .ExtendCfgInfo("opFile.value", "threshold_grad_v2_d"); // 这里制定的值会对应到kernel入口文件名.cpp
46+ this->AICore().AddConfig("ascend910b", aicoreConfig); // 其他的soc版本补充部分配置项
47+ }
48+};
49+OP_ADD(ThresholdGradV2D); // 添加算子信息库
50+} // namespace ops
Aexperimental/activation/threshold_grad_v2_d/op_host/threshold_grad_v2_d_infershape.cpp+43-0
@@ -0,0 +1,43 @@
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 threshold_grad_v2_d_infershape.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+static constexpr int64_t IDX_0 = 0;
22+ 
23+static ge::graphStatus InferShapeThresholdGradV2D(gert::InferShapeContext* context)
24+{
25+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
26+ OP_LOGD(context->GetNodeName(), "Begin to do InferShapeThresholdGradV2D");
27+ 
28+ // get input shapes
29+ const gert::Shape* xShape = context->GetInputShape(IDX_0);
30+ OP_CHECK_NULL_WITH_CONTEXT(context, xShape);
31+ 
32+ // get output shapes
33+ gert::Shape* yShape = context->GetOutputShape(IDX_0);
34+ OP_CHECK_NULL_WITH_CONTEXT(context, yShape);
35+ 
36+ // 填充输出shape大小
37+ *yShape = *xShape;
38+ OP_LOGD(context->GetNodeName(), "End to do InferShapeThresholdGradV2D");
39+ return GRAPH_SUCCESS;
40+}
41+ 
42+IMPL_OP_INFERSHAPE(ThresholdGradV2D).InferShape(InferShapeThresholdGradV2D);
43+}
Aexperimental/activation/threshold_grad_v2_d/op_host/threshold_grad_v2_d_tiling.cpp+201-0
@@ -0,0 +1,201 @@
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 threshold_grad_v2_d_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 "tiling/platform/platform_ascendc.h"
22+#include "register/op_impl_registry.h"
23+#include "../op_kernel/threshold_grad_v2_d_tiling_data.h"
24+#include "../op_kernel/threshold_grad_v2_d_tiling_key.h"
25+ 
26+namespace optiling {
27+ 
28+using namespace Ops::NN::OpTiling;
29+ 
30+#define UB_NUM_F32_F16 7U
31+#define UB_NUM_BF16 10U
32+#define UB_NUM_INT8_UINT8 14U
33+#define UB_NUM_INT32 8U
34+#define BLOCK_SIZE 256U
35+constexpr uint32_t BUFFER_NUM = 2;
36+constexpr uint32_t WS_SYS_SIZE = 0;
37+ 
38+struct ThresholdGradV2DCompileInfo {};
39+ 
40+static ge::graphStatus TilingParseForThresholdGradV2D([[maybe_unused]] gert::TilingParseContext* context)
41+{
42+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
43+ return ge::GRAPH_SUCCESS;
44+}
45+ 
46+// 获取平台信息如ubSize, coreNum
47+static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t& ubSize, int64_t& coreNum)
48+{
49+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
50+ // 获取ubsize coreNum
51+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
52+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
53+ coreNum = ascendcPlatform.GetCoreNumAiv();
54+ OP_CHECK_IF(coreNum == 0, OP_LOGE(context, "coreNum is 0"), return ge::GRAPH_FAILED);
55+ OP_CHECK_IF(ubSize == 0, OP_LOGE(context, "ubSize is 0"), return ge::GRAPH_FAILED);
56+ return ge::GRAPH_SUCCESS;
57+}
58+ 
59+static ge::graphStatus GetWorkspaceSize(gert::TilingContext* context)
60+{
61+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
62+ size_t usrSize = 0;
63+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
64+ uint32_t sysWorkspaceSize = ascendcPlatform.GetLibApiWorkSpaceSize();
65+ size_t* currentWorkspace = context->GetWorkspaceSizes(
66+ 1); // 通过框架获取workspace的指针，GetWorkspaceSizes入参为所需workspace的块数。当前限制使用一块。
67+ currentWorkspace[0] = usrSize + sysWorkspaceSize;
68+ return ge::GRAPH_SUCCESS;
69+}
70+ 
71+static ge::graphStatus GetShapeAttrsInfo(
72+ gert::TilingContext* context, uint64_t ubSize, uint64_t& inputNum, uint64_t& inputBytes,
73+ uint64_t& tileBlockNum, uint64_t& tileDataNum, uint64_t& inputLengthAlgin)
74+{
75+ OP_CHECK_IF(
76+ context == nullptr || context->GetInputShape(0) == nullptr, OP_LOGE(context, "context is nullptr"),
77+ return ge::GRAPH_FAILED);
78+ inputNum = context->GetInputShape(0)->GetStorageShape().GetShapeSize();
79+ uint32_t typeLength = 0;
80+ ge::TypeUtils::GetDataTypeLength(context->GetInputDesc(0)->GetDataType(), typeLength);
81+ uint64_t inputLength = inputNum * typeLength;
82+ if (inputNum == 0) {
83+ OP_LOGE(context, "inputNum is 0");
84+ return ge::GRAPH_FAILED;
85+ }
86+ inputBytes = inputLength / inputNum;
87+ uint64_t ubDataNumber;
88+ if (context->GetInputDesc(0)->GetDataType() == ge::DT_FLOAT || context->GetInputDesc(0)->GetDataType() == ge::DT_FLOAT16) {
89+ ubDataNumber = UB_NUM_F32_F16;
90+ } else if (context->GetInputDesc(0)->GetDataType() == ge::DT_INT32) {
91+ ubDataNumber = UB_NUM_INT32;
92+ } else if (context->GetInputDesc(0)->GetDataType() == ge::DT_BF16) {
93+ ubDataNumber = UB_NUM_BF16;
94+ } else {
95+ ubDataNumber = UB_NUM_INT8_UINT8;
96+ }
97+ if (ubDataNumber == 0 || BLOCK_SIZE == 0) {
98+ OP_LOGE(context, "ubDataNumber or BLOCK_SIZE is 0");
99+ return ge::GRAPH_FAILED;
100+ }
101+ tileBlockNum = (ubSize / BLOCK_SIZE) / ubDataNumber;
102+ if (inputBytes == 0) {
103+ OP_LOGE(context, "inputBytes is 0");
104+ return ge::GRAPH_FAILED;
105+ }
106+ tileDataNum = (tileBlockNum * BLOCK_SIZE) / inputBytes;
107+ inputLengthAlgin = (((inputLength + BLOCK_SIZE - 1) / BLOCK_SIZE) * BLOCK_SIZE);
108+ return ge::GRAPH_SUCCESS;
109+}
110+ 
111+static ge::graphStatus CalculateCoreBlockNums(
112+ gert::TilingContext* context, uint64_t inputLengthAlgin, int64_t coreNum, uint64_t tileBlockNum, uint64_t inputBytes, uint64_t tileDataNum,
113+ uint64_t& smallCoreDataNum, uint64_t& bigCoreDataNum, uint64_t& smallTailDataNum, uint64_t& bigTailDataNum, uint64_t& finalSmallTileNum, uint64_t& finalBigTileNum,
114+ uint64_t& tailBlockNum)
115+{
116+ if (0 == BLOCK_SIZE || 0 == coreNum || 0 == tileBlockNum || 0 == inputBytes) {
117+ OP_LOGE(context, "BLOCK_SIZE or coreNum or tileBlockNum or inputBytes is 0");
118+ return ge::GRAPH_FAILED;
119+ }
120+ uint64_t everyCoreInputBlockNum = inputLengthAlgin / BLOCK_SIZE / coreNum;
121+ tailBlockNum = (inputLengthAlgin / BLOCK_SIZE) % coreNum;
122+ smallCoreDataNum = everyCoreInputBlockNum * BLOCK_SIZE / inputBytes;
123+ uint64_t smallTileNum = everyCoreInputBlockNum / tileBlockNum;
124+ finalSmallTileNum = (everyCoreInputBlockNum % tileBlockNum) == 0 ? smallTileNum : smallTileNum + 1;
125+ smallTailDataNum = smallCoreDataNum - (tileDataNum * smallTileNum);
126+ smallTailDataNum = smallTailDataNum == 0 ? tileDataNum : smallTailDataNum;
127+ 
128+ everyCoreInputBlockNum += 1;
129+ bigCoreDataNum = everyCoreInputBlockNum * BLOCK_SIZE / inputBytes;
130+ uint64_t bigTileNum = everyCoreInputBlockNum / tileBlockNum;
131+ finalBigTileNum = (everyCoreInputBlockNum % tileBlockNum) == 0 ? bigTileNum : bigTileNum + 1;
132+ bigTailDataNum = bigCoreDataNum - tileDataNum * bigTileNum;
133+ bigTailDataNum = bigTailDataNum == 0 ? tileDataNum : bigTailDataNum;
134+ 
135+ return ge::GRAPH_SUCCESS;
136+}
137+ 
138+// tiling 分发入口
139+static ge::graphStatus ThresholdGradV2DTilingFunc(gert::TilingContext* context)
140+{
141+ // 1、获取平台运行信息
142+ uint64_t ubSize;
143+ int64_t coreNum;
144+ ge::graphStatus ret = GetPlatformInfo(context, ubSize, coreNum);
145+ OP_CHECK_IF(ret != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetPlatformInfo error"), return ge::GRAPH_FAILED);
146+ // 2、获取shape、属性信息
147+ uint64_t inputNum, inputBytes, tileBlockNum, tileDataNum, inputLengthAlgin;
148+ ret = GetShapeAttrsInfo(context, ubSize, inputNum, inputBytes, tileBlockNum, tileDataNum, inputLengthAlgin);
149+ OP_CHECK_IF(ret != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetShapeAttrsInfo error"), return ge::GRAPH_FAILED);
150+ // 3、获取WorkspaceSize信息
151+ OP_CHECK_IF(
152+ GetWorkspaceSize(context) != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetWorkspaceSize error"),
153+ return ge::GRAPH_FAILED);
154+ // 4、设置tiling信息
155+ ThresholdGradV2DTilingData* tiling = context->GetTilingData<ThresholdGradV2DTilingData>();
156+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
157+ OP_CHECK_IF(
158+ memset_s(tiling, sizeof(ThresholdGradV2DTilingData), 0, sizeof(ThresholdGradV2DTilingData)) != EOK,
159+ OP_LOGE(context, "set tiling data error"), return ge::GRAPH_FAILED);
160+ 
161+ if (tileDataNum >= inputNum) {
162+ coreNum = 1;
163+ } else {
164+ coreNum = (static_cast<uint64_t>(coreNum) < inputLengthAlgin / BLOCK_SIZE) ? coreNum : inputLengthAlgin / BLOCK_SIZE;
165+ }
166+ // 计算每个core处理的数据块数
167+ uint64_t smallCoreDataNum, bigCoreDataNum, smallTailDataNum, bigTailDataNum, finalSmallTileNum, finalBigTileNum, tailBlockNum;
168+ ret = CalculateCoreBlockNums(
169+ context, inputLengthAlgin, coreNum, tileBlockNum, inputBytes, tileDataNum, smallCoreDataNum, bigCoreDataNum,
170+ smallTailDataNum, bigTailDataNum, finalSmallTileNum, finalBigTileNum, tailBlockNum);
171+ OP_CHECK_IF(ret != ge::GRAPH_SUCCESS, OP_LOGE(context, "CalculateCoreBlockNums error"), return ge::GRAPH_FAILED);
172+ // 设置tiling数据
173+ tiling->smallCoreDataNum = static_cast<uint64_t>(smallCoreDataNum);
174+ tiling->bigCoreDataNum = static_cast<uint64_t>(bigCoreDataNum);
175+ tiling->tileDataNum = static_cast<uint64_t>(tileDataNum);
176+ tiling->smallTailDataNum = static_cast<uint64_t>(smallTailDataNum);
177+ tiling->bigTailDataNum = static_cast<uint64_t>(bigTailDataNum);
178+ tiling->finalSmallTileNum = static_cast<uint64_t>(finalSmallTileNum);
179+ tiling->finalBigTileNum = static_cast<uint64_t>(finalBigTileNum);
180+ tiling->tailBlockNum = static_cast<uint64_t>(tailBlockNum);
181+ 
182+ float threshold = 1.0f;
183+ auto attrs = context->GetAttrs();
184+ if (attrs) {
185+ const float* attrS = attrs->GetFloat(0);
186+ if (attrS != nullptr) {
187+ threshold = *attrS;
188+ }
189+ }
190+ tiling->threshold = threshold;
191+ 
192+ context->SetBlockDim(coreNum);
193+ uint64_t tilingKey = 0;
194+ tilingKey = GET_TPL_TILING_KEY(ELEMENTWISE_TPL_SCH_MODE_0);
195+ context->SetTilingKey(tilingKey);
196+ return ge::GRAPH_SUCCESS;
197+}
198+ 
199+// tiling注册入口.
200+IMPL_OP_OPTILING(ThresholdGradV2D).Tiling(ThresholdGradV2DTilingFunc).TilingParse<ThresholdGradV2DCompileInfo>(TilingParseForThresholdGradV2D);
201+} // namespace optiling
Aexperimental/activation/threshold_grad_v2_d/op_kernel/threshold_grad_v2_d.cpp+28-0
@@ -0,0 +1,28 @@
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 threshold_grad_v2_d.cpp
13+ * \brief
14+ */
15+ 
16+#include "threshold_grad_v2_d.h"
17+ 
18+ 
19+template <uint32_t schMode>
20+__global__ __aicore__ void threshold_grad_v2_d(GM_ADDR input_gradient, GM_ADDR input_feature, GM_ADDR output_backprops, GM_ADDR workspace, GM_ADDR tiling)
21+{
22+ REGISTER_TILING_DEFAULT(ThresholdGradV2DTilingData);
23+ GET_TILING_DATA_WITH_STRUCT(ThresholdGradV2DTilingData, tilingData, tiling);
24+ NsThresholdGradV2D::KernelThresholdGradV2D<DTYPE_INPUT_GRADIENT> op; // 算子kernel实例获取
25+ op.Init(input_gradient, input_feature, output_backprops, tilingData.smallCoreDataNum, tilingData.bigCoreDataNum, tilingData.finalBigTileNum, tilingData.finalSmallTileNum, tilingData.tileDataNum,
26+ tilingData.smallTailDataNum, tilingData.bigTailDataNum, tilingData.tailBlockNum, tilingData.threshold);
27+ op.Process();
28+}
Aexperimental/activation/threshold_grad_v2_d/op_kernel/threshold_grad_v2_d.h+221-0
@@ -0,0 +1,221 @@
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 threshold_grad_v2_d.h
13+ * \brief
14+ */
15+#ifndef THRESHOLD_GRAD_V2_D_H
16+#define THRESHOLD_GRAD_V2_D_H
17+ 
18+#include "kernel_operator.h"
19+#include "kernel_tiling/kernel_tiling.h"
20+#include "threshold_grad_v2_d_tiling_data.h"
21+#include "threshold_grad_v2_d_tiling_key.h"
22+ 
23+#include "kernel_operator.h"
24+ 
25+namespace NsThresholdGradV2D {
26+ 
27+using namespace AscendC;
28+ 
29+constexpr int32_t BUFFER_NUM = 2;
30+ 
31+template <typename TYPE_INPUT_GRADIENT>
32+class KernelThresholdGradV2D {
33+public:
34+ __aicore__ inline KernelThresholdGradV2D(){};
35+ 
36+ __aicore__ inline void Init(GM_ADDR input_gradient, GM_ADDR input_feature, GM_ADDR output_backprops, uint64_t smallCoreDataNum, uint64_t bigCoreDataNum, uint64_t finalBigTileNum,
37+ uint64_t finalSmallTileNum, uint64_t tileDataNum, uint64_t smallTailDataNum, uint64_t bigTailDataNum, uint64_t tailBlockNum, float threshold);
38+ __aicore__ inline void Process();
39+ 
40+private:
41+ __aicore__ inline void CopyIn(int32_t progress);
42+ __aicore__ inline void CopyOut(int32_t progress);
43+ __aicore__ inline void Compute(int32_t progress);
44+ 
45+private:
46+ AscendC::TPipe pipe;
47+ AscendC::TQue<AscendC::TPosition::VECIN, BUFFER_NUM> inQueueG, inQueueF;
48+ AscendC::TQue<AscendC::TPosition::VECOUT, BUFFER_NUM> outQueueout;
49+ AscendC::TBuf<AscendC::TPosition::VECCALC> tmpQueue1, tmpQueue2, tmpQueueMask;
50+ 
51+ AscendC::GlobalTensor<TYPE_INPUT_GRADIENT> input_gradientGm, input_featureGm, output_backpropsGm;
52+ uint64_t coreDataNum = 0;
53+ uint64_t tileNum = 0;
54+ uint64_t tileDataNum = 0;
55+ uint64_t tailDataNum = 0;
56+ uint64_t processDataNum = 0;
57+ float thresholdValue = 0.0f;
58+};
59+ 
60+template <typename TYPE_INPUT_GRADIENT>
61+__aicore__ inline void KernelThresholdGradV2D<TYPE_INPUT_GRADIENT>::Init(GM_ADDR input_gradient, GM_ADDR input_feature, GM_ADDR output_backprops, uint64_t smallCoreDataNum, uint64_t bigCoreDataNum, uint64_t finalBigTileNum,
62+ uint64_t finalSmallTileNum, uint64_t tileDataNum, uint64_t smallTailDataNum, uint64_t bigTailDataNum, uint64_t tailBlockNum, float threshold)
63+{
64+ ASSERT(AscendC::GetBlockNum() != 0 && "block dim can not be zero!");
65+ uint64_t coreId = AscendC::GetBlockIdx();
66+ uint64_t globalBufferIndex = bigCoreDataNum * coreId;
67+ this->tileDataNum = tileDataNum;
68+ this->thresholdValue = threshold;
69+ if (coreId < tailBlockNum) {
70+ this->coreDataNum = bigCoreDataNum;
71+ this->tileNum = finalBigTileNum;
72+ this->tailDataNum = bigTailDataNum;
73+ } else {
74+ this->coreDataNum = smallCoreDataNum;
75+ this->tileNum = finalSmallTileNum;
76+ this->tailDataNum = smallTailDataNum;
77+ globalBufferIndex -= (bigCoreDataNum - smallCoreDataNum) * (coreId - tailBlockNum);
78+ }
79+ input_gradientGm.SetGlobalBuffer((__gm__ TYPE_INPUT_GRADIENT *)input_gradient + globalBufferIndex, this->coreDataNum);
80+ input_featureGm.SetGlobalBuffer((__gm__ TYPE_INPUT_GRADIENT *)input_feature + globalBufferIndex, this->coreDataNum);
81+ output_backpropsGm.SetGlobalBuffer((__gm__ TYPE_INPUT_GRADIENT *)output_backprops + globalBufferIndex, this->coreDataNum);
82+ 
83+ pipe.InitBuffer(inQueueG, BUFFER_NUM, this->tileDataNum * sizeof(TYPE_INPUT_GRADIENT));
84+ pipe.InitBuffer(inQueueF, BUFFER_NUM, this->tileDataNum * sizeof(TYPE_INPUT_GRADIENT));
85+ pipe.InitBuffer(outQueueout, BUFFER_NUM, this->tileDataNum * sizeof(TYPE_INPUT_GRADIENT));
86+ pipe.InitBuffer(tmpQueueMask, this->tileDataNum * sizeof(uint8_t));
87+ if (std::is_same_v<TYPE_INPUT_GRADIENT, int32_t> || std::is_same_v<TYPE_INPUT_GRADIENT, bfloat16_t>) {
88+ pipe.InitBuffer(tmpQueue1, this->tileDataNum * sizeof(float));
89+ } else if (std::is_same_v<TYPE_INPUT_GRADIENT, uint8_t> || std::is_same_v<TYPE_INPUT_GRADIENT, int8_t>) {
90+ pipe.InitBuffer(tmpQueue1, this->tileDataNum * sizeof(half));
91+ pipe.InitBuffer(tmpQueue2, this->tileDataNum * sizeof(half));
92+ }
93+}
94+ 
95+template <typename TYPE_INPUT_GRADIENT>
96+__aicore__ inline void KernelThresholdGradV2D<TYPE_INPUT_GRADIENT>::CopyIn(int32_t progress)
97+{
98+ AscendC::LocalTensor<TYPE_INPUT_GRADIENT> gLocal = inQueueG.AllocTensor<TYPE_INPUT_GRADIENT>();
99+ AscendC::LocalTensor<TYPE_INPUT_GRADIENT> fLocal = inQueueF.AllocTensor<TYPE_INPUT_GRADIENT>();
100+ AscendC::DataCopy(gLocal, input_gradientGm[progress * this->tileDataNum], this->processDataNum);
101+ AscendC::DataCopy(fLocal, input_featureGm[progress * this->tileDataNum], this->processDataNum);
102+ inQueueG.EnQue(gLocal);
103+ inQueueF.EnQue(fLocal);
104+}
105+ 
106+template <typename TYPE_INPUT_GRADIENT>
107+__aicore__ inline void KernelThresholdGradV2D<TYPE_INPUT_GRADIENT>::CopyOut(int32_t progress)
108+{
109+ AscendC::LocalTensor<TYPE_INPUT_GRADIENT> outLocal = outQueueout.DeQue<TYPE_INPUT_GRADIENT>();
110+ AscendC::DataCopy(output_backpropsGm[progress * this->tileDataNum], outLocal, this->processDataNum);
111+ outQueueout.FreeTensor(outLocal);
112+}
113+ 
114+template <typename TYPE_INPUT_GRADIENT>
115+__aicore__ inline void KernelThresholdGradV2D<TYPE_INPUT_GRADIENT>::Compute(int32_t progress)
116+{
117+ if (std::is_same_v<TYPE_INPUT_GRADIENT, int8_t>) {
118+ AscendC::LocalTensor<int8_t> gLocal = inQueueG.DeQue<int8_t>();
119+ AscendC::LocalTensor<int8_t> fLocal = inQueueF.DeQue<int8_t>();
120+ AscendC::LocalTensor<int8_t> outLocal = outQueueout.AllocTensor<int8_t>();
121+ AscendC::LocalTensor<half> tmp1Local = tmpQueue1.AllocTensor<half>();
122+ AscendC::LocalTensor<float> tmp2Local = tmpQueue2.AllocTensor<float>();
123+ AscendC::LocalTensor<uint8_t> maskLocal = tmpQueueMask.AllocTensor<uint8_t>();
124+ AscendC::Cast(tmp1Local, fLocal, AscendC::RoundMode::CAST_NONE, this->processDataNum);
125+ AscendC::Cast(tmp2Local, tmp1Local, AscendC::RoundMode::CAST_NONE, this->processDataNum);
126+ AscendC::CompareScalar(maskLocal, tmp2Local, static_cast<float>(this->thresholdValue), AscendC::CMPMODE::GT, this->processDataNum);
127+ AscendC::Cast(tmp1Local, gLocal, AscendC::RoundMode::CAST_NONE, this->processDataNum);
128+ AscendC::Cast(tmp2Local, tmp1Local, AscendC::RoundMode::CAST_NONE, this->processDataNum);
129+ AscendC::Select(tmp2Local, maskLocal, tmp2Local, static_cast<float>(0.0), AscendC::SELMODE::VSEL_TENSOR_SCALAR_MODE, this->processDataNum);
130+ AscendC::Cast(tmp1Local, tmp2Local, AscendC::RoundMode::CAST_NONE, this->processDataNum);
131+ AscendC::Cast(outLocal, tmp1Local, AscendC::RoundMode::CAST_TRUNC, this->processDataNum);
132+ outQueueout.EnQue<int8_t>(outLocal);
133+ inQueueG.FreeTensor(gLocal);
134+ inQueueF.FreeTensor(fLocal);
135+ } else if (std::is_same_v<TYPE_INPUT_GRADIENT, uint8_t>) {
136+ AscendC::LocalTensor<uint8_t> gLocal = inQueueG.DeQue<uint8_t>();
137+ AscendC::LocalTensor<uint8_t> fLocal = inQueueF.DeQue<uint8_t>();
138+ AscendC::LocalTensor<uint8_t> outLocal = outQueueout.AllocTensor<uint8_t>();
139+ AscendC::LocalTensor<half> tmp1Local = tmpQueue1.AllocTensor<half>();
140+ AscendC::LocalTensor<float> tmp2Local = tmpQueue2.AllocTensor<float>();
141+ AscendC::LocalTensor<uint8_t> maskLocal = tmpQueueMask.AllocTensor<uint8_t>();
142+ AscendC::Cast(tmp1Local, fLocal, AscendC::RoundMode::CAST_NONE, this->processDataNum);
143+ AscendC::Cast(tmp2Local, tmp1Local, AscendC::RoundMode::CAST_NONE, this->processDataNum);
144+ AscendC::CompareScalar(maskLocal, tmp2Local, static_cast<float>(this->thresholdValue), AscendC::CMPMODE::GT, this->processDataNum);
145+ AscendC::Cast(tmp1Local, gLocal, AscendC::RoundMode::CAST_NONE, this->processDataNum);
146+ AscendC::Cast(tmp2Local, tmp1Local, AscendC::RoundMode::CAST_NONE, this->processDataNum);
147+ AscendC::Select(tmp2Local, maskLocal, tmp2Local, static_cast<float>(0.0), AscendC::SELMODE::VSEL_TENSOR_SCALAR_MODE, this->processDataNum);
148+ AscendC::Cast(tmp1Local, tmp2Local, AscendC::RoundMode::CAST_NONE, this->processDataNum);
149+ AscendC::Cast(outLocal, tmp1Local, AscendC::RoundMode::CAST_TRUNC, this->processDataNum);
150+ outQueueout.EnQue<uint8_t>(outLocal);
151+ inQueueG.FreeTensor(gLocal);
152+ inQueueF.FreeTensor(fLocal);
153+ } else if (std::is_same_v<TYPE_INPUT_GRADIENT, int32_t>) {
154+ AscendC::LocalTensor<int32_t> gLocal = inQueueG.DeQue<int32_t>();
155+ AscendC::LocalTensor<int32_t> fLocal = inQueueF.DeQue<int32_t>();
156+ AscendC::LocalTensor<int32_t> outLocal = outQueueout.AllocTensor<int32_t>();
157+ AscendC::LocalTensor<float> tmp1Local = tmpQueue1.AllocTensor<float>();
158+ AscendC::LocalTensor<uint8_t> maskLocal = tmpQueueMask.AllocTensor<uint8_t>();
159+ AscendC::Cast(tmp1Local, fLocal, AscendC::RoundMode::CAST_NONE, this->processDataNum);
160+ AscendC::CompareScalar(maskLocal, tmp1Local, static_cast<float>(this->thresholdValue), AscendC::CMPMODE::GT, this->processDataNum);
161+ AscendC::Cast(tmp1Local, gLocal, AscendC::RoundMode::CAST_NONE, this->processDataNum);
162+ AscendC::Select(tmp1Local, maskLocal, tmp1Local, static_cast<float>(0.0), AscendC::SELMODE::VSEL_TENSOR_SCALAR_MODE, this->processDataNum);
163+ AscendC::Cast(outLocal, tmp1Local, AscendC::RoundMode::CAST_TRUNC, this->processDataNum);
164+ outQueueout.EnQue<int32_t>(outLocal);
165+ inQueueG.FreeTensor(gLocal);
166+ inQueueF.FreeTensor(fLocal);
167+ } else if (std::is_same_v<TYPE_INPUT_GRADIENT, bfloat16_t>) {
168+ AscendC::LocalTensor<bfloat16_t> gLocal = inQueueG.DeQue<bfloat16_t>();
169+ AscendC::LocalTensor<bfloat16_t> fLocal = inQueueF.DeQue<bfloat16_t>();
170+ AscendC::LocalTensor<bfloat16_t> outLocal = outQueueout.AllocTensor<bfloat16_t>();
171+ AscendC::LocalTensor<float> tmp1Local = tmpQueue1.AllocTensor<float>();
172+ AscendC::LocalTensor<uint8_t> maskLocal = tmpQueueMask.AllocTensor<uint8_t>();
173+ AscendC::Cast(tmp1Local, fLocal, AscendC::RoundMode::CAST_NONE, this->processDataNum);
174+ AscendC::CompareScalar(maskLocal, tmp1Local, static_cast<float>(this->thresholdValue), AscendC::CMPMODE::GT, this->processDataNum);
175+ AscendC::Cast(tmp1Local, gLocal, AscendC::RoundMode::CAST_NONE, this->processDataNum);
176+ AscendC::Select(tmp1Local, maskLocal, tmp1Local, static_cast<float>(0.0), AscendC::SELMODE::VSEL_TENSOR_SCALAR_MODE, this->processDataNum);
177+ AscendC::Cast(outLocal, tmp1Local, AscendC::RoundMode::CAST_RINT, this->processDataNum);
178+ outQueueout.EnQue<bfloat16_t>(outLocal);
179+ inQueueG.FreeTensor(gLocal);
180+ inQueueF.FreeTensor(fLocal);
181+ } else if (std::is_same_v<TYPE_INPUT_GRADIENT, float>) {
182+ AscendC::LocalTensor<float> gLocal = inQueueG.DeQue<float>();
183+ AscendC::LocalTensor<float> fLocal = inQueueF.DeQue<float>();
184+ AscendC::LocalTensor<float> outLocal = outQueueout.AllocTensor<float>();
185+ AscendC::LocalTensor<uint8_t> maskLocal = tmpQueueMask.AllocTensor<uint8_t>();
186+ AscendC::CompareScalar(maskLocal, fLocal, static_cast<float>(this->thresholdValue), AscendC::CMPMODE::GT, this->processDataNum);
187+ AscendC::Select(outLocal, maskLocal, gLocal, static_cast<float>(0.0), AscendC::SELMODE::VSEL_TENSOR_SCALAR_MODE, this->processDataNum);
188+ outQueueout.EnQue<float>(outLocal);
189+ inQueueG.FreeTensor(gLocal);
190+ inQueueF.FreeTensor(fLocal);
191+ } else {
192+ AscendC::LocalTensor<half> gLocal = inQueueG.DeQue<half>();
193+ AscendC::LocalTensor<half> fLocal = inQueueF.DeQue<half>();
194+ AscendC::LocalTensor<half> outLocal = outQueueout.AllocTensor<half>();
195+ AscendC::LocalTensor<uint8_t> maskLocal = tmpQueueMask.AllocTensor<uint8_t>();
196+ AscendC::CompareScalar(maskLocal, fLocal, static_cast<half>(this->thresholdValue), AscendC::CMPMODE::GT, this->processDataNum);
197+ AscendC::Select(outLocal, maskLocal, gLocal, static_cast<half>(0.0), AscendC::SELMODE::VSEL_TENSOR_SCALAR_MODE, this->processDataNum);
198+ outQueueout.EnQue<half>(outLocal);
199+ inQueueG.FreeTensor(gLocal);
200+ inQueueF.FreeTensor(fLocal);
201+ }
202+}
203+ 
204+template <typename TYPE_INPUT_GRADIENT>
205+__aicore__ inline void KernelThresholdGradV2D<TYPE_INPUT_GRADIENT>::Process()
206+{
207+ int32_t loopCount = this->tileNum;
208+ this->processDataNum = this->tileDataNum;
209+ for (int32_t i = 0; i < loopCount - 1; i++) {
210+ CopyIn(i);
211+ Compute(i);
212+ CopyOut(i);
213+ }
214+ this->processDataNum = this->tailDataNum;
215+ CopyIn(loopCount - 1);
216+ Compute(loopCount - 1);
217+ CopyOut(loopCount - 1);
218+}
219+ 
220+} // namespace NsThresholdGradV2D
221+#endif // THRESHOLD_GRAD_V2_D_H
Aexperimental/activation/threshold_grad_v2_d/op_kernel/threshold_grad_v2_d_tiling_data.h+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+ * \file threshold_grad_v2_d_tiling_data.h
13+ * \brief tiling data struct
14+ */
15+ 
16+#ifndef THRESHOLD_GRAD_V2_D_TILING_DATA_H_
17+#define THRESHOLD_GRAD_V2_D_TILING_DATA_H_
18+ 
19+struct ThresholdGradV2DTilingData {
20+ uint64_t smallCoreDataNum;
21+ uint64_t bigCoreDataNum;
22+ uint64_t finalBigTileNum;
23+ uint64_t finalSmallTileNum;
24+ uint64_t tileDataNum;
25+ uint64_t smallTailDataNum;
26+ uint64_t bigTailDataNum;
27+ uint64_t tailBlockNum;
28+ float threshold;
29+};
30+#endif
Aexperimental/activation/threshold_grad_v2_d/op_kernel/threshold_grad_v2_d_tiling_key.h+31-0
@@ -0,0 +1,31 @@
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 threshold_grad_v2_d_tiling_key.h
13+ * \brief threshold_grad_v2_d tiling key declare
14+ */
15+ 
16+#ifndef __THRESHOLD_GRAD_V2_D_TILING_KEY_H__
17+#define __THRESHOLD_GRAD_V2_D_TILING_KEY_H__
18+ 
19+#include "ascendc/host_api/tiling/template_argument.h"
20+ 
21+#define ELEMENTWISE_TPL_SCH_MODE_0 0
22+#define ELEMENTWISE_TPL_SCH_MODE_1 1
23+ 
24+ASCENDC_TPL_ARGS_DECL(
25+ ThresholdGradV2D,
26+ ASCENDC_TPL_UINT_DECL(schMode, 1, ASCENDC_TPL_UI_LIST, ELEMENTWISE_TPL_SCH_MODE_0, ELEMENTWISE_TPL_SCH_MODE_1));
27+ 
28+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(
29+ ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST, ELEMENTWISE_TPL_SCH_MODE_0, ELEMENTWISE_TPL_SCH_MODE_1)));
30+ 
31+#endif
Aexperimental/activation/threshold_grad_v2_d/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.activation.threshold_grad_v2_d.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/activation/threshold_grad_v2_d/tests/ut/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+# 每个目录下需要生成的可执行文件，具体参考：ops/built-in/test/CMakeLists.txt： 50~124
12+message(STATUS "=== Debug: start ops.activation.threshold_grad_v2_d.tests.ut.CMakeLists.txt ")
13+file(GLOB CURRENT_SOURCE_DIRS LIST_DIRECTORIES true ${CMAKE_CURRENT_SOURCE_DIR}/*)
14+message(STATUS "=== Debug: CURRENT_SOURCE_DIRS =${CURRENT_SOURCE_DIRS} ")
15+foreach(SUB_DIR ${CURRENT_SOURCE_DIRS})
16+ if(EXISTS "${SUB_DIR}/CMakeLists.txt")
17+ add_subdirectory(${SUB_DIR})
18+ endif()
19+endforeach()
Aexperimental/activation/threshold_grad_v2_d/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_threshold_grad_v2_d")
12+add_modules_llt_sources(HOSTNAME ${OPTEST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
Aexperimental/activation/threshold_grad_v2_d/tests/ut/op_api/test_aclnn_threshold_grad_v2_d.cpp+242-0
@@ -0,0 +1,242 @@
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 "gtest/gtest.h"
12+#include "level2/aclnn_threshold_backward.h"
13+#include "op_api_ut_common/tensor_desc.h"
14+#include "op_api_ut_common/op_api_ut.h"
15+#include "../../../op_host/op_api/aclnn_threshold_backward.h"
16+ 
17+class l2_threshold_backward_test : public testing::Test {
18+protected:
19+ static void SetUpTestCase() {
20+ std::cout << "l2_threshold_backward_test SetUp" << std::endl;
21+ }
22+ 
23+ static void TearDownTestCase() { std::cout << "l2_threshold_backward_test TearDown" << std::endl; }
24+};
25+ 
26+TEST_F(l2_threshold_backward_test, case_01_float)
27+{
28+ op::SetPlatformSocVersion(op::SocVersion::ASCEND910B);
29+ 
30+ auto gradOutputDesc = TensorDesc({2, 2}, ACL_FLOAT, ACL_FORMAT_ND).Value(vector<float>{1, 2, 3, 4});
31+ auto selfDesc = TensorDesc({2, 2}, ACL_FLOAT, ACL_FORMAT_ND).Value(vector<float>{1, 2, 3, 4});
32+ auto outDesc = TensorDesc({2, 2}, ACL_FLOAT, ACL_FORMAT_ND).Value(vector<float>{1, 2, 3, 4});
33+ auto scalar_desc = ScalarDesc(0.0f);
34+ auto outDesc = TensorDesc({2, 2}, ACL_FLOAT, ACL_FORMAT_ND);
35+ 
36+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, scalar_desc), OUTPUT(outDesc));
37+ 
38+ uint64_t workspaceSize = 0;
39+ aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
40+ EXPECT_EQ(getWorkspaceResult, ACLNN_ERR_PARAM_INVALID);
41+}
42+ 
43+// 计算的数据类型不在支持范围内
44+TEST_F(l2_threshold_backward_test, l2_test_unsupport_dtype) {
45+ auto gradOutputDesc = TensorDesc({2, 3}, ACL_INT64, ACL_FORMAT_ND);
46+ auto selfDesc = TensorDesc({2, 3}, ACL_INT64, ACL_FORMAT_ND);
47+ auto outDesc = TensorDesc({2, 3}, ACL_INT64, ACL_FORMAT_ND);
48+ auto scalar_desc = ScalarDesc(0.0f);
49+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, scalar_desc), OUTPUT(outDesc));
50+ 
51+ uint64_t workspaceSize = 0;
52+ aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
53+ EXPECT_EQ(getWorkspaceResult, ACLNN_ERR_PARAM_INVALID);
54+}
55+ 
56+// 输入的数据类型不匹配
57+TEST_F(l2_threshold_backward_test, l2_test_unmatch_dtype) {
58+ auto gradOutputDesc = TensorDesc({2, 3}, ACL_INT64, ACL_FORMAT_ND);
59+ auto selfDesc = TensorDesc({2, 3}, ACL_INT32, ACL_FORMAT_ND);
60+ auto outDesc = TensorDesc({2, 3}, ACL_INT32, ACL_FORMAT_ND);
61+ auto scalar_desc = ScalarDesc(0.0f);
62+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, scalar_desc), OUTPUT(outDesc));
63+ 
64+ uint64_t workspaceSize = 0;
65+ aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
66+ EXPECT_EQ(getWorkspaceResult, ACLNN_ERR_PARAM_INVALID);
67+}
68+ 
69+// 输入的shape不一致场景
70+TEST_F(l2_threshold_backward_test, l2_test_unmatch_shape) {
71+ auto gradOutputDesc = TensorDesc({2, 1}, ACL_INT32, ACL_FORMAT_ND);
72+ auto selfDesc = TensorDesc({2, 3}, ACL_INT32, ACL_FORMAT_ND);
73+ auto outDesc = TensorDesc({2, 3}, ACL_INT32, ACL_FORMAT_ND);
74+ auto scalar_desc = ScalarDesc(0.0f);
75+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, scalar_desc), OUTPUT(outDesc));
76+ 
77+ uint64_t workspaceSize = 0;
78+ aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
79+ EXPECT_EQ(getWorkspaceResult, ACLNN_SUCCESS);
80+}
81+ 
82+// 输出shape不可用场景
83+TEST_F(l2_threshold_backward_test, l2_test_invalid_out_shape) {
84+ auto gradOutputDesc = TensorDesc({2, 1}, ACL_INT32, ACL_FORMAT_ND);
85+ auto selfDesc = TensorDesc({2, 3}, ACL_INT32, ACL_FORMAT_ND);
86+ auto outDesc = TensorDesc({2, 4}, ACL_INT32, ACL_FORMAT_ND);
87+ auto scalar_desc = ScalarDesc(0.0f);
88+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, scalar_desc), OUTPUT(outDesc));
89+ 
90+ uint64_t workspaceSize = 0;
91+ aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
92+ EXPECT_EQ(getWorkspaceResult, ACLNN_ERR_PARAM_INVALID);
93+}
94+ 
95+// 空指针
96+TEST_F(l2_threshold_backward_test, l2_test_nullptr) {
97+ auto gradOutputDesc = TensorDesc({2, 0}, ACL_INT32, ACL_FORMAT_ND);
98+ auto selfDesc = TensorDesc({2, 0}, ACL_INT32, ACL_FORMAT_ND);
99+ auto outDesc = TensorDesc({2, 0}, ACL_INT32, ACL_FORMAT_ND);
100+ auto scalar_desc = ScalarDesc(0.0f);
101+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(nullptr, selfDesc, scalar_desc), OUTPUT(outDesc));
102+ 
103+ uint64_t workspaceSize = 0;
104+ aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
105+ EXPECT_EQ(getWorkspaceResult, ACLNN_ERR_PARAM_NULLPTR);
106+}
107+ 
108+// 空tensor
109+TEST_F(l2_threshold_backward_test, l2_test_empty_tensor) {
110+ auto gradOutputDesc = TensorDesc({2, 0}, ACL_INT32, ACL_FORMAT_ND);
111+ auto selfDesc = TensorDesc({2, 0}, ACL_INT32, ACL_FORMAT_ND);
112+ auto outDesc = TensorDesc({2, 0}, ACL_INT32, ACL_FORMAT_ND);
113+ auto scalar_desc = ScalarDesc(0.0f);
114+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, scalar_desc), OUTPUT(outDesc));
115+ 
116+ uint64_t workspaceSize = 0;
117+ aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
118+ EXPECT_EQ(getWorkspaceResult, ACLNN_SUCCESS);
119+}
120+ 
121+TEST_F(l2_threshold_backward_test, ascend910_9589_l2_test_relu_grad_int32_success) {
122+ auto gradOutputDesc = TensorDesc({10,}, ACL_INT32, ACL_FORMAT_ND);
123+ auto selfDesc = TensorDesc({10,}, ACL_INT32, ACL_FORMAT_ND);
124+ auto scalarSesc = ScalarDesc(0.0f);
125+ auto outDesc = TensorDesc(selfDesc);
126+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, scalarSesc), OUTPUT(outDesc));
127+ 
128+ uint64_t workspaceSize = 0;
129+ aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
130+ EXPECT_EQ(getWorkspaceResult, ACLNN_SUCCESS);
131+ 
132+ ut.TestPrecision();
133+}
134+ 
135+TEST_F(l2_threshold_backward_test, ascend910_9589_l2_test_relu_grad_int8_success) {
136+ auto gradOutputDesc = TensorDesc({10,}, ACL_INT8, ACL_FORMAT_ND);
137+ auto selfDesc = TensorDesc({10,}, ACL_INT8, ACL_FORMAT_ND);
138+ auto scalarSesc = ScalarDesc(0.0f);
139+ auto outDesc = TensorDesc(selfDesc);
140+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, scalarSesc), OUTPUT(outDesc));
141+ 
142+ uint64_t workspaceSize = 0;
143+ aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
144+ EXPECT_EQ(getWorkspaceResult, ACLNN_SUCCESS);
145+ 
146+ ut.TestPrecision();
147+}
148+ 
149+TEST_F(l2_threshold_backward_test, ascend910_9589_l2_test_relu_grad_fp16_success) {
150+ auto gradOutputDesc = TensorDesc({10,}, ACL_FLOAT16, ACL_FORMAT_ND);
151+ auto selfDesc = TensorDesc({10,}, ACL_FLOAT16, ACL_FORMAT_ND);
152+ auto scalarSesc = ScalarDesc(0.0f);
153+ auto outDesc = TensorDesc(selfDesc);
154+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, scalarSesc), OUTPUT(outDesc));
155+ 
156+ uint64_t workspaceSize = 0;
157+ aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
158+ EXPECT_EQ(getWorkspaceResult, ACLNN_SUCCESS);
159+ 
160+ ut.TestPrecision();
161+}
162+ 
163+TEST_F(l2_threshold_backward_test, ascend910_9589_l2_test_relu_grad_fp32_success) {
164+ auto gradOutputDesc = TensorDesc({10,}, ACL_FLOAT, ACL_FORMAT_ND);
165+ auto selfDesc = TensorDesc({10,}, ACL_FLOAT, ACL_FORMAT_ND);
166+ auto scalarSesc = ScalarDesc(0.0f);
167+ auto outDesc = TensorDesc(selfDesc);
168+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, scalarSesc), OUTPUT(outDesc));
169+ 
170+ uint64_t workspaceSize = 0;
171+ aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
172+ EXPECT_EQ(getWorkspaceResult, ACLNN_SUCCESS);
173+ 
174+ ut.TestPrecision();
175+}
176+ 
177+TEST_F(l2_threshold_backward_test, ascend910_9589_l2_test_relu_grad_int64_success) {
178+ auto gradOutputDesc = TensorDesc({10,}, ACL_INT64, ACL_FORMAT_ND);
179+ auto selfDesc = TensorDesc({10,}, ACL_INT64, ACL_FORMAT_ND);
180+ auto scalarSesc = ScalarDesc(0.0f);
181+ auto outDesc = TensorDesc(selfDesc);
182+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, scalarSesc), OUTPUT(outDesc));
183+ 
184+ uint64_t workspaceSize = 0;
185+ aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
186+ EXPECT_EQ(getWorkspaceResult, ACLNN_SUCCESS);
187+ 
188+ ut.TestPrecision();
189+}
190+ 
191+TEST_F(l2_threshold_backward_test, ascend910B2_l2_test_threshold_grad_v2_d_positive_threshold_success) {
192+ auto gradOutputDesc = TensorDesc({10,}, ACL_INT32, ACL_FORMAT_ND);
193+ auto selfDesc = TensorDesc({10,}, ACL_INT32, ACL_FORMAT_ND);
194+ auto scalarSesc = ScalarDesc(1.0f);
195+ auto outDesc = TensorDesc(selfDesc);
196+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, scalarSesc), OUTPUT(outDesc));
197+ 
198+ uint64_t workspaceSize = 0;
199+ aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
200+ EXPECT_EQ(getWorkspaceResult, ACLNN_SUCCESS);
201+ 
202+ ut.TestPrecision();
203+}
204+ 
205+TEST_F(l2_threshold_backward_test, ascend910B2_l2_test_threshold_grad_v2_d_negative_threshold_success) {
206+ auto gradOutputDesc = TensorDesc({10,}, ACL_INT32, ACL_FORMAT_ND);
207+ auto selfDesc = TensorDesc({10,}, ACL_INT32, ACL_FORMAT_ND);
208+ auto scalarSesc = ScalarDesc(-1.0f);
209+ auto outDesc = TensorDesc(selfDesc);
210+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, scalarSesc), OUTPUT(outDesc));
211+ 
212+ uint64_t workspaceSize = 0;
213+ aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
214+ EXPECT_EQ(getWorkspaceResult, ACLNN_SUCCESS);
215+ 
216+ ut.TestPrecision();
217+}
218+ 
219+TEST_F(l2_threshold_backward_test, ascend910B2_l2_test_relu_grad_bfp16_success) {
220+ auto gradOutputDesc = TensorDesc({10,}, ACL_BF16, ACL_FORMAT_ND);
221+ auto selfDesc = TensorDesc({10,}, ACL_BF16, ACL_FORMAT_ND);
222+ auto scalarSesc = ScalarDesc(0.0f);
223+ auto outDesc = TensorDesc(selfDesc);
224+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, scalarSesc), OUTPUT(outDesc));
225+ 
226+ uint64_t workspaceSize = 0;
227+ aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
228+ EXPECT_EQ(getWorkspaceResult, ACLNN_SUCCESS);
229+}
230+ 
231+//校验维度超过8维的场景
232+TEST_F(l2_threshold_backward_test, ascend910B2_l2_test_relu_grad_check_max_dim) {
233+ auto gradOutputDesc = TensorDesc({2,2,2,2,2,2,2,2,2,2}, ACL_FLOAT, ACL_FORMAT_ND);
234+ auto selfDesc = TensorDesc({2,2,2,2,2,2,2,2,2,2}, ACL_FLOAT, ACL_FORMAT_ND);
235+ auto scalarSesc = ScalarDesc(0.0f);
236+ auto outDesc = TensorDesc(selfDesc);
237+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, scalarSesc), OUTPUT(outDesc));
238+ 
239+ uint64_t workspaceSize = 0;
240+ aclnnStatus getWorkspaceResult = ut.TestGetWorkspaceSize(&workspaceSize);
241+ EXPECT_EQ(getWorkspaceResult, ACLNN_ERR_PARAM_INVALID);
242+}
Aexperimental/activation/threshold_grad_v2_d/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/activation/threshold_grad_v2_d/tests/ut/op_host/test_threshold_grad_v2_d_tiling.cpp+105-0
@@ -0,0 +1,105 @@
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 ThresholdGradV2DTiling : public testing::Test {
29+protected:
30+ static void SetUpTestCase()
31+ {
32+ std::cout << "ThresholdGradV2DTiling SetUp" << std::endl;
33+ }
34+ 
35+ static void TearDownTestCase()
36+ {
37+ std::cout << "ThresholdGradV2DTiling TearDown" << std::endl;
38+ }
39+};
40+ 
41+TEST_F(ThresholdGradV2DTiling, threshold_grad_v2_float32_success) {
42+ // input
43+ gert::StorageShape x1_shape = {{1, 2, 8, 16}, {1, 2, 8, 16}};
44+ gert::StorageShape x2_shape = {{1, 2, 8, 16}, {1, 2, 8, 16}};
45+ // output
46+ gert::StorageShape y_shape = {{1, 2, 8, 16}, {1, 2, 8, 16}};
47+ string compile_info_string = R"({
48+ "hardware_info": {"BT_SIZE": 0, "load3d_constraints": "1",
49+ "Intrinsic_fix_pipe_l0c2out": false, "Intrinsic_data_move_l12ub": true, "Intrinsic_data_move_l0c2ub": true, "Intrinsic_data_move_out2l1_nd2nz": false,
50+ "UB_SIZE": 196608, "L2_SIZE": 33554432, "L1_SIZE": 524288,
51+ "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072,
52+ "CORE_NUM": 48}
53+ })";
54+ map<string, string> soc_infos;
55+ map<string, string> aicore_spec;
56+ map<string, string> intrinsics;
57+ GetPlatFormInfos(compile_info_string.c_str(), soc_infos, aicore_spec, intrinsics);
58+ 
59+ // platform info
60+ fe::PlatFormInfos platform_info;
61+ platform_info.Init();
62+ 
63+ // compile info
64+ struct ThresholdGradV2DTilingCompileInfo {};
65+ ThresholdGradV2DTilingCompileInfo compile_info;
66+ 
67+ std::string op_type("ThresholdGradV2D");
68+ auto tiling_func = gert::OpImplRegistry::GetInstance().GetOpImpl(op_type.c_str())->tiling;
69+ 
70+ // tilingParseFunc simulate
71+ auto kernel_holder =
72+ gert::KernelRunContextFaker()
73+ .KernelIONum(2, 1)
74+ .Inputs({const_cast<char*>(compile_info_string.c_str()), reinterpret_cast<void*>(&platform_info)})
75+ .Outputs({&compile_info})
76+ .Build();
77+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", soc_infos);
78+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicore_spec);
79+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
80+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap",
81+ intrinsics);
82+ 
83+ // tilingFunc simulate
84+ auto param = gert::TilingData::CreateCap(4096);
85+ auto workspace_size_holer = gert::ContinuousVector::Create<size_t>(4096);
86+ auto ws_size = reinterpret_cast<gert::ContinuousVector*>(workspace_size_holer.get());
87+ ASSERT_NE(param, nullptr);
88+ auto holder = gert::TilingContextFaker()
89+ .SetOpType("ThresholdGradV2D")
90+ .NodeIoNum(2, 1)
91+ .IrInstanceNum({1, 1})
92+ .InputShapes({&x1_shape, &x2_shape})
93+ .OutputShapes({&y_shape})
94+ .CompileInfo(&compile_info)
95+ .PlatformInfo(reinterpret_cast<char*>(&platform_info))
96+ .NodeInputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
97+ .NodeInputTd(1, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
98+ .NodeOutputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
99+ .TilingData(param.get())
100+ .Workspace(ws_size)
101+ .Build();
102+ gert::TilingContext* tiling_context = holder.GetContext<gert::TilingContext>();
103+ ASSERT_NE(tiling_context, nullptr);
104+ EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_SUCCESS);
105+}
Aexperimental/activation/threshold_grad_v2_d/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(threshold_grad_v2_d_tiling_files
19+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/threshold_grad_v2_d_tiling.cpp
20+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/threshold_grad_v2_d_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(threshold_grad_v2_d "ascend910b" "${threshold_grad_v2_d_tiling_files}")
29+ AddOpTestCase(threshold_grad_v2_d "ascend910B1" "-DDTYPE_INPUT_GRADIENT=float" "${threshold_grad_v2_d_tiling_files}")
30+endif()
Aexperimental/activation/threshold_grad_v2_d/tests/ut/op_kernel/test_threshold_grad_v2_d.cpp+84-0
@@ -0,0 +1,84 @@
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/threshold_grad_v2_d.cpp"
23+#include "../../../op_kernel/threshold_grad_v2_d_tiling_data.h"
24+#include <cstdint>
25+ 
26+using namespace std;
27+ 
28+class threshold_grad_v2_d_test : public testing::Test {
29+protected:
30+ static void SetUpTestCase()
31+ {
32+ cout << "threshold_grad_v2_d_test SetUp\n" << endl;
33+ }
34+ static void TearDownTestCase()
35+ {
36+ cout << "threshold_grad_v2_d_test TearDown\n" << endl;
37+ }
38+};
39+ 
40+TEST_F(threshold_grad_v2_d_test, test_case_0)
41+{
42+ size_t input_gradientByteSize = 32 * 4 * 4 * 4 * sizeof(float);
43+ size_t input_featureByteSize = 32 * 4 * 4 * 4 * sizeof(float);
44+ size_t output_backpropsByteSize = 32 * 4 * 4 * 4 * sizeof(float);
45+ size_t tiling_data_size = sizeof(ThresholdGradV2DTilingData);
46+ uint32_t blockDim = 1;
47+ 
48+ uint8_t* input_gradient = (uint8_t*)AscendC::GmAlloc(input_gradientByteSize);
49+ uint8_t* input_feature = (uint8_t*)AscendC::GmAlloc(input_featureByteSize);
50+ uint8_t* output_backprops = (uint8_t*)AscendC::GmAlloc(output_backpropsByteSize);
51+ 
52+ uint8_t* workspace = (uint8_t*)AscendC::GmAlloc(1024 * 1024 * 16);
53+ uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(tiling_data_size);
54+ 
55+ char* path_ = get_current_dir_name();
56+ string path(path_);
57+ 
58+ ThresholdGradV2DTilingData* tilingDatafromBin = reinterpret_cast<ThresholdGradV2DTilingData*>(tiling);
59+ 
60+ tilingDatafromBin->smallCoreDataNum = 2048;
61+ tilingDatafromBin->bigCoreDataNum = 2112;
62+ tilingDatafromBin->tileDataNum = 4032;
63+ tilingDatafromBin->smallTailDataNum = 2048;
64+ tilingDatafromBin->bigTailDataNum = 2112;
65+ tilingDatafromBin->finalSmallTileNum = 1;
66+ tilingDatafromBin->finalBigTileNum = 1;
67+ tilingDatafromBin->tailBlockNum = 0;
68+ tilingDatafromBin->threshold = 1.0;
69+ 
70+ auto ThresholdGradV2DKernel = [](GM_ADDR input_gradient, GM_ADDR input_feature, GM_ADDR output_backprops, GM_ADDR workspace, GM_ADDR tiling) {
71+ ::threshold_grad_v2_d<0>(input_gradient, input_feature, output_backprops, workspace, tiling);
72+ };
73+ 
74+ ICPU_SET_TILING_KEY(0);
75+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
76+ ICPU_RUN_KF(ThresholdGradV2DKernel, blockDim, input_gradient, input_feature, output_backprops, workspace, (uint8_t *)(tilingDatafromBin));
77+ 
78+ AscendC::GmFree(input_gradient);
79+ AscendC::GmFree(input_feature);
80+ AscendC::GmFree(output_backprops);
81+ AscendC::GmFree(workspace);
82+ AscendC::GmFree(tiling);
83+ free(path_);
84+}