Star908
1826
代码
代码
Issues226
Pull Requests308
Actions
项目讨论
Wiki
分析
Star908
1826
已合并
qbmm int4 非对称量化 #2281
bynshard创建于 3月3日
qbmm int4 非对称量化 #2281
已合并
bynshard创建于 3月3日
从bynshard/ops-nnqbmm-int4-kcg合入到cann/ops-nnmaster
共 27 个文件变更+2561-312
Mmatmul/quant_batch_matmul_v3/op_api/quant_matmul_checker.cpp+180-77
@@ -184,6 +184,79 @@ static inline bool IsMicroScaling(const aclTensor *x1Scale, const aclTensor *x2S
184 x2Scale->GetDataType() == op::DataType::DT_FLOAT8_E8M0;184 x2Scale->GetDataType() == op::DataType::DT_FLOAT8_E8M0;
185}185}
186 186 
187+struct GroupSizeMNK {
188+ uint64_t m = 0;
189+ uint64_t n = 0;
190+ uint64_t k = 0;
191+};
192+ 
193+static inline GroupSizeMNK DecodeGroupSizeMnk(int64_t groupSize)
194+{
195+ GroupSizeMNK groupSizeMnk;
196+ groupSizeMnk.k = static_cast<uint64_t>(groupSize) & GROUP_MNK_BIT_SIZE;
197+ groupSizeMnk.n = (static_cast<uint64_t>(groupSize) >> GROUP_N_OFFSET) & GROUP_MNK_BIT_SIZE;
198+ groupSizeMnk.m = (static_cast<uint64_t>(groupSize) >> GROUP_M_OFFSET) & GROUP_MNK_BIT_SIZE;
199+ return groupSizeMnk;
200+}
201+ 
202+static inline int64_t EncodeGroupSizeMnk(const GroupSizeMNK &groupSizeMnk)
203+{
204+ return static_cast<int64_t>((groupSizeMnk.m << GROUP_M_OFFSET) | (groupSizeMnk.n << GROUP_N_OFFSET) | groupSizeMnk.k);
205+}
206+ 
207+static inline bool CheckMxGroupSize(int64_t groupSize, const GroupSizeMNK &groupSizeMnk)
208+{
209+ if (groupSizeMnk.k != static_cast<uint64_t>(PERGROUP_GROUP_SIZE) ||
210+ groupSizeMnk.m != 1ULL || groupSizeMnk.n != 1ULL) {
211+ OP_LOGE(ACLNN_ERR_PARAM_INVALID,
212+ "Unsupported groupSize. In mx quantification, input or infered groupSize should be \
213+4295032864(for torch api, group_sizes should be [1, 1, 32]). Actual groupSize: %lu(for torch api \
214+group_sizes is [%lu, %lu, %lu]).",
215+ groupSize, groupSizeMnk.m, groupSizeMnk.n, groupSizeMnk.k);
216+ return false;
217+ }
218+ return true;
219+}
220+ 
221+static inline bool CheckA4W4PergroupNonSymmetricGroupSize(const GroupSizeMNK &groupSizeMnk)
222+{
223+ if (groupSizeMnk.k != static_cast<uint64_t>(PERGROUP_GROUPSIZEK_SIZE)) {
224+ OP_LOGE(
225+ ACLNN_ERR_PARAM_INVALID,
226+ "Unsupported groupSize. In A4W4 pertoken-pergroup non-symmetric quantification, groupSize K should be "
227+ "256 . Actual groupSize is [%lu].",
228+ groupSizeMnk.k);
229+ return false;
230+ }
231+ return true;
232+}
233+ 
234+static inline bool CheckPerblockGroupSize(const GroupSizeMNK &groupSizeMnk)
235+{
236+ if (groupSizeMnk.k != PERBLOCK_BLOCK_SIZE) {
237+ OP_LOGE(ACLNN_ERR_PARAM_INVALID,
238+ "Unsupported groupSize. When quantification mode is G-B or B-B quantification, input or infered \
239+groupSizeK(for torch api, group_size[2]) should be 128. Actual groupSizeK: %lu, groupSizeK = groupSize & 0xFFFF.",
240+ groupSizeMnk.k);
241+ return false;
242+ }
243+ if (groupSizeMnk.n != PERBLOCK_BLOCK_SIZE) {
244+ OP_LOGE(ACLNN_ERR_PARAM_INVALID,
245+ "Unsupported groupSize. When quantification mode is G-B or B-B quantification, input or infered \
246+groupSizeN(for torch api, group_size[1]) should be 128. Actual groupSizeN: %lu, groupSizeN = (groupSize >> 16) & 0xFFFF.",
247+ groupSizeMnk.n);
248+ return false;
249+ }
250+ if (groupSizeMnk.m != PERBLOCK_BLOCK_SIZE && groupSizeMnk.m != 1UL) {
251+ OP_LOGE(ACLNN_ERR_PARAM_INVALID,
252+ "Unsupported groupSize. When quantification mode is G-B or B-B quantification, input or infered \
253+groupSizeM(for torch api, group_size[0]) \
254+should be 128 or 1. Actual groupSizeM: %lu, groupSizeM = (groupSize >> 32) & 0xFFFF.", groupSizeMnk.m);
255+ return false;
256+ }
257+ return true;
258+}
259+ 
187static inline bool IsFp8Input(const aclTensor *x1, const aclTensor *x2) {260static inline bool IsFp8Input(const aclTensor *x1, const aclTensor *x2) {
188 return (std::find(EIGHT_BIT_FLOAT_INPUT_LIST.begin(), EIGHT_BIT_FLOAT_INPUT_LIST.end(), x1->GetDataType()) !=261 return (std::find(EIGHT_BIT_FLOAT_INPUT_LIST.begin(), EIGHT_BIT_FLOAT_INPUT_LIST.end(), x1->GetDataType()) !=
189 EIGHT_BIT_FLOAT_INPUT_LIST.end()) &&262 EIGHT_BIT_FLOAT_INPUT_LIST.end()) &&
@@ -511,54 +584,86 @@ the %s dimension of %s [%ld], the real groupSize in %s dimension can not be infe
511 return true;584 return true;
512}585}
513 586 
514-bool QuantMatmulChecker::InferGroupSize(int64_t &groupSize)587+bool QuantMatmulChecker::InferGroupSizeM(
588+ const aclTensor* x1, const aclTensor* x1Scale, const aclTensor* x2Scale, bool transX1, uint64_t& groupSizeM) const
515{589{
516- auto &x1 = std::get<INDEX_X1_IN_INPUT_TUPLE>(inputTensors_);590+ auto x1DimNum = x1->GetViewShape().GetDimNum();
517- auto &x2 = std::get<INDEX_X2_IN_INPUT_TUPLE>(inputTensors_);591+ auto x1ScaleDimNum = x1Scale->GetViewShape().GetDimNum();
518- auto &x1Scale = std::get<INDEX_X1_SCALE_IN_QUANT_TUPLE>(quantTensors_);592+ auto inputSizeM =
519- auto &x2Scale = std::get<INDEX_X2_SCALE_IN_QUANT_TUPLE>(quantTensors_);593+ transX1 ? x1->GetViewShape().GetDim(x1DimNum - 1) : x1->GetViewShape().GetDim(x1DimNum - PENULTIMATE_DIM);
594+ auto scaleSizeM = 0L;
595+ if (IsMicroScaling(x1Scale, x2Scale)) {
596+ scaleSizeM = x1Scale->GetViewShape().GetDim(transX1 ? 1 : 0);
597+ } else {
598+ scaleSizeM = transX1 ? x1Scale->GetViewShape().GetDim(x1ScaleDimNum - 1) :
599+ x1Scale->GetViewShape().GetDim(x1ScaleDimNum - PENULTIMATE_DIM);
600+ }
601+ return ReCalcGroupSize(inputSizeM, scaleSizeM, groupSizeM, "m");
602+}
603+ 
604+bool QuantMatmulChecker::InferGroupSizeK(
605+ const aclTensor* x1, const aclTensor* x1Scale, const aclTensor* x2Scale, bool transX1, uint64_t& groupSizeK) const
606+{
607+ auto x1DimNum = x1->GetViewShape().GetDimNum();
608+ auto x1ScaleDimNum = x1Scale->GetViewShape().GetDimNum();
609+ auto inputSizeK =
610+ transX1 ? x1->GetViewShape().GetDim(x1DimNum - PENULTIMATE_DIM) : x1->GetViewShape().GetDim(x1DimNum - 1);
611+ auto scaleSizeK = 0L;
612+ if (IsMicroScaling(x1Scale, x2Scale)) {
613+ // when scale type is e8m0, scalex1 shape is [m, k/2, 2] or [k/2, m, 2]
614+ scaleSizeK = x1Scale->GetViewShape().GetDim(transX1 ? 0 : 1) * MICRO_SCALING_ALIGN_NUM;
615+ } else {
616+ scaleSizeK = transX1 ? x1Scale->GetViewShape().GetDim(x1ScaleDimNum - PENULTIMATE_DIM) :
617+ x1Scale->GetViewShape().GetDim(x1ScaleDimNum - 1);
618+ }
619+ return ReCalcGroupSize(inputSizeK, scaleSizeK, groupSizeK, "k");
620+}
621+ 
622+bool QuantMatmulChecker::InferGroupSizeN(
623+ const aclTensor* x2, const aclTensor* x1Scale, const aclTensor* x2Scale, bool transX2, uint64_t groupSizeK,
624+ uint64_t& groupSizeN) const
625+{
626+ auto x2DimNum = x2->GetViewShape().GetDimNum();
627+ auto x2ScaleDimNum = x2Scale->GetViewShape().GetDimNum();
628+ auto inputSizeN =
629+ transX2 ? x2->GetViewShape().GetDim(x2DimNum - PENULTIMATE_DIM) : x2->GetViewShape().GetDim(x2DimNum - 1);
630+ auto scaleSizeN = 0L;
631+ if (IsMicroScaling(x1Scale, x2Scale)) {
632+ scaleSizeN = x2Scale->GetViewShape().GetDim(transX2 ? 0 : 1);
633+ } else {
634+ if (IsA4W4PergroupNonSymmetric(groupSizeK)) {
635+ scaleSizeN = transX2 ? x2Scale->GetViewShape().GetDim(x2ScaleDimNum - 1) :
636+ x2Scale->GetViewShape().GetDim(x2ScaleDimNum - PENULTIMATE_DIM);
637+ } else {
638+ scaleSizeN = transX2 ? x2Scale->GetViewShape().GetDim(x2ScaleDimNum - PENULTIMATE_DIM) :
639+ x2Scale->GetViewShape().GetDim(x2ScaleDimNum - 1);
640+ }
641+ }
642+ return ReCalcGroupSize(inputSizeN, scaleSizeN, groupSizeN, "n");
643+}
644+ 
645+bool QuantMatmulChecker::InferGroupSize(int64_t& groupSize)
646+{
647+ auto& x1 = std::get<INDEX_X1_IN_INPUT_TUPLE>(inputTensors_);
648+ auto& x2 = std::get<INDEX_X2_IN_INPUT_TUPLE>(inputTensors_);
649+ auto& x1Scale = std::get<INDEX_X1_SCALE_IN_QUANT_TUPLE>(quantTensors_);
650+ auto& x2Scale = std::get<INDEX_X2_SCALE_IN_QUANT_TUPLE>(quantTensors_);
520 // when x1Scale and x2Scale dim num is less than 2, groupsize not used651 // when x1Scale and x2Scale dim num is less than 2, groupsize not used
521 if (x1Scale == nullptr || x1Scale->GetViewShape().GetDimNum() < 2 || x2Scale->GetViewShape().GetDimNum() < 2) {652 if (x1Scale == nullptr || x1Scale->GetViewShape().GetDimNum() < 2 || x2Scale->GetViewShape().GetDimNum() < 2) {
522 return true;653 return true;
523 }654 }
524- auto x1DimNum = x1->GetViewShape().GetDimNum();
525- auto x2DimNum = x2->GetViewShape().GetDimNum();
526- auto x1ScaleDimNum = x1Scale->GetViewShape().GetDimNum();
527- auto x2ScaleDimNum = x2Scale->GetViewShape().GetDimNum();
528 auto transX1 = std::get<INDEX_X1_IN_INPUT_TUPLE>(boolsTrans_);655 auto transX1 = std::get<INDEX_X1_IN_INPUT_TUPLE>(boolsTrans_);
529 auto transX2 = std::get<INDEX_X2_IN_INPUT_TUPLE>(boolsTrans_);656 auto transX2 = std::get<INDEX_X2_IN_INPUT_TUPLE>(boolsTrans_);
530- uint64_t groupSizeK = static_cast<uint64_t>(groupSize) & GROUP_MNK_BIT_SIZE;657+ auto groupSizeMnk = DecodeGroupSizeMnk(groupSize);
531- uint64_t groupSizeN = (static_cast<uint64_t>(groupSize) >> GROUP_N_OFFSET) & GROUP_MNK_BIT_SIZE;658+ 
532- uint64_t groupSizeM = (static_cast<uint64_t>(groupSize) >> GROUP_M_OFFSET) & GROUP_MNK_BIT_SIZE;659+ CHECK_RET(InferGroupSizeM(x1, x1Scale, x2Scale, transX1, groupSizeMnk.m), false);
533- auto inputSizeM = transX1 ? x1->GetViewShape().GetDim(x1DimNum - 1) : x1->GetViewShape().GetDim(x1DimNum - PENULTIMATE_DIM);660+ CHECK_RET(InferGroupSizeK(x1, x1Scale, x2Scale, transX1, groupSizeMnk.k), false);
534- auto scaleSizeM = 0;661+ CHECK_RET(InferGroupSizeN(x2, x1Scale, x2Scale, transX2, groupSizeMnk.k, groupSizeMnk.n), false);
535- if (IsMicroScaling(x1Scale, x2Scale)) {662+ 
536- scaleSizeM = x1Scale->GetViewShape().GetDim(transX1 ? 1 : 0);663+ OP_LOGD(
537- } else {664+ "Infered groupSize: groupSizeM: %lu, groupSizeN: %lu, groupSizeK: %lu.", groupSizeMnk.m, groupSizeMnk.n,
538- scaleSizeM = transX1 ? x1Scale->GetViewShape().GetDim(x1ScaleDimNum - 1)665+ groupSizeMnk.k);
539- : x1Scale->GetViewShape().GetDim(x1ScaleDimNum - PENULTIMATE_DIM);666+ groupSize = EncodeGroupSizeMnk(groupSizeMnk);
540- }
541- CHECK_RET(ReCalcGroupSize(inputSizeM, scaleSizeM, groupSizeM, "m"), false);
542- auto inputSizeK = transX1 ? x1->GetViewShape().GetDim(x1DimNum - PENULTIMATE_DIM) : x1->GetViewShape().GetDim(x1DimNum - 1);
543- auto scaleSizeK = 0;
544- if (IsMicroScaling(x1Scale, x2Scale)) {
545- scaleSizeK = x1Scale->GetViewShape().GetDim(transX1 ? 0 : 1) * 2; //when scale type is e8m0, scalex1 shape is [m, k/2, 2] or [k/2, m, 2]
546- } else {
547- scaleSizeK = transX1 ? x1Scale->GetViewShape().GetDim(x1ScaleDimNum - PENULTIMATE_DIM)
548- : x1Scale->GetViewShape().GetDim(x1ScaleDimNum - 1);
549- }
550- CHECK_RET(ReCalcGroupSize(inputSizeK, scaleSizeK, groupSizeK, "k"), false);
551- auto inputSizeN = transX2 ? x2->GetViewShape().GetDim(x2DimNum - PENULTIMATE_DIM) : x2->GetViewShape().GetDim(x2DimNum - 1);
552- auto scaleSizeN = 0;
553- if (IsMicroScaling(x1Scale, x2Scale)) {
554- scaleSizeN = x2Scale->GetViewShape().GetDim(transX2 ? 0 : 1);
555- } else {
556- scaleSizeN = transX2 ? x2Scale->GetViewShape().GetDim(x2ScaleDimNum - PENULTIMATE_DIM)
557- : x2Scale->GetViewShape().GetDim(x2ScaleDimNum - 1);
558- }
559- CHECK_RET(ReCalcGroupSize(inputSizeN, scaleSizeN, groupSizeN, "n"), false);
560- OP_LOGD("Infered groupSize: groupSizeM: %lu, groupSizeN: %lu, groupSizeK: %lu.", groupSizeM, groupSizeN, groupSizeK);
561- groupSize = static_cast<int64_t>((groupSizeM << GROUP_M_OFFSET) | (groupSizeN << GROUP_N_OFFSET) | groupSizeK);
562 return true;667 return true;
563}668}
564 669 
@@ -877,46 +982,20 @@ bool QuantMatmulChecker::CheckGroupSize() const
877 if (npuArch_ != NpuArch::DAV_3510) {982 if (npuArch_ != NpuArch::DAV_3510) {
878 return true;983 return true;
879 }984 }
880- uint64_t groupSizeM = (static_cast<uint64_t>(groupSize_) >> GROUP_M_OFFSET) & GROUP_MNK_BIT_SIZE;985+ const GroupSizeMNK groupSizeMnk = DecodeGroupSizeMnk(groupSize_);
881- uint64_t groupSizeN = (static_cast<uint64_t>(groupSize_) >> GROUP_N_OFFSET) & GROUP_MNK_BIT_SIZE;986+ if (IsA4W4PergroupNonSymmetric(groupSizeMnk.k)) {
882- uint64_t groupSizeK = static_cast<uint64_t>(groupSize_) & GROUP_MNK_BIT_SIZE;987+ CHECK_RET(CheckA4W4PergroupNonSymmetricGroupSize(groupSizeMnk), false);
883- if (IsMicroScaling(x1Scale_, x2Scale_)) {988+ } else if (IsMicroScaling(x1Scale_, x2Scale_)) {
884- if (groupSizeK != static_cast<uint64_t>(PERGROUP_GROUP_SIZE) || groupSizeM != 1ULL || groupSizeN != 1ULL) {989+ CHECK_RET(CheckMxGroupSize(groupSize_, groupSizeMnk), false);
885- OP_LOGE(ACLNN_ERR_PARAM_INVALID,
886- "Unsupported groupSize. In mx quantification, input or infered groupSize should be \
887-4295032864(for torch api, group_sizes should be [1, 1, 32]). Actual groupSize: %lu(for torch api \
888-group_sizes is [%lu, %lu, %lu]).",
889- groupSize_, groupSizeM, groupSizeN, groupSizeK);
890- return false;
891- }
892 } else if (IsPerblock(x1_, x2_, x1Scale_, x2Scale_)) {990 } else if (IsPerblock(x1_, x2_, x1Scale_, x2Scale_)) {
893- if (groupSizeK != PERBLOCK_BLOCK_SIZE) {991+ CHECK_RET(CheckPerblockGroupSize(groupSizeMnk), false);
894- OP_LOGE(ACLNN_ERR_PARAM_INVALID,
895- "Unsupported groupSize. When quantification mode is G-B or B-B quantification, input or infered \
896-groupSizeK(for torch api, group_size[2]) should be 128. Actual groupSizeK: %lu, groupSizeK = groupSize & 0xFFFF.",
897- groupSizeK);
898- return false;
899- }
900- if (groupSizeN != PERBLOCK_BLOCK_SIZE) {
901- OP_LOGE(ACLNN_ERR_PARAM_INVALID,
902- "Unsupported groupSize. When quantification mode is G-B or B-B quantification, input or infered \
903-groupSizeN(for torch api, group_size[1]) should be 128. Actual groupSizeN: %lu, groupSizeN = (groupSize >> 16) & 0xFFFF.",
904- groupSizeN);
905- return false;
906- }
907- if (groupSizeM != PERBLOCK_BLOCK_SIZE && groupSizeM != 1UL) {
908- OP_LOGE(ACLNN_ERR_PARAM_INVALID,
909- "Unsupported groupSize. When quantification mode is G-B or B-B quantification, input or infered \
910-groupSizeM(for torch api, group_size[0]) \
911-should be 128 or 1. Actual groupSizeM: %lu, groupSizeM = (groupSize >> 32) & 0xFFFF.", groupSizeM);
912- return false;
913- }
914 } else if (groupSize_ != 0UL) {992 } else if (groupSize_ != 0UL) {
915- OP_LOGE(ACLNN_ERR_PARAM_INVALID,993+ OP_LOGE(
916- "Unsupported groupSize. When quantification mode is not G-B or B-B or mx quantification, \994+ ACLNN_ERR_PARAM_INVALID,
995+ "Unsupported groupSize. When quantification mode is not G-B or B-B or mx quantification, \
917groupSize should be 0(torch api group_sizes should be [0, 0, 0] or None). \996groupSize should be 0(torch api group_sizes should be [0, 0, 0] or None). \
918Actual groupSize: %lu(torch api group_sizes is [%lu, %lu, %lu]).",997Actual groupSize: %lu(torch api group_sizes is [%lu, %lu, %lu]).",
919- groupSize_, groupSizeM, groupSizeN, groupSizeK);998+ groupSize_, groupSizeMnk.m, groupSizeMnk.n, groupSizeMnk.k);
920 return false;999 return false;
921 }1000 }
922 OP_LOGD("QuantMatmul check group_size success.");1001 OP_LOGD("QuantMatmul check group_size success.");
@@ -1348,6 +1427,28 @@ bool QuantMatmulChecker::CheckL0c2outOrL0c2ubPertoken() const
1348 return true;1427 return true;
1349}1428}
1350 1429 
1430+bool QuantMatmulChecker::CheckL0C2outOrL0C2ubPertokenPergroup() const {
1431+ if (x1_ == nullptr || x2_ == nullptr || out_ == nullptr ||
1432+ x1Scale_ == nullptr || x2Scale_ == nullptr || x2Offset_ == nullptr) {
1433+ return false;
1434+ }
1435+ CHECK_RET(OpCheckDtypeNotSupport(interfaceType_, X1_NAME, x1_, {op::DataType::DT_INT8, op::DataType::DT_INT4}), false);
1436+ CHECK_RET(OpCheckDtypeNotSupport(interfaceType_, X2_NAME, x2_, {op::DataType::DT_INT8, op::DataType::DT_INT4}), false);
1437+ CHECK_RET(OpCheckDtypeNotSupport(interfaceType_, OUT_NAME, out_, {op::DataType::DT_FLOAT16, op::DataType::DT_BF16}), false);
1438+ CHECK_RET(OpCheckDtypeNotSupport(interfaceType_, X1SCALE_NAME, x1Scale_, {op::DataType::DT_FLOAT}), false);
1439+ CHECK_RET(OpCheckDtypeNotSupport(interfaceType_, X2SCALE_NAME, x2Scale_, {op::DataType::DT_FLOAT}), false);
1440+ CHECK_RET(OpCheckDtypeNotSupport(interfaceType_, X2OFFSET_NAME, x2Offset_, {op::DataType::DT_FLOAT16}), false);
1441+ 
1442+ if (transposeX1_ != false || transposeX2_ != true) {
1443+ OP_LOGE(
1444+ ACLNN_ERR_PARAM_INVALID,
1445+ "In a4w4 scenario, only support transA=false and transB=true, but got transA=%s, transB=%s.",
1446+ transposeX1_ ? "true" : "false", transposeX2_ ? "true" : "false");
1447+ return false;
1448+ }
1449+ return true;
1450+}
1451+ 
1351bool QuantMatmulChecker::CheckDoubleScaleAndFp8Hif8PertokenPerblock() const1452bool QuantMatmulChecker::CheckDoubleScaleAndFp8Hif8PertokenPerblock() const
1352{1453{
1353 CHECK_RET(OpCheckDtypeNotMatch(interfaceType_, X1SCALE_NAME, x1Scale_, op::DataType::DT_FLOAT), false);1454 CHECK_RET(OpCheckDtypeNotMatch(interfaceType_, X1SCALE_NAME, x1Scale_, op::DataType::DT_FLOAT), false);
@@ -1390,6 +1491,8 @@ aclnnStatus QuantMatmulChecker::CheckDtypeL0c2outOrL0c2ub() const
1390 CHECK_RET(CheckL0c2outOrL0c2ubPertensorPerchannel(), ACLNN_ERR_PARAM_INVALID);1491 CHECK_RET(CheckL0c2outOrL0c2ubPertensorPerchannel(), ACLNN_ERR_PARAM_INVALID);
1391 } else if (IsMicroScaling(x1Scale_, x2Scale_)) { // micro scaling1492 } else if (IsMicroScaling(x1Scale_, x2Scale_)) { // micro scaling
1392 CHECK_RET(CheckMicroScaling(), ACLNN_ERR_PARAM_INVALID);1493 CHECK_RET(CheckMicroScaling(), ACLNN_ERR_PARAM_INVALID);
1494+ } else if (IsInt4Input(x1_, x2_) && x2Offset_ != nullptr) { // pertoken
1495+ CHECK_RET(CheckL0C2outOrL0C2ubPertokenPergroup(), ACLNN_ERR_PARAM_INVALID);
1393 } else if (IsInt8Input(x1_, x2_) || IsInt4Input(x1_, x2_)) { // pertoken1496 } else if (IsInt8Input(x1_, x2_) || IsInt4Input(x1_, x2_)) { // pertoken
1394 CHECK_RET(CheckL0c2outOrL0c2ubPertoken(), ACLNN_ERR_PARAM_INVALID);1497 CHECK_RET(CheckL0c2outOrL0c2ubPertoken(), ACLNN_ERR_PARAM_INVALID);
1395 } else if (IsHif8Input(x1_, x2_) ||1498 } else if (IsHif8Input(x1_, x2_) ||
Mmatmul/quant_batch_matmul_v3/op_api/quant_matmul_checker.h+8-1
@@ -44,6 +44,7 @@ private:
44 aclnnStatus CheckDtypeL0c2outOrL0c2ub() const;44 aclnnStatus CheckDtypeL0c2outOrL0c2ub() const;
45 bool CheckDoubleScaleAndFp8Hif8PertokenPerblock() const;45 bool CheckDoubleScaleAndFp8Hif8PertokenPerblock() const;
46 bool CheckL0c2outOrL0c2ubPertoken() const;46 bool CheckL0c2outOrL0c2ubPertoken() const;
47+ bool CheckL0C2outOrL0C2ubPertokenPergroup() const;
47 bool CheckMicroScaling() const;48 bool CheckMicroScaling() const;
48 bool CheckL0c2outOrL0c2ubPertensorPerchannel() const;49 bool CheckL0c2outOrL0c2ubPertensorPerchannel() const;
49 bool CheckL0c2outOrL0c2ubPertensorPerchannel4Int8Input() const;50 bool CheckL0c2outOrL0c2ubPertensorPerchannel4Int8Input() const;
@@ -74,6 +75,12 @@ private:
74 std::string GetX1ScaleName() const;75 std::string GetX1ScaleName() const;
75 std::string GetX2ScaleName() const;76 std::string GetX2ScaleName() const;
76 std::string GetX2OffsetName() const;77 std::string GetX2OffsetName() const;
78+ bool InferGroupSizeM(const aclTensor *x1, const aclTensor *x1Scale, const aclTensor *x2Scale,
79+ bool transX1, uint64_t &groupSizeM) const;
80+ bool InferGroupSizeK(const aclTensor *x1, const aclTensor *x1Scale, const aclTensor *x2Scale,
81+ bool transX1, uint64_t &groupSizeK) const;
82+ bool InferGroupSizeN(const aclTensor *x2, const aclTensor *x1Scale, const aclTensor *x2Scale,
83+ bool transX2, uint64_t groupSizeK, uint64_t &groupSizeN) const;
77 bool ReCalcGroupSize(int64_t inputSize, int64_t scaleSize, uint64_t &groupSize, const char *dimensionName) const;84 bool ReCalcGroupSize(int64_t inputSize, int64_t scaleSize, uint64_t &groupSize, const char *dimensionName) const;
78 85 
79public:86public:
@@ -109,4 +116,4 @@ private:
109};116};
110}117}
111 118 
112-#endif // OP_API_SRC_QUANT_MATMUL_CHECKER_H_119+#endif // OP_API_SRC_QUANT_MATMUL_CHECKER_H_
Mmatmul/quant_batch_matmul_v3/op_kernel/arch35/qbmm_int4_to_int8_preprocess.h+178-178
@@ -1,179 +1,179 @@
1-/**1+/**
2- * Copyright (c) 2026 Huawei Technologies Co., Ltd.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 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.
6- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,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.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- 10+ 
11-/*!11+/*!
12- * \file qbmm_int4_to_int8_preprocess.h12+ * \file qbmm_int4_to_int8_preprocess.h
13- * \brief Preprocess class for converting int4 inputs (x1, x2) to int8 before matrix multiplication.13+ * \brief Preprocess class for converting int4 inputs (x1, x2) to int8 before matrix multiplication.
14- * Core splitting: cores are split between x1 and x2 by m:n ratio.14+ * Core splitting: cores are split between x1 and x2 by m:n ratio.
15- * Each core processes ONLY x1 or x2.15+ * Each core processes ONLY x1 or x2.
16- * Pipeline: CopyIn -> Compute (int4->half->int8) -> CopyOut16+ * Pipeline: CopyIn -> Compute (int4->half->int8) -> CopyOut
17- */17+ */
18- 18+ 
19-#ifndef QBMM_INT4_TO_INT8_PREPROCESS_H19+#ifndef QBMM_INT4_TO_INT8_PREPROCESS_H
20-#define QBMM_INT4_TO_INT8_PREPROCESS_H20+#define QBMM_INT4_TO_INT8_PREPROCESS_H
21- 21+ 
22-#include "../quant_batch_matmul_v3_base.h"22+#include "../quant_batch_matmul_v3_base.h"
23-#include "quant_batch_matmul_v3_tiling_data.h"23+#include "quant_batch_matmul_v3_tiling_data.h"
24- 24+ 
25-using namespace AscendC;25+using namespace AscendC;
26- 26+ 
27-namespace {27+namespace {
28-constexpr uint64_t ALIGN_SIZE_128 = 128u;28+constexpr uint64_t ALIGN_SIZE_128 = 128u;
29-constexpr uint64_t TILE_ELEMS_16K = 16 * 1024u;29+constexpr uint64_t TILE_ELEMS_16K = 16 * 1024u;
30-constexpr uint32_t ELEM_ALIGN_64 = 64u;30+constexpr uint32_t ELEM_ALIGN_64 = 64u;
31-constexpr uint64_t NUM_2 = 2u;31+constexpr uint64_t NUM_2 = 2u;
32-}32+}
33- 33+ 
34-class QbmmInt4ToInt8Preprocess {34+class QbmmInt4ToInt8Preprocess {
35-public:35+public:
36- __aicore__ inline QbmmInt4ToInt8Preprocess() {}36+ __aicore__ inline QbmmInt4ToInt8Preprocess() {}
37- 37+ 
38- __aicore__ inline void Init(GM_ADDR x1In, GM_ADDR x2In, GM_ADDR workspace,38+ __aicore__ inline void Init(GM_ADDR x1In, GM_ADDR x2In, GM_ADDR workspace,
39- TPipe& pipe, uint64_t m, uint64_t n,39+ TPipe& pipe, uint64_t m, uint64_t n,
40- uint64_t k, uint64_t batchC);40+ uint64_t k, uint64_t batchC);
41- __aicore__ inline void Process();41+ __aicore__ inline void Process();
42- 42+ 
43-private:43+private:
44- __aicore__ inline void CopyIn(uint32_t progress, uint32_t currentNum);44+ __aicore__ inline void CopyIn(uint32_t progress, uint32_t currentNum);
45- __aicore__ inline void Compute(uint32_t currentNum);45+ __aicore__ inline void Compute(uint32_t currentNum);
46- __aicore__ inline void CopyOut(uint32_t progress, uint32_t currentNum);46+ __aicore__ inline void CopyOut(uint32_t progress, uint32_t currentNum);
47- 47+ 
48-private:48+private:
49- // Queues49+ // Queues
50- TQue<QuePosition::VECIN, BUFFER_NUM> inQueueInt4_;50+ TQue<QuePosition::VECIN, BUFFER_NUM> inQueueInt4_;
51- TQue<QuePosition::VECCALC, BUFFER_NUM> computeQueueHalf_;51+ TQue<QuePosition::VECCALC, BUFFER_NUM> computeQueueHalf_;
52- TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueInt8_;52+ TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueInt8_;
53- 53+ 
54- GlobalTensor<int8_t> srcInt4Global_;54+ GlobalTensor<int8_t> srcInt4Global_;
55- GlobalTensor<int8_t> dstInt8Global_;55+ GlobalTensor<int8_t> dstInt8Global_;
56- 56+ 
57- GM_ADDR x1Out_;57+ GM_ADDR x1Out_;
58- GM_ADDR x2Out_;58+ GM_ADDR x2Out_;
59- 59+ 
60- bool isX1Core_;60+ bool isX1Core_;
61- uint64_t blockLength_ = 0;61+ uint64_t blockLength_ = 0;
62- uint32_t ubLength_ = 0;62+ uint32_t ubLength_ = 0;
63-};63+};
64- 64+ 
65-__aicore__ inline void QbmmInt4ToInt8Preprocess::Init(GM_ADDR x1In, GM_ADDR x2In, GM_ADDR workspace,65+__aicore__ inline void QbmmInt4ToInt8Preprocess::Init(GM_ADDR x1In, GM_ADDR x2In, GM_ADDR workspace,
66- TPipe& pipe, uint64_t m, uint64_t n, uint64_t k,66+ TPipe& pipe, uint64_t m, uint64_t n, uint64_t k,
67- uint64_t batchC)67+ uint64_t batchC)
68-{68+{
69- uint64_t x1TotalElems = batchC * m * k;69+ uint64_t x1TotalElems = batchC * m * k;
70- uint64_t x2TotalElems = k * n;70+ uint64_t x2TotalElems = k * n;
71- 71+ 
72- x1Out_ = workspace;72+ x1Out_ = workspace;
73- x2Out_ = workspace + DequantBmm::Align(x1TotalElems * sizeof(int8_t), ALIGN_SIZE_128);73+ x2Out_ = workspace + DequantBmm::Align(x1TotalElems * sizeof(int8_t), ALIGN_SIZE_128);
74- 74+ 
75- // ---- core assignment by m:n ratio ----75+ // ---- core assignment by m:n ratio ----
76- uint64_t totalCores = GetBlockNum();76+ uint64_t totalCores = GetBlockNum();
77- if ASCEND_IS_AIV {77+ if ASCEND_IS_AIV {
78- totalCores = totalCores * NUM_2;78+ totalCores = totalCores * NUM_2;
79- }79+ }
80- uint64_t coresForX1 = (totalCores * batchC * m + (batchC * m + n) / NUM_2) / (batchC * m + n);80+ uint64_t coresForX1 = (totalCores * batchC * m + (batchC * m + n) / NUM_2) / (batchC * m + n);
81- if (coresForX1 < 1) coresForX1 = 1;81+ if (coresForX1 < 1) coresForX1 = 1;
82- if (coresForX1 >= totalCores) coresForX1 = totalCores - 1;82+ if (coresForX1 >= totalCores) coresForX1 = totalCores - 1;
83- uint64_t coresForX2 = totalCores - coresForX1;83+ uint64_t coresForX2 = totalCores - coresForX1;
84- 84+ 
85- uint64_t coreIdx = GetBlockIdx();85+ uint64_t coreIdx = GetBlockIdx();
86- if (coreIdx >= totalCores) {86+ if (coreIdx >= totalCores) {
87- return;87+ return;
88- }88+ }
89- isX1Core_ = (coreIdx < coresForX1);89+ isX1Core_ = (coreIdx < coresForX1);
90- 90+ 
91- // ---- compute this core's element range ----91+ // ---- compute this core's element range ----
92- uint64_t totalElems = isX1Core_ ? x1TotalElems : x2TotalElems;92+ uint64_t totalElems = isX1Core_ ? x1TotalElems : x2TotalElems;
93- uint64_t groupCores = isX1Core_ ? coresForX1 : coresForX2;93+ uint64_t groupCores = isX1Core_ ? coresForX1 : coresForX2;
94- uint64_t localId = isX1Core_ ? coreIdx : coreIdx - coresForX1;94+ uint64_t localId = isX1Core_ ? coreIdx : coreIdx - coresForX1;
95- 95+ 
96- uint64_t totalBlocks = DequantBmm::Align(totalElems, static_cast<uint64_t>(ELEM_ALIGN_64));96+ uint64_t totalBlocks = DequantBmm::Align(totalElems, static_cast<uint64_t>(ELEM_ALIGN_64));
97- uint64_t numChunks = totalBlocks / ELEM_ALIGN_64;97+ uint64_t numChunks = totalBlocks / ELEM_ALIGN_64;
98- uint64_t baseChunks = numChunks / groupCores;98+ uint64_t baseChunks = numChunks / groupCores;
99- uint64_t remainChunks = numChunks % groupCores;99+ uint64_t remainChunks = numChunks % groupCores;
100- 100+ 
101- uint64_t chunkStart = localId * baseChunks + DequantBmm::Min(localId, remainChunks);101+ uint64_t chunkStart = localId * baseChunks + DequantBmm::Min(localId, remainChunks);
102- uint64_t chunkEnd = chunkStart + baseChunks + (localId < remainChunks ? 1u : 0u);102+ uint64_t chunkEnd = chunkStart + baseChunks + (localId < remainChunks ? 1u : 0u);
103- 103+ 
104- uint64_t elemStart = chunkStart * ELEM_ALIGN_64;104+ uint64_t elemStart = chunkStart * ELEM_ALIGN_64;
105- uint64_t elemEnd = chunkEnd * ELEM_ALIGN_64;105+ uint64_t elemEnd = chunkEnd * ELEM_ALIGN_64;
106- 106+ 
107- blockLength_ = elemEnd - elemStart;107+ blockLength_ = elemEnd - elemStart;
108- 108+ 
109- uint64_t tileSize = DequantBmm::Min(blockLength_, TILE_ELEMS_16K);109+ uint64_t tileSize = DequantBmm::Min(blockLength_, TILE_ELEMS_16K);
110- uint64_t alignUb = DequantBmm::Align(tileSize, static_cast<uint64_t>(ELEM_ALIGN_64));110+ uint64_t alignUb = DequantBmm::Align(tileSize, static_cast<uint64_t>(ELEM_ALIGN_64));
111- ubLength_ = DequantBmm::Max(ELEM_ALIGN_64, static_cast<uint32_t>(alignUb));111+ ubLength_ = DequantBmm::Max(ELEM_ALIGN_64, static_cast<uint32_t>(alignUb));
112- 112+ 
113- // ---- bind global buffers to this core's slice ----113+ // ---- bind global buffers to this core's slice ----
114- if (isX1Core_) {114+ if (isX1Core_) {
115- srcInt4Global_.SetGlobalBuffer((__gm__ int8_t*)x1In + elemStart / NUM_2, blockLength_ / NUM_2);115+ srcInt4Global_.SetGlobalBuffer((__gm__ int8_t*)x1In + elemStart / NUM_2, blockLength_ / NUM_2);
116- dstInt8Global_.SetGlobalBuffer((__gm__ int8_t*)x1Out_ + elemStart, blockLength_);116+ dstInt8Global_.SetGlobalBuffer((__gm__ int8_t*)x1Out_ + elemStart, blockLength_);
117- } else {117+ } else {
118- srcInt4Global_.SetGlobalBuffer((__gm__ int8_t*)x2In + elemStart / NUM_2, blockLength_ / NUM_2);118+ srcInt4Global_.SetGlobalBuffer((__gm__ int8_t*)x2In + elemStart / NUM_2, blockLength_ / NUM_2);
119- dstInt8Global_.SetGlobalBuffer((__gm__ int8_t*)x2Out_ + elemStart, blockLength_);119+ dstInt8Global_.SetGlobalBuffer((__gm__ int8_t*)x2Out_ + elemStart, blockLength_);
120- }120+ }
121- 121+ 
122- // ---- init queues ----122+ // ---- init queues ----
123- pipe.InitBuffer(inQueueInt4_, BUFFER_NUM, static_cast<uint32_t>(ubLength_) / NUM_2);123+ pipe.InitBuffer(inQueueInt4_, BUFFER_NUM, static_cast<uint32_t>(ubLength_) / NUM_2);
124- pipe.InitBuffer(computeQueueHalf_, BUFFER_NUM, static_cast<uint32_t>(ubLength_) * sizeof(half));124+ pipe.InitBuffer(computeQueueHalf_, BUFFER_NUM, static_cast<uint32_t>(ubLength_) * sizeof(half));
125- pipe.InitBuffer(outQueueInt8_, BUFFER_NUM, static_cast<uint32_t>(ubLength_) * sizeof(int8_t));125+ pipe.InitBuffer(outQueueInt8_, BUFFER_NUM, static_cast<uint32_t>(ubLength_) * sizeof(int8_t));
126-}126+}
127- 127+ 
128-__aicore__ inline void QbmmInt4ToInt8Preprocess::Process()128+__aicore__ inline void QbmmInt4ToInt8Preprocess::Process()
129-{129+{
130- if (blockLength_ == 0) return;130+ if (blockLength_ == 0) return;
131- 131+ 
132- uint64_t loopCount = DequantBmm::CeilDiv(blockLength_, static_cast<uint64_t>(ubLength_));132+ uint64_t loopCount = DequantBmm::CeilDiv(blockLength_, static_cast<uint64_t>(ubLength_));
133- for (uint32_t i = 0; i < loopCount; i++) {133+ for (uint32_t i = 0; i < loopCount; i++) {
134- uint64_t remaining = blockLength_ - ubLength_ * i;134+ uint64_t remaining = blockLength_ - ubLength_ * i;
135- uint32_t currentNum = DequantBmm::Min(static_cast<uint32_t>(remaining), ubLength_);135+ uint32_t currentNum = DequantBmm::Min(static_cast<uint32_t>(remaining), ubLength_);
136- 136+ 
137- currentNum = DequantBmm::FloorAlign(currentNum, ELEM_ALIGN_64);137+ currentNum = DequantBmm::FloorAlign(currentNum, ELEM_ALIGN_64);
138- if (currentNum == 0) break;138+ if (currentNum == 0) break;
139- 139+ 
140- CopyIn(i, currentNum);140+ CopyIn(i, currentNum);
141- Compute(currentNum);141+ Compute(currentNum);
142- CopyOut(i, currentNum);142+ CopyOut(i, currentNum);
143- }143+ }
144-}144+}
145- 145+ 
146-__aicore__ inline void QbmmInt4ToInt8Preprocess::CopyIn(uint32_t progress, uint32_t currentNum)146+__aicore__ inline void QbmmInt4ToInt8Preprocess::CopyIn(uint32_t progress, uint32_t currentNum)
147-{147+{
148- LocalTensor<int8_t> int4Local = inQueueInt4_.AllocTensor<int8_t>();148+ LocalTensor<int8_t> int4Local = inQueueInt4_.AllocTensor<int8_t>();
149- DataCopy(int4Local, srcInt4Global_[progress * ubLength_ / NUM_2], static_cast<uint32_t>(currentNum) / NUM_2);149+ DataCopy(int4Local, srcInt4Global_[progress * ubLength_ / NUM_2], static_cast<uint32_t>(currentNum) / NUM_2);
150- inQueueInt4_.EnQue<int8_t>(int4Local);150+ inQueueInt4_.EnQue<int8_t>(int4Local);
151-}151+}
152- 152+ 
153-__aicore__ inline void QbmmInt4ToInt8Preprocess::Compute(uint32_t currentNum)153+__aicore__ inline void QbmmInt4ToInt8Preprocess::Compute(uint32_t currentNum)
154-{154+{
155- LocalTensor<int8_t> int4Local = inQueueInt4_.DeQue<int8_t>();155+ LocalTensor<int8_t> int4Local = inQueueInt4_.DeQue<int8_t>();
156- LocalTensor<int4b_t> int4View = int4Local.ReinterpretCast<int4b_t>();156+ LocalTensor<int4b_t> int4View = int4Local.ReinterpretCast<int4b_t>();
157- 157+ 
158- // ---- int4 -> half ----158+ // ---- int4 -> half ----
159- LocalTensor<half> halfLocal = computeQueueHalf_.AllocTensor<half>();159+ LocalTensor<half> halfLocal = computeQueueHalf_.AllocTensor<half>();
160- Cast<half, int4b_t>(halfLocal, int4View, RoundMode::CAST_NONE, static_cast<uint32_t>(currentNum));160+ Cast<half, int4b_t>(halfLocal, int4View, RoundMode::CAST_NONE, static_cast<uint32_t>(currentNum));
161- inQueueInt4_.FreeTensor(int4Local);161+ inQueueInt4_.FreeTensor(int4Local);
162- computeQueueHalf_.EnQue<half>(halfLocal);162+ computeQueueHalf_.EnQue<half>(halfLocal);
163- 163+ 
164- // ---- half -> int8 ----164+ // ---- half -> int8 ----
165- LocalTensor<half> halfSrc = computeQueueHalf_.DeQue<half>();165+ LocalTensor<half> halfSrc = computeQueueHalf_.DeQue<half>();
166- LocalTensor<int8_t> int8Local = outQueueInt8_.AllocTensor<int8_t>();166+ LocalTensor<int8_t> int8Local = outQueueInt8_.AllocTensor<int8_t>();
167- Cast<int8_t, half>(int8Local, halfSrc, RoundMode::CAST_ROUND, static_cast<uint32_t>(currentNum));167+ Cast<int8_t, half>(int8Local, halfSrc, RoundMode::CAST_ROUND, static_cast<uint32_t>(currentNum));
168- computeQueueHalf_.FreeTensor(halfSrc);168+ computeQueueHalf_.FreeTensor(halfSrc);
169- outQueueInt8_.EnQue<int8_t>(int8Local);169+ outQueueInt8_.EnQue<int8_t>(int8Local);
170-}170+}
171- 171+ 
172-__aicore__ inline void QbmmInt4ToInt8Preprocess::CopyOut(uint32_t progress, uint32_t currentNum)172+__aicore__ inline void QbmmInt4ToInt8Preprocess::CopyOut(uint32_t progress, uint32_t currentNum)
173-{173+{
174- LocalTensor<int8_t> int8Local = outQueueInt8_.DeQue<int8_t>();174+ LocalTensor<int8_t> int8Local = outQueueInt8_.DeQue<int8_t>();
175- DataCopy(dstInt8Global_[progress * ubLength_], int8Local, static_cast<uint32_t>(currentNum * sizeof(int8_t)));175+ DataCopy(dstInt8Global_[progress * ubLength_], int8Local, static_cast<uint32_t>(currentNum * sizeof(int8_t)));
176- outQueueInt8_.FreeTensor(int8Local);176+ outQueueInt8_.FreeTensor(int8Local);
177-}177+}
178- 178+ 
179#endif // QBMM_INT4_TO_INT8_PREPROCESS_H179#endif // QBMM_INT4_TO_INT8_PREPROCESS_H
Mmatmul/quant_batch_matmul_v4/docs/aclnnQuantMatmulV5.md+39-1
@@ -130,7 +130,19 @@
130 130 
131 - <term>Ascend 950PR/Ascend 950DT</term>：131 - <term>Ascend 950PR/Ascend 950DT</term>：
132 132 
133- 支持T-C && T-T、K-C && K-T、G-B、B-B、mx、T-CG [量化模式](../../../docs/zh/context/量化介绍.md)，不同量化模式对应的输入输出数据类型组合参见[约束说明](#约束说明)。133+ 支持T-C && T-T、K-C && K-T、G-B、B-B、mx、T-CG、K-G [量化模式](../../../docs/zh/context/量化介绍.md)，不同量化模式对应的输入输出数据类型组合参见[约束说明](#约束说明)。
kknan
kknankknan3月25日

aclnnQuantMatmulV5的.h文件也有描述，也需要适配新需求

likedislike
134+ 
135+ <details>
136+
137+ <summary><term><strong>K-G量化模式</strong></term></summary>
138+
139+ - x1，x2为INT4，x1Scale，x2Scale为FLOAT32，x2Offset为FLOAT16，out为FLOAT16/BFLOAT16 (pertoken-pergroup非对称量化)：
140+ 
141+ $$
142+ out = x1Scale * x2Scale @ (x1 @ x2 - x1 @ x2Offset)
143+ $$
144+ 
145+ </details>
134 146 
135 <details>147 <details>
136 148 
@@ -889,6 +901,32 @@ aclnnStatus aclnnQuantMatmulV5(
889 901 
890 </details>902 </details>
891 903 
904+ <details>
905+ 
906+ <summary><term><strong>K-G量化场景约束：</strong></term></summary>
907+ <a id="K-G量化"></a>
908+ 
909+ - 输入和输出支持以下数据类型组合：
910+ <a id="输入和输出支持以下数据类型组合K-G"></a>
911+ 
912+ | x1 | x2 | x1Scale | x2Scale | x2Offset | yScale | bias | yOffset | out |
913+ | ------------------------- | ------------------------- | ----------- | ----------- | ----------- | ------- | ------------ | -----------| -------------------------------------- |
914+ | INT4 | INT4 | FLOAT32 | FLOAT32 | FLOAT16 | null | null | null | BFLOAT16/FLOAT16 |
915+ 
916+ - x1、x2、x1Scale、x2Scale和groupSize的取值关系：
917+ 
918+ |量化类型| x1数据类型 | x2数据类型 | x1Scale数据类型| x2Scale数据类型| x1 shape | x2 shape| x1Scale shape| x2Scale shape| yOffset shape| [gsM，gsN，gsK]|
919+ | ----- | ------------------------- | ------------------------- | -------------- | ------------- | -------- | ------- | ------------ | ------------ | ------------ | ------------ |
920+ | K-G量化 | INT4 |INT4 |FLOAT32 |FLOAT32 |(m, k)|(n, k)|(m, 1)|(ceil(k / 256), n)| null | [0, 0, 256]|
921+ - x1的约束：
922+ - 当数据类型为INT4时，k需与1024对齐。transposeX1为false。
923+ - x2的约束：
924+ - 当数据类型为INT4时，k需与1024对称，n需与256对齐。transposeX2为true，
925+ - x2Scale的约束：
926+ - 当x1、x2为INT4时，x2Scale的shape为(ceil(k / 256), n)。
927+ 
928+ </details>
929+ 
892</details>930</details>
893 931 
894## 调用示例932## 调用示例
Amatmul/quant_batch_matmul_v4/examples/arch35/test_aclnn_quant_matmul_v5_int4_int4.cpp+281-0
@@ -0,0 +1,281 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include <iostream>
12+#include <memory>
13+#include <vector>
14+#include <cstdint>
15+ 
16+#include "acl/acl.h"
17+#include "aclnnop/aclnn_quant_matmul_v5.h"
18+ 
19+#define CHECK_RET(cond, return_expr) \
20+ do { \
21+ if (!(cond)) { \
22+ return_expr; \
23+ } \
24+ } while (0)
25+ 
26+#define CHECK_FREE_RET(cond, return_expr) \
27+ do { \
28+ if (!(cond)) { \
29+ Finalize(deviceId, stream); \
30+ return_expr; \
31+ } \
32+ } while (0)
33+ 
34+#define LOG_PRINT(message, ...) \
35+ do { \
36+ printf(message, ##__VA_ARGS__); \
37+ } while (0)
38+ 
39+int64_t GetShapeSize(const std::vector<int64_t>& shape)
40+{
41+ int64_t shapeSize = 1;
42+ for (auto i : shape) {
43+ shapeSize *= i;
44+ }
45+ return shapeSize;
46+}
47+ 
48+float Bf16ToFloat(uint16_t h)
49+{
50+ uint32_t bits = static_cast<uint32_t>(h) << 16;
51+ return *reinterpret_cast<float*>(&bits);
52+}
53+ 
54+int Init(int32_t deviceId, aclrtStream* stream)
55+{
56+ // 固定写法，资源初始化
57+ auto ret = aclInit(nullptr);
58+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
59+ ret = aclrtSetDevice(deviceId);
60+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
61+ ret = aclrtCreateStream(stream);
62+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
63+ return 0;
64+}
65+ 
66+template <typename T>
67+int CreateAclTensor(
68+ const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr, aclDataType dataType,
69+ aclTensor** tensor)
70+{
71+ auto size = hostData.size() * sizeof(T);
72+ // 调用aclrtMalloc申请device侧内存
73+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
74+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
75+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
76+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
77+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
78+ 
79+ // 计算连续tensor的strides
80+ std::vector<int64_t> strides(shape.size(), 1);
81+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
82+ strides[i] = shape[i + 1] * strides[i + 1];
83+ }
84+ 
85+ // 调用aclCreateTensor接口创建aclTensor
86+ *tensor = aclCreateTensor(
87+ shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND, shape.data(), shape.size(),
88+ *deviceAddr);
89+ return 0;
90+}
91+ 
92+void Finalize(int32_t deviceId, aclrtStream stream)
93+{
94+ aclrtDestroyStream(stream);
95+ aclrtResetDevice(deviceId);
96+ aclFinalize();
97+}
98+ 
99+int AclnnQuantMatmulV5Int4Int4Test(int32_t deviceId, aclrtStream& stream)
100+{
101+ auto ret = Init(deviceId, &stream);
102+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
103+ 
104+ // 2. 构造输入与输出，需要根据API的接口自定义构造
105+ int64_t m = 1024;
106+ int64_t k = 1024;
107+ int64_t n = 1024;
108+ 
109+ // K-CG量化模式：x1和x2为INT4，x1Scale和x2Scale为FLOAT32，x2Offset为FLOAT16，out为BFLOAT16
110+ aclDataType x1Dtype = aclDataType::ACL_INT4;
111+ aclDataType x2Dtype = aclDataType::ACL_INT4;
112+ aclDataType x1ScaleDtype = aclDataType::ACL_FLOAT;
113+ aclDataType x2ScaleDtype = aclDataType::ACL_FLOAT;
114+ aclDataType x2OffsetDtype = aclDataType::ACL_FLOAT16;
115+ aclDataType outDtype = aclDataType::ACL_BF16;
116+ 
117+ // 形状设置
118+ std::vector<int64_t> x1Shape = {m, k};
119+ std::vector<int64_t> x2Shape = {n, k}; // transposeX2为true时是(n, k)
120+ std::vector<int64_t> x1ScaleShape = {m, 1};
121+ std::vector<int64_t> x2ScaleShape = {k / 256, n};
122+ std::vector<int64_t> x2OffsetShape = {k / 256, n};
123+ std::vector<int64_t> outShape = {m, n};
124+ 
125+ // 设备内存地址
126+ void* x1DeviceAddr = nullptr;
127+ void* x2DeviceAddr = nullptr;
128+ void* x1ScaleDeviceAddr = nullptr;
129+ void* x2ScaleDeviceAddr = nullptr;
130+ void* x2OffsetDeviceAddr = nullptr;
131+ void* outDeviceAddr = nullptr;
132+ 
133+ // 张量指针
134+ aclTensor* x1 = nullptr;
135+ aclTensor* x2 = nullptr;
136+ aclTensor* x1Scale = nullptr;
137+ aclTensor* x2Scale = nullptr;
138+ aclTensor* x2Offset = nullptr;
139+ aclTensor* yScale = nullptr; // 不使用
140+ aclTensor* x1Offset = nullptr; // 不使用
141+ aclTensor* yOffset = nullptr; // 不使用
142+ aclTensor* bias = nullptr; // 不使用
143+ aclTensor* out = nullptr;
144+ 
145+ // 构造输入数据
146+ // x1/x2: INT4，每个int8存储2个int4值（低4bit + 高4bit）
147+ auto PackInt4Pair = [](int8_t low, int8_t high) -> int8_t {
148+ uint8_t lowNibble = static_cast<uint8_t>(low) & 0x0F;
149+ uint8_t highNibble = (static_cast<uint8_t>(high) & 0x0F) << 4;
150+ return static_cast<int8_t>(highNibble | lowNibble);
151+ };
152+ 
153+ std::vector<int8_t> x1HostData(m * k / 2, 0);
154+ std::vector<int8_t> x2HostData(n * k / 2, 0);
155+ 
156+ // x1 按 0 -> 7 循环递增
157+ for (int64_t i = 0; i < m * k / 2; ++i) {
158+ int64_t idx0 = 2 * i;
159+ int64_t idx1 = idx0 + 1;
160+ int8_t v0 = static_cast<int8_t>(idx0 % 8);
161+ int8_t v1 = static_cast<int8_t>(idx1 % 8);
162+ x1HostData[i] = PackInt4Pair(v0, v1);
163+ }
164+ 
165+ // x2 按 7 -> 0 循环递减
166+ for (int64_t i = 0; i < n * k / 2; ++i) {
167+ int64_t idx0 = 2 * i;
168+ int64_t idx1 = idx0 + 1;
169+ int8_t v0 = static_cast<int8_t>(7 - (idx0 % 8));
170+ int8_t v1 = static_cast<int8_t>(7 - (idx1 % 8));
171+ x2HostData[i] = PackInt4Pair(v0, v1);
172+ }
173+ // x1Scale: FLOAT32，全1
174+ std::vector<float> x1ScaleHostData(m, 1.0f);
175+ // x2Scale: FLOAT32，全1
176+ std::vector<float> x2ScaleHostData((k / 256) * n, 1.0f);
177+ // x2Offset: FLOAT16，全0
178+ std::vector<uint16_t> x2OffsetHostData((k / 256) * n, 0);
179+ // out: BFLOAT16，用于接收结果
180+ std::vector<uint16_t> outHostData(m * n, 0);
181+ 
182+ // 创建x1 aclTensor
183+ ret = CreateAclTensor(x1HostData, x1Shape, &x1DeviceAddr, x1Dtype, &x1);
184+ std::unique_ptr<aclTensor, aclnnStatus (*)(const aclTensor*)> x1TensorPtr(x1, aclDestroyTensor);
185+ std::unique_ptr<void, aclError (*)(void*)> x1DeviceAddrPtr(x1DeviceAddr, aclrtFree);
186+ CHECK_RET(ret == ACL_SUCCESS, return ret);
187+ 
188+ // 创建x2 aclTensor
189+ ret = CreateAclTensor(x2HostData, x2Shape, &x2DeviceAddr, x2Dtype, &x2);
190+ std::unique_ptr<aclTensor, aclnnStatus (*)(const aclTensor*)> x2TensorPtr(x2, aclDestroyTensor);
191+ std::unique_ptr<void, aclError (*)(void*)> x2DeviceAddrPtr(x2DeviceAddr, aclrtFree);
192+ CHECK_RET(ret == ACL_SUCCESS, return ret);
193+ 
194+ // 创建x1Scale aclTensor
195+ ret = CreateAclTensor(x1ScaleHostData, x1ScaleShape, &x1ScaleDeviceAddr, x1ScaleDtype, &x1Scale);
196+ std::unique_ptr<aclTensor, aclnnStatus (*)(const aclTensor*)> x1ScaleTensorPtr(x1Scale, aclDestroyTensor);
197+ std::unique_ptr<void, aclError (*)(void*)> x1ScaleDeviceAddrPtr(x1ScaleDeviceAddr, aclrtFree);
198+ CHECK_RET(ret == ACL_SUCCESS, return ret);
199+ 
200+ // 创建x2Scale aclTensor
201+ ret = CreateAclTensor(x2ScaleHostData, x2ScaleShape, &x2ScaleDeviceAddr, x2ScaleDtype, &x2Scale);
202+ std::unique_ptr<aclTensor, aclnnStatus (*)(const aclTensor*)> x2ScaleTensorPtr(x2Scale, aclDestroyTensor);
203+ std::unique_ptr<void, aclError (*)(void*)> x2ScaleDeviceAddrPtr(x2ScaleDeviceAddr, aclrtFree);
204+ CHECK_RET(ret == ACL_SUCCESS, return ret);
205+ 
206+ // 创建x2Offset aclTensor
207+ ret = CreateAclTensor(x2OffsetHostData, x2OffsetShape, &x2OffsetDeviceAddr, x2OffsetDtype, &x2Offset);
208+ std::unique_ptr<aclTensor, aclnnStatus (*)(const aclTensor*)> x2OffsetTensorPtr(x2Offset, aclDestroyTensor);
209+ std::unique_ptr<void, aclError (*)(void*)> x2OffsetDeviceAddrPtr(x2OffsetDeviceAddr, aclrtFree);
210+ CHECK_RET(ret == ACL_SUCCESS, return ret);
211+ 
212+ // 创建out aclTensor
213+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, outDtype, &out);
214+ std::unique_ptr<aclTensor, aclnnStatus (*)(const aclTensor*)> outTensorPtr(out, aclDestroyTensor);
215+ std::unique_ptr<void, aclError (*)(void*)> outDeviceAddrPtr(outDeviceAddr, aclrtFree);
216+ CHECK_RET(ret == ACL_SUCCESS, return ret);
217+ 
218+ // 转置设置
219+ bool transposeX1 = false;
220+ bool transposeX2 = true; // K-CG量化模式下x2需要转置
221+ 
222+ // groupSize设置：K-CG量化模式下为[0, 0, 256]
223+ int64_t groupSize = 256; // groupSizeK
224+ 
225+ // 3. 调用CANN算子库API
226+ uint64_t workspaceSize = 0;
227+ aclOpExecutor* executor = nullptr;
228+ 
229+ // 调用aclnnQuantMatmulV5第一段接口
230+ ret = aclnnQuantMatmulV5GetWorkspaceSize(
231+ x1, x2, x1Scale, x2Scale, yScale, x1Offset, x2Offset, yOffset, bias, transposeX1, transposeX2, groupSize, out,
232+ &workspaceSize, &executor);
233+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnQuantMatmulV5GetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
234+ 
235+ // 根据第一段接口计算出的workspaceSize申请device内存
236+ void* workspaceAddr = nullptr;
237+ std::unique_ptr<void, aclError (*)(void*)> workspaceAddrPtr(nullptr, aclrtFree);
238+ if (workspaceSize > 0) {
239+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
240+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
241+ workspaceAddrPtr.reset(workspaceAddr);
242+ }
243+ 
244+ // 调用aclnnQuantMatmulV5第二段接口
245+ ret = aclnnQuantMatmulV5(workspaceAddr, workspaceSize, executor, stream);
246+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnQuantMatmulV5 failed. ERROR: %d\n", ret); return ret);
247+ 
248+ // 4. （固定写法）同步等待任务执行结束
249+ ret = aclrtSynchronizeStream(stream);
250+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
251+ 
252+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧
253+ auto size = GetShapeSize(outShape);
254+ std::vector<uint16_t> resultData(size, 0); // C语言中无法直接打印bfloat16的数据，需要用uint16读出来
255+ ret = aclrtMemcpy(
256+ resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr, size * sizeof(resultData[0]),
257+ ACL_MEMCPY_DEVICE_TO_HOST);
258+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
259+ 
260+ // 打印部分结果
261+ LOG_PRINT("First 10 results:\n");
262+ for (int64_t i = 0; i < std::min<int64_t>(10, size); i++) {
263+ LOG_PRINT("result[%ld] is: %u\n", i, resultData[i]);
264+ }
265+ 
266+ return ACL_SUCCESS;
267+}
268+ 
269+int main(int argc, char* argv[])
270+{
271+ // 1. （固定写法）device/stream初始化，参考acl API手册
272+ // 根据自己的实际device填写deviceId
273+ int32_t deviceId = 0;
274+ aclrtStream stream;
275+ auto ret = AclnnQuantMatmulV5Int4Int4Test(deviceId, stream);
276+ CHECK_FREE_RET(ret == ACL_SUCCESS, LOG_PRINT("AclnnQuantMatmulV5Int4Int4Test failed. ERROR: %d\n", ret);
277+ return ret);
278+ 
279+ Finalize(deviceId, stream);
280+ return 0;
281+}
Mmatmul/quant_batch_matmul_v4/op_host/config/ascend950/quant_batch_matmul_v4_binary.json+310-0文件内容审核中，请稍后刷新重试
Mmatmul/quant_batch_matmul_v4/op_host/op_api/aclnn_quant_matmul_v5.h+1-1
@@ -29,7 +29,7 @@ extern "C" {
29 * @param [in] x2Scale: x2量化参数，数据类型支持：float8_e8m0, bfloat16, float32, int64, uint64。29 * @param [in] x2Scale: x2量化参数，数据类型支持：float8_e8m0, bfloat16, float32, int64, uint64。
30 * @param [in] yScale: y量化参数，数据类型支持：int64、uint64。30 * @param [in] yScale: y量化参数，数据类型支持：int64、uint64。
31 * @param [in] x1Offset: 预留参数，当前接口不支持该参数。31 * @param [in] x1Offset: 预留参数，当前接口不支持该参数。
32- * @param [in] x2Offset: 量化参数，数据类型支持：float32。32+ * @param [in] x2Offset: 量化参数，数据类型支持：float32, float16。
33 * @param [in] yOffset: 预留参数，当前接口不支持该参数。33 * @param [in] yOffset: 预留参数，当前接口不支持该参数。
34 * @param [in] bias: 偏置，数据类型支持：int32, bfloat16, float16, float32。34 * @param [in] bias: 偏置，数据类型支持：int32, bfloat16, float16, float32。
35 * @param [in] transposeX1: x1矩阵是否转置。35 * @param [in] transposeX1: x1矩阵是否转置。
Mmatmul/quant_batch_matmul_v4/op_host/op_tiling/arch35/quant_batch_matmul_v4_basic_block_tiling.h+1-1
@@ -23,7 +23,7 @@
23#include <vector>23#include <vector>
24 24 
25#include "error_util.h"25#include "error_util.h"
26-#include "../../../op_kernel/arch35/quant_batch_matmul_v4_tiling_data.h"26+#include "../../../op_kernel/arch35/quant_batch_matmul_v4_tiling_data_apt.h"
27 27 
28using std::ignore;28using std::ignore;
29using std::make_tuple;29using std::make_tuple;
Amatmul/quant_batch_matmul_v4/op_host/op_tiling/arch35/quant_batch_matmul_v4_pergroup_arch35_tiling.cpp+517-0
@@ -0,0 +1,517 @@
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 quant_batch_matmul_v4_pergroup_arch35_tiling.cpp
12+ * \brief
13+ */
14+ 
15+#include "quant_batch_matmul_v4_pergroup_arch35_tiling.h"
16+#include "../../../op_kernel/arch35/quant_batch_matmul_v4_tiling_key.h"
17+#include <algorithm>
18+ 
19+#include "error_util.h"
20+#include "graph/utils/type_utils.h"
21+#include "matmul/common/op_host/op_tiling/debug_tiling.h"
22+#include "platform/platform_infos_def.h"
23+ 
24+namespace optiling {
25+ 
26+namespace {
27+constexpr uint64_t GROUP_MKN_BIT_SIZE = 0xFFFF;
28+constexpr int ALIGN1024 = 1024;
29+constexpr int ALIGN256 = 256;
30+constexpr int64_t GROUP_SIZE_K = 256;
31+constexpr uint64_t L2_REAL_SIZE = 168; // B4真实的L2Size大小
32+constexpr uint64_t L2_FAKE_SIZE = 96; // B4被上层修改后的L2Size大小
33+constexpr const char* INT4_KG_QUANT_MODE = "int4 K-G quantification";
34+} // namespace
35+ 
36+const gert::Shape QuantBatchMatmulV4PergroupArch35Tiling::GetShape(const size_t index)
37+{
38+ return context_->GetInputShape(index)->GetStorageShape();
39+}
40+ 
41+const gert::Shape QuantBatchMatmulV4PergroupArch35Tiling::GetOptionShape(const size_t index)
42+{
43+ return context_->GetOptionalInputShape(index)->GetStorageShape();
44+}
45+ 
46+ge::graphStatus QuantBatchMatmulV4PergroupArch35Tiling::CalcDequantTiling(
47+ uint32_t baseM, uint32_t baseN, uint32_t groupSizeK)
48+{
49+ uint64_t ubSize = aicoreParams_.ubSize;
50+ uint64_t elesize = ubSize / sizeof(float);
51+ uint64_t aivM = ops::CeilDiv(baseM, 2U);
52+ constexpr uint64_t ALIGN8 = 8UL;
53+ 
54+ elesize -= aivM * baseN;
55+ elesize -= ops::CeilAlign(aivM, ALIGN8);
56+ elesize -= ops::CeilAlign(aivM, ALIGN8) * ALIGN8;
57+ 
58+ uint32_t ubCalcN = baseN;
59+ constexpr double NUM_UB_CALC_N_FACTOR = 3.5;
60+ elesize -= static_cast<uint64_t>(NUM_UB_CALC_N_FACTOR * ubCalcN);
61+ elesize -= ops::CeilAlign(aivM, ALIGN8);
62+ 
63+ constexpr double NUM_GROUP_SIZE_K_FACTOR = 1.5;
64+ uint32_t ubCalcM = elesize / (4U * ubCalcN + static_cast<uint32_t>(NUM_GROUP_SIZE_K_FACTOR * groupSizeK));
65+ tilingData_.params.ubCalcN = ubCalcN;
66+ tilingData_.params.ubCalcM = ubCalcM;
67+ OP_LOGD(inputParams_.opName, "UbTiling ubCalcM: %u, ubCalcN: %u, groupSizeK: %u", ubCalcM, ubCalcN, groupSizeK);
68+ return ge::GRAPH_SUCCESS;
69+}
70+ 
71+ge::graphStatus QuantBatchMatmulV4PergroupArch35Tiling::DoOpTiling()
72+{
73+ isUbQuant_ = true;
74+ InitCompileInfo();
75+ SetTransAttr(trans_);
76+ OP_LOGE_IF(!SetPlatformInfoForTiling(), ge::GRAPH_FAILED, inputParams_.opName, "SetPlatformInfoForTiling fail");
77+ 
78+ auto* platformInfoPtr = context_->GetPlatformInfo();
79+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
80+ 
81+ auto savedADtype = inputParams_.aDtype;
82+ auto savedBDtype = inputParams_.bDtype;
83+ 
84+ inputParams_.aDtype = ge::DT_INT8;
85+ inputParams_.bDtype = ge::DT_INT8;
86+ 
87+ OP_LOGE_IF(!DoBasicTiling(), ge::GRAPH_FAILED, inputParams_.opName, "DoBasicTiling failed.");
88+ 
89+ constexpr uint32_t BASE_M = 128;
90+ constexpr uint32_t BASE_N = 256;
91+ constexpr uint32_t DEPTH_A1 = 8;
92+ constexpr uint32_t DEPTH_B1 = 8;
93+ constexpr uint32_t DB_L0C = 1;
94+ constexpr uint32_t DB_L0A = 2;
95+ constexpr uint32_t DB_L0B = 1;
96+ 
97+ basicTiling_.baseM = BASE_M;
98+ basicTiling_.baseN = BASE_N;
99+ basicTiling_.depthA1 = DEPTH_A1;
100+ basicTiling_.depthB1 = DEPTH_B1;
101+ basicTiling_.dbL0c = DB_L0C;
102+ tilingData_.matmulTiling.dbL0A = DB_L0A;
103+ tilingData_.matmulTiling.dbL0B = DB_L0B;
104+ 
105+ uint32_t l1Size = aicoreParams_.l1Size;
106+ uint32_t stepK = l1Size / (2U * (BASE_M + BASE_N) * inputParams_.groupSizeK);
107+ stepK = std::max(1U, stepK);
108+ basicTiling_.stepKa = stepK;
109+ basicTiling_.stepKb = stepK;
110+ OP_LOGD(
111+ inputParams_.opName, "arch35 int8 tiling: groupSizeK=%u, stepK=%u, l1Size=%u", inputParams_.groupSizeK, stepK,
112+ l1Size);
113+ 
114+ basicTiling_.baseK = inputParams_.groupSizeK;
115+ QuantBatchMatmulV3BasicTiling::DoL2CacheTiling();
116+ 
117+ inputParams_.aDtype = savedADtype;
118+ inputParams_.bDtype = savedBDtype;
119+ if (basicTiling_.mTileCntl2 == 1UL && basicTiling_.nTileCntl2 == 1UL) {
120+ basicTiling_.mTileBlock = ops::CeilDiv(inputParams_.mSize, basicTiling_.baseM);
121+ basicTiling_.nTileBlock = ops::CeilDiv(inputParams_.nSize, basicTiling_.baseN);
122+ }
123+ basicTiling_.usedCoreNum = std::min(basicTiling_.usedCoreNum, basicTiling_.mTileBlock * basicTiling_.nTileBlock);
124+ tilingData_.params.groupSizeK = inputParams_.groupSizeK;
125+ return CalcDequantTiling(basicTiling_.baseM, basicTiling_.baseN, inputParams_.groupSizeK);
126+}
127+ 
128+ge::graphStatus QuantBatchMatmulV4PergroupArch35Tiling::PostTiling()
129+{
130+ ge::graphStatus ret = QuantBatchMatmulV3BasicTiling::PostTiling();
131+ if (ret != ge::GRAPH_SUCCESS) {
132+ return ret;
133+ }
134+ context_->SetScheduleMode(1);
135+ return ge::GRAPH_SUCCESS;
136+}
137+ 
138+ge::graphStatus QuantBatchMatmulV4PergroupArch35Tiling::CheckContext()
139+{
140+ auto x1Shape = context_->GetInputShape(X1_IDX);
141+ auto x1Desc = context_->GetInputDesc(X1_IDX);
142+ auto x2Shape = context_->GetInputShape(X2_IDX);
143+ auto x2Desc = context_->GetInputDesc(X2_IDX);
144+ auto outputShape = context_->GetOutputShape(Y_OUTPUT_IDX);
145+ auto outputDesc = context_->GetOutputDesc(Y_OUTPUT_IDX);
146+ auto attrs = context_->GetAttrs();
147+ 
148+ OP_TILING_CHECK(
149+ attrs == nullptr, VECTOR_INNER_ERR_REPORT_TILIING(inputParams_.opName, "Function context_.GetAttrs() failed!"),
150+ return ge::GRAPH_FAILED);
151+ 
152+ OPS_CHECK_NULL_WITH_CONTEXT(context_, x1Shape);
153+ OPS_CHECK_NULL_WITH_CONTEXT(context_, x1Desc);
154+ OPS_CHECK_NULL_WITH_CONTEXT(context_, x2Shape);
155+ OPS_CHECK_NULL_WITH_CONTEXT(context_, x2Desc);
156+ OPS_CHECK_NULL_WITH_CONTEXT(context_, outputShape);
157+ OPS_CHECK_NULL_WITH_CONTEXT(context_, outputDesc);
158+ OPS_CHECK_NULL_WITH_CONTEXT(context_, context_->GetRawTilingData());
159+ OPS_CHECK_NULL_WITH_CONTEXT(context_, context_->GetRawTilingData()->GetData());
160+ OP_TILING_CHECK(
161+ context_->GetRawTilingData()->GetCapacity() < tilingDataSize_,
162+ CUBE_INNER_ERR_REPORT(
163+ inputParams_.opName, "context tiling data capacity %zu < actual tiling data size %zu.",
164+ context_->GetRawTilingData()->GetCapacity(), tilingDataSize_),
165+ return ge::GRAPH_FAILED);
166+ return ge::GRAPH_SUCCESS;
167+}
168+ 
169+ge::graphStatus QuantBatchMatmulV4PergroupArch35Tiling::GetShapeAttrsInfo()
170+{
171+ tilingDataSize_ = sizeof(QuantBatchMatmulV3TilingData);
172+ if (context_ == nullptr) {
173+ return ge::GRAPH_FAILED;
174+ }
175+ 
176+ inputParams_.opName = context_->GetNodeName();
177+ OPS_LOG_D(inputParams_.opName, "TilingContext: %s", Ops::NN::DebugTilingContext(context_).c_str());
178+ 
179+ if (CheckContext() != ge::GRAPH_SUCCESS) {
180+ CUBE_INNER_ERR_REPORT(inputParams_.opName, "Invalid context.");
181+ return ge::GRAPH_FAILED;
182+ }
183+ 
184+ auto x1ScaleShape = context_->GetOptionalInputShape(X1_SCALE_IDX);
185+ auto x1ScaleDesc = context_->GetOptionalInputDesc(X1_SCALE_IDX);
186+ auto x2ScaleShape = context_->GetOptionalInputShape(X2_SCALE_IDX);
187+ auto x2ScaleDesc = context_->GetOptionalInputDesc(X2_SCALE_IDX);
188+ auto x2OffsetShape = context_->GetOptionalInputShape(X2_OFFSET_IDX);
189+ auto x2OffsetDesc = context_->GetOptionalInputDesc(X2_OFFSET_IDX);
190+ if (x1ScaleShape == nullptr || x1ScaleDesc == nullptr || x2ScaleShape == nullptr || x2ScaleDesc == nullptr ||
191+ x2OffsetShape == nullptr || x2OffsetDesc == nullptr) {
192+ OP_LOGD(inputParams_.opName, "x1Scale/x2Scale/x2Offset is nullptr, skip this tiling template.");
193+ return ge::GRAPH_PARAM_INVALID;
194+ }
195+ 
196+ if (!AnalyzeAttrs() || !AnalyzeDtype() || !AnalyzeInputs()) {
197+ OP_LOGD(inputParams_.opName, "Fail to analyze context info.");
198+ return ge::GRAPH_PARAM_INVALID;
199+ }
200+ 
201+ return ge::GRAPH_SUCCESS;
202+}
203+ 
204+bool QuantBatchMatmulV4PergroupArch35Tiling::IsCapable()
205+{
206+ auto* platformInfo = context_->GetPlatformInfo();
207+ if (platformInfo == nullptr) {
208+ return false;
209+ }
210+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo);
211+ if (ascendcPlatform.GetCurNpuArch() != NpuArch::DAV_3510) {
212+ return false;
213+ }
214+ if (!CheckPergroupAttrs()) {
215+ return false;
216+ }
217+ if (!CheckPergroupDtype()) {
218+ return false;
219+ }
220+ if (!CheckPergroupShape()) {
221+ return false;
222+ }
223+ return true;
224+}
225+ 
226+bool QuantBatchMatmulV4PergroupArch35Tiling::AnalyzeAttrs()
227+{
228+ auto attrs = context_->GetAttrs();
229+ auto transposeX1 = attrs->GetAttrPointer<bool>(TRANSPOSE_X1_IDX);
230+ auto transposeX2 = attrs->GetAttrPointer<bool>(TRANSPOSE_X2_IDX);
231+ const int64_t* groupSizePtr = attrs->GetAttrPointer<int64_t>(GROUP_SIZE_IDX);
232+ 
233+ inputParams_.groupSize = groupSizePtr != nullptr ? static_cast<uint64_t>(*groupSizePtr) : 0UL;
234+ inputParams_.groupSizeK = inputParams_.groupSize & GROUP_MKN_BIT_SIZE;
235+ inputParams_.groupSizeN = (inputParams_.groupSize >> 16U) & GROUP_MKN_BIT_SIZE;
236+ inputParams_.groupSizeM = (inputParams_.groupSize >> 32U) & GROUP_MKN_BIT_SIZE;
237+ inputParams_.transA = transposeX1 != nullptr && *transposeX1;
238+ inputParams_.transB = transposeX2 != nullptr && *transposeX2;
239+ return true;
240+}
241+ 
242+bool QuantBatchMatmulV4PergroupArch35Tiling::AnalyzeDtype()
243+{
244+ inputParams_.aDtype = context_->GetInputDesc(X1_IDX)->GetDataType();
245+ inputParams_.bDtype = context_->GetInputDesc(X2_IDX)->GetDataType();
246+ 
247+ auto x1ScaleDesc = context_->GetOptionalInputDesc(X1_SCALE_IDX);
248+ auto x2ScaleDesc = context_->GetOptionalInputDesc(X2_SCALE_IDX);
249+ auto x2OffsetDesc = context_->GetOptionalInputDesc(X2_OFFSET_IDX);
250+ 
251+ inputParams_.perTokenScaleDtype = x1ScaleDesc != nullptr ? x1ScaleDesc->GetDataType() : ge::DT_UNDEFINED;
252+ inputParams_.scaleDtype = x2ScaleDesc != nullptr ? x2ScaleDesc->GetDataType() : ge::DT_UNDEFINED;
253+ inputParamsPergroup_.x2OffsetDtype = x2OffsetDesc != nullptr ? x2OffsetDesc->GetDataType() : ge::DT_UNDEFINED;
254+ inputParams_.cDtype = context_->GetOutputDesc(Y_OUTPUT_IDX)->GetDataType();
255+ return true;
256+}
257+ 
258+bool QuantBatchMatmulV4PergroupArch35Tiling::AnalyzeInputs()
259+{
260+ auto x1Shape = GetShape(X1_IDX);
261+ auto x2Shape = GetShape(X2_IDX);
262+ 
263+ if (x1Shape.GetDimNum() != 2 || x2Shape.GetDimNum() != 2) {
264+ OP_LOGD(inputParams_.opName, "Input dims of x1/x2 should be 2 for this pergroup template.");
265+ return false;
266+ }
267+ 
268+ auto x1Inner = x1Shape.GetDim(1);
269+ auto x1Outer = x1Shape.GetDim(0);
270+ auto x2Inner = x2Shape.GetDim(1);
271+ auto x2Outer = x2Shape.GetDim(0);
272+ 
273+ inputParams_.mSize = static_cast<uint64_t>(inputParams_.transA ? x1Inner : x1Outer);
274+ inputParams_.kSize = static_cast<uint64_t>(inputParams_.transA ? x1Outer : x1Inner);
275+ inputParams_.nSize = static_cast<uint64_t>(inputParams_.transB ? x2Outer : x2Inner);
276+ inputParams_.batchA = GetBatchSize(x1Shape);
277+ inputParams_.batchB = GetBatchSize(x2Shape);
278+ 
279+ auto biasShapePtr = GetBiasShape(BIAS_IDX);
280+ inputParams_.hasBias = (biasShapePtr != nullptr);
281+ 
282+ AnalyzeBatchInfo(context_->GetInputShape(0)->GetOriginShape(), context_->GetInputShape(1)->GetOriginShape());
283+ (void)InferOutBatchDim(x1Shape, x2Shape);
284+ return true;
285+}
286+ 
287+bool QuantBatchMatmulV4PergroupArch35Tiling::SetPlatformInfoForTiling()
288+{
289+ if (!compileInfoInit_) {
290+ InitCompileInfo();
291+ }
292+ OP_LOGE_IF(compileInfo_.aicNum <= 0, false, inputParams_.opName, "coreNum <= 0");
293+ aicoreParams_.aicNum = compileInfo_.aicNum;
294+ OP_LOGE_IF(compileInfo_.l2Size <= 0, false, inputParams_.opName, "l2Size <= 0");
295+ compileInfo_.l2Size = compileInfo_.l2Size == L2_FAKE_SIZE * MB_SIZE ? L2_REAL_SIZE * MB_SIZE : compileInfo_.l2Size;
296+ inputParams_.libApiWorkSpaceSize = compileInfo_.workspaceNum;
297+ aicoreParams_.ubSize = compileInfo_.ubSize;
298+ aicoreParams_.l1Size = compileInfo_.l1Size;
299+ aicoreParams_.l0aSize = compileInfo_.l0aSize;
300+ aicoreParams_.l0cSize = compileInfo_.l0cSize;
301+ aicoreParams_.blockDim = 0;
302+ return true;
303+}
304+ 
305+bool QuantBatchMatmulV4PergroupArch35Tiling::CheckPergroupAttrs() const
306+{
307+ auto attrs = context_->GetAttrs();
308+ const int64_t* groupSizePtr = attrs->GetAttrPointer<int64_t>(GROUP_SIZE_IDX);
309+ if (groupSizePtr == nullptr) {
310+ OP_LOGD(inputParams_.opName, "In %s, groupSize is required, but got nullptr.", INT4_KG_QUANT_MODE);
311+ return false;
312+ }
313+ if (inputParams_.groupSizeK != GROUP_SIZE_K) {
314+ OP_LOGD(
315+ inputParams_.opName,
316+ "In %s, expected groupSizeK == %ld (groupSizeK = groupSize & 0xFFFF), but got %lu, full groupSize=%lu.",
317+ INT4_KG_QUANT_MODE, GROUP_SIZE_K, inputParams_.groupSizeK, inputParams_.groupSize);
318+ return false;
319+ }
320+ if (inputParams_.transA != false || inputParams_.transB != true) {
321+ OP_LOGD(
322+ inputParams_.opName,
323+ "In %s, only support transA=false and transB=true, but got transA=%s, transB=%s.",
324+ INT4_KG_QUANT_MODE, inputParams_.transA ? "true" : "false", inputParams_.transB ? "true" : "false");
325+ return false;
326+ }
327+ return true;
328+}
329+ 
330+bool QuantBatchMatmulV4PergroupArch35Tiling::CheckPergroupDtype() const
331+{
332+ if (inputParams_.aDtype != ge::DT_INT4) {
333+ OP_LOGD(
334+ inputParams_.opName, "In %s, expected x1 dtype=int4, but got %s.", INT4_KG_QUANT_MODE,
335+ ge::TypeUtils::DataTypeToSerialString(inputParams_.aDtype).c_str());
336+ return false;
337+ }
338+ if (inputParams_.bDtype != ge::DT_INT4) {
339+ OP_LOGD(
340+ inputParams_.opName, "In %s, expected x2 dtype=int4, but got %s.", INT4_KG_QUANT_MODE,
341+ ge::TypeUtils::DataTypeToSerialString(inputParams_.bDtype).c_str());
342+ return false;
343+ }
344+ if (inputParams_.perTokenScaleDtype != ge::DT_FLOAT) {
345+ OP_LOGD(
346+ inputParams_.opName, "In %s, expected x1Scale dtype=float32, but got %s.", INT4_KG_QUANT_MODE,
347+ ge::TypeUtils::DataTypeToSerialString(inputParams_.perTokenScaleDtype).c_str());
348+ return false;
349+ }
350+ if (inputParams_.scaleDtype != ge::DT_FLOAT) {
351+ OP_LOGD(
352+ inputParams_.opName, "In %s, expected x2Scale dtype=float32, but got %s.", INT4_KG_QUANT_MODE,
353+ ge::TypeUtils::DataTypeToSerialString(inputParams_.scaleDtype).c_str());
354+ return false;
355+ }
356+ if (!(inputParams_.cDtype == ge::DT_FLOAT16 || inputParams_.cDtype == ge::DT_BF16)) {
357+ OP_LOGD(
358+ inputParams_.opName, "In %s, expected output dtype in {float16, bfloat16}, but got %s.",
359+ INT4_KG_QUANT_MODE,
360+ ge::TypeUtils::DataTypeToSerialString(inputParams_.cDtype).c_str());
361+ return false;
362+ }
363+ if (inputParamsPergroup_.x2OffsetDtype != ge::DT_FLOAT16) {
364+ OP_LOGD(
365+ inputParams_.opName, "In %s, expected x2Offset dtype=float16, but got %s.", INT4_KG_QUANT_MODE,
366+ ge::TypeUtils::DataTypeToSerialString(inputParamsPergroup_.x2OffsetDtype).c_str());
367+ return false;
368+ }
369+ return true;
370+}
371+ 
372+bool QuantBatchMatmulV4PergroupArch35Tiling::CheckPergroupShape()
373+{
374+ // This template only handles int4 K-G quantification path on arch35.
375+ if (!CheckPergroupBasicShapeConstraints()) {
376+ return false;
377+ }
378+ 
379+ auto x1Shape = GetShape(X1_IDX);
380+ auto x2Shape = GetShape(X2_IDX);
381+ auto x1ScaleShape = GetOptionShape(X1_SCALE_IDX);
382+ auto x2ScaleShape = GetOptionShape(X2_SCALE_IDX);
383+ auto x2OffsetShape = GetOptionShape(X2_OFFSET_IDX);
384+ 
385+ return CheckPergroupDimAndOutput(x1Shape, x2Shape, x1ScaleShape, x2ScaleShape, x2OffsetShape) &&
386+ CheckPergroupScaleShape(x1ScaleShape, x2ScaleShape, x2OffsetShape) && CheckPergroupInputFormat();
387+}
388+ 
389+bool QuantBatchMatmulV4PergroupArch35Tiling::CheckPergroupBasicShapeConstraints() const
390+{
391+ if (inputParams_.hasBias) {
392+ OP_LOGD(inputParams_.opName, "In %s, expected bias=null, but got non-null bias.", INT4_KG_QUANT_MODE);
393+ return false;
394+ }
395+ if (inputParams_.kSize % ALIGN1024 != 0) {
396+ OP_LOGD(
397+ inputParams_.opName, "In %s, expected k aligned to %d, but got k=%lu.", INT4_KG_QUANT_MODE, ALIGN1024,
398+ inputParams_.kSize);
399+ return false;
400+ }
401+ if (inputParams_.nSize % ALIGN256 != 0) {
402+ OP_LOGD(
403+ inputParams_.opName, "In %s, expected n aligned to %d, but got n=%lu.", INT4_KG_QUANT_MODE, ALIGN256,
404+ inputParams_.nSize);
405+ return false;
406+ }
407+ return true;
408+}
409+ 
410+bool QuantBatchMatmulV4PergroupArch35Tiling::CheckPergroupDimAndOutput(
411+ const gert::Shape& x1Shape, const gert::Shape& x2Shape, const gert::Shape& x1ScaleShape,
412+ const gert::Shape& x2ScaleShape, const gert::Shape& x2OffsetShape)
413+{
414+ if (x1Shape.GetDimNum() != 2 || x2Shape.GetDimNum() != 2 || x1ScaleShape.GetDimNum() != 2 ||
415+ x2ScaleShape.GetDimNum() != 2 || x2OffsetShape.GetDimNum() != 2) {
416+ OP_LOGD(
417+ inputParams_.opName,
418+ "In %s, expected dims=2 for x1/x2/x1Scale/x2Scale/x2Offset, but got [%zu, %zu, %zu, %zu, %zu].",
419+ INT4_KG_QUANT_MODE, x1Shape.GetDimNum(), x2Shape.GetDimNum(), x1ScaleShape.GetDimNum(),
420+ x2ScaleShape.GetDimNum(), x2OffsetShape.GetDimNum());
421+ return false;
422+ }
423+ if (!InferOutBatchDim(x1Shape, x2Shape)) {
424+ OP_LOGD(
425+ inputParams_.opName, "In %s, expected batch dims of x1 and x2 to be broadcastable, but got non-broadcastable.",
426+ INT4_KG_QUANT_MODE);
427+ return false;
428+ }
429+ if (!CheckOutputShapeAvailable()) {
430+ OP_LOGD(
431+ inputParams_.opName, "In %s, expected output shape product within INT64_MAX, but got overflow risk.",
432+ INT4_KG_QUANT_MODE);
433+ return false;
434+ }
435+ return true;
436+}
437+ 
438+bool QuantBatchMatmulV4PergroupArch35Tiling::CheckPergroupScaleShape(
439+ const gert::Shape& x1ScaleShape, const gert::Shape& x2ScaleShape, const gert::Shape& x2OffsetShape) const
440+{
441+ int64_t m = static_cast<int64_t>(inputParams_.mSize);
442+ int64_t k = static_cast<int64_t>(inputParams_.kSize);
443+ int64_t n = static_cast<int64_t>(inputParams_.nSize);
444+ int64_t nkgroup = ops::CeilDiv(k, GROUP_SIZE_K);
445+ 
446+ if (x1ScaleShape.GetDim(0) != m || x1ScaleShape.GetDim(1) != 1) {
447+ OP_LOGD(
448+ inputParams_.opName, "In %s, expected x1Scale shape=[%ld, 1], but got [%ld, %ld].", INT4_KG_QUANT_MODE, m,
449+ x1ScaleShape.GetDim(0), x1ScaleShape.GetDim(1));
450+ return false;
451+ }
452+ if (x2ScaleShape.GetDim(0) != nkgroup || x2ScaleShape.GetDim(1) != n) {
453+ OP_LOGD(
454+ inputParams_.opName, "In %s, expected x2Scale shape=[%ld, %ld], but got [%ld, %ld].", INT4_KG_QUANT_MODE,
455+ nkgroup, n, x2ScaleShape.GetDim(0), x2ScaleShape.GetDim(1));
456+ return false;
457+ }
458+ if (x2OffsetShape.GetDim(0) != nkgroup || x2OffsetShape.GetDim(1) != n) {
459+ OP_LOGD(
460+ inputParams_.opName, "In %s, expected x2Offset shape=[%ld, %ld], but got [%ld, %ld].", INT4_KG_QUANT_MODE,
461+ nkgroup, n, x2OffsetShape.GetDim(0), x2OffsetShape.GetDim(1));
462+ return false;
463+ }
464+ return true;
465+}
466+ 
467+bool QuantBatchMatmulV4PergroupArch35Tiling::CheckPergroupInputFormat() const
468+{
469+ auto x1Format = static_cast<ge::Format>(ge::GetPrimaryFormat(context_->GetInputDesc(X1_IDX)->GetStorageFormat()));
470+ auto x2Format = static_cast<ge::Format>(ge::GetPrimaryFormat(context_->GetInputDesc(X2_IDX)->GetStorageFormat()));
471+ auto x1ScaleFormat =
472+ static_cast<ge::Format>(ge::GetPrimaryFormat(context_->GetOptionalInputDesc(X1_SCALE_IDX)->GetStorageFormat()));
473+ auto x2ScaleFormat =
474+ static_cast<ge::Format>(ge::GetPrimaryFormat(context_->GetOptionalInputDesc(X2_SCALE_IDX)->GetStorageFormat()));
475+ auto x2OffsetFormat = static_cast<ge::Format>(
476+ ge::GetPrimaryFormat(context_->GetOptionalInputDesc(X2_OFFSET_IDX)->GetStorageFormat()));
477+ 
478+ if (x1Format != ge::FORMAT_ND || x2Format != ge::FORMAT_ND || x1ScaleFormat != ge::FORMAT_ND ||
479+ x2ScaleFormat != ge::FORMAT_ND || x2OffsetFormat != ge::FORMAT_ND) {
480+ OP_LOGD(
481+ inputParams_.opName,
482+ "In %s, expected storage format=ND for x1/x2/x1Scale/x2Scale/x2Offset, but got [%s, %s, %s, %s, %s].",
483+ INT4_KG_QUANT_MODE, ge::TypeUtils::FormatToSerialString(x1Format).c_str(),
484+ ge::TypeUtils::FormatToSerialString(x2Format).c_str(),
485+ ge::TypeUtils::FormatToSerialString(x1ScaleFormat).c_str(),
486+ ge::TypeUtils::FormatToSerialString(x2ScaleFormat).c_str(),
487+ ge::TypeUtils::FormatToSerialString(x2OffsetFormat).c_str());
488+ return false;
489+ }
490+ return true;
491+}
492+ 
493+uint64_t QuantBatchMatmulV4PergroupArch35Tiling::GetTilingKey() const
494+{
495+ uint64_t trans = (static_cast<uint64_t>(inputParams_.transA) << 1) | static_cast<uint64_t>(inputParams_.transB);
496+ matmul_v4::KernelTemplateType kernelType = matmul_v4::KernelTemplateType::LUT_ASW;
497+ return GET_TPL_TILING_KEY(
498+ trans, static_cast<uint64_t>(matmul_v4::QuantType::INT4_ASYMMETRICAL), static_cast<uint64_t>(false),
499+ static_cast<uint64_t>(false), static_cast<uint64_t>(kernelType));
500+}
501+ 
502+ge::graphStatus QuantBatchMatmulV4PergroupArch35Tiling::GetWorkspaceSize()
503+{
504+ workspaceSize_ = inputParams_.libApiWorkSpaceSize;
505+ constexpr uint64_t BASIC_BLOCK_SIZE_128 = 128;
506+ uint64_t aInt8Size = inputParams_.mSize * inputParams_.kSize;
507+ uint64_t bInt8Size = inputParams_.kSize * inputParams_.nSize;
508+ uint64_t convertWorkspaceSize = ops::CeilAlign(aInt8Size, (uint64_t)BASIC_BLOCK_SIZE_128) +
509+ ops::CeilAlign(bInt8Size, (uint64_t)BASIC_BLOCK_SIZE_128);
510+ workspaceSize_ += convertWorkspaceSize;
511+ workspaceSize_ += sizeof(int32_t) * static_cast<uint64_t>(tilingData_.matmulTiling.baseM) *
512+ tilingData_.matmulTiling.baseN * tilingData_.matmulTiling.usedCoreNum * NUM_DB;
513+ 
514+ return ge::GRAPH_SUCCESS;
515+}
516+ 
517+} // namespace optiling
Amatmul/quant_batch_matmul_v4/op_host/op_tiling/arch35/quant_batch_matmul_v4_pergroup_arch35_tiling.h+69-0
@@ -0,0 +1,69 @@
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 quant_batch_matmul_v4_pergroup_arch35_tiling.h
12+ * \brief
13+ */
14+ 
15+#ifndef QUANT_BATCH_MATMUL_V4_PERGROUP_ARCH35_TILING_H
16+#define QUANT_BATCH_MATMUL_V4_PERGROUP_ARCH35_TILING_H
17+ 
18+#include "../quant_batch_matmul_v4_pergroup_tiling.h"
19+#include "quant_batch_matmul_v4_tiling.h"
20+#include "../quant_batch_matmul_v4_tiling_info.h"
21+ 
22+namespace optiling {
23+ 
24+struct QuantBatchMatmulPergroupArch35Info : public optiling::QuantBatchMatmulInfo {
25+ // Extra metadata for int4 K-G quantification path on arch35.
26+ ge::DataType x2OffsetDtype = ge::DT_UNDEFINED;
27+};
28+ 
29+class QuantBatchMatmulV4PergroupArch35Tiling : public QuantBatchMatmulV4PergroupTiling {
30+public:
31+ explicit QuantBatchMatmulV4PergroupArch35Tiling(gert::TilingContext* contextIn)
32+ : QuantBatchMatmulV4PergroupTiling(contextIn)
33+ {}
34+ QuantBatchMatmulV4PergroupArch35Tiling(gert::TilingContext* contextIn, QuantBatchMatmulV3TilingData* out)
35+ : QuantBatchMatmulV4PergroupTiling(contextIn, out)
36+ {}
37+ ~QuantBatchMatmulV4PergroupArch35Tiling() override = default;
38+ 
39+protected:
40+ const gert::Shape GetShape(const size_t index);
41+ const gert::Shape GetOptionShape(const size_t index);
42+ ge::graphStatus DoOpTiling() override;
43+ ge::graphStatus PostTiling() override;
44+ ge::graphStatus CheckContext() override;
45+ ge::graphStatus GetShapeAttrsInfo() override;
46+ bool IsCapable() override;
47+ bool AnalyzeAttrs() override;
48+ bool AnalyzeDtype() override;
49+ bool AnalyzeInputs() override;
50+ bool SetPlatformInfoForTiling() override;
51+ ge::graphStatus CalcDequantTiling(uint32_t baseM, uint32_t baseN, uint32_t groupSizeK);
52+ uint64_t GetTilingKey() const override;
53+ ge::graphStatus GetWorkspaceSize() override;
54+ 
55+private:
56+ bool CheckPergroupAttrs() const;
W
Wwuyi_513月25日

S4S4 K-G量化只支持A不转置B转置，tiling，aclnn里要有对应校验。

likedislike
57+ bool CheckPergroupDtype() const;
58+ bool CheckPergroupShape();
59+ bool CheckPergroupBasicShapeConstraints() const;
60+ bool CheckPergroupDimAndOutput(
61+ const gert::Shape& x1Shape, const gert::Shape& x2Shape, const gert::Shape& x1ScaleShape,
62+ const gert::Shape& x2ScaleShape, const gert::Shape& x2OffsetShape);
63+ bool CheckPergroupScaleShape(const gert::Shape& x1ScaleShape, const gert::Shape& x2ScaleShape, const gert::Shape& x2OffsetShape) const;
64+ bool CheckPergroupInputFormat() const;
65+ QuantBatchMatmulPergroupArch35Info inputParamsPergroup_;
66+};
67+} // namespace optiling
68+ 
69+#endif // QUANT_BATCH_MATMUL_V4_PERGROUP_ARCH35_TILING_H
Mmatmul/quant_batch_matmul_v4/op_host/op_tiling/arch35/quant_batch_matmul_v4_tiling.cpp+1-1
@@ -31,7 +31,7 @@
31#include "matmul/common/op_host/math_util.h"31#include "matmul/common/op_host/math_util.h"
32#include "platform/platform_infos_def.h"32#include "platform/platform_infos_def.h"
33#include "matmul/common/op_host/op_tiling/debug_tiling.h"33#include "matmul/common/op_host/op_tiling/debug_tiling.h"
34-#include "../../../op_kernel/arch35/quant_batch_matmul_v4_tiling_data.h"34+#include "../../../op_kernel/arch35/quant_batch_matmul_v4_tiling_data_apt.h"
35 35 
36using AscendC::BLOCK_CUBE;36using AscendC::BLOCK_CUBE;
37using namespace Ops::NN;37using namespace Ops::NN;
Mmatmul/quant_batch_matmul_v4/op_host/op_tiling/arch35/quant_batch_matmul_v4_tiling.h+3-2
@@ -27,7 +27,7 @@
27#include "quant_batch_matmul_v4_basic_block_tiling.h"27#include "quant_batch_matmul_v4_basic_block_tiling.h"
28#include "../../../../weight_quant_batch_matmul_v2/op_host/op_tiling/weight_quant_batch_matmul_v2_tiling_tool.h"28#include "../../../../weight_quant_batch_matmul_v2/op_host/op_tiling/weight_quant_batch_matmul_v2_tiling_tool.h"
29#include "../quant_batch_matmul_v4_compile_info.h"29#include "../quant_batch_matmul_v4_compile_info.h"
30-#include "../../../op_kernel/arch35/quant_batch_matmul_v4_tiling_data.h"30+#include "../../../op_kernel/arch35/quant_batch_matmul_v4_tiling_data_apt.h"
31 31 
32namespace optiling {32namespace optiling {
33using matmul_tiling::MatrixTraverse;33using matmul_tiling::MatrixTraverse;
@@ -99,7 +99,8 @@ enum class QuantType : uint32_t {
99 PER_CHANNEL = 2,99 PER_CHANNEL = 2,
100 PER_GROUP = 3,100 PER_GROUP = 3,
101 MX = 4,101 MX = 4,
102- PER_TILE = 5102+ PER_TILE = 5,
103+ INT4_ASYMMETRICAL = 6
103};104};
104 105 
105enum class KernelTemplateType : uint32_t {106enum class KernelTemplateType : uint32_t {
Mmatmul/quant_batch_matmul_v4/op_host/op_tiling/quant_batch_matmul_v4_tiling_registry.cpp+5-2
@@ -18,6 +18,7 @@
18#include "quant_batch_matmul_v4_pergroup_tiling.h"18#include "quant_batch_matmul_v4_pergroup_tiling.h"
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 "quant_batch_matmul_v4_compile_info.h"22#include "quant_batch_matmul_v4_compile_info.h"
22#include "error_util.h"23#include "error_util.h"
23#include "platform/platform_infos_def.h"24#include "platform/platform_infos_def.h"
@@ -31,12 +32,14 @@ constexpr int32_t MSD_PRIORITY = 2;
31constexpr int32_t PERBLOCK_PRIORITY = 3;32constexpr int32_t PERBLOCK_PRIORITY = 3;
32constexpr int32_t PERGROUP_PRIORITY = 4;33constexpr int32_t PERGROUP_PRIORITY = 4;
33constexpr int32_t LUT_PRIORITY = 5;34constexpr int32_t LUT_PRIORITY = 5;
35+constexpr int32_t PERGROUP_ARCH35_PRIORITY = 6;
34 36 
35REGISTER_TILING_TEMPLATE("QuantBatchMatmulV4", AdaptiveSlidingWindowBasicTilingV4, BASIC_PERBLOCK_PRIORITY);37REGISTER_TILING_TEMPLATE("QuantBatchMatmulV4", AdaptiveSlidingWindowBasicTilingV4, BASIC_PERBLOCK_PRIORITY);
36REGISTER_TILING_TEMPLATE("QuantBatchMatmulV4", QuantBatchMatmulV4MsdTiling, MSD_PRIORITY);38REGISTER_TILING_TEMPLATE("QuantBatchMatmulV4", QuantBatchMatmulV4MsdTiling, MSD_PRIORITY);
37REGISTER_TILING_TEMPLATE("QuantBatchMatmulV4", QuantBatchMatmulV4PerblockTiling, PERBLOCK_PRIORITY);39REGISTER_TILING_TEMPLATE("QuantBatchMatmulV4", QuantBatchMatmulV4PerblockTiling, PERBLOCK_PRIORITY);
38REGISTER_TILING_TEMPLATE("QuantBatchMatmulV4", QuantBatchMatmulV4PergroupTiling, PERGROUP_PRIORITY);40REGISTER_TILING_TEMPLATE("QuantBatchMatmulV4", QuantBatchMatmulV4PergroupTiling, PERGROUP_PRIORITY);
39REGISTER_TILING_TEMPLATE("QuantBatchMatmulV4", AdaptiveSlidingWindowTilingV4, LUT_PRIORITY);41REGISTER_TILING_TEMPLATE("QuantBatchMatmulV4", AdaptiveSlidingWindowTilingV4, LUT_PRIORITY);
42+REGISTER_TILING_TEMPLATE("QuantBatchMatmulV4", QuantBatchMatmulV4PergroupArch35Tiling, PERGROUP_ARCH35_PRIORITY);
40 43 
41 44 
42ge::graphStatus QuantBatchMatmulV4TilingFunc(gert::TilingContext *context)45ge::graphStatus QuantBatchMatmulV4TilingFunc(gert::TilingContext *context)
@@ -65,8 +68,8 @@ ge::graphStatus QuantBatchMatmulV4TilingFunc(gert::TilingContext *context)
65 } else if (supportMmadS8S4) {68 } else if (supportMmadS8S4) {
66 vector<int32_t> regitserList = {LUT_PRIORITY};69 vector<int32_t> regitserList = {LUT_PRIORITY};
67 return TilingRegistry::GetInstance().DoTilingImpl(context, regitserList);70 return TilingRegistry::GetInstance().DoTilingImpl(context, regitserList);
68- } 71+ }
69- std::vector<int32_t> registerList = {BASIC_PERBLOCK_PRIORITY, BASIC_PRIORITY};72+ std::vector<int32_t> registerList = {BASIC_PERBLOCK_PRIORITY, PERGROUP_ARCH35_PRIORITY, optiling::BASIC_PRIORITY};
70 return TilingRegistry::GetInstance().DoTilingImpl(context, registerList);73 return TilingRegistry::GetInstance().DoTilingImpl(context, registerList);
71}74}
72 75 
Mmatmul/quant_batch_matmul_v4/op_host/quant_batch_matmul_v4_def.cpp+70-33
@@ -47,7 +47,9 @@ public:
47 ge::DT_FLOAT8_E4M3FN, ge::DT_FLOAT8_E5M2,47 ge::DT_FLOAT8_E4M3FN, ge::DT_FLOAT8_E5M2,
48 ge::DT_FLOAT8_E4M3FN, ge::DT_FLOAT8_E5M2,48 ge::DT_FLOAT8_E4M3FN, ge::DT_FLOAT8_E5M2,
49 ge::DT_FLOAT8_E4M3FN, ge::DT_FLOAT8_E5M2,49 ge::DT_FLOAT8_E4M3FN, ge::DT_FLOAT8_E5M2,
50- ge::DT_FLOAT8_E4M3FN, ge::DT_FLOAT8_E4M3FN, ge::DT_FLOAT8_E4M3FN, ge::DT_FLOAT8_E4M3FN, ge::DT_FLOAT8_E4M3FN})50+ ge::DT_FLOAT8_E4M3FN, ge::DT_FLOAT8_E4M3FN, ge::DT_FLOAT8_E4M3FN, ge::DT_FLOAT8_E4M3FN, ge::DT_FLOAT8_E4M3FN,
51+ ge::DT_INT4, ge::DT_INT4
52+ })
51 .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,53 .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
52 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,54 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
53 ge::FORMAT_ND, ge::FORMAT_ND,55 ge::FORMAT_ND, ge::FORMAT_ND,
@@ -73,7 +75,8 @@ public:
73 ge::FORMAT_ND, ge::FORMAT_ND,75 ge::FORMAT_ND, ge::FORMAT_ND,
74 ge::FORMAT_ND, ge::FORMAT_ND,76 ge::FORMAT_ND, ge::FORMAT_ND,
75 ge::FORMAT_ND, ge::FORMAT_ND,77 ge::FORMAT_ND, ge::FORMAT_ND,
76- ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND78+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
79+ ge::FORMAT_ND, ge::FORMAT_ND
77 })80 })
78 .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,81 .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
79 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,82 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
@@ -100,7 +103,9 @@ public:
100 ge::FORMAT_ND, ge::FORMAT_ND,103 ge::FORMAT_ND, ge::FORMAT_ND,
101 ge::FORMAT_ND, ge::FORMAT_ND,104 ge::FORMAT_ND, ge::FORMAT_ND,
102 ge::FORMAT_ND, ge::FORMAT_ND,105 ge::FORMAT_ND, ge::FORMAT_ND,
103- ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});106+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
107+ ge::FORMAT_ND, ge::FORMAT_ND
108+ });
104 this->Input("x2")109 this->Input("x2")
105 .ParamType(REQUIRED)110 .ParamType(REQUIRED)
106 .DataType({ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8,111 .DataType({ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8,
@@ -128,7 +133,9 @@ public:
128 ge::DT_FLOAT4_E2M1, ge::DT_FLOAT4_E2M1,133 ge::DT_FLOAT4_E2M1, ge::DT_FLOAT4_E2M1,
129 ge::DT_FLOAT4_E2M1, ge::DT_FLOAT4_E2M1,134 ge::DT_FLOAT4_E2M1, ge::DT_FLOAT4_E2M1,
130 ge::DT_FLOAT4_E2M1, ge::DT_FLOAT4_E2M1,135 ge::DT_FLOAT4_E2M1, ge::DT_FLOAT4_E2M1,
131- ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT8_E4M3FN, ge::DT_FLOAT8_E4M3FN, ge::DT_FLOAT8_E4M3FN})136+ ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT8_E4M3FN, ge::DT_FLOAT8_E4M3FN, ge::DT_FLOAT8_E4M3FN,
137+ ge::DT_INT4, ge::DT_INT4
138+ })
132 .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,139 .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
133 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,140 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
134 ge::FORMAT_ND, ge::FORMAT_ND,141 ge::FORMAT_ND, ge::FORMAT_ND,
@@ -154,7 +161,9 @@ public:
154 ge::FORMAT_ND, ge::FORMAT_ND,161 ge::FORMAT_ND, ge::FORMAT_ND,
155 ge::FORMAT_FRACTAL_NZ, ge::FORMAT_FRACTAL_NZ,162 ge::FORMAT_FRACTAL_NZ, ge::FORMAT_FRACTAL_NZ,
156 ge::FORMAT_FRACTAL_NZ, ge::FORMAT_FRACTAL_NZ,163 ge::FORMAT_FRACTAL_NZ, ge::FORMAT_FRACTAL_NZ,
157- ge::FORMAT_FRACTAL_NZ_C0_4, ge::FORMAT_FRACTAL_NZ_C0_4, ge::FORMAT_FRACTAL_NZ, ge::FORMAT_FRACTAL_NZ, ge::FORMAT_FRACTAL_NZ})164+ ge::FORMAT_FRACTAL_NZ_C0_4, ge::FORMAT_FRACTAL_NZ_C0_4, ge::FORMAT_FRACTAL_NZ, ge::FORMAT_FRACTAL_NZ, ge::FORMAT_FRACTAL_NZ,
165+ ge::FORMAT_ND, ge::FORMAT_ND
166+ })
158 .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,167 .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
159 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,168 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
160 ge::FORMAT_ND, ge::FORMAT_ND,169 ge::FORMAT_ND, ge::FORMAT_ND,
@@ -180,7 +189,8 @@ public:
180 ge::FORMAT_ND, ge::FORMAT_ND,189 ge::FORMAT_ND, ge::FORMAT_ND,
181 ge::FORMAT_FRACTAL_NZ, ge::FORMAT_FRACTAL_NZ,190 ge::FORMAT_FRACTAL_NZ, ge::FORMAT_FRACTAL_NZ,
182 ge::FORMAT_FRACTAL_NZ, ge::FORMAT_FRACTAL_NZ,191 ge::FORMAT_FRACTAL_NZ, ge::FORMAT_FRACTAL_NZ,
183- ge::FORMAT_FRACTAL_NZ_C0_4, ge::FORMAT_FRACTAL_NZ_C0_4, ge::FORMAT_FRACTAL_NZ, ge::FORMAT_FRACTAL_NZ, ge::FORMAT_FRACTAL_NZ});192+ ge::FORMAT_FRACTAL_NZ_C0_4, ge::FORMAT_FRACTAL_NZ_C0_4, ge::FORMAT_FRACTAL_NZ, ge::FORMAT_FRACTAL_NZ, ge::FORMAT_FRACTAL_NZ,
193+ ge::FORMAT_ND, ge::FORMAT_ND});
184 this->Input("bias")194 this->Input("bias")
185 .ParamType(OPTIONAL)195 .ParamType(OPTIONAL)
186 .DataType({ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32,196 .DataType({ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32,
@@ -208,7 +218,8 @@ public:
208 ge::DT_FLOAT16, ge::DT_FLOAT16,218 ge::DT_FLOAT16, ge::DT_FLOAT16,
209 ge::DT_BF16, ge::DT_BF16,219 ge::DT_BF16, ge::DT_BF16,
210 ge::DT_FLOAT16, ge::DT_FLOAT16,220 ge::DT_FLOAT16, ge::DT_FLOAT16,
211- ge::DT_BF16, ge::DT_BF16, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT})221+ ge::DT_BF16, ge::DT_BF16, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT,
222+ ge::DT_BF16, ge::DT_BF16})
212 .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,223 .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
213 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,224 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
214 ge::FORMAT_ND, ge::FORMAT_ND,225 ge::FORMAT_ND, ge::FORMAT_ND,
@@ -234,7 +245,8 @@ public:
234 ge::FORMAT_ND, ge::FORMAT_ND,245 ge::FORMAT_ND, ge::FORMAT_ND,
235 ge::FORMAT_ND, ge::FORMAT_ND,246 ge::FORMAT_ND, ge::FORMAT_ND,
236 ge::FORMAT_ND, ge::FORMAT_ND,247 ge::FORMAT_ND, ge::FORMAT_ND,
237- ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})248+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
249+ ge::FORMAT_ND, ge::FORMAT_ND})
238 .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,250 .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
239 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,251 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
240 ge::FORMAT_ND, ge::FORMAT_ND,252 ge::FORMAT_ND, ge::FORMAT_ND,
@@ -260,7 +272,8 @@ public:
260 ge::FORMAT_ND, ge::FORMAT_ND,272 ge::FORMAT_ND, ge::FORMAT_ND,
261 ge::FORMAT_ND, ge::FORMAT_ND,273 ge::FORMAT_ND, ge::FORMAT_ND,
262 ge::FORMAT_ND, ge::FORMAT_ND,274 ge::FORMAT_ND, ge::FORMAT_ND,
263- ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});275+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
276+ ge::FORMAT_ND, ge::FORMAT_ND});
264 this->Input("x1_scale")277 this->Input("x1_scale")
265 .ParamType(OPTIONAL)278 .ParamType(OPTIONAL)
266 .DataType({ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT,279 .DataType({ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT,
@@ -288,7 +301,8 @@ public:
288 ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E8M0,301 ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E8M0,
289 ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E8M0,302 ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E8M0,
290 ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E8M0,303 ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E8M0,
291- ge::DT_FLOAT8_E8M0, ge::DT_FLOAT, ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E8M0})304+ ge::DT_FLOAT8_E8M0, ge::DT_FLOAT, ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E8M0,
305+ ge::DT_FLOAT, ge::DT_FLOAT})
292 .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,306 .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
293 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,307 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
294 ge::FORMAT_ND, ge::FORMAT_ND,308 ge::FORMAT_ND, ge::FORMAT_ND,
@@ -314,7 +328,8 @@ public:
314 ge::FORMAT_ND, ge::FORMAT_ND,328 ge::FORMAT_ND, ge::FORMAT_ND,
315 ge::FORMAT_ND, ge::FORMAT_ND,329 ge::FORMAT_ND, ge::FORMAT_ND,
316 ge::FORMAT_ND, ge::FORMAT_ND,330 ge::FORMAT_ND, ge::FORMAT_ND,
317- ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})331+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
332+ ge::FORMAT_ND, ge::FORMAT_ND})
318 .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,333 .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
319 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,334 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
320 ge::FORMAT_ND, ge::FORMAT_ND,335 ge::FORMAT_ND, ge::FORMAT_ND,
@@ -340,7 +355,8 @@ public:
340 ge::FORMAT_ND, ge::FORMAT_ND,355 ge::FORMAT_ND, ge::FORMAT_ND,
341 ge::FORMAT_ND, ge::FORMAT_ND,356 ge::FORMAT_ND, ge::FORMAT_ND,
342 ge::FORMAT_ND, ge::FORMAT_ND,357 ge::FORMAT_ND, ge::FORMAT_ND,
343- ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});358+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
359+ ge::FORMAT_ND, ge::FORMAT_ND});
344 this->Input("x2_scale")360 this->Input("x2_scale")
345 .ParamType(OPTIONAL)361 .ParamType(OPTIONAL)
346 .DataType({ge::DT_UINT64, ge::DT_UINT64, ge::DT_UINT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_FLOAT, ge::DT_FLOAT,362 .DataType({ge::DT_UINT64, ge::DT_UINT64, ge::DT_UINT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_FLOAT, ge::DT_FLOAT,
@@ -368,7 +384,8 @@ public:
368 ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E8M0,384 ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E8M0,
369 ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E8M0,385 ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E8M0,
370 ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E8M0,386 ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E8M0,
371- ge::DT_FLOAT8_E8M0, ge::DT_BF16, ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E8M0})387+ ge::DT_FLOAT8_E8M0, ge::DT_BF16, ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E8M0,
388+ ge::DT_FLOAT, ge::DT_FLOAT})
372 .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,389 .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
373 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,390 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
374 ge::FORMAT_ND, ge::FORMAT_ND,391 ge::FORMAT_ND, ge::FORMAT_ND,
@@ -394,7 +411,8 @@ public:
394 ge::FORMAT_ND, ge::FORMAT_ND,411 ge::FORMAT_ND, ge::FORMAT_ND,
395 ge::FORMAT_ND, ge::FORMAT_ND,412 ge::FORMAT_ND, ge::FORMAT_ND,
396 ge::FORMAT_ND, ge::FORMAT_ND,413 ge::FORMAT_ND, ge::FORMAT_ND,
397- ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})414+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
415+ ge::FORMAT_ND, ge::FORMAT_ND})
398 .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,416 .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
399 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,417 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
400 ge::FORMAT_ND, ge::FORMAT_ND,418 ge::FORMAT_ND, ge::FORMAT_ND,
@@ -420,7 +438,8 @@ public:
420 ge::FORMAT_ND, ge::FORMAT_ND,438 ge::FORMAT_ND, ge::FORMAT_ND,
421 ge::FORMAT_ND, ge::FORMAT_ND,439 ge::FORMAT_ND, ge::FORMAT_ND,
422 ge::FORMAT_ND, ge::FORMAT_ND,440 ge::FORMAT_ND, ge::FORMAT_ND,
423- ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});441+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
442+ ge::FORMAT_ND, ge::FORMAT_ND});
424 this->Input("y_scale")443 this->Input("y_scale")
425 .ParamType(OPTIONAL)444 .ParamType(OPTIONAL)
426 .DataType({ge::DT_UINT64, ge::DT_UINT64, ge::DT_UINT64, ge::DT_UINT64, ge::DT_UINT64, ge::DT_UINT64, ge::DT_UINT64, ge::DT_UINT64,445 .DataType({ge::DT_UINT64, ge::DT_UINT64, ge::DT_UINT64, ge::DT_UINT64, ge::DT_UINT64, ge::DT_UINT64, ge::DT_UINT64, ge::DT_UINT64,
@@ -448,7 +467,8 @@ public:
448 ge::DT_UINT64, ge::DT_UINT64,467 ge::DT_UINT64, ge::DT_UINT64,
449 ge::DT_UINT64, ge::DT_UINT64,468 ge::DT_UINT64, ge::DT_UINT64,
450 ge::DT_UINT64, ge::DT_UINT64,469 ge::DT_UINT64, ge::DT_UINT64,
451- ge::DT_UINT64, ge::DT_UINT64, ge::DT_UINT64, ge::DT_UINT64, ge::DT_UINT64})470+ ge::DT_UINT64, ge::DT_UINT64, ge::DT_UINT64, ge::DT_UINT64, ge::DT_UINT64,
471+ ge::DT_UINT64, ge::DT_UINT64})
452 .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,472 .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
453 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,473 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
454 ge::FORMAT_ND, ge::FORMAT_ND,474 ge::FORMAT_ND, ge::FORMAT_ND,
@@ -474,7 +494,8 @@ public:
474 ge::FORMAT_ND, ge::FORMAT_ND,494 ge::FORMAT_ND, ge::FORMAT_ND,
475 ge::FORMAT_ND, ge::FORMAT_ND,495 ge::FORMAT_ND, ge::FORMAT_ND,
476 ge::FORMAT_ND, ge::FORMAT_ND,496 ge::FORMAT_ND, ge::FORMAT_ND,
477- ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})497+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
498+ ge::FORMAT_ND, ge::FORMAT_ND})
478 .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,499 .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
479 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,500 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
480 ge::FORMAT_ND, ge::FORMAT_ND,501 ge::FORMAT_ND, ge::FORMAT_ND,
@@ -500,7 +521,8 @@ public:
500 ge::FORMAT_ND, ge::FORMAT_ND,521 ge::FORMAT_ND, ge::FORMAT_ND,
501 ge::FORMAT_ND, ge::FORMAT_ND,522 ge::FORMAT_ND, ge::FORMAT_ND,
502 ge::FORMAT_ND, ge::FORMAT_ND,523 ge::FORMAT_ND, ge::FORMAT_ND,
503- ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});524+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
525+ ge::FORMAT_ND, ge::FORMAT_ND});
504 this->Input("x1_offset")526 this->Input("x1_offset")
505 .ParamType(OPTIONAL)527 .ParamType(OPTIONAL)
506 .DataType({ge::DT_BF16, ge::DT_BF16, ge::DT_BF16, ge::DT_BF16, ge::DT_BF16, ge::DT_BF16, ge::DT_BF16, ge::DT_BF16,528 .DataType({ge::DT_BF16, ge::DT_BF16, ge::DT_BF16, ge::DT_BF16, ge::DT_BF16, ge::DT_BF16, ge::DT_BF16, ge::DT_BF16,
@@ -528,7 +550,8 @@ public:
528 ge::DT_FLOAT16, ge::DT_FLOAT16,550 ge::DT_FLOAT16, ge::DT_FLOAT16,
529 ge::DT_BF16, ge::DT_BF16,551 ge::DT_BF16, ge::DT_BF16,
530 ge::DT_FLOAT16, ge::DT_FLOAT16,552 ge::DT_FLOAT16, ge::DT_FLOAT16,
531- ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT16})553+ ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT16,
554+ ge::DT_FLOAT16, ge::DT_FLOAT16})
532 .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,555 .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
533 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,556 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
534 ge::FORMAT_ND, ge::FORMAT_ND,557 ge::FORMAT_ND, ge::FORMAT_ND,
@@ -554,7 +577,8 @@ public:
554 ge::FORMAT_ND, ge::FORMAT_ND,577 ge::FORMAT_ND, ge::FORMAT_ND,
555 ge::FORMAT_ND, ge::FORMAT_ND,578 ge::FORMAT_ND, ge::FORMAT_ND,
556 ge::FORMAT_ND, ge::FORMAT_ND,579 ge::FORMAT_ND, ge::FORMAT_ND,
557- ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})580+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
581+ ge::FORMAT_ND, ge::FORMAT_ND})
558 .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,582 .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
559 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,583 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
560 ge::FORMAT_ND, ge::FORMAT_ND,584 ge::FORMAT_ND, ge::FORMAT_ND,
@@ -580,7 +604,8 @@ public:
580 ge::FORMAT_ND, ge::FORMAT_ND,604 ge::FORMAT_ND, ge::FORMAT_ND,
581 ge::FORMAT_ND, ge::FORMAT_ND,605 ge::FORMAT_ND, ge::FORMAT_ND,
582 ge::FORMAT_ND, ge::FORMAT_ND,606 ge::FORMAT_ND, ge::FORMAT_ND,
583- ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});607+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
608+ ge::FORMAT_ND, ge::FORMAT_ND});
584 this->Input("x2_offset")609 this->Input("x2_offset")
585 .ParamType(OPTIONAL)610 .ParamType(OPTIONAL)
586 .DataType({ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT,611 .DataType({ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT,
@@ -608,7 +633,8 @@ public:
608 ge::DT_FLOAT16, ge::DT_FLOAT16,633 ge::DT_FLOAT16, ge::DT_FLOAT16,
609 ge::DT_BF16, ge::DT_BF16,634 ge::DT_BF16, ge::DT_BF16,
610 ge::DT_FLOAT16, ge::DT_FLOAT16,635 ge::DT_FLOAT16, ge::DT_FLOAT16,
611- ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT16})636+ ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT16,
637+ ge::DT_FLOAT16, ge::DT_FLOAT16})
612 .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,638 .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
613 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,639 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
614 ge::FORMAT_ND, ge::FORMAT_ND,640 ge::FORMAT_ND, ge::FORMAT_ND,
@@ -634,7 +660,8 @@ public:
634 ge::FORMAT_ND, ge::FORMAT_ND,660 ge::FORMAT_ND, ge::FORMAT_ND,
635 ge::FORMAT_ND, ge::FORMAT_ND,661 ge::FORMAT_ND, ge::FORMAT_ND,
636 ge::FORMAT_ND, ge::FORMAT_ND,662 ge::FORMAT_ND, ge::FORMAT_ND,
637- ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})663+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
664+ ge::FORMAT_ND, ge::FORMAT_ND})
638 .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,665 .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
639 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,666 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
640 ge::FORMAT_ND, ge::FORMAT_ND,667 ge::FORMAT_ND, ge::FORMAT_ND,
@@ -660,7 +687,8 @@ public:
660 ge::FORMAT_ND, ge::FORMAT_ND,687 ge::FORMAT_ND, ge::FORMAT_ND,
661 ge::FORMAT_ND, ge::FORMAT_ND,688 ge::FORMAT_ND, ge::FORMAT_ND,
662 ge::FORMAT_ND, ge::FORMAT_ND,689 ge::FORMAT_ND, ge::FORMAT_ND,
663- ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});690+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
691+ ge::FORMAT_ND, ge::FORMAT_ND});
664 this->Input("y_offset")692 this->Input("y_offset")
665 .ParamType(OPTIONAL)693 .ParamType(OPTIONAL)
666 .DataType({ge::DT_BF16, ge::DT_BF16, ge::DT_BF16, ge::DT_BF16, ge::DT_BF16, ge::DT_BF16, ge::DT_BF16, ge::DT_BF16,694 .DataType({ge::DT_BF16, ge::DT_BF16, ge::DT_BF16, ge::DT_BF16, ge::DT_BF16, ge::DT_BF16, ge::DT_BF16, ge::DT_BF16,
@@ -688,7 +716,8 @@ public:
688 ge::DT_FLOAT16, ge::DT_FLOAT16,716 ge::DT_FLOAT16, ge::DT_FLOAT16,
689 ge::DT_BF16, ge::DT_BF16,717 ge::DT_BF16, ge::DT_BF16,
690 ge::DT_FLOAT16, ge::DT_FLOAT16,718 ge::DT_FLOAT16, ge::DT_FLOAT16,
691- ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT16})719+ ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT16,
720+ ge::DT_FLOAT, ge::DT_FLOAT})
692 .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,721 .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
693 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,722 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
694 ge::FORMAT_ND, ge::FORMAT_ND,723 ge::FORMAT_ND, ge::FORMAT_ND,
@@ -714,7 +743,8 @@ public:
714 ge::FORMAT_ND, ge::FORMAT_ND,743 ge::FORMAT_ND, ge::FORMAT_ND,
715 ge::FORMAT_ND, ge::FORMAT_ND,744 ge::FORMAT_ND, ge::FORMAT_ND,
716 ge::FORMAT_ND, ge::FORMAT_ND,745 ge::FORMAT_ND, ge::FORMAT_ND,
717- ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})746+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
747+ ge::FORMAT_ND, ge::FORMAT_ND})
718 .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,748 .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
719 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,749 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
720 ge::FORMAT_ND, ge::FORMAT_ND,750 ge::FORMAT_ND, ge::FORMAT_ND,
@@ -740,7 +770,8 @@ public:
740 ge::FORMAT_ND, ge::FORMAT_ND,770 ge::FORMAT_ND, ge::FORMAT_ND,
741 ge::FORMAT_ND, ge::FORMAT_ND,771 ge::FORMAT_ND, ge::FORMAT_ND,
742 ge::FORMAT_ND, ge::FORMAT_ND,772 ge::FORMAT_ND, ge::FORMAT_ND,
743- ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});773+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
774+ ge::FORMAT_ND, ge::FORMAT_ND});
744 this->Input("x2_table")775 this->Input("x2_table")
745 .ParamType(OPTIONAL)776 .ParamType(OPTIONAL)
746 .DataType({ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8,777 .DataType({ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8,
@@ -768,7 +799,8 @@ public:
768 ge::DT_INT8, ge::DT_INT8,799 ge::DT_INT8, ge::DT_INT8,
769 ge::DT_INT8, ge::DT_INT8,800 ge::DT_INT8, ge::DT_INT8,
770 ge::DT_INT8, ge::DT_INT8,801 ge::DT_INT8, ge::DT_INT8,
771- ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8})802+ ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8,
803+ ge::DT_INT8, ge::DT_INT8})
772 .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,804 .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
773 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,805 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
774 ge::FORMAT_ND, ge::FORMAT_ND,806 ge::FORMAT_ND, ge::FORMAT_ND,
@@ -794,7 +826,8 @@ public:
794 ge::FORMAT_ND, ge::FORMAT_ND,826 ge::FORMAT_ND, ge::FORMAT_ND,
795 ge::FORMAT_ND, ge::FORMAT_ND,827 ge::FORMAT_ND, ge::FORMAT_ND,
796 ge::FORMAT_ND, ge::FORMAT_ND,828 ge::FORMAT_ND, ge::FORMAT_ND,
797- ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})829+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
830+ ge::FORMAT_ND, ge::FORMAT_ND})
798 .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,831 .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
799 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,832 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
800 ge::FORMAT_ND, ge::FORMAT_ND,833 ge::FORMAT_ND, ge::FORMAT_ND,
@@ -820,7 +853,8 @@ public:
820 ge::FORMAT_ND, ge::FORMAT_ND,853 ge::FORMAT_ND, ge::FORMAT_ND,
821 ge::FORMAT_ND, ge::FORMAT_ND,854 ge::FORMAT_ND, ge::FORMAT_ND,
822 ge::FORMAT_ND, ge::FORMAT_ND,855 ge::FORMAT_ND, ge::FORMAT_ND,
823- ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});856+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
857+ ge::FORMAT_ND, ge::FORMAT_ND});
824 this->Output("y")858 this->Output("y")
825 .ParamType(REQUIRED)859 .ParamType(REQUIRED)
826 .DataType({ge::DT_FLOAT16, ge::DT_BF16, ge::DT_INT8, ge::DT_FLOAT16, ge::DT_BF16, ge::DT_INT8, ge::DT_FLOAT16, ge::DT_BF16,860 .DataType({ge::DT_FLOAT16, ge::DT_BF16, ge::DT_INT8, ge::DT_FLOAT16, ge::DT_BF16, ge::DT_INT8, ge::DT_FLOAT16, ge::DT_BF16,
@@ -848,7 +882,8 @@ public:
848 ge::DT_FLOAT16, ge::DT_FLOAT16,882 ge::DT_FLOAT16, ge::DT_FLOAT16,
849 ge::DT_BF16, ge::DT_BF16,883 ge::DT_BF16, ge::DT_BF16,
850 ge::DT_FLOAT16, ge::DT_FLOAT16,884 ge::DT_FLOAT16, ge::DT_FLOAT16,
851- ge::DT_BF16, ge::DT_BF16, ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT})885+ ge::DT_BF16, ge::DT_BF16, ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT,
886+ ge::DT_BF16, ge::DT_FLOAT16})
852 .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,887 .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
853 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,888 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
854 ge::FORMAT_ND, ge::FORMAT_ND,889 ge::FORMAT_ND, ge::FORMAT_ND,
@@ -874,7 +909,8 @@ public:
874 ge::FORMAT_ND, ge::FORMAT_ND,909 ge::FORMAT_ND, ge::FORMAT_ND,
875 ge::FORMAT_ND, ge::FORMAT_ND,910 ge::FORMAT_ND, ge::FORMAT_ND,
876 ge::FORMAT_ND, ge::FORMAT_ND,911 ge::FORMAT_ND, ge::FORMAT_ND,
877- ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})912+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
913+ ge::FORMAT_ND, ge::FORMAT_ND})
878 .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,914 .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
879 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,915 ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
880 ge::FORMAT_ND, ge::FORMAT_ND,916 ge::FORMAT_ND, ge::FORMAT_ND,
@@ -900,7 +936,8 @@ public:
900 ge::FORMAT_ND, ge::FORMAT_ND,936 ge::FORMAT_ND, ge::FORMAT_ND,
901 ge::FORMAT_ND, ge::FORMAT_ND,937 ge::FORMAT_ND, ge::FORMAT_ND,
902 ge::FORMAT_ND, ge::FORMAT_ND,938 ge::FORMAT_ND, ge::FORMAT_ND,
903- ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});939+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
W
Wwangyongguang3月26日

