Star1012
1937
代码
代码
Issues194
Pull Requests332
Actions
项目讨论
Wiki
分析
Star1012
1937
已合并
新增 ReLU 实验算子 #4478
cwzhang创建于 4月30日
新增 ReLU 实验算子 #4478
已合并
cwzhang创建于 4月30日
从cwzhang/ops-nnrelu合入到cann/ops-nnmaster
共 24 个文件变更+2394-0
Aexperimental/activation/relu/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/README.md+176-0
@@ -0,0 +1,176 @@
1+# Relu
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+- 算子功能：对输入 Tensor 完成 ReLU 运算。
17+ 
18+计算公式：
19+ 
20+$$
21+\operatorname{relu}(x) = \max(x, 0)
22+$$
23+ 
24+- 目录 `experimental/activation/relu` 对外导出 `aclnnRelu` 和 `aclnnInplaceRelu` 两段式 ACLNN 接口。
25+- `op_host/op_api/aclnn_relu.cpp` 是对外 ACLNN 接口入口。
26+- `op_host/op_api/relu.h` 和 `op_host/op_api/relu.cpp` 提供内部 `l0op::Relu` 封装，当前由 ACLNN 接口直接调用。
27+ 
28+## 调用方式
29+ 
30+| 调用方式 | 是否支持 |
31+| :------- | :------: |
32+| ACLNN 调用 | 是 |
33+| ACLNN 原地调用 | 是 |
34+ 
35+## ACLNN 接口
36+ 
37+### 函数原型
38+ 
39+当前 experimental Relu 提供两段式 ACLNN 接口：
40+ 
41+```cpp
42+aclnnStatus aclnnReluGetWorkspaceSize(
43+ const aclTensor *self,
44+ const aclTensor *out,
45+ uint64_t *workspaceSize,
46+ aclOpExecutor **executor);
47+ 
48+aclnnStatus aclnnRelu(
49+ void *workspace,
50+ uint64_t workspaceSize,
51+ aclOpExecutor *executor,
52+ const aclrtStream stream);
53+```
54+ 
55+原地版本接口如下：
56+ 
57+```cpp
58+aclnnStatus aclnnInplaceReluGetWorkspaceSize(
59+ aclTensor *selfRef,
60+ uint64_t *workspaceSize,
61+ aclOpExecutor **executor);
62+ 
63+aclnnStatus aclnnInplaceRelu(
64+ void *workspace,
65+ uint64_t workspaceSize,
66+ aclOpExecutor *executor,
67+ const aclrtStream stream);
68+```
69+ 
70+详细参数说明见 [docs/aclnnRelu&aclnnInplaceRelu.md](docs/aclnnRelu&aclnnInplaceRelu.md)。
71+ 
72+## 参数说明
73+ 
74+<table style="undefined;table-layout: fixed; width: 1393px"><colgroup>
75+<col style="width: 171px">
76+<col style="width: 115px">
77+<col style="width: 260px">
78+<col style="width: 220px">
79+<col style="width: 200px">
80+<col style="width: 104px">
81+</colgroup>
82+<thead>
83+ <tr>
84+ <th>参数名</th>
85+ <th>输入/输出</th>
86+ <th>描述</th>
87+ <th>数据类型</th>
88+ <th>数据格式</th>
89+ </tr>
90+</thead>
91+<tbody>
92+ <tr>
93+ <td>self</td>
94+ <td>输入</td>
95+ <td>待进行 ReLU 计算的入参。</td>
96+ <td>FLOAT、FLOAT16、BFLOAT16、INT8、INT32、INT64</td>
97+ <td>ND</td>
98+ </tr>
99+ <tr>
100+ <td>out</td>
101+ <td>输出</td>
102+ <td>非原地接口的计算出参。</td>
103+ <td>FLOAT、FLOAT16、BFLOAT16、INT8、INT32、INT64</td>
104+ <td>ND</td>
105+ </tr>
106+ <tr>
107+ <td>selfRef</td>
108+ <td>输入/输出</td>
109+ <td>原地接口的输入输出张量。</td>
110+ <td>FLOAT、FLOAT16、BFLOAT16、INT8、INT32、INT64</td>
111+ <td>ND</td>
112+ </tr>
113+</tbody>
114+</table>
115+ 
116+## 约束说明
117+ 
118+- 输入仅支持 `FLOAT`、`FLOAT16`、`BFLOAT16`、`INT8`、`INT32`、`INT64`。
119+- `BFLOAT16` 仅在 `Ascend910B` 及后续同代 SoC 上支持。
120+- 输入和输出的 dtype 必须一致。
121+- 输入和输出的 shape 必须完全一致。
122+- 支持 0 到 8 维 Tensor。
123+- 支持空 Tensor。
124+- 支持非连续 Tensor，接口内部会在需要时做 `Contiguous` 和 `ViewCopy`。
125+- `FLOAT`、`FLOAT16`、`INT32` 路径直接使用 AscendC `Relu` 计算。
126+- `BFLOAT16` 路径在 kernel 中升精度到 `float32` 计算后回写。
127+- `INT8` 路径在 kernel 中升精度到 `float16` 计算后回写。
128+- `INT64` 路径按已验证基线走逐元素标量语义。
129+ 
130+## 目录说明
131+ 
132+| 路径 | 说明 |
133+| :--- | :--- |
134+| [examples/test_aclnn_relu.cpp](examples/test_aclnn_relu.cpp) | `aclnnRelu` 两段式调用示例。 |
135+| [examples/test_aclnn_inplace_relu.cpp](examples/test_aclnn_inplace_relu.cpp) | `aclnnInplaceRelu` 两段式调用示例。 |
136+| [examples/run.sh](examples/run.sh) | 编译并运行 example 的脚本。 |
137+| [docs/aclnnRelu&aclnnInplaceRelu.md](docs/aclnnRelu&aclnnInplaceRelu.md) | `aclnnRelu` / `aclnnInplaceRelu` 接口文档。 |
138+| [tests/ut/op_api/test_aclnn_relu.cpp](tests/ut/op_api/test_aclnn_relu.cpp) | `op_api` 单元测试。 |
139+| [tests/st/aclnnRelu/all_aclnnRelu.json](tests/st/aclnnRelu/all_aclnnRelu.json) | 适用于 ATK 的小规模标准化测试集。 |
140+| [tests/st/aclnnRelu/executor_aclnnRelu.py](tests/st/aclnnRelu/executor_aclnnRelu.py) | ATK CPU benchmark 执行器。 |
141+ 
142+## Example 运行
143+ 
144+先确保 custom run 包已经安装，并加载 CANN 环境：
145+ 
146+```bash
147+source /usr/local/Ascend/cann-8.5.0-beta.1/set_env.sh
148+export LD_LIBRARY_PATH=/usr/local/Ascend/cann-8.5.0-beta.1/opp/vendors/customize_nn/op_api/lib:${LD_LIBRARY_PATH}
149+cd /root/src/cann/ops-nn/experimental/activation/relu/examples
150+bash run.sh
151+```
152+ 
153+示例会编译并运行 `test_aclnn_relu.cpp` 与 `test_aclnn_inplace_relu.cpp`。
154+ 
155+## Tests 运行
156+ 
157+### 1. op_api 单元测试
158+ 
159+```bash
160+source /usr/local/Ascend/cann/set_env.sh
161+cd /root/src/cann/ops-nn
162+bash build.sh --experimental --ops=relu -u --opapi -j8 -O2
163+```
164+ 
165+### 2. ATK 小规模标准化测试
166+ 
167+```bash
168+export ATK_BIND_CPU_TYPE=2
169+source /usr/local/Ascend/cann/set_env.sh
170+source /root/src/kernel/ascend-kernel/.venv/bin/activate
171+cd /root/src/testcase
172+atk node --backend npu --devices 2 \
173+ node --backend cpu task --task accuracy \
174+ -c /root/src/cann/ops-nn/experimental/activation/relu/tests/st/aclnnRelu/all_aclnnRelu.json \
175+ -p /root/src/cann/ops-nn/experimental/activation/relu/tests/st/aclnnRelu/executor_aclnnRelu.py
176+```
Aexperimental/activation/relu/docs/aclnnRelu&aclnnInplaceRelu.md+370-0
@@ -0,0 +1,370 @@
1+# aclnnRelu 与 aclnnInplaceRelu
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+- `aclnnRelu`：对输入 Tensor 执行 ReLU 计算，并将结果写入独立输出 Tensor。
17+- `aclnnInplaceRelu`：对输入 Tensor 原地执行 ReLU 计算，结果直接覆盖输入内存。
18+- `experimental/activation/relu` 目录对外导出的 ACLNN 接口名与原实现保持一致。
19+ 
20+计算公式：
21+ 
22+$$
23+\operatorname{relu}(x) = \max(x, 0)
24+$$
25+ 
26+## 函数原型
27+ 
28+每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md)，必须先调用 `aclnnReluGetWorkspaceSize` 或 `aclnnInplaceReluGetWorkspaceSize` 获取执行器和 workspace 大小，再调用第二段接口执行计算。
29+ 
30+```cpp
31+aclnnStatus aclnnReluGetWorkspaceSize(
32+ const aclTensor *self,
33+ const aclTensor *out,
34+ uint64_t *workspaceSize,
35+ aclOpExecutor **executor);
36+```
37+ 
38+```cpp
39+aclnnStatus aclnnRelu(
40+ void *workspace,
41+ uint64_t workspaceSize,
42+ aclOpExecutor *executor,
43+ const aclrtStream stream);
44+```
45+ 
46+```cpp
47+aclnnStatus aclnnInplaceReluGetWorkspaceSize(
48+ aclTensor *selfRef,
49+ uint64_t *workspaceSize,
50+ aclOpExecutor **executor);
51+```
52+ 
53+```cpp
54+aclnnStatus aclnnInplaceRelu(
55+ void *workspace,
56+ uint64_t workspaceSize,
57+ aclOpExecutor *executor,
58+ const aclrtStream stream);
59+```
60+ 
61+## aclnnReluGetWorkspaceSize
62+ 
63+- **参数说明：**
64+ 
65+ <table style="undefined;table-layout: fixed; width: 1497px"><colgroup>
66+ <col style="width: 271px">
67+ <col style="width: 115px">
68+ <col style="width: 247px">
69+ <col style="width: 300px">
70+ <col style="width: 177px">
71+ <col style="width: 104px">
72+ <col style="width: 138px">
73+ <col style="width: 145px">
74+ </colgroup>
75+ <thead>
76+ <tr>
77+ <th>参数名</th>
78+ <th>输入/输出</th>
79+ <th>描述</th>
80+ <th>使用说明</th>
81+ <th>数据类型</th>
82+ <th>数据格式</th>
83+ <th>维度(shape)</th>
84+ <th>非连续Tensor</th>
85+ </tr>
86+ </thead>
87+ <tbody>
88+ <tr>
89+ <td>self（aclTensor*）</td>
90+ <td>输入</td>
91+ <td>待进行 ReLU 计算的输入张量。</td>
92+ <td><ul><li>支持空Tensor。</li><li>shape 必须与 out 完全一致。</li><li>数据类型必须与 out 完全一致。</li></ul></td>
93+ <td>Ascend910B 及同代 SoC：BFLOAT16、FLOAT16、FLOAT32、INT8、INT32、INT64；其他支持产品：FLOAT16、FLOAT32、INT8、INT32、INT64</td>
94+ <td>ND</td>
95+ <td>0-8</td>
96+ <td>√</td>
97+ </tr>
98+ <tr>
99+ <td>out（aclTensor*）</td>
100+ <td>输出</td>
101+ <td>计算的出参。</td>
102+ <td><ul><li>支持空Tensor。</li><li>shape 必须与 self 完全一致。</li><li>数据类型必须与 self 完全一致。</li></ul></td>
103+ <td>Ascend910B 及同代 SoC：BFLOAT16、FLOAT16、FLOAT32、INT8、INT32、INT64；其他支持产品：FLOAT16、FLOAT32、INT8、INT32、INT64</td>
104+ <td>ND</td>
105+ <td>0-8</td>
106+ <td>√</td>
107+ </tr>
108+ <tr>
109+ <td>workspaceSize（uint64_t*）</td>
110+ <td>输出</td>
111+ <td>返回需要在 Device 侧申请的 workspace 大小。</td>
112+ <td>-</td>
113+ <td>-</td>
114+ <td>-</td>
115+ <td>-</td>
116+ <td>-</td>
117+ </tr>
118+ <tr>
119+ <td>executor（aclOpExecutor**）</td>
120+ <td>输出</td>
121+ <td>返回 op 执行器，包含算子计算流程。</td>
122+ <td>-</td>
123+ <td>-</td>
124+ <td>-</td>
125+ <td>-</td>
126+ <td>-</td>
127+ </tr>
128+ </tbody>
129+ </table>
130+ 
131+- **返回值：**
132+ 
133+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
134+ 
135+ 第一段接口会完成入参校验，出现以下场景时报错：
136+ 
137+ <table style="undefined;table-layout: fixed;width: 979px"><colgroup>
138+ <col style="width: 272px">
139+ <col style="width: 103px">
140+ <col style="width: 604px">
141+ </colgroup>
142+ <thead>
143+ <tr>
144+ <th>返回码</th>
145+ <th>错误码</th>
146+ <th>描述</th>
147+ </tr>
148+ </thead>
149+ <tbody>
150+ <tr>
151+ <td>ACLNN_ERR_PARAM_NULLPTR</td>
152+ <td>161001</td>
153+ <td>传入的 self 或 out 是空指针。</td>
154+ </tr>
155+ <tr>
156+ <td rowspan="4">ACLNN_ERR_PARAM_INVALID</td>
157+ <td rowspan="4">161002</td>
158+ <td>self 或 out 的数据类型不在支持范围内。</td>
159+ </tr>
160+ <tr>
161+ <td>self 和 out 的数据类型不一致。</td>
162+ </tr>
163+ <tr>
164+ <td>self 和 out 的 shape 不一致。</td>
165+ </tr>
166+ <tr>
167+ <td>self 或 out 的维度大于 8。</td>
168+ </tr>
169+ </tbody></table>
170+ 
171+## aclnnRelu
172+ 
173+- **参数说明：**
174+ 
175+ <table style="undefined;table-layout: fixed; width: 953px"><colgroup>
176+ <col style="width: 173px">
177+ <col style="width: 112px">
178+ <col style="width: 668px">
179+ </colgroup>
180+ <thead>
181+ <tr>
182+ <th>参数名</th>
183+ <th>输入/输出</th>
184+ <th>描述</th>
185+ </tr>
186+ </thead>
187+ <tbody>
188+ <tr>
189+ <td>workspace</td>
190+ <td>输入</td>
191+ <td>在 Device 侧申请的 workspace 内存地址。</td>
192+ </tr>
193+ <tr>
194+ <td>workspaceSize</td>
195+ <td>输入</td>
196+ <td>在 Device 侧申请的 workspace 大小，由第一段接口 aclnnReluGetWorkspaceSize 获取。</td>
197+ </tr>
198+ <tr>
199+ <td>executor</td>
200+ <td>输入</td>
201+ <td>op 执行器，包含算子计算流程。</td>
202+ </tr>
203+ <tr>
204+ <td>stream</td>
205+ <td>输入</td>
206+ <td>指定执行任务的 Stream。</td>
207+ </tr>
208+ </tbody>
209+ </table>
210+ 
211+- **返回值：**
212+ 
213+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
214+ 
215+## aclnnInplaceReluGetWorkspaceSize
216+ 
217+- **参数说明：**
218+ 
219+ <table style="undefined;table-layout: fixed; width: 1355px"><colgroup>
220+ <col style="width: 271px">
221+ <col style="width: 115px">
222+ <col style="width: 297px">
223+ <col style="width: 108px">
224+ <col style="width: 177px">
225+ <col style="width: 104px">
226+ <col style="width: 138px">
227+ <col style="width: 145px">
228+ </colgroup>
229+ <thead>
230+ <tr>
231+ <th>参数名</th>
232+ <th>输入/输出</th>
233+ <th>描述</th>
234+ <th>使用说明</th>
235+ <th>数据类型</th>
236+ <th>数据格式</th>
237+ <th>维度(shape)</th>
238+ <th>非连续Tensor</th>
239+ </tr>
240+ </thead>
241+ <tbody>
242+ <tr>
243+ <td>selfRef（aclTensor*）</td>
244+ <td>输入/输出</td>
245+ <td>原地计算的输入输出张量。</td>
246+ <td><ul><li>支持空Tensor。</li><li>原地计算后 shape 和 dtype 不变。</li></ul></td>
247+ <td>BFLOAT16、FLOAT16、FLOAT32、INT8、INT32、INT64</td>
248+ <td>ND</td>
249+ <td>0-8</td>
250+ <td>√</td>
251+ </tr>
252+ <tr>
253+ <td>workspaceSize（uint64_t*）</td>
254+ <td>输出</td>
255+ <td>返回需要在 Device 侧申请的 workspace 大小。</td>
256+ <td>-</td>
257+ <td>-</td>
258+ <td>-</td>
259+ <td>-</td>
260+ <td>-</td>
261+ </tr>
262+ <tr>
263+ <td>executor（aclOpExecutor**）</td>
264+ <td>输出</td>
265+ <td>返回 op 执行器，包含算子计算流程。</td>
266+ <td>-</td>
267+ <td>-</td>
268+ <td>-</td>
269+ <td>-</td>
270+ <td>-</td>
271+ </tr>
272+ </tbody>
273+ </table>
274+ 
275+- **返回值：**
276+ 
277+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
278+ 
279+ 第一段接口会完成入参校验，出现以下场景时报错：
280+ 
281+ <table style="undefined;table-layout: fixed;width: 979px"><colgroup>
282+ <col style="width: 272px">
283+ <col style="width: 103px">
284+ <col style="width: 604px">
285+ </colgroup>
286+ <thead>
287+ <tr>
288+ <th>返回码</th>
289+ <th>错误码</th>
290+ <th>描述</th>
291+ </tr>
292+ </thead>
293+ <tbody>
294+ <tr>
295+ <td>ACLNN_ERR_PARAM_NULLPTR</td>
296+ <td>161001</td>
297+ <td>传入的 selfRef 是空指针。</td>
298+ </tr>
299+ <tr>
300+ <td rowspan="2">ACLNN_ERR_PARAM_INVALID</td>
301+ <td rowspan="2">161002</td>
302+ <td>selfRef 的数据类型不在支持范围内。</td>
303+ </tr>
304+ <tr>
305+ <td>selfRef 的维度大于 8。</td>
306+ </tr>
307+ </tbody></table>
308+ 
309+## aclnnInplaceRelu
310+ 
311+- **参数说明：**
312+ 
313+ <table style="undefined;table-layout: fixed; width: 953px"><colgroup>
314+ <col style="width: 173px">
315+ <col style="width: 112px">
316+ <col style="width: 668px">
317+ </colgroup>
318+ <thead>
319+ <tr>
320+ <th>参数名</th>
321+ <th>输入/输出</th>
322+ <th>描述</th>
323+ </tr>
324+ </thead>
325+ <tbody>
326+ <tr>
327+ <td>workspace</td>
328+ <td>输入</td>
329+ <td>在 Device 侧申请的 workspace 内存地址。</td>
330+ </tr>
331+ <tr>
332+ <td>workspaceSize</td>
333+ <td>输入</td>
334+ <td>在 Device 侧申请的 workspace 大小，由第一段接口 aclnnInplaceReluGetWorkspaceSize 获取。</td>
335+ </tr>
336+ <tr>
337+ <td>executor</td>
338+ <td>输入</td>
339+ <td>op 执行器，包含算子计算流程。</td>
340+ </tr>
341+ <tr>
342+ <td>stream</td>
343+ <td>输入</td>
344+ <td>指定执行任务的 Stream。</td>
345+ </tr>
346+ </tbody>
347+ </table>
348+ 
349+- **返回值：**
350+ 
351+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
352+ 
353+## 约束说明
354+ 
355+- 输入 dtype 仅支持 `FLOAT`、`FLOAT16`、`BFLOAT16`、`INT8`、`INT32`、`INT64`。
356+- `BFLOAT16` 仅在 `Ascend910B` 及后续同代 SoC 上支持。
357+- `aclnnRelu` 中，`self` 和 `out` 的 dtype 必须一致。
358+- `aclnnRelu` 中，`self` 和 `out` 的 shape 必须一致。
359+- 维度范围为 0 到 8。
360+- 支持空 Tensor。
361+- 支持非连续 Tensor，内部会执行 `Contiguous` 和 `ViewCopy`。
362+ 
363+## 实现说明
364+ 
365+- `op_host/op_api/aclnn_relu.cpp` 提供对外 ACLNN 两段式接口。
366+- `op_host/op_api/relu.h` 和 `op_host/op_api/relu.cpp` 提供内部 `l0op::Relu` 封装，当前由 `aclnn_relu.cpp` 直接调用。
367+ 
368+## 调用示例
369+ 
370+完整示例见 [../examples/test_aclnn_relu.cpp](../examples/test_aclnn_relu.cpp) 和 [../examples/test_aclnn_inplace_relu.cpp](../examples/test_aclnn_inplace_relu.cpp)。
Aexperimental/activation/relu/examples/run.sh+45-0
@@ -0,0 +1,45 @@
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-8.5.0-beta.1}
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.cpp" \
25+ -I"${CANN_ROOT}/aarch64-linux/include" \
26+ -I"${CANN_ROOT}/opp/vendors/customize_nn/op_api/include" \
27+ -L"${CANN_ROOT}/lib64" \
28+ -L"${CANN_ROOT}/opp/vendors/customize_nn/op_api/lib" \
29+ -Wl,-rpath,"${CANN_ROOT}/lib64:${CANN_ROOT}/opp/vendors/customize_nn/op_api/lib" \
30+ -lcust_opapi -lnnopbase -lascendcl \
31+ -o "${BUILD_DIR}/test_aclnn_relu"
32+ 
33+"${BUILD_DIR}/test_aclnn_relu"
34+ 
35+g++ -std=c++17 -O2 \
36+ "${SCRIPT_DIR}/test_aclnn_inplace_relu.cpp" \
37+ -I"${CANN_ROOT}/aarch64-linux/include" \
38+ -I"${CANN_ROOT}/opp/vendors/customize_nn/op_api/include" \
39+ -L"${CANN_ROOT}/lib64" \
40+ -L"${CANN_ROOT}/opp/vendors/customize_nn/op_api/lib" \
41+ -Wl,-rpath,"${CANN_ROOT}/lib64:${CANN_ROOT}/opp/vendors/customize_nn/op_api/lib" \
42+ -lcust_opapi -lnnopbase -lascendcl \
43+ -o "${BUILD_DIR}/test_aclnn_inplace_relu"
44+ 
45+"${BUILD_DIR}/test_aclnn_inplace_relu"
Aexperimental/activation/relu/examples/test_aclnn_inplace_relu.cpp+138-0
@@ -0,0 +1,138 @@
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 <cmath>
12+#include <cstdint>
13+#include <cstdio>
14+#include <cstdlib>
15+#include <cstring>
16+#include <vector>
17+ 
18+#include "acl/acl.h"
19+#include "aclnnop/aclnn_relu.h"
20+ 
21+#define CHECK_RET(cond, expr) \
22+ do { \
23+ if (!(cond)) { \
24+ expr; \
25+ } \
26+ } while (0)
27+ 
28+namespace {
29+int64_t GetShapeSize(const std::vector<int64_t> &shape)
30+{
31+ int64_t shape_size = 1;
32+ for (int64_t dim : shape) {
33+ shape_size *= dim;
34+ }
35+ return shape_size;
36+}
37+ 
38+std::vector<int64_t> MakeStrides(const std::vector<int64_t> &shape)
39+{
40+ if (shape.empty()) {
41+ return {};
42+ }
43+ std::vector<int64_t> strides(shape.size(), 1);
44+ for (int64_t i = static_cast<int64_t>(shape.size()) - 2; i >= 0; --i) {
45+ strides[static_cast<size_t>(i)] = shape[static_cast<size_t>(i + 1)] * strides[static_cast<size_t>(i + 1)];
46+ }
47+ return strides;
48+}
49+ 
50+aclError CreateAclTensor(const std::vector<int64_t> &shape, aclDataType dtype, void *device_addr, aclTensor **tensor)
51+{
52+ std::vector<int64_t> strides = MakeStrides(shape);
53+ const int64_t *shape_ptr = shape.empty() ? nullptr : shape.data();
54+ const int64_t *strides_ptr = strides.empty() ? nullptr : strides.data();
55+ *tensor = aclCreateTensor(
56+ shape_ptr, shape.size(), dtype, strides_ptr, 0, ACL_FORMAT_ND, shape_ptr, shape.size(), device_addr);
57+ return *tensor == nullptr ? ACL_ERROR_FAILURE : ACL_SUCCESS;
58+}
59+} // namespace
60+ 
61+int main()
62+{
63+ const int32_t device_id = 0;
64+ std::vector<int64_t> shape = {2, 4};
65+ std::vector<float> host_data = {-4.0f, -1.0f, 0.0f, 2.0f, 3.0f, -5.0f, 6.0f, -7.0f};
66+ const size_t bytes = host_data.size() * sizeof(float);
67+ std::vector<float> result(host_data.size(), 0.0f);
68+ const std::vector<float> expected = {0.0f, 0.0f, 0.0f, 2.0f, 3.0f, 0.0f, 6.0f, 0.0f};
69+ uint64_t workspace_size = 0;
70+ aclrtStream stream = nullptr;
71+ void *device_addr = nullptr;
72+ void *workspace = nullptr;
73+ aclTensor *self = nullptr;
74+ aclOpExecutor *executor = nullptr;
75+ bool acl_initialized = false;
76+ bool device_set = false;
77+ 
78+ auto ret = aclInit(nullptr);
79+ CHECK_RET(ret == ACL_SUCCESS, return ret);
80+ acl_initialized = true;
81+ ret = aclrtSetDevice(device_id);
82+ CHECK_RET(ret == ACL_SUCCESS, goto cleanup);
83+ device_set = true;
84+ ret = aclrtCreateStream(&stream);
85+ CHECK_RET(ret == ACL_SUCCESS, goto cleanup);
86+ 
87+ ret = aclrtMalloc(&device_addr, bytes, ACL_MEM_MALLOC_HUGE_FIRST);
88+ CHECK_RET(ret == ACL_SUCCESS, goto cleanup);
89+ ret = aclrtMemcpy(device_addr, bytes, host_data.data(), bytes, ACL_MEMCPY_HOST_TO_DEVICE);
90+ CHECK_RET(ret == ACL_SUCCESS, goto cleanup);
91+ ret = CreateAclTensor(shape, ACL_FLOAT, device_addr, &self);
92+ CHECK_RET(ret == ACL_SUCCESS, goto cleanup);
93+ 
94+ ret = aclnnInplaceReluGetWorkspaceSize(self, &workspace_size, &executor);
95+ CHECK_RET(ret == ACL_SUCCESS, goto cleanup);
96+ if (workspace_size > 0) {
97+ ret = aclrtMalloc(&workspace, workspace_size, ACL_MEM_MALLOC_HUGE_FIRST);
98+ CHECK_RET(ret == ACL_SUCCESS, goto cleanup);
99+ }
100+ 
101+ ret = aclnnInplaceRelu(workspace, workspace_size, executor, stream);
102+ CHECK_RET(ret == ACL_SUCCESS, goto cleanup);
103+ ret = aclrtSynchronizeStream(stream);
104+ CHECK_RET(ret == ACL_SUCCESS, goto cleanup);
105+ 
106+ ret = aclrtMemcpy(result.data(), bytes, device_addr, bytes, ACL_MEMCPY_DEVICE_TO_HOST);
107+ CHECK_RET(ret == ACL_SUCCESS, goto cleanup);
108+ 
109+ for (size_t i = 0; i < expected.size(); ++i) {
110+ if (std::fabs(result[i] - expected[i]) > 1e-6f) {
111+ std::fprintf(stderr, "check failed at %zu: got %.6f expected %.6f\n", i, result[i], expected[i]);
112+ ret = ACL_ERROR_FAILURE;
113+ goto cleanup;
114+ }
115+ }
116+ std::printf("inplace relu example passed\n");
117+ 
118+cleanup:
119+ if (self != nullptr) {
120+ aclDestroyTensor(self);
121+ }
122+ if (device_addr != nullptr) {
123+ aclrtFree(device_addr);
124+ }
125+ if (workspace != nullptr) {
126+ aclrtFree(workspace);
127+ }
128+ if (stream != nullptr) {
129+ aclrtDestroyStream(stream);
130+ }
131+ if (device_set) {
132+ aclrtResetDevice(device_id);
133+ }
134+ if (acl_initialized) {
135+ aclFinalize();
136+ }
137+ return ret;
138+}
Aexperimental/activation/relu/examples/test_aclnn_relu.cpp+298-0
@@ -0,0 +1,298 @@
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.cpp
13+ * @brief ACLNN invocation example for experimental Relu operator
14+ */
15+ 
16+#include <algorithm>
17+#include <cstdint>
18+#include <cstdio>
19+#include <cstdlib>
20+#include <cstring>
21+#include <fstream>
22+#include <string>
23+#include <vector>
24+ 
25+#include "acl/acl.h"
26+#include "aclnnop/aclnn_relu.h"
27+ 
28+#define CHECK_RET(cond, expr) \
29+ do { \
30+ if (!(cond)) { \
31+ expr; \
32+ } \
33+ } while (0)
34+ 
35+namespace {
36+struct ReluConfig {
37+ aclDataType acl_dtype;
38+ size_t element_size;
39+ std::string name;
40+};
41+ 
42+int64_t GetShapeSize(const std::vector<int64_t> &shape)
43+{
44+ int64_t shape_size = 1;
45+ for (int64_t dim : shape) {
46+ shape_size *= dim;
47+ }
48+ return shape_size;
49+}
50+ 
51+std::vector<int64_t> MakeStrides(const std::vector<int64_t> &shape)
52+{
53+ if (shape.empty()) {
54+ return {};
55+ }
56+ std::vector<int64_t> strides(shape.size(), 1);
57+ for (int64_t i = static_cast<int64_t>(shape.size()) - 2; i >= 0; --i) {
58+ strides[static_cast<size_t>(i)] = shape[static_cast<size_t>(i + 1)] * strides[static_cast<size_t>(i + 1)];
59+ }
60+ return strides;
61+}
62+ 
63+bool ParseDtype(const std::string &dtype_name, ReluConfig *config)
64+{
65+ if (dtype_name == "fp16" || dtype_name == "float16") {
66+ *config = {ACL_FLOAT16, sizeof(uint16_t), "fp16"};
67+ return true;
68+ }
69+ if (dtype_name == "fp32" || dtype_name == "float32") {
70+ *config = {ACL_FLOAT, sizeof(float), "fp32"};
71+ return true;
72+ }
73+ if (dtype_name == "bf16" || dtype_name == "bfloat16") {
74+ *config = {ACL_BF16, sizeof(uint16_t), "bf16"};
75+ return true;
76+ }
77+ if (dtype_name == "int8") {
78+ *config = {ACL_INT8, sizeof(int8_t), "int8"};
79+ return true;
80+ }
81+ if (dtype_name == "int32") {
82+ *config = {ACL_INT32, sizeof(int32_t), "int32"};
83+ return true;
84+ }
85+ if (dtype_name == "int64") {
86+ *config = {ACL_INT64, sizeof(int64_t), "int64"};
87+ return true;
88+ }
89+ return false;
90+}
91+ 
92+bool ParseShape(const std::string &shape_text, std::vector<int64_t> *shape)
93+{
94+ shape->clear();
95+ if (shape_text.empty() || shape_text == "scalar") {
96+ return true;
97+ }
98+ 
99+ size_t start = 0;
100+ while (start < shape_text.size()) {
101+ size_t end = shape_text.find(',', start);
102+ std::string token = shape_text.substr(start, end == std::string::npos ? std::string::npos : end - start);
103+ if (token.empty()) {
104+ return false;
105+ }
106+ shape->push_back(std::stoll(token));
107+ if (end == std::string::npos) {
108+ break;
109+ }
110+ start = end + 1;
111+ }
112+ return true;
113+}
114+ 
115+bool ReadFile(const std::string &path, std::vector<char> *buffer)
116+{
117+ std::ifstream stream(path, std::ios::binary);
118+ if (!stream.is_open()) {
119+ return false;
120+ }
121+ stream.seekg(0, std::ios::end);
122+ std::streamsize size = stream.tellg();
123+ stream.seekg(0, std::ios::beg);
124+ if (size < 0) {
125+ return false;
126+ }
127+ buffer->resize(static_cast<size_t>(size));
128+ return size == 0 || stream.read(buffer->data(), size).good();
129+}
130+ 
131+bool WriteFile(const std::string &path, const std::vector<char> &buffer)
132+{
133+ std::ofstream stream(path, std::ios::binary);
134+ if (!stream.is_open()) {
135+ return false;
136+ }
137+ stream.write(buffer.data(), static_cast<std::streamsize>(buffer.size()));
138+ return stream.good();
139+}
140+ 
141+aclError CreateAclTensor(
142+ const std::vector<int64_t> &shape, aclDataType dtype, void *device_addr, aclTensor **tensor)
143+{
144+ std::vector<int64_t> strides = MakeStrides(shape);
145+ const int64_t *shape_ptr = shape.empty() ? nullptr : shape.data();
146+ const int64_t *strides_ptr = strides.empty() ? nullptr : strides.data();
147+ *tensor = aclCreateTensor(
148+ shape_ptr, shape.size(), dtype, strides_ptr, 0, ACL_FORMAT_ND, shape_ptr, shape.size(), device_addr);
149+ return *tensor == nullptr ? ACL_ERROR_FAILURE : ACL_SUCCESS;
150+}
151+ 
152+int RunRelu(
153+ const std::vector<char> &input_host, const std::vector<int64_t> &shape, const ReluConfig &config,
154+ std::vector<char> *output_host, int32_t device_id)
155+{
156+ aclrtStream stream = nullptr;
157+ void *input_device = nullptr;
158+ void *output_device = nullptr;
159+ void *workspace = nullptr;
160+ aclTensor *input_tensor = nullptr;
161+ aclTensor *output_tensor = nullptr;
162+ aclOpExecutor *executor = nullptr;
163+ uint64_t workspace_size = 0;
164+ const size_t bytes = static_cast<size_t>(GetShapeSize(shape)) * config.element_size;
165+ bool acl_initialized = false;
166+ bool device_set = false;
167+ 
168+ auto ret = aclInit(nullptr);
169+ CHECK_RET(ret == ACL_SUCCESS, goto cleanup);
170+ acl_initialized = true;
171+ ret = aclrtSetDevice(device_id);
172+ CHECK_RET(ret == ACL_SUCCESS, goto cleanup);
173+ device_set = true;
174+ ret = aclrtCreateStream(&stream);
175+ CHECK_RET(ret == ACL_SUCCESS, goto cleanup);
176+ 
177+ if (bytes > 0) {
178+ ret = aclrtMalloc(&input_device, bytes, ACL_MEM_MALLOC_HUGE_FIRST);
179+ CHECK_RET(ret == ACL_SUCCESS, goto cleanup);
180+ ret = aclrtMalloc(&output_device, bytes, ACL_MEM_MALLOC_HUGE_FIRST);
181+ CHECK_RET(ret == ACL_SUCCESS, goto cleanup);
182+ 
183+ ret = aclrtMemcpy(input_device, bytes, input_host.data(), bytes, ACL_MEMCPY_HOST_TO_DEVICE);
184+ CHECK_RET(ret == ACL_SUCCESS, goto cleanup);
185+ std::vector<char> zero_buffer(bytes, 0);
186+ ret = aclrtMemcpy(output_device, bytes, zero_buffer.data(), bytes, ACL_MEMCPY_HOST_TO_DEVICE);
187+ CHECK_RET(ret == ACL_SUCCESS, goto cleanup);
188+ }
189+ 
190+ ret = CreateAclTensor(shape, config.acl_dtype, input_device, &input_tensor);
191+ CHECK_RET(ret == ACL_SUCCESS, goto cleanup);
192+ ret = CreateAclTensor(shape, config.acl_dtype, output_device, &output_tensor);
193+ CHECK_RET(ret == ACL_SUCCESS, goto cleanup);
194+ 
195+ ret = aclnnReluGetWorkspaceSize(input_tensor, output_tensor, &workspace_size, &executor);
196+ CHECK_RET(ret == ACL_SUCCESS, goto cleanup);
197+ if (workspace_size > 0) {
198+ ret = aclrtMalloc(&workspace, workspace_size, ACL_MEM_MALLOC_HUGE_FIRST);
199+ CHECK_RET(ret == ACL_SUCCESS, goto cleanup);
200+ }
201+ 
202+ ret = aclnnRelu(workspace, workspace_size, executor, stream);
203+ CHECK_RET(ret == ACL_SUCCESS, goto cleanup);
204+ ret = aclrtSynchronizeStream(stream);
205+ CHECK_RET(ret == ACL_SUCCESS, goto cleanup);
206+ 
207+ output_host->resize(bytes);
208+ if (bytes > 0) {
209+ ret = aclrtMemcpy(output_host->data(), bytes, output_device, bytes, ACL_MEMCPY_DEVICE_TO_HOST);
210+ CHECK_RET(ret == ACL_SUCCESS, goto cleanup);
211+ }
212+ 
213+cleanup:
214+ if (input_tensor != nullptr) {
215+ aclDestroyTensor(input_tensor);
216+ }
217+ if (output_tensor != nullptr) {
218+ aclDestroyTensor(output_tensor);
219+ }
220+ if (workspace != nullptr) {
221+ aclrtFree(workspace);
222+ }
223+ if (input_device != nullptr) {
224+ aclrtFree(input_device);
225+ }
226+ if (output_device != nullptr) {
227+ aclrtFree(output_device);
228+ }
229+ if (stream != nullptr) {
230+ aclrtDestroyStream(stream);
231+ }
232+ if (device_set) {
233+ aclrtResetDevice(device_id);
234+ }
235+ if (acl_initialized) {
236+ aclFinalize();
237+ }
238+ return ret;
239+}
240+ 
241+template <typename T>
242+std::vector<char> ToBytes(const std::vector<T> &values)
243+{
244+ std::vector<char> buffer(values.size() * sizeof(T));
245+ if (!buffer.empty()) {
246+ std::memcpy(buffer.data(), values.data(), buffer.size());
247+ }
248+ return buffer;
249+}
250+ 
251+int RunDefaultExample()
252+{
253+ ReluConfig config{ACL_FLOAT, sizeof(float), "fp32"};
254+ std::vector<int64_t> shape = {4, 2};
255+ std::vector<float> input = {-4.0f, -3.0f, -2.0f, 0.0f, 1.0f, 2.0f, 4.0f, 5.0f};
256+ std::vector<char> output;
257+ auto ret = RunRelu(ToBytes(input), shape, config, &output, 0);
258+ CHECK_RET(ret == ACL_SUCCESS, return ret);
259+ 
260+ const float *result = reinterpret_cast<const float *>(output.data());
261+ const std::vector<float> expected = {0.0f, 0.0f, 0.0f, 0.0f, 1.0f, 2.0f, 4.0f, 5.0f};
262+ for (size_t i = 0; i < expected.size(); ++i) {
263+ if (result[i] != expected[i]) {
264+ std::fprintf(stderr, "default check failed at %zu: got %.8f expected %.8f\n", i, result[i], expected[i]);
265+ return 1;
266+ }
267+ }
268+ std::printf("default example passed\n");
269+ return 0;
270+}
271+} // namespace
272+ 
273+int main(int argc, char **argv)
274+{
275+ if (argc == 1) {
276+ return RunDefaultExample();
277+ }
278+ 
279+ if (argc != 6) {
280+ std::fprintf(stderr, "Usage: %s <dtype> <shape|scalar> <input.bin> <output.bin> <device_id>\n", argv[0]);
281+ return 2;
282+ }
283+ 
284+ ReluConfig config{};
285+ CHECK_RET(ParseDtype(argv[1], &config), return 2);
286+ 
287+ std::vector<int64_t> shape;
288+ CHECK_RET(ParseShape(argv[2], &shape), return 2);
289+ 
290+ std::vector<char> input_host;
291+ CHECK_RET(ReadFile(argv[3], &input_host), return 3);
292+ 
293+ std::vector<char> output_host;
294+ int ret = RunRelu(input_host, shape, config, &output_host, std::atoi(argv[5]));
295+ CHECK_RET(ret == ACL_SUCCESS, return ret);
296+ CHECK_RET(WriteFile(argv[4], output_host), return 4);
297+ return 0;
298+}
Aexperimental/activation/relu/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 ACLNNTYPE aclnn_exclude DISABLE_IN_OPP TRUE)
Aexperimental/activation/relu/op_host/op_api/aclnn_relu.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_relu.h"
12+#include "aclnn_kernels/common/op_error_check.h"
13+#include "aclnn_kernels/contiguous.h"
14+#include "op_api/aclnn_util.h"
15+#include "opdev/make_op_executor.h"
16+#include "opdev/op_dfx.h"
17+#include "opdev/op_executor.h"
18+#include "relu.h"
19+ 
20+using namespace op;
21+ 
22+#ifdef __cplusplus
23+extern "C" {
24+#endif
25+ 
26+constexpr size_t MAX_DIM_LEN = 8;
27+ 
28+static const std::initializer_list<op::DataType> ASCEND910_DTYPE_SUPPORT_LIST = {
29+ op::DataType::DT_FLOAT,
30+ op::DataType::DT_FLOAT16,
31+ op::DataType::DT_INT8,
32+ op::DataType::DT_INT32,
33+ op::DataType::DT_INT64};
34+static const std::initializer_list<op::DataType> ASCEND910B_DTYPE_SUPPORT_LIST = {
35+ op::DataType::DT_FLOAT,
36+ op::DataType::DT_FLOAT16,
37+ op::DataType::DT_BF16,
38+ op::DataType::DT_INT8,
39+ op::DataType::DT_INT32,
40+ op::DataType::DT_INT64};
41+ 
42+static inline const std::initializer_list<op::DataType> &GetDtypeSupportList()
43+{
44+ if (GetCurrentPlatformInfo().GetSocVersion() >= SocVersion::ASCEND910B &&
45+ GetCurrentPlatformInfo().GetSocVersion() <= SocVersion::ASCEND910E) {
46+ return ASCEND910B_DTYPE_SUPPORT_LIST;
47+ }
48+ if (Ops::NN::AclnnUtil::IsRegbase()) {
49+ return ASCEND910B_DTYPE_SUPPORT_LIST;
50+ }
51+ return ASCEND910_DTYPE_SUPPORT_LIST;
52+}
53+ 
54+static bool CheckNotNull(const aclTensor *self, const aclTensor *out)
55+{
56+ OP_CHECK_NULL(self, return false);
57+ OP_CHECK_NULL(out, return false);
58+ return true;
59+}
60+ 
61+static bool CheckDtypeValid(const aclTensor *self, const aclTensor *out)
62+{
63+ auto supportList = GetDtypeSupportList();
64+ OP_CHECK_DTYPE_NOT_SUPPORT(self, supportList, return false);
65+ OP_CHECK_DTYPE_NOT_SUPPORT(out, supportList, return false);
66+ OP_CHECK_DTYPE_NOT_MATCH(self, out->GetDataType(), return false);
67+ return true;
68+}
69+ 
70+static bool CheckShape(const aclTensor *self, const aclTensor *out)
71+{
72+ OP_CHECK_SHAPE_NOT_EQUAL(self, out, return false);
73+ OP_CHECK_MAX_DIM(self, MAX_DIM_LEN, return false);
74+ OP_CHECK_MAX_DIM(out, MAX_DIM_LEN, return false);
75+ return true;
76+}
77+ 
78+static aclnnStatus CheckParams(const aclTensor *self, const aclTensor *out)
79+{
80+ CHECK_RET(CheckNotNull(self, out), ACLNN_ERR_PARAM_NULLPTR);
81+ CHECK_RET(CheckDtypeValid(self, out), ACLNN_ERR_PARAM_INVALID);
82+ CHECK_RET(CheckShape(self, out), ACLNN_ERR_PARAM_INVALID);
83+ return ACLNN_SUCCESS;
84+}
85+ 
86+static aclnnStatus CheckInplaceParams(const aclTensor *self)
87+{
88+ OP_CHECK_NULL(self, return ACLNN_ERR_PARAM_NULLPTR);
89+ auto supportList = GetDtypeSupportList();
90+ OP_CHECK_DTYPE_NOT_SUPPORT(self, supportList, return ACLNN_ERR_PARAM_INVALID);
91+ OP_CHECK_MAX_DIM(self, MAX_DIM_LEN, return ACLNN_ERR_PARAM_INVALID);
92+ return ACLNN_SUCCESS;
93+}
94+ 
95+static aclnnStatus ExecReluGetWorkspaceSize(
96+ const aclTensor *self, const aclTensor *out, uint64_t *workspaceSize, aclOpExecutor **executor)
97+{
98+ auto uniqueExecutor = CREATE_EXECUTOR();
99+ CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR);
100+ 
101+ auto ret = CheckParams(self, out);
102+ CHECK_RET(ret == ACLNN_SUCCESS, ret);
103+ 
104+ if (self->IsEmpty()) {
105+ *workspaceSize = 0;
106+ uniqueExecutor.ReleaseTo(executor);
107+ return ACLNN_SUCCESS;
108+ }
109+ 
110+ auto selfContiguous = l0op::Contiguous(self, uniqueExecutor.get());
111+ CHECK_RET(selfContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
112+ 
113+ auto reluOpOut = l0op::Relu(selfContiguous, uniqueExecutor.get());
114+ CHECK_RET(reluOpOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
115+ 
116+ auto viewCopyResult = l0op::ViewCopy(reluOpOut, out, uniqueExecutor.get());
117+ CHECK_RET(viewCopyResult != nullptr, ACLNN_ERR_INNER_NULLPTR);
118+ 
119+ *workspaceSize = uniqueExecutor->GetWorkspaceSize();
120+ uniqueExecutor.ReleaseTo(executor);
121+ return ACLNN_SUCCESS;
122+}
123+ 
124+aclnnStatus aclnnReluGetWorkspaceSize(
125+ const aclTensor *self, const aclTensor *out, uint64_t *workspaceSize, aclOpExecutor **executor)
126+{
127+ OP_CHECK_COMM_INPUT(workspaceSize, executor);
128+ L2_DFX_PHASE_1(aclnnRelu, DFX_IN(self), DFX_OUT(out));
129+ return ExecReluGetWorkspaceSize(self, out, workspaceSize, executor);
130+}
131+ 
132+aclnnStatus aclnnInplaceReluGetWorkspaceSize(
133+ aclTensor *selfRef, uint64_t *workspaceSize, aclOpExecutor **executor)
134+{
135+ OP_CHECK_COMM_INPUT(workspaceSize, executor);
136+ L2_DFX_PHASE_1(aclnnInplaceRelu, DFX_IN(selfRef), DFX_OUT(selfRef));
137+ 
138+ auto ret = CheckInplaceParams(selfRef);
139+ CHECK_RET(ret == ACLNN_SUCCESS, ret);
140+ 
141+ return ExecReluGetWorkspaceSize(selfRef, selfRef, workspaceSize, executor);
142+}
143+ 
144+aclnnStatus aclnnRelu(void *workspace, uint64_t workspaceSize, aclOpExecutor *executor, const aclrtStream stream)
145+{
146+ L2_DFX_PHASE_2(aclnnRelu);
147+ return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);
148+}
149+ 
150+aclnnStatus aclnnInplaceRelu(
151+ void *workspace, uint64_t workspaceSize, aclOpExecutor *executor, const aclrtStream stream)
152+{
153+ L2_DFX_PHASE_2(aclnnInplaceRelu);
154+ return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);
155+}
156+ 
157+#ifdef __cplusplus
158+}
159+#endif
Aexperimental/activation/relu/op_host/op_api/aclnn_relu.h+37-0
@@ -0,0 +1,37 @@
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_RELU_H_
12+#define ACLNN_RELU_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 aclnnReluGetWorkspaceSize(
22+ const aclTensor *self, const aclTensor *out, uint64_t *workspaceSize, aclOpExecutor **executor);
23+ 
24+ACLNN_API aclnnStatus aclnnRelu(
25+ void *workspace, uint64_t workspaceSize, aclOpExecutor *executor, const aclrtStream stream);
26+ 
27+ACLNN_API aclnnStatus aclnnInplaceReluGetWorkspaceSize(
28+ aclTensor *selfRef, uint64_t *workspaceSize, aclOpExecutor **executor);
29+ 
30+ACLNN_API aclnnStatus aclnnInplaceRelu(
31+ void *workspace, uint64_t workspaceSize, aclOpExecutor *executor, const aclrtStream stream);
32+ 
33+#ifdef __cplusplus
34+}
35+#endif
36+ 
37+#endif
Aexperimental/activation/relu/op_host/op_api/relu.cpp+37-0
@@ -0,0 +1,37 @@
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.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(Relu);
24+ 
25+const aclTensor *Relu(const aclTensor *self, aclOpExecutor *executor)
26+{
27+ L0_DFX(Relu, self);
28+ auto reluOut = executor->AllocTensor(self->GetStorageShape(), self->GetDataType(), self->GetStorageFormat());
29+ CHECK_RET(reluOut != nullptr, nullptr);
30+ 
31+ auto ret = ADD_TO_LAUNCHER_LIST_AICORE(Relu, OP_INPUT(self), OP_OUTPUT(reluOut));
32+ OP_CHECK(ret == ACLNN_SUCCESS, OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "Relu ADD_TO_LAUNCHER_LIST_AICORE failed."),
33+ return nullptr);
34+ return reluOut;
35+}
36+ 
37+} // namespace l0op
Aexperimental/activation/relu/op_host/op_api/relu.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_H_
12+#define OP_API_OP_API_COMMON_INC_LEVEL0_OP_RELU_H_
13+ 
14+#include "opdev/make_op_executor.h"
15+#include "opdev/op_executor.h"
16+ 
17+namespace l0op {
18+const aclTensor *Relu(const aclTensor *self, aclOpExecutor *executor);
19+}
20+ 
21+#endif
Aexperimental/activation/relu/op_host/relu_def.cpp+47-0
@@ -0,0 +1,47 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include "register/op_def_registry.h"
12+ 
13+namespace ops {
14+class Relu : public OpDef {
15+public:
16+ explicit Relu(const char *name) : OpDef(name)
17+ {
18+ this->Input("x")
19+ .ParamType(REQUIRED)
20+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16, ge::DT_INT8, ge::DT_INT32, ge::DT_INT64})
21+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
22+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
23+ ge::FORMAT_ND})
24+ .AutoContiguous();
25+ 
26+ this->Output("y")
27+ .ParamType(REQUIRED)
28+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16, ge::DT_INT8, ge::DT_INT32, ge::DT_INT64})
29+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
30+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
31+ ge::FORMAT_ND})
32+ .AutoContiguous();
33+ 
34+ OpAICoreConfig aicoreConfig;
35+ aicoreConfig.DynamicCompileStaticFlag(true)
36+ .DynamicFormatFlag(false)
37+ .DynamicRankSupportFlag(true)
38+ .DynamicShapeSupportFlag(true)
39+ .NeedCheckSupportFlag(false)
40+ .PrecisionReduceFlag(true)
41+ .ExtendCfgInfo("opFile.value", "relu");
42+ this->AICore().AddConfig("ascend910b", aicoreConfig);
43+ }
44+};
45+ 
46+OP_ADD(Relu);
47+} // namespace ops
Aexperimental/activation/relu/op_host/relu_infershape.cpp+22-0
@@ -0,0 +1,22 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include "register/op_impl_registry.h"
12+#include "op_common/op_host/infershape_elewise_util.h"
13+ 
14+namespace ops {
15+ 
16+static ge::graphStatus InferShape4Relu(gert::InferShapeContext *context)
17+{
18+ return Ops::Base::InferShape4Elewise(context);
19+}
20+ 
21+IMPL_OP_INFERSHAPE(Relu).InferShape(InferShape4Relu);
22+} // namespace ops
Aexperimental/activation/relu/op_host/relu_tiling.cpp+180-0
@@ -0,0 +1,180 @@
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_tiling_data.h"
19+#include "../op_kernel/relu_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 MIN_BYTES_PER_CORE = 16 * 1024;
27+ 
28+struct ReluCompileInfo {};
29+ 
30+static const gert::Shape SCALAR_SHAPE = {1};
31+ 
32+static const gert::Shape &EnsureNotScalar(const gert::Shape &shape)
33+{
34+ if (shape.GetDimNum() == 0) {
35+ return SCALAR_SHAPE;
36+ }
37+ return shape;
38+}
39+ 
40+static ge::graphStatus GetPlatformInfo(gert::TilingContext *context, uint64_t &ubSize, int64_t &coreNum)
41+{
42+ auto *platformInfoPtr = context->GetPlatformInfo();
43+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr);
44+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
45+ coreNum = static_cast<int64_t>(ascendcPlatform.GetCoreNumAiv());
46+ OP_CHECK_IF(coreNum <= 0, OP_LOGE(context, "coreNum is invalid"), return ge::GRAPH_FAILED);
47+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
48+ OP_CHECK_IF(ubSize == 0, OP_LOGE(context, "ubSize is 0"), return ge::GRAPH_FAILED);
49+ return ge::GRAPH_SUCCESS;
50+}
51+ 
52+static ge::graphStatus GetShapeAttrsInfo(gert::TilingContext *context, int64_t &totalLength, ge::DataType &dataType)
53+{
54+ auto *inputX = context->GetInputShape(0);
55+ OP_CHECK_NULL_WITH_CONTEXT(context, inputX);
56+ totalLength = EnsureNotScalar(inputX->GetStorageShape()).GetShapeSize();
57+ 
58+ auto *inputDesc = context->GetInputDesc(0);
59+ OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc);
60+ dataType = inputDesc->GetDataType();
61+ const std::set<ge::DataType> supportedDtype = {
62+ ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16, ge::DT_INT8, ge::DT_INT32, ge::DT_INT64};
63+ OP_CHECK_IF(supportedDtype.count(dataType) == 0, OP_LOGE(context, "invalid dtype"), return ge::GRAPH_FAILED);
64+ 
65+ return ge::GRAPH_SUCCESS;
66+}
67+ 
68+static ge::graphStatus GetWorkspaceSize(gert::TilingContext *context)
69+{
70+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
71+ uint32_t sysWorkspaceSize = ascendcPlatform.GetLibApiWorkSpaceSize();
72+ size_t *currentWorkspace = context->GetWorkspaceSizes(1);
73+ OP_CHECK_NULL_WITH_CONTEXT(context, currentWorkspace);
74+ currentWorkspace[0] = sysWorkspaceSize;
75+ return ge::GRAPH_SUCCESS;
76+}
77+ 
78+static void FillDefaultTiling(ReluTilingData *tiling, uint32_t dTypeX, gert::TilingContext *context)
79+{
80+ tiling->formerNum = 0;
81+ tiling->formerLength = 0;
82+ tiling->tailLength = 0;
83+ tiling->tileLength = 1;
84+ context->SetBlockDim(1);
85+ ASCENDC_TPL_SEL_PARAM(context, dTypeX);
86+}
87+ 
88+static void FillNormalTiling(ReluTilingData *tiling, int64_t totalLength, uint64_t ubSize, int64_t coreNum,
89+ ge::DataType dataType, uint32_t dTypeX, gert::TilingContext *context)
90+{
91+ uint32_t typeLength = 0;
92+ ge::TypeUtils::GetDataTypeLength(dataType, typeLength);
93+ OP_CHECK_IF(typeLength == 0, OP_LOGE(context, "typeLength is 0"), return);
94+ 
95+ int64_t dtypeSize = static_cast<int64_t>(typeLength);
96+ int64_t totalBytes = totalLength * dtypeSize;
97+ int64_t targetCoreNum = 1;
98+ if (dataType == ge::DT_INT64) {
99+ targetCoreNum = coreNum;
100+ } else {
101+ targetCoreNum = (totalBytes + MIN_BYTES_PER_CORE - 1) / MIN_BYTES_PER_CORE;
102+ targetCoreNum = std::max<int64_t>(1, targetCoreNum);
103+ targetCoreNum = std::min<int64_t>(targetCoreNum, coreNum);
104+ }
105+ 
106+ int64_t cacheLineElements = std::max<int64_t>(1, CACHE_LINE_BYTE_LENGTH / dtypeSize);
107+ int64_t totalLengthCore = (totalLength + targetCoreNum - 1) / targetCoreNum;
108+ int64_t totalLengthCoreAlign =
109+ ((totalLengthCore + cacheLineElements - 1) / cacheLineElements) * cacheLineElements;
110+ 
111+ int64_t usedCoreNum = (totalLength + totalLengthCoreAlign - 1) / totalLengthCoreAlign;
112+ usedCoreNum = std::max<int64_t>(1, usedCoreNum);
113+ int64_t formerNum = usedCoreNum - 1;
114+ int64_t formerLength = totalLengthCoreAlign;
115+ int64_t tailLength = totalLength - formerNum * formerLength;
116+ 
117+ int64_t bufferCoefficient = 8;
118+ if (dataType == ge::DT_FLOAT16) {
119+ bufferCoefficient = 4;
120+ } else if (dataType == ge::DT_BF16) {
121+ bufferCoefficient = 8;
122+ } else if (dataType == ge::DT_INT8) {
123+ bufferCoefficient = 4;
124+ } else if (dataType == ge::DT_INT64) {
125+ bufferCoefficient = 32;
126+ }
127+ 
128+ int64_t maxTileElements = static_cast<int64_t>(ubSize) / bufferCoefficient;
129+ int64_t alignElements = std::max<int64_t>(1, 32 / dtypeSize);
130+ int64_t tileLength = (maxTileElements / alignElements) * alignElements;
131+ if (tileLength <= 0) {
132+ tileLength = alignElements;
133+ }
134+ 
135+ tiling->formerNum = formerNum;
136+ tiling->formerLength = formerLength;
137+ tiling->tailLength = tailLength;
138+ tiling->tileLength = tileLength;
139+ 
140+ context->SetBlockDim(static_cast<uint32_t>(usedCoreNum));
141+ ASCENDC_TPL_SEL_PARAM(context, dTypeX);
142+}
143+ 
144+static ge::graphStatus ReluTilingFunc(gert::TilingContext *context)
145+{
146+ uint64_t ubSize = 0;
147+ int64_t coreNum = 0;
148+ OP_CHECK_IF(GetPlatformInfo(context, ubSize, coreNum) != ge::GRAPH_SUCCESS,
149+ OP_LOGE(context, "GetPlatformInfo error"), return ge::GRAPH_FAILED);
150+ 
151+ int64_t totalLength = 0;
152+ ge::DataType dataType;
153+ OP_CHECK_IF(GetShapeAttrsInfo(context, totalLength, dataType) != ge::GRAPH_SUCCESS,
154+ OP_LOGE(context, "GetShapeAttrsInfo error"), return ge::GRAPH_FAILED);
155+ 
156+ OP_CHECK_IF(GetWorkspaceSize(context) != ge::GRAPH_SUCCESS,
157+ OP_LOGE(context, "GetWorkspaceSize error"), return ge::GRAPH_FAILED);
158+ 
159+ auto *tiling = context->GetTilingData<ReluTilingData>();
160+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
161+ OP_CHECK_IF(memset_s(tiling, sizeof(ReluTilingData), 0, sizeof(ReluTilingData)) != EOK,
162+ OP_LOGE(context, "set tiling data error"), return ge::GRAPH_FAILED);
163+ 
164+ uint32_t dTypeX = static_cast<uint32_t>(dataType);
165+ if (totalLength <= 0) {
166+ FillDefaultTiling(tiling, dTypeX, context);
167+ return ge::GRAPH_SUCCESS;
168+ }
169+ 
170+ FillNormalTiling(tiling, totalLength, ubSize, coreNum, dataType, dTypeX, context);
171+ return ge::GRAPH_SUCCESS;
172+}
173+ 
174+static ge::graphStatus TilingParseForRelu([[maybe_unused]] gert::TilingParseContext *context)
175+{
176+ return ge::GRAPH_SUCCESS;
177+}
178+ 
179+IMPL_OP_OPTILING(Relu).Tiling(ReluTilingFunc).TilingParse<ReluCompileInfo>(TilingParseForRelu);
180+} // namespace optiling
Aexperimental/activation/relu/op_kernel/relu.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 <type_traits>
12+ 
13+#include "relu.h"
14+ 
15+template <typename D_T_X>
16+__global__ __aicore__ void relu(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
17+{
18+ REGISTER_TILING_DEFAULT(ReluTilingData);
19+ GET_TILING_DATA_WITH_STRUCT(ReluTilingData, tilingData, tiling);
20+ 
21+ if constexpr (std::is_same_v<D_T_X, float> || std::is_same_v<D_T_X, half> || std::is_same_v<D_T_X, int32_t>) {
22+ NsRelu::KernelRelu<D_T_X> op;
23+ op.Init(x, y, &tilingData);
24+ op.Process();
25+ } else if constexpr (std::is_same_v<D_T_X, bfloat16_t>) {
26+ NsRelu::KernelReluUpcast<D_T_X, float> op;
27+ op.Init(x, y, &tilingData);
28+ op.Process();
29+ } else if constexpr (std::is_same_v<D_T_X, int8_t>) {
30+ NsRelu::KernelReluUpcast<D_T_X, half> op;
31+ op.Init(x, y, &tilingData);
32+ op.Process();
33+ } else {
34+ NsRelu::KernelReluScalarInt64<D_T_X> op;
35+ op.Init(x, y, &tilingData);
36+ op.Process();
37+ }
38+}
Aexperimental/activation/relu/op_kernel/relu.h+245-0
@@ -0,0 +1,245 @@
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_H_
12+#define EXPERIMENTAL_ACTIVATION_RELU_H_
13+ 
14+#include "kernel_operator.h"
15+#include "kernel_tiling/kernel_tiling.h"
16+#include "relu_tiling_data.h"
17+#include "relu_tiling_key.h"
18+ 
19+namespace NsRelu {
20+ 
21+using namespace AscendC;
22+ 
23+constexpr int32_t BUFFER_NUM = 2;
24+constexpr int32_t QUEUE_DEPTH = 1;
25+ 
26+template <typename T>
27+__aicore__ inline void CopyInFromGm(TQue<TPosition::VECIN, BUFFER_NUM> &inQueueX, GlobalTensor<T> &xGm,
28+ int64_t progress, int64_t tileLength, int64_t curTileLength)
29+{
30+ LocalTensor<T> xLocal = inQueueX.AllocTensor<T>();
31+ DataCopyExtParams copyParams{1, static_cast<uint32_t>(curTileLength * sizeof(T)), 0, 0, 0};
32+ DataCopyPadExtParams<T> padParams{false, 0, 0, 0};
33+ DataCopyPad(xLocal, xGm[progress * tileLength], copyParams, padParams);
34+ inQueueX.EnQue(xLocal);
35+}
36+ 
37+template <typename T>
38+__aicore__ inline void CopyOutToGm(TQue<TPosition::VECOUT, BUFFER_NUM> &outQueueY, GlobalTensor<T> &yGm,
39+ int64_t progress, int64_t tileLength, int64_t curTileLength)
40+{
41+ LocalTensor<T> yLocal = outQueueY.DeQue<T>();
42+ DataCopyExtParams copyParams{1, static_cast<uint32_t>(curTileLength * sizeof(T)), 0, 0, 0};
43+ DataCopyPad(yGm[progress * tileLength], yLocal, copyParams);
44+ outQueueY.FreeTensor(yLocal);
45+}
46+ 
47+template <typename T>
48+__aicore__ inline void CopyInFromGm(
49+ TQueBind<TPosition::VECIN, TPosition::VECOUT, QUEUE_DEPTH> &inOutQueue, GlobalTensor<T> &xGm, int64_t progress,
50+ int64_t tileLength, int64_t curTileLength)
51+{
52+ LocalTensor<T> xLocal = inOutQueue.AllocTensor<T>();
53+ DataCopyExtParams copyParams{1, static_cast<uint32_t>(curTileLength * sizeof(T)), 0, 0, 0};
54+ DataCopyPadExtParams<T> padParams{false, 0, 0, 0};
55+ DataCopyPad(xLocal, xGm[progress * tileLength], copyParams, padParams);
56+ inOutQueue.template EnQue<QuePosition::GM, QuePosition::VECIN, T>(xLocal);
57+}
58+ 
59+template <typename T>
60+__aicore__ inline void CopyOutToGm(
61+ TQueBind<TPosition::VECIN, TPosition::VECOUT, QUEUE_DEPTH> &inOutQueue, GlobalTensor<T> &yGm,
62+ int64_t progress, int64_t tileLength, int64_t curTileLength)
63+{
64+ LocalTensor<T> yLocal = inOutQueue.template DeQue<QuePosition::VECOUT, QuePosition::GM, T>();
65+ DataCopyExtParams copyParams{1, static_cast<uint32_t>(curTileLength * sizeof(T)), 0, 0, 0};
66+ DataCopyPad(yGm[progress * tileLength], yLocal, copyParams);
67+ inOutQueue.FreeTensor(yLocal);
68+}
69+ 
70+template <typename T>
71+__aicore__ inline void InitBlockGm(GlobalTensor<T> &xGm, GlobalTensor<T> &yGm, int64_t &blockLength,
72+ int64_t &tileLength, GM_ADDR x, GM_ADDR y, const ReluTilingData *tilingData)
73+{
74+ int64_t blockIdx = GetBlockIdx();
75+ if (blockIdx < tilingData->formerNum) {
76+ blockLength = tilingData->formerLength;
77+ int64_t offset = tilingData->formerLength * blockIdx;
78+ xGm.SetGlobalBuffer((__gm__ T *)x + offset, tilingData->formerLength);
79+ yGm.SetGlobalBuffer((__gm__ T *)y + offset, tilingData->formerLength);
80+ } else {
81+ blockLength = tilingData->tailLength;
82+ int64_t offset = tilingData->formerLength * tilingData->formerNum;
83+ xGm.SetGlobalBuffer((__gm__ T *)x + offset, tilingData->tailLength);
84+ yGm.SetGlobalBuffer((__gm__ T *)y + offset, tilingData->tailLength);
85+ }
86+ 
87+ tileLength = tilingData->tileLength;
88+}
89+ 
90+template <typename Kernel>
91+__aicore__ inline void ProcessTiles(Kernel &kernel, int64_t blockLength, int64_t tileLength)
92+{
93+ int64_t tileNum = (blockLength + tileLength - 1) / tileLength;
94+ if (tileNum == 0) {
95+ return;
96+ }
97+ 
98+ int64_t tailTileLength = blockLength - (tileNum - 1) * tileLength;
99+ for (int64_t i = 0; i < tileNum - 1; ++i) {
100+ kernel.CopyIn(i, tileLength);
101+ kernel.Compute(tileLength);
102+ kernel.CopyOut(i, tileLength);
103+ }
104+ kernel.CopyIn(tileNum - 1, tailTileLength);
105+ kernel.Compute(tailTileLength);
106+ kernel.CopyOut(tileNum - 1, tailTileLength);
107+}
108+ 
109+template <typename Derived, typename T>
110+class KernelReluBase {
111+public:
112+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, const ReluTilingData *tilingData)
113+ {
114+ InitBlockGm(xGm_, yGm_, blockLength_, tileLength_, x, y, tilingData);
115+ pipe_.InitBuffer(inOutQueue_, BUFFER_NUM, tileLength_ * sizeof(T));
116+ }
117+ 
118+ __aicore__ inline void Process()
119+ {
120+ ProcessTiles(static_cast<Derived &>(*this), blockLength_, tileLength_);
121+ }
122+ 
123+ __aicore__ inline void CopyIn(int64_t progress, int64_t curTileLength)
124+ {
125+ CopyInFromGm(inOutQueue_, xGm_, progress, tileLength_, curTileLength);
126+ }
127+ 
128+ __aicore__ inline void CopyOut(int64_t progress, int64_t curTileLength)
129+ {
130+ CopyOutToGm(inOutQueue_, yGm_, progress, tileLength_, curTileLength);
131+ }
132+ 
133+protected:
134+ TPipe pipe_;
135+ TQueBind<TPosition::VECIN, TPosition::VECOUT, QUEUE_DEPTH> inOutQueue_;
136+ GlobalTensor<T> xGm_;
137+ GlobalTensor<T> yGm_;
138+ int64_t blockLength_ = 0;
139+ int64_t tileLength_ = 0;
140+};
141+ 
142+template <typename T>
143+class KernelRelu : public KernelReluBase<KernelRelu<T>, T> {
144+public:
145+ __aicore__ inline KernelRelu() {}
146+ 
147+ __aicore__ inline void Compute(int64_t curTileLength)
148+ {
149+ LocalTensor<T> xLocal = this->inOutQueue_.template DeQue<QuePosition::GM, QuePosition::VECIN, T>();
150+ Relu(xLocal, xLocal, curTileLength);
151+ this->inOutQueue_.template EnQue<QuePosition::VECOUT, QuePosition::GM, T>(xLocal);
152+ }
153+};
154+ 
155+template <typename T, typename MidT>
156+class KernelReluUpcast : public KernelReluBase<KernelReluUpcast<T, MidT>, T> {
157+public:
158+ __aicore__ inline KernelReluUpcast() {}
159+ 
160+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, const ReluTilingData *tilingData)
161+ {
162+ KernelReluBase<KernelReluUpcast<T, MidT>, T>::Init(x, y, tilingData);
163+ this->pipe_.InitBuffer(tmpBufX_, this->tileLength_ * sizeof(MidT));
164+ }
165+ 
166+ __aicore__ inline void Compute(int64_t curTileLength)
167+ {
168+ LocalTensor<T> xLocal = this->inOutQueue_.template DeQue<QuePosition::GM, QuePosition::VECIN, T>();
169+ LocalTensor<MidT> xMid = tmpBufX_.Get<MidT>();
170+ Cast(xMid, xLocal, RoundMode::CAST_NONE, curTileLength);
171+ Relu(xMid, xMid, curTileLength);
172+ Cast(xLocal, xMid, RoundMode::CAST_RINT, curTileLength);
173+ this->inOutQueue_.template EnQue<QuePosition::VECOUT, QuePosition::GM, T>(xLocal);
174+ }
175+ 
176+private:
177+ TBuf<TPosition::VECCALC> tmpBufX_;
178+};
179+ 
180+template <typename T>
181+class KernelReluScalarInt64 {
182+public:
183+ __aicore__ inline KernelReluScalarInt64() {}
184+ 
185+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, const ReluTilingData *tilingData)
186+ {
187+ InitBlockGm(xGm_, yGm_, blockLength_, tileLength_, x, y, tilingData);
188+ pipe_.InitBuffer(inQueueX_, BUFFER_NUM, tileLength_ * sizeof(T));
189+ pipe_.InitBuffer(outQueueY_, BUFFER_NUM, tileLength_ * sizeof(T));
190+ }
191+ 
192+ __aicore__ inline void Process()
193+ {
194+ int64_t tileNum = (blockLength_ + tileLength_ - 1) / tileLength_;
195+ if (tileNum == 0) {
196+ return;
197+ }
198+ 
199+ int64_t tailTileLength = blockLength_ - (tileNum - 1) * tileLength_;
200+ for (int64_t i = 0; i < tileNum - 1; ++i) {
201+ CopyIn(i, tileLength_);
202+ Compute(tileLength_);
203+ CopyOut(i, tileLength_);
204+ }
205+ CopyIn(tileNum - 1, tailTileLength);
206+ Compute(tailTileLength);
207+ CopyOut(tileNum - 1, tailTileLength);
208+ }
209+ 
210+private:
211+ __aicore__ inline void CopyIn(int64_t progress, int64_t curTileLength)
212+ {
213+ CopyInFromGm(inQueueX_, xGm_, progress, tileLength_, curTileLength);
214+ }
215+ 
216+ __aicore__ inline void Compute(int64_t curTileLength)
217+ {
218+ LocalTensor<T> xLocal = inQueueX_.DeQue<T>();
219+ LocalTensor<T> yLocal = outQueueY_.AllocTensor<T>();
220+ for (int64_t index = 0; index < curTileLength; ++index) {
221+ T value = xLocal.GetValue(index);
222+ yLocal.SetValue(index, value > static_cast<T>(0) ? value : static_cast<T>(0));
223+ }
224+ outQueueY_.EnQue<T>(yLocal);
225+ inQueueX_.FreeTensor(xLocal);
226+ }
227+ 
228+ __aicore__ inline void CopyOut(int64_t progress, int64_t curTileLength)
229+ {
230+ CopyOutToGm(outQueueY_, yGm_, progress, tileLength_, curTileLength);
231+ }
232+ 
233+private:
234+ TPipe pipe_;
235+ TQue<TPosition::VECIN, BUFFER_NUM> inQueueX_;
236+ TQue<TPosition::VECOUT, BUFFER_NUM> outQueueY_;
237+ GlobalTensor<T> xGm_;
238+ GlobalTensor<T> yGm_;
239+ int64_t blockLength_ = 0;
240+ int64_t tileLength_ = 0;
241+};
242+ 
243+} // namespace NsRelu
244+ 
245+#endif
Aexperimental/activation/relu/op_kernel/relu_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_TILING_DATA_H_
12+#define RELU_TILING_DATA_H_
13+ 
14+#include <cstdint>
15+ 
16+struct ReluTilingData {
17+ int64_t formerNum;
18+ int64_t formerLength;
19+ int64_t tailLength;
20+ int64_t tileLength;
21+};
22+ 
23+#endif
Aexperimental/activation/relu/op_kernel/relu_tiling_key.h+30-0
@@ -0,0 +1,30 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#ifndef __RELU_TILING_KEY_H__
12+#define __RELU_TILING_KEY_H__
13+ 
14+#include "ascendc/host_api/tiling/template_argument.h"
15+ 
16+ASCENDC_TPL_ARGS_DECL(
17+ Relu,
18+ ASCENDC_TPL_DATATYPE_DECL(D_T_X, C_DT_FLOAT, C_DT_FLOAT16, C_DT_BF16, C_DT_INT8, C_DT_INT32, C_DT_INT64,
19+ 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_INT32)),
27+ ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_INT64)),
28+);
29+ 
30+#endif
Aexperimental/activation/relu/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/tests/st/aclnnRelu/all_aclnnRelu.json+75-0
@@ -0,0 +1,75 @@
1+[
2+ {
3+ "id": 1,
4+ "name": "torch.relu fp16",
5+ "aclnn_name": "aclnnRelu",
6+ "version": "v1.0",
7+ "expected_error_msg": null,
8+ "api": "pytorch",
9+ "api_type": "aclnn_relu",
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": "x",
20+ "type": "tensor",
21+ "required": true,
22+ "dtype": "fp16",
23+ "shape": [
24+ 128
25+ ],
26+ "range_values": [
27+ -10,
28+ 10
29+ ],
30+ "backward": false,
31+ "align_32B": null
32+ }
33+ ],
34+ "acl_json": "",
35+ "method_inputs": null,
36+ "tensor_input": null
37+ },
38+ {
39+ "id": 2,
40+ "name": "torch.relu bf16",
41+ "aclnn_name": "aclnnRelu",
42+ "version": "v1.0",
43+ "expected_error_msg": null,
44+ "api": "pytorch",
45+ "api_type": "aclnn_relu",
46+ "aclnn_api_type": "aclnn_function",
47+ "backward": false,
48+ "standard": {
49+ "acc": "cv_fused_double_benchmark",
50+ "perf": "not_key"
51+ },
52+ "outputs": null,
53+ "inputs": [
54+ {
55+ "name": "x",
56+ "type": "tensor",
57+ "required": true,
58+ "dtype": "bf16",
59+ "shape": [
60+ 64,
61+ 256
62+ ],
63+ "range_values": [
64+ -10,
65+ 10
66+ ],
67+ "backward": false,
68+ "align_32B": null
69+ }
70+ ],
71+ "acl_json": "",
72+ "method_inputs": null,
73+ "tensor_input": null
74+ }
75+]
Aexperimental/activation/relu/tests/st/aclnnRelu/executor_aclnnRelu.py+29-0
@@ -0,0 +1,29 @@
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 relu_reference(x: torch.Tensor) -> torch.Tensor:
21+ zeros = torch.zeros_like(x)
22+ return torch.where(x > 0, x, zeros)
23+ 
24+ 
25+@register("aclnn_relu")
26+class TorchRelu(BaseApi):
27+ def __call__(self, input_data: InputDataset, with_output: bool = False):
28+ x = input_data.kwargs["x"]
29+ return relu_reference(x)
Aexperimental/activation/relu/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/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/tests/ut/op_api/test_aclnn_relu.cpp+348-0
@@ -0,0 +1,348 @@
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_relu.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_relu_test : public testing::Test {
21+ protected:
22+ static void SetUpTestCase()
23+ {
24+ std::cout << "relu_test SetUp" << std::endl;
25+ }
26+ 
27+ static void TearDownTestCase()
28+ {
29+ std::cout << "relu_test TearDown" << std::endl;
30+ }
31+};
32+ 
33+TEST_F(l2_relu_test, case_001_float)
34+{
35+ auto selfDesc = TensorDesc({2, 4}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-10, 10);
36+ auto outDesc = TensorDesc({2, 4}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.0001, 0.0001);
37+ 
38+ auto ut = OP_API_UT(aclnnRelu, INPUT(selfDesc), OUTPUT(outDesc));
39+ 
40+ uint64_t workspaceSize = 0;
41+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
42+ EXPECT_EQ(aclRet, ACL_SUCCESS);
43+ 
44+ ut.TestPrecision();
45+}
46+ 
47+TEST_F(l2_relu_test, case_002_float16)
48+{
49+ auto selfDesc = TensorDesc({2, 4}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-10, 10);
50+ auto outDesc = TensorDesc({2, 4}, ACL_FLOAT16, ACL_FORMAT_ND).Precision(0.001, 0.001);
51+ 
52+ auto ut = OP_API_UT(aclnnRelu, INPUT(selfDesc), OUTPUT(outDesc));
53+ 
54+ uint64_t workspaceSize = 0;
55+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
56+ EXPECT_EQ(aclRet, ACL_SUCCESS);
57+ 
58+ ut.TestPrecision();
59+}
60+ 
61+TEST_F(l2_relu_test, case_003_bfloat16)
62+{
63+ auto selfDesc = TensorDesc({2, 4}, ACL_BF16, ACL_FORMAT_ND).ValueRange(-10, 10);
64+ auto outDesc = TensorDesc({2, 4}, ACL_BF16, ACL_FORMAT_ND).Precision(0.01, 0.01);
65+ 
66+ auto ut = OP_API_UT(aclnnRelu, INPUT(selfDesc), OUTPUT(outDesc));
67+ 
68+ uint64_t workspaceSize = 0;
69+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
S
Ssunday5月8日

PR 描述里把这个用例写成 BFLOAT16 精度验证，但这里只调用了 TestGetWorkspaceSize，没有调用 TestPrecision 做数值比对。如果确实要覆盖 BF16 精度，建议像 fp16、fp32 那样补上 TestPrecision；如果暂时不打算做精度校验，建议同步改 PR 描述或测试清单里的措辞，避免对外宣称的测试内容和代码不一致。

likedislike
C
cwzhang
5月11日 评论：
70+ EXPECT_EQ(aclRet, ACL_SUCCESS);
71+ 
72+ ut.TestPrecision();
73+}
74+ 
75+TEST_F(l2_relu_test, case_004_int8)
76+{
77+ auto selfDesc = TensorDesc({1024}, ACL_INT8, ACL_FORMAT_ND).ValueRange(-20, 20);
78+ auto outDesc = TensorDesc({1024}, ACL_INT8, ACL_FORMAT_ND);
79+ 
80+ auto ut = OP_API_UT(aclnnRelu, INPUT(selfDesc), OUTPUT(outDesc));
81+ 
82+ uint64_t workspaceSize = 0;
83+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
S
Ssunday5月8日

case_004 到 case_006 目前都只断言 GetWorkspaceSize 成功，没有对输出张量做 ReLU 结果或等价性检查。如果测试清单里写的是整型类型的精度或正确性验证，建议至少补一层与参考实现或期望向量的比对；如果只是验证接口参数合法，可以把用例说明改成接口可用性测试，避免名称和实际断言力度脱节。

likedislike
C
cwzhang
5月11日 评论：
84+ EXPECT_EQ(aclRet, ACL_SUCCESS);
85+ 
86+ ut.TestPrecision();
87+}
88+ 
89+TEST_F(l2_relu_test, case_005_int32)
90+{
91+ auto selfDesc = TensorDesc({512}, ACL_INT32, ACL_FORMAT_ND).ValueRange(-20, 20);
92+ auto outDesc = TensorDesc({512}, ACL_INT32, ACL_FORMAT_ND);
93+ 
94+ auto ut = OP_API_UT(aclnnRelu, INPUT(selfDesc), OUTPUT(outDesc));
95+ 
96+ uint64_t workspaceSize = 0;
97+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
98+ EXPECT_EQ(aclRet, ACL_SUCCESS);
99+ 
100+ ut.TestPrecision();
101+}
102+ 
103+TEST_F(l2_relu_test, case_006_int64)
104+{
105+ auto selfDesc = TensorDesc({257}, ACL_INT64, ACL_FORMAT_ND).ValueRange(-20, 20);
106+ auto outDesc = TensorDesc({257}, ACL_INT64, ACL_FORMAT_ND);
107+ 
108+ auto ut = OP_API_UT(aclnnRelu, INPUT(selfDesc), OUTPUT(outDesc));
109+ 
110+ uint64_t workspaceSize = 0;
111+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
112+ EXPECT_EQ(aclRet, ACL_SUCCESS);
113+ 
114+ ut.TestPrecision();
115+}
116+ 
117+TEST_F(l2_relu_test, case_007_empty_tensor)
118+{
119+ auto selfDesc = TensorDesc({2, 0}, ACL_FLOAT, ACL_FORMAT_ND);
120+ auto outDesc = TensorDesc({2, 0}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.0001, 0.0001);
121+ 
122+ auto ut = OP_API_UT(aclnnRelu, INPUT(selfDesc), OUTPUT(outDesc));
123+ 
124+ uint64_t workspaceSize = 0;
125+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
126+ EXPECT_EQ(aclRet, ACL_SUCCESS);
127+ 
128+ ut.TestPrecision();
129+}
130+ 
131+TEST_F(l2_relu_test, case_008_not_contiguous)
132+{
133+ auto selfDesc = TensorDesc({5, 4}, ACL_FLOAT, ACL_FORMAT_ND, {1, 5}, 0, {4, 5}).ValueRange(-10, 10);
134+ auto outDesc = TensorDesc({5, 4}, ACL_FLOAT, ACL_FORMAT_ND, {1, 5}, 0, {4, 5}).Precision(0.0001, 0.0001);
135+ 
136+ auto ut = OP_API_UT(aclnnRelu, INPUT(selfDesc), 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_relu_test, case_009_invalid_input_dtype)
146+{
147+ auto selfDesc = TensorDesc({2, 3}, ACL_BOOL, ACL_FORMAT_ND);
148+ auto outDesc = TensorDesc({2, 3}, ACL_FLOAT, ACL_FORMAT_ND);
149+ 
150+ auto ut = OP_API_UT(aclnnRelu, INPUT(selfDesc), OUTPUT(outDesc));
151+ 
152+ uint64_t workspaceSize = 0;
153+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
154+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
155+}
156+ 
157+TEST_F(l2_relu_test, case_010_invalid_output_dtype)
158+{
159+ auto selfDesc = TensorDesc({2, 3}, ACL_FLOAT, ACL_FORMAT_ND);
160+ auto outDesc = TensorDesc({2, 3}, ACL_BOOL, ACL_FORMAT_ND);
161+ 
162+ auto ut = OP_API_UT(aclnnRelu, INPUT(selfDesc), OUTPUT(outDesc));
163+ 
164+ uint64_t workspaceSize = 0;
165+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
166+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
167+}
168+ 
169+TEST_F(l2_relu_test, case_011_dtype_mismatch)
170+{
171+ auto selfDesc = TensorDesc({2, 3}, ACL_FLOAT, ACL_FORMAT_ND);
172+ auto outDesc = TensorDesc({2, 3}, ACL_FLOAT16, ACL_FORMAT_ND);
173+ 
174+ auto ut = OP_API_UT(aclnnRelu, INPUT(selfDesc), OUTPUT(outDesc));
175+ 
176+ uint64_t workspaceSize = 0;
177+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
178+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
179+}
180+ 
181+TEST_F(l2_relu_test, case_012_nullptr)
182+{
183+ auto outDesc = TensorDesc({2, 4}, ACL_INT8, ACL_FORMAT_ND);
184+ 
185+ auto ut = OP_API_UT(aclnnRelu, INPUT(nullptr), OUTPUT(outDesc));
186+ 
187+ uint64_t workspaceSize = 0;
188+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
189+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_NULLPTR);
190+}
191+ 
192+TEST_F(l2_relu_test, case_013_shape_mismatch)
193+{
194+ auto selfDesc = TensorDesc({1, 2, 3, 4}, ACL_INT8, ACL_FORMAT_ND);
195+ auto outDesc = TensorDesc({1, 2, 3, 3}, ACL_INT8, ACL_FORMAT_ND);
196+ 
197+ auto ut = OP_API_UT(aclnnRelu, INPUT(selfDesc), OUTPUT(outDesc));
198+ 
199+ uint64_t workspaceSize = 0;
200+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
201+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
202+}
203+ 
204+TEST_F(l2_relu_test, case_014_max_dim)
205+{
206+ auto selfDesc = TensorDesc({1, 2, 3, 4, 5, 6, 7, 8, 9}, ACL_INT8, ACL_FORMAT_ND);
207+ auto outDesc = TensorDesc({1, 2, 3, 4, 5, 6, 7, 8, 9}, ACL_INT8, ACL_FORMAT_ND);
208+ 
209+ auto ut = OP_API_UT(aclnnRelu, INPUT(selfDesc), OUTPUT(outDesc));
210+ 
211+ uint64_t workspaceSize = 0;
212+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
213+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
214+}
215+ 
216+TEST_F(l2_relu_test, case_015_inplace_float)
217+{
218+ auto selfDesc = TensorDesc({2, 4}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-10, 10);
219+ 
220+ auto ut = OP_API_UT(aclnnInplaceRelu, INPUT(selfDesc), OUTPUT());
221+ 
222+ uint64_t workspaceSize = 0;
223+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
224+ EXPECT_EQ(aclRet, ACL_SUCCESS);
225+ 
226+ ut.TestPrecision();
227+}
228+ 
229+TEST_F(l2_relu_test, case_016_inplace_int64)
230+{
231+ auto selfDesc = TensorDesc({257}, ACL_INT64, ACL_FORMAT_ND).ValueRange(-20, 20);
232+ 
233+ auto ut = OP_API_UT(aclnnInplaceRelu, INPUT(selfDesc), OUTPUT());
234+ 
235+ uint64_t workspaceSize = 0;
236+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
237+ EXPECT_EQ(aclRet, ACL_SUCCESS);
238+ 
239+ ut.TestPrecision();
240+}
241+ 
242+TEST_F(l2_relu_test, case_017_inplace_invalid_dtype)
243+{
244+ auto selfDesc = TensorDesc({2, 3}, ACL_BOOL, ACL_FORMAT_ND);
245+ 
246+ auto ut = OP_API_UT(aclnnInplaceRelu, INPUT(selfDesc), OUTPUT());
247+ 
248+ uint64_t workspaceSize = 0;
249+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
250+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
251+}
252+ 
253+TEST_F(l2_relu_test, case_018_inplace_nullptr)
254+{
255+ uint64_t workspaceSize = 0;
256+ auto ut = OP_API_UT(aclnnInplaceRelu, INPUT(static_cast<aclTensor *>(nullptr)), OUTPUT());
257+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
258+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_NULLPTR);
259+}
260+ 
261+TEST_F(l2_relu_test, case_019_inplace_max_dim)
262+{
263+ auto selfDesc = TensorDesc({1, 2, 3, 4, 5, 6, 7, 8, 9}, ACL_INT8, ACL_FORMAT_ND);
264+ 
265+ auto ut = OP_API_UT(aclnnInplaceRelu, INPUT(selfDesc), OUTPUT());
266+ 
267+ uint64_t workspaceSize = 0;
268+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
269+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
270+}
271+ 
272+TEST_F(l2_relu_test, case_020_inplace_float16)
273+{
274+ auto selfDesc = TensorDesc({2, 4}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-10, 10);
275+ 
276+ auto ut = OP_API_UT(aclnnInplaceRelu, INPUT(selfDesc), OUTPUT());
277+ 
278+ uint64_t workspaceSize = 0;
279+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
280+ EXPECT_EQ(aclRet, ACL_SUCCESS);
281+ 
282+ ut.TestPrecision();
283+}
284+ 
285+TEST_F(l2_relu_test, case_021_inplace_bfloat16)
286+{
287+ auto selfDesc = TensorDesc({2, 4}, ACL_BF16, ACL_FORMAT_ND).ValueRange(-10, 10);
288+ 
289+ auto ut = OP_API_UT(aclnnInplaceRelu, INPUT(selfDesc), OUTPUT());
290+ 
291+ uint64_t workspaceSize = 0;
292+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
293+ EXPECT_EQ(aclRet, ACL_SUCCESS);
294+ 
295+ ut.TestPrecision();
296+}
297+ 
298+TEST_F(l2_relu_test, case_022_inplace_not_contiguous)
299+{
300+ auto selfDesc = TensorDesc({5, 4}, ACL_FLOAT, ACL_FORMAT_ND, {1, 5}, 0, {4, 5}).ValueRange(-10, 10);
301+ 
302+ auto ut = OP_API_UT(aclnnInplaceRelu, INPUT(selfDesc), OUTPUT());
303+ 
304+ uint64_t workspaceSize = 0;
305+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
306+ EXPECT_EQ(aclRet, ACL_SUCCESS);
307+ 
308+ ut.TestPrecision();
309+}
310+ 
311+TEST_F(l2_relu_test, case_023_inplace_empty_tensor)
312+{
313+ auto selfDesc = TensorDesc({2, 0}, ACL_FLOAT, ACL_FORMAT_ND);
314+ 
315+ auto ut = OP_API_UT(aclnnInplaceRelu, INPUT(selfDesc), OUTPUT());
316+ 
317+ uint64_t workspaceSize = 0;
318+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
319+ EXPECT_EQ(aclRet, ACL_SUCCESS);
320+ 
321+ ut.TestPrecision();
322+}
323+ 
324+TEST_F(l2_relu_test, case_024_inplace_int8)
325+{
326+ auto selfDesc = TensorDesc({1024}, ACL_INT8, ACL_FORMAT_ND).ValueRange(-20, 20);
327+ 
328+ auto ut = OP_API_UT(aclnnInplaceRelu, INPUT(selfDesc), OUTPUT());
329+ 
330+ uint64_t workspaceSize = 0;
331+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
332+ EXPECT_EQ(aclRet, ACL_SUCCESS);
333+ 
334+ ut.TestPrecision();
335+}
336+ 
337+TEST_F(l2_relu_test, case_025_inplace_int32)
338+{
339+ auto selfDesc = TensorDesc({512}, ACL_INT32, ACL_FORMAT_ND).ValueRange(-20, 20);
340+ 
341+ auto ut = OP_API_UT(aclnnInplaceRelu, INPUT(selfDesc), OUTPUT());
342+ 
343+ uint64_t workspaceSize = 0;
344+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
345+ EXPECT_EQ(aclRet, ACL_SUCCESS);
346+ 
347+ ut.TestPrecision();
348+}