feat: 新增SwigluGroup算子实现-ops-nn-AtomGit

Aactivation/swiglu_group/CMakeLists.txt+16-0

@@ -0,0 +1,16 @@
	1	+# ----------------------------------------------------------------------------
	2	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+# CANN Open Software License Agreement Version 2.0 (the "License").
	5	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+# See LICENSE in the root of the software repository for the full text of the License.
	9	+# ----------------------------------------------------------------------------
	10	+
	11	+set(SUPPORT_COMPUTE_UNIT "ascend950")
	12	+set(SUPPORT_TILING_DIR "arch35")
	13	+add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR}
	14	+ OPTYPE swiglu_group ACLNNTYPE aclnn
	15	+ COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR}
	16	+ DISABLE_IN_OPP TRUE)

Aactivation/swiglu_group/README.md+63-0

@@ -0,0 +1,63 @@
	1	+# SwigluGroup
	2	+
	3	+## 产品支持情况
	4	+
	5	+\| 产品 \| 是否支持 \|
	W W wangyongguang6月25日补充PR描述和issue
	6	+\| :--- \| :------: \|
	7	+\| <term>Ascend 950PR/Ascend 950DT</term> \| √ \|
	8	+\| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> \| × \|
	9	+\| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> \| × \|
	10	+\| <term>Atlas 200I/500 A2 推理产品</term> \| × \|
	11	+\| <term>Atlas 推理系列产品</term> \| × \|
	12	+\| <term>Atlas 训练系列产品</term> \| × \|
	13	+
	14	+## 功能说明
	15	+
	16	+- 算子功能：实现SwiGLU激活。输入`x`的最后一维被均分为`A`和`B`，计算`silu(A) * B`，直接输出与输入`x`相同数据类型的激活结果。
	17	+
	18	+- 计算公式：
	19	+
	20	+ $$
	21	+ y=silu(A) \times B
	22	+ $$
	23	+
	24	+ 当传入`clamp_limit`时：
	25	+
	26	+ $$
	27	+ A=min(A, clamp\_limit)
	28	+ $$
	29	+
	30	+ $$
	31	+ B=min(max(B, -clamp\_limit), clamp\_limit)
	32	+ $$
	33	+
	34	+ 当传入`weight`时：
	35	+
	36	+ $$
	37	+ y=y \times weight
	38	+ $$
	39	+
	40	+## 参数说明
	41	+
	42	+\| 参数名 \| 输入/输出/属性 \| 描述 \| 数据类型 \| 数据格式 \|
	43	+\| :--- \| :--- \| :--- \| :--- \| :--- \|
	44	+\| x \| 输入 \| 待计算的输入张量，最后一维被均分为两部分用于SwiGLU，需为正且能被2整除。 \| FLOAT16、BFLOAT16、FLOAT32 \| ND \|
	45	+\| weight \| 可选输入 \| 每个token的权重，乘到SwiGLU结果上。元素个数需等于`x`除最后一维外的维度乘积。 \| FLOAT32 \| ND \|
	46	+\| group_index \| 可选输入 \| count模式下的分组token数量。 \| INT64 \| ND \|
	47	+\| y \| 输出 \| SwiGLU计算结果，数据类型与`x`一致，最后一维为`x`最后一维的一半。 \| FLOAT16、BFLOAT16、FLOAT32 \| ND \|
	48	+\| clamp_limit \| 可选属性 \| 默认值-1.0，表示不进行clamp；若设置为正数，则在激活前对SwiGLU输入做clamp。 \| FLOAT \| - \|
	49	+
	50	+## 约束说明
	51	+
	52	+- 输入`x`的最后一维需为正，且能被2整除（被均分为`A`、`B`两部分，输出最后一维为输入的一半）。
	53	+- 输入`x`的数据类型仅支持FLOAT16、BFLOAT16、FLOAT32，且`y`的数据类型需与`x`一致。
	54	+- 当传入`weight`时，数据类型为FLOAT32，且元素个数需等于`x`除最后一维外的维度乘积。
	55	+- 当传入`group_index`时，数据类型为INT64，为count模式分组token数，实际处理行数为`group_index`所有元素之和与`bs`两者中的较小值（`bs`为`x`除最后一维外的维度乘积）。
	56	+- 当传入`group_index`时，调用者需保证`group_index`所有元素之和非负。
	57	+
	58	+## 调用说明
	59	+
	60	+\| 调用方式 \| 调用样例 \| 说明 \|
	61	+\| :--- \| :--- \| :--- \|
	62	+\| aclnn调用 \| [test_aclnn_swiglu_group](./examples/test_aclnn_swiglu_group.cpp) \| 通过[aclnnSwigluGroup](./docs/aclnnSwigluGroup.md)接口方式调用SwigluGroup算子。 \|
	63	+\| 图模式调用 \| - \| 通过[算子IR](./op_graph/swiglu_group_proto.h)构图方式调用SwigluGroup算子。 \|

Aactivation/swiglu_group/docs/aclnnSwigluGroup.md+405-0

@@ -0,0 +1,405 @@
	1	+# aclnnSwigluGroup
	2	+
	3	+[📄 查看源码](https://gitcode.com/cann/ops-nn/tree/master/activation/swiglu_group)
	4	+
	5	+## 产品支持情况
	6	+
	7	+\| 产品 \| 是否支持 \|
	8	+\| :--- \| :------: \|
	9	+\| <term>Ascend 950PR/Ascend 950DT</term> \| √ \|
	10	+\| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> \| × \|
	11	+\| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> \| × \|
	12	+\| <term>Atlas 200I/500 A2 推理产品</term> \| × \|
	13	+\| <term>Atlas 推理系列产品</term> \| × \|
	14	+\| <term>Atlas 训练系列产品</term> \| × \|
	W W wang-xing0016月26日确认一下这个算子是否只支持950，不支持A2和A3 L luwei_1104 6月26日评论： 6月26日评论：
	15	+
	16	+## 功能说明
	17	+
	18	+- 接口功能：执行SwiGLU激活，输出与输入x相同数据类型的激活结果。
	19	+
	20	+- 计算公式：
	21	+
	22	+ 令输入x的最后一维大小为D，左半部分为A，右半部分为B（即`h = D // 2`，`A = x[..., :h]`，`B = x[..., h:]`）。
	23	+
	24	+ 当`clampLimit`大于0时，先对A、B进行截断：
	25	+
	26	+ $$
	27	+ A = min(A, clampLimit)
	28	+ $$
	29	+
	30	+ $$
	31	+ B = min(max(B, -clampLimit), clampLimit)
	32	+ $$
	33	+
	34	+ 计算SwiGLU：
	35	+
	36	+ $$
	37	+ y = silu(A) \times B = \frac{A}{1 + e^{-A}} \times B
	38	+ $$
	39	+
	40	+ 当`weightOptional`非空时，按token乘以权重：
	41	+
	42	+ $$
	43	+ y = y \times weightOptional
	44	+ $$
	45	+
	46	+ 当`groupIndexOptional`非空时，按count模式控制实际处理的token数：实际处理的token数为`groupIndexOptional`所有元素之和与`bs`两者中的较小值，其中`bs`为x除最后一维外的维度乘积。
	47	+
	48	+## 函数原型
	49	+
	50	+每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md)，必须先调用“aclnnSwigluGroupGetWorkspaceSize”接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用“aclnnSwigluGroup”接口执行计算。
	51	+
	52	+```Cpp
	53	+aclnnStatus aclnnSwigluGroupGetWorkspaceSize(
	54	+ const aclTensor *x,
	55	+ const aclTensor *weightOptional,
	56	+ const aclTensor *groupIndexOptional,
	57	+ double clampLimit,
	58	+ const aclTensor *out,
	59	+ uint64_t *workspaceSize,
	60	+ aclOpExecutor **executor);
	61	+```
	62	+
	63	+```Cpp
	64	+aclnnStatus aclnnSwigluGroup(
	65	+ void *workspace,
	66	+ uint64_t workspaceSize,
	67	+ aclOpExecutor *executor,
	68	+ aclrtStream stream);
	69	+```
	70	+
	71	+## aclnnSwigluGroupGetWorkspaceSize
	72	+
	73	+- 参数说明：
	74	+
	75	+ <table style="undefined;table-layout: fixed; width: 1567px"><colgroup>
	76	+ <col style="width: 200px">
	77	+ <col style="width: 100px">
	78	+ <col style="width: 300px">
	79	+ <col style="width: 360px">
	80	+ <col style="width: 230px">
	81	+ <col style="width: 100px">
	82	+ <col style="width: 100px">
	83	+ <col style="width: 120px">
	84	+ </colgroup>
	85	+ <thead>
	86	+ <tr>
	87	+ <th>参数名</th>
	88	+ <th>输入/输出</th>
	89	+ <th>描述</th>
	90	+ <th>使用说明</th>
	91	+ <th>数据类型</th>
	92	+ <th>数据格式</th>
	93	+ <th>维度(shape)</th>
	94	+ <th>非连续Tensor</th>
	95	+ </tr></thead>
	96	+ <tbody>
	97	+ <tr>
	98	+ <td>x（aclTensor*）</td>
	99	+ <td>输入</td>
	100	+ <td>公式中的输入x，最后一维被均分为A、B两部分。</td>
	101	+ <td>不支持空指针，不支持空Tensor。维度必须大于0，最后一维必须为偶数。</td>
	102	+ <td>FLOAT16、BFLOAT16、FLOAT32</td>
	103	+ <td>ND</td>
	104	+ <td>1-8</td>
	105	+ <td>×</td>
	106	+ </tr>
	107	+ <tr>
	108	+ <td>weightOptional（aclTensor*）</td>
	109	+ <td>输入</td>
	110	+ <td>公式中的输入weightOptional，逐token权重。</td>
	111	+ <td>支持空指针，不支持空Tensor。不为空指针时，元素个数需等于x除最后一维外的维度乘积（bs）。</td>
	112	+ <td>FLOAT32</td>
	113	+ <td>ND</td>
	114	+ <td>1-8</td>
	115	+ <td>×</td>
	116	+ </tr>
	117	+ <tr>
	118	+ <td>groupIndexOptional（aclTensor*）</td>
	119	+ <td>输入</td>
	120	+ <td>公式中的输入groupIndexOptional，count模式下的分组token数量。</td>
	121	+ <td>支持空指针，不支持空Tensor。不为空指针时，维度要求为1维，元素需大于等于0。第i个元素代表第i组需要处理x的token数量，所有元素之和经截断后为实际处理的token数。</td>
	122	+ <td>INT64</td>
	123	+ <td>ND</td>
	124	+ <td>1</td>
	125	+ <td>×</td>
	126	+ </tr>
	127	+ <tr>
	128	+ <td>clampLimit（double）</td>
	129	+ <td>输入</td>
	130	+ <td>公式中的输入clampLimit，表示激活前对A、B进行截断的上界。</td>
	131	+ <td>传入-1.0表示不进行截断；传入大于0的值时，在激活前对A、B进行截断。</td>
	132	+ <td>-</td>
	133	+ <td>-</td>
	134	+ <td>-</td>
	135	+ <td>不涉及</td>
	136	+ </tr>
	137	+ <tr>
	138	+ <td>out（aclTensor*）</td>
	139	+ <td>输出</td>
	140	+ <td>公式中的输出y，SwiGLU计算结果。</td>
	141	+ <td>数据类型与x一致，最后一维为x最后一维的一半，其余维度与x一致。</td>
	142	+ <td>FLOAT16、BFLOAT16、FLOAT32</td>
	143	+ <td>ND</td>
	144	+ <td>1-8</td>
	145	+ <td>×</td>
	146	+ </tr>
	147	+ <tr>
	148	+ <td>workspaceSize（uint64_t*）</td>
	149	+ <td>输出</td>
	150	+ <td>返回需要在Device侧申请的workspace大小。</td>
	151	+ <td>-</td>
	152	+ <td>-</td>
	153	+ <td>-</td>
	154	+ <td>-</td>
	155	+ <td>不涉及</td>
	156	+ </tr>
	157	+ <tr>
	158	+ <td>executor（aclOpExecutor**）</td>
	159	+ <td>输出</td>
	160	+ <td>返回op执行器，包含了算子计算流程。</td>
	161	+ <td>-</td>
	162	+ <td>-</td>
	163	+ <td>-</td>
	164	+ <td>-</td>
	165	+ <td>不涉及</td>
	166	+ </tr>
	167	+ </tbody>
	168	+ </table>
	169	+
	170	+- 返回值：
	171	+
	172	+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
	173	+
	174	+ 第一段接口会完成入参校验，出现以下场景时报错：
	175	+
	176	+ <table style="undefined;table-layout: fixed;width: 979px"><colgroup>
	177	+ <col style="width: 272px">
	178	+ <col style="width: 103px">
	179	+ <col style="width: 604px">
	180	+ </colgroup>
	181	+ <thead>
	182	+ <tr>
	183	+ <th>返回码</th>
	184	+ <th>错误码</th>
	185	+ <th>描述</th>
	186	+ </tr>
	187	+ </thead>
	188	+ <tbody>
	189	+ <tr>
	190	+ <td>ACLNN_ERR_PARAM_NULLPTR</td>
	191	+ <td>161001</td>
	192	+ <td>传入的x、out是空指针。</td>
	193	+ </tr>
	194	+ <tr>
	195	+ <td rowspan="3">ACLNN_ERR_PARAM_INVALID</td>
	196	+ <td rowspan="3">161002</td>
	197	+ <td>x、out的数据类型不在支持的范围内，或out与x的数据类型不一致。</td>
	198	+ </tr>
	199	+ <tr>
	200	+ <td>x的最后一维不是偶数，或维度不在支持的范围内。</td>
	201	+ </tr>
	202	+ <tr>
	203	+ <td>weightOptional的元素个数与x除最后一维外的维度乘积不一致。</td>
	204	+ </tr>
	205	+ </tbody>
	206	+ </table>
	207	+
	208	+## aclnnSwigluGroup
	209	+
	210	+- 参数说明：
	211	+
	212	+ <table style="undefined;table-layout: fixed; width: 953px"><colgroup>
	213	+ <col style="width: 173px">
	214	+ <col style="width: 112px">
	215	+ <col style="width: 668px">
	216	+ </colgroup>
	217	+ <thead>
	218	+ <tr>
	219	+ <th>参数名</th>
	220	+ <th>输入/输出</th>
	221	+ <th>描述</th>
	222	+ </tr></thead>
	223	+ <tbody>
	224	+ <tr>
	225	+ <td>workspace</td>
	226	+ <td>输入</td>
	227	+ <td>在Device侧申请的workspace内存地址。</td>
	228	+ </tr>
	229	+ <tr>
	230	+ <td>workspaceSize</td>
	231	+ <td>输入</td>
	232	+ <td>在Device侧申请的workspace大小，由第一段接口aclnnSwigluGroupGetWorkspaceSize获取。</td>
	233	+ </tr>
	234	+ <tr>
	235	+ <td>executor</td>
	236	+ <td>输入</td>
	237	+ <td>op执行器，包含了算子计算流程。</td>
	238	+ </tr>
	239	+ <tr>
	240	+ <td>stream</td>
	241	+ <td>输入</td>
	242	+ <td>指定执行任务的Stream。</td>
	243	+ </tr>
	244	+ </tbody>
	245	+ </table>
	246	+
	247	+- 返回值：
	248	+
	249	+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
	250	+
	251	+## 约束说明
	252	+
	253	+- 不支持空Tensor：输入x、weightOptional、groupIndexOptional均不支持空Tensor（其中weightOptional、groupIndexOptional可传入空指针，但传入时不可为空Tensor）。
	254	+- 不支持非连续Tensor：输入x、weightOptional、groupIndexOptional及输出out均需为连续Tensor。
	255	+- 输入x的最后一维必须为偶数（被均分为A、B两部分，输出最后一维为输入的一半）。
	256	+- 输入x的数据类型仅支持FLOAT16、BFLOAT16、FLOAT32，且out的数据类型需与x一致。
	257	+- 确定性计算：aclnnSwigluGroup默认为确定性实现，暂不支持非确定性实现。
	258	+
	259	+## 调用示例
	260	+
	261	+示例代码如下，仅供参考，具体编译和执行过程请参考[编译与运行样例](../../../docs/zh/context/编译与运行样例.md)。
	262	+
	263	+```Cpp
	264	+#include <cstdint>
	265	+#include <cstring>
	266	+#include <iostream>
	267	+#include <vector>
	268	+#include "acl/acl.h"
	269	+#include "aclnnop/aclnn_swiglu_group.h"
	270	+
	271	+#define CHECK_RET(cond, return_expr) \
	272	+ do { \
	273	+ if (!(cond)) { \
	274	+ return_expr; \
	275	+ } \
	276	+ } while (0)
	277	+
	278	+#define LOG_PRINT(message, ...) \
	279	+ do { \
	280	+ printf(message, ##__VA_ARGS__); \
	281	+ } while (0)
	282	+
	283	+int64_t GetShapeSize(const std::vector<int64_t>& shape)
	284	+{
	285	+ int64_t shapeSize = 1;
	286	+ for (auto i : shape) {
	287	+ shapeSize *= i;
	288	+ }
	289	+ return shapeSize;
	290	+}
	291	+
	292	+int Init(int32_t deviceId, aclrtStream* stream)
	293	+{
	294	+ // 固定写法，acl初始化
	295	+ auto ret = aclInit(nullptr);
	296	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
	297	+ ret = aclrtSetDevice(deviceId);
	298	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
	299	+ ret = aclrtCreateStream(stream);
	300	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
	301	+ return ACL_SUCCESS;
	302	+}
	303	+
	304	+template <typename T>
	305	+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
	306	+ aclDataType dataType, aclTensor** tensor)
	307	+{
	308	+ auto size = GetShapeSize(shape) * sizeof(T);
	309	+ // 调用aclrtMalloc申请Device侧内存
	310	+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
	311	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
	312	+ // 调用aclrtMemcpy将Host侧数据拷贝到Device侧内存上
	313	+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
	314	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
	315	+
	316	+ // 计算连续tensor的strides
	317	+ std::vector<int64_t> strides(shape.size(), 1);
	318	+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
	319	+ strides[i] = shape[i + 1] * strides[i + 1];
	320	+ }
	321	+
	322	+ // 调用aclCreateTensor接口创建aclTensor
	323	+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, ACL_FORMAT_ND,
	324	+ shape.data(), shape.size(), *deviceAddr);
	325	+ return ACL_SUCCESS;
	326	+}
	327	+
	328	+int main()
	329	+{
	330	+ // 1. （固定写法）device/stream初始化，参考acl对外接口列表
	331	+ int32_t deviceId = 0;
	332	+ aclrtStream stream;
	333	+ auto ret = Init(deviceId, &stream);
	334	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
	335	+
	336	+ // 2. 构造输入与输出，需要根据API的接口自定义构造
	337	+ std::vector<int64_t> xShape = {2, 256};
	338	+ std::vector<int64_t> yShape = {2, 128};
	339	+
	340	+ std::vector<uint16_t> xHostData(GetShapeSize(xShape), 0);
	341	+ for (size_t i = 0; i < xHostData.size(); ++i) {
	342	+ xHostData[i] = static_cast<uint16_t>(i % 23);
	343	+ }
	344	+ std::vector<uint16_t> yHostData(GetShapeSize(yShape), 0);
	345	+
	346	+ void* xDeviceAddr = nullptr;
	347	+ void* yDeviceAddr = nullptr;
	348	+ aclTensor* x = nullptr;
	349	+ aclTensor* y = nullptr;
	350	+
	351	+ // 创建x aclTensor
	352	+ ret = CreateAclTensor(xHostData, xShape, &xDeviceAddr, ACL_FLOAT16, &x);
	353	+ CHECK_RET(ret == ACL_SUCCESS, return ret);
	354	+ // 创建y aclTensor
	355	+ ret = CreateAclTensor(yHostData, yShape, &yDeviceAddr, ACL_FLOAT16, &y);
	356	+ CHECK_RET(ret == ACL_SUCCESS, return ret);
	357	+
	358	+ // weightOptional与groupIndexOptional为可选输入，此处传入nullptr表示不使用
	359	+ double clampLimit = -1.0;
	360	+
	361	+ // 3. 调用CANN算子库API，需要修改为具体的Api名称
	362	+ uint64_t workspaceSize = 0;
	363	+ aclOpExecutor* executor = nullptr;
	364	+ // 调用aclnnSwigluGroup第一段接口
	365	+ ret = aclnnSwigluGroupGetWorkspaceSize(x, nullptr, nullptr, clampLimit, y, &workspaceSize, &executor);
	366	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSwigluGroupGetWorkspaceSize failed. ERROR: %d\n", ret);
	367	+ return ret);
	368	+ // 根据第一段接口计算出的workspaceSize申请Device内存
	369	+ void* workspaceAddr = nullptr;
	370	+ if (workspaceSize > 0) {
	371	+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
	372	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
	373	+ }
	374	+ // 调用aclnnSwigluGroup第二段接口
	375	+ ret = aclnnSwigluGroup(workspaceAddr, workspaceSize, executor, stream);
	376	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSwigluGroup failed. ERROR: %d\n", ret); return ret);
	377	+
	378	+ // 4. （固定写法）同步等待任务执行结束
	379	+ ret = aclrtSynchronizeStream(stream);
	380	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
	381	+
	382	+ // 5. 获取输出的值，将Device侧内存上的结果拷贝至Host侧
	383	+ std::vector<uint16_t> resultData(GetShapeSize(yShape), 0);
	384	+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), yDeviceAddr,
	385	+ resultData.size() * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
	386	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret);
	387	+ return ret);
	388	+ LOG_PRINT("result[0] is: %u\n", resultData[0]);
	389	+
	390	+ // 6. 释放aclTensor和aclScalar，需要根据具体API的接口定义修改
	391	+ aclDestroyTensor(x);
	392	+ aclDestroyTensor(y);
	393	+
	394	+ // 7. 释放Device资源，需要根据具体API的接口定义修改
	395	+ aclrtFree(xDeviceAddr);
	396	+ aclrtFree(yDeviceAddr);
	397	+ if (workspaceSize > 0) {
	398	+ aclrtFree(workspaceAddr);
	399	+ }
	400	+ aclrtDestroyStream(stream);
	401	+ aclrtResetDevice(deviceId);
	402	+ aclFinalize();
	403	+ return 0;
	404	+}
	405	+```

