QuantBatchMatmul伪量化MXA8W4算子迁移blaze框架-ops-nn-AtomGit

Mcmake/third_party/ops-tensor.cmake+1-1

	@@ -7,7 +7,7 @@
7	# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.	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.	8	# See LICENSE in the root of the software repository for the full text of the License.
9	# ----------------------------------------------------------------------------	9	# ----------------------------------------------------------------------------
10	-set(OPTENSOR_TAG_ID 8badbc6e5b4b1f9f91d177568abe4e27d4acabfa)	10	+set(OPTENSOR_TAG_ID 7e9400095a9c6fc9806daa23496b04eb2d52c726)
11		11
12	if(EXISTS "${PROJECT_SOURCE_DIR}/../ops-tensor")	12	if(EXISTS "${PROJECT_SOURCE_DIR}/../ops-tensor")
13	get_filename_component(OPTENSOR_SOURCE_PATH	13	get_filename_component(OPTENSOR_SOURCE_PATH

Mmatmul/quant_batch_matmul_v4/op_host/op_tiling/arch35/quant_batch_matmul_v4_reg_base_tiling.cpp+83-194

	@@ -18,132 +18,18 @@
18	#include "graph/utils/type_utils.h"	18	#include "graph/utils/type_utils.h"
19	#include "log/log.h"	19	#include "log/log.h"
20	#include "register/op_impl_registry.h"	20	#include "register/op_impl_registry.h"
21	-#include "quant_batch_matmul_v4_tiling.h"	21	+#include "quant_batch_matmul_v4_reg_base_tiling.h"
22	#include "matmul/common/op_host/math_util.h"	22	#include "matmul/common/op_host/math_util.h"
23	#include "error_util.h"	23	#include "error_util.h"
24	-#include "../../../op_kernel/arch35/quant_batch_matmul_v4_tiling_key.h"
25		24
26	using AscendC::BLOCK_CUBE; // uint32_t	25	using AscendC::BLOCK_CUBE; // uint32_t
27	using AscendC::ONE_BLK_SIZE; // uint32_t	26	using AscendC::ONE_BLK_SIZE; // uint32_t
28		27
29	-namespace {
30	-// aiv和aic核数比例
31	-constexpr uint32_t CORE_RATIO = 2U;
32	-} // namespace
33	-
34	namespace optiling {	28	namespace optiling {
35	using namespace matmul_v4;	29	using namespace matmul_v4;
36		30
37	bool QuantBatchMatmulV4RegBase::IsCapable() { return true; }	31	bool QuantBatchMatmulV4RegBase::IsCapable() { return true; }
38		32
39	-bool QuantBatchMatmulV4RegBase::CheckA8W4Params() const
40	-{
41	- OP_CHECK_IF(inputParams_.transA,
42	- OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(inputParams_.opName, "transposeX1",
43	- (inputParams_.transA ? "true" : "false"),
44	- "The value of transposeX1 must be false"),
45	- return false);
46	- OP_CHECK_IF(inputParams_.bFormat == ge::FORMAT_ND && !inputParams_.transB,
47	- OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(inputParams_.opName, "transposeX2", "false",
48	- "When the format of x2 is ND, transposeX2 must be true"),
49	- return false);
50	- OP_CHECK_IF(
51	- inputParams_.antiQuantType == QuantType::PER_GROUP && inputParams_.bFormat == ge::FORMAT_FRACTAL_NZ &&
52	- inputParams_.transB,
53	- OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(
54	- inputParams_.opName, "transposeX2", "true",
55	- "When the quantization mode is pergroup and the format of x2 is FRACTAL_NZ, transposeX2 must be false"),
56	- return false);
57	-
58	- if (inputParams_.antiQuantType == QuantType::MX) {
59	- OP_CHECK_IF(inputParams_.groupSize != MX_GROUP_SIZE,
60	- OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(inputParams_.opName, "groupSize",
61	- std::to_string(inputParams_.groupSize).c_str(),
62	- "groupSize must be 32 when the quantization mode is MX"),
63	- return false);
64	- } else {
65	- OP_CHECK_IF(
66	- inputParams_.groupSize <= 0 \|\| inputParams_.kSize < inputParams_.groupSize,
67	- OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(
68	- inputParams_.opName, "groupSize", std::to_string(inputParams_.groupSize).c_str(),
69	- "groupSize must be greater than 0 and less than kSize(" + std::to_string(inputParams_.kSize) + ")"),
70	- return false);
71	- }
72	-
73	- OP_CHECK_IF(inputParams_.groupSize % GROUP_ALIGN_SIZE > 0,
74	- OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(inputParams_.opName, "groupSize",
75	- std::to_string(inputParams_.groupSize).c_str(),
76	- "groupSize must be aligned to 32"),
77	- return false);
78	- // A8W4 Nz场景要求n为32B对齐
79	- OP_CHECK_IF(
80	- inputParams_.bFormat == ge::FORMAT_FRACTAL_NZ && inputParams_.nSize % N_ALIGN_SIZE > 0,
81	- OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(inputParams_.opName, "nSize", std::to_string(inputParams_.nSize).c_str(),
82	- "nSize must be aligned to 8 when the format of x2 is FRACTAL_NZ"),
83	- return false);
84	- return true;
85	-}
86	-
87	-bool QuantBatchMatmulV4RegBase::CustomCheck() const
88	-{
89	- if (inputParams_.antiQuantType == QuantType::MX) {
90	- OP_CHECK_IF(inputParams_.kSize % K_ALIGN_SIZE_MX > 0,
91	- OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(
92	- inputParams_.opName, "kSize", std::to_string(inputParams_.kSize).c_str(),
93	- "kSize must be aligned to 8 when the quantization mode is MX"),
94	- return false);
95	- } else {
96	- OP_CHECK_IF(inputParams_.kSize % K_ALIGN_SIZE > 0 \|\| inputParams_.kSize <= K_ALIGN_SIZE,
97	- OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(inputParams_.opName, "kSize",
98	- std::to_string(inputParams_.kSize).c_str(),
99	- "kSize must be aligned to 32 and greater than 32"),
100	- return false);
101	- }
102	-
103	- OP_CHECK_IF((inputParams_.cDtype != ge::DT_BF16) && (inputParams_.cDtype != ge::DT_FLOAT16),
104	- OP_LOGE_FOR_INVALID_DTYPE_WITH_REASON(
105	- inputParams_.opName, "y", ge::TypeUtils::DataTypeToSerialString(inputParams_.cDtype).c_str(),
106	- "The dtype of y must be BF16 or FLOAT16"),
107	- return false);
108	-
109	- bool a8w4Flag = (inputParams_.aDtype == ge::DT_HIFLOAT8 \|\| inputParams_.aDtype == ge::DT_FLOAT8_E5M2 \|\|
110	- inputParams_.aDtype == ge::DT_FLOAT8_E4M3FN) &&
111	- (inputParams_.bDtype == ge::DT_FLOAT4_E2M1 \|\| inputParams_.bDtype == ge::DT_FLOAT);
112	- if (a8w4Flag) {
113	- return CheckA8W4Params();
114	- } else {
115	- std::string incorrectVals = std::string("x1:") + ge::TypeUtils::DataTypeToSerialString(inputParams_.aDtype) +
116	- ", x2:" + ge::TypeUtils::DataTypeToSerialString(inputParams_.bDtype) +
117	- ", y:" + ge::TypeUtils::DataTypeToSerialString(inputParams_.cDtype) +
118	- ", groupSize:" + std::to_string(inputParams_.groupSize) +
119	- ", transA:" + (inputParams_.transA ? "true" : "false") +
120	- ", transB:" + (inputParams_.transB ? "true" : "false");
121	- OP_LOGE_FOR_INVALID_VALUES_WITH_REASON(
122	- inputParams_.opName, "x1, x2, y, groupSize, transposeX1, transposeX2", incorrectVals.c_str(),
123	- "The dtype of x1 must be HIFLOAT8, FLOAT8_E5M2, or FLOAT8_E4M3FN, the dtype of x2 must be FLOAT4_E2M1 or "
124	- "FLOAT, and the dtype of y must be BF16 or FLOAT16");
125	- return false;
126	- }
127	- return true;
128	-}
129	-
130	-bool QuantBatchMatmulV4RegBase::CheckCoreNum() const
131	-{
132	- if (aivNum_ == 0 \|\| aicNum_ == 0) {
133	- OP_LOGE_FOR_INVALID_VALUES_WITH_REASON(inputParams_.opName, "aicNum, aivNum",
134	- std::to_string(aicNum_) + ", " + std::to_string(aivNum_),
135	- "aicNum and aivNum must be greater than 0");
136	- return false;
137	- }
138	- if (aivNum_ != CORE_RATIO * aicNum_) {
139	- OP_LOGE_FOR_INVALID_VALUES_WITH_REASON(inputParams_.opName, "aicNum, aivNum",
140	- std::to_string(aicNum_) + ", " + std::to_string(aivNum_),
141	- "aicNum:aivNum must be 1:2");
142	- return false;
143	- }
144	- return true;
145	-}
146	-
147	ge::graphStatus QuantBatchMatmulV4RegBase::DoOpTiling()	33	ge::graphStatus QuantBatchMatmulV4RegBase::DoOpTiling()
148	{	34	{
149	OP_TILING_CHECK(InstantiateTilingData() == ge::GRAPH_FAILED,	35	OP_TILING_CHECK(InstantiateTilingData() == ge::GRAPH_FAILED,
	@@ -157,11 +43,11 @@ ge::graphStatus QuantBatchMatmulV4RegBase::DoOpTiling()
157		43
158	uint64_t weightBlockAlignSize = GetBlockAlignSizeByDataType(inputParams_.bDtype);	44	uint64_t weightBlockAlignSize = GetBlockAlignSizeByDataType(inputParams_.bDtype);
159	// transB的场景	45	// transB的场景
160	- tilingData_->kAlign = ops::CeilAlign(inputParams_.kSize, weightBlockAlignSize);	46	+ tilingData_.kAlign = ops::CeilAlign(inputParams_.kSize, weightBlockAlignSize);
161	- tilingData_->nAlign = inputParams_.nSize;	47	+ tilingData_.nAlign = inputParams_.nSize;
162	- tilingData_->kSize = inputParams_.kSize;	48	+ tilingData_.kSize = inputParams_.kSize;
163	- tilingData_->nSize = inputParams_.nSize;	49	+ tilingData_.nSize = inputParams_.nSize;
164	- tilingData_->mSize = inputParams_.mSize;	50	+ tilingData_.mSize = inputParams_.mSize;
165		51
166	PlatformParam platformParam = {aicNum_,	52	PlatformParam platformParam = {aicNum_,
167	aicNum_,	53	aicNum_,
	@@ -192,42 +78,19 @@ ge::graphStatus QuantBatchMatmulV4RegBase::DoOpTiling()
192	return ge::GRAPH_SUCCESS;	78	return ge::GRAPH_SUCCESS;
193	}	79	}
194		80
195	-uint64_t QuantBatchMatmulV4RegBase::GetTilingKey() const
196	-{
197	- uint64_t trans = (static_cast<uint64_t>(inputParams_.transA) << 1) \| static_cast<uint64_t>(inputParams_.transB);
198	- return GET_TPL_TILING_KEY(trans, static_cast<uint64_t>(inputParams_.antiQuantType),
199	- static_cast<uint64_t>(inputParams_.hasAntiQuantOffset),
200	- static_cast<uint64_t>(inputParams_.weightNz),
201	- static_cast<uint64_t>(KernelTemplateType::BASIS));
202	-}
203	-
204	ge::graphStatus QuantBatchMatmulV4RegBase::GetWorkspaceSize()	81	ge::graphStatus QuantBatchMatmulV4RegBase::GetWorkspaceSize()
205	{	82	{
206	workspaceSize_ = WORKSPACE_SIZE;	83	workspaceSize_ = WORKSPACE_SIZE;
207	- workspaceSize_ += static_cast<uint64_t>(tilingData_->cubeNumBlocksN) * tilingData_->cubeNumBlocksM *	84	+ workspaceSize_ += tilingData_.cubeNumBlocksN * tilingData_.cubeNumBlocksM * sizeof(uintptr_t);
208	- sizeof(uintptr_t);
209	return ge::GRAPH_SUCCESS;	85	return ge::GRAPH_SUCCESS;
210	}	86	}
211		87
212	ge::graphStatus QuantBatchMatmulV4RegBase::PostTiling()	88	ge::graphStatus QuantBatchMatmulV4RegBase::PostTiling()
213	{	89	{
214	- OP_LOGD(inputParams_.opName, "final tiling data size: %zu", tilingDataSize_);	90	+ uint32_t usedCoreNum = tilingData_.cubeNumBlocksM * tilingData_.cubeNumBlocksN;
215	-	91	+ auto status = SerializeTilingData(&tilingData_, sizeof(tilingData_), usedCoreNum);
216	- OP_TILING_CHECK(tilingDataSize_ % sizeof(uint64_t) != 0,	92	+ if (status != ge::GRAPH_SUCCESS) {
217	- OP_LOGE(inputParams_.opName, "tiling data size[%zu] not aligned to 8", tilingDataSize_),	93	+ return status;
218	- return ge::GRAPH_FAILED);
219	- context_->GetRawTilingData()->SetDataSize(tilingDataSize_);
220	- context_->SetBlockDim(tilingData_->cubeNumBlocksM * tilingData_->cubeNumBlocksN);
221	-
222	- size_t* workspaces = context_->GetWorkspaceSizes(1); // set workspace
223	- OPS_CHECK_NULL_WITH_CONTEXT(context_, workspaces);
224	- workspaces[0] = workspaceSize_;
225	-
226	- errno_t ret = memcpy_s(context_->GetRawTilingData()->GetData(), context_->GetRawTilingData()->GetCapacity(),
227	- reinterpret_cast<void*>(tilingData_), tilingDataSize_);
228	- if (ret != EOK) {
229	- OP_LOGE(context_->GetNodeName(), "memcpy_s failed, ret=%d", ret);
230	- return ge::GRAPH_FAILED;
231	}	94	}
232	PrintCVTilingData(true);	95	PrintCVTilingData(true);
233	return ge::GRAPH_SUCCESS;	96	return ge::GRAPH_SUCCESS;
	@@ -238,57 +101,57 @@ void QuantBatchMatmulV4RegBase::SetBubTiling()
238	int64_t nBubSize;	101	int64_t nBubSize;
239	int64_t kBubSize;	102	int64_t kBubSize;
240	GetBubTilingA8W4(nBubSize, kBubSize);	103	GetBubTilingA8W4(nBubSize, kBubSize);
241	- tilingData_->nBubSize = nBubSize;	104	+ tilingData_.nBubSize = nBubSize;
242	- tilingData_->kBubSize = kBubSize;	105	+ tilingData_.kBubSize = kBubSize;
243	}	106	}
244		107
245	void QuantBatchMatmulV4RegBase::SetMatmulTiling()	108	void QuantBatchMatmulV4RegBase::SetMatmulTiling()
246	{	109	{
247	const BasicBlockParam& tilingRes = tilingSolver_.GetTilingResult();	110	const BasicBlockParam& tilingRes = tilingSolver_.GetTilingResult();
248	- tilingData_->cubeNumBlocksM = static_cast<uint8_t>(tilingRes.mDim);	111	+ tilingData_.cubeNumBlocksM = static_cast<uint8_t>(tilingRes.mDim);
249	- tilingData_->cubeNumBlocksN = static_cast<uint8_t>(tilingRes.nDim);	112	+ tilingData_.cubeNumBlocksN = static_cast<uint8_t>(tilingRes.nDim);
250		113
251	- tilingData_->matmulTiling.M = tilingRes.mSize;	114	+ tilingData_.matmulTiling.M = tilingRes.mSize;
252	- tilingData_->matmulTiling.Ka = tilingRes.kSize;	115	+ tilingData_.matmulTiling.Ka = tilingRes.kSize;
253	- tilingData_->matmulTiling.N = tilingRes.nSize;	116	+ tilingData_.matmulTiling.N = tilingRes.nSize;
254	- tilingData_->matmulTiling.Kb = tilingRes.kSize;	117	+ tilingData_.matmulTiling.Kb = tilingRes.kSize;
255	- tilingData_->matmulTiling.singleCoreM = tilingRes.l1Param.stepM * tilingRes.basicBlock.baseM;	118	+ tilingData_.matmulTiling.singleCoreM = tilingRes.l1Param.stepM * tilingRes.basicBlock.baseM;
256	- tilingData_->matmulTiling.singleCoreK = tilingRes.l1Param.stepKa * tilingRes.basicBlock.baseK;	119	+ tilingData_.matmulTiling.singleCoreK = tilingRes.l1Param.stepKa * tilingRes.basicBlock.baseK;
257	- tilingData_->matmulTiling.singleCoreN = tilingRes.l1Param.stepN * tilingRes.basicBlock.baseN;	120	+ tilingData_.matmulTiling.singleCoreN = tilingRes.l1Param.stepN * tilingRes.basicBlock.baseN;
258		121
259	- tilingData_->matmulTiling.baseM = tilingRes.basicBlock.baseM;	122	+ tilingData_.matmulTiling.baseM = tilingRes.basicBlock.baseM;
260	- tilingData_->matmulTiling.baseN = tilingRes.basicBlock.baseN;	123	+ tilingData_.matmulTiling.baseN = tilingRes.basicBlock.baseN;
261	- tilingData_->matmulTiling.baseK = tilingRes.basicBlock.baseK;	124	+ tilingData_.matmulTiling.baseK = tilingRes.basicBlock.baseK;
262	- tilingData_->matmulTiling.dbL0A = DB_BUFFER;	125	+ tilingData_.matmulTiling.dbL0A = DB_BUFFER;
263	- tilingData_->matmulTiling.dbL0B = DB_BUFFER;	126	+ tilingData_.matmulTiling.dbL0B = DB_BUFFER;
264	- tilingData_->matmulTiling.dbL0C = 1;	127	+ tilingData_.matmulTiling.dbL0C = 1;
265		128
266	- tilingData_->matmulTiling.stepM = tilingRes.l1Param.stepM;	129	+ tilingData_.matmulTiling.stepM = tilingRes.l1Param.stepM;
267	- tilingData_->matmulTiling.stepN = tilingRes.l1Param.stepN;	130	+ tilingData_.matmulTiling.stepN = tilingRes.l1Param.stepN;
268	- tilingData_->matmulTiling.stepKa = tilingRes.l1Param.stepKa;	131	+ tilingData_.matmulTiling.stepKa = tilingRes.l1Param.stepKa;
269	- tilingData_->matmulTiling.stepKb = tilingRes.l1Param.stepKb;	132	+ tilingData_.matmulTiling.stepKb = tilingRes.l1Param.stepKb;
270	- tilingData_->matmulTiling.depthA1 = tilingRes.l1Param.A1BufferNum * tilingRes.l1Param.stepM *	133	+ tilingData_.matmulTiling.depthA1 = tilingRes.l1Param.A1BufferNum * tilingRes.l1Param.stepM *
271	- tilingRes.l1Param.stepKa;	134	+ tilingRes.l1Param.stepKa;
272	- tilingData_->matmulTiling.depthB1 = tilingRes.l1Param.B1BufferNum * tilingRes.l1Param.stepN *	135	+ tilingData_.matmulTiling.depthB1 = tilingRes.l1Param.B1BufferNum * tilingRes.l1Param.stepN *
273	- tilingRes.l1Param.stepKb;	136	+ tilingRes.l1Param.stepKb;
274	- tilingData_->matmulTiling.iterateOrder = tilingRes.l1Param.iterateOrder;	137	+ tilingData_.matmulTiling.iterateOrder = tilingRes.l1Param.iterateOrder;
275		138
276	- tilingData_->matmulTiling.isBias = static_cast<int32_t>(inputParams_.hasBias);	139	+ tilingData_.matmulTiling.isBias = static_cast<int32_t>(inputParams_.hasBias);
277	- tilingData_->hasX1Scale = static_cast<int32_t>(inputParams_.hasX1Scale);	140	+ tilingData_.hasX1Scale = static_cast<int32_t>(inputParams_.hasX1Scale);
278	- tilingData_->hasX2Scale = static_cast<int32_t>(inputParams_.hasX2Scale);	141	+ tilingData_.hasX2Scale = static_cast<int32_t>(inputParams_.hasX2Scale);
279	- tilingData_->matmulTiling.shareL1Size = 0;	142	+ tilingData_.matmulTiling.shareL1Size = 0;
280	- tilingData_->matmulTiling.shareL0CSize = 0;	143	+ tilingData_.matmulTiling.shareL0CSize = 0;
281		144
282	uint32_t scaleFactorA = static_cast<uint32_t>(tilingRes.l1Param.scaleFactor);	145	uint32_t scaleFactorA = static_cast<uint32_t>(tilingRes.l1Param.scaleFactor);
283	uint32_t scaleFactorB = static_cast<uint32_t>(tilingRes.l1Param.scaleFactor);	146	uint32_t scaleFactorB = static_cast<uint32_t>(tilingRes.l1Param.scaleFactor);
284	- tilingData_->matmulTiling.mxTypePara = (scaleFactorB << B8_BITS) + scaleFactorA;	147	+ tilingData_.matmulTiling.mxTypePara = (scaleFactorB << B8_BITS) + scaleFactorA;
285	- tilingData_->AL1Pingpong = tilingRes.l1Param.A1BufferNum;	148	+ tilingData_.AL1Pingpong = tilingRes.l1Param.A1BufferNum;
286	- tilingData_->BL1Pingpong = tilingRes.l1Param.B1BufferNum;	149	+ tilingData_.BL1Pingpong = tilingRes.l1Param.B1BufferNum;
287	- tilingData_->mAL1Size = tilingRes.l1Param.stepM * tilingRes.basicBlock.baseM;	150	+ tilingData_.mAL1Size = tilingRes.l1Param.stepM * tilingRes.basicBlock.baseM;
288	- tilingData_->kAL1Size = std::min(tilingRes.l1Param.stepKa * tilingRes.basicBlock.baseK, tilingRes.singleK);	151	+ tilingData_.kAL1Size = std::min(tilingRes.l1Param.stepKa * tilingRes.basicBlock.baseK, tilingRes.singleK);
289	- tilingData_->nBL1Size = tilingRes.l1Param.stepN * tilingRes.basicBlock.baseN;	152	+ tilingData_.nBL1Size = tilingRes.l1Param.stepN * tilingRes.basicBlock.baseN;
290	- tilingData_->kBL1Size = std::min(tilingRes.l1Param.stepKb * tilingRes.basicBlock.baseK, tilingRes.singleK);	153	+ tilingData_.kBL1Size = std::min(tilingRes.l1Param.stepKb * tilingRes.basicBlock.baseK, tilingRes.singleK);
291	- tilingData_->groupSize = inputParams_.groupSize;	154	+ tilingData_.groupSize = inputParams_.groupSize;
292	}	155	}
293		156
294	uint64_t QuantBatchMatmulV4RegBase::GetGroupNumBub(uint64_t kDimSzie) const	157	uint64_t QuantBatchMatmulV4RegBase::GetGroupNumBub(uint64_t kDimSzie) const
	@@ -420,13 +283,13 @@ void QuantBatchMatmulV4RegBase::PrintCVTilingData(const bool debugLevel) const
420	int64_t QuantBatchMatmulV4RegBase::DumpCVTilingDataToLog(const bool debugLevel) const	283	int64_t QuantBatchMatmulV4RegBase::DumpCVTilingDataToLog(const bool debugLevel) const
421	{	284	{
422	std::stringstream ss;	285	std::stringstream ss;
423	- ss << "kAlign: " << tilingData_->kAlign << " kSize: " << tilingData_->kSize << " nSize: " << tilingData_->nSize	286	+ ss << "kAlign: " << tilingData_.kAlign << " kSize: " << tilingData_.kSize << " nSize: " << tilingData_.nSize
424	- << " mSize: " << tilingData_->mSize << " cubeNumBlocksN: " << static_cast<uint32_t>(tilingData_->cubeNumBlocksN)	287	+ << " mSize: " << tilingData_.mSize << " cubeNumBlocksN: " << static_cast<uint32_t>(tilingData_.cubeNumBlocksN)
425	- << " cubeNumBlocksM: " << static_cast<uint32_t>(tilingData_->cubeNumBlocksM)	288	+ << " cubeNumBlocksM: " << static_cast<uint32_t>(tilingData_.cubeNumBlocksM)
426	- << " nBubSize: " << tilingData_->nBubSize << " kBubSize: " << tilingData_->kBubSize	289	+ << " nBubSize: " << tilingData_.nBubSize << " kBubSize: " << tilingData_.kBubSize
427	- << " mAL1Size: " << tilingData_->mAL1Size << " kAL1Size: " << tilingData_->kAL1Size	290	+ << " mAL1Size: " << tilingData_.mAL1Size << " kAL1Size: " << tilingData_.kAL1Size
428	- << " nBL1Size: " << tilingData_->nBL1Size << " kBL1Size: " << tilingData_->kBL1Size	291	+ << " nBL1Size: " << tilingData_.nBL1Size << " kBL1Size: " << tilingData_.kBL1Size
429	- << " AL1Pingpong: " << tilingData_->AL1Pingpong << " BL1Pingpong: " << tilingData_->BL1Pingpong;	292	+ << " AL1Pingpong: " << tilingData_.AL1Pingpong << " BL1Pingpong: " << tilingData_.BL1Pingpong;
430	if (debugLevel) {	293	if (debugLevel) {
431	OPS_LOG_D(inputParams_.opName, "tiling data: %s", ss.str().c_str());	294	OPS_LOG_D(inputParams_.opName, "tiling data: %s", ss.str().c_str());
432	} else {	295	} else {
	@@ -435,4 +298,30 @@ int64_t QuantBatchMatmulV4RegBase::DumpCVTilingDataToLog(const bool debugLevel)
435	PrintMatMulTiling();	298	PrintMatMulTiling();
436	return 0;	299	return 0;
437	}	300	}
		301	+
		302	+ge::graphStatus QuantBatchMatmulV4RegBase::InstantiateTilingData()
		303	+{
		304	+ return CheckTilingDataCapacity(&tilingData_, sizeof(tilingData_));
		305	+}
		306	+
		307	+void QuantBatchMatmulV4RegBase::PrintTilingData(bool debugLevel) { PrintCVTilingData(debugLevel); }
		308	+
		309	+void QuantBatchMatmulV4RegBase::PrintMatMulTiling() const
		310	+{
		311	+ std::stringstream ss;
		312	+ auto& matmulTiling = tilingData_.matmulTiling;
		313	+ ss << "usedCoreNum " << matmulTiling.usedCoreNum << " M " << matmulTiling.M << " N " << matmulTiling.N << " Ka "
		314	+ << matmulTiling.Ka << " Kb " << matmulTiling.Kb << " singleCoreM " << matmulTiling.singleCoreM << " singleCoreN "
		315	+ << matmulTiling.singleCoreN << " singleCoreK " << matmulTiling.singleCoreK << " baseM " << matmulTiling.baseM
		316	+ << " baseN " << matmulTiling.baseN << " baseK " << matmulTiling.baseK << " depthA1 " << matmulTiling.depthA1
		317	+ << " depthB1 " << matmulTiling.depthB1 << " stepM " << matmulTiling.stepM << " stepN " << matmulTiling.stepN
		318	+ << " isBias " << matmulTiling.isBias << " transLength " << matmulTiling.transLength << " iterateOrder "
		319	+ << matmulTiling.iterateOrder << " shareMode " << matmulTiling.shareMode << " shareL1Size "
		320	+ << matmulTiling.shareL1Size << " shareL0CSize " << matmulTiling.shareL0CSize << " shareUbSize "
		321	+ << matmulTiling.shareUbSize << " batchM " << matmulTiling.batchM << " batchN " << matmulTiling.batchN
		322	+ << " stepKa " << matmulTiling.stepKa << " stepKb " << matmulTiling.stepKb << " dbL0A " << matmulTiling.dbL0A
		323	+ << " dbL0B " << matmulTiling.dbL0B << " dbL0C " << matmulTiling.dbL0C;
		324	+
		325	+ OPS_LOG_I(inputParams_.opName, "matmul tiling: %s", ss.str().c_str());
		326	+}
438	} // namespace optiling	327	} // namespace optiling

Amatmul/quant_batch_matmul_v4/op_host/op_tiling/arch35/quant_batch_matmul_v4_reg_base_tiling.h+68-0

@@ -0,0 +1,68 @@
	1	+/**
	2	+ * Copyright (c) 2025-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 quant_batch_matmul_v4_reg_base_tiling.h
	13	+ * \brief
	14	+ */
	15	+
	16	+#pragma once
	17	+
	18	+#include "quant_batch_matmul_v4_basic_block_tiling.h"
	19	+#include "quant_batch_matmul_v4_tiling.h"
	20	+
	21	+namespace optiling {
	22	+class QuantBatchMatmulV4RegBase : public QuantBatchMatmulV4TilingBase {
	23	+public:
	24	+ explicit QuantBatchMatmulV4RegBase(gert::TilingContext* context) : QuantBatchMatmulV4TilingBase(context)
	25	+ {
	26	+ tilingSolver_.Init();
	27	+ }
	28	+ ~QuantBatchMatmulV4RegBase() override = default;
	29	+
	30	+protected:
	31	+ bool IsCapable() override;
	32	+ ge::graphStatus DoOpTiling() override;
	33	+ ge::graphStatus DoLibApiTiling() override { return ge::GRAPH_SUCCESS; }
	34	+ ge::graphStatus GetWorkspaceSize() override;
	35	+ ge::graphStatus PostTiling() override;
	36	+
	37	+ bool CalcUBSize(uint64_t vecSingleN, uint64_t vecSingleK) const override
	38	+ {
	39	+ (void)vecSingleN;
	40	+ (void)vecSingleK;
	41	+ return true;
	42	+ }
	43	+
	44	+ bool SetQuantType(const gert::StorageShape* quantScaleShape, const gert::StorageShape* quantOffsetShape) override
	45	+ {
	46	+ (void)quantScaleShape;
	47	+ (void)quantOffsetShape;
	48	+ return true;
	49	+ }
	50	+
	51	+ QuantBatchMatmulV4BasicBlockTiling tilingSolver_;
	52	+
	53	+private:
	54	+ ge::graphStatus InstantiateTilingData();
	55	+ void PrintTilingData(bool debugLevel);
	56	+ void SetBubTiling();
	57	+ void GetBubTilingA8W4(int64_t& nBubSize, int64_t& kBubSize) const;
	58	+ void GetBubTilingA8W4BySize(int64_t& nBubSize, int64_t& kBubSize, int64_t& kBl1Size, int64_t& nBl1Size) const;
	59	+ void SetMatmulTiling();
	60	+ uint64_t GetGroupNumBub(uint64_t kDimSzie) const;
	61	+ uint64_t GetBubSize(uint64_t bubN, uint64_t bubD, bool isWeightNz) const;
	62	+ void PrintCVTilingData(const bool debugLevel) const;
	63	+ int64_t DumpCVTilingDataToLog(const bool debugLevel) const;
	64	+ void PrintMatMulTiling() const;
	65	+
	66	+ qbmmv4_tiling::QuantBatchMatmulV4TilingDataParams tilingData_;
	67	+};
	68	+} // namespace optiling

Mmatmul/quant_batch_matmul_v4/op_host/op_tiling/arch35/quant_batch_matmul_v4_tiling.cpp+163-108

	@@ -19,11 +19,13 @@
19	#include <mutex>	19	#include <mutex>
20	#include <numeric>	20	#include <numeric>
21	#include <set>	21	#include <set>
		22	+#include <sstream>
22		23
23	#include "util/math_util.h"	24	#include "util/math_util.h"
24	#include "graph/utils/type_utils.h"	25	#include "graph/utils/type_utils.h"
25	#include "log/log.h"	26	#include "log/log.h"
26	#include "error_util.h"	27	#include "error_util.h"
		28	+#include "securec.h"
27	#include "register/op_impl_registry.h"	29	#include "register/op_impl_registry.h"
28	#include "log/log.h"	30	#include "log/log.h"
29	#include "error_util.h"	31	#include "error_util.h"
	@@ -31,6 +33,7 @@
31	#include "platform/platform_infos_def.h"	33	#include "platform/platform_infos_def.h"
32	#include "matmul/common/op_host/op_tiling/debug_tiling.h"	34	#include "matmul/common/op_host/op_tiling/debug_tiling.h"
33	#include "../../../op_kernel/arch35/quant_batch_matmul_v4_tiling_data_apt.h"	35	#include "../../../op_kernel/arch35/quant_batch_matmul_v4_tiling_data_apt.h"
		36	+#include "../../../op_kernel/arch35/quant_batch_matmul_v4_tiling_key.h"
34		37
35	using AscendC::BLOCK_CUBE;	38	using AscendC::BLOCK_CUBE;
36	using namespace Ops::NN;	39	using namespace Ops::NN;
	@@ -38,6 +41,7 @@ using namespace Ops::NN;
38	namespace optiling {	41	namespace optiling {
39	constexpr uint64_t B4_IN_B32_NUMS = 8UL;	42	constexpr uint64_t B4_IN_B32_NUMS = 8UL;
40	constexpr uint64_t GROUP_MKN_BIT_SIZE = 0xFFFF;	43	constexpr uint64_t GROUP_MKN_BIT_SIZE = 0xFFFF;
		44	+constexpr uint32_t CORE_RATIO = 2U;
41	using namespace matmul_v4;	45	using namespace matmul_v4;
42		46
43	inline bool IsNotEmptyShape(const gert::StorageShape* storageShape)	47	inline bool IsNotEmptyShape(const gert::StorageShape* storageShape)
	@@ -90,6 +94,8 @@ void QuantBatchMatmulV4TilingBase::InitCompileInfo()
90	compileInfoPtr_->aicNum = ascendcPlatform.GetCoreNumAic();	94	compileInfoPtr_->aicNum = ascendcPlatform.GetCoreNumAic();
91	ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, compileInfoPtr_->ubSize);	95	ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, compileInfoPtr_->ubSize);
92	ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::L1, compileInfoPtr_->l1Size);	96	ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::L1, compileInfoPtr_->l1Size);
		97	+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::L0_A, compileInfoPtr_->l0aSize);
		98	+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::L0_B, compileInfoPtr_->l0bSize);
93	ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::L0_C, compileInfoPtr_->l0cSize);	99	ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::L0_C, compileInfoPtr_->l0cSize);
94	compileInfoPtr_->workspaceNum = ascendcPlatform.GetLibApiWorkSpaceSize();	100	compileInfoPtr_->workspaceNum = ascendcPlatform.GetLibApiWorkSpaceSize();
95		101
	@@ -126,7 +132,6 @@ void QuantBatchMatmulV4TilingBase::Reset()
126	inputParams_.mSize = 0L;	132	inputParams_.mSize = 0L;
127	inputParams_.kSize = 0L;	133	inputParams_.kSize = 0L;
128	inputParams_.nSize = 0L;	134	inputParams_.nSize = 0L;
129	- cubeBaseN_ = static_cast<uint64_t>(BLOCK_CUBE);
130	inputParams_.vecInnerAxisAlignUnit = VEC_INNER_AXIS_ALIGN_UINT;	135	inputParams_.vecInnerAxisAlignUnit = VEC_INNER_AXIS_ALIGN_UINT;
131	inputParams_.aDtype = ge::DT_FLOAT16;	136	inputParams_.aDtype = ge::DT_FLOAT16;
132	inputParams_.bDtype = ge::DT_INT8;	137	inputParams_.bDtype = ge::DT_INT8;
	@@ -134,19 +139,11 @@ void QuantBatchMatmulV4TilingBase::Reset()
134	inputParams_.x1ScaleDtype = ge::DT_BF16;	139	inputParams_.x1ScaleDtype = ge::DT_BF16;
135	inputParams_.x2ScaleDtype = ge::DT_BF16;	140	inputParams_.x2ScaleDtype = ge::DT_BF16;
136	inputParams_.biasDtype = ge::DT_FLOAT16;	141	inputParams_.biasDtype = ge::DT_FLOAT16;
137	- aFormat = ge::FORMAT_ND;
138	- bFormat = ge::FORMAT_ND;
139	- cFormat = ge::FORMAT_ND;
140	inputParams_.templateDtype = DtypeEnum::FLOAT16;	142	inputParams_.templateDtype = DtypeEnum::FLOAT16;
141	inputParams_.antiQuantType = QuantType::PER_GROUP;	143	inputParams_.antiQuantType = QuantType::PER_GROUP;
142	- mmInputDtype_ = matmul_tiling::DataType::DT_FLOAT16;	144	+ inputParams_.opName = nullptr;
143	- mmOutputDtype_ = matmul_tiling::DataType::DT_FLOAT16;
144	- mmBiasDtype_ = matmul_tiling::DataType::DT_FLOAT16;
145	- mmScaleADtype_ = matmul_tiling::DataType::DT_BF16;
146	- mmScaleBDtype_ = matmul_tiling::DataType::DT_BF16;
147	aivNum_ = 0;	145	aivNum_ = 0;
148	aicNum_ = 0;	146	aicNum_ = 0;
149	- inputParams_.opName = nullptr;
150	}	147	}
151		148
152	ge::graphStatus QuantBatchMatmulV4TilingBase::GetShapeAttrsInfo()	149	ge::graphStatus QuantBatchMatmulV4TilingBase::GetShapeAttrsInfo()
	@@ -155,7 +152,7 @@ ge::graphStatus QuantBatchMatmulV4TilingBase::GetShapeAttrsInfo()
155	OPS_LOG_D(inputParams_.opName, "TilingContext: %s", Ops::NN::DebugTilingContext(context_).c_str());	152	OPS_LOG_D(inputParams_.opName, "TilingContext: %s", Ops::NN::DebugTilingContext(context_).c_str());
156	auto compileInfoPtr = compileInfoPtr_ ?	153	auto compileInfoPtr = compileInfoPtr_ ?
157	compileInfoPtr_.get() :	154	compileInfoPtr_.get() :
158	- reinterpret_cast<const QuantBatchMatmulV4CompileInfo*>(context_->GetCompileInfo());	155	+ static_cast<const QuantBatchMatmulV4CompileInfo*>(context_->GetCompileInfo());
159	OPS_CHECK_NULL_WITH_CONTEXT(context_, compileInfoPtr);	156	OPS_CHECK_NULL_WITH_CONTEXT(context_, compileInfoPtr);
160	inputParams_.supportL0c2Out = compileInfoPtr->supportL0c2Out;	157	inputParams_.supportL0c2Out = compileInfoPtr->supportL0c2Out;
161	inputParams_.supportL12BtBf16 = compileInfoPtr->supportL12BtBf16;	158	inputParams_.supportL12BtBf16 = compileInfoPtr->supportL12BtBf16;
	@@ -223,6 +220,120 @@ ge::graphStatus QuantBatchMatmulV4TilingBase::CheckContext() const
223	return ge::GRAPH_SUCCESS;	220	return ge::GRAPH_SUCCESS;
224	}	221	}
225		222
		223	+bool QuantBatchMatmulV4TilingBase::CheckA8W4Params() const
		224	+{
		225	+ OP_CHECK_IF(inputParams_.transA,
		226	+ OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(inputParams_.opName, "transposeX1", "true",
		227	+ "The value of transposeX1 must be false"),
		228	+ return false);
		229	+ OP_CHECK_IF(inputParams_.bFormat == ge::FORMAT_ND && !inputParams_.transB,
		230	+ OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(inputParams_.opName, "transposeX2", "false",
		231	+ "When the format of x2 is ND, transposeX2 must be true"),
		232	+ return false);
		233	+ OP_CHECK_IF(
		234	+ inputParams_.antiQuantType == QuantType::PER_GROUP && inputParams_.bFormat == ge::FORMAT_FRACTAL_NZ &&
		235	+ inputParams_.transB,
		236	+ OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(
		237	+ inputParams_.opName, "transposeX2", "true",
		238	+ "When the quantization mode is pergroup and the format of x2 is FRACTAL_NZ, transposeX2 must be false"),
		239	+ return false);
		240	+
		241	+ if (inputParams_.antiQuantType == QuantType::MX) {
		242	+ OP_CHECK_IF(inputParams_.groupSize != MX_GROUP_SIZE,
		243	+ OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(inputParams_.opName, "groupSize",
		244	+ std::to_string(inputParams_.groupSize).c_str(),
		245	+ "groupSize must be 32 when the quantization mode is MX"),
		246	+ return false);
		247	+ } else {
		248	+ OP_CHECK_IF(
		249	+ inputParams_.groupSize <= 0 \|\| inputParams_.kSize < inputParams_.groupSize,
		250	+ OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(
		251	+ inputParams_.opName, "groupSize", std::to_string(inputParams_.groupSize).c_str(),
		252	+ "groupSize must be greater than 0 and less than kSize(" + std::to_string(inputParams_.kSize) + ")"),
		253	+ return false);
		254	+ }
		255	+
		256	+ OP_CHECK_IF(inputParams_.groupSize % GROUP_ALIGN_SIZE > 0,
		257	+ OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(inputParams_.opName, "groupSize",
		258	+ std::to_string(inputParams_.groupSize).c_str(),
		259	+ "groupSize must be aligned to 32"),
		260	+ return false);
		261	+ OP_CHECK_IF(
		262	+ inputParams_.bFormat == ge::FORMAT_FRACTAL_NZ && inputParams_.nSize % N_ALIGN_SIZE > 0,
		263	+ OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(inputParams_.opName, "nSize", std::to_string(inputParams_.nSize).c_str(),
		264	+ "nSize must be aligned to 8 when the format of x2 is FRACTAL_NZ"),
		265	+ return false);
		266	+ return true;
		267	+}
		268	+
		269	+bool QuantBatchMatmulV4TilingBase::CustomCheck() const
		270	+{
		271	+ if (inputParams_.antiQuantType == QuantType::MX) {
		272	+ OP_CHECK_IF(inputParams_.kSize % K_ALIGN_SIZE_MX > 0,
		273	+ OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(
		274	+ inputParams_.opName, "kSize", std::to_string(inputParams_.kSize).c_str(),
		275	+ "kSize must be aligned to 8 when the quantization mode is MX"),
		276	+ return false);
		277	+ } else {
		278	+ OP_CHECK_IF(inputParams_.kSize % K_ALIGN_SIZE > 0 \|\| inputParams_.kSize <= K_ALIGN_SIZE,
		279	+ OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(inputParams_.opName, "kSize",
		280	+ std::to_string(inputParams_.kSize).c_str(),
		281	+ "kSize must be aligned to 32 and greater than 32"),
		282	+ return false);
		283	+ }
		284	+
		285	+ OP_CHECK_IF((inputParams_.cDtype != ge::DT_BF16) && (inputParams_.cDtype != ge::DT_FLOAT16),
		286	+ OP_LOGE_FOR_INVALID_DTYPE_WITH_REASON(
		287	+ inputParams_.opName, "y", ge::TypeUtils::DataTypeToSerialString(inputParams_.cDtype).c_str(),
		288	+ "The dtype of y must be BF16 or FLOAT16"),
		289	+ return false);
		290	+
		291	+ bool a8w4Flag = (inputParams_.aDtype == ge::DT_HIFLOAT8 \|\| inputParams_.aDtype == ge::DT_FLOAT8_E5M2 \|\|
		292	+ inputParams_.aDtype == ge::DT_FLOAT8_E4M3FN) &&
		293	+ (inputParams_.bDtype == ge::DT_FLOAT4_E2M1 \|\| inputParams_.bDtype == ge::DT_FLOAT);
		294	+ if (a8w4Flag) {
		295	+ return CheckA8W4Params();
		296	+ }
		297	+
		298	+ std::string incorrectVals = std::string("x1:") + ge::TypeUtils::DataTypeToSerialString(inputParams_.aDtype) +
		299	+ ", x2:" + ge::TypeUtils::DataTypeToSerialString(inputParams_.bDtype) +
		300	+ ", y:" + ge::TypeUtils::DataTypeToSerialString(inputParams_.cDtype) +
		301	+ ", groupSize:" + std::to_string(inputParams_.groupSize) +
		302	+ ", transA:" + (inputParams_.transA ? "true" : "false") +
		303	+ ", transB:" + (inputParams_.transB ? "true" : "false");
		304	+ OP_LOGE_FOR_INVALID_VALUES_WITH_REASON(
		305	+ inputParams_.opName, "x1, x2, y, groupSize, transposeX1, transposeX2", incorrectVals.c_str(),
		306	+ "The dtype of x1 must be HIFLOAT8, FLOAT8_E5M2, or FLOAT8_E4M3FN, the dtype of x2 must be FLOAT4_E2M1 or "
		307	+ "FLOAT, and the dtype of y must be BF16 or FLOAT16");
		308	+ return false;
		309	+}
		310	+
		311	+bool QuantBatchMatmulV4TilingBase::CheckCoreNum() const
		312	+{
		313	+ if (aivNum_ == 0 \|\| aicNum_ == 0) {
		314	+ OP_LOGE_FOR_INVALID_VALUES_WITH_REASON(inputParams_.opName, "aicNum, aivNum",
		315	+ std::to_string(aicNum_) + ", " + std::to_string(aivNum_),
		316	+ "aicNum and aivNum must be greater than 0");
		317	+ return false;
		318	+ }
		319	+ if (aivNum_ != CORE_RATIO * aicNum_) {
		320	+ OP_LOGE_FOR_INVALID_VALUES_WITH_REASON(inputParams_.opName, "aicNum, aivNum",
		321	+ std::to_string(aicNum_) + ", " + std::to_string(aivNum_),
		322	+ "aicNum:aivNum must be 1:2");
		323	+ return false;
		324	+ }
		325	+ return true;
		326	+}
		327	+
		328	+uint64_t QuantBatchMatmulV4TilingBase::GetTilingKey() const
		329	+{
		330	+ uint64_t trans = (static_cast<uint64_t>(inputParams_.transA) << 1) \| static_cast<uint64_t>(inputParams_.transB);
		331	+ return GET_TPL_TILING_KEY(trans, static_cast<uint64_t>(inputParams_.antiQuantType),
		332	+ static_cast<uint64_t>(inputParams_.hasAntiQuantOffset),
		333	+ static_cast<uint64_t>(inputParams_.weightNz),
		334	+ static_cast<uint64_t>(KernelTemplateType::BASIS));
		335	+}
		336	+
