A2A3 TStore Conv3d-pto-isa-AtomGit

已合并

A2A3 TStore Conv3d #912

zc1110创建于 5月13日从zhengchao0613/pto-isa： conv3d_tstore 合入到 cann/pto-isa： master

A2A3 TStore Conv3d #912

已合并

zc1110创建于 5月13日

从zhengchao0613/pto-isa conv3d_tstore合入到cann/pto-isa master

共 4 个文件变更+732-9

Minclude/pto/npu/a2a3/TStore.hpp+61-4

	@@ -394,6 +394,55 @@ PTO_INTERNAL void TStoreAccNz2nz(typename GlobalData::DType *dstAddr, __cc__ typ
394	copy_matrix_cc_to_gm(dstAddr, srcAddr, xmReg, xtReg);	394	copy_matrix_cc_to_gm(dstAddr, srcAddr, xmReg, xtReg);
395	}	395	}
396		396
		397	+template <typename GlobalData, typename TileData, QuantMode_t quantizationMode = QuantMode_t::NoQuant,
		398	+ ReluPreMode reluPreMode = ReluPreMode::NoRelu, STPhase Phase = STPhase::Unspecified>
		399	+PTO_INTERNAL void TStoreAccNz2NDC1HWC0(typename GlobalData::DType dstAddr, __cc__ typename TileData::DType srcAddr,
		400	+ int gShape0, int gShape1, int gShape2, int gShape3, int gShape4, int gStride2,
		401	+ int gStride4, int validRow, int validCol)
		402	+{
		403	+ constexpr uint32_t c0ElemCount = std::is_same_v<typename GlobalData::DType, __gm__ int32_t> ?
		404	+ 16 :
		405	+ C0_SIZE_BYTE / sizeof(typename GlobalData::DType);
		406	+ PTO_ASSERT(validRow == gShape0 * gShape3 * gShape4,
		407	+ "The validRow of TileData must be equal to Shape0 * Shape3 * Shape4 of NDC1HWC0 shape!");
		408	+ PTO_ASSERT(validCol == gShape1 * gShape2 * c0ElemCount,
		409	+ "The validCol of TileData must be equal to Shape1 * Shape2 * c0ElemCount of NDC1HWC0 shape!");
		410	+ PTO_ASSERT(validRow >= 1 && validRow <= 65535,
		411	+ "When GlobalData is NDC1HWC0 format, the range of validRow is [1, 65535].");
		412	+
		413	+ static_assert(std::is_same_v<typename GlobalData::DType, __gm__ float> \|\|
		414	+ std::is_same_v<typename GlobalData::DType, __gm__ int32_t> \|\|
		415	+ std::is_same_v<typename GlobalData::DType, __gm__ half> \|\|
		416	+ std::is_same_v<typename GlobalData::DType, __gm__ bfloat16_t> \|\|
		417	+ std::is_same_v<typename GlobalData::DType, __gm__ int8_t> \|\|
		418	+ std::is_same_v<typename GlobalData::DType, __gm__ uint8_t>,
		419	+ "GlobalData::DType must be float/int32_t/half/bfloat16_t/int8_t/uint8_t.");
		420	+
		421	+ uint8_t channelSplitEn = 0;
		422	+ if (std::is_same_v<typename TileData::DType, float> && std::is_same_v<typename GlobalData::DType, __gm__ float>) {
		423	+ channelSplitEn = 1;
		424	+ }
		425	+
		426	+ uint16_t mSize = validRow;
		427	+ uint16_t nSize = validCol;
		428	+ uint16_t srcStride = TileData::Rows;
		429	+ if constexpr (CompactMode::Normal == TileData::Compact) {
		430	+ srcStride = CeilAlignment(validRow, FRACTAL_NZ_ROW);
		431	+ }
		432	+
		433	+ uint32_t c0Size = sizeof(typename GlobalData::DType) * c0ElemCount;
		434	+ uint32_t dstStride = gShape0 * gStride2 / gStride4 * c0Size >> SHIFT_BLOCK_BYTE;
		435	+ constexpr uint8_t unitFlagCtrl = static_cast<uint8_t>(Phase);
		436	+ uint64_t xtReg = srcStride \| (static_cast<uint64_t>(unitFlagCtrl & 0x3) << 32) \|
		437	+ (static_cast<uint64_t>(quantizationMode & 0x1f) << 34) \|
		438	+ ((static_cast<uint64_t>(reluPreMode) & 0x7) << 39) \|
		439	+ (static_cast<uint64_t>(channelSplitEn & 0x1) << 42);
		440	+ uint64_t xmReg = (static_cast<uint64_t>(nSize & 0xfff) << 4) \| (static_cast<uint64_t>(mSize & 0xffff) << 16) \|
		441	+ (static_cast<uint64_t>(dstStride & 0xffffffff) << 32);
		442	+
		443	+ copy_matrix_cc_to_gm(dstAddr, srcAddr, xmReg, xtReg);
		444	+}
		445	+
397	template <typename GlobalData, typename TileData, QuantMode_t quantizationMode = QuantMode_t::NoQuant,	446	template <typename GlobalData, typename TileData, QuantMode_t quantizationMode = QuantMode_t::NoQuant,
398	ReluPreMode reluPreMode = ReluPreMode::NoRelu, STPhase Phase = STPhase::Unspecified>	447	ReluPreMode reluPreMode = ReluPreMode::NoRelu, STPhase Phase = STPhase::Unspecified>
399	PTO_INTERNAL void TStoreAccNz2NC1HWC0(typename GlobalData::DType dstAddr, __cc__ typename TileData::DType srcAddr,	448	PTO_INTERNAL void TStoreAccNz2NC1HWC0(typename GlobalData::DType dstAddr, __cc__ typename TileData::DType srcAddr,
	@@ -462,6 +511,9 @@ __tf__ AICORE void TStoreAcc(typename GlobalData::DType __out__ *dst, typename T
462	} else if constexpr (GlobalData::layout == pto::Layout::NC1HWC0) {	511	} else if constexpr (GlobalData::layout == pto::Layout::NC1HWC0) {
463	TStoreAccNz2NC1HWC0<GlobalData, TileData, quantizationMode, reluPreMode, Phase>(	512	TStoreAccNz2NC1HWC0<GlobalData, TileData, quantizationMode, reluPreMode, Phase>(
464	dstAddr, srcAddr, gShape0, gShape1, gShape2, gShape3, gShape4, gStride1, gStride3, validRow, validCol);	513	dstAddr, srcAddr, gShape0, gShape1, gShape2, gShape3, gShape4, gStride1, gStride3, validRow, validCol);
		514	+ } else if constexpr (GlobalData::layout == pto::Layout::NDC1HWC0) {
		515	+ TStoreAccNz2NDC1HWC0<GlobalData, TileData, quantizationMode, reluPreMode, Phase>(
		516	+ dstAddr, srcAddr, gShape0, gShape1, gShape2, gShape3, gShape4, gStride2, gStride4, validRow, validCol);
465	}	517	}
466	}	518	}
467		519
	@@ -488,6 +540,10 @@ __tf__ AICORE void TStoreAccFp(typename GlobalData::DType __out__ *dst, typename
488	TStoreAccNz2NC1HWC0<GlobalData, TileData, quantizationMode, reluPreMode>(	540	TStoreAccNz2NC1HWC0<GlobalData, TileData, quantizationMode, reluPreMode>(
489	dst, __cce_get_tile_ptr(src), gShape0, gShape1, gShape2, gShape3, gShape4, gStride1, gStride3, validRow,	541	dst, __cce_get_tile_ptr(src), gShape0, gShape1, gShape2, gShape3, gShape4, gStride1, gStride3, validRow,
490	validCol);	542	validCol);
		543	+ } else if constexpr (GlobalData::layout == pto::Layout::NDC1HWC0) {
		544	+ TStoreAccNz2NDC1HWC0<GlobalData, TileData, quantizationMode, reluPreMode>(
		545	+ dst, __cce_get_tile_ptr(src), gShape0, gShape1, gShape2, gShape3, gShape4, gStride2, gStride4, validRow,
		546	+ validCol);
491	}	547	}
492	}	548	}
493		549
	@@ -526,8 +582,8 @@ template <typename TileData, typename GlobalData, bool isQuant>
