TT量化场景支持StreamK-ops-nn-AtomGit

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

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

Mmatmul/quant_batch_matmul_v3/op_host/op_tiling/arch35/base_block_calculator.cpp+9-1

	@@ -27,6 +27,8 @@ constexpr uint64_t LOAD_BALANCE_BASE_N_128_ALIGN_K_THRESHOLD = 2560UL;
27	constexpr uint64_t BASEM_BASEN_RATIO = 2UL;	27	constexpr uint64_t BASEM_BASEN_RATIO = 2UL;
		Y Y yangyang0168月3日 PR 描述目前仍是空模板，未关联 Issue，也没有填写测试范围、类型标签和性能数据
28	// Oversized baseK candidates are halved above this supported tiling range.	28	// Oversized baseK candidates are halved above this supported tiling range.
29	constexpr uint64_t BASEK_LIMIT = 4095UL;	29	constexpr uint64_t BASEK_LIMIT = 4095UL;
		30	+// Use one 256-byte K-split alignment for every StreamK layout and input type.
		31	+constexpr uint64_t STREAMK_INNER_K_GM_ALIGN_SIZE = 256UL;
30	constexpr uint32_t DOUBLE_CORE_NUM = 2U;	32	constexpr uint32_t DOUBLE_CORE_NUM = 2U;
31	// Epsilon for comparing score ratios during base-block search.	33	// Epsilon for comparing score ratios during base-block search.
32	constexpr double SCORE_COMPARE_EPS = 1e-12;	34	constexpr double SCORE_COMPARE_EPS = 1e-12;
	@@ -59,6 +61,11 @@ namespace optiling {
59		61
60	using Ops::NN::MathUtil;	62	using Ops::NN::MathUtil;
61		63
		64	+uint64_t GetStreamKSingleCoreKAlignSize(ge::DataType inputDtype)
		65	+{
		66	+ return GetShapeWithDataType(STREAMK_INNER_K_GM_ALIGN_SIZE, inputDtype);
		67	+}
		68	+
62	BaseBlockCalculator::BaseBlockCalculator(const QuantBatchMatmulInfo& inputParams,	69	BaseBlockCalculator::BaseBlockCalculator(const QuantBatchMatmulInfo& inputParams,
63	const QuantBatchMatmulV3CompileInfo& compileInfo, uint64_t batchCoreCnt)	70	const QuantBatchMatmulV3CompileInfo& compileInfo, uint64_t batchCoreCnt)
64	: inputParams_(inputParams), compileInfo_(compileInfo), batchCoreCnt_(batchCoreCnt)	71	: inputParams_(inputParams), compileInfo_(compileInfo), batchCoreCnt_(batchCoreCnt)
	@@ -249,7 +256,8 @@ void BaseBlockCalculator::UpdateTailStreamKBase()
249	bool BaseBlockCalculator::FinalizeStreamKBaseK()	256	bool BaseBlockCalculator::FinalizeStreamKBaseK()
250	{	257	{
251	uint64_t baseKAlignValue = GetBaseKAlignSize();	258	uint64_t baseKAlignValue = GetBaseKAlignSize();
252	- baseBlockRes_.singleCoreK = ops::CeilAlign(baseBlockRes_.singleCoreK, baseKAlignValue);	259	+ uint64_t singleCoreKAlignValue = GetStreamKSingleCoreKAlignSize(inputParams_.aDtype);
		260	+ baseBlockRes_.singleCoreK = ops::CeilAlign(baseBlockRes_.singleCoreK, singleCoreKAlignValue);
253	OP_TILING_CHECK(baseBlockRes_.singleCoreK == 0UL,	261	OP_TILING_CHECK(baseBlockRes_.singleCoreK == 0UL,
254	CUBE_INNER_ERR_REPORT(inputParams_.opName, "Invalid StreamK singleCoreK should be greater than 0."),	262	CUBE_INNER_ERR_REPORT(inputParams_.opName, "Invalid StreamK singleCoreK should be greater than 0."),
255	return false);	263	return false);

Mmatmul/quant_batch_matmul_v3/op_host/op_tiling/arch35/base_block_calculator.h+2-0

	@@ -34,6 +34,8 @@ struct BaseBlockRes {
34		34
35	enum class BaseBlockMode { DEFAULT = 0, PERBLOCK, MMAD_S8S4, STREAMK };	35	enum class BaseBlockMode { DEFAULT = 0, PERBLOCK, MMAD_S8S4, STREAMK };
36		36
		37	+uint64_t GetStreamKSingleCoreKAlignSize(ge::DataType inputDtype);
		38	+
37	class BaseBlockCalculator {	39	class BaseBlockCalculator {
38	public:	40	public:
39	BaseBlockCalculator(const QuantBatchMatmulInfo& inputParams, const QuantBatchMatmulV3CompileInfo& compileInfo,	41	BaseBlockCalculator(const QuantBatchMatmulInfo& inputParams, const QuantBatchMatmulV3CompileInfo& compileInfo,

Mmatmul/quant_batch_matmul_v3/op_host/op_tiling/arch35/qbmm_streamk_tiling.cpp+112-17

	@@ -255,22 +255,30 @@ bool AddMte2TensorBytes(uint64_t repeatCnt, uint64_t outerDim, uint64_t innerDim
255		255
256	uint64_t EstimateMte2Bytes(const optiling::QuantBatchMatmulInfo& inputParams, uint64_t mCnt, uint64_t nCnt)	256	uint64_t EstimateMte2Bytes(const optiling::QuantBatchMatmulInfo& inputParams, uint64_t mCnt, uint64_t nCnt)
257	{	257	{
258	- uint64_t scaleBytes = GetDtypeBytes(inputParams.scaleDtype);	258	+ if (inputParams.isMxPerGroup &&
259	- uint64_t pertokenScaleBytes = GetDtypeBytes(inputParams.perTokenScaleDtype);	259	+ (GetDtypeBytes(inputParams.scaleDtype) == 0UL \|\| GetDtypeBytes(inputParams.perTokenScaleDtype) == 0UL)) {
260	- if (scaleBytes == 0UL \|\| pertokenScaleBytes == 0UL) {
261	return 0UL;	260	return 0UL;
262	}	261	}
263		262
		263	+ uint64_t totalBytes = 0UL;
		264	+ bool ok = AddMte2TensorBytes(nCnt, inputParams.mSize, inputParams.kSize, inputParams.aDtype, totalBytes) &&
		265	+ AddMte2TensorBytes(mCnt, inputParams.nSize, inputParams.kSize, inputParams.bDtype, totalBytes);
		266	+ if (!ok) {
		267	+ return UINT64_SATURATED;
		268	+ }
		269	+
		270	+ if (!inputParams.isMxPerGroup) {
		271	+ // The one or two scalar scales used by non-MX per-tensor input are negligible in this traffic model.
		272	+ return totalBytes;
		273	+ }
		274	+
264	uint64_t scaleK = SaturatingMul(SafeCeilDiv(inputParams.kSize, MXFP_DIVISOR_SIZE), MXFP_MULTI_BASE_SIZE);	275	uint64_t scaleK = SaturatingMul(SafeCeilDiv(inputParams.kSize, MXFP_DIVISOR_SIZE), MXFP_MULTI_BASE_SIZE);
265	if (scaleK == UINT64_SATURATED) {	276	if (scaleK == UINT64_SATURATED) {
266	return UINT64_SATURATED;	277	return UINT64_SATURATED;
267	}	278	}
268	-	279	+ // MX scale tensors use [M, ceil(K / 64) * 2] and [ceil(K / 64) * 2, N] layouts.
269	- uint64_t totalBytes = 0UL;	280	+ ok = AddMte2TensorBytes(nCnt, inputParams.mSize, scaleK, inputParams.perTokenScaleDtype, totalBytes) &&
270	- bool ok = AddMte2TensorBytes(nCnt, inputParams.mSize, inputParams.kSize, inputParams.aDtype, totalBytes) &&	281	+ AddMte2TensorBytes(mCnt, inputParams.nSize, scaleK, inputParams.scaleDtype, totalBytes);
271	- AddMte2TensorBytes(mCnt, inputParams.nSize, inputParams.kSize, inputParams.bDtype, totalBytes) &&
272	- AddMte2TensorBytes(nCnt, inputParams.mSize, scaleK, inputParams.perTokenScaleDtype, totalBytes) &&
273	- AddMte2TensorBytes(mCnt, inputParams.nSize, scaleK, inputParams.scaleDtype, totalBytes);
274	return ok ? totalBytes : UINT64_SATURATED;	282	return ok ? totalBytes : UINT64_SATURATED;
275	}	283	}
276		284
	@@ -415,7 +423,7 @@ bool UpdateActualStreamKSchedule(const optiling::QuantBatchMatmulInfo& inputPara
415	shape.streamKCnt = SafeCeilDiv(inputParams.kSize, shape.singleCoreK);	423	shape.streamKCnt = SafeCeilDiv(inputParams.kSize, shape.singleCoreK);
416	}	424	}
417	}	425	}
418	- shape.singleCoreK = ops::CeilAlign(shape.singleCoreK, shape.baseKAlign);	426	+ shape.singleCoreK = ops::CeilAlign(shape.singleCoreK, optiling::GetStreamKSingleCoreKAlignSize(inputParams.aDtype));
419	return shape.singleCoreK != 0UL;	427	return shape.singleCoreK != 0UL;
420	}	428	}
421		429
	@@ -865,14 +873,91 @@ bool QBMMV3StreamKTiling::IsMxInput() const
865	return isMxfp8 \|\| isMxfp4;	873	return isMxfp8 \|\| isMxfp4;
866	}	874	}
867		875
		876	+bool QBMMV3StreamKTiling::IsPostDequantBiasInput() const
		877	+{
		878	+ const bool isInt8 = inputParams_.aDtype == ge::DT_INT8 && inputParams_.bDtype == ge::DT_INT8;
		879	+ const bool isInt8SingleScale = isInt8 && !inputParams_.isDoubleScale && inputParams_.hasBias &&
		880	+ (inputParams_.scaleDtype == ge::DT_FLOAT \|\|
		881	+ inputParams_.scaleDtype == ge::DT_BF16) &&
		882	+ inputParams_.biasDtype == inputParams_.scaleDtype;
		883	+ const bool isDoubleFp32Scale = inputParams_.isDoubleScale && inputParams_.scaleDtype == ge::DT_FLOAT &&
		884	+ inputParams_.perTokenScaleDtype == ge::DT_FLOAT && inputParams_.hasBias &&
		885	+ inputParams_.biasDtype == ge::DT_FLOAT;
		886	+ return isInt8SingleScale \|\| isDoubleFp32Scale;
		887	+}
		888	+
		889	+bool QBMMV3StreamKTiling::IsAllSkScheduleSupported(uint64_t mnCnt) const
		890	+{
		891	+ // DP cannot add a bias after dequantization, so post-dequant bias must use a uniform all-SK schedule. Without
		892	+ // post-dequant bias, both DP and SK combine the two per-tensor scales before applying the fixpipe mask.
		893	+ // The device scheduler uses usedCoreNum == aicNum and computes DP tiles as mnCnt - mnCnt % usedCoreNum;
		894	+ // requiring mnCnt < aicNum makes that value exactly zero.
		895	+ return !IsPostDequantBiasInput() \|\| (compileInfo_.aicNum != 0UL && mnCnt < compileInfo_.aicNum);
		896	+}
		897	+
		898	+bool QBMMV3StreamKTiling::IsPertensorStreamKInput() const
		ddssz8月3日 [性能][建议改进] 💡 为 per-tensor 重算收益模型这个新增入口让非 MX 输入继续复用 EvaluateStreamKBenefitGate，但 EstimateMte2Bytes 固定按 MX 的 ceil(K/64)*2 分组规模展开两路 scale 流量；per-tensor 实际只有一到两个标量，单 scale 时 perTokenScale 还不存在。该偏差会改变 savedMte2Bytes/permille 门限，可能误选或漏选 StreamK；建议按 quant mode 分支估算并用边界 shape/benchmark 校准。
		899	+{
		900	+ const bool isSupportedFormat = inputParams_.aFormat == ge::FORMAT_ND && inputParams_.cFormat == ge::FORMAT_ND &&
		901	+ (inputParams_.bFormat == ge::FORMAT_ND \|\|
		902	+ inputParams_.bFormat == ge::FORMAT_FRACTAL_NZ);
		903	+ // 本 StreamK 扩展只覆盖非 MX 的 per-tensor 量化：x2Scale 为 {1}；
		904	+ // x1Scale(接口名 pertokenScaleOptional)若存在，也必须是 {1}，对应 isDoubleScale。
		905	+ // shape 为 {M} 的真正 per-token 行广播场景暂不纳入该模板，避免误走 StreamK。
		906	+ const bool isSupportedQuantMode = inputParams_.isPerTensor && !inputParams_.isPertoken &&
		907	+ !inputParams_.isPerChannel && !inputParams_.isMxPerGroup &&
		908	+ !inputParams_.isPerBlock && !inputParams_.isPerBlockPerToken;
		909	+ if (!isSupportedFormat \|\| !isSupportedQuantMode) {
		910	+ return false;
		911	+ }
		912	+
		913	+ // 非 MX StreamK 的 SK block 把 raw accumulator 写入 workspace，最终在 AIV/UB 上做 scale/bias 反量化。
		914	+ // X2 scale 按 Fixpipe 乘法字段掩码后再乘；uint64/int64 输入从低 32bit 的 deq_scale 编码中取该字段。
		915	+ const bool isIntScale = inputParams_.scaleDtype == ge::DT_UINT64 \|\| inputParams_.scaleDtype == ge::DT_INT64;
		916	+ const bool isFloatScale = inputParams_.scaleDtype == ge::DT_FLOAT \|\| inputParams_.scaleDtype == ge::DT_BF16;
		917	+
		918	+ const bool isInt8 = inputParams_.aDtype == ge::DT_INT8 && inputParams_.bDtype == ge::DT_INT8;
		919	+ if (isInt8) {
		920	+ const bool isSupportedScaleAndOutput = !inputParams_.isDoubleScale &&
		921	+ ((isIntScale && (inputParams_.cDtype == ge::DT_FLOAT16 \|\|
		922	+ inputParams_.cDtype == ge::DT_BF16)) \|\|
		923	+ (isFloatScale && inputParams_.cDtype == ge::DT_BF16));
		924	+ // INT32 bias stays in the MMAD accumulation domain. Matching FP32/BF16 scale and bias are
		925	+ // applied by the AIV epilogue; IsCapable limits that combination to an all-SK schedule.
		926	+ const bool isSupportedBias = !inputParams_.hasBias \|\| inputParams_.biasDtype == ge::DT_INT32 \|\|
		927	+ IsPostDequantBiasInput();
		928	+ return isSupportedScaleAndOutput && isSupportedBias;
		929	+ }
		930	+
		931	+ const auto isFp8 = [](ge::DataType dtype) { return dtype == ge::DT_FLOAT8_E4M3FN \|\| dtype == ge::DT_FLOAT8_E5M2; };
		932	+ const bool isHif8Pair = inputParams_.aDtype == ge::DT_HIFLOAT8 && inputParams_.bDtype == ge::DT_HIFLOAT8;
		933	+ const bool isFp8Pair = isFp8(inputParams_.aDtype) && isFp8(inputParams_.bDtype);
		934	+ if (!isHif8Pair && !isFp8Pair) {
		935	+ return false;
		936	+ }
		937	+
		938	+ const bool isSupportedOutput = inputParams_.cDtype == ge::DT_FLOAT16 \|\| inputParams_.cDtype == ge::DT_BF16 \|\|
		939	+ inputParams_.cDtype == ge::DT_FLOAT;
		940	+ const bool isSupportedScale = (!inputParams_.isDoubleScale && isIntScale) \|\|
		941	+ (inputParams_.isDoubleScale && inputParams_.scaleDtype == ge::DT_FLOAT &&
		942	+ inputParams_.perTokenScaleDtype == ge::DT_FLOAT);
		943	+ // Encoded integer scale keeps FP32 bias in MMAD. Double-FP32 scale applies FP32 bias after both scale
		944	+ // multiplications in the dedicated AIV epilogue; IsCapable limits that case to an all-SK schedule.
		945	+ const bool isSupportedBias = !inputParams_.hasBias \|\| (isIntScale && inputParams_.biasDtype == ge::DT_FLOAT) \|\|
		946	+ IsPostDequantBiasInput();
		947	+ return isSupportedScale && isSupportedOutput && isSupportedBias;
		948	+}
		949	+
868	bool QBMMV3StreamKTiling::IsCapable()	950	bool QBMMV3StreamKTiling::IsCapable()
869	{	951	{
870	- if (!IsMxInput()) {	952	+ bool isMxInput = IsMxInput();
871	- OP_LOGD(inputParams_.opName, "QBMM StreamK only supports MX per-group input.");	953	+ bool isPertensorStreamKInput = IsPertensorStreamKInput();
		Z Z zhaoyingchao28月3日这里把非 MX per-tensor 输入送进了原先的 MX benefit gate，但 `EstimateMte2Bytes()` 仍按 MX per-group 布局计算：`scaleK = ceil(K/64) * 2`，并按 M/N tile 重复搬运 x1/x2 scale。TT 场景的 scale 实际各只有 1 个元素（单 scale 场景甚至没有 x1Scale），因此 `savedMte2Bytes/savedMte2Permille` 会被虚构的 scale 流量显著抬高，可能把实际不受益的 case 错误选进 StreamK，造成性能回退。请在收益模型中区分 `isMxPerGroup`，按标量 scale 的真实访问量计算 TT。
		954	+ if (!isMxInput && !isPertensorStreamKInput) {
		955	+ OP_LOGD(inputParams_.opName,
		956	+ "QBMM StreamK only supports MX per-group or non-MX per-tensor vector-dequant input.");
872	return false;	957	return false;
873	}	958	}
874	if (inputParams_.batchC != 1UL) {	959	if (inputParams_.batchC != 1UL) {
875	- OP_LOGD(inputParams_.opName, "QBMM StreamK only supports no-batch MX input, batchC=%lu.", inputParams_.batchC);	960	+ OP_LOGD(inputParams_.opName, "QBMM StreamK only supports no-batch input, batchC=%lu.", inputParams_.batchC);
876	return false;	961	return false;
877	}	962	}
878	if (compileInfo_.aivNum == 0UL) {	963	if (compileInfo_.aivNum == 0UL) {
	@@ -904,8 +989,14 @@ bool QBMMV3StreamKTiling::IsCapable()
904	LogBenefitGateEval(inputParams_.opName, inputParams_, benefitGate);	989	LogBenefitGateEval(inputParams_.opName, inputParams_, benefitGate);
905	if (!benefitGate.admit) {	990	if (!benefitGate.admit) {
906	OP_LOGD(inputParams_.opName, "QBMM StreamK capability gate result: reject reason=%s.", benefitGate.reason);	991	OP_LOGD(inputParams_.opName, "QBMM StreamK capability gate result: reject reason=%s.", benefitGate.reason);
		992	+ return false;
907	}	993	}
908	- return benefitGate.admit;	994	+ if (!IsAllSkScheduleSupported(benefitGate.skMnCnt)) {
		995	+ OP_LOGD(inputParams_.opName, "QBMM StreamK post-dequant bias requires all-SK, mnCnt=%lu aicNum=%u.",
		996	+ benefitGate.skMnCnt, compileInfo_.aicNum);
		997	+ return false;
		998	+ }
		999	+ return true;
909	}	1000	}
910		1001
911	bool QBMMV3StreamKTiling::CalcBaseBlock()	1002	bool QBMMV3StreamKTiling::CalcBaseBlock()
	@@ -998,8 +1089,6 @@ void QBMMV3StreamKTiling::SetTilingData()
998	QuantBatchMatMulV3TilingUtil::SetCommonTilingData(inputParams_, tilingData_);	1089	QuantBatchMatMulV3TilingUtil::SetCommonTilingData(inputParams_, tilingData_);
999	tilingData_.matmulTiling.weightMustHitL2 = static_cast<uint8_t>(	1090	tilingData_.matmulTiling.weightMustHitL2 = static_cast<uint8_t>(
1000	IsWeightMustHitL2(inputParams_, basicTiling_.baseM));	1091	IsWeightMustHitL2(inputParams_, basicTiling_.baseM));
1001	- tilingData_.params.x1QuantMode = static_cast<uint32_t>(BasicQuantMode::MX_PERGROUP_MODE);
1002	- tilingData_.params.x2QuantMode = static_cast<uint32_t>(BasicQuantMode::MX_PERGROUP_MODE);
1003	tilingData_.matmulTiling.m = static_cast<uint32_t>(inputParams_.mSize);	1092	tilingData_.matmulTiling.m = static_cast<uint32_t>(inputParams_.mSize);
1004	tilingData_.matmulTiling.n = static_cast<uint32_t>(inputParams_.nSize);	1093	tilingData_.matmulTiling.n = static_cast<uint32_t>(inputParams_.nSize);
1005	tilingData_.matmulTiling.k = static_cast<uint32_t>(inputParams_.kSize);	1094	tilingData_.matmulTiling.k = static_cast<uint32_t>(inputParams_.kSize);
	@@ -1015,7 +1104,13 @@ void QBMMV3StreamKTiling::SetTilingData()
1015	tilingData_.matmulTiling.scaleKL1 = static_cast<uint32_t>(scaleKL1_);	1104	tilingData_.matmulTiling.scaleKL1 = static_cast<uint32_t>(scaleKL1_);
1016	// Current StreamK BlockMmad uses fixed double L1 buffers. nBufferNum is retained for BasicAPI tiling/log	1105	// Current StreamK BlockMmad uses fixed double L1 buffers. nBufferNum is retained for BasicAPI tiling/log
1017	// compatibility and is not used by the StreamK kernel as a runtime tuning knob.	1106	// compatibility and is not used by the StreamK kernel as a runtime tuning knob.
1018	- CalculateNBufferNum4MX();	1107	+ if (IsMxInput()) {
		1108	+ tilingData_.params.x1QuantMode = static_cast<uint32_t>(BasicQuantMode::MX_PERGROUP_MODE);
		1109	+ tilingData_.params.x2QuantMode = static_cast<uint32_t>(BasicQuantMode::MX_PERGROUP_MODE);
		1110	+ CalculateNBufferNum4MX();
		1111	+ } else {
		1112	+ tilingData_.matmulTiling.nBufferNum = L1_TWO_BUFFER;
		1113	+ }
1019	// adaptiveSlidingWin is kept only to preserve the shared BasicAPI tiling data layout. StreamK scheduling uses	1114	// adaptiveSlidingWin is kept only to preserve the shared BasicAPI tiling data layout. StreamK scheduling uses
1020	// streamKTiling fields instead of ASW tail/window parameters, so fill neutral placeholders here.	1115	// streamKTiling fields instead of ASW tail/window parameters, so fill neutral placeholders here.
1021	tilingData_.adaptiveSlidingWin.mTailTile = 1U;	1116	tilingData_.adaptiveSlidingWin.mTailTile = 1U;

Mmatmul/quant_batch_matmul_v3/op_host/op_tiling/arch35/qbmm_streamk_tiling.h+3-0

	@@ -44,6 +44,9 @@ protected:
44	private:	44	private:
45	void Reset();	45	void Reset();
46	bool IsMxInput() const;	46	bool IsMxInput() const;
		47	+ bool IsPertensorStreamKInput() const;
		48	+ bool IsPostDequantBiasInput() const;
		49	+ bool IsAllSkScheduleSupported(uint64_t mnCnt) const;
47	bool CalcBaseBlock();	50	bool CalcBaseBlock();
48	bool CalcL1Tiling();	51	bool CalcL1Tiling();
49	void SetTilingData();	52	void SetTilingData();

Amatmul/quant_batch_matmul_v3/op_kernel/arch35/qbmm_pertensor_streamk_tensor_api_blaze.h+75-0

@@ -0,0 +1,75 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file qbmm_pertensor_streamk_tensor_api_blaze.h
	13	+ * \brief Quantized batch matrix multiplication using StreamK and AIV vector dequant.
	14	+ */
	15	+#pragma once
	16	+
	17	+#include "blaze/gemm/kernel/kernel_qbmm_pertensor_streamk.h"
	18	+#include "blaze/epilogue/block/block_epilogue_qbmm_pertensor_streamk.h"
	19	+#include "blaze/gemm/block/block_scheduler_matmul_streamk.h"
	20	+#include "blaze/gemm/block/block_mmad_a8w8_fixpipe_quant.h"
	21	+
	22	+template <class A_TYPE, class B_TYPE, class SCALE_TYPE, class C_TYPE, class BIAS_TYPE, class aLayout, class bLayout,
	23	+ class cLayout, uint64_t FULL_LOAD_MODE = 0>
	24	+__aicore__ inline void QbmmPertensorStreamKTensorApiKernel(GM_ADDR aGM, GM_ADDR bGM, GM_ADDR scale, GM_ADDR bias,
	ddssz8月3日 [测试][必须修复] 🔴 补充新 Kernel 的精度与调度覆盖本 PR 新增的三条用例都是 op_host tiling 用例，且 tilingData=0 会在校验 key/blockDim 后跳过 tiling 内容；没有新增 op_kernel 或 ST 用例，因此这个新 Kernel 未被结果精度覆盖。请至少覆盖 ND/Weight-NZ、转置、single/double scale、MMAD/AIV bias，以及 mnCnt>=core 的 DP+SK 和全 SK 两类调度。
	25	+ GM_ADDR perTokenScale, GM_ADDR cGM, GM_ADDR workspaceGm,
	26	+ const void* tilingData)
	27	+{
	28	+ using AType = A_TYPE;
	29	+ using BType = B_TYPE;
	30	+ using X2ScaleType = SCALE_TYPE;
	31	+ using BiasType = BIAS_TYPE;
	32	+ using OutType = C_TYPE;
	33	+
	34	+ using ProblemShape = AscendC::Te::Shape<int64_t, int64_t, int64_t, int64_t>;
	35	+
	36	+ using DispatchPolicy = Blaze::Gemm::MatmulWithScaleFixpipeQuant<FULL_LOAD_MODE, false,
	37	+ Blaze::Gemm::KernelQbmmPertensorMultiBlockStreamK>;
	38	+ using EpilogueDispatchPolicy = DispatchPolicy;
	39	+ using BlockScheduler = Blaze::Gemm::Block::BlockSchedulerMatmulStreamK<ProblemShape>;
	40	+ using BlockMmad = Blaze::Gemm::Block::BlockMmad<DispatchPolicy, AType, aLayout,
	41	+ AscendC::Std::tuple<BType, X2ScaleType>, bLayout, OutType, cLayout,
	42	+ BiasType, cLayout>;
	43	+ using WorkspaceType = typename BlockMmad::WorkspaceType;
	44	+ using BlockEpilogue = Blaze::Epilogue::Block::BlockEpilogueQbmmPertensorStreamK<
	45	+ WorkspaceType, OutType, EpilogueDispatchPolicy, X2ScaleType, float>;
	46	+ using MatmulKernel = Blaze::Gemm::Kernel::GemmUniversal<ProblemShape, BlockMmad, BlockEpilogue, BlockScheduler>;
	47	+ using Params = typename MatmulKernel::Params;
	48	+
	49	+ const DequantBmm::QuantBatchMatmulV3StreamKBasicAPITilingData*
	50	+ quantBmmTilingData = static_cast<const DequantBmm::QuantBatchMatmulV3StreamKBasicAPITilingData*>(tilingData);
	51	+ DequantBmm::BasicAPICubeTiling matmulTiling = quantBmmTilingData->matmulTiling;
	52	+ bool hasBias = matmulTiling.isBias != 0U;
	53	+ constexpr bool isIntScale = AscendC::IsSameType<X2ScaleType, uint64_t>::value \|\|
	54	+ AscendC::IsSameType<X2ScaleType, int64_t>::value;
	55	+ constexpr bool isFp8Input = !AscendC::IsSameType<AType, int8_t>::value;
	56	+ bool biasInMmad = hasBias && ((quantBmmTilingData->params.biasDtype == DT_INT32 &&
	57	+ AscendC::IsSameType<BiasType, int32_t>::value) \|\|
	58	+ (isFp8Input && isIntScale && quantBmmTilingData->params.biasDtype == DT_FLOAT &&
	59	+ AscendC::IsSameType<BiasType, float>::value));
	60	+ GM_ADDR biasMmadGm = biasInMmad ? bias : nullptr;
	61	+ GM_ADDR biasEpilogueGm = hasBias && !biasInMmad ? bias : nullptr;
	62	+ bool isBiasEpilogue = biasEpilogueGm != nullptr;
	63	+
	64	+ uint64_t kernelKL1 = quantBmmTilingData->streamKTiling.kL1;
	65	+
	66	+ Params params{
	67	+ {matmulTiling.m, matmulTiling.n, matmulTiling.k, quantBmmTilingData->params.batchC},
	68	+ {aGM, bGM, cGM, biasMmadGm, perTokenScale, scale},
	69	+ {cGM, workspaceGm, scale, perTokenScale, biasEpilogueGm, isBiasEpilogue, quantBmmTilingData->params.biasDtype},
	70	+ {AscendC::GetBlockNum(), matmulTiling.baseM, matmulTiling.baseN, matmulTiling.baseK,
	71	+ quantBmmTilingData->streamKTiling.singleCoreK, static_cast<int64_t>(kernelKL1)}};
	72	+
	73	+ MatmulKernel qbmm;
	74	+ qbmm(params);
	75	+}

Mmatmul/quant_batch_matmul_v3/op_kernel/arch35/quant_batch_matmul_v3.cpp+39-0

	@@ -51,6 +51,7 @@
51	#include "qbmm_cube_tensor_api_blaze.h"	51	#include "qbmm_cube_tensor_api_blaze.h"
52	#include "qbmm_mix_tensor_api_blaze.h"	52	#include "qbmm_mix_tensor_api_blaze.h"
53	#include "qbmm_mix_without_batch_tensor_api_blaze.h"	53	#include "qbmm_mix_without_batch_tensor_api_blaze.h"
		54	+#include "qbmm_pertensor_streamk_tensor_api_blaze.h"
54	#if (ORIG_DTYPE_SCALE == DT_FLOAT8_E8M0)	55	#if (ORIG_DTYPE_SCALE == DT_FLOAT8_E8M0)
55	#include "qbmm_mx_tensor_api_blaze.h"	56	#include "qbmm_mx_tensor_api_blaze.h"
56	#include "qbmm_mx_without_batch_tensor_api_blaze.h"	57	#include "qbmm_mx_without_batch_tensor_api_blaze.h"
	@@ -240,6 +241,27 @@ constexpr CubeFormat format_y = CubeFormat::ND;
240	QbmmMixWithoutBatchTensorApiKernel<DTYPE_X1, DTYPE_X2, DTYPE_SCALE, DTYPE_Y, DTYPE_BIAS, aLayout, bLayout, \	241	QbmmMixWithoutBatchTensorApiKernel<DTYPE_X1, DTYPE_X2, DTYPE_SCALE, DTYPE_Y, DTYPE_BIAS, aLayout, bLayout, \
241	cLayout, fullLoadMode>(x1, x2, scale, bias, pertokenScale, y, &tilingData); \	242	cLayout, fullLoadMode>(x1, x2, scale, bias, pertokenScale, y, &tilingData); \
242	} while (0)	243	} while (0)
		244	+
		245	+// Non-MX per-tensor StreamK template dtype combinations are selected once by
		246	+// SUPPORT_NON_MX_STREAMK_TILING_KEY in quant_batch_matmul_v3_apt_tiling_key.h (bias is optional):
		247	+// 1. x1/x2: int8, scale: uint64/int64, perTokenScale: null, bias: int32, y: fp16/bf16;
		248	+// 2. x1/x2: int8, scale: fp32, perTokenScale: null, bias: int32/fp32, y: bf16;
		249	+// 3. x1/x2: int8, scale: bf16, perTokenScale: null, bias: int32/bf16, y: bf16;
		250	+// 4. x1/x2: both FP8 (e4m3fn/e5m2 may be mixed) or both hifloat8, scale: uint64/int64,
		251	+// perTokenScale: null, bias: fp32, y: fp16/bf16/fp32;
		252	+// 5. x1/x2: both FP8 (e4m3fn/e5m2 may be mixed) or both hifloat8, scale/perTokenScale: fp32/fp32,
		253	+// bias: fp32, y: fp16/bf16/fp32. Hifloat8 and FP8 matrix inputs cannot be mixed.
		254	+// MMAD writes unscaled partials to workspace; AIV reduces them, applies scale and optional post-dequant bias,
		255	+// then casts and writes the final C tile. INT32 bias and FP32 bias paired with encoded scale are accumulated by
		256	+// MMAD; matching floating bias is applied by the AIV epilogue. Host tiling checks the bias/scale pairing,
		257	+// per-tensor mode, single batch, and all-SK post-dequant-bias schedule.
		258	+#define QUANT_BMMV3_PERTENSOR_STREAMK_BLAZE_IMPL_CLASS(aLayout, bLayout, cLayout, fullLoadMode) \
		259	+ do { \
		260	+ GET_TILING_DATA_WITH_STRUCT(DequantBmm::QuantBatchMatmulV3StreamKBasicAPITilingData, tilingData, tiling); \
		261	+ QbmmPertensorStreamKTensorApiKernel<DTYPE_X1, DTYPE_X2, DTYPE_SCALE, DTYPE_Y, DTYPE_BIAS, aLayout, bLayout, \
		262	+ cLayout, fullLoadMode>(x1, x2, scale, bias, pertokenScale, y, user1, \
		263	+ &tilingData); \
		264	+ } while (0)
243	#endif	265	#endif
244		266
245	// ASCEND_IS_NOT_AIV 等价于 (分离架构ASCEND_IS_AIC OR 耦合架构)	267	// ASCEND_IS_NOT_AIV 等价于 (分离架构ASCEND_IS_AIC OR 耦合架构)
	@@ -712,6 +734,23 @@ UT_STATIC __global__ __aicore__ void quant_batch_matmul_v3(GM_ADDR x1, GM_ADDR x
712	#endif	734	#endif
713	#endif	735	#endif
714	} else {	736	} else {
		737	+#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 3510) && IS_BLAZE && SUPPORT_NON_MX_STREAMK_TILING_KEY
		738	+ // Non-MX per-tensor StreamK uses the no-batch tiling key; keep it outside the legacy with-batch dispatch.
		739	+ if constexpr (TPL_BATCHMODE == TPL_WITHOUT_BATCH && TPL_KERNELTYPE == TPL_VEC_EPILOGUE_STREAMK_WITH_MMAPI &&
		740	+ TPL_APILEVEL == TPL_API_LEVEL_BLAZE) {
		741	+ using StreamKALayout = typename AscendC::Conditional<
		742	+ static_cast<bool>(TPL_ATRANS), AscendC::Te::DNExtLayoutPtn, AscendC::Te::NDExtLayoutPtn>::type;
		743	+#if CUBE_TEMPLATE_ND
		744	+ using StreamKBLayout = typename AscendC::Conditional<
		745	+ static_cast<bool>(TPL_BTRANS), AscendC::Te::DNExtLayoutPtn, AscendC::Te::NDExtLayoutPtn>::type;
		746	+#elif defined(FORMAT_X2) && FORMAT_X2 == FORMAT_FRACTAL_NZ
		747	+ using StreamKBLayout = typename AscendC::Conditional<
		748	+ static_cast<bool>(TPL_BTRANS), AscendC::Te::ZNLayoutPtn, AscendC::Te::NZLayoutPtn>::type;
		749	+#endif
		750	+ QUANT_BMMV3_PERTENSOR_STREAMK_BLAZE_IMPL_CLASS(StreamKALayout, StreamKBLayout, AscendC::Te::NDExtLayoutPtn,
		751	+ 0);
		752	+ }
		753	+#endif
715	if constexpr (TPL_BATCHMODE == TPL_WITH_BATCH) { // Batch Mode = WITH_BATCH	754	if constexpr (TPL_BATCHMODE == TPL_WITH_BATCH) { // Batch Mode = WITH_BATCH
716	#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 3510)	755	#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 3510)
717	#if CUBE_TEMPLATE_ND	756	#if CUBE_TEMPLATE_ND

Mmatmul/quant_batch_matmul_v3/op_kernel/arch35/quant_batch_matmul_v3_apt_tiling_key.h+41-3

	@@ -77,6 +77,43 @@ namespace QuantBatchMatmulV3Arch35TilingKey {
77	#define SUPPORT_MX_WITHOUT_BATCH_TILING_KEY false	77	#define SUPPORT_MX_WITHOUT_BATCH_TILING_KEY false
78	#endif	78	#endif
79		79
		80	+// This is the compile-time upper bound for non-MX per-tensor StreamK. Host tiling further validates optional
		81	+// scale/bias inputs, batch, and the all-SK requirement. Keep the dtype-family guards independent because some
		82	+// INT8 compile environments do not define FP8 dtype macros.
		83	+#if defined(ORIG_DTYPE_X1) && defined(ORIG_DTYPE_X2) && defined(ORIG_DTYPE_SCALE) && defined(ORIG_DTYPE_Y) && \
		84	+ defined(DT_INT8) && defined(DT_UINT64) && defined(DT_INT64) && defined(DT_FLOAT) && defined(DT_BF16) && \
		85	+ defined(DT_FLOAT16)
		86	+#define QBMMV3_IS_INT8_PERTENSOR_STREAMK_TPL \
		87	+ ((ORIG_DTYPE_X1 == DT_INT8) && (ORIG_DTYPE_X2 == DT_INT8) && \
		88	+ ((((ORIG_DTYPE_SCALE == DT_UINT64) \|\| (ORIG_DTYPE_SCALE == DT_INT64)) && \
		89	+ ((ORIG_DTYPE_Y == DT_FLOAT16) \|\| (ORIG_DTYPE_Y == DT_BF16))) \|\| \
		90	+ (((ORIG_DTYPE_SCALE == DT_FLOAT) \|\| (ORIG_DTYPE_SCALE == DT_BF16)) && ORIG_DTYPE_Y == DT_BF16)))
		Z Z zhaoyingchao28月3日是否只需要感知输出和scale的数据类型
		91	+#else
		92	+#define QBMMV3_IS_INT8_PERTENSOR_STREAMK_TPL false
		93	+#endif
		94	+
		95	+#if defined(ORIG_DTYPE_X1) && defined(ORIG_DTYPE_X2) && defined(ORIG_DTYPE_SCALE) && defined(ORIG_DTYPE_Y) && \
		96	+ defined(DT_UINT64) && defined(DT_INT64) && defined(DT_FLOAT) && defined(DT_BF16) && defined(DT_FLOAT16) && \
		97	+ defined(DT_HIFLOAT8) && defined(DT_FLOAT8_E4M3FN) && defined(DT_FLOAT8_E5M2)
		98	+#define QBMMV3_IS_FP8_PERTENSOR_STREAMK_TPL \
		99	+ ((((ORIG_DTYPE_X1 == DT_HIFLOAT8) && (ORIG_DTYPE_X2 == DT_HIFLOAT8)) \|\| \
		100	+ (((ORIG_DTYPE_X1 == DT_FLOAT8_E4M3FN) \|\| (ORIG_DTYPE_X1 == DT_FLOAT8_E5M2)) && \
		101	+ ((ORIG_DTYPE_X2 == DT_FLOAT8_E4M3FN) \|\| (ORIG_DTYPE_X2 == DT_FLOAT8_E5M2)))) && \
		102	+ ((ORIG_DTYPE_SCALE == DT_UINT64) \|\| (ORIG_DTYPE_SCALE == DT_INT64) \|\| (ORIG_DTYPE_SCALE == DT_FLOAT)) && \
		103	+ ((ORIG_DTYPE_Y == DT_FLOAT16) \|\| (ORIG_DTYPE_Y == DT_BF16) \|\| (ORIG_DTYPE_Y == DT_FLOAT)))
		104	+#else
		105	+#define QBMMV3_IS_FP8_PERTENSOR_STREAMK_TPL false
		106	+#endif
		107	+
		108	+#if defined(FORMAT_X2) && defined(FORMAT_ND) && defined(FORMAT_FRACTAL_NZ)
		109	+#define SUPPORT_NON_MX_STREAMK_TILING_KEY \
		110	+ (!QBMMV3_IS_MX_DTYPE_TPL && (QBMMV3_IS_INT8_PERTENSOR_STREAMK_TPL \|\| QBMMV3_IS_FP8_PERTENSOR_STREAMK_TPL) && \
		111	+ (FORMAT_X2 == FORMAT_ND \|\| FORMAT_X2 == FORMAT_FRACTAL_NZ))
		112	+#else
		113	+#define SUPPORT_NON_MX_STREAMK_TILING_KEY false
		114	+#endif
		115	+
		116	+// Keep the existing non-MX Cube ND/API-level classification independent from StreamK key selection.
80	#if defined(__CCE_AICORE__) && defined(ORIG_DTYPE_SCALE) && defined(FORMAT_X2) && defined(FORMAT_ND) && \	117	#if defined(__CCE_AICORE__) && defined(ORIG_DTYPE_SCALE) && defined(FORMAT_X2) && defined(FORMAT_ND) && \
81	defined(DT_UINT64) && defined(DT_INT64) && defined(DT_FLOAT) && defined(DT_BF16)	118	defined(DT_UINT64) && defined(DT_INT64) && defined(DT_FLOAT) && defined(DT_BF16)
82	#define QBMMV3_IS_NON_MX_CUBE_ND_TPL \	119	#define QBMMV3_IS_NON_MX_CUBE_ND_TPL \
	@@ -121,7 +158,8 @@ namespace QuantBatchMatmulV3Arch35TilingKey {
121	#define QBMMV3_IS_NON_MX_WEIGHT_NZ_TPL false	158	#define QBMMV3_IS_NON_MX_WEIGHT_NZ_TPL false
122	#endif	159	#endif
123		160
124	-#define SUPPORT_MX_STREAMK_TILING_KEY SUPPORT_MX_WITHOUT_BATCH_TILING_KEY	161	+// Kernel type 11 is shared by the original MX StreamK path and the non-MX per-tensor StreamK path above.
		162	+#define SUPPORT_STREAMK_TILING_KEY (SUPPORT_MX_WITHOUT_BATCH_TILING_KEY \|\| SUPPORT_NON_MX_STREAMK_TILING_KEY)
125		163
126	// Batch Mode	164	// Batch Mode
127	#define TPL_WITH_BATCH 0	165	#define TPL_WITH_BATCH 0
	@@ -198,8 +236,8 @@ ASCENDC_TPL_SEL(
198	TPL_NO_VEC_EPILOGUE_CUSTOM_GMTOAL1_WITH_MMAPI),	236	TPL_NO_VEC_EPILOGUE_CUSTOM_GMTOAL1_WITH_MMAPI),
199	ASCENDC_TPL_UINT_SEL(APILEVEL, ASCENDC_TPL_UI_LIST, TPL_API_LEVEL_BLAZE)),	237	ASCENDC_TPL_UINT_SEL(APILEVEL, ASCENDC_TPL_UI_LIST, TPL_API_LEVEL_BLAZE)),
200	#endif	238	#endif
201	-#if ((!defined(__CCE_AICORE__)) \|\| (SUPPORT_MX_STREAMK_TILING_KEY))	239	+#if ((!defined(__CCE_AICORE__)) \|\| (SUPPORT_STREAMK_TILING_KEY))
202	- ASCENDC_TPL_ARGS_SEL( // kernel type {11} * ATRANS {0, 1} * BTRANS {0, 1}	240	+ ASCENDC_TPL_ARGS_SEL( // kernel type {11}: AIC split-K workspace + AIV scale/bias epilogue
203	ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2), ASCENDC_TPL_UINT_SEL(ATRANS, ASCENDC_TPL_UI_LIST, 0, 1),	241	ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIC_1_2), ASCENDC_TPL_UINT_SEL(ATRANS, ASCENDC_TPL_UI_LIST, 0, 1),
204	ASCENDC_TPL_UINT_SEL(BTRANS, ASCENDC_TPL_UI_LIST, 0, 1),	242	ASCENDC_TPL_UINT_SEL(BTRANS, ASCENDC_TPL_UI_LIST, 0, 1),
205	ASCENDC_TPL_UINT_SEL(BATCHMODE, ASCENDC_TPL_UI_LIST, TPL_WITHOUT_BATCH),	243	ASCENDC_TPL_UINT_SEL(BATCHMODE, ASCENDC_TPL_UI_LIST, TPL_WITHOUT_BATCH),

Mmatmul/quant_batch_matmul_v3/tests/st/arch35/ttk_kernel_quant_batch_matmul_v3_st.csv+4-1

	@@ -24,4 +24,7 @@ qbmm_KT_case0011,quant_batch_matmul_v3,"((1, 1), (1, 1), (1,), None, (1,), (1,))
24	qbmm_KC_case0003,quant_batch_matmul_v3,"((679, 149), (238, 149), (238,), None, None, (679,))","((679, 149), (238, 149), (238,), None, None, (679,))","((679, 238),)","((679, 238),)","('int8', 'int8', 'bfloat16', 'int8', 'int8', 'float32')",'bfloat16',"('ND', 'ND', 'ND', 'ND', 'ND', 'ND')","('ND', 'ND', 'ND', 'ND', 'ND', 'ND')","('ND',)","('ND',)","{'dtype': 1, 'transpose_x1': False, 'transpose_x2': True}","((-10, 10),)","((0.001, 0.001),)",,1,random,random,random	24	qbmm_KC_case0003,quant_batch_matmul_v3,"((679, 149), (238, 149), (238,), None, None, (679,))","((679, 149), (238, 149), (238,), None, None, (679,))","((679, 238),)","((679, 238),)","('int8', 'int8', 'bfloat16', 'int8', 'int8', 'float32')",'bfloat16',"('ND', 'ND', 'ND', 'ND', 'ND', 'ND')","('ND', 'ND', 'ND', 'ND', 'ND', 'ND')","('ND',)","('ND',)","{'dtype': 1, 'transpose_x1': False, 'transpose_x2': True}","((-10, 10),)","((0.001, 0.001),)",,1,random,random,random
25	qbmm_KC_case0014,quant_batch_matmul_v3,"((1232, 1348), (1104, 1348), (1104,), None, (1104,), (1232,))","((1232, 1348), (1104, 1348), (1104,), None, (1104,), (1232,))","((1232, 1104),)","((1232, 1104),)","('int8', 'int8', 'float32', 'int8', 'float32', 'float32')",'bfloat16',"('ND', 'ND', 'ND', 'ND', 'ND', 'ND')","('ND', 'ND', 'ND', 'ND', 'ND', 'ND')","('ND',)","('ND',)","{'dtype': 1, 'transpose_x1': False, 'transpose_x2': True}","((-10, 10),)","((0.001, 0.001),)",,1,multiple_of_16,multiple_of_16,random	25	qbmm_KC_case0014,quant_batch_matmul_v3,"((1232, 1348), (1104, 1348), (1104,), None, (1104,), (1232,))","((1232, 1348), (1104, 1348), (1104,), None, (1104,), (1232,))","((1232, 1104),)","((1232, 1104),)","('int8', 'int8', 'float32', 'int8', 'float32', 'float32')",'bfloat16',"('ND', 'ND', 'ND', 'ND', 'ND', 'ND')","('ND', 'ND', 'ND', 'ND', 'ND', 'ND')","('ND',)","('ND',)","{'dtype': 1, 'transpose_x1': False, 'transpose_x2': True}","((-10, 10),)","((0.001, 0.001),)",,1,multiple_of_16,multiple_of_16,random
26	qbmm_TC_case0027,quant_batch_matmul_v3,"((240, 176), (176, 304), (304,), None, (304,), None)","((240, 176), (176, 304), (304,), None, (304,), None)","((240, 304),)","((240, 304),)","('int8', 'int8', 'uint64', 'int8', 'int32', 'int8')",'int8',"('ND', 'ND', 'ND', 'ND', 'ND', 'ND')","('ND', 'ND', 'ND', 'ND', 'ND', 'ND')","('ND',)","('ND',)","{'dtype': 1, 'transpose_x1': False, 'transpose_x2': False}","((-10, 10),)","((0.001, 0.001),)",,1,multiple_of_16,multiple_of_16,multiple_of_16	26	qbmm_TC_case0027,quant_batch_matmul_v3,"((240, 176), (176, 304), (304,), None, (304,), None)","((240, 176), (176, 304), (304,), None, (304,), None)","((240, 304),)","((240, 304),)","('int8', 'int8', 'uint64', 'int8', 'int32', 'int8')",'int8',"('ND', 'ND', 'ND', 'ND', 'ND', 'ND')","('ND', 'ND', 'ND', 'ND', 'ND', 'ND')","('ND',)","('ND',)","{'dtype': 1, 'transpose_x1': False, 'transpose_x2': False}","((-10, 10),)","((0.001, 0.001),)",,1,multiple_of_16,multiple_of_16,multiple_of_16
27	-qbmm_TC_case0020,quant_batch_matmul_v3,"((443, 128), (443, 320), (320,), None, None, None)","((443, 128), (443, 320), (320,), None, None, None)","((128, 320),)","((128, 320),)","('float8_e4m3fn', 'float8_e5m2', 'uint64', 'int8', 'int8', 'int8')",'bfloat16',"('ND', 'ND', 'ND', 'ND', 'ND', 'ND')","('ND', 'ND', 'ND', 'ND', 'ND', 'ND')","('ND',)","('ND',)","{'dtype': 1, 'transpose_x1': True, 'transpose_x2': False}","((-10, 10),)","((0.001, 0.001),)",,1,multiple_of_32,multiple_of_32,prime	27	+qbmm_TC_case0020,quant_batch_matmul_v3,"((443, 128), (443, 320), (320,), None, None, None)","((443, 128), (443, 320), (320,), None, None, None)","((128, 320),)","((128, 320),)","('float8_e4m3fn', 'float8_e5m2', 'uint64', 'int8', 'int8', 'int8')",'bfloat16',"('ND', 'ND', 'ND', 'ND', 'ND', 'ND')","('ND', 'ND', 'ND', 'ND', 'ND', 'ND')","('ND',)","('ND',)","{'dtype': 1, 'transpose_x1': True, 'transpose_x2': False}","((-10, 10),)","((0.001, 0.001),)",,1,multiple_of_32,multiple_of_32,prime
		28	+qbmm_streamk_tt_int8_u64_int32_bias_nd,quant_batch_matmul_v3,"((44, 8192), (8192, 144), (1,), None, (144,), None)","((44, 8192), (8192, 144), (1,), None, (144,), None)","((44, 144),)","((44, 144),)","('int8', 'int8', 'uint64', 'int8', 'int32', 'float32')",'bfloat16',"('ND', 'ND', 'ND', 'ND', 'ND', 'ND')","('ND', 'ND', 'ND', 'ND', 'ND', 'ND')","('ND',)","('ND',)","{'dtype': 1, 'transpose_x1': False, 'transpose_x2': False}","((-3, 3),)","((0.001, 0.001),)",,1,multiple_of_16,multiple_of_16,multiple_of_16
		29	+qbmm_streamk_tt_int8_fp32_fp32_bias_nd,quant_batch_matmul_v3,"((44, 8192), (8192, 144), (1,), None, (144,), None)","((44, 8192), (8192, 144), (1,), None, (144,), None)","((44, 144),)","((44, 144),)","('int8', 'int8', 'float32', 'int8', 'float32', 'float32')",'bfloat16',"('ND', 'ND', 'ND', 'ND', 'ND', 'ND')","('ND', 'ND', 'ND', 'ND', 'ND', 'ND')","('ND',)","('ND',)","{'dtype': 1, 'transpose_x1': False, 'transpose_x2': False}","((-3, 3),)","((0.001, 0.001),)",,1,multiple_of_16,multiple_of_16,multiple_of_16
		30	+qbmm_streamk_tt_fp8_double_fp32_bias_nd,quant_batch_matmul_v3,"((44, 8192), (8192, 144), (1,), None, (144,), (1,))","((44, 8192), (8192, 144), (1,), None, (144,), (1,))","((44, 144),)","((44, 144),)","('float8_e4m3fn', 'float8_e5m2', 'float32', 'int8', 'float32', 'float32')",'float16',"('ND', 'ND', 'ND', 'ND', 'ND', 'ND')","('ND', 'ND', 'ND', 'ND', 'ND', 'ND')","('ND',)","('ND',)","{'dtype': 1, 'transpose_x1': False, 'transpose_x2': False}","((-3, 3),)","((0.001, 0.001),)",,1,multiple_of_16,multiple_of_16,multiple_of_16

Mmatmul/quant_batch_matmul_v3/tests/ut/op_host/test_quant_batch_matmul_v3.csv+10-0

	@@ -245,6 +245,16 @@ Ascend950,QUANTMM950/TestQuantBatchMatmulV3Tiling.generalTest/402,,mxfp4-streamk
245	Ascend950,QUANTMM950/TestQuantBatchMatmulV3Tiling.generalTest/403,,mxfp4-streamk-ta0tb1-key-only,,,2,2,2,0,0,0,4096,16384,256,0,1,0,0,1,2,FLOAT4-E2M1,FLOAT4-E2M1,FLOAT8-E8M0,FLOAT8-E8M0,FLOAT,FLOAT16,ND,ND,TRUE,32,133908,,FALSE,	245	Ascend950,QUANTMM950/TestQuantBatchMatmulV3Tiling.generalTest/403,,mxfp4-streamk-ta0tb1-key-only,,,2,2,2,0,0,0,4096,16384,256,0,1,0,0,1,2,FLOAT4-E2M1,FLOAT4-E2M1,FLOAT8-E8M0,FLOAT8-E8M0,FLOAT,FLOAT16,ND,ND,TRUE,32,133908,,FALSE,
246	Ascend950,QUANTMM950/TestQuantBatchMatmulV3Tiling.generalTest/404,,mxfp4-streamk-ta0tb1-small-gate-stub,,,2,2,2,0,0,0,4,8192,128,0,1,0,0,1,2,FLOAT4-E2M1,FLOAT4-E2M1,FLOAT8-E8M0,FLOAT8-E8M0,FLOAT,FLOAT16,ND,ND,TRUE,32,0,,TRUE,	246	Ascend950,QUANTMM950/TestQuantBatchMatmulV3Tiling.generalTest/404,,mxfp4-streamk-ta0tb1-small-gate-stub,,,2,2,2,0,0,0,4,8192,128,0,1,0,0,1,2,FLOAT4-E2M1,FLOAT4-E2M1,FLOAT8-E8M0,FLOAT8-E8M0,FLOAT,FLOAT16,ND,ND,TRUE,32,0,,TRUE,
247	Ascend950,QUANTMM950/TestQuantBatchMatmulV3Tiling.generalTest/405,,mxfp8-streamk-deterministic-level-fallback-aswt,,,2,2,2,0,0,0,256,8192,256,0,1,0,0,0,2,FLOAT8-E4M3,FLOAT8-E5M2,FLOAT8-E8M0,FLOAT8-E8M0,FLOAT,FLOAT16,ND,ND,TRUE,32,131088,,FALSE,,2	247	Ascend950,QUANTMM950/TestQuantBatchMatmulV3Tiling.generalTest/405,,mxfp8-streamk-deterministic-level-fallback-aswt,,,2,2,2,0,0,0,256,8192,256,0,1,0,0,0,2,FLOAT8-E4M3,FLOAT8-E5M2,FLOAT8-E8M0,FLOAT8-E8M0,FLOAT,FLOAT16,ND,ND,TRUE,32,131088,,FALSE,,2
		248	+Ascend950,QUANTMM950/TestQuantBatchMatmulV3Tiling.generalTest/406,,fp8-doublescale-streamk-postbias,,,2,2,2,0,0,0,256,8192,256,0,1,1,0,0,6,FLOAT8-E4M3,FLOAT8-E4M3,FLOAT,FLOAT,FLOAT,FLOAT16,ND,ND,TRUE,32,133904,0,FALSE,double FP32 scale with post-dequant FP32 bias
		249	+Ascend950,QUANTMM950/TestQuantBatchMatmulV3Tiling.generalTest/407,,int8-fp32-scale-bias-allsk,,,2,2,2,0,0,0,256,8192,256,0,0,1,0,0,0,INT8,INT8,FLOAT,FLOAT,FLOAT,BF16,ND,ND,TRUE,32,133904,0,FALSE,INT8 single FP32 scale with matching post-dequant bias; all-SK only
		250	+Ascend950,QUANTMM950/TestQuantBatchMatmulV3Tiling.generalTest/408,,int8-bf16-scale-bias-allsk,,,2,2,2,0,0,0,256,8192,256,0,0,1,0,0,0,INT8,INT8,BF16,FLOAT,BF16,BF16,ND,ND,TRUE,32,133904,0,FALSE,INT8 single BF16 scale with matching post-dequant bias; all-SK only
		251	+Ascend950,QUANTMM950/TestQuantBatchMatmulV3Tiling.generalTest/409,,fp8-doublescale-streamk-nobias-ta0tb0,,,2,2,2,0,0,0,256,8192,256,0,1,0,0,0,6,FLOAT8-E4M3,FLOAT8-E5M2,FLOAT,FLOAT,FLOAT,FLOAT16,ND,ND,TRUE,32,133904,0,FALSE,double FP32 scale without bias; DP and SK share merged masked scale
		252	+Ascend950,QUANTMM950/TestQuantBatchMatmulV3Tiling.generalTest/410,,fp8-doublescale-streamk-nobias-ta1tb0,,,2,2,2,0,0,0,256,8192,256,0,1,0,1,0,6,FLOAT8-E4M3,FLOAT8-E5M2,FLOAT,FLOAT,FLOAT,FLOAT16,ND,ND,TRUE,32,133905,0,FALSE,double FP32 scale transpose x1; DP and SK share merged masked scale
		253	+Ascend950,QUANTMM950/TestQuantBatchMatmulV3Tiling.generalTest/411,,fp8-doublescale-streamk-nobias-ta0tb1,,,2,2,2,0,0,0,4096,16384,256,0,1,0,0,1,6,FLOAT8-E4M3,FLOAT8-E5M2,FLOAT,FLOAT,FLOAT,FLOAT16,ND,ND,TRUE,32,133908,0,FALSE,double FP32 scale transpose x2; DP and SK share merged masked scale
		254	+Ascend950,QUANTMM950/TestQuantBatchMatmulV3Tiling.generalTest/412,,fp8-doublescale-streamk-nobias-ta1tb1,,,2,2,2,0,0,0,256,8192,256,0,1,0,1,1,6,FLOAT8-E4M3,FLOAT8-E5M2,FLOAT,FLOAT,FLOAT,FLOAT16,ND,ND,TRUE,32,133909,0,FALSE,double FP32 scale transpose x1/x2; DP and SK share merged masked scale
		255	+Ascend950,QUANTMM950/TestQuantBatchMatmulV3Tiling.generalTest/413,,fp8-doublescale-weightnz-unsupported,,,2,2,2,0,0,0,256,8192,256,0,1,0,0,0,6,FLOAT8-E4M3,FLOAT8-E4M3,FLOAT,FLOAT,FLOAT,FLOAT16,ND,NZ,FALSE,32,0,,FALSE,Weight-NZ FP8 double per-tensor scale remains unsupported
		256	+Ascend950,QUANTMM950/TestQuantBatchMatmulV3Tiling.generalTest/414,,hif8-doublescale-weightnz-unsupported,,,2,2,2,0,0,0,256,8192,256,0,1,0,0,0,6,HIFLOAT8,HIFLOAT8,FLOAT,FLOAT,FLOAT,FLOAT,ND,NZ,FALSE,32,0,,FALSE,Weight-NZ HIFLOAT8 double per-tensor scale remains unsupported
		257	+Ascend950,QUANTMM950/TestQuantBatchMatmulV3Tiling.generalTest/415,,int8-encoded-scale-int32-bias-streamk,,,2,2,2,0,0,0,256,8192,256,0,0,1,0,0,0,INT8,INT8,UINT64,FLOAT,INT32,BF16,ND,ND,TRUE,32,133904,0,FALSE,INT32 bias is accumulated by MMAD
248	Ascend950,QUANTMM950/TestQuantBatchMatmulV3Tiling.generalTest/43,,hi8-llama2-0010,,,3,3,3,1,1,1,4,8192,1024,0,0,1,0,1,0,HIFLOAT8,HIFLOAT8,UINT64,FLOAT,FLOAT,FLOAT16,ND,ND,TRUE,32,131076,1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 0 1 0 0 0 0 0 4 1024 8192 16 32 1024 6144 3072 0 131330 1 1 1 1 0 0,FALSE,	258	Ascend950,QUANTMM950/TestQuantBatchMatmulV3Tiling.generalTest/43,,hi8-llama2-0010,,,3,3,3,1,1,1,4,8192,1024,0,0,1,0,1,0,HIFLOAT8,HIFLOAT8,UINT64,FLOAT,FLOAT,FLOAT16,ND,ND,TRUE,32,131076,1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 0 1 0 0 0 0 0 4 1024 8192 16 32 1024 6144 3072 0 131330 1 1 1 1 0 0,FALSE,
249	Ascend950,QUANTMM950/TestQuantBatchMatmulV3Tiling.generalTest/44,,mxfp4-llama2-0001,,,3,3,3,1,1,1,4096,8192,128,0,1,1,0,1,2,FLOAT4-E2M1,FLOAT4-E2M1,FLOAT8-E8M0,FLOAT8-E8M0,FLOAT,FLOAT16,ND,ND,TRUE,32,131092,4096 128 8192 1024 1024 8192 8388736 66048 2097153 65537 65537 0 67592 66050 0 0,FALSE,	259	Ascend950,QUANTMM950/TestQuantBatchMatmulV3Tiling.generalTest/44,,mxfp4-llama2-0001,,,3,3,3,1,1,1,4096,8192,128,0,1,1,0,1,2,FLOAT4-E2M1,FLOAT4-E2M1,FLOAT8-E8M0,FLOAT8-E8M0,FLOAT,FLOAT16,ND,ND,TRUE,32,131092,4096 128 8192 1024 1024 8192 8388736 66048 2097153 65537 65537 0 67592 66050 0 0,FALSE,
250	Ascend950,QUANTMM950/TestQuantBatchMatmulV3Tiling.generalTest/45,,Key-llama2-case29,,,3,3,3,1,1,1,256,6656,4992,0,1,0,0,1,1,INT8,INT8,FLOAT,FLOAT,INT32,FLOAT16,ND,ND,TRUE,32,516,1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 0 0 1 0 0 127 160 0 0 0 3 0 0 0 0 0 0 0 0 32 256 4992 6656 6656 256 160 6656 256 160 128 6 12 1 1 0 0 0 0 0 0 0 0 0 0 0 3 6 0 0 2 2 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 0 0,FALSE,	260	Ascend950,QUANTMM950/TestQuantBatchMatmulV3Tiling.generalTest/45,,Key-llama2-case29,,,3,3,3,1,1,1,256,6656,4992,0,1,0,0,1,1,INT8,INT8,FLOAT,FLOAT,INT32,FLOAT16,ND,ND,TRUE,32,516,1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 0 0 1 0 0 127 160 0 0 0 3 0 0 0 0 0 0 0 0 32 256 4992 6656 6656 256 160 6656 256 160 128 6 12 1 1 0 0 0 0 0 0 0 0 0 0 0 3 6 0 0 2 2 1 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1 1 1 1 0 0,FALSE,

Mmatmul/quant_batch_matmul_v3/tests/ut/op_host/test_quant_batch_matmul_v3_tiling.cpp+191-0

	@@ -36,6 +36,8 @@
36	#include "../../../op_host/op_tiling/quant_batch_matmul_v3_basic_tiling.h"	36	#include "../../../op_host/op_tiling/quant_batch_matmul_v3_basic_tiling.h"
37	#include "../../../op_host/op_tiling/quant_batch_matmul_v3_tiling.h"	37	#include "../../../op_host/op_tiling/quant_batch_matmul_v3_tiling.h"
38	#include "../../../op_host/op_tiling/arch35/quant_batch_matmul_v3_tiling_util.h"	38	#include "../../../op_host/op_tiling/arch35/quant_batch_matmul_v3_tiling_util.h"
		39	+#include "../../../op_host/op_tiling/arch35/base_block_calculator.h"
		40	+#include "../../../op_host/op_tiling/arch35/qbmm_streamk_tiling.h"
39	#include "../../../op_kernel/arch35/quant_batch_matmul_v3_tiling_data.h"	41	#include "../../../op_kernel/arch35/quant_batch_matmul_v3_tiling_data.h"
40	#include "platform/platform_infos_def.h"	42	#include "platform/platform_infos_def.h"
41	#include "ut_string_utils.h"	43	#include "ut_string_utils.h"
	@@ -478,6 +480,9 @@ void QuantBatchMatmulV3TilingTestParam::Prepare(QuantBatchMatmulV3CompileInfo& c
478	} else if (quantMode == 5) { // dynamic T-C: x1Scale is per-tensor, x2Scale is per-channel.	480	} else if (quantMode == 5) { // dynamic T-C: x1Scale is per-tensor, x2Scale is per-channel.
479	pertokenShape.MutableStorageShape() = gert::Shape({1});	481	pertokenShape.MutableStorageShape() = gert::Shape({1});
480	scaleShape.MutableStorageShape() = gert::Shape({n});	482	scaleShape.MutableStorageShape() = gert::Shape({n});
		483	+ } else if (quantMode == 6) { // double per-tensor scale: x1Scale {1}, x2Scale {1}.
		484	+ pertokenShape.MutableStorageShape() = gert::Shape({1});
		485	+ scaleShape.MutableStorageShape() = gert::Shape({1});
481	}	486	}
482		487
483	biasShape.MutableStorageShape() = gert::Shape({n});	488	biasShape.MutableStorageShape() = gert::Shape({n});
	@@ -684,6 +689,9 @@ void QuantBatchMatmulV3TilingTestParam::InvokeTilingFunc(QuantBatchMatmulV3Compi
684	} else if (quantMode == 5) { // dynamic T-C: x1Scale is per-tensor, x2Scale is per-channel.	689	} else if (quantMode == 5) { // dynamic T-C: x1Scale is per-tensor, x2Scale is per-channel.
685	pertokenShape.MutableStorageShape() = gert::Shape({1});	690	pertokenShape.MutableStorageShape() = gert::Shape({1});
686	scaleShape.MutableStorageShape() = gert::Shape({n});	691	scaleShape.MutableStorageShape() = gert::Shape({n});
		692	+ } else if (quantMode == 6) { // double per-tensor scale: x1Scale {1}, x2Scale {1}.
		693	+ pertokenShape.MutableStorageShape() = gert::Shape({1});
		694	+ scaleShape.MutableStorageShape() = gert::Shape({1});
687	}	695	}
688		696
689	biasShape.MutableStorageShape() = gert::Shape({n});	697	biasShape.MutableStorageShape() = gert::Shape({n});
	@@ -883,6 +891,189 @@ TEST(QuantBatchMatmulV3TilingCsv, ShouldLoadValidCases)
