Star917
1836
代码
代码
Issues222
Pull Requests329
Actions
项目讨论
Wiki
分析
Star917
1836
已合并
Added clipped_swiglu op with clamp_mode #8534
clwsy创建于 27 天前
Added clipped_swiglu op with clamp_mode #8534
已合并
clwsy创建于 27 天前
从clwsy1/ops-nndev合入到cann/ops-nnmaster
共 24 个文件变更+1448-118
Mactivation/clipped_swiglu/README.md+39-6
@@ -13,7 +13,7 @@
13 13 
14## 功能说明14## 功能说明
15 15 
16-- 接口功能：带截断的Swish门控线性单元激活函数，实现x的SwiGlu计算。本算子相较于SwiGlu算子，新增了部分输入参数：groupIndex、alpha、limit、bias、interleaved，用于支持GPT-OSS模型使用的变体SwiGlu以及MoE模型使用的分组场景。16+- 接口功能：带截断的Swish门控线性单元激活函数，实现x的SwiGlu计算。本算子相较于SwiGlu算子，新增了部分输入参数：groupIndex、alpha、limit、bias、interleaved，clamp_mode，用于支持GPT-OSS模型使用的变体SwiGlu、MoE模型使用的分组场景以及部分新模型需要将clamp操作后移的场景。
17 17 
18- 计算流程：18- 计算流程：
19 19 
@@ -58,7 +58,9 @@
58 $$58 $$
59 B = x[ : , h : ]59 B = x[ : , h : ]
60 $$60 $$
61- 4. 根据输入参数alpha、limit、bias进行变体SwiGlu计算，公式如下：61+ 4. 根据输入参数alpha、limit、bias、clamp_mode进行变体SwiGlu计算，公式如下：
62+ 
63+ 当clamp_mode为0时，表示clamp操作在silu之前：
62 64 
63 $$65 $$
64 A = A.clamp(min=None, max=limit)66 A = A.clamp(min=None, max=limit)
@@ -75,6 +77,25 @@
75 $$77 $$
76 y = y\_glu * (B + bias)78 y = y\_glu * (B + bias)
77 $$79 $$
80+ 
81+ 当clamp_mode为1时，表示clamp操作在silu之后：
82+ 
83+ $$
84+ y\_glu = A * sigmoid(A)
85+ $$
86+ 
87+ $$
88+ y\_glu = y\_glu.clamp(min=None, max=limit)
89+ $$
90+ 
91+ $$
92+ B = B.clamp(min=-limit, max=limit)
93+ $$
94+ 
95+ $$
96+ y = y\_glu * B
97+ $$
98+ 
78 5. 重塑输出张量y的维度数量与合轴前的x的维度数量一致，dim轴上的大小为x的一半，其他维度与x相同。99 5. 重塑输出张量y的维度数量与合轴前的x的维度数量一致，dim轴上的大小为x的一半，其他维度与x相同。
79 100 
80## 参数说明101## 参数说明
@@ -98,16 +119,16 @@
98 <tr>119 <tr>
99 <td>x</td>120 <td>x</td>
100 <td>输入</td>121 <td>输入</td>
101- <td>不支持空Tensor。公式中的输入x。维度必须大于0且必须在入参dim对应维度上是偶数。输入不支持包含±inf或nan。</td>122+ <td>公式中的输入x。维度必须大于0且必须在入参dim对应维度上是偶数。</td>
102 <td>FLOAT、FLOAT16、BFLOAT16</td>123 <td>FLOAT、FLOAT16、BFLOAT16</td>
103 <td>ND</td>124 <td>ND</td>
104 </tr>125 </tr>
105 <tr>126 <tr>
106 <td>group_index</td>127 <td>group_index</td>
107 <td>可选输入</td>128 <td>可选输入</td>
108- <td>不支持空Tensor。公式中的输入group_index。维度必须是1维，且元素个数必须小于等于8192。输入不支持包含±inf或nan。</td>129+ <td>公式中的输入group_index。维度必须是1维，且元素个数必须小于等于8192。</td>
109 <td>INT64</td>130 <td>INT64</td>
110- <td>-</td>131+ <td>ND</td>
111 </tr>132 </tr>
112 <tr>133 <tr>
113 <td>dim</td>134 <td>dim</td>
@@ -144,15 +165,25 @@
144 <td>BOOL</td>165 <td>BOOL</td>
145 <td>-</td>166 <td>-</td>
146 </tr>167 </tr>
168+ <tr>
169+ <td>clamp_mode</td>
170+ <td>可选属性</td>
171+ <td>公式中的输入clamp_mode，设置为1表示将A的clamp操作后移至silu激活之后，即对silu(A)的结果做clamp；设置为0表示A的clamp操作在silu激活之前。默认为0。</td>
172+ <td>INT64</td>
173+ <td>-</td>
174+ </tr>
147 <tr>175 <tr>
148 <td>y</td>176 <td>y</td>
149 <td>输出</td>177 <td>输出</td>
150- <td>不支持空Tensor。公式中的输出y。维度数需与输入x一致，在入参dim对应维度上为x的一半，其他维度上与x一致。</td>178+ <td>公式中的输出y。维度数需与输入x一致，在入参dim对应维度上为x的一半，其他维度上与x一致。</td>
151 <td>FLOAT、FLOAT16、BFLOAT16</td>179 <td>FLOAT、FLOAT16、BFLOAT16</td>
152 <td>ND</td>180 <td>ND</td>
153 </tr>181 </tr>
154 </tbody></table>182 </tbody></table>
155 183 
184+- <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品/Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：
185+ - 属性`clamp_mode`为无效值，计算逻辑同clamp_mode=0。
186+ 
156## 约束说明187## 约束说明
157 188 
158无。189无。
@@ -162,3 +193,5 @@
162| 调用方式 | 调用样例 | 说明 |193| 调用方式 | 调用样例 | 说明 |
163|--------------|------------------------------------------------------------------------|----------------------------------------------------------------|194|--------------|------------------------------------------------------------------------|----------------------------------------------------------------|
164| aclnn调用 | [test_aclnn_clipped_swiglu](./examples/test_aclnn_clipped_swiglu.cpp) | 通过[aclnnClippedSwiglu](./docs/aclnnClippedSwiglu.md)接口方式调用ClippedSwiglu算子。 |195| aclnn调用 | [test_aclnn_clipped_swiglu](./examples/test_aclnn_clipped_swiglu.cpp) | 通过[aclnnClippedSwiglu](./docs/aclnnClippedSwiglu.md)接口方式调用ClippedSwiglu算子。 |
196+| aclnn调用 | [test_aclnn_clipped_swiglu_v2](./examples/test_aclnn_clipped_swiglu_v2.cpp) | 通过[aclnnClippedSwigluV2](./docs/aclnnClippedSwigluV2.md)接口方式调用ClippedSwiglu算子。 |
197+| 图模式调用 | - | 通过[算子IR](./op_graph/clipped_swiglu_proto.h)构图方式调用ClippedSwiglu算子。 |
Mactivation/clipped_swiglu/docs/aclnnClippedSwiglu.md+12-3
@@ -144,7 +144,7 @@ aclnnStatus aclnnClippedSwiglu(
144 <td>x（aclTensor*）</td>144 <td>x（aclTensor*）</td>
145 <td>输入</td>145 <td>输入</td>
146 <td>公式中的输入x。</td>146 <td>公式中的输入x。</td>
147- <td>不支持空指针，维度必须大于0且shape必须在入参dim对应维度上是偶数。不支持空Tensor。输入不支持包含±inf或nan。</td>147+ <td>不支持空指针，维度必须大于0且shape必须在入参dim对应维度上是偶数。支持空Tensor。</td>
148 <td>FLOAT、FLOAT16、BFLOAT16</td>148 <td>FLOAT、FLOAT16、BFLOAT16</td>
149 <td>ND</td>149 <td>ND</td>
150 <td>1-8</td>150 <td>1-8</td>
@@ -154,7 +154,7 @@ aclnnStatus aclnnClippedSwiglu(
154 <td>groupIndexOptional（aclTensor*）</td>154 <td>groupIndexOptional（aclTensor*）</td>
155 <td>输入</td>155 <td>输入</td>
156 <td>公式中的输入group_index，表示分组的情况。</td>156 <td>公式中的输入group_index，表示分组的情况。</td>
157- <td>支持空指针。不为空指针时，维度要求为1维，长度不超过8192，且元素需大于等于0。第i个元素代表第i组需要处理x的batch数量。不支持空Tensor。输入不支持包含±inf或nan。</td>157+ <td>支持空指针。不为空指针时，维度要求为1维，长度不超过8192，且元素需大于等于0。第i个元素代表第i组需要处理x的batch数量。</td>
158 <td>INT64</td>158 <td>INT64</td>
159 <td>ND</td>159 <td>ND</td>
160 <td>1</td>160 <td>1</td>
@@ -214,7 +214,7 @@ aclnnStatus aclnnClippedSwiglu(
214 <td>out（aclTensor*）</td>214 <td>out（aclTensor*）</td>
215 <td>输出</td>215 <td>输出</td>
216 <td>公式中的输出y。</td>216 <td>公式中的输出y。</td>
217- <td>不支持空指针。shape在入参dim对应的维度上为x的一半，其他维度上与x一致。不支持空Tensor。</td>217+ <td>不支持空指针。shape在入参dim对应的维度上为x的一半，其他维度上与x一致。支持空Tensor。</td>
218 <td>FLOAT、FLOAT16、BFLOAT16</td>218 <td>FLOAT、FLOAT16、BFLOAT16</td>
219 <td>ND</td>219 <td>ND</td>
220 <td>1-8</td>220 <td>1-8</td>
@@ -243,6 +243,15 @@ aclnnStatus aclnnClippedSwiglu(
243 </tbody>243 </tbody>
244 </table>244 </table>
245 245 
246+ <!-- npu="950" id7 -->
247+ - <term>Ascend 950PR/Ascend 950DT</term>：
248+ - 入参`groupIndexOptional`，不支持空tensor。
249+ <!-- end id7 -->
250+ <!-- npu="A3,910b" id8 -->
251+ - <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>、<term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term>：
252+ - 入参`groupIndexOptional`，支持空tensor。
253+ <!-- end id8 -->
254+ 
246- **返回值**255- **返回值**
247 256 
248 aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn_return_code.md)。257 aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn_return_code.md)。
Aactivation/clipped_swiglu/docs/aclnnClippedSwigluV2.md+518-0
@@ -0,0 +1,518 @@
1+# aclnnClippedSwigluV2
2+ 
C
Cchenjiao18 天前

新增api时，文档交付件不全

likedislike
3+[📄 查看源码](https://gitcode.com/cann/ops-nn/tree/master/activation/clipped_swiglu)
4+ 
5+## 产品支持情况
6+ 
7+<!-- npu="950" id1 -->
8+- <term>Ascend 950PR/Ascend 950DT</term>：支持
9+<!-- end id1 -->
10+<!-- npu="A3" id2 -->
11+- <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：不支持
12+<!-- end id2 -->
13+<!-- npu="910b" id3 -->
14+- <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term>：不支持
15+<!-- end id3 -->
16+<!-- npu="310b" id4 -->
17+- <term>Atlas 200I/500 A2 推理产品</term>：不支持
18+<!-- end id4 -->
19+<!-- npu="310p" id5 -->
20+- <term>Atlas 推理系列产品</term>：不支持
21+<!-- end id5 -->
22+<!-- npu="910" id6 -->
23+- <term>Atlas 训练系列产品</term>：不支持
24+<!-- end id6 -->
25+ 
26+## 功能说明
27+ 
28+- 接口功能：带截断的Swish门控线性单元激活函数，实现x的SwiGlu计算。本接口相较于aclnnClippedSwiglu，新增了一个输入参数：clamp_mode，用于支持部分新模型需要将clamp操作后移至silu激活之后的场景。
29+ 
30+- 计算公式：
31+ 
32+ 对给定的输入张量x，其维度为[a,b,c,d,e,f,g…]，aclnnClippedSwigluV2对其进行以下计算：
33+ 
34+ 1. 将x基于输入参数dim进行合轴，合轴后维度为[pre,cut,after]。其中cut轴为合轴之后需要切分为两个张量的轴，切分方式分为前后切分或者奇偶切分；pre，after可以等于1。例如当dim为3，合轴后x的维度为[a*b*c,d,e*f*g*…]。此外，由于after轴的元素为连续存放，且计算操作为逐元素的，因此将cut轴与after轴合并，得到x的维度为[pre,cut]。
35+ 
36+ 2. 根据输入参数group_index，对x的pre轴进行过滤处理，公式如下：
37+ 
38+ $$
39+ sum = \text{Sum}(group\_index)
40+ $$
41+ 
42+ $$
43+ x = x[ : sum, : ]
44+ $$
45+ 
46+ 其中sum表示group_index的所有元素之和。当不输入group_index时，跳过该步骤。
47+ 
48+ 3. 根据输入参数interleaved，对x进行切分，公式如下：
49+ 
50+ 当interleaved为true时，表示奇偶切分：
51+ 
52+ $$
53+ A = x[ : , : : 2]
54+ $$
55+ 
56+ $$
57+ B = x[ : , 1 : : 2]
58+ $$
59+ 
60+ 当interleaved为false时，表示前后切分：
61+ 
62+ $$
63+ h = x.shape[1] // 2
64+ $$
65+ 
66+ $$
67+ A = x[ : , : h]
68+ $$
69+ 
70+ $$
71+ B = x[ : , h : ]
72+ $$
73+ 4. 根据输入参数alpha、limit、bias、clamp_mode进行变体SwiGlu计算，公式如下：
74+ 
75+ 当clampMode为0时，表示clamp操作在silu之前：
76+ 
77+ $$
78+ A = A.clamp(min=None, max=limit)
79+ $$
80+ 
81+ $$
82+ B = B.clamp(min=-limit, max=limit)
83+ $$
84+ 
85+ $$
86+ y\_glu = A * sigmoid(alpha * A)
87+ $$
88+ 
89+ $$
90+ y = y\_glu * (B + bias)
91+ $$
92+ 
93+ 当clampMode为1时，表示clamp操作在silu之后：
94+ 
95+ $$
96+ y\_glu = A * sigmoid(A)
97+ $$
98+ 
99+ $$
100+ y\_glu = y\_glu.clamp(min=None, max=limit)
101+ $$
102+ 
103+ $$
104+ B = B.clamp(min=-limit, max=limit)
105+ $$
106+ 
107+ $$
108+ y = y\_glu * B
109+ $$
110+ 
111+ 5. 重塑输出张量y的维度数量与合轴前的x的维度数量一致，dim轴上的大小为x的一半，其他维度与x相同。
112+ 
113+## 函数原型
114+ 
115+每个算子分为[两段式接口](../../../docs/zh/context/two_phase_api.md)，必须先调用“aclnnClippedSwigluV2GetWorkspaceSize”接口获取计算所需workspace大小以及包含了算子计算流程的执行器，再调用“aclnnClippedSwigluV2”接口执行计算。
116+ 
117+```Cpp
118+aclnnStatus aclnnClippedSwigluV2GetWorkspaceSize(
119+ const aclTensor *x,
120+ const aclTensor *groupIndexOptional,
121+ int64_t dim,
122+ double alpha,
123+ double limit,
124+ double bias,
125+ bool interleaved,
126+ int64_t clampMode,
127+ const aclTensor *out,
128+ uint64_t *workspaceSize,
129+ aclOpExecutor **executor)
130+```
131+ 
132+```Cpp
133+aclnnStatus aclnnClippedSwigluV2(
134+ void *workspace,
135+ uint64_t workspaceSize,
136+ aclOpExecutor *executor,
137+ aclrtStream stream)
138+```
139+ 
140+## aclnnClippedSwigluV2GetWorkspaceSize
141+ 
142+- **参数说明**
143+ <table style="undefined;table-layout: fixed; width: 1567px"><colgroup>
144+ <col style="width: 270px">
145+ <col style="width: 120px">
146+ <col style="width: 300px">
147+ <col style="width: 330px">
148+ <col style="width: 212px">
149+ <col style="width: 100px">
150+ <col style="width: 90px">
151+ <col style="width: 145px">
152+ </colgroup>
153+ <thead>
154+ <tr>
155+ <th>参数名</th>
156+ <th>输入/输出</th>
157+ <th>描述</th>
158+ <th>使用说明</th>
159+ <th>数据类型</th>
160+ <th>数据格式</th>
161+ <th>维度(shape)</th>
162+ <th>非连续Tensor</th>
163+ </tr></thead>
164+ <tbody>
165+ <tr>
166+ <td>x（aclTensor*）</td>
167+ <td>输入</td>
168+ <td>公式中的输入x。</td>
169+ <td><ul><li>不支持空指针，维度必须大于0且shape必须在入参dim对应维度上是偶数。</li><li>支持空Tensor。</li></ul></td>
170+ <td>FLOAT、FLOAT16、BFLOAT16</td>
171+ <td>ND</td>
172+ <td>1-8</td>
173+ <td>-</td>
174+ </tr>
175+ <tr>
176+ <td>groupIndexOptional（aclTensor*）</td>
177+ <td>输入</td>
178+ <td>公式中的输入group_index，表示分组的情况。</td>
179+ <td><ul><li>支持空指针。</li><li>不为空指针时，维度要求为1维，长度不超过8192，且元素需大于等于0。第i个元素代表第i组需要处理x的batch数量。</li><li>不支持空Tensor。</li></ul></td>
180+ <td>INT64</td>
181+ <td>ND</td>
182+ <td>1</td>
183+ <td>-</td>
184+ </tr>
185+ <tr>
186+ <td>dim（int64_t）</td>
187+ <td>输入</td>
188+ <td>公式中的输入dim，表示对x进行合轴以及切分的维度序号。</td>
189+ <td>取值范围为[-x.dim(), x.dim()-1]。</td>
190+ <td>-</td>
191+ <td>-</td>
192+ <td>-</td>
193+ <td>-</td>
194+ </tr>
195+ <tr>
196+ <td>alpha（double）</td>
197+ <td>输入</td>
198+ <td>公式中的输入alpha，表示变体SwiGlu使用的参数。</td>
199+ <td>建议为1.702。</td>
200+ <td>-</td>
201+ <td>-</td>
202+ <td>-</td>
203+ <td>-</td>
204+ </tr>
205+ <tr>
206+ <td>limit（double）</td>
207+ <td>输入</td>
208+ <td>公式中的输入limit，表示变体SwiGlu使用的门限值。</td>
209+ <td>必须大于0，建议为7.0。</td>
210+ <td>-</td>
211+ <td>-</td>
212+ <td>-</td>
213+ <td>-</td>
214+ </tr>
215+ <tr>
216+ <td>bias（double）</td>
217+ <td>输入</td>
218+ <td>公式中的输入bias，表示变体SwiGlu使用的偏差参数。</td>
219+ <td>建议为1.0。</td>
220+ <td>-</td>
221+ <td>-</td>
222+ <td>-</td>
223+ <td>-</td>
224+ </tr>
225+ <tr>
226+ <td>interleaved（bool）</td>
227+ <td>输入</td>
228+ <td>公式中的输入interleaved，表示切分x时是否按奇偶方式切分</td>
229+ <td>设置为true表示对x进行奇偶切分，设置为false表示对x进行前后切分。</td>
230+ <td>-</td>
231+ <td>-</td>
232+ <td>-</td>
233+ <td>-</td>
234+ </tr>
235+ <tr>
236+ <td>clampMode（int64_t）</td>
237+ <td>输入</td>
238+ <td>公式中的输入clampMode，表示clamp操作与silu操作执行的先后顺序</td>
239+ <td>设置为1表示将A的clamp操作后移至silu激活之后，即对silu(A)的结果做clamp；设置为0表示A的clamp操作在silu激活之前。</td>
240+ <td>-</td>
241+ <td>-</td>
242+ <td>-</td>
243+ <td>-</td>
244+ </tr>
245+ <tr>
246+ <td>out（aclTensor*）</td>
247+ <td>输出</td>
248+ <td>公式中的输出y。</td>
249+ <td><ul><li>不支持空指针，shape在入参dim对应的维度上为x的一半，其他维度上与x一致。</li><li>支持空Tensor。</li></ul></td>
250+ <td>FLOAT、FLOAT16、BFLOAT16</td>
251+ <td>ND</td>
252+ <td>1-8</td>
253+ <td>-</td>
254+ </tr>
255+ <tr>
256+ <td>workspaceSize（uint64_t*）</td>
257+ <td>输出</td>
258+ <td>返回需要在Device侧申请的workspace大小。</td>
259+ <td>-</td>
260+ <td>-</td>
261+ <td>-</td>
262+ <td>-</td>
263+ <td>-</td>
264+ </tr>
265+ <tr>
266+ <td>executor（aclOpExecutor**）</td>
267+ <td>输出</td>
268+ <td>返回op执行器，包含了算子计算流程。</td>
269+ <td>-</td>
270+ <td>-</td>
271+ <td>-</td>
272+ <td>-</td>
273+ <td>-</td>
274+ </tr>
275+ </tbody>
276+ </table>
277+ 
278+- **返回值**
279+ 
280+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn_return_code.md)。
281+ 
282+ 第一段接口会完成入参校验，出现以下场景时报错：
283+ 
284+ <table style="undefined;table-layout: fixed;width: 979px"><colgroup>
285+ <col style="width: 272px">
286+ <col style="width: 103px">
287+ <col style="width: 604px">
288+ </colgroup>
289+ <thead>
290+ <tr>
291+ <th>返回码</th>
292+ <th>错误码</th>
293+ <th>描述</th>
294+ </tr>
295+ </thead>
296+ <tbody>
297+ <tr>
298+ <td>ACLNN_ERR_PARAM_NULLPTR</td>
299+ <td>161001</td>
300+ <td>传入的x、out是空指针。</td>
301+ </tr>
302+ <tr>
303+ <td rowspan="3">ACLNN_ERR_PARAM_INVALID</td>
304+ <td rowspan="3">161002</td>
305+ <td>输入或输出的数据类型不在支持的范围内。</td>
306+ </tr>
307+ <tr>
308+ <td>输入或输出的参数维度不在支持的范围内。</td>
309+ </tr>
310+ <tr>
311+ <td>dim或clamp_mode不在指定的取值范围内。</td>
312+ </tr>
313+ </tbody>
314+ </table>
315+ 
316+## aclnnClippedSwigluV2
317+ 
318+- **参数说明：**
319+ <table style="undefined;table-layout: fixed; width: 953px"><colgroup>
320+ <col style="width: 173px">
321+ <col style="width: 112px">
322+ <col style="width: 668px">
323+ </colgroup>
324+ <thead>
325+ <tr>
326+ <th>参数名</th>
327+ <th>输入/输出</th>
328+ <th>描述</th>
329+ </tr></thead>
330+ <tbody>
331+ <tr>
332+ <td>workspace</td>
333+ <td>输入</td>
334+ <td>在Device侧申请的workspace内存地址。</td>
335+ </tr>
336+ <tr>
337+ <td>workspaceSize</td>
338+ <td>输入</td>
339+ <td>在Device侧申请的workspace大小，由第一段接口aclnnClippedSwigluV2GetWorkspaceSize获取。</td>
340+ </tr>
341+ <tr>
342+ <td>executor</td>
343+ <td>输入</td>
344+ <td>op执行器，包含了算子计算流程。</td>
345+ </tr>
346+ <tr>
347+ <td>stream</td>
348+ <td>输入</td>
349+ <td>指定执行任务的Stream。</td>
350+ </tr>
351+ </tbody>
352+ </table>
353+ 
354+- **返回值**：
355+ 
356+ aclnnStatus：返回状态码，具体参见[aclnn返回码](../../../docs/zh/context/aclnn_return_code.md)。
357+ 
358+## 约束说明
359+ 
360+确定性计算：aclnnClippedSwigluV2默认为确定性实现，暂不支持非确定性实现，即便通过确定性计算配置也不会生效。
C
Cchenjiao18 天前

请确认是否涉及batch一致性

likedislike
361+ 
362+## 调用示例
363+ 
364+示例代码如下，仅供参考，具体编译和执行过程请参考[编译与运行样例](../../../docs/zh/context/compile_and_run_sample.md)。
365+ 
366+```Cpp
367+#include <iostream>
368+#include <vector>
369+#include "acl/acl.h"
370+#include "aclnnop/aclnn_clipped_swiglu_v2.h"
371+ 
372+#define CHECK_RET(cond, return_expr) \
373+ do { \
374+ if (!(cond)) { \
375+ return_expr; \
376+ } \
377+ } while (0)
378+ 
379+#define LOG_PRINT(message, ...) \
380+ do { \
381+ printf(message, ##__VA_ARGS__); \
382+ } while (0)
383+ 
384+int64_t GetShapeSize(const std::vector<int64_t>& shape)
385+{
386+ int64_t shapeSize = 1;
387+ for (auto i : shape) {
388+ shapeSize *= i;
389+ }
390+ return shapeSize;
391+}
392+ 
393+int Init(int32_t deviceId, aclrtStream* stream)
394+{
395+ // 固定写法，acl初始化
396+ auto ret = aclInit(nullptr);
397+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
398+ ret = aclrtSetDevice(deviceId);
399+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
400+ ret = aclrtCreateStream(stream);
401+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
402+ return 0;
403+}
404+ 
405+template <typename T>
406+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
407+ aclDataType dataType, aclTensor** tensor)
408+{
409+ auto size = GetShapeSize(shape) * sizeof(T);
410+ // 调用aclrtMalloc申请device侧内存
411+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
412+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
413+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
414+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
415+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
416+ 
417+ // 计算连续tensor的strides
418+ std::vector<int64_t> strides(shape.size(), 1);
419+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
420+ strides[i] = shape[i + 1] * strides[i + 1];
421+ }
422+ 
423+ // 调用aclCreateTensor接口创建aclTensor
424+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
425+ shape.data(), shape.size(), *deviceAddr);
426+ return 0;
427+}
428+ 
429+int main()
430+{
431+ // 1. （固定写法）device/stream初始化，参考acl API手册
432+ // 根据自己的实际device填写deviceId
433+ int32_t deviceId = 0;
434+ aclrtStream stream;
435+ auto ret = Init(deviceId, &stream);
436+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
437+ 
438+ // 2. 构造输入与输出，需要根据API的接口自定义构造
439+ std::vector<int64_t> xShape = {2, 32};
440+ std::vector<int64_t> groupIndexShape = {1};
441+ std::vector<int64_t> outShape = {2, 16};
442+ void* xDeviceAddr = nullptr;
443+ void* groupIndexDeviceAddr = nullptr;
444+ void* outDeviceAddr = nullptr;
445+ aclTensor* x = nullptr;
446+ aclTensor* groupIndex = nullptr;
447+ aclTensor* out = nullptr;
448+ std::vector<float> xHostData = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15,
449+ 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28, 29, 30, 31,
450+ 32, 33, 34, 35, 36, 37, 38, 39, 40, 41, 42, 43, 44, 45, 46, 47,
451+ 48, 49, 50, 51, 52, 53, 54, 55, 56, 57, 58, 59, 60, 61, 62, 63};
452+ std::vector<int64_t> groupIndexData = {1};
453+ std::vector<float> outHostData = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
454+ 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
455+ 
456+ int dim = -1;
457+ float alpha = 1.0;
458+ float limit = 7.0;
459+ float bias = 1.702;
460+ bool interleaved = true;
461+ int64_t clampMode = 1;
462+ // 创建x aclTensor
463+ ret = CreateAclTensor(xHostData, xShape, &xDeviceAddr, aclDataType::ACL_FLOAT, &x);
464+ CHECK_RET(ret == ACL_SUCCESS, return ret);
465+ // 创建groupIndex aclTensor
466+ ret = CreateAclTensor(groupIndexData, groupIndexShape, &groupIndexDeviceAddr, aclDataType::ACL_INT64, &groupIndex);
467+ CHECK_RET(ret == ACL_SUCCESS, return ret);
468+ // 创建out aclTensor
469+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
470+ CHECK_RET(ret == ACL_SUCCESS, return ret);
471+ // 3. 调用CANN算子库API，需要修改为具体的Api名称
472+ uint64_t workspaceSize = 0;
473+ aclOpExecutor* executor;
474+ // 调用aclnnClippedSwigluV2第一段接口
475+ ret = aclnnClippedSwigluV2GetWorkspaceSize(x, groupIndex, dim, alpha, limit, bias, interleaved, clampMode, out,
476+ &workspaceSize, &executor);
477+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnClippedSwigluV2GetWorkspaceSize failed. ERROR: %d\n", ret);
478+ return ret);
479+ // 根据第一段接口计算出的workspaceSize申请device内存
480+ void* workspaceAddr = nullptr;
481+ if (workspaceSize > 0) {
482+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
483+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
484+ }
485+ // 调用aclnnClippedSwigluV2第二段接口
486+ ret = aclnnClippedSwigluV2(workspaceAddr, workspaceSize, executor, stream);
487+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnClippedSwigluV2 failed. ERROR: %d\n", ret); return ret);
488+ 
489+ // 4. （固定写法）同步等待任务执行结束
490+ ret = aclrtSynchronizeStream(stream);
491+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
492+ 
493+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
494+ auto size = GetShapeSize(outShape);
495+ std::vector<float> resultData(size, 0);
496+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr,
497+ size * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
498+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
499+ for (int64_t i = 0; i < size; i++) {
500+ LOG_PRINT("result[%ld] is: %f\n", i, resultData[i]);
501+ }
502+ // 6. 释放aclTensor和aclScalar，需要根据具体API的接口定义修改
503+ aclDestroyTensor(x);
504+ aclDestroyTensor(groupIndex);
505+ aclDestroyTensor(out);
506+ // 7. 释放device资源，需要根据具体API的接口定义修改
507+ aclrtFree(xDeviceAddr);
508+ aclrtFree(groupIndexDeviceAddr);
509+ aclrtFree(outDeviceAddr);
510+ if (workspaceSize > 0) {
511+ aclrtFree(workspaceAddr);
512+ }
513+ aclrtDestroyStream(stream);
514+ aclrtResetDevice(deviceId);
515+ aclFinalize();
516+ return 0;
517+}
518+```
Aactivation/clipped_swiglu/docs/torchapi_clipped_swiglu.md+211-0
@@ -0,0 +1,211 @@
1+# clipped_swiglu
2+ 
3+## 产品支持情况
4+ 
5+<!-- npu="950" id1 -->
6+- <term>Ascend 950PR/Ascend 950DT</term>：支持
7+<!-- end id1 -->
8+<!-- npu="A3" id2 -->
9+- <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term>：支持
10+<!-- end id2 -->
11+<!-- npu="910b" id3 -->
12+- <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term>：支持
13+<!-- end id3 -->
14+<!-- npu="310b" id4 -->
15+- <term>Atlas 200I/500 A2 推理产品</term>：不支持
16+<!-- end id4 -->
17+<!-- npu="310p" id5 -->
18+- <term>Atlas 推理系列产品</term>：不支持
19+<!-- end id5 -->
20+<!-- npu="910" id6 -->
21+- <term>Atlas 训练系列产品</term>：不支持
22+<!-- end id6 -->
23+ 
24+## 功能说明
25+ 
26+- 接口功能：
27+ 
28+ 带截断的 Swish 门控线性单元（ClippedSwiGLU）激活函数。相较于标准 SwiGLU，新增 `group_index`、`alpha`、`limit`、`bias`、`interleaved`、`clamp_mode` 等参数，用于支持 GPT-OSS 模型使用的变体 SwiGlu、MoE 模型的分组场景以及部分新模型需要将clamp操作后移至silu激活之后的场景。
29+ 
30+- 计算公式：
31+ 
32+ 对给定的输入张量 `x`，其维度为 `[a, b, c, d, e, f, g, ...]`，`clipped_swiglu` 进行以下计算：
33+ 
34+ 1. 将 `x` 基于输入参数 `dim` 进行合轴，合轴后维度为 `[pre, cut, after]`。其中 `cut` 轴为合轴之后需要切分为两个张量的轴，切分方式分为前后切分或者奇偶切分；`pre`、`after` 可以等于 1。例如当 `dim` 为 3 时，合轴后 `x` 的维度为 `[a*b*c, d, e*f*g*...]`。由于 `after` 轴元素连续存放且计算为逐元素的，将 `cut` 轴与 `after` 轴合并，得到 `x` 的维度为 `[pre, cut]`。
35+ 
36+ 2. 根据输入参数 `group_index`，对 `x` 的 `pre` 轴进行过滤处理：
37+ 
38+ $$
39+ sum = \text{Sum}(group\_index)
40+ $$
41+ 
42+ $$
43+ x = x[ : sum, : ]
44+ $$
45+ 
46+ 其中 `sum` 表示 `group_index` 所有元素之和。当不输入 `group_index` 时，跳过该步骤。
47+ 
48+ 3. 根据输入参数 `interleaved`，对 `x` 进行切分：
49+ 
50+ 当 `interleaved=True`（奇偶切分）：
51+ 
52+ $$
53+ A = x[ : , ::2], \quad B = x[ : , 1::2]
54+ $$
55+ 
56+ 当 `interleaved=False`（前后切分）：
57+ 
58+ $$
59+ h = x.shape[1] // 2
60+ $$
61+ 
62+ $$
63+ A = x[ : , : h], \quad B = x[ : , h : ]
64+ $$
65+ 
66+ 4. 根据 `alpha`、`limit`、`bias`、`clamp_mode` 进行变体 SwiGlu 计算：
67+ 
68+ 当 `clamp_mode=0`（clamp 操作在 silu 之前）：
69+ 
70+ $$
71+ A = A.clamp(min=\text{None}, max=limit)
72+ $$
73+ 
74+ $$
75+ B = B.clamp(min=-limit, max=limit)
76+ $$
77+ 
78+ $$
79+ y\_glu = A \cdot sigmoid(\alpha \cdot A)
80+ $$
81+ 
82+ $$
83+ y = y\_glu \cdot (B + bias)
84+ $$
85+ 
86+ 当 `clamp_mode=1`（clamp 操作在 silu 之后，需底层支持 V2 接口）：
87+ 
88+ $$
89+ y\_glu = A \cdot sigmoid(A)
90+ $$
91+ 
92+ $$
93+ y\_glu = y\_glu.clamp(min=\text{None}, max=limit)
94+ $$
95+ 
96+ $$
97+ B = B.clamp(min=-limit, max=limit)
98+ $$
99+ 
100+ $$
101+ y = y\_glu \cdot B
102+ $$
103+ 
104+ 5. 重塑输出张量 `y` 的维度数量与合轴前的 `x` 一致，`dim` 轴上的大小为 `x` 的一半，其他维度与 `x` 相同。
105+ 
106+## 函数原型
107+ 
108+```python
109+cann_ops_nn.clipped_swiglu(
110+ x,
111+ *,
112+ group_index=None,
113+ dim=-1,
114+ alpha=1.702,
115+ limit=7.0,
116+ bias=1.0,
117+ interleaved=True,
118+ clamp_mode=0,
119+) -> Tensor
120+```
121+ 
122+## 参数说明
123+ 
124+| 参数名 | 参数类型 | 可选/必选 | 描述 | 数据类型 | 维度(shape) |
125+| --- | --- | --- | --- | --- | --- |
126+| `x` | Tensor | 必选 | 公式中的输入 `x`，在 `dim` 对应维度上必须为偶数。 | `torch.float16`、`torch.bfloat16`、`torch.float32` | 1-8 维 |
127+| `group_index` | Tensor | 可选 | 公式中的 `group_index`。第 `i` 个元素代表第 `i` 组需要处理 `x` 的 batch 数量。传入 `None` 表示不分组。 | `torch.int64` | 1 维，长度不超过 8192，元素需大于等于 0 |
128+| `dim` | int | 可选 | 对 `x` 进行合轴以及切分的维度序号，取值范围 `[-x.dim(), x.dim()-1]`。默认值 `-1`。 | - | - |
129+| `alpha` | float | 可选 | 变体 SwiGlu 的缩放参数，建议值 `1.702`。默认值 `1.702`。 | - | - |
130+| `limit` | float | 可选 | 变体 SwiGlu 的门限值，必须大于 0，建议值 `7.0`。默认值 `7.0`。 | - | - |
131+| `bias` | float | 可选 | 变体 SwiGlu 的偏差参数，建议值 `1.0`。默认值 `1.0`。 | - | - |
132+| `interleaved` | bool | 可选 | 切分 `x` 的方式。`True` 表示奇偶切分，`False` 表示前后切分。默认值 `True`。 | - | - |
133+| `clamp_mode` | int | 可选 | clamp 操作与 silu 操作的先后顺序。`0` 表示 clamp 在 silu 之前，`1` 表示 clamp 在 silu 之后。默认值 `0`。 | - | - |
134+ 
135+## 返回值说明
136+ 
137+| 参数名 | 参数类型 | 描述 | 数据类型 | 维度(shape) |
138+| --- | --- | --- | --- | --- |
139+| `y` | Tensor | ClippedSwiglu 激活结果。 | 与 `x` 相同 | 与 `x` 相同，但 `dim` 对应维度上为 `x.shape[dim] // 2` |
140+ 
141+## 约束说明
142+ 
143+- 该接口支持单算子模式和 TorchAir 图模式调用。
144+- `x`、`group_index` 均需为 NPU Tensor；可选 Tensor 可以传 `None`。
145+- `clamp_mode` 仅支持取值 `0` 或 `1`；当目标芯片未注册 `aclnnClippedSwigluV2` 内核时，会调用 `aclnnClippedSwiglu`，传入 `clamp_mode=1` 会被忽略并按 `clamp_mode=0` 执行。
146+ 
147+## 确定性计算
148+ 
149+默认支持确定性计算。
150+ 
151+## 调用示例
152+ 
153+- 单算子模式调用：
154+ 
155+ ```python
156+ import torch
157+ import torch_npu
158+ import cann_ops_nn
159+ 
160+ x = torch.randn(32, 128, dtype=torch.float16).npu()
161+ group_index = torch.randint(1, 10, (2, ), dtype=torch.int64).npu()
162+ 
163+ y = cann_ops_nn.clipped_swiglu(x, group_index=group_index, dim=-1, alpha=1.702, limit=7.0, bias=1.0, interleaved=True, clamp_mode=1)
164+ 
165+ print("y:============", y.shape, y.cpu())
166+ ```
167+ 
168+- 图模式（torchair）调用：
169+ 
170+ ```python
171+ import torch
172+ import torch_npu
173+ import torchair
174+ import cann_ops_nn
175+ 
176+ npu_backend = "npu"
177+ 
178+ class NetModel(torch.nn.Module):
179+ def __init__(self):
180+ super().__init__()
181+ 
182+ def forward(
183+ self,
184+ x,
185+ group_index,
186+ dim,
187+ alpha,
188+ limit,
189+ bias,
190+ interleaved,
191+ clamp_mode
192+ ):
193+ return cann_ops_nn.clipped_swiglu(
194+ x, group_index=group_index, dim=dim, alpha=alpha, limit=limit, bias=bias, interleaved=interleaved, clamp_mode=clamp_mode
195+ )
196+ 
197+ def clipped_swiglu_test():
198+ x = torch.randn(32, 128, dtype=torch.float16).npu()
199+ group_index = torch.randint(1, 10, (2, ), dtype=torch.int64).npu()
200+ 
201+ model = NetModel()
202+ config = torchair.CompilerConfig()
203+ npu_backend = torchair.get_npu_backend(compiler_config=config)
204+ model = torch.compile(model, backend=npu_backend, dynamic=True, fullgraph=True)
205+ 
206+ y = model(x, group_index=group_index, dim=-1, alpha=1.702, limit=7.0, bias=1.0, interleaved=True, clamp_mode=1)
207+ print("y:============", y.shape, y.cpu())
208+ 
209+ if __name__ == "__main__":
210+ clipped_swiglu_test()
211+ ```
Aactivation/clipped_swiglu/examples/test_aclnn_clipped_swiglu_v2.cpp+161-0
@@ -0,0 +1,161 @@
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_clipped_swiglu_v2.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+{
30+ int64_t shapeSize = 1;
31+ for (auto i : shape) {
32+ shapeSize *= i;
33+ }
34+ return shapeSize;
35+}
36+ 
37+int Init(int32_t deviceId, aclrtStream* stream)
38+{
39+ // 固定写法，acl初始化
40+ auto ret = aclInit(nullptr);
41+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
42+ ret = aclrtSetDevice(deviceId);
43+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
44+ ret = aclrtCreateStream(stream);
45+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
46+ return 0;
47+}
48+ 
49+template <typename T>
50+int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
51+ aclDataType dataType, aclTensor** tensor)
52+{
53+ auto size = GetShapeSize(shape) * sizeof(T);
54+ // 调用aclrtMalloc申请device侧内存
55+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
56+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
57+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
58+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
59+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
60+ 
61+ // 计算连续tensor的strides
62+ std::vector<int64_t> strides(shape.size(), 1);
63+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
64+ strides[i] = shape[i + 1] * strides[i + 1];
65+ }
66+ 
67+ // 调用aclCreateTensor接口创建aclTensor
68+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
69+ shape.data(), shape.size(), *deviceAddr);
70+ return 0;
71+}
72+ 
73+int main()
74+{
75+ // 1. （固定写法）device/stream初始化，参考acl API手册
76+ // 根据自己的实际device填写deviceId
77+ int32_t deviceId = 0;
78+ aclrtStream stream;
79+ auto ret = Init(deviceId, &stream);
80+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
81+ 
82+ // 2. 构造输入与输出，需要根据API的接口自定义构造
83+ std::vector<int64_t> xShape = {2, 32};
84+ std::vector<int64_t> groupIndexShape = {1};
85+ std::vector<int64_t> outShape = {2, 16};
86+ void* xDeviceAddr = nullptr;
87+ void* groupIndexDeviceAddr = nullptr;
88+ void* outDeviceAddr = nullptr;
89+ aclTensor* x = nullptr;
90+ aclTensor* groupIndex = nullptr;
91+ aclTensor* out = nullptr;
92+ std::vector<float> xHostData = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15,
93+ 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28, 29, 30, 31,
94+ 32, 33, 34, 35, 36, 37, 38, 39, 40, 41, 42, 43, 44, 45, 46, 47,
95+ 48, 49, 50, 51, 52, 53, 54, 55, 56, 57, 58, 59, 60, 61, 62, 63};
96+ std::vector<int64_t> groupIndexData = {1};
97+ std::vector<float> outHostData = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
98+ 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
99+ 
100+ int dim = -1;
101+ float alpha = 1.0;
102+ float limit = 7.0;
103+ float bias = 1.702;
104+ bool interleaved = true;
105+ int64_t clampMode = 1;
106+ // 创建x aclTensor
107+ ret = CreateAclTensor(xHostData, xShape, &xDeviceAddr, aclDataType::ACL_FLOAT, &x);
108+ CHECK_RET(ret == ACL_SUCCESS, return ret);
109+ // 创建groupIndex aclTensor
110+ ret = CreateAclTensor(groupIndexData, groupIndexShape, &groupIndexDeviceAddr, aclDataType::ACL_INT64, &groupIndex);
atomgit-bot
atomgit-botatomgit-bot27 天前

