support ndc1hwc0 of tload in a2-pto-isa-AtomGit

Minclude/pto/npu/a2a3/TLoad.hpp+67-3

	@@ -593,6 +593,60 @@ __tf__ PTO_INTERNAL void TLoadFractalZ(typename TileData::TileDType __out__ dst,
593	}	593	}
594	}	594	}
595		595
		596	+template <typename TileData, typename GlobalData>
		597	+__tf__ PTO_INTERNAL void TLoadNDC1HWC0(typename TileData::TileDType __out__ dst, typename GlobalData::DType __in__ *src,
		598	+ int srcN, int srcD, int srcC1, int srcH, int srcW, int gStride0, int gStride1,
		599	+ int gStride2, int gStride3, int gStride4, int dstN, int dstD, int dstC1,
		600	+ int dstH, int dstW)
		601	+{
		602	+ __cbuf__ typename TileData::DType dstAddr = (__cbuf__ typename TileData::DType )__cce_get_tile_ptr(dst);
		603	+ typename GlobalData::DType *srcAddr = src;
		604	+
		605	+ constexpr uint32_t c0ElemCount = C0_SIZE_BYTE / sizeof(typename TileData::DType);
		606	+ typename GlobalData::DType *srcAddrP = srcAddr;
		607	+ __cbuf__ typename TileData::DType *dstAddrP = dstAddr;
		608	+ constexpr uint32_t maxSupportBurst = 4095;
		609	+ // gmGap unit is 32B
		610	+ uint32_t gmGap = ((gStride2 - dstH * dstW * c0ElemCount) * sizeof(typename TileData::DType)) >> SHIFT_BLOCK_BYTE;
		611	+
		612	+ if ((gStride3 == dstW * c0ElemCount \|\| dstH == 1) && // process for W direction all load or H=1
		613	+ gmGap <= UINT16_MAX && dstC1 <= maxSupportBurst && dstH * dstW <= UINT16_MAX) {
		614	+ uint16_t nBurst = dstC1;
		615	+ uint16_t srcGap = gmGap;
		616	+ uint16_t lenBurst = dstH * dstW;
		617	+ for (uint32_t i = 0; i < dstN; i++) {
		618	+ int64_t srcAddr1 = i * gStride0;
		619	+ int64_t dstAddr1 = i * dstD * dstH * dstW * dstC1 * c0ElemCount;
		620	+ for (uint32_t j = 0; j < dstD; j++) {
		621	+ srcAddrP = srcAddr + srcAddr1 + j * gStride1;
		622	+ dstAddrP = dstAddr + dstAddr1 + j * dstH * dstW * dstC1 * c0ElemCount;
		623	+ TLoadInstrGm2L1<TileData, GlobalData>(dstAddrP, srcAddrP, nBurst, lenBurst, srcGap, 0);
		624	+ }
		625	+ }
		626	+ } else {
		627	+ PTO_ASSERT(dstH <= maxSupportBurst, "Fix: max support dstH is 4095!");
		628	+ PTO_ASSERT(dstW <= UINT16_MAX, "Fix: max support dstW is UINT16_MAX!");
		629	+
		630	+ uint16_t nBurst = dstH;
		631	+ uint16_t lenBurst = dstW;
		632	+ uint16_t srcGap = ((gStride3 - srcW * c0ElemCount) * sizeof(typename TileData::DType)) >> SHIFT_BLOCK_BYTE;
		633	+ uint16_t l1Gap = 0;
		634	+ for (uint32_t i = 0; i < dstN; i++) {
		635	+ int64_t srcAddr1 = i * gStride0;
		636	+ int64_t dstAddr1 = i * dstD * dstH * dstW * dstC1 * c0ElemCount;
		637	+ for (uint32_t j = 0; j < dstD; j++) {
		638	+ int64_t srcAddr2 = j * gStride1;
		639	+ int64_t dstAddr2 = j * dstH * dstW * dstC1 * c0ElemCount;
		640	+ for (uint32_t k = 0; k < dstC1; k++) {
		641	+ srcAddrP = srcAddr + srcAddr1 + srcAddr2 + k * gStride2;
		642	+ dstAddrP = dstAddr + dstAddr1 + dstAddr2 + k * dstH * dstW * c0ElemCount;
		643	+ TLoadInstrGm2L1<TileData, GlobalData>(dstAddrP, srcAddrP, nBurst, lenBurst, srcGap, l1Gap);
		644	+ }
		645	+ }
		646	+ }
		647	+ }
		648	+}
		649	+