补充PR描述

likedislike
B
bynshard
3月26日 评论：
940+ ge::FORMAT_ND, ge::FORMAT_ND});
904 941 
905 this->Attr("dtype").AttrType(REQUIRED).Int(-1);942 this->Attr("dtype").AttrType(REQUIRED).Int(-1);
906 this->Attr("compute_type").AttrType(OPTIONAL).Int(-1);943 this->Attr("compute_type").AttrType(OPTIONAL).Int(-1);
Mmatmul/quant_batch_matmul_v4/op_kernel/arch35/cmct/block/block_mmad_mx_weight_from_ub.h+1-1
@@ -16,7 +16,7 @@
16#define QUANT_BATCH_MATMUL_V4_CMCT_BLOCK_BLOCK_MMAD_MX_WEIGHT_FROM_UB_H16#define QUANT_BATCH_MATMUL_V4_CMCT_BLOCK_BLOCK_MMAD_MX_WEIGHT_FROM_UB_H
17 17 
18#include "cmct/block/block_mmad.h"18#include "cmct/block/block_mmad.h"
19-#include "../../quant_batch_matmul_v4_tiling_data.h"19+#include "../../quant_batch_matmul_v4_tiling_data_apt.h"
20 20 
21namespace QuantBatchMatmulV4 {21namespace QuantBatchMatmulV4 {
22namespace Block {22namespace Block {
Mmatmul/quant_batch_matmul_v4/op_kernel/arch35/cmct/block/block_scheduler_swizzle_in_mn_core_nn.h+1-1
@@ -16,7 +16,7 @@
16#ifndef QUANT_BATCH_MATMUL_V4_ARCH35_CMCT_BLOCK_BLOCK_SCHEDULER_SWIZZLE_IN_MN_CORE_NN_H16#ifndef QUANT_BATCH_MATMUL_V4_ARCH35_CMCT_BLOCK_BLOCK_SCHEDULER_SWIZZLE_IN_MN_CORE_NN_H
17#define QUANT_BATCH_MATMUL_V4_ARCH35_CMCT_BLOCK_BLOCK_SCHEDULER_SWIZZLE_IN_MN_CORE_NN_H17#define QUANT_BATCH_MATMUL_V4_ARCH35_CMCT_BLOCK_BLOCK_SCHEDULER_SWIZZLE_IN_MN_CORE_NN_H
18#include "cmct/block/block_scheduler_swizzle_in_mn_core.h"18#include "cmct/block/block_scheduler_swizzle_in_mn_core.h"
19-#include "../../quant_batch_matmul_v4_tiling_data.h"19+#include "../../quant_batch_matmul_v4_tiling_data_apt.h"
20/*20/*
21iterateOrder = 021iterateOrder = 0
22scheduler diagram c:core b:block22scheduler diagram c:core b:block
Mmatmul/quant_batch_matmul_v4/op_kernel/arch35/cmct/kernel/kernel_matmul_mix_with_weight_prologue_nn.h+1-1
@@ -17,7 +17,7 @@
17#define QUANT_BATCH_MATMUL_V4_ARCH35_CMCT_KERNEL_KERNEL_MATMUL_MIX_WITH_WEIGHT_PROLOGUE_NN_H17#define QUANT_BATCH_MATMUL_V4_ARCH35_CMCT_KERNEL_KERNEL_MATMUL_MIX_WITH_WEIGHT_PROLOGUE_NN_H
18 18 
19#include "cmct/kernel/kernel_matmul_mix_with_weight_prologue.h"19#include "cmct/kernel/kernel_matmul_mix_with_weight_prologue.h"
20-#include "../../quant_batch_matmul_v4_tiling_data.h"20+#include "../../quant_batch_matmul_v4_tiling_data_apt.h"
21 21 
22namespace QuantBatchMatmulV4 {22namespace QuantBatchMatmulV4 {
23namespace Kernel {23namespace Kernel {
Mmatmul/quant_batch_matmul_v4/op_kernel/arch35/cmct/prologue/block_prologue_b_cast_scsc_nn.h+1-1
@@ -16,7 +16,7 @@
16#ifndef QUANT_BATCH_MATMUL_V4_ARCH35_CMCT_PROLOGUE_BLOCK_PROLOGUE_B_CAST_SCSC_NN_H16#ifndef QUANT_BATCH_MATMUL_V4_ARCH35_CMCT_PROLOGUE_BLOCK_PROLOGUE_B_CAST_SCSC_NN_H
17#define QUANT_BATCH_MATMUL_V4_ARCH35_CMCT_PROLOGUE_BLOCK_PROLOGUE_B_CAST_SCSC_NN_H17#define QUANT_BATCH_MATMUL_V4_ARCH35_CMCT_PROLOGUE_BLOCK_PROLOGUE_B_CAST_SCSC_NN_H
18#include "cmct/prologue/block_prologue_b_cast_scsc.h"18#include "cmct/prologue/block_prologue_b_cast_scsc.h"
19-#include "../../quant_batch_matmul_v4_tiling_data.h"19+#include "../../quant_batch_matmul_v4_tiling_data_apt.h"
20 20 
21namespace QuantBatchMatmulV4 {21namespace QuantBatchMatmulV4 {
22namespace Prologue {22namespace Prologue {
Mmatmul/quant_batch_matmul_v4/op_kernel/arch35/cmct_convertor.h+1-1
@@ -28,7 +28,7 @@
28#include "kernel_operator.h"28#include "kernel_operator.h"
29#endif29#endif
30#include "lib/std/type_traits.h"30#include "lib/std/type_traits.h"
31-#include "quant_batch_matmul_v4_tiling_data.h"31+#include "quant_batch_matmul_v4_tiling_data_apt.h"
32 32 
33namespace QuantBatchMatmulV4 {33namespace QuantBatchMatmulV4 {
34using AscendC::fp8_e8m0_t;34using AscendC::fp8_e8m0_t;
Mmatmul/quant_batch_matmul_v4/op_kernel/arch35/quant_batch_matmul_v4_perchannel.h+1-1
@@ -17,7 +17,7 @@
17#define QUANT_BATCH_MATMUL_V4_PERCHANNEL_H17#define QUANT_BATCH_MATMUL_V4_PERCHANNEL_H
18 18 
19#include "quant_batch_matmul_v4_reg_base_common.h"19#include "quant_batch_matmul_v4_reg_base_common.h"
20-#include "quant_batch_matmul_v4_tiling_data.h"20+#include "quant_batch_matmul_v4_tiling_data_apt.h"
21 21 
22using matmul::MatmulImpl;22using matmul::MatmulImpl;
23 23 
Amatmul/quant_batch_matmul_v4/op_kernel/arch35/quant_batch_matmul_v4_pertoken_pergroup.h+564-0
@@ -0,0 +1,564 @@
1+/**
2+ * Copyright (c) 2025 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_pergroup.h
13+ * \brief
14+ */
15+ 
16+#ifndef QUANT_BATCH_MATMUL_V4_PERTOKEN_PERGROUP_H
17+#define QUANT_BATCH_MATMUL_V4_PERTOKEN_PERGROUP_H
18+ 
19+#include "../quant_batch_matmul_v4_common.h"
20+ 
21+namespace AscendC {
22+ 
23+template <typename xType, typename wType, typename biasType, typename scaleType, typename yType>
24+class QuantBatchMatmulV4Pergroup : public QuantBatchMatmulV4Common {
25+public:
26+ __aicore__ inline QuantBatchMatmulV4Pergroup(){};
27+ __aicore__ inline void Init(
28+ GM_ADDR x1, GM_ADDR x2, GM_ADDR bias, GM_ADDR x1_scale, GM_ADDR x2_scale, GM_ADDR y_scale, GM_ADDR x1_offset,
29+ GM_ADDR x2_offset, GM_ADDR y_offset, GM_ADDR y, GM_ADDR workspace,
30+ const QuantBatchMatmulV3TilingData* tilingData, TPipe* tPipe)
31+ {
32+ commonInit(tPipe, &(tilingData->matmulTiling));
33+ groupSizeK_ = tilingData->params.groupSizeK;
34+ 
35+ x1Global_.SetGlobalBuffer(reinterpret_cast<__gm__ int8_t*>(x1));
36+ x2Global_.SetGlobalBuffer(reinterpret_cast<__gm__ int8_t*>(x2));
37+ biasGlobal_.SetGlobalBuffer(reinterpret_cast<__gm__ biasType*>(bias));
38+ mmOutGlobal_.SetGlobalBuffer(reinterpret_cast<__gm__ int32_t*>(workspace));
39+ x1ScaleGlobal_.SetGlobalBuffer(reinterpret_cast<__gm__ scaleType*>(x1_scale));
40+ x2ScaleGlobal_.SetGlobalBuffer(reinterpret_cast<__gm__ scaleType*>(x2_scale));
41+ x2OffsetGlobal_.SetGlobalBuffer(reinterpret_cast<__gm__ half*>(x2_offset));
42+ yGlobal_.SetGlobalBuffer(reinterpret_cast<__gm__ yType*>(y));
43+ 
44+ ubCalcM_ = tilingData->params.ubCalcM;
45+ ubCalcN_ = tilingData->params.ubCalcN;
46+ maxAivM_ = CeilDiv(baseM_, VECCORE_NUM);
47+ 
48+ block_.Init(tilingData);
49+ update_.template Init<FORMAT_X1, FORMAT_X2, false, true>(&tilingData->matmulTiling, block_.params_);
50+ nKgroups_ = CeilDiv(block_.matmulTilingData_->Ka, groupSizeK_);
51+ mmObj_.SetSubBlockIdx(0);
52+ mmObj_.Init(&(tilingData->matmulTiling), pipe_);
53+ mmObj_.SetOrgShape(baseM_, baseN_, block_.matmulTilingData_->Ka);
54+ InitLocalBuffers();
55+ };
56+ 
57+ __aicore__ inline void Process()
58+ {
59+ if (blockIdx_ >= usedCoreNum_) {
60+ return;
61+ }
62+ if ASCEND_IS_AIV {
63+ SetFlag<HardEvent::MTE3_V>(EVENT_ID7);
64+ SetFlag<HardEvent::V_MTE2>(eX2Scale_);
65+ SetFlag<HardEvent::V_MTE2>(eX2Offset_);
66+ SetFlag<HardEvent::V_MTE2>(eX1Scale_);
67+ CrossCoreSetFlag<CV_SYNC_FLAG, PIPE_MTE2>(V2C_PING_FLAG);
68+ CrossCoreSetFlag<CV_SYNC_FLAG, PIPE_MTE2>(V2C_PONG_FLAG);
69+ }
70+ // 首块计算，兼容无L2cache切分场景，减少scalar计算
71+ // 反向flag
72+ if ASCEND_IS_AIC {
73+ initEventID();
74+ eAL1_12_ = eAL1Ping12_;
75+ eAL1_21_ = eAL1Ping21_;
76+ eBL1_12_ = eBL1Ping12_;
77+ eBL1_21_ = eBL1Ping21_;
78+ SetFlag<HardEvent::MTE1_MTE2>(eAL1Ping12_);
79+ SetFlag<HardEvent::MTE1_MTE2>(eAL1Pong12_);
80+ SetFlag<HardEvent::MTE1_MTE2>(eBL1Ping12_);
81+ SetFlag<HardEvent::MTE1_MTE2>(eBL1Pong12_);
82+ }
83+ 
84+ bool reverse = true;
85+ block_.InitFirstTileBlockIndex();
86+ OneTileCompute(0, 0);
87+ 
88+ for (uint64_t mTileIndex = 0; mTileIndex < block_.params_.mTileCntL2; mTileIndex++) {
89+ reverse = !reverse;
90+ for (uint64_t nTileIndexTemp = 0; nTileIndexTemp < block_.params_.nTileCntL2; nTileIndexTemp++) {
91+ uint64_t nTileIndex = reverse ? (block_.params_.nTileCntL2 - nTileIndexTemp - 1) : nTileIndexTemp;
92+ if (mTileIndex > 0 || nTileIndex > 0) { // 跳过首块
93+ block_.UpdateBlockCnt(mTileIndex, nTileIndex);
94+ block_.InitBlockIndex();
95+ OneTileCompute(mTileIndex, nTileIndex);
96+ }
97+ }
98+ }
99+ 
100+ // 反向flag
101+ if ASCEND_IS_AIC {
102+ WaitFlag<HardEvent::MTE1_MTE2>(eAL1Ping12_);
103+ WaitFlag<HardEvent::MTE1_MTE2>(eAL1Pong12_);
104+ WaitFlag<HardEvent::MTE1_MTE2>(eBL1Ping12_);
105+ WaitFlag<HardEvent::MTE1_MTE2>(eBL1Pong12_);
106+ CrossCoreWaitFlag(V2C_PING_FLAG);
107+ CrossCoreWaitFlag(V2C_PONG_FLAG);
108+ mmObj_.End();
109+ releaseEventID();
110+ }
111+ if ASCEND_IS_AIV {
112+ WaitFlag<HardEvent::MTE3_V>(EVENT_ID7);
113+ WaitFlag<HardEvent::V_MTE2>(eX2Scale_);
114+ WaitFlag<HardEvent::V_MTE2>(eX2Offset_);
115+ WaitFlag<HardEvent::V_MTE2>(eX1Scale_);
116+ }
117+ }
118+ 
119+ __aicore__ inline void InitLocalBuffers()
120+ {
121+ if ASCEND_IS_AIC {
122+ // Note: assume xtype == wtype
123+ pipe_->InitBuffer(L1Buf_, L1_MAX_SIZE_910B);
124+ uint32_t a_size = baseM_ * groupSizeK_ * stepka_;
125+ uint32_t b_size = baseN_ * groupSizeK_ * stepkb_;
126+ aL1Ping_ = L1Buf_.Get<xType>()[0];
127+ aL1Pong_ = aL1Ping_[a_size];
128+ bL1Ping_ = aL1Pong_[a_size];
129+ bL1Pong_ = bL1Ping_[b_size];
130+ aL1_ = aL1Ping_;
131+ bL1_ = bL1Ping_;
132+ }
133+ 
134+ if ASCEND_IS_AIV {
135+ // mmOutFp32 和 x2scale需要避免bank冲突
136+ // mmOutFp32 和 dequantAccu需要避免bank冲突
137+ pipe_->InitBuffer(ubBuf_, TOTAL_UB_SIZE); // 192KB
138+ uint64_t offset = 0;
139+ dequantAccu = ubBuf_.Get<float>()[0];
140+ dequantOut = ubBuf_.Get<yType>()[0];
141+ offset += maxAivM_ * ubCalcN_;
142+ mmOut = ubBuf_.Get<int32_t>()[offset]; // [ubCalcM_, ubCalcN_]
143+ offset += ubCalcM_ * ubCalcN_;
144+ // 解bank冲突：mmOutFp32 和 dequantAccu 需要隔开256B
145+ // offset += 64;
146+ mmOutFp32 = ubBuf_.Get<float>()[offset];
147+ offset += ubCalcM_ * ubCalcN_;
148+ // ubCalcN_ 32B对齐
149+ X2Scale = ubBuf_.Get<float>()[offset];
150+ offset += ubCalcN_;
151+ bias = ubBuf_.Get<float>()[offset];
152+ offset += ubCalcN_;
153+ X1Scale = ubBuf_.Get<float>()[offset];
154+ offset += CeilAlign(maxAivM_, ALIGN_UNIT_8);
155+ X1BrcbScale = ubBuf_.Get<float>()[offset];
156+ offset += (CeilAlign(maxAivM_, ALIGN_UNIT_8) * ALIGN_UNIT_8);
157+ 
158+ X1Fp32 = ubBuf_.Get<float>()[offset];
159+ X1INT8 = X1Fp32.ReinterpretCast<int8_t>();
160+ offset += ubCalcM_ * groupSizeK_;
161+ X1Half = ubBuf_.Get<float>()[offset].ReinterpretCast<half>();
162+ offset += ubCalcM_ * groupSizeK_ >> 1;
163+ X1RowSum = ubBuf_.Get<float>()[offset];
164+ offset += CeilAlign(maxAivM_, ALIGN_UNIT_8);
165+ X1BCast = ubBuf_.Get<float>()[offset];
166+ offset += ubCalcM_ * ubCalcN_;
167+ X2Offset = ubBuf_.Get<float>()[offset];
168+ offset += ubCalcN_;
169+ X2OffsetFP16 = ubBuf_.Get<float>()[offset].ReinterpretCast<half>();
170+ offset += ubCalcN_ >> 1;
171+ X1BCastMulX2Offset = ubBuf_.Get<float>()[offset];
172+ offset += ubCalcM_ * ubCalcN_;
173+ }
174+ };
175+ 
176+private:
177+ // Continuous-process repeat stride: one hardware iteration covers 256B, with 32B granularity => 256 / 32 = 8.
178+ static constexpr uint8_t CONTINUOUS_PROCESS_STRIDE = 8U;
179+ 
180+ QBmmBlockOffset preload_offset_;
181+ 
182+ TEventID eAL1_12_;
183+ TEventID eAL1_21_;
184+ TEventID eBL1_12_;
185+ TEventID eBL1_21_;
186+ 
187+ bool aL1PingFlag_ = true;
188+ bool bL1PingFlag_ = true;
189+ bool wsPingFlag_ = true;
190+ bool ubPingFlag_ = true;
191+ bool x2ScalePingFlag_ = true;
192+ 
193+ GlobalTensor<int8_t> x1Global_;
194+ GlobalTensor<int8_t> x2Global_;
195+ GlobalTensor<biasType> biasGlobal_;
196+ GlobalTensor<int32_t> mmOutGlobal_;
197+ // GlobalTensor<float16_t> mmOutGlobal_;
198+ GlobalTensor<scaleType> x1ScaleGlobal_;
199+ GlobalTensor<scaleType> x2ScaleGlobal_;
200+ GlobalTensor<half> x2OffsetGlobal_;
201+ GlobalTensor<yType> yGlobal_;
202+ // define pingpong buf
203+ TBuf<TPosition::A1> L1Buf_;
204+ LocalTensor<xType> aL1_;
205+ LocalTensor<wType> bL1_;
206+ LocalTensor<xType> aL1Ping_;
207+ LocalTensor<xType> aL1Pong_;
208+ LocalTensor<wType> bL1Ping_;
209+ LocalTensor<wType> bL1Pong_;
210+ 
211+ // Vec buffer
212+ // 不切N ubCalcN_ = BaseN
213+ // mmOut, mmOutFp32, dequantOut可共享内存
214+ TBuf<> ubBuf_;
215+ LocalTensor<int32_t> mmOut; // [ubCalcM_, ubCalcN_]
216+ LocalTensor<float> mmOutFp32; // [ubCalcM_, ubCalcN_]
217+ LocalTensor<float> dequantAccu; // [AivM_, ubCalcN_]
218+ LocalTensor<yType> dequantOut; // [AivM_, ubCalcN_]
219+ LocalTensor<float> X1BrcbScale; // [AivM_, 8]
220+ LocalTensor<float> X1Scale; // [AivM_]
221+ LocalTensor<float> X2Scale; // [ubCalcN_] // 32B align
222+ LocalTensor<float> bias; // [ubCalcN_]
223+ 
224+ LocalTensor<int8_t> X1INT8; // [ubCalcM_, groupsizeK]
225+ LocalTensor<float> X2Offset; // [ubCalcN_]
226+ LocalTensor<half> X2OffsetFP16; // [ubCalcN_]
227+ LocalTensor<float> X1Fp32; // [ubCalcM_, groupsizeK]
228+ LocalTensor<half> X1Half; // [ubCalcM_, groupsizeK]
229+ LocalTensor<float> X1RowSum; // [ubCalcM_]
230+ LocalTensor<float> X1BCast; // [ubCalcM_, ubCalcN_]
231+ LocalTensor<float> X1BCastMulX2Offset; // [ubCalcM_, ubCalcN_]
232+ 
233+ TEventID eX1Scale_ = EVENT_ID0;
234+ TEventID eX2Scale_ = EVENT_ID1;
235+ TEventID eMmOut_ = EVENT_ID2;
236+ TEventID eX2Offset_ = EVENT_ID3;
237+ TEventID eX1_ = EVENT_ID4;
238+ 
239+ uint32_t maxAivM_;
240+ 
241+ // tilingData
242+ using inputX1Type = MatmulType<TPosition::TSCM, CubeFormat::NZ, xType, false>;
243+ using inputX2Type = MatmulType<TPosition::TSCM, CubeFormat::NZ, wType, true>;
244+ using inputBiasType = MatmulType<TPosition::GM, CubeFormat::ND, biasType, false>;
245+ using outputYType = MatmulType<TPosition::GM, CubeFormat::ND, int32_t, false>;
246+ MatmulImpl<inputX1Type, inputX2Type, outputYType, inputBiasType, CFG_MDL> mmObj_;
247+ 
248+ __aicore__ inline void CopyInX1Scale(uint32_t curAivM)
249+ {
250+ // X1Scale [AivM_] // 32B align
251+ uint32_t offset = offset_.offsetPertoken + subBlockIdx_ * maxAivM_;
252+ WaitFlag<HardEvent::V_MTE2>(eX1Scale_);
253+ DataCopyPad(
254+ X1Scale, x1ScaleGlobal_[offset], {/*blk_count*/ 1, /*blk_len*/ uint32_t(curAivM * sizeof(float)), 0, 0, 0},
255+ {false, 0, 0, 0});
256+ SetFlag<HardEvent::MTE2_V>(eX1Scale_);
257+ }
258+ 
259+ __aicore__ inline void CopyInX2Scale(uint32_t kidx, uint32_t curAivN)
260+ {
261+ // X2Scale; // [ubCalcN_] // 32B align
262+ // x2ScaleGlobal_ [nKgroups_, n]
263+ uint32_t offset = kidx * block_.matmulTilingData_->N + offset_.offsetScale;
264+ WaitFlag<HardEvent::V_MTE2>(eX2Scale_);
265+ DataCopyPad(
266+ X2Scale, x2ScaleGlobal_[offset], {/*blk_count*/ 1, /*blk_len*/ uint32_t(curAivN * sizeof(float)), 0, 0, 0},
267+ {false, 0, 0, 0});
268+ SetFlag<HardEvent::MTE2_V>(eX2Scale_);
269+ }
270+ 
271+ __aicore__ inline void CopyInX2Offset(uint32_t kidx, uint32_t curAivN)
272+ {
273+ // X2Offset; // [ubCalcN_] // 32B align
274+ // x2SOffset_ [nKgroups_, n]
275+ WaitFlag<HardEvent::V_MTE2>(eX2Offset_);
276+ uint32_t offset = kidx * block_.matmulTilingData_->N + offset_.offsetScale;
277+ DataCopyPad(
278+ X2OffsetFP16, x2OffsetGlobal_[offset],
279+ {/*blk_count*/ 1, /*blk_len*/ uint32_t(curAivN * sizeof(half)), 0, 0, 0}, {false, 0, 0, 0});
280+ SetFlag<HardEvent::MTE2_V>(eX2Offset_);
281+ }
282+ 
283+ __aicore__ inline void CopyInX1INT8(uint32_t kidx, uint32_t midx, uint32_t curAivM)
284+ {
285+ // CopyIn X1
286+ // [m, k]
287+ uint32_t X1_offset = offset_.offsetA + kidx * groupSizeK_ +
288+ (subBlockIdx_ * maxAivM_ + ubCalcM_ * midx) * block_.matmulTilingData_->Ka;
289+ uint16_t blockCount = curAivM;
290+ uint16_t dstGap = 0;
291+ // INT8: 1 byte per element, no packing
292+ uint16_t blockLen = groupSizeK_ / DATA_BLOCK_LEN;
293+ uint16_t srcGap = (block_.matmulTilingData_->Ka - groupSizeK_) / DATA_BLOCK_LEN;
294+ DataCopy(X1INT8, x1Global_[X1_offset], {blockCount, blockLen, srcGap, dstGap});
295+ SetFlag<HardEvent::MTE2_V>(eX1_);
296+ WaitFlag<HardEvent::MTE2_V>(eX1_);
297+ }
298+ 
299+ __aicore__ inline void X2OffsetProcess(uint32_t curAivM, uint32_t curAivN, uint32_t kidx, uint32_t midx)
300+ {
301+ CopyInX1INT8(kidx, midx, curAivM);
302+ Cast<half, int8_t>(X1Half, X1INT8, RoundMode::CAST_NONE, curAivM * groupSizeK_);
303+ PipeBarrier<PIPE_V>();
304+ Cast<float, half>(X1Fp32, X1Half, RoundMode::CAST_NONE, groupSizeK_ * curAivM);
305+ PipeBarrier<PIPE_V>();
306+ // // use X1BCastMulX2Offset as sharedTmpBuffer
307+ // // 文档未说明此处tmpbuffer会用多大，但已X1BCastMulX2Offset为起点的ub空间不会踩踏前面有用的空间
308+ // // 后面再用到X1BCastMulX2Offset时通过pipe_V隔离
309+ uint32_t shape[] = {curAivM, groupSizeK_};
310+ ReduceSum<float, Pattern::Reduce::AR, false>(
311+ X1RowSum, X1Fp32, /*tmp*/ X1BCastMulX2Offset.ReinterpretCast<uint8_t>(), shape, true);
312+ PipeBarrier<PIPE_V>();
313+ uint32_t dstShape_[] = {curAivM, curAivN};
314+ uint32_t srcShape_[] = {curAivM, 1};
315+ LocalTensor<uint8_t> sharedTmpBuffer = X1BCastMulX2Offset.ReinterpretCast<uint8_t>();
316+ constexpr int32_t BROADCAST_DIM = 2;
317+ constexpr int32_t BROADCAST_AXIS = 1;
318+ Broadcast<float, BROADCAST_DIM, BROADCAST_AXIS, false>(X1BCast, X1RowSum, dstShape_, srcShape_,
319+ sharedTmpBuffer);
320+ PipeBarrier<PIPE_V>();
321+ int32_t resN = curAivN;
322+ uint8_t repeatStride = CeilDiv(curAivN, ALIGN_UNIT_8);
323+ if (midx == 0) {
324+ WaitFlag<HardEvent::MTE2_V>(eX2Offset_);
325+ Cast<float, half>(X2Offset, X2OffsetFP16, RoundMode::CAST_NONE, curAivN);
326+ PipeBarrier<PIPE_V>();
327+ }
328+ while (resN >= NUM_ELEMENTS_PER_ITER) {
329+ Mul<float, false>(
330+ X1BCastMulX2Offset[curAivN - resN], X2Offset[curAivN - resN], X1BCast[curAivN - resN], 0UL, curAivM,
331+ {1, 1, 1, repeatStride, 0, repeatStride});
332+ resN = resN - NUM_ELEMENTS_PER_ITER;
333+ }
334+ if (resN > 0) {
335+ Mul<float, true>(
336+ X1BCastMulX2Offset[curAivN - resN], X2Offset[curAivN - resN], X1BCast[curAivN - resN], resN, curAivM,
337+ {1, 1, 1, repeatStride, 0, repeatStride});
338+ }
339+ PipeBarrier<PIPE_V>();
340+ }
341+ 
342+ __aicore__ inline void DequantAndAccu(uint32_t curAivM, uint32_t curAivN, uint32_t kidx, uint32_t midx)
343+ {
344+ uint8_t dstRepeatStride = CONTINUOUS_PROCESS_STRIDE;
345+ Cast<float, int32_t>(mmOutFp32, mmOut, RoundMode::CAST_NONE, curAivM * curAivN);
346+ PipeBarrier<PIPE_V>();
347+ X2OffsetProcess(curAivM, curAivN, kidx, midx);
348+ uint8_t src0RepeatStride = CONTINUOUS_PROCESS_STRIDE;
349+ uint8_t src1RepeatStride = CONTINUOUS_PROCESS_STRIDE;
350+ Sub<float, false>(
351+ mmOutFp32, mmOutFp32, X1BCastMulX2Offset, 0UL, FOLD4 * curAivM,
352+ {1, 1, 1, dstRepeatStride, src0RepeatStride, src1RepeatStride});
353+ PipeBarrier<PIPE_V>();
354+ 
355+ if (midx == 0) {
356+ WaitFlag<HardEvent::MTE2_V>(eX2Scale_);
357+ }
358+ for (uint32_t i = 0; i < curAivM; ++i) {
359+ Mul(mmOutFp32[i * curAivN], X2Scale, mmOutFp32[i * curAivN], curAivN);
360+ PipeBarrier<PIPE_V>();
361+ }
362+ uint32_t moffset = ubCalcM_ * midx * curAivN;
363+ Add(dequantAccu[moffset], mmOutFp32, dequantAccu[moffset], curAivM * curAivN);
364+ PipeBarrier<PIPE_ALL>();
365+ }
366+ 
367+ __aicore__ inline void CopyInMMOut(uint32_t midx, uint32_t realM, uint32_t CurAicN)
368+ {
369+ GlobalTensor<int32_t> src = mmOutGlobal_[wsPingFlag_ ? offsetWorkspaceC_ : offsetWorkspacePong_];
370+ uint32_t moffset = (subBlockIdx_ * maxAivM_ + ubCalcM_ * midx) * CurAicN;
371+ src = src[moffset];
372+ PipeBarrier<PIPE_MTE2>();
373+ DataCopy(mmOut, src, realM * CurAicN);
374+ SetFlag<HardEvent::MTE2_V>(eMmOut_);
375+ WaitFlag<HardEvent::MTE2_V>(eMmOut_);
376+ }
377+ 
378+ __aicore__ inline void CastCopyOut(uint32_t curAivM, uint32_t curAivN)
379+ {
380+ WaitFlag<HardEvent::MTE2_V>(eX1Scale_);
381+ uint8_t dstRepeatStride = CONTINUOUS_PROCESS_STRIDE;
382+ Brcb(X1BrcbScale, X1Scale, CeilDiv(curAivM, ALIGN_UNIT_8), {1, dstRepeatStride});
383+ PipeBarrier<PIPE_V>();
384+ int32_t resN = curAivN;
385+ uint8_t repeatStride = CeilDiv(curAivN, ALIGN_UNIT_8);
386+ while (resN >= NUM_ELEMENTS_PER_ITER) {
387+ Mul(dequantAccu[curAivN - resN], X1BrcbScale, dequantAccu[curAivN - resN], NUM_ELEMENTS_PER_ITER, curAivM,
388+ {1, 0, 1, repeatStride, 1, repeatStride});
389+ resN = resN - NUM_ELEMENTS_PER_ITER;
390+ }
391+ if (resN > 0) {
392+ Mul(dequantAccu[curAivN - resN], X1BrcbScale, dequantAccu[curAivN - resN], resN, curAivM,
393+ {1, 0, 1, repeatStride, 1, repeatStride});
394+ }
395+ PipeBarrier<PIPE_V>();
396+ SetFlag<HardEvent::V_MTE2>(eX1Scale_);
397+ Cast(dequantOut, dequantAccu, RoundMode::CAST_RINT, curAivM * curAivN);
398+ SetFlag<HardEvent::V_MTE3>(EVENT_ID7);
399+ uint64_t stride = DATA_BLOCK / sizeof(yType);
400+ DataCopyParams copyParams{
401+ uint16_t(curAivM), uint16_t(curAivN / stride), // * sizeof(bfloat16) / 32B
402+ 0, uint16_t((block_.matmulTilingData_->N - curAivN) / stride)};
403+ uint64_t offset = offset_.offsetC;
404+ offset += subBlockIdx_ * maxAivM_ * block_.matmulTilingData_->N;
405+ WaitFlag<HardEvent::V_MTE3>(EVENT_ID7);
406+ DataCopy(yGlobal_[offset], dequantOut, copyParams);
407+ SetFlag<HardEvent::MTE3_V>(EVENT_ID7);
408+ return;
409+ }
410+ 
411+ __aicore__ inline void shiftFlagAL1()
412+ {
413+ aL1PingFlag_ = !aL1PingFlag_;
414+ aL1_ = aL1PingFlag_ ? aL1Ping_ : aL1Pong_;
415+ eAL1_12_ = aL1PingFlag_ ? eAL1Ping12_ : eAL1Pong12_;
416+ eAL1_21_ = aL1PingFlag_ ? eAL1Ping21_ : eAL1Pong21_;
417+ }
418+ 
419+ __aicore__ inline void shiftFlagBL1()
420+ {
421+ bL1PingFlag_ = !bL1PingFlag_;
422+ bL1_ = bL1PingFlag_ ? bL1Ping_ : bL1Pong_;
423+ eBL1_12_ = bL1PingFlag_ ? eBL1Ping12_ : eBL1Pong12_;
424+ eBL1_21_ = bL1PingFlag_ ? eBL1Ping21_ : eBL1Pong21_;
425+ }
426+ 
427+ __aicore__ inline void CubeProcess(uint32_t CurAicM, uint32_t CurAicN, uint32_t kidx)
428+ {
429+ if (kidx % stepka_ == 0) {
430+ WaitFlag<HardEvent::MTE1_MTE2>(eAL1_12_);
431+ Nd2NzParams nd2nzParamsA;
432+ nd2nzParamsA.ndNum = 1;
433+ nd2nzParamsA.srcNdMatrixStride = 0;
434+ nd2nzParamsA.dstNzNStride = 1;
435+ nd2nzParamsA.dstNzMatrixStride = 0;
436+ nd2nzParamsA.dstNzC0Stride = CeilAlign(CurAicM, ALIGN_UNIT_16);
437+ nd2nzParamsA.nValue = CurAicM;
438+ // INT8: 1 byte per element
439+ nd2nzParamsA.srcDValue = block_.matmulTilingData_->Ka;
440+ nd2nzParamsA.dValue = stepka_ * groupSizeK_;
441+ DataCopy(
442+ aL1_.template ReinterpretCast<int8_t>(), x1Global_[offset_.offsetA + kidx * groupSizeK_], nd2nzParamsA);
443+ SetFlag<HardEvent::MTE2_MTE1>(eAL1_21_);
444+ WaitFlag<HardEvent::MTE2_MTE1>(eAL1_21_);
445+ }
446+ if (kidx % stepkb_ == 0) {
447+ WaitFlag<HardEvent::MTE1_MTE2>(eBL1_12_);
448+ Nd2NzParams nd2nzParamsB;
449+ nd2nzParamsB.ndNum = 1;
450+ nd2nzParamsB.srcNdMatrixStride = 0;
451+ nd2nzParamsB.dstNzNStride = 1;
452+ nd2nzParamsB.dstNzMatrixStride = 0;
453+ nd2nzParamsB.dstNzC0Stride = CeilAlign(CurAicN, ALIGN_UNIT_16);
454+ nd2nzParamsB.nValue = CurAicN;
455+ // INT8: 1 byte per element
456+ nd2nzParamsB.srcDValue = block_.matmulTilingData_->Ka;
457+ nd2nzParamsB.dValue = stepkb_ * groupSizeK_;
458+ DataCopy(
459+ bL1_.template ReinterpretCast<int8_t>(), x2Global_[offset_.offsetB + kidx * groupSizeK_], nd2nzParamsB);
460+ SetFlag<HardEvent::MTE2_MTE1>(eBL1_21_);
461+ WaitFlag<HardEvent::MTE2_MTE1>(eBL1_21_);
462+ }
463+ 
464+ mmObj_.SetTensorA(aL1_[kidx % stepka_ * groupSizeK_ * CeilAlign(CurAicM, ALIGN_UNIT_16)], /*transa*/ false);
465+ mmObj_.SetTensorB(bL1_[kidx % stepkb_ * groupSizeK_ * baseN_], /*transb*/ true);
466+ 
467+ mmObj_.SetSingleShape(block_.params_.singleCoreM, block_.params_.singleCoreN, groupSizeK_);
468+ mmObj_.Iterate(false);
469+ const bool isLastKGroup = (kidx + 1 == nKgroups_);
470+ if (((kidx + 1) % stepka_ == 0) || isLastKGroup) {
471+ SetFlag<HardEvent::MTE1_MTE2>(eAL1_12_);
472+ shiftFlagAL1();
473+ }
474+ if (((kidx + 1) % stepkb_ == 0) || isLastKGroup) {
475+ SetFlag<HardEvent::MTE1_MTE2>(eBL1_12_);
476+ shiftFlagBL1();
477+ }
478+ const uint16_t v2cFlag = V2C_PING_FLAG | wsPingFlag_;
479+ CrossCoreWaitFlag(v2cFlag);
480+ PipeBarrier<PIPE_FIX>();
481+ mmObj_.GetTensorC(mmOutGlobal_[wsPingFlag_ ? offsetWorkspaceC_ : offsetWorkspacePong_], 0, true);
482+ PipeBarrier<PIPE_FIX>();
483+ const uint16_t c2vFlag = C2V_PING_FLAG | wsPingFlag_;
484+ CrossCoreSetFlag<CV_SYNC_FLAG, PIPE_FIX>(c2vFlag);
485+ }
486+ 
487+ __aicore__ inline void BasicMMDequantCompute(uint32_t CurAicM, uint32_t CurAicN, uint32_t kidx, int32_t CurAivM)
488+ {
489+ if ASCEND_IS_AIC {
490+ CubeProcess(CurAicM, CurAicN, kidx);
491+ }
492+ if ASCEND_IS_AIV {
493+ if (CurAivM > 0) {
494+ uint64_t mloop = CeilDiv(CurAivM, ubCalcM_);
495+ uint32_t tailM = CurAivM - ubCalcM_ * (mloop - 1);
496+ for (uint32_t mUbLoopIdx = 0; mUbLoopIdx < mloop; ++mUbLoopIdx) {
497+ if (mUbLoopIdx == 0) {
498+ CrossCoreWaitFlag(C2V_PING_FLAG | wsPingFlag_);
499+ }
500+ CopyInMMOut(mUbLoopIdx, mUbLoopIdx == mloop - 1 ? tailM : ubCalcM_, CurAicN);
501+ DequantAndAccu(mUbLoopIdx == mloop - 1 ? tailM : ubCalcM_, CurAicN, kidx, mUbLoopIdx);
502+ if (mUbLoopIdx == mloop - 1) {
503+ SetFlag<HardEvent::V_MTE2>(eX2Scale_);
504+ SetFlag<HardEvent::V_MTE2>(eX2Offset_);
505+ CrossCoreSetFlag<CV_SYNC_FLAG, PIPE_MTE2>(V2C_PING_FLAG | wsPingFlag_);
506+ }
507+ }
508+ } else {
509+ CrossCoreWaitFlag(C2V_PING_FLAG | wsPingFlag_);
510+ CrossCoreSetFlag<CV_SYNC_FLAG, PIPE_MTE2>(V2C_PING_FLAG | wsPingFlag_);
511+ }
512+ }
513+ wsPingFlag_ = !wsPingFlag_;
514+ }
515+ 
516+ __aicore__ inline uint32_t GetCurAivM()
517+ { // 计算m轴尾�?
518+ uint32_t curAivM = maxAivM_;
519+ if (block_.params_.singleCoreM != baseM_) {
520+ if (subBlockIdx_ == 0) {
521+ curAivM = block_.params_.singleCoreM > curAivM ? curAivM : block_.params_.singleCoreM;
522+ } else {
523+ curAivM = block_.params_.singleCoreM > curAivM ? block_.params_.singleCoreM - curAivM : 0UL;
524+ }
525+ }
526+ return curAivM;
527+ }
528+ 
529+ __aicore__ inline void OneTileCompute(uint64_t mTileIndex, uint64_t nTileIndex)
530+ {
531+ for (uint64_t j = 0; j < block_.realRound_; j++) {
532+ // 更新此次基本块的大小和输入输出地址
533+ update_.template UpdateBlockParamsAndCalcGmOffset<FORMAT_X1, FORMAT_X2, /*atrans*/ false, /*btrans*/ true>(
534+ block_.params_, offset_, mTileIndex, nTileIndex);
535+ uint32_t aivM = GetCurAivM();
536+ if ASCEND_IS_AIV {
537+ if (aivM > 0) {
538+ WaitFlag<HardEvent::MTE3_V>(EVENT_ID7);
539+ Duplicate(dequantAccu, float(0), aivM* ubCalcN_);
540+ PipeBarrier<PIPE_V>();
541+ CopyInX1Scale(aivM);
542+ }
543+ }
544+ for (uint64_t k = 0; k < nKgroups_; k++) {
545+ if ASCEND_IS_AIV {
546+ if (aivM > 0) {
547+ CopyInX2Scale(k, block_.params_.singleCoreN);
548+ CopyInX2Offset(k, block_.params_.singleCoreN);
549+ }
550+ }
551+ BasicMMDequantCompute(block_.params_.singleCoreM, block_.params_.singleCoreN, k, aivM);
552+ }
553+ block_.UpdateBlockIndex();
554+ if ASCEND_IS_AIV {
555+ if (aivM > 0) {
556+ CastCopyOut(aivM, block_.params_.singleCoreN);
557+ }
558+ }
559+ }
560+ }
561+};
562+ 
563+} // namespace AscendC
564+#endif
Mmatmul/quant_batch_matmul_v4/op_kernel/arch35/quant_batch_matmul_v4_reg_base_common.h+1-1
@@ -18,7 +18,7 @@
18 18 
19#include "lib/matmul_intf.h"19#include "lib/matmul_intf.h"
20#include "quant_batch_matmul_v4_constant.h"20#include "quant_batch_matmul_v4_constant.h"
21-#include "quant_batch_matmul_v4_tiling_data.h"21+#include "quant_batch_matmul_v4_tiling_data_apt.h"
22#include "quant_batch_matmul_v4_vf.h"22#include "quant_batch_matmul_v4_vf.h"
23 23 
24using AscendC::BLOCK_CUBE;24using AscendC::BLOCK_CUBE;
Rmatmul/quant_batch_matmul_v4/op_kernel/arch35/quant_batch_matmul_v4_tiling_data.h→matmul/quant_batch_matmul_v4/op_kernel/arch35/quant_batch_matmul_v4_tiling_data_apt.h+6-4
@@ -9,12 +9,12 @@
9 */9 */
10 10 
11/*!11/*!
12- * \file quant_batch_matmul_v4_tiling_data.h12+ * \file quant_batch_matmul_v4_tiling_data_apt.h
13 * \brief13 * \brief
14 */14 */
15 15 
16-#ifndef QUANT_BATCH_MATMUL_V4_TILING_DATA_H16+#ifndef QUANT_BATCH_MATMUL_V4_TILING_DATA_APT_H
17-#define QUANT_BATCH_MATMUL_V4_TILING_DATA_H17+#define QUANT_BATCH_MATMUL_V4_TILING_DATA_APT_H
18#include "kernel_tiling/kernel_tiling.h" // TCubeTiling结构体通过C++语法定义18#include "kernel_tiling/kernel_tiling.h" // TCubeTiling结构体通过C++语法定义
19 19 
20#ifndef __CCE_AICORE__20#ifndef __CCE_AICORE__
@@ -174,4 +174,6 @@ struct QuantBatchMatmulV4TilingDataParams {
174};174};
175#pragma pack(pop)175#pragma pack(pop)
176} // namespace qbmmv4_tiling176} // namespace qbmmv4_tiling
177-#endif // QUANT_BATCH_MATMUL_V4_TILING_DATA_H177+#endif // QUANT_BATCH_MATMUL_V4_TILING_DATA_APT_H
178+ 
179+ 
Mmatmul/quant_batch_matmul_v4/op_kernel/arch35/quant_batch_matmul_v4_tiling_key.h+10-2文件内容审核中，请稍后刷新重试
Mmatmul/quant_batch_matmul_v4/op_kernel/quant_batch_matmul_v4_apt.cpp+56-2
@@ -30,6 +30,15 @@
30#define CMCT_PRETILE_INT8_INT8_BF16 030#define CMCT_PRETILE_INT8_INT8_BF16 0
31#endif31#endif
32 32 
33+#if (defined(ORIG_DTYPE_X1) && defined(DT_INT4) && (ORIG_DTYPE_X1 == DT_INT4)) && \
34+ (defined(ORIG_DTYPE_X2) && defined(DT_INT4) && (ORIG_DTYPE_X2 == DT_INT4)) && \
35+ (defined(ORIG_DTYPE_X1_SCALE) && defined(DT_FLOAT) && (ORIG_DTYPE_X1_SCALE == DT_FLOAT)) && \
36+ (defined(ORIG_DTYPE_X2_SCALE) && defined(DT_FLOAT) && (ORIG_DTYPE_X2_SCALE == DT_FLOAT))
37+#define CMCT_PRETILE_INT4_INT4_ASYMMETRICAL 1
38+#else
39+#define CMCT_PRETILE_INT4_INT4_ASYMMETRICAL 0
40+#endif
41+ 
33// if run with ttk without bias, can't get DTYPE_BIAS macro42// if run with ttk without bias, can't get DTYPE_BIAS macro
34#ifndef DTYPE_BIAS43#ifndef DTYPE_BIAS
35#if CMCT_PRETILE_INT8_INT8_BF1644#if CMCT_PRETILE_INT8_INT8_BF16
@@ -51,20 +60,29 @@
51#endif60#endif
52 61 
53#include "arch35/quant_batch_matmul_v4_tiling_key.h"62#include "arch35/quant_batch_matmul_v4_tiling_key.h"
54-#include "arch35/quant_batch_matmul_v4_tiling_data.h"63+#include "arch35/quant_batch_matmul_v4_tiling_data_apt.h"
55#if !(defined(__NPU_ARCH__) && (__NPU_ARCH__ == 5102))64#if !(defined(__NPU_ARCH__) && (__NPU_ARCH__ == 5102))
56// define DTYPE_X2 should before cmct65// define DTYPE_X2 should before cmct
66+#if CMCT_PRETILE_INT4_INT4_ASYMMETRICAL
67+#include "quant_batch_matmul_v4_tiling_data.h"
68+#include "../quant_batch_matmul_v3/arch35/qbmm_int4_to_int8_preprocess.h"
69+#include "quant_batch_matmul_v4_constant.h"
70+#include "arch35/quant_batch_matmul_v4_pertoken_pergroup.h"
71+#else
57#include "arch35/cmct_convertor.h"72#include "arch35/cmct_convertor.h"
58#include "arch35/quant_batch_matmul_v4_constant.h"73#include "arch35/quant_batch_matmul_v4_constant.h"
59#include "arch35/quant_batch_matmul_v4_perchannel.h"74#include "arch35/quant_batch_matmul_v4_perchannel.h"
60#include "../quant_batch_matmul_v3/arch35/qbmm_mix_pertile_cmct.h"75#include "../quant_batch_matmul_v3/arch35/qbmm_mix_pertile_cmct.h"
76+#endif
61#else77#else
62#include "../quant_batch_matmul_v3/quant_batch_matmul_v3_base.h"78#include "../quant_batch_matmul_v3/quant_batch_matmul_v3_base.h"
63#include "../quant_batch_matmul_v3/arch35/qbmm_cube_on_the_fly.h"79#include "../quant_batch_matmul_v3/arch35/qbmm_cube_on_the_fly.h"
64#include "../quant_batch_matmul_v3/arch35/qbmm_cube_on_the_fly_al1_full_load.h"80#include "../quant_batch_matmul_v3/arch35/qbmm_cube_on_the_fly_al1_full_load.h"
65using namespace AscendC;81using namespace AscendC;
66#endif82#endif
83+ 
67#if !(defined(__NPU_ARCH__) && (__NPU_ARCH__ == 5102))84#if !(defined(__NPU_ARCH__) && (__NPU_ARCH__ == 5102))
85+#if !CMCT_PRETILE_INT4_INT4_ASYMMETRICAL
68using namespace QuantBatchMatmulV4;86using namespace QuantBatchMatmulV4;
69namespace QuantBatchMatmulV4 {87namespace QuantBatchMatmulV4 {
70namespace Arch35 {88namespace Arch35 {
@@ -89,6 +107,7 @@ __aicore__ inline void InvokeWeightQuantBmmOpImpl(GM_ADDR x1, GM_ADDR x2, GM_ADD
89} // namespace Arch35107} // namespace Arch35
90} // namespace QuantBatchMatmulV4108} // namespace QuantBatchMatmulV4
91#endif109#endif
110+#endif
92 111 
93#define QBMM_QUANT_GB_IMPL_CLASS(xLayout, wLayout, yLayout) \112#define QBMM_QUANT_GB_IMPL_CLASS(xLayout, wLayout, yLayout) \
94 do { \113 do { \
@@ -122,7 +141,41 @@ __global__ __aicore__ void quant_batch_matmul_v4(
122 Cmct::Gemm::layout::RowMajor, Cmct::Gemm::layout::ColumnMajor, Cmct::Gemm::layout::RowMajorAlign);141 Cmct::Gemm::layout::RowMajor, Cmct::Gemm::layout::ColumnMajor, Cmct::Gemm::layout::RowMajorAlign);
123 }142 }
124#else143#else
125- REGISTER_TILING_DEFAULT(qbmmv4_tiling::QuantBatchMatmulV4TilingDataParams);144+#if CMCT_PRETILE_INT4_INT4_ASYMMETRICAL
145+ REGISTER_TILING_DEFAULT(QuantBatchMatmulV3TilingData);
146+ if (QUANT_TYPE == QBMMV4_INT4_ASYMMETRICAL) {
147+ GET_TILING_DATA(tilingData, tiling);
148+ AscendC::TPipe tPipe;
149+ 
150+ GM_ADDR userWS = AscendC::GetUserWorkspace(workspace);
151+ auto* tilingData_ = static_cast<QuantBatchMatmulV3TilingData*>(&tilingData);
152+ uint64_t m = tilingData_->matmulTiling.M;
153+ uint64_t n = tilingData_->matmulTiling.N;
154+ uint64_t k = tilingData_->matmulTiling.Ka;
155+ uint64_t batchC = tilingData_->params.batchC;
156+ QbmmInt4ToInt8Preprocess preprocessOp;
157+ if ASCEND_IS_AIV {
158+ preprocessOp.Init(x1, x2, userWS, tPipe, m, n, k, batchC);
159+ preprocessOp.Process();
160+ tPipe.Reset();
161+ }
162+ SyncAll<false>();
163+ 
164+ constexpr uint64_t ALIGN_SIZE_128 = 128;
165+ uint64_t x1TotalElems = batchC * m * k;
166+ uint64_t x2TotalElems = batchC * k * n;
167+ uint64_t offsetA = 0;
168+ uint64_t offsetB = DequantBmm::Align(x1TotalElems * sizeof(int8_t), ALIGN_SIZE_128);
169+ uint64_t offsetMMOut = offsetB + DequantBmm::Align(x2TotalElems * sizeof(int8_t), ALIGN_SIZE_128);
170+ AscendC::QuantBatchMatmulV4Pergroup<int8_t, int8_t, float, float, DTYPE_Y> op;
171+ op.Init(
172+ userWS + offsetA, userWS + offsetB, bias, x1_scale, x2_scale, y_scale, x1_offset, x2_offset, y_offset, y,
173+ userWS + offsetMMOut, tilingData_, &tPipe);
174+ op.Process();
175+ tPipe.Destroy();
176+ }
177+#else
178+ REGISTER_TILING_DEFAULT(DequantBmm::QuantBatchMatmulV3TilingDataParams);
126 if (QUANT_TYPE == QBMMV4_PER_GROUP) {179 if (QUANT_TYPE == QBMMV4_PER_GROUP) {
127 constexpr bool isTransA = TRANS == QBMMV4_A_TRANS || TRANS == QBMMV4_ALL_TRANS;180 constexpr bool isTransA = TRANS == QBMMV4_A_TRANS || TRANS == QBMMV4_ALL_TRANS;
128 constexpr bool isTransB = TRANS == QBMMV4_B_TRANS || TRANS == QBMMV4_ALL_TRANS;181 constexpr bool isTransB = TRANS == QBMMV4_B_TRANS || TRANS == QBMMV4_ALL_TRANS;
@@ -134,6 +187,7 @@ __global__ __aicore__ void quant_batch_matmul_v4(
134 QuantBatchMatmulV4::InvokeKernel<WEIGHTNZ>(KERNEL_PARAMS);187 QuantBatchMatmulV4::InvokeKernel<WEIGHTNZ>(KERNEL_PARAMS);
135 }188 }
136#endif189#endif
190+#endif
137 191 
138#else192#else
139 GM_ADDR userWS = AscendC::GetUserWorkspace(workspace);193 GM_ADDR userWS = AscendC::GetUserWorkspace(workspace);
Mmatmul/quant_batch_matmul_v4/tests/ut/op_host/test_aclnn_quant_matmul_v5_api.cpp+34-0
@@ -1246,3 +1246,37 @@ TEST_F(l2_QuantBatchMatmulV5_test_910B2, ascend910B2_test_case_A4W4_PERGROUP_4)
1246 aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);1246 aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
1247 EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);1247 EXPECT_EQ(aclRet, ACLNN_ERR_PARAM_INVALID);
1248}1248}
1249+ 
1250+TEST_F(l2_QuantBatchMatmulV5_test_950, ascend910D_A4W4_test_case_PERTOKEN_PERGROUP_0)
1251+{
1252+ // A4W4 pertoken-pergroup_scale out: bf16
1253+ op::NpuArchManager archManager(NpuArch::DAV_3510);
1254+ TensorDesc x1_desc = TensorDesc({9, 1024}, ACL_INT4, ACL_FORMAT_ND).ValueRange(-1, 1);
1255+ TensorDesc x2_desc = TensorDesc({256, 1024}, ACL_INT4, ACL_FORMAT_ND).ValueRange(-1, 1);
1256+ TensorDesc x1scale_desc = TensorDesc({9,1}, ACL_FLOAT, ACL_FORMAT_ND);
1257+ TensorDesc x2scale_desc = TensorDesc({4, 256}, ACL_FLOAT, ACL_FORMAT_ND);
1258+ TensorDesc x2offset_desc = TensorDesc({4, 256}, ACL_FLOAT16, ACL_FORMAT_ND);
1259+ TensorDesc out_desc = TensorDesc({9, 256}, ACL_BF16, ACL_FORMAT_ND);
1260+ int64_t groupSize = 256L;
1261+ auto ut = OP_API_UT(aclnnQuantMatmulV5, INPUT(x1_desc, x2_desc, x1scale_desc, x2scale_desc, nullptr, nullptr, x2offset_desc, nullptr, nullptr, false, true, groupSize), OUTPUT(out_desc));
1262+ uint64_t workspace_size = 0;
1263+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
1264+ EXPECT_EQ(aclRet, ACLNN_SUCCESS);
1265+}
1266+ 
1267+TEST_F(l2_QuantBatchMatmulV5_test_950, ascend910D_A4W4_test_case_PERTOKEN_PERGROUP_1)
1268+{
1269+ // A4W4 pertoken-pergroup_scale out: float16
1270+ op::NpuArchManager archManager(NpuArch::DAV_3510);
1271+ TensorDesc x1_desc = TensorDesc({9, 1024}, ACL_INT4, ACL_FORMAT_ND).ValueRange(-1, 1);
1272+ TensorDesc x2_desc = TensorDesc({256, 1024}, ACL_INT4, ACL_FORMAT_ND).ValueRange(-1, 1);
1273+ TensorDesc x1scale_desc = TensorDesc({9,1}, ACL_FLOAT, ACL_FORMAT_ND);
1274+ TensorDesc x2scale_desc = TensorDesc({4, 256}, ACL_FLOAT, ACL_FORMAT_ND);
1275+ TensorDesc x2offset_desc = TensorDesc({4, 256}, ACL_FLOAT16, ACL_FORMAT_ND);
1276+ TensorDesc out_desc = TensorDesc({9, 256}, ACL_FLOAT16, ACL_FORMAT_ND);
1277+ int64_t groupSize = 256L;
1278+ auto ut = OP_API_UT(aclnnQuantMatmulV5, INPUT(x1_desc, x2_desc, x1scale_desc, x2scale_desc, nullptr, nullptr, x2offset_desc, nullptr, nullptr, false, true, groupSize), OUTPUT(out_desc));
1279+ uint64_t workspace_size = 0;
1280+ aclnnStatus aclRet = ut.TestGetWorkspaceSize(&workspace_size);
1281+ EXPECT_EQ(aclRet, ACLNN_SUCCESS);
1282+}
Amatmul/quant_batch_matmul_v4/tests/ut/op_host/test_quant_batch_matmul_v4_pergroup_basic_api_tiling.cpp+221-0
@@ -0,0 +1,221 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include <gtest/gtest.h>
12+ 
13+#include <cstdint>
14+#include <map>
15+#include <string>
16+#include <vector>
17+ 
18+#include "ut_op_util.h"
19+#include "exe_graph/runtime/storage_shape.h"
20+#include "kernel_run_context_facker.h"
21+#include "test_cube_util.h"
22+#include "platform/platform_infos_def.h"
23+#include "../../../op_host/op_tiling/quant_batch_matmul_v4_compile_info.h"
24+ 
25+using namespace ge;
26+using namespace ut_util;
27+ 
28+namespace {
29+struct PergroupBasicApiShapeGuardCase {
30+ std::string caseName;
31+ std::string socVersion;
32+ std::vector<int64_t> x1Dims;
33+ std::vector<int64_t> x2Dims;
34+ ge::graphStatus expectedResult;
35+ uint64_t expectedTilingKey;
36+ uint32_t expectedNumBlocks;
37+ bool checkTilingMeta;
38+};
39+ 
40+struct PergroupBasicApiCaseResult {
41+ ge::graphStatus result = ge::GRAPH_FAILED;
42+ uint64_t tilingKey = 0;
43+ uint64_t numBlocks = 0;
44+};
45+ 
46+class TestQuantBatchMatmulV4PergroupBasicApiTiling
47+ : public testing::TestWithParam<PergroupBasicApiShapeGuardCase> {};
48+ 
49+static gert::StorageShape MakeStorageShape(const std::vector<int64_t> &dims)
50+{
51+ gert::StorageShape shape;
52+ auto &storageShape = shape.MutableStorageShape();
53+ for (const auto dim : dims) {
54+ storageShape.AppendDim(dim);
55+ }
56+ shape.MutableOriginShape() = shape.MutableStorageShape();
57+ return shape;
58+}
59+ 
60+static PergroupBasicApiCaseResult RunPergroupBasicApiCase(const PergroupBasicApiShapeGuardCase &param)
61+{
62+ PergroupBasicApiCaseResult caseResult;
63+ const int64_t m = param.x1Dims[0];
64+ const int64_t k = param.x1Dims[1];
65+ const int64_t n = param.x2Dims[0];
66+ constexpr int64_t groupSize = 256;
67+ auto x1Shape = MakeStorageShape(param.x1Dims);
68+ auto x2Shape = MakeStorageShape(param.x2Dims);
69+ auto x1ScaleShape = MakeStorageShape({m, 1});
70+ auto x2ScaleShape = MakeStorageShape({k / groupSize, n});
71+ auto x2OffsetShape = MakeStorageShape({k / groupSize, n});
72+ gert::StorageShape outputShape({m, n}, {m, n});
73+ 
74+ std::string compileInfoStr;
75+ if (param.socVersion == "Ascend910B") {
76+ compileInfoStr = R"({
77+ "hardware_info": {"BT_SIZE": 1024, "load3d_constraints": "0",
78+ "Intrinsic_fix_pipe_l0c2out": true, "Intrinsic_data_move_l12ub": true,
79+ "Intrinsic_data_move_l0c2ub": true, "Intrinsic_data_move_out2l1_nd2nz": false,
80+ "UB_SIZE": 196352, "L2_SIZE": 201326592, "L1_SIZE": 524032,
81+ "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072, "CORE_NUM": 24,
82+ "cube_core_cnt": 24, "vector_core_cnt": 48, "core_type_list": "CubeCore,VectorCore"}
83+ })";
84+ } else {
85+ // Ascend950 platform setup. Keep Intrinsic_mmad as s8s4 to trigger supportMmadS8S4 branch.
86+ compileInfoStr = R"({
87+ "hardware_info" : {
88+ "BT_SIZE" : 4096,
89+ "load3d_constraints" : "unknown",
90+ "Intrinsic_fix_pipe_l0c2out" : true,
91+ "Intrinsic_data_move_l12ub" : false,
92+ "Intrinsic_data_move_l0c2ub" : false,
93+ "Intrinsic_data_move_l12bt" : true,
94+ "Intrinsic_data_move_out2l1_nd2nz" : true,
95+ "UB_SIZE" : 253952,
96+ "L2_SIZE" : 134217728,
97+ "L1_SIZE" : 524288,
98+ "L0A_SIZE" : 65536,
99+ "L0B_SIZE" : 65536,
100+ "L0C_SIZE" : 262144,
101+ "CORE_NUM" : 32,
102+ "cube_core_cnt": 32,
103+ "vector_core_cnt": 64,
104+ "core_type_list": "CubeCore,VectorCore",
105+ "socVersion" : "Ascend950",
106+ "NpuArch" : "3510"
107+ }})";
108+ }
109+ 
110+ std::map<std::string, std::string> socInfos;
111+ std::map<std::string, std::string> aicoreSpec;
112+ std::map<std::string, std::string> intrinsics;
113+ std::map<std::string, std::string> version;
114+ GetPlatFormInfos(compileInfoStr.c_str(), socInfos, aicoreSpec, intrinsics, version);
115+ socInfos["socVersion"] = param.socVersion;
116+ aicoreSpec["cube_freq"] = "1800";
117+ if (param.socVersion == "Ascend910B") {
118+ version["Short_SoC_version"] = "Ascend910B";
119+ }
120+ 
121+ fe::PlatFormInfos platformInfo;
122+ platformInfo.Init();
123+ optiling::QuantBatchMatmulV4CompileInfo compileInfo;
124+ 
125+ auto kernelHold = gert::KernelRunContextFaker()
126+ .KernelIONum(2, 1)
127+ .Inputs({const_cast<char *>(compileInfoStr.c_str()), reinterpret_cast<void *>(&platformInfo)})
128+ .Outputs({&compileInfo})
129+ .Build();
130+ 
131+ const std::string opType("QuantBatchMatmulV4");
132+ auto *opImpl = gert::OpImplRegistry::GetInstance().GetOpImpl(opType.c_str());
133+ if (opImpl == nullptr) {
134+ return caseResult;
135+ }
136+ 
137+ auto rawTilingData = gert::TilingData::CreateCap(4096);
138+ if (rawTilingData == nullptr) {
139+ return caseResult;
140+ }
141+ auto workspaceHolder = gert::ContinuousVector::Create<size_t>(4096);
142+ auto workspace = reinterpret_cast<gert::ContinuousVector *>(workspaceHolder.get());
143+ 
144+ auto holder = gert::TilingContextFaker()
145+ .NodeIoNum(10, 1)
146+ .IrInstanceNum({1, 1, 1, 1, 1, 1, 1, 1, 1, 1})
147+ .InputShapes({&x1Shape, &x2Shape, nullptr, &x1ScaleShape, &x2ScaleShape, nullptr,
148+ nullptr, &x2OffsetShape, nullptr, nullptr})
149+ .OutputShapes({&outputShape})
150+ .CompileInfo(&compileInfo)
151+ .PlatformInfo(reinterpret_cast<char *>(&platformInfo))
152+ .NodeInputTd(0, ge::DT_INT4, ge::FORMAT_ND, ge::FORMAT_ND)
153+ .NodeInputTd(1, ge::DT_INT4, ge::FORMAT_ND, ge::FORMAT_ND)
154+ .NodeInputTd(2, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
155+ .NodeInputTd(3, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
156+ .NodeInputTd(4, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
157+ .NodeInputTd(5, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
158+ .NodeInputTd(6, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
159+ .NodeInputTd(7, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
160+ .NodeInputTd(8, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
161+ .NodeInputTd(9, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
162+ .NodeOutputTd(0, ge::DT_BF16, ge::FORMAT_ND, ge::FORMAT_ND)
163+ .NodeAttrs({{"dtype", Ops::NN::AnyValue::CreateFrom<int64_t>(groupSize)},
164+ {"compute_type", Ops::NN::AnyValue::CreateFrom<bool>(false)},
165+ {"transpose_x1", Ops::NN::AnyValue::CreateFrom<bool>(false)},
166+ {"transpose_x2", Ops::NN::AnyValue::CreateFrom<bool>(true)},
167+ {"group_size", Ops::NN::AnyValue::CreateFrom<int64_t>(groupSize)}})
168+ .TilingData(rawTilingData.get())
169+ .Workspace(workspace)
170+ .SetOpType(opType)
171+ .Build();
172+ 
173+ auto *tilingContext = holder.GetContext<gert::TilingContext>();
174+ if (tilingContext == nullptr || tilingContext->GetPlatformInfo() == nullptr) {
175+ return caseResult;
176+ }
177+ tilingContext->GetPlatformInfo()->SetPlatformRes("SoCInfo", socInfos);
178+ tilingContext->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicoreSpec);
179+ tilingContext->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
180+ tilingContext->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics);
181+ tilingContext->GetPlatformInfo()->SetPlatformRes("version", version);
182+ 
183+ if (opImpl->tiling_parse == nullptr || opImpl->tiling == nullptr) {
184+ return caseResult;
185+ }
186+ if (opImpl->tiling_parse(kernelHold.GetContext<gert::KernelContext>()) != ge::GRAPH_SUCCESS) {
187+ return caseResult;
188+ }
189+ caseResult.result = opImpl->tiling(tilingContext);
190+ if (caseResult.result == ge::GRAPH_SUCCESS) {
191+ caseResult.tilingKey = tilingContext->GetTilingKey();
192+ caseResult.numBlocks = tilingContext->GetBlockDim();
193+ }
194+ return caseResult;
195+}
196+ 
197+TEST_P(TestQuantBatchMatmulV4PergroupBasicApiTiling, ShapeGuard)
198+{
199+ const auto &param = GetParam();
200+ const auto result = RunPergroupBasicApiCase(param);
201+ ASSERT_EQ(result.result, param.expectedResult) << "case=" << param.caseName;
202+ if (param.checkTilingMeta) {
203+ ASSERT_EQ(result.tilingKey, param.expectedTilingKey) << "case=" << param.caseName;
204+ ASSERT_EQ(result.numBlocks, param.expectedNumBlocks) << "case=" << param.caseName;
205+ }
206+}
207+ 
208+// Guard framework for Ascend950/Ascend910B:
209+// 1) Ascend950 keeps Intrinsic_mmad=s8s4.
210+// 2) Ascend910B maps Short_SoC_version=Ascend910B.
211+// 3) Verify by only changing x1/x2 dims below.
212+// 4) Add future guard cases in this table.
213+static PergroupBasicApiShapeGuardCase g_cases[] = {
214+ // Legal cases for pergroup basic api path:
215+ // m=128, k=1024 (k%1024==0), n=256 (n%256==0), all inputs are 2D.
216+ {"legal_rank2_should_succeed_950", "Ascend950", {1024, 1024}, {1024, 1024}, ge::GRAPH_SUCCESS, 537, 32, true},
217+ {"legal_rank2_should_succeed_910b", "Ascend910B", {128, 1024}, {256, 1024}, ge::GRAPH_SUCCESS, 1040, 1, true},
218+};
219+ 
220+INSTANTIATE_TEST_CASE_P(MM, TestQuantBatchMatmulV4PergroupBasicApiTiling, testing::ValuesIn(g_cases));
221+} // namespace