chunk_cat算子支持950芯片-ops-math-AtomGit

已合并

chunk_cat算子支持950芯片 #3165

wy519创建于 6月7日从wy519/ops-math： chunk_cat_a5 合入到 cann/ops-math： master

chunk_cat算子支持950芯片 #3165

已合并

wy519创建于 6月7日

从wy519/ops-math chunk_cat_a5合入到cann/ops-math master

共 13 个文件变更+1006-359

Mconversion/chunk_cat/README.md+1-1

	@@ -4,7 +4,7 @@
4		4
5	\| 产品 \| 是否支持 \|	5	\| 产品 \| 是否支持 \|
6	\| :----------------------------------------------------------- \| :------: \|	6	\| :----------------------------------------------------------- \| :------: \|
7	-\| Ascend 950PR/Ascend 950DT \| × \|	7	+\| Ascend 950PR/Ascend 950DT \| √ \|
8	\| Atlas A3 训练系列产品/Atlas A3 推理系列产品 \| √ \|	8	\| Atlas A3 训练系列产品/Atlas A3 推理系列产品 \| √ \|
9	\| Atlas A2 训练系列产品/Atlas A2 推理系列产品 \| √ \|	9	\| Atlas A2 训练系列产品/Atlas A2 推理系列产品 \| √ \|
10	\| Atlas 200I/500 A2 推理产品 \| × \|	10	\| Atlas 200I/500 A2 推理产品 \| × \|

Mconversion/chunk_cat/docs/aclnnChunkCat.md+1-1

	@@ -6,7 +6,7 @@
6		6
7	\| 产品 \| 是否支持 \|	7	\| 产品 \| 是否支持 \|
8	\| :------------------------------------------------------- \| :------: \|	8	\| :------------------------------------------------------- \| :------: \|
9	-\| <term>Ascend 950PR/Ascend 950DT</term> \| × \|	9	+\| <term>Ascend 950PR/Ascend 950DT</term> \| √ \|
10	\| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> \| √ \|	10	\| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> \| √ \|
11	\| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> \| √ \|	11	\| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> \| √ \|
12	\| <term>Atlas 200I/500 A2 推理产品</term> \| × \|	12	\| <term>Atlas 200I/500 A2 推理产品</term> \| × \|

Mconversion/chunk_cat/op_api/aclnn_chunk_cat.cpp+6-5

	@@ -44,8 +44,7 @@ static const std::initializer_list<op::DataType> ASCEND910B_DTYPE_SUPPORT_LIST =
44	static const inline std::initializer_list<DataType>& GetSupportDtypeList(NpuArch npuArch)	44	static const inline std::initializer_list<DataType>& GetSupportDtypeList(NpuArch npuArch)
45	{	45	{
46	static const std::initializer_list<DataType> emptyDtypes = {};	46	static const std::initializer_list<DataType> emptyDtypes = {};
47	- if (	47	+ if (npuArch == NpuArch::DAV_2201 \|\| IsRegBase(npuArch)) {
48	- npuArch == NpuArch::DAV_2201) {
49	return ASCEND910B_DTYPE_SUPPORT_LIST;	48	return ASCEND910B_DTYPE_SUPPORT_LIST;
50	} else {	49	} else {
51	return emptyDtypes;	50	return emptyDtypes;
	@@ -54,11 +53,13 @@ static const inline std::initializer_list<DataType>& GetSupportDtypeList(NpuArch
54		53
55	static bool CheckDtypeValid(const aclTensorList* tensors, const aclTensor* out)	54	static bool CheckDtypeValid(const aclTensorList* tensors, const aclTensor* out)
56	{	55	{
		56	+ auto npuArch = op::GetCurrentPlatformInfo().GetCurNpuArch();
		57	+ const auto& dTypeSupportList = GetSupportDtypeList(npuArch);
57	op::DataType inputType = (*tensors)[0]->GetDataType();	58	op::DataType inputType = (*tensors)[0]->GetDataType();
58	- if (!CheckType(inputType, ASCEND910B_DTYPE_SUPPORT_LIST)) {	59	+ if (!CheckType(inputType, dTypeSupportList)) {
59	OP_LOGE(	60	OP_LOGE(
60	ACLNN_ERR_PARAM_INVALID, "tensor %lu not implemented for %s, should be in dtype support list %s.", 0,	61	ACLNN_ERR_PARAM_INVALID, "tensor %lu not implemented for %s, should be in dtype support list %s.", 0,
61	- op::ToString(inputType).GetString(), op::ToString(ASCEND910B_DTYPE_SUPPORT_LIST).GetString());	62	+ op::ToString(inputType).GetString(), op::ToString(dTypeSupportList).GetString());
62	return false;	63	return false;
63	}	64	}
64	for (uint64_t i = 1; i < tensors->Size(); i++) {	65	for (uint64_t i = 1; i < tensors->Size(); i++) {
	@@ -67,7 +68,7 @@ static bool CheckDtypeValid(const aclTensorList* tensors, const aclTensor* out)
67	return false;	68	return false;
68	}	69	}
69	}	70	}
70	- OP_CHECK_DTYPE_NOT_SUPPORT(out, ASCEND910B_DTYPE_SUPPORT_LIST, return false);	71	+ OP_CHECK_DTYPE_NOT_SUPPORT(out, dTypeSupportList, return false);
71	if (inputType == DataType::DT_FLOAT && out->GetDataType() != DataType::DT_FLOAT) {	72	if (inputType == DataType::DT_FLOAT && out->GetDataType() != DataType::DT_FLOAT) {
72	OP_LOGE(ACLNN_ERR_PARAM_INVALID, "output dtype must be float when input dtype is float.");	73	OP_LOGE(ACLNN_ERR_PARAM_INVALID, "output dtype must be float when input dtype is float.");
73	return false;	74	return false;

Mconversion/chunk_cat/op_host/chunk_cat_def.cpp+1-0

	@@ -37,6 +37,7 @@ public:
37		37
38	this->AICore().AddConfig("ascend910b");	38	this->AICore().AddConfig("ascend910b");
39	this->AICore().AddConfig("ascend910_93");	39	this->AICore().AddConfig("ascend910_93");
		40	+ this->AICore().AddConfig("ascend950");
40	}	41	}
41	};	42	};
42		43

Mconversion/chunk_cat/op_host/chunk_cat_tiling.cpp+19-4

	@@ -15,6 +15,7 @@
15		15
16	#include "chunk_cat_tiling.h"	16	#include "chunk_cat_tiling.h"
17	#include "tiling/platform/platform_ascendc.h"	17	#include "tiling/platform/platform_ascendc.h"
		18	+#include "op_host/tiling_base_util.h"
18		19
19	namespace optiling {	20	namespace optiling {
20		21
	@@ -52,6 +53,7 @@ ge::graphStatus ChunkCatTiling::GetPlatformInfo()
52	ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize_);	53	ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize_);
53	OP_CHECK_IF(ubSize_ == 0, OP_LOGE(context_, "ubSize is 0"), return ge::GRAPH_FAILED);	54	OP_CHECK_IF(ubSize_ == 0, OP_LOGE(context_, "ubSize is 0"), return ge::GRAPH_FAILED);
54	sysWorkspaceSize_ = ascendcPlatform.GetLibApiWorkSpaceSize();	55	sysWorkspaceSize_ = ascendcPlatform.GetLibApiWorkSpaceSize();
		56	+ isRegBase = Ops::Base::IsRegbaseSocVersion(context_);
55	return ge::GRAPH_SUCCESS;	57	return ge::GRAPH_SUCCESS;
56	}	58	}
57		59
	@@ -90,12 +92,13 @@ ge::graphStatus ChunkCatTiling::GetInputInfo()
90	srcDtypeSize_ = ge::GetSizeByDataType(inputDataType);	92	srcDtypeSize_ = ge::GetSizeByDataType(inputDataType);
91	OP_CHECK_IF(srcDtypeSize_ == 0, OP_LOGE(context_, "input dtype size can not be 0"), return ge::GRAPH_FAILED);	93	OP_CHECK_IF(srcDtypeSize_ == 0, OP_LOGE(context_, "input dtype size can not be 0"), return ge::GRAPH_FAILED);
92	srcEleUbBlock_ = UB_BLOCK_SIZE / srcDtypeSize_;	94	srcEleUbBlock_ = UB_BLOCK_SIZE / srcDtypeSize_;
		95	+ uint32_t reserveUb = isRegBase ? 0 : RESERVE_UB;
93	if (inputDataType != outputDataType) {	96	if (inputDataType != outputDataType) {
94	- inUbSize_ = (ubSize_ + RESERVE_UB) / ONETHIRD;	97	+ inUbSize_ = (ubSize_ + reserveUb) / ONETHIRD;
95	} else {	98	} else {
96	- inUbSize_ = (ubSize_ + RESERVE_UB) / HALF;	99	+ inUbSize_ = (ubSize_ + reserveUb) / HALF;
97	}	100	}
98	- outUbSize_ = (ubSize_ + RESERVE_UB) - inUbSize_ ;	101	+ outUbSize_ = (ubSize_ + reserveUb) - inUbSize_;
99	return ge::GRAPH_SUCCESS;	102	return ge::GRAPH_SUCCESS;
100	}	103	}
101		104
	@@ -138,7 +141,19 @@ ge::graphStatus ChunkCatTiling::CalculateOutputInfo()
138		141
139	void ChunkCatTiling::DoUbSplit()	142	void ChunkCatTiling::DoUbSplit()
140	{	143	{
141	- if (isAllAlign_ \|\| isOneConcat_) {	144	+ if (isRegBase) {
		145	+ // 列切
		146	+ uint32_t colLimit = inUbSize_ / srcDtypeSize_;
		147	+ colLimit = colLimit - 32 * srcEleUbBlock_;
		148	+ int64_t ubColLoop = (outputCol_ + colLimit - 1) / colLimit ;
		149	+ ubColFactor_ = (outputCol_ + ubColLoop - 1) / ubColLoop;
		150	+ ubColFactor_ = (ubColFactor_ + srcEleUbBlock_ - 1) / srcEleUbBlock_ * srcEleUbBlock_;
		151	+ ubColFactor_ = ubColFactor_ > colLimit ? colLimit : ubColFactor_;
		152	+ // 行切
		153	+ uint32_t rowLimit = ubColLoop == 1 ? colLimit / ubColFactor_ : 1;
		154	+ int64_t ubRowLoop = (outputRow_ + rowLimit - 1) / rowLimit ;
		155	+ ubRowFactor_ = (outputRow_ + ubRowLoop - 1) / ubRowLoop;
		156	+ } else if (isAllAlign_ \|\| isOneConcat_) {
142	// 列切	157	// 列切
143	uint32_t colLimit = inUbSize_ / srcDtypeSize_;	158	uint32_t colLimit = inUbSize_ / srcDtypeSize_;
144	colLimit = (isOneConcat_) ? colLimit - 32 * (srcEleUbBlock_ - 1) : colLimit;	159	colLimit = (isOneConcat_) ? colLimit - 32 * (srcEleUbBlock_ - 1) : colLimit;

Mconversion/chunk_cat/op_host/chunk_cat_tiling.h+1-0

	@@ -40,6 +40,7 @@ public:
40		40
41	private:	41	private:
42	gert::TilingContext* context_;	42	gert::TilingContext* context_;
		43	+ bool isRegBase;
43	int32_t coreNum_{0};	44	int32_t coreNum_{0};
44	uint64_t ubSize_{0};	45	uint64_t ubSize_{0};
45	uint64_t sysWorkspaceSize_{0};	46	uint64_t sysWorkspaceSize_{0};

Aconversion/chunk_cat/op_host/config/ascend950/chunk_cat_binary.json+292-0

@@ -0,0 +1,292 @@
	1	+{
	2	+ "op_type": "ChunkCat",
	3	+ "op_list": [
	4	+ {
	5	+ "bin_filename": "ChunkCat_fp16_fp16",
	6	+ "inputs": [
	7	+ [
	8	+ {
	9	+ "name": "x",
	10	+ "index": 0,
	11	+ "dtype": "float16",
	12	+ "format": "ND",
	13	+ "paramType": "dynamic",
	14	+ "shape": [
	15	+ -2
	16	+ ]
	17	+ }
	18	+ ]
	19	+ ],
	20	+ "outputs": [
	21	+ {
	22	+ "name": "y",
	23	+ "index": 0,
	24	+ "dtype": "float16",
	25	+ "format": "ND",
	26	+ "paramType": "required",
	27	+ "shape": [
	28	+ -2
	29	+ ]
	30	+ }
	31	+ ],
	32	+ "attrs": [
	33	+ {
	34	+ "name": "dim",
	35	+ "dtype": "int",
	36	+ "value": null
	37	+ },
	38	+ {
	39	+ "name": "num_chunks",
	40	+ "dtype": "int",
	41	+ "value": null
	42	+ }
	43	+ ]
	44	+ },
	45	+ {
	46	+ "bin_filename": "ChunkCat_fp16_bf16",
	47	+ "inputs": [
	48	+ [
	49	+ {
	50	+ "name": "x",
	51	+ "index": 0,
	52	+ "dtype": "float16",
	53	+ "format": "ND",
	54	+ "paramType": "dynamic",
	55	+ "shape": [
	56	+ -2
	57	+ ]
	58	+ }
	59	+ ]
	60	+ ],
	61	+ "outputs": [
	62	+ {
	63	+ "name": "y",
	64	+ "index": 0,
	65	+ "dtype": "bfloat16",
	66	+ "format": "ND",
	67	+ "paramType": "required",
	68	+ "shape": [
	69	+ -2
	70	+ ]
	71	+ }
	72	+ ],
	73	+ "attrs": [
	74	+ {
	75	+ "name": "dim",
	76	+ "dtype": "int",
	77	+ "value": null
	78	+ },
	79	+ {
	80	+ "name": "num_chunks",
	81	+ "dtype": "int",
	82	+ "value": null
	83	+ }
	84	+ ]
	85	+ },
	86	+ {
	87	+ "bin_filename": "ChunkCat_fp16_fp32",
	88	+ "inputs": [
	89	+ [
	90	+ {
	91	+ "name": "x",
	92	+ "index": 0,
	93	+ "dtype": "float16",
	94	+ "format": "ND",
	95	+ "paramType": "dynamic",
	96	+ "shape": [
	97	+ -2
	98	+ ]
	99	+ }
	100	+ ]
	101	+ ],
	102	+ "outputs": [
	103	+ {
	104	+ "name": "y",
	105	+ "index": 0,
	106	+ "dtype": "float32",
	107	+ "format": "ND",
	108	+ "paramType": "required",
	109	+ "shape": [
	110	+ -2
	111	+ ]
	112	+ }
	113	+ ],
	114	+ "attrs": [
	115	+ {
	116	+ "name": "dim",
	117	+ "dtype": "int",
	118	+ "value": null
	119	+ },
	120	+ {
	121	+ "name": "num_chunks",
	122	+ "dtype": "int",
	123	+ "value": null
	124	+ }
	125	+ ]
	126	+ },
	127	+ {
	128	+ "bin_filename": "ChunkCat_bf16_fp16",
	129	+ "inputs": [
	130	+ [
	131	+ {
	132	+ "name": "x",
	133	+ "index": 0,
	134	+ "dtype": "bfloat16",
	135	+ "format": "ND",
	136	+ "paramType": "dynamic",
	137	+ "shape": [
	138	+ -2
	139	+ ]
	140	+ }
	141	+ ]
	142	+ ],
	143	+ "outputs": [
	144	+ {
	145	+ "name": "y",
	146	+ "index": 0,
	147	+ "dtype": "float16",
	148	+ "format": "ND",
	149	+ "paramType": "required",
	150	+ "shape": [
	151	+ -2
	152	+ ]
	153	+ }
	154	+ ],
	155	+ "attrs": [
	156	+ {
	157	+ "name": "dim",
	158	+ "dtype": "int",
	159	+ "value": null
	160	+ },
	161	+ {
	162	+ "name": "num_chunks",
	163	+ "dtype": "int",
	164	+ "value": null
	165	+ }
	166	+ ]
	167	+ },
	168	+ {
	169	+ "bin_filename": "ChunkCat_bf16_bf16",
	170	+ "inputs": [
	171	+ [
	172	+ {
	173	+ "name": "x",
	174	+ "index": 0,
	175	+ "dtype": "bfloat16",
	176	+ "format": "ND",
	177	+ "paramType": "dynamic",
	178	+ "shape": [
	179	+ -2
	180	+ ]
	181	+ }
	182	+ ]
	183	+ ],
	184	+ "outputs": [
	185	+ {
	186	+ "name": "y",
	187	+ "index": 0,
	188	+ "dtype": "bfloat16",
	189	+ "format": "ND",
	190	+ "paramType": "required",
	191	+ "shape": [
	192	+ -2
	193	+ ]
	194	+ }
	195	+ ],
	196	+ "attrs": [
	197	+ {
	198	+ "name": "dim",
	199	+ "dtype": "int",
	200	+ "value": null
	201	+ },
	202	+ {
	203	+ "name": "num_chunks",
	204	+ "dtype": "int",
	205	+ "value": null
	206	+ }
	207	+ ]
	208	+ },
	209	+ {
	210	+ "bin_filename": "ChunkCat_bf16_fp32",
	211	+ "inputs": [
	212	+ [
	213	+ {
	214	+ "name": "x",
	215	+ "index": 0,
	216	+ "dtype": "bfloat16",
	217	+ "format": "ND",
	218	+ "paramType": "dynamic",
	219	+ "shape": [
	220	+ -2
	221	+ ]
	222	+ }
	223	+ ]
	224	+ ],
	225	+ "outputs": [
	226	+ {
	227	+ "name": "y",
	228	+ "index": 0,
	229	+ "dtype": "float32",
	230	+ "format": "ND",
	231	+ "paramType": "required",
	232	+ "shape": [
	233	+ -2
	234	+ ]
	235	+ }
	236	+ ],
	237	+ "attrs": [
	238	+ {
	239	+ "name": "dim",
	240	+ "dtype": "int",
	241	+ "value": null
	242	+ },
	243	+ {
	244	+ "name": "num_chunks",
	245	+ "dtype": "int",
	246	+ "value": null
	247	+ }
	248	+ ]
	249	+ },
	250	+ {
	251	+ "bin_filename": "ChunkCat_fp32_fp32",
	252	+ "inputs": [
	253	+ [
	254	+ {
	255	+ "name": "x",
	256	+ "index": 0,
	257	+ "dtype": "float32",
	258	+ "format": "ND",
	259	+ "paramType": "dynamic",
	260	+ "shape": [
	261	+ -2
	262	+ ]
	263	+ }
	264	+ ]
	265	+ ],
	266	+ "outputs": [
	267	+ {
	268	+ "name": "y",
	269	+ "index": 0,
	270	+ "dtype": "float32",
	271	+ "format": "ND",
	272	+ "paramType": "required",
	273	+ "shape": [
	274	+ -2
	275	+ ]
	276	+ }
	277	+ ],
	278	+ "attrs": [
	279	+ {
	280	+ "name": "dim",
	281	+ "dtype": "int",
	282	+ "value": null
	283	+ },
	284	+ {
	285	+ "name": "num_chunks",
	286	+ "dtype": "int",
	287	+ "value": null
	288	+ }
	289	+ ]
	290	+ }
	291	+ ]
	292	+}

Aconversion/chunk_cat/op_host/config/ascend950/chunk_cat_simplified_key.ini+13-0

@@ -0,0 +1,13 @@
	1	+; 该文件主要影响 opc 工具编译二进制kernel时， --simplified_key_mode 选项中填写的值，格式如下所示：
	2	+; [某算子]
	3	+; default=xx
	4	+; ascendxx=xx
	5	+; 其中，default为默认mode，ascendxx为可选mode，如果不同芯片有差异化要求时，需要配置；
	6	+; 1)如果没有配置：非ascendC算子继续按空处理，即opc编译命令中不添加 --simplified_key_mode 选项，AscendC算子按照 simplified_key_mode=0 处理
	7	+; 2)如果仅有default配置：各个版本按default配置
	8	+; 3)如果仅有某些平台的配置，没有default配置：对应平台的按照配置的值传递，非对应平台的：非AscendC算子继续按空处理，AscendC算子按照 simplified_key_mode=0 处理
	9	+; 4)如果default配置和平台配置都有：对应平台的使用平台的配置，非对应的平台的以default值配置。
	10	+; 5)对于自定义simplified key的情况，需要在binary_simplified_key_mode.ini 文件中显式配置为None，不传入 --simplified_key_mode 选项，由opc工具和FE框架自行判断使用何种模式
	11	+; 6)是否是AscendC算子，由 ops/build-in/tbe/op_info_cfg/parser/ascendc_config.json 中配置的算子名字和对于的平台决定
	12	+[ChunkCat]
	13	+default=0

