swiglu_group_quant正反向文档、用例、校验代码同步-ops-nn-AtomGit

已合并

swiglu_group_quant正反向文档、用例、校验代码同步 #7491

shilulu创建于 7月14日从shilulu/ops-nn_3710： 9.1.0 合入到 cann/ops-nn： 9.1.0

swiglu_group_quant正反向文档、用例、校验代码同步 #7491

已合并

shilulu创建于 7月14日

从shilulu/ops-nn_3710 9.1.0合入到cann/ops-nn 9.1.0

共 11 个文件变更+1600-496

Mactivation/swiglu_group_quant/README.md+387-54

	@@ -1,5 +1,7 @@
1	# SwigluGroupQuant	1	# SwigluGroupQuant
2		2
		3	+[📄 查看源码](https://gitcode.com/cann/ops-nn/tree/master/activation/swiglu_group_quant)
		4	+
3	## 产品支持情况	5	## 产品支持情况
4		6
5	\| 产品 \| 是否支持 \|	7	\| 产品 \| 是否支持 \|
	@@ -13,41 +15,370 @@
13		15
14	## 功能说明	16	## 功能说明
15		17
16	-- 算子功能：融合实现SwiGLU激活和分组低比特量化，支持FP8和FP4输出。输入`x`的最后一维被均分为`A`和`B`，先计算`silu(A) * B`，再执行量化。	18	+### 接口功能
17		19
18	-- 计算公式：	20	+SwigluGroupQuant算子实现SwiGLU激活函数与分组量化融合计算。支持四种量化模式：
		21	+- quant_mode=0: Block Quant（FP8块量化，固定128元素分组）
		22	+- quant_mode=1: MX Quant（FP8 MX量化，固定32元素分组）
		23	+- quant_mode=2: HiFp8 Static Quant（HiFp8静态量化）
		24	+- quant_mode=3: HiFp8 Dynamic Quant（HiFp8动态量化）
19		25
20	- $$
21	- y_{tmp}=silu(A) \times B
22	- $$
23		26
24	- 当传入`clamp_limit`时：	27	+### 计算公式
25		28
26	- $$	29	+#### 基础计算流程
27	- A=min(A, clamp\_limit)
28	- $$
29		30
30	- $$	31	+```
31	- B=min(max(B, -clamp\_limit), clamp\_limit)	32	+步骤〇：GroupIndex处理（可选）→ 计算real_bs
32	- $$	33	+步骤一：输入切分（仅处理前real_bs行）
		34	+步骤二：Clamp处理（可选，仅处理前real_bs行）
		35	+步骤三：SwiGLU激活（仅处理前real_bs行）
		36	+步骤四：Weight加权（可选，仅处理前real_bs行）
		37	+步骤五：量化计算（仅处理前real_bs行）
		38	+```
33		39
34	- 当传入`weight`时，量化前执行：	40	+#### 步骤〇：GroupIndex处理（可选）
35		41
36	- $$	42	+当提供 `group_index` 时，用于动态计算实际处理的token数量：
37	- y_{tmp}=y_{tmp} \times weight
38	- $$
39		43
40	- 进行量化：	44	+$$
		45	+\text{group\_sum} = \sum_{g=0}^{G-1} \text{group\_index}[g]
		46	+$$
41		47
42	- $$	48	+$$
43	- scale=row\_max(abs(y_{tmp}))/dstTypeScale	49	+\text{real\_bs} = \min(\text{group\_sum}, N)
44	- $$	50	+$$
45		51
46	- $$	52	+其中：
47	- y = Cast(Mul(y_{tmp}, 1/scale))	53	+- $G$ 为MoE专家分组数
48	- $$	54	+- $N$ 为输入张量的第一维（预设batch size）
		55	+- 后续所有步骤仅处理前 $\text{real\_bs}$ 行数据
49		56
50	- quant_mode为0时输出FP8类型的y和FLOAT32类型的y_scale；quant_mode为1时输出FP8/FP4类型的y和FLOAT8_E8M0类型的y_scale；quant_mode为2、3时输出HIFP8类型的y和FLOAT32类型的y_scale。	57	+MoE场景说明：在MoE推理中，不同专家可能处理不同数量的token，group_index允许动态调整处理范围，避免处理空数据。
		58	+
		59	+#### 步骤一：输入切分
		60	+
		61	+输入张量 $\mathbf{x} \in \mathbb{R}^{N \times D}$ 沿最后一维切分为两部分：
		62	+
		63	+$$
		64	+\mathbf{x}_0[n, d] = \mathbf{x}[n, d], \quad d \in [0, D/2)
		65	+$$
		66	+
		67	+$$
		68	+\mathbf{x}_1[n, d] = \mathbf{x}[n, d + D/2], \quad d \in [0, D/2)
		69	+$$
		70	+
		71	+#### 步骤二：Clamp处理（可选）
		72	+
		73	+当 `clamp_limit > 0` 时，对输入进行限制：
		74	+
		75	+$$
		76	+\mathbf{x}_0'[n, d] = \min(\mathbf{x}_0[n, d], c)
		77	+$$
		78	+
		79	+$$
		80	+\mathbf{x}_1'[n, d] = \min(\max(\mathbf{x}_1[n, d], -c), c)
		81	+$$
		82	+
		83	+其中 $c$ 为 `clamp_limit`。
		84	+
		85	+Clamp的作用：
		86	+- $\mathbf{x}_0$（门控分支）限制为正值范围 $[0, c]$，防止sigmoid梯度消失
		87	+- $\mathbf{x}_1$（线性分支）限制为对称范围 $[-c, c]$，防止数值溢出
		88	+
		89	+#### 步骤三：SwiGLU激活
		90	+
		91	+SwiGLU激活函数定义（逐元素计算）：
		92	+
		93	+$$
		94	+\mathbf{y}_{\text{swiglu}}[n, d] = \text{Swish}(\mathbf{x}_0'[n, d]) \cdot \mathbf{x}_1'[n, d]
		95	+$$
		96	+
		97	+其中Swish函数：
		98	+
		99	+$$
		100	+\text{Swish}(z) = z \cdot \sigma(z) = z \cdot \frac{1}{1 + e^{-z}}
		101	+$$
		102	+
		103	+完整计算步骤分解：
		104	+
		105	+$$
		106	+\begin{aligned}
		107	+t_1[n, d] &= -\mathbf{x}_0'[n, d] \quad \text{(neg)} \\
		108	+t_2[n, d] &= e^{t_1[n, d]} = e^{-\mathbf{x}_0'[n, d]} \quad \text{(exp)} \\
		109	+t_3[n, d] &= t_2[n, d] + 1 = 1 + e^{-\mathbf{x}_0'[n, d]} \quad \text{(add)} \\
		110	+t_4[n, d] &= \frac{\mathbf{x}_0'[n, d]}{t_3[n, d]} = \text{Swish}(\mathbf{x}_0'[n, d]) \quad \text{(div)} \\
		111	+\mathbf{y}_{\text{swiglu}}[n, d] &= t_4[n, d] \cdot \mathbf{x}_1'[n, d] \quad \text{(mul)}
		112	+\end{aligned}
		113	+$$
		114	+
		115	+#### 步骤四：Weight加权（可选）
		116	+
		117	+当提供 `weight` 时，对SwiGLU输出进行加权：
		118	+
		119	+$$
		120	+\mathbf{y}_{\text{weighted}}[n, d] = \mathbf{y}_{\text{swiglu}}[n, d] \cdot w[n]
		121	+$$
		122	+
		123	+其中 $w[n]$ 为第 $n$ 个token的weight值。
		124	+
		125	+MoE场景：weight来自专家路由器的softmax输出，表示该token对当前专家的权重。
		126	+
		127	+#### 步骤五：量化计算
		128	+
		129	+---
		130	+
		131	+#### quant_mode=0 (Block Quant)
		132	+
		133	+分组划分：将输出沿最后一维按128元素为一组划分：
		134	+
		135	+$$
		136	+\mathbf{y} = [\mathbf{g}_0, \mathbf{g}_1, \ldots, \mathbf{g}_K], \quad K = \lceil D/2 / 128 \rceil
		137	+$$
		138	+
		139	+每个组 $\mathbf{g}_i \in \mathbb{R}^{N \times 128}$。
		140	+
		141	+非有限值屏蔽与绝对值计算：
		142	+
		143	+$$
		144	+\begin{aligned}
		145	+\mathbf{z}[n, j] &= \mathbf{y}_{\text{weighted}}[n, j] \cdot 0 \quad \text{(生成零张量)} \\
		146	+\mathbf{m}_{\text{finite}}[n, j] &= (\mathbf{z}[n, j] = \mathbf{z}[n, j]) \\
		147	+\mathbf{y}_{\text{abs}}[n, j] &=
		148	+\begin{cases}
		149	+\|\mathbf{y}_{\text{weighted}}[n, j]\|, & \mathbf{m}_{\text{finite}}[n, j] \\
		150	+0, & \text{otherwise}
		151	+\end{cases}
		152	+\end{aligned}
		153	+$$
		154	+
		155	+屏蔽原理：NaN的特性是 `NaN != NaN`；同时 `Inf * 0` 也会得到NaN，因此该步骤在计算amax时屏蔽NaN和Inf。
		156	+
		157	+Scale计算：
		158	+
		159	+对于第 $i$ 个组（包含128个连续元素）：
		160	+
		161	+$$
		162	+a_i = \max_{j=0}^{127} \mathbf{y}_{\text{abs}}[j]
		163	+$$
		164	+
		165	+$$
		166	+\hat{a}_i = \max(a_i, 10^{-4})
		167	+$$
		168	+
		169	+$$
		170	+s_i^{\text{raw}} = \frac{\hat{a}_i}{M_{\text{fp8}}}
		171	+$$
		172	+
		173	+其中 $M_{\text{fp8}}$ 取值：
		174	+- FP8 E4M3FN：$M_{\text{fp8}} = 448.0$
		175	+- FP8 E5M2：$M_{\text{fp8}} = 57344.0$
		176	+
		177	+Scale输出与InvScale计算：
		178	+
		179	+当 `round_scale=false` 时：
		180	+
		181	+$$
		182	+s_i = s_i^{\text{raw}}, \quad \text{InvScale}_i = \frac{M_{\text{fp8}}}{\hat{a}_i} = \frac{1}{s_i}
		183	+$$
		184	+
		185	+当 `round_scale=true` 时，将scale向上取整到2的幂：
		186	+
		187	+$$
		188	+e_i = \lceil \log_2(s_i^{\text{raw}}) \rceil
		189	+$$
		190	+
		191	+$$
		192	+s_i = 2^{e_i}, \quad \text{InvScale}_i = 2^{-e_i}
		193	+$$
		194	+
		195	+其中 $s_i$ 写入FLOAT32类型的scale输出。
		196	+
		197	+量化计算：
		198	+
		199	+$$
		200	+\mathbf{y}_{\text{scaled}}[n, j] = \mathbf{y}_{\text{weighted}}[n, j] \cdot \text{InvScale}_i, \quad j \in \text{group } i
		201	+$$
		202	+
		203	+若 $\mathbf{y}_{\text{scaled}}[n,j]$ 为NaN或Inf，实现会使用原始 $\mathbf{y}_{\text{weighted}}[n,j]$ 作为FP8 cast输入：
		204	+
		205	+$$
		206	+\mathbf{y}_{\text{cast\_in}}[n, j] =
		207	+\begin{cases}
		208	+\mathbf{y}_{\text{scaled}}[n, j], & \mathbf{y}_{\text{scaled}}[n, j] \text{ is finite} \\
		209	+\mathbf{y}_{\text{weighted}}[n, j], & \text{otherwise}
		210	+\end{cases}
		211	+$$
		212	+
		213	+$$
		214	+\mathbf{y}_{\text{quant}}[n, j] = \text{cast\_fp8\_rint}(\mathbf{y}_{\text{cast\_in}}[n, j])
		215	+$$
		216	+
		217	+其中 `cast_fp8_rint` 为FP32到FP8的类型转换，采用RINT（就近舍入）模式。
		218	+
		219	+---
		220	+
		221	+#### quant_mode=1 (MX Quant)
		222	+
		223	+MX量化原理：采用E8M0 Scale + FP8 Data的组合。
		224	+
		225	+分组方式：每32元素为一组：
		226	+
		227	+$$
		228	+\mathbf{y} = [\mathbf{g}_0, \mathbf{g}_1, \ldots, \mathbf{g}_K], \quad \mathbf{g}_i \in \mathbb{R}^{32}
		229	+$$
		230	+
		231	+Amax计算：
		232	+
		233	+$$
		234	+a_i = \max_{j=0}^{31} \|\mathbf{g}_i[j]\|
		235	+$$
		236	+
		237	+$$
		238	+\hat{a}_i = \max(a_i, 10^{-4})
		239	+$$
		240	+
		241	+原始Scale计算：
		242	+
		243	+$$
		244	+s_i^{\text{raw}} = \frac{\hat{a}_i}{M_{\text{fp8}}}
		245	+$$
		246	+
		247	+其中 $M_{\text{fp8}}$ 取值：
		248	+- FP8 E4M3FN：$M_{\text{fp8}} = 448.0$
		249	+- FP8 E5M2：$M_{\text{fp8}} = 57344.0$
		250	+
		251	+quant_mode=1仅支持 `round_scale=true`，将原始scale向上取整到2的幂：
		252	+
		253	+$$
		254	+e_i = \lceil \log_2(s_i^{\text{raw}}) \rceil
		255	+$$
		256	+
		257	+等价于基于FP32位模式计算：
		258	+
		259	+$$
		260	+e_i = E(s_i^{\text{raw}}) - 127 + \mathbf{1}_{\text{mantissa}(s_i^{\text{raw}}) \ne 0}
		261	+$$
		262	+
		263	+E8M0 Scale编码：
		264	+
		265	+$$
		266	+s_i^{\text{e8m0}} = e_i + 127
		267	+$$
		268	+
		269	+其中 $s_i^{\text{e8m0}}$ 写入FLOAT8_E8M0类型的scale输出，表示的实际scale值为 $2^{e_i}$。
		270	+
		271	+InvScale计算：
		272	+
		273	+$$
		274	+\text{InvScale}_i = 2^{-e_i}
		275	+$$
		276	+
		277	+量化计算：
		278	+
		279	+$$
		280	+\mathbf{y}_{\text{quant}}[j] = \text{cast\_fp8\_rint}\left(\mathbf{y}_{\text{weighted}}[j] \cdot \text{InvScale}_i\right), \quad j \in \text{group } i
		281	+$$
		282	+
		283	+---
		284	+
		285	+#### quant_mode=2 (HiFp8 Static Quant)
		286	+
		287	+静态量化说明：使用预先提供的 `invScale` 对加权后的SwiGLU输出进行缩放量化。
		288	+
		289	+情况1：无GroupIndex（groupIndex为空）：
		290	+
		291	+$$
		292	+\mathbf{y}_{\text{quant}}[n, d] = \text{hif8\_cast}\left(\mathbf{y}_{\text{weighted}}[n, d] \cdot \text{invScale}[0]\right), \quad n \in [0, N), \quad d \in [0, D/2)
		293	+$$
		294	+
		295	+其中 `hif8_cast` 为HiFloat8类型转换函数。
		296	+
		297	+情况2：有GroupIndex（groupIndex非空）：
		298	+
		299	+设 $G$ 为MoE专家分组数，$\text{groupIndex}[g]$ 表示第 $g$ 个专家处理的token数量。
		300	+
		301	+计算每个group的起止索引：
		302	+
		303	+$$
		304	+\text{start}^{(0)} = 0, \quad \text{end}^{(g)} = \sum_{k=0}^{g} \text{groupIndex}[k], \quad \text{start}^{(g)} = \text{end}^{(g-1)}
		305	+$$
		306	+
		307	+对于第 $g$ 个group，使用对应的缩放因子 $\text{invScale}[g]$ 进行量化：
		308	+
		309	+$$
		310	+\mathbf{y}_{\text{quant}}[n, d] = \text{hif8\_cast}\left(\mathbf{y}_{\text{weighted}}[n, d] \cdot \text{invScale}[g]\right), \quad n \in [\text{start}^{(g)}, \text{end}^{(g)}), \quad d \in [0, D/2)
		311	+$$
		312	+
		313	+MoE场景说明：在MoE推理中，不同专家处理不同数量的token，groupIndex用于标识每个专家处理的token范围，invScale为每个专家预先计算的静态缩放因子。
		314	+
		315	+---
		316	+
		317	+#### quant_mode=3 (HiFp8 Dynamic Quant)
		318	+
		319	+动态量化说明：根据加权后的SwiGLU输出动态计算缩放因子进行量化。
		320	+
		321	+情况1：无GroupIndex（groupIndex为空）：
		322	+
		323	+计算全局绝对值最大值：
		324	+
		325	+$$
		326	+a_{\max} = \max\left(\max_{n \in [0, N), d \in [0, D/2)} \|\mathbf{y}_{\text{weighted}}[n, d]\|, \epsilon\right)
		327	+$$
		328	+
		329	+其中 $\epsilon$ 为数值稳定性常数。
		330	+
		331	+计算缩放因子：
		332	+
		333	+$$
		334	+s = \frac{a_{\max}}{M_{\text{hif8}}}
		335	+$$
		336	+
		337	+其中 $M_{\text{hif8}}$ 为 `dstTypeMax`，表示HiFloat8类型的最大有限值。
		338	+
		339	+量化计算：
		340	+
		341	+$$
		342	+\mathbf{y}_{\text{quant}}[n, d] = \text{hif8\_cast}\left(\frac{\mathbf{y}_{\text{weighted}}[n, d]}{s}\right), \quad n \in [0, N), \quad d \in [0, D/2)
		343	+$$
		344	+
		345	+其中 `hif8_cast` 为HiFloat8类型转换函数。
		346	+
		347	+情况2：有GroupIndex（groupIndex非空）：
		348	+
		349	+设 $G$ 为MoE专家分组数，$\text{groupIndex}[g]$ 表示第 $g$ 个专家处理的token数量。
		350	+
		351	+计算每个group的起止索引：
		352	+
		353	+$$
		354	+\text{start}^{(0)} = 0, \quad \text{end}^{(g)} = \sum_{k=0}^{g} \text{groupIndex}[k], \quad \text{start}^{(g)} = \text{end}^{(g-1)}
		355	+$$
		356	+
		357	+对于第 $g$ 个group，提取对应的数据：
		358	+
		359	+$$
		360	+\mathbf{y}^{(g)} = \mathbf{y}_{\text{weighted}}[\text{start}^{(g)}:\text{end}^{(g)}, :]
		361	+$$
		362	+
		363	+计算该group的绝对值最大值：
		364	+
		365	+$$
		366	+a_{\max}^{(g)} = \max\left(\max_{n \in [\text{start}^{(g)}, \text{end}^{(g)}), d \in [0, D/2)} \|\mathbf{y}^{(g)}[n, d]\|, \epsilon\right)
		367	+$$
		368	+
		369	+计算该group的缩放因子：
		370	+
		371	+$$
		372	+s^{(g)} = \frac{a_{\max}^{(g)}}{M_{\text{hif8}}}
		373	+$$
		374	+
		375	+对该group进行量化：
		376	+
		377	+$$
		378	+\mathbf{y}_{\text{quant}}[n, d] = \text{hif8\_cast}\left(\frac{\mathbf{y}_{\text{weighted}}[n, d]}{s^{(g)}}\right), \quad n \in [\text{start}^{(g)}, \text{end}^{(g)}), \quad d \in [0, D/2)
		379	+$$
		380	+
		381	+MoE场景说明：在MoE推理中，不同专家处理不同数量的token，groupIndex用于标识每个专家处理的token范围，每个group独立计算缩放因子以适应不同数据分布。
51		382
52	## 参数说明	383	## 参数说明
53		384
	@@ -70,116 +401,118 @@
70	<tr>	401	<tr>
71	<td>x</td>	402	<td>x</td>
72	<td>输入</td>	403	<td>输入</td>
73	- <td>输入张量，shape为[...,D]，最后一维D按左右两半做SwiGLU。</td>	404	+ <td>SwiGLU输入。shape为[...,D]，D必须大于等于256，且能被256整除。不支持空Tensor。</td>
74	<td>FLOAT、FLOAT16、BFLOAT16</td>	405	<td>FLOAT、FLOAT16、BFLOAT16</td>
75	<td>ND</td>	406	<td>ND</td>
76	</tr>	407	</tr>
77	<tr>	408	<tr>
78	<td>weight</td>	409	<td>weight</td>
79	- <td>可选输入</td>	410	+ <td>输入（可选）</td>
80	- <td>量化前按token乘到SwiGLU输出上的权重。</td>	411	+ <td>MOE权重张量，用于SwiGLU输出的加权计算。不支持空Tensor。不为空时，数据类型为FLOAT32，元素个数需等于x除最后一维外的元素个数之积。</td>
81	<td>FLOAT32</td>	412	<td>FLOAT32</td>
82	<td>ND</td>	413	<td>ND</td>
83	</tr>	414	</tr>
84	<tr>	415	<tr>
85	<td>group_index</td>	416	<td>group_index</td>
86	- <td>可选输入</td>	417	+ <td>输入（可选）</td>
87	- <td>count模式的group token数。</td>	418	+ <td>count模式的group token数。不支持空Tensor。不为空时，数据类型为INT64，shape为[G]。</td>
88	<td>INT64</td>	419	<td>INT64</td>
89	<td>ND</td>	420	<td>ND</td>
90	</tr>	421	</tr>
91	<tr>	422	<tr>
92	<td>scale</td>	423	<td>scale</td>
93	- <td>可选输入</td>	424	+ <td>输入（可选）</td>
94	- <td>静态量化输入的scale张量。仅quant_mode=2时使用，quant_mode=3时不使用。</td>	425	+ <td>quantMode=2时静态量化输入的invScale张量。仅quant_mode为2时使用。groupIndex存在的话，shape=[G]，不存在的话shape=[1]。</td>
95	<td>FLOAT32</td>	426	<td>FLOAT32</td>
96	<td>ND</td>	427	<td>ND</td>
97	</tr>	428	</tr>
98	<tr>	429	<tr>
99	<td>dst_type</td>	430	<td>dst_type</td>
100	<td>属性</td>	431	<td>属性</td>
101	- <td>目标量化类型：27=HIFLOAT8，35=FLOAT8_E5M2，36=FLOAT8_E4M3FN，40=FLOAT4_E2M1，41=FLOAT4_E1M2。</td>	432	+ <td>目标量化类型。仅quantMode为0或1时，该参数生效。支持取值35、36、40、41，分别表示FLOAT8_E5M2、FLOAT8_E4M3FN、FLOAT4_E2M1、FLOAT4_E1M2。dstType为40或41时，quantMode必须为1。</td>
102	<td>INT64</td>	433	<td>INT64</td>
103	<td>-</td>	434	<td>-</td>
104	</tr>	435	</tr>
105	<tr>	436	<tr>
106	<td>quant_mode</td>	437	<td>quant_mode</td>
107	<td>属性</td>	438	<td>属性</td>
108	- <td>量化模式。0表示Block FP8，1表示MX，2表示HIFP8静态量化，3表示HIFP8动态量化。</td>	439	+ <td>量化模式。支持取值0、1、2、3。0表示Block FP8模式。1表示MX模式。2表示HIFP8静态量化模式。3表示HIFP8动态量化模式。</td>
109	<td>INT64</td>	440	<td>INT64</td>
110	<td>-</td>	441	<td>-</td>
111	</tr>	442	</tr>
112	<tr>	443	<tr>
113	<td>block_size</td>	444	<td>block_size</td>
114	<td>属性</td>	445	<td>属性</td>
115	- <td>量化块大小。0表示使用默认值；Block FP8支持128，MX支持32；quant_mode=2或3时不生效。</td>	446	+ <td>量化块大小。0表示使用当前量化模式的默认block大小。quantMode为0时，支持0或128。quantMode为1时，支持0或32。quantMode为2或3时，该参数不生效，默认0。</td>
116	<td>INT64</td>	447	<td>INT64</td>
117	<td>-</td>	448	<td>-</td>
118	</tr>	449	</tr>
119	<tr>	450	<tr>
120	<td>round_scale</td>	451	<td>round_scale</td>
121	<td>属性</td>	452	<td>属性</td>
122	- <td>是否将scale取整为2的幂。MX模式必须为true；quant_mode=2或3时不生效。</td>	453	+ <td>是否将scale取整为2的幂。quantMode为1时，roundScale必须为true。quantMode为2或3时，该参数不生效。</td>
123	<td>BOOL</td>	454	<td>BOOL</td>
124	<td>-</td>	455	<td>-</td>
125	</tr>	456	</tr>
126	<tr>	457	<tr>
127	<td>clamp_limit</td>	458	<td>clamp_limit</td>
128	<td>属性</td>	459	<td>属性</td>
129	- <td>SwiGLU计算前的clamp阈值。默认不启用clamp。</td>	460	+ <td>SwiGLU计算前的clamp阈值。-1.0表示不启用clamp。启用clamp时，clampLimit必须大于0。</td>
130	<td>FLOAT</td>	461	<td>FLOAT</td>
131	<td>-</td>	462	<td>-</td>
132	</tr>	463	</tr>
133	<tr>	464	<tr>
134	<td>dst_type_max</td>	465	<td>dst_type_max</td>
135	<td>属性</td>	466	<td>属性</td>
136	- <td>目标量化类型的最大有限值。quant_mode=2或3时用于计算scale = amax / dst_type_max，默认值为15.0。</td>	467	+ <td>目标量化类型的最大有限值。仅quantMode为3时，该参数生效。默认值为15.0。</td>
137	<td>FLOAT</td>	468	<td>FLOAT</td>
138	<td>-</td>	469	<td>-</td>
139	</tr>	470	</tr>
140	<tr>	471	<tr>
141	<td>output_origin</td>	472	<td>output_origin</td>
142	<td>属性</td>	473	<td>属性</td>
143	- <td>是否输出量化前的SwiGLU结果。MX FP4模式下该输出仅作占位。</td>	474	+ <td>是否输出量化前的SwiGLU结果。true表示支持输出原始激活值yOrigin，false表示不支持输出原始激活值yOrigin。quantMode为0或1时支持false。quantMode为2或3时true/false都支持。</td>
144	<td>BOOL</td>	475	<td>BOOL</td>
145	<td>-</td>	476	<td>-</td>
146	</tr>	477	</tr>
147	<tr>	478	<tr>
148	<td>y</td>	479	<td>y</td>
149	<td>输出</td>	480	<td>输出</td>
150	- <td>量化输出。FP8 shape为[...,D/2]；FP4每字节打包2个4-bit值，shape为[...,D/4]；HIFP8 shape为[...,D/2]。</td>	481	+ <td>量化输出。quantMode为0或1时，数据类型需与dstType一致，dstType为35或36时，shape为[...,D/2]；dstType为40或41时，shape为[...,D/4]；dstType为27时，shape为[...,D/2]。quantMode为2或3时，数据类型默认为HIFLOAT8，shape为[...,D/2]。不支持空Tensor。</td>
151	- <td>HIFLOAT8、FLOAT8_E4M3FN、FLOAT8_E5M2、FLOAT4_E2M1、FLOAT4_E1M2</td>	482	+ <td>HIFLOAT8、FLOAT8_E5M2、FLOAT8_E4M3FN、FLOAT4_E2M1、FLOAT4_E1M2</td>
152	<td>ND</td>	483	<td>ND</td>
153	</tr>	484	</tr>
154	<tr>	485	<tr>
155	<td>y_scale</td>	486	<td>y_scale</td>
156	<td>输出</td>	487	<td>输出</td>
157	- <td>量化scale。Block FP8输出FLOAT32，shape为[...,ceil((D/2)/128)]；MX输出FLOAT8_E8M0，shape为[...,ceil(ceil((D/2)/32)/2),2]；HIFP8输出FLOAT32，无group_index时shape为[1]，有group_index时shape为[G]。</td>	488	+ <td>量化scale输出。quantMode为0时，shape为[...,ceil((D/2)/128)]，数据类型为FLOAT32。quantMode为1时，shape为[...,ceil(ceil((D/2)/32)/2),2]，数据类型为FLOAT8_E8M0。quantMode为2或3时，无groupIndex时shape为[1]，有groupIndex时shape为[G]，数据类型为FLOAT32。不支持空Tensor。</td>
158	<td>FLOAT32、FLOAT8_E8M0</td>	489	<td>FLOAT32、FLOAT8_E8M0</td>
159	<td>ND</td>	490	<td>ND</td>
160	</tr>	491	</tr>
161	<tr>	492	<tr>
162	<td>y_origin</td>	493	<td>y_origin</td>
163	<td>输出</td>	494	<td>输出</td>
164	- <td>量化前的SwiGLU结果，shape为[...,D/2]。</td>	495	+ <td>量化前的SwiGLU结果。shape为[...,D/2]。数据类型需与x一致。不支持空指针。</td>
165	- <td>与x相同</td>	496	+ <td>FLOAT、FLOAT16、BFLOAT16</td>
166	<td>ND</td>	497	<td>ND</td>
167	</tr>	498	</tr>
168	</tbody></table>	499	</tbody></table>
169		500
170	## 约束说明	501	## 约束说明
171		502
172	-- 输入`x`的rank必须大于0，最后一维`D`必须大于等于256且能被256整除。	503	+- 确定性计算：aclnnSwigluGroupQuant默认确定性实现。
173	-- `dst_type`支持`HIFLOAT8`、`FLOAT8_E4M3FN`、`FLOAT8_E5M2`、`FLOAT4_E2M1`、`FLOAT4_E1M2`。	504	+- quantMode为0时，仅支持FP8输出，blockSize支持0或128。
174	-- `quant_mode=0`时仅支持FP8输出，`block_size`支持0或128。	505	+- quantMode为1时，支持FP8/FP4输出，blockSize支持0或32，roundScale必须为true。
175	-- `quant_mode=1`时支持FP8/FP4输出，`block_size`支持0或32，`round_scale`必须为true。	506	+- quantMode为2或3时，支持HIFP8量化输出，dstType, blockSize和roundScale不生效。输入x的维度为[T, D]或[B, S, D]，需满足以下规格约束：
176	-- `quant_mode=2`时支持HIFP8静态量化输出，`dst_type`、`block_size`和`round_scale`不生效。	507	+
177	-- `quant_mode=3`时支持HIFP8动态量化输出，`dst_type`、`block_size`和`round_scale`不生效。	508	+ \| 规格项 \| 规格 \| 规格说明 \|
178	-- `dst_type`为`FLOAT4_E2M1`或`FLOAT4_E1M2`时，必须使用`quant_mode=1`。	509	+ \| :--- \| :--- \| :--- \|
179	-- `dst_type`为`HIFLOAT8`时，必须使用`quant_mode=2`或`3`。	510	+ \| B \| 1~31 \| - \|
180	-- `y_scale`的数据类型必须与`quant_mode`匹配：Block FP8为FLOAT32，MX为FLOAT8_E8M0，HIFP8为FLOAT32。	511	+ \| S \| 0~128K \| - \|
181	-- `quant_mode=2`或`3`时，`group_index`可用于MoE场景的分组量化，`y_scale`的shape为[G]（G为group数量）。	512	+ \| D/2 \| 512, 768, 1024, 1536, 1792, 2048, 2560, 4096 \| - \|
182	-- `clamp_limit`不启用时使用默认占位值-1.0；启用时必须大于0。	513	+ \| dstTypeMax \| 15, 56, 224, 32768 \| - \|
		514	+- dstType为FLOAT4_E2M1或FLOAT4_E1M2时，必须使用quantMode=1。
		515	+- yScale的数据类型必须与quantMode匹配：Block FP8为FLOAT32，MX为FLOAT8_E8M0，HIFP8为FLOAT32。
183		516
184	## 调用说明	517	## 调用说明
185		518

Mactivation/swiglu_group_quant/docs/aclnnSwigluGroupQuant.md+619-188

	@@ -1,4 +1,4 @@
1	-# aclnnSwigluGroupQuant	1	+# aclnnSwigluGroupQuant
2		2
3	[📄 查看源码](https://gitcode.com/cann/ops-nn/tree/master/activation/swiglu_group_quant)	3	[📄 查看源码](https://gitcode.com/cann/ops-nn/tree/master/activation/swiglu_group_quant)
4		4
	@@ -15,31 +15,370 @@
15		15
16	## 功能说明	16	## 功能说明
17		17
18	-- 接口功能：在SwiGLU激活后执行分组低比特量化，支持FP8和FP4输出。	18	+### 接口功能
19	-- 计算公式：
20	- 令输入x的最后一维为D，左半部分为A，右半部分为B，计算：
21		19
22	- $$	20	+SwigluGroupQuant算子实现SwiGLU激活函数与分组量化融合计算。支持四种量化模式：
23	- y_{tmp}=silu(A) \times B	21	+- quant_mode=0: Block Quant（FP8块量化，固定128元素分组）
24	- $$	22	+- quant_mode=1: MX Quant（FP8 MX量化，固定32元素分组）
		23	+- quant_mode=2: HiFp8 Static Quant（HiFp8静态量化）
		24	+- quant_mode=3: HiFp8 Dynamic Quant（HiFp8动态量化）
25		25
26	- 若传入weight，则量化前执行：
27		26
28	- $$	27	+### 计算公式
29	- y_{tmp}=y_{tmp} \times weight
30	- $$
31		28
32	- 进行量化：	29	+#### 基础计算流程
33		30
34	- $$	31	+```
35	- scale=row\_max(abs(y_{tmp}))/dstTypeScale	32	+步骤〇：GroupIndex处理（可选）→ 计算real_bs
36	- $$	33	+步骤一：输入切分（仅处理前real_bs行）
		34	+步骤二：Clamp处理（可选，仅处理前real_bs行）
		35	+步骤三：SwiGLU激活（仅处理前real_bs行）
		36	+步骤四：Weight加权（可选，仅处理前real_bs行）
		37	+步骤五：量化计算（仅处理前real_bs行）
		38	+```
37		39
38	- $$	40	+#### 步骤〇：GroupIndex处理（可选）
39	- y = Cast(Mul(y_{tmp}, 1/scale))
40	- $$
41		41
42	- quant_mode为0时输出FP8类型的yOut和FLOAT32类型的yScaleOut；quant_mode为1时输出FP8/FP4类型的yOut和FLOAT8_E8M0类型的yScaleOut；quant_mode为2、3时输出HIFP8类型的yOut和FLOAT32类型的yScaleOut。	42	+当提供 `group_index` 时，用于动态计算实际处理的token数量：
		43	+
		44	+$$
		45	+\text{group\_sum} = \sum_{g=0}^{G-1} \text{group\_index}[g]
		46	+$$
		47	+
		48	+$$
		49	+\text{real\_bs} = \min(\text{group\_sum}, N)
		50	+$$
		51	+
		52	+其中：
		53	+- $G$ 为MoE专家分组数
		54	+- $N$ 为输入张量的第一维（预设batch size）
		55	+- 后续所有步骤仅处理前 $\text{real\_bs}$ 行数据
		56	+
		57	+MoE场景说明：在MoE推理中，不同专家可能处理不同数量的token，group_index允许动态调整处理范围，避免处理空数据。
		58	+
		59	+#### 步骤一：输入切分
		60	+
		61	+输入张量 $\mathbf{x} \in \mathbb{R}^{N \times D}$ 沿最后一维切分为两部分：
		62	+
		63	+$$
		64	+\mathbf{x}_0[n, d] = \mathbf{x}[n, d], \quad d \in [0, D/2)
		65	+$$
		66	+
		67	+$$
		68	+\mathbf{x}_1[n, d] = \mathbf{x}[n, d + D/2], \quad d \in [0, D/2)
		69	+$$
		70	+
		71	+#### 步骤二：Clamp处理（可选）
		72	+
		73	+当 `clamp_limit > 0` 时，对输入进行限制：
		74	+
		75	+$$
		76	+\mathbf{x}_0'[n, d] = \min(\mathbf{x}_0[n, d], c)
		77	+$$
		78	+
		79	+$$
		80	+\mathbf{x}_1'[n, d] = \min(\max(\mathbf{x}_1[n, d], -c), c)
		81	+$$
		82	+
		83	+其中 $c$ 为 `clamp_limit`。
		84	+
		85	+Clamp的作用：
		86	+- $\mathbf{x}_0$（门控分支）限制为正值范围 $[0, c]$，防止sigmoid梯度消失
		87	+- $\mathbf{x}_1$（线性分支）限制为对称范围 $[-c, c]$，防止数值溢出
		88	+
		89	+#### 步骤三：SwiGLU激活
		90	+
		91	+SwiGLU激活函数定义（逐元素计算）：
		92	+
		93	+$$
		94	+\mathbf{y}_{\text{swiglu}}[n, d] = \text{Swish}(\mathbf{x}_0'[n, d]) \cdot \mathbf{x}_1'[n, d]
		95	+$$
		96	+
		97	+其中Swish函数：
		98	+
		99	+$$
		100	+\text{Swish}(z) = z \cdot \sigma(z) = z \cdot \frac{1}{1 + e^{-z}}
		101	+$$
		102	+
		103	+完整计算步骤分解：
		104	+
		105	+$$
		106	+\begin{aligned}
		107	+t_1[n, d] &= -\mathbf{x}_0'[n, d] \quad \text{(neg)} \\
		108	+t_2[n, d] &= e^{t_1[n, d]} = e^{-\mathbf{x}_0'[n, d]} \quad \text{(exp)} \\
		109	+t_3[n, d] &= t_2[n, d] + 1 = 1 + e^{-\mathbf{x}_0'[n, d]} \quad \text{(add)} \\
		110	+t_4[n, d] &= \frac{\mathbf{x}_0'[n, d]}{t_3[n, d]} = \text{Swish}(\mathbf{x}_0'[n, d]) \quad \text{(div)} \\
		111	+\mathbf{y}_{\text{swiglu}}[n, d] &= t_4[n, d] \cdot \mathbf{x}_1'[n, d] \quad \text{(mul)}
		112	+\end{aligned}
		113	+$$
		114	+
		115	+#### 步骤四：Weight加权（可选）
		116	+
		117	+当提供 `weight` 时，对SwiGLU输出进行加权：
		118	+
		119	+$$
		120	+\mathbf{y}_{\text{weighted}}[n, d] = \mathbf{y}_{\text{swiglu}}[n, d] \cdot w[n]
		121	+$$
		122	+
		123	+其中 $w[n]$ 为第 $n$ 个token的weight值。
		124	+
		125	+MoE场景：weight来自专家路由器的softmax输出，表示该token对当前专家的权重。
		126	+
		127	+#### 步骤五：量化计算
		128	+
		129	+---
		130	+
		131	+#### quant_mode=0 (Block Quant)
		132	+
		133	+分组划分：将输出沿最后一维按128元素为一组划分：
		134	+
		135	+$$
		136	+\mathbf{y} = [\mathbf{g}_0, \mathbf{g}_1, \ldots, \mathbf{g}_K], \quad K = \lceil D/2 / 128 \rceil
		137	+$$
		138	+
		139	+每个组 $\mathbf{g}_i \in \mathbb{R}^{N \times 128}$。
		140	+
		141	+非有限值屏蔽与绝对值计算：
		142	+
		143	+$$
		144	+\begin{aligned}
		145	+\mathbf{z}[n, j] &= \mathbf{y}_{\text{weighted}}[n, j] \cdot 0 \quad \text{(生成零张量)} \\
		146	+\mathbf{m}_{\text{finite}}[n, j] &= (\mathbf{z}[n, j] = \mathbf{z}[n, j]) \\
		147	+\mathbf{y}_{\text{abs}}[n, j] &=
		148	+\begin{cases}
		149	+\|\mathbf{y}_{\text{weighted}}[n, j]\|, & \mathbf{m}_{\text{finite}}[n, j] \\
		150	+0, & \text{otherwise}
		151	+\end{cases}
		152	+\end{aligned}
		153	+$$
		154	+
		155	+屏蔽原理：NaN的特性是 `NaN != NaN`；同时 `Inf * 0` 也会得到NaN，因此该步骤在计算amax时屏蔽NaN和Inf。
		156	+
		157	+Scale计算：
		158	+
		159	+对于第 $i$ 个组（包含128个连续元素）：
		160	+
		161	+$$
		162	+a_i = \max_{j=0}^{127} \mathbf{y}_{\text{abs}}[j]
		163	+$$
		164	+
		165	+$$
		166	+\hat{a}_i = \max(a_i, 10^{-4})
		167	+$$
		168	+
		169	+$$
		170	+s_i^{\text{raw}} = \frac{\hat{a}_i}{M_{\text{fp8}}}
		171	+$$
		172	+
		173	+其中 $M_{\text{fp8}}$ 取值：
		174	+- FP8 E4M3FN：$M_{\text{fp8}} = 448.0$
		175	+- FP8 E5M2：$M_{\text{fp8}} = 57344.0$
		176	+
		177	+Scale输出与InvScale计算：
		178	+
		179	+当 `round_scale=false` 时：
		180	+
		181	+$$
		182	+s_i = s_i^{\text{raw}}, \quad \text{InvScale}_i = \frac{M_{\text{fp8}}}{\hat{a}_i} = \frac{1}{s_i}
		183	+$$
		184	+
		185	+当 `round_scale=true` 时，将scale向上取整到2的幂：
		186	+
		187	+$$
		188	+e_i = \lceil \log_2(s_i^{\text{raw}}) \rceil
		189	+$$
		190	+
		191	+$$
		192	+s_i = 2^{e_i}, \quad \text{InvScale}_i = 2^{-e_i}
		193	+$$
		194	+
		195	+其中 $s_i$ 写入FLOAT32类型的scale输出。
		196	+
		197	+量化计算：
		198	+
		199	+$$
		200	+\mathbf{y}_{\text{scaled}}[n, j] = \mathbf{y}_{\text{weighted}}[n, j] \cdot \text{InvScale}_i, \quad j \in \text{group } i
		201	+$$
		202	+
		203	+若 $\mathbf{y}_{\text{scaled}}[n,j]$ 为NaN或Inf，实现会使用原始 $\mathbf{y}_{\text{weighted}}[n,j]$ 作为FP8 cast输入：
		204	+
		205	+$$
		206	+\mathbf{y}_{\text{cast\_in}}[n, j] =
		207	+\begin{cases}
		208	+\mathbf{y}_{\text{scaled}}[n, j], & \mathbf{y}_{\text{scaled}}[n, j] \text{ is finite} \\
		209	+\mathbf{y}_{\text{weighted}}[n, j], & \text{otherwise}
		210	+\end{cases}
		211	+$$
		212	+
		213	+$$
		214	+\mathbf{y}_{\text{quant}}[n, j] = \text{cast\_fp8\_rint}(\mathbf{y}_{\text{cast\_in}}[n, j])
		215	+$$
		216	+
		217	+其中 `cast_fp8_rint` 为FP32到FP8的类型转换，采用RINT（就近舍入）模式。
		218	+
		219	+---
		220	+
		221	+#### quant_mode=1 (MX Quant)
		222	+
		223	+MX量化原理：采用E8M0 Scale + FP8 Data的组合。
		224	+
		225	+分组方式：每32元素为一组：
		226	+
		227	+$$
		228	+\mathbf{y} = [\mathbf{g}_0, \mathbf{g}_1, \ldots, \mathbf{g}_K], \quad \mathbf{g}_i \in \mathbb{R}^{32}
		229	+$$
		230	+
		231	+Amax计算：
		232	+
		233	+$$
		234	+a_i = \max_{j=0}^{31} \|\mathbf{g}_i[j]\|
		235	+$$
		236	+
		237	+$$
		238	+\hat{a}_i = \max(a_i, 10^{-4})
		239	+$$
		240	+
		241	+原始Scale计算：
		242	+
		243	+$$
		244	+s_i^{\text{raw}} = \frac{\hat{a}_i}{M_{\text{fp8}}}
		245	+$$
		246	+
		247	+其中 $M_{\text{fp8}}$ 取值：
		248	+- FP8 E4M3FN：$M_{\text{fp8}} = 448.0$
		249	+- FP8 E5M2：$M_{\text{fp8}} = 57344.0$
		250	+
		251	+quant_mode=1仅支持 `round_scale=true`，将原始scale向上取整到2的幂：
		252	+
		253	+$$
		254	+e_i = \lceil \log_2(s_i^{\text{raw}}) \rceil
		255	+$$
		256	+
		257	+等价于基于FP32位模式计算：
		258	+
		259	+$$
		260	+e_i = E(s_i^{\text{raw}}) - 127 + \mathbf{1}_{\text{mantissa}(s_i^{\text{raw}}) \ne 0}
		261	+$$
		262	+
		263	+E8M0 Scale编码：
		264	+
		265	+$$
		266	+s_i^{\text{e8m0}} = e_i + 127
		267	+$$
		268	+
		269	+其中 $s_i^{\text{e8m0}}$ 写入FLOAT8_E8M0类型的scale输出，表示的实际scale值为 $2^{e_i}$。
		270	+
		271	+InvScale计算：
		272	+
		273	+$$
		274	+\text{InvScale}_i = 2^{-e_i}
		275	+$$
		276	+
		277	+量化计算：
		278	+
		279	+$$
		280	+\mathbf{y}_{\text{quant}}[j] = \text{cast\_fp8\_rint}\left(\mathbf{y}_{\text{weighted}}[j] \cdot \text{InvScale}_i\right), \quad j \in \text{group } i
		281	+$$
		282	+
		283	+---
		284	+
		285	+#### quant_mode=2 (HiFp8 Static Quant)
		286	+
		287	+静态量化说明：使用预先提供的 `invScale` 对加权后的SwiGLU输出进行缩放量化。
		288	+
		289	+情况1：无GroupIndex（groupIndex为空）：
		290	+
		291	+$$
		292	+\mathbf{y}_{\text{quant}}[n, d] = \text{hif8\_cast}\left(\mathbf{y}_{\text{weighted}}[n, d] \cdot \text{invScale}[0]\right), \quad n \in [0, N), \quad d \in [0, D/2)
		293	+$$
		294	+
		295	+其中 `hif8_cast` 为HiFloat8类型转换函数。
		296	+
		297	+情况2：有GroupIndex（groupIndex非空）：
		298	+
		299	+设 $G$ 为MoE专家分组数，$\text{groupIndex}[g]$ 表示第 $g$ 个专家处理的token数量。
		300	+
		301	+计算每个group的起止索引：
		302	+
		303	+$$
		304	+\text{start}^{(0)} = 0, \quad \text{end}^{(g)} = \sum_{k=0}^{g} \text{groupIndex}[k], \quad \text{start}^{(g)} = \text{end}^{(g-1)}
		305	+$$
		306	+
		307	+对于第 $g$ 个group，使用对应的缩放因子 $\text{invScale}[g]$ 进行量化：
		308	+
		309	+$$
		310	+\mathbf{y}_{\text{quant}}[n, d] = \text{hif8\_cast}\left(\mathbf{y}_{\text{weighted}}[n, d] \cdot \text{invScale}[g]\right), \quad n \in [\text{start}^{(g)}, \text{end}^{(g)}), \quad d \in [0, D/2)
		311	+$$
		312	+
		313	+MoE场景说明：在MoE推理中，不同专家处理不同数量的token，groupIndex用于标识每个专家处理的token范围，invScale为每个专家预先计算的静态缩放因子。
		314	+
		315	+---
		316	+
		317	+#### quant_mode=3 (HiFp8 Dynamic Quant)
		318	+
		319	+动态量化说明：根据加权后的SwiGLU输出动态计算缩放因子进行量化。
		320	+
		321	+情况1：无GroupIndex（groupIndex为空）：
		322	+
		323	+计算全局绝对值最大值：
		324	+
		325	+$$
		326	+a_{\max} = \max\left(\max_{n \in [0, N), d \in [0, D/2)} \|\mathbf{y}_{\text{weighted}}[n, d]\|, \epsilon\right)
		327	+$$
		328	+
		329	+其中 $\epsilon$ 为数值稳定性常数。
		330	+
		331	+计算缩放因子：
		332	+
		333	+$$
		334	+s = \frac{a_{\max}}{M_{\text{hif8}}}
		335	+$$
		336	+
		337	+其中 $M_{\text{hif8}}$ 为 `dstTypeMax`，表示HiFloat8类型的最大有限值。
		338	+
		339	+量化计算：
		340	+
		341	+$$
		342	+\mathbf{y}_{\text{quant}}[n, d] = \text{hif8\_cast}\left(\frac{\mathbf{y}_{\text{weighted}}[n, d]}{s}\right), \quad n \in [0, N), \quad d \in [0, D/2)
		343	+$$
		344	+
		345	+其中 `hif8_cast` 为HiFloat8类型转换函数。
		346	+
		347	+情况2：有GroupIndex（groupIndex非空）：
		348	+
		349	+设 $G$ 为MoE专家分组数，$\text{groupIndex}[g]$ 表示第 $g$ 个专家处理的token数量。
		350	+
		351	+计算每个group的起止索引：
		352	+
		353	+$$
		354	+\text{start}^{(0)} = 0, \quad \text{end}^{(g)} = \sum_{k=0}^{g} \text{groupIndex}[k], \quad \text{start}^{(g)} = \text{end}^{(g-1)}
		355	+$$
		356	+
		357	+对于第 $g$ 个group，提取对应的数据：
		358	+
		359	+$$
		360	+\mathbf{y}^{(g)} = \mathbf{y}_{\text{weighted}}[\text{start}^{(g)}:\text{end}^{(g)}, :]
		361	+$$
		362	+
		363	+计算该group的绝对值最大值：
		364	+
		365	+$$
		366	+a_{\max}^{(g)} = \max\left(\max_{n \in [\text{start}^{(g)}, \text{end}^{(g)}), d \in [0, D/2)} \|\mathbf{y}^{(g)}[n, d]\|, \epsilon\right)
		367	+$$
		368	+
		369	+计算该group的缩放因子：
		370	+
		371	+$$
		372	+s^{(g)} = \frac{a_{\max}^{(g)}}{M_{\text{hif8}}}
		373	+$$
		374	+
		375	+对该group进行量化：
		376	+
		377	+$$
		378	+\mathbf{y}_{\text{quant}}[n, d] = \text{hif8\_cast}\left(\frac{\mathbf{y}_{\text{weighted}}[n, d]}{s^{(g)}}\right), \quad n \in [\text{start}^{(g)}, \text{end}^{(g)}), \quad d \in [0, D/2)
		379	+$$
		380	+
		381	+MoE场景说明：在MoE推理中，不同专家处理不同数量的token，groupIndex用于标识每个专家处理的token范围，每个group独立计算缩放因子以适应不同数据分布。
43		382
44	## 函数原型	383	## 函数原型
45		384
	@@ -55,8 +394,8 @@ aclnnStatus aclnnSwigluGroupQuantGetWorkspaceSize(
55	int64_t quantMode,	394	int64_t quantMode,
56	int64_t blockSize,	395	int64_t blockSize,
57	bool roundScale,	396	bool roundScale,
58	- float clampLimit,	397	+ double clampLimit,
59	- float dstTypeMax,	398	+ double dstTypeMax,
60	bool outputOrigin,	399	bool outputOrigin,
61	const aclTensor *yOut,	400	const aclTensor *yOut,
62	const aclTensor *yScaleOut,	401	const aclTensor *yScaleOut,
	@@ -106,24 +445,24 @@ aclnnStatus aclnnSwigluGroupQuant(
106	<td><ul><li>shape为[...,D]。</li><li>D必须大于等于256，且能被256整除。</li><li>不支持空Tensor。</li></ul></td>	445	<td><ul><li>shape为[...,D]。</li><li>D必须大于等于256，且能被256整除。</li><li>不支持空Tensor。</li></ul></td>
107	<td>FLOAT、FLOAT16、BFLOAT16</td>	446	<td>FLOAT、FLOAT16、BFLOAT16</td>
108	<td>ND</td>	447	<td>ND</td>
109	- <td>1-7</td>	448	+ <td>2-8</td>
110	<td>×</td>	449	<td>×</td>
111	</tr>	450	</tr>
112	<tr>	451	<tr>
113	<td>weightOptional（aclTensor*）</td>	452	<td>weightOptional（aclTensor*）</td>
114	- <td>输入</td>	453	+ <td>输入（可选）</td>
115	- <td>量化前按token乘到SwiGLU输出上的权重。</td>	454	+ <td>MOE权重张量，用于SwiGLU输出的加权计算。</td>
116	<td><ul><li>可选参数，不支持空Tensor。</li><li>不为空时，数据类型为FLOAT32，元素个数需等于x除最后一维外的元素个数之积。</li></ul></td>	455	<td><ul><li>可选参数，不支持空Tensor。</li><li>不为空时，数据类型为FLOAT32，元素个数需等于x除最后一维外的元素个数之积。</li></ul></td>
117	<td>FLOAT32</td>	456	<td>FLOAT32</td>
118	<td>ND</td>	457	<td>ND</td>
119	- <td>1-2</td>	458	+ <td>2-8</td>
120	<td>×</td>	459	<td>×</td>
121	</tr>	460	</tr>
122	<tr>	461	<tr>
123	<td>groupIndexOptional（aclTensor*）</td>	462	<td>groupIndexOptional（aclTensor*）</td>
124	- <td>输入</td>	463	+ <td>输入（可选）</td>
125	<td>count模式的group token数。</td>	464	<td>count模式的group token数。</td>
126	- <td><ul><li>可选参数，不支持空Tensor。</li><li>不为空时，数据类型为INT64，shape为[groupNum]。</li></ul></td>	465	+ <td><ul><li>可选参数，不支持空Tensor。</li><li>不为空时，数据类型为INT64，shape为[G]。</li></ul></td>
127	<td>INT64</td>	466	<td>INT64</td>
128	<td>ND</td>	467	<td>ND</td>
129	<td>1</td>	468	<td>1</td>
	@@ -131,9 +470,9 @@ aclnnStatus aclnnSwigluGroupQuant(
131	</tr>	470	</tr>
132	<tr>	471	<tr>
133	<td>scaleOptional（aclTensor*）</td>	472	<td>scaleOptional（aclTensor*）</td>
134	- <td>输入</td>	473	+ <td>输入（可选）</td>
135	- <td>静态量化输入的scale张量。</td>	474	+ <td>quantMode=2时静态量化输入的invScale张量。</td>
136	- <td><ul><li>可选参数，仅quant_mode为2时使用。</li><li>quant_mode为3时不使用，可传空。</li></ul></td>	475	+ <td><ul><li>可选参数，仅quant_mode为2时使用。</li><li>groupIndex存在的话，shape=[G]，不存在的话shape=[1]。</li></ul></td>
137	<td>FLOAT32</td>	476	<td>FLOAT32</td>
138	<td>ND</td>	477	<td>ND</td>
139	<td>1</td>	478	<td>1</td>
	@@ -143,7 +482,7 @@ aclnnStatus aclnnSwigluGroupQuant(
143	<td>dstType（int64_t）</td>	482	<td>dstType（int64_t）</td>
144	<td>输入</td>	483	<td>输入</td>
145	<td>目标量化类型。</td>	484	<td>目标量化类型。</td>
146	- <td><ul><li>支持取值35、36、40、41，分别表示FLOAT8_E5M2、FLOAT8_E4M3FN、FLOAT4_E2M1、FLOAT4_E1M2。</li><li>dstType为40或41时，quantMode必须为1。</li></ul></td>	485	+ <td><ul><li>仅quantMode为0或1时，该参数生效。</li><li>支持取值35、36、40、41，分别表示FLOAT8_E5M2、FLOAT8_E4M3FN、FLOAT4_E2M1、FLOAT4_E1M2。</li><li>dstType为40或41时，quantMode必须为1。</li></ul></td>
147	<td>-</td>	486	<td>-</td>
148	<td>-</td>	487	<td>-</td>
149	<td>-</td>	488	<td>-</td>
	@@ -180,7 +519,7 @@ aclnnStatus aclnnSwigluGroupQuant(
180	<td>-</td>	519	<td>-</td>
181	</tr>	520	</tr>
182	<tr>	521	<tr>
183	- <td>clampLimit（float）</td>	522	+ <td>clampLimit（double）</td>
184	<td>输入</td>	523	<td>输入</td>
185	<td>SwiGLU计算前的clamp阈值。</td>	524	<td>SwiGLU计算前的clamp阈值。</td>
186	<td><ul><li>-1.0表示不启用clamp。</li><li>启用clamp时，clampLimit必须大于0。</li></ul></td>	525	<td><ul><li>-1.0表示不启用clamp。</li><li>启用clamp时，clampLimit必须大于0。</li></ul></td>
	@@ -190,10 +529,10 @@ aclnnStatus aclnnSwigluGroupQuant(
190	<td>-</td>	529	<td>-</td>
191	</tr>	530	</tr>
192	<tr>	531	<tr>
193	- <td>dstTypeMax（float）</td>	532	+ <td>dstTypeMax（double）</td>
194	<td>输入</td>	533	<td>输入</td>
195	<td>目标量化类型的最大有限值。</td>	534	<td>目标量化类型的最大有限值。</td>
196	- <td><ul><li>quant_mode为2或3时，用于计算scale = amax / dstTypeMax。</li><li>默认值为15.0（HIFLOAT8的最大值）。</li></ul></td>	535	+ <td><ul><li>仅quantMode为3时，该参数生效。</li><li>默认值为15.0。</li></ul></td>
197	<td>-</td>	536	<td>-</td>
198	<td>-</td>	537	<td>-</td>
199	<td>-</td>	538	<td>-</td>
	@@ -203,40 +542,40 @@ aclnnStatus aclnnSwigluGroupQuant(
203	<td>outputOrigin（bool）</td>	542	<td>outputOrigin（bool）</td>
204	<td>输入</td>	543	<td>输入</td>
205	<td>是否输出量化前的SwiGLU结果。</td>	544	<td>是否输出量化前的SwiGLU结果。</td>
206	- <td><ul><li>MX FP4模式下，yOriginOut仅作占位。</li><li>quant_mode为2或3时，支持输出原始激活值。</li></ul></td>	545	+ <td><ul><li>true表示支持输出原始激活值yOrigin，false表示不支持输出原始激活值yOrigin。</li><li>quantMode为0或1时支持false。</li><li>quantMode为2或3时true/false都支持。</li></ul></td>
207	<td>-</td>	546	<td>-</td>
208	<td>-</td>	547	<td>-</td>
209	<td>-</td>	548	<td>-</td>
210	<td>-</td>	549	<td>-</td>
211	</tr>	550	</tr>
212	<tr>	551	<tr>
213	- <td>y（aclTensor*）</td>	552	+ <td>yOut（aclTensor*）</td>
214	<td>输出</td>	553	<td>输出</td>
215	<td>量化输出。</td>	554	<td>量化输出。</td>
216	- <td><ul><li>dstType为35或36时，shape为[...,D/2]。</li><li>dstType为40或41时，shape为[...,D/4]。</li><li>dstType为27时，shape为[...,D/2]。</li><li>数据类型需与dstType一致。</li><li>不支持空Tensor。</li></ul></td>	555	+ <td><ul><li>quantMode为0或1时，数据类型需与dstType一致，dstType为35或36时，shape为[...,D/2]；dstType为40或41时，shape为[...,D/4]；dstType为27时，shape为[...,D/2]。</li><li>quantMode为2或3时，数据类型默认为HIFLOAT8，shape为[...,D/2]。</li><li>不支持空Tensor。</li></ul></td>
217	<td>HIFLOAT8、FLOAT8_E5M2、FLOAT8_E4M3FN、FLOAT4_E2M1、FLOAT4_E1M2</td>	556	<td>HIFLOAT8、FLOAT8_E5M2、FLOAT8_E4M3FN、FLOAT4_E2M1、FLOAT4_E1M2</td>
218	<td>ND</td>	557	<td>ND</td>
219	- <td>1-7</td>	558	+ <td>2-8</td>
220	<td>×</td>	559	<td>×</td>
221	</tr>	560	</tr>
222	<tr>	561	<tr>
223	- <td>yScale（aclTensor*）</td>	562	+ <td>yScaleOut（aclTensor*）</td>
224	<td>输出</td>	563	<td>输出</td>
225	<td>量化scale输出。</td>	564	<td>量化scale输出。</td>
226	- <td><ul><li>quantMode为0时，shape为[...,ceil((D/2)/128)]，数据类型为FLOAT32。</li><li>quantMode为1时，shape为[...,ceil(ceil((D/2)/32)/2),2]，数据类型为FLOAT8_E8M0。</li><li>quantMode为2或3时，无groupIndex时shape为[1]，有groupIndex时shape为[G]，数据类型为FLOAT32。</li><li>不支持空Tensor。</li></ul></td>	565	+ <td><ul><li>quantMode为0时，shape为[...,ceil((D/2)/128)]，数据类型为FLOAT32。</li><li>quantMode为1时，shape为[...,ceil(ceil((D/2)/32)/2),2]，数据类型为FLOAT8_E8M0。</li><li>quantMode为3时，无groupIndex时shape为[1]，有groupIndex时shape为[G]，数据类型为FLOAT32。</li><li>不支持空Tensor。</li></ul></td>
227	<td>FLOAT32、FLOAT8_E8M0</td>	566	<td>FLOAT32、FLOAT8_E8M0</td>
228	<td>ND</td>	567	<td>ND</td>
229	<td>1-8</td>	568	<td>1-8</td>
230	<td>×</td>	569	<td>×</td>
231	</tr>	570	</tr>
232	<tr>	571	<tr>
233	- <td>yOrigin（aclTensor*）</td>	572	+ <td>yOriginOut（aclTensor*）</td>
234	<td>输出</td>	573	<td>输出</td>
235	<td>量化前的SwiGLU结果。</td>	574	<td>量化前的SwiGLU结果。</td>
236	- <td><ul><li>shape为[...,D/2]。</li><li>数据类型需与x一致。</li><li>不支持空Tensor。</li></ul></td>	575	+ <td><ul><li>shape为[...,D/2]。</li><li>数据类型需与x一致。</li><li>不支持空指针。</li></ul></td>
237	<td>FLOAT、FLOAT16、BFLOAT16</td>	576	<td>FLOAT、FLOAT16、BFLOAT16</td>
238	<td>ND</td>	577	<td>ND</td>
239	- <td>1-7</td>	578	+ <td>2-8</td>
240	<td>×</td>	579	<td>×</td>
241	</tr>	580	</tr>
242	<tr>	581	<tr>
	@@ -353,15 +692,19 @@ aclnnStatus aclnnSwigluGroupQuant(
353		692
354	## 约束说明	693	## 约束说明
355		694
		695	+- 确定性计算：aclnnSwigluGroupQuant默认确定性实现。
356	- quantMode为0时，仅支持FP8输出，blockSize支持0或128。	696	- quantMode为0时，仅支持FP8输出，blockSize支持0或128。
357	- quantMode为1时，支持FP8/FP4输出，blockSize支持0或32，roundScale必须为true。	697	- quantMode为1时，支持FP8/FP4输出，blockSize支持0或32，roundScale必须为true。
358	-- quantMode为2时，支持HIFP8静态量化输出，dstType, blockSize和roundScale不生效。	698	+- quantMode为2或3时，支持HIFP8量化输出，dstType, blockSize和roundScale不生效。输入x的维度为[T, D]或[B, S, D]，需满足以下规格约束：
359	-- quantMode为3时，支持HIFP8动态量化输出，dstType, blockSize和roundScale不生效。	699	+
		700	+ \| 规格项 \| 规格 \| 规格说明 \|
		701	+ \| :--- \| :--- \| :--- \|
		702	+ \| B \| 1~31 \| - \|
		703	+ \| S \| 0~128K \| - \|
		704	+ \| D/2 \| 512, 768, 1024, 1536, 1792, 2048, 2560, 4096 \| - \|
		705	+ \| dstTypeMax \| 15, 56, 224, 32768 \| - \|
360	- dstType为FLOAT4_E2M1或FLOAT4_E1M2时，必须使用quantMode=1。	706	- dstType为FLOAT4_E2M1或FLOAT4_E1M2时，必须使用quantMode=1。
361	-- yScale的数据类型必须与quantMode匹配：Block FP8为FLOAT32，MX为FLOAT8_E8M0，HIFP8为FLOAT32。	707	+- yScale的数据类型必须与quantMode匹配：quantMode=0或3时数据类型为FLOAT32，quantMode=1时数据类型为FLOAT8_E8M0。
362	-- quantMode为2或3时，groupIndexOptional可用于MoE场景的分组量化，yScaleOut的shape为[G]（G为group数量）。
363	-- 确定性计算：
364	- - aclnnSwigluGroupQuant默认确定性实现。
365		708
366	## 调用示例	709	## 调用示例
367		710
	@@ -376,166 +719,254 @@ aclnnStatus aclnnSwigluGroupQuant(
376	#include "aclnnop/aclnn_swiglu_group_quant.h"	719	#include "aclnnop/aclnn_swiglu_group_quant.h"
377		720
378	#define CHECK_RET(cond, return_expr) \	721	#define CHECK_RET(cond, return_expr) \
379	- do { \	722	+ do { \
380	- if (!(cond)) { \	723	+ if (!(cond)) { \
381	- return_expr; \	724	+ return_expr; \
382	- } \	725	+ } \
383	- } while (0)	726	+ } while (0)
384		727
385	-#define LOG_PRINT(message, ...) \	728	+#define LOG_PRINT(message, ...) \
386	- do { \	729	+ do { \
387	- printf(message, ##__VA_ARGS__); \	730	+ printf(message, ##__VA_ARGS__); \
388	- } while (0)	731	+ } while (0)
389		732
390	-int64_t GetShapeSize(const std::vector<int64_t>& shape) {	733	+int64_t GetShapeSize(const std::vector<int64_t>& shape)
391	- int64_t shapeSize = 1;	734	+{
392	- for (auto dim : shape) {	735	+ int64_t shapeSize = 1;
393	- shapeSize *= dim;	736	+ for (auto dim : shape) {
394	- }	737	+ shapeSize *= dim;
395	- return shapeSize;	738	+ }
		739	+ return shapeSize;
396	}	740	}
397		741
398	-int Init(int32_t deviceId, aclrtStream* stream) {	742	+int Init(int32_t deviceId, aclrtStream* stream)
399	- auto ret = aclInit(nullptr);	743	+{
400	- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);	744	+ auto ret = aclInit(nullptr);
401	- ret = aclrtSetDevice(deviceId);	745	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
402	- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);	746	+ ret = aclrtSetDevice(deviceId);
403	- ret = aclrtCreateStream(stream);	747	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
404	- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);	748	+ ret = aclrtCreateStream(stream);
405	- return ACL_SUCCESS;	749	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
		750	+ return ACL_SUCCESS;
406	}	751	}
407		752
408	-bool CheckHardwareSupport() {	753	+bool CheckHardwareSupport()
409	- const char* socName = aclrtGetSocName();	754	+{
410	- if (socName == nullptr) {	755	+ const char* socName = aclrtGetSocName();
411	- LOG_PRINT("Warning: Cannot get SOC name, skip hardware check\n");	756	+ if (socName == nullptr) {
412	- return true;	757	+ LOG_PRINT("Warning: Cannot get SOC name, skip hardware check\n");
413	- }	758	+ return true;
		759	+ }
414		760
415	- LOG_PRINT("Current SOC: %s\n", socName);	761	+ LOG_PRINT("Current SOC: %s\n", socName);
416	- if (strstr(socName, "Ascend950") != nullptr \|\| strstr(socName, "ascend950") != nullptr) {	762	+ if (strstr(socName, "Ascend950") != nullptr \|\| strstr(socName, "ascend950") != nullptr) {
417	- return true;	763	+ return true;
418	- }	764	+ }
419		765
420	- LOG_PRINT("Warning: SwigluGroupQuant only supports Ascend950, current SOC '%s' is not supported. Skip test.\n",	766	+ LOG_PRINT("Warning: SwigluGroupQuant only supports Ascend950, current SOC '%s' is not supported. Skip test.\n",
421	- socName);	767	+ socName);
422	- return false;	768	+ return false;
423	}	769	}
424		770
425	-void Finalize(int32_t deviceId, aclrtStream stream) {	771	+void Finalize(int32_t deviceId, aclrtStream stream)
426	- (void)aclrtDestroyStream(stream);	772	+{
427	- (void)aclrtResetDevice(deviceId);	773	+ (void)aclrtDestroyStream(stream);
428	- (void)aclFinalize();	774	+ (void)aclrtResetDevice(deviceId);
		775	+ (void)aclFinalize();
		776	+}
		777	+
		778	+struct AclTensorResource {
		779	+ void* deviceAddr = nullptr;
		780	+ aclTensor* tensor = nullptr;
		781	+};
		782	+
		783	+void DestroyAclTensorResource(AclTensorResource& resource)
		784	+{
		785	+ if (resource.tensor != nullptr) {
		786	+ aclDestroyTensor(resource.tensor);
		787	+ resource.tensor = nullptr;
		788	+ }
		789	+ if (resource.deviceAddr != nullptr) {
		790	+ aclrtFree(resource.deviceAddr);
		791	+ resource.deviceAddr = nullptr;
		792	+ }
429	}	793	}
430		794
431	template <typename T>	795	template <typename T>
432	int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,	796	int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
433	- aclDataType dataType, aclTensor** tensor) {	797	+ aclDataType dataType, aclTensor** tensor)
434	- auto size = GetShapeSize(shape) * sizeof(T);	798	+{
435	- auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);	799	+ auto size = GetShapeSize(shape) * sizeof(T);
436	- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);	800	+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
		801	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
437		802
438	- ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);	803	+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
439	- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);	804	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
440		805
441	- std::vector<int64_t> strides(shape.size(), 1);	806	+ std::vector<int64_t> strides(shape.size(), 1);
442	- for (int64_t i = static_cast<int64_t>(shape.size()) - 2; i >= 0; --i) {	807	+ for (int64_t i = static_cast<int64_t>(shape.size()) - 2; i >= 0; --i) {
443	- strides[i] = shape[i + 1] * strides[i + 1];	808	+ strides[i] = shape[i + 1] * strides[i + 1];
444	- }	809	+ }
445		810
446	- *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, ACL_FORMAT_ND,	811	+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, ACL_FORMAT_ND, shape.data(),
447	- shape.data(), shape.size(), *deviceAddr);	812	+ shape.size(), *deviceAddr);
448	- return ACL_SUCCESS;	813	+ return ACL_SUCCESS;
449	}	814	}
450		815
451	-int main() {	816	+struct SwigluGroupQuantCase {
452	- int32_t deviceId = 0;	817	+ const char* name;
453	- aclrtStream stream;	818	+ int64_t quantMode;
454	- auto ret = Init(deviceId, &stream);	819	+ int64_t blockSize;
455	- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);	820	+ bool roundScale;
		821	+ std::vector<int64_t> yScaleShape;
		822	+ aclDataType yScaleDataType;
		823	+ bool yScaleIsE8M0;
		824	+ int64_t dstType;
		825	+ double dstTypeMax;
		826	+ aclDataType yDataType;
		827	+ int64_t yElementSize;
		828	+ bool hasScaleInput;
		829	+ std::vector<int64_t> scaleShape;
		830	+};
		831	+
		832	+int RunSwigluGroupQuantCase(const SwigluGroupQuantCase& testCase, aclrtStream stream)
		833	+{
		834	+ std::vector<int64_t> xShape = {2, 256};
		835	+ std::vector<int64_t> yShape = {2, 128};
		836	+ std::vector<int64_t> yOriginShape = {2, 128};
		837	+
		838	+ std::vector<uint16_t> xHostData(GetShapeSize(xShape), 0);
		839	+ for (size_t i = 0; i < xHostData.size(); ++i) {
		840	+ xHostData[i] = static_cast<uint16_t>(i % 23);
		841	+ }
		842	+ std::vector<uint8_t> yHostData(GetShapeSize(yShape) * testCase.yElementSize, 0);
		843	+ std::vector<uint8_t> yScaleE8M0HostData(GetShapeSize(testCase.yScaleShape), 0);
		844	+ std::vector<float> yScaleFp32HostData(GetShapeSize(testCase.yScaleShape), 0.0f);
		845	+ std::vector<uint16_t> yOriginHostData(GetShapeSize(yOriginShape), 0);
		846	+ std::vector<float> scaleHostData;
		847	+ if (testCase.hasScaleInput) {
		848	+ scaleHostData.resize(GetShapeSize(testCase.scaleShape), 1.0f);
		849	+ }
		850	+
		851	+ AclTensorResource xResource;
		852	+ AclTensorResource yResource;
		853	+ AclTensorResource yScaleResource;
		854	+ AclTensorResource yOriginResource;
		855	+ AclTensorResource scaleResource;
		856	+
		857	+ auto ret = CreateAclTensor(xHostData, xShape, &xResource.deviceAddr, ACL_FLOAT16, &xResource.tensor);
		858	+ CHECK_RET(ret == ACL_SUCCESS, return ret);
		859	+ ret = CreateAclTensor(yHostData, yShape, &yResource.deviceAddr, testCase.yDataType, &yResource.tensor);
		860	+ CHECK_RET(ret == ACL_SUCCESS, DestroyAclTensorResource(xResource); return ret);
		861	+ if (testCase.yScaleIsE8M0) {
		862	+ ret = CreateAclTensor(yScaleE8M0HostData, testCase.yScaleShape, &yScaleResource.deviceAddr,
		863	+ testCase.yScaleDataType, &yScaleResource.tensor);
		864	+ } else {
		865	+ ret = CreateAclTensor(yScaleFp32HostData, testCase.yScaleShape, &yScaleResource.deviceAddr,
		866	+ testCase.yScaleDataType, &yScaleResource.tensor);
		867	+ }
		868	+ CHECK_RET(ret == ACL_SUCCESS, DestroyAclTensorResource(xResource); DestroyAclTensorResource(yResource);
		869	+ DestroyAclTensorResource(yScaleResource); return ret);
		870	+ ret = CreateAclTensor(yOriginHostData, yOriginShape, &yOriginResource.deviceAddr, ACL_FLOAT16,
		871	+ &yOriginResource.tensor);
		872	+ CHECK_RET(ret == ACL_SUCCESS, DestroyAclTensorResource(xResource); DestroyAclTensorResource(yResource);
		873	+ DestroyAclTensorResource(yScaleResource); DestroyAclTensorResource(yOriginResource); return ret);
		874	+ if (testCase.hasScaleInput) {
		875	+ ret = CreateAclTensor(scaleHostData, testCase.scaleShape, &scaleResource.deviceAddr,
		876	+ ACL_FLOAT, &scaleResource.tensor);
		877	+ CHECK_RET(ret == ACL_SUCCESS, DestroyAclTensorResource(xResource); DestroyAclTensorResource(yResource);
		878	+ DestroyAclTensorResource(yScaleResource); DestroyAclTensorResource(yOriginResource);
		879	+ DestroyAclTensorResource(scaleResource); return ret);
		880	+ }
		881	+
		882	+ double clampLimit = -1.0;
		883	+ bool outputOrigin = false;
		884	+ const aclTensor* scaleTensor = testCase.hasScaleInput ? scaleResource.tensor : nullptr;
		885	+
		886	+ LOG_PRINT("Run %s: quant_mode=%ld\n", testCase.name, testCase.quantMode);
		887	+
		888	+ uint64_t workspaceSize = 0;
		889	+ aclOpExecutor* executor = nullptr;
		890	+ ret = aclnnSwigluGroupQuantGetWorkspaceSize(xResource.tensor, nullptr, nullptr, scaleTensor, testCase.dstType,
		891	+ testCase.quantMode, testCase.blockSize, testCase.roundScale, clampLimit,
		892	+ testCase.dstTypeMax, outputOrigin, yResource.tensor,
		893	+ yScaleResource.tensor, yOriginResource.tensor, &workspaceSize,
		894	+ &executor);
		895	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSwigluGroupQuantGetWorkspaceSize failed. ERROR: %d\n", ret);
		896	+ DestroyAclTensorResource(xResource); DestroyAclTensorResource(yResource);
		897	+ DestroyAclTensorResource(yScaleResource); DestroyAclTensorResource(yOriginResource);
		898	+ DestroyAclTensorResource(scaleResource); return ret);
		899	+
		900	+ void* workspaceAddr = nullptr;
		901	+ if (workspaceSize > 0) {
		902	+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
		903	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret);
		904	+ DestroyAclTensorResource(xResource); DestroyAclTensorResource(yResource);
		905	+ DestroyAclTensorResource(yScaleResource); DestroyAclTensorResource(yOriginResource);
		906	+ DestroyAclTensorResource(scaleResource); return ret);
		907	+ }
		908	+
		909	+ ret = aclnnSwigluGroupQuant(workspaceAddr, workspaceSize, executor, stream);
		910	+ CHECK_RET(
		911	+ ret == ACL_SUCCESS, LOG_PRINT("aclnnSwigluGroupQuant failed. ERROR: %d\n", ret);
		912	+ if (workspaceAddr != nullptr) { aclrtFree(workspaceAddr); } DestroyAclTensorResource(xResource);
		913	+ DestroyAclTensorResource(yResource); DestroyAclTensorResource(yScaleResource);
		914	+ DestroyAclTensorResource(yOriginResource); DestroyAclTensorResource(scaleResource); return ret);
		915	+
		916	+ ret = aclrtSynchronizeStream(stream);
		917	+ CHECK_RET(
		918	+ ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret);
		919	+ if (workspaceAddr != nullptr) { aclrtFree(workspaceAddr); } DestroyAclTensorResource(xResource);
		920	+ DestroyAclTensorResource(yResource); DestroyAclTensorResource(yScaleResource);
		921	+ DestroyAclTensorResource(yOriginResource); DestroyAclTensorResource(scaleResource); return ret);
		922	+
		923	+ std::vector<uint8_t> resultData(GetShapeSize(yShape) * testCase.yElementSize, 0);
		924	+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), yResource.deviceAddr,
		925	+ resultData.size() * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
		926	+ CHECK_RET(
		927	+ ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret);
		928	+ if (workspaceAddr != nullptr) { aclrtFree(workspaceAddr); } DestroyAclTensorResource(xResource);
		929	+ DestroyAclTensorResource(yResource); DestroyAclTensorResource(yScaleResource);
		930	+ DestroyAclTensorResource(yOriginResource); DestroyAclTensorResource(scaleResource); return ret);
		931	+ LOG_PRINT("%s result[0] is: %d\n", testCase.name, resultData[0]);
		932	+
		933	+ DestroyAclTensorResource(xResource);
		934	+ DestroyAclTensorResource(yResource);
		935	+ DestroyAclTensorResource(yScaleResource);
		936	+ DestroyAclTensorResource(yOriginResource);
		937	+ DestroyAclTensorResource(scaleResource);
		938	+ if (workspaceAddr != nullptr) {
		939	+ aclrtFree(workspaceAddr);
		940	+ }
		941	+ return ACL_SUCCESS;
		942	+}
		943	+
		944	+int main()
		945	+{
		946	+ int32_t deviceId = 0;
		947	+ aclrtStream stream;
		948	+ auto ret = Init(deviceId, &stream);
		949	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
		950	+
		951	+ if (!CheckHardwareSupport()) {
		952	+ LOG_PRINT("\n=== Test SKIPPED (hardware not supported) ===\n");
		953	+ Finalize(deviceId, stream);
		954	+ return ACL_SUCCESS;
		955	+ }
		956	+
		957	+ std::vector<SwigluGroupQuantCase> testCases = {
		958	+ {"block_fp8", 0, 0, false, {2, 1}, ACL_FLOAT, false, 36, 448.0, ACL_FLOAT8_E4M3FN, 1, false, {}},
		959	+ {"mx_fp8", 1, 0, true, {2, 2, 2}, ACL_FLOAT8_E8M0, true, 36, 448.0, ACL_FLOAT8_E4M3FN, 1, false, {}},
		960	+ {"hifp8_static", 2, 0, false, {1}, ACL_FLOAT, false, 27, 448.0, ACL_HIFLOAT8, 1, true, {1}},
		961	+ {"hifp8_dynamic", 3, 0, false, {1}, ACL_FLOAT, false, 27, 15.0, ACL_HIFLOAT8, 1, false, {}},
		962	+ };
		963	+
		964	+ for (const auto& testCase : testCases) {
		965	+ ret = RunSwigluGroupQuantCase(testCase, stream);
		966	+ CHECK_RET(ret == ACL_SUCCESS, Finalize(deviceId, stream); return ret);
		967	+ }
456		968
457	- if (!CheckHardwareSupport()) {
458	- LOG_PRINT("\n=== Test SKIPPED (hardware not supported) ===\n");
459	Finalize(deviceId, stream);	969	Finalize(deviceId, stream);
460	return ACL_SUCCESS;	970	return ACL_SUCCESS;
461	- }
462	-
463	- std::vector<int64_t> xShape = {2, 256};
464	- std::vector<int64_t> yShape = {2, 128};
465	- std::vector<int64_t> yScaleShape = {2, 1};
466	- std::vector<int64_t> yOriginShape = {2, 128};
467	-
468	- std::vector<uint16_t> xHostData(GetShapeSize(xShape), 0);
469	- for (size_t i = 0; i < xHostData.size(); ++i) {
470	- xHostData[i] = static_cast<uint16_t>(i % 23);
471	- }
472	- std::vector<uint8_t> yHostData(GetShapeSize(yShape), 0);
473	- std::vector<float> yScaleHostData(GetShapeSize(yScaleShape), 0.0f);
474	- std::vector<uint16_t> yOriginHostData(GetShapeSize(yOriginShape), 0);
475	-
476	- void* xDeviceAddr = nullptr;
477	- void* yDeviceAddr = nullptr;
478	- void* yScaleDeviceAddr = nullptr;
479	- void* yOriginDeviceAddr = nullptr;
480	- aclTensor* x = nullptr;
481	- aclTensor* y = nullptr;
482	- aclTensor* yScale = nullptr;
483	- aclTensor* yOrigin = nullptr;
484	-
485	- ret = CreateAclTensor(xHostData, xShape, &xDeviceAddr, ACL_FLOAT16, &x);
486	- CHECK_RET(ret == ACL_SUCCESS, return ret);
487	- ret = CreateAclTensor(yHostData, yShape, &yDeviceAddr, ACL_FLOAT8_E4M3FN, &y);
488	- CHECK_RET(ret == ACL_SUCCESS, return ret);
489	- ret = CreateAclTensor(yScaleHostData, yScaleShape, &yScaleDeviceAddr, ACL_FLOAT, &yScale);
490	- CHECK_RET(ret == ACL_SUCCESS, return ret);
491	- ret = CreateAclTensor(yOriginHostData, yOriginShape, &yOriginDeviceAddr, ACL_FLOAT16, &yOrigin);
492	- CHECK_RET(ret == ACL_SUCCESS, return ret);
493	-
494	- int64_t dstType = 36;
495	- int64_t quantMode = 0;
496	- int64_t blockSize = 0;
497	- bool roundScale = false;
498	- double clampLimit = -1.0;
499	- bool outputOrigin = false;
500	-
501	- uint64_t workspaceSize = 0;
502	- aclOpExecutor* executor = nullptr;
503	- ret = aclnnSwigluGroupQuantGetWorkspaceSize(x, nullptr, nullptr, dstType, quantMode, blockSize, roundScale,
504	- clampLimit, outputOrigin, y, yScale, yOrigin, &workspaceSize,
505	- &executor);
506	- CHECK_RET(ret == ACL_SUCCESS,
507	- LOG_PRINT("aclnnSwigluGroupQuantGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
508	-
509	- void* workspaceAddr = nullptr;
510	- if (workspaceSize > 0) {
511	- ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
512	- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
513	- }
514	-
515	- ret = aclnnSwigluGroupQuant(workspaceAddr, workspaceSize, executor, stream);
516	- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSwigluGroupQuant failed. ERROR: %d\n", ret); return ret);
517	-
518	- ret = aclrtSynchronizeStream(stream);
519	- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
520	-
521	- std::vector<uint8_t> resultData(GetShapeSize(yShape), 0);
522	- ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), yDeviceAddr,
523	- resultData.size() * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
524	- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
525	- LOG_PRINT("result[0] is: %d\n", resultData[0]);
526	-
527	- aclDestroyTensor(x);
528	- aclDestroyTensor(y);
529	- aclDestroyTensor(yScale);
530	- aclDestroyTensor(yOrigin);
531	- aclrtFree(xDeviceAddr);
532	- aclrtFree(yDeviceAddr);
533	- aclrtFree(yScaleDeviceAddr);
534	- aclrtFree(yOriginDeviceAddr);
535	- if (workspaceSize > 0) {
536	- aclrtFree(workspaceAddr);
537	- }
538	- Finalize(deviceId, stream);
539	- return ACL_SUCCESS;
540	}	971	}
541	```	972	```

Mactivation/swiglu_group_quant/examples/test_aclnn_swiglu_group_quant.cpp+51-16

	@@ -1,3 +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 details.
		9	+ */
		10	+
1	#include <cstdint>	11	#include <cstdint>
2	#include <cstring>	12	#include <cstring>
3	#include <iostream>	13	#include <iostream>
	@@ -102,6 +112,12 @@ struct SwigluGroupQuantCase {
102	std::vector<int64_t> yScaleShape;	112	std::vector<int64_t> yScaleShape;
103	aclDataType yScaleDataType;	113	aclDataType yScaleDataType;
104	bool yScaleIsE8M0;	114	bool yScaleIsE8M0;
		115	+ int64_t dstType;
		116	+ double dstTypeMax;
		117	+ aclDataType yDataType;
		118	+ int64_t yElementSize;
		119	+ bool hasScaleInput;
		120	+ std::vector<int64_t> scaleShape;
105	};	121	};
106		122
107	int RunSwigluGroupQuantCase(const SwigluGroupQuantCase& testCase, aclrtStream stream) {	123	int RunSwigluGroupQuantCase(const SwigluGroupQuantCase& testCase, aclrtStream stream) {
	@@ -113,19 +129,24 @@ int RunSwigluGroupQuantCase(const SwigluGroupQuantCase& testCase, aclrtStream st
113	for (size_t i = 0; i < xHostData.size(); ++i) {	129	for (size_t i = 0; i < xHostData.size(); ++i) {
114	xHostData[i] = static_cast<uint16_t>(i % 23);	130	xHostData[i] = static_cast<uint16_t>(i % 23);
115	}	131	}
116	- std::vector<uint8_t> yHostData(GetShapeSize(yShape), 0);	132	+ std::vector<uint8_t> yHostData(GetShapeSize(yShape) * testCase.yElementSize, 0);
117	std::vector<uint8_t> yScaleE8M0HostData(GetShapeSize(testCase.yScaleShape), 0);	133	std::vector<uint8_t> yScaleE8M0HostData(GetShapeSize(testCase.yScaleShape), 0);
118	std::vector<float> yScaleFp32HostData(GetShapeSize(testCase.yScaleShape), 0.0f);	134	std::vector<float> yScaleFp32HostData(GetShapeSize(testCase.yScaleShape), 0.0f);
119	std::vector<uint16_t> yOriginHostData(GetShapeSize(yOriginShape), 0);	135	std::vector<uint16_t> yOriginHostData(GetShapeSize(yOriginShape), 0);
		136	+ std::vector<float> scaleHostData;
		137	+ if (testCase.hasScaleInput) {
		138	+ scaleHostData.resize(GetShapeSize(testCase.scaleShape), 1.0f);
		139	+ }
120		140
121	AclTensorResource xResource;	141	AclTensorResource xResource;
122	AclTensorResource yResource;	142	AclTensorResource yResource;
123	AclTensorResource yScaleResource;	143	AclTensorResource yScaleResource;
124	AclTensorResource yOriginResource;	144	AclTensorResource yOriginResource;
		145	+ AclTensorResource scaleResource;
125		146
126	auto ret = CreateAclTensor(xHostData, xShape, &xResource.deviceAddr, ACL_FLOAT16, &xResource.tensor);	147	auto ret = CreateAclTensor(xHostData, xShape, &xResource.deviceAddr, ACL_FLOAT16, &xResource.tensor);
127	CHECK_RET(ret == ACL_SUCCESS, return ret);	148	CHECK_RET(ret == ACL_SUCCESS, return ret);
128	- ret = CreateAclTensor(yHostData, yShape, &yResource.deviceAddr, ACL_FLOAT8_E4M3FN, &yResource.tensor);	149	+ ret = CreateAclTensor(yHostData, yShape, &yResource.deviceAddr, testCase.yDataType, &yResource.tensor);
129	CHECK_RET(ret == ACL_SUCCESS, DestroyAclTensorResource(xResource); return ret);	150	CHECK_RET(ret == ACL_SUCCESS, DestroyAclTensorResource(xResource); return ret);
130	if (testCase.yScaleIsE8M0) {	151	if (testCase.yScaleIsE8M0) {
131	ret = CreateAclTensor(yScaleE8M0HostData, testCase.yScaleShape, &yScaleResource.deviceAddr,	152	ret = CreateAclTensor(yScaleE8M0HostData, testCase.yScaleShape, &yScaleResource.deviceAddr,
	@@ -142,24 +163,31 @@ int RunSwigluGroupQuantCase(const SwigluGroupQuantCase& testCase, aclrtStream st
142	CHECK_RET(ret == ACL_SUCCESS,	163	CHECK_RET(ret == ACL_SUCCESS,
143	DestroyAclTensorResource(xResource); DestroyAclTensorResource(yResource);	164	DestroyAclTensorResource(xResource); DestroyAclTensorResource(yResource);
144	DestroyAclTensorResource(yScaleResource); DestroyAclTensorResource(yOriginResource); return ret);	165	DestroyAclTensorResource(yScaleResource); DestroyAclTensorResource(yOriginResource); return ret);
		166	+ if (testCase.hasScaleInput) {
		167	+ ret = CreateAclTensor(scaleHostData, testCase.scaleShape, &scaleResource.deviceAddr, ACL_FLOAT,
		168	+ &scaleResource.tensor);
		169	+ CHECK_RET(ret == ACL_SUCCESS, DestroyAclTensorResource(xResource); DestroyAclTensorResource(yResource);
		170	+ DestroyAclTensorResource(yScaleResource); DestroyAclTensorResource(yOriginResource);
		171	+ DestroyAclTensorResource(scaleResource); return ret);
		172	+ }
145		173
146	- int64_t dstType = 36;
147	double clampLimit = -1.0;	174	double clampLimit = -1.0;
148	- double dstTypeMax = 448.0;
149	bool outputOrigin = false;	175	bool outputOrigin = false;
		176	+ const aclTensor* scaleTensor = testCase.hasScaleInput ? scaleResource.tensor : nullptr;
150		177
151	LOG_PRINT("Run %s: quant_mode=%ld\n", testCase.name, testCase.quantMode);	178	LOG_PRINT("Run %s: quant_mode=%ld\n", testCase.name, testCase.quantMode);
152		179
153	uint64_t workspaceSize = 0;	180	uint64_t workspaceSize = 0;
154	aclOpExecutor* executor = nullptr;	181	aclOpExecutor* executor = nullptr;
155	- ret = aclnnSwigluGroupQuantGetWorkspaceSize(xResource.tensor, nullptr, nullptr, nullptr, dstType, testCase.quantMode,	182	+ ret = aclnnSwigluGroupQuantGetWorkspaceSize(xResource.tensor, nullptr, nullptr, scaleTensor, testCase.dstType,
156	- testCase.blockSize, testCase.roundScale, clampLimit, dstTypeMax,	183	+ testCase.quantMode, testCase.blockSize, testCase.roundScale, clampLimit,
157	- outputOrigin, yResource.tensor, yScaleResource.tensor,	184	+ testCase.dstTypeMax, outputOrigin, yResource.tensor,
158	- yOriginResource.tensor, &workspaceSize, &executor);	185	+ yScaleResource.tensor, yOriginResource.tensor, &workspaceSize, &executor);
159	CHECK_RET(ret == ACL_SUCCESS,	186	CHECK_RET(ret == ACL_SUCCESS,
160	LOG_PRINT("aclnnSwigluGroupQuantGetWorkspaceSize failed. ERROR: %d\n", ret);	187	LOG_PRINT("aclnnSwigluGroupQuantGetWorkspaceSize failed. ERROR: %d\n", ret);
161	DestroyAclTensorResource(xResource); DestroyAclTensorResource(yResource);	188	DestroyAclTensorResource(xResource); DestroyAclTensorResource(yResource);
162	- DestroyAclTensorResource(yScaleResource); DestroyAclTensorResource(yOriginResource); return ret);	189	+ DestroyAclTensorResource(yScaleResource); DestroyAclTensorResource(yOriginResource);
		190	+ DestroyAclTensorResource(scaleResource); return ret);
163		191
164	void* workspaceAddr = nullptr;	192	void* workspaceAddr = nullptr;
165	if (workspaceSize > 0) {	193	if (workspaceSize > 0) {
	@@ -167,7 +195,8 @@ int RunSwigluGroupQuantCase(const SwigluGroupQuantCase& testCase, aclrtStream st
167	CHECK_RET(ret == ACL_SUCCESS,	195	CHECK_RET(ret == ACL_SUCCESS,
168	LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret);	196	LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret);
169	DestroyAclTensorResource(xResource); DestroyAclTensorResource(yResource);	197	DestroyAclTensorResource(xResource); DestroyAclTensorResource(yResource);
170	- DestroyAclTensorResource(yScaleResource); DestroyAclTensorResource(yOriginResource); return ret);	198	+ DestroyAclTensorResource(yScaleResource); DestroyAclTensorResource(yOriginResource);
		199	+ DestroyAclTensorResource(scaleResource); return ret);
171	}	200	}
172		201
173	ret = aclnnSwigluGroupQuant(workspaceAddr, workspaceSize, executor, stream);	202	ret = aclnnSwigluGroupQuant(workspaceAddr, workspaceSize, executor, stream);
	@@ -175,29 +204,33 @@ int RunSwigluGroupQuantCase(const SwigluGroupQuantCase& testCase, aclrtStream st
175	LOG_PRINT("aclnnSwigluGroupQuant failed. ERROR: %d\n", ret);	204	LOG_PRINT("aclnnSwigluGroupQuant failed. ERROR: %d\n", ret);
176	if (workspaceAddr != nullptr) { aclrtFree(workspaceAddr); }	205	if (workspaceAddr != nullptr) { aclrtFree(workspaceAddr); }
177	DestroyAclTensorResource(xResource); DestroyAclTensorResource(yResource);	206	DestroyAclTensorResource(xResource); DestroyAclTensorResource(yResource);
178	- DestroyAclTensorResource(yScaleResource); DestroyAclTensorResource(yOriginResource); return ret);	207	+ DestroyAclTensorResource(yScaleResource); DestroyAclTensorResource(yOriginResource);
		208	+ DestroyAclTensorResource(scaleResource); return ret);
179		209
180	ret = aclrtSynchronizeStream(stream);	210	ret = aclrtSynchronizeStream(stream);
181	CHECK_RET(ret == ACL_SUCCESS,	211	CHECK_RET(ret == ACL_SUCCESS,
182	LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret);	212	LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret);
183	if (workspaceAddr != nullptr) { aclrtFree(workspaceAddr); }	213	if (workspaceAddr != nullptr) { aclrtFree(workspaceAddr); }
184	DestroyAclTensorResource(xResource); DestroyAclTensorResource(yResource);	214	DestroyAclTensorResource(xResource); DestroyAclTensorResource(yResource);
185	- DestroyAclTensorResource(yScaleResource); DestroyAclTensorResource(yOriginResource); return ret);	215	+ DestroyAclTensorResource(yScaleResource); DestroyAclTensorResource(yOriginResource);
		216	+ DestroyAclTensorResource(scaleResource); return ret);
186		217
187	- std::vector<uint8_t> resultData(GetShapeSize(yShape), 0);	218	+ std::vector<uint8_t> resultData(GetShapeSize(yShape) * testCase.yElementSize, 0);
188	ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), yResource.deviceAddr,	219	ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), yResource.deviceAddr,
189	resultData.size() * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST);	220	resultData.size() * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
190	CHECK_RET(ret == ACL_SUCCESS,	221	CHECK_RET(ret == ACL_SUCCESS,
191	LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret);	222	LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret);
192	if (workspaceAddr != nullptr) { aclrtFree(workspaceAddr); }	223	if (workspaceAddr != nullptr) { aclrtFree(workspaceAddr); }
193	DestroyAclTensorResource(xResource); DestroyAclTensorResource(yResource);	224	DestroyAclTensorResource(xResource); DestroyAclTensorResource(yResource);
194	- DestroyAclTensorResource(yScaleResource); DestroyAclTensorResource(yOriginResource); return ret);	225	+ DestroyAclTensorResource(yScaleResource); DestroyAclTensorResource(yOriginResource);
		226	+ DestroyAclTensorResource(scaleResource); return ret);
