Star926
1847
代码
代码
Issues233
Pull Requests357
Actions
项目讨论
Wiki
分析
Star926
1847
已合并
新增SwigluGroupQuantGrad算子 #6281
shilulu创建于 6月18日
新增SwigluGroupQuantGrad算子 #6281
已合并
shilulu创建于 6月18日
从shilulu/ops-nn_40249.1.0合入到cann/ops-nn9.1.0
共 26 个文件变更+3378-0
Aactivation/swiglu_group_quant_grad/CMakeLists.txt+22-0
@@ -0,0 +1,22 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------
10+ 
11+message(STATUS "=== Debug: start activation.swiglu_group_quant_grad.CMakeLists.txt ")
12+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
atomgit-bot
atomgit-botatomgit-bot6月18日

🔵 Low Priority

顶层 CMakeLists.txt 第12行使用 file(GLOB CURRENT_DIRS ...) 收集子目录，tests/CMakeLists.txt 第12行同样使用 file(GLOB ...)。CMake 官方文档明确指出 file(GLOB ...) 不保证在新增或删除文件时自动更新，会导致增量构建时遗漏新添加的子目录（如新增的测试目录），需要手动重新运行 cmake 才能生效。在 CI 或开发者增量构建场景下可能造成构建不完整。

建议：考虑显式列出子目录（如 add_subdirectory(op_graph)、add_subdirectory(op_host) 等），而非依赖 file(GLOB)，确保增量构建一致性。如果项目已有明确约定使用 GLOB 模式且可接受此限制，可忽略此建议。

likedislike
不准确？
13+ 
14+if(NOT ENABLE_TEST AND NOT BENCHMARK)
15+ list(REMOVE_ITEM CURRENT_DIRS tests)
16+endif()
17+ 
18+foreach(SUB_DIR ${CURRENT_DIRS})
19+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
20+ add_subdirectory(${SUB_DIR})
21+ endif()
22+endforeach()
Aactivation/swiglu_group_quant_grad/README.md+201-0
@@ -0,0 +1,201 @@
1+# SwigluGroupQuantGrad
2+ 
3+## 产品支持情况
4+ 
5+| 产品 | 是否支持 |
6+| :----------------------------------------------------------- |:-------:|
7+| <term>Ascend 950PR/Ascend 950DT</term> | √ |
8+| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | × |
9+| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | × |
10+| <term>Atlas 200I/500 A2 推理产品</term> | × |
11+| <term>Atlas 推理系列产品</term> | × |
12+| <term>Atlas 训练系列产品</term> | × |
13+ 
14+## 功能说明
15+ 
16+- 算子功能：SwigluGroupQuantGrad算子实现SwiGLU激活函数分组量化的反向梯度计算，用于计算输入梯度`grad_x`和权重梯度`grad_weight`。
17+- 算子支持范围：支持MoE场景（传入groupIndex）和非MoE场景（groupIndex传空），支持可选的Clamp反向传播掩码，支持可选的Weight梯度计算。
18+- 计算流程：
19+ - 步骤〇：GroupIndex处理（可选）→ 计算trunc
20+ - 步骤一：输入切分（将x切分为x0和x1）
21+ - 步骤二：Clamp处理（可选）
22+ - 步骤三：SwiGLU反向传播计算
23+ - 步骤四：Weight梯度计算（可选）
24+ - 步骤五：梯度拼接输出
25+ 
26+- MoE场景GroupIndex处理公式：
27+ $$
28+ \text{trunc} = \sum_{g=0}^{G-1} \text{groupIndex}[g]
29+ $$
30+ 其中：$G$ 为MoE专家分组数，后续所有步骤仅处理前 $\text{trunc}$ 行数据。
31+ 
32+- 输入切分公式：
33+ $$
34+ \mathbf{x}_0[t, h] = \mathbf{x}[t, h], \quad h \in [0, H)
35+ $$
36+
37+ $$
38+ \mathbf{x}_1[t, h] = \mathbf{x}[t, h + H], \quad h \in [0, H)
39+ $$
40+ 
41+- Clamp处理公式（当clamp_limit > 0时）：
42+ $$
43+ \mathbf{x}_0'[t, h] = \min(\mathbf{x}_0[t, h], c)
44+ $$
45+
46+ $$
47+ \mathbf{x}_1'[t, h] = \min(\max(\mathbf{x}_1[t, h], -c), c)
48+ $$
49+ 其中 $c$ 为 `clamp_limit`。
50+ 
51+- SiLU梯度公式：
52+ $$
53+ \frac{d\text{SiLU}}{d\mathbf{x}_0'} = \sigma(\mathbf{x}_0') \cdot \left(1 + \mathbf{x}_0' \cdot (1 - \sigma(\mathbf{x}_0'))\right)
54+ $$
55+ 其中：$\sigma(\mathbf{x}_0') = \frac{1}{1 + e^{-\mathbf{x}_0'}}$
56+ 
57+- 输入梯度计算公式：
58+ $$
59+ \mathbf{grad}_{x_0}[t, h] = \mathbf{grad}_{y_0}[t, h] \cdot \mathbf{x}_1'[t, h] \cdot \frac{d\text{SiLU}}{d\mathbf{x}_0'}[t, h]
60+ $$
61+
62+ $$
63+ \mathbf{grad}_{x_1}[t, h] = \mathbf{grad}_{y_0}[t, h] \cdot \text{SiLU}(\mathbf{x}_0'[t, h])
64+ $$
65+ 其中：如果提供了weight，则 $\mathbf{grad}_{y_0} = \mathbf{grad}_{\text{output}} \cdot \mathbf{weight}$；如果未提供weight，则 $\mathbf{grad}_{y_0} = \mathbf{grad}_{\text{output}}$
66+ 
67+- Weight梯度计算公式（可选）：
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]
70+ $$
71+ 其中：$\mathbf{y}_{\text{origin}}$ 为SwiGLU前向传播的原始激活值输出，沿最后一维（H维度）求和。
72+ 
73+- Clamp反向传播掩码公式（当clamp_limit > 0时）：
74+ $$
75+ \mathbf{grad}_{x_0}[t, h] = \mathbf{grad}_{x_0}[t, h] \cdot \mathbb{I}(\mathbf{x}_0[t, h] < c)
76+ $$
77+
78+ $$
79+ \mathbf{grad}_{x_1}[t, h] = \mathbf{grad}_{x_1}[t, h] \cdot \mathbb{I}(-c < \mathbf{x}_1[t, h] < c)
80+ $$
81+ 其中 $\mathbb{I}$ 为指示函数。
82+ 
83+- 梯度拼接与GroupIndex处理公式：
84+ $$
85+ \mathbf{grad}_x[t, h] = \begin{cases}
86+ \mathbf{grad}_{x_0}[t, h] & h \in [0, H) \\
87+ \mathbf{grad}_{x_1}[t, h-H] & h \in [H, 2H)
88+ \end{cases}
89+ $$
90+
91+ $$
92+ \mathbf{grad}_x[t, :] = \mathbf{grad}_x[t, :] \cdot \mathbb{I}(t < \text{trunc})
93+ $$
94+ 
95+## 参数说明
96+ 
97+<table style="undefined;table-layout: fixed; width: 970px"><colgroup>
98+ <col style="width: 181px">
99+ <col style="width: 144px">
100+ <col style="width: 273px">
101+ <col style="width: 256px">
102+ <col style="width: 116px">
103+ </colgroup>
104+ <thead>
105+ <tr>
106+ <th>参数名</th>
107+ <th>输入/输出/属性</th>
108+ <th>描述</th>
109+ <th>数据类型</th>
110+ <th>数据格式</th>
111+ </tr></thead>
112+ <tbody>
113+ <tr>
114+ <td>gradY</td>
115+ <td>输入</td>
116+ <td>梯度输出张量，来自下游层的梯度。</td>
117+ <td>BFLOAT16、FLOAT16、FLOAT</td>
118+ <td>ND</td>
119+ </tr>
120+ <tr>
121+ <td>x</td>
122+ <td>输入</td>
123+ <td>前向传播的输入张量。</td>
124+ <td>BFLOAT16、FLOAT16、FLOAT</td>
125+ <td>ND</td>
126+ </tr>
127+ <tr>
128+ <td>weightOptional</td>
129+ <td>输入</td>
130+ <td>MoE权重张量。</td>
131+ <td>FLOAT</td>
132+ <td>ND</td>
133+ </tr>
134+ <tr>
135+ <td>yOriginOptional</td>
136+ <td>输入</td>
137+ <td>SwiGLU前向传播的原始激活值输出。</td>
138+ <td>BFLOAT16、FLOAT16、FLOAT</td>
139+ <td>ND</td>
140+ </tr>
141+ <tr>
142+ <td>groupIndexOptional</td>
S
Ssunday6月18日

README参数表中groupIndexOptional标注为输入（非可选），但proto中group_index为OPTIONAL_INPUT。建议标注为输入（可选）。

likedislike
S
shilulu
6月18日 评论：
143+ <td>输入</td>
144+ <td>GroupIndex张量，动态核分配。</td>
145+ <td>INT64</td>
146+ <td>ND</td>
147+ </tr>
148+ <tr>
149+ <td>clampLimit</td>
150+ <td>属性</td>
151+ <td><ul><li>Clamp阈值。</li><li>取值范围≥0.0。</li><li>clampLimit=0表示不启用Clamp反向传播掩码。</li></ul></td>
152+ <td>FLOAT</td>
153+ <td>-</td>
154+ </tr>
155+ <tr>
156+ <td>gradXOut</td>
157+ <td>输出</td>
158+ <td>输入梯度张量。</td>
159+ <td>BFLOAT16、FLOAT16、FLOAT</td>
160+ <td>ND</td>
161+ </tr>
162+ <tr>
163+ <td>gradWeightOutOptional</td>
164+ <td>输出</td>
165+ <td>权重梯度张量。</td>
166+ <td>FLOAT</td>
167+ <td>ND</td>
168+ </tr>
169+ </tbody></table>
170+ 
171+## 约束说明
172+ 
173+- 确定性计算：
174+ - 当提供 `groupIndex` 参数时：前 trunc 行保证计算结果确定性，后 T-trunc 行保证确定性（填充0）
175+ - 当未提供 `groupIndex` 参数时：所有行数据保证计算结果确定性
176+ 
177+- 输入shape约束：
178+ - x最后一维必须为偶数（$2H$）
179+ - gradY最后一维为 $H$，与x最后一维的一半对应
180+ - gradY与x的前n-1维shape必须一致
181+ 
182+- 可选参数约束：
183+ - weight提供时，必须同时提供yOrigin才能计算gradWeight
184+ - weight的shape需与gradY的第一维一致
185+ - yOrigin的shape需与gradY一致
186+ 
187+- 数据类型约束：
188+ - gradY、x、yOrigin、gradXOut数据类型必须一致（FLOAT、FLOAT16或BFLOAT16）
189+ - weight、gradWeightOutOptional必须为FLOAT类型
190+ - groupIndex必须为INT64类型
191+ 
192+- Clamp约束：
193+ - clampLimit必须 ≥ 0.0
194+ - clampLimit=0表示不启用Clamp反向传播掩码
195+ 
196+## 调用说明
197+ 
198+| 调用方式 | 调用样例 | 说明 |
199+|--------------|------------------------------------------------------------------------|----------------------------------------------------------------|
200+| aclnn调用 | [test_aclnn_swiglu_group_quant_grad](./examples/arch35/test_aclnn_swiglu_group_quant_grad.cpp) | 通过[aclnnSwigluGroupQuantGrad](./docs/aclnnSwigluGroupQuantGrad.md)接口方式调用SwigluGroupQuantGrad算子。 |
201+| 图模式调用 | - | 通过[算子IR](./op_graph/swiglu_group_quant_grad_proto.h)构图方式调用SwigluGroupQuantGrad算子。 |
Aactivation/swiglu_group_quant_grad/docs/aclnnSwigluGroupQuantGrad.md+572-0
@@ -0,0 +1,572 @@
1+# aclnnSwigluGroupQuantGrad
2+ 
3+[📄 查看源码](https://gitcode.com/cann/ops-nn/tree/master/activation/swiglu_group_quant_grad)
4+ 
5+## 产品支持情况
6+ 
7+| 产品 | 是否支持 |
8+| :----------------------------------------------------------- |:-------:|
9+| <term>Ascend 950PR/Ascend 950DT</term> | √ |
10+| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | × |
11+| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | × |
12+| <term>Atlas 200I/500 A2 推理产品</term> | × |
13+| <term>Atlas 推理系列产品</term> | × |
14+| <term>Atlas 训练系列产品</term> | × |
15+ 
16+## 功能说明
17+ 
18+- 接口功能：SwigluGroupQuantGrad算子实现SwiGLU激活函数分组量化的反向梯度计算。用于计算输入梯度`grad_x`和权重梯度`grad_weight`。
19+- 算子支持范围：支持MoE场景（传入groupIndex）和非MoE场景（groupIndex传空），支持可选的Clamp反向传播掩码，支持可选的Weight梯度计算。
20+- 计算流程：
21+ - 步骤〇：GroupIndex处理（可选）→ 计算trunc
22+ - 步骤一：输入切分（将x切分为x0和x1）
23+ - 步骤二：Clamp处理（可选）
24+ - 步骤三：SwiGLU反向传播计算
25+ - 步骤四：Weight梯度计算（可选）
26+ - 步骤五：梯度拼接输出
27+ 
28+- MoE场景GroupIndex处理公式：
29+ $$
30+ \text{trunc} = \sum_{g=0}^{G-1} \text{groupIndex}[g]
31+ $$
32+ 其中：$G$ 为MoE专家分组数，后续所有步骤仅处理前 $\text{trunc}$ 行数据。
33+ 
34+- 输入切分公式：
35+ $$
36+ \mathbf{x}_0[t, h] = \mathbf{x}[t, h], \quad h \in [0, H)
37+ $$
38+
39+ $$
40+ \mathbf{x}_1[t, h] = \mathbf{x}[t, h + H], \quad h \in [0, H)
41+ $$
42+ 
43+- Clamp处理公式（当clamp_limit > 0时）：
44+ $$
45+ \mathbf{x}_0'[t, h] = \min(\mathbf{x}_0[t, h], c)
46+ $$
47+
48+ $$
49+ \mathbf{x}_1'[t, h] = \min(\max(\mathbf{x}_1[t, h], -c), c)
50+ $$
51+ 其中 $c$ 为 `clamp_limit`。
52+ 
53+- SiLU梯度公式：
54+ $$
55+ \frac{d\text{SiLU}}{d\mathbf{x}_0'} = \sigma(\mathbf{x}_0') \cdot \left(1 + \mathbf{x}_0' \cdot (1 - \sigma(\mathbf{x}_0'))\right)
56+ $$
57+ 其中：$\sigma(\mathbf{x}_0') = \frac{1}{1 + e^{-\mathbf{x}_0'}}$
58+ 
59+- 输入梯度计算公式：
60+ $$
61+ \mathbf{grad}_{x_0}[t, h] = \mathbf{grad}_{y_0}[t, h] \cdot \mathbf{x}_1'[t, h] \cdot \frac{d\text{SiLU}}{d\mathbf{x}_0'}[t, h]
62+ $$
atomgit-bot
atomgit-botatomgit-bot6月18日

🟠 High Priority

文档中的示例代码（第482、485、491、498行）与示例文件 test_aclnn_swiglu_group_quant_grad.cpp 存在相同的 CreateAclTensor 数据类型/内存不匹配问题：host 数据为 std::vector<float>（T=float，4字节/元素），但 dataType 传入 ACL_FLOAT16（2字节/元素），导致 device 上分配和拷贝的内存大小与 tensor 声明类型不一致。

建议：同示例文件的修复方案：将 hostData 的类型与 dataType 对齐，或在 CreateAclTensor 内部根据 dataType 计算正确 sizeof。

likedislike
不准确？
63+
64+ $$
65+ \mathbf{grad}_{x_1}[t, h] = \mathbf{grad}_{y_0}[t, h] \cdot \text{SiLU}(\mathbf{x}_0'[t, h])
66+ $$
67+ 其中：如果提供了weight，则 $\mathbf{grad}_{y_0} = \mathbf{grad}_{\text{output}} \cdot \mathbf{weight}$；如果未提供weight，则 $\mathbf{grad}_{y_0} = \mathbf{grad}_{\text{output}}$
68+ 
69+- Weight梯度计算公式（可选）：
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]
72+ $$
73+ 其中：$\mathbf{y}_{\text{origin}}$ 为SwiGLU前向传播的原始激活值输出，沿最后一维（H维度）求和。
74+ 
75+- Clamp反向传播掩码公式（当clamp_limit > 0时）：
76+ $$
77+ \mathbf{grad}_{x_0}[t, h] = \mathbf{grad}_{x_0}[t, h] \cdot \mathbb{I}(\mathbf{x}_0[t, h] < c)
78+ $$
79+
80+ $$
81+ \mathbf{grad}_{x_1}[t, h] = \mathbf{grad}_{x_1}[t, h] \cdot \mathbb{I}(-c < \mathbf{x}_1[t, h] < c)
82+ $$
83+ 其中 $\mathbb{I}$ 为指示函数。
84+ 
85+- 梯度拼接与GroupIndex处理公式：
86+ $$
87+ \mathbf{grad}_x[t, h] = \begin{cases}
88+ \mathbf{grad}_{x_0}[t, h] & h \in [0, H) \\
89+ \mathbf{grad}_{x_1}[t, h-H] & h \in [H, 2H)
90+ \end{cases}
91+ $$
92+
93+ $$
94+ \mathbf{grad}_x[t, :] = \mathbf{grad}_x[t, :] \cdot \mathbb{I}(t < \text{trunc})
95+ $$
96+ 
97+## 函数原型
98+ 
99+每个算子分为[两段式接口](../../../docs/zh/context/两段式接口.md)，必须先调用"aclnnSwigluGroupQuantGradGetWorkspaceSize"接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用"aclnnSwigluGroupQuantGrad"接口执行计算。
100+ 
101+```Cpp
102+aclnnStatus aclnnSwigluGroupQuantGradGetWorkspaceSize(
103+ const aclTensor *gradY,
104+ const aclTensor *x,
105+ const aclTensor *weightOptional,
106+ const aclTensor *yOriginOptional,
107+ const aclIntArray *groupIndexOptional,
108+ double clampLimit,
S
Ssunday6月18日

函数原型声明clampLimit为double，参数表标注为float，示例代码使用1.0f，类型不一致。建议统一为一种类型。

likedislike
S
shilulu
6月18日 评论：
109+ const aclTensor *gradXOut,
110+ const aclTensor *gradWeightOutOptional,
111+ uint64_t *workspaceSize,
112+ aclOpExecutor **executor)
113+```
114+ 
115+```Cpp
116+aclnnStatus aclnnSwigluGroupQuantGrad(
117+ void *workspace,
118+ uint64_t workspaceSize,
119+ aclOpExecutor *executor,
120+ aclrtStream stream)
121+```
122+ 
123+## aclnnSwigluGroupQuantGradGetWorkspaceSize
124+ 
125+- **参数说明：**
126+ 
127+ <table style="undefined;table-layout: fixed; width: 1547px"><colgroup>
128+ <col style="width: 200px">
129+ <col style="width: 120px">
130+ <col style="width: 250px">
131+ <col style="width: 330px">
132+ <col style="width: 212px">
133+ <col style="width: 100px">
134+ <col style="width: 190px">
135+ <col style="width: 145px">
136+ </colgroup>
137+ <thead>
138+ <tr>
139+ <th>参数名</th>
140+ <th>输入/输出</th>
141+ <th>描述</th>
142+ <th>使用说明</th>
143+ <th>数据类型</th>
144+ <th>数据格式</th>
145+ <th>维度(shape)</th>
146+ <th>非连续Tensor</th>
147+ </tr></thead>
148+ <tbody>
149+ <tr>
150+ <td>gradY（aclTensor*）</td>
151+ <td>输入</td>
152+ <td>梯度输出张量，来自下游层的梯度。</td>
153+ <td><ul><li>shape=[T, H]或[B, S, H]。</li><li>T为token数量，B为batch size，S为sequence length，H为hidden size。</li></ul></td>
154+ <td>BFLOAT16、FLOAT16、FLOAT</td>
155+ <td>ND</td>
156+ <td>2-3</td>
157+ <td>√</td>
158+ </tr>
159+ <tr>
160+ <td>x（aclTensor*）</td>
161+ <td>输入</td>
162+ <td>前向传播的输入张量。</td>
163+ <td><ul><li>shape=[T, 2H]或[B, S, 2H]，最后一维必须为偶数。</li><li>最后一维的H与gradY的H对应。</li></ul></td>
164+ <td>BFLOAT16、FLOAT16、FLOAT</td>
165+ <td>ND</td>
166+ <td>2-3</td>
167+ <td>√</td>
168+ </tr>
169+ <tr>
170+ <td>weightOptional（aclTensor*）</td>
171+ <td>输入（可选）</td>
172+ <td>MoE权重张量。</td>
173+ <td><ul><li>shape=[T, 1]或[B, S, 1]，需与gradY的第一维一致。</li><li>当提供weight时，必须同时提供yOrigin才能计算gradWeight。</li></ul></td>
174+ <td>FLOAT</td>
175+ <td>ND</td>
176+ <td>2-3</td>
177+ <td>√</td>
178+ </tr>
179+ <tr>
180+ <td>yOriginOptional（aclTensor*）</td>
181+ <td>输入（可选）</td>
182+ <td>SwiGLU前向传播的原始激活值输出。</td>
183+ <td><ul><li>shape=[T, H]或[B, S, H]，需与gradY的shape一致。</li><li>当提供weight时，必须同时提供yOrigin才能计算gradWeight。</li></ul></td>
184+ <td>BFLOAT16、FLOAT16、FLOAT</td>
185+ <td>ND</td>
186+ <td>2-3</td>
187+ <td>√</td>
188+ </tr>
189+ <tr>
190+ <td>groupIndexOptional（aclTensor*）</td>
S
Ssunday6月18日

