aipp新增CSC色域转换-ops-cv-AtomGit

已合并

aipp新增CSC色域转换 #765

jinyijian创建于 4月21日从jinyijian/ops-cv： aipp_csc 合入到 cann/ops-cv： master

aipp新增CSC色域转换 #765

已合并

jinyijian创建于 4月21日

从jinyijian/ops-cv aipp_csc合入到cann/ops-cv master

共 26 个文件变更+2493-394

Mimage/aipp/README.md+1-1

	@@ -171,7 +171,7 @@
171	<td rowspan="3">静态AIPP参数设置</td>	171	<td rowspan="3">静态AIPP参数设置</td>
172	<td>input_format</td>	172	<td>input_format</td>
173	<td>必选属性</td>	173	<td>必选属性</td>
174	- <td>输入图片类型，目前只支持RGB888_U8、XRGB8888_U8、YUV420SP_U8。</td>	174	+ <td>输入图片类型，目前只支持RGB888_U8、XRGB8888_U8、YUV420SP_U8、YUV400_U8。</td>
175	<td>ENUM</td>	175	<td>ENUM</td>
176	</tr>	176	</tr>
177	<tr>	177	<tr>

Mimage/aipp/op_host/arch35/aipp_tiling.cpp+70-49

	@@ -73,33 +73,31 @@ ge::graphStatus AippTiling::ParseNumAndValidateRange(
73		73
74	uint64_t AippTiling::GetTilingKey() const	74	uint64_t AippTiling::GetTilingKey() const
75	{	75	{
76	- uint64_t tilingKey = 0;	76	+ const bool cscSwitch = static_cast<bool>(tilingData.cscParam.cscSwitch);
77	- if ((tilingData.imageFormat == IMAGE_FORMART_MAP.at(IMAGE_FORMAT_RGB888_U8) \|\|	77	+ const bool isRgbFormat = (tilingData.imageFormat == IMAGE_FORMAT_MAP.at(IMAGE_FORMAT_RGB888_U8) \|\|
78	- tilingData.imageFormat == IMAGE_FORMART_MAP.at(IMAGE_FORMAT_XRGB8888_U8)) &&	78	+ tilingData.imageFormat == IMAGE_FORMAT_MAP.at(IMAGE_FORMAT_XRGB8888_U8));
79	- !static_cast<bool>(tilingData.cscParam.cscSwitch)) {	79	+ const bool isYuvFormat = (tilingData.imageFormat == IMAGE_FORMAT_MAP.at(IMAGE_FORMAT_YUV420SP_U8) \|\|
80	- tilingKey = FORMAT_RGB_INDICE_UINT32;	80	+ tilingData.imageFormat == IMAGE_FORMAT_MAP.at(IMAGE_FORMAT_YUV400_U8));
		81	+
		82	+ if (isGray) {
		83	+ if (isRgbFormat) {
		84	+ return AIPP_RGB_TO_GRAY;
		85	+ }
		86	+ if (isYuvFormat) {
		87	+ return AIPP_YUV_TO_GRAY;
		88	+ }
81	}	89	}
82	- if (tilingData.imageFormat == IMAGE_FORMART_MAP.at(IMAGE_FORMAT_YUV420SP_U8) &&	90	+
83	- static_cast<bool>(tilingData.cscParam.cscSwitch) && !isGray) {	91	+ if (isRgbFormat) {
84	- tilingKey = FORMAT_YUV_INDICE_UINT32;	92	+ return cscSwitch ? AIPP_RGB_TO_YUV : AIPP_RGB_PASS_THROUGH;
85	}	93	}
86	- if ((tilingData.imageFormat == IMAGE_FORMART_MAP.at(IMAGE_FORMAT_RGB888_U8) \|\|	94	+
87	- tilingData.imageFormat == IMAGE_FORMART_MAP.at(IMAGE_FORMAT_XRGB8888_U8)) &&	95	+ if (isYuvFormat) {
88	- static_cast<bool>(tilingData.cscParam.cscSwitch) && !isGray) {	96	+ return cscSwitch ? AIPP_YUV_TO_RGB : AIPP_YUV_PASS_THROUGH;
89	- tilingKey = FORMAT_RGB_SWITCH_OPEN_UINT32;
90	}	97	}
91	- if (tilingData.imageFormat == IMAGE_FORMART_MAP.at(IMAGE_FORMAT_RGB888_U8) &&	98	+
92	- static_cast<bool>(tilingData.cscParam.cscSwitch) && isGray) {	99	+ OP_LOGE(context_->GetNodeName(), "tilingKey is:0, please check aipp.cfg's inputformat and csc_switch");
93	- tilingKey = FORMAT_RGB_SWITCH_OPEN_UINT32_TO_GRAY;	100	+ return AIPP_ERROR_TILINGKEY;
94	- }
95	- if (tilingData.imageFormat == IMAGE_FORMART_MAP.at(IMAGE_FORMAT_YUV420SP_U8) &&
96	- static_cast<bool>(tilingData.cscParam.cscSwitch) && isGray) {
97	- tilingKey = FORMAT_YUV_INDICE_UINT32_TO_GRAY;
98	- }
99	- if (tilingKey == 0) {
100	- OP_LOGE(context_->GetNodeName(), "tilingKey is:0, please check aipp.cfg's inputformat and csc_switch");
101	- }
102	- return tilingKey;
103	}	101	}
104		102
105	ge::graphStatus AippTiling::GetShapeAttrsInfo()	103	ge::graphStatus AippTiling::GetShapeAttrsInfo()
	@@ -207,17 +205,11 @@ ge::graphStatus AippTiling::SetImagesValue()
207	{	205	{
208	if (aippCfg.find(AIPP_INPUT_FORMAT) != aippCfg.end()) {	206	if (aippCfg.find(AIPP_INPUT_FORMAT) != aippCfg.end()) {
209	OP_CHECK_IF(	207	OP_CHECK_IF(
210	- aippCfg.at(AIPP_INPUT_FORMAT) != IMAGE_FORMAT_RGB888_U8 &&	208	+ IMAGE_FORMAT_MAP.find(aippCfg.at(AIPP_INPUT_FORMAT)) == IMAGE_FORMAT_MAP.end(),
211	- aippCfg.at(AIPP_INPUT_FORMAT) != IMAGE_FORMAT_YUV420SP_U8 &&	209	+ OP_LOGE(context_->GetNodeName(), "aipp input_format only support RGB888_U8 or XRGB8888_U8 or YUV420SP_U8 or YUV400_U8."),
212	- aippCfg.at(AIPP_INPUT_FORMAT) != IMAGE_FORMAT_XRGB8888_U8,
213	- OP_LOGE(context_->GetNodeName(), "aipp input_format only support RGB888_U8 or XRGB8888_U8 or YUV420SP_U8."),
214	- return ge::GRAPH_FAILED);
215	- OP_CHECK_IF(
216	- IMAGE_FORMART_MAP.find(aippCfg.at(AIPP_INPUT_FORMAT)) == IMAGE_FORMART_MAP.end(),
217	- OP_LOGE(context_->GetNodeName(), "can not find image format in IMAGE_FORMART_MAP."),
218	return ge::GRAPH_FAILED);	210	return ge::GRAPH_FAILED);
219		211
220	- tilingData.imageFormat = IMAGE_FORMART_MAP.at(aippCfg.at(AIPP_INPUT_FORMAT));	212	+ tilingData.imageFormat = IMAGE_FORMAT_MAP.at(aippCfg.at(AIPP_INPUT_FORMAT));
221	} else {	213	} else {
222	OP_LOGE(context_->GetNodeName(), "can not find input_format in aipp config.");	214	OP_LOGE(context_->GetNodeName(), "can not find input_format in aipp config.");
223	return ge::GRAPH_FAILED;	215	return ge::GRAPH_FAILED;
	@@ -248,7 +240,7 @@ ge::graphStatus AippTiling::CheckInputImage()
248	stringstream errorCheckLog;	240	stringstream errorCheckLog;
249	errorCheckLog << "when input_format is " << aippCfg.at(AIPP_INPUT_FORMAT)	241	errorCheckLog << "when input_format is " << aippCfg.at(AIPP_INPUT_FORMAT)
250	<< ", input image size should be bigger than N * src_image_size_w * src_image_size_h * ";	242	<< ", input image size should be bigger than N * src_image_size_w * src_image_size_h * ";
251	- if (tilingData.imageFormat == IMAGE_FORMART_MAP.at(IMAGE_FORMAT_YUV420SP_U8)) {	243	+ if (tilingData.imageFormat == IMAGE_FORMAT_MAP.at(IMAGE_FORMAT_YUV420SP_U8)) {
252	errorCheckLog << IMAGE_FORMAT_YUV420SP_U8_SIZE_LIMIT;	244	errorCheckLog << IMAGE_FORMAT_YUV420SP_U8_SIZE_LIMIT;
253	OP_CHECK_IF((inputImageSize * CONST_VALUE_TWO <	245	OP_CHECK_IF((inputImageSize * CONST_VALUE_TWO <
254	tilingData.batchNum * tilingData.inputSizeW * tilingData.inputSizeH * CONST_VALUE_THREE),	246	tilingData.batchNum * tilingData.inputSizeW * tilingData.inputSizeH * CONST_VALUE_THREE),
	@@ -257,7 +249,12 @@ ge::graphStatus AippTiling::CheckInputImage()
257	OP_CHECK_IF(tilingData.inputSizeW % EVEN_NUMBER_BASE != 0 \|\| tilingData.inputSizeH % EVEN_NUMBER_BASE != 0,	249	OP_CHECK_IF(tilingData.inputSizeW % EVEN_NUMBER_BASE != 0 \|\| tilingData.inputSizeH % EVEN_NUMBER_BASE != 0,
258	OP_LOGE(context_->GetNodeName(), \	250	OP_LOGE(context_->GetNodeName(), \
259	"When input_format is YUV420SP_U8, src_image_size_h/w must be even number"), return ge::GRAPH_FAILED);	251	"When input_format is YUV420SP_U8, src_image_size_h/w must be even number"), return ge::GRAPH_FAILED);
260	- } else if (tilingData.imageFormat == IMAGE_FORMART_MAP.at(IMAGE_FORMAT_RGB888_U8)) {	252	+ } else if (tilingData.imageFormat == IMAGE_FORMAT_MAP.at(IMAGE_FORMAT_YUV400_U8)) {
		253	+ errorCheckLog << static_cast<int>(IMAGE_FORMAT_YUV400_U8_SIZE_LIMIT);
		254	+ OP_CHECK_IF((inputImageSize <
		255	+ tilingData.batchNum * tilingData.inputSizeW * tilingData.inputSizeH * IMAGE_FORMAT_YUV400_U8_SIZE_LIMIT),
		256	+ OP_LOGE(context_->GetNodeName(), "%s", errorCheckLog.str().c_str()), return ge::GRAPH_FAILED);
		257	+ } else if (tilingData.imageFormat == IMAGE_FORMAT_MAP.at(IMAGE_FORMAT_RGB888_U8)) {
261	errorCheckLog << static_cast<int>(IMAGE_FORMAT_RGB888_U8_SIZE_LIMIT);	258	errorCheckLog << static_cast<int>(IMAGE_FORMAT_RGB888_U8_SIZE_LIMIT);
262	OP_CHECK_IF((inputImageSize <	259	OP_CHECK_IF((inputImageSize <
263	tilingData.batchNum * tilingData.inputSizeW * tilingData.inputSizeH * IMAGE_FORMAT_RGB888_U8_SIZE_LIMIT),	260	tilingData.batchNum * tilingData.inputSizeW * tilingData.inputSizeH * IMAGE_FORMAT_RGB888_U8_SIZE_LIMIT),
	@@ -464,7 +461,7 @@ ge::graphStatus AippTiling::GetWorkspaceSize()
464		461
465	ge::graphStatus AippTiling::DoOpTiling()	462	ge::graphStatus AippTiling::DoOpTiling()
466	{	463	{
467	- IsGrayForCSC();	464	+ SetGrayFlag();
468	SwapChannelForCSC();	465	SwapChannelForCSC();
469	PrintTilingData();	466	PrintTilingData();
470	return ge::GRAPH_SUCCESS;	467	return ge::GRAPH_SUCCESS;
	@@ -472,14 +469,14 @@ ge::graphStatus AippTiling::DoOpTiling()
472		469
473	void AippTiling::SwapChannelForCSC()	470	void AippTiling::SwapChannelForCSC()
474	{	471	{
475	- if ((tilingData.imageFormat == IMAGE_FORMART_MAP.at(IMAGE_FORMAT_RGB888_U8) \|\|	472	+ if ((tilingData.imageFormat == IMAGE_FORMAT_MAP.at(IMAGE_FORMAT_RGB888_U8) \|\|
476	- tilingData.imageFormat == IMAGE_FORMART_MAP.at(IMAGE_FORMAT_XRGB8888_U8)) &&	473	+ tilingData.imageFormat == IMAGE_FORMAT_MAP.at(IMAGE_FORMAT_XRGB8888_U8)) &&
477	tilingData.cscParam.rbuvSwapSwitch == 1) {	474	tilingData.cscParam.rbuvSwapSwitch == 1) {
478	swap(tilingData.cscParam.cscMatrix00, tilingData.cscParam.cscMatrix02);	475	swap(tilingData.cscParam.cscMatrix00, tilingData.cscParam.cscMatrix02);
479	swap(tilingData.cscParam.cscMatrix10, tilingData.cscParam.cscMatrix12);	476	swap(tilingData.cscParam.cscMatrix10, tilingData.cscParam.cscMatrix12);
480	swap(tilingData.cscParam.cscMatrix20, tilingData.cscParam.cscMatrix22);	477	swap(tilingData.cscParam.cscMatrix20, tilingData.cscParam.cscMatrix22);
481	}	478	}
482	- if (tilingData.imageFormat == IMAGE_FORMART_MAP.at(IMAGE_FORMAT_YUV420SP_U8) &&	479	+ if (tilingData.imageFormat == IMAGE_FORMAT_MAP.at(IMAGE_FORMAT_YUV420SP_U8) &&
483	tilingData.cscParam.rbuvSwapSwitch == 1) {	480	tilingData.cscParam.rbuvSwapSwitch == 1) {
484	swap(tilingData.cscParam.cscMatrix01, tilingData.cscParam.cscMatrix02);	481	swap(tilingData.cscParam.cscMatrix01, tilingData.cscParam.cscMatrix02);
485	swap(tilingData.cscParam.cscMatrix11, tilingData.cscParam.cscMatrix12);	482	swap(tilingData.cscParam.cscMatrix11, tilingData.cscParam.cscMatrix12);
	@@ -487,8 +484,14 @@ void AippTiling::SwapChannelForCSC()
487	}	484	}
488	}	485	}
489		486
490	-void AippTiling::IsGrayForCSC()	487	+void AippTiling::SetGrayFlag()
491	{	488	{
		489	+ if (tilingData.imageFormat == IMAGE_FORMAT_MAP.at(IMAGE_FORMAT_YUV400_U8) &&
		490	+ !static_cast<bool>(tilingData.cscParam.cscSwitch)) {
		491	+ isGray = true;
		492	+ return;
		493	+ }
		494	+
492	bool anyMatrix1NotZero = (tilingData.cscParam.cscMatrix10 != 0) \|\|	495	bool anyMatrix1NotZero = (tilingData.cscParam.cscMatrix10 != 0) \|\|
493	(tilingData.cscParam.cscMatrix11 != 0) \|\|	496	(tilingData.cscParam.cscMatrix11 != 0) \|\|
494	(tilingData.cscParam.cscMatrix12 != 0);	497	(tilingData.cscParam.cscMatrix12 != 0);
	@@ -498,14 +501,15 @@ void AippTiling::IsGrayForCSC()
498	if (anyMatrix1NotZero \|\| anyMatrix2NotZero) {	501	if (anyMatrix1NotZero \|\| anyMatrix2NotZero) {
499	return;	502	return;
500	}	503	}
501	- if (tilingData.imageFormat == IMAGE_FORMART_MAP.at(IMAGE_FORMAT_RGB888_U8)) {	504	+ if (tilingData.imageFormat == IMAGE_FORMAT_MAP.at(IMAGE_FORMAT_RGB888_U8) \|\|
		505	+ tilingData.imageFormat == IMAGE_FORMAT_MAP.at(IMAGE_FORMAT_XRGB8888_U8)) {
502	if ((tilingData.cscParam.outBias0 == 0) &&	506	if ((tilingData.cscParam.outBias0 == 0) &&
503	(tilingData.cscParam.outBias1 == 0) &&	507	(tilingData.cscParam.outBias1 == 0) &&
504	(tilingData.cscParam.outBias2 == 0)) {	508	(tilingData.cscParam.outBias2 == 0)) {
505	isGray = true;	509	isGray = true;
506	}	510	}
507	}	511	}
508	- if (tilingData.imageFormat == IMAGE_FORMART_MAP.at(IMAGE_FORMAT_YUV420SP_U8) &&	512	+ if (tilingData.imageFormat == IMAGE_FORMAT_MAP.at(IMAGE_FORMAT_YUV420SP_U8) &&
509	tilingData.cscParam.cscMatrix01 == 0 && tilingData.cscParam.cscMatrix02 == 0) {	513	tilingData.cscParam.cscMatrix01 == 0 && tilingData.cscParam.cscMatrix02 == 0) {
510	if ((tilingData.cscParam.inBias0 == 0) &&	514	if ((tilingData.cscParam.inBias0 == 0) &&
511	(tilingData.cscParam.inBias1 == 0) &&	515	(tilingData.cscParam.inBias1 == 0) &&
	@@ -568,7 +572,7 @@ ge::graphStatus AippTiling::CheckCropSize()
568	"aipp.cfg's crop_start_pos_h + crop_size_h should smaller than src_image_size_h."),	572	"aipp.cfg's crop_start_pos_h + crop_size_h should smaller than src_image_size_h."),
569	return ge::GRAPH_FAILED);	573	return ge::GRAPH_FAILED);
570		574
571	- if (tilingData.imageFormat == IMAGE_FORMART_MAP.at(IMAGE_FORMAT_YUV420SP_U8)) {	575	+ if (tilingData.imageFormat == IMAGE_FORMAT_MAP.at(IMAGE_FORMAT_YUV420SP_U8)) {
572	OP_CHECK_IF( tilingData.cropParam.cropStartPosW % EVEN_NUMBER_BASE != 0 \|\|	576	OP_CHECK_IF( tilingData.cropParam.cropStartPosW % EVEN_NUMBER_BASE != 0 \|\|
573	tilingData.cropParam.cropStartPosH % EVEN_NUMBER_BASE != 0, \	577	tilingData.cropParam.cropStartPosH % EVEN_NUMBER_BASE != 0, \
574	OP_LOGE( context_->GetNodeName(), \	578	OP_LOGE( context_->GetNodeName(), \
	@@ -584,6 +588,10 @@ ge::graphStatus AippTiling::SetCscValue()
584	{	588	{
585	if ((aippCfg.find(AIPP_CFG_CSC_SWITCH) != aippCfg.end()) && (aippCfg.at(AIPP_CFG_CSC_SWITCH) == "true")) {	589	if ((aippCfg.find(AIPP_CFG_CSC_SWITCH) != aippCfg.end()) && (aippCfg.at(AIPP_CFG_CSC_SWITCH) == "true")) {
586	tilingData.cscParam.cscSwitch = 1;	590	tilingData.cscParam.cscSwitch = 1;
		591	+ OP_CHECK_IF(
		592	+ tilingData.imageFormat == IMAGE_FORMAT_MAP.at(IMAGE_FORMAT_YUV400_U8), OP_LOGE(context_->GetNodeName(),
		593	+ "When input format is YUV400_U8, it doesn't make sense to convert to RGB, csc_switch : true."),
		594	+ return ge::GRAPH_FAILED);
587	const std::map<string, int16_t*> CSC_MATRIX_MAP = {	595	const std::map<string, int16_t*> CSC_MATRIX_MAP = {
588	{AIPP_CFG_MATRIX_R0C0, &tilingData.cscParam.cscMatrix00},	596	{AIPP_CFG_MATRIX_R0C0, &tilingData.cscParam.cscMatrix00},
589	{AIPP_CFG_MATRIX_R0C1, &tilingData.cscParam.cscMatrix01},	597	{AIPP_CFG_MATRIX_R0C1, &tilingData.cscParam.cscMatrix01},
	@@ -611,13 +619,21 @@ ge::graphStatus AippTiling::SetCscValue()
611	OP_CHECK_IF(ParseNumAndValidateRange(CSC_BIAS_MAP, 0, MAX_RGB_BOUND) != ge::GRAPH_SUCCESS,	619	OP_CHECK_IF(ParseNumAndValidateRange(CSC_BIAS_MAP, 0, MAX_RGB_BOUND) != ge::GRAPH_SUCCESS,
612	OP_LOGE(context_->GetNodeName(), "CSC_BIAS_MAP parse and validate range failed."),	620	OP_LOGE(context_->GetNodeName(), "CSC_BIAS_MAP parse and validate range failed."),
613	return ge::GRAPH_FAILED);	621	return ge::GRAPH_FAILED);
		622	+ } else {
		623	+ tilingData.cscParam.cscMatrix00 = 256; tilingData.cscParam.cscMatrix01 = 0; tilingData.cscParam.cscMatrix02 = 0;
		624	+ tilingData.cscParam.cscMatrix10 = 0; tilingData.cscParam.cscMatrix11 = 256; tilingData.cscParam.cscMatrix12 = 0;
		625	+ tilingData.cscParam.cscMatrix20 = 0; tilingData.cscParam.cscMatrix21 = 0; tilingData.cscParam.cscMatrix22 = 256;
		626	+ tilingData.cscParam.inBias0 = 0; tilingData.cscParam.inBias1 = 0; tilingData.cscParam.inBias2 = 0;
		627	+ tilingData.cscParam.outBias0 = 0; tilingData.cscParam.outBias1 = 0; tilingData.cscParam.outBias2 = 0;
614	}	628	}
615	- OP_CHECK_IF(
616	- SetSwapSwitch() != ge::GRAPH_SUCCESS,
617	- OP_LOGE(context_->GetNodeName(), "Only XRGB888_U8 supports ax_swap_switch being true."),
618	- return ge::GRAPH_FAILED);
619		629
620	- return ge::GRAPH_SUCCESS;	630	+ if (tilingData.imageFormat == IMAGE_FORMAT_MAP.at(IMAGE_FORMAT_YUV420SP_U8)) {
		631	+ tilingData.cscParam.outBias0 = 0; tilingData.cscParam.outBias1 = 0; tilingData.cscParam.outBias2 = 0;
		632	+ } else {
		633	+ tilingData.cscParam.inBias0 = 0; tilingData.cscParam.inBias1 = 0; tilingData.cscParam.inBias2 = 0;
		634	+ }
		635	+
		636	+ return SetSwapSwitch();
621	}	637	}
622		638
623	ge::graphStatus AippTiling::SetSwapSwitch()	639	ge::graphStatus AippTiling::SetSwapSwitch()
	@@ -627,12 +643,17 @@ ge::graphStatus AippTiling::SetSwapSwitch()
627	}	643	}
628		644
629	if (aippCfg.find(AIPP_AX_SWAP_SWITCH) != aippCfg.end() && aippCfg.at(AIPP_AX_SWAP_SWITCH) == "true") {	645	if (aippCfg.find(AIPP_AX_SWAP_SWITCH) != aippCfg.end() && aippCfg.at(AIPP_AX_SWAP_SWITCH) == "true") {
630	- if (tilingData.imageFormat != IMAGE_FORMART_MAP.at(IMAGE_FORMAT_XRGB8888_U8)) {	646	+ if (tilingData.imageFormat != IMAGE_FORMAT_MAP.at(IMAGE_FORMAT_XRGB8888_U8)) {
		647	+ OP_LOGE(context_->GetNodeName(), "Only XRGB888_U8 supports ax_swap_switch being true.");
631	return ge::GRAPH_FAILED;	648	return ge::GRAPH_FAILED;
632	}	649	}
633	tilingData.cscParam.axSwapSwitch = 1;	650	tilingData.cscParam.axSwapSwitch = 1;
634	}	651	}
635		652
		653	+ tilingData.srcChannelOffset =
		654	+ (tilingData.imageFormat == IMAGE_FORMAT_MAP.at(IMAGE_FORMAT_XRGB8888_U8) &&
		655	+ tilingData.cscParam.axSwapSwitch == 1) ? 1 : 0;
		656	+
636	return ge::GRAPH_SUCCESS;	657	return ge::GRAPH_SUCCESS;
637	}	658	}
638		659

Mimage/aipp/op_host/arch35/aipp_tiling.h+14-9

	@@ -60,11 +60,13 @@ constexpr uint8_t CONST_VALUE_FOUR = 4;
60	constexpr uint8_t NCHW_FORMAT_INDEX = 1;	60	constexpr uint8_t NCHW_FORMAT_INDEX = 1;
61	constexpr uint8_t NHWC_FORMAT_INDEX = 2;	61	constexpr uint8_t NHWC_FORMAT_INDEX = 2;
62		62
63	-constexpr int64_t FORMAT_RGB_INDICE_UINT32 = 1;	63	+constexpr int64_t AIPP_ERROR_TILINGKEY = 0;
64	-constexpr int64_t FORMAT_YUV_INDICE_UINT32 = 2;	64	+constexpr int64_t AIPP_RGB_PASS_THROUGH = 1;
65	-constexpr int64_t FORMAT_RGB_SWITCH_OPEN_UINT32 = 3;	65	+constexpr int64_t AIPP_YUV_PASS_THROUGH = 2;
66	-constexpr int64_t FORMAT_RGB_SWITCH_OPEN_UINT32_TO_GRAY = 4;	66	+constexpr int64_t AIPP_RGB_TO_YUV = 3;
67	-constexpr int64_t FORMAT_YUV_INDICE_UINT32_TO_GRAY = 6;	67	+constexpr int64_t AIPP_RGB_TO_GRAY = 4;
		68	+constexpr int64_t AIPP_YUV_TO_RGB = 5;
		69	+constexpr int64_t AIPP_YUV_TO_GRAY = 6;
68		70
69	constexpr size_t INPUT_IMAGES_IDX = 0;	71	constexpr size_t INPUT_IMAGES_IDX = 0;
70	constexpr size_t INPUT_PARAMS_IDX = 1;	72	constexpr size_t INPUT_PARAMS_IDX = 1;
	@@ -86,9 +88,11 @@ constexpr int64_t MAX_PADDING_OUT_W = 1080;
86		88
87	const string IMAGE_FORMAT_RGB888_U8 = "RGB888_U8";	89	const string IMAGE_FORMAT_RGB888_U8 = "RGB888_U8";
88	const string IMAGE_FORMAT_YUV420SP_U8 = "YUV420SP_U8";	90	const string IMAGE_FORMAT_YUV420SP_U8 = "YUV420SP_U8";
		91	+const string IMAGE_FORMAT_YUV400_U8 = "YUV400_U8";
89	const string IMAGE_FORMAT_XRGB8888_U8 = "XRGB8888_U8";	92	const string IMAGE_FORMAT_XRGB8888_U8 = "XRGB8888_U8";
90	constexpr uint8_t IMAGE_FORMAT_RGB888_U8_SIZE_LIMIT = (3);	93	constexpr uint8_t IMAGE_FORMAT_RGB888_U8_SIZE_LIMIT = (3);
91	constexpr float IMAGE_FORMAT_YUV420SP_U8_SIZE_LIMIT = (1.5f);	94	constexpr float IMAGE_FORMAT_YUV420SP_U8_SIZE_LIMIT = (1.5f);
		95	+constexpr uint8_t IMAGE_FORMAT_YUV400_U8_SIZE_LIMIT = (1);
92	constexpr uint8_t IMAGE_FORMAT_XRGB8888_U8_SIZE_LIMIT = (4);	96	constexpr uint8_t IMAGE_FORMAT_XRGB8888_U8_SIZE_LIMIT = (4);
93		97
94	const string AIPP_MODE = "aipp_mode";	98	const string AIPP_MODE = "aipp_mode";
	@@ -141,9 +145,10 @@ const string AIPP_TOP_PADDING_SIZE = "top_padding_size";
141	const string AIPP_BOTTOM_PADDING_SIZE = "bottom_padding_size";	145	const string AIPP_BOTTOM_PADDING_SIZE = "bottom_padding_size";
142	const string AIPP_PADDING_VALUE = "padding_value";	146	const string AIPP_PADDING_VALUE = "padding_value";
143		147
144	-const std::map<string, uint8_t> IMAGE_FORMART_MAP = {{IMAGE_FORMAT_YUV420SP_U8, 1},	148	+const std::map<string, uint8_t> IMAGE_FORMAT_MAP = {{IMAGE_FORMAT_YUV420SP_U8, 1},
		CANN-robot4月21日命名规范与可读性: 映射变量名从 `IMAGE_FORMART_MAP` 更正为 `IMAGE_FORMAT_MAP`，修正了拼写错误（`FORMART` -> `FORMAT`），这是一个积极的改进，提高了代码的准确性。同时，在映射表中新增了 `IMAGE_FORMAT_YUV400_U8` 的条目，确保了数据结构的完整性。问题类型: 命名规范与可读性文件路径: `image/aipp/op_host/arch35/aipp_tiling.h` 行号: 147 问题代码: const std::map IMAGE_FORMAT_MAP = {{IMAGE_FORMAT_YUV420SP_U8, 1}, {IMAGE_FORMAT_RGB888_U8, 2}, {IMAGE_FORMAT_XRGB8888_U8, 3}, {IMAGE_FORMAT_YUV400_U8, 4}}; 修改建议: 无。此项修改是正确且必要的。 --- 此评论由代码审查工具自动生成
145	- {IMAGE_FORMAT_RGB888_U8, 2},	149	+ {IMAGE_FORMAT_RGB888_U8, 2},
146	- {IMAGE_FORMAT_XRGB8888_U8, 3}};	150	+ {IMAGE_FORMAT_XRGB8888_U8, 3},
		151	+ {IMAGE_FORMAT_YUV400_U8, 4}};
147		152
148	struct AippCompileInfo {	153	struct AippCompileInfo {
149	int64_t coreNum = 0;	154	int64_t coreNum = 0;
	@@ -183,7 +188,7 @@ private:
183	ge::graphStatus SetCscValue();	188	ge::graphStatus SetCscValue();
184	ge::graphStatus SetDTCValue();	189	ge::graphStatus SetDTCValue();
185	ge::graphStatus SetSwapSwitch();	190	ge::graphStatus SetSwapSwitch();
186	- void IsGrayForCSC();	191	+ void SetGrayFlag();
187	void SwapChannelForCSC();	192	void SwapChannelForCSC();
188	ge::graphStatus SetPaddingValue();	193	ge::graphStatus SetPaddingValue();
189	ge::graphStatus CheckPaddingSize();	194	ge::graphStatus CheckPaddingSize();

Mimage/aipp/op_kernel/aipp_apt.cpp+16-10

	@@ -15,15 +15,17 @@
15		15
16	#include "arch35/aipp_rgb.h"	16	#include "arch35/aipp_rgb.h"
17	#include "arch35/aipp_yuv.h"	17	#include "arch35/aipp_yuv.h"
		18	+#include "arch35/aipp_yuv_rgb.h"
18	#include "arch35/aipp_rgb_yuv.h"	19	#include "arch35/aipp_rgb_yuv.h"
19	#include "arch35/aipp_yuv_gray.h"	20	#include "arch35/aipp_yuv_gray.h"
20	#include "arch35/aipp_rgb_gray.h"	21	#include "arch35/aipp_rgb_gray.h"
21		22
22	-#define FORMAT_RGB_INDICE_UINT32 1	23	+#define AIPP_RGB_PASS_THROUGH 1
23	-#define FORMAT_YUV_INDICE_UINT32 2	24	+#define AIPP_YUV_PASS_THROUGH 2
24	-#define FORMAT_RGB_2_YUV_INDICE_UINT32 3	25	+#define AIPP_RGB_TO_YUV 3
25	-#define FORMAT_RGB_2_GRAY_INDICE_UINT32 4	26	+#define AIPP_RGB_TO_GRAY 4
26	-#define FORMAT_YUV_2_GRAY_INDICE_UINT32 6	27	+#define AIPP_YUV_TO_RGB 5
		28	+#define AIPP_YUV_TO_GRAY 6
27	using namespace Aipp_Kernel;	29	using namespace Aipp_Kernel;
28		30
29	extern "C" __global__ __aicore__ void Aipp(	31	extern "C" __global__ __aicore__ void Aipp(
	@@ -41,25 +43,29 @@ extern "C" __global__ __aicore__ void Aipp(
41	REGISTER_TILING_DEFAULT(AippTilingData);	43	REGISTER_TILING_DEFAULT(AippTilingData);
42	GET_TILING_DATA(tilingData, tiling);	44	GET_TILING_DATA(tilingData, tiling);
43	KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);	45	KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
44	- if (TILING_KEY_IS(FORMAT_RGB_INDICE_UINT32)) {	46	+ if (TILING_KEY_IS(AIPP_RGB_PASS_THROUGH)) {
45	Aipp_Kernel::AippRgb<DTYPE_FEATURES, uint32_t> op;	47	Aipp_Kernel::AippRgb<DTYPE_FEATURES, uint32_t> op;
46	op.Init(tilingData);	48	op.Init(tilingData);
47	op.Process(images, features);	49	op.Process(images, features);
48	- } else if (TILING_KEY_IS(FORMAT_YUV_INDICE_UINT32)) {	50	+ } else if (TILING_KEY_IS(AIPP_YUV_PASS_THROUGH)) {
49	Aipp_Kernel::AippYuv<DTYPE_FEATURES, uint32_t> op;	51	Aipp_Kernel::AippYuv<DTYPE_FEATURES, uint32_t> op;
50	op.Init(tilingData);	52	op.Init(tilingData);
51	op.Process(images, features);	53	op.Process(images, features);
52	- } else if (TILING_KEY_IS(FORMAT_RGB_2_YUV_INDICE_UINT32)) {	54	+ } else if (TILING_KEY_IS(AIPP_RGB_TO_YUV)) {
53	Aipp_Kernel::AippRgbYuv<DTYPE_FEATURES, uint32_t> op;	55	Aipp_Kernel::AippRgbYuv<DTYPE_FEATURES, uint32_t> op;
54	op.Init(tilingData);	56	op.Init(tilingData);
55	op.Process(images, features);	57	op.Process(images, features);
56	- } else if (TILING_KEY_IS(FORMAT_RGB_2_GRAY_INDICE_UINT32)) {	58	+ } else if (TILING_KEY_IS(AIPP_RGB_TO_GRAY)) {
57	Aipp_Kernel::AippRgbGray<DTYPE_FEATURES, uint32_t> op;	59	Aipp_Kernel::AippRgbGray<DTYPE_FEATURES, uint32_t> op;
58	op.Init(tilingData);	60	op.Init(tilingData);
59	op.Process(images, features);	61	op.Process(images, features);
60	- } else if (TILING_KEY_IS(FORMAT_YUV_2_GRAY_INDICE_UINT32)) {	62	+ } else if (TILING_KEY_IS(AIPP_YUV_TO_GRAY)) {
61	Aipp_Kernel::AippYuvGray<DTYPE_FEATURES, uint32_t> op;	63	Aipp_Kernel::AippYuvGray<DTYPE_FEATURES, uint32_t> op;
62	op.Init(tilingData);	64	op.Init(tilingData);
63	op.Process(images, features);	65	op.Process(images, features);
		66	+ } else if (TILING_KEY_IS(AIPP_YUV_TO_RGB)) {
		67	+ Aipp_Kernel::AippYuvRgb<DTYPE_FEATURES, uint32_t> op;
		68	+ op.Init(tilingData);
		69	+ op.Process(images, features);
64	}	70	}
65	}	71	}

Mimage/aipp/op_kernel/arch35/aipp_base.h+43-2

	@@ -17,6 +17,8 @@
17		17
18	#include "kernel_operator.h"	18	#include "kernel_operator.h"
19	#include "aipp_struct.h"	19	#include "aipp_struct.h"
		20	+#include "simt_api/asc_simt.h"
		21	+#include "simt_api/math_functions.h"
20		22
21	#define CLIP3(x, min, max) ((x) < (min) ? (min) : ((x) > (max) ? (max) : (x)))	23	#define CLIP3(x, min, max) ((x) < (min) ? (min) : ((x) > (max) ? (max) : (x)))
22		24
	@@ -39,6 +41,9 @@ constexpr uint8_t DIGIT_3 = 3;
39	constexpr uint8_t NCHW_FORMAT_INDEX = 1;	41	constexpr uint8_t NCHW_FORMAT_INDEX = 1;
40	constexpr uint8_t NHWC_FORMAT_INDEX = 2;	42	constexpr uint8_t NHWC_FORMAT_INDEX = 2;
41	constexpr uint8_t XRGB8888_U8_FORMAT = 3;	43	constexpr uint8_t XRGB8888_U8_FORMAT = 3;
		44	+constexpr uint8_t YUV420SP_U8_FORMAT = 1;
		45	+constexpr uint8_t RGB888_U8_FORMAT = 2;
		46	+constexpr uint8_t YUV400_U8_FORMAT = 4;
42		47
43	template <typename T, typename DataType>	48	template <typename T, typename DataType>
44	class AippBase {	49	class AippBase {
	@@ -59,8 +64,8 @@ __aicore__ inline void AippBase<T, DataType>::BaseInit(const AippTilingData& til
59	{	64	{
60	tilingData_ = tilingData;	65	tilingData_ = tilingData;
61		66
62	- blockNum_ = GetBlockNum();	67	+ blockNum_ = gridDim.x;
63	- blockIdx_ = GetBlockIdx();	68	+ blockIdx_ = blockIdx.x;
64	totalNum_ = tilingData_.batchNum * tilingData_.outputSizeH * tilingData_.outputSizeW;	69	totalNum_ = tilingData_.batchNum * tilingData_.outputSizeH * tilingData_.outputSizeW;
65	}	70	}
66		71
	@@ -143,6 +148,42 @@ __aicore__ __attribute__((always_inline)) inline bool IsPixelInPadding(
143	(wIdx >= leftPaddingSize + cropSizeW);	148	(wIdx >= leftPaddingSize + cropSizeW);
144	}	149	}
145		150
		151	+__aicore__ __attribute__((always_inline)) inline bool IsPixelInPaddingForYuv(
		152	+ uint32_t pixelH, uint32_t pixelW, const AippTilingData& tD,
		153	+ bool allEvenPadding, bool blockAllInPadding)
		154	+{
		155	+ if (tD.paddingParam.paddingSwitch == 0) {
		156	+ return false;
		157	+ }
		158	+ if (allEvenPadding) {
		159	+ return blockAllInPadding;
		160	+ }
		161	+ return (pixelH < (uint32_t)tD.paddingParam.topPaddingSize) \|\|
		162	+ (pixelH >= (uint32_t)tD.paddingParam.topPaddingSize + tD.cropParam.cropSizeH) \|\|
		163	+ (pixelW < (uint32_t)tD.paddingParam.leftPaddingSize) \|\|
		164	+ (pixelW >= (uint32_t)tD.paddingParam.leftPaddingSize + tD.cropParam.cropSizeW);
		165	+}
		166	+
		167	+__aicore__ __attribute__((always_inline)) inline void ApplyCscMatrix(
		168	+ RgbPack<uint8_t>& dst, uint8_t ch0, uint8_t ch1, uint8_t ch2, const CscParam& cscParam)
		169	+{
		170	+ auto t0 = static_cast<int16_t>(ch0) - cscParam.inBias0;
		171	+ auto t1 = static_cast<int16_t>(ch1) - cscParam.inBias1;
		172	+ auto t2 = static_cast<int16_t>(ch2) - cscParam.inBias2;
		173	+ auto r = static_cast<int16_t>(roundf(static_cast<float>(
		174	+ cscParam.cscMatrix00 * t0 * 2 + cscParam.cscMatrix01 * t1 * 2 +
		175	+ cscParam.cscMatrix02 * t2 * 2 + 1) / CSC_MATRIX_SCALE));
		176	+ auto g = static_cast<int16_t>(roundf(static_cast<float>(
		177	+ cscParam.cscMatrix10 * t0 * 2 + cscParam.cscMatrix11 * t1 * 2 +
		178	+ cscParam.cscMatrix12 * t2 * 2 + 1) / CSC_MATRIX_SCALE));
		179	+ auto b = static_cast<int16_t>(roundf(static_cast<float>(
		180	+ cscParam.cscMatrix20 * t0 * 2 + cscParam.cscMatrix21 * t1 * 2 +
		181	+ cscParam.cscMatrix22 * t2 * 2 + 1) / CSC_MATRIX_SCALE));
		182	+ dst.r = static_cast<uint8_t>(CLIP3(r + cscParam.outBias0, 0, MAX_UINT8));
		183	+ dst.g = static_cast<uint8_t>(CLIP3(g + cscParam.outBias1, 0, MAX_UINT8));
		184	+ dst.b = static_cast<uint8_t>(CLIP3(b + cscParam.outBias2, 0, MAX_UINT8));
		185	+}
		186	+
146	template <typename DataType>	187	template <typename DataType>
147	__aicore__ __attribute__((always_inline)) inline void ComputeCoordFromIndex(	188	__aicore__ __attribute__((always_inline)) inline void ComputeCoordFromIndex(
148	DataType idx, uint32_t outputSizeH, uint32_t outputSizeW,	189	DataType idx, uint32_t outputSizeH, uint32_t outputSizeW,

Mimage/aipp/op_kernel/arch35/aipp_rgb.h+9-15

	@@ -47,10 +47,8 @@ __simt_vf__ LAUNCH_BOUND(MAX_THREAD_NUM) __aicore__ void SimtComputeRgb(
47	uint32_t outputSizeW = tD.outputSizeW;	47	uint32_t outputSizeW = tD.outputSizeW;
48	float padValue = tD.paddingParam.padValue;	48	float padValue = tD.paddingParam.padValue;
49		49
50	- DataType xrgbChannelOffset =	50	+ for (DataType idx = threadIdx.x + blockIdx * blockDim.x; idx < batchSize;
51	- (tD.imageFormat == XRGB8888_U8_FORMAT && tD.cscParam.axSwapSwitch) ? 1 : 0;	51	+ idx += blockNum * blockDim.x) {
52	- for (DataType idx = Simt::GetThreadIdx() + blockIdx * Simt::GetThreadNum(); idx < batchSize;
53	- idx += blockNum * Simt::GetThreadNum()) {
54		52
55	CoordPack<DataType> coord;	53	CoordPack<DataType> coord;
56	ComputeCoordFromIndex(idx, outputSizeH, outputSizeW, coord);	54	ComputeCoordFromIndex(idx, outputSizeH, outputSizeW, coord);
	@@ -65,17 +63,13 @@ __simt_vf__ LAUNCH_BOUND(MAX_THREAD_NUM) __aicore__ void SimtComputeRgb(
65	AssignPadValue(outputGM[dstRgbIdx.b], padValue);	63	AssignPadValue(outputGM[dstRgbIdx.b], padValue);
66	} else {	64	} else {
67	RgbPack<DataType> srcRgbIdx;	65	RgbPack<DataType> srcRgbIdx;
68	- RgbComputeSrcIdx(srcRgbIdx, coord, tD, xrgbChannelOffset);	66	+ RgbComputeSrcIdx(srcRgbIdx, coord, tD, (DataType)tD.srcChannelOffset);
69		67
70	- if (tD.cscParam.rbuvSwapSwitch) {	68	+ RgbPack<uint8_t> result;
71	- DataConversion(outputGM[dstRgbIdx.b], rgbGM[srcRgbIdx.r], tD.dtcParam, CHANNEL_NUM_2);	69	+ ApplyCscMatrix(result, rgbGM[srcRgbIdx.r], rgbGM[srcRgbIdx.g], rgbGM[srcRgbIdx.b], tD.cscParam);
72	- DataConversion(outputGM[dstRgbIdx.g], rgbGM[srcRgbIdx.g], tD.dtcParam, CHANNEL_NUM_1);	70	+ DataConversion(outputGM[dstRgbIdx.r], result.r, tD.dtcParam, CHANNEL_NUM_0);
73	- DataConversion(outputGM[dstRgbIdx.r], rgbGM[srcRgbIdx.b], tD.dtcParam, CHANNEL_NUM_0);	71	+ DataConversion(outputGM[dstRgbIdx.g], result.g, tD.dtcParam, CHANNEL_NUM_1);
74	- } else {	72	+ DataConversion(outputGM[dstRgbIdx.b], result.b, tD.dtcParam, CHANNEL_NUM_2);
75	- DataConversion(outputGM[dstRgbIdx.r], rgbGM[srcRgbIdx.r], tD.dtcParam, CHANNEL_NUM_0);
76	- DataConversion(outputGM[dstRgbIdx.g], rgbGM[srcRgbIdx.g], tD.dtcParam, CHANNEL_NUM_1);
77	- DataConversion(outputGM[dstRgbIdx.b], rgbGM[srcRgbIdx.b], tD.dtcParam, CHANNEL_NUM_2);
78	- }
79	}	73	}
80	}	74	}
81	}	75	}
	@@ -83,7 +77,7 @@ __simt_vf__ LAUNCH_BOUND(MAX_THREAD_NUM) __aicore__ void SimtComputeRgb(
83	template <typename T, typename DataType>	77	template <typename T, typename DataType>
84	__aicore__ inline void AippRgb<T, DataType>::Process(GM_ADDR x, GM_ADDR y)	78	__aicore__ inline void AippRgb<T, DataType>::Process(GM_ADDR x, GM_ADDR y)
85	{	79	{
86	- Simt::VF_CALL<Aipp_Kernel::SimtComputeRgb<T, DataType>>(Simt::Dim3(this->blockDimX_),	80	+ asc_vf_call<Aipp_Kernel::SimtComputeRgb<T, DataType>>(dim3(this->blockDimX_),
87	(__gm__ uint8_t)x, (__gm__ T)y, this->tilingData_, this->blockIdx_, this->blockNum_, this->totalNum_);	81	(__gm__ uint8_t)x, (__gm__ T)y, this->tilingData_, this->blockIdx_, this->blockNum_, this->totalNum_);
88	}	82	}
89		83

Mimage/aipp/op_kernel/arch35/aipp_rgb_gray.h+7-16

	@@ -39,16 +39,6 @@ __aicore__ inline void AippRgbGray<T, DataType>::Init(const AippTilingData& tili
39	this->BaseInit(tilingData);	39	this->BaseInit(tilingData);
40	}	40	}
41		41
42	-__aicore__ __attribute__((always_inline)) inline uint8_t Rgb2Gray(uint8_t r, uint8_t g, uint8_t b,
43	- const CscParam& cscParam)
44	-{
45	- auto gTmp = static_cast<int16_t>(roundf(static_cast<float>(
46	- cscParam.cscMatrix00 * static_cast<int16_t>(r) * DIGIT_2 +
47	- cscParam.cscMatrix01 * static_cast<int16_t>(g) * DIGIT_2 +
48	- cscParam.cscMatrix02 * static_cast<int16_t>(b) * DIGIT_2 + 1) / CSC_MATRIX_SCALE));
49	- return CLIP3(gTmp, 0, MAX_UINT8);
50	-}
51	-
52	template <typename T, typename DataType>	42	template <typename T, typename DataType>
53	__simt_vf__ LAUNCH_BOUND(MAX_THREAD_NUM) __aicore__ void SimtComputeRgb2Gray(	43	__simt_vf__ LAUNCH_BOUND(MAX_THREAD_NUM) __aicore__ void SimtComputeRgb2Gray(
54	__gm__ uint8_t* rgbGM, __gm__ T* grayGM, const AippTilingData tD,	44	__gm__ uint8_t* rgbGM, __gm__ T* grayGM, const AippTilingData tD,
	@@ -58,8 +48,8 @@ __simt_vf__ LAUNCH_BOUND(MAX_THREAD_NUM) __aicore__ void SimtComputeRgb2Gray(
58	uint32_t outputSizeH = tD.outputSizeH;	48	uint32_t outputSizeH = tD.outputSizeH;
59	uint32_t outputSizeW = tD.outputSizeW;	49	uint32_t outputSizeW = tD.outputSizeW;
60		50
61	- for (DataType idx = Simt::GetThreadIdx() + blockIdx * Simt::GetThreadNum(); idx < batchSize;	51	+ for (DataType idx = threadIdx.x + blockIdx * blockDim.x; idx < batchSize;
62	- idx += blockNum * Simt::GetThreadNum()) {	52	+ idx += blockNum * blockDim.x) {
63	CoordPack<DataType> coord;	53	CoordPack<DataType> coord;
64	ComputeCoordFromIndex(idx, outputSizeH, outputSizeW, coord);	54	ComputeCoordFromIndex(idx, outputSizeH, outputSizeW, coord);
65		55
	@@ -74,10 +64,11 @@ __simt_vf__ LAUNCH_BOUND(MAX_THREAD_NUM) __aicore__ void SimtComputeRgb2Gray(
74	AssignPadValue(grayGM[dstGrayIdx.b], padValue);	64	AssignPadValue(grayGM[dstGrayIdx.b], padValue);
75	} else {	65	} else {
76	RgbPack<DataType> srcRgbIdx;	66	RgbPack<DataType> srcRgbIdx;
77	- RgbComputeSrcIdx(srcRgbIdx, coord, tD);	67	+ RgbComputeSrcIdx(srcRgbIdx, coord, tD, (DataType)tD.srcChannelOffset);
78		68
79	- uint8_t grayValue = Rgb2Gray(rgbGM[srcRgbIdx.r], rgbGM[srcRgbIdx.g], rgbGM[srcRgbIdx.b], tD.cscParam);	69	+ RgbPack<uint8_t> result;
80	- DataConversion(grayGM[dstGrayIdx.r], grayValue, tD.dtcParam, CHANNEL_NUM_0);	70	+ ApplyCscMatrix(result, rgbGM[srcRgbIdx.r], rgbGM[srcRgbIdx.g], rgbGM[srcRgbIdx.b], tD.cscParam);
		71	+ DataConversion(grayGM[dstGrayIdx.r], result.r, tD.dtcParam, CHANNEL_NUM_0);
81	DataConversion(grayGM[dstGrayIdx.g], 0, tD.dtcParam, CHANNEL_NUM_1);	72	DataConversion(grayGM[dstGrayIdx.g], 0, tD.dtcParam, CHANNEL_NUM_1);
82	DataConversion(grayGM[dstGrayIdx.b], 0, tD.dtcParam, CHANNEL_NUM_2);	73	DataConversion(grayGM[dstGrayIdx.b], 0, tD.dtcParam, CHANNEL_NUM_2);
83	}	74	}
	@@ -87,7 +78,7 @@ __simt_vf__ LAUNCH_BOUND(MAX_THREAD_NUM) __aicore__ void SimtComputeRgb2Gray(
87	template <typename T, typename DataType>	78	template <typename T, typename DataType>
88	__aicore__ inline void AippRgbGray<T, DataType>::Process(GM_ADDR x, GM_ADDR y)	79	__aicore__ inline void AippRgbGray<T, DataType>::Process(GM_ADDR x, GM_ADDR y)
89	{	80	{
90	- Simt::VF_CALL<Aipp_Kernel::SimtComputeRgb2Gray<T, DataType>>(Simt::Dim3(this->blockDimX_),	81	+ asc_vf_call<Aipp_Kernel::SimtComputeRgb2Gray<T, DataType>>(dim3(this->blockDimX_),
91	(__gm__ uint8_t)x, (__gm__ T)y, this->tilingData_, this->blockIdx_, this->blockNum_, this->totalNum_);	82	(__gm__ uint8_t)x, (__gm__ T)y, this->tilingData_, this->blockIdx_, this->blockNum_, this->totalNum_);
92	}	83	}
93	} // namespace Aipp_Kernel	84	} // namespace Aipp_Kernel

Mimage/aipp/op_kernel/arch35/aipp_rgb_yuv.h+9-24

	@@ -39,21 +39,6 @@ __aicore__ inline void AippRgbYuv<T, DataType>::Init(const AippTilingData& tilin
39	this->BaseInit(tilingData);	39	this->BaseInit(tilingData);
40	}	40	}
41		41
42	-__aicore__ __attribute__((always_inline)) inline void Rgb2Yuv(
43	- YuvPack<uint8_t>& dstYuv, uint8_t r, uint8_t g, uint8_t b, const CscParam& cscParam)
44	-{
45	- auto yTmp = static_cast<int16_t>(roundf(static_cast<float>(cscParam.cscMatrix00 * r * DIGIT_2 +
46	- cscParam.cscMatrix01 * g * DIGIT_2 + cscParam.cscMatrix02 * b * DIGIT_2 + 1) / CSC_MATRIX_SCALE));
47	- auto uTmp = static_cast<int16_t>(roundf(static_cast<float>(cscParam.cscMatrix10 * r * DIGIT_2 +
48	- cscParam.cscMatrix11 * g * DIGIT_2 + cscParam.cscMatrix12 * b * DIGIT_2 + 1) / CSC_MATRIX_SCALE));
49	- auto vTmp = static_cast<int16_t>(roundf(static_cast<float>(cscParam.cscMatrix20 * r * DIGIT_2 +
50	- cscParam.cscMatrix21 * g * DIGIT_2 + cscParam.cscMatrix22 * b * DIGIT_2 + 1) / CSC_MATRIX_SCALE));
51	-
52	- dstYuv.y = CLIP3(yTmp + cscParam.outBias0, 0, MAX_UINT8);
53	- dstYuv.u = CLIP3(uTmp + cscParam.outBias1, 0, MAX_UINT8);
54	- dstYuv.v = CLIP3(vTmp + cscParam.outBias2, 0, MAX_UINT8);
55	-}
56	-
57	template <typename T, typename DataType>	42	template <typename T, typename DataType>
58	__simt_vf__ LAUNCH_BOUND(MAX_THREAD_NUM) __aicore__ void SimtComputeRgb2Yuv(	43	__simt_vf__ LAUNCH_BOUND(MAX_THREAD_NUM) __aicore__ void SimtComputeRgb2Yuv(
59	__gm__ uint8_t* rgbGM, __gm__ T* yuvGM, const AippTilingData tD,	44	__gm__ uint8_t* rgbGM, __gm__ T* yuvGM, const AippTilingData tD,
	@@ -63,8 +48,8 @@ __simt_vf__ LAUNCH_BOUND(MAX_THREAD_NUM) __aicore__ void SimtComputeRgb2Yuv(
63	uint32_t outputSizeW = tD.outputSizeW;	48	uint32_t outputSizeW = tD.outputSizeW;
64	float padValue = tD.paddingParam.padValue;	49	float padValue = tD.paddingParam.padValue;
65		50
66	- for (DataType idx = Simt::GetThreadIdx() + blockIdx * Simt::GetThreadNum(); idx < batchSize;	51	+ for (DataType idx = threadIdx.x + blockIdx * blockDim.x; idx < batchSize;
67	- idx += blockNum * Simt::GetThreadNum()) {	52	+ idx += blockNum * blockDim.x) {
68	CoordPack<DataType> coord;	53	CoordPack<DataType> coord;
69	ComputeCoordFromIndex(idx, outputSizeH, outputSizeW, coord);	54	ComputeCoordFromIndex(idx, outputSizeH, outputSizeW, coord);
70		55
	@@ -78,13 +63,13 @@ __simt_vf__ LAUNCH_BOUND(MAX_THREAD_NUM) __aicore__ void SimtComputeRgb2Yuv(
78	AssignPadValue(yuvGM[dstYuvIdx.b], padValue);	63	AssignPadValue(yuvGM[dstYuvIdx.b], padValue);
79	} else {	64	} else {
80	RgbPack<DataType> srcRgbIdx;	65	RgbPack<DataType> srcRgbIdx;
81	- RgbComputeSrcIdx(srcRgbIdx, coord, tD);	66	+ RgbComputeSrcIdx(srcRgbIdx, coord, tD, (DataType)tD.srcChannelOffset);
82		67
83	- YuvPack<uint8_t> dstYuv;	68	+ RgbPack<uint8_t> dstYuv;
84	- Rgb2Yuv(dstYuv, rgbGM[srcRgbIdx.r], rgbGM[srcRgbIdx.g], rgbGM[srcRgbIdx.b], tD.cscParam);	69	+ ApplyCscMatrix(dstYuv, rgbGM[srcRgbIdx.r], rgbGM[srcRgbIdx.g], rgbGM[srcRgbIdx.b], tD.cscParam);
85	- DataConversion(yuvGM[dstYuvIdx.r], dstYuv.y, tD.dtcParam, CHANNEL_NUM_0);	70	+ DataConversion(yuvGM[dstYuvIdx.r], dstYuv.r, tD.dtcParam, CHANNEL_NUM_0);
86	- DataConversion(yuvGM[dstYuvIdx.g], dstYuv.u, tD.dtcParam, CHANNEL_NUM_1);	71	+ DataConversion(yuvGM[dstYuvIdx.g], dstYuv.g, tD.dtcParam, CHANNEL_NUM_1);
87	- DataConversion(yuvGM[dstYuvIdx.b], dstYuv.v, tD.dtcParam, CHANNEL_NUM_2);	72	+ DataConversion(yuvGM[dstYuvIdx.b], dstYuv.b, tD.dtcParam, CHANNEL_NUM_2);
88	}	73	}
89	}	74	}
90	}	75	}
	@@ -92,7 +77,7 @@ __simt_vf__ LAUNCH_BOUND(MAX_THREAD_NUM) __aicore__ void SimtComputeRgb2Yuv(
92	template <typename T, typename DataType>	77	template <typename T, typename DataType>
93	__aicore__ inline void AippRgbYuv<T, DataType>::Process(GM_ADDR x, GM_ADDR y)	78	__aicore__ inline void AippRgbYuv<T, DataType>::Process(GM_ADDR x, GM_ADDR y)
94	{	79	{
95	- Simt::VF_CALL<Aipp_Kernel::SimtComputeRgb2Yuv<T, DataType>>(Simt::Dim3(this->blockDimX_),	80	+ asc_vf_call<Aipp_Kernel::SimtComputeRgb2Yuv<T, DataType>>(dim3(this->blockDimX_),
96	(__gm__ uint8_t)x, (__gm__ T)y, this->tilingData_, this->blockIdx_, this->blockNum_, this->totalNum_);	81	(__gm__ uint8_t)x, (__gm__ T)y, this->tilingData_, this->blockIdx_, this->blockNum_, this->totalNum_);
97	}	82	}
98	} // namespace Aipp_Kernel	83	} // namespace Aipp_Kernel

Mimage/aipp/op_kernel/arch35/aipp_struct.h+2-0

	@@ -84,6 +84,8 @@ struct AippTilingData {
84	uint32_t outputSizeW = 0;	84	uint32_t outputSizeW = 0;
85	uint32_t outputSizeH = 0;	85	uint32_t outputSizeH = 0;
86		86
		87	+ uint8_t srcChannelOffset = 0;
		88	+
87	CscParam cscParam;	89	CscParam cscParam;
88	CropParam cropParam;	90	CropParam cropParam;
89	DtcParam dtcParam;	91	DtcParam dtcParam;

Mimage/aipp/op_kernel/arch35/aipp_yuv.h+150-142

	@@ -1,23 +1,33 @@
1	/**	1	/**
2	- * Copyright (c) 2025 Huawei Technologies Co., Ltd.	2	+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3	- * This program is free software, you can redistribute it and/or modify it under the terms and conditions of	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4	- * CANN Open Software License Agreement Version 2.0 (the "License").	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
5	- * Please refer to the License for details. You may not use this file except in compliance with the License.	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6	- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7	- * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8	- * See LICENSE in the root of the software repository for the full text of the License.	8	+ * See LICENSE in the root of the software repository for the full text of the License.
9	- */	9	+ */
10		10
11	/*!	11	/*!
12	* \file aipp_yuv.h	12	* \file aipp_yuv.h
13	- * \brief aipp yuv	13	+ * \brief AIPP YUV420SP to YUV444/YVU444 upsampling kernel.
		14	+ *
		15	+ * TilingKey = FORMAT_YUV_INDICE_UINT32 (2), csc_switch = false.
		16	+ *
		17	+ * Algorithm (nearest-neighbour UV upsampling):
		18	+ * Y444(x, y) = Y420(x, y)
		19	+ * U444(x, y) = U420(x/2, y/2)
		20	+ * V444(x, y) = V420(x/2, y/2)
		21	+ *
		22	+ * YVU444: rbuv_swap_switch=true → tiling side swaps CSC matrix columns 1&2,
		23	+ * kernel is format-agnostic.
		24	+ * Thread model: one thread per 2x2 block (4 pixels share one UV pair).
14	*/	25	*/
15		26
16	-#ifndef AIPP_OP_KERNEL_ARCH35_YUV_RGB_H	27	+#ifndef AIPP_OP_KERNEL_ARCH35_YUV_H
17	-#define AIPP_OP_KERNEL_ARCH35_YUV_RGB_H	28	+#define AIPP_OP_KERNEL_ARCH35_YUV_H
18		29
19	#include "aipp_base.h"	30	#include "aipp_base.h"
20	-#include "simt_api/math_functions.h"
21		31
22	namespace Aipp_Kernel {	32	namespace Aipp_Kernel {
23	using namespace AscendC;	33	using namespace AscendC;
	@@ -40,180 +50,178 @@ __aicore__ inline void AippYuv<T, DataType>::Init(const AippTilingData& tilingDa
40	}	50	}
41		51
42	template <typename DataType>	52	template <typename DataType>
43	-__aicore__ __attribute__((always_inline)) inline void ComputeDstIdx(	53	+__aicore__ __attribute__((always_inline)) inline void ComputeDstYuv444Idx(
44	- RgbPack<DataType>* dstRgbIdx, const CoordPack<DataType>& coord, const AippTilingData& tD)	54	+ DataType dstYIdx[YUV_PER_DEAL_NUM],
		55	+ DataType& dstUBase,
		56	+ DataType& dstVBase,
		57	+ const CoordPack<DataType>& coord,
		58	+ const AippTilingData& tD)
45	{	59	{
		60	+ const DataType outputSizeH = tD.outputSizeH;
		61	+ const DataType outputSizeW = tD.outputSizeW;
		62	+ const DataType h0 = coord.hIdx * DIGIT_2;
		63	+ const DataType h1 = h0 + 1;
		64	+ const DataType w0 = coord.wIdx * DIGIT_2;
		65	+ const DataType w1 = w0 + 1;
		66	+
46	if (tD.outputFormat == NCHW_FORMAT_INDEX) {	67	if (tD.outputFormat == NCHW_FORMAT_INDEX) {
47	- dstRgbIdx[YUV_DEAL_NUM_0].r = coord.nIdx * (tD.outputSizeH * tD.outputSizeW * CHANNEL_THREE) +	68	+ const DataType yPlaneBase = coord.nIdx * CHANNEL_THREE * outputSizeH * outputSizeW;
48	- DIGIT_2 * coord.hIdx * tD.outputSizeW + DIGIT_2 * coord.wIdx;	69	+ dstYIdx[YUV_DEAL_NUM_0] = yPlaneBase + h0 * outputSizeW + w0;
49	- dstRgbIdx[YUV_DEAL_NUM_0].g = dstRgbIdx[YUV_DEAL_NUM_0].r + tD.outputSizeH * tD.outputSizeW;	70	+ dstYIdx[YUV_DEAL_NUM_1] = yPlaneBase + h0 * outputSizeW + w1;
50	- dstRgbIdx[YUV_DEAL_NUM_0].b = dstRgbIdx[YUV_DEAL_NUM_0].g + tD.outputSizeH * tD.outputSizeW;	71	+ dstYIdx[YUV_DEAL_NUM_2] = yPlaneBase + h1 * outputSizeW + w0;
51	-	72	+ dstYIdx[YUV_DEAL_NUM_3] = yPlaneBase + h1 * outputSizeW + w1;
52	- dstRgbIdx[YUV_DEAL_NUM_1].r = dstRgbIdx[YUV_DEAL_NUM_0].r + 1;	73	+ dstUBase = yPlaneBase + outputSizeH * outputSizeW + h0 * outputSizeW + w0;
53	- dstRgbIdx[YUV_DEAL_NUM_1].g = dstRgbIdx[YUV_DEAL_NUM_0].g + 1;	74	+ dstVBase = yPlaneBase + DIGIT_2 * outputSizeH * outputSizeW + h0 * outputSizeW + w0;
54	- dstRgbIdx[YUV_DEAL_NUM_1].b = dstRgbIdx[YUV_DEAL_NUM_0].b + 1;
55	-
56	- dstRgbIdx[YUV_DEAL_NUM_2].r = coord.nIdx * (tD.outputSizeH * tD.outputSizeW * CHANNEL_THREE) +
57	- (DIGIT_2 * coord.hIdx + 1) * tD.outputSizeW + DIGIT_2 * coord.wIdx;
58	- dstRgbIdx[YUV_DEAL_NUM_2].g = dstRgbIdx[YUV_DEAL_NUM_2].r + tD.outputSizeH * tD.outputSizeW;
59	- dstRgbIdx[YUV_DEAL_NUM_2].b = dstRgbIdx[YUV_DEAL_NUM_2].g + tD.outputSizeH * tD.outputSizeW;
60	-
61	- dstRgbIdx[YUV_DEAL_NUM_3].r = dstRgbIdx[YUV_DEAL_NUM_2].r + 1;
62	- dstRgbIdx[YUV_DEAL_NUM_3].g = dstRgbIdx[YUV_DEAL_NUM_2].g + 1;
63	- dstRgbIdx[YUV_DEAL_NUM_3].b = dstRgbIdx[YUV_DEAL_NUM_2].b + 1;
64	} else {	75	} else {
65	- dstRgbIdx[YUV_DEAL_NUM_0].r = coord.nIdx * (tD.outputSizeH * tD.outputSizeW * CHANNEL_THREE) +	76	+ // NHWC: each pixel triplet is [Y, ch1, ch2]
66	- DIGIT_2 * coord.hIdx * tD.outputSizeW * CHANNEL_THREE +	77	+ const DataType pixelBase = coord.nIdx * outputSizeH * outputSizeW * CHANNEL_THREE;
67	- DIGIT_2 * coord.wIdx * CHANNEL_THREE;	78	+ dstYIdx[YUV_DEAL_NUM_0] = pixelBase + (h0 * outputSizeW + w0) * CHANNEL_THREE;
68	- dstRgbIdx[YUV_DEAL_NUM_0].g = dstRgbIdx[YUV_DEAL_NUM_0].r + 1;	79	+ dstYIdx[YUV_DEAL_NUM_1] = pixelBase + (h0 * outputSizeW + w1) * CHANNEL_THREE;
69	- dstRgbIdx[YUV_DEAL_NUM_0].b = dstRgbIdx[YUV_DEAL_NUM_0].g + 1;	80	+ dstYIdx[YUV_DEAL_NUM_2] = pixelBase + (h1 * outputSizeW + w0) * CHANNEL_THREE;
70	-	81	+ dstYIdx[YUV_DEAL_NUM_3] = pixelBase + (h1 * outputSizeW + w1) * CHANNEL_THREE;
71	- dstRgbIdx[YUV_DEAL_NUM_1].r = dstRgbIdx[YUV_DEAL_NUM_0].r + CHANNEL_THREE;	82	+ dstUBase = dstYIdx[YUV_DEAL_NUM_0] + 1;
72	- dstRgbIdx[YUV_DEAL_NUM_1].g = dstRgbIdx[YUV_DEAL_NUM_0].g + CHANNEL_THREE;	83	+ dstVBase = dstYIdx[YUV_DEAL_NUM_0] + DIGIT_2;
73	- dstRgbIdx[YUV_DEAL_NUM_1].b = dstRgbIdx[YUV_DEAL_NUM_0].b + CHANNEL_THREE;
74	-
75	- dstRgbIdx[YUV_DEAL_NUM_2].r = coord.nIdx * (tD.outputSizeH * tD.outputSizeW * CHANNEL_THREE) +
76	- (DIGIT_2 * coord.hIdx + 1) * tD.outputSizeW * CHANNEL_THREE +
77	- DIGIT_2 * coord.wIdx * CHANNEL_THREE;
78	- dstRgbIdx[YUV_DEAL_NUM_2].g = dstRgbIdx[YUV_DEAL_NUM_2].r + 1;
79	- dstRgbIdx[YUV_DEAL_NUM_2].b = dstRgbIdx[YUV_DEAL_NUM_2].g + 1;
80	-
81	- dstRgbIdx[YUV_DEAL_NUM_3].r = dstRgbIdx[YUV_DEAL_NUM_2].r + CHANNEL_THREE;
82	- dstRgbIdx[YUV_DEAL_NUM_3].g = dstRgbIdx[YUV_DEAL_NUM_2].g + CHANNEL_THREE;
83	- dstRgbIdx[YUV_DEAL_NUM_3].b = dstRgbIdx[YUV_DEAL_NUM_2].b + CHANNEL_THREE;
84	}	84	}
85	}	85	}
86		86
87	-template <typename T>
88	-__aicore__ __attribute__((always_inline)) inline void Yuv2Rgb(
89	- RgbPack<T>& dstRgb, uint8_t y, uint8_t u, uint8_t v, const CscParam& cscParam)
90	-{
91	- auto rTmp = static_cast<int16_t>(y) - cscParam.inBias0;
92	- auto gTmp = static_cast<int16_t>(u) - cscParam.inBias1;
93	- auto bTmp = static_cast<int16_t>(v) - cscParam.inBias2;
94	-
95	- dstRgb.r = CLIP3(roundf(static_cast<float>(cscParam.cscMatrix00 * rTmp * 2 +
96	- cscParam.cscMatrix01 * gTmp * 2 + cscParam.cscMatrix02 * bTmp * 2 + 1) / CSC_MATRIX_SCALE), 0, 255);
97	- dstRgb.g = CLIP3(roundf(static_cast<float>(cscParam.cscMatrix10 * rTmp * 2 +
98	- cscParam.cscMatrix11 * gTmp * 2 + cscParam.cscMatrix12 * bTmp * 2 + 1) / CSC_MATRIX_SCALE), 0, 255);
99	- dstRgb.b = CLIP3(roundf(static_cast<float>(cscParam.cscMatrix20 * rTmp * 2 +
100	- cscParam.cscMatrix21 * gTmp * 2 + cscParam.cscMatrix22 * bTmp * 2 + 1) / CSC_MATRIX_SCALE), 0, 255);
101	-}
102	-
103	-__aicore__ __attribute__((always_inline)) inline bool IsPixelInPaddingForYuv(
104	- uint32_t pixelH, uint32_t pixelW, const AippTilingData& tD, bool allEvenPadding, bool blockAllInPadding)
105	-{
106	- if (tD.paddingParam.paddingSwitch == 0) {
107	- return false;
108	- }
109	- if (allEvenPadding) {
110	- return blockAllInPadding;
111	- }
112	- return (pixelH < tD.paddingParam.topPaddingSize) \|\|
113	- (pixelH >= tD.paddingParam.topPaddingSize + tD.cropParam.cropSizeH) \|\|
114	- (pixelW < tD.paddingParam.leftPaddingSize) \|\|
115	- (pixelW >= tD.paddingParam.leftPaddingSize + tD.cropParam.cropSizeW);
116	-}
117	-
118	template <typename T, typename DataType>	87	template <typename T, typename DataType>
119	-__aicore__ __attribute__((always_inline)) inline void ProcessYuvPixel(	88	+__aicore__ __attribute__((always_inline)) inline void ProcessYuv444Pixel(
120	- __gm__ uint8_t* yuvGM, __gm__ T* outputGM, const RgbPack<DataType> dstIdx,	89	+ __gm__ uint8_t* inputGM, __gm__ T* outputGM,
121	- uint32_t nIdx, uint32_t croodH, uint32_t croodW, const AippTilingData& tD)	90	+ DataType dstYIdx, DataType dstUIdx, DataType dstVIdx,
		91	+ uint32_t nIdx, uint32_t croodH, uint32_t croodW,
		92	+ const AippTilingData& tD)
122	{	93	{
123	- uint32_t srcYIdx = nIdx * tD.inputSizeH * tD.inputSizeW * 3 / 2 +	94	+ const uint32_t yuvPlaneSize = tD.inputSizeH * tD.inputSizeW * 3 / 2;
124	- (tD.cropParam.cropStartPosH + croodH) * tD.inputSizeW +	95	+
125	- (tD.cropParam.cropStartPosW + croodW);	96	+ uint32_t srcYIdx = nIdx * yuvPlaneSize +
126	- uint32_t srcUIdx = nIdx * tD.inputSizeH * tD.inputSizeW * 3 / 2 + tD.inputSizeH * tD.inputSizeW +	97	+ (tD.cropParam.cropStartPosH + croodH) * tD.inputSizeW +
127	- ((tD.cropParam.cropStartPosH + (croodH & ~1)) >> 1) * tD.inputSizeW +	98	+ (tD.cropParam.cropStartPosW + croodW);
128	- (tD.cropParam.cropStartPosW + (croodW & ~1));	99	+ uint32_t srcUIdx = nIdx * yuvPlaneSize + tD.inputSizeH * tD.inputSizeW +
		100	+ ((tD.cropParam.cropStartPosH + (croodH & ~1u)) >> 1) * tD.inputSizeW +
		101	+ (tD.cropParam.cropStartPosW + (croodW & ~1u));
129	uint32_t srcVIdx = srcUIdx + 1;	102	uint32_t srcVIdx = srcUIdx + 1;
130	- RgbPack<uint8_t> dstRgb;
131	- Yuv2Rgb(dstRgb, yuvGM[srcYIdx], yuvGM[srcUIdx], yuvGM[srcVIdx], tD.cscParam);
132		103
133	- DataConversion(outputGM[dstIdx.r], dstRgb.r, tD.dtcParam, CHANNEL_NUM_0);	104	+ RgbPack<uint8_t> dstYuv;
134	- DataConversion(outputGM[dstIdx.g], dstRgb.g, tD.dtcParam, CHANNEL_NUM_1);	105	+ ApplyCscMatrix(dstYuv, inputGM[srcYIdx], inputGM[srcUIdx], inputGM[srcVIdx], tD.cscParam);
135	- DataConversion(outputGM[dstIdx.b], dstRgb.b, tD.dtcParam, CHANNEL_NUM_2);	106	+
		107	+ DataConversion(outputGM[dstYIdx], dstYuv.r, tD.dtcParam, CHANNEL_NUM_0);
		108	+ DataConversion(outputGM[dstUIdx], dstYuv.g, tD.dtcParam, CHANNEL_NUM_1);
		109	+ DataConversion(outputGM[dstVIdx], dstYuv.b, tD.dtcParam, CHANNEL_NUM_2);
136	}	110	}
137		111
138	template <typename T, typename DataType>	112	template <typename T, typename DataType>
139	-__aicore__ __attribute__((always_inline)) inline void ProcessYuvBlock(	113	+__aicore__ __attribute__((always_inline)) inline void ProcessYuv444Block(
140	- __gm__ uint8_t* yuvGM, __gm__ T* outputGM,	114	+ __gm__ uint8_t* inputGM, __gm__ T* outputGM,
141	- const RgbPack<DataType>* dstRgbIdx, const CoordPack<DataType>& coord,	115	+ const DataType dstYIdx[YUV_PER_DEAL_NUM],
142	- const AippTilingData& tD, float padValue, bool allEvenPadding, bool blockAllInPadding)	116	+ DataType dstUBase, DataType dstVBase,
		117	+ const CoordPack<DataType>& coord,
		118	+ const AippTilingData& tD, float padValue,
		119	+ bool allEvenPadding, bool blockAllInPadding)
143	{	120	{
144	- uint32_t actualH = coord.hIdx * DIGIT_2;	121	+ const uint32_t actualH = coord.hIdx * DIGIT_2;
145	- uint32_t actualW = coord.wIdx * DIGIT_2;	122	+ const uint32_t actualW = coord.wIdx * DIGIT_2;
		123	+
		124	+ DataType dstUIdx[YUV_PER_DEAL_NUM];
		125	+ DataType dstVIdx[YUV_PER_DEAL_NUM];
		126	+ if (tD.outputFormat == NCHW_FORMAT_INDEX) {
		127	+ const DataType outputW = tD.outputSizeW;
		128	+ dstUIdx[YUV_DEAL_NUM_0] = dstUBase;
		129	+ dstUIdx[YUV_DEAL_NUM_1] = dstUBase + 1;
		130	+ dstUIdx[YUV_DEAL_NUM_2] = dstUBase + outputW;
		131	+ dstUIdx[YUV_DEAL_NUM_3] = dstUBase + outputW + 1;
		132	+ dstVIdx[YUV_DEAL_NUM_0] = dstVBase;
		133	+ dstVIdx[YUV_DEAL_NUM_1] = dstVBase + 1;
		134	+ dstVIdx[YUV_DEAL_NUM_2] = dstVBase + outputW;
		135	+ dstVIdx[YUV_DEAL_NUM_3] = dstVBase + outputW + 1;
		136	+ } else {
		137	+ const DataType rowStride = tD.outputSizeW * CHANNEL_THREE;
		138	+ dstUIdx[YUV_DEAL_NUM_0] = dstUBase;
		139	+ dstUIdx[YUV_DEAL_NUM_1] = dstUBase + CHANNEL_THREE;
		140	+ dstUIdx[YUV_DEAL_NUM_2] = dstUBase + rowStride;
		141	+ dstUIdx[YUV_DEAL_NUM_3] = dstUBase + rowStride + CHANNEL_THREE;
		142	+ dstVIdx[YUV_DEAL_NUM_0] = dstVBase;
		143	+ dstVIdx[YUV_DEAL_NUM_1] = dstVBase + CHANNEL_THREE;
		144	+ dstVIdx[YUV_DEAL_NUM_2] = dstVBase + rowStride;
		145	+ dstVIdx[YUV_DEAL_NUM_3] = dstVBase + rowStride + CHANNEL_THREE;
		146	+ }
146		147
147	for (uint8_t i = 0; i < YUV_PER_DEAL_NUM; i++) {	148	for (uint8_t i = 0; i < YUV_PER_DEAL_NUM; i++) {
148	- RgbPack<DataType> dstRgbIdxIndex = dstRgbIdx[i];
149	uint32_t pixelH = (i >= YUV_DEAL_NUM_2) ? actualH + 1 : actualH;	149	uint32_t pixelH = (i >= YUV_DEAL_NUM_2) ? actualH + 1 : actualH;
150	uint32_t pixelW = (i == YUV_DEAL_NUM_1 \|\| i == YUV_DEAL_NUM_3) ? actualW + 1 : actualW;	150	uint32_t pixelW = (i == YUV_DEAL_NUM_1 \|\| i == YUV_DEAL_NUM_3) ? actualW + 1 : actualW;
151	if (IsPixelInPaddingForYuv(pixelH, pixelW, tD, allEvenPadding, blockAllInPadding)) {	151	if (IsPixelInPaddingForYuv(pixelH, pixelW, tD, allEvenPadding, blockAllInPadding)) {
152	- AssignPadValue(outputGM[dstRgbIdxIndex.r], padValue);	152	+ AssignPadValue(outputGM[dstYIdx[i]], padValue);
153	- AssignPadValue(outputGM[dstRgbIdxIndex.g], padValue);	153	+ AssignPadValue(outputGM[dstUIdx[i]], padValue);
154	- AssignPadValue(outputGM[dstRgbIdxIndex.b], padValue);	154	+ AssignPadValue(outputGM[dstVIdx[i]], padValue);
155	} else {	155	} else {
156	uint32_t cropCoordH = pixelH - tD.paddingParam.topPaddingSize;	156	uint32_t cropCoordH = pixelH - tD.paddingParam.topPaddingSize;
157	uint32_t cropCoordW = pixelW - tD.paddingParam.leftPaddingSize;	157	uint32_t cropCoordW = pixelW - tD.paddingParam.leftPaddingSize;
158	- ProcessYuvPixel<T, DataType>(yuvGM, outputGM, dstRgbIdxIndex, coord.nIdx,	158	+ ProcessYuv444Pixel<T, DataType>(inputGM, outputGM,
159	- cropCoordH, cropCoordW, tD);	159	+ dstYIdx[i], dstUIdx[i], dstVIdx[i],
		160	+ coord.nIdx, cropCoordH, cropCoordW, tD);
160	}	161	}
161	}	162	}
162	}	163	}
163		164
164	template <typename T, typename DataType>	165	template <typename T, typename DataType>
165	-__simt_vf__ LAUNCH_BOUND(MAX_THREAD_NUM) __aicore__ void SimtComputeYuv(	166	+__simt_vf__ LAUNCH_BOUND(MAX_THREAD_NUM) __aicore__ void SimtComputeYuv420ToYuv444(
166	- __gm__ uint8_t* yuvGM, __gm__ T* outputGM, const AippTilingData tD,	167	+ __gm__ uint8_t* inputGM, __gm__ T* outputGM, const AippTilingData tD,
167	uint32_t blockIdx, uint32_t blockNum, uint64_t batchSize)	168	uint32_t blockIdx, uint32_t blockNum, uint64_t batchSize)
168	{	169	{
169	- uint32_t outputSizeH = tD.outputSizeH;	170	+ const uint32_t outputSizeH = tD.outputSizeH;
170	- uint32_t outputSizeW = tD.outputSizeW;	171	+ const uint32_t outputSizeW = tD.outputSizeW;
171	- int32_t paddingSwitch = tD.paddingParam.paddingSwitch;	172	+ const int32_t topPaddingSize = tD.paddingParam.topPaddingSize;
172	- int32_t leftPaddingSize = tD.paddingParam.leftPaddingSize;	173	+ const int32_t bottomPaddingSize = tD.paddingParam.bottomPaddingSize;
173	- int32_t rightPaddingSize = tD.paddingParam.rightPaddingSize;	174	+ const int32_t leftPaddingSize = tD.paddingParam.leftPaddingSize;
174	- int32_t topPaddingSize = tD.paddingParam.topPaddingSize;	175	+ const int32_t rightPaddingSize = tD.paddingParam.rightPaddingSize;
175	- int32_t bottomPaddingSize = tD.paddingParam.bottomPaddingSize;	176	+ const float padValue = tD.paddingParam.padValue;
176	- float padValue = tD.paddingParam.padValue;
177		177
178	- bool allEvenPadding = (topPaddingSize % DIGIT_2 == 0) &&	178	+ const bool allEvenPadding = (topPaddingSize % DIGIT_2 == 0) &&
179	- (bottomPaddingSize % DIGIT_2 == 0) &&	179	+ (bottomPaddingSize % DIGIT_2 == 0) &&
180	- (leftPaddingSize % DIGIT_2 == 0) &&	180	+ (leftPaddingSize % DIGIT_2 == 0) &&
181	- (rightPaddingSize % DIGIT_2 == 0);	181	+ (rightPaddingSize % DIGIT_2 == 0);
182		182
183	- for (DataType idx = Simt::GetThreadIdx() + blockIdx * Simt::GetThreadNum(); idx < batchSize;	183	+ for (DataType idx = threadIdx.x + blockIdx * blockDim.x;
184	- idx += blockNum * Simt::GetThreadNum()) {	184	+ idx < batchSize;
		185	+ idx += blockNum * blockDim.x) {
		186	+
		187	+ // 1. Decode idx → 2x2-block coordinate (nIdx, hIdx, wIdx)
185	CoordPack<DataType> coord;	188	CoordPack<DataType> coord;
186	coord.nIdx = idx / ((outputSizeH >> 1) * (outputSizeW >> 1));	189	coord.nIdx = idx / ((outputSizeH >> 1) * (outputSizeW >> 1));
187	- DataType newIdx = idx - coord.nIdx * ((outputSizeH >> 1) * (outputSizeW >> 1));	190	+ DataType rem = idx - coord.nIdx * ((outputSizeH >> 1) * (outputSizeW >> 1));
188	- coord.hIdx = newIdx / (outputSizeW >> 1);	191	+ coord.hIdx = rem / (outputSizeW >> 1);
189	- coord.wIdx = newIdx - coord.hIdx * (outputSizeW >> 1);	192	+ coord.wIdx = rem - coord.hIdx * (outputSizeW >> 1);
190		193
191	- RgbPack<DataType> dstRgbIdx[YUV_PER_DEAL_NUM];	194	+ // 2. Compute destination indices for this 2x2 block
192	- ComputeDstIdx(dstRgbIdx, coord, tD);	195	+ DataType dstYIdx[YUV_PER_DEAL_NUM];
		196	+ DataType dstUBase, dstVBase;
		197	+ ComputeDstYuv444Idx<DataType>(dstYIdx, dstUBase, dstVBase, coord, tD);
193		198
194	- uint32_t actualH = coord.hIdx * DIGIT_2;	199	+ // 3. Fast-path: check if entire 2x2 block is in padding
195	- uint32_t actualW = coord.wIdx * DIGIT_2;	200	+ const uint32_t actualH = coord.hIdx * DIGIT_2;
		201	+ const uint32_t actualW = coord.wIdx * DIGIT_2;
		202	+ const bool blockAllInPadding = (tD.paddingParam.paddingSwitch != 0 && allEvenPadding) &&
		203	+ ((actualH < topPaddingSize) \|\|
		204	+ (actualH >= topPaddingSize + tD.cropParam.cropSizeH) \|\|
		205	+ (actualW < leftPaddingSize) \|\|
		206	+ (actualW >= leftPaddingSize + tD.cropParam.cropSizeW));
196		207
197	- bool blockAllInPadding = (paddingSwitch != 0 && allEvenPadding) &&	208	+ // 4. Process block (padding check + CSC + DataConversion per pixel)
198	- ((actualH < topPaddingSize) \|\|	209	+ ProcessYuv444Block<T, DataType>(inputGM, outputGM,
199	- (actualH >= topPaddingSize + tD.cropParam.cropSizeH) \|\|	210	+ dstYIdx, dstUBase, dstVBase, coord, tD,
200	- (actualW < leftPaddingSize) \|\|	211	+ padValue, allEvenPadding, blockAllInPadding);
201	- (actualW >= leftPaddingSize + tD.cropParam.cropSizeW));
202	-
203	- ProcessYuvBlock<T, DataType>(yuvGM, outputGM, dstRgbIdx, coord, tD,
204	- padValue, allEvenPadding, blockAllInPadding);
205	}	212	}
206	}	213	}
207		214
208	template <typename T, typename DataType>	215	template <typename T, typename DataType>
209	__aicore__ inline void AippYuv<T, DataType>::Process(GM_ADDR x, GM_ADDR y)	216	__aicore__ inline void AippYuv<T, DataType>::Process(GM_ADDR x, GM_ADDR y)
210	{	217	{
211	- uint64_t batchSize = this->tilingData_.batchNum *	218	+ uint64_t batchSize = (uint64_t)this->tilingData_.batchNum *
212	- ((this->tilingData_.outputSizeH) >> 1) * ((this->tilingData_.outputSizeW) >> 1);	219	+ ((uint64_t)this->tilingData_.outputSizeH >> 1) *
		220	+ ((uint64_t)this->tilingData_.outputSizeW >> 1);
213		221
214	- Simt::VF_CALL<Aipp_Kernel::SimtComputeYuv<T, DataType>>(Simt::Dim3(this->blockDimX_),	222	+ asc_vf_call<Aipp_Kernel::SimtComputeYuv420ToYuv444<T, DataType>>(dim3(this->blockDimX_),
215	(__gm__ uint8_t)x, (__gm__ T)y, this->tilingData_, this->blockIdx_, this->blockNum_, batchSize);	223	(__gm__ uint8_t)x, (__gm__ T)y, this->tilingData_, this->blockIdx_, this->blockNum_, batchSize);
216	}	224	}
217		225
218	} // namespace Aipp_Kernel	226	} // namespace Aipp_Kernel
219	-#endif // AIPP_OP_KERNEL_ARCH35_YUV_RGB_H	227	+#endif // AIPP_OP_KERNEL_ARCH35_YUV_H

Mimage/aipp/op_kernel/arch35/aipp_yuv_gray.h+21-15

	@@ -11,6 +11,10 @@
11	/*!	11	/*!
12	* \file aipp_yuv_gray.h	12	* \file aipp_yuv_gray.h
13	* \brief aipp yuv to gray	13	* \brief aipp yuv to gray
		14	+ *
		15	+ * Supports YUV420SP_U8 -> GRAY and YUV400_U8 -> GRAY
		16	+ * YUV420SP: YUV420SP format with size HW1.5, extracts Y component from interleaved storage
		17	+ * YUV400: Single-channel grayscale format with size H*W, direct access
14	*/	18	*/
15		19
16	#ifndef AIPP_OP_KERNEL_ARCH35_YUV_GRAY_H	20	#ifndef AIPP_OP_KERNEL_ARCH35_YUV_GRAY_H
	@@ -39,13 +43,6 @@ __aicore__ inline void AippYuvGray<T, DataType>::Init(const AippTilingData& tili
39	this->BaseInit(tilingData);	43	this->BaseInit(tilingData);
40	}	44	}
41		45
42	-__aicore__ __attribute__((always_inline)) inline uint8_t Yuv2Gray(uint8_t y, const CscParam& cscParam)
43	-{
44	- auto gTmp = static_cast<int16_t>(roundf(static_cast<float>(
45	- cscParam.cscMatrix00 * static_cast<int16_t>(y) * DIGIT_2 + 1) / CSC_MATRIX_SCALE));
46	- return CLIP3(gTmp, 0, MAX_UINT8);
47	-}
48	-
49	template <typename T, typename DataType>	46	template <typename T, typename DataType>
50	__simt_vf__ LAUNCH_BOUND(MAX_THREAD_NUM) __aicore__ void SimtComputeYuvGray(	47	__simt_vf__ LAUNCH_BOUND(MAX_THREAD_NUM) __aicore__ void SimtComputeYuvGray(
51	__gm__ uint8_t* yuvGM, __gm__ T* outputGM, const AippTilingData tD,	48	__gm__ uint8_t* yuvGM, __gm__ T* outputGM, const AippTilingData tD,
	@@ -54,9 +51,10 @@ __simt_vf__ LAUNCH_BOUND(MAX_THREAD_NUM) __aicore__ void SimtComputeYuvGray(
54	uint32_t outputSizeH = tD.outputSizeH;	51	uint32_t outputSizeH = tD.outputSizeH;
55	uint32_t outputSizeW = tD.outputSizeW;	52	uint32_t outputSizeW = tD.outputSizeW;
56	float padValue = tD.paddingParam.padValue;	53	float padValue = tD.paddingParam.padValue;
		54	+ bool isYuv400 = (tD.imageFormat == YUV400_U8_FORMAT);
57		55
58	- for (DataType idx = Simt::GetThreadIdx() + blockIdx * Simt::GetThreadNum(); idx < batchSize;	56	+ for (DataType idx = threadIdx.x + blockIdx * blockDim.x; idx < batchSize;
59	- idx += blockNum * Simt::GetThreadNum()) {	57	+ idx += blockNum * blockDim.x) {
60	CoordPack<DataType> coord;	58	CoordPack<DataType> coord;
61	RgbPack<DataType> dstGrayIdx;	59	RgbPack<DataType> dstGrayIdx;
62	ComputeCoordFromIndex(idx, outputSizeH, outputSizeW, coord);	60	ComputeCoordFromIndex(idx, outputSizeH, outputSizeW, coord);
	@@ -68,11 +66,19 @@ __simt_vf__ LAUNCH_BOUND(MAX_THREAD_NUM) __aicore__ void SimtComputeYuvGray(
68	AssignPadValue(outputGM[dstGrayIdx.g], padValue);	66	AssignPadValue(outputGM[dstGrayIdx.g], padValue);
69	AssignPadValue(outputGM[dstGrayIdx.b], padValue);	67	AssignPadValue(outputGM[dstGrayIdx.b], padValue);
70	} else {	68	} else {
71	- DataType srcYuvYIdx = coord.nIdx * tD.inputSizeH * tD.inputSizeW * DIGIT_3 / DIGIT_2 +	69	+ DataType srcIdx;
72	- (tD.cropParam.cropStartPosH + coord.hIdx - tD.paddingParam.topPaddingSize) * tD.inputSizeW +	70	+ if (isYuv400) {
73	- tD.cropParam.cropStartPosW + coord.wIdx - tD.paddingParam.leftPaddingSize;	71	+ srcIdx = coord.nIdx * tD.inputSizeH * tD.inputSizeW +
74	- uint8_t grayValue = Yuv2Gray(yuvGM[srcYuvYIdx], tD.cscParam);	72	+ (tD.cropParam.cropStartPosH + coord.hIdx - tD.paddingParam.topPaddingSize) * tD.inputSizeW +
75	- DataConversion(outputGM[dstGrayIdx.r], grayValue, tD.dtcParam, CHANNEL_NUM_0);	73	+ tD.cropParam.cropStartPosW + coord.wIdx - tD.paddingParam.leftPaddingSize;
		74	+ } else {
		75	+ srcIdx = coord.nIdx * tD.inputSizeH * tD.inputSizeW * DIGIT_3 / DIGIT_2 +
		76	+ (tD.cropParam.cropStartPosH + coord.hIdx - tD.paddingParam.topPaddingSize) * tD.inputSizeW +
		77	+ tD.cropParam.cropStartPosW + coord.wIdx - tD.paddingParam.leftPaddingSize;
		78	+ }
		79	+ RgbPack<uint8_t> result;
		80	+ ApplyCscMatrix(result, yuvGM[srcIdx], 0, 0, tD.cscParam);
		81	+ DataConversion(outputGM[dstGrayIdx.r], result.r, tD.dtcParam, CHANNEL_NUM_0);
76	DataConversion(outputGM[dstGrayIdx.g], 0, tD.dtcParam, CHANNEL_NUM_1);	82	DataConversion(outputGM[dstGrayIdx.g], 0, tD.dtcParam, CHANNEL_NUM_1);
77	DataConversion(outputGM[dstGrayIdx.b], 0, tD.dtcParam, CHANNEL_NUM_2);	83	DataConversion(outputGM[dstGrayIdx.b], 0, tD.dtcParam, CHANNEL_NUM_2);
78	}	84	}
	@@ -85,7 +91,7 @@ __aicore__ inline void AippYuvGray<T, DataType>::Process(GM_ADDR x, GM_ADDR y)
85	uint64_t batchSize = static_cast<uint64_t>(this->tilingData_.batchNum) *	91	uint64_t batchSize = static_cast<uint64_t>(this->tilingData_.batchNum) *
86	this->tilingData_.outputSizeH * this->tilingData_.outputSizeW;	92	this->tilingData_.outputSizeH * this->tilingData_.outputSizeW;
87		93
88	- Simt::VF_CALL<Aipp_Kernel::SimtComputeYuvGray<T, DataType>>(Simt::Dim3(this->blockDimX_),	94	+ asc_vf_call<Aipp_Kernel::SimtComputeYuvGray<T, DataType>>(dim3(this->blockDimX_),
89	(__gm__ uint8_t)x, (__gm__ T)y, this->tilingData_, this->blockIdx_, this->blockNum_, batchSize);	95	(__gm__ uint8_t)x, (__gm__ T)y, this->tilingData_, this->blockIdx_, this->blockNum_, batchSize);
90	}	96	}
91		97

Aimage/aipp/op_kernel/arch35/aipp_yuv_rgb.h+195-0

@@ -0,0 +1,195 @@
	1	+/**
	2	+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file aipp_yuv_rgb.h
	13	+ * \brief AIPP YUV420SP to RGB/BGR conversion kernel (with CSC matrix).
	14	+ *
	15	+ * Moved from aipp_yuv.h. Handles YUV420SP_U8 → RGB/BGR output via a
	16	+ * configurable 3x3 CSC matrix (ApplyCscMatrix).
	17	+ *
	18	+ * TilingKey = FORMAT_YUV_2_RGB_UINT32 (5)
	19	+ */
	20	+
	21	+#ifndef AIPP_OP_KERNEL_ARCH35_YUV_RGB_H
	22	+#define AIPP_OP_KERNEL_ARCH35_YUV_RGB_H
	23	+
	24	+#include "aipp_base.h"
	25	+#include "simt_api/math_functions.h"
	26	+
	27	+namespace Aipp_Kernel {
	28	+using namespace AscendC;
	29	+
	30	+template <typename T, typename DataType>
	31	+class AippYuvRgb : public AippBase<T, DataType> {
	32	+public:
	33	+ __aicore__ inline AippYuvRgb(){};
	34	+ __aicore__ inline void Init(const AippTilingData& tilingData);
	35	+ __aicore__ inline void Process(GM_ADDR x, GM_ADDR y);
	36	+
	37	+private:
	38	+ uint16_t blockDimX_ = MAX_THREAD_NUM;
	39	+};
	40	+
	41	+template <typename T, typename DataType>
	42	+__aicore__ inline void AippYuvRgb<T, DataType>::Init(const AippTilingData& tilingData)
	43	+{
	44	+ this->BaseInit(tilingData);
	45	+}
	46	+
	47	+template <typename DataType>
	48	+__aicore__ __attribute__((always_inline)) inline void ComputeYuvRgbDstIdx(
	49	+ RgbPack<DataType>* dstRgbIdx, const CoordPack<DataType>& coord, const AippTilingData& tD)
	50	+{
	51	+ if (tD.outputFormat == NCHW_FORMAT_INDEX) {
	52	+ dstRgbIdx[YUV_DEAL_NUM_0].r = coord.nIdx * (tD.outputSizeH * tD.outputSizeW * CHANNEL_THREE) +
	53	+ DIGIT_2 * coord.hIdx * tD.outputSizeW + DIGIT_2 * coord.wIdx;
	54	+ dstRgbIdx[YUV_DEAL_NUM_0].g = dstRgbIdx[YUV_DEAL_NUM_0].r + tD.outputSizeH * tD.outputSizeW;
	55	+ dstRgbIdx[YUV_DEAL_NUM_0].b = dstRgbIdx[YUV_DEAL_NUM_0].g + tD.outputSizeH * tD.outputSizeW;
	56	+
	57	+ dstRgbIdx[YUV_DEAL_NUM_1].r = dstRgbIdx[YUV_DEAL_NUM_0].r + 1;
	58	+ dstRgbIdx[YUV_DEAL_NUM_1].g = dstRgbIdx[YUV_DEAL_NUM_0].g + 1;
	59	+ dstRgbIdx[YUV_DEAL_NUM_1].b = dstRgbIdx[YUV_DEAL_NUM_0].b + 1;
	60	+
	61	+ dstRgbIdx[YUV_DEAL_NUM_2].r = coord.nIdx * (tD.outputSizeH * tD.outputSizeW * CHANNEL_THREE) +
	62	+ (DIGIT_2 * coord.hIdx + 1) * tD.outputSizeW + DIGIT_2 * coord.wIdx;
	63	+ dstRgbIdx[YUV_DEAL_NUM_2].g = dstRgbIdx[YUV_DEAL_NUM_2].r + tD.outputSizeH * tD.outputSizeW;
	64	+ dstRgbIdx[YUV_DEAL_NUM_2].b = dstRgbIdx[YUV_DEAL_NUM_2].g + tD.outputSizeH * tD.outputSizeW;
	65	+
	66	+ dstRgbIdx[YUV_DEAL_NUM_3].r = dstRgbIdx[YUV_DEAL_NUM_2].r + 1;
	67	+ dstRgbIdx[YUV_DEAL_NUM_3].g = dstRgbIdx[YUV_DEAL_NUM_2].g + 1;
	68	+ dstRgbIdx[YUV_DEAL_NUM_3].b = dstRgbIdx[YUV_DEAL_NUM_2].b + 1;
	69	+ } else {
	70	+ dstRgbIdx[YUV_DEAL_NUM_0].r = coord.nIdx * (tD.outputSizeH * tD.outputSizeW * CHANNEL_THREE) +
	71	+ DIGIT_2 * coord.hIdx * tD.outputSizeW * CHANNEL_THREE +
	72	+ DIGIT_2 * coord.wIdx * CHANNEL_THREE;
	73	+ dstRgbIdx[YUV_DEAL_NUM_0].g = dstRgbIdx[YUV_DEAL_NUM_0].r + 1;
	74	+ dstRgbIdx[YUV_DEAL_NUM_0].b = dstRgbIdx[YUV_DEAL_NUM_0].g + 1;
	75	+
	76	+ dstRgbIdx[YUV_DEAL_NUM_1].r = dstRgbIdx[YUV_DEAL_NUM_0].r + CHANNEL_THREE;
	77	+ dstRgbIdx[YUV_DEAL_NUM_1].g = dstRgbIdx[YUV_DEAL_NUM_0].g + CHANNEL_THREE;
	78	+ dstRgbIdx[YUV_DEAL_NUM_1].b = dstRgbIdx[YUV_DEAL_NUM_0].b + CHANNEL_THREE;
	79	+
	80	+ dstRgbIdx[YUV_DEAL_NUM_2].r = coord.nIdx * (tD.outputSizeH * tD.outputSizeW * CHANNEL_THREE) +
	81	+ (DIGIT_2 * coord.hIdx + 1) * tD.outputSizeW * CHANNEL_THREE +
	82	+ DIGIT_2 * coord.wIdx * CHANNEL_THREE;
	83	+ dstRgbIdx[YUV_DEAL_NUM_2].g = dstRgbIdx[YUV_DEAL_NUM_2].r + 1;
	84	+ dstRgbIdx[YUV_DEAL_NUM_2].b = dstRgbIdx[YUV_DEAL_NUM_2].g + 1;
	85	+
	86	+ dstRgbIdx[YUV_DEAL_NUM_3].r = dstRgbIdx[YUV_DEAL_NUM_2].r + CHANNEL_THREE;
	87	+ dstRgbIdx[YUV_DEAL_NUM_3].g = dstRgbIdx[YUV_DEAL_NUM_2].g + CHANNEL_THREE;
	88	+ dstRgbIdx[YUV_DEAL_NUM_3].b = dstRgbIdx[YUV_DEAL_NUM_2].b + CHANNEL_THREE;
	89	+ }
	90	+}
	91	+
	92	+// IsPixelInPaddingForYuv is defined in aipp_base.h
	93	+
	94	+template <typename T, typename DataType>
	95	+__aicore__ __attribute__((always_inline)) inline void ProcessYuvRgbPixel(
	96	+ __gm__ uint8_t* yuvGM, __gm__ T* outputGM, const RgbPack<DataType> dstIdx,
	97	+ uint32_t nIdx, uint32_t croodH, uint32_t croodW, const AippTilingData& tD)
	98	+{
	99	+ uint32_t srcYIdx = nIdx * tD.inputSizeH * tD.inputSizeW * 3 / 2 +
	100	+ (tD.cropParam.cropStartPosH + croodH) * tD.inputSizeW +
	101	+ (tD.cropParam.cropStartPosW + croodW);
	102	+ uint32_t srcUIdx = nIdx * tD.inputSizeH * tD.inputSizeW * 3 / 2 + tD.inputSizeH * tD.inputSizeW +
	S S sunday4月24日这里 `srcUIdx/srcVIdx` 固定按 YUV420SP 的 1.5 倍平面去取 UV，但本 PR 已支持 `YUV400_U8`。当 `YUV400_U8 + csc_switch=true` 走到 YUV->RGB 分支时，这里会访问不存在的 UV 区域，存在越界读风险。建议在 tiling 侧禁止该组合（直接报错），或新增 YUV400->RGB 专用实现并避免读取 UV。 J jinyijian 4月24日评论： 4月24日评论： J jinyijian 4月24日评论： 4月24日评论：
	103	+ ((tD.cropParam.cropStartPosH + (croodH & ~1)) >> 1) * tD.inputSizeW +
	104	+ (tD.cropParam.cropStartPosW + (croodW & ~1));
	105	+ uint32_t srcVIdx = srcUIdx + 1;
	106	+ RgbPack<uint8_t> dstRgb;
	107	+ ApplyCscMatrix(dstRgb, yuvGM[srcYIdx], yuvGM[srcUIdx], yuvGM[srcVIdx], tD.cscParam);
	108	+
	109	+ DataConversion(outputGM[dstIdx.r], dstRgb.r, tD.dtcParam, CHANNEL_NUM_0);
	110	+ DataConversion(outputGM[dstIdx.g], dstRgb.g, tD.dtcParam, CHANNEL_NUM_1);
	111	+ DataConversion(outputGM[dstIdx.b], dstRgb.b, tD.dtcParam, CHANNEL_NUM_2);
	112	+}
	113	+
	114	+template <typename T, typename DataType>
	115	+__aicore__ __attribute__((always_inline)) inline void ProcessYuvRgbBlock(
	116	+ __gm__ uint8_t* yuvGM, __gm__ T* outputGM,
	117	+ const RgbPack<DataType>* dstRgbIdx, const CoordPack<DataType>& coord,
	118	+ const AippTilingData& tD, float padValue, bool allEvenPadding, bool blockAllInPadding)
	119	+{
	120	+ uint32_t actualH = coord.hIdx * DIGIT_2;
	121	+ uint32_t actualW = coord.wIdx * DIGIT_2;
	122	+
	123	+ for (uint8_t i = 0; i < YUV_PER_DEAL_NUM; i++) {
	124	+ RgbPack<DataType> dstRgbIdxIndex = dstRgbIdx[i];
	125	+ uint32_t pixelH = (i >= YUV_DEAL_NUM_2) ? actualH + 1 : actualH;
	126	+ uint32_t pixelW = (i == YUV_DEAL_NUM_1 \|\| i == YUV_DEAL_NUM_3) ? actualW + 1 : actualW;
	127	+ if (IsPixelInPaddingForYuv(pixelH, pixelW, tD, allEvenPadding, blockAllInPadding)) {
	128	+ AssignPadValue(outputGM[dstRgbIdxIndex.r], padValue);
	129	+ AssignPadValue(outputGM[dstRgbIdxIndex.g], padValue);
	130	+ AssignPadValue(outputGM[dstRgbIdxIndex.b], padValue);
	131	+ } else {
	132	+ uint32_t cropCoordH = pixelH - tD.paddingParam.topPaddingSize;
	133	+ uint32_t cropCoordW = pixelW - tD.paddingParam.leftPaddingSize;
	134	+ ProcessYuvRgbPixel<T, DataType>(yuvGM, outputGM, dstRgbIdxIndex, coord.nIdx,
	135	+ cropCoordH, cropCoordW, tD);
	136	+ }
	137	+ }
	138	+}
	139	+
	140	+template <typename T, typename DataType>
	141	+__simt_vf__ LAUNCH_BOUND(MAX_THREAD_NUM) __aicore__ void SimtComputeYuvRgb(
	142	+ __gm__ uint8_t* yuvGM, __gm__ T* outputGM, const AippTilingData tD,
	143	+ uint32_t blockIdx, uint32_t blockNum, uint64_t batchSize)
	144	+{
	145	+ uint32_t outputSizeH = tD.outputSizeH;
	146	+ uint32_t outputSizeW = tD.outputSizeW;
	147	+ int32_t paddingSwitch = tD.paddingParam.paddingSwitch;
	148	+ int32_t leftPaddingSize = tD.paddingParam.leftPaddingSize;
	149	+ int32_t rightPaddingSize = tD.paddingParam.rightPaddingSize;
	150	+ int32_t topPaddingSize = tD.paddingParam.topPaddingSize;
	151	+ int32_t bottomPaddingSize = tD.paddingParam.bottomPaddingSize;
	152	+ float padValue = tD.paddingParam.padValue;
	153	+
	154	+ bool allEvenPadding = (topPaddingSize % DIGIT_2 == 0) &&
	155	+ (bottomPaddingSize % DIGIT_2 == 0) &&
	156	+ (leftPaddingSize % DIGIT_2 == 0) &&
	157	+ (rightPaddingSize % DIGIT_2 == 0);
	158	+
	159	+ for (DataType idx = threadIdx.x + blockIdx * blockDim.x; idx < batchSize;
	160	+ idx += blockNum * blockDim.x) {
	161	+ CoordPack<DataType> coord;
	162	+ coord.nIdx = idx / ((outputSizeH >> 1) * (outputSizeW >> 1));
	163	+ DataType newIdx = idx - coord.nIdx * ((outputSizeH >> 1) * (outputSizeW >> 1));
	164	+ coord.hIdx = newIdx / (outputSizeW >> 1);
	165	+ coord.wIdx = newIdx - coord.hIdx * (outputSizeW >> 1);
	166	+
	167	+ RgbPack<DataType> dstRgbIdx[YUV_PER_DEAL_NUM];
	168	+ ComputeYuvRgbDstIdx(dstRgbIdx, coord, tD);
	169	+
	170	+ uint32_t actualH = coord.hIdx * DIGIT_2;
	171	+ uint32_t actualW = coord.wIdx * DIGIT_2;
	172	+
	173	+ bool blockAllInPadding = (paddingSwitch != 0 && allEvenPadding) &&
	174	+ ((actualH < topPaddingSize) \|\|
	175	+ (actualH >= topPaddingSize + tD.cropParam.cropSizeH) \|\|
	176	+ (actualW < leftPaddingSize) \|\|
	177	+ (actualW >= leftPaddingSize + tD.cropParam.cropSizeW));
	178	+
	179	+ ProcessYuvRgbBlock<T, DataType>(yuvGM, outputGM, dstRgbIdx, coord, tD,
	180	+ padValue, allEvenPadding, blockAllInPadding);
	181	+ }
	182	+}
	183	+
	184	+template <typename T, typename DataType>
	185	+__aicore__ inline void AippYuvRgb<T, DataType>::Process(GM_ADDR x, GM_ADDR y)
	186	+{
	187	+ uint64_t batchSize = this->tilingData_.batchNum *
	188	+ ((this->tilingData_.outputSizeH) >> 1) * ((this->tilingData_.outputSizeW) >> 1);
	189	+
	190	+ asc_vf_call<Aipp_Kernel::SimtComputeYuvRgb<T, DataType>>(dim3(this->blockDimX_),
	191	+ (__gm__ uint8_t)x, (__gm__ T)y, this->tilingData_, this->blockIdx_, this->blockNum_, batchSize);
	192	+}
	193	+
	194	+} // namespace Aipp_Kernel
	195	+#endif // AIPP_OP_KERNEL_ARCH35_YUV_RGB_H

Aimage/aipp/tests/ut/op_host/aipp_ut_test_12.cfg+9-0

@@ -0,0 +1,9 @@
	1	+# YUV420SP_U8转YUV444 - 不带swap
	2	+aipp_op {
	3	+ aipp_mode: static
	4	+ input_format : YUV420SP_U8
	5	+ csc_switch : false
	6	+ rbuv_swap_switch : false
	7	+ src_image_size_w : 1920
	8	+ src_image_size_h : 1080
	9	+}

Aimage/aipp/tests/ut/op_host/aipp_ut_test_13.cfg+9-0

@@ -0,0 +1,9 @@
	1	+# YUV420SP_U8转YVU444 - 带rbuv swap
	2	+aipp_op {
	3	+ aipp_mode: static
	4	+ input_format : YUV420SP_U8
	5	+ csc_switch : false
	6	+ rbuv_swap_switch : true
	7	+ src_image_size_w : 1920
	8	+ src_image_size_h : 1080
	9	+}

Aimage/aipp/tests/ut/op_host/aipp_ut_test_14.cfg+9-0

@@ -0,0 +1,9 @@
	1	+# RGB888_U8转RGB - 不带csc和swap
	2	+aipp_op {
	3	+ aipp_mode: static
	4	+ input_format : RGB888_U8
	5	+ csc_switch : false
	6	+ rbuv_swap_switch : false
	7	+ src_image_size_w : 1920
	8	+ src_image_size_h : 1080
	9	+}

Aimage/aipp/tests/ut/op_host/aipp_ut_test_15.cfg+21-0

@@ -0,0 +1,21 @@
	1	+# RGB888_U8转YVU444 - 带csc矩阵转换
	2	+aipp_op {
	3	+ aipp_mode: static
	4	+ input_format : RGB888_U8
	5	+ csc_switch : true
	6	+ rbuv_swap_switch : false
	7	+ matrix_r0c0 : 77
	8	+ matrix_r0c1 : 150
	9	+ matrix_r0c2 : 29
	10	+ matrix_r1c0 : 128
	11	+ matrix_r1c1 : -107
	12	+ matrix_r1c2 : -21
	13	+ matrix_r2c0 : -43
	14	+ matrix_r2c1 : -85
	15	+ matrix_r2c2 : 128
	16	+ output_bias_0 : 0
	17	+ output_bias_1 : 128
	18	+ output_bias_2 : 128
	19	+ src_image_size_w : 1920
	20	+ src_image_size_h : 1080
	21	+}

Aimage/aipp/tests/ut/op_host/aipp_ut_test_16.cfg+9-0

@@ -0,0 +1,9 @@
	1	+# BGR888_U8转RGB - 带rbuv swap
	2	+aipp_op {
	3	+ aipp_mode: static
	4	+ input_format : RGB888_U8
	5	+ csc_switch : false
	6	+ rbuv_swap_switch : true
	7	+ src_image_size_w : 1920
	8	+ src_image_size_h : 1080
	9	+}

Aimage/aipp/tests/ut/op_host/aipp_ut_test_17.cfg+9-0

@@ -0,0 +1,9 @@
	1	+# BGR888_U8转BGR - 不带swap
	2	+aipp_op {
	3	+ aipp_mode: static
	4	+ input_format : RGB888_U8
	5	+ csc_switch : false
	6	+ rbuv_swap_switch : false
	7	+ src_image_size_w : 1920
	8	+ src_image_size_h : 1080
	9	+}

Aimage/aipp/tests/ut/op_host/aipp_ut_test_18.cfg+22-0

@@ -0,0 +1,22 @@
	1	+# XRGB8888_U8转YUV444 - 带ax swap和csc矩阵
	2	+aipp_op {
	3	+ aipp_mode: static
	4	+ input_format : XRGB8888_U8
	5	+ csc_switch : true
	6	+ rbuv_swap_switch : false
	7	+ ax_swap_switch : true
	8	+ matrix_r0c0 : 77
	9	+ matrix_r0c1 : 150
	10	+ matrix_r0c2 : 29
	11	+ matrix_r1c0 : -43
	12	+ matrix_r1c1 : -85
	13	+ matrix_r1c2 : 128
	14	+ matrix_r2c0 : 128
	15	+ matrix_r2c1 : -107
	16	+ matrix_r2c2 : -21
	17	+ output_bias_0 : 0
	18	+ output_bias_1 : 128
	19	+ output_bias_2 : 128
	20	+ src_image_size_w : 1920
	21	+ src_image_size_h : 1080
	22	+}

Aimage/aipp/tests/ut/op_host/aipp_ut_test_19.cfg+22-0

@@ -0,0 +1,22 @@
	1	+# XRGB8888_U8转YVU444 - 带ax swap和csc矩阵(UV互换)
	2	+aipp_op {
	3	+ aipp_mode: static
	4	+ input_format : XRGB8888_U8
	5	+ csc_switch : true
	6	+ rbuv_swap_switch : false
	7	+ ax_swap_switch : true
	8	+ matrix_r0c0 : 77
	9	+ matrix_r0c1 : 150
	10	+ matrix_r0c2 : 29
	11	+ matrix_r1c0 : 128
	12	+ matrix_r1c1 : -107
	13	+ matrix_r1c2 : -21
	14	+ matrix_r2c0 : -43
	15	+ matrix_r2c1 : -85
	16	+ matrix_r2c2 : 128
	17	+ output_bias_0 : 0
	18	+ output_bias_1 : 128
	19	+ output_bias_2 : 128
	20	+ src_image_size_w : 1920
	21	+ src_image_size_h : 1080
	22	+}

Aimage/aipp/tests/ut/op_host/aipp_ut_test_20.cfg+22-0

@@ -0,0 +1,22 @@
	1	+# XRGB8888_U8转GRAY - 带ax swap和灰度矩阵
	2	+aipp_op {
	3	+ aipp_mode: static
	4	+ input_format : XRGB8888_U8
	5	+ csc_switch : true
	6	+ rbuv_swap_switch : false
	7	+ ax_swap_switch : true
	8	+ matrix_r0c0 : 76
	9	+ matrix_r0c1 : 150
	10	+ matrix_r0c2 : 30
	11	+ matrix_r1c0 : 0
	12	+ matrix_r1c1 : 0
	13	+ matrix_r1c2 : 0
	14	+ matrix_r2c0 : 0
	15	+ matrix_r2c1 : 0
	16	+ matrix_r2c2 : 0
	17	+ output_bias_0 : 0
	18	+ output_bias_1 : 0
	19	+ output_bias_2 : 0
	20	+ src_image_size_w : 1920
	21	+ src_image_size_h : 1080
	22	+}

Aimage/aipp/tests/ut/op_host/aipp_ut_test_21.cfg+22-0

@@ -0,0 +1,22 @@
	1	+# XBGR8888_U8转GRAY - 带ax swap和rbuv swap
	2	+aipp_op {
	3	+ aipp_mode: static
	4	+ input_format : XRGB8888_U8
	5	+ csc_switch : true
	6	+ rbuv_swap_switch : true
	7	+ ax_swap_switch : true
	8	+ matrix_r0c0 : 76
	9	+ matrix_r0c1 : 150
	10	+ matrix_r0c2 : 30
	11	+ matrix_r1c0 : 0
	12	+ matrix_r1c1 : 0
	13	+ matrix_r1c2 : 0
	14	+ matrix_r2c0 : 0
	15	+ matrix_r2c1 : 0
	16	+ matrix_r2c2 : 0
	17	+ output_bias_0 : 0
	18	+ output_bias_1 : 0
	19	+ output_bias_2 : 0
	20	+ src_image_size_w : 1920
	21	+ src_image_size_h : 1080
	22	+}

Aimage/aipp/tests/ut/op_host/aipp_ut_test_22.cfg+22-0

@@ -0,0 +1,22 @@
	1	+# RGBX8888_U8转GRAY - 带ax swap关闭
	2	+aipp_op {
	3	+ aipp_mode: static
	4	+ input_format : XRGB8888_U8
	5	+ csc_switch : true
	6	+ rbuv_swap_switch : false
	7	+ ax_swap_switch : false
	8	+ matrix_r0c0 : 76
	9	+ matrix_r0c1 : 150
	10	+ matrix_r0c2 : 30
	11	+ matrix_r1c0 : 0
	12	+ matrix_r1c1 : 0
	13	+ matrix_r1c2 : 0
	14	+ matrix_r2c0 : 0
	15	+ matrix_r2c1 : 0
	16	+ matrix_r2c2 : 0
	17	+ output_bias_0 : 0
	18	+ output_bias_1 : 0
	19	+ output_bias_2 : 0
	20	+ src_image_size_w : 1920
	21	+ src_image_size_h : 1080
	22	+}

Aimage/aipp/tests/ut/op_host/aipp_ut_test_23.cfg+22-0

@@ -0,0 +1,22 @@
	1	+# BGRX8888_U8转GRAY - 带rbuv swap和ax swap关闭
	2	+aipp_op {
	3	+ aipp_mode: static
	4	+ input_format : XRGB8888_U8
	5	+ csc_switch : true
	6	+ rbuv_swap_switch : true
	7	+ ax_swap_switch : false
	8	+ matrix_r0c0 : 76
	9	+ matrix_r0c1 : 150
	10	+ matrix_r0c2 : 30
	11	+ matrix_r1c0 : 0
	12	+ matrix_r1c1 : 0
	13	+ matrix_r1c2 : 0
	14	+ matrix_r2c0 : 0
	15	+ matrix_r2c1 : 0
	16	+ matrix_r2c2 : 0
	17	+ output_bias_0 : 0
	18	+ output_bias_1 : 0
	19	+ output_bias_2 : 0
	20	+ src_image_size_w : 1920
	21	+ src_image_size_h : 1080
	22	+}

Aimage/aipp/tests/ut/op_host/aipp_ut_test_24.cfg+8-0

@@ -0,0 +1,8 @@
	1	+# YUV400_U8转GRAY - 不带csc
	2	+aipp_op {
	3	+ aipp_mode: static
	4	+ input_format : YUV400_U8
	5	+ csc_switch : false
	6	+ src_image_size_w : 1920
	7	+ src_image_size_h : 1080
	8	+}

Mimage/aipp/tests/ut/op_host/test_aipp_tiling.cpp+1750-111

	@@ -237,11 +237,74 @@ static void GetPlatFormInfos(const char* compileInfoStr, map<string, string>& so
237	}	237	}
238	}	238	}
239		239
240	-void AippExecuteTestCase(const gert::TilingContextPara& tilingContextPara,	240	+void CompareAippTilingData(const Aipp_Kernel::AippTilingData& actual,
241	- ge::graphStatus expectResult,	241	+ const Aipp_Kernel::AippTilingData& expected)
242	- uint64_t expectTilingKey,	242	+{
243	- const string& expectTilingData,	243	+ EXPECT_EQ(actual.imageFormat, expected.imageFormat);
244	- const std::vector<size_t>& expectWorkspaces)	244	+ EXPECT_EQ(actual.outputFormat, expected.outputFormat);
		245	+ EXPECT_EQ(actual.batchNum, expected.batchNum);
		246	+ EXPECT_EQ(actual.channelNum, expected.channelNum);
		247	+ EXPECT_EQ(actual.inputSizeH, expected.inputSizeH);
		248	+ EXPECT_EQ(actual.inputSizeW, expected.inputSizeW);
		249	+ EXPECT_EQ(actual.outputSizeH, expected.outputSizeH);
		250	+ EXPECT_EQ(actual.outputSizeW, expected.outputSizeW);
		251	+ EXPECT_EQ(actual.srcChannelOffset, expected.srcChannelOffset);
		252	+
		253	+ // CscParam
		254	+ EXPECT_EQ(actual.cscParam.cscSwitch, expected.cscParam.cscSwitch);
		255	+ EXPECT_EQ(actual.cscParam.rbuvSwapSwitch, expected.cscParam.rbuvSwapSwitch);
		256	+ EXPECT_EQ(actual.cscParam.axSwapSwitch, expected.cscParam.axSwapSwitch);
		257	+ EXPECT_EQ(actual.cscParam.cscMatrix00, expected.cscParam.cscMatrix00);
		258	+ EXPECT_EQ(actual.cscParam.cscMatrix01, expected.cscParam.cscMatrix01);
		259	+ EXPECT_EQ(actual.cscParam.cscMatrix02, expected.cscParam.cscMatrix02);
		260	+ EXPECT_EQ(actual.cscParam.cscMatrix10, expected.cscParam.cscMatrix10);
		261	+ EXPECT_EQ(actual.cscParam.cscMatrix11, expected.cscParam.cscMatrix11);
		262	+ EXPECT_EQ(actual.cscParam.cscMatrix12, expected.cscParam.cscMatrix12);
		263	+ EXPECT_EQ(actual.cscParam.cscMatrix20, expected.cscParam.cscMatrix20);
		264	+ EXPECT_EQ(actual.cscParam.cscMatrix21, expected.cscParam.cscMatrix21);
		265	+ EXPECT_EQ(actual.cscParam.cscMatrix22, expected.cscParam.cscMatrix22);
		266	+ EXPECT_EQ(actual.cscParam.outBias0, expected.cscParam.outBias0);
		267	+ EXPECT_EQ(actual.cscParam.outBias1, expected.cscParam.outBias1);
		268	+ EXPECT_EQ(actual.cscParam.outBias2, expected.cscParam.outBias2);
		269	+ EXPECT_EQ(actual.cscParam.inBias0, expected.cscParam.inBias0);
		270	+ EXPECT_EQ(actual.cscParam.inBias1, expected.cscParam.inBias1);
		271	+ EXPECT_EQ(actual.cscParam.inBias2, expected.cscParam.inBias2);
		272	+
		273	+ // CropParam
		274	+ EXPECT_EQ(actual.cropParam.cropSwitch, expected.cropParam.cropSwitch);
		275	+ EXPECT_EQ(actual.cropParam.cropStartPosH, expected.cropParam.cropStartPosH);
		276	+ EXPECT_EQ(actual.cropParam.cropStartPosW, expected.cropParam.cropStartPosW);
		277	+ EXPECT_EQ(actual.cropParam.cropSizeH, expected.cropParam.cropSizeH);
		278	+ EXPECT_EQ(actual.cropParam.cropSizeW, expected.cropParam.cropSizeW);
		279	+
		280	+ // DtcParam
		281	+ EXPECT_EQ(actual.dtcParam.dtcPixelMeanChn0, expected.dtcParam.dtcPixelMeanChn0);
		282	+ EXPECT_EQ(actual.dtcParam.dtcPixelMeanChn1, expected.dtcParam.dtcPixelMeanChn1);
		283	+ EXPECT_EQ(actual.dtcParam.dtcPixelMeanChn2, expected.dtcParam.dtcPixelMeanChn2);
		284	+ EXPECT_EQ(actual.dtcParam.dtcPixelMeanChn3, expected.dtcParam.dtcPixelMeanChn3);
		285	+ EXPECT_FLOAT_EQ(actual.dtcParam.dtcPixelMinChn0, expected.dtcParam.dtcPixelMinChn0);
		286	+ EXPECT_FLOAT_EQ(actual.dtcParam.dtcPixelMinChn1, expected.dtcParam.dtcPixelMinChn1);
		287	+ EXPECT_FLOAT_EQ(actual.dtcParam.dtcPixelMinChn2, expected.dtcParam.dtcPixelMinChn2);
		288	+ EXPECT_FLOAT_EQ(actual.dtcParam.dtcPixelMinChn3, expected.dtcParam.dtcPixelMinChn3);
		289	+ EXPECT_FLOAT_EQ(actual.dtcParam.dtcPixelVarReciChn0, expected.dtcParam.dtcPixelVarReciChn0);
		290	+ EXPECT_FLOAT_EQ(actual.dtcParam.dtcPixelVarReciChn1, expected.dtcParam.dtcPixelVarReciChn1);
		291	+ EXPECT_FLOAT_EQ(actual.dtcParam.dtcPixelVarReciChn2, expected.dtcParam.dtcPixelVarReciChn2);
		292	+ EXPECT_FLOAT_EQ(actual.dtcParam.dtcPixelVarReciChn3, expected.dtcParam.dtcPixelVarReciChn3);
		293	+
		294	+ // PaddingParam
		295	+ EXPECT_EQ(actual.paddingParam.paddingSwitch, expected.paddingParam.paddingSwitch);
		296	+ EXPECT_EQ(actual.paddingParam.leftPaddingSize, expected.paddingParam.leftPaddingSize);
		297	+ EXPECT_EQ(actual.paddingParam.rightPaddingSize, expected.paddingParam.rightPaddingSize);
		298	+ EXPECT_EQ(actual.paddingParam.topPaddingSize, expected.paddingParam.topPaddingSize);
		299	+ EXPECT_EQ(actual.paddingParam.bottomPaddingSize, expected.paddingParam.bottomPaddingSize);
		300	+ EXPECT_FLOAT_EQ(actual.paddingParam.padValue, expected.paddingParam.padValue);
		301	+}
		302	+
		303	+void AippExecuteTestCase(const gert::TilingContextPara& tilingContextPara,
		304	+ ge::graphStatus expectResult,
		305	+ uint64_t expectTilingKey,
		306	+ const Aipp_Kernel::AippTilingData& expectTiling,
		307	+ const std::vector<size_t>& expectWorkspaces)
245	{	308	{
246	DO_TILING(tilingContextPara)	309	DO_TILING(tilingContextPara)
247		310
	@@ -262,50 +325,12 @@ void AippExecuteTestCase(const gert::TilingContextPara& tilingContextPara,
262		325
263	// check tiling key	326	// check tiling key
264	auto tilingKeyResult = tilingContext->GetTilingKey();	327	auto tilingKeyResult = tilingContext->GetTilingKey();
265	- // fixme
266	ASSERT_EQ(tilingKeyResult, expectTilingKey);	328	ASSERT_EQ(tilingKeyResult, expectTilingKey);
267		329
268	// check tiling data	330	// check tiling data
269	- auto buf = (Aipp_Kernel::AippTilingData*)tilingContext->GetRawTilingData()->GetData();	331	+ auto buf = reinterpret_cast<const Aipp_Kernel::AippTilingData*>(tilingContext->GetRawTilingData()->GetData());
270	- Aipp_Kernel::AippTilingData tilingParam = *buf;	332	+ Aipp_Kernel::AippTilingData actualTiling = *buf;
271	- std::stringstream ss;	333	+ CompareAippTilingData(actualTiling, expectTiling);
272	- ss << (int)tilingParam.imageFormat << " " << (int)tilingParam.outputFormat
273	- << " " << tilingParam.batchNum << " " << tilingParam.channelNum
274	- << " " << tilingParam.cscParam.rbuvSwapSwitch << " " << tilingParam.cscParam.axSwapSwitch
275	- << " " << tilingParam.inputSizeH << " " << tilingParam.inputSizeW
276	- << " " << tilingParam.cropParam.cropSwitch << " " << tilingParam.cropParam.cropStartPosH
277	- << " " << tilingParam.cropParam.cropStartPosW << " " << tilingParam.cropParam.cropSizeH
278	- << " " << tilingParam.cropParam.cropSizeW << " " << tilingParam.cscParam.cscSwitch
279	- << " " << tilingParam.cscParam.cscMatrix00 << " " << tilingParam.cscParam.cscMatrix01
280	- << " " << tilingParam.cscParam.cscMatrix02 << " " << tilingParam.cscParam.cscMatrix10
281	- << " " << tilingParam.cscParam.cscMatrix11 << " " << tilingParam.cscParam.cscMatrix12
282	- << " " << tilingParam.cscParam.cscMatrix20 << " " << tilingParam.cscParam.cscMatrix21
283	- << " " << tilingParam.cscParam.cscMatrix22 << " " << tilingParam.cscParam.outBias0
284	- << " " << tilingParam.cscParam.outBias1 << " " << tilingParam.cscParam.outBias2
285	- << " " << tilingParam.cscParam.inBias0 << " " << tilingParam.cscParam.inBias1
286	- << " " << tilingParam.cscParam.inBias2
287	- << " " << tilingParam.dtcParam.dtcPixelMeanChn0
288	- << " " << tilingParam.dtcParam.dtcPixelMeanChn1
289	- << " " << tilingParam.dtcParam.dtcPixelMeanChn2
290	- << " " << tilingParam.dtcParam.dtcPixelMeanChn3
291	- << " " << tilingParam.dtcParam.dtcPixelMinChn0
292	- << " " << tilingParam.dtcParam.dtcPixelMinChn1
293	- << " " << tilingParam.dtcParam.dtcPixelMinChn2
294	- << " " << tilingParam.dtcParam.dtcPixelMinChn3
295	- << " " << tilingParam.dtcParam.dtcPixelVarReciChn0
296	- << " " << tilingParam.dtcParam.dtcPixelVarReciChn1
297	- << " " << tilingParam.dtcParam.dtcPixelVarReciChn2
298	- << " " << tilingParam.dtcParam.dtcPixelVarReciChn3
299	- << " " << tilingParam.paddingParam.paddingSwitch
300	- << " " << tilingParam.paddingParam.leftPaddingSize
301	- << " " << tilingParam.paddingParam.rightPaddingSize
302	- << " " << tilingParam.paddingParam.topPaddingSize
303	- << " " << tilingParam.paddingParam.bottomPaddingSize
304	- << " " << tilingParam.paddingParam.padValue;
305	-
306	- auto tilingDataResult = ss.str();
307	- // fixme
308	- EXPECT_EQ(tilingDataResult, expectTilingData);
309	}	334	}
310		335
311	TEST_F(AippTiling, aipp_tiling_test_0)	336	TEST_F(AippTiling, aipp_tiling_test_0)
	@@ -319,10 +344,69 @@ TEST_F(AippTiling, aipp_tiling_test_0)
319	(R"({"aipp_mode":"static","input_format":"RGB888_U8"})"))},	344	(R"({"aipp_mode":"static","input_format":"RGB888_U8"})"))},
320	&compileInfo);	345	&compileInfo);
321		346
322	- uint64_t expectTilingKey = 1;	347	+ int64_t expectTilingKey = optiling::AIPP_RGB_PASS_THROUGH;
323	- string expectTilingData = "2 1 10 3 0 0 256 224 0 0 0 256 224 0 298 516 0 298 -100 -208 298 0 409 16 128 128 16 128 128 0 0 0 0 0 0 0 0 1 1 1 1 0 0 0 0 0 0";	348	+ Aipp_Kernel::AippTilingData expectTiling = {
		349	+ .imageFormat = 2,
		350	+ .outputFormat = 1,
		351	+ .batchNum = 10,
		352	+ .channelNum = 3,
		353	+ .inputSizeW = 224,
		354	+ .inputSizeH = 256,
		355	+ .outputSizeW = 224,
		356	+ .outputSizeH = 256,
		357	+ .srcChannelOffset = 0,
		358	+ .cscParam = {
		359	+ .cscSwitch = 0,
		360	+ .rbuvSwapSwitch = 0,
		361	+ .axSwapSwitch = 0,
		362	+ .cscMatrix00 = 256,
		363	+ .cscMatrix01 = 0,
		364	+ .cscMatrix02 = 0,
		365	+ .cscMatrix10 = 0,
		366	+ .cscMatrix11 = 256,
		367	+ .cscMatrix12 = 0,
		368	+ .cscMatrix20 = 0,
		369	+ .cscMatrix21 = 0,
		370	+ .cscMatrix22 = 256,
		371	+ .outBias0 = 0,
		372	+ .outBias1 = 0,
		373	+ .outBias2 = 0,
		374	+ .inBias0 = 0,
		375	+ .inBias1 = 0,
		376	+ .inBias2 = 0
		377	+ },
		378	+ .cropParam = {
		379	+ .cropSwitch = 0,
		380	+ .cropStartPosH = 0,
		381	+ .cropStartPosW = 0,
		382	+ .cropSizeH = 256,
		383	+ .cropSizeW = 224
		384	+ },
		385	+ .dtcParam = {
		386	+ .dtcPixelMeanChn0 = 0,
		387	+ .dtcPixelMeanChn1 = 0,
		388	+ .dtcPixelMeanChn2 = 0,
		389	+ .dtcPixelMeanChn3 = 0,
		390	+ .dtcPixelMinChn0 = 0,
		391	+ .dtcPixelMinChn1 = 0,
		392	+ .dtcPixelMinChn2 = 0,
		393	+ .dtcPixelMinChn3 = 0,
		394	+ .dtcPixelVarReciChn0 = 1,
		395	+ .dtcPixelVarReciChn1 = 1,
		396	+ .dtcPixelVarReciChn2 = 1,
		397	+ .dtcPixelVarReciChn3 = 1
		398	+ },
		399	+ .paddingParam = {
		400	+ .paddingSwitch = 0,
		401	+ .leftPaddingSize = 0,
		402	+ .rightPaddingSize = 0,
		403	+ .topPaddingSize = 0,
		404	+ .bottomPaddingSize = 0,
		405	+ .padValue = 0.0f
		406	+ }
		407	+ };
324	std::vector<size_t> expectWorkspaces = {16777216};	408	std::vector<size_t> expectWorkspaces = {16777216};
325	- AippExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);	409	+ AippExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTiling, expectWorkspaces);
326	}	410	}
327		411
328	TEST_F(AippTiling, aipp_tiling_test_1)	412	TEST_F(AippTiling, aipp_tiling_test_1)
	@@ -331,7 +415,7 @@ TEST_F(AippTiling, aipp_tiling_test_1)
331	gert::TilingContextPara tilingContextPara("Aipp",	415	gert::TilingContextPara tilingContextPara("Aipp",
332	{{{{100, 2560, 2240, 3}, {100, 2560, 2240, 3}}, ge::DT_UINT8, ge::FORMAT_NHWC}},	416	{{{{100, 2560, 2240, 3}, {100, 2560, 2240, 3}}, ge::DT_UINT8, ge::FORMAT_NHWC}},
333	{{{{100, 256, 224, 3}, {100, 256, 224, 3}}, ge::DT_FLOAT16, ge::FORMAT_NHWC}},	417	{{{{100, 256, 224, 3}, {100, 256, 224, 3}}, ge::DT_FLOAT16, ge::FORMAT_NHWC}},
334	- {gert::TilingContextPara::OpAttr("aipp_config_path",	418	+ {gert::TilingContextPara::OpAttr("aipp_config_path",
335	Ops::Cv::AnyValue::CreateFrom<string>\	419	Ops::Cv::AnyValue::CreateFrom<string>\
336	(R"({"aipp_mode":"static","crop":true,"crop_size_h":256,"crop_size_w":224,	420	(R"({"aipp_mode":"static","crop":true,"crop_size_h":256,"crop_size_w":224,
337	"input_format":"RGB888_U8","load_start_pos_h":16,"load_start_pos_w":16,	421	"input_format":"RGB888_U8","load_start_pos_h":16,"load_start_pos_w":16,
	@@ -344,11 +428,69 @@ TEST_F(AippTiling, aipp_tiling_test_1)
344	"output_bias_0":110, "output_bias_1":120, "output_bias_2":83})"))},	428	"output_bias_0":110, "output_bias_1":120, "output_bias_2":83})"))},
345	&compileInfo);	429	&compileInfo);
346		430
347	- uint64_t expectTilingKey = 3;	431	+ int64_t expectTilingKey = optiling::AIPP_RGB_TO_YUV;
348	- string expectTilingData = \	432	+ Aipp_Kernel::AippTilingData expectTiling = {
349	- "2 2 100 3 0 0 2560 2240 1 16 16 256 224 1 256 101 -202 301 -256 402 503 601 256 110 120 83 16 128 128 0 0 0 0 123.675 116.28 103.53 0 0.0171248 0.017507 0.0174292 1 0 0 0 0 0 0";	433	+ .imageFormat = 2,
		434	+ .outputFormat = 2,
		435	+ .batchNum = 100,
		436	+ .channelNum = 3,
		437	+ .inputSizeW = 2240,
		438	+ .inputSizeH = 2560,
		439	+ .outputSizeW = 224,
		440	+ .outputSizeH = 256,
		441	+ .srcChannelOffset = 0,
		442	+ .cscParam = {
		443	+ .cscSwitch = 1,
		444	+ .rbuvSwapSwitch = 0,
		445	+ .axSwapSwitch = 0,
		446	+ .cscMatrix00 = 256,
		447	+ .cscMatrix01 = 101,
		448	+ .cscMatrix02 = -202,
		449	+ .cscMatrix10 = 301,
		450	+ .cscMatrix11 = -256,
		451	+ .cscMatrix12 = 402,
		452	+ .cscMatrix20 = 503,
		453	+ .cscMatrix21 = 601,
		454	+ .cscMatrix22 = 256,
		455	+ .outBias0 = 110,
		456	+ .outBias1 = 120,
		457	+ .outBias2 = 83,
		458	+ .inBias0 = 0,
		459	+ .inBias1 = 0,
		460	+ .inBias2 = 0
		461	+ },
		462	+ .cropParam = {
		463	+ .cropSwitch = 1,
		464	+ .cropStartPosH = 16,
		465	+ .cropStartPosW = 16,
		466	+ .cropSizeH = 256,
		467	+ .cropSizeW = 224
		468	+ },
		469	+ .dtcParam = {
		470	+ .dtcPixelMeanChn0 = 0,
		471	+ .dtcPixelMeanChn1 = 0,
		472	+ .dtcPixelMeanChn2 = 0,
		473	+ .dtcPixelMeanChn3 = 0,
		474	+ .dtcPixelMinChn0 = 123.675f,
		475	+ .dtcPixelMinChn1 = 116.28f,
		476	+ .dtcPixelMinChn2 = 103.53f,
		477	+ .dtcPixelMinChn3 = 0,
		478	+ .dtcPixelVarReciChn0 = 0.017124753f,
		479	+ .dtcPixelVarReciChn1 = 0.017507f,
		480	+ .dtcPixelVarReciChn2 = 0.0174292f,
		481	+ .dtcPixelVarReciChn3 = 1
		482	+ },
		483	+ .paddingParam = {
		484	+ .paddingSwitch = 0,
		485	+ .leftPaddingSize = 0,
		486	+ .rightPaddingSize = 0,
		487	+ .topPaddingSize = 0,
		488	+ .bottomPaddingSize = 0,
		489	+ .padValue = 0.0f
		490	+ }
		491	+ };
350	std::vector<size_t> expectWorkspaces = {16777216};	492	std::vector<size_t> expectWorkspaces = {16777216};
351	- AippExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);	493	+ AippExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTiling, expectWorkspaces);
352	}	494	}
353		495
354	TEST_F(AippTiling, aipp_tiling_test_2)	496	TEST_F(AippTiling, aipp_tiling_test_2)
	@@ -363,10 +505,69 @@ TEST_F(AippTiling, aipp_tiling_test_2)
363	{gert::TilingContextPara::OpAttr("aipp_config_path", Ops::Cv::AnyValue::CreateFrom<string>(filePath))},	505	{gert::TilingContextPara::OpAttr("aipp_config_path", Ops::Cv::AnyValue::CreateFrom<string>(filePath))},
364	&compileInfo);	506	&compileInfo);
365		507
366	- uint64_t expectTilingKey = 1;	508	+ int64_t expectTilingKey = optiling::AIPP_RGB_PASS_THROUGH;
367	- string expectTilingData = "2 1 1 3 0 0 224 224 0 0 0 224 224 0 298 516 0 298 -100 -208 298 0 409 16 128 128 16 128 128 128 128 128 0 0 0 0 0 0.00781 0.00781 0.00781 1 0 0 0 0 0 0";	509	+ Aipp_Kernel::AippTilingData expectTiling = {
		510	+ .imageFormat = 2,
		511	+ .outputFormat = 1,
		512	+ .batchNum = 1,
		513	+ .channelNum = 3,
		514	+ .inputSizeW = 224,
		515	+ .inputSizeH = 224,
		516	+ .outputSizeW = 224,
		517	+ .outputSizeH = 224,
		518	+ .srcChannelOffset = 0,
		519	+ .cscParam = {
		520	+ .cscSwitch = 0,
		521	+ .rbuvSwapSwitch = 0,
		522	+ .axSwapSwitch = 0,
		523	+ .cscMatrix00 = 256,
		524	+ .cscMatrix01 = 0,
		525	+ .cscMatrix02 = 0,
		526	+ .cscMatrix10 = 0,
		527	+ .cscMatrix11 = 256,
		528	+ .cscMatrix12 = 0,
		529	+ .cscMatrix20 = 0,
		530	+ .cscMatrix21 = 0,
		531	+ .cscMatrix22 = 256,
		532	+ .outBias0 = 0,
		533	+ .outBias1 = 0,
		534	+ .outBias2 = 0,
		535	+ .inBias0 = 0,
		536	+ .inBias1 = 0,
		537	+ .inBias2 = 0
		538	+ },
		539	+ .cropParam = {
		540	+ .cropSwitch = 0,
		541	+ .cropStartPosH = 0,
		542	+ .cropStartPosW = 0,
		543	+ .cropSizeH = 224,
		544	+ .cropSizeW = 224
		545	+ },
		546	+ .dtcParam = {
		547	+ .dtcPixelMeanChn0 = 128,
		548	+ .dtcPixelMeanChn1 = 128,
		549	+ .dtcPixelMeanChn2 = 128,
		550	+ .dtcPixelMeanChn3 = 0,
		551	+ .dtcPixelMinChn0 = 0,
		552	+ .dtcPixelMinChn1 = 0,
		553	+ .dtcPixelMinChn2 = 0,
		554	+ .dtcPixelMinChn3 = 0,
		555	+ .dtcPixelVarReciChn0 = 0.00781f,
		556	+ .dtcPixelVarReciChn1 = 0.00781f,
		557	+ .dtcPixelVarReciChn2 = 0.00781f,
		558	+ .dtcPixelVarReciChn3 = 1
		559	+ },
		560	+ .paddingParam = {
		561	+ .paddingSwitch = 0,
		562	+ .leftPaddingSize = 0,
		563	+ .rightPaddingSize = 0,
		564	+ .topPaddingSize = 0,
		565	+ .bottomPaddingSize = 0,
		566	+ .padValue = 0.0f
		567	+ }
		568	+ };
368	std::vector<size_t> expectWorkspaces = {16777216};	569	std::vector<size_t> expectWorkspaces = {16777216};
369	- AippExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);	570	+ AippExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTiling, expectWorkspaces);
370	}	571	}
371		572
372	TEST_F(AippTiling, aipp_tiling_test_3)	573	TEST_F(AippTiling, aipp_tiling_test_3)
	@@ -381,10 +582,69 @@ TEST_F(AippTiling, aipp_tiling_test_3)
381	{gert::TilingContextPara::OpAttr("aipp_config_path", Ops::Cv::AnyValue::CreateFrom<string>(filePath))},	582	{gert::TilingContextPara::OpAttr("aipp_config_path", Ops::Cv::AnyValue::CreateFrom<string>(filePath))},
382	&compileInfo);	583	&compileInfo);
383		584
384	- uint64_t expectTilingKey = 3;	585	+ int64_t expectTilingKey = optiling::AIPP_RGB_TO_YUV;
385	- string expectTilingData = "2 1 1 3 0 0 1080 256 1 2 16 256 224 1 298 10 409 298 -100 -208 298 516 20 110 120 83 16 128 128 10 20 30 0 123.675 116.28 103.53 0 0.0171248 0.017507 0.0174292 1 0 0 0 0 0 0";	586	+ Aipp_Kernel::AippTilingData expectTiling = {
		587	+ .imageFormat = 2,
		588	+ .outputFormat = 1,
		589	+ .batchNum = 1,
		590	+ .channelNum = 3,
		591	+ .inputSizeW = 256,
		592	+ .inputSizeH = 1080,
		593	+ .outputSizeW = 256,
		594	+ .outputSizeH = 224,
		595	+ .srcChannelOffset = 0,
		596	+ .cscParam = {
		597	+ .cscSwitch = 1,
		598	+ .rbuvSwapSwitch = 0,
		599	+ .axSwapSwitch = 0,
		600	+ .cscMatrix00 = 298,
		601	+ .cscMatrix01 = 10,
		602	+ .cscMatrix02 = 409,
		603	+ .cscMatrix10 = 298,
		604	+ .cscMatrix11 = -100,
		605	+ .cscMatrix12 = -208,
		606	+ .cscMatrix20 = 298,
		607	+ .cscMatrix21 = 516,
		608	+ .cscMatrix22 = 20,
		609	+ .outBias0 = 110,
		610	+ .outBias1 = 120,
		611	+ .outBias2 = 83,
		612	+ .inBias0 = 0,
		613	+ .inBias1 = 0,
		614	+ .inBias2 = 0
		615	+ },
		616	+ .cropParam = {
		617	+ .cropSwitch = 1,
		618	+ .cropStartPosH = 2,
		619	+ .cropStartPosW = 16,
		620	+ .cropSizeH = 256,
		621	+ .cropSizeW = 224
		622	+ },
		623	+ .dtcParam = {
		624	+ .dtcPixelMeanChn0 = 10,
		625	+ .dtcPixelMeanChn1 = 20,
		626	+ .dtcPixelMeanChn2 = 30,
		627	+ .dtcPixelMeanChn3 = 0,
		628	+ .dtcPixelMinChn0 = 123.675f,
		629	+ .dtcPixelMinChn1 = 116.28f,
		630	+ .dtcPixelMinChn2 = 103.53f,
		631	+ .dtcPixelMinChn3 = 0,
		632	+ .dtcPixelVarReciChn0 = 0.017124753f,
		633	+ .dtcPixelVarReciChn1 = 0.017507004f,
		634	+ .dtcPixelVarReciChn2 = 0.017429193f,
		635	+ .dtcPixelVarReciChn3 = 1
		636	+ },
		637	+ .paddingParam = {
		638	+ .paddingSwitch = 0,
		639	+ .leftPaddingSize = 0,
		640	+ .rightPaddingSize = 0,
		641	+ .topPaddingSize = 0,
		642	+ .bottomPaddingSize = 0,
		643	+ .padValue = 0.0f
		644	+ }
		645	+ };
386	std::vector<size_t> expectWorkspaces = {16777216};	646	std::vector<size_t> expectWorkspaces = {16777216};
387	- AippExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);	647	+ AippExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTiling, expectWorkspaces);
388	}	648	}
389		649
390	TEST_F(AippTiling, aipp_tiling_test_4)	650	TEST_F(AippTiling, aipp_tiling_test_4)
	@@ -399,10 +659,69 @@ TEST_F(AippTiling, aipp_tiling_test_4)
399	{gert::TilingContextPara::OpAttr("aipp_config_path", Ops::Cv::AnyValue::CreateFrom<string>(filePath))},	659	{gert::TilingContextPara::OpAttr("aipp_config_path", Ops::Cv::AnyValue::CreateFrom<string>(filePath))},
400	&compileInfo);	660	&compileInfo);
401		661
402	- uint64_t expectTilingKey = 2;	662	+ int64_t expectTilingKey = optiling::AIPP_YUV_TO_RGB;
403	- string expectTilingData = "1 2 1 3 1 0 256 256 0 0 0 256 256 1 298 409 0 298 -208 -100 298 0 516 16 128 128 110 120 83 10 20 30 0 123.675 116.28 103.53 0 0.0171248 0.017507 0.0174292 1 0 0 0 0 0 0";	663	+ Aipp_Kernel::AippTilingData expectTiling = {
		664	+ .imageFormat = 1,
		665	+ .outputFormat = 2,
		666	+ .batchNum = 1,
		667	+ .channelNum = 3,
		668	+ .inputSizeW = 256,
		669	+ .inputSizeH = 256,
		670	+ .outputSizeW = 256,
		671	+ .outputSizeH = 256,
		672	+ .srcChannelOffset = 0,
		673	+ .cscParam = {
		674	+ .cscSwitch = 1,
		675	+ .rbuvSwapSwitch = 1,
		676	+ .axSwapSwitch = 0,
		677	+ .cscMatrix00 = 298,
		678	+ .cscMatrix01 = 409,
		679	+ .cscMatrix02 = 0,
		680	+ .cscMatrix10 = 298,
		681	+ .cscMatrix11 = -208,
		682	+ .cscMatrix12 = -100,
		683	+ .cscMatrix20 = 298,
		684	+ .cscMatrix21 = 0,
		685	+ .cscMatrix22 = 516,
		686	+ .outBias0 = 0,
		687	+ .outBias1 = 0,
		688	+ .outBias2 = 0,
		689	+ .inBias0 = 110,
		690	+ .inBias1 = 120,
		691	+ .inBias2 = 83
		692	+ },
		693	+ .cropParam = {
		694	+ .cropSwitch = 0,
		695	+ .cropStartPosH = 0,
		696	+ .cropStartPosW = 0,
		697	+ .cropSizeH = 256,
		698	+ .cropSizeW = 256
		699	+ },
		700	+ .dtcParam = {
		701	+ .dtcPixelMeanChn0 = 10,
		702	+ .dtcPixelMeanChn1 = 20,
		703	+ .dtcPixelMeanChn2 = 30,
		704	+ .dtcPixelMeanChn3 = 0,
		705	+ .dtcPixelMinChn0 = 123.675f,
		706	+ .dtcPixelMinChn1 = 116.28f,
		707	+ .dtcPixelMinChn2 = 103.53f,
		708	+ .dtcPixelMinChn3 = 0,
		709	+ .dtcPixelVarReciChn0 = 0.017124753f,
		710	+ .dtcPixelVarReciChn1 = 0.017507f,
		711	+ .dtcPixelVarReciChn2 = 0.0174292f,
		712	+ .dtcPixelVarReciChn3 = 1
		713	+ },
		714	+ .paddingParam = {
		715	+ .paddingSwitch = 0,
		716	+ .leftPaddingSize = 0,
		717	+ .rightPaddingSize = 0,
		718	+ .topPaddingSize = 0,
		719	+ .bottomPaddingSize = 0,
		720	+ .padValue = 0.0f
		721	+ }
		722	+ };