Mconversion/chunk_cat/op_kernel/chunk_cat.cpp+10-1

	@@ -12,14 +12,22 @@
12	* \file chunk_cat.cpp	12	* \file chunk_cat.cpp
13	* \brief	13	* \brief
14	*/	14	*/
15	-	15	+#if __CCE_AICORE__ == 310
		16	+#include "chunk_cat_arch35.h"
		17	+#else
16	#include "chunk_cat.h"	18	#include "chunk_cat.h"
		19	+#endif
17		20
18	extern "C" __global__ __aicore__ void chunk_cat(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) {	21	extern "C" __global__ __aicore__ void chunk_cat(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) {
19	KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);	22	KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
20	AscendC::TPipe pipe;	23	AscendC::TPipe pipe;
21	REGISTER_TILING_DEFAULT(ChunkCatTilingData);	24	REGISTER_TILING_DEFAULT(ChunkCatTilingData);
22	GET_TILING_DATA_WITH_STRUCT(ChunkCatTilingData, tilingData, tiling);	25	GET_TILING_DATA_WITH_STRUCT(ChunkCatTilingData, tilingData, tiling);
		26	+#if __CCE_AICORE__ == 310
		27	+ ChunkCatArch35<DTYPE_X, DTYPE_Y> op(&pipe);
		28	+ op.Init(x, y, tilingData);
		29	+ op.Process();
		30	+#else
23	#if (ORIG_DTYPE_X == ORIG_DTYPE_Y)	31	#if (ORIG_DTYPE_X == ORIG_DTYPE_Y)
24	ChunkCat<DTYPE_X, DTYPE_Y> op(&pipe);	32	ChunkCat<DTYPE_X, DTYPE_Y> op(&pipe);
25	op.Init(x, y, tilingData);	33	op.Init(x, y, tilingData);
	@@ -29,4 +37,5 @@ extern "C" __global__ __aicore__ void chunk_cat(GM_ADDR x, GM_ADDR y, GM_ADDR wo
29	op.Init(x, y, tilingData);	37	op.Init(x, y, tilingData);
30	op.Process();	38	op.Process();
31	#endif	39	#endif
		40	+#endif
32	}	41	}

Mconversion/chunk_cat/op_kernel/chunk_cat.h+161-346

	@@ -13,67 +13,41 @@