🟡 Medium Priority

在 test_aclnn_clipped_swiglu_v2.cpp 第 110 行，调用 CreateAclTensor(groupIndexData, ...) 创建 groupIndex 张量之后，缺少 CHECK_RET 宏对返回值的检查。

证据链：

  • 第 113 行的 CHECK_RET 实际只校验了 out 张量的创建结果

失败模式： 若 groupIndex 张量创建失败（如内存不足），但 out 张量创建成功，则 ret 被覆盖为 ACL_SUCCESS，CHECK_RET 通过。代码随后将无效的 groupIndex 指针传入 aclnnClippedSwigluV2GetWorkspaceSize，可能导致崩溃或未定义行为。

建议：在 CreateAclTensor(groupIndexData, ...) 调用之后立即添加 CHECK_RET(ret == ACL_SUCCESS, return ret); 以校验 groupIndex 张量创建是否成功。

改动建议
110
- ret = CreateAclTensor(groupIndexData, groupIndexShape, &groupIndexDeviceAddr, aclDataType::ACL_INT64, &groupIndex);
110
+ ret = CreateAclTensor(groupIndexData, groupIndexShape, &groupIndexDeviceAddr, aclDataType::ACL_INT64, &groupIndex);
111
+ CHECK_RET(ret == ACL_SUCCESS, return ret);
应用建议
likedislike
不准确？
clwsy
clwsy
22 天前 评论：
111+ CHECK_RET(ret == ACL_SUCCESS, return ret);
112+ // 创建out aclTensor
atomgit-bot
atomgit-botatomgit-bot27 天前