404	std::vector<size_t> expectWorkspaces = {16777216};	723	std::vector<size_t> expectWorkspaces = {16777216};
405	- AippExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);	724	+ AippExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTiling, expectWorkspaces);
406	}	725	}
407		726
408	TEST_F(AippTiling, aipp_tiling_test_5)	727	TEST_F(AippTiling, aipp_tiling_test_5)
	@@ -416,10 +735,10 @@ TEST_F(AippTiling, aipp_tiling_test_5)
416	(R"({"aipp_mode":"static","input_format":"RGB888_U8"})"))},	735	(R"({"aipp_mode":"static","input_format":"RGB888_U8"})"))},
417	&compileInfo);	736	&compileInfo);
418		737
419	- uint64_t expectTilingKey = 0;	738	+ int64_t expectTilingKey = optiling::AIPP_ERROR_TILINGKEY;
420	- string expectTilingData = "";	739	+ Aipp_Kernel::AippTilingData expectTiling = {}; // 不检查
421	std::vector<size_t> expectWorkspaces = {16777216};	740	std::vector<size_t> expectWorkspaces = {16777216};
422	- AippExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, expectTilingKey, expectTilingData, expectWorkspaces);	741	+ AippExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, expectTilingKey, expectTiling, expectWorkspaces);
423	}	742	}
424		743
425	TEST_F(AippTiling, aipp_tiling_test_6)	744	TEST_F(AippTiling, aipp_tiling_test_6)
	@@ -433,10 +752,10 @@ TEST_F(AippTiling, aipp_tiling_test_6)