参数表中groupIndexOptional类型标注为aclTensor*，但函数原型声明为aclIntArray*，类型矛盾。aclIntArray不支持非连续Tensor标注。建议参数表改为aclIntArray*，并移除非连续列的勾选。

likedislike
S
shilulu
6月18日 评论：
191+ <td>输入（可选）</td>
192+ <td>GroupIndex张量，动态核分配。</td>
193+ <td><ul><li>shape=[G]，dtype=INT64。</li><li>G为MoE专家分组数。</li><li>groupIndex内元素要求为非递减。</li></ul></td>
194+ <td>INT64</td>
195+ <td>ND</td>
196+ <td>1</td>
197+ <td>√</td>
198+ </tr>
199+ <tr>
200+ <td>clampLimit（float）</td>
201+ <td>输入</td>
202+ <td>Clamp阈值。</td>
203+ <td><ul><li>取值范围≥0.0。</li><li>clampLimit=0表示不启用Clamp反向传播掩码。</li></ul></td>
204+ <td>FLOAT</td>
205+ <td>-</td>
206+ <td>-</td>
207+ <td>-</td>
208+ </tr>
209+ <tr>
210+ <td>gradXOut（aclTensor*）</td>
211+ <td>输出</td>
212+ <td>输入梯度张量。</td>
213+ <td><ul><li>shape=[T, 2H]或[B, S, 2H]，与x一致。</li><li>数据类型与gradY/x保持一致。</li></ul></td>
214+ <td>BFLOAT16、FLOAT16、FLOAT</td>
215+ <td>ND</td>
216+ <td>2-3</td>
217+ <td>√</td>
218+ </tr>
219+ <tr>
220+ <td>gradWeightOutOptional（aclTensor*）</td>
221+ <td>输出（可选）</td>
222+ <td>权重梯度张量。</td>
223+ <td><ul><li>当提供weight时输出，shape=[T, 1]或[B, S, 1]。</li><li>数据类型为FLOAT。</li></ul></td>
224+ <td>FLOAT</td>
225+ <td>ND</td>
226+ <td>2-3</td>
227+ <td>√</td>
228+ </tr>
229+ <tr>
230+ <td>workspaceSize（uint64_t*）</td>
231+ <td>输出</td>
232+ <td>返回需要在Device侧申请的workspace大小。</td>
233+ <td>-</td>
234+ <td>-</td>
235+ <td>-</td>
236+ <td>-</td>
237+ <td>-</td>
238+ </tr>
239+ <tr>
240+ <td>executor（aclOpExecutor**）</td>
241+ <td>输出</td>
242+ <td>返回op执行器，包含了算子计算流程。</td>
243+ <td>-</td>
244+ <td>-</td>
245+ <td>-</td>
246+ <td>-</td>
247+ <td>-</td>
248+ </tr>
249+ </tbody>
250+ </table>
251+ 
252+- **返回值：**
253+ 
254+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
255+ 
256+ 第一段接口会完成入参校验，出现以下场景时报错：
257+ 
258+ <table style="undefined;table-layout: fixed; width: 1155px"><colgroup>
259+ <col style="width: 253px">
260+ <col style="width: 140px">
261+ <col style="width: 762px">
262+ </colgroup>
263+ <thead>
264+ <tr>
265+ <th>返回码</th>
266+ <th>错误码</th>
267+ <th>描述</th>
268+ </tr></thead>
269+ <tbody>
270+ <tr>
271+ <td>ACLNN_ERR_PARAM_NULLPTR</td>
272+ <td>161001</td>
273+ <td>必选参数gradY/x/gradXOut为nullptr。</td>
274+ </tr>
275+ <tr>
276+ <td>ACLNN_ERR_INNER_TILING_ERROR</td>
277+ <td>161002</td>
278+ <td>gradY、x、weight等输入变量的数据类型和数据格式不在支持的范围内。</td>
279+ </tr>
280+ <tr>
281+ <td>ACLNN_ERR_INNER_TILING_ERROR</td>
282+ <td>561002</td>
283+ <td>多个输入tensor之间的shape信息不匹配、输入属性不在取值范围（详见参数说明）。</td>
284+ </tr>
285+ </tbody></table>
286+ 
287+## aclnnSwigluGroupQuantGrad
288+ 
289+- **参数说明：**
290+ 
291+ <table style="undefined;table-layout: fixed; width: 1149px"><colgroup>
292+ <col style="width: 173px">
293+ <col style="width: 124px">
294+ <col style="width: 852px">
295+ </colgroup>
296+ <thead>
297+ <tr>
298+ <th>参数名</th>
299+ <th>输入/输出</th>
300+ <th>描述</th>
301+ </tr></thead>
302+ <tbody>
303+ <tr>
304+ <td>workspace</td>
305+ <td>输入</td>
306+ <td>在Device侧申请的workspace内存地址。</td>
307+ </tr>
308+ <tr>
309+ <td>workspaceSize</td>
310+ <td>输入</td>
311+ <td>在Device侧申请的workspace大小，由第一段接口aclnnSwigluGroupQuantGradGetWorkspaceSize获取。</td>
312+ </tr>
313+ <tr>
314+ <td>executor</td>
315+ <td>输入</td>
316+ <td>op执行器，包含了算子计算流程。</td>
317+ </tr>
318+ <tr>
319+ <td>stream</td>
320+ <td>输入</td>
321+ <td>指定执行任务的Stream。</td>
322+ </tr>
323+ </tbody>
324+ </table>
325+ 
326+- **返回值：**
327+ 
328+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn返回码.md)。
329+ 
330+## 约束说明
331+ 
332+- 确定性计算：
333+ - aclnnSwigluGroupQuantGrad默认确定性实现。
334+ 
335+- 输入shape约束：
336+ - x最后一维必须为偶数（$2H$）
337+ - gradY最后一维为 $H$，与x最后一维的一半对应
338+ - gradY与x的前n-1维shape必须一致
339+ 
340+- 可选参数约束：
341+ - weight提供时，必须同时提供yOrigin才能计算gradWeight
342+ - weight的shape需与gradY的第一维一致
343+ - yOrigin的shape需与gradY一致
344+ 
345+- 数据类型约束：
346+ - gradY、x、yOrigin、gradX数据类型必须一致（FLOAT、FLOAT16或BFLOAT16）
347+ - weight、gradWeight必须为FLOAT类型
348+ - groupIndex必须为INT64类型
349+ 
350+- Clamp约束：
351+ - clampLimit必须 ≥ 0.0
352+ - clampLimit=0表示不启用Clamp反向传播掩码
353+ 
354+## 调用示例
355+ 
356+示例代码如下，仅供参考，具体编译和执行过程请参考[编译与运行样例](../../../docs/zh/context/编译与运行样例.md)。
357+ 
358+```Cpp
359+#include <iostream>
360+#include <vector>
361+#include "acl/acl.h"
362+#include "aclnnop/aclnn_swiglu_group_quant_grad.h"
363+ 
364+#define CHECK_RET(cond, return_expr) \
365+ do { \
366+ if (!(cond)) { \
367+ return_expr; \
368+ } \
369+ } while (0)
370+ 
371+#define LOG_PRINT(message, ...) \
372+ do { \
373+ printf(message, ##__VA_ARGS__); \
374+ } while (0)
375+ 
376+int64_t GetShapeSize(const std::vector<int64_t>& shape) {
377+ int64_t shapeSize = 1;
378+ for (auto i : shape) {
379+ shapeSize *= i;
380+ }
381+ return shapeSize;
382+}
383+ 
384+int Init(int32_t deviceId, aclrtStream* stream) {
385+ auto ret = aclInit(nullptr);
386+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
387+ ret = aclrtSetDevice(deviceId);
388+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
389+ ret = aclrtCreateStream(stream);
390+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
391+ return 0;
392+}
393+ 
394+template <typename T>
395+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
396+ aclDataType dataType, aclTensor** tensor) {
397+ auto size = GetShapeSize(shape) * sizeof(T);
398+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
399+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
400+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
401+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
402+ 
403+ std::vector<int64_t> strides(shape.size(), 1);
404+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
405+ strides[i] = shape[i + 1] * strides[i + 1];
406+ }
407+ 
408+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
409+ shape.data(), shape.size(), *deviceAddr);
410+ return 0;
411+}
412+ 
413+template <typename T>
414+int CreateAclTensorWithValue(const std::vector<int64_t>& shape, void** deviceAddr,
415+ aclDataType dataType, aclTensor** tensor, T value) {
416+ int64_t shapeSize = GetShapeSize(shape);
417+ std::vector<T> hostData(shapeSize, value);
418+ return CreateAclTensor(hostData, shape, deviceAddr, dataType, tensor);
419+}
420+ 
421+int main() {
422+ int32_t deviceId = 0;
423+ aclrtStream stream;
424+ auto ret = Init(deviceId, &stream);
425+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
426+ 
427+ std::vector<int64_t> gradYShape = {512, 512};
428+ std::vector<int64_t> xShape = {512, 1024};
429+ std::vector<int64_t> weightShape = {512, 1};
430+ std::vector<int64_t> yOriginShape = {512, 512};
431+ std::vector<int64_t> groupIndexShape = {256};
432+ std::vector<int64_t> gradXShape = {512, 1024};
433+ std::vector<int64_t> gradWeightShape = {512, 1};
434+ 
435+ void* gradYDeviceAddr = nullptr;
436+ void* xDeviceAddr = nullptr;
437+ void* weightDeviceAddr = nullptr;
438+ void* yOriginDeviceAddr = nullptr;
439+ void* groupIndexDeviceAddr = nullptr;
440+ void* gradXDeviceAddr = nullptr;
441+ void* gradWeightDeviceAddr = nullptr;
442+ 
443+ aclTensor* gradYTensor = nullptr;
444+ aclTensor* xTensor = nullptr;
445+ aclTensor* weightTensor = nullptr;
446+ aclTensor* yOriginTensor = nullptr;
447+ aclIntArray* groupIndexArray = nullptr;
448+ aclTensor* gradXTensor = nullptr;
449+ aclTensor* gradWeightTensor = nullptr;
450+ 
451+ int64_t gradYSize = GetShapeSize(gradYShape);
452+ std::vector<float> gradYHostData(gradYSize, 1.0f);
453+ for (int64_t i = 0; i < gradYSize; i++) {
454+ gradYHostData[i] = static_cast<float>(i % 10) * 0.1f;
455+ }
456+ 
457+ int64_t xSize = GetShapeSize(xShape);
458+ std::vector<float> xHostData(xSize, 1.0f);
459+ for (int64_t i = 0; i < xSize; i++) {
460+ xHostData[i] = static_cast<float>((i % 20) - 10) * 0.5f;
461+ }
462+ 
463+ int64_t weightSize = GetShapeSize(weightShape);
464+ std::vector<float> weightHostData(weightSize, 1.0f);
465+ for (int64_t i = 0; i < weightSize; i++) {
466+ weightHostData[i] = static_cast<float>((i % 5) + 1) * 0.2f;
467+ }
468+ 
469+ int64_t yOriginSize = GetShapeSize(yOriginShape);
470+ std::vector<float> yOriginHostData(yOriginSize, 1.0f);
471+ for (int64_t i = 0; i < yOriginSize; i++) {
472+ yOriginHostData[i] = static_cast<float>((i % 8) + 1) * 0.3f;
473+ }
474+ 
475+ int64_t groupIndexSize = GetShapeSize(groupIndexShape);
476+ std::vector<int64_t> groupIndexHostData(groupIndexSize, 0);
477+ int64_t groupStride = 512 / 256;
478+ for (int64_t i = 0; i < groupIndexSize; i++) {
479+ groupIndexHostData[i] = i * groupStride;
480+ }
481+ 
482+ ret = CreateAclTensor(gradYHostData, gradYShape, &gradYDeviceAddr, aclDataType::ACL_FLOAT16, &gradYTensor);
483+ CHECK_RET(ret == ACL_SUCCESS, return ret);
484+ 
485+ ret = CreateAclTensor(xHostData, xShape, &xDeviceAddr, aclDataType::ACL_FLOAT16, &xTensor);
486+ CHECK_RET(ret == ACL_SUCCESS, return ret);
487+ 
488+ ret = CreateAclTensor(weightHostData, weightShape, &weightDeviceAddr, aclDataType::ACL_FLOAT, &weightTensor);
489+ CHECK_RET(ret == ACL_SUCCESS, return ret);
490+ 
491+ ret = CreateAclTensor(yOriginHostData, yOriginShape, &yOriginDeviceAddr, aclDataType::ACL_FLOAT16, &yOriginTensor);
492+ CHECK_RET(ret == ACL_SUCCESS, return ret);
493+ 
494+ std::vector<int64_t> groupArray = {256, 256};
495+ groupIndexArray = aclCreateIntArray(groupArray.data(), groupArray.size());
496+ CHECK_RET(ret == ACL_SUCCESS, return ret);
atomgit-bot
atomgit-botatomgit-bot6月18日

🟠 High Priority

文档示例代码第494-496行：aclCreateIntArray 之后 CHECK_RET(ret == ACL_SUCCESS, return ret) 检查的是第491行上一个 CreateAclTensor 调用的残留 ret，而非 aclCreateIntArray 的返回值。失败时无法检测。

建议：将 CHECK_RET 改为检查 groupIndexArray 指针非空。

改动建议
496
- CHECK_RET(ret == ACL_SUCCESS, return ret);
496
+ std::vector<int64_t> groupArray = {256, 256};
497
+ groupIndexArray = aclCreateIntArray(groupArray.data(), groupArray.size());
498
+ CHECK_RET(groupIndexArray != nullptr, LOG_PRINT("aclCreateIntArray failed.\n"); return -1);
应用建议
likedislike
不准确？
497+ 
498+ ret = CreateAclTensorWithValue<float>(gradXShape, &gradXDeviceAddr, aclDataType::ACL_FLOAT16, &gradXTensor, 0.0f);
499+ CHECK_RET(ret == ACL_SUCCESS, return ret);
500+ 
501+ ret = CreateAclTensorWithValue<float>(gradWeightShape, &gradWeightDeviceAddr, aclDataType::ACL_FLOAT, &gradWeightTensor, 0.0f);
502+ CHECK_RET(ret == ACL_SUCCESS, return ret);
503+ 
504+ float clampLimit = 1.0f;
505+ 
506+ uint64_t workspaceSize = 0;
507+ aclOpExecutor* executor;
508+ 
509+ ret = aclnnSwigluGroupQuantGradGetWorkspaceSize(gradYTensor, xTensor, weightTensor, yOriginTensor,
510+ groupIndexArray, clampLimit, gradXTensor, gradWeightTensor,
511+ &workspaceSize, &executor);
512+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSwigluGroupQuantGradGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
513+ 
514+ void* workspaceAddr = nullptr;
515+ if (workspaceSize > 0) {
516+ 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);
518+ }
519+ 
520+ ret = aclnnSwigluGroupQuantGrad(workspaceAddr, workspaceSize, executor, stream);
521+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSwigluGroupQuantGrad failed. ERROR: %d\n", ret); return ret);
522+ 
523+ ret = aclrtSynchronizeStream(stream);
524+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
525+ 
526+ auto gradXResultSize = GetShapeSize(gradXShape);
527+ std::vector<float> gradXResultData(gradXResultSize, 0);
528+ ret = aclrtMemcpy(gradXResultData.data(), gradXResultData.size() * sizeof(float),
529+ gradXDeviceAddr, 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);
531+ 
532+ LOG_PRINT("gradX output (first 10 elements):\n");
533+ for (int64_t i = 0; i < 10 && i < gradXResultSize; i++) {
534+ LOG_PRINT("gradX[%ld] = %f\n", i, gradXResultData[i]);
535+ }
536+ 
537+ auto gradWeightResultSize = GetShapeSize(gradWeightShape);
538+ std::vector<float> gradWeightResultData(gradWeightResultSize, 0);
539+ ret = aclrtMemcpy(gradWeightResultData.data(), gradWeightResultData.size() * sizeof(float),
540+ gradWeightDeviceAddr, 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);
542+ 
543+ LOG_PRINT("gradWeight output (first 10 elements):\n");
544+ for (int64_t i = 0; i < 10 && i < gradWeightResultSize; i++) {
545+ LOG_PRINT("gradWeight[%ld] = %f\n", i, gradWeightResultData[i]);
546+ }
547+ 
548+ aclDestroyTensor(gradYTensor);
549+ aclDestroyTensor(xTensor);
550+ aclDestroyTensor(weightTensor);
551+ aclDestroyTensor(yOriginTensor);
552+ aclDestroyTensor(gradXTensor);
553+ aclDestroyTensor(gradWeightTensor);
atomgit-bot
atomgit-botatomgit-bot6月18日

🟡 Medium Priority

文档示例代码中 groupIndexArray（通过 aclCreateIntArray 创建）在函数退出前从未调用 aclDestroyIntArray 销毁，造成资源泄漏。

建议：在清理代码段添加 if (groupIndexArray != nullptr) { aclDestroyIntArray(groupIndexArray); }

likedislike
不准确？
554+ 
555+ aclrtFree(gradYDeviceAddr);
556+ aclrtFree(xDeviceAddr);
557+ aclrtFree(weightDeviceAddr);
558+ aclrtFree(yOriginDeviceAddr);
559+ aclrtFree(groupIndexDeviceAddr);
560+ aclrtFree(gradXDeviceAddr);
561+ aclrtFree(gradWeightDeviceAddr);
562+ if (workspaceSize > 0) {
563+ aclrtFree(workspaceAddr);
564+ }
565+ 
566+ aclrtDestroyStream(stream);
567+ aclrtResetDevice(deviceId);
568+ aclFinalize();
569+ 
570+ return 0;
571+}
572+```
Aactivation/swiglu_group_quant_grad/examples/arch35/test_aclnn_swiglu_group_quant_grad.cpp+223-0
@@ -0,0 +1,223 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include <iostream>
12+#include <vector>
13+#include "acl/acl.h"
14+#include "aclnnop/aclnn_swiglu_group_quant_grad.h"
15+ 
16+#define CHECK_RET(cond, return_expr) \
17+ do { \
18+ if (!(cond)) { \
19+ return_expr; \
20+ } \
21+ } while (0)
22+ 
23+#define LOG_PRINT(message, ...) \
24+ do { \
25+ printf(message, ##__VA_ARGS__); \
26+ } while (0)
27+ 
28+int64_t GetShapeSize(const std::vector<int64_t>& shape) {
29+ int64_t shapeSize = 1;
30+ for (auto i : shape) {
31+ shapeSize *= i;
32+ }
33+ return shapeSize;
34+}
35+ 
36+int Init(int32_t deviceId, aclrtStream* stream) {
37+ auto ret = aclInit(nullptr);
38+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
39+ ret = aclrtSetDevice(deviceId);
40+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
41+ ret = aclrtCreateStream(stream);
42+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
43+ return 0;
44+}
45+ 
46+template <typename T>
47+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
48+ aclDataType dataType, aclTensor** tensor) {
49+ auto size = GetShapeSize(shape) * sizeof(T);
50+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
51+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
52+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
53+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
54+ 
55+ std::vector<int64_t> strides(shape.size(), 1);
56+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
57+ strides[i] = shape[i + 1] * strides[i + 1];
58+ }
59+ 
60+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
61+ shape.data(), shape.size(), *deviceAddr);
62+ return 0;
atomgit-bot
atomgit-botatomgit-bot6月18日

