npu arch adapt.-ops-nn-AtomGit

Mconv/common/op_graph/fusion_pass/conv_fusion_utils_pass.cpp+37-8

	@@ -21,6 +21,8 @@ namespace ConvFusionUtils {
21	using namespace ge;	21	using namespace ge;
22	using namespace fe;	22	using namespace fe;
23	using namespace fusion;	23	using namespace fusion;
		24	+using conv_arch::GetNpuArchKey;
		25	+using conv_arch::IsCubeVectorFuseSoc;
24		26
25	bool ConvFusionUtilsPass::AddSubgraphInput(std::unique_ptr<SubgraphBoundary>& boundary, const GNode& node,	27	bool ConvFusionUtilsPass::AddSubgraphInput(std::unique_ptr<SubgraphBoundary>& boundary, const GNode& node,
26	const int64_t subgraphIndex, const int64_t boundaryIndex)	28	const int64_t subgraphIndex, const int64_t boundaryIndex)
	@@ -48,7 +50,8 @@ bool ConvFusionUtilsPass::AddSubgraphOutput(std::unique_ptr<SubgraphBoundary>& b
48	return true;	50	return true;
49	}	51	}
50		52
51	-bool ConvFusionUtilsPass::CheckSocList(const std::map<std::string, NpuArch>& socList, NpuArch& npuArch)	53	+bool ConvFusionUtilsPass::CheckSocList(const std::map<std::string, NpuArch>& socList, NpuArch& npuArch,
		54	+ bool supportFuse)