🟡 Medium Priority

第 110 行 CreateAclTensor(groupIndexData, ...) 的返回值赋给了 ret，但紧接着第 112 行 CreateAclTensor(outHostData, ...) 再次覆写了 ret，而第 113 行的 CHECK_RET 只检查了最后一次（out tensor）的返回值。groupIndex tensor 创建失败时（ret != ACL_SUCCESS），groupIndex 仍为 nullptr，后续第 118 行将其传入 aclnnClippedSwigluV2GetWorkspaceSize 会导致空指针解引用或未定义行为。

对比第 107-108 行对 x tensor 正确使用了 CHECK_RET，此处明显遗漏。

建议：在第 111 行之后添加 CHECK_RET(ret == ACL_SUCCESS, return ret); 对 groupIndex tensor 创建的返回值进行检查。

改动建议
112
+ ret = CreateAclTensor(groupIndexData, groupIndexShape, &groupIndexDeviceAddr, aclDataType::ACL_INT64, &groupIndex);
112
- // 创建out aclTensor
113
+ CHECK_RET(ret == ACL_SUCCESS, return ret);
应用建议
likedislike
不准确？
clwsy
clwsy
22 天前 评论：
113+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
114+ CHECK_RET(ret == ACL_SUCCESS, return ret);
115+ // 3. 调用CANN算子库API，需要修改为具体的Api名称
116+ uint64_t workspaceSize = 0;
117+ aclOpExecutor* executor;
118+ // 调用aclnnClippedSwigluV2第一段接口
119+ ret = aclnnClippedSwigluV2GetWorkspaceSize(x, groupIndex, dim, alpha, limit, bias, interleaved, clampMode, out,
120+ &workspaceSize, &executor);
121+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnClippedSwigluV2GetWorkspaceSize failed. ERROR: %d\n", ret);
122+ return ret);
123+ // 根据第一段接口计算出的workspaceSize申请device内存
124+ void* workspaceAddr = nullptr;
125+ if (workspaceSize > 0) {
126+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
127+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
128+ }
129+ // 调用aclnnClippedSwigluV2第二段接口
130+ ret = aclnnClippedSwigluV2(workspaceAddr, workspaceSize, executor, stream);
131+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnClippedSwigluV2 failed. ERROR: %d\n", ret); return ret);
132+ 
133+ // 4. （固定写法）同步等待任务执行结束
134+ ret = aclrtSynchronizeStream(stream);
135+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
136+ 
137+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
138+ auto size = GetShapeSize(outShape);
139+ std::vector<float> resultData(size, 0);
140+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr,
141+ size * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
142+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
143+ for (int64_t i = 0; i < size; i++) {
144+ LOG_PRINT("result[%ld] is: %f\n", i, resultData[i]);
145+ }
146+ // 6. 释放aclTensor和aclScalar，需要根据具体API的接口定义修改
147+ aclDestroyTensor(x);
148+ aclDestroyTensor(groupIndex);
149+ aclDestroyTensor(out);
150+ // 7. 释放device资源，需要根据具体API的接口定义修改
151+ aclrtFree(xDeviceAddr);
152+ aclrtFree(groupIndexDeviceAddr);
153+ aclrtFree(outDeviceAddr);
154+ if (workspaceSize > 0) {
155+ aclrtFree(workspaceAddr);
156+ }
157+ aclrtDestroyStream(stream);
158+ aclrtResetDevice(deviceId);
159+ aclFinalize();
160+ return 0;
161+}
Mactivation/clipped_swiglu/op_graph/clipped_swiglu_proto.h+5-2
@@ -37,8 +37,10 @@ namespace ge {
37* @li alpha: An optional float. The activation coefficient for the GLU activation function, default is 1.702.37* @li alpha: An optional float. The activation coefficient for the GLU activation function, default is 1.702.
38* @li limit: An optional float. The threshold limit for SWIGLU input, default is 7.0.38* @li limit: An optional float. The threshold limit for SWIGLU input, default is 7.0.
39* @li bias: An optional float. The bias applied during SWIGLU linear computation, default is 1.0.39* @li bias: An optional float. The bias applied during SWIGLU linear computation, default is 1.0.
40-* @li interleaved: An optional bool. The way of splitting x: true for interleaved splitting, false for front-back40+* @li interleaved: An optional bool. Indicates the way of splitting x, default is true.
41-splitting, default is true.41+* When it's true, it means interleaved splitting; when it's false, it means front-back splitting.
42+* @li clamp_mode: An optional int. Indicates the order in which the clamp and silu operations are executed.
43+* When it's 0, the clamp operation happens before silu; when it's 1, it happens after silu, default is 0.
42 44 
43* @attention Constraints:45* @attention Constraints:
44* The dim dimension of x must be divisible by 2, and the dim dimension of y must be equal to the dim dimension of x46* The dim dimension of x must be divisible by 2, and the dim dimension of y must be equal to the dim dimension of x
@@ -53,6 +55,7 @@ REG_OP(ClippedSwiglu)
53 .ATTR(limit, Float, 7.0)55 .ATTR(limit, Float, 7.0)
54 .ATTR(bias, Float, 1.0)56 .ATTR(bias, Float, 1.0)
55 .ATTR(interleaved, Bool, true)57 .ATTR(interleaved, Bool, true)
58+ .ATTR(clamp_mode, Int, 0)
56 .OP_END_FACTORY_REG(ClippedSwiglu)59 .OP_END_FACTORY_REG(ClippedSwiglu)
57} // namespace ge60} // namespace ge
58#endif // OPS_BUILT_IN_OP_PROTO_INC_CLIPPED_SWIGLU_PROTO_H_61#endif // OPS_BUILT_IN_OP_PROTO_INC_CLIPPED_SWIGLU_PROTO_H_
Mactivation/clipped_swiglu/op_host/clipped_swiglu_def.cpp+2-0
@@ -15,6 +15,7 @@
15#include <register/op_def_registry.h>15#include <register/op_def_registry.h>
16 16 
17namespace ops {17namespace ops {
18+constexpr uint32_t CLIPPED_SWIGLU_VERSION_TWO = 2;
18constexpr float DEFAULT_ALPHA = 1.702;19constexpr float DEFAULT_ALPHA = 1.702;
19constexpr float DEFAULT_LIMIT = 7.0;20constexpr float DEFAULT_LIMIT = 7.0;
20 21 
@@ -38,6 +39,7 @@ public:
38 this->Attr("limit").AttrType(OPTIONAL).Float(DEFAULT_LIMIT);39 this->Attr("limit").AttrType(OPTIONAL).Float(DEFAULT_LIMIT);
39 this->Attr("bias").AttrType(OPTIONAL).Float(1.0);40 this->Attr("bias").AttrType(OPTIONAL).Float(1.0);
40 this->Attr("interleaved").AttrType(OPTIONAL).Bool(true);41 this->Attr("interleaved").AttrType(OPTIONAL).Bool(true);
42+ this->Attr("clamp_mode").AttrType(OPTIONAL).Version(CLIPPED_SWIGLU_VERSION_TWO).Int(0);
41 43 
42 this->AICore().AddConfig("ascend910b");44 this->AICore().AddConfig("ascend910b");
43 this->AICore().AddConfig("ascend910_93");45 this->AICore().AddConfig("ascend910_93");
Mactivation/clipped_swiglu/op_host/clipped_swiglu_tiling_arch35.cpp+19-7
@@ -36,6 +36,7 @@ constexpr int64_t ALPHA_INDEX = 1;
36constexpr int64_t LIMIT_INDEX = 2;36constexpr int64_t LIMIT_INDEX = 2;
37constexpr int64_t BIAS_INDEX = 3;37constexpr int64_t BIAS_INDEX = 3;
38constexpr int64_t INTERLEAVED_INDEX = 4;38constexpr int64_t INTERLEAVED_INDEX = 4;
39+constexpr int64_t CLAMPMODE_INDEX = 5;
39 40 
40constexpr int64_t CONST_2 = 2;41constexpr int64_t CONST_2 = 2;
41constexpr int64_t CONST_4 = 4;42constexpr int64_t CONST_4 = 4;
@@ -87,6 +88,7 @@ private:
87 int64_t dtypeSize_ = CONST_2;88 int64_t dtypeSize_ = CONST_2;
88 int64_t isGroup_ = 0;89 int64_t isGroup_ = 0;
89 int64_t isInterleaved_ = 1;90 int64_t isInterleaved_ = 1;
91+ int64_t clampMode_ = 0;
90 float gluLimit_ = 0.0;92 float gluLimit_ = 0.0;
91 float gluAlpha_ = 0.0;93 float gluAlpha_ = 0.0;
92 float gluBias_ = 0.0;94 float gluBias_ = 0.0;
@@ -215,6 +217,13 @@ ge::graphStatus ClippedSwigluArch35Tiling::CheckAndGetXAndAttrs()
215 bool interleaved = attrInterleaved == nullptr ? true : *attrInterleaved;217 bool interleaved = attrInterleaved == nullptr ? true : *attrInterleaved;
216 isInterleaved_ = interleaved ? 1 : 0;218 isInterleaved_ = interleaved ? 1 : 0;
217 219 
220+ auto* attrClampMode = attrs->GetAttrPointer<int64_t>(CLAMPMODE_INDEX);
221+ clampMode_ = attrClampMode == nullptr ? 0 : *attrClampMode;
222+ OP_CHECK_IF(
223+ clampMode_ != 0 && clampMode_ != 1,
224+ OP_LOGE_WITH_INVALID_ATTR(context_->GetNodeName(), "clamp_mode", std::to_string(clampMode_), "must be 1 or 0"),
225+ return ge::GRAPH_FAILED);
226+ 
218 auto shapeX = context_->GetInputShape(X_INDEX);227 auto shapeX = context_->GetInputShape(X_INDEX);
219 OP_CHECK_NULL_WITH_CONTEXT(context_, shapeX);228 OP_CHECK_NULL_WITH_CONTEXT(context_, shapeX);
220 const gert::Shape& inputShapeX = shapeX->GetStorageShape();229 const gert::Shape& inputShapeX = shapeX->GetStorageShape();
@@ -252,11 +261,11 @@ ge::graphStatus ClippedSwigluArch35Tiling::CheckAndGetGroupIndex()
252 ge::TypeUtils::DataTypeToSerialString(groupIndexDtype).c_str(), "int64"),261 ge::TypeUtils::DataTypeToSerialString(groupIndexDtype).c_str(), "int64"),
253 return ge::GRAPH_FAILED);262 return ge::GRAPH_FAILED);
254 groupNum_ = inputShapeGroupIndex.GetDim(0);263 groupNum_ = inputShapeGroupIndex.GetDim(0);
255- OP_CHECK_IF(264+ OP_CHECK_IF(groupNum_ > MAX_GROUP_NUM || groupNum_ <= 0,
256- groupNum_ > MAX_GROUP_NUM,265+ OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(
257- OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(context_->GetNodeName(), "group_index", std::to_string(groupNum_),266+ context_->GetNodeName(), "group_index", std::to_string(groupNum_),
258- "the number of elements of group_index should be no more than 8192"),267+ "the number of elements of group_index should be greater than 0 and no more than 8192"),
259- return ge::GRAPH_FAILED);268+ return ge::GRAPH_FAILED);
260 }269 }
261 return ge::GRAPH_SUCCESS;270 return ge::GRAPH_SUCCESS;
262}271}
@@ -383,8 +392,10 @@ void ClippedSwigluArch35Tiling::SetTilingKey()
383{392{
384 uint64_t isInterleavedKey = (isInterleaved_ != 0) ? TPL_INTERLEAVED_TRUE : TPL_INTERLEAVED_FALSE;393 uint64_t isInterleavedKey = (isInterleaved_ != 0) ? TPL_INTERLEAVED_TRUE : TPL_INTERLEAVED_FALSE;
385 uint64_t isGroupKey = (isGroup_ != 0) ? TPL_GROUP_INDEX : TPL_NO_GROUP_INDEX;394 uint64_t isGroupKey = (isGroup_ != 0) ? TPL_GROUP_INDEX : TPL_NO_GROUP_INDEX;
386- OP_LOGI(context_->GetNodeName(), "isInterleavedKey = %lu, isGroupKey = %lu", isInterleavedKey, isGroupKey);395+ uint64_t clampMode = (clampMode_ == 0) ? TPL_NO_CLAMP_MODE : TPL_CLAMP_MODE;
387- tilingKey_ = GET_TPL_TILING_KEY(isInterleavedKey, isGroupKey);396+ OP_LOGI(context_->GetNodeName(), "isInterleavedKey = %lu, isGroupKey = %lu, clampMode = %lu", isInterleavedKey,
397+ isGroupKey, clampMode);
398+ tilingKey_ = GET_TPL_TILING_KEY(isInterleavedKey, isGroupKey, clampMode);
388}399}
389 400 
390void ClippedSwigluArch35Tiling::FillTilingData()401void ClippedSwigluArch35Tiling::FillTilingData()
@@ -418,6 +429,7 @@ void ClippedSwigluArch35Tiling::PrintTilingInfo()
418 info << ", bUbFactor: " << bUbFactor_;429 info << ", bUbFactor: " << bUbFactor_;
419 info << ", groupNum: " << groupNum_;430 info << ", groupNum: " << groupNum_;
420 info << ", realCoreNum: " << realCoreNum_;431 info << ", realCoreNum: " << realCoreNum_;
432+ info << ", clampMode: " << clampMode_;
421 OP_LOGI(context_->GetNodeName(), "%s", info.str().c_str());433 OP_LOGI(context_->GetNodeName(), "%s", info.str().c_str());
422}434}
423 435 
Mactivation/clipped_swiglu/op_host/config/ascend910_93/clipped_swiglu_binary.json+16-1
@@ -62,6 +62,11 @@
62 "name": "interleaved",62 "name": "interleaved",
63 "dtype": "bool",63 "dtype": "bool",
64 "value": true64 "value": true
65+ },
66+ {
67+ "name": "clamp_mode",
68+ "dtype": "int64",
69+ "value": 0
65 }70 }
66 ]71 ]
67 },72 },
@@ -126,6 +131,11 @@
126 "name": "interleaved",131 "name": "interleaved",
127 "dtype": "bool",132 "dtype": "bool",
128 "value": true133 "value": true
134+ },
135+ {
136+ "name": "clamp_mode",
137+ "dtype": "int64",
138+ "value": 0
129 }139 }
130 ]140 ]
131 },141 },
@@ -190,8 +200,13 @@
190 "name": "interleaved",200 "name": "interleaved",
191 "dtype": "bool",201 "dtype": "bool",
192 "value": true202 "value": true
203+ },
204+ {
205+ "name": "clamp_mode",
206+ "dtype": "int64",
207+ "value": 0
193 }208 }
194 ]209 ]
195 }210 }
196 ]211 ]
197-}212+}
Mactivation/clipped_swiglu/op_host/config/ascend910b/clipped_swiglu_binary.json+16-1
@@ -62,6 +62,11 @@
62 "name": "interleaved",62 "name": "interleaved",
63 "dtype": "bool",63 "dtype": "bool",
64 "value": true64 "value": true
65+ },
66+ {
67+ "name": "clamp_mode",
68+ "dtype": "int64",
69+ "value": 0
65 }70 }
66 ]71 ]
67 },72 },
@@ -126,6 +131,11 @@
126 "name": "interleaved",131 "name": "interleaved",
127 "dtype": "bool",132 "dtype": "bool",
128 "value": true133 "value": true
134+ },
135+ {
136+ "name": "clamp_mode",
137+ "dtype": "int64",
138+ "value": 0
129 }139 }
130 ]140 ]
131 },141 },
@@ -190,8 +200,13 @@
190 "name": "interleaved",200 "name": "interleaved",
191 "dtype": "bool",201 "dtype": "bool",
192 "value": true202 "value": true
203+ },
204+ {
205+ "name": "clamp_mode",
206+ "dtype": "int64",
207+ "value": 0
193 }208 }
194 ]209 ]
195 }210 }
196 ]211 ]
197-}212+}
Mactivation/clipped_swiglu/op_host/config/ascend950/clipped_swiglu_binary.json+15-0
@@ -65,6 +65,11 @@
65 "name": "interleaved",65 "name": "interleaved",
66 "dtype": "bool",66 "dtype": "bool",
67 "value": null67 "value": null
68+ },
69+ {
70+ "name": "clamp_mode",
71+ "dtype": "int64",
72+ "value": null
68 }73 }
69 ]74 ]
70 },75 },
@@ -132,6 +137,11 @@
132 "name": "interleaved",137 "name": "interleaved",
133 "dtype": "bool",138 "dtype": "bool",
134 "value": null139 "value": null
140+ },
141+ {
142+ "name": "clamp_mode",
143+ "dtype": "int64",
144+ "value": null
135 }145 }
136 ]146 ]
137 },147 },
@@ -199,6 +209,11 @@
199 "name": "interleaved",209 "name": "interleaved",
200 "dtype": "bool",210 "dtype": "bool",
201 "value": null211 "value": null
212+ },
213+ {
214+ "name": "clamp_mode",
215+ "dtype": "int64",
216+ "value": null
202 }217 }
203 ]218 ]
204 }219 }
Mactivation/clipped_swiglu/op_kernel/arch35/clipped_swiglu_kernel.h+66-37
@@ -75,7 +75,7 @@ __aicore__ inline void ReduceAllVf(LocalTensor<int64_t>& reduceSumUb, LocalTenso
75 }75 }
76}76}
77 77 
78-template <typename T, bool isInterleaved, bool isGroup>78+template <typename T, bool isInterleaved, bool isGroup, uint64_t clampMode>
79class ClippedSwigluKernel {79class ClippedSwigluKernel {
80public:80public:
81 __aicore__ inline ClippedSwigluKernel(const ClippedSwigluArch35TilingData* tilingData, TPipe* pipe)81 __aicore__ inline ClippedSwigluKernel(const ClippedSwigluArch35TilingData* tilingData, TPipe* pipe)
@@ -128,8 +128,9 @@ private:
128 float bias_ = 0.0f;128 float bias_ = 0.0f;
129};129};
130 130 
131-template <typename T, bool isInterleaved, bool isGroup>131+template <typename T, bool isInterleaved, bool isGroup, uint64_t clampMode>
132-__aicore__ inline void ClippedSwigluKernel<T, isInterleaved, isGroup>::Init(GM_ADDR x, GM_ADDR groupIndex, GM_ADDR y)132+__aicore__ inline void ClippedSwigluKernel<T, isInterleaved, isGroup, clampMode>::Init(GM_ADDR x, GM_ADDR groupIndex,
133+ GM_ADDR y)
133{134{
134 blockIdx_ = GetBlockIdx();135 blockIdx_ = GetBlockIdx();
135 realCoreNum_ = tiling_->realCoreNum;136 realCoreNum_ = tiling_->realCoreNum;
@@ -155,8 +156,8 @@ __aicore__ inline void ClippedSwigluKernel<T, isInterleaved, isGroup>::Init(GM_A
155 }156 }
156}157}
157 158 
158-template <typename T, bool isInterleaved, bool isGroup>159+template <typename T, bool isInterleaved, bool isGroup, uint64_t clampMode>
159-__aicore__ inline void ClippedSwigluKernel<T, isInterleaved, isGroup>::ComputeTiling()160+__aicore__ inline void ClippedSwigluKernel<T, isInterleaved, isGroup, clampMode>::ComputeTiling()
160{161{
161 int64_t dimBatchSize = tiling_->dimBatchSize;162 int64_t dimBatchSize = tiling_->dimBatchSize;
162 if constexpr (isGroup) {163 if constexpr (isGroup) {
@@ -212,8 +213,8 @@ __aicore__ inline void ClippedSwigluKernel<T, isInterleaved, isGroup>::ComputeTi
212 }213 }
213}214}
214 215 
215-template <typename T, bool isInterleaved, bool isGroup>216+template <typename T, bool isInterleaved, bool isGroup, uint64_t clampMode>
216-__aicore__ inline void ClippedSwigluKernel<T, isInterleaved, isGroup>::Process()217+__aicore__ inline void ClippedSwigluKernel<T, isInterleaved, isGroup, clampMode>::Process()
217{218{
218 ComputeTiling();219 ComputeTiling();
219 if (blockIdx_ >= realCoreNum_) {220 if (blockIdx_ >= realCoreNum_) {
@@ -269,9 +270,10 @@ __aicore__ inline void ClippedSwigluKernel<T, isInterleaved, isGroup>::Process()
269 }270 }
270}271}
271 272 
272-template <typename T, bool isInterleaved, bool isGroup>273+template <typename T, bool isInterleaved, bool isGroup, uint64_t clampMode>
273-__aicore__ inline void ClippedSwigluKernel<T, isInterleaved, isGroup>::CopyIn(int64_t gmOffset, int64_t count,274+__aicore__ inline void ClippedSwigluKernel<T, isInterleaved, isGroup, clampMode>::CopyIn(int64_t gmOffset,
274- int64_t blockLen)275+ int64_t count,
276+ int64_t blockLen)
275{277{
276 LocalTensor<T> xDTypeUb = inQueX_.AllocTensor<T>();278 LocalTensor<T> xDTypeUb = inQueX_.AllocTensor<T>();
277 DataCopyPadExtParams<T> padParams = {false, 0, 0, 0};279 DataCopyPadExtParams<T> padParams = {false, 0, 0, 0};
@@ -293,10 +295,10 @@ __aicore__ inline void ClippedSwigluKernel<T, isInterleaved, isGroup>::CopyIn(in
293 inQueX_.EnQue(xDTypeUb);295 inQueX_.EnQue(xDTypeUb);
294}296}
295 297 
296-template <typename T, bool isInterleaved, bool isGroup>298+template <typename T, bool isInterleaved, bool isGroup, uint64_t clampMode>
297-__aicore__ inline void ClippedSwigluKernel<T, isInterleaved, isGroup>::UbStoreAlign(__ubuf__ T* inAddr,299+__aicore__ inline void ClippedSwigluKernel<T, isInterleaved, isGroup, clampMode>::UbStoreAlign(__ubuf__ T* inAddr,
298- __ubuf__ T* outAddr,300+ __ubuf__ T* outAddr,
299- int64_t onceNum)301+ int64_t onceNum)
300{302{
301 uint32_t size = onceNum;303 uint32_t size = onceNum;
302 uint32_t vfLen = vfLenT_;304 uint32_t vfLen = vfLenT_;
@@ -314,8 +316,8 @@ __aicore__ inline void ClippedSwigluKernel<T, isInterleaved, isGroup>::UbStoreAl
314 }316 }
315}317}
316 318 
317-template <typename T, bool isInterleaved, bool isGroup>319+template <typename T, bool isInterleaved, bool isGroup, uint64_t clampMode>
318-__aicore__ inline void ClippedSwigluKernel<T, isInterleaved, isGroup>::ComputeVfSwiglu(320+__aicore__ inline void ClippedSwigluKernel<T, isInterleaved, isGroup, clampMode>::ComputeVfSwiglu(
319 __ubuf__ T* x1UbAddr, __ubuf__ T* x2UbAddr, __ubuf__ T* swigluUbAddr, int64_t dim0OnceSize, int64_t dim1OnceSize,321 __ubuf__ T* x1UbAddr, __ubuf__ T* x2UbAddr, __ubuf__ T* swigluUbAddr, int64_t dim0OnceSize, int64_t dim1OnceSize,
320 int64_t alignDim1In)322 int64_t alignDim1In)
321{323{
@@ -325,6 +327,7 @@ __aicore__ inline void ClippedSwigluKernel<T, isInterleaved, isGroup>::ComputeVf
325 float gluBias = bias_;327 float gluBias = bias_;
326 uint16_t dim0VfTimes = 1;328 uint16_t dim0VfTimes = 1;
327 float scalarOne = 1.0f;329 float scalarOne = 1.0f;
330+ float negScalarOne = -1.0f;
328 uint32_t vfLen = VF_LEN_FP32 * DIM_HALVE;331 uint32_t vfLen = VF_LEN_FP32 * DIM_HALVE;
329 if constexpr (!isInterleaved) {332 if constexpr (!isInterleaved) {
330 dim0VfTimes = dim0OnceSize;333 dim0VfTimes = dim0OnceSize;
@@ -361,6 +364,7 @@ __aicore__ inline void ClippedSwigluKernel<T, isInterleaved, isGroup>::ComputeVf
361 AscendC::MicroAPI::RegTensor<float> addsReg;364 AscendC::MicroAPI::RegTensor<float> addsReg;
362 AscendC::MicroAPI::RegTensor<float> sigmoidReg;365 AscendC::MicroAPI::RegTensor<float> sigmoidReg;
363 AscendC::MicroAPI::RegTensor<float> outFReg;366 AscendC::MicroAPI::RegTensor<float> outFReg;
367+ AscendC::MicroAPI::RegTensor<float> vregX1NegDeF;
364 AscendC::MicroAPI::RegTensor<T> outTReg;368 AscendC::MicroAPI::RegTensor<T> outTReg;
365 MicroAPI::MaskReg mask = MicroAPI::CreateMask<float, MicroAPI::MaskPattern::ALL>();369 MicroAPI::MaskReg mask = MicroAPI::CreateMask<float, MicroAPI::MaskPattern::ALL>();
366 MicroAPI::MaskReg maskT = MicroAPI::UpdateMask<float>(tail);370 MicroAPI::MaskReg maskT = MicroAPI::UpdateMask<float>(tail);
@@ -396,17 +400,29 @@ __aicore__ inline void ClippedSwigluKernel<T, isInterleaved, isGroup>::ComputeVf
396 AscendC::MicroAPI::LoadAlign((MicroAPI::RegTensor<T>&)vregX2DeF, x2UbAddr, srcIdxOffset);400 AscendC::MicroAPI::LoadAlign((MicroAPI::RegTensor<T>&)vregX2DeF, x2UbAddr, srcIdxOffset);
397 }401 }
398 }402 }
399- AscendC::MicroAPI::Mins(minsReg, vregX1DeF, clampLimit, mask);403+ if constexpr (clampMode == 0) {
400- AscendC::MicroAPI::Muls(mulsReg, minsReg, negAlpha, mask);404+ AscendC::MicroAPI::Mins(minsReg, vregX1DeF, clampLimit, mask);
401- AscendC::MicroAPI::Exp(expReg, mulsReg, mask);405+ AscendC::MicroAPI::Muls(mulsReg, minsReg, negAlpha, mask);
402- AscendC::MicroAPI::Adds(addsReg, expReg, scalarOne, mask);406+ AscendC::MicroAPI::Exp(expReg, mulsReg, mask);
403- AscendC::MicroAPI::Div(sigmoidReg, minsReg, addsReg, mask);407+ AscendC::MicroAPI::Adds(addsReg, expReg, scalarOne, mask);
408+ AscendC::MicroAPI::Div(sigmoidReg, minsReg, addsReg, mask);
404 409 
405- AscendC::MicroAPI::Mins(vregX2DeF, vregX2DeF, clampLimit, mask);410+ AscendC::MicroAPI::Mins(vregX2DeF, vregX2DeF, clampLimit, mask);
406- AscendC::MicroAPI::Maxs(vregX2DeF, vregX2DeF, negClampLimit, mask);411+ AscendC::MicroAPI::Maxs(vregX2DeF, vregX2DeF, negClampLimit, mask);
407- AscendC::MicroAPI::Adds(vregX2DeF, vregX2DeF, gluBias, mask);412+ AscendC::MicroAPI::Adds(vregX2DeF, vregX2DeF, gluBias, mask);
408 413 
409- AscendC::MicroAPI::Mul(outFReg, sigmoidReg, vregX2DeF, mask);414+ AscendC::MicroAPI::Mul(outFReg, sigmoidReg, vregX2DeF, mask);
415+ } else {
416+ AscendC::MicroAPI::Muls(vregX1NegDeF, vregX1DeF, negScalarOne, mask);
417+ AscendC::MicroAPI::Exp(expReg, vregX1NegDeF, mask);
418+ AscendC::MicroAPI::Adds(addsReg, expReg, scalarOne, mask);
419+ AscendC::MicroAPI::Div(sigmoidReg, vregX1DeF, addsReg, mask);
420+ AscendC::MicroAPI::Mins(minsReg, sigmoidReg, clampLimit, mask);
421+ 
422+ AscendC::MicroAPI::Mins(vregX2DeF, vregX2DeF, clampLimit, mask);
423+ AscendC::MicroAPI::Maxs(vregX2DeF, vregX2DeF, negClampLimit, mask);
424+ AscendC::MicroAPI::Mul(outFReg, minsReg, vregX2DeF, mask);
425+ }
410 AscendC::MicroAPI::AddrReg outOffset = AscendC::MicroAPI::CreateAddrReg<T>(dim0vfLoopIdx, alignDim1Out,426 AscendC::MicroAPI::AddrReg outOffset = AscendC::MicroAPI::CreateAddrReg<T>(dim0vfLoopIdx, alignDim1Out,
411 dim1vfLoopIdx, VF_LEN_FP32);427 dim1vfLoopIdx, VF_LEN_FP32);
412 if constexpr (sizeof(T) == sizeof(half)) {428 if constexpr (sizeof(T) == sizeof(half)) {
@@ -447,17 +463,29 @@ __aicore__ inline void ClippedSwigluKernel<T, isInterleaved, isGroup>::ComputeVf
447 AscendC::MicroAPI::LoadAlign((MicroAPI::RegTensor<T>&)vregX2DeF, x2UbAddrT, srcIdxOffset1);463 AscendC::MicroAPI::LoadAlign((MicroAPI::RegTensor<T>&)vregX2DeF, x2UbAddrT, srcIdxOffset1);
448 }464 }
449 }465 }
450- AscendC::MicroAPI::Mins(minsReg, vregX1DeF, clampLimit, maskT);466+ if constexpr (clampMode == 0) {
451- AscendC::MicroAPI::Muls(mulsReg, minsReg, negAlpha, maskT);467+ AscendC::MicroAPI::Mins(minsReg, vregX1DeF, clampLimit, maskT);
452- AscendC::MicroAPI::Exp(expReg, mulsReg, maskT);468+ AscendC::MicroAPI::Muls(mulsReg, minsReg, negAlpha, maskT);
453- AscendC::MicroAPI::Adds(addsReg, expReg, scalarOne, maskT);469+ AscendC::MicroAPI::Exp(expReg, mulsReg, maskT);
454- AscendC::MicroAPI::Div(sigmoidReg, minsReg, addsReg, maskT);470+ AscendC::MicroAPI::Adds(addsReg, expReg, scalarOne, maskT);
471+ AscendC::MicroAPI::Div(sigmoidReg, minsReg, addsReg, maskT);
455 472 
456- AscendC::MicroAPI::Mins(vregX2DeF, vregX2DeF, clampLimit, maskT);473+ AscendC::MicroAPI::Mins(vregX2DeF, vregX2DeF, clampLimit, maskT);
457- AscendC::MicroAPI::Maxs(vregX2DeF, vregX2DeF, negClampLimit, maskT);474+ AscendC::MicroAPI::Maxs(vregX2DeF, vregX2DeF, negClampLimit, maskT);
458- AscendC::MicroAPI::Adds(vregX2DeF, vregX2DeF, gluBias, maskT);475+ AscendC::MicroAPI::Adds(vregX2DeF, vregX2DeF, gluBias, maskT);
459 476 
460- AscendC::MicroAPI::Mul(outFReg, sigmoidReg, vregX2DeF, maskT);477+ AscendC::MicroAPI::Mul(outFReg, sigmoidReg, vregX2DeF, maskT);
478+ } else {
479+ AscendC::MicroAPI::Muls(vregX1NegDeF, vregX1DeF, negScalarOne, maskT);
480+ AscendC::MicroAPI::Exp(expReg, vregX1NegDeF, maskT);
481+ AscendC::MicroAPI::Adds(addsReg, expReg, scalarOne, maskT);
482+ AscendC::MicroAPI::Div(sigmoidReg, vregX1DeF, addsReg, maskT);
483+ AscendC::MicroAPI::Mins(minsReg, sigmoidReg, clampLimit, maskT);
484+ 
485+ AscendC::MicroAPI::Mins(vregX2DeF, vregX2DeF, clampLimit, maskT);
486+ AscendC::MicroAPI::Maxs(vregX2DeF, vregX2DeF, negClampLimit, maskT);
487+ AscendC::MicroAPI::Mul(outFReg, minsReg, vregX2DeF, maskT);
488+ }
atomgit-bot
atomgit-botatomgit-bot27 天前