Aactivation/swiglu_group/examples/test_aclnn_swiglu_group.cpp+147-0

@@ -0,0 +1,147 @@
	1	+#include <cstdint>
	2	+#include <cstring>
	3	+#include <iostream>
	4	+#include <vector>
	5	+#include "acl/acl.h"
	6	+#include "aclnnop/aclnn_swiglu_group.h"
	7	+
	8	+#define CHECK_RET(cond, return_expr) \
	9	+ do { \
	10	+ if (!(cond)) { \
	11	+ return_expr; \
	12	+ } \
	13	+ } while (0)
	14	+
	15	+#define LOG_PRINT(message, ...) \
	16	+ do { \
	17	+ printf(message, ##__VA_ARGS__); \
	18	+ } while (0)
	19	+
	20	+int64_t GetShapeSize(const std::vector<int64_t>& shape) {
	21	+ int64_t shapeSize = 1;
	22	+ for (auto dim : shape) {
	23	+ shapeSize *= dim;
	24	+ }
	25	+ return shapeSize;
	26	+}
	27	+
	28	+int Init(int32_t deviceId, aclrtStream* stream) {
	29	+ auto ret = aclInit(nullptr);
	30	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
	31	+ ret = aclrtSetDevice(deviceId);
	32	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
	33	+ ret = aclrtCreateStream(stream);
	34	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
	35	+ return ACL_SUCCESS;
	36	+}
	37	+
	38	+bool CheckHardwareSupport() {
	39	+ const char* socName = aclrtGetSocName();
	40	+ if (socName == nullptr) {
	41	+ LOG_PRINT("Warning: Cannot get SOC name, skip hardware check\n");
	42	+ return true;
	43	+ }
	44	+
	45	+ LOG_PRINT("Current SOC: %s\n", socName);
	46	+ if (strstr(socName, "Ascend950") != nullptr \|\| strstr(socName, "ascend950") != nullptr) {
	47	+ return true;
	48	+ }
	49	+
	50	+ LOG_PRINT("Warning: SwigluGroup only supports Ascend950, current SOC '%s' is not supported. Skip test.\n",
	51	+ socName);
	52	+ return false;
	53	+}
	54	+
	55	+void Finalize(int32_t deviceId, aclrtStream stream) {
	56	+ (void)aclrtDestroyStream(stream);
	57	+ (void)aclrtResetDevice(deviceId);
	58	+ (void)aclFinalize();
	59	+}
	60	+
	61	+template <typename T>
	62	+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
	63	+ aclDataType dataType, aclTensor** tensor) {
	64	+ auto size = GetShapeSize(shape) * sizeof(T);
	65	+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
	66	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
	67	+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
	68	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
	69	+
	70	+ std::vector<int64_t> strides(shape.size(), 1);
	71	+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
	72	+ strides[i] = shape[i + 1] * strides[i + 1];
	73	+ }
	74	+
	75	+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, ACL_FORMAT_ND,
	76	+ shape.data(), shape.size(), *deviceAddr);
	77	+ return ACL_SUCCESS;
	78	+}
	79	+
	80	+int main() {
	81	+ int32_t deviceId = 0;
	82	+ aclrtStream stream;
	83	+ auto ret = Init(deviceId, &stream);
	84	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
	85	+
	86	+ if (!CheckHardwareSupport()) {
	87	+ LOG_PRINT("\n=== Test SKIPPED (hardware not supported) ===\n");
	88	+ Finalize(deviceId, stream);
	89	+ return ACL_SUCCESS;
	90	+ }
	91	+
	92	+ std::vector<int64_t> xShape = {2, 256};
	93	+ std::vector<int64_t> yShape = {2, 128};
	94	+
	95	+ std::vector<uint16_t> xHostData(GetShapeSize(xShape), 0);
	96	+ for (size_t i = 0; i < xHostData.size(); ++i) {
	97	+ xHostData[i] = static_cast<uint16_t>(i % 23);
	98	+ }
	99	+ std::vector<uint16_t> yHostData(GetShapeSize(yShape), 0);
	100	+
	101	+ void* xDeviceAddr = nullptr;
	102	+ void* yDeviceAddr = nullptr;
	103	+ aclTensor* x = nullptr;
	104	+ aclTensor* y = nullptr;
	105	+
	106	+ ret = CreateAclTensor(xHostData, xShape, &xDeviceAddr, ACL_FLOAT16, &x);
	107	+ CHECK_RET(ret == ACL_SUCCESS, return ret);
	108	+ ret = CreateAclTensor(yHostData, yShape, &yDeviceAddr, ACL_FLOAT16, &y);
	109	+ CHECK_RET(ret == ACL_SUCCESS, return ret);
	atomgit-bot6月23日 🟡 Medium Priority 在 `main()` 函数第106-109行中：如果 `CreateAclTensor(xHostData, ...)` 成功但 `CreateAclTensor(yHostData, ...)` 失败，代码执行 `CHECK_RET(ret == ACL_SUCCESS, return ret)` 直接返回，此时 `xDeviceAddr` 和 `x`（已分配）未释放，造成 device 内存和 aclTensor 泄漏。此外，`Init()` 已成功执行（分配了 stream），但在错误路径上 `Finalize()` 未被调用，导致 stream 和 device 资源未清理。建议：将 `CHECK_RET` 改为跳转到统一的清理标签（如 `goto cleanup`），在清理标签中检查并释放已分配的资源，并调用 `Finalize`。或者在每个 CreateAclTensor 失败时手动释放已分配资源再返回。 atomgit-bot6月24日 🟡 Medium Priority 在 `main()` 函数第106-109行中：第106行 `CreateAclTensor(xHostData, ...)` 成功创建了 `xDeviceAddr` 和 `x`（aclTensor），但如果第108行 `CreateAclTensor(yHostData, ...)` 失败，`CHECK_RET` 宏执行 `return ret` 直接退出，此时 `xDeviceAddr`（通过 `aclrtMalloc` 分配）和 `x`（通过 `aclCreateTensor` 创建）均未被释放，造成 device 内存和 tensor 资源泄漏。类似地，后续任何失败（如 workspace malloc 失败）也存在同样的泄漏风险。建议：将 `CHECK_RET` 改为跳转到统一的清理标签（如 `goto cleanup`），在清理标签中检查并释放已分配的资源（xDeviceAddr、x、yDeviceAddr、y、workspaceAddr），并调用 `Finalize`。或者在每个 CreateAclTensor 失败时手动释放已分配资源再返回。
	110	+
	111	+ // weight and group_index are optional; pass nullptr to skip them.
	112	+ double clampLimit = -1.0;
	113	+
	114	+ uint64_t workspaceSize = 0;
	115	+ aclOpExecutor* executor = nullptr;
	116	+ ret = aclnnSwigluGroupGetWorkspaceSize(x, nullptr, nullptr, clampLimit, y, &workspaceSize, &executor);
	117	+ CHECK_RET(ret == ACL_SUCCESS,
	118	+ LOG_PRINT("aclnnSwigluGroupGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
	119	+
	120	+ void* workspaceAddr = nullptr;
	121	+ if (workspaceSize > 0) {
	122	+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
	123	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
	124	+ }
	125	+
	126	+ ret = aclnnSwigluGroup(workspaceAddr, workspaceSize, executor, stream);
	127	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSwigluGroup failed. ERROR: %d\n", ret); return ret);
	128	+
	129	+ ret = aclrtSynchronizeStream(stream);
	130	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
	131	+
	132	+ std::vector<uint16_t> resultData(GetShapeSize(yShape), 0);
	133	+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), yDeviceAddr,
	134	+ resultData.size() * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
	135	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
	136	+ LOG_PRINT("result[0] is: %u\n", resultData[0]);
	137	+
	138	+ aclDestroyTensor(x);
	139	+ aclDestroyTensor(y);
	140	+ aclrtFree(xDeviceAddr);
	141	+ aclrtFree(yDeviceAddr);
	142	+ if (workspaceSize > 0) {
	143	+ aclrtFree(workspaceAddr);
	144	+ }
	145	+ Finalize(deviceId, stream);
	146	+ return ACL_SUCCESS;
	147	+}

Aactivation/swiglu_group/op_graph/CMakeLists.txt+11-0

@@ -0,0 +1,11 @@
	1	+# ----------------------------------------------------------------------------
	2	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+# CANN Open Software License Agreement Version 2.0 (the "License").
	5	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+# See LICENSE in the root of the software repository for the full text of the License.
	9	+# ----------------------------------------------------------------------------
	10	+
	11	+add_graph_plugin_sources()

Aactivation/swiglu_group/op_graph/swiglu_group_proto.h+52-0

@@ -0,0 +1,52 @@
	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.
	W W wangyongguang6月25日注释找陈萍萍看下
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file swiglu_group_proto.h
	13	+ * \brief SwiGLU activation with optional per-token weight and grouped tokens.
	14	+ */
	15	+
	16	+#ifndef ACTIVATION_SWIGLU_GROUP_PROTO_H_
	17	+#define ACTIVATION_SWIGLU_GROUP_PROTO_H_
	18	+
	19	+#include "graph/operator_reg.h"
	20	+
	21	+namespace ge {
	22	+
	23	+/**
	24	+* @brief Performs SwiGLU activation.
	25	+*
	26	+* @par Inputs:
	27	+* @li x: Required tensor of type float16, bfloat16 or float32. The last dimension is split into two
	28	+* equal parts for SwiGLU and must be divisible by 2.
	29	+* @li weight: Optional float32 tensor. Per-token weight multiplied into the SwiGLU result.
	30	+* @li group_index: Optional int64 tensor. Count-mode group token numbers.
	31	+*
	32	+* @par Attributes:
	33	+* @li clamp_limit: Optional float. Defaults to -1.0, which disables clamp. If set to a positive value,
	34	+* clamps SwiGLU inputs before activation.
	35	+*
	36	+* @par Outputs:
	37	+* @li y: SwiGLU result tensor with the same dtype as x and last dimension halved.
	38	+*
	39	+* @par Third-party framework compatibility
	40	+* It is a custom operator. It has no corresponding operator in Caffe, ONNX, TensorFlow, or PyTorch.
	41	+*/
	42	+REG_OP(SwigluGroup)
	43	+ .INPUT(x, TensorType({DT_FLOAT16, DT_BF16, DT_FLOAT}))
	44	+ .OPTIONAL_INPUT(weight, TensorType({DT_FLOAT}))
	45	+ .OPTIONAL_INPUT(group_index, TensorType({DT_INT64}))
	46	+ .OUTPUT(y, TensorType({DT_FLOAT16, DT_BF16, DT_FLOAT}))
	47	+ .ATTR(clamp_limit, Float, -1.0f)
	48	+ .OP_END_FACTORY_REG(SwigluGroup)
	49	+
	50	+} // namespace ge
	51	+
	52	+#endif // ACTIVATION_SWIGLU_GROUP_PROTO_H_

Aactivation/swiglu_group/op_host/arch35/swiglu_group_tiling.cpp+466-0

@@ -0,0 +1,466 @@
	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 swiglu_group_tiling.cpp
	13	+ * \brief
	14	+ */
	15	+
	16	+#include <algorithm>
	17	+#include <cmath>
	18	+#include <sstream>
	19	+#include "swiglu_group_tiling.h"
	20	+
	21	+using namespace ge;
	22	+namespace optiling {
	23	+namespace {
	24	+constexpr uint64_t WORKSPACE_SIZE = 32;
	25	+int64_t CeilDiv(int64_t x, int64_t y)
	26	+{
	27	+ if (y != 0) {
	28	+ return (x + y - 1) / y;
	29	+ }
	30	+ return x;
	31	+}
	32	+int64_t DownAlign(int64_t x, int64_t y) {
	33	+ if (y == 0) {
	34	+ return x;
	35	+ }
	36	+ return (x / y) * y;
	37	+}
	38	+int64_t RoundUp(int64_t x, int64_t y) {
	39	+ return CeilDiv(x, y) * y;
	40	+}
	41	+
	42	+constexpr int64_t BLOCK_SIZE = 32;
	43	+constexpr int64_t NUM_TWO = 2;
	44	+constexpr int64_t DOUBLE_BUFFER = 2;
	45	+constexpr int64_t B16_BYTES = 2;
	46	+constexpr int64_t B32_BYTES = 4;
	47	+constexpr int64_t B16_ALIGN_NUM = BLOCK_SIZE / B16_BYTES;
	48	+constexpr int64_t B32_ALIGN_NUM = BLOCK_SIZE / B32_BYTES;
	49	+// Ascend950 cacheline size is 512B. Split d-factor by one cacheline so each chunk is cacheline
	50	+// aligned; the element count per cacheline depends on the x dtype size (2 bytes for fp16/bf16,
	51	+// 4 bytes for float32).
	52	+constexpr int64_t ASCEND950_CACHE_LINE_BYTES = 512;
	53	+constexpr size_t ATTR_INDEX_CLAMP_LIMIT = 0;
	54	+constexpr size_t INPUT_INDEX_X = 0;
	55	+constexpr size_t INPUT_INDEX_WEIGHT = 1;
	56	+constexpr size_t INPUT_INDEX_GROUP_INDEX = 2;
	57	+constexpr size_t OUTPUT_INDEX_Y = 0;
	58	+constexpr size_t CACHE_LINE_SIZE = 128;
	59	+constexpr float DEFAULT_CLAMP_LIMIT = -1.0f;
	60	+constexpr int64_t SWIGLU_GROUP_TILING_KEY = 1000;
	61	+
	62	+int64_t ShapeElementNum(const gert::Shape& shape)
	63	+{
	64	+ int64_t elementNum = 1;
	65	+ for (size_t i = 0; i < shape.GetDimNum(); i++) {
	66	+ elementNum *= shape.GetDim(i);
	67	+ }
	68	+ return elementNum;
	69	+}
	70	+}
	71	+
	72	+ge::graphStatus SwigluGroupTiling::GetPlatformInfo()
	73	+{
	74	+ auto platformInfo = context_->GetPlatformInfo();
	75	+ if (platformInfo == nullptr) {
	76	+ auto compileInfoPtr = context_->GetCompileInfo<SwigluGroupCompileInfo>();
	77	+ OP_CHECK_IF(compileInfoPtr == nullptr, OP_LOGE(context_->GetNodeName(), "compile info is null"),
	78	+ return ge::GRAPH_FAILED);
	79	+ coreNum_ = compileInfoPtr->coreNum;
	80	+ ubSize_ = compileInfoPtr->ubSize;
	81	+ } else {
	82	+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo);
	83	+ coreNum_ = ascendcPlatform.GetCoreNumAiv();
	84	+ uint64_t ubSizePlatForm;
	85	+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSizePlatForm);
	86	+ ubSize_ = ubSizePlatForm;
	87	+ }
	88	+ return ge::GRAPH_SUCCESS;
	89	+}
	90	+
	91	+ge::graphStatus SwigluGroupTiling::GetClampLimitAttr(const gert::RuntimeAttrs* attrs)
	92	+{
	93	+ auto clampLimitAttr = attrs->GetAttrPointer<float>(ATTR_INDEX_CLAMP_LIMIT);
	94	+ if (clampLimitAttr != nullptr) {
	95	+ // DEFAULT_CLAMP_LIMIT means user did not pass clamp_limit.
	96	+ if (*clampLimitAttr != DEFAULT_CLAMP_LIMIT) {
	97	+ OP_CHECK_IF(!(*clampLimitAttr > 0.0f),
	98	+ OP_LOGE(context_->GetNodeName(), "attr clamp_limit should be greater than 0.0, got %f.",
	99	+ *clampLimitAttr),
	100	+ return ge::GRAPH_FAILED);
	101	+ clampLimit_ = *clampLimitAttr;
	102	+ hasClampLimit_ = 1;
	103	+ }
	104	+ }
	105	+ return ge::GRAPH_SUCCESS;
	106	+}
	107	+
	108	+ge::graphStatus SwigluGroupTiling::GetAttr()
	109	+{
	110	+ auto* attrs = context_->GetAttrs();
	111	+ OP_CHECK_NULL_WITH_CONTEXT(context_, attrs);
	112	+
	113	+ if (GetClampLimitAttr(attrs) == ge::GRAPH_FAILED) {
	114	+ return ge::GRAPH_FAILED;
	115	+ }
	116	+
	117	+ return ge::GRAPH_SUCCESS;
	118	+}
	119	+
	120	+ge::graphStatus SwigluGroupTiling::CheckWeightInfo()
	121	+{
	122	+ auto weightDesc = context_->GetOptionalInputDesc(INPUT_INDEX_WEIGHT);
	123	+ if (weightDesc != nullptr) {
	124	+ auto weightDtype = weightDesc->GetDataType();
	125	+ OP_CHECK_IF((weightDtype != ge::DT_FLOAT),
	126	+ OP_LOGE(context_->GetNodeName(), "input weight dtype should be FLOAT, got %d.",
	127	+ static_cast<int>(weightDtype)),
	128	+ return ge::GRAPH_FAILED);
	129	+ auto weightShape = context_->GetOptionalInputShape(INPUT_INDEX_WEIGHT);
	130	+ if (weightShape != nullptr) {
	131	+ auto weightStorageShape = weightShape->GetStorageShape();
	132	+ auto weightElementNum = ShapeElementNum(weightStorageShape);
	133	+ OP_CHECK_IF((weightElementNum != bs_),
	134	+ OP_LOGE(context_->GetNodeName(),
	135	+ "input weight element num should be equal to input x outer dim product, got %ld, "
	136	+ "expected %ld.",
	137	+ weightElementNum, bs_),
	138	+ return ge::GRAPH_FAILED);
	139	+ hasWeight_ = true;
	140	+ }
	141	+ }
	142	+ return ge::GRAPH_SUCCESS;
	143	+}
	144	+
	145	+ge::graphStatus SwigluGroupTiling::CheckGroupIndexInfo()
	146	+{
	147	+ auto groupIndexDesc = context_->GetOptionalInputDesc(INPUT_INDEX_GROUP_INDEX);
	148	+ if (groupIndexDesc != nullptr) {
	149	+ auto groupIndexDtype = groupIndexDesc->GetDataType();
	150	+ OP_CHECK_IF((groupIndexDtype != ge::DT_INT64),
	151	+ OP_LOGE(context_->GetNodeName(), "input group_index dtype should be INT64, got %d.",
	152	+ static_cast<int>(groupIndexDtype)),
	153	+ return ge::GRAPH_FAILED);
	154	+ auto groupIndexShape = context_->GetOptionalInputShape(INPUT_INDEX_GROUP_INDEX);
	155	+ if (groupIndexShape != nullptr) {
	156	+ auto groupIndexStorageShape = groupIndexShape->GetStorageShape();
	157	+ g_ = 1;
	158	+ for (size_t i = 0; i < groupIndexStorageShape.GetDimNum(); i++) {
	159	+ g_ = g_ * groupIndexStorageShape.GetDim(i);
	160	+ }
	161	+ // Empty tensor is not supported: a passed group_index must have a positive element count.
	162	+ OP_CHECK_IF((g_ <= 0),
	163	+ OP_LOGE(context_->GetNodeName(),
	164	+ "input group_index is empty tensor, which is not supported, got element num %ld.", g_),
	165	+ return ge::GRAPH_FAILED);
	166	+ hasGroupIndex_ = true;
	167	+ }
	168	+ }
	169	+ return ge::GRAPH_SUCCESS;
	170	+}
	171	+
	172	+ge::graphStatus SwigluGroupTiling::CheckOutputInfo(ge::DataType xDtype)
	173	+{
	174	+ auto yDesc = context_->GetOutputDesc(OUTPUT_INDEX_Y);
	175	+ OP_CHECK_NULL_WITH_CONTEXT(context_, yDesc);
	176	+ auto yDtype = yDesc->GetDataType();
	177	+ OP_CHECK_IF((yDtype != xDtype),
	178	+ OP_LOGE(context_->GetNodeName(), "output y dtype should be same as input x, got y dtype %d, x dtype %d.",
	179	+ static_cast<int>(yDtype), static_cast<int>(xDtype)),
	180	+ return ge::GRAPH_FAILED);
	181	+ return ge::GRAPH_SUCCESS;
	182	+}
	183	+
	184	+ge::graphStatus SwigluGroupTiling::GetShapeAttrsInfoInner()
	185	+{
	186	+ // (b, s, hc_mix)
	187	+ auto shapeX = context_->GetInputShape(INPUT_INDEX_X);
	188	+ OP_CHECK_NULL_WITH_CONTEXT(context_, shapeX);
	189	+
	190	+ auto xStorageShape = shapeX->GetStorageShape();
	191	+ auto xDesc = context_->GetInputDesc(INPUT_INDEX_X);
	192	+ OP_CHECK_NULL_WITH_CONTEXT(context_, xDesc);
	193	+ auto xDtype = xDesc->GetDataType();
	194	+ OP_CHECK_IF((xDtype != ge::DT_FLOAT16 && xDtype != ge::DT_BF16 && xDtype != ge::DT_FLOAT),
	195	+ OP_LOGE(context_->GetNodeName(), "input x dtype only support FLOAT16, BFLOAT16 or FLOAT, got %d.",
	196	+ static_cast<int>(xDtype)),
	197	+ return ge::GRAPH_FAILED);
	198	+ xElemBytes_ = (xDtype == ge::DT_FLOAT) ? B32_BYTES : B16_BYTES;
	199	+ auto xDimNum = xStorageShape.GetDimNum();
	200	+ OP_CHECK_IF((xDimNum == 0),
	201	+ OP_LOGE(context_->GetNodeName(), "input x dim num should be greater than 0."),
	202	+ return ge::GRAPH_FAILED);
	203	+ bs_ = 1;
	204	+ for (size_t i = 0; i < xDimNum - 1; i++) {
	205	+ bs_ = bs_ * xStorageShape.GetDim(i);
	206	+ }
	207	+ // Empty tensor is not supported, so every dim must be positive. The last dim is checked below;
	208	+ // bs_ is the product of the remaining dims, which is positive only when none of them is 0.
	209	+ OP_CHECK_IF((bs_ <= 0),
	210	+ OP_LOGE(context_->GetNodeName(),
	211	+ "input x is empty tensor, which is not supported, got outer dim product %ld.", bs_),
	212	+ return ge::GRAPH_FAILED);
	213	+ d_ = xStorageShape.GetDim(xDimNum - 1);
	214	+ OP_CHECK_IF((d_ <= 0 \|\| d_ % NUM_TWO != 0),
	215	+ OP_LOGE(context_->GetNodeName(),
	216	+ "input x last dim should be positive and divisible by %ld, got %ld.",
	217	+ NUM_TWO, d_),
	218	+ return ge::GRAPH_FAILED);
	219	+
	220	+ if (CheckWeightInfo() == ge::GRAPH_FAILED \|\| CheckGroupIndexInfo() == ge::GRAPH_FAILED) {
	221	+ return ge::GRAPH_FAILED;
	222	+ }
	223	+
	224	+ // Get Attrs
	225	+ if (GetAttr() == ge::GRAPH_FAILED) {
	226	+ OP_LOGE(context_->GetNodeName(), "Get attr failed.");
	227	+ return ge::GRAPH_FAILED;
	228	+ }
	229	+
	230	+ if (CheckOutputInfo(xDtype) == ge::GRAPH_FAILED) {
	231	+ return ge::GRAPH_FAILED;
	232	+ }
	233	+
	234	+ splitD_ = d_ / 2;
	235	+ return ge::GRAPH_SUCCESS;
	236	+}
	237	+
	238	+ge::graphStatus SwigluGroupTiling::CalcGroupIndexTiling()
	239	+{
	240	+ if (!hasGroupIndex_) {
	241	+ return ge::GRAPH_SUCCESS;
	242	+ }
	243	+
	244	+ gFactor_ = g_;
	245	+ int64_t groupIndexSize = RoundUp(gFactor_, BLOCK_SIZE / sizeof(int64_t)) * DOUBLE_BUFFER * sizeof(int64_t);
	246	+ int64_t groupIndexSumSize = BLOCK_SIZE;
	247	+ if (groupIndexSize + groupIndexSumSize <= static_cast<int64_t>(ubSize_)) {
	248	+ gLoop_ = 1;
	249	+ tailGFactor_ = gFactor_;
	250	+ return ge::GRAPH_SUCCESS;
	251	+ }
	252	+
	253	+ int64_t base = 2;
	254	+ int64_t maxBase = std::max(g_, base);
	255	+ while (base <= maxBase) {
	256	+ gFactor_ = CeilDiv(g_, base);
	257	+ groupIndexSize = RoundUp(gFactor_, BLOCK_SIZE / sizeof(int64_t)) * DOUBLE_BUFFER * sizeof(int64_t);
	258	+ if (groupIndexSize + groupIndexSumSize < static_cast<int64_t>(ubSize_)) {
	259	+ break;
	260	+ }
	261	+ base++;
	262	+ }
	263	+ gFactor_ = std::max(gFactor_, static_cast<int64_t>(1));
	264	+ if (gFactor_ > static_cast<int64_t>(CACHE_LINE_SIZE / sizeof(int64_t))) {
	265	+ gFactor_ = DownAlign(gFactor_, CACHE_LINE_SIZE / sizeof(int64_t));
	266	+ gFactor_ = std::max(gFactor_, static_cast<int64_t>(1));
	267	+ }
	268	+ gLoop_ = CeilDiv(g_, gFactor_);
	269	+ tailGFactor_ = g_ % gFactor_ == 0 ? gFactor_ : g_ % gFactor_;
	atomgit-bot6月23日 🔵 Low Priority 在 `CheckGroupIndexInfo`（tiling.cpp 第146-166行）中，`g_` 初始化为1后乘以 `groupIndexStorageShape` 的各维度。若某个维度为0，`g_` 变为0。随后 `CalcGroupIndexTiling`（第227-260行）计算 `gLoop_ = CeilDiv(g_, gFactor_)`，当 `g_==0` 且 `gFactor_>=1` 时，`gLoop_` 为0。在 kernel 端 `ProcessGroupIndexTiling`（base.h 第377-389行）中，`for (int64_t idx = 0; idx < 0; idx++)` 循环体不执行，`groupSumLocal` 从未被 `VFProcessGroupIndex` 写入，但第393行仍然读取 `groupSumLocal.GetValue(0)` 来决定 `realBs`，导致未初始化数据被用于后续 block tiling 计算——这可能产生错误的 tile 划分和越界访存。建议：在 `CheckGroupIndexInfo` 中添加 `g_ <= 0` 的校验并返回 GRAPH_FAILED，或在 `CalcGroupIndexTiling` 中若 `g_ <= 0` 则设置 `hasGroupIndex_ = false` 并提前返回。这样确保 kernel 端不会在 `gLoop_ == 0` 时读取未初始化的 `groupSumLocal`。 atomgit-bot6月24日 🟡 Medium Priority 在 `CheckGroupIndexInfo`（tiling.cpp 第146-166行）中，`g_` 初始化为1后乘以 `groupIndexStorageShape` 的各维度。若某个维度为0，`g_` 变为0。随后 `CalcGroupIndexTiling`（第227-260行）计算 `gLoop_ = CeilDiv(0, gFactor_)`：第一次尝试时 `gFactor_ = g_ = 0`，`gLoop_ = CeilDiv(0, 0) = 0`。这意味着 kernel 中 `ProcessGroupIndexTiling` 的 `for (idx = 0; idx < 0; ...)` 循环不执行，`groupSumLocal` 保持未初始化状态，随后 `groupSumLocal.GetValue(0)` 读取未初始化数据。此外 `tBufPool.InitBuffer(groupIndexQue, ...)` 以 `gFactor_=0` 分配0字节 buffer 也可能导致问题。建议：在 `CheckGroupIndexInfo` 中添加 `g_ <= 0` 的校验并返回 `GRAPH_FAILED`，或在 `CalcGroupIndexTiling` 开头若 `g_ <= 0` 则设置 `hasGroupIndex_ = false` 提前返回。
	270	+ return ge::GRAPH_SUCCESS;
	271	+}
	272	+
	273	+void SwigluGroupTiling::InitCoreTiling()
	274	+{
	275	+ rowOfFormerBlock_ = CeilDiv(bs_, static_cast<int64_t>(coreNum_));
	276	+ usedCoreNums_ = std::min(CeilDiv(bs_, rowOfFormerBlock_), static_cast<int64_t>(coreNum_));
	277	+ rowOfTailBlock_ = bs_ - (usedCoreNums_ - 1) * rowOfFormerBlock_;
	278	+}
	279	+
	280	+void SwigluGroupTiling::SetFullDTiling()
	281	+{
	282	+ dLoop_ = 1;
	283	+ dFactor_ = splitD_;
	284	+ tailDFactor_ = dFactor_;
	285	+}
	286	+
	287	+void SwigluGroupTiling::SetSplitDTiling()
	288	+{
	289	+ dLoop_ = CeilDiv(splitD_, dFactor_);
	290	+ tailDFactor_ = splitD_ % dFactor_ == 0 ? dFactor_ : splitD_ % dFactor_;
	291	+}
	292	+
	293	+void SwigluGroupTiling::SetRowLoopTiling()
	294	+{
	295	+ rowLoopOfFormerBlock_ = CeilDiv(rowOfFormerBlock_, rowFactor_);
	296	+ rowLoopOfTailBlock_ = CeilDiv(rowOfTailBlock_, rowFactor_);
	297	+ tailRowFactorOfFormerBlock_ = rowOfFormerBlock_ % rowFactor_ == 0 ? rowFactor_ : rowOfFormerBlock_ % rowFactor_;
	298	+ tailRowFactorOfTailBlock_ = rowOfTailBlock_ % rowFactor_ == 0 ? rowFactor_ : rowOfTailBlock_ % rowFactor_;
	299	+}
	300	+
	301	+int64_t SwigluGroupTiling::AddWeightSize(int64_t totalSize, int64_t rowFactor) const
	302	+{
	303	+ return hasWeight_ ? totalSize + RoundUp(rowFactor, B32_ALIGN_NUM) * B32_BYTES * DOUBLE_BUFFER : totalSize;
	304	+}
	305	+
	306	+int64_t SwigluGroupTiling::CalcTotalSize(int64_t rowFactor, int64_t dFactor) const
	307	+{
	308	+ // x0, x1 and y share the x dtype (2 bytes for fp16/bf16, 4 bytes for float32).
	309	+ int64_t alignNum = BLOCK_SIZE / xElemBytes_;
	310	+ int64_t bufBytes = rowFactor * RoundUp(dFactor, alignNum) * xElemBytes_ * DOUBLE_BUFFER;
	311	+ int64_t totalSize = bufBytes + bufBytes + bufBytes; // x0 + x1 + y
	312	+ return AddWeightSize(totalSize, rowFactor);
	313	+}
	314	+
	315	+void SwigluGroupTiling::CalcDAndRowFactorTiling(int64_t rowOnceLoop, int64_t dStep)
	316	+{
	317	+ rowFactor_ = rowOnceLoop;
	318	+ if (CalcTotalSize(rowOnceLoop, splitD_) <= static_cast<int64_t>(ubSize_)) {
	319	+ SetFullDTiling();
	320	+ } else {
	321	+ int64_t base = 1;
	322	+ int64_t maxBase = CeilDiv(splitD_, dStep);
	323	+ dFactor_ = dStep;
	324	+ while (base <= maxBase) {
	325	+ int64_t tryDFactor = base * dStep;
	326	+ if (tryDFactor > splitD_) {
	327	+ tryDFactor = splitD_;
	328	+ }
	329	+ if (CalcTotalSize(rowOnceLoop, tryDFactor) > static_cast<int64_t>(ubSize_)) {
	330	+ break;
	331	+ }
	332	+ dFactor_ = tryDFactor;
	333	+ if (tryDFactor == splitD_) {
	334	+ break;
	335	+ }
	336	+ base++;
	337	+ }
	338	+ SetSplitDTiling();
	339	+ }
	340	+
	341	+ if (dFactor_ == splitD_) {
	342	+ int64_t lo = 1;
	343	+ int64_t hi = rowOfFormerBlock_;
	344	+ while (lo < hi) {
	345	+ int64_t mid = lo + (hi - lo + 1) / 2;
	346	+ if (CalcTotalSize(mid, dFactor_) <= static_cast<int64_t>(ubSize_)) {
	347	+ lo = mid;
	348	+ } else {
	349	+ hi = mid - 1;
	350	+ }
	351	+ }
	352	+ rowFactor_ = lo;
	353	+ }
	354	+
	355	+ SetRowLoopTiling();
	356	+}
	357	+
	358	+ge::graphStatus SwigluGroupTiling::CalcOpTiling()
	359	+{
	360	+ ge::graphStatus status = CalcGroupIndexTiling();
	361	+ if (status == ge::GRAPH_FAILED) {
	362	+ return status;
	363	+ }
	364	+ InitCoreTiling();
	365	+ int64_t rowOnceLoop = std::min(rowOfFormerBlock_, static_cast<int64_t>(1));
	366	+ int64_t dChunk = ASCEND950_CACHE_LINE_BYTES / xElemBytes_; // elements per cacheline
	367	+ CalcDAndRowFactorTiling(rowOnceLoop, dChunk);
	368	+ SetTilingData();
	369	+ return ge::GRAPH_SUCCESS;
	370	+}
	371	+
	372	+void SwigluGroupTiling::SetTilingData()
	373	+{
	374	+ tilingData_.set_bs(bs_);
	375	+ tilingData_.set_d(d_);
	376	+ tilingData_.set_splitD(splitD_);
	377	+ tilingData_.set_rowOfFormerBlock(rowOfFormerBlock_);
	378	+ tilingData_.set_rowOfTailBlock(rowOfTailBlock_);
	379	+ tilingData_.set_rowLoopOfFormerBlock(rowLoopOfFormerBlock_);
	380	+ tilingData_.set_rowLoopOfTailBlock(rowLoopOfTailBlock_);
	381	+ tilingData_.set_rowFactor(rowFactor_);
	382	+ tilingData_.set_tailRowFactorOfFormerBlock(tailRowFactorOfFormerBlock_);
	383	+ tilingData_.set_tailRowFactorOfTailBlock(tailRowFactorOfTailBlock_);
	384	+ tilingData_.set_dLoop(dLoop_);
	385	+ tilingData_.set_dFactor(dFactor_);
	386	+ tilingData_.set_tailDFactor(tailDFactor_);
	387	+ tilingData_.set_clampLimit(clampLimit_);
	388	+ tilingData_.set_g(g_);
	389	+ tilingData_.set_ubSize(ubSize_);
	390	+ tilingData_.set_gLoop(gLoop_);
	391	+ tilingData_.set_gFactor(gFactor_);
	392	+ tilingData_.set_tailGFactor(tailGFactor_);
	393	+ tilingData_.set_coreNum(coreNum_);
	394	+ tilingData_.set_hasClampLimit(hasClampLimit_);
	395	+}
	396	+
	397	+void SwigluGroupTiling::SetTilingKey()
	398	+{
	399	+ tilingKey_ = SWIGLU_GROUP_TILING_KEY;
	400	+ context_->SetTilingKey(tilingKey_);
	401	+}
	402	+
	403	+ge::graphStatus SwigluGroupTiling::DoOpTiling()
	404	+{
	405	+ if (GetPlatformInfo() == ge::GRAPH_FAILED) {
	406	+ return ge::GRAPH_FAILED;
	407	+ }
	408	+
	409	+ if (GetShapeAttrsInfoInner() == ge::GRAPH_FAILED) {
	410	+ return ge::GRAPH_FAILED;
	411	+ }
	412	+
	413	+ if (CalcOpTiling() == ge::GRAPH_FAILED) {
	414	+ return ge::GRAPH_FAILED;
	415	+ }
	416	+
	417	+ if (GetWorkspaceSize() == ge::GRAPH_FAILED) {
	418	+ return ge::GRAPH_FAILED;
	419	+ }
	420	+
	421	+ if (PostTiling() == ge::GRAPH_FAILED) {
	422	+ return ge::GRAPH_FAILED;
	423	+ }
	424	+ SetTilingKey();
	425	+
	426	+ return ge::GRAPH_SUCCESS;
	427	+}
	428	+
	429	+ge::graphStatus SwigluGroupTiling::GetWorkspaceSize()
	430	+{
	431	+ workspaceSize_ = WORKSPACE_SIZE;
	432	+ return ge::GRAPH_SUCCESS;
	433	+}
	434	+
	435	+ge::graphStatus SwigluGroupTiling::PostTiling()
	436	+{
	437	+ if (hasGroupIndex_) {
	438	+ context_->SetBlockDim(coreNum_);
	439	+ } else {
	440	+ context_->SetBlockDim(usedCoreNums_);
	441	+ }
	442	+ size_t* workspaces = context_->GetWorkspaceSizes(1);
	443	+ workspaces[0] = workspaceSize_;
	444	+ tilingData_.SaveToBuffer(context_->GetRawTilingData()->GetData(), context_->GetRawTilingData()->GetCapacity());
	445	+ context_->GetRawTilingData()->SetDataSize(tilingData_.GetDataSize());
	446	+ return ge::GRAPH_SUCCESS;
	447	+}
	448	+
	449	+ge::graphStatus TilingPrepareForSwigluGroup(gert::TilingParseContext *context)
	450	+{
	451	+ (void)context;
	452	+ return ge::GRAPH_SUCCESS;
	453	+}
	454	+
	455	+ge::graphStatus TilingForSwigluGroup(gert::TilingContext *context)
	456	+{
	457	+ OP_CHECK_IF(context == nullptr, OP_LOGE("SwigluGroup", "Tiling context is null"),
	458	+ return ge::GRAPH_FAILED);
	459	+ SwigluGroupTiling swigluGroupTiling(context);
	460	+ return swigluGroupTiling.DoOpTiling();
	461	+}
	462	+
	463	+IMPL_OP_OPTILING(SwigluGroup)
	464	+ .Tiling(TilingForSwigluGroup)
	465	+ .TilingParse<SwigluGroupCompileInfo>(TilingPrepareForSwigluGroup);
	466	+} // namespace optiling