883	}	891	}
884	}	892	}
885		893
		894	+static BaseBlockRes ComputeStreamKBaseBlock(bool isMxPerGroup, bool transA, bool transB, ge::DataType aDtype,
		895	+ ge::DataType bDtype, uint64_t mSize = 256UL, uint64_t nSize = 256UL,
		896	+ uint64_t kSize = 1000UL)
		897	+{
		898	+ QuantBatchMatmulInfo inputParams{};
		899	+ inputParams.opName = "QuantBatchMatmulV3StreamKSingleCoreKAlignUt";
		900	+ inputParams.mSize = mSize;
		901	+ inputParams.nSize = nSize;
		902	+ inputParams.kSize = kSize;
		903	+ inputParams.batchC = 1UL;
		904	+ inputParams.transA = transA;
		905	+ inputParams.transB = transB;
		906	+ inputParams.aDtype = aDtype;
		907	+ inputParams.bDtype = bDtype;
		908	+ inputParams.isMxPerGroup = isMxPerGroup;
		909	+ inputParams.isPerTensor = !isMxPerGroup;
		910	+
		911	+ QuantBatchMatmulV3CompileInfo compileInfo{};
		912	+ compileInfo.aicNum = 24U;
		913	+ compileInfo.l0aSize = 65536UL;
		914	+ compileInfo.l0bSize = 65536UL;
		915	+ compileInfo.npuArch = NpuArch::DAV_3510;
		916	+
		917	+ BaseBlockCalculator calculator(inputParams, compileInfo);
		918	+ EXPECT_TRUE(calculator.Compute(BaseBlockMode::STREAMK));
		919	+ return calculator.GetOutput();
		920	+}
		921	+
		922	+TEST(QuantBatchMatmulV3StreamKSingleCoreKAlign, CubeStreamKAlignsEveryTransposeTo256Bytes)
		923	+{
		924	+ for (bool transA : {false, true}) {
		925	+ for (bool transB : {false, true}) {
		926	+ const auto result = ComputeStreamKBaseBlock(false, transA, transB, ge::DT_INT8, ge::DT_INT8);
		927	+ EXPECT_EQ(result.singleCoreK, 256UL) << "transA=" << transA << ", transB=" << transB;
		928	+ EXPECT_EQ(GetSizeWithDataType(result.singleCoreK, ge::DT_INT8) % 256UL, 0UL);
		929	+ }
		930	+ }
		931	+}
		932	+
		933	+TEST(QuantBatchMatmulV3StreamKSingleCoreKAlign, MxStreamKAlignsEveryTransposeTo256Bytes)
		934	+{
		935	+ for (bool transA : {false, true}) {
		936	+ for (bool transB : {false, true}) {
		937	+ const auto result = ComputeStreamKBaseBlock(true, transA, transB, ge::DT_FLOAT4_E2M1, ge::DT_FLOAT4_E2M1);
		938	+ EXPECT_EQ(result.singleCoreK, 512UL) << "transA=" << transA << ", transB=" << transB;
		939	+ EXPECT_EQ(GetSizeWithDataType(result.singleCoreK, ge::DT_FLOAT4_E2M1) % 256UL, 0UL);
		940	+ }
		941	+ }
		942	+}
		943	+
		944	+TEST(QuantBatchMatmulV3StreamKSingleCoreKAlign, CubeStreamKKeepsByteAlignmentWhenBaseKIsNotFactor)
		945	+{
		946	+ const auto result = ComputeStreamKBaseBlock(false, false, false, ge::DT_INT8, ge::DT_INT8, 270UL, 16UL, 8192UL);
		947	+
		948	+ ASSERT_EQ(result.baseK, 224UL);
		949	+ EXPECT_EQ(result.singleCoreK, 768UL);
		950	+ EXPECT_EQ(GetSizeWithDataType(result.singleCoreK, ge::DT_INT8) % 256UL, 0UL);
		951	+ EXPECT_NE(result.singleCoreK % result.baseK, 0UL);
		952	+}
		953	+
		954	+TEST(QuantBatchMatmulV3StreamKAllSk, DoubleFp32ScaleRequiresAllSkOnlyWithPostBias)
		955	+{
		956	+ QBMMV3StreamKTiling tiling(nullptr);
		957	+ auto& input = tiling.inputParams_;
		958	+ input.aFormat = ge::FORMAT_ND;
		959	+ input.bFormat = ge::FORMAT_ND;
		960	+ input.cFormat = ge::FORMAT_ND;
		961	+ input.aDtype = ge::DT_FLOAT8_E4M3FN;
		962	+ input.bDtype = ge::DT_FLOAT8_E4M3FN;
		963	+ input.cDtype = ge::DT_FLOAT16;
		964	+ input.scaleDtype = ge::DT_FLOAT;
		965	+ input.perTokenScaleDtype = ge::DT_FLOAT;
		966	+ input.biasDtype = ge::DT_FLOAT;
		967	+ input.isPerTensor = true;
		968	+ input.isDoubleScale = true;
		969	+ input.hasBias = true;
		970	+
		971	+ tiling.compileInfo_.aicNum = 32U;
		972	+
		973	+ EXPECT_TRUE(tiling.IsPostDequantBiasInput());
		974	+ EXPECT_TRUE(tiling.IsPertensorStreamKInput());
		975	+ EXPECT_TRUE(tiling.IsAllSkScheduleSupported(1UL));
		976	+ EXPECT_FALSE(tiling.IsAllSkScheduleSupported(32UL));
		977	+ EXPECT_FALSE(tiling.IsAllSkScheduleSupported(33UL));
		978	+
		979	+ input.hasBias = false;
		980	+ EXPECT_FALSE(tiling.IsPostDequantBiasInput());
		981	+ EXPECT_TRUE(tiling.IsPertensorStreamKInput());
		982	+ EXPECT_TRUE(tiling.IsAllSkScheduleSupported(31UL));
		983	+ EXPECT_TRUE(tiling.IsAllSkScheduleSupported(32UL));
		984	+ EXPECT_TRUE(tiling.IsAllSkScheduleSupported(33UL));
		985	+}
		986	+
		987	+TEST(QuantBatchMatmulV3StreamKPostDequantBias, SupportsInt8MatchingFloatingBiasOnlyForAllSk)
		988	+{
		989	+ QBMMV3StreamKTiling tiling(nullptr);
		990	+ auto& input = tiling.inputParams_;
		991	+ input.aFormat = ge::FORMAT_ND;
		992	+ input.bFormat = ge::FORMAT_ND;
		993	+ input.cFormat = ge::FORMAT_ND;
		994	+ input.aDtype = ge::DT_INT8;
		995	+ input.bDtype = ge::DT_INT8;
		996	+ input.cDtype = ge::DT_BF16;
		997	+ input.scaleDtype = ge::DT_FLOAT;
		998	+ input.biasDtype = ge::DT_FLOAT;
		999	+ input.isPerTensor = true;
		1000	+ input.isDoubleScale = false;
		1001	+ input.hasBias = true;
		1002	+
		1003	+ tiling.compileInfo_.aicNum = 32U;
		1004	+
		1005	+ EXPECT_TRUE(tiling.IsPostDequantBiasInput());
		1006	+ EXPECT_TRUE(tiling.IsPertensorStreamKInput());
		1007	+ EXPECT_TRUE(tiling.IsAllSkScheduleSupported(31UL));
		1008	+ EXPECT_FALSE(tiling.IsAllSkScheduleSupported(32UL));
		1009	+
		1010	+ input.scaleDtype = ge::DT_BF16;
		1011	+ input.biasDtype = ge::DT_BF16;
		1012	+ EXPECT_TRUE(tiling.IsPostDequantBiasInput());
		1013	+ EXPECT_TRUE(tiling.IsPertensorStreamKInput());
		1014	+
		1015	+ input.biasDtype = ge::DT_FLOAT;
		1016	+ EXPECT_FALSE(tiling.IsPostDequantBiasInput());
		1017	+ EXPECT_FALSE(tiling.IsPertensorStreamKInput());
		1018	+
		1019	+ input.scaleDtype = ge::DT_FLOAT;
		1020	+ input.biasDtype = ge::DT_BF16;
		1021	+ EXPECT_FALSE(tiling.IsPostDequantBiasInput());
		1022	+ EXPECT_FALSE(tiling.IsPertensorStreamKInput());
		1023	+}
		1024	+
		1025	+TEST(QuantBatchMatmulV3StreamKDtype, RejectsHifloat8AndFp8MixedPair)
		1026	+{
		1027	+ QBMMV3StreamKTiling tiling(nullptr);
		1028	+ auto& input = tiling.inputParams_;
		1029	+ input.aFormat = ge::FORMAT_ND;
		1030	+ input.bFormat = ge::FORMAT_ND;
		1031	+ input.cFormat = ge::FORMAT_ND;
		1032	+ input.cDtype = ge::DT_FLOAT;
		1033	+ input.scaleDtype = ge::DT_FLOAT;
		1034	+ input.perTokenScaleDtype = ge::DT_FLOAT;
		1035	+ input.isPerTensor = true;
		1036	+ input.isDoubleScale = true;
		1037	+ input.hasBias = false;
		1038	+
		1039	+ input.aDtype = ge::DT_FLOAT8_E4M3FN;
		1040	+ input.bDtype = ge::DT_FLOAT8_E5M2;
		1041	+ EXPECT_TRUE(tiling.IsPertensorStreamKInput());
		1042	+
		1043	+ input.aDtype = ge::DT_HIFLOAT8;
		1044	+ input.bDtype = ge::DT_HIFLOAT8;
		1045	+ EXPECT_TRUE(tiling.IsPertensorStreamKInput());
		1046	+
		1047	+ input.bDtype = ge::DT_FLOAT8_E4M3FN;
		1048	+ EXPECT_FALSE(tiling.IsPertensorStreamKInput());
		1049	+}
		1050	+
		1051	+TEST(QuantBatchMatmulV3StreamKCapability, RejectsBatchBeforeBenefitEvaluation)
		1052	+{
		1053	+ QBMMV3StreamKTiling tiling(nullptr);
		1054	+ auto& input = tiling.inputParams_;
		1055	+ input.aFormat = ge::FORMAT_ND;
		1056	+ input.bFormat = ge::FORMAT_ND;
		1057	+ input.cFormat = ge::FORMAT_ND;
		1058	+ input.aDtype = ge::DT_INT8;
		1059	+ input.bDtype = ge::DT_INT8;
		1060	+ input.cDtype = ge::DT_BF16;
		1061	+ input.scaleDtype = ge::DT_FLOAT;
		1062	+ input.biasDtype = ge::DT_FLOAT;
		1063	+ input.isPerTensor = true;
		1064	+ input.isDoubleScale = false;
		1065	+ input.isPertoken = false;
		1066	+ input.isPerChannel = false;
		1067	+ input.isMxPerGroup = false;
		1068	+ input.isPerBlock = false;
		1069	+ input.isPerBlockPerToken = false;
		1070	+ input.hasBias = true;
		1071	+ input.batchC = 2UL;
		1072	+
		1073	+ EXPECT_TRUE(tiling.IsPertensorStreamKInput());
		1074	+ EXPECT_FALSE(tiling.IsCapable());
		1075	+}
		1076	+
886	TEST_P(TestQuantBatchMatmulV3Tiling, generalTest) { GetParam().Test(); }	1077	TEST_P(TestQuantBatchMatmulV3Tiling, generalTest) { GetParam().Test(); }
887		1078
888	static const std::vector<QuantBatchMatmulV3TilingTestParam> kCasesParams910B2 = GetParams("Ascend910B2");	1079	static const std::vector<QuantBatchMatmulV3TilingTestParam> kCasesParams910B2 = GetParams("Ascend910B2");