596	template <typename TileData, typename GlobalData>	650	template <typename TileData, typename GlobalData>
597	PTO_INTERNAL void CheckConvTileData(TileData &dst, GlobalData &src)	651	PTO_INTERNAL void CheckConvTileData(TileData &dst, GlobalData &src)
598	{	652	{
	@@ -609,8 +663,11 @@ PTO_INTERNAL void CheckConvTileData(TileData &dst, GlobalData &src)
609		663
610	constexpr bool isSameLayout =	664	constexpr bool isSameLayout =
611	(GlobalData::layout == pto::Layout::NC1HWC0 && TileData::layout == pto::Layout::NC1HWC0) \|\|	665	(GlobalData::layout == pto::Layout::NC1HWC0 && TileData::layout == pto::Layout::NC1HWC0) \|\|
612	- (GlobalData::layout == pto::Layout::FRACTAL_Z && TileData::layout == pto::Layout::FRACTAL_Z);	666	+ (GlobalData::layout == pto::Layout::FRACTAL_Z && TileData::layout == pto::Layout::FRACTAL_Z) \|\|
613	- static_assert(isSameLayout == true, "Fix: Src Dst layout must be NC1HWC0 or FRACTAL_Z!");	667	+ (GlobalData::layout == pto::Layout::FRACTAL_Z_3D && TileData::layout == pto::Layout::FRACTAL_Z_3D) \|\|
		668	+ (GlobalData::layout == pto::Layout::NDC1HWC0 && TileData::layout == pto::Layout::NDC1HWC0);
		669	+ static_assert(isSameLayout == true,
		670	+ "Fix: Src Dst layout must be NC1HWC0 or FRACTAL_Z or FRACTAL_Z_3D or NDC1HWC0!");
614	}	671	}
615		672
616	template <typename TileData, typename GlobalData>	673	template <typename TileData, typename GlobalData>
	@@ -622,11 +679,18 @@ PTO_INTERNAL void TLOAD_CONVTILE_IMPL(TileData &dst, GlobalData &src)
622	src.GetShape(3), src.GetStride(0), src.GetStride(1), src.GetStride(2),	679	src.GetShape(3), src.GetStride(0), src.GetStride(1), src.GetStride(2),
623	src.GetStride(3), src.GetStride(4), dst.GetShape(0), dst.GetShape(1),	680	src.GetStride(3), src.GetStride(4), dst.GetShape(0), dst.GetShape(1),
624	dst.GetShape(2), dst.GetShape(3));	681	dst.GetShape(2), dst.GetShape(3));
625	- } else if constexpr (GlobalData::layout == pto::Layout::FRACTAL_Z) { // C1HWNC0, dst dim4 is c0Size	682	+ } else if constexpr (GlobalData::layout == pto::Layout::FRACTAL_Z \|\|
		683	+ GlobalData::layout == pto::Layout::FRACTAL_Z_3D) {
		684	+ // [C1HW,N/16,16,C0] or [C1DHW,N/16,16,C0], dst dim4 is c0Size
626	TLoadFractalZ<TileData, GlobalData>(dst.data(), src.data(), src.GetShape(0), src.GetShape(1), src.GetShape(2),	685	TLoadFractalZ<TileData, GlobalData>(dst.data(), src.data(), src.GetShape(0), src.GetShape(1), src.GetShape(2),
627	src.GetShape(3), src.GetShape(4), src.GetStride(0), src.GetStride(1),	686	src.GetShape(3), src.GetShape(4), src.GetStride(0), src.GetStride(1),
628	src.GetStride(2), src.GetStride(3), src.GetStride(4), dst.GetShape(0),	687	src.GetStride(2), src.GetStride(3), src.GetStride(4), dst.GetShape(0),
629	dst.GetShape(1), dst.GetShape(2), dst.GetShape(3));	688	dst.GetShape(1), dst.GetShape(2), dst.GetShape(3));
		689	+ } else if constexpr (GlobalData::layout == pto::Layout::NDC1HWC0) { // NDC1HWC0, globaltensor is NDC1HW
		690	+ TLoadNDC1HWC0<TileData, GlobalData>(dst.data(), src.data(), src.GetShape(0), src.GetShape(1), src.GetShape(2),
		691	+ src.GetShape(3), src.GetShape(4), src.GetStride(0), src.GetStride(1),
		692	+ src.GetStride(2), src.GetStride(3), src.GetStride(4), dst.GetShape(0),
		693	+ dst.GetShape(1), dst.GetShape(2), dst.GetShape(3), dst.GetShape(4));
630	}	694	}
631	}	695	}
632		696