13	* \brief	13	* \brief
14	*/	14	*/
15		15
16	-#ifndef _CHUNK_CAT_DATA_H_	16	+#ifndef _CHUNK_CAT_H_
17	-#define _CHUNK_CAT_DATA_H_	17	+#define _CHUNK_CAT_H_
18		18
19	#include "chunk_cat_common.h"	19	#include "chunk_cat_common.h"
20		20
21	using namespace AscendC;	21	using namespace AscendC;
22	-template <typename T1, typename T2, bool NEAD_CAST=false>	22	+template <typename T1, typename T2, bool NEED_CAST=false>
23	-class ChunkCat	23	+class ChunkCat : public ChunkCatCommon<T1, T2>
24	{	24	{
25	public:	25	public:
26	- __aicore__ inline ChunkCat(TPipe *pipe) : pipe_(pipe) {}	26	+ __aicore__ inline ChunkCat(TPipe *pipe) : ChunkCatCommon<T1, T2>(pipe) {}
27		27
28	__aicore__ inline void Init(GM_ADDR x, GM_ADDR y, const ChunkCatTilingData& tilingData)	28	__aicore__ inline void Init(GM_ADDR x, GM_ADDR y, const ChunkCatTilingData& tilingData)
29	{	29	{
30	- blockIdx_ = GetBlockIdx();	30	+ this->InitCommon(x, y, tilingData);
31	- int32_t usedCoreNum = GetBlockNum();
32	// 获取tiling信息	31	// 获取tiling信息
33	isAllAlign_ = tilingData.isAllAlign;	32	isAllAlign_ = tilingData.isAllAlign;
34	isHalfAlign_ = tilingData.isHalfAlign;	33	isHalfAlign_ = tilingData.isHalfAlign;
35	isOneConcat_ = tilingData.isOneConcat;	34	isOneConcat_ = tilingData.isOneConcat;
36	- dim_ = tilingData.dim;	35	+ colRepeatNum_ = isHalfAlign_ ? HALF : this->srcEleUbBlock_;
37	- numChunk_ = tilingData.numChunk;
38	- outputRow_ = tilingData.outputRow;
39	- outputCol_ = tilingData.outputCol;
40	- blockRowFactor_ = tilingData.blockRowFactor;
41	- blockColFactor_ = tilingData.blockColFactor;
42	- tailBlockRowFactor_ = tilingData.tailBlockRowFactor;
43	- tailBlockColFactor_ = tilingData.tailBlockColFactor;
44	- ubRowFactor_ = tilingData.ubRowFactor;
45	- ubColFactor_ = tilingData.ubColFactor;
46	- inputNum_ = tilingData.inputNum;
47	- srcEleUbBlock_ = UB_BLOCK_SIZE / sizeof(T1);
48	- dstEleUbBlock_ = UB_BLOCK_SIZE / sizeof(T2);
49	- colRepeatNum_ = isHalfAlign_ ? HALF : srcEleUbBlock_;
50		36
51	- blockRowGroup_ = blockIdx_ / tilingData.blockColNum;	37	+ srcLocalT2_ = this->srcLocal_.template ReinterpretCast<T2>();
52	- blockColGroup_ = blockIdx_ % tilingData.blockColNum;	38	+ dstLocalT1_ = this->dstLocal_.template ReinterpretCast<T1>();
53	- currentBlockRowFactor_ = blockRowGroup_ == tilingData.blockRowNum - 1 ? tailBlockRowFactor_ : blockRowFactor_;	39	+ dstLocalFP32_ = this->dstLocal_.template ReinterpretCast<float>();
54	- currentBlockColFactor_ = blockColGroup_ == tilingData.blockColNum - 1 ? tailBlockColFactor_ : blockColFactor_;
55	- int64_t dstGmOffset = blockRowGroup_ * blockRowFactor_ * outputCol_ + blockColGroup_ * blockColFactor_;
56	- dstGlobal_.SetGlobalBuffer((__gm__ T2*)y + dstGmOffset);
57	- inputList_ = ListTensorDesc(reinterpret_cast<__gm__ void*>(x));
58	-
59	- pipe_->InitBuffer(srcBuf_, tilingData.inUbSize);
60	- pipe_->InitBuffer(dstBuf_, tilingData.outUbSize);
61	- srcLocal_ = srcBuf_.Get<T1>();
62	- srcLocalT2_ = srcLocal_.template ReinterpretCast<T2>();
63	- dstLocal_ = dstBuf_.Get<T2>();
64	- dstLocalT1_ = dstLocal_.template ReinterpretCast<T1>();
65	- dstLocalFP32_ = dstLocal_.template ReinterpretCast<float>();
66	}	40	}
67		41
68	__aicore__ inline void Process()	42	__aicore__ inline void Process()
69	{	43	{
70	- int64_t rowLoop = GetAlign(currentBlockRowFactor_, ubRowFactor_) / ubRowFactor_;	44	+ int64_t rowLoop = this->GetAlign(this->currentBlockRowFactor_, this->ubRowFactor_) / this->ubRowFactor_;
71	- int64_t colLoop = GetAlign(currentBlockColFactor_, ubColFactor_) / ubColFactor_;	45	+ int64_t colLoop = this->GetAlign(this->currentBlockColFactor_, this->ubColFactor_) / this->ubColFactor_;
72	- int64_t rowTail = currentBlockRowFactor_ % ubRowFactor_;	46	+ int64_t rowTail = this->currentBlockRowFactor_ % this->ubRowFactor_;
73	- int64_t colTail = currentBlockColFactor_ % ubColFactor_;	47	+ int64_t colTail = this->currentBlockColFactor_ % this->ubColFactor_;
74		48
75	uint64_t buf[10];	49	uint64_t buf[10];
76	- desc_.SetShapeAddr(buf); // 用于获取shape信息	50	+ this->desc_.SetShapeAddr(buf); // 用于获取shape信息
77	int64_t inputCol[32];	51	int64_t inputCol[32];
78		52
79	for (int64_t i = 0; i < rowLoop * colLoop; i++) {	53	for (int64_t i = 0; i < rowLoop * colLoop; i++) {
	@@ -82,80 +56,61 @@ public:
82	ubLoopInfo.ubRowGroup = i / colLoop;	56	ubLoopInfo.ubRowGroup = i / colLoop;
83	ubLoopInfo.ubColGroup = i % colLoop;	57	ubLoopInfo.ubColGroup = i % colLoop;
84	ubLoopInfo.currentUbRowFactor = (rowTail != 0 && ubLoopInfo.ubRowGroup == rowLoop - 1) ?	58	ubLoopInfo.currentUbRowFactor = (rowTail != 0 && ubLoopInfo.ubRowGroup == rowLoop - 1) ?
85	- rowTail : ubRowFactor_;	59	+ rowTail : this->ubRowFactor_;
86	ubLoopInfo.currentUbColFactor = (colTail != 0 && ubLoopInfo.ubColGroup == colLoop - 1) ?	60	ubLoopInfo.currentUbColFactor = (colTail != 0 && ubLoopInfo.ubColGroup == colLoop - 1) ?
87	- colTail : ubColFactor_;	61	+ colTail : this->ubColFactor_;
		62	+ // 1、清零ub
		63	+ dupToZero();
88	// 搬入	64	// 搬入
89	CopyIn(ubLoopInfo);	65	CopyIn(ubLoopInfo);
90	// 计算	66	// 计算
91	Compute(ubLoopInfo);	67	Compute(ubLoopInfo);
92	// 搬出	68	// 搬出
93	- CopyCout(ubLoopInfo);	69	+ CopyOut(ubLoopInfo);
94	}	70	}
95	}	71	}
96		72
97	private:	73	private:
98	__aicore__ inline void CopyIn(UbLoopInfo& ubLoopInfo)	74	__aicore__ inline void CopyIn(UbLoopInfo& ubLoopInfo)
99	{	75	{
100	- // 1、清零ub
101	- dupToZero();
102	// 2、遍历tensor搬运	76	// 2、遍历tensor搬运
103	int64_t totalCol = 0;	77	int64_t totalCol = 0;
104	int64_t localOffset = 0;	78	int64_t localOffset = 0;
105	- ubLoopInfo.colStart = blockColGroup_ * blockColFactor_ + ubLoopInfo.ubColGroup * ubColFactor_;	79	+ ubLoopInfo.colStart = this->blockColGroup_ * this->blockColFactor_ + ubLoopInfo.ubColGroup * this->ubColFactor_;
106	- ubLoopInfo.rowStart = blockRowGroup_ * blockRowFactor_ + ubLoopInfo.ubRowGroup * ubRowFactor_;	80	+ ubLoopInfo.rowStart = this->blockRowGroup_ * this->blockRowFactor_ + ubLoopInfo.ubRowGroup * this->ubRowFactor_;
107		81
108	- for (uint32_t i = 0; i < inputNum_; i++) {	82	+ for (uint32_t i = 0; i < this->inputNum_; i++) {
109	if (ubLoopInfo.totalUbCol >= ubLoopInfo.currentUbColFactor) {	83	if (ubLoopInfo.totalUbCol >= ubLoopInfo.currentUbColFactor) {
110	break;	84	break;
111	}	85	}
112	- srcGlobal_.SetGlobalBuffer(inputList_.GetDataPtr<T1>(i));	86	+ this->srcGlobal_.SetGlobalBuffer(this->inputList_.template GetDataPtr<T1>(i));
113	TensorInfo tensorInfo{};	87	TensorInfo tensorInfo{};
114	- inputList_.GetDesc(desc_, i); // scalar很大(将buf改为局部变量有改善)	88	+ this->GetChunkInfo(i, tensorInfo);
115	- // 获取chunk相关信息	89	+ // 判断当前核是否处理当前tensor
116	- tensorInfo.chunkDimSize = desc_.GetShape(dim_);	90	+ if (!this->IsTensorInRange(totalCol, ubLoopInfo, tensorInfo)) {
117	- tensorInfo.chunkCol = (tensorInfo.chunkDimSize + numChunk_ - 1) / numChunk_;	91	+ totalCol += tensorInfo.tensorCol;
118	- CopyInChunk(totalCol, localOffset, ubLoopInfo, tensorInfo);	92	+ } else {
119	-	93	+ this->SplitTensorDim0(totalCol, ubLoopInfo, tensorInfo);
120	- if (isOneConcat_ && ubLoopInfo.count > 31) {	94	+ ubLoopInfo.inputCol[ubLoopInfo.count] = (!isOneConcat_ && !isAllAlign_ && tensorInfo.isSplit) ? -tensorInfo.splitCol : tensorInfo.splitCol;
121	- // 计算	95	+ tensorInfo.chunkRow = tensorInfo.chunkDimSize / tensorInfo.chunkCol;
122	- Compute(ubLoopInfo);	96	+ tensorInfo.chunkRowAlign = this->GetAlign(tensorInfo.chunkDimSize, tensorInfo.chunkCol) / tensorInfo.chunkCol;
123	- // 搬出	97	+ int64_t localOffsetIncrement = (isOneConcat_ \|\| isAllAlign_) ? ubLoopInfo.currentUbRowFactor :
124	- CopyCout(ubLoopInfo);	98	+ (isHalfAlign_ ? TRANS_BLOCK * HALF : TRANS_BLOCK * this->srcEleUbBlock_);
125	- localOffset = 0;	99	+ if (ubLoopInfo.rowStart >= tensorInfo.chunkRowAlign) {
126	- ubLoopInfo.preCatCol += ubLoopInfo.totalUbCol;	100	+ ubLoopInfo.inputCol[ubLoopInfo.count] = tensorInfo.splitCol;
127	- ubLoopInfo.count = 0;	101	+ ubLoopInfo.totalUbColAlign += tensorInfo.splitCol;
128	- ubLoopInfo.totalUbCol = 0;	102	+ localOffsetIncrement *= tensorInfo.splitCol;
129	- ubLoopInfo.totalUbColAlign = 0;	103	+ } else {
130	- }	104	+ this->CopyInChunk(totalCol, localOffset, ubLoopInfo, tensorInfo);
131	- else if (ubLoopInfo.count > 31) {	105	+ localOffsetIncrement *= (isOneConcat_ \|\| tensorInfo.isSplit) ? tensorInfo.splitColAlign : tensorInfo.splitCol;
132	- // 提前做部分concat	106	+ ubLoopInfo.totalUbColAlign += (isOneConcat_ \|\| (!isAllAlign_ && tensorInfo.isSplit)) ? tensorInfo.splitColAlign : tensorInfo.splitCol;
133	- if (!ubLoopInfo.isAllZero) {
134	- SetFlag<HardEvent::MTE2_V>(event_);
135	- WaitFlag<HardEvent::MTE2_V>(event_);
136	- if (isAllAlign_) {
137	- UBRearrange4Concat(ubLoopInfo, srcLocal_, dstLocalT1_);
138	- DataCopy(srcLocal_, dstLocalT1_, ubLoopInfo.currentUbRowFactor * ubLoopInfo.totalUbCol);
139	- } else {
140	- // 3、ub重排
141	- UBRearrange4Trans(ubLoopInfo, srcLocal_, dstLocalT1_);
142	- // 4、跨block对齐转置
143	- Trans1(ubLoopInfo, dstLocalT1_, srcLocal_);
144	- // 5、ub重排
145	- UBRearrange4TransConcat<true>(ubLoopInfo, srcLocal_, dstLocalT1_);
146	- // 6、跨block对齐转置
147	- Trans2<true>(ubLoopInfo, dstLocalT1_, srcLocal_);
148	- }
149	- SetFlag<HardEvent::V_MTE2>(event_);
150	- WaitFlag<HardEvent::V_MTE2>(event_);
151	}	107	}
152	- ubLoopInfo.inputCol[0] = ubLoopInfo.totalUbCol;	108	+ ubLoopInfo.totalUbCol += tensorInfo.splitCol;
153	- ubLoopInfo.count = 1;	109	+ ubLoopInfo.count++;
154	- ubLoopInfo.totalUbColAlign = ubLoopInfo.totalUbCol;	110	+ totalCol += tensorInfo.tensorCol;
155	- localOffset = isAllAlign_ ? ubLoopInfo.currentUbRowFactor * ubLoopInfo.totalUbCol :	111	+ localOffset += localOffsetIncrement;
156	- (isHalfAlign_ ? TRANS_BLOCK * HALF * ubLoopInfo.totalUbCol :
157	- TRANS_BLOCK * srcEleUbBlock_ * ubLoopInfo.totalUbCol);
158	}	112	}
		113	+ ComputeOver32(totalCol, localOffset, ubLoopInfo, tensorInfo);
159	}	114	}
160	}	115	}
161		116
	@@ -166,23 +121,23 @@ private:
166	}	121	}
167	else if (ubLoopInfo.isAllZero) {	122	else if (ubLoopInfo.isAllZero) {
168	PipeBarrier<PIPE_V>();	123	PipeBarrier<PIPE_V>();
169	- SetFlag<HardEvent::MTE3_V>(event_);	124	+ SetFlag<HardEvent::MTE3_V>(this->event_);
170	- WaitFlag<HardEvent::MTE3_V>(event_);	125	+ WaitFlag<HardEvent::MTE3_V>(this->event_);
171	- if constexpr (NEAD_CAST) {	126	+ if constexpr (NEED_CAST) {
172	- uint32_t castCount = ubLoopInfo.currentUbRowFactor * GetAlign(ubLoopInfo.currentUbColFactor, dstEleUbBlock_);	127	+ uint32_t castCount = ubLoopInfo.currentUbRowFactor * this->GetAlign(ubLoopInfo.currentUbColFactor, this->dstEleUbBlock_);
173	DoCast(ubLoopInfo, castCount);	128	DoCast(ubLoopInfo, castCount);
174	} else {	129	} else {
175	- DataCopy(dstLocalT1_, srcLocal_, ubLoopInfo.currentUbRowFactor * GetAlign(ubLoopInfo.currentUbColFactor, dstEleUbBlock_));	130	+ DataCopy(dstLocalT1_, this->srcLocal_, ubLoopInfo.currentUbRowFactor * this->GetAlign(ubLoopInfo.currentUbColFactor, this->dstEleUbBlock_));
176	}	131	}
177	}	132	}
178	- else if (ubLoopInfo.count == 1 && ubLoopInfo.currentUbColFactor % srcEleUbBlock_ == 0) {	133	+ else if (ubLoopInfo.count == 1 && ubLoopInfo.currentUbColFactor % this->srcEleUbBlock_ == 0) {
179	- SetFlag<HardEvent::MTE2_V>(event_);	134	+ SetFlag<HardEvent::MTE2_V>(this->event_);
180	- WaitFlag<HardEvent::MTE2_V>(event_);	135	+ WaitFlag<HardEvent::MTE2_V>(this->event_);
181	- if constexpr (NEAD_CAST) {	136	+ if constexpr (NEED_CAST) {
182	uint32_t castCount = ubLoopInfo.currentUbRowFactor * ubLoopInfo.currentUbColFactor;	137	uint32_t castCount = ubLoopInfo.currentUbRowFactor * ubLoopInfo.currentUbColFactor;
183	DoCast(ubLoopInfo, castCount);	138	DoCast(ubLoopInfo, castCount);
184	} else {	139	} else {
185	- DataCopy(dstLocalT1_, srcLocal_, ubLoopInfo.currentUbRowFactor * ubLoopInfo.currentUbColFactor);	140	+ DataCopy(dstLocalT1_, this->srcLocal_, ubLoopInfo.currentUbRowFactor * ubLoopInfo.currentUbColFactor);
186	}	141	}
187	}	142	}
188	else if (isAllAlign_) {	143	else if (isAllAlign_) {
	@@ -190,199 +145,87 @@ private:
190	} else {	145	} else {
191	ComputeNotAlign(ubLoopInfo);	146	ComputeNotAlign(ubLoopInfo);
192	}	147	}
193	- SetFlag<HardEvent::V_MTE3>(event_);	148	+ SetFlag<HardEvent::V_MTE3>(this->event_);
194	- WaitFlag<HardEvent::V_MTE3>(event_);	149	+ WaitFlag<HardEvent::V_MTE3>(this->event_);
195	}	150	}
196		151
197	- __aicore__ inline void CopyCout(const UbLoopInfo& ubLoopInfo)	152	+ __aicore__ inline void CopyOut(const UbLoopInfo& ubLoopInfo)
198	{	153	{
199	if (isOneConcat_) {	154	if (isOneConcat_) {
200	int64_t localOffset = 0;	155	int64_t localOffset = 0;
201	- int64_t globalOffset = ubLoopInfo.ubRowGroup * ubRowFactor_ * outputCol_ + ubLoopInfo.ubColGroup * ubColFactor_ + ubLoopInfo.preCatCol;	156	+ int64_t globalOffset = ubLoopInfo.ubRowGroup * this->ubRowFactor_ * this->outputCol_ + ubLoopInfo.ubColGroup * this->ubColFactor_ + ubLoopInfo.preCatCol;
202	for (int i = 0; i < ubLoopInfo.count; i++) {	157	for (int i = 0; i < ubLoopInfo.count; i++) {
203	uint16_t blockCount = ubLoopInfo.currentUbRowFactor;	158	uint16_t blockCount = ubLoopInfo.currentUbRowFactor;
204	uint32_t blockLen = ubLoopInfo.inputCol[i] * sizeof(T2);	159	uint32_t blockLen = ubLoopInfo.inputCol[i] * sizeof(T2);
205	- uint32_t dstStride = (outputCol_ - ubLoopInfo.inputCol[i]) * sizeof(T2);	160	+ uint32_t dstStride = (this->outputCol_ - ubLoopInfo.inputCol[i]) * sizeof(T2);
206	DataCopyExtParams copyParamsOut{blockCount, blockLen, 0, dstStride, 0};	161	DataCopyExtParams copyParamsOut{blockCount, blockLen, 0, dstStride, 0};
207	- DataCopyPad(dstGlobal_[globalOffset], dstLocal_[localOffset], copyParamsOut);	162	+ DataCopyPad(this->dstGlobal_[globalOffset], this->dstLocal_[localOffset], copyParamsOut);
208	- localOffset += GetAlign(ubLoopInfo.inputCol[i], srcEleUbBlock_);	163	+ localOffset += this->GetAlign(ubLoopInfo.inputCol[i], this->srcEleUbBlock_);
209	globalOffset += ubLoopInfo.inputCol[i];	164	globalOffset += ubLoopInfo.inputCol[i];
210	}	165	}
211	- SetFlag<HardEvent::MTE3_MTE2>(event_);	166	+ SetFlag<HardEvent::MTE3_MTE2>(this->event_);
212	- WaitFlag<HardEvent::MTE3_MTE2>(event_);	167	+ WaitFlag<HardEvent::MTE3_MTE2>(this->event_);
213	return;	168	return;
214	}	169	}
215	uint16_t blockCount = ubLoopInfo.currentUbRowFactor;	170	uint16_t blockCount = ubLoopInfo.currentUbRowFactor;
216	uint32_t blockLen = ubLoopInfo.currentUbColFactor * sizeof(T2);	171	uint32_t blockLen = ubLoopInfo.currentUbColFactor * sizeof(T2);
217	- uint32_t dstStride = (outputCol_ - ubLoopInfo.currentUbColFactor)* sizeof(T2);	172	+ uint32_t dstStride = (this->outputCol_ - ubLoopInfo.currentUbColFactor)* sizeof(T2);
218	DataCopyExtParams copyParamsOut{blockCount, blockLen, 0, dstStride, 0};	173	DataCopyExtParams copyParamsOut{blockCount, blockLen, 0, dstStride, 0};
219	- int64_t dstOffset = ubLoopInfo.ubRowGroup * ubRowFactor_ * outputCol_ + ubLoopInfo.ubColGroup * ubColFactor_;	174	+ int64_t dstOffset = ubLoopInfo.ubRowGroup * this->ubRowFactor_ * this->outputCol_ + ubLoopInfo.ubColGroup * this->ubColFactor_;
220	- DataCopyPad(dstGlobal_[dstOffset], dstLocal_, copyParamsOut);	175	+ DataCopyPad(this->dstGlobal_[dstOffset], this->dstLocal_, copyParamsOut);
221	- SetFlag<HardEvent::MTE3_MTE2>(event_);	176	+ SetFlag<HardEvent::MTE3_MTE2>(this->event_);
222	- WaitFlag<HardEvent::MTE3_MTE2>(event_);	177	+ WaitFlag<HardEvent::MTE3_MTE2>(this->event_);
223	}	178	}
224		179
225	- __aicore__ inline int64_t GetAlign(int64_t value, int64_t align)	180	+ __aicore__ inline void ComputeOver32(int64_t& totalCol, int64_t& localOffset, UbLoopInfo& ubLoopInfo, TensorInfo& tensorInfo)
226	{	181	{
227	- return align == 0 ? value : (value + align - 1) / align * align;	182	+ if (isOneConcat_ && ubLoopInfo.count > 31) {
		183	+ // 计算
		184	+ Compute(ubLoopInfo);
		185	+ // 搬出
		186	+ CopyOut(ubLoopInfo);
		187	+ localOffset = 0;
		188	+ ubLoopInfo.preCatCol += ubLoopInfo.totalUbCol;
		189	+ ubLoopInfo.count = 0;
		190	+ ubLoopInfo.totalUbCol = 0;
		191	+ ubLoopInfo.totalUbColAlign = 0;
		192	+ }
		193	+ else if (ubLoopInfo.count > 31) {
		194	+ // 提前做部分concat
		195	+ if (!ubLoopInfo.isAllZero) {
		196	+ SetFlag<HardEvent::MTE2_V>(this->event_);
		197	+ WaitFlag<HardEvent::MTE2_V>(this->event_);
		198	+ if (isAllAlign_) {
		199	+ UBRearrange4Concat(ubLoopInfo, this->srcLocal_, dstLocalT1_);
		200	+ DataCopy(this->srcLocal_, dstLocalT1_, ubLoopInfo.currentUbRowFactor * ubLoopInfo.totalUbCol);
		201	+ } else {
		202	+ // 3、ub重排
		203	+ UBRearrange4Trans(ubLoopInfo, this->srcLocal_, dstLocalT1_);
		204	+ // 4、跨block对齐转置
		205	+ Trans1(ubLoopInfo, dstLocalT1_, this->srcLocal_);
		206	+ // 5、ub重排
		207	+ UBRearrange4TransConcat<true>(ubLoopInfo, this->srcLocal_, dstLocalT1_);
		208	+ // 6、跨block对齐转置
		209	+ Trans2<true>(ubLoopInfo, dstLocalT1_, this->srcLocal_);
		210	+ }
		211	+ SetFlag<HardEvent::V_MTE2>(this->event_);
		212	+ WaitFlag<HardEvent::V_MTE2>(this->event_);
		213	+ }
		214	+ ubLoopInfo.inputCol[0] = ubLoopInfo.totalUbCol;
		215	+ ubLoopInfo.count = 1;
		216	+ ubLoopInfo.totalUbColAlign = ubLoopInfo.totalUbCol;
		217	+ localOffset = isAllAlign_ ? ubLoopInfo.currentUbRowFactor * ubLoopInfo.totalUbCol :
		218	+ (isHalfAlign_ ? TRANS_BLOCK * HALF * ubLoopInfo.totalUbCol :
		219	+ TRANS_BLOCK * this->srcEleUbBlock_ * ubLoopInfo.totalUbCol);
		220	+ }
228	}	221	}
229		222
230	__aicore__ inline void dupToZero()	223	__aicore__ inline void dupToZero()
231	{	224	{
232	T1 inputVal(0.0);	225	T1 inputVal(0.0);
233	- Duplicate<T1>(srcLocal_, inputVal, srcLocal_.GetSize());	226	+ Duplicate<T1>(this->srcLocal_, inputVal, this->srcLocal_.GetSize());
234	- SetFlag<HardEvent::V_MTE2>(event_);	227	+ SetFlag<HardEvent::V_MTE2>(this->event_);
235	- WaitFlag<HardEvent::V_MTE2>(event_);	228	+ WaitFlag<HardEvent::V_MTE2>(this->event_);
236	- }
237	-
238	- __aicore__ inline bool IsTensorInRange(int64_t totalCol, const UbLoopInfo& ubLoopInfo, const TensorInfo& tensorInfo)
239	- {
240	- return (totalCol < ubLoopInfo.colStart + ubLoopInfo.currentUbColFactor) &&
241	- (totalCol + tensorInfo.tensorCol > ubLoopInfo.colStart);
242	- }
243	-
244	- __aicore__ inline void SplitTensorDim0(int64_t& totalCol, const UbLoopInfo& ubLoopInfo, TensorInfo& tensorInfo)
245	- {
246	- // tensor是否被切分
247	- tensorInfo.splitCol = tensorInfo.tensorCol;
248	- int64_t colEnd = ubLoopInfo.colStart + ubLoopInfo.currentUbColFactor;
249	- if (totalCol < ubLoopInfo.colStart && (totalCol + tensorInfo.tensorCol) > colEnd) {
250	- // 中间部分
251	- tensorInfo.isSplit = true;
252	- tensorInfo.splitCol = ubLoopInfo.currentUbColFactor;
253	- tensorInfo.startOffset = ubLoopInfo.colStart - totalCol;
254	- } else if (totalCol < ubLoopInfo.colStart) {
255	- // 被切分的后半部分
256	- tensorInfo.isSplit = true;
257	- tensorInfo.splitCol = totalCol + tensorInfo.tensorCol - ubLoopInfo.colStart;
258	- tensorInfo.startOffset = ubLoopInfo.colStart - totalCol;
259	- } else if ((totalCol + tensorInfo.tensorCol) > colEnd) {
260	- // 被切分的前半部分
261	- tensorInfo.isSplit = true;
262	- tensorInfo.splitCol = colEnd - totalCol;
263	- }
264	- tensorInfo.splitColAlign = GetAlign(tensorInfo.splitCol, srcEleUbBlock_);
265	- }
266	-
267	- __aicore__ inline void ExecuteDataCopy(int64_t localOffset, int64_t gmOffset, uint16_t blockCount,
268	- uint32_t blockLen, uint32_t srcStride)
269	- {
270	- AscendC::DataCopyExtParams copyParams{blockCount, blockLen, srcStride, 0, 0};
271	- uint8_t rightPadValue = (GetAlign(blockLen, UB_BLOCK_SIZE) - blockLen) / sizeof(T1);
272	- AscendC::DataCopyPadExtParams<T1> padParams{true, 0, rightPadValue, 0};
273	- AscendC::DataCopyPad(srcLocal_[localOffset], srcGlobal_[gmOffset], copyParams, padParams);
274	- }
275	-
276	- __aicore__ inline void DoRowsCopy(int64_t localOffset, const UbLoopInfo& ubLoopInfo, const TensorInfo& tensorInfo)
277	- {
278	- uint16_t blockCount = tensorInfo.isSplit ? static_cast<uint16_t>(ubLoopInfo.currentUbRowFactor) : 1;
279	- uint32_t blockLen = tensorInfo.isSplit ?
280	- static_cast<uint32_t>(tensorInfo.splitCol * sizeof(T1)) :
281	- static_cast<uint32_t>(ubLoopInfo.currentUbRowFactor * tensorInfo.splitCol * sizeof(T1));
282	- uint32_t srcStride = (tensorInfo.tensorCol - tensorInfo.splitCol) * sizeof(T1);
283	- int64_t gmOffset = ubLoopInfo.rowStart * tensorInfo.tensorCol + tensorInfo.startOffset;
284	- ExecuteDataCopy(localOffset, gmOffset, blockCount, blockLen, srcStride);
285	- }
286	-
287	- __aicore__ inline void DoLastRowsCopy(int64_t localOffset, const UbLoopInfo& ubLoopInfo, const TensorInfo& tensorInfo)
288	- {
289	- // 0 无切分
290	- int64_t srcGmOffset = ubLoopInfo.rowStart * tensorInfo.tensorCol;
291	- uint32_t srcStride = (tensorInfo.tensorCol - tensorInfo.splitCol) * sizeof(T1);
292	- if (!tensorInfo.isSplit) {
293	- uint32_t blockLen = static_cast<uint32_t>(
294	- (tensorInfo.chunkDimSize * tensorInfo.originCol - ubLoopInfo.rowStart * tensorInfo.tensorCol) * sizeof(T1));
295	- ExecuteDataCopy(localOffset, srcGmOffset + tensorInfo.startOffset, 1, blockLen, srcStride);
296	- return;
297	- }
298	-
299	- uint16_t blockCount = 0;
300	- uint32_t blockLen = 0;
301	- int64_t remainderCol = (tensorInfo.chunkDimSize % tensorInfo.chunkCol) * tensorInfo.originCol;
302	- // 1 有切分
303	- // 1.0 remainder等于0
304	- if (remainderCol == 0) {
305	- blockCount = static_cast<uint16_t>(tensorInfo.chunkRow - ubLoopInfo.rowStart);
306	- blockLen = static_cast<uint32_t>(tensorInfo.splitCol * sizeof(T1));
307	- }
308	- // 1.1 切分+偏移值小于等于remainder
309	- else if (tensorInfo.startOffset + tensorInfo.splitCol <= remainderCol) {
310	- blockCount = static_cast<uint16_t>(tensorInfo.chunkRowAlign - ubLoopInfo.rowStart);
311	- blockLen = static_cast<uint32_t>(tensorInfo.splitCol * sizeof(T1));
312	- }
313	- // 1.2 偏移值大于等于remainder
314	- else if (tensorInfo.startOffset >= remainderCol) {
315	- blockCount = static_cast<uint16_t>(tensorInfo.chunkRow - ubLoopInfo.rowStart);
316	- blockLen = static_cast<uint32_t>(tensorInfo.splitCol * sizeof(T1));
317	- }
318	- // 1.3 偏移值小于remainder，且切分+偏移值大于remainder
319	- else {
320	- // 1.3.1
321	- blockLen = (remainderCol - tensorInfo.startOffset) * sizeof(T1);
322	- int64_t localOffsetPart = localOffset + (tensorInfo.chunkRow - ubLoopInfo.rowStart) * tensorInfo.splitColAlign;
323	- int64_t gmOffsetPart = srcGmOffset + tensorInfo.startOffset +
324	- (tensorInfo.chunkRow - ubLoopInfo.rowStart) * tensorInfo.tensorCol;
325	- ExecuteDataCopy(localOffsetPart, gmOffsetPart, 1, blockLen, srcStride);
326	- // 1.3.2
327	- blockCount = static_cast<uint16_t>(tensorInfo.chunkRow - ubLoopInfo.rowStart);
328	- blockLen = static_cast<uint32_t>(tensorInfo.splitCol * sizeof(T1));
329	- }
330	- ExecuteDataCopy(localOffset, srcGmOffset + tensorInfo.startOffset, blockCount, blockLen, srcStride);
331	- }
332	-
333	- __aicore__ inline void CopyInChunk(int64_t& totalCol, int64_t& localOffset, UbLoopInfo& ubLoopInfo, TensorInfo& tensorInfo)
334	- {
335	- // 获取concat阶段输入的col
336	- for (uint32_t j = 1; j < desc_.GetDim(); j++) {
337	- tensorInfo.originCol *= desc_.GetShape(j);
338	- }
339	- tensorInfo.tensorCol = tensorInfo.chunkCol * tensorInfo.originCol;
340	- // 判断当前核是否处理当前tensor
341	- if (!IsTensorInRange(totalCol, ubLoopInfo, tensorInfo)) {
342	- totalCol += tensorInfo.tensorCol;
343	- return;
344	- }
345	- SplitTensorDim0(totalCol, ubLoopInfo, tensorInfo);
346	- tensorInfo.chunkRow = tensorInfo.chunkDimSize / tensorInfo.chunkCol;
347	- tensorInfo.chunkRowAlign = GetAlign(tensorInfo.chunkDimSize, tensorInfo.chunkCol) / tensorInfo.chunkCol;
348	- int64_t localOffsetIncrement = (isOneConcat_ \|\| isAllAlign_) ? ubLoopInfo.currentUbRowFactor :
349	- (isHalfAlign_ ? TRANS_BLOCK * HALF : TRANS_BLOCK * srcEleUbBlock_);
350	- if (ubLoopInfo.rowStart >= tensorInfo.chunkRowAlign) {
351	- ubLoopInfo.inputCol[ubLoopInfo.count] = tensorInfo.splitCol;
352	- ubLoopInfo.totalUbColAlign += tensorInfo.splitCol;
353	- ubLoopInfo.totalUbCol += tensorInfo.splitCol;
354	- ubLoopInfo.count++;
355	- totalCol += tensorInfo.tensorCol;
356	- localOffset += localOffsetIncrement * tensorInfo.splitCol;
357	- return;
358	- }
359	- ubLoopInfo.isAllZero = false;
360	-
361	- int64_t rowEnd = ubLoopInfo.rowStart + ubLoopInfo.currentUbRowFactor;
362	- if (rowEnd < tensorInfo.chunkRowAlign) {
363	- DoRowsCopy(localOffset, ubLoopInfo, tensorInfo);
364	- } else {
365	- DoLastRowsCopy(localOffset, ubLoopInfo, tensorInfo);
366	- }
367	-
368	- localOffsetIncrement *= (isOneConcat_ \|\| tensorInfo.isSplit) ? tensorInfo.splitColAlign : tensorInfo.splitCol;
369	- if (isOneConcat_) {
370	- ubLoopInfo.inputCol[ubLoopInfo.count] = tensorInfo.splitCol;
371	- ubLoopInfo.totalUbColAlign += tensorInfo.splitColAlign;
372	- } else if (isAllAlign_) {
373	- ubLoopInfo.inputCol[ubLoopInfo.count] = tensorInfo.splitCol;
374	- ubLoopInfo.totalUbColAlign += tensorInfo.splitCol;
375	- } else if (tensorInfo.isSplit) {
376	- ubLoopInfo.inputCol[ubLoopInfo.count] = -tensorInfo.splitCol;
377	- ubLoopInfo.totalUbColAlign += tensorInfo.splitColAlign;
378	- } else {
379	- ubLoopInfo.inputCol[ubLoopInfo.count] = tensorInfo.splitCol;
380	- ubLoopInfo.totalUbColAlign += tensorInfo.splitCol;
381	- }
382	- ubLoopInfo.totalUbCol += tensorInfo.splitCol;
383	- ubLoopInfo.count++;
384	- totalCol += tensorInfo.tensorCol;
385	- localOffset += localOffsetIncrement;
386	}	229	}
387		230
388	__aicore__ inline void UBRearrange4Trans(const UbLoopInfo& ubLoopInfo, LocalTensor<T1>& srcLocal, LocalTensor<T1>& dstLocal)	231	__aicore__ inline void UBRearrange4Trans(const UbLoopInfo& ubLoopInfo, LocalTensor<T1>& srcLocal, LocalTensor<T1>& dstLocal)
	@@ -392,20 +235,20 @@ private:
392	for (int64_t i = 0; i < ubLoopInfo.count; i++) {	235	for (int64_t i = 0; i < ubLoopInfo.count; i++) {
393	uint16_t blockCount = TRANS_BLOCK;	236	uint16_t blockCount = TRANS_BLOCK;
394	uint16_t actualCol = ubLoopInfo.inputCol[i] > 0 ? ubLoopInfo.inputCol[i] :	237	uint16_t actualCol = ubLoopInfo.inputCol[i] > 0 ? ubLoopInfo.inputCol[i] :
395	- GetAlign(-ubLoopInfo.inputCol[i], srcEleUbBlock_);	238	+ this->GetAlign(-ubLoopInfo.inputCol[i], this->srcEleUbBlock_);
396	- uint16_t blockLen = actualCol * colRepeatNum_ / srcEleUbBlock_;	239	+ uint16_t blockLen = actualCol * colRepeatNum_ / this->srcEleUbBlock_;
397	- uint16_t dstGap = ubLoopInfo.totalUbColAlign * colRepeatNum_ / srcEleUbBlock_ - blockLen;	240	+ uint16_t dstGap = ubLoopInfo.totalUbColAlign * colRepeatNum_ / this->srcEleUbBlock_ - blockLen;
398	DataCopyParams copyParams{blockCount, blockLen, 0, dstGap};	241	DataCopyParams copyParams{blockCount, blockLen, 0, dstGap};
399	DataCopy(dstLocal[dstOffset], srcLocal[srcOffset], copyParams);	242	DataCopy(dstLocal[dstOffset], srcLocal[srcOffset], copyParams);
400	- srcOffset += blockCount * blockLen * srcEleUbBlock_;	243	+ srcOffset += blockCount * blockLen * this->srcEleUbBlock_;
401	- dstOffset += blockLen * srcEleUbBlock_;	244	+ dstOffset += blockLen * this->srcEleUbBlock_;
402	}	245	}
403	PipeBarrier<PIPE_V>();	246	PipeBarrier<PIPE_V>();
404	}	247	}
405		248
406	__aicore__ inline void Trans1(const UbLoopInfo& ubLoopInfo, LocalTensor<T1>& srcLocal, LocalTensor<T1>& dstLocal)	249	__aicore__ inline void Trans1(const UbLoopInfo& ubLoopInfo, LocalTensor<T1>& srcLocal, LocalTensor<T1>& dstLocal)
407	{	250	{
408	- uint8_t repeatTimes = ubLoopInfo.totalUbColAlign * colRepeatNum_ / srcEleUbBlock_;	251	+ uint8_t repeatTimes = ubLoopInfo.totalUbColAlign * colRepeatNum_ / this->srcEleUbBlock_;
409	uint16_t srcRepStride = repeatTimes == 1 ? 0 : 1;	252	uint16_t srcRepStride = repeatTimes == 1 ? 0 : 1;
410	uint16_t dstRepStride = repeatTimes == 1 ? 0 : TRANS_BLOCK;	253	uint16_t dstRepStride = repeatTimes == 1 ? 0 : TRANS_BLOCK;
411	TransDataTo5HDParams transDataParams{false, false, repeatTimes, dstRepStride, srcRepStride};	254	TransDataTo5HDParams transDataParams{false, false, repeatTimes, dstRepStride, srcRepStride};
	@@ -429,10 +272,10 @@ private:
429	uint64_t offset = i * ubLoopInfo.totalUbColAlign * colRepeatNum_;	272	uint64_t offset = i * ubLoopInfo.totalUbColAlign * colRepeatNum_;
430	srcLocalList[i] = reinterpret_cast<uint64_t>(srcLocal[offset].GetPhyAddr());	273	srcLocalList[i] = reinterpret_cast<uint64_t>(srcLocal[offset].GetPhyAddr());
431	}	274	}
432	- for (uint64_t i = 0; i < srcEleUbBlock_; i++) {	275	+ for (uint64_t i = 0; i < this->srcEleUbBlock_; i++) {
433	- for (uint64_t j = 0; j < TRANS_BLOCK / srcEleUbBlock_; j++) {	276	+ for (uint64_t j = 0; j < TRANS_BLOCK / this->srcEleUbBlock_; j++) {
434	- uint64_t offset = i * TRANS_BLOCK + j * srcEleUbBlock_;	277	+ uint64_t offset = i * TRANS_BLOCK + j * this->srcEleUbBlock_;
435	- dstLocalList[i * TRANS_BLOCK / srcEleUbBlock_ + j] =	278	+ dstLocalList[i * TRANS_BLOCK / this->srcEleUbBlock_ + j] =
436	reinterpret_cast<uint64_t>(dstLocal[offset].GetPhyAddr());	279	reinterpret_cast<uint64_t>(dstLocal[offset].GetPhyAddr());
437	}	280	}
438	}	281	}
	@@ -449,17 +292,17 @@ private:
449	for (int64_t i = 0; i < ubLoopInfo.count; i++) {	292	for (int64_t i = 0; i < ubLoopInfo.count; i++) {
450	uint16_t blockCount = colRepeatNum_;	293	uint16_t blockCount = colRepeatNum_;
451	uint16_t actualCol = ubLoopInfo.inputCol[i] > 0 ? ubLoopInfo.inputCol[i] : -ubLoopInfo.inputCol[i];	294	uint16_t actualCol = ubLoopInfo.inputCol[i] > 0 ? ubLoopInfo.inputCol[i] : -ubLoopInfo.inputCol[i];
452	- uint16_t blockLen = actualCol * TRANS_BLOCK / srcEleUbBlock_;	295	+ uint16_t blockLen = actualCol * TRANS_BLOCK / this->srcEleUbBlock_;
453	uint16_t srcGap = ubLoopInfo.inputCol[i] > 0 ? 0 :	296	uint16_t srcGap = ubLoopInfo.inputCol[i] > 0 ? 0 :
454	- (GetAlign(-ubLoopInfo.inputCol[i], srcEleUbBlock_) + ubLoopInfo.inputCol[i]) * TRANS_BLOCK / srcEleUbBlock_;	297	+ (this->GetAlign(-ubLoopInfo.inputCol[i], this->srcEleUbBlock_) + ubLoopInfo.inputCol[i]) * TRANS_BLOCK / this->srcEleUbBlock_;
455	- uint16_t dstGap = GetAlign(ubLoopInfo.totalUbCol, dstEleUbBlock_) * TRANS_BLOCK / srcEleUbBlock_ - blockLen;	298	+ uint16_t dstGap = this->GetAlign(ubLoopInfo.totalUbCol, this->dstEleUbBlock_) * TRANS_BLOCK / this->srcEleUbBlock_ - blockLen;
456	if constexpr (NO_NEED_ALIGN) {	299	if constexpr (NO_NEED_ALIGN) {
457	- dstGap = ubLoopInfo.totalUbCol * TRANS_BLOCK / srcEleUbBlock_ - blockLen;	300	+ dstGap = ubLoopInfo.totalUbCol * TRANS_BLOCK / this->srcEleUbBlock_ - blockLen;
458	}	301	}
459	DataCopyParams copyParams{blockCount, blockLen, srcGap, dstGap};	302	DataCopyParams copyParams{blockCount, blockLen, srcGap, dstGap};
460	DataCopy(dstLocal[dstOffset], srcLocal[srcOffset], copyParams);	303	DataCopy(dstLocal[dstOffset], srcLocal[srcOffset], copyParams);
461	srcOffset += ubLoopInfo.inputCol[i] > 0 ? (colRepeatNum_ * actualCol * TRANS_BLOCK) :	304	srcOffset += ubLoopInfo.inputCol[i] > 0 ? (colRepeatNum_ * actualCol * TRANS_BLOCK) :
462	- (colRepeatNum_ * GetAlign(actualCol, srcEleUbBlock_) * TRANS_BLOCK);	305	+ (colRepeatNum_ * this->GetAlign(actualCol, this->srcEleUbBlock_) * TRANS_BLOCK);
463	dstOffset += (actualCol * TRANS_BLOCK);	306	dstOffset += (actualCol * TRANS_BLOCK);
464	}	307	}
465	PipeBarrier<PIPE_V>();	308	PipeBarrier<PIPE_V>();
	@@ -468,11 +311,11 @@ private:
468	template <bool NO_NEED_ALIGN=false>	311	template <bool NO_NEED_ALIGN=false>
469	__aicore__ inline void Trans2(const UbLoopInfo& ubLoopInfo, LocalTensor<T1>& srcLocal, LocalTensor<T1>& dstLocal)	312	__aicore__ inline void Trans2(const UbLoopInfo& ubLoopInfo, LocalTensor<T1>& srcLocal, LocalTensor<T1>& dstLocal)
470	{	313	{
471	- int64_t actualTotalUbCol = GetAlign(ubLoopInfo.totalUbCol, dstEleUbBlock_);	314	+ int64_t actualTotalUbCol = this->GetAlign(ubLoopInfo.totalUbCol, this->dstEleUbBlock_);
472	if constexpr (NO_NEED_ALIGN) {	315	if constexpr (NO_NEED_ALIGN) {
473	actualTotalUbCol = ubLoopInfo.totalUbCol;	316	actualTotalUbCol = ubLoopInfo.totalUbCol;
474	}	317	}
475	- uint8_t repeatTimes = actualTotalUbCol * colRepeatNum_ / srcEleUbBlock_;	318	+ uint8_t repeatTimes = actualTotalUbCol * colRepeatNum_ / this->srcEleUbBlock_;
476	uint16_t srcRepStride = repeatTimes == 1 ? 0 : TRANS_BLOCK;	319	uint16_t srcRepStride = repeatTimes == 1 ? 0 : TRANS_BLOCK;
477	uint16_t dstRepStride = repeatTimes == 1 ? 0 : 1;	320	uint16_t dstRepStride = repeatTimes == 1 ? 0 : 1;
478	TransDataTo5HDParams transDataParams = {false, false, repeatTimes, dstRepStride, srcRepStride};	321	TransDataTo5HDParams transDataParams = {false, false, repeatTimes, dstRepStride, srcRepStride};
	@@ -491,10 +334,10 @@ private:
491	}	334	}
492	TransDataTo5HD<half>(dstLocalList, srcLocalList, transDataParams);	335	TransDataTo5HD<half>(dstLocalList, srcLocalList, transDataParams);
493	} else {	336	} else {
494	- for (uint64_t i = 0; i < TRANS_BLOCK / srcEleUbBlock_; i++) {	337	+ for (uint64_t i = 0; i < TRANS_BLOCK / this->srcEleUbBlock_; i++) {
495	- for (uint64_t j = 0; j < srcEleUbBlock_; j++) {	338	+ for (uint64_t j = 0; j < this->srcEleUbBlock_; j++) {
496	- uint64_t offset = i * srcEleUbBlock_ + j * TRANS_BLOCK;	339	+ uint64_t offset = i * this->srcEleUbBlock_ + j * TRANS_BLOCK;
497	- srcLocalList[i * srcEleUbBlock_ + j] =	340	+ srcLocalList[i * this->srcEleUbBlock_ + j] =
498	reinterpret_cast<uint64_t>(srcLocal[offset].GetPhyAddr());	341	reinterpret_cast<uint64_t>(srcLocal[offset].GetPhyAddr());
499	}	342	}
500	}	343	}
	@@ -503,7 +346,7 @@ private:
503	dstLocalList[i] = reinterpret_cast<uint64_t>(dstLocal[offset].GetPhyAddr());	346	dstLocalList[i] = reinterpret_cast<uint64_t>(dstLocal[offset].GetPhyAddr());
504	}	347	}
505	for (uint64_t i = 1; i < TRANS_BLOCK; i += 2) { // 2 is stride	348	for (uint64_t i = 1; i < TRANS_BLOCK; i += 2) { // 2 is stride
506	- uint64_t offset = (i / 2 + srcEleUbBlock_) * actualTotalUbCol * colRepeatNum_; // 2 is stride	349	+ uint64_t offset = (i / 2 + this->srcEleUbBlock_) * actualTotalUbCol * colRepeatNum_; // 2 is stride
507	dstLocalList[i] = reinterpret_cast<uint64_t>(dstLocal[offset].GetPhyAddr());	350	dstLocalList[i] = reinterpret_cast<uint64_t>(dstLocal[offset].GetPhyAddr());
508	}	351	}
509	TransDataTo5HD<T1>(dstLocalList, srcLocalList, transDataParams);	352	TransDataTo5HD<T1>(dstLocalList, srcLocalList, transDataParams);
	@@ -518,9 +361,9 @@ private:
518	int64_t dstOffset = 0;	361	int64_t dstOffset = 0;
519	for (int64_t i = 0; i < ubLoopInfo.count; i++) {	362	for (int64_t i = 0; i < ubLoopInfo.count; i++) {
520	uint16_t blockCount = ubLoopInfo.currentUbRowFactor;	363	uint16_t blockCount = ubLoopInfo.currentUbRowFactor;
521	- uint16_t blockLen = ubLoopInfo.inputCol[i] / srcEleUbBlock_;	364	+ uint16_t blockLen = ubLoopInfo.inputCol[i] / this->srcEleUbBlock_;
522	uint16_t srcGap = 0;	365	uint16_t srcGap = 0;
523	- uint16_t dstGap = (ubLoopInfo.totalUbColAlign - ubLoopInfo.inputCol[i]) / srcEleUbBlock_;	366	+ uint16_t dstGap = (ubLoopInfo.totalUbColAlign - ubLoopInfo.inputCol[i]) / this->srcEleUbBlock_;
524	DataCopyParams copyParams{blockCount, blockLen, srcGap, dstGap};	367	DataCopyParams copyParams{blockCount, blockLen, srcGap, dstGap};
525	DataCopy(dstLocal[dstOffset], srcLocal[srcOffset], copyParams);	368	DataCopy(dstLocal[dstOffset], srcLocal[srcOffset], copyParams);
526	srcOffset += blockCount * ubLoopInfo.inputCol[i];	369	srcOffset += blockCount * ubLoopInfo.inputCol[i];
	@@ -532,36 +375,36 @@ private:
532	__aicore__ inline void DoCast(const UbLoopInfo& ubLoopInfo, uint32_t castCount)	375	__aicore__ inline void DoCast(const UbLoopInfo& ubLoopInfo, uint32_t castCount)
533	{	376	{
534	if constexpr (sizeof(T1) == sizeof(T2)) {	377	if constexpr (sizeof(T1) == sizeof(T2)) {
535	- Cast(dstLocalFP32_, srcLocal_, RoundMode::CAST_NONE, castCount);	378	+ Cast(dstLocalFP32_, this->srcLocal_, RoundMode::CAST_NONE, castCount);
536	Cast(srcLocalT2_, dstLocalFP32_, RoundMode::CAST_RINT, castCount);	379	Cast(srcLocalT2_, dstLocalFP32_, RoundMode::CAST_RINT, castCount);
537	- DataCopy(dstLocal_, srcLocalT2_, castCount);	380	+ DataCopy(this->dstLocal_, srcLocalT2_, castCount);
538	} else {	381	} else {
539	- Cast(dstLocal_, srcLocal_, RoundMode::CAST_NONE, castCount);	382	+ Cast(this->dstLocal_, this->srcLocal_, RoundMode::CAST_NONE, castCount);
540	}	383	}
541	}	384	}
542		385
543	__aicore__ inline void ComputeOneConcat(const UbLoopInfo& ubLoopInfo)	386	__aicore__ inline void ComputeOneConcat(const UbLoopInfo& ubLoopInfo)
544	{	387	{
545	- SetFlag<HardEvent::MTE2_V>(event_);	388	+ SetFlag<HardEvent::MTE2_V>(this->event_);
546	- WaitFlag<HardEvent::MTE2_V>(event_);	389	+ WaitFlag<HardEvent::MTE2_V>(this->event_);
547	if (ubLoopInfo.totalUbColAlign == 0) {	390	if (ubLoopInfo.totalUbColAlign == 0) {
548	return;	391	return;
549	}	392	}
550	- if constexpr (NEAD_CAST) {	393	+ if constexpr (NEED_CAST) {
551	uint32_t castCount = ubLoopInfo.currentUbRowFactor * ubLoopInfo.totalUbColAlign;	394	uint32_t castCount = ubLoopInfo.currentUbRowFactor * ubLoopInfo.totalUbColAlign;
552	DoCast(ubLoopInfo, castCount);	395	DoCast(ubLoopInfo, castCount);
553	} else {	396	} else {
554	- DataCopy(dstLocalT1_, srcLocal_, ubLoopInfo.totalUbColAlign);	397	+ DataCopy(dstLocalT1_, this->srcLocal_, ubLoopInfo.totalUbColAlign);
555	}	398	}
556	}	399	}
557		400
558	__aicore__ inline void ComputeAllAlign(const UbLoopInfo& ubLoopInfo)	401	__aicore__ inline void ComputeAllAlign(const UbLoopInfo& ubLoopInfo)
559	{	402	{
560	- SetFlag<HardEvent::MTE2_V>(event_);	403	+ SetFlag<HardEvent::MTE2_V>(this->event_);
561	- WaitFlag<HardEvent::MTE2_V>(event_);	404	+ WaitFlag<HardEvent::MTE2_V>(this->event_);
562	- UBRearrange4Concat(ubLoopInfo, srcLocal_, dstLocalT1_);	405	+ UBRearrange4Concat(ubLoopInfo, this->srcLocal_, dstLocalT1_);
563	- if constexpr (NEAD_CAST) {	406	+ if constexpr (NEED_CAST) {
564	- DataCopy(srcLocal_, dstLocalT1_, ubLoopInfo.currentUbRowFactor * ubLoopInfo.currentUbColFactor);	407	+ DataCopy(this->srcLocal_, dstLocalT1_, ubLoopInfo.currentUbRowFactor * ubLoopInfo.currentUbColFactor);
565	PipeBarrier<PIPE_V>();	408	PipeBarrier<PIPE_V>();
566	uint32_t castCount = ubLoopInfo.currentUbRowFactor * ubLoopInfo.currentUbColFactor;	409	uint32_t castCount = ubLoopInfo.currentUbRowFactor * ubLoopInfo.currentUbColFactor;
567	DoCast(ubLoopInfo, castCount);	410	DoCast(ubLoopInfo, castCount);
	@@ -570,23 +413,23 @@ private:
570		413
571	__aicore__ inline void ComputeNotAlign(const UbLoopInfo& ubLoopInfo)	414	__aicore__ inline void ComputeNotAlign(const UbLoopInfo& ubLoopInfo)
572	{	415	{
573	- SetFlag<HardEvent::MTE2_V>(event_);	416	+ SetFlag<HardEvent::MTE2_V>(this->event_);
574	- WaitFlag<HardEvent::MTE2_V>(event_);	417	+ WaitFlag<HardEvent::MTE2_V>(this->event_);
575	// 3、ub重排	418	// 3、ub重排
576	- UBRearrange4Trans(ubLoopInfo, srcLocal_, dstLocalT1_);	419	+ UBRearrange4Trans(ubLoopInfo, this->srcLocal_, dstLocalT1_);
577	// 4、跨block对齐转置	420	// 4、跨block对齐转置
578	- Trans1(ubLoopInfo, dstLocalT1_, srcLocal_);	421	+ Trans1(ubLoopInfo, dstLocalT1_, this->srcLocal_);
579	// 5、ub重排	422	// 5、ub重排
580	- UBRearrange4TransConcat(ubLoopInfo, srcLocal_, dstLocalT1_);	423	+ UBRearrange4TransConcat(ubLoopInfo, this->srcLocal_, dstLocalT1_);
581	// 6、跨block对齐转置	424	// 6、跨block对齐转置
582	- Trans2(ubLoopInfo, dstLocalT1_, srcLocal_);	425	+ Trans2(ubLoopInfo, dstLocalT1_, this->srcLocal_);
583		426
584	// 7、cast or ubToub	427	// 7、cast or ubToub
585	- if constexpr (NEAD_CAST) {	428	+ if constexpr (NEED_CAST) {
586	- uint32_t castCount = ubLoopInfo.currentUbRowFactor * GetAlign(ubLoopInfo.currentUbColFactor, dstEleUbBlock_);	429	+ uint32_t castCount = ubLoopInfo.currentUbRowFactor * this->GetAlign(ubLoopInfo.currentUbColFactor, this->dstEleUbBlock_);
587	DoCast(ubLoopInfo, castCount);	430	DoCast(ubLoopInfo, castCount);
588	} else {	431	} else {
589	- DataCopy(dstLocalT1_, srcLocal_, ubLoopInfo.currentUbRowFactor * GetAlign(ubLoopInfo.currentUbColFactor, dstEleUbBlock_));	432	+ DataCopy(dstLocalT1_, this->srcLocal_, ubLoopInfo.currentUbRowFactor * this->GetAlign(ubLoopInfo.currentUbColFactor, this->dstEleUbBlock_));
590	}	433	}
591	}	434	}
592		435
	@@ -594,38 +437,10 @@ private:
594	bool isAllAlign_{false};	437	bool isAllAlign_{false};
595	bool isHalfAlign_{false};	438	bool isHalfAlign_{false};
596	bool isOneConcat_{false};	439	bool isOneConcat_{false};
597	- int64_t blockIdx_{0};
598	- int64_t inputNum_{0};
599	- int64_t ubRowFactor_{0};
600	- int64_t ubColFactor_{0};
601	- int64_t srcEleUbBlock_{0};
602	- int64_t dstEleUbBlock_{0};
603	- int64_t dim_{0};
604	- int64_t numChunk_{0};
605	- int64_t outputRow_{0};
606	- int64_t outputCol_{0};
607	- int64_t blockRowFactor_{0};
608	- int64_t blockColFactor_{0};
609	- int64_t tailBlockRowFactor_{0};
610	- int64_t tailBlockColFactor_{0};
611	- int64_t blockRowGroup_{0};
612	- int64_t blockColGroup_{0};
613	- int64_t currentBlockRowFactor_{0};
614	- int64_t currentBlockColFactor_{0};
615	int64_t colRepeatNum_{0};	440	int64_t colRepeatNum_{0};
616		441
617	- TPipe *pipe_;
618	- TEventID event_{0};
619	- TensorDesc<T1> desc_;
620	- ListTensorDesc inputList_;
621	- GlobalTensor<T2> dstGlobal_;
622	- GlobalTensor<T1> srcGlobal_;
623	- TBuf<AscendC::TPosition::VECCALC> srcBuf_;
624	- TBuf<AscendC::TPosition::VECCALC> dstBuf_;
625	- LocalTensor<T1> srcLocal_;
626	LocalTensor<T2> srcLocalT2_;	442	LocalTensor<T2> srcLocalT2_;
627	- LocalTensor<T2> dstLocal_;
628	LocalTensor<T1> dstLocalT1_;	443	LocalTensor<T1> dstLocalT1_;
629	LocalTensor<float> dstLocalFP32_;	444	LocalTensor<float> dstLocalFP32_;
630	};	445	};
631	-#endif // _CHUNK_CAT_DATA_H_	446	+#endif // _CHUNK_CAT_H_

Aconversion/chunk_cat/op_kernel/chunk_cat_arch35.h+293-0

@@ -0,0 +1,293 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+ /*!
	12	+ * \file chunk_cat_arch35.h
	13	+ * \brief
	14	+ */
	15	+
	16	+#ifndef _CHUNK_CAT_ARCH35_H_
	17	+#define _CHUNK_CAT_ARCH35_H_
	18	+
	19	+#include "chunk_cat_common.h"
	20	+#include "op_kernel/platform_util.h"
	21	+
	22	+using namespace AscendC;
	23	+using namespace Ops::Base;
	24	+
	25	+static constexpr MicroAPI::CastTrait castTraitZero = {
	26	+ MicroAPI::RegLayout::ZERO, MicroAPI::SatMode::UNKNOWN, MicroAPI::MaskMergeMode::ZEROING,
	27	+ RoundMode::UNKNOWN
	28	+};
	29	+static constexpr MicroAPI::CastTrait castTraitOne = {
	30	+ MicroAPI::RegLayout::ONE, MicroAPI::SatMode::UNKNOWN, MicroAPI::MaskMergeMode::ZEROING,
	31	+ RoundMode::UNKNOWN
	32	+};
	33	+static constexpr MicroAPI::CastTrait castTraitbf2half = {
	34	+ MicroAPI::RegLayout::UNKNOWN, MicroAPI::SatMode::NO_SAT, MicroAPI::MaskMergeMode::ZEROING,
	35	+ RoundMode::CAST_RINT
	36	+};
	37	+static constexpr MicroAPI::CastTrait castTraithalf2bf = {
	38	+ MicroAPI::RegLayout::UNKNOWN, MicroAPI::SatMode::UNKNOWN, MicroAPI::MaskMergeMode::ZEROING,
	39	+ RoundMode::CAST_RINT
	40	+};
	41	+
	42	+template <typename T1, typename T2, bool NEED_CAST = false>
	43	+class ChunkCatArch35 : public ChunkCatCommon<T1, T2>
	44	+{
	45	+public:
	46	+ __aicore__ inline ChunkCatArch35(TPipe *pipe) : ChunkCatCommon<T1, T2>(pipe) {}
	47	+
	48	+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, const ChunkCatTilingData& tilingData)
	49	+ {
	50	+ this->InitCommon(x, y, tilingData);
	51	+ }
	52	+
	53	+ __aicore__ inline void Process()
	54	+ {
	55	+ int64_t rowLoop = this->GetAlign(this->currentBlockRowFactor_, this->ubRowFactor_) / this->ubRowFactor_;
	56	+ int64_t colLoop = this->GetAlign(this->currentBlockColFactor_, this->ubColFactor_) / this->ubColFactor_;
	57	+ int64_t rowTail = this->currentBlockRowFactor_ % this->ubRowFactor_;
	58	+ int64_t colTail = this->currentBlockColFactor_ % this->ubColFactor_;
	59	+
	60	+ uint64_t buf[10];
	61	+ this->desc_.SetShapeAddr(buf); // 用于获取shape信息
	62	+
	63	+ for (int64_t i = 0; i < rowLoop * colLoop; i++) {
	64	+ UbLoopInfo ubLoopInfo{};
	65	+ ubLoopInfo.ubRowGroup = i / colLoop;
	66	+ ubLoopInfo.ubColGroup = i % colLoop;
	67	+ ubLoopInfo.currentUbRowFactor = (rowTail != 0 && ubLoopInfo.ubRowGroup == rowLoop - 1) ?
	68	+ rowTail : this->ubRowFactor_;
	69	+ ubLoopInfo.currentUbColFactor = (colTail != 0 && ubLoopInfo.ubColGroup == colLoop - 1) ?
	70	+ colTail : this->ubColFactor_;
	71	+ // 搬入&&计算
	72	+ CopyInAndCompute(ubLoopInfo);
	73	+ SetFlag<HardEvent::V_MTE3>(this->event_);
	74	+ WaitFlag<HardEvent::V_MTE3>(this->event_);
	75	+ // 搬出
	76	+ CopyOut(ubLoopInfo);
	77	+ }
	78	+ }
	79	+
	80	+private:
	81	+ __aicore__ inline void CopyInAndCompute(UbLoopInfo& ubLoopInfo)
	82	+ {
	83	+ int64_t localOffset = 0;
	84	+ int64_t totalCol = 0;
	85	+ ubLoopInfo.colStart = this->blockColGroup_ * this->blockColFactor_ + ubLoopInfo.ubColGroup * this->ubColFactor_;
	86	+ ubLoopInfo.rowStart = this->blockRowGroup_ * this->blockRowFactor_ + ubLoopInfo.ubRowGroup * this->ubRowFactor_;
	87	+ for (uint32_t i = 0; i < this->inputNum_; i++) {
	88	+ if (ubLoopInfo.totalUbCol >= ubLoopInfo.currentUbColFactor) {
	89	+ break;
	90	+ }
	91	+ TensorInfo tensorInfo{};
	92	+ this->srcGlobal_.SetGlobalBuffer(this->inputList_.template GetDataPtr<T1>(i));
	93	+ this->GetChunkInfo(i, tensorInfo);
	94	+ if (!this->IsTensorInRange(totalCol, ubLoopInfo, tensorInfo)) {
	95	+ totalCol += tensorInfo.tensorCol;
	96	+ } else {
	97	+ tensorInfo.chunkRow = tensorInfo.chunkDimSize / tensorInfo.chunkCol;
	98	+ tensorInfo.chunkRowAlign = this->GetAlign(tensorInfo.chunkDimSize, tensorInfo.chunkCol) / tensorInfo.chunkCol;
	99	+ this->SplitTensorDim0(totalCol, ubLoopInfo, tensorInfo);
	100	+ if (ubLoopInfo.rowStart >= tensorInfo.chunkRowAlign) {
	101	+ tensorInfo.isZero = true;
	102	+ } else {
	103	+ this->CopyInChunk(totalCol, localOffset, ubLoopInfo, tensorInfo);
	104	+ }
	105	+ // 计算
	106	+ Compute(ubLoopInfo, tensorInfo, totalCol, localOffset);
	107	+ ubLoopInfo.count++;
	108	+ ubLoopInfo.totalUbCol += tensorInfo.splitCol;
	109	+ totalCol += tensorInfo.tensorCol;
	110	+ localOffset += this->GetAlign(ubLoopInfo.currentUbRowFactor * tensorInfo.splitCol, this->srcEleUbBlock_);
	111	+ }
	112	+
	113	+ if (ubLoopInfo.count > 31) {
	114	+ // 32个tensor处理一次
	115	+ CopyOut(ubLoopInfo);
	116	+ localOffset = 0;
	117	+ ubLoopInfo.preCatCol += ubLoopInfo.totalUbCol;
	118	+ }
	119	+ }
	120	+ }
	121	+
	122	+ __aicore__ inline void Compute(const UbLoopInfo& ubLoopInfo, const TensorInfo& tensorInfo, int64_t totalCol, int64_t localOffset)
	123	+ {
	124	+ if (!tensorInfo.isZero) {
	125	+ SetFlag<HardEvent::MTE2_V>(this->event_);
	126	+ WaitFlag<HardEvent::MTE2_V>(this->event_);
	127	+ }
	128	+ CopyCatVF(ubLoopInfo, tensorInfo, totalCol, localOffset);
	129	+ SetFlag<HardEvent::V_MTE3>(this->event_);
	130	+ WaitFlag<HardEvent::V_MTE3>(this->event_);
	131	+ }
	132	+
	133	+ __aicore__ inline void CopyOut(const UbLoopInfo& ubLoopInfo)
	134	+ {
	135	+ uint16_t blockCount = ubLoopInfo.currentUbRowFactor;
	136	+ uint32_t blockLen = ubLoopInfo.totalUbCol * sizeof(T2);
	137	+ uint32_t dstStride = (this->outputCol_ - ubLoopInfo.totalUbCol)* sizeof(T2);
	138	+ uint32_t srcStride = (this->GetAlign(ubLoopInfo.currentUbColFactor, this->dstEleUbBlock_) -
	139	+ this->GetAlign(ubLoopInfo.totalUbCol, this->dstEleUbBlock_)) / this->dstEleUbBlock_;
	140	+ DataCopyExtParams copyParamsOut{blockCount, blockLen, srcStride, dstStride, 0};
	141	+ int64_t dstOffset = ubLoopInfo.ubRowGroup * this->ubRowFactor_ * this->outputCol_ +
	142	+ ubLoopInfo.ubColGroup * this->ubColFactor_ + ubLoopInfo.preCatCol;
	143	+
	144	+ DataCopyPad(this->dstGlobal_[dstOffset], this->dstLocal_, copyParamsOut);
	145	+ SetFlag<HardEvent::MTE3_MTE2>(this->event_);
	146	+ WaitFlag<HardEvent::MTE3_MTE2>(this->event_);
	147	+ SetFlag<HardEvent::MTE3_V>(this->event_);
	148	+ WaitFlag<HardEvent::MTE3_V>(this->event_);
	149	+ }
	150	+ // Compute
	151	+ __aicore__ inline void DoCopyCatVF(__ubuf__ T2* dstAddr, __ubuf__ T1* srcAddr, uint16_t rowLoop, uint16_t colLoop,
	152	+ uint32_t tail, uint32_t rowStride)
	153	+ {
	154	+ uint32_t main = GetVRegSize() / sizeof(T1);
	155	+ uint32_t mainFP32 = GetVRegSize() / sizeof(T2);
	156	+ uint32_t tailFP32Fir = tail > mainFP32 ? mainFP32 : tail;
	157	+ uint32_t tailFP32Sec = tail - tailFP32Fir;
	158	+ AscendC::MicroAPI::UnalignReg u0;
	159	+ AscendC::MicroAPI::UnalignReg uReg;
	160	+ AscendC::MicroAPI::RegTensor<T1> srcReg0;
	161	+ AscendC::MicroAPI::RegTensor<T2> dstReg0;
	162	+ AscendC::MicroAPI::RegTensor<T2> dstReg1;
	163	+ AscendC::MicroAPI::MaskReg mask = AscendC::MicroAPI::CreateMask<T1, AscendC::MicroAPI:: MaskPattern::ALL>();
	164	+
	165	+ AscendC::MicroAPI::DataCopyUnAlignPre(u0, srcAddr);
	166	+ for (uint16_t i = 0; i < rowLoop; i++) {
	167	+ auto curDstAddr = dstAddr + i * rowStride;
	168	+ for (uint16_t j = 0; j < colLoop; j++) {
	169	+ AscendC::MicroAPI::DataCopyUnAlign(srcReg0, u0, srcAddr, main);
	170	+ if constexpr (std::is_same_v<T1, T2>) {
	171	+ AscendC::MicroAPI::DataCopyUnAlign(curDstAddr, srcReg0, uReg, main);
	172	+ } else if constexpr (std::is_same_v<T1, half> && std::is_same_v<T2, bfloat16_t>) {
	173	+ AscendC::MicroAPI::Cast<T2, T1, castTraithalf2bf>(dstReg0, srcReg0, mask);
	174	+ AscendC::MicroAPI::DataCopyUnAlign(curDstAddr, dstReg0, uReg, main);
	175	+ } else if constexpr (std::is_same_v<T1, bfloat16_t> && std::is_same_v<T2, half>) {
	176	+ AscendC::MicroAPI::Cast<T2, T1, castTraitbf2half>(dstReg0, srcReg0, mask);
	177	+ AscendC::MicroAPI::DataCopyUnAlign(curDstAddr, dstReg0, uReg, main);
	178	+ } else {
	179	+ AscendC::MicroAPI::Cast<T2, T1, castTraitZero>(dstReg0, srcReg0, mask);
	180	+ AscendC::MicroAPI::Cast<T2, T1, castTraitOne>(dstReg1, srcReg0, mask);
	181	+ AscendC::MicroAPI::Interleave(dstReg0, dstReg1, dstReg0, dstReg1);
	182	+ AscendC::MicroAPI::DataCopyUnAlign(curDstAddr, dstReg0, uReg, mainFP32);
	183	+ AscendC::MicroAPI::LocalMemBar<AscendC::MicroAPI::MemType::VEC_STORE, AscendC::MicroAPI::MemType::VEC_STORE>();
	184	+ AscendC::MicroAPI::DataCopyUnAlign(curDstAddr, dstReg1, uReg, mainFP32);
	185	+ }
	186	+ }
	187	+ AscendC::MicroAPI::DataCopyUnAlign(srcReg0, u0, srcAddr, tail);
	188	+ if constexpr (std::is_same_v<T1, T2>) {
	189	+ AscendC::MicroAPI::DataCopyUnAlign(curDstAddr, srcReg0, uReg, tail);
	190	+ } else if constexpr (std::is_same_v<T1, half> && std::is_same_v<T2, bfloat16_t>) {
	191	+ AscendC::MicroAPI::Cast<T2, T1, castTraithalf2bf>(dstReg0, srcReg0, mask);
	192	+ AscendC::MicroAPI::DataCopyUnAlign(curDstAddr, dstReg0, uReg, tail);
	193	+ } else if constexpr (std::is_same_v<T1, bfloat16_t> && std::is_same_v<T2, half>) {
	194	+ AscendC::MicroAPI::Cast<T2, T1, castTraitbf2half>(dstReg0, srcReg0, mask);
	195	+ AscendC::MicroAPI::DataCopyUnAlign(curDstAddr, dstReg0, uReg, tail);
	196	+ } else {
	197	+ AscendC::MicroAPI::Cast<T2, T1, castTraitZero>(dstReg0, srcReg0, mask);
	198	+ AscendC::MicroAPI::Cast<T2, T1, castTraitOne>(dstReg1, srcReg0, mask);
	199	+ AscendC::MicroAPI::Interleave(dstReg0, dstReg1, dstReg0, dstReg1);
	200	+ AscendC::MicroAPI::DataCopyUnAlign(curDstAddr, dstReg0, uReg, tailFP32Fir);
	201	+ AscendC::MicroAPI::LocalMemBar<AscendC::MicroAPI::MemType::VEC_STORE, AscendC::MicroAPI::MemType::VEC_STORE>();
	202	+ AscendC::MicroAPI::DataCopyUnAlign(curDstAddr, dstReg1, uReg, tailFP32Sec);
	203	+ }
	204	+ AscendC::MicroAPI::DataCopyUnAlignPost(curDstAddr, uReg, 0);
	205	+ }
	206	+ }
	207	+
	208	+ __aicore__ inline void DoPadCatVF(__ubuf__ T2* dstAddr, uint16_t rowLoop, uint16_t colLoop,
	209	+ uint32_t main, uint32_t tail, uint32_t rowStride)
	210	+ {
	211	+ AscendC::MicroAPI::UnalignReg uReg;
	212	+ AscendC::MicroAPI::RegTensor<T2> dstReg0;
	213	+
	214	+ // // 纯pad
	215	+ T2 scalarValue = 0;
	216	+ AscendC::MicroAPI::Duplicate(dstReg0, scalarValue);
	217	+ for (uint16_t i = 0; i < rowLoop; i++) {
	218	+ auto curDstAddr = dstAddr + i * rowStride;
	219	+ for (uint16_t j = 0; j < colLoop; j++) {
	220	+ AscendC::MicroAPI::DataCopyUnAlign(curDstAddr, dstReg0, uReg, main);
	221	+ }
	222	+ AscendC::MicroAPI::DataCopyUnAlign(curDstAddr, dstReg0, uReg, tail);
	223	+ AscendC::MicroAPI::DataCopyUnAlignPost(curDstAddr, uReg, 0);
	224	+ }
	225	+ }
	226	+
	227	+ __aicore__ inline void CopyCatVF(const UbLoopInfo& ubLoopInfo, const TensorInfo& tensorInfo, int64_t totalCol, int64_t localOffset)
	228	+ {
	229	+ uint32_t srcLen = tensorInfo.splitCol;
	230	+ bool isSplit = tensorInfo.isSplit;
	231	+
	232	+ // 无搬运&&完整搬运阶段参数
	233	+ uint32_t main = GetVRegSize() / sizeof(T1);
	234	+ uint32_t mainFP32 = GetVRegSize() / sizeof(T2);
	235	+ uint16_t colLoop = srcLen / main; // 一行需要几次循环
	236	+ uint32_t tail = srcLen - colLoop * main;
	237	+ // fp16/bf16 -> fp32
	238	+ uint16_t colLoopFP32 = srcLen / mainFP32; // 一行需要几次循环
	239	+ uint32_t tailFP32 = srcLen - colLoopFP32 * mainFP32;
	240	+ uint32_t tailFP32Fir = tail > mainFP32 ? mainFP32 : tail;
	241	+ uint32_t tailFP32Sec = tail - tailFP32Fir;
	242	+ uint32_t rowStride = this->GetAlign(ubLoopInfo.currentUbColFactor, this->dstEleUbBlock_);
	243	+
	244	+ // 三个部分的行数
	245	+ uint16_t rowLoop0 = 0;
	246	+ uint16_t rowLoop1 = 0;
	247	+ uint16_t rowLoop2 = 0;
	248	+ // 部分pad情况下pad部分
	249	+ uint16_t colLen0 = 0;
	250	+ uint16_t colLen1 = 0;
	251	+ uint16_t colLoop1 = 0;
	252	+ uint32_t tailPad = 0;
	253	+
	254	+ if (ubLoopInfo.rowStart >= tensorInfo.chunkRowAlign) {
	255	+ // 无搬运，纯pad
	256	+ rowLoop2 = ubLoopInfo.currentUbRowFactor;
	257	+ } else if (ubLoopInfo.rowStart + ubLoopInfo.currentUbRowFactor < tensorInfo.chunkRowAlign) {
	258	+ // 完整搬运，无需pad
	259	+ rowLoop0 = ubLoopInfo.currentUbRowFactor;
	260	+ } else {
	261	+ // 部分搬运，需要pad
	262	+ rowLoop0 = tensorInfo.chunkRow - ubLoopInfo.rowStart;
	263	+ rowLoop1 = tensorInfo.chunkRowAlign - tensorInfo.chunkRow;
	264	+ rowLoop2 = ubLoopInfo.rowStart + ubLoopInfo.currentUbRowFactor - tensorInfo.chunkRowAlign;
	265	+ if (rowLoop1 > 0) {
	266	+ uint32_t remainderCol = (tensorInfo.chunkDimSize % tensorInfo.chunkCol) * tensorInfo.originCol;
	267	+ uint32_t blockLen = remainderCol > tensorInfo.startOffset ? remainderCol - tensorInfo.startOffset : 0;
	268	+ colLen0 = blockLen > srcLen ? srcLen : blockLen;
	269	+ colLen1 = srcLen - colLen0;
	270	+ colLoop1 = colLen1 / mainFP32;
	271	+ tailPad = colLen1 - colLoop1 * mainFP32;
	272	+ }
	273	+ }
	274	+ uint32_t padLen = isSplit ? this->GetAlign(srcLen * rowLoop0, this->srcEleUbBlock_) - srcLen * rowLoop0 : 0;
	275	+ uint32_t dstOffset = totalCol + tensorInfo.startOffset - ubLoopInfo.colStart - ubLoopInfo.preCatCol;
	276	+
	277	+ auto dstAddr = (__ubuf__ T2*)this->dstLocal_.GetPhyAddr() + dstOffset;
	278	+ auto srcAddr = (__ubuf__ T1*)this->srcLocal_.GetPhyAddr() + localOffset;
	279	+
	280	+ __VEC_SCOPE__
	281	+ {
	282	+ // rowLoop0
	283	+ DoCopyCatVF(dstAddr, srcAddr, rowLoop0, colLoop, tail, rowStride);
	284	+ // rowLoop2
	285	+ DoPadCatVF(dstAddr + (rowLoop0 + rowLoop1) * rowStride, rowLoop2, colLoopFP32, mainFP32, tailFP32, rowStride);
	286	+ // rowLoop1
	287	+ DoCopyCatVF(dstAddr + rowLoop0 * rowStride, srcAddr + rowLoop0 * srcLen + padLen, rowLoop1, colLoop, tail, rowStride);
	288	+ AscendC::MicroAPI::LocalMemBar<AscendC::MicroAPI::MemType::VEC_STORE, AscendC::MicroAPI::MemType::VEC_STORE>();
	289	+ DoPadCatVF(dstAddr + rowLoop0 * rowStride + colLen0, rowLoop1, colLoop1, mainFP32, tailPad, rowStride);
	290	+ }
	291	+ }
	292	+};
	293	+#endif // _CHUNK_CAT_ARCH35_H_

Mconversion/chunk_cat/op_kernel/chunk_cat_common.h+207-0

	@@ -26,6 +26,7 @@ constexpr uint32_t HALF = 2; // 半对齐/UB对半切分
26		26
27	struct TensorInfo {	27	struct TensorInfo {
28	bool isSplit{false};	28	bool isSplit{false};
		29	+ bool isZero{false};
29	int64_t chunkDimSize{0};	30	int64_t chunkDimSize{0};
30	int64_t chunkCol{0};	31	int64_t chunkCol{0};
31	int64_t chunkRow{0};	32	int64_t chunkRow{0};
	@@ -52,4 +53,210 @@ struct UbLoopInfo {
52	int64_t* inputCol;	53	int64_t* inputCol;
53	};	54	};
54		55
		56	+using namespace AscendC;
		57	+template <typename T1, typename T2>
		58	+class ChunkCatCommon
		59	+{
		60	+public:
		61	+ __aicore__ inline ChunkCatCommon(TPipe *pipe) : pipe_(pipe) {}
		62	+
		63	+ __aicore__ inline void InitCommon(GM_ADDR x, GM_ADDR y, const ChunkCatTilingData& tilingData)
		64	+ {
		65	+ blockIdx_ = GetBlockIdx();
		66	+ // 获取tiling信息
		67	+ dim_ = tilingData.dim;
		68	+ numChunk_ = tilingData.numChunk;
		69	+ outputRow_ = tilingData.outputRow;
		70	+ outputCol_ = tilingData.outputCol;
		71	+ blockRowFactor_ = tilingData.blockRowFactor;
		72	+ blockColFactor_ = tilingData.blockColFactor;
		73	+ tailBlockRowFactor_ = tilingData.tailBlockRowFactor;
		74	+ tailBlockColFactor_ = tilingData.tailBlockColFactor;
		75	+ ubRowFactor_ = tilingData.ubRowFactor;
		76	+ ubColFactor_ = tilingData.ubColFactor;
		77	+ inputNum_ = tilingData.inputNum;
		78	+ srcEleUbBlock_ = UB_BLOCK_SIZE / sizeof(T1);
		79	+ dstEleUbBlock_ = UB_BLOCK_SIZE / sizeof(T2);
		80	+
		81	+ blockRowGroup_ = blockIdx_ / tilingData.blockColNum;
		82	+ blockColGroup_ = blockIdx_ % tilingData.blockColNum;
		83	+ currentBlockRowFactor_ = blockRowGroup_ == tilingData.blockRowNum - 1 ? tailBlockRowFactor_ : blockRowFactor_;
		84	+ currentBlockColFactor_ = blockColGroup_ == tilingData.blockColNum - 1 ? tailBlockColFactor_ : blockColFactor_;
		85	+ int64_t dstGmOffset = blockRowGroup_ * blockRowFactor_ * outputCol_ + blockColGroup_ * blockColFactor_;
		D D DaiHuina16月8日确定是否存在数值溢出风险 wy519 6月8日评论： 6月8日评论：
		86	+ dstGlobal_.SetGlobalBuffer((__gm__ T2*)y + dstGmOffset);
		87	+ inputList_ = ListTensorDesc(reinterpret_cast<__gm__ void*>(x));
		88	+
		89	+ pipe_->InitBuffer(srcBuf_, tilingData.inUbSize);
		90	+ pipe_->InitBuffer(dstBuf_, tilingData.outUbSize);
		91	+ srcLocal_ = srcBuf_.Get<T1>();
		92	+ dstLocal_ = dstBuf_.Get<T2>();
		93	+ }
		94	+
		95	+ __aicore__ inline int64_t GetAlign(int64_t value, int64_t align)
		96	+ {
		97	+ return align == 0 ? value : (value + align - 1) / align * align;
		98	+ }
		99	+
		100	+ __aicore__ inline void GetChunkInfo(int32_t idx, TensorInfo& tensorInfo)
		101	+ {
		102	+ inputList_.GetDesc(desc_, idx); // scalar很大(将buf改为局部变量有改善)
		103	+ // 获取chunk相关信息
		104	+ tensorInfo.chunkDimSize = desc_.GetShape(dim_);
		105	+ tensorInfo.chunkCol = (tensorInfo.chunkDimSize + numChunk_ - 1) / numChunk_;
		106	+ // 获取concat阶段输入的col
		107	+ for (uint32_t j = 1; j < desc_.GetDim(); j++) {
		108	+ tensorInfo.originCol *= desc_.GetShape(j);
		D D DaiHuina16月8日确定是否存在数值溢出风险 wy519 6月8日评论： 6月8日评论：
		109	+ }
		110	+ tensorInfo.tensorCol = tensorInfo.chunkCol * tensorInfo.originCol;
		111	+ }
		112	+
		113	+ __aicore__ inline bool IsTensorInRange(int64_t totalCol, const UbLoopInfo& ubLoopInfo, const TensorInfo& tensorInfo)
		114	+ {
		115	+ return (totalCol < ubLoopInfo.colStart + ubLoopInfo.currentUbColFactor) &&
		116	+ (totalCol + tensorInfo.tensorCol > ubLoopInfo.colStart);
		117	+ }
		118	+
		119	+ __aicore__ inline void SplitTensorDim0(int64_t& totalCol, const UbLoopInfo& ubLoopInfo, TensorInfo& tensorInfo)
		120	+ {
		121	+ // tensor是否被切分
		122	+ tensorInfo.splitCol = tensorInfo.tensorCol;
		123	+ int64_t colEnd = ubLoopInfo.colStart + ubLoopInfo.currentUbColFactor;
		124	+ if (totalCol < ubLoopInfo.colStart && (totalCol + tensorInfo.tensorCol) > colEnd) {
		125	+ // 中间部分
		126	+ tensorInfo.isSplit = true;
		127	+ tensorInfo.splitCol = ubLoopInfo.currentUbColFactor;
		128	+ tensorInfo.startOffset = ubLoopInfo.colStart - totalCol;
		129	+ } else if (totalCol < ubLoopInfo.colStart) {
		130	+ // 被切分的后半部分
		131	+ tensorInfo.isSplit = true;
		132	+ tensorInfo.splitCol = totalCol + tensorInfo.tensorCol - ubLoopInfo.colStart;
		133	+ tensorInfo.startOffset = ubLoopInfo.colStart - totalCol;
		134	+ } else if ((totalCol + tensorInfo.tensorCol) > colEnd) {
		135	+ // 被切分的前半部分
		136	+ tensorInfo.isSplit = true;
		137	+ tensorInfo.splitCol = colEnd - totalCol;
		138	+ }
		139	+ tensorInfo.splitColAlign = GetAlign(tensorInfo.splitCol, srcEleUbBlock_);
		140	+ }
		141	+
		142	+ __aicore__ inline void ExecuteDataCopy(int64_t localOffset, int64_t gmOffset, uint16_t blockCount,
		143	+ uint32_t blockLen, uint32_t srcStride)
		144	+ {
		145	+ AscendC::DataCopyExtParams copyParams{blockCount, blockLen, srcStride, 0, 0};
		146	+ uint8_t rightPadValue = (GetAlign(blockLen, UB_BLOCK_SIZE) - blockLen) / sizeof(T1);
		147	+ AscendC::DataCopyPadExtParams<T1> padParams{true, 0, rightPadValue, 0};
		148	+ #if __CCE_AICORE__ == 310
		149	+ AscendC::DataCopyPad<T1, PaddingMode::Compact>(srcLocal_[localOffset], srcGlobal_[gmOffset], copyParams, padParams);
		150	+ #else
		151	+ AscendC::DataCopyPad(srcLocal_[localOffset], srcGlobal_[gmOffset], copyParams, padParams);
		152	+ #endif
		153	+ }
		154	+
		155	+ __aicore__ inline void DoRowsCopy(int64_t localOffset, const UbLoopInfo& ubLoopInfo, const TensorInfo& tensorInfo)
		156	+ {
		157	+ uint16_t blockCount = tensorInfo.isSplit ? static_cast<uint16_t>(ubLoopInfo.currentUbRowFactor) : 1;
		158	+ uint32_t blockLen = tensorInfo.isSplit ?
		159	+ static_cast<uint32_t>(tensorInfo.splitCol * sizeof(T1)) :
		160	+ static_cast<uint32_t>(ubLoopInfo.currentUbRowFactor * tensorInfo.splitCol * sizeof(T1));
		161	+ uint32_t srcStride = (tensorInfo.tensorCol - tensorInfo.splitCol) * sizeof(T1);
		162	+ int64_t gmOffset = ubLoopInfo.rowStart * tensorInfo.tensorCol + tensorInfo.startOffset;
		163	+ ExecuteDataCopy(localOffset, gmOffset, blockCount, blockLen, srcStride);
		164	+ }
		165	+
		166	+ __aicore__ inline void DoLastRowsCopy(int64_t localOffset, const UbLoopInfo& ubLoopInfo, const TensorInfo& tensorInfo)
		167	+ {
		168	+ // 0 无切分
		169	+ int64_t srcGmOffset = ubLoopInfo.rowStart * tensorInfo.tensorCol;
		170	+ uint32_t srcStride = (tensorInfo.tensorCol - tensorInfo.splitCol) * sizeof(T1);
		171	+ if (!tensorInfo.isSplit) {
		172	+ uint32_t blockLen = static_cast<uint32_t>(
		173	+ (tensorInfo.chunkDimSize * tensorInfo.originCol - ubLoopInfo.rowStart * tensorInfo.tensorCol) * sizeof(T1));
		174	+ ExecuteDataCopy(localOffset, srcGmOffset + tensorInfo.startOffset, 1, blockLen, srcStride);
		175	+ return;
		176	+ }
		177	+
		178	+ uint16_t blockCount = 0;
		179	+ uint32_t blockLen = 0;
		180	+ int64_t remainderCol = (tensorInfo.chunkDimSize % tensorInfo.chunkCol) * tensorInfo.originCol;
		181	+ // 1 有切分
		182	+ // 1.0 remainder等于0
		183	+ if (remainderCol == 0) {
		184	+ blockCount = static_cast<uint16_t>(tensorInfo.chunkRow - ubLoopInfo.rowStart);
		185	+ blockLen = static_cast<uint32_t>(tensorInfo.splitCol * sizeof(T1));
		186	+ }
		187	+ // 1.1 切分+偏移值小于等于remainder
		188	+ else if (tensorInfo.startOffset + tensorInfo.splitCol <= remainderCol) {
		189	+ blockCount = static_cast<uint16_t>(tensorInfo.chunkRowAlign - ubLoopInfo.rowStart);
		190	+ blockLen = static_cast<uint32_t>(tensorInfo.splitCol * sizeof(T1));
		191	+ }
		192	+ // 1.2 偏移值大于等于remainder
		193	+ else if (tensorInfo.startOffset >= remainderCol) {
		194	+ blockCount = static_cast<uint16_t>(tensorInfo.chunkRow - ubLoopInfo.rowStart);
		195	+ blockLen = static_cast<uint32_t>(tensorInfo.splitCol * sizeof(T1));
		196	+ }
		197	+ // 1.3 偏移值小于remainder，且切分+偏移值大于remainder
		198	+ else {
		199	+ // 1.3.1
		200	+ blockLen = (remainderCol - tensorInfo.startOffset) * sizeof(T1);
		201	+ int64_t localOffsetPart = 0;
		202	+ #if __CCE_AICORE__ == 310
		203	+ localOffsetPart = localOffset + GetAlign((tensorInfo.chunkRow - ubLoopInfo.rowStart) *
		204	+ tensorInfo.splitCol, srcEleUbBlock_);
		205	+ #else
		206	+ localOffsetPart = localOffset + (tensorInfo.chunkRow - ubLoopInfo.rowStart) * tensorInfo.splitColAlign;
		207	+ #endif
		208	+
		209	+ int64_t gmOffsetPart = srcGmOffset + tensorInfo.startOffset +
		210	+ (tensorInfo.chunkRow - ubLoopInfo.rowStart) * tensorInfo.tensorCol;
		211	+ ExecuteDataCopy(localOffsetPart, gmOffsetPart, 1, blockLen, srcStride);
		212	+ // 1.3.2
		213	+ blockCount = static_cast<uint16_t>(tensorInfo.chunkRow - ubLoopInfo.rowStart);
		214	+ blockLen = static_cast<uint32_t>(tensorInfo.splitCol * sizeof(T1));
		215	+ }
		216	+ ExecuteDataCopy(localOffset, srcGmOffset + tensorInfo.startOffset, blockCount, blockLen, srcStride);
		217	+ }
		218	+
		219	+ __aicore__ inline void CopyInChunk(int64_t& totalCol, int64_t& localOffset, UbLoopInfo& ubLoopInfo, TensorInfo& tensorInfo)
		220	+ {
		221	+ ubLoopInfo.isAllZero = false;
		222	+ int64_t rowEnd = ubLoopInfo.rowStart + ubLoopInfo.currentUbRowFactor;
		223	+ if (rowEnd < tensorInfo.chunkRowAlign) {
		224	+ DoRowsCopy(localOffset, ubLoopInfo, tensorInfo);
		225	+ } else {
		226	+ DoLastRowsCopy(localOffset, ubLoopInfo, tensorInfo);
		227	+ }
		228	+ }
		229	+
		230	+protected:
		231	+ int64_t blockIdx_{0};
		232	+ int64_t inputNum_{0};
		233	+ int64_t ubRowFactor_{0};
		234	+ int64_t ubColFactor_{0};
		235	+ int64_t srcEleUbBlock_{0};
		236	+ int64_t dstEleUbBlock_{0};
		237	+ int64_t dim_{0};
		238	+ int64_t numChunk_{0};
		239	+ int64_t outputRow_{0};
		240	+ int64_t outputCol_{0};
		241	+ int64_t blockRowFactor_{0};
		242	+ int64_t blockColFactor_{0};
		243	+ int64_t tailBlockRowFactor_{0};
		244	+ int64_t tailBlockColFactor_{0};
		245	+ int64_t blockRowGroup_{0};
		246	+ int64_t blockColGroup_{0};
		247	+ int64_t currentBlockRowFactor_{0};
		248	+ int64_t currentBlockColFactor_{0};
		249	+
		250	+ TPipe *pipe_;
		251	+ TEventID event_{0};
		252	+ TensorDesc<T1> desc_;
		253	+ ListTensorDesc inputList_;
		254	+ GlobalTensor<T2> dstGlobal_;
		255	+ GlobalTensor<T1> srcGlobal_;
		256	+ TBuf<AscendC::TPosition::VECCALC> srcBuf_;
		257	+ TBuf<AscendC::TPosition::VECCALC> dstBuf_;
		258	+ LocalTensor<T1> srcLocal_;
		259	+ LocalTensor<T2> dstLocal_;
		260	+};
		261	+
55	#endif // _CHUNK_CAT_COMMON_DATA_H_	262	#endif // _CHUNK_CAT_COMMON_DATA_H_

Mdocs/zh/op_api_list.md+1-1

	@@ -81,7 +81,7 @@
81	\| [aclnnCdistBackward](../../math/cdist_grad/docs/aclnnCdistBackward.md) \| 完成aclnnCdist的反向 \| 默认确定性实现\| - \|	81	\| [aclnnCdistBackward](../../math/cdist_grad/docs/aclnnCdistBackward.md) \| 完成aclnnCdist的反向 \| 默认确定性实现\| - \|
82	\| [aclnnCeil&aclnnInplaceCeil](../../math/ceil/docs/aclnnCeil&aclnnInplaceCeil.md) \| 返回输入tensor中每个元素向上取整的结果。 \| 默认确定性实现\| 默认确定性实现\|	82	\| [aclnnCeil&aclnnInplaceCeil](../../math/ceil/docs/aclnnCeil&aclnnInplaceCeil.md) \| 返回输入tensor中每个元素向上取整的结果。 \| 默认确定性实现\| 默认确定性实现\|
83	\| [aclnnChannelShuffle](../../conversion/transpose/docs/aclnnChannelShuffle.md) \| 将(\*, C, H, W)张量的channels分成g个组，然后将每个通道组中的通道进行随机重排，最后将所有通道合并输出，同时保持最终输出张量的shape和输入张量保持一致。 \| 默认确定性实现\| 默认确定性实现\|	83	\| [aclnnChannelShuffle](../../conversion/transpose/docs/aclnnChannelShuffle.md) \| 将(\*, C, H, W)张量的channels分成g个组，然后将每个通道组中的通道进行随机重排，最后将所有通道合并输出，同时保持最终输出张量的shape和输入张量保持一致。 \| 默认确定性实现\| 默认确定性实现\|
84	-\| [aclnnChunkCat](../../conversion/chunk_cat/docs/aclnnChunkCat.md) \| 将tensors中所有tensor先按照维度dim切分为numChunks块，再按照dim后一维进行级联，最后转换为out的数据类型。 \| 默认确定性实现\| - \|	84	+\| [aclnnChunkCat](../../conversion/chunk_cat/docs/aclnnChunkCat.md) \| 将tensors中所有tensor先按照维度dim切分为numChunks块，再按照dim后一维进行级联，最后转换为out的数据类型。 \| 默认确定性实现\| 默认确定性实现\|
85	\| [aclnnCircularPad2d](../../conversion/circular_pad/docs/aclnnCircularPad2d.md) \| 使用输入循环填充输入tensor的最后两维。 \| 默认确定性实现\| 默认确定性实现\|	85	\| [aclnnCircularPad2d](../../conversion/circular_pad/docs/aclnnCircularPad2d.md) \| 使用输入循环填充输入tensor的最后两维。 \| 默认确定性实现\| 默认确定性实现\|
86	\| [aclnnCircularPad2dBackward](../../conversion/circular_pad_grad/docs/aclnnCircularPad2dBackward.md) \| circular_pad2d的反向传播, 前向计算参考aclnnCircularPad2d。 \| 默认确定性实现\| 默认确定性实现\|	86	\| [aclnnCircularPad2dBackward](../../conversion/circular_pad_grad/docs/aclnnCircularPad2dBackward.md) \| circular_pad2d的反向传播, 前向计算参考aclnnCircularPad2d。 \| 默认确定性实现\| 默认确定性实现\|
87	\| [aclnnCircularPad3d](../../conversion/circular_pad/docs/aclnnCircularPad3d.md) \| 使用输入循环填充输入tensor的最后三维。 \| 默认确定性实现\| 默认确定性实现\|	87	\| [aclnnCircularPad3d](../../conversion/circular_pad/docs/aclnnCircularPad3d.md) \| 使用输入循环填充输入tensor的最后三维。 \| 默认确定性实现\| 默认确定性实现\|

	@@ -4,7 +4,7 @@
4		4
5	\| 产品 \| 是否支持 \|	5	\| 产品 \| 是否支持 \|
6	\| :----------------------------------------------------------- \| :------: \|	6	\| :----------------------------------------------------------- \| :------: \|
7	-\| Ascend 950PR/Ascend 950DT \| × \|	7	+\| Ascend 950PR/Ascend 950DT \| √ \|
8	\| Atlas A3 训练系列产品/Atlas A3 推理系列产品 \| √ \|	8	\| Atlas A3 训练系列产品/Atlas A3 推理系列产品 \| √ \|
9	\| Atlas A2 训练系列产品/Atlas A2 推理系列产品 \| √ \|	9	\| Atlas A2 训练系列产品/Atlas A2 推理系列产品 \| √ \|
10	\| Atlas 200I/500 A2 推理产品 \| × \|	10	\| Atlas 200I/500 A2 推理产品 \| × \|

	@@ -6,7 +6,7 @@
6		6
7	\| 产品 \| 是否支持 \|	7	\| 产品 \| 是否支持 \|
8	\| :------------------------------------------------------- \| :------: \|	8	\| :------------------------------------------------------- \| :------: \|
9	-\| <term>Ascend 950PR/Ascend 950DT</term> \| × \|	9	+\| <term>Ascend 950PR/Ascend 950DT</term> \| √ \|
10	\| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> \| √ \|	10	\| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> \| √ \|
11	\| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> \| √ \|	11	\| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> \| √ \|
12	\| <term>Atlas 200I/500 A2 推理产品</term> \| × \|	12	\| <term>Atlas 200I/500 A2 推理产品</term> \| × \|