226	ge::graphStatus QuantBatchMatmulV4TilingBase::CheckInputParams() const	337	ge::graphStatus QuantBatchMatmulV4TilingBase::CheckInputParams() const
227	{	338	{
228	bool maxDimCheck = inputParams_.kSize > MAX_SHAPE_DIM \|\| inputParams_.nSize > MAX_SHAPE_DIM;	339	bool maxDimCheck = inputParams_.kSize > MAX_SHAPE_DIM \|\| inputParams_.nSize > MAX_SHAPE_DIM;
	@@ -282,8 +393,6 @@ bool QuantBatchMatmulV4TilingBase::AnalyzeDtype()
282	auto x2ScaleDesc = context_->GetOptionalInputDesc(X2_SCALE_INDEX);	393	auto x2ScaleDesc = context_->GetOptionalInputDesc(X2_SCALE_INDEX);
283	auto yScaleDesc = context_->GetOptionalInputDesc(Y_SCALE_INDEX);	394	auto yScaleDesc = context_->GetOptionalInputDesc(Y_SCALE_INDEX);
284	inputParams_.cDtype = context_->GetOutputDesc(Y_OUTPUT_INDEX)->GetDataType();	395	inputParams_.cDtype = context_->GetOutputDesc(Y_OUTPUT_INDEX)->GetDataType();
285	- mmInputDtype_ = GetMatmulTilingDtype(inputParams_.aDtype);
286	- mmOutputDtype_ = GetMatmulTilingDtype(inputParams_.cDtype);
287	inputParams_.templateDtype = inputParams_.cDtype == ge::DT_FLOAT16 ? DtypeEnum::FLOAT16 : DtypeEnum::BFLOAT16;	396	inputParams_.templateDtype = inputParams_.cDtype == ge::DT_FLOAT16 ? DtypeEnum::FLOAT16 : DtypeEnum::BFLOAT16;
288	// check x1 dtype	397	// check x1 dtype
289	OP_TILING_CHECK(inputParams_.aDtype != ge::DT_FLOAT8_E5M2 && inputParams_.aDtype != ge::DT_FLOAT8_E4M3FN,	398	OP_TILING_CHECK(inputParams_.aDtype != ge::DT_FLOAT8_E5M2 && inputParams_.aDtype != ge::DT_FLOAT8_E4M3FN,
	@@ -333,7 +442,6 @@ bool QuantBatchMatmulV4TilingBase::AnalyzeBiasDtype(const gert::CompileTimeTenso
333	ge::TypeUtils::DataTypeToSerialString(inputParams_.biasDtype).c_str(),	442	ge::TypeUtils::DataTypeToSerialString(inputParams_.biasDtype).c_str(),
334	"The dtype of bias must be BF16 or FLOAT16"),	443	"The dtype of bias must be BF16 or FLOAT16"),
335	return false);	444	return false);
336	- mmBiasDtype_ = GetMatmulTilingDtype(inputParams_.biasDtype);
337	}	445	}
338		446
339	return true;	447	return true;
	@@ -440,7 +548,6 @@ bool QuantBatchMatmulV4TilingBase::AnalyzeAttrs()