Mtests/npu/a2a3/src/st/testcase/tload_gm2mat/gen_data.py+22-1

	@@ -35,6 +35,12 @@ def gen_golden_data(case_name, gInfo):
35	gWholeShape2, gWholeShape3, gWholeShape4)).astype(data_type)	35	gWholeShape2, gWholeShape3, gWholeShape4)).astype(data_type)
36	output_arr = np.zeros(shape=(gShape0, gShape1, gShape2, gShape3, gShape4), dtype=data_type)	36	output_arr = np.zeros(shape=(gShape0, gShape1, gShape2, gShape3, gShape4), dtype=data_type)
37	output_arr = input_arr[0:gShape0, 0:gShape1, 0:gShape2, 0:gShape3, 0:gShape4]	37	output_arr = input_arr[0:gShape0, 0:gShape1, 0:gShape2, 0:gShape3, 0:gShape4]
		38	+ elif gInfo.format == "NDC1HWC02NDC1HWC0":
		39	+ c0_size = 32 // np.dtype(data_type).itemsize
		40	+ input_arr = np.random.randint(-5, 5, size=(gWholeShape0, gWholeShape1,
		41	+ gWholeShape2, gWholeShape3, gWholeShape4, c0_size)).astype(data_type)
		42	+ output_arr = np.zeros(shape=(gShape0, gShape1, gShape2, gShape3, gShape4, c0_size), dtype=data_type)
		43	+ output_arr = input_arr[0:gShape0, 0:gShape1, 0:gShape2, 0:gShape3, 0:gShape4, :]
38	elif gInfo.format == "DN":	44	elif gInfo.format == "DN":
39	input_arr = np.random.randint(-5, 5, size=(gWholeShape0, gWholeShape1,	45	input_arr = np.random.randint(-5, 5, size=(gWholeShape0, gWholeShape1,
40	gWholeShape2, gWholeShape4, gWholeShape3)).astype(data_type)	46	gWholeShape2, gWholeShape4, gWholeShape3)).astype(data_type)
	@@ -139,6 +145,13 @@ if __name__ == "__main__":
139	"TLoadGM2L1Test.FZ4D2FZ4D_int8_t_1_125_3_16_32_1_250_5_16_32", # cut C1HW N	145	"TLoadGM2L1Test.FZ4D2FZ4D_int8_t_1_125_3_16_32_1_250_5_16_32", # cut C1HW N
140	"TLoadGM2L1Test.FZ4D2FZ4D_float_1_256_3_16_8_1_4704_7_16_8", # cut C1HW N	146	"TLoadGM2L1Test.FZ4D2FZ4D_float_1_256_3_16_8_1_4704_7_16_8", # cut C1HW N
141		147
		148	+ "TLoadGM2L1Test.NDC1HWC02NDC1HWC0_int8_t_1_2_3_16_128_2_3_4_1024_1024", # cut N H
		149	+ "TLoadGM2L1Test.NDC1HWC02NDC1HWC0_int8_t_1_3_4_128_8_2_3_4_128_128", # cut W
		150	+ "TLoadGM2L1Test.NDC1HWC02NDC1HWC0_int8_t_1_3_4_8_128_2_3_8_8_128", # cut C1
		151	+ "TLoadGM2L1Test.NDC1HWC02NDC1HWC0_bfloat16_1_1_16_10_100_2_1_16_100_100", # cut H
		152	+ "TLoadGM2L1Test.NDC1HWC02NDC1HWC0_bfloat16_1_1_10_16_2_2_2_256_16_100", # cut N D C1 W
		153	+ "TLoadGM2L1Test.NDC1HWC02NDC1HWC0_bfloat16_1_1_1_1_8192_2_8_16_16_8192", # cut N D C1 H
		154	+ "TLoadGM2L1Test.NDC1HWC02NDC1HWC0_float_1_1_1_112_112_2_2_3_224_224", # cut N D C1 H W