52	{	55	{
53	PlatformInfo platformInfo;	56	PlatformInfo platformInfo;
54	OptionalInfo optionalInfo;	57	OptionalInfo optionalInfo;
	@@ -57,15 +60,41 @@ bool ConvFusionUtilsPass::CheckSocList(const std::map<std::string, NpuArch>& soc
57	OP_LOGW(UTIL_NAME, "Get platform_info failed."), return false);	60	OP_LOGW(UTIL_NAME, "Get platform_info failed."), return false);
58	const std::string soc = platformInfo.str_info.short_soc_version;	61	const std::string soc = platformInfo.str_info.short_soc_version;
59		62
60	- FUSION_PASS_CHECK(	63	+ auto it = socList.find(soc);
61	- socList.find(soc) == socList.end(),	64	+ if (it != socList.end()) {
62	- OP_LOGD(UTIL_NAME, "Current soc %s not in check list %s.", soc.c_str(), SocListToString(socList).c_str()),	65	+ npuArch = it->second;
63	- return false);	66	+ OP_LOGD(UTIL_NAME, "Current NpuArch is DAV_%u.", npuArch);
		67	+ return true;
		68	+ }
64		69
65	- npuArch = socList.at(soc);	70	+ if (supportFuse) {
66	- OP_LOGD(UTIL_NAME, "Current NpuArch is DAV_%u.", npuArch);	71	+ fe::PlatFormInfos platFormInfos;
		72	+ fe::OptionalInfos optionalInfos;
		73	+ if (PlatformInfoManager::Instance().GetPlatformInfoWithOutSocVersion(platFormInfos, optionalInfos) !=
		74	+ GRAPH_SUCCESS) {
		75	+ OP_LOGW(UTIL_NAME, "Get PlatFormInfos failed.");
		76	+ return false;
		77	+ }
		78	+ if (IsCubeVectorFuseSoc(platFormInfos)) {
		79	+ OP_LOGD(UTIL_NAME, "Current platform is cube_vector_combine=fuse.");
		80	+ return true;
		81	+ }
		82	+ }
67		83
68	- return true;	84	+ OP_LOGD(UTIL_NAME, "Current soc %s not in check list %s.", soc.c_str(), SocListToString(socList).c_str());
		85	+ return false;
		86	+}
		87	+
		88	+const std::string& ConvFusionUtilsPass::GetArchKey()
		89	+{
		90	+ fe::PlatFormInfos platFormInfos;
		91	+ fe::OptionalInfos optionalInfos;
		92	+ if (PlatformInfoManager::Instance().GetPlatformInfoWithOutSocVersion(platFormInfos, optionalInfos) !=
		93	+ GRAPH_SUCCESS) {
		94	+ OP_LOGW(UTIL_NAME, "Get PlatFormInfos failed, fallback to 3510.");
		95	+ return NPU_ARCH_KEY_3510;
		96	+ }
		97	+ return GetNpuArchKey(platFormInfos);
69	}	98	}
70		99
71	bool ConvFusionUtilsPass::GetConvBaseAttr(const GNode& convNode, ConvBaseAttrs& baseAttrs,	100	bool ConvFusionUtilsPass::GetConvBaseAttr(const GNode& convNode, ConvBaseAttrs& baseAttrs,

Mconv/common/op_graph/fusion_pass/conv_fusion_utils_pass.h+8-1

	@@ -23,8 +23,11 @@
23	#include "graph/operator.h"	23	#include "graph/operator.h"
24	#include "graph/utils/type_utils.h"	24	#include "graph/utils/type_utils.h"
25	#include "log/log.h"	25	#include "log/log.h"
		26	+#include "platform/platform_infos_def.h"
26	#include "platform/soc_spec.h"	27	#include "platform/soc_spec.h"
27		28
		29	+#include "../../op_host/conv_npu_arch_resolver.h"
		30	+
28	namespace Ops {	31	namespace Ops {
29	namespace NN {	32	namespace NN {
30	namespace Conv {	33	namespace Conv {
	@@ -66,6 +69,9 @@ const std::set<ge::AscendString> SPECIFIC_PAD_LIST = {"NOTSET", "EXPLICIT"};
66	const std::set<ge::AscendString> CONV_OP_LIST = {CONV2D, CONV2DV2, CONV3D, CONV3DV2, DEPTHWISE_CONV2D};	69	const std::set<ge::AscendString> CONV_OP_LIST = {CONV2D, CONV2DV2, CONV3D, CONV3DV2, DEPTHWISE_CONV2D};
67	const std::vector<int64_t> HF32_PRECISION_MODES_INT = {0x1, 0x2, 0x40};	70	const std::vector<int64_t> HF32_PRECISION_MODES_INT = {0x1, 0x2, 0x40};
68		71
		72	+const std::string NPU_ARCH_KEY_3510 = "3510";
		73	+const std::string NPU_ARCH_KEY_FUSE = "FUSE";
		74	+
69	#define FUSION_PASS_CHECK(condition, log_func, return_expr) \	75	#define FUSION_PASS_CHECK(condition, log_func, return_expr) \
70	static_assert(std::is_same<bool, std::decay<decltype(condition)>::type>::value, "condition should be bool"); \	76	static_assert(std::is_same<bool, std::decay<decltype(condition)>::type>::value, "condition should be bool"); \
71	do { \	77	do { \
	@@ -172,7 +178,8 @@ public:
172	const int64_t subgraphIndex, const int64_t boundaryIndex);	178	const int64_t subgraphIndex, const int64_t boundaryIndex);
173	template <typename T>	179	template <typename T>
174	static bool CheckSupportList(const std::vector<std::vector<T>>& supportLists, const std::vector<T>& curList);	180	static bool CheckSupportList(const std::vector<std::vector<T>>& supportLists, const std::vector<T>& curList);
175	- static bool CheckSocList(const std::map<std::string, NpuArch>& socList, NpuArch& npuArch);	181	+ static bool CheckSocList(const std::map<std::string, NpuArch>& socList, NpuArch& npuArch, bool supportFuse = false);
		182	+ static const std::string& GetArchKey();
176	static bool GetConvBaseAttr(const ge::GNode& convNode, ConvBaseAttrs& baseAttrs, const ConvDescInfo& convDescInfo);	183	static bool GetConvBaseAttr(const ge::GNode& convNode, ConvBaseAttrs& baseAttrs, const ConvDescInfo& convDescInfo);
177	static bool GetConvDescInfo(const ge::GNode& convNode, ConvDescInfo& convDescInfo);	184	static bool GetConvDescInfo(const ge::GNode& convNode, ConvDescInfo& convDescInfo);
178	static bool GetMatchedNodes(const ge::GraphPtr& graph, std::vector<ge::GNode>& matchedNodes,	185	static bool GetMatchedNodes(const ge::GraphPtr& graph, std::vector<ge::GNode>& matchedNodes,

Aconv/common/op_host/conv_npu_arch_resolver.h+66-0

@@ -0,0 +1,66 @@
	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	+#ifndef OPS_NN_CONV_COMMON_OP_HOST_CONV_NPU_ARCH_RESOLVER_H
	12	+#define OPS_NN_CONV_COMMON_OP_HOST_CONV_NPU_ARCH_RESOLVER_H
	13	+
	14	+#include <string>
	15	+#include "platform/platform_infos_def.h"
	16	+
	17	+namespace conv_arch {
	18	+
	19	+inline bool IsCubeVectorFuseSoc(fe::PlatFormInfos& platformInfo)
	20	+{
	21	+ std::string cubeVecState;
	22	+ platformInfo.GetPlatformRes("SoCInfo", "cube_vector_combine", cubeVecState);
	23	+ return cubeVecState == "fuse";
	24	+}
	25	+
	26	+class NpuArchResolver {
	27	+public:
	28	+ virtual ~NpuArchResolver() = default;
	29	+ virtual const std::string& GetArchKey() const = 0;
	30	+ static const NpuArchResolver& GetInstance(fe::PlatFormInfos& platformInfo);
	31	+};
	32	+
	33	+class FuseNpuArchResolver : public NpuArchResolver {
	34	+public:
	35	+ const std::string& GetArchKey() const override
	36	+ {
	37	+ static const std::string key = "FUSE";
	38	+ return key;
	39	+ }
	40	+};
	41	+
	42	+class DefaultNpuArchResolver : public NpuArchResolver {
	43	+public:
	44	+ const std::string& GetArchKey() const override
	45	+ {
	46	+ static const std::string key = "3510";
	47	+ return key;
	48	+ }
	49	+};
	50	+
	51	+inline const NpuArchResolver& NpuArchResolver::GetInstance(fe::PlatFormInfos& platformInfo)
	52	+{
	53	+ static const FuseNpuArchResolver fuseResolver;
	54	+ static const DefaultNpuArchResolver defaultResolver;
	55	+ return IsCubeVectorFuseSoc(platformInfo) ? static_cast<const NpuArchResolver&>(fuseResolver) :
	56	+ static_cast<const NpuArchResolver&>(defaultResolver);
	57	+}
	58	+
	59	+inline const std::string& GetNpuArchKey(fe::PlatFormInfos& platformInfo)
	60	+{
	61	+ return NpuArchResolver::GetInstance(platformInfo).GetArchKey();
	62	+}
	63	+
	64	+} // namespace conv_arch
	65	+
	66	+#endif // OPS_NN_CONV_COMMON_OP_HOST_CONV_NPU_ARCH_RESOLVER_H

Mconv/common/op_host/op_tiling/arch35/conv_api_tiling_base.cpp+3-2

	@@ -24,6 +24,7 @@ namespace conv_tiling {
24	ConvTilingBase::ConvTilingBase(const PlatformInfo& platform)	24	ConvTilingBase::ConvTilingBase(const PlatformInfo& platform)
25	{	25	{
26	platformInfo.npuArch = platform.npuArch;	26	platformInfo.npuArch = platform.npuArch;
		27	+ platformInfo.isCubeVectorFuse = platform.isCubeVectorFuse;
27	platformInfo.l1Size = platform.l1Size;	28	platformInfo.l1Size = platform.l1Size;
28	platformInfo.l0ASize = platform.l0ASize;	29	platformInfo.l0ASize = platform.l0ASize;
29	platformInfo.l0BSize = platform.l0BSize;	30	platformInfo.l0BSize = platform.l0BSize;
	@@ -188,7 +189,7 @@ vector<vector<ConvDtype>> ConvTilingBase::GetSupportedDataTypes() const
188	bool ConvTilingBase::CheckLoad3DLimits()	189	bool ConvTilingBase::CheckLoad3DLimits()
189	{	190	{
190	auto LogHelper = [this](const std::string& paramName, const std::string& actualValue, const std::string& reason) {	191	auto LogHelper = [this](const std::string& paramName, const std::string& actualValue, const std::string& reason) {
191	- if (platformInfo.npuArch == NpuArch::DAV_5102) {	192	+ if (platformInfo.isCubeVectorFuse) {
192	OP_LOGE_FOR_INVALID_SHAPE_WITH_REASON(nodeType.c_str(), paramName.c_str(), actualValue.c_str(),	193	OP_LOGE_FOR_INVALID_SHAPE_WITH_REASON(nodeType.c_str(), paramName.c_str(), actualValue.c_str(),
193	reason.c_str());	194	reason.c_str());
194	} else {	195	} else {
	@@ -248,4 +249,4 @@ bool ConvTilingBase::CheckLoad3DLimits()
248	}	249	}
249	return true;	250	return true;
250	}	251	}
251	-} // namespace conv_tiling	252	+} // namespace conv_tiling

Mconv/common/op_host/op_tiling/arch35/conv_api_tiling_util.h+2-1

	@@ -68,6 +68,7 @@ struct PlatformInfo {
68	uint64_t fbSize = 0;	68	uint64_t fbSize = 0;
69	uint64_t aivPerAic = 0;	69	uint64_t aivPerAic = 0;
70	NpuArch npuArch = NpuArch::DAV_RESV;	70	NpuArch npuArch = NpuArch::DAV_RESV;
		71	+ bool isCubeVectorFuse = false;
71	};	72	};
72		73
73	enum class BoundType { CUBE_BOUND = 0, MEMORY_BOUND, INVALID };	74	enum class BoundType { CUBE_BOUND = 0, MEMORY_BOUND, INVALID };
	@@ -309,4 +310,4 @@ uint64_t DivideAndAlign(uint64_t num, uint64_t b, uint64_t c);
309	uint64_t Lcm(const uint64_t valueA, const uint64_t valueB);	310	uint64_t Lcm(const uint64_t valueA, const uint64_t valueB);
310	} // namespace conv_tiling	311	} // namespace conv_tiling
311		312
312	-#endif // ASCENDC_TILING_CONV_API_TILING_UTIL_H	313	+#endif // ASCENDC_TILING_CONV_API_TILING_UTIL_H

Mconv/common/op_host/op_tiling/arch35/conv_base.cpp+12-15

	@@ -190,23 +190,19 @@ void GetSupportedDataTypes(bool hasBias, bool quantFlag, std::vector<std::vector
190	}	190	}
191	}	191	}
192		192
193	-void GetSupportedDataTypes(const NpuArch& socVersion, bool quantFlag, ge::Format fMapFormat, bool exendConvFlag,	193	+void GetSupportedDataTypes(const std::string& archKey, bool quantFlag, ge::Format fMapFormat, bool exendConvFlag,
194	std::vector<std::vector<ge::DataType>>& supportTypes)	194	std::vector<std::vector<ge::DataType>>& supportTypes)
195	{	195	{
196	if (exendConvFlag) {	196	if (exendConvFlag) {
197	- if (fMapFormat == ge::Format::FORMAT_NCHW &&	197	+ if (fMapFormat == ge::Format::FORMAT_NCHW) {
198	- SOC_EXTENDCONV_SUPPORTED_TYPES_NCHW.find(socVersion) != SOC_EXTENDCONV_SUPPORTED_TYPES_NCHW.end()) {	198	+ supportTypes = EXTENDCONV_SUPPORTED_TYPES_NCHW_MAP.at(archKey);
199	- supportTypes = SOC_EXTENDCONV_SUPPORTED_TYPES_NCHW.at(socVersion);	199	+ } else if (fMapFormat == ge::Format::FORMAT_NHWC) {
200	- } else if (fMapFormat == ge::Format::FORMAT_NHWC &&	200	+ supportTypes = EXTENDCONV_SUPPORTED_TYPES_NHWC_MAP.at(archKey);
201	- SOC_EXTENDCONV_SUPPORTED_TYPES_NHWC.find(socVersion) != SOC_EXTENDCONV_SUPPORTED_TYPES_NHWC.end()) {
202	- supportTypes = SOC_EXTENDCONV_SUPPORTED_TYPES_NHWC.at(socVersion);
203	}	201	}
204	} else if (quantFlag) {	202	} else if (quantFlag) {
205	supportTypes = QUANTCONV_SUPPORTED_TYPES;	203	supportTypes = QUANTCONV_SUPPORTED_TYPES;
206	} else {	204	} else {
207	- if (SOC_CONV_SUPPORTED_TYPES.find(socVersion) != SOC_CONV_SUPPORTED_TYPES.end()) {	205	+ supportTypes = CONV_SUPPORTED_TYPES_MAP.at(archKey);
208	- supportTypes = SOC_CONV_SUPPORTED_TYPES.at(socVersion);
209	- }
210	}	206	}
211	}	207	}
212		208
	@@ -378,20 +374,21 @@ void ConvBase::GetSupportedFormats(bool quantFlag, bool is2dFlag, std::stringstr
378	auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);	374	auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
379	if (ascendcPlatform.GetCurNpuArch() == NpuArch::DAV_3510) {	375	if (ascendcPlatform.GetCurNpuArch() == NpuArch::DAV_3510) {
380	if (extendConvFlag) {	376	if (extendConvFlag) {
381	- supportFormats = EXTENDCONV2D_SUPPORT_FORMAT_LIST;	377	+ supportFormats = EXTENDCONV2D_SUPPORT_FORMAT_LIST_MAP.at(NPU_ARCH_KEY_3510);
382	} else if (quantFlag) {	378	} else if (quantFlag) {
383	supportFormats = is2dFlag ? SUPPORT_QUANT_CONV2D_FORMAT_LIST : SUPPORT_QUANT_CONV3D_FORMAT_LIST;	379	supportFormats = is2dFlag ? SUPPORT_QUANT_CONV2D_FORMAT_LIST : SUPPORT_QUANT_CONV3D_FORMAT_LIST;
384	} else if (!quantFlag && (descInfo_.fMapDtype != ge::DataType::DT_HIFLOAT8 \|\|	380	} else if (!quantFlag && (descInfo_.fMapDtype != ge::DataType::DT_HIFLOAT8 \|\|
385	descInfo_.weightDtype != ge::DataType::DT_HIFLOAT8)) {	381	descInfo_.weightDtype != ge::DataType::DT_HIFLOAT8)) {
386	- supportFormats = is2dFlag ? SUPPORT_CONV2D_FORMAT_LIST : SUPPORT_CONV3D_FORMAT_LIST;	382	+ supportFormats = is2dFlag ? SUPPORT_CONV2D_FORMAT_LIST_MAP.at(NPU_ARCH_KEY_3510) :
		383	+ SUPPORT_CONV3D_FORMAT_LIST;
387	} else {	384	} else {
388	supportFormats = is2dFlag ? SUPPORT_CONV2D_DEFAULT_FORMAT_LIST : SUPPORT_CONV3D_DEFAULT_FORMAT_LIST;	385	supportFormats = is2dFlag ? SUPPORT_CONV2D_DEFAULT_FORMAT_LIST : SUPPORT_CONV3D_DEFAULT_FORMAT_LIST;
389	}	386	}
390	- } else if (ascendcPlatform.GetCurNpuArch() == NpuArch::DAV_5102) {	387	+ } else if (platformInfoPtr != nullptr && IsCubeVectorFuseSoc(*platformInfoPtr)) {
391	if (extendConvFlag) {	388	if (extendConvFlag) {
392	- supportFormats = EXTENDCONV2D_SUPPORT_FORMAT_LIST_MDC;	389	+ supportFormats = EXTENDCONV2D_SUPPORT_FORMAT_LIST_MAP.at(NPU_ARCH_KEY_FUSE);
393	} else {	390	} else {
394	- supportFormats = SUPPORT_CONV2D_FORMAT_LIST_MDC;	391	+ supportFormats = SUPPORT_CONV2D_FORMAT_LIST_MAP.at(NPU_ARCH_KEY_FUSE);
395	}	392	}
396	} else {	393	} else {
397	supportFormats = is2dFlag ? SUPPORT_CONV2D_DEFAULT_FORMAT_LIST : SUPPORT_CONV3D_DEFAULT_FORMAT_LIST;	394	supportFormats = is2dFlag ? SUPPORT_CONV2D_DEFAULT_FORMAT_LIST : SUPPORT_CONV3D_DEFAULT_FORMAT_LIST;

Mconv/common/op_host/op_tiling/arch35/conv_base.h+13-12

	@@ -97,11 +97,12 @@ const std::vector<std::vector<ge::DataType>> CONV_SUPPORTED_TYPES_WITHOUT_BIAS_D
97	{ge::DataType::DT_HIFLOAT8, ge::DataType::DT_HIFLOAT8, ge::DataType::DT_HIFLOAT8}};	97	{ge::DataType::DT_HIFLOAT8, ge::DataType::DT_HIFLOAT8, ge::DataType::DT_HIFLOAT8}};
98		98
99	// [fmap, weight, output, bias]	99	// [fmap, weight, output, bias]
100	-const std::vector<std::vector<ge::DataType>> CONV_SUPPORTED_TYPES_MDC = {	100	+const std::vector<std::vector<ge::DataType>> CONV_SUPPORTED_TYPES_FUSE = {
101	{ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16}};	101	{ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16}};
102		102
103	-const std::map<NpuArch, std::vector<std::vector<ge::DataType>>> SOC_CONV_SUPPORTED_TYPES = {	103	+// arch-keyed dtype support list maps (extend by adding new arch keys)
104	- {NpuArch::DAV_3510, CONV_SUPPORTED_TYPES_DAV}, {NpuArch::DAV_5102, CONV_SUPPORTED_TYPES_MDC}};	104	+const std::map<std::string, std::vector<std::vector<ge::DataType>>> CONV_SUPPORTED_TYPES_MAP = {
		105	+ {NPU_ARCH_KEY_3510, CONV_SUPPORTED_TYPES_DAV}, {NPU_ARCH_KEY_FUSE, CONV_SUPPORTED_TYPES_FUSE}};
105		106
106	// [fmap, weight, output, bias]	107	// [fmap, weight, output, bias]
107	const std::vector<std::vector<ge::DataType>> QUANTCONV_SUPPORTED_TYPES_WITH_BIAS = {	108	const std::vector<std::vector<ge::DataType>> QUANTCONV_SUPPORTED_TYPES_WITH_BIAS = {
	@@ -142,7 +143,7 @@ const std::vector<std::vector<ge::DataType>> QUANTCONV_SUPPORTED_TYPES = {
142	ge::DataType::DT_FLOAT}};	143	ge::DataType::DT_FLOAT}};
143		144
144	// [fmap, weight, output, bias]	145	// [fmap, weight, output, bias]
145	-const std::vector<std::vector<ge::DataType>> EXTENDCONV2D_SUPPORTED_TYPES_MDC = {	146	+const std::vector<std::vector<ge::DataType>> EXTENDCONV2D_SUPPORTED_TYPES_FUSE = {
146	{ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16},	147	{ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16},
147	{ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16, ge::DataType::DT_INT8, ge::DataType::DT_FLOAT16},	148	{ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16, ge::DataType::DT_INT8, ge::DataType::DT_FLOAT16},
148	{ge::DataType::DT_INT8, ge::DataType::DT_INT8, ge::DataType::DT_INT8, ge::DataType::DT_INT32},	149	{ge::DataType::DT_INT8, ge::DataType::DT_INT8, ge::DataType::DT_INT8, ge::DataType::DT_INT32},
	@@ -173,11 +174,11 @@ const std::vector<std::vector<ge::DataType>> EXTENDCONV_SUPPORTED_TYPES_NHWC = {
173	{ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16},	174	{ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16},
174	{ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16, ge::DataType::DT_INT8, ge::DataType::DT_FLOAT16}};	175	{ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16, ge::DataType::DT_INT8, ge::DataType::DT_FLOAT16}};
175		176
176	-const std::map<NpuArch, std::vector<std::vector<ge::DataType>>> SOC_EXTENDCONV_SUPPORTED_TYPES_NCHW = {	177	+const std::map<std::string, std::vector<std::vector<ge::DataType>>> EXTENDCONV_SUPPORTED_TYPES_NCHW_MAP = {
177	- {NpuArch::DAV_5102, EXTENDCONV2D_SUPPORTED_TYPES_MDC}, {NpuArch::DAV_3510, EXTENDCONV_SUPPORTED_TYPES_NCHW}};	178	+ {NPU_ARCH_KEY_3510, EXTENDCONV_SUPPORTED_TYPES_NCHW}, {NPU_ARCH_KEY_FUSE, EXTENDCONV2D_SUPPORTED_TYPES_FUSE}};
178		179
179	-const std::map<NpuArch, std::vector<std::vector<ge::DataType>>> SOC_EXTENDCONV_SUPPORTED_TYPES_NHWC = {	180	+const std::map<std::string, std::vector<std::vector<ge::DataType>>> EXTENDCONV_SUPPORTED_TYPES_NHWC_MAP = {
180	- {NpuArch::DAV_5102, EXTENDCONV2D_SUPPORTED_TYPES_MDC}, {NpuArch::DAV_3510, EXTENDCONV_SUPPORTED_TYPES_NHWC}};	181	+ {NPU_ARCH_KEY_3510, EXTENDCONV_SUPPORTED_TYPES_NHWC}, {NPU_ARCH_KEY_FUSE, EXTENDCONV2D_SUPPORTED_TYPES_FUSE}};
181		182
182	struct ShapeBound {	183	struct ShapeBound {
183	std::map<ge::DataType, uint64_t> boundTab;	184	std::map<ge::DataType, uint64_t> boundTab;
	@@ -267,7 +268,7 @@ ge::graphStatus ShapeAttrSynthesisCheck(const ConvAscendcOriginShapeAttrInfo& or
267	ge::graphStatus ShapeAttrSynthesisCheckAux(const ConvAscendcOriginShapeAttrInfo& oriShapeAttrInfo,	268	ge::graphStatus ShapeAttrSynthesisCheckAux(const ConvAscendcOriginShapeAttrInfo& oriShapeAttrInfo,
268	ConvParamInfo paramInfo, const gert::TilingContext* context);	269	ConvParamInfo paramInfo, const gert::TilingContext* context);
269	void GetSupportedDataTypes(bool hasBias, bool quantFlag, std::vector<std::vector<ge::DataType>>& supportTypes);	270	void GetSupportedDataTypes(bool hasBias, bool quantFlag, std::vector<std::vector<ge::DataType>>& supportTypes);
270	-void GetSupportedDataTypes(const NpuArch& socVersion, bool quantFlag, ge::Format fMapFormat, bool exendConvFlag,	271	+void GetSupportedDataTypes(const std::string& archKey, bool quantFlag, ge::Format fMapFormat, bool exendConvFlag,
271	std::vector<std::vector<ge::DataType>>& supportTypes);	272	std::vector<std::vector<ge::DataType>>& supportTypes);
272	bool GetConvParamsIdx(const std::vector<ge::Format> formatVec, std::vector<std::vector<std::size_t>>& idxVec);	273	bool GetConvParamsIdx(const std::vector<ge::Format> formatVec, std::vector<std::vector<std::size_t>>& idxVec);
273	bool IsWeightNZFormat(ge::Format weightFormat);	274	bool IsWeightNZFormat(ge::Format weightFormat);
	@@ -308,8 +309,8 @@ bool ConvArrMatchWithSize(T& arr1, const T& arr2, size_t size)
308		309
309	class ConvBase : public ConvBaseDeci {	310	class ConvBase : public ConvBaseDeci {
310	public:	311	public:
311	- ConvBase(){};	312	+ ConvBase() {};
312	- explicit ConvBase(gert::TilingContext* context) : context_(context){};	313	+ explicit ConvBase(gert::TilingContext* context) : context_(context) {};
313	void ConvBaseInit(ConvAscendcShapesInfo shapeInfo, ConvAscendcDescInfo descInfo, ConvAscendcTilingFlag flagInfo,	314	void ConvBaseInit(ConvAscendcShapesInfo shapeInfo, ConvAscendcDescInfo descInfo, ConvAscendcTilingFlag flagInfo,
314	ConvParamInfo paramInfo, gert::TilingContext* context);	315	ConvParamInfo paramInfo, gert::TilingContext* context);
315	void ConvBaseInitOpInfo(const ConvTilingParseInfo* opInfo);	316	void ConvBaseInitOpInfo(const ConvTilingParseInfo* opInfo);
	@@ -339,4 +340,4 @@ private:
339	};	340	};
340	} // namespace conv_ops_tiling	341	} // namespace conv_ops_tiling
341	} // namespace optiling	342	} // namespace optiling
342	-#endif	343	+#endif

Mconv/common/op_host/op_tiling/arch35/conv_base_utils.h+26-2

	@@ -16,9 +16,24 @@
16	#define OPS_BUILT_IN_OP_TILING_RUNTIME_CONV_BASE_UTILS_H	16	#define OPS_BUILT_IN_OP_TILING_RUNTIME_CONV_BASE_UTILS_H
17	#include "../cube_tiling.h"	17	#include "../cube_tiling.h"
18	#include "conv_template_utils.h"	18	#include "conv_template_utils.h"
		19	+#include "../../conv_npu_arch_resolver.h"
		20	+#include "platform/platform_infos_def.h"
19	namespace optiling {	21	namespace optiling {
20	namespace conv_ops_tiling {	22	namespace conv_ops_tiling {
21		23
		24	+inline bool IsCubeVectorFuseSoc(fe::PlatFormInfos& platformInfo)
		25	+{
		26	+ return conv_arch::IsCubeVectorFuseSoc(platformInfo);
		27	+}
		28	+
		29	+const std::string NPU_ARCH_KEY_3510 = "3510";
		30	+const std::string NPU_ARCH_KEY_FUSE = "FUSE";
		31	+
		32	+inline const std::string& GetNpuArchKey(fe::PlatFormInfos& platformInfo)
		33	+{
		34	+ return conv_arch::GetNpuArchKey(platformInfo);
		35	+}
		36	+
22	enum class QuantMode : std::uint8_t { NO_QUANT = 0, SCALAR_QUANT, VECTOR_QUANT, UNDEFINED };	37	enum class QuantMode : std::uint8_t { NO_QUANT = 0, SCALAR_QUANT, VECTOR_QUANT, UNDEFINED };
23		38
24	enum class ReluMode : std::uint8_t { NORELU = 0, NORMALRELU = 1, SCALARRELU = 2, VECTORRELU = 3, UNDEFINED };	39	enum class ReluMode : std::uint8_t { NORELU = 0, NORMALRELU = 1, SCALARRELU = 2, VECTORRELU = 3, UNDEFINED };
	@@ -40,6 +55,7 @@ struct ConvTilingParseInfo : CubeTilingCommonParseInfo {
40	NpuArch npuArch = NpuArch::DAV_RESV;	55	NpuArch npuArch = NpuArch::DAV_RESV;
41	uint32_t aivNum = 0;	56	uint32_t aivNum = 0;
42	uint64_t fbSize = 0;	57	uint64_t fbSize = 0;
		58	+ bool isCubeVectorFuse = false;
43	ConvTilingParseInfo& operator=(const ConvTilingParseInfo* other)	59	ConvTilingParseInfo& operator=(const ConvTilingParseInfo* other)
44	{	60	{
45	if (this != other) { // 防止自赋值	61	if (this != other) { // 防止自赋值
	@@ -58,6 +74,7 @@ struct ConvTilingParseInfo : CubeTilingCommonParseInfo {
58	socVersion = other->socVersion;	74	socVersion = other->socVersion;
59	shortSocVersion = other->shortSocVersion;	75	shortSocVersion = other->shortSocVersion;
60	npuArch = other->npuArch;	76	npuArch = other->npuArch;
		77	+ isCubeVectorFuse = other->isCubeVectorFuse;
61	}	78	}
62	return *this;	79	return *this;
63	}	80	}
	@@ -113,7 +130,7 @@ const std::vector<std::vector<ge::Format>> SUPPORT_CONV2D_FORMAT_LIST = {
113	{ge::Format::FORMAT_NCHW, ge::Format::FORMAT_NCHW, ge::Format::FORMAT_NCHW},	130	{ge::Format::FORMAT_NCHW, ge::Format::FORMAT_NCHW, ge::Format::FORMAT_NCHW},
114	{ge::Format::FORMAT_NHWC, ge::Format::FORMAT_HWCN, ge::Format::FORMAT_NHWC}};	131	{ge::Format::FORMAT_NHWC, ge::Format::FORMAT_HWCN, ge::Format::FORMAT_NHWC}};
115		132
116	-const std::vector<std::vector<ge::Format>> SUPPORT_CONV2D_FORMAT_LIST_MDC = {	133	+const std::vector<std::vector<ge::Format>> SUPPORT_CONV2D_FORMAT_LIST_FUSE = {
117	{ge::Format::FORMAT_NCHW, ge::Format::FORMAT_NCHW, ge::Format::FORMAT_NCHW},	134	{ge::Format::FORMAT_NCHW, ge::Format::FORMAT_NCHW, ge::Format::FORMAT_NCHW},
118	{ge::Format::FORMAT_NCHW, ge::Format::FORMAT_HWCN, ge::Format::FORMAT_NCHW},	135	{ge::Format::FORMAT_NCHW, ge::Format::FORMAT_HWCN, ge::Format::FORMAT_NCHW},
119	{ge::Format::FORMAT_NHWC, ge::Format::FORMAT_HWCN, ge::Format::FORMAT_NHWC},	136	{ge::Format::FORMAT_NHWC, ge::Format::FORMAT_HWCN, ge::Format::FORMAT_NHWC},
	@@ -146,7 +163,7 @@ const std::vector<std::vector<ge::Format>> EXTENDCONV2D_SUPPORT_FORMAT_LIST = {
146	{ge::Format::FORMAT_NHWC, ge::Format::FORMAT_HWCN, ge::Format::FORMAT_NHWC}};	163	{ge::Format::FORMAT_NHWC, ge::Format::FORMAT_HWCN, ge::Format::FORMAT_NHWC}};
147		164
148	// ExtendConv2D fmap, weight, output supprot format list	165	// ExtendConv2D fmap, weight, output supprot format list
149	-const std::vector<std::vector<ge::Format>> EXTENDCONV2D_SUPPORT_FORMAT_LIST_MDC = {	166	+const std::vector<std::vector<ge::Format>> EXTENDCONV2D_SUPPORT_FORMAT_LIST_FUSE = {
150	{ge::Format::FORMAT_NCHW, ge::Format::FORMAT_NCHW, ge::Format::FORMAT_NCHW},	167	{ge::Format::FORMAT_NCHW, ge::Format::FORMAT_NCHW, ge::Format::FORMAT_NCHW},
151	{ge::Format::FORMAT_NCHW, ge::Format::FORMAT_HWCN, ge::Format::FORMAT_NCHW},	168	{ge::Format::FORMAT_NCHW, ge::Format::FORMAT_HWCN, ge::Format::FORMAT_NCHW},
152	{ge::Format::FORMAT_NHWC, ge::Format::FORMAT_HWCN, ge::Format::FORMAT_NHWC},	169	{ge::Format::FORMAT_NHWC, ge::Format::FORMAT_HWCN, ge::Format::FORMAT_NHWC},
	@@ -156,6 +173,13 @@ const std::vector<std::vector<ge::Format>> EXTENDCONV2D_SUPPORT_FORMAT_LIST_MDC
156	{ge::Format::FORMAT_NHWC, ge::Format::FORMAT_NCHW, ge::Format::FORMAT_NCHW},	173	{ge::Format::FORMAT_NHWC, ge::Format::FORMAT_NCHW, ge::Format::FORMAT_NCHW},
157	{ge::Format::FORMAT_NHWC, ge::Format::FORMAT_HWCN, ge::Format::FORMAT_NCHW}};	174	{ge::Format::FORMAT_NHWC, ge::Format::FORMAT_HWCN, ge::Format::FORMAT_NCHW}};
158		175
		176	+// arch-keyed format support list maps (extend by adding new arch keys)
		177	+const std::map<std::string, std::vector<std::vector<ge::Format>>> SUPPORT_CONV2D_FORMAT_LIST_MAP = {
		178	+ {NPU_ARCH_KEY_3510, SUPPORT_CONV2D_FORMAT_LIST}, {NPU_ARCH_KEY_FUSE, SUPPORT_CONV2D_FORMAT_LIST_FUSE}};
		179	+
		180	+const std::map<std::string, std::vector<std::vector<ge::Format>>> EXTENDCONV2D_SUPPORT_FORMAT_LIST_MAP = {
		181	+ {NPU_ARCH_KEY_3510, EXTENDCONV2D_SUPPORT_FORMAT_LIST}, {NPU_ARCH_KEY_FUSE, EXTENDCONV2D_SUPPORT_FORMAT_LIST_FUSE}};
		182	+
159	struct ConvParamInfo {	183	struct ConvParamInfo {
160	// Fmap, Weight, Output, FmapOri(for attr) param info	184	// Fmap, Weight, Output, FmapOri(for attr) param info
161	std::vector<ge::Format> paramsFormat = {ge::Format::FORMAT_MAX, ge::Format::FORMAT_MAX, ge::Format::FORMAT_MAX};	185	std::vector<ge::Format> paramsFormat = {ge::Format::FORMAT_MAX, ge::Format::FORMAT_MAX, ge::Format::FORMAT_MAX};

Mconv/common/op_kernel/arch35/conv_config.h+3-2

	@@ -20,6 +20,7 @@
20	#include "kernel_tiling/kernel_tiling.h"	20	#include "kernel_tiling/kernel_tiling.h"
21	#include "kernel_common.h"	21	#include "kernel_common.h"
22	#include "conv_util.h"	22	#include "conv_util.h"
		23	+#include "../../inc/macro.h"
23		24
24	namespace conv {	25	namespace conv {
25	using namespace AscendC;	26	using namespace AscendC;
	@@ -130,7 +131,7 @@ struct GetDstType<float> {
130		131
131	template <>	132	template <>
132	struct GetDstType<half> {	133	struct GetDstType<half> {
133	-#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 5102)	134	+#if defined(__DAV_35_FAMILY__)
134	using Type = int32_t;	135	using Type = int32_t;
135	#else	136	#else
136	using Type = float;	137	using Type = float;
	@@ -412,4 +413,4 @@ public:
412	};	413	};
413		414
414	} // namespace conv	415	} // namespace conv
415	-#endif	416	+#endif

Mconv/common/op_kernel/arch35/conv_instr_hw_mode_impl.h+3-3

Mconv/common/op_kernel/arch35/conv_instr_impl.h+1-1

	@@ -101,7 +101,7 @@ public:
101		101
102	// fixed-point multiplication should set cvt_mode = 2 and fix_val, which is encapsulated by basic api.	102	// fixed-point multiplication should set cvt_mode = 2 and fix_val, which is encapsulated by basic api.
103	DataCopyParams biasBtCopyParams(1, currentNL0_ * Intf::sizeOfBias / BT_BLOCK_SIZE, 0, 0);	103	DataCopyParams biasBtCopyParams(1, currentNL0_ * Intf::sizeOfBias / BT_BLOCK_SIZE, 0, 0);
104	-#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 5102)	104	+#if defined(__DAV_35_FAMILY__)
105	if constexpr (Intf::isFixedPoint) {	105	if constexpr (Intf::isFixedPoint) {
106	biasBtCopyParams.fixShiftVal = FIX_SHIFT_VAL_LEN_A16W16 - self_->ctx.convTilingData->fixedShiftValue;	106	biasBtCopyParams.fixShiftVal = FIX_SHIFT_VAL_LEN_A16W16 - self_->ctx.convTilingData->fixedShiftValue;
107	}	107	}

Mconv/common/op_kernel/arch35/conv_instr_m_mode_impl.h+4-4

Mconv/common/op_kernel/arch35/conv_iterate_impl.h+1-1

	@@ -890,7 +890,7 @@ template <class Intf, uint32_t ImplType>
890	__aicore__ void Iterate<Intf, ImplType>::IterateK(Intf* self)	890	__aicore__ void Iterate<Intf, ImplType>::IterateK(Intf* self)
891	{	891	{
892	MmadParams mmadParams;	892	MmadParams mmadParams;
893	-#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 5102)	893	+#if defined(__DAV_35_FAMILY__)
894	if constexpr (AscendC::IsSameType<typename Intf::FmapT, half>::value) {	894	if constexpr (AscendC::IsSameType<typename Intf::FmapT, half>::value) {
895	mmadParams.fixShiftVal = self->ctx.convTilingData->fixedShiftValue;	895	mmadParams.fixShiftVal = self->ctx.convTilingData->fixedShiftValue;
896	}	896	}

Mconv/common/op_kernel/arch35/conv_util.h+4-3

	@@ -17,6 +17,7 @@
17	#define CONV_UTIL_H	17	#define CONV_UTIL_H
18		18
19	#include "kernel_common.h"	19	#include "kernel_common.h"
		20	+#include "../../inc/macro.h"
20		21
21	namespace conv {	22	namespace conv {
22	using namespace AscendC;	23	using namespace AscendC;
	@@ -152,18 +153,18 @@ static constexpr IsResetLoad3dConfig CONV_LOAD3DV2_DEFAULT_CONFIG = {false, fals
152	const static uint8_t FIX_SHIFT_VAL_LEN_A16W16 = 58;	153	const static uint8_t FIX_SHIFT_VAL_LEN_A16W16 = 58;
153		154
154	constexpr FixpipeConfig CFG_COLUMN_MAJOR_FIXED_POINT = {CO2Layout::COLUMN_MAJOR, false,	155	constexpr FixpipeConfig CFG_COLUMN_MAJOR_FIXED_POINT = {CO2Layout::COLUMN_MAJOR, false,
155	-#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 5102)	156	+#if defined(__DAV_35_FAMILY__)
156	true	157	true
157	#endif	158	#endif
158	};	159	};
159	constexpr FixpipeConfig CFG_ROW_MAJOR_FIXED_POINT = {CO2Layout::ROW_MAJOR, false,	160	constexpr FixpipeConfig CFG_ROW_MAJOR_FIXED_POINT = {CO2Layout::ROW_MAJOR, false,
160	-#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 5102)	161	+#if defined(__DAV_35_FAMILY__)
161	true	162	true
162	#endif	163	#endif
163	};	164	};
164	constexpr FixpipeConfig CFG_ROW_MAJOR_UB = {CO2Layout::ROW_MAJOR, true};	165	constexpr FixpipeConfig CFG_ROW_MAJOR_UB = {CO2Layout::ROW_MAJOR, true};
165		166
166	-#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 5102)	167	+#if defined(__DAV_35_FAMILY__)
167	#define ASCEND_IS_AIC_CONV constexpr(true)	168	#define ASCEND_IS_AIC_CONV constexpr(true)
168	#define ASCEND_IS_AIV_CONV constexpr(true)	169	#define ASCEND_IS_AIV_CONV constexpr(true)
169	#else	170	#else

Mconv/common/tests/ut/op_tiling_test/test_conv_base_utils_tiling.cpp+139-4

	@@ -14,7 +14,10 @@
14	*/	14	*/
15		15
16	#include <gtest/gtest.h>	16	#include <gtest/gtest.h>
		17	+#include <map>
		18	+#include <string>
17	#include "conv/common/op_host/op_tiling/arch35/conv_base_utils.h"	19	#include "conv/common/op_host/op_tiling/arch35/conv_base_utils.h"
		20	+#include "conv/common/op_host/op_tiling/arch35/conv_base.h"
18	#include "conv/common/op_host/op_tiling/cube_tiling.h"	21	#include "conv/common/op_host/op_tiling/cube_tiling.h"
19		22
20	using namespace optiling;	23	using namespace optiling;
	@@ -53,6 +56,7 @@ TEST(ConvBaseUtilsTest, ConvTilingParseInfoDefaults)
53	EXPECT_EQ(info.npuArch, NpuArch::DAV_RESV);	56	EXPECT_EQ(info.npuArch, NpuArch::DAV_RESV);
54	EXPECT_EQ(info.aivNum, 0u);	57	EXPECT_EQ(info.aivNum, 0u);
55	EXPECT_EQ(info.fbSize, 0u);	58	EXPECT_EQ(info.fbSize, 0u);
		59	+ EXPECT_FALSE(info.isCubeVectorFuse);
56	}	60	}
57		61
58	TEST(ConvBaseUtilsTest, ConvTilingParseInfoOperatorAssign)	62	TEST(ConvBaseUtilsTest, ConvTilingParseInfoOperatorAssign)
	@@ -62,12 +66,14 @@ TEST(ConvBaseUtilsTest, ConvTilingParseInfoOperatorAssign)
62	src.l1Size = 524288;	66	src.l1Size = 524288;
63	src.socVersion = "ascend910b";	67	src.socVersion = "ascend910b";
64	src.npuArch = NpuArch::DAV_3510;	68	src.npuArch = NpuArch::DAV_3510;
		69	+ src.isCubeVectorFuse = true;
65	ConvTilingParseInfo dst;	70	ConvTilingParseInfo dst;
66	dst.operator=(&src);	71	dst.operator=(&src);
67	EXPECT_EQ(dst.aicoreNum, 32u);	72	EXPECT_EQ(dst.aicoreNum, 32u);
68	EXPECT_EQ(dst.l1Size, 524288u);	73	EXPECT_EQ(dst.l1Size, 524288u);
69	EXPECT_EQ(dst.socVersion, "ascend910b");	74	EXPECT_EQ(dst.socVersion, "ascend910b");
70	EXPECT_EQ(dst.npuArch, NpuArch::DAV_3510);	75	EXPECT_EQ(dst.npuArch, NpuArch::DAV_3510);
		76	+ EXPECT_TRUE(dst.isCubeVectorFuse);
71	}	77	}
72		78
73	TEST(ConvBaseUtilsTest, ConvAscendcOriginShapeAttrInfoDefaults)	79	TEST(ConvBaseUtilsTest, ConvAscendcOriginShapeAttrInfoDefaults)
	@@ -130,10 +136,10 @@ TEST(ConvBaseUtilsTest, SupportConv2dFormatListNotEmpty)
130	EXPECT_EQ(SUPPORT_CONV2D_FORMAT_LIST.size(), 2u);	136	EXPECT_EQ(SUPPORT_CONV2D_FORMAT_LIST.size(), 2u);
131	}	137	}
132		138
133	-TEST(ConvBaseUtilsTest, SupportConv2dFormatListMdcNotEmpty)	139	+TEST(ConvBaseUtilsTest, SupportConv2dFormatListFuseNotEmpty)
134	{	140	{
135	- EXPECT_FALSE(SUPPORT_CONV2D_FORMAT_LIST_MDC.empty());	141	+ EXPECT_FALSE(SUPPORT_CONV2D_FORMAT_LIST_FUSE.empty());
136	- EXPECT_EQ(SUPPORT_CONV2D_FORMAT_LIST_MDC.size(), 8u);	142	+ EXPECT_EQ(SUPPORT_CONV2D_FORMAT_LIST_FUSE.size(), 8u);
137	}	143	}
138		144
139	TEST(ConvBaseUtilsTest, SupportConv3dFormatListNotEmpty)	145	TEST(ConvBaseUtilsTest, SupportConv3dFormatListNotEmpty)
	@@ -151,5 +157,134 @@ TEST(ConvBaseUtilsTest, SupportQuantConvFormatListNotEmpty)
151	TEST(ConvBaseUtilsTest, ExtendConv2dFormatListNotEmpty)	157	TEST(ConvBaseUtilsTest, ExtendConv2dFormatListNotEmpty)
152	{	158	{
153	EXPECT_FALSE(EXTENDCONV2D_SUPPORT_FORMAT_LIST.empty());	159	EXPECT_FALSE(EXTENDCONV2D_SUPPORT_FORMAT_LIST.empty());
154	- EXPECT_FALSE(EXTENDCONV2D_SUPPORT_FORMAT_LIST_MDC.empty());	160	+ EXPECT_FALSE(EXTENDCONV2D_SUPPORT_FORMAT_LIST_FUSE.empty());
		161	+}
		162	+
		163	+// ============================================================================
		164	+// IsCubeVectorFuseSoc: check whether cube_vector_combine == "fuse"
		165	+// ============================================================================
		166	+TEST(ConvBaseUtilsTest, IsCubeVectorFuseSocFuseReturnsTrue)
		167	+{
		168	+ fe::PlatFormInfos platformInfo;
		169	+ platformInfo.Init();
		170	+ std::map<std::string, std::string> socInfos = {{"cube_vector_combine", "fuse"}};
		171	+ platformInfo.SetPlatformRes("SoCInfo", socInfos);
		172	+ EXPECT_TRUE(IsCubeVectorFuseSoc(platformInfo));
		173	+}
		174	+
		175	+TEST(ConvBaseUtilsTest, IsCubeVectorFuseSocSplitReturnsFalse)
		176	+{
		177	+ fe::PlatFormInfos platformInfo;
		178	+ platformInfo.Init();
		179	+ std::map<std::string, std::string> socInfos = {{"cube_vector_combine", "split"}};
		180	+ platformInfo.SetPlatformRes("SoCInfo", socInfos);
		181	+ EXPECT_FALSE(IsCubeVectorFuseSoc(platformInfo));
		182	+}
		183	+
		184	+TEST(ConvBaseUtilsTest, IsCubeVectorFuseSocEmptyValueReturnsFalse)
		185	+{
		186	+ fe::PlatFormInfos platformInfo;
		187	+ platformInfo.Init();
		188	+ std::map<std::string, std::string> socInfos = {{"cube_vector_combine", ""}};
		189	+ platformInfo.SetPlatformRes("SoCInfo", socInfos);
		190	+ EXPECT_FALSE(IsCubeVectorFuseSoc(platformInfo));
		191	+}
		192	+
		193	+TEST(ConvBaseUtilsTest, IsCubeVectorFuseSocMissingKeyReturnsFalse)
		194	+{
		195	+ fe::PlatFormInfos platformInfo;
		196	+ platformInfo.Init();
		197	+ std::map<std::string, std::string> socInfos = {{"ai_core_cnt", "32"}};
		198	+ platformInfo.SetPlatformRes("SoCInfo", socInfos);
		199	+ EXPECT_FALSE(IsCubeVectorFuseSoc(platformInfo));
		200	+}
		201	+
		202	+// ============================================================================
		203	+// ConvTilingParseInfo::isCubeVectorFuse field default and operator=
		204	+// ============================================================================
		205	+TEST(ConvBaseUtilsTest, ConvTilingParseInfoIsCubeVectorFuseDefault)
		206	+{
		207	+ ConvTilingParseInfo info;
		208	+ EXPECT_FALSE(info.isCubeVectorFuse);
		209	+}
		210	+
		211	+TEST(ConvBaseUtilsTest, ConvTilingParseInfoIsCubeVectorFuseOperatorAssign)
		212	+{
		213	+ ConvTilingParseInfo src;
		214	+ src.isCubeVectorFuse = true;
		215	+ src.npuArch = NpuArch::DAV_3510;
		216	+ ConvTilingParseInfo dst;
		217	+ dst.operator=(&src);
		218	+ EXPECT_TRUE(dst.isCubeVectorFuse);
		219	+ EXPECT_EQ(dst.npuArch, NpuArch::DAV_3510);
		220	+
		221	+ src.isCubeVectorFuse = false;
		222	+ dst.operator=(&src);
		223	+ EXPECT_FALSE(dst.isCubeVectorFuse);
		224	+}
		225	+
		226	+// ============================================================================
		227	+// conv_tiling::PlatformInfo::isCubeVectorFuse field default
		228	+// ============================================================================
		229	+TEST(ConvBaseUtilsTest, PlatformInfoIsCubeVectorFuseDefault)
		230	+{
		231	+ conv_tiling::PlatformInfo info;
		232	+ EXPECT_FALSE(info.isCubeVectorFuse);
		233	+ EXPECT_EQ(info.npuArch, NpuArch::DAV_RESV);
		234	+}
		235	+
		236	+// ============================================================================
		237	+// GetSupportedDataTypes: arch key selects FUSE vs DAV type lists
		238	+// ============================================================================
		239	+TEST(ConvBaseUtilsTest, GetSupportedDataTypesFuseNonQuantNonExtend)
		240	+{
		241	+ std::vector<std::vector<ge::DataType>> result;
		242	+ GetSupportedDataTypes("FUSE", false, ge::FORMAT_NCHW, false, result);
		243	+ EXPECT_EQ(result, CONV_SUPPORTED_TYPES_FUSE);
		244	+}
		245	+
		246	+TEST(ConvBaseUtilsTest, GetSupportedDataTypesDavNonQuantNonExtend)
		247	+{
		248	+ std::vector<std::vector<ge::DataType>> result;
		249	+ GetSupportedDataTypes("3510", false, ge::FORMAT_NCHW, false, result);
		250	+ EXPECT_EQ(result, CONV_SUPPORTED_TYPES_DAV);
		251	+}
		252	+
		253	+TEST(ConvBaseUtilsTest, GetSupportedDataTypesFuseExtendNchw)
		254	+{
		255	+ std::vector<std::vector<ge::DataType>> result;
		256	+ GetSupportedDataTypes("FUSE", false, ge::FORMAT_NCHW, true, result);
		257	+ EXPECT_EQ(result, EXTENDCONV2D_SUPPORTED_TYPES_FUSE);
		258	+}
		259	+
		260	+TEST(ConvBaseUtilsTest, GetSupportedDataTypesDavExtendNchw)
		261	+{
		262	+ std::vector<std::vector<ge::DataType>> result;
		263	+ GetSupportedDataTypes("3510", false, ge::FORMAT_NCHW, true, result);
		264	+ EXPECT_EQ(result, EXTENDCONV_SUPPORTED_TYPES_NCHW);
		265	+}
		266	+
		267	+TEST(ConvBaseUtilsTest, GetSupportedDataTypesFuseExtendNhwc)
		268	+{
		269	+ std::vector<std::vector<ge::DataType>> result;
		270	+ GetSupportedDataTypes("FUSE", false, ge::FORMAT_NHWC, true, result);
		271	+ EXPECT_EQ(result, EXTENDCONV2D_SUPPORTED_TYPES_FUSE);
		272	+}
		273	+
		274	+TEST(ConvBaseUtilsTest, GetSupportedDataTypesDavExtendNhwc)
		275	+{
		276	+ std::vector<std::vector<ge::DataType>> result;
		277	+ GetSupportedDataTypes("3510", false, ge::FORMAT_NHWC, true, result);
		278	+ EXPECT_EQ(result, EXTENDCONV_SUPPORTED_TYPES_NHWC);
		279	+}
		280	+
		281	+TEST(ConvBaseUtilsTest, GetSupportedDataTypesQuantReturnsQuantTypesRegardlessArch)
		282	+{
		283	+ std::vector<std::vector<ge::DataType>> resultFuse;
		284	+ GetSupportedDataTypes("FUSE", true, ge::FORMAT_NCHW, false, resultFuse);
		285	+ EXPECT_EQ(resultFuse, QUANTCONV_SUPPORTED_TYPES);
		286	+
		287	+ std::vector<std::vector<ge::DataType>> resultDav;
		288	+ GetSupportedDataTypes("3510", true, ge::FORMAT_NCHW, false, resultDav);
		289	+ EXPECT_EQ(resultDav, QUANTCONV_SUPPORTED_TYPES);
155	}	290	}

Mconv/conv2d_v2/op_graph/fusion_pass/conv2d_to_conv2dv2_fusion_pass.cpp+3-4

	@@ -32,7 +32,7 @@ bool Conv2dToConv2dV2FusionPass::MeetRequirements(const GNode& convNode)
32	{	32	{
33	InitMember();	33	InitMember();
34		34
35	- FUSION_PASS_CHECK(!ConvFusionUtilsPass::CheckSocList(SUPPORT_SOC_LIST, npuArch),	35	+ FUSION_PASS_CHECK(!ConvFusionUtilsPass::CheckSocList(SUPPORT_SOC_LIST, npuArch, true),
36	OP_LOGD(FUSION_NAME, "Current soc not supported, no fusion."), return false);	36	OP_LOGD(FUSION_NAME, "Current soc not supported, no fusion."), return false);
37		37
38	FUSION_PASS_CHECK_NOLOG(!ConvFusionUtilsPass::GetConvDescInfo(convNode, convDescInfo), return false);	38	FUSION_PASS_CHECK_NOLOG(!ConvFusionUtilsPass::GetConvDescInfo(convNode, convDescInfo), return false);
	@@ -42,8 +42,7 @@ bool Conv2dToConv2dV2FusionPass::MeetRequirements(const GNode& convNode)
42	if (convDescInfo.hasBias) {	42	if (convDescInfo.hasBias) {
43	convDtypes.emplace_back(convDescInfo.biasDtype);	43	convDtypes.emplace_back(convDescInfo.biasDtype);
44	}	44	}
45	- const auto& convSupportList = (npuArch == NpuArch::DAV_3510) ? CONV_SUPPORT_DTYPES_DAV_3510 :	45	+ const auto& convSupportList = CONV_SUPPORT_DTYPES_MAP.at(ConvFusionUtilsPass::GetArchKey());
46	- CONV_SUPPORT_DTYPES_DAV_5102;
47	FUSION_PASS_CHECK(!ConvFusionUtilsPass::CheckSupportList<DataType>(convSupportList, convDtypes),	46	FUSION_PASS_CHECK(!ConvFusionUtilsPass::CheckSupportList<DataType>(convSupportList, convDtypes),
48	OP_LOGD(convDescInfo.nodeNameStr, "Conv2D dtype not supported, no fusion."), return false);	47	OP_LOGD(convDescInfo.nodeNameStr, "Conv2D dtype not supported, no fusion."), return false);
49		48
	@@ -74,4 +73,4 @@ GraphUniqPtr Conv2dToConv2dV2FusionPass::Replacement(const GNode& convNode)
74	return graphBuilder.BuildAndReset({conv2dV2});	73	return graphBuilder.BuildAndReset({conv2dV2});
75	}	74	}
76		75
77	-} // namespace Ops	76	+} // namespace Ops

Mconv/conv2d_v2/op_graph/fusion_pass/conv2d_to_conv2dv2_fusion_pass.h+8-3

	@@ -25,7 +25,7 @@ namespace Conv {
25	namespace Conv2dToConv2dV2Fusion {	25	namespace Conv2dToConv2dV2Fusion {
26	const std::string FUSION_NAME = "Conv2dToConv2dV2FusionPass";	26	const std::string FUSION_NAME = "Conv2dToConv2dV2FusionPass";
27		27
28	-const std::map<std::string, NpuArch> SUPPORT_SOC_LIST = {{"Ascend950", NpuArch::DAV_3510}, {"MC62", NpuArch::DAV_5102}};	28	+const std::map<std::string, NpuArch> SUPPORT_SOC_LIST = {{"Ascend950", NpuArch::DAV_3510}};
29		29
30	// Fmap Filter Output Bias	30	// Fmap Filter Output Bias
31	const std::vector<std::vector<ge::DataType>> CONV_SUPPORT_DTYPES_DAV_3510 = {	31	const std::vector<std::vector<ge::DataType>> CONV_SUPPORT_DTYPES_DAV_3510 = {
	@@ -35,9 +35,14 @@ const std::vector<std::vector<ge::DataType>> CONV_SUPPORT_DTYPES_DAV_3510 = {
35	{ge::DataType::DT_HIFLOAT8, ge::DataType::DT_HIFLOAT8, ge::DataType::DT_HIFLOAT8, ge::DataType::DT_FLOAT}};	35	{ge::DataType::DT_HIFLOAT8, ge::DataType::DT_HIFLOAT8, ge::DataType::DT_HIFLOAT8, ge::DataType::DT_FLOAT}};
36		36
37	// Fmap Filter Output Bias	37	// Fmap Filter Output Bias
38	-const std::vector<std::vector<ge::DataType>> CONV_SUPPORT_DTYPES_DAV_5102 = {	38	+const std::vector<std::vector<ge::DataType>> CONV_SUPPORT_DTYPES_FUSE = {
39	{ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16},	39	{ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16},
40	{ge::DataType::DT_INT8, ge::DataType::DT_INT8, ge::DataType::DT_INT32, ge::DataType::DT_INT32}};	40	{ge::DataType::DT_INT8, ge::DataType::DT_INT8, ge::DataType::DT_INT32, ge::DataType::DT_INT32}};
		41	+
		42	+// arch-keyed dtype support list map (extend by adding new arch keys)
		43	+const std::map<std::string, std::vector<std::vector<ge::DataType>>> CONV_SUPPORT_DTYPES_MAP = {
		44	+ {ConvFusionUtils::NPU_ARCH_KEY_3510, CONV_SUPPORT_DTYPES_DAV_3510},
		45	+ {ConvFusionUtils::NPU_ARCH_KEY_FUSE, CONV_SUPPORT_DTYPES_FUSE}};
41	} // namespace Conv2dToConv2dV2Fusion	46	} // namespace Conv2dToConv2dV2Fusion
42		47
43	class __attribute__((visibility("default"))) Conv2dToConv2dV2FusionPass : public ge::fusion::DecomposePass {	48	class __attribute__((visibility("default"))) Conv2dToConv2dV2FusionPass : public ge::fusion::DecomposePass {
	@@ -58,4 +63,4 @@ private:
58	} // namespace Conv	63	} // namespace Conv
59	} // namespace NN	64	} // namespace NN
60	} // namespace Ops	65	} // namespace Ops
61	-#endif // CONV2D_TO_CONV2DV2_FUSION_PASS_H	66	+#endif // CONV2D_TO_CONV2DV2_FUSION_PASS_H

Mconv/conv2d_v2/op_host/op_tiling/arch35/conv2d_v2_api_tiling.cpp+3-3

	@@ -1042,7 +1042,7 @@ bool Conv2dTiling::CheckL1SizeLimitsKernelFullLoad(bool isC04)
1042	std::vector<int64_t> xShape = {shapeInfo.singleBatch, shapeInfo.orgCi, shapeInfo.orgHi, shapeInfo.orgWi};	1042	std::vector<int64_t> xShape = {shapeInfo.singleBatch, shapeInfo.orgCi, shapeInfo.orgHi, shapeInfo.orgWi};
1043	std::vector<int64_t> filterShape = {shapeInfo.orgCo, shapeInfo.orgCi, shapeInfo.orgkH, shapeInfo.orgkW};	1043	std::vector<int64_t> filterShape = {shapeInfo.orgCo, shapeInfo.orgCi, shapeInfo.orgkH, shapeInfo.orgkW};
1044	if (minL1LoadSize > platformInfo.l1Size) {	1044	if (minL1LoadSize > platformInfo.l1Size) {
1045	- if (platformInfo.npuArch == NpuArch::DAV_5102) {	1045	+ if (platformInfo.isCubeVectorFuse) {
1046	OP_LOGE_FOR_INVALID_SHAPES_WITH_REASON(	1046	OP_LOGE_FOR_INVALID_SHAPES_WITH_REASON(
1047	nodeType.c_str(), "x, filter",	1047	nodeType.c_str(), "x, filter",
1048	VectorsToString(std::vector<std::vector<int64_t>>{xShape, filterShape}, IntToString<int64_t>).c_str(),	1048	VectorsToString(std::vector<std::vector<int64_t>>{xShape, filterShape}, IntToString<int64_t>).c_str(),
	@@ -1062,7 +1062,7 @@ bool Conv2dTiling::CheckInstructionLimits()
1062	{	1062	{
1063	if (!CheckLoad3DLimits() \|\| !CheckL1SizeLimitsKernelFullLoad(isC04Flag)) {	1063	if (!CheckLoad3DLimits() \|\| !CheckL1SizeLimitsKernelFullLoad(isC04Flag)) {
1064	this->isDmaFlag = true;	1064	this->isDmaFlag = true;
1065	- if (platformInfo.npuArch == NpuArch::DAV_5102) {	1065	+ if (platformInfo.isCubeVectorFuse) {
1066	OP_LOGE(nodeType, "Conv2d does not support DMA mode. "	1066	OP_LOGE(nodeType, "Conv2d does not support DMA mode. "
1067	"Please adjust the parameters to satisfy Load3D constraints (see above error details).");	1067	"Please adjust the parameters to satisfy Load3D constraints (see above error details).");
1068	return false;	1068	return false;
	@@ -1268,4 +1268,4 @@ bool Conv2dTiling::GetCoreBindingDecisionFactor(Conv2DBasicBlockInfo& conv2DBasi
1268	algoBBPtr->CalcCoreUtilization();	1268	algoBBPtr->CalcCoreUtilization();
1269	return true;	1269	return true;
1270	}	1270	}
1271	-} // namespace conv_tiling	1271	+} // namespace conv_tiling

Mconv/conv2d_v2/op_host/op_tiling/arch35/conv2d_v2_base_tiling.cpp+10-6

	@@ -31,8 +31,11 @@ ge::graphStatus Conv2dBaseTiling::GetPlatformInfoInner()
31	Conv2dTilingCache& tilingCache = Conv2dTilingCache::GetInstance();	31	Conv2dTilingCache& tilingCache = Conv2dTilingCache::GetInstance();
32	opInfo_ = tilingCache.GetPlatFormInfo();	32	opInfo_ = tilingCache.GetPlatFormInfo();
33	npuArch = tilingCache.GetSocVersion();	33	npuArch = tilingCache.GetSocVersion();
		34	+ fe::PlatFormInfos* platformInfoPtr = context_->GetPlatformInfo();
		35	+ if (platformInfoPtr != nullptr) {
		36	+ opInfo_->isCubeVectorFuse = IsCubeVectorFuseSoc(*platformInfoPtr);
		37	+ }
34	if (npuArch == NpuArch::DAV_RESV) {	38	if (npuArch == NpuArch::DAV_RESV) {
35	- fe::PlatFormInfos* platformInfoPtr = context_->GetPlatformInfo();
36	OPS_CHECK_NULL_WITH_CONTEXT(context_, platformInfoPtr);	39	OPS_CHECK_NULL_WITH_CONTEXT(context_, platformInfoPtr);
37	auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);	40	auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
38	opInfo_->aicoreNum = ascendcPlatform.GetCoreNumAic();	41	opInfo_->aicoreNum = ascendcPlatform.GetCoreNumAic();
	@@ -55,6 +58,7 @@ ge::graphStatus Conv2dBaseTiling::GetPlatformInfoInner()
55	void Conv2dBaseTiling::SetApiInputPlatformInfo()	58	void Conv2dBaseTiling::SetApiInputPlatformInfo()
56	{	59	{
57	apiInputPlatformInfo.npuArch = opInfo_->npuArch;	60	apiInputPlatformInfo.npuArch = opInfo_->npuArch;
		61	+ apiInputPlatformInfo.isCubeVectorFuse = opInfo_->isCubeVectorFuse;
58	apiInputPlatformInfo.l1Size = opInfo_->l1Size;	62	apiInputPlatformInfo.l1Size = opInfo_->l1Size;
59	apiInputPlatformInfo.l0CSize = opInfo_->l0cSize;	63	apiInputPlatformInfo.l0CSize = opInfo_->l0cSize;
60	apiInputPlatformInfo.l0ASize = opInfo_->l0aSize;	64	apiInputPlatformInfo.l0ASize = opInfo_->l0aSize;
	@@ -64,9 +68,9 @@ void Conv2dBaseTiling::SetApiInputPlatformInfo()
64	apiInputPlatformInfo.fbSize = opInfo_->fbSize;	68	apiInputPlatformInfo.fbSize = opInfo_->fbSize;
65	OP_LOGD(context_->GetNodeName(),	69	OP_LOGD(context_->GetNodeName(),
66	"%s AscendC: Tiling get platformInfo: l1Size: %ld, l0CSize: %ld, l0ASize: %ld, l0BSize: %ld, ubSize: %ld, "	70	"%s AscendC: Tiling get platformInfo: l1Size: %ld, l0CSize: %ld, l0ASize: %ld, l0BSize: %ld, ubSize: %ld, "
67	- "btSize: %ld, fbSize: %ld.",	71	+ "btSize: %ld, fbSize: %ld, isCubeVectorFuse: %d.",
68	paramInfo_.nodeType.c_str(), opInfo_->l1Size, opInfo_->l0cSize, opInfo_->l0aSize, opInfo_->l0bSize,	72	paramInfo_.nodeType.c_str(), opInfo_->l1Size, opInfo_->l0cSize, opInfo_->l0aSize, opInfo_->l0bSize,
69	- opInfo_->ubSize, opInfo_->btSize, opInfo_->fbSize);	73	+ opInfo_->ubSize, opInfo_->btSize, opInfo_->fbSize, opInfo_->isCubeVectorFuse);
70	}	74	}
71		75
72	ge::graphStatus Conv2dBaseTiling::InitConv2dApiTiling()	76	ge::graphStatus Conv2dBaseTiling::InitConv2dApiTiling()
	@@ -206,7 +210,7 @@ void Conv2dBaseTiling::GetDescInfo()
206	descInfo_.biasFormat = static_cast<ge::Format>(	210	descInfo_.biasFormat = static_cast<ge::Format>(
207	GetPrimaryFormat(context_->GetOptionalInputDesc(biasIndex)->GetStorageFormat()));	211	GetPrimaryFormat(context_->GetOptionalInputDesc(biasIndex)->GetStorageFormat()));
208	}	212	}
209	- if (IsMdcSoc(opInfo_->npuArch)) {	213	+ if (opInfo_->isCubeVectorFuse) {
210	paramInfo_.paramsFormat = {descInfo_.fMapFormat, GetWeightFormat(), descInfo_.outFormat};	214	paramInfo_.paramsFormat = {descInfo_.fMapFormat, GetWeightFormat(), descInfo_.outFormat};
211	} else {	215	} else {
212	paramInfo_.paramsFormat = {descInfo_.fMapFormat, descInfo_.weightFormat, descInfo_.outFormat};	216	paramInfo_.paramsFormat = {descInfo_.fMapFormat, descInfo_.weightFormat, descInfo_.outFormat};
	@@ -220,7 +224,7 @@ void Conv2dBaseTiling::GetDescInfo()
220		224
221	bool Conv2dBaseTiling::IsEnableC04()	225	bool Conv2dBaseTiling::IsEnableC04()
222	{	226	{
223	- if (IsMdcSoc(opInfo_->npuArch)) {	227	+ if (opInfo_->isCubeVectorFuse) {
224	return descInfo_.weightFormat == ge::Format::FORMAT_FRACTAL_Z_C04;	228	return descInfo_.weightFormat == ge::Format::FORMAT_FRACTAL_Z_C04;
225	}	229	}
226		230
	@@ -535,4 +539,4 @@ ge::graphStatus Conv2dBaseTiling::GetWorkspaceSize()
535	return ge::GRAPH_SUCCESS;	539	return ge::GRAPH_SUCCESS;
536	}	540	}
537	} // namespace conv_ops_tiling	541	} // namespace conv_ops_tiling
538	-} // namespace optiling	542	+} // namespace optiling

Mconv/conv2d_v2/op_host/op_tiling/arch35/conv2d_v2_base_tiling_check_attrs.cpp+4-6

Mconv/conv2d_v2/op_host/op_tiling/arch35/conv2d_v2_base_tiling_check_input.cpp+7-2

	@@ -497,14 +497,19 @@ ge::graphStatus Conv2dBaseTiling::CheckParamsDtypeWithoutBias(std::vector<std::v
497	ge::graphStatus Conv2dBaseTiling::CheckParamsDtype()	497	ge::graphStatus Conv2dBaseTiling::CheckParamsDtype()
498	{	498	{
499	// check int8 input not support c04	499	// check int8 input not support c04
500	- if (IsMdcSoc(opInfo_->npuArch) && descInfo_.weightFormat == ge::Format::FORMAT_FRACTAL_Z_C04 &&	500	+ if (opInfo_->isCubeVectorFuse && descInfo_.weightFormat == ge::Format::FORMAT_FRACTAL_Z_C04 &&
501	dtypeMap.at(descInfo_.fMapDtype) == ConvDtype::INT8) {	501	dtypeMap.at(descInfo_.fMapDtype) == ConvDtype::INT8) {
502	OP_LOGE(context_->GetNodeName(), "%s AscendC: int8 input not support C04.", context_->GetNodeType());	502	OP_LOGE(context_->GetNodeName(), "%s AscendC: int8 input not support C04.", context_->GetNodeType());
503	return ge::GRAPH_FAILED;	503	return ge::GRAPH_FAILED;
504	}	504	}
505		505
506	std::vector<std::vector<ge::DataType>> supportedTypesList;	506	std::vector<std::vector<ge::DataType>> supportedTypesList;
507	- GetSupportedDataTypes(apiInputPlatformInfo.npuArch, flagInfo_.quantFlag, descInfo_.fMapFormat,	507	+ fe::PlatFormInfos* platformInfoPtr = context_->GetPlatformInfo();
		508	+ if (platformInfoPtr == nullptr) {
		509	+ OP_LOGE(context_->GetNodeName(), "%s AscendC: GetPlatformInfo return nullptr.", paramInfo_.nodeType.c_str());
		510	+ return ge::GRAPH_FAILED;
		511	+ }
		512	+ GetSupportedDataTypes(GetNpuArchKey(*platformInfoPtr), flagInfo_.quantFlag, descInfo_.fMapFormat,
508	flagInfo_.extendConvFlag, supportedTypesList);	513	flagInfo_.extendConvFlag, supportedTypesList);
509	OP_TILING_CHECK(	514	OP_TILING_CHECK(
510	supportedTypesList.size() == 0,	515	supportedTypesList.size() == 0,

Mconv/conv2d_v2/op_host/op_tiling/arch35/conv2d_v2_base_tiling_check_limits.cpp+1-1

	@@ -19,7 +19,7 @@ namespace conv_ops_tiling {
19		19
20	ge::graphStatus Conv2dBaseTiling::CheckC04Mdc()	20	ge::graphStatus Conv2dBaseTiling::CheckC04Mdc()
21	{	21	{
22	- if (!IsMdcSoc(opInfo_->npuArch)) {	22	+ if (!opInfo_->isCubeVectorFuse) {
23	return ge::GRAPH_SUCCESS;	23	return ge::GRAPH_SUCCESS;
24	}	24	}
25		25

Mconv/conv2d_v2/op_host/op_tiling/arch35/conv2d_v2_tiling_utils.h+0-3

	@@ -120,9 +120,6 @@ struct Conv2dOriginFormatAixsPosInfo {
120	uint32_t wIndex = 0;	120	uint32_t wIndex = 0;
121	};	121	};
122		122
123	-// the function used by new MDC chip supporting fix-point operation
124	-inline bool IsMdcSoc(const NpuArch shortSoc) { return shortSoc == NpuArch::DAV_5102; }
125	-
126	// the function used by judgeing operation type	123	// the function used by judgeing operation type
127	inline bool isQuantConv2D(const string& nodeType) { return nodeType == "QuantConv2D"; }	124	inline bool isQuantConv2D(const string& nodeType) { return nodeType == "QuantConv2D"; }
128		125

Mconv/conv2d_v2/op_host/op_tiling/conv2d_v2_mdc_tiling.cpp+1-1

	@@ -24,4 +24,4 @@ using namespace optiling::conv_ops_tiling;
24	namespace optiling {	24	namespace optiling {
25	// using op_tiling register capability in "tiling_templates_registry" for AscendC conv2d operator	25	// using op_tiling register capability in "tiling_templates_registry" for AscendC conv2d operator
26	CONV_REGISTER_TILING_TEMPLATE(Conv2DV2, Conv2dBaseTiling, static_cast<int32_t>(NpuArch::DAV_5102), 1);	26	CONV_REGISTER_TILING_TEMPLATE(Conv2DV2, Conv2dBaseTiling, static_cast<int32_t>(NpuArch::DAV_5102), 1);
27	-} // namespace optiling	27	+} // namespace optiling

Mconv/conv2d_v2/op_kernel/arch35/conv2d_small_kernel.h+6-6

	@@ -54,7 +54,7 @@ public:
54	using BiasT = biasType;	54	using BiasT = biasType;
55	using Output0T = out0Type;	55	using Output0T = out0Type;
56	using Output1T = out1Type;	56	using Output1T = out1Type;
57	-#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 5102)	57	+#if defined(__DAV_35_FAMILY__)
58	using L0cT = int32_t;	58	using L0cT = int32_t;
59	#else	59	#else
60	using L0cT = float;	60	using L0cT = float;
	@@ -349,7 +349,7 @@ __aicore__ inline void Conv2dSmallKernel<FmapType, weightType, biasType, out0Typ
349	uint32_t kL0MaxIter)	349	uint32_t kL0MaxIter)
350	{	350	{
351	MmadParams mp;	351	MmadParams mp;
352	-#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 5102)	352	+#if defined(__DAV_35_FAMILY__)
353	if constexpr (AscendC::IsSameType<FmapType, half>::value) {	353	if constexpr (AscendC::IsSameType<FmapType, half>::value) {
354	mp.fixShiftVal = tiling_->fixedShiftValue;	354	mp.fixShiftVal = tiling_->fixedShiftValue;
355	}	355	}
	@@ -932,7 +932,7 @@ __aicore__ inline void Conv2dSmallKernel<FmapType, weightType, biasType, out0Typ
932	LocalTensor<BiasT> biasL1src(TPosition::A1, biasL1OffBytes_, tiling_->singleCoreCo);	932	LocalTensor<BiasT> biasL1src(TPosition::A1, biasL1OffBytes_, tiling_->singleCoreCo);
933	uint32_t blkCnt = AlignB(actualCo_ * sizeof(BiasT), BT_ALIGN) / 32;	933	uint32_t blkCnt = AlignB(actualCo_ * sizeof(BiasT), BT_ALIGN) / 32;
934	DataCopyParams cp(1, static_cast<uint16_t>(blkCnt), 0, 0);	934	DataCopyParams cp(1, static_cast<uint16_t>(blkCnt), 0, 0);
935	-#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 5102)	935	+#if defined(__DAV_35_FAMILY__)
936	if constexpr (AscendC::IsSameType<weightType, half>::value) {	936	if constexpr (AscendC::IsSameType<weightType, half>::value) {
937	cp.fixShiftVal = FIX_SHIFT_VAL_LEN_A16W16 - tiling_->fixedShiftValue;	937	cp.fixShiftVal = FIX_SHIFT_VAL_LEN_A16W16 - tiling_->fixedShiftValue;
938	}	938	}
	@@ -1017,7 +1017,7 @@ template <typename OutputT, uint64_t FixpipeIdx>
1017	__aicore__ inline QuantMode_t Conv2dSmallKernel<FmapType, weightType, biasType, out0Type, out1Type, isNHWCin, isNHWCout,	1017	__aicore__ inline QuantMode_t Conv2dSmallKernel<FmapType, weightType, biasType, out0Type, out1Type, isNHWCin, isNHWCout,
1018	WeightFmt, IsHwMode>::GetQuantPreInt32()	1018	WeightFmt, IsHwMode>::GetQuantPreInt32()
1019	{	1019	{
1020	- // l0c (int32) -> ddr(fp16/int8) — for NPU_ARCH 5102	1020	+ // l0c (int32) -> ddr(fp16/int8) — quant pre-cast path
1021	if constexpr (AscendC::IsSameType<OutputT, half>::value) {	1021	if constexpr (AscendC::IsSameType<OutputT, half>::value) {
1022	if constexpr (AscendC::IsSameType<WeightT, int8_t>::value) {	1022	if constexpr (AscendC::IsSameType<WeightT, int8_t>::value) {
1023	uint8_t quantMode = (FixpipeIdx == 0) ? tiling_->quantMode0 : tiling_->quantMode1;	1023	uint8_t quantMode = (FixpipeIdx == 0) ? tiling_->quantMode0 : tiling_->quantMode1;
	@@ -1106,7 +1106,7 @@ __aicore__ inline void Conv2dSmallKernel<FmapType, weightType, biasType, out0Typ
1106	outputGm.SetGlobalBuffer(reinterpret_cast<__gm__ OutputT*>(yAddr) + batchOutOff + nOutOff);	1106	outputGm.SetGlobalBuffer(reinterpret_cast<__gm__ OutputT*>(yAddr) + batchOutOff + nOutOff);
1107		1107
1108	FixpipeParamsC310<Layout> fp;	1108	FixpipeParamsC310<Layout> fp;
1109	-#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 5102)	1109	+#if defined(__DAV_35_FAMILY__)
1110	if constexpr (AscendC::IsSameType<weightType, half>::value) {	1110	if constexpr (AscendC::IsSameType<weightType, half>::value) {
1111	fp.fixShiftVal = FIX_SHIFT_VAL_LEN_A16W16 - tiling_->fixedShiftValue;	1111	fp.fixShiftVal = FIX_SHIFT_VAL_LEN_A16W16 - tiling_->fixedShiftValue;
1112	}	1112	}
	@@ -1136,7 +1136,7 @@ __aicore__ inline void Conv2dSmallKernel<FmapType, weightType, biasType, out0Typ
1136	fp.params.srcNzC0Stride = 1;	1136	fp.params.srcNzC0Stride = 1;
1137	}	1137	}
1138		1138
1139	-#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 5102)	1139	+#if defined(__DAV_35_FAMILY__)
1140	fp.preReluMode = static_cast<ReluMode>(reluMode);	1140	fp.preReluMode = static_cast<ReluMode>(reluMode);
1141	if (reluMode == static_cast<uint8_t>(ReluMode::SCALAR_RELU)) {	1141	if (reluMode == static_cast<uint8_t>(ReluMode::SCALAR_RELU)) {
1142	float m2 = IsOutput0 ? reluWeight0Gm_.GetValue(0) : reluWeight1Gm_.GetValue(0);	1142	float m2 = IsOutput0 ? reluWeight0Gm_.GetValue(0) : reluWeight1Gm_.GetValue(0);

Mconv/conv2d_v2/op_kernel/arch35/conv2d_small_kernel_parallelism.h+2-2

	@@ -683,7 +683,7 @@ Conv2dSmallKernelParallelism<FmapType, weightType, biasType, out0Type, out1Type,
683		683
684	LocalTensor<L0cT> cl0(TPosition::CO1, 0, this->L0C_ELEMS);	684	LocalTensor<L0cT> cl0(TPosition::CO1, 0, this->L0C_ELEMS);
685	MmadParams mp;	685	MmadParams mp;
686	-#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 5102)	686	+#if defined(__DAV_35_FAMILY__)
687	if constexpr (AscendC::IsSameType<FmapType, half>::value) {	687	if constexpr (AscendC::IsSameType<FmapType, half>::value) {
688	mp.fixShiftVal = this->tiling_->fixedShiftValue;	688	mp.fixShiftVal = this->tiling_->fixedShiftValue;
689	}	689	}
	@@ -777,7 +777,7 @@ __aicore__ inline void Conv2dSmallKernelParallelism<FmapType, weightType, biasTy
777	LocalTensor<L0cT> cl0(TPosition::CO1, 0, this->L0C_ELEMS);	777	LocalTensor<L0cT> cl0(TPosition::CO1, 0, this->L0C_ELEMS);
778		778
779	MmadParams mp;	779	MmadParams mp;
780	-#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 5102)	780	+#if defined(__DAV_35_FAMILY__)
781	if constexpr (AscendC::IsSameType<FmapType, half>::value) {	781	if constexpr (AscendC::IsSameType<FmapType, half>::value) {
782	mp.fixShiftVal = this->tiling_->fixedShiftValue;	782	mp.fixShiftVal = this->tiling_->fixedShiftValue;
783	}	783	}

Mconv/conv2d_v2/op_kernel/arch35/conv2d_v2.h+2-2

	@@ -145,7 +145,7 @@ public:
145	(IsSameType<FMAP_T, hifloat8_t>::value) \|\|	145	(IsSameType<FMAP_T, hifloat8_t>::value) \|\|
146	(IsSameType<FMAP_T, fp8_e4m3fn_t>::value);	146	(IsSameType<FMAP_T, fp8_e4m3fn_t>::value);
147	constexpr static int8_t IS_EXTEND_CONV2D = CONV_CFG::isExtendConv2d;	147	constexpr static int8_t IS_EXTEND_CONV2D = CONV_CFG::isExtendConv2d;
148	-#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 5102)	148	+#if defined(__DAV_35_FAMILY__)
149	constexpr static bool IS_OPTGROUP_PRELOAD = false;	149	constexpr static bool IS_OPTGROUP_PRELOAD = false;
150	#else	150	#else
151	constexpr static bool IS_OPTGROUP_PRELOAD = (CONV_CFG::fmapTiling ==	151	constexpr static bool IS_OPTGROUP_PRELOAD = (CONV_CFG::fmapTiling ==
	@@ -323,4 +323,4 @@ Conv2dBase<FMAP_TYPE, WEIGHT_TYPE, OUTPUT_TYPE, BIAS_TYPE, SCALE_TYPE, CONV_CFG>
323	conv.End();	323	conv.End();
324	}	324	}
325		325
326	-#endif // CONV_2D_H	326	+#endif // CONV_2D_H

Mconv/conv2d_v2/op_kernel/arch35/conv2d_v2_api_impl.h+2-2

	@@ -108,7 +108,7 @@ public:
108	Conv2dFunc::C04LoadUB2L1Tools<Intf> c04LoadUB2L1Tools;	108	Conv2dFunc::C04LoadUB2L1Tools<Intf> c04LoadUB2L1Tools;
109	Conv2dFunc::C04ProcessTools<Intf> c04ProcessTools;	109	Conv2dFunc::C04ProcessTools<Intf> c04ProcessTools;
110		110
111	-#if defined(__NPU_ARCH__) && (__NPU_ARCH__ != 5102)	111	+#if !defined(__DAV_35_FAMILY__)
112	// Used in weight ub trans mode	112	// Used in weight ub trans mode
113	Conv2dFunc::WeightLoadGM2UBTools<Intf> weightUbLoadGM2UBTools;	113	Conv2dFunc::WeightLoadGM2UBTools<Intf> weightUbLoadGM2UBTools;
114	Conv2dFunc::WeightND2NZTools<Intf> weightUbTransND2NZTools;	114	Conv2dFunc::WeightND2NZTools<Intf> weightUbTransND2NZTools;
	@@ -191,4 +191,4 @@ public:
191	};	191	};
192	} // namespace conv2d	192	} // namespace conv2d
193		193
194	-#endif // CONV2D_V2_API_IMPL_H	194	+#endif // CONV2D_V2_API_IMPL_H

Mconv/conv2d_v2/op_kernel/arch35/conv2d_v2_intf.h+2-2

	@@ -81,7 +81,7 @@ struct Conv2dIntf {
81	constexpr static bool kPreLoadABFlag = ConvParam::l1PingPong == static_cast<int8_t>(ConvL1PingPong::ALL_OPEN) &&	81	constexpr static bool kPreLoadABFlag = ConvParam::l1PingPong == static_cast<int8_t>(ConvL1PingPong::ALL_OPEN) &&
82	ConvParam::groupType == static_cast<int8_t>(ConvGroupType::NORMAL_CONV);	82	ConvParam::groupType == static_cast<int8_t>(ConvGroupType::NORMAL_CONV);
83	constexpr static bool kPreLoadFlag = kPreLoadAFlag \|\| kPreLoadBFlag \|\| kPreLoadABFlag;	83	constexpr static bool kPreLoadFlag = kPreLoadAFlag \|\| kPreLoadBFlag \|\| kPreLoadABFlag;
84	-#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 5102)	84	+#if defined(__DAV_35_FAMILY__)
85	constexpr static bool groupOptPreloadFlag = false;	85	constexpr static bool groupOptPreloadFlag = false;
86	#else	86	#else
87	constexpr static bool groupOptPreloadFlag = ConvParam::groupType ==	87	constexpr static bool groupOptPreloadFlag = ConvParam::groupType ==
	@@ -352,4 +352,4 @@ public:
352		352
353	} // namespace conv2d	353	} // namespace conv2d
354		354
355	-#endif // CONV2D_V2_INTF_H	355	+#endif // CONV2D_V2_INTF_H

Mconv/conv2d_v2/tests/ut/op_graph/test_conv2d_to_conv2dv2_fusion_pass.cpp+10-10

	@@ -30,10 +30,10 @@ using namespace test_conv_fusion_framework;
30		30
31	class Conv2dToConv2dV2FusionPassTest : public testing::Test {	31	class Conv2dToConv2dV2FusionPassTest : public testing::Test {
32	public:	32	public:
33	- GraphPtr BuildSingleConvGraph(const char* graphName, bool useDav5102, const Conv2DConfig& convCfg)	33	+ GraphPtr BuildSingleConvGraph(const char* graphName, bool useFuseSoc, const Conv2DConfig& convCfg)
34	{	34	{
35	TestGraph builder(graphName);	35	TestGraph builder(graphName);
36	- if (useDav5102) {	36	+ if (useFuseSoc) {
37	builder.SetSocMC62();	37	builder.SetSocMC62();
38	} else {	38	} else {
39	builder.SetSocAscend950();	39	builder.SetSocAscend950();
	@@ -88,13 +88,13 @@ protected:
88	};	88	};
89		89
90	// ==========================================================================================	90	// ==========================================================================================
91	-// MeetRequirements dtype matrix success: DAV_3510 / DAV_5102 × dtype × bias	91	+// MeetRequirements dtype matrix success: Ascend950 / MC62 × dtype × bias
92	// ==========================================================================================	92	// ==========================================================================================
93	TEST_F(Conv2dToConv2dV2FusionPassTest, conv2d_to_conv2dv2_fusion_success)	93	TEST_F(Conv2dToConv2dV2FusionPassTest, conv2d_to_conv2dv2_fusion_success)
94	{	94	{
95	struct {	95	struct {
96	const char* pointName;	96	const char* pointName;
97	- bool useDav5102;	97	+ bool useFuseSoc;
98	DataType ioDtype;	98	DataType ioDtype;
99	DataType outputDtype;	99	DataType outputDtype;
100	bool hasBias;	100	bool hasBias;
	@@ -108,14 +108,14 @@ TEST_F(Conv2dToConv2dV2FusionPassTest, conv2d_to_conv2dv2_fusion_success)
108	{"dav3510_bf16_bias", false, DT_BF16, DT_BF16, true, DT_BF16},	108	{"dav3510_bf16_bias", false, DT_BF16, DT_BF16, true, DT_BF16},
109	{"dav3510_hifloat8", false, DT_HIFLOAT8, DT_HIFLOAT8, false, DT_FLOAT},	109	{"dav3510_hifloat8", false, DT_HIFLOAT8, DT_HIFLOAT8, false, DT_FLOAT},
110	{"dav3510_hifloat8_bias", false, DT_HIFLOAT8, DT_HIFLOAT8, true, DT_FLOAT},	110	{"dav3510_hifloat8_bias", false, DT_HIFLOAT8, DT_HIFLOAT8, true, DT_FLOAT},
111	- {"dav5102_fp16", true, DT_FLOAT16, DT_FLOAT16, false, DT_FLOAT16},	111	+ {"fuse_fp16", true, DT_FLOAT16, DT_FLOAT16, false, DT_FLOAT16},
112	- {"dav5102_fp16_bias", true, DT_FLOAT16, DT_FLOAT16, true, DT_FLOAT16},	112	+ {"fuse_fp16_bias", true, DT_FLOAT16, DT_FLOAT16, true, DT_FLOAT16},
113	};	113	};
114		114
115	for (const auto& p : points) {	115	for (const auto& p : points) {
116	SCOPED_TRACE(p.pointName);	116	SCOPED_TRACE(p.pointName);
117	std::string name = std::string("conv2d_to_conv2dv2_fusion_success_") + p.pointName;	117	std::string name = std::string("conv2d_to_conv2dv2_fusion_success_") + p.pointName;
118	- auto graph = BuildSingleConvGraph(name.c_str(), p.useDav5102,	118	+ auto graph = BuildSingleConvGraph(name.c_str(), p.useFuseSoc,
119	MakeConvCfg(p.ioDtype, p.outputDtype, p.hasBias, p.biasDtype));	119	MakeConvCfg(p.ioDtype, p.outputDtype, p.hasBias, p.biasDtype));
120	EXPECT_TRUE(GraphChecker::HasNode(graph, "Conv2D"));	120	EXPECT_TRUE(GraphChecker::HasNode(graph, "Conv2D"));
121	TestTotalPass(name, graph, SUCCESS);	121	TestTotalPass(name, graph, SUCCESS);
	@@ -141,9 +141,9 @@ TEST_F(Conv2dToConv2dV2FusionPassTest, conv2d_to_conv2dv2_no_fusion)
141	return BuildSingleConvGraph("conv2d_to_conv2dv2_no_fusion_bad_bias_dtype", false,	141	return BuildSingleConvGraph("conv2d_to_conv2dv2_no_fusion_bad_bias_dtype", false,
142	Conv2DConfig::Basic("Conv2D", DT_FLOAT16, DT_FLOAT16).WithBias(DT_INT32));	142	Conv2DConfig::Basic("Conv2D", DT_FLOAT16, DT_FLOAT16).WithBias(DT_INT32));
143	}},	143	}},
144	- {"dav5102_int8_proto_unsupported",	144	+ {"fuse_int8_proto_unsupported",
145	[this]() {	145	[this]() {
146	- return BuildSingleConvGraph("conv2d_to_conv2dv2_no_fusion_dav5102_int8_proto_unsupported", true,	146	+ return BuildSingleConvGraph("conv2d_to_conv2dv2_no_fusion_fuse_int8_proto_unsupported", true,
147	Conv2DConfig::Basic("Conv2D", DT_INT8, DT_INT32));	147	Conv2DConfig::Basic("Conv2D", DT_INT8, DT_INT32));
148	}},	148	}},
149	{"filter_dtype_mismatch",	149	{"filter_dtype_mismatch",
	@@ -416,4 +416,4 @@ TEST_F(Conv2dToConv2dV2FusionPassTest, conv2d_to_conv2dv2_graph_topology)
416	}	416	}
417	}	417	}
418		418
419	-#endif	419	+#endif

Mconv/conv3d_v2/op_kernel/conv3d_v2_template.h+5-2

	@@ -15,14 +15,17 @@
15		15
16	#ifndef CONV3D_V2_TEMPLATE_H	16	#ifndef CONV3D_V2_TEMPLATE_H
17	#define CONV3D_V2_TEMPLATE_H	17	#define CONV3D_V2_TEMPLATE_H
18	-#if defined(__NPU_ARCH__) && ((__NPU_ARCH__ == 3510) \|\| (__NPU_ARCH__ == 5102))	18	+#if defined(__NPU_ARCH__) && ((__NPU_ARCH__ == 3510) \|\| defined(__DAV_35_FAMILY__))
19	#include "arch35/conv3d_v2.h"	19	#include "arch35/conv3d_v2.h"
20	#endif	20	#endif
21	#include "kernel_operator.h"	21	#include "kernel_operator.h"
22	#include "conv3d_v2_tiling_data.h"	22	#include "conv3d_v2_tiling_data.h"
23	#include "../common/arch35/conv_config.h"	23	#include "../common/arch35/conv_config.h"
		24	+#include "../inc/macro.h"
		25	+
24	using namespace AscendC;	26	using namespace AscendC;
25	using namespace conv;	27	using namespace conv;
		28	+
26	constexpr ConvFormat fmapFormat = ConvFormat::NCDHW;	29	constexpr ConvFormat fmapFormat = ConvFormat::NCDHW;
27	constexpr ConvFormat filterFormat = ConvFormat::NCDHW;	30	constexpr ConvFormat filterFormat = ConvFormat::NCDHW;
28	constexpr ConvFormat outputFormat = ConvFormat::NCDHW;	31	constexpr ConvFormat outputFormat = ConvFormat::NCDHW;
	@@ -49,4 +52,4 @@ __global__ __aicore__ void conv3dv2_template(GM_ADDR x, GM_ADDR filter, GM_ADDR
49	return;	52	return;
50	}	53	}
51		54
52	-#endif // CONV3D_V2_TEMPLATE_H	55	+#endif // CONV3D_V2_TEMPLATE_H

Mconv/convolution_forward/op_graph/fusion_pass/depthwise_to_conv2d_fusion_pass.cpp+40-31

	@@ -37,7 +37,8 @@ bool DepthwiseToConv2dFusionPass::CheckDynamicShape() const
37	return true;	37	return true;
38	}	38	}
39		39
40	- FUSION_PASS_CHECK(!ConvFusionUtilsPass::IsUnknownShape(convDescInfo.fmapDesc),	40	+ FUSION_PASS_CHECK(
		41	+ !ConvFusionUtilsPass::IsUnknownShape(convDescInfo.fmapDesc),
41	OP_LOGD(FUSION_NAME, "%s only support dynamic mode on non-ND soc.", convDescInfo.nodeNameStr.c_str()),	42	OP_LOGD(FUSION_NAME, "%s only support dynamic mode on non-ND soc.", convDescInfo.nodeNameStr.c_str()),
42	return false);	43	return false);
43		44
	@@ -48,7 +49,8 @@ bool DepthwiseToConv2dFusionPass::GetFmapChannel()
48	{	49	{
49	auto originFormat = convDescInfo.fmapDesc.GetOriginFormat();	50	auto originFormat = convDescInfo.fmapDesc.GetOriginFormat();
50	auto inputShape = convDescInfo.fmapDesc.GetOriginShape().GetDims();	51	auto inputShape = convDescInfo.fmapDesc.GetOriginShape().GetDims();
51	- FUSION_PASS_CHECK(inputShape.size() != MAX_DIM_NUM,	52	+ FUSION_PASS_CHECK(
		53	+ inputShape.size() != MAX_DIM_NUM,
52	OP_LOGE(FUSION_NAME, "%s fmap origin shape dim not equal to 4.", convDescInfo.nodeNameStr.c_str()),	54	OP_LOGE(FUSION_NAME, "%s fmap origin shape dim not equal to 4.", convDescInfo.nodeNameStr.c_str()),
53	return false);	55	return false);
54		56
	@@ -58,42 +60,45 @@ bool DepthwiseToConv2dFusionPass::GetFmapChannel()
58	fmapChannel = inputShape[FMAP_CHANNEL_NHWC_INDEX];	60	fmapChannel = inputShape[FMAP_CHANNEL_NHWC_INDEX];
59	} else {	61	} else {
60	OP_LOGE(FUSION_NAME, "%s not support format [%s].", convDescInfo.nodeNameStr.c_str(),	62	OP_LOGE(FUSION_NAME, "%s not support format [%s].", convDescInfo.nodeNameStr.c_str(),
61	- TypeUtils::FormatToAscendString(originFormat).GetString());	63	+ TypeUtils::FormatToAscendString(originFormat).GetString());
62	return false;	64	return false;
63	}	65	}
64		66
65	FUSION_PASS_CHECK(fmapChannel == -1,	67	FUSION_PASS_CHECK(fmapChannel == -1,
66	- OP_LOGD(FUSION_NAME, "%s not support fmapChannel is -1.", convDescInfo.nodeNameStr.c_str()), return false);	68	+ OP_LOGD(FUSION_NAME, "%s not support fmapChannel is -1.", convDescInfo.nodeNameStr.c_str()),
		69	+ return false);
67		70
68	return true;	71	return true;
69	}	72	}
70		73
71	-bool DepthwiseToConv2dFusionPass::GetDepthwiseConvAttrs(const GNode &depthwiseNode)	74	+bool DepthwiseToConv2dFusionPass::GetDepthwiseConvAttrs(const GNode& depthwiseNode)
72	{	75	{
73	depthwiseAttrs = ConvBaseAttrs();	76	depthwiseAttrs = ConvBaseAttrs();
74	FUSION_PASS_CHECK(depthwiseNode.GetAttr(STRIDES, depthwiseAttrs.strides) != GRAPH_SUCCESS,	77	FUSION_PASS_CHECK(depthwiseNode.GetAttr(STRIDES, depthwiseAttrs.strides) != GRAPH_SUCCESS,
75	- OP_LOGE(FUSION_NAME, "%s get strides failed.", convDescInfo.nodeNameStr.c_str()), return false);	78	+ OP_LOGE(FUSION_NAME, "%s get strides failed.", convDescInfo.nodeNameStr.c_str()), return false);
76	FUSION_PASS_CHECK(depthwiseNode.GetAttr(PADS, depthwiseAttrs.pads) != GRAPH_SUCCESS,	79	FUSION_PASS_CHECK(depthwiseNode.GetAttr(PADS, depthwiseAttrs.pads) != GRAPH_SUCCESS,
77	- OP_LOGE(FUSION_NAME, "%s get pads failed.", convDescInfo.nodeNameStr.c_str()), return false);	80	+ OP_LOGE(FUSION_NAME, "%s get pads failed.", convDescInfo.nodeNameStr.c_str()), return false);
78	FUSION_PASS_CHECK(depthwiseNode.GetAttr(DILATIONS, depthwiseAttrs.dilations) != GRAPH_SUCCESS,	81	FUSION_PASS_CHECK(depthwiseNode.GetAttr(DILATIONS, depthwiseAttrs.dilations) != GRAPH_SUCCESS,
79	- OP_LOGE(FUSION_NAME, "%s get dilations failed.", convDescInfo.nodeNameStr.c_str()), return false);	82	+ OP_LOGE(FUSION_NAME, "%s get dilations failed.", convDescInfo.nodeNameStr.c_str()), return false);
80	FUSION_PASS_CHECK(depthwiseNode.GetAttr(DATA_FORMAT, depthwiseAttrs.dataFormat) != GRAPH_SUCCESS,	83	FUSION_PASS_CHECK(depthwiseNode.GetAttr(DATA_FORMAT, depthwiseAttrs.dataFormat) != GRAPH_SUCCESS,
81	- OP_LOGE(FUSION_NAME, "%s get data_format failed.", convDescInfo.nodeNameStr.c_str()), return false);	84	+ OP_LOGE(FUSION_NAME, "%s get data_format failed.", convDescInfo.nodeNameStr.c_str()),
		85	+ return false);
82		86
83	if (isNdSoc) {	87	if (isNdSoc) {
84	FUSION_PASS_CHECK(depthwiseNode.GetAttr(OFFSET_X, depthwiseAttrs.offsetX) != GRAPH_SUCCESS,	88	FUSION_PASS_CHECK(depthwiseNode.GetAttr(OFFSET_X, depthwiseAttrs.offsetX) != GRAPH_SUCCESS,
85	- OP_LOGE(FUSION_NAME, "%s get offset_x failed.", convDescInfo.nodeNameStr.c_str()), return false);	89	+ OP_LOGE(FUSION_NAME, "%s get offset_x failed.", convDescInfo.nodeNameStr.c_str()),
		90	+ return false);
86	}	91	}
87		92
88	depthwiseNode.GetAttr(PADDING, depthwiseAttrs.padding);	93	depthwiseNode.GetAttr(PADDING, depthwiseAttrs.padding);
89	return true;	94	return true;
90	}	95	}
91		96
92	-bool DepthwiseToConv2dFusionPass::MeetRequirements(const GNode &depthwiseNode)	97	+bool DepthwiseToConv2dFusionPass::MeetRequirements(const GNode& depthwiseNode)
93	{	98	{
94	InitMember();	99	InitMember();
95		100
96	- isNdSoc = ConvFusionUtilsPass::CheckSocList(ND_SOC_LIST, npuArch);	101	+ isNdSoc = ConvFusionUtilsPass::CheckSocList(ND_SOC_LIST, npuArch, true);
97	FUSION_PASS_CHECK_NOLOG(!ConvFusionUtilsPass::GetConvDescInfo(depthwiseNode, convDescInfo), return false);	102	FUSION_PASS_CHECK_NOLOG(!ConvFusionUtilsPass::GetConvDescInfo(depthwiseNode, convDescInfo), return false);
98	OP_LOGD(convDescInfo.nodeNameStr, "Begin to do DepthwiseToConv2dFusionPass.");	103	OP_LOGD(convDescInfo.nodeNameStr, "Begin to do DepthwiseToConv2dFusionPass.");
99		104
	@@ -103,7 +108,7 @@ bool DepthwiseToConv2dFusionPass::MeetRequirements(const GNode &depthwiseNode)
103	return true;	108	return true;
104	}	109	}
105		110
106	-GraphUniqPtr DepthwiseToConv2dFusionPass::Replacement(const GNode &depthwiseNode)	111	+GraphUniqPtr DepthwiseToConv2dFusionPass::Replacement(const GNode& depthwiseNode)
107	{	112	{
108	FUSION_PASS_CHECK_NOLOG(!GetDepthwiseConvAttrs(depthwiseNode), return nullptr);	113	FUSION_PASS_CHECK_NOLOG(!GetDepthwiseConvAttrs(depthwiseNode), return nullptr);
109		114
	@@ -114,41 +119,45 @@ GraphUniqPtr DepthwiseToConv2dFusionPass::Replacement(const GNode &depthwiseNode
114	inputs.emplace_back(graphBuilder.CreateInput(static_cast<int64_t>(INPUT_BIAS_INDEX)));	119	inputs.emplace_back(graphBuilder.CreateInput(static_cast<int64_t>(INPUT_BIAS_INDEX)));
115	}	120	}
116		121
117	- auto *replaceGraph = graphBuilder.GetCGraphBuilder()->GetGraph();	122	+ auto* replaceGraph = graphBuilder.GetCGraphBuilder()->GetGraph();
118	GNode conv2dNode;	123	GNode conv2dNode;
119	- FUSION_PASS_CHECK(!ConvFusionUtilsPass::BuildConv2dNode(	124	+ FUSION_PASS_CHECK(!ConvFusionUtilsPass::BuildConv2dNode(replaceGraph, convDescInfo.nodeNameStr + "_To_Conv2D",
120	- replaceGraph, convDescInfo.nodeNameStr + "_To_Conv2D", inputs, conv2dNode),	125	+ inputs, conv2dNode),
121	- OP_LOGE(FUSION_NAME, "%s build Conv2D node failed.", convDescInfo.nodeNameStr.c_str()), return nullptr);	126	+ OP_LOGE(FUSION_NAME, "%s build Conv2D node failed.", convDescInfo.nodeNameStr.c_str()),
		127	+ return nullptr);
122		128
123	FUSION_PASS_CHECK_NOLOG(!SetConv2dAttrs(conv2dNode, depthwiseNode), return nullptr);	129	FUSION_PASS_CHECK_NOLOG(!SetConv2dAttrs(conv2dNode, depthwiseNode), return nullptr);
124	FUSION_PASS_CHECK_NOLOG(!UpdateConv2dDesc(conv2dNode), return nullptr);	130	FUSION_PASS_CHECK_NOLOG(!UpdateConv2dDesc(conv2dNode), return nullptr);
125		131
126	- auto *yHolder = graphBuilder.GetCGraphBuilder()->GetTensorHolderFromNode(conv2dNode, OUTPUT_INDEX);	132	+ auto* yHolder = graphBuilder.GetCGraphBuilder()->GetTensorHolderFromNode(conv2dNode, OUTPUT_INDEX);
127	- FUSION_PASS_CHECK(yHolder == nullptr,	133	+ FUSION_PASS_CHECK(
		134	+ yHolder == nullptr,
128	OP_LOGE(FUSION_NAME, "%s get Conv2D output tensor holder failed.", convDescInfo.nodeNameStr.c_str()),	135	OP_LOGE(FUSION_NAME, "%s get Conv2D output tensor holder failed.", convDescInfo.nodeNameStr.c_str()),
129	return nullptr);	136	return nullptr);
130		137
131	return graphBuilder.BuildAndReset({es::EsTensorHolder(yHolder)});	138	return graphBuilder.BuildAndReset({es::EsTensorHolder(yHolder)});
132	}	139	}
133		140
134	-bool DepthwiseToConv2dFusionPass::SetConv2dAttrs(GNode &conv2dNode, const GNode &depthwiseNode)	141	+bool DepthwiseToConv2dFusionPass::SetConv2dAttrs(GNode& conv2dNode, const GNode& depthwiseNode)
135	{	142	{
136	FUSION_PASS_CHECK(conv2dNode.SetAttr(STRIDES, depthwiseAttrs.strides) != GRAPH_SUCCESS,	143	FUSION_PASS_CHECK(conv2dNode.SetAttr(STRIDES, depthwiseAttrs.strides) != GRAPH_SUCCESS,
137	- OP_LOGE(FUSION_NAME, "%s set strides failed.", convDescInfo.nodeNameStr.c_str()), return false);	144	+ OP_LOGE(FUSION_NAME, "%s set strides failed.", convDescInfo.nodeNameStr.c_str()), return false);
138	FUSION_PASS_CHECK(conv2dNode.SetAttr(PADS, depthwiseAttrs.pads) != GRAPH_SUCCESS,	145	FUSION_PASS_CHECK(conv2dNode.SetAttr(PADS, depthwiseAttrs.pads) != GRAPH_SUCCESS,
139	- OP_LOGE(FUSION_NAME, "%s set pads failed.", convDescInfo.nodeNameStr.c_str()), return false);	146	+ OP_LOGE(FUSION_NAME, "%s set pads failed.", convDescInfo.nodeNameStr.c_str()), return false);
140	FUSION_PASS_CHECK(conv2dNode.SetAttr(DILATIONS, depthwiseAttrs.dilations) != GRAPH_SUCCESS,	147	FUSION_PASS_CHECK(conv2dNode.SetAttr(DILATIONS, depthwiseAttrs.dilations) != GRAPH_SUCCESS,
141	- OP_LOGE(FUSION_NAME, "%s set dilations failed.", convDescInfo.nodeNameStr.c_str()), return false);	148	+ OP_LOGE(FUSION_NAME, "%s set dilations failed.", convDescInfo.nodeNameStr.c_str()), return false);
142	FUSION_PASS_CHECK(conv2dNode.SetAttr(DATA_FORMAT, depthwiseAttrs.dataFormat) != GRAPH_SUCCESS,	149	FUSION_PASS_CHECK(conv2dNode.SetAttr(DATA_FORMAT, depthwiseAttrs.dataFormat) != GRAPH_SUCCESS,
143	- OP_LOGE(FUSION_NAME, "%s set data_format failed.", convDescInfo.nodeNameStr.c_str()), return false);	150	+ OP_LOGE(FUSION_NAME, "%s set data_format failed.", convDescInfo.nodeNameStr.c_str()),
		151	+ return false);
144	FUSION_PASS_CHECK(conv2dNode.SetAttr(GROUPS, fmapChannel) != GRAPH_SUCCESS,	152	FUSION_PASS_CHECK(conv2dNode.SetAttr(GROUPS, fmapChannel) != GRAPH_SUCCESS,
145	- OP_LOGE(FUSION_NAME, "%s set groups failed.", convDescInfo.nodeNameStr.c_str()), return false);	153	+ OP_LOGE(FUSION_NAME, "%s set groups failed.", convDescInfo.nodeNameStr.c_str()), return false);
146	if (isNdSoc) {	154	if (isNdSoc) {
147	FUSION_PASS_CHECK(conv2dNode.SetAttr(OFFSET_X, depthwiseAttrs.offsetX) != GRAPH_SUCCESS,	155	FUSION_PASS_CHECK(conv2dNode.SetAttr(OFFSET_X, depthwiseAttrs.offsetX) != GRAPH_SUCCESS,
148	- OP_LOGE(FUSION_NAME, "%s set offset_x failed.", convDescInfo.nodeNameStr.c_str()), return false);	156	+ OP_LOGE(FUSION_NAME, "%s set offset_x failed.", convDescInfo.nodeNameStr.c_str()),
		157	+ return false);
149	}	158	}
150	FUSION_PASS_CHECK(conv2dNode.SetAttr(PADDING, depthwiseAttrs.padding) != GRAPH_SUCCESS,	159	FUSION_PASS_CHECK(conv2dNode.SetAttr(PADDING, depthwiseAttrs.padding) != GRAPH_SUCCESS,
151	- OP_LOGE(FUSION_NAME, "%s set padding failed.", convDescInfo.nodeNameStr.c_str()), return false);	160	+ OP_LOGE(FUSION_NAME, "%s set padding failed.", convDescInfo.nodeNameStr.c_str()), return false);
152		161
153	int64_t opImplModeEnum = 0;	162	int64_t opImplModeEnum = 0;
154	if (depthwiseNode.GetAttr(OP_IMPL_MODE_ENUM, opImplModeEnum) == GRAPH_SUCCESS) {	163	if (depthwiseNode.GetAttr(OP_IMPL_MODE_ENUM, opImplModeEnum) == GRAPH_SUCCESS) {
	@@ -158,10 +167,11 @@ bool DepthwiseToConv2dFusionPass::SetConv2dAttrs(GNode &conv2dNode, const GNode
158	return true;	167	return true;
159	}	168	}
160		169
161	-bool DepthwiseToConv2dFusionPass::UpdateConv2dDesc(GNode &conv2dNode)	170	+bool DepthwiseToConv2dFusionPass::UpdateConv2dDesc(GNode& conv2dNode)
162	{	171	{
163	FUSION_PASS_CHECK_NOLOG(!ConvFusionUtilsPass::UpdateInputDesc(&conv2dNode, convDescInfo), return false);	172	FUSION_PASS_CHECK_NOLOG(!ConvFusionUtilsPass::UpdateInputDesc(&conv2dNode, convDescInfo), return false);
164	- FUSION_PASS_CHECK(conv2dNode.UpdateOutputDesc(OUTPUT_INDEX, convDescInfo.outputDesc) != GRAPH_SUCCESS,	173	+ FUSION_PASS_CHECK(
		174	+ conv2dNode.UpdateOutputDesc(OUTPUT_INDEX, convDescInfo.outputDesc) != GRAPH_SUCCESS,
165	OP_LOGE(FUSION_NAME, "%s update Conv2D output tensor desc failed.", convDescInfo.nodeNameStr.c_str()),	175	OP_LOGE(FUSION_NAME, "%s update Conv2D output tensor desc failed.", convDescInfo.nodeNameStr.c_str()),
166	return false);	176	return false);
167		177
	@@ -169,8 +179,7 @@ bool DepthwiseToConv2dFusionPass::UpdateConv2dDesc(GNode &conv2dNode)
169	}	179	}
170		180
171	#if GE_COMPILER_VERSION_NUM >= 90000000U	181	#if GE_COMPILER_VERSION_NUM >= 90000000U
172	-REG_DECOMPOSE_PASS(DepthwiseToConv2dFusionPass, {DEPTHWISE_CONV2D})	182	+REG_DECOMPOSE_PASS(DepthwiseToConv2dFusionPass, {DEPTHWISE_CONV2D}).Stage(CustomPassStage::kCompatibleInherited);
173	- .Stage(CustomPassStage::kCompatibleInherited);
174	#endif	183	#endif
175		184
176	} // namespace Ops	185	} // namespace Ops

Mconv/convolution_forward/op_graph/fusion_pass/depthwise_to_conv2d_fusion_pass.h+8-12

	@@ -29,28 +29,24 @@ constexpr int32_t MAX_DIM_NUM = 4;
29	constexpr int32_t FMAP_CHANNEL_NCHW_INDEX = 1;	29	constexpr int32_t FMAP_CHANNEL_NCHW_INDEX = 1;
30	constexpr int32_t FMAP_CHANNEL_NHWC_INDEX = 3;	30	constexpr int32_t FMAP_CHANNEL_NHWC_INDEX = 3;
31		31
32	-const std::map<std::string, NpuArch> ND_SOC_LIST = {	32	+const std::map<std::string, NpuArch> ND_SOC_LIST = {{"Ascend950", NpuArch::DAV_3510}};
33	- {"Ascend950", NpuArch::DAV_3510},
34	- {"MC62", NpuArch::DAV_5102}
35	-};
36	} // namespace DepthwiseToConv2dFusion	33	} // namespace DepthwiseToConv2dFusion
37		34
38	class __attribute__((visibility("default"))) DepthwiseToConv2dFusionPass : public ge::fusion::DecomposePass {	35	class __attribute__((visibility("default"))) DepthwiseToConv2dFusionPass : public ge::fusion::DecomposePass {
39	public:	36	public:
40	- explicit DepthwiseToConv2dFusionPass(const std::vector<ge::AscendString> &opTypes)	37	+ explicit DepthwiseToConv2dFusionPass(const std::vector<ge::AscendString>& opTypes) : DecomposePass(opTypes) {}
41	- : DecomposePass(opTypes) {}
42		38
43	protected:	39	protected:
44	- bool MeetRequirements(const ge::GNode &depthwiseNode) override;	40	+ bool MeetRequirements(const ge::GNode& depthwiseNode) override;
45	- ge::fusion::GraphUniqPtr Replacement(const ge::GNode &depthwiseNode) override;	41	+ ge::fusion::GraphUniqPtr Replacement(const ge::GNode& depthwiseNode) override;
46		42
47	private:	43	private:
48	void InitMember();	44	void InitMember();
49	bool CheckDynamicShape() const;	45	bool CheckDynamicShape() const;
50	bool GetFmapChannel();	46	bool GetFmapChannel();
51	- bool GetDepthwiseConvAttrs(const ge::GNode &depthwiseNode);	47	+ bool GetDepthwiseConvAttrs(const ge::GNode& depthwiseNode);
52	- bool SetConv2dAttrs(ge::GNode &conv2dNode, const ge::GNode &depthwiseNode);	48	+ bool SetConv2dAttrs(ge::GNode& conv2dNode, const ge::GNode& depthwiseNode);
53	- bool UpdateConv2dDesc(ge::GNode &conv2dNode);	49	+ bool UpdateConv2dDesc(ge::GNode& conv2dNode);
54		50
55	NpuArch npuArch = NpuArch::DAV_RESV;	51	NpuArch npuArch = NpuArch::DAV_RESV;
56	bool isNdSoc = false;	52	bool isNdSoc = false;
	@@ -62,4 +58,4 @@ private:
62	} // namespace Conv	58	} // namespace Conv
63	} // namespace NN	59	} // namespace NN
64	} // namespace Ops	60	} // namespace Ops
65	-#endif // DEPTHWISE_TO_CONV2D_FUSION_PASS_H	61	+#endif // DEPTHWISE_TO_CONV2D_FUSION_PASS_H

Mconv/convolution_forward/tests/ut/op_graph/test_depthwise_to_conv2d_fusion_pass.cpp+184-134

	@@ -30,8 +30,9 @@ using namespace test_conv_fusion_framework;
30	#define CONV_DEBUG false	30	#define CONV_DEBUG false
31		31
32	namespace {	32	namespace {
33	-GraphPtr BuildDepthwiseGraph(const char *graphName, const DepthwiseConv2DConfig &depthwiseCfg,	33	+GraphPtr BuildDepthwiseGraph(const char* graphName, const DepthwiseConv2DConfig& depthwiseCfg,
34	- const SocConfig *socConfig = nullptr) {	34	+ const SocConfig* socConfig = nullptr)
		35	+{
35	TestGraph builder(graphName);	36	TestGraph builder(graphName);
36	if (socConfig != nullptr) {	37	if (socConfig != nullptr) {
37	builder.SetSoc(*socConfig);	38	builder.SetSoc(*socConfig);
	@@ -47,16 +48,12 @@ const SocConfig kNonNdSocConfig("Ascend910B", "Ascend910B1");
47		48
48	class DepthwiseToConv2dFusionPassTest : public testing::Test {	49	class DepthwiseToConv2dFusionPassTest : public testing::Test {
49	protected:	50	protected:
50	- static void SetUpTestCase() {	51	+ static void SetUpTestCase() { std::cout << "DepthwiseToConv2dFusionPassTest SetUp" << std::endl; }
51	- std::cout << "DepthwiseToConv2dFusionPassTest SetUp" << std::endl;
52	- }
53		52
54	- static void TearDownTestCase() {	53	+ static void TearDownTestCase() { std::cout << "DepthwiseToConv2dFusionPassTest TearDown" << std::endl; }
55	- std::cout << "DepthwiseToConv2dFusionPassTest TearDown" << std::endl;
56	- }
57		54
58	- void TestTotalPass(const std::string &passName, GraphPtr &graph, Status expectRes,	55	+ void TestTotalPass(const std::string& passName, GraphPtr& graph, Status expectRes, int32_t depthwiseCountBefore = 1)
59	- int32_t depthwiseCountBefore = 1) {	56	+ {
60	CustomPassContext passContext;	57	CustomPassContext passContext;
61	passContext.SetPassName(passName.c_str());	58	passContext.SetPassName(passName.c_str());
62	DepthwiseToConv2dFusionPass pass({DEPTHWISE_CONV2D});	59	DepthwiseToConv2dFusionPass pass({DEPTHWISE_CONV2D});
	@@ -77,38 +74,40 @@ protected:
77	}	74	}
78	}	75	}
79		76
80	- GNode GetFirstConv2dNode(GraphPtr &graph) {	77	+ GNode GetFirstConv2dNode(GraphPtr& graph)
		78	+ {
81	GNode fused;	79	GNode fused;
82	EXPECT_TRUE(GraphChecker::FindFirstNodeByOpType(graph, "Conv2D", fused));	80	EXPECT_TRUE(GraphChecker::FindFirstNodeByOpType(graph, "Conv2D", fused));
83	return fused;	81	return fused;
84	}	82	}
85	};	83	};
86		84
87	-TEST_F(DepthwiseToConv2dFusionPassTest, depthwise_to_conv2d_fusion_success) {	85	+TEST_F(DepthwiseToConv2dFusionPassTest, depthwise_to_conv2d_fusion_success)
		86	+{
88	struct Point {	87	struct Point {
89	- const char *pointName;	88	+ const char* pointName;
90	- const SocConfig *socConfig;	89	+ const SocConfig* socConfig;
91	DepthwiseConv2DConfig cfg;	90	DepthwiseConv2DConfig cfg;
92	} const points[] = {	91	} const points[] = {
93	{"dynamic_nchw_dav3510", nullptr,	92	{"dynamic_nchw_dav3510", nullptr,
94	- DepthwiseConv2DConfig::Basic("depthwise_conv2d", DT_FLOAT, FORMAT_NCHW, {-1, 16, 256, 256},	93	+ DepthwiseConv2DConfig::Basic("depthwise_conv2d", DT_FLOAT, FORMAT_NCHW, {-1, 16, 256, 256}, {1, 1, 1, 64},
95	- {1, 1, 1, 64}, {-1, 4, 256, 256})},	94	+ {-1, 4, 256, 256})},
96	{"static_nchw_dav3510", nullptr,	95	{"static_nchw_dav3510", nullptr,
97	- DepthwiseConv2DConfig::Basic("depthwise_conv2d", DT_FLOAT, FORMAT_NCHW, {1, 16, 256, 256},	96	+ DepthwiseConv2DConfig::Basic("depthwise_conv2d", DT_FLOAT, FORMAT_NCHW, {1, 16, 256, 256}, {1, 1, 1, 64},
98	- {1, 1, 1, 64}, {1, 4, 256, 256})},	97	+ {1, 4, 256, 256})},
99	- {"dynamic_nchw_dav5102", &kMc62SocConfig,	98	+ {"dynamic_nchw_mc62", &kMc62SocConfig,
100	- DepthwiseConv2DConfig::Basic("depthwise_conv2d", DT_FLOAT16, FORMAT_NCHW, {-1, 16, 256, 256},	99	+ DepthwiseConv2DConfig::Basic("depthwise_conv2d", DT_FLOAT16, FORMAT_NCHW, {-1, 16, 256, 256}, {1, 1, 1, 64},
101	- {1, 1, 1, 64}, {-1, 4, 256, 256})},	100	+ {-1, 4, 256, 256})},
102	{"dynamic_nhwc_dav3510", nullptr,	101	{"dynamic_nhwc_dav3510", nullptr,
103	- DepthwiseConv2DConfig::Basic("depthwise_conv2d", DT_FLOAT, FORMAT_NHWC, {-1, 256, 256, 16},	102	+ DepthwiseConv2DConfig::Basic("depthwise_conv2d", DT_FLOAT, FORMAT_NHWC, {-1, 256, 256, 16}, {1, 1, 1, 64},
104	- {1, 1, 1, 64}, {-1, 256, 256, 4})},	103	+ {-1, 256, 256, 4})},
105	{"dynamic_nchw_with_bias", nullptr, DepthwiseConv2DConfig::Basic("depthwise_conv2d").WithBias()},	104	{"dynamic_nchw_with_bias", nullptr, DepthwiseConv2DConfig::Basic("depthwise_conv2d").WithBias()},
106	{"static_nchw_mc62", &kMc62SocConfig,	105	{"static_nchw_mc62", &kMc62SocConfig,
107	- DepthwiseConv2DConfig::Basic("depthwise_conv2d", DT_FLOAT16, FORMAT_NCHW, {1, 16, 256, 256},	106	+ DepthwiseConv2DConfig::Basic("depthwise_conv2d", DT_FLOAT16, FORMAT_NCHW, {1, 16, 256, 256}, {1, 1, 1, 64},
108	- {1, 1, 1, 64}, {1, 4, 256, 256})},	107	+ {1, 4, 256, 256})},
109	};	108	};
110		109
111	- for (const auto &p : points) {	110	+ for (const auto& p : points) {
112	SCOPED_TRACE(p.pointName);	111	SCOPED_TRACE(p.pointName);
113	std::string name = std::string("depthwise_to_conv2d_fusion_success_") + p.pointName;	112	std::string name = std::string("depthwise_to_conv2d_fusion_success_") + p.pointName;
114	auto graph = BuildDepthwiseGraph(name.c_str(), p.cfg, p.socConfig);	113	auto graph = BuildDepthwiseGraph(name.c_str(), p.cfg, p.socConfig);
	@@ -122,55 +121,70 @@ TEST_F(DepthwiseToConv2dFusionPassTest, depthwise_to_conv2d_fusion_success) {
122	}	121	}
123	}	122	}
124		123
125	-TEST_F(DepthwiseToConv2dFusionPassTest, depthwise_to_conv2d_no_fusion) {	124	+TEST_F(DepthwiseToConv2dFusionPassTest, depthwise_to_conv2d_no_fusion)
		125	+{
126	struct Point {	126	struct Point {
127	- const char *pointName;	127	+ const char* pointName;
128	std::function<GraphPtr()> build;	128	std::function<GraphPtr()> build;
129	Status expectRes;	129	Status expectRes;
130	int32_t depthwiseCountBefore;	130	int32_t depthwiseCountBefore;
131	} const points[] = {	131	} const points[] = {
132	- {"wrong_op_type", []() {	132	+ {"wrong_op_type",
		133	+ []() {
133	return TestGraph("depthwise_to_conv2d_no_fusion_wrong_op_type")	134	return TestGraph("depthwise_to_conv2d_no_fusion_wrong_op_type")
134	.SetSocAscend950()	135	.SetSocAscend950()
135	.AddConv2D(Conv2DConfig::Basic("conv2d", DT_FLOAT, DT_FLOAT, FORMAT_NCHW, {1, 16, 256, 256},	136	.AddConv2D(Conv2DConfig::Basic("conv2d", DT_FLOAT, DT_FLOAT, FORMAT_NCHW, {1, 16, 256, 256},
136	- {1, 1, 1, 64}, {1, 4, 256, 256}))	137	+ {1, 1, 1, 64}, {1, 4, 256, 256}))
137	.SetOutput("conv2d")	138	.SetOutput("conv2d")
138	.Build();	139	.Build();
139	- }, GRAPH_NOT_CHANGED, 0},	140	+ },
140	- {"static_shape_non_nd_soc", []() {	141	+ GRAPH_NOT_CHANGED, 0},
141	- return BuildDepthwiseGraph("depthwise_to_conv2d_no_fusion_static_non_nd",	142	+ {"static_shape_non_nd_soc",
		143	+ []() {
		144	+ return BuildDepthwiseGraph(
		145	+ "depthwise_to_conv2d_no_fusion_static_non_nd",
142	DepthwiseConv2DConfig::Basic("depthwise_conv2d", DT_FLOAT, FORMAT_NCHW, {1, 16, 256, 256},	146	DepthwiseConv2DConfig::Basic("depthwise_conv2d", DT_FLOAT, FORMAT_NCHW, {1, 16, 256, 256},
143	- {1, 1, 1, 64}, {1, 4, 256, 256}),	147	+ {1, 1, 1, 64}, {1, 4, 256, 256}),
144	&kNonNdSocConfig);	148	&kNonNdSocConfig);
145	- }, GRAPH_NOT_CHANGED, 1},	149	+ },
146	- {"unsupported_format", []() {	150	+ GRAPH_NOT_CHANGED, 1},
		151	+ {"unsupported_format",
		152	+ []() {
147	DepthwiseConv2DConfig cfg = DepthwiseConv2DConfig::Basic("depthwise_conv2d");	153	DepthwiseConv2DConfig cfg = DepthwiseConv2DConfig::Basic("depthwise_conv2d");
148	cfg.inputs[0].format = FORMAT_NC1HWC0;	154	cfg.inputs[0].format = FORMAT_NC1HWC0;
149	- cfg.inputs[0].tensorDesc = BuildTensorDesc(DT_FLOAT, FORMAT_NC1HWC0, {-1, 16, 256, 256},	155	+ cfg.inputs[0].tensorDesc = BuildTensorDesc(DT_FLOAT, FORMAT_NC1HWC0, {-1, 16, 256, 256}, FORMAT_NC1HWC0,
150	- FORMAT_NC1HWC0, {-1, 16, 256, 256});	156	+ {-1, 16, 256, 256});
151	return BuildDepthwiseGraph("depthwise_to_conv2d_no_fusion_unsupported_format", cfg);	157	return BuildDepthwiseGraph("depthwise_to_conv2d_no_fusion_unsupported_format", cfg);
152	- }, GRAPH_NOT_CHANGED, 1},	158	+ },
153	- {"unknown_fmap_channel", []() {	159	+ GRAPH_NOT_CHANGED, 1},
154	- return BuildDepthwiseGraph("depthwise_to_conv2d_no_fusion_unknown_channel",	160	+ {"unknown_fmap_channel",
		161	+ []() {
		162	+ return BuildDepthwiseGraph(
		163	+ "depthwise_to_conv2d_no_fusion_unknown_channel",
155	DepthwiseConv2DConfig::Basic("depthwise_conv2d", DT_FLOAT, FORMAT_NCHW, {-1, -1, 256, 256}));	164	DepthwiseConv2DConfig::Basic("depthwise_conv2d", DT_FLOAT, FORMAT_NCHW, {-1, -1, 256, 256}));
156	- }, GRAPH_NOT_CHANGED, 1},	165	+ },
157	- {"fmap_not_4d", []() {	166	+ GRAPH_NOT_CHANGED, 1},
		167	+ {"fmap_not_4d",
		168	+ []() {
158	DepthwiseConv2DConfig cfg = DepthwiseConv2DConfig::Basic("depthwise_conv2d");	169	DepthwiseConv2DConfig cfg = DepthwiseConv2DConfig::Basic("depthwise_conv2d");
159	cfg.inputs[0].shape = {-1, 16, 256};	170	cfg.inputs[0].shape = {-1, 16, 256};
160	- cfg.inputs[0].tensorDesc =	171	+ cfg.inputs[0].tensorDesc = BuildTensorDesc(DT_FLOAT, FORMAT_NCHW, {-1, 16, 256}, FORMAT_NCHW,
161	- BuildTensorDesc(DT_FLOAT, FORMAT_NCHW, {-1, 16, 256}, FORMAT_NCHW, {-1, 16, 256});	172	+ {-1, 16, 256});
162	return BuildDepthwiseGraph("depthwise_to_conv2d_no_fusion_fmap_not_4d", cfg);	173	return BuildDepthwiseGraph("depthwise_to_conv2d_no_fusion_fmap_not_4d", cfg);
163	- }, GRAPH_NOT_CHANGED, 1},	174	+ },
164	- {"empty_graph_no_match", []() {	175	+ GRAPH_NOT_CHANGED, 1},
		176	+ {"empty_graph_no_match",
		177	+ []() {
165	return TestGraph("depthwise_to_conv2d_no_fusion_empty_graph")	178	return TestGraph("depthwise_to_conv2d_no_fusion_empty_graph")
166	.SetSocAscend950()	179	.SetSocAscend950()
167	.AddRelu(ReluConfig::Basic("relu", DT_FLOAT))	180	.AddRelu(ReluConfig::Basic("relu", DT_FLOAT))
168	.SetOutput("relu")	181	.SetOutput("relu")
169	.Build();	182	.Build();
170	- }, GRAPH_NOT_CHANGED, 0},	183	+ },
		184	+ GRAPH_NOT_CHANGED, 0},
171	};	185	};
172		186
173	- for (const auto &p : points) {	187	+ for (const auto& p : points) {
174	SCOPED_TRACE(p.pointName);	188	SCOPED_TRACE(p.pointName);
175	auto graph = p.build();	189	auto graph = p.build();
176	std::string name = std::string("depthwise_to_conv2d_no_fusion_") + p.pointName;	190	std::string name = std::string("depthwise_to_conv2d_no_fusion_") + p.pointName;
	@@ -178,23 +192,28 @@ TEST_F(DepthwiseToConv2dFusionPassTest, depthwise_to_conv2d_no_fusion) {
178	}	192	}
179	}	193	}
180		194
181	-TEST_F(DepthwiseToConv2dFusionPassTest, depthwise_to_conv2d_fusion_dynamic_non_nd_soc) {	195	+TEST_F(DepthwiseToConv2dFusionPassTest, depthwise_to_conv2d_fusion_dynamic_non_nd_soc)
		196	+{
182	auto graph = BuildDepthwiseGraph("depthwise_to_conv2d_fusion_dynamic_non_nd",	197	auto graph = BuildDepthwiseGraph("depthwise_to_conv2d_fusion_dynamic_non_nd",
183	- DepthwiseConv2DConfig::Basic("depthwise_conv2d"), &kNonNdSocConfig);	198	+ DepthwiseConv2DConfig::Basic("depthwise_conv2d"), &kNonNdSocConfig);
184	TestTotalPass("depthwise_to_conv2d_fusion_dynamic_non_nd", graph, SUCCESS);	199	TestTotalPass("depthwise_to_conv2d_fusion_dynamic_non_nd", graph, SUCCESS);
185	}	200	}
186		201
187	-TEST_F(DepthwiseToConv2dFusionPassTest, depthwise_to_conv2d_fusion_desc_and_structure) {	202	+TEST_F(DepthwiseToConv2dFusionPassTest, depthwise_to_conv2d_fusion_desc_and_structure)
		203	+{
188	struct Point {	204	struct Point {
189	- const char *pointName;	205	+ const char* pointName;
190	std::function<GraphPtr()> build;	206	std::function<GraphPtr()> build;
191	- std::function<void(GraphPtr &)> verify;	207	+ std::function<void(GraphPtr&)> verify;
192	} const points[] = {	208	} const points[] = {
193	- {"desc_naming_and_origin_shape", []() {	209	+ {"desc_naming_and_origin_shape",
		210	+ []() {
194	return BuildDepthwiseGraph("depthwise_to_conv2d_desc_naming",	211	return BuildDepthwiseGraph("depthwise_to_conv2d_desc_naming",
195	- DepthwiseConv2DConfig::Basic("depthwise_conv2d", DT_FLOAT16, FORMAT_NCHW, {1, 16, 240, 352},	212	+ DepthwiseConv2DConfig::Basic("depthwise_conv2d", DT_FLOAT16, FORMAT_NCHW,
196	- {1, 1, 1, 64}, {1, 4, 240, 352}).WithBias(DT_FLOAT));	213	+ {1, 16, 240, 352}, {1, 1, 1, 64}, {1, 4, 240, 352})
197	- }, [](GraphPtr &graph) {	214	+ .WithBias(DT_FLOAT));
		215	+ },
		216	+ [](GraphPtr& graph) {
198	GNode fused;	217	GNode fused;
199	ASSERT_TRUE(GraphChecker::FindFirstNodeByOpType(graph, "Conv2D", fused));	218	ASSERT_TRUE(GraphChecker::FindFirstNodeByOpType(graph, "Conv2D", fused));
200	std::string nodeName;	219	std::string nodeName;
	@@ -216,7 +235,8 @@ TEST_F(DepthwiseToConv2dFusionPassTest, depthwise_to_conv2d_fusion_desc_and_stru
216	EXPECT_EQ(fmapDesc.GetDataType(), DT_FLOAT16);	235	EXPECT_EQ(fmapDesc.GetDataType(), DT_FLOAT16);
217	EXPECT_EQ(outDesc.GetDataType(), DT_FLOAT16);	236	EXPECT_EQ(outDesc.GetDataType(), DT_FLOAT16);
218	}},	237	}},
219	- {"replacement_subgraph_structure", []() {	238	+ {"replacement_subgraph_structure",
		239	+ []() {
220	return TestGraph("depthwise_to_conv2d_replacement_structure")	240	return TestGraph("depthwise_to_conv2d_replacement_structure")
221	.SetSocAscend950()	241	.SetSocAscend950()
222	.AddDepthwiseConv2D(DepthwiseConv2DConfig::Basic("dw"))	242	.AddDepthwiseConv2D(DepthwiseConv2DConfig::Basic("dw"))
	@@ -224,7 +244,8 @@ TEST_F(DepthwiseToConv2dFusionPassTest, depthwise_to_conv2d_fusion_desc_and_stru
224	.Connect("dw", 0, "Relu", 0)	244	.Connect("dw", 0, "Relu", 0)
225	.SetOutput("Relu")	245	.SetOutput("Relu")
226	.Build();	246	.Build();
227	- }, [](GraphPtr &graph) {	247	+ },
		248	+ [](GraphPtr& graph) {
228	EXPECT_EQ(GraphChecker::CountNodes(graph, "Conv2D"), 1);	249	EXPECT_EQ(GraphChecker::CountNodes(graph, "Conv2D"), 1);
229	EXPECT_EQ(GraphChecker::CountNodes(graph, "DepthwiseConv2D"), 0);	250	EXPECT_EQ(GraphChecker::CountNodes(graph, "DepthwiseConv2D"), 0);
230	EXPECT_TRUE(GraphChecker::HasNode(graph, "Relu"));	251	EXPECT_TRUE(GraphChecker::HasNode(graph, "Relu"));
	@@ -236,7 +257,7 @@ TEST_F(DepthwiseToConv2dFusionPassTest, depthwise_to_conv2d_fusion_desc_and_stru
236	}},	257	}},
237	};	258	};
238		259
239	- for (const auto &p : points) {	260	+ for (const auto& p : points) {
240	SCOPED_TRACE(p.pointName);	261	SCOPED_TRACE(p.pointName);
241	auto graph = p.build();	262	auto graph = p.build();
242	TestTotalPass(std::string("depthwise_to_conv2d_fusion_desc_and_structure_") + p.pointName, graph, SUCCESS);	263	TestTotalPass(std::string("depthwise_to_conv2d_fusion_desc_and_structure_") + p.pointName, graph, SUCCESS);
	@@ -246,19 +267,22 @@ TEST_F(DepthwiseToConv2dFusionPassTest, depthwise_to_conv2d_fusion_desc_and_stru
246	}	267	}
247	}	268	}
248		269
249	-TEST_F(DepthwiseToConv2dFusionPassTest, depthwise_to_conv2d_replacement_attr_and_desc) {	270	+TEST_F(DepthwiseToConv2dFusionPassTest, depthwise_to_conv2d_replacement_attr_and_desc)
		271	+{
250	struct Point {	272	struct Point {
251	- const char *pointName;	273	+ const char* pointName;
252	std::function<GraphPtr()> build;	274	std::function<GraphPtr()> build;
253	- std::function<void(GNode &)> verify;	275	+ std::function<void(GNode&)> verify;
254	} const points[] = {	276	} const points[] = {
255	- {"attr_strides_pads_dilations_passthrough", []() {	277	+ {"attr_strides_pads_dilations_passthrough",
		278	+ []() {
256	DepthwiseConv2DConfig cfg = DepthwiseConv2DConfig::Basic("depthwise_conv2d");	279	DepthwiseConv2DConfig cfg = DepthwiseConv2DConfig::Basic("depthwise_conv2d");
257	cfg.SetAttr("strides", std::vector<int64_t>{2, 2, 2, 2});	280	cfg.SetAttr("strides", std::vector<int64_t>{2, 2, 2, 2});
258	cfg.SetAttr("pads", std::vector<int64_t>{2, 2, 2, 2});	281	cfg.SetAttr("pads", std::vector<int64_t>{2, 2, 2, 2});
259	cfg.SetAttr("dilations", std::vector<int64_t>{2, 2, 2, 2});	282	cfg.SetAttr("dilations", std::vector<int64_t>{2, 2, 2, 2});
260	return BuildDepthwiseGraph("depthwise_to_conv2d_attr_strides_pads_dilations", cfg);	283	return BuildDepthwiseGraph("depthwise_to_conv2d_attr_strides_pads_dilations", cfg);
261	- }, [](GNode &fused) {	284	+ },
		285	+ [](GNode& fused) {
262	std::vector<int64_t> strides;	286	std::vector<int64_t> strides;
263	std::vector<int64_t> pads;	287	std::vector<int64_t> pads;
264	std::vector<int64_t> dilations;	288	std::vector<int64_t> dilations;
	@@ -269,73 +293,92 @@ TEST_F(DepthwiseToConv2dFusionPassTest, depthwise_to_conv2d_replacement_attr_and
269	EXPECT_EQ(pads, (std::vector<int64_t>{2, 2, 2, 2}));	293	EXPECT_EQ(pads, (std::vector<int64_t>{2, 2, 2, 2}));
270	EXPECT_EQ(dilations, (std::vector<int64_t>{2, 2, 2, 2}));	294	EXPECT_EQ(dilations, (std::vector<int64_t>{2, 2, 2, 2}));
271	}},	295	}},
272	- {"attr_groups_nchw", []() {	296	+ {"attr_groups_nchw",
273	- return BuildDepthwiseGraph("depthwise_to_conv2d_attr_groups_nchw",	297	+ []() {
		298	+ return BuildDepthwiseGraph(
		299	+ "depthwise_to_conv2d_attr_groups_nchw",
274	DepthwiseConv2DConfig::Basic("depthwise_conv2d", DT_FLOAT, FORMAT_NCHW, {-1, 24, 256, 256}));	300	DepthwiseConv2DConfig::Basic("depthwise_conv2d", DT_FLOAT, FORMAT_NCHW, {-1, 24, 256, 256}));
275	- }, [](GNode &fused) {	301	+ },
		302	+ [](GNode& fused) {
276	int64_t groups = 0;	303	int64_t groups = 0;
277	ASSERT_EQ(fused.GetAttr(GROUPS, groups), GRAPH_SUCCESS);	304	ASSERT_EQ(fused.GetAttr(GROUPS, groups), GRAPH_SUCCESS);
278	EXPECT_EQ(groups, int64_t{24});	305	EXPECT_EQ(groups, int64_t{24});
279	}},	306	}},
280	- {"attr_groups_nhwc", []() {	307	+ {"attr_groups_nhwc",
281	- return BuildDepthwiseGraph("depthwise_to_conv2d_attr_groups_nhwc",	308	+ []() {
		309	+ return BuildDepthwiseGraph(
		310	+ "depthwise_to_conv2d_attr_groups_nhwc",
282	DepthwiseConv2DConfig::Basic("depthwise_conv2d", DT_FLOAT, FORMAT_NHWC, {-1, 256, 256, 32},	311	DepthwiseConv2DConfig::Basic("depthwise_conv2d", DT_FLOAT, FORMAT_NHWC, {-1, 256, 256, 32},
283	- {1, 1, 1, 64}, {-1, 256, 256, 8}));	312	+ {1, 1, 1, 64}, {-1, 256, 256, 8}));
284	- }, [](GNode &fused) {	313	+ },
		314	+ [](GNode& fused) {
285	int64_t groups = 0;	315	int64_t groups = 0;
286	ASSERT_EQ(fused.GetAttr(GROUPS, groups), GRAPH_SUCCESS);	316	ASSERT_EQ(fused.GetAttr(GROUPS, groups), GRAPH_SUCCESS);
287	EXPECT_EQ(groups, int64_t{32});	317	EXPECT_EQ(groups, int64_t{32});
288	}},	318	}},
289	- {"attr_data_format_nhwc", []() {	319	+ {"attr_data_format_nhwc",
		320	+ []() {
290	DepthwiseConv2DConfig cfg = DepthwiseConv2DConfig::Basic("depthwise_conv2d", DT_FLOAT, FORMAT_NHWC,	321	DepthwiseConv2DConfig cfg = DepthwiseConv2DConfig::Basic("depthwise_conv2d", DT_FLOAT, FORMAT_NHWC,
291	- {-1, 256, 256, 16});	322	+ {-1, 256, 256, 16});
292	cfg.SetAttr("data_format", std::string("NHWC"));	323	cfg.SetAttr("data_format", std::string("NHWC"));
293	return BuildDepthwiseGraph("depthwise_to_conv2d_attr_data_format_nhwc", cfg);	324	return BuildDepthwiseGraph("depthwise_to_conv2d_attr_data_format_nhwc", cfg);
294	- }, [](GNode &fused) {	325	+ },
		326	+ [](GNode& fused) {
295	AscendString dataFormat;	327	AscendString dataFormat;
296	ASSERT_EQ(fused.GetAttr(DATA_FORMAT, dataFormat), GRAPH_SUCCESS);	328	ASSERT_EQ(fused.GetAttr(DATA_FORMAT, dataFormat), GRAPH_SUCCESS);
297	EXPECT_STREQ(dataFormat.GetString(), "NHWC");	329	EXPECT_STREQ(dataFormat.GetString(), "NHWC");
298	}},	330	}},
299	- {"attr_padding_passthrough", []() {	331	+ {"attr_padding_passthrough",
		332	+ []() {
300	DepthwiseConv2DConfig cfg = DepthwiseConv2DConfig::Basic("depthwise_conv2d");	333	DepthwiseConv2DConfig cfg = DepthwiseConv2DConfig::Basic("depthwise_conv2d");
301	cfg.SetAttr("padding", std::string("SAME"));	334	cfg.SetAttr("padding", std::string("SAME"));
302	return BuildDepthwiseGraph("depthwise_to_conv2d_attr_padding_passthrough", cfg);	335	return BuildDepthwiseGraph("depthwise_to_conv2d_attr_padding_passthrough", cfg);
303	- }, [](GNode &fused) {	336	+ },
		337	+ [](GNode& fused) {
304	AscendString paddingVal;	338	AscendString paddingVal;
305	ASSERT_EQ(fused.GetAttr(PADDING, paddingVal), GRAPH_SUCCESS);	339	ASSERT_EQ(fused.GetAttr(PADDING, paddingVal), GRAPH_SUCCESS);
306	EXPECT_STREQ(paddingVal.GetString(), "SAME");	340	EXPECT_STREQ(paddingVal.GetString(), "SAME");
307	}},	341	}},
308	- {"offset_x_on_nd_soc", []() {	342	+ {"offset_x_on_nd_soc",
		343	+ []() {
309	DepthwiseConv2DConfig cfg = DepthwiseConv2DConfig::Basic("depthwise_conv2d");	344	DepthwiseConv2DConfig cfg = DepthwiseConv2DConfig::Basic("depthwise_conv2d");
310	cfg.SetAttr("offset_x", int64_t{2});	345	cfg.SetAttr("offset_x", int64_t{2});
311	return BuildDepthwiseGraph("depthwise_to_conv2d_offset_x_on_nd_soc", cfg);	346	return BuildDepthwiseGraph("depthwise_to_conv2d_offset_x_on_nd_soc", cfg);
312	- }, [](GNode &fused) {	347	+ },
		348	+ [](GNode& fused) {
313	int64_t offsetX = 0;	349	int64_t offsetX = 0;
314	ASSERT_EQ(fused.GetAttr(OFFSET_X, offsetX), GRAPH_SUCCESS);	350	ASSERT_EQ(fused.GetAttr(OFFSET_X, offsetX), GRAPH_SUCCESS);
315	EXPECT_EQ(offsetX, int64_t{2});	351	EXPECT_EQ(offsetX, int64_t{2});
316	}},	352	}},
317	- {"offset_x_not_set_on_non_nd_soc", []() {	353	+ {"offset_x_not_set_on_non_nd_soc",
		354	+ []() {
318	DepthwiseConv2DConfig cfg = DepthwiseConv2DConfig::Basic("depthwise_conv2d");	355	DepthwiseConv2DConfig cfg = DepthwiseConv2DConfig::Basic("depthwise_conv2d");
319	cfg.SetAttr("offset_x", int64_t{2});	356	cfg.SetAttr("offset_x", int64_t{2});
320	return BuildDepthwiseGraph("depthwise_to_conv2d_offset_x_not_set_on_non_nd_soc", cfg, &kNonNdSocConfig);	357	return BuildDepthwiseGraph("depthwise_to_conv2d_offset_x_not_set_on_non_nd_soc", cfg, &kNonNdSocConfig);
321	- }, [](GNode &fused) {	358	+ },
		359	+ [](GNode& fused) {
322	int64_t offsetX = 0;	360	int64_t offsetX = 0;
323	EXPECT_NE(fused.GetAttr(OFFSET_X, offsetX), GRAPH_SUCCESS);	361	EXPECT_NE(fused.GetAttr(OFFSET_X, offsetX), GRAPH_SUCCESS);
324	}},	362	}},
325	- {"op_impl_mode_enum_passthrough", []() {	363	+ {"op_impl_mode_enum_passthrough",
		364	+ []() {
326	DepthwiseConv2DConfig cfg = DepthwiseConv2DConfig::Basic("depthwise_conv2d");	365	DepthwiseConv2DConfig cfg = DepthwiseConv2DConfig::Basic("depthwise_conv2d");
327	cfg.SetAttr("_op_impl_mode_enum", int64_t{0x40});	366	cfg.SetAttr("_op_impl_mode_enum", int64_t{0x40});
328	return BuildDepthwiseGraph("depthwise_to_conv2d_op_impl_mode_enum", cfg);	367	return BuildDepthwiseGraph("depthwise_to_conv2d_op_impl_mode_enum", cfg);
329	- }, [](GNode &fused) {	368	+ },
		369	+ [](GNode& fused) {
330	int64_t implMode = 0;	370	int64_t implMode = 0;
331	ASSERT_EQ(fused.GetAttr(OP_IMPL_MODE_ENUM, implMode), GRAPH_SUCCESS);	371	ASSERT_EQ(fused.GetAttr(OP_IMPL_MODE_ENUM, implMode), GRAPH_SUCCESS);
332	EXPECT_EQ(implMode, int64_t{0x40});	372	EXPECT_EQ(implMode, int64_t{0x40});
333	}},	373	}},
334	- {"desc_input_output_preserved", []() {	374	+ {"desc_input_output_preserved",
335	- return BuildDepthwiseGraph("depthwise_to_conv2d_desc_input_output_preserved",	375	+ []() {
		376	+ return BuildDepthwiseGraph(
		377	+ "depthwise_to_conv2d_desc_input_output_preserved",
336	DepthwiseConv2DConfig::Basic("depthwise_conv2d", DT_FLOAT16, FORMAT_NCHW, {1, 16, 240, 352},	378	DepthwiseConv2DConfig::Basic("depthwise_conv2d", DT_FLOAT16, FORMAT_NCHW, {1, 16, 240, 352},
337	- {1, 1, 1, 64}, {1, 4, 240, 352}));	379	+ {1, 1, 1, 64}, {1, 4, 240, 352}));
338	- }, [](GNode &fused) {	380	+ },
		381	+ [](GNode& fused) {
339	TensorDesc fmapDesc;	382	TensorDesc fmapDesc;
340	TensorDesc filterDesc;	383	TensorDesc filterDesc;
341	TensorDesc outDesc;	384	TensorDesc outDesc;
	@@ -348,16 +391,18 @@ TEST_F(DepthwiseToConv2dFusionPassTest, depthwise_to_conv2d_replacement_attr_and
348	EXPECT_EQ(fmapDesc.GetFormat(), FORMAT_NCHW);	391	EXPECT_EQ(fmapDesc.GetFormat(), FORMAT_NCHW);
349	EXPECT_EQ(outDesc.GetFormat(), FORMAT_NCHW);	392	EXPECT_EQ(outDesc.GetFormat(), FORMAT_NCHW);
350	}},	393	}},
351	- {"bias_input_absent", []() {	394	+ {"bias_input_absent",
		395	+ []() {
352	return BuildDepthwiseGraph("depthwise_to_conv2d_bias_input_absent",	396	return BuildDepthwiseGraph("depthwise_to_conv2d_bias_input_absent",
353	- DepthwiseConv2DConfig::Basic("depthwise_conv2d"));	397	+ DepthwiseConv2DConfig::Basic("depthwise_conv2d"));
354	- }, [](GNode &fused) {	398	+ },
355	- EXPECT_EQ(fused.GetInputsSize(), size_t{2});	399	+ [](GNode& fused) { EXPECT_EQ(fused.GetInputsSize(), size_t{2}); }},
356	- }},	400	+ {"bias_input_present",
357	- {"bias_input_present", []() {	401	+ []() {
358	return BuildDepthwiseGraph("depthwise_to_conv2d_bias_input_present",	402	return BuildDepthwiseGraph("depthwise_to_conv2d_bias_input_present",
359	- DepthwiseConv2DConfig::Basic("depthwise_conv2d").WithBias(DT_FLOAT));	403	+ DepthwiseConv2DConfig::Basic("depthwise_conv2d").WithBias(DT_FLOAT));
360	- }, [](GNode &fused) {	404	+ },
		405	+ [](GNode& fused) {
361	TensorDesc biasDesc;	406	TensorDesc biasDesc;
362	EXPECT_EQ(fused.GetInputsSize(), size_t{3});	407	EXPECT_EQ(fused.GetInputsSize(), size_t{3});
363	ASSERT_EQ(fused.GetInputDesc(INPUT_BIAS_INDEX, biasDesc), GRAPH_SUCCESS);	408	ASSERT_EQ(fused.GetInputDesc(INPUT_BIAS_INDEX, biasDesc), GRAPH_SUCCESS);
	@@ -366,7 +411,7 @@ TEST_F(DepthwiseToConv2dFusionPassTest, depthwise_to_conv2d_replacement_attr_and
366	}},	411	}},
367	};	412	};
368		413
369	- for (const auto &p : points) {	414	+ for (const auto& p : points) {
370	SCOPED_TRACE(p.pointName);	415	SCOPED_TRACE(p.pointName);
371	auto graph = p.build();	416	auto graph = p.build();
372	TestTotalPass(std::string("depthwise_to_conv2d_replacement_attr_and_desc_") + p.pointName, graph, SUCCESS);	417	TestTotalPass(std::string("depthwise_to_conv2d_replacement_attr_and_desc_") + p.pointName, graph, SUCCESS);
	@@ -377,45 +422,49 @@ TEST_F(DepthwiseToConv2dFusionPassTest, depthwise_to_conv2d_replacement_attr_and
377	}	422	}
378	}	423	}
379		424
380	-TEST_F(DepthwiseToConv2dFusionPassTest, depthwise_to_conv2d_graph_topology) {	425	+TEST_F(DepthwiseToConv2dFusionPassTest, depthwise_to_conv2d_graph_topology)
		426	+{
381	struct Point {	427	struct Point {
382	- const char *pointName;	428	+ const char* pointName;
383	- std::function<void(DepthwiseToConv2dFusionPassTest &)> run;	429	+ std::function<void(DepthwiseToConv2dFusionPassTest&)> run;
384	} const points[] = {	430	} const points[] = {
385	- {"multi_depthwise_same_graph", [](DepthwiseToConv2dFusionPassTest &self) {	431	+ {"multi_depthwise_same_graph",
		432	+ [](DepthwiseToConv2dFusionPassTest& self) {
386	auto graph = TestGraph("depthwise_to_conv2d_multi_depthwise_same_graph")	433	auto graph = TestGraph("depthwise_to_conv2d_multi_depthwise_same_graph")
387	- .SetSocAscend950()	434	+ .SetSocAscend950()
388	- .AddDepthwiseConv2D(DepthwiseConv2DConfig::Basic("dw1"))	435	+ .AddDepthwiseConv2D(DepthwiseConv2DConfig::Basic("dw1"))
389	- .AddDepthwiseConv2D(DepthwiseConv2DConfig::Basic("dw2"))	436	+ .AddDepthwiseConv2D(DepthwiseConv2DConfig::Basic("dw2"))
390	- .SetOutput("dw1")	437	+ .SetOutput("dw1")
391	- .SetOutput("dw2")	438	+ .SetOutput("dw2")
392	- .Build();	439	+ .Build();
393	self.TestTotalPass("depthwise_to_conv2d_multi_depthwise_same_graph", graph, SUCCESS);	440	self.TestTotalPass("depthwise_to_conv2d_multi_depthwise_same_graph", graph, SUCCESS);
394	EXPECT_EQ(GraphChecker::CountNodes(graph, "Conv2D"), 2);	441	EXPECT_EQ(GraphChecker::CountNodes(graph, "Conv2D"), 2);
395	EXPECT_EQ(GraphChecker::CountNodes(graph, "DepthwiseConv2D"), 0);	442	EXPECT_EQ(GraphChecker::CountNodes(graph, "DepthwiseConv2D"), 0);
396	}},	443	}},
397	- {"multi_depthwise_serial", [](DepthwiseToConv2dFusionPassTest &self) {	444	+ {"multi_depthwise_serial",
		445	+ [](DepthwiseToConv2dFusionPassTest& self) {
398	auto graph = TestGraph("depthwise_to_conv2d_multi_depthwise_serial")	446	auto graph = TestGraph("depthwise_to_conv2d_multi_depthwise_serial")
399	- .SetSocAscend950()	447	+ .SetSocAscend950()
400	- .AddDepthwiseConv2D(DepthwiseConv2DConfig::Basic("dw1"))	448	+ .AddDepthwiseConv2D(DepthwiseConv2DConfig::Basic("dw1"))
401	- .AddDepthwiseConv2D(DepthwiseConv2DConfig::Basic("dw2", DT_FLOAT, FORMAT_NCHW,	449	+ .AddDepthwiseConv2D(DepthwiseConv2DConfig::Basic(
402	- {1, 4, 256, 256}, {1, 1, 1, 64}, {1, 2, 256, 256}))	450	+ "dw2", DT_FLOAT, FORMAT_NCHW, {1, 4, 256, 256}, {1, 1, 1, 64}, {1, 2, 256, 256}))
403	- .Connect("dw1", 0, "dw2", 0)	451	+ .Connect("dw1", 0, "dw2", 0)
404	- .SetOutput("dw2")	452	+ .SetOutput("dw2")
405	- .Build();	453	+ .Build();
406	self.TestTotalPass("depthwise_to_conv2d_multi_depthwise_serial", graph, SUCCESS);	454	self.TestTotalPass("depthwise_to_conv2d_multi_depthwise_serial", graph, SUCCESS);
407	EXPECT_EQ(GraphChecker::CountNodes(graph, "Conv2D"), 2);	455	EXPECT_EQ(GraphChecker::CountNodes(graph, "Conv2D"), 2);
408	EXPECT_EQ(GraphChecker::CountNodes(graph, "DepthwiseConv2D"), 0);	456	EXPECT_EQ(GraphChecker::CountNodes(graph, "DepthwiseConv2D"), 0);
409	}},	457	}},
410	- {"multi_depthwise_mixed_eligibility", [](DepthwiseToConv2dFusionPassTest &) {	458	+ {"multi_depthwise_mixed_eligibility",
		459	+ [](DepthwiseToConv2dFusionPassTest&) {
411	auto graph = TestGraph("depthwise_to_conv2d_multi_depthwise_mixed_eligibility")	460	auto graph = TestGraph("depthwise_to_conv2d_multi_depthwise_mixed_eligibility")
412	- .SetSocAscend950()	461	+ .SetSocAscend950()
413	- .AddDepthwiseConv2D(DepthwiseConv2DConfig::Basic("dw_ok"))	462	+ .AddDepthwiseConv2D(DepthwiseConv2DConfig::Basic("dw_ok"))
414	- .AddDepthwiseConv2D(DepthwiseConv2DConfig::Basic("dw_bad", DT_FLOAT, FORMAT_NCHW,	463	+ .AddDepthwiseConv2D(
415	- {-1, -1, 256, 256}))	464	+ DepthwiseConv2DConfig::Basic("dw_bad", DT_FLOAT, FORMAT_NCHW, {-1, -1, 256, 256}))
416	- .SetOutput("dw_ok")	465	+ .SetOutput("dw_ok")
417	- .SetOutput("dw_bad")	466	+ .SetOutput("dw_bad")
418	- .Build();	467	+ .Build();
419	CustomPassContext passContext;	468	CustomPassContext passContext;
420	passContext.SetPassName("depthwise_to_conv2d_multi_depthwise_mixed_eligibility");	469	passContext.SetPassName("depthwise_to_conv2d_multi_depthwise_mixed_eligibility");
421	DepthwiseToConv2dFusionPass pass({DEPTHWISE_CONV2D});	470	DepthwiseToConv2dFusionPass pass({DEPTHWISE_CONV2D});
	@@ -423,29 +472,30 @@ TEST_F(DepthwiseToConv2dFusionPassTest, depthwise_to_conv2d_graph_topology) {
423	EXPECT_EQ(GraphChecker::CountNodes(graph, "Conv2D"), 1);	472	EXPECT_EQ(GraphChecker::CountNodes(graph, "Conv2D"), 1);
424	EXPECT_EQ(GraphChecker::CountNodes(graph, "DepthwiseConv2D"), 1);	473	EXPECT_EQ(GraphChecker::CountNodes(graph, "DepthwiseConv2D"), 1);
425	}},	474	}},
426	- {"downstream_consumer_preserved", [](DepthwiseToConv2dFusionPassTest &self) {	475	+ {"downstream_consumer_preserved",
		476	+ [](DepthwiseToConv2dFusionPassTest& self) {
427	auto graph = TestGraph("depthwise_to_conv2d_downstream_consumer_preserved")	477	auto graph = TestGraph("depthwise_to_conv2d_downstream_consumer_preserved")
428	- .SetSocAscend950()	478	+ .SetSocAscend950()
429	- .AddDepthwiseConv2D(DepthwiseConv2DConfig::Basic("dw"))	479	+ .AddDepthwiseConv2D(DepthwiseConv2DConfig::Basic("dw"))
430	- .AddRelu(ReluConfig::Basic("Relu", DT_FLOAT))	480	+ .AddRelu(ReluConfig::Basic("Relu", DT_FLOAT))
431	- .Connect("dw", 0, "Relu", 0)	481	+ .Connect("dw", 0, "Relu", 0)
432	- .SetOutput("Relu")	482	+ .SetOutput("Relu")
433	- .Build();	483	+ .Build();
434	self.TestTotalPass("depthwise_to_conv2d_downstream_consumer_preserved", graph, SUCCESS);	484	self.TestTotalPass("depthwise_to_conv2d_downstream_consumer_preserved", graph, SUCCESS);
435	EXPECT_TRUE(GraphChecker::HasNode(graph, "Relu"));	485	EXPECT_TRUE(GraphChecker::HasNode(graph, "Relu"));
436	EXPECT_TRUE(GraphChecker::HasNode(graph, "Conv2D"));	486	EXPECT_TRUE(GraphChecker::HasNode(graph, "Conv2D"));
437	}},	487	}},
438	};	488	};
439		489
440	- for (const auto &p : points) {	490	+ for (const auto& p : points) {
441	SCOPED_TRACE(p.pointName);	491	SCOPED_TRACE(p.pointName);
442	p.run(*this);	492	p.run(*this);
443	}	493	}
444	}	494	}
445		495
446	-TEST_F(DepthwiseToConv2dFusionPassTest, depthwise_to_conv2d_reentrant) {	496	+TEST_F(DepthwiseToConv2dFusionPassTest, depthwise_to_conv2d_reentrant)
447	- auto graph = BuildDepthwiseGraph("depthwise_to_conv2d_reentrant",	497	+{
448	- DepthwiseConv2DConfig::Basic("depthwise_conv2d"));	498	+ auto graph = BuildDepthwiseGraph("depthwise_to_conv2d_reentrant", DepthwiseConv2DConfig::Basic("depthwise_conv2d"));
449		499
450	TestTotalPass("depthwise_to_conv2d_reentrant_1", graph, SUCCESS);	500	TestTotalPass("depthwise_to_conv2d_reentrant_1", graph, SUCCESS);
451		501

Mconv/extend_conv2d/op_graph/fusion_pass/conv2d_postcube_to_extendconv2d_fusion_pass.cpp+3-5

Mconv/extend_conv2d/op_graph/fusion_pass/conv2d_postcube_to_extendconv2d_fusion_pass.h+12-3

Mconv/extend_conv2d/op_host/op_tiling/extend_conv2d_mdc_tiling.cpp+1-1

	@@ -23,4 +23,4 @@ using namespace optiling::conv_ops_tiling;
23	namespace optiling {	23	namespace optiling {
24	// using op_tiling register capability in "tiling_templates_registry" for AscendC extendconv2d operator	24	// using op_tiling register capability in "tiling_templates_registry" for AscendC extendconv2d operator
25	CONV_REGISTER_TILING_TEMPLATE(ExtendConv2D, Conv2dBaseTiling, static_cast<int32_t>(NpuArch::DAV_5102), 1);	25	CONV_REGISTER_TILING_TEMPLATE(ExtendConv2D, Conv2dBaseTiling, static_cast<int32_t>(NpuArch::DAV_5102), 1);
26	-} // namespace optiling	26	+} // namespace optiling

Mconv/extend_conv2d/tests/ut/op_graph/test_conv2d_postcube_to_extendconv2d_fusion_pass.cpp+1-1

	@@ -41,7 +41,7 @@ protected:
41	GNode convNode = testGraphBuilder.GetNode(convNodeName);	41	GNode convNode = testGraphBuilder.GetNode(convNodeName);
42	pass.InitMember();	42	pass.InitMember();
43	ConvFusionUtilsPass::GetConvDescInfo(convNode, pass.convDescInfo);	43	ConvFusionUtilsPass::GetConvDescInfo(convNode, pass.convDescInfo);
44	- ConvFusionUtilsPass::CheckSocList(Conv2DPostCubeToExtendConv2DFusion::SUPPORT_SOC_LIST, pass.npuArch);	44	+ ConvFusionUtilsPass::CheckSocList(Conv2DPostCubeToExtendConv2DFusion::SUPPORT_SOC_LIST, pass.npuArch, true);
45	}	45	}
46		46
47	void TestTotalPass(const std::string& passName, GraphPtr& graph, Status expectRes)	47	void TestTotalPass(const std::string& passName, GraphPtr& graph, Status expectRes)

	@@ -323,7 +323,7 @@ public:
323	if constexpr (Intf::isExtendConv2d) {	323	if constexpr (Intf::isExtendConv2d) {
324	if constexpr (FixpipeIdx == 0) {	324	if constexpr (FixpipeIdx == 0) {
325	intriParams.reluEn = self_->ctx.convTilingData->reluMode0 != 0;	325	intriParams.reluEn = self_->ctx.convTilingData->reluMode0 != 0;
326	-#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 5102)	326	+#if defined(__DAV_35_FAMILY__)
327	intriParams.preReluMode = static_cast<ReluMode>(self_->ctx.convTilingData->reluMode0);	327	intriParams.preReluMode = static_cast<ReluMode>(self_->ctx.convTilingData->reluMode0);
328	if (self_->ctx.convTilingData->reluMode0 == static_cast<uint8_t>(ReluMode::SCALAR_RELU)) {	328	if (self_->ctx.convTilingData->reluMode0 == static_cast<uint8_t>(ReluMode::SCALAR_RELU)) {
329	intriParams.reluScalar = self_->ctx.preReluScalar0;	329	intriParams.reluScalar = self_->ctx.preReluScalar0;
	@@ -334,7 +334,7 @@ public:
334	intriParams.deqScalar = self_->ctx.deqScalar0;	334	intriParams.deqScalar = self_->ctx.deqScalar0;
335	} else {	335	} else {
336	intriParams.reluEn = self_->ctx.convTilingData->reluMode1 != 0;	336	intriParams.reluEn = self_->ctx.convTilingData->reluMode1 != 0;
337	-#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 5102)	337	+#if defined(__DAV_35_FAMILY__)
338	intriParams.preReluMode = static_cast<ReluMode>(self_->ctx.convTilingData->reluMode1);	338	intriParams.preReluMode = static_cast<ReluMode>(self_->ctx.convTilingData->reluMode1);
339	if (self_->ctx.convTilingData->reluMode1 == static_cast<uint8_t>(ReluMode::SCALAR_RELU)) {	339	if (self_->ctx.convTilingData->reluMode1 == static_cast<uint8_t>(ReluMode::SCALAR_RELU)) {
340	intriParams.reluScalar = self_->ctx.preReluScalar1;	340	intriParams.reluScalar = self_->ctx.preReluScalar1;
	@@ -447,7 +447,7 @@ public:
447	}	447	}
448		448
449	FixpipeParamsC310<config.format> intriParams;	449	FixpipeParamsC310<config.format> intriParams;
450	-#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 5102)	450	+#if defined(__DAV_35_FAMILY__)
451	if constexpr (Intf::isFixedPoint) {	451	if constexpr (Intf::isFixedPoint) {
452	intriParams.fixShiftVal = FIX_SHIFT_VAL_LEN_A16W16 - self_->ctx.convTilingData->fixedShiftValue;	452	intriParams.fixShiftVal = FIX_SHIFT_VAL_LEN_A16W16 - self_->ctx.convTilingData->fixedShiftValue;
453	}	453	}

	@@ -187,7 +187,7 @@ public:
187	CopyOutInnerBatch<TensorTypeT, config.format, config>(output, offset, ubInfo);	187	CopyOutInnerBatch<TensorTypeT, config.format, config>(output, offset, ubInfo);
188	} else {	188	} else {
189	FixpipeParamsC310<config.format> intriParams;	189	FixpipeParamsC310<config.format> intriParams;
190	-#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 5102)	190	+#if defined(__DAV_35_FAMILY__)
191	if constexpr (Intf::isFixedPoint) {	191	if constexpr (Intf::isFixedPoint) {
192	intriParams.fixShiftVal = FIX_SHIFT_VAL_LEN_A16W16 - self_->ctx.convTilingData->fixedShiftValue;	192	intriParams.fixShiftVal = FIX_SHIFT_VAL_LEN_A16W16 - self_->ctx.convTilingData->fixedShiftValue;
193	}	193	}
	@@ -224,7 +224,7 @@ private:
224	CopyUbInfo* ubInfo = nullptr)	224	CopyUbInfo* ubInfo = nullptr)
225	{	225	{
226	FixpipeParamsC310<format> intriParams;	226	FixpipeParamsC310<format> intriParams;
227	-#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 5102)	227	+#if defined(__DAV_35_FAMILY__)
228	if constexpr (Intf::isFixedPoint) {	228	if constexpr (Intf::isFixedPoint) {
229	intriParams.fixShiftVal = FIX_SHIFT_VAL_LEN_A16W16 - self_->ctx.convTilingData->fixedShiftValue;	229	intriParams.fixShiftVal = FIX_SHIFT_VAL_LEN_A16W16 - self_->ctx.convTilingData->fixedShiftValue;
230	}	230	}
	@@ -364,7 +364,7 @@ private:
364	if constexpr (Intf::isExtendConv2d) {	364	if constexpr (Intf::isExtendConv2d) {
365	if constexpr (FixpipeIdx == 0) {	365	if constexpr (FixpipeIdx == 0) {
366	intriParams.reluEn = self_->ctx.convTilingData->reluMode0 != 0;	366	intriParams.reluEn = self_->ctx.convTilingData->reluMode0 != 0;
367	-#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 5102)	367	+#if defined(__DAV_35_FAMILY__)
368	intriParams.preReluMode = static_cast<ReluMode>(self_->ctx.convTilingData->reluMode0);	368	intriParams.preReluMode = static_cast<ReluMode>(self_->ctx.convTilingData->reluMode0);
369	if (self_->ctx.convTilingData->reluMode0 == static_cast<uint8_t>(ReluMode::SCALAR_RELU)) {	369	if (self_->ctx.convTilingData->reluMode0 == static_cast<uint8_t>(ReluMode::SCALAR_RELU)) {
370	intriParams.reluScalar = self_->ctx.preReluScalar0;	370	intriParams.reluScalar = self_->ctx.preReluScalar0;
	@@ -375,7 +375,7 @@ private:
375	intriParams.deqScalar = self_->ctx.deqScalar0;	375	intriParams.deqScalar = self_->ctx.deqScalar0;
376	} else {	376	} else {
377	intriParams.reluEn = self_->ctx.convTilingData->reluMode1 != 0;	377	intriParams.reluEn = self_->ctx.convTilingData->reluMode1 != 0;
378	-#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 5102)	378	+#if defined(__DAV_35_FAMILY__)
379	intriParams.preReluMode = static_cast<ReluMode>(self_->ctx.convTilingData->reluMode1);	379	intriParams.preReluMode = static_cast<ReluMode>(self_->ctx.convTilingData->reluMode1);
380	if (self_->ctx.convTilingData->reluMode1 == static_cast<uint8_t>(ReluMode::SCALAR_RELU)) {	380	if (self_->ctx.convTilingData->reluMode1 == static_cast<uint8_t>(ReluMode::SCALAR_RELU)) {
381	intriParams.reluScalar = self_->ctx.preReluScalar1;	381	intriParams.reluScalar = self_->ctx.preReluScalar1;

	@@ -32,8 +32,7 @@ ge::graphStatus Conv2dBaseTiling::CheckStrideLegal()
32	oriShapeAttrInfo_.oriStrideC = stridePtr->GetData()[conv2dOriginFormatAixsPosInfo_.cIndex];	32	oriShapeAttrInfo_.oriStrideC = stridePtr->GetData()[conv2dOriginFormatAixsPosInfo_.cIndex];
33	oriShapeAttrInfo_.oriStrideH = stridePtr->GetData()[conv2dOriginFormatAixsPosInfo_.hIndex];	33	oriShapeAttrInfo_.oriStrideH = stridePtr->GetData()[conv2dOriginFormatAixsPosInfo_.hIndex];
34	oriShapeAttrInfo_.oriStrideW = stridePtr->GetData()[conv2dOriginFormatAixsPosInfo_.wIndex];	34	oriShapeAttrInfo_.oriStrideW = stridePtr->GetData()[conv2dOriginFormatAixsPosInfo_.wIndex];
35	- uint64_t maxStrideHW = (apiInputPlatformInfo.npuArch == NpuArch::DAV_5102) ? LOAD3D_MAX_STRIDE_H_W :	35	+ uint64_t maxStrideHW = (apiInputPlatformInfo.isCubeVectorFuse) ? LOAD3D_MAX_STRIDE_H_W : MAX_ATTRS_SHAPE;
36	- MAX_ATTRS_SHAPE;
37	if (oriShapeAttrInfo_.oriStrideH <= 0 \|\| oriShapeAttrInfo_.oriStrideW <= 0 \|\|	36	if (oriShapeAttrInfo_.oriStrideH <= 0 \|\| oriShapeAttrInfo_.oriStrideW <= 0 \|\|
38	static_cast<uint64_t>(oriShapeAttrInfo_.oriStrideH) > maxStrideHW \|\|	37	static_cast<uint64_t>(oriShapeAttrInfo_.oriStrideH) > maxStrideHW \|\|
39	static_cast<uint64_t>(oriShapeAttrInfo_.oriStrideW) > maxStrideHW) {	38	static_cast<uint64_t>(oriShapeAttrInfo_.oriStrideW) > maxStrideHW) {
	@@ -73,8 +72,7 @@ ge::graphStatus Conv2dBaseTiling::CheckDilationLegal()
73	oriShapeAttrInfo_.oriDilationC = dilationPtr->GetData()[conv2dOriginFormatAixsPosInfo_.cIndex];	72	oriShapeAttrInfo_.oriDilationC = dilationPtr->GetData()[conv2dOriginFormatAixsPosInfo_.cIndex];
74	oriShapeAttrInfo_.oriDilationH = dilationPtr->GetData()[conv2dOriginFormatAixsPosInfo_.hIndex];	73	oriShapeAttrInfo_.oriDilationH = dilationPtr->GetData()[conv2dOriginFormatAixsPosInfo_.hIndex];
75	oriShapeAttrInfo_.oriDilationW = dilationPtr->GetData()[conv2dOriginFormatAixsPosInfo_.wIndex];	74	oriShapeAttrInfo_.oriDilationW = dilationPtr->GetData()[conv2dOriginFormatAixsPosInfo_.wIndex];
76	- uint64_t maxDilationHW = (apiInputPlatformInfo.npuArch == NpuArch::DAV_5102) ? LOAD3D_MAX_DILATION_H_W :	75	+ uint64_t maxDilationHW = (apiInputPlatformInfo.isCubeVectorFuse) ? LOAD3D_MAX_DILATION_H_W : MAX_ATTRS_SHAPE;
77	- MAX_ATTRS_SHAPE;
78	if (oriShapeAttrInfo_.oriDilationH <= 0 \|\| oriShapeAttrInfo_.oriDilationW <= 0 \|\|	76	if (oriShapeAttrInfo_.oriDilationH <= 0 \|\| oriShapeAttrInfo_.oriDilationW <= 0 \|\|
79	static_cast<uint64_t>(oriShapeAttrInfo_.oriDilationH) > maxDilationHW \|\|	77	static_cast<uint64_t>(oriShapeAttrInfo_.oriDilationH) > maxDilationHW \|\|
80	static_cast<uint64_t>(oriShapeAttrInfo_.oriDilationW) > maxDilationHW) {	78	static_cast<uint64_t>(oriShapeAttrInfo_.oriDilationW) > maxDilationHW) {
	@@ -115,7 +113,7 @@ ge::graphStatus Conv2dBaseTiling::CheckPadLegal()
115	OP_LOGE_IF(!UpdateOriPadFromPadMode(), ge::GRAPH_FAILED, context_->GetNodeName(),	113	OP_LOGE_IF(!UpdateOriPadFromPadMode(), ge::GRAPH_FAILED, context_->GetNodeName(),
116	"%s AscendC: UpdateOriPadFromPadMode Failed.", paramInfo_.nodeType.c_str());	114	"%s AscendC: UpdateOriPadFromPadMode Failed.", paramInfo_.nodeType.c_str());
117		115
118	- uint64_t maxPad = (apiInputPlatformInfo.npuArch == NpuArch::DAV_5102) ? LOAD3D_MAX_PAD : MAX_ATTRS_SHAPE;	116	+ uint64_t maxPad = (apiInputPlatformInfo.isCubeVectorFuse) ? LOAD3D_MAX_PAD : MAX_ATTRS_SHAPE;
119	if (oriShapeAttrInfo_.oriPadTop < 0 \|\| oriShapeAttrInfo_.oriPadBottom < 0 \|\| oriShapeAttrInfo_.oriPadLeft < 0 \|\|	117	if (oriShapeAttrInfo_.oriPadTop < 0 \|\| oriShapeAttrInfo_.oriPadBottom < 0 \|\| oriShapeAttrInfo_.oriPadLeft < 0 \|\|
120	oriShapeAttrInfo_.oriPadRight < 0 \|\| static_cast<uint64_t>(oriShapeAttrInfo_.oriPadTop) > maxPad \|\|	118	oriShapeAttrInfo_.oriPadRight < 0 \|\| static_cast<uint64_t>(oriShapeAttrInfo_.oriPadTop) > maxPad \|\|
121	static_cast<uint64_t>(oriShapeAttrInfo_.oriPadBottom) > maxPad \|\|	119	static_cast<uint64_t>(oriShapeAttrInfo_.oriPadBottom) > maxPad \|\|
	@@ -478,7 +476,7 @@ ge::graphStatus Conv2dBaseTiling::CheckExtendDtypeLegal()
478		476
479	ge::graphStatus Conv2dBaseTiling::CheckFixedShiftValueLegal()	477	ge::graphStatus Conv2dBaseTiling::CheckFixedShiftValueLegal()
480	{	478	{
481	- if (!IsMdcSoc(opInfo_->npuArch) \|\| descInfo_.fMapDtype != ge::DataType::DT_FLOAT16) {	479	+ if (!opInfo_->isCubeVectorFuse \|\| descInfo_.fMapDtype != ge::DataType::DT_FLOAT16) {
482	return ge::GRAPH_SUCCESS;	480	return ge::GRAPH_SUCCESS;
483	}	481	}
484		482

	@@ -39,7 +39,7 @@ void Conv2DPostCubeToExtendConv2DFusionPass::InitMember()
39		39
40	bool Conv2DPostCubeToExtendConv2DFusionPass::MeetRequirements(const GNode& convNode)	40	bool Conv2DPostCubeToExtendConv2DFusionPass::MeetRequirements(const GNode& convNode)
41	{	41	{
42	- FUSION_PASS_CHECK(!ConvFusionUtilsPass::CheckSocList(SUPPORT_SOC_LIST, npuArch),	42	+ FUSION_PASS_CHECK(!ConvFusionUtilsPass::CheckSocList(SUPPORT_SOC_LIST, npuArch, true),
43	OP_LOGD(FUSION_NAME, "Current soc not supported, no fusion."), return false);	43	OP_LOGD(FUSION_NAME, "Current soc not supported, no fusion."), return false);
44		44
45	OP_LOGD(convDescInfo.nodeNameStr, "Begin to do Conv2DPostCubeToExtendConv2DFusionPass.");	45	OP_LOGD(convDescInfo.nodeNameStr, "Begin to do Conv2DPostCubeToExtendConv2DFusionPass.");
	@@ -58,8 +58,7 @@ bool Conv2DPostCubeToExtendConv2DFusionPass::MeetRequirements(const GNode& convN
58		58
59	// Check cur node's formats whether it is supported.	59	// Check cur node's formats whether it is supported.
60	std::vector<Format> convFormats = {convDescInfo.fmapFormat, convDescInfo.filterFormat, convDescInfo.outputFormat};	60	std::vector<Format> convFormats = {convDescInfo.fmapFormat, convDescInfo.filterFormat, convDescInfo.outputFormat};
61	- auto convSupportFormats = npuArch == NpuArch::DAV_5102 ? CONV_SUPPORT_FORMATS_DAV_5102 :	61	+ auto convSupportFormats = CONV_SUPPORT_FORMATS_MAP.at(ConvFusionUtilsPass::GetArchKey());
62	- CONV_SUPPORT_FORMATS_DAV_3510;
63	FUSION_PASS_CHECK(!ConvFusionUtilsPass::CheckSupportList<Format>(convSupportFormats, convFormats),	62	FUSION_PASS_CHECK(!ConvFusionUtilsPass::CheckSupportList<Format>(convSupportFormats, convFormats),
64	OP_LOGD(convDescInfo.nodeNameStr, "Conv2D format not supported, no fusion."), return false);	63	OP_LOGD(convDescInfo.nodeNameStr, "Conv2D format not supported, no fusion."), return false);
65		64
	@@ -275,8 +274,7 @@ bool Conv2DPostCubeToExtendConv2DFusionPass::CheckConvPostCubeDtype(const GNodeP
275	TypeUtils::DataTypeToSerialString(postCubeInDtype).c_str(),	274	TypeUtils::DataTypeToSerialString(postCubeInDtype).c_str(),
276	TypeUtils::DataTypeToSerialString(postCubeOutDtype).c_str());	275	TypeUtils::DataTypeToSerialString(postCubeOutDtype).c_str());
277		276
278	- auto supportedDtypes = npuArch == NpuArch::DAV_5102 ? SUPPORTED_DTYPES_WITH_POST_CUBE_DAV_5102 :	277	+ auto supportedDtypes = SUPPORTED_DTYPES_WITH_POST_CUBE_MAP.at(ConvFusionUtilsPass::GetArchKey());
279	- SUPPORTED_DTYPES_WITH_POST_CUBE_DAV_3510;
280	if (!ConvFusionUtilsPass::CheckSupportList<DataType>(supportedDtypes, checkDtypes)) {	278	if (!ConvFusionUtilsPass::CheckSupportList<DataType>(supportedDtypes, checkDtypes)) {
281	std::string incorrectDtypes = VectorToString(checkDtypes);	279	std::string incorrectDtypes = VectorToString(checkDtypes);
282	std::string reason = "The dtypes of these parameters support only the following combinations: " +	280	std::string reason = "The dtypes of these parameters support only the following combinations: " +

	@@ -39,7 +39,7 @@ const ge::AscendString SCALE_1 = "scale1";
39	const ge::AscendString RELU_WEIGHT_0 = "relu_weight0";	39	const ge::AscendString RELU_WEIGHT_0 = "relu_weight0";
40	const ge::AscendString RELU_WEIGHT_1 = "relu_weight1";	40	const ge::AscendString RELU_WEIGHT_1 = "relu_weight1";
41		41
42	-const std::map<std::string, NpuArch> SUPPORT_SOC_LIST = {{"Ascend950", NpuArch::DAV_3510}, {"MC62", NpuArch::DAV_5102}};	42	+const std::map<std::string, NpuArch> SUPPORT_SOC_LIST = {{"Ascend950", NpuArch::DAV_3510}};
43	const std::string FUSION_NAME = "Conv2DPostCubeToExtendConv2DFusionPass";	43	const std::string FUSION_NAME = "Conv2DPostCubeToExtendConv2DFusionPass";
44		44
45	constexpr int32_t EXTENDCONV2D_QUANT_SCALE_0_INDEX = 4;	45	constexpr int32_t EXTENDCONV2D_QUANT_SCALE_0_INDEX = 4;
	@@ -64,7 +64,7 @@ const std::vector<std::vector<ge::Format>> CONV_SUPPORT_FORMATS_DAV_3510 = {
64	{ge::FORMAT_NCHW, ge::FORMAT_NCHW, ge::FORMAT_NCHW}, {ge::FORMAT_NHWC, ge::FORMAT_HWCN, ge::FORMAT_NHWC}};	64	{ge::FORMAT_NCHW, ge::FORMAT_NCHW, ge::FORMAT_NCHW}, {ge::FORMAT_NHWC, ge::FORMAT_HWCN, ge::FORMAT_NHWC}};
65		65
66	// Fmap Filter Output	66	// Fmap Filter Output
67	-const std::vector<std::vector<ge::Format>> CONV_SUPPORT_FORMATS_DAV_5102 = {	67	+const std::vector<std::vector<ge::Format>> CONV_SUPPORT_FORMATS_FUSE = {
68	{ge::FORMAT_NCHW, ge::FORMAT_FRACTAL_Z, ge::FORMAT_NCHW},	68	{ge::FORMAT_NCHW, ge::FORMAT_FRACTAL_Z, ge::FORMAT_NCHW},
69	{ge::FORMAT_NCHW, ge::FORMAT_FRACTAL_Z_C04, ge::FORMAT_NCHW},	69	{ge::FORMAT_NCHW, ge::FORMAT_FRACTAL_Z_C04, ge::FORMAT_NCHW},
70	{ge::FORMAT_NHWC, ge::FORMAT_FRACTAL_Z, ge::FORMAT_NHWC},	70	{ge::FORMAT_NHWC, ge::FORMAT_FRACTAL_Z, ge::FORMAT_NHWC},
	@@ -78,7 +78,7 @@ const std::vector<std::vector<ge::DataType>> SUPPORTED_DTYPES_WITH_POST_CUBE_DAV
78	{ge::DataType::DT_INT8, ge::DataType::DT_INT8, ge::DataType::DT_INT32, ge::DataType::DT_INT8}};	78	{ge::DataType::DT_INT8, ge::DataType::DT_INT8, ge::DataType::DT_INT32, ge::DataType::DT_INT8}};
79		79
80	// Fmap Filter PostCubeIn PostCubeOut	80	// Fmap Filter PostCubeIn PostCubeOut
81	-const std::vector<std::vector<ge::DataType>> SUPPORTED_DTYPES_WITH_POST_CUBE_DAV_5102 = {	81	+const std::vector<std::vector<ge::DataType>> SUPPORTED_DTYPES_WITH_POST_CUBE_FUSE = {
82	{ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16},	82	{ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16},
83	{ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16, ge::DataType::DT_INT8},	83	{ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16, ge::DataType::DT_FLOAT16, ge::DataType::DT_INT8},
84	{ge::DataType::DT_INT8, ge::DataType::DT_INT8, ge::DataType::DT_INT32, ge::DataType::DT_FLOAT16},	84	{ge::DataType::DT_INT8, ge::DataType::DT_INT8, ge::DataType::DT_INT32, ge::DataType::DT_FLOAT16},
	@@ -86,6 +86,15 @@ const std::vector<std::vector<ge::DataType>> SUPPORTED_DTYPES_WITH_POST_CUBE_DAV
86	{ge::DataType::DT_FLOAT16, ge::DataType::DT_INT8, ge::DataType::DT_INT32, ge::DataType::DT_FLOAT16},	86	{ge::DataType::DT_FLOAT16, ge::DataType::DT_INT8, ge::DataType::DT_INT32, ge::DataType::DT_FLOAT16},
87	{ge::DataType::DT_FLOAT16, ge::DataType::DT_INT8, ge::DataType::DT_INT32, ge::DataType::DT_INT8}};	87	{ge::DataType::DT_FLOAT16, ge::DataType::DT_INT8, ge::DataType::DT_INT32, ge::DataType::DT_INT8}};
88		88
		89	+// arch-keyed support list maps (extend by adding new arch keys)
		90	+const std::map<std::string, std::vector<std::vector<ge::Format>>> CONV_SUPPORT_FORMATS_MAP = {
		91	+ {ConvFusionUtils::NPU_ARCH_KEY_3510, CONV_SUPPORT_FORMATS_DAV_3510},
		92	+ {ConvFusionUtils::NPU_ARCH_KEY_FUSE, CONV_SUPPORT_FORMATS_FUSE}};
		93	+
		94	+const std::map<std::string, std::vector<std::vector<ge::DataType>>> SUPPORTED_DTYPES_WITH_POST_CUBE_MAP = {
		95	+ {ConvFusionUtils::NPU_ARCH_KEY_3510, SUPPORTED_DTYPES_WITH_POST_CUBE_DAV_3510},
		96	+ {ConvFusionUtils::NPU_ARCH_KEY_FUSE, SUPPORTED_DTYPES_WITH_POST_CUBE_FUSE}};
		97	+
89	const std::vector<ge::AscendString> SUPPORTED_NODE_TYPES = {"Conv2D", "AscendDequant", "AscendRequant",	98	const std::vector<ge::AscendString> SUPPORTED_NODE_TYPES = {"Conv2D", "AscendDequant", "AscendRequant",
90	"AscendQuant", "Relu", "LeakyRelu"};	99	"AscendQuant", "Relu", "LeakyRelu"};
91	const std::vector<ge::AscendString> POST_CUBE_NODE_TYPES = {"AscendDequant", "AscendRequant", "AscendQuant", "Relu",	100	const std::vector<ge::AscendString> POST_CUBE_NODE_TYPES = {"AscendDequant", "AscendRequant", "AscendQuant", "Relu",