440	inputParams_.groupSize = static_cast<uint64_t>(*groupSizePtr);	548	inputParams_.groupSize = static_cast<uint64_t>(*groupSizePtr);
441	inputParams_.vecInnerAxisAlignUnit = inputParams_.groupSize;	549	inputParams_.vecInnerAxisAlignUnit = inputParams_.groupSize;
442	return AnalyzeTranspose();	550	return AnalyzeTranspose();
443	- ;
444	}	551	}
445		552
446	bool QuantBatchMatmulV4TilingBase::AnalyzeX2InputDim(const gert::StorageShape* x2Shape)	553	bool QuantBatchMatmulV4TilingBase::AnalyzeX2InputDim(const gert::StorageShape* x2Shape)
	@@ -542,20 +649,27 @@ bool QuantBatchMatmulV4TilingBase::AnalyzeShapeSize(const gert::StorageShape* x1
542	kBSize = static_cast<uint64_t>(inputParams_.transB ?	649	kBSize = static_cast<uint64_t>(inputParams_.transB ?
543	x2Shape->GetStorageShape().GetDim(x2ShapeDimSize - 1) :	650	x2Shape->GetStorageShape().GetDim(x2ShapeDimSize - 1) :
544	x2Shape->GetStorageShape().GetDim(x2ShapeDimSize - MATMUL_SHAPE_DIM_NUM));	651	x2Shape->GetStorageShape().GetDim(x2ShapeDimSize - MATMUL_SHAPE_DIM_NUM));
545	- OP_TILING_CHECK(inputParams_.kSize != kBSize,
546	- OP_LOGE_FOR_INVALID_SHAPES_WITH_REASON(inputParams_.opName, "x1, x2", "kSize mismatch",
547	- "The k dimension sizes of x1 and x2 must be equal"),
548	- return false);
549	} else if (x2ShapeDimSize == VALID_WEIGHT_NZ_DIM_NUM) {	652	} else if (x2ShapeDimSize == VALID_WEIGHT_NZ_DIM_NUM) {
550	auto x2OriginShape = x2Shape->GetOriginShape();	653	auto x2OriginShape = x2Shape->GetOriginShape();
551	- auto x2ShapeDimSize = x2OriginShape.GetDimNum();	654	+ auto x2OriginShapeDimSize = x2OriginShape.GetDimNum();
552	- inputParams_.nSize = static_cast<uint64_t>(inputParams_.transB ?	655	+ inputParams_.nSize = static_cast<uint64_t>(
553	- x2OriginShape.GetDim(x2ShapeDimSize - MATMUL_SHAPE_DIM_NUM) :	656	+ inputParams_.transB ? x2OriginShape.GetDim(x2OriginShapeDimSize - MATMUL_SHAPE_DIM_NUM) :
554	- x2OriginShape.GetDim(x2ShapeDimSize - 1)); // - 1: 表示尾轴为n轴	657	+ x2OriginShape.GetDim(x2OriginShapeDimSize - 1));
555	- if (context_->GetInputDesc(X2_INDEX)->GetDataType() == ge::DT_FLOAT && !inputParams_.transB) {	658	+ kBSize = static_cast<uint64_t>(inputParams_.transB ?
		659	+ x2OriginShape.GetDim(x2OriginShapeDimSize - 1) :
		660	+ x2OriginShape.GetDim(x2OriginShapeDimSize - MATMUL_SHAPE_DIM_NUM));
		661	+ }
		662	+ if (context_->GetInputDesc(X2_INDEX)->GetDataType() == ge::DT_FLOAT) {
		663	+ if (inputParams_.transB) {
		664	+ kBSize *= B4_IN_B32_NUMS;
		665	+ } else {
556	inputParams_.nSize *= B4_IN_B32_NUMS;	666	inputParams_.nSize *= B4_IN_B32_NUMS;
557	}	667	}
558	}	668	}
		669	+ OP_TILING_CHECK(inputParams_.kSize != kBSize,
		670	+ OP_LOGE_FOR_INVALID_SHAPES_WITH_REASON(inputParams_.opName, "x1, x2", "kSize mismatch",
		671	+ "The k dimension sizes of x1 and x2 must be equal"),
		672	+ return false);
559	return ValidateShapeDimensions();	673	return ValidateShapeDimensions();
560	}	674	}
561		675
	@@ -750,7 +864,7 @@ ge::graphStatus QuantBatchMatmulV4TilingBase::GetPlatformInfo()
750	{	864	{
751	auto compileInfoPtr = compileInfoPtr_ ?	865	auto compileInfoPtr = compileInfoPtr_ ?
752	compileInfoPtr_.get() :	866	compileInfoPtr_.get() :
753	- reinterpret_cast<const QuantBatchMatmulV4CompileInfo*>(context_->GetCompileInfo());	867	+ static_cast<const QuantBatchMatmulV4CompileInfo*>(context_->GetCompileInfo());
754	OPS_CHECK_NULL_WITH_CONTEXT(context_, compileInfoPtr);	868	OPS_CHECK_NULL_WITH_CONTEXT(context_, compileInfoPtr);
755		869
756	aivNum_ = compileInfoPtr->aivNum;	870	aivNum_ = compileInfoPtr->aivNum;
	@@ -758,11 +872,15 @@ ge::graphStatus QuantBatchMatmulV4TilingBase::GetPlatformInfo()
758	aicoreParams_.blockDim = 0;	872	aicoreParams_.blockDim = 0;
759	aicoreParams_.ubSize = compileInfoPtr->ubSize;	873	aicoreParams_.ubSize = compileInfoPtr->ubSize;
760	aicoreParams_.l1Size = compileInfoPtr->l1Size;	874	aicoreParams_.l1Size = compileInfoPtr->l1Size;
		875	+ aicoreParams_.l0aSize = compileInfoPtr->l0aSize;
		876	+ aicoreParams_.l0bSize = compileInfoPtr->l0bSize;
761	aicoreParams_.l0cSize = compileInfoPtr->l0cSize;	877	aicoreParams_.l0cSize = compileInfoPtr->l0cSize;
762	inputParams_.libApiWorkSpaceSize = compileInfoPtr->workspaceNum;	878	inputParams_.libApiWorkSpaceSize = compileInfoPtr->workspaceNum;
763		879
764	- OP_LOGI(inputParams_.opName, "get platform: aivNum(%u) aicNum(%u) ubSize(%lu) l1Size(%lu) l0cSize(%lu)", aivNum_,	880	+ OP_LOGI(inputParams_.opName,
765	- aicNum_, aicoreParams_.ubSize, aicoreParams_.l1Size, aicoreParams_.l0cSize);	881	+ "get platform: aivNum(%u) aicNum(%u) ubSize(%lu) l1Size(%lu) l0aSize(%lu) l0bSize(%lu) l0cSize(%lu)",
		882	+ aivNum_, aicNum_, aicoreParams_.ubSize, aicoreParams_.l1Size, aicoreParams_.l0aSize, aicoreParams_.l0bSize,
		883	+ aicoreParams_.l0cSize);
766		884
767	if (inputParams_.bDtype == ge::DT_INT4) {	885	if (inputParams_.bDtype == ge::DT_INT4) {
768	OP_TILING_CHECK(!CalcUBSize(1UL, inputParams_.groupSize),	886	OP_TILING_CHECK(!CalcUBSize(1UL, inputParams_.groupSize),
	@@ -772,99 +890,36 @@ ge::graphStatus QuantBatchMatmulV4TilingBase::GetPlatformInfo()
772	return ge::GRAPH_SUCCESS;	890	return ge::GRAPH_SUCCESS;
773	}	891	}
774		892
775	-bool QuantBatchMatmulV4TilingBase::GetTilingFromCache() { return false; }	893	+ge::graphStatus QuantBatchMatmulV4TilingBase::CheckTilingDataCapacity(const void* data, size_t size) const
776	-
777	-ge::graphStatus QuantBatchMatmulV4TilingBase::PostTiling()
778	{	894	{
779	-#ifdef A8W4_TILING	895	+ OPS_CHECK_NULL_WITH_CONTEXT(context_, data);
780	- OP_LOGD(inputParams_.opName, "final tiling data size: %zu", tilingDataSize_);	896	+ OP_TILING_CHECK(context_->GetRawTilingData()->GetCapacity() < size,
781	-	897	+ OP_LOGE(inputParams_.opName, "tiling data capacity %zu < actual tiling data size %zu",
782	- OP_TILING_CHECK(tilingDataSize_ % sizeof(uint64_t) != 0,	898	+ context_->GetRawTilingData()->GetCapacity(), size),
783	- OP_LOGE(inputParams_.opName, "tiling data size[%zu] not aligned to 8", tilingDataSize_),
784	return ge::GRAPH_FAILED);	899	return ge::GRAPH_FAILED);
785	- context_->GetRawTilingData()->SetDataSize(tilingDataSize_);	900	+ return ge::GRAPH_SUCCESS;
786	- uint32_t usedAicNum = tilingData_->cubeNumBlocksM * tilingData_->cubeNumBlocksN;	901	+}
787	- uint32_t usedAivNum = tilingData_->vecNumBlocksK * tilingData_->vecNumBlocksN;
788	- context_->SetBlockDim(std::max(usedAicNum, CalcTschNumBlocks(usedAivNum, aicNum_, aivNum_)));
789		902
790	- OP_TILING_CHECK(!CheckFinalTilingData(), PrintTilingData(false);	903	+ge::graphStatus QuantBatchMatmulV4TilingBase::SerializeTilingData(const void* data, size_t size, uint32_t usedCoreNum)
791	- OP_LOGE(inputParams_.opName, "get invalid tiling data, check above validate rule"),	904	+{
		905	+ OPS_CHECK_NULL_WITH_CONTEXT(context_, data);
		906	+ OP_TILING_CHECK(size % sizeof(uint64_t) != 0,
		907	+ OP_LOGE(inputParams_.opName, "tiling data size[%zu] not aligned to 8", size),
792	return ge::GRAPH_FAILED);	908	return ge::GRAPH_FAILED);
793	- size_t* workspaces = context_->GetWorkspaceSizes(1); // set workspace	909	+ context_->GetRawTilingData()->SetDataSize(size);
		910	+ context_->SetSimdNumBlocks(usedCoreNum);
		911	+
		912	+ size_t* workspaces = context_->GetWorkspaceSizes(1);
794	OPS_CHECK_NULL_WITH_CONTEXT(context_, workspaces);	913	OPS_CHECK_NULL_WITH_CONTEXT(context_, workspaces);
795	workspaces[0] = workspaceSize_;	914	workspaces[0] = workspaceSize_;
796		915
797	- errno_t ret = memcpy_s(context_->GetRawTilingData()->GetData(), context_->GetRawTilingData()->GetCapacity(),	916	+ errno_t ret = memcpy_s(context_->GetRawTilingData()->GetData(), context_->GetRawTilingData()->GetCapacity(), data,
798	- reinterpret_cast<void*>(tilingData_), tilingDataSize_);	917	+ size);
799	if (ret != EOK) {	918	if (ret != EOK) {
800	OP_LOGE(context_->GetNodeName(), "memcpy_s failed, ret=%d", ret);	919	OP_LOGE(context_->GetNodeName(), "memcpy_s failed, ret=%d", ret);
801	return ge::GRAPH_FAILED;	920	return ge::GRAPH_FAILED;
802	}	921	}
803	- PrintTilingData(true);
804	-#endif
805	return ge::GRAPH_SUCCESS;	922	return ge::GRAPH_SUCCESS;
806	}	923	}
807		924
808	-void QuantBatchMatmulV4TilingBase::PrintTilingData(bool debugLevel)
809	-{
810	- if (debugLevel) {
811	- OPS_LOG_D(inputParams_.opName, "%ld", DumpTilingDataToLog(debugLevel));
812	- } else {
813	- OPS_LOG_E(inputParams_.opName, "%ld", DumpTilingDataToLog(debugLevel));
814	- }
815	-}
816	-
817	-int64_t QuantBatchMatmulV4TilingBase::DumpTilingDataToLog(bool debugLevel)
818	-{
819	- std::stringstream ss;
820	- ss << "kAlign: " << tilingData_->kAlign << " nAlign: " << tilingData_->nAlign << " kSize: " << tilingData_->kSize
821	- << " nSize: " << tilingData_->nSize << " mSize: " << tilingData_->mSize
822	- << " groupSize: " << tilingData_->groupSize
823	- << " cubeNumBlocksN: " << static_cast<uint32_t>(tilingData_->cubeNumBlocksN)
824	- << " cubeNumBlocksM: " << static_cast<uint32_t>(tilingData_->cubeNumBlocksM);
825	- if (debugLevel) {
826	- OPS_LOG_D(inputParams_.opName, "tiling data: %s", ss.str().c_str());
827	- } else {
828	- OPS_LOG_E(inputParams_.opName, "tiling data: %s", ss.str().c_str());
829	- }
830	- PrintMatMulTiling();
831	- return 0;
832	-}
833	-
834	-void QuantBatchMatmulV4TilingBase::PrintMatMulTiling() const
835	-{
836	- std::stringstream ss;
837	- auto& matmulTiling = tilingData_->matmulTiling;
838	- ss << "usedCoreNum " << matmulTiling.usedCoreNum << " M " << matmulTiling.M << " N " << matmulTiling.N << " Ka "
839	- << matmulTiling.Ka << " Kb " << matmulTiling.Kb << " singleCoreM " << matmulTiling.singleCoreM << " singleCoreN "
840	- << matmulTiling.singleCoreN << " singleCoreK " << matmulTiling.singleCoreK << " baseM " << matmulTiling.baseM
841	- << " baseN " << matmulTiling.baseN << " baseK " << matmulTiling.baseK << " depthA1 " << matmulTiling.depthA1
842	- << " depthB1 " << matmulTiling.depthB1 << " stepM " << matmulTiling.stepM << " stepN " << matmulTiling.stepN
843	- << " isBias " << matmulTiling.isBias << " transLength " << matmulTiling.transLength << " iterateOrder "
844	- << matmulTiling.iterateOrder << " shareMode " << matmulTiling.shareMode << " shareL1Size "
845	- << matmulTiling.shareL1Size << " shareL0CSize " << matmulTiling.shareL0CSize << " shareUbSize "
846	- << matmulTiling.shareUbSize << " batchM " << matmulTiling.batchM << " batchN " << matmulTiling.batchN
847	- << " stepKa " << matmulTiling.stepKa << " stepKb " << matmulTiling.stepKb << " dbL0A " << matmulTiling.dbL0A
848	- << " dbL0B " << matmulTiling.dbL0B << " dbL0C " << matmulTiling.dbL0C;
849	-
850	- OPS_LOG_I(inputParams_.opName, "matmul tiling: %s", ss.str().c_str());
851	-}
852	-
853	-ge::graphStatus QuantBatchMatmulV4TilingBase::InstantiateTilingData()
854	-{
855	- if (tilingData_ == nullptr) {
856	- tilingDataManager_ = std::make_unique<qbmmv4_tiling::QuantBatchMatmulV4TilingDataParams>();
857	- OPS_CHECK_NULL_WITH_CONTEXT(context_, tilingDataManager_);
858	- tilingData_ = tilingDataManager_.get();
859	- }
860	- OPS_CHECK_NULL_WITH_CONTEXT(context_, tilingData_);
861	- OP_TILING_CHECK(context_->GetRawTilingData()->GetCapacity() < tilingDataSize_,
862	- OP_LOGE(inputParams_.opName, "tiling data capacity %zu < actual tiling data size %zu",
863	- context_->GetRawTilingData()->GetCapacity(), tilingDataSize_),
864	- return ge::GRAPH_FAILED);
865	- return ge::GRAPH_SUCCESS;
866	-}
867	-
868	-REGISTER_TILING_TEMPLATE("QuantBatchMatmulV4", QuantBatchMatmulV4RegBase, BASIC_PRIORITY);
869	-
870	} // namespace optiling	925	} // namespace optiling

Mmatmul/quant_batch_matmul_v4/op_host/op_tiling/arch35/quant_batch_matmul_v4_tiling.h+19-93

	@@ -16,20 +16,18 @@
16	#pragma once	16	#pragma once
17		17
18	#include <cstdint>	18	#include <cstdint>
		19	+#include <memory>
19	#include <vector>	20	#include <vector>
20		21
21	#include "op_host/tiling_templates_registry.h"	22	#include "op_host/tiling_templates_registry.h"
22	#include "../../../../common/op_host/op_tiling/tiling_type.h"	23	#include "../../../../common/op_host/op_tiling/tiling_type.h"
23	#include "op_cache_tiling.h"	24	#include "op_cache_tiling.h"
24		25
25	-#include "quant_batch_matmul_v4_basic_block_tiling.h"
26	#include "../../../../weight_quant_batch_matmul_v2/op_host/op_tiling/weight_quant_batch_matmul_v2_tiling_tool.h"	26	#include "../../../../weight_quant_batch_matmul_v2/op_host/op_tiling/weight_quant_batch_matmul_v2_tiling_tool.h"
27	#include "../quant_batch_matmul_v4_compile_info.h"	27	#include "../quant_batch_matmul_v4_compile_info.h"
28	#include "../../../op_kernel/arch35/quant_batch_matmul_v4_tiling_data_apt.h"	28	#include "../../../op_kernel/arch35/quant_batch_matmul_v4_tiling_data_apt.h"
29		29
30	namespace optiling {	30	namespace optiling {
31	-using matmul_tiling::MatrixTraverse;
32	-using namespace matmul_v4;
33	using Ops::NN::Optiling::TilingBaseClass;	31	using Ops::NN::Optiling::TilingBaseClass;
34	namespace matmul_v4 {	32	namespace matmul_v4 {
35	// dim index	33	// dim index
	@@ -139,6 +137,7 @@ struct QuantBatchMatmulInfo {
139	ge::Format bFormat = ge::FORMAT_ND;	137	ge::Format bFormat = ge::FORMAT_ND;
140	};	138	};
141	} // namespace matmul_v4	139	} // namespace matmul_v4
		140	+using namespace matmul_v4;
142		141
143	class QuantBatchMatmulV4TilingBase : public TilingBaseClass {	142	class QuantBatchMatmulV4TilingBase : public TilingBaseClass {
144	public:	143	public:
	@@ -149,14 +148,7 @@ public:
149	InitCompileInfo();	148	InitCompileInfo();
150	}	149	}
151	}	150	}
152	- explicit QuantBatchMatmulV4TilingBase(gert::TilingContext* context,	151	+
153	- qbmmv4_tiling::QuantBatchMatmulV4TilingDataParams* out)
154	- : TilingBaseClass(context)
155	- {
156	- Reset();
157	- tilingData_ = out;
158	- InitCompileInfo();
159	- }
160	~QuantBatchMatmulV4TilingBase() override = default;	152	~QuantBatchMatmulV4TilingBase() override = default;
161		153
162	void Reset(gert::TilingContext* context) override	154	void Reset(gert::TilingContext* context) override
	@@ -169,10 +161,23 @@ protected:
169	bool IsCapable() override { return true; }	161	bool IsCapable() override { return true; }
170	ge::graphStatus GetPlatformInfo() override;	162	ge::graphStatus GetPlatformInfo() override;
171	ge::graphStatus GetShapeAttrsInfo() override;	163	ge::graphStatus GetShapeAttrsInfo() override;
172	- ge::graphStatus InstantiateTilingData();	164	+ virtual bool SetQuantType(const gert::StorageShape* antiQuantScaleShape,
173	- ge::graphStatus PostTiling() override;	165	+ const gert::StorageShape* antiQuantOffsetShape) = 0;
174	- virtual bool CheckFinalTilingData() { return true; }	166	+ virtual bool CalcUBSize(uint64_t vecSingleN, uint64_t vecSingleK) const = 0;
		167	+ virtual bool CheckCoreNum() const;
		168	+ uint64_t GetTilingKey() const override;
175		169
		170	+ ge::graphStatus SerializeTilingData(const void* data, size_t size, uint32_t usedCoreNum);
		171	+ ge::graphStatus CheckTilingDataCapacity(const void* data, size_t size) const;
		172	+ bool CheckA8W4Params() const;
		173	+ bool CustomCheck() const;
		174	+
		175	+ matmul_v4::QuantBatchMatmulInfo inputParams_;
		176	+ uint32_t aivNum_;
		177	+ uint32_t aicNum_;
		178	+ std::unique_ptr<QuantBatchMatmulV4CompileInfo> compileInfoPtr_;
		179	+
		180	+private:
176	void Reset();	181	void Reset();
177	void InitCompileInfo();	182	void InitCompileInfo();
178	ge::graphStatus CheckContext() const;	183	ge::graphStatus CheckContext() const;
	@@ -181,11 +186,7 @@ protected:
181	bool AnalyzeBiasDtype(const gert::CompileTimeTensorDesc* biasDesc);	186	bool AnalyzeBiasDtype(const gert::CompileTimeTensorDesc* biasDesc);
182	bool AnalyzeX1scaleDtype(const gert::CompileTimeTensorDesc* x1ScaleDesc);	187	bool AnalyzeX1scaleDtype(const gert::CompileTimeTensorDesc* x1ScaleDesc);
183	bool AnalyzeX2scaleDtype(const gert::CompileTimeTensorDesc* x2ScaleDesc);	188	bool AnalyzeX2scaleDtype(const gert::CompileTimeTensorDesc* x2ScaleDesc);
184	- bool AnalyzeAntiQuantDtype(ge::DataType antiQuantScaleDtype,
185	- const gert::CompileTimeTensorDesc* antiQuantOffsetDesc) const;
186	bool AnalyzeYScaleOffsetShape(const gert::StorageShape* yScaleShape, const gert::StorageShape* yOffsetShape) const;	189	bool AnalyzeYScaleOffsetShape(const gert::StorageShape* yScaleShape, const gert::StorageShape* yOffsetShape) const;
187	- bool AnalyzeAntiQuantShape(const gert::StorageShape* antiQuantScaleShape,
188	- const gert::StorageShape* antiQuantOffsetShape);
189	bool AnalyzeTranspose();	190	bool AnalyzeTranspose();
190	bool AnalyzeAttrs();	191	bool AnalyzeAttrs();
191	bool AnalyzeX2InputDim(const gert::StorageShape* x2Shape);	192	bool AnalyzeX2InputDim(const gert::StorageShape* x2Shape);
	@@ -197,81 +198,6 @@ protected:
197	bool AnalyzeX1ScaleShape(const gert::StorageShape* x1ScaleShape);	198	bool AnalyzeX1ScaleShape(const gert::StorageShape* x1ScaleShape);
198	bool AnalyzeX2ScaleShape(const gert::StorageShape* x2ScaleShape);	199	bool AnalyzeX2ScaleShape(const gert::StorageShape* x2ScaleShape);
199	bool AnalyzeQuantType();	200	bool AnalyzeQuantType();
200	- virtual bool SetQuantType(const gert::StorageShape* antiQuantScaleShape,
201	- const gert::StorageShape* antiQuantOffsetShape) = 0;
202	- virtual bool CalcUBSize(uint64_t vecSingleN, uint64_t vecSingleK) const = 0;
203	- void PrintTilingData(bool debugLevel);
204	- int64_t DumpTilingDataToLog(bool debugLevel);
205	- virtual void PrintMatMulTiling() const;
206	- virtual bool GetTilingFromCache();
207	-
208	- uint32_t CalcAntiQuantTmpSize(uint64_t vecSingleN, uint64_t vecSingleK) const;
209	- void Convert2AscendCTiling(const CacheTilingData& tbeTiling, TCubeTiling& matmulTiling);
210	- MatrixTraverse GetIteratorOrder(const CacheTilingData& tbeTiling, int32_t singleCoreM, int32_t singleCoreN,
211	- int32_t singleCoreK) const;
212	- virtual bool CheckCoreNum() const { return true; }
213	- matmul_v4::QuantBatchMatmulInfo inputParams_;
214	- uint64_t cubeBaseN_;
215	- int32_t templateId_ = -1;
216	- ge::Format aFormat;
217	- ge::Format bFormat;
218	- ge::Format cFormat;
219	-
220	- matmul_tiling::DataType mmInputDtype_;
221	- matmul_tiling::DataType mmOutputDtype_;
222	- matmul_tiling::DataType mmBiasDtype_;
223	- matmul_tiling::DataType mmScaleADtype_;
224	- matmul_tiling::DataType mmScaleBDtype_;
225	- uint32_t aivNum_;
226	- uint32_t aicNum_;
227	- std::unique_ptr<qbmmv4_tiling::QuantBatchMatmulV4TilingDataParams> tilingDataManager_;
228	- qbmmv4_tiling::QuantBatchMatmulV4TilingDataParams* tilingData_ = nullptr;
229	- size_t tilingDataSize_ = sizeof(qbmmv4_tiling::QuantBatchMatmulV4TilingDataParams);
230	- std::unique_ptr<QuantBatchMatmulV4CompileInfo> compileInfoPtr_;
231	};	201	};
232		202
233	-class QuantBatchMatmulV4RegBase : public QuantBatchMatmulV4TilingBase {
234	-public:
235	- explicit QuantBatchMatmulV4RegBase(gert::TilingContext* context) : QuantBatchMatmulV4TilingBase(context)
236	- {
237	- tilingSolver_.Init();
238	- }
239	- ~QuantBatchMatmulV4RegBase() override = default;
240	-
241	-protected:
242	- bool IsCapable() override;
243	- void SetBubTiling();
244	- void GetBubTilingA8W4(int64_t& nBubSize, int64_t& kBubSize) const;
245	- void GetBubTilingA8W4BySize(int64_t& nBubSize, int64_t& kBubSize, int64_t& kBl1Size, int64_t& nBl1Size) const;
246	- bool CustomCheck() const;
247	- bool CheckA8W4Params() const;
248	- ge::graphStatus DoOpTiling() override;
249	- ge::graphStatus DoLibApiTiling() override { return ge::GRAPH_SUCCESS; }
250	- uint64_t GetTilingKey() const override;
251	- ge::graphStatus GetWorkspaceSize() override;
252	- void SetMatmulTiling();
253	- bool CalcUBSize(uint64_t vecSingleN, uint64_t vecSingleK) const override
254	- {
255	- (void)vecSingleN;
256	- (void)vecSingleK;
257	- return true;
258	- }
259	-
260	- bool SetQuantType(const gert::StorageShape* quantScaleShape, const gert::StorageShape* quantOffsetShape) override
261	- {
262	- (void)quantScaleShape;
263	- (void)quantOffsetShape;
264	- return true;
265	- }
266	- void UpdateL1Tiling(uint64_t minKL1AL1Size, uint64_t minKL1BL1Size, uint64_t fullLoadAl1Size,
267	- uint64_t fullLoadBl1Size, uint64_t minKL1);
268	- uint64_t GetGroupNumBub(uint64_t kDimSzie) const;
269	- uint64_t GetBubSize(uint64_t bubN, uint64_t bubD, bool isWeightNz) const;
270	- void PrintCVTilingData(const bool debugLevel) const;
271	- int64_t DumpCVTilingDataToLog(const bool debugLevel) const;
272	- ge::graphStatus PostTiling() override;
273	- bool CheckCoreNum() const override;
274	-
275	- QuantBatchMatmulV4BasicBlockTiling tilingSolver_;
276	-};
277	} // namespace optiling	203	} // namespace optiling