433	(R"({"aipp_mode":"static","input_format":"YUV420SP_U8"})"))},	752	(R"({"aipp_mode":"static","input_format":"YUV420SP_U8"})"))},
434	&compileInfo);	753	&compileInfo);
435		754
436	- uint64_t expectTilingKey = 0;	755	+ int64_t expectTilingKey = optiling::AIPP_ERROR_TILINGKEY;
437	- string expectTilingData = "";	756	+ Aipp_Kernel::AippTilingData expectTiling = {}; // 不检查
438	std::vector<size_t> expectWorkspaces = {16777216};	757	std::vector<size_t> expectWorkspaces = {16777216};
439	- AippExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, expectTilingKey, expectTilingData, expectWorkspaces);	758	+ AippExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, expectTilingKey, expectTiling, expectWorkspaces);
440	}	759	}
441		760
442	TEST_F(AippTiling, aipp_tiling_test_7)	761	TEST_F(AippTiling, aipp_tiling_test_7)
	@@ -451,10 +770,10 @@ TEST_F(AippTiling, aipp_tiling_test_7)
451	"crop":true,"load_start_pos_h":222,"load_start_pos_w":21})"))},	770	"crop":true,"load_start_pos_h":222,"load_start_pos_w":21})"))},
452	&compileInfo);	771	&compileInfo);
453		772
454	- uint64_t expectTilingKey = 0;	773	+ int64_t expectTilingKey = optiling::AIPP_ERROR_TILINGKEY;
455	- string expectTilingData = "";	774	+ Aipp_Kernel::AippTilingData expectTiling = {}; // 不检查
456	std::vector<size_t> expectWorkspaces = {16777216};	775	std::vector<size_t> expectWorkspaces = {16777216};
457	- AippExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, expectTilingKey, expectTilingData, expectWorkspaces);	776	+ AippExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, expectTilingKey, expectTiling, expectWorkspaces);
458	}	777	}
459		778
460	TEST_F(AippTiling, aipp_tiling_test_8)	779	TEST_F(AippTiling, aipp_tiling_test_8)
	@@ -468,10 +787,10 @@ TEST_F(AippTiling, aipp_tiling_test_8)