Aactivation/swiglu_group/op_host/arch35/swiglu_group_tiling.h+129-0

@@ -0,0 +1,129 @@
	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 swiglu_group_tiling.h
	13	+ * \brief
	14	+ */
	15	+
	16	+#ifndef SWIGLU_GROUP_TILING_H
	17	+#define SWIGLU_GROUP_TILING_H
	18	+
	19	+
	20	+#include <vector>
	21	+#include <iostream>
	22	+#include "register/op_impl_registry.h"
	23	+#include "platform/platform_infos_def.h"
	24	+#include "exe_graph/runtime/tiling_context.h"
	25	+#include "tiling/platform/platform_ascendc.h"
	26	+#include "register/op_def_registry.h"
	27	+#include "register/tilingdata_base.h"
	28	+#include "tiling/tiling_api.h"
	29	+#include "log/log.h"
	30	+#include "platform/platform_info.h"
	31	+
	32	+namespace optiling {
	33	+// ----------算子TilingData定义----------
	34	+BEGIN_TILING_DATA_DEF(SwigluGroupTilingData)
	35	+TILING_DATA_FIELD_DEF(int64_t, bs);
	36	+TILING_DATA_FIELD_DEF(int64_t, d);
	37	+TILING_DATA_FIELD_DEF(int64_t, splitD);
	38	+TILING_DATA_FIELD_DEF(int64_t, rowOfFormerBlock);
	39	+TILING_DATA_FIELD_DEF(int64_t, rowOfTailBlock);
	40	+TILING_DATA_FIELD_DEF(int64_t, rowLoopOfFormerBlock);
	41	+TILING_DATA_FIELD_DEF(int64_t, rowLoopOfTailBlock);
	42	+TILING_DATA_FIELD_DEF(int64_t, rowFactor);
	43	+TILING_DATA_FIELD_DEF(int64_t, tailRowFactorOfFormerBlock);
	atomgit-bot6月24日 🔵 Low Priority `swiglu_group_tiling.h` 第43-44行的 `TILING_DATA_FIELD_DEF` 宏字段名 `tailRowFactorOfFormerBlock` 拼写为 "Former"（应为 "Former"），而同一文件其余所有引用均使用 "Former"（如第42行 `rowOfFormerBlock`、第108行 `rowOfFormerBlock_`、第112行 `tailRowFactorOfFormerBlock_`）。虽然 tiling data 的字段名不影响运行时行为（由宏生成的 setter/getter 基于字面量字符串），但拼写不一致会在代码搜索、文档和代码审查中造成混淆。C 代码中的拼写错误在声明处属于需报告的问题。建议：这是一个可选改进项：将 `tailRowFactorOfFormerBlock` 改为 `tailRowFactorOfFormerBlock` 以与其余代码保持一致。如框架强制要求字段名与代码中使用的名称匹配（如序列化依赖），则此为必须修复项。
	44	+TILING_DATA_FIELD_DEF(int64_t, tailRowFactorOfTailBlock);
	45	+TILING_DATA_FIELD_DEF(int64_t, dLoop);
	46	+TILING_DATA_FIELD_DEF(int64_t, dFactor);
	47	+TILING_DATA_FIELD_DEF(int64_t, tailDFactor);
	48	+TILING_DATA_FIELD_DEF(float, clampLimit);
	49	+TILING_DATA_FIELD_DEF(int64_t, hasClampLimit);
	50	+TILING_DATA_FIELD_DEF(int64_t, g);
	51	+TILING_DATA_FIELD_DEF(int64_t, ubSize);
	52	+TILING_DATA_FIELD_DEF(int64_t, gLoop);
	53	+TILING_DATA_FIELD_DEF(int64_t, gFactor);
	54	+TILING_DATA_FIELD_DEF(int64_t, tailGFactor);
	55	+TILING_DATA_FIELD_DEF(int64_t, coreNum);
	56	+END_TILING_DATA_DEF;
	57	+
	58	+REGISTER_TILING_DATA_CLASS(SwigluGroup, SwigluGroupTilingData)
	59	+
	60	+// ----------算子CompileInfo定义----------
	61	+struct SwigluGroupCompileInfo {
	62	+ uint64_t coreNum = 0;
	63	+ uint64_t ubSize = 0;
	64	+};
	65	+
	66	+// ----------算子Tiling入参信息解析及check类----------
	67	+class SwigluGroupTiling {
	68	+public:
	69	+ explicit SwigluGroupTiling(gert::TilingContext* tilingContext) : context_(tilingContext)
	70	+ {
	71	+ }
	72	+ ~SwigluGroupTiling() = default;
	73	+
	74	+ ge::graphStatus GetPlatformInfo();
	75	+ ge::graphStatus DoOpTiling();
	76	+ ge::graphStatus GetWorkspaceSize();
	77	+ ge::graphStatus PostTiling();
	78	+ ge::graphStatus GetAttr();
	79	+ ge::graphStatus GetShapeAttrsInfoInner();
	80	+ ge::graphStatus CalcOpTiling();
	81	+ ge::graphStatus CalcGroupIndexTiling();
	82	+ void SetTilingData();
	83	+ void SetTilingKey();
	84	+private:
	85	+ ge::graphStatus GetClampLimitAttr(const gert::RuntimeAttrs* attrs);
	86	+ ge::graphStatus CheckWeightInfo();
	87	+ ge::graphStatus CheckGroupIndexInfo();
	88	+ ge::graphStatus CheckOutputInfo(ge::DataType xDtype);
	89	+ void InitCoreTiling();
	90	+ void CalcDAndRowFactorTiling(int64_t rowOnceLoop, int64_t dStep);
	91	+ void SetFullDTiling();
	92	+ void SetSplitDTiling();
	93	+ void SetRowLoopTiling();
	94	+ int64_t AddWeightSize(int64_t totalSize, int64_t rowFactor) const;
	95	+ int64_t CalcTotalSize(int64_t rowFactor, int64_t dFactor) const;
	96	+
	97	+ gert::TilingContext *context_ = nullptr;
	98	+ uint64_t tilingKey_ = 0;
	99	+ SwigluGroupTilingData tilingData_;
	100	+ uint64_t coreNum_ = 0;
	101	+ uint64_t workspaceSize_ = 0;
	102	+ uint64_t usedCoreNums_ = 0;
	103	+ uint64_t ubSize_ = 0;
	104	+ int64_t bs_ = 0;
	105	+ int64_t d_ = 0;
	106	+ int64_t splitD_ = 0;
	107	+ int64_t xElemBytes_ = 2;
	108	+ int64_t rowOfFormerBlock_ = 0;
	109	+ int64_t rowOfTailBlock_ = 0;
	110	+ int64_t rowLoopOfFormerBlock_ = 0;
	111	+ int64_t rowLoopOfTailBlock_ = 0;
	112	+ int64_t rowFactor_ = 0;
	113	+ int64_t tailRowFactorOfFormerBlock_ = 0;
	114	+ int64_t tailRowFactorOfTailBlock_= 0;
	115	+ int64_t dLoop_ = 0;
	116	+ int64_t dFactor_ = 0;
	117	+ int64_t tailDFactor_ = 0;
	118	+ float clampLimit_ = 0.0f;
	119	+ int64_t hasClampLimit_ = 0;
	120	+ bool hasWeight_ = false;
	121	+ int64_t g_ = 0;
	122	+ int64_t gLoop_ = 0;
	123	+ int64_t gFactor_ = 0;
	124	+ int64_t tailGFactor_ = 0;
	125	+ bool hasGroupIndex_ = false;
	126	+};
	127	+
	128	+} // namespace optiling
	129	+#endif // SWIGLU_GROUP_TILING_H