Amatmul/quant_batch_matmul_v4/op_host/op_tiling/arch35/quant_batch_matmul_v4_weight_quant_mx_swat_tiling.cpp+174-0

@@ -0,0 +1,174 @@
	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 quant_batch_matmul_v4_weight_quant_mx_swat_tiling.cpp
	13	+ * \brief
	14	+ */
	15	+
	16	+#include "quant_batch_matmul_v4_weight_quant_mx_swat_tiling.h"
	17	+
	18	+#include <sstream>
	19	+#include <string>
	20	+
	21	+#include "error_util.h"
	22	+#include "log/log.h"
	23	+#include "quant_batch_matmul_v4_weight_quant_mx_swat_tiling_solver.h"
	24	+
	25	+namespace optiling {
	26	+using namespace matmul_v4;
	27	+
	28	+namespace {
	29	+constexpr uint64_t DB_SIZE = 2UL;
	30	+constexpr uint64_t L1_FOUR_BUFFER = 4UL;
	31	+constexpr const char* OP_TYPE = "QuantBatchMatmulV4";
	32	+
	33	+struct WeightQuantMxSwatScenario {
	34	+ ge::DataType aDtype;
	35	+ ge::DataType bDtype;
	36	+ ge::DataType x1ScaleDtype;
	37	+ ge::DataType x2ScaleDtype;
	38	+};
	39	+
	40	+bool MatchWeightQuantMxSwatScenario(const WeightQuantMxSwatScenario& scenario)
	41	+{
	42	+ if (scenario.aDtype == ge::DT_FLOAT8_E4M3FN &&
	43	+ (scenario.bDtype == ge::DT_FLOAT4_E2M1 \|\| scenario.bDtype == ge::DT_FLOAT) &&
	44	+ scenario.x1ScaleDtype == ge::DT_FLOAT8_E8M0 && scenario.x2ScaleDtype == ge::DT_FLOAT8_E8M0) {
	45	+ return true;
	46	+ }
	47	+ return false;
	48	+}
	49	+
	50	+ge::graphStatus GetWeightQuantMxSwatScenario(const gert::TilingContext* context, WeightQuantMxSwatScenario& scenario)
	51	+{
	52	+ if (context == nullptr) {
	53	+ OP_LOGE(OP_TYPE, "Tiling context is null.");
	54	+ return ge::GRAPH_FAILED;
	55	+ }
	56	+
	57	+ const auto* x1Desc = context->GetInputDesc(X1_INDEX);
	58	+ const auto* x2Desc = context->GetInputDesc(X2_INDEX);
	59	+ const auto* x1ScaleDesc = context->GetOptionalInputDesc(X1_SCALE_INDEX);
	60	+ const auto* x2ScaleDesc = context->GetOptionalInputDesc(X2_SCALE_INDEX);
	61	+ if (unlikely(x1Desc == nullptr \|\| x2Desc == nullptr)) {
	62	+ OP_LOGE(context, "Get required x1/x2 desc failed");
	63	+ return ge::GRAPH_FAILED;
	64	+ }
	65	+ if (unlikely(x1ScaleDesc == nullptr \|\| x2ScaleDesc == nullptr)) {
	66	+ const bool hasMxScale = (x1ScaleDesc != nullptr && x1ScaleDesc->GetDataType() == ge::DT_FLOAT8_E8M0) \|\|
	67	+ (x2ScaleDesc != nullptr && x2ScaleDesc->GetDataType() == ge::DT_FLOAT8_E8M0);
	68	+ if (hasMxScale) {
	69	+ OP_LOGE(context, "MX x1Scale/x2Scale desc is incomplete");
	70	+ return ge::GRAPH_FAILED;
	71	+ }
	72	+ OP_LOGD(context, "x1Scale/x2Scale desc is missing, skip MX SWAT template");
	73	+ return ge::GRAPH_PARAM_INVALID;
	74	+ }
	75	+ scenario = {x1Desc->GetDataType(), x2Desc->GetDataType(), x1ScaleDesc->GetDataType(), x2ScaleDesc->GetDataType()};
	76	+ return ge::GRAPH_SUCCESS;
	77	+}
	78	+} // namespace
	79	+
	80	+ge::graphStatus QuantBatchMatmulV4WeightQuantMxSwatTiling::GetShapeAttrsInfo()
	81	+{
	82	+ WeightQuantMxSwatScenario scenario{};
	83	+ auto status = GetWeightQuantMxSwatScenario(context_, scenario);
	84	+ if (status != ge::GRAPH_SUCCESS) {
	85	+ return status;
	86	+ }
	87	+ if (!MatchWeightQuantMxSwatScenario(scenario)) {
	88	+ return ge::GRAPH_PARAM_INVALID;
	89	+ }
	90	+ return QuantBatchMatmulV4TilingBase::GetShapeAttrsInfo();
	91	+}
	92	+
	93	+bool QuantBatchMatmulV4WeightQuantMxSwatTiling::IsCapable() { return IsWeightQuantMxSwatScenario(); }
	94	+
	95	+bool QuantBatchMatmulV4WeightQuantMxSwatTiling::IsWeightQuantMxSwatScenario() const
	96	+{
	97	+ WeightQuantMxSwatScenario scenario{inputParams_.aDtype, inputParams_.bDtype, inputParams_.x1ScaleDtype,
	98	+ inputParams_.x2ScaleDtype};
	99	+ return MatchWeightQuantMxSwatScenario(scenario);
	100	+}
	101	+
	102	+ge::graphStatus QuantBatchMatmulV4WeightQuantMxSwatTiling::DoOpTiling()
	103	+{
	104	+ OP_TILING_CHECK(CheckTilingDataCapacity(&tilingData_, sizeof(tilingData_)) != ge::GRAPH_SUCCESS,
	105	+ OP_LOGE(inputParams_.opName, "unable to get pointer of SWAT tiling data"), return ge::GRAPH_FAILED);
	106	+ OP_CHECK_IF(!CustomCheck(), OP_LOGE(inputParams_.opName, "Custom check failed."), return ge::GRAPH_FAILED);
	107	+ OP_CHECK_IF(!CheckCoreNum(), OP_LOGE(inputParams_.opName, "Check CoreNum failed."), return ge::GRAPH_FAILED);
	108	+
	109	+ WeightQuantMxSwatPlatformParam platform = {aicNum_,
	110	+ aicoreParams_.ubSize,
	111	+ aicoreParams_.l1Size,
	112	+ aicoreParams_.l0aSize,
	113	+ aicoreParams_.l0bSize,
	114	+ aicoreParams_.l0cSize};
	115	+ WeightQuantMxSwatShapeParam shape = {inputParams_.mSize, inputParams_.nSize, inputParams_.kSize};
	116	+ std::string reason;
	117	+ WeightQuantMxSwatTilingSolver doubleBufferSolver(DB_SIZE);
	118	+ if (!doubleBufferSolver.Solve(platform, shape, inputParams_.groupSize, inputParams_.hasBias,
	119	+ inputParams_.hasX1Scale, inputParams_.hasX2Scale, inputParams_.weightNz,
	120	+ inputParams_.cDtype, tilingData_, reason)) {
	121	+ OP_LOGD(inputParams_.opName, "2-buffer SWAT tiling skipped: %s", reason.c_str());
	122	+ reason.clear();
	123	+ WeightQuantMxSwatTilingSolver fourBufferSolver(L1_FOUR_BUFFER, true);
	124	+ OP_CHECK_IF(!fourBufferSolver.Solve(platform, shape, inputParams_.groupSize, inputParams_.hasBias,
	125	+ inputParams_.hasX1Scale, inputParams_.hasX2Scale, inputParams_.weightNz,
	126	+ inputParams_.cDtype, tilingData_, reason),
	127	+ OP_LOGE(inputParams_.opName, "Unable to get SWAT tiling for mnk[%lu, %lu, %lu]: %s",
	128	+ inputParams_.mSize, inputParams_.nSize, inputParams_.kSize, reason.c_str()),
	129	+ return ge::GRAPH_FAILED);
	130	+ }
	131	+ return ge::GRAPH_SUCCESS;
	132	+}
	133	+
	134	+ge::graphStatus QuantBatchMatmulV4WeightQuantMxSwatTiling::GetWorkspaceSize()
	135	+{
	136	+ workspaceSize_ = 0UL;
	137	+ return ge::GRAPH_SUCCESS;
	138	+}
	139	+
	140	+ge::graphStatus QuantBatchMatmulV4WeightQuantMxSwatTiling::PostTiling()
	141	+{
	142	+ auto status = SerializeTilingData(&tilingData_, sizeof(tilingData_), tilingData_.usedCoreNum);
	143	+ if (status != ge::GRAPH_SUCCESS) {
	144	+ return status;
	145	+ }
	146	+ PrintSwatTilingData(true);
	147	+ return ge::GRAPH_SUCCESS;
	148	+}
	149	+
	150	+void QuantBatchMatmulV4WeightQuantMxSwatTiling::PrintSwatTilingData(bool debugLevel) const
	151	+{
	152	+ DumpSwatTilingDataToLog(debugLevel);
	153	+}
	154	+
	155	+void QuantBatchMatmulV4WeightQuantMxSwatTiling::DumpSwatTilingDataToLog(bool debugLevel) const
	156	+{
	157	+ std::stringstream ss;
	158	+ ss << "m/n/k: " << tilingData_.m << "/" << tilingData_.n << "/" << tilingData_.k
	159	+ << " baseM/baseN/baseK: " << tilingData_.baseM << "/" << tilingData_.baseN << "/" << tilingData_.baseK
	160	+ << " tileShapeKL1: " << tilingData_.tileShapeKL1 << " tileShapeScaleKL1: " << tilingData_.tileShapeScaleKL1
	161	+ << " usedCoreNum: " << tilingData_.usedCoreNum << " cubeNumBlocksM/N: " << tilingData_.cubeNumBlocksM << "/"
	162	+ << tilingData_.cubeNumBlocksN << " mTailTile/nTailTile: " << tilingData_.mTailTile << "/"
	163	+ << tilingData_.nTailTile << " mBaseTailSplitCnt/nBaseTailSplitCnt: " << tilingData_.mBaseTailSplitCnt << "/"
	164	+ << tilingData_.nBaseTailSplitCnt << " mTailMain/nTailMain: " << tilingData_.mTailMain << "/"
	165	+ << tilingData_.nTailMain << " nBubSize/kBubSize: " << tilingData_.nBubSize << "/" << tilingData_.kBubSize
	166	+ << " groupSize: " << tilingData_.groupSize << " weightNz: " << tilingData_.weightNz
	167	+ << " hasBias: " << tilingData_.hasBias << " l1BufferNum: " << tilingData_.l1BufferNum;
	168	+ if (debugLevel) {
	169	+ OPS_LOG_D(inputParams_.opName, "SWAT tiling data: %s", ss.str().c_str());
	170	+ } else {
	171	+ OPS_LOG_E(inputParams_.opName, "SWAT tiling data: %s", ss.str().c_str());
	172	+ }
	173	+}
	174	+} // namespace optiling

Rmatmul/quant_batch_matmul_v4/op_kernel/arch35/cmct/kernel/kernel_matmul_mix_with_weight_prologue_nn.h→matmul/quant_batch_matmul_v4/op_host/op_tiling/arch35/quant_batch_matmul_v4_weight_quant_mx_swat_tiling.h+34-25

	@@ -9,39 +9,48 @@