468	(R"({"aipp_mode":"static","input_format":"RGB888_U8"})"))},	787	(R"({"aipp_mode":"static","input_format":"RGB888_U8"})"))},
469	&compileInfo);	788	&compileInfo);
470		789
471	- uint64_t expectTilingKey = 1;	790	+ int64_t expectTilingKey = optiling::AIPP_RGB_PASS_THROUGH;
472	- string expectTilingData = "";	791	+ Aipp_Kernel::AippTilingData expectTiling = {}; // 不检查
473	std::vector<size_t> expectWorkspaces = {16777216};	792	std::vector<size_t> expectWorkspaces = {16777216};
474	- AippExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, expectTilingKey, expectTilingData, expectWorkspaces);	793	+ AippExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, expectTilingKey, expectTiling, expectWorkspaces);
475	}	794	}
476		795
477	TEST_F(AippTiling, aipp_tiling_test_9)	796	TEST_F(AippTiling, aipp_tiling_test_9)
	@@ -486,10 +805,69 @@ TEST_F(AippTiling, aipp_tiling_test_9)
486	{gert::TilingContextPara::OpAttr("aipp_config_path", Ops::Cv::AnyValue::CreateFrom<string>(filePath))},	805	{gert::TilingContextPara::OpAttr("aipp_config_path", Ops::Cv::AnyValue::CreateFrom<string>(filePath))},
487	&compileInfo);	806	&compileInfo);
488		807
489	- uint64_t expectTilingKey = 3;	808	+ int64_t expectTilingKey = optiling::AIPP_RGB_TO_YUV;
490	- string expectTilingData = "3 1 1 4 1 1 224 224 0 0 0 224 224 1 -202 101 298 402 -256 301 256 601 503 110 120 83 16 128 128 128 128 128 0 0 0 0 0 0.00781 0.00781 0.00781 1 0 0 0 0 0 0";	809	+ Aipp_Kernel::AippTilingData expectTiling = {
		810	+ .imageFormat = 3,
		811	+ .outputFormat = 1,
		812	+ .batchNum = 1,
		813	+ .channelNum = 4,
		814	+ .inputSizeW = 224,
		815	+ .inputSizeH = 224,
		816	+ .outputSizeW = 224,
		817	+ .outputSizeH = 224,
		818	+ .srcChannelOffset = 1,
		819	+ .cscParam = {
		820	+ .cscSwitch = 1,
		821	+ .rbuvSwapSwitch = 1,
		822	+ .axSwapSwitch = 1,
		823	+ .cscMatrix00 = -202,
		824	+ .cscMatrix01 = 101,
		825	+ .cscMatrix02 = 298,
		826	+ .cscMatrix10 = 402,
		827	+ .cscMatrix11 = -256,
		828	+ .cscMatrix12 = 301,
		829	+ .cscMatrix20 = 256,
		830	+ .cscMatrix21 = 601,
		831	+ .cscMatrix22 = 503,
		832	+ .outBias0 = 110,
		833	+ .outBias1 = 120,
		834	+ .outBias2 = 83,
		835	+ .inBias0 = 0,
		836	+ .inBias1 = 0,
		837	+ .inBias2 = 0
		838	+ },
		839	+ .cropParam = {
		840	+ .cropSwitch = 0,
		841	+ .cropStartPosH = 0,
		842	+ .cropStartPosW = 0,
		843	+ .cropSizeH = 224,
		844	+ .cropSizeW = 224
		845	+ },
		846	+ .dtcParam = {
		847	+ .dtcPixelMeanChn0 = 128,
		848	+ .dtcPixelMeanChn1 = 128,
		849	+ .dtcPixelMeanChn2 = 128,
		850	+ .dtcPixelMeanChn3 = 0,
		851	+ .dtcPixelMinChn0 = 0,
		852	+ .dtcPixelMinChn1 = 0,
		853	+ .dtcPixelMinChn2 = 0,
		854	+ .dtcPixelMinChn3 = 0,
		855	+ .dtcPixelVarReciChn0 = 0.00781f,
		856	+ .dtcPixelVarReciChn1 = 0.00781f,
		857	+ .dtcPixelVarReciChn2 = 0.00781f,
		858	+ .dtcPixelVarReciChn3 = 1
		859	+ },
		860	+ .paddingParam = {
		861	+ .paddingSwitch = 0,
		862	+ .leftPaddingSize = 0,
		863	+ .rightPaddingSize = 0,
		864	+ .topPaddingSize = 0,
		865	+ .bottomPaddingSize = 0,
		866	+ .padValue = 0.0f
		867	+ }
		868	+ };