Aactivation/swiglu_group/op_host/config/ascend950/swiglu_group_binary.json+47-0

@@ -0,0 +1,47 @@
	1	+{
	2	+ "op_type": "SwigluGroup",
	3	+ "op_list": [
	4	+ {
	5	+ "bin_filename": "SwigluGroup_A950_f16_q7m4x9p2r6b1",
	6	+ "inputs": [
	7	+ { "name": "x", "index": 0, "dtype": "float16", "format": "ND", "paramType": "required", "shape": [-2], "format_match_mode": "FormatAgnostic" },
	8	+ { "name": "weight", "index": 1, "dtype": "float32", "format": "ND", "paramType": "optional", "shape": [-2], "format_match_mode": "FormatAgnostic" },
	9	+ { "name": "group_index", "index": 2, "dtype": "int64", "format": "ND", "paramType": "optional", "shape": [-1], "format_match_mode": "FormatAgnostic" }
	10	+ ],
	11	+ "outputs": [
	12	+ { "name": "y", "index": 0, "dtype": "float16", "format": "ND", "paramType": "required", "shape": [-2], "format_match_mode": "FormatAgnostic" }
	13	+ ],
	14	+ "attrs": [
	15	+ { "name": "clamp_limit", "dtype": "float", "value": null }
	16	+ ]
	17	+ },
	18	+ {
	19	+ "bin_filename": "SwigluGroup_A950_bf16_s8n2v5c7t3w0",
	20	+ "inputs": [
	21	+ { "name": "x", "index": 0, "dtype": "bfloat16", "format": "ND", "paramType": "required", "shape": [-2], "format_match_mode": "FormatAgnostic" },
	22	+ { "name": "weight", "index": 1, "dtype": "float32", "format": "ND", "paramType": "optional", "shape": [-2], "format_match_mode": "FormatAgnostic" },
	23	+ { "name": "group_index", "index": 2, "dtype": "int64", "format": "ND", "paramType": "optional", "shape": [-1], "format_match_mode": "FormatAgnostic" }
	24	+ ],
	25	+ "outputs": [
	26	+ { "name": "y", "index": 0, "dtype": "bfloat16", "format": "ND", "paramType": "required", "shape": [-2], "format_match_mode": "FormatAgnostic" }
	27	+ ],
	28	+ "attrs": [
	29	+ { "name": "clamp_limit", "dtype": "float", "value": null }
	30	+ ]
	31	+ },
	32	+ {
	33	+ "bin_filename": "SwigluGroup_A950_f32_t6w9b2x5q1n4",
	34	+ "inputs": [
	35	+ { "name": "x", "index": 0, "dtype": "float32", "format": "ND", "paramType": "required", "shape": [-2], "format_match_mode": "FormatAgnostic" },
	36	+ { "name": "weight", "index": 1, "dtype": "float32", "format": "ND", "paramType": "optional", "shape": [-2], "format_match_mode": "FormatAgnostic" },
	37	+ { "name": "group_index", "index": 2, "dtype": "int64", "format": "ND", "paramType": "optional", "shape": [-1], "format_match_mode": "FormatAgnostic" }
	38	+ ],
	39	+ "outputs": [
	40	+ { "name": "y", "index": 0, "dtype": "float32", "format": "ND", "paramType": "required", "shape": [-2], "format_match_mode": "FormatAgnostic" }
	41	+ ],
	42	+ "attrs": [
	43	+ { "name": "clamp_limit", "dtype": "float", "value": null }
	44	+ ]
	45	+ }
	46	+ ]
	47	+}

Aactivation/swiglu_group/op_host/swiglu_group_def.cpp+61-0

@@ -0,0 +1,61 @@
	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 swiglu_group_def.cpp
	13	+ * \brief
	14	+ */
	15	+#include "register/op_def_registry.h"
	16	+
	17	+namespace ops {
	18	+namespace {
	19	+constexpr float DEFAULT_CLAMP_LIMIT = -1.0f;
	20	+} // namespace
	21	+
	22	+class SwigluGroup : public OpDef {
	23	+public:
	24	+ explicit SwigluGroup(const char *name) : OpDef(name)
	25	+ {
	26	+ this->Input("x")
	27	+ .ParamType(REQUIRED)
	28	+ .DataType({ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT})
	29	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	30	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
	31	+ this->Input("weight")
	32	+ .ParamType(OPTIONAL)
	33	+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT})
	34	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	35	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
	36	+ this->Input("group_index")
	37	+ .ParamType(OPTIONAL)
	38	+ .DataType({ge::DT_INT64, ge::DT_INT64, ge::DT_INT64})
	39	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	40	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
	41	+ this->Output("y")
	42	+ .ParamType(REQUIRED)
	43	+ .DataType({ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT})
	44	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	45	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
	46	+ this->Attr("clamp_limit").AttrType(OPTIONAL).Float(DEFAULT_CLAMP_LIMIT);
	47	+
	48	+ OpAICoreConfig aicoreConfig;
	49	+ aicoreConfig.DynamicCompileStaticFlag(true)
	50	+ .DynamicFormatFlag(false)
	51	+ .DynamicRankSupportFlag(true)
	52	+ .DynamicShapeSupportFlag(true)
	53	+ .NeedCheckSupportFlag(false)
	54	+ .PrecisionReduceFlag(true)
	55	+ .ExtendCfgInfo("opFile.value", "swiglu_group");
	56	+ this->AICore().AddConfig("ascend950", aicoreConfig);
	57	+ }
	58	+};
	59	+
	60	+OP_ADD(SwigluGroup);
	61	+} // namespace ops

Aactivation/swiglu_group/op_host/swiglu_group_infershape.cpp+77-0

@@ -0,0 +1,77 @@
	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 swiglu_group_infershape.cpp
	13	+ * \brief Shape and dtype inference for SwigluGroup.
	14	+ */
	15	+
	16	+#include "graph/utils/type_utils.h"
	17	+#include "log/log.h"
	18	+#include "register/op_impl_registry.h"
	19	+#include "util/shape_util.h"
	20	+
	21	+using namespace ge;
	22	+namespace ops {
	23	+namespace {
	24	+constexpr size_t INPUT_IDX_X = 0;
	25	+constexpr size_t OUTPUT_IDX_Y = 0;
	26	+constexpr int64_t NUM_TWO = 2;
	27	+} // namespace
	28	+
	29	+graphStatus InferShape4SwigluGroup(gert::InferShapeContext* context)
	30	+{
	31	+ OP_LOGD(context->GetNodeName(), "Begin to do InferShape4SwigluGroup.");
	32	+ const gert::Shape* xShape = context->GetInputShape(INPUT_IDX_X);
	33	+ OP_CHECK_NULL_WITH_CONTEXT(context, xShape);
	34	+ gert::Shape* yShape = context->GetOutputShape(OUTPUT_IDX_Y);
	35	+ OP_CHECK_NULL_WITH_CONTEXT(context, yShape);
	36	+
	37	+ if (Ops::Base::IsUnknownRank(*xShape)) {
	38	+ Ops::Base::SetUnknownRank(*yShape);
	39	+ return ge::GRAPH_SUCCESS;
	40	+ }
	41	+
	42	+ int64_t xRank = static_cast<int64_t>(xShape->GetDimNum());
	43	+ OP_CHECK_IF(xRank < 1,
	44	+ OP_LOGE(context->GetNodeName(), "The rank of x should be greater than 0, but is %ld.", xRank),
	45	+ return ge::GRAPH_FAILED);
	46	+ int64_t splitDim = xRank - 1;
	47	+
	48	+ yShape = xShape;
	49	+ if (xShape->GetDim(splitDim) == -1) {
	50	+ return ge::GRAPH_SUCCESS;
	51	+ }
	52	+
	53	+ OP_CHECK_IF(xShape->GetDim(splitDim) < 0 \|\| xShape->GetDim(splitDim) % NUM_TWO != 0,
	54	+ OP_LOGE(context->GetNodeName(),
	55	+ "The last dimension of x should be non-negative and divisible by 2, but got %ld.",
	56	+ xShape->GetDim(splitDim)),
	57	+ return ge::GRAPH_FAILED);
	58	+
	59	+ yShape->SetDim(splitDim, xShape->GetDim(splitDim) / NUM_TWO);
	60	+
	61	+ OP_LOGD(context->GetNodeName(), "End to do InferShape4SwigluGroup.");
	62	+ return ge::GRAPH_SUCCESS;
	63	+}
	64	+
	65	+graphStatus InferDtype4SwigluGroup(gert::InferDataTypeContext* context)
	66	+{
	67	+ OP_LOGD(context->GetNodeName(), "Begin to do InferDtype4SwigluGroup.");
	68	+ auto xDtype = context->GetInputDataType(INPUT_IDX_X);
	69	+ context->SetOutputDataType(OUTPUT_IDX_Y, xDtype);
	70	+ OP_LOGD(context->GetNodeName(), "End to do InferDtype4SwigluGroup.");
	71	+ return GRAPH_SUCCESS;
	72	+}
	73	+
	74	+IMPL_OP_INFERSHAPE(SwigluGroup)
	75	+ .InferShape(InferShape4SwigluGroup)
	76	+ .InferDataType(InferDtype4SwigluGroup);
	77	+} // namespace ops

Aactivation/swiglu_group/op_kernel/CMakeLists.txt+16-0

@@ -0,0 +1,16 @@
	1	+# ----------------------------------------------------------------------------
	2	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+# CANN Open Software License Agreement Version 2.0 (the "License").
	5	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+# See LICENSE in the root of the software repository for the full text of the License.
	9	+# ----------------------------------------------------------------------------
	10	+
	11	+add_kernel_sources(
	12	+ KERNEL_SRC swiglu_group.cpp
	13	+ COMPUTE_UNITS ascend950
	14	+ AUTO_SYNC false
	15	+ OPTIONS -Wno-deprecated-declarations -Werror
	16	+)

Aactivation/swiglu_group/op_kernel/arch35/swiglu_group_base.h+403-0