9	*/	9	*/
10		10
11	/*!	11	/*!
12	- * \file kernel_matmul_mix_with_weight_prologue_nn.h	12	+ * \file quant_batch_matmul_v4_weight_quant_mx_swat_tiling.h
13	* \brief	13	* \brief
14	*/	14	*/
15		15
16	#pragma once	16	#pragma once
17		17
18	-#include "cmct/kernel/kernel_matmul_mix_with_weight_prologue.h"	18	+#include "quant_batch_matmul_v4_tiling.h"
19	-#include "../../quant_batch_matmul_v4_tiling_data_apt.h"
20		19
21	-namespace QuantBatchMatmulV4 {	20	+namespace optiling {
22	-namespace Kernel {	21	+class QuantBatchMatmulV4WeightQuantMxSwatTiling : public QuantBatchMatmulV4TilingBase {
23	-using Cmct::Gemm::Kernel::KernelMatmulMixWeightPrologue;
24	-template <class ProblemShape_, class BlockMmad_, class BlockScheduler_, class BlockPrologue_>
25	-class KernelMatmulMixWeightPrologueNN
26	- : public KernelMatmulMixWeightPrologue<ProblemShape_, BlockMmad_, BlockScheduler_, BlockPrologue_> {
27	public:	22	public:
28	- using BlockMmad = BlockMmad_;	23	+ explicit QuantBatchMatmulV4WeightQuantMxSwatTiling(gert::TilingContext* context)
29	- using BlockScheduler = BlockScheduler_;	24	+ : QuantBatchMatmulV4TilingBase(context)
30	- using BlockPrologue = BlockPrologue_;	25	+ {}
31	- using KernelCmct = KernelMatmulMixWeightPrologue<ProblemShape_, BlockMmad_, BlockScheduler_, BlockPrologue_>;	26	+ ~QuantBatchMatmulV4WeightQuantMxSwatTiling() override = default;
32	- using Arguments = typename KernelCmct::Arguments;
33	- using Params = typename KernelCmct::Params;
34		27
35	- __aicore__ inline KernelMatmulMixWeightPrologueNN() = default;	28	+protected:
36	- __aicore__ inline KernelMatmulMixWeightPrologueNN(const Params& params) : KernelCmct(params) {}	29	+ ge::graphStatus GetShapeAttrsInfo() override;
37	- __host_aicore__ static Params ToUnderlyingArguments(Arguments const& args,	30	+ bool IsCapable() override;
38	- const qbmmv4_tiling::QuantBatchMatmulV4TilingDataParams* tiling)	31	+ ge::graphStatus DoOpTiling() override;
		32	+ ge::graphStatus DoLibApiTiling() override { return ge::GRAPH_SUCCESS; }
		33	+ ge::graphStatus GetWorkspaceSize() override;
		34	+ ge::graphStatus PostTiling() override;
		35	+
		36	+ bool CalcUBSize(uint64_t vecSingleN, uint64_t vecSingleK) const override
39	{	37	{
40	- return {.problemShape = args.problemShape,	38	+ (void)vecSingleN;
41	- .mmad = BlockMmad::ToUnderlyingArguments(args.problemShape, args.mmad, tiling),	39	+ (void)vecSingleK;
42	- .prologue = BlockPrologue::ToUnderlyingArguments(args.problemShape, args.prologue, tiling),	40	+ return true;
43	- .scheduler = BlockScheduler::ToUnderlyingArguments(args.problemShape, args.scheduler, tiling)};
44	}	41	}
		42	+ bool SetQuantType(const gert::StorageShape* quantScaleShape, const gert::StorageShape* quantOffsetShape) override
		43	+ {
		44	+ (void)quantScaleShape;
		45	+ (void)quantOffsetShape;
		46	+ return true;
		47	+ }
		48	+
		49	+private:
		50	+ bool IsWeightQuantMxSwatScenario() const;
		51	+ void PrintSwatTilingData(bool debugLevel) const;
		52	+ void DumpSwatTilingDataToLog(bool debugLevel) const;
		53	+
		54	+ qbmmv4_tiling::QuantBatchMatmulV4WeightQuantMxSwatTilingData tilingData_;
45	};	55	};
46	-} // namespace Kernel	56	+} // namespace optiling
47	-} // namespace QuantBatchMatmulV4

Amatmul/quant_batch_matmul_v4/op_host/op_tiling/arch35/quant_batch_matmul_v4_weight_quant_mx_swat_tiling_solver.cpp+577-0

@@ -0,0 +1,577 @@
	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 quant_batch_matmul_v4_weight_quant_mx_swat_tiling_solver.cpp
	13	+ * \brief
	14	+ */
	15	+
	16	+#include "quant_batch_matmul_v4_weight_quant_mx_swat_tiling_solver.h"
	17	+
	18	+#include <algorithm>
	19	+#include <limits>
	20	+
	21	+#include "quant_batch_matmul_v4_tiling.h"
	22	+
	23	+namespace {
	24	+constexpr uint64_t BLOCK_CUBE_SIZE = 16UL;
	25	+constexpr uint64_t NUM_TWO = 2UL;
	26	+constexpr uint64_t DB_SIZE = 2UL;
	27	+constexpr uint64_t L1_FOUR_BUFFER = 4UL;
	28	+constexpr uint64_t STEPK_THRESHOLD = 4UL;
	29	+constexpr uint64_t BASEM_BASEN_RATIO = 2UL;
	30	+constexpr uint64_t BASEK_LIMIT = 4095UL;
	31	+constexpr uint64_t BASIC_BLOCK_SIZE_128 = 128UL;
	32	+constexpr uint64_t BASIC_BLOCK_SIZE_256 = 256UL;
	33	+constexpr uint64_t MTE2_CACHELINE_SIZE = 128UL;
	34	+constexpr uint64_t TAIL_WINDOW_LEN = 4UL;
	35	+constexpr uint64_t SCALE_FACTOR_MAX = 4UL;
	36	+constexpr uint64_t L1_HALF_SIZE = 256UL * 1024UL;
	37	+constexpr uint64_t DATA_SIZE_UINT8 = 1UL;
	38	+constexpr uint64_t DATA_SIZE_FP32 = 4UL;
	39	+constexpr uint64_t VECTOR_REG_BYTES = 256UL;
	40	+constexpr uint64_t INT4_PACK_NUM = 2UL;
	41	+constexpr uint64_t K_ALIGN_SIZE_MX_L1 = 64UL;
	42	+constexpr uint64_t K_ALIGN_SIZE_MX_BLOCK = 32UL;
	43	+constexpr uint32_t ORDER_N = 1U;
	44	+
	45	+static uint64_t CeilDiv(uint64_t lhs, uint64_t rhs) { return rhs == 0UL ? 0UL : (lhs + rhs - 1UL) / rhs; }
	46	+
	47	+static uint64_t Align(uint64_t value, uint64_t align) { return align == 0UL ? value : CeilDiv(value, align) * align; }
	48	+
	49	+static uint64_t FloorAlign(uint64_t value, uint64_t align) { return align == 0UL ? value : value / align * align; }
	50	+} // namespace
	51	+
	52	+namespace optiling {
	53	+using SwatTilingData = qbmmv4_tiling::QuantBatchMatmulV4WeightQuantMxSwatTilingData;
	54	+
	55	+WeightQuantMxSwatTilingSolver::WeightQuantMxSwatTilingSolver(uint64_t targetL1BufferNum, bool allowConservativeFallback)
	56	+ : targetL1BufferNum_(targetL1BufferNum), allowConservativeFallback_(allowConservativeFallback)
	57	+{}
	58	+
	59	+bool WeightQuantMxSwatTilingSolver::Solve(const WeightQuantMxSwatPlatformParam& platform,
	60	+ const WeightQuantMxSwatShapeParam& shape, uint64_t groupSize, bool hasBias,
	61	+ bool hasX1Scale, bool hasX2Scale, bool weightNz, ge::DataType yDtype,
	62	+ SwatTilingData& tilingData, std::string& reason)
	63	+{
	64	+ if (!Init(platform, shape, reason)) {
	65	+ return false;
	66	+ }
	67	+ hasBias_ = hasBias;
	68	+ weightNz_ = weightNz;
	69	+ biasDataSize_ = GetBiasDataSize(yDtype);
	70	+ bool foundTiling = CalcBasicBlock(reason);
	71	+ if (foundTiling) {
	72	+ OptimizeEdgeBasicBlock();
	73	+ CalcTailBasicBlock();
	74	+ foundTiling = CalcPathSpecificL1(reason) && ValidateTilingResult();
	75	+ }
	76	+ if (!foundTiling && (!allowConservativeFallback_ \|\| !CalcConservativeTiling(reason))) {
	77	+ return false;
	78	+ }
	79	+ BuildTilingData(groupSize, hasBias, hasX1Scale, hasX2Scale, weightNz, yDtype, tilingData);
	80	+ return true;
	81	+}
	82	+
	83	+bool WeightQuantMxSwatTilingSolver::CalcConservativeTiling(std::string& reason)
	84	+{
	85	+ runInfo_ = {};
	86	+ runInfo_.baseM = BLOCK_CUBE_SIZE;
	87	+ runInfo_.baseN = BLOCK_CUBE_SIZE;
	88	+ runInfo_.baseK = std::min(shape_.k, K_ALIGN_SIZE_MX_L1);
	89	+ if (!IsL0Feasible(runInfo_.baseM, runInfo_.baseN, runInfo_.baseK)) {
	90	+ reason = "minimum SWAT base shape exceeds L0A/L0B/L0C capacity";
	91	+ return false;
	92	+ }
	93	+
	94	+ runInfo_.mBlockCnt = CeilDiv(shape_.m, runInfo_.baseM);
	95	+ runInfo_.nBlockCnt = CeilDiv(shape_.n, runInfo_.baseN);
	96	+ runInfo_.totalBlockCnt = runInfo_.mBlockCnt * runInfo_.nBlockCnt;
	97	+ runInfo_.tailBlockCnt = runInfo_.totalBlockCnt % platform_.aicNum;
	98	+ runInfo_.mTailSize = shape_.m - (runInfo_.mBlockCnt - 1UL) * runInfo_.baseM;
	99	+ runInfo_.nTailSize = shape_.n - (runInfo_.nBlockCnt - 1UL) * runInfo_.baseN;
	100	+ CalcTailBasicBlock();
	101	+ if (!CalcPathSpecificL1(reason) \|\| !ValidateTilingResult()) {
	102	+ reason = "minimum SWAT base shape cannot satisfy L1/UB capacity";
	103	+ return false;
	104	+ }
	105	+ return true;
	106	+}
	107	+
	108	+bool WeightQuantMxSwatTilingSolver::Init(const WeightQuantMxSwatPlatformParam& platform,
	109	+ const WeightQuantMxSwatShapeParam& shape, std::string& reason)
	110	+{
	111	+ if (targetL1BufferNum_ != DB_SIZE && targetL1BufferNum_ != L1_FOUR_BUFFER) {
	112	+ reason = "target L1 buffer count must be 2 or 4";
	113	+ return false;
	114	+ }
	115	+ if (shape.m == 0UL \|\| shape.n == 0UL \|\| shape.k == 0UL) {
	116	+ reason = "m, n, and k must be greater than zero";
	117	+ return false;
	118	+ }
	119	+ if ((shape.k % matmul_v4::K_ALIGN_SIZE_MX) != 0UL) {
	120	+ reason = "k must be aligned to 8";
	121	+ return false;
	122	+ }
	123	+ if (shape.m > std::numeric_limits<uint32_t>::max() \|\| shape.n > std::numeric_limits<uint32_t>::max() \|\|
	124	+ shape.k > std::numeric_limits<uint32_t>::max()) {
	125	+ reason = "m, n, and k must not exceed UINT32_MAX";
	126	+ return false;
	127	+ }
	128	+
	129	+ if (platform.aicNum == 0UL) {
	130	+ reason = "AI Core count must be greater than zero";
	131	+ return false;
	132	+ }
	133	+ if (platform.ubSize == 0UL \|\| platform.l1Size == 0UL \|\| platform.l0aSize == 0UL \|\| platform.l0bSize == 0UL \|\|
	134	+ platform.l0cSize == 0UL) {
	135	+ reason = "UB, L1, L0A, L0B, and L0C sizes must be greater than zero";
	136	+ return false;
	137	+ }
	138	+
	139	+ shape_ = shape;
	140	+ platform_ = platform;
	141	+ runInfo_ = {};
	142	+ return true;
	143	+}
	144	+
	145	+bool WeightQuantMxSwatTilingSolver::CalcBasicBlock(std::string& reason)
	146	+{
	147	+ runInfo_.baseM = Align(std::min(shape_.m, BASIC_BLOCK_SIZE_256), BLOCK_CUBE_SIZE);
	148	+ runInfo_.baseN = Align(std::min(shape_.n, BASIC_BLOCK_SIZE_256), BLOCK_CUBE_SIZE);
	149	+ runInfo_.baseK = Align(std::min(shape_.k, BASIC_BLOCK_SIZE_128), matmul_v4::K_ALIGN_SIZE_MX);
	150	+
	151	+ uint64_t blockNum = CeilDiv(shape_.m, runInfo_.baseM) * CeilDiv(shape_.n, runInfo_.baseN);
	152	+ if (blockNum < platform_.aicNum) {
	153	+ AdjustBasicBlock();
	154	+ }
	155	+
	156	+ if (runInfo_.baseM == 0UL \|\| runInfo_.baseN == 0UL \|\| runInfo_.baseK == 0UL) {
	157	+ reason = "baseM, baseN, and baseK must be non-zero";
	158	+ return false;
	159	+ }
	160	+ if (!IsL0Feasible(runInfo_.baseM, runInfo_.baseN, runInfo_.baseK)) {
	161	+ reason = "base shape exceeds L0A/L0B/L0C capacity";
	162	+ return false;
	163	+ }
	164	+
	165	+ runInfo_.mBlockCnt = CeilDiv(shape_.m, runInfo_.baseM);
	166	+ runInfo_.nBlockCnt = CeilDiv(shape_.n, runInfo_.baseN);
	167	+ runInfo_.totalBlockCnt = runInfo_.mBlockCnt * runInfo_.nBlockCnt;
	168	+ runInfo_.tailBlockCnt = runInfo_.totalBlockCnt % platform_.aicNum;
	169	+ runInfo_.mTailSize = shape_.m - (runInfo_.mBlockCnt - 1UL) * runInfo_.baseM;
	170	+ runInfo_.nTailSize = shape_.n - (runInfo_.nBlockCnt - 1UL) * runInfo_.baseN;
	171	+ return true;
	172	+}
	173	+
	174	+void WeightQuantMxSwatTilingSolver::AdjustBasicBlock()
	175	+{
	176	+ uint64_t mMaxTile = CeilDiv(shape_.m, BLOCK_CUBE_SIZE);
	177	+ uint64_t nMaxTile = CeilDiv(shape_.n, BLOCK_CUBE_SIZE);
	178	+ uint64_t tempBaseM = runInfo_.baseM;
	179	+ uint64_t tempBaseN = runInfo_.baseN;
	180	+
	181	+ uint64_t mCnt = std::max<uint64_t>(1UL, CeilDiv(shape_.m, runInfo_.baseM));
	182	+ uint64_t nCnt = std::max<uint64_t>(1UL, CeilDiv(shape_.n, runInfo_.baseN));
	183	+ if (mMaxTile > nMaxTile) {
	184	+ tempBaseN = Align(CeilDiv(shape_.n, nCnt), BLOCK_CUBE_SIZE);
	185	+ nCnt = std::max<uint64_t>(1UL, CeilDiv(shape_.n, tempBaseN));
	186	+ mCnt = std::max<uint64_t>(1UL, platform_.aicNum / nCnt);
	187	+ tempBaseM = Align(CeilDiv(shape_.m, mCnt), BLOCK_CUBE_SIZE);
	188	+ } else {
	189	+ tempBaseM = Align(CeilDiv(shape_.m, mCnt), BLOCK_CUBE_SIZE);
	190	+ mCnt = std::max<uint64_t>(1UL, CeilDiv(shape_.m, tempBaseM));
	191	+ nCnt = std::max<uint64_t>(1UL, platform_.aicNum / mCnt);
	192	+ tempBaseN = Align(CeilDiv(shape_.n, nCnt), BLOCK_CUBE_SIZE);
	193	+ }
	194	+
	195	+ mCnt = std::max<uint64_t>(1UL, CeilDiv(shape_.m, tempBaseM));
	196	+ nCnt = std::max<uint64_t>(1UL, CeilDiv(shape_.n, tempBaseN));
	197	+ while (tempBaseN > tempBaseM * BASEM_BASEN_RATIO && nCnt < platform_.aicNum / NUM_TWO &&
	198	+ tempBaseN != BLOCK_CUBE_SIZE) {
	199	+ nCnt *= NUM_TWO;
	200	+ mCnt = std::max<uint64_t>(1UL, platform_.aicNum / nCnt);
	201	+ tempBaseM = Align(CeilDiv(shape_.m, mCnt), BLOCK_CUBE_SIZE);
	202	+ tempBaseN = Align(CeilDiv(shape_.n, nCnt), BLOCK_CUBE_SIZE);
	203	+ mCnt = std::max<uint64_t>(1UL, CeilDiv(shape_.m, tempBaseM));
	204	+ nCnt = std::max<uint64_t>(1UL, CeilDiv(shape_.n, tempBaseN));
	205	+ }
	206	+ while (tempBaseM >= tempBaseN * BASEM_BASEN_RATIO && mCnt < platform_.aicNum / NUM_TWO &&
	207	+ tempBaseM != BLOCK_CUBE_SIZE) {
	208	+ mCnt *= NUM_TWO;
	209	+ nCnt = std::max<uint64_t>(1UL, platform_.aicNum / mCnt);
	210	+ tempBaseM = Align(CeilDiv(shape_.m, mCnt), BLOCK_CUBE_SIZE);
	211	+ tempBaseN = Align(CeilDiv(shape_.n, nCnt), BLOCK_CUBE_SIZE);
	212	+ mCnt = std::max<uint64_t>(1UL, CeilDiv(shape_.m, tempBaseM));
	213	+ nCnt = std::max<uint64_t>(1UL, CeilDiv(shape_.n, tempBaseN));
	214	+ }
	215	+
	216	+ uint64_t kAlignValue = Align(shape_.k, BASIC_BLOCK_SIZE_128);
	217	+ uint64_t kMaxValue = (platform_.l0aSize / DB_SIZE) / std::max(tempBaseM, tempBaseN);
	218	+ kMaxValue = FloorAlign(kMaxValue, BASIC_BLOCK_SIZE_128);
	219	+ if (kMaxValue >= BASIC_BLOCK_SIZE_128 && IsL0Feasible(tempBaseM, tempBaseN, std::min(kAlignValue, kMaxValue))) {
	220	+ runInfo_.baseM = tempBaseM;
	221	+ runInfo_.baseN = tempBaseN;
	222	+ runInfo_.baseK = std::min(kAlignValue, kMaxValue);
	223	+ runInfo_.baseK = runInfo_.baseK > BASEK_LIMIT ? Align(runInfo_.baseK / NUM_TWO, BASIC_BLOCK_SIZE_256) :
	224	+ runInfo_.baseK;
	225	+ }
	226	+}
	227	+
	228	+void WeightQuantMxSwatTilingSolver::OptimizeEdgeBasicBlock()
	229	+{
	230	+ if (runInfo_.mBlockCnt == 1UL && runInfo_.nBlockCnt == 1UL) {
	231	+ return;
	232	+ }
	233	+
	234	+ bool isInnerAxisAlign = (shape_.k * DATA_SIZE_UINT8) % MTE2_CACHELINE_SIZE == 0UL;
	235	+ uint64_t mTailSize = shape_.m % runInfo_.baseM;
	236	+ if (runInfo_.mBlockCnt > 1UL && mTailSize > 0UL && isInnerAxisAlign) {
	237	+ uint64_t baseTailCntMax = std::min((runInfo_.baseM - mTailSize) / BLOCK_CUBE_SIZE, runInfo_.mBlockCnt);
	238	+ uint64_t windowSize = std::min(TAIL_WINDOW_LEN, runInfo_.mBlockCnt);
	239	+ uint64_t mainWindowNum = runInfo_.mBlockCnt / windowSize - 1UL;
	240	+ uint64_t tailWindowSize = runInfo_.mBlockCnt - mainWindowNum * windowSize;
	241	+ uint64_t perfRes = (mainWindowNum + 1UL) * runInfo_.baseM;
	242	+ uint64_t mergeWindowNum = 1UL;
	243	+ for (uint64_t mergeLen = tailWindowSize - 1UL; mergeLen < baseTailCntMax;
	244	+ mergeLen += windowSize, ++mergeWindowNum) {
	245	+ uint64_t newTailMain = Align(CeilDiv(mergeLen * runInfo_.baseM + mTailSize, mergeLen + 1UL),
	246	+ BLOCK_CUBE_SIZE);
	247	+ uint64_t curPerf = (mainWindowNum + 1UL - mergeWindowNum) * runInfo_.baseM + mergeWindowNum * newTailMain;
	248	+ if (curPerf <= perfRes) {
	249	+ perfRes = curPerf;
	250	+ runInfo_.mTailMain = newTailMain;
	251	+ runInfo_.mBaseTailSplitCnt = mergeLen + 1UL;
	252	+ }
	253	+ }
	254	+ }
	255	+
	256	+ uint64_t nTailSize = shape_.n % runInfo_.baseN;
	257	+ if (runInfo_.nBlockCnt > 1UL && nTailSize > 0UL && isInnerAxisAlign) {
	258	+ uint64_t baseTailCntMax = std::min((runInfo_.baseN - nTailSize) / BLOCK_CUBE_SIZE, runInfo_.nBlockCnt);
	259	+ uint64_t windowSize = std::min(TAIL_WINDOW_LEN, runInfo_.nBlockCnt);
	260	+ uint64_t mainWindowNum = runInfo_.nBlockCnt / windowSize - 1UL;
	261	+ uint64_t tailWindowSize = runInfo_.nBlockCnt - mainWindowNum * windowSize;
	262	+ uint64_t perfRes = (mainWindowNum + 1UL) * runInfo_.baseN;
	263	+ uint64_t mergeWindowNum = 1UL;
	264	+ for (uint64_t mergeLen = tailWindowSize - 1UL; mergeLen < baseTailCntMax;
	265	+ mergeLen += windowSize, ++mergeWindowNum) {
	266	+ uint64_t newTailMain = Align(CeilDiv(mergeLen * runInfo_.baseN + nTailSize, mergeLen + 1UL),
	267	+ BLOCK_CUBE_SIZE);
	268	+ uint64_t curPerf = (mainWindowNum + 1UL - mergeWindowNum) * runInfo_.baseN + mergeWindowNum * newTailMain;
	269	+ if (curPerf <= perfRes) {
	270	+ perfRes = curPerf;
	271	+ runInfo_.nTailMain = newTailMain;
	272	+ runInfo_.nBaseTailSplitCnt = mergeLen + 1UL;
	273	+ }
	274	+ }
	275	+ }
	276	+}
	277	+
	278	+void WeightQuantMxSwatTilingSolver::CalcTailBasicBlock()
	279	+{
	280	+ if (runInfo_.tailBlockCnt == 0UL) {
	281	+ return;
	282	+ }
	283	+
	284	+ uint64_t mTile = 1UL;
	285	+ uint64_t nTile = 1UL;
	286	+ uint64_t preSplit = 1UL;
	287	+ uint64_t secSplit = 1UL;
	288	+ uint64_t& preSplitValid = runInfo_.mTailSize >= runInfo_.nTailSize ? mTile : nTile;
	289	+ uint64_t& secSplitValid = runInfo_.mTailSize >= runInfo_.nTailSize ? nTile : mTile;
	290	+ uint64_t mTileMax = CeilDiv(runInfo_.baseM, BLOCK_CUBE_SIZE);
	291	+ uint64_t nTileMax = CeilDiv(runInfo_.baseN, BLOCK_CUBE_SIZE);
	292	+ uint64_t preSplitMax = runInfo_.mTailSize >= runInfo_.nTailSize ? mTileMax : nTileMax;
	293	+ uint64_t secSplitMax = runInfo_.mTailSize >= runInfo_.nTailSize ? nTileMax : mTileMax;
	294	+ bool splitMFirst = runInfo_.mTailSize >= runInfo_.nTailSize;
	295	+ bool updated = true;
	296	+ while (updated) {
	297	+ updated = false;
	298	+ uint64_t currentUsedCoreNum = CalUsedCoreNum(mTile, nTile);
	299	+ uint64_t preCandidateM = splitMFirst ? preSplit + 1UL : secSplit;
	300	+ uint64_t preCandidateN = splitMFirst ? secSplit : preSplit + 1UL;
	301	+ uint64_t preCandidateUsedCoreNum = CalUsedCoreNum(preCandidateM, preCandidateN);
	302	+ if (preSplit < preSplitMax && preCandidateUsedCoreNum <= platform_.aicNum &&
	303	+ preCandidateUsedCoreNum > currentUsedCoreNum) {
	304	+ preSplitValid = ++preSplit;
	305	+ updated = true;
	306	+ currentUsedCoreNum = preCandidateUsedCoreNum;
	307	+ }
	308	+ uint64_t secCandidateM = splitMFirst ? preSplit : secSplit + 1UL;
	309	+ uint64_t secCandidateN = splitMFirst ? secSplit + 1UL : preSplit;
	310	+ uint64_t secCandidateUsedCoreNum = CalUsedCoreNum(secCandidateM, secCandidateN);
	311	+ if (secSplit < secSplitMax && secCandidateUsedCoreNum <= platform_.aicNum &&
	312	+ secCandidateUsedCoreNum > currentUsedCoreNum) {
	313	+ secSplitValid = ++secSplit;
	314	+ updated = true;
	315	+ }
	316	+ }
	317	+
	318	+ runInfo_.mTailTile = mTile;
	319	+ runInfo_.nTailTile = nTile;
	320	+}
	321	+
	322	+bool WeightQuantMxSwatTilingSolver::CalcPathSpecificL1(std::string& reason)
	323	+{
	324	+ uint64_t maxStepK = std::min(STEPK_THRESHOLD, CeilDiv(shape_.k, runInfo_.baseK));
	325	+ for (uint64_t stepK = maxStepK; stepK > 0UL; --stepK) {
	326	+ uint64_t kBl1Size = std::min(shape_.k, stepK * runInfo_.baseK);
	327	+ uint64_t nBl1Size = std::min(shape_.n, runInfo_.baseN);
	328	+ uint64_t nBubSize = 0UL;
	329	+ uint64_t kBubSize = 0UL;
	330	+ if (!FindBubSize(nBl1Size, kBl1Size, nBubSize, kBubSize)) {
	331	+ continue;
	332	+ }
	333	+
	334	+ uint64_t maxScaleFactor = CalcMaxScaleFactor(stepK);
	335	+ for (uint64_t scaleFactor = maxScaleFactor; scaleFactor > 0UL; --scaleFactor) {
	336	+ uint64_t tileShapeKL1 = stepK * runInfo_.baseK;
	337	+ uint64_t tileShapeScaleKL1 = tileShapeKL1 * scaleFactor;
	338	+ if (IsL1Feasible(tileShapeKL1, tileShapeScaleKL1)) {
	339	+ runInfo_.tileShapeKL1 = tileShapeKL1;
	340	+ runInfo_.tileShapeScaleKL1 = tileShapeScaleKL1;
	341	+ runInfo_.nBubSize = nBubSize;
	342	+ runInfo_.kBubSize = kBubSize;
	343	+ return true;
	344	+ }
	345	+ }
	346	+ }
	347	+ reason = "cannot satisfy L1 and path-specific UB capacity constraints";
	348	+ return false;
	349	+}
	350	+
	351	+uint64_t WeightQuantMxSwatTilingSolver::CalcMaxScaleFactor(uint64_t stepK) const
	352	+{
	353	+ uint64_t kL1Size = stepK * runInfo_.baseK;
	354	+ return std::max<uint64_t>(1UL, std::min(SCALE_FACTOR_MAX, CeilDiv(shape_.k, kL1Size)));
	355	+}
	356	+
	357	+bool WeightQuantMxSwatTilingSolver::ValidateTilingResult() const
	358	+{
	359	+ bool hasValidTileShape = runInfo_.baseM > 0UL && runInfo_.baseN > 0UL && runInfo_.baseK > 0UL &&
	360	+ runInfo_.baseK % matmul_v4::K_ALIGN_SIZE_MX == 0UL && runInfo_.tileShapeKL1 > 0UL &&
	361	+ runInfo_.tileShapeKL1 % runInfo_.baseK == 0UL && runInfo_.tileShapeScaleKL1 > 0UL &&
	362	+ runInfo_.tileShapeScaleKL1 % runInfo_.tileShapeKL1 == 0UL;
	363	+ if (!hasValidTileShape) {
	364	+ return false;
	365	+ }
	366	+ uint64_t nBl1Size = std::min(shape_.n, runInfo_.baseN);
	367	+ uint64_t kBl1Size = std::min(shape_.k, runInfo_.tileShapeKL1);
	368	+ uint64_t expectedN = 0UL;
	369	+ uint64_t expectedK = 0UL;
	370	+ bool hasValidBub = FindBubSize(nBl1Size, kBl1Size, expectedN, expectedK);
	371	+ return IsL0Feasible(runInfo_.baseM, runInfo_.baseN, runInfo_.baseK) &&
	372	+ IsL1Feasible(runInfo_.tileShapeKL1, runInfo_.tileShapeScaleKL1) && hasValidBub &&
	373	+ runInfo_.nBubSize == expectedN && runInfo_.kBubSize == expectedK;
	374	+}
	375	+
	376	+bool WeightQuantMxSwatTilingSolver::IsL0Feasible(uint64_t baseM, uint64_t baseN, uint64_t baseK) const
	377	+{
	378	+ uint64_t a2Size = baseM * baseK * DB_SIZE;
	379	+ uint64_t b2Size = baseN * baseK * DB_SIZE;
	380	+ uint64_t cSize = baseM * baseN * DATA_SIZE_FP32;
	381	+ return a2Size <= platform_.l0aSize && b2Size <= platform_.l0bSize && cSize <= platform_.l0cSize;
	382	+}
	383	+
	384	+bool WeightQuantMxSwatTilingSolver::IsL1Feasible(uint64_t tileShapeKL1, uint64_t tileShapeScaleKL1) const
	385	+{
	386	+ if (platform_.l1Size < L1_HALF_SIZE) {
	387	+ return false;
	388	+ }
	389	+ uint64_t kL1SizeAligned = Align(tileShapeKL1, K_ALIGN_SIZE_MX_L1);
	390	+ uint64_t scaleKL1SizeAligned = Align(tileShapeScaleKL1, K_ALIGN_SIZE_MX_L1);
	391	+ uint64_t aL1Size = runInfo_.baseM * kL1SizeAligned * DATA_SIZE_UINT8;
	392	+ uint64_t bL1Size = runInfo_.baseN * kL1SizeAligned * DATA_SIZE_UINT8;
	393	+ uint64_t scaleAL1Size = runInfo_.baseM * scaleKL1SizeAligned * DATA_SIZE_UINT8 / matmul_v4::MX_GROUP_SIZE;
	394	+ uint64_t scaleBL1Size = runInfo_.baseN * scaleKL1SizeAligned * DATA_SIZE_UINT8 / matmul_v4::MX_GROUP_SIZE;
	395	+ uint64_t biasL1Size = hasBias_ ? Align(runInfo_.baseN, BLOCK_CUBE_SIZE) * biasDataSize_ : 0UL;
	396	+ uint64_t buffersPerHalf = targetL1BufferNum_ / NUM_TWO;
	397	+ uint64_t halfL1Use = buffersPerHalf * (aL1Size + bL1Size + biasL1Size) + scaleAL1Size + scaleBL1Size;
	398	+ // Kernel buffer 1 starts at the fixed 256 KiB boundary, so the shorter physical half sets the limit.
	399	+ uint64_t halfL1Limit = std::min(L1_HALF_SIZE, platform_.l1Size - L1_HALF_SIZE);
	400	+ return halfL1Use <= halfL1Limit;
	401	+}
	402	+
	403	+bool WeightQuantMxSwatTilingSolver::IsBubTilingValid(uint64_t nBubSize, uint64_t kBubSize) const
	404	+{
	405	+ return nBubSize > 0UL && kBubSize > 0UL && kBubSize % matmul_v4::K_ALIGN_SIZE_MX == 0UL &&
	406	+ GetBubSize(targetL1BufferNum_, nBubSize, kBubSize) <= platform_.ubSize;
	407	+}
	408	+
	409	+void WeightQuantMxSwatTilingSolver::BuildTilingData(uint64_t groupSize, bool hasBias, bool hasX1Scale, bool hasX2Scale,
	410	+ bool weightNz, ge::DataType yDtype,
	411	+ SwatTilingData& tilingData) const
	412	+{
	413	+ tilingData = {};
	414	+ tilingData.m = static_cast<uint32_t>(shape_.m);
	415	+ tilingData.n = static_cast<uint32_t>(shape_.n);
	416	+ tilingData.k = static_cast<uint32_t>(shape_.k);
	417	+ tilingData.baseM = static_cast<uint32_t>(runInfo_.baseM);
	418	+ tilingData.baseN = static_cast<uint32_t>(runInfo_.baseN);
	419	+ tilingData.baseK = static_cast<uint32_t>(runInfo_.baseK);
	420	+ tilingData.tileShapeKL1 = static_cast<uint32_t>(runInfo_.tileShapeKL1);
	421	+ tilingData.tileShapeScaleKL1 = static_cast<uint32_t>(runInfo_.tileShapeScaleKL1);
	422	+ tilingData.usedCoreNum = static_cast<uint32_t>(runInfo_.totalBlockCnt >= platform_.aicNum ?
	423	+ platform_.aicNum :
	424	+ CalUsedCoreNum(runInfo_.mTailTile, runInfo_.nTailTile));
	425	+ tilingData.cubeNumBlocksM = static_cast<uint32_t>(runInfo_.mBlockCnt);
	426	+ tilingData.cubeNumBlocksN = static_cast<uint32_t>(runInfo_.nBlockCnt);
	427	+ tilingData.iterateOrder = ORDER_N;
	428	+ tilingData.mTailTile = static_cast<uint32_t>(runInfo_.mTailTile);
	429	+ tilingData.nTailTile = static_cast<uint32_t>(runInfo_.nTailTile);
	430	+ tilingData.mBaseTailSplitCnt = static_cast<uint32_t>(runInfo_.mBaseTailSplitCnt);
	431	+ tilingData.nBaseTailSplitCnt = static_cast<uint32_t>(runInfo_.nBaseTailSplitCnt);
	432	+ tilingData.mTailMain = static_cast<uint32_t>(runInfo_.mTailMain);
	433	+ tilingData.nTailMain = static_cast<uint32_t>(runInfo_.nTailMain);
	434	+ tilingData.nBubSize = static_cast<uint32_t>(runInfo_.nBubSize);
	435	+ tilingData.kBubSize = static_cast<uint32_t>(runInfo_.kBubSize);
	436	+ tilingData.groupSize = static_cast<uint32_t>(groupSize);
	437	+ tilingData.hasBias = static_cast<uint32_t>(hasBias);
	438	+ tilingData.hasX1Scale = static_cast<uint32_t>(hasX1Scale);
	439	+ tilingData.hasX2Scale = static_cast<uint32_t>(hasX2Scale);
	440	+ tilingData.weightNz = static_cast<uint32_t>(weightNz);
	441	+ tilingData.yDtype = static_cast<uint32_t>(yDtype);
	442	+ tilingData.l1BufferNum = static_cast<uint32_t>(targetL1BufferNum_);
	443	+}
	444	+
	445	+bool WeightQuantMxSwatTilingSolver::FindBubSize(uint64_t nBl1Size, uint64_t kBl1Size, uint64_t& nBubSize,
	446	+ uint64_t& kBubSize) const
	447	+{
	448	+ nBubSize = nBl1Size;
	449	+ if (weightNz_) {
	450	+ kBubSize = FindKOnlyBubSize(nBubSize, kBl1Size);
	451	+ return kBubSize > 0UL;
	452	+ }
	453	+
	454	+ kBubSize = kBl1Size;
	455	+ if (targetL1BufferNum_ == L1_FOUR_BUFFER && nBl1Size > BLOCK_CUBE_SIZE) {
	456	+ nBubSize = Align(CeilDiv(nBl1Size, NUM_TWO), BLOCK_CUBE_SIZE);
	457	+ }
	458	+ return IsBubTilingValid(nBubSize, kBubSize);
	459	+}
	460	+
	461	+uint64_t WeightQuantMxSwatTilingSolver::FindKOnlyBubSize(uint64_t nBubSize, uint64_t kBl1Size) const
	462	+{
	463	+ if (kBl1Size <= K_ALIGN_SIZE_MX_L1) {
	464	+ return IsBubTilingValid(nBubSize, kBl1Size) ? kBl1Size : 0UL;
	465	+ }
	466	+ uint64_t minK = Align(CeilDiv(kBl1Size, NUM_TWO), K_ALIGN_SIZE_MX_L1);
	467	+ if (minK >= kBl1Size \|\| (minK % matmul_v4::K_ALIGN_SIZE_MX) != 0UL) {
	468	+ return 0UL;
	469	+ }
	470	+ return IsBubTilingValid(nBubSize, minK) ? minK : 0UL;
	471	+}
	472	+
	473	+uint64_t WeightQuantMxSwatTilingSolver::GetBubSize(uint64_t bufferNum, uint64_t nDimSize, uint64_t kDimSize) const
	474	+{
	475	+ uint64_t nDimAlign = Align(nDimSize, BLOCK_CUBE_SIZE);
	476	+ uint64_t kDimBlockAlign = Align(kDimSize, K_ALIGN_SIZE_MX_BLOCK);
	477	+ uint64_t kDimL1Align = Align(kDimSize, K_ALIGN_SIZE_MX_L1);
	478	+ uint64_t sizeWeightIn = 0UL;
	479	+ uint64_t sizeWeightOut = 0UL;
	480	+ if (weightNz_) {
	481	+ sizeWeightIn = bufferNum * DATA_SIZE_UINT8 * nDimAlign * kDimBlockAlign / INT4_PACK_NUM;
	482	+ sizeWeightOut = bufferNum * DATA_SIZE_UINT8 * nDimAlign * kDimL1Align;
	483	+ } else {
	484	+ sizeWeightIn = bufferNum * DATA_SIZE_UINT8 * nDimSize * kDimL1Align / INT4_PACK_NUM;
	485	+ sizeWeightOut = bufferNum * DATA_SIZE_UINT8 * (nDimAlign + 1UL) * kDimL1Align;
	486	+ }
	487	+ uint64_t sizeBias = 0UL;
	488	+ if (hasBias_) {
	489	+ uint64_t biasVectorElements = VECTOR_REG_BYTES / biasDataSize_;
	490	+ // The prologue keeps a complete baseN bias tile even when ND weight conversion is split along N.
	491	+ uint64_t singleBiasBufferSize = Align(runInfo_.baseN, biasVectorElements) * biasDataSize_;
	492	+ sizeBias = NUM_TWO * bufferNum * singleBiasBufferSize;
	493	+ }
	494	+ return sizeWeightIn + sizeWeightOut + sizeBias;
	495	+}
	496	+
	497	+uint64_t WeightQuantMxSwatTilingSolver::GetBiasDataSize(ge::DataType yDtype)
	498	+{
	499	+ switch (yDtype) {
	500	+ case ge::DT_FLOAT:
	501	+ return DATA_SIZE_FP32;
	502	+ case ge::DT_FLOAT16:
	503	+ case ge::DT_BF16:
	504	+ return DATA_SIZE_UINT8 * NUM_TWO;
	505	+ default:
	506	+ return DATA_SIZE_FP32;
	507	+ }
	508	+}
	509	+
	510	+uint64_t WeightQuantMxSwatTilingSolver::CalUsedCoreNum(uint64_t mTile, uint64_t nTile) const
	511	+{
	512	+ uint64_t usedCoreNum = 0UL;
	513	+ uint64_t baseRoundTileNum = runInfo_.totalBlockCnt - runInfo_.tailBlockCnt;
	514	+ for (uint64_t tailIdx = 0UL; tailIdx < runInfo_.tailBlockCnt; ++tailIdx) {
	515	+ uint64_t mTileIdx = 0UL;
	516	+ uint64_t nTileIdx = 0UL;
	517	+ GetLogicalTileCoord(baseRoundTileNum + tailIdx, mTileIdx, nTileIdx);
	518	+ usedCoreNum += CalcValidSplitCount(GetSingleCoreM(mTileIdx), GetSingleCoreN(nTileIdx), mTile, nTile);
	519	+ }
	520	+ return usedCoreNum;
	521	+}
	522	+
	523	+uint64_t WeightQuantMxSwatTilingSolver::CalcValidSplitCount(uint64_t singleCoreM, uint64_t singleCoreN, uint64_t mTile,
	524	+ uint64_t nTile) const
	525	+{
	526	+ uint64_t singleCoreMSplit = Align(CeilDiv(singleCoreM, mTile), BLOCK_CUBE_SIZE);
	527	+ uint64_t singleCoreNSplit = Align(CeilDiv(singleCoreN, nTile), BLOCK_CUBE_SIZE);
	528	+ uint64_t validM = std::min(mTile, CeilDiv(singleCoreM, singleCoreMSplit));
	529	+ uint64_t validN = std::min(nTile, CeilDiv(singleCoreN, singleCoreNSplit));
	530	+ return validM * validN;
	531	+}
	532	+
	533	+void WeightQuantMxSwatTilingSolver::GetLogicalTileCoord(uint64_t tileIdx, uint64_t& mTileIdx, uint64_t& nTileIdx) const
	534	+{
	535	+ uint64_t mCoreNum = std::min(TAIL_WINDOW_LEN, runInfo_.mBlockCnt);
	536	+ uint64_t mainRow = runInfo_.mBlockCnt / mCoreNum - 1UL;
	537	+ uint64_t mTailCoreNum = runInfo_.mBlockCnt - mCoreNum * mainRow;
	538	+ uint64_t rowIdx = tileIdx / (mCoreNum * runInfo_.nBlockCnt);
	539	+ if (rowIdx < mainRow) {
	540	+ uint64_t localTileIdx = tileIdx - rowIdx * mCoreNum * runInfo_.nBlockCnt;
	541	+ mTileIdx = rowIdx * mCoreNum + localTileIdx % mCoreNum;
	542	+ nTileIdx = (localTileIdx / mCoreNum) % runInfo_.nBlockCnt;
	543	+ } else {
	544	+ rowIdx = mainRow;
	545	+ uint64_t tailIdx = tileIdx - mainRow * mCoreNum * runInfo_.nBlockCnt;
	546	+ mTileIdx = mainRow * mCoreNum + tailIdx % mTailCoreNum;
	547	+ nTileIdx = (tailIdx / mTailCoreNum) % runInfo_.nBlockCnt;
	548	+ }
	549	+ if ((rowIdx & 1UL) != 0UL) {
	550	+ nTileIdx = runInfo_.nBlockCnt - 1UL - nTileIdx;
	551	+ }
	552	+}
	553	+
	554	+uint64_t WeightQuantMxSwatTilingSolver::GetSingleCoreM(uint64_t mTileIdx) const
	555	+{
	556	+ uint64_t mBaseNormCnt = runInfo_.mBlockCnt - runInfo_.mBaseTailSplitCnt;
	557	+ if (mTileIdx >= mBaseNormCnt) {
	558	+ uint64_t mMergeSize = shape_.m - mBaseNormCnt * runInfo_.baseM;
	559	+ uint64_t mBaseTailMain = runInfo_.mBaseTailSplitCnt == 1UL ? mMergeSize : runInfo_.mTailMain;
	560	+ uint64_t mBaseTailLast = mMergeSize - (runInfo_.mBaseTailSplitCnt - 1UL) * mBaseTailMain;
	561	+ return mTileIdx < runInfo_.mBlockCnt - 1UL ? mBaseTailMain : mBaseTailLast;
	562	+ }
	563	+ return runInfo_.baseM;
	564	+}
	565	+
	566	+uint64_t WeightQuantMxSwatTilingSolver::GetSingleCoreN(uint64_t nTileIdx) const
	567	+{
	568	+ uint64_t nBaseNormCnt = runInfo_.nBlockCnt - runInfo_.nBaseTailSplitCnt;
	569	+ if (nTileIdx >= nBaseNormCnt) {
	570	+ uint64_t nMergeSize = shape_.n - nBaseNormCnt * runInfo_.baseN;
	571	+ uint64_t nBaseTailMain = runInfo_.nBaseTailSplitCnt == 1UL ? nMergeSize : runInfo_.nTailMain;
	572	+ uint64_t nBaseTailLast = nMergeSize - (runInfo_.nBaseTailSplitCnt - 1UL) * nBaseTailMain;
	573	+ return nTileIdx < runInfo_.nBlockCnt - 1UL ? nBaseTailMain : nBaseTailLast;
	574	+ }
	575	+ return runInfo_.baseN;
	576	+}
	577	+} // namespace optiling

Amatmul/quant_batch_matmul_v4/op_host/op_tiling/arch35/quant_batch_matmul_v4_weight_quant_mx_swat_tiling_solver.h+106-0

@@ -0,0 +1,106 @@
	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 quant_batch_matmul_v4_weight_quant_mx_swat_tiling_solver.h
	13	+ * \brief
	14	+ */
	15	+
	16	+#pragma once
	17	+
	18	+#include <cstdint>
	19	+#include <string>
	20	+
	21	+#include "graph/types.h"
	22	+#include "../../../op_kernel/arch35/quant_batch_matmul_v4_tiling_data_apt.h"
	23	+
	24	+namespace optiling {
	25	+struct WeightQuantMxSwatPlatformParam {
	26	+ uint64_t aicNum = 0UL;
	27	+ uint64_t ubSize = 0UL;
	28	+ uint64_t l1Size = 0UL;
	29	+ uint64_t l0aSize = 0UL;
	30	+ uint64_t l0bSize = 0UL;
	31	+ uint64_t l0cSize = 0UL;
	32	+};
	33	+
	34	+struct WeightQuantMxSwatShapeParam {
	35	+ uint64_t m = 0UL;
	36	+ uint64_t n = 0UL;
	37	+ uint64_t k = 0UL;
	38	+};
	39	+
	40	+class WeightQuantMxSwatTilingSolver {
	41	+public:
	42	+ explicit WeightQuantMxSwatTilingSolver(uint64_t targetL1BufferNum, bool allowConservativeFallback = false);
	43	+
	44	+ bool Solve(const WeightQuantMxSwatPlatformParam& platform, const WeightQuantMxSwatShapeParam& shape,
	45	+ uint64_t groupSize, bool hasBias, bool hasX1Scale, bool hasX2Scale, bool weightNz, ge::DataType yDtype,
	46	+ qbmmv4_tiling::QuantBatchMatmulV4WeightQuantMxSwatTilingData& tilingData, std::string& reason);
	47	+
	48	+private:
	49	+ struct RunInfo {
	50	+ uint64_t baseM = 0UL;
	51	+ uint64_t baseN = 0UL;
	52	+ uint64_t baseK = 0UL;
	53	+ uint64_t tileShapeKL1 = 0UL;
	54	+ uint64_t tileShapeScaleKL1 = 0UL;
	55	+ uint64_t nBubSize = 0UL;
	56	+ uint64_t kBubSize = 0UL;
	57	+ uint64_t mBlockCnt = 0UL;
	58	+ uint64_t nBlockCnt = 0UL;
	59	+ uint64_t totalBlockCnt = 0UL;
	60	+ uint64_t tailBlockCnt = 0UL;
	61	+ uint64_t mTailSize = 0UL;
	62	+ uint64_t nTailSize = 0UL;
	63	+ uint64_t mTailTile = 1UL;
	64	+ uint64_t nTailTile = 1UL;
	65	+ uint64_t mBaseTailSplitCnt = 1UL;
	66	+ uint64_t nBaseTailSplitCnt = 1UL;
	67	+ uint64_t mTailMain = 0UL;
	68	+ uint64_t nTailMain = 0UL;
	69	+ };
	70	+
	71	+ bool Init(const WeightQuantMxSwatPlatformParam& platform, const WeightQuantMxSwatShapeParam& shape,
	72	+ std::string& reason);
	73	+ bool CalcBasicBlock(std::string& reason);
	74	+ bool CalcConservativeTiling(std::string& reason);
	75	+ void AdjustBasicBlock();
	76	+ void OptimizeEdgeBasicBlock();
	77	+ void CalcTailBasicBlock();
	78	+ bool CalcPathSpecificL1(std::string& reason);
	79	+ uint64_t CalcMaxScaleFactor(uint64_t stepK) const;
	80	+ bool ValidateTilingResult() const;
	81	+ bool IsL0Feasible(uint64_t baseM, uint64_t baseN, uint64_t baseK) const;
	82	+ bool IsL1Feasible(uint64_t tileShapeKL1, uint64_t tileShapeScaleKL1) const;
	83	+ bool IsBubTilingValid(uint64_t nBubSize, uint64_t kBubSize) const;
	84	+ void BuildTilingData(uint64_t groupSize, bool hasBias, bool hasX1Scale, bool hasX2Scale, bool weightNz,
	85	+ ge::DataType yDtype,
	86	+ qbmmv4_tiling::QuantBatchMatmulV4WeightQuantMxSwatTilingData& tilingData) const;
	87	+ bool FindBubSize(uint64_t nBl1Size, uint64_t kBl1Size, uint64_t& nBubSize, uint64_t& kBubSize) const;
	88	+ uint64_t FindKOnlyBubSize(uint64_t nBubSize, uint64_t kBl1Size) const;
	89	+ uint64_t GetBubSize(uint64_t bufferNum, uint64_t nDimSize, uint64_t kDimSize) const;
	90	+ uint64_t CalUsedCoreNum(uint64_t mTile, uint64_t nTile) const;
	91	+ uint64_t CalcValidSplitCount(uint64_t singleCoreM, uint64_t singleCoreN, uint64_t mTile, uint64_t nTile) const;
	92	+ void GetLogicalTileCoord(uint64_t tileIdx, uint64_t& mTileIdx, uint64_t& nTileIdx) const;
	93	+ uint64_t GetSingleCoreM(uint64_t mTileIdx) const;
	94	+ uint64_t GetSingleCoreN(uint64_t nTileIdx) const;
	95	+ static uint64_t GetBiasDataSize(ge::DataType yDtype);
	96	+
	97	+ WeightQuantMxSwatShapeParam shape_;
	98	+ WeightQuantMxSwatPlatformParam platform_;
	99	+ RunInfo runInfo_;
	100	+ uint64_t targetL1BufferNum_ = 2UL;
	101	+ bool allowConservativeFallback_ = false;
	102	+ bool hasBias_ = false;
	103	+ bool weightNz_ = false;
	104	+ uint64_t biasDataSize_ = 0UL;
	105	+};
	106	+} // namespace optiling

Mmatmul/quant_batch_matmul_v4/op_host/op_tiling/quant_batch_matmul_v4_tiling_registry.cpp+9-2

Dmatmul/quant_batch_matmul_v4/op_kernel/arch35/cmct/block/block_mmad_mx_weight_from_ub.h+0-59

	@@ -1,59 +0,0 @@
1	-/**
2	- * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3	- * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4	- * CANN Open Software License Agreement Version 2.0 (the "License").
5	- * Please refer to the License for details. You may not use this file except in compliance with the License.
6	- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7	- * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8	- * See LICENSE in the root of the software repository for the full text of the License.
9	- */
10	-/*!
11	- * \file block_mmad_mx_weight_from_ub.h
12	- * \brief
13	- */
14	-
15	-#pragma once
16	-
17	-#include "cmct/block/block_mmad.h"
18	-#include "../../quant_batch_matmul_v4_tiling_data_apt.h"
19	-
20	-namespace QuantBatchMatmulV4 {
21	-namespace Block {
22	-using Cmct::Gemm::Block::BlockMmad;
23	-template <class DispatchPolicy, class L1TileShape_, class L0TileShape_, class ATypeTuple_, class BType_, class CType_,
24	- class BiasType_, class TileCopy_, class TileMmad_>
25	-class BlockMmadNN : public BlockMmad<DispatchPolicy, L1TileShape_, L0TileShape_, ATypeTuple_, BType_, CType_, BiasType_,
26	- TileCopy_, TileMmad_> {
27	-public:
28	- using MmadCmct = BlockMmad<DispatchPolicy, L1TileShape_, L0TileShape_, ATypeTuple_, BType_, CType_, BiasType_,
29	- TileCopy_, TileMmad_>;
30	- using Arguments = typename MmadCmct::Arguments;
31	- using Params = typename MmadCmct::Params;
32	- __aicore__ inline BlockMmadNN() = delete;
33	- __aicore__ inline BlockMmadNN(const Params& params) : MmadCmct(params) {}
34	- template <typename ProblemShape>
35	- __aicore__ inline static Params ToUnderlyingArguments(
36	- ProblemShape const& problemShape, Arguments const& args,
37	- qbmmv4_tiling::QuantBatchMatmulV4TilingDataParams const* tiling)
38	- {
39	- auto stepKa = static_cast<uint32_t>(tiling->matmulTiling.stepKa);
40	- auto stepKb = static_cast<uint32_t>(tiling->matmulTiling.stepKb);
41	- auto baseK = static_cast<uint32_t>(tiling->matmulTiling.baseK);
42	- auto baseM = static_cast<uint32_t>(tiling->matmulTiling.baseM);
43	- auto baseN = static_cast<uint32_t>(tiling->matmulTiling.baseN);
44	- return {.ptrA = args.ptrA,
45	- .ptrC = args.ptrC,
46	- .ptrAScale = args.ptrAScale,
47	- .ptrBScale = args.ptrBScale,
48	- .layoutA = args.layoutA,
49	- .layoutC = args.layoutC,
50	- .layoutScale = args.layoutScale,
51	- .tileShapeL1 = AscendC::MakeShape(baseM, baseN, stepKa * baseK, stepKb * baseK),
52	- .tileShapeL0 = AscendC::MakeShape(baseM, baseN, baseK),
53	- .scaleFactor = tiling->matmulTiling.mxTypePara & 0xff, // 0xff：to obtain the lower 8 bits
54	- .aL1BufNum = tiling->AL1Pingpong,
55	- .isBias = bool(tiling->matmulTiling.isBias)};
56	- }
57	-};
58	-} // namespace Block
59	-} // namespace QuantBatchMatmulV4

Dmatmul/quant_batch_matmul_v4/op_kernel/arch35/cmct/block/block_scheduler_swizzle_in_mn_core_nn.h+0-63

	@@ -1,63 +0,0 @@
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 block_scheduler_swizzle_in_mn_core_nn.h
13	- * \brief
14	- */
15	-
16	-#pragma once
17	-#include "cmct/block/block_scheduler_swizzle_in_mn_core.h"
18	-#include "../../quant_batch_matmul_v4_tiling_data_apt.h"
19	-/*
20	-iterateOrder = 0
21	-scheduler diagram c:core b:block
22	-\| c0b0 \| c0b1 \| c0b2 \| c2b0 \| c2b1 \|
23	-------------------------------------
24	-\| c0b3 \| c0b4 \| c0b5 \| c2b2 \| c2b3 \|
25	-------------------------------------
26	-\| c1b0 \| c1b1 \| c1b2 \| c3b0 \| c3b1 \|f
27	-
28	-iterateOrder = 1
29	-\| c0b0 \| c0b2 \| c0b4 \| c2b0 \| c2b2 \|
30	-------------------------------------
31	-\| c0b1 \| c0b3 \| c0b5 \| c2b1 \| c2b3 \|
32	-------------------------------------
33	-\| c1b0 \| c1b1 \| c1b2 \| c3b0 \| c3b1 \|
34	-*/
35	-namespace QuantBatchMatmulV4 {
36	-namespace Block {
37	-using Cmct::Gemm::Block::BlockSchedulerSwizzleInMnCore;
38	-template <class ProblemShape_, class TileShape_, class BlockShape_>
39	-class BlockSchedulerSwizzleInMnCoreNN : public BlockSchedulerSwizzleInMnCore<ProblemShape_, TileShape_, BlockShape_> {
40	-public:
41	- using ProblemShape = ProblemShape_;
42	- using KernelCmct = BlockSchedulerSwizzleInMnCore<ProblemShape_, TileShape_, BlockShape_>;
43	- using Arguments = typename KernelCmct::Arguments;
44	- using Params = typename KernelCmct::Params;
45	- using BlockSchedulerSwizzleInMnCoreCmct = BlockSchedulerSwizzleInMnCore<ProblemShape_, TileShape_, BlockShape_>;
46	- __aicore__ inline BlockSchedulerSwizzleInMnCoreNN() = delete;
47	- __aicore__ inline BlockSchedulerSwizzleInMnCoreNN(const Params& params) : BlockSchedulerSwizzleInMnCoreCmct(params)
48	- {}
49	-
50	- __aicore__ inline static Params ToUnderlyingArguments(
51	- ProblemShape const& problemShape, [[maybe_unused]] Arguments const& args,
52	- const qbmmv4_tiling::QuantBatchMatmulV4TilingDataParams* tiling)
53	- {
54	- auto baseM = static_cast<uint32_t>(tiling->matmulTiling.baseM);
55	- auto baseN = static_cast<uint32_t>(tiling->matmulTiling.baseN);
56	- return {.iterateOrder = tiling->matmulTiling.iterateOrder,
57	- .problemShape = problemShape,
58	- .tileShape = AscendC::MakeShape(baseM, baseN),
59	- .blockShape = AscendC::MakeShape(tiling->cubeNumBlocksM, tiling->cubeNumBlocksN)};
60	- }
61	-};
62	-} // namespace Block
63	-} // namespace QuantBatchMatmulV4

Dmatmul/quant_batch_matmul_v4/op_kernel/arch35/cmct/prologue/block_prologue_b_cast_scsc_nn.h+0-55

	@@ -1,55 +0,0 @@
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 block_prologue_b_cast_scsc_nn.h
13	- * \brief
14	- */
15	-
16	-#pragma once
17	-#include "cmct/prologue/block_prologue_b_cast_scsc.h"
18	-#include "../../quant_batch_matmul_v4_tiling_data_apt.h"
19	-
20	-namespace QuantBatchMatmulV4 {
21	-namespace Prologue {
22	-using Cmct::Prologue::BlockPrologue;
23	-template <class DispatchPolicy, class InType, class OutType, class BiasType, class TileShapeL1>
24	-class BlockPrologueNN : public BlockPrologue<DispatchPolicy, InType, OutType, BiasType, TileShapeL1> {
25	-public:
26	- using PrologueCmct = BlockPrologue<DispatchPolicy, InType, OutType, BiasType, TileShapeL1>;
27	- using Arguments = typename PrologueCmct::Arguments;
28	- using Params = typename PrologueCmct::Params;
29	-
30	- __aicore__ inline BlockPrologueNN() = default;
31	- __aicore__ inline BlockPrologueNN(const Params& params) : PrologueCmct(params) {}
32	-
33	- template <class ProblemShape>
34	- __aicore__ inline static Params ToUnderlyingArguments(
35	- ProblemShape const& problemShape, Arguments const& args,
36	- const qbmmv4_tiling::QuantBatchMatmulV4TilingDataParams* tiling)
37	- {
38	- auto stepKa = static_cast<uint32_t>(tiling->matmulTiling.stepKa);
39	- auto stepKb = static_cast<uint32_t>(tiling->matmulTiling.stepKb);
40	- auto baseK = static_cast<uint32_t>(tiling->matmulTiling.baseK);
41	- auto baseM = static_cast<uint32_t>(tiling->matmulTiling.baseM);
42	- auto baseN = static_cast<uint32_t>(tiling->matmulTiling.baseN);
43	- return {.ptrB = args.ptrB,
44	- .ptrBias = args.ptrBias,
45	- .tileShapeL1 = AscendC::MakeShape(baseM, baseN, stepKa * baseK, stepKb * baseK),
46	- .layoutB = args.layoutB,
47	- .layoutBias = args.layoutBias,
48	- .l1BufNum = tiling->BL1Pingpong,
49	- .nUbSize = static_cast<int32_t>(tiling->nBubSize),
50	- .kUbSize = static_cast<int32_t>(tiling->kBubSize),
51	- .hasBias = bool(tiling->matmulTiling.isBias)};
52	- }
53	-};
54	-} // namespace Prologue
55	-} // namespace QuantBatchMatmulV4

Dmatmul/quant_batch_matmul_v4/op_kernel/arch35/cmct_convertor.h+0-163

	@@ -1,163 +0,0 @@
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 cmct_convertor.h
13	- * \brief
14	- */
15	-
16	-#pragma once
17	-#include "cmct/block/block_mmad_mx_weight_from_ub.h"
18	-#include "cmct/block/block_scheduler_swizzle_in_mn_core_nn.h"
19	-#include "cmct/kernel/kernel_matmul_mix_with_weight_prologue_nn.h"
20	-#include "cmct/prologue/block_prologue_b_cast_scsc_nn.h"
21	-#include "cmct/policy/dispatch_policy.h"
22	-#include "cmct/utils/gemm_type.h"
23	-#include "cmct/utils/integral_constant.h"
24	-#if ASC_DEVKIT_MAJOR >= 9
25	-#include "kernel_basic_intf.h"
26	-#else
27	-#include "kernel_operator.h"
28	-#endif
29	-#include "lib/std/type_traits.h"
30	-#include "quant_batch_matmul_v4_tiling_data_apt.h"
31	-
32	-namespace QuantBatchMatmulV4 {
33	-using AscendC::fp8_e8m0_t;
34	-using ProblemShape = AscendC::Std::tuple<uint64_t, uint64_t, uint64_t>; // m, n, k
35	-using TileShapeL1 = AscendC::Std::tuple<uint32_t, uint32_t, uint32_t, uint32_t>; // m, n, ka, kb
36	-using TileShapeL0 = AscendC::Std::tuple<uint32_t, uint32_t, uint32_t>; // m, n, k
37	-using LayoutA = AscendC::Layout<AscendC::Std::tuple<uint64_t, uint64_t>, AscendC::Std::tuple<uint64_t, Cmct::Gemm::_1>>;
38	-using LayoutC = AscendC::Layout<AscendC::Std::tuple<uint64_t, uint64_t>, AscendC::Std::tuple<uint64_t, Cmct::Gemm::_1>>;
39	-using LayoutBias = AscendC::Layout<AscendC::Std::tuple<uint64_t>, AscendC::Std::tuple<Cmct::Gemm::_1>>;
40	-using LayoutScale = AscendC::Layout<AscendC::Std::tuple<uint64_t, uint64_t>,
41	- AscendC::Std::tuple<uint64_t, Cmct::Gemm::_1>>;
42	-using AType = Cmct::Gemm::GemmType<DTYPE_X1, LayoutA>;
43	-using CType = Cmct::Gemm::GemmType<DTYPE_Y, LayoutC>;
44	-using BiasType = Cmct::Gemm::GemmType<DTYPE_Y, LayoutBias>;
45	-using ScaleType = Cmct::Gemm::GemmType<fp8_e8m0_t, LayoutScale>;
46	-
47	-// 不要使用 AscendC namespace下的CeilDiv和CeilAlign函数!
48	-using Cmct::CeilAlign;
49	-using Cmct::CeilDiv;
50	-
51	-constexpr uint64_t MX_GROUP_SIZE = 32UL;
52	-constexpr uint64_t MX_K_ALIGN_SIZE = 64UL;
53	-
54	-template <bool weightNz>
55	-struct StrideWeight {
56	- static_assert(AscendC::Std::always_false_v<decltype(weightNz)>,
57	- "StrideWeight should be specialized by values (true or false)");
58	-};
59	-
60	-template <>
61	-struct StrideWeight<true> {
62	- using type = AscendC::Std::tuple<AscendC::Std::tuple<Cmct::Gemm::_32, Cmct::Gemm::_512>,
63	- AscendC::Std::tuple<Cmct::Gemm::_1, uint64_t>>;
64	-};
65	-
66	-template <>
67	-struct StrideWeight<false> {
68	- using type = AscendC::Std::tuple<uint64_t, Cmct::Gemm::_1>;
69	-};
70	-
71	-template <bool weightNz>
72	-struct ShapeWeight {
73	- static_assert(AscendC::Std::always_false_v<decltype(weightNz)>,
74	- "ShapeWeight should be specialized by values (true or false)");
75	-};
76	-
77	-template <>
78	-struct ShapeWeight<true> {
79	- using type = AscendC::Std::tuple<AscendC::Std::tuple<Cmct::Gemm::_16, uint64_t>,
80	- AscendC::Std::tuple<Cmct::Gemm::_32, uint64_t>>;
81	-};
82	-
83	-template <>
84	-struct ShapeWeight<false> {
85	- using type = AscendC::Std::tuple<uint64_t, uint64_t>;
86	-};
87	-
88	-template <bool isNz>
89	-struct CreateLayoutB {};
90	-
91	-template <>
92	-struct CreateLayoutB<false> {
93	- __aicore__ inline decltype(auto) operator()(uint64_t n, uint64_t k)
94	- {
95	- return AscendC::MakeLayout(AscendC::MakeShape(n, k), AscendC::MakeStride(k, Cmct::Gemm::_1{}));
96	- }
97	-};
98	-
99	-template <>
100	-struct CreateLayoutB<true> {
101	- __aicore__ inline decltype(auto) operator()(uint64_t n, uint64_t k)
102	- {
103	- return AscendC::MakeLayout(
104	- AscendC::MakeShape(AscendC::MakeShape(Cmct::Gemm::_16{},
105	- static_cast<uint64_t>(Cmct::CeilDiv<uint64_t>(n, Cmct::Gemm::_16{}))),
106	- AscendC::MakeShape(Cmct::Gemm::_32{}, static_cast<uint64_t>(Cmct::CeilDiv<uint64_t>(
107	- k, Cmct::Gemm::_32{})))),
108	- AscendC::MakeStride(AscendC::MakeStride(Cmct::Gemm::_32{}, Cmct::Gemm::_512{}),
109	- AscendC::MakeStride(Cmct::Gemm::_1{}, Cmct::CeilAlign<uint64_t>(n, Cmct::Gemm::_16{}) *
110	- Cmct::Gemm::_32{})));
111	- }
112	-};
113	-
114	-template <bool IS_WEIGHT_NZ>
115	-__aicore__ inline void InvokeKernel(GM_ADDR x1, GM_ADDR x2, GM_ADDR bias, GM_ADDR x1_scale, GM_ADDR x2_scale,
116	- [[maybe_unused]] GM_ADDR y_scale, [[maybe_unused]] GM_ADDR x1_offset,
117	- [[maybe_unused]] GM_ADDR x2_offset, [[maybe_unused]] GM_ADDR y_offset,
118	- [[maybe_unused]] GM_ADDR x2_table, GM_ADDR y, [[maybe_unused]] GM_ADDR workspace,
119	- const GM_ADDR tiling)
120	-{
121	- KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIC_1_2);
122	- GET_TILING_DATA_WITH_STRUCT(qbmmv4_tiling::QuantBatchMatmulV4TilingDataParams, tilingDataIn, tiling);
123	- using LayoutB = AscendC::Layout<typename ShapeWeight<IS_WEIGHT_NZ>::type,
124	- typename StrideWeight<IS_WEIGHT_NZ>::type>;
125	- using BType = Cmct::Gemm::GemmType<DTYPE_X2, LayoutB>;
126	- using DispatchPolicy = Cmct::Gemm::UbAntiquantWithScSc;
127	- using BlockMmad = Block::BlockMmadNN<DispatchPolicy, TileShapeL1, TileShapeL0,
128	- AscendC::Std::tuple<AType, ScaleType>, BType, CType, BiasType, void, void>;
129	- using BlockPrologue = Prologue::BlockPrologueNN<Cmct::Prologue::BCastScsc, BType, AType, BiasType, TileShapeL1>;
130	- using BlockScheduler = Block::BlockSchedulerSwizzleInMnCoreNN<ProblemShape, AscendC::Std::tuple<uint32_t, uint32_t>,
131	- AscendC::Std::tuple<uint8_t, uint8_t>>;
132	- using KernelMmad = Kernel::KernelMatmulMixWeightPrologueNN<ProblemShape, BlockMmad, BlockScheduler, BlockPrologue>;
133	- auto problemShape = AscendC::MakeShape(tilingDataIn.mSize, tilingDataIn.nSize, tilingDataIn.kSize);
134	- uint64_t kAlign = Cmct::CeilAlign<uint64_t>(tilingDataIn.kSize, MX_K_ALIGN_SIZE);
135	- typename BlockMmad::Arguments mmad{
136	- .ptrA = x1,
137	- .ptrC = y,
138	- .ptrAScale = x1_scale,
139	- .ptrBScale = x2_scale,
140	- .layoutA = AscendC::MakeLayout(AscendC::MakeShape(tilingDataIn.mSize, tilingDataIn.kSize),
141	- AscendC::MakeStride(tilingDataIn.kSize, Cmct::Gemm::_1{})),
142	- .layoutC = AscendC::MakeLayout(AscendC::MakeShape(tilingDataIn.mSize, tilingDataIn.nSize),
143	- AscendC::MakeStride(tilingDataIn.nSize, Cmct::Gemm::_1{})),
144	- .layoutScale = AscendC::MakeLayout(
145	- AscendC::MakeShape(tilingDataIn.nSize, Cmct::CeilDiv<uint64_t>(kAlign, MX_GROUP_SIZE)),
146	- AscendC::MakeStride(Cmct::CeilDiv<uint64_t>(kAlign, MX_GROUP_SIZE), Cmct::Gemm::_1{}))};
147	- typename BlockPrologue::Arguments prologue{
148	- .ptrB = x2,
149	- .ptrBias = bias,
150	- .layoutB = CreateLayoutB<IS_WEIGHT_NZ>{}(tilingDataIn.nSize, tilingDataIn.kSize),
151	- .layoutBias = AscendC::MakeLayout(AscendC::MakeShape(tilingDataIn.nSize),
152	- AscendC::MakeStride(Cmct::Gemm::_1{}))};
153	- typename BlockScheduler::Arguments scheduler{};
154	- typename KernelMmad::Arguments args{
155	- .problemShape = problemShape, .mmad = mmad, .prologue = prologue, .scheduler = scheduler};
156	- auto params = KernelMmad::ToUnderlyingArguments(args, &tilingDataIn);
157	- KernelMmad op;
158	- op(params);
159	-}
160	-} // namespace QuantBatchMatmulV4
161	-
162	-#define KERNEL_PARAMS \
163	- x1, x2, bias, x1_scale, x2_scale, y_scale, x1_offset, x2_offset, y_offset, x2_table, y, workspace, tiling