🟠 High Priority

在 CreateAclTensor 模板函数中，size 由 GetShapeSize(shape) * sizeof(T) 计算，模板参数 T 被推导为 float（因为 hostData 是 std::vector<float>），因此 device 上分配了 float（4字节/元素）大小的内存并拷贝了 float 数据。但 dataType 参数传入的是 ACL_FLOAT16（2字节/元素），导致 tensor 元数据声明的数据类型与实际 device 数据不匹配。kernel 按 FLOAT16 读取时将得到错误数据。

此问题影响示例文件 test_aclnn_swiglu_group_quant_grad.cpp 中 gradY、x、yOrigin、gradX 等 tensor 的创建（第134、137、143、150行），以及文档 aclnnSwigluGroupQuantGrad.md 中对应的相同代码（第482、485、491、498行）。

likedislike
不准确？
63+}
64+ 
65+template <typename T>
66+int CreateAclTensorWithValue(const std::vector<int64_t>& shape, void** deviceAddr,
67+ aclDataType dataType, aclTensor** tensor, T value) {
68+ int64_t shapeSize = GetShapeSize(shape);
69+ std::vector<T> hostData(shapeSize, value);
70+ return CreateAclTensor(hostData, shape, deviceAddr, dataType, tensor);
71+}
72+ 
73+int main() {
74+ int32_t deviceId = 0;
75+ aclrtStream stream;
76+ auto ret = Init(deviceId, &stream);
77+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
78+ 
79+ std::vector<int64_t> gradYShape = {512, 512};
80+ std::vector<int64_t> xShape = {512, 1024};
81+ std::vector<int64_t> weightShape = {512, 1};
82+ std::vector<int64_t> yOriginShape = {512, 512};
83+ std::vector<int64_t> groupIndexShape = {256};
84+ std::vector<int64_t> gradXShape = {512, 1024};
85+ std::vector<int64_t> gradWeightShape = {512, 1};
86+ 
87+ void* gradYDeviceAddr = nullptr;
88+ void* xDeviceAddr = nullptr;
89+ void* weightDeviceAddr = nullptr;
90+ void* yOriginDeviceAddr = nullptr;
91+ void* groupIndexDeviceAddr = nullptr;
S
Ssunday6月18日

groupIndex作为aclIntArray传入API不需要device内存，但代码声明了groupIndexDeviceAddr并在line 211调用aclrtFree释放，实际从未aclrtMalloc分配，释放nullptr是无效操作。建议移除该变量和对应的aclrtFree。另外释放资源时缺少aclDestroyIntArray(groupIndexArray)，导致aclIntArray内存泄漏。

likedislike
S
shilulu
6月18日 评论：
92+ void* gradXDeviceAddr = nullptr;
93+ void* gradWeightDeviceAddr = nullptr;
94+ 
95+ aclTensor* gradYTensor = nullptr;
96+ aclTensor* xTensor = nullptr;
97+ aclTensor* weightTensor = nullptr;
98+ aclTensor* yOriginTensor = nullptr;
99+ aclIntArray* groupIndexArray = nullptr;
100+ aclTensor* gradXTensor = nullptr;
101+ aclTensor* gradWeightTensor = nullptr;
102+ 
103+ int64_t gradYSize = GetShapeSize(gradYShape);
104+ std::vector<float> gradYHostData(gradYSize, 1.0f);
105+ for (int64_t i = 0; i < gradYSize; i++) {
106+ gradYHostData[i] = static_cast<float>(i % 10) * 0.1f;
107+ }
108+ 
109+ int64_t xSize = GetShapeSize(xShape);
110+ std::vector<float> xHostData(xSize, 1.0f);
111+ for (int64_t i = 0; i < xSize; i++) {
112+ xHostData[i] = static_cast<float>((i % 20) - 10) * 0.5f;
113+ }
114+ 
115+ int64_t weightSize = GetShapeSize(weightShape);
116+ std::vector<float> weightHostData(weightSize, 1.0f);
117+ for (int64_t i = 0; i < weightSize; i++) {
118+ weightHostData[i] = static_cast<float>((i % 5) + 1) * 0.2f;
119+ }
120+ 
121+ int64_t yOriginSize = GetShapeSize(yOriginShape);
122+ std::vector<float> yOriginHostData(yOriginSize, 1.0f);
123+ for (int64_t i = 0; i < yOriginSize; i++) {
124+ yOriginHostData[i] = static_cast<float>((i % 8) + 1) * 0.3f;
125+ }
126+ 
127+ int64_t groupIndexSize = GetShapeSize(groupIndexShape);
128+ std::vector<int64_t> groupIndexHostData(groupIndexSize, 0);
129+ int64_t groupStride = 512 / 256;
130+ for (int64_t i = 0; i < groupIndexSize; i++) {
131+ groupIndexHostData[i] = i * groupStride;
132+ }
133+ 
134+ ret = CreateAclTensor(gradYHostData, gradYShape, &gradYDeviceAddr, aclDataType::ACL_FLOAT16, &gradYTensor);
135+ CHECK_RET(ret == ACL_SUCCESS, return ret);
136+ 
137+ ret = CreateAclTensor(xHostData, xShape, &xDeviceAddr, aclDataType::ACL_FLOAT16, &xTensor);
138+ CHECK_RET(ret == ACL_SUCCESS, return ret);
139+ 
140+ ret = CreateAclTensor(weightHostData, weightShape, &weightDeviceAddr, aclDataType::ACL_FLOAT, &weightTensor);
141+ CHECK_RET(ret == ACL_SUCCESS, return ret);
142+ 
143+ ret = CreateAclTensor(yOriginHostData, yOriginShape, &yOriginDeviceAddr, aclDataType::ACL_FLOAT16, &yOriginTensor);
144+ CHECK_RET(ret == ACL_SUCCESS, return ret);
145+ 
146+ std::vector<int64_t> groupArray = {256, 256};
S
Ssunday6月18日

使用{256,256}数据创建groupIndex数组，但前面已正确计算了groupIndexHostData（每个元素=i*groupStride），此处应使用groupIndexHostData而非硬编码的{256,256}。另外line 148的CHECK_RET(ret==ACL_SUCCESS)检查的是上一个CreateAclTensor的返回值，aclCreateIntArray的返回值未被检查。

likedislike
S
shilulu
6月18日 评论：
147+ groupIndexArray = aclCreateIntArray(groupArray.data(), groupArray.size());
148+ CHECK_RET(ret == ACL_SUCCESS, return ret);
atomgit-bot
atomgit-botatomgit-bot6月18日

🟠 High Priority

第146-148行：groupIndexArray = aclCreateIntArray(groupArray.data(), groupArray.size()); 之后紧跟 CHECK_RET(ret == ACL_SUCCESS, return ret);。这里检查的 ret 是第143行 CreateAclTensor(... yOrigin ...) 的返回值，而非 aclCreateIntArray 的返回值。如果 aclCreateIntArray 因任何原因返回 nullptr，该错误不会被捕获，后续将 nullptr 的 groupIndexArray 传入 aclnnSwigluGroupQuantGradGetWorkspaceSize 会导致空指针解引用。

文档 aclnnSwigluGroupQuantGrad.md 第494-496行存在相同问题。

建议：将 CHECK_RET 改为检查 aclCreateIntArray 的返回值：CHECK_RET(groupIndexArray != nullptr, LOG_PRINT("aclCreateIntArray failed.\n"); return -1);

改动建议
148
- CHECK_RET(ret == ACL_SUCCESS, return ret);
148
+ std::vector<int64_t> groupArray = {256, 256};
149
+ groupIndexArray = aclCreateIntArray(groupArray.data(), groupArray.size());
150
+ CHECK_RET(groupIndexArray != nullptr, LOG_PRINT("aclCreateIntArray failed.\n"); return -1);
应用建议
likedislike
不准确？
149+ 
150+ ret = CreateAclTensorWithValue<float>(gradXShape, &gradXDeviceAddr, aclDataType::ACL_FLOAT16, &gradXTensor, 0.0f);
151+ CHECK_RET(ret == ACL_SUCCESS, return ret);
152+ 
153+ ret = CreateAclTensorWithValue<float>(gradWeightShape, &gradWeightDeviceAddr, aclDataType::ACL_FLOAT, &gradWeightTensor, 0.0f);
154+ CHECK_RET(ret == ACL_SUCCESS, return ret);
155+ 
156+ float clampLimit = 1.0f;
157+ 
158+ uint64_t workspaceSize = 0;
159+ aclOpExecutor* executor;
160+ 
161+ ret = aclnnSwigluGroupQuantGradGetWorkspaceSize(gradYTensor, xTensor, weightTensor, yOriginTensor,
162+ groupIndexArray, clampLimit, gradXTensor, gradWeightTensor,
163+ &workspaceSize, &executor);
164+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSwigluGroupQuantGradGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
165+ 
166+ void* workspaceAddr = nullptr;
167+ if (workspaceSize > 0) {
168+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
169+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
170+ }
171+ 
172+ ret = aclnnSwigluGroupQuantGrad(workspaceAddr, workspaceSize, executor, stream);
173+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSwigluGroupQuantGrad failed. ERROR: %d\n", ret); return ret);
174+ 
175+ ret = aclrtSynchronizeStream(stream);
176+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
177+ 
178+ auto gradXResultSize = GetShapeSize(gradXShape);
179+ std::vector<float> gradXResultData(gradXResultSize, 0);
180+ ret = aclrtMemcpy(gradXResultData.data(), gradXResultData.size() * sizeof(float),
181+ gradXDeviceAddr, gradXResultSize * sizeof(float), ACL_MEMCPY_DEVICE_TO_HOST);
182+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy gradX result from device to host failed. ERROR: %d\n", ret); return ret);
183+ 
184+ LOG_PRINT("gradX output (first 10 elements):\n");
185+ for (int64_t i = 0; i < 10 && i < gradXResultSize; i++) {
186+ LOG_PRINT("gradX[%ld] = %f\n", i, gradXResultData[i]);
187+ }
188+ 
189+ auto gradWeightResultSize = GetShapeSize(gradWeightShape);
190+ std::vector<float> gradWeightResultData(gradWeightResultSize, 0);
191+ ret = aclrtMemcpy(gradWeightResultData.data(), gradWeightResultData.size() * sizeof(float),
192+ gradWeightDeviceAddr, gradWeightResultSize * sizeof(float), ACL_MEMCPY_DEVICE_TO_HOST);
193+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy gradWeight result from device to host failed. ERROR: %d\n", ret); return ret);
194+ 
195+ LOG_PRINT("gradWeight output (first 10 elements):\n");
196+ for (int64_t i = 0; i < 10 && i < gradWeightResultSize; i++) {
197+ LOG_PRINT("gradWeight[%ld] = %f\n", i, gradWeightResultData[i]);
198+ }
199+ 
200+ aclDestroyTensor(gradYTensor);
201+ aclDestroyTensor(xTensor);
202+ aclDestroyTensor(weightTensor);
203+ aclDestroyTensor(yOriginTensor);
204+ aclDestroyTensor(gradXTensor);
205+ aclDestroyTensor(gradWeightTensor);
206+ 
atomgit-bot
atomgit-botatomgit-bot6月18日

🟡 Medium Priority

第147行通过 aclCreateIntArray 创建的 groupIndexArray 在函数退出前从未调用 aclDestroyIntArray 销毁。对比其他资源（aclTensor 有 aclDestroyTensor，deviceAddr 有 aclrtFree），groupIndexArray 缺少对应的释放调用。这是一个资源泄漏。

文档 aclnnSwigluGroupQuantGrad.md 存在相同问题。

建议：在清理代码段（在 aclDestroyTensor 调用附近）添加 if (groupIndexArray != nullptr) { aclDestroyIntArray(groupIndexArray); }

likedislike
不准确？
207+ aclrtFree(gradYDeviceAddr);
208+ aclrtFree(xDeviceAddr);
209+ aclrtFree(weightDeviceAddr);
210+ aclrtFree(yOriginDeviceAddr);
211+ aclrtFree(groupIndexDeviceAddr);
212+ aclrtFree(gradXDeviceAddr);
213+ aclrtFree(gradWeightDeviceAddr);
214+ if (workspaceSize > 0) {
215+ aclrtFree(workspaceAddr);
216+ }
217+ 
218+ aclrtDestroyStream(stream);
219+ aclrtResetDevice(deviceId);
220+ aclFinalize();
221+ 
222+ return 0;
223+}
Aactivation/swiglu_group_quant_grad/op_graph/CMakeLists.txt+12-0
@@ -0,0 +1,12 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------
10+ 
11+message(STATUS "=== Debug: start ops.activation.swiglu_group_quant_grad.graph_plugin.CMakeLists.txt ")
12+add_graph_plugin_sources()
Aactivation/swiglu_group_quant_grad/op_graph/swiglu_group_quant_grad_proto.h+51-0
@@ -0,0 +1,51 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file swiglu_group_quant_grad_proto.h
13+ * \brief SwiGLU Group Dynamic Quant Backward operator prototype
14+ */
15+#ifndef OPS_QUANT_SWIGLU_GROUP_QUANT_GRAD_PROTO_H_
16+#define OPS_QUANT_SWIGLU_GROUP_QUANT_GRAD_PROTO_H_
17+ 
18+#include "graph/operator_reg.h"
19+ 
20+namespace ge {
21+/**
22+* @brief SwiGLU Group Dynamic Quant Backward operator.
23+ 
24+* @par Inputs:
25+* @li grad_y: Gradient input tensor. Must be one of the following types: float32,float16,bfloat16, has format ND.
26+* @li x: Forward pass input tensor. Must be one of the following types: float32,float16,bfloat16, has format ND.
27+* @li weight: Optional tensor. topk weight tensor. Type is float32, has format ND.
28+* @li y_origin: Optional tensor. Forward pass output before quantization.
29+ Must be one of the following types: float32,float16,bfloat16, has format ND.
30+* @li group_index: Optional tensor. Group index tensor for dynamic quantization. Type is int64, has format ND.
31+ 
32+* @par Attributes:
33+* @li clamp_limit: Optional float. Clamp value for gradient mask, default is 0.0.
34+ 
35+* @par Outputs:
36+* @li grad_x: Gradient of x tensor. Same type as input x, has format ND.
37+* @li grad_weight: Optional output. Gradient of weight tensor. Type is float32, has format ND.
38+*/
39+REG_OP(SwigluGroupQuantGrad)
40+ .INPUT(grad_y, TensorType({DT_FLOAT16, DT_BF16, DT_FLOAT}))
41+ .INPUT(x, TensorType({DT_FLOAT16, DT_BF16, DT_FLOAT}))
42+ .OPTIONAL_INPUT(weight, TensorType({DT_FLOAT}))
43+ .OPTIONAL_INPUT(y_origin, TensorType({DT_FLOAT16, DT_BF16, DT_FLOAT}))
44+ .OPTIONAL_INPUT(group_index, TensorType({DT_INT64}))
45+ .OUTPUT(grad_x, TensorType({DT_FLOAT16, DT_BF16, DT_FLOAT}))
46+ .OUTPUT(grad_weight, TensorType({DT_FLOAT}))
Satomgit-bot
Ssunday6月18日

proto中grad_weight声明为 .OUTPUT() 而非 .OPTIONAL_OUTPUT()，但def.cpp中grad_weight的ParamType是OPTIONAL，两者语义矛盾。框架需通过OPTIONAL_OUTPUT()才能正确区分可选输出。建议proto改为 .OPTIONAL_OUTPUT(grad_weight, TensorType({DT_FLOAT}))。

likedislike
S
shilulu
6月18日 评论：
atomgit-botatomgit-bot6月18日

🟡 Medium Priority

swiglu_group_quant_grad_proto.h 第46行将 grad_weight 声明为 .OUTPUT(grad_weight, ...)（必选输出），但 swiglu_group_quant_grad_def.cpp 第53-57行将其声明为 ParamType(OPTIONAL)（可选输出），且算子规格说明中 gradWeightOutOptional 是可选输出。图模式下如果 grad_weight 不提供输出 tensor，框架可能因为 OUTPUT 声明而报错或行为不符合预期。

建议：将 .OUTPUT(grad_weight, ...) 改为 .OPTIONAL_OUTPUT(grad_weight, ...) 或 .DYNAMIC_OUTPUT(grad_weight, ...) 以与 def.cpp 及算子系统规格保持一致。

改动建议
46
- .OUTPUT(grad_weight, TensorType({DT_FLOAT}))
46
+ .OPTIONAL_OUTPUT(grad_weight, TensorType({DT_FLOAT}))
应用建议
likedislike
不准确？
47+ .ATTR(clamp_limit, Float, 0.0)
48+ .OP_END_FACTORY_REG(SwigluGroupQuantGrad)
49+} // namespace ge
50+ 
51+#endif // OPS_QUANT_SWIGLU_GROUP_QUANT_GRAD_PROTO_H_
Aactivation/swiglu_group_quant_grad/op_host/CMakeLists.txt+12-0
@@ -0,0 +1,12 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------
10+ 
11+message(STATUS "=== Debug: start activation.swiglu_group_quant_grad.op_host.CMakeLists.txt ")
12+add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE swiglu_group_quant_grad ACLNNTYPE aclnn)
Aactivation/swiglu_group_quant_grad/op_host/config/ascend950/swiglu_group_quant_grad_binary.json+257-0
@@ -0,0 +1,257 @@
1+{
2+ "op_type": "SwigluGroupQuantGrad",
3+ "op_list": [
4+ {
5+ "bin_filename": "SwigluGroupQuantGrad_float16",
6+ "inputs": [
7+ {
8+ "name": "grad_y",
9+ "index": 0,
10+ "dtype": "float16",
11+ "format": "ND",
12+ "paramType": "required",
13+ "shape": [
14+ -2
15+ ]
16+ },
17+ {
18+ "name": "x",
19+ "index": 1,
20+ "dtype": "float16",
21+ "format": "ND",
22+ "paramType": "required",
23+ "shape": [
24+ -2
25+ ]
26+ },
27+ {
28+ "name": "weight",
29+ "index": 2,
30+ "dtype": "float32",
31+ "format": "ND",
32+ "paramType": "optional",
33+ "shape": [
34+ -2
35+ ]
36+ },
37+ {
38+ "name": "y_origin",
39+ "index": 3,
40+ "dtype": "float16",
41+ "format": "ND",
42+ "paramType": "optional",
43+ "shape": [
44+ -2
45+ ]
46+ },
47+ {
48+ "name": "group_index",
49+ "index": 4,
50+ "dtype": "int64",
51+ "format": "ND",
52+ "paramType": "optional",
53+ "shape": [
54+ -2
55+ ]
56+ }
57+ ],
58+ "outputs": [
59+ {
60+ "name": "grad_x",
61+ "index": 0,
62+ "dtype": "float16",
63+ "format": "ND",
64+ "paramType": "required",
65+ "shape": [
66+ -2
67+ ]
68+ },
69+ {
70+ "name": "grad_weight",
71+ "index": 1,
72+ "dtype": "float32",
73+ "format": "ND",
74+ "paramType": "optional",
75+ "shape": [
76+ -2
77+ ]
78+ }
79+ ],
80+ "attrs": [
81+ {
82+ "name": "clamp_limit",
83+ "dtype": "float",
84+ "value": 0.0
85+ }
86+ ]
87+ },
88+ {
89+ "bin_filename": "SwigluGroupQuantGrad_bfloat16",
90+ "inputs": [
91+ {
92+ "name": "grad_y",
93+ "index": 0,
94+ "dtype": "bfloat16",
95+ "format": "ND",
96+ "paramType": "required",
97+ "shape": [
98+ -2
99+ ]
100+ },
101+ {
102+ "name": "x",
103+ "index": 1,
104+ "dtype": "bfloat16",
105+ "format": "ND",
106+ "paramType": "required",
107+ "shape": [
108+ -2
109+ ]
110+ },
111+ {
112+ "name": "weight",
113+ "index": 2,
114+ "dtype": "float32",
115+ "format": "ND",
116+ "paramType": "optional",
117+ "shape": [
118+ -2
119+ ]
120+ },
121+ {
122+ "name": "y_origin",
123+ "index": 3,
124+ "dtype": "bfloat16",
125+ "format": "ND",
126+ "paramType": "optional",
127+ "shape": [
128+ -2
129+ ]
130+ },
131+ {
132+ "name": "group_index",
133+ "index": 4,
134+ "dtype": "int64",
135+ "format": "ND",
136+ "paramType": "optional",
137+ "shape": [
138+ -2
139+ ]
140+ }
141+ ],
142+ "outputs": [
143+ {
144+ "name": "grad_x",
145+ "index": 0,
146+ "dtype": "bfloat16",
147+ "format": "ND",
148+ "paramType": "required",
149+ "shape": [
150+ -2
151+ ]
152+ },
153+ {
154+ "name": "grad_weight",
155+ "index": 1,
156+ "dtype": "float32",
157+ "format": "ND",
158+ "paramType": "optional",
159+ "shape": [
160+ -2
161+ ]
162+ }
163+ ],
164+ "attrs": [
165+ {
166+ "name": "clamp_limit",
167+ "dtype": "float",
168+ "value": 0.0
169+ }
170+ ]
171+ },
172+ {
173+ "bin_filename": "SwigluGroupQuantGrad_float32",
174+ "inputs": [
175+ {
176+ "name": "grad_y",
177+ "index": 0,
178+ "dtype": "float32",
179+ "format": "ND",
180+ "paramType": "required",
181+ "shape": [
182+ -2
183+ ]
184+ },
185+ {
186+ "name": "x",
187+ "index": 1,
188+ "dtype": "float32",
189+ "format": "ND",
190+ "paramType": "required",
191+ "shape": [
192+ -2
193+ ]
194+ },
195+ {
196+ "name": "weight",
197+ "index": 2,
198+ "dtype": "float32",
199+ "format": "ND",
200+ "paramType": "optional",
201+ "shape": [
202+ -2
203+ ]
204+ },
205+ {
206+ "name": "y_origin",
207+ "index": 3,
208+ "dtype": "float32",
209+ "format": "ND",
210+ "paramType": "optional",
211+ "shape": [
212+ -2
213+ ]
214+ },
215+ {
216+ "name": "group_index",
217+ "index": 4,
218+ "dtype": "int64",
219+ "format": "ND",
220+ "paramType": "optional",
221+ "shape": [
222+ -2
223+ ]
224+ }
225+ ],
226+ "outputs": [
227+ {
228+ "name": "grad_x",
229+ "index": 0,
230+ "dtype": "float32",
231+ "format": "ND",
232+ "paramType": "required",
233+ "shape": [
234+ -2
235+ ]
236+ },
237+ {
238+ "name": "grad_weight",
239+ "index": 1,
240+ "dtype": "float32",
241+ "format": "ND",
242+ "paramType": "optional",
243+ "shape": [
244+ -2
245+ ]
246+ }
247+ ],
248+ "attrs": [
249+ {
250+ "name": "clamp_limit",
251+ "dtype": "float",
252+ "value": 0.0
253+ }
254+ ]
255+ }
256+ ]
257+}
Aactivation/swiglu_group_quant_grad/op_host/swiglu_group_quant_grad_def.cpp+65-0
@@ -0,0 +1,65 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file swiglu_group_quant_grad_def.cpp
13+ * \brief SwiGLU Group Dynamic Quant Backward operator definition
14+ */
15+#include "register/op_def_registry.h"
16+ 
17+namespace ops {
18+ class SwigluGroupQuantGrad : public OpDef {
19+ public:
20+ explicit SwigluGroupQuantGrad(const char* name) : OpDef(name)
21+ {
22+ this->Input("grad_y")
23+ .ParamType(REQUIRED)
24+ .DataType({ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT})
25+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
26+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
27+ this->Input("x")
28+ .ParamType(REQUIRED)
29+ .DataType({ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT})
30+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
31+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
32+ this->Input("weight")
33+ .ParamType(OPTIONAL)
34+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT})
35+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
36+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
37+ this->Input("y_origin")
38+ .ParamType(OPTIONAL)
39+ .DataType({ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT})
40+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
41+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
42+ this->Input("group_index")
43+ .ParamType(OPTIONAL)
44+ .DataType({ge::DT_INT64, ge::DT_INT64, ge::DT_INT64})
45+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
46+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
47+ .ValueDepend(OPTIONAL);
48+ this->Output("grad_x")
49+ .ParamType(REQUIRED)
50+ .DataType({ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT})
51+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
52+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
53+ this->Output("grad_weight")
54+ .ParamType(OPTIONAL)
55+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT})
56+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
57+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
58+ this->Attr("clamp_limit").AttrType(OPTIONAL).Float(0.0);
59+ 
60+ this->AICore().AddConfig("ascend950");
61+ }
62+ };
63+ 
64+ OP_ADD(SwigluGroupQuantGrad);
65+} // namespace ops
Aactivation/swiglu_group_quant_grad/op_host/swiglu_group_quant_grad_infershape.cpp+200-0
@@ -0,0 +1,200 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file swiglu_group_quant_grad_infershape.cpp
13+ * \brief SwiGLU Group Dynamic Quant Backward shape inference
14+ */
15+ 
16+#include "register/op_impl_registry.h"
17+#include "log/log.h"
18+ 
19+using namespace ge;
20+ 
21+namespace {
22+constexpr size_t INPUT_GRAD_Y = 0;
23+constexpr size_t INPUT_X = 1;
24+constexpr size_t INPUT_WEIGHT = 2;
25+constexpr size_t INPUT_Y_ORIGIN = 3;
26+constexpr size_t INPUT_GROUP_INDEX = 4;
27+constexpr size_t OUTPUT_GRAD_X = 0;
28+constexpr size_t OUTPUT_GRAD_WEIGHT = 1;
29+constexpr size_t SPLIT_NUM = 2;
30+} // namespace
31+ 
32+namespace ops {
33+ 
34+static ge::graphStatus CheckGradYAndXShape(const gert::Shape* gradYShape, const gert::Shape* xShape)
35+{
36+ int64_t gradYDimNum = gradYShape->GetDimNum();
37+ int64_t xDimNum = xShape->GetDimNum();
38+
39+ if (gradYDimNum < 1 || xDimNum < 1) {
40+ D_OP_LOGE("SwigluGroupQuantGrad", "Input shape dimension must >= 1.");
41+ return GRAPH_FAILED;
42+ }
43+
44+ if (gradYDimNum != xDimNum) {
45+ D_OP_LOGE("SwigluGroupQuantGrad", "gradY and x shape dimension must be same.");
46+ return GRAPH_FAILED;
47+ }
48+
49+ for (int64_t i = 0; i < xDimNum; i++) {
50+ if (i < xDimNum - 1) {
51+ if (gradYShape->GetDim(i) != xShape->GetDim(i)) {
52+ D_OP_LOGE("SwigluGroupQuantGrad", "gradY and x shape must be same except last dim.");
53+ return GRAPH_FAILED;
54+ }
55+ } else {
56+ int64_t gradYDimLast = gradYShape->GetDim(i);
57+ int64_t xDimLast = xShape->GetDim(i);
58+
59+ if (xDimLast % SPLIT_NUM != 0) {
60+ D_OP_LOGE("SwigluGroupQuantGrad", "Input x last dim must be divisible by 2.");
61+ return GRAPH_FAILED;
62+ }
63+
64+ if (gradYDimLast != xDimLast / SPLIT_NUM) {
65+ D_OP_LOGE("SwigluGroupQuantGrad", "Input gradY last dim must be half of x last dim.");
66+ return GRAPH_FAILED;
67+ }
68+ }
69+ }
70+
71+ return GRAPH_SUCCESS;
72+}
73+ 
74+static ge::graphStatus CheckWeightShape(const gert::Shape* weightShape, const gert::Shape* gradYShape,
75+ int64_t gradYDimNum)
76+{
77+ int64_t weightDimNum = weightShape->GetDimNum();
78+ if (weightDimNum != gradYDimNum) {
79+ D_OP_LOGE("SwigluGroupQuantGrad", "weight and gradY shape dimension must be same.");
80+ return GRAPH_FAILED;
81+ }
82+
83+ for (int64_t i = 0; i < gradYDimNum; i++) {
84+ if (i < gradYDimNum - 1) {
85+ if (weightShape->GetDim(i) != gradYShape->GetDim(i)) {
86+ D_OP_LOGE("SwigluGroupQuantGrad", "weight and gradY shape must be same except last dim.");
87+ return GRAPH_FAILED;
88+ }
89+ } else {
90+ int64_t weightDimLast = weightShape->GetDim(i);
91+ if (weightDimLast != 1) {
92+ D_OP_LOGE("SwigluGroupQuantGrad", "weight last dim must be 1.");
93+ return GRAPH_FAILED;
94+ }
95+ }
96+ }
97+
98+ return GRAPH_SUCCESS;
99+}
100+ 
101+static ge::graphStatus CheckYOriginShape(const gert::Shape* yOriginShape, const gert::Shape* gradYShape,
102+ int64_t gradYDimNum)
103+{
104+ int64_t yOriginDimNum = yOriginShape->GetDimNum();
105+ if (yOriginDimNum != gradYDimNum) {
106+ D_OP_LOGE("SwigluGroupQuantGrad", "yOrigin and gradY shape dimension must be same.");
107+ return GRAPH_FAILED;
108+ }
109+
110+ for (int64_t i = 0; i < gradYDimNum; i++) {
111+ if (yOriginShape->GetDim(i) != gradYShape->GetDim(i)) {
112+ D_OP_LOGE("SwigluGroupQuantGrad", "yOrigin shape must be same as gradY.");
113+ return GRAPH_FAILED;
114+ }
115+ }
116+
117+ return GRAPH_SUCCESS;
118+}
119+ 
120+static ge::graphStatus CheckGroupIndexShape(const gert::Shape* groupIndexShape)
121+{
122+ int64_t groupIndexDimNum = groupIndexShape->GetDimNum();
123+ if (groupIndexDimNum != 1) {
124+ D_OP_LOGE("SwigluGroupQuantGrad", "groupIndex must be 1D tensor.");
125+ return GRAPH_FAILED;
126+ }
127+
128+ return GRAPH_SUCCESS;
129+}
130+ 
131+static ge::graphStatus InferShapeForSwigluGroupQuantGrad(gert::InferShapeContext* context)
132+{
133+ OP_LOGD(context, "Enter SwigluGroupQuantGrad InferShape impl.");
134+
135+ auto gradYShape = context->GetInputShape(INPUT_GRAD_Y);
136+ OP_CHECK_NULL_WITH_CONTEXT(context, gradYShape);
137+
138+ auto xShape = context->GetInputShape(INPUT_X);
139+ OP_CHECK_NULL_WITH_CONTEXT(context, xShape);
140+
141+ auto gradXShape = context->GetOutputShape(OUTPUT_GRAD_X);
142+ OP_CHECK_NULL_WITH_CONTEXT(context, gradXShape);
143+
144+ auto gradWeightShape = context->GetOutputShape(OUTPUT_GRAD_WEIGHT);
145+ OP_CHECK_NULL_WITH_CONTEXT(context, gradWeightShape);
146+
147+ if (CheckGradYAndXShape(gradYShape, xShape) != GRAPH_SUCCESS) {
148+ return GRAPH_FAILED;
149+ }
150+
151+ *gradXShape = *xShape;
152+
153+ int64_t gradYDimNum = gradYShape->GetDimNum();
154+
155+ auto weightShape = context->GetOptionalInputShape(INPUT_WEIGHT);
156+ if (weightShape != nullptr) {
157+ if (CheckWeightShape(weightShape, gradYShape, gradYDimNum) != GRAPH_SUCCESS) {
158+ return GRAPH_FAILED;
159+ }
160+ *gradWeightShape = *weightShape;
161+ }
162+
163+ auto yOriginShape = context->GetOptionalInputShape(INPUT_Y_ORIGIN);
164+ if (yOriginShape != nullptr) {
165+ if (CheckYOriginShape(yOriginShape, gradYShape, gradYDimNum) != GRAPH_SUCCESS) {
166+ return GRAPH_FAILED;
167+ }
168+ }
169+
170+ auto groupIndexShape = context->GetOptionalInputShape(INPUT_GROUP_INDEX);
171+ if (groupIndexShape != nullptr) {
172+ if (CheckGroupIndexShape(groupIndexShape) != GRAPH_SUCCESS) {
173+ return GRAPH_FAILED;
174+ }
175+ }
176+
177+ OP_LOGD(context, "SwigluGroupQuantGrad InferShape impl end.");
178+ return ge::GRAPH_SUCCESS;
179+}
180+ 
181+static ge::graphStatus InferDataTypeForSwigluGroupQuantGrad(gert::InferDataTypeContext *context) {
182+ OP_LOGD(context, "Enter SwigluGroupQuantGrad inferDataType impl.");
183+
184+ auto xDtype = context->GetInputDataType(INPUT_X);
185+ context->SetOutputDataType(OUTPUT_GRAD_X, xDtype);
186+
187+ auto weightDesc = context->GetOptionalInputDesc(INPUT_WEIGHT);
188+ if (weightDesc != nullptr) {
189+ context->SetOutputDataType(OUTPUT_GRAD_WEIGHT, ge::DT_FLOAT);
190+ }
191+
192+ OP_LOGD(context, "SwigluGroupQuantGrad inferDataType impl end.");
193+ return ge::GRAPH_SUCCESS;
194+}
195+ 
196+IMPL_OP_INFERSHAPE(SwigluGroupQuantGrad)
197+ .InferShape(InferShapeForSwigluGroupQuantGrad)
198+ .InferDataType(InferDataTypeForSwigluGroupQuantGrad);
199+ 
200+} // namespace ops
Aactivation/swiglu_group_quant_grad/op_host/swiglu_group_quant_grad_tiling.cpp+68-0
@@ -0,0 +1,68 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file swiglu_group_quant_grad_tiling.cpp
13+ * \brief SwiGLU Group Dynamic Quant Backward tiling implementation
14+ */
15+ 
16+#include "register/op_def_registry.h"
17+#include "swiglu_group_quant_grad_tiling_utils.h"
18+ 
19+namespace optiling {
20+ 
21+constexpr uint32_t BATCH_MODE = 1;
22+ 
23+static ge::graphStatus Tiling4SwigluGroupQuantGrad(gert::TilingContext *context)
24+{
25+ OP_LOGD(context, "Tiling4SwigluGroupQuantGrad start.");
26+ context->SetScheduleMode(BATCH_MODE);
27+
28+ SwigluGroupQuantGradCompileInfo compileInfo;
29+ SwigluGroupQuantGradTilingData tilingData;
30+ if (GetCompileInfo(context, compileInfo) != ge::GRAPH_SUCCESS) {
31+ return ge::GRAPH_FAILED;
32+ }
33+ if (CheckOpParams(context, compileInfo) != ge::GRAPH_SUCCESS) {
34+ return ge::GRAPH_FAILED;
35+ }
36+
37+ SetBasicTilingData(context, compileInfo, tilingData);
38+ CalculateTilingParams(context, compileInfo, tilingData);
39+ if (SetTilingDataToContext(context, tilingData) != ge::GRAPH_SUCCESS) {
40+ return ge::GRAPH_FAILED;
41+ }
42+
43+ context->SetBlockDim(tilingData.get_usedCoreNum());
44+
45+ size_t *workspaces = context->GetWorkspaceSizes(1);
46+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
47+ uint32_t sysWorkspaceSize = ascendcPlatform.GetLibApiWorkSpaceSize();
48+ workspaces[0] = sysWorkspaceSize + tilingData.get_usedCoreNum() * BLOCK_SIZE;
49+
50+ OP_LOGD(context,
51+ "Tiling4SwigluGroupQuantGrad end. usedCoreNum=%u, totalTokens=%u, truncValue=%u, tileH=%u, tileTokens=%u",
52+ tilingData.get_usedCoreNum(), tilingData.get_totalTokens(), tilingData.get_truncValue(),
53+ tilingData.get_tileH(), tilingData.get_tileTokens());
54+
55+ return ge::GRAPH_SUCCESS;
56+}
57+ 
58+static ge::graphStatus TilingPrepare4SwigluGroupQuantGrad(gert::TilingParseContext *context)
59+{
60+ OP_LOGD(context, "TilingPrepare4SwigluGroupQuantGrad start and end.");
61+ return ge::GRAPH_SUCCESS;
62+}
63+ 
64+IMPL_OP_OPTILING(SwigluGroupQuantGrad)
65+ .Tiling(Tiling4SwigluGroupQuantGrad)
66+ .TilingParse<CoreCompileInfo>(TilingPrepare4SwigluGroupQuantGrad);
67+ 
68+} // namespace optiling
Aactivation/swiglu_group_quant_grad/op_host/swiglu_group_quant_grad_tiling.h+72-0
@@ -0,0 +1,72 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file swiglu_group_quant_grad_tiling.h
13+ * \brief SwiGLU Group Dynamic Quant Backward tiling data definition
14+ */
15+ 
16+#ifndef SWIGLU_GROUP_QUANT_GRAD_TILING_H
17+#define SWIGLU_GROUP_QUANT_GRAD_TILING_H
18+ 
19+#include "register/tilingdata_base.h"
20+ 
21+namespace optiling {
22+ 
23+BEGIN_TILING_DATA_DEF(SwigluGroupQuantGradTilingData)
24+ TILING_DATA_FIELD_DEF(uint32_t, coreNumAll);
25+ TILING_DATA_FIELD_DEF(uint32_t, ubSize);
26+
27+ TILING_DATA_FIELD_DEF(uint32_t, totalTokens);
28+ TILING_DATA_FIELD_DEF(uint32_t, dim2H);
29+ TILING_DATA_FIELD_DEF(uint32_t, dimH);
30+
31+ TILING_DATA_FIELD_DEF(uint32_t, hasWeight);
32+ TILING_DATA_FIELD_DEF(uint32_t, hasYOrigin);
33+ TILING_DATA_FIELD_DEF(uint32_t, hasGroupIndex);
34+ TILING_DATA_FIELD_DEF(uint32_t, hasClampLimit);
35+ TILING_DATA_FIELD_DEF(uint32_t, needSplitH);
36+
37+ TILING_DATA_FIELD_DEF(float, clampLimit);
38+
39+ TILING_DATA_FIELD_DEF(uint32_t, groupNum);
40+ TILING_DATA_FIELD_DEF(uint32_t, truncValue);
41+
42+ TILING_DATA_FIELD_DEF(uint32_t, tileTokens);
43+ TILING_DATA_FIELD_DEF(uint32_t, tileH);
44+ TILING_DATA_FIELD_DEF(uint32_t, numHTiles);
45+ TILING_DATA_FIELD_DEF(uint32_t, totalTiles);
46+
47+ TILING_DATA_FIELD_DEF(uint32_t, usedCoreNum);
48+ TILING_DATA_FIELD_DEF(uint32_t, tokensPerCore);
49+ TILING_DATA_FIELD_DEF(uint32_t, coreTokenStart);
S
Ssunday6月18日

coreTokenStart字段在tiling data中声明但从未赋值（SetBasicTilingData中不设置），kernel base.h中直接计算tokenStart=blockIdx*tokensPerCore而非从tiling data读取。这是死代码，建议删除。

likedislike
S
shilulu
6月18日 评论：
50+END_TILING_DATA_DEF;
51+ 
52+REGISTER_TILING_DATA_CLASS(SwigluGroupQuantGrad, SwigluGroupQuantGradTilingData)
53+ 
54+struct CoreCompileInfo {
55+};
56+ 
57+struct SwigluGroupQuantGradCompileInfo {
58+ uint32_t totalCore = 1;
59+ uint32_t ubSize = 0;
60+ uint32_t inputDataByte = 4;
61+ float clampLimit = 0.0f;
62+ uint32_t hasWeight = 0;
63+ uint32_t hasYOrigin = 0;
64+ uint32_t hasGroupIndex = 0;
65+
66+ uint32_t dataNumSingleUb = 1;
67+ uint32_t blockNum = 8;
68+};
69+ 
70+} // namespace optiling
71+ 
72+#endif // SWIGLU_GROUP_QUANT_GRAD_TILING_H
Aactivation/swiglu_group_quant_grad/op_host/swiglu_group_quant_grad_tiling_utils.h+529-0
@@ -0,0 +1,529 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file swiglu_group_quant_grad_tiling_utils.h
13+ * \brief SwiGLU Group Dynamic Quant Backward tiling utility functions
14+ */
15+ 
16+#ifndef SWIGLU_GROUP_QUANT_GRAD_TILING_UTILS_H
17+#define SWIGLU_GROUP_QUANT_GRAD_TILING_UTILS_H
18+ 
19+#include "register/op_impl_registry.h"
20+#include "util/math_util.h"
21+#include "log/log.h"
22+#include "tiling/platform/platform_ascendc.h"
23+#include "platform/platform_infos_def.h"
24+#include "swiglu_group_quant_grad_tiling.h"
25+ 
26+namespace optiling {
27+ 
28+constexpr uint32_t INPUT_GRAD_Y_INDEX = 0;
29+constexpr uint32_t INPUT_X_INDEX = 1;
30+constexpr uint32_t INPUT_WEIGHT_INDEX = 2;
31+constexpr uint32_t INPUT_Y_ORIGIN_INDEX = 3;
32+constexpr uint32_t INPUT_GROUP_INDEX_INDEX = 4;
33+constexpr uint32_t OUTPUT_GRAD_X_INDEX = 0;
34+constexpr uint32_t OUTPUT_GRAD_WEIGHT_INDEX = 1;
35+ 
36+constexpr uint32_t TMP_DATA_UB_SIZE = 8 * 1024;
37+constexpr uint32_t BLOCK_SIZE = 32;
38+constexpr uint32_t ALIGN_EIGHT = 8;
39+constexpr uint32_t FP16_BFP16_32B_ALIGN_NUM = 16;
40+constexpr uint32_t ZERO = 0;
41+constexpr uint32_t ONE = 1;
42+constexpr uint32_t TWO = 2;
43+ 
44+constexpr uint32_t UB_BASE_FACTOR = 12;
45+constexpr uint32_t UB_CLAMP_EXTRA_FACTOR = 8;
46+constexpr uint32_t UB_WEIGHT_EXTRA_FACTOR = 4;
47+constexpr uint32_t UB_WEIGHT_EXTRA_TOKENS = 16;
48+ 
49+constexpr uint32_t MAX_H = 4096;
50+constexpr uint32_t MIN_H = 512;
51+ 
52+template <typename T>
53+inline auto AlignUp(T num, T div) -> decltype(num)
54+{
55+ return (div == 0) ? 0 : (num + div - 1) / div * div;
56+}
57+ 
58+template <typename T>
59+inline auto AlignDown(T num, T div) -> decltype(num)
60+{
61+ return (div == 0) ? 0 : num / div * div;
62+}
63+ 
64+template <typename T>
65+inline auto CeilDiv(T num, T div) -> decltype(num)
66+{
67+ return div == 0 ? 0 : (num + div - 1) / div;
68+}
69+ 
70+inline ge::graphStatus CheckInputDtype(const gert::TilingContext *context)
71+{
72+ auto gradYDtype = context->GetInputDesc(INPUT_GRAD_Y_INDEX)->GetDataType();
73+ auto xDtype = context->GetInputDesc(INPUT_X_INDEX)->GetDataType();
74+
75+ if (gradYDtype != ge::DT_FLOAT16 && gradYDtype != ge::DT_BF16 && gradYDtype != ge::DT_FLOAT) {
76+ OP_LOGE(context->GetNodeName(), "input grad_y dtype is only support fp16/bf16/fp32.");
77+ return ge::GRAPH_FAILED;
78+ }
79+
80+ if (xDtype != ge::DT_FLOAT16 && xDtype != ge::DT_BF16 && xDtype != ge::DT_FLOAT) {
81+ OP_LOGE(context->GetNodeName(), "input x dtype is only support fp16/bf16/fp32.");
82+ return ge::GRAPH_FAILED;
83+ }
84+
85+ if (gradYDtype != xDtype) {
86+ OP_LOGE(context->GetNodeName(), "input grad_y and x dtype must be same.");
87+ return ge::GRAPH_FAILED;
88+ }
89+
90+ auto weightDesc = context->GetOptionalInputDesc(INPUT_WEIGHT_INDEX);
91+ if (weightDesc != nullptr) {
92+ auto weightDtype = weightDesc->GetDataType();
93+ if (weightDtype != ge::DataType::DT_FLOAT) {
94+ OP_LOGE(context->GetNodeName(), "input weight dtype is only support fp32.");
95+ return ge::GRAPH_FAILED;
96+ }
97+ }
98+
99+ auto yOriginDesc = context->GetOptionalInputDesc(INPUT_Y_ORIGIN_INDEX);
100+ if (yOriginDesc != nullptr) {
101+ auto yOriginDtype = yOriginDesc->GetDataType();
102+ if (yOriginDtype != ge::DT_FLOAT16 && yOriginDtype != ge::DT_BF16 && yOriginDtype != ge::DT_FLOAT) {
103+ OP_LOGE(context->GetNodeName(), "input y_origin dtype is only support fp16/bf16/fp32.");
104+ return ge::GRAPH_FAILED;
105+ }
106+ }
atomgit-bot
atomgit-botatomgit-bot6月18日

🟡 Medium Priority

CheckInputDtype（tiling_utils.h 第70-118行）在校验 yOrigin dtype 时（第99-106行）仅检查了是否属于 {FLOAT16, BF16, FLOAT} 三种类型之一，但未检查其与 gradY 的 dtype 是否一致。根据算子规格约束："gradY、x、yOrigin、gradXOut数据类型必须一致"。如果 yOrigin 的 dtype 与 gradY 不同（例如 gradY 为 FLOAT16 而 yOrigin 为 FLOAT），缺少此校验会导致 tiling 阶段通过，但 kernel 计算时可能产生错误结果。

建议：在 yOriginDesc != nullptr 分支中，增加 if (yOriginDtype != gradYDtype) 的检查并返回 GRAPH_FAILED。

likedislike
不准确？
107+
108+ auto groupIndexDesc = context->GetOptionalInputDesc(INPUT_GROUP_INDEX_INDEX);
109+ if (groupIndexDesc != nullptr) {
110+ auto groupIndexDtype = groupIndexDesc->GetDataType();
111+ if (groupIndexDtype != ge::DataType::DT_INT64) {
112+ OP_LOGE(context->GetNodeName(), "input group_index dtype is only support int64.");
113+ return ge::GRAPH_FAILED;
114+ }
115+ }
116+
117+ return ge::GRAPH_SUCCESS;
118+}
119+ 
120+inline ge::graphStatus CheckOutputDtype(const gert::TilingContext *context)
121+{
122+ auto gradXDtype = context->GetOutputDesc(OUTPUT_GRAD_X_INDEX)->GetDataType();
123+ auto xDtype = context->GetInputDesc(INPUT_X_INDEX)->GetDataType();
124+
125+ if (gradXDtype != xDtype) {
126+ OP_LOGE(context->GetNodeName(), "output grad_x dtype must be same as input x.");
127+ return ge::GRAPH_FAILED;
128+ }
129+
130+ auto gradWeightDesc = context->GetOutputDesc(OUTPUT_GRAD_WEIGHT_INDEX);
131+ if (gradWeightDesc != nullptr) {
132+ auto gradWeightDtype = gradWeightDesc->GetDataType();
133+ if (gradWeightDtype != ge::DataType::DT_FLOAT) {
134+ OP_LOGE(context->GetNodeName(), "output grad_weight dtype is only support fp32.");
135+ return ge::GRAPH_FAILED;
136+ }
137+ }
138+
139+ return ge::GRAPH_SUCCESS;
140+}
141+ 
142+inline ge::graphStatus CheckAttrs(const gert::TilingContext *context, SwigluGroupQuantGradCompileInfo &compileInfo)
143+{
144+ const gert::RuntimeAttrs *attrs = context->GetAttrs();
145+ if (attrs != nullptr) {
146+ auto clampLimitPtr = attrs->GetFloat(0);
147+ if (clampLimitPtr != nullptr) {
148+ compileInfo.clampLimit = *clampLimitPtr;
149+ }
150+ }
151+ return ge::GRAPH_SUCCESS;
152+}
atomgit-bot
atomgit-botatomgit-bot6月18日

🟡 Medium Priority

算子规格明确约束 "clampLimit 必须 ≥ 0.0"。但 CheckAttrs（tiling_utils.h 第142-152行）仅读取 clampLimit 值而不校验其合法性；InferShapeForSwigluGroupQuantGrad（infershape.cpp）完全不检查 clampLimit 属性。如果用户传入负的 clampLimit（如 -1.0），tiling 阶段不会报错，kernel 会将 hasClampLimit = (clampLimit != 0.0f) 设为 1，然后以负阈值执行 Clamp 操作（如 Mins(x, x, -1.0) 和 Maxs(x, x, 1.0)），产生非预期的计算结果且无错误提示。

建议：在 CheckAttrs 中，读取 clampLimit 后增加 if (compileInfo.clampLimit < 0.0f) 的检查并返回 GRAPH_FAILED。同时在 InferShapeForSwigluGroupQuantGrad 中也增加相同的属性校验。

likedislike
不准确？
153+ 
154+inline ge::graphStatus CheckGradYAndXShapeDim(const gert::TilingContext *context)
155+{
156+ auto gradYShape = context->GetInputShape(INPUT_GRAD_Y_INDEX);
157+ auto xShape = context->GetInputShape(INPUT_X_INDEX);
158+
159+ size_t gradYDimNum = gradYShape->GetStorageShape().GetDimNum();
160+ size_t xDimNum = xShape->GetStorageShape().GetDimNum();
161+
162+ if (gradYDimNum < ONE || xDimNum < ONE) {
163+ OP_LOGE(context->GetNodeName(), "Input shape dimension must >= 1.");
164+ return ge::GRAPH_FAILED;
165+ }
166+
167+ if (gradYDimNum != xDimNum) {
168+ OP_LOGE(context->GetNodeName(), "grad_y and x shape dimension must be same.");
169+ return ge::GRAPH_FAILED;
170+ }
171+
172+ for (size_t i = 0; i < xDimNum; i++) {
173+ if (i < xDimNum - 1) {
174+ if (gradYShape->GetStorageShape().GetDim(i) != xShape->GetStorageShape().GetDim(i)) {
175+ OP_LOGE(context->GetNodeName(), "grad_y and x shape must be same except last dim.");
176+ return ge::GRAPH_FAILED;
177+ }
178+ } else {
179+ int64_t gradYDimLast = gradYShape->GetStorageShape().GetDim(i);
180+ int64_t xDimLast = xShape->GetStorageShape().GetDim(i);
181+
182+ if (xDimLast % TWO != 0) {
183+ OP_LOGE(context->GetNodeName(), "Input x last dim must be divisible by 2.");
184+ return ge::GRAPH_FAILED;
185+ }
186+
187+ if (gradYDimLast != xDimLast / TWO) {
188+ OP_LOGE(context->GetNodeName(), "Input grad_y last dim must be half of x last dim.");
189+ return ge::GRAPH_FAILED;
190+ }
191+ }
192+ }
193+
194+ return ge::GRAPH_SUCCESS;
195+}
196+ 
197+inline ge::graphStatus CheckWeightShapeDim(const gert::TilingContext *context,
198+ SwigluGroupQuantGradCompileInfo &compileInfo)
199+{
200+ auto weightShape = context->GetOptionalInputShape(INPUT_WEIGHT_INDEX);
201+ if (weightShape == nullptr) {
202+ return ge::GRAPH_SUCCESS;
203+ }
204+
205+ compileInfo.hasWeight = 1;
206+ auto gradYShape = context->GetInputShape(INPUT_GRAD_Y_INDEX);
207+ size_t gradYDimNum = gradYShape->GetStorageShape().GetDimNum();
208+ size_t weightDimNum = weightShape->GetStorageShape().GetDimNum();
209+
210+ if (weightDimNum != gradYDimNum) {
211+ OP_LOGE(context->GetNodeName(), "weight and grad_y shape dimension must be same.");
212+ return ge::GRAPH_FAILED;
213+ }
214+
215+ for (size_t i = 0; i < gradYDimNum; i++) {
216+ if (i < gradYDimNum - 1) {
217+ if (weightShape->GetStorageShape().GetDim(i) != gradYShape->GetStorageShape().GetDim(i)) {
218+ OP_LOGE(context->GetNodeName(), "weight and grad_y shape must be same except last dim.");
219+ return ge::GRAPH_FAILED;
220+ }
221+ } else {
222+ int64_t weightDimLast = weightShape->GetStorageShape().GetDim(i);
223+ if (weightDimLast != ONE) {
224+ OP_LOGE(context->GetNodeName(), "weight last dim must be 1.");
225+ return ge::GRAPH_FAILED;
226+ }
227+ }
228+ }
229+
230+ return ge::GRAPH_SUCCESS;
231+}
232+ 
233+inline ge::graphStatus CheckYOriginShapeDim(const gert::TilingContext *context,
234+ SwigluGroupQuantGradCompileInfo &compileInfo)
235+{
236+ auto yOriginShape = context->GetOptionalInputShape(INPUT_Y_ORIGIN_INDEX);
237+ if (yOriginShape == nullptr) {
238+ return ge::GRAPH_SUCCESS;
239+ }
240+
241+ compileInfo.hasYOrigin = 1;
242+ auto gradYShape = context->GetInputShape(INPUT_GRAD_Y_INDEX);
243+ size_t gradYDimNum = gradYShape->GetStorageShape().GetDimNum();
244+ size_t yOriginDimNum = yOriginShape->GetStorageShape().GetDimNum();
245+
246+ if (yOriginDimNum != gradYDimNum) {
247+ OP_LOGE(context->GetNodeName(), "y_origin and grad_y shape dimension must be same.");
248+ return ge::GRAPH_FAILED;
249+ }
250+
251+ for (size_t i = 0; i < gradYDimNum; i++) {
252+ if (yOriginShape->GetStorageShape().GetDim(i) != gradYShape->GetStorageShape().GetDim(i)) {
253+ OP_LOGE(context->GetNodeName(), "y_origin shape must be same as grad_y.");
254+ return ge::GRAPH_FAILED;
255+ }
256+ }
257+
258+ return ge::GRAPH_SUCCESS;
259+}
260+ 
261+inline ge::graphStatus CheckGroupIndexShapeDim(const gert::TilingContext *context,
262+ SwigluGroupQuantGradCompileInfo &compileInfo)
263+{
264+ auto groupIndexShape = context->GetOptionalInputShape(INPUT_GROUP_INDEX_INDEX);
265+ if (groupIndexShape == nullptr) {
266+ return ge::GRAPH_SUCCESS;
267+ }
268+
269+ compileInfo.hasGroupIndex = 1;
270+ size_t groupIndexDimNum = groupIndexShape->GetStorageShape().GetDimNum();
271+ if (groupIndexDimNum != ONE) {
272+ OP_LOGE(context->GetNodeName(), "group_index must be 1D tensor.");
273+ return ge::GRAPH_FAILED;
274+ }
275+
276+ return ge::GRAPH_SUCCESS;
277+}
278+ 
279+inline ge::graphStatus CheckInputShape(const gert::TilingContext *context, SwigluGroupQuantGradCompileInfo &compileInfo)
280+{
281+ if (CheckGradYAndXShapeDim(context) != ge::GRAPH_SUCCESS) {
282+ return ge::GRAPH_FAILED;
283+ }
284+
285+ if (CheckWeightShapeDim(context, compileInfo) != ge::GRAPH_SUCCESS) {
286+ return ge::GRAPH_FAILED;
287+ }
288+
289+ if (CheckYOriginShapeDim(context, compileInfo) != ge::GRAPH_SUCCESS) {
290+ return ge::GRAPH_FAILED;
291+ }
292+
293+ if (compileInfo.hasWeight && !compileInfo.hasYOrigin) {
294+ OP_LOGE(context->GetNodeName(), "When weight exists, y_origin must also exist.");
295+ return ge::GRAPH_FAILED;
296+ }
297+
298+ if (CheckGroupIndexShapeDim(context, compileInfo) != ge::GRAPH_SUCCESS) {
299+ return ge::GRAPH_FAILED;
300+ }
301+
302+ return ge::GRAPH_SUCCESS;
303+}
304+ 
305+inline ge::graphStatus CheckOutputShape(const gert::TilingContext *context)
306+{
307+ auto xShape = context->GetInputShape(INPUT_X_INDEX);
308+ auto gradXShape = context->GetOutputShape(OUTPUT_GRAD_X_INDEX);
309+
310+ size_t xDimNum = xShape->GetStorageShape().GetDimNum();
311+ size_t gradXDimNum = gradXShape->GetStorageShape().GetDimNum();
312+
313+ if (xDimNum != gradXDimNum) {
314+ OP_LOGE(context->GetNodeName(), "x and grad_x shape dimension must be same.");
315+ return ge::GRAPH_FAILED;
316+ }
317+
318+ for (size_t i = 0; i < xDimNum; i++) {
319+ if (xShape->GetStorageShape().GetDim(i) != gradXShape->GetStorageShape().GetDim(i)) {
320+ OP_LOGE(context->GetNodeName(), "x and grad_x shape must be same.");
321+ return ge::GRAPH_FAILED;
322+ }
323+ }
324+
325+ auto weightShape = context->GetOptionalInputShape(INPUT_WEIGHT_INDEX);
326+ if (weightShape != nullptr) {
327+ auto gradWeightShape = context->GetOutputShape(OUTPUT_GRAD_WEIGHT_INDEX);
328+
329+ size_t gradWeightDimNum = gradWeightShape->GetStorageShape().GetDimNum();
330+ size_t weightDimNum = weightShape->GetStorageShape().GetDimNum();
331+
332+ if (gradWeightDimNum != weightDimNum) {
333+ OP_LOGE(context->GetNodeName(), "grad_weight and weight shape dimension must be same.");
334+ return ge::GRAPH_FAILED;
335+ }
336+
337+ for (size_t i = 0; i < weightDimNum; i++) {
338+ if (gradWeightShape->GetStorageShape().GetDim(i) != weightShape->GetStorageShape().GetDim(i)) {
339+ OP_LOGE(context->GetNodeName(), "grad_weight shape must be same as weight.");
340+ return ge::GRAPH_FAILED;
341+ }
342+ }
343+ }
344+
345+ return ge::GRAPH_SUCCESS;
346+}
347+ 
348+inline ge::graphStatus CheckOpParams(gert::TilingContext *context, SwigluGroupQuantGradCompileInfo &compileInfo)
349+{
350+ if (CheckInputDtype(context) != ge::GRAPH_SUCCESS) {
351+ OP_LOGE(context->GetNodeName(), "Check input dtype failed.");
352+ return ge::GRAPH_FAILED;
353+ }
354+ if (CheckOutputDtype(context) != ge::GRAPH_SUCCESS) {
355+ OP_LOGE(context->GetNodeName(), "Check output dtype failed.");
356+ return ge::GRAPH_FAILED;
357+ }
358+ if (CheckAttrs(context, compileInfo) != ge::GRAPH_SUCCESS) {
359+ OP_LOGE(context->GetNodeName(), "Check attrs failed.");
360+ return ge::GRAPH_FAILED;
361+ }
362+ if (CheckInputShape(context, compileInfo) != ge::GRAPH_SUCCESS) {
363+ OP_LOGE(context->GetNodeName(), "Check input shape failed.");
364+ return ge::GRAPH_FAILED;
365+ }
366+ if (CheckOutputShape(context) != ge::GRAPH_SUCCESS) {
367+ OP_LOGE(context->GetNodeName(), "Check output shape failed.");
368+ return ge::GRAPH_FAILED;
369+ }
370+ return ge::GRAPH_SUCCESS;
371+}
372+ 
373+inline uint32_t GetTotalTokens(const gert::StorageShape *shape, gert::TilingContext *context)
374+{
375+ uint32_t totalTokens = 1;
376+ size_t dimNum = shape->GetStorageShape().GetDimNum();
377+ for (size_t i = 0; i < dimNum - 1; i++) {
378+ totalTokens *= static_cast<uint32_t>(shape->GetStorageShape().GetDim(i));
379+ }
380+ return totalTokens;
381+}
382+ 
383+inline uint32_t GetTruncValue(const gert::TilingContext *context, uint32_t totalTokens)
384+{
385+ auto groupIndexTensor = context->GetOptionalInputTensor(INPUT_GROUP_INDEX_INDEX);
386+ if (groupIndexTensor == nullptr) {
387+ return totalTokens;
388+ }
389+
390+ uint32_t truncValue = 0;
391+ const int64_t* groupIndexData = groupIndexTensor->GetData<int64_t>();
392+ if (groupIndexData == nullptr) {
393+ return totalTokens;
394+ }
395+ size_t groupNum = groupIndexTensor->GetStorageShape().GetDim(0);
396+ for (size_t i = 0; i < groupNum; i++) {
397+ truncValue += static_cast<uint32_t>(groupIndexData[i]);
398+ }
399+ return std::min(truncValue, totalTokens);
400+}
401+ 
402+inline void CalculateTilingParams(const gert::TilingContext *context,
403+ SwigluGroupQuantGradCompileInfo &compileInfo,
404+ SwigluGroupQuantGradTilingData &tilingData)
405+{
406+ uint32_t dimH = tilingData.get_dimH();
407+ uint32_t truncValue = tilingData.get_truncValue();
408+
409+ uint32_t hasClampLimit = (compileInfo.clampLimit != 0.0f) ? 1 : 0;
410+ uint32_t hasWeight = compileInfo.hasWeight;
411+
412+ uint32_t ubFactor = UB_BASE_FACTOR;
413+ if (hasClampLimit) {
414+ ubFactor += UB_CLAMP_EXTRA_FACTOR;
415+ }
416+ if (hasWeight) {
417+ ubFactor += UB_WEIGHT_EXTRA_FACTOR;
418+ }
419+
420+ uint32_t ubAvailable = compileInfo.ubSize - BLOCK_SIZE - TMP_DATA_UB_SIZE;
421+ uint32_t ubPerTokenFullH = ubFactor * dimH * sizeof(float);
422+
423+ if (hasWeight) {
424+ ubPerTokenFullH += UB_WEIGHT_EXTRA_TOKENS * sizeof(float);
425+ }
426+
427+ uint32_t needSplitH = (ubPerTokenFullH > ubAvailable) ? 1 : 0;
428+ tilingData.set_needSplitH(needSplitH);
429+
430+ uint32_t tileH = dimH;
431+ uint32_t tileTokens = 0;
432+
433+ if (needSplitH == 0) {
434+ tileTokens = ubAvailable / ubPerTokenFullH;
435+ tileTokens = std::max(tileTokens, ONE);
436+ } else {
437+ uint32_t ubPerTokenMinH = ubFactor * sizeof(float);
438+ if (hasWeight) {
439+ ubAvailable -= UB_WEIGHT_EXTRA_TOKENS * sizeof(float);
440+ }
441+ tileH = ubAvailable / ubPerTokenMinH;
442+ tileH = AlignDown(tileH, FP16_BFP16_32B_ALIGN_NUM);
443+ tileH = std::max(tileH, FP16_BFP16_32B_ALIGN_NUM);
444+ tileTokens = ONE;
445+ }
446+
447+ uint32_t numHTiles = CeilDiv(dimH, tileH);
448+
449+ uint32_t usedCoreNum = std::min(truncValue, compileInfo.totalCore);
450+ usedCoreNum = std::max(usedCoreNum, ONE);
451+
452+ uint32_t tokensPerCore = CeilDiv(truncValue, usedCoreNum);
453+
454+ uint32_t totalTiles = truncValue * numHTiles;
455+ OP_LOGD(context, "ubAvailable %u needSplitH %u tileTokens %u tileH %u compileInfo.totalCore %u \n",
456+ ubAvailable, needSplitH, tileTokens, tileH, compileInfo.totalCore);
457+ tilingData.set_tileTokens(tileTokens);
458+ tilingData.set_tileH(tileH);
459+ tilingData.set_numHTiles(numHTiles);
460+ tilingData.set_totalTiles(totalTiles);
461+ tilingData.set_usedCoreNum(usedCoreNum);
462+ tilingData.set_tokensPerCore(tokensPerCore);
463+}
464+ 
465+inline ge::graphStatus SetTilingDataToContext(gert::TilingContext *context, SwigluGroupQuantGradTilingData &tilingData)
466+{
467+ tilingData.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity());
468+ context->GetRawTilingData()->SetDataSize(tilingData.GetDataSize());
469+ return ge::GRAPH_SUCCESS;
470+}
471+ 
472+inline ge::graphStatus GetCompileInfo(gert::TilingContext *context, SwigluGroupQuantGradCompileInfo &compileInfo)
473+{
474+ auto platformInfo = context->GetPlatformInfo();
475+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo);
476+ uint32_t totalCoreNum = ascendcPlatform.GetCoreNumAiv();
477+ uint64_t ubSizePlatform;
478+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSizePlatform);
479+ uint32_t ubSize = static_cast<uint32_t>(ubSizePlatform);
480+
481+ if (totalCoreNum == 0 || ubSize <= 0) {
482+ OP_LOGE(context->GetNodeName(), "GetCompileInfo failed, coreNum:%u, ubSize:%u.", totalCoreNum, ubSize);
483+ return ge::GRAPH_FAILED;
484+ }
atomgit-bot
atomgit-botatomgit-bot6月18日