491	std::vector<size_t> expectWorkspaces = {16777216};	869	std::vector<size_t> expectWorkspaces = {16777216};
492	- AippExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);	870	+ AippExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTiling, expectWorkspaces);
493	}	871	}
494		872
495	TEST_F(AippTiling, aipp_tiling_test_10)	873	TEST_F(AippTiling, aipp_tiling_test_10)
	@@ -504,10 +882,10 @@ TEST_F(AippTiling, aipp_tiling_test_10)
504	{gert::TilingContextPara::OpAttr("aipp_config_path", Ops::Cv::AnyValue::CreateFrom<string>(filePath))},	882	{gert::TilingContextPara::OpAttr("aipp_config_path", Ops::Cv::AnyValue::CreateFrom<string>(filePath))},
505	&compileInfo);	883	&compileInfo);
506		884
507	- uint64_t expectTilingKey = 0;	885	+ int64_t expectTilingKey = optiling::AIPP_ERROR_TILINGKEY;
508	- string expectTilingData = "";	886	+ Aipp_Kernel::AippTilingData expectTiling = {}; // 不检查
509	std::vector<size_t> expectWorkspaces = {16777216};	887	std::vector<size_t> expectWorkspaces = {16777216};
510	- AippExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, expectTilingKey, expectTilingData, expectWorkspaces);	888	+ AippExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, expectTilingKey, expectTiling, expectWorkspaces);
511	}	889	}
512		890
513	TEST_F(AippTiling, aipp_tiling_test_11)	891	TEST_F(AippTiling, aipp_tiling_test_11)
	@@ -521,10 +899,10 @@ TEST_F(AippTiling, aipp_tiling_test_11)