Mmatmul/quant_batch_matmul_v4/op_kernel/arch35/quant_batch_matmul_v4.cpp+5-4

	@@ -71,9 +71,9 @@
71	#include "../quant_batch_matmul_v4_constant.h"	71	#include "../quant_batch_matmul_v4_constant.h"
72	#include "quant_batch_matmul_v4_pertoken_pergroup.h"	72	#include "quant_batch_matmul_v4_pertoken_pergroup.h"
73	#else	73	#else
74	-#include "cmct_convertor.h"
75	#include "quant_batch_matmul_v4_constant.h"	74	#include "quant_batch_matmul_v4_constant.h"
76	#include "quant_batch_matmul_v4_perchannel.h"	75	#include "quant_batch_matmul_v4_perchannel.h"
		76	+#include "quant_batch_matmul_v4_weight_quant_mx_blaze.h"
77	#include "../../quant_batch_matmul_v3/arch35/qbmm_mix_pertile_cmct.h"	77	#include "../../quant_batch_matmul_v3/arch35/qbmm_mix_pertile_cmct.h"
78	#endif	78	#endif
79	#else	79	#else
	@@ -177,15 +177,16 @@ __global__ __aicore__ void quant_batch_matmul_v4(GM_ADDR x1, GM_ADDR x2, GM_ADDR
177	}	177	}
178	#else	178	#else
179	REGISTER_TILING_DEFAULT(qbmmv4_tiling::QuantBatchMatmulV4TilingDataParams);	179	REGISTER_TILING_DEFAULT(qbmmv4_tiling::QuantBatchMatmulV4TilingDataParams);
180	- if (QUANT_TYPE == QBMMV4_PER_GROUP) {	180	+ if constexpr (QUANT_TYPE == QBMMV4_PER_GROUP) {
181	constexpr bool isTransA = TRANS == QBMMV4_A_TRANS \|\| TRANS == QBMMV4_ALL_TRANS;	181	constexpr bool isTransA = TRANS == QBMMV4_A_TRANS \|\| TRANS == QBMMV4_ALL_TRANS;
182	constexpr bool isTransB = TRANS == QBMMV4_B_TRANS \|\| TRANS == QBMMV4_ALL_TRANS;	182	constexpr bool isTransB = TRANS == QBMMV4_B_TRANS \|\| TRANS == QBMMV4_ALL_TRANS;
183	QuantBatchMatmulV4::Arch35::InvokeWeightQuantBmmOpImpl<	183	QuantBatchMatmulV4::Arch35::InvokeWeightQuantBmmOpImpl<
184	QuantBatchMatmulV4PerChannelKernel<DTYPE_X1, DTYPE_X2, DTYPE_BIAS, DTYPE_Y, isTransA, isTransB, false,	184	QuantBatchMatmulV4PerChannelKernel<DTYPE_X1, DTYPE_X2, DTYPE_BIAS, DTYPE_Y, isTransA, isTransB, false,
185	QuantType::PER_GROUP, DTYPE_Y, WEIGHTNZ> >(	185	QuantType::PER_GROUP, DTYPE_Y, WEIGHTNZ> >(
186	x1, x2, bias, x1_scale, x2_scale, y_scale, x1_offset, x2_offset, y_offset, y, workspace, tiling);	186	x1, x2, bias, x1_scale, x2_scale, y_scale, x1_offset, x2_offset, y_offset, y, workspace, tiling);
187	- } else if (QUANT_TYPE == QBMMV4_MX) {	187	+ } else if constexpr (QUANT_TYPE == QBMMV4_MX) {
188	- QuantBatchMatmulV4::InvokeKernel<WEIGHTNZ>(KERNEL_PARAMS);	188	+ QuantBatchMatmulV4::Arch35::InvokeWeightQuantMxBlazeSwat<WEIGHTNZ>(
		189	+ x1, x2, bias, x1_scale, x2_scale, y_scale, x1_offset, x2_offset, y_offset, x2_table, y, workspace, tiling);
189	}	190	}
190	#endif	191	#endif
191	#endif	192	#endif