195	LOG_PRINT("%s result[0] is: %d\n", testCase.name, resultData[0]);	227	LOG_PRINT("%s result[0] is: %d\n", testCase.name, resultData[0]);
196		228
197	DestroyAclTensorResource(xResource);	229	DestroyAclTensorResource(xResource);
198	DestroyAclTensorResource(yResource);	230	DestroyAclTensorResource(yResource);
199	DestroyAclTensorResource(yScaleResource);	231	DestroyAclTensorResource(yScaleResource);
200	DestroyAclTensorResource(yOriginResource);	232	DestroyAclTensorResource(yOriginResource);
		233	+ DestroyAclTensorResource(scaleResource);
201	if (workspaceAddr != nullptr) {	234	if (workspaceAddr != nullptr) {
202	aclrtFree(workspaceAddr);	235	aclrtFree(workspaceAddr);
203	}	236	}
	@@ -217,8 +250,10 @@ int main() {
217	}	250	}
218		251
219	std::vector<SwigluGroupQuantCase> testCases = {	252	std::vector<SwigluGroupQuantCase> testCases = {
220	- {"block_fp8", 0, 0, false, {2, 1}, ACL_FLOAT, false},	253	+ {"block_fp8", 0, 0, false, {2, 1}, ACL_FLOAT, false, 36, 448.0, ACL_FLOAT8_E4M3FN, 1, false, {}},
221	- {"mx_fp8", 1, 0, true, {2, 2, 2}, ACL_FLOAT8_E8M0, true},	254	+ {"mx_fp8", 1, 0, true, {2, 2, 2}, ACL_FLOAT8_E8M0, true, 36, 448.0, ACL_FLOAT8_E4M3FN, 1, false, {}},
		255	+ {"hifp8_static", 2, 0, false, {1}, ACL_FLOAT, false, 27, 448.0, ACL_HIFLOAT8, 1, true, {1}},
		256	+ {"hifp8_dynamic", 3, 0, false, {1}, ACL_FLOAT, false, 27, 15.0, ACL_HIFLOAT8, 1, false, {}},
222	};	257	};
223		258
224	for (const auto& testCase : testCases) {	259	for (const auto& testCase : testCases) {

Mactivation/swiglu_group_quant/op_host/arch35/swiglu_group_quant_hifp8_tiling.cpp+18-26

	@@ -92,18 +92,6 @@ ge::graphStatus SwigluGroupQuantHifp8Tiling::CheckInputDtype()
92	return ge::GRAPH_FAILED);	92	return ge::GRAPH_FAILED);
93	}	93	}
94		94
95	- if (quantMode_ == QUANT_MODE_STATIC) {
96	- auto scaleDesc = context_->GetOptionalInputDesc(INPUT_INDEX_SCALE);
97	- OP_CHECK_IF((scaleDesc == nullptr),
98	- OP_LOGE(context_->GetNodeName(), "scale input is required for static quant mode 2."),
99	- return ge::GRAPH_FAILED);
100	- auto scaleDtype = scaleDesc->GetDataType();
101	- OP_CHECK_IF((scaleDtype != ge::DT_FLOAT),
102	- OP_LOGE(context_->GetNodeName(), "scale dtype only support fp32, got %d.",
103	- static_cast<int>(scaleDtype)),
104	- return ge::GRAPH_FAILED);
105	- }
106	-
107	return ge::GRAPH_SUCCESS;	95	return ge::GRAPH_SUCCESS;
108	}	96	}
109		97
	@@ -115,14 +103,6 @@ ge::graphStatus SwigluGroupQuantHifp8Tiling::CheckOutputDtype()