521	(R"({"aipp_mode":"static","input_format":"RGB888_U8"})"))},	899	(R"({"aipp_mode":"static","input_format":"RGB888_U8"})"))},
522	&compileInfo);	900	&compileInfo);
523		901
524	- uint64_t expectTilingKey = 0;	902	+ int64_t expectTilingKey = optiling::AIPP_ERROR_TILINGKEY;
525	- string expectTilingData = "";	903	+ Aipp_Kernel::AippTilingData expectTiling = {}; // 不检查
526	std::vector<size_t> expectWorkspaces = {16777216};	904	std::vector<size_t> expectWorkspaces = {16777216};
527	- AippExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, expectTilingKey, expectTilingData, expectWorkspaces);	905	+ AippExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, expectTilingKey, expectTiling, expectWorkspaces);
528	}	906	}
529		907
530	TEST_F(AippTiling, aipp_tiling_padding_test_1)	908	TEST_F(AippTiling, aipp_tiling_padding_test_1)
	@@ -539,10 +917,69 @@ TEST_F(AippTiling, aipp_tiling_padding_test_1)
539	{gert::TilingContextPara::OpAttr("aipp_config_path", Ops::Cv::AnyValue::CreateFrom<string>(filePath))},	917	{gert::TilingContextPara::OpAttr("aipp_config_path", Ops::Cv::AnyValue::CreateFrom<string>(filePath))},
540	&compileInfo);	918	&compileInfo);
541		919
542	- uint64_t expectTilingKey = 1;	920	+ int64_t expectTilingKey = optiling::AIPP_RGB_PASS_THROUGH;
543	- string expectTilingData = "2 1 1 3 0 0 224 224 0 0 0 224 224 0 298 516 0 298 -100 -208 298 0 409 16 128 128 16 128 128 128 128 128 0 0 0 0 0 0.00781 0.00781 0.00781 1 1 10 10 10 10 0";	921	+ Aipp_Kernel::AippTilingData expectTiling = {
		922	+ .imageFormat = 2,
		923	+ .outputFormat = 1,
		924	+ .batchNum = 1,
		925	+ .channelNum = 3,
		926	+ .inputSizeW = 224,
		927	+ .inputSizeH = 224,
		928	+ .outputSizeW = 244,
		929	+ .outputSizeH = 244,
		930	+ .srcChannelOffset = 0,
		931	+ .cscParam = {
		932	+ .cscSwitch = 0,
		933	+ .rbuvSwapSwitch = 0,
		934	+ .axSwapSwitch = 0,
		935	+ .cscMatrix00 = 256,
		936	+ .cscMatrix01 = 0,
		937	+ .cscMatrix02 = 0,
		938	+ .cscMatrix10 = 0,
		939	+ .cscMatrix11 = 256,
		940	+ .cscMatrix12 = 0,
		941	+ .cscMatrix20 = 0,
		942	+ .cscMatrix21 = 0,
		943	+ .cscMatrix22 = 256,
		944	+ .outBias0 = 0,
		945	+ .outBias1 = 0,
		946	+ .outBias2 = 0,
		947	+ .inBias0 = 0,
		948	+ .inBias1 = 0,
		949	+ .inBias2 = 0
		950	+ },
		951	+ .cropParam = {
		952	+ .cropSwitch = 0,
		953	+ .cropStartPosH = 0,
		954	+ .cropStartPosW = 0,
		955	+ .cropSizeH = 224,
		956	+ .cropSizeW = 224
		957	+ },
		958	+ .dtcParam = {
		959	+ .dtcPixelMeanChn0 = 128,
		960	+ .dtcPixelMeanChn1 = 128,
		961	+ .dtcPixelMeanChn2 = 128,
		962	+ .dtcPixelMeanChn3 = 0,
		963	+ .dtcPixelMinChn0 = 0,
		964	+ .dtcPixelMinChn1 = 0,
		965	+ .dtcPixelMinChn2 = 0,
		966	+ .dtcPixelMinChn3 = 0,
		967	+ .dtcPixelVarReciChn0 = 0.00781f,
		968	+ .dtcPixelVarReciChn1 = 0.00781f,
		969	+ .dtcPixelVarReciChn2 = 0.00781f,
		970	+ .dtcPixelVarReciChn3 = 1
		971	+ },
		972	+ .paddingParam = {
		973	+ .paddingSwitch = 1,
		974	+ .leftPaddingSize = 10,
		975	+ .rightPaddingSize = 10,
		976	+ .topPaddingSize = 10,
		977	+ .bottomPaddingSize = 10,
		978	+ .padValue = 0.0f
		979	+ }
		980	+ };
544	std::vector<size_t> expectWorkspaces = {16777216};	981	std::vector<size_t> expectWorkspaces = {16777216};
545	- AippExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);	982	+ AippExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTiling, expectWorkspaces);
546	}	983	}
547		984
548	TEST_F(AippTiling, aipp_tiling_padding_test_2)	985	TEST_F(AippTiling, aipp_tiling_padding_test_2)
	@@ -557,10 +994,69 @@ TEST_F(AippTiling, aipp_tiling_padding_test_2)
557	{gert::TilingContextPara::OpAttr("aipp_config_path", Ops::Cv::AnyValue::CreateFrom<string>(filePath))},	994	{gert::TilingContextPara::OpAttr("aipp_config_path", Ops::Cv::AnyValue::CreateFrom<string>(filePath))},
558	&compileInfo);	995	&compileInfo);
559		996
560	- uint64_t expectTilingKey = 1;	997	+ int64_t expectTilingKey = optiling::AIPP_RGB_PASS_THROUGH;
561	- string expectTilingData = "2 1 1 3 0 0 224 224 0 0 0 224 224 0 298 516 0 298 -100 -208 298 0 409 16 128 128 16 128 128 128 128 128 0 0 0 0 0 0.00781 0.00781 0.00781 1 1 20 0 0 0 0";	998	+ Aipp_Kernel::AippTilingData expectTiling = {
		999	+ .imageFormat = 2,
		1000	+ .outputFormat = 1,
		1001	+ .batchNum = 1,
		1002	+ .channelNum = 3,
		1003	+ .inputSizeW = 224,
		1004	+ .inputSizeH = 224,
		1005	+ .outputSizeW = 244,
		1006	+ .outputSizeH = 224,
		1007	+ .srcChannelOffset = 0,
		1008	+ .cscParam = {
		1009	+ .cscSwitch = 0,
		1010	+ .rbuvSwapSwitch = 0,
		1011	+ .axSwapSwitch = 0,
		1012	+ .cscMatrix00 = 256,
		1013	+ .cscMatrix01 = 0,
		1014	+ .cscMatrix02 = 0,
		1015	+ .cscMatrix10 = 0,
		1016	+ .cscMatrix11 = 256,
		1017	+ .cscMatrix12 = 0,
		1018	+ .cscMatrix20 = 0,
		1019	+ .cscMatrix21 = 0,
		1020	+ .cscMatrix22 = 256,
		1021	+ .outBias0 = 0,
		1022	+ .outBias1 = 0,
		1023	+ .outBias2 = 0,
		1024	+ .inBias0 = 0,
		1025	+ .inBias1 = 0,
		1026	+ .inBias2 = 0
		1027	+ },
		1028	+ .cropParam = {
		1029	+ .cropSwitch = 0,
		1030	+ .cropStartPosH = 0,
		1031	+ .cropStartPosW = 0,
		1032	+ .cropSizeH = 224,
		1033	+ .cropSizeW = 224
		1034	+ },
		1035	+ .dtcParam = {
		1036	+ .dtcPixelMeanChn0 = 128,
		1037	+ .dtcPixelMeanChn1 = 128,
		1038	+ .dtcPixelMeanChn2 = 128,
		1039	+ .dtcPixelMeanChn3 = 0,
		1040	+ .dtcPixelMinChn0 = 0,
		1041	+ .dtcPixelMinChn1 = 0,
		1042	+ .dtcPixelMinChn2 = 0,
		1043	+ .dtcPixelMinChn3 = 0,
		1044	+ .dtcPixelVarReciChn0 = 0.00781f,
		1045	+ .dtcPixelVarReciChn1 = 0.00781f,
		1046	+ .dtcPixelVarReciChn2 = 0.00781f,
		1047	+ .dtcPixelVarReciChn3 = 1
		1048	+ },
		1049	+ .paddingParam = {
		1050	+ .paddingSwitch = 1,
		1051	+ .leftPaddingSize = 20,
		1052	+ .rightPaddingSize = 0,
		1053	+ .topPaddingSize = 0,
		1054	+ .bottomPaddingSize = 0,
		1055	+ .padValue = 0.0f
		1056	+ }
		1057	+ };
562	std::vector<size_t> expectWorkspaces = {16777216};	1058	std::vector<size_t> expectWorkspaces = {16777216};
563	- AippExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);	1059	+ AippExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTiling, expectWorkspaces);
564	}	1060	}
565		1061
566	TEST_F(AippTiling, aipp_tiling_padding_test_3)	1062	TEST_F(AippTiling, aipp_tiling_padding_test_3)
	@@ -575,10 +1071,69 @@ TEST_F(AippTiling, aipp_tiling_padding_test_3)
575	{gert::TilingContextPara::OpAttr("aipp_config_path", Ops::Cv::AnyValue::CreateFrom<string>(filePath))},	1071	{gert::TilingContextPara::OpAttr("aipp_config_path", Ops::Cv::AnyValue::CreateFrom<string>(filePath))},
576	&compileInfo);	1072	&compileInfo);
577		1073
578	- uint64_t expectTilingKey = 2;	1074	+ int64_t expectTilingKey = optiling::AIPP_YUV_TO_RGB;
579	- string expectTilingData = "1 1 1 3 0 0 256 256 0 0 0 256 256 1 298 0 409 298 -100 -208 298 516 0 16 128 128 110 120 83 128 128 128 0 123.675 116.28 103.53 0 0.0171248 0.017507 0.0174292 1 1 20 0 0 0 0";	1075	+ Aipp_Kernel::AippTilingData expectTiling = {
		1076	+ .imageFormat = 1,
		1077	+ .outputFormat = 1,
		1078	+ .batchNum = 1,
		1079	+ .channelNum = 3,
		1080	+ .inputSizeW = 256,
		1081	+ .inputSizeH = 256,
		1082	+ .outputSizeW = 276,
		1083	+ .outputSizeH = 256,
		1084	+ .srcChannelOffset = 0,
		1085	+ .cscParam = {
		1086	+ .cscSwitch = 1,
		1087	+ .rbuvSwapSwitch = 0,
		1088	+ .axSwapSwitch = 0,
		1089	+ .cscMatrix00 = 298,
		1090	+ .cscMatrix01 = 0,
		1091	+ .cscMatrix02 = 409,
		1092	+ .cscMatrix10 = 298,
		1093	+ .cscMatrix11 = -100,
		1094	+ .cscMatrix12 = -208,
		1095	+ .cscMatrix20 = 298,
		1096	+ .cscMatrix21 = 516,
		1097	+ .cscMatrix22 = 0,
		1098	+ .outBias0 = 0,
		1099	+ .outBias1 = 0,
		1100	+ .outBias2 = 0,
		1101	+ .inBias0 = 110,
		1102	+ .inBias1 = 120,
		1103	+ .inBias2 = 83
		1104	+ },
		1105	+ .cropParam = {
		1106	+ .cropSwitch = 0,
		1107	+ .cropStartPosH = 0,
		1108	+ .cropStartPosW = 0,
		1109	+ .cropSizeH = 256,
		1110	+ .cropSizeW = 256
		1111	+ },
		1112	+ .dtcParam = {
		1113	+ .dtcPixelMeanChn0 = 128,
		1114	+ .dtcPixelMeanChn1 = 128,
		1115	+ .dtcPixelMeanChn2 = 128,
		1116	+ .dtcPixelMeanChn3 = 0,
		1117	+ .dtcPixelMinChn0 = 123.675f,
		1118	+ .dtcPixelMinChn1 = 116.28f,
		1119	+ .dtcPixelMinChn2 = 103.53f,
		1120	+ .dtcPixelMinChn3 = 0,
		1121	+ .dtcPixelVarReciChn0 = 0.017124753f,
		1122	+ .dtcPixelVarReciChn1 = 0.017507f,
		1123	+ .dtcPixelVarReciChn2 = 0.0174292f,
		1124	+ .dtcPixelVarReciChn3 = 1
		1125	+ },
		1126	+ .paddingParam = {
		1127	+ .paddingSwitch = 1,
		1128	+ .leftPaddingSize = 20,
		1129	+ .rightPaddingSize = 0,
		1130	+ .topPaddingSize = 0,
		1131	+ .bottomPaddingSize = 0,
		1132	+ .padValue = 0.0f
		1133	+ }
		1134	+ };
580	std::vector<size_t> expectWorkspaces = {16777216};	1135	std::vector<size_t> expectWorkspaces = {16777216};
581	- AippExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);	1136	+ AippExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTiling, expectWorkspaces);
582	}	1137	}
583		1138
584	TEST_F(AippTiling, aipp_tiling_padding_test_4)	1139	TEST_F(AippTiling, aipp_tiling_padding_test_4)
	@@ -593,10 +1148,69 @@ TEST_F(AippTiling, aipp_tiling_padding_test_4)
593	{gert::TilingContextPara::OpAttr("aipp_config_path", Ops::Cv::AnyValue::CreateFrom<string>(filePath))},	1148	{gert::TilingContextPara::OpAttr("aipp_config_path", Ops::Cv::AnyValue::CreateFrom<string>(filePath))},
594	&compileInfo);	1149	&compileInfo);
595		1150
596	- uint64_t expectTilingKey = 3;	1151	+ int64_t expectTilingKey = optiling::AIPP_RGB_TO_YUV;
597	- string expectTilingData = "2 1 1 3 0 0 1080 256 1 2 16 224 224 1 298 10 409 298 -100 -208 298 516 20 16 128 128 110 120 83 10 20 30 0 123.675 116.28 103.53 0 0.0171248 0.017507 0.0174292 1 1 10 10 5 5 10.5";	1152	+ Aipp_Kernel::AippTilingData expectTiling = {
		1153	+ .imageFormat = 2,
		1154	+ .outputFormat = 1,
		1155	+ .batchNum = 1,
		1156	+ .channelNum = 3,
		1157	+ .inputSizeW = 256,
		1158	+ .inputSizeH = 1080,
		1159	+ .outputSizeW = 244,
		1160	+ .outputSizeH = 234,
		1161	+ .srcChannelOffset = 0,
		1162	+ .cscParam = {
		1163	+ .cscSwitch = 1,
		1164	+ .rbuvSwapSwitch = 0,
		1165	+ .axSwapSwitch = 0,
		1166	+ .cscMatrix00 = 298,
		1167	+ .cscMatrix01 = 10,
		1168	+ .cscMatrix02 = 409,
		1169	+ .cscMatrix10 = 298,
		1170	+ .cscMatrix11 = -100,
		1171	+ .cscMatrix12 = -208,
		1172	+ .cscMatrix20 = 298,
		1173	+ .cscMatrix21 = 516,
		1174	+ .cscMatrix22 = 20,
		1175	+ .outBias0 = 16,
		1176	+ .outBias1 = 128,
		1177	+ .outBias2 = 128,
		1178	+ .inBias0 = 0,
		1179	+ .inBias1 = 0,
		1180	+ .inBias2 = 0
		1181	+ },
		1182	+ .cropParam = {
		1183	+ .cropSwitch = 1,
		1184	+ .cropStartPosH = 2,
		1185	+ .cropStartPosW = 16,
		1186	+ .cropSizeH = 224,
		1187	+ .cropSizeW = 224
		1188	+ },
		1189	+ .dtcParam = {
		1190	+ .dtcPixelMeanChn0 = 10,
		1191	+ .dtcPixelMeanChn1 = 20,
		1192	+ .dtcPixelMeanChn2 = 30,
		1193	+ .dtcPixelMeanChn3 = 0,
		1194	+ .dtcPixelMinChn0 = 123.675f,
		1195	+ .dtcPixelMinChn1 = 116.28f,
		1196	+ .dtcPixelMinChn2 = 103.53f,
		1197	+ .dtcPixelMinChn3 = 0,
		1198	+ .dtcPixelVarReciChn0 = 0.017124753f,
		1199	+ .dtcPixelVarReciChn1 = 0.017507004f,
		1200	+ .dtcPixelVarReciChn2 = 0.017429193f,
		1201	+ .dtcPixelVarReciChn3 = 1
		1202	+ },
		1203	+ .paddingParam = {
		1204	+ .paddingSwitch = 1,
		1205	+ .leftPaddingSize = 10,
		1206	+ .rightPaddingSize = 10,
		1207	+ .topPaddingSize = 5,
		1208	+ .bottomPaddingSize = 5,
		1209	+ .padValue = 10.5f
		1210	+ }
		1211	+ };
598	std::vector<size_t> expectWorkspaces = {16777216};	1212	std::vector<size_t> expectWorkspaces = {16777216};
599	- AippExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);	1213	+ AippExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTiling, expectWorkspaces);
600	}	1214	}
601		1215
602	TEST_F(AippTiling, aipp_tiling_padding_test_5)	1216	TEST_F(AippTiling, aipp_tiling_padding_test_5)
	@@ -611,10 +1225,69 @@ TEST_F(AippTiling, aipp_tiling_padding_test_5)