142	]	155	]
143		156
144	case_params_list = [	157	case_params_list = [
	@@ -195,6 +208,14 @@ if __name__ == "__main__":
195	GlobalTensorInfo(np.int8, "FZ4D2FZ4D", 1, 125, 3, 16, 32, 1, 250, 5, 16, 32),	208	GlobalTensorInfo(np.int8, "FZ4D2FZ4D", 1, 125, 3, 16, 32, 1, 250, 5, 16, 32),
196	GlobalTensorInfo(np.float32, "FZ4D2FZ4D", 1, 256, 3, 16, 8, 1, 4704, 7, 16, 8),	209	GlobalTensorInfo(np.float32, "FZ4D2FZ4D", 1, 256, 3, 16, 8, 1, 4704, 7, 16, 8),
197		210
		211	+ GlobalTensorInfo(np.int8, "NDC1HWC02NDC1HWC0", 1, 2, 3, 16, 128, 2, 3, 4, 1024, 1024),
		212	+ GlobalTensorInfo(np.int8, "NDC1HWC02NDC1HWC0", 1, 3, 4, 128, 8, 2, 3, 4, 128, 128),
		213	+ GlobalTensorInfo(np.int8, "NDC1HWC02NDC1HWC0", 1, 3, 4, 8, 128, 2, 3, 8, 8, 128),
		214	+ GlobalTensorInfo(np.float16, "NDC1HWC02NDC1HWC0", 1, 1, 16, 10, 100, 2, 1, 16, 100, 100),
		215	+ GlobalTensorInfo(np.float16, "NDC1HWC02NDC1HWC0", 1, 1, 10, 16, 2, 2, 2, 256, 16, 100),
		216	+ GlobalTensorInfo(np.float16, "NDC1HWC02NDC1HWC0", 1, 1, 1, 1, 8192, 2, 8, 16, 16, 8192),
		217	+ GlobalTensorInfo(np.float32, "NDC1HWC02NDC1HWC0", 1, 1, 1, 112, 112, 2, 2, 3, 224, 224),
		218	+
198	]	219	]
199		220
200	for i, case_name in enumerate(case_name_list):	221	for i, case_name in enumerate(case_name_list):
	@@ -203,4 +224,4 @@ if __name__ == "__main__":
203	original_dir = os.getcwd()	224	original_dir = os.getcwd()
204	os.chdir(case_name)	225	os.chdir(case_name)
205	gen_golden_data(case_name, case_params_list[i])	226	gen_golden_data(case_name, case_params_list[i])
206	- os.chdir(original_dir)	227	+ os.chdir(original_dir)

Mtests/npu/a2a3/src/st/testcase/tload_gm2mat/main.cpp+49-5

	@@ -47,16 +47,30 @@ void TestTload()
47	// format = 4: DN2ZN	47	// format = 4: DN2ZN
48	// format = 5: NC1HWC02NC1HWC0	48	// format = 5: NC1HWC02NC1HWC0
49	// format = 6: FRACTALZ2FRACTALZ	49	// format = 6: FRACTALZ2FRACTALZ
50	- size_t srcDataSize = gWholeShape0 * gWholeShape1 * gWholeShape2 * gWholeShape3 * gWholeShape4 * sizeof(DataType);	50	+ // format = 7: FRACTALZ4D2FRACTALZ4D
51	- size_t dstDataSize = gShape0 * gShape1 * gShape2 * gShape3 * gShape4 * sizeof(DataType);	51	+ // format = 8: NDC1HWC02NDC1HWC0
		52	+ size_t srcDataSize = static_cast<size_t>(gWholeShape0) * static_cast<size_t>(gWholeShape1) *
		53	+ static_cast<size_t>(gWholeShape2) * static_cast<size_t>(gWholeShape3) *
		54	+ static_cast<size_t>(gWholeShape4) * sizeof(DataType);
52		55
		56	+ size_t dstDataSize = static_cast<size_t>(gShape0) * static_cast<size_t>(gShape1) * static_cast<size_t>(gShape2) *
		57	+ static_cast<size_t>(gShape3) * static_cast<size_t>(gShape4) * sizeof(DataType);