Mmatmul/quant_batch_matmul_v4/op_kernel/arch35/quant_batch_matmul_v4_tiling_data_apt.h+37-1

	@@ -49,7 +49,7 @@ struct QuantBatchMatmulV3DataParams {
49	uint32_t needUbBuffer = 0;	49	uint32_t needUbBuffer = 0;
50	uint32_t realSingleCoreM = 0;	50	uint32_t realSingleCoreM = 0;
51	uint32_t realSingleCoreN = 0;	51	uint32_t realSingleCoreN = 0;
52	- uint32_t biasDtype = 0; //代替原来的isBiasBf16	52	+ uint32_t biasDtype = 0; // 代替原来的isBiasBf16
53	uint32_t ubSize = 0;	53	uint32_t ubSize = 0;
54	uint32_t isMClash = 0;	54	uint32_t isMClash = 0;
55	uint32_t isNClash = 0;	55	uint32_t isNClash = 0;
	@@ -172,4 +172,40 @@ struct QuantBatchMatmulV4TilingDataParams {
172	TCubeTiling matmulTiling;	172	TCubeTiling matmulTiling;
173	};	173	};
174	#pragma pack(pop)	174	#pragma pack(pop)
		175	+
		176	+#pragma pack(push, 8)
		177	+struct alignas(8) QuantBatchMatmulV4WeightQuantMxSwatTilingData {
		178	+ uint32_t m = 0;
		179	+ uint32_t n = 0;
		180	+ uint32_t k = 0;
		181	+
		182	+ uint32_t baseM = 0;
		183	+ uint32_t baseN = 0;
		184	+ uint32_t baseK = 0;
		185	+ uint32_t tileShapeKL1 = 0;
		186	+ uint32_t tileShapeScaleKL1 = 0;
		187	+
		188	+ uint32_t usedCoreNum = 1;
		189	+ uint32_t cubeNumBlocksM = 1;
		190	+ uint32_t cubeNumBlocksN = 1;
		191	+ uint32_t iterateOrder = 1;
		192	+ uint32_t mTailTile = 1;
		193	+ uint32_t nTailTile = 1;
		194	+ uint32_t mBaseTailSplitCnt = 1;
		195	+ uint32_t nBaseTailSplitCnt = 1;
		196	+ uint32_t mTailMain = 0;
		197	+ uint32_t nTailMain = 0;
		198	+
		199	+ uint32_t nBubSize = 0;
		200	+ uint32_t kBubSize = 0;
		201	+ uint32_t groupSize = 0;
		202	+ uint32_t hasBias = 0;
		203	+ uint32_t hasX1Scale = 0;
		204	+ uint32_t hasX2Scale = 0;
		205	+ uint32_t weightNz = 0;
		206	+ uint32_t yDtype = 0;
		207	+ uint32_t l1BufferNum = 0;
		208	+ uint32_t reserved = 0;
		209	+};
		210	+#pragma pack(pop)
175	} // namespace qbmmv4_tiling	211	} // namespace qbmmv4_tiling

Mmatmul/quant_batch_matmul_v4/op_kernel/arch35/quant_batch_matmul_v4_tiling_key.h+1-1

	@@ -81,7 +81,7 @@ ASCENDC_TPL_SEL(
81	ASCENDC_TPL_UINT_SEL(OPTION_ATTRS, ASCENDC_TPL_UI_LIST, QBMMV4_OPTION_ATTR_NONE),	81	ASCENDC_TPL_UINT_SEL(OPTION_ATTRS, ASCENDC_TPL_UI_LIST, QBMMV4_OPTION_ATTR_NONE),
82	ASCENDC_TPL_UINT_SEL(WEIGHTNZ, ASCENDC_TPL_UI_LIST, QBMMV4_NOT_WEIGHT_NZ, QBMMV4_IS_WEIGHT_NZ),	82	ASCENDC_TPL_UINT_SEL(WEIGHTNZ, ASCENDC_TPL_UI_LIST, QBMMV4_NOT_WEIGHT_NZ, QBMMV4_IS_WEIGHT_NZ),
83	ASCENDC_TPL_UINT_SEL(KERNEL_TEMPLATE_TYPE, ASCENDC_TPL_UI_LIST, QBMMV4_BASIS),	83	ASCENDC_TPL_UINT_SEL(KERNEL_TEMPLATE_TYPE, ASCENDC_TPL_UI_LIST, QBMMV4_BASIS),
84	- ASCENDC_TPL_TILING_STRUCT_SEL(qbmmv4_tiling::QuantBatchMatmulV4TilingDataParams), ),	84	+ ASCENDC_TPL_TILING_STRUCT_SEL(qbmmv4_tiling::QuantBatchMatmulV4WeightQuantMxSwatTilingData), ),
85	ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),	85	ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),
86	ASCENDC_TPL_UINT_SEL(TRANS, ASCENDC_TPL_UI_LIST, QBMMV4_NOT_TRANS, QBMMV4_B_TRANS,	86	ASCENDC_TPL_UINT_SEL(TRANS, ASCENDC_TPL_UI_LIST, QBMMV4_NOT_TRANS, QBMMV4_B_TRANS,
87	QBMMV4_A_TRANS, QBMMV4_ALL_TRANS),	87	QBMMV4_A_TRANS, QBMMV4_ALL_TRANS),

Amatmul/quant_batch_matmul_v4/op_kernel/arch35/quant_batch_matmul_v4_weight_quant_mx_blaze.h+87-0

@@ -0,0 +1,87 @@
	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 quant_batch_matmul_v4_weight_quant_mx_blaze.h
	13	+ * \brief
	14	+ */
	15	+
	16	+#pragma once
	17	+
	18	+#if ASC_DEVKIT_MAJOR >= 9
	19	+#include "kernel_basic_intf.h"
	20	+#else
	21	+#include "kernel_operator.h"
	22	+#endif
	23	+
	24	+#include "blaze/gemm/kernel/kernel_matmul_mix_weight_prologue.h"
	25	+#include "quant_batch_matmul_v4_tiling_data_apt.h"
	26	+
	27	+namespace QuantBatchMatmulV4 {
	28	+namespace Arch35 {
	29	+
	30	+template <bool IS_WEIGHT_NZ>
	31	+__aicore__ inline void RunWeightQuantMxBlazeSwat(
	32	+ GM_ADDR x1, GM_ADDR x2, GM_ADDR bias, GM_ADDR x1Scale, GM_ADDR x2Scale, GM_ADDR y,
	33	+ const qbmmv4_tiling::QuantBatchMatmulV4WeightQuantMxSwatTilingData& tilingData)
	34	+{
	35	+ using AType = DTYPE_X1;
	36	+ using BType = DTYPE_X2;
	37	+ using ScaleAType = AscendC::fp8_e8m0_t;
	38	+ using ScaleBType = AscendC::fp8_e8m0_t;
	39	+ using CType = DTYPE_Y;
	40	+ using BiasType = DTYPE_BIAS;
	41	+
	42	+ using LayoutA = AscendC::Te::NDExtLayoutPtn;
	43	+ using LayoutB = AscendC::Std::conditional_t<IS_WEIGHT_NZ, AscendC::Te::ZNLayoutPtn, AscendC::Te::DNExtLayoutPtn>;
	44	+ using LayoutC = AscendC::Te::NDExtLayoutPtn;
	45	+ using LayoutScaleA = AscendC::Te::ScaleANDLayoutPtn;
	46	+ using LayoutScaleB = AscendC::Te::ScaleBDNLayoutPtn;
	47	+ using LayoutBias = LayoutC;
	48	+
	49	+ using ProblemShape = AscendC::Te::Shape<int64_t, int64_t, int64_t>;
	50	+ using DispatchPolicy = Blaze::Gemm::MatmulWithWeightQuantMx;
	51	+ using BlockMmad = Blaze::Gemm::Block::BlockMmad<
	52	+ DispatchPolicy, AscendC::Std::tuple<AType, ScaleAType>, AscendC::Std::tuple<LayoutA, LayoutScaleA>,
	53	+ AscendC::Std::tuple<BType, ScaleBType>, AscendC::Std::tuple<LayoutB, LayoutScaleB>, CType, LayoutC, BiasType,
	54	+ LayoutBias>;
	55	+ using BlockScheduler = Blaze::Gemm::Block::BlockSchedulerMatmulSwatWithTailSplit<ProblemShape>;
	56	+ using KernelImpl = Blaze::Gemm::Kernel::GemmUniversal<ProblemShape, BlockMmad, void, BlockScheduler>;
	57	+
	58	+ typename KernelImpl::Params params{
	59	+ AscendC::Te::MakeShape(static_cast<int64_t>(tilingData.m), static_cast<int64_t>(tilingData.n),
	60	+ static_cast<int64_t>(tilingData.k)),
	61	+ {x1, x1Scale, x2Scale, y,
	62	+ AscendC::Te::MakeShape(static_cast<int64_t>(tilingData.baseM), static_cast<int64_t>(tilingData.baseN),
	63	+ static_cast<int64_t>(tilingData.tileShapeKL1),
	64	+ static_cast<int64_t>(tilingData.tileShapeScaleKL1)),
	65	+ AscendC::Te::MakeShape(static_cast<int64_t>(tilingData.baseM), static_cast<int64_t>(tilingData.baseN),
	66	+ static_cast<int64_t>(tilingData.baseK)),
	67	+ tilingData.l1BufferNum, tilingData.hasBias != 0U},
	68	+ {x2, bias, tilingData.kBubSize, tilingData.nBubSize},
	69	+ {tilingData.baseM, tilingData.baseN, tilingData.mTailTile, tilingData.nTailTile, tilingData.mBaseTailSplitCnt,
	70	+ tilingData.nBaseTailSplitCnt, tilingData.mTailMain, tilingData.nTailMain}};
	71	+ KernelImpl kernel;
	72	+ kernel(params);
	73	+}
	74	+
	75	+template <bool IS_WEIGHT_NZ>
	76	+__aicore__ inline void InvokeWeightQuantMxBlazeSwat(
	77	+ GM_ADDR x1, GM_ADDR x2, GM_ADDR bias, GM_ADDR x1_scale, GM_ADDR x2_scale, [[maybe_unused]] GM_ADDR y_scale,
	78	+ [[maybe_unused]] GM_ADDR x1_offset, [[maybe_unused]] GM_ADDR x2_offset, [[maybe_unused]] GM_ADDR y_offset,
	79	+ [[maybe_unused]] GM_ADDR x2_table, GM_ADDR y, [[maybe_unused]] GM_ADDR workspace, const GM_ADDR tiling)
	80	+{
	81	+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIC_1_2);
	82	+ GET_TILING_DATA_WITH_STRUCT(qbmmv4_tiling::QuantBatchMatmulV4WeightQuantMxSwatTilingData, tilingDataIn, tiling);
	83	+ RunWeightQuantMxBlazeSwat<IS_WEIGHT_NZ>(x1, x2, bias, x1_scale, x2_scale, y, tilingDataIn);
	84	+}
	85	+
	86	+} // namespace Arch35
	87	+} // namespace QuantBatchMatmulV4

Mmatmul/quant_batch_matmul_v4/tests/ut/op_host/test_quant_batch_matmul_v4_tiling.cpp+30-16

	@@ -50,6 +50,10 @@ using namespace ge;
50	using namespace ut_util;	50	using namespace ut_util;
51	using namespace optiling;	51	using namespace optiling;
52		52
		53	+namespace {
		54	+constexpr int64_t PACKED_FP4_PER_FLOAT32 = 8;
		55	+}
		56	+
