Conv2DTranspose支持非对称量化场景-ops-nn-AtomGit

Mconv/conv3d_backprop_input_v2/op_host/op_tiling/arch35/conv3d_backprop_input_v2_base_tiling.cpp+5-1

	@@ -214,6 +214,9 @@ ge::graphStatus Conv3DBackpropInputV2TilingArch35::DoOpTiling()
214	}	214	}
215	}	215	}
216		216
		217	+ const auto offset = context_->GetAttrs()->GetAttrPointer<int64_t>(OFFSET_X_INDEX);
		218	+ runInfo_.offsetX = (offset != nullptr) ? static_cast<int8_t>(*offset) : 0;
		219	+
217	blockSize_ = BYTE_BLOCK / runInfo_.b_dtype_bytes;	220	blockSize_ = BYTE_BLOCK / runInfo_.b_dtype_bytes;
218	dtypeByteL0a_ = runInfo_.a_dtype_bytes;	221	dtypeByteL0a_ = runInfo_.a_dtype_bytes;
219	dtypeByteL0b_ = runInfo_.b_dtype_bytes;	222	dtypeByteL0b_ = runInfo_.b_dtype_bytes;
	@@ -379,7 +382,7 @@ DtypeFlags Conv3DBackpropInputV2TilingArch35::ComputeDtypeFlags(const ge::DataTy
379	flags.bf16flag = outputBackpropDtype == ge::DT_BF16 && filterDtype == ge::DT_BF16 && yDtype == ge::DT_BF16;	382	flags.bf16flag = outputBackpropDtype == ge::DT_BF16 && filterDtype == ge::DT_BF16 && yDtype == ge::DT_BF16;
380	flags.f16flag = outputBackpropDtype == ge::DT_FLOAT16 && filterDtype == ge::DT_FLOAT16 && yDtype == ge::DT_FLOAT16;	383	flags.f16flag = outputBackpropDtype == ge::DT_FLOAT16 && filterDtype == ge::DT_FLOAT16 && yDtype == ge::DT_FLOAT16;
381	flags.f32flag = outputBackpropDtype == ge::DT_FLOAT && filterDtype == ge::DT_FLOAT && yDtype == ge::DT_FLOAT;	384	flags.f32flag = outputBackpropDtype == ge::DT_FLOAT && filterDtype == ge::DT_FLOAT && yDtype == ge::DT_FLOAT;
382	- flags.int8flag = outputBackpropDtype == ge::DT_INT8 && filterDtype == ge::DT_INT8 && yDtype == ge::DT_FLOAT16;	385	+ flags.int8flag = outputBackpropDtype == ge::DT_INT8 && filterDtype == ge::DT_INT8 && (yDtype == ge::DT_FLOAT16 \|\| yDtype == ge::DT_INT8);
383	return flags;	386	return flags;
384	}	387	}
385		388
	@@ -1047,6 +1050,7 @@ void Conv3DBackpropInputV2TilingArch35::SetRunInfoTiling(conv_bp_v2_kernel::TCon
1047	dxt.singleIterateDk = singleIterateDk_;	1050	dxt.singleIterateDk = singleIterateDk_;
1048	dxt.enRelu = runInfo_.enRelu;	1051	dxt.enRelu = runInfo_.enRelu;
1049	dxt.quantMode = runInfo_.quantMode;	1052	dxt.quantMode = runInfo_.quantMode;
		1053	+ dxt.offsetX = runInfo_.offsetX;
1050	}	1054	}
1051		1055
1052	void Conv3DBackpropInputV2TilingArch35::SetDxTilingFromTbeTiling()	1056	void Conv3DBackpropInputV2TilingArch35::SetDxTilingFromTbeTiling()

Mconv/conv3d_backprop_input_v2/op_host/op_tiling/arch35/conv3d_backprop_input_v2_inner_product_tiling.cpp+3-0

	@@ -107,6 +107,9 @@ ge::graphStatus Conv3DDXV2InnerProductTiling::GetPublicShapeAttrsInfo()