115	static_cast<int>(yDtype)),	103	static_cast<int>(yDtype)),
116	return ge::GRAPH_FAILED);	104	return ge::GRAPH_FAILED);
117		105
118	- if (quantMode_ == QUANT_MODE_DYNAMIC) {
119	- auto yScaleDtype = context_->GetOutputDesc(OUTPUT_INDEX_Y_SCALE)->GetDataType();
120	- OP_CHECK_IF((yScaleDtype != ge::DT_FLOAT),
121	- OP_LOGE(context_->GetNodeName(), "y_scale dtype must be fp32, got %d.",
122	- static_cast<int>(yScaleDtype)),
123	- return ge::GRAPH_FAILED);
124	- }
125	-
126	return ge::GRAPH_SUCCESS;	106	return ge::GRAPH_SUCCESS;
127	}	107	}
128		108
	@@ -195,10 +175,16 @@ ge::graphStatus SwigluGroupQuantHifp8Tiling::CheckScaleInfo()
195	if (quantMode_ != QUANT_MODE_STATIC) {	175	if (quantMode_ != QUANT_MODE_STATIC) {
196	return ge::GRAPH_SUCCESS;	176	return ge::GRAPH_SUCCESS;
197	}	177	}
198	- auto scaleShape = context_->GetOptionalInputShape(INPUT_INDEX_SCALE);	178	+ auto scaleDesc = context_->GetOptionalInputDesc(INPUT_INDEX_SCALE);
199	- OP_CHECK_IF((scaleShape == nullptr),	179	+ OP_CHECK_IF((scaleDesc == nullptr), OP_LOGE(context_->GetNodeName(), "scale input is required for quant_mode=2."),
200	- OP_LOGE(context_->GetNodeName(), "scale input is required for quant_mode=2."),	180	+ return ge::GRAPH_FAILED);
		181	+ auto scaleDtype = scaleDesc->GetDataType();
		182	+ OP_CHECK_IF((scaleDtype != ge::DT_FLOAT),
		183	+ OP_LOGE(context_->GetNodeName(), "scale dtype only support fp32, got %d.", static_cast<int>(scaleDtype)),
201	return ge::GRAPH_FAILED);	184	return ge::GRAPH_FAILED);
		185	+ auto scaleShape = context_->GetOptionalInputShape(INPUT_INDEX_SCALE);
		186	+ OP_CHECK_IF((scaleShape == nullptr), OP_LOGE(context_->GetNodeName(), "scale shape is null."),
		187	+ return ge::GRAPH_FAILED);