@@ -0,0 +1,403 @@
	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 swiglu_group_base.h
	13	+ * \brief
	14	+ */
	15	+
	16	+#ifndef SWIGLU_GROUP_BASE_H
	17	+#define SWIGLU_GROUP_BASE_H
	18	+
	19	+#include "kernel_operator.h"
	20	+
	21	+namespace SwigluGroup {
	22	+using namespace AscendC;
	23	+using namespace AscendC::Reg;
	24	+using AscendC::Reg::MaskReg;
	25	+using AscendC::Reg::RegTensor;
	26	+constexpr int32_t BLOCK_SIZE = 32;
	27	+constexpr int32_t DOUBLE_BUFFER_NUM = 2;
	28	+constexpr int32_t VL_FP32 = 64;
	29	+constexpr uint32_t REPEAT_SIZE = 256;
	30	+constexpr uint16_t FOUR_UNFOLD = 4;
	31	+
	32	+__aicore__ inline int32_t CeilDiv(int32_t a, int b)
	33	+{
	34	+ if (b == 0) {
	35	+ return a;
	36	+ }
	37	+ return (a + b - 1) / b;
	38	+}
	39	+
	40	+__aicore__ inline int32_t CeilAlign(int32_t a, int b)
	41	+{
	42	+ return CeilDiv(a, b) * b;
	43	+}
	44	+
	45	+template <typename T>
	46	+__aicore__ inline int32_t RoundUp(int32_t num)
	47	+{
	48	+ int32_t elemNum = BLOCK_SIZE / sizeof(T);
	49	+ return CeilAlign(num, elemNum);
	50	+}
	51	+
	52	+constexpr AscendC::Reg::CastTrait castTraitB162B32Even = {
	53	+ AscendC::Reg::RegLayout::ZERO,
	54	+ AscendC::Reg::SatMode::UNKNOWN,
	55	+ AscendC::Reg::MaskMergeMode::ZEROING,
	56	+ AscendC::RoundMode::UNKNOWN,
	57	+};
	58	+
	59	+constexpr AscendC::Reg::CastTrait castTraitB322B16Even = {
	60	+ AscendC::Reg::RegLayout::ZERO,
	61	+ AscendC::Reg::SatMode::NO_SAT,
	62	+ AscendC::Reg::MaskMergeMode::ZEROING,
	63	+ AscendC::RoundMode::CAST_RINT,
	64	+};
	65	+
	66	+template <typename T>
	67	+__simd_callee__ inline void LoadInputData(RegTensor<float>& dst, __ubuf__ T* src, MaskReg pregLoop, uint32_t srcOffset)
	68	+{
	69	+ if constexpr (IsSameType<T, float>::value) {
	70	+ LoadAlign(dst, src + srcOffset);
	71	+ } else if constexpr (IsSameType<T, half>::value \|\| IsSameType<T, bfloat16_t>::value) {
	72	+ RegTensor<T> tmp;
	73	+ LoadAlign<T, AscendC::Reg::LoadDist::DIST_UNPACK_B16>(tmp, src + srcOffset);
	74	+ Cast<float, T, castTraitB162B32Even>(dst, tmp, pregLoop);
	75	+ }
	76	+}
	77	+
	78	+template <typename T>
	79	+__simd_callee__ inline void StoreOutputData(
	80	+ __ubuf__ T* dst, RegTensor<float>& src, MaskReg pregLoop, uint32_t dstOffset)
	81	+{
	82	+ if constexpr (IsSameType<T, float>::value) {
	83	+ StoreAlign(dst + dstOffset, src, pregLoop);
	84	+ } else if constexpr (IsSameType<T, half>::value \|\| IsSameType<T, bfloat16_t>::value) {
	85	+ RegTensor<T> tmp;
	86	+ Cast<T, float, castTraitB322B16Even>(tmp, src, pregLoop);
	87	+ StoreAlign<T, AscendC::Reg::StoreDist::DIST_PACK_B32>(dst + dstOffset, tmp, pregLoop);
	88	+ }
	89	+}
	90	+
	91	+__simd_callee__ inline void VFSwiGlu(RegTensor<float>& y, RegTensor<float>& x0, RegTensor<float>& x1,
	92	+ RegTensor<float>& vreg, MaskReg pregLoop)
	93	+{
	94	+ Muls(vreg, x0, static_cast<float>(-1.0f), pregLoop);
	95	+ Exp(vreg, vreg, pregLoop);
	96	+ Adds(vreg, vreg, static_cast<float>(1.0f), pregLoop);
	97	+ Div(vreg, x0, vreg, pregLoop);
	98	+ Mul(y, vreg, x1, pregLoop);
	99	+}
	100	+
	101	+template <typename T, bool hasTopkWeight = false, bool hasClampValue = false, bool singleLoop = false>
	102	+__simd_vf__ inline void VFProcessSwigluVf(__ubuf__ T* yLocalAddr, __ubuf__ T* x0LocalAddr,
	103	+ __ubuf__ T* x1LocalAddr, __ubuf__ float* topkWeightLocalAddr, uint16_t loopCount,
	104	+ uint32_t sregNum, uint32_t curColNumAlign, const uint16_t curRowNum, float clampValue)
	105	+{
	106	+ RegTensor<float> weight;
	107	+ RegTensor<float> x0;
	108	+ RegTensor<float> x1;
	109	+ RegTensor<float> y;
	110	+ RegTensor<float> tmp;
	111	+ MaskReg pregLoop = CreateMask<float>();
	112	+ if constexpr (singleLoop) {
	113	+ uint32_t sreg = sregNum;
	114	+ MaskReg pregFixed = UpdateMask<float>(sreg);
	115	+ for (uint16_t i = 0; i < curRowNum; i++) {
	116	+ if constexpr (hasTopkWeight) {
	117	+ LoadAlign<float, AscendC::Reg::LoadDist::DIST_BRC_B32>(weight, topkWeightLocalAddr + i);
	118	+ }
	119	+ LoadInputData<T>(x0, x0LocalAddr, pregFixed, i * curColNumAlign);
	120	+ LoadInputData<T>(x1, x1LocalAddr, pregFixed, i * curColNumAlign);
	121	+ if constexpr (hasClampValue) {
	122	+ Mins(x0, x0, clampValue, pregFixed);
	123	+ Maxs(x1, x1, -clampValue, pregFixed);
	124	+ Mins(x1, x1, clampValue, pregFixed);
	125	+ }
	126	+ VFSwiGlu(y, x0, x1, tmp, pregFixed);
	127	+ if constexpr (hasTopkWeight) {
	128	+ Mul(y, y, weight, pregFixed);
	129	+ }
	130	+ StoreOutputData<T>(yLocalAddr, y, pregFixed, i * curColNumAlign);
	131	+ }
	132	+ } else {
	133	+ for (uint16_t i = 0; i < curRowNum; i++) {
	134	+ if constexpr (hasTopkWeight) {
	135	+ LoadAlign<float, AscendC::Reg::LoadDist::DIST_BRC_B32>(weight, topkWeightLocalAddr + i);
	136	+ }
	137	+ uint32_t sreg = sregNum;
	138	+ for (uint16_t j = 0; j < loopCount; j++) {
	139	+ pregLoop = UpdateMask<float>(sreg);
	140	+ LoadInputData<T>(x0, x0LocalAddr, pregLoop, j * VL_FP32 + i * curColNumAlign);
	141	+ LoadInputData<T>(x1, x1LocalAddr, pregLoop, j * VL_FP32 + i * curColNumAlign);
	142	+ if constexpr (hasClampValue) {
	143	+ Mins(x0, x0, clampValue, pregLoop);
	144	+ Maxs(x1, x1, -clampValue, pregLoop);
	145	+ Mins(x1, x1, clampValue, pregLoop);
	146	+ }
	147	+ VFSwiGlu(y, x0, x1, tmp, pregLoop);
	148	+ if constexpr (hasTopkWeight) {
	149	+ Mul(y, y, weight, pregLoop);
	150	+ }
	151	+ StoreOutputData<T>(yLocalAddr, y, pregLoop, j * VL_FP32 + i * curColNumAlign);
	152	+ }
	153	+ }
	154	+ }
	155	+}
	156	+
	157	+template <typename T, bool hasTopkWeight = false, bool hasClampValue = false>
	158	+__aicore__ inline void VFProcessSwiglu(
	159	+ const LocalTensor<T>& yLocal, const LocalTensor<T>& x0Local, const LocalTensor<T>& x1Local,
	160	+ const LocalTensor<float>& topkWeightLocal,
	161	+ const uint16_t curRowNum, const uint32_t curColNum, float clampValue)
	162	+{
	163	+ __ubuf__ T* yLocalAddr = (__ubuf__ T*)yLocal.GetPhyAddr();
	164	+ __ubuf__ T* x0LocalAddr = (__ubuf__ T*)x0Local.GetPhyAddr();
	165	+ __ubuf__ T* x1LocalAddr = (__ubuf__ T*)x1Local.GetPhyAddr();
	166	+ __ubuf__ float* topkWeightLocalAddr =
	167	+ hasTopkWeight ? (__ubuf__ float*)topkWeightLocal.GetPhyAddr() : nullptr;
	168	+ uint16_t loopCount = CeilDiv(curColNum, VL_FP32);
	169	+ uint32_t sregNum = curColNum;
	170	+ uint32_t curColNumAlign = RoundUp<T>(curColNum);
	171	+ if (loopCount == 1) {
	172	+ AscendC::VF_CALL<VFProcessSwigluVf<T, hasTopkWeight, hasClampValue, true>>(
	173	+ yLocalAddr, x0LocalAddr, x1LocalAddr, topkWeightLocalAddr, loopCount, sregNum, curColNumAlign, curRowNum,
	174	+ clampValue);
	175	+ } else {
	176	+ AscendC::VF_CALL<VFProcessSwigluVf<T, hasTopkWeight, hasClampValue, false>>(
	177	+ yLocalAddr, x0LocalAddr, x1LocalAddr, topkWeightLocalAddr, loopCount, sregNum, curColNumAlign, curRowNum,
	178	+ clampValue);
	179	+ }
	180	+}
	181	+
	182	+template <typename T>
	183	+__aicore__ inline void SwigluGroupDispatcher(const LocalTensor<T>& yLocal, const LocalTensor<T>& x0Local,
	184	+ const LocalTensor<T>& x1Local, const LocalTensor<float>& topkWeightLocal, float clampValue,
	185	+ const uint16_t curRowNum, const uint32_t curColNum, int32_t maskBit)
	186	+{
	187	+ if (maskBit == 0b00) {
	188	+ VFProcessSwiglu<T, false, false>(yLocal, x0Local, x1Local, topkWeightLocal, curRowNum, curColNum, clampValue);
	189	+ } else if (maskBit == 0b01) {
	190	+ VFProcessSwiglu<T, true, false>(yLocal, x0Local, x1Local, topkWeightLocal, curRowNum, curColNum, clampValue);
	191	+ } else if (maskBit == 0b10) {
	192	+ VFProcessSwiglu<T, false, true>(yLocal, x0Local, x1Local, topkWeightLocal, curRowNum, curColNum, clampValue);
	193	+ } else if (maskBit == 0b11) {
	194	+ VFProcessSwiglu<T, true, true>(yLocal, x0Local, x1Local, topkWeightLocal, curRowNum, curColNum, clampValue);
	195	+ }
	196	+}
	197	+
	198	+template <typename T, bool withUbReduce = false>
	199	+__simd_vf__ inline void VFProcessGroupIndexSmallVf(
	200	+ __ubuf__ T* yLocalAddr, __ubuf__ T* xLocalAddr, uint16_t curColNum, uint16_t vlLen, uint16_t loopCount)
	201	+{
	202	+ RegTensor<T> x;
	203	+ RegTensor<T> sum;
	204	+ MaskReg pregMain = CreateMask<T, AscendC::Reg::MaskPattern::ALL>();
	205	+ MaskReg pregMerge = CreateMask<T, AscendC::Reg::MaskPattern::VL1>();
	206	+ Duplicate(sum, static_cast<T>(0), pregMain);
	207	+ uint32_t sreg = curColNum;
	208	+ MaskReg pregLoop;
	209	+ for (uint16_t i = 0; i < loopCount; i++) {
	210	+ pregLoop = UpdateMask<T>(sreg);
	atomgit-bot6月24日 🟡 Medium Priority `VFProcessGroupIndexSmallVf`（第204-227行）对所有循环迭代使用同一个 `sreg = curColNum` 作为 mask 参数。当 `loopCount >= 2` 且 `curColNum % vlLen != 0` 时，最后一个迭代是部分块（仅 `curColNum % vlLen` 个有效元素），但 `sreg = curColNum > vlLen` 创建满 mask，导致 `Adds(x, x, 0, pregLoop)` 无法将无效 lane 清零，随后 `Add(sum, sum, x, pregMain)` 将 UB 垃圾数据加到求和中。此问题与 `VFProcessGroupIndexLargeVf` 中的 bug 同一根因但位于不同函数。触发条件：`loopCount < 4`（进入 small 路径）且 `loopCount >= 2` 且 `curColNum` 不是 `vlLen` 整数倍。建议：在循环内部为每次迭代计算正确的 mask：对于非最后一次迭代使用满 mask（sreg >= vlLen），最后一次迭代若为部分块则使用 `curColNum % vlLen` 作为 sreg。
	211	+ LoadAlign(x, xLocalAddr + i * vlLen);
	212	+ Adds(x, x, static_cast<T>(0), pregLoop);
	213	+ Add(sum, sum, x, pregMain);
	214	+ }
	215	+ ReduceSum(sum, sum, pregMain);
	216	+ if (withUbReduce) {
	217	+ RegTensor<T> origin;
	218	+ LoadAlign(origin, yLocalAddr);
	219	+ Add(sum, sum, origin, pregMerge);
	220	+ }
	221	+ StoreAlign(yLocalAddr, sum, pregMerge);
	222	+}
	223	+
	224	+template <typename T, bool withUbReduce = false>
	225	+__simd_vf__ inline void VFProcessGroupIndexLargeVf(__ubuf__ T* yLocalAddr, __ubuf__ T* xLocalAddr, uint16_t vlLen,
	226	+ uint16_t fourLoopCount, uint16_t tailLoopNum, uint32_t tailReminder)
	227	+{
	228	+ RegTensor<T> x0;
	229	+ RegTensor<T> x1;
	230	+ RegTensor<T> x2;
	231	+ RegTensor<T> x3;
	232	+ RegTensor<T> sum0;
	233	+ RegTensor<T> sum1;
	234	+ RegTensor<T> sum2;
	235	+ RegTensor<T> sum3;
	236	+ MaskReg pregMain = CreateMask<T, AscendC::Reg::MaskPattern::ALL>();
	237	+ MaskReg pregMerge = CreateMask<T, AscendC::Reg::MaskPattern::VL1>();
	238	+ Duplicate(sum0, static_cast<T>(0), pregMain);
	239	+ Duplicate(sum1, static_cast<T>(0), pregMain);
	240	+ Duplicate(sum2, static_cast<T>(0), pregMain);
	241	+ Duplicate(sum3, static_cast<T>(0), pregMain);
	242	+ MaskReg pregLoop;
	243	+ for (uint16_t i = 0; i < fourLoopCount; i++) {
	244	+ LoadAlign(x0, xLocalAddr + i * FOUR_UNFOLD * vlLen);
	245	+ Add(sum0, sum0, x0, pregMain);
	246	+ LoadAlign(x1, xLocalAddr + (i * FOUR_UNFOLD + 1) * vlLen);
	247	+ Add(sum1, sum1, x1, pregMain);
	248	+ LoadAlign(x2, xLocalAddr + (i * FOUR_UNFOLD + 2) * vlLen);
	249	+ Add(sum2, sum2, x2, pregMain);
	250	+ LoadAlign(x3, xLocalAddr + (i * FOUR_UNFOLD + 3) * vlLen);
	251	+ Add(sum3, sum3, x3, pregMain);
	252	+ }
	253	+ uint32_t sreg = tailReminder;
	254	+ for (uint16_t i = 0; i < tailLoopNum; i++) {
	255	+ pregLoop = UpdateMask<T>(sreg);
	256	+ LoadAlign(x0, xLocalAddr + (fourLoopCount * FOUR_UNFOLD + i) * vlLen);
	257	+ Adds(x0, x0, static_cast<T>(0), pregLoop);
	258	+ Add(sum0, sum0, x0, pregMain);
	259	+ }
	atomgit-bot6月24日 🟡 Medium Priority `VFProcessGroupIndexLargeVf`（第230-275行）对所有 tail 迭代使用同一个 `sreg = tailReminder` 作为 mask 参数。当 `tailLoopNum >= 2` 且存在不完整的尾块（`curColNum % vlLen != 0`）时，最后一个迭代处理的不是完整的 `vlLen` 个元素，但 `sreg = tailReminder`（此时 tailReminder > 该迭代的有效元素数）会创建一个满 mask，导致 `Adds(x0, x0, 0, pregLoop)` 无法将无效 lane 清零，随后 `Add(sum0, sum0, x0, pregMain)` 将 UB 中的垃圾数据加到了 group_index 求和中。这会污染 `groupSumLocal`，进而污染 `realBs`（虽然 `realBs` 最终通过 `min(sum, bs)` 被 clamp，但求和不正确仍会导致分块行为偏差）。触发条件：`loopCount >= 4`（即 `curColNum >= 4 * vlLen`）且 `curColNum` 不是 `vlLen` 的整数倍。例如 `g=97`（`vlLen=32`）时此 bug 被触发。建议：在 tail 循环内部，每次迭代根据已处理的元素个数计算正确的 mask 值：完整块用 >= vlLen 的满 mask，部分块用 `curColNum - alreadyProcessedElements` 作为 sreg。例如：`uint32_t sreg = (i == tailLoopNum - 1 && tailReminder % vlLen != 0) ? (tailReminder % vlLen) : vlLen;` 或类似逻辑。
	260	+ Add(sum0, sum0, sum1, pregMain);
	261	+ Add(sum2, sum2, sum3, pregMain);
	262	+ Add(sum0, sum0, sum2, pregMain);
	263	+ ReduceSum(sum0, sum0, pregMain);
	264	+ if (withUbReduce) {
	265	+ RegTensor<T> origin;
	266	+ LoadAlign(origin, yLocalAddr);
	267	+ Add(sum0, sum0, origin, pregMerge);
	268	+ }
	269	+ StoreAlign(yLocalAddr, sum0, pregMerge);
	270	+}
	271	+
	272	+template <typename T, bool withUbReduce = false>
	273	+__aicore__ inline void VFProcessGroupIndex(const LocalTensor<T>& yLocal, const LocalTensor<T>& xLocal,
	274	+ uint16_t curColNum)
	275	+{
	276	+ __ubuf__ T* yLocalAddr = (__ubuf__ T*)yLocal.GetPhyAddr();
	277	+ __ubuf__ T* xLocalAddr = (__ubuf__ T*)xLocal.GetPhyAddr();
	278	+ uint16_t vlLen = REPEAT_SIZE / sizeof(T);
	279	+ uint16_t loopCount = CeilDiv(curColNum, vlLen);
	280	+ uint16_t fullBlocks = curColNum / vlLen;
	281	+ uint16_t fourLoopCount = fullBlocks / FOUR_UNFOLD;
	282	+ uint16_t tailLoopNum = loopCount - fourLoopCount * FOUR_UNFOLD;
	283	+ uint32_t tailReminder = curColNum - fourLoopCount * vlLen * FOUR_UNFOLD;
	284	+ if (loopCount < FOUR_UNFOLD) {
	285	+ AscendC::VF_CALL<VFProcessGroupIndexSmallVf<T, withUbReduce>>(
	286	+ yLocalAddr, xLocalAddr, curColNum, vlLen, loopCount);
	287	+ } else {
	288	+ AscendC::VF_CALL<VFProcessGroupIndexLargeVf<T, withUbReduce>>(
	289	+ yLocalAddr, xLocalAddr, vlLen, fourLoopCount, tailLoopNum, tailReminder);
	290	+ }
	291	+}
	292	+
	293	+template <typename T>
	294	+__aicore__ inline void CopyIn(
	295	+ const GlobalTensor<T>& inputGm, const LocalTensor<T>& inputTensor, const uint16_t nBurst, const uint32_t copyLen,
	296	+ uint32_t srcStride = 0)
	297	+{
	298	+ DataCopyPadExtParams<T> dataCopyPadExtParams;
	299	+ dataCopyPadExtParams.isPad = false;
	300	+ dataCopyPadExtParams.leftPadding = 0;
	301	+ dataCopyPadExtParams.rightPadding = 0;
	302	+ dataCopyPadExtParams.paddingValue = 0;
	303	+
	304	+ DataCopyExtParams dataCoptExtParams;
	305	+ dataCoptExtParams.blockCount = nBurst;
	306	+ dataCoptExtParams.blockLen = copyLen * sizeof(T);
	307	+ dataCoptExtParams.srcStride = srcStride * sizeof(T);
	308	+ dataCoptExtParams.dstStride = 0;
	309	+ DataCopyPad(inputTensor, inputGm, dataCoptExtParams, dataCopyPadExtParams);
	310	+}
	311	+
	312	+__aicore__ inline void SetDefaultBlockTiling(
	313	+ const SwigluGroupTilingData* tilingData, int64_t& usedCoreNums, int64_t& rowOfFormerBlock,
	314	+ int64_t& rowOfTailBlock, int64_t& rowLoopOfFormerBlock, int64_t& rowLoopOfTailBlock,
	315	+ int64_t& tailRowFactorOfFormerBlock, int64_t& tailRowFactorOfTailBlock)
	316	+{
	317	+ rowOfFormerBlock = tilingData->rowOfFormerBlock;
	318	+ rowOfTailBlock = tilingData->rowOfTailBlock;
	319	+ rowLoopOfFormerBlock = tilingData->rowLoopOfFormerBlock;
	320	+ rowLoopOfTailBlock = tilingData->rowLoopOfTailBlock;
	321	+ tailRowFactorOfFormerBlock = tilingData->tailRowFactorOfFormerBlock;
	322	+ tailRowFactorOfTailBlock = tilingData->tailRowFactorOfTailBlock;
	323	+ usedCoreNums = GetBlockNum();
	324	+}
	325	+
	326	+__aicore__ inline void SetGroupIndexBlockTiling(
	327	+ const SwigluGroupTilingData* tilingData, int64_t realBs, int64_t& usedCoreNums,
	328	+ int64_t& rowOfFormerBlock, int64_t& rowOfTailBlock, int64_t& rowLoopOfFormerBlock,
	329	+ int64_t& rowLoopOfTailBlock, int64_t& tailRowFactorOfFormerBlock, int64_t& tailRowFactorOfTailBlock)
	330	+{
	331	+ rowOfFormerBlock = CeilDiv(realBs, static_cast<int64_t>(tilingData->coreNum));
	332	+ usedCoreNums = CeilDiv(realBs, rowOfFormerBlock) < tilingData->coreNum
	333	+ ? CeilDiv(realBs, rowOfFormerBlock)
	334	+ : tilingData->coreNum;
	335	+ rowOfTailBlock = realBs - (usedCoreNums - 1) * rowOfFormerBlock;
	336	+
	337	+ rowLoopOfFormerBlock = CeilDiv(rowOfFormerBlock, tilingData->rowFactor);
	338	+ rowLoopOfTailBlock = CeilDiv(rowOfTailBlock, tilingData->rowFactor);
	339	+ tailRowFactorOfFormerBlock = rowOfFormerBlock % tilingData->rowFactor == 0
	340	+ ? tilingData->rowFactor
	341	+ : rowOfFormerBlock % tilingData->rowFactor;
	342	+ tailRowFactorOfTailBlock = rowOfTailBlock % tilingData->rowFactor == 0
	343	+ ? tilingData->rowFactor
	344	+ : rowOfTailBlock % tilingData->rowFactor;
	345	+}
	346	+
	347	+template <typename TBufPoolType>
	348	+__aicore__ inline void ProcessGroupIndexTiling(
	349	+ GM_ADDR groupIndex, const SwigluGroupTilingData* tilingData, TBufPoolType& tBufPool,
	350	+ TQue<QuePosition::VECIN, 1>& groupIndexQue, TBuf<QuePosition::VECCALC>& groupIndexSumBuf,
	351	+ GlobalTensor<int64_t>& groupIndexGm, LocalTensor<int64_t>& groupSumLocal, bool& hasGroupIndex,
	352	+ int64_t& usedCoreNums, int64_t& rowOfFormerBlock, int64_t& rowOfTailBlock, int64_t& rowLoopOfFormerBlock,
	353	+ int64_t& rowLoopOfTailBlock, int64_t& tailRowFactorOfFormerBlock, int64_t& tailRowFactorOfTailBlock)
	354	+{
	355	+ if (groupIndex == nullptr) {
	356	+ SetDefaultBlockTiling(tilingData, usedCoreNums, rowOfFormerBlock, rowOfTailBlock, rowLoopOfFormerBlock,
	357	+ rowLoopOfTailBlock, tailRowFactorOfFormerBlock, tailRowFactorOfTailBlock);
	358	+ return;
	359	+ }
	360	+
	361	+ hasGroupIndex = true;
	362	+ groupIndexGm.SetGlobalBuffer((__gm__ int64_t*)groupIndex);
	363	+ tBufPool.InitBuffer(groupIndexQue, DOUBLE_BUFFER_NUM, RoundUp<int64_t>(tilingData->gFactor) * sizeof(int64_t));
	364	+ tBufPool.InitBuffer(groupIndexSumBuf, BLOCK_SIZE);
	365	+ groupSumLocal = groupIndexSumBuf.Get<int64_t>();
	366	+ for (int64_t idx = 0; idx < tilingData->gLoop; idx++) {
	367	+ int64_t curGFactor = (idx == tilingData->gLoop - 1) ? tilingData->tailGFactor : tilingData->gFactor;
	368	+ LocalTensor<int64_t> groupIndexLocal = groupIndexQue.template AllocTensor<int64_t>();
	369	+ CopyIn(groupIndexGm[idx * tilingData->gFactor], groupIndexLocal, 1, curGFactor);
	370	+ groupIndexQue.template EnQue(groupIndexLocal);
	371	+ groupIndexLocal = groupIndexQue.template DeQue<int64_t>();
	372	+ if (idx == 0) {
	373	+ VFProcessGroupIndex<int64_t, false>(groupSumLocal, groupIndexLocal, curGFactor);
	374	+ } else {
	375	+ VFProcessGroupIndex<int64_t, true>(groupSumLocal, groupIndexLocal, curGFactor);
	376	+ }
	377	+ groupIndexQue.template FreeTensor(groupIndexLocal);
	378	+ }
	379	+ event_t eventId = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::V_S));
	380	+ SetFlag<HardEvent::V_S>(eventId);
	381	+ WaitFlag<HardEvent::V_S>(eventId);
	382	+ int64_t groupSum = groupSumLocal.GetValue(0);
	383	+ int64_t realBs = groupSum > tilingData->bs ? tilingData->bs : groupSum;
	384	+ SetGroupIndexBlockTiling(tilingData, realBs, usedCoreNums, rowOfFormerBlock, rowOfTailBlock,
	385	+ rowLoopOfFormerBlock, rowLoopOfTailBlock, tailRowFactorOfFormerBlock, tailRowFactorOfTailBlock);
	386	+ tBufPool.Reset();
	387	+}
	388	+
	389	+template <typename T, AscendC::PaddingMode mode = AscendC::PaddingMode::Normal>
	390	+__aicore__ inline void CopyOut(
	391	+ const LocalTensor<T>& outputTensor, const GlobalTensor<T>& outputGm, const uint16_t nBurst, const uint32_t copyLen,
	392	+ uint32_t dstStride = 0)
	393	+{
	394	+ DataCopyExtParams dataCopyParams;
	395	+ dataCopyParams.blockCount = nBurst;
	396	+ dataCopyParams.blockLen = copyLen * sizeof(T);
	397	+ dataCopyParams.srcStride = 0;
	398	+ dataCopyParams.dstStride = dstStride * sizeof(T);
	399	+ DataCopyPad<T, mode>(outputGm, outputTensor, dataCopyParams);
	400	+}
	401	+} // namespace SwigluGroup
	402	+
	403	+#endif

Aactivation/swiglu_group/op_kernel/arch35/swiglu_group_perf.h+167-0

@@ -0,0 +1,167 @@

1 +/**

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 swiglu_group_perf.h

13 + * \brief

14 + */

15 +

16 +#ifndef SWIGLU_GROUP_PERF_H

17 +#define SWIGLU_GROUP_PERF_H

18 +

19 +#include "kernel_operator.h"

20 +#include "swiglu_group_base.h"

21 +