107	}	107	}
108	}	108	}
109		109
		110	+ const auto offset = context_->GetAttrs()->GetAttrPointer<int64_t>(OFFSET_X_INDEX);
		111	+ runInfo_.offsetX = (offset != nullptr) ? static_cast<int8_t>(*offset) : 0;
		112	+
110	blockSize_ = BYTE_BLOCK / runInfo_.b_dtype_bytes;	113	blockSize_ = BYTE_BLOCK / runInfo_.b_dtype_bytes;
111	dtypeByteL0a_ = runInfo_.a_dtype_bytes;	114	dtypeByteL0a_ = runInfo_.a_dtype_bytes;
112	dtypeByteL0b_ = runInfo_.b_dtype_bytes;	115	dtypeByteL0b_ = runInfo_.b_dtype_bytes;

Mconv/conv3d_backprop_input_v2/op_host/op_tiling/arch35/conv3d_backprop_input_v2_kernel_split_tiling.cpp+5-0

	@@ -423,6 +423,11 @@ bool Conv3DDXV2KernelSplitTiling::CheckKernelSplitEnable()
423	return false;	423	return false;
424	}	424	}
425		425
		426	+ // kernel拆分不支持offsetX
		427	+ if (runInfo_.offsetX != 0) {
		428	+ return false;
		429	+ }
		430	+