53	constexpr int c0Size = 32 / sizeof(DataType);	58	constexpr int c0Size = 32 / sizeof(DataType);
54	if (format == 3) {	59	if (format == 3) {
55	int gShape4Align = (gShape4 + c0Size - 1) / c0Size * c0Size;	60	int gShape4Align = (gShape4 + c0Size - 1) / c0Size * c0Size;
56	- dstDataSize = gShape0 * gShape1 * gShape2 * gShape3 * gShape4Align * sizeof(DataType);	61	+ dstDataSize = static_cast<size_t>(gShape0) * static_cast<size_t>(gShape1) * static_cast<size_t>(gShape2) *
		62	+ static_cast<size_t>(gShape3) * static_cast<size_t>(gShape4Align) * sizeof(DataType);
57	} else if (format == 4) {	63	} else if (format == 4) {
58	int gShape3Align = (gShape3 + c0Size - 1) / c0Size * c0Size;	64	int gShape3Align = (gShape3 + c0Size - 1) / c0Size * c0Size;
59	- dstDataSize = gShape0 * gShape1 * gShape2 * gShape3Align * gShape4 * sizeof(DataType);	65	+ dstDataSize = static_cast<size_t>(gShape0) * static_cast<size_t>(gShape1) * static_cast<size_t>(gShape2) *
		66	+ static_cast<size_t>(gShape3Align) * static_cast<size_t>(gShape4) * sizeof(DataType);
		67	+ } else if (format == 8) {
		68	+ srcDataSize = static_cast<size_t>(gWholeShape0) * static_cast<size_t>(gWholeShape1) *
		69	+ static_cast<size_t>(gWholeShape2) * static_cast<size_t>(gWholeShape3) *
		70	+ static_cast<size_t>(gWholeShape4) * static_cast<size_t>(c0Size) * sizeof(DataType);
		71	+ dstDataSize = static_cast<size_t>(gShape0) * static_cast<size_t>(gShape1) * static_cast<size_t>(gShape2) *
		72	+ static_cast<size_t>(gShape3) * static_cast<size_t>(gShape4) * static_cast<size_t>(c0Size) *
		73	+ sizeof(DataType);
60	}	74	}
61		75
62	aclInit(nullptr);	76	aclInit(nullptr);
	@@ -331,4 +345,34 @@ TEST_F(TLoadGM2L1Test, FZ4D2FZ4D_int8_t_1_125_3_16_32_1_250_5_16_32)
331	TEST_F(TLoadGM2L1Test, FZ4D2FZ4D_float_1_256_3_16_8_1_4704_7_16_8)	345	TEST_F(TLoadGM2L1Test, FZ4D2FZ4D_float_1_256_3_16_8_1_4704_7_16_8)
332	{	346	{
333	TestTload<7, float, 1, 256, 3, 16, 8, 1, 4704, 7, 16, 8>();	347	TestTload<7, float, 1, 256, 3, 16, 8, 1, 4704, 7, 16, 8>();
334	-}	348	+}
		349	+
		350	+TEST_F(TLoadGM2L1Test, NDC1HWC02NDC1HWC0_int8_t_1_2_3_16_128_2_3_4_1024_1024)
		351	+{
		352	+ TestTload<8, int8_t, 1, 2, 3, 16, 128, 2, 3, 4, 1024, 1024>();
		353	+}
		354	+TEST_F(TLoadGM2L1Test, NDC1HWC02NDC1HWC0_int8_t_1_3_4_128_8_2_3_4_128_128)
		355	+{
		356	+ TestTload<8, int8_t, 1, 3, 4, 128, 8, 2, 3, 4, 128, 128>();
		357	+}
		358	+TEST_F(TLoadGM2L1Test, NDC1HWC02NDC1HWC0_int8_t_1_3_4_8_128_2_3_8_8_128)
		359	+{
		360	+ TestTload<8, int8_t, 1, 3, 4, 8, 128, 2, 3, 8, 8, 128>();
		361	+}
		362	+
		363	+TEST_F(TLoadGM2L1Test, NDC1HWC02NDC1HWC0_bfloat16_1_1_16_10_100_2_1_16_100_100)
		364	+{
		365	+ TestTload<8, uint16_t, 1, 1, 16, 10, 100, 2, 1, 16, 100, 100>();
		366	+}
		367	+TEST_F(TLoadGM2L1Test, NDC1HWC02NDC1HWC0_bfloat16_1_1_10_16_2_2_2_256_16_100)
		368	+{
		369	+ TestTload<8, uint16_t, 1, 1, 10, 16, 2, 2, 2, 256, 16, 100>();
		370	+}
		371	+TEST_F(TLoadGM2L1Test, NDC1HWC02NDC1HWC0_bfloat16_1_1_1_1_8192_2_8_16_16_8192)
		372	+{
		373	+ TestTload<8, uint16_t, 1, 1, 1, 1, 8192, 2, 8, 16, 16, 8192>();
		374	+}
		375	+TEST_F(TLoadGM2L1Test, NDC1HWC02NDC1HWC0_float_1_1_1_112_112_2_2_3_224_224)
		376	+{
		377	+ TestTload<8, float, 1, 1, 1, 112, 112, 2, 2, 3, 224, 224>();
		378	+}