611	{gert::TilingContextPara::OpAttr("aipp_config_path", Ops::Cv::AnyValue::CreateFrom<string>(filePath))},	1225	{gert::TilingContextPara::OpAttr("aipp_config_path", Ops::Cv::AnyValue::CreateFrom<string>(filePath))},
612	&compileInfo);	1226	&compileInfo);
613		1227
614	- uint64_t expectTilingKey = 1;	1228	+ int64_t expectTilingKey = optiling::AIPP_RGB_PASS_THROUGH;
615	- string expectTilingData = "2 1 1 3 0 0 224 224 0 0 0 224 224 0 298 516 0 298 -100 -208 298 0 409 16 128 128 16 128 128 128 128 128 0 0 0 0 0 0.00781 0.00781 0.00781 1 1 32 32 32 32 255";	1229	+ Aipp_Kernel::AippTilingData expectTiling = {
		1230	+ .imageFormat = 2,
		1231	+ .outputFormat = 1,
		1232	+ .batchNum = 1,
		1233	+ .channelNum = 3,
		1234	+ .inputSizeW = 224,
		1235	+ .inputSizeH = 224,
		1236	+ .outputSizeW = 288,
		1237	+ .outputSizeH = 288,
		1238	+ .srcChannelOffset = 0,
		1239	+ .cscParam = {
		1240	+ .cscSwitch = 0,
		1241	+ .rbuvSwapSwitch = 0,
		1242	+ .axSwapSwitch = 0,
		1243	+ .cscMatrix00 = 256,
		1244	+ .cscMatrix01 = 0,
		1245	+ .cscMatrix02 = 0,
		1246	+ .cscMatrix10 = 0,
		1247	+ .cscMatrix11 = 256,
		1248	+ .cscMatrix12 = 0,
		1249	+ .cscMatrix20 = 0,
		1250	+ .cscMatrix21 = 0,
		1251	+ .cscMatrix22 = 256,
		1252	+ .outBias0 = 0,
		1253	+ .outBias1 = 0,
		1254	+ .outBias2 = 0,
		1255	+ .inBias0 = 0,
		1256	+ .inBias1 = 0,
		1257	+ .inBias2 = 0
		1258	+ },
		1259	+ .cropParam = {
		1260	+ .cropSwitch = 0,
		1261	+ .cropStartPosH = 0,
		1262	+ .cropStartPosW = 0,
		1263	+ .cropSizeH = 224,
		1264	+ .cropSizeW = 224
		1265	+ },
		1266	+ .dtcParam = {
		1267	+ .dtcPixelMeanChn0 = 128,
		1268	+ .dtcPixelMeanChn1 = 128,
		1269	+ .dtcPixelMeanChn2 = 128,
		1270	+ .dtcPixelMeanChn3 = 0,
		1271	+ .dtcPixelMinChn0 = 0,
		1272	+ .dtcPixelMinChn1 = 0,
		1273	+ .dtcPixelMinChn2 = 0,
		1274	+ .dtcPixelMinChn3 = 0,
		1275	+ .dtcPixelVarReciChn0 = 0.00781f,
		1276	+ .dtcPixelVarReciChn1 = 0.00781f,
		1277	+ .dtcPixelVarReciChn2 = 0.00781f,
		1278	+ .dtcPixelVarReciChn3 = 1
		1279	+ },
		1280	+ .paddingParam = {
		1281	+ .paddingSwitch = 1,
		1282	+ .leftPaddingSize = 32,
		1283	+ .rightPaddingSize = 32,
		1284	+ .topPaddingSize = 32,
		1285	+ .bottomPaddingSize = 32,
		1286	+ .padValue = 255.0f
		1287	+ }
		1288	+ };
616	std::vector<size_t> expectWorkspaces = {16777216};	1289	std::vector<size_t> expectWorkspaces = {16777216};
617	- AippExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);	1290	+ AippExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTiling, expectWorkspaces);
618	}	1291	}
619		1292
620	TEST_F(AippTiling, aipp_tiling_padding_test_6)	1293	TEST_F(AippTiling, aipp_tiling_padding_test_6)
	@@ -628,10 +1301,10 @@ TEST_F(AippTiling, aipp_tiling_padding_test_6)
628	(R"({"aipp_mode":"static","input_format":"RGB888_U8", "padding":"aaa"})"))},	1301	(R"({"aipp_mode":"static","input_format":"RGB888_U8", "padding":"aaa"})"))},
629	&compileInfo);	1302	&compileInfo);
630		1303
631	- uint64_t expectTilingKey = 1;	1304	+ int64_t expectTilingKey = optiling::AIPP_RGB_PASS_THROUGH;
632	- string expectTilingData = "";	1305	+ Aipp_Kernel::AippTilingData expectTiling = {}; // 不检查
633	std::vector<size_t> expectWorkspaces = {16777216};	1306	std::vector<size_t> expectWorkspaces = {16777216};
634	- AippExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, expectTilingKey, expectTilingData, expectWorkspaces);	1307	+ AippExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, expectTilingKey, expectTiling, expectWorkspaces);
635	}	1308	}
636		1309
637	TEST_F(AippTiling, aipp_tiling_padding_test_7)	1310	TEST_F(AippTiling, aipp_tiling_padding_test_7)
	@@ -645,10 +1318,10 @@ TEST_F(AippTiling, aipp_tiling_padding_test_7)
645	(R"({"aipp_mode":"static","input_format":"RGB888_U8","padding":"true","left_padding_size":"33"})"))},	1318	(R"({"aipp_mode":"static","input_format":"RGB888_U8","padding":"true","left_padding_size":"33"})"))},
646	&compileInfo);	1319	&compileInfo);
647		1320
648	- uint64_t expectTilingKey = 1;	1321	+ int64_t expectTilingKey = optiling::AIPP_RGB_PASS_THROUGH;
649	- string expectTilingData = "";	1322	+ Aipp_Kernel::AippTilingData expectTiling = {}; // 不检查
650	std::vector<size_t> expectWorkspaces = {16777216};	1323	std::vector<size_t> expectWorkspaces = {16777216};
651	- AippExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, expectTilingKey, expectTilingData, expectWorkspaces);	1324	+ AippExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, expectTilingKey, expectTiling, expectWorkspaces);
652	}	1325	}
653		1326
654	TEST_F(AippTiling, aipp_tiling_padding_test_8)	1327	TEST_F(AippTiling, aipp_tiling_padding_test_8)
	@@ -662,10 +1335,10 @@ TEST_F(AippTiling, aipp_tiling_padding_test_8)
662	(R"({"aipp_mode":"static","input_format":"RGB888_U8","padding":"true","left_padding_size":"10","padding_value":"65506"})"))},	1335	(R"({"aipp_mode":"static","input_format":"RGB888_U8","padding":"true","left_padding_size":"10","padding_value":"65506"})"))},
663	&compileInfo);	1336	&compileInfo);
664		1337
665	- uint64_t expectTilingKey = 1;	1338	+ int64_t expectTilingKey = optiling::AIPP_RGB_PASS_THROUGH;
666	- string expectTilingData = "";	1339	+ Aipp_Kernel::AippTilingData expectTiling = {}; // 不检查
667	std::vector<size_t> expectWorkspaces = {16777216};	1340	std::vector<size_t> expectWorkspaces = {16777216};
668	- AippExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, expectTilingKey, expectTilingData, expectWorkspaces);	1341	+ AippExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, expectTilingKey, expectTiling, expectWorkspaces);
669	}	1342	}
670		1343
671	TEST_F(AippTiling, aipp_tiling_padding_test_9)	1344	TEST_F(AippTiling, aipp_tiling_padding_test_9)
	@@ -680,8 +1353,974 @@ TEST_F(AippTiling, aipp_tiling_padding_test_9)
