TTrans add NC1HWC0 to NC1C0HW-pto-isa-AtomGit

已合并

TTrans add NC1HWC0 to NC1C0HW #950

小呀小石头创建于 5月19日从denglei0610/pto-isa： master 合入到 cann/pto-isa： master

TTrans add NC1HWC0 to NC1C0HW #950

已合并

小呀小石头创建于 5月19日

从denglei0610/pto-isa master合入到cann/pto-isa master

共 4 个文件变更+421-57

Minclude/pto/npu/a2a3/TTrans.hpp+95-55

	@@ -390,21 +390,30 @@ PTO_INTERNAL void TTransRepeatXOperation(__ubuf__ T dstPtr, __ubuf__ T srcPtr,
390		390
391	///////////////////	391	///////////////////
392		392
393	-template <typename T, unsigned blockSizeElem>	393	+template <typename T, unsigned blockSizeElem, bool reverse = false>
394	PTO_INTERNAL void ConvNCHW2NC1HWC0Unalign(__ubuf__ T dst, __ubuf__ T src, unsigned srcN, unsigned srcC, unsigned srcH,	394	PTO_INTERNAL void ConvNCHW2NC1HWC0Unalign(__ubuf__ T dst, __ubuf__ T src, unsigned srcN, unsigned srcC, unsigned srcH,
395	- unsigned srcW, unsigned dstC0)	395	+ unsigned srcW, unsigned validC0)
396	{	396	{
397	- unsigned srcStride = srcH * srcW;	397	+ unsigned srcStride, dstStride, validCol, validRow, validC1;
398	- unsigned dstStride = dstC0;	398	+ if constexpr (reverse) {
399	- unsigned validCol = srcH * srcW;	399	+ srcStride = validC0;
400	- unsigned validRow = dstC0;	400	+ dstStride = srcH * srcW;
401	- unsigned dstC1 = (srcC + dstC0 - 1) / dstC0;	401	+ validCol = validC0;
402	- unsigned nStride = dstC1 * dstC0 * srcH * srcW;	402	+ validRow = srcH * srcW;
403	- unsigned cStride = dstC0 * srcH * srcW;	403	+ validC1 = srcC;
		404	+ } else {
		405	+ srcStride = srcH * srcW;
		406	+ dstStride = validC0;
		407	+ validCol = srcH * srcW;
		408	+ validRow = validC0;
		409	+ validC1 = (srcC + validC0 - 1) / validC0;
		410	+ }
		411	+ unsigned cStride = validC0 * srcH * srcW;
		412	+ unsigned nStride = validC1 * cStride;
404	constexpr unsigned yTileSizeElem = (sizeof(T) == 1) ? Y_ELEM_B8 : Y_ELEM_OTHER;	413	constexpr unsigned yTileSizeElem = (sizeof(T) == 1) ? Y_ELEM_B8 : Y_ELEM_OTHER;
405	- // N C1 C0 HW -> N C1 HW C0	414	+ // N C1 C0 HW -> N C1 HW C0 or N C1 HW C0 -> N C1 C0 HW
406	for (int n = 0; n < srcN; n++) {	415	for (int n = 0; n < srcN; n++) {
407	- for (int c = 0; c < dstC1; c++) {	416	+ for (int c = 0; c < validC1; c++) {
408	__ubuf__ T srcPtr = src + n nStride + c * cStride;	417	__ubuf__ T srcPtr = src + n nStride + c * cStride;
409	__ubuf__ T dstPtr = dst + n nStride + c * cStride;	418	__ubuf__ T dstPtr = dst + n nStride + c * cStride;
410	// tmpStride should computed in static way	419	// tmpStride should computed in static way
	@@ -415,11 +424,12 @@ PTO_INTERNAL void ConvNCHW2NC1HWC0Unalign(__ubuf__ T dst, __ubuf__ T src, unsi
415	}	424	}
416	}	425	}
417		426
418	-template <typename TileDataDst, typename TileDataSrc, typename TileDataTmp, unsigned blockSizeElem>	427	+template <typename TileDataDst, typename TileDataSrc, typename TileDataTmp, unsigned blockSizeElem,
		428	+ bool reverse = false>
419	__tf__ PTO_INTERNAL void TTransConvNCHW2NC1HWC0(typename TileDataDst::TileDType __out__ dst,	429	__tf__ PTO_INTERNAL void TTransConvNCHW2NC1HWC0(typename TileDataDst::TileDType __out__ dst,
420	typename TileDataSrc::TileDType __in__ src,	430	typename TileDataSrc::TileDType __in__ src,
421	typename TileDataTmp::TileDType __in__ tmp, unsigned srcN,	431	typename TileDataTmp::TileDType __in__ tmp, unsigned srcN,
422	- unsigned srcC, unsigned srcH, unsigned srcW, unsigned dstC0)	432	+ unsigned srcC, unsigned srcH, unsigned srcW, unsigned validC0)
423	{	433	{
424	using Tdst = typename TileDataDst::DType;	434	using Tdst = typename TileDataDst::DType;
425	using Tsrc = typename TileDataSrc::DType;	435	using Tsrc = typename TileDataSrc::DType;
	@@ -428,20 +438,31 @@ __tf__ PTO_INTERNAL void TTransConvNCHW2NC1HWC0(typename TileDataDst::TileDType
428	__ubuf__ Tdst dstPtrOrig = (__ubuf__ Tdst )__cce_get_tile_ptr(dst);	438	__ubuf__ Tdst dstPtrOrig = (__ubuf__ Tdst )__cce_get_tile_ptr(dst);
429	__ubuf__ Tsrc srcPtrOrig = (__ubuf__ Tsrc )__cce_get_tile_ptr(src);	439	__ubuf__ Tsrc srcPtrOrig = (__ubuf__ Tsrc )__cce_get_tile_ptr(src);
430	__ubuf__ Ttmp tmpPtr = (__ubuf__ Ttmp )__cce_get_tile_ptr(tmp);	440	__ubuf__ Ttmp tmpPtr = (__ubuf__ Ttmp )__cce_get_tile_ptr(tmp);
431	- unsigned srcStride = srcH * srcW;	441	+
432	- unsigned dstStride = dstC0;	442	+ unsigned srcStride, dstStride, validCol, validRow, validC1;
		443	+ if constexpr (reverse) {
		444	+ validCol = validC0;
		445	+ validRow = srcH * srcW;
		446	+ srcStride = validC0;
		447	+ dstStride = srcH * srcW;
		448	+ validC1 = srcC;
		449	+ } else {
		450	+ validCol = srcH * srcW;
		451	+ validRow = validC0;
		452	+ srcStride = srcH * srcW;
		453	+ dstStride = validC0;
		454	+ validC1 = (srcC + validC0 - 1) / validC0;
		455	+ }
433	if (((dstStride % blockSizeElem) != 0) \|\| ((srcStride % blockSizeElem) != 0) \|\| srcStride / blockSizeElem > 255) {	456	if (((dstStride % blockSizeElem) != 0) \|\| ((srcStride % blockSizeElem) != 0) \|\| srcStride / blockSizeElem > 255) {
434	- ConvNCHW2NC1HWC0Unalign<Tsrc, blockSizeElem>(dstPtrOrig, srcPtrOrig, srcN, srcC, srcH, srcW, dstC0);	457	+ ConvNCHW2NC1HWC0Unalign<Tsrc, blockSizeElem, reverse>(dstPtrOrig, srcPtrOrig, srcN, srcC, srcH, srcW, validC0);
435	return;	458	return;
436	}	459	}
437	- unsigned validCol = srcH * srcW;	460	+
438	- unsigned validRow = dstC0;	461	+ unsigned cStride = validC0 * srcH * srcW;
439	- unsigned dstC1 = (srcC + dstC0 - 1) / dstC0;	462	+ unsigned nStride = validC1 * cStride;
440	- unsigned nStride = dstC1 * dstC0 * srcH * srcW;
441	- unsigned cStride = dstC0 * srcH * srcW;
442	// N C1 C0 HW -> N C1 HW C0	463	// N C1 C0 HW -> N C1 HW C0
443	for (int n = 0; n < srcN; n++) {	464	for (int n = 0; n < srcN; n++) {
444	- for (int c = 0; c < dstC1; c++) {	465	+ for (int c = 0; c < validC1; c++) {
445	__ubuf__ Tsrc srcPtr = srcPtrOrig + n nStride + c * cStride;	466	__ubuf__ Tsrc srcPtr = srcPtrOrig + n nStride + c * cStride;
446	__ubuf__ Tdst dstPtr = dstPtrOrig + n nStride + c * cStride;	467	__ubuf__ Tdst dstPtr = dstPtrOrig + n nStride + c * cStride;
447	TTransRepeatXOperation<Tsrc, blockSizeElem>(dstPtr, srcPtr, tmpPtr, validRow, validCol, dstStride,	468	TTransRepeatXOperation<Tsrc, blockSizeElem>(dstPtr, srcPtr, tmpPtr, validRow, validCol, dstStride,
	@@ -503,38 +524,29 @@ __tf__ PTO_INTERNAL void TTransConvNC1HWC02C1HWNC0(typename TileDataDst::TileDTy
503	}	524	}
504	}	525	}
505		526
506	-template <typename T, unsigned blockSizeElem>	527	+template <typename T, unsigned blockSizeElem, bool reverse = false>
507	PTO_INTERNAL void ConvGNCHW2GNC1HWC0Unalign(__ubuf__ T dst, __ubuf__ T src, unsigned srcG, unsigned srcN,	528	PTO_INTERNAL void ConvGNCHW2GNC1HWC0Unalign(__ubuf__ T dst, __ubuf__ T src, unsigned srcG, unsigned srcN,
508	- unsigned srcC, unsigned srcH, unsigned srcW, unsigned dstC0)	529	+ unsigned srcC, unsigned srcH, unsigned srcW, unsigned validC0)
509	{	530	{
510	- unsigned srcStride = srcH * srcW;	531	+ unsigned validC1 = reverse ? srcC : ((srcC + validC0 - 1) / validC0);
511	- unsigned dstStride = dstC0;	532	+ unsigned cStride = validC0 * srcH * srcW;
512	- unsigned validCol = srcH * srcW;	533	+ unsigned nStride = validC1 * cStride;
513	- unsigned validRow = dstC0;	534	+ unsigned gStride = srcN * nStride;
514	- unsigned dstC1 = (srcC + dstC0 - 1) / dstC0;
515	- unsigned gStride = srcN * dstC1 * dstC0 * srcH * srcW;
516	- unsigned nStride = dstC1 * dstC0 * srcH * srcW;
517	- unsigned cStride = dstC0 * srcH * srcW;
518	constexpr unsigned yTileSizeElem = (sizeof(T) == 1) ? Y_ELEM_B8 : Y_ELEM_OTHER;	535	constexpr unsigned yTileSizeElem = (sizeof(T) == 1) ? Y_ELEM_B8 : Y_ELEM_OTHER;
519	for (unsigned g = 0; g < srcG; g++) {	536	for (unsigned g = 0; g < srcG; g++) {
520	- for (unsigned n = 0; n < srcN; n++) {	537	+ __ubuf__ T srcPtr = src + g gStride;
521	- for (unsigned c = 0; c < dstC1; c++) {	538	+ __ubuf__ T dstPtr = dst + g gStride;
522	- __ubuf__ T srcPtr = src + g gStride + n * nStride + c * cStride;	539	+ ConvNCHW2NC1HWC0Unalign<T, blockSizeElem, reverse>(dstPtr, srcPtr, srcN, srcC, srcH, srcW, validC0);
523	- __ubuf__ T dstPtr = dst + g gStride + n * nStride + c * cStride;
524	- unsigned tmpStride = (validRow + yTileSizeElem - 1) / yTileSizeElem * yTileSizeElem;
525	- TransTailTiles<T, blockSizeElem, yTileSizeElem>(dstPtr, srcPtr, tmpStride, validRow, validCol,
526	- dstStride, srcStride);
527	- }
528	- }
529	}	540	}
530	}	541	}
531		542
532	-template <typename TileDataDst, typename TileDataSrc, typename TileDataTmp, unsigned blockSizeElem>	543	+template <typename TileDataDst, typename TileDataSrc, typename TileDataTmp, unsigned blockSizeElem,
		544	+ bool reverse = false>
533	__tf__ PTO_INTERNAL void TTransConvGNCHW2GNC1HWC0(typename TileDataDst::TileDType __out__ dst,	545	__tf__ PTO_INTERNAL void TTransConvGNCHW2GNC1HWC0(typename TileDataDst::TileDType __out__ dst,
534	typename TileDataSrc::TileDType __in__ src,	546	typename TileDataSrc::TileDType __in__ src,
535	typename TileDataTmp::TileDType __in__ tmp, unsigned srcG,	547	typename TileDataTmp::TileDType __in__ tmp, unsigned srcG,
536	unsigned srcN, unsigned srcC, unsigned srcH, unsigned srcW,	548	unsigned srcN, unsigned srcC, unsigned srcH, unsigned srcW,
537	- unsigned dstC0)	549	+ unsigned validC0)
538	{	550	{
539	using Tdst = typename TileDataDst::DType;	551	using Tdst = typename TileDataDst::DType;
540	using Tsrc = typename TileDataSrc::DType;	552	using Tsrc = typename TileDataSrc::DType;
	@@ -543,21 +555,30 @@ __tf__ PTO_INTERNAL void TTransConvGNCHW2GNC1HWC0(typename TileDataDst::TileDTyp
543	__ubuf__ Tdst dstPtrOrig = (__ubuf__ Tdst )__cce_get_tile_ptr(dst);	555	__ubuf__ Tdst dstPtrOrig = (__ubuf__ Tdst )__cce_get_tile_ptr(dst);
544	__ubuf__ Tsrc srcPtrOrig = (__ubuf__ Tsrc )__cce_get_tile_ptr(src);	556	__ubuf__ Tsrc srcPtrOrig = (__ubuf__ Tsrc )__cce_get_tile_ptr(src);
545	__ubuf__ Ttmp tmpPtr = (__ubuf__ Ttmp )__cce_get_tile_ptr(tmp);	557	__ubuf__ Ttmp tmpPtr = (__ubuf__ Ttmp )__cce_get_tile_ptr(tmp);
546	- unsigned srcStride = srcH * srcW;	558	+
547	- unsigned dstStride = dstC0;	559	+ unsigned srcStride, dstStride, validCol, validRow, validC1;
		560	+ if constexpr (reverse) {
		561	+ validCol = validC0;
		562	+ validRow = srcH * srcW;
		563	+ validC1 = srcC;
		564	+ } else {
		565	+ validCol = srcH * srcW;
		566	+ validRow = validC0;
		567	+ validC1 = (srcC + validC0 - 1) / validC0;
		568	+ }
		569	+ srcStride = validCol;
		570	+ dstStride = validRow;
548	if (((dstStride % blockSizeElem) != 0) \|\| ((srcStride % blockSizeElem) != 0) \|\| srcStride / blockSizeElem > 255) {	571	if (((dstStride % blockSizeElem) != 0) \|\| ((srcStride % blockSizeElem) != 0) \|\| srcStride / blockSizeElem > 255) {
549	- ConvGNCHW2GNC1HWC0Unalign<Tsrc, blockSizeElem>(dstPtrOrig, srcPtrOrig, srcG, srcN, srcC, srcH, srcW, dstC0);	572	+ ConvGNCHW2GNC1HWC0Unalign<Tsrc, blockSizeElem, reverse>(dstPtrOrig, srcPtrOrig, srcG, srcN, srcC, srcH, srcW,
		573	+ validC0);
550	return;	574	return;
551	}	575	}
552	- unsigned validCol = srcH * srcW;	576	+ unsigned cStride = validC0 * srcH * srcW;
553	- unsigned validRow = dstC0;	577	+ unsigned nStride = validC1 * cStride;
554	- unsigned dstC1 = (srcC + dstC0 - 1) / dstC0;	578	+ unsigned gStride = srcN * nStride;
555	- unsigned gStride = srcN * dstC1 * dstC0 * srcH * srcW;
556	- unsigned nStride = dstC1 * dstC0 * srcH * srcW;
557	- unsigned cStride = dstC0 * srcH * srcW;
558	for (unsigned g = 0; g < srcG; g++) {	579	for (unsigned g = 0; g < srcG; g++) {
559	for (unsigned n = 0; n < srcN; n++) {	580	for (unsigned n = 0; n < srcN; n++) {
560	- for (unsigned c = 0; c < dstC1; c++) {	581	+ for (unsigned c = 0; c < validC1; c++) {
561	__ubuf__ Tsrc srcPtr = srcPtrOrig + g gStride + n * nStride + c * cStride;	582	__ubuf__ Tsrc srcPtr = srcPtrOrig + g gStride + n * nStride + c * cStride;
562	__ubuf__ Tdst dstPtr = dstPtrOrig + g gStride + n * nStride + c * cStride;	583	__ubuf__ Tdst dstPtr = dstPtrOrig + g gStride + n * nStride + c * cStride;
563	TTransRepeatXOperation<Tsrc, blockSizeElem>(dstPtr, srcPtr, tmpPtr, validRow, validCol, dstStride,	584	TTransRepeatXOperation<Tsrc, blockSizeElem>(dstPtr, srcPtr, tmpPtr, validRow, validCol, dstStride,
	@@ -1045,7 +1066,7 @@ PTO_INTERNAL void TTransImplConvTile(TileDataDst &dst, TileDataSrc &src, TileDat
1045	unsigned dstC0 = dst.GetShape(GlobalTensorDim::TOTAL_DIM);	1066	unsigned dstC0 = dst.GetShape(GlobalTensorDim::TOTAL_DIM);
1046	TTransConvGNCHW2GNC1HWC0<TileDataDst, TileDataSrc, TileDataTmp, blockSizeElem>(	1067	TTransConvGNCHW2GNC1HWC0<TileDataDst, TileDataSrc, TileDataTmp, blockSizeElem>(
1047	dst.data(), src.data(), tmp.data(), srcG, srcN, srcC, srcH, srcW, dstC0);	1068	dst.data(), src.data(), tmp.data(), srcG, srcN, srcC, srcH, srcW, dstC0);
1048	- } else if (TileDataSrc::layout == Layout::GNC1HWC0 && TileDataDst::layout == Layout::FRACTAL_Z) {	1069	+ } else if constexpr (TileDataSrc::layout == Layout::GNC1HWC0 && TileDataDst::layout == Layout::FRACTAL_Z) {
1049	CheckGroupConvTile<TileDataDst, TileDataSrc, TileDataTmp>(dst, src, tmp);	1070	CheckGroupConvTile<TileDataDst, TileDataSrc, TileDataTmp>(dst, src, tmp);
1050	unsigned dstN1 = dst.GetShape(GlobalTensorDim::DIM_1);	1071	unsigned dstN1 = dst.GetShape(GlobalTensorDim::DIM_1);
1051	unsigned dstN0 = dst.GetShape(GlobalTensorDim::DIM_2);	1072	unsigned dstN0 = dst.GetShape(GlobalTensorDim::DIM_2);
	@@ -1057,7 +1078,7 @@ PTO_INTERNAL void TTransImplConvTile(TileDataDst &dst, TileDataSrc &src, TileDat
1057	unsigned srcC0 = src.GetShape(GlobalTensorDim::TOTAL_DIM);	1078	unsigned srcC0 = src.GetShape(GlobalTensorDim::TOTAL_DIM);
1058	TTransConvGNC1HWC02GC1HWNC0<TileDataDst, TileDataSrc, TileDataTmp, blockSizeElem>(	1079	TTransConvGNC1HWC02GC1HWNC0<TileDataDst, TileDataSrc, TileDataTmp, blockSizeElem>(
1059	dst.data(), src.data(), tmp.data(), dstN0 * dstN1, srcG, srcN, srcC1 * srcH * srcW, srcC0);	1080	dst.data(), src.data(), tmp.data(), dstN0 * dstN1, srcG, srcN, srcC1 * srcH * srcW, srcC0);
1060	- } else if (TileDataSrc::layout == Layout::NCHW && TileDataDst::layout == Layout::NC1HWC0) {	1081	+ } else if constexpr (TileDataSrc::layout == Layout::NCHW && TileDataDst::layout == Layout::NC1HWC0) {
1061	CheckConvTile<TileDataDst, TileDataSrc, TileDataTmp>(dst, src, tmp);	1082	CheckConvTile<TileDataDst, TileDataSrc, TileDataTmp>(dst, src, tmp);
1062	unsigned srcN = src.GetShape(GlobalTensorDim::DIM_0);	1083	unsigned srcN = src.GetShape(GlobalTensorDim::DIM_0);
1063	unsigned srcC = src.GetShape(GlobalTensorDim::DIM_1);	1084	unsigned srcC = src.GetShape(GlobalTensorDim::DIM_1);
	@@ -1077,6 +1098,25 @@ PTO_INTERNAL void TTransImplConvTile(TileDataDst &dst, TileDataSrc &src, TileDat
1077	unsigned dstC0 = dst.GetShape(GlobalTensorDim::DIM_3);	1098	unsigned dstC0 = dst.GetShape(GlobalTensorDim::DIM_3);
1078	TTransConvNCDHW2FractalZ3D<TileDataDst, TileDataSrc, TileDataTmp, blockSizeElem>(	1099	TTransConvNCDHW2FractalZ3D<TileDataDst, TileDataSrc, TileDataTmp, blockSizeElem>(
1079	dst.data(), src.data(), tmp.data(), srcN, srcC, srcD, srcH, srcW, dstN0, dstC0);	1100	dst.data(), src.data(), tmp.data(), srcN, srcC, srcD, srcH, srcW, dstN0, dstC0);
		1101	+ } else if constexpr (TileDataSrc::layout == Layout::NC1HWC0 && TileDataDst::layout == Layout::NCHW) {
		1102	+ CheckConvTile<TileDataDst, TileDataSrc, TileDataTmp>(dst, src, tmp);
		1103	+ unsigned srcN = src.GetShape(GlobalTensorDim::DIM_0);
		1104	+ unsigned srcC1 = src.GetShape(GlobalTensorDim::DIM_1);
		1105	+ unsigned srcH = src.GetShape(GlobalTensorDim::DIM_2);
		1106	+ unsigned srcW = src.GetShape(GlobalTensorDim::DIM_3);
		1107	+ unsigned srcC0 = src.GetShape(GlobalTensorDim::DIM_4);
		1108	+ TTransConvNCHW2NC1HWC0<TileDataDst, TileDataSrc, TileDataTmp, blockSizeElem, true>(
		1109	+ dst.data(), src.data(), tmp.data(), srcN, srcC1, srcH, srcW, srcC0);
		1110	+ } else if constexpr (TileDataSrc::layout == Layout::GNC1HWC0 && TileDataDst::layout == Layout::GNCHW) {
		1111	+ CheckGroupConvTile<TileDataDst, TileDataSrc, TileDataTmp>(dst, src, tmp);
		1112	+ unsigned srcG = src.GetShape(GlobalTensorDim::DIM_0);
		1113	+ unsigned srcN = src.GetShape(GlobalTensorDim::DIM_1);
		1114	+ unsigned srcC1 = src.GetShape(GlobalTensorDim::DIM_2);
		1115	+ unsigned srcH = src.GetShape(GlobalTensorDim::DIM_3);
		1116	+ unsigned srcW = src.GetShape(GlobalTensorDim::DIM_4);
		1117	+ constexpr unsigned srcC0 = BLOCK_BYTE_SIZE / sizeof(T);
		1118	+ TTransConvGNCHW2GNC1HWC0<TileDataDst, TileDataSrc, TileDataTmp, blockSizeElem, true>(
		1119	+ dst.data(), src.data(), tmp.data(), srcG, srcN, srcC1, srcH, srcW, srcC0);
1080	}	1120	}
1081	}	1121	}
1082		1122

Mtests/npu/a2a3/src/st/testcase/ttrans_conv/gen_data.py+145-0

	@@ -23,6 +23,8 @@ class DataFormat(Enum):
23	NC1HWC02C1HWN1N0C0 = 2	23	NC1HWC02C1HWN1N0C0 = 2
24	GNCHW2GNC1HWC0 = 3	24	GNCHW2GNC1HWC0 = 3
25	GNC1HWC02C1HWN1N0C0 = 4	25	GNC1HWC02C1HWN1N0C0 = 4
		26	+ NC1HWC02NC1C0HW = 5
		27	+ GNC1HWC02GNC1C0HW = 6
26		28
27		29
28	def nchw_to_nc1hwc0(nchw_tensor: np.ndarray, c0: int) -> np.ndarray:	30	def nchw_to_nc1hwc0(nchw_tensor: np.ndarray, c0: int) -> np.ndarray:
	@@ -175,6 +177,43 @@ def _golden_gnc1hwc0_to_c1hwn1n0c0(g_info):
175	return input_arr, output_arr	177	return input_arr, output_arr
176		178
177		179
		180	+def _golden_nc1hwc0_to_nc1c0hw(g_info):
		181	+ """[N, C1, H, W, C0] -> [N, C1, C0, H, W]; pads ??? when needed."""
		182	+ dtype = g_info.data_type
		183	+ src_n = g_info.g_whole_shape0
		184	+ src_c1 = g_info.g_whole_shape1
		185	+ src_h = g_info.g_whole_shape2
		186	+ src_w = g_info.g_whole_shape3
		187	+ src_c0 = g_info.g_whole_shape4
		188	+
		189	+ input_arr = np.random.randint(1, 5, size=(src_n, src_c1, src_h, src_w, src_c0)).astype(dtype)
		190	+
		191	+ # NC1HWC0 → NC1C0HW
		192	+ # origin index：0(n),1(c1),2(h),3(w),4(C0) → new index ：0,1,4,2,3
		193	+ output_arr = np.transpose(input_arr, axes=(0, 1, 4, 2, 3))
		194	+ output_arr = np.reshape(output_arr, (src_n, src_c1 * src_c0, src_h, src_w))
		195	+ return input_arr, output_arr
		196	+
		197	+
		198	+def _golden_gnc1hwc0_to_gnc1c0hw(g_info):
		199	+ """[G, N, C1, H, W, C0] -> [G, N, C1, C0, H, W]; pads ??? when needed."""
		200	+ dtype = g_info.data_type
		201	+ src_g = g_info.g_whole_shape0
		202	+ src_n = g_info.g_whole_shape1
		203	+ src_c1 = g_info.g_whole_shape2
		204	+ src_h = g_info.g_whole_shape3
		205	+ src_w = g_info.g_whole_shape4
		206	+ src_c0 = g_info.g_whole_shape5
		207	+
		208	+ input_arr = np.random.randint(1, 5, size=(src_g, src_n, src_c1, src_h, src_w, src_c0)).astype(dtype)
		209	+
		210	+ # GNC1HWC0 → GNC1C0HW
		211	+ # origin index：0(g),1(n),2(c1),3(h),4(w),5(C0) → new index ：0,1,2,5,3,4
		212	+ output_arr = np.transpose(input_arr, axes=(0, 1, 2, 5, 3, 4))
		213	+ output_arr = np.reshape(output_arr, (src_g, src_n, src_c1 * src_c0, src_h, src_w))
		214	+ return input_arr, output_arr
		215	+
		216	+
178	def gen_golden_data(g_info):	217	def gen_golden_data(g_info):
179	shape1 = g_info.shape	218	shape1 = g_info.shape
180	if shape1 == DataFormat["NCHW2NC1HWC0"].value:	219	if shape1 == DataFormat["NCHW2NC1HWC0"].value:
	@@ -185,6 +224,10 @@ def gen_golden_data(g_info):
185	input_arr, output_arr = _golden_gnchw2_gnc1hwc0(g_info)	224	input_arr, output_arr = _golden_gnchw2_gnc1hwc0(g_info)
186	elif shape1 == DataFormat["GNC1HWC02C1HWN1N0C0"].value:	225	elif shape1 == DataFormat["GNC1HWC02C1HWN1N0C0"].value:
187	input_arr, output_arr = _golden_gnc1hwc0_to_c1hwn1n0c0(g_info)	226	input_arr, output_arr = _golden_gnc1hwc0_to_c1hwn1n0c0(g_info)
		227	+ elif shape1 == DataFormat["NC1HWC02NC1C0HW"].value:
		228	+ input_arr, output_arr = _golden_nc1hwc0_to_nc1c0hw(g_info)
		229	+ elif shape1 == DataFormat["GNC1HWC02GNC1C0HW"].value:
		230	+ input_arr, output_arr = _golden_gnc1hwc0_to_gnc1c0hw(g_info)
188	else:	231	else:
189	data_type = g_info.data_type	232	data_type = g_info.data_type
190	g_shape3 = g_info.g_shape3	233	g_shape3 = g_info.g_shape3
	@@ -711,6 +754,108 @@ if __name__ == "__main__":
711	4,	754	4,
712	4,	755	4,
713	),	756	),
		757	+ TTRANSParams(
		758	+ "TTRANSConvTest.float32_NC1HWC02NC1C0HW_0",
		759	+ np.float32,
		760	+ DataFormat["NC1HWC02NC1C0HW"].value,
		761	+ 1,
		762	+ 1,
		763	+ 2,
		764	+ 4,
		765	+ 8,
		766	+ 1,
		767	+ 1,
		768	+ 1,
		769	+ 2,
		770	+ 4,
		771	+ 8
		772	+ ),
		773	+ TTRANSParams(
		774	+ "TTRANSConvTest.float32_NC1HWC02NC1C0HW_1",
		775	+ np.float32,
		776	+ DataFormat["NC1HWC02NC1C0HW"].value,
		777	+ 2,
		778	+ 2,
		779	+ 2,
		780	+ 2,
		781	+ 4,
		782	+ 8,
		783	+ 2,
		784	+ 2,
		785	+ 2,
		786	+ 4,
		787	+ 8
		788	+ ),
		789	+ TTRANSParams(
		790	+ "TTRANSConvTest.float32_NC1HWC02NC1C0HW_2",
		791	+ np.float32,
		792	+ DataFormat["NC1HWC02NC1C0HW"].value,
		793	+ 2,
		794	+ 2,
		795	+ 2,
		796	+ 3,
		797	+ 4,
		798	+ 8,
		799	+ 2,
		800	+ 2,
		801	+ 3,
		802	+ 4,
		803	+ 8
		804	+ ),
		805	+ TTRANSParams(
		806	+ "TTRANSConvTest.float32_GNC1HWC02GNC1C0HW_0",
		807	+ np.float32,
		808	+ DataFormat["GNC1HWC02GNC1C0HW"].value,
		809	+ 1,
		810	+ 1,
		811	+ 1,
		812	+ 2,
		813	+ 4,
		814	+ 8,
		815	+ 1,
		816	+ 1,
		817	+ 1,
		818	+ 2,
		819	+ 4,
		820	+ 1,
		821	+ 8
		822	+ ),
		823	+ TTRANSParams(
		824	+ "TTRANSConvTest.float32_GNC1HWC02GNC1C0HW_1",
		825	+ np.float32,
		826	+ DataFormat["GNC1HWC02GNC1C0HW"].value,
		827	+ 2,
		828	+ 2,
		829	+ 2,
		830	+ 2,
		831	+ 4,
		832	+ 8,
		833	+ 2,
		834	+ 2,
		835	+ 2,
		836	+ 2,
		837	+ 4,
		838	+ 1,
		839	+ 8
		840	+ ),
		841	+ TTRANSParams(
		842	+ "TTRANSConvTest.float32_GNC1HWC02GNC1C0HW_2",
		843	+ np.float32,
		844	+ DataFormat["GNC1HWC02GNC1C0HW"].value,
		845	+ 2,
		846	+ 2,
		847	+ 2,
		848	+ 3,
		849	+ 4,
		850	+ 8,
		851	+ 2,
		852	+ 2,
		853	+ 2,
		854	+ 3,
		855	+ 4,
		856	+ 1,
		857	+ 8
		858	+ ),
714	]	859	]
715		860
716	for case_params in case_params_list:	861	for case_params in case_params_list:

Mtests/npu/a2a3/src/st/testcase/ttrans_conv/main.cpp+31-1

	@@ -289,4 +289,34 @@ TEST_F(TTRANSConvTest, float16_1_7_2_1_8_16)
289	TEST_F(TTRANSConvTest, float16_4_7_2_1_8_4)	289	TEST_F(TTRANSConvTest, float16_4_7_2_1_8_4)
290	{	290	{
291	test_ttrans_group<aclFloat16, 1, 4, 2, 1, 8, 1, 16, 4, 4, 7, 2, 1, 8, 4>();	291	test_ttrans_group<aclFloat16, 1, 4, 2, 1, 8, 1, 16, 4, 4, 7, 2, 1, 8, 4>();
292	-}	292	+}
		293	+
		294	+TEST_F(TTRANSConvTest, float32_NC1HWC02NC1C0HW_0)
		295	+{
		296	+ test_ttrans<float, 2, 1, 1, 2, 4, 8, 1, 1, 1, 2, 4, 8>();
		297	+}
		298	+
		299	+TEST_F(TTRANSConvTest, float32_NC1HWC02NC1C0HW_1)
		300	+{
		301	+ test_ttrans<float, 2, 2, 2, 2, 4, 8, 1, 2, 2, 2, 4, 8>();
		302	+}
		303	+
		304	+TEST_F(TTRANSConvTest, float32_NC1HWC02NC1C0HW_2)
		305	+{
		306	+ test_ttrans<float, 2, 2, 2, 3, 4, 8, 1, 2, 2, 3, 4, 8>();
		307	+}
		308	+
		309	+TEST_F(TTRANSConvTest, float32_GNC1HWC02GNC1C0HW_0)
		310	+{
		311	+ test_ttrans_group<float, 2, 1, 1, 1, 2, 4, 8, 1, 1, 1, 1, 2, 4, 8>();
		312	+}
		313	+
		314	+TEST_F(TTRANSConvTest, float32_GNC1HWC02GNC1C0HW_1)
		315	+{
		316	+ test_ttrans_group<float, 2, 2, 2, 2, 2, 4, 8, 1, 2, 2, 2, 2, 4, 8>();
		317	+}
		318	+
		319	+TEST_F(TTRANSConvTest, float32_GNC1HWC02GNC1C0HW_2)
		320	+{
		321	+ test_ttrans_group<float, 2, 2, 2, 2, 3, 4, 8, 1, 2, 2, 2, 3, 4, 8>();
		322	+}

Mtests/npu/a2a3/src/st/testcase/ttrans_conv/ttrans_conv_kernel.cpp+150-1

	@@ -160,6 +160,67 @@ __global__ AICORE void runTTRANSConv2(__gm__ T __out__ out, __gm__ T __in__ sr
160	TSTORE(dstGlobal, dst0Tile);	160	TSTORE(dstGlobal, dst0Tile);
161	}	161	}
162		162
		163	+// NC1HWC0 -> NCHW
		164	+template <typename T, int srcN, int srcC1, int srcH, int srcW, int srcC0, int gWholeShape0, int gWholeShape1,
		165	+ int gWholeShape2, int gWholeShape3, int gWholeShape4>
		166	+__global__ AICORE void runTTRANSConv3(__gm__ T __out__ out, __gm__ T __in__ src)
		167	+{
		168	+ constexpr int elemNum = srcN * srcC1 * srcH * srcW * srcC0;
		169	+ constexpr int bufferSize = elemNum * sizeof(T);
		170	+
		171	+ using ShapeDim5 = Shape<1, 1, 1, 1, elemNum>;
		172	+ using StrideDim5 = pto::Stride<elemNum, elemNum, elemNum, elemNum, 1>;
		173	+ using GlobalDataIn = GlobalTensor<T, ShapeDim5, StrideDim5>;
		174	+
		175	+ using SrcTileData = Tile<TileType::Vec, T, 1, elemNum, BLayout::RowMajor, 1, elemNum>;
		176	+ SrcTileData src0Tile;
		177	+ TASSIGN(src0Tile, 0x0);
		178	+ using TileData =
		179	+ ConvTile<TileType::Vec, T, elemNum, Layout::NC1HWC0, ConvTileShape<srcN, srcC1, srcH, srcW, srcC0>>;
		180	+ TileData srcTile;
		181	+ static_assert(srcTile.totalDimCount == 5);
		182	+ TASSIGN(srcTile, 0x0);
		183	+#ifdef __PTO_AUTO__
		184	+ TRESHAPE(src0Tile, srcTile);
		185	+#endif
		186	+
		187	+ using DstTileData =
		188	+ ConvTile<TileType::Vec, T, elemNum, Layout::NCHW, ConvTileShape<srcN, srcC1 * srcC0, srcH, srcW>>;
		189	+ DstTileData dstTile;
		190	+ static_assert(dstTile.totalDimCount == 4);
		191	+ TASSIGN(dstTile, 0x0 + bufferSize);
		192	+ SrcTileData dst0Tile;
		193	+#ifdef __PTO_AUTO__
		194	+ TRESHAPE(dst0Tile, dstTile);
		195	+#endif
		196	+ TASSIGN(dst0Tile, 0x0 + bufferSize);
		197	+
		198	+ constexpr int tmpTileH = srcH * srcW;
		199	+ constexpr unsigned yTileSizeElem = (sizeof(T) == 1) ? 32 : 16;
		200	+ constexpr int tmpTileW = (srcC0 + yTileSizeElem - 1) / yTileSizeElem * yTileSizeElem;
		201	+ using TmpTileData = Tile<TileType::Vec, T, tmpTileH, tmpTileW, BLayout::RowMajor, tmpTileH, tmpTileW>;
		202	+ TmpTileData tmpTile;
		203	+ TASSIGN(tmpTile, 0x0 + bufferSize * 2);
		204	+
		205	+ GlobalDataIn srcGlobal(src);
		206	+ GlobalDataIn dstGlobal(out);
		207	+ TLOAD(src0Tile, srcGlobal);
		208	+#ifndef __PTO_AUTO__
		209	+ set_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);
		210	+ wait_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);
		211	+ set_flag(PIPE_MTE2, PIPE_S, EVENT_ID0);
		212	+ wait_flag(PIPE_MTE2, PIPE_S, EVENT_ID0);
		213	+#endif
		214	+ TTRANS(dstTile, srcTile, tmpTile);
		215	+#ifndef __PTO_AUTO__
		216	+ set_flag(PIPE_S, PIPE_MTE3, EVENT_ID0);
		217	+ wait_flag(PIPE_S, PIPE_MTE3, EVENT_ID0);
		218	+ set_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);
		219	+ wait_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);
		220	+#endif
		221	+ TSTORE(dstGlobal, dst0Tile);
		222	+}
		223	+
163	template <typename T, int format, int gShape0, int gShape1, int gShape2, int gShape3, int gShape4, int gShape5,	224	template <typename T, int format, int gShape0, int gShape1, int gShape2, int gShape3, int gShape4, int gShape5,
164	int gWholeShape0, int gWholeShape1, int gWholeShape2, int gWholeShape3, int gWholeShape4>	225	int gWholeShape0, int gWholeShape1, int gWholeShape2, int gWholeShape3, int gWholeShape4>
165	void LaunchTTRANSConv(T out, T src, void *stream)	226	void LaunchTTRANSConv(T out, T src, void *stream)
	@@ -172,6 +233,9 @@ void LaunchTTRANSConv(T out, T src, void *stream)
172	runTTRANSConv2<half, gShape0, gShape1, gShape2, gShape3, gShape4, gShape5, gWholeShape0, gWholeShape1,	233	runTTRANSConv2<half, gShape0, gShape1, gShape2, gShape3, gShape4, gShape5, gWholeShape0, gWholeShape1,
173	gWholeShape2, gWholeShape3, gWholeShape4>	234	gWholeShape2, gWholeShape3, gWholeShape4>
174	<<<1, nullptr, stream>>>((half )(out), (half )(src));	235	<<<1, nullptr, stream>>>((half )(out), (half )(src));
		236	+ } else if constexpr (format == 2) {
		237	+ runTTRANSConv3<half, gShape0, gShape1, gShape2, gShape3, gShape4, gWholeShape0, gWholeShape1, gWholeShape2,
		238	+ gWholeShape3, gWholeShape4><<<1, nullptr, stream>>>((half )(out), (half )(src));
175	}	239	}
176	} else {	240	} else {
177	if constexpr (format == 0) {	241	if constexpr (format == 0) {
	@@ -180,6 +244,9 @@ void LaunchTTRANSConv(T out, T src, void *stream)
180	} else if constexpr (format == 1) {	244	} else if constexpr (format == 1) {
181	runTTRANSConv2<T, gShape0, gShape1, gShape2, gShape3, gShape4, gShape5, gWholeShape0, gWholeShape1,	245	runTTRANSConv2<T, gShape0, gShape1, gShape2, gShape3, gShape4, gShape5, gWholeShape0, gWholeShape1,
182	gWholeShape2, gWholeShape3, gWholeShape4><<<1, nullptr, stream>>>(out, src);	246	gWholeShape2, gWholeShape3, gWholeShape4><<<1, nullptr, stream>>>(out, src);
		247	+ } else if constexpr (format == 2) {
		248	+ runTTRANSConv3<T, gShape0, gShape1, gShape2, gShape3, gShape4, gWholeShape0, gWholeShape1, gWholeShape2,
		249	+ gWholeShape3, gWholeShape4><<<1, nullptr, stream>>>(out, src);
183	}	250	}
184	}	251	}
185	}	252	}
	@@ -218,6 +285,11 @@ template void LaunchTTRANSConv<int8_t, 1, 5, 1, 6, 2, 16, 32, 25, 5, 1, 6, 32>(i
218	template void LaunchTTRANSConv<uint8_t, 1, 2, 7, 7, 1, 16, 32, 11, 2, 7, 7, 32>(uint8_t out, uint8_t src,	285	template void LaunchTTRANSConv<uint8_t, 1, 2, 7, 7, 1, 16, 32, 11, 2, 7, 7, 32>(uint8_t out, uint8_t src,
219	void *stream);	286	void *stream);
220		287
		288	+// NC1HWC0 -> NCHW
		289	+template void LaunchTTRANSConv<float, 2, 1, 1, 2, 4, 8, 1, 1, 1, 2, 4, 8>(float out, float src, void *stream);
		290	+template void LaunchTTRANSConv<float, 2, 2, 2, 2, 4, 8, 1, 2, 2, 2, 4, 8>(float out, float src, void *stream);
		291	+template void LaunchTTRANSConv<float, 2, 2, 2, 3, 4, 8, 1, 2, 2, 3, 4, 8>(float out, float src, void *stream);
		292	+
221	// GNCHW -> GNC1HWC0	293	// GNCHW -> GNC1HWC0
222	template <typename T, int dstG, int dstN, int dstC1, int dstH, int dstW, int dstC0, int gWholeShape0, int gWholeShape1,	294	template <typename T, int dstG, int dstN, int dstC1, int dstH, int dstW, int dstC0, int gWholeShape0, int gWholeShape1,
223	int gWholeShape2, int gWholeShape3, int gWholeShape4, int gWholeShape5>	295	int gWholeShape2, int gWholeShape3, int gWholeShape4, int gWholeShape5>
	@@ -358,6 +430,67 @@ __global__ AICORE void runTTRANSGroupConv2(__gm__ T __out__ *out, __gm__ T __in_
358	TSTORE(dstGlobal, dst0Tile);	430	TSTORE(dstGlobal, dst0Tile);
359	}	431	}
360		432
		433	+// GNC1HWC0 -> GNCHW
		434	+template <typename T, int dstG, int dstC1, int dstH, int dstW, int dstN1, int dstN0, int dstC0, int srcG, int srcN,
		435	+ int srcC1, int srcH, int srcW, int srcC0>
		436	+__global__ AICORE void runTTRANSGroupConv3(__gm__ T __out__ out, __gm__ T __in__ src)
		437	+{
		438	+ constexpr int elemNum = srcG * srcN * srcC1 * srcH * srcW * srcC0;
		439	+ constexpr int bufferSize = elemNum * sizeof(T);
		440	+
		441	+ using ShapeDim5 = Shape<1, 1, 1, 1, elemNum>;
		442	+ using StrideDim5 = pto::Stride<elemNum, elemNum, elemNum, elemNum, 1>;
		443	+ using GlobalDataIn = GlobalTensor<T, ShapeDim5, StrideDim5>;
		444	+
		445	+ using SrcTileData = Tile<TileType::Vec, T, 1, elemNum, BLayout::RowMajor, 1, elemNum>;
		446	+ SrcTileData src0Tile;
		447	+ TASSIGN(src0Tile, 0x0);
		448	+ using TileData =
		449	+ ConvTile<TileType::Vec, T, elemNum, Layout::GNC1HWC0, ConvTileShape<srcG, srcN, srcC1, srcH, srcW, srcC0>>;
		450	+ TileData srcTile;
		451	+ static_assert(srcTile.totalDimCount == 6);
		452	+ TASSIGN(srcTile, 0x0);
		453	+#ifdef __PTO_AUTO__
		454	+ TRESHAPE(src0Tile, srcTile);
		455	+#endif
		456	+
		457	+ using DstTileData =
		458	+ ConvTile<TileType::Vec, T, elemNum, Layout::GNCHW, ConvTileShape<srcG, srcN, srcC1 * srcC0, srcH, srcW>>;
		459	+ DstTileData dstTile;
		460	+ static_assert(dstTile.totalDimCount == 5);
		461	+ TASSIGN(dstTile, 0x0 + elemNum * sizeof(T));
		462	+ SrcTileData dst0Tile;
		463	+#ifdef __PTO_AUTO__
		464	+ TRESHAPE(dst0Tile, dstTile);
		465	+#endif
		466	+ TASSIGN(dst0Tile, 0x0 + elemNum * sizeof(T));
		467	+
		468	+ constexpr int tmpTileH = srcH * srcW;
		469	+ constexpr unsigned yTileSizeElem = (sizeof(T) == 1) ? 32 : 16;
		470	+ constexpr int tmpTileW = (srcC0 + yTileSizeElem - 1) / yTileSizeElem * yTileSizeElem;
		471	+ using TmpTileData = Tile<TileType::Vec, T, tmpTileH, tmpTileW, BLayout::RowMajor, tmpTileH, tmpTileW>;
		472	+ TmpTileData tmpTile;
		473	+ TASSIGN(tmpTile, 0x0 + elemNum * sizeof(T) * 2);
		474	+
		475	+ GlobalDataIn srcGlobal(src);
		476	+ GlobalDataIn dstGlobal(out);
		477	+ TLOAD(src0Tile, srcGlobal);
		478	+#ifndef __PTO_AUTO__
		479	+ set_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);
		480	+ wait_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);
		481	+ set_flag(PIPE_MTE2, PIPE_S, EVENT_ID0);
		482	+ wait_flag(PIPE_MTE2, PIPE_S, EVENT_ID0);
		483	+#endif
		484	+ TTRANS(dstTile, srcTile, tmpTile);
		485	+#ifndef __PTO_AUTO__
		486	+ set_flag(PIPE_S, PIPE_MTE3, EVENT_ID0);
		487	+ wait_flag(PIPE_S, PIPE_MTE3, EVENT_ID0);
		488	+ set_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);
		489	+ wait_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);
		490	+#endif
		491	+ TSTORE(dstGlobal, dst0Tile);
		492	+}
		493	+
361	template <typename T, int format, int gShape0, int gShape1, int gShape2, int gShape3, int gShape4, int gShape5,	494	template <typename T, int format, int gShape0, int gShape1, int gShape2, int gShape3, int gShape4, int gShape5,
362	int gShape6, int gWholeShape0, int gWholeShape1, int gWholeShape2, int gWholeShape3, int gWholeShape4,	495	int gShape6, int gWholeShape0, int gWholeShape1, int gWholeShape2, int gWholeShape3, int gWholeShape4,
363	int gWholeShape5>	496	int gWholeShape5>
	@@ -372,6 +505,10 @@ void LaunchTTRANSGroupConv(T out, T src, void *stream)
372	runTTRANSGroupConv2<half, gShape0, gShape1, gShape2, gShape3, gShape4, gShape5, gShape6, gWholeShape0,	505	runTTRANSGroupConv2<half, gShape0, gShape1, gShape2, gShape3, gShape4, gShape5, gShape6, gWholeShape0,
373	gWholeShape1, gWholeShape2, gWholeShape3, gWholeShape4, gWholeShape5>	506	gWholeShape1, gWholeShape2, gWholeShape3, gWholeShape4, gWholeShape5>
374	<<<1, nullptr, stream>>>((half )(out), (half )(src));	507	<<<1, nullptr, stream>>>((half )(out), (half )(src));
		508	+ } else if constexpr (format == 2) {
		509	+ runTTRANSGroupConv3<half, gShape0, gShape1, gShape2, gShape3, gShape4, gShape5, gShape6, gWholeShape0,
		510	+ gWholeShape1, gWholeShape2, gWholeShape3, gWholeShape4, gWholeShape5>
		511	+ <<<1, nullptr, stream>>>((half )(out), (half )(src));
375	}	512	}
376	} else {	513	} else {
377	if constexpr (format == 0) {	514	if constexpr (format == 0) {
	@@ -382,6 +519,10 @@ void LaunchTTRANSGroupConv(T out, T src, void *stream)
382	runTTRANSGroupConv2<T, gShape0, gShape1, gShape2, gShape3, gShape4, gShape5, gShape6, gWholeShape0,	519	runTTRANSGroupConv2<T, gShape0, gShape1, gShape2, gShape3, gShape4, gShape5, gShape6, gWholeShape0,
383	gWholeShape1, gWholeShape2, gWholeShape3, gWholeShape4, gWholeShape5>	520	gWholeShape1, gWholeShape2, gWholeShape3, gWholeShape4, gWholeShape5>
384	<<<1, nullptr, stream>>>(out, src);	521	<<<1, nullptr, stream>>>(out, src);
		522	+ } else if constexpr (format == 2) {
		523	+ runTTRANSGroupConv3<T, gShape0, gShape1, gShape2, gShape3, gShape4, gShape5, gShape6, gWholeShape0,
		524	+ gWholeShape1, gWholeShape2, gWholeShape3, gWholeShape4, gWholeShape5>
		525	+ <<<1, nullptr, stream>>>(out, src);
385	}	526	}
386	}	527	}
387	}	528	}
	@@ -411,4 +552,12 @@ template void LaunchTTRANSGroupConv<aclFloat16, 1, 1, 2, 1, 8, 1, 16, 16, 1, 7,
411	void *stream);	552	void *stream);
412	template void LaunchTTRANSGroupConv<aclFloat16, 1, 4, 2, 1, 8, 1, 16, 4, 4, 7, 2, 1, 8, 4>(aclFloat16 *out,	553	template void LaunchTTRANSGroupConv<aclFloat16, 1, 4, 2, 1, 8, 1, 16, 4, 4, 7, 2, 1, 8, 4>(aclFloat16 *out,
413	aclFloat16 *src,	554	aclFloat16 *src,
414	- void *stream);	555	+ void *stream);
		556	+
		557	+// GNC1HWC0 -> GNCHW
		558	+template void LaunchTTRANSGroupConv<float, 2, 1, 1, 1, 2, 4, 8, 1, 1, 1, 1, 2, 4, 8>(float out, float src,
		559	+ void *stream);
		560	+template void LaunchTTRANSGroupConv<float, 2, 2, 2, 2, 2, 4, 8, 1, 2, 2, 2, 2, 4, 8>(float out, float src,
		561	+ void *stream);
		562	+template void LaunchTTRANSGroupConv<float, 2, 2, 2, 2, 3, 4, 8, 1, 2, 2, 2, 3, 4, 8>(float out, float src,
		563	+ void *stream);