Mtests/npu/a2a3/src/st/testcase/tload_gm2mat/tload_gm2mat_kernel.cpp+54-1

	@@ -281,6 +281,47 @@ AICORE inline void RunTLoad5HD(__gm__ T __out__ out, __gm__ T __in__ src)
281	tf_copy_cbuf_to_gm<T, OutTileData>(out, outTile.data(), (uint8_t)0, 1, validRow, 0, 0);	281	tf_copy_cbuf_to_gm<T, OutTileData>(out, outTile.data(), (uint8_t)0, 1, validRow, 0, 0);
282	}	282	}
283		283
		284	+template <typename T, int dstN, int dstD, int dstC1, int dstH, int dstW, int gWholeShape0, int gWholeShape1,
		285	+ int gWholeShape2, int gWholeShape3, int gWholeShape4>
		286	+AICORE inline void RunTLoad6HD(__gm__ T __out__ out, __gm__ T __in__ src)
		287	+{
		288	+ constexpr int dstC0 = 32 / sizeof(T);
		289	+ constexpr int gStride[5] = {gWholeShape1 * gWholeShape2 * gWholeShape3 * gWholeShape4 * dstC0,
		290	+ gWholeShape2 * gWholeShape3 * gWholeShape4 * dstC0, gWholeShape3 * gWholeShape4 * dstC0,
		291	+ gWholeShape4 * dstC0, dstC0};
		292	+ constexpr int blockSize = 32 / sizeof(T);
		293	+ // for auto mode, bufferSize is a misleading variable name in convTile, it shouldn't be number of bytes it should be
		294	+ // the number of elements
		295	+ constexpr int bufferSize = dstN * dstD * dstC1 * dstH * dstW * dstC0; // * sizeof(T);
		296	+ constexpr int validRow = dstN * dstD * dstC1 * dstH * dstW;
		297	+ constexpr int validCol = dstC0;
		298	+ constexpr int Rows = dstN * dstD * dstC1 * dstH * dstW;
		299	+ constexpr int Cols = (dstC0 + blockSize - 1) / blockSize * blockSize;
		300	+
		301	+ using ShapeDim5 = pto::Shape<dstN, dstD, dstC1, dstH, dstW>;
		302	+ using StridDim5 = pto::Stride<gStride[0], gStride[1], gStride[2], gStride[3], gStride[4]>;
		303	+ using GlobalDataIn = GlobalTensor<T, ShapeDim5, StridDim5, Layout::NDC1HWC0>;
		304	+ using TileData = ConvTile<TileType::Mat, T, bufferSize, Layout::NDC1HWC0,
		305	+ pto::ConvTileShape<dstN, dstD, dstC1, dstH, dstW, dstC0>>;
		306	+
		307	+ TileData srcTile;
		308	+ static_assert(srcTile.totalDimCount == 6);
		309	+ TASSIGN(srcTile, 0x0);
		310	+
		311	+ GlobalDataIn srcGlobal(src);
		312	+ TLOAD(srcTile, srcGlobal);
		313	+#ifndef __PTO_AUTO__
		314	+ set_flag(PIPE_MTE2, PIPE_MTE3, EVENT_ID0);
		315	+ wait_flag(PIPE_MTE2, PIPE_MTE3, EVENT_ID0);
		316	+#endif
		317	+
		318	+ using OutTileData = Tile<TileType::Mat, T, Rows, Cols, BLayout::RowMajor, validRow, validCol>;
		319	+ OutTileData outTile;
		320	+ TASSIGN(outTile, 0x0);
		321	+ // __cbuf__ typename TileData::DType srcAddr = (__cbuf__ typename TileData::DType )outTile.data();
		322	+ tf_copy_cbuf_to_gm<T, OutTileData>(out, outTile.data(), (uint8_t)0, 1, validRow, 0, 0);
		323	+}
		324	+
