Star859
1731
代码
代码
Issues170
Pull Requests306
Actions
项目讨论
Wiki
分析
Star859
1731
已合并
Swish算子Ascend C实现 #4020
fulltower创建于 4月20日
Swish算子Ascend C实现 #4020
已合并
fulltower创建于 4月20日
从fullt/ops-nnswish合入到cann/ops-nnmaster
共 28 个文件变更+2461-0
Aexperimental/activation/swish/CMakeLists.txt+19-0
@@ -0,0 +1,19 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------
10+ 
11+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+if(NOT ENABLE_TEST AND NOT BENCHMARK)
13+ list(REMOVE_ITEM CURRENT_DIRS tests)
14+endif()
15+foreach(SUB_DIR ${CURRENT_DIRS})
16+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
17+ add_subdirectory(${SUB_DIR})
18+ endif()
19+endforeach()
Aexperimental/activation/swish/README.md+51-0
@@ -0,0 +1,51 @@
1+# Swish
2+## 贡献说明
3+| 贡献者 | 贡献方 | 贡献算子 | 贡献时间 | 贡献内容 |
4+|--------|------------------|-------|-----------|-----------|
5+| skywang2 | 个人开发者 | Swish | 2025/12/31 | 新增Swish算子 |
6+ 
7+## 支持的产品型号
8+ 
9+- Atlas A2训练系列产品
10+ 
11+产品形态详细说明请参见[昇腾产品形态说明](http://www.hiascend.com/document/redirect/CannCommunityProductForm)
12+ 
13+## 算子描述
14+- 功能描述
15+ 
16+ `Swish`算子实现Swish激活函数，是一种由输入与其经过Sigmoid函数结果相乘得到的平滑、非线性函数。
17+ 
18+- 原型信息
19+ 
20+ <table>
21+ <tr><th align="center">算子类型(OpType)</th><th colspan="4" align="center">Swish</th></tr>
22+ <tr><td align="center"> </td><td align="center">name</td><td align="center">Type</td><td align="center">data type</td><td align="center">format</td></tr>
23+ <tr><td rowspan="2" align="center">算子输入</td>
24+
25+ <tr><td align="center">x</td><td align="center">tensor</td><td align="center">float32,float16,bfloat16</td><td align="center">ND</td></tr>
26+
27+ <tr><td rowspan="1" align="center">算子输出</td>
28+ <td align="center">y</td><td align="center">tensor</td><td align="center">float32,float16,bfloat16</td><td align="center">ND</td></tr>
29+ <tr><td rowspan="1" align="center">算子属性</td>
30+ <td align="center">scale</td><td align="center">scalar</td><td align="center">float</td><td align="center">-</td></tr>
31+ <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">swish</td></tr>
32+ </table>
33+ 
34+## 约束与限制
35+- x,y的数据类型仅支持float32,float16,bfloat1，数据格式仅支持ND
36+ 
37+### 运行验证
38+编译
39+```bash
40+bash build.sh --pkg --soc=ascend910b --experimental --ops=swish
41+```
42+运行
43+```bash
44+bash build.sh --run_example swish eager cust --vendor_name=custom --experimental
45+```
46+<table>
47+ <th>目录</th><th>描述</th>
48+ <tr>
49+ <td><a href="./examples/test_aclnn_swish.cpp">test_aclnn_swish.cpp</td><td>通过aclnn调用的方式调用Swish算子。</td>
50+ </tr>
51+</table>
Aexperimental/activation/swish/docs/aclnnSilu.md+327-0
@@ -0,0 +1,327 @@
1+# aclnnSilu
2+ 
3+## 产品支持情况
4+ 
5+|产品 | 是否支持 |
6+|:-------------------------|:----------:|
7+| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |
8+ 
9+## 功能说明
10+ 
11+- 接口功能：该算子也被称为Swish函数（$\beta = 1$），对输入tensor逐元素进行Silu函数运算并输出结果tensor。
12+ 
13+- 计算公式：
14+ 
15+ $$
16+ out_i = \frac {self_i} {1 + e^{-self_i}}
17+ $$
18+ 
19+## 函数原型
20+ 
21+ 每个算子分为[两段式接口](../../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnSiluGetWorkspaceSize”接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用“aclnnSilu”接口执行计算。
22+ 
23+```Cpp
24+aclnnStatus aclnnSiluGetWorkspaceSize(
25+ const aclTensor* self,
26+ aclTensor* out,
27+ uint64_t* workspaceSize,
28+ aclOpExecutor** executor)
29+```
30+ 
31+```Cpp
32+aclnnStatus aclnnSilu(
33+ void* workspace,
34+ uint64_t workspaceSize,
35+ aclOpExecutor* executor,
36+ aclrtStream stream)
37+```
38+ 
39+## aclnnSiluGetWorkspaceSize
40+ 
41+- **参数说明：**
42+ 
43+ <table style="undefined;table-layout: fixed; width: 1390px"><colgroup>
44+ <col style="width: 271px">
45+ <col style="width: 115px">
46+ <col style="width: 200px">
47+ <col style="width: 240px">
48+ <col style="width: 177px">
49+ <col style="width: 104px">
50+ <col style="width: 138px">
51+ <col style="width: 145px">
52+ </colgroup>
53+ <thead>
54+ <tr>
55+ <th>参数名</th>
56+ <th>输入/输出</th>
57+ <th>描述</th>
58+ <th>使用说明</th>
59+ <th>数据类型</th>
60+ <th>数据格式</th>
61+ <th>维度(shape)</th>
62+ <th>非连续Tensor</th>
63+ </tr></thead>
64+ <tbody>
65+ <tr>
66+ <td>self（aclTensor*）</td>
67+ <td>输入</td>
68+ <td>待进行Silu计算的入参，公式中的输入self。</td>
69+ <td><ul><li>支持空Tensor。</li><li>数据类型和shape需要与out一致。</li></ul></td>
70+ <td>FLOAT16、FLOAT、BFLOAT16</td>
71+ <td>ND</td>
72+ <td>0-8</td>
73+ <td>√</td>
74+ </tr>
75+ <tr>
76+ <td>out（aclTensor*）</td>
77+ <td>输出</td>
78+ <td>计算的出参，公式中的out。</td>
79+ <td>数据类型和shape需要与self一致。</td>
80+ <td>FLOAT16、FLOAT、BFLOAT16</td>
81+ <td>ND</td>
82+ <td>0-8</td>
83+ <td>√</td>
84+ </tr>
85+ <tr>
86+ <td>workspaceSize（uint64_t*）</td>
87+ <td>输出</td>
88+ <td>返回需要在Device侧申请的workspace大小。</td>
89+ <td>-</td>
90+ <td>-</td>
91+ <td>-</td>
92+ <td>-</td>
93+ <td>-</td>
94+ </tr>
95+ <tr>
96+ <td>executor（aclOpExecutor**）</td>
97+ <td>输出</td>
98+ <td>返回op执行器，包含了算子计算流程。</td>
99+ <td>-</td>
100+ <td>-</td>
101+ <td>-</td>
102+ <td>-</td>
103+ <td>-</td>
104+ </tr>
105+ </tbody>
106+ </table>
107+
108+ - <term>Atlas 推理系列产品</term>、<term>Atlas 训练系列产品</term>：数据类型支持FLOAT16、FLOAT。
109+ 
110+- **返回值：**
111+ 
112+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../../docs/zh/context/aclnn返回码.md)。
113+ 
114+ 第一段接口会完成入参校验，出现以下场景时报错：
115+ 
116+ <table style="undefined;table-layout: fixed;width: 979px"><colgroup>
117+ <col style="width: 272px">
118+ <col style="width: 103px">
119+ <col style="width: 604px">
120+ </colgroup>
121+ <thead>
122+ <tr>
123+ <th>返回码</th>
124+ <th>错误码</th>
125+ <th>描述</th>
126+ </tr>
127+ </thead>
128+ <tbody>
129+ <tr>
130+ <td>ACLNN_ERR_PARAM_NULLPTR</td>
131+ <td>161001</td>
132+ <td>传入的self或out是空指针。</td>
133+ </tr>
134+ <tr>
135+ <td rowspan="8">ACLNN_ERR_PARAM_INVALID</td>
136+ <td rowspan="8">161002</td>
137+ <td>self和out的数据类型不在支持的范围内。</td>
138+ </tr>
139+ <tr>
140+ <td>self和out的数据类型不一致。</td>
141+ </tr>
142+ <tr>
143+ <td>self和out的shape不一致。</td>
144+ </tr>
145+ </tbody></table>
146+ 
147+## aclnnSilu
148+ 
149+- **参数说明：**
150+ 
151+ <table style="undefined;table-layout: fixed; width: 953px"><colgroup>
152+ <col style="width: 173px">
153+ <col style="width: 112px">
154+ <col style="width: 668px">
155+ </colgroup>
156+ <thead>
157+ <tr>
158+ <th>参数名</th>
159+ <th>输入/输出</th>
160+ <th>描述</th>
161+ </tr></thead>
162+ <tbody>
163+ <tr>
164+ <td>workspace</td>
165+ <td>输入</td>
166+ <td>在Device侧申请的workspace内存地址。</td>
167+ </tr>
168+ <tr>
169+ <td>workspaceSize</td>
170+ <td>输入</td>
171+ <td>在Device侧申请的workspace大小，由第一段接口aclnnSiluGetWorkspaceSize获取。</td>
172+ </tr>
173+ <tr>
174+ <td>executor</td>
175+ <td>输入</td>
176+ <td>op执行器，包含了算子计算流程。</td>
177+ </tr>
178+ <tr>
179+ <td>stream</td>
180+ <td>输入</td>
181+ <td>指定执行任务的Stream。</td>
182+ </tr>
183+ </tbody>
184+ </table>
185+ 
186+- **返回值：**
187+ 
188+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../../docs/zh/context/aclnn返回码.md)。
189+ 
190+## 约束说明
191+ 
192+- 确定性计算：
193+ - aclnnSilu默认确定性实现。
194+ 
195+## 调用示例
196+ 
197+示例代码如下，仅供参考，具体编译和执行过程请参考[编译与运行样例](../../../../docs/zh/context/编译与运行样例.md)。
198+ 
199+```Cpp
200+#include <iostream>
201+#include <vector>
202+#include "acl/acl.h"
203+#include "aclnnop/aclnn_silu.h"
204+ 
205+#define CHECK_RET(cond, return_expr) \
206+ do { \
207+ if (!(cond)) { \
208+ return_expr; \
209+ } \
210+ } while (0)
211+ 
212+#define LOG_PRINT(message, ...) \
213+ do { \
214+ printf(message, ##__VA_ARGS__); \
215+ } while (0)
216+ 
217+int64_t GetShapeSize(const std::vector<int64_t>& shape) {
218+ int64_t shape_size = 1;
219+ for (auto i : shape) {
220+ shape_size *= i;
221+ }
222+ return shape_size;
223+}
224+ 
225+int Init(int32_t deviceId, aclrtStream* stream) {
226+ // 固定写法，资源初始化
227+ auto ret = aclInit(nullptr);
228+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
229+ ret = aclrtSetDevice(deviceId);
230+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
231+ ret = aclrtCreateStream(stream);
232+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
233+ return 0;
234+}
235+ 
236+template <typename T>
237+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
238+ aclDataType dataType, aclTensor** tensor) {
239+ auto size = GetShapeSize(shape) * sizeof(T);
240+ // 调用aclrtMalloc申请device侧内存
241+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
242+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
243+ 
244+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
245+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
246+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
247+ 
248+ // 计算连续tensor的strides
249+ std::vector<int64_t> strides(shape.size(), 1);
250+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
251+ strides[i] = shape[i + 1] * strides[i + 1];
252+ }
253+ 
254+ // 调用aclCreateTensor接口创建aclTensor
255+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
256+ shape.data(), shape.size(), *deviceAddr);
257+ return 0;
258+}
259+ 
260+int main() {
261+ // 1. （固定写法）device/stream初始化, 参考acl API手册
262+ // 根据自己的实际device填写deviceId
263+ int32_t deviceId = 0;
264+ aclrtStream stream;
265+ auto ret = Init(deviceId, &stream);
266+ // check根据自己的需要处理
267+ CHECK_RET(ret == 0, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
268+ // 2. 构造输入与输出，需要根据API的接口自定义构造
269+ std::vector<int64_t> selfShape = {4, 2};
270+ std::vector<int64_t> outShape = {4, 2};
271+ void* selfDeviceAddr = nullptr;
272+ void* outDeviceAddr = nullptr;
273+ aclTensor* self = nullptr;
274+ aclTensor* out = nullptr;
275+ std::vector<float> selfHostData = {0, 1, 2, 3, 4, 5, 6, 7};
276+ std::vector<float> outHostData = {0, 0, 0, 0, 0, 0, 0, 0};
277+ // 创建self aclTensor
278+ ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self);
279+ CHECK_RET(ret == ACL_SUCCESS, return ret);
280+ // 创建out aclTensor
281+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
282+ CHECK_RET(ret == ACL_SUCCESS, return ret);
283+ 
284+ // 3. 调用CANN算子库API，需要修改为具体的API
285+ uint64_t workspaceSize = 0;
286+ aclOpExecutor* executor;
287+ // 调用aclnnSilu第一段接口
288+ ret = aclnnSiluGetWorkspaceSize(self, out, &workspaceSize, &executor);
289+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSiluGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
290+ // 根据第一段接口计算出的workspaceSize申请device内存
291+ void* workspaceAddr = nullptr;
292+ if (workspaceSize > 0) {
293+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
294+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret;);
295+ }
296+ // 调用aclnnSilu第二段接口
297+ ret = aclnnSilu(workspaceAddr, workspaceSize, executor, stream);
298+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSilu failed. ERROR: %d\n", ret); return ret);
299+ // 4. （固定写法）同步等待任务执行结束
300+ ret = aclrtSynchronizeStream(stream);
301+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
302+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
303+ auto size = GetShapeSize(outShape);
304+ std::vector<float> resultData(size, 0);
305+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr, size * sizeof(float),
306+ ACL_MEMCPY_DEVICE_TO_HOST);
307+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
308+ for (int64_t i = 0; i < size; i++) {
309+ LOG_PRINT("result[%ld] is: %f\n", i, resultData[i]);
310+ }
311+ 
312+ // 6. 释放aclTensor和aclScalar，需要根据具体API的接口定义修改
313+ aclDestroyTensor(self);
314+ aclDestroyTensor(out);
315+ 
316+ // 7. 释放device资源，需要根据具体API的接口定义修改
317+ aclrtFree(selfDeviceAddr);
318+ aclrtFree(outDeviceAddr);
319+ if (workspaceSize > 0) {
320+ aclrtFree(workspaceAddr);
321+ }
322+ aclrtDestroyStream(stream);
323+ aclrtResetDevice(deviceId);
324+ aclFinalize();
325+ return 0;
326+}
327+```
Aexperimental/activation/swish/docs/aclnnSwish.md+353-0
@@ -0,0 +1,353 @@
1+# aclnnSwish
2+ 
3+## 产品支持情况
4+ 
5+|产品 | 是否支持 |
6+|:-------------------------|:----------:|
7+| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |
8+ 
9+## 功能说明
10+ 
11+- 接口功能：Swish激活函数，对输入Tensor逐元素进行Swish函数运算并输出结果Tensor。
12+ 
13+- 计算公式：
14+ 
15+ $$
16+ s(x) = x*\sigma(\beta x)
17+ $$
18+ 
19+ $$
20+ \sigma(x) = {\frac{1} {1+{e}^{-x}}}
21+ $$
22+ 
23+ 其中$\sigma(x)$为sigmoid函数。
24+ 
25+## 函数原型
26+ 
27+ 每个算子分为[两段式接口](../../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnSwishGetWorkspaceSize”接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用“aclnnSwish”接口执行计算。
28+ 
29+```Cpp
30+aclnnStatus aclnnSwishGetWorkspaceSize(
31+ const aclTensor* self,
32+ const aclScalar* betaOptional,
33+ aclTensor* out,
34+ uint64_t* workspaceSize,
35+ aclOpExecutor** executor)
36+```
37+ 
38+```Cpp
39+aclnnStatus aclnnSwish(
40+ void* workspace,
41+ uint64_t workspaceSize,
42+ aclOpExecutor* executor,
43+ aclrtStream stream)
44+```
45+ 
46+## aclnnSwishGetWorkspaceSize
47+ 
48+- **参数说明：**
49+ 
50+ <table style="undefined;table-layout: fixed; width: 1420px"><colgroup>
51+ <col style="width: 271px">
52+ <col style="width: 115px">
53+ <col style="width: 230px">
54+ <col style="width: 240px">
55+ <col style="width: 177px">
56+ <col style="width: 104px">
57+ <col style="width: 138px">
58+ <col style="width: 145px">
59+ </colgroup>
60+ <thead>
61+ <tr>
62+ <th>参数名</th>
63+ <th>输入/输出</th>
64+ <th>描述</th>
65+ <th>使用说明</th>
66+ <th>数据类型</th>
67+ <th>数据格式</th>
68+ <th>维度(shape)</th>
69+ <th>非连续Tensor</th>
70+ </tr></thead>
71+ <tbody>
72+ <tr>
73+ <td>self（aclTensor*）</td>
74+ <td>输入</td>
75+ <td>表示用于计算激活函数的张量，公式中的x。</td>
76+ <td><ul><li>支持空Tensor。</li><li>self的shape和数据类型与out的一致。</li></ul></td>
77+ <td>BFLOAT16、FLOAT16、FLOAT</td>
78+ <td>ND</td>
79+ <td>0-8</td>
80+ <td>√</td>
81+ </tr>
82+ <tr>
83+ <td>betaOptional（aclScalar*）</td>
84+ <td>输入</td>
85+ <td>表示可调节参数，用于控制Swish函数的形状和斜率的标量，公式中的β。</td>
86+ <td><ul><li>数据类型需要是可转换为FLOAT的数据类型（参见<a href="../../../../docs/zh/context/互转换关系.md" target="_blank">互转换关系</a>）。</li><li>当betaOptional为空指针时，接口以1.0进行计算。</li></ul></td>
87+ <td>-</td>
88+ <td>-</td>
89+ <td>-</td>
90+ <td>-</td>
91+ </tr>
92+ <tr>
93+ <td>out（aclTensor*）</td>
94+ <td>输出</td>
95+ <td>表示Swish函数的输出，公式中的s(x)。</td>
96+ <td><ul><li>支持空Tensor。</li><li>out的shape和数据类型与self的一致。</li></ul></td>
97+ <td>BFLOAT16、FLOAT16、FLOAT</td>
98+ <td>ND</td>
99+ <td>0-8</td>
100+ <td>√</td>
101+ </tr>
102+ <tr>
103+ <td>workspaceSize（uint64_t*）</td>
104+ <td>输出</td>
105+ <td>返回需要在Device侧申请的workspace大小。</td>
106+ <td>-</td>
107+ <td>-</td>
108+ <td>-</td>
109+ <td>-</td>
110+ <td>-</td>
111+ </tr>
112+ <tr>
113+ <td>executor（aclOpExecutor**）</td>
114+ <td>输出</td>
115+ <td>返回op执行器，包含了算子计算流程。</td>
116+ <td>-</td>
117+ <td>-</td>
118+ <td>-</td>
119+ <td>-</td>
120+ <td>-</td>
121+ </tr>
122+ </tbody>
123+ </table>
124+
125+ - <term>Atlas 推理系列产品</term>、<term>Atlas 训练系列产品</term>：数据类型支持FLOAT16、FLOAT。
126+ 
127+- **返回值：**
128+ 
129+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../../docs/zh/context/aclnn返回码.md)。
130+ 
131+ 第一段接口会完成入参校验，出现以下场景时报错：
132+ 
133+ <table style="undefined;table-layout: fixed;width: 979px"><colgroup>
134+ <col style="width: 272px">
135+ <col style="width: 103px">
136+ <col style="width: 604px">
137+ </colgroup>
138+ <thead>
139+ <tr>
140+ <th>返回码</th>
141+ <th>错误码</th>
142+ <th>描述</th>
143+ </tr>
144+ </thead>
145+ <tbody>
146+ <tr>
147+ <td>ACLNN_ERR_PARAM_NULLPTR</td>
148+ <td>161001</td>
149+ <td>传入的self或out是空指针。</td>
150+ </tr>
151+ <tr>
152+ <td rowspan="8">ACLNN_ERR_PARAM_INVALID</td>
153+ <td rowspan="8">161002</td>
154+ <td>self、betaOptional或out的数据类型不在支持的范围内。</td>
155+ </tr>
156+ <tr>
157+ <td>self和out的数据类型不一致。</td>
158+ </tr>
159+ <tr>
160+ <td>self和out的shape不一致。</td>
161+ </tr>
162+ </tbody></table>
163+ 
164+## aclnnSwish
165+ 
166+- **参数说明：**
167+ 
168+ <table style="undefined;table-layout: fixed; width: 953px"><colgroup>
169+ <col style="width: 173px">
170+ <col style="width: 112px">
171+ <col style="width: 668px">
172+ </colgroup>
173+ <thead>
174+ <tr>
175+ <th>参数名</th>
176+ <th>输入/输出</th>
177+ <th>描述</th>
178+ </tr></thead>
179+ <tbody>
180+ <tr>
181+ <td>workspace</td>
182+ <td>输入</td>
183+ <td>在Device侧申请的workspace内存地址。</td>
184+ </tr>
185+ <tr>
186+ <td>workspaceSize</td>
187+ <td>输入</td>
188+ <td>在Device侧申请的workspace大小，由第一段接口aclnnSwishGetWorkspaceSize获取。</td>
189+ </tr>
190+ <tr>
191+ <td>executor</td>
192+ <td>输入</td>
193+ <td>op执行器，包含了算子计算流程。</td>
194+ </tr>
195+ <tr>
196+ <td>stream</td>
197+ <td>输入</td>
198+ <td>指定执行任务的Stream。</td>
199+ </tr>
200+ </tbody>
201+ </table>
202+ 
203+- **返回值：**
204+ 
205+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../../docs/zh/context/aclnn返回码.md)。
206+ 
207+## 约束说明
208+ 
209+- 确定性计算：
210+ - aclnnSwish默认确定性实现。
211+ 
212+## 调用示例
213+ 
214+示例代码如下，仅供参考，具体编译和执行过程请参考[编译与运行样例](../../../../docs/zh/context/编译与运行样例.md)。
215+ 
216+```Cpp
217+#include <iostream>
218+#include <vector>
219+#include "acl/acl.h"
220+#include "aclnnop/aclnn_swish.h"
221+ 
222+#define CHECK_RET(cond, return_expr) \
223+ do { \
224+ if (!(cond)) { \
225+ return_expr; \
226+ } \
227+ } while (0)
228+ 
229+#define LOG_PRINT(message, ...) \
230+ do { \
231+ printf(message, ##__VA_ARGS__); \
232+ } while (0)
233+ 
234+int64_t GetShapeSize(const std::vector<int64_t>& shape) {
235+ int64_t shape_size = 1;
236+ for (auto i : shape) {
237+ shape_size *= i;
238+ }
239+ return shape_size;
240+}
241+ 
242+int Init(int32_t deviceId, aclrtStream* stream) {
243+ // 固定写法，资源初始化
244+ auto ret = aclInit(nullptr);
245+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
246+ ret = aclrtSetDevice(deviceId);
247+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
248+ ret = aclrtCreateStream(stream);
249+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
250+ return 0;
251+}
252+ 
253+template <typename T>
254+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
255+ aclDataType dataType, aclTensor** tensor) {
256+ auto size = GetShapeSize(shape) * sizeof(T);
257+ // 调用aclrtMalloc申请device侧内存
258+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
259+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
260+ 
261+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
262+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
263+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
264+ 
265+ // 计算连续tensor的strides
266+ std::vector<int64_t> strides(shape.size(), 1);
267+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
268+ strides[i] = shape[i + 1] * strides[i + 1];
269+ }
270+ 
271+ // 调用aclCreateTensor接口创建aclTensor
272+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
273+ shape.data(), shape.size(), *deviceAddr);
274+ return 0;
275+}
276+ 
277+int main() {
278+ // 1. （固定写法）device/stream初始化, 参考acl API手册
279+ // 根据自己的实际device填写deviceId
280+ int32_t deviceId = 0;
281+ aclrtStream stream;
282+ auto ret = Init(deviceId, &stream);
283+ CHECK_RET(ret == 0, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
284+ 
285+ // 2. 构造输入与输出，需要根据API的接口自定义构造
286+ std::vector<int64_t> selfShape = {4, 2};
287+ std::vector<int64_t> outShape = {4, 2};
288+ void* selfDeviceAddr = nullptr;
289+ void* outDeviceAddr = nullptr;
290+ aclTensor* self = nullptr;
291+ aclScalar* betaOptional = nullptr;
292+ aclTensor* out = nullptr;
293+ std::vector<float> selfHostData = {0, 1, 2, 3, 4, 5, 6, 7};
294+ std::vector<float> outHostData = {0, 0, 0, 0, 0, 0, 0, 0};
295+ float betaValue = 1.1f;
296+ 
297+ // 创建self aclTensor
298+ ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self);
299+ CHECK_RET(ret == ACL_SUCCESS, return ret);
300+ 
301+ // 创建betaOptional aclScalar
302+ betaOptional = aclCreateScalar(&betaValue, aclDataType::ACL_FLOAT);
303+ CHECK_RET(betaOptional != nullptr, return ret);
304+ 
305+ // 创建out aclTensor
306+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
307+ CHECK_RET(ret == ACL_SUCCESS, return ret);
308+ 
309+ // 3. 调用CANN算子库API，需要修改为具体的API
310+ uint64_t workspaceSize = 0;
311+ aclOpExecutor* executor;
312+ // 调用aclnnSwish第一段接口
313+ ret = aclnnSwishGetWorkspaceSize(self, betaOptional, out, &workspaceSize, &executor);
314+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSwishGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
315+ // 根据第一段接口计算出的workspaceSize申请device内存
316+ void* workspaceAddr = nullptr;
317+ if (workspaceSize > 0) {
318+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
319+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret;);
320+ }
321+ // 调用aclnnSwish第二段接口
322+ ret = aclnnSwish(workspaceAddr, workspaceSize, executor, stream);
323+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSwish failed. ERROR: %d\n", ret); return ret);
324+ // 4. （固定写法）同步等待任务执行结束
325+ ret = aclrtSynchronizeStream(stream);
326+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
327+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
328+ auto size = GetShapeSize(outShape);
329+ std::vector<float> resultData(size, 0);
330+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr, size * sizeof(float),
331+ ACL_MEMCPY_DEVICE_TO_HOST);
332+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
333+ for (int64_t i = 0; i < size; i++) {
334+ LOG_PRINT("result[%ld] is: %f\n", i, resultData[i]);
335+ }
336+ 
337+ // 6. 释放aclTensor和aclScalar，需要根据具体API的接口定义修改
338+ aclDestroyTensor(self);
339+ aclDestroyScalar(betaOptional);
340+ aclDestroyTensor(out);
341+ 
342+ // 7. 释放device资源，需要根据具体API的接口定义修改
343+ aclrtFree(selfDeviceAddr);
344+ aclrtFree(outDeviceAddr);
345+ if (workspaceSize > 0) {
346+ aclrtFree(workspaceAddr);
347+ }
348+ aclrtDestroyStream(stream);
349+ aclrtResetDevice(deviceId);
350+ aclFinalize();
351+ return 0;
352+}
353+```
Aexperimental/activation/swish/examples/test_aclnn_swish.cpp+141-0
@@ -0,0 +1,141 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+#include <iostream>
11+#include <vector>
12+#include "acl/acl.h"
13+#include "aclnnop/aclnn_swish.h"
14+ 
15+#define CHECK_RET(cond, return_expr) \
16+ do { \
17+ if (!(cond)) { \
18+ return_expr; \
19+ } \
20+ } while (0)
21+ 
22+#define LOG_PRINT(message, ...) \
23+ do { \
24+ printf(message, ##__VA_ARGS__); \
25+ } while (0)
26+ 
27+int64_t GetShapeSize(const std::vector<int64_t>& shape) {
28+ int64_t shape_size = 1;
29+ for (auto i : shape) {
30+ shape_size *= i;
31+ }
32+ return shape_size;
33+}
34+ 
35+int Init(int32_t deviceId, aclrtStream* stream) {
36+ // 固定写法，资源初始化
37+ auto ret = aclInit(nullptr);
38+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
39+ ret = aclrtSetDevice(deviceId);
40+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
41+ ret = aclrtCreateStream(stream);
42+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
43+ return 0;
44+}
45+ 
46+template <typename T>
47+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
48+ aclDataType dataType, aclTensor** tensor) {
49+ auto size = GetShapeSize(shape) * sizeof(T);
50+ // 调用aclrtMalloc申请device侧内存
51+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
52+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
53+ 
54+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
55+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
56+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
57+ 
58+ // 计算连续tensor的strides
59+ std::vector<int64_t> strides(shape.size(), 1);
60+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
61+ strides[i] = shape[i + 1] * strides[i + 1];
62+ }
63+ 
64+ // 调用aclCreateTensor接口创建aclTensor
65+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
66+ shape.data(), shape.size(), *deviceAddr);
67+ return 0;
68+}
69+ 
70+int main() {
71+ // 1. （固定写法）device/stream初始化, 参考acl API手册
72+ // 根据自己的实际device填写deviceId
73+ int32_t deviceId = 0;
74+ aclrtStream stream;
75+ auto ret = Init(deviceId, &stream);
76+ CHECK_RET(ret == 0, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
77+ 
78+ // 2. 构造输入与输出，需要根据API的接口自定义构造
79+ std::vector<int64_t> selfShape = {4, 2};
80+ std::vector<int64_t> outShape = {4, 2};
81+ void* selfDeviceAddr = nullptr;
82+ void* outDeviceAddr = nullptr;
83+ aclTensor* self = nullptr;
84+ aclTensor* out = nullptr;
85+ std::vector<float> selfHostData = {0, 1, 2, 3, 4, 5, 6, 7};
86+ std::vector<float> outHostData = {0, 0, 0, 0, 0, 0, 0, 0};
87+ float scaleValue = 1.1f;
88+ aclScalar* scale = aclCreateScalar(&scaleValue, aclDataType::ACL_FLOAT);
89+ 
90+ // 创建self aclTensor
91+ ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self);
92+ CHECK_RET(ret == ACL_SUCCESS, return ret);
93+ 
94+ // 创建out aclTensor
95+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
96+ CHECK_RET(ret == ACL_SUCCESS, return ret);
97+ 
98+ // 3. 调用CANN算子库API，需要修改为具体的API
99+ uint64_t workspaceSize = 0;
100+ aclOpExecutor* executor;
101+ // 调用aclnnSwish第一段接口
102+ ret = aclnnSwishGetWorkspaceSize(self, scale, out, &workspaceSize, &executor);
103+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSwishGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
104+ // 根据第一段接口计算出的workspaceSize申请device内存
105+ void* workspaceAddr = nullptr;
106+ if (workspaceSize > 0) {
107+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
108+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret;);
109+ }
110+ // 调用aclnnSwish第二段接口
111+ ret = aclnnSwish(workspaceAddr, workspaceSize, executor, stream);
112+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSwish failed. ERROR: %d\n", ret); return ret);
113+ // 4. （固定写法）同步等待任务执行结束
114+ ret = aclrtSynchronizeStream(stream);
115+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
116+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
117+ auto size = GetShapeSize(outShape);
118+ std::vector<float> resultData(size, 0);
119+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr, size * sizeof(float),
120+ ACL_MEMCPY_DEVICE_TO_HOST);
121+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
122+ for (int64_t i = 0; i < size; i++) {
123+ LOG_PRINT("aclnnSwish result[%ld] is: %f\n", i, resultData[i]);
124+ }
125+ 
126+ // 6. 释放aclTensor，需要根据具体API的接口定义修改
127+ aclDestroyTensor(self);
128+ aclDestroyTensor(out);
129+ aclDestroyScalar(scale);
130+ 
131+ // 7. 释放device资源，需要根据具体API的接口定义修改
132+ aclrtFree(selfDeviceAddr);
133+ aclrtFree(outDeviceAddr);
134+ if (workspaceSize > 0) {
135+ aclrtFree(workspaceAddr);
136+ }
137+ aclrtDestroyStream(stream);
138+ aclrtResetDevice(deviceId);
139+ aclFinalize();
140+ return 0;
141+}
Aexperimental/activation/swish/op_graph/swish_proto.h+44-0
@@ -0,0 +1,44 @@
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+/*!
CANN-robot
CANN-robotCANN-robot4月20日
代码结构与可维护性: 文件头注释中的文件名与实际文件名不一致。注释中描述的文件名为'nonlinear_fuc_ops.h'，但实际文件名为'swish_proto.h'。这种不一致会导致开发者在查找文件时产生困惑，降低代码的可读性和可维护性。
问题类型: 代码结构与可维护性
文件路径: experimental/activation/swish/op_graph/swish_proto.h
行号: 12
问题代码:
/*!
 * \file nonlinear_fuc_ops.h
 * \brief
 */