526	PTO_INTERNAL void CheckAcc2gm(GlobalData &dst, TileData &src)	582	PTO_INTERNAL void CheckAcc2gm(GlobalData &dst, TileData &src)
527	{	583	{
528	static_assert((GlobalData::layout == pto::Layout::ND \|\| GlobalData::layout == pto::Layout::NZ \|\|	584	static_assert((GlobalData::layout == pto::Layout::ND \|\| GlobalData::layout == pto::Layout::NZ \|\|
529	- GlobalData::layout == pto::Layout::NC1HWC0),	585	+ GlobalData::layout == pto::Layout::NC1HWC0 \|\| GlobalData::layout == pto::Layout::NDC1HWC0),
530	- "The output data layout must be ND, NZ or NC1HWC0.");	586	+ "The output data layout must be ND, NZ, NC1HWC0 or NDC1HWC0.");
531	static_assert(std::is_same_v<typename TileData::DType, int32_t> \|\| std::is_same_v<typename TileData::DType, float>,	587	static_assert(std::is_same_v<typename TileData::DType, int32_t> \|\| std::is_same_v<typename TileData::DType, float>,
532	"The input data type must be restricted to int32_t/float!");	588	"The input data type must be restricted to int32_t/float!");
533	if constexpr (!isQuant) {	589	if constexpr (!isQuant) {
	@@ -550,10 +606,11 @@ PTO_INTERNAL void CheckAcc2gm(GlobalData &dst, TileData &src)
550	}	606	}
551	static_assert(TileData::Cols >= 1 && TileData::Cols <= 4095, "The range of Cols is [1, 4095].");	607	static_assert(TileData::Cols >= 1 && TileData::Cols <= 4095, "The range of Cols is [1, 4095].");
552	static_assert((GlobalData::layout == pto::Layout::ND && TileData::Rows >= 1 && TileData::Rows <= 8192) \|\|	608	static_assert((GlobalData::layout == pto::Layout::ND && TileData::Rows >= 1 && TileData::Rows <= 8192) \|\|
553	- ((GlobalData::layout == pto::Layout::NZ \|\| GlobalData::layout == pto::Layout::NC1HWC0) &&	609	+ ((GlobalData::layout == pto::Layout::NZ \|\| GlobalData::layout == pto::Layout::NC1HWC0 \|\|
		610	+ GlobalData::layout == pto::Layout::NDC1HWC0) &&
554	TileData::Rows >= 1 && TileData::Rows <= 65535 && TileData::Cols % 16 == 0),	611	TileData::Rows >= 1 && TileData::Rows <= 65535 && TileData::Cols % 16 == 0),
555	"When GlobalData is ND format, the range of Rows is [1, 8192]."	612	"When GlobalData is ND format, the range of Rows is [1, 8192]."
556	- "When GlobalData is NZ or NC1HWC0 format, the range of Rows is [1, 65535] and Cols "	613	+ "When GlobalData is NZ or NC1HWC0 or NDC1HWC0 format, the range of Rows is [1, 65535] and Cols "
557	"must be an integer multiple of 16.");	614	"must be an integer multiple of 16.");
558	PTO_ASSERT(src.GetValidCol() >= 1 && src.GetValidCol() <= 4095, "The range of validCol is [1, 4095].");	615	PTO_ASSERT(src.GetValidCol() >= 1 && src.GetValidCol() <= 4095, "The range of validCol is [1, 4095].");
559	PTO_ASSERT(dst.GetShape(pto::GlobalTensorDim::DIM_0) > 0 && dst.GetShape(pto::GlobalTensorDim::DIM_1) > 0 &&	616	PTO_ASSERT(dst.GetShape(pto::GlobalTensorDim::DIM_0) > 0 && dst.GetShape(pto::GlobalTensorDim::DIM_1) > 0 &&

Mtests/npu/a2a3/src/st/testcase/tstore_acc2gm/gen_data.py+82-4

	@@ -28,6 +28,12 @@ def ceil_div(num_1, num_2):
28	return (num_1 + num_2 - 1) // num_2	28	return (num_1 + num_2 - 1) // num_2
29		29
30		30
		31	+def get_c0_size(data_type):
		32	+ if np.dtype(data_type) == np.dtype(np.int32):
		33	+ return 16
		34	+ return 32 // np.dtype(data_type).itemsize
		35	+
		36	+
31	def saturation(arr, min_val, max_val, dtype):	37	def saturation(arr, min_val, max_val, dtype):
32	arr = np.clip(arr, min_val, max_val)	38	arr = np.clip(arr, min_val, max_val)
33	return arr.astype(dtype)	39	return arr.astype(dtype)
	@@ -111,6 +117,23 @@ def get_quant_golden(dst_data_type, m, n, quant_type, golden):
111	return quant_golden	117	return quant_golden
112		118
113		119
		120	+def nz_matrix_to_ndc1hwc0(golden, g_info):
		121	+ n_dim, d_dim, c1_dim, h_dim, w_dim, c0_dim = g_info.ndc1hwc0_shape
		122	+ c0_size = get_c0_size(g_info.dst_data_type)
		123	+ if c0_dim != c0_size:
		124	+ raise ValueError(f"NDC1HWC0 C0 dim {c0_dim} does not match c0_size {c0_size}.")
		125	+ if g_info.m != n_dim * h_dim * w_dim:
		126	+ raise ValueError("NDC1HWC0 requires m == N * H * W.")
		127	+ if g_info.n != d_dim * c1_dim * c0_dim:
		128	+ raise ValueError("NDC1HWC0 requires n == D * C1 * C0.")
		129	+
		130	+ return (
		131	+ golden.reshape(n_dim, h_dim, w_dim, d_dim, c1_dim, c0_dim)
		132	+ .transpose(0, 3, 4, 1, 2, 5)
		133	+ .astype(g_info.dst_data_type)
		134	+ )
		135	+
		136	+
114	def gen_golden_data(case_name, g_info):	137	def gen_golden_data(case_name, g_info):
115	src_data_type = g_info.src_data_type	138	src_data_type = g_info.src_data_type
116	dst_data_type = g_info.dst_data_type	139	dst_data_type = g_info.dst_data_type
	@@ -150,6 +173,8 @@ def gen_golden_data(case_name, g_info):
150	elif format == 3:	173	elif format == 3:
151	c0_size = 8	174	c0_size = 8
152	golden = golden.reshape(int(m / 16), 16, int(n / c0_size), c0_size).transpose(2, 0, 1, 3).astype(dst_data_type)	175	golden = golden.reshape(int(m / 16), 16, int(n / c0_size), c0_size).transpose(2, 0, 1, 3).astype(dst_data_type)
		176	+ elif format == 4:
		177	+ golden = nz_matrix_to_ndc1hwc0(golden, g_info)
153		178
154	if relu_mode == 1:	179	if relu_mode == 1:
155	golden = np.maximum(golden, 0)	180	golden = np.maximum(golden, 0)
	@@ -160,7 +185,20 @@ def gen_golden_data(case_name, g_info):
160		185
161		186
162	class TStoreAcc2gmParams:	187	class TStoreAcc2gmParams:
163	- def __init__(self, dst_data_type, src_data_type, format, m, n, k, quant_mode=0, scalar=1, quant_type=None, relu_mode=0):	188	+ def __init__(
		189	+ self,
		190	+ dst_data_type,
		191	+ src_data_type,
		192	+ format,
		193	+ m,
		194	+ n,
		195	+ k,
		196	+ quant_mode=0,
		197	+ scalar=1,
		198	+ quant_type=None,
		199	+ relu_mode=0,
		200	+ ndc1hwc0_shape=(0, 0, 0, 0, 0, 0),
		201	+ ):
164	self.src_data_type = src_data_type	202	self.src_data_type = src_data_type
165	self.dst_data_type = dst_data_type	203	self.dst_data_type = dst_data_type
166	self.format = format	204	self.format = format
	@@ -171,6 +209,7 @@ class TStoreAcc2gmParams:
171	self.scalar = scalar	209	self.scalar = scalar
172	self.quant_type = quant_type	210	self.quant_type = quant_type
173	self.relu_mode = relu_mode	211	self.relu_mode = relu_mode
		212	+ self.ndc1hwc0_shape = ndc1hwc0_shape
174		213
175	if __name__ == "__main__":	214	if __name__ == "__main__":
176	# 用例名称	215	# 用例名称
	@@ -214,7 +253,15 @@ if __name__ == "__main__":
214	"TStoreAcc2gmTest.case_relu_21",	253	"TStoreAcc2gmTest.case_relu_21",
215	"TStoreAcc2gmTest.case_relu_31",	254	"TStoreAcc2gmTest.case_relu_31",
216	"TStoreAcc2gmTest.case_relu_41",	255	"TStoreAcc2gmTest.case_relu_41",
217	- "TStoreAcc2gmTest.case_relu_51"	256	+ "TStoreAcc2gmTest.case_relu_51",
		257	+ "TStoreAcc2gmTest.case_ndc1hwc0_1",
		258	+ "TStoreAcc2gmTest.case_ndc1hwc0_2",
		259	+ "TStoreAcc2gmTest.case_ndc1hwc0_3",
		260	+ "TStoreAcc2gmTest.case_ndc1hwc0_relu_1",
		261	+ "TStoreAcc2gmTest.case_ndc1hwc0_scalar_1",
		262	+ "TStoreAcc2gmTest.case_ndc1hwc0_scalar_2",
		263	+ "TStoreAcc2gmTest.case_ndc1hwc0_vector_1",
		264	+ "TStoreAcc2gmTest.case_ndc1hwc0_vector_relu_1"
218	]	265	]
219		266
220	case_params_list = [	267	case_params_list = [
	@@ -264,7 +311,38 @@ if __name__ == "__main__":
264	TStoreAcc2gmParams(np.int8, np.float16, 1, 55, 27, 33, quant_mode=1, scalar=2, relu_mode=1),	311	TStoreAcc2gmParams(np.int8, np.float16, 1, 55, 27, 33, quant_mode=1, scalar=2, relu_mode=1),
265	TStoreAcc2gmParams(np.int8, np.int8, 2, 80, 96, 114, quant_mode=1, scalar=2, relu_mode=1),	312	TStoreAcc2gmParams(np.int8, np.int8, 2, 80, 96, 114, quant_mode=1, scalar=2, relu_mode=1),
266	TStoreAcc2gmParams(np.int8, np.float16, 1, 79, 63, 33, quant_mode=2, quant_type=np.uint64, relu_mode=1),	313	TStoreAcc2gmParams(np.int8, np.float16, 1, 79, 63, 33, quant_mode=2, quant_type=np.uint64, relu_mode=1),
267	- TStoreAcc2gmParams(np.int8, np.int8, 2, 80, 128, 90, quant_mode=2, quant_type=np.uint64, relu_mode=1)	314	+ TStoreAcc2gmParams(np.int8, np.int8, 2, 80, 128, 90, quant_mode=2, quant_type=np.uint64, relu_mode=1),
		315	+
		316	+ # NDC1HWC0
		317	+ TStoreAcc2gmParams(
		318	+ np.float32, np.float32, 4, 128, 64, 31, ndc1hwc0_shape=(1, 2, 4, 16, 8, 8)
		319	+ ),
		320	+ TStoreAcc2gmParams(
		321	+ np.int32, np.int8, 4, 256, 48, 27, ndc1hwc0_shape=(1, 3, 1, 16, 16, 16)
		322	+ ),
		323	+ TStoreAcc2gmParams(
		324	+ bfloat16, bfloat16, 4, 40, 96, 23, ndc1hwc0_shape=(1, 2, 3, 10, 4, 16)
		325	+ ),
		326	+ TStoreAcc2gmParams(
		327	+ np.float32, np.float16, 4, 63, 64, 32, relu_mode=1,
		328	+ ndc1hwc0_shape=(1, 1, 8, 9, 7, 8)
		329	+ ),
		330	+ TStoreAcc2gmParams(
		331	+ np.int8, np.float16, 4, 100, 64, 33, quant_mode=1, scalar=2, relu_mode=1,
		332	+ ndc1hwc0_shape=(1, 2, 1, 20, 5, 32)
		333	+ ),
		334	+ TStoreAcc2gmParams(
		335	+ np.uint8, np.float32, 4, 70, 64, 25, quant_mode=1, scalar=1.5,
		336	+ ndc1hwc0_shape=(1, 1, 2, 7, 10, 32)
		337	+ ),
		338	+ TStoreAcc2gmParams(
		339	+ np.float16, np.int8, 4, 55, 64, 32, quant_mode=2, quant_type=np.uint64,
		340	+ ndc1hwc0_shape=(1, 2, 2, 11, 5, 16)
		341	+ ),
		342	+ TStoreAcc2gmParams(
		343	+ np.int8, np.int8, 4, 52, 64, 19, quant_mode=2, quant_type=np.uint64, relu_mode=1,
		344	+ ndc1hwc0_shape=(1, 1, 2, 13, 4, 32)
		345	+ )
268		346
269	]	347	]
270		348
	@@ -274,4 +352,4 @@ if __name__ == "__main__":
274	original_dir = os.getcwd()	352	original_dir = os.getcwd()
275	os.chdir(case_name)	353	os.chdir(case_name)
276	gen_golden_data(case_name, case_params_list[i])	354	gen_golden_data(case_name, case_params_list[i])
277	- os.chdir(original_dir)	355	+ os.chdir(original_dir)

Mtests/npu/a2a3/src/st/testcase/tstore_acc2gm/main.cpp+252-0

	@@ -15,24 +15,40 @@ See LICENSE in the root of the software repository for the full text of the Lice
15	using namespace std;	15	using namespace std;
16	using namespace PtoTestCommon;	16	using namespace PtoTestCommon;
17		17
		18	+template <typename T>
		19	+constexpr int GetNDC1HWC0C0Size()
		20	+{
		21	+ return std::is_same_v<T, int32_t> ? 16 : 32 / sizeof(T);
		22	+}
		23	+
18	template <int tilingKey>	24	template <int tilingKey>
19	void LaunchTStoreAcc2gmNz2nd(uint8_t out, uint8_t src0, uint8_t src1, void stream);	25	void LaunchTStoreAcc2gmNz2nd(uint8_t out, uint8_t src0, uint8_t src1, void stream);
20		26
21	template <int tilingKey>	27	template <int tilingKey>
22	void LaunchTStoreAcc2gmNz2nz(uint8_t out, uint8_t src0, uint8_t src1, void stream);	28	void LaunchTStoreAcc2gmNz2nz(uint8_t out, uint8_t src0, uint8_t src1, void stream);
23		29
		30	+template <int tilingKey>
		31	+void LaunchTStoreAcc2gmNz2NDC1HWC0(uint8_t out, uint8_t src0, uint8_t src1, void stream);
		32	+
24	template <int tilingKey>	33	template <int tilingKey>
25	void LaunchTStoreAcc2gmScalarNz2nd(uint8_t out, uint8_t src0, uint8_t src1, void stream, float scalarQuant);	34	void LaunchTStoreAcc2gmScalarNz2nd(uint8_t out, uint8_t src0, uint8_t src1, void stream, float scalarQuant);
26		35
27	template <int tilingKey>	36	template <int tilingKey>
28	void LaunchTStoreAcc2gmScalarNz2nz(uint8_t out, uint8_t src0, uint8_t src1, void stream, float scalarQuant);	37	void LaunchTStoreAcc2gmScalarNz2nz(uint8_t out, uint8_t src0, uint8_t src1, void stream, float scalarQuant);
29		38
		39	+template <int tilingKey>
		40	+void LaunchTStoreAcc2gmScalarNz2NDC1HWC0(uint8_t out, uint8_t src0, uint8_t src1, void stream, float scalarQuant);
		41	+
30	template <int tilingKey>	42	template <int tilingKey>
31	void LaunchTStoreAcc2gmVectorNz2nd(uint8_t out, uint8_t src0, uint8_t src1, uint8_t quantTensor, void *stream);	43	void LaunchTStoreAcc2gmVectorNz2nd(uint8_t out, uint8_t src0, uint8_t src1, uint8_t quantTensor, void *stream);
32		44
33	template <int tilingKey>	45	template <int tilingKey>
34	void LaunchTStoreAcc2gmVectorNz2nz(uint8_t out, uint8_t src0, uint8_t src1, uint8_t quantTensor, void *stream);	46	void LaunchTStoreAcc2gmVectorNz2nz(uint8_t out, uint8_t src0, uint8_t src1, uint8_t quantTensor, void *stream);
35		47
		48	+template <int tilingKey>
		49	+void LaunchTStoreAcc2gmVectorNz2NDC1HWC0(uint8_t out, uint8_t src0, uint8_t src1, uint8_t quantTensor,
		50	+ void *stream);
		51	+
36	class TStoreAcc2gmTest : public testing::Test {	52	class TStoreAcc2gmTest : public testing::Test {
37	protected:	53	protected:
38	void SetUp() override	54	void SetUp() override
	@@ -163,6 +179,127 @@ void test_tstore_acc2gm_nz2nz()
163	EXPECT_TRUE(ret);	179	EXPECT_TRUE(ret);
164	}	180	}
165		181
		182	+template <int tilingKey, typename dstDataType, typename srcDataType, int validM, int validN, int validK, int dstN,
		183	+ int dstD, int dstC1, int dstH, int dstW>
		184	+void test_tstore_acc2gm_nz2ndc1hwc0()
		185	+{
		186	+ constexpr int c0Size = GetNDC1HWC0C0Size<dstDataType>();
		187	+ static_assert(validM == dstN * dstH * dstW, "NDC1HWC0 requires validM == N * H * W.");
		188	+ static_assert(validN == dstD * dstC1 * c0Size, "NDC1HWC0 requires validN == D * C1 * C0.");
		189	+ size_t aFileSize = validM * validK * sizeof(srcDataType);
		190	+ size_t bFileSize = validK * validN * sizeof(srcDataType);
		191	+ size_t cFileSize = dstN * dstD * dstC1 * dstH * dstW * c0Size * sizeof(dstDataType);
		192	+
		193	+ aclInit(nullptr);
		194	+ aclrtSetDevice(0);
		195	+
		196	+ aclrtStream stream;
		197	+ aclrtCreateStream(&stream);
		198	+
		199	+ uint8_t dstHost, src0Host, *src1Host;
		200	+ uint8_t dstDevice, src0Device, *src1Device;
		201	+
		202	+ aclrtMallocHost((void **)(&dstHost), cFileSize);
		203	+ aclrtMallocHost((void **)(&src0Host), aFileSize);
		204	+ aclrtMallocHost((void **)(&src1Host), bFileSize);
		205	+
		206	+ aclrtMalloc((void **)&dstDevice, cFileSize, ACL_MEM_MALLOC_HUGE_FIRST);
		207	+ aclrtMalloc((void **)&src0Device, aFileSize, ACL_MEM_MALLOC_HUGE_FIRST);
		208	+ aclrtMalloc((void **)&src1Device, bFileSize, ACL_MEM_MALLOC_HUGE_FIRST);
		209	+
		210	+ ReadFile(GetGoldenDir() + "/x1_gm.bin", aFileSize, src0Host, aFileSize);
		211	+ ReadFile(GetGoldenDir() + "/x2_gm.bin", bFileSize, src1Host, bFileSize);
		212	+
		213	+ aclrtMemcpy(src0Device, aFileSize, src0Host, aFileSize, ACL_MEMCPY_HOST_TO_DEVICE);
		214	+ aclrtMemcpy(src1Device, bFileSize, src1Host, bFileSize, ACL_MEMCPY_HOST_TO_DEVICE);
		215	+ LaunchTStoreAcc2gmNz2NDC1HWC0<tilingKey>(dstDevice, src0Device, src1Device, stream);
		216	+
		217	+ aclrtSynchronizeStream(stream);
		218	+ aclrtMemcpy(dstHost, cFileSize, dstDevice, cFileSize, ACL_MEMCPY_DEVICE_TO_HOST);
		219	+
		220	+ WriteFile(GetGoldenDir() + "/output_z.bin", dstHost, cFileSize);
		221	+
		222	+ aclrtFree(dstDevice);
		223	+ aclrtFree(src0Device);
		224	+ aclrtFree(src1Device);
		225	+
		226	+ aclrtFreeHost(dstHost);
		227	+ aclrtFreeHost(src0Host);
		228	+ aclrtFreeHost(src1Host);
		229	+
		230	+ aclrtDestroyStream(stream);
		231	+ aclrtResetDevice(0);
		232	+ aclFinalize();
		233	+
		234	+ std::vector<dstDataType> golden(cFileSize);
		235	+ std::vector<dstDataType> devFinal(cFileSize);
		236	+ ReadFile(GetGoldenDir() + "/golden.bin", cFileSize, golden.data(), cFileSize);
		237	+ ReadFile(GetGoldenDir() + "/output_z.bin", cFileSize, devFinal.data(), cFileSize);
		238	+
		239	+ bool ret = ResultCmp<dstDataType>(golden, devFinal, 0.001f);
		240	+ EXPECT_TRUE(ret);
		241	+}
		242	+
		243	+template <int tilingKey, typename dstDataType, typename srcDataType, int validM, int validN, int validK, int dstN,
		244	+ int dstD, int dstC1, int dstH, int dstW>
		245	+void test_tstore_acc2gm_scalar_nz2ndc1hwc0(float scalarQuant)
		246	+{
		247	+ constexpr int c0Size = GetNDC1HWC0C0Size<dstDataType>();
		248	+ static_assert(validM == dstN * dstH * dstW, "NDC1HWC0 requires validM == N * H * W.");
		249	+ static_assert(validN == dstD * dstC1 * c0Size, "NDC1HWC0 requires validN == D * C1 * C0.");
		250	+ size_t aFileSize = validM * validK * sizeof(srcDataType);
		251	+ size_t bFileSize = validK * validN * sizeof(srcDataType);
		252	+ size_t cFileSize = dstN * dstD * dstC1 * dstH * dstW * c0Size * sizeof(dstDataType);
		253	+
		254	+ aclInit(nullptr);
		255	+ aclrtSetDevice(0);
		256	+
		257	+ aclrtStream stream;
		258	+ aclrtCreateStream(&stream);
		259	+
		260	+ uint8_t dstHost, src0Host, *src1Host;
		261	+ uint8_t dstDevice, src0Device, *src1Device;
		262	+
		263	+ aclrtMallocHost((void **)(&dstHost), cFileSize);
		264	+ aclrtMallocHost((void **)(&src0Host), aFileSize);
		265	+ aclrtMallocHost((void **)(&src1Host), bFileSize);
		266	+
		267	+ aclrtMalloc((void **)&dstDevice, cFileSize, ACL_MEM_MALLOC_HUGE_FIRST);
		268	+ aclrtMalloc((void **)&src0Device, aFileSize, ACL_MEM_MALLOC_HUGE_FIRST);
		269	+ aclrtMalloc((void **)&src1Device, bFileSize, ACL_MEM_MALLOC_HUGE_FIRST);
		270	+
		271	+ ReadFile(GetGoldenDir() + "/x1_gm.bin", aFileSize, src0Host, aFileSize);
		272	+ ReadFile(GetGoldenDir() + "/x2_gm.bin", bFileSize, src1Host, bFileSize);
		273	+
		274	+ aclrtMemcpy(src0Device, aFileSize, src0Host, aFileSize, ACL_MEMCPY_HOST_TO_DEVICE);
		275	+ aclrtMemcpy(src1Device, bFileSize, src1Host, bFileSize, ACL_MEMCPY_HOST_TO_DEVICE);
		276	+ LaunchTStoreAcc2gmScalarNz2NDC1HWC0<tilingKey>(dstDevice, src0Device, src1Device, stream, scalarQuant);
		277	+ aclrtSynchronizeStream(stream);
		278	+ aclrtMemcpy(dstHost, cFileSize, dstDevice, cFileSize, ACL_MEMCPY_DEVICE_TO_HOST);
		279	+
		280	+ WriteFile(GetGoldenDir() + "/output_z.bin", dstHost, cFileSize);
		281	+
		282	+ aclrtFree(dstDevice);
		283	+ aclrtFree(src0Device);
		284	+ aclrtFree(src1Device);
		285	+
		286	+ aclrtFreeHost(dstHost);
		287	+ aclrtFreeHost(src0Host);
		288	+ aclrtFreeHost(src1Host);
		289	+
		290	+ aclrtDestroyStream(stream);
		291	+ aclrtResetDevice(0);
		292	+ aclFinalize();
		293	+
		294	+ std::vector<dstDataType> golden(cFileSize);
		295	+ std::vector<dstDataType> devFinal(cFileSize);
		296	+ ReadFile(GetGoldenDir() + "/golden.bin", cFileSize, golden.data(), cFileSize);
		297	+ ReadFile(GetGoldenDir() + "/output_z.bin", cFileSize, devFinal.data(), cFileSize);
		298	+
		299	+ bool ret = ResultCmp<dstDataType>(golden, devFinal, 0.001f);
		300	+ EXPECT_TRUE(ret);
		301	+}
		302	+
166	template <int tilingKey, typename dstDataType, typename srcDataType, int validM, int validN, int validK>	303	template <int tilingKey, typename dstDataType, typename srcDataType, int validM, int validN, int validK>
167	void test_tstore_acc2gm_scalar_nz2nd(float scalarQuant)	304	void test_tstore_acc2gm_scalar_nz2nd(float scalarQuant)
168	{	305	{
	@@ -219,6 +356,81 @@ void test_tstore_acc2gm_scalar_nz2nd(float scalarQuant)
219	EXPECT_TRUE(ret);	356	EXPECT_TRUE(ret);
220	}	357	}
221		358
		359	+template <int tilingKey, typename dstDataType, typename srcDataType, int validM, int validN, int validK, int dstN,
		360	+ int dstD, int dstC1, int dstH, int dstW>
		361	+void test_tstore_acc2gm_vector_nz2ndc1hwc0()
		362	+{
		363	+ using ScalingT = uint64_t;
		364	+ constexpr int c0Size = GetNDC1HWC0C0Size<dstDataType>();
		365	+ constexpr int alignFbN = (validN * sizeof(ScalingT) + 127) / 128 * 128 / sizeof(ScalingT);
		366	+ static_assert(validM == dstN * dstH * dstW, "NDC1HWC0 requires validM == N * H * W.");
		367	+ static_assert(validN == dstD * dstC1 * c0Size, "NDC1HWC0 requires validN == D * C1 * C0.");
		368	+ size_t aFileSize = validM * validK * sizeof(srcDataType);
		369	+ size_t bFileSize = validK * validN * sizeof(srcDataType);
		370	+ size_t cFileSize = dstN * dstD * dstC1 * dstH * dstW * c0Size * sizeof(dstDataType);
		371	+ size_t fbFileSize = alignFbN * sizeof(ScalingT);
		372	+
		373	+ aclInit(nullptr);
		374	+ aclrtSetDevice(0);
		375	+
		376	+ aclrtStream stream;
		377	+ aclrtCreateStream(&stream);
		378	+
		379	+ uint8_t *dstHost;
		380	+ uint8_t *src0Host;
		381	+ uint8_t *src1Host;
		382	+ uint8_t *quantTensorHost;
		383	+ uint8_t *dstDevice;
		384	+ uint8_t *src0Device;
		385	+ uint8_t *src1Device;
		386	+ uint8_t *quantTensorDevice;
		387	+
		388	+ aclrtMallocHost((void **)(&dstHost), cFileSize);
		389	+ aclrtMallocHost((void **)(&src0Host), aFileSize);
		390	+ aclrtMallocHost((void **)(&src1Host), bFileSize);
		391	+ aclrtMallocHost((void **)(&quantTensorHost), fbFileSize);
		392	+
		393	+ aclrtMalloc((void **)&dstDevice, cFileSize, ACL_MEM_MALLOC_HUGE_FIRST);
		394	+ aclrtMalloc((void **)&src0Device, aFileSize, ACL_MEM_MALLOC_HUGE_FIRST);
		395	+ aclrtMalloc((void **)&src1Device, bFileSize, ACL_MEM_MALLOC_HUGE_FIRST);
		396	+ aclrtMalloc((void **)&quantTensorDevice, fbFileSize, ACL_MEM_MALLOC_HUGE_FIRST);
		397	+
		398	+ ReadFile(GetGoldenDir() + "/x1_gm.bin", aFileSize, src0Host, aFileSize);
		399	+ ReadFile(GetGoldenDir() + "/x2_gm.bin", bFileSize, src1Host, bFileSize);
		400	+ ReadFile(GetGoldenDir() + "/quant_vector_gm.bin", fbFileSize, quantTensorHost, fbFileSize);
		401	+
		402	+ aclrtMemcpy(src0Device, aFileSize, src0Host, aFileSize, ACL_MEMCPY_HOST_TO_DEVICE);
		403	+ aclrtMemcpy(src1Device, bFileSize, src1Host, bFileSize, ACL_MEMCPY_HOST_TO_DEVICE);
		404	+ aclrtMemcpy(quantTensorDevice, fbFileSize, quantTensorHost, fbFileSize, ACL_MEMCPY_HOST_TO_DEVICE);
		405	+ LaunchTStoreAcc2gmVectorNz2NDC1HWC0<tilingKey>(dstDevice, src0Device, src1Device, quantTensorDevice, stream);
		406	+ aclrtSynchronizeStream(stream);
		407	+ aclrtMemcpy(dstHost, cFileSize, dstDevice, cFileSize, ACL_MEMCPY_DEVICE_TO_HOST);
		408	+
		409	+ WriteFile(GetGoldenDir() + "/output_z.bin", dstHost, cFileSize);
		410	+
		411	+ aclrtFree(dstDevice);
		412	+ aclrtFree(src0Device);
		413	+ aclrtFree(src1Device);
		414	+ aclrtFree(quantTensorDevice);
		415	+
		416	+ aclrtFreeHost(dstHost);
		417	+ aclrtFreeHost(src0Host);
		418	+ aclrtFreeHost(src1Host);
		419	+ aclrtFreeHost(quantTensorHost);
		420	+
		421	+ aclrtDestroyStream(stream);
		422	+ aclrtResetDevice(0);
		423	+ aclFinalize();
		424	+
		425	+ std::vector<dstDataType> golden(cFileSize);
		426	+ std::vector<dstDataType> devFinal(cFileSize);
		427	+ ReadFile(GetGoldenDir() + "/golden.bin", cFileSize, golden.data(), cFileSize);
		428	+ ReadFile(GetGoldenDir() + "/output_z.bin", cFileSize, devFinal.data(), cFileSize);
		429	+
		430	+ bool ret = ResultCmp<dstDataType>(golden, devFinal, 0.001f);
		431	+ EXPECT_TRUE(ret);
		432	+}
		433	+
222	template <int tilingKey, typename dstDataType, typename srcDataType, int validM, int validN, int validK>	434	template <int tilingKey, typename dstDataType, typename srcDataType, int validM, int validN, int validK>
223	void test_tstore_acc2gm_scalar_nz2nz(float scalarQuant)	435	void test_tstore_acc2gm_scalar_nz2nz(float scalarQuant)
224	{	436	{
	@@ -616,3 +828,43 @@ TEST_F(TStoreAcc2gmTest, case_relu_51)
616	{	828	{
617	test_tstore_acc2gm_vector_nz2nz<21, uint8_t, uint8_t, 80, 128, 90>();	829	test_tstore_acc2gm_vector_nz2nz<21, uint8_t, uint8_t, 80, 128, 90>();
618	}	830	}
		831	+
		832	+TEST_F(TStoreAcc2gmTest, case_ndc1hwc0_1)
		833	+{
		834	+ test_tstore_acc2gm_nz2ndc1hwc0<1, float, float, 128, 64, 31, 1, 2, 4, 16, 8>();
		835	+}
		836	+
		837	+TEST_F(TStoreAcc2gmTest, case_ndc1hwc0_2)
		838	+{
		839	+ test_tstore_acc2gm_nz2ndc1hwc0<2, int32_t, int8_t, 256, 48, 27, 1, 3, 1, 16, 16>();
		840	+}
		841	+
		842	+TEST_F(TStoreAcc2gmTest, case_ndc1hwc0_3)
		843	+{
		844	+ test_tstore_acc2gm_nz2ndc1hwc0<3, uint16_t, uint16_t, 40, 96, 23, 1, 2, 3, 10, 4>();
		845	+}
		846	+
		847	+TEST_F(TStoreAcc2gmTest, case_ndc1hwc0_relu_1)
		848	+{
		849	+ test_tstore_acc2gm_nz2ndc1hwc0<21, float, uint16_t, 63, 64, 32, 1, 1, 8, 9, 7>();
		850	+}
		851	+
		852	+TEST_F(TStoreAcc2gmTest, case_ndc1hwc0_scalar_1)
		853	+{
		854	+ test_tstore_acc2gm_scalar_nz2ndc1hwc0<1, int8_t, uint16_t, 100, 64, 33, 1, 2, 1, 20, 5>(2);
		855	+}
		856	+
		857	+TEST_F(TStoreAcc2gmTest, case_ndc1hwc0_scalar_2)
		858	+{
		859	+ test_tstore_acc2gm_scalar_nz2ndc1hwc0<2, uint8_t, float, 70, 64, 25, 1, 1, 2, 7, 10>(1.5);
		860	+}
		861	+
		862	+TEST_F(TStoreAcc2gmTest, case_ndc1hwc0_vector_1)
		863	+{
		864	+ test_tstore_acc2gm_vector_nz2ndc1hwc0<1, uint16_t, int8_t, 55, 64, 32, 1, 2, 2, 11, 5>();
		865	+}
		866	+
		867	+TEST_F(TStoreAcc2gmTest, case_ndc1hwc0_vector_relu_1)
		868	+{
		869	+ test_tstore_acc2gm_vector_nz2ndc1hwc0<21, int8_t, int8_t, 52, 64, 19, 1, 1, 2, 13, 4>();
		870	+}

Mtests/npu/a2a3/src/st/testcase/tstore_acc2gm/tstore_acc2gm_kernel.cpp+337-1

	@@ -181,6 +181,280 @@ __global__ AICORE void TStoreAcc2gmNz2nz(__gm__ dstDataType *out, __gm__ srcData
181	out = dstGlobal.data();	181	out = dstGlobal.data();
182	}	182	}
183		183
		184	+template <int atomicType, typename accDataType, typename dstDataType, typename srcDataType, int dstN, int dstD,
		185	+ int dstC1, int dstH, int dstW, int gWholeShape0, int gWholeShape1, int gWholeShape2, int gWholeShape3,
		186	+ int gWholeShape4, int validM, int validN, int validK, int reluMode = 0>
		187	+__global__ AICORE void TStoreAcc2gmNz2NDC1HWC0(__gm__ dstDataType out, __gm__ srcDataType src0,
		188	+ __gm__ srcDataType *src1)
		189	+{
		190	+ constexpr int dstC0 = std::is_same_v<dstDataType, int32_t> ? 16 : 32 / sizeof(dstDataType);
		191	+ constexpr int gStride[5] = {gWholeShape1 * gWholeShape2 * gWholeShape3 * gWholeShape4 * dstC0,
		192	+ gWholeShape2 * gWholeShape3 * gWholeShape4 * dstC0, gWholeShape3 * gWholeShape4 * dstC0,
		193	+ gWholeShape4 * dstC0, dstC0};
		194	+ constexpr int nAlign = BLOCK_CUBE_M_N > (32 / sizeof(srcDataType)) ? BLOCK_CUBE_M_N : (32 / sizeof(srcDataType));
		195	+ constexpr int M = (validM + BLOCK_CUBE_M_N - 1) / BLOCK_CUBE_M_N * BLOCK_CUBE_M_N;
		196	+ constexpr int N = (validN + nAlign - 1) / nAlign * nAlign;
		197	+ constexpr int K = (validK + BLOCK_CUBE_M_N - 1) / BLOCK_CUBE_M_N * BLOCK_CUBE_M_N;
		198	+
		199	+ constexpr int validRow = dstN * dstH * dstW;
		200	+ constexpr int validCol = dstD * dstC1 * dstC0;
		201	+ constexpr int Rows = M;
		202	+ constexpr int Cols = N;
		203	+ static_assert(validM == validRow, "NDC1HWC0 requires validM == N * H * W.");
		204	+ static_assert(validN == validCol, "NDC1HWC0 requires validN == D * C1 * C0.");
		205	+
		206	+ using DynShapeDim5 = pto::Shape<dstN, dstD, dstC1, dstH, dstW>;
		207	+ using DynStridDim5 = pto::Stride<gStride[0], gStride[1], gStride[2], gStride[3], gStride[4]>;
		208	+ using GlobalDataOut = GlobalTensor<dstDataType, DynShapeDim5, DynStridDim5, Layout::NDC1HWC0>;
		209	+ using GlobalDataSrc0 =
		210	+ GlobalTensor<srcDataType, pto::Shape<1, 1, 1, validM, validK>,
		211	+ pto::Stride<1 * validM * validK, 1 * validM * validK, validM * validK, validK, 1>>;
		212	+ using GlobalDataSrc1 =
		213	+ GlobalTensor<srcDataType, pto::Shape<1, 1, 1, validK, validN>,
		214	+ pto::Stride<1 * validK * validN, 1 * validK * validN, validK * validN, validN, 1>>;
		215	+
		216	+ GlobalDataSrc0 src0Global(src0);
		217	+ GlobalDataSrc1 src1Global(src1);
		218	+ GlobalDataOut dstGlobal(out);
		219	+
		220	+ using TileMatAData = Tile<TileType::Mat, srcDataType, M, K, BLayout::ColMajor, M, K, SLayout::RowMajor, 512>;
		221	+ using TileMatBData = Tile<TileType::Mat, srcDataType, K, N, BLayout::ColMajor, K, N, SLayout::RowMajor, 512>;
		222	+ using LeftTile = Tile<TileType::Left, srcDataType, M, K, BLayout::RowMajor, validM, validK, SLayout::RowMajor, 512>;
		223	+ using RightTile = TileRight<srcDataType, K, N, validK, validN>;
		224	+ using AccTile = Tile<TileType::Acc, accDataType, Rows, Cols, BLayout::ColMajor, -1, -1, SLayout::RowMajor, 1024>;
		225	+
		226	+ uint32_t aMatSize = M * K * sizeof(srcDataType);
		227	+ TileMatAData aMatTile;
		228	+ TileMatBData bMatTile;
		229	+ TASSIGN(aMatTile, 0x0);
		230	+ TASSIGN(bMatTile, aMatSize);
		231	+
		232	+ LeftTile aTile;
		233	+ RightTile bTile;
		234	+ AccTile cTile(validRow, validCol);
		235	+ TASSIGN(aTile, 0x0);
		236	+ TASSIGN(bTile, 0x0);
		237	+ TASSIGN(cTile, 0x0);
		238	+
		239	+ TLOAD(aMatTile, src0Global);
		240	+ TLOAD(bMatTile, src1Global);
		241	+
		242	+#ifndef __PTO_AUTO__
		243	+ set_flag(PIPE_MTE2, PIPE_MTE1, EVENT_ID0);
		244	+ wait_flag(PIPE_MTE2, PIPE_MTE1, EVENT_ID0);
		245	+#endif
		246	+ TMOV(aTile, aMatTile);
		247	+ TMOV(bTile, bMatTile);
		248	+#ifndef __PTO_AUTO__
		249	+ set_flag(PIPE_MTE1, PIPE_M, EVENT_ID0);
		250	+ wait_flag(PIPE_MTE1, PIPE_M, EVENT_ID0);
		251	+#endif
		252	+ TMATMUL(cTile, aTile, bTile);
		253	+#ifndef __PTO_AUTO__
		254	+ set_flag(PIPE_M, PIPE_FIX, EVENT_ID0);
		255	+ wait_flag(PIPE_M, PIPE_FIX, EVENT_ID0);
		256	+#endif
		257	+ constexpr AtomicType atomicTypeEnum = atomicType == 1 ? AtomicType::AtomicAdd : AtomicType::AtomicNone;
		258	+ if constexpr (reluMode == 0) {
		259	+ TSTORE<AccTile, GlobalDataOut, atomicTypeEnum>(dstGlobal, cTile);
		260	+ } else if constexpr (reluMode == 1) {
		261	+ constexpr ReluPreMode reluPreMode = ReluPreMode::NormalRelu;
		262	+ TSTORE<AccTile, GlobalDataOut, atomicTypeEnum, reluPreMode>(dstGlobal, cTile);
		263	+ }
		264	+#ifndef __PTO_AUTO__
		265	+ set_flag(PIPE_FIX, PIPE_M, EVENT_ID0);
		266	+ wait_flag(PIPE_FIX, PIPE_M, EVENT_ID0);
		267	+#endif
		268	+ out = dstGlobal.data();
		269	+}
		270	+
		271	+template <int atomicType, typename accDataType, typename dstDataType, typename srcDataType, int dstN, int dstD,
		272	+ int dstC1, int dstH, int dstW, int gWholeShape0, int gWholeShape1, int gWholeShape2, int gWholeShape3,
		273	+ int gWholeShape4, int validM, int validN, int validK, int reluMode = 0>
		274	+__global__ AICORE void TStoreAcc2gmScalarNz2NDC1HWC0(__gm__ dstDataType out, __gm__ srcDataType src0,
		275	+ __gm__ srcDataType *src1, float scalarQuant)
		276	+{
		277	+ constexpr int dstC0 = std::is_same_v<dstDataType, int32_t> ? 16 : 32 / sizeof(dstDataType);
		278	+ constexpr int gStride[5] = {gWholeShape1 * gWholeShape2 * gWholeShape3 * gWholeShape4 * dstC0,
		279	+ gWholeShape2 * gWholeShape3 * gWholeShape4 * dstC0, gWholeShape3 * gWholeShape4 * dstC0,
		280	+ gWholeShape4 * dstC0, dstC0};
		281	+ constexpr int nAlign = BLOCK_CUBE_M_N > (32 / sizeof(srcDataType)) ? BLOCK_CUBE_M_N : (32 / sizeof(srcDataType));
		282	+ constexpr int M = (validM + BLOCK_CUBE_M_N - 1) / BLOCK_CUBE_M_N * BLOCK_CUBE_M_N;
		283	+ constexpr int N = (validN + nAlign - 1) / nAlign * nAlign;
		284	+ constexpr int K = (validK + BLOCK_CUBE_M_N - 1) / BLOCK_CUBE_M_N * BLOCK_CUBE_M_N;
		285	+
		286	+ constexpr int validRow = dstN * dstH * dstW;
		287	+ constexpr int validCol = dstD * dstC1 * dstC0;
		288	+ static_assert(validM == validRow, "NDC1HWC0 requires validM == N * H * W.");
		289	+ static_assert(validN == validCol, "NDC1HWC0 requires validN == D * C1 * C0.");
		290	+
		291	+ using DynShapeDim5 = pto::Shape<dstN, dstD, dstC1, dstH, dstW>;
		292	+ using DynStridDim5 = pto::Stride<gStride[0], gStride[1], gStride[2], gStride[3], gStride[4]>;
		293	+ using GlobalDataOut = GlobalTensor<dstDataType, DynShapeDim5, DynStridDim5, Layout::NDC1HWC0>;
		294	+ using GlobalDataSrc0 =
		295	+ GlobalTensor<srcDataType, pto::Shape<1, 1, 1, validM, validK>,
		296	+ pto::Stride<1 * validM * validK, 1 * validM * validK, validM * validK, validK, 1>>;
		297	+ using GlobalDataSrc1 =
		298	+ GlobalTensor<srcDataType, pto::Shape<1, 1, 1, validK, validN>,
		299	+ pto::Stride<1 * validK * validN, 1 * validK * validN, validK * validN, validN, 1>>;
		300	+
		301	+ GlobalDataSrc0 src0Global(src0);
		302	+ GlobalDataSrc1 src1Global(src1);
		303	+ GlobalDataOut dstGlobal(out);
		304	+
		305	+ using TileMatAData = Tile<TileType::Mat, srcDataType, M, K, BLayout::ColMajor, M, K, SLayout::RowMajor, 512>;
		306	+ using TileMatBData = Tile<TileType::Mat, srcDataType, K, N, BLayout::ColMajor, K, N, SLayout::RowMajor, 512>;
		307	+ using LeftTile = Tile<TileType::Left, srcDataType, M, K, BLayout::RowMajor, validM, validK, SLayout::RowMajor, 512>;
		308	+ using RightTile = TileRight<srcDataType, K, N, validK, validN>;
		309	+ using AccTile = Tile<TileType::Acc, accDataType, M, N, BLayout::ColMajor, -1, -1, SLayout::RowMajor, 1024>;
		310	+
		311	+ uint32_t aMatSize = M * K * sizeof(srcDataType);
		312	+ TileMatAData aMatTile;
		313	+ TileMatBData bMatTile;
		314	+ TASSIGN(aMatTile, 0x0);
		315	+ TASSIGN(bMatTile, aMatSize);
		316	+
		317	+ LeftTile aTile;
		318	+ RightTile bTile;
		319	+ AccTile cTile(validRow, validCol);
		320	+ TASSIGN(aTile, 0x0);
		321	+ TASSIGN(bTile, 0x0);
		322	+ TASSIGN(cTile, 0x0);
		323	+
		324	+ TLOAD(aMatTile, src0Global);
		325	+ TLOAD(bMatTile, src1Global);
		326	+#ifndef __PTO_AUTO__
		327	+ set_flag(PIPE_MTE2, PIPE_MTE1, EVENT_ID0);
		328	+ wait_flag(PIPE_MTE2, PIPE_MTE1, EVENT_ID0);
		329	+#endif
		330	+ TMOV(aTile, aMatTile);
		331	+ TMOV(bTile, bMatTile);
		332	+#ifndef __PTO_AUTO__
		333	+ set_flag(PIPE_MTE1, PIPE_M, EVENT_ID0);
		334	+ wait_flag(PIPE_MTE1, PIPE_M, EVENT_ID0);
		335	+#endif
		336	+ TMATMUL(cTile, aTile, bTile);
		337	+#ifndef __PTO_AUTO__
		338	+ set_flag(PIPE_M, PIPE_FIX, EVENT_ID0);
		339	+ wait_flag(PIPE_M, PIPE_FIX, EVENT_ID0);
		340	+#endif
		341	+ uint64_t preQuantScalar = static_cast<uint64_t>(reinterpret_cast<int32_t >(&scalarQuant));
		342	+ if (sizeof(dstDataType) == 1) {
		343	+ constexpr bool sign = (std::is_same_v<dstDataType, int8_t>) ? true : false;
		344	+ preQuantScalar = (preQuantScalar & ~(static_cast<uint64_t>(1) << 46)) \| (static_cast<uint64_t>(sign) << 46);
		345	+ }
		346	+ constexpr AtomicType atomicTypeEnum = atomicType == 1 ? AtomicType::AtomicAdd : AtomicType::AtomicNone;
		347	+ if constexpr (reluMode == 0) {
		348	+ TSTORE<AccTile, GlobalDataOut, atomicTypeEnum>(dstGlobal, cTile, preQuantScalar);
		349	+ } else if constexpr (reluMode == 1) {
		350	+ constexpr ReluPreMode reluPreMode = ReluPreMode::NormalRelu;
		351	+ TSTORE<AccTile, GlobalDataOut, atomicTypeEnum, reluPreMode>(dstGlobal, cTile, preQuantScalar);
		352	+ }
		353	+#ifndef __PTO_AUTO__
		354	+ set_flag(PIPE_FIX, PIPE_M, EVENT_ID0);
		355	+ wait_flag(PIPE_FIX, PIPE_M, EVENT_ID0);
		356	+#endif
		357	+ out = dstGlobal.data();
		358	+}
		359	+
		360	+template <int atomicType, typename accDataType, typename dstDataType, typename srcDataType, int dstN, int dstD,
		361	+ int dstC1, int dstH, int dstW, int gWholeShape0, int gWholeShape1, int gWholeShape2, int gWholeShape3,
		362	+ int gWholeShape4, int validM, int validN, int validK, int reluMode = 0>
		363	+__global__ AICORE void TStoreAcc2gmVectorNz2NDC1HWC0(__gm__ dstDataType out, __gm__ srcDataType src0,
		364	+ __gm__ srcDataType src1, __gm__ uint64_t quantTensor)
		365	+{
		366	+ constexpr int dstC0 = std::is_same_v<dstDataType, int32_t> ? 16 : 32 / sizeof(dstDataType);
		367	+ constexpr int gStride[5] = {gWholeShape1 * gWholeShape2 * gWholeShape3 * gWholeShape4 * dstC0,
		368	+ gWholeShape2 * gWholeShape3 * gWholeShape4 * dstC0, gWholeShape3 * gWholeShape4 * dstC0,
		369	+ gWholeShape4 * dstC0, dstC0};
		370	+ constexpr int nAlign = BLOCK_CUBE_M_N > (32 / sizeof(srcDataType)) ? BLOCK_CUBE_M_N : (32 / sizeof(srcDataType));
		371	+ constexpr int M = (validM + BLOCK_CUBE_M_N - 1) / BLOCK_CUBE_M_N * BLOCK_CUBE_M_N;
		372	+ constexpr int N = (validN + nAlign - 1) / nAlign * nAlign;
		373	+ constexpr int K = (validK + BLOCK_CUBE_M_N - 1) / BLOCK_CUBE_M_N * BLOCK_CUBE_M_N;
		374	+ constexpr int alignScalingN = ((validN * sizeof(uint64_t) + 127) / 128) * 128 / sizeof(uint64_t);
		375	+
		376	+ constexpr int validRow = dstN * dstH * dstW;
		377	+ constexpr int validCol = dstD * dstC1 * dstC0;
		378	+ static_assert(validM == validRow, "NDC1HWC0 requires validM == N * H * W.");
		379	+ static_assert(validN == validCol, "NDC1HWC0 requires validN == D * C1 * C0.");
		380	+
		381	+ using DynShapeDim5 = pto::Shape<dstN, dstD, dstC1, dstH, dstW>;
		382	+ using DynStridDim5 = pto::Stride<gStride[0], gStride[1], gStride[2], gStride[3], gStride[4]>;
		383	+ using GlobalDataOut = GlobalTensor<dstDataType, DynShapeDim5, DynStridDim5, Layout::NDC1HWC0>;
		384	+ using GlobalDataSrc0 =
		385	+ GlobalTensor<srcDataType, pto::Shape<1, 1, 1, validM, validK>,
		386	+ pto::Stride<1 * validM * validK, 1 * validM * validK, validM * validK, validK, 1>>;
		387	+ using GlobalDataSrc1 =
		388	+ GlobalTensor<srcDataType, pto::Shape<1, 1, 1, validK, validN>,
		389	+ pto::Stride<1 * validK * validN, 1 * validK * validN, validK * validN, validN, 1>>;
		390	+ using GlobalDataSrc2 = GlobalTensor<uint64_t, Shape<1, 1, 1, 1, alignScalingN>,
		391	+ Stride<1 * alignScalingN, 1 * alignScalingN, alignScalingN, alignScalingN, 1>>;
		392	+
		393	+ GlobalDataSrc0 src0Global(src0);
		394	+ GlobalDataSrc1 src1Global(src1);
		395	+ GlobalDataSrc2 src2Global(quantTensor);
		396	+ GlobalDataOut dstGlobal(out);
		397	+
		398	+ using TileMatAData = Tile<TileType::Mat, srcDataType, M, K, BLayout::ColMajor, M, K, SLayout::RowMajor, 512>;
		399	+ using TileMatBData = Tile<TileType::Mat, srcDataType, K, N, BLayout::ColMajor, K, N, SLayout::RowMajor, 512>;
		400	+ using TileMatScalingData =
		401	+ Tile<TileType::Mat, uint64_t, 1, alignScalingN, BLayout::RowMajor, 1, -1, SLayout::NoneBox>;
		402	+ using LeftTile = Tile<TileType::Left, srcDataType, M, K, BLayout::RowMajor, validM, validK, SLayout::RowMajor, 512>;
		403	+ using RightTile = TileRight<srcDataType, K, N, validK, validN>;
		404	+ using AccTile = Tile<TileType::Acc, accDataType, M, N, BLayout::ColMajor, -1, -1, SLayout::RowMajor, 1024>;
		405	+ using ScalingTile = Tile<TileType::Scaling, uint64_t, 1, alignScalingN, BLayout::RowMajor, 1, -1, SLayout::NoneBox>;
		406	+
		407	+ uint32_t aMatSize = M * K * sizeof(srcDataType);
		408	+ uint32_t bMatSize = K * N * sizeof(srcDataType);
		409	+ TileMatAData aMatTile;
		410	+ TileMatBData bMatTile;
		411	+ TileMatScalingData scalingMatTile(alignScalingN);
		412	+ TASSIGN(aMatTile, 0x0);
		413	+ TASSIGN(bMatTile, aMatSize);
		414	+ TASSIGN(scalingMatTile, aMatSize + bMatSize);
		415	+
		416	+ LeftTile aTile;
		417	+ RightTile bTile;
		418	+ AccTile cTile(validRow, validCol);
		419	+ ScalingTile scalingTile(validN);
		420	+ TASSIGN(aTile, 0x0);
		421	+ TASSIGN(bTile, 0x0);
		422	+ TASSIGN(cTile, 0x0);
		423	+ TASSIGN(scalingTile, 0x0);
		424	+
		425	+ TLOAD(aMatTile, src0Global);
		426	+ TLOAD(bMatTile, src1Global);
		427	+ TLOAD(scalingMatTile, src2Global);
		428	+#ifndef __PTO_AUTO__
		429	+ set_flag(PIPE_MTE2, PIPE_MTE1, EVENT_ID0);
		430	+ wait_flag(PIPE_MTE2, PIPE_MTE1, EVENT_ID0);
		431	+#endif
		432	+ TMOV(aTile, aMatTile);
		433	+ TMOV(bTile, bMatTile);
		434	+#ifndef __PTO_AUTO__
		435	+ set_flag(PIPE_MTE1, PIPE_M, EVENT_ID0);
		436	+ wait_flag(PIPE_MTE1, PIPE_M, EVENT_ID0);
		437	+#endif
		438	+ TMATMUL(cTile, aTile, bTile);
		439	+#ifndef __PTO_AUTO__
		440	+ set_flag(PIPE_M, PIPE_FIX, EVENT_ID0);
		441	+ wait_flag(PIPE_M, PIPE_FIX, EVENT_ID0);
		442	+#endif
		443	+ TMOV(scalingTile, scalingMatTile);
		444	+ constexpr AtomicType atomicTypeEnum = atomicType == 1 ? AtomicType::AtomicAdd : AtomicType::AtomicNone;
		445	+ if constexpr (reluMode == 0) {
		446	+ TSTORE_FP<AccTile, GlobalDataOut, ScalingTile, atomicTypeEnum>(dstGlobal, cTile, scalingTile);
		447	+ } else if constexpr (reluMode == 1) {
		448	+ constexpr ReluPreMode reluPreMode = ReluPreMode::NormalRelu;
		449	+ TSTORE_FP<AccTile, GlobalDataOut, ScalingTile, atomicTypeEnum, reluPreMode>(dstGlobal, cTile, scalingTile);
		450	+ }
		451	+#ifndef __PTO_AUTO__
		452	+ set_flag(PIPE_FIX, PIPE_M, EVENT_ID0);
		453	+ wait_flag(PIPE_FIX, PIPE_M, EVENT_ID0);
		454	+#endif
		455	+ out = dstGlobal.data();
		456	+}
		457	+
184	template <int atomicType, typename accDataType, typename dstDataType, typename srcDataType, int gShape0, int gShape1,	458	template <int atomicType, typename accDataType, typename dstDataType, typename srcDataType, int gShape0, int gShape1,
185	int gShape2, int gShape3, int gShape4, int gWholeShape0, int gWholeShape1, int gWholeShape2, int gWholeShape3,	459	int gShape2, int gShape3, int gShape4, int gWholeShape0, int gWholeShape1, int gWholeShape2, int gWholeShape3,
186	int gWholeShape4, int validM, int validN, int validK, int reluMode = 0>	460	int gWholeShape4, int validM, int validN, int validK, int reluMode = 0>
	@@ -630,6 +904,56 @@ void LaunchTStoreAcc2gmNz2nz(uint8_t out, uint8_t src0, uint8_t src1, void s
630	}	904	}
631	}	905	}
632		906
		907	+template <int tilingKey>
		908	+void LaunchTStoreAcc2gmNz2NDC1HWC0(uint8_t out, uint8_t src0, uint8_t src1, void stream)
		909	+{
		910	+ if constexpr (tilingKey == 1) {
		911	+ TStoreAcc2gmNz2NDC1HWC0<0, float, float, float, 1, 2, 4, 16, 8, 1, 2, 4, 16, 8, 128, 64, 31>
		912	+ <<<1, nullptr, stream>>>(reinterpret_cast<float >(out), reinterpret_cast<float >(src0),
		913	+ reinterpret_cast<float *>(src1));
		914	+ } else if constexpr (tilingKey == 2) {
		915	+ TStoreAcc2gmNz2NDC1HWC0<0, int32_t, int32_t, int8_t, 1, 3, 1, 16, 16, 1, 3, 1, 16, 16, 256, 48, 27>
		916	+ <<<1, nullptr, stream>>>(reinterpret_cast<int32_t >(out), reinterpret_cast<int8_t >(src0),
		917	+ reinterpret_cast<int8_t *>(src1));
		918	+ } else if constexpr (tilingKey == 3) {
		919	+ TStoreAcc2gmNz2NDC1HWC0<0, float, bfloat16_t, bfloat16_t, 1, 2, 3, 10, 4, 1, 2, 3, 10, 4, 40, 96, 23>
		920	+ <<<1, nullptr, stream>>>(reinterpret_cast<bfloat16_t >(out), reinterpret_cast<bfloat16_t >(src0),
		921	+ reinterpret_cast<bfloat16_t *>(src1));
		922	+ } else if constexpr (tilingKey == 21) {
		923	+ TStoreAcc2gmNz2NDC1HWC0<0, float, float, half, 1, 1, 8, 9, 7, 1, 1, 8, 9, 7, 63, 64, 32, 1>
		924	+ <<<1, nullptr, stream>>>(reinterpret_cast<float >(out), reinterpret_cast<half >(src0),
		925	+ reinterpret_cast<half *>(src1));
		926	+ }
		927	+}
		928	+
		929	+template <int tilingKey>
		930	+void LaunchTStoreAcc2gmScalarNz2NDC1HWC0(uint8_t out, uint8_t src0, uint8_t src1, void stream, float scalarQuant)
		931	+{
		932	+ if constexpr (tilingKey == 1) {
		933	+ TStoreAcc2gmScalarNz2NDC1HWC0<0, float, int8_t, half, 1, 2, 1, 20, 5, 1, 2, 1, 20, 5, 100, 64, 33, 1>
		934	+ <<<1, nullptr, stream>>>(reinterpret_cast<int8_t >(out), reinterpret_cast<half >(src0),
		935	+ reinterpret_cast<half *>(src1), scalarQuant);
		936	+ } else if constexpr (tilingKey == 2) {
		937	+ TStoreAcc2gmScalarNz2NDC1HWC0<0, float, uint8_t, float, 1, 1, 2, 7, 10, 1, 1, 2, 7, 10, 70, 64, 25>
		938	+ <<<1, nullptr, stream>>>(reinterpret_cast<uint8_t >(out), reinterpret_cast<float >(src0),
		939	+ reinterpret_cast<float *>(src1), scalarQuant);
		940	+ }
		941	+}
		942	+
		943	+template <int tilingKey>
		944	+void LaunchTStoreAcc2gmVectorNz2NDC1HWC0(uint8_t out, uint8_t src0, uint8_t src1, uint8_t quantTensor, void *stream)
		945	+{
		946	+ if constexpr (tilingKey == 1) {
		947	+ TStoreAcc2gmVectorNz2NDC1HWC0<0, int32_t, half, int8_t, 1, 2, 2, 11, 5, 1, 2, 2, 11, 5, 55, 64, 32>
		948	+ <<<1, nullptr, stream>>>(reinterpret_cast<half >(out), reinterpret_cast<int8_t >(src0),
		949	+ reinterpret_cast<int8_t >(src1), reinterpret_cast<uint64_t >(quantTensor));
		950	+ } else if constexpr (tilingKey == 21) {
		951	+ TStoreAcc2gmVectorNz2NDC1HWC0<0, int32_t, int8_t, int8_t, 1, 1, 2, 13, 4, 1, 1, 2, 13, 4, 52, 64, 19, 1>
		952	+ <<<1, nullptr, stream>>>(reinterpret_cast<int8_t >(out), reinterpret_cast<int8_t >(src0),
		953	+ reinterpret_cast<int8_t >(src1), reinterpret_cast<uint64_t >(quantTensor));
		954	+ }
		955	+}
		956	+