🔵 Low Priority

GetCompileInfo（tiling_utils.h 第472-491行）中 ubSizePlatform 为 uint64_t 类型，第479行通过 static_cast<uint32_t>(ubSizePlatform) 窄化为 uint32_t。第481行检查 ubSize <= 0 时使用的是窄化后的值。虽然当前硬件 UB 大小远小于 4GB，窄化不会导致实际截断，但对 ubSizePlatform 的原值做零值检查更为稳健，避免未来平台变化时窄化为零而绕过检查。

建议：将空值检查提前到窄化之前：先检查 ubSizePlatform == 0，然后再 static_cast<uint32_t>，或直接使用 ubSizePlatform 进行比较。

likedislike
不准确？
485+
486+ compileInfo.totalCore = totalCoreNum;
487+ compileInfo.ubSize = ubSize;
488+ compileInfo.blockNum = BLOCK_SIZE / sizeof(float);
489+
490+ return ge::GRAPH_SUCCESS;
491+}
492+ 
493+inline void SetBasicTilingData(gert::TilingContext *context,
494+ const SwigluGroupQuantGradCompileInfo &compileInfo,
495+ SwigluGroupQuantGradTilingData &tilingData)
496+{
497+ auto gradYShape = context->GetInputShape(INPUT_GRAD_Y_INDEX);
498+ auto xShape = context->GetInputShape(INPUT_X_INDEX);
499+ uint32_t totalTokens = GetTotalTokens(gradYShape, context);
500+ uint32_t dimH = static_cast<uint32_t>(gradYShape->GetStorageShape()
501+ .GetDim(gradYShape->GetStorageShape().GetDimNum() - 1));
502+ uint32_t dim2H = static_cast<uint32_t>(xShape->GetStorageShape()
503+ .GetDim(xShape->GetStorageShape().GetDimNum() - 1));
504+ uint32_t truncValue = GetTruncValue(context, totalTokens);
505+
506+ tilingData.set_coreNumAll(compileInfo.totalCore);
507+ tilingData.set_ubSize(compileInfo.ubSize);
508+ tilingData.set_totalTokens(totalTokens);
509+ tilingData.set_dim2H(dim2H);
510+ tilingData.set_dimH(dimH);
511+ tilingData.set_hasWeight(compileInfo.hasWeight);
512+ tilingData.set_hasYOrigin(compileInfo.hasYOrigin);
513+ tilingData.set_hasGroupIndex(compileInfo.hasGroupIndex);
514+ tilingData.set_hasClampLimit(compileInfo.clampLimit != 0.0f ? 1 : 0);
515+ tilingData.set_clampLimit(compileInfo.clampLimit);
516+ tilingData.set_truncValue(truncValue);
517+
518+ auto groupIndexShape = context->GetOptionalInputShape(INPUT_GROUP_INDEX_INDEX);
519+ if (groupIndexShape != nullptr) {
520+ uint32_t groupNum = static_cast<uint32_t>(groupIndexShape->GetStorageShape().GetDim(0));
521+ tilingData.set_groupNum(groupNum);
522+ } else {
523+ tilingData.set_groupNum(0);
524+ }
525+}
526+ 
527+}
528+ 
529+#endif // SWIGLU_GROUP_QUANT_GRAD_TILING_UTILS_H
Aactivation/swiglu_group_quant_grad/op_kernel/swiglu_group_quant_grad.cpp+27-0
@@ -0,0 +1,27 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file swiglu_group_quant_grad.cpp
13+ * \brief SwiGLU Group Dynamic Quant Backward kernel entry for Ascend 950 (A5)
14+ */
15+ 
16+#include "swiglu_group_quant_grad.h"
17+ 
18+using namespace SwigluGroupQuantGradOp;
19+ 
20+extern "C" __global__ __aicore__ void swiglu_group_quant_grad(GM_ADDR gradY, GM_ADDR x, GM_ADDR weight,
21+ GM_ADDR yOrigin, GM_ADDR groupIndex, GM_ADDR gradX,
22+ GM_ADDR gradWeight, GM_ADDR workspace, GM_ADDR tiling)
23+{
24+ SwigluGroupQuantGrad<DTYPE_X> op;
25+ op.Init(gradY, x, weight, yOrigin, groupIndex, gradX, gradWeight, workspace, tiling);
26+ op.Process();
27+}
Aactivation/swiglu_group_quant_grad/op_kernel/swiglu_group_quant_grad.h+548-0
@@ -0,0 +1,548 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file swiglu_group_quant_grad.h
13+ * \brief SwiGLU Group Dynamic Quant Backward kernel for Ascend 950 (A5)
14+ */
15+ 
16+#ifndef SWIGLU_GROUP_QUANT_GRAD_H
17+#define SWIGLU_GROUP_QUANT_GRAD_H
18+ 
19+#include "kernel_operator.h"
20+#include "swiglu_group_quant_grad_base.h"
21+ 
22+namespace SwigluGroupQuantGradOp {
23+using namespace AscendC;
24+ 
25+template <typename T>
26+class SwigluGroupQuantGrad : public SwigluGroupQuantGradBase {
27+public:
28+ __aicore__ inline SwigluGroupQuantGrad() {}
29+
30+ __aicore__ inline void Init(GM_ADDR gradY, GM_ADDR x, GM_ADDR weight, GM_ADDR yOrigin,
31+ GM_ADDR groupIndex, GM_ADDR gradX, GM_ADDR gradWeight,
32+ GM_ADDR workspace, GM_ADDR tiling);
33+
34+ __aicore__ inline void Process();
35+ 
36+private:
37+ __aicore__ inline void CopyInGradY(uint32_t tokenIdx, uint32_t hTileIdx, uint32_t computeSize);
38+ __aicore__ inline void CopyInX(uint32_t tokenIdx, uint32_t hTileIdx, uint32_t actualTokens, uint32_t currentTileH);
39+ __aicore__ inline void CopyInTopkWeight(uint32_t tokenIdx, uint32_t currentTileTokens);
40+ __aicore__ inline void CopyInYOrigin(uint32_t tokenIdx, uint32_t hTileIdx, uint32_t computeSize);
41+
42+ __aicore__ inline void ClampX(LocalTensor<float>& xFloatLocalTensor, uint32_t computeSize);
43+ __aicore__ inline void ComputeSiLUGrad(LocalTensor<float>& xFloatLocalTensor, uint32_t computeSize);
44+ __aicore__ inline void ComputeGradX(LocalTensor<float>& xFloatLocalTensor,
45+ LocalTensor<float>& gradYFloatLocalTensor, uint32_t computeSize);
46+ __aicore__ inline void ApplyClampMask(LocalTensor<float>& xFloatLocalTensor, uint32_t computeSize);
47+ __aicore__ inline void AccumulateGradWeight(LocalTensor<float>& xFloatLocalTensor,
48+ LocalTensor<float>& gradYFloatLocalTensor,
49+ LocalTensor<float>& yOriginFloatLocalTensor,
50+ LocalTensor<float>& weightLocalTensor,
51+ uint32_t currentTileTokens, uint32_t currentTileH, uint32_t hIdx);
52+ __aicore__ inline void UpdateGradY(LocalTensor<float>& gradYFloatLocalTensor,
53+ LocalTensor<float>& weightLocalTensor,
54+ uint32_t currentTileTokens,
55+ uint32_t currentTileH);
56+ __aicore__ inline void CopyOutGradWeight(LocalTensor<float>& weightLocalTensor,
57+ uint32_t tokenIdx,
58+ uint32_t currentTileTokens);
59+ __aicore__ inline void CopyOutGradX(LocalTensor<float>& xFloatLocalTensor,
60+ uint32_t tokenIdx, uint32_t hTileIdx, uint32_t currentTileTokens,
61+ uint32_t currentTileH);
62+ __aicore__ inline void ZeroOutTrunc();
63+ 
64+ __aicore__ inline void ProcessTile(LocalTensor<float>& weightLocalTensor, uint32_t tokenIdx, uint32_t hTileIdx,
65+ uint32_t currentTileTokens, uint32_t currentTileH);
66+
67+ GlobalTensor<T> gradYGm;
68+ GlobalTensor<T> xGm;
69+ GlobalTensor<float> weightGm;
70+ GlobalTensor<T> yOriginGm;
71+ GlobalTensor<int64_t> groupIndexGm;
72+ GlobalTensor<T> gradXGm;
73+ GlobalTensor<float> gradWeightGm;
74+};
75+ 
76+template <typename T>
77+__aicore__ inline void SwigluGroupQuantGrad<T>::Init(GM_ADDR gradY, GM_ADDR x, GM_ADDR weight,
78+ GM_ADDR yOrigin, GM_ADDR groupIndex, GM_ADDR gradX,
79+ GM_ADDR gradWeight, GM_ADDR workspace, GM_ADDR tiling)
80+{
81+ ParseTilingData(tiling);
82+
83+ gradYGm.SetGlobalBuffer((__gm__ T *)gradY, totalTokens * dimH);
84+ xGm.SetGlobalBuffer((__gm__ T *)x, totalTokens * dim2H);
85+ gradXGm.SetGlobalBuffer((__gm__ T *)gradX, totalTokens * dim2H);
86+
87+ if (hasWeight) {
88+ weightGm.SetGlobalBuffer((__gm__ float *)weight, totalTokens);
89+ gradWeightGm.SetGlobalBuffer((__gm__ float *)gradWeight, totalTokens);
90+ }
91+
92+ if (hasYOrigin) {
93+ yOriginGm.SetGlobalBuffer((__gm__ T *)yOrigin, totalTokens * dimH);
94+ }
95+
96+ if (hasGroupIndex) {
97+ groupIndexGm.SetGlobalBuffer((__gm__ int64_t *)groupIndex, groupNum);
98+ }
99+
100+ InitBuffer();
101+}
102+ 
103+template <typename T>
104+__aicore__ inline void SwigluGroupQuantGrad<T>::CopyInGradY(uint32_t tokenIdx, uint32_t hTileIdx, uint32_t computeSize)
105+{
106+ LocalTensor<T> gradYTLocalTensor = gradYQueue.AllocTensor<T>();
107+ uint32_t gmOffset = tokenIdx * dimH + hTileIdx * tileH;
108+ DataCopyParams gradYCopyParams;
109+ gradYCopyParams.blockCount = 1;
110+ gradYCopyParams.blockLen = computeSize * sizeof(T);
111+ gradYCopyParams.srcStride = 0;
112+ gradYCopyParams.dstStride = 0;
113+ DataCopyPadParams padParams{false, 0, 0, 0};
114+ if constexpr (std::is_same_v<T, float>) {
115+ DataCopyPad(gradYTLocalTensor, gradYGm[gmOffset], gradYCopyParams, padParams);
116+ gradYQueue.EnQue<float>(gradYTLocalTensor);
117+ } else {
118+ DataCopyPad(gradYTLocalTensor, gradYGm[gmOffset], gradYCopyParams, padParams);
119+ gradYQueue.EnQue<T>(gradYTLocalTensor);
120+ gradYTLocalTensor = gradYQueue.DeQue<T>();
121+ LocalTensor<float> gradYFloatLocalTensor = gradYTLocalTensor.template ReinterpretCast<float>();
122+ Cast(gradYFloatLocalTensor, gradYTLocalTensor, RoundMode::CAST_NONE, computeSize);
123+ PipeBarrier<PIPE_V>();
124+ gradYQueue.EnQue<float>(gradYFloatLocalTensor);
125+ }
126+}
127+ 
128+template <typename T>
129+__aicore__ inline void SwigluGroupQuantGrad<T>::CopyInX(uint32_t tokenIdx, uint32_t hTileIdx,
130+ uint32_t currentTileTokens, uint32_t currentTileH)
131+{
132+ LocalTensor<T> xTLocalTensor = xQueue.AllocTensor<T>();
133+ uint32_t copySize = currentTileTokens * currentTileH;
134+ DataCopyParams copyParams;
135+ copyParams.blockCount = currentTileTokens;
136+ copyParams.blockLen = currentTileH * sizeof(T);
137+ copyParams.srcStride = (dim2H - currentTileH) * sizeof(T);
138+ copyParams.dstStride = 0;
139+ DataCopyPadParams padParams{false, 0, 0, 0};
140+ if constexpr (std::is_same_v<T, float>) {
141+ uint32_t x0GmOffset = tokenIdx * dim2H + hTileIdx * tileH;
142+ DataCopyPad(xTLocalTensor, xGm[x0GmOffset], copyParams, padParams);
143+ uint32_t x1GmOffset = tokenIdx * dim2H + dimH + hTileIdx * tileH;
144+ DataCopyPad(xTLocalTensor[tileLength], xGm[x1GmOffset], copyParams, padParams);
145+ xQueue.EnQue<float>(xTLocalTensor);
146+ } else {
147+ uint32_t x0GmOffset = tokenIdx * dim2H + hTileIdx * tileH;
148+ DataCopyPad(xTLocalTensor, xGm[x0GmOffset], copyParams, padParams);
149+ uint32_t x1GmOffset = tokenIdx * dim2H + dimH + hTileIdx * tileH;
150+ DataCopyPad(xTLocalTensor[tileLength * sizeof(float) / sizeof(T)], xGm[x1GmOffset], copyParams, padParams);
151+ xQueue.EnQue<T>(xTLocalTensor);
152+ xTLocalTensor = xQueue.DeQue<T>();
153+ LocalTensor<float> xFloatLocalTensor = xTLocalTensor.template ReinterpretCast<float>();
154+ Cast(xFloatLocalTensor, xTLocalTensor, RoundMode::CAST_NONE, copySize);
155+ Cast(xFloatLocalTensor[tileLength], xTLocalTensor[tileLength * sizeof(float) / sizeof(T)],
156+ RoundMode::CAST_NONE, copySize);
157+ PipeBarrier<PIPE_V>();
158+ xQueue.EnQue<float>(xFloatLocalTensor);
159+ }
160+}
161+ 
162+template <typename T>
163+__aicore__ inline void SwigluGroupQuantGrad<T>::ClampX(LocalTensor<float>& xFloatLocalTensor, uint32_t computeSize)
164+{
165+ LocalTensor<float> x0FloatLocalTensor = xFloatLocalTensor;
166+ LocalTensor<float> x1FloatLocalTensor = xFloatLocalTensor[tileLength];
167+ LocalTensor<float> x0TruncatedLocalTensor = xFloatLocalTensor[tileLength * CLAMP_BUFFER_INDEX];
168+ LocalTensor<float> x1TruncatedLocalTensor = xFloatLocalTensor[tileLength * CLAMP_BUFFER_INDEX + tileLength];
169+ Copy(x0TruncatedLocalTensor, x0FloatLocalTensor, computeSize);
170+ PipeBarrier<PIPE_V>();
171+ Copy(x1TruncatedLocalTensor, x1FloatLocalTensor, computeSize);
172+ PipeBarrier<PIPE_V>();
173+ 
174+ Mins(x0FloatLocalTensor, x0FloatLocalTensor, clampLimit, computeSize);
175+ PipeBarrier<PIPE_V>();
176+ Maxs(x1FloatLocalTensor, x1FloatLocalTensor, -clampLimit, computeSize);
177+ PipeBarrier<PIPE_V>();
178+ Mins(x1FloatLocalTensor, x1FloatLocalTensor, clampLimit, computeSize);
179+ PipeBarrier<PIPE_V>();
180+}
181+ 
182+template <typename T>
183+__aicore__ inline void SwigluGroupQuantGrad<T>::CopyInTopkWeight(uint32_t tokenIdx, uint32_t currentTileTokens)
184+{
185+ LocalTensor<float> weightLocalTensor = weightQueue.AllocTensor<float>();
186+ DataCopyParams copyParams;
187+ copyParams.blockCount = 1;
188+ copyParams.blockLen = currentTileTokens * sizeof(float);
189+ copyParams.srcStride = 0;
190+ copyParams.dstStride = 0;
191+ DataCopyPadParams padParams{false, 0, 0, 0};
192+ DataCopyPad(weightLocalTensor, weightGm[tokenIdx], copyParams, padParams);
193+ weightQueue.EnQue<float>(weightLocalTensor);
194+}
195+ 
196+template <typename T>
197+__aicore__ inline void SwigluGroupQuantGrad<T>::CopyInYOrigin(uint32_t tokenIdx, uint32_t hTileIdx,
198+ uint32_t computeSize)
199+{
200+ LocalTensor<T> yOriginTLocalTensor = yOriginQueue.AllocTensor<T>();
201+ uint32_t gmOffset = tokenIdx * dimH + hTileIdx * tileH;
202+ DataCopyParams copyParams;
203+ copyParams.blockCount = 1;
204+ copyParams.blockLen = computeSize * sizeof(T);
205+ copyParams.srcStride = 0;
206+ copyParams.dstStride = 0;
207+ DataCopyPadParams padParams{false, 0, 0, 0};
208+ 
209+ if constexpr (std::is_same_v<T, float>) {
210+ DataCopyPad(yOriginTLocalTensor, yOriginGm[gmOffset], copyParams, padParams);
211+ yOriginQueue.EnQue<float>(yOriginTLocalTensor);
212+ } else {
213+ DataCopyPad(yOriginTLocalTensor, yOriginGm[gmOffset], copyParams, padParams);
214+ yOriginQueue.EnQue<T>(yOriginTLocalTensor);
215+ yOriginTLocalTensor = yOriginQueue.DeQue<T>();
216+ LocalTensor<float> yOriginFloatLocalTensor = yOriginTLocalTensor.template ReinterpretCast<float>();
217+ Cast(yOriginFloatLocalTensor, yOriginTLocalTensor, RoundMode::CAST_NONE, computeSize);
218+ PipeBarrier<PIPE_V>();
219+ yOriginQueue.EnQue<float>(yOriginFloatLocalTensor);
220+ }
221+}
222+ 
223+template <typename T>
224+__aicore__ inline void SwigluGroupQuantGrad<T>::ComputeSiLUGrad(LocalTensor<float>& xFloatLocalTensor,
225+ uint32_t computeSize)
S
Ssunday6月18日

ComputeSiLUGrad中使用 exp(x)/(1+exp(x)) 计算sigmoid，当x0>88时exp溢出产生NaN。仓库同类算子swiglu_mx_quant使用 1/(1+exp(-x)) 稳定公式，建议改为一致。具体做法：先Neg取负，再Exp，再Adds(1.0)，最后Div(1.0, ...)，即sigmoid=1/(1+exp(-x))。

likedislike
S
shilulu
6月18日 评论：
226+{
227+ LocalTensor<float> x0FloatLocalTensor = xFloatLocalTensor;
228+ LocalTensor<float> x1FloatLocalTensor = xFloatLocalTensor[tileLength];
229+ 
230+ LocalTensor<float> sigmoidX0LocalTensor = xFloatLocalTensor[tileLength * SILU_GRAD_BUFFER_INDEX];
231+ LocalTensor<float> siluX0LocalTensor = xFloatLocalTensor[tileLength * SILU_GRAD_BUFFER_INDEX + tileLength];
232+ LocalTensor<float> siluGradX0LocalTensor = xFloatLocalTensor[tileLength * SILU_GRAD_BUFFER_INDEX + 2 * tileLength];
233+ LocalTensor<float> tmpLocalTensor = xFloatLocalTensor[tileLength * TMP_BUFFER_INDEX];
234+ 
235+ Exp(tmpLocalTensor, x0FloatLocalTensor, computeSize);
236+ PipeBarrier<PIPE_V>();
237+ Copy(sigmoidX0LocalTensor, tmpLocalTensor, computeSize);
238+ PipeBarrier<PIPE_V>();
239+ Adds(tmpLocalTensor, tmpLocalTensor, 1.0f, computeSize);
240+ PipeBarrier<PIPE_V>();
241+ Div(sigmoidX0LocalTensor, sigmoidX0LocalTensor, tmpLocalTensor, computeSize);
242+ PipeBarrier<PIPE_V>();
243+
244+ Mul(siluX0LocalTensor, x0FloatLocalTensor, sigmoidX0LocalTensor, computeSize);
245+ PipeBarrier<PIPE_V>();
246+
247+ Subs(tmpLocalTensor, (float)1.0, sigmoidX0LocalTensor, computeSize);
248+ PipeBarrier<PIPE_V>();
249+ Mul(tmpLocalTensor, tmpLocalTensor, x0FloatLocalTensor, computeSize);
250+ PipeBarrier<PIPE_V>();
251+ Adds(tmpLocalTensor, tmpLocalTensor, 1.0f, computeSize);
252+ PipeBarrier<PIPE_V>();
253+ Mul(siluGradX0LocalTensor, sigmoidX0LocalTensor, tmpLocalTensor, computeSize);
254+ PipeBarrier<PIPE_V>();
255+}
256+ 
257+template <typename T>
258+__aicore__ inline void SwigluGroupQuantGrad<T>::ComputeGradX(LocalTensor<float>& xFloatLocalTensor,
259+ LocalTensor<float>& gradYFloatLocalTensor,
260+ uint32_t computeSize)
261+{
262+ LocalTensor<float> x0FloatLocalTensor = xFloatLocalTensor;
263+ LocalTensor<float> x1FloatLocalTensor = xFloatLocalTensor[tileLength];
264+ LocalTensor<float> siluX0LocalTensor = xFloatLocalTensor[tileLength * SILU_GRAD_BUFFER_INDEX + tileLength];
265+ LocalTensor<float> siluGradX0LocalTensor = xFloatLocalTensor[tileLength * SILU_GRAD_BUFFER_INDEX + 2 * tileLength];
266+ 
267+ // update x0 to x0Grad
268+ Mul(x0FloatLocalTensor, gradYFloatLocalTensor, x1FloatLocalTensor, computeSize);
269+ PipeBarrier<PIPE_V>();
270+ Mul(x0FloatLocalTensor, x0FloatLocalTensor, siluGradX0LocalTensor, computeSize);
271+ PipeBarrier<PIPE_V>();
272+ // update x1 to x1Grad
273+ Mul(x1FloatLocalTensor, gradYFloatLocalTensor, siluX0LocalTensor, computeSize);
274+ PipeBarrier<PIPE_V>();
275+}
276+ 
277+template <typename T>
278+__aicore__ inline void SwigluGroupQuantGrad<T>::ApplyClampMask(LocalTensor<float>& xFloatLocalTensor,
279+ uint32_t computeSize)
280+{
281+ LocalTensor<float> tmpLocalTensor = xFloatLocalTensor[tileLength * TMP_BUFFER_INDEX];
282+ LocalTensor<uint8_t> maskX0U8Local = tmpLocalTensor.template ReinterpretCast<uint8_t>();
283+ LocalTensor<uint8_t> maskX1LeftU8Local = tmpLocalTensor[tileLength].template ReinterpretCast<uint8_t>();
284+ LocalTensor<uint8_t> maskX1RightU8Local = tmpLocalTensor[tileLength * 2].template ReinterpretCast<uint8_t>();
285+ LocalTensor<uint8_t> maskX1U8Local = maskX1LeftU8Local;
286+ 
287+ LocalTensor<float> x0FloatLocalTensor = xFloatLocalTensor;
288+ LocalTensor<float> x1FloatLocalTensor = xFloatLocalTensor[tileLength];
289+ LocalTensor<float> x0TruncatedLocalTensor = xFloatLocalTensor[tileLength * CLAMP_BUFFER_INDEX];
290+ LocalTensor<float> x1TruncatedLocalTensor = xFloatLocalTensor[tileLength * CLAMP_BUFFER_INDEX + tileLength];
291+ 
292+ // reuse x0TruncatedLocalTensor/x1TruncatedLocalTensor
293+ LocalTensor<float> maskX0Local = x0TruncatedLocalTensor;
294+ LocalTensor<float> maskX1Local = x1TruncatedLocalTensor;
295+
296+ CompareScalar(maskX0U8Local, x0TruncatedLocalTensor, clampLimit, CMPMODE::LT, computeSize);
297+ PipeBarrier<PIPE_V>();
298+ CompareScalar(maskX1LeftU8Local, x1TruncatedLocalTensor, -clampLimit, CMPMODE::GT, computeSize);
299+ PipeBarrier<PIPE_V>();
300+ CompareScalar(maskX1RightU8Local, x1TruncatedLocalTensor, clampLimit, CMPMODE::LT, computeSize);
301+ PipeBarrier<PIPE_V>();
302+ And(maskX1U8Local, maskX1LeftU8Local, maskX1RightU8Local, computeSize);
303+ PipeBarrier<PIPE_V>();
304+ 
305+ LocalTensor<float> onesTensor = tmpLocalTensor[tileLength * 2];
306+ Duplicate(onesTensor, (float)1.0, computeSize);
307+ PipeBarrier<PIPE_V>();
308+ 
309+ Select(maskX0Local, maskX0U8Local, onesTensor, static_cast<float>(0), SELMODE::VSEL_TENSOR_SCALAR_MODE,
310+ computeSize);
311+ PipeBarrier<PIPE_V>();
312+ Select(maskX1Local, maskX1U8Local, onesTensor, static_cast<float>(0), SELMODE::VSEL_TENSOR_SCALAR_MODE,
313+ computeSize);
314+ PipeBarrier<PIPE_V>();
315+ 
316+ Mul(x0FloatLocalTensor, x0FloatLocalTensor, maskX0Local, computeSize);
317+ PipeBarrier<PIPE_V>();
318+ Mul(x1FloatLocalTensor, x1FloatLocalTensor, maskX1Local, computeSize);
319+ PipeBarrier<PIPE_V>();
320+}
321+ 
322+template <typename T>
323+__aicore__ inline void SwigluGroupQuantGrad<T>::AccumulateGradWeight(LocalTensor<float>& xFloatLocalTensor,
324+ LocalTensor<float>& gradYFloatLocalTensor, LocalTensor<float>& yOriginFloatLocalTensor,
325+ LocalTensor<float>& weightLocalTensor, uint32_t currentTileTokens, uint32_t currentTileH, uint32_t hIdx)
326+{
327+ LocalTensor<float> gradWeightAccumLocalTensor = weightLocalTensor[AlignUp(tileTokens, FP32_32B_ALIGN_NUM)];
328+ LocalTensor<float> tmpLocalTensor = xFloatLocalTensor[tileLength * TMP_BUFFER_INDEX];
329+ uint32_t computeSize = currentTileTokens * currentTileH;
330+ Mul(yOriginFloatLocalTensor, yOriginFloatLocalTensor, gradYFloatLocalTensor, computeSize);
331+ PipeBarrier<PIPE_V>();
332+
333+ if (hIdx == 0) {
334+ Copy(gradWeightAccumLocalTensor, yOriginFloatLocalTensor, computeSize);
335+ PipeBarrier<PIPE_V>();
336+ } else {
337+ Add(gradWeightAccumLocalTensor, gradWeightAccumLocalTensor, yOriginFloatLocalTensor, computeSize);
338+ PipeBarrier<PIPE_V>();
339+ }
340+ 
341+ if (hIdx == numHTiles - 1) {
342+ for (uint32_t t = 0; t < currentTileTokens; t++) {
343+ ReduceSum<float>(gradWeightAccumLocalTensor[t * tileH], gradWeightAccumLocalTensor[t * tileH],
344+ tmpLocalTensor, tileH);
345+ PipeBarrier<PIPE_V>();
346+ }
347+ }
348+}
349+ 
350+template <typename T>
351+__aicore__ inline void SwigluGroupQuantGrad<T>::UpdateGradY(LocalTensor<float>& gradYFloatLocalTensor,
352+ LocalTensor<float>& weightLocalTensor,
353+ uint32_t currentTileTokens, uint32_t currentTileH)
354+{
355+ for (uint32_t t = 0; t < currentTileTokens; t++) {
356+ float weightVal = weightLocalTensor.GetValue(t);
S
Ssunday6月18日

UpdateGradY逐token标量读取weight值+Muls循环，每token一次vector指令调用开销大。建议改为向量级broadcast模式：先Duplicate weight到UB buffer，再用Mul整体完成gradY*weight。

likedislike
S
shilulu
6月18日 评论：
357+ Muls(gradYFloatLocalTensor[t * currentTileH], gradYFloatLocalTensor[t * currentTileH], weightVal, currentTileH);
358+ PipeBarrier<PIPE_V>();
359+ }
360+}
361+ 
362+template <typename T>
363+__aicore__ inline void SwigluGroupQuantGrad<T>::CopyOutGradWeight(LocalTensor<float>& weightLocalTensor,
364+ uint32_t tokenIdx, uint32_t currentTileTokens)
365+{
366+ event_t vToMte3 = static_cast<event_t>(GetTPipePtr()->AllocEventID<HardEvent::V_MTE3>());
367+ SetFlag<HardEvent::V_MTE3>(vToMte3);
368+ WaitFlag<HardEvent::V_MTE3>(vToMte3);
369+ LocalTensor<float> gradWeightAccumLocalTensor = weightLocalTensor[AlignUp(tileTokens, FP32_32B_ALIGN_NUM)];
370+ DataCopyParams copyParams;
371+ copyParams.blockCount = 1;
372+ copyParams.blockLen = 1 * sizeof(float);
373+ copyParams.srcStride = 0;
S
Ssunday6月18日

CopyOutGradWeight per-token循环调用DataCopyPad，每次仅拷4字节(float)，DMA效率极低。建议改为批量一次搬运，或至少合并到一个连续buffer后一次DMA写出。

likedislike
S
shilulu
6月18日 评论：
374+ copyParams.dstStride = 0;
375+ for (uint32_t t = 0; t < currentTileTokens; t++) {
376+ DataCopyPad(gradWeightGm[tokenIdx + t], gradWeightAccumLocalTensor[t * tileH], copyParams);
377+ }
378+ GetTPipePtr()->ReleaseEventID<AscendC::HardEvent::V_MTE3>(vToMte3);
379+}
380+ 
381+template <typename T>
382+__aicore__ inline void SwigluGroupQuantGrad<T>::ZeroOutTrunc()
383+{
384+ if (truncValue >= totalTokens) {
385+ return;
386+ }
387+
388+ uint32_t zeroTokenStart = truncValue + blockIdx;
389+ uint32_t zeroTokenStep = usedCoreNum;
390+ 
391+ LocalTensor<T> zeroXLocal = zeroQueue.AllocTensor<T>();
392+ LocalTensor<float> zeroXFloatLocal;
393+ LocalTensor<float> zeroWeightLocal;
394+ if constexpr (std::is_same_v<T, float>) {
395+ zeroWeightLocal = zeroXLocal[dim2H];
396+ } else {
397+ zeroXFloatLocal = zeroXLocal.template ReinterpretCast<float>();
398+ zeroWeightLocal = zeroXFloatLocal[dim2H];
399+ }
400+ Duplicate(zeroXLocal, (T)0, dim2H);
401+ Duplicate(zeroWeightLocal, (float)0.0, 1);
402+ event_t vToMte3 = static_cast<event_t>(GetTPipePtr()->AllocEventID<HardEvent::V_MTE3>());
403+ SetFlag<HardEvent::V_MTE3>(vToMte3);
404+ WaitFlag<HardEvent::V_MTE3>(vToMte3);
405+ for (uint32_t t = zeroTokenStart; t < totalTokens; t += zeroTokenStep) {
406+ DataCopyParams gradXCopyParams;
407+ gradXCopyParams.blockCount = 1;
408+ gradXCopyParams.blockLen = dim2H * sizeof(T);
409+ gradXCopyParams.srcStride = 0;
410+ gradXCopyParams.dstStride = 0;
411+ DataCopyPad(gradXGm[t * dim2H], zeroXLocal, gradXCopyParams);
412+ if (hasWeight) {
413+ DataCopyParams gradWightCopyParams{1, 1 * sizeof(float), 0, 0};
414+ DataCopyPad(gradWeightGm[t], zeroWeightLocal, gradWightCopyParams);
415+ }
416+ }
417+ zeroQueue.FreeTensor<T>(zeroXLocal);
418+ GetTPipePtr()->ReleaseEventID<AscendC::HardEvent::V_MTE3>(vToMte3);
419+}
420+ 
421+template <typename T>
422+__aicore__ inline void SwigluGroupQuantGrad<T>::CopyOutGradX(LocalTensor<float>& xFloatLocalTensor, uint32_t tokenIdx,
423+ uint32_t hTileIdx, uint32_t currentTileTokens,
424+ uint32_t currentTileH)
425+{
426+ uint32_t gmOffset0 = tokenIdx * dim2H + hTileIdx * tileH;
427+ uint32_t gmOffset1 = tokenIdx * dim2H + dimH + hTileIdx * tileH;
428+ uint32_t currentTileLength = currentTileTokens * currentTileH;
429+ DataCopyParams outCopyParams;
430+ outCopyParams.blockCount = currentTileTokens;
431+ outCopyParams.blockLen = currentTileH * sizeof(T);
432+ outCopyParams.srcStride = 0;
433+ outCopyParams.dstStride = (dim2H - currentTileH) * sizeof(T);
434+ 
435+ event_t vToMte3 = static_cast<event_t>(GetTPipePtr()->AllocEventID<HardEvent::V_MTE3>());
436+ SetFlag<HardEvent::V_MTE3>(vToMte3);
437+ WaitFlag<HardEvent::V_MTE3>(vToMte3);
438+ 
439+ if constexpr (std::is_same_v<T, float>) {
440+ LocalTensor<float> x0FloatLocalTensor = xFloatLocalTensor;
441+ LocalTensor<float> x1FloatLocalTensor = xFloatLocalTensor[tileLength];
442+ DataCopyPad(gradXGm[gmOffset0], x0FloatLocalTensor, outCopyParams);
443+ DataCopyPad(gradXGm[gmOffset1], x1FloatLocalTensor, outCopyParams);
444+ } else {
445+ LocalTensor<float> x0FloatLocalTensor = xFloatLocalTensor;
446+ LocalTensor<float> x1FloatLocalTensor = xFloatLocalTensor[tileLength];
447+ LocalTensor<T> x0TLocalTensor = x0FloatLocalTensor.template ReinterpretCast<T>();
448+ Cast(x0TLocalTensor, x0FloatLocalTensor, RoundMode::CAST_RINT, currentTileLength);
449+ PipeBarrier<PIPE_V>();
450+ LocalTensor<T> x1TLocalTensor = x1FloatLocalTensor.template ReinterpretCast<T>();
451+ Cast(x1TLocalTensor, x1FloatLocalTensor, RoundMode::CAST_RINT, currentTileLength);
452+ PipeBarrier<PIPE_V>();
453+ SetFlag<HardEvent::V_MTE3>(vToMte3);
454+ WaitFlag<HardEvent::V_MTE3>(vToMte3);
455+ DataCopyPad(gradXGm[gmOffset0], x0TLocalTensor, outCopyParams);
456+ DataCopyPad(gradXGm[gmOffset1], x1TLocalTensor, outCopyParams);
457+ }
458+ GetTPipePtr()->ReleaseEventID<AscendC::HardEvent::V_MTE3>(vToMte3);
459+}
460+ 
461+template <typename T>
462+__aicore__ inline void SwigluGroupQuantGrad<T>::ProcessTile(LocalTensor<float>& weightLocalTensor, uint32_t tokenIdx,
463+ uint32_t hTileIdx, uint32_t currentTileTokens,
464+ uint32_t currentTileH)
465+{
466+ uint32_t computeSize = currentTileTokens * currentTileH;
467+ uint32_t copySize = currentTileTokens * tileH;
468+ // 1.copy in grad_y, x
469+ CopyInGradY(tokenIdx, hTileIdx, computeSize);
470+ CopyInX(tokenIdx, hTileIdx, currentTileTokens, currentTileH);
471+ LocalTensor<float> gradYFloatLocalTensor = gradYQueue.DeQue<float>();
472+ LocalTensor<float> xFloatLocalTensor = xQueue.DeQue<float>();
473+ 
474+ // 2.update grad_y with weight and compute grad_weight
475+ if (hasWeight) {
476+ CopyInYOrigin(tokenIdx, hTileIdx, computeSize);
477+ LocalTensor<float> yOriginFloatLocalTensor = yOriginQueue.DeQue<float>();
478+ AccumulateGradWeight(xFloatLocalTensor, gradYFloatLocalTensor, yOriginFloatLocalTensor, weightLocalTensor,
479+ currentTileTokens, currentTileH, hTileIdx);
480+ yOriginQueue.FreeTensor<float>(yOriginFloatLocalTensor);
481+ UpdateGradY(gradYFloatLocalTensor, weightLocalTensor, currentTileTokens, currentTileH);
482+ }
483+ 
484+ // 3.clamp x0 x1
485+ if (hasClampLimit) {
486+ ClampX(xFloatLocalTensor, computeSize);
487+ }
488+ 
489+ // 4.compute silugrad
490+ ComputeSiLUGrad(xFloatLocalTensor, computeSize);
491+ // 5.compute grad_x
492+ ComputeGradX(xFloatLocalTensor, gradYFloatLocalTensor, computeSize);
493+ // 6.grad_x clamp mask
494+ if (hasClampLimit) {
495+ ApplyClampMask(xFloatLocalTensor, computeSize);
496+ }
497+ 
498+ // 7.copy out grad_x
499+ CopyOutGradX(xFloatLocalTensor, tokenIdx, hTileIdx, currentTileTokens, currentTileH);
500+ gradYQueue.FreeTensor<float>(gradYFloatLocalTensor);
501+ xQueue.FreeTensor<float>(xFloatLocalTensor);
502+}
503+ 
504+template <typename T>
505+__aicore__ inline void SwigluGroupQuantGrad<T>::Process()
506+{
507+ if (blockIdx >= usedCoreNum) {
508+ return;
509+ }
510+
511+ uint32_t tokenEnd = tokenStart + tokensPerCore;
512+ if (tokenEnd > truncValue) {
513+ tokenEnd = truncValue;
514+ }
515+ uint32_t tokenIdx = tokenStart;
516+ while (tokenIdx < tokenEnd) {
517+ uint32_t currentTileTokens = tileTokens;
518+ if (tokenIdx + tileTokens > tokenEnd) {
519+ currentTileTokens = tokenEnd - tokenIdx;
520+ }
521+ 
522+ LocalTensor<float> weightLocalTensor;
523+ if (hasWeight) {
524+ CopyInTopkWeight(tokenIdx, currentTileTokens);
525+ weightLocalTensor = weightQueue.DeQue<float>();
526+ }
527+ 
528+ for (uint32_t hTileIdx = 0; hTileIdx < numHTiles; hTileIdx++) {
529+ uint32_t currentTileH = tileH;
530+ if (hTileIdx == numHTiles - 1) {
531+ currentTileH = dimH - hTileIdx * tileH;
532+ }
533+ ProcessTile(weightLocalTensor, tokenIdx, hTileIdx, currentTileTokens, currentTileH);
534+ }
535+
536+ if (hasWeight) {
537+ CopyOutGradWeight(weightLocalTensor, tokenIdx, currentTileTokens);
538+ weightQueue.FreeTensor<float>(weightLocalTensor);
539+ }
540+ tokenIdx += currentTileTokens;
541+ }
542+ SyncAll();
543+ ZeroOutTrunc();
544+}
545+ 
546+} // namespace SwigluGroupQuantGradOp
547+ 
548+#endif // SWIGLU_GROUP_QUANT_GRAD_H
Aactivation/swiglu_group_quant_grad/op_kernel/swiglu_group_quant_grad_base.h+138-0
@@ -0,0 +1,138 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file swiglu_group_quant_grad_base.h
13+ * \brief SwiGLU Group Dynamic Quant Backward base class
14+ */
15+ 
16+#ifndef SWIGLU_GROUP_QUANT_GRAD_BASE_H
17+#define SWIGLU_GROUP_QUANT_GRAD_BASE_H
18+ 
19+#include "kernel_operator.h"
20+ 
21+namespace SwigluGroupQuantGradOp {
22+using namespace AscendC;
23+ 
24+constexpr uint32_t BUFFER_NUM = 2;
25+constexpr uint32_t SPLIT_NUM = 2;
26+constexpr uint32_t BLOCK_SIZE = 32;
27+constexpr uint32_t FP32_32B_ALIGN_NUM = 8;
28+constexpr uint32_t CLAMP_BUFFER_INDEX = 8;
29+constexpr uint32_t SILU_GRAD_BUFFER_INDEX = 2;
30+constexpr uint32_t TMP_BUFFER_INDEX = 5;
31+constexpr uint32_t GRAD_WEIGHT_ACCUM_BUFFER_INDEX = 1;
32+constexpr uint32_t HAS_CLAMP_SCENE_TILE_NUM = 10;
33+constexpr uint32_t BASE_SCENE_TILE_NUM = 6;
34+ 
35+class SwigluGroupQuantGradBase {
36+public:
37+ __aicore__ inline SwigluGroupQuantGradBase() {}
38+ 
39+ __aicore__ inline void ParseTilingData(GM_ADDR tiling)
40+ {
41+ GET_TILING_DATA_WITH_STRUCT(SwigluGroupQuantGradTilingData, tilingData, tiling);
42+
43+ blockIdx = GetBlockIdx();
44+ totalTokens = tilingData.totalTokens;
45+ dimH = tilingData.dimH;
46+ dim2H = tilingData.dim2H;
47+ groupNum = tilingData.groupNum;
48+ truncValue = tilingData.truncValue;
49+ tileTokens = tilingData.tileTokens;
50+ tileH = tilingData.tileH;
51+ numHTiles = tilingData.numHTiles;
52+ usedCoreNum = tilingData.usedCoreNum;
53+ tokensPerCore = tilingData.tokensPerCore;
54+ hasWeight = tilingData.hasWeight;
55+ hasYOrigin = tilingData.hasYOrigin;
56+ hasGroupIndex = tilingData.hasGroupIndex;
57+ hasClampLimit = tilingData.hasClampLimit;
58+ clampLimit = tilingData.clampLimit;
59+
60+ tokenStart = blockIdx * tokensPerCore;
61+
62+ tileLength = tileTokens * tileH;
63+ tileDataSize = tileLength * sizeof(float);
64+ }
65+ 
66+ __aicore__ inline void InitBuffer()
67+ {
68+ pipe.InitBuffer(gradYQueue, BUFFER_NUM, tileLength * sizeof(float));
69+ if (hasClampLimit) {
70+ pipe.InitBuffer(xQueue, BUFFER_NUM, tileLength * HAS_CLAMP_SCENE_TILE_NUM * sizeof(float));
71+ } else {
72+ pipe.InitBuffer(xQueue, BUFFER_NUM, tileLength * BASE_SCENE_TILE_NUM * sizeof(float));
73+ }
74+
75+ if (hasWeight) {
76+ pipe.InitBuffer(weightQueue, BUFFER_NUM, AlignUp(tileTokens, FP32_32B_ALIGN_NUM) * sizeof(float) +
77+ tileLength * sizeof(float));
78+ pipe.InitBuffer(yOriginQueue, BUFFER_NUM, tileLength * sizeof(float));
79+ }
80+ 
81+ if (truncValue < totalTokens) {
82+ uint32_t zeroOutBufSize = (AlignUp(dim2H, FP32_32B_ALIGN_NUM) + FP32_32B_ALIGN_NUM) * sizeof(float);
83+ pipe.InitBuffer(zeroQueue, BUFFER_NUM, zeroOutBufSize);
84+ }
85+ }
86+ 
87+ template <typename T>
88+ __aicore__ inline T CeilDiv(T x, T y)
89+ {
90+ return y == 0 ? 0 : (x + y - 1) / y;
91+ }
92+ 
93+ template <typename T>
94+ __aicore__ inline T AlignUp(T num, T div)
95+ {
96+ return (div == 0) ? 0 : (num + div - 1) / div * div;
97+ }
98+ 
99+ template <typename T>
100+ __aicore__ inline T AlignDown(T num, T div)
101+ {
102+ return (div == 0) ? 0 : num / div * div;
103+ }
104+ 
105+protected:
106+ TPipe pipe;
107+
108+ TQue<TPosition::VECIN, BUFFER_NUM> gradYQueue;
109+ TQue<TPosition::VECIN, BUFFER_NUM> xQueue;
110+ TQue<TPosition::VECIN, BUFFER_NUM> weightQueue;
111+ TQue<TPosition::VECIN, BUFFER_NUM> yOriginQueue;
112+ TQue<TPosition::VECIN, BUFFER_NUM> zeroQueue;
113+
114+ uint32_t blockIdx = 0;
115+ uint32_t totalTokens = 0;
116+ uint32_t dimH = 0;
117+ uint32_t dim2H = 0;
118+ uint32_t groupNum = 0;
119+ uint32_t truncValue = 0;
120+ uint32_t tileTokens = 0;
121+ uint32_t tileH = 0;
122+ uint32_t numHTiles = 0;
123+ uint32_t usedCoreNum = 0;
124+ uint32_t tokensPerCore = 0;
125+ uint32_t tokenStart = 0;
126+ uint32_t hasWeight = 0;
127+ uint32_t hasYOrigin = 0;
128+ uint32_t hasGroupIndex = 0;
129+ uint32_t hasClampLimit = 0;
130+ float clampLimit = 0.0f;
131+
132+ uint32_t tileLength = 0;
133+ uint32_t tileDataSize = 0;
134+};
135+ 
136+} // namespace SwigluGroupQuantGradOp
137+ 
138+#endif // SWIGLU_GROUP_QUANT_GRAD_BASE_H
Mdocs/zh/op_list.md+10-0