修改建议:
将文件头注释中的文件名修改为实际文件名'swish_proto.h'，保持注释与代码的一致性。
---
此评论由代码审查工具自动生成
likedislike
13+ * \file swish_proto.h
14+ * \brief use for graph
15+ */
16+#ifndef OPS_ACTIVATION_SWISH_OPS_H_
17+#define OPS_ACTIVATION_SWISH_OPS_H_
18+ 
19+#include "graph/operator_reg.h"
20+ 
21+namespace ge {
22+/**
23+*@brief Computes the Swish activation function for input tensor x .
24+ 
25+*@par Inputs:
26+*One input, including:
27+* x: A tensor, which supports 1D-8D defaultly and must be one of the following types: float16, bfloat16, float32. \n
28+ 
29+*@par Outputs:
30+* y: A tensor of the same type, shape and format as "x", and y = x / (1 + exp(-scale * x)). \n
31+ 
32+*@par Attributes:
33+* scale: scalar parameter, the multiplier of x. Must be of type float. Default value = 1.0. \n
34+ 
35+*@par Third-party framework compatibility
36+*Compatible with the Torch operator Swish
37+*/
38+REG_OP(Swish)
39+ .INPUT(x, TensorType({DT_FLOAT16, DT_FLOAT, DT_BF16}))
40+ .OUTPUT(y, TensorType({DT_FLOAT16, DT_FLOAT, DT_BF16}))
41+ .ATTR(scale, Float, 1.0)
42+ .OP_END_FACTORY_REG(Swish)
43+} // namespace ge
44+#endif // OPS_ACTIVATION_SWISH_OPS_H_
Aexperimental/activation/swish/op_host/CMakeLists.txt+10-0
@@ -0,0 +1,10 @@
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+add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE swish ACLNNTYPE aclnn_exclude)
Aexperimental/activation/swish/op_host/op_api/aclnn_silu.cpp+120-0
@@ -0,0 +1,120 @@
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_silu.h"
12+ 
13+#include "aclnn/aclnn_base.h"
14+#include "aclnn_kernels/cast.h"
15+#include "aclnn_kernels/contiguous.h"
16+#include "aclnn_kernels/reshape.h"
17+#include "swish.h"
18+#include "opdev/common_types.h"
19+#include "opdev/data_type_utils.h"
20+#include "opdev/format_utils.h"
21+#include "opdev/op_dfx.h"
22+#include "opdev/op_executor.h"
23+#include "opdev/op_log.h"
24+#include "opdev/platform.h"
25+#include "opdev/shape_utils.h"
26+#include "opdev/tensor_view_utils.h"
27+#include "op_api/op_api_def.h"
28+#include "op_api/level2_base_caculation.h"
29+ 
30+using namespace op;
31+ 
32+#ifdef __cplusplus
33+extern "C" {
34+#endif
35+ 
36+static const std::initializer_list<op::DataType> ASCEND910_DTYPE_SUPPORT_LIST = {
37+ op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16};
38+ 
39+static const std::initializer_list<op::DataType> ASCEND910B_DTYPE_SUPPORT_LIST = {
40+ op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16, op::DataType::DT_BF16};
41+ 
42+static inline bool CheckNotNull(const aclTensor* self, const aclTensor* out)
43+{
44+ OP_CHECK_NULL(self, return false);
45+ OP_CHECK_NULL(out, return false);
46+ return true;
47+}
48+ 
49+static aclnnStatus CheckParams(const aclTensor *self, const aclTensor *out) {
50+ CHECK_RET(CheckNotNull(self, out), ACLNN_ERR_PARAM_NULLPTR);
51+ 
52+ auto supportList = GetDtypeSupportListV2(ASCEND910B_DTYPE_SUPPORT_LIST, ASCEND910_DTYPE_SUPPORT_LIST);
53+ CHECK_RET(CheckDtypeValidActivation(self, out, supportList), ACLNN_ERR_PARAM_INVALID);
54+ 
55+ CHECK_RET(CheckSameShapeNotlimit1In1Out(self, out), ACLNN_ERR_PARAM_INVALID);
56+ 
57+ return ACLNN_SUCCESS;
58+}
59+ 
60+aclnnStatus aclnnSiluGetWorkspaceSize(const aclTensor* self, aclTensor* out, uint64_t* workspaceSize,
61+ aclOpExecutor** executor) {
62+ OP_CHECK_COMM_INPUT(workspaceSize, executor);
63+ 
64+ L2_DFX_PHASE_1(aclnnSilu, DFX_IN(self), DFX_OUT(out));
65+ 
66+ auto uniqueExecutor = CREATE_EXECUTOR();
67+ CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR);
68+ 
69+ auto ret = CheckParams(self, out);
70+ CHECK_RET(ret == ACLNN_SUCCESS, ret);
71+ 
72+ if (self->IsEmpty() || out->IsEmpty()) {
73+ *workspaceSize = 0;
74+ uniqueExecutor.ReleaseTo(executor);
75+ return ACLNN_SUCCESS;
76+ }
77+ 
78+ auto selfContiguous = l0op::Contiguous(self, uniqueExecutor.get());
79+ CHECK_RET(selfContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
80+ 
81+ auto shapeOri = self->GetViewShape();
82+ int64_t dimSize = self->GetViewShape().GetDimNum();
83+ auto shapeOriDetail = GetTensorShapeActivation(selfContiguous, uniqueExecutor.get());
84+ auto reshapeSelf = selfContiguous;
85+ 
86+ if (dimSize > (int64_t)MAX_SUPPORT_DIMS_NUMS) {
87+ int64_t allDimValue = 1;
88+ for (int i = 0; i < dimSize; i++) {
89+ allDimValue *= shapeOri[i];
90+ }
91+ int64_t AllDim[1] = {allDimValue};
92+ auto shape1d = (uniqueExecutor)->AllocIntArray(AllDim, 1);
93+ reshapeSelf = ReshapeLongTensorActivation(selfContiguous, uniqueExecutor.get(), dimSize, shape1d);
94+ }
95+ 
96+ float scale = 1.0;
97+ auto siluOut = l0op::Swish(reshapeSelf, scale, uniqueExecutor.get());
98+ CHECK_RET(siluOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
99+ 
100+ auto reshapeSiluOut = siluOut;
101+ if (dimSize > (int64_t)MAX_SUPPORT_DIMS_NUMS) {
102+ reshapeSiluOut = ReshapeLongTensorActivation(siluOut, uniqueExecutor.get(), dimSize, shapeOriDetail);
103+ }
104+ 
105+ auto viewCopyOut = l0op::ViewCopy(reshapeSiluOut, out, uniqueExecutor.get());
106+ CHECK_RET(viewCopyOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
107+ 
108+ *workspaceSize = uniqueExecutor->GetWorkspaceSize();
109+ uniqueExecutor.ReleaseTo(executor);
110+ return ACLNN_SUCCESS;
111+}
112+ 
113+aclnnStatus aclnnSilu(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream) {
114+ L2_DFX_PHASE_2(aclnnSilu);
115+ return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);
116+}
117+ 
118+#ifdef __cplusplus
119+}
120+#endif
Aexperimental/activation/swish/op_host/op_api/aclnn_silu.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 OP_API_INC_LEVEL2_ACLNN_SILU_H_
12+#define OP_API_INC_LEVEL2_ACLNN_SILU_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 aclnnSiluGetWorkspaceSize(const aclTensor* self, aclTensor* out, uint64_t* workspaceSize,
22+ aclOpExecutor** executor);
23+ 
24+ACLNN_API aclnnStatus aclnnSilu(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor,
25+ aclrtStream stream);
26+ 
27+#ifdef __cplusplus
28+}
29+#endif
30+ 
31+#endif
Aexperimental/activation/swish/op_host/op_api/aclnn_swish.cpp+134-0
@@ -0,0 +1,134 @@
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_swish.h"
12+ 
13+#include "aclnn/aclnn_base.h"
14+#include "aclnn_kernels/cast.h"
15+#include "aclnn_kernels/contiguous.h"
16+#include "aclnn_kernels/reshape.h"
17+#include "swish.h"
18+#include "opdev/common_types.h"
19+#include "opdev/data_type_utils.h"
20+#include "opdev/format_utils.h"
21+#include "opdev/op_dfx.h"
22+#include "opdev/op_executor.h"
23+#include "opdev/op_log.h"
24+#include "opdev/platform.h"
25+#include "opdev/shape_utils.h"
26+#include "opdev/tensor_view_utils.h"
27+#include "op_api/op_api_def.h"
28+#include "op_api/level2_base_caculation.h"
29+ 
30+using namespace op;
31+ 
32+#ifdef __cplusplus
33+extern "C" {
34+#endif
35+ 
36+static const std::initializer_list<op::DataType> ASCEND910_DTYPE_SUPPORT_LIST = {
37+ op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16};
38+ 
39+static const std::initializer_list<op::DataType> ASCEND910B_DTYPE_SUPPORT_LIST = {
40+ op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16, op::DataType::DT_BF16};
41+ 
42+static bool CheckDtypeValidBetaToFloat(const aclScalar* betaOptional) {
43+ if (betaOptional != nullptr && !CanCast(betaOptional->GetDataType(), DataType::DT_FLOAT)) {
44+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "betaOptional dtype %s can not cast to float32.",
45+ ToString(betaOptional->GetDataType()).GetString());
46+ return false;
47+ }
48+ return true;
49+}
50+ 
51+static bool CheckDim(const aclTensor *self, const aclTensor *out) {
52+ OP_CHECK_MAX_DIM(self, MAX_SUPPORT_DIMS_NUMS, return false);
53+ OP_CHECK_MAX_DIM(out, MAX_SUPPORT_DIMS_NUMS, return false);
54+ return true;
55+}
56+ 
57+static aclnnStatus CheckParams(const aclTensor *self, const aclScalar* betaOptional, const aclTensor *out) {
58+ CHECK_RET(CheckNotNull2Tensor(self, out), ACLNN_ERR_PARAM_NULLPTR);
59+ 
60+ auto supportList = GetDtypeSupportListV1(ASCEND910B_DTYPE_SUPPORT_LIST, ASCEND910_DTYPE_SUPPORT_LIST);
61+ CHECK_RET(CheckDtypeValidActivation(self, out, supportList), ACLNN_ERR_PARAM_INVALID);
62+ CHECK_RET(CheckDtypeValidBetaToFloat(betaOptional), ACLNN_ERR_PARAM_INVALID);
63+ 
64+ CHECK_RET(CheckDim(self, out), ACLNN_ERR_PARAM_INVALID);
65+ CHECK_RET(CheckSameShapeNotlimit1In1Out(self, out), ACLNN_ERR_PARAM_INVALID);
66+ 
67+ return ACLNN_SUCCESS;
68+}
69+ 
70+static const aclTensor *reshapeLongTensor(const aclTensor *x, aclOpExecutor *executor, size_t originalDimSize,
71+ aclIntArray *valuePerm = nullptr) {
72+ size_t dimSize = x->GetViewShape().GetDimNum();
73+ if (originalDimSize == dimSize && dimSize <= MAX_SUPPORT_DIMS_NUMS) {
74+ return x;
75+ }
76+ 
77+ auto reshapeSelf = l0op::Reshape(x, valuePerm, executor);
78+ return reshapeSelf;
79+}
80+ 
81+aclnnStatus aclnnSwishGetWorkspaceSize(const aclTensor* self, const aclScalar* betaOptional, aclTensor* out, uint64_t* workspaceSize,
82+ aclOpExecutor** executor) {
83+ OP_CHECK_COMM_INPUT(workspaceSize, executor);
84+ 
85+ L2_DFX_PHASE_1(aclnnSwish, DFX_IN(self, betaOptional), DFX_OUT(out));
86+ 
87+ auto uniqueExecutor = CREATE_EXECUTOR();
88+ CHECK_RET(uniqueExecutor.get() != nullptr, ACLNN_ERR_INNER_CREATE_EXECUTOR);
89+ 
90+ auto ret = CheckParams(self, betaOptional, out);
91+ CHECK_RET(ret == ACLNN_SUCCESS, ret);
92+ 
93+ if (self->IsEmpty() || out->IsEmpty()) {
94+ *workspaceSize = 0;
95+ uniqueExecutor.ReleaseTo(executor);
96+ return ACLNN_SUCCESS;
97+ }
98+ 
99+ auto selfContiguous = l0op::Contiguous(self, uniqueExecutor.get());
100+ CHECK_RET(selfContiguous != nullptr, ACLNN_ERR_INNER_NULLPTR);
101+ 
102+ size_t dimSize = self->GetViewShape().GetDimNum();
103+ auto shapeOriDetial = GetTensorShapeActivation(selfContiguous, uniqueExecutor.get());
104+ auto reshapeSelf = ReshapeSelfValueGetActivation(self, dimSize, selfContiguous, uniqueExecutor);
105+ 
106+ float scale = 1.0f;
107+ if (betaOptional != nullptr) {
108+ scale = betaOptional->ToFloat();
109+ }
110+ 
111+ auto swishOut = l0op::Swish(reshapeSelf, scale, uniqueExecutor.get());
112+ CHECK_RET(swishOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
113+ 
114+ auto reshapeSwishOut = swishOut;
115+ if (dimSize > MAX_SUPPORT_DIMS_NUMS) {
116+ reshapeSwishOut = reshapeLongTensor(swishOut, uniqueExecutor.get(), dimSize, shapeOriDetial);
117+ }
118+ 
119+ auto viewCopyOut = l0op::ViewCopy(reshapeSwishOut, out, uniqueExecutor.get());
120+ CHECK_RET(viewCopyOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
121+ 
122+ *workspaceSize = uniqueExecutor->GetWorkspaceSize();
123+ uniqueExecutor.ReleaseTo(executor);
124+ return ACLNN_SUCCESS;
125+}
126+ 
127+aclnnStatus aclnnSwish(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream) {
128+ L2_DFX_PHASE_2(aclnnSwish);
129+ return CommonOpExecutorRun(workspace, workspaceSize, executor, stream);
130+}
131+ 
132+#ifdef __cplusplus
133+}
134+#endif
Aexperimental/activation/swish/op_host/op_api/aclnn_swish.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 OP_API_INC_LEVEL2_ACLNN_SWISH_H_
12+#define OP_API_INC_LEVEL2_ACLNN_SWISH_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 aclnnSwishGetWorkspaceSize(const aclTensor* self, const aclScalar* betaOptional, aclTensor* out,
22+ uint64_t* workspaceSize, aclOpExecutor** executor);
23+ 
24+ACLNN_API aclnnStatus aclnnSwish(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor,
25+ aclrtStream stream);
26+ 
27+#ifdef __cplusplus
28+}
29+#endif
30+ 
31+#endif
Aexperimental/activation/swish/op_host/op_api/silu.cpp+35-0
@@ -0,0 +1,35 @@
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 "swish.h"
12+#include "opdev/make_op_executor.h"
13+#include "opdev/op_def.h"
14+#include "opdev/op_dfx.h"
15+#include "opdev/op_executor.h"
16+#include "opdev/op_log.h"
17+#include "opdev/shape_utils.h"
18+ 
19+using namespace op;
20+namespace l0op {
21+OP_TYPE_REGISTER(Swish);
22+ 
23+static const aclTensor *SwishAiCore(const aclTensor *x, aclTensor *y, float scale, aclOpExecutor *executor) {
24+ L0_DFX(SwishAiCore, x, y);
25+ auto ret = ADD_TO_LAUNCHER_LIST_AICORE(Swish, OP_INPUT(x), OP_OUTPUT(y), OP_ATTR(scale));
26+ OP_CHECK(ret == ACLNN_SUCCESS, OP_LOGE(ACLNN_ERR_INNER_NULLPTR, "SwishAiCore ADD_TO_LAUNCHER_LIST_AICORE failed."),
27+ return nullptr);
28+ return y;
29+}
30+ 
31+const aclTensor *Swish(const aclTensor *x, float scale, aclOpExecutor *executor) {
32+ auto out = executor->AllocTensor(x->GetViewShape(), x->GetDataType());
33+ return SwishAiCore(x, out, scale, executor);
34+}
35+} // namespace l0op
Aexperimental/activation/swish/op_host/op_api/swish.h+20-0
@@ -0,0 +1,20 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#ifndef OP_API_INC_LEVEL0_SWISH_H_
12+#define OP_API_INC_LEVEL0_SWISH_H_
13+ 
14+#include "opdev/op_executor.h"
15+ 
16+namespace l0op {
17+const aclTensor *Swish(const aclTensor *self, float scale, aclOpExecutor *executor);
18+} // namespace l0op
19+ 
20+#endif
Aexperimental/activation/swish/op_host/swish_def.cpp+46-0
@@ -0,0 +1,46 @@
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 swish.cpp
13+ * \brief
14+ */
15+#include "register/op_def_registry.h"
16+ 
17+namespace ops {
18+class Swish : public OpDef {
19+public:
20+ explicit Swish(const char* name) : OpDef(name)
21+ {
22+ this->Input("x")
23+ .ParamType(REQUIRED)
24+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})
25+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
26+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
27+ this->Output("y")
28+ .ParamType(REQUIRED)
29+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})
30+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
31+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
32+ this->Attr("scale").AttrType(OPTIONAL).Float(1.0);
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", "swish");
42+ this->AICore().AddConfig("ascend910b", aicoreConfig);
43+ }
44+};
45+OP_ADD(Swish);
46+} // namespace ops
Aexperimental/activation/swish/op_host/swish_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+/*!
12+ * \file swish_infer.cpp
13+ * \brief
14+ */
15+#include "register/op_impl_registry.h"
16+#include "log/log.h"
17+ 
18+using namespace ge;
19+ 
20+namespace ops {
21+ 
22+static ge::graphStatus InferShapeSwish(gert::InferShapeContext* context)
23+{
24+ const gert::Shape* x1_shape = context->GetInputShape(0);
25+ gert::Shape* y_shape = context->GetOutputShape(0);
26+ *y_shape = *x1_shape;
27+ return GRAPH_SUCCESS;
28+}
29+ 
30+IMPL_OP_INFERSHAPE(Swish).InferShape(InferShapeSwish);
31+} // namespace ops
Aexperimental/activation/swish/op_host/swish_tiling.cpp+142-0
@@ -0,0 +1,142 @@
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 swish_tiling.cpp
13+ * \brief
14+ */
15+ 
16+#include "log/log.h"
17+#include "util/math_util.h"
18+#include "util/platform_util.h"
19+#include "op_host/tiling_util.h"
20+#include "op_host/tiling_templates_registry.h"
21+#include "tiling/platform/platform_ascendc.h"
22+#include "register/op_impl_registry.h"
23+#include "swish/op_kernel/swish_tiling_data.h"
24+#include "swish/op_kernel/swish_tiling_key.h"
25+ 
26+namespace optiling {
27+ 
28+struct SwishCompileInfo {};
29+ 
30+const uint64_t BUFFER_NUM = 2;
31+ 
32+// tiling 分发入口
33+static ge::graphStatus SwishTilingFunc(gert::TilingContext* context)
34+{
35+ uint64_t blockSize = 0;
36+ uint64_t ubSize = 0;
37+ 
38+ blockSize = Ops::Base::GetUbBlockSize(context);
39+ OP_CHECK_IF(blockSize == 0, OP_LOGE(context, "blockSize is 0"), return ge::GRAPH_FAILED);
40+ 
41+ SwishTilingData* tiling = context->GetTilingData<SwishTilingData>();
42+ fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo();
43+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr);
44+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
45+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
46+ auto coreNum = ascendcPlatform.GetCoreNum();
47+ auto socVersion = ascendcPlatform.GetSocVersion();
48+ if (socVersion != platform_ascendc::SocVersion::ASCEND910B && socVersion != platform_ascendc::SocVersion::ASCEND310B && context->GetInputDesc(0)->GetDataType() == ge::DT_BF16) {
49+ OP_LOGE(context, "socVersion error.");
50+ return ge::GRAPH_FAILED;
51+ }
52+ 
53+ auto inputShape = context->GetInputShape(0);
54+ OP_CHECK_NULL_WITH_CONTEXT(context, inputShape);
55+ uint64_t inputNum = inputShape->GetStorageShape().GetShapeSize();
56+ OP_CHECK_IF(inputNum == 0, OP_LOGE(context, "inputNum is 0"), return ge::GRAPH_FAILED);
57+ 
58+ auto inputDesc = context->GetInputDesc(0);
59+ OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc);
60+ ge::DataType dataType = inputDesc->GetDataType();
61+ uint32_t typeLength = 0;
62+ ge::TypeUtils::GetDataTypeLength(dataType, typeLength);
63+ 
64+ uint64_t inputLength = inputNum * typeLength;
65+ uint64_t inputBytes = inputLength / inputNum;
66+ OP_CHECK_IF(inputNum == 0, OP_LOGE(context, "inputNum is 0"), return ge::GRAPH_FAILED);
67+ 
68+ uint64_t ubDataNumber = (dataType == ge::DT_FLOAT) ? 2 : 4;
69+ uint64_t tileBlockNum = (ubSize / blockSize / BUFFER_NUM) / ubDataNumber;
70+ uint64_t tileDataNum = (tileBlockNum * blockSize) / inputBytes;
71+ 
72+ uint64_t inputLengthAlgin32 = (((inputLength + blockSize - 1) / blockSize) * blockSize);
73+ if(tileDataNum >= inputNum)
74+ {
75+ coreNum=1;
76+ }
77+ else
78+ {
79+ // There is at least 32B of data on each core, satisfying several settings for several cores. The maximum number of audits is the actual number of audits
80+ coreNum = (coreNum < inputLengthAlgin32 / blockSize) ? coreNum : inputLengthAlgin32 / blockSize;
81+ }
82+ OP_CHECK_IF(coreNum == 0, OP_LOGE(context, "coreNum is 0"), return ge::GRAPH_FAILED);
83+ OP_CHECK_IF(inputBytes == 0, OP_LOGE(context, "inputBytes is 0"), return ge::GRAPH_FAILED);
84+ 
85+ uint64_t everyCoreInputBlockNum = inputLengthAlgin32 / blockSize / coreNum;
86+ uint64_t tailBlockNum = (inputLengthAlgin32 / blockSize) % coreNum;
87+
88+ uint64_t smallCoreDataNum = everyCoreInputBlockNum * blockSize / inputBytes;
89+ uint64_t smallTileNum = everyCoreInputBlockNum / tileBlockNum;
90+ uint64_t finalSmallTileNum = (everyCoreInputBlockNum % tileBlockNum) == 0 ? smallTileNum : smallTileNum + 1;
91+ uint64_t smallTailDataNum = smallCoreDataNum - (tileDataNum * smallTileNum);
92+ smallTailDataNum = smallTailDataNum == 0 ? tileDataNum : smallTailDataNum;
93+
94+ everyCoreInputBlockNum += 1;
95+ uint64_t bigCoreDataNum = everyCoreInputBlockNum * blockSize / inputBytes;
96+ uint64_t bigTileNum = everyCoreInputBlockNum / tileBlockNum;
97+ uint64_t finalBigTileNum = (everyCoreInputBlockNum % tileBlockNum) == 0 ? bigTileNum : bigTileNum + 1;
98+ uint64_t bigTailDataNum = bigCoreDataNum - tileDataNum * bigTileNum;
99+ bigTailDataNum = bigTailDataNum == 0 ? tileDataNum : bigTailDataNum;
100+
101+ tiling->smallCoreDataNum = (uint32_t)smallCoreDataNum;
102+ tiling->bigCoreDataNum = (uint32_t)bigCoreDataNum;
103+ tiling->tileDataNum = (uint32_t)tileDataNum;
104+ tiling->smallTailDataNum = (uint32_t)smallTailDataNum;
105+ tiling->bigTailDataNum = (uint32_t)bigTailDataNum;
106+ tiling->finalSmallTileNum = (uint32_t)finalSmallTileNum;
107+ tiling->finalBigTileNum = (uint32_t)finalBigTileNum;
108+ tiling->tailBlockNum = (uint32_t)tailBlockNum;
109+ 
110+ auto attrs = context->GetAttrs();
111+ OP_CHECK_NULL_WITH_CONTEXT(context, attrs);
112+ const float* scaleValueAttr = attrs->GetAttrPointer<float>(0);
113+ float scale = scaleValueAttr == nullptr ? 1.0f : *scaleValueAttr;
114+ tiling->scale = scale;
115+ 
116+ uint64_t attrWork = TPL_SCALE_OTHER;
117+ constexpr float NEG_ONE = -1.0f;
118+ constexpr float ZERO = 0.0f;
119+ if (scale == NEG_ONE) {
120+ attrWork = TPL_SCALE_NEG_ONE;
121+ } else if (scale == ZERO) {
122+ attrWork = TPL_SCALE_ZERO;
123+ } else {
124+ attrWork = TPL_SCALE_OTHER;
125+ }
126+ 
127+ context->SetBlockDim(coreNum);
128+ context->SetTilingKey(GET_TPL_TILING_KEY(TPL_SCH_MODE_0, attrWork));
129+ uint32_t sysWorkspaceSize = ascendcPlatform.GetLibApiWorkSpaceSize();
130+ size_t *currentWorkspace = context->GetWorkspaceSizes(1);
131+ currentWorkspace[0] = sysWorkspaceSize;
132+ return ge::GRAPH_SUCCESS;
133+}
134+ 
135+// tiling注册入口.
136+static ge::graphStatus TilingPrepare4Swish([[maybe_unused]] gert::TilingParseContext* context)
137+{
138+ return ge::GRAPH_SUCCESS;
139+}
140+ 
141+IMPL_OP_OPTILING(Swish).Tiling(SwishTilingFunc).TilingParse<SwishCompileInfo>(TilingPrepare4Swish);
142+} // namespace optiling
Aexperimental/activation/swish/op_kernel/swish.cpp+33-0
@@ -0,0 +1,33 @@
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 swish.cpp
13+ * \brief
14+ */
15+ 
16+#include "swish.h"
17+ 
18+using namespace NsSwish;
19+ 
20+template <uint64_t schMode, uint64_t attrWork>
21+__global__ __aicore__ void swish(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
22+{
23+ REGISTER_TILING_DEFAULT(SwishTilingData);
24+ GET_TILING_DATA_WITH_STRUCT(SwishTilingData, tiling_data, tiling);
25+ 
26+ KernelSwish<DTYPE_X, DTYPE_X> op;
27+ op.Init(x, tiling_data.scale, y, tiling_data.smallCoreDataNum,
28+ tiling_data.bigCoreDataNum, tiling_data.finalBigTileNum,
29+ tiling_data.finalSmallTileNum, tiling_data.tileDataNum,
30+ tiling_data.smallTailDataNum, tiling_data.bigTailDataNum,
31+ tiling_data.tailBlockNum);
32+ op.Process();
33+}
Aexperimental/activation/swish/op_kernel/swish.h+139-0
@@ -0,0 +1,139 @@
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 swish.h
13+ * \brief
14+ */
15+#ifndef __SWISH_H__
16+#define __SWISH_H__
17+ 
18+#include "kernel_operator.h"
19+#include "kernel_tiling/kernel_tiling.h"
20+#include "swish_tiling_data.h"
21+#include "swish_tiling_key.h"
22+ 
23+namespace NsSwish {
24+ 
25+using namespace AscendC;
26+ 
27+constexpr int32_t BUFFER_NUM = 2;
28+ 
29+template <typename TYPE_X, typename TYPE_Y>
30+class KernelSwish
31+{
32+public:
33+ __aicore__ inline KernelSwish() {}
34+ __aicore__ inline void Init(GM_ADDR x, float scale, GM_ADDR y, uint32_t smallCoreDataNum,
35+ uint32_t bigCoreDataNum, uint32_t finalBigTileNum,
36+ uint32_t finalSmallTileNum, uint32_t tileDataNum,
37+ uint32_t smallTailDataNum, uint32_t bigTailDataNum,
38+ uint32_t tailBlockNum)
39+ {
40+ ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");
41+ uint32_t coreNum = GetBlockIdx();
42+ uint32_t globalBufferIndex = bigCoreDataNum * GetBlockIdx();
43+ this->tileDataNum = tileDataNum;
44+ if (coreNum < tailBlockNum)
45+ {
46+ this->coreDataNum = bigCoreDataNum;
47+ this->tileNum = finalBigTileNum;
48+ this->tailDataNum = bigTailDataNum;
49+ }
50+ else
51+ {
52+ this->coreDataNum = smallCoreDataNum;
53+ this->tileNum = finalSmallTileNum;
54+ this->tailDataNum = smallTailDataNum;
55+ globalBufferIndex -= (bigCoreDataNum - smallCoreDataNum) * (GetBlockIdx() - tailBlockNum);
56+ }
57+ xGm.SetGlobalBuffer((__gm__ TYPE_X *)x + globalBufferIndex, this->coreDataNum);
58+ yGm.SetGlobalBuffer((__gm__ TYPE_Y *)y + globalBufferIndex, this->coreDataNum);
59+ 
60+ this->scale = -1.0f * scale;
61+ pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileDataNum * sizeof(TYPE_X));
62+ pipe.InitBuffer(outQueueY, BUFFER_NUM, this->tileDataNum * sizeof(TYPE_Y));
63+ if constexpr (!std::is_same_v<TYPE_X, float>)
64+ {
65+ pipe.InitBuffer(calcBuf1, this->tileDataNum * sizeof(float));
66+ pipe.InitBuffer(calcBuf2, this->tileDataNum * sizeof(float));
67+ }
68+ }
69+ __aicore__ inline void Process()
70+ {
71+ int32_t loopCount = this->tileNum;
72+ this->processDataNum = this->tileDataNum;
73+ for (int32_t i = 0; i < loopCount-1; i++)
74+ {
75+ CopyIn(i);
76+ Compute(i);
77+ CopyOut(i);
78+ }
79+ this->processDataNum = this->tailDataNum;
80+ CopyIn(loopCount-1);
81+ Compute(loopCount-1);
82+ CopyOut(loopCount-1);
83+ }
84+ 
85+private:
86+ __aicore__ inline void CopyIn(int32_t progress)
87+ {
88+ LocalTensor<TYPE_X> xLocal = inQueueX.AllocTensor<TYPE_X>();
89+ DataCopy(xLocal, xGm[progress * this->tileDataNum], this->processDataNum);
90+ inQueueX.EnQue(xLocal);
91+ }
92+ __aicore__ inline void Compute(int32_t progress)
93+ {
94+ LocalTensor<TYPE_X> xLocal = inQueueX.DeQue<TYPE_X>();
95+ LocalTensor<TYPE_Y> yLocal = outQueueY.AllocTensor<TYPE_Y>();
96+ if constexpr (std::is_same_v<TYPE_X, float>){
97+ Muls(yLocal, xLocal, this->scale, this->processDataNum);
98+ Exp(yLocal, yLocal, this->processDataNum);
99+ Adds(yLocal, yLocal, (TYPE_X)1.0f, this->processDataNum);
100+ Div(yLocal, xLocal, yLocal, this->processDataNum);
101+ }else if constexpr (!std::is_same_v<TYPE_X, float>){
102+ LocalTensor<float> xLocalFp32 = calcBuf1.Get<float>();
103+ LocalTensor<float> yLocalFp32 = calcBuf2.Get<float>();
104+
105+ Cast(xLocalFp32, xLocal, RoundMode::CAST_NONE, this->processDataNum);
106+ Muls(yLocalFp32, xLocalFp32, this->scale, this->processDataNum);
107+ Exp(yLocalFp32, yLocalFp32, this->processDataNum);
108+ Adds(yLocalFp32, yLocalFp32, 1.0f, this->processDataNum);
109+ Div(yLocalFp32, xLocalFp32, yLocalFp32, this->processDataNum);
110+ Cast(yLocal, yLocalFp32, RoundMode::CAST_ROUND, this->processDataNum);
111+ }
112+ outQueueY.EnQue<TYPE_Y>(yLocal);
113+ inQueueX.FreeTensor(xLocal);
114+ }
115+ __aicore__ inline void CopyOut(int32_t progress)
116+ {
117+ LocalTensor<TYPE_Y> yLocal = outQueueY.DeQue<TYPE_Y>();
118+ DataCopy(yGm[progress * this->tileDataNum], yLocal, this->processDataNum);
119+ outQueueY.FreeTensor(yLocal);
120+ }
121+ 
122+private:
123+ TPipe pipe;
124+ TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX;
125+ TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueY;
126+ TBuf<QuePosition::VECCALC> calcBuf1;
127+ TBuf<QuePosition::VECCALC> calcBuf2;
128+ GlobalTensor<TYPE_X> xGm;
129+ GlobalTensor<TYPE_Y> yGm;
130+ uint32_t coreDataNum;
131+ uint32_t tileNum;
132+ uint32_t tileDataNum;
133+ uint32_t tailDataNum;
134+ uint32_t processDataNum;
135+ float scale;
136+};
137+ 
138+} // namespace NsSwish
139+#endif // SWISH_H
Aexperimental/activation/swish/op_kernel/swish_tiling_data.h+30-0
@@ -0,0 +1,30 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+*/
10+ 
11+/*!
12+ * \file swish_tiling_data.h
13+ * \brief tiling data struct
14+ */
15+ 
16+#ifndef __SWISH_TILLING_DATA_H__
17+#define __SWISH_TILLING_DATA_H__
18+ 
19+struct SwishTilingData {
20+ uint32_t smallCoreDataNum;
21+ uint32_t bigCoreDataNum;
22+ uint32_t finalBigTileNum;
23+ uint32_t finalSmallTileNum;
24+ uint32_t tileDataNum;
25+ uint32_t smallTailDataNum;
26+ uint32_t bigTailDataNum;
27+ uint32_t tailBlockNum;
28+ float scale;
29+};
30+#endif
Aexperimental/activation/swish/op_kernel/swish_tiling_key.h+44-0
@@ -0,0 +1,44 @@
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 swish_tiling_key.h
13+ * \brief swish tiling key declare
14+ */
15+ 
16+#ifndef __SWISH_TILING_KEY_H__
17+#define __SWISH_TILING_KEY_H__
18+ 
19+#include "ascendc/host_api/tiling/template_argument.h"
20+ 
21+namespace SwishOp {
22+#define TPL_SCALE_NEG_ONE 1
23+#define TPL_SCALE_ZERO 2
24+#define TPL_SCALE_OTHER 3
25+ 
26+#define TPL_SCH_MODE_0 0
27+#define TPL_SCH_MODE_1 1
28+ 
29+ASCENDC_TPL_ARGS_DECL(Swish,
30+ ASCENDC_TPL_UINT_DECL(schMode, 1, ASCENDC_TPL_UI_LIST, TPL_SCH_MODE_0, TPL_SCH_MODE_1),
31+ ASCENDC_TPL_DTYPE_DECL(dType, TPL_SCALE_NEG_ONE, TPL_SCALE_ZERO, TPL_SCALE_OTHER)
32+);
33+ 
34+ASCENDC_TPL_SEL(
35+ ASCENDC_TPL_ARGS_SEL(
36+ ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST, TPL_SCH_MODE_0, TPL_SCH_MODE_1),
37+ ASCENDC_TPL_DTYPE_SEL(dType, TPL_SCALE_NEG_ONE, TPL_SCALE_ZERO, TPL_SCALE_OTHER)
38+ )
39+);
40+} // namespace SwishOp
41+ 
42+using namespace SwishOp;
43+ 
44+#endif
Aexperimental/activation/swish/tests/CMakeLists.txt+17-0
@@ -0,0 +1,17 @@
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+message(STATUS "=== Debug: CURRENT_SOURCE_DIRS =${CURRENT_SOURCE_DIRS} ")
13+foreach(SUB_DIR ${CURRENT_SOURCE_DIRS})
14+ if(EXISTS "${SUB_DIR}/CMakeLists.txt")
15+ add_subdirectory(${SUB_DIR})
16+ endif()
17+endforeach()
Aexperimental/activation/swish/tests/ut/CMakeLists.txt+17-0
@@ -0,0 +1,17 @@
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+message(STATUS "=== Debug: CURRENT_SOURCE_DIRS =${CURRENT_SOURCE_DIRS} ")
13+foreach(SUB_DIR ${CURRENT_SOURCE_DIRS})
14+ if(EXISTS "${SUB_DIR}/CMakeLists.txt")
15+ add_subdirectory(${SUB_DIR})
16+ endif()
17+endforeach()
Aexperimental/activation/swish/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/swish/tests/ut/op_api/test_aclnn_silu.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+#include <float.h>
11+ 
12+#include <array>
13+#include <iostream>
14+#include <vector>
15+ 
16+#include "gtest/gtest.h"
17+#include "../../../op_host/op_api/aclnn_silu.h"
18+#include "op_api_ut_common/op_api_ut.h"
19+#include "op_api_ut_common/scalar_desc.h"
20+#include "op_api_ut_common/tensor_desc.h"
21+ 
22+using namespace std;
23+ 
24+class silu_test : public testing::Test {
25+ protected:
26+ static void SetUpTestCase() { cout << "sign_test SetUp" << endl; }
27+ 
28+ static void TearDownTestCase() { cout << "sign_test TeastDown" << endl; }
29+};
30+ 
31+TEST_F(silu_test, test_silu_dataType_error) {
32+ vector<aclDataType> ValidList = {ACL_DT_UNDEFINED, ACL_INT8, ACL_INT16, ACL_INT32, ACL_INT64,
33+ ACL_UINT8, ACL_UINT16, ACL_UINT32, ACL_UINT64, ACL_DOUBLE,
34+ ACL_BOOL, ACL_STRING, ACL_COMPLEX64, ACL_COMPLEX128, ACL_BF16};
35+ 
36+ int length = ValidList.size();
37+ vector<int64_t> input_dim = {2, 16, 32, 16};
38+ vector<int64_t> result_dim = {2, 16, 32, 16};
39+ 
40+ for (int i = 0; i < length; i++) {
41+ auto inputDesc = TensorDesc(input_dim, ValidList[i], ACL_FORMAT_ND).ValueRange(-2, 2);
42+ auto outDesc = TensorDesc(result_dim, ValidList[i], ACL_FORMAT_ND).Precision(0.0001, 0.0001);
43+ 
44+ auto ut = OP_API_UT(aclnnSilu, INPUT(inputDesc), OUTPUT(outDesc));
45+ uint64_t workspaceSize = 0;
46+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
47+ }
48+}
49+ 
50+TEST_F(silu_test, test_silu_format) {
51+ vector<aclFormat> ValidList = {
52+ ACL_FORMAT_UNDEFINED, ACL_FORMAT_NCHW, ACL_FORMAT_NHWC, ACL_FORMAT_ND, ACL_FORMAT_NC1HWC0,
53+ ACL_FORMAT_FRACTAL_Z, ACL_FORMAT_NC1HWC0_C04, ACL_FORMAT_HWCN, ACL_FORMAT_NDHWC, ACL_FORMAT_FRACTAL_NZ,
54+ ACL_FORMAT_NCDHW, ACL_FORMAT_NDC1HWC0, ACL_FRACTAL_Z_3D, ACL_FORMAT_NC, ACL_FORMAT_NCL};
55+ 
56+ int length = ValidList.size();
57+ vector<int64_t> input_dim = {2, 16, 32, 16};
58+ vector<int64_t> result_dim = {2, 16, 32, 16};
59+ 
60+ for (int i = 0; i < length; i++) {
61+ auto inputDesc = TensorDesc(input_dim, ACL_FLOAT, ValidList[i]).ValueRange(-1, 1);
62+ auto outDesc = TensorDesc(result_dim, ACL_FLOAT, ValidList[i]).Precision(0.0001, 0.0001);
63+ 
64+ auto ut = OP_API_UT(aclnnSilu, INPUT(inputDesc), OUTPUT(outDesc));
65+ uint64_t workspaceSize = 0;
66+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
67+ EXPECT_EQ(aclRet, ACLNN_SUCCESS);
68+ ut.TestPrecision();
69+ }
70+}
71+ 
72+TEST_F(silu_test, test_silu_inconsistent_shape) {
73+ auto inputDesc = TensorDesc({2, 16, 32, 16}, ACL_FLOAT, ACL_FORMAT_ND);
74+ auto outDesc = TensorDesc({2, 16, 32, 18}, ACL_FLOAT, ACL_FORMAT_ND);
75+ 
76+ auto ut = OP_API_UT(aclnnSilu, INPUT(inputDesc), OUTPUT(outDesc));
77+ 
78+ uint64_t workspaceSize = 0;
79+ aclnnStatus getWorkspaceout = ut.TestGetWorkspaceSize(&workspaceSize);
80+ EXPECT_EQ(getWorkspaceout, ACLNN_ERR_PARAM_INVALID);
81+}
82+ 
83+TEST_F(silu_test, test_silu_inconsistent_dtype) {
84+ auto inputDesc = TensorDesc({2, 16, 32, 16}, ACL_FLOAT, ACL_FORMAT_ND);
85+ auto outDesc = TensorDesc({2, 16, 32, 16}, ACL_FLOAT16, ACL_FORMAT_ND);
86+ 
87+ auto ut = OP_API_UT(aclnnSilu, INPUT(inputDesc), OUTPUT(outDesc));
88+ 
89+ uint64_t workspaceSize = 0;
90+ aclnnStatus getWorkspaceout = ut.TestGetWorkspaceSize(&workspaceSize);
91+ EXPECT_EQ(getWorkspaceout, ACLNN_ERR_PARAM_INVALID);
92+}
93+ 
94+TEST_F(silu_test, test_silu_empty_input) {
95+ auto inputDesc = TensorDesc({2, 16, 32, 16}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
96+ auto outDesc = TensorDesc({2, 16, 32, 16}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.0001, 0.0001);
97+ 
98+ auto ut = OP_API_UT(aclnnSilu, INPUT(inputDesc), OUTPUT(outDesc));
99+ 
100+ uint64_t workspaceSize = 0;
101+ aclnnStatus getWorkspaceout = ut.TestGetWorkspaceSize(&workspaceSize);
102+ EXPECT_EQ(getWorkspaceout, ACLNN_SUCCESS);
103+ ut.TestPrecision();
104+}
105+ 
106+TEST_F(silu_test, test_silu_nullptr_input) {
107+ auto outDesc = TensorDesc({2, 16, 32, 16}, ACL_FLOAT, ACL_FORMAT_ND);
108+ 
109+ auto ut = OP_API_UT(aclnnSilu, INPUT((aclTensor *)nullptr), OUTPUT(outDesc));
110+ 
111+ uint64_t workspaceSize = 0;
112+ aclnnStatus getWorkspaceout = ut.TestGetWorkspaceSize(&workspaceSize);
113+ EXPECT_EQ(getWorkspaceout, ACLNN_ERR_PARAM_NULLPTR);
114+}
115+ 
116+TEST_F(silu_test, test_silu_nullptr_out) {
117+ auto inputDesc = TensorDesc({2, 16, 32, 16}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
118+ 
119+ auto ut = OP_API_UT(aclnnSilu, INPUT(inputDesc), OUTPUT((aclTensor *)nullptr));
120+ 
121+ uint64_t workspaceSize = 0;
122+ aclnnStatus getWorkspaceout = ut.TestGetWorkspaceSize(&workspaceSize);
123+ EXPECT_EQ(getWorkspaceout, ACLNN_ERR_PARAM_NULLPTR);
124+}
125+ 
126+TEST_F(silu_test, test_silu_FP32) {
127+ auto inputDesc = TensorDesc({2, 16, 32, 16}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
128+ auto outDesc = TensorDesc({2, 16, 32, 16}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.0001, 0.0001);
129+ 
130+ auto ut = OP_API_UT(aclnnSilu, INPUT(inputDesc), OUTPUT(outDesc));
131+ 
132+ uint64_t workspaceSize = 0;
133+ aclnnStatus getWorkspaceout = ut.TestGetWorkspaceSize(&workspaceSize);
134+ EXPECT_EQ(getWorkspaceout, ACLNN_SUCCESS);
135+ ut.TestPrecision();
136+}
137+ 
138+TEST_F(silu_test, test_silu_FP16) {
139+ auto inputDesc = TensorDesc({2, 16, 32, 16}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-1, 1);
140+ auto outDesc = TensorDesc({2, 16, 32, 16}, ACL_FLOAT16, ACL_FORMAT_ND).Precision(0.0001, 0.0001);
141+ 
142+ auto ut = OP_API_UT(aclnnSilu, INPUT(inputDesc), OUTPUT(outDesc));
143+ 
144+ uint64_t workspaceSize = 0;
145+ aclnnStatus getWorkspaceout = ut.TestGetWorkspaceSize(&workspaceSize);
146+ EXPECT_EQ(getWorkspaceout, ACLNN_SUCCESS);
147+ ut.TestPrecision();
148+}
149+ 
150+TEST_F(silu_test, test_silu_uncontiguous) {
151+ auto inputDesc = TensorDesc({2, 16}, ACL_FLOAT, ACL_FORMAT_ND, {1, 2}, 0, {16, 2}).ValueRange(-1, 1);
152+ auto outDesc = TensorDesc({2, 16}, ACL_FLOAT, ACL_FORMAT_ND, {1, 2}, 0, {16, 2}).Precision(0.0001, 0.0001);
153+ 
154+ auto ut = OP_API_UT(aclnnSilu, INPUT(inputDesc), OUTPUT(outDesc));
155+ 
156+ uint64_t workspaceSize = 0;
157+ aclnnStatus getWorkspaceout = ut.TestGetWorkspaceSize(&workspaceSize);
158+ EXPECT_EQ(getWorkspaceout, ACLNN_SUCCESS);
159+ ut.TestPrecision();
160+}
161+ 
162+TEST_F(silu_test, test_silu_shape_larger_8) {
163+ vector<aclFormat> ValidList = {
164+ ACL_FORMAT_UNDEFINED, ACL_FORMAT_NCHW, ACL_FORMAT_NHWC, ACL_FORMAT_ND, ACL_FORMAT_NC1HWC0,
165+ ACL_FORMAT_FRACTAL_Z, ACL_FORMAT_NC1HWC0_C04, ACL_FORMAT_HWCN, ACL_FORMAT_NDHWC, ACL_FORMAT_FRACTAL_NZ,
166+ ACL_FORMAT_NCDHW, ACL_FORMAT_NDC1HWC0, ACL_FRACTAL_Z_3D, ACL_FORMAT_NC, ACL_FORMAT_NCL};
167+ 
168+ int length = ValidList.size();
169+ 
170+ for (int i = 0; i < length; i++) {
171+ auto inputDesc = TensorDesc({1, 1, 1, 1, 2, 1, 1, 1, 2}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
172+ auto resultDesc = TensorDesc({1, 1, 1, 1, 2, 1, 1, 1, 2}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.0001, 0.0001);
173+ auto ut = OP_API_UT(aclnnSilu, INPUT(inputDesc), OUTPUT(resultDesc));
174+ 
175+ uint64_t workspaceSize = 0;
176+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
177+ EXPECT_EQ(aclRet, ACLNN_SUCCESS);
178+ // ut.TestPrecision(); // comment bcz of timeout in model tests (986327 ms)
179+ }
180+}
Aexperimental/activation/swish/tests/ut/op_api/test_aclnn_swish.cpp+216-0
@@ -0,0 +1,216 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+#include <float.h>
11+ 
12+#include <array>
13+#include <iostream>
14+#include <vector>
15+ 
16+#include "gtest/gtest.h"
17+#include "../../../op_host/op_api/aclnn_swish.h"
18+#include "op_api_ut_common/op_api_ut.h"
19+#include "op_api_ut_common/scalar_desc.h"
20+#include "op_api_ut_common/tensor_desc.h"
21+ 
22+using namespace std;
23+ 
24+class swish_test : public testing::Test {
25+ protected:
26+ static void SetUpTestCase() { cout << "sign_test SetUp" << endl; }
27+ 
28+ static void TearDownTestCase() { cout << "sign_test TeastDown" << endl; }
29+};
30+ 
31+TEST_F(swish_test, test_swish_dataType_error) {
32+ vector<aclDataType> ValidList = {ACL_DT_UNDEFINED, ACL_INT8, ACL_INT16, ACL_INT32, ACL_INT64,
33+ ACL_UINT8, ACL_UINT16, ACL_UINT32, ACL_UINT64, ACL_DOUBLE,
34+ ACL_BOOL, ACL_STRING, ACL_COMPLEX64, ACL_COMPLEX128, ACL_BF16};
35+ 
36+ int length = ValidList.size();
37+ vector<int64_t> input_dim = {2, 16, 32, 16};
38+ vector<int64_t> result_dim = {2, 16, 32, 16};
39+ 
40+ for (int i = 0; i < length; i++) {
41+ auto inputDesc = TensorDesc(input_dim, ValidList[i], ACL_FORMAT_ND).ValueRange(-2, 2);
42+ auto betaDesc = ScalarDesc(static_cast<bool>(false));
43+ auto outDesc = TensorDesc(result_dim, ValidList[i], ACL_FORMAT_ND).Precision(0.0001, 0.0001);
44+ auto ut = OP_API_UT(aclnnSwish, INPUT(inputDesc, betaDesc), OUTPUT(outDesc));
45+ 
46+ uint64_t workspaceSize = 0;
47+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
48+ }
49+}
50+ 
51+TEST_F(swish_test, test_swish_format) {
52+ vector<aclFormat> ValidList = {
53+ ACL_FORMAT_UNDEFINED, ACL_FORMAT_NCHW, ACL_FORMAT_NHWC, ACL_FORMAT_ND, ACL_FORMAT_NC1HWC0,
54+ ACL_FORMAT_FRACTAL_Z, ACL_FORMAT_NC1HWC0_C04, ACL_FORMAT_HWCN, ACL_FORMAT_NDHWC, ACL_FORMAT_FRACTAL_NZ,
55+ ACL_FORMAT_NCDHW, ACL_FORMAT_NDC1HWC0, ACL_FRACTAL_Z_3D, ACL_FORMAT_NC, ACL_FORMAT_NCL};
56+ 
57+ int length = ValidList.size();
58+ vector<int64_t> input_dim = {2, 16, 32, 16};
59+ vector<int64_t> result_dim = {2, 16, 32, 16};
60+ 
61+ for (int i = 0; i < length; i++) {
62+ auto inputDesc = TensorDesc(input_dim, ACL_FLOAT, ValidList[i]).ValueRange(-1, 1);
63+ auto betaDesc = ScalarDesc(static_cast<double>(0.0));
64+ auto outDesc = TensorDesc(result_dim, ACL_FLOAT, ValidList[i]).Precision(0.0001, 0.0001);
65+ auto ut = OP_API_UT(aclnnSwish, INPUT(inputDesc, betaDesc), OUTPUT(outDesc));
66+ 
67+ uint64_t workspaceSize = 0;
68+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
69+ EXPECT_EQ(aclRet, ACLNN_SUCCESS);
70+ ut.TestPrecision();
71+ }
72+}
73+ 
74+TEST_F(swish_test, test_swish_inconsistent_shape) {
75+ auto inputDesc = TensorDesc({2, 16, 32, 16}, ACL_FLOAT, ACL_FORMAT_ND);
76+ auto betaDesc = ScalarDesc(1.1f, ACL_FLOAT);
77+ auto outDesc = TensorDesc({2, 16, 32, 18}, ACL_FLOAT, ACL_FORMAT_ND);
78+ auto ut = OP_API_UT(aclnnSwish, INPUT(inputDesc, betaDesc), OUTPUT(outDesc));
79+ 
80+ uint64_t workspaceSize = 0;
81+ aclnnStatus getWorkspaceout = ut.TestGetWorkspaceSize(&workspaceSize);
82+ EXPECT_EQ(getWorkspaceout, ACLNN_ERR_PARAM_INVALID);
83+}
84+ 
85+TEST_F(swish_test, test_swish_inconsistent_dtype) {
86+ auto inputDesc = TensorDesc({2, 16, 32, 16}, ACL_FLOAT, ACL_FORMAT_ND);
87+ auto betaDesc = ScalarDesc(1.1f, ACL_FLOAT);
88+ auto outDesc = TensorDesc({2, 16, 32, 16}, ACL_FLOAT16, ACL_FORMAT_ND);
89+ auto ut = OP_API_UT(aclnnSwish, INPUT(inputDesc, betaDesc), OUTPUT(outDesc));
90+ 
91+ uint64_t workspaceSize = 0;
92+ aclnnStatus getWorkspaceout = ut.TestGetWorkspaceSize(&workspaceSize);
93+ EXPECT_EQ(getWorkspaceout, ACLNN_ERR_PARAM_INVALID);
94+}
95+ 
96+TEST_F(swish_test, test_swish_empty_input) {
97+ auto inputDesc = TensorDesc({2, 16, 32, 16}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
98+ auto betaDesc = ScalarDesc(static_cast<uint8_t>(0));
99+ auto outDesc = TensorDesc({2, 16, 32, 16}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.0001, 0.0001);
100+ auto ut = OP_API_UT(aclnnSwish, INPUT(inputDesc, betaDesc), OUTPUT(outDesc));
101+ 
102+ uint64_t workspaceSize = 0;
103+ aclnnStatus getWorkspaceout = ut.TestGetWorkspaceSize(&workspaceSize);
104+ EXPECT_EQ(getWorkspaceout, ACLNN_SUCCESS);
105+ ut.TestPrecision();
106+}
107+ 
108+TEST_F(swish_test, test_swish_nullptr_input) {
109+ auto betaDesc = ScalarDesc(1.1f, ACL_FLOAT);
110+ auto outDesc = TensorDesc({2, 16, 32, 16}, ACL_FLOAT, ACL_FORMAT_ND);
111+ auto ut = OP_API_UT(aclnnSwish, INPUT((aclTensor *)nullptr, betaDesc), OUTPUT(outDesc));
112+ 
113+ uint64_t workspaceSize = 0;
114+ aclnnStatus getWorkspaceout = ut.TestGetWorkspaceSize(&workspaceSize);
115+ EXPECT_EQ(getWorkspaceout, ACLNN_ERR_PARAM_NULLPTR);
116+}
117+ 
118+TEST_F(swish_test, test_swish_nullptr_out) {
119+ auto betaDesc = ScalarDesc(1.1f, ACL_FLOAT);
120+ auto inputDesc = TensorDesc({2, 16, 32, 16}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
121+ auto ut = OP_API_UT(aclnnSwish, INPUT(inputDesc, betaDesc), OUTPUT((aclTensor *)nullptr));
122+ 
123+ uint64_t workspaceSize = 0;
124+ aclnnStatus getWorkspaceout = ut.TestGetWorkspaceSize(&workspaceSize);
125+ EXPECT_EQ(getWorkspaceout, ACLNN_ERR_PARAM_NULLPTR);
126+}
127+ 
128+TEST_F(swish_test, test_swish_nullptr_beta) {
129+ auto inputDesc = TensorDesc({2, 16, 32, 16}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
130+ auto betaDesc = (aclScalar*)nullptr;
131+ auto outDesc = TensorDesc({2, 16, 32, 16}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.0001, 0.0001);
132+ auto ut = OP_API_UT(aclnnSwish, INPUT(inputDesc, betaDesc), OUTPUT(outDesc));
133+ 
134+ uint64_t workspaceSize = 0;
135+ aclnnStatus getWorkspaceout = ut.TestGetWorkspaceSize(&workspaceSize);
136+ EXPECT_EQ(getWorkspaceout, ACLNN_SUCCESS);
137+}
138+ 
139+TEST_F(swish_test, test_swish_FP32) {
140+ auto inputDesc = TensorDesc({2, 16, 32, 16}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
141+ auto betaDesc = ScalarDesc(1.1f, ACL_FLOAT);
142+ auto outDesc = TensorDesc({2, 16, 32, 16}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.0001, 0.0001);
143+ auto ut = OP_API_UT(aclnnSwish, INPUT(inputDesc, betaDesc), OUTPUT(outDesc));
144+ 
145+ uint64_t workspaceSize = 0;
146+ aclnnStatus getWorkspaceout = ut.TestGetWorkspaceSize(&workspaceSize);
147+ EXPECT_EQ(getWorkspaceout, ACLNN_SUCCESS);
148+ ut.TestPrecision();
149+}
150+ 
151+TEST_F(swish_test, test_swish_FP16) {
152+ auto inputDesc = TensorDesc({2, 16, 32, 16}, ACL_FLOAT16, ACL_FORMAT_ND).ValueRange(-1, 1);
153+ auto betaDesc = ScalarDesc(0.01f, ACL_FLOAT16);
154+ auto outDesc = TensorDesc({2, 16, 32, 16}, ACL_FLOAT16, ACL_FORMAT_ND).Precision(0.0001, 0.0001);
155+ auto ut = OP_API_UT(aclnnSwish, INPUT(inputDesc, betaDesc), OUTPUT(outDesc));
156+ 
157+ uint64_t workspaceSize = 0;
158+ aclnnStatus getWorkspaceout = ut.TestGetWorkspaceSize(&workspaceSize);
159+ EXPECT_EQ(getWorkspaceout, ACLNN_SUCCESS);
160+ ut.TestPrecision();
161+}
162+ 
163+TEST_F(swish_test, test_swish_uncontiguous) {
164+ auto inputDesc = TensorDesc({2, 16}, ACL_FLOAT, ACL_FORMAT_ND, {1, 2}, 0, {16, 2}).ValueRange(-1, 1);
165+ auto betaDesc = ScalarDesc(static_cast<int8_t>(-1));
166+ auto outDesc = TensorDesc({2, 16}, ACL_FLOAT, ACL_FORMAT_ND, {1, 2}, 0, {16, 2}).Precision(0.0001, 0.0001);
167+ auto ut = OP_API_UT(aclnnSwish, INPUT(inputDesc, betaDesc), OUTPUT(outDesc));
168+ 
169+ uint64_t workspaceSize = 0;
170+ aclnnStatus getWorkspaceout = ut.TestGetWorkspaceSize(&workspaceSize);
171+ EXPECT_EQ(getWorkspaceout, ACLNN_SUCCESS);
172+ ut.TestPrecision();
173+}
174+ 
175+TEST_F(swish_test, test_swish_shape_larger_8) {
176+ vector<aclFormat> ValidList = {
177+ ACL_FORMAT_UNDEFINED, ACL_FORMAT_NCHW, ACL_FORMAT_NHWC, ACL_FORMAT_ND, ACL_FORMAT_NC1HWC0,
178+ ACL_FORMAT_FRACTAL_Z, ACL_FORMAT_NC1HWC0_C04, ACL_FORMAT_HWCN, ACL_FORMAT_NDHWC, ACL_FORMAT_FRACTAL_NZ,
179+ ACL_FORMAT_NCDHW, ACL_FORMAT_NDC1HWC0, ACL_FRACTAL_Z_3D, ACL_FORMAT_NC, ACL_FORMAT_NCL};
180+ 
181+ int length = ValidList.size();
182+ 
183+ for (int i = 0; i < length; i++) {
184+ auto inputDesc = TensorDesc({1, 1, 1, 1, 2, 1, 1, 1, 2}, ACL_FLOAT, ACL_FORMAT_ND).ValueRange(-1, 1);
185+ auto betaDesc = ScalarDesc(1.1f, ACL_FLOAT);
186+ auto resultDesc = TensorDesc({1, 1, 1, 1, 2, 1, 1, 1, 2}, ACL_FLOAT, ACL_FORMAT_ND).Precision(0.0001, 0.0001);
187+ auto ut = OP_API_UT(aclnnSwish, INPUT(inputDesc, betaDesc), OUTPUT(resultDesc));
188+ 
189+ uint64_t workspaceSize = 0;
190+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspaceSize);
191+ EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
192+ // ut.TestPrecision(); // comment bcz of timeout in model tests (986327 ms)
193+ }
194+}
195+ 
196+TEST_F(swish_test, test_swish_input_invalid_dim) {
197+ auto inputDesc = TensorDesc({1, 2, 3, 4, 5, 6, 7, 8, 9}, ACL_FLOAT, ACL_FORMAT_ND);
198+ auto betaDesc = ScalarDesc(1.1f, ACL_FLOAT);
199+ auto outDesc = TensorDesc({1, 2, 3, 4}, ACL_FLOAT, ACL_FORMAT_ND);
200+ auto ut = OP_API_UT(aclnnSwish, INPUT(inputDesc, betaDesc), OUTPUT(outDesc));
201+ 
202+ uint64_t workspaceSize = 0;
203+ aclnnStatus getWorkspaceout = ut.TestGetWorkspaceSize(&workspaceSize);
204+ EXPECT_EQ(getWorkspaceout, ACLNN_ERR_PARAM_INVALID);
205+}
206+ 
207+TEST_F(swish_test, test_swish_output_invalid_dim) {
208+ auto inputDesc = TensorDesc({1, 2, 3, 4}, ACL_FLOAT, ACL_FORMAT_ND);
209+ auto betaDesc = ScalarDesc(1.1f, ACL_FLOAT);
210+ auto outDesc = TensorDesc({1, 2, 3, 4, 5, 6, 7, 8, 9}, ACL_FLOAT, ACL_FORMAT_ND);
211+ auto ut = OP_API_UT(aclnnSwish, INPUT(inputDesc, betaDesc), OUTPUT(outDesc));
212+ 
213+ uint64_t workspaceSize = 0;
214+ aclnnStatus getWorkspaceout = ut.TestGetWorkspaceSize(&workspaceSize);
215+ EXPECT_EQ(getWorkspaceout, ACLNN_ERR_PARAM_INVALID);
216+}
Aexperimental/activation/swish/tests/ut/op_host/CMakeLists.txt+17-0
@@ -0,0 +1,17 @@
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+ 
13+if(UT_TEST_ALL OR OP_HOST_UT)
14+ add_modules_ut_sources(HOSTNAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
15+ add_modules_ut_sources(HOSTNAME ${OP_INFERSHAPE_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
16+endif()
17+ 
Aexperimental/activation/swish/tests/ut/op_host/test_swish_infershape.cpp+62-0
@@ -0,0 +1,62 @@
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_swish_infershape.cpp
13+* \brief
14+*/
15+#include <gtest/gtest.h>
16+#include <iostream>
17+#include "log/log.h"
18+#include "ut_op_common.h"
19+#include "infershape_test_util.h"
20+#include "platform/platform_info.h"
21+ 
22+#include "../../../op_graph/swish_proto.h"
23+ 
24+class Swish : public testing::Test {
25+protected:
26+ static void SetUpTestCase()
27+ {
28+ std::cout << "Swish Proto Test SetUp" << std::endl;
29+ }
30+ 
31+ static void TearDownTestCase() {
32+ std::cout << "Swish Proto Test TearDown" << std::endl;
33+ }
34+};
35+ 
36+TEST_F(Swish, swish_infershape_test0)
37+{
38+ fe::PlatformInfo platformInfo;
39+ fe::OptionalInfo optiCompilationInfo;
40+ platformInfo.soc_info.ai_core_cnt = 64;
41+ platformInfo.str_info.short_soc_version = "Ascend910_95";
42+ optiCompilationInfo.soc_version = "Ascend910_95";
43+ fe::PlatformInfoManager::Instance().platform_info_map_["Ascend910_95"] = platformInfo;
44+ fe::PlatformInfoManager::Instance().SetOptionalCompilationInfo(optiCompilationInfo);
45+ 
46+ auto inferShapeFunc = gert::OpImplRegistry::GetInstance().GetOpImpl("Swish")->infer_shape;
47+ gert::Shape xShape = {3, 4, 5};
48+ gert::Shape output_shape = {};
49+ 
50+ auto holder = gert::InferShapeContextFaker()
51+ .NodeIoNum(1, 1)
52+ .IrInstanceNum({1})
53+ .InputShapes({&xShape})
54+ .OutputShapes({&output_shape})
55+ .NodeAttrs(
56+ {{"scale", Ops::NN::AnyValue::CreateFrom<float>(1.0f)}})
57+ .NodeInputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
58+ .NodeOutputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
59+ .Build();
60+ 
61+ ASSERT_EQ(inferShapeFunc(holder.GetContext<gert::InferShapeContext>()), ge::GRAPH_SUCCESS);
62+}
Aexperimental/activation/swish/tests/ut/op_host/test_swish_tiling.cpp+157-0
@@ -0,0 +1,157 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include <iostream>
12+#include <vector>
13+#include <gtest/gtest.h>
14+#include "log/log.h"
15+#include "kernel_run_context_facker.h"
16+#include "exe_graph/runtime/storage_format.h"
17+#include "exe_graph/runtime/storage_shape.h"
18+#include "test_cube_util.h"
19+#include "register/op_impl_registry.h"
20+#include "ut_op_util.h"
21+#include "ut_op_common.h"
22+#include "platform/platform_infos_def.h"
23+ 
24+using namespace ut_util;
25+using namespace std;
26+using namespace ge;
27+ 
28+class SwishTiling : public testing::Test {
29+protected:
30+ static void SetUpTestCase()
31+ {
32+ std::cout << "SwishTiling SetUp" << std::endl;
33+ }
34+ 
35+ static void TearDownTestCase()
36+ {
37+ std::cout << "SwishTiling TearDown" << std::endl;
38+ }
39+};
40+ 
41+TEST_F(SwishTiling, swish_float32_success) {
42+ gert::StorageShape x_shape = {{1, 2, 8, 16}, {1, 2, 8, 16}};
43+ gert::StorageShape y_shape = {{1, 2, 8, 16}, {1, 2, 8, 16}};
44+ string compile_info_string = R"({
45+ "hardware_info": {"BT_SIZE": 0, "load3d_constraints": "1",
46+ "Intrinsic_fix_pipe_l0c2out": false, "Intrinsic_data_move_l12ub": true, "Intrinsic_data_move_l0c2ub": true, "Intrinsic_data_move_out2l1_nd2nz": false,
47+ "UB_SIZE": 196608, "L2_SIZE": 33554432, "L1_SIZE": 524288,
48+ "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072,
49+ "CORE_NUM": 48}
50+ })";
51+ map<string, string> soc_infos;
52+ map<string, string> aicore_spec;
53+ map<string, string> intrinsics;
54+ GetPlatFormInfos(compile_info_string.c_str(), soc_infos, aicore_spec, intrinsics);
55+ 
56+ fe::PlatFormInfos platform_info;
57+ platform_info.Init();
58+ 
59+ struct SwishCompileInfo {};
60+ SwishCompileInfo compile_info;
61+ 
62+ std::string op_type("Swish");
63+ auto tiling_func = gert::OpImplRegistry::GetInstance().GetOpImpl(op_type.c_str())->tiling;
64+ 
65+ auto kernel_holder =
66+ gert::KernelRunContextFaker()
67+ .KernelIONum(3, 1)
68+ .Inputs({const_cast<char*>(compile_info_string.c_str()), reinterpret_cast<void*>(&platform_info)})
69+ .Outputs({&compile_info})
70+ .Build();
71+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", soc_infos);
72+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicore_spec);
73+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
74+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap",
75+ intrinsics);
76+ 
77+ auto param = gert::TilingData::CreateCap(4096);
78+ auto workspace_size_holer = gert::ContinuousVector::Create<size_t>(4096);
79+ auto ws_size = reinterpret_cast<gert::ContinuousVector*>(workspace_size_holer.get());
80+ ASSERT_NE(param, nullptr);
81+ auto holder = gert::TilingContextFaker()
82+ .SetOpType("Swish")
83+ .NodeIoNum(1, 1)
84+ .IrInstanceNum({1})
85+ .InputShapes({&x_shape})
86+ .OutputShapes({&y_shape})
87+ .CompileInfo(&compile_info)
88+ .PlatformInfo(reinterpret_cast<char*>(&platform_info))
89+ .NodeAttrs({{"scale", Ops::NN::AnyValue::CreateFrom<float>(1.0f)}})
90+ .NodeInputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
91+ .NodeOutputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
92+ .TilingData(param.get())
93+ .Workspace(ws_size)
94+ .Build();
95+ gert::TilingContext* tiling_context = holder.GetContext<gert::TilingContext>();
96+ ASSERT_NE(tiling_context, nullptr);
97+ EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_SUCCESS);
98+}
99+ 
100+TEST_F(SwishTiling, swish_float16_success) {
101+ gert::StorageShape x_shape = {{1, 2, 8, 16}, {1, 2, 8, 16}};
102+ gert::StorageShape y_shape = {{1, 2, 8, 16}, {1, 2, 8, 16}};
103+ string compile_info_string = R"({
104+ "hardware_info": {"BT_SIZE": 0, "load3d_constraints": "1",
105+ "Intrinsic_fix_pipe_l0c2out": false, "Intrinsic_data_move_l12ub": true, "Intrinsic_data_move_l0c2ub": true, "Intrinsic_data_move_out2l1_nd2nz": false,
106+ "UB_SIZE": 196608, "L2_SIZE": 33554432, "L1_SIZE": 524288,
107+ "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072,
108+ "CORE_NUM": 48}
109+ })";
110+ map<string, string> soc_infos;
111+ map<string, string> aicore_spec;
112+ map<string, string> intrinsics;
113+ GetPlatFormInfos(compile_info_string.c_str(), soc_infos, aicore_spec, intrinsics);
114+ 
115+ fe::PlatFormInfos platform_info;
116+ platform_info.Init();
117+ 
118+ struct SwishCompileInfo {};
119+ SwishCompileInfo compile_info;
120+ 
121+ std::string op_type("Swish");
122+ auto tiling_func = gert::OpImplRegistry::GetInstance().GetOpImpl(op_type.c_str())->tiling;
123+ 
124+ auto kernel_holder =
125+ gert::KernelRunContextFaker()
126+ .KernelIONum(3, 1)
127+ .Inputs({const_cast<char*>(compile_info_string.c_str()), reinterpret_cast<void*>(&platform_info)})
128+ .Outputs({&compile_info})
129+ .Build();
130+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", soc_infos);
131+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicore_spec);
132+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
133+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap",
134+ intrinsics);
135+ 
136+ auto param = gert::TilingData::CreateCap(4096);
137+ auto workspace_size_holer = gert::ContinuousVector::Create<size_t>(4096);
138+ auto ws_size = reinterpret_cast<gert::ContinuousVector*>(workspace_size_holer.get());
139+ ASSERT_NE(param, nullptr);
140+ auto holder = gert::TilingContextFaker()
141+ .SetOpType("Swish")
142+ .NodeIoNum(1, 1)
143+ .IrInstanceNum({1})
144+ .InputShapes({&x_shape})
145+ .OutputShapes({&y_shape})
146+ .CompileInfo(&compile_info)
147+ .PlatformInfo(reinterpret_cast<char*>(&platform_info))
148+ .NodeAttrs({{"scale", Ops::NN::AnyValue::CreateFrom<float>(1.0f)}})
149+ .NodeInputTd(0, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
150+ .NodeOutputTd(0, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
151+ .TilingData(param.get())
152+ .Workspace(ws_size)
153+ .Build();
154+ gert::TilingContext* tiling_context = holder.GetContext<gert::TilingContext>();
155+ ASSERT_NE(tiling_context, nullptr);
156+ EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_SUCCESS);
157+}