633	template <int tilingKey>	957	template <int tilingKey>
634	void LaunchTStoreAcc2gmScalarNz2nd(uint8_t out, uint8_t src0, uint8_t src1, void stream, float scalarQuant)	958	void LaunchTStoreAcc2gmScalarNz2nd(uint8_t out, uint8_t src0, uint8_t src1, void stream, float scalarQuant)
635	{	959	{
	@@ -761,6 +1085,9 @@ template void LaunchTStoreAcc2gmNz2nz<5>(uint8_t out, uint8_t src0, uint8_t *s
761	template void LaunchTStoreAcc2gmNz2nz<6>(uint8_t out, uint8_t src0, uint8_t src1, void stream);	1085	template void LaunchTStoreAcc2gmNz2nz<6>(uint8_t out, uint8_t src0, uint8_t src1, void stream);
762	template void LaunchTStoreAcc2gmNz2nz<7>(uint8_t out, uint8_t src0, uint8_t src1, void stream);	1086	template void LaunchTStoreAcc2gmNz2nz<7>(uint8_t out, uint8_t src0, uint8_t src1, void stream);
763	template void LaunchTStoreAcc2gmNz2nz<8>(uint8_t out, uint8_t src0, uint8_t src1, void stream);	1087	template void LaunchTStoreAcc2gmNz2nz<8>(uint8_t out, uint8_t src0, uint8_t src1, void stream);
		1088	+template void LaunchTStoreAcc2gmNz2NDC1HWC0<1>(uint8_t out, uint8_t src0, uint8_t src1, void stream);
		1089	+template void LaunchTStoreAcc2gmNz2NDC1HWC0<2>(uint8_t out, uint8_t src0, uint8_t src1, void stream);
		1090	+template void LaunchTStoreAcc2gmNz2NDC1HWC0<3>(uint8_t out, uint8_t src0, uint8_t src1, void stream);
764		1091
765	template void LaunchTStoreAcc2gmScalarNz2nd<1>(uint8_t out, uint8_t src0, uint8_t src1, void stream,	1092	template void LaunchTStoreAcc2gmScalarNz2nd<1>(uint8_t out, uint8_t src0, uint8_t src1, void stream,
766	float scalarQuant);	1093	float scalarQuant);
	@@ -782,6 +1109,10 @@ template void LaunchTStoreAcc2gmScalarNz2nz<4>(uint8_t out, uint8_t src0, uint
782	float scalarQuant);	1109	float scalarQuant);
783	template void LaunchTStoreAcc2gmScalarNz2nz<5>(uint8_t out, uint8_t src0, uint8_t src1, void stream,	1110	template void LaunchTStoreAcc2gmScalarNz2nz<5>(uint8_t out, uint8_t src0, uint8_t src1, void stream,
784	float scalarQuant);	1111	float scalarQuant);
		1112	+template void LaunchTStoreAcc2gmScalarNz2NDC1HWC0<1>(uint8_t out, uint8_t src0, uint8_t src1, void stream,
		1113	+ float scalarQuant);
		1114	+template void LaunchTStoreAcc2gmScalarNz2NDC1HWC0<2>(uint8_t out, uint8_t src0, uint8_t src1, void stream,
		1115	+ float scalarQuant);
785		1116
786	template void LaunchTStoreAcc2gmVectorNz2nd<1>(uint8_t out, uint8_t src0, uint8_t src1, uint8_t quantTensor,	1117	template void LaunchTStoreAcc2gmVectorNz2nd<1>(uint8_t out, uint8_t src0, uint8_t src1, uint8_t quantTensor,
787	void *stream);	1118	void *stream);
	@@ -801,10 +1132,13 @@ template void LaunchTStoreAcc2gmVectorNz2nz<3>(uint8_t out, uint8_t src0, uint
801	void *stream);	1132	void *stream);
802	template void LaunchTStoreAcc2gmVectorNz2nz<4>(uint8_t out, uint8_t src0, uint8_t src1, uint8_t quantTensor,	1133	template void LaunchTStoreAcc2gmVectorNz2nz<4>(uint8_t out, uint8_t src0, uint8_t src1, uint8_t quantTensor,
803	void *stream);	1134	void *stream);
		1135	+template void LaunchTStoreAcc2gmVectorNz2NDC1HWC0<1>(uint8_t out, uint8_t src0, uint8_t src1, uint8_t quantTensor,
		1136	+ void *stream);
804		1137
805	// relu	1138	// relu
806	template void LaunchTStoreAcc2gmNz2nd<21>(uint8_t out, uint8_t src0, uint8_t src1, void stream);	1139	template void LaunchTStoreAcc2gmNz2nd<21>(uint8_t out, uint8_t src0, uint8_t src1, void stream);
807	template void LaunchTStoreAcc2gmNz2nz<21>(uint8_t out, uint8_t src0, uint8_t src1, void stream);	1140	template void LaunchTStoreAcc2gmNz2nz<21>(uint8_t out, uint8_t src0, uint8_t src1, void stream);
		1141	+template void LaunchTStoreAcc2gmNz2NDC1HWC0<21>(uint8_t out, uint8_t src0, uint8_t src1, void stream);
808		1142
809	template void LaunchTStoreAcc2gmScalarNz2nd<21>(uint8_t out, uint8_t src0, uint8_t src1, void stream,	1143	template void LaunchTStoreAcc2gmScalarNz2nd<21>(uint8_t out, uint8_t src0, uint8_t src1, void stream,
810	float scalarQuant);	1144	float scalarQuant);
	@@ -814,4 +1148,6 @@ template void LaunchTStoreAcc2gmScalarNz2nz<21>(uint8_t out, uint8_t src0, uin
814	template void LaunchTStoreAcc2gmVectorNz2nd<21>(uint8_t out, uint8_t src0, uint8_t src1, uint8_t quantTensor,	1148	template void LaunchTStoreAcc2gmVectorNz2nd<21>(uint8_t out, uint8_t src0, uint8_t src1, uint8_t quantTensor,
815	void *stream);	1149	void *stream);
816	template void LaunchTStoreAcc2gmVectorNz2nz<21>(uint8_t out, uint8_t src0, uint8_t src1, uint8_t quantTensor,	1150	template void LaunchTStoreAcc2gmVectorNz2nz<21>(uint8_t out, uint8_t src0, uint8_t src1, uint8_t quantTensor,
817	- void *stream);	1151	+ void *stream);
		1152	+template void LaunchTStoreAcc2gmVectorNz2NDC1HWC0<21>(uint8_t out, uint8_t src0, uint8_t src1, uint8_t quantTensor,
		1153	+ void *stream);