53	static std::vector<QuantBatchMatmulV4TilingTestParam> GetParams()	57	static std::vector<QuantBatchMatmulV4TilingTestParam> GetParams()
54	{	58	{
55	std::vector<QuantBatchMatmulV4TilingTestParam> params;	59	std::vector<QuantBatchMatmulV4TilingTestParam> params;
	@@ -102,11 +106,7 @@ static void TestOneParamCase(const QuantBatchMatmulV4TilingTestParam& param)
102	string socVersion = testParam[idx++];	106	string socVersion = testParam[idx++];
103	int64_t m = stol(testParam[idx++]);	107	int64_t m = stol(testParam[idx++]);
104	int64_t k = stol(testParam[idx++]);	108	int64_t k = stol(testParam[idx++]);
105	- int64_t k0 = 16;
106	- int64_t k1 = ops::CeilDiv(k, k0);
107	int64_t n = stol(testParam[idx++]);	109	int64_t n = stol(testParam[idx++]);
108	- int64_t n0 = 32;
109	- int64_t n1 = ops::CeilDiv(n, n0);
110	int64_t transA = stol(testParam[idx++]);	110	int64_t transA = stol(testParam[idx++]);
111	int64_t transB = stol(testParam[idx++]);	111	int64_t transB = stol(testParam[idx++]);
112	int64_t group = stol(testParam[idx++]);	112	int64_t group = stol(testParam[idx++]);
	@@ -114,12 +114,6 @@ static void TestOneParamCase(const QuantBatchMatmulV4TilingTestParam& param)
114	ge::Format x2Format = ParseFormat(testParam[idx++]);	114	ge::Format x2Format = ParseFormat(testParam[idx++]);
115	ge::DataType x1Dtype = ParseDtype(testParam[idx++]);	115	ge::DataType x1Dtype = ParseDtype(testParam[idx++]);
116	ge::DataType x2Dtype = ParseDtype(testParam[idx++]);	116	ge::DataType x2Dtype = ParseDtype(testParam[idx++]);
117	- if (transB) {
118	- k0 = 32;
119	- k1 = ops::CeilDiv(k, k0);
120	- n0 = 16;
121	- n1 = ops::CeilDiv(n, n0);
122	- }
123	bool hasBias = true;	117	bool hasBias = true;
124	ge::DataType biasDtype = ge::DT_FLOAT;	118	ge::DataType biasDtype = ge::DT_FLOAT;
125	string biasDtypeStr = testParam[idx++];	119	string biasDtypeStr = testParam[idx++];
	@@ -171,6 +165,25 @@ static void TestOneParamCase(const QuantBatchMatmulV4TilingTestParam& param)
171	ge::DataType yDtype = ParseDtype(testParam[idx++]);	165	ge::DataType yDtype = ParseDtype(testParam[idx++]);
172	uint32_t aicNum = stoul(testParam[idx++]);	166	uint32_t aicNum = stoul(testParam[idx++]);
173	uint32_t aivNum = stoul(testParam[idx++]);	167	uint32_t aivNum = stoul(testParam[idx++]);
		168	+ int64_t weightK = k;
		169	+ int64_t weightN = n;
		170	+ if (x2Dtype == ge::DT_FLOAT) {
		171	+ ASSERT_EQ((transB ? weightK : weightN) % PACKED_FP4_PER_FLOAT32, 0)
		172	+ << "packed FP4 logical axis must be divisible by " << PACKED_FP4_PER_FLOAT32
		173	+ << ", caseName: " << param.caseName;
		174	+ if (transB) {
		175	+ weightK /= PACKED_FP4_PER_FLOAT32;
		176	+ } else {
		177	+ weightN /= PACKED_FP4_PER_FLOAT32;
		178	+ }
		179	+ }
		180	+ if (testParam.size() > idx) {
		181	+ weightK = stol(testParam[idx]);
		182	+ }
		183	+ int64_t k0 = transB ? (x2Dtype == ge::DT_FLOAT ? 4 : 32) : 16;
		184	+ int64_t n0 = transB ? 16 : (x2Dtype == ge::DT_FLOAT ? 4 : 32);
		185	+ int64_t k1 = ops::CeilDiv(weightK, k0);
		186	+ int64_t n1 = ops::CeilDiv(weightN, n0);
174	string compileInfoStr = R"({	187	string compileInfoStr = R"({
175	"hardware_info": {"BT_SIZE": 1024, "load3d_constraints": "0",	188	"hardware_info": {"BT_SIZE": 1024, "load3d_constraints": "0",
176	"Intrinsic_fix_pipe_l0c2out": true, "Intrinsic_data_move_l12ub": true,	189	"Intrinsic_fix_pipe_l0c2out": true, "Intrinsic_data_move_l12ub": true,
	@@ -210,17 +223,17 @@ static void TestOneParamCase(const QuantBatchMatmulV4TilingTestParam& param)
210	x1Shape.MutableOriginShape() = x1Shape.MutableStorageShape();	223	x1Shape.MutableOriginShape() = x1Shape.MutableStorageShape();
211	if (x2Format == ge::FORMAT_ND) {	224	if (x2Format == ge::FORMAT_ND) {
212	if (transB) {	225	if (transB) {
213	- x2Shape.MutableStorageShape() = gert::Shape({n, k});	226	+ x2Shape.MutableStorageShape() = gert::Shape({weightN, weightK});
214	} else {	227	} else {
215	- x2Shape.MutableStorageShape() = gert::Shape({k, n});	228	+ x2Shape.MutableStorageShape() = gert::Shape({weightK, weightN});
216	}	229	}
217	x2Shape.MutableOriginShape() = x2Shape.MutableStorageShape();	230	x2Shape.MutableOriginShape() = x2Shape.MutableStorageShape();
218	} else if (x2Format == ge::FORMAT_FRACTAL_NZ) {	231	} else if (x2Format == ge::FORMAT_FRACTAL_NZ) {
219	if (transB) {	232	if (transB) {
220	- x2Shape.MutableOriginShape() = gert::Shape({n, k});	233	+ x2Shape.MutableOriginShape() = gert::Shape({weightN, weightK});
221	x2Shape.MutableStorageShape() = gert::Shape({k1, n1, n0, k0});	234	x2Shape.MutableStorageShape() = gert::Shape({k1, n1, n0, k0});
222	} else {	235	} else {
223	- x2Shape.MutableOriginShape() = gert::Shape({k, n});	236	+ x2Shape.MutableOriginShape() = gert::Shape({weightK, weightN});
224	x2Shape.MutableStorageShape() = gert::Shape({n1, k1, k0, n0});	237	x2Shape.MutableStorageShape() = gert::Shape({n1, k1, k0, n0});
225	}	238	}
226	}	239	}
	@@ -237,12 +250,13 @@ static void TestOneParamCase(const QuantBatchMatmulV4TilingTestParam& param)
237	groupM = static_cast<int64_t>((group & 0xFFFF00000000) >> 32); // 32-47bit group_m	250	groupM = static_cast<int64_t>((group & 0xFFFF00000000) >> 32); // 32-47bit group_m
238	int64_t groupNum = (k + group - 1) / group;	251	int64_t groupNum = (k + group - 1) / group;
239	if (!hasX2Table) {	252	if (!hasX2Table) {
		253	+ int64_t mxScaleGroupNum = (k + 63) / 64; // ceil(k/64) for MX E8M0 scale
240	x1ScaleShape.MutableStorageShape() = x1ScaleDtype == ge::DT_FLOAT8_E8M0 ?	254	x1ScaleShape.MutableStorageShape() = x1ScaleDtype == ge::DT_FLOAT8_E8M0 ?
241	- gert::Shape({m, groupNum / 2, 2}) :	255	+ gert::Shape({m, mxScaleGroupNum, 2}) :
242	gert::Shape({m, groupNum});	256	gert::Shape({m, groupNum});
243	if (transB) {	257	if (transB) {
244	x2ScaleShape.MutableStorageShape() = x2ScaleDtype == ge::DT_FLOAT8_E8M0 ?	258	x2ScaleShape.MutableStorageShape() = x2ScaleDtype == ge::DT_FLOAT8_E8M0 ?
245	- gert::Shape({n, groupNum / 2, 2}) :	259	+ gert::Shape({n, mxScaleGroupNum, 2}) :
246	gert::Shape({n, groupNum});	260	gert::Shape({n, groupNum});
247	} else {	261	} else {
248	x2ScaleShape.MutableStorageShape() = gert::Shape({groupNum, n});	262	x2ScaleShape.MutableStorageShape() = gert::Shape({groupNum, n});

Mmatmul/quant_batch_matmul_v4/tests/ut/op_host/test_quant_batch_matmul_v4_tiling.csv+46-3

	@@ -1,16 +1,59 @@
1	caseName,numBlocks,tilingResult,tilingKey	1	caseName,numBlocks,tilingResult,tilingKey
2	mx-1_Ascend950_128_512_128_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_UINT64_NULL_BF16_32_64,32,GRAPH_SUCCESS,17	2	mx-1_Ascend950_128_512_128_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_UINT64_NULL_BF16_32_64,32,GRAPH_SUCCESS,17
		3	+mx-2_Ascend950_81_24_32_0_1_32_ND_NZ_FP8-E4M3_FP4-E2M1_NULL_FP8-E8M0_FP8-E8M0_UINT64_NULL_BF16_32_64,12,GRAPH_SUCCESS,273
		4	+mx-2_Ascend950_81_24_32_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_NULL_FP8-E8M0_FP8-E8M0_UINT64_NULL_BF16_32_64,12,GRAPH_SUCCESS,17
3	mx-menkan40_Ascend950_944_7680_256_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_UINT64_NULL_BF16_32_64,30,GRAPH_SUCCESS,17	5	mx-menkan40_Ascend950_944_7680_256_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_UINT64_NULL_BF16_32_64,30,GRAPH_SUCCESS,17
4	mx-menkan18_Ascend950_736_1536_2800_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_UINT64_NULL_BF16_32_64,32,GRAPH_SUCCESS,17	6	mx-menkan18_Ascend950_736_1536_2800_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_UINT64_NULL_BF16_32_64,32,GRAPH_SUCCESS,17
5	mx-menkan17_Ascend950_320_1536_224_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_UINT64_NULL_BF16_32_64,28,GRAPH_SUCCESS,17	7	mx-menkan17_Ascend950_320_1536_224_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_UINT64_NULL_BF16_32_64,28,GRAPH_SUCCESS,17
6	-mx-menkan12_Ascend950_48_7680_80_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_UINT64_NULL_BF16_32_64,9,GRAPH_SUCCESS,17	8	+mx-menkan12_Ascend950_48_7680_80_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_UINT64_NULL_BF16_32_64,15,GRAPH_SUCCESS,17
7	mx-random0001_Ascend950_608_1024_704_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_UINT64_NULL_BF16_32_64,30,GRAPH_SUCCESS,17	9	mx-random0001_Ascend950_608_1024_704_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_UINT64_NULL_BF16_32_64,30,GRAPH_SUCCESS,17
8	mx-random0003_Ascend950_3840_512_3200_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_UINT64_NULL_BF16_32_64,32,GRAPH_SUCCESS,17	10	mx-random0003_Ascend950_3840_512_3200_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_UINT64_NULL_BF16_32_64,32,GRAPH_SUCCESS,17
9	mx-random0013_Ascend950_2256_512_192_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_UINT64_NULL_BF16_32_64,32,GRAPH_SUCCESS,17	11	mx-random0013_Ascend950_2256_512_192_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_UINT64_NULL_BF16_32_64,32,GRAPH_SUCCESS,17
10	-mx-random0022_Ascend950_32_3072_64_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_UINT64_NULL_BF16_32_64,4,GRAPH_SUCCESS,17	12	+mx-random0022_Ascend950_32_3072_64_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_UINT64_NULL_BF16_32_64,8,GRAPH_SUCCESS,17
11	-mx-random0025_Ascend950_2720_512_192_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_UINT64_NULL_BF16_32_64,32,GRAPH_SUCCESS,17	13	+mx-random0025_Ascend950_2720_512_192_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_UINT64_NULL_BF16_32_64,29,GRAPH_SUCCESS,17
12	mx-error-x1ScaleDtype_Ascend950_2720_512_192_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_BF16_BF16_FP8-E8M0_UINT64_NULL_BF16_32_64,32,GRAPH_FAILED,17	14	mx-error-x1ScaleDtype_Ascend950_2720_512_192_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_BF16_BF16_FP8-E8M0_UINT64_NULL_BF16_32_64,32,GRAPH_FAILED,17
13	mx-error-x2ScaleDtype_Ascend950_2720_512_192_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_BF16_UINT64_NULL_BF16_32_64,32,GRAPH_FAILED,17	15	mx-error-x2ScaleDtype_Ascend950_2720_512_192_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_BF16_UINT64_NULL_BF16_32_64,32,GRAPH_FAILED,17
		16	+mx-nz-1_Ascend950_128_512_128_0_1_32_ND_NZ_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_NULL_NULL_BF16_32_64,32,GRAPH_SUCCESS,273
		17	+mx-packed-fp32-k-valid_Ascend950_128_512_128_0_1_32_ND_ND_FP8-E4M3_FP32_BF16_FP8-E8M0_FP8-E8M0_NULL_NULL_BF16_32_64,32,GRAPH_SUCCESS,17
		18	+mx-packed-fp32-k-mismatch_Ascend950_128_512_128_0_1_32_ND_ND_FP8-E4M3_FP32_BF16_FP8-E8M0_FP8-E8M0_NULL_NULL_BF16_32_64_512,0,GRAPH_FAILED,17
		19	+mx-packed-fp32-k-valid-nz_Ascend950_128_512_128_0_1_32_ND_NZ_FP8-E4M3_FP32_BF16_FP8-E8M0_FP8-E8M0_NULL_NULL_BF16_32_64,32,GRAPH_SUCCESS,273
		20	+mx-packed-fp32-k-mismatch-nz_Ascend950_128_512_128_0_1_32_ND_NZ_FP8-E4M3_FP32_BF16_FP8-E8M0_FP8-E8M0_NULL_NULL_BF16_32_64_512,0,GRAPH_FAILED,273
		21	+pergroup-packed-fp32-n-valid-nz_Ascend950_128_512_128_0_0_32_ND_NZ_FP8-E4M3_FP32_NULL_NULL_BF16_UINT64_NULL_BF16_32_64,16,GRAPH_SUCCESS,268
		22	+mx-native-fp4-k-mismatch-nz_Ascend950_128_512_128_0_1_32_ND_NZ_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_NULL_NULL_BF16_32_64_256,0,GRAPH_FAILED,273
		23	+# TUT-S: SWAT success cases
		24	+mx-s02_Ascend950_128_512_128_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_NULL_FP8-E8M0_FP8-E8M0_NULL_NULL_BF16_32_64,32,GRAPH_SUCCESS,17
		25	+mx-s03_Ascend950_128_512_128_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_FP16_FP8-E8M0_FP8-E8M0_NULL_NULL_FP16_32_64,32,GRAPH_SUCCESS,17
		26	+mx-s04_Ascend950_128_512_128_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_NULL_FP8-E8M0_FP8-E8M0_NULL_NULL_FP16_32_64,32,GRAPH_SUCCESS,17
		27	+mx-s07_Ascend950_128_512_128_0_1_32_ND_NZ_FP8-E4M3_FP4-E2M1_FP16_FP8-E8M0_FP8-E8M0_NULL_NULL_FP16_32_64,32,GRAPH_SUCCESS,273
		28	+mx-s08_Ascend950_128_512_128_0_1_32_ND_NZ_FP8-E4M3_FP4-E2M1_NULL_FP8-E8M0_FP8-E8M0_NULL_NULL_FP16_32_64,32,GRAPH_SUCCESS,273
		29	+mx-s09_Ascend950_1_8_1_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_NULL_NULL_BF16_32_64,1,GRAPH_SUCCESS,17
		30	+mx-s10_Ascend950_1_8_8_0_1_32_ND_NZ_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_NULL_NULL_BF16_32_64,1,GRAPH_SUCCESS,273
		31	+mx-s11_Ascend950_17_24_7_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_NULL_NULL_BF16_32_64,2,GRAPH_SUCCESS,17
		32	+mx-s12_Ascend950_17_40_24_0_1_32_ND_NZ_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_NULL_NULL_BF16_32_64,4,GRAPH_SUCCESS,273
		33	+mx-s15_Ascend950_128_512_128_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_NULL_NULL_BF16_1_2,1,GRAPH_SUCCESS,17
		34	+mx-long-k-boundary_Ascend950_16_65536_64_0_1_32_ND_NZ_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_NULL_NULL_BF16_32_64,4,GRAPH_SUCCESS,273
		35	+# TUT-E: SWAT exception interception — dtype errors
		36	+mx-e01_Ascend950_128_512_128_0_1_32_ND_ND_BF16_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_NULL_NULL_BF16_32_64,0,GRAPH_FAILED,17
		37	+mx-e02_Ascend950_128_512_128_0_1_32_ND_ND_FP8-E4M3_BF16_BF16_FP8-E8M0_FP8-E8M0_NULL_NULL_BF16_32_64,0,GRAPH_FAILED,17
		38	+mx-e05_Ascend950_128_512_128_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_BF16_NULL_FP8-E8M0_NULL_NULL_BF16_32_64,0,GRAPH_FAILED,17
		39	+mx-e06_Ascend950_128_512_128_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_NULL_NULL_NULL_BF16_32_64,0,GRAPH_FAILED,17
		40	+mx-e07_Ascend950_128_512_128_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_FP32_FP8-E8M0_FP8-E8M0_NULL_NULL_BF16_32_64,0,GRAPH_FAILED,17
		41	+mx-e10_Ascend950_128_512_128_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_NULL_NULL_FP32_32_64,0,GRAPH_FAILED,17
		42	+# TUT-E: SWAT exception interception — format and alignment errors
		43	+mx-e11_Ascend950_128_512_128_0_1_32_NZ_ND_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_NULL_NULL_BF16_32_64,0,GRAPH_FAILED,17
		44	+mx-e33_Ascend950_128_510_128_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_NULL_NULL_BF16_32_64,0,GRAPH_FAILED,17
		45	+mx-e34_Ascend950_128_510_128_0_1_32_ND_NZ_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_NULL_NULL_BF16_32_64,0,GRAPH_FAILED,273
		46	+mx-e35_Ascend950_128_512_124_0_1_32_ND_NZ_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_NULL_NULL_BF16_32_64,0,GRAPH_FAILED,273
		47	+# TUT-E: SWAT exception interception — attr errors
		48	+mx-e36_Ascend950_128_512_128_1_1_32_ND_ND_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_NULL_NULL_BF16_32_64,0,GRAPH_FAILED,17
		49	+mx-e37_Ascend950_128_512_128_0_0_32_ND_ND_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_NULL_NULL_BF16_32_64,0,GRAPH_FAILED,17
		50	+mx-e38_Ascend950_128_512_128_0_0_32_ND_NZ_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_NULL_NULL_BF16_32_64,0,GRAPH_FAILED,273
		51	+mx-e39_Ascend950_128_512_128_0_1_0_ND_ND_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_NULL_NULL_BF16_32_64,0,GRAPH_FAILED,17
		52	+mx-e40_Ascend950_128_512_128_0_1_64_ND_ND_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_NULL_NULL_BF16_32_64,0,GRAPH_FAILED,17
		53	+mx-e41_Ascend950_128_512_128_0_1_33_ND_ND_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_NULL_NULL_BF16_32_64,0,GRAPH_FAILED,17
		54	+# TUT-E: SWAT exception interception — core num errors
		55	+mx-e42_Ascend950_128_512_128_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_NULL_NULL_BF16_0_0,0,GRAPH_FAILED,17
		56	+mx-e43_Ascend950_128_512_128_0_1_32_ND_ND_FP8-E4M3_FP4-E2M1_BF16_FP8-E8M0_FP8-E8M0_NULL_NULL_BF16_32_32,0,GRAPH_FAILED,17
14	UT-A8W4-PerGroup-NZ-Testcase-0_Ascend950_848_640_896_0_0_32_ND_NZ_FP8-E4M3_FP4-E2M1_NULL_NULL_BF16_UINT64_NULL_BF16_32_64,30,GRAPH_SUCCESS,268	57	UT-A8W4-PerGroup-NZ-Testcase-0_Ascend950_848_640_896_0_0_32_ND_NZ_FP8-E4M3_FP4-E2M1_NULL_NULL_BF16_UINT64_NULL_BF16_32_64,30,GRAPH_SUCCESS,268
15	UT-A8W4-PerGroup-NZ-Testcase-1_Ascend950_96_8960_8384_0_0_32_ND_NZ_FP8-E4M3_FP4-E2M1_NULL_NULL_BF16_UINT64_NULL_BF16_32_64,32,GRAPH_SUCCESS,268	58	UT-A8W4-PerGroup-NZ-Testcase-1_Ascend950_96_8960_8384_0_0_32_ND_NZ_FP8-E4M3_FP4-E2M1_NULL_NULL_BF16_UINT64_NULL_BF16_32_64,32,GRAPH_SUCCESS,268
16	UT-A8W4-PerGroup-NZ-Testcase-2_Ascend950_176_64_128_0_0_32_ND_NZ_FP8-E4M3_FP4-E2M1_NULL_NULL_BF16_UINT64_NULL_BF16_32_64,22,GRAPH_SUCCESS,268	59	UT-A8W4-PerGroup-NZ-Testcase-2_Ascend950_176_64_128_0_0_32_ND_NZ_FP8-E4M3_FP4-E2M1_NULL_NULL_BF16_UINT64_NULL_BF16_32_64,22,GRAPH_SUCCESS,268

	@@ -19,6 +19,8 @@
19	#include "arch35/adaptive_sliding_window_basic_api_v4_tiling.h"	19	#include "arch35/adaptive_sliding_window_basic_api_v4_tiling.h"
20	#include "arch35/quant_batch_matmul_v4_asw_tiling.h"	20	#include "arch35/quant_batch_matmul_v4_asw_tiling.h"
21	#include "arch35/quant_batch_matmul_v4_pergroup_arch35_tiling.h"	21	#include "arch35/quant_batch_matmul_v4_pergroup_arch35_tiling.h"
		22	+#include "arch35/quant_batch_matmul_v4_reg_base_tiling.h"
		23	+#include "arch35/quant_batch_matmul_v4_weight_quant_mx_swat_tiling.h"
22	#include "quant_batch_matmul_v4_compile_info.h"	24	#include "quant_batch_matmul_v4_compile_info.h"
23	#include "error_util.h"	25	#include "error_util.h"
24	#include "platform/platform_infos_def.h"	26	#include "platform/platform_infos_def.h"
	@@ -27,19 +29,23 @@ namespace optiling {
27	using Ops::NN::Optiling::TilingRegistry;	29	using Ops::NN::Optiling::TilingRegistry;
28		30
29	constexpr int32_t BASIC_PERBLOCK_PRIORITY = 0;	31	constexpr int32_t BASIC_PERBLOCK_PRIORITY = 0;
30	-constexpr int32_t BASIC_PRIORITY = 1;	32	+constexpr int32_t REG_BASE_PRIORITY = 1;
31	constexpr int32_t MSD_PRIORITY = 2;	33	constexpr int32_t MSD_PRIORITY = 2;
32	constexpr int32_t PERBLOCK_PRIORITY = 3;	34	constexpr int32_t PERBLOCK_PRIORITY = 3;
33	constexpr int32_t PERGROUP_PRIORITY = 4;	35	constexpr int32_t PERGROUP_PRIORITY = 4;
34	constexpr int32_t LUT_PRIORITY = 5;	36	constexpr int32_t LUT_PRIORITY = 5;
35	constexpr int32_t PERGROUP_ARCH35_PRIORITY = 6;	37	constexpr int32_t PERGROUP_ARCH35_PRIORITY = 6;
		38	+constexpr int32_t WEIGHT_QUANT_MX_SWAT_PRIORITY = 7;
36		39
37	REGISTER_TILING_TEMPLATE("QuantBatchMatmulV4", AdaptiveSlidingWindowBasicTilingV4, BASIC_PERBLOCK_PRIORITY);	40	REGISTER_TILING_TEMPLATE("QuantBatchMatmulV4", AdaptiveSlidingWindowBasicTilingV4, BASIC_PERBLOCK_PRIORITY);
		41	+REGISTER_TILING_TEMPLATE("QuantBatchMatmulV4", QuantBatchMatmulV4RegBase, REG_BASE_PRIORITY);
38	REGISTER_TILING_TEMPLATE("QuantBatchMatmulV4", QuantBatchMatmulV4MsdTiling, MSD_PRIORITY);	42	REGISTER_TILING_TEMPLATE("QuantBatchMatmulV4", QuantBatchMatmulV4MsdTiling, MSD_PRIORITY);
39	REGISTER_TILING_TEMPLATE("QuantBatchMatmulV4", QuantBatchMatmulV4PerblockTiling, PERBLOCK_PRIORITY);	43	REGISTER_TILING_TEMPLATE("QuantBatchMatmulV4", QuantBatchMatmulV4PerblockTiling, PERBLOCK_PRIORITY);
40	REGISTER_TILING_TEMPLATE("QuantBatchMatmulV4", QuantBatchMatmulV4PergroupTiling, PERGROUP_PRIORITY);	44	REGISTER_TILING_TEMPLATE("QuantBatchMatmulV4", QuantBatchMatmulV4PergroupTiling, PERGROUP_PRIORITY);
41	REGISTER_TILING_TEMPLATE("QuantBatchMatmulV4", AdaptiveSlidingWindowTilingV4, LUT_PRIORITY);	45	REGISTER_TILING_TEMPLATE("QuantBatchMatmulV4", AdaptiveSlidingWindowTilingV4, LUT_PRIORITY);
42	REGISTER_TILING_TEMPLATE("QuantBatchMatmulV4", QuantBatchMatmulV4PergroupArch35Tiling, PERGROUP_ARCH35_PRIORITY);	46	REGISTER_TILING_TEMPLATE("QuantBatchMatmulV4", QuantBatchMatmulV4PergroupArch35Tiling, PERGROUP_ARCH35_PRIORITY);
		47	+REGISTER_TILING_TEMPLATE("QuantBatchMatmulV4", QuantBatchMatmulV4WeightQuantMxSwatTiling,
		48	+ WEIGHT_QUANT_MX_SWAT_PRIORITY);
43		49
44	ge::graphStatus QuantBatchMatmulV4TilingFunc(gert::TilingContext* context)	50	ge::graphStatus QuantBatchMatmulV4TilingFunc(gert::TilingContext* context)
45	{	51	{
	@@ -68,7 +74,8 @@ ge::graphStatus QuantBatchMatmulV4TilingFunc(gert::TilingContext* context)
68	vector<int32_t> regitserList = {LUT_PRIORITY};	74	vector<int32_t> regitserList = {LUT_PRIORITY};
69	return TilingRegistry::GetInstance().DoTilingImpl(context, regitserList);	75	return TilingRegistry::GetInstance().DoTilingImpl(context, regitserList);
70	}	76	}
71	- std::vector<int32_t> registerList = {BASIC_PERBLOCK_PRIORITY, PERGROUP_ARCH35_PRIORITY, optiling::BASIC_PRIORITY};	77	+ std::vector<int32_t> registerList = {BASIC_PERBLOCK_PRIORITY, PERGROUP_ARCH35_PRIORITY,
		78	+ WEIGHT_QUANT_MX_SWAT_PRIORITY, REG_BASE_PRIORITY};
72	return TilingRegistry::GetInstance().DoTilingImpl(context, registerList);	79	return TilingRegistry::GetInstance().DoTilingImpl(context, registerList);
73	}	80	}
74		81