Star900
1814
代码
代码
Issues233
Pull Requests310
Actions
项目讨论
Wiki
分析
Star900
1814
已合并
新增ThresholdBackward实验算子（relu_grad_v2） #4479
cwzhang创建于 4月30日
新增ThresholdBackward实验算子（relu_grad_v2） #4479
已合并
cwzhang创建于 4月30日
从cwzhang/ops-nnrelu_grad合入到cann/ops-nnmaster
共 23 个文件变更+2570-0
Aexperimental/activation/relu_grad_v2/CMakeLists.txt+19-0
@@ -0,0 +1,19 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------
10+ 
11+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+if(NOT ENABLE_TEST)
13+ list(REMOVE_ITEM CURRENT_DIRS tests)
14+endif()
15+foreach(SUB_DIR ${CURRENT_DIRS})
16+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
17+ add_subdirectory(${SUB_DIR})
18+ endif()
19+endforeach()
Aexperimental/activation/relu_grad_v2/README.md+171-0
@@ -0,0 +1,171 @@
1+# ThresholdBackward
2+ 
3+## 产品支持情况
4+ 
5+| 产品 | 是否支持 |
6+| :--- | :---: |
7+| <term>Ascend 950PR/Ascend 950DT</term> | √ |
8+| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |
9+| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |
10+| <term>Atlas 200I/500 A2 推理产品</term> | × |
11+| <term>Atlas 推理系列产品</term> | √ |
12+| <term>Atlas 训练系列产品</term> | √ |
13+ 
14+## 功能说明
15+ 
16+- 算子功能：对 `gradOutput` 和 `self` 执行 ReLU 反向梯度计算。
17+ 
18+计算公式：
19+ 
20+$$
21+\operatorname{threshold\_backward}(gradOutput, self, threshold) =
22+\begin{cases}
23+gradOutput, & self > 0 \\
24+gradOutput, & self = \mathrm{NaN}\ \text{且 dtype 为浮点类型} \\
25+0, & self \le 0
26+\end{cases}
27+$$
28+ 
29+- 目录 `experimental/activation/relu_grad_v2` 对外导出 `aclnnThresholdBackward` 两段式 ACLNN 接口。
30+- `op_host/op_api/aclnn_threshold_backward.cpp` 是对外 ACLNN 接口入口。
31+- `op_host/op_api/relu_grad_v2.h` 和 `op_host/op_api/relu_grad_v2.cpp` 提供内部 `l0op::ReluGradV2` 封装，当前由 ACLNN 接口直接调用。
32+- 当前实现仅接受 `threshold == 0`，与 ReLU backward 语义保持一致。
33+ 
34+## 调用方式
35+ 
36+| 调用方式 | 是否支持 |
37+| :------- | :------: |
38+| ACLNN 调用 | 是 |
39+ 
40+## ACLNN 接口
41+ 
42+### 函数原型
43+ 
44+当前 experimental ThresholdBackward 提供两段式 ACLNN 接口：
45+ 
46+```cpp
47+aclnnStatus aclnnThresholdBackwardGetWorkspaceSize(
48+ const aclTensor *gradOutput,
49+ const aclTensor *self,
50+ const aclScalar *threshold,
51+ aclTensor *out,
52+ uint64_t *workspaceSize,
53+ aclOpExecutor **executor);
54+ 
55+aclnnStatus aclnnThresholdBackward(
56+ void *workspace,
57+ uint64_t workspaceSize,
58+ aclOpExecutor *executor,
59+ const aclrtStream stream);
60+```
61+ 
62+详细参数和返回值说明见 [docs/aclnnThresholdBackward.md](docs/aclnnThresholdBackward.md)。
63+ 
64+## 参数说明
65+ 
66+<table style="undefined;table-layout: fixed; width: 1393px"><colgroup>
67+<col style="width: 171px">
68+<col style="width: 115px">
69+<col style="width: 260px">
70+<col style="width: 220px">
71+<col style="width: 200px">
72+<col style="width: 104px">
73+</colgroup>
74+<thead>
75+ <tr>
76+ <th>参数名</th>
77+ <th>输入/输出</th>
78+ <th>描述</th>
79+ <th>数据类型</th>
80+ <th>数据格式</th>
81+ </tr>
82+</thead>
83+<tbody>
84+ <tr>
85+ <td>gradOutput</td>
86+ <td>输入</td>
87+ <td>上游梯度张量。</td>
88+ <td>FLOAT、FLOAT16、BFLOAT16（仅 Ascend910B 及后续同代 SoC 支持）、INT8、UINT8、INT32、INT64</td>
89+ <td>ND</td>
90+ </tr>
91+ <tr>
92+ <td>self</td>
93+ <td>输入</td>
94+ <td>前向输入张量，用于生成 ReLU 掩码。</td>
95+ <td>FLOAT、FLOAT16、BFLOAT16（仅 Ascend910B 及后续同代 SoC 支持）、INT8、UINT8、INT32、INT64</td>
96+ <td>ND</td>
97+ </tr>
98+ <tr>
99+ <td>threshold</td>
100+ <td>输入</td>
101+ <td>阈值标量。当前实现仅接受值为 `0` 的 `INT32` 标量。</td>
102+ <td>INT32</td>
103+ <td>Scalar</td>
104+ </tr>
105+ <tr>
106+ <td>out</td>
107+ <td>输出</td>
108+ <td>计算得到的输出梯度张量。</td>
109+ <td>FLOAT、FLOAT16、BFLOAT16（仅 Ascend910B 及后续同代 SoC 支持）、INT8、UINT8、INT32、INT64</td>
110+ <td>ND</td>
111+ </tr>
112+</tbody>
113+</table>
114+ 
115+## 约束说明
116+ 
117+- `gradOutput`、`self` 和 `out` 的 dtype 必须完全一致。
118+- `gradOutput`、`self` 和 `out` 的 shape 必须完全一致。
119+- `threshold` 必须是值为 `0` 的标量。
120+- 输入仅支持 `FLOAT`、`FLOAT16`、`BFLOAT16`、`INT8`、`UINT8`、`INT32`、`INT64`。
121+- 支持 0 到 8 维 Tensor。
122+- 支持空 Tensor。
123+- 支持非连续 Tensor，接口内部会在需要时做 `Contiguous` 和 `ViewCopy`。
124+- `FLOAT` 路径遵循 PyTorch `threshold_backward` 的标量语义，当 `self` 为 `NaN` 时保留 `gradOutput`。
125+- `FLOAT16` 和 `BFLOAT16` 路径在 kernel 中升精度到 `float32` 计算后回写。
126+- `INT8`、`UINT8`、`INT32` 和 `INT64` 路径遵循 `self > 0 ? gradOutput : 0`。
127+ 
128+## 目录说明
129+ 
130+| 路径 | 说明 |
131+| :--- | :--- |
132+| [examples/test_aclnn_relu_grad_v2.cpp](examples/test_aclnn_relu_grad_v2.cpp) | `aclnnThresholdBackward` 两段式调用示例。 |
133+| [examples/run.sh](examples/run.sh) | 编译并运行 example 的脚本。 |
134+| [docs/aclnnThresholdBackward.md](docs/aclnnThresholdBackward.md) | `aclnnThresholdBackward` 接口文档。 |
135+| [tests/ut/op_api/test_aclnn_threshold_backward.cpp](tests/ut/op_api/test_aclnn_threshold_backward.cpp) | `op_api` 单元测试。 |
136+| [tests/st/aclnnThresholdBackward/all_aclnnThresholdBackward.json](tests/st/aclnnThresholdBackward/all_aclnnThresholdBackward.json) | 适用于 ATK 的小规模标准化测试集。 |
137+| [tests/st/aclnnThresholdBackward/executor_aclnnThresholdBackward.py](tests/st/aclnnThresholdBackward/executor_aclnnThresholdBackward.py) | ATK CPU benchmark 执行器。 |
138+ 
139+## Example 运行
140+ 
141+先确保 custom run 包已经安装，并加载 CANN 环境：
142+ 
143+```bash
144+source /usr/local/Ascend/cann/set_env.sh
145+export LD_LIBRARY_PATH=/usr/local/Ascend/cann/opp/vendors/customize_nn/op_api/lib:${LD_LIBRARY_PATH}
146+cd <ops-nn-repo>/experimental/activation/relu_grad_v2/examples
147+bash run.sh
148+```
149+ 
150+## Tests 运行
151+ 
152+### 1. op_api 单元测试
153+ 
154+```bash
155+source /usr/local/Ascend/cann/set_env.sh
156+cd <ops-nn-repo>
157+bash build.sh --experimental --ops=relu_grad_v2 -u --opapi -j8 -O2
158+```
159+ 
160+### 2. ATK 小规模标准化测试
161+ 
162+```bash
163+export ATK_BIND_CPU_TYPE=2
164+source /usr/local/Ascend/cann/set_env.sh
165+source /root/src/kernel/ascend-kernel/.venv/bin/activate
166+cd /root/src/testcase
167+atk node --backend npu --devices 2 \
168+ node --backend cpu task --task accuracy \
169+ -c ./experimental/activation/relu_grad_v2/tests/st/aclnnThresholdBackward/all_aclnnThresholdBackward.json \
170+ -p ./experimental/activation/relu_grad_v2/tests/st/aclnnThresholdBackward/executor_aclnnThresholdBackward.py
171+```
Aexperimental/activation/relu_grad_v2/docs/aclnnThresholdBackward.md+269-0
@@ -0,0 +1,269 @@
1+# aclnnThresholdBackward
2+ 
3+## 产品支持情况
4+ 
5+|产品|是否支持|
6+|:---|:---:|
7+|<term>Ascend 950PR/Ascend 950DT</term>|√|
8+|<term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>|√|
9+|<term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term>|√|
10+|<term>Atlas 200I/500 A2 推理产品</term>|×|
11+|<term>Atlas 推理系列产品</term>|√|
12+|<term>Atlas 训练系列产品</term>|√|
13+ 
14+## 功能说明
15+ 
16+- `aclnnThresholdBackward`：对输入 `gradOutput`、`self` 和 `threshold` 执行 ReLU 反向梯度计算，并将结果写入独立输出 Tensor。
17+- `experimental/activation/relu_grad_v2` 目录对外导出的 ACLNN 接口名与当前实现保持一致。
18+- 当前实现仅接受 `threshold == 0`，用于对齐 ReLU backward 语义。
19+ 
20+计算公式：
21+ 
22+$$
23+\operatorname{threshold\_backward}(gradOutput, self, threshold) =
24+\begin{cases}
25+gradOutput, & self > 0 \\
26+gradOutput, & self = \mathrm{NaN}\ \text{且 dtype 为浮点类型} \\
27+0, & self \le 0
28+\end{cases}
29+$$
30+ 
31+## 函数原型
32+ 
33+每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md)，必须先调用 `aclnnThresholdBackwardGetWorkspaceSize` 获取执行器和 workspace 大小，再调用第二段接口执行计算。
34+ 
35+```cpp
36+aclnnStatus aclnnThresholdBackwardGetWorkspaceSize(
37+ const aclTensor *gradOutput,
38+ const aclTensor *self,
39+ const aclScalar *threshold,
40+ aclTensor *out,
41+ uint64_t *workspaceSize,
42+ aclOpExecutor **executor);
43+```
44+ 
45+```cpp
46+aclnnStatus aclnnThresholdBackward(
47+ void *workspace,
48+ uint64_t workspaceSize,
49+ aclOpExecutor *executor,
50+ const aclrtStream stream);
51+```
52+ 
53+## aclnnThresholdBackwardGetWorkspaceSize
54+ 
55+- **参数说明：**
56+ 
57+ <table style="undefined;table-layout: fixed; width: 1497px"><colgroup>
58+ <col style="width: 271px">
59+ <col style="width: 115px">
60+ <col style="width: 247px">
61+ <col style="width: 300px">
62+ <col style="width: 177px">
63+ <col style="width: 104px">
64+ <col style="width: 138px">
65+ <col style="width: 145px">
66+ </colgroup>
67+ <thead>
68+ <tr>
69+ <th>参数名</th>
70+ <th>输入/输出</th>
71+ <th>描述</th>
72+ <th>使用说明</th>
73+ <th>数据类型</th>
74+ <th>数据格式</th>
75+ <th>维度(shape)</th>
76+ <th>非连续Tensor</th>
77+ </tr>
78+ </thead>
79+ <tbody>
80+ <tr>
81+ <td>gradOutput（aclTensor*）</td>
82+ <td>输入</td>
83+ <td>上游梯度张量。</td>
84+ <td><ul><li>支持空Tensor。</li><li>shape 必须与 self、out 完全一致。</li><li>数据类型必须与 self、out 完全一致。</li></ul></td>
85+ <td>BFLOAT16（仅 Ascend910B 及后续同代 SoC 支持）、FLOAT16、FLOAT32、INT8、UINT8、INT32、INT64</td>
86+ <td>ND</td>
87+ <td>0-8</td>
88+ <td>√</td>
89+ </tr>
90+ <tr>
91+ <td>self（aclTensor*）</td>
92+ <td>输入</td>
93+ <td>前向输入张量，用于生成 ReLU 掩码。</td>
94+ <td><ul><li>支持空Tensor。</li><li>shape 必须与 gradOutput、out 完全一致。</li><li>数据类型必须与 gradOutput、out 完全一致。</li></ul></td>
95+ <td>BFLOAT16（仅 Ascend910B 及后续同代 SoC 支持）、FLOAT16、FLOAT32、INT8、UINT8、INT32、INT64</td>
96+ <td>ND</td>
97+ <td>0-8</td>
98+ <td>√</td>
99+ </tr>
100+ <tr>
101+ <td>threshold（aclScalar*）</td>
102+ <td>输入</td>
103+ <td>阈值标量。</td>
104+ <td><ul><li>当前实现仅支持数值为 0。</li><li>建议传入 `INT32` 标量。</li></ul></td>
105+ <td>INT32</td>
106+ <td>Scalar</td>
107+ <td>-</td>
108+ <td>-</td>
109+ </tr>
110+ <tr>
111+ <td>out（aclTensor*）</td>
112+ <td>输出</td>
113+ <td>计算的出参。</td>
114+ <td><ul><li>支持空Tensor。</li><li>shape 必须与 gradOutput、self 完全一致。</li><li>数据类型必须与 gradOutput、self 完全一致。</li></ul></td>
115+ <td>BFLOAT16（仅 Ascend910B 及后续同代 SoC 支持）、FLOAT16、FLOAT32、INT8、UINT8、INT32、INT64</td>
116+ <td>ND</td>
117+ <td>0-8</td>
118+ <td>√</td>
119+ </tr>
120+ <tr>
121+ <td>workspaceSize（uint64_t*）</td>
122+ <td>输出</td>
123+ <td>返回需要在 Device 侧申请的 workspace 大小。</td>
124+ <td>-</td>
125+ <td>-</td>
126+ <td>-</td>
127+ <td>-</td>
128+ <td>-</td>
129+ </tr>
130+ <tr>
131+ <td>executor（aclOpExecutor**）</td>
132+ <td>输出</td>
133+ <td>返回 op 执行器，包含算子计算流程。</td>
134+ <td>-</td>
135+ <td>-</td>
136+ <td>-</td>
137+ <td>-</td>
138+ <td>-</td>
139+ </tr>
140+ </tbody>
141+ </table>
142+ 
143+- **返回值：**
144+ 
145+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
146+ 
147+ 第一段接口会完成入参校验，出现以下场景时报错：
148+ 
149+ <table style="undefined;table-layout: fixed;width: 979px"><colgroup>
150+ <col style="width: 272px">
151+ <col style="width: 103px">
152+ <col style="width: 604px">
153+ </colgroup>
154+ <thead>
155+ <tr>
156+ <th>返回码</th>
157+ <th>错误码</th>
158+ <th>描述</th>
159+ </tr>
160+ </thead>
161+ <tbody>
162+ <tr>
163+ <td>ACLNN_ERR_PARAM_NULLPTR</td>
164+ <td>161001</td>
165+ <td>传入的 gradOutput、self、threshold 或 out 是空指针。</td>
166+ </tr>
167+ <tr>
168+ <td rowspan="5">ACLNN_ERR_PARAM_INVALID</td>
169+ <td rowspan="5">161002</td>
170+ <td>gradOutput、self 或 out 的数据类型不在支持范围内。</td>
171+ </tr>
172+ <tr>
173+ <td>gradOutput、self 和 out 的数据类型不一致。</td>
174+ </tr>
175+ <tr>
176+ <td>gradOutput、self 和 out 的 shape 不一致。</td>
177+ </tr>
178+ <tr>
179+ <td>gradOutput、self 或 out 的维度大于 8。</td>
180+ </tr>
181+ <tr>
182+ <td>threshold 不等于 0。</td>
183+ </tr>
184+ </tbody></table>
185+ 
186+## aclnnThresholdBackward
187+ 
188+- **参数说明：**
189+ 
190+ <table style="undefined;table-layout: fixed; width: 953px"><colgroup>
191+ <col style="width: 173px">
192+ <col style="width: 112px">
193+ <col style="width: 668px">
194+ </colgroup>
195+ <thead>
196+ <tr>
197+ <th>参数名</th>
198+ <th>输入/输出</th>
199+ <th>描述</th>
200+ </tr>
201+ </thead>
202+ <tbody>
203+ <tr>
204+ <td>workspace</td>
205+ <td>输入</td>
206+ <td>在 Device 侧申请的 workspace 内存地址。</td>
207+ </tr>
208+ <tr>
209+ <td>workspaceSize</td>
210+ <td>输入</td>
211+ <td>在 Device 侧申请的 workspace 大小，由第一段接口 aclnnThresholdBackwardGetWorkspaceSize 获取。</td>
212+ </tr>
213+ <tr>
214+ <td>executor</td>
215+ <td>输入</td>
216+ <td>op 执行器，包含算子计算流程。</td>
217+ </tr>
218+ <tr>
219+ <td>stream</td>
220+ <td>输入</td>
221+ <td>指定执行任务的 Stream。</td>
222+ </tr>
223+ </tbody>
224+ </table>
225+ 
226+- **返回值：**
227+ 
228+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
229+ 
230+## 实现说明
231+ 
232+- 当前实现调用 `l0op::ReluGradV2` AscendC kernel 完成主体计算。
233+- 接口内部会先对输入执行 `Contiguous`，再通过 `ViewCopy` 将结果写回 `out`，因此支持非连续 Tensor。
234+- 浮点类型遵循 PyTorch `threshold_backward` 在 `threshold == 0` 时的 `NaN` 处理语义；整型路径遵循 `self > 0 ? gradOutput : 0`。
235+ 
236+## 调用示例
237+ 
238+```cpp
239+#include "aclnnop/aclnn_threshold_backward.h"
240+ 
241+aclnnStatus RunThresholdBackward(const aclTensor *gradOutput,
242+ const aclTensor *self,
243+ const aclScalar *threshold,
244+ aclTensor *out,
245+ aclrtStream stream)
246+{
247+ uint64_t workspaceSize = 0;
248+ aclOpExecutor *executor = nullptr;
249+ auto ret = aclnnThresholdBackwardGetWorkspaceSize(
250+ gradOutput, self, threshold, out, &workspaceSize, &executor);
251+ if (ret != ACL_SUCCESS) {
252+ return ret;
253+ }
254+ 
255+ void *workspace = nullptr;
256+ if (workspaceSize > 0) {
257+ ret = aclrtMalloc(&workspace, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
258+ if (ret != ACL_SUCCESS) {
259+ return ret;
260+ }
261+ }
262+ 
263+ ret = aclnnThresholdBackward(workspace, workspaceSize, executor, stream);
264+ if (workspace != nullptr) {
265+ aclrtFree(workspace);
266+ }
267+ return ret;
268+}
269+```
Aexperimental/activation/relu_grad_v2/examples/run.sh+34-0
@@ -0,0 +1,34 @@
1+#!/usr/bin/env bash
2+# ----------------------------------------------------------------------------
3+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
4+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
5+# CANN Open Software License Agreement Version 2.0 (the "License").
6+# Please refer to the License for details. You may not use this file except in compliance with the License.
7+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
8+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9+# See LICENSE in the root of the software repository for the full text of the License.
10+# ----------------------------------------------------------------------------
11+ 
12+set -euo pipefail
13+ 
14+SCRIPT_DIR=$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)
15+CANN_ROOT=${ASCEND_HOME_PATH:-/usr/local/Ascend/cann}
16+ 
17+source "${CANN_ROOT}/set_env.sh"
18+export LD_LIBRARY_PATH="${CANN_ROOT}/opp/vendors/customize_nn/op_api/lib:${LD_LIBRARY_PATH:-}"
19+ 
20+BUILD_DIR="${SCRIPT_DIR}/build"
21+mkdir -p "${BUILD_DIR}"
22+ 
23+g++ -std=c++17 -O2 \
24+ "${SCRIPT_DIR}/test_aclnn_relu_grad_v2.cpp" \
25+ -I"${CANN_ROOT}/aarch64-linux/include" \
26+ -I"${CANN_ROOT}/aarch64-linux/include/aclnnop" \
27+ -I"${CANN_ROOT}/opp/vendors/customize_nn/op_api/include" \
28+ -L"${CANN_ROOT}/lib64" \
29+ -L"${CANN_ROOT}/opp/vendors/customize_nn/op_api/lib" \
30+ -Wl,-rpath,"${CANN_ROOT}/lib64:${CANN_ROOT}/opp/vendors/customize_nn/op_api/lib" \
31+ -lcust_opapi -lnnopbase -lascendcl \
32+ -o "${BUILD_DIR}/test_aclnn_relu_grad_v2"
33+ 
34+"${BUILD_DIR}/test_aclnn_relu_grad_v2"
Aexperimental/activation/relu_grad_v2/examples/test_aclnn_relu_grad_v2.cpp+393-0
@@ -0,0 +1,393 @@
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 test_aclnn_relu_grad_v2.cpp
13+ * @brief ACLNN invocation example for experimental ThresholdBackward(ReLU grad) operator
14+ */
15+ 
16+#include <cstdint>
17+#include <cstdio>
18+#include <cstdlib>
19+#include <cstring>
20+#include <fstream>
21+#include <string>
22+#include <vector>
23+ 
24+#include "acl/acl.h"
25+#include "aclnnop/aclnn_threshold_backward.h"
26+ 
27+#define CHECK_RET(cond, expr) \
28+ do { \
29+ if (!(cond)) { \
30+ expr; \
31+ } \
32+ } while (0)
33+ 
34+namespace {
35+struct ReluGradConfig {
36+ aclDataType acl_dtype;
37+ size_t element_size;
38+ std::string name;
39+};
40+ 
41+int ReportAclError(const char *stage, int ret)
42+{
43+ std::fprintf(stderr, "%s failed, ret=%d, msg=%s\n", stage, ret, aclGetRecentErrMsg());
44+ return ret;
45+}
46+ 
47+int64_t GetShapeSize(const std::vector<int64_t> &shape)
48+{
49+ int64_t shape_size = 1;
50+ for (int64_t dim : shape) {
51+ shape_size *= dim;
52+ }
53+ return shape_size;
54+}
55+ 
56+std::vector<int64_t> MakeStrides(const std::vector<int64_t> &shape)
57+{
58+ if (shape.empty()) {
59+ return {};
60+ }
61+ std::vector<int64_t> strides(shape.size(), 1);
62+ for (int64_t i = static_cast<int64_t>(shape.size()) - 2; i >= 0; --i) {
63+ strides[static_cast<size_t>(i)] = shape[static_cast<size_t>(i + 1)] * strides[static_cast<size_t>(i + 1)];
64+ }
65+ return strides;
66+}
67+ 
68+bool ParseDtype(const std::string &dtype_name, ReluGradConfig *config)
69+{
70+ if (dtype_name == "fp16" || dtype_name == "float16") {
71+ *config = {ACL_FLOAT16, sizeof(uint16_t), "fp16"};
72+ return true;
73+ }
74+ if (dtype_name == "fp32" || dtype_name == "float32") {
75+ *config = {ACL_FLOAT, sizeof(float), "fp32"};
76+ return true;
77+ }
78+ if (dtype_name == "bf16" || dtype_name == "bfloat16") {
79+ *config = {ACL_BF16, sizeof(uint16_t), "bf16"};
80+ return true;
81+ }
82+ if (dtype_name == "int8") {
83+ *config = {ACL_INT8, sizeof(int8_t), "int8"};
84+ return true;
85+ }
86+ if (dtype_name == "uint8") {
87+ *config = {ACL_UINT8, sizeof(uint8_t), "uint8"};
88+ return true;
89+ }
90+ if (dtype_name == "int32") {
91+ *config = {ACL_INT32, sizeof(int32_t), "int32"};
92+ return true;
93+ }
94+ if (dtype_name == "int64") {
95+ *config = {ACL_INT64, sizeof(int64_t), "int64"};
96+ return true;
97+ }
98+ return false;
99+}
100+ 
101+bool ParseShape(const std::string &shape_text, std::vector<int64_t> *shape)
102+{
103+ shape->clear();
104+ if (shape_text.empty() || shape_text == "scalar") {
105+ return true;
106+ }
107+ 
108+ size_t start = 0;
109+ while (start < shape_text.size()) {
110+ size_t end = shape_text.find(',', start);
111+ std::string token = shape_text.substr(start, end == std::string::npos ? std::string::npos : end - start);
112+ if (token.empty()) {
113+ return false;
114+ }
115+ shape->push_back(std::stoll(token));
116+ if (end == std::string::npos) {
117+ break;
118+ }
119+ start = end + 1;
120+ }
121+ return true;
122+}
123+ 
124+bool ReadFile(const std::string &path, std::vector<char> *buffer)
125+{
126+ std::ifstream stream(path, std::ios::binary);
127+ if (!stream.is_open()) {
128+ return false;
129+ }
130+ stream.seekg(0, std::ios::end);
131+ std::streamsize size = stream.tellg();
132+ stream.seekg(0, std::ios::beg);
133+ if (size < 0) {
134+ return false;
135+ }
136+ buffer->resize(static_cast<size_t>(size));
137+ return size == 0 || stream.read(buffer->data(), size).good();
138+}
139+ 
140+bool WriteFile(const std::string &path, const std::vector<char> &buffer)
141+{
142+ std::ofstream stream(path, std::ios::binary);
143+ if (!stream.is_open()) {
144+ return false;
145+ }
146+ stream.write(buffer.data(), static_cast<std::streamsize>(buffer.size()));
147+ return stream.good();
148+}
149+ 
150+aclError CreateAclTensor(
151+ const std::vector<int64_t> &shape, aclDataType dtype, void *device_addr, aclTensor **tensor)
152+{
153+ std::vector<int64_t> strides = MakeStrides(shape);
154+ const int64_t *shape_ptr = shape.empty() ? nullptr : shape.data();
155+ const int64_t *strides_ptr = strides.empty() ? nullptr : strides.data();
156+ *tensor = aclCreateTensor(
157+ shape_ptr, shape.size(), dtype, strides_ptr, 0, ACL_FORMAT_ND, shape_ptr, shape.size(), device_addr);
158+ return *tensor == nullptr ? ACL_ERROR_FAILURE : ACL_SUCCESS;
159+}
160+ 
161+int RunReluGradV2(const std::vector<char> &gradients_host, const std::vector<char> &features_host,
162+ const std::vector<int64_t> &shape, const ReluGradConfig &config, std::vector<char> *output_host,
163+ int32_t device_id)
164+{
165+ int final_ret = ACL_SUCCESS;
166+ bool acl_initialized = false;
167+ bool device_set = false;
168+ aclrtStream stream = nullptr;
169+ void *gradients_device = nullptr;
170+ void *features_device = nullptr;
171+ void *output_device = nullptr;
172+ void *workspace = nullptr;
173+ aclTensor *gradients_tensor = nullptr;
174+ aclTensor *features_tensor = nullptr;
175+ aclTensor *output_tensor = nullptr;
176+ aclScalar *threshold_scalar = nullptr;
177+ aclOpExecutor *executor = nullptr;
178+ uint64_t workspace_size = 0;
179+ const size_t bytes = static_cast<size_t>(GetShapeSize(shape)) * config.element_size;
180+ std::vector<char> zero_buffer(bytes, 0);
181+ 
182+ auto cleanup = [&]() -> int {
183+ if (gradients_tensor != nullptr) {
184+ aclDestroyTensor(gradients_tensor);
185+ }
186+ if (features_tensor != nullptr) {
187+ aclDestroyTensor(features_tensor);
188+ }
189+ if (output_tensor != nullptr) {
190+ aclDestroyTensor(output_tensor);
191+ }
192+ if (threshold_scalar != nullptr) {
193+ aclDestroyScalar(threshold_scalar);
194+ }
195+ if (workspace != nullptr) {
196+ aclrtFree(workspace);
197+ }
198+ if (gradients_device != nullptr) {
199+ aclrtFree(gradients_device);
200+ }
201+ if (features_device != nullptr) {
202+ aclrtFree(features_device);
203+ }
204+ if (output_device != nullptr) {
205+ aclrtFree(output_device);
206+ }
207+ if (stream != nullptr) {
208+ aclrtDestroyStream(stream);
209+ }
210+ if (device_set) {
211+ aclrtResetDevice(device_id);
212+ }
213+ if (acl_initialized) {
214+ aclFinalize();
215+ }
216+ return final_ret;
217+ };
218+ 
219+ auto ret = aclInit(nullptr);
220+ if (ret != ACL_SUCCESS) {
221+ final_ret = ReportAclError("aclInit", ret);
222+ return cleanup();
223+ }
224+ acl_initialized = true;
225+ ret = aclrtSetDevice(device_id);
226+ if (ret != ACL_SUCCESS) {
227+ final_ret = ReportAclError("aclrtSetDevice", ret);
228+ return cleanup();
229+ }
230+ device_set = true;
231+ ret = aclrtCreateStream(&stream);
232+ if (ret != ACL_SUCCESS) {
233+ final_ret = ReportAclError("aclrtCreateStream", ret);
234+ return cleanup();
235+ }
236+ 
237+ if (bytes > 0) {
238+ ret = aclrtMalloc(&gradients_device, bytes, ACL_MEM_MALLOC_HUGE_FIRST);
239+ if (ret != ACL_SUCCESS) {
240+ final_ret = ReportAclError("aclrtMalloc(gradients)", ret);
241+ return cleanup();
242+ }
243+ ret = aclrtMalloc(&features_device, bytes, ACL_MEM_MALLOC_HUGE_FIRST);
244+ if (ret != ACL_SUCCESS) {
245+ final_ret = ReportAclError("aclrtMalloc(features)", ret);
246+ return cleanup();
247+ }
248+ ret = aclrtMalloc(&output_device, bytes, ACL_MEM_MALLOC_HUGE_FIRST);
249+ if (ret != ACL_SUCCESS) {
250+ final_ret = ReportAclError("aclrtMalloc(output)", ret);
251+ return cleanup();
252+ }
253+ 
254+ ret = aclrtMemcpy(gradients_device, bytes, gradients_host.data(), bytes, ACL_MEMCPY_HOST_TO_DEVICE);
255+ if (ret != ACL_SUCCESS) {
256+ final_ret = ReportAclError("aclrtMemcpy(gradients H2D)", ret);
257+ return cleanup();
258+ }
259+ ret = aclrtMemcpy(features_device, bytes, features_host.data(), bytes, ACL_MEMCPY_HOST_TO_DEVICE);
260+ if (ret != ACL_SUCCESS) {
261+ final_ret = ReportAclError("aclrtMemcpy(features H2D)", ret);
262+ return cleanup();
263+ }
264+ ret = aclrtMemcpy(output_device, bytes, zero_buffer.data(), bytes, ACL_MEMCPY_HOST_TO_DEVICE);
265+ if (ret != ACL_SUCCESS) {
266+ final_ret = ReportAclError("aclrtMemcpy(output H2D)", ret);
267+ return cleanup();
268+ }
269+ }
270+ 
271+ ret = CreateAclTensor(shape, config.acl_dtype, gradients_device, &gradients_tensor);
272+ if (ret != ACL_SUCCESS) {
273+ final_ret = ReportAclError("CreateAclTensor(gradients)", ret);
274+ return cleanup();
275+ }
276+ ret = CreateAclTensor(shape, config.acl_dtype, features_device, &features_tensor);
277+ if (ret != ACL_SUCCESS) {
278+ final_ret = ReportAclError("CreateAclTensor(features)", ret);
279+ return cleanup();
280+ }
281+ ret = CreateAclTensor(shape, config.acl_dtype, output_device, &output_tensor);
282+ if (ret != ACL_SUCCESS) {
283+ final_ret = ReportAclError("CreateAclTensor(output)", ret);
284+ return cleanup();
285+ }
286+ int32_t threshold_value = 0;
287+ threshold_scalar = aclCreateScalar(&threshold_value, ACL_INT32);
288+ if (threshold_scalar == nullptr) {
289+ final_ret = ReportAclError("aclCreateScalar(threshold)", ACL_ERROR_FAILURE);
290+ return cleanup();
291+ }
292+ 
293+ ret = aclnnThresholdBackwardGetWorkspaceSize(
294+ gradients_tensor, features_tensor, threshold_scalar, output_tensor, &workspace_size, &executor);
295+ if (ret != ACL_SUCCESS) {
296+ final_ret = ReportAclError("aclnnThresholdBackwardGetWorkspaceSize", ret);
297+ return cleanup();
298+ }
299+ if (workspace_size > 0) {
300+ ret = aclrtMalloc(&workspace, workspace_size, ACL_MEM_MALLOC_HUGE_FIRST);
301+ if (ret != ACL_SUCCESS) {
302+ final_ret = ReportAclError("aclrtMalloc(workspace)", ret);
303+ return cleanup();
304+ }
305+ }
306+ 
307+ ret = aclnnThresholdBackward(workspace, workspace_size, executor, stream);
308+ if (ret != ACL_SUCCESS) {
309+ final_ret = ReportAclError("aclnnThresholdBackward", ret);
310+ return cleanup();
311+ }
312+ ret = aclrtSynchronizeStream(stream);
313+ if (ret != ACL_SUCCESS) {
314+ final_ret = ReportAclError("aclrtSynchronizeStream", ret);
315+ return cleanup();
316+ }
317+ 
318+ output_host->resize(bytes);
319+ if (bytes > 0) {
320+ ret = aclrtMemcpy(output_host->data(), bytes, output_device, bytes, ACL_MEMCPY_DEVICE_TO_HOST);
321+ if (ret != ACL_SUCCESS) {
322+ final_ret = ReportAclError("aclrtMemcpy(output D2H)", ret);
323+ return cleanup();
324+ }
325+ }
326+ 
327+ final_ret = ret;
328+ return cleanup();
329+}
330+ 
331+template <typename T>
332+std::vector<char> ToBytes(const std::vector<T> &values)
333+{
334+ std::vector<char> buffer(values.size() * sizeof(T));
335+ if (!buffer.empty()) {
336+ std::memcpy(buffer.data(), values.data(), buffer.size());
337+ }
338+ return buffer;
339+}
340+ 
341+int RunDefaultExample()
342+{
343+ ReluGradConfig config{ACL_FLOAT, sizeof(float), "fp32"};
344+ std::vector<int64_t> shape = {4, 2};
345+ std::vector<float> gradients = {1.0f, 2.0f, 3.0f, 4.0f, 5.0f, 6.0f, 7.0f, 8.0f};
346+ std::vector<float> features = {-4.0f, -3.0f, -2.0f, 0.0f, 1.0f, 2.0f, 4.0f, 5.0f};
347+ std::vector<char> output;
348+ auto ret = RunReluGradV2(ToBytes(gradients), ToBytes(features), shape, config, &output, 0);
349+ CHECK_RET(ret == ACL_SUCCESS, return ret);
350+ 
351+ const float *result = reinterpret_cast<const float *>(output.data());
352+ const std::vector<float> expected = {0.0f, 0.0f, 0.0f, 0.0f, 5.0f, 6.0f, 7.0f, 8.0f};
353+ for (size_t i = 0; i < expected.size(); ++i) {
354+ if (result[i] != expected[i]) {
355+ std::fprintf(stderr, "default check failed at %zu: got %.8f expected %.8f\n", i, result[i], expected[i]);
356+ return 1;
357+ }
358+ }
359+ std::printf("default example passed\n");
360+ return 0;
361+}
362+} // namespace
363+ 
364+int main(int argc, char **argv)
365+{
366+ if (argc == 1) {
367+ return RunDefaultExample();
368+ }
369+ 
370+ if (argc != 7) {
371+ std::fprintf(stderr,
372+ "Usage: %s <dtype> <shape|scalar> <gradients.bin> <features.bin> <output.bin> <device_id>\n",
373+ argv[0]);
374+ return 2;
375+ }
376+ 
377+ ReluGradConfig config{};
378+ CHECK_RET(ParseDtype(argv[1], &config), return 2);
379+ 
380+ std::vector<int64_t> shape;
381+ CHECK_RET(ParseShape(argv[2], &shape), return 2);
382+ 
383+ std::vector<char> gradients_host;
384+ CHECK_RET(ReadFile(argv[3], &gradients_host), return 3);
385+ std::vector<char> features_host;
386+ CHECK_RET(ReadFile(argv[4], &features_host), return 3);
387+ 
388+ std::vector<char> output_host;
389+ int ret = RunReluGradV2(gradients_host, features_host, shape, config, &output_host, std::atoi(argv[6]));
390+ CHECK_RET(ret == ACL_SUCCESS, return ret);
391+ CHECK_RET(WriteFile(argv[5], output_host), return 4);
392+ return 0;
393+}
Aexperimental/activation/relu_grad_v2/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 relu_grad_v2 ACLNNTYPE aclnn_exclude DISABLE_IN_OPP TRUE)
Aexperimental/activation/relu_grad_v2/op_host/op_api/aclnn_threshold_backward.cpp+159-0
@@ -0,0 +1,159 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include "aclnn_threshold_backward.h"
12+ 
13+#include <cmath>
14+#include <limits>
15+ 
16+#include "aclnn_kernels/common/op_error_check.h"
17+#include "aclnn_kernels/contiguous.h"
18+#include "op_api/aclnn_util.h"
19+#include "opdev/make_op_executor.h"
20+#include "opdev/op_dfx.h"
21+#include "opdev/op_executor.h"
22+#include "relu_grad_v2.h"
23+ 
24+using namespace op;
25+ 
26+#ifdef __cplusplus
27+extern "C" {
28+#endif
29+ 
30+namespace {
31+constexpr size_t MAX_DIM_LEN = 8;
32+constexpr float RELU_THRESHOLD = 0.0f;
33+ 
34+static const std::initializer_list<op::DataType> ASCEND910_DTYPE_SUPPORT_LIST = {
35+ op::DataType::DT_FLOAT,
36+ op::DataType::DT_FLOAT16,
37+ op::DataType::DT_INT8,
38+ op::DataType::DT_UINT8,
39+ op::DataType::DT_INT32,
40+ op::DataType::DT_INT64};
41+static const std::initializer_list<op::DataType> ASCEND910B_DTYPE_SUPPORT_LIST = {
42+ op::DataType::DT_FLOAT,
43+ op::DataType::DT_FLOAT16,
44+ op::DataType::DT_BF16,
45+ op::DataType::DT_INT8,
46+ op::DataType::DT_UINT8,
47+ op::DataType::DT_INT32,
48+ op::DataType::DT_INT64};
49+ 
50+static inline const std::initializer_list<op::DataType> &GetDtypeSupportList()
51+{
52+ if (GetCurrentPlatformInfo().GetSocVersion() >= SocVersion::ASCEND910B &&
53+ GetCurrentPlatformInfo().GetSocVersion() <= SocVersion::ASCEND910E) {
54+ return ASCEND910B_DTYPE_SUPPORT_LIST;
55+ }
56+ if (Ops::NN::AclnnUtil::IsRegbase()) {
57+ return ASCEND910B_DTYPE_SUPPORT_LIST;
58+ }
59+ return ASCEND910_DTYPE_SUPPORT_LIST;
60+}
61+ 
62+static bool CheckNotNull(const aclTensor *gradOutput, const aclTensor *self, const aclScalar *threshold,
63+ const aclTensor *out)
64+{
65+ OP_CHECK_NULL(gradOutput, return false);
66+ OP_CHECK_NULL(self, return false);
67+ OP_CHECK_NULL(threshold, return false);
68+ OP_CHECK_NULL(out, return false);
69+ return true;
70+}
71+ 
72+static bool CheckDtypeValid(const aclTensor *gradOutput, const aclTensor *self, const aclTensor *out)
73+{
74+ auto supportList = GetDtypeSupportList();
75+ OP_CHECK_DTYPE_NOT_SUPPORT(gradOutput, supportList, return false);
76+ OP_CHECK_DTYPE_NOT_SUPPORT(self, supportList, return false);
77+ OP_CHECK_DTYPE_NOT_SUPPORT(out, supportList, return false);
78+ OP_CHECK_DTYPE_NOT_MATCH(gradOutput, self->GetDataType(), return false);
79+ OP_CHECK_DTYPE_NOT_MATCH(gradOutput, out->GetDataType(), return false);
80+ return true;
81+}
82+ 
83+static bool CheckShape(const aclTensor *gradOutput, const aclTensor *self, const aclTensor *out)
84+{
85+ OP_CHECK_MAX_DIM(gradOutput, MAX_DIM_LEN, return false);
86+ OP_CHECK_MAX_DIM(self, MAX_DIM_LEN, return false);
87+ OP_CHECK_MAX_DIM(out, MAX_DIM_LEN, return false);
88+ OP_CHECK_SHAPE_NOT_EQUAL(gradOutput, self, return false);
89+ OP_CHECK_SHAPE_NOT_EQUAL(gradOutput, out, return false);
90+ return true;
91+}
92+ 
93+static bool CheckThresholdValue(const aclScalar *threshold)
94+{
95+ return std::fabs(threshold->ToFloat() - RELU_THRESHOLD) <= std::numeric_limits<float>::epsilon();
96+}
97+ 
98+static aclnnStatus CheckParams(const aclTensor *gradOutput, const aclTensor *self, const aclScalar *threshold,
99+ const aclTensor *out)
100+{
101+ CHECK_RET(CheckNotNull(gradOutput, self, threshold, out), ACLNN_ERR_PARAM_NULLPTR);
102+ CHECK_RET(CheckDtypeValid(gradOutput, self, out), ACLNN_ERR_PARAM_INVALID);
103+ CHECK_RET(CheckShape(gradOutput, self, out), ACLNN_ERR_PARAM_INVALID);
104+ CHECK_RET(CheckThresholdValue(threshold), ACLNN_ERR_PARAM_INVALID);
105+ return ACLNN_SUCCESS;
106+}
107+ 
108+static aclnnStatus ExecThresholdBackwardGetWorkspaceSize(const aclTensor *gradOutput, const aclTensor *self,
109+ const aclScalar *threshold, aclTensor *out,
110+ uint64_t *workspaceSize, aclOpExecutor **executor)
111+{
112+ auto uniqueExecutor = CREATE_EXECUTOR();
113+ CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR);
114+ 
115+ auto ret = CheckParams(gradOutput, self, threshold, out);
116+ CHECK_RET(ret == ACLNN_SUCCESS, ret);
117+ 
118+ if (gradOutput->IsEmpty() || self->IsEmpty()) {
119+ *workspaceSize = 0;
120+ uniqueExecutor.ReleaseTo(executor);
121+ return ACLNN_SUCCESS;
122+ }
123+ 
124+ auto gradOutputContiguous = l0op::Contiguous(gradOutput, uniqueExecutor.get());
125+ CHECK_RET(gradOutputContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
126+ auto selfContiguous = l0op::Contiguous(self, uniqueExecutor.get());
127+ CHECK_RET(selfContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
128+ 
129+ auto thresholdBackwardOut = l0op::ReluGradV2(gradOutputContiguous, selfContiguous, uniqueExecutor.get());
130+ CHECK_RET(thresholdBackwardOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
131+ 
132+ auto viewCopyResult = l0op::ViewCopy(thresholdBackwardOut, out, uniqueExecutor.get());
133+ CHECK_RET(viewCopyResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
134+ 
135+ *workspaceSize = uniqueExecutor->GetWorkspaceSize();
136+ uniqueExecutor.ReleaseTo(executor);
137+ return ACLNN_SUCCESS;
138+}
139+} // namespace
140+ 
141+aclnnStatus aclnnThresholdBackwardGetWorkspaceSize(const aclTensor *gradOutput, const aclTensor *self,
142+ const aclScalar *threshold, aclTensor *out,
143+ uint64_t *workspaceSize, aclOpExecutor **executor)
144+{
145+ OP_CHECK_COMM_INPUT(workspaceSize, executor);
146+ L2_DFX_PHASE_1(aclnnThresholdBackward, DFX_IN(gradOutput, self, threshold), DFX_OUT(out));
147+ return ExecThresholdBackwardGetWorkspaceSize(gradOutput, self, threshold, out, workspaceSize, executor);
148+}
149+ 
150+aclnnStatus aclnnThresholdBackward(void *workspace, uint64_t workspaceSize, aclOpExecutor *executor,
151+ const aclrtStream stream)
152+{
153+ L2_DFX_PHASE_2(aclnnThresholdBackward);
154+ return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);
155+}
156+ 
157+#ifdef __cplusplus
158+}
159+#endif
Aexperimental/activation/relu_grad_v2/op_host/op_api/aclnn_threshold_backward.h+32-0
@@ -0,0 +1,32 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#ifndef ACLNN_THRESHOLD_BACKWARD_H_
12+#define ACLNN_THRESHOLD_BACKWARD_H_
13+ 
14+#include "aclnn/aclnn_base.h"
15+#include "aclnn_util.h"
16+ 
17+#ifdef __cplusplus
18+extern "C" {
19+#endif
20+ 
21+ACLNN_API aclnnStatus aclnnThresholdBackwardGetWorkspaceSize(const aclTensor *gradOutput, const aclTensor *self,
22+ const aclScalar *threshold, aclTensor *out,
23+ uint64_t *workspaceSize, aclOpExecutor **executor);
24+ 
25+ACLNN_API aclnnStatus aclnnThresholdBackward(void *workspace, uint64_t workspaceSize, aclOpExecutor *executor,
26+ const aclrtStream stream);
27+ 
28+#ifdef __cplusplus
29+}
30+#endif
31+ 
32+#endif
Aexperimental/activation/relu_grad_v2/op_host/op_api/relu_grad_v2.cpp+38-0
@@ -0,0 +1,38 @@
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 "relu_grad_v2.h"
12+#include "aclnn_kernels/common/op_error_check.h"
13+#include "opdev/make_op_executor.h"
14+#include "opdev/op_def.h"
15+#include "opdev/op_dfx.h"
16+#include "opdev/op_executor.h"
17+#include "opdev/op_log.h"
18+ 
19+using namespace op;
20+ 
21+namespace l0op {
22+ 
23+OP_TYPE_REGISTER(ReluGradV2);
24+ 
25+const aclTensor *ReluGradV2(const aclTensor *gradients, const aclTensor *features, aclOpExecutor *executor)
26+{
27+ L0_DFX(ReluGradV2, gradients, features);
28+ auto out = executor->AllocTensor(
29+ gradients->GetStorageShape(), gradients->GetDataType(), gradients->GetStorageFormat());
30+ CHECK_RET(out != nullptr, nullptr);
31+ 
32+ auto ret = ADD_TO_LAUNCHER_LIST_AICORE(ReluGradV2, OP_INPUT(gradients, features), OP_OUTPUT(out));
33+ OP_CHECK(ret == ACLNN_SUCCESS, OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "ReluGradV2 ADD_TO_LAUNCHER_LIST_AICORE failed."),
34+ return nullptr);
35+ return out;
36+}
37+ 
38+} // namespace l0op
Aexperimental/activation/relu_grad_v2/op_host/op_api/relu_grad_v2.h+21-0
@@ -0,0 +1,21 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#ifndef OP_API_OP_API_COMMON_INC_LEVEL0_OP_RELU_GRAD_V2_H_
12+#define OP_API_OP_API_COMMON_INC_LEVEL0_OP_RELU_GRAD_V2_H_
13+ 
14+#include "opdev/make_op_executor.h"
15+#include "opdev/op_executor.h"
16+ 
17+namespace l0op {
18+const aclTensor *ReluGradV2(const aclTensor *gradients, const aclTensor *features, aclOpExecutor *executor);
19+}
20+ 
21+#endif
Aexperimental/activation/relu_grad_v2/op_host/relu_grad_v2_def.cpp+59-0
@@ -0,0 +1,59 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include "register/op_def_registry.h"
12+ 
13+namespace ops {
14+class ReluGradV2 : public OpDef {
15+public:
16+ explicit ReluGradV2(const char *name) : OpDef(name)
17+ {
18+ this->Input("gradients")
19+ .ParamType(REQUIRED)
20+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16, ge::DT_INT8, ge::DT_UINT8, ge::DT_INT32,
21+ ge::DT_INT64})
22+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
23+ ge::FORMAT_ND})
24+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
25+ ge::FORMAT_ND, ge::FORMAT_ND})
26+ .AutoContiguous();
27+ this->Input("features")
28+ .ParamType(REQUIRED)
29+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16, ge::DT_INT8, ge::DT_UINT8, ge::DT_INT32,
30+ ge::DT_INT64})
31+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
32+ ge::FORMAT_ND})
33+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
34+ ge::FORMAT_ND, ge::FORMAT_ND})
35+ .AutoContiguous();
36+ this->Output("backprops")
37+ .ParamType(REQUIRED)
38+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16, ge::DT_INT8, ge::DT_UINT8, ge::DT_INT32,
39+ ge::DT_INT64})
40+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
41+ ge::FORMAT_ND})
42+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
43+ ge::FORMAT_ND, ge::FORMAT_ND})
44+ .AutoContiguous();
45+ 
46+ OpAICoreConfig aicoreConfig;
47+ aicoreConfig.DynamicCompileStaticFlag(true)
48+ .DynamicFormatFlag(false)
49+ .DynamicRankSupportFlag(true)
50+ .DynamicShapeSupportFlag(true)
51+ .NeedCheckSupportFlag(false)
52+ .PrecisionReduceFlag(true)
53+ .ExtendCfgInfo("opFile.value", "relu_grad_v2");
54+ this->AICore().AddConfig("ascend910b", aicoreConfig);
55+ }
56+};
57+ 
58+OP_ADD(ReluGradV2);
59+} // namespace ops
Aexperimental/activation/relu_grad_v2/op_host/relu_grad_v2_infershape.cpp+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+#include "register/op_impl_registry.h"
12+#include "log/log.h"
13+ 
14+namespace ops {
15+ 
16+static ge::graphStatus InferShapeReluGradV2(gert::InferShapeContext *context)
17+{
18+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
19+ 
20+ const gert::Shape *gradientsShape = context->GetInputShape(0);
21+ OP_CHECK_NULL_WITH_CONTEXT(context, gradientsShape);
22+ 
23+ gert::Shape *backpropsShape = context->GetOutputShape(0);
24+ OP_CHECK_NULL_WITH_CONTEXT(context, backpropsShape);
25+ 
26+ *backpropsShape = *gradientsShape;
27+ return ge::GRAPH_SUCCESS;
28+}
29+ 
30+IMPL_OP_INFERSHAPE(ReluGradV2).InferShape(InferShapeReluGradV2);
31+} // namespace ops
Aexperimental/activation/relu_grad_v2/op_host/relu_grad_v2_tiling.cpp+236-0
@@ -0,0 +1,236 @@
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 <algorithm>
12+#include <set>
13+ 
14+#include "log/log.h"
15+#include "op_host/tiling_templates_registry.h"
16+#include "op_host/tiling_util.h"
17+#include "tiling/platform/platform_ascendc.h"
18+#include "../op_kernel/relu_grad_v2_tiling_data.h"
19+#include "../op_kernel/relu_grad_v2_tiling_key.h"
20+ 
21+namespace optiling {
22+ 
23+using namespace Ops::NN::OpTiling;
24+ 
25+constexpr int64_t CACHE_LINE_BYTE_LENGTH = 512;
26+constexpr int64_t COMPARE_ALIGN_BYTES = 256;
27+constexpr int64_t SCALAR_FP32_BUFFER_COEFFICIENT = 25;
28+constexpr int64_t CAST_SCALAR_FP16_BF16_BUFFER_COEFFICIENT = 24;
29+ 
30+struct ReluGradV2CompileInfo {};
31+ 
32+struct ReluGradV2TilingInfo {
33+ int64_t formerNum = 0;
34+ int64_t formerLength = 0;
35+ int64_t tailLength = 0;
36+ int64_t tileLength = 1;
37+ uint32_t blockDim = 1;
38+};
39+ 
40+static const gert::Shape SCALAR_SHAPE = {1};
41+ 
42+static const gert::Shape &EnsureNotScalar(const gert::Shape &shape)
43+{
44+ if (shape.GetDimNum() == 0) {
45+ return SCALAR_SHAPE;
46+ }
47+ return shape;
48+}
49+ 
50+static ge::graphStatus GetPlatformInfo(gert::TilingContext *context, uint64_t &ubSize, int64_t &coreNum)
51+{
52+ auto *platformInfoPtr = context->GetPlatformInfo();
53+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr);
54+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
55+ coreNum = static_cast<int64_t>(ascendcPlatform.GetCoreNumAiv());
56+ OP_CHECK_IF(coreNum <= 0, OP_LOGE(context, "coreNum is invalid"), return ge::GRAPH_FAILED);
57+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
58+ OP_CHECK_IF(ubSize == 0, OP_LOGE(context, "ubSize is 0"), return ge::GRAPH_FAILED);
59+ return ge::GRAPH_SUCCESS;
60+}
61+ 
62+static ge::graphStatus GetShapeAttrsInfo(gert::TilingContext *context, int64_t &totalLength, ge::DataType &dataType)
63+{
64+ auto *gradientsShape = context->GetInputShape(0);
65+ auto *featuresShape = context->GetInputShape(1);
66+ OP_CHECK_NULL_WITH_CONTEXT(context, gradientsShape);
67+ OP_CHECK_NULL_WITH_CONTEXT(context, featuresShape);
68+ 
69+ const gert::Shape &gradientsStorageShape = EnsureNotScalar(gradientsShape->GetStorageShape());
70+ const gert::Shape &featuresStorageShape = EnsureNotScalar(featuresShape->GetStorageShape());
71+ OP_CHECK_IF(gradientsStorageShape != featuresStorageShape, OP_LOGE(context, "shape mismatch"),
72+ return ge::GRAPH_FAILED);
73+ totalLength = gradientsStorageShape.GetShapeSize();
74+ 
75+ auto *gradientsDesc = context->GetInputDesc(0);
76+ auto *featuresDesc = context->GetInputDesc(1);
77+ OP_CHECK_NULL_WITH_CONTEXT(context, gradientsDesc);
78+ OP_CHECK_NULL_WITH_CONTEXT(context, featuresDesc);
79+ dataType = gradientsDesc->GetDataType();
80+ 
81+ const std::set<ge::DataType> supportedDtype = {
82+ ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16, ge::DT_INT8, ge::DT_UINT8, ge::DT_INT32, ge::DT_INT64};
83+ OP_CHECK_IF(supportedDtype.count(dataType) == 0, OP_LOGE(context, "invalid dtype"), return ge::GRAPH_FAILED);
84+ OP_CHECK_IF(featuresDesc->GetDataType() != dataType, OP_LOGE(context, "dtype mismatch"), return ge::GRAPH_FAILED);
85+ return ge::GRAPH_SUCCESS;
86+}
87+ 
88+static ge::graphStatus GetWorkspaceSize(gert::TilingContext *context)
89+{
90+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
91+ uint32_t sysWorkspaceSize = ascendcPlatform.GetLibApiWorkSpaceSize();
92+ size_t *currentWorkspace = context->GetWorkspaceSizes(1);
93+ OP_CHECK_NULL_WITH_CONTEXT(context, currentWorkspace);
94+ currentWorkspace[0] = sysWorkspaceSize;
95+ return ge::GRAPH_SUCCESS;
96+}
97+ 
98+static void SetEmptyTiling(ReluGradV2TilingData *tiling)
99+{
100+ tiling->formerNum = 0;
101+ tiling->formerLength = 0;
102+ tiling->tailLength = 0;
103+ tiling->tileLength = 1;
104+}
105+ 
106+static bool IsFloatFamily(ge::DataType dataType)
107+{
108+ return dataType == ge::DT_FLOAT || dataType == ge::DT_FLOAT16 || dataType == ge::DT_BF16;
109+}
110+ 
111+static int64_t GetBufferCoefficient(ge::DataType dataType, int64_t dtypeSize)
112+{
113+ if (dataType == ge::DT_FLOAT) {
114+ return SCALAR_FP32_BUFFER_COEFFICIENT;
115+ }
116+ if (dataType == ge::DT_FLOAT16 || dataType == ge::DT_BF16) {
117+ return CAST_SCALAR_FP16_BF16_BUFFER_COEFFICIENT;
118+ }
119+ return dtypeSize * 6;
120+}
121+ 
122+static int64_t GetAlignElements(ge::DataType dataType, int64_t dtypeSize)
123+{
124+ if (dtypeSize == 0) {
125+ return 1;
126+ }
127+ int64_t alignElements = 32 / dtypeSize;
128+ if (alignElements <= 0) {
129+ alignElements = 1;
130+ }
131+ if (IsFloatFamily(dataType)) {
132+ int64_t compareAlignElements = COMPARE_ALIGN_BYTES / static_cast<int64_t>(sizeof(float));
133+ alignElements = std::max<int64_t>(alignElements, compareAlignElements);
134+ }
135+ return alignElements;
136+}
137+ 
138+static ge::graphStatus BuildTilingData(int64_t totalLength, ge::DataType dataType, uint64_t ubSize, int64_t coreNum,
139+ ReluGradV2TilingInfo &info)
140+{
141+ uint32_t typeLength = 0;
142+ ge::TypeUtils::GetDataTypeLength(dataType, typeLength);
143+ if (typeLength == 0) {
144+ return ge::GRAPH_FAILED;
145+ }
146+ int64_t dtypeSize = static_cast<int64_t>(typeLength);
147+ 
148+ if (dtypeSize == 0) {
149+ return ge::GRAPH_FAILED;
150+ }
151+ int64_t cacheLineElements = std::max<int64_t>(1, CACHE_LINE_BYTE_LENGTH / dtypeSize);
152+ if (coreNum == 0) {
153+ return ge::GRAPH_FAILED;
154+ }
155+ int64_t totalLengthCore = (totalLength + coreNum - 1) / coreNum;
156+ int64_t totalLengthCoreAlign = ((totalLengthCore + cacheLineElements - 1) / cacheLineElements) * cacheLineElements;
157+ int64_t usedCoreNum = std::max<int64_t>(1, (totalLength + totalLengthCoreAlign - 1) / totalLengthCoreAlign);
158+ 
159+ info.formerNum = usedCoreNum - 1;
160+ info.formerLength = totalLengthCoreAlign;
161+ info.tailLength = totalLength - info.formerNum * info.formerLength;
162+ info.blockDim = static_cast<uint32_t>(usedCoreNum);
163+ 
164+ int64_t bufferCoefficient = GetBufferCoefficient(dataType, dtypeSize);
165+ if (bufferCoefficient == 0) {
166+ return ge::GRAPH_FAILED;
167+ }
168+ int64_t maxTileElements = static_cast<int64_t>(ubSize) / bufferCoefficient;
169+ int64_t alignElements = GetAlignElements(dataType, dtypeSize);
170+ if (alignElements == 0) {
171+ return ge::GRAPH_FAILED;
172+ }
173+ info.tileLength = (maxTileElements / alignElements) * alignElements;
174+ if (info.tileLength <= 0) {
175+ info.tileLength = alignElements;
176+ }
177+ return ge::GRAPH_SUCCESS;
178+}
179+ 
180+static ge::graphStatus SetTilingData(gert::TilingContext *context, const ReluGradV2TilingInfo &info,
181+ ReluGradV2TilingData *tiling)
182+{
183+ tiling->formerNum = info.formerNum;
184+ tiling->formerLength = info.formerLength;
185+ tiling->tailLength = info.tailLength;
186+ tiling->tileLength = info.tileLength;
187+ context->SetBlockDim(info.blockDim);
188+ return ge::GRAPH_SUCCESS;
189+}
190+ 
191+static ge::graphStatus ReluGradV2TilingFunc(gert::TilingContext *context)
192+{
193+ uint64_t ubSize = 0;
194+ int64_t coreNum = 0;
195+ OP_CHECK_IF(GetPlatformInfo(context, ubSize, coreNum) != ge::GRAPH_SUCCESS,
196+ OP_LOGE(context, "GetPlatformInfo error"), return ge::GRAPH_FAILED);
197+ 
198+ int64_t totalLength = 0;
199+ ge::DataType dataType;
200+ OP_CHECK_IF(GetShapeAttrsInfo(context, totalLength, dataType) != ge::GRAPH_SUCCESS,
201+ OP_LOGE(context, "GetShapeAttrsInfo error"), return ge::GRAPH_FAILED);
202+ 
203+ OP_CHECK_IF(GetWorkspaceSize(context) != ge::GRAPH_SUCCESS,
204+ OP_LOGE(context, "GetWorkspaceSize error"), return ge::GRAPH_FAILED);
205+ 
206+ auto *tiling = context->GetTilingData<ReluGradV2TilingData>();
207+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
208+ OP_CHECK_IF(memset_s(tiling, sizeof(ReluGradV2TilingData), 0, sizeof(ReluGradV2TilingData)) != EOK,
209+ OP_LOGE(context, "set tiling data error"), return ge::GRAPH_FAILED);
210+ 
211+ uint32_t dTypeX = static_cast<uint32_t>(dataType);
212+ if (totalLength <= 0) {
213+ SetEmptyTiling(tiling);
214+ context->SetBlockDim(1);
215+ ASCENDC_TPL_SEL_PARAM(context, dTypeX);
216+ return ge::GRAPH_SUCCESS;
217+ }
218+ 
219+ ReluGradV2TilingInfo info;
220+ OP_CHECK_IF(BuildTilingData(totalLength, dataType, ubSize, coreNum, info) != ge::GRAPH_SUCCESS,
221+ OP_LOGE(context, "BuildTilingData error"), return ge::GRAPH_FAILED);
222+ OP_CHECK_IF(SetTilingData(context, info, tiling) != ge::GRAPH_SUCCESS,
223+ OP_LOGE(context, "SetTilingData error"), return ge::GRAPH_FAILED);
224+ ASCENDC_TPL_SEL_PARAM(context, dTypeX);
225+ return ge::GRAPH_SUCCESS;
226+}
227+ 
228+static ge::graphStatus TilingParseForReluGradV2([[maybe_unused]] gert::TilingParseContext *context)
229+{
230+ return ge::GRAPH_SUCCESS;
231+}
232+ 
233+IMPL_OP_OPTILING(ReluGradV2)
234+ .Tiling(ReluGradV2TilingFunc)
235+ .TilingParse<ReluGradV2CompileInfo>(TilingParseForReluGradV2);
236+} // namespace optiling
Aexperimental/activation/relu_grad_v2/op_kernel/relu_grad_v2.cpp+40-0
@@ -0,0 +1,40 @@
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 <type_traits>
12+ 
13+#include "relu_grad_v2.h"
14+ 
15+template <typename D_T_X>
16+__global__ __aicore__ void relu_grad_v2(
17+ GM_ADDR gradients, GM_ADDR features, GM_ADDR backprops, GM_ADDR workspace, GM_ADDR tiling)
18+{
19+ REGISTER_TILING_DEFAULT(ReluGradV2TilingData);
20+ GET_TILING_DATA_WITH_STRUCT(ReluGradV2TilingData, tilingData, tiling);
21+ AscendC::TPipe pipe;
22+ 
23+ if constexpr (std::is_same_v<D_T_X, float>) {
24+ NsReluGradV2::KernelReluGradScalar<D_T_X> op;
25+ op.Init(gradients, features, backprops, &tilingData, &pipe);
26+ op.Process();
27+ } else if constexpr (std::is_same_v<D_T_X, half>) {
28+ NsReluGradV2::KernelReluGradCastScalar<D_T_X> op;
29+ op.Init(gradients, features, backprops, &tilingData, &pipe);
30+ op.Process();
31+ } else if constexpr (std::is_same_v<D_T_X, bfloat16_t>) {
32+ NsReluGradV2::KernelReluGradCastScalar<D_T_X> op;
33+ op.Init(gradients, features, backprops, &tilingData, &pipe);
34+ op.Process();
35+ } else {
36+ NsReluGradV2::KernelReluGradScalar<D_T_X> op;
37+ op.Init(gradients, features, backprops, &tilingData, &pipe);
38+ op.Process();
39+ }
40+}
Aexperimental/activation/relu_grad_v2/op_kernel/relu_grad_v2.h+230-0
@@ -0,0 +1,230 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#ifndef EXPERIMENTAL_ACTIVATION_RELU_GRAD_V2_H_
12+#define EXPERIMENTAL_ACTIVATION_RELU_GRAD_V2_H_
13+ 
14+#include <type_traits>
15+ 
16+#include "kernel_operator.h"
17+#include "kernel_tiling/kernel_tiling.h"
18+#include "relu_grad_v2_tiling_data.h"
19+#include "relu_grad_v2_tiling_key.h"
20+ 
21+namespace NsReluGradV2 {
22+ 
23+using namespace AscendC;
24+ 
25+constexpr int32_t BUFFER_NUM = 2;
26+constexpr int64_t DATA_COPY_ALIGN_BYTES = 32;
27+ 
28+__aicore__ inline int64_t AlignUp(int64_t value, int64_t align)
29+{
30+ if (align == 0) {
31+ return value;
32+ }
33+ if (align < 0) {
34+ return value;
35+ }
36+ int64_t remainder = value % align;
37+ if (remainder == 0) {
38+ return value;
39+ }
40+ return value + align - remainder;
41+}
42+ 
43+template <typename T, typename Derived>
44+class KernelReluGradBase {
45+public:
46+ __aicore__ inline KernelReluGradBase() = default;
47+ 
48+ __aicore__ inline void Init(GM_ADDR gradients, GM_ADDR features, GM_ADDR backprops,
49+ const ReluGradV2TilingData *tilingData, TPipe *pipe)
50+ {
51+ pipe_ = pipe;
52+ InitGlobalTensors(gradients, features, backprops, tilingData);
53+ tileLength_ = tilingData->tileLength;
54+ pipe_->InitBuffer(inQueueGradients_, BUFFER_NUM, tileLength_ * sizeof(T));
55+ pipe_->InitBuffer(inQueueFeatures_, BUFFER_NUM, tileLength_ * sizeof(T));
56+ pipe_->InitBuffer(outQueueBackprops_, BUFFER_NUM, tileLength_ * sizeof(T));
57+ static_cast<Derived *>(this)->InitExtraBuffers();
58+ }
59+ 
60+ __aicore__ inline void Process()
61+ {
62+ int64_t tileNum = (blockLength_ + tileLength_ - 1) / tileLength_;
63+ if (tileNum == 0) {
64+ return;
65+ }
66+ 
67+ Derived *derived = static_cast<Derived *>(this);
68+ int64_t alignNum = derived->GetAlignNum();
69+ for (int64_t i = 0; i < tileNum; ++i) {
70+ int64_t validLength = tileLength_;
71+ if (i == tileNum - 1) {
72+ validLength = blockLength_ - (tileNum - 1) * tileLength_;
73+ }
74+ int64_t paddedLength = AlignUp(validLength, alignNum);
75+ CopyIn(i, validLength, paddedLength);
76+ derived->Compute(derived->GetComputeLength(validLength, paddedLength));
77+ CopyOut(i, validLength);
78+ }
79+ }
80+ 
81+ __aicore__ inline int64_t GetAlignNum() const
82+ {
83+ return DATA_COPY_ALIGN_BYTES / static_cast<int64_t>(sizeof(T));
84+ }
85+ 
86+ __aicore__ inline int64_t GetComputeLength(int64_t validLength, int64_t paddedLength) const
87+ {
88+ (void)paddedLength;
89+ return validLength;
90+ }
91+ 
92+protected:
93+ __aicore__ inline void CopyIn(int64_t progress, int64_t validLength, int64_t paddedLength)
94+ {
95+ DataCopyExtParams copyParams{1, static_cast<uint32_t>(validLength * sizeof(T)), 0, 0, 0};
96+ DataCopyPadExtParams<T> padParams{true, 0, static_cast<uint8_t>(paddedLength - validLength), static_cast<T>(0)};
97+ 
98+ LocalTensor<T> gradientsLocal = inQueueGradients_.AllocTensor<T>();
99+ LocalTensor<T> featuresLocal = inQueueFeatures_.AllocTensor<T>();
100+ 
101+ DataCopyPad(gradientsLocal, gradientsGm_[progress * tileLength_], copyParams, padParams);
102+ DataCopyPad(featuresLocal, featuresGm_[progress * tileLength_], copyParams, padParams);
103+ 
104+ inQueueGradients_.EnQue(gradientsLocal);
105+ inQueueFeatures_.EnQue(featuresLocal);
106+ }
107+ 
108+ __aicore__ inline void CopyOut(int64_t progress, int64_t validLength)
109+ {
110+ LocalTensor<T> backpropsLocal = outQueueBackprops_.DeQue<T>();
111+ DataCopyExtParams copyParams{1, static_cast<uint32_t>(validLength * sizeof(T)), 0, 0, 0};
112+ DataCopyPad(backpropsGm_[progress * tileLength_], backpropsLocal, copyParams);
113+ outQueueBackprops_.FreeTensor(backpropsLocal);
114+ }
115+ 
116+ __aicore__ inline void FinishCompute(LocalTensor<T> backpropsLocal, LocalTensor<T> gradientsLocal,
117+ LocalTensor<T> featuresLocal)
118+ {
119+ outQueueBackprops_.EnQue<T>(backpropsLocal);
120+ inQueueGradients_.FreeTensor(gradientsLocal);
121+ inQueueFeatures_.FreeTensor(featuresLocal);
122+ }
123+ 
124+protected:
125+ TQue<TPosition::VECIN, BUFFER_NUM> inQueueGradients_;
126+ TQue<TPosition::VECIN, BUFFER_NUM> inQueueFeatures_;
127+ TQue<TPosition::VECOUT, BUFFER_NUM> outQueueBackprops_;
128+ 
129+private:
130+ __aicore__ inline void InitGlobalTensors(GM_ADDR gradients, GM_ADDR features, GM_ADDR backprops,
131+ const ReluGradV2TilingData *tilingData)
132+ {
133+ int64_t blockIdx = GetBlockIdx();
134+ int64_t offset = 0;
135+ int64_t blockSize = tilingData->tailLength;
136+ if (blockIdx < tilingData->formerNum) {
137+ offset = tilingData->formerLength * blockIdx;
138+ blockSize = tilingData->formerLength;
139+ } else {
140+ offset = tilingData->formerLength * tilingData->formerNum;
141+ }
142+ 
143+ blockLength_ = blockSize;
144+ gradientsGm_.SetGlobalBuffer((__gm__ T *)gradients + offset, blockSize);
145+ featuresGm_.SetGlobalBuffer((__gm__ T *)features + offset, blockSize);
146+ backpropsGm_.SetGlobalBuffer((__gm__ T *)backprops + offset, blockSize);
147+ }
148+ 
149+protected:
150+ TPipe *pipe_ = nullptr;
151+ GlobalTensor<T> gradientsGm_;
152+ GlobalTensor<T> featuresGm_;
153+ GlobalTensor<T> backpropsGm_;
154+ int64_t blockLength_ = 0;
155+ int64_t tileLength_ = 0;
156+};
157+ 
158+template <typename T>
159+class KernelReluGradScalar : public KernelReluGradBase<T, KernelReluGradScalar<T>> {
160+public:
161+ __aicore__ inline KernelReluGradScalar() = default;
162+ 
163+ __aicore__ inline void InitExtraBuffers() {}
164+ 
165+ __aicore__ inline void Compute(int64_t validLength)
166+ {
167+ LocalTensor<T> gradientsLocal = this->inQueueGradients_.template DeQue<T>();
168+ LocalTensor<T> featuresLocal = this->inQueueFeatures_.template DeQue<T>();
169+ LocalTensor<T> backpropsLocal = this->outQueueBackprops_.template AllocTensor<T>();
170+ 
171+ for (int64_t idx = 0; idx < validLength; ++idx) {
172+ bool keepGradient = false;
173+ if constexpr (std::is_same_v<T, float> || std::is_same_v<T, half> || std::is_same_v<T, bfloat16_t>) {
174+ float featureValue = static_cast<float>(featuresLocal.GetValue(idx));
175+ keepGradient = featureValue > 0.0f || featureValue != featureValue;
176+ } else {
177+ T featureValue = featuresLocal.GetValue(idx);
178+ keepGradient = featureValue > static_cast<T>(0);
179+ }
180+ backpropsLocal.SetValue(idx, keepGradient ? gradientsLocal.GetValue(idx) : static_cast<T>(0));
181+ }
182+ 
183+ this->FinishCompute(backpropsLocal, gradientsLocal, featuresLocal);
184+ }
185+};
186+ 
187+template <typename T>
188+class KernelReluGradCastScalar : public KernelReluGradBase<T, KernelReluGradCastScalar<T>> {
189+public:
190+ __aicore__ inline KernelReluGradCastScalar() = default;
191+ 
192+ __aicore__ inline void InitExtraBuffers()
193+ {
194+ this->pipe_->InitBuffer(gradientsFp32Buf_, this->tileLength_ * sizeof(float));
195+ this->pipe_->InitBuffer(featuresFp32Buf_, this->tileLength_ * sizeof(float));
196+ this->pipe_->InitBuffer(backpropsFp32Buf_, this->tileLength_ * sizeof(float));
197+ }
198+ 
199+ __aicore__ inline void Compute(int64_t validLength)
200+ {
201+ LocalTensor<T> gradientsLocal = this->inQueueGradients_.template DeQue<T>();
202+ LocalTensor<T> featuresLocal = this->inQueueFeatures_.template DeQue<T>();
203+ LocalTensor<T> backpropsLocal = this->outQueueBackprops_.template AllocTensor<T>();
204+ LocalTensor<float> gradientsFp32 = gradientsFp32Buf_.Get<float>();
205+ LocalTensor<float> featuresFp32 = featuresFp32Buf_.Get<float>();
206+ LocalTensor<float> backpropsFp32 = backpropsFp32Buf_.Get<float>();
207+ 
208+ Cast(gradientsFp32, gradientsLocal, RoundMode::CAST_NONE, static_cast<uint32_t>(validLength));
209+ Cast(featuresFp32, featuresLocal, RoundMode::CAST_NONE, static_cast<uint32_t>(validLength));
210+ PipeBarrier<PIPE_V>();
211+ 
212+ for (int64_t idx = 0; idx < validLength; ++idx) {
213+ float featureValue = featuresFp32.GetValue(idx);
214+ backpropsFp32.SetValue(idx, featureValue > 0.0f || featureValue != featureValue ? gradientsFp32.GetValue(idx)
215+ : 0.0f);
216+ }
217+ 
218+ Cast(backpropsLocal, backpropsFp32, RoundMode::CAST_ROUND, static_cast<uint32_t>(validLength));
219+ this->FinishCompute(backpropsLocal, gradientsLocal, featuresLocal);
220+ }
221+ 
222+private:
223+ TBuf<TPosition::VECCALC> gradientsFp32Buf_;
224+ TBuf<TPosition::VECCALC> featuresFp32Buf_;
225+ TBuf<TPosition::VECCALC> backpropsFp32Buf_;
226+};
227+ 
228+} // namespace NsReluGradV2
229+ 
230+#endif
Aexperimental/activation/relu_grad_v2/op_kernel/relu_grad_v2_tiling_data.h+23-0
@@ -0,0 +1,23 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#ifndef RELU_GRAD_V2_TILING_DATA_H_
12+#define RELU_GRAD_V2_TILING_DATA_H_
13+ 
14+#include <cstdint>
15+ 
16+struct ReluGradV2TilingData {
17+ int64_t formerNum;
18+ int64_t formerLength;
19+ int64_t tailLength;
20+ int64_t tileLength;
21+};
22+ 
23+#endif
Aexperimental/activation/relu_grad_v2/op_kernel/relu_grad_v2_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+#ifndef __RELU_GRAD_V2_TILING_KEY_H__
12+#define __RELU_GRAD_V2_TILING_KEY_H__
13+ 
14+#include "ascendc/host_api/tiling/template_argument.h"
15+ 
16+ASCENDC_TPL_ARGS_DECL(
17+ ReluGradV2,
18+ ASCENDC_TPL_DATATYPE_DECL(D_T_X, C_DT_FLOAT, C_DT_FLOAT16, C_DT_BF16, C_DT_INT8, C_DT_UINT8, C_DT_INT32,
19+ C_DT_INT64, ASCENDC_TPL_INPUT(0)));
20+ 
21+ASCENDC_TPL_SEL(
22+ ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_FLOAT)),
23+ ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_FLOAT16)),
24+ ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_BF16)),
25+ ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_INT8)),
26+ ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_UINT8)),
27+ ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_INT32)),
28+ ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_INT64)),
29+);
30+ 
31+#endif
Aexperimental/activation/relu_grad_v2/tests/CMakeLists.txt+16-0
@@ -0,0 +1,16 @@
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+foreach(SUB_DIR ${CURRENT_DIRS})
13+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
14+ add_subdirectory(${SUB_DIR})
15+ endif()
16+endforeach()
Aexperimental/activation/relu_grad_v2/tests/st/aclnnThresholdBackward/all_aclnnThresholdBackward.json+405-0
@@ -0,0 +1,405 @@
1+[
2+ {
3+ "id": 1,
4+ "name": "torch.threshold_backward fp16",
5+ "aclnn_name": "aclnnThresholdBackward",
6+ "version": "v1.0",
7+ "expected_error_msg": null,
8+ "api": "pytorch",
9+ "api_type": "aclnn_threshold_backward",
10+ "aclnn_api_type": "aclnn_function",
11+ "backward": false,
12+ "standard": {
13+ "acc": "cv_fused_double_benchmark",
14+ "perf": "not_key"
15+ },
16+ "outputs": null,
17+ "inputs": [
18+ {
19+ "name": "gradOutput",
20+ "type": "tensor",
21+ "required": true,
22+ "dtype": "fp16",
23+ "shape": [
24+ 128
25+ ],
26+ "range_values": 1,
27+ "backward": false,
28+ "align_32B": null
29+ },
30+ {
31+ "name": "self",
32+ "type": "tensor",
33+ "required": true,
34+ "dtype": "fp16",
35+ "shape": [
36+ 128
37+ ],
38+ "range_values": [
39+ -10,
40+ 10
41+ ],
42+ "backward": false,
43+ "align_32B": null
44+ },
45+ {
46+ "name": "threshold",
47+ "type": "scalar",
48+ "required": true,
49+ "dtype": "int32",
50+ "shape": null,
51+ "range_values": 0,
52+ "backward": false,
53+ "align_32B": null
54+ }
55+ ],
56+ "acl_json": "",
57+ "method_inputs": null,
58+ "tensor_input": null
59+ },
60+ {
61+ "id": 2,
62+ "name": "torch.threshold_backward bf16",
63+ "aclnn_name": "aclnnThresholdBackward",
64+ "version": "v1.0",
65+ "expected_error_msg": null,
66+ "api": "pytorch",
67+ "api_type": "aclnn_threshold_backward",
68+ "aclnn_api_type": "aclnn_function",
69+ "backward": false,
70+ "standard": {
71+ "acc": "cv_fused_double_benchmark",
72+ "perf": "not_key"
73+ },
74+ "outputs": null,
75+ "inputs": [
76+ {
77+ "name": "gradOutput",
78+ "type": "tensor",
79+ "required": true,
80+ "dtype": "bf16",
81+ "shape": [
82+ 32,
83+ 256
84+ ],
85+ "range_values": 1,
86+ "backward": false,
87+ "align_32B": null
88+ },
89+ {
90+ "name": "self",
91+ "type": "tensor",
92+ "required": true,
93+ "dtype": "bf16",
94+ "shape": [
95+ 32,
96+ 256
97+ ],
98+ "range_values": [
99+ -10,
100+ 10
101+ ],
102+ "backward": false,
103+ "align_32B": null
104+ },
105+ {
106+ "name": "threshold",
107+ "type": "scalar",
108+ "required": true,
109+ "dtype": "int32",
110+ "shape": null,
111+ "range_values": 0,
112+ "backward": false,
113+ "align_32B": null
114+ }
115+ ],
116+ "acl_json": "",
117+ "method_inputs": null,
118+ "tensor_input": null
119+ },
120+ {
121+ "id": 3,
122+ "name": "torch.threshold_backward fp32 scalar",
123+ "aclnn_name": "aclnnThresholdBackward",
124+ "version": "v1.0",
125+ "expected_error_msg": null,
126+ "api": "pytorch",
127+ "api_type": "aclnn_threshold_backward",
128+ "aclnn_api_type": "aclnn_function",
129+ "backward": false,
130+ "standard": {
131+ "acc": "cv_fused_double_benchmark",
132+ "perf": "not_key"
133+ },
134+ "outputs": null,
135+ "inputs": [
136+ {
137+ "name": "gradOutput",
138+ "type": "tensor",
139+ "required": true,
140+ "dtype": "fp32",
141+ "shape": [],
142+ "range_values": 1,
143+ "backward": false,
144+ "align_32B": null
145+ },
146+ {
147+ "name": "self",
148+ "type": "tensor",
149+ "required": true,
150+ "dtype": "fp32",
151+ "shape": [],
152+ "range_values": [
153+ 1.5
154+ ],
155+ "backward": false,
156+ "align_32B": null
157+ },
158+ {
159+ "name": "threshold",
160+ "type": "scalar",
161+ "required": true,
162+ "dtype": "int32",
163+ "shape": null,
164+ "range_values": 0,
165+ "backward": false,
166+ "align_32B": null
167+ }
168+ ],
169+ "acl_json": "",
170+ "method_inputs": null,
171+ "tensor_input": null
172+ },
173+ {
174+ "id": 4,
175+ "name": "torch.threshold_backward int8",
176+ "aclnn_name": "aclnnThresholdBackward",
177+ "version": "v1.0",
178+ "expected_error_msg": null,
179+ "api": "pytorch",
180+ "api_type": "aclnn_threshold_backward",
181+ "aclnn_api_type": "aclnn_function",
182+ "backward": false,
183+ "standard": {
184+ "acc": "cv_fused_double_benchmark",
185+ "perf": "not_key"
186+ },
187+ "outputs": null,
188+ "inputs": [
189+ {
190+ "name": "gradOutput",
191+ "type": "tensor",
192+ "required": true,
193+ "dtype": "int8",
194+ "shape": [
195+ 1024
196+ ],
197+ "range_values": 1,
198+ "backward": false,
199+ "align_32B": null
200+ },
201+ {
202+ "name": "self",
203+ "type": "tensor",
204+ "required": true,
205+ "dtype": "int8",
206+ "shape": [
207+ 1024
208+ ],
209+ "range_values": [
210+ -20,
211+ 20
212+ ],
213+ "backward": false,
214+ "align_32B": null
215+ },
216+ {
217+ "name": "threshold",
218+ "type": "scalar",
219+ "required": true,
220+ "dtype": "int32",
221+ "shape": null,
222+ "range_values": 0,
223+ "backward": false,
224+ "align_32B": null
225+ }
226+ ],
227+ "acl_json": "",
228+ "method_inputs": null,
229+ "tensor_input": null
230+ },
231+ {
232+ "id": 5,
233+ "name": "torch.threshold_backward uint8",
234+ "aclnn_name": "aclnnThresholdBackward",
235+ "version": "v1.0",
236+ "expected_error_msg": null,
237+ "api": "pytorch",
238+ "api_type": "aclnn_threshold_backward",
239+ "aclnn_api_type": "aclnn_function",
240+ "backward": false,
241+ "standard": {
242+ "acc": "cv_fused_double_benchmark",
243+ "perf": "not_key"
244+ },
245+ "outputs": null,
246+ "inputs": [
247+ {
248+ "name": "gradOutput",
249+ "type": "tensor",
250+ "required": true,
251+ "dtype": "uint8",
252+ "shape": [
253+ 257
254+ ],
255+ "range_values": 1,
256+ "backward": false,
257+ "align_32B": null
258+ },
259+ {
260+ "name": "self",
261+ "type": "tensor",
262+ "required": true,
263+ "dtype": "uint8",
264+ "shape": [
265+ 257
266+ ],
267+ "range_values": [
268+ 0,
269+ 20
270+ ],
271+ "backward": false,
272+ "align_32B": null
273+ },
274+ {
275+ "name": "threshold",
276+ "type": "scalar",
277+ "required": true,
278+ "dtype": "int32",
279+ "shape": null,
280+ "range_values": 0,
281+ "backward": false,
282+ "align_32B": null
283+ }
284+ ],
285+ "acl_json": "",
286+ "method_inputs": null,
287+ "tensor_input": null
288+ },
289+ {
290+ "id": 6,
291+ "name": "torch.threshold_backward int32",
292+ "aclnn_name": "aclnnThresholdBackward",
293+ "version": "v1.0",
294+ "expected_error_msg": null,
295+ "api": "pytorch",
296+ "api_type": "aclnn_threshold_backward",
297+ "aclnn_api_type": "aclnn_function",
298+ "backward": false,
299+ "standard": {
300+ "acc": "cv_fused_double_benchmark",
301+ "perf": "not_key"
302+ },
303+ "outputs": null,
304+ "inputs": [
305+ {
306+ "name": "gradOutput",
307+ "type": "tensor",
308+ "required": true,
309+ "dtype": "int32",
310+ "shape": [
311+ 513
312+ ],
313+ "range_values": 1,
314+ "backward": false,
315+ "align_32B": null
316+ },
317+ {
318+ "name": "self",
319+ "type": "tensor",
320+ "required": true,
321+ "dtype": "int32",
322+ "shape": [
323+ 513
324+ ],
325+ "range_values": [
326+ -20,
327+ 20
328+ ],
329+ "backward": false,
330+ "align_32B": null
331+ },
332+ {
333+ "name": "threshold",
334+ "type": "scalar",
335+ "required": true,
336+ "dtype": "int32",
337+ "shape": null,
338+ "range_values": 0,
339+ "backward": false,
340+ "align_32B": null
341+ }
342+ ],
343+ "acl_json": "",
344+ "method_inputs": null,
345+ "tensor_input": null
346+ },
347+ {
348+ "id": 7,
349+ "name": "torch.threshold_backward int64",
350+ "aclnn_name": "aclnnThresholdBackward",
351+ "version": "v1.0",
352+ "expected_error_msg": null,
353+ "api": "pytorch",
354+ "api_type": "aclnn_threshold_backward",
355+ "aclnn_api_type": "aclnn_function",
356+ "backward": false,
357+ "standard": {
358+ "acc": "cv_fused_double_benchmark",
359+ "perf": "not_key"
360+ },
361+ "outputs": null,
362+ "inputs": [
363+ {
364+ "name": "gradOutput",
365+ "type": "tensor",
366+ "required": true,
367+ "dtype": "int64",
368+ "shape": [
369+ 257
370+ ],
371+ "range_values": 1,
372+ "backward": false,
373+ "align_32B": null
374+ },
375+ {
376+ "name": "self",
377+ "type": "tensor",
378+ "required": true,
379+ "dtype": "int64",
380+ "shape": [
381+ 257
382+ ],
383+ "range_values": [
384+ -20,
385+ 20
386+ ],
387+ "backward": false,
388+ "align_32B": null
389+ },
390+ {
391+ "name": "threshold",
392+ "type": "scalar",
393+ "required": true,
394+ "dtype": "int32",
395+ "shape": null,
396+ "range_values": 0,
397+ "backward": false,
398+ "align_32B": null
399+ }
400+ ],
401+ "acl_json": "",
402+ "method_inputs": null,
403+ "tensor_input": null
404+ }
405+]
Aexperimental/activation/relu_grad_v2/tests/st/aclnnThresholdBackward/executor_aclnnThresholdBackward.py+34-0
@@ -0,0 +1,34 @@
1+#!/usr/bin/env python3
2+# -*- coding: utf-8 -*-
3+# ----------------------------------------------------------------------------
4+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
5+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
6+# CANN Open Software License Agreement Version 2.0 (the "License").
7+# Please refer to the License for details. You may not use this file except in compliance with the License.
8+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
9+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
10+# See LICENSE in the root of the software repository for the full text of the License.
11+#
12+ 
13+ 
14+import torch
15+from atk.configs.dataset_config import InputDataset
16+from atk.tasks.api_execute import register
17+from atk.tasks.api_execute.base_api import BaseApi
18+ 
19+ 
20+def threshold_backward_reference(grad_output: torch.Tensor, self_tensor: torch.Tensor) -> torch.Tensor:
21+ if self_tensor.is_floating_point():
22+ return torch.ops.aten.threshold_backward.default(grad_output, self_tensor, 0)
23+ return torch.where(self_tensor > 0, grad_output, torch.zeros_like(grad_output))
24+ 
25+ 
26+@register("aclnn_threshold_backward")
27+class TorchThresholdBackward(BaseApi):
28+ def __call__(self, input_data: InputDataset, with_output: bool = False):
29+ grad_output = input_data.kwargs["gradOutput"]
30+ self_tensor = input_data.kwargs["self"]
31+ threshold = input_data.kwargs["threshold"]
32+ if int(threshold) != 0:
33+ raise ValueError(f"aclnnThresholdBackward expects threshold == 0, got {threshold}")
34+ return threshold_backward_reference(grad_output, self_tensor)
Aexperimental/activation/relu_grad_v2/tests/ut/CMakeLists.txt+16-0
@@ -0,0 +1,16 @@
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_SOURCE_DIRS LIST_DIRECTORIES true ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+foreach(SUB_DIR ${CURRENT_SOURCE_DIRS})
13+ if(EXISTS "${SUB_DIR}/CMakeLists.txt")
14+ add_subdirectory(${SUB_DIR})
15+ endif()
16+endforeach()
Aexperimental/activation/relu_grad_v2/tests/ut/op_api/CMakeLists.txt+14-0
@@ -0,0 +1,14 @@
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_DIR RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+if(UT_TEST_ALL OR OP_API_UT)
13+ add_modules_ut_sources(HOSTNAME ${OP_API_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
14+endif()
Aexperimental/activation/relu_grad_v2/tests/ut/op_api/test_aclnn_threshold_backward.cpp+288-0
@@ -0,0 +1,288 @@
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+ 
13+#include "../../../op_host/op_api/aclnn_threshold_backward.h"
14+ 
15+#include "op_api_ut_common/op_api_ut.h"
16+#include "op_api_ut_common/tensor_desc.h"
17+ 
18+using namespace op;
19+ 
20+class l2_threshold_backward_test : public testing::Test {
21+ protected:
22+ static void SetUpTestCase()
23+ {
24+ std::cout << "threshold_backward_test SetUp" << std::endl;
25+ }
26+ 
27+ static void TearDownTestCase()
28+ {
29+ std::cout << "threshold_backward_test TearDown" << std::endl;
30+ }
31+};
32+ 
33+TEST_F(l2_threshold_backward_test, case_001_float)
34+{
35+ auto gradOutputDesc = TensorDesc({2, 4}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-10, 10);
36+ auto selfDesc = TensorDesc({2, 4}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-10, 10);
37+ auto thresholdDesc = ScalarDesc(0);
38+ auto outDesc = TensorDesc({2, 4}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.0001, 0.0001);
39+ 
40+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, thresholdDesc), OUTPUT(outDesc));
41+ 
42+ uint64_t workspaceSize = 0;
43+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
44+ EXPECT_EQ(aclRet, ACL_SUCCESS);
45+ 
46+ ut.TestPrecision();
47+}
48+ 
49+TEST_F(l2_threshold_backward_test, case_002_float16)
50+{
51+ auto gradOutputDesc = TensorDesc({2, 4}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-10, 10);
52+ auto selfDesc = TensorDesc({2, 4}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-10, 10);
53+ auto thresholdDesc = ScalarDesc(0);
54+ auto outDesc = TensorDesc({2, 4}, ACL_FLOAT16, ACL_FORMAT_ND).Precision(0.001, 0.001);
55+ 
56+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, thresholdDesc), OUTPUT(outDesc));
57+ 
58+ uint64_t workspaceSize = 0;
59+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
60+ EXPECT_EQ(aclRet, ACL_SUCCESS);
61+ 
62+ ut.TestPrecision();
63+}
64+ 
65+TEST_F(l2_threshold_backward_test, case_003_bfloat16)
66+{
67+ auto gradOutputDesc = TensorDesc({2, 4}, ACL_BF16, ACL_FORMAT_ND).ValueRange(-10, 10);
68+ auto selfDesc = TensorDesc({2, 4}, ACL_BF16, ACL_FORMAT_ND).ValueRange(-10, 10);
69+ auto thresholdDesc = ScalarDesc(0);
70+ auto outDesc = TensorDesc({2, 4}, ACL_BF16, ACL_FORMAT_ND).Precision(0.01, 0.01);
71+ 
72+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, thresholdDesc), OUTPUT(outDesc));
73+ 
74+ uint64_t workspaceSize = 0;
75+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
76+ EXPECT_EQ(aclRet, ACL_SUCCESS);
77+ 
78+ ut.TestPrecision();
79+}
80+ 
81+TEST_F(l2_threshold_backward_test, case_004_int8)
82+{
83+ auto gradOutputDesc = TensorDesc({1024}, ACL_INT8, ACL_FORMAT_ND).ValueRange(-20, 20);
84+ auto selfDesc = TensorDesc({1024}, ACL_INT8, ACL_FORMAT_ND).ValueRange(-20, 20);
85+ auto thresholdDesc = ScalarDesc(0);
86+ auto outDesc = TensorDesc({1024}, ACL_INT8, ACL_FORMAT_ND);
87+ 
88+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, thresholdDesc), OUTPUT(outDesc));
89+ 
90+ uint64_t workspaceSize = 0;
91+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
92+ EXPECT_EQ(aclRet, ACL_SUCCESS);
93+ 
94+ ut.TestPrecision();
95+}
96+ 
97+TEST_F(l2_threshold_backward_test, case_005_uint8)
98+{
99+ auto gradOutputDesc = TensorDesc({257}, ACL_UINT8, ACL_FORMAT_ND).ValueRange(0, 20);
100+ auto selfDesc = TensorDesc({257}, ACL_UINT8, ACL_FORMAT_ND).ValueRange(0, 20);
101+ auto thresholdDesc = ScalarDesc(0);
102+ auto outDesc = TensorDesc({257}, ACL_UINT8, ACL_FORMAT_ND);
103+ 
104+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, thresholdDesc), OUTPUT(outDesc));
105+ 
106+ uint64_t workspaceSize = 0;
107+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
108+ EXPECT_EQ(aclRet, ACL_SUCCESS);
109+ 
110+ ut.TestPrecision();
111+}
112+ 
113+TEST_F(l2_threshold_backward_test, case_006_int32)
114+{
115+ auto gradOutputDesc = TensorDesc({512}, ACL_INT32, ACL_FORMAT_ND).ValueRange(-20, 20);
116+ auto selfDesc = TensorDesc({512}, ACL_INT32, ACL_FORMAT_ND).ValueRange(-20, 20);
117+ auto thresholdDesc = ScalarDesc(0);
118+ auto outDesc = TensorDesc({512}, ACL_INT32, ACL_FORMAT_ND);
119+ 
120+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, thresholdDesc), OUTPUT(outDesc));
121+ 
122+ uint64_t workspaceSize = 0;
123+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
124+ EXPECT_EQ(aclRet, ACL_SUCCESS);
125+ 
126+ ut.TestPrecision();
127+}
128+ 
129+TEST_F(l2_threshold_backward_test, case_007_int64)
130+{
131+ auto gradOutputDesc = TensorDesc({257}, ACL_INT64, ACL_FORMAT_ND).ValueRange(-20, 20);
132+ auto selfDesc = TensorDesc({257}, ACL_INT64, ACL_FORMAT_ND).ValueRange(-20, 20);
133+ auto thresholdDesc = ScalarDesc(0);
134+ auto outDesc = TensorDesc({257}, ACL_INT64, ACL_FORMAT_ND);
135+ 
136+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, thresholdDesc), OUTPUT(outDesc));
137+ 
138+ uint64_t workspaceSize = 0;
139+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
140+ EXPECT_EQ(aclRet, ACL_SUCCESS);
141+ 
142+ ut.TestPrecision();
143+}
144+ 
145+TEST_F(l2_threshold_backward_test, case_008_empty_tensor)
146+{
147+ auto gradOutputDesc = TensorDesc({2, 0}, ACL_FLOAT, ACL_FORMAT_ND);
148+ auto selfDesc = TensorDesc({2, 0}, ACL_FLOAT, ACL_FORMAT_ND);
149+ auto thresholdDesc = ScalarDesc(0);
150+ auto outDesc = TensorDesc({2, 0}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.0001, 0.0001);
151+ 
152+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, thresholdDesc), OUTPUT(outDesc));
153+ 
154+ uint64_t workspaceSize = 0;
155+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
156+ EXPECT_EQ(aclRet, ACL_SUCCESS);
157+ 
158+ ut.TestPrecision();
159+}
160+ 
161+TEST_F(l2_threshold_backward_test, case_009_not_contiguous)
162+{
163+ auto gradOutputDesc = TensorDesc({5, 4}, ACL_FLOAT, ACL_FORMAT_ND, {1, 5}, 0, {4, 5}).ValueRange(-10, 10);
164+ auto selfDesc = TensorDesc({5, 4}, ACL_FLOAT, ACL_FORMAT_ND, {1, 5}, 0, {4, 5}).ValueRange(-10, 10);
165+ auto thresholdDesc = ScalarDesc(0);
166+ auto outDesc = TensorDesc({5, 4}, ACL_FLOAT, ACL_FORMAT_ND, {1, 5}, 0, {4, 5}).Precision(0.0001, 0.0001);
167+ 
168+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, thresholdDesc), OUTPUT(outDesc));
169+ 
170+ uint64_t workspaceSize = 0;
171+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
172+ EXPECT_EQ(aclRet, ACL_SUCCESS);
173+ 
174+ ut.TestPrecision();
175+}
176+ 
177+TEST_F(l2_threshold_backward_test, case_010_invalid_input_dtype)
178+{
179+ auto gradOutputDesc = TensorDesc({2, 3}, ACL_BOOL, ACL_FORMAT_ND);
180+ auto selfDesc = TensorDesc({2, 3}, ACL_FLOAT, ACL_FORMAT_ND);
181+ auto thresholdDesc = ScalarDesc(0);
182+ auto outDesc = TensorDesc({2, 3}, ACL_FLOAT, ACL_FORMAT_ND);
183+ 
184+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, thresholdDesc), OUTPUT(outDesc));
185+ 
186+ uint64_t workspaceSize = 0;
187+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
188+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
189+}
190+ 
191+TEST_F(l2_threshold_backward_test, case_011_invalid_output_dtype)
192+{
193+ auto gradOutputDesc = TensorDesc({2, 3}, ACL_FLOAT, ACL_FORMAT_ND);
194+ auto selfDesc = TensorDesc({2, 3}, ACL_FLOAT, ACL_FORMAT_ND);
195+ auto thresholdDesc = ScalarDesc(0);
196+ auto outDesc = TensorDesc({2, 3}, ACL_BOOL, ACL_FORMAT_ND);
197+ 
198+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, thresholdDesc), OUTPUT(outDesc));
199+ 
200+ uint64_t workspaceSize = 0;
201+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
202+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
203+}
204+ 
205+TEST_F(l2_threshold_backward_test, case_012_dtype_mismatch)
206+{
207+ auto gradOutputDesc = TensorDesc({2, 3}, ACL_FLOAT, ACL_FORMAT_ND);
208+ auto selfDesc = TensorDesc({2, 3}, ACL_FLOAT16, ACL_FORMAT_ND);
209+ auto thresholdDesc = ScalarDesc(0);
210+ auto outDesc = TensorDesc({2, 3}, ACL_FLOAT16, ACL_FORMAT_ND);
211+ 
212+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, thresholdDesc), OUTPUT(outDesc));
213+ 
214+ uint64_t workspaceSize = 0;
215+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
216+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
217+}
218+ 
219+TEST_F(l2_threshold_backward_test, case_013_nullptr)
220+{
221+ auto selfDesc = TensorDesc({2, 4}, ACL_FLOAT, ACL_FORMAT_ND);
222+ auto thresholdDesc = ScalarDesc(0);
223+ auto outDesc = TensorDesc({2, 4}, ACL_FLOAT, ACL_FORMAT_ND);
224+ 
225+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(nullptr, selfDesc, thresholdDesc), OUTPUT(outDesc));
226+ 
227+ uint64_t workspaceSize = 0;
228+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
229+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_NULLPTR);
230+}
231+ 
232+TEST_F(l2_threshold_backward_test, case_014_shape_mismatch)
233+{
234+ auto gradOutputDesc = TensorDesc({1, 2, 3, 4}, ACL_INT8, ACL_FORMAT_ND);
235+ auto selfDesc = TensorDesc({1, 2, 3, 3}, ACL_INT8, ACL_FORMAT_ND);
236+ auto thresholdDesc = ScalarDesc(0);
237+ auto outDesc = TensorDesc({1, 2, 3, 3}, ACL_INT8, ACL_FORMAT_ND);
238+ 
239+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, thresholdDesc), OUTPUT(outDesc));
240+ 
241+ uint64_t workspaceSize = 0;
242+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
243+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
244+}
245+ 
246+TEST_F(l2_threshold_backward_test, case_015_max_dim)
247+{
248+ auto gradOutputDesc = TensorDesc({1, 2, 3, 4, 5, 6, 7, 8, 9}, ACL_INT8, ACL_FORMAT_ND);
249+ auto selfDesc = TensorDesc({1, 2, 3, 4, 5, 6, 7, 8, 9}, ACL_INT8, ACL_FORMAT_ND);
250+ auto thresholdDesc = ScalarDesc(0);
251+ auto outDesc = TensorDesc({1, 2, 3, 4, 5, 6, 7, 8, 9}, ACL_INT8, ACL_FORMAT_ND);
252+ 
253+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, thresholdDesc), OUTPUT(outDesc));
254+ 
255+ uint64_t workspaceSize = 0;
256+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
257+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
258+}
259+ 
260+TEST_F(l2_threshold_backward_test, case_016_invalid_threshold)
261+{
262+ auto gradOutputDesc = TensorDesc({2, 4}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-10, 10);
263+ auto selfDesc = TensorDesc({2, 4}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-10, 10);
264+ auto thresholdDesc = ScalarDesc(1);
265+ auto outDesc = TensorDesc({2, 4}, ACL_FLOAT, ACL_FORMAT_ND);
266+ 
267+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, thresholdDesc), OUTPUT(outDesc));
268+ 
269+ uint64_t workspaceSize = 0;
270+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
271+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
272+}
273+ 
274+TEST_F(l2_threshold_backward_test, case_017_scalar_float)
275+{
276+ auto gradOutputDesc = TensorDesc({}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-10, 10);
277+ auto selfDesc = TensorDesc({}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-10, 10);
278+ auto thresholdDesc = ScalarDesc(0);
279+ auto outDesc = TensorDesc({}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.0001, 0.0001);
280+ 
281+ auto ut = OP_API_UT(aclnnThresholdBackward, INPUT(gradOutputDesc, selfDesc, thresholdDesc), OUTPUT(outDesc));
282+ 
283+ uint64_t workspaceSize = 0;
284+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
285+ EXPECT_EQ(aclRet, ACL_SUCCESS);
286+ 
287+ ut.TestPrecision();
288+}