284	// C1HWNC0	325	// C1HWNC0
285	template <typename T, int dstC1, int dstH, int dstW, int dstN, int dstC0, int gWholeShape0, int gWholeShape1,	326	template <typename T, int dstC1, int dstH, int dstW, int dstN, int dstC0, int gWholeShape0, int gWholeShape1,
286	int gWholeShape2, int gWholeShape3, int gWholeShape4>	327	int gWholeShape2, int gWholeShape3, int gWholeShape4>
	@@ -390,6 +431,9 @@ __global__ AICORE void TLoadKernel(__gm__ T out, __gm__ T src)
390	} else if constexpr (format == 7) { // format = 7: FRACTALZ4D2FRACTALZ4D	431	} else if constexpr (format == 7) { // format = 7: FRACTALZ4D2FRACTALZ4D
391	RunTLoadFractalZ4D<T, gShape0, gShape1, gShape2, gShape3, gShape4, gWholeShape0, gWholeShape1, gWholeShape2,	432	RunTLoadFractalZ4D<T, gShape0, gShape1, gShape2, gShape3, gShape4, gWholeShape0, gWholeShape1, gWholeShape2,
392	gWholeShape3, gWholeShape4>(out, src);	433	gWholeShape3, gWholeShape4>(out, src);
		434	+ } else if constexpr (format == 8) { // format = 8: NDC1HWC02NDC1HWC0
		435	+ RunTLoad6HD<T, gShape0, gShape1, gShape2, gShape3, gShape4, gWholeShape0, gWholeShape1, gWholeShape2,
		436	+ gWholeShape3, gWholeShape4>(out, src);
393	}	437	}
394	}	438	}
395		439
	@@ -471,4 +515,13 @@ template void LaunchTLoad<7, uint16_t, 1, 49, 7, 16, 16, 1, 980, 32, 16, 16>(uin
471	template void LaunchTLoad<7, uint16_t, 1, 81, 3, 16, 16, 1, 90, 3, 16, 16>(uint16_t out, uint16_t src, void *stream);	515	template void LaunchTLoad<7, uint16_t, 1, 81, 3, 16, 16, 1, 90, 3, 16, 16>(uint16_t out, uint16_t src, void *stream);
472	template void LaunchTLoad<7, int8_t, 1, 63, 3, 16, 32, 1, 63, 9, 16, 32>(int8_t out, int8_t src, void *stream);	516	template void LaunchTLoad<7, int8_t, 1, 63, 3, 16, 32, 1, 63, 9, 16, 32>(int8_t out, int8_t src, void *stream);
473	template void LaunchTLoad<7, int8_t, 1, 125, 3, 16, 32, 1, 250, 5, 16, 32>(int8_t out, int8_t src, void *stream);	517	template void LaunchTLoad<7, int8_t, 1, 125, 3, 16, 32, 1, 250, 5, 16, 32>(int8_t out, int8_t src, void *stream);
474	-template void LaunchTLoad<7, float, 1, 256, 3, 16, 8, 1, 4704, 7, 16, 8>(float out, float src, void *stream);	518	+template void LaunchTLoad<7, float, 1, 256, 3, 16, 8, 1, 4704, 7, 16, 8>(float out, float src, void *stream);
		519	+template void LaunchTLoad<8, int8_t, 1, 2, 3, 16, 128, 2, 3, 4, 1024, 1024>(int8_t out, int8_t src, void *stream);
		520	+template void LaunchTLoad<8, int8_t, 1, 3, 4, 128, 8, 2, 3, 4, 128, 128>(int8_t out, int8_t src, void *stream);
		521	+template void LaunchTLoad<8, int8_t, 1, 3, 4, 8, 128, 2, 3, 8, 8, 128>(int8_t out, int8_t src, void *stream);
		522	+template void LaunchTLoad<8, uint16_t, 1, 1, 16, 10, 100, 2, 1, 16, 100, 100>(uint16_t out, uint16_t src,
		523	+ void *stream);
		524	+template void LaunchTLoad<8, uint16_t, 1, 1, 10, 16, 2, 2, 2, 256, 16, 100>(uint16_t out, uint16_t src, void *stream);
		525	+template void LaunchTLoad<8, uint16_t, 1, 1, 1, 1, 8192, 2, 8, 16, 16, 8192>(uint16_t out, uint16_t src,
		526	+ void *stream);
		527	+template void LaunchTLoad<8, float, 1, 1, 1, 112, 112, 2, 2, 3, 224, 224>(float out, float src, void *stream);