22 +namespace SwigluGroup {

23 +using namespace AscendC;

24 +template <typename T>

25 +class SwigluGroupPerf {

26 +public:

27 + __aicore__ inline SwigluGroupPerf()

28 + {}

29 +

30 + __aicore__ inline void Init(

31 + GM_ADDR x, GM_ADDR weight, GM_ADDR groupIndex, GM_ADDR y,

32 + GM_ADDR workspace, const SwigluGroupTilingData* tilingDataPtr, TPipe* pipePtr)

33 + {

34 + pipe = pipePtr;

35 + tilingData = tilingDataPtr;

36 +

37 + xGm.SetGlobalBuffer((__gm__ T*)x);

38 + yGm.SetGlobalBuffer((__gm__ T*)y);

39 +

40 + pipe->InitBufPool(tBufPool, tilingData->ubSize);

41 + ProcessGroupIndexTiling(groupIndex, tilingData, tBufPool, groupIndexQue, groupIndexSumBuf, groupIndexGm,

42 + groupSumLocal, hasGroupIndex_, usedCoreNums, rowOfFormerBlock, rowOfTailBlock, rowLoopOfFormerBlock,

43 + rowLoopOfTailBlock, tailRowFactorOfFormerBlock, tailRowFactorOfTailBlock);

44 +

45 + if (weight != nullptr) {

46 + hasWeight_ = true;

47 + weightGm.SetGlobalBuffer((__gm__ float*)weight);

48 + tBufPool.InitBuffer(weightQue, DOUBLE_BUFFER_NUM, RoundUp<float>(tilingData->rowFactor) * sizeof(float));

49 + }

50 +

51 + tBufPool.InitBuffer(

52 + x0Que, DOUBLE_BUFFER_NUM, tilingData->rowFactor * RoundUp<T>(tilingData->dFactor) * sizeof(T));

53 + tBufPool.InitBuffer(

54 + x1Que, DOUBLE_BUFFER_NUM, tilingData->rowFactor * RoundUp<T>(tilingData->dFactor) * sizeof(T));

55 + tBufPool.InitBuffer(

56 + yQue, DOUBLE_BUFFER_NUM, tilingData->rowFactor * RoundUp<T>(tilingData->dFactor) * sizeof(T));

57 + hasClampValue_ = (tilingData->hasClampLimit == 1);

58 + clampValue_ = tilingData->clampLimit;

59 + }

60 +

61 + __aicore__ inline void Process()

62 + {

63 + if (GetBlockIdx() >= usedCoreNums) {

64 + return;

65 + }

66 + int64_t curBlockIdx = GetBlockIdx();

67 + int64_t rowOuterLoop =

68 + (curBlockIdx == usedCoreNums - 1) ? rowLoopOfTailBlock : rowLoopOfFormerBlock;

69 + int64_t tailRowFactor = (curBlockIdx == usedCoreNums - 1) ? tailRowFactorOfTailBlock :

70 + tailRowFactorOfFormerBlock;

71 + int64_t x0GmBaseOffset = curBlockIdx * rowOfFormerBlock * tilingData->d;

72 + int64_t x1GmBaseOffset = x0GmBaseOffset + tilingData->splitD;

73 + int64_t yGmBaseOffset = curBlockIdx * rowOfFormerBlock * tilingData->splitD;

74 + int64_t weightGmBaseOffset = curBlockIdx * rowOfFormerBlock;

75 + for (int64_t rowOuterIdx = 0; rowOuterIdx < rowOuterLoop; rowOuterIdx++) {

76 + int64_t curRowFactor = (rowOuterIdx == rowOuterLoop - 1) ? tailRowFactor : tilingData->rowFactor;

77 + if (hasWeight_) {

78 + weightLocal = weightQue.template AllocTensor<float>();

79 + CopyIn(weightGm[weightGmBaseOffset + rowOuterIdx * tilingData->rowFactor],

80 + weightLocal, 1, curRowFactor);

81 + weightQue.template EnQue(weightLocal);

82 + weightLocal = weightQue.template DeQue<float>();

83 + }

84 +

85 + for (int64_t dLoopIdx = 0; dLoopIdx < tilingData->dLoop; dLoopIdx++) {

86 + int64_t curDFactor =

87 + (dLoopIdx == tilingData->dLoop - 1) ? tilingData->tailDFactor : tilingData->dFactor;

88 + int64_t xBaseOffset =

89 + rowOuterIdx * tilingData->rowFactor * tilingData->d + dLoopIdx * tilingData->dFactor;

90 + x0Local = x0Que.template AllocTensor<T>();

91 + CopyIn(

92 + xGm[x0GmBaseOffset + xBaseOffset],

93 + x0Local, curRowFactor, curDFactor, tilingData->d - curDFactor);

94 + x0Que.template EnQue(x0Local);

95 + x0Local = x0Que.template DeQue<T>();

96 +

97 + x1Local = x1Que.template AllocTensor<T>();

98 + CopyIn(

99 + xGm[x1GmBaseOffset + xBaseOffset],

100 + x1Local, curRowFactor, curDFactor, tilingData->d - curDFactor);

101 + x1Que.template EnQue(x1Local);

102 + x1Local = x1Que.template DeQue<T>();

103 +

104 + yLocal = yQue.template AllocTensor<T>();

105 +

106 + int32_t maskBit = (hasClampValue_ << 1) | hasWeight_;

107 + SwigluGroupDispatcher<T>(yLocal, x0Local, x1Local, weightLocal, clampValue_,

108 + curRowFactor, curDFactor, maskBit);

109 +

110 + x0Que.template FreeTensor(x0Local);

111 + x1Que.template FreeTensor(x1Local);

112 +

113 + yQue.template EnQue(yLocal);

114 + yLocal = yQue.template DeQue<T>();

115 + CopyOut(yLocal, yGm[yGmBaseOffset + rowOuterIdx * tilingData->rowFactor *

116 + tilingData->splitD + dLoopIdx * tilingData->dFactor],

117 + curRowFactor, curDFactor, tilingData->splitD - curDFactor);

118 + yQue.template FreeTensor(yLocal);

119 + }

120 + if (hasWeight_) {

121 + weightQue.template FreeTensor(weightLocal);

122 + }

123 + }

124 + }

125 +

126 +private:

127 + TPipe* pipe;

128 + const SwigluGroupTilingData* tilingData;

129 + GlobalTensor<T> xGm;

130 + GlobalTensor<T> yGm;

131 + GlobalTensor<float> weightGm;

132 + GlobalTensor<int64_t> groupIndexGm;

133 +

134 + TQue<QuePosition::VECIN, 1> x0Que;

135 + TQue<QuePosition::VECIN, 1> x1Que;

136 + TQue<QuePosition::VECOUT, 1> yQue;

137 + TQue<QuePosition::VECIN, 1> weightQue;

138 +

139 + TQue<QuePosition::VECIN, 1> groupIndexQue;

140 + TBuf<QuePosition::VECCALC> groupIndexSumBuf;

141 + TBufPool<QuePosition::VECCALC, 12> tBufPool;

142 +

143 + LocalTensor<T> x0Local;

144 + LocalTensor<T> x1Local;

145 + LocalTensor<T> yLocal;

146 + LocalTensor<float> weightLocal;

147 +

148 + LocalTensor<int64_t> groupIndexLocal;

149 + LocalTensor<int64_t> groupSumLocal;

150 +

151 + float clampValue_ = 448.0f;

atomgit-bot6月24日

🔵 Low Priority

在 swiglu_group_perf.h 第152行，成员变量 clampValue_ 的类内初始值为 448.0f。该值在 Init() 方法（第58行）中始终被 clampValue_ = tilingData->clampLimit; 覆盖，因此 448.0f 永远不会被实际使用。然而 448.0f 是一个特定且不寻常的数值（疑似测试残留），如果未来有人重构 Init() 但遗漏了 clampValue_ 的赋值，kernel 将悄无声息地使用这个错误的 clamp 值，导致计算结果错误。

此问题的触发条件是：未来代码变更使得 Init() 中的 clampValue_ 赋值被跳过或移除。在当前代码中该值不会被实际读取，属于死代码/代码异味。

建议：将 clampValue_ 的默认值改为 0.0f 或 -1.0f（与 DEFAULT_CLAMP_LIMIT = -1.0f 保持一致），以消除误导。"

改动建议

151	- float clampValue_ = ~~448~~.0f;
151	+ float clampValue_ = 0.0f;

应用建议

152 + bool hasWeight_ = false;

153 + bool hasClampValue_ = false;

154 +

155 + bool hasGroupIndex_ = false;

156 + int64_t tailRowFactorOfTailBlock = 0;

157 + int64_t tailRowFactorOfFormerBlock = 0;

158 + int64_t rowLoopOfTailBlock = 0;

159 + int64_t rowLoopOfFormerBlock = 0;

160 + int64_t usedCoreNums = 0;

161 + int64_t rowOfFormerBlock = 0;

162 + int64_t rowOfTailBlock = 0;

163 +};

164 +

165 +} // namespace SwigluGroup

166 +

167 +#endif

Aactivation/swiglu_group/op_kernel/swiglu_group.cpp+38-0

@@ -0,0 +1,38 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file swiglu_group.cpp
	13	+ * \brief
	14	+ */
	15	+
	16	+#include "arch35/swiglu_group_perf.h"
	17	+#define SWIGLU_GROUP_TILING_KEY 1000
	18	+using namespace AscendC;
	19	+
	20	+extern "C" __global__ __aicore__ void swiglu_group(GM_ADDR x, GM_ADDR weight, GM_ADDR groupIndex, GM_ADDR y,
	21	+ GM_ADDR workspace, GM_ADDR tiling)
	22	+{
	23	+ if (workspace == nullptr) {
	24	+ return;
	25	+ }
	26	+
	27	+ GM_ADDR userWs = GetUserWorkspace(workspace);
	28	+ if (userWs == nullptr) {
	29	+ return;
	30	+ }
	31	+ GET_TILING_DATA(tilingData, tiling);
	32	+ TPipe pipe;
	33	+ if (TILING_KEY_IS(SWIGLU_GROUP_TILING_KEY)) {
	34	+ SwigluGroup::SwigluGroupPerf<DTYPE_X> op;
	35	+ op.Init(x, weight, groupIndex, y, userWs, &tilingData, &pipe);
	36	+ op.Process();
	37	+ }
	38	+}

Aactivation/swiglu_group/tests/CMakeLists.txt+16-0

@@ -0,0 +1,16 @@
	1	+# ----------------------------------------------------------------------------
	2	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+# CANN Open Software License Agreement Version 2.0 (the "License").
	5	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+# See LICENSE in the root of the software repository for the full text of the License.
	9	+# ----------------------------------------------------------------------------
	10	+
	11	+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
	atomgit-bot6月24日 🔵 Low Priority `tests/CMakeLists.txt`（第11行）和 `tests/ut/CMakeLists.txt`（第11行）使用 `file(GLOB ...)` 发现子目录以执行 `add_subdirectory`。虽然这里 GLOB 的是目录而非源文件，但当新增测试子目录时，若未重新运行 CMake 配置步骤，新增的测试子目录不会被自动发现和加入构建，导致新测试不被编译执行。这是 CMake 官方文档明确指出的反模式。建议：将 `file(GLOB ...)` 替换为显式列出已知子目录（`add_subdirectory(st)`、`add_subdirectory(ut)` 等），或者在新增子目录时提醒开发者需重新运行 CMake。考虑到当前子目录结构固定（st、ut），建议使用显式列举。
	12	+foreach(SUB_DIR ${CURRENT_DIRS})
	13	+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
	14	+ add_subdirectory(${SUB_DIR})
	15	+ endif()
	16	+endforeach()

Aactivation/swiglu_group/tests/assets/golden.py+183-0

@@ -0,0 +1,183 @@

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 +"""Golden / input plugin for the SwigluGroup operator (plain SwiGLU, no quant).

10 +

11 + y = silu(A) * B

12 +

13 +where the last dim of ``x`` is evenly split into ``A`` and ``B``. ``weight`` is an

14 +optional per-token scale; ``clamp_limit`` (> 0) clamps the SwiGLU inputs before

15 +activation. ``group_index`` (count mode) only bounds how many rows are processed

16 +(realBs = min(sum(group_index), bs)); the input plugin makes its sum equal the

17 +token count so every row is produced and the all-rows golden matches.

18 +"""

19 +

20 +import numpy as np

21 +

22 +from ttk.utilities.dtypes import (

23 + numpy_bfloat16,

24 + numpy_to_torch_tensor,

25 + torch_to_numpy_tensor,

26 +)

27 +

28 +

29 +__golden__ = {

30 + "kernel": {"swiglu_group": "swiglu_group_golden"},

31 + "aclnn": {"aclnnSwigluGroup": "aclnn_swiglu_group_golden"},

32 +}

33 +

34 +__input__ = {

35 + "kernel": {"swiglu_group": "swiglu_group_input"},

36 + "aclnn": {"aclnnSwigluGroup": "aclnn_swiglu_group_input"},

37 +}

38 +

39 +DEFAULT_CLAMP_LIMIT = -1.0

40 +

41 +# inf / nan / boundary pool used when the testcase name contains "edge". 65504 = fp16 max,

42 +# 6.10e-5 = fp16 smallest normal, 5.96e-8 = fp16 smallest denormal. The golden computes

43 +# SwiGLU in numpy, which propagates inf/nan IEEE-consistently with the kernel.

44 +EDGE_VALUES = np.array(

45 + [np.inf, -np.inf, np.nan, 0.0, -0.0, 1.0, -1.0, 2.0, -2.0,

46 + 65504.0, -65504.0, 6.10e-5, -6.10e-5, 5.96e-8, 1.0e4, -1.0e4],

47 + dtype=np.float32,

48 +)

49 +

50 +

51 +def _is_edge_case(kwargs):

52 + return "edge" in str(kwargs.get("testcase_name", "")).lower()

53 +

54 +

55 +def _fill_edge(x, weight):

56 + """Fill x entirely with the inf/nan/boundary pool; set weight (if any) to ones."""

57 + x_np = _to_numpy(x)

58 + _write_tensor(x, np.resize(EDGE_VALUES, x_np.size).reshape(x_np.shape))

59 + if weight is not None:

60 + _write_tensor(weight, np.ones(_to_numpy(weight).shape, dtype=np.float32))

61 +

62 +

63 +def _to_numpy(tensor):

64 + if tensor is None:

65 + return None

66 + if isinstance(tensor, np.ndarray):

67 + return tensor

68 + if hasattr(tensor, "detach"):

69 + return torch_to_numpy_tensor(tensor.detach().cpu())

70 + if hasattr(tensor, "cpu"):

71 + return tensor.cpu().numpy()

72 + return np.asarray(tensor)

73 +

74 +

75 +def _write_tensor(tensor, value):

76 + if tensor is None:

77 + return

78 + if isinstance(tensor, np.ndarray):

79 + tensor[...] = value.astype(tensor.dtype, copy=False)

80 + return

81 + import torch

82 +

83 + src = torch.as_tensor(value, device=tensor.device)

84 + if src.dtype != tensor.dtype:

85 + src = src.to(tensor.dtype)

86 + tensor.copy_(src.reshape(tensor.shape))

87 +

88 +

89 +def _sigmoid(x):

90 + return 1.0 / (1.0 + np.exp(-x))

91 +

92 +

93 +def _silu(x):

94 + return x * _sigmoid(x)

95 +

96 +

97 +def _use_clamp(clamp_limit):

98 + return clamp_limit is not None and float(clamp_limit) != DEFAULT_CLAMP_LIMIT

99 +

100 +

101 +def _compute_swiglu(x, weight=None, clamp_limit=None):

102 + x = _to_numpy(x).astype(np.float32)

103 + orig_shape = x.shape

104 + x = x.reshape(-1, orig_shape[-1])

105 + hidden = orig_shape[-1] // 2

106 + x0 = x[:, :hidden]

107 + x1 = x[:, hidden:]

108 +

109 + if _use_clamp(clamp_limit):

110 + limit = float(clamp_limit)

111 + x0 = np.minimum(x0, limit)

112 + x1 = np.minimum(limit, np.maximum(x1, -limit))

113 +

114 + y = _silu(x0) * x1

115 + weight = _to_numpy(weight)

116 + if weight is not None:

117 + y *= weight.reshape(-1, 1).astype(np.float32)

118 + return y.reshape(*orig_shape[:-1], hidden)

119 +

120 +

121 +def _cast_like(y, x):

122 + x_np = _to_numpy(x)

123 + if x_np.dtype.name == "bfloat16":

124 + return y.astype(numpy_bfloat16())

125 + return y.astype(x_np.dtype, copy=False)

126 +

127 +

128 +def _maybe_to_torch(value, use_torch):

129 + if not use_torch or value is None or not isinstance(value, np.ndarray):

130 + return value

131 + return numpy_to_torch_tensor(value)

132 +

133 +

134 +def _fill_group_index(x, group_index):

135 + """Make sum(group_index) == token count so realBs == bs and all rows are produced."""

136 + if group_index is None:

137 + return

138 + token_num = int(np.prod(_to_numpy(x).shape[:-1]))

139 + gi = _to_numpy(group_index).reshape(-1).copy()

140 + gi[...] = 0

141 + n = gi.size

142 + if n == 1:

143 + gi[0] = token_num

144 + else:

145 + base = token_num // n

146 + gi[:] = base

147 + gi[-1] = token_num - base * (n - 1)

148 + _write_tensor(group_index, gi)

149 +

150 +

151 +# --------------------------------------------------------------------------- #

152 +# Kernel level (numpy)

153 +# --------------------------------------------------------------------------- #

154 +def swiglu_group_golden(x, weight, group_index, clamp_limit=DEFAULT_CLAMP_LIMIT, **kwargs):

155 + """Golden for swiglu_group. Parameters follow swiglu_group_def.cpp (no outputs)."""

156 + del group_index, kwargs

157 + return _cast_like(_compute_swiglu(x, weight, clamp_limit), x)

158 +

159 +

160 +def swiglu_group_input(x, weight, group_index, clamp_limit=DEFAULT_CLAMP_LIMIT, **kwargs):

161 + del clamp_limit

162 + if _is_edge_case(kwargs):

163 + _fill_edge(x, weight)

164 + _fill_group_index(x, group_index)

165 + return [x, weight, group_index]

166 +

167 +

168 +# --------------------------------------------------------------------------- #

169 +# ACLNN level (torch)

170 +# --------------------------------------------------------------------------- #

171 +def aclnn_swiglu_group_golden(x, weightOptional, groupIndexOptional, clampLimit, yOut, **kwargs):

172 + """Golden for aclnnSwigluGroup. Parameters follow aclnn_swiglu_group.h."""

173 + del groupIndexOptional, yOut

174 + use_torch = kwargs.get("use_torch", False)

175 + return _maybe_to_torch(_cast_like(_compute_swiglu(x, weightOptional, clampLimit), x), use_torch)

176 +

177 +

178 +def aclnn_swiglu_group_input(x, weightOptional, groupIndexOptional, clampLimit, yOut, **kwargs):

179 + del clampLimit, yOut

180 + if _is_edge_case(kwargs):

181 + _fill_edge(x, weightOptional)

182 + _fill_group_index(x, groupIndexOptional)

atomgit-bot6月23日

🟠 High Priority

函数 aclnn_swiglu_group_input 在修改输入张量后缺少 return 语句，因此隐式返回 None。对比同文件中的 swiglu_group_input 函数（第160-165行），后者正确返回了 [x, weight, group_index]。TTK 测试框架的 input plugin 依赖于返回值来获取修改后的输入；返回 None 将导致 aclnn 级别的 ST 测试在调用 input plugin 时失败（TypeError 或 None unpacking），使所有 aclnnSwigluGroup 的 ST 用例无法正确执行。

建议：在函数末尾添加 return [x, weightOptional, groupIndexOptional, clampLimit, yOut]，与 kernel 级别 swiglu_group_input 的返回模式保持一致。

改动建议

+def aclnn_swiglu_group_input(x, weightOptional, groupIndexOptional, clampLimit, yOut, **kwargs):
+    del clampLimit, yOut
+    if _is_edge_case(kwargs):
+        _fill_edge(x, weightOptional)
-    _fill_group_index(x, groupIndexOptional)
+    _fill_group_index(x, groupIndexOptional)
+    return [x, weightOptional, groupIndexOptional, clampLimit, yOut]

应用建议

atomgit-bot6月24日

🟠 High Priority

函数 aclnn_swiglu_group_input（第178-182行）在修改输入张量后缺少 return 语句，Python 隐式返回 None。对比同文件中的 swiglu_group_input 函数（第160-165行），后者正确返回了 [x, weight, group_index]。TTK 测试框架的 input plugin 依赖于返回值来获取修改后的输入张量；返回 None 将导致 aclnn 级别的所有测试（包括 sg_fp16、sg_bf16、sg_weight_fp16、sg_clamp_bf16、sg_group_fp16、sg_group_large_fp16）在运行时因无法解包 None 而崩溃（TypeError）。

建议：在函数末尾添加 return [x, weightOptional, groupIndexOptional, clampLimit, yOut]，与 kernel 级别 swiglu_group_input 的返回模式保持一致。

改动建议

+def aclnn_swiglu_group_input(x, weightOptional, groupIndexOptional, clampLimit, yOut, **kwargs):
+    del clampLimit, yOut
+    if _is_edge_case(kwargs):
+        _fill_edge(x, weightOptional)
-    _fill_group_index(x, groupIndexOptional)
+    _fill_group_index(x, groupIndexOptional)
+    return [x, weightOptional, groupIndexOptional, clampLimit, yOut]

应用建议

183 + return [x, weightOptional, groupIndexOptional]

Aactivation/swiglu_group/tests/st/ttk_aclnn_swiglu_group.csv+9-0

@@ -0,0 +1,9 @@
	1	+testcase_name,api_name,tensor_view_shapes,tensor_dtypes,tensor_formats,attributes,output_tensor_indexes,input_data_ranges,precision_tolerances,absolute_precision,soc_series
	2	+sg_fp16,aclnnSwigluGroup,"((2,256),None,None,(2,128))","('float16','float32','int64','float16')","('ND',)",{'clampLimit':-1.0},"(3,)","((-2,2),None,None,(0,0))","((0.001,0.001),)",0.001,"('Ascend950',)"
	3	+sg_bf16,aclnnSwigluGroup,"((3,512),None,None,(3,256))","('bfloat16','float32','int64','bfloat16')","('ND',)",{'clampLimit':-1.0},"(3,)","((-2,2),None,None,(0,0))","((0.004,0.004),)",0.004,"('Ascend950',)"
	4	+sg_weight_fp16,aclnnSwigluGroup,"((8,512),(8,1),None,(8,256))","('float16','float32','int64','float16')","('ND',)",{'clampLimit':-1.0},"(3,)","((-2,2),(-1,1),None,(0,0))","((0.001,0.001),)",0.001,"('Ascend950',)"
	5	+sg_clamp_bf16,aclnnSwigluGroup,"((4,512),None,None,(4,256))","('bfloat16','float32','int64','bfloat16')","('ND',)",{'clampLimit':2.0},"(3,)","((-2,2),None,None,(0,0))","((0.004,0.004),)",0.004,"('Ascend950',)"
	6	+sg_group_fp16,aclnnSwigluGroup,"((8,512),None,(2,),(8,256))","('float16','float32','int64','float16')","('ND',)",{'clampLimit':-1.0},"(3,)","((-2,2),None,(0,8),(0,0))","((0.001,0.001),)",0.001,"('Ascend950',)"
	7	+sg_group_large_fp16,aclnnSwigluGroup,"((128,512),None,(128,),(128,256))","('float16','float32','int64','float16')","('ND',)",{'clampLimit':-1.0},"(3,)","((-2,2),None,(0,128),(0,0))","((0.001,0.001),)",0.001,"('Ascend950',)"
	8	+sg_f32,aclnnSwigluGroup,"((2,256),None,None,(2,128))","('float32','float32','int64','float32')","('ND',)",{'clampLimit':-1.0},"(3,)","((-2,2),None,None,(0,0))","((0.0001,0.0001),)",0.0001,"('Ascend950',)"
	9	+sg_f32_group,aclnnSwigluGroup,"((8,512),None,(2,),(8,256))","('float32','float32','int64','float32')","('ND',)",{'clampLimit':-1.0},"(3,)","((-2,2),None,(0,8),(0,0))","((0.0001,0.0001),)",0.0001,"('Ascend950',)"

Aactivation/swiglu_group/tests/st/ttk_kernel_swiglu_group.csv+18-0