🟠 High Priority

在 ComputeVfSwiglu 的尾块（tail section，处理 dim1OnceSize % VF_LEN_FP32 的剩余元素）中，原始代码在所有计算操作（Mins、Muls、Exp、Adds、Div、Mins、Maxs、Adds、Mul）上均使用 maskT（MaskReg maskT = MicroAPI::UpdateMask<float>(tail)），以仅对有效的尾块元素进行计算。

本 diff 在包装 if constexpr (clampMode == 0) / else 分支时，将尾块中两个分支的计算操作都改为了 mask（全量掩码 MaskPattern::ALL）。这导致在尾块中对全部 VF_LEN_FP32 个 lane 进行计算，其中超出 tail 的 lane 可能包含脏数据（未初始化的寄存器值），可能触发：

虽然最终写入使用 maskT 仅写回有效元素（第 491-494 行），但中间计算的 NaN/Inf 可能在 NPU 上触发浮点异常，影响结果正确性。这是对原始行为的回归（原始代码明确使用 maskT），属于正确性缺陷。

触发条件：当 dim1OnceSize 不能被 VF_LEN_FP32 整除且 tail > 0 时触发。

建议：将尾块中所有计算操作的掩码从 mask 改为 maskT，与原始代码保持一致。主循环中保持 mask 不变（因为主循环处理完整 VF 块，mask 是正确的）。

