Star901
1819
代码
代码
Issues224
Pull Requests313
Actions
项目讨论
Wiki
分析
Star901
1819
已合并
feat: clipped_swiglu 算子适配 Ascend 950 平台 #5483
chenfeng创建于 5月30日
feat: clipped_swiglu 算子适配 Ascend 950 平台 #5483
已合并
chenfeng创建于 5月30日
从chenfeng61/ops-nnascend950_clippen_swiglu合入到cann/ops-nnmaster
共 14 个文件变更+1683-41
Mactivation/clipped_swiglu/README.md+5-4
@@ -4,7 +4,7 @@
4 4 
5|产品 | 是否支持 |5|产品 | 是否支持 |
6|:-------------------------|:----------:|6|:-------------------------|:----------:|
7-| <term>Ascend 950PR/Ascend 950DT</term> | × |7+| <term>Ascend 950PR/Ascend 950DT</term> | √ |
8| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |8| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |
9| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |9| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |
10| <term>Atlas 200I/500 A2 推理产品</term> | × |10| <term>Atlas 200I/500 A2 推理产品</term> | × |
@@ -105,7 +105,7 @@
105 <tr>105 <tr>
106 <td>group_index</td>106 <td>group_index</td>
107 <td>可选输入</td>107 <td>可选输入</td>
108- <td>公式中的输入group_index。维度必须是1维。</td>108+ <td>公式中的输入group_index。维度必须是1维，且元素个数必须小于等于1024。</td>
109 <td>INT64</td>109 <td>INT64</td>
110 <td>-</td>110 <td>-</td>
111 </tr>111 </tr>
@@ -126,7 +126,7 @@
126 <tr>126 <tr>
127 <td>limit</td>127 <td>limit</td>
128 <td>可选属性</td>128 <td>可选属性</td>
129- <td>公式中的输入limit，表示变体SwiGlu使用的门限值。默认为7.0。</td>129+ <td>公式中的输入limit，表示变体SwiGlu使用的门限值。必须大于0。默认为7.0。</td>
130 <td>FLOAT</td>130 <td>FLOAT</td>
131 <td>-</td>131 <td>-</td>
132 </tr>132 </tr>
@@ -155,7 +155,8 @@
155 155 
156## 约束说明156## 约束说明
157 157 
158-无158+- 可选输入 group_index 为1维，且元素个数必须小于等于1024。
159+- 可选属性 limit 必须大于0。
159 160 
160## 调用说明161## 调用说明
161 162 
Mactivation/clipped_swiglu/docs/aclnnClippedSwiglu.md+3-3
@@ -6,7 +6,7 @@
6 6 
7|产品 | 是否支持 |7|产品 | 是否支持 |
8|:-------------------------|:----------:|8|:-------------------------|:----------:|
C
Cchenjiao6月1日
已过期

需同步修改op_api_list.md中确定性