@@ -0,0 +1,18 @@
	1	+testcase_name,network_name,op_name,input_shapes,input_dtypes,input_formats,output_shapes,output_dtypes,output_formats,input_ori_shapes,input_ori_formats,output_ori_shapes,output_ori_formats,attributes,input_data_ranges,precision_tolerances,absolute_precision,output_inplace_indexes,output_shape_unknown_indexes,is_enabled,remark,soc_series,priority,dump_file_prefix,manual_input_binaries,manual_golden_binaries
	2	+sg_block_fp16,UNKNOWN,swiglu_group,"((2,256),None,None)","('float16','float32','int64')","('ND','ND','ND')","((2,128),)","('float16',)","('ND',)","((2,256),None,None)","('ND','ND','ND')","((2,128),)","('ND',)",{'clamp_limit':-1.0},"((-2,2),None,None)","((0.001,0.001),)",0.001,(),(),True,,"('Ascend950',)",0,,(),()
	3	+sg_block_bf16,UNKNOWN,swiglu_group,"((3,512),None,None)","('bfloat16','float32','int64')","('ND','ND','ND')","((3,256),)","('bfloat16',)","('ND',)","((3,512),None,None)","('ND','ND','ND')","((3,256),)","('ND',)",{'clamp_limit':-1.0},"((-2,2),None,None)","((0.004,0.004),)",0.004,(),(),True,,"('Ascend950',)",0,,(),()
	4	+sg_block_fp16_3d,UNKNOWN,swiglu_group,"((2,3,768),None,None)","('float16','float32','int64')","('ND','ND','ND')","((2,3,384),)","('float16',)","('ND',)","((2,3,768),None,None)","('ND','ND','ND')","((2,3,384),)","('ND',)",{'clamp_limit':-1.0},"((-2,2),None,None)","((0.001,0.001),)",0.001,(),(),True,,"('Ascend950',)",0,,(),()
	5	+sg_block_bf16_3d,UNKNOWN,swiglu_group,"((1,4,1024),None,None)","('bfloat16','float32','int64')","('ND','ND','ND')","((1,4,512),)","('bfloat16',)","('ND',)","((1,4,1024),None,None)","('ND','ND','ND')","((1,4,512),)","('ND',)",{'clamp_limit':-1.0},"((-2,2),None,None)","((0.004,0.004),)",0.004,(),(),True,,"('Ascend950',)",0,,(),()
	6	+sg_weight_fp16,UNKNOWN,swiglu_group,"((8,512),(8,1),None)","('float16','float32','int64')","('ND','ND','ND')","((8,256),)","('float16',)","('ND',)","((8,512),(8,1),None)","('ND','ND','ND')","((8,256),)","('ND',)",{'clamp_limit':-1.0},"((-2,2),(-1,1),None)","((0.001,0.001),)",0.001,(),(),True,,"('Ascend950',)",0,,(),()
	7	+sg_weight_clamp_bf16,UNKNOWN,swiglu_group,"((4,512),(4,1),None)","('bfloat16','float32','int64')","('ND','ND','ND')","((4,256),)","('bfloat16',)","('ND',)","((4,512),(4,1),None)","('ND','ND','ND')","((4,256),)","('ND',)",{'clamp_limit':2.0},"((-2,2),(-1,1),None)","((0.004,0.004),)",0.004,(),(),True,,"('Ascend950',)",0,,(),()
	8	+sg_clamp_fp16,UNKNOWN,swiglu_group,"((3,512),None,None)","('float16','float32','int64')","('ND','ND','ND')","((3,256),)","('float16',)","('ND',)","((3,512),None,None)","('ND','ND','ND')","((3,256),)","('ND',)",{'clamp_limit':1.0},"((-2,2),None,None)","((0.001,0.001),)",0.001,(),(),True,,"('Ascend950',)",0,,(),()
	9	+sg_group_small_fp16,UNKNOWN,swiglu_group,"((8,512),None,(2,))","('float16','float32','int64')","('ND','ND','ND')","((8,256),)","('float16',)","('ND',)","((8,512),None,(2,))","('ND','ND','ND')","((8,256),)","('ND',)",{'clamp_limit':-1.0},"((-2,2),None,(0,8))","((0.001,0.001),)",0.001,(),(),True,,"('Ascend950',)",0,,(),()
	10	+sg_group_weight_fp16,UNKNOWN,swiglu_group,"((6,512),(6,1),(3,))","('float16','float32','int64')","('ND','ND','ND')","((6,256),)","('float16',)","('ND',)","((6,512),(6,1),(3,))","('ND','ND','ND')","((6,256),)","('ND',)",{'clamp_limit':-1.0},"((-2,2),(-1,1),(0,6))","((0.001,0.001),)",0.001,(),(),True,,"('Ascend950',)",0,,(),()
	11	+sg_group_large_fp16,UNKNOWN,swiglu_group,"((128,512),None,(128,))","('float16','float32','int64')","('ND','ND','ND')","((128,256),)","('float16',)","('ND',)","((128,512),None,(128,))","('ND','ND','ND')","((128,256),)","('ND',)",{'clamp_limit':-1.0},"((-2,2),None,(0,128))","((0.001,0.001),)",0.001,(),(),True,,"('Ascend950',)",0,,(),()
	12	+sg_group_large_tail_bf16,UNKNOWN,swiglu_group,"((27,8,804),None,(97,))","('bfloat16','float32','int64')","('ND','ND','ND')","((27,8,402),)","('bfloat16',)","('ND',)","((27,8,804),None,(97,))","('ND','ND','ND')","((27,8,402),)","('ND',)",{'clamp_limit':-1.0},"((-2,2),None,(0,216))","((0.004,0.004),)",0.004,(),(),True,,"('Ascend950',)",0,,(),()
	13	+sg_even_384_fp16,UNKNOWN,swiglu_group,"((4,384),None,None)","('float16','float32','int64')","('ND','ND','ND')","((4,192),)","('float16',)","('ND',)","((4,384),None,None)","('ND','ND','ND')","((4,192),)","('ND',)",{'clamp_limit':-1.0},"((-2,2),None,None)","((0.001,0.001),)",0.001,(),(),True,,"('Ascend950',)",0,,(),()
	14	+sg_f32_2d,UNKNOWN,swiglu_group,"((4,512),None,None)","('float32','float32','int64')","('ND','ND','ND')","((4,256),)","('float32',)","('ND',)","((4,512),None,None)","('ND','ND','ND')","((4,256),)","('ND',)",{'clamp_limit':-1.0},"((-2,2),None,None)","((0.0001,0.0001),)",0.0001,(),(),True,,"('Ascend950',)",0,,(),()
	15	+sg_f32_3d,UNKNOWN,swiglu_group,"((2,3,768),None,None)","('float32','float32','int64')","('ND','ND','ND')","((2,3,384),)","('float32',)","('ND',)","((2,3,768),None,None)","('ND','ND','ND')","((2,3,384),)","('ND',)",{'clamp_limit':-1.0},"((-2,2),None,None)","((0.0001,0.0001),)",0.0001,(),(),True,,"('Ascend950',)",0,,(),()
	16	+sg_f32_weight_clamp,UNKNOWN,swiglu_group,"((8,512),(8,1),None)","('float32','float32','int64')","('ND','ND','ND')","((8,256),)","('float32',)","('ND',)","((8,512),(8,1),None)","('ND','ND','ND')","((8,256),)","('ND',)",{'clamp_limit':2.0},"((-2,2),(-1,1),None)","((0.0001,0.0001),)",0.0001,(),(),True,,"('Ascend950',)",0,,(),()
	17	+sg_f32_group_large,UNKNOWN,swiglu_group,"((128,512),None,(128,))","('float32','float32','int64')","('ND','ND','ND')","((128,256),)","('float32',)","('ND',)","((128,512),None,(128,))","('ND','ND','ND')","((128,256),)","('ND',)",{'clamp_limit':-1.0},"((-2,2),None,(0,128))","((0.0001,0.0001),)",0.0001,(),(),True,,"('Ascend950',)",0,,(),()
	18	+sg_f32_even_384,UNKNOWN,swiglu_group,"((4,384),None,None)","('float32','float32','int64')","('ND','ND','ND')","((4,192),)","('float32',)","('ND',)","((4,384),None,None)","('ND','ND','ND')","((4,192),)","('ND',)",{'clamp_limit':-1.0},"((-2,2),None,None)","((0.0001,0.0001),)",0.0001,(),(),True,,"('Ascend950',)",0,,(),()

Aactivation/swiglu_group/tests/st/ttk_kernel_swiglu_group_edge.csv+10-0

@@ -0,0 +1,10 @@
	1	+testcase_name,network_name,op_name,input_shapes,input_dtypes,input_formats,output_shapes,output_dtypes,output_formats,input_ori_shapes,input_ori_formats,output_ori_shapes,output_ori_formats,attributes,input_data_ranges,precision_tolerances,absolute_precision,output_inplace_indexes,output_shape_unknown_indexes,is_enabled,remark,soc_series,priority,dump_file_prefix,manual_input_binaries,manual_golden_binaries
	2	+edge_inf_nan_fp16,UNKNOWN,swiglu_group,"((2,256),None,None)","('float16','float32','int64')","('ND','ND','ND')","((2,128),)","('float16',)","('ND',)","((2,256),None,None)","('ND','ND','ND')","((2,128),)","('ND',)",{'clamp_limit':-1.0},"((-2,2),None,None)","((0.001,0.001),)",0.001,(),(),True,inf/nan/boundary,"('Ascend950',)",0,,(),()
	3	+edge_inf_nan_bf16,UNKNOWN,swiglu_group,"((2,256),None,None)","('bfloat16','float32','int64')","('ND','ND','ND')","((2,128),)","('bfloat16',)","('ND',)","((2,256),None,None)","('ND','ND','ND')","((2,128),)","('ND',)",{'clamp_limit':-1.0},"((-2,2),None,None)","((0.004,0.004),)",0.004,(),(),True,inf/nan/boundary,"('Ascend950',)",0,,(),()
	4	+edge_3d_fp16,UNKNOWN,swiglu_group,"((2,3,512),None,None)","('float16','float32','int64')","('ND','ND','ND')","((2,3,256),)","('float16',)","('ND',)","((2,3,512),None,None)","('ND','ND','ND')","((2,3,256),)","('ND',)",{'clamp_limit':-1.0},"((-2,2),None,None)","((0.001,0.001),)",0.001,(),(),True,inf/nan/boundary,"('Ascend950',)",0,,(),()
	5	+edge_clamp_fp16,UNKNOWN,swiglu_group,"((2,256),None,None)","('float16','float32','int64')","('ND','ND','ND')","((2,128),)","('float16',)","('ND',)","((2,256),None,None)","('ND','ND','ND')","((2,128),)","('ND',)",{'clamp_limit':2.0},"((-2,2),None,None)","((0.001,0.001),)",0.001,(),(),True,inf/nan/boundary,"('Ascend950',)",0,,(),()
	6	+edge_clamp_bf16,UNKNOWN,swiglu_group,"((2,256),None,None)","('bfloat16','float32','int64')","('ND','ND','ND')","((2,128),)","('bfloat16',)","('ND',)","((2,256),None,None)","('ND','ND','ND')","((2,128),)","('ND',)",{'clamp_limit':2.0},"((-2,2),None,None)","((0.004,0.004),)",0.004,(),(),True,inf/nan/boundary,"('Ascend950',)",0,,(),()
	7	+edge_weight_fp16,UNKNOWN,swiglu_group,"((4,256),(4,1),None)","('float16','float32','int64')","('ND','ND','ND')","((4,128),)","('float16',)","('ND',)","((4,256),(4,1),None)","('ND','ND','ND')","((4,128),)","('ND',)",{'clamp_limit':-1.0},"((-2,2),(-1,1),None)","((0.001,0.001),)",0.001,(),(),True,inf/nan/boundary,"('Ascend950',)",0,,(),()
	8	+edge_group_bf16,UNKNOWN,swiglu_group,"((4,512),None,(2,))","('bfloat16','float32','int64')","('ND','ND','ND')","((4,256),)","('bfloat16',)","('ND',)","((4,512),None,(2,))","('ND','ND','ND')","((4,256),)","('ND',)",{'clamp_limit':-1.0},"((-2,2),None,(0,4))","((0.004,0.004),)",0.004,(),(),True,inf/nan/boundary,"('Ascend950',)",0,,(),()
	9	+edge_inf_nan_f32,UNKNOWN,swiglu_group,"((2,256),None,None)","('float32','float32','int64')","('ND','ND','ND')","((2,128),)","('float32',)","('ND',)","((2,256),None,None)","('ND','ND','ND')","((2,128),)","('ND',)",{'clamp_limit':-1.0},"((-2,2),None,None)","((0.0001,0.0001),)",0.0001,(),(),True,inf/nan/boundary,"('Ascend950',)",0,,(),()
	10	+edge_clamp_f32,UNKNOWN,swiglu_group,"((2,256),None,None)","('float32','float32','int64')","('ND','ND','ND')","((2,128),)","('float32',)","('ND',)","((2,256),None,None)","('ND','ND','ND')","((2,128),)","('ND',)",{'clamp_limit':2.0},"((-2,2),None,None)","((0.0001,0.0001),)",0.0001,(),(),True,inf/nan/boundary,"('Ascend950',)",0,,(),()

Aactivation/swiglu_group/tests/ut/CMakeLists.txt+16-0

@@ -0,0 +1,16 @@
	1	+# ----------------------------------------------------------------------------
	2	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+# CANN Open Software License Agreement Version 2.0 (the "License").
	5	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+# See LICENSE in the root of the software repository for the full text of the License.
	9	+# ----------------------------------------------------------------------------
	10	+
	11	+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
	12	+foreach(SUB_DIR ${CURRENT_DIRS})
	13	+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
	14	+ add_subdirectory(${SUB_DIR})
	15	+ endif()
	16	+endforeach()

Aactivation/swiglu_group/tests/ut/op_host/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	+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()

Aactivation/swiglu_group/tests/ut/op_host/arch35/test_swiglu_group_tiling.cpp+262-0

@@ -0,0 +1,262 @@
	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 <map>
	13	+#include <string>
	14	+#include "exe_graph/runtime/storage_format.h"
	15	+#include "exe_graph/runtime/storage_shape.h"
	16	+#include <gtest/gtest.h>
	17	+#include "kernel_run_context_facker.h"
	18	+#include "log/log.h"
	19	+#include "platform/platform_infos_def.h"
	20	+#include "register/op_impl_registry.h"
	21	+#include "test_cube_util.h"
	22	+#include "ut_op_common.h"
	23	+#include "ut_op_util.h"
	24	+#include "../../../../op_host/arch35/swiglu_group_tiling.h"
	25	+
	26	+namespace {
	27	+constexpr float DEFAULT_CLAMP_LIMIT = -1.0f;
	28	+
	29	+struct TilingCase {
	30	+ ge::DataType xDtype = ge::DT_FLOAT16;
	31	+ ge::DataType yDtype = ge::DT_FLOAT16;
	32	+ ge::DataType weightDtype = ge::DT_FLOAT;
	33	+ ge::DataType groupIndexDtype = ge::DT_INT64;
	34	+ gert::StorageShape xShape = {{8, 128, 8192}, {8, 128, 8192}};
	35	+ gert::StorageShape weightShape = {{1024}, {1024}};
	36	+ gert::StorageShape groupIndexShape = {{2}, {2}};
	37	+ gert::StorageShape yShape = {{8, 128, 4096}, {8, 128, 4096}};
	38	+ float clampLimit = DEFAULT_CLAMP_LIMIT;
	39	+ bool hasWeight = false;
	40	+ bool hasGroupIndex = false;
	41	+ ge::graphStatus status = ge::GRAPH_SUCCESS;
	42	+};
	43	+
	44	+class SwigluGroupTilingTest : public testing::Test {
	45	+protected:
	46	+ static void SetUpTestCase()
	47	+ {
	48	+ std::cout << "SwigluGroupTilingTest SetUp" << std::endl;
	49	+ }
	50	+
	51	+ static void TearDownTestCase()
	52	+ {
	53	+ std::cout << "SwigluGroupTilingTest TearDown" << std::endl;
	54	+ }
	55	+};
	56	+
	57	+void ExecuteTilingCase(const TilingCase& tc)
	58	+{
	59	+ const std::string compileInfoString = R"({
	60	+ "hardware_info": {"BT_SIZE": 0, "load3d_constraints": "1",
	61	+ "Intrinsic_fix_pipe_l0c2out": false,
	62	+ "Intrinsic_data_move_l12ub": true,
	63	+ "Intrinsic_data_move_l0c2ub": true,
	64	+ "Intrinsic_data_move_out2l1_nd2nz": false,
	65	+ "UB_SIZE": 253952, "L2_SIZE": 33554432, "L1_SIZE": 524288,
	66	+ "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072,
	67	+ "CORE_NUM": 64}
	68	+ })";
	69	+ std::map<std::string, std::string> socInfos;
	70	+ std::map<std::string, std::string> aicoreSpec;
	71	+ std::map<std::string, std::string> intrinsics;
	72	+ std::map<std::string, std::string> socVersions = {{"Short_SoC_version", "Ascend950"}, {"NpuArch", "3510"}};
	73	+ GetPlatFormInfos(compileInfoString.c_str(), socInfos, aicoreSpec, intrinsics);
	74	+
	75	+ fe::PlatFormInfos platformInfo;
	76	+ platformInfo.Init();
	77	+ optiling::SwigluGroupCompileInfo compileInfo;
	78	+
	79	+ const std::string opType("SwigluGroup");
	80	+ ASSERT_NE(gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str()), nullptr);
	81	+ auto tilingFunc = gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str())->tiling;
	82	+ auto tilingParseFunc = gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str())->tiling_parse;
	83	+
	84	+ auto kernelHolder = gert::KernelRunContextFaker()
	85	+ .KernelIONum(2, 1)
	86	+ .Inputs({const_cast<char>(compileInfoString.c_str()), reinterpret_cast<void>(&platformInfo)})
	87	+ .Outputs({&compileInfo})
	88	+ .Build();
	89	+
	90	+ auto* parsePlatform = kernelHolder.GetContext<gert::TilingParseContext>()->GetPlatformInfo();
	91	+ ASSERT_TRUE(parsePlatform->Init());
	92	+ parsePlatform->SetPlatformRes("SoCInfo", socInfos);
	93	+ parsePlatform->SetPlatformRes("AICoreSpec", aicoreSpec);
	94	+ parsePlatform->SetCoreNumByCoreType("AICore");
	95	+ parsePlatform->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics);
	96	+ parsePlatform->SetPlatformRes("version", socVersions);
	97	+ ASSERT_EQ(tilingParseFunc(kernelHolder.GetContext<gert::KernelContext>()), ge::GRAPH_SUCCESS);
	98	+
	99	+ auto tilingData = gert::TilingData::CreateCap(4096);
	100	+ auto workspaceSizeHolder = gert::ContinuousVector::Create<size_t>(4096);
	101	+ auto* workspaceSizes = reinterpret_cast<gert::ContinuousVector*>(workspaceSizeHolder.get());
	102	+ ASSERT_NE(tilingData, nullptr);
	103	+ gert::StorageShape xShape = tc.xShape;
	104	+ gert::StorageShape weightShape = tc.weightShape;
	105	+ gert::StorageShape groupIndexShape = tc.groupIndexShape;
	106	+ gert::StorageShape yShape = tc.yShape;
	107	+
	108	+ std::vector<uint32_t> inputInstanceNum = {1, tc.hasWeight ? 1U : 0U, tc.hasGroupIndex ? 1U : 0U};
	109	+ std::vector<gert::StorageShape*> inputShapes = {&xShape};
	110	+ if (tc.hasWeight \|\| tc.hasGroupIndex) {
	111	+ inputShapes.emplace_back(tc.hasWeight ? &weightShape : nullptr);
	112	+ }
	113	+ if (tc.hasGroupIndex) {
	114	+ inputShapes.emplace_back(&groupIndexShape);
	115	+ }
	116	+
	117	+ gert::TilingContextFaker contextFaker;
	118	+ contextFaker.SetOpType(opType)
	119	+ .NodeIoNum(3, 1)
	120	+ .IrInstanceNum(inputInstanceNum)
	121	+ .InputShapes(inputShapes)
	122	+ .OutputShapes({&yShape})
	123	+ .CompileInfo(&compileInfo)
	124	+ .PlatformInfo(reinterpret_cast<char*>(&platformInfo))
	125	+ .NodeInputTd(0, tc.xDtype, ge::FORMAT_ND, ge::FORMAT_ND)
	126	+ .NodeOutputTd(0, tc.yDtype, ge::FORMAT_ND, ge::FORMAT_ND)
	127	+ .NodeAttrs({{"clamp_limit", Ops::NN::AnyValue::CreateFrom<float>(tc.clampLimit)}})
	128	+ .TilingData(tilingData.get())
	129	+ .Workspace(workspaceSizes);
	130	+ int32_t inputTdIndex = 1;
	131	+ if (tc.hasWeight) {
	132	+ contextFaker.NodeInputTd(inputTdIndex++, tc.weightDtype, ge::FORMAT_ND, ge::FORMAT_ND);
	133	+ }
	134	+ if (tc.hasGroupIndex) {
	135	+ contextFaker.NodeInputTd(inputTdIndex++, tc.groupIndexDtype, ge::FORMAT_ND, ge::FORMAT_ND);
	136	+ }
	137	+
	138	+ auto holder = contextFaker.Build();
	139	+
	140	+ gert::TilingContext* tilingContext = holder.GetContext<gert::TilingContext>();
	141	+ ASSERT_NE(tilingContext, nullptr);
	142	+ ASSERT_NE(tilingContext->GetPlatformInfo(), nullptr);
	143	+ tilingContext->GetPlatformInfo()->SetPlatformRes("SoCInfo", socInfos);
	144	+ tilingContext->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicoreSpec);
	145	+ tilingContext->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
	146	+ tilingContext->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics);
	147	+ tilingContext->GetPlatformInfo()->SetPlatformRes("version", socVersions);
	148	+
	149	+ EXPECT_EQ(tilingFunc(tilingContext), tc.status);
	150	+}
	151	+
	152	+TEST_F(SwigluGroupTilingTest, tiling_fp16)
	153	+{
	154	+ TilingCase tc;
	155	+ ExecuteTilingCase(tc);
	156	+}
	157	+
	158	+TEST_F(SwigluGroupTilingTest, tiling_bf16)
	159	+{
	160	+ TilingCase tc;
	161	+ tc.xDtype = ge::DT_BF16;
	162	+ tc.yDtype = ge::DT_BF16;
	163	+ tc.xShape = {{4, 64, 2048}, {4, 64, 2048}};
	164	+ tc.yShape = {{4, 64, 1024}, {4, 64, 1024}};
	165	+ ExecuteTilingCase(tc);
	166	+}
	167	+
	168	+TEST_F(SwigluGroupTilingTest, tiling_f32)
	169	+{
	170	+ TilingCase tc;
	171	+ tc.xDtype = ge::DT_FLOAT;
	172	+ tc.yDtype = ge::DT_FLOAT;
	173	+ tc.xShape = {{4, 64, 2048}, {4, 64, 2048}};
	174	+ tc.yShape = {{4, 64, 1024}, {4, 64, 1024}};
	175	+ ExecuteTilingCase(tc);
	176	+}
	177	+
	178	+TEST_F(SwigluGroupTilingTest, tiling_clamp_limit)
	179	+{
	180	+ TilingCase tc;
	181	+ tc.clampLimit = 7.0f;
	182	+ ExecuteTilingCase(tc);
	183	+}
	184	+
	185	+TEST_F(SwigluGroupTilingTest, tiling_weight_valid)
	186	+{
	187	+ TilingCase tc;
	188	+ tc.hasWeight = true;
	189	+ ExecuteTilingCase(tc);
	190	+}
	191	+
	192	+TEST_F(SwigluGroupTilingTest, tiling_group_index_valid)
	193	+{
	194	+ TilingCase tc;
	195	+ tc.hasGroupIndex = true;
	196	+ ExecuteTilingCase(tc);
	197	+}
	198	+
	199	+TEST_F(SwigluGroupTilingTest, tiling_error_invalid_last_dim)
	200	+{
	201	+ // Last dim must be divisible by 2; an odd last dim is invalid.
	202	+ TilingCase tc;
	203	+ tc.xShape = {{4, 64, 1023}, {4, 64, 1023}};
	204	+ tc.yShape = {{4, 64, 511}, {4, 64, 511}};
	205	+ tc.status = ge::GRAPH_FAILED;
	206	+ ExecuteTilingCase(tc);
	207	+}
	208	+
	209	+TEST_F(SwigluGroupTilingTest, tiling_even_non256_aligned)
	210	+{
	211	+ // Last dim only needs to be divisible by 2 (relaxed from 256); 384 is valid.
	212	+ TilingCase tc;
	213	+ tc.xShape = {{4, 384}, {4, 384}};
	214	+ tc.yShape = {{4, 192}, {4, 192}};
	215	+ ExecuteTilingCase(tc);
	216	+}
	217	+
	218	+TEST_F(SwigluGroupTilingTest, tiling_error_invalid_x_dtype)
	219	+{
	220	+ // x only supports float16/bfloat16/float32; int32 is invalid.
	221	+ TilingCase tc;
	222	+ tc.xDtype = ge::DT_INT32;
	223	+ tc.yDtype = ge::DT_INT32;
	224	+ tc.status = ge::GRAPH_FAILED;
	225	+ ExecuteTilingCase(tc);
	226	+}
	227	+
	228	+TEST_F(SwigluGroupTilingTest, tiling_error_invalid_weight_dtype)
	229	+{
	230	+ TilingCase tc;
	231	+ tc.hasWeight = true;
	232	+ tc.weightDtype = ge::DT_FLOAT16;
	233	+ tc.status = ge::GRAPH_FAILED;
	234	+ ExecuteTilingCase(tc);
	235	+}
	236	+
	237	+TEST_F(SwigluGroupTilingTest, tiling_error_invalid_weight_shape)
	238	+{
	239	+ TilingCase tc;
	240	+ tc.hasWeight = true;
	241	+ tc.weightShape = {{1023}, {1023}};
	242	+ tc.status = ge::GRAPH_FAILED;
	243	+ ExecuteTilingCase(tc);
	244	+}
	245	+
	246	+TEST_F(SwigluGroupTilingTest, tiling_error_invalid_group_index_dtype)
	247	+{
	248	+ TilingCase tc;
	249	+ tc.hasGroupIndex = true;
	250	+ tc.groupIndexDtype = ge::DT_INT32;
	251	+ tc.status = ge::GRAPH_FAILED;
	252	+ ExecuteTilingCase(tc);
	253	+}
	254	+
	255	+TEST_F(SwigluGroupTilingTest, tiling_error_zero_clamp_limit)
	256	+{
	257	+ TilingCase tc;
	258	+ tc.clampLimit = 0.0f;
	259	+ tc.status = ge::GRAPH_FAILED;
	260	+ ExecuteTilingCase(tc);
	261	+}
	262	+} // namespace