改动建议
488
+ if constexpr (clampMode == 0) {
489
+ AscendC::MicroAPI::Mins(minsReg, vregX1DeF, clampLimit, maskT);
490
+ AscendC::MicroAPI::Muls(mulsReg, minsReg, negAlpha, maskT);
491
+ AscendC::MicroAPI::Exp(expReg, mulsReg, maskT);
492
+ AscendC::MicroAPI::Adds(addsReg, expReg, scalarOne, maskT);
493
+ AscendC::MicroAPI::Div(sigmoidReg, minsReg, addsReg, maskT);
494
+
495
+ AscendC::MicroAPI::Mins(vregX2DeF, vregX2DeF, clampLimit, maskT);
496
+ AscendC::MicroAPI::Maxs(vregX2DeF, vregX2DeF, negClampLimit, maskT);
497
+ AscendC::MicroAPI::Adds(vregX2DeF, vregX2DeF, gluBias, maskT);
498
+
499
+ AscendC::MicroAPI::Mul(outFReg, sigmoidReg, vregX2DeF, maskT);
500
+ } else {
501
+ AscendC::MicroAPI::Muls(vregX1NegDeF, vregX1DeF, negScalarOne, maskT);
502
+ AscendC::MicroAPI::Exp(expReg, vregX1NegDeF, maskT);
503
+ AscendC::MicroAPI::Adds(addsReg, expReg, scalarOne, maskT);
504
+ AscendC::MicroAPI::Div(sigmoidReg, vregX1DeF, addsReg, maskT);
505
+ AscendC::MicroAPI::Mins(minsReg, sigmoidReg, clampLimit, maskT);
506
+
507
+ AscendC::MicroAPI::Mins(vregX2DeF, vregX2DeF, clampLimit, maskT);
508
+ AscendC::MicroAPI::Maxs(vregX2DeF, vregX2DeF, negClampLimit, maskT);
509
+ AscendC::MicroAPI::Mul(outFReg, minsReg, vregX2DeF, maskT);
488
510
  }