680	{gert::TilingContextPara::OpAttr("aipp_config_path", Ops::Cv::AnyValue::CreateFrom<string>(filePath))},	1353	{gert::TilingContextPara::OpAttr("aipp_config_path", Ops::Cv::AnyValue::CreateFrom<string>(filePath))},
681	&compileInfo);	1354	&compileInfo);
682		1355
683	- uint64_t expectTilingKey = 1;	1356	+ int64_t expectTilingKey = optiling::AIPP_RGB_PASS_THROUGH;
684	- string expectTilingData = "";	1357	+ Aipp_Kernel::AippTilingData expectTiling = {}; // 不检查
685	std::vector<size_t> expectWorkspaces = {16777216};	1358	std::vector<size_t> expectWorkspaces = {16777216};
686	- AippExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, expectTilingKey, expectTilingData, expectWorkspaces);	1359	+ AippExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, expectTilingKey, expectTiling, expectWorkspaces);
		1360	+}
		1361	+
		1362	+// test_12: YUV420SP_U8转YUV444 - 不带swap
		1363	+TEST_F(AippTiling, aipp_tiling_csc_test_12)
		1364	+{
		1365	+ AippCompileInfo compileInfo = {56, 253952};
		1366	+ std::string filePath = std::filesystem::current_path() / "aipp_ut_test_12.cfg";
		1367	+ std::string command = "cp ../../../../image/aipp/tests/ut/op_host/aipp_ut_test_12.cfg " + filePath;
		1368	+ system(command.c_str());
		1369	+ gert::TilingContextPara tilingContextPara("Aipp",
		1370	+ {{{{1, 1080, 1920, 3}, {1, 1080, 1920, 3}}, ge::DT_UINT8, ge::FORMAT_NHWC}},
		1371	+ {{{{1, 3, 1080, 1920}, {1, 3, 1080, 1920}}, ge::DT_FLOAT16, ge::FORMAT_NCHW}},
		1372	+ {gert::TilingContextPara::OpAttr("aipp_config_path", Ops::Cv::AnyValue::CreateFrom<string>(filePath))},
		1373	+ &compileInfo);
		1374	+
		1375	+ int64_t expectTilingKey = optiling::AIPP_YUV_PASS_THROUGH;
		1376	+ Aipp_Kernel::AippTilingData expectTiling = {
		1377	+ .imageFormat = 1, // YUV420SP_U8
		1378	+ .outputFormat = 1,
		1379	+ .batchNum = 1,
		1380	+ .channelNum = 3,
		1381	+ .inputSizeW = 1920,
		1382	+ .inputSizeH = 1080,
		1383	+ .outputSizeW = 1920,
		1384	+ .outputSizeH = 1080,
		1385	+ .srcChannelOffset = 0,
		1386	+ .cscParam = {
		1387	+ .cscSwitch = 0,
		1388	+ .rbuvSwapSwitch = 0,
		1389	+ .axSwapSwitch = 0,
		1390	+ .cscMatrix00 = 256,
		1391	+ .cscMatrix01 = 0,
		1392	+ .cscMatrix02 = 0,
		1393	+ .cscMatrix10 = 0,
		1394	+ .cscMatrix11 = 256,
		1395	+ .cscMatrix12 = 0,
		1396	+ .cscMatrix20 = 0,
		1397	+ .cscMatrix21 = 0,
		1398	+ .cscMatrix22 = 256,
		1399	+ .outBias0 = 0,
		1400	+ .outBias1 = 0,
		1401	+ .outBias2 = 0,
		1402	+ .inBias0 = 0,
		1403	+ .inBias1 = 0,
		1404	+ .inBias2 = 0
		1405	+ },
		1406	+ .cropParam = {
		1407	+ .cropSwitch = 0,
		1408	+ .cropStartPosH = 0,
		1409	+ .cropStartPosW = 0,
		1410	+ .cropSizeH = 1080,
		1411	+ .cropSizeW = 1920
		1412	+ },
		1413	+ .dtcParam = {
		1414	+ .dtcPixelMeanChn0 = 0,
		1415	+ .dtcPixelMeanChn1 = 0,
		1416	+ .dtcPixelMeanChn2 = 0,
		1417	+ .dtcPixelMeanChn3 = 0,
		1418	+ .dtcPixelMinChn0 = 0,
		1419	+ .dtcPixelMinChn1 = 0,
		1420	+ .dtcPixelMinChn2 = 0,
		1421	+ .dtcPixelMinChn3 = 0,
		1422	+ .dtcPixelVarReciChn0 = 1,
		1423	+ .dtcPixelVarReciChn1 = 1,
		1424	+ .dtcPixelVarReciChn2 = 1,
		1425	+ .dtcPixelVarReciChn3 = 1
		1426	+ },
		1427	+ .paddingParam = {
		1428	+ .paddingSwitch = 0,
		1429	+ .leftPaddingSize = 0,
		1430	+ .rightPaddingSize = 0,
		1431	+ .topPaddingSize = 0,
		1432	+ .bottomPaddingSize = 0,
		1433	+ .padValue = 0.0f
		1434	+ }
		1435	+ };
		1436	+ std::vector<size_t> expectWorkspaces = {16777216};
		1437	+ AippExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTiling, expectWorkspaces);
		1438	+}
		1439	+
		1440	+// test_13: YUV420SP_U8转YVU444 - 带rbuv swap (UV互换)
		1441	+TEST_F(AippTiling, aipp_tiling_csc_test_13)
		1442	+{
		1443	+ AippCompileInfo compileInfo = {56, 253952};
		1444	+ std::string filePath = std::filesystem::current_path() / "aipp_ut_test_13.cfg";
		1445	+ std::string command = "cp ../../../../image/aipp/tests/ut/op_host/aipp_ut_test_13.cfg " + filePath;
		1446	+ system(command.c_str());
		1447	+ gert::TilingContextPara tilingContextPara("Aipp",
		1448	+ {{{{1, 1080, 1920, 3}, {1, 1080, 1920, 3}}, ge::DT_UINT8, ge::FORMAT_NHWC}},
		1449	+ {{{{1, 3, 1080, 1920}, {1, 3, 1080, 1920}}, ge::DT_FLOAT16, ge::FORMAT_NCHW}},
		1450	+ {gert::TilingContextPara::OpAttr("aipp_config_path", Ops::Cv::AnyValue::CreateFrom<string>(filePath))},
		1451	+ &compileInfo);
		1452	+
		1453	+ int64_t expectTilingKey = optiling::AIPP_YUV_PASS_THROUGH;
		1454	+ Aipp_Kernel::AippTilingData expectTiling = {
		1455	+ .imageFormat = 1, .outputFormat = 1, .batchNum = 1, .channelNum = 3,
		1456	+ .inputSizeW = 1920, .inputSizeH = 1080, .outputSizeW = 1920, .outputSizeH = 1080, .srcChannelOffset = 0,
		1457	+ .cscParam = {.cscSwitch = 0, .rbuvSwapSwitch = 1, .axSwapSwitch = 0,
		1458	+ .cscMatrix00 = 256, .cscMatrix01 = 0, .cscMatrix02 = 0,
		1459	+ .cscMatrix10 = 0, .cscMatrix11 = 0, .cscMatrix12 = 256,
		1460	+ .cscMatrix20 = 0, .cscMatrix21 = 256, .cscMatrix22 = 0,
		1461	+ .outBias0 = 0, .outBias1 = 0, .outBias2 = 0, .inBias0 = 0, .inBias1 = 0, .inBias2 = 0},
		1462	+ .cropParam = {.cropSwitch = 0, .cropStartPosH = 0, .cropStartPosW = 0, .cropSizeH = 1080, .cropSizeW = 1920},
		1463	+ .dtcParam = {.dtcPixelMeanChn0 = 0, .dtcPixelMeanChn1 = 0, .dtcPixelMeanChn2 = 0, .dtcPixelMeanChn3 = 0,
		1464	+ .dtcPixelMinChn0 = 0, .dtcPixelMinChn1 = 0, .dtcPixelMinChn2 = 0, .dtcPixelMinChn3 = 0,
		1465	+ .dtcPixelVarReciChn0 = 1, .dtcPixelVarReciChn1 = 1, .dtcPixelVarReciChn2 = 1, .dtcPixelVarReciChn3 = 1},
		1466	+ .paddingParam = {.paddingSwitch = 0, .leftPaddingSize = 0, .rightPaddingSize = 0, .topPaddingSize = 0, .bottomPaddingSize = 0, .padValue = 0.0f}
		1467	+ };
		1468	+ std::vector<size_t> expectWorkspaces = {16777216};
		1469	+ AippExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTiling, expectWorkspaces);
		1470	+}
		1471	+
		1472	+// test_14: RGB888_U8转RGB - 不带csc和swap (直通模式)
		1473	+TEST_F(AippTiling, aipp_tiling_csc_test_14)
		1474	+{
		1475	+ AippCompileInfo compileInfo = {56, 253952};
		1476	+ std::string filePath = std::filesystem::current_path() / "aipp_ut_test_14.cfg";
		1477	+ std::string command = "cp ../../../../image/aipp/tests/ut/op_host/aipp_ut_test_14.cfg " + filePath;
		1478	+ system(command.c_str());
		1479	+ gert::TilingContextPara tilingContextPara("Aipp",
		1480	+ {{{{1, 1080, 1920, 3}, {1, 1080, 1920, 3}}, ge::DT_UINT8, ge::FORMAT_NHWC}},
		1481	+ {{{{1, 3, 1080, 1920}, {1, 3, 1080, 1920}}, ge::DT_FLOAT16, ge::FORMAT_NCHW}},
		1482	+ {gert::TilingContextPara::OpAttr("aipp_config_path", Ops::Cv::AnyValue::CreateFrom<string>(filePath))},
		1483	+ &compileInfo);
		1484	+
		1485	+ int64_t expectTilingKey = optiling::AIPP_RGB_PASS_THROUGH;
		1486	+ Aipp_Kernel::AippTilingData expectTiling = {
		1487	+ .imageFormat = 2,
		1488	+ .outputFormat = 1,
		1489	+ .batchNum = 1,
		1490	+ .channelNum = 3,
		1491	+ .inputSizeW = 1920,
		1492	+ .inputSizeH = 1080,
		1493	+ .outputSizeW = 1920,
		1494	+ .outputSizeH = 1080,
		1495	+ .srcChannelOffset = 0,
		1496	+ .cscParam = {
		1497	+ .cscSwitch = 0,
		1498	+ .rbuvSwapSwitch = 0,
		1499	+ .axSwapSwitch = 0,
		1500	+ .cscMatrix00 = 256,
		1501	+ .cscMatrix01 = 0,
		1502	+ .cscMatrix02 = 0,
		1503	+ .cscMatrix10 = 0,
		1504	+ .cscMatrix11 = 256,
		1505	+ .cscMatrix12 = 0,
		1506	+ .cscMatrix20 = 0,
		1507	+ .cscMatrix21 = 0,
		1508	+ .cscMatrix22 = 256,
		1509	+ .outBias0 = 0,
		1510	+ .outBias1 = 0,
		1511	+ .outBias2 = 0,
		1512	+ .inBias0 = 0,
		1513	+ .inBias1 = 0,
		1514	+ .inBias2 = 0
		1515	+ },
		1516	+ .cropParam = {
		1517	+ .cropSwitch = 0,
		1518	+ .cropStartPosH = 0,
		1519	+ .cropStartPosW = 0,
		1520	+ .cropSizeH = 1080,
		1521	+ .cropSizeW = 1920
		1522	+ },
		1523	+ .dtcParam = {
		1524	+ .dtcPixelMeanChn0 = 0,
		1525	+ .dtcPixelMeanChn1 = 0,
		1526	+ .dtcPixelMeanChn2 = 0,
		1527	+ .dtcPixelMeanChn3 = 0,
		1528	+ .dtcPixelMinChn0 = 0,
		1529	+ .dtcPixelMinChn1 = 0,
		1530	+ .dtcPixelMinChn2 = 0,
		1531	+ .dtcPixelMinChn3 = 0,
		1532	+ .dtcPixelVarReciChn0 = 1,
		1533	+ .dtcPixelVarReciChn1 = 1,
		1534	+ .dtcPixelVarReciChn2 = 1,
		1535	+ .dtcPixelVarReciChn3 = 1
		1536	+ },
		1537	+ .paddingParam = {
		1538	+ .paddingSwitch = 0,
		1539	+ .leftPaddingSize = 0,
		1540	+ .rightPaddingSize = 0,
		1541	+ .topPaddingSize = 0,
		1542	+ .bottomPaddingSize = 0,
		1543	+ .padValue = 0.0f
		1544	+ }
		1545	+ };
		1546	+ std::vector<size_t> expectWorkspaces = {16777216};
		1547	+ AippExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTiling, expectWorkspaces);
		1548	+}
		1549	+
		1550	+// test_15: RGB888_U8转YVU444 - 带csc矩阵转换(BT601标准YUV转换)
		1551	+TEST_F(AippTiling, aipp_tiling_csc_test_15)
		1552	+{
		1553	+ AippCompileInfo compileInfo = {56, 253952};
		1554	+ std::string filePath = std::filesystem::current_path() / "aipp_ut_test_15.cfg";
		1555	+ std::string command = "cp ../../../../image/aipp/tests/ut/op_host/aipp_ut_test_15.cfg " + filePath;
		1556	+ system(command.c_str());
		1557	+ gert::TilingContextPara tilingContextPara("Aipp",
		1558	+ {{{{1, 1080, 1920, 3}, {1, 1080, 1920, 3}}, ge::DT_UINT8, ge::FORMAT_NHWC}},
		1559	+ {{{{1, 3, 1080, 1920}, {1, 3, 1080, 1920}}, ge::DT_FLOAT16, ge::FORMAT_NCHW}},
		1560	+ {gert::TilingContextPara::OpAttr("aipp_config_path", Ops::Cv::AnyValue::CreateFrom<string>(filePath))},
		1561	+ &compileInfo);
		1562	+
		1563	+ int64_t expectTilingKey = optiling::AIPP_RGB_TO_YUV;
		1564	+ Aipp_Kernel::AippTilingData expectTiling = {
		1565	+ .imageFormat = 2,
		1566	+ .outputFormat = 1,
		1567	+ .batchNum = 1,
		1568	+ .channelNum = 3,
		1569	+ .inputSizeW = 1920,
		1570	+ .inputSizeH = 1080,
		1571	+ .outputSizeW = 1920,
		1572	+ .outputSizeH = 1080,
		1573	+ .srcChannelOffset = 0,
		1574	+ .cscParam = {
		1575	+ .cscSwitch = 1,
		1576	+ .rbuvSwapSwitch = 0,
		1577	+ .axSwapSwitch = 0,
		1578	+ .cscMatrix00 = 77,
		1579	+ .cscMatrix01 = 150,
		1580	+ .cscMatrix02 = 29,
		1581	+ .cscMatrix10 = 128,
		1582	+ .cscMatrix11 = -107,
		1583	+ .cscMatrix12 = -21,
		1584	+ .cscMatrix20 = -43,
		1585	+ .cscMatrix21 = -85,
		1586	+ .cscMatrix22 = 128,
		1587	+ .outBias0 = 0,
		1588	+ .outBias1 = 128,
		1589	+ .outBias2 = 128,
		1590	+ .inBias0 = 0,
		1591	+ .inBias1 = 0,
		1592	+ .inBias2 = 0},
		1593	+ .cropParam = {
		1594	+ .cropSwitch = 0,
		1595	+ .cropStartPosH = 0,
		1596	+ .cropStartPosW = 0,
		1597	+ .cropSizeH = 1080,
		1598	+ .cropSizeW = 1920},
		1599	+ .dtcParam = {
		1600	+ .dtcPixelMeanChn0 = 0,
		1601	+ .dtcPixelMeanChn1 = 0,
		1602	+ .dtcPixelMeanChn2 = 0,
		1603	+ .dtcPixelMeanChn3 = 0,
		1604	+ .dtcPixelMinChn0 = 0,
		1605	+ .dtcPixelMinChn1 = 0,
		1606	+ .dtcPixelMinChn2 = 0,
		1607	+ .dtcPixelMinChn3 = 0,
		1608	+ .dtcPixelVarReciChn0 = 1,
		1609	+ .dtcPixelVarReciChn1 = 1,
		1610	+ .dtcPixelVarReciChn2 = 1,
		1611	+ .dtcPixelVarReciChn3 = 1
		1612	+ },
		1613	+ .paddingParam = {
		1614	+ .paddingSwitch = 0,
		1615	+ .leftPaddingSize = 0,
		1616	+ .rightPaddingSize = 0,
		1617	+ .topPaddingSize = 0,
		1618	+ .bottomPaddingSize = 0,
		1619	+ .padValue = 0.0f
		1620	+ }
		1621	+ };
		1622	+ std::vector<size_t> expectWorkspaces = {16777216};
		1623	+ AippExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTiling, expectWorkspaces);
		1624	+}
		1625	+
		1626	+// test_16: BGR888_U8转RGB - 带rbuv swap (BGR通道互换)
		1627	+TEST_F(AippTiling, aipp_tiling_csc_test_16)
		1628	+{
		1629	+ AippCompileInfo compileInfo = {56, 253952};
		1630	+ std::string filePath = std::filesystem::current_path() / "aipp_ut_test_16.cfg";
		1631	+ std::string command = "cp ../../../../image/aipp/tests/ut/op_host/aipp_ut_test_16.cfg " + filePath;
		1632	+ system(command.c_str());
		1633	+ gert::TilingContextPara tilingContextPara("Aipp",
		1634	+ {{{{1, 1080, 1920, 3}, {1, 1080, 1920, 3}}, ge::DT_UINT8, ge::FORMAT_NHWC}},
		1635	+ {{{{1, 3, 1080, 1920}, {1, 3, 1080, 1920}}, ge::DT_FLOAT16, ge::FORMAT_NCHW}},
		1636	+ {gert::TilingContextPara::OpAttr("aipp_config_path", Ops::Cv::AnyValue::CreateFrom<string>(filePath))},
		1637	+ &compileInfo);
		1638	+
		1639	+ int64_t expectTilingKey = optiling::AIPP_RGB_PASS_THROUGH;
		1640	+ Aipp_Kernel::AippTilingData expectTiling = {
		1641	+ .imageFormat = 2,
		1642	+ .outputFormat = 1,
		1643	+ .batchNum = 1,
		1644	+ .channelNum = 3,
		1645	+ .inputSizeW = 1920,
		1646	+ .inputSizeH = 1080,
		1647	+ .outputSizeW = 1920,
		1648	+ .outputSizeH = 1080,
		1649	+ .srcChannelOffset = 0,
		1650	+ .cscParam = {
		1651	+ .cscSwitch = 0,
		1652	+ .rbuvSwapSwitch = 1,
		1653	+ .axSwapSwitch = 0,
		1654	+ .cscMatrix00 = 0,
		1655	+ .cscMatrix01 = 0,
		1656	+ .cscMatrix02 = 256,
		1657	+ .cscMatrix10 = 0,
		1658	+ .cscMatrix11 = 256,
		1659	+ .cscMatrix12 = 0,
		1660	+ .cscMatrix20 = 256,
		1661	+ .cscMatrix21 = 0,
		1662	+ .cscMatrix22 = 0,
		1663	+ .outBias0 = 0,
		1664	+ .outBias1 = 0,
		1665	+ .outBias2 = 0,
		1666	+ .inBias0 = 0,
		1667	+ .inBias1 = 0,
		1668	+ .inBias2 = 0
		1669	+ },
		1670	+ .cropParam = {
		1671	+ .cropSwitch = 0,
		1672	+ .cropStartPosH = 0,
		1673	+ .cropStartPosW = 0,
		1674	+ .cropSizeH = 1080,
		1675	+ .cropSizeW = 1920
		1676	+ },
		1677	+ .dtcParam = {
		1678	+ .dtcPixelMeanChn0 = 0,
		1679	+ .dtcPixelMeanChn1 = 0,
		1680	+ .dtcPixelMeanChn2 = 0,
		1681	+ .dtcPixelMeanChn3 = 0,
		1682	+ .dtcPixelMinChn0 = 0,
		1683	+ .dtcPixelMinChn1 = 0,
		1684	+ .dtcPixelMinChn2 = 0,
		1685	+ .dtcPixelMinChn3 = 0,
		1686	+ .dtcPixelVarReciChn0 = 1,
		1687	+ .dtcPixelVarReciChn1 = 1,
		1688	+ .dtcPixelVarReciChn2 = 1,
		1689	+ .dtcPixelVarReciChn3 = 1
		1690	+ },
		1691	+ .paddingParam = {
		1692	+ .paddingSwitch = 0,
		1693	+ .leftPaddingSize = 0,
		1694	+ .rightPaddingSize = 0,
		1695	+ .topPaddingSize = 0,
		1696	+ .bottomPaddingSize = 0,
		1697	+ .padValue = 0.0f
		1698	+ }
		1699	+ };
		1700	+ std::vector<size_t> expectWorkspaces = {16777216};
		1701	+ AippExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTiling, expectWorkspaces);
		1702	+}
		1703	+
		1704	+// test_17: BGR888_U8转BGR - 不带swap (直通)
		1705	+TEST_F(AippTiling, aipp_tiling_csc_test_17)
		1706	+{
		1707	+ AippCompileInfo compileInfo = {56, 253952};
		1708	+ std::string filePath = std::filesystem::current_path() / "aipp_ut_test_17.cfg";
		1709	+ std::string command = "cp ../../../../image/aipp/tests/ut/op_host/aipp_ut_test_17.cfg " + filePath;
		1710	+ system(command.c_str());
		1711	+ gert::TilingContextPara tilingContextPara("Aipp",
		1712	+ {{{{1, 1080, 1920, 3}, {1, 1080, 1920, 3}}, ge::DT_UINT8, ge::FORMAT_NHWC}},
		1713	+ {{{{1, 3, 1080, 1920}, {1, 3, 1080, 1920}}, ge::DT_FLOAT16, ge::FORMAT_NCHW}},
		1714	+ {gert::TilingContextPara::OpAttr("aipp_config_path", Ops::Cv::AnyValue::CreateFrom<string>(filePath))},
		1715	+ &compileInfo);
		1716	+
		1717	+ int64_t expectTilingKey = optiling::AIPP_RGB_PASS_THROUGH;
		1718	+ Aipp_Kernel::AippTilingData expectTiling = {
		1719	+ .imageFormat = 2,
		1720	+ .outputFormat = 1,
		1721	+ .batchNum = 1,
		1722	+ .channelNum = 3,
		1723	+ .inputSizeW = 1920,
		1724	+ .inputSizeH = 1080,
		1725	+ .outputSizeW = 1920,
		1726	+ .outputSizeH = 1080,
		1727	+ .srcChannelOffset = 0,
		1728	+ .cscParam = {
		1729	+ .cscSwitch = 0,
		1730	+ .rbuvSwapSwitch = 0,
		1731	+ .axSwapSwitch = 0,
		1732	+ .cscMatrix00 = 256,
		1733	+ .cscMatrix01 = 0,
		1734	+ .cscMatrix02 = 0,
		1735	+ .cscMatrix10 = 0,
		1736	+ .cscMatrix11 = 256,
		1737	+ .cscMatrix12 = 0,
		1738	+ .cscMatrix20 = 0,
		1739	+ .cscMatrix21 = 0,
		1740	+ .cscMatrix22 = 256,
		1741	+ .outBias0 = 0,
		1742	+ .outBias1 = 0,
		1743	+ .outBias2 = 0,
		1744	+ .inBias0 = 0,
		1745	+ .inBias1 = 0,
		1746	+ .inBias2 = 0
		1747	+ },
		1748	+ .cropParam = {
		1749	+ .cropSwitch = 0,
		1750	+ .cropStartPosH = 0,
		1751	+ .cropStartPosW = 0,
		1752	+ .cropSizeH = 1080,
		1753	+ .cropSizeW = 1920
		1754	+ },
		1755	+ .dtcParam = {
		1756	+ .dtcPixelMeanChn0 = 0,
		1757	+ .dtcPixelMeanChn1 = 0,
		1758	+ .dtcPixelMeanChn2 = 0,
		1759	+ .dtcPixelMeanChn3 = 0,
		1760	+ .dtcPixelMinChn0 = 0,
		1761	+ .dtcPixelMinChn1 = 0,
		1762	+ .dtcPixelMinChn2 = 0,
		1763	+ .dtcPixelMinChn3 = 0,
		1764	+ .dtcPixelVarReciChn0 = 1,
		1765	+ .dtcPixelVarReciChn1 = 1,
		1766	+ .dtcPixelVarReciChn2 = 1,
		1767	+ .dtcPixelVarReciChn3 = 1
		1768	+ },
		1769	+ .paddingParam = {
		1770	+ .paddingSwitch = 0,
		1771	+ .leftPaddingSize = 0,
		1772	+ .rightPaddingSize = 0,
		1773	+ .topPaddingSize = 0,
		1774	+ .bottomPaddingSize = 0,
		1775	+ .padValue = 0.0f
		1776	+ }
		1777	+ };
		1778	+ std::vector<size_t> expectWorkspaces = {16777216};
		1779	+ AippExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTiling, expectWorkspaces);
		1780	+}
		1781	+
		1782	+// test_18: XRGB8888_U8转YUV444 - 带ax swap和csc矩阵
		1783	+TEST_F(AippTiling, aipp_tiling_csc_test_18)
		1784	+{
		1785	+ AippCompileInfo compileInfo = {56, 253952};
		1786	+ std::string filePath = std::filesystem::current_path() / "aipp_ut_test_18.cfg";
		1787	+ std::string command = "cp ../../../../image/aipp/tests/ut/op_host/aipp_ut_test_18.cfg " + filePath;
		1788	+ system(command.c_str());
		1789	+ gert::TilingContextPara tilingContextPara("Aipp",
		1790	+ {{{{1, 1080, 1920, 4}, {1, 1080, 1920, 4}}, ge::DT_UINT8, ge::FORMAT_NHWC}},
		1791	+ {{{{1, 3, 1080, 1920}, {1, 3, 1080, 1920}}, ge::DT_FLOAT16, ge::FORMAT_NCHW}},
		1792	+ {gert::TilingContextPara::OpAttr("aipp_config_path", Ops::Cv::AnyValue::CreateFrom<string>(filePath))},
		1793	+ &compileInfo);
		1794	+
		1795	+ int64_t expectTilingKey = optiling::AIPP_RGB_TO_YUV;
		1796	+ Aipp_Kernel::AippTilingData expectTiling = {
		1797	+ .imageFormat = 3,
		1798	+ .outputFormat = 1,
		1799	+ .batchNum = 1,
		1800	+ .channelNum = 4,
		1801	+ .inputSizeW = 1920,
		1802	+ .inputSizeH = 1080,
		1803	+ .outputSizeW = 1920,
		1804	+ .outputSizeH = 1080,
		1805	+ .srcChannelOffset = 1,
		1806	+ .cscParam = {
		1807	+ .cscSwitch = 1,
		1808	+ .rbuvSwapSwitch = 0,
		1809	+ .axSwapSwitch = 1,
		1810	+ .cscMatrix00 = 77,
		1811	+ .cscMatrix01 = 150,
		1812	+ .cscMatrix02 = 29,
		1813	+ .cscMatrix10 = -43,
		1814	+ .cscMatrix11 = -85,
		1815	+ .cscMatrix12 = 128,
		1816	+ .cscMatrix20 = 128,
		1817	+ .cscMatrix21 = -107,
		1818	+ .cscMatrix22 = -21,
		1819	+ .outBias0 = 0,
		1820	+ .outBias1 = 128,
		1821	+ .outBias2 = 128,
		1822	+ .inBias0 = 0,
		1823	+ .inBias1 = 0,
		1824	+ .inBias2 = 0
		1825	+ },
		1826	+ .cropParam = {
		1827	+ .cropSwitch = 0,
		1828	+ .cropStartPosH = 0,
		1829	+ .cropStartPosW = 0,
		1830	+ .cropSizeH = 1080,
		1831	+ .cropSizeW = 1920
		1832	+ },
		1833	+ .dtcParam = {
		1834	+ .dtcPixelMeanChn0 = 0,
		1835	+ .dtcPixelMeanChn1 = 0,
		1836	+ .dtcPixelMeanChn2 = 0,
		1837	+ .dtcPixelMeanChn3 = 0,
		1838	+ .dtcPixelMinChn0 = 0,
		1839	+ .dtcPixelMinChn1 = 0,
		1840	+ .dtcPixelMinChn2 = 0,
		1841	+ .dtcPixelMinChn3 = 0,
		1842	+ .dtcPixelVarReciChn0 = 1,
		1843	+ .dtcPixelVarReciChn1 = 1,
		1844	+ .dtcPixelVarReciChn2 = 1,
		1845	+ .dtcPixelVarReciChn3 = 1
		1846	+ },
		1847	+ .paddingParam = {
		1848	+ .paddingSwitch = 0,
		1849	+ .leftPaddingSize = 0,
		1850	+ .rightPaddingSize = 0,
		1851	+ .topPaddingSize = 0,
		1852	+ .bottomPaddingSize = 0,
		1853	+ .padValue = 0.0f
		1854	+ }
		1855	+ };
		1856	+ std::vector<size_t> expectWorkspaces = {16777216};
		1857	+ AippExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTiling, expectWorkspaces);
		1858	+}
		1859	+
		1860	+// test_19: XRGB8888_U8转YVU444 - 带ax swap和csc矩阵(UV互换)
		1861	+TEST_F(AippTiling, aipp_tiling_csc_test_19)
		1862	+{
		1863	+ AippCompileInfo compileInfo = {56, 253952};
		1864	+ std::string filePath = std::filesystem::current_path() / "aipp_ut_test_19.cfg";
		1865	+ std::string command = "cp ../../../../image/aipp/tests/ut/op_host/aipp_ut_test_19.cfg " + filePath;
		1866	+ system(command.c_str());
		1867	+ gert::TilingContextPara tilingContextPara("Aipp",
		1868	+ {{{{1, 1080, 1920, 4}, {1, 1080, 1920, 4}}, ge::DT_UINT8, ge::FORMAT_NHWC}},
		1869	+ {{{{1, 3, 1080, 1920}, {1, 3, 1080, 1920}}, ge::DT_FLOAT16, ge::FORMAT_NCHW}},
		1870	+ {gert::TilingContextPara::OpAttr("aipp_config_path", Ops::Cv::AnyValue::CreateFrom<string>(filePath))},
		1871	+ &compileInfo);
		1872	+
		1873	+ int64_t expectTilingKey = optiling::AIPP_RGB_TO_YUV;
		1874	+ Aipp_Kernel::AippTilingData expectTiling = {
		1875	+ .imageFormat = 3,
		1876	+ .outputFormat = 1,
		1877	+ .batchNum = 1,
		1878	+ .channelNum = 4,
		1879	+ .inputSizeW = 1920,
		1880	+ .inputSizeH = 1080,
		1881	+ .outputSizeW = 1920,
		1882	+ .outputSizeH = 1080,
		1883	+ .srcChannelOffset = 1,
		1884	+ .cscParam = {
		1885	+ .cscSwitch = 1,
		1886	+ .rbuvSwapSwitch = 0,
		1887	+ .axSwapSwitch = 1,
		1888	+ .cscMatrix00 = 77,
		1889	+ .cscMatrix01 = 150,
		1890	+ .cscMatrix02 = 29,
		1891	+ .cscMatrix10 = 128,
		1892	+ .cscMatrix11 = -107,
		1893	+ .cscMatrix12 = -21,
		1894	+ .cscMatrix20 = -43,
		1895	+ .cscMatrix21 = -85,
		1896	+ .cscMatrix22 = 128,
		1897	+ .outBias0 = 0,
		1898	+ .outBias1 = 128,
		1899	+ .outBias2 = 128,
		1900	+ .inBias0 = 0,
		1901	+ .inBias1 = 0,
		1902	+ .inBias2 = 0
		1903	+ },
		1904	+ .cropParam = {
		1905	+ .cropSwitch = 0,
		1906	+ .cropStartPosH = 0,
		1907	+ .cropStartPosW = 0,
		1908	+ .cropSizeH = 1080,
		1909	+ .cropSizeW = 1920
		1910	+ },
		1911	+ .dtcParam = {
		1912	+ .dtcPixelMeanChn0 = 0,
		1913	+ .dtcPixelMeanChn1 = 0,
		1914	+ .dtcPixelMeanChn2 = 0,
		1915	+ .dtcPixelMeanChn3 = 0,
		1916	+ .dtcPixelMinChn0 = 0,
		1917	+ .dtcPixelMinChn1 = 0,
		1918	+ .dtcPixelMinChn2 = 0,
		1919	+ .dtcPixelMinChn3 = 0,
		1920	+ .dtcPixelVarReciChn0 = 1,
		1921	+ .dtcPixelVarReciChn1 = 1,
		1922	+ .dtcPixelVarReciChn2 = 1,
		1923	+ .dtcPixelVarReciChn3 = 1
		1924	+ },
		1925	+ .paddingParam = {
		1926	+ .paddingSwitch = 0,
		1927	+ .leftPaddingSize = 0,
		1928	+ .rightPaddingSize = 0,
		1929	+ .topPaddingSize = 0,
		1930	+ .bottomPaddingSize = 0,
		1931	+ .padValue = 0.0f
		1932	+ }
		1933	+ };
		1934	+ std::vector<size_t> expectWorkspaces = {16777216};
		1935	+ AippExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTiling, expectWorkspaces);
		1936	+}
		1937	+
		1938	+// test_20: XRGB8888_U8转GRAY - 带ax swap和灰度矩阵
		1939	+TEST_F(AippTiling, aipp_tiling_csc_test_20)
		1940	+{
		1941	+ AippCompileInfo compileInfo = {56, 253952};
		1942	+ std::string filePath = std::filesystem::current_path() / "aipp_ut_test_20.cfg";
		1943	+ std::string command = "cp ../../../../image/aipp/tests/ut/op_host/aipp_ut_test_20.cfg " + filePath;
		1944	+ system(command.c_str());
		1945	+ gert::TilingContextPara tilingContextPara("Aipp",
		1946	+ {{{{1, 1080, 1920, 4}, {1, 1080, 1920, 4}}, ge::DT_UINT8, ge::FORMAT_NHWC}},
		1947	+ {{{{1, 1, 1080, 1920}, {1, 1, 1080, 1920}}, ge::DT_FLOAT16, ge::FORMAT_NCHW}},
		1948	+ {gert::TilingContextPara::OpAttr("aipp_config_path", Ops::Cv::AnyValue::CreateFrom<string>(filePath))},
		1949	+ &compileInfo);
		1950	+
		1951	+ int64_t expectTilingKey = optiling::AIPP_RGB_TO_GRAY;
		1952	+ Aipp_Kernel::AippTilingData expectTiling = {
		1953	+ .imageFormat = 3,
		1954	+ .outputFormat = 1,
		1955	+ .batchNum = 1,
		1956	+ .channelNum = 4,
		1957	+ .inputSizeW = 1920,
		1958	+ .inputSizeH = 1080,
		1959	+ .outputSizeW = 1920,
		1960	+ .outputSizeH = 1080,
		1961	+ .srcChannelOffset = 1,
		1962	+ .cscParam = {
		1963	+ .cscSwitch = 1,
		1964	+ .rbuvSwapSwitch = 0,
		1965	+ .axSwapSwitch = 1,
		1966	+ .cscMatrix00 = 76,
		1967	+ .cscMatrix01 = 150,
		1968	+ .cscMatrix02 = 30,
		1969	+ .cscMatrix10 = 0,
		1970	+ .cscMatrix11 = 0,
		1971	+ .cscMatrix12 = 0,
		1972	+ .cscMatrix20 = 0,
		1973	+ .cscMatrix21 = 0,
		1974	+ .cscMatrix22 = 0,
		1975	+ .outBias0 = 0,
		1976	+ .outBias1 = 0,
		1977	+ .outBias2 = 0,
		1978	+ .inBias0 = 0,
		1979	+ .inBias1 = 0,
		1980	+ .inBias2 = 0
		1981	+ },
		1982	+ .cropParam = {
		1983	+ .cropSwitch = 0,
		1984	+ .cropStartPosH = 0,
		1985	+ .cropStartPosW = 0,
		1986	+ .cropSizeH = 1080,
		1987	+ .cropSizeW = 1920
		1988	+ },
		1989	+ .dtcParam = {
		1990	+ .dtcPixelMeanChn0 = 0,
		1991	+ .dtcPixelMeanChn1 = 0,
		1992	+ .dtcPixelMeanChn2 = 0,
		1993	+ .dtcPixelMeanChn3 = 0,
		1994	+ .dtcPixelMinChn0 = 0,
		1995	+ .dtcPixelMinChn1 = 0,
		1996	+ .dtcPixelMinChn2 = 0,
		1997	+ .dtcPixelMinChn3 = 0,
		1998	+ .dtcPixelVarReciChn0 = 1,
		1999	+ .dtcPixelVarReciChn1 = 1,
		2000	+ .dtcPixelVarReciChn2 = 1,
		2001	+ .dtcPixelVarReciChn3 = 1
		2002	+ },
		2003	+ .paddingParam = {
		2004	+ .paddingSwitch = 0,
		2005	+ .leftPaddingSize = 0,
		2006	+ .rightPaddingSize = 0,
		2007	+ .topPaddingSize = 0,
		2008	+ .bottomPaddingSize = 0,
		2009	+ .padValue = 0.0f
		2010	+ }
		2011	+ };
		2012	+ std::vector<size_t> expectWorkspaces = {16777216};
		2013	+ AippExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTiling, expectWorkspaces);
		2014	+}
		2015	+
		2016	+// test_21: XBGR8888_U8转GRAY - 带ax swap和rbuv swap
		2017	+TEST_F(AippTiling, aipp_tiling_csc_test_21)
		2018	+{
		2019	+ AippCompileInfo compileInfo = {56, 253952};
		2020	+ std::string filePath = std::filesystem::current_path() / "aipp_ut_test_21.cfg";
		2021	+ std::string command = "cp ../../../../image/aipp/tests/ut/op_host/aipp_ut_test_21.cfg " + filePath;
		2022	+ system(command.c_str());
		2023	+ gert::TilingContextPara tilingContextPara("Aipp",
		2024	+ {{{{1, 1080, 1920, 4}, {1, 1080, 1920, 4}}, ge::DT_UINT8, ge::FORMAT_NHWC}},
		2025	+ {{{{1, 1, 1080, 1920}, {1, 1, 1080, 1920}}, ge::DT_FLOAT16, ge::FORMAT_NCHW}},
		2026	+ {gert::TilingContextPara::OpAttr("aipp_config_path", Ops::Cv::AnyValue::CreateFrom<string>(filePath))},
		2027	+ &compileInfo);
		2028	+
		2029	+ int64_t expectTilingKey = optiling::AIPP_RGB_TO_GRAY;
		2030	+ Aipp_Kernel::AippTilingData expectTiling = {
		2031	+ .imageFormat = 3,
		2032	+ .outputFormat = 1,
		2033	+ .batchNum = 1,
		2034	+ .channelNum = 4,
		2035	+ .inputSizeW = 1920,
		2036	+ .inputSizeH = 1080,
		2037	+ .outputSizeW = 1920,
		2038	+ .outputSizeH = 1080,
		2039	+ .srcChannelOffset = 1,
		2040	+ .cscParam = {
		2041	+ .cscSwitch = 1,
		2042	+ .rbuvSwapSwitch = 1,
		2043	+ .axSwapSwitch = 1,
		2044	+ .cscMatrix00 = 30,
		2045	+ .cscMatrix01 = 150,
		2046	+ .cscMatrix02 = 76,
		2047	+ .cscMatrix10 = 0,
		2048	+ .cscMatrix11 = 0,
		2049	+ .cscMatrix12 = 0,
		2050	+ .cscMatrix20 = 0,
		2051	+ .cscMatrix21 = 0,
		2052	+ .cscMatrix22 = 0,
		2053	+ .outBias0 = 0,
		2054	+ .outBias1 = 0,
		2055	+ .outBias2 = 0,
		2056	+ .inBias0 = 0,
		2057	+ .inBias1 = 0,
		2058	+ .inBias2 = 0
		2059	+ },
		2060	+ .cropParam = {
		2061	+ .cropSwitch = 0,
		2062	+ .cropStartPosH = 0,
		2063	+ .cropStartPosW = 0,
		2064	+ .cropSizeH = 1080,
		2065	+ .cropSizeW = 1920
		2066	+ },
		2067	+ .dtcParam = {
		2068	+ .dtcPixelMeanChn0 = 0,
		2069	+ .dtcPixelMeanChn1 = 0,
		2070	+ .dtcPixelMeanChn2 = 0,
		2071	+ .dtcPixelMeanChn3 = 0,
		2072	+ .dtcPixelMinChn0 = 0,
		2073	+ .dtcPixelMinChn1 = 0,
		2074	+ .dtcPixelMinChn2 = 0,
		2075	+ .dtcPixelMinChn3 = 0,
		2076	+ .dtcPixelVarReciChn0 = 1,
		2077	+ .dtcPixelVarReciChn1 = 1,
		2078	+ .dtcPixelVarReciChn2 = 1,
		2079	+ .dtcPixelVarReciChn3 = 1
		2080	+ },
		2081	+ .paddingParam = {
		2082	+ .paddingSwitch = 0,
		2083	+ .leftPaddingSize = 0,
		2084	+ .rightPaddingSize = 0,
		2085	+ .topPaddingSize = 0,
		2086	+ .bottomPaddingSize = 0,
		2087	+ .padValue = 0.0f
		2088	+ }
		2089	+ };
		2090	+ std::vector<size_t> expectWorkspaces = {16777216};
		2091	+ AippExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTiling, expectWorkspaces);
		2092	+}
		2093	+
		2094	+// test_22: RGBX8888_U8转GRAY - 带ax swap关闭
		2095	+TEST_F(AippTiling, aipp_tiling_csc_test_22)
		2096	+{
		2097	+ AippCompileInfo compileInfo = {56, 253952};
		2098	+ std::string filePath = std::filesystem::current_path() / "aipp_ut_test_22.cfg";
		2099	+ std::string command = "cp ../../../../image/aipp/tests/ut/op_host/aipp_ut_test_22.cfg " + filePath;
		2100	+ system(command.c_str());
		2101	+ gert::TilingContextPara tilingContextPara("Aipp",
		2102	+ {{{{1, 1080, 1920, 4}, {1, 1080, 1920, 4}}, ge::DT_UINT8, ge::FORMAT_NHWC}},
		2103	+ {{{{1, 1, 1080, 1920}, {1, 1, 1080, 1920}}, ge::DT_FLOAT16, ge::FORMAT_NCHW}},
		2104	+ {gert::TilingContextPara::OpAttr("aipp_config_path", Ops::Cv::AnyValue::CreateFrom<string>(filePath))},
		2105	+ &compileInfo);
		2106	+
		2107	+ int64_t expectTilingKey = optiling::AIPP_RGB_TO_GRAY;
		2108	+ Aipp_Kernel::AippTilingData expectTiling = {
		2109	+ .imageFormat = 3,
		2110	+ .outputFormat = 1,
		2111	+ .batchNum = 1,
		2112	+ .channelNum = 4,
		2113	+ .inputSizeW = 1920,
		2114	+ .inputSizeH = 1080,
		2115	+ .outputSizeW = 1920,
		2116	+ .outputSizeH = 1080,
		2117	+ .srcChannelOffset = 0,
		2118	+ .cscParam = {
		2119	+ .cscSwitch = 1,
		2120	+ .rbuvSwapSwitch = 0,
		2121	+ .axSwapSwitch = 0,
		2122	+ .cscMatrix00 = 76,
		2123	+ .cscMatrix01 = 150,
		2124	+ .cscMatrix02 = 30,
		2125	+ .cscMatrix10 = 0,
		2126	+ .cscMatrix11 = 0,
		2127	+ .cscMatrix12 = 0,
		2128	+ .cscMatrix20 = 0,
		2129	+ .cscMatrix21 = 0,
		2130	+ .cscMatrix22 = 0,
		2131	+ .outBias0 = 0,
		2132	+ .outBias1 = 0,
		2133	+ .outBias2 = 0,
		2134	+ .inBias0 = 0,
		2135	+ .inBias1 = 0,
		2136	+ .inBias2 = 0
		2137	+ },
		2138	+ .cropParam = {
		2139	+ .cropSwitch = 0,
		2140	+ .cropStartPosH = 0,
		2141	+ .cropStartPosW = 0,
		2142	+ .cropSizeH = 1080,
		2143	+ .cropSizeW = 1920
		2144	+ },
		2145	+ .dtcParam = {
		2146	+ .dtcPixelMeanChn0 = 0,
		2147	+ .dtcPixelMeanChn1 = 0,
		2148	+ .dtcPixelMeanChn2 = 0,
		2149	+ .dtcPixelMeanChn3 = 0,
		2150	+ .dtcPixelMinChn0 = 0,
		2151	+ .dtcPixelMinChn1 = 0,
		2152	+ .dtcPixelMinChn2 = 0,
		2153	+ .dtcPixelMinChn3 = 0,
		2154	+ .dtcPixelVarReciChn0 = 1,
		2155	+ .dtcPixelVarReciChn1 = 1,
		2156	+ .dtcPixelVarReciChn2 = 1,
		2157	+ .dtcPixelVarReciChn3 = 1
		2158	+ },
		2159	+ .paddingParam = {
		2160	+ .paddingSwitch = 0,
		2161	+ .leftPaddingSize = 0,
		2162	+ .rightPaddingSize = 0,
		2163	+ .topPaddingSize = 0,
		2164	+ .bottomPaddingSize = 0,
		2165	+ .padValue = 0.0f
		2166	+ }
		2167	+ };
		2168	+ std::vector<size_t> expectWorkspaces = {16777216};
		2169	+ AippExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTiling, expectWorkspaces);
		2170	+}
		2171	+
		2172	+// test_23: BGRX8888_U8转GRAY - 带rbuv swap和ax swap关闭
		2173	+TEST_F(AippTiling, aipp_tiling_csc_test_23)
		2174	+{
		2175	+ AippCompileInfo compileInfo = {56, 253952};
		2176	+ std::string filePath = std::filesystem::current_path() / "aipp_ut_test_23.cfg";
		2177	+ std::string command = "cp ../../../../image/aipp/tests/ut/op_host/aipp_ut_test_23.cfg " + filePath;
		2178	+ system(command.c_str());
		2179	+ gert::TilingContextPara tilingContextPara("Aipp",
		2180	+ {{{{1, 1080, 1920, 4}, {1, 1080, 1920, 4}}, ge::DT_UINT8, ge::FORMAT_NHWC}},
		2181	+ {{{{1, 1, 1080, 1920}, {1, 1, 1080, 1920}}, ge::DT_FLOAT16, ge::FORMAT_NCHW}},
		2182	+ {gert::TilingContextPara::OpAttr("aipp_config_path", Ops::Cv::AnyValue::CreateFrom<string>(filePath))},
		2183	+ &compileInfo);
		2184	+
		2185	+ int64_t expectTilingKey = optiling::AIPP_RGB_TO_GRAY;
		2186	+ Aipp_Kernel::AippTilingData expectTiling = {
		2187	+ .imageFormat = 3,
		2188	+ .outputFormat = 1,
		2189	+ .batchNum = 1,
		2190	+ .channelNum = 4,
		2191	+ .inputSizeW = 1920,
		2192	+ .inputSizeH = 1080,
		2193	+ .outputSizeW = 1920,
		2194	+ .outputSizeH = 1080,
		2195	+ .srcChannelOffset = 0,
		2196	+ .cscParam = {
		2197	+ .cscSwitch = 1,
		2198	+ .rbuvSwapSwitch = 1,
		2199	+ .axSwapSwitch = 0,
		2200	+ .cscMatrix00 = 30,
		2201	+ .cscMatrix01 = 150,
		2202	+ .cscMatrix02 = 76,
		2203	+ .cscMatrix10 = 0,
		2204	+ .cscMatrix11 = 0,
		2205	+ .cscMatrix12 = 0,
		2206	+ .cscMatrix20 = 0,
		2207	+ .cscMatrix21 = 0,
		2208	+ .cscMatrix22 = 0,
		2209	+ .outBias0 = 0,
		2210	+ .outBias1 = 0,
		2211	+ .outBias2 = 0,
		2212	+ .inBias0 = 0,
		2213	+ .inBias1 = 0,
		2214	+ .inBias2 = 0
		2215	+ },
		2216	+ .cropParam = {
		2217	+ .cropSwitch = 0,
		2218	+ .cropStartPosH = 0,
		2219	+ .cropStartPosW = 0,
		2220	+ .cropSizeH = 1080,
		2221	+ .cropSizeW = 1920
		2222	+ },
		2223	+ .dtcParam = {
		2224	+ .dtcPixelMeanChn0 = 0,
		2225	+ .dtcPixelMeanChn1 = 0,
		2226	+ .dtcPixelMeanChn2 = 0,
		2227	+ .dtcPixelMeanChn3 = 0,
		2228	+ .dtcPixelMinChn0 = 0,
		2229	+ .dtcPixelMinChn1 = 0,
		2230	+ .dtcPixelMinChn2 = 0,
		2231	+ .dtcPixelMinChn3 = 0,
		2232	+ .dtcPixelVarReciChn0 = 1,
		2233	+ .dtcPixelVarReciChn1 = 1,
		2234	+ .dtcPixelVarReciChn2 = 1,
		2235	+ .dtcPixelVarReciChn3 = 1
		2236	+ },
		2237	+ .paddingParam = {
		2238	+ .paddingSwitch = 0,
		2239	+ .leftPaddingSize = 0,
		2240	+ .rightPaddingSize = 0,
		2241	+ .topPaddingSize = 0,
		2242	+ .bottomPaddingSize = 0,
		2243	+ .padValue = 0.0f
		2244	+ }
		2245	+ };
		2246	+ std::vector<size_t> expectWorkspaces = {16777216};
		2247	+ AippExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTiling, expectWorkspaces);
		2248	+}
		2249	+
		2250	+// test_24: YUV400_U8转GRAY - 不带csc (YUV400单通道直通)
		2251	+TEST_F(AippTiling, aipp_tiling_csc_test_24)
		2252	+{
		2253	+ AippCompileInfo compileInfo = {56, 253952};
		2254	+ std::string filePath = std::filesystem::current_path() / "aipp_ut_test_24.cfg";
		2255	+ std::string command = "cp ../../../../image/aipp/tests/ut/op_host/aipp_ut_test_24.cfg " + filePath;
		2256	+ system(command.c_str());
		2257	+ gert::TilingContextPara tilingContextPara("Aipp",
		2258	+ {{{{1, 1080, 1920, 1}, {1, 1080, 1920, 1}}, ge::DT_UINT8, ge::FORMAT_NHWC}},
		2259	+ {{{{1, 1, 1080, 1920}, {1, 1, 1080, 1920}}, ge::DT_FLOAT16, ge::FORMAT_NCHW}},
		2260	+ {gert::TilingContextPara::OpAttr("aipp_config_path", Ops::Cv::AnyValue::CreateFrom<string>(filePath))},
		2261	+ &compileInfo);
		2262	+
		2263	+ int64_t expectTilingKey = optiling::AIPP_YUV_TO_GRAY;
		2264	+ Aipp_Kernel::AippTilingData expectTiling = {
		2265	+ .imageFormat = 4,
		2266	+ .outputFormat = 1,
		2267	+ .batchNum = 1,
		2268	+ .channelNum = 1,
		2269	+ .inputSizeW = 1920,
		2270	+ .inputSizeH = 1080,
		2271	+ .outputSizeW = 1920,
		2272	+ .outputSizeH = 1080,
		2273	+ .srcChannelOffset = 0,
		2274	+ .cscParam = {
		2275	+ .cscSwitch = 0,
		2276	+ .rbuvSwapSwitch = 0,
		2277	+ .axSwapSwitch = 0,
		2278	+ .cscMatrix00 = 256,
		2279	+ .cscMatrix01 = 0,
		2280	+ .cscMatrix02 = 0,
		2281	+ .cscMatrix10 = 0,
		2282	+ .cscMatrix11 = 256,
		2283	+ .cscMatrix12 = 0,
		2284	+ .cscMatrix20 = 0,
		2285	+ .cscMatrix21 = 0,
		2286	+ .cscMatrix22 = 256,
		2287	+ .outBias0 = 0,
		2288	+ .outBias1 = 0,
		2289	+ .outBias2 = 0,
		2290	+ .inBias0 = 0,
		2291	+ .inBias1 = 0,
		2292	+ .inBias2 = 0
		2293	+ },
		2294	+ .cropParam = {
		2295	+ .cropSwitch = 0,
		2296	+ .cropStartPosH = 0,
		2297	+ .cropStartPosW = 0,
		2298	+ .cropSizeH = 1080,
		2299	+ .cropSizeW = 1920
		2300	+ },
		2301	+ .dtcParam = {
		2302	+ .dtcPixelMeanChn0 = 0,
		2303	+ .dtcPixelMeanChn1 = 0,
		2304	+ .dtcPixelMeanChn2 = 0,
		2305	+ .dtcPixelMeanChn3 = 0,
		2306	+ .dtcPixelMinChn0 = 0,
		2307	+ .dtcPixelMinChn1 = 0,
		2308	+ .dtcPixelMinChn2 = 0,
		2309	+ .dtcPixelMinChn3 = 0,
		2310	+ .dtcPixelVarReciChn0 = 1,
		2311	+ .dtcPixelVarReciChn1 = 1,
		2312	+ .dtcPixelVarReciChn2 = 1,
		2313	+ .dtcPixelVarReciChn3 = 1
		2314	+ },
		2315	+ .paddingParam = {
		2316	+ .paddingSwitch = 0,
		2317	+ .leftPaddingSize = 0,
		2318	+ .rightPaddingSize = 0,
		2319	+ .topPaddingSize = 0,
		2320	+ .bottomPaddingSize = 0,
		2321	+ .padValue = 0.0f
		2322	+ }
		2323	+ };
		2324	+ std::vector<size_t> expectWorkspaces = {16777216};
		2325	+ AippExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTiling, expectWorkspaces);
687	}	2326	}