Star996
1926
代码
代码
Issues248
Pull Requests353
Actions
项目讨论
Wiki
分析
Star996
1926
已合并
SiluMul 算子AscendC实现贡献 #567
wuxs68创建于 2025年12月30日
SiluMul 算子AscendC实现贡献 #567
已合并
wuxs68创建于 2025年12月30日
从wuxs68/ops-nnmaster合入到cann/ops-nnmaster
共 27 个文件变更+2229-0
Aexperimental/activation/silu_mul/CMakeLists.txt+18-0
@@ -0,0 +1,18 @@
1+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
2+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
3+# CANN Open Software License Agreement Version 2.0 (the "License").
4+# Please refer to the License for details. You may not use this file except in compliance with the License.
5+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
6+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
7+# See LICENSE in the root of the software repository for the full text of the License.
8+#/
9+message(STATUS "=== Debug: start ops.activation.silu_mul.CMakeLists.txt ")
10+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
11+if(NOT ENABLE_TEST AND NOT BENCHMARK)
12+ list(REMOVE_ITEM CURRENT_DIRS tests)
13+endif()
14+foreach(SUB_DIR ${CURRENT_DIRS})
15+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
16+ add_subdirectory(${SUB_DIR})
17+ endif()
18+endforeach()
Aexperimental/activation/silu_mul/README.md+73-0
@@ -0,0 +1,73 @@
1+# SiluMul
2+ 
3+## 产品支持情况
4+ 
5+|产品 | 是否支持 |
6+|:-------------------------|:----------:|
7+| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |
8+| <term>Atlas A2 训练系列产品/Atlas 800I A2 推理产品/A200I A2 Box 异构组件</term> | √ |
9+ 
10+## 功能说明
11+ 
12+- 算子功能：
13+ 
14+ 对输入Tensor x进行Silu激活计算，将计算结果与输入Tensor y相乘。
15+ 
16+- 计算公式：
17+
18+ 给定输入张量 `x` 和 `y`，函数 `SiluMul` 进行以下计算：
19+ 
20+ 1. 对 `x` 应用 SiLU (Sigmoid Linear Unit) 激活函数：
21+ $$
22+ \text{SiLU}(x) = x \cdot \text{Sigmoid}(x) = \frac{x}{1 + e^{-x}}
23+ $$
24+ 
25+ 2. 最终输出是 SiLU(x) 和 y 的逐元素乘积：
26+ $$
27+ \text{out} = \text{SiLU}(x) \times y
28+ $$
29+ 
30+## 参数说明
31+ 
32+<table style="undefined;table-layout: fixed; width: 970px"><colgroup>
33+ <col style="width: 181px">
34+ <col style="width: 144px">
35+ <col style="width: 273px">
36+ <col style="width: 256px">
37+ <col style="width: 116px">
38+ </colgroup>
39+ <thead>
40+ <tr>
41+ <th>参数名</th>
42+ <th>输入/输出/属性</th>
43+ <th>描述</th>
44+ <th>数据类型</th>
45+ <th>数据格式</th>
46+ </tr></thead>
47+ <tbody>
48+ <tr>
49+ <td>input</td>
50+ <td>输入</td>
51+ <td>公式中的输入input。</td>
52+ <td>BFLOAT16、FLOAT16、FLOAT</td>
53+ <td>ND</td>
54+ </tr>
55+ <tr>
56+ <td>out</td>
57+ <td>输出</td>
58+ <td>公式中的out。</td>
59+ <td>BFLOAT16、FLOAT16、FLOAT</td>
60+ <td>ND</td>
61+ </tr>
62+ </tbody></table>
63+ 
64+## 约束说明
65+ 
66+典型场景尾轴为16的倍数，当尾轴为非32B对齐时，建议走小算子拼接逻辑。
67+ 
68+## 调用说明
69+ 
70+| 调用方式 | 调用样例 | 说明 |
71+|--------------|------------------------------------------------------------------------|----------------------------------------------------------------|
72+| aclnn调用 | [test_aclnn_silu_mul](./examples/test_aclnn_silu_mul.cpp) | 通过[aclnnSiluMul](./docs/aclnnSiluMul.md)接口方式调用SiluMul算子。 |
73+| 图模式调用 | - | 通过[算子IR](./op_graph/silu_mul_proto.h)构图方式调用SiluMul算子。 |
Aexperimental/activation/silu_mul/docs/aclnnSiluMul.md+359-0
@@ -0,0 +1,359 @@
1+# aclnnSiluMul
2+ 
3+## 产品支持情况
4+ 
5+|产品 | 是否支持 |
6+|:-------------------------|:----------:|
7+| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |
8+| <term>Atlas A2 训练系列产品/Atlas 800I A2 推理产品/A200I A2 Box 异构组件</term> | √ |
9+ 
10+## 功能说明
11+ 
12+- 算子功能：
13+ 
14+ 对输入Tensor x进行Silu激活计算，将计算结果与输入Tensor y相乘。
15+ 
16+- 计算公式：
17+
18+ 给定输入张量 `x` 和 `y`，函数 `SiluMul` 进行以下计算：
19+ 
20+ 1. 对 `x` 应用 SiLU (Sigmoid Linear Unit) 激活函数：
21+ $$
22+ \text{SiLU}(x) = x \cdot \text{Sigmoid}(x) = \frac{x}{1 + e^{-x}}
23+ $$
24+ 
25+ 2. 最终输出是 SiLU(x) 和 y 的逐元素乘积：
26+ $$
27+ \text{out} = \text{SiLU}(x) \times y
28+ $$
29+ 
30+## 函数原型
31+ 
32+每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnSiluMulGetWorkspaceSize”接口获取入参并根据计算流程计算所需workspace大小，再调用“aclnnSiluMul”接口执行计算。
33+```Cpp
34+aclnnStatus aclnnSiluMulGetWorkspaceSize(
35+ const aclTensor *x,
36+ const aclTensor *y,
37+ const aclTensor *out,
38+ uint64_t *workspaceSize,
39+ aclOpExecutor **executor)
40+```
41+ 
42+```cpp
43+aclnnStatus aclnnSiluMul(
44+ void *workspace,
45+ uint64_t workspaceSize,
46+ aclOpExecutor *executor,
47+ aclrtStream stream)
48+```
49+ 
50+## aclnnSiluMulGetWorkspaceSize
51+ 
52+- **参数说明：**
53+ 
54+ <table style="undefined;table-layout: fixed; width: 1458px"><colgroup>
55+ <col style="width: 154px">
56+ <col style="width: 120px">
57+ <col style="width: 276px">
58+ <col style="width: 308px">
59+ <col style="width: 212px">
60+ <col style="width: 107px">
61+ <col style="width: 136px">
62+ <col style="width: 145px">
63+ </colgroup>
64+ <thead>
65+ <tr>
66+ <th>参数名</th>
67+ <th>输入/输出</th>
68+ <th>描述</th>
69+ <th>使用说明</th>
70+ <th>数据类型</th>
71+ <th>数据格式</th>
72+ <th>维度(shape)</th>
73+ <th>非连续Tensor</th>
74+ </tr></thead>
75+ <tbody>
76+ <tr>
77+ <td>x</td>
78+ <td>输入</td>
79+ <td>输入的张量，公式中的x。</td>
80+ <td><ul><li>支持空Tensor。</li><li>数据类型需与y一致。</li><li>shape需与y一致。</li></ul></td>
81+ <td>BFLOAT16、FLOAT16、FLOAT</td>
82+ <td>ND</td>
83+ <td>2-8</td>
84+ <td>√</td>
85+ </tr>
86+ <tr>
87+ <td>y</td>
88+ <td>输入</td>
89+ <td>输入的张量，公式中的y。</td>
90+ <td><ul><li>支持空Tensor。</li><li>数据类型需与x一致。</li><li>shape需与x一致。</li></ul></td>
91+ <td>BFLOAT16、FLOAT16、FLOAT</td>
92+ <td>ND</td>
93+ <td>2-8</td>
94+ <td>√</td>
95+ </tr>
96+ <tr>
97+ <td>out</td>
98+ <td>输出</td>
99+ <td>输出的张量，公式中的out。</td>
100+ <td><ul><li>输出的数据类型与输入保持一致。</li><li>输出shape与输入shape一致。</li></ul></td>
101+ <td>BFLOAT16、FLOAT16、FLOAT</td>
102+ <td>ND</td>
103+ <td>2-8</td>
104+ <td>√</td>
105+ </tr>
106+ <tr>
107+ <td>workspaceSize</td>
108+ <td>输出</td>
109+ <td>返回需要在Device侧申请的workspace大小。</td>
110+ <td>-</td>
111+ <td>-</td>
112+ <td>-</td>
113+ <td>-</td>
114+ <td>-</td>
115+ </tr>
116+ <tr>
117+ <td>executor</td>
118+ <td>输出</td>
119+ <td>返回op执行器，包含了算子计算流程。</td>
120+ <td>-</td>
121+ <td>-</td>
122+ <td>-</td>
123+ <td>-</td>
124+ <td>-</td>
125+ </tr>
126+ </tbody>
127+ </table>
128+ 
129+ 
130+- **返回值：**
131+ 
132+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
133+ 第一段接口会完成入参校验，出现以下场景时报错：
134+ <table style="undefined;table-layout: fixed;width: 979px"><colgroup>
135+ <col style="width: 272px">
136+ <col style="width: 103px">
137+ <col style="width: 604px">
138+ </colgroup>
139+ <thead>
140+ <tr>
141+ <th>返回码</th>
142+ <th>错误码</th>
143+ <th>描述</th>
144+ </tr>
145+ </thead>
146+ <tbody>
147+ <tr>
148+ <td>ACLNN_ERR_PARAM_NULLPTR</td>
149+ <td>161001</td>
150+ <td>传入的x、y或out是空指针。</td>
151+ </tr>
152+ <tr>
153+ <td rowspan="8">ACLNN_ERR_PARAM_INVALID</td>
154+ <td rowspan="8">161002</td>
155+ <td>input的数据类型不在支持的范围之内。</td>
156+ </tr>
157+ </tbody></table>
158+ 
159+## aclnnSiluMul
160+ 
161+- **参数说明：**
162+ 
163+ <table style="undefined;table-layout: fixed; width: 953px"><colgroup>
164+ <col style="width: 173px">
165+ <col style="width: 112px">
166+ <col style="width: 668px">
167+ </colgroup>
168+ <thead>
169+ <tr>
170+ <th>参数名</th>
171+ <th>输入/输出</th>
172+ <th>描述</th>
173+ </tr></thead>
174+ <tbody>
175+ <tr>
176+ <td>workspace</td>
177+ <td>输入</td>
178+ <td>在Device侧申请的workspace内存地址。</td>
179+ </tr>
180+ <tr>
181+ <td>workspaceSize</td>
182+ <td>输入</td>
183+ <td>在Device侧申请的workspace大小，由第一段接口aclnnSiluMulGetWorkspaceSize获取。</td>
184+ </tr>
185+ <tr>
186+ <td>executor</td>
187+ <td>输入</td>
188+ <td>op执行器，包含了算子计算流程。</td>
189+ </tr>
190+ <tr>
191+ <td>stream</td>
192+ <td>输入</td>
193+ <td>指定执行任务的Stream。</td>
194+ </tr>
195+ </tbody>
196+ </table>
197+ 
198+- **返回值：**
199+aclnnStatus：返回状态码，具体参见[可疑链接已删除]。
200+ 
201+## 约束说明
202+ 
203+无。
204+ 
205+## 调用示例
206+ 
207+示例代码如下，仅供参考，具体编译和执行过程请参考[编译与运行样例](../../../docs/zh/context/编译与运行样例.md)。
208+ 
209+```cpp
210+#include <iostream>
211+#include <vector>
212+#include "acl/acl.h"
213+#include "aclnnop/aclnn_silu_mul.h"
214+ 
215+#define CHECK_RET(cond, return_expr) \
216+ do { \
217+ if (!(cond)) { \
218+ return_expr; \
219+ } \
220+ } while (0)
221+ 
222+#define LOG_PRINT(message, ...) \
223+ do { \
224+ printf(message, ##__VA_ARGS__); \
225+ } while (0)
226+ 
227+int64_t GetShapeSize(const std::vector<int64_t>& shape) {
228+ int64_t shapeSize = 1;
229+ for (auto i : shape) {
230+ shapeSize *= i;
231+ }
232+ return shapeSize;
233+}
234+ 
235+void PrintOutResult(std::vector<int64_t> &shape, void** deviceAddr) {
236+ auto size = GetShapeSize(shape);
237+ std::vector<float> resultData(size, 0);
238+ auto ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]),
239+ *deviceAddr, size * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
240+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return);
241+ for (int64_t i = 0; i < size; i++) {
242+ LOG_PRINT("mean result[%ld] is: %f\n", i, resultData[i]);
243+ }
244+}
245+ 
246+int Init(int32_t deviceId, aclrtStream* stream) {
247+ // 固定写法，资源初始化
248+ auto ret = aclInit(nullptr);
249+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
250+ ret = aclrtSetDevice(deviceId);
251+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
252+ ret = aclrtCreateStream(stream);
253+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
254+ return 0;
255+}
256+ 
257+template <typename T>
258+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
259+ aclDataType dataType, aclTensor** tensor) {
260+ auto size = GetShapeSize(shape) * sizeof(T);
261+ // 调用aclrtMalloc申请device侧内存
262+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
263+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
264+ // 调用aclrtMemcpy将host侧数据复制到device侧内存上
265+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
266+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
267+ 
268+ // 计算连续tensor的strides
269+ std::vector<int64_t> strides(shape.size(), 1);
270+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
271+ strides[i] = shape[i + 1] * strides[i + 1];
272+ }
273+ 
274+ // 调用aclCreateTensor接口创建aclTensor
275+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
276+ shape.data(), shape.size(), *deviceAddr);
277+ return 0;
278+}
279+ 
280+int main() {
281+ // 1. （固定写法）device/stream初始化，参考acl API手册
282+ // 根据自己的实际device填写deviceId
283+ int32_t deviceId = 0;
284+ aclrtStream stream;
285+ auto ret = Init(deviceId, &stream);
286+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
287+ 
288+ // 2. 构造输入与输出，需要根据API的接口自定义构造
289+ std::vector<int64_t> inputXShape = {2, 4};
290+ std::vector<int64_t> inputYShape = {2, 4};
291+ std::vector<int64_t> outShape = {2, 4};
292+ 
293+ std::vector<float> inputXHostData = {0, 1, 2, 3, 4, 5, 6, 7};
294+ std::vector<float> inputYHostData = {1, 1, 1, 1, 1, 1, 1, 1};
295+ std::vector<float> outHostData(8, 0);
296+ 
297+ void* inputXDeviceAddr = nullptr;
298+ void* inputYDeviceAddr = nullptr;
299+ void* outDeviceAddr = nullptr;
300+ 
301+ aclTensor* inputX = nullptr;
302+ aclTensor* inputY = nullptr;
303+ aclTensor* out = nullptr;
304+ 
305+ // 创建inputX aclTensor
306+ ret = CreateAclTensor(inputXHostData, inputXShape, &inputXDeviceAddr, aclDataType::ACL_FLOAT, &inputX);
307+ CHECK_RET(ret == ACL_SUCCESS, return ret);
308+ // 创建inputY aclTensor
309+ ret = CreateAclTensor(inputYHostData, inputYShape, &inputYDeviceAddr, aclDataType::ACL_FLOAT, &inputY);
310+ CHECK_RET(ret == ACL_SUCCESS, return ret);
311+ // 创建out aclTensor
312+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
313+ CHECK_RET(ret == ACL_SUCCESS, return ret);
314+ 
315+ // 3. 调用CANN算子库API，需要修改为具体的Api名称
316+ uint64_t workspaceSize = 0;
317+ aclOpExecutor* executor;
318+ 
319+ // 调用aclnnSiluMul第一段接口
320+ ret = aclnnSiluMulGetWorkspaceSize(inputX, inputY, out, &workspaceSize, &executor);
321+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSiluMulGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
322+ 
323+ // 根据第一段接口计算出的workspaceSize申请device内存
324+ void* workspaceAddr = nullptr;
325+ if (workspaceSize > 0) {
326+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
327+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
328+ }
329+ 
330+ // 调用aclnnSiluMul第二段接口
331+ ret = aclnnSiluMul(workspaceAddr, workspaceSize, executor, stream);
332+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSiluMul failed. ERROR: %d\n", ret); return ret);
333+ 
334+ // 4. （固定写法）同步等待任务执行结束
335+ ret = aclrtSynchronizeStream(stream);
336+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
337+ 
338+ // 5. 获取输出的值，将device侧内存上的结果复制至host侧，需要根据具体API的接口定义修改
339+ PrintOutResult(outShape, &outDeviceAddr);
340+ 
341+ // 6. 释放aclTensor和aclTensor，需要根据具体API的接口定义修改
342+ aclDestroyTensor(inputX);
343+ aclDestroyTensor(inputY);
344+ aclDestroyTensor(out);
345+ 
346+ // 7.释放device资源，需要根据具体API的接口定义修改
347+ aclrtFree(inputXDeviceAddr);
348+ aclrtFree(inputYDeviceAddr);
349+ aclrtFree(outDeviceAddr);
350+ if (workspaceSize > 0) {
351+ aclrtFree(workspaceAddr);
352+ }
353+ aclrtDestroyStream(stream);
354+ aclrtResetDevice(deviceId);
355+ aclFinalize();
356+ 
357+ return 0;
358+}
359+```
Aexperimental/activation/silu_mul/examples/test_aclnn_silu_mul.cpp+153-0
@@ -0,0 +1,153 @@
1+/*
2+ * Copyright (c) 2025 联通（广东）产业互联网有限公司.
3+ *
4+ * Licensed under the Apache License, Version 2.0 (the "License");
5+ * you may not use this file except in compliance with the License.
6+ * You may obtain a copy of the License at
7+ *
8+ * http://www.apache.org/licenses/LICENSE-2.0
9+ *
10+ * Unless required by applicable law or agreed to in writing, software
11+ * distributed under the License is distributed on an "AS IS" BASIS,
12+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13+ * See the License for the specific language governing permissions and
14+ * limitations under the License.
15+ */
16+#include <iostream>
17+#include <vector>
18+#include "acl/acl.h"
19+#include "aclnnop/aclnn_silu_mul.h"
20+ 
21+#define CHECK_RET(cond, return_expr) \
22+ do { \
23+ if (!(cond)) { \
24+ return_expr; \
25+ } \
26+ } while (0)
27+ 
28+#define LOG_PRINT(message, ...) \
29+ do { \
30+ printf(message, ##__VA_ARGS__); \
31+ } while (0)
32+ 
33+int64_t GetShapeSize(const std::vector<int64_t>& shape)
34+{
35+ int64_t shapeSize = 1;
36+ for (auto i : shape) {
37+ shapeSize *= i;
38+ }
39+ return shapeSize;
40+}
41+ 
42+void PrintOutResult(std::vector<int64_t>& shape, void** deviceAddr)
43+{
44+ auto size = GetShapeSize(shape);
45+ std::vector<float> resultData(size, 0);
46+ auto ret = aclrtMemcpy(
47+ resultData.data(), resultData.size() * sizeof(resultData[0]), *deviceAddr, size * sizeof(resultData[0]),
48+ ACL_MEMCPY_DEVICE_TO_HOST);
49+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return);
50+ for (int64_t i = 0; i < size; i++) {
51+ LOG_PRINT("mean result[%ld] is: %f\n", i, resultData[i]);
52+ }
53+}
54+ 
55+int Init(int32_t deviceId, aclrtStream* stream)
56+{
57+ auto ret = aclInit(nullptr);
58+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
59+ ret = aclrtSetDevice(deviceId);
60+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
61+ ret = aclrtCreateStream(stream);
62+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
63+ return 0;
64+}
65+ 
66+template <typename T>
67+int CreateAclTensor(
68+ const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr, aclDataType dataType,
69+ aclTensor** tensor)
70+{
71+ auto size = GetShapeSize(shape) * sizeof(T);
72+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
73+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
74+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
75+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
76+ 
77+ std::vector<int64_t> strides(shape.size(), 1);
78+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
79+ strides[i] = shape[i + 1] * strides[i + 1];
80+ }
81+ 
82+ *tensor = aclCreateTensor(
83+ shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND, shape.data(), shape.size(),
84+ *deviceAddr);
85+ return 0;
86+}
87+ 
88+int main()
89+{
90+ int32_t deviceId = 0;
91+ aclrtStream stream;
92+ auto ret = Init(deviceId, &stream);
93+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
94+ 
95+ std::vector<int64_t> inputXShape = {2, 4};
96+ std::vector<int64_t> inputYShape = {2, 4};
97+ std::vector<int64_t> outShape = {2, 4};
98+ 
99+ std::vector<float> inputXHostData = {0, 1, 2, 3, 4, 5, 6, 7};
100+ std::vector<float> inputYHostData = {1, 1, 1, 1, 1, 1, 1, 1};
101+ std::vector<float> outHostData(8, 0);
102+ 
103+ void* inputXDeviceAddr = nullptr;
104+ void* inputYDeviceAddr = nullptr;
105+ void* outDeviceAddr = nullptr;
C
Cchenqi3172025年12月30日

补充测试结果

likedislike
W
wuxs68
2025年12月30日 评论：
106+ 
107+ aclTensor* inputX = nullptr;
108+ aclTensor* inputY = nullptr;
109+ aclTensor* out = nullptr;
110+ 
111+ ret = CreateAclTensor(inputXHostData, inputXShape, &inputXDeviceAddr, aclDataType::ACL_FLOAT, &inputX);
112+ CHECK_RET(ret == ACL_SUCCESS, return ret);
113+ ret = CreateAclTensor(inputYHostData, inputYShape, &inputYDeviceAddr, aclDataType::ACL_FLOAT, &inputY);
114+ CHECK_RET(ret == ACL_SUCCESS, return ret);
115+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
116+ CHECK_RET(ret == ACL_SUCCESS, return ret);
117+ 
118+ uint64_t workspaceSize = 0;
119+ aclOpExecutor* executor;
120+ 
121+ ret = aclnnSiluMulGetWorkspaceSize(inputX, inputY, out, &workspaceSize, &executor);
122+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSiluMulGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
123+ 
124+ void* workspaceAddr = nullptr;
125+ if (workspaceSize > 0) {
126+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
127+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
128+ }
129+ 
130+ ret = aclnnSiluMul(workspaceAddr, workspaceSize, executor, stream);
131+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSiluMul failed. ERROR: %d\n", ret); return ret);
132+ 
133+ ret = aclrtSynchronizeStream(stream);
134+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
135+ 
136+ PrintOutResult(outShape, &outDeviceAddr);
137+ 
138+ aclDestroyTensor(inputX);
139+ aclDestroyTensor(inputY);
140+ aclDestroyTensor(out);
141+ 
142+ aclrtFree(inputXDeviceAddr);
143+ aclrtFree(inputYDeviceAddr);
144+ aclrtFree(outDeviceAddr);
145+ if (workspaceSize > 0) {
146+ aclrtFree(workspaceAddr);
147+ }
148+ aclrtDestroyStream(stream);
149+ aclrtResetDevice(deviceId);
150+ aclFinalize();
151+ 
152+ return 0;
153+}
Aexperimental/activation/silu_mul/op_graph/CMakeLists.txt+10-0
@@ -0,0 +1,10 @@
1+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
2+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
3+# CANN Open Software License Agreement Version 2.0 (the "License").
4+# Please refer to the License for details. You may not use this file except in compliance with the License.
5+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
6+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
7+# See LICENSE in the root of the software repository for the full text of the License.
8+ 
9+message(STATUS "=== Debug: start ops.activation.silu_mul.graph_plugin.CMakeLists.txt ")
10+add_graph_plugin_sources()
Aexperimental/activation/silu_mul/op_graph/silu_mul_proto.h+48-0
@@ -0,0 +1,48 @@
1+/*
2+ * Copyright (c) 2025 联通（广东）产业互联网有限公司.
3+ *
4+ * Licensed under the Apache License, Version 2.0 (the "License");
5+ * you may not use this file except in compliance with the License.
6+ * You may obtain a copy of the License at
7+ *
8+ * http://www.apache.org/licenses/LICENSE-2.0
9+ *
10+ * Unless required by applicable law or agreed to in writing, software
11+ * distributed under the License is distributed on an "AS IS" BASIS,
12+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13+ * See the License for the specific language governing permissions and
14+ * limitations under the License.
15+ */
16+ 
17+/*!
18+ * \file nn_activation.h
19+ * \brief
20+ */
21+ 
22+#ifndef OPS_BUILT_IN_OP_PROTO_INC_NN_SILU_MUL_H_
23+#define OPS_BUILT_IN_OP_PROTO_INC_NN_SILU_MUL_H_
24+ 
25+#include "graph/operator_reg.h"
26+ 
27+namespace ge {
28+/**
29+ * @brief Performs Silu multiplication. z = silu(x) * y. \n
30+ 
31+ * @par Inputs:
32+ * x: A tensor of type float, float16 or bfloat16. Shape support 2D ~ 8D.
33+ * The format must be ND.
34+ * y: A tensor of type float, float16 or bfloat16. Shape support 2D ~ 8D.
35+ * The format must be ND.
36+ 
37+ * @par Outputs:
38+ * z: A tensor has the same type and format as "x".
39+ * Other dimensions of its shape are the same as those of "x". \n
40+ */
41+REG_OP(SiluMul)
42+ .INPUT(x, TensorType({DT_BF16, DT_FLOAT16, DT_FLOAT}))
43+ .INPUT(y, TensorType({DT_BF16, DT_FLOAT16, DT_FLOAT}))
44+ .OUTPUT(z, TensorType({DT_BF16, DT_FLOAT16, DT_FLOAT}))
45+ .OP_END_FACTORY_REG(SiluMul)
46+ 
47+} // namespace ge
48+#endif // OPS_BUILT_IN_OP_PROTO_INC_NN_SILU_MUL_H_
Aexperimental/activation/silu_mul/op_host/CMakeLists.txt+11-0
@@ -0,0 +1,11 @@
1+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
2+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
3+# CANN Open Software License Agreement Version 2.0 (the "License").
4+# Please refer to the License for details. You may not use this file except in compliance with the License.
5+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
6+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
7+# See LICENSE in the root of the software repository for the full text of the License.
8+#/
9+ 
10+message(STATUS "=== Debug: start ops.activation.op_host.CMakeLists.txt ")
11+add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE silu_mul ACLNNTYPE aclnn)
Aexperimental/activation/silu_mul/op_host/config/ascend910_93/silu_mul_binary.json+116-0
@@ -0,0 +1,116 @@
1+{
2+ "op_type": "SiluMul",
3+ "op_list": [
4+ {
5+ "bin_filename": "SiluMul_float32",
6+ "inputs": [
7+ {
8+ "name": "x",
9+ "index": 0,
10+ "dtype": "float32",
11+ "format": "ND",
12+ "paramType": "required",
13+ "shape": [
14+ -2
15+ ]
16+ },
17+ {
18+ "name": "y",
19+ "index": 1,
20+ "dtype": "float32",
21+ "format": "ND",
22+ "paramType": "required",
23+ "shape": [
24+ -2
25+ ]
26+ }
27+ ],
28+ "outputs": [
29+ {
30+ "name": "z",
31+ "index": 0,
32+ "dtype": "float32",
33+ "format": "ND",
34+ "paramType": "required",
35+ "shape": [
36+ -2
37+ ]
38+ }
39+ ]
40+ },
41+ {
42+ "bin_filename": "SiluMul_bfloat16",
43+ "inputs": [
44+ {
45+ "name": "x",
46+ "index": 0,
47+ "dtype": "bfloat16",
48+ "format": "ND",
49+ "paramType": "required",
50+ "shape": [
51+ -2
52+ ]
53+ },
54+ {
55+ "name": "y",
56+ "index": 1,
57+ "dtype": "bfloat16",
58+ "format": "ND",
59+ "paramType": "required",
60+ "shape": [
61+ -2
62+ ]
63+ }
64+ ],
65+ "outputs": [
66+ {
67+ "name": "z",
68+ "index": 0,
69+ "dtype": "bfloat16",
70+ "format": "ND",
71+ "paramType": "required",
72+ "shape": [
73+ -2
74+ ]
75+ }
76+ ]
77+ },
78+ {
79+ "bin_filename": "SiluMul_float16",
80+ "inputs": [
81+ {
82+ "name": "x",
83+ "index": 0,
84+ "dtype": "float16",
85+ "format": "ND",
86+ "paramType": "required",
87+ "shape": [
88+ -2
89+ ]
90+ },
91+ {
92+ "name": "y",
93+ "index": 1,
94+ "dtype": "float16",
95+ "format": "ND",
96+ "paramType": "required",
97+ "shape": [
98+ -2
99+ ]
100+ }
101+ ],
102+ "outputs": [
103+ {
104+ "name": "z",
105+ "index": 0,
106+ "dtype": "float16",
107+ "format": "ND",
108+ "paramType": "required",
109+ "shape": [
110+ -2
111+ ]
112+ }
113+ ]
114+ }
115+ ]
116+}
Aexperimental/activation/silu_mul/op_host/config/ascend910_93/silu_mul_simplified_key.ini+14-0
@@ -0,0 +1,14 @@
1+; 该文件主要影响 opc 工具 编译二进制kernel时， --simplified_key_mode 选项中填写的值，格式如下所示：
2+; [某算子]
3+; default=xx
4+; ascendxx=xx
5+; 其中，default为默认mode，ascendxx为可选mode，如果不同芯片有差异化要求时，需要配置；
6+; 1)如果没有配置：非ascendC算子继续按空处理，即opc编译命令中不添加 --simplified_key_mode 选项，AscendC算子按照 simplified_key_mode=0 处理
7+; 2)如果仅有default配置：各个版本按default配置
8+; 3)如果仅有某些平台的配置，没有default配置：对应平台的按照配置的值传递，非对应平台的：非AscendC算子继续按空处理，AscendC算子按照 simplified_key_mode=0 处理
9+; 4)如果default配置和平台配置都有：对应平台的使用平台的配置，非对应的平台的以default值配置。
10+; 5)对于自定义simplified key的情况，需要在binary_simplified_key_mode.ini 文件中显式配置为None，不传入 --simplified_key_mode 选项，由opc工具和FE框架自行判断使用何种模式
11+; 6)是否是AscendC算子，由 ops/build-in/tbe/op_info_cfg/parser/ascendc_config.json 中配置的算子名字和对于的平台决定
12+ 
13+[SiluMul]
14+default=0
Aexperimental/activation/silu_mul/op_host/config/ascend910b/silu_mul_binary.json+116-0
@@ -0,0 +1,116 @@
1+{
2+ "op_type": "SiluMul",
3+ "op_list": [
4+ {
5+ "bin_filename": "SiluMul_float32",
6+ "inputs": [
7+ {
8+ "name": "x",
9+ "index": 0,
10+ "dtype": "float32",
11+ "format": "ND",
12+ "paramType": "required",
13+ "shape": [
14+ -2
15+ ]
16+ },
17+ {
18+ "name": "y",
19+ "index": 1,
20+ "dtype": "float32",
21+ "format": "ND",
22+ "paramType": "required",
23+ "shape": [
24+ -2
25+ ]
26+ }
27+ ],
28+ "outputs": [
29+ {
30+ "name": "z",
31+ "index": 0,
32+ "dtype": "float32",
33+ "format": "ND",
34+ "paramType": "required",
35+ "shape": [
36+ -2
37+ ]
38+ }
39+ ]
40+ },
41+ {
42+ "bin_filename": "SiluMul_bfloat16",
43+ "inputs": [
44+ {
45+ "name": "x",
46+ "index": 0,
47+ "dtype": "bfloat16",
48+ "format": "ND",
49+ "paramType": "required",
50+ "shape": [
51+ -2
52+ ]
53+ },
54+ {
55+ "name": "y",
56+ "index": 1,
57+ "dtype": "bfloat16",
58+ "format": "ND",
59+ "paramType": "required",
60+ "shape": [
61+ -2
62+ ]
63+ }
64+ ],
65+ "outputs": [
66+ {
67+ "name": "z",
68+ "index": 0,
69+ "dtype": "bfloat16",
70+ "format": "ND",
71+ "paramType": "required",
72+ "shape": [
73+ -2
74+ ]
75+ }
76+ ]
77+ },
78+ {
79+ "bin_filename": "SiluMul_float16",
80+ "inputs": [
81+ {
82+ "name": "x",
83+ "index": 0,
84+ "dtype": "float16",
85+ "format": "ND",
86+ "paramType": "required",
87+ "shape": [
88+ -2
89+ ]
90+ },
91+ {
92+ "name": "y",
93+ "index": 1,
94+ "dtype": "float16",
95+ "format": "ND",
96+ "paramType": "required",
97+ "shape": [
98+ -2
99+ ]
100+ }
101+ ],
102+ "outputs": [
103+ {
104+ "name": "z",
105+ "index": 0,
106+ "dtype": "float16",
107+ "format": "ND",
108+ "paramType": "required",
109+ "shape": [
110+ -2
111+ ]
112+ }
113+ ]
114+ }
115+ ]
116+}
Aexperimental/activation/silu_mul/op_host/config/ascend910b/silu_mul_simplified_key.ini+14-0
@@ -0,0 +1,14 @@
1+; 该文件主要影响 opc 工具 编译二进制kernel时， --simplified_key_mode 选项中填写的值，格式如下所示：
2+; [某算子]
3+; default=xx
4+; ascendxx=xx
5+; 其中，default为默认mode，ascendxx为可选mode，如果不同芯片有差异化要求时，需要配置；
6+; 1)如果没有配置：非ascendC算子继续按空处理，即opc编译命令中不添加 --simplified_key_mode 选项，AscendC算子按照 simplified_key_mode=0 处理
7+; 2)如果仅有default配置：各个版本按default配置
8+; 3)如果仅有某些平台的配置，没有default配置：对应平台的按照配置的值传递，非对应平台的：非AscendC算子继续按空处理，AscendC算子按照 simplified_key_mode=0 处理
9+; 4)如果default配置和平台配置都有：对应平台的使用平台的配置，非对应的平台的以default值配置。
10+; 5)对于自定义simplified key的情况，需要在binary_simplified_key_mode.ini 文件中显式配置为None，不传入 --simplified_key_mode 选项，由opc工具和FE框架自行判断使用何种模式
11+; 6)是否是AscendC算子，由 ops/build-in/tbe/op_info_cfg/parser/ascendc_config.json 中配置的算子名字和对于的平台决定
12+ 
13+[SiluMul]
14+default=0
Aexperimental/activation/silu_mul/op_host/silu_mul_def.cpp+57-0
@@ -0,0 +1,57 @@
1+/*
2+ * Copyright (c) 2025 联通（广东）产业互联网有限公司.
3+ *
4+ * Licensed under the Apache License, Version 2.0 (the "License");
5+ * you may not use this file except in compliance with the License.
6+ * You may obtain a copy of the License at
7+ *
8+ * http://www.apache.org/licenses/LICENSE-2.0
9+ *
10+ * Unless required by applicable law or agreed to in writing, software
11+ * distributed under the License is distributed on an "AS IS" BASIS,
12+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13+ * See the License for the specific language governing permissions and
14+ * limitations under the License.
15+ */
16+ 
17+/*!
18+ * \file silu_mul.cpp
19+ * \brief aicore info for silu mul op
20+ */
21+ 
22+#include "register/op_def_registry.h"
23+ 
24+namespace ops {
25+ 
26+class SiluMul : public OpDef {
27+public:
28+ explicit SiluMul(const char* name) : OpDef(name)
29+ {
30+ this->Input("x")
31+ .ParamType(REQUIRED)
32+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})
33+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
34+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
35+ .AutoContiguous();
36+ 
37+ this->Input("y")
38+ .ParamType(REQUIRED)
39+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})
40+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
41+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
42+ .AutoContiguous();
43+ 
44+ this->Output("z")
45+ .ParamType(REQUIRED)
46+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})
47+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
48+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
49+ .AutoContiguous();
50+ 
51+ this->AICore().AddConfig("ascend910b");
52+ this->AICore().AddConfig("ascend910_93");
53+ }
54+};
55+ 
56+OP_ADD(SiluMul);
57+} // namespace ops
Aexperimental/activation/silu_mul/op_host/silu_mul_infershape.cpp+60-0
@@ -0,0 +1,60 @@
1+/*
2+ * Copyright (c) 2025 联通（广东）产业互联网有限公司.
3+ *
4+ * Licensed under the Apache License, Version 2.0 (the "License");
5+ * you may not use this file except in compliance with the License.
6+ * You may obtain a copy of the License at
7+ *
8+ * http://www.apache.org/licenses/LICENSE-2.0
9+ *
10+ * Unless required by applicable law or agreed to in writing, software
11+ * distributed under the License is distributed on an "AS IS" BASIS,
12+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13+ * See the License for the specific language governing permissions and
14+ * limitations under the License.
15+ */
16+ 
17+/*!
18+ * \file silu_mul.cc
19+ * \brief
20+ */
21+#include "log/log.h"
22+#include "register/op_impl_registry.h"
23+ 
24+using namespace ge;
25+ 
26+namespace ops {
27+ 
28+static constexpr int64_t IDX_0 = 0;
29+ 
30+static ge::graphStatus InferShape4SiluMul(gert::InferShapeContext* context)
31+{
32+ OP_LOGD(context->GetNodeName(), "Begin to do InferShape4SiluMul");
33+ 
34+ auto xShape = context->GetInputShape(IDX_0);
35+ OP_CHECK_NULL_WITH_CONTEXT(context, xShape);
36+ 
37+ auto zShape = context->GetOutputShape(IDX_0);
38+ OP_CHECK_NULL_WITH_CONTEXT(context, zShape);
39+ 
40+ *zShape = *xShape;
41+ 
42+ OP_LOGD(context->GetNodeName(), "End to do InferShape4SiluMul");
43+ return GRAPH_SUCCESS;
44+}
45+ 
46+static graphStatus InferDataType4SiluMul(gert::InferDataTypeContext* context)
47+{
48+ OP_LOGD(context->GetNodeName(), "Begin to do InferDataType4SiluMul");
49+ 
50+ auto input_dtype = context->GetInputDataType(IDX_0);
51+ 
52+ context->SetOutputDataType(IDX_0, input_dtype);
53+ 
54+ OP_LOGD(context->GetNodeName(), "End to do InferDataType4SiluMul");
55+ 
56+ return GRAPH_SUCCESS;
57+}
58+ 
59+IMPL_OP_INFERSHAPE(SiluMul).InferShape(InferShape4SiluMul).InferDataType(InferDataType4SiluMul);
60+} // namespace ops
Aexperimental/activation/silu_mul/op_host/silu_mul_tiling.cpp+191-0
@@ -0,0 +1,191 @@
1+/*
2+ * Copyright (c) 2025 联通（广东）产业互联网有限公司.
3+ *
4+ * Licensed under the Apache License, Version 2.0 (the "License");
5+ * you may not use this file except in compliance with the License.
6+ * You may obtain a copy of the License at
7+ *
8+ * http://www.apache.org/licenses/LICENSE-2.0
9+ *
10+ * Unless required by applicable law or agreed to in writing, software
11+ * distributed under the License is distributed on an "AS IS" BASIS,
12+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13+ * See the License for the specific language governing permissions and
14+ * limitations under the License.
15+ */
16+ 
17+/*!
18+ * \file silu_mul_tiling.cpp
19+ * \brief silu_mul_tiling source file
20+ */
21+#include "silu_mul_tiling.h"
22+#include <vector>
23+#include <iostream>
24+#include "register/tilingdata_base.h"
25+#include "register/op_impl_registry.h"
26+#include "log/log.h"
27+#include "tiling_base/tiling_base.h"
28+#include "tiling_base/tiling_templates_registry.h"
29+ 
30+namespace optiling {
31+ 
32+static constexpr int32_t UB_SIZE = 184 * 1024;
33+static constexpr int32_t ONE_BLOCK_SIZE = 32;
34+static constexpr int32_t CALC_BUF_NUM = 8;
35+static constexpr int32_t HALF_SIZE = 2;
36+static constexpr int32_t BF16_SIZE = 2;
37+static constexpr int32_t SIZE_2 = 2;
38+// static constexpr int32_t TILING_KEY_HALF = 1;
39+// static constexpr int32_t TILING_KEY_FLOAT = 2;
40+// static constexpr int32_t TILING_KEY_BFLOAT16 = 3;
41+ 
42+static constexpr int32_t SIZE_16 = 16;
43+static constexpr int32_t LENGTH_1024 = 1024;
44+static constexpr int32_t LENGTH_LIMIT = 200000;
45+ 
46+class SiluMulTiling {
47+public:
48+ explicit SiluMulTiling(gert::TilingContext* context) : tilingContext(context) {};
49+ ge::graphStatus RunBigKernelTiling();
50+ ge::graphStatus FillTilingKey();
51+ 
52+private:
53+ ge::graphStatus ShapeCheck();
54+ ge::DataType dataType = ge::DT_UNDEFINED;
55+ gert::TilingContext* tilingContext = nullptr;
56+ gert::Shape inputShape;
57+ SiluMulTilingData tilingData;
58+ int32_t batchSize = 0;
59+ int32_t inputShapeSize = 0;
60+ int32_t lastDimSize = 0;
61+ int32_t oneBlockNum = 0;
62+ int32_t PPMaxCalNum = 0;
63+ const int32_t workspaceSize_ = SIZE_16 * LENGTH_1024 * LENGTH_1024;
64+ 
65+ static inline int32_t CeilA2B(const int32_t a, const int32_t b)
66+ {
67+ if (b != 0) {
68+ return (a + b - 1) / b;
69+ } else {
70+ return a;
71+ }
72+ }
73+ 
74+ int32_t GetNeedCoreNum(const int32_t coreNumPlatform)
75+ {
76+ int32_t needCoreNum = 1;
77+ if (lastDimSize / SIZE_2 > PPMaxCalNum) {
78+ needCoreNum = batchSize;
79+ } else {
80+ const int32_t d = lastDimSize / SIZE_2;
81+ auto dAlign = (d + oneBlockNum - 1) / oneBlockNum * oneBlockNum;
82+ const int32_t n = PPMaxCalNum / dAlign;
83+ needCoreNum = CeilA2B(batchSize, n);
84+ }
85+ if (needCoreNum == 0) {
86+ needCoreNum = 1;
87+ }
88+ if (needCoreNum >= coreNumPlatform) {
89+ return coreNumPlatform;
90+ } else {
91+ return needCoreNum;
92+ }
93+ }
94+};
95+ 
96+ge::graphStatus SiluMulTiling::ShapeCheck()
Z
Zzhajianqing1232025年12月30日

*_tiling.cpp, *_infershape.cpp, *_graph_infer.cpp中的函数要加上static

likedislike
W
wuxs68
2025年12月30日 评论：
97+{
98+ OP_CHECK_IF(
99+ (lastDimSize > LENGTH_1024),
100+ OP_LOGE(tilingContext->GetNodeName(), "Last dim size should be no more than 1024."), return ge::GRAPH_FAILED);
101+ 
102+ OP_CHECK_IF(
103+ (lastDimSize % SIZE_2 == 1), OP_LOGE(tilingContext->GetNodeName(), "Last dim size should be even."),
104+ return ge::GRAPH_FAILED);
105+ 
106+ OP_CHECK_IF(
107+ (batchSize > LENGTH_LIMIT),
108+ OP_LOGE(tilingContext->GetNodeName(), "Batch dim size should be no more than 200000."),
109+ return ge::GRAPH_FAILED);
110+ 
111+ return ge::GRAPH_SUCCESS;
112+}
113+ 
114+ge::graphStatus SiluMulTiling::RunBigKernelTiling()
Z
Zzhajianqing1232025年12月30日

*_tiling.cpp, *_infershape.cpp, *_graph_infer.cpp中的函数要加上static

likedislike
W
wuxs68
2025年12月30日 评论：
115+{
116+ auto srcTensor = tilingContext->GetInputTensor(0);
117+ OP_CHECK_NULL_WITH_CONTEXT(tilingContext, srcTensor);
118+ 
119+ PPMaxCalNum = UB_SIZE / CALC_BUF_NUM / static_cast<int32_t>(sizeof(float));
120+ 
121+ FillTilingKey();
122+ 
123+ auto srcShape = tilingContext->GetInputShape(0);
124+ inputShape = srcShape->GetOriginShape();
125+ size_t inputShapeDim = inputShape.GetDimNum();
126+ OP_CHECK_IF(
127+ (inputShapeDim < static_cast<size_t>(SIZE_2)),
128+ OP_LOGE(tilingContext->GetNodeName(), "Input shape dim should be no less than 2."), return ge::GRAPH_FAILED);
129+ lastDimSize = inputShape.GetDim(inputShapeDim - 1);
130+ inputShapeSize = inputShape.GetShapeSize();
131+ 
132+ if (lastDimSize == 0) {
133+ OP_LOGE(tilingContext->GetNodeName(),
134+ "Last dim elements can not be zero.");
135+ return ge::GRAPH_FAILED;
136+ }
137+ 
138+ batchSize = inputShapeSize / lastDimSize;
139+ 
140+ auto platformInfo = platform_ascendc::PlatformAscendC(tilingContext->GetPlatformInfo());
141+ int32_t needCoreNum = GetNeedCoreNum(platformInfo.GetCoreNumAiv());
142+ 
143+ size_t* currentWorkspace = tilingContext->GetWorkspaceSizes(1);
144+ currentWorkspace[0] = static_cast<size_t>(workspaceSize_);
145+ OP_CHECK_IF(
146+ (ShapeCheck() == ge::GRAPH_FAILED), OP_LOGE(tilingContext->GetNodeName(), "ShapeCheck failed!"),
147+ return ge::GRAPH_FAILED);
148+ tilingData.set_lastDimSize(lastDimSize);
149+ tilingData.set_batchSize(batchSize);
150+ tilingData.set_PPMaxCalNum(PPMaxCalNum);
151+ tilingData.set_needCoreNum(needCoreNum);
152+ 
153+ tilingData.SaveToBuffer(
154+ tilingContext->GetRawTilingData()->GetData(), tilingContext->GetRawTilingData()->GetCapacity());
155+ tilingContext->GetRawTilingData()->SetDataSize(tilingData.GetDataSize());
156+ 
157+ tilingContext->SetBlockDim(needCoreNum);
158+ return ge::GRAPH_SUCCESS;
159+}
160+ 
161+ge::graphStatus SiluMulTiling::FillTilingKey()
Z
Zzhajianqing1232025年12月30日

*_tiling.cpp, *_infershape.cpp, *_graph_infer.cpp中的函数要加上static

likedislike
W
wuxs68
2025年12月30日 评论：
162+{
163+ auto temp = tilingContext->GetInputDesc(0);
164+ OP_CHECK_NULL_WITH_CONTEXT(tilingContext, temp);
165+ dataType = tilingContext->GetInputDesc(0)->GetDataType();
166+ if (dataType == ge::DT_FLOAT16) {
167+ oneBlockNum = ONE_BLOCK_SIZE / HALF_SIZE;
168+ } else if (dataType == ge::DT_FLOAT) {
169+ oneBlockNum = ONE_BLOCK_SIZE / static_cast<int32_t>(sizeof(float));
170+ } else if (dataType == ge::DT_BF16) {
171+ oneBlockNum = ONE_BLOCK_SIZE / BF16_SIZE;
172+ } else {
173+ return ge::GRAPH_FAILED;
174+ }
175+ tilingContext->SetTilingKey(0);
176+ return ge::GRAPH_SUCCESS;
177+}
178+ 
179+static ge::graphStatus TilingPrepare4SiluMulTiling([[maybe_unused]] gert::TilingParseContext* context)
180+{
181+ return ge::GRAPH_SUCCESS;
182+}
183+ 
184+static ge::graphStatus TilingSiluMulTiling(gert::TilingContext* context)
185+{
186+ SiluMulTiling tilingObject(context);
187+ return tilingObject.RunBigKernelTiling();
188+}
189+ 
190+IMPL_OP_OPTILING(SiluMul).Tiling(TilingSiluMulTiling).TilingParse<SiluMulCompileInfo>(TilingPrepare4SiluMulTiling);
191+} // namespace optiling
Aexperimental/activation/silu_mul/op_host/silu_mul_tiling.h+43-0
@@ -0,0 +1,43 @@
1+/*
2+ * Copyright (c) 2025 联通（广东）产业互联网有限公司.
3+ *
4+ * Licensed under the Apache License, Version 2.0 (the "License");
5+ * you may not use this file except in compliance with the License.
6+ * You may obtain a copy of the License at
7+ *
8+ * http://www.apache.org/licenses/LICENSE-2.0
9+ *
10+ * Unless required by applicable law or agreed to in writing, software
11+ * distributed under the License is distributed on an "AS IS" BASIS,
12+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13+ * See the License for the specific language governing permissions and
14+ * limitations under the License.
15+ */
16+ 
17+/*!
18+ * \file silu_mul_tiling_def.h
19+ * \brief silu_mul_tiling_def
20+ */
21+ 
22+#ifndef SILU_MUL_TILING_DEF_H
23+#define SILU_MUL_TILING_DEF_H
24+ 
25+#include "register/tilingdata_base.h"
26+#include "tiling/tiling_api.h"
27+ 
28+namespace optiling {
29+struct SiluMulCompileInfo {};
30+ 
31+BEGIN_TILING_DATA_DEF(SiluMulTilingData)
32+ 
33+TILING_DATA_FIELD_DEF(int32_t, lastDimSize);
34+TILING_DATA_FIELD_DEF(int32_t, batchSize);
35+TILING_DATA_FIELD_DEF(int32_t, PPMaxCalNum);
36+TILING_DATA_FIELD_DEF(uint32_t, needCoreNum);
37+ 
38+END_TILING_DATA_DEF;
39+ 
40+REGISTER_TILING_DATA_CLASS(SiluMul, SiluMulTilingData)
41+} // namespace optiling
42+ 
43+#endif
Aexperimental/activation/silu_mul/op_kernel/silu_mul.cpp+40-0
@@ -0,0 +1,40 @@
1+/*
2+ * Copyright (c) 2025 联通（广东）产业互联网有限公司.
3+ *
4+ * Licensed under the Apache License, Version 2.0 (the "License");
5+ * you may not use this file except in compliance with the License.
6+ * You may obtain a copy of the License at
7+ *
8+ * http://www.apache.org/licenses/LICENSE-2.0
9+ *
10+ * Unless required by applicable law or agreed to in writing, software
11+ * distributed under the License is distributed on an "AS IS" BASIS,
12+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13+ * See the License for the specific language governing permissions and
14+ * limitations under the License.
15+ */
16+ 
17+/*!
18+ * \file silu_mul.cpp
19+ * \brief silu_mul kernel
20+ */
21+ 
22+#include "silu_mul.h"
23+ 
24+using namespace AscendC;
25+ 
26+using namespace SiluMul;
27+ 
28+extern "C" __global__ __aicore__ void silu_mul(GM_ADDR input, GM_ADDR output, GM_ADDR workspace, GM_ADDR tiling)
29+{
30+ GET_TILING_DATA(tilingData, tiling);
31+ 
32+ if (TILING_KEY_IS(0)) {
33+ GM_ADDR userWs = nullptr;
34+ 
35+ SiluMulND<DTYPE_X> op;
36+ 
37+ op.Init(input, output, userWs, &tilingData);
38+ op.Process();
39+ }
40+}
Aexperimental/activation/silu_mul/op_kernel/silu_mul.h+341-0
@@ -0,0 +1,341 @@
1+/*
2+ * Copyright (c) 2025 联通（广东）产业互联网有限公司.
3+ *
4+ * Licensed under the Apache License, Version 2.0 (the "License");
5+ * you may not use this file except in compliance with the License.
6+ * You may obtain a copy of the License at
7+ *
8+ * http://www.apache.org/licenses/LICENSE-2.0
9+ *
10+ * Unless required by applicable law or agreed to in writing, software
11+ * distributed under the License is distributed on an "AS IS" BASIS,
12+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13+ * See the License for the specific language governing permissions and
14+ * limitations under the License.
15+ */
16+ 
17+/*!
18+ * \file silu_mul.h
19+ * \brief silu_mul head file
20+ */
21+ 
22+#ifndef SILU_MUL_H
23+#define SILU_MUL_H
24+ 
25+#include <type_traits>
26+#include "kernel_operator.h"
27+ 
28+namespace SiluMul {
29+ 
30+using namespace AscendC;
31+ 
32+constexpr int32_t MAX_UB_SIZE = 184 * 1024;
33+constexpr int32_t ONE_BLOCK_SIZE = 32;
34+ 
35+template <typename T>
36+class SiluMulND {
37+public:
38+ TPipe pipe;
39+ __aicore__ inline SiluMulND(){};
40+ __aicore__ inline void Init(GM_ADDR input, GM_ADDR output, GM_ADDR workspace, const SiluMulTilingData* tilingData);
41+ __aicore__ inline void Process();
42+ 
43+private:
44+ __aicore__ inline void BigTailProcess();
45+ __aicore__ inline void SmallTailProcess();
46+ __aicore__ inline void CopyIn(int64_t inputOffset, DataCopyExtParams dataCopyParams);
47+ __aicore__ inline void Compute(int64_t dataCount);
48+ __aicore__ inline void CopyOut(int64_t outputOffset, int64_t dataCount, DataCopyExtParams dataCopyParams);
49+ __aicore__ inline void SmallTailCopyIn(
50+ int32_t gmOffset, uint32_t tmpCalNum, uint32_t dAlign, LocalTensor<T>& x1Pad, LocalTensor<T>& x2Pad);
51+ __aicore__ inline void SmallTailCompute(
52+ uint32_t calcLen, uint32_t maxN, uint32_t dAlign, LocalTensor<T>& x1Pad, LocalTensor<T>& x2Pad,
53+ LocalTensor<float>& floatBuf);
54+ __aicore__ inline void SmallTailCopyOut(
55+ int32_t gmOffset, uint32_t tmpCalNum, uint32_t dAlign, LocalTensor<T>& x1Pad);
56+ 
57+private:
58+ TBuf<QuePosition::VECCALC> ubTBuf;
59+ LocalTensor<uint8_t> tmpTensor;
60+ 
61+ LocalTensor<T> x1Tmp;
62+ LocalTensor<T> x2Tmp;
63+ 
64+ LocalTensor<T> x1Tensor;
65+ LocalTensor<T> x2Tensor;
66+ 
67+ LocalTensor<float> tempResTensor;
68+ 
69+ LocalTensor<float> x1TensorFp32;
70+ LocalTensor<float> x2TensorFp32;
71+ 
72+ GlobalTensor<T> inputGm;
73+ GlobalTensor<T> outputGm;
74+ 
75+ int64_t lastDimSize;
76+ int64_t batchSize;
77+ uint32_t d;
78+ 
79+ int64_t PPMaxCalNum;
80+ 
81+ uint32_t needCoreNumber;
82+ int32_t blockIdx;
83+ 
84+ event_t eventId = EVENT_ID0;
85+ int32_t pingPongFlag = 0;
86+};
87+ 
88+template <typename T>
89+__aicore__ inline void SiluMulND<T>::Init(
90+ GM_ADDR input, GM_ADDR output, GM_ADDR workspace, const SiluMulTilingData* tilingData)
91+{
92+ inputGm.SetGlobalBuffer((__gm__ T*)input);
93+ outputGm.SetGlobalBuffer((__gm__ T*)output);
94+ 
95+ batchSize = tilingData->batchSize;
96+ lastDimSize = tilingData->lastDimSize;
97+ needCoreNumber = tilingData->needCoreNum;
98+ PPMaxCalNum = tilingData->PPMaxCalNum;
99+ 
100+ d = lastDimSize / 2;
101+ 
102+ blockIdx = GetBlockIdx();
103+ pipe.InitBuffer(ubTBuf, MAX_UB_SIZE);
104+ tmpTensor = ubTBuf.Get<uint8_t>();
105+}
106+ 
107+template <typename T>
108+__aicore__ inline void SiluMulND<T>::Process()
109+{
110+ if (blockIdx >= needCoreNumber) {
111+ return;
112+ }
113+ if (d > PPMaxCalNum) {
114+ BigTailProcess();
115+ } else {
116+ SmallTailProcess();
117+ }
118+}
119+ 
120+template <typename T>
121+__aicore__ inline void SiluMulND<T>::BigTailProcess()
122+{
123+ int32_t loopNum = batchSize / needCoreNumber;
124+ int32_t loopRemain = batchSize % needCoreNumber;
125+ if (loopRemain > 0 && blockIdx < loopRemain) {
126+ loopNum++;
127+ }
128+ for (int32_t i = 0; i < loopNum; i++) {
129+ int32_t totalOffset = i * needCoreNumber * lastDimSize + blockIdx * lastDimSize;
130+ int32_t outOffset = i * needCoreNumber * d + blockIdx * d;
131+ int32_t eachLineLoop = d / PPMaxCalNum;
132+ uint32_t remain = d % PPMaxCalNum;
133+ if (remain > 0) {
134+ eachLineLoop++;
135+ }
136+ pingPongFlag = 0;
137+ SetFlag<HardEvent::MTE3_MTE2>(EVENT_ID0);
138+ SetFlag<HardEvent::MTE3_MTE2>(EVENT_ID1);
139+ for (int32_t j = 0; j < eachLineLoop; j++) {
140+ uint32_t dataCount = PPMaxCalNum;
141+ if (j == eachLineLoop - 1 && remain > 0) {
142+ dataCount = remain;
143+ }
144+ int32_t localOffset = j * PPMaxCalNum;
145+ eventId = pingPongFlag ? EVENT_ID1 : EVENT_ID0;
146+ DataCopyExtParams dataCopyParams{1, static_cast<uint32_t>(dataCount * sizeof(T)), 0, 0, 0};
147+ CopyIn(totalOffset + localOffset, dataCopyParams);
148+ Compute(dataCount);
149+ CopyOut(outOffset + localOffset, dataCount, dataCopyParams);
150+ pingPongFlag = 1 - pingPongFlag;
151+ }
152+ WaitFlag<HardEvent::MTE3_MTE2>(EVENT_ID0);
153+ WaitFlag<HardEvent::MTE3_MTE2>(EVENT_ID1);
154+ }
155+}
156+ 
157+template <typename T>
158+__aicore__ inline void SiluMulND<T>::CopyIn(int64_t inputOffset, DataCopyExtParams dataCopyParams)
159+{
160+ x1Tensor = pingPongFlag ? tmpTensor[MAX_UB_SIZE / 2].ReinterpretCast<T>() : tmpTensor[0].ReinterpretCast<T>();
161+ x2Tensor = pingPongFlag ? tmpTensor[PPMaxCalNum * sizeof(float) + MAX_UB_SIZE / 2].ReinterpretCast<T>() :
162+ tmpTensor[PPMaxCalNum * sizeof(float)].ReinterpretCast<T>();
163+ WaitFlag<HardEvent::MTE3_MTE2>(eventId);
164+ 
165+ DataCopyPadExtParams<T> padParams{false, 0, 0, 0};
166+ if (std::is_same_v<T, bfloat16_t> || std::is_same_v<T, half>) {
167+ int32_t elementByte = PPMaxCalNum * sizeof(T);
168+ x1Tmp = pingPongFlag ? tmpTensor[elementByte + MAX_UB_SIZE / 2].ReinterpretCast<T>() :
169+ tmpTensor[elementByte].ReinterpretCast<T>();
170+ x2Tmp = pingPongFlag ?
171+ tmpTensor[elementByte + PPMaxCalNum * sizeof(float) + MAX_UB_SIZE / 2].ReinterpretCast<T>() :
172+ tmpTensor[elementByte + PPMaxCalNum * sizeof(float)].ReinterpretCast<T>();
173+ DataCopyPad(x1Tmp, inputGm[inputOffset], dataCopyParams, padParams);
174+ DataCopyPad(x2Tmp, inputGm[inputOffset + d], dataCopyParams, padParams);
175+ } else {
176+ DataCopyPad(x1Tensor, inputGm[inputOffset], dataCopyParams, padParams);
177+ DataCopyPad(x2Tensor, inputGm[inputOffset + d], dataCopyParams, padParams);
178+ }
179+ 
180+ SetFlag<HardEvent::MTE2_V>(eventId);
181+ WaitFlag<HardEvent::MTE2_V>(eventId);
182+}
183+ 
184+template <typename T>
185+__aicore__ inline void SiluMulND<T>::Compute(int64_t dataCount)
186+{
187+ x1TensorFp32 = x1Tensor.template ReinterpretCast<float>();
188+ x2TensorFp32 = x2Tensor.template ReinterpretCast<float>();
189+ if (std::is_same_v<T, bfloat16_t> || std::is_same_v<T, half>) {
190+ Cast(x1TensorFp32, x1Tmp, RoundMode::CAST_NONE, dataCount);
191+ PipeBarrier<PIPE_V>();
192+ Cast(x2TensorFp32, x2Tmp, RoundMode::CAST_NONE, dataCount);
193+ PipeBarrier<PIPE_V>();
194+ }
195+ tempResTensor = pingPongFlag ?
196+ tmpTensor[PPMaxCalNum * 2 * sizeof(float) + MAX_UB_SIZE / 2].ReinterpretCast<float>() :
197+ tmpTensor[PPMaxCalNum * 2 * sizeof(float)].ReinterpretCast<float>();
198+ 
199+ // Silu Calculation: x1 * Sigmoid(x1)
200+ Sigmoid(tempResTensor, x1TensorFp32, dataCount);
201+ 
202+ PipeBarrier<PIPE_V>();
203+ 
204+ Mul(x1TensorFp32, x1TensorFp32, tempResTensor, dataCount);
范其瑞
范其瑞范其瑞2025年12月30日

缺少两个PIPE_V同步，可能有精度问题

likedislike
W
wuxs68
2025年12月30日 评论：
范其瑞
范其瑞
2025年12月31日 评论：
W
wuxs68
2025年12月31日 评论：
205+ 
206+ PipeBarrier<PIPE_V>();
207+ 
208+ // Silu * x2
209+ Mul(x1TensorFp32, x1TensorFp32, x2TensorFp32, dataCount);
210+ PipeBarrier<PIPE_V>();
211+}
212+ 
213+template <typename T>
214+__aicore__ inline void SiluMulND<T>::CopyOut(int64_t outputOffset, int64_t dataCount, DataCopyExtParams dataCopyParams)
215+{
216+ if (std::is_same_v<T, half>) {
217+ Cast(x1Tensor, x1TensorFp32, RoundMode::CAST_NONE, dataCount);
218+ PipeBarrier<PIPE_V>();
219+ } else if (std::is_same_v<T, bfloat16_t>) {
220+ Cast(x1Tensor, x1TensorFp32, RoundMode::CAST_RINT, dataCount);
221+ PipeBarrier<PIPE_V>();
222+ }
223+ SetFlag<HardEvent::V_MTE3>(eventId);
224+ WaitFlag<HardEvent::V_MTE3>(eventId);
225+ DataCopyPad(outputGm[outputOffset], x1Tensor, dataCopyParams);
226+ SetFlag<HardEvent::MTE3_MTE2>(eventId);
227+}
228+template <typename T>
229+__aicore__ inline void SiluMulND<T>::SmallTailCopyIn(
230+ int32_t gmOffset, uint32_t tmpCalNum, uint32_t dAlign, LocalTensor<T>& x1Pad, LocalTensor<T>& x2Pad)
231+{
232+ for (int32_t j = 0; j < tmpCalNum; j++) {
233+ int32_t batchOffset = gmOffset + j * lastDimSize;
234+ for (int32_t k = 0; k < d; k++) {
235+ x1Pad.SetValue(j * dAlign + k, inputGm.GetValue(batchOffset + k));
236+ x2Pad.SetValue(j * dAlign + k, inputGm.GetValue(batchOffset + d + k));
237+ }
238+ }
239+}
240+ 
241+template <typename T>
242+__aicore__ inline void SiluMulND<T>::SmallTailCompute(
243+ uint32_t calcLen, uint32_t maxN, uint32_t dAlign, LocalTensor<T>& x1Pad, LocalTensor<T>& x2Pad,
244+ LocalTensor<float>& floatBuf)
245+{
246+ LocalTensor<float> compRes = floatBuf;
247+ LocalTensor<float> compX1 = floatBuf[maxN * dAlign];
248+ LocalTensor<float> compX2 = floatBuf[2 * maxN * dAlign];
249+ 
250+ if constexpr (std::is_same_v<T, float>) {
251+ Sigmoid(compRes, x1Pad, calcLen);
252+ PipeBarrier<PIPE_V>();
253+ Mul(x1Pad, x1Pad, compRes, calcLen);
254+ PipeBarrier<PIPE_V>();
255+ Mul(x1Pad, x1Pad, x2Pad, calcLen);
256+ } else {
257+ Cast(compX1, x1Pad, RoundMode::CAST_NONE, calcLen);
258+ Cast(compX2, x2Pad, RoundMode::CAST_NONE, calcLen);
259+ PipeBarrier<PIPE_V>();
260+ 
261+ Sigmoid(compRes, compX1, calcLen);
262+ PipeBarrier<PIPE_V>();
263+ Mul(compX1, compX1, compRes, calcLen);
264+ PipeBarrier<PIPE_V>();
265+ Mul(compX1, compX1, compX2, calcLen);
266+ 
267+ if constexpr (std::is_same_v<T, half>) {
268+ Cast(x1Pad, compX1, RoundMode::CAST_NONE, calcLen);
269+ } else {
270+ Cast(x1Pad, compX1, RoundMode::CAST_RINT, calcLen);
271+ }
272+ }
273+}
274+ 
275+template <typename T>
276+__aicore__ inline void SiluMulND<T>::SmallTailCopyOut(
277+ int32_t gmOffset, uint32_t tmpCalNum, uint32_t dAlign, LocalTensor<T>& x1Pad)
278+{
279+ int32_t outOffset = gmOffset / 2; // 注意：这里假设输入是 x,y 拼接，输出是 z，且 lastDimSize = 2*d
280+ for (int32_t j = 0; j < tmpCalNum; j++) {
281+ for (int32_t k = 0; k < d; k++) {
282+ T val = x1Pad.GetValue(j * dAlign + k);
283+ outputGm.SetValue(outOffset + j * d + k, val);
284+ }
285+ }
286+}
287+ 
288+template <typename T>
289+__aicore__ inline void SiluMulND<T>::SmallTailProcess()
290+{
291+ uint32_t sizeOfT = sizeof(T);
292+ uint32_t oneBlockNum = 32 / sizeOfT;
293+ uint32_t dAlign = (d + oneBlockNum - 1) / oneBlockNum * oneBlockNum;
294+ 
295+ uint32_t bytesPerRow = 2 * dAlign * sizeOfT + 3 * dAlign * sizeof(float);
296+ uint32_t n = (MAX_UB_SIZE) / bytesPerRow;
297+ if (n == 0)
298+ n = 1;
299+ if (n > PPMaxCalNum / dAlign)
300+ n = PPMaxCalNum / dAlign;
301+ 
302+ int32_t eachCoreNum = batchSize / needCoreNumber;
303+ int32_t remain = batchSize % needCoreNumber;
304+ if (remain > 0 && blockIdx < remain)
305+ eachCoreNum++;
306+ 
307+ int32_t loopNum = eachCoreNum / n;
308+ int32_t loopRemain = eachCoreNum % n;
309+ if (loopRemain > 0)
310+ loopNum++;
311+ 
312+ int32_t totalOffset = eachCoreNum * blockIdx * lastDimSize;
313+ if (remain > 0) {
314+ totalOffset = (blockIdx < remain) ? (eachCoreNum * blockIdx * lastDimSize) :
315+ ((blockIdx * eachCoreNum + remain) * lastDimSize);
316+ }
317+ 
318+ LocalTensor<T> x1Pad = tmpTensor[0].ReinterpretCast<T>();
319+ LocalTensor<T> x2Pad = tmpTensor[n * dAlign * sizeOfT].ReinterpretCast<T>();
320+ LocalTensor<float> floatBuf = tmpTensor[2 * n * dAlign * sizeOfT].ReinterpretCast<float>();
321+ 
322+ for (int32_t i = 0; i < loopNum; i++) {
323+ uint32_t tmpCalNum = (loopRemain > 0 && i == loopNum - 1) ? loopRemain : n;
324+ int32_t gmOffset = totalOffset + i * n * lastDimSize;
325+ 
326+ SmallTailCopyIn(gmOffset, tmpCalNum, dAlign, x1Pad, x2Pad);
327+ 
328+ PipeBarrier<PIPE_ALL>();
329+ 
330+ SmallTailCompute(tmpCalNum * dAlign, n, dAlign, x1Pad, x2Pad, floatBuf);
331+ 
332+ PipeBarrier<PIPE_ALL>();
333+ 
334+ SmallTailCopyOut(gmOffset, tmpCalNum, dAlign, x1Pad);
335+ 
336+ PipeBarrier<PIPE_ALL>();
337+ }
338+}
339+ 
340+} // namespace SiluMul
341+#endif
Aexperimental/activation/silu_mul/tests/CMakeLists.txt+25-0
@@ -0,0 +1,25 @@
1+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
2+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
3+# CANN Open Software License Agreement Version 2.0 (the "License").
4+# Please refer to the License for details. You may not use this file except in compliance with the License.
5+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
6+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
7+# See LICENSE in the root of the software repository for the full text of the License.
8+#/
9+ 
10+message(STATUS "=== Debug: start ops.activation.silu_mul.tests.CMakeLists.txt ")
11+if(NOT DEFINED TEST_DTYPE)
12+ set(TEST_DTYPE bfloat16_t) # float / half / bfloat16_t
13+endif()
14+ 
15+add_compile_definitions(DTYPE_X=${TEST_DTYPE})
16+ 
17+message(STATUS ">>> UT is compiling for DTYPE_X = ${TEST_DTYPE}")
18+ 
19+file(GLOB CURRENT_SOURCE_DIRS LIST_DIRECTORIES true ${CMAKE_CURRENT_SOURCE_DIR}/*)
20+message(STATUS "=== Debug: CURRENT_SOURCE_DIRS =${CURRENT_SOURCE_DIRS} ")
21+foreach(SUB_DIR ${CURRENT_SOURCE_DIRS})
22+ if(EXISTS "${SUB_DIR}/CMakeLists.txt")
23+ add_subdirectory(${SUB_DIR})
24+ endif()
25+endforeach()
Aexperimental/activation/silu_mul/tests/ut/CMakeLists.txt+18-0
@@ -0,0 +1,18 @@
1+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
2+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
3+# CANN Open Software License Agreement Version 2.0 (the "License").
4+# Please refer to the License for details. You may not use this file except in compliance with the License.
5+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
6+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
7+# See LICENSE in the root of the software repository for the full text of the License.
8+#/
9+ 
10+# 每个目录下需要生成的可执行文件，具体参考：ops/built-in/test/CMakeLists.txt： 50~124
11+message(STATUS "=== Debug: start ops.activation.silu_mul.tests.ut.CMakeLists.txt ")
12+file(GLOB CURRENT_SOURCE_DIRS LIST_DIRECTORIES true ${CMAKE_CURRENT_SOURCE_DIR}/*)
13+message(STATUS "=== Debug: CURRENT_SOURCE_DIRS =${CURRENT_SOURCE_DIRS} ")
14+foreach(SUB_DIR ${CURRENT_SOURCE_DIRS})
15+ if(EXISTS "${SUB_DIR}/CMakeLists.txt")
16+ add_subdirectory(${SUB_DIR})
17+ endif()
18+endforeach()
Aexperimental/activation/silu_mul/tests/ut/op_host/CMakeLists.txt+14-0
@@ -0,0 +1,14 @@
1+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
2+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
3+# CANN Open Software License Agreement Version 2.0 (the "License").
4+# Please refer to the License for details. You may not use this file except in compliance with the License.
5+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
6+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
7+# See LICENSE in the root of the software repository for the full text of the License.
8+#/
9+ 
10+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
11+if(UT_TEST_ALL OR OP_HOST_UT)
12+ add_modules_ut_sources(HOSTNAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
13+ add_modules_ut_sources(HOSTNAME ${OP_INFERSHAPE_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
14+endif()
Aexperimental/activation/silu_mul/tests/ut/op_host/test_silu_mul_infershape.cpp+44-0
@@ -0,0 +1,44 @@
1+/*
2+ * Copyright (c) 2025 联通（广东）产业互联网有限公司.
3+ *
4+ * Licensed under the Apache License, Version 2.0 (the "License");
5+ * you may not use this file except in compliance with the License.
6+ * You may obtain a copy of the License at
7+ *
8+ * http://www.apache.org/licenses/LICENSE-2.0
9+ *
10+ * Unless required by applicable law or agreed to in writing, software
11+ * distributed under the License is distributed on an "AS IS" BASIS,
12+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13+ * See the License for the specific language governing permissions and
14+ * limitations under the License.
15+ */
16+ 
17+#include <gtest/gtest.h> // NOLINT
18+#include <iostream>
19+#include "infershape_test_util.h"
20+#include "ut_op_common.h"
21+#include "../../../op_graph/silu_mul_proto.h"
22+#include "log/log.h"
23+class SiluMul : public testing::Test {
24+protected:
25+ static void SetUpTestCase()
26+ {
27+ std::cout << "SiluMul SetUp" << std::endl;
28+ }
29+ 
30+ static void TearDownTestCase()
31+ {
32+ std::cout << "SiluMul TearDown" << std::endl;
33+ }
34+};
35+ 
36+TEST_F(SiluMul, SiluMul_infershape_case_0)
37+{
38+ ge::op::SiluMul op;
39+ op.UpdateInputDesc("x", create_desc({4, 1, 1280}, ge::DT_FLOAT16));
40+ op.UpdateInputDesc("y", create_desc({4, 1, 1280}, ge::DT_FLOAT16));
41+ 
42+ EXPECT_EQ(InferShapeTest(op), ge::GRAPH_SUCCESS);
43+ EXPECT_EQ(InferDataTypeTest(op), ge::GRAPH_SUCCESS);
44+}
Aexperimental/activation/silu_mul/tests/ut/op_host/test_silu_mul_tiling.cpp+141-0
@@ -0,0 +1,141 @@
1+/*
2+ * Copyright (c) 2025 联通（广东）产业互联网有限公司.
3+ *
4+ * Licensed under the Apache License, Version 2.0 (the "License");
5+ * you may not use this file except in compliance with the License.
6+ * You may obtain a copy of the License at
7+ *
8+ * http://www.apache.org/licenses/LICENSE-2.0
9+ *
10+ * Unless required by applicable law or agreed to in writing, software
11+ * distributed under the License is distributed on an "AS IS" BASIS,
12+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13+ * See the License for the specific language governing permissions and
14+ * limitations under the License.
15+ */
16+ 
17+#include <iostream>
18+#include <fstream>
19+#include <vector>
20+#include <gtest/gtest.h>
21+#include "log/log.h"
22+#include "../../../op_host/silu_mul_tiling.h"
23+#include "kernel_run_context_facker.h"
24+#include "test_cube_util.h"
25+#include "exe_graph/runtime/storage_format.h"
26+#include "exe_graph/runtime/storage_shape.h"
27+#include "platform/platform_infos_def.h"
28+#include "register/op_impl_registry.h"
29+#include "ut_op_util.h"
30+ 
31+using namespace ut_util;
32+using namespace std;
33+using namespace ge;
34+ 
35+class SiluMulTiling : public testing::Test {
36+protected:
37+ static void SetUpTestCase()
38+ {
39+ std::cout << "SiluMulTiling SetUp" << std::endl;
40+ }
41+ 
42+ static void TearDownTestCase()
43+ {
44+ std::cout << "SiluMulTiling TearDown" << std::endl;
45+ }
46+};
47+ 
48+static string TilingData2Str(const gert::TilingData* tiling_data)
49+{
50+ auto data = tiling_data->GetData();
51+ string result;
52+ for (size_t i = 0; i < tiling_data->GetDataSize(); i += sizeof(int64_t)) {
53+ result += std::to_string((reinterpret_cast<const int64_t*>(tiling_data->GetData())[i / sizeof(int64_t)]));
54+ result += " ";
55+ }
56+ 
57+ return result;
58+}
59+ 
60+TEST_F(SiluMulTiling, SiluMulTiling_01)
61+{
62+ size_t M = 2;
63+ size_t N = 4;
64+ 
65+ gert::StorageShape x_shape = {{M, N}, {M, N}};
66+ gert::StorageShape y_shape = {{M, N}, {M, N}};
67+ gert::StorageShape z_shape = {{M, N}, {M, N}};
68+ 
69+ string compile_info_string = R"({
70+ "hardware_info": {"BT_SIZE": 0, "load3d_constraints": "1",
71+ "Intrinsic_fix_pipe_l0c2out": false, "Intrinsic_data_move_l12ub": true, "Intrinsic_data_move_l0c2ub": true, "Intrinsic_data_move_out2l1_nd2nz": false,
72+ "UB_SIZE": 196608, "L2_SIZE": 33554432, "L1_SIZE": 524288,
73+ "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072,
74+ "CORE_NUM": 40}
75+ })";
76+ map<string, string> soc_infos;
77+ map<string, string> aicore_spec;
78+ map<string, string> intrinsics;
79+ GetPlatFormInfos(compile_info_string.c_str(), soc_infos, aicore_spec, intrinsics);
80+ 
81+ // platform info
82+ fe::PlatFormInfos platform_info;
83+ platform_info.Init();
84+ // compile info
85+ optiling::SiluMulCompileInfo compile_info;
86+ 
87+ std::string op_type("SiluMul");
88+ ASSERT_NE(gert::OpImplRegistry::GetInstance().GetOpImpl(op_type.c_str()), nullptr);
89+ auto tiling_func = gert::OpImplRegistry::GetInstance().GetOpImpl(op_type.c_str())->tiling;
90+ auto tiling_parse_func = gert::OpImplRegistry::GetInstance().GetOpImpl(op_type.c_str())->tiling_parse;
91+ 
92+ // tilingParseFunc simulate
93+ auto kernel_holder =
94+ gert::KernelRunContextFaker()
95+ .KernelIONum(2, 1)
96+ .Inputs({const_cast<char*>(compile_info_string.c_str()), reinterpret_cast<void*>(&platform_info)})
97+ .Outputs({&compile_info})
98+ .Build();
99+ 
100+ ASSERT_TRUE(kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->Init());
101+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", soc_infos);
102+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicore_spec);
103+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
104+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes(
105+ "AICoreintrinsicDtypeMap", intrinsics);
106+ 
107+ ASSERT_EQ(tiling_parse_func(kernel_holder.GetContext<gert::KernelContext>()), ge::GRAPH_SUCCESS);
108+ 
109+ // tilingFunc simulate
110+ auto param = gert::TilingData::CreateCap(4096);
111+ auto workspace_size_holer = gert::ContinuousVector::Create<size_t>(4096);
112+ auto ws_size = reinterpret_cast<gert::ContinuousVector*>(workspace_size_holer.get());
113+ ASSERT_NE(param, nullptr);
114+ 
115+ auto holder = gert::TilingContextFaker()
116+ .NodeIoNum(2, 1)
117+ .IrInstanceNum({1})
118+ .InputShapes({&x_shape, &y_shape})
119+ .OutputShapes({&z_shape})
120+ .CompileInfo(&compile_info)
121+ .PlatformInfo(reinterpret_cast<char*>(&platform_info))
122+ .NodeInputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
123+ .NodeInputTd(1, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
124+ .NodeOutputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
125+ .TilingData(param.get())
126+ .Workspace(ws_size)
127+ .Build();
128+ 
129+ gert::TilingContext* tiling_context = holder.GetContext<gert::TilingContext>();
130+ ASSERT_NE(tiling_context->GetPlatformInfo(), nullptr);
131+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", soc_infos);
132+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicore_spec);
133+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
134+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics);
135+ 
136+ // workspaces nullptr return failed
137+ EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_SUCCESS);
138+ 
139+ auto tiling_key = tiling_context->GetTilingKey();
140+ ASSERT_EQ(tiling_key, 0);
141+}
Aexperimental/activation/silu_mul/tests/ut/op_kernel/CMakeLists.txt+12-0
@@ -0,0 +1,12 @@
1+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
2+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
3+# CANN Open Software License Agreement Version 2.0 (the "License").
4+# Please refer to the License for details. You may not use this file except in compliance with the License.
5+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
6+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
7+# See LICENSE in the root of the software repository for the full text of the License.
8+#/
9+ 
10+if ((UT_TEST_ALL OR OP_KERNEL_UT) AND NOT UT_DONE)
11+ AddOpTestCase(silu_mul "ascend910B1" "" )
12+endif()
Aexperimental/activation/silu_mul/tests/ut/op_kernel/silu_mul_data/compare_data.py+92-0
@@ -0,0 +1,92 @@
1+#!/usr/bin/env python3
2+# -*- coding: utf-8 -*-
3+# ----------------------------------------------------------------------------
4+# Copyright (c) 2025 联通（广东）产业互联网有限公司.
5+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
6+# CANN Open Software License Agreement Version 2.0 (the "License").
7+# Please refer to the License for details. You may not use this file except in compliance with the License.
8+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
9+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
10+# See LICENSE in the root of the software repository for the full text of the License.
11+# ----------------------------------------------------------------------------
12+ 
13+import sys
14+import numpy as np
15+import glob
16+import os
17+ 
18+curr_dir = os.path.dirname(os.path.realpath(__file__))
19+ 
20+def compare_data(golden_file_lists, output_file_lists, d_type):
21+ if d_type == "float16":
22+ np_dtype = np.float16
23+ rtol = 1.e-3
24+ atol = 1.e-3
25+ elif d_type == "float32":
26+ np_dtype = np.float32
27+ rtol = 1.e-4
28+ atol = 1.e-4
29+ elif d_type == "bfloat16":
30+ # numpy does not support bfloat16, using float32 as placeholder if converted beforehand
31+ np_dtype = np.float32
32+ rtol = 1.e-1
33+ atol = 1.e-1
34+ else:
35+ np_dtype = np.float32
36+ rtol = 1.e-4
37+ atol = 1.e-4
38+
39+ data_same = True
40+ for gold, out in zip(golden_file_lists, output_file_lists):
41+ try:
42+ tmp_out = np.fromfile(out, np_dtype)
43+ tmp_gold = np.fromfile(gold, np_dtype)
44+ except Exception as e:
45+ print(f"Read file failed: {e}")
46+ continue
47+ 
48+ if tmp_out.shape != tmp_gold.shape:
49+ print(f"Shape mismatch in {os.path.basename(out)}: output {tmp_out.shape} vs golden {tmp_gold.shape}")
50+ data_same = False
51+ continue
52+ 
53+ diff_res = np.isclose(tmp_out, tmp_gold, rtol=rtol, atol=atol, equal_nan=True)
54+
55+ print(f"=== Preview first 5 elements for {os.path.basename(out)} ===")
56+ for idx in range(min(5, tmp_out.size)):
57+ print(f"index: {idx}, output: {tmp_out[idx]:.30f}, golden: {tmp_gold[idx]:.30f}")
58+ 
59+ if not np.all(diff_res):
60+ data_same = False
61+ diff_idx = np.where(diff_res == False)[0]
62+ print(f"Comparison failed for {os.path.basename(out)}")
63+ for idx in diff_idx[:5]:
64+ print(f"index: {idx}, output: {tmp_out[idx]}, golden: {tmp_gold[idx]}")
65+ else:
66+ print(f"Comparison passed for {os.path.basename(out)}")
67+ 
68+ if data_same:
69+ print("PASSED!")
70+ else:
71+ print("FAILED!")
72+ return data_same
73+ 
74+def get_file_lists():
75+ golden_file_lists = sorted(glob.glob(os.path.join(curr_dir, "*golden*.bin")))
76+ output_file_lists = sorted(glob.glob(os.path.join(curr_dir, "*output*.bin")))
77+ return golden_file_lists, output_file_lists
78+ 
79+def process(d_type):
80+ golden_file_lists, output_file_lists = get_file_lists()
81+ if len(golden_file_lists) != len(output_file_lists) or len(golden_file_lists) == 0:
82+ print("Error: File count mismatch or no files found.")
83+ print(f"Golden files: {len(golden_file_lists)}")
84+ print(f"Output files: {len(output_file_lists)}")
85+ return
86+ compare_data(golden_file_lists, output_file_lists, d_type)
87+ 
88+if __name__ == '__main__':
89+ if len(sys.argv) < 2:
90+ print("Usage: python3 verify_result.py <dtype>")
91+ else:
92+ process(sys.argv[1])
Aexperimental/activation/silu_mul/tests/ut/op_kernel/silu_mul_data/gen_data.py+56-0
@@ -0,0 +1,56 @@
1+#!/usr/bin/env python3
2+# -*- coding: utf-8 -*-
3+# ----------------------------------------------------------------------------
4+# Copyright (c) 2025 联通（广东）产业互联网有限公司.
5+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
6+# CANN Open Software License Agreement Version 2.0 (the "License").
7+# Please refer to the License for details. You may not use this file except in compliance with the License.
8+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
9+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
10+# See LICENSE in the root of the software repository for the full text of the License.
11+# ----------------------------------------------------------------------------
12+ 
13+import sys
14+import os
15+import numpy as np
16+import re
17+import torch
18+import tensorflow as tf
19+ 
20+ 
21+def parse_str_to_shape_list(shape_str):
22+ shape_str = shape_str.strip('(').strip(')')
23+ shape_list = [int(x) for x in shape_str.split(",")]
24+ return np.array(shape_list), shape_list
25+ 
26+def gen_data_and_golden(input_shape_str, output_size_str, d_type="float32"):
27+ d_type_dict = {
28+ "float32": np.float32,
29+ "float16": np.float16,
30+ "bfloat16_t": tf.bfloat16.as_numpy_dtype
31+ }
32+ np_type = d_type_dict[d_type]
33+ input_shape, _ = parse_str_to_shape_list(input_shape_str)
34+ _, output_size = parse_str_to_shape_list(output_size_str)
35+ 
36+ size = np.prod(input_shape)
37+ tmp_input = np.random.random(size).reshape(input_shape).astype(np_type)
38+ x_tensor = torch.tensor(tmp_input.astype(np.float32), dtype=torch.float32)
39+ d = input_shape[-1] // 2
40+ x1 = x_tensor[..., :d]
41+ x2 = x_tensor[..., d:]
42+ m = torch.nn.SiLU("none")
43+ x1 = m(x1)
44+ y_golden = x1 * x2
45+ tmp_golden = np.array(y_golden).astype(np_type)
46+ 
47+ tmp_input.astype(np_type).tofile(f"{d_type}_input_silu_mul.bin")
48+ tmp_golden.astype(np_type).tofile(f"{d_type}_golden_silu_mul.bin")
49+ 
50+if __name__ == "__main__":
51+ if len(sys.argv) != 4:
52+ print("Param num must be 4, actually is ", len(sys.argv))
53+ exit(1)
54+ # 清理bin文件
55+ os.system("rm -rf *.bin")
56+ gen_data_and_golden(sys.argv[1], sys.argv[2], sys.argv[3])
Aexperimental/activation/silu_mul/tests/ut/op_kernel/silu_mul_tiling_def.h+43-0
@@ -0,0 +1,43 @@
1+/*
2+ * Copyright (c) 2025 联通（广东）产业互联网有限公司.
3+ *
4+ * Licensed under the Apache License, Version 2.0 (the "License");
5+ * you may not use this file except in compliance with the License.
6+ * You may obtain a copy of the License at
7+ *
8+ * http://www.apache.org/licenses/LICENSE-2.0
9+ *
10+ * Unless required by applicable law or agreed to in writing, software
11+ * distributed under the License is distributed on an "AS IS" BASIS,
12+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13+ * See the License for the specific language governing permissions and
14+ * limitations under the License.
15+ */
16+#ifndef SILU_MUL_TILING_DEF_H
17+#define SILU_MUL_TILING_DEF_H
18+ 
19+#include "kernel_tiling/kernel_tiling.h"
20+ 
21+#include <cstdint>
22+#include <cstring>
23+ 
24+#define __CCE_UT_TEST__
25+ 
26+#define __aicore__
27+ 
28+struct SiluMulTilingData {
29+ int32_t lastDimSize = 4;
30+ int32_t batchSize = 2;
31+ int32_t PPMaxCalNum = 6144;
32+ uint32_t needCoreNum = 1;
33+};
34+ 
35+inline void ISiluMulTilingData(uint8_t* tiling, SiluMulTilingData* const_data)
36+{
37+ memcpy(const_data, tiling, sizeof(SiluMulTilingData));
38+}
39+ 
40+#define GET_TILING_DATA(tilingData, tilingPointer) \
41+ SiluMulTilingData tilingData; \
42+ ISiluMulTilingData(tilingPointer, &tilingData)
43+#endif
Aexperimental/activation/silu_mul/tests/ut/op_kernel/test_silu_mul.cpp+120-0
@@ -0,0 +1,120 @@
1+/*
2+ * Copyright (c) 2025 联通（广东）产业互联网有限公司.
3+ *
4+ * Licensed under the Apache License, Version 2.0 (the "License");
5+ * you may not use this file except in compliance with the License.
6+ * You may obtain a copy of the License at
7+ *
8+ * http://www.apache.org/licenses/LICENSE-2.0
9+ *
10+ * Unless required by applicable law or agreed to in writing, software
11+ * distributed under the License is distributed on an "AS IS" BASIS,
12+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13+ * See the License for the specific language governing permissions and
14+ * limitations under the License.
15+ */
16+#include <array>
17+#include <vector>
18+#include <iostream>
19+#include <string>
20+#include <cstdint>
21+#include <type_traits>
22+#include "gtest/gtest.h"
23+#include "tikicpulib.h"
24+#include "silu_mul_tiling_def.h"
25+#include "data_utils.h"
26+ 
27+using namespace std;
28+ 
29+#ifndef DTYPE_X
30+#define DTYPE_X float
31+#endif
32+ 
33+extern "C" __global__ __aicore__ void silu_mul(GM_ADDR input, GM_ADDR output, GM_ADDR workspace, GM_ADDR tiling);
34+ 
35+class silu_mul_test : public testing::Test {
36+protected:
37+ static void SetUpTestCase()
38+ {
39+ cout << "silu_mul SetUp\n" << endl;
40+ }
41+ static void TearDownTestCase()
42+ {
43+ cout << "silu_mul TearDown\n" << endl;
44+ }
45+};
46+ 
47+// 【新增】简单的类型映射辅助结构体
48+template <typename T>
49+struct DataTypeName {
50+ static constexpr const char* val = "unknown";
51+};
52+template <>
53+struct DataTypeName<float> {
54+ static constexpr const char* val = "float32";
55+};
56+template <>
57+struct DataTypeName<half> {
58+ static constexpr const char* val = "float16";
59+};
60+#if !(defined(__NPU_ARCH__) && __NPU_ARCH__ == 3003)
61+template <>
62+struct DataTypeName<bfloat16_t> {
63+ static constexpr const char* val = "bfloat16_t";
64+};
65+#endif
66+ 
67+TEST_F(silu_mul_test, test_silu_mul_dynamic)
68+{
69+ const char* dtypeStr = DataTypeName<DTYPE_X>::val;
70+ std::cout << ">>> Current Test Type: " << dtypeStr << std::endl;
71+ 
72+ system(
73+ "cp -rf "
74+ "../../../../activation/silu_mul/tests/ut/op_kernel/silu_mul_data ./");
75+ system("chmod -R 755 ./silu_mul_data/");
76+ 
77+ std::string genCmd = std::string("cd ./silu_mul_data/ && python3 gen_data.py '(2, 4)' '(2, 2)' '") + dtypeStr + "'";
78+ system(genCmd.c_str());
79+ 
80+ size_t M = 2;
81+ size_t N = 4;
82+ size_t D = N / 2;
83+ 
84+ size_t xFileSize = M * N * sizeof(DTYPE_X);
85+ size_t yFileSize = M * D * sizeof(DTYPE_X);
86+ 
87+ uint8_t* x = (uint8_t*)AscendC::GmAlloc(xFileSize);
88+ uint8_t* y = (uint8_t*)AscendC::GmAlloc(yFileSize);
89+ 
90+ uint64_t tilingKey = 0;
91+ uint32_t blockDim = 1;
92+ size_t workspaceFileSize = 16781184;
93+ size_t tilingDataSize = sizeof(SiluMulTilingData);
94+ uint8_t* workspace = (uint8_t*)AscendC::GmAlloc(workspaceFileSize);
95+ uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(tilingDataSize);
96+ 
97+ std::string fileName = std::string("./silu_mul_data/") + dtypeStr + "_input_silu_mul.bin";
98+ 
99+ ReadFile(fileName, xFileSize, x, xFileSize);
100+ 
101+ SiluMulTilingData* tilingDatafromBin = reinterpret_cast<SiluMulTilingData*>(tiling);
102+ tilingDatafromBin->lastDimSize = 4;
103+ tilingDatafromBin->batchSize = 2;
104+ tilingDatafromBin->PPMaxCalNum = 5888;
105+ tilingDatafromBin->needCoreNum = 1;
106+ 
107+ ICPU_SET_TILING_KEY(tilingKey);
108+ ICPU_RUN_KF(silu_mul, blockDim, x, y, workspace, (uint8_t*)tilingDatafromBin);
109+ 
110+ fileName = std::string("./silu_mul_data/") + dtypeStr + "_output_silu_mul.bin";
111+ WriteFile(fileName, y, yFileSize);
112+ 
113+ AscendC::GmFree((void*)x);
114+ AscendC::GmFree((void*)y);
115+ AscendC::GmFree((void*)workspace);
116+ AscendC::GmFree((void*)tiling);
117+ 
118+ std::string compareCmd = std::string("cd ./silu_mul_data/ && python3 compare_data.py '") + dtypeStr + "'";
119+ system(compareCmd.c_str());
120+}