202	auto scaleStorageShape = scaleShape->GetStorageShape();	188	auto scaleStorageShape = scaleShape->GetStorageShape();
203	size_t scaleDimNum = scaleStorageShape.GetDimNum();	189	size_t scaleDimNum = scaleStorageShape.GetDimNum();
204	OP_CHECK_IF((scaleDimNum != 1),	190	OP_CHECK_IF((scaleDimNum != 1),
	@@ -247,10 +233,16 @@ ge::graphStatus SwigluGroupQuantHifp8Tiling::CheckYScaleShape()
247	if (quantMode_ == QUANT_MODE_STATIC) {	233	if (quantMode_ == QUANT_MODE_STATIC) {
248	return ge::GRAPH_SUCCESS;	234	return ge::GRAPH_SUCCESS;
249	}	235	}
		236	+ auto yScaleDesc = context_->GetOutputDesc(OUTPUT_INDEX_Y_SCALE);
		237	+ OP_CHECK_IF((yScaleDesc == nullptr), OP_LOGE(context_->GetNodeName(), "y_scale desc is null."),
		238	+ return ge::GRAPH_FAILED);
		239	+ auto yScaleDtype = yScaleDesc->GetDataType();
		240	+ OP_CHECK_IF((yScaleDtype != ge::DT_FLOAT),
		241	+ OP_LOGE(context_->GetNodeName(), "y_scale dtype must be fp32, got %d.", static_cast<int>(yScaleDtype)),
		242	+ return ge::GRAPH_FAILED);
250	auto yScaleShape = context_->GetOutputShape(OUTPUT_INDEX_Y_SCALE);	243	auto yScaleShape = context_->GetOutputShape(OUTPUT_INDEX_Y_SCALE);
251	- OP_CHECK_IF((yScaleShape == nullptr),	244	+ OP_CHECK_IF((yScaleShape == nullptr), OP_LOGE(context_->GetNodeName(), "y_scale shape is null."),
252	- OP_LOGE(context_->GetNodeName(), "y_scale shape is null."),	245	+ return ge::GRAPH_FAILED);
253	- return ge::GRAPH_FAILED);
254	const gert::Shape &yScaleShapeStorage = yScaleShape->GetStorageShape();	246	const gert::Shape &yScaleShapeStorage = yScaleShape->GetStorageShape();
255	OP_CHECK_IF((yScaleShapeStorage.GetDimNum() != 1),	247	OP_CHECK_IF((yScaleShapeStorage.GetDimNum() != 1),
256	OP_LOGE(context_->GetNodeName(), "y_scale must be 1D, got %zu dims.", yScaleShapeStorage.GetDimNum()),	248	OP_LOGE(context_->GetNodeName(), "y_scale must be 1D, got %zu dims.", yScaleShapeStorage.GetDimNum()),

Mactivation/swiglu_group_quant/op_host/arch35/swiglu_group_quant_tiling.cpp+2-0

	@@ -188,6 +188,8 @@ ge::graphStatus SwigluGroupQuantTiling::GetAttr()
188	if (outputOriginAttr != nullptr) {	188	if (outputOriginAttr != nullptr) {
189	outputOrigin_ = (*outputOriginAttr) ? 1 : 0;	189	outputOrigin_ = (*outputOriginAttr) ? 1 : 0;
190	}	190	}
		191	+ OP_CHECK_IF((outputOrigin_ != 0), OP_LOGE(context_->GetNodeName(), "attr output_origin must be false."),
		192	+ return ge::GRAPH_FAILED);
191		193
192	if (GetClampLimitAttr(attrs) == ge::GRAPH_FAILED) {	194	if (GetClampLimitAttr(attrs) == ge::GRAPH_FAILED) {
193	return ge::GRAPH_FAILED;	195	return ge::GRAPH_FAILED;

Mactivation/swiglu_group_quant/op_kernel/swiglu_group_quant.cpp+5-5

	@@ -41,27 +41,27 @@ extern "C" __global__ __aicore__ void swiglu_group_quant(GM_ADDR x, GM_ADDR weig
41	TPipe pipe;	41	TPipe pipe;
42	int64_t oriOverflowMode = AscendC::GetCtrlSpr<FLOAT_OVERFLOW_MODE_CTRL, FLOAT_OVERFLOW_MODE_CTRL>();	42	int64_t oriOverflowMode = AscendC::GetCtrlSpr<FLOAT_OVERFLOW_MODE_CTRL, FLOAT_OVERFLOW_MODE_CTRL>();
43	if (TILING_KEY_IS(BLOCK_QUANT_TILING_KEY)) {	43	if (TILING_KEY_IS(BLOCK_QUANT_TILING_KEY)) {
44	- GET_TILING_DATA(tilingData, tiling);	44	+ GET_TILING_DATA_WITH_STRUCT(SwigluGroupQuantTilingData, tilingData, tiling);
45	SwigluGroupQuant::SwigluGroupQuantPerf<DTYPE_X, DTYPE_Y, DTYPE_Y_SCALE, false> op;	45	SwigluGroupQuant::SwigluGroupQuantPerf<DTYPE_X, DTYPE_Y, DTYPE_Y_SCALE, false> op;
46	op.Init(x, weight, groupIndex, y, yScale, yOrigin, userWs, &tilingData, &pipe);	46	op.Init(x, weight, groupIndex, y, yScale, yOrigin, userWs, &tilingData, &pipe);
47	op.Process();	47	op.Process();
48	} else if (TILING_KEY_IS(BLOCK_QUANT_YORIGIN_TILING_KEY)) {	48	} else if (TILING_KEY_IS(BLOCK_QUANT_YORIGIN_TILING_KEY)) {
49	- GET_TILING_DATA(tilingData, tiling);	49	+ GET_TILING_DATA_WITH_STRUCT(SwigluGroupQuantTilingData, tilingData, tiling);
50	SwigluGroupQuant::SwigluGroupQuantPerf<DTYPE_X, DTYPE_Y, DTYPE_Y_SCALE, true> op;	50	SwigluGroupQuant::SwigluGroupQuantPerf<DTYPE_X, DTYPE_Y, DTYPE_Y_SCALE, true> op;
51	op.Init(x, weight, groupIndex, y, yScale, yOrigin, userWs, &tilingData, &pipe);	51	op.Init(x, weight, groupIndex, y, yScale, yOrigin, userWs, &tilingData, &pipe);
52	op.Process();	52	op.Process();
53	} else if (TILING_KEY_IS(MX_QUANT_TILING_KEY)) {	53	} else if (TILING_KEY_IS(MX_QUANT_TILING_KEY)) {
54	- GET_TILING_DATA(tilingData, tiling);	54	+ GET_TILING_DATA_WITH_STRUCT(SwigluGroupQuantTilingData, tilingData, tiling);
55	SwigluGroupQuant::SwigluMxQuantPerf<DTYPE_X, DTYPE_Y, DTYPE_Y_SCALE, false> op;	55	SwigluGroupQuant::SwigluMxQuantPerf<DTYPE_X, DTYPE_Y, DTYPE_Y_SCALE, false> op;
56	op.Init(x, weight, groupIndex, y, yScale, yOrigin, userWs, &tilingData, &pipe);	56	op.Init(x, weight, groupIndex, y, yScale, yOrigin, userWs, &tilingData, &pipe);
57	op.Process();	57	op.Process();
58	} else if (TILING_KEY_IS(MX_QUANT_YORIGIN_TILING_KEY)) {	58	} else if (TILING_KEY_IS(MX_QUANT_YORIGIN_TILING_KEY)) {
59	- GET_TILING_DATA(tilingData, tiling);	59	+ GET_TILING_DATA_WITH_STRUCT(SwigluGroupQuantTilingData, tilingData, tiling);
60	SwigluGroupQuant::SwigluMxQuantPerf<DTYPE_X, DTYPE_Y, DTYPE_Y_SCALE, true> op;	60	SwigluGroupQuant::SwigluMxQuantPerf<DTYPE_X, DTYPE_Y, DTYPE_Y_SCALE, true> op;
61	op.Init(x, weight, groupIndex, y, yScale, yOrigin, userWs, &tilingData, &pipe);	61	op.Init(x, weight, groupIndex, y, yScale, yOrigin, userWs, &tilingData, &pipe);
62	op.Process();	62	op.Process();
63	} else if (TILING_KEY_IS(MXFP4_QUANT_TILING_KEY)) {	63	} else if (TILING_KEY_IS(MXFP4_QUANT_TILING_KEY)) {
64	- GET_TILING_DATA(tilingData, tiling);	64	+ GET_TILING_DATA_WITH_STRUCT(SwigluGroupQuantTilingData, tilingData, tiling);
65	SwigluGroupQuant::SwigluMxFp4QuantPerf<DTYPE_X, DTYPE_Y, DTYPE_Y_SCALE> op;	65	SwigluGroupQuant::SwigluMxFp4QuantPerf<DTYPE_X, DTYPE_Y, DTYPE_Y_SCALE> op;
66	op.Init(x, weight, groupIndex, y, yScale, userWs, &tilingData, &pipe);	66	op.Init(x, weight, groupIndex, y, yScale, userWs, &tilingData, &pipe);
67	op.Process();	67	op.Process();

Mactivation/swiglu_group_quant/tests/assets/golden.py+46-12

	@@ -1,3 +1,4 @@
		1	+# -- coding: utf-8 --
1	# Copyright (c) 2026 Huawei Technologies Co., Ltd.	2	# Copyright (c) 2026 Huawei Technologies Co., Ltd.
2	# This program is free software, you can redistribute it and/or modify it under the terms and conditions of	3	# 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	# CANN Open Software License Agreement Version 2.0 (the "License").
	@@ -209,7 +210,9 @@ def _amax_block_fp8(y_block):
209		210
210		211
211	def _amax_mx_fp8(y_block):	212	def _amax_mx_fp8(y_block):
212	- return np.maximum(np.max(np.abs(y_block.astype(np.float32)), axis=-1).astype(np.float32), EPS)	213	+ return np.maximum(
		214	+ np.max(np.abs(y_block.astype(np.float32)), axis=-1).astype(np.float32), EPS
		215	+ )
213		216
214		217
215	def _quantize_fp8(y_origin, dst_type, quant_mode, round_scale):	218	def _quantize_fp8(y_origin, dst_type, quant_mode, round_scale):
	@@ -245,7 +248,9 @@ def _quantize_fp8(y_origin, dst_type, quant_mode, round_scale):
245	else:	248	else:
246	if num_blocks % 2 != 0:	249	if num_blocks % 2 != 0:
247	y_scale = np.pad(y_scale, ((0, 0), (0, 1)), mode="constant")	250	y_scale = np.pad(y_scale, ((0, 0), (0, 1)), mode="constant")
248	- scale = y_scale.view(numpy_float8_e8m0()).reshape(*y_shape[:-1], _ceil_div(num_blocks, MX_SCALE_ALIGN), 2)	251	+ scale = y_scale.view(numpy_float8_e8m0()).reshape(
		252	+ *y_shape[:-1], _ceil_div(num_blocks, MX_SCALE_ALIGN), 2
		253	+ )
249	return y_fp8.reshape(y_shape), scale	254	return y_fp8.reshape(y_shape), scale
250		255
251		256
	@@ -279,17 +284,27 @@ def _quantize_fp4(y_origin, dst_type, pack=True):
279	y_scaled = y_block * inv_scale.reshape(num_tokens, num_blocks, 1)	284	y_scaled = y_block * inv_scale.reshape(num_tokens, num_blocks, 1)
280		285
281	if dst_type == DT_FLOAT4_E2M1:	286	if dst_type == DT_FLOAT4_E2M1:
282	- fp4_values = y_scaled.reshape(num_tokens, hidden).astype(np.float32).astype(_fp4_dtype(dst_type))	287	+ fp4_values = (
		288	+ y_scaled.reshape(num_tokens, hidden)
		289	+ .astype(np.float32)
		290	+ .astype(_fp4_dtype(dst_type))
		291	+ )
283	nibble = fp4_values.view(np.uint8).reshape(num_tokens, hidden)	292	nibble = fp4_values.view(np.uint8).reshape(num_tokens, hidden)
284	elif dst_type == DT_FLOAT4_E1M2:	293	elif dst_type == DT_FLOAT4_E1M2:
285	- nibble = _quantize_e1m2_nibble(y_scaled.astype(np.float32)).reshape(num_tokens, hidden)	294	+ nibble = _quantize_e1m2_nibble(y_scaled.astype(np.float32)).reshape(
		295	+ num_tokens, hidden
		296	+ )
286	fp4_values = nibble.view(_fp4_dtype(dst_type))	297	fp4_values = nibble.view(_fp4_dtype(dst_type))
287	else:	298	else:
288	- raise RuntimeError(f"MxFp4 golden only implements FLOAT4_E2M1/E1M2, got {dst_type}")	299	+ raise RuntimeError(
		300	+ f"MxFp4 golden only implements FLOAT4_E2M1/E1M2, got {dst_type}"
		301	+ )
289		302
290	if num_blocks % 2 != 0:	303	if num_blocks % 2 != 0:
291	e8m0 = np.pad(e8m0, ((0, 0), (0, 1)), mode="constant")	304	e8m0 = np.pad(e8m0, ((0, 0), (0, 1)), mode="constant")
292	- scale = e8m0.view(numpy_float8_e8m0()).reshape(*y_shape[:-1], _ceil_div(num_blocks, 2), 2)	305	+ scale = e8m0.view(numpy_float8_e8m0()).reshape(
		306	+ *y_shape[:-1], _ceil_div(num_blocks, 2), 2
		307	+ )
293	if not pack:	308	if not pack:
294	return fp4_values.reshape(y_shape), scale	309	return fp4_values.reshape(y_shape), scale
295		310
	@@ -351,8 +366,16 @@ def aclnn_swiglu_group_quant_golden(
351	y_origin = _cast_y_origin(y_origin, x)	366	y_origin = _cast_y_origin(y_origin, x)
352	y = _merge_output_buffer(y, yOut, token_num, real_bs)	367	y = _merge_output_buffer(y, yOut, token_num, real_bs)
353	y_scale = _merge_output_buffer(y_scale, yScaleOut, token_num, real_bs)	368	y_scale = _merge_output_buffer(y_scale, yScaleOut, token_num, real_bs)
354	- y_origin = _merge_output_buffer(y_origin, yOriginOut, token_num, real_bs) if outputOrigin else None	369	+ y_origin = (
355	- return _maybe_to_torch(y, use_torch), _maybe_to_torch(y_scale, use_torch), _maybe_to_torch(y_origin, use_torch)	370	+ _merge_output_buffer(y_origin, yOriginOut, token_num, real_bs)
		371	+ if outputOrigin
		372	+ else None
		373	+ )
		374	+ return (
		375	+ _maybe_to_torch(y, use_torch),
		376	+ _maybe_to_torch(y_scale, use_torch),
		377	+ _maybe_to_torch(y_origin, use_torch),
		378	+ )
356		379
357		380
358	def aclnn_swiglu_group_quant_input(	381	def aclnn_swiglu_group_quant_input(
	@@ -379,11 +402,16 @@ def aclnn_swiglu_group_quant_input(
379	x_np = _to_numpy(x)	402	x_np = _to_numpy(x)
380	hidden = x_np.shape[-1] // 2	403	hidden = x_np.shape[-1] // 2
381	stable_x = np.ones(x_np.shape, dtype=np.float32)	404	stable_x = np.ones(x_np.shape, dtype=np.float32)
382	- sign_pattern = np.where(np.arange(hidden) % 2 == 0, 1.0, -1.0).astype(np.float32)	405	+ sign_pattern = np.where(np.arange(hidden) % 2 == 0, 1.0, -1.0).astype(
		406	+ np.float32
		407	+ )
383	stable_x.reshape(-1, x_np.shape[-1])[:, hidden:] = sign_pattern	408	stable_x.reshape(-1, x_np.shape[-1])[:, hidden:] = sign_pattern
384	_write_tensor(x, stable_x)	409	_write_tensor(x, stable_x)
385	if weightOptional is not None:	410	if weightOptional is not None:
386	- _write_tensor(weightOptional, np.ones(_to_numpy(weightOptional).shape, dtype=np.float32))	411	+ _write_tensor(
		412	+ weightOptional,
		413	+ np.ones(_to_numpy(weightOptional).shape, dtype=np.float32),
		414	+ )
387	if groupIndexOptional is None:	415	if groupIndexOptional is None:
388	return	416	return
389	token_num = int(np.prod(_to_numpy(x).shape[:-1]))	417	token_num = int(np.prod(_to_numpy(x).shape[:-1]))
	@@ -398,7 +426,9 @@ def aclnn_swiglu_group_quant_input(
398	return	426	return
399	import torch	427	import torch
400		428
401	- groupIndexOptional.copy_(torch.as_tensor(group_index, device=groupIndexOptional.device))	429	+ groupIndexOptional.copy_(
		430	+ torch.as_tensor(group_index, device=groupIndexOptional.device)
		431	+ )
402		432
403		433
404	def swiglu_group_quant_golden(	434	def swiglu_group_quant_golden(
	@@ -429,7 +459,11 @@ def swiglu_group_quant_golden(
429	y_origin = _cast_y_origin(y_origin, x) if output_origin else None	459	y_origin = _cast_y_origin(y_origin, x) if output_origin else None
430	y = _merge_output_buffer(y, None, token_num, real_bs)	460	y = _merge_output_buffer(y, None, token_num, real_bs)
431	y_scale = _merge_output_buffer(y_scale, None, token_num, real_bs)	461	y_scale = _merge_output_buffer(y_scale, None, token_num, real_bs)
432	- y_origin = _merge_output_buffer(y_origin, None, token_num, real_bs) if output_origin else None	462	+ y_origin = (
		463	+ _merge_output_buffer(y_origin, None, token_num, real_bs)
		464	+ if output_origin
		465	+ else None
		466	+ )
433	return y, y_scale, y_origin	467	return y, y_scale, y_origin
434		468
435		469

Mactivation/swiglu_group_quant/tests/ut/op_host/arch35/test_swiglu_group_quant_tiling.cpp+164-2

	@@ -174,11 +174,12 @@ TEST_F(SwigluGroupQuantTilingTest, tiling_block_fp8)
174	ExecuteTilingCase(tc);	174	ExecuteTilingCase(tc);
175	}	175	}
176		176
177	-TEST_F(SwigluGroupQuantTilingTest, tiling_block_fp8_y_origin)	177	+TEST_F(SwigluGroupQuantTilingTest, tiling_error_block_fp8_y_origin)
178	{	178	{
179	TilingCase tc;	179	TilingCase tc;
180	tc.outputOrigin = true;	180	tc.outputOrigin = true;
181	tc.roundScale = true;	181	tc.roundScale = true;
		182	+ tc.status = ge::GRAPH_FAILED;
182	ExecuteTilingCase(tc);	183	ExecuteTilingCase(tc);
183	}	184	}
184		185
	@@ -199,7 +200,7 @@ TEST_F(SwigluGroupQuantTilingTest, tiling_mx_fp8)
199	ExecuteTilingCase(tc);	200	ExecuteTilingCase(tc);
200	}	201	}
201		202
202	-TEST_F(SwigluGroupQuantTilingTest, tiling_mx_fp8_y_origin)	203	+TEST_F(SwigluGroupQuantTilingTest, tiling_error_mx_fp8_y_origin)
203	{	204	{
204	TilingCase tc;	205	TilingCase tc;
205	tc.scaleDtype = ge::DT_FLOAT8_E8M0;	206	tc.scaleDtype = ge::DT_FLOAT8_E8M0;
	@@ -207,6 +208,7 @@ TEST_F(SwigluGroupQuantTilingTest, tiling_mx_fp8_y_origin)
207	tc.quantMode = 1;	208	tc.quantMode = 1;
208	tc.roundScale = true;	209	tc.roundScale = true;
209	tc.outputOrigin = true;	210	tc.outputOrigin = true;
		211	+ tc.status = ge::GRAPH_FAILED;
210	ExecuteTilingCase(tc);	212	ExecuteTilingCase(tc);
211	}	213	}
212		214
	@@ -314,4 +316,164 @@ TEST_F(SwigluGroupQuantTilingTest, tiling_error_zero_clamp_limit)
314	tc.status = ge::GRAPH_FAILED;	316	tc.status = ge::GRAPH_FAILED;
315	ExecuteTilingCase(tc);	317	ExecuteTilingCase(tc);
316	}	318	}
		319	+
		320	+TEST_F(SwigluGroupQuantTilingTest, tiling_hifp8_static)
		321	+{
		322	+ TilingCase tc;
		323	+ tc.xDtype = ge::DT_FLOAT16;
		324	+ tc.yDtype = ge::DT_HIFLOAT8;
		325	+ tc.scaleDtype = ge::DT_FLOAT;
		326	+ tc.yOriginDtype = ge::DT_FLOAT16;
		327	+ tc.xShape = {{8, 128, 8192}, {8, 128, 8192}};
		328	+ tc.yShape = {{8, 128, 4096}, {8, 128, 4096}};
		329	+ tc.scaleShape = {{1}, {1}};
		330	+ tc.yOriginShape = {{8, 128, 4096}, {8, 128, 4096}};
		331	+ tc.dstType = ge::DT_HIFLOAT8;
		332	+ tc.quantMode = 2;
		333	+ tc.hasScale = true;
		334	+ ExecuteTilingCase(tc);
		335	+}
		336	+
		337	+TEST_F(SwigluGroupQuantTilingTest, tiling_hifp8_static_group)
		338	+{
		339	+ TilingCase tc;
		340	+ tc.xDtype = ge::DT_BF16;
		341	+ tc.yDtype = ge::DT_HIFLOAT8;
		342	+ tc.scaleDtype = ge::DT_FLOAT;
		343	+ tc.yOriginDtype = ge::DT_BF16;
		344	+ tc.xShape = {{8, 128, 8192}, {8, 128, 8192}};
		345	+ tc.groupIndexShape = {{2}, {2}};
		346	+ tc.yShape = {{8, 128, 4096}, {8, 128, 4096}};
		347	+ tc.scaleShape = {{2}, {2}};
		348	+ tc.yOriginShape = {{8, 128, 4096}, {8, 128, 4096}};
		349	+ tc.dstType = ge::DT_HIFLOAT8;
		350	+ tc.quantMode = 2;
		351	+ tc.hasGroupIndex = true;
		352	+ tc.hasScale = true;
		353	+ ExecuteTilingCase(tc);
		354	+}
		355	+
		356	+TEST_F(SwigluGroupQuantTilingTest, tiling_hifp8_dynamic)
		357	+{
		358	+ TilingCase tc;
		359	+ tc.xDtype = ge::DT_FLOAT16;
		360	+ tc.yDtype = ge::DT_HIFLOAT8;
		361	+ tc.scaleDtype = ge::DT_FLOAT;
		362	+ tc.yOriginDtype = ge::DT_FLOAT16;
		363	+ tc.xShape = {{8, 128, 8192}, {8, 128, 8192}};
		364	+ tc.yShape = {{8, 128, 4096}, {8, 128, 4096}};
		365	+ tc.scaleShape = {{1}, {1}};
		366	+ tc.yOriginShape = {{8, 128, 4096}, {8, 128, 4096}};
		367	+ tc.dstType = ge::DT_HIFLOAT8;
		368	+ tc.quantMode = 3;
		369	+ tc.dstTypeMax = 15.0f;
		370	+ ExecuteTilingCase(tc);
		371	+}
		372	+
		373	+TEST_F(SwigluGroupQuantTilingTest, tiling_hifp8_dynamic_group)
		374	+{
		375	+ TilingCase tc;
		376	+ tc.xDtype = ge::DT_BF16;
		377	+ tc.yDtype = ge::DT_HIFLOAT8;
		378	+ tc.scaleDtype = ge::DT_FLOAT;
		379	+ tc.yOriginDtype = ge::DT_BF16;
		380	+ tc.xShape = {{8, 128, 8192}, {8, 128, 8192}};
		381	+ tc.groupIndexShape = {{4}, {4}};
		382	+ tc.yShape = {{8, 128, 4096}, {8, 128, 4096}};
		383	+ tc.scaleShape = {{4}, {4}};
		384	+ tc.yOriginShape = {{8, 128, 4096}, {8, 128, 4096}};
		385	+ tc.dstType = ge::DT_HIFLOAT8;
		386	+ tc.quantMode = 3;
		387	+ tc.dstTypeMax = 15.0f;
		388	+ tc.hasGroupIndex = true;
		389	+ ExecuteTilingCase(tc);
		390	+}
		391	+
		392	+TEST_F(SwigluGroupQuantTilingTest, tiling_hifp8_dynamic_output_origin)
		393	+{
		394	+ TilingCase tc;
		395	+ tc.xDtype = ge::DT_FLOAT16;
		396	+ tc.yDtype = ge::DT_HIFLOAT8;
		397	+ tc.scaleDtype = ge::DT_FLOAT;
		398	+ tc.yOriginDtype = ge::DT_FLOAT16;
		399	+ tc.xShape = {{8, 128, 8192}, {8, 128, 8192}};
		400	+ tc.yShape = {{8, 128, 4096}, {8, 128, 4096}};
		401	+ tc.scaleShape = {{1}, {1}};
		402	+ tc.yOriginShape = {{8, 128, 4096}, {8, 128, 4096}};
		403	+ tc.dstType = ge::DT_HIFLOAT8;
		404	+ tc.quantMode = 3;
		405	+ tc.dstTypeMax = 15.0f;
		406	+ tc.outputOrigin = true;
		407	+ ExecuteTilingCase(tc);
		408	+}
		409	+
		410	+TEST_F(SwigluGroupQuantTilingTest, tiling_error_hifp8_static_without_scale)
		411	+{
		412	+ TilingCase tc;
		413	+ tc.xDtype = ge::DT_FLOAT16;
		414	+ tc.yDtype = ge::DT_HIFLOAT8;
		415	+ tc.scaleDtype = ge::DT_FLOAT;
		416	+ tc.yOriginDtype = ge::DT_FLOAT16;
		417	+ tc.xShape = {{8, 128, 8192}, {8, 128, 8192}};
		418	+ tc.yShape = {{8, 128, 4096}, {8, 128, 4096}};
		419	+ tc.scaleShape = {{1}, {1}};
		420	+ tc.yOriginShape = {{8, 128, 4096}, {8, 128, 4096}};
		421	+ tc.dstType = ge::DT_HIFLOAT8;
		422	+ tc.quantMode = 2;
		423	+ tc.hasScale = false;
		424	+ tc.status = ge::GRAPH_FAILED;
		425	+ ExecuteTilingCase(tc);
		426	+}
		427	+
		428	+TEST_F(SwigluGroupQuantTilingTest, tiling_error_hifp8_invalid_y_dtype)
		429	+{
		430	+ TilingCase tc;
		431	+ tc.xDtype = ge::DT_FLOAT16;
		432	+ tc.yDtype = ge::DT_FLOAT8_E4M3FN;
		433	+ tc.scaleDtype = ge::DT_FLOAT;
		434	+ tc.yOriginDtype = ge::DT_FLOAT16;
		435	+ tc.xShape = {{8, 128, 8192}, {8, 128, 8192}};
		436	+ tc.yShape = {{8, 128, 4096}, {8, 128, 4096}};
		437	+ tc.scaleShape = {{1}, {1}};
		438	+ tc.yOriginShape = {{8, 128, 4096}, {8, 128, 4096}};
		439	+ tc.dstType = ge::DT_HIFLOAT8;
		440	+ tc.quantMode = 3;
		441	+ tc.status = ge::GRAPH_FAILED;
		442	+ ExecuteTilingCase(tc);
		443	+}
		444	+
		445	+TEST_F(SwigluGroupQuantTilingTest, tiling_error_hifp8_static_wrong_scale_dtype)
		446	+{
		447	+ TilingCase tc;
		448	+ tc.xDtype = ge::DT_FLOAT16;
		449	+ tc.yDtype = ge::DT_HIFLOAT8;
		450	+ tc.scaleDtype = ge::DT_FLOAT16;
		451	+ tc.yOriginDtype = ge::DT_FLOAT16;
		452	+ tc.xShape = {{8, 128, 8192}, {8, 128, 8192}};
		453	+ tc.yShape = {{8, 128, 4096}, {8, 128, 4096}};
		454	+ tc.scaleShape = {{1}, {1}};
		455	+ tc.yOriginShape = {{8, 128, 4096}, {8, 128, 4096}};
		456	+ tc.dstType = ge::DT_HIFLOAT8;
		457	+ tc.quantMode = 2;
		458	+ tc.hasScale = true;
		459	+ tc.status = ge::GRAPH_FAILED;
		460	+ ExecuteTilingCase(tc);
		461	+}
		462	+
		463	+TEST_F(SwigluGroupQuantTilingTest, tiling_error_hifp8_dynamic_wrong_yscale_dtype)
		464	+{
		465	+ TilingCase tc;
		466	+ tc.xDtype = ge::DT_FLOAT16;
		467	+ tc.yDtype = ge::DT_HIFLOAT8;
		468	+ tc.scaleDtype = ge::DT_FLOAT8_E8M0;
		469	+ tc.yOriginDtype = ge::DT_FLOAT16;
		470	+ tc.xShape = {{8, 128, 8192}, {8, 128, 8192}};
		471	+ tc.yShape = {{8, 128, 4096}, {8, 128, 4096}};
		472	+ tc.scaleShape = {{1}, {1}};
		473	+ tc.yOriginShape = {{8, 128, 4096}, {8, 128, 4096}};
		474	+ tc.dstType = ge::DT_HIFLOAT8;
		475	+ tc.quantMode = 3;
		476	+ tc.status = ge::GRAPH_FAILED;
		477	+ ExecuteTilingCase(tc);
		478	+}
317	} // namespace	479	} // namespace

Mactivation/swiglu_group_quant/tests/ut/op_kernel/test_swiglu_group_quant.cpp+72-0

	@@ -103,4 +103,76 @@ TEST_F(SwigluGroupQuantKernelTest, block_fp8_y_origin)
103	{	103	{
104	RunKernelWithTilingKey(1100, true);	104	RunKernelWithTilingKey(1100, true);
105	}	105	}
		106	+
		107	+void RunHifp8KernelWithTilingKey(uint64_t tilingKey, bool hasScale, bool outputOrigin)
		108	+{
		109	+ constexpr int64_t totalTokens = 4;
		110	+ constexpr int64_t dimH = 128;
		111	+ constexpr int64_t dim2H = 2 * dimH;
		112	+ constexpr int64_t groupNum = 1;
		113	+ constexpr int64_t usedCoreNum = 2;
		114	+ constexpr int64_t tokensPerCore = totalTokens / usedCoreNum;
		115	+ constexpr int64_t tileLength = tokensPerCore * dimH;
		116	+ constexpr uint32_t blockDim = 2;
		117	+
		118	+ const size_t inputSize = totalTokens * dim2H * sizeof(half);
		119	+ const size_t outputYSize = totalTokens * dimH * sizeof(uint8_t);
		120	+ const size_t outputScaleSize = hasScale ? groupNum * sizeof(float) : 32;
		121	+ const size_t yOriginSize = outputOrigin ? totalTokens * dimH * sizeof(half) : 32;
		122	+ const size_t tilingDataSize = sizeof(SwigluGroupQuantHifp8TilingData);
		123	+
		124	+ const size_t weightSize = totalTokens * sizeof(float);
		125	+
		126	+ uint8_t* x = reinterpret_cast<uint8_t*>(AscendC::GmAlloc(inputSize));
		127	+ uint8_t* weight = reinterpret_cast<uint8_t*>(AscendC::GmAlloc(weightSize));
		128	+ uint8_t* scale = hasScale ? reinterpret_cast<uint8_t*>(AscendC::GmAlloc(outputScaleSize)) : nullptr;
		129	+ uint8_t* y = reinterpret_cast<uint8_t*>(AscendC::GmAlloc(outputYSize));
		130	+ uint8_t* yScale = reinterpret_cast<uint8_t*>(AscendC::GmAlloc(outputScaleSize));
		131	+ uint8_t* yOrigin = outputOrigin ? reinterpret_cast<uint8_t*>(AscendC::GmAlloc(yOriginSize)) : nullptr;
		132	+ uint8_t* workspace = reinterpret_cast<uint8_t*>(AscendC::GmAlloc(32));
		133	+ uint8_t* tiling = reinterpret_cast<uint8_t*>(AscendC::GmAlloc(tilingDataSize));
		134	+
		135	+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
		136	+ auto* tilingData = reinterpret_cast<SwigluGroupQuantHifp8TilingData*>(tiling);
		137	+ tilingData->totalTokens = totalTokens;
		138	+ tilingData->dim2H = dim2H;
		139	+ tilingData->dimH = dimH;
		140	+ tilingData->isGroup = 0;
		141	+ tilingData->hasWeight = 1;
		142	+ tilingData->hasClamp = 0;
		143	+ tilingData->outputOrigin = outputOrigin ? 1 : 0;
		144	+ tilingData->clampLimit = 0.0f;
		145	+ tilingData->dstTypeMax = 15.0f;
		146	+ tilingData->tileTokens = tokensPerCore;
		147	+ tilingData->usedCoreNum = usedCoreNum;
		148	+ tilingData->tokensPerCore = tokensPerCore;
		149	+ tilingData->groupNum = groupNum;
		150	+ tilingData->tileLength = tileLength;
		151	+
		152	+ ICPU_SET_TILING_KEY(tilingKey);
		153	+ auto swigluGroupQuantKernel = [](GM_ADDR x, GM_ADDR weight, GM_ADDR groupIndex, GM_ADDR scale, GM_ADDR y,
		154	+ GM_ADDR yScale, GM_ADDR yOrigin, GM_ADDR workspace, GM_ADDR tiling) {
		155	+ ::swiglu_group_quant(x, weight, groupIndex, scale, y, yScale, yOrigin, workspace, tiling);
		156	+ };
		157	+ ICPU_RUN_KF(swigluGroupQuantKernel, blockDim, x, weight, nullptr, scale, y, yScale, yOrigin, workspace, tiling);
		158	+
		159	+ AscendC::GmFree(x);
		160	+ AscendC::GmFree(weight);
		161	+ if (hasScale)
		162	+ AscendC::GmFree(scale);
		163	+ AscendC::GmFree(y);
		164	+ AscendC::GmFree(yScale);
		165	+ if (outputOrigin)
		166	+ AscendC::GmFree(yOrigin);
		167	+ AscendC::GmFree(workspace);
		168	+ AscendC::GmFree(tiling);
		169	+}
		170	+
		171	+TEST_F(SwigluGroupQuantKernelTest, hifp8_static) { RunHifp8KernelWithTilingKey(4100, true, false); }
		172	+
		173	+TEST_F(SwigluGroupQuantKernelTest, hifp8_static_output_origin) { RunHifp8KernelWithTilingKey(4100, true, true); }
		174	+
		175	+TEST_F(SwigluGroupQuantKernelTest, hifp8_dynamic) { RunHifp8KernelWithTilingKey(4000, false, false); }
		176	+
		177	+TEST_F(SwigluGroupQuantKernelTest, hifp8_dynamic_output_origin) { RunHifp8KernelWithTilingKey(4000, false, true); }
106	} // namespace	178	} // namespace

Mactivation/swiglu_group_quant_grad/README.md+22-12

	@@ -62,13 +62,17 @@
62	$$	62	$$
63	\mathbf{grad}_{x_1}[t, h] = \mathbf{grad}_{y_0}[t, h] \cdot \text{SiLU}(\mathbf{x}_0'[t, h])	63	\mathbf{grad}_{x_1}[t, h] = \mathbf{grad}_{y_0}[t, h] \cdot \text{SiLU}(\mathbf{x}_0'[t, h])
64	$$	64	$$
65	- 其中：如果提供了weight，则 $\mathbf{grad}_{y_0} = \mathbf{grad}_{\text{output}} \cdot \mathbf{weight}$；如果未提供weight，则 $\mathbf{grad}_{y_0} = \mathbf{grad}_{\text{output}}$	65	+ 其中：如果提供了weight，则 $\mathbf{grad}_{y_0} = \mathbf{grad}_{y} \cdot \mathbf{weight}$；如果未提供weight，则 $\mathbf{grad}_{y_0} = \mathbf{grad}_{y}$
66		66
67	- Weight梯度计算公式（可选）：	67	- Weight梯度计算公式（可选）：
68	- $$	68	+ $$
69	- \mathbf{grad}_{\text{weight}}[t] = \sum_{h=0}^{H-1} \mathbf{grad}_{\text{output}}[t, h] \cdot \mathbf{y}_{\text{origin}}[t, h]	69	+ \mathbf{grad}_{\text{weight}}[t] = \sum_{h=0}^{H-1} \mathbf{grad}_{y}[t, h] \cdot \mathbf{y}_{\text{origin}}[t, h]
70	- $$	70	+ $$
71	- 其中：$\mathbf{y}_{\text{origin}}$ 为SwiGLU前向传播的原始激活值输出，沿最后一维（H维度）求和。	71	+ 其中：$\mathbf{y}_{\text{origin}}$ 为SwiGLU前向传播的原始激活值输出，沿最后一维（H维度）求和。
		72	+
		73	+ $$
		74	+ \mathbf{grad}_{\text{weight}}[t] = \mathbf{grad}_{\text{weight}}[t] \cdot \mathbb{I}(t < \text{trunc})
		75	+ $$
72		76
73	- Clamp反向传播掩码公式（当clamp_limit > 0时）：	77	- Clamp反向传播掩码公式（当clamp_limit > 0时）：
74	$$	78	$$
	@@ -148,7 +152,7 @@
148	<tr>	152	<tr>
149	<td>clampLimit</td>	153	<td>clampLimit</td>
150	<td>属性</td>	154	<td>属性</td>
151	- <td><ul><li>Clamp阈值。</li><li>取值范围≥0.0。</li><li>clampLimit=0表示不启用Clamp反向传播掩码。</li></ul></td>	155	+ <td>Clamp阈值。</td>
152	<td>FLOAT</td>	156	<td>FLOAT</td>
153	<td>-</td>	157	<td>-</td>
154	</tr>	158	</tr>
	@@ -170,9 +174,7 @@
170		174
171	## 约束说明	175	## 约束说明
172		176
173	-- 确定性计算：	177	+- 确定性计算：默认确定性实现。
174	- - 当提供 `groupIndex` 参数时：前 trunc 行保证计算结果确定性，后 T-trunc 行保证确定性（填充0）
175	- - 当未提供 `groupIndex` 参数时：所有行数据保证计算结果确定性
176		178
177	- 输入shape约束：	179	- 输入shape约束：
178	- x最后一维必须为偶数（$2H$）	180	- x最后一维必须为偶数（$2H$）
	@@ -181,7 +183,7 @@
181		183
182	- 可选参数约束：	184	- 可选参数约束：
183	- weight提供时，必须同时提供yOrigin才能计算gradWeight	185	- weight提供时，必须同时提供yOrigin才能计算gradWeight
184	- - weight的shape需与gradY的第一维一致	186	+ - weight的shape需与gradY的前n-1维一致
185	- yOrigin的shape需与gradY一致	187	- yOrigin的shape需与gradY一致
186		188
187	- 数据类型约束：	189	- 数据类型约束：
	@@ -190,8 +192,16 @@
190	- groupIndex必须为INT64类型	192	- groupIndex必须为INT64类型
191		193
192	- Clamp约束：	194	- Clamp约束：
193	- - clampLimit必须 ≥ 0.0	195	+ - clampLimit取值范围为-1.0或>0.0
194	- - clampLimit=0表示不启用Clamp反向传播掩码	196	+ - clampLimit=-1.0表示不启用Clamp反向传播掩码，启用时clampLimit必须>0.0
		197	+
		198	+- 规格约束：
		199	+
		200	+ \| 规格项 \| 规格 \| 规格说明 \|
		201	+ \|--------\|------\|----------\|
		202	+ \| B \| 1~31 \| - \|
		203	+ \| S \| 0~128K \| - \|
		204	+ \| H \| 512, 768, 1024, 1536, 1792, 2048, 2560, 4096 \| - \|
195		205
196	## 调用说明	206	## 调用说明
197		207

Mactivation/swiglu_group_quant_grad/docs/aclnnSwigluGroupQuantGrad.md+214-181

	@@ -64,13 +64,17 @@
64	$$	64	$$
65	\mathbf{grad}_{x_1}[t, h] = \mathbf{grad}_{y_0}[t, h] \cdot \text{SiLU}(\mathbf{x}_0'[t, h])	65	\mathbf{grad}_{x_1}[t, h] = \mathbf{grad}_{y_0}[t, h] \cdot \text{SiLU}(\mathbf{x}_0'[t, h])
66	$$	66	$$
67	- 其中：如果提供了weight，则 $\mathbf{grad}_{y_0} = \mathbf{grad}_{\text{output}} \cdot \mathbf{weight}$；如果未提供weight，则 $\mathbf{grad}_{y_0} = \mathbf{grad}_{\text{output}}$	67	+ 其中：如果提供了weight，则 $\mathbf{grad}_{y_0} = \mathbf{grad}_{y} \cdot \mathbf{weight}$；如果未提供weight，则 $\mathbf{grad}_{y_0} = \mathbf{grad}_{y}$
68		68
69	- Weight梯度计算公式（可选）：	69	- Weight梯度计算公式（可选）：
70	- $$	70	+ $$
71	- \mathbf{grad}_{\text{weight}}[t] = \sum_{h=0}^{H-1} \mathbf{grad}_{\text{output}}[t, h] \cdot \mathbf{y}_{\text{origin}}[t, h]	71	+ \mathbf{grad}_{\text{weight}}[t] = \sum_{h=0}^{H-1} \mathbf{grad}_{y}[t, h] \cdot \mathbf{y}_{\text{origin}}[t, h]
72	- $$	72	+ $$
73	- 其中：$\mathbf{y}_{\text{origin}}$ 为SwiGLU前向传播的原始激活值输出，沿最后一维（H维度）求和。	73	+ 其中：$\mathbf{y}_{\text{origin}}$ 为SwiGLU前向传播的原始激活值输出，沿最后一维（H维度）求和。
		74	+
		75	+ $$
		76	+ \mathbf{grad}_{\text{weight}}[t] = \mathbf{grad}_{\text{weight}}[t] \cdot \mathbb{I}(t < \text{trunc})
		77	+ $$
74		78
75	- Clamp反向传播掩码公式（当clamp_limit > 0时）：	79	- Clamp反向传播掩码公式（当clamp_limit > 0时）：
76	$$	80	$$
	@@ -104,7 +108,7 @@ aclnnStatus aclnnSwigluGroupQuantGradGetWorkspaceSize(
104	const aclTensor *x,	108	const aclTensor *x,
105	const aclTensor *weightOptional,	109	const aclTensor *weightOptional,
106	const aclTensor *yOriginOptional,	110	const aclTensor *yOriginOptional,
107	- const aclIntArray *groupIndexOptional,	111	+ const aclTensor *groupIndexOptional,
108	double clampLimit,	112	double clampLimit,
109	const aclTensor *gradXOut,	113	const aclTensor *gradXOut,
110	const aclTensor *gradWeightOutOptional,	114	const aclTensor *gradWeightOutOptional,
	@@ -154,7 +158,7 @@ aclnnStatus aclnnSwigluGroupQuantGrad(
154	<td>BFLOAT16、FLOAT16、FLOAT</td>	158	<td>BFLOAT16、FLOAT16、FLOAT</td>
155	<td>ND</td>	159	<td>ND</td>
156	<td>2-3</td>	160	<td>2-3</td>
157	- <td>√</td>	161	+ <td>×</td>
158	</tr>	162	</tr>
159	<tr>	163	<tr>
160	<td>x（aclTensor*）</td>	164	<td>x（aclTensor*）</td>
	@@ -164,17 +168,17 @@ aclnnStatus aclnnSwigluGroupQuantGrad(
164	<td>BFLOAT16、FLOAT16、FLOAT</td>	168	<td>BFLOAT16、FLOAT16、FLOAT</td>
165	<td>ND</td>	169	<td>ND</td>
166	<td>2-3</td>	170	<td>2-3</td>
167	- <td>√</td>	171	+ <td>×</td>
168	</tr>	172	</tr>
169	<tr>	173	<tr>
170	<td>weightOptional（aclTensor*）</td>	174	<td>weightOptional（aclTensor*）</td>
171	<td>输入（可选）</td>	175	<td>输入（可选）</td>
172	<td>MoE权重张量。</td>	176	<td>MoE权重张量。</td>
173	- <td><ul><li>shape=[T, 1]或[B, S, 1]，需与gradY的第一维一致。</li><li>当提供weight时，必须同时提供yOrigin才能计算gradWeight。</li></ul></td>	177	+ <td><ul><li>shape=[T, 1]或[B, S, 1]，需与gradY的第一维或前两维一致。</li><li>当提供weight时，必须同时提供yOrigin才能计算gradWeight。</li></ul></td>
174	<td>FLOAT</td>	178	<td>FLOAT</td>
175	<td>ND</td>	179	<td>ND</td>
176	<td>2-3</td>	180	<td>2-3</td>
177	- <td>√</td>	181	+ <td>×</td>
178	</tr>	182	</tr>
179	<tr>	183	<tr>
180	<td>yOriginOptional（aclTensor*）</td>	184	<td>yOriginOptional（aclTensor*）</td>
	@@ -184,23 +188,23 @@ aclnnStatus aclnnSwigluGroupQuantGrad(
184	<td>BFLOAT16、FLOAT16、FLOAT</td>	188	<td>BFLOAT16、FLOAT16、FLOAT</td>
185	<td>ND</td>	189	<td>ND</td>
186	<td>2-3</td>	190	<td>2-3</td>
187	- <td>√</td>	191	+ <td>×</td>
188	</tr>	192	</tr>
189	<tr>	193	<tr>
190	<td>groupIndexOptional（aclTensor*）</td>	194	<td>groupIndexOptional（aclTensor*）</td>
191	<td>输入（可选）</td>	195	<td>输入（可选）</td>
192	<td>GroupIndex张量，动态核分配。</td>	196	<td>GroupIndex张量，动态核分配。</td>
193	- <td><ul><li>shape=[G]，dtype=INT64。</li><li>G为MoE专家分组数。</li><li>groupIndex内元素要求为非递减。</li></ul></td>	197	+ <td><ul><li>shape=[G]，dtype=INT64。</li><li>G为MoE专家分组数。</li><li>groupIndex内元素为count模式的group token数。</li></ul></td>
194	<td>INT64</td>	198	<td>INT64</td>
195	<td>ND</td>	199	<td>ND</td>
196	<td>1</td>	200	<td>1</td>
197	- <td>√</td>	201	+ <td>×</td>
198	</tr>	202	</tr>
199	<tr>	203	<tr>
200	<td>clampLimit（float）</td>	204	<td>clampLimit（float）</td>
201	<td>输入</td>	205	<td>输入</td>
202	<td>Clamp阈值。</td>	206	<td>Clamp阈值。</td>
203	- <td><ul><li>取值范围≥0.0。</li><li>clampLimit=0表示不启用Clamp反向传播掩码。</li></ul></td>	207	+ <td><ul><li>取值范围为-1.0或>0.0。</li><li>clampLimit=-1.0表示不启用Clamp反向传播掩码，启用时clampLimit必须>0.0。</li></ul></td>
204	<td>FLOAT</td>	208	<td>FLOAT</td>
205	<td>-</td>	209	<td>-</td>
206	<td>-</td>	210	<td>-</td>
	@@ -214,7 +218,7 @@ aclnnStatus aclnnSwigluGroupQuantGrad(
214	<td>BFLOAT16、FLOAT16、FLOAT</td>	218	<td>BFLOAT16、FLOAT16、FLOAT</td>
215	<td>ND</td>	219	<td>ND</td>
216	<td>2-3</td>	220	<td>2-3</td>
217	- <td>√</td>	221	+ <td>×</td>
218	</tr>	222	</tr>
219	<tr>	223	<tr>
220	<td>gradWeightOutOptional（aclTensor*）</td>	224	<td>gradWeightOutOptional（aclTensor*）</td>
	@@ -224,7 +228,7 @@ aclnnStatus aclnnSwigluGroupQuantGrad(
224	<td>FLOAT</td>	228	<td>FLOAT</td>
225	<td>ND</td>	229	<td>ND</td>
226	<td>2-3</td>	230	<td>2-3</td>
227	- <td>√</td>	231	+ <td>×</td>
228	</tr>	232	</tr>
229	<tr>	233	<tr>
230	<td>workspaceSize（uint64_t*）</td>	234	<td>workspaceSize（uint64_t*）</td>
	@@ -339,7 +343,7 @@ aclnnStatus aclnnSwigluGroupQuantGrad(
339		343
340	- 可选参数约束：	344	- 可选参数约束：
341	- weight提供时，必须同时提供yOrigin才能计算gradWeight	345	- weight提供时，必须同时提供yOrigin才能计算gradWeight
342	- - weight的shape需与gradY的第一维一致	346	+ - weight的shape需与gradY的前n-1维一致
343	- yOrigin的shape需与gradY一致	347	- yOrigin的shape需与gradY一致
344		348
345	- 数据类型约束：	349	- 数据类型约束：
	@@ -348,225 +352,254 @@ aclnnStatus aclnnSwigluGroupQuantGrad(
348	- groupIndex必须为INT64类型	352	- groupIndex必须为INT64类型
349		353
350	- Clamp约束：	354	- Clamp约束：
351	- - clampLimit必须 ≥ 0.0	355	+ - clampLimit取值范围为-1.0或>0.0
352	- - clampLimit=0表示不启用Clamp反向传播掩码	356	+ - clampLimit=-1.0表示不启用Clamp反向传播掩码，启用时clampLimit必须>0.0
		357	+
		358	+- 规格约束：
		359	+
		360	+ \| 规格项 \| 规格 \| 规格说明 \|
		361	+ \|--------\|------\|----------\|
		362	+ \| B \| 1~31 \| - \|
		363	+ \| S \| 0~128K \| - \|
		364	+ \| H \| 512, 768, 1024, 1536, 1792, 2048, 2560, 4096 \| - \|
353		365
354	## 调用示例	366	## 调用示例
355		367
356	示例代码如下，仅供参考，具体编译和执行过程请参考[编译与运行样例](../../../docs/zh/context/编译与运行样例.md)。	368	示例代码如下，仅供参考，具体编译和执行过程请参考[编译与运行样例](../../../docs/zh/context/编译与运行样例.md)。
357		369
358	```Cpp	370	```Cpp
		371	+/**
		372	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
		373	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
		374	+ * CANN Open Software License Agreement Version 2.0 (the "License").
		375	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
		376	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
		377	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
		378	+ * See LICENSE in the root of the software repository for the full text of the License.
		379	+ */
		380	+
359	#include <iostream>	381	#include <iostream>
360	#include <vector>	382	#include <vector>
361	#include "acl/acl.h"	383	#include "acl/acl.h"
362	#include "aclnnop/aclnn_swiglu_group_quant_grad.h"	384	#include "aclnnop/aclnn_swiglu_group_quant_grad.h"
363		385
364	#define CHECK_RET(cond, return_expr) \	386	#define CHECK_RET(cond, return_expr) \
365	- do { \	387	+ do { \
366	- if (!(cond)) { \	388	+ if (!(cond)) { \
367	- return_expr; \	389	+ return_expr; \
368	- } \	390	+ } \
369	- } while (0)	391	+ } while (0)
370		392
371	-#define LOG_PRINT(message, ...) \	393	+#define LOG_PRINT(message, ...) \
372	- do { \	394	+ do { \
373	- printf(message, ##__VA_ARGS__); \	395	+ printf(message, ##__VA_ARGS__); \
374	- } while (0)	396	+ } while (0)
375		397
376	-int64_t GetShapeSize(const std::vector<int64_t>& shape) {	398	+int64_t GetShapeSize(const std::vector<int64_t>& shape)
377	- int64_t shapeSize = 1;	399	+{
378	- for (auto i : shape) {	400	+ int64_t shapeSize = 1;
379	- shapeSize *= i;	401	+ for (auto i : shape) {
380	- }	402	+ shapeSize *= i;
381	- return shapeSize;	403	+ }
		404	+ return shapeSize;
382	}	405	}
383		406
384	-int Init(int32_t deviceId, aclrtStream* stream) {	407	+int Init(int32_t deviceId, aclrtStream* stream)
385	- auto ret = aclInit(nullptr);	408	+{
386	- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);	409	+ auto ret = aclInit(nullptr);
387	- ret = aclrtSetDevice(deviceId);	410	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
388	- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);	411	+ ret = aclrtSetDevice(deviceId);
389	- ret = aclrtCreateStream(stream);	412	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
390	- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);	413	+ ret = aclrtCreateStream(stream);
391	- return 0;	414	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
		415	+ return 0;
392	}	416	}
393		417
394	template <typename T>	418	template <typename T>
395	int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,	419	int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
396	- aclDataType dataType, aclTensor** tensor) {	420	+ aclDataType dataType, aclTensor** tensor)
397	- auto size = GetShapeSize(shape) * sizeof(T);	421	+{
398	- auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);	422	+ auto size = GetShapeSize(shape) * sizeof(T);
399	- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);	423	+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
400	- ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);	424	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
401	- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);	425	+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
		426	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
402		427
403	- std::vector<int64_t> strides(shape.size(), 1);	428	+ std::vector<int64_t> strides(shape.size(), 1);
404	- for (int64_t i = shape.size() - 2; i >= 0; i--) {	429	+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
405	- strides[i] = shape[i + 1] * strides[i + 1];	430	+ strides[i] = shape[i + 1] * strides[i + 1];
406	- }	431	+ }
407		432
408	- *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,	433	+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
409	- shape.data(), shape.size(), *deviceAddr);	434	+ shape.data(), shape.size(), *deviceAddr);
410	- return 0;	435	+ return 0;
411	}	436	}
412		437
413	template <typename T>	438	template <typename T>
414	-int CreateAclTensorWithValue(const std::vector<int64_t>& shape, void** deviceAddr,	439	+int CreateAclTensorWithValue(const std::vector<int64_t>& shape, void** deviceAddr, aclDataType dataType,
415	- aclDataType dataType, aclTensor** tensor, T value) {	440	+ aclTensor** tensor, T value)
416	- int64_t shapeSize = GetShapeSize(shape);	441	+{
417	- std::vector<T> hostData(shapeSize, value);	442	+ int64_t shapeSize = GetShapeSize(shape);
418	- return CreateAclTensor(hostData, shape, deviceAddr, dataType, tensor);	443	+ std::vector<T> hostData(shapeSize, value);
		444	+ return CreateAclTensor(hostData, shape, deviceAddr, dataType, tensor);
419	}	445	}
420		446
421	-int main() {	447	+int main()
422	- int32_t deviceId = 0;	448	+{
423	- aclrtStream stream;	449	+ int32_t deviceId = 0;
424	- auto ret = Init(deviceId, &stream);	450	+ aclrtStream stream;
425	- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);	451	+ auto ret = Init(deviceId, &stream);
		452	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
426		453
427	- std::vector<int64_t> gradYShape = {512, 512};	454	+ std::vector<int64_t> gradYShape = {512, 512};
428	- std::vector<int64_t> xShape = {512, 1024};	455	+ std::vector<int64_t> xShape = {512, 1024};
429	- std::vector<int64_t> weightShape = {512, 1};	456	+ std::vector<int64_t> weightShape = {512, 1};
430	- std::vector<int64_t> yOriginShape = {512, 512};	457	+ std::vector<int64_t> yOriginShape = {512, 512};
431	- std::vector<int64_t> groupIndexShape = {256};	458	+ std::vector<int64_t> groupIndexShape = {256};
432	- std::vector<int64_t> gradXShape = {512, 1024};	459	+ std::vector<int64_t> gradXShape = {512, 1024};
433	- std::vector<int64_t> gradWeightShape = {512, 1};	460	+ std::vector<int64_t> gradWeightShape = {512, 1};
434		461
435	- void* gradYDeviceAddr = nullptr;	462	+ void* gradYDeviceAddr = nullptr;
436	- void* xDeviceAddr = nullptr;	463	+ void* xDeviceAddr = nullptr;
437	- void* weightDeviceAddr = nullptr;	464	+ void* weightDeviceAddr = nullptr;
438	- void* yOriginDeviceAddr = nullptr;	465	+ void* yOriginDeviceAddr = nullptr;
439	- void* groupIndexDeviceAddr = nullptr;	466	+ void* groupIndexDeviceAddr = nullptr;
440	- void* gradXDeviceAddr = nullptr;	467	+ void* gradXDeviceAddr = nullptr;
441	- void* gradWeightDeviceAddr = nullptr;	468	+ void* gradWeightDeviceAddr = nullptr;
442		469
443	- aclTensor* gradYTensor = nullptr;	470	+ aclTensor* gradYTensor = nullptr;
444	- aclTensor* xTensor = nullptr;	471	+ aclTensor* xTensor = nullptr;
445	- aclTensor* weightTensor = nullptr;	472	+ aclTensor* weightTensor = nullptr;
446	- aclTensor* yOriginTensor = nullptr;	473	+ aclTensor* yOriginTensor = nullptr;
447	- aclIntArray* groupIndexArray = nullptr;	474	+ aclTensor* groupIndexTensor = nullptr;
448	- aclTensor* gradXTensor = nullptr;	475	+ aclTensor* gradXTensor = nullptr;
449	- aclTensor* gradWeightTensor = nullptr;	476	+ aclTensor* gradWeightTensor = nullptr;
450		477
451	- int64_t gradYSize = GetShapeSize(gradYShape);	478	+ int64_t gradYSize = GetShapeSize(gradYShape);
452	- std::vector<float> gradYHostData(gradYSize, 1.0f);	479	+ std::vector<float> gradYHostData(gradYSize, 1.0f);
453	- for (int64_t i = 0; i < gradYSize; i++) {	480	+ for (int64_t i = 0; i < gradYSize; i++) {
454	- gradYHostData[i] = static_cast<float>(i % 10) * 0.1f;	481	+ gradYHostData[i] = static_cast<float>(i % 10) * 0.1f;
455	- }	482	+ }
456		483
457	- int64_t xSize = GetShapeSize(xShape);	484	+ int64_t xSize = GetShapeSize(xShape);
458	- std::vector<float> xHostData(xSize, 1.0f);	485	+ std::vector<float> xHostData(xSize, 1.0f);
459	- for (int64_t i = 0; i < xSize; i++) {	486	+ for (int64_t i = 0; i < xSize; i++) {
460	- xHostData[i] = static_cast<float>((i % 20) - 10) * 0.5f;	487	+ xHostData[i] = static_cast<float>((i % 20) - 10) * 0.5f;
461	- }	488	+ }
462		489
463	- int64_t weightSize = GetShapeSize(weightShape);	490	+ int64_t weightSize = GetShapeSize(weightShape);
464	- std::vector<float> weightHostData(weightSize, 1.0f);	491	+ std::vector<float> weightHostData(weightSize, 1.0f);
465	- for (int64_t i = 0; i < weightSize; i++) {	492	+ for (int64_t i = 0; i < weightSize; i++) {
466	- weightHostData[i] = static_cast<float>((i % 5) + 1) * 0.2f;	493	+ weightHostData[i] = static_cast<float>((i % 5) + 1) * 0.2f;
467	- }	494	+ }
468		495
469	- int64_t yOriginSize = GetShapeSize(yOriginShape);	496	+ int64_t yOriginSize = GetShapeSize(yOriginShape);
470	- std::vector<float> yOriginHostData(yOriginSize, 1.0f);	497	+ std::vector<float> yOriginHostData(yOriginSize, 1.0f);
471	- for (int64_t i = 0; i < yOriginSize; i++) {	498	+ for (int64_t i = 0; i < yOriginSize; i++) {
472	- yOriginHostData[i] = static_cast<float>((i % 8) + 1) * 0.3f;	499	+ yOriginHostData[i] = static_cast<float>((i % 8) + 1) * 0.3f;
473	- }	500	+ }
474		501
475	- int64_t groupIndexSize = GetShapeSize(groupIndexShape);	502	+ int64_t groupIndexSize = GetShapeSize(groupIndexShape);
476	- std::vector<int64_t> groupIndexHostData(groupIndexSize, 0);	503	+ std::vector<int64_t> groupIndexHostData(groupIndexSize, 0);
477	- int64_t groupStride = 512 / 256;	504	+ int64_t groupStride = 512 / 256;
478	- for (int64_t i = 0; i < groupIndexSize; i++) {	505	+ for (int64_t i = 0; i < groupIndexSize; i++) {
479	- groupIndexHostData[i] = i * groupStride;	506	+ groupIndexHostData[i] = i * groupStride;
480	- }	507	+ }
481		508
482	- ret = CreateAclTensor(gradYHostData, gradYShape, &gradYDeviceAddr, aclDataType::ACL_FLOAT16, &gradYTensor);	509	+ ret = CreateAclTensor(gradYHostData, gradYShape, &gradYDeviceAddr, aclDataType::ACL_FLOAT16, &gradYTensor);
483	- CHECK_RET(ret == ACL_SUCCESS, return ret);	510	+ CHECK_RET(ret == ACL_SUCCESS, return ret);
484		511
485	- ret = CreateAclTensor(xHostData, xShape, &xDeviceAddr, aclDataType::ACL_FLOAT16, &xTensor);	512	+ ret = CreateAclTensor(xHostData, xShape, &xDeviceAddr, aclDataType::ACL_FLOAT16, &xTensor);
486	- CHECK_RET(ret == ACL_SUCCESS, return ret);	513	+ CHECK_RET(ret == ACL_SUCCESS, return ret);
487		514
488	- ret = CreateAclTensor(weightHostData, weightShape, &weightDeviceAddr, aclDataType::ACL_FLOAT, &weightTensor);	515	+ ret = CreateAclTensor(weightHostData, weightShape, &weightDeviceAddr, aclDataType::ACL_FLOAT, &weightTensor);
489	- CHECK_RET(ret == ACL_SUCCESS, return ret);	516	+ CHECK_RET(ret == ACL_SUCCESS, return ret);
490		517
491	- ret = CreateAclTensor(yOriginHostData, yOriginShape, &yOriginDeviceAddr, aclDataType::ACL_FLOAT16, &yOriginTensor);	518	+ ret = CreateAclTensor(yOriginHostData, yOriginShape, &yOriginDeviceAddr, aclDataType::ACL_FLOAT16, &yOriginTensor);
492	- CHECK_RET(ret == ACL_SUCCESS, return ret);	519	+ CHECK_RET(ret == ACL_SUCCESS, return ret);
493		520
494	- std::vector<int64_t> groupArray = {256, 256};	521	+ ret = CreateAclTensor(groupIndexHostData, groupIndexShape, &groupIndexDeviceAddr, aclDataType::ACL_INT64,
495	- groupIndexArray = aclCreateIntArray(groupArray.data(), groupArray.size());	522	+ &groupIndexTensor);
496	- CHECK_RET(ret == ACL_SUCCESS, return ret);	523	+ CHECK_RET(ret == ACL_SUCCESS, return ret);
497		524
498	- ret = CreateAclTensorWithValue<float>(gradXShape, &gradXDeviceAddr, aclDataType::ACL_FLOAT16, &gradXTensor, 0.0f);	525	+ ret = CreateAclTensorWithValue<float>(gradXShape, &gradXDeviceAddr, aclDataType::ACL_FLOAT16, &gradXTensor, 0.0f);
499	- CHECK_RET(ret == ACL_SUCCESS, return ret);	526	+ CHECK_RET(ret == ACL_SUCCESS, return ret);
500		527
501	- ret = CreateAclTensorWithValue<float>(gradWeightShape, &gradWeightDeviceAddr, aclDataType::ACL_FLOAT, &gradWeightTensor, 0.0f);	528	+ ret = CreateAclTensorWithValue<float>(gradWeightShape, &gradWeightDeviceAddr, aclDataType::ACL_FLOAT,
502	- CHECK_RET(ret == ACL_SUCCESS, return ret);	529	+ &gradWeightTensor, 0.0f);
		530	+ CHECK_RET(ret == ACL_SUCCESS, return ret);
503		531
504	- float clampLimit = 1.0f;	532	+ float clampLimit = 1.0f;
505		533
506	- uint64_t workspaceSize = 0;	534	+ uint64_t workspaceSize = 0;
507	- aclOpExecutor* executor;	535	+ aclOpExecutor* executor;
508		536
509	- ret = aclnnSwigluGroupQuantGradGetWorkspaceSize(gradYTensor, xTensor, weightTensor, yOriginTensor,	537	+ ret = aclnnSwigluGroupQuantGradGetWorkspaceSize(gradYTensor, xTensor, weightTensor, yOriginTensor, groupIndexTensor,
510	- groupIndexArray, clampLimit, gradXTensor, gradWeightTensor,	538	+ clampLimit, gradXTensor, gradWeightTensor, &workspaceSize,
511	- &workspaceSize, &executor);	539	+ &executor);
512	- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSwigluGroupQuantGradGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);	540	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSwigluGroupQuantGradGetWorkspaceSize failed. ERROR: %d\n", ret);
		541	+ return ret);
513		542
514	- void* workspaceAddr = nullptr;	543	+ void* workspaceAddr = nullptr;
515	- if (workspaceSize > 0) {	544	+ if (workspaceSize > 0) {
516	- ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);	545	+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
517	- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);	546	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
518	- }	547	+ }
519		548
520	- ret = aclnnSwigluGroupQuantGrad(workspaceAddr, workspaceSize, executor, stream);	549	+ ret = aclnnSwigluGroupQuantGrad(workspaceAddr, workspaceSize, executor, stream);
521	- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSwigluGroupQuantGrad failed. ERROR: %d\n", ret); return ret);	550	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSwigluGroupQuantGrad failed. ERROR: %d\n", ret); return ret);
522		551
523	- ret = aclrtSynchronizeStream(stream);	552	+ ret = aclrtSynchronizeStream(stream);
524	- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);	553	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
525		554
526	- auto gradXResultSize = GetShapeSize(gradXShape);	555	+ auto gradXResultSize = GetShapeSize(gradXShape);
527	- std::vector<float> gradXResultData(gradXResultSize, 0);	556	+ std::vector<float> gradXResultData(gradXResultSize, 0);
528	- ret = aclrtMemcpy(gradXResultData.data(), gradXResultData.size() * sizeof(float),	557	+ ret = aclrtMemcpy(gradXResultData.data(), gradXResultData.size() * sizeof(float), gradXDeviceAddr,
529	- gradXDeviceAddr, gradXResultSize * sizeof(float), ACL_MEMCPY_DEVICE_TO_HOST);	558	+ gradXResultSize * sizeof(float), ACL_MEMCPY_DEVICE_TO_HOST);
530	- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy gradX result from device to host failed. ERROR: %d\n", ret); return ret);	559	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy gradX result from device to host failed. ERROR: %d\n", ret);
		560	+ return ret);