Aactivation/swiglu_group/tests/ut/op_host/test_swiglu_group_infershape.cpp+127-0

@@ -0,0 +1,127 @@
	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 "exe_graph/runtime/storage_format.h"
	14	+#include "exe_graph/runtime/storage_shape.h"
	15	+#include <gtest/gtest.h>
	16	+#include "infershape_test_util.h"
	17	+#include "ut_op_common.h"
	18	+#include "register/op_impl_registry.h"
	19	+#include "log/log.h"
	20	+#include "../../../op_graph/swiglu_group_proto.h"
	21	+
	22	+namespace {
	23	+const Runtime2TestParam kRuntimeParam{{"clamp_limit"}};
	24	+
	25	+void UpdateInputX(ge::op::SwigluGroup& op, const std::vector<int64_t>& dims, ge::DataType dtype)
	26	+{
	27	+ ge::TensorDesc xDesc;
	28	+ ge::Shape xShape(dims);
	29	+ xDesc.SetDataType(dtype);
	30	+ xDesc.SetShape(xShape);
	31	+ xDesc.SetOriginShape(xShape);
	32	+ op.UpdateInputDesc("x", xDesc);
	33	+}
	34	+
	35	+class SwigluGroupInferShapeTest : public testing::Test {
	36	+protected:
	37	+ static void SetUpTestCase()
	38	+ {
	39	+ std::cout << "SwigluGroupInferShapeTest SetUp" << std::endl;
	40	+ }
	41	+
	42	+ static void TearDownTestCase()
	43	+ {
	44	+ std::cout << "SwigluGroupInferShapeTest TearDown" << std::endl;
	45	+ }
	46	+};
	47	+
	48	+TEST_F(SwigluGroupInferShapeTest, infer_shape_fp16)
	49	+{
	50	+ ge::op::SwigluGroup op;
	51	+ UpdateInputX(op, {8, 128, 8192}, ge::DT_FLOAT16);
	52	+
	53	+ EXPECT_EQ(InferShapeTest(op, kRuntimeParam), ge::GRAPH_SUCCESS);
	54	+ EXPECT_EQ(op.GetOutputDesc(0).GetShape().GetDims(), std::vector<int64_t>({8, 128, 4096}));
	55	+}
	56	+
	57	+TEST_F(SwigluGroupInferShapeTest, infer_shape_bf16)
	58	+{
	59	+ ge::op::SwigluGroup op;
	60	+ UpdateInputX(op, {4, 64, 2048}, ge::DT_BF16);
	61	+
	62	+ EXPECT_EQ(InferShapeTest(op, kRuntimeParam), ge::GRAPH_SUCCESS);
	63	+ EXPECT_EQ(op.GetOutputDesc(0).GetShape().GetDims(), std::vector<int64_t>({4, 64, 1024}));
	64	+}
	65	+
	66	+TEST_F(SwigluGroupInferShapeTest, infer_shape_fp32)
	67	+{
	68	+ ge::op::SwigluGroup op;
	69	+ UpdateInputX(op, {4, 512}, ge::DT_FLOAT);
	70	+
	71	+ EXPECT_EQ(InferShapeTest(op, kRuntimeParam), ge::GRAPH_SUCCESS);
	72	+ EXPECT_EQ(op.GetOutputDesc(0).GetShape().GetDims(), std::vector<int64_t>({4, 256}));
	73	+}
	74	+
	75	+TEST_F(SwigluGroupInferShapeTest, infer_shape_unknown_rank)
	76	+{
	77	+ ge::op::SwigluGroup op;
	78	+ UpdateInputX(op, {-2}, ge::DT_FLOAT16);
	79	+
	80	+ EXPECT_EQ(InferShapeTest(op, kRuntimeParam), ge::GRAPH_SUCCESS);
	81	+ EXPECT_EQ(op.GetOutputDesc(0).GetShape().GetDims(), std::vector<int64_t>({-2}));
	82	+}
	83	+
	84	+TEST_F(SwigluGroupInferShapeTest, infer_shape_unknown_last_dim)
	85	+{
	86	+ ge::op::SwigluGroup op;
	87	+ UpdateInputX(op, {4, 64, -1}, ge::DT_BF16);
	88	+
	89	+ EXPECT_EQ(InferShapeTest(op, kRuntimeParam), ge::GRAPH_SUCCESS);
	90	+ EXPECT_EQ(op.GetOutputDesc(0).GetShape().GetDims(), std::vector<int64_t>({4, 64, -1}));
	91	+}
	92	+
	93	+TEST_F(SwigluGroupInferShapeTest, infer_shape_error_invalid_last_dim)
	94	+{
	95	+ ge::op::SwigluGroup op;
	96	+ UpdateInputX(op, {4, 64, 1023}, ge::DT_FLOAT16);
	97	+
	98	+ EXPECT_EQ(InferShapeTest(op, kRuntimeParam), ge::GRAPH_FAILED);
	99	+}
	100	+
	101	+TEST_F(SwigluGroupInferShapeTest, infer_dtype_fp16)
	102	+{
	103	+ ge::op::SwigluGroup op;
	104	+ UpdateInputX(op, {2, 8, 1024}, ge::DT_FLOAT16);
	105	+
	106	+ EXPECT_EQ(InferDataTypeTest(op, kRuntimeParam), ge::GRAPH_SUCCESS);
	107	+ EXPECT_EQ(op.GetOutputDesc(0).GetDataType(), ge::DT_FLOAT16);
	108	+}
	109	+
	110	+TEST_F(SwigluGroupInferShapeTest, infer_dtype_bf16)
	111	+{
	112	+ ge::op::SwigluGroup op;
	113	+ UpdateInputX(op, {2, 8, 1024}, ge::DT_BF16);
	114	+
	115	+ EXPECT_EQ(InferDataTypeTest(op, kRuntimeParam), ge::GRAPH_SUCCESS);
	116	+ EXPECT_EQ(op.GetOutputDesc(0).GetDataType(), ge::DT_BF16);
	117	+}
	118	+
	119	+TEST_F(SwigluGroupInferShapeTest, infer_dtype_fp32)
	120	+{
	121	+ ge::op::SwigluGroup op;
	122	+ UpdateInputX(op, {2, 8, 1024}, ge::DT_FLOAT);
	123	+
	124	+ EXPECT_EQ(InferDataTypeTest(op, kRuntimeParam), ge::GRAPH_SUCCESS);
	125	+ EXPECT_EQ(op.GetOutputDesc(0).GetDataType(), ge::DT_FLOAT);
	126	+}
	127	+} // namespace

Aactivation/swiglu_group/tests/ut/op_kernel/CMakeLists.txt+13-0

@@ -0,0 +1,13 @@
	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	+if((UT_TEST_ALL OR OP_KERNEL_UT) AND NOT UT_DONE)
	12	+ AddOpTestCase(swiglu_group "ascend950pr_9599" "-DDTYPE_X=half")
	13	+endif()

Aactivation/swiglu_group/tests/ut/op_kernel/test_swiglu_group.cpp+143-0

@@ -0,0 +1,143 @@
	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 <cstdint>
	12	+#include <iostream>
	13	+#include "gtest/gtest.h"
	14	+#include "tikicpulib.h"
	15	+
	16	+extern "C" __global__ __aicore__ void swiglu_group(
	17	+ GM_ADDR x, GM_ADDR weight, GM_ADDR groupIndex, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling);
	18	+
	19	+namespace {
	20	+class SwigluGroupKernelTest : public testing::Test {
	21	+protected:
	22	+ static void SetUpTestCase()
	23	+ {
	24	+ std::cout << "SwigluGroupKernelTest SetUp" << std::endl;
	25	+ }
	26	+
	27	+ static void TearDownTestCase()
	28	+ {
	29	+ std::cout << "SwigluGroupKernelTest TearDown" << std::endl;
	30	+ }
	31	+};
	32	+
	33	+void FillCommonTiling(SwigluGroupTilingData* tilingData, int64_t bs, int64_t d, uint32_t coreNum, float clampLimit)
	34	+{
	35	+ int64_t splitD = d / 2;
	36	+ tilingData->bs = bs;
	37	+ tilingData->d = d;
	38	+ tilingData->splitD = splitD;
	39	+ tilingData->rowOfFormerBlock = 1;
	40	+ tilingData->rowOfTailBlock = 1;
	41	+ tilingData->rowLoopOfFormerBlock = 1;
	42	+ tilingData->rowLoopOfTailBlock = 1;
	43	+ tilingData->rowFactor = 1;
	44	+ tilingData->tailRowFactorOfFormerBlock = 1;
	45	+ tilingData->tailRowFactorOfTailBlock = 1;
	46	+ tilingData->dLoop = 1;
	47	+ tilingData->dFactor = splitD;
	48	+ tilingData->tailDFactor = splitD;
	49	+ tilingData->clampLimit = clampLimit;
	50	+ tilingData->hasClampLimit = (clampLimit > 0.0f) ? 1 : 0;
	51	+ tilingData->g = 0;
	52	+ tilingData->ubSize = 253952;
	53	+ tilingData->gLoop = 0;
	54	+ tilingData->gFactor = 0;
	55	+ tilingData->tailGFactor = 0;
	56	+ tilingData->coreNum = coreNum;
	57	+}
	58	+
	59	+// Run swiglu_group without group_index (optionally with clamp).
	60	+void RunKernelBasic(float clampLimit)
	61	+{
	62	+ constexpr int64_t bs = 2;
	63	+ constexpr int64_t d = 256;
	64	+ constexpr int64_t splitD = d / 2;
	65	+ constexpr uint32_t blockDim = 2;
	66	+
	67	+ uint8_t* x = reinterpret_cast<uint8_t>(AscendC::GmAlloc(bs d * sizeof(half)));
	68	+ uint8_t* y = reinterpret_cast<uint8_t>(AscendC::GmAlloc(bs splitD * sizeof(half)));
	69	+ uint8_t* workspace = reinterpret_cast<uint8_t*>(AscendC::GmAlloc(32));
	70	+ uint8_t* tiling = reinterpret_cast<uint8_t*>(AscendC::GmAlloc(sizeof(SwigluGroupTilingData)));
	71	+
	72	+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
	73	+ FillCommonTiling(reinterpret_cast<SwigluGroupTilingData*>(tiling), bs, d, blockDim, clampLimit);
	74	+
	75	+ ICPU_SET_TILING_KEY(1000);
	76	+ auto kernel = [](GM_ADDR x, GM_ADDR weight, GM_ADDR groupIndex, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) {
	77	+ ::swiglu_group(x, weight, groupIndex, y, workspace, tiling);
	78	+ };
	79	+ ICPU_RUN_KF(kernel, blockDim, x, nullptr, nullptr, y, workspace, tiling);
	80	+
	81	+ AscendC::GmFree(x);
	82	+ AscendC::GmFree(y);
	83	+ AscendC::GmFree(workspace);
	84	+ AscendC::GmFree(tiling);
	85	+}
	86	+
	87	+// Run swiglu_group with a large group_index (g >= 97) to exercise VFProcessGroupIndexLargeVf.
	88	+void RunKernelGroupLarge(int64_t g)
	89	+{
	90	+ constexpr int64_t bs = 4;
	91	+ constexpr int64_t d = 256;
	92	+ constexpr int64_t splitD = d / 2;
	93	+ constexpr uint32_t blockDim = 4;
	94	+
	95	+ uint8_t* x = reinterpret_cast<uint8_t>(AscendC::GmAlloc(bs d * sizeof(half)));
	96	+ uint8_t* y = reinterpret_cast<uint8_t>(AscendC::GmAlloc(bs splitD * sizeof(half)));
	97	+ uint8_t* groupIndex = reinterpret_cast<uint8_t>(AscendC::GmAlloc(g sizeof(int64_t)));
	98	+ uint8_t* workspace = reinterpret_cast<uint8_t*>(AscendC::GmAlloc(32));
	99	+ uint8_t* tiling = reinterpret_cast<uint8_t*>(AscendC::GmAlloc(sizeof(SwigluGroupTilingData)));
	100	+
	101	+ // group_index counts sum to bs so realBs == bs (all rows processed).
	102	+ auto* gi = reinterpret_cast<int64_t*>(groupIndex);
	103	+ for (int64_t i = 0; i < g; i++) {
	104	+ gi[i] = 0;
	105	+ }
	106	+ gi[g - 1] = bs;
	107	+
	108	+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
	109	+ auto* tilingData = reinterpret_cast<SwigluGroupTilingData*>(tiling);
	110	+ FillCommonTiling(tilingData, bs, d, blockDim, 0.0f);
	111	+ tilingData->g = g;
	112	+ tilingData->gLoop = 1;
	113	+ tilingData->gFactor = g;
	114	+ tilingData->tailGFactor = g;
	115	+
	116	+ ICPU_SET_TILING_KEY(1000);
	117	+ auto kernel = [](GM_ADDR x, GM_ADDR weight, GM_ADDR groupIndex, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) {
	118	+ ::swiglu_group(x, weight, groupIndex, y, workspace, tiling);
	119	+ };
	120	+ ICPU_RUN_KF(kernel, blockDim, x, nullptr, groupIndex, y, workspace, tiling);
	121	+
	122	+ AscendC::GmFree(x);
	123	+ AscendC::GmFree(y);
	124	+ AscendC::GmFree(groupIndex);
	125	+ AscendC::GmFree(workspace);
	126	+ AscendC::GmFree(tiling);
	127	+}
	128	+
	129	+TEST_F(SwigluGroupKernelTest, block_fp16)
	130	+{
	131	+ RunKernelBasic(-1.0f);
	132	+}
	133	+
	134	+TEST_F(SwigluGroupKernelTest, block_fp16_clamp)
	135	+{
	136	+ RunKernelBasic(1.0f);
	137	+}
	138	+
	139	+TEST_F(SwigluGroupKernelTest, group_index_large)
	140	+{
	141	+ RunKernelGroupLarge(97);
	142	+}
	143	+} // namespace

Mdocs/zh/op_api_list.md+1-0

	@@ -114,6 +114,7 @@
114	\| [aclnnDequantBias](../../quant/dequant_bias/docs/aclnnDequantBias.md) \| 对输入x反量化操作，将输入的int32的数据转化为FLOAT16/BFLOAT16输出。 \| 默认确定性实现 \| - \|	114	\| [aclnnDequantBias](../../quant/dequant_bias/docs/aclnnDequantBias.md) \| 对输入x反量化操作，将输入的int32的数据转化为FLOAT16/BFLOAT16输出。 \| 默认确定性实现 \| - \|
115	\| [aclnnDequantSwigluQuant](../../quant/dequant_swiglu_quant/docs/aclnnDequantSwigluQuant.md) \| 在Swish门控线性单元激活函数前后添加dequant和quant操作，实现x的DequantSwigluQuant计算。 \| 默认确定性实现 \| 默认确定性实现 \|	115	\| [aclnnDequantSwigluQuant](../../quant/dequant_swiglu_quant/docs/aclnnDequantSwigluQuant.md) \| 在Swish门控线性单元激活函数前后添加dequant和quant操作，实现x的DequantSwigluQuant计算。 \| 默认确定性实现 \| 默认确定性实现 \|
116	\| [aclnnDequantSwigluQuantV2](../../quant/dequant_swiglu_quant/docs/aclnnDequantSwigluQuantV2.md) \| 在Swish门控线性单元激活函数前后添加dequant和quant操作，实现x的DequantSwigluQuant计算。 \| 默认确定性实现 \| 默认确定性实现 \|	116	\| [aclnnDequantSwigluQuantV2](../../quant/dequant_swiglu_quant/docs/aclnnDequantSwigluQuantV2.md) \| 在Swish门控线性单元激活函数前后添加dequant和quant操作，实现x的DequantSwigluQuant计算。 \| 默认确定性实现 \| 默认确定性实现 \|
		117	+\| [aclnnSwigluGroup](../../activation/swiglu_group/docs/aclnnSwigluGroup.md) \| 执行SwiGLU激活，实现x的SwiGlu计算，输出与x相同数据类型的激活结果。 \| - \| 默认确定性实现 \|
117	\| [aclnnSwigluGroupQuant](../../quant/swiglu_group_quant/docs/aclnnSwigluGroupQuant.md) \| 在Swish门控线性单元激活函数后执行分组低比特量化，支持FP8和FP4量化输出。 \| - \| 默认确定性实现 \|	118	\| [aclnnSwigluGroupQuant](../../quant/swiglu_group_quant/docs/aclnnSwigluGroupQuant.md) \| 在Swish门控线性单元激活函数后执行分组低比特量化，支持FP8和FP4量化输出。 \| - \| 默认确定性实现 \|
118	\| [aclnnSwigluMxQuant](../../quant/swiglu_mx_quant/docs/aclnnSwigluMxQuant.md) \| 在Swish门控线性单元激活函数后添加DynamicMxQuant操作，实现x的SwigluMxQuant计算。 \| - \| 默认确定性实现 \|	119	\| [aclnnSwigluMxQuant](../../quant/swiglu_mx_quant/docs/aclnnSwigluMxQuant.md) \| 在Swish门控线性单元激活函数后添加DynamicMxQuant操作，实现x的SwigluMxQuant计算。 \| - \| 默认确定性实现 \|
119	\| [aclnnSwigluMxQuantWithDualAxis](../../quant/swiglu_mx_quant_with_dual_axis/docs/aclnnSwigluMxQuantWithDualAxis.md) \| 在Swish门控线性单元激活函数后添加GroupDynamicMxQuant操作，根据传入的分组索引的起始值，对swiglu后的数据进行分组的-1轴和-2轴同时动态量化。 \| - \| 默认确定性实现 \|	120	\| [aclnnSwigluMxQuantWithDualAxis](../../quant/swiglu_mx_quant_with_dual_axis/docs/aclnnSwigluMxQuantWithDualAxis.md) \| 在Swish门控线性单元激活函数后添加GroupDynamicMxQuant操作，根据传入的分组索引的起始值，对swiglu后的数据进行分组的-1轴和-2轴同时动态量化。 \| - \| 默认确定性实现 \|

Mdocs/zh/op_list.md+10-0

	@@ -687,6 +687,16 @@
687	<td>AI Core</td>	687	<td>AI Core</td>
688	<td>完成aclnnSwiGlu的反向计算，完成x的SwiGlu反向梯度计算。</td>	688	<td>完成aclnnSwiGlu的反向计算，完成x的SwiGlu反向梯度计算。</td>
689	</tr>	689	</tr>
		690	+ <tr>
		691	+ <td>activation</td>
		692	+ <td><a href="../../activation/swiglu_group/README.md">swiglu_group</a></td>
		693	+ <td>✓</td>
		694	+ <td>✓</td>
		695	+ <td>✓</td>
		696	+ <td>✓</td>
		697	+ <td>AI Core</td>
		698	+ <td>执行SwiGLU激活，实现输入x的SwigluGroup计算，输出与x相同数据类型的激活结果。</td>
		699	+ </tr>
690	<tr>	700	<tr>
691	<td>activation</td>	701	<td>activation</td>
692	<td><a href="../../activation/swiglu_group_quant_grad/README.md">swiglu_group_quant_grad</a></td>	702	<td><a href="../../activation/swiglu_group_quant_grad/README.md">swiglu_group_quant_grad</a></td>