Star871
1756
代码
代码
Issues231
Pull Requests325
Actions
项目讨论
Wiki
分析
Star871
1756
已合并
qbmm mxA8W4 supports k % 8 == 0 #5439
yuanshuai000创建于 5月29日
qbmm mxA8W4 supports k % 8 == 0 #5439
已合并
yuanshuai000创建于 5月29日
从yuanshuai000/ops-nnqbmm_mxa8w4_k_8_align合入到cann/ops-nnmaster
共 15 个文件变更+198-138
Mmatmul/common/cmct/block/block_mmad_b_prologue_mx.h+7-4
@@ -30,6 +30,7 @@
30 30 
31namespace Cmct::Gemm::Block {31namespace Cmct::Gemm::Block {
32using AscendC::BLOCK_CUBE;32using AscendC::BLOCK_CUBE;
33+using AscendC::ONE_BLK_SIZE;
33using AscendC::CrossCoreSetFlag;34using AscendC::CrossCoreSetFlag;
34using AscendC::CrossCoreWaitFlag;35using AscendC::CrossCoreWaitFlag;
35using AscendC::GlobalTensor;36using AscendC::GlobalTensor;
@@ -317,7 +318,7 @@ private:
317 WaitFlag<HardEvent::MTE1_MTE2>(eventIdsScaleBMte1ToMte2_[scaleBBufIdx_]);318 WaitFlag<HardEvent::MTE1_MTE2>(eventIdsScaleBMte1ToMte2_[scaleBBufIdx_]);
318 CopyScaleB2L1(tensorBlockScaleB);319 CopyScaleB2L1(tensorBlockScaleB);
319 }320 }
320- if (kL1Len_ % K_ALIGN_SIZE != 0) {321+ if (CeilAlign(kL1Len_, ONE_BLK_SIZE) % K_ALIGN_SIZE != 0) {
321 LocalTensor<ElementA> bL1Tensor;322 LocalTensor<ElementA> bL1Tensor;
322 if (l1BufIdx_ == IDX_0) {323 if (l1BufIdx_ == IDX_0) {
323 bL1Tensor = bL1LocalBuf0_;324 bL1Tensor = bL1LocalBuf0_;
@@ -328,7 +329,8 @@ private:
328 } else {329 } else {
329 bL1Tensor = bL1LocalBuf3_;330 bL1Tensor = bL1LocalBuf3_;
330 }331 }
331- FillL1WithZero(bL1Tensor[CeilAlign(nL1Len_, BLOCK_CUBE) * kL1Len_].template ReinterpretCast<uint32_t>(), nL1Len_);332+ uint64_t offset = CeilAlign(nL1Len_, BLOCK_CUBE) * CeilAlign(kL1Len_, ONE_BLK_SIZE);
333+ FillL1WithZero(bL1Tensor[offset].template ReinterpretCast<uint32_t>(), nL1Len_);
332 }334 }
333 WaitForVector(l1BufIdx_);335 WaitForVector(l1BufIdx_);
334 }336 }
@@ -590,8 +592,9 @@ private:
590 }592 }
591 DataCopy(aL1LocalBuf, srcTensor, nd2nzParams);593 DataCopy(aL1LocalBuf, srcTensor, nd2nzParams);
592 594 
593- if (kL1Len_ % K_ALIGN_SIZE != 0) {595+ if (CeilAlign(kL1Len_, ONE_BLK_SIZE) % K_ALIGN_SIZE != 0) {
594- FillL1WithZero(aL1LocalBuf[CeilAlign(mL1Len_, BLOCK_CUBE) * kL1Len_].template ReinterpretCast<uint32_t>(), mL1Len_);596+ uint64_t offset = CeilAlign(mL1Len_, BLOCK_CUBE) * CeilAlign(kL1Len_, ONE_BLK_SIZE);
597+ FillL1WithZero(aL1LocalBuf[offset].template ReinterpretCast<uint32_t>(), mL1Len_);
595 }598 }
596 }599 }
597 600 
Mmatmul/common/cmct/prologue/block_prologue_b_cast_scsc.h+9-6
@@ -376,7 +376,7 @@ private:
376 if constexpr (weightNz) {376 if constexpr (weightNz) {
377 int64_t nAlignSize = CeilAlign(nSize_, static_cast<uint64_t>(BLOCK_CUBE));377 int64_t nAlignSize = CeilAlign(nSize_, static_cast<uint64_t>(BLOCK_CUBE));
378 int64_t nUbAlignSize = CeilAlign(nUbLen_, BLOCK_CUBE);378 int64_t nUbAlignSize = CeilAlign(nUbLen_, BLOCK_CUBE);
379- intriParams.blockCount = kUbLen_ / C0_SIZE_B8;379+ intriParams.blockCount = CeilDiv(kUbLen_, C0_SIZE_B8);
380 intriParams.blockLen = nUbAlignSize * C0_SIZE_B8;380 intriParams.blockLen = nUbAlignSize * C0_SIZE_B8;
381 intriParams.srcStride = (nAlignSize - nUbAlignSize) * C0_SIZE_B8;381 intriParams.srcStride = (nAlignSize - nUbAlignSize) * C0_SIZE_B8;
382 } else {382 } else {
@@ -422,7 +422,9 @@ private:
422 AscendC::DataCopyParams params;422 AscendC::DataCopyParams params;
423 if constexpr (weightNz) {423 if constexpr (weightNz) {
424 params.blockLen = BLOCK_NUM_REG;424 params.blockLen = BLOCK_NUM_REG;
425- params.blockCount = CeilAlign(nUbLen_, BLOCK_CUBE) * kUbLen_ * sizeof(ElementOut) / VECTOR_REG_WIDTH;425+ params.blockCount = CeilAlign(nUbLen_, BLOCK_CUBE) *
426+ CeilAlign(kUbLen_, static_cast<int32_t>(ONE_BLK_SIZE)) *
427+ sizeof(ElementOut) / VECTOR_REG_WIDTH;
426 params.srcStride = (l1BufNum_ - 1) * BLOCK_NUM_REG;428 params.srcStride = (l1BufNum_ - 1) * BLOCK_NUM_REG;
427 params.dstStride = 0;429 params.dstStride = 0;
428 DataCopy(l1Local_[l1Offset], ubLocal, params);430 DataCopy(l1Local_[l1Offset], ubLocal, params);
@@ -466,11 +468,12 @@ private:
466 wParams.dataBlockStride = CeilAlign(nUbLen_, BLOCK_CUBE) + 1;468 wParams.dataBlockStride = CeilAlign(nUbLen_, BLOCK_CUBE) + 1;
467 wParams.repeatStride = wParams.dataBlockStride * BLOCK_CUBE;469 wParams.repeatStride = wParams.dataBlockStride * BLOCK_CUBE;
468 wParams.outDimOffset = ONE_BLOCK_SIZE - wParams.innerExtend * wParams.repeatStride * ONE_BLOCK_SIZE;470 wParams.outDimOffset = ONE_BLOCK_SIZE - wParams.innerExtend * wParams.repeatStride * ONE_BLOCK_SIZE;
469- wParams.maskB8Tail0 = Min(kUbLen_ % VECTOR_REG_WIDTH_FOR_4BITS, static_cast<int32_t>(VECTOR_REG_WIDTH)) +471+ int32_t kUbLenAlign = CeilAlign(kUbLen_, static_cast<int32_t>(ONE_BLOCK_SIZE));
470- kUbLen_ / VECTOR_REG_WIDTH_FOR_4BITS * VECTOR_REG_WIDTH;472+ wParams.maskB8Tail0 = Min(kUbLenAlign % VECTOR_REG_WIDTH_FOR_4BITS, static_cast<int32_t>(VECTOR_REG_WIDTH)) +
473+ kUbLenAlign / VECTOR_REG_WIDTH_FOR_4BITS * VECTOR_REG_WIDTH;
471 wParams.maskB8Tail1 =474 wParams.maskB8Tail1 =
472- Cmct::Gemm::Max(kUbLen_ % VECTOR_REG_WIDTH_FOR_4BITS - static_cast<int32_t>(VECTOR_REG_WIDTH), 0) +475+ Cmct::Gemm::Max(kUbLenAlign % VECTOR_REG_WIDTH_FOR_4BITS - static_cast<int32_t>(VECTOR_REG_WIDTH), 0) +
473- kUbLen_ / VECTOR_REG_WIDTH_FOR_4BITS * VECTOR_REG_WIDTH;476+ kUbLenAlign / VECTOR_REG_WIDTH_FOR_4BITS * VECTOR_REG_WIDTH;
474 wParams.weightInUbBaseAddr = weightInUbBaseAddr_;477 wParams.weightInUbBaseAddr = weightInUbBaseAddr_;
475 wParams.weightOutUbAddr = weightOutUbAddr_;478 wParams.weightOutUbAddr = weightOutUbAddr_;
476 wParams.weightOutUbAddr1 = weightOutUbAddr1_;479 wParams.weightOutUbAddr1 = weightOutUbAddr1_;
Mmatmul/quant_batch_matmul_v3/docs/aclnnQuantMatmulWeightNz.md+4-3
@@ -579,13 +579,14 @@ aclnnStatus aclnnQuantMatmulWeightNz(
579 579
580 |量化类型|x1 dtype |x2 dtype | x1Scale dtype |x2Scale dtype |bias dtype| yScale dtype | out dtype | x1 shape | x2 shape| x1Scale shape | x2Scale shape |bias shape | yScale shape| [groupSizeM, groupSizeN, groupSizeK]|580 |量化类型|x1 dtype |x2 dtype | x1Scale dtype |x2Scale dtype |bias dtype| yScale dtype | out dtype | x1 shape | x2 shape| x1Scale shape | x2Scale shape |bias shape | yScale shape| [groupSizeM, groupSizeN, groupSizeK]|
581 |---------------| ------------| -------------- |--------------|-------------|--------- | -------- |--------| ---------- | --------------| ------------ |---------- | ------------| ---------------------------------------|-------|581 |---------------| ------------| -------------- |--------------|-------------|--------- | -------- |--------| ---------- | --------------| ------------ |---------- | ------------| ---------------------------------------|-------|
582- | mx量化 |FLOAT8_E4M3FN |FLOAT4_E2M1 |FLOAT8_E8M0 |FLOAT8_E8M0 |null/BFLOAT16/FLOAT16| null |BFLOAT16/FLOAT16|(m, k) |(n, k) |(m, k/64, 2) |(n, k/64, 2) |(1, n) | null | [0, 0, 32] / [1, 1, 32] |582+ | mx量化 |FLOAT8_E4M3FN |FLOAT4_E2M1 |FLOAT8_E8M0 |FLOAT8_E8M0 |null/BFLOAT16/FLOAT16| null |BFLOAT16/FLOAT16|(m, k) |(n, k) |(m, ceil(k/64), 2) |(n, ceil(k/64), 2) |(1, n) | null | [0, 0, 32] / [1, 1, 32] |
583- | mx量化 |FLOAT8_E4M3FN |FLOAT32 |FLOAT8_E8M0 |FLOAT8_E8M0 |null/BFLOAT16/FLOAT16| null |BFLOAT16/FLOAT16|(m, k) |(n, k/8)|(m, k/64, 2) |(n, k/64, 2) |(1, n) | null | [0, 0, 32] / [1, 1, 32] |583+ | mx量化 |FLOAT8_E4M3FN |FLOAT32 |FLOAT8_E8M0 |FLOAT8_E8M0 |null/BFLOAT16/FLOAT16| null |BFLOAT16/FLOAT16|(m, k) |(n, k/8)|(m, ceil(k/64), 2) |(n, ceil(k/64), 2) |(1, n) | null | [0, 0, 32] / [1, 1, 32] |
584 | T-CG量化 |FLOAT8_E4M3FN |FLOAT4_E2M1 |null |BFLOAT16/FLOAT16 |null| uint64/int64 |BFLOAT16/FLOAT16|(m, k) |(k, n) |null |(k/32, n) |null |(1, n) | [0, 0, 32] / [1, 1, 32] |584 | T-CG量化 |FLOAT8_E4M3FN |FLOAT4_E2M1 |null |BFLOAT16/FLOAT16 |null| uint64/int64 |BFLOAT16/FLOAT16|(m, k) |(k, n) |null |(k/32, n) |null |(1, n) | [0, 0, 32] / [1, 1, 32] |
585 | T-CG量化 |FLOAT8_E4M3FN |FLOAT32 |null |BFLOAT16/FLOAT16 |null| uint64/int64 |BFLOAT16/FLOAT16|(m, k) |(k, n/8)|null |(k/32, n) |null |(1, n) | [0, 0, 32] / [1, 1, 32] |585 | T-CG量化 |FLOAT8_E4M3FN |FLOAT32 |null |BFLOAT16/FLOAT16 |null| uint64/int64 |BFLOAT16/FLOAT16|(m, k) |(k, n/8)|null |(k/32, n) |null |(1, n) | [0, 0, 32] / [1, 1, 32] |
586 586 
587 - 约束说明：587 - 约束说明：
588- - k要求32对齐, n要求8对齐。MX量化模式下k要求大于32，n大于等于8；T-CG量化模式下k要求大于等于64，n的要求大于等于8。588+ - MX量化模式下, k要求8对齐, n要求8对齐。
589+ - T-CG量化模式下，k要求32对齐且大于等于64，n要求8对齐。
589 - x1是FLOAT8_E4M3FN，x2是FLOAT32时, x2表示一个FLOAT32存储8个FLOAT4_E2M1的紧密排布的数据格式。590 - x1是FLOAT8_E4M3FN，x2是FLOAT32时, x2表示一个FLOAT32存储8个FLOAT4_E2M1的紧密排布的数据格式。
590 - MX量化模式下，bias数据类型与out数据类型需要一致；T-CG量化模式下，x2Scale的数据类型和out的数据类型需要一致。591 - MX量化模式下，bias数据类型与out数据类型需要一致；T-CG量化模式下，x2Scale的数据类型和out的数据类型需要一致。
591 592 
Mmatmul/quant_batch_matmul_v3/op_api/aclnn_quant_matmul_v4.cpp+31-35
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.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 of3 * 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").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.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -85,7 +85,8 @@ static constexpr int64_t M_RANGE1_RIGHT = 512;
85static constexpr int32_t CORE_NUM_20 = 20;85static constexpr int32_t CORE_NUM_20 = 20;
86static constexpr int64_t SUPPORTED_GROUP_SIZE = 32;86static constexpr int64_t SUPPORTED_GROUP_SIZE = 32;
87static constexpr uint64_t B4_PER_B32 = 8UL;87static constexpr uint64_t B4_PER_B32 = 8UL;
88-static constexpr int64_t SUPPORTED_K_ALIGN_NUM = 32;88+static constexpr int64_t SUPPORTED_TCG_A8W4_K_ALIGN_NUM = 32;
89+static constexpr int64_t SUPPORTED_MX_A8W4_K_ALIGN_NUM = 8;
89static constexpr int64_t SUPPORTED_N_ALIGN_NUM = 8;90static constexpr int64_t SUPPORTED_N_ALIGN_NUM = 8;
90static constexpr size_t MAX_DIM_VALUE = 2;91static constexpr size_t MAX_DIM_VALUE = 2;
91static constexpr size_t MX_SCALE_DIM_VALUE = 3;92static constexpr size_t MX_SCALE_DIM_VALUE = 3;
@@ -640,8 +641,7 @@ static inline bool CheckDimValue(const aclTensor *scale, const aclTensor *offset
640 if (offset != nullptr) {641 if (offset != nullptr) {
641 OP_CHECK_WRONG_DIMENSION(offset, 1, return false);642 OP_CHECK_WRONG_DIMENSION(offset, 1, return false);
642 if (offset->GetViewShape().GetDim(0) != x2NDim && offset->GetViewShape().GetDim(0) != 1) {643 if (offset->GetViewShape().GetDim(0) != x2NDim && offset->GetViewShape().GetDim(0) != 1) {
643- OP_LOGE(ACLNN_ERR_PARAM_INVALID,644+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "Offset 1st dim should equal to x2 n dim %ld or 1, but actual is %ld.",
644- "Offset 1st dim should equal to x2 n dim %ld or 1, but actual is %ld.",
645 x2NDim, offset->GetViewShape().GetDim(0));645 x2NDim, offset->GetViewShape().GetDim(0));
646 return false;646 return false;
647 }647 }
@@ -948,8 +948,7 @@ static inline bool CheckA8W4ScaleX1Shape(
948 "aclnnQuantMatmulWeightNzGetWorkspaceSize", "x1Scale",948 "aclnnQuantMatmulWeightNzGetWorkspaceSize", "x1Scale",
949 FormatString("%ld, %ld, %ld", x1Scale->GetViewShape().GetDim(0),949 FormatString("%ld, %ld, %ld", x1Scale->GetViewShape().GetDim(0),
950 x1Scale->GetViewShape().GetDim(1), x1Scale->GetViewShape().GetDim(2)).c_str(),950 x1Scale->GetViewShape().GetDim(1), x1Scale->GetViewShape().GetDim(2)).c_str(),
951- FormatString("the shape of x1Scale must be [%ld, %ld, 2]", groupDimM,951+ FormatString("the shape of x1Scale must be [%ld, %ld, 2]", groupDimM, CeilDiv(groupDimK, 2L)).c_str());
952- CeilDiv(groupDimK, 2L)).c_str());
953 return false;952 return false;
954 }953 }
955 }954 }
@@ -971,8 +970,7 @@ static inline bool CheckA8W4ScaleX2Shape(
971 x2Scale->GetViewShape().GetDim(2) != 2) {970 x2Scale->GetViewShape().GetDim(2) != 2) {
972 OP_LOGE_FOR_INVALID_SHAPE_WITH_REASON(971 OP_LOGE_FOR_INVALID_SHAPE_WITH_REASON(
973 "aclnnQuantMatmulWeightNzGetWorkspaceSize", "x2Scale",972 "aclnnQuantMatmulWeightNzGetWorkspaceSize", "x2Scale",
974- FormatString("%ld, %ld, %ld", x2ScaleNDim, x2ScaleGroupDim,973+ FormatString("%ld, %ld, %ld", x2ScaleNDim, x2ScaleGroupDim, x2Scale->GetViewShape().GetDim(2)).c_str(),
975- x2Scale->GetViewShape().GetDim(2)).c_str(),
976 FormatString("the shape of x2Scale must be [%ld, %ld, 2]", groupDimN,974 FormatString("the shape of x2Scale must be [%ld, %ld, 2]", groupDimN,
977 CeilDiv(groupDimK, x2ScaleReshapeFactor)).c_str());975 CeilDiv(groupDimK, x2ScaleReshapeFactor)).c_str());
978 return false;976 return false;
@@ -1014,8 +1012,7 @@ static inline bool CheckA8W4OutAndBiasShape(const TupleOptional& optionalTensors
1014 if (yScale->GetViewShape().GetDim(1) != x2NDim || yScale->GetViewShape().GetDim(0) != 1) {1012 if (yScale->GetViewShape().GetDim(1) != x2NDim || yScale->GetViewShape().GetDim(0) != 1) {
1015 OP_LOGE_FOR_INVALID_SHAPE_WITH_REASON(1013 OP_LOGE_FOR_INVALID_SHAPE_WITH_REASON(
1016 "aclnnQuantMatmulWeightNzGetWorkspaceSize", "yScale",1014 "aclnnQuantMatmulWeightNzGetWorkspaceSize", "yScale",
1017- FormatString("%ld, %ld", yScale->GetViewShape().GetDim(0), yScale->GetViewShape().GetDim(1))1015+ FormatString("%ld, %ld", yScale->GetViewShape().GetDim(0), yScale->GetViewShape().GetDim(1)).c_str(),
1018- .c_str(),
1019 FormatString("the shape of yScale must be [1, %ld]", x2NDim).c_str());1016 FormatString("the shape of yScale must be [1, %ld]", x2NDim).c_str());
1020 return false;1017 return false;
1021 }1018 }
@@ -1023,7 +1020,7 @@ static inline bool CheckA8W4OutAndBiasShape(const TupleOptional& optionalTensors
1023 return true;1020 return true;
1024}1021}
1025 1022 
1026-static inline bool CheckA8W4X1X2Shape(int64_t x1KDim, int64_t x2KDim, int64_t x2NDim) {1023+static inline bool CheckA8W4X1X2Shape(int64_t x1KDim, int64_t x2KDim, int64_t x2NDim, bool isMx) {
1027 // CHECK x1KDim1024 // CHECK x1KDim
1028 if (x1KDim <= 0) {1025 if (x1KDim <= 0) {
1029 OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(1026 OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(
@@ -1037,19 +1034,24 @@ static inline bool CheckA8W4X1X2Shape(int64_t x1KDim, int64_t x2KDim, int64_t x2
1037 "the n dimension of x2 must be greater than 0");1034 "the n dimension of x2 must be greater than 0");
1038 return false;1035 return false;
1039 }1036 }
1040- 
1041 if (x1KDim != x2KDim) {1037 if (x1KDim != x2KDim) {
1042 OP_LOGE_FOR_INVALID_VALUES_WITH_REASON(1038 OP_LOGE_FOR_INVALID_VALUES_WITH_REASON(
1043- "aclnnQuantMatmulWeightNzGetWorkspaceSize", "x1 K, x2 K",1039+ "aclnnQuantMatmulWeightNzGetWorkspaceSize", "x1 K, x2 K", FormatString("%ld, %ld", x1KDim, x2KDim).c_str(),
1044- FormatString("%ld, %ld", x1KDim, x2KDim).c_str(),
1045 "the k dimension of x1 and x2 must be equal");1040 "the k dimension of x1 and x2 must be equal");
1046 return false;1041 return false;
1047 }1042 }
1048- if (x1KDim % SUPPORTED_K_ALIGN_NUM != 0 || x1KDim <= SUPPORTED_K_ALIGN_NUM) {1043+ if (isMx && (x1KDim % SUPPORTED_MX_A8W4_K_ALIGN_NUM != 0)) { // Mx量化k方向8对齐
1044+ OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(
1045+ "aclnnQuantMatmulWeightNzGetWorkspaceSize", "x1", std::to_string(x1KDim).c_str(),
1046+ FormatString("the k dimension of x1 must be aligned to %ld for MX quantization",
1047+ SUPPORTED_MX_A8W4_K_ALIGN_NUM).c_str());
1048+ return false;
1049+ }
1050+ if (!isMx && (x1KDim % SUPPORTED_TCG_A8W4_K_ALIGN_NUM != 0 || x1KDim <= SUPPORTED_TCG_A8W4_K_ALIGN_NUM)) {
1049 OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(1051 OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(
1050 "aclnnQuantMatmulWeightNzGetWorkspaceSize", "x1", std::to_string(x1KDim).c_str(),1052 "aclnnQuantMatmulWeightNzGetWorkspaceSize", "x1", std::to_string(x1KDim).c_str(),
1051 FormatString("the k dimension of x1 must be aligned to %ld and greater than %ld",1053 FormatString("the k dimension of x1 must be aligned to %ld and greater than %ld",
1052- SUPPORTED_K_ALIGN_NUM, SUPPORTED_K_ALIGN_NUM).c_str());1054+ SUPPORTED_TCG_A8W4_K_ALIGN_NUM, SUPPORTED_TCG_A8W4_K_ALIGN_NUM).c_str());
1053 return false;1055 return false;
1054 }1056 }
1055 if (x2NDim % SUPPORTED_N_ALIGN_NUM != 0) {1057 if (x2NDim % SUPPORTED_N_ALIGN_NUM != 0) {
@@ -1084,7 +1086,9 @@ static inline bool CheckA8W4Shape(const TupleTensor &mandatoryTensors, const Tup
1084 "the m dimension of x1 must be greater than 0");1086 "the m dimension of x1 must be greater than 0");
1085 return false;1087 return false;
1086 }1088 }
1087- CHECK_RET(CheckA8W4X1X2Shape(x1KDim, x2KDim, x2NDim), false);1089+ auto x1Scale = std::get<INDEX_PERTOKEN_IN_OPTIONAL_TUPLE>(optionalTensors);
1090+ auto x2Scale = std::get<INDEX_SCALE_IN_MANDTORY_TUPLE>(mandatoryTensors);
1091+ CHECK_RET(CheckA8W4X1X2Shape(x1KDim, x2KDim, x2NDim, IsMicroScaling(x1Scale, x2Scale)), false);
1088 int64_t groupDimK = (x2KDim + SUPPORTED_GROUP_SIZE - 1) / SUPPORTED_GROUP_SIZE;1092 int64_t groupDimK = (x2KDim + SUPPORTED_GROUP_SIZE - 1) / SUPPORTED_GROUP_SIZE;
1089 int64_t groupDimM = x1MDim;1093 int64_t groupDimM = x1MDim;
1090 int64_t groupDimN = x2NDim;1094 int64_t groupDimN = x2NDim;
@@ -1497,8 +1501,8 @@ static inline bool CheckInputAttrExistence(const TupleAttr &boolsTrans, const Tu
1497 "aclnnQuantMatmulWeightNzGetWorkspaceSize", "transposeX2", transposeX2 ? "true" : "false",1501 "aclnnQuantMatmulWeightNzGetWorkspaceSize", "transposeX2", transposeX2 ? "true" : "false",
1498 "in A8W4 scenario with NZ format, when the quant mode is t-cg, transposeX2 must be false");1502 "in A8W4 scenario with NZ format, when the quant mode is t-cg, transposeX2 must be false");
1499 return false;1503 return false;
1500- } else if (IsMicroScaling(x1Scale, x2Scale) && !transposeX2) {1504+ } else if (IsMicroScaling(x1Scale, x2Scale) && !transposeX2) {
1501- // A8W4 scenario with mx quant mode1505+ // A8W4 scenario with mx quant mode
1502 OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(1506 OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(
1503 "aclnnQuantMatmulWeightNzGetWorkspaceSize", "transposeX2", transposeX2 ? "true" : "false",1507 "aclnnQuantMatmulWeightNzGetWorkspaceSize", "transposeX2", transposeX2 ? "true" : "false",
1504 "in A8W4 scenario with NZ format, when the quant mode is mx, transposeX2 must be true");1508 "in A8W4 scenario with NZ format, when the quant mode is mx, transposeX2 must be true");
@@ -1521,31 +1525,23 @@ static inline bool CheckDimRangeA8W4(const TupleTensor& mandatoryTensors, const
1521 auto bias = std::get<INDEX_BIAS_IN_OPTIONAL_TUPLE>(optionalTensors);1525 auto bias = std::get<INDEX_BIAS_IN_OPTIONAL_TUPLE>(optionalTensors);
1522 1526 
1523 if (x1->GetViewShape().GetDimNum() != MAX_DIM_VALUE) {1527 if (x1->GetViewShape().GetDimNum() != MAX_DIM_VALUE) {
1524- OP_LOGE_FOR_INVALID_SHAPEDIM_WITH_REASON(1528+ OP_LOGE_FOR_INVALID_SHAPEDIM_WITH_REASON("aclnnQuantMatmulWeightNzGetWorkspaceSize", "x1",
1525- "aclnnQuantMatmulWeightNzGetWorkspaceSize", "x1",1529+ FormatString("%zuD", x1->GetViewShape().GetDimNum()).c_str(), "the shape dim of x1 must be 2");
1526- FormatString("%zuD", x1->GetViewShape().GetDimNum()).c_str(),
1527- "the shape dim of x1 must be 2");
1528 return false;1530 return false;
1529 }1531 }
1530 if (x2->GetViewShape().GetDimNum() != MAX_DIM_VALUE) {1532 if (x2->GetViewShape().GetDimNum() != MAX_DIM_VALUE) {
1531- OP_LOGE_FOR_INVALID_SHAPEDIM_WITH_REASON(1533+ OP_LOGE_FOR_INVALID_SHAPEDIM_WITH_REASON("aclnnQuantMatmulWeightNzGetWorkspaceSize", "x2",
1532- "aclnnQuantMatmulWeightNzGetWorkspaceSize", "x2",1534+ FormatString("%zuD", x2->GetViewShape().GetDimNum()).c_str(), "the shape dim of x2 must be 2");
1533- FormatString("%zuD", x2->GetViewShape().GetDimNum()).c_str(),
1534- "the shape dim of x2 must be 2");
1535 return false;1535 return false;
1536 }1536 }
1537 if (bias != nullptr && bias->GetViewShape().GetDimNum() != MAX_DIM_VALUE) {1537 if (bias != nullptr && bias->GetViewShape().GetDimNum() != MAX_DIM_VALUE) {
1538- OP_LOGE_FOR_INVALID_SHAPEDIM_WITH_REASON(1538+ OP_LOGE_FOR_INVALID_SHAPEDIM_WITH_REASON("aclnnQuantMatmulWeightNzGetWorkspaceSize", "bias",
1539- "aclnnQuantMatmulWeightNzGetWorkspaceSize", "bias",1539+ FormatString("%zuD", bias->GetViewShape().GetDimNum()).c_str(), "the shape dim of bias must be 2");
1540- FormatString("%zuD", bias->GetViewShape().GetDimNum()).c_str(),
1541- "the shape dim of bias must be 2");
1542 return false;1540 return false;
1543 }1541 }
1544 if (out->GetViewShape().GetDimNum() != MAX_DIM_VALUE) {1542 if (out->GetViewShape().GetDimNum() != MAX_DIM_VALUE) {
1545- OP_LOGE_FOR_INVALID_SHAPEDIM_WITH_REASON(1543+ OP_LOGE_FOR_INVALID_SHAPEDIM_WITH_REASON("aclnnQuantMatmulWeightNzGetWorkspaceSize", "out",
1546- "aclnnQuantMatmulWeightNzGetWorkspaceSize", "out",1544+ FormatString("%zuD", out->GetViewShape().GetDimNum()).c_str(), "the shape dim of out must be 2");
1547- FormatString("%zuD", out->GetViewShape().GetDimNum()).c_str(),
1548- "the shape dim of out must be 2");
1549 return false;1545 return false;
1550 }1546 }
1551 OP_LOGD("QuantMatmul check dimension range success.");1547 OP_LOGD("QuantMatmul check dimension range success.");
Mmatmul/quant_batch_matmul_v4/docs/aclnnQuantMatmulV5.md+1-1
@@ -902,7 +902,7 @@ aclnnStatus aclnnQuantMatmulV5(
902 - 注：上表中gsM、gsK和gsN分别表示groupSizeM、groupSizeK和groupSizeN。gsM、gsK和gsN为0的维度会自动推导，上表中是不用自动推导的情况。902 - 注：上表中gsM、gsK和gsN分别表示groupSizeM、groupSizeK和groupSizeN。gsM、gsK和gsN为0的维度会自动推导，上表中是不用自动推导的情况。
903 - mx全量化场景下，当x2数据类型为FLOAT8_E4M3FN/FLOAT8_E5M2时，x1和x1Scale的转置属性需要保持一致，x2和x2Scale的转置属性需要保持一致。903 - mx全量化场景下，当x2数据类型为FLOAT8_E4M3FN/FLOAT8_E5M2时，x1和x1Scale的转置属性需要保持一致，x2和x2Scale的转置属性需要保持一致。
904 - mx全量化场景下，当x2数据类型为FLOAT4_E2M1时，x1和x2的内轴必须为偶数，且k必须大于2。904 - mx全量化场景下，当x2数据类型为FLOAT4_E2M1时，x1和x2的内轴必须为偶数，且k必须大于2。
905- - mx伪量化场景下，当x2数据类型为FLOAT4_E2M1时，transposeX1为false且transposeX2为true，不支持batch轴。数据格式支持ND格式。要求支持k是32的倍数。905+ - mx伪量化场景下，当x2数据类型为FLOAT4_E2M1时，transposeX1为false且transposeX2为true，不支持batch轴。数据格式支持ND格式。要求支持k是8的倍数。
906 - mx伪量化场景下，bias为可选参数。数据类型支持BFLOAT16或FLOAT16，数据类型要求与输出类型保持一致。数据格式支持ND，shape支持2维，shape表示(1，n)。如不需要使用该参数，传入nullptr。906 - mx伪量化场景下，bias为可选参数。数据类型支持BFLOAT16或FLOAT16，数据类型要求与输出类型保持一致。数据格式支持ND，shape支持2维，shape表示(1，n)。如不需要使用该参数，传入nullptr。
907 907 
908 </details>908 </details>
Mmatmul/quant_batch_matmul_v4/op_host/op_api/aclnn_quant_matmul_v5.cpp+39-26
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.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 of3 * 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").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.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -46,7 +46,8 @@ static const int64_t SUPPORTED_GROUP_SIZE = 32;
46static const int64_t MAX_SHAPE_SIZE_A8W4_INT = 29576;46static const int64_t MAX_SHAPE_SIZE_A8W4_INT = 29576;
47static const int64_t SUPPORTED_GROUP_SIZE_A8W4_INT = 256;47static const int64_t SUPPORTED_GROUP_SIZE_A8W4_INT = 256;
48static const int64_t SUPPORTED_K_ALIGN_NUM = 64;48static const int64_t SUPPORTED_K_ALIGN_NUM = 64;
49-static const int64_t SUPPORTED_A8W4_K_ALIGN_NUM = 32;49+static const int64_t SUPPORTED_TCG_A8W4_K_ALIGN_NUM = 32;
50+static const int64_t SUPPORTED_MX_A8W4_K_ALIGN_NUM = 8;
50static const int64_t SUPPORTED_N_ALIGN_NUM = 8;51static const int64_t SUPPORTED_N_ALIGN_NUM = 8;
51static const uint64_t GROUP_M_OFFSET = 32;52static const uint64_t GROUP_M_OFFSET = 32;
52static const uint64_t GROUP_N_OFFSET = 16;53static const uint64_t GROUP_N_OFFSET = 16;
@@ -591,18 +592,44 @@ static inline bool CheckOutAndOffsetShape(const TupleQuant& quantTensors, int64_
591 return true;592 return true;
592}593}
593 594 
594-static inline bool CheckKDimAndBasicShape(595+static inline bool CheckA8W4KDim(const TupleInput& inputTensors, const TupleQuant& quantTensors, int64_t kDim)
595- const TupleInput& inputTensors, const TupleQuant& quantTensors, int64_t x1KDim, int64_t x2KDim, int64_t x2NDim)
596{596{
597 auto x1 = std::get<INDEX_X1_IN_INPUT_TUPLE>(inputTensors);597 auto x1 = std::get<INDEX_X1_IN_INPUT_TUPLE>(inputTensors);
598 auto x2 = std::get<INDEX_X2_IN_INPUT_TUPLE>(inputTensors);598 auto x2 = std::get<INDEX_X2_IN_INPUT_TUPLE>(inputTensors);
599 auto x1Scale = std::get<INDEX_X1_SCALE_IN_QUANT_TUPLE>(quantTensors);599 auto x1Scale = std::get<INDEX_X1_SCALE_IN_QUANT_TUPLE>(quantTensors);
600 auto x2Scale = std::get<INDEX_X2_SCALE_IN_QUANT_TUPLE>(quantTensors);600 auto x2Scale = std::get<INDEX_X2_SCALE_IN_QUANT_TUPLE>(quantTensors);
601- bool isPerChannel = x2Scale->GetViewShape().GetDimNum() == 1;601+ if (isA8W4IntAfterPre(x1, x2)) {
602- bool isA8W4INT = isA8W4IntAfterPre(x1, x2);602+ bool isPerChannel = x2Scale->GetViewShape().GetDimNum() == 1;
603- bool isA8W4TCG = isA8W4FloatTCG(x1, x2, x1Scale);603+ size_t kAlign = isPerChannel ? SUPPORTED_K_ALIGN_NUM_INT4 : SUPPORTED_GROUP_SIZE_A8W4_INT;
604- bool isA8W4Mx = isA8W4FloatMx(x1, x2, x1Scale, x2Scale);604+ if (kDim % kAlign != 0) {
605- // CHECK x1KDim605+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "the k dim must be align to %ld, which is %ld", kAlign, kDim);
606+ return false;
607+ }
608+ } else if (isA8W4FloatTCG(x1, x2, x1Scale)) {
609+ if (kDim % SUPPORTED_TCG_A8W4_K_ALIGN_NUM != 0 || kDim <= SUPPORTED_TCG_A8W4_K_ALIGN_NUM) {
610+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "the k dim must to be aligned to %ld and more than %ld, which is %ld",
611+ SUPPORTED_TCG_A8W4_K_ALIGN_NUM, SUPPORTED_TCG_A8W4_K_ALIGN_NUM, kDim);
612+ return false;
613+ }
614+ } else if (isA8W4FloatMx(x1, x2, x1Scale, x2Scale)) {
615+ if (kDim % SUPPORTED_MX_A8W4_K_ALIGN_NUM != 0) {
616+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "The k dim must be aligned to %ld for MX quantization, which is %ld",
617+ SUPPORTED_MX_A8W4_K_ALIGN_NUM, kDim);
618+ return false;
619+ }
620+ } else {
621+ if (kDim % SUPPORTED_K_ALIGN_NUM != 0) {
622+ OP_LOGE(ACLNN_ERR_PARAM_INVALID, "the k dim must be align to %ld, which is %ld", SUPPORTED_K_ALIGN_NUM, kDim);
623+ return false;
624+ }
625+ }
626+ return true;
627+}
628+ 
629+static inline bool CheckKDimAndBasicShape(
630+ const TupleInput& inputTensors, const TupleQuant& quantTensors, int64_t x1KDim, int64_t x2KDim, int64_t x2NDim)
631+{
632+ // CHECK basic shape
606 auto npuArch = op::GetCurrentPlatformInfo().GetCurNpuArch();633 auto npuArch = op::GetCurrentPlatformInfo().GetCurNpuArch();
607 if (npuArch == NpuArch::DAV_2201) { // A8W4INT A2 A3634 if (npuArch == NpuArch::DAV_2201) { // A8W4INT A2 A3
608 if (x1KDim <= 0 || x1KDim > MAX_SHAPE_SIZE_A8W4_INT) {635 if (x1KDim <= 0 || x1KDim > MAX_SHAPE_SIZE_A8W4_INT) {
@@ -624,23 +651,9 @@ static inline bool CheckKDimAndBasicShape(
624 x1KDim, x2KDim);651 x1KDim, x2KDim);
625 return false;652 return false;
626 }653 }
627- if (isA8W4INT) {654+ // CHECK kDim
628- size_t kAlign = isPerChannel ? SUPPORTED_K_ALIGN_NUM_INT4 : SUPPORTED_GROUP_SIZE_A8W4_INT;655+ if (!CheckA8W4KDim(inputTensors, quantTensors, x1KDim)) {
629- if (x1KDim % kAlign != 0) {656+ return false;
630- OP_LOGE(ACLNN_ERR_PARAM_INVALID, "the k dim must be align to %ld, which is %ld", kAlign, x1KDim);
631- return false;
632- }
633- } else if (isA8W4TCG || isA8W4Mx) {
634- if (x1KDim % SUPPORTED_A8W4_K_ALIGN_NUM != 0 || x1KDim <= SUPPORTED_A8W4_K_ALIGN_NUM) {
635- OP_LOGE(ACLNN_ERR_PARAM_INVALID,"the k dim must to be aligned to %ld and more than %ld, which is %ld",
636- SUPPORTED_A8W4_K_ALIGN_NUM, SUPPORTED_A8W4_K_ALIGN_NUM, x1KDim);
637- return false;
638- }
639- } else {
640- if (x1KDim % SUPPORTED_K_ALIGN_NUM != 0) {
641- OP_LOGE(ACLNN_ERR_PARAM_INVALID, "the k dim must be align to %ld, which is %ld", SUPPORTED_K_ALIGN_NUM, x1KDim);
642- return false;
643- }
644 }657 }
645 return true;658 return true;
646}659}
Mmatmul/quant_batch_matmul_v4/op_host/op_tiling/arch35/quant_batch_matmul_v4_basic_block_tiling.cpp+4-4
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.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 of3 * 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").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.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -121,7 +121,7 @@ void QuantBatchMatmulV4BasicBlockTiling::SetShape(int64_t mSize, int64_t nSize,
121 basicBlockParam_.mSize = mSize;121 basicBlockParam_.mSize = mSize;
122 basicBlockParam_.nSize = nSize;122 basicBlockParam_.nSize = nSize;
123 basicBlockParam_.kSize = kSize;123 basicBlockParam_.kSize = kSize;
124- basicBlockParam_.singleK = kSize;124+ basicBlockParam_.singleK = ops::CeilAlign(kSize, ONE_BLOCK_SIZE); // k方向单核处理粒度需32对齐
125 basicBlockParam_.groupSize = groupSize;125 basicBlockParam_.groupSize = groupSize;
126 OP_LOGI(opName_, "Init shape param, mSize: %ld, nSize: %ld, kSize: %ld, groupSize: %ld", basicBlockParam_.mSize,126 OP_LOGI(opName_, "Init shape param, mSize: %ld, nSize: %ld, kSize: %ld, groupSize: %ld", basicBlockParam_.mSize,
127 basicBlockParam_.nSize, basicBlockParam_.kSize, basicBlockParam_.groupSize);127 basicBlockParam_.nSize, basicBlockParam_.kSize, basicBlockParam_.groupSize);
@@ -183,7 +183,7 @@ bool QuantBatchMatmulV4BasicBlockTiling::ValidateInputParam() const
183 183 
184 OP_TILING_CHECK(basicBlockParam_.groupSize < 0,184 OP_TILING_CHECK(basicBlockParam_.groupSize < 0,
185 VECTOR_INNER_ERR_REPORT_TILIING(185 VECTOR_INNER_ERR_REPORT_TILIING(
186- opName_, "Invalid param, groupSize must be greater than or equalt to 0, groupSize: %ld",186+ opName_, "Invalid param, groupSize must be greater than or equal to 0, groupSize: %ld",
187 basicBlockParam_.groupSize),187 basicBlockParam_.groupSize),
188 return false);188 return false);
189 189 
@@ -313,7 +313,7 @@ void QuantBatchMatmulV4BasicBlockTiling::GetMte2DataSizeMx(BasicBlockParam& basi
313 CeilDiv(basicBlockParam.singleK, basicBlockParam.basicBlock.baseK * basicBlockParam.l1Param.stepKa *313 CeilDiv(basicBlockParam.singleK, basicBlockParam.basicBlock.baseK * basicBlockParam.l1Param.stepKa *
314 basicBlockParam.l1Param.scaleFactor);314 basicBlockParam.l1Param.scaleFactor);
315 // 每行搬运量都小于cacheline, 按照cacheline大小计算315 // 每行搬运量都小于cacheline, 按照cacheline大小计算
316- int64_t scaleSingleK = basicBlockParam.singleK / basicBlockParam.groupSize;316+ int64_t scaleSingleK = CeilDiv(basicBlockParam.singleK, basicBlockParam.groupSize);
317 int64_t scaleKLoopNumHbm = min(CeilDiv(scaleSingleK, platformParam_.cacheLine), scaleKloopNumTotal);317 int64_t scaleKLoopNumHbm = min(CeilDiv(scaleSingleK, platformParam_.cacheLine), scaleKloopNumTotal);
318 int64_t scaleKLoopNumL2 = scaleKloopNumTotal - scaleKLoopNumHbm;318 int64_t scaleKLoopNumL2 = scaleKloopNumTotal - scaleKLoopNumHbm;
319 int64_t scaleASizeHbm = basicBlockParam.singleM * platformParam_.cacheLine * scaleKLoopNumHbm;319 int64_t scaleASizeHbm = basicBlockParam.singleM * platformParam_.cacheLine * scaleKLoopNumHbm;
Mmatmul/quant_batch_matmul_v4/op_host/op_tiling/arch35/quant_batch_matmul_v4_basic_block_tiling.h+2-1
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.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 of3 * 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").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.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -42,6 +42,7 @@ constexpr double BYTE_BITS = 8;
42constexpr int64_t BITS_16 = 16;42constexpr int64_t BITS_16 = 16;
43constexpr int64_t BITS_4 = 4;43constexpr int64_t BITS_4 = 4;
44constexpr int64_t BLOCK_CUBE = 16;44constexpr int64_t BLOCK_CUBE = 16;
45+constexpr int64_t ONE_BLOCK_SIZE = 32;
45constexpr int64_t BUFF_NUM_1 = 1;46constexpr int64_t BUFF_NUM_1 = 1;
46constexpr int64_t BUFF_NUM_2 = 2;47constexpr int64_t BUFF_NUM_2 = 2;
47constexpr int64_t BUFF_NUM_3 = 3;48constexpr int64_t BUFF_NUM_3 = 3;
Mmatmul/quant_batch_matmul_v4/op_host/op_tiling/arch35/quant_batch_matmul_v4_pergroup_arch35_tiling.cpp+7-9
@@ -71,6 +71,11 @@ ge::graphStatus QuantBatchMatmulV4PergroupArch35Tiling::CalcDequantTiling(
71 71 
72ge::graphStatus QuantBatchMatmulV4PergroupArch35Tiling::DoOpTiling()72ge::graphStatus QuantBatchMatmulV4PergroupArch35Tiling::DoOpTiling()
73{73{
74+ OPS_CHECK_NULL_WITH_CONTEXT(context_, context_->GetRawTilingData());
75+ OP_TILING_CHECK(context_->GetRawTilingData()->GetCapacity() < tilingDataSize_,
76+ CUBE_INNER_ERR_REPORT(inputParams_.opName, "context tiling data capacity %zu < actual tiling data size %zu.",
77+ context_->GetRawTilingData()->GetCapacity(), tilingDataSize_),
78+ return ge::GRAPH_FAILED);
74 isUbQuant_ = true;79 isUbQuant_ = true;
75 InitCompileInfo();80 InitCompileInfo();
76 SetTransAttr(trans_);81 SetTransAttr(trans_);
@@ -108,9 +113,8 @@ ge::graphStatus QuantBatchMatmulV4PergroupArch35Tiling::DoOpTiling()
108 stepK = std::max(1U, stepK);113 stepK = std::max(1U, stepK);
109 basicTiling_.stepKa = stepK;114 basicTiling_.stepKa = stepK;
110 basicTiling_.stepKb = stepK;115 basicTiling_.stepKb = stepK;
111- OP_LOGD(116+ OP_LOGD(inputParams_.opName, "arch35 int8 tiling: groupSizeK=%u, stepK=%u, l1Size=%u",
112- inputParams_.opName, "arch35 int8 tiling: groupSizeK=%u, stepK=%u, l1Size=%u", inputParams_.groupSizeK, stepK,117+ inputParams_.groupSizeK, stepK, l1Size);
113- l1Size);
114 118 
115 basicTiling_.baseK = inputParams_.groupSizeK;119 basicTiling_.baseK = inputParams_.groupSizeK;
116 QuantBatchMatmulV3BasicTiling::DoL2CacheTiling();120 QuantBatchMatmulV3BasicTiling::DoL2CacheTiling();
@@ -158,12 +162,6 @@ ge::graphStatus QuantBatchMatmulV4PergroupArch35Tiling::CheckContext()
158 OPS_CHECK_NULL_WITH_CONTEXT(context_, outputDesc);162 OPS_CHECK_NULL_WITH_CONTEXT(context_, outputDesc);
159 OPS_CHECK_NULL_WITH_CONTEXT(context_, context_->GetRawTilingData());163 OPS_CHECK_NULL_WITH_CONTEXT(context_, context_->GetRawTilingData());
160 OPS_CHECK_NULL_WITH_CONTEXT(context_, context_->GetRawTilingData()->GetData());164 OPS_CHECK_NULL_WITH_CONTEXT(context_, context_->GetRawTilingData()->GetData());
161- OP_TILING_CHECK(
162- context_->GetRawTilingData()->GetCapacity() < tilingDataSize_,
163- CUBE_INNER_ERR_REPORT(
164- inputParams_.opName, "context tiling data capacity %zu < actual tiling data size %zu.",
165- context_->GetRawTilingData()->GetCapacity(), tilingDataSize_),
166- return ge::GRAPH_FAILED);
167 return ge::GRAPH_SUCCESS;165 return ge::GRAPH_SUCCESS;
168}166}
169 167 
Mmatmul/quant_batch_matmul_v4/op_host/op_tiling/arch35/quant_batch_matmul_v4_reg_base_tiling.cpp+35-15
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.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 of3 * 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").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.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -54,12 +54,23 @@ bool QuantBatchMatmulV4RegBase::CheckA8W4Params() const
54 VECTOR_INNER_ERR_REPORT_TILIING(inputParams_.opName,54 VECTOR_INNER_ERR_REPORT_TILIING(inputParams_.opName,
55 "Invalid params, only support x2 not transpose FORMAT_FRACTAL_NZ."),55 "Invalid params, only support x2 not transpose FORMAT_FRACTAL_NZ."),
56 return false);56 return false);
57- OP_CHECK_IF(57+ 
58- inputParams_.groupSize <= 0 || inputParams_.kSize < inputParams_.groupSize,58+ if (inputParams_.antiQuantType == QuantType::MX) {
59- VECTOR_INNER_ERR_REPORT_TILIING(59+ OP_CHECK_IF(
60- inputParams_.opName,60+ inputParams_.groupSize != MX_GROUP_SIZE,
61- "Invalid params, groupSize must be greater than 0 and less than kSize, kSize: %lu, groupSize: %lu.",61+ VECTOR_INNER_ERR_REPORT_TILIING(
62- inputParams_.kSize, inputParams_.groupSize), return false);62+ inputParams_.opName,
63+ "Invalid params, groupSize must be 32 for MX quantization, but got %lu.",
64+ inputParams_.groupSize), return false);
65+ } else {
66+ OP_CHECK_IF(
67+ inputParams_.groupSize <= 0 || inputParams_.kSize < inputParams_.groupSize,
68+ VECTOR_INNER_ERR_REPORT_TILIING(
69+ inputParams_.opName,
70+ "Invalid params, groupSize must be greater than 0 and less than kSize, kSize: %lu, groupSize: %lu.",
71+ inputParams_.kSize, inputParams_.groupSize), return false);
72+ }
73+ 
63 OP_CHECK_IF(inputParams_.groupSize % GROUP_ALIGN_SIZE > 0,74 OP_CHECK_IF(inputParams_.groupSize % GROUP_ALIGN_SIZE > 0,
64 VECTOR_INNER_ERR_REPORT_TILIING(inputParams_.opName,75 VECTOR_INNER_ERR_REPORT_TILIING(inputParams_.opName,
65 "Invalid params, groupSize must be 32 aligned, groupSize: %lu.",76 "Invalid params, groupSize must be 32 aligned, groupSize: %lu.",
@@ -67,19 +78,28 @@ bool QuantBatchMatmulV4RegBase::CheckA8W4Params() const
67 // A8W4 Nz场景要求n为32B对齐78 // A8W4 Nz场景要求n为32B对齐
68 OP_CHECK_IF(inputParams_.bFormat == ge::FORMAT_FRACTAL_NZ && inputParams_.nSize % N_ALIGN_SIZE > 0,79 OP_CHECK_IF(inputParams_.bFormat == ge::FORMAT_FRACTAL_NZ && inputParams_.nSize % N_ALIGN_SIZE > 0,
69 VECTOR_INNER_ERR_REPORT_TILIING(inputParams_.opName,80 VECTOR_INNER_ERR_REPORT_TILIING(inputParams_.opName,
70- "Invalid params, nSize only support aligned to 64 when weight format is NZ, but nSize is %lu.",81+ "Invalid params, nSize only support aligned to 8 when weight format is NZ, but nSize is %lu.",
71 inputParams_.nSize), return false);82 inputParams_.nSize), return false);
72 return true;83 return true;
73}84}
74 85 
75bool QuantBatchMatmulV4RegBase::CustomCheck() const86bool QuantBatchMatmulV4RegBase::CustomCheck() const
76{87{
77- OP_CHECK_IF(88+ if (inputParams_.antiQuantType == QuantType::MX) {
78- inputParams_.kSize % K_ALIGN_SIZE > 0 || inputParams_.kSize <= K_ALIGN_SIZE,89+ OP_CHECK_IF(
79- VECTOR_INNER_ERR_REPORT_TILIING(90+ inputParams_.kSize % K_ALIGN_SIZE_MX > 0,
80- inputParams_.opName, "Invalid params, kSize must be aligned to 32 and greater than 32, but got %lu.",91+ VECTOR_INNER_ERR_REPORT_TILIING(
81- inputParams_.kSize),92+ inputParams_.opName, "Invalid params, kSize must be aligned to 8 for MX quantization, but got %lu.",
82- return false);93+ inputParams_.kSize),
94+ return false);
95+ } else {
96+ OP_CHECK_IF(
97+ inputParams_.kSize % K_ALIGN_SIZE > 0 || inputParams_.kSize <= K_ALIGN_SIZE,
98+ VECTOR_INNER_ERR_REPORT_TILIING(
99+ inputParams_.opName, "Invalid params, kSize must be aligned to 32 and greater than 32, but got %lu.",
100+ inputParams_.kSize),
101+ return false);
102+ }
83 103 
84 OP_CHECK_IF((inputParams_.cDtype != ge::DT_BF16) && (inputParams_.cDtype != ge::DT_FLOAT16),104 OP_CHECK_IF((inputParams_.cDtype != ge::DT_BF16) && (inputParams_.cDtype != ge::DT_FLOAT16),
85 VECTOR_INNER_ERR_REPORT_TILIING(inputParams_.opName, "Invalid params, output only support DT_BF16 or DT_FLOAT16."),105 VECTOR_INNER_ERR_REPORT_TILIING(inputParams_.opName, "Invalid params, output only support DT_BF16 or DT_FLOAT16."),
@@ -399,7 +419,7 @@ void QuantBatchMatmulV4RegBase::PrintCVTilingData(const bool debugLevel) const
399 << " AL1Pingpong: " << tilingData_->AL1Pingpong << " BL1Pingpong: " << tilingData_->BL1Pingpong;419 << " AL1Pingpong: " << tilingData_->AL1Pingpong << " BL1Pingpong: " << tilingData_->BL1Pingpong;
400 if (debugLevel) {420 if (debugLevel) {
401 OPS_LOG_D(inputParams_.opName, "tiling data: %s", ss.str().c_str());421 OPS_LOG_D(inputParams_.opName, "tiling data: %s", ss.str().c_str());
402- }else {422+ } else {
403 OPS_LOG_E(inputParams_.opName, "tiling data: %s", ss.str().c_str());423 OPS_LOG_E(inputParams_.opName, "tiling data: %s", ss.str().c_str());
404 }424 }
405 PrintMatMulTiling();425 PrintMatMulTiling();
Mmatmul/quant_batch_matmul_v4/op_host/op_tiling/arch35/quant_batch_matmul_v4_tiling.cpp+40-25
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.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 of3 * 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").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.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -149,30 +149,8 @@ ge::graphStatus QuantBatchMatmulV4TilingBase::GetShapeAttrsInfo()
149 inputParams_.weightNz = inputParams_.bFormat == ge::FORMAT_FRACTAL_NZ;149 inputParams_.weightNz = inputParams_.bFormat == ge::FORMAT_FRACTAL_NZ;
150 OP_TILING_CHECK(!AnalyzeQuantType() || !AnalyzeAttrs() || !AnalyzeInputs() || !AnalyzeDtype(),150 OP_TILING_CHECK(!AnalyzeQuantType() || !AnalyzeAttrs() || !AnalyzeInputs() || !AnalyzeDtype(),
151 VECTOR_INNER_ERR_REPORT_TILIING(inputParams_.opName, "Fail to analyze context info"), return ge::GRAPH_FAILED);151 VECTOR_INNER_ERR_REPORT_TILIING(inputParams_.opName, "Fail to analyze context info"), return ge::GRAPH_FAILED);
152- bool maxDimCheck = inputParams_.kSize > MAX_SHAPE_DIM || inputParams_.nSize > MAX_SHAPE_DIM;152+ OP_TILING_CHECK(CheckInputParams() != ge::GRAPH_SUCCESS,
153- if (inputParams_.transA) {153+ VECTOR_INNER_ERR_REPORT_TILIING(inputParams_.opName, "invalid input parameters"),
154- maxDimCheck |= inputParams_.mSize > MAX_SHAPE_DIM;
155- }
156- OP_TILING_CHECK(inputParams_.supportL0c2Out && maxDimCheck,
157- VECTOR_INNER_ERR_REPORT_TILIING(inputParams_.opName,
158- "only support MKN in range [1, %lu], get actual value[%lu, %lu, %lu]",
159- MAX_SHAPE_DIM, inputParams_.mSize, inputParams_.kSize, inputParams_.nSize),
160- return ge::GRAPH_FAILED);
161- OP_TILING_CHECK(inputParams_.groupSize > inputParams_.kSize || inputParams_.groupSize % MIN_GROUP_SIZE != 0,
162- VECTOR_INNER_ERR_REPORT_TILIING(inputParams_.opName,
163- "Only support group size greater than %lu, less than K and align to %lu, get K[%lu] and group size[%lu]",
164- MIN_GROUP_SIZE, MIN_GROUP_SIZE, inputParams_.kSize, inputParams_.groupSize),
165- return ge::GRAPH_FAILED);
166- OP_TILING_CHECK(inputParams_.supportL0c2Out && !inputParams_.supportL12BtBf16 &&
167- inputParams_.bFormat == ge::FORMAT_FRACTAL_NZ &&
168- (inputParams_.bDtype != ge::DT_INT8 || inputParams_.antiQuantType != QuantType::PER_CHANNEL ||
169- inputParams_.cDtype == ge::DT_INT8),
170- VECTOR_INNER_ERR_REPORT_TILIING(inputParams_.opName,
171- "weight Nz only support weight dtype INT8 per-channel scene, and not support quant scale, "
172- "current input bDtype[%s], antiquantType[%d], cDtype[%s]",
173- ge::TypeUtils::DataTypeToSerialString(inputParams_.bDtype).c_str(),
174- static_cast<int>(inputParams_.antiQuantType),
175- ge::TypeUtils::DataTypeToSerialString(inputParams_.bDtype).c_str()),
176 return ge::GRAPH_FAILED);154 return ge::GRAPH_FAILED);
177 auto transA_str = inputParams_.transA ? "true" : "false";155 auto transA_str = inputParams_.transA ? "true" : "false";
178 auto transB_str = inputParams_.transB ? "true" : "false";156 auto transB_str = inputParams_.transB ? "true" : "false";
@@ -226,6 +204,43 @@ ge::graphStatus QuantBatchMatmulV4TilingBase::CheckContext() const
226 return ge::GRAPH_SUCCESS;204 return ge::GRAPH_SUCCESS;
227}205}
228 206 
207+ge::graphStatus QuantBatchMatmulV4TilingBase::CheckInputParams() const
208+{
209+ bool maxDimCheck = inputParams_.kSize > MAX_SHAPE_DIM || inputParams_.nSize > MAX_SHAPE_DIM;
210+ if (inputParams_.transA) {
211+ maxDimCheck |= inputParams_.mSize > MAX_SHAPE_DIM;
212+ }
213+ OP_TILING_CHECK(inputParams_.supportL0c2Out && maxDimCheck,
214+ VECTOR_INNER_ERR_REPORT_TILIING(inputParams_.opName,
215+ "only support MKN in range [1, %lu], get actual value[%lu, %lu, %lu]",
216+ MAX_SHAPE_DIM, inputParams_.mSize, inputParams_.kSize, inputParams_.nSize),
217+ return ge::GRAPH_FAILED);
218+ if (inputParams_.antiQuantType == QuantType::MX) {
219+ OP_TILING_CHECK(inputParams_.groupSize != MX_GROUP_SIZE,
220+ VECTOR_INNER_ERR_REPORT_TILIING(inputParams_.opName,
221+ "Group size must be 32 for MX quantization, get group size[%lu]", inputParams_.groupSize),
222+ return ge::GRAPH_FAILED);
223+ } else {
224+ OP_TILING_CHECK(inputParams_.groupSize > inputParams_.kSize || inputParams_.groupSize % MIN_GROUP_SIZE != 0,
225+ VECTOR_INNER_ERR_REPORT_TILIING(inputParams_.opName,
226+ "Only support group size greater than %lu, less than K and align to %lu, get K[%lu] and group size[%lu]",
227+ MIN_GROUP_SIZE, MIN_GROUP_SIZE, inputParams_.kSize, inputParams_.groupSize),
228+ return ge::GRAPH_FAILED);
229+ }
230+ OP_TILING_CHECK(inputParams_.supportL0c2Out && !inputParams_.supportL12BtBf16 &&
231+ inputParams_.bFormat == ge::FORMAT_FRACTAL_NZ &&
232+ (inputParams_.bDtype != ge::DT_INT8 || inputParams_.antiQuantType != QuantType::PER_CHANNEL ||
233+ inputParams_.cDtype == ge::DT_INT8),
234+ VECTOR_INNER_ERR_REPORT_TILIING(inputParams_.opName,
235+ "weight Nz only support weight dtype INT8 per-channel scene, and not support quant scale, "
236+ "current input bDtype[%s], antiquantType[%d], cDtype[%s]",
237+ ge::TypeUtils::DataTypeToSerialString(inputParams_.bDtype).c_str(),
238+ static_cast<int>(inputParams_.antiQuantType),
239+ ge::TypeUtils::DataTypeToSerialString(inputParams_.bDtype).c_str()),
240+ return ge::GRAPH_FAILED);
241+ return ge::GRAPH_SUCCESS;
242+}
243+ 
229bool QuantBatchMatmulV4TilingBase::AnalyzeDtype()244bool QuantBatchMatmulV4TilingBase::AnalyzeDtype()
230{245{
231 inputParams_.aDtype = context_->GetInputDesc(X1_INDEX)->GetDataType();246 inputParams_.aDtype = context_->GetInputDesc(X1_INDEX)->GetDataType();
Mmatmul/quant_batch_matmul_v4/op_host/op_tiling/arch35/quant_batch_matmul_v4_tiling.h+4-1
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.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 of3 * 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").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.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -64,12 +64,14 @@ constexpr size_t VALID_BIAS_MAX_DIM = 2;
64constexpr uint64_t VEC_INNER_AXIS_ALIGN_UINT = 128UL;64constexpr uint64_t VEC_INNER_AXIS_ALIGN_UINT = 128UL;
65constexpr uint64_t MAX_SHAPE_DIM = 0x7fffffffUL;65constexpr uint64_t MAX_SHAPE_DIM = 0x7fffffffUL;
66constexpr uint64_t MIN_GROUP_SIZE = 32UL;66constexpr uint64_t MIN_GROUP_SIZE = 32UL;
67+constexpr uint64_t MX_GROUP_SIZE = 32UL;
67constexpr int32_t BASIC_PRIORITY = 1;68constexpr int32_t BASIC_PRIORITY = 1;
68constexpr uint64_t INT4_DTYPE_PARAM = 2;69constexpr uint64_t INT4_DTYPE_PARAM = 2;
69constexpr uint32_t WORKSPACE_SIZE = 16777216; // 16 * 1024 * 102470constexpr uint32_t WORKSPACE_SIZE = 16777216; // 16 * 1024 * 1024
70constexpr int32_t DB_BUFFER = 2;71constexpr int32_t DB_BUFFER = 2;
71constexpr int32_t EXTRA_GROUP_NUM = 2;72constexpr int32_t EXTRA_GROUP_NUM = 2;
72constexpr uint64_t K_ALIGN_SIZE = 32;73constexpr uint64_t K_ALIGN_SIZE = 32;
74+constexpr uint64_t K_ALIGN_SIZE_MX = 8;
73constexpr uint64_t N_ALIGN_SIZE = 8;75constexpr uint64_t N_ALIGN_SIZE = 8;
74 76 
75constexpr int64_t B64_BITS = 64;77constexpr int64_t B64_BITS = 64;
@@ -185,6 +187,7 @@ protected:
185 void Reset();187 void Reset();
186 void InitCompileInfo();188 void InitCompileInfo();
187 ge::graphStatus CheckContext() const;189 ge::graphStatus CheckContext() const;
190+ ge::graphStatus CheckInputParams() const;
188 bool AnalyzeDtype();191 bool AnalyzeDtype();
189 bool AnalyzeBiasDtype(const gert::CompileTimeTensorDesc *biasDesc);192 bool AnalyzeBiasDtype(const gert::CompileTimeTensorDesc *biasDesc);
190 bool AnalyzeX1scaleDtype(const gert::CompileTimeTensorDesc *x1ScaleDesc);193 bool AnalyzeX1scaleDtype(const gert::CompileTimeTensorDesc *x1ScaleDesc);
Mmatmul/quant_batch_matmul_v4/op_kernel/arch35/quant_batch_matmul_v4.cpp+1-1
@@ -175,7 +175,7 @@ __global__ __aicore__ void quant_batch_matmul_v4(
175 tPipe.Destroy();175 tPipe.Destroy();
176 }176 }
177#else177#else
178- REGISTER_TILING_DEFAULT(DequantBmm::QuantBatchMatmulV3TilingDataParams);178+ REGISTER_TILING_DEFAULT(qbmmv4_tiling::QuantBatchMatmulV4TilingDataParams);
179 if (QUANT_TYPE == QBMMV4_PER_GROUP) {179 if (QUANT_TYPE == QBMMV4_PER_GROUP) {
180 constexpr bool isTransA = TRANS == QBMMV4_A_TRANS || TRANS == QBMMV4_ALL_TRANS;180 constexpr bool isTransA = TRANS == QBMMV4_A_TRANS || TRANS == QBMMV4_ALL_TRANS;
181 constexpr bool isTransB = TRANS == QBMMV4_B_TRANS || TRANS == QBMMV4_ALL_TRANS;181 constexpr bool isTransB = TRANS == QBMMV4_B_TRANS || TRANS == QBMMV4_ALL_TRANS;
Mmatmul/quant_batch_matmul_v4/op_kernel/arch35/quant_batch_matmul_v4_reg_base_common.h+6-3
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.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 of3 * 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").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.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -628,8 +628,11 @@ QuantBatchMatmulV4RegBaseCommonKernel<xType, wType, biasType, yType, aTrans, bTr
628 ubOffset = ubBufIdx_ * (vecScaleOffsetLen_ / sizeof(scaleType) / tiling_->BL1Pingpong);628 ubOffset = ubBufIdx_ * (vecScaleOffsetLen_ / sizeof(scaleType) / tiling_->BL1Pingpong);
629 }629 }
630 } else { // B 矩阵非转置630 } else { // B 矩阵非转置
631- intriParams.blockLen = bubNLen * sizeof(scaleType);631+ int32_t bubNLenReal = (bubNOffset + bubNLen) > tiling_->nSize ? tiling_->nSize - bubNOffset : bubNLen;
632- intriParams.srcStride = (tiling_->nSize - bubNLen) * sizeof(scaleType);632+ intriParams.blockLen = bubNLenReal * sizeof(scaleType);
633+ intriParams.srcStride = (tiling_->nSize - bubNLenReal) * sizeof(scaleType);
634+ intriParams.dstStride = (bubNLen * sizeof(scaleType) - CeilAlign(intriParams.blockLen, ALIGNED_32_SIZE))
635+ / ALIGNED_32_SIZE;
633 if constexpr (antiQuantType == QuantType::PER_GROUP) {636 if constexpr (antiQuantType == QuantType::PER_GROUP) {
634 // k_offset + n_offset637 // k_offset + n_offset
635 gmOffset = bubKOffset / tiling_->groupSize * tiling_->nSize + bubNOffset;638 gmOffset = bubKOffset / tiling_->groupSize * tiling_->nSize + bubNOffset;
Mmatmul/quant_batch_matmul_v4/op_kernel/arch35/quant_batch_matmul_v4_tiling_key.h+8-4
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.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 of3 * 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").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.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -16,6 +16,7 @@
16#pragma once16#pragma once
17 17 
18#include "ascendc/host_api/tiling/template_argument.h"18#include "ascendc/host_api/tiling/template_argument.h"
19+#include "quant_batch_matmul_v4_tiling_data_apt.h"
19 20 
20#define QBMMV4_NOT_TRANS 021#define QBMMV4_NOT_TRANS 0
21#define QBMMV4_B_TRANS 122#define QBMMV4_B_TRANS 1
@@ -61,21 +62,24 @@ ASCENDC_TPL_SEL(
61 ASCENDC_TPL_UINT_SEL(QUANT_TYPE, ASCENDC_TPL_UI_LIST, QBMMV4_PER_GROUP),62 ASCENDC_TPL_UINT_SEL(QUANT_TYPE, ASCENDC_TPL_UI_LIST, QBMMV4_PER_GROUP),
62 ASCENDC_TPL_UINT_SEL(OPTION_ATTRS, ASCENDC_TPL_UI_LIST, QBMMV4_OPTION_ATTR_NONE),63 ASCENDC_TPL_UINT_SEL(OPTION_ATTRS, ASCENDC_TPL_UI_LIST, QBMMV4_OPTION_ATTR_NONE),
63 ASCENDC_TPL_UINT_SEL(WEIGHTNZ, ASCENDC_TPL_UI_LIST, QBMMV4_NOT_WEIGHT_NZ),64 ASCENDC_TPL_UINT_SEL(WEIGHTNZ, ASCENDC_TPL_UI_LIST, QBMMV4_NOT_WEIGHT_NZ),
64- ASCENDC_TPL_UINT_SEL(KERNEL_TEMPLATE_TYPE, ASCENDC_TPL_UI_LIST, QBMMV4_BASIS), ),65+ ASCENDC_TPL_UINT_SEL(KERNEL_TEMPLATE_TYPE, ASCENDC_TPL_UI_LIST, QBMMV4_BASIS),
66+ ASCENDC_TPL_TILING_STRUCT_SEL(qbmmv4_tiling::QuantBatchMatmulV4TilingDataParams), ),
65 ASCENDC_TPL_ARGS_SEL(67 ASCENDC_TPL_ARGS_SEL(
66 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),68 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),
67 ASCENDC_TPL_UINT_SEL(TRANS, ASCENDC_TPL_UI_LIST, QBMMV4_NOT_TRANS),69 ASCENDC_TPL_UINT_SEL(TRANS, ASCENDC_TPL_UI_LIST, QBMMV4_NOT_TRANS),
68 ASCENDC_TPL_UINT_SEL(QUANT_TYPE, ASCENDC_TPL_UI_LIST, QBMMV4_PER_GROUP),70 ASCENDC_TPL_UINT_SEL(QUANT_TYPE, ASCENDC_TPL_UI_LIST, QBMMV4_PER_GROUP),
69 ASCENDC_TPL_UINT_SEL(OPTION_ATTRS, ASCENDC_TPL_UI_LIST, QBMMV4_OPTION_ATTR_NONE),71 ASCENDC_TPL_UINT_SEL(OPTION_ATTRS, ASCENDC_TPL_UI_LIST, QBMMV4_OPTION_ATTR_NONE),
70 ASCENDC_TPL_UINT_SEL(WEIGHTNZ, ASCENDC_TPL_UI_LIST, QBMMV4_IS_WEIGHT_NZ),72 ASCENDC_TPL_UINT_SEL(WEIGHTNZ, ASCENDC_TPL_UI_LIST, QBMMV4_IS_WEIGHT_NZ),
71- ASCENDC_TPL_UINT_SEL(KERNEL_TEMPLATE_TYPE, ASCENDC_TPL_UI_LIST, QBMMV4_BASIS), ),73+ ASCENDC_TPL_UINT_SEL(KERNEL_TEMPLATE_TYPE, ASCENDC_TPL_UI_LIST, QBMMV4_BASIS),
74+ ASCENDC_TPL_TILING_STRUCT_SEL(qbmmv4_tiling::QuantBatchMatmulV4TilingDataParams), ),
72 ASCENDC_TPL_ARGS_SEL(75 ASCENDC_TPL_ARGS_SEL(
73 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),76 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),
74 ASCENDC_TPL_UINT_SEL(TRANS, ASCENDC_TPL_UI_LIST, QBMMV4_B_TRANS),77 ASCENDC_TPL_UINT_SEL(TRANS, ASCENDC_TPL_UI_LIST, QBMMV4_B_TRANS),
75 ASCENDC_TPL_UINT_SEL(QUANT_TYPE, ASCENDC_TPL_UI_LIST, QBMMV4_MX),78 ASCENDC_TPL_UINT_SEL(QUANT_TYPE, ASCENDC_TPL_UI_LIST, QBMMV4_MX),
76 ASCENDC_TPL_UINT_SEL(OPTION_ATTRS, ASCENDC_TPL_UI_LIST, QBMMV4_OPTION_ATTR_NONE),79 ASCENDC_TPL_UINT_SEL(OPTION_ATTRS, ASCENDC_TPL_UI_LIST, QBMMV4_OPTION_ATTR_NONE),
77 ASCENDC_TPL_UINT_SEL(WEIGHTNZ, ASCENDC_TPL_UI_LIST, QBMMV4_NOT_WEIGHT_NZ, QBMMV4_IS_WEIGHT_NZ),80 ASCENDC_TPL_UINT_SEL(WEIGHTNZ, ASCENDC_TPL_UI_LIST, QBMMV4_NOT_WEIGHT_NZ, QBMMV4_IS_WEIGHT_NZ),
78- ASCENDC_TPL_UINT_SEL(KERNEL_TEMPLATE_TYPE, ASCENDC_TPL_UI_LIST, QBMMV4_BASIS), ),81+ ASCENDC_TPL_UINT_SEL(KERNEL_TEMPLATE_TYPE, ASCENDC_TPL_UI_LIST, QBMMV4_BASIS),
82+ ASCENDC_TPL_TILING_STRUCT_SEL(qbmmv4_tiling::QuantBatchMatmulV4TilingDataParams), ),
79 ASCENDC_TPL_ARGS_SEL(83 ASCENDC_TPL_ARGS_SEL(
80 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),84 ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2),
81 ASCENDC_TPL_UINT_SEL(85 ASCENDC_TPL_UINT_SEL(