应用建议
likedislike
不准确？
clwsy
clwsy
22 天前 评论：
461 if constexpr (sizeof(T) == sizeof(half)) {489 if constexpr (sizeof(T) == sizeof(half)) {
462 AscendC::MicroAPI::Cast<T, float, CAST_FP32_TO_FP16_BF16>(outTReg, outFReg, maskT);490 AscendC::MicroAPI::Cast<T, float, CAST_FP32_TO_FP16_BF16>(outTReg, outFReg, maskT);
463 StoreAlign<T, AscendC::MicroAPI::StoreDist::DIST_PACK_B32>(swigluUbAddrT, outTReg, outOffset1,491 StoreAlign<T, AscendC::MicroAPI::StoreDist::DIST_PACK_B32>(swigluUbAddrT, outTReg, outOffset1,
@@ -470,9 +498,10 @@ __aicore__ inline void ClippedSwigluKernel<T, isInterleaved, isGroup>::ComputeVf
470 }498 }
471}499}
472 500 
473-template <typename T, bool isInterleaved, bool isGroup>501+template <typename T, bool isInterleaved, bool isGroup, uint64_t clampMode>
474-__aicore__ inline void ClippedSwigluKernel<T, isInterleaved, isGroup>::CopyOut(int64_t gmOffset, int64_t count,502+__aicore__ inline void ClippedSwigluKernel<T, isInterleaved, isGroup, clampMode>::CopyOut(int64_t gmOffset,
475- int64_t blockLen)503+ int64_t count,
504+ int64_t blockLen)
476{505{
477 LocalTensor<T> outputUb = outQueY_.DeQue<T>();506 LocalTensor<T> outputUb = outQueY_.DeQue<T>();
478 outQueY_.EnQue(outputUb);507 outQueY_.EnQue(outputUb);
Mactivation/clipped_swiglu/op_kernel/arch35/clipped_swiglu_tiling_key.h+7-2
@@ -24,15 +24,20 @@
24#define TPL_NO_GROUP_INDEX 024#define TPL_NO_GROUP_INDEX 0
25#define TPL_GROUP_INDEX 125#define TPL_GROUP_INDEX 1
26 26 
27+#define TPL_NO_CLAMP_MODE 0
28+#define TPL_CLAMP_MODE 1
29+ 
27namespace ClippedSwigluOp {30namespace ClippedSwigluOp {
28ASCENDC_TPL_ARGS_DECL(ClippedSwiglu,31ASCENDC_TPL_ARGS_DECL(ClippedSwiglu,
29 ASCENDC_TPL_UINT_DECL(isInterleaved, 1, ASCENDC_TPL_UI_LIST, TPL_INTERLEAVED_FALSE,32 ASCENDC_TPL_UINT_DECL(isInterleaved, 1, ASCENDC_TPL_UI_LIST, TPL_INTERLEAVED_FALSE,
30 TPL_INTERLEAVED_TRUE),33 TPL_INTERLEAVED_TRUE),
31- ASCENDC_TPL_UINT_DECL(isGroup, 1, ASCENDC_TPL_UI_LIST, TPL_NO_GROUP_INDEX, TPL_GROUP_INDEX));34+ ASCENDC_TPL_UINT_DECL(isGroup, 1, ASCENDC_TPL_UI_LIST, TPL_NO_GROUP_INDEX, TPL_GROUP_INDEX),
35+ ASCENDC_TPL_UINT_DECL(clampMode, 1, ASCENDC_TPL_UI_LIST, TPL_NO_CLAMP_MODE, TPL_CLAMP_MODE));
32 36 
33ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(37ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(
34 ASCENDC_TPL_UINT_SEL(isInterleaved, ASCENDC_TPL_UI_LIST, TPL_INTERLEAVED_FALSE, TPL_INTERLEAVED_TRUE),38 ASCENDC_TPL_UINT_SEL(isInterleaved, ASCENDC_TPL_UI_LIST, TPL_INTERLEAVED_FALSE, TPL_INTERLEAVED_TRUE),
35- ASCENDC_TPL_UINT_SEL(isGroup, ASCENDC_TPL_UI_LIST, TPL_NO_GROUP_INDEX, TPL_GROUP_INDEX, TPL_FLOOR)));39+ ASCENDC_TPL_UINT_SEL(isGroup, ASCENDC_TPL_UI_LIST, TPL_NO_GROUP_INDEX, TPL_GROUP_INDEX, TPL_FLOOR),
40+ ASCENDC_TPL_UINT_SEL(clampMode, ASCENDC_TPL_UI_LIST, TPL_NO_CLAMP_MODE, TPL_CLAMP_MODE)));
36} // namespace ClippedSwigluOp41} // namespace ClippedSwigluOp
37 42 
38#endif // CLIPPED_SWIGLU_TILING_KEY_H43#endif // CLIPPED_SWIGLU_TILING_KEY_H
Mactivation/clipped_swiglu/op_kernel/clipped_swiglu_apt.cpp+5-5
@@ -22,7 +22,7 @@
22using namespace AscendC;22using namespace AscendC;
23using namespace ClippedSwigluOp;23using namespace ClippedSwigluOp;
24 24 
25-template <uint64_t isInterleaved, uint64_t isGroup>25+template <uint64_t isInterleaved, uint64_t isGroup, uint64_t clampMode>
26__global__ __aicore__ void clipped_swiglu(GM_ADDR x, GM_ADDR groupIndex, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)26__global__ __aicore__ void clipped_swiglu(GM_ADDR x, GM_ADDR groupIndex, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
27{27{
28 KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);28 KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
@@ -33,21 +33,21 @@ __global__ __aicore__ void clipped_swiglu(GM_ADDR x, GM_ADDR groupIndex, GM_ADDR
33 33 
34 if constexpr (isInterleaved == 1) {34 if constexpr (isInterleaved == 1) {
35 if constexpr (isGroup == 1) {35 if constexpr (isGroup == 1) {
36- ClippedSwigluKernel<DTYPE_X, true, true> op(&tilingData, &pipe);36+ ClippedSwigluKernel<DTYPE_X, true, true, clampMode> op(&tilingData, &pipe);
37 op.Init(x, groupIndex, y);37 op.Init(x, groupIndex, y);
38 op.Process();38 op.Process();
39 } else {39 } else {
40- ClippedSwigluKernel<DTYPE_X, true, false> op(&tilingData, &pipe);40+ ClippedSwigluKernel<DTYPE_X, true, false, clampMode> op(&tilingData, &pipe);
41 op.Init(x, groupIndex, y);41 op.Init(x, groupIndex, y);
42 op.Process();42 op.Process();
43 }43 }
44 } else {44 } else {
45 if constexpr (isGroup == 1) {45 if constexpr (isGroup == 1) {
46- ClippedSwigluKernel<DTYPE_X, false, true> op(&tilingData, &pipe);46+ ClippedSwigluKernel<DTYPE_X, false, true, clampMode> op(&tilingData, &pipe);
47 op.Init(x, groupIndex, y);47 op.Init(x, groupIndex, y);
48 op.Process();48 op.Process();
49 } else {49 } else {
50- ClippedSwigluKernel<DTYPE_X, false, false> op(&tilingData, &pipe);50+ ClippedSwigluKernel<DTYPE_X, false, false, clampMode> op(&tilingData, &pipe);
51 op.Init(x, groupIndex, y);51 op.Init(x, groupIndex, y);
52 op.Process();52 op.Process();
53 }53 }
Mactivation/clipped_swiglu/tests/ut/op_kernel/clipped_swiglu_data/gen_data.py+6-6
@@ -51,12 +51,12 @@ def do_clippedSwiglu(x, group_index, dim, alpha, limit, bias, interleaved):
51 51 
52 52 
53params_info = {53params_info = {
54- "test_case_bf16_shortH": {"x_shape": [3200, 5760], "x_dtype": torch.bfloat16},54+ "test_case_bf16_shortH": {"x_shape": [40, 5760], "x_dtype": torch.bfloat16},
55- "test_case_fp16_shortH": {"x_shape": [3200, 5760], "x_dtype": torch.float16},55+ "test_case_fp16_shortH": {"x_shape": [40, 5760], "x_dtype": torch.float16},
56- "test_case_fp32_shortH": {"x_shape": [3200, 5760], "x_dtype": torch.float32},56+ "test_case_fp32_shortH": {"x_shape": [40, 5760], "x_dtype": torch.float32},
57- "test_case_bf16_longH": {"x_shape": [3200, 23040], "x_dtype": torch.bfloat16},57+ "test_case_bf16_longH": {"x_shape": [40, 23040], "x_dtype": torch.bfloat16},
58- "test_case_fp16_longH": {"x_shape": [3200, 23040], "x_dtype": torch.float16},58+ "test_case_fp16_longH": {"x_shape": [40, 23040], "x_dtype": torch.float16},
59- "test_case_fp32_longH": {"x_shape": [3200, 23040], "x_dtype": torch.float32},59+ "test_case_fp32_longH": {"x_shape": [40, 23040], "x_dtype": torch.float32},
60}60}
61 61 
62 62 
Mactivation/clipped_swiglu/tests/ut/op_kernel/clipped_swiglu_data/gen_tiling.py+47-10
@@ -2,16 +2,17 @@
2# -*- coding: utf-8 -*-2# -*- coding: utf-8 -*-
3# ----------------------------------------------------------------------------3# ----------------------------------------------------------------------------
4# Copyright (c) 2025 Huawei Technologies Co., Ltd.4# Copyright (c) 2025 Huawei Technologies Co., Ltd.
5-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of 5+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
6# CANN Open Software License Agreement Version 2.0 (the "License").6# CANN Open Software License Agreement Version 2.0 (the "License").
7# Please refer to the License for details. You may not use this file except in compliance with the License.7# Please refer to the License for details. You may not use this file except in compliance with the License.
8-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, 8+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
9# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.9# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
10# See LICENSE in the root of the software repository for the full text of the License.10# See LICENSE in the root of the software repository for the full text of the License.
11# ----------------------------------------------------------------------------11# ----------------------------------------------------------------------------
12 12 
13import numpy as np13import numpy as np
14import sys14import sys
15+ 
15"""16"""
16BEGIN_TILING_DATA_DEF(ClippedSwigluTilingData)17BEGIN_TILING_DATA_DEF(ClippedSwigluTilingData)
17TILING_DATA_FIELD_DEF(int64_t, coreNumAll);18TILING_DATA_FIELD_DEF(int64_t, coreNumAll);
@@ -28,14 +29,51 @@ TILING_DATA_FIELD_DEF(int64_t, groupNum);
28"""29"""
29 30 
30params_info = {31params_info = {
31- "case_half_ungrouped_shortH": [40, 3200, 5760, 0, 0, 0, 1.702, 7.0, 1.0, 5424, 0],32+ "case_half_ungrouped_shortH": [40, 40, 5760, 0, 0, 0, 1.702, 7.0, 1.0, 5424, 0],
32- "case_interleaved_ungrouped_shortH": [40, 3200, 5760, 0, 0, 1, 1.702, 7.0, 1.0, 5424, 0],33+ "case_interleaved_ungrouped_shortH": [
33- "case_half_grouped_shortH": [40, 3200, 5760, 0, 1, 0, 1.702, 7.0, 1.0, 5424, 4],34+ 40,
34- "case_interleaved_grouped_shortH": [40, 3200, 5760, 0, 1, 1, 1.702, 7.0, 1.0, 5424, 4],35+ 40,
35- "case_half_grouped_longH": [40, 3200, 23040, 1, 1, 0, 1.702, 7.0, 1.0, 5424, 4],36+ 5760,
36- "case_interleaved_grouped_longH": [40, 3200, 23040, 1, 1, 1, 1.702, 7.0, 1.0, 5424, 4]37+ 0,
38+ 0,
39+ 1,
40+ 1.702,
41+ 7.0,
42+ 1.0,
43+ 5424,
44+ 0,
45+ ],
46+ "case_half_grouped_shortH": [40, 40, 5760, 0, 1, 0, 1.702, 7.0, 1.0, 5424, 4],
47+ "case_interleaved_grouped_shortH": [
48+ 40,
49+ 40,
50+ 5760,
51+ 0,
52+ 1,
53+ 1,
54+ 1.702,
55+ 7.0,
56+ 1.0,
57+ 5424,
58+ 4,
59+ ],
60+ "case_half_grouped_longH": [40, 40, 23040, 1, 1, 0, 1.702, 7.0, 1.0, 5424, 4],
61+ "case_interleaved_grouped_longH": [
62+ 40,
63+ 40,
64+ 23040,
65+ 1,
66+ 1,
67+ 1,
68+ 1.702,
69+ 7.0,
70+ 1.0,
71+ 5424,
72+ 4,
73+ ],
37}74}
38 75 
76+ 
39def main():77def main():
40 params_list = params_info[sys.argv[1]]78 params_list = params_info[sys.argv[1]]
41 79 
@@ -48,6 +86,5 @@ def main():
48 r3.tofile(tiling_file)86 r3.tofile(tiling_file)
49 87 
50 88 
51- 89+if __name__ == "__main__":
52-if __name__ == '__main__':
53 main()90 main()
Mactivation/clipped_swiglu/tests/ut/op_kernel/test_clipped_swiglu.cpp+37-37
@@ -32,8 +32,8 @@ protected:
32 32 
33TEST_F(clipped_swiglu_test, test_case_bf16_half_ungrouped_shortH)33TEST_F(clipped_swiglu_test, test_case_bf16_half_ungrouped_shortH)
34{34{
35- size_t inputByteSize = 3200 * 2880 * 2 * sizeof(half);35+ size_t inputByteSize = 40 * 2880 * 2 * sizeof(half);
36- size_t outputByteSize = 3200 * 2880 * sizeof(half);36+ size_t outputByteSize = 40 * 2880 * sizeof(half);
37 size_t groupIndexByteSize = 10 * sizeof(int64_t);37 size_t groupIndexByteSize = 10 * sizeof(int64_t);
38 size_t tiling_data_size = sizeof(ClippedSwigluTilingData);38 size_t tiling_data_size = sizeof(ClippedSwigluTilingData);
39 39 
@@ -76,8 +76,8 @@ TEST_F(clipped_swiglu_test, test_case_bf16_half_ungrouped_shortH)
76 76 
77TEST_F(clipped_swiglu_test, test_case_fp16_half_ungrouped_shortH)77TEST_F(clipped_swiglu_test, test_case_fp16_half_ungrouped_shortH)
78{78{
79- size_t inputByteSize = 3200 * 2880 * 2 * sizeof(half);79+ size_t inputByteSize = 40 * 2880 * 2 * sizeof(half);
80- size_t outputByteSize = 3200 * 2880 * sizeof(half);80+ size_t outputByteSize = 40 * 2880 * sizeof(half);
81 size_t groupIndexByteSize = 10 * sizeof(int64_t);81 size_t groupIndexByteSize = 10 * sizeof(int64_t);
82 size_t tiling_data_size = sizeof(ClippedSwigluTilingData);82 size_t tiling_data_size = sizeof(ClippedSwigluTilingData);
83 83 
@@ -120,8 +120,8 @@ TEST_F(clipped_swiglu_test, test_case_fp16_half_ungrouped_shortH)
120 120 
121TEST_F(clipped_swiglu_test, test_case_fp32_half_ungrouped_shortH)121TEST_F(clipped_swiglu_test, test_case_fp32_half_ungrouped_shortH)
122{122{
123- size_t inputByteSize = 3200 * 2880 * 2 * sizeof(float);123+ size_t inputByteSize = 40 * 2880 * 2 * sizeof(float);
124- size_t outputByteSize = 3200 * 2880 * sizeof(float);124+ size_t outputByteSize = 40 * 2880 * sizeof(float);
125 size_t groupIndexByteSize = 10 * sizeof(int64_t);125 size_t groupIndexByteSize = 10 * sizeof(int64_t);
126 size_t tiling_data_size = sizeof(ClippedSwigluTilingData);126 size_t tiling_data_size = sizeof(ClippedSwigluTilingData);
127 127 
@@ -164,8 +164,8 @@ TEST_F(clipped_swiglu_test, test_case_fp32_half_ungrouped_shortH)
164 164 
165TEST_F(clipped_swiglu_test, test_case_bf16_interleaved_ungrouped_shortH)165TEST_F(clipped_swiglu_test, test_case_bf16_interleaved_ungrouped_shortH)
166{166{
167- size_t inputByteSize = 3200 * 2880 * 2 * sizeof(half);167+ size_t inputByteSize = 40 * 2880 * 2 * sizeof(half);
168- size_t outputByteSize = 3200 * 2880 * sizeof(half);168+ size_t outputByteSize = 40 * 2880 * sizeof(half);
169 size_t groupIndexByteSize = 10 * sizeof(int64_t);169 size_t groupIndexByteSize = 10 * sizeof(int64_t);
170 size_t tiling_data_size = sizeof(ClippedSwigluTilingData);170 size_t tiling_data_size = sizeof(ClippedSwigluTilingData);
171 171 
@@ -208,8 +208,8 @@ TEST_F(clipped_swiglu_test, test_case_bf16_interleaved_ungrouped_shortH)
208 208 
209TEST_F(clipped_swiglu_test, test_case_fp16_interleaved_ungrouped_shortH)209TEST_F(clipped_swiglu_test, test_case_fp16_interleaved_ungrouped_shortH)
210{210{
211- size_t inputByteSize = 3200 * 2880 * 2 * sizeof(half);211+ size_t inputByteSize = 40 * 2880 * 2 * sizeof(half);
212- size_t outputByteSize = 3200 * 2880 * sizeof(half);212+ size_t outputByteSize = 40 * 2880 * sizeof(half);
213 size_t groupIndexByteSize = 10 * sizeof(int64_t);213 size_t groupIndexByteSize = 10 * sizeof(int64_t);
214 size_t tiling_data_size = sizeof(ClippedSwigluTilingData);214 size_t tiling_data_size = sizeof(ClippedSwigluTilingData);
215 215 
@@ -252,8 +252,8 @@ TEST_F(clipped_swiglu_test, test_case_fp16_interleaved_ungrouped_shortH)
252 252 
253TEST_F(clipped_swiglu_test, test_case_fp32_interleaved_ungrouped_shortH)253TEST_F(clipped_swiglu_test, test_case_fp32_interleaved_ungrouped_shortH)
254{254{
255- size_t inputByteSize = 3200 * 2880 * 2 * sizeof(float);255+ size_t inputByteSize = 40 * 2880 * 2 * sizeof(float);
256- size_t outputByteSize = 3200 * 2880 * sizeof(float);256+ size_t outputByteSize = 40 * 2880 * sizeof(float);
257 size_t groupIndexByteSize = 10 * sizeof(int64_t);257 size_t groupIndexByteSize = 10 * sizeof(int64_t);
258 size_t tiling_data_size = sizeof(ClippedSwigluTilingData);258 size_t tiling_data_size = sizeof(ClippedSwigluTilingData);
259 259 
@@ -296,8 +296,8 @@ TEST_F(clipped_swiglu_test, test_case_fp32_interleaved_ungrouped_shortH)
296 296 
297TEST_F(clipped_swiglu_test, test_case_bf16_half_grouped_shortH)297TEST_F(clipped_swiglu_test, test_case_bf16_half_grouped_shortH)
298{298{
299- size_t inputByteSize = 3200 * 2880 * 2 * sizeof(half);299+ size_t inputByteSize = 40 * 2880 * 2 * sizeof(half);
300- size_t outputByteSize = 3200 * 2880 * sizeof(half);300+ size_t outputByteSize = 40 * 2880 * sizeof(half);
301 size_t groupIndexByteSize = 10 * sizeof(int64_t);301 size_t groupIndexByteSize = 10 * sizeof(int64_t);
302 size_t tiling_data_size = sizeof(ClippedSwigluTilingData);302 size_t tiling_data_size = sizeof(ClippedSwigluTilingData);
303 303 
@@ -340,8 +340,8 @@ TEST_F(clipped_swiglu_test, test_case_bf16_half_grouped_shortH)
340 340 
341TEST_F(clipped_swiglu_test, test_case_fp16_half_grouped_shortH)341TEST_F(clipped_swiglu_test, test_case_fp16_half_grouped_shortH)
342{342{
343- size_t inputByteSize = 3200 * 2880 * 2 * sizeof(half);343+ size_t inputByteSize = 40 * 2880 * 2 * sizeof(half);
344- size_t outputByteSize = 3200 * 2880 * sizeof(half);344+ size_t outputByteSize = 40 * 2880 * sizeof(half);
345 size_t groupIndexByteSize = 10 * sizeof(int64_t);345 size_t groupIndexByteSize = 10 * sizeof(int64_t);
346 size_t tiling_data_size = sizeof(ClippedSwigluTilingData);346 size_t tiling_data_size = sizeof(ClippedSwigluTilingData);
347 347 
@@ -384,8 +384,8 @@ TEST_F(clipped_swiglu_test, test_case_fp16_half_grouped_shortH)
384 384 
385TEST_F(clipped_swiglu_test, test_case_fp32_half_grouped_shortH)385TEST_F(clipped_swiglu_test, test_case_fp32_half_grouped_shortH)
386{386{
387- size_t inputByteSize = 3200 * 2880 * 2 * sizeof(float);387+ size_t inputByteSize = 40 * 2880 * 2 * sizeof(float);
388- size_t outputByteSize = 3200 * 2880 * sizeof(float);388+ size_t outputByteSize = 40 * 2880 * sizeof(float);
389 size_t groupIndexByteSize = 10 * sizeof(int64_t);389 size_t groupIndexByteSize = 10 * sizeof(int64_t);
390 size_t tiling_data_size = sizeof(ClippedSwigluTilingData);390 size_t tiling_data_size = sizeof(ClippedSwigluTilingData);
391 391 
@@ -428,8 +428,8 @@ TEST_F(clipped_swiglu_test, test_case_fp32_half_grouped_shortH)
428 428 
429TEST_F(clipped_swiglu_test, test_case_bf16_interleaved_grouped_shortH)429TEST_F(clipped_swiglu_test, test_case_bf16_interleaved_grouped_shortH)
430{430{
431- size_t inputByteSize = 3200 * 2880 * 2 * sizeof(half);431+ size_t inputByteSize = 40 * 2880 * 2 * sizeof(half);
432- size_t outputByteSize = 3200 * 2880 * sizeof(half);432+ size_t outputByteSize = 40 * 2880 * sizeof(half);
433 size_t groupIndexByteSize = 10 * sizeof(int64_t);433 size_t groupIndexByteSize = 10 * sizeof(int64_t);
434 size_t tiling_data_size = sizeof(ClippedSwigluTilingData);434 size_t tiling_data_size = sizeof(ClippedSwigluTilingData);
435 435 
@@ -472,8 +472,8 @@ TEST_F(clipped_swiglu_test, test_case_bf16_interleaved_grouped_shortH)
472 472 
473TEST_F(clipped_swiglu_test, test_case_fp16_interleaved_grouped_shortH)473TEST_F(clipped_swiglu_test, test_case_fp16_interleaved_grouped_shortH)
474{474{
475- size_t inputByteSize = 3200 * 2880 * 2 * sizeof(half);475+ size_t inputByteSize = 40 * 2880 * 2 * sizeof(half);
476- size_t outputByteSize = 3200 * 2880 * sizeof(half);476+ size_t outputByteSize = 40 * 2880 * sizeof(half);
477 size_t groupIndexByteSize = 10 * sizeof(int64_t);477 size_t groupIndexByteSize = 10 * sizeof(int64_t);
478 size_t tiling_data_size = sizeof(ClippedSwigluTilingData);478 size_t tiling_data_size = sizeof(ClippedSwigluTilingData);
479 479 
@@ -516,8 +516,8 @@ TEST_F(clipped_swiglu_test, test_case_fp16_interleaved_grouped_shortH)
516 516 
517TEST_F(clipped_swiglu_test, test_case_fp32_interleaved_grouped_shortH)517TEST_F(clipped_swiglu_test, test_case_fp32_interleaved_grouped_shortH)
518{518{
519- size_t inputByteSize = 3200 * 2880 * 2 * sizeof(float);519+ size_t inputByteSize = 40 * 2880 * 2 * sizeof(float);
520- size_t outputByteSize = 3200 * 2880 * sizeof(float);520+ size_t outputByteSize = 40 * 2880 * sizeof(float);
521 size_t groupIndexByteSize = 10 * sizeof(int64_t);521 size_t groupIndexByteSize = 10 * sizeof(int64_t);
522 size_t tiling_data_size = sizeof(ClippedSwigluTilingData);522 size_t tiling_data_size = sizeof(ClippedSwigluTilingData);
523 523 
@@ -560,8 +560,8 @@ TEST_F(clipped_swiglu_test, test_case_fp32_interleaved_grouped_shortH)
560 560 
561TEST_F(clipped_swiglu_test, test_case_bf16_half_grouped_longH)561TEST_F(clipped_swiglu_test, test_case_bf16_half_grouped_longH)
562{562{
563- size_t inputByteSize = 3200 * 11520 * 2 * sizeof(half);563+ size_t inputByteSize = 40 * 11520 * 2 * sizeof(half);
564- size_t outputByteSize = 3200 * 11520 * sizeof(half);564+ size_t outputByteSize = 40 * 11520 * sizeof(half);
565 size_t groupIndexByteSize = 10 * sizeof(int64_t);565 size_t groupIndexByteSize = 10 * sizeof(int64_t);
566 size_t tiling_data_size = sizeof(ClippedSwigluTilingData);566 size_t tiling_data_size = sizeof(ClippedSwigluTilingData);
567 567 
@@ -604,8 +604,8 @@ TEST_F(clipped_swiglu_test, test_case_bf16_half_grouped_longH)
604 604 
605TEST_F(clipped_swiglu_test, test_case_fp16_half_grouped_longH)605TEST_F(clipped_swiglu_test, test_case_fp16_half_grouped_longH)
606{606{
607- size_t inputByteSize = 3200 * 11520 * 2 * sizeof(half);607+ size_t inputByteSize = 40 * 11520 * 2 * sizeof(half);
608- size_t outputByteSize = 3200 * 11520 * sizeof(half);608+ size_t outputByteSize = 40 * 11520 * sizeof(half);
609 size_t groupIndexByteSize = 10 * sizeof(int64_t);609 size_t groupIndexByteSize = 10 * sizeof(int64_t);
610 size_t tiling_data_size = sizeof(ClippedSwigluTilingData);610 size_t tiling_data_size = sizeof(ClippedSwigluTilingData);
611 611 
@@ -648,8 +648,8 @@ TEST_F(clipped_swiglu_test, test_case_fp16_half_grouped_longH)
648 648 
649TEST_F(clipped_swiglu_test, test_case_fp32_half_grouped_longH)649TEST_F(clipped_swiglu_test, test_case_fp32_half_grouped_longH)
650{650{
651- size_t inputByteSize = 3200 * 11520 * 2 * sizeof(float);651+ size_t inputByteSize = 40 * 11520 * 2 * sizeof(float);
652- size_t outputByteSize = 3200 * 11520 * sizeof(float);652+ size_t outputByteSize = 40 * 11520 * sizeof(float);
653 size_t groupIndexByteSize = 10 * sizeof(int64_t);653 size_t groupIndexByteSize = 10 * sizeof(int64_t);
654 size_t tiling_data_size = sizeof(ClippedSwigluTilingData);654 size_t tiling_data_size = sizeof(ClippedSwigluTilingData);
655 655 
@@ -692,8 +692,8 @@ TEST_F(clipped_swiglu_test, test_case_fp32_half_grouped_longH)
692 692 
693TEST_F(clipped_swiglu_test, test_case_bf16_interleaved_grouped_longH)693TEST_F(clipped_swiglu_test, test_case_bf16_interleaved_grouped_longH)
694{694{
695- size_t inputByteSize = 3200 * 11520 * 2 * sizeof(half);695+ size_t inputByteSize = 40 * 11520 * 2 * sizeof(half);
696- size_t outputByteSize = 3200 * 11520 * sizeof(half);696+ size_t outputByteSize = 40 * 11520 * sizeof(half);
697 size_t groupIndexByteSize = 10 * sizeof(int64_t);697 size_t groupIndexByteSize = 10 * sizeof(int64_t);
698 size_t tiling_data_size = sizeof(ClippedSwigluTilingData);698 size_t tiling_data_size = sizeof(ClippedSwigluTilingData);
699 699 
@@ -736,8 +736,8 @@ TEST_F(clipped_swiglu_test, test_case_bf16_interleaved_grouped_longH)
736 736 
737TEST_F(clipped_swiglu_test, test_case_fp16_interleaved_grouped_longH)737TEST_F(clipped_swiglu_test, test_case_fp16_interleaved_grouped_longH)
738{738{
739- size_t inputByteSize = 3200 * 11520 * 2 * sizeof(half);739+ size_t inputByteSize = 40 * 11520 * 2 * sizeof(half);
740- size_t outputByteSize = 3200 * 11520 * sizeof(half);740+ size_t outputByteSize = 40 * 11520 * sizeof(half);
741 size_t groupIndexByteSize = 10 * sizeof(int64_t);741 size_t groupIndexByteSize = 10 * sizeof(int64_t);
742 size_t tiling_data_size = sizeof(ClippedSwigluTilingData);742 size_t tiling_data_size = sizeof(ClippedSwigluTilingData);
743 743 
@@ -780,8 +780,8 @@ TEST_F(clipped_swiglu_test, test_case_fp16_interleaved_grouped_longH)
780 780 
781TEST_F(clipped_swiglu_test, test_case_fp32_interleaved_grouped_longH)781TEST_F(clipped_swiglu_test, test_case_fp32_interleaved_grouped_longH)
782{782{
783- size_t inputByteSize = 3200 * 11520 * 2 * sizeof(float);783+ size_t inputByteSize = 40 * 11520 * 2 * sizeof(float);
784- size_t outputByteSize = 3200 * 11520 * sizeof(float);784+ size_t outputByteSize = 40 * 11520 * sizeof(float);
785 size_t groupIndexByteSize = 10 * sizeof(int64_t);785 size_t groupIndexByteSize = 10 * sizeof(int64_t);
786 size_t tiling_data_size = sizeof(ClippedSwigluTilingData);786 size_t tiling_data_size = sizeof(ClippedSwigluTilingData);
787 787 
@@ -820,4 +820,4 @@ TEST_F(clipped_swiglu_test, test_case_fp32_interleaved_grouped_longH)
820 AscendC::GmFree((void*)workspace);820 AscendC::GmFree((void*)workspace);
821 AscendC::GmFree((void*)tiling);821 AscendC::GmFree((void*)tiling);
822 free(path_);822 free(path_);
823-}823+}
Aactivation/clipped_swiglu/torch_extension/__init__.py+14-0
@@ -0,0 +1,14 @@
1+# -----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# -----------------------------------------------------------------------------------------------------------
10+ 
11+from .clipped_swiglu import clipped_swiglu
12+from .graph_convert_clipped_swiglu import convert_clipped_swiglu
13+ 
14+__all__ = ["clipped_swiglu", "convert_clipped_swiglu"]
Aactivation/clipped_swiglu/torch_extension/clipped_swiglu.py+81-0
@@ -0,0 +1,81 @@
1+#!/usr/bin/env python3
2+# -*- coding: UTF-8 -*-
3+# -----------------------------------------------------------------------------------------------------------
4+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
5+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
6+# CANN Open Software License Agreement Version 2.0 (the "License").
7+# Please refer to the License for details. You may not use this file except in compliance with the License.
8+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
9+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
10+# See LICENSE in the root of the software repository for the full text of the License.
11+# -----------------------------------------------------------------------------------------------------------
12+ 
13+import torch
14+from torch.library import impl
15+from cann_ops_nn.op_builder import OpBuilder, get_as_library
16+ 
17+ 
18+class ClippedSwigluOpBuilder(OpBuilder):
19+ def __init__(self):
20+ super().__init__("clipped_swiglu")
21+ 
22+ def sources(self) -> list:
23+ return [self.resolve_source("clipped_swiglu.cpp")]
24+ 
25+ def schema(self) -> str:
26+ return (
27+ "clipped_swiglu("
28+ "Tensor x, *, Tensor? group_index=None, "
29+ "int dim=-1, float alpha=1.702, float limit=7.0, float bias=1.0, bool interleaved=True, int clamp_mode=0"
30+ ") -> Tensor "
31+ )
32+ 
33+ def register_meta(self):
34+ @impl(get_as_library(), self.name, "Meta")
35+ def clipped_swiglu_meta(
36+ x: torch.Tensor,
37+ *,
38+ group_index=None,
39+ dim=-1,
40+ alpha=1.702,
41+ limit=7.0,
42+ bias=1.0,
43+ interleaved=True,
44+ clamp_mode=0,
45+ ):
46+ real_dim = dim if dim >= 0 else dim + x.dim()
47+ if x.dim() == 0:
48+ raise RuntimeError("x must be at least 1-D")
49+ if real_dim < 0 or real_dim >= x.dim():
50+ raise RuntimeError(
51+ f"dim out of range [-{x.dim()}, {x.dim() - 1}], got {dim}"
52+ )
53+ if x.size(real_dim) % 2 != 0:
54+ raise RuntimeError(
55+ f"x size at dim {real_dim} must be even, but got {x.size(real_dim)}"
56+ )
57+ output_size = list(x.shape)
58+ output_size[real_dim] = output_size[real_dim] // 2
59+ return torch.empty(output_size, dtype=x.dtype, device=x.device)
60+ 
61+ 
62+clipped_swiglu_builder = ClippedSwigluOpBuilder()
63+clipped_swiglu_builder._ensure_initialized()
64+ 
65+ 
66+@impl(get_as_library(), clipped_swiglu_builder.name, "PrivateUse1")
67+def clipped_swiglu(
68+ x: torch.Tensor,
69+ *,
70+ group_index=None,
71+ dim=-1,
72+ alpha=1.702,
73+ limit=7.0,
74+ bias=1.0,
75+ interleaved=True,
76+ clamp_mode=0,
77+):
78+ op_module = clipped_swiglu_builder.load()
79+ return op_module.clipped_swiglu(
80+ x, group_index, dim, alpha, limit, bias, interleaved, clamp_mode
81+ )
Aactivation/clipped_swiglu/torch_extension/csrc/clipped_swiglu.cpp+73-0
@@ -0,0 +1,73 @@
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 <torch/extension.h>
12+#include "aclnn_common.h"
13+ 
14+namespace cann_ops_nn {
15+namespace activation {
16+namespace {
17+constexpr int64_t NUM_TWO = 2;
18+ 
19+void CheckNpuTensor(const at::Tensor& tensor, const char* name)
20+{
21+ TORCH_CHECK(tensor.defined(), name, " must be defined");
22+ TORCH_CHECK(torch_npu::utils::is_npu(tensor), name, " must be on NPU device");
23+}
24+ 
25+void CheckOptionalNpuTensor(const c10::optional<at::Tensor>& tensor, const char* name)
26+{
27+ if (tensor.has_value() && tensor.value().defined()) {
28+ CheckNpuTensor(tensor.value(), name);
29+ }
30+}
31+ 
32+bool check_aclnn_kernel_available(std::string aclnn_name)
33+{
34+ std::string workspace_name = aclnn_name + "GetWorkspaceSize";
35+ if (GetOpApiFuncAddr(aclnn_name.c_str()) == nullptr || GetOpApiFuncAddr(workspace_name.c_str()) == nullptr) {
36+ return false;
37+ }
38+ return true;
39+}
40+} // namespace
41+ 
42+at::Tensor clipped_swiglu(const at::Tensor& x, const c10::optional<at::Tensor>& group_index, int64_t dim, double alpha,
43+ double limit, double bias, bool interleaved, int64_t clamp_mode)
44+{
45+ CheckNpuTensor(x, "x");
46+ CheckOptionalNpuTensor(group_index, "group_index");
47+ TORCH_CHECK(clamp_mode == 0 || clamp_mode == 1, "clamp_mode should be 0 or 1, but got ", clamp_mode);
48+ 
49+ if (dim < 0) {
50+ dim += static_cast<int64_t>(x.sizes().size());
51+ }
52+ TORCH_CHECK(dim >= 0 && dim < static_cast<int64_t>(x.sizes().size()), "dim out of range, got ", dim);
53+ TORCH_CHECK(x.size(dim) % NUM_TWO == 0, "x size at dim ", dim, " must be even, but got ", x.size(dim));
54+ auto y_shape = op_infer::array_to_small_vector(x.sizes());
55+ y_shape[dim] /= NUM_TWO;
56+ at::Tensor y = at::empty(y_shape, x.options());
57+ 
58+ static bool npu_support_v2 = check_aclnn_kernel_available("aclnnClippedSwigluV2");
59+ if (npu_support_v2) {
60+ ACLNN_CMD(aclnnClippedSwigluV2, x, group_index, dim, alpha, limit, bias, interleaved, clamp_mode, y);
61+ } else {
62+ ACLNN_CMD(aclnnClippedSwiglu, x, group_index, dim, alpha, limit, bias, interleaved, y);
63+ }
64+ return y;
65+}
66+ 
67+} // namespace activation
68+} // namespace cann_ops_nn
69+ 
70+PYBIND11_MODULE(TORCH_EXTENSION_NAME, m)
71+{
72+ m.def("clipped_swiglu", &cann_ops_nn::activation::clipped_swiglu, "ClippedSwiglu operator on NPU");
73+}
Aactivation/clipped_swiglu/torch_extension/graph_convert_clipped_swiglu.py+95-0
@@ -0,0 +1,95 @@
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+# GE Converter for Graph Mode
11+ 
12+try:
13+ from typing import Optional
14+ 
15+ import torch
16+ from torchair.ge import attr
17+ from torchair.ge._ge_graph import Tensor, TensorSpec
18+ from torchair._ge_concrete_graph.compat_ir import ge_op, IrDef
19+ from torchair._ge_concrete_graph.fx2ge_converter import (
20+ register_fx_node_ge_converter,
21+ )
22+ 
23+ _TORCHAIR_AVAILABLE = True
24+except ImportError:
25+ _TORCHAIR_AVAILABLE = False
26+ 
27+ 
28+if _TORCHAIR_AVAILABLE:
29+ 
30+ def ClippedSwiglu(
31+ x: Tensor,
32+ *,
33+ group_index: Optional[Tensor],
34+ dim: int = -1,
35+ alpha: float = 1.702,
36+ limit: float = 7.0,
37+ bias: float = 1.0,
38+ interleaved: bool = True,
39+ clamp_mode: int = 0,
40+ ):
41+ inputs = {"x": x}
42+ if group_index is not None:
43+ inputs["group_index"] = group_index
44+ 
45+ return ge_op(
46+ op_type="ClippedSwiglu",
47+ inputs=inputs,
48+ attrs={
49+ "dim": attr.Int(dim),
50+ "alpha": attr.Float(alpha),
51+ "limit": attr.Float(limit),
52+ "bias": attr.Float(bias),
53+ "interleaved": attr.Bool(interleaved),
54+ "clamp_mode": attr.Int(clamp_mode),
55+ },
56+ outputs=["y"],
57+ ir=IrDef("ClippedSwiglu")
58+ .input("x", "DT_BF16, DT_FLOAT16, DT_FLOAT")
59+ .optional_input("group_index", "DT_INT64")
60+ .attr("dim", attr.Int(-1))
61+ .attr("alpha", attr.Float(1.702))
62+ .attr("limit", attr.Float(7.0))
63+ .attr("bias", attr.Float(1.0))
64+ .attr("interleaved", attr.Bool(True))
65+ .attr("clamp_mode", attr.Int(0))
66+ .output("y", "DT_BF16, DT_FLOAT16, DT_FLOAT"),
67+ )
68+ 
69+ @register_fx_node_ge_converter(torch.ops.cann_ops_nn.clipped_swiglu.default)
70+ def convert_clipped_swiglu(
71+ x: Tensor,
72+ *,
73+ group_index: Optional[Tensor] = None,
74+ dim: int = -1,
75+ alpha: float = 1.702,
76+ limit: float = 7.0,
77+ bias: float = 1.0,
78+ interleaved: bool = True,
79+ clamp_mode: int = 0,
80+ meta_outputs: TensorSpec = None,
81+ ):
82+ return ClippedSwiglu(
83+ x,
84+ group_index=group_index,
85+ dim=dim,
86+ alpha=alpha,
87+ limit=limit,
88+ bias=bias,
89+ interleaved=interleaved,
90+ clamp_mode=clamp_mode,
91+ )
92+else:
93+ 
94+ def convert_clipped_swiglu(*args, **kwargs):
95+ raise RuntimeError("ClippedSwiglu graph converter: torchair is not available.")
Mdocs/zh/op_api_list.md+1-0
@@ -97,6 +97,7 @@
97| [aclnnCelu&aclnnInplaceCelu](../../activation/celu_v2/docs/aclnnCelu&aclnnInplaceCelu.md) | aclnnCelu对输入张量self中的每个元素x调用连续可微指数线性单元激活函数CELU，并将得到的结果存入输出张量out中。 | 默认确定性实现 | 默认确定性实现 |97| [aclnnCelu&aclnnInplaceCelu](../../activation/celu_v2/docs/aclnnCelu&aclnnInplaceCelu.md) | aclnnCelu对输入张量self中的每个元素x调用连续可微指数线性单元激活函数CELU，并将得到的结果存入输出张量out中。 | 默认确定性实现 | 默认确定性实现 |
98| [aclnnChamferDistanceBackward](../../loss/chamfer_distance_grad/docs/aclnnChamferDistanceBackward.md) | ChamferDistance（倒角距离）的反向算子，根据正向的输入对输出的贡献及初始梯度求出输入对应的梯度。 | 默认非确定性实现，支持配置开启 | - |98| [aclnnChamferDistanceBackward](../../loss/chamfer_distance_grad/docs/aclnnChamferDistanceBackward.md) | ChamferDistance（倒角距离）的反向算子，根据正向的输入对输出的贡献及初始梯度求出输入对应的梯度。 | 默认非确定性实现，支持配置开启 | - |
99| [aclnnClippedSwiglu](../../activation/clipped_swiglu/docs/aclnnClippedSwiglu.md) | 带截断的Swish门控线性单元激活函数，实现x的SwiGlu计算。 | 默认确定性实现 | 默认确定性实现 |99| [aclnnClippedSwiglu](../../activation/clipped_swiglu/docs/aclnnClippedSwiglu.md) | 带截断的Swish门控线性单元激活函数，实现x的SwiGlu计算。 | 默认确定性实现 | 默认确定性实现 |
100+| [aclnnClippedSwigluV2](../../activation/clipped_swiglu/docs/aclnnClippedSwigluV2.md) | 带截断的Swish门控线性单元激活函数，实现x的SwiGlu计算，新增clamp_mode用于控制clamp和silu计算的顺序。 | 默认确定性实现 | 默认确定性实现 |
100| [aclnnClippedSwigluGrad](../../activation/clipped_swiglu_grad/docs/aclnnClippedSwigluGrad.md) | aclnnClippedSwiglu的反向算子，计算输入x的梯度gradXOut。 | 默认确定性实现 | 默认确定性实现 |101| [aclnnClippedSwigluGrad](../../activation/clipped_swiglu_grad/docs/aclnnClippedSwigluGrad.md) | aclnnClippedSwiglu的反向算子，计算输入x的梯度gradXOut。 | 默认确定性实现 | 默认确定性实现 |
101| [aclnnConvolution](../../conv/convolution_forward/docs/aclnnConvolution.md) | 实现卷积功能，支持1D/2D/3D、转置卷积、空洞卷积、分组卷积。 | 默认确定性实现 | 默认确定性实现 |102| [aclnnConvolution](../../conv/convolution_forward/docs/aclnnConvolution.md) | 实现卷积功能，支持1D/2D/3D、转置卷积、空洞卷积、分组卷积。 | 默认确定性实现 | 默认确定性实现 |
102| [aclnnConvolutionBackward](../../conv/convolution_backward/docs/aclnnConvolutionBackward.md) | 实现卷积的反向传播。 | 默认非确定性实现，支持配置开启 | 默认确定性实现 |103| [aclnnConvolutionBackward](../../conv/convolution_backward/docs/aclnnConvolutionBackward.md) | 实现卷积的反向传播。 | 默认非确定性实现，支持配置开启 | 默认确定性实现 |
Mdocs/zh/op_list.md+1-1
@@ -52,7 +52,7 @@
52 <td><a href="../../activation/clipped_swiglu/README.md">clipped_swiglu</a></td>52 <td><a href="../../activation/clipped_swiglu/README.md">clipped_swiglu</a></td>
53 <td>✓</td>53 <td>✓</td>
54 <td>✓</td>54 <td>✓</td>
55- <td>✗</td>55+ <td>✓</td>
56 <td>✓</td>56 <td>✓</td>
57 <td>AI Core</td>57 <td>AI Core</td>
58 <td>带截断的Swish门控线性单元激活函数，实现x的SwiGlu计算。本算子相较于SwiGlu算子，新增了部分输入参数：groupIndex、alpha、limit、bias、interleaved，用于支持GPT-OSS模型使用的变体SwiGlu以及MoE模型使用的分组场景。</td>58 <td>带截断的Swish门控线性单元激活函数，实现x的SwiGlu计算。本算子相较于SwiGlu算子，新增了部分输入参数：groupIndex、alpha、limit、bias、interleaved，用于支持GPT-OSS模型使用的变体SwiGlu以及MoE模型使用的分组场景。</td>
Mdocs/zh/torch_api_list.md+1-0
@@ -27,6 +27,7 @@
27|[situ_glu](../../activation/situ_glu/docs/torchapi_situ_glu.md)|SiTU门控线性单元激活函数，对输入张量沿指定维度切分为门控与上路径两半，按SiTU公式计算输出。|默认支持确定性计算|默认支持确定性计算|27|[situ_glu](../../activation/situ_glu/docs/torchapi_situ_glu.md)|SiTU门控线性单元激活函数，对输入张量沿指定维度切分为门控与上路径两半，按SiTU公式计算输出。|默认支持确定性计算|默认支持确定性计算|
28|[situ_glu_grad](../../activation/situ_glu_grad/docs/torchapi_situ_glu_grad.md)|SiTU门控线性单元激活函数的反向梯度计算，根据上游梯度和前向输入计算输入梯度。|默认支持确定性计算|默认支持确定性计算|28|[situ_glu_grad](../../activation/situ_glu_grad/docs/torchapi_situ_glu_grad.md)|SiTU门控线性单元激活函数的反向梯度计算，根据上游梯度和前向输入计算输入梯度。|默认支持确定性计算|默认支持确定性计算|
29|[swiglu_group](../../activation/swiglu_group/docs/torchapi_swiglu_group.md)|SwiGLU分组激活算子，对输入张量按最后一维拆分为两部分，分别进行clamp和sigmoid操作后相乘，支持可选的权重和分组索引。|-|默认支持确定性计算|29|[swiglu_group](../../activation/swiglu_group/docs/torchapi_swiglu_group.md)|SwiGLU分组激活算子，对输入张量按最后一维拆分为两部分，分别进行clamp和sigmoid操作后相乘，支持可选的权重和分组索引。|-|默认支持确定性计算|
30+|[clipped_swiglu](../../activation/clipped_swiglu/docs/torchapi_clipped_swiglu.md)|带截断的Swish门控线性单元激活函数，实现x的SwiGlu计算，输出SwiGlu的结果。|-|默认支持确定性计算|
30|[clipped_swiglu_grad](../../activation/clipped_swiglu_grad/docs/torchapi_clipped_swiglu_grad.md)|ClippedSwiglu的反向梯度算子，根据上游梯度grad_y和前向输入x重算clamp mask与sigmoid，输出grad_x。|-|默认支持确定性计算|31|[clipped_swiglu_grad](../../activation/clipped_swiglu_grad/docs/torchapi_clipped_swiglu_grad.md)|ClippedSwiglu的反向梯度算子，根据上游梯度grad_y和前向输入x重算clamp mask与sigmoid，输出grad_x。|-|默认支持确定性计算|
31|[swiglu_group_backward](../../quant/swiglu_group_grad/docs/torchapi_swiglu_group_backward.md)|ClampedSwiglu激活函数的反向梯度算子，从上游梯度grad_output和前向输入x重算clamp mask与sigmoid，输出grad_x与可选grad_weight。|-|-|32|[swiglu_group_backward](../../quant/swiglu_group_grad/docs/torchapi_swiglu_group_backward.md)|ClampedSwiglu激活函数的反向梯度算子，从上游梯度grad_output和前向输入x重算clamp mask与sigmoid，输出grad_x与可选grad_weight。|-|-|
32|[swiglu_group_quant](../../activation/swiglu_group_quant/docs/torchapi_swiglu_group_quant.md)|融合SwiGLU分组激活与量化的算子，在SwiGLU计算基础上支持FP8/MXFP4等多种量化模式输出。|-|默认支持确定性计算|33|[swiglu_group_quant](../../activation/swiglu_group_quant/docs/torchapi_swiglu_group_quant.md)|融合SwiGLU分组激活与量化的算子，在SwiGLU计算基础上支持FP8/MXFP4等多种量化模式输出。|-|默认支持确定性计算|