531		561
532	- LOG_PRINT("gradX output (first 10 elements):\n");	562	+ LOG_PRINT("gradX output (first 10 elements):\n");
533	- for (int64_t i = 0; i < 10 && i < gradXResultSize; i++) {	563	+ for (int64_t i = 0; i < 10 && i < gradXResultSize; i++) {
534	- LOG_PRINT("gradX[%ld] = %f\n", i, gradXResultData[i]);	564	+ LOG_PRINT("gradX[%ld] = %f\n", i, gradXResultData[i]);
535	- }	565	+ }
536		566
537	- auto gradWeightResultSize = GetShapeSize(gradWeightShape);	567	+ auto gradWeightResultSize = GetShapeSize(gradWeightShape);
538	- std::vector<float> gradWeightResultData(gradWeightResultSize, 0);	568	+ std::vector<float> gradWeightResultData(gradWeightResultSize, 0);
539	- ret = aclrtMemcpy(gradWeightResultData.data(), gradWeightResultData.size() * sizeof(float),	569	+ ret = aclrtMemcpy(gradWeightResultData.data(), gradWeightResultData.size() * sizeof(float), gradWeightDeviceAddr,
540	- gradWeightDeviceAddr, gradWeightResultSize * sizeof(float), ACL_MEMCPY_DEVICE_TO_HOST);	570	+ gradWeightResultSize * sizeof(float), ACL_MEMCPY_DEVICE_TO_HOST);
541	- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy gradWeight result from device to host failed. ERROR: %d\n", ret); return ret);	571	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy gradWeight result from device to host failed. ERROR: %d\n", ret);
		572	+ return ret);