426	constexpr uint32_t bestBaseMN = 256;	431	constexpr uint32_t bestBaseMN = 256;
427		432
428	if (TryKernelSplitHW(bestBaseMN)) {	433	if (TryKernelSplitHW(bestBaseMN)) {

Mconv/conv3d_backprop_input_v2/op_host/op_tiling/common/conv_backprop_input_context_utils.cpp+14-11

	@@ -180,6 +180,11 @@ static bool IsArchAfter35(const gert::TilingContext *context) {
180	NpuArch::DAV_3510;	180	NpuArch::DAV_3510;
181	}	181	}
182		182
		183	+static bool IsSupportedDtypeForOutputPadding(const ge::DataType dtype) {
		184	+ return dtype == ge::DT_BF16 \|\| dtype == ge::DT_FLOAT16 \|\|
		185	+ dtype == ge::DT_FLOAT \|\| dtype == ge::DT_INT8;
		186	+}
		187	+
183	bool ValidateConvBackpropContext(const gert::TilingContext *context) {	188	bool ValidateConvBackpropContext(const gert::TilingContext *context) {
184	// 校验输入张量描述是否获取成功	189	// 校验输入张量描述是否获取成功
185	auto input_size_desc = context->GetInputDesc(INPUT_SIZE_INDEX);	190	auto input_size_desc = context->GetInputDesc(INPUT_SIZE_INDEX);
	@@ -421,7 +426,9 @@ bool CheckTransposeAttr(gert::TilingContext *context, OtherParams& otherParams)
421	if (attrs->GetAttrNum() > K_OFFSET_X_CONV3D_TRANSPOSE_IDX) {	426	if (attrs->GetAttrNum() > K_OFFSET_X_CONV3D_TRANSPOSE_IDX) {
422	const auto offsetX = attrs->GetAttrPointer<int64_t>(K_OFFSET_X_CONV3D_TRANSPOSE_IDX);	427	const auto offsetX = attrs->GetAttrPointer<int64_t>(K_OFFSET_X_CONV3D_TRANSPOSE_IDX);
423	OP_CHECK_IF(offsetX == nullptr, OP_LOGE(context, "failed to get offsetX attrs"), return false);	428	OP_CHECK_IF(offsetX == nullptr, OP_LOGE(context, "failed to get offsetX attrs"), return false);
424	- OP_CHECK_IF(offsetX != 0, OP_LOGE(context, "offsetX:%ld is invalid, it should be 0", offsetX), return false);	429	+ if (!IsArchAfter35(context) && !IsSocVersionFuse(context)) {
		430	+ OP_CHECK_IF(offsetX != 0, OP_LOGE(context, "offsetX:%ld is invalid, it should be 0", offsetX), return false);
		431	+ }
425	}	432	}
426	if (IsSocVersionFuse(context)) {	433	if (IsSocVersionFuse(context)) {
427	if (attrs->GetAttrNum() > K_FUSION_MODE_CONV3D_TRANSPOSE_IDX) {	434	if (attrs->GetAttrNum() > K_FUSION_MODE_CONV3D_TRANSPOSE_IDX) {
	@@ -527,7 +534,7 @@ static bool UpdateDtypeParams(const gert::TilingContext *context, Conv3dBpInputV
527	otherParams.b_dtype == ge::DT_FLOAT8_E4M3FN &&	534	otherParams.b_dtype == ge::DT_FLOAT8_E4M3FN &&
528	otherParams.c_dtype == ge::DT_FLOAT8_E4M3FN;	535	otherParams.c_dtype == ge::DT_FLOAT8_E4M3FN;
529	bool isInt8Flag = otherParams.a_dtype == ge::DT_INT8 &&	536	bool isInt8Flag = otherParams.a_dtype == ge::DT_INT8 &&
530	- otherParams.b_dtype == ge::DT_INT8 && otherParams.c_dtype == ge::DT_FLOAT16;	537	+ otherParams.b_dtype == ge::DT_INT8 && (otherParams.c_dtype == ge::DT_FLOAT16 \|\| otherParams.c_dtype == ge::DT_INT8);
531	dtypeSupportFlag = isHiF8Flag \|\| isFp8E4M3Flag \|\| isFp16Flag \|\| isFp32Flag \|\| isInt8Flag;	538	dtypeSupportFlag = isHiF8Flag \|\| isFp8E4M3Flag \|\| isFp16Flag \|\| isFp32Flag \|\| isInt8Flag;
532	dtypeCheckLog = "hifloat8, float8_e4m3, int8, fp16 and fp32";	539	dtypeCheckLog = "hifloat8, float8_e4m3, int8, fp16 and fp32";
533	}	540	}
	@@ -1868,14 +1875,10 @@ bool CheckTranspose(const char* opName, const gert::TilingContext* context) {
1868	outputPaddingValue.push_back(outputPaddingData[index]);	1875	outputPaddingValue.push_back(outputPaddingData[index]);
1869	}	1876	}
1870	OP_CHECK_IF((!outputPaddingAllzero) &&	1877	OP_CHECK_IF((!outputPaddingAllzero) &&
1871	- ((context->GetInputDesc(FILTER_INDEX)->GetDataType() != ge::DT_BF16 &&	1878	+ (!IsSupportedDtypeForOutputPadding(context->GetInputDesc(FILTER_INDEX)->GetDataType()) \|\|
1872	- context->GetInputDesc(FILTER_INDEX)->GetDataType() != ge::DT_FLOAT16 &&	1879	+ !IsSupportedDtypeForOutputPadding(context->GetInputDesc(OUT_BACKPROP_INDEX)->GetDataType())),
1873	- context->GetInputDesc(FILTER_INDEX)->GetDataType() != ge::DT_FLOAT) \|\|
1874	- (context->GetInputDesc(OUT_BACKPROP_INDEX)->GetDataType() != ge::DT_BF16 &&
1875	- context->GetInputDesc(OUT_BACKPROP_INDEX)->GetDataType() != ge::DT_FLOAT16 &&
1876	- context->GetInputDesc(OUT_BACKPROP_INDEX)->GetDataType() != ge::DT_FLOAT)),
1877	CUBE_INNER_ERR_REPORT(opName,	1880	CUBE_INNER_ERR_REPORT(opName,
1878	- "when output_padding[%s] is not all zero, op only supports bfloat16, float16 and float32 for all inputs, get filter dtype[%s], output backprop dtype[%s]",	1881	+ "when output_padding[%s] is not all zero, op only supports bfloat16, float16, float32 and int8 for all inputs, get filter dtype[%s], output backprop dtype[%s]",
1879	DebugString(outputPaddingValue).c_str(),	1882	DebugString(outputPaddingValue).c_str(),
1880	ge::TypeUtils::DataTypeToSerialString(context->GetInputDesc(FILTER_INDEX)->GetDataType()).c_str(),	1883	ge::TypeUtils::DataTypeToSerialString(context->GetInputDesc(FILTER_INDEX)->GetDataType()).c_str(),
1881	ge::TypeUtils::DataTypeToSerialString(context->GetInputDesc(OUT_BACKPROP_INDEX)->GetDataType()).c_str()),	1884	ge::TypeUtils::DataTypeToSerialString(context->GetInputDesc(OUT_BACKPROP_INDEX)->GetDataType()).c_str()),
	@@ -1885,9 +1888,9 @@ bool CheckTranspose(const char* opName, const gert::TilingContext* context) {
1885	"output_padding[%s] contains negative values, op only supports all non-negative inputs.",	1888	"output_padding[%s] contains negative values, op only supports all non-negative inputs.",
1886	DebugString(outputPaddingValue).c_str()), return false);	1889	DebugString(outputPaddingValue).c_str()), return false);
1887	}	1890	}
1888	- OP_CHECK_IF(offsetX != nullptr && *offsetX != 0,
1889	- OP_LOGE(opName, "cannot support offset_x attribute parameters"), return false);
1890	if (!IsSocVersionFuse(context) && !IsArchAfter35(context)) {	1891	if (!IsSocVersionFuse(context) && !IsArchAfter35(context)) {
		1892	+ OP_CHECK_IF(offsetX != nullptr && *offsetX != 0,
		1893	+ OP_LOGE(opName, "cannot support offset_x attribute parameters"), return false);
1891	OP_CHECK_IF(offsetWShape != nullptr && offsetWShape->GetStorageShape().GetShapeSize() != 0,	1894	OP_CHECK_IF(offsetWShape != nullptr && offsetWShape->GetStorageShape().GetShapeSize() != 0,
1892	OP_LOGE(opName,"cannot support offset_w input parameters"), return false);	1895	OP_LOGE(opName,"cannot support offset_w input parameters"), return false);
1893	auto biasShape = context->GetOptionalInputShape(BAIS_INDEX);	1896	auto biasShape = context->GetOptionalInputShape(BAIS_INDEX);

Mconv/conv3d_backprop_input_v2/op_host/op_tiling/common/conv_backprop_input_context_utils.h+1-0

	@@ -91,6 +91,7 @@ struct Conv3dBpInputV2RunInfo {
91	int32_t initOutputFlag = 0;	91	int32_t initOutputFlag = 0;
92	uint8_t enRelu = 0;	92	uint8_t enRelu = 0;
93	uint8_t quantMode = 0;	93	uint8_t quantMode = 0;
		94	+ int8_t offsetX = 0;
94		95
95	ge::Format outBackpropFormat = ge::FORMAT_NCDHW;	96	ge::Format outBackpropFormat = ge::FORMAT_NCDHW;
96	ge::Format filterFormat = ge::FORMAT_NCDHW;	97	ge::Format filterFormat = ge::FORMAT_NCDHW;

Mconv/conv3d_backprop_input_v2/op_kernel/arch35/conv3d_backprop_input_v2/conv3d_backprop_input_v2_tiling_data.h+1-0

	@@ -88,6 +88,7 @@ struct TConv3DInputV2Tiling {
88	uint64_t kValueSegment = 0;	88	uint64_t kValueSegment = 0;
89	bool enableSplitK = false;	89	bool enableSplitK = false;
90	bool useUbAccumForSplitK = false;	90	bool useUbAccumForSplitK = false;
		91	+ int8_t offsetX = 0;
91	};	92	};
92		93
93	struct Conv3DBackpropInputV2Params {	94	struct Conv3DBackpropInputV2Params {

Mconv/conv3d_backprop_input_v2/op_kernel/arch35/convolution_3d_backprop/impl/conv_bp_sub_func.h+5-2

	@@ -120,6 +120,9 @@ static __aicore__ inline void InitLoadToA2Params(Intf *self)
120	self->ctx.load3d_.fMatrixCtrl = 0;	120	self->ctx.load3d_.fMatrixCtrl = 0;
121	// l1 only cut cout in k	121	// l1 only cut cout in k
122	self->ctx.load3d_.channelSize = self->ctx.channelSize_;	122	self->ctx.load3d_.channelSize = self->ctx.channelSize_;
		123	+ if constexpr (std::is_same<typename Intf::L0cT, int32_t>::value) {
		124	+ self->ctx.load3d_.padValue = self->ctx.tiling_->offsetX;
		125	+ }
123	}	126	}
124		127
125	template <class Intf>	128	template <class Intf>
	@@ -469,8 +472,8 @@ static __aicore__ inline void LoadL0c2GmForNz2Dn(Intf *self, const GlobalTensor<
469	fixPipeParams.nSize = self->ctx.baseUseN_; // N: cin	472	fixPipeParams.nSize = self->ctx.baseUseN_; // N: cin
470	// loop1_src_stride, c0_size, cin1	473	// loop1_src_stride, c0_size, cin1
471	fixPipeParams.srcStride = AlignUp16(self->ctx.baseUseM_); // src N stride, loop1_src_stride (unit: 32B)	474	fixPipeParams.srcStride = AlignUp16(self->ctx.baseUseM_); // src N stride, loop1_src_stride (unit: 32B)
472	-#if (__NPU_ARCH__ == 5102)
473	fixPipeParams.reluEn = self->ctx.tiling_->enRelu;	475	fixPipeParams.reluEn = self->ctx.tiling_->enRelu;
		476	+#if (__NPU_ARCH__ == 5102)
474	fixPipeParams.preReluMode = static_cast<ReluMode>(self->ctx.tiling_->enRelu);	477	fixPipeParams.preReluMode = static_cast<ReluMode>(self->ctx.tiling_->enRelu);
475	#endif	478	#endif
476	uint64_t dstOffset = ComputeDstOffset(self, fixPipeParams);	479	uint64_t dstOffset = ComputeDstOffset(self, fixPipeParams);
	@@ -618,8 +621,8 @@ static __aicore__ inline void LoadL0c2GmForKernelSplitH(Intf *self, const Global
618	fixPipeParams.srcStride = AlignUp16(self->ctx.baseUseM_); // src N stride, loop1_src_stride (unit: 32B)	621	fixPipeParams.srcStride = AlignUp16(self->ctx.baseUseM_); // src N stride, loop1_src_stride (unit: 32B)
619	// loop2_dst_stride, element, c	622	// loop2_dst_stride, element, c
620	fixPipeParams.dstStride = self->ctx.diHiWi_; // dst N stride, loop2_dst_stride (unit: element)	623	fixPipeParams.dstStride = self->ctx.diHiWi_; // dst N stride, loop2_dst_stride (unit: element)
621	-#if (__NPU_ARCH__ == 5102)
622	fixPipeParams.reluEn = self->ctx.tiling_->enRelu;	624	fixPipeParams.reluEn = self->ctx.tiling_->enRelu;
		625	+#if (__NPU_ARCH__ == 5102)
623	fixPipeParams.preReluMode = static_cast<ReluMode>(self->ctx.tiling_->enRelu);	626	fixPipeParams.preReluMode = static_cast<ReluMode>(self->ctx.tiling_->enRelu);
624	#endif	627	#endif
625	int64_t srcOffset = 0;	628	int64_t srcOffset = 0;

Mconv/conv3d_backprop_input_v2/op_kernel/arch35/convolution_3d_backprop/impl/conv_bp_sub_func_load_gm_to_l1.h+17-10

	@@ -26,21 +26,28 @@ using AscendC::Nd2NzParams;
26		26
27	namespace Convolution3DBackpropFunc {	27	namespace Convolution3DBackpropFunc {
28	template <class Intf, typename SrcType>	28	template <class Intf, typename SrcType>
29	-__aicore__ inline void InitZeroValue(Intf *self, const LocalTensor<SrcType> &buf)	29	+__aicore__ inline void InitZeroValue(Intf *self, const LocalTensor<SrcType> &buf, bool useOffsetX = false)
30	{	30	{
31	uint32_t len = buf.GetSize() * sizeof(SrcType);	31	uint32_t len = buf.GetSize() * sizeof(SrcType);
		32	+ uint16_t padValue = 0;
		33	+ if constexpr(std::is_same<SrcType, int8_t>::value) {
		34	+ if (useOffsetX) {
		35	+ uint8_t offsetX = static_cast<uint8_t>(self->ctx.tiling_->offsetX);
		36	+ padValue = (static_cast<uint16_t>(offsetX)) << 8 \| (static_cast<uint16_t>(offsetX));
		37	+ }
		38	+ }
32	if constexpr(std::is_same<SrcType, hifloat8_t>::value \|\|	39	if constexpr(std::is_same<SrcType, hifloat8_t>::value \|\|
33	std::is_same<SrcType, fp8_e4m3fn_t>::value \|\|	40	std::is_same<SrcType, fp8_e4m3fn_t>::value \|\|
34	std::is_same<SrcType, int8_t>::value) {	41	std::is_same<SrcType, int8_t>::value) {
35	- InitConstValue(buf.template ReinterpretCast<uint16_t>(), {1, static_cast<uint16_t>(len >> 5), 0, 0});	42	+ InitConstValue(buf.template ReinterpretCast<uint16_t>(), {1, static_cast<uint16_t>(len >> 5), 0, padValue});
36	- } else {	43	+ } else {
37	- AscendC::InitConstValueParams<SrcType> initConstValueParams;	44	+ AscendC::InitConstValueParams<SrcType> initConstValueParams;
38	- initConstValueParams.repeatTimes = 1;	45	+ initConstValueParams.repeatTimes = 1;
39	- initConstValueParams.blockNum = len >> 5; // 除以blockSize	46	+ initConstValueParams.blockNum = len >> 5;
40	- initConstValueParams.dstGap = 0;	47	+ initConstValueParams.dstGap = 0;
41	- initConstValueParams.initValue = (SrcType)(0);	48	+ initConstValueParams.initValue = (SrcType)(0);
42	- InitConstValue(buf, initConstValueParams);	49	+ InitConstValue(buf, initConstValueParams);
43	- }	50	+ }
44	PipeBarrier<PIPE_MTE2>();	51	PipeBarrier<PIPE_MTE2>();
45	}	52	}
46		53

Mconv/conv3d_backprop_input_v2/op_kernel/arch35/convolution_3d_backprop/impl/conv_bp_sub_func_load_gm_to_l1a.h+2-2

	@@ -395,7 +395,7 @@ __aicore__ inline void LoadToA1ForDn2Nz(Intf *self, LocalTensor<typename Intf::S
395	strideH = 1;	395	strideH = 1;
396	}	396	}
397	if (unlikely(self->ctx.tiling_->strideW * strideH > 1)) {	397	if (unlikely(self->ctx.tiling_->strideW * strideH > 1)) {
398	- InitZeroValue<Intf, typename Intf::SrcAT>(self, useA1Buf);	398	+ InitZeroValue<Intf, typename Intf::SrcAT>(self, useA1Buf, true);
399	}	399	}
400	CalcLoadToA1Dn2NzParams<Intf, typename Intf::SrcAT>(self, dn2NzParams, out2A1DstAddrOffset, curCoutIdx, kIdx);	400	CalcLoadToA1Dn2NzParams<Intf, typename Intf::SrcAT>(self, dn2NzParams, out2A1DstAddrOffset, curCoutIdx, kIdx);
401	if (strideH > 1) {	401	if (strideH > 1) {
	@@ -442,7 +442,7 @@ __aicore__ inline void LoadToA1ForNd2Nz(Intf *self, LocalTensor<typename Intf::S
442	strideH = 1;	442	strideH = 1;
443	}	443	}
444	if (unlikely(self->ctx.tiling_->strideW * strideH > 1)) {	444	if (unlikely(self->ctx.tiling_->strideW * strideH > 1)) {
445	- InitZeroValue<Intf, typename Intf::SrcAT>(self, useA1Buf);	445	+ InitZeroValue<Intf, typename Intf::SrcAT>(self, useA1Buf, true);
446	}	446	}
447	uint32_t curCoutIdx = 0;	447	uint32_t curCoutIdx = 0;
448	if constexpr (!Intf::conv3dConfig.enableC04Flag) {	448	if constexpr (!Intf::conv3dConfig.enableC04Flag) {

Mconv/conv3d_backprop_input_v2/op_kernel/arch35/convolution_3d_backprop/impl/conv_bp_sub_func_mix.h+1-1

	@@ -474,8 +474,8 @@ static __aicore__ inline void LoadL0c2OutForKernelSplitHW(Intf *self, const Loca
474	SetFixPipeQuantVal<Intf>(self, fixPipeParams);	474	SetFixPipeQuantVal<Intf>(self, fixPipeParams);
475	fixPipeParams.params.srcNzC0Stride = 1; // src M stride, loop0_src_stride (unit: 32B)	475	fixPipeParams.params.srcNzC0Stride = 1; // src M stride, loop0_src_stride (unit: 32B)
476	fixPipeParams.nSize = self->ctx.baseUseN_; // N: cin	476	fixPipeParams.nSize = self->ctx.baseUseN_; // N: cin
477	-#if (__NPU_ARCH__ == 5102)
478	fixPipeParams.reluEn = self->ctx.tiling_->enRelu;	477	fixPipeParams.reluEn = self->ctx.tiling_->enRelu;
		478	+#if (__NPU_ARCH__ == 5102)
479	fixPipeParams.preReluMode = static_cast<ReluMode>(self->ctx.tiling_->enRelu);	479	fixPipeParams.preReluMode = static_cast<ReluMode>(self->ctx.tiling_->enRelu);
480	#endif	480	#endif
481	// loop1_src_stride, c0_size, cin1	481	// loop1_src_stride, c0_size, cin1

Mconv/extend_conv_transpose/op_host/extend_conv_transpose_def.cpp+3-3

	@@ -53,8 +53,8 @@ public:
53	aicore_config_950.Input("filter")	53	aicore_config_950.Input("filter")
54	.ParamType(REQUIRED)	54	.ParamType(REQUIRED)
55	.DataType({ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8})	55	.DataType({ge::DT_INT8, ge::DT_INT8, ge::DT_INT8, ge::DT_INT8})
56	- .Format({ge::FORMAT_NCDHW, ge::FORMAT_NCDHW, ge::FORMAT_NDHWC, ge::FORMAT_NDHWC})	56	+ .Format({ge::FORMAT_NCDHW, ge::FORMAT_NCDHW, ge::FORMAT_NCDHW, ge::FORMAT_NCDHW})
57	- .UnknownShapeFormat({ge::FORMAT_NCDHW, ge::FORMAT_NCDHW, ge::FORMAT_NDHWC, ge::FORMAT_NDHWC});	57	+ .UnknownShapeFormat({ge::FORMAT_NCDHW, ge::FORMAT_NCDHW, ge::FORMAT_NCDHW, ge::FORMAT_NCDHW});
58	aicore_config_950.Input("bias")	58	aicore_config_950.Input("bias")
59	.ParamType(OPTIONAL)	59	.ParamType(OPTIONAL)
60	.DataType({ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32})	60	.DataType({ge::DT_INT32, ge::DT_INT32, ge::DT_INT32, ge::DT_INT32})
	@@ -67,7 +67,7 @@ public:
67	.UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});	67	.UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
68	aicore_config_950.Output("y")	68	aicore_config_950.Output("y")
69	.ParamType(REQUIRED)	69	.ParamType(REQUIRED)
70	- .DataType({ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_FLOAT16})	70	+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT16, ge::DT_INT8, ge::DT_INT8})
71	.Format({ge::FORMAT_NCDHW, ge::FORMAT_NCDHW, ge::FORMAT_NCDHW, ge::FORMAT_NCDHW})	71	.Format({ge::FORMAT_NCDHW, ge::FORMAT_NCDHW, ge::FORMAT_NCDHW, ge::FORMAT_NCDHW})
72	.UnknownShapeFormat({ge::FORMAT_NCDHW, ge::FORMAT_NCDHW, ge::FORMAT_NCDHW, ge::FORMAT_NCDHW});	72	.UnknownShapeFormat({ge::FORMAT_NCDHW, ge::FORMAT_NCDHW, ge::FORMAT_NCDHW, ge::FORMAT_NCDHW});
73	this->AICore().AddConfig("ascend950", aicore_config_950);	73	this->AICore().AddConfig("ascend950", aicore_config_950);

Mconv/extend_conv_transpose/op_kernel/extend_conv_transpose.cpp+10-1

	@@ -51,7 +51,7 @@ __global__ __aicore__ void extend_conv_transpose(GM_ADDR input_size, GM_ADDR x,
51		51
52	#if (__NPU_ARCH__ == 5102)	52	#if (__NPU_ARCH__ == 5102)
53	KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIC_ONLY);	53	KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIC_ONLY);
54	-#elif (__NPU_ARCH__ == 3510)	54	+#else
55	KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIC_1_2);	55	KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIC_1_2);
56	#endif	56	#endif
57		57
	@@ -62,10 +62,19 @@ __global__ __aicore__ void extend_conv_transpose(GM_ADDR input_size, GM_ADDR x,
62	opInitOutput.Destroy();	62	opInitOutput.Destroy();
63	}	63	}
64		64
		65	+#if (__NPU_ARCH__ == 5102)
65	if (tilingData.conv3DDxTiling.enableVecTrans) {	66	if (tilingData.conv3DDxTiling.enableVecTrans) {
66	// VecTranspose	67	// VecTranspose
67	EXTEND_CONV_TRANSPOSE_RUN_OP_VECTRANSPOSE(DxVecTranspose::Conv3dDxVecTranspose<DTYPE_FILTER>);	68	EXTEND_CONV_TRANSPOSE_RUN_OP_VECTRANSPOSE(DxVecTranspose::Conv3dDxVecTranspose<DTYPE_FILTER>);
68	}	69	}
		70	+#else
		71	+ if ASCEND_IS_AIV {
		72	+ if (tilingData.conv3DDxTiling.enableVecTrans) {
		73	+ // VecTranspose
		74	+ EXTEND_CONV_TRANSPOSE_RUN_OP_VECTRANSPOSE(DxVecTranspose::Conv3dDxVecTranspose<DTYPE_FILTER>);
		75	+ }
		76	+ }
		77	+#endif
69		78
70	if constexpr (kernelSplitMode != TPL_NO_SPLIT_KERNEL) {	79	if constexpr (kernelSplitMode != TPL_NO_SPLIT_KERNEL) {
71	EXTEND_CONV_TRANSPOSE_RUN_OP(Conv3dDxKsBlock<DTYPE_FILTER, FORMAT_FILTER, DTYPE_X, FORMAT_X, DTYPE_Y, FORMAT_Y,	80	EXTEND_CONV_TRANSPOSE_RUN_OP(Conv3dDxKsBlock<DTYPE_FILTER, FORMAT_FILTER, DTYPE_X, FORMAT_X, DTYPE_Y, FORMAT_Y,