likedislike
9-| <term>Ascend 950PR/Ascend 950DT</term> | × |9+| <term>Ascend 950PR/Ascend 950DT</term> | √ |
10| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |10| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ |
11| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |11| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ |
12| <term>Atlas 200I/500 A2 推理产品</term> | × |12| <term>Atlas 200I/500 A2 推理产品</term> | × |
@@ -144,7 +144,7 @@ aclnnStatus aclnnClippedSwiglu(
144 <td>groupIndexOptional（aclTensor*）</td>144 <td>groupIndexOptional（aclTensor*）</td>
145 <td>输入</td>145 <td>输入</td>
146 <td>公式中的输入group_index，表示分组的情况。</td>146 <td>公式中的输入group_index，表示分组的情况。</td>
147- <td>支持空指针。不为空指针时，维度要求为1维，且元素需大于等于0。第i个元素代表第i组需要处理x的batch数量。</td>147+ <td>支持空指针。不为空指针时，维度要求为1维，长度不超过1024，且元素需大于等于0。第i个元素代表第i组需要处理x的batch数量。</td>
148 <td>INT64</td>148 <td>INT64</td>
149 <td>ND</td>149 <td>ND</td>
150 <td>1</td>150 <td>1</td>
@@ -174,7 +174,7 @@ aclnnStatus aclnnClippedSwiglu(
174 <td>limit（double）</td>174 <td>limit（double）</td>
175 <td>输入</td>175 <td>输入</td>
176 <td>公式中的输入limit，表示变体SwiGlu使用的门限值。</td>176 <td>公式中的输入limit，表示变体SwiGlu使用的门限值。</td>
177- <td>建议为7.0。</td>177+ <td>必须大于0，建议为7.0。</td>
178 <td>-</td>178 <td>-</td>
179 <td>-</td>179 <td>-</td>
180 <td>-</td>180 <td>-</td>
Mactivation/clipped_swiglu/op_host/clipped_swiglu_def.cpp+39-9
@@ -17,32 +17,62 @@
17namespace ops {17namespace ops {
18constexpr float DEFAULT_ALPHA = 1.702;18constexpr float DEFAULT_ALPHA = 1.702;
19constexpr float DEFAULT_LIMIT = 7.0;19constexpr float DEFAULT_LIMIT = 7.0;
20+ 
21+static const std::vector<ge::DataType> xDtype = {ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16};
22+static const std::vector<ge::DataType> groupIndexDtype = {ge::DT_INT64, ge::DT_INT64, ge::DT_INT64};
23+static const std::vector<ge::Format> xFormat = {ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND};
24+ 
20class ClippedSwiglu : public OpDef {25class ClippedSwiglu : public OpDef {
21public:26public:
22 explicit ClippedSwiglu(const char* name) : OpDef(name)27 explicit ClippedSwiglu(const char* name) : OpDef(name)
23 {28 {
24 this->Input("x")29 this->Input("x")
25 .ParamType(REQUIRED)30 .ParamType(REQUIRED)
26- .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})31+ .DataType(xDtype)
27- .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})32+ .Format(xFormat)
28- .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});33+ .UnknownShapeFormat(xFormat);
29 this->Input("group_index")34 this->Input("group_index")
30 .ParamType(OPTIONAL)35 .ParamType(OPTIONAL)
31- .DataType({ge::DT_INT64, ge::DT_INT64, ge::DT_INT64})36+ .DataType(groupIndexDtype)
32- .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})37+ .Format(xFormat)
33- .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});38+ .UnknownShapeFormat(xFormat);
34 this->Output("y")39 this->Output("y")
35 .ParamType(REQUIRED)40 .ParamType(REQUIRED)
36- .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})41+ .DataType(xDtype)
37- .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})42+ .Format(xFormat)
38- .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});43+ .UnknownShapeFormat(xFormat);
39 this->Attr("dim").AttrType(OPTIONAL).Int(-1);44 this->Attr("dim").AttrType(OPTIONAL).Int(-1);
40 this->Attr("alpha").AttrType(OPTIONAL).Float(DEFAULT_ALPHA);45 this->Attr("alpha").AttrType(OPTIONAL).Float(DEFAULT_ALPHA);
41 this->Attr("limit").AttrType(OPTIONAL).Float(DEFAULT_LIMIT);46 this->Attr("limit").AttrType(OPTIONAL).Float(DEFAULT_LIMIT);
42 this->Attr("bias").AttrType(OPTIONAL).Float(1.0);47 this->Attr("bias").AttrType(OPTIONAL).Float(1.0);
43 this->Attr("interleaved").AttrType(OPTIONAL).Bool(true);48 this->Attr("interleaved").AttrType(OPTIONAL).Bool(true);
49+ 
44 this->AICore().AddConfig("ascend910b");50 this->AICore().AddConfig("ascend910b");
45 this->AICore().AddConfig("ascend910_93");51 this->AICore().AddConfig("ascend910_93");
52+ 
53+ OpAICoreConfig regbaseConfig;
54+ regbaseConfig.Input("x")
55+ .ParamType(REQUIRED)
56+ .DataType(xDtype)
57+ .Format(xFormat)
58+ .UnknownShapeFormat(xFormat)
59+ .AutoContiguous();
60+ regbaseConfig.Input("group_index")
61+ .ParamType(OPTIONAL)
62+ .DataType(groupIndexDtype)
63+ .Format(xFormat)
64+ .UnknownShapeFormat(xFormat)
65+ .AutoContiguous();
66+ regbaseConfig.Output("y")
67+ .ParamType(REQUIRED)
68+ .DataType(xDtype)
69+ .Format(xFormat)
70+ .UnknownShapeFormat(xFormat);
71+ regbaseConfig.DynamicCompileStaticFlag(true)
72+ .DynamicRankSupportFlag(true)
73+ .DynamicShapeSupportFlag(true)
74+ .ExtendCfgInfo("opFile.value", "clipped_swiglu_apt");
75+ this->AICore().AddConfig("ascend950", regbaseConfig);
46 }76 }
47};77};
48OP_ADD(ClippedSwiglu);78OP_ADD(ClippedSwiglu);
Mactivation/clipped_swiglu/op_host/clipped_swiglu_infershape.cpp+27-24
@@ -16,7 +16,7 @@
16#include "log/log.h"16#include "log/log.h"
17#include "register/op_impl_registry.h"17#include "register/op_impl_registry.h"
18#include "error_util.h"18#include "error_util.h"
19- 19+#include "util/shape_util.h"
20using namespace ge;20using namespace ge;
21 21 
22namespace {22namespace {
@@ -29,37 +29,40 @@ const size_t SPLIT_NUM = 2;
29namespace ops {29namespace ops {
30static ge::graphStatus InferShapeForClippedSwiglu(gert::InferShapeContext* context) {30static ge::graphStatus InferShapeForClippedSwiglu(gert::InferShapeContext* context) {
31 OP_LOGD(context->GetNodeName(), "Begin to do InferShapeForClippedSwiglu");31 OP_LOGD(context->GetNodeName(), "Begin to do InferShapeForClippedSwiglu");
32- auto x_shape = context->GetInputShape(GLU_IN_X);32+ auto xShape = context->GetInputShape(GLU_IN_X);
33- OPS_CHECK_NULL_WITH_CONTEXT(context, x_shape);33+ OPS_CHECK_NULL_WITH_CONTEXT(context, xShape);
34- auto y_shape = context->GetOutputShape(GLU_OUT_Y);34+ auto yShape = context->GetOutputShape(GLU_OUT_Y);
35- OPS_CHECK_NULL_WITH_CONTEXT(context, y_shape);35+ OPS_CHECK_NULL_WITH_CONTEXT(context, yShape);
36 auto attrs = context->GetAttrs();36 auto attrs = context->GetAttrs();
37 OPS_CHECK_NULL_WITH_CONTEXT(context, attrs);37 OPS_CHECK_NULL_WITH_CONTEXT(context, attrs);
38 38 
39- auto split_dim_ptr = attrs->GetAttrPointer<int64_t>(GLU_ATTR_DIM);39+ auto splitDimPtr = attrs->GetAttrPointer<int64_t>(GLU_ATTR_DIM);
40- OPS_CHECK_NULL_WITH_CONTEXT(context, split_dim_ptr);40+ OPS_CHECK_NULL_WITH_CONTEXT(context, splitDimPtr);
41- 41+ if (Ops::Base::IsUnknownRank(*xShape)) {
42- auto split_dim = *split_dim_ptr;42+ Ops::Base::SetUnknownRank(*yShape);
43- if (split_dim < 0) {
44- split_dim += x_shape->GetDimNum();
45- }
46- if (split_dim < 0 || split_dim >= static_cast<int64_t>(x_shape->GetDimNum())) {
47- OP_LOGE("ClippedSwiglu", "The value of attr [dim] must be in the range [-%zu, %zu], but got [%ld].",
48- x_shape->GetDimNum(), x_shape->GetDimNum() - 1, split_dim);
49- return GRAPH_FAILED;
50- }
51- OP_LOGD(context->GetNodeName(), "Begin to generate y_shape");
52- *y_shape = *x_shape;
53- // dynamic shape
54- if (x_shape->GetDim(split_dim) == -1) {
55 return ge::GRAPH_SUCCESS;43 return ge::GRAPH_SUCCESS;
56 }44 }
57- if (x_shape->GetDim(split_dim) < 0 || x_shape->GetDim(split_dim) % SPLIT_NUM != 0) {45+ auto splitDim = *splitDimPtr;
58- OP_LOGE("ClippedSwiglu", "The shape [%s] is not divisible by 2.", Ops::Base::ToString(*x_shape).c_str());46+ if (splitDim < 0) {
47+ splitDim += xShape->GetDimNum();
48+ }
49+ if (splitDim < 0 || splitDim >= static_cast<int64_t>(xShape->GetDimNum())) {
50+ OP_LOGE("ClippedSwiglu", "The value of attr [dim] must be in the range [-%zu, %zu], but got [%ld].",
51+ xShape->GetDimNum(), xShape->GetDimNum() - 1, splitDim);
52+ return GRAPH_FAILED;
53+ }
54+ OP_LOGD(context->GetNodeName(), "Begin to generate yShape");
55+ *yShape = *xShape;
56+ // dynamic shape
57+ if (xShape->GetDim(splitDim) == -1) {
58+ return ge::GRAPH_SUCCESS;
59+ }
60+ if (xShape->GetDim(splitDim) < 0 || xShape->GetDim(splitDim) % SPLIT_NUM != 0) {
61+ OP_LOGE("ClippedSwiglu", "The shape [%s] is not divisible by 2.", Ops::Base::ToString(*xShape).c_str());
59 return ge::GRAPH_FAILED;62 return ge::GRAPH_FAILED;
60 }63 }
61 64 
62- y_shape->SetDim(split_dim, x_shape->GetDim(split_dim) / SPLIT_NUM);65+ yShape->SetDim(splitDim, xShape->GetDim(splitDim) / SPLIT_NUM);
63 OP_LOGD(context->GetNodeName(), "End to do InferShapeForClippedSwiglu");66 OP_LOGD(context->GetNodeName(), "End to do InferShapeForClippedSwiglu");
64 return ge::GRAPH_SUCCESS;67 return ge::GRAPH_SUCCESS;
65}68}
Mactivation/clipped_swiglu/op_host/clipped_swiglu_tiling.cpp+4-0
@@ -298,6 +298,10 @@ REGISTER_TILING_TEMPLATE("ClippedSwiglu", ClippedSwigluTiling, 20000);
298 298 
299ge::graphStatus TilingForClippedSwiglu(gert::TilingContext* context)299ge::graphStatus TilingForClippedSwiglu(gert::TilingContext* context)
300{300{
301+ bool regBase = Ops::NN::OpTiling::IsRegbaseSocVersion(context);
302+ if (regBase) {
303+ return Tiling4ClippedSwigluArch35(context);
304+ }
301 return TilingRegistry::GetInstance().DoTilingImpl(context);305 return TilingRegistry::GetInstance().DoTilingImpl(context);
302}306}
303 307 
Mactivation/clipped_swiglu/op_host/clipped_swiglu_tiling.h+3-0
@@ -25,6 +25,7 @@
25#include "../op_graph/clipped_swiglu_proto.h"25#include "../op_graph/clipped_swiglu_proto.h"
26#include "op_host/tiling_base.h"26#include "op_host/tiling_base.h"
27#include "op_host/tiling_templates_registry.h"27#include "op_host/tiling_templates_registry.h"
28+#include "op_host/tiling_util.h"
28 29 
29namespace optiling {30namespace optiling {
30using Ops::NN::Optiling::TilingBaseClass;31using Ops::NN::Optiling::TilingBaseClass;
@@ -95,5 +96,7 @@ private:
95 int64_t groupNum_ = 0;96 int64_t groupNum_ = 0;
96};97};
97 98 
99+ge::graphStatus Tiling4ClippedSwigluArch35(gert::TilingContext* context);
100+ 
98} // namespace optiling101} // namespace optiling
99#endif // AIR_CXX_RUNTIME_V2_OP_IMPL_CLIPPED_SWIGLU_H_102#endif // AIR_CXX_RUNTIME_V2_OP_IMPL_CLIPPED_SWIGLU_H_
Aactivation/clipped_swiglu/op_host/clipped_swiglu_tiling_arch35.cpp+447-0
@@ -0,0 +1,447 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file clipped_swiglu_tiling_arch35.cpp
13+ * \brief Tiling implementation for ClippedSwiglu Arch35 (Ascend 950)
14+ */
15+ 
16+#include <iostream>
17+#include <cstring>
18+#include "register/tilingdata_base.h"
19+#include "register/op_def_registry.h"
20+#include "tiling/tiling_api.h"
21+#include "clipped_swiglu_tiling.h"
22+#include "op_common/op_host/util/platform_util.h"
23+#include "../op_kernel/arch35/clipped_swiglu_tiling_data.h"
24+#include "../op_kernel/arch35/clipped_swiglu_tiling_key.h"
25+ 
26+using namespace ge;
27+using namespace ClippedSwigluOp;
28+ 
29+namespace optiling {
30+ 
31+constexpr int64_t X_INDEX = 0;
32+constexpr int64_t GROUP_INDEX_INDEX = 1;
33+constexpr int64_t Y_INDEX = 0;
34+constexpr int64_t DIM_INDEX = 0;
35+constexpr int64_t ALPHA_INDEX = 1;
36+constexpr int64_t LIMIT_INDEX = 2;
37+constexpr int64_t BIAS_INDEX = 3;
38+constexpr int64_t INTERLEAVED_INDEX = 4;
39+ 
40+constexpr int64_t CONST_2 = 2;
41+constexpr int64_t CONST_4 = 4;
42+constexpr int64_t CONST_7 = 7; // inputUb = 2x*db outputUb = x*db vectorUb = x
43+constexpr int64_t CONST_8 = 8; // int64 size is 8
44+constexpr int64_t DB_BUFFER = 2;
45+ 
46+constexpr float CLAMP_LIMIT_DEFAULT = 7.0;
47+constexpr float GLU_ALPHA_DEFAULT = 1.702;
48+constexpr float GLU_BIAS_DEFAULT = 1.0;
49+ 
50+static const std::set<ge::DataType> SUPPORT_DTYPE = {ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16};
51+ 
52+class ClippedSwigluArch35Tiling {
53+public:
54+ explicit ClippedSwigluArch35Tiling(gert::TilingContext* context) : context_(context) {}
55+ 
56+ ge::graphStatus Init();
57+ ge::graphStatus DoTiling();
58+ 
59+private:
60+ ge::graphStatus GetPlatformInfo();
61+ ge::graphStatus CheckAndGetXAndAttrs();
62+ ge::graphStatus CheckInputX(const gert::Shape& inputShapeX, int64_t xSize);
63+ ge::graphStatus CheckAfterDim(const gert::Shape &xShape, const gert::Shape &inputShapeY, std::string shapeMsg);
64+ ge::graphStatus CheckAndGetGroupIndex();
65+ ge::graphStatus CheckY();
66+ ge::graphStatus CountUbFactor();
67+ void ComputeCoreSplit();
68+ void SetTilingKey();
69+ void FillTilingData();
70+ void PrintTilingInfo();
71+ 
72+private:
73+ gert::TilingContext* context_;
74+ ClippedSwigluArch35TilingData* tilingData_ = nullptr;
75+ 
76+ uint64_t coreNumAll_ = 0;
77+ uint64_t ubSize_ = 0;
78+ int64_t blockSize_ = 0;
79+ int64_t xDims_ = 0;
80+ int64_t cutDim_ = 0;
81+ int64_t dimBatchSize_ = 1;
82+ int64_t dim2H_ = 1;
83+ int64_t dimH_ = 1;
84+ int64_t xCutDimNum_ = 0;
85+ ge::DataType xDtype_ = ge::DT_FLOAT;
86+ int64_t dtypeSize_ = CONST_2;
87+ int64_t isGroup_ = 0;
88+ int64_t isInterleaved_ = 1;
89+ float gluLimit_ = 0.0;
90+ float gluAlpha_ = 0.0;
91+ float gluBias_ = 0.0;
92+ int64_t hUbFactor_ = 1;
93+ int64_t bUbFactor_ = 1;
94+ int64_t groupNum_ = 0;
95+ int64_t realCoreNum_ = 0;
96+ uint64_t tilingKey_ = 0;
97+ int64_t workspaceSize_ = 0;
98+};
99+ 
100+ge::graphStatus ClippedSwigluArch35Tiling::Init()
101+{
102+ tilingData_ = context_->GetTilingData<ClippedSwigluArch35TilingData>();
103+ OP_CHECK_IF(tilingData_ == nullptr, OP_LOGE(context_, "get tilingdata ptr failed"), return ge::GRAPH_FAILED);
104+ OP_CHECK_IF(
105+ (memset_s(tilingData_, sizeof(ClippedSwigluArch35TilingData), 0, sizeof(ClippedSwigluArch35TilingData)) != EOK),
106+ OP_LOGE(context_, "memset tilingdata failed"), return ge::GRAPH_FAILED);
107+ if (GetPlatformInfo() != ge::GRAPH_SUCCESS) {
108+ OP_LOGE(context_, "GetPlatformInfo failed.");
109+ return ge::GRAPH_FAILED;
110+ }
111+ return ge::GRAPH_SUCCESS;
112+}
113+ 
114+ge::graphStatus ClippedSwigluArch35Tiling::GetPlatformInfo()
115+{
116+ auto platformInfo = context_->GetPlatformInfo();
117+ OP_CHECK_NULL_WITH_CONTEXT(context_, platformInfo);
118+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo);
119+ coreNumAll_ = ascendcPlatform.GetCoreNumAiv();
120+ OP_CHECK_IF((coreNumAll_ <= 0), OP_LOGE(context_, "core num must > 0"), return ge::GRAPH_FAILED);
121+ uint64_t ubSize;
122+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
123+ ubSize_ = ubSize;
124+ OP_CHECK_IF((ubSize_ == 0), OP_LOGE(context_, "ubSize must > 0"), return ge::GRAPH_FAILED);
125+ blockSize_ = Ops::Base::GetUbBlockSize(context_);
126+ OP_CHECK_IF((blockSize_ <= 0), OP_LOGE(context_, "block size is invalid."), return ge::GRAPH_FAILED);
127+ workspaceSize_ = ascendcPlatform.GetLibApiWorkSpaceSize();
128+ return ge::GRAPH_SUCCESS;
129+}
130+ 
131+ge::graphStatus ClippedSwigluArch35Tiling::DoTiling()
132+{
133+ OP_CHECK_IF(
134+ CheckAndGetXAndAttrs() != ge::GRAPH_SUCCESS, OP_LOGE(context_, "check x and attrs failed."),
135+ return ge::GRAPH_FAILED);
136+ OP_CHECK_IF(
137+ CheckAndGetGroupIndex() != ge::GRAPH_SUCCESS, OP_LOGE(context_, "check group_index failed."),
138+ return ge::GRAPH_FAILED);
139+ OP_CHECK_IF(CheckY() != ge::GRAPH_SUCCESS, OP_LOGE(context_, "check y failed."), return ge::GRAPH_FAILED);
140+ OP_CHECK_IF(
141+ CountUbFactor() != ge::GRAPH_SUCCESS, OP_LOGE(context_, "CountUbFactor failed."), return ge::GRAPH_FAILED);
142+ ComputeCoreSplit();
143+ SetTilingKey();
144+ FillTilingData();
145+ PrintTilingInfo();
146+ 
147+ context_->SetTilingKey(tilingKey_);
148+ context_->SetBlockDim(realCoreNum_);
149+ size_t* workspaces = context_->GetWorkspaceSizes(1);
150+ OP_CHECK_NULL_WITH_CONTEXT(context_, workspaces);
151+ workspaces[0] = workspaceSize_;
152+ return ge::GRAPH_SUCCESS;
153+}
154+ 
155+ge::graphStatus ClippedSwigluArch35Tiling::CheckInputX(const gert::Shape& inputShapeX, int64_t xSize)
156+{
157+ std::string reasonMsg = "in [" + std::to_string(0) + ", " + std::to_string(xDims_ - 1) + "]";
158+ OP_CHECK_IF(
159+ (cutDim_ > (xDims_ - 1) || cutDim_ < 0),
160+ OP_LOGE_FOR_INVALID_VALUE(context_->GetNodeName(), "dim", std::to_string(cutDim_), reasonMsg),
161+ return ge::GRAPH_FAILED);
162+ OP_CHECK_IF(
163+ xSize <= 0,
164+ OP_LOGE_FOR_INVALID_SHAPESIZES_WITH_REASON(
165+ context_->GetNodeName(), "x", std::to_string(xSize), "x shape size must > 0"),
166+ return ge::GRAPH_FAILED);
167+ xCutDimNum_ = inputShapeX.GetDim(cutDim_);
168+ dimBatchSize_ = 1;
169+ dim2H_ = 1;
170+ if (xDims_ == 1) {
171+ dimBatchSize_ = 1;
172+ dim2H_ = inputShapeX.GetDim(0);
173+ } else {
174+ for (int64_t i = 0; i < cutDim_; i++) {
175+ dimBatchSize_ *= inputShapeX.GetDim(i);
176+ }
177+ for (int64_t j = cutDim_; j < xDims_; j++) {
178+ dim2H_ *= inputShapeX.GetDim(j);
179+ }
180+ }
181+ dimH_ = dim2H_ / CONST_2;
182+ std::string reason = "xShape[ " + std::to_string(cutDim_) + "] must be divisible by 2";
183+ OP_CHECK_IF(
184+ xCutDimNum_ % 2 != 0,
185+ OP_LOGE_FOR_INVALID_SHAPE_WITH_REASON(
186+ context_->GetNodeName(), "x", Ops::Base::ToString(inputShapeX).c_str(), reason.c_str()),
187+ return ge::GRAPH_FAILED);
188+ auto descX = context_->GetInputDesc(X_INDEX);
189+ OP_CHECK_NULL_WITH_CONTEXT(context_, descX);
190+ xDtype_ = descX->GetDataType();
191+ OP_CHECK_IF(
192+ (SUPPORT_DTYPE.find(xDtype_) == SUPPORT_DTYPE.end()),
193+ OP_LOGE_FOR_INVALID_DTYPE(
194+ context_->GetNodeName(), "x", ge::TypeUtils::DataTypeToSerialString(xDtype_).c_str(),
195+ "float16, bfloat16, float32"),
196+ return ge::GRAPH_FAILED);
197+ if (xDtype_ == ge::DT_FLOAT) {
198+ dtypeSize_ = CONST_4;
199+ }
200+ return ge::GRAPH_SUCCESS;
201+}
202+ 
203+ge::graphStatus ClippedSwigluArch35Tiling::CheckAndGetXAndAttrs()
204+{
205+ auto* attrs = context_->GetAttrs();
206+ OP_CHECK_NULL_WITH_CONTEXT(context_, attrs);
207+ auto* attrDim = attrs->GetAttrPointer<int64_t>(DIM_INDEX);
208+ cutDim_ = attrDim == nullptr ? -1 : *attrDim;
209+ auto* attrAlpha = attrs->GetAttrPointer<float>(ALPHA_INDEX);
210+ gluAlpha_ = attrAlpha == nullptr ? GLU_ALPHA_DEFAULT : *attrAlpha;
211+ auto* attrLimit = attrs->GetAttrPointer<float>(LIMIT_INDEX);
212+ gluLimit_ = attrLimit == nullptr ? CLAMP_LIMIT_DEFAULT : *attrLimit;
213+ 
214+ OP_CHECK_IF(
215+ gluLimit_ <= 0.0f,
216+ OP_LOGE_WITH_INVALID_ATTR(context_->GetNodeName(), "limit", std::to_string(gluLimit_), "> 0"),
217+ return ge::GRAPH_FAILED);
218+ auto* attrBias = attrs->GetAttrPointer<float>(BIAS_INDEX);
219+ gluBias_ = attrBias == nullptr ? GLU_BIAS_DEFAULT : *attrBias;
220+ auto* attrInterleaved = attrs->GetAttrPointer<bool>(INTERLEAVED_INDEX);
221+ bool interleaved = attrInterleaved == nullptr ? true : *attrInterleaved;
222+ isInterleaved_ = interleaved ? 1 : 0;
223+ 
224+ auto shapeX = context_->GetInputShape(X_INDEX);
225+ OP_CHECK_NULL_WITH_CONTEXT(context_, shapeX);
226+ const gert::Shape& inputShapeX = shapeX->GetStorageShape();
227+ xDims_ = inputShapeX.GetDimNum();
228+ OP_CHECK_IF(
229+ xDims_ <= 0, OP_LOGE_FOR_INVALID_SHAPEDIM(context_->GetNodeName(), "x", std::to_string(xDims_), "> 0"),
230+ return ge::GRAPH_FAILED);
231+ if (cutDim_ < 0) {
232+ cutDim_ = cutDim_ + xDims_;
233+ }
234+ int64_t xSize = inputShapeX.GetShapeSize();
235+ OP_CHECK_IF(
236+ CheckInputX(inputShapeX, xSize) != ge::GRAPH_SUCCESS, OP_LOGE(context_->GetNodeName(), "check input x failed"),
237+ return ge::GRAPH_FAILED);
238+ return ge::GRAPH_SUCCESS;
239+}
240+ 
241+ge::graphStatus ClippedSwigluArch35Tiling::CheckAndGetGroupIndex()
242+{
243+ auto shapeGroupIndex = context_->GetOptionalInputShape(GROUP_INDEX_INDEX);
244+ if (shapeGroupIndex == nullptr) {
245+ isGroup_ = 0;
246+ } else {
247+ isGroup_ = 1;
248+ const gert::Shape& inputShapeGroupIndex = shapeGroupIndex->GetStorageShape();
249+ int64_t groupIndexDim = inputShapeGroupIndex.GetDimNum();
250+ auto descGroupIndex = context_->GetOptionalInputDesc(GROUP_INDEX_INDEX);
251+ OP_CHECK_NULL_WITH_CONTEXT(context_, descGroupIndex);
252+ auto groupIndexDtype = descGroupIndex->GetDataType();
253+ OP_CHECK_IF(
254+ groupIndexDim != 1,
255+ OP_LOGE_FOR_INVALID_SHAPEDIM(context_->GetNodeName(), "group_index", std::to_string(groupIndexDim), "1D"),
256+ return ge::GRAPH_FAILED);
257+ 
258+ OP_CHECK_IF(
259+ groupIndexDtype != ge::DT_INT64,
260+ OP_LOGE_FOR_INVALID_DTYPE(
261+ context_->GetNodeName(), "group_index", ge::TypeUtils::DataTypeToSerialString(groupIndexDtype).c_str(),
262+ "int64"),
263+ return ge::GRAPH_FAILED);
264+ groupNum_ = inputShapeGroupIndex.GetDim(0);
265+ }
266+ return ge::GRAPH_SUCCESS;
267+}
268+ 
269+ge::graphStatus ClippedSwigluArch35Tiling::CheckAfterDim(const gert::Shape &xShape, const gert::Shape &inputShapeY, std::string shapeMsg)
270+{
271+ if (cutDim_ < xDims_ - 1) {
272+ for (int64_t i = cutDim_ + 1; i < xDims_; i++) {
273+ int64_t xShapeValue = xShape.GetDim(i);
274+ int64_t yShapeValue = inputShapeY.GetDim(i);
275+ std::string shapeMsgValue =
276+ "xShape[" + std::to_string(i) + "] should be equal yShape[" + std::to_string(i) + "].";
277+ OP_CHECK_IF(
278+ xShapeValue != yShapeValue,
279+ OP_LOGE_FOR_INVALID_SHAPES_WITH_REASON(context_->GetNodeName(), "x and y", shapeMsg, shapeMsgValue),
280+ return ge::GRAPH_FAILED);
281+ }
282+ }
283+ return ge::GRAPH_SUCCESS;
284+}
285+ 
286+ge::graphStatus ClippedSwigluArch35Tiling::CheckY()
287+{
288+ auto shapeY = context_->GetOutputShape(Y_INDEX);
289+ OP_CHECK_NULL_WITH_CONTEXT(context_, shapeY);
290+ const gert::Shape& inputShapeY = shapeY->GetStorageShape();
291+ int64_t yDims = inputShapeY.GetDimNum();
292+ auto descY = context_->GetInputDesc(Y_INDEX);
293+ OP_CHECK_NULL_WITH_CONTEXT(context_, descY);
294+ auto yDtype = descY->GetDataType();
295+ auto xShape = context_->GetInputShape(0)->GetStorageShape();
296+ OP_CHECK_IF(
297+ yDims != xDims_,
298+ OP_LOGE_FOR_INVALID_SHAPEDIM(context_->GetNodeName(), "y", std::to_string(yDims), std::to_string(xDims_)),
299+ return ge::GRAPH_FAILED);
300+ 
301+ std::string reasonMsg = "x shape is " + Ops::Base::ToString(xShape) + "xShape[" + std::to_string(cutDim_) +
302+ "] / 2 must be equal yShape[" + std::to_string(cutDim_) + "]";
303+ OP_CHECK_IF(
304+ inputShapeY.GetDim(cutDim_) != (xCutDimNum_ / CONST_2),
305+ OP_LOGE_FOR_INVALID_SHAPE_WITH_REASON(
306+ context_->GetNodeName(), "y", Ops::Base::ToString(inputShapeY).c_str(), reasonMsg.c_str()),
307+ return ge::GRAPH_FAILED);
308+ 
309+ std::string dtypeMsg = "x dtype is " + ge::TypeUtils::DataTypeToSerialString(xDtype_) + " , y dtype is " +
310+ ge::TypeUtils::DataTypeToSerialString(yDtype);
311+ OP_CHECK_IF(
312+ yDtype != xDtype_,
313+ OP_LOGE_FOR_INVALID_DTYPES_WITH_REASON(
314+ context_->GetNodeName(), "x and y", dtypeMsg.c_str(), "y dtype must be same as x dtype"),
315+ return ge::GRAPH_FAILED);
316+ std::string shapeMsg =
317+ "x shape is " + Ops::Base::ToString(xShape) + " y shape is " + Ops::Base::ToString(inputShapeY);
318+ for (int64_t i = 0; i < cutDim_; i++) {
319+ int64_t xShapeValue = xShape.GetDim(i);
320+ int64_t yShapeValue = inputShapeY.GetDim(i);
321+ std::string shapeMsgValue =
322+ "xShape[" + std::to_string(i) + "] should be equal yShape[" + std::to_string(i) + "].";
323+ OP_CHECK_IF(
324+ xShapeValue != yShapeValue,
325+ OP_LOGE_FOR_INVALID_SHAPES_WITH_REASON(context_->GetNodeName(), "x and y", shapeMsg, shapeMsgValue),
326+ return ge::GRAPH_FAILED);
327+ }
328+ OP_CHECK_IF(CheckAfterDim(xShape, inputShapeY, shapeMsg) != ge::GRAPH_SUCCESS, OP_LOGE(context_, "CheckAfterDim failed."), return ge::GRAPH_FAILED);
329+ return ge::GRAPH_SUCCESS;
330+}
331+ 
332+ge::graphStatus ClippedSwigluArch35Tiling::CountUbFactor()
333+{
334+ hUbFactor_ = 1;
335+ int64_t groupIndexBuf = 0;
336+ if (isGroup_ != 0) {
337+ groupIndexBuf = blockSize_;
338+ }
339+ int64_t oneBlockNum = blockSize_ / dtypeSize_;
340+ int64_t oneBlockNumG = blockSize_ / CONST_8;
341+ int64_t allUbNum = (static_cast<int64_t>(ubSize_) - groupIndexBuf);
342+ int64_t ubFactor = allUbNum / (CONST_7 * dtypeSize_);
343+ ubFactor = Ops::Base::FloorDiv(ubFactor, oneBlockNum) * oneBlockNum;
344+ hUbFactor_ = ubFactor;
345+ int64_t groupNumAlign = Ops::Base::CeilDiv(groupNum_, oneBlockNumG) * oneBlockNumG;
346+ int64_t groupUb = groupNumAlign * CONST_8;
347+ OP_CHECK_IF(
348+ (groupUb > (ubFactor * CONST_2) || ubFactor <= 0),
349+ OP_LOGE(
350+ context_, "ubFactor must > 0 and groupUb <= (ubFactor*CONST_2), but ubFactor is %ld, groupUb is %ld",
351+ ubFactor, groupUb),
352+ return ge::GRAPH_FAILED);
353+ if ((dimH_ < hUbFactor_) && (isInterleaved_ == 0)) {
354+ hUbFactor_ = Ops::Base::FloorDiv(dimH_, oneBlockNum) * oneBlockNum;
355+ if (hUbFactor_ <= 0) {
356+ hUbFactor_ = oneBlockNum;
357+ }
358+ bUbFactor_ = ubFactor / hUbFactor_;
359+ } else {
360+ bUbFactor_ = 1;
361+ }
362+ return ge::GRAPH_SUCCESS;
363+}
364+ 
365+void ClippedSwigluArch35Tiling::ComputeCoreSplit()
366+{
367+ if (isInterleaved_ != 0) {
368+ int64_t pairTotal = dim2H_ * dimBatchSize_ / CONST_2;
369+ int64_t blockFactor = (pairTotal + coreNumAll_ - 1) / coreNumAll_;
370+ realCoreNum_ = (pairTotal + blockFactor - 1) / blockFactor;
371+ } else {
372+ int64_t hCore = 1;
373+ int64_t bBlockFactor = (dimBatchSize_ + coreNumAll_ - 1) / coreNumAll_;
374+ int64_t bCore = (dimBatchSize_ + bBlockFactor - 1) / bBlockFactor;
375+ int64_t core = coreNumAll_ / bCore;
376+ if (core > 1) {
377+ hCore = core;
378+ int64_t hBlockFactor = (dimH_ + hCore - 1) / hCore;
379+ hCore = (dimH_ + hBlockFactor - 1) / hBlockFactor;
380+ }
381+ realCoreNum_ = bCore * hCore;
382+ if (isGroup_ == 1) {
383+ realCoreNum_ = coreNumAll_;
384+ }
385+ }
386+ return;
387+}
388+ 
389+void ClippedSwigluArch35Tiling::SetTilingKey()
390+{
391+ uint64_t isInterleavedKey = (isInterleaved_ != 0) ? TPL_INTERLEAVED_TRUE : TPL_INTERLEAVED_FALSE;
392+ uint64_t isGroupKey = (isGroup_ != 0) ? TPL_GROUP_INDEX : TPL_NO_GROUP_INDEX;
393+ OP_LOGI(context_->GetNodeName(), "isInterleavedKey = %lu, isGroupKey = %lu", isInterleavedKey, isGroupKey);
394+ tilingKey_ = GET_TPL_TILING_KEY(isInterleavedKey, isGroupKey);
395+}
396+ 
397+void ClippedSwigluArch35Tiling::FillTilingData()
398+{
399+ tilingData_->dimBatchSize = dimBatchSize_;
400+ tilingData_->dimH = dimH_;
401+ tilingData_->gluAlpha = gluAlpha_;
402+ tilingData_->gluLimit = gluLimit_;
403+ tilingData_->gluBias = gluBias_;
404+ tilingData_->hUbFactor = hUbFactor_;
405+ tilingData_->bUbFactor = bUbFactor_;
406+ tilingData_->groupNum = groupNum_;
407+ tilingData_->realCoreNum = realCoreNum_;
408+}
409+ 
410+void ClippedSwigluArch35Tiling::PrintTilingInfo()
411+{
412+ std::ostringstream info;
413+ info << "Print tilingData: tilingKey_: " << tilingKey_;
414+ info << ", coreNumAll: " << coreNumAll_;
415+ info << ", ubSize_: " << ubSize_;
416+ info << ", dimBatchSize: " << dimBatchSize_;
417+ info << ", dim2H: " << dim2H_;
418+ info << ", dimH: " << dimH_;
419+ info << ", isGroup: " << isGroup_;
420+ info << ", isInterleaved: " << isInterleaved_;
421+ info << ", gluLimit: " << gluLimit_;
422+ info << ", gluAlpha: " << gluAlpha_;
423+ info << ", gluBias: " << gluBias_;
424+ info << ", hUbFactor: " << hUbFactor_;
425+ info << ", bUbFactor: " << bUbFactor_;
426+ info << ", groupNum: " << groupNum_;
427+ info << ", realCoreNum: " << realCoreNum_;
428+ OP_LOGI(context_->GetNodeName(), "%s", info.str().c_str());
429+}
430+ 
431+ge::graphStatus Tiling4ClippedSwigluArch35(gert::TilingContext* context)
432+{
433+ OP_LOGI("ClippedSwigluArch35Tiling", "Enter Tiling4ClippedSwigluArch35");
434+ ClippedSwigluArch35Tiling tilingImpl = ClippedSwigluArch35Tiling(context);
435+ if (tilingImpl.Init() != ge::GRAPH_SUCCESS) {
436+ OP_LOGE(context, "Tiling4ClippedSwigluArch35 init failed.");
437+ return ge::GRAPH_FAILED;
438+ }
439+ if (tilingImpl.DoTiling() != ge::GRAPH_SUCCESS) {
440+ OP_LOGE(context, "Tiling4ClippedSwigluArch35 do tiling failed.");
441+ return ge::GRAPH_FAILED;
442+ }
443+ OP_LOGI("ClippedSwigluArch35Tiling", "Tiling4ClippedSwigluArch35 done.");
444+ return ge::GRAPH_SUCCESS;
445+}
446+ 
447+} // namespace optiling
Aactivation/clipped_swiglu/op_host/config/ascend950/clipped_swiglu_binary.json+206-0
@@ -0,0 +1,206 @@
1+{
2+ "op_type": "ClippedSwiglu",
3+ "op_list": [
4+ {
5+ "bin_filename": "ClippedSwiglu_fp16",
6+ "inputs": [
7+ {
8+ "name": "x",
9+ "index": 0,
10+ "dtype": "float16",
11+ "format": "ND",
12+ "paramType": "required",
13+ "shape": [
14+ -2
15+ ],
16+ "format_match_mode": "FormatAgnostic"
17+ },
18+ {
19+ "name": "group_index",
20+ "index": 1,
21+ "dtype": "int64",
22+ "format": "ND",
23+ "paramType": "optional",
24+ "shape": [
25+ -2
26+ ],
27+ "format_match_mode": "FormatAgnostic"
28+ }
29+ ],
30+ "outputs": [
31+ {
32+ "name": "y",
33+ "index": 0,
34+ "dtype": "float16",
35+ "format": "ND",
36+ "paramType": "required",
37+ "shape": [
38+ -2
39+ ],
40+ "format_match_mode": "FormatAgnostic"
41+ }
42+ ],
43+ "attrs": [
44+ {
45+ "name": "dim",
46+ "dtype": "int64",
47+ "value": null
48+ },
49+ {
50+ "name": "alpha",
51+ "dtype": "float32",
52+ "value": null
53+ },
54+ {
55+ "name": "limit",
56+ "dtype": "float32",
57+ "value": null
58+ },
59+ {
60+ "name": "bias",
61+ "dtype": "float32",
62+ "value": null
63+ },
64+ {
65+ "name": "interleaved",
66+ "dtype": "bool",
67+ "value": null
68+ }
69+ ]
70+ },
71+ {
72+ "bin_filename": "ClippedSwiglu_fp32",
73+ "inputs": [
74+ {
75+ "name": "x",
76+ "index": 0,
77+ "dtype": "float32",
78+ "format": "ND",
79+ "paramType": "required",
80+ "shape": [
81+ -2
82+ ],
83+ "format_match_mode": "FormatAgnostic"
84+ },
85+ {
86+ "name": "group_index",
87+ "index": 1,
88+ "dtype": "int64",
89+ "format": "ND",
90+ "paramType": "optional",
91+ "shape": [
92+ -2
93+ ],
94+ "format_match_mode": "FormatAgnostic"
95+ }
96+ ],
97+ "outputs": [
98+ {
99+ "name": "y",
100+ "index": 0,
101+ "dtype": "float32",
102+ "format": "ND",
103+ "paramType": "required",
104+ "shape": [
105+ -2
106+ ],
107+ "format_match_mode": "FormatAgnostic"
108+ }
109+ ],
110+ "attrs": [
111+ {
112+ "name": "dim",
113+ "dtype": "int64",
114+ "value": null
115+ },
116+ {
117+ "name": "alpha",
118+ "dtype": "float32",
119+ "value": null
120+ },
121+ {
122+ "name": "limit",
123+ "dtype": "float32",
124+ "value": null
125+ },
126+ {
127+ "name": "bias",
128+ "dtype": "float32",
129+ "value": null
130+ },
131+ {
132+ "name": "interleaved",
133+ "dtype": "bool",
134+ "value": null
135+ }
136+ ]
137+ },
138+ {
139+ "bin_filename": "ClippedSwiglu_bf16",
140+ "inputs": [
141+ {
142+ "name": "x",
143+ "index": 0,
144+ "dtype": "bfloat16",
145+ "format": "ND",
146+ "paramType": "required",
147+ "shape": [
148+ -2
149+ ],
150+ "format_match_mode": "FormatAgnostic"
151+ },
152+ {
153+ "name": "group_index",
154+ "index": 1,
155+ "dtype": "int64",
156+ "format": "ND",
157+ "paramType": "optional",
158+ "shape": [
159+ -2
160+ ],
161+ "format_match_mode": "FormatAgnostic"
162+ }
163+ ],
164+ "outputs": [
165+ {
166+ "name": "y",
167+ "index": 0,
168+ "dtype": "bfloat16",
169+ "format": "ND",
170+ "paramType": "required",
171+ "shape": [
172+ -2
173+ ],
174+ "format_match_mode": "FormatAgnostic"
175+ }
176+ ],
177+ "attrs": [
178+ {
179+ "name": "dim",
180+ "dtype": "int64",
181+ "value": null
182+ },
183+ {
184+ "name": "alpha",
185+ "dtype": "float32",
186+ "value": null
187+ },
188+ {
189+ "name": "limit",
190+ "dtype": "float32",
191+ "value": null
192+ },
193+ {
194+ "name": "bias",
195+ "dtype": "float32",
196+ "value": null
197+ },
198+ {
199+ "name": "interleaved",
200+ "dtype": "bool",
201+ "value": null
202+ }
203+ ]
204+ }
205+ ]
206+}
Aactivation/clipped_swiglu/op_kernel/arch35/clipped_swiglu_kernel.h+495-0
@@ -0,0 +1,495 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file clipped_swiglu_kernel.h
13+ * \brief Regbase kernel for ClippedSwiglu (Ascend 950 / DAV_3510)
14+ */
15+ 
16+#ifndef CLIPPED_SWIGLU_KERNEL_H
17+#define CLIPPED_SWIGLU_KERNEL_H
18+ 
19+#include "kernel_operator.h"
20+#include "op_kernel/math_util.h"
21+#include "op_kernel/platform_util.h"
22+#include "kernel_tiling/kernel_tiling.h"
23+#include "clipped_swiglu_tiling_key.h"
24+#include "clipped_swiglu_tiling_data.h"
25+ 
26+namespace ClippedSwigluOp {
27+ 
28+using namespace AscendC;
29+using namespace AscendC::MicroAPI;
30+ 
31+constexpr int64_t DB_BUFFER = 2;
32+constexpr int64_t BLOCK_SIZE = Ops::Base::GetUbBlockSize(); // 32
33+constexpr int64_t DIM_HALVE = 2;
34+constexpr uint32_t VF_LEN_FP32 = Ops::Base::GetVRegSize() / sizeof(float);
35+ 
36+static constexpr AscendC::MicroAPI::CastTrait CAST_BF16_FP16_TO_FP32 = {
37+ AscendC::MicroAPI::RegLayout::ZERO, AscendC::MicroAPI::SatMode::UNKNOWN, AscendC::MicroAPI::MaskMergeMode::ZEROING,
38+ AscendC::RoundMode::UNKNOWN};
39+constexpr static AscendC::MicroAPI::CastTrait CAST_FP32_TO_FP16_BF16 = {
40+ AscendC::MicroAPI::RegLayout::ZERO, AscendC::MicroAPI::SatMode::NO_SAT, AscendC::MicroAPI::MaskMergeMode::ZEROING,
41+ AscendC::RoundMode::CAST_RINT};
42+ 
43+__aicore__ inline void ReduceAllVf(
44+ LocalTensor<int64_t>& reduceSumUb, LocalTensor<int64_t>& groupIndexUb, int64_t groupIndexNum)
45+{
46+ uint32_t vfTidx = Ops::Base::GetVRegSize() / sizeof(int64_t);
47+ uint16_t times = groupIndexNum / vfTidx;
48+ uint32_t tailNum = groupIndexNum % vfTidx;
49+ uint16_t tailTimes = tailNum != 0 ? 1 : 0;
50+ auto dstAddr = (__ubuf__ int64_t*)reduceSumUb.GetPhyAddr();
51+ auto srcAddr = (__ubuf__ int64_t*)groupIndexUb.GetPhyAddr();
52+ auto srcAddr1 = (__ubuf__ int64_t*)groupIndexUb[times * vfTidx].GetPhyAddr();
53+ __VEC_SCOPE__
54+ {
55+ AscendC::MicroAPI::RegTensor<int64_t> addReg;
56+ AscendC::MicroAPI::RegTensor<int64_t> reduceSumReg;
57+ AscendC::MicroAPI::RegTensor<int64_t> reduceSumTReg;
58+ AscendC::MicroAPI::RegTensor<int64_t> srcReg;
59+ AscendC::MicroAPI::Duplicate(addReg, 0);
60+ AscendC::MicroAPI::MaskReg mask = AscendC::MicroAPI::CreateMask<int64_t, MicroAPI::MaskPattern::ALL>();
61+ for (uint16_t i = 0; i < times; i++) {
62+ AscendC::MicroAPI::AddrReg srcIdxOffset = AscendC::MicroAPI::CreateAddrReg<int64_t>(i, vfTidx);
63+ AscendC::MicroAPI::DataCopy(srcReg, srcAddr, srcIdxOffset);
64+ AscendC::MicroAPI::Add(addReg, addReg, srcReg, mask);
65+ }
66+ AscendC::MicroAPI::ReduceSum(reduceSumReg, addReg, mask);
67+ for (uint16_t j = 0; j < tailTimes; j++) {
68+ AscendC::MicroAPI::MaskReg maskT = AscendC::MicroAPI::UpdateMask<int64_t>(tailNum);
69+ AscendC::MicroAPI::DataCopy(srcReg, srcAddr1);
70+ AscendC::MicroAPI::ReduceSum(reduceSumTReg, srcReg, maskT);
71+ AscendC::MicroAPI::Add(reduceSumReg, reduceSumTReg, reduceSumReg, maskT);
72+ }
73+ AscendC::MicroAPI::MaskReg maskOne = AscendC::MicroAPI::CreateMask<int64_t, MicroAPI::MaskPattern::VL1>();
74+ AscendC::MicroAPI::DataCopy(dstAddr, reduceSumReg, maskOne);
75+ }
76+}
77+ 
78+template <typename T, bool isInterleaved, bool isGroup>
79+class ClippedSwigluKernel {
80+public:
81+ __aicore__ inline ClippedSwigluKernel(const ClippedSwigluArch35TilingData* tilingData, TPipe* pipe)
82+ : tiling_(tilingData), pipe_(pipe){};
83+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR groupIndex, GM_ADDR y);
84+ __aicore__ inline void Process();
85+ 
86+private:
87+ __aicore__ inline void ComputeTiling();
88+ __aicore__ inline void UbDataCopy(__local_mem__ T* inAddr, __local_mem__ T* outAddr, int64_t onceNum);
89+ __aicore__ inline void CopyIn(int64_t gmOffset, int64_t count, int64_t blockLen);
90+ __aicore__ inline void ComputeVfSwiglu(
91+ __local_mem__ T* x1UbAddr, __local_mem__ T* x2UbAddr, __local_mem__ T* swigluUbAddr, int64_t dim0OnceSize,
92+ int64_t dim1OnceSize, int64_t alignDim1In);
93+ __aicore__ inline void CopyOut(int64_t gmOffset, int64_t count, int64_t blockLen);
94+ 
95+private:
96+ GlobalTensor<T> xGm_;
97+ GlobalTensor<int64_t> groupIndexGm_;
98+ GlobalTensor<T> yGm_;
99+ const ClippedSwigluArch35TilingData* tiling_;
100+ TPipe* pipe_;
101+ 
102+ TQue<QuePosition::VECIN, 1> inQueX_;
103+ TQue<QuePosition::VECOUT, 1> outQueY_;
104+ TBuf<QuePosition::VECCALC> vectorBuf_;
105+ TBuf<QuePosition::VECCALC> reduceSumBuf_;
106+ 
107+ uint32_t blockIdx_ = 0;
108+ int64_t dimH_ = 0;
109+ int64_t dimB_ = 0;
110+ int64_t hUbFactor_ = 0;
111+ int64_t bUbFactor_ = 0;
112+ 
113+ int64_t realCoreNum_ = 0;
114+ int64_t hPreBlockNum_ = 0; // 每个核处理多少个数，注意尾核
115+ int64_t hLoopTimes_ = 0; // 每个核 h方向循环多少次
116+ int64_t hTailNum_ = 0; // 最后一次ub循环处理的数据量
117+ int64_t bPreBlockNum_ = 0; // 每个核处理多少个数，注意尾核
118+ int64_t bLoopTimes_ = 0; // 每个核 h方向循环多少次
119+ int64_t bTailNum_ = 0; // 最后一次ub循环处理的数据量
120+ int64_t hBlockFactor_ = 0;
121+ int64_t bBlockFactor_ = 0;
122+ int64_t hCore_ = 1;
123+ int64_t bCore_ = 0;
124+ int64_t vfLenT_ = 0;
125+ int64_t bBlockIdx_ = 0;
126+ int64_t hBlockIdx_ = 0;
127+ float limit_ = 0.0f;
128+ float alpha_ = 0.0f;
129+ float bias_ = 0.0f;
130+};
131+ 
132+template <typename T, bool isInterleaved, bool isGroup>
133+__aicore__ inline void ClippedSwigluKernel<T, isInterleaved, isGroup>::Init(GM_ADDR x, GM_ADDR groupIndex, GM_ADDR y)
134+{
135+ blockIdx_ = GetBlockIdx();
136+ realCoreNum_ = tiling_->realCoreNum;
137+ dimH_ = tiling_->dimH;
138+ hBlockFactor_ = dimH_;
139+ vfLenT_ = Ops::Base::GetVRegSize() / sizeof(T);
140+ hUbFactor_ = tiling_->hUbFactor;
141+ bUbFactor_ = tiling_->bUbFactor;
142+ limit_ = tiling_->gluLimit;
143+ alpha_ = tiling_->gluAlpha;
144+ bias_ = tiling_->gluBias;
145+ xGm_.SetGlobalBuffer(reinterpret_cast<__gm__ T*>(x));
146+ yGm_.SetGlobalBuffer(reinterpret_cast<__gm__ T*>(y));
147+ if constexpr (isGroup) {
148+ groupIndexGm_.SetGlobalBuffer(reinterpret_cast<__gm__ int64_t*>(groupIndex));
149+ }
150+ int64_t ub = hUbFactor_ * bUbFactor_ * sizeof(T);
151+ pipe_->InitBuffer(inQueX_, DB_BUFFER, ub * DIM_HALVE);
152+ pipe_->InitBuffer(outQueY_, DB_BUFFER, ub);
153+ pipe_->InitBuffer(vectorBuf_, ub);
154+ if constexpr (isGroup) {
155+ pipe_->InitBuffer(reduceSumBuf_, BLOCK_SIZE);
156+ }
157+}
158+ 
159+template <typename T, bool isInterleaved, bool isGroup>
160+__aicore__ inline void ClippedSwigluKernel<T, isInterleaved, isGroup>::ComputeTiling()
161+{
162+ int64_t dimBatchSize = tiling_->dimBatchSize;
163+ if constexpr (isGroup) {
164+ int64_t groupNum = tiling_->groupNum;
165+ LocalTensor<int64_t> reduceSumUb = reduceSumBuf_.Get<int64_t>();
166+ LocalTensor<int64_t> groupUb = inQueX_.AllocTensor<int64_t>();
167+ DataCopyExtParams copyParams = {1, 0, 0, 0, 0};
168+ DataCopyPadExtParams<int64_t> padParams = {false, 0, 0, 0};
169+ copyParams.blockLen = groupNum * sizeof(int64_t);
170+ DataCopyPad(groupUb, groupIndexGm_, copyParams, padParams);
171+ inQueX_.EnQue(groupUb);
172+ groupUb = inQueX_.DeQue<int64_t>();
173+ ReduceAllVf(reduceSumUb, groupUb, groupNum);
174+ inQueX_.FreeTensor(groupUb);
175+ event_t eventVS = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::V_S));
176+ SetFlag<HardEvent::V_S>(eventVS);
177+ WaitFlag<HardEvent::V_S>(eventVS);
178+ dimB_ = static_cast<int64_t>(reduceSumUb.GetValue(0));
179+ dimB_ = dimB_ < dimBatchSize ? dimB_ : dimBatchSize;
180+ } else {
181+ dimB_ = dimBatchSize;
182+ }
183+ 
184+ if constexpr (isInterleaved) {
185+ int64_t pairTotal = dimH_ * dimB_;
186+ hBlockFactor_ = (pairTotal + realCoreNum_ - 1) / realCoreNum_;
187+ realCoreNum_ = (pairTotal + hBlockFactor_ - 1) / hBlockFactor_;
188+ int64_t tailBlockNum = pairTotal - hBlockFactor_ * (realCoreNum_ - 1);
189+ hPreBlockNum_ = blockIdx_ == (realCoreNum_ - 1) ? tailBlockNum : hBlockFactor_;
190+ hLoopTimes_ = (hPreBlockNum_ + hUbFactor_ - 1) / hUbFactor_;
191+ hTailNum_ = hPreBlockNum_ - hUbFactor_ * (hLoopTimes_ - 1);
192+ } else {
193+ bBlockFactor_ = (dimB_ + realCoreNum_ - 1) / realCoreNum_;
194+ bCore_ = (dimB_ + bBlockFactor_ - 1) / bBlockFactor_;
195+ int64_t core = realCoreNum_ / bCore_;
196+ if (core > 1) {
197+ hBlockFactor_ = (dimH_ + core - 1) / core;
198+ hCore_ = (dimH_ + hBlockFactor_ - 1) / hBlockFactor_;
199+ }
200+ realCoreNum_ = bCore_ * hCore_;
201+ bBlockIdx_ = blockIdx_ / hCore_;
202+ hBlockIdx_ = blockIdx_ % hCore_;
203+ 
204+ int64_t tailHcoreNum = dimH_ - hBlockFactor_ * (hCore_ - 1);
205+ hPreBlockNum_ = hBlockIdx_ == (hCore_ - 1) ? tailHcoreNum : hBlockFactor_;
206+ hLoopTimes_ = (hPreBlockNum_ + hUbFactor_ - 1) / hUbFactor_;
207+ hTailNum_ = hPreBlockNum_ - hUbFactor_ * (hLoopTimes_ - 1);
208+ 
209+ int64_t tailBcoreNum = dimB_ - bBlockFactor_ * (bCore_ - 1);
210+ bPreBlockNum_ = bBlockIdx_ == (bCore_ - 1) ? tailBcoreNum : bBlockFactor_;
211+ bLoopTimes_ = (bPreBlockNum_ + bUbFactor_ - 1) / bUbFactor_;
212+ bTailNum_ = bPreBlockNum_ - bUbFactor_ * (bLoopTimes_ - 1);
213+ }
214+}
215+ 
216+template <typename T, bool isInterleaved, bool isGroup>
217+__aicore__ inline void ClippedSwigluKernel<T, isInterleaved, isGroup>::Process()
218+{
219+ ComputeTiling();
220+ if (blockIdx_ >= realCoreNum_) {
221+ return;
222+ }
223+ LocalTensor<T> swigluUb = vectorBuf_.Get<T>();
224+ int64_t oneBlockNum = BLOCK_SIZE / sizeof(T);
225+ if constexpr (isInterleaved) {
226+ int64_t blockOffset = blockIdx_ * hBlockFactor_;
227+ for (int64_t loopIdx = 0; loopIdx < hLoopTimes_; ++loopIdx) {
228+ int64_t onceNum = loopIdx == (hLoopTimes_ - 1) ? hTailNum_ : hUbFactor_;
229+ int64_t outGmOffset = blockOffset + loopIdx * hUbFactor_;
230+ int64_t inGmOffset = outGmOffset * DIM_HALVE;
231+ CopyIn(inGmOffset, 1, onceNum * DIM_HALVE);
232+ LocalTensor<T> inputUb = inQueX_.DeQue<T>();
233+ auto actAddr = (__ubuf__ T*)inputUb.GetPhyAddr();
234+ auto gateAddr = (__ubuf__ T*)inputUb[VF_LEN_FP32].GetPhyAddr();
235+ auto swigluAddr = (__ubuf__ T*)swigluUb.GetPhyAddr();
236+ int64_t alignDim1In = ((onceNum * DIM_HALVE + oneBlockNum - 1) / oneBlockNum) * oneBlockNum;
237+ ComputeVfSwiglu(actAddr, gateAddr, swigluAddr, 1, onceNum, alignDim1In);
238+ inQueX_.FreeTensor(inputUb);
239+ LocalTensor<T> outUb = outQueY_.AllocTensor<T>();
240+ auto outUbAddr = (__ubuf__ T*)outUb.GetPhyAddr();
241+ UbDataCopy(swigluAddr, outUbAddr, onceNum);
242+ outQueY_.EnQue(outUb);
243+ CopyOut(outGmOffset, 1, onceNum);
244+ }
245+ } else {
246+ int64_t bCoreOffset = bBlockFactor_ * bBlockIdx_;
247+ int64_t hCoreOffset = hBlockFactor_ * hBlockIdx_;
248+ for (int64_t bi = 0; bi < bLoopTimes_; bi++) {
249+ int64_t onceBNum = bi == (bLoopTimes_ - 1) ? bTailNum_ : bUbFactor_;
250+ for (int64_t hi = 0; hi < hLoopTimes_; hi++) {
251+ int64_t onceHNum = hi == (hLoopTimes_ - 1) ? hTailNum_ : hUbFactor_;
252+ int64_t alignDim1In = ((onceHNum + oneBlockNum - 1) / oneBlockNum) * oneBlockNum;
253+ int64_t inGmOffset =
254+ (bCoreOffset + bi * bUbFactor_) * dimH_ * DIM_HALVE + hCoreOffset + hi * hUbFactor_;
255+ CopyIn(inGmOffset, onceBNum, onceHNum);
256+ LocalTensor<T> inputUb = inQueX_.DeQue<T>();
257+ auto actAddr = (__ubuf__ T*)inputUb.GetPhyAddr();
258+ auto gateAddr = (__ubuf__ T*)inputUb[hUbFactor_ * bUbFactor_].GetPhyAddr();
259+ auto swigluAddr = (__ubuf__ T*)swigluUb.GetPhyAddr();
260+ ComputeVfSwiglu(actAddr, gateAddr, swigluAddr, onceBNum, onceHNum, alignDim1In);
261+ inQueX_.FreeTensor(inputUb);
262+ LocalTensor<T> outUb = outQueY_.AllocTensor<T>();
263+ auto outUbAddr = (__ubuf__ T*)outUb.GetPhyAddr();
264+ UbDataCopy(swigluAddr, outUbAddr, onceBNum * alignDim1In);
265+ outQueY_.EnQue(outUb);
266+ int64_t outGmOffset = (bCoreOffset + bi * bUbFactor_) * dimH_ + hCoreOffset + hi * hUbFactor_;
267+ CopyOut(outGmOffset, onceBNum, onceHNum);
268+ }
269+ }
270+ }
271+}
272+ 
273+template <typename T, bool isInterleaved, bool isGroup>
274+__aicore__ inline void ClippedSwigluKernel<T, isInterleaved, isGroup>::CopyIn(
275+ int64_t gmOffset, int64_t count, int64_t blockLen)
276+{
277+ LocalTensor<T> xDTypeUb = inQueX_.AllocTensor<T>();
278+ DataCopyPadExtParams<T> padParams = {false, 0, 0, 0};
279+ DataCopyExtParams copyParams = {1, 0, 0, 0, 0};
280+ if constexpr (isInterleaved) {
281+ copyParams.blockCount = 1;
282+ copyParams.blockLen = blockLen * sizeof(T);
283+ copyParams.srcStride = 0;
284+ copyParams.dstStride = 0;
285+ DataCopyPad(xDTypeUb, xGm_[gmOffset], copyParams, padParams);
286+ } else {
287+ copyParams.blockCount = count;
288+ copyParams.blockLen = blockLen * sizeof(T);
289+ copyParams.srcStride = (dimH_ * DIM_HALVE - blockLen) * sizeof(T);
290+ copyParams.dstStride = 0;
291+ DataCopyPad(xDTypeUb, xGm_[gmOffset], copyParams, padParams);
292+ DataCopyPad(xDTypeUb[bUbFactor_ * hUbFactor_], xGm_[gmOffset + dimH_], copyParams, padParams);
293+ }
294+ inQueX_.EnQue(xDTypeUb);
295+}
296+ 
297+template <typename T, bool isInterleaved, bool isGroup>
298+__aicore__ inline void ClippedSwigluKernel<T, isInterleaved, isGroup>::UbDataCopy(
299+ __local_mem__ T* inAddr, __local_mem__ T* outAddr, int64_t onceNum)
300+{
301+ uint32_t size = onceNum;
302+ uint32_t vfLen = vfLenT_;
303+ uint16_t times = CeilDivision(size, vfLen);
304+ __VEC_SCOPE__
305+ {
306+ AscendC::MicroAPI::RegTensor<T> xReg;
307+ AscendC::MicroAPI::MaskReg mask;
308+ for (uint16_t i = 0; i < times; i++) {
309+ mask = MicroAPI::UpdateMask<T>(size);
310+ AscendC::MicroAPI::AddrReg srcIdxOffset = AscendC::MicroAPI::CreateAddrReg<T>(i, vfLen);
311+ AscendC::MicroAPI::DataCopy(xReg, inAddr, srcIdxOffset);
312+ AscendC::MicroAPI::DataCopy(outAddr, xReg, srcIdxOffset, mask);
313+ }
314+ }
315+}
316+ 
317+template <typename T, bool isInterleaved, bool isGroup>
318+__aicore__ inline void ClippedSwigluKernel<T, isInterleaved, isGroup>::ComputeVfSwiglu(
319+ __local_mem__ T* x1UbAddr, __local_mem__ T* x2UbAddr, __local_mem__ T* swigluUbAddr, int64_t dim0OnceSize,
320+ int64_t dim1OnceSize, int64_t alignDim1In)
321+{
322+ float clampLimit = limit_;
323+ float negClampLimit = -limit_;
324+ float negAlpha = -alpha_;
325+ float gluBias = bias_;
326+ uint16_t dim0VfTimes = 1;
327+ float scalarOne = 1.0f;
328+ uint32_t vfLen = VF_LEN_FP32 * DIM_HALVE;
329+ if constexpr (!isInterleaved) {
330+ dim0VfTimes = dim0OnceSize;
331+ vfLen = VF_LEN_FP32;
332+ }
333+ int64_t oneBlockNum = BLOCK_SIZE / sizeof(T);
334+ int64_t alignDim1Out = 0;
335+ if constexpr (isInterleaved) {
336+ alignDim1Out = ((dim1OnceSize + oneBlockNum - 1) / oneBlockNum) * oneBlockNum;
337+ } else {
338+ alignDim1Out = alignDim1In;
339+ }
340+ uint16_t dim1VfTimes = dim1OnceSize / VF_LEN_FP32;
341+ uint32_t tail = dim1OnceSize % VF_LEN_FP32;
342+ uint16_t tailTimes = 0;
343+ if (tail > 0) {
344+ tailTimes = 1;
345+ }
346+ __local_mem__ T* x1UbAddrT = x1UbAddr + dim1VfTimes * vfLen;
347+ __local_mem__ T* x2UbAddrT = x2UbAddr + dim1VfTimes * vfLen;
348+ __local_mem__ T* swigluUbAddrT = swigluUbAddr + dim1VfTimes * VF_LEN_FP32;
349+ __VEC_SCOPE__
350+ {
351+ AscendC::MicroAPI::RegTensor<T> vregX1;
352+ AscendC::MicroAPI::RegTensor<T> vregX2;
353+ AscendC::MicroAPI::RegTensor<float> vregX1F;
354+ AscendC::MicroAPI::RegTensor<float> vregX2F;
355+ 
356+ AscendC::MicroAPI::RegTensor<float> vregX1DeF;
357+ AscendC::MicroAPI::RegTensor<float> vregX2DeF;
358+ AscendC::MicroAPI::RegTensor<float> minsReg;
359+ AscendC::MicroAPI::RegTensor<float> mulsReg;
360+ AscendC::MicroAPI::RegTensor<float> expReg;
361+ AscendC::MicroAPI::RegTensor<float> addsReg;
362+ AscendC::MicroAPI::RegTensor<float> sigmoidReg;
363+ AscendC::MicroAPI::RegTensor<float> outFReg;
364+ AscendC::MicroAPI::RegTensor<T> outTReg;
365+ MicroAPI::MaskReg mask = MicroAPI::CreateMask<float, MicroAPI::MaskPattern::ALL>();
366+ MicroAPI::MaskReg maskT = MicroAPI::UpdateMask<float>(tail);
367+ for (uint16_t dim0vfLoopIdx = 0; dim0vfLoopIdx < dim0VfTimes; dim0vfLoopIdx++) {
368+ for (uint16_t dim1vfLoopIdx = 0; dim1vfLoopIdx < dim1VfTimes; dim1vfLoopIdx++) {
369+ AscendC::MicroAPI::AddrReg srcIdxOffset =
370+ AscendC::MicroAPI::CreateAddrReg<T>(dim0vfLoopIdx, alignDim1In, dim1vfLoopIdx, vfLen);
371+ if constexpr (isInterleaved) {
372+ if constexpr (sizeof(T) == sizeof(half)) {
373+ AscendC::MicroAPI::DataCopy<T, AscendC::MicroAPI::LoadDist::DIST_UNPACK_B16>(
374+ vregX1, x1UbAddr, srcIdxOffset);
375+ AscendC::MicroAPI::DataCopy<T, AscendC::MicroAPI::LoadDist::DIST_UNPACK_B16>(
376+ vregX2, x2UbAddr, srcIdxOffset);
377+ AscendC::MicroAPI::Cast<float, T, CAST_BF16_FP16_TO_FP32>(vregX1F, vregX1, mask);
378+ AscendC::MicroAPI::Cast<float, T, CAST_BF16_FP16_TO_FP32>(vregX2F, vregX2, mask);
379+ } else {
380+ // float
381+ AscendC::MicroAPI::DataCopy((MicroAPI::RegTensor<T>&)vregX1F, x1UbAddr, srcIdxOffset);
382+ AscendC::MicroAPI::DataCopy((MicroAPI::RegTensor<T>&)vregX2F, x2UbAddr, srcIdxOffset);
383+ }
384+ AscendC::MicroAPI::DeInterleave(vregX1DeF, vregX2DeF, vregX1F, vregX2F);
385+ } else {
386+ if constexpr (sizeof(T) == sizeof(half)) {
387+ AscendC::MicroAPI::DataCopy<T, AscendC::MicroAPI::LoadDist::DIST_UNPACK_B16>(
388+ vregX1, x1UbAddr, srcIdxOffset);
389+ AscendC::MicroAPI::DataCopy<T, AscendC::MicroAPI::LoadDist::DIST_UNPACK_B16>(
390+ vregX2, x2UbAddr, srcIdxOffset);
391+ AscendC::MicroAPI::Cast<float, T, CAST_BF16_FP16_TO_FP32>(vregX1DeF, vregX1, mask);
392+ AscendC::MicroAPI::Cast<float, T, CAST_BF16_FP16_TO_FP32>(vregX2DeF, vregX2, mask);
393+ } else {
394+ // float
395+ AscendC::MicroAPI::DataCopy((MicroAPI::RegTensor<T>&)vregX1DeF, x1UbAddr, srcIdxOffset);
396+ AscendC::MicroAPI::DataCopy((MicroAPI::RegTensor<T>&)vregX2DeF, x2UbAddr, srcIdxOffset);
397+ }
398+ }
399+ AscendC::MicroAPI::Mins(minsReg, vregX1DeF, clampLimit, mask);
400+ AscendC::MicroAPI::Muls(mulsReg, minsReg, negAlpha, mask);
401+ AscendC::MicroAPI::Exp(expReg, mulsReg, mask);
402+ AscendC::MicroAPI::Adds(addsReg, expReg, scalarOne, mask);
403+ AscendC::MicroAPI::Div(sigmoidReg, minsReg, addsReg, mask);
404+ 
405+ AscendC::MicroAPI::Mins(vregX2DeF, vregX2DeF, clampLimit, mask);
406+ AscendC::MicroAPI::Maxs(vregX2DeF, vregX2DeF, negClampLimit, mask);
407+ AscendC::MicroAPI::Adds(vregX2DeF, vregX2DeF, gluBias, mask);
408+ 
409+ AscendC::MicroAPI::Mul(outFReg, sigmoidReg, vregX2DeF, mask);
410+ AscendC::MicroAPI::AddrReg outOffset =
411+ AscendC::MicroAPI::CreateAddrReg<T>(dim0vfLoopIdx, alignDim1Out, dim1vfLoopIdx, VF_LEN_FP32);
412+ if constexpr (sizeof(T) == sizeof(half)) {
413+ AscendC::MicroAPI::Cast<T, float, CAST_FP32_TO_FP16_BF16>(outTReg, outFReg, mask);
414+ DataCopy<T, AscendC::MicroAPI::StoreDist::DIST_PACK_B32>(swigluUbAddr, outTReg, outOffset, mask);
415+ } else {
416+ DataCopy(swigluUbAddr, (MicroAPI::RegTensor<T>&)outFReg, outOffset, mask);
417+ }
418+ }
419+ AscendC::MicroAPI::AddrReg srcIdxOffset1 = AscendC::MicroAPI::CreateAddrReg<T>(dim0vfLoopIdx, alignDim1In);
420+ AscendC::MicroAPI::AddrReg outOffset1 = AscendC::MicroAPI::CreateAddrReg<T>(dim0vfLoopIdx, alignDim1Out);
421+ for (uint16_t ti = 0; ti < tailTimes; ti++) {
422+ if constexpr (isInterleaved) {
423+ if constexpr (sizeof(T) == sizeof(half)) {
424+ AscendC::MicroAPI::DataCopy<T, AscendC::MicroAPI::LoadDist::DIST_UNPACK_B16>(
425+ vregX1, x1UbAddrT, srcIdxOffset1);
426+ AscendC::MicroAPI::DataCopy<T, AscendC::MicroAPI::LoadDist::DIST_UNPACK_B16>(
427+ vregX2, x2UbAddrT, srcIdxOffset1);
428+ AscendC::MicroAPI::Cast<float, T, CAST_BF16_FP16_TO_FP32>(vregX1F, vregX1, mask);
429+ AscendC::MicroAPI::Cast<float, T, CAST_BF16_FP16_TO_FP32>(vregX2F, vregX2, mask);
430+ } else {
431+ // float
432+ AscendC::MicroAPI::DataCopy((MicroAPI::RegTensor<T>&)vregX1F, x1UbAddrT, srcIdxOffset1);
433+ AscendC::MicroAPI::DataCopy((MicroAPI::RegTensor<T>&)vregX2F, x2UbAddrT, srcIdxOffset1);
434+ }
435+ AscendC::MicroAPI::DeInterleave(vregX1DeF, vregX2DeF, vregX1F, vregX2F);
436+ } else {
437+ if constexpr (sizeof(T) == sizeof(half)) {
438+ AscendC::MicroAPI::DataCopy<T, AscendC::MicroAPI::LoadDist::DIST_UNPACK_B16>(
439+ vregX1, x1UbAddrT, srcIdxOffset1);
440+ AscendC::MicroAPI::DataCopy<T, AscendC::MicroAPI::LoadDist::DIST_UNPACK_B16>(
441+ vregX2, x2UbAddrT, srcIdxOffset1);
442+ AscendC::MicroAPI::Cast<float, T, CAST_BF16_FP16_TO_FP32>(vregX1DeF, vregX1, mask);
443+ AscendC::MicroAPI::Cast<float, T, CAST_BF16_FP16_TO_FP32>(vregX2DeF, vregX2, mask);
444+ } else {
445+ // float
446+ AscendC::MicroAPI::DataCopy((MicroAPI::RegTensor<T>&)vregX1DeF, x1UbAddrT, srcIdxOffset1);
447+ AscendC::MicroAPI::DataCopy((MicroAPI::RegTensor<T>&)vregX2DeF, x2UbAddrT, srcIdxOffset1);
448+ }
449+ }
450+ AscendC::MicroAPI::Mins(minsReg, vregX1DeF, clampLimit, maskT);
451+ AscendC::MicroAPI::Muls(mulsReg, minsReg, negAlpha, maskT);
452+ AscendC::MicroAPI::Exp(expReg, mulsReg, maskT);
453+ AscendC::MicroAPI::Adds(addsReg, expReg, scalarOne, maskT);
454+ AscendC::MicroAPI::Div(sigmoidReg, minsReg, addsReg, maskT);
455+ 
456+ AscendC::MicroAPI::Mins(vregX2DeF, vregX2DeF, clampLimit, maskT);
457+ AscendC::MicroAPI::Maxs(vregX2DeF, vregX2DeF, negClampLimit, maskT);
458+ AscendC::MicroAPI::Adds(vregX2DeF, vregX2DeF, gluBias, maskT);
459+ 
460+ AscendC::MicroAPI::Mul(outFReg, sigmoidReg, vregX2DeF, maskT);
461+ if constexpr (sizeof(T) == sizeof(half)) {
462+ AscendC::MicroAPI::Cast<T, float, CAST_FP32_TO_FP16_BF16>(outTReg, outFReg, maskT);
463+ DataCopy<T, AscendC::MicroAPI::StoreDist::DIST_PACK_B32>(swigluUbAddrT, outTReg, outOffset1, maskT);
464+ } else {
465+ DataCopy(swigluUbAddrT, (MicroAPI::RegTensor<T>&)outFReg, outOffset1, maskT);
466+ }
467+ }
468+ }
469+ }
470+}
471+ 
472+template <typename T, bool isInterleaved, bool isGroup>
473+__aicore__ inline void ClippedSwigluKernel<T, isInterleaved, isGroup>::CopyOut(
474+ int64_t gmOffset, int64_t count, int64_t blockLen)
475+{
476+ LocalTensor<T> outputUb = outQueY_.DeQue<T>();
477+ outQueY_.EnQue(outputUb);
478+ 
479+ DataCopyExtParams copyParams = {1, 0, 0, 0, 0};
480+ if constexpr (isInterleaved) {
481+ copyParams.blockCount = 1;
482+ copyParams.dstStride = 0;
483+ } else {
484+ copyParams.blockCount = count;
485+ copyParams.dstStride = (dimH_ - blockLen) * sizeof(T);
486+ }
487+ copyParams.blockLen = blockLen * sizeof(T);
488+ copyParams.srcStride = 0;
489+ 
490+ DataCopyPad(yGm_[gmOffset], outputUb, copyParams);
491+ outQueY_.FreeTensor(outputUb);
492+}
493+} // namespace ClippedSwigluOp
494+ 
495+#endif // CLIPPED_SWIGLU_KERNEL_H
Aactivation/clipped_swiglu/op_kernel/arch35/clipped_swiglu_tiling_data.h+32-0
@@ -0,0 +1,32 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file clipped_swiglu_tiling_data.h
13+ * \brief TilingData struct for ClippedSwiglu Arch35
14+ */
15+ 
16+#ifndef CLIPPED_SWIGLU_TILING_DATA_H
17+#define CLIPPED_SWIGLU_TILING_DATA_H
18+ 
19+#include "kernel_tiling/kernel_tiling.h"
20+ 
21+struct ClippedSwigluArch35TilingData {
22+ int64_t realCoreNum;
23+ int64_t dimBatchSize;
24+ int64_t dimH;
25+ int64_t groupNum;
26+ int64_t bUbFactor;
27+ int64_t hUbFactor;
28+ float gluAlpha;
29+ float gluLimit;
30+ float gluBias;
31+};
32+#endif // CLIPPED_SWIGLU_TILING_DATA_H
Aactivation/clipped_swiglu/op_kernel/arch35/clipped_swiglu_tiling_key.h+38-0
@@ -0,0 +1,38 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file clipped_swiglu_tiling_key.h
13+ * \brief TPL tiling key for ClippedSwiglu (Ascend 950)
14+ */
15+ 
16+#ifndef CLIPPED_SWIGLU_TILING_KEY_H
17+#define CLIPPED_SWIGLU_TILING_KEY_H
18+ 
19+#include "ascendc/host_api/tiling/template_argument.h"
20+ 
21+#define TPL_INTERLEAVED_FALSE 0
22+#define TPL_INTERLEAVED_TRUE 1
23+ 
24+#define TPL_NO_GROUP_INDEX 0
25+#define TPL_GROUP_INDEX 1
26+ 
27+namespace ClippedSwigluOp {
28+ASCENDC_TPL_ARGS_DECL(
29+ ClippedSwiglu,
30+ ASCENDC_TPL_UINT_DECL(isInterleaved, 1, ASCENDC_TPL_UI_LIST, TPL_INTERLEAVED_FALSE, TPL_INTERLEAVED_TRUE),
31+ ASCENDC_TPL_UINT_DECL(isGroup, 1, ASCENDC_TPL_UI_LIST, TPL_NO_GROUP_INDEX, TPL_GROUP_INDEX));
32+ 
33+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(
34+ 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)));
36+} // namespace ClippedSwigluOp
37+ 
38+#endif // CLIPPED_SWIGLU_TILING_KEY_H
Aactivation/clipped_swiglu/op_kernel/clipped_swiglu_apt.cpp+56-0
@@ -0,0 +1,56 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file clipped_swiglu_apt.cpp
13+ * \brief Kernel entry for ClippedSwiglu Arch35 (Ascend 950)
14+ */
15+ 
16+#include "kernel_tiling/kernel_tiling.h"
17+#include "kernel_operator.h"
18+#include "arch35/clipped_swiglu_tiling_key.h"
19+#include "arch35/clipped_swiglu_tiling_data.h"
20+#include "arch35/clipped_swiglu_kernel.h"
21+ 
22+using namespace AscendC;
23+using namespace ClippedSwigluOp;
24+ 
25+template <uint64_t isInterleaved, uint64_t isGroup>
26+__global__ __aicore__ void clipped_swiglu(GM_ADDR x, GM_ADDR groupIndex, GM_ADDR y,
27+ GM_ADDR workspace, GM_ADDR tiling)
28+{
29+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
30+ REGISTER_TILING_DEFAULT(ClippedSwigluArch35TilingData);
31+ GET_TILING_DATA_WITH_STRUCT(ClippedSwigluArch35TilingData, tilingData, tiling);
32+ GM_ADDR usrWorkspace = AscendC::GetUserWorkspace(workspace);
33+ TPipe pipe;
34+ 
35+ if constexpr (isInterleaved == 1) {
36+ if constexpr (isGroup == 1) {
37+ ClippedSwigluKernel<DTYPE_X, true, true> op(&tilingData, &pipe);
38+ op.Init(x, groupIndex, y);
39+ op.Process();
40+ } else {
41+ ClippedSwigluKernel<DTYPE_X, true, false> op(&tilingData, &pipe);
42+ op.Init(x, groupIndex, y);
43+ op.Process();
44+ }
45+ } else {
46+ if constexpr (isGroup == 1) {
47+ ClippedSwigluKernel<DTYPE_X, false, true> op(&tilingData, &pipe);
48+ op.Init(x, groupIndex, y);
49+ op.Process();
50+ } else {
51+ ClippedSwigluKernel<DTYPE_X, false, false> op(&tilingData, &pipe);
52+ op.Init(x, groupIndex, y);
53+ op.Process();
54+ }
55+ }
56+}
Aactivation/clipped_swiglu/tests/ut/op_host/arch35/test_clipped_swiglu_tiling_arch35.cpp+327-0
@@ -0,0 +1,327 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file test_clipped_swiglu_tiling_arch35.cpp
13+ * \brief Arch35 (Ascend 950) tiling UT for ClippedSwiglu
14+ */
15+ 
16+#include <iostream>
17+#include <vector>
18+#include <gtest/gtest.h>
19+#include "log/log.h"
20+#include "array_ops.h"
21+#include "../../../../op_host/clipped_swiglu_tiling.h"
22+#include "kernel_run_context_facker.h"
23+#include "test_cube_util.h"
24+#include "exe_graph/runtime/storage_format.h"
25+#include "exe_graph/runtime/storage_shape.h"
26+#include "ut_op_util.h"
27+#include "platform/platform_infos_def.h"
28+#include "register/op_impl_registry.h"
29+#include "error_util.h"
30+ 
31+using namespace ut_util;
32+using namespace std;
33+using namespace ge;
34+ 
35+namespace optiling {
36+struct ClippedSwigluCompileInfo {};
37+} // namespace optiling
38+ 
39+class ClippedSwigluArch35TilingTest : public testing::Test {
40+protected:
41+ static void SetUpTestCase()
42+ {
43+ std::cout << "ClippedSwigluArch35TilingTest SetUp" << std::endl;
44+ }
45+ 
46+ static void TearDownTestCase()
47+ {
48+ std::cout << "ClippedSwigluArch35TilingTest TearDown" << std::endl;
49+ }
50+};
51+ 
52+static const string kCompileInfoStr = R"({
53+ "hardware_info": {
54+ "BT_SIZE": 0, "load3d_constraints": "1",
55+ "Intrinsic_fix_pipe_l0c2out": false, "Intrinsic_data_move_l12ub": true,
56+ "Intrinsic_data_move_l0c2ub": true, "Intrinsic_data_move_out2l1_nd2nz": false,
57+ "UB_SIZE": 245760, "L2_SIZE": 33554432, "L1_SIZE": 524288,
58+ "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072,
59+ "CORE_NUM": 64, "socVersion": "Ascend950"}
60+})";
61+ 
62+static const map<string, string> kSocVersion = {
63+ {"NpuArch", "3510"}, {"Short_SoC_version", "ASCEND950"}};
64+ 
65+struct Arch35TilingTestParam {
66+ gert::StorageShape xShape;
67+ gert::StorageShape* groupShape = nullptr;
68+ gert::StorageShape yShape;
69+ ge::DataType xDtype = ge::DT_FLOAT16;
70+ ge::DataType groupDtype = ge::DT_INT64;
71+ int64_t dim = -1;
72+ float alpha = 1.702f;
73+ float limit = 7.0f;
74+ float bias = 1.0f;
75+ bool interleaved = false;
76+ ge::graphStatus expectedStatus = ge::GRAPH_SUCCESS;
77+};
78+ 
79+static void RunArch35TilingTest(const Arch35TilingTestParam& tc)
80+{
81+ map<string, string> soc_infos;
82+ map<string, string> aicore_spec;
83+ map<string, string> intrinsics;
84+ GetPlatFormInfos(kCompileInfoStr.c_str(), soc_infos, aicore_spec, intrinsics, kSocVersion);
85+ 
86+ fe::PlatFormInfos platform_info;
87+ platform_info.Init();
88+ 
89+ optiling::ClippedSwigluCompileInfo compile_info;
90+ string op_type("ClippedSwiglu");
91+ ASSERT_NE(gert::OpImplRegistry::GetInstance().GetOpImpl(op_type.c_str()), nullptr);
92+ auto tiling_func = gert::OpImplRegistry::GetInstance().GetOpImpl(op_type.c_str())->tiling;
93+ auto tiling_parse_func =
94+ gert::OpImplRegistry::GetInstance().GetOpImpl(op_type.c_str())->tiling_parse;
95+ 
96+ auto kernel_holder = gert::KernelRunContextFaker()
97+ .KernelIONum(2, 1)
98+ .Inputs({const_cast<char*>(kCompileInfoStr.c_str()),
99+ reinterpret_cast<void*>(&platform_info)})
100+ .Outputs({&compile_info})
101+ .Build();
102+ 
103+ ASSERT_TRUE(kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->Init());
104+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes(
105+ "version", kSocVersion);
106+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes(
107+ "SoCInfo", soc_infos);
108+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes(
109+ "AICoreSpec", aicore_spec);
110+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetCoreNumByCoreType(
111+ "AICore");
112+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes(
113+ "AICoreintrinsicDtypeMap", intrinsics);
114+ 
115+ ASSERT_EQ(tiling_parse_func(kernel_holder.GetContext<gert::KernelContext>()), ge::GRAPH_SUCCESS);
116+ 
117+ auto param = gert::TilingData::CreateCap(4096);
118+ auto workspace_size_holer = gert::ContinuousVector::Create<size_t>(4096);
119+ auto ws_size = reinterpret_cast<gert::ContinuousVector*>(workspace_size_holer.get());
120+ ASSERT_NE(param, nullptr);
121+ 
122+ auto holder = gert::TilingContextFaker()
123+ .SetOpType("ClippedSwiglu")
124+ .NodeIoNum(2, 1)
125+ .IrInstanceNum({1, 1})
126+ .InputShapes({&tc.xShape, tc.groupShape})
127+ .OutputShapes({&tc.yShape})
128+ .CompileInfo(&compile_info)
129+ .PlatformInfo(reinterpret_cast<char*>(&platform_info))
130+ .NodeInputTd(0, tc.xDtype, ge::FORMAT_ND, ge::FORMAT_ND)
131+ .NodeInputTd(1, tc.groupDtype, ge::FORMAT_ND, ge::FORMAT_ND)
132+ .NodeOutputTd(0, tc.xDtype, ge::FORMAT_ND, ge::FORMAT_ND)
133+ .NodeAttrs({{"dim", Ops::NN::AnyValue::CreateFrom<int64_t>(tc.dim)},
134+ {"alpha", Ops::NN::AnyValue::CreateFrom<float>(tc.alpha)},
135+ {"limit", Ops::NN::AnyValue::CreateFrom<float>(tc.limit)},
136+ {"bias", Ops::NN::AnyValue::CreateFrom<float>(tc.bias)},
137+ {"interleaved", Ops::NN::AnyValue::CreateFrom<bool>(tc.interleaved)}})
138+ .TilingData(param.get())
139+ .Workspace(ws_size)
140+ .Build();
141+ 
142+ gert::TilingContext* tiling_context = holder.GetContext<gert::TilingContext>();
143+ ASSERT_NE(tiling_context, nullptr);
144+ ASSERT_NE(tiling_context->GetPlatformInfo(), nullptr);
145+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", soc_infos);
146+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes(
147+ "AICoreSpec", aicore_spec);
148+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
149+ holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes(
150+ "AICoreintrinsicDtypeMap", intrinsics);
151+ 
152+ EXPECT_EQ(tiling_func(tiling_context), tc.expectedStatus);
153+}
154+ 
155+// ========== Normal cases: half (interleaved=false), ungrouped ==========
156+ 
157+TEST_F(ClippedSwigluArch35TilingTest, arch35_fp16_half_ungrouped)
158+{
159+ RunArch35TilingTest({{5760}, nullptr, {2880}});
160+}
161+ 
162+TEST_F(ClippedSwigluArch35TilingTest, arch35_fp32_half_ungrouped)
163+{
164+ RunArch35TilingTest({{3200, 5760}, nullptr, {1600, 5760},
165+ ge::DT_FLOAT, ge::DT_INT64, 0});
166+}
167+ 
168+TEST_F(ClippedSwigluArch35TilingTest, arch35_bf16_half_ungrouped)
169+{
170+ RunArch35TilingTest({{5760}, nullptr, {2880}, ge::DT_BF16});
171+}
172+ 
173+// ========== Normal cases: interleaved, ungrouped ==========
174+ 
175+TEST_F(ClippedSwigluArch35TilingTest, arch35_fp16_interleaved_ungrouped)
176+{
177+ RunArch35TilingTest({{5760}, nullptr, {2880}, ge::DT_FLOAT16, ge::DT_INT64, -1, 1.702f, 7.0f, 1.0f, true});
178+}
179+ 
180+TEST_F(ClippedSwigluArch35TilingTest, arch35_fp32_interleaved_ungrouped)
181+{
182+ RunArch35TilingTest({{3200, 5760}, nullptr, {1600, 5760},
183+ ge::DT_FLOAT, ge::DT_INT64, 0, 1.702f, 7.0f, 1.0f, true});
184+}
185+ 
186+TEST_F(ClippedSwigluArch35TilingTest, arch35_bf16_interleaved_ungrouped)
187+{
188+ RunArch35TilingTest({{5760}, nullptr, {2880}, ge::DT_BF16, ge::DT_INT64, -1, 1.702f, 7.0f, 1.0f, true});
189+}
190+ 
191+// ========== Normal cases: half (interleaved=false), grouped ==========
192+ 
193+TEST_F(ClippedSwigluArch35TilingTest, arch35_fp16_half_grouped)
194+{
195+ gert::StorageShape gs = {{200}};
196+ RunArch35TilingTest({{3200, 5760}, &gs, {1600, 5760}, ge::DT_FLOAT16, ge::DT_INT64, 0});
197+}
198+ 
199+TEST_F(ClippedSwigluArch35TilingTest, arch35_fp32_half_grouped)
200+{
201+ gert::StorageShape gs = {{200}};
202+ RunArch35TilingTest({{3200, 5760}, &gs, {1600, 5760}, ge::DT_FLOAT, ge::DT_INT64, 0});
203+}
204+ 
205+TEST_F(ClippedSwigluArch35TilingTest, arch35_bf16_half_grouped)
206+{
207+ gert::StorageShape gs = {{200}};
208+ RunArch35TilingTest({{3200, 5760}, &gs, {1600, 5760}, ge::DT_BF16, ge::DT_INT64, 0});
209+}
210+ 
211+// ========== Normal cases: interleaved, grouped ==========
212+ 
213+TEST_F(ClippedSwigluArch35TilingTest, arch35_fp16_interleaved_grouped)
214+{
215+ gert::StorageShape gs = {{200}};
216+ RunArch35TilingTest({{3200, 5760}, &gs, {1600, 5760},
217+ ge::DT_FLOAT16, ge::DT_INT64, 0, 1.702f, 7.0f, 1.0f, true});
218+}
219+ 
220+TEST_F(ClippedSwigluArch35TilingTest, arch35_fp32_interleaved_grouped)
221+{
222+ gert::StorageShape gs = {{200}};
223+ RunArch35TilingTest({{3200, 5760}, &gs, {1600, 5760},
224+ ge::DT_FLOAT, ge::DT_INT64, 0, 1.702f, 7.0f, 1.0f, true});
225+}
226+ 
227+TEST_F(ClippedSwigluArch35TilingTest, arch35_bf16_interleaved_grouped)
228+{
229+ gert::StorageShape gs = {{200}};
230+ RunArch35TilingTest({{3200, 5760}, &gs, {1600, 5760},
231+ ge::DT_BF16, ge::DT_INT64, 0, 1.702f, 7.0f, 1.0f, true});
232+}
233+ 
234+// ========== Normal cases: custom attrs ==========
235+ 
236+TEST_F(ClippedSwigluArch35TilingTest, arch35_fp16_custom_alpha)
237+{
238+ RunArch35TilingTest({{5760}, nullptr, {2880}, ge::DT_FLOAT16, ge::DT_INT64, -1, 1.5f});
239+}
240+ 
241+TEST_F(ClippedSwigluArch35TilingTest, arch35_fp16_custom_bias)
242+{
243+ RunArch35TilingTest({{5760}, nullptr, {2880}, ge::DT_FLOAT16, ge::DT_INT64, -1, 1.702f, 7.0f, 0.5f});
244+}
245+ 
246+TEST_F(ClippedSwigluArch35TilingTest, arch35_fp16_custom_limit)
247+{
248+ RunArch35TilingTest({{5760}, nullptr, {2880}, ge::DT_FLOAT16, ge::DT_INT64, -1, 1.702f, 10.0f});
249+}
250+ 
251+TEST_F(ClippedSwigluArch35TilingTest, arch35_fp16_positive_dim)
252+{
253+ RunArch35TilingTest({{2, 4, 5760}, nullptr, {2, 4, 2880}, ge::DT_FLOAT16, ge::DT_INT64, 2});
254+}
255+ 
256+// ========== Normal: large shape ==========
257+ 
258+TEST_F(ClippedSwigluArch35TilingTest, arch35_fp16_large_shape)
259+{
260+ RunArch35TilingTest({{64, 4096}, nullptr, {32, 4096}, ge::DT_FLOAT16, ge::DT_INT64, 0});
261+}
262+ 
263+// ========== Error cases ==========
264+ 
265+TEST_F(ClippedSwigluArch35TilingTest, arch35_x_is_none_wrong)
266+{
267+ // To test x shape is nullptr, we set xShape with 0 dims and expect failure
268+ // when xDims_ <= 0 check triggers
269+ RunArch35TilingTest({{}, nullptr, {}, ge::DT_FLOAT16, ge::DT_INT64, -1, 1.702f, 7.0f, 1.0f,
270+ false, ge::GRAPH_FAILED});
271+}
272+ 
273+TEST_F(ClippedSwigluArch35TilingTest, arch35_x_dim_div2_wrong)
274+{
275+ RunArch35TilingTest({{3200, 5761}, nullptr, {3200, 2880}, ge::DT_FLOAT16, ge::DT_INT64, 0, 1.702f,
276+ 7.0f, 1.0f, false, ge::GRAPH_FAILED});
277+}
278+ 
279+TEST_F(ClippedSwigluArch35TilingTest, arch35_x_dtype_wrong)
280+{
281+ RunArch35TilingTest({{5760}, nullptr, {2880}, ge::DT_INT8, ge::DT_INT64, -1, 1.702f, 7.0f, 1.0f,
282+ false, ge::GRAPH_FAILED});
283+}
284+ 
285+TEST_F(ClippedSwigluArch35TilingTest, arch35_limit_zero_wrong)
286+{
287+ RunArch35TilingTest({{5760}, nullptr, {2880}, ge::DT_FLOAT16, ge::DT_INT64, -1, 1.702f, 0.0f, 1.0f,
288+ false, ge::GRAPH_FAILED});
289+}
290+ 
291+TEST_F(ClippedSwigluArch35TilingTest, arch35_limit_negative_wrong)
292+{
293+ RunArch35TilingTest({{5760}, nullptr, {2880}, ge::DT_FLOAT16, ge::DT_INT64, -1, 1.702f, -1.0f, 1.0f,
294+ false, ge::GRAPH_FAILED});
295+}
296+ 
297+TEST_F(ClippedSwigluArch35TilingTest, arch35_groupindex_dtype_wrong)
298+{
299+ gert::StorageShape gs = {{200}};
300+ RunArch35TilingTest({{3200, 5760}, &gs, {1600, 5760}, ge::DT_FLOAT16, ge::DT_FLOAT, 0, 1.702f,
301+ 7.0f, 1.0f, false, ge::GRAPH_FAILED});
302+}
303+ 
304+TEST_F(ClippedSwigluArch35TilingTest, arch35_groupindex_dims_wrong)
305+{
306+ gert::StorageShape gs = {{200, 2}};
307+ RunArch35TilingTest({{3200, 5760}, &gs, {1600, 5760}, ge::DT_FLOAT16, ge::DT_INT64, 0, 1.702f,
308+ 7.0f, 1.0f, false, ge::GRAPH_FAILED});
309+}
310+ 
311+TEST_F(ClippedSwigluArch35TilingTest, arch35_y_dims_diff_wrong)
312+{
313+ RunArch35TilingTest({{3200, 5760}, nullptr, {3200, 2880, 2}, ge::DT_FLOAT16, ge::DT_INT64, 0,
314+ 1.702f, 7.0f, 1.0f, false, ge::GRAPH_FAILED});
315+}
316+ 
317+TEST_F(ClippedSwigluArch35TilingTest, arch35_y_dim_equal_x_dim_div2_wrong)
318+{
319+ RunArch35TilingTest({{3200, 5760}, nullptr, {3200, 2800}, ge::DT_FLOAT16, ge::DT_INT64, 0, 1.702f,
320+ 7.0f, 1.0f, false, ge::GRAPH_FAILED});
321+}
322+ 
323+TEST_F(ClippedSwigluArch35TilingTest, arch35_dim_value_wrong)
324+{
325+ RunArch35TilingTest({{3200, 5760}, nullptr, {1600, 5760}, ge::DT_FLOAT16, ge::DT_INT64, 10,
326+ 1.702f, 7.0f, 1.0f, false, ge::GRAPH_FAILED});
327+}
Mscripts/kernel/binary_config/ascendc_config.json+1-1
@@ -347,7 +347,7 @@
347 {"name": "GroupNormV2", "compute_units": ["ascend950"], "auto_sync": false, "impl_mode": "", "compile_options": {"ascend950": ["--cce-no-dcache-flush"]}},347 {"name": "GroupNormV2", "compute_units": ["ascend950"], "auto_sync": false, "impl_mode": "", "compile_options": {"ascend950": ["--cce-no-dcache-flush"]}},
348 {"name":"AddRmsNormDynamicMxQuant", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode": "", "compile_options": {"ascend950": ["--cce-no-dcache-flush", "--cce-simd-vf-fusion=false"]}},348 {"name":"AddRmsNormDynamicMxQuant", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode": "", "compile_options": {"ascend950": ["--cce-no-dcache-flush", "--cce-simd-vf-fusion=false"]}},
349 {"name":"SwiGlu", "compute_units": ["ascend910_93", "ascend910b", "ascend950", "kirinx90", "kirin9030", "ascend310p"], "auto_sync": true, "compile_options": {"ascend950": ["--cce-no-dcache-flush"]}},349 {"name":"SwiGlu", "compute_units": ["ascend910_93", "ascend910b", "ascend950", "kirinx90", "kirin9030", "ascend310p"], "auto_sync": true, "compile_options": {"ascend950": ["--cce-no-dcache-flush"]}},
350- {"name":"ClippedSwiglu", "compute_units": ["ascend910_93", "ascend910b"], "auto_sync": true},350+ {"name":"ClippedSwiglu", "compute_units": ["ascend910_93", "ascend910b", "ascend950"], "auto_sync": {"ascend910_93": true, "ascend910b": true, "ascend950": false}, "compile_options": {"ascend950": ["--cce-no-dcache-flush"]}},
351 {"name":"Axpy", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},351 {"name":"Axpy", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},
352 {"name":"AxpyV2", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},352 {"name":"AxpyV2", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},
353 {"name":"SwiGluGrad", "compute_units": ["ascend910_93", "ascend910b", "ascend950", "ascend310p"], "auto_sync": true, "compile_options": {"ascend950": ["--cce-no-dcache-flush"]}},353 {"name":"SwiGluGrad", "compute_units": ["ascend910_93", "ascend910b", "ascend950", "ascend310p"], "auto_sync": true, "compile_options": {"ascend950": ["--cce-no-dcache-flush"]}},