542		573
543	- LOG_PRINT("gradWeight output (first 10 elements):\n");	574	+ LOG_PRINT("gradWeight output (first 10 elements):\n");
544	- for (int64_t i = 0; i < 10 && i < gradWeightResultSize; i++) {	575	+ for (int64_t i = 0; i < 10 && i < gradWeightResultSize; i++) {
545	- LOG_PRINT("gradWeight[%ld] = %f\n", i, gradWeightResultData[i]);	576	+ LOG_PRINT("gradWeight[%ld] = %f\n", i, gradWeightResultData[i]);
546	- }	577	+ }
547		578
548	- aclDestroyTensor(gradYTensor);	579	+ aclDestroyTensor(gradYTensor);
549	- aclDestroyTensor(xTensor);	580	+ aclDestroyTensor(xTensor);
550	- aclDestroyTensor(weightTensor);	581	+ aclDestroyTensor(weightTensor);
551	- aclDestroyTensor(yOriginTensor);	582	+ aclDestroyTensor(yOriginTensor);
552	- aclDestroyTensor(gradXTensor);	583	+ aclDestroyTensor(groupIndexTensor);
553	- aclDestroyTensor(gradWeightTensor);	584	+ aclDestroyTensor(gradXTensor);
		585	+ aclDestroyTensor(gradWeightTensor);
554		586
555	- aclrtFree(gradYDeviceAddr);	587	+ aclrtFree(gradYDeviceAddr);
556	- aclrtFree(xDeviceAddr);	588	+ aclrtFree(xDeviceAddr);
557	- aclrtFree(weightDeviceAddr);	589	+ aclrtFree(weightDeviceAddr);
558	- aclrtFree(yOriginDeviceAddr);	590	+ aclrtFree(yOriginDeviceAddr);
559	- aclrtFree(groupIndexDeviceAddr);	591	+ aclrtFree(groupIndexDeviceAddr);
560	- aclrtFree(gradXDeviceAddr);	592	+ aclrtFree(gradXDeviceAddr);
561	- aclrtFree(gradWeightDeviceAddr);	593	+ aclrtFree(gradWeightDeviceAddr);
562	- if (workspaceSize > 0) {	594	+ if (workspaceSize > 0) {
563	- aclrtFree(workspaceAddr);	595	+ aclrtFree(workspaceAddr);
564	- }	596	+ }
565		597
566	- aclrtDestroyStream(stream);	598	+ aclrtDestroyStream(stream);
567	- aclrtResetDevice(deviceId);	599	+ aclrtResetDevice(deviceId);
568	- aclFinalize();	600	+ aclFinalize();
569		601
570	- return 0;	602	+ return 0;
571	}	603	}
572	```	604	```
		605	+