fix(tile_with_axis): Tiling溢出/CeilDiv回绕与标量UB越界加固，补动态shape范围推导及golden-ops-math-AtomGit

Mconversion/tile_with_axis/README.md+13-11

	@@ -4,12 +4,12 @@
4		4
5	\| 产品 \| 是否支持 \|	5	\| 产品 \| 是否支持 \|
6	\| :----------------------------------------------------------- \| :------: \|	6	\| :----------------------------------------------------------- \| :------: \|
7	-\| Ascend 950PR&950DT系列产品 \| √ \|	7	+\| <term>Ascend 950PR&950DT系列产品</term> \| √ \|
8	-\| Atlas A3系列产品 \| √ \|	8	+\| <term>Atlas A3系列产品</term> \| √ \|
9	-\| Atlas A2系列产品 \| √ \|	9	+\| <term>Atlas A2系列产品</term> \| √ \|
10	-\| Atlas 200I/500 A2推理产品 \| √ \|	10	+\| <term>Atlas 200I/500 A2推理产品</term> \| √ \|
11	-\| Atlas推理系列产品 \| √ \|	11	+\| <term>Atlas推理系列产品</term> \| √ \|
12	-\| Atlas训练系列产品 \| √ \|	12	+\| <term>Atlas训练系列产品</term> \| √ \|
13		13
14	## 功能说明	14	## 功能说明
15		15
	@@ -27,10 +27,10 @@
27		27
28	\| 参数名 \| 输入/输出/属性 \| 描述 \| 数据类型 \| 数据格式 \|	28	\| 参数名 \| 输入/输出/属性 \| 描述 \| 数据类型 \| 数据格式 \|
29	\|--------\|---------------\|------\|----------\|----------\|	29	\|--------\|---------------\|------\|----------\|----------\|
30	-\| x \| 输入 \| 输入张量 \| FLOAT16、FLOAT、INT8、INT16、INT32、INT64、UINT8、UINT16、UINT32、UINT64 \| ND \|	30	+\| x \| 输入 \| 输入张量。 \| BFLOAT16、FLOAT16、FLOAT、INT32、INT8、UINT8 \| ND \|
31	-\| axis \| 属性（可选） \| 指定复制的维度，默认为1 \| INT \| - \|	31	+\| y \| 输出 \| 输出张量，与x具有相同的数据类型和格式。 \| 与x一致 \| ND \|
32	-\| tiles \| 属性（必选） \| 复制次数，必须大于0 \| INT \| - \|	32	+\| axis \| 可选属性 \| 指定复制的维度，默认为1。 \| INT \| - \|
33	-\| y \| 输出 \| 输出张量，与x具有相同的数据类型和格式 \| 与x一致 \| ND \|	33	+\| tiles \| 属性 \| 复制次数，必须大于0。 \| INT \| - \|
34		34
35	## 约束说明	35	## 约束说明
36		36
	@@ -38,9 +38,11 @@
38	- tiles 必须大于0。	38	- tiles 必须大于0。
39	- 输出Tensor在 axis 维度的大小必须等于输入Tensor在 axis 维度的大小乘以 tiles。	39	- 输出Tensor在 axis 维度的大小必须等于输入Tensor在 axis 维度的大小乘以 tiles。
40	- 输入Tensor维度不超过8维。	40	- 输入Tensor维度不超过8维。
		41	+- 支持标量输入（rank=0）：输出 shape 为 [tiles]，即输入标量值复制 tiles 份；此时 axis 不生效。
		42	+- 输入Tensor任一维允许为 0（空 Tensor）：输出对应维仍为 0，输出Tensor为空。
41		43
42	## 调用说明	44	## 调用说明
43		45
44	\| 调用方式 \| 样例代码 \| 说明 \|	46	\| 调用方式 \| 样例代码 \| 说明 \|
45	\|------------\|----------\|------\|	47	\|------------\|----------\|------\|
46	-\| 图模式调用 \| [test_geir_tile_with_axis](./examples/test_geir_tile_with_axis.cpp) \| 通过[算子IR](./op_graph/tile_with_axis_proto.h)构图方式调用TileWithAxis算子。 \|	48	+\| GE图模式 \| [test_geir_tile_with_axis](./examples/test_geir_tile_with_axis.cpp) \| 通过[算子IR](./op_graph/tile_with_axis_proto.h)构图方式调用TileWithAxis算子。 \|

Mconversion/tile_with_axis/examples/test_geir_tile_with_axis.cpp+28-8

	@@ -184,7 +184,7 @@ int32_t WriteDataToFile(string bin_file, uint64_t data_size, uint8_t* inputData)
184	return SUCCESS;	184	return SUCCESS;
185	}	185	}
186		186
187	-void ProcessInputData(vector<Tensor>& input)	187	+int ProcessInputData(vector<Tensor>& input)
188	{	188	{
189	for (size_t i = 0; i < input.size(); i++) {	189	for (size_t i = 0; i < input.size(); i++) {
190	string input_file = "./tc_ge_irrun_test_0008_npu_input_" + std::to_string(i) + ".bin";	190	string input_file = "./tc_ge_irrun_test_0008_npu_input_" + std::to_string(i) + ".bin";
	@@ -193,14 +193,18 @@ void ProcessInputData(vector<Tensor>& input)
193	uint32_t type_size = GetDataTypeSize(input[i].GetTensorDesc().GetDataType());	193	uint32_t type_size = GetDataTypeSize(input[i].GetTensorDesc().GetDataType());
194	if (type_size == 0U) {	194	if (type_size == 0U) {
195	printf("ERROR: input %zu has unsupported dtype\n", i);	195	printf("ERROR: input %zu has unsupported dtype\n", i);
196	- continue;	196	+ return FAILED;
197	}	197	}
198	uint32_t data_size = static_cast<uint32_t>(input_shape * type_size);	198	uint32_t data_size = static_cast<uint32_t>(input_shape * type_size);
199	- WriteDataToFile(input_file.c_str(), data_size, input_data_i);	199	+ if (WriteDataToFile(input_file.c_str(), data_size, input_data_i) != SUCCESS) {
		200	+ printf("ERROR: write input bin %s failed\n", input_file.c_str());
		201	+ return FAILED;
		202	+ }
200	}	203	}
		204	+ return SUCCESS;
201	}	205	}
202		206
203	-void ProcessOutputData(vector<Tensor>& output)	207	+int ProcessOutputData(vector<Tensor>& output)
204	{	208	{
205	for (size_t i = 0; i < output.size(); i++) {	209	for (size_t i = 0; i < output.size(); i++) {
206	string output_file = "./tc_ge_irrun_test_0008_npu_output_" + std::to_string(i) + ".bin";	210	string output_file = "./tc_ge_irrun_test_0008_npu_output_" + std::to_string(i) + ".bin";
	@@ -209,11 +213,15 @@ void ProcessOutputData(vector<Tensor>& output)
209	uint32_t type_size = GetDataTypeSize(output[i].GetTensorDesc().GetDataType());	213	uint32_t type_size = GetDataTypeSize(output[i].GetTensorDesc().GetDataType());
210	if (type_size == 0U) {	214	if (type_size == 0U) {
211	printf("ERROR: output %zu has unsupported dtype\n", i);	215	printf("ERROR: output %zu has unsupported dtype\n", i);
212	- continue;	216	+ return FAILED;
213	}	217	}
214	uint32_t data_size = static_cast<uint32_t>(output_shape * type_size);	218	uint32_t data_size = static_cast<uint32_t>(output_shape * type_size);
215	- WriteDataToFile(output_file.c_str(), data_size, output_data_i);	219	+ if (WriteDataToFile(output_file.c_str(), data_size, output_data_i) != SUCCESS) {
		220	+ printf("ERROR: write output bin %s failed\n", output_file.c_str());
		221	+ return FAILED;
		222	+ }
216	}	223	}
		224	+ return SUCCESS;
217	}	225	}
218		226
219	int CreateOppInGraph(DataType inDtype, vector<Tensor>& input, vector<Operator>& inputs, vector<Operator>& outputs,	227	int CreateOppInGraph(DataType inDtype, vector<Tensor>& input, vector<Operator>& inputs, vector<Operator>& outputs,
	@@ -280,6 +288,7 @@ int main(int argc, char* argv[])
280	ret = CreateOppInGraph(inDtype, input, inputs, outputs, graph);	288	ret = CreateOppInGraph(inDtype, input, inputs, outputs, graph);
281	if (ret != SUCCESS) {	289	if (ret != SUCCESS) {
282	printf("%s - ERROR - [XIR]: Create graph failed\n", GetTime().c_str());	290	printf("%s - ERROR - [XIR]: Create graph failed\n", GetTime().c_str());
		291	+ GEFinalize();
283	return FAILED;	292	return FAILED;
284	}	293	}
285		294
	@@ -291,6 +300,7 @@ int main(int argc, char* argv[])
291	ge::Session* session = new Session(build_options);	300	ge::Session* session = new Session(build_options);
292	if (session == nullptr) {	301	if (session == nullptr) {
293	printf("%s - ERROR - [XIR]: Create session failed\n", GetTime().c_str());	302	printf("%s - ERROR - [XIR]: Create session failed\n", GetTime().c_str());
		303	+ GEFinalize();
294	return FAILED;	304	return FAILED;
295	}	305	}
296		306
	@@ -314,8 +324,18 @@ int main(int argc, char* argv[])
314	}	324	}
315	printf("%s - INFO - [XIR]: Run graph success\n", GetTime().c_str());	325	printf("%s - INFO - [XIR]: Run graph success\n", GetTime().c_str());
316		326
317	- ProcessInputData(input);	327	+ if (ProcessInputData(input) != SUCCESS) {
318	- ProcessOutputData(output);	328	+ printf("%s - ERROR - [XIR]: Process input data failed\n", GetTime().c_str());
		329	+ delete session;
		330	+ GEFinalize();
		331	+ return FAILED;
		332	+ }
		333	+ if (ProcessOutputData(output) != SUCCESS) {
		334	+ printf("%s - ERROR - [XIR]: Process output data failed\n", GetTime().c_str());
		335	+ delete session;
		336	+ GEFinalize();
		337	+ return FAILED;
		338	+ }
319		339
320	delete session;	340	delete session;
321	ret = ge::GEFinalize();	341	ret = ge::GEFinalize();

Mconversion/tile_with_axis/op_graph/tile_with_axis_proto.h+1-1

	@@ -25,7 +25,7 @@ namespace ge {
25		25
26	* @par Inputs:	26	* @par Inputs:
27	* One input:	27	* One input:
28	-* x: A Tensor with any format. Must be one of the following types:	28	+* x: A Tensor with ND format. Must be one of the following types:
29	* bfloat16, float16, float32, int8, int16, int32, int64, uint8, uint16, uint32, uint64 . \n	29	* bfloat16, float16, float32, int8, int16, int32, int64, uint8, uint16, uint32, uint64 . \n
30		30
31	* @par Attributes:	31	* @par Attributes:

Mconversion/tile_with_axis/op_host/arch35/tile_with_axis_tiling.cpp+94-21

	@@ -46,6 +46,22 @@ static constexpr int32_t TILING_KEY_ELEMENT_COUNT = 2;
46	// ============================================================================	46	// ============================================================================
47	// Helper: Get platform info (ubSize, coreNum)	47	// Helper: Get platform info (ubSize, coreNum)
48	// ============================================================================	48	// ============================================================================
		49	+// ============================================================================
		50	+// Helper: guard an int64 multiplication against signed overflow
		51	+// ============================================================================
		52	+static ge::graphStatus CheckMulNoOverflow(gert::TilingContext* context, int64_t a, int64_t b, const char* product_name)
		53	+{
		54	+ if (a < 0 \|\| b < 0) {
		55	+ OP_LOGE(context, "%s operand is negative: %ld * %ld", product_name, a, b);
		56	+ return ge::GRAPH_FAILED;
		57	+ }
		58	+ if (a != 0 && b != 0 && a > INT64_MAX / b) {
		59	+ OP_LOGE(context, "%s overflows int64: %ld * %ld", product_name, a, b);
		60	+ return ge::GRAPH_FAILED;
		61	+ }
		62	+ return ge::GRAPH_SUCCESS;
		63	+}
		64	+
49	static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t* ubSize, int64_t* coreNum)	65	static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t* ubSize, int64_t* coreNum)
50	{	66	{
51	fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo();	67	fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo();
	@@ -82,7 +98,8 @@ static void LogTilingData(gert::TilingContext* context, const TileWithAxisTiling
82	tiling->bufferSize, tiling->tiles, tiling->rowLength);	98	tiling->bufferSize, tiling->tiles, tiling->rowLength);
83	}	99	}
84		100
85	-static ge::graphStatus HandleScalarInput(gert::TilingContext* context, int64_t tiles, uint64_t ubSize)	101	+static ge::graphStatus HandleScalarInput(gert::TilingContext* context, int64_t tiles, uint64_t ubSize,
		102	+ int64_t dtypeSize, int64_t coreNum)
86	{	103	{
87	TileWithAxisTilingData* tiling = context->GetTilingData<TileWithAxisTilingData>();	104	TileWithAxisTilingData* tiling = context->GetTilingData<TileWithAxisTilingData>();
88	OP_CHECK_NULL_WITH_CONTEXT(context, tiling);	105	OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
	@@ -94,16 +111,38 @@ static ge::graphStatus HandleScalarInput(gert::TilingContext* context, int64_t t
94	tiling->outShape[0] = 1;	111	tiling->outShape[0] = 1;
95	tiling->outShape[1] = tiles;	112	tiling->outShape[1] = tiles;
96	tiling->outShape[2] = 1;	113	tiling->outShape[2] = 1;
97	- tiling->totalCount = 1;
98	- tiling->perCoreCount = 1;
99	- tiling->ubAxis = 0;
100	- tiling->ubFactor = 1;
101	tiling->bufferSize = ubSize / 2;	114	tiling->bufferSize = ubSize / 2;
102	tiling->tiles = tiles;	115	tiling->tiles = tiles;
103	tiling->rowLength = 1;	116	tiling->rowLength = 1;
104	- LogTilingData(context, tiling);	117	+
105	- context->SetBlockDim(1);	118	+ // 标量输出共 tiles 个元素；单缓冲容量 bufferSize/dtypeSize 个元素。
106	- ASCENDC_TPL_SEL_PARAM(context, static_cast<uint32_t>(0));	119	+ // tiles 超过单缓冲容量时沿 tiles 维分块（ubAxis=1: NDDMA dim1 源 stride=0 广播读，见 kernel
		120	+ // InitNddmaParams），每块不超过单缓冲，消除原实现的 UB 越界写。
		121	+ int64_t bufferSizeElements = static_cast<int64_t>(tiling->bufferSize) / dtypeSize;
		122	+ OP_CHECK_IF(bufferSizeElements <= 0,
		123	+ OP_LOGE(context, "UB half buffer %u bytes too small for dtype size %ld", tiling->bufferSize, dtypeSize),
		124	+ return ge::GRAPH_FAILED);
		125	+ if (tiles <= bufferSizeElements) {
		126	+ tiling->totalCount = 1;
		127	+ tiling->perCoreCount = 1;
		128	+ tiling->ubAxis = 0;
		129	+ tiling->ubFactor = 1;
		130	+ LogTilingData(context, tiling);
		131	+ context->SetBlockDim(1);
		132	+ ASCENDC_TPL_SEL_PARAM(context, static_cast<uint32_t>(0));
		133	+ } else {
		134	+ uint32_t ubFactor = static_cast<uint32_t>(bufferSizeElements);
		135	+ uint64_t totalCount = static_cast<uint64_t>(Ops::Base::CeilDiv(tiles, static_cast<int64_t>(ubFactor)));
		136	+ uint64_t perCoreCount = static_cast<uint64_t>(Ops::Base::CeilDiv(static_cast<int64_t>(totalCount), coreNum));
		137	+ int64_t realCoreNum = Ops::Base::CeilDiv(static_cast<int64_t>(totalCount), static_cast<int64_t>(perCoreCount));
		138	+ tiling->totalCount = totalCount;
		139	+ tiling->perCoreCount = perCoreCount;
		140	+ tiling->ubAxis = 1;
		141	+ tiling->ubFactor = ubFactor;
		142	+ LogTilingData(context, tiling);
		143	+ context->SetBlockDim(realCoreNum);
		144	+ ASCENDC_TPL_SEL_PARAM(context, static_cast<uint32_t>(1));
		145	+ }
107	OP_CHECK_IF(GetWorkspaceSize(context) != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetWorkspaceSize error"),	146	OP_CHECK_IF(GetWorkspaceSize(context) != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetWorkspaceSize error"),
108	return ge::GRAPH_FAILED);	147	return ge::GRAPH_FAILED);
109	return ge::GRAPH_SUCCESS;	148	return ge::GRAPH_SUCCESS;
	@@ -131,9 +170,12 @@ static int64_t MaxBlocksOnAxis(uint8_t ax, int64_t outerDim, int64_t tiles, int6
131	{	170	{
132	if (ax == 0)	171	if (ax == 0)
133	return outerDim;	172	return outerDim;
134	- if (ax == 1)	173	+ if (ax == 1) {
135	- return outerDim * tiles;	174	+ // 仅 rowLength>0 时可达（此时 outerDim*tiles 已被入口总乘积检查约束）
136	- return outerDim * tiles * rowLength;	175	+ return (rowLength == 0) ? 0 : outerDim * tiles;
		176	+ }
		177	+ // ax == 2
		178	+ return (rowLength == 0) ? 0 : outerDim * tiles * rowLength;
137	}	179	}
138		180
139	// ============================================================================	181	// ============================================================================
	@@ -148,6 +190,10 @@ static uint32_t ComputeUbFactor(uint8_t ax, int64_t target, int64_t outerDim, in
148	int64_t bpo = CeilDiv(target, outerDim);	190	int64_t bpo = CeilDiv(target, outerDim);
149	return static_cast<uint32_t>(std::max(static_cast<int64_t>(1), CeilDiv(tiles, bpo)));	191	return static_cast<uint32_t>(std::max(static_cast<int64_t>(1), CeilDiv(tiles, bpo)));
150	}	192	}
		193	+ // ax == 2；rowLength==0 为空张量，无需按块计算 ubFactor
		194	+ if (rowLength == 0) {
		195	+ return 1;
		196	+ }
151	int64_t bpot = CeilDiv(target, outerDim * tiles);	197	int64_t bpot = CeilDiv(target, outerDim * tiles);
152	return static_cast<uint32_t>(std::max(static_cast<int64_t>(1), CeilDiv(rowLength, bpot)));	198	return static_cast<uint32_t>(std::max(static_cast<int64_t>(1), CeilDiv(rowLength, bpot)));
153	}	199	}
	@@ -161,7 +207,8 @@ static uint64_t ComputeTotalCount(uint8_t ax, int64_t uf, int64_t outerDim, int6
161	return static_cast<uint64_t>(CeilDiv(outerDim, uf));	207	return static_cast<uint64_t>(CeilDiv(outerDim, uf));
162	if (ax == 1)	208	if (ax == 1)
163	return static_cast<uint64_t>(outerDim * CeilDiv(tiles, uf));	209	return static_cast<uint64_t>(outerDim * CeilDiv(tiles, uf));
164	- return static_cast<uint64_t>(outerDim * tiles * CeilDiv(rowLength, uf));	210	+ // ax == 2；rowLength==0 为空张量，块数为 0
		211	+ return (rowLength == 0) ? 0U : static_cast<uint64_t>(outerDim * tiles * CeilDiv(rowLength, uf));
165	}	212	}
166		213
167	// ============================================================================	214	// ============================================================================
	@@ -182,7 +229,8 @@ static void SelectInitialUbAxis(int64_t ubFactor0, int64_t ubFactor1, int64_t ub
182	} else {	229	} else {
183	ubAxis = 2;	230	ubAxis = 2;
184	ubFactor = static_cast<uint32_t>(std::max(ubFactor2, static_cast<int64_t>(1)));	231	ubFactor = static_cast<uint32_t>(std::max(ubFactor2, static_cast<int64_t>(1)));
185	- totalCount = static_cast<uint64_t>(outerDim * tiles * CeilDiv(rowLength, ubFactor2));	232	+ // rowLength==0 为空张量，块数为 0
		233	+ totalCount = (rowLength == 0) ? 0U : static_cast<uint64_t>(outerDim * tiles * CeilDiv(rowLength, ubFactor2));
186	}	234	}
187	}	235	}
188		236
	@@ -274,8 +322,11 @@ static ge::graphStatus TileWithAxisTilingFunc(gert::TilingContext* context)
274	OP_CHECK_NULL_WITH_CONTEXT(context, attrs);	322	OP_CHECK_NULL_WITH_CONTEXT(context, attrs);
275	const int64_t* axis_ptr = attrs->GetAttrPointer<int64_t>(ATTR_AXIS_IDX);	323	const int64_t* axis_ptr = attrs->GetAttrPointer<int64_t>(ATTR_AXIS_IDX);
276	const int64_t* tiles_ptr = attrs->GetAttrPointer<int64_t>(ATTR_TILES_IDX);	324	const int64_t* tiles_ptr = attrs->GetAttrPointer<int64_t>(ATTR_TILES_IDX);
277	- int64_t axis = (axis_ptr != nullptr) ? *axis_ptr : 0;	325	+ // tiles 为必选属性（proto REQUIRED_ATTR / def REQUIRED / README）：缺失时拒绝，不静默回退
278	- int64_t tiles = (tiles_ptr != nullptr) ? *tiles_ptr : 1;	326	+ OP_CHECK_IF(tiles_ptr == nullptr, OP_LOGE(context, "tiles is required but missing"), return ge::GRAPH_FAILED);
		327	+ int64_t tiles = *tiles_ptr;
		328	+ // axis 默认 1（与 TBE/AICPU 缺省一致）；属性缺失时兜底 1
		329	+ int64_t axis = (axis_ptr != nullptr) ? *axis_ptr : 1;
279		330
280	OP_CHECK_IF(tiles <= 0, OP_LOGE(context, "tiles must be > 0, got %ld", tiles), return ge::GRAPH_FAILED);	331	OP_CHECK_IF(tiles <= 0, OP_LOGE(context, "tiles must be > 0, got %ld", tiles), return ge::GRAPH_FAILED);
281		332
	@@ -283,7 +334,11 @@ static ge::graphStatus TileWithAxisTilingFunc(gert::TilingContext* context)
283		334
284	// 3. 标量输入特殊处理	335	// 3. 标量输入特殊处理
285	if (rank == 0) {	336	if (rank == 0) {
286	- return HandleScalarInput(context, tiles, ubSize);	337	+ int64_t scalarDtypeSize = static_cast<int64_t>(ge::GetSizeByDataType(dataType));
		338	+ OP_CHECK_IF(scalarDtypeSize <= 0,
		339	+ OP_LOGE(context, "unsupported dtype, GetSizeByDataType returned %ld", scalarDtypeSize),
		340	+ return ge::GRAPH_FAILED);
		341	+ return HandleScalarInput(context, tiles, ubSize, scalarDtypeSize, coreNum);
287	}	342	}
288		343
289	// 4. 校验并归一化 axis	344	// 4. 校验并归一化 axis
	@@ -293,13 +348,21 @@ static ge::graphStatus TileWithAxisTilingFunc(gert::TilingContext* context)
293	// 5. 计算展平参数 (v1.6 折叠 axisDim 进 rowLength)	348	// 5. 计算展平参数 (v1.6 折叠 axisDim 进 rowLength)
294	int64_t outerDim = 1;	349	int64_t outerDim = 1;
295	for (int64_t i = 0; i < axis; i++) {	350	for (int64_t i = 0; i < axis; i++) {
296	- outerDim *= inputShape.GetDim(i);	351	+ int64_t dim = inputShape.GetDim(i);
		352	+ OP_CHECK_IF(CheckMulNoOverflow(context, outerDim, dim, "outerDim") != ge::GRAPH_SUCCESS,
		353	+ OP_LOGE(context, "outerDim overflow"), return ge::GRAPH_FAILED);
		354	+ outerDim *= dim;
297	}	355	}
298	int64_t axisDim = inputShape.GetDim(axis);	356	int64_t axisDim = inputShape.GetDim(axis);
299	int64_t innerDim = 1;	357	int64_t innerDim = 1;
300	for (int64_t i = axis + 1; i < rank; i++) {	358	for (int64_t i = axis + 1; i < rank; i++) {
301	- innerDim *= inputShape.GetDim(i);	359	+ int64_t dim = inputShape.GetDim(i);
		360	+ OP_CHECK_IF(CheckMulNoOverflow(context, innerDim, dim, "innerDim") != ge::GRAPH_SUCCESS,
		361	+ OP_LOGE(context, "innerDim overflow"), return ge::GRAPH_FAILED);
		362	+ innerDim *= dim;
302	}	363	}
		364	+ OP_CHECK_IF(CheckMulNoOverflow(context, axisDim, innerDim, "rowLength") != ge::GRAPH_SUCCESS,
		365	+ OP_LOGE(context, "rowLength overflow"), return ge::GRAPH_FAILED);
303	int64_t rowLength = axisDim * innerDim;	366	int64_t rowLength = axisDim * innerDim;
304		367
305	// 6. TilingData 空间分配	368	// 6. TilingData 空间分配
	@@ -313,7 +376,13 @@ static ge::graphStatus TileWithAxisTilingFunc(gert::TilingContext* context)
313	OP_CHECK_IF(dtypeSize <= 0, OP_LOGE(context, "unsupported dtype, GetSizeByDataType returned %ld", dtypeSize),	376	OP_CHECK_IF(dtypeSize <= 0, OP_LOGE(context, "unsupported dtype, GetSizeByDataType returned %ld", dtypeSize),
314	return ge::GRAPH_FAILED);	377	return ge::GRAPH_FAILED);
315	int64_t bufferSizeElements = (static_cast<int64_t>(ubSize) / 2) / dtypeSize;	378	int64_t bufferSizeElements = (static_cast<int64_t>(ubSize) / 2) / dtypeSize;
		379	+ OP_CHECK_IF(CheckMulNoOverflow(context, tiles, rowLength, "rowOutElements") != ge::GRAPH_SUCCESS,
		380	+ OP_LOGE(context, "rowOutElements overflow"), return ge::GRAPH_FAILED);
316	int64_t rowOutElements = tiles * rowLength;	381	int64_t rowOutElements = tiles * rowLength;
		382	+ // 总输出元素数 outerDimtilesrowLength：下游 AdjustForCoreUtilization 的 totalOutElements、
		383	+ // MaxBlocksOnAxis 与 ComputeTotalCount 均为这三个因子的子集/缩放，验证后不再溢出。
		384	+ OP_CHECK_IF(CheckMulNoOverflow(context, outerDim, rowOutElements, "totalOutElements") != ge::GRAPH_SUCCESS,
		385	+ OP_LOGE(context, "totalOutElements overflow"), return ge::GRAPH_FAILED);
317		386
318	int64_t ubFactor0 = (rowOutElements > 0) ? std::min(outerDim, bufferSizeElements / rowOutElements) : 0;	387	int64_t ubFactor0 = (rowOutElements > 0) ? std::min(outerDim, bufferSizeElements / rowOutElements) : 0;
319	int64_t ubFactor1 = (rowLength > 0) ? std::min(tiles, bufferSizeElements / rowLength) : 0;	388	int64_t ubFactor1 = (rowLength > 0) ? std::min(tiles, bufferSizeElements / rowLength) : 0;
	@@ -328,9 +397,13 @@ static ge::graphStatus TileWithAxisTilingFunc(gert::TilingContext* context)
328	AdjustForCoreUtilization(ubAxis, ubFactor, totalCount, ubFactor0, ubFactor1, ubFactor2, outerDim, tiles, rowLength,	397	AdjustForCoreUtilization(ubAxis, ubFactor, totalCount, ubFactor0, ubFactor1, ubFactor2, outerDim, tiles, rowLength,
329	coreNum);	398	coreNum);
330		399
331	- // 8. 多核切分	400	+ // 8. 多核切分（空张量兜底：totalCount==0 时 SetBlockDim(1)，kernel 空循环，见 REQUIREMENTS §3.5.2）
332	- uint64_t perCoreCount = static_cast<uint64_t>(CeilDiv(static_cast<int64_t>(totalCount), coreNum));	401	+ uint64_t perCoreCount = 1;
333	- int64_t realCoreNum = CeilDiv(static_cast<int64_t>(totalCount), static_cast<int64_t>(perCoreCount));	402	+ int64_t realCoreNum = 1;
		403	+ if (totalCount > 0) {
		404	+ perCoreCount = static_cast<uint64_t>(CeilDiv(static_cast<int64_t>(totalCount), coreNum));
		405	+ realCoreNum = CeilDiv(static_cast<int64_t>(totalCount), static_cast<int64_t>(perCoreCount));
		406	+ }
334		407
335	// 9. 填充 TilingData	408	// 9. 填充 TilingData
336	tiling->inShape[0] = outerDim;	409	tiling->inShape[0] = outerDim;

Mconversion/tile_with_axis/op_host/config/ascend350/tile_with_axis_binary.json+12-24

	@@ -12,8 +12,7 @@
12	"paramType": "required",	12	"paramType": "required",
13	"shape": [	13	"shape": [
14	-2	14	-2
15	- ],	15	+ ]
16	- "dtype_match_mode": "DtypeByte"
17	}	16	}
18	],	17	],
19	"outputs": [	18	"outputs": [
	@@ -25,8 +24,7 @@
25	"paramType": "required",	24	"paramType": "required",
26	"shape": [	25	"shape": [
27	-2	26	-2
28	- ],	27	+ ]
29	- "dtype_match_mode": "DtypeByte"
30	}	28	}
31	],	29	],
32	"attrs": [	30	"attrs": [
	@@ -53,8 +51,7 @@
53	"paramType": "required",	51	"paramType": "required",
54	"shape": [	52	"shape": [
55	-2	53	-2
56	- ],	54	+ ]
57	- "dtype_match_mode": "DtypeByte"
58	}	55	}
59	],	56	],
60	"outputs": [	57	"outputs": [
	@@ -66,8 +63,7 @@
66	"paramType": "required",	63	"paramType": "required",
67	"shape": [	64	"shape": [
68	-2	65	-2
69	- ],	66	+ ]
70	- "dtype_match_mode": "DtypeByte"
71	}	67	}
72	],	68	],
73	"attrs": [	69	"attrs": [
	@@ -94,8 +90,7 @@
94	"paramType": "required",	90	"paramType": "required",
95	"shape": [	91	"shape": [
96	-2	92	-2
97	- ],	93	+ ]
98	- "dtype_match_mode": "DtypeByte"
99	}	94	}
100	],	95	],
101	"outputs": [	96	"outputs": [
	@@ -107,8 +102,7 @@
107	"paramType": "required",	102	"paramType": "required",
108	"shape": [	103	"shape": [
109	-2	104	-2
110	- ],	105	+ ]
111	- "dtype_match_mode": "DtypeByte"
112	}	106	}
113	],	107	],
114	"attrs": [	108	"attrs": [
	@@ -135,8 +129,7 @@
135	"paramType": "required",	129	"paramType": "required",
136	"shape": [	130	"shape": [
137	-2	131	-2
138	- ],	132	+ ]
139	- "dtype_match_mode": "DtypeByte"
140	}	133	}
141	],	134	],
142	"outputs": [	135	"outputs": [
	@@ -148,8 +141,7 @@
148	"paramType": "required",	141	"paramType": "required",
149	"shape": [	142	"shape": [
150	-2	143	-2
151	- ],	144	+ ]
152	- "dtype_match_mode": "DtypeByte"
153	}	145	}
154	],	146	],
155	"attrs": [	147	"attrs": [
	@@ -176,8 +168,7 @@
176	"paramType": "required",	168	"paramType": "required",
177	"shape": [	169	"shape": [
178	-2	170	-2
179	- ],	171	+ ]
180	- "dtype_match_mode": "DtypeByte"
181	}	172	}
182	],	173	],
183	"outputs": [	174	"outputs": [
	@@ -189,8 +180,7 @@
189	"paramType": "required",	180	"paramType": "required",
190	"shape": [	181	"shape": [
191	-2	182	-2
192	- ],	183	+ ]
193	- "dtype_match_mode": "DtypeByte"
194	}	184	}
195	],	185	],
196	"attrs": [	186	"attrs": [
	@@ -217,8 +207,7 @@
217	"paramType": "required",	207	"paramType": "required",
218	"shape": [	208	"shape": [
219	-2	209	-2
220	- ],	210	+ ]
221	- "dtype_match_mode": "DtypeByte"
222	}	211	}
223	],	212	],
224	"outputs": [	213	"outputs": [
	@@ -230,8 +219,7 @@
230	"paramType": "required",	219	"paramType": "required",
231	"shape": [	220	"shape": [
232	-2	221	-2
233	- ],	222	+ ]
234	- "dtype_match_mode": "DtypeByte"
235	}	223	}
236	],	224	],
237	"attrs": [	225	"attrs": [

Mconversion/tile_with_axis/op_host/config/ascend950/tile_with_axis_binary.json+12-24

	@@ -12,8 +12,7 @@
12	"paramType": "required",	12	"paramType": "required",
13	"shape": [	13	"shape": [
14	-2	14	-2
15	- ],	15	+ ]
16	- "dtype_match_mode": "DtypeByte"
17	}	16	}
18	],	17	],
19	"outputs": [	18	"outputs": [
	@@ -25,8 +24,7 @@
25	"paramType": "required",	24	"paramType": "required",
26	"shape": [	25	"shape": [
27	-2	26	-2
28	- ],	27	+ ]
29	- "dtype_match_mode": "DtypeByte"
30	}	28	}
31	],	29	],
32	"attrs": [	30	"attrs": [
	@@ -53,8 +51,7 @@
53	"paramType": "required",	51	"paramType": "required",
54	"shape": [	52	"shape": [
55	-2	53	-2
56	- ],	54	+ ]
57	- "dtype_match_mode": "DtypeByte"
58	}	55	}
59	],	56	],
60	"outputs": [	57	"outputs": [
	@@ -66,8 +63,7 @@
66	"paramType": "required",	63	"paramType": "required",
67	"shape": [	64	"shape": [
68	-2	65	-2
69	- ],	66	+ ]
70	- "dtype_match_mode": "DtypeByte"
71	}	67	}
72	],	68	],
73	"attrs": [	69	"attrs": [
	@@ -94,8 +90,7 @@
94	"paramType": "required",	90	"paramType": "required",
95	"shape": [	91	"shape": [
96	-2	92	-2
97	- ],	93	+ ]
98	- "dtype_match_mode": "DtypeByte"
99	}	94	}
100	],	95	],
101	"outputs": [	96	"outputs": [
	@@ -107,8 +102,7 @@
107	"paramType": "required",	102	"paramType": "required",
108	"shape": [	103	"shape": [
109	-2	104	-2
110	- ],	105	+ ]
111	- "dtype_match_mode": "DtypeByte"
112	}	106	}
113	],	107	],
114	"attrs": [	108	"attrs": [
	@@ -135,8 +129,7 @@
135	"paramType": "required",	129	"paramType": "required",
136	"shape": [	130	"shape": [
137	-2	131	-2
138	- ],	132	+ ]
139	- "dtype_match_mode": "DtypeByte"
140	}	133	}
141	],	134	],
142	"outputs": [	135	"outputs": [
	@@ -148,8 +141,7 @@
148	"paramType": "required",	141	"paramType": "required",
149	"shape": [	142	"shape": [
150	-2	143	-2
151	- ],	144	+ ]
152	- "dtype_match_mode": "DtypeByte"
153	}	145	}
154	],	146	],
155	"attrs": [	147	"attrs": [
	@@ -176,8 +168,7 @@
176	"paramType": "required",	168	"paramType": "required",
177	"shape": [	169	"shape": [
178	-2	170	-2
179	- ],	171	+ ]
180	- "dtype_match_mode": "DtypeByte"
181	}	172	}
182	],	173	],
183	"outputs": [	174	"outputs": [
	@@ -189,8 +180,7 @@
189	"paramType": "required",	180	"paramType": "required",
190	"shape": [	181	"shape": [
191	-2	182	-2
192	- ],	183	+ ]
193	- "dtype_match_mode": "DtypeByte"
194	}	184	}
195	],	185	],
196	"attrs": [	186	"attrs": [
	@@ -217,8 +207,7 @@
217	"paramType": "required",	207	"paramType": "required",
218	"shape": [	208	"shape": [
219	-2	209	-2
220	- ],	210	+ ]
221	- "dtype_match_mode": "DtypeByte"
222	}	211	}
223	],	212	],
224	"outputs": [	213	"outputs": [
	@@ -230,8 +219,7 @@
230	"paramType": "required",	219	"paramType": "required",
231	"shape": [	220	"shape": [
232	-2	221	-2
233	- ],	222	+ ]
234	- "dtype_match_mode": "DtypeByte"
235	}	223	}
236	],	224	],
237	"attrs": [	225	"attrs": [

Mconversion/tile_with_axis/op_host/tile_with_axis_def.cpp+4-4

	@@ -37,11 +37,11 @@ public:
37	{ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})	37	{ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
38	.AutoContiguous();	38	.AutoContiguous();
39		39
40	- // 属性: axis (int64_t) -- 复制维度	40	+ // 属性: axis (int64_t) -- 复制维度（缺省按实现默认 1 处理，与 TBE/AICPU 一致）
41	- this->Attr("axis").AttrType(OPTIONAL).Int();	41	+ this->Attr("axis").AttrType(OPTIONAL).Int(1);
42		42
43	- // 属性: tiles (int64_t) -- 复制次数	43	+ // 属性: tiles (int64_t) -- 复制次数（必选，与 proto REQUIRED_ATTR / README / AICPU def 一致）
44	- this->Attr("tiles").AttrType(OPTIONAL).Int();	44	+ this->Attr("tiles").AttrType(REQUIRED).Int();
45		45
46	// Ascend950 (DAV_3510) 配置	46	// Ascend950 (DAV_3510) 配置
47	OpAICoreConfig aicoreConfig950;	47	OpAICoreConfig aicoreConfig950;

Mconversion/tile_with_axis/op_host/tile_with_axis_infershape.cpp+91-8

	@@ -14,18 +14,44 @@
14	*	14	*
15	* 设计依据: DESIGN.md v1.6	15	* 设计依据: DESIGN.md v1.6
16	* 输出 y 的 shape: y.shape[axis] = x.shape[axis] * tiles, 其余维度与 x 相同	16	* 输出 y 的 shape: y.shape[axis] = x.shape[axis] * tiles, 其余维度与 x 相同
17	- * rank=0 时 y.shape = [tiles]	17	+ * rank=0 时 y.shape = [tiles]（人工裁决：支持标量输入）
		18	+ * 动态 shape: axis 维为未知维标记（-1/-2）时原样传播，不参与乘法
18	*/	19	*/
19		20
20	#include "register/op_impl_registry.h"	21	#include "register/op_impl_registry.h"
21	#include "exe_graph/runtime/infer_shape_context.h"	22	#include "exe_graph/runtime/infer_shape_context.h"
		23	+#include "exe_graph/runtime/infer_shape_range_context.h"
22	#include "op_common/log/log.h"	24	#include "op_common/log/log.h"
23	#include "op_host/util/shape_util.h"	25	#include "op_host/util/shape_util.h"
24		26
		27	+#include <cstdint>
		28	+
25	using namespace ge;	29	using namespace ge;
26		30
27	namespace ops {	31	namespace ops {
28		32
		33	+namespace {
		34	+constexpr int64_t ATTR_AXIS_IDX = 0;
		35	+constexpr int64_t ATTR_TILES_IDX = 1;
		36	+// 仓库约定（对齐 circular_pad_infershape.cpp）：-1 为未知维标记；
		37	+// -2 为未知 rank 标记，known-rank shape 中出现时同样原样传播（对齐 canndev legacy 基线）。
		38	+constexpr int64_t UNKNOWN_DIM = -1;
		39	+constexpr int64_t UNKNOWN_RANK_DIM = -2;
		40	+
		41	+// axis 维输出值：未知维标记原样传播，其余乘 tiles 并做 int64 溢出检查
		42	+ge::graphStatus ScaleAxisDim(int64_t xDim, int64_t tiles, int64_t& yDim, const char* node_name)
		43	+{
		44	+ if (xDim == UNKNOWN_DIM \|\| xDim == UNKNOWN_RANK_DIM) {
		45	+ yDim = xDim;
		46	+ return ge::GRAPH_SUCCESS;
		47	+ }
		48	+ OP_CHECK_IF(xDim > 0 && xDim > INT64_MAX / tiles,
		49	+ OP_LOGE(node_name, "axis dim %ld * tiles %ld overflows int64", xDim, tiles), return ge::GRAPH_FAILED);
		50	+ yDim = xDim * tiles;
		51	+ return ge::GRAPH_SUCCESS;
		52	+}
		53	+} // namespace
		54	+
29	static ge::graphStatus InferShape4TileWithAxis(gert::InferShapeContext* context)	55	static ge::graphStatus InferShape4TileWithAxis(gert::InferShapeContext* context)
30	{	56	{
31	const gert::Shape* input_shape = context->GetInputShape(0);	57	const gert::Shape* input_shape = context->GetInputShape(0);
	@@ -35,13 +61,14 @@ static ge::graphStatus InferShape4TileWithAxis(gert::InferShapeContext* context)
35	OP_CHECK_NULL_WITH_CONTEXT(context, output_shape);	61	OP_CHECK_NULL_WITH_CONTEXT(context, output_shape);
36		62
37	// 获取属性 axis 和 tiles	63	// 获取属性 axis 和 tiles
38	- constexpr size_t ATTR_AXIS_IDX = 0;
39	- constexpr size_t ATTR_TILES_IDX = 1;
40	auto attrs = context->GetAttrs();	64	auto attrs = context->GetAttrs();
41	const int64_t* axis_ptr = attrs->GetAttrPointer<int64_t>(ATTR_AXIS_IDX);	65	const int64_t* axis_ptr = attrs->GetAttrPointer<int64_t>(ATTR_AXIS_IDX);
42	const int64_t* tiles_ptr = attrs->GetAttrPointer<int64_t>(ATTR_TILES_IDX);	66	const int64_t* tiles_ptr = attrs->GetAttrPointer<int64_t>(ATTR_TILES_IDX);
43	- int64_t axis = (axis_ptr != nullptr) ? *axis_ptr : 0;	67	+ // tiles 为必选属性（proto REQUIRED_ATTR / def REQUIRED / README）：缺失时拒绝，不静默回退
44	- int64_t tiles = (tiles_ptr != nullptr) ? *tiles_ptr : 1;	68	+ OP_CHECK_IF(tiles_ptr == nullptr, OP_LOGE(context, "tiles is required but missing"), return ge::GRAPH_FAILED);
		69	+ int64_t tiles = *tiles_ptr;
		70	+ // axis 默认 1（与 TBE/AICPU 缺省一致）；属性缺失时兜底 1
		71	+ int64_t axis = (axis_ptr != nullptr) ? *axis_ptr : 1;
45		72
46	OP_CHECK_IF(tiles <= 0, OP_LOGE(context, "tiles must be > 0, got %ld", tiles), return ge::GRAPH_FAILED);	73	OP_CHECK_IF(tiles <= 0, OP_LOGE(context, "tiles must be > 0, got %ld", tiles), return ge::GRAPH_FAILED);
47		74
	@@ -71,11 +98,14 @@ static ge::graphStatus InferShape4TileWithAxis(gert::InferShapeContext* context)
71		98
72	OP_CHECK_IF(axis < 0 \|\| axis >= rank, OP_LOGE(context, "axis is out of range"), return ge::GRAPH_FAILED);	99	OP_CHECK_IF(axis < 0 \|\| axis >= rank, OP_LOGE(context, "axis is out of range"), return ge::GRAPH_FAILED);
73		100
74	- // 复制 input shape 到 output shape	101	+ // 复制 input shape 到 output shape（axis 维: 未知维传播 + 溢出检查）
75	output_shape->SetDimNum(rank);	102	output_shape->SetDimNum(rank);
76	for (int64_t i = 0; i < rank; i++) {	103	for (int64_t i = 0; i < rank; i++) {
77	if (i == axis) {	104	if (i == axis) {
78	- output_shape->SetDim(i, input_shape->GetDim(i) * tiles);	105	+ int64_t yDim = 0;
		106	+ OP_CHECK_IF(ScaleAxisDim(input_shape->GetDim(i), tiles, yDim, context->GetNodeName()) != ge::GRAPH_SUCCESS,
		107	+ OP_LOGE(context, "scale axis dim failed"), return ge::GRAPH_FAILED);
		108	+ output_shape->SetDim(i, yDim);
79	} else {	109	} else {
80	output_shape->SetDim(i, input_shape->GetDim(i));	110	output_shape->SetDim(i, input_shape->GetDim(i));
81	}	111	}
	@@ -85,12 +115,65 @@ static ge::graphStatus InferShape4TileWithAxis(gert::InferShapeContext* context)
85	return ge::GRAPH_SUCCESS;	115	return ge::GRAPH_SUCCESS;
86	}	116	}
87		117
		118	+// 动态 shape 的 range 推导：输出 range = 输入 range，axis 维乘 tiles（未知维标记原样传播）
		119	+static ge::graphStatus InferShapeRange4TileWithAxis(gert::InferShapeRangeContext* context)
		120	+{
		121	+ auto attrs = context->GetAttrs();
		122	+ OP_CHECK_NULL_WITH_CONTEXT(context, attrs);
		123	+ const int64_t* axis_ptr = attrs->GetAttrPointer<int64_t>(ATTR_AXIS_IDX);
		124	+ const int64_t* tiles_ptr = attrs->GetAttrPointer<int64_t>(ATTR_TILES_IDX);
		125	+ OP_CHECK_IF(tiles_ptr == nullptr \|\| *tiles_ptr <= 0, OP_LOGE(context, "tiles is required and must be > 0"),
		126	+ return ge::GRAPH_FAILED);
		127	+ const int64_t tiles = *tiles_ptr;
		128	+ int64_t axis = (axis_ptr != nullptr) ? *axis_ptr : 1;
		129	+
		130	+ const auto* xRange = context->GetInputShapeRange(0);
		131	+ OP_CHECK_NULL_WITH_CONTEXT(context, xRange);
		132	+ auto* xRangeMin = xRange->GetMin();
		133	+ auto* xRangeMax = xRange->GetMax();
		134	+ OP_CHECK_NULL_WITH_CONTEXT(context, xRangeMin);
		135	+ OP_CHECK_NULL_WITH_CONTEXT(context, xRangeMax);
		136	+
		137	+ auto* yRange = context->GetOutputShapeRange(0);
		138	+ OP_CHECK_NULL_WITH_CONTEXT(context, yRange);
		139	+ auto* yRangeMin = yRange->GetMin();
		140	+ auto* yRangeMax = yRange->GetMax();
		141	+ OP_CHECK_NULL_WITH_CONTEXT(context, yRangeMin);
		142	+ OP_CHECK_NULL_WITH_CONTEXT(context, yRangeMax);
		143	+
		144	+ for (auto* yBound : {yRangeMin, yRangeMax}) {
		145	+ const auto* xBound = (yBound == yRangeMin) ? xRangeMin : xRangeMax;
		146	+ yBound = xBound;
		147	+ int64_t boundRank = static_cast<int64_t>(xBound->GetDimNum());
		148	+ if (boundRank == 0) {
		149	+ // 标量: 输出 range = [tiles, tiles]
		150	+ yBound->SetDimNum(1);
		151	+ yBound->SetDim(0, tiles);
		152	+ continue;
		153	+ }
		154	+ int64_t boundAxis = axis;
		155	+ if (boundAxis < 0) {
		156	+ boundAxis += boundRank;
		157	+ }
		158	+ OP_CHECK_IF(boundAxis < 0 \|\| boundAxis >= boundRank, OP_LOGE(context, "axis is out of range"),
		159	+ return ge::GRAPH_FAILED);
		160	+ int64_t yDim = 0;
		161	+ OP_CHECK_IF(ScaleAxisDim(xBound->GetDim(boundAxis), tiles, yDim, context->GetNodeName()) != ge::GRAPH_SUCCESS,
		162	+ OP_LOGE(context, "scale axis range dim failed"), return ge::GRAPH_FAILED);
		163	+ yBound->SetDim(boundAxis, yDim);
		164	+ }
		165	+ return ge::GRAPH_SUCCESS;
		166	+}
		167	+
88	static ge::graphStatus InferDataType4TileWithAxis(gert::InferDataTypeContext* context)	168	static ge::graphStatus InferDataType4TileWithAxis(gert::InferDataTypeContext* context)
89	{	169	{
90	context->SetOutputDataType(0, context->GetInputDataType(0));	170	context->SetOutputDataType(0, context->GetInputDataType(0));
91	return ge::GRAPH_SUCCESS;	171	return ge::GRAPH_SUCCESS;
92	}	172	}
93		173
94	-IMPL_OP_INFERSHAPE(TileWithAxis).InferShape(InferShape4TileWithAxis).InferDataType(InferDataType4TileWithAxis);	174	+IMPL_OP_INFERSHAPE(TileWithAxis)
		175	+ .InferShape(InferShape4TileWithAxis)
		176	+ .InferDataType(InferDataType4TileWithAxis)
		177	+ .InferShapeRange(InferShapeRange4TileWithAxis);
95		178
96	} // namespace ops	179	} // namespace ops

Mconversion/tile_with_axis/op_kernel/arch35/tile_with_axis.h+88-103

	@@ -42,26 +42,29 @@
42		42
43	using namespace AscendC;	43	using namespace AscendC;
44		44
45	-// CeilDiv: ceil(a / b) for positive integers	45	+// CeilDiv: ceil(a / b) for positive integers, without a+b-1 overflow
46	template <typename IntT>	46	template <typename IntT>
47	-__aicore__ inline IntT CeilDiv(IntT a, IntT b) { return (a + b - 1) / b; }	47	+__aicore__ inline IntT CeilDiv(IntT a, IntT b)
		48	+{
		49	+ return a / b + (a % b != 0);
		50	+}
48		51
49	// NDDMA 维度: 3 维 (ubAxis=0 用满 3 维; ubAxis=1 退化为 2 维; ubAxis=2 不使用)	52	// NDDMA 维度: 3 维 (ubAxis=0 用满 3 维; ubAxis=1 退化为 2 维; ubAxis=2 不使用)
50	static constexpr uint8_t NDDMA_DIM = 3;	53	static constexpr uint8_t NDDMA_DIM = 3;
51	static constexpr int AXIS_COUNT = 3;	54	static constexpr int AXIS_COUNT = 3;
52		55
53	struct AxisLayout {	56	struct AxisLayout {
54	- int64_t inStart; // 输入起始坐标	57	+ int64_t inStart; // 输入起始坐标
55	- int64_t outStart; // 输出起始坐标	58	+ int64_t outStart; // 输出起始坐标
56	- int64_t length; // 搬运长度	59	+ int64_t length; // 搬运长度
57	- int64_t inStride; // 输入 stride (InitLayout 一次性计算)	60	+ int64_t inStride; // 输入 stride (InitLayout 一次性计算)
58	- int64_t outStride; // 输出 stride (InitLayout 一次性计算)	61	+ int64_t outStride; // 输出 stride (InitLayout 一次性计算)
59	};	62	};
60		63
61	struct TileLayout {	64	struct TileLayout {
62	AxisLayout axes[AXIS_COUNT];	65	AxisLayout axes[AXIS_COUNT];
63	- int64_t inOffset; // = SUM(inStart[axis] * inStride[axis])	66	+ int64_t inOffset; // = SUM(inStart[axis] * inStride[axis])
64	- int64_t outOffset; // = SUM(outStart[axis] * outStride[axis])	67	+ int64_t outOffset; // = SUM(outStart[axis] * outStride[axis])
65	};	68	};
66		69
67	template <typename T, int UB_AXIS>	70	template <typename T, int UB_AXIS>
	@@ -70,8 +73,7 @@ public:
70	// ================================================================	73	// ================================================================
71	// Init: 一次性完成 stride 计算、buffer 申请、NDDMA 参数预计算	74	// Init: 一次性完成 stride 计算、buffer 申请、NDDMA 参数预计算
72	// ================================================================	75	// ================================================================
73	- __aicore__ inline void Init(GM_ADDR x, GM_ADDR y,	76	+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, const TileWithAxisTilingData* tilingData)
74	- const TileWithAxisTilingData* tilingData)
75	{	77	{
76	td_ = tilingData;	78	td_ = tilingData;
77		79
	@@ -83,9 +85,9 @@ public:
83	InitLayout();	85	InitLayout();
84	InitNddmaParams();	86	InitNddmaParams();
85		87
86	- uint64_t blockIdx = GetBlockIdx();	88	+ uint64_t blockIdx = GetBlockIdx();
87	startBlock_ = blockIdx * td_->perCoreCount;	89	startBlock_ = blockIdx * td_->perCoreCount;
88	- endBlock_ = min(startBlock_ + td_->perCoreCount, td_->totalCount);	90	+ endBlock_ = min(startBlock_ + td_->perCoreCount, td_->totalCount);
89	}	91	}
90		92
91	// ================================================================	93	// ================================================================
	@@ -120,13 +122,13 @@ private:
120	uint64_t endBlock_;	122	uint64_t endBlock_;
121		123
122	// 预计算常量	124	// 预计算常量
123	- int64_t rowInLength_; // = axisDim * innerDim = rowLength	125	+ int64_t rowInLength_; // = axisDim * innerDim = rowLength
124	- int64_t rowOutLength_; // = tiles * rowLength	126	+ int64_t rowOutLength_; // = tiles * rowLength
125		127
126	// NDDMA 预计算参数 (仅 ubAxis=0/1 使用)	128	// NDDMA 预计算参数 (仅 ubAxis=0/1 使用)
127	AscendC::NdDmaLoopInfo<NDDMA_DIM> nddmaLoopInfo_;	129	AscendC::NdDmaLoopInfo<NDDMA_DIM> nddmaLoopInfo_;
128	- int64_t nddmaBlockLen_; // NDDMA 每次连续搬运的元素数	130	+ int64_t nddmaBlockLen_; // NDDMA 每次连续搬运的元素数
129	- int64_t nddmaRepeat_; // NDDMA dim1 重复次数	131	+ int64_t nddmaRepeat_; // NDDMA dim1 重复次数
130		132
131	// ================================================================	133	// ================================================================
132	// InitLayout: 基于 3D 模型 inShape=[outerDim,1,rowLength] outShape=[outerDim,tiles,rowLength]	134	// InitLayout: 基于 3D 模型 inShape=[outerDim,1,rowLength] outShape=[outerDim,tiles,rowLength]
	@@ -134,16 +136,16 @@ private:
134	__aicore__ inline void InitLayout()	136	__aicore__ inline void InitLayout()
135	{	137	{
136	// 输入 stride: inShape = [outerDim, 1, rowLength]	138	// 输入 stride: inShape = [outerDim, 1, rowLength]
137	- layout_.axes[2].inStride = 1; // rowLength 维 stride=1	139	+ layout_.axes[2].inStride = 1; // rowLength 维 stride=1
138	- layout_.axes[1].inStride = td_->inShape[2]; // 1 维 stride = rowLength	140	+ layout_.axes[1].inStride = td_->inShape[2]; // 1 维 stride = rowLength
139	- layout_.axes[0].inStride = td_->inShape[1] * td_->inShape[2]; // outer 维 stride = 1 * rowLength	141	+ layout_.axes[0].inStride = td_->inShape[1] * td_->inShape[2]; // outer 维 stride = 1 * rowLength
140		142
141	// 输出 stride: outShape = [outerDim, tiles, rowLength]	143	// 输出 stride: outShape = [outerDim, tiles, rowLength]
142	- layout_.axes[2].outStride = 1; // rowLength 维 stride=1	144	+ layout_.axes[2].outStride = 1; // rowLength 维 stride=1
143	- layout_.axes[1].outStride = td_->outShape[2]; // tiles 维 stride = rowLength	145	+ layout_.axes[1].outStride = td_->outShape[2]; // tiles 维 stride = rowLength
144	- layout_.axes[0].outStride = td_->outShape[1] * td_->outShape[2]; // outer 维 stride = tiles * rowLength	146	+ layout_.axes[0].outStride = td_->outShape[1] * td_->outShape[2]; // outer 维 stride = tiles * rowLength
145		147
146	- rowInLength_ = td_->rowLength;	148	+ rowInLength_ = td_->rowLength;
147	rowOutLength_ = td_->tiles * td_->rowLength;	149	rowOutLength_ = td_->tiles * td_->rowLength;
148	}	150	}
149		151
	@@ -156,11 +158,11 @@ private:
156	__aicore__ inline void InitNddmaParams()	158	__aicore__ inline void InitNddmaParams()
157	{	159	{
158	if constexpr (UB_AXIS > 1) {	160	if constexpr (UB_AXIS > 1) {
159	- return; // UB_AXIS=2 不使用 NDDMA	161	+ return; // UB_AXIS=2 不使用 NDDMA
160	}	162	}
161	// ---- 公共部分 ----	163	// ---- 公共部分 ----
162	nddmaBlockLen_ = rowInLength_;	164	nddmaBlockLen_ = rowInLength_;
163	- nddmaRepeat_ = td_->ubFactor; // UB_AXIS=1 默认值	165	+ nddmaRepeat_ = td_->ubFactor; // UB_AXIS=1 默认值
164	nddmaLoopInfo_.loopSrcStride[0] = 1;	166	nddmaLoopInfo_.loopSrcStride[0] = 1;
165	nddmaLoopInfo_.loopSrcStride[1] = 0;	167	nddmaLoopInfo_.loopSrcStride[1] = 0;
166	nddmaLoopInfo_.loopSrcStride[2] = 0;	168	nddmaLoopInfo_.loopSrcStride[2] = 0;
	@@ -170,14 +172,18 @@ private:
170	nddmaLoopInfo_.loopSize[0] = nddmaBlockLen_;	172	nddmaLoopInfo_.loopSize[0] = nddmaBlockLen_;
171	nddmaLoopInfo_.loopSize[1] = nddmaRepeat_;	173	nddmaLoopInfo_.loopSize[1] = nddmaRepeat_;
172	nddmaLoopInfo_.loopSize[2] = 1;	174	nddmaLoopInfo_.loopSize[2] = 1;
173	- nddmaLoopInfo_.loopLpSize[0] = 0; nddmaLoopInfo_.loopLpSize[1] = 0; nddmaLoopInfo_.loopLpSize[2] = 0;	175	+ nddmaLoopInfo_.loopLpSize[0] = 0;
174	- nddmaLoopInfo_.loopRpSize[0] = 0; nddmaLoopInfo_.loopRpSize[1] = 0; nddmaLoopInfo_.loopRpSize[2] = 0;	176	+ nddmaLoopInfo_.loopLpSize[1] = 0;
		177	+ nddmaLoopInfo_.loopLpSize[2] = 0;
		178	+ nddmaLoopInfo_.loopRpSize[0] = 0;
		179	+ nddmaLoopInfo_.loopRpSize[1] = 0;
		180	+ nddmaLoopInfo_.loopRpSize[2] = 0;
175	// ---- UB_AXIS=0 覆盖 dim2 参数（注意 loopSize[1] 需同步 nddmaRepeat_）----	181	// ---- UB_AXIS=0 覆盖 dim2 参数（注意 loopSize[1] 需同步 nddmaRepeat_）----
176	if constexpr (UB_AXIS == 0) {	182	if constexpr (UB_AXIS == 0) {
177	- nddmaRepeat_ = td_->tiles;	183	+ nddmaRepeat_ = td_->tiles;
178	nddmaLoopInfo_.loopSrcStride[2] = nddmaBlockLen_;	184	nddmaLoopInfo_.loopSrcStride[2] = nddmaBlockLen_;
179	nddmaLoopInfo_.loopDstStride[2] = nddmaBlockLen_ * nddmaRepeat_;	185	nddmaLoopInfo_.loopDstStride[2] = nddmaBlockLen_ * nddmaRepeat_;
180	- nddmaLoopInfo_.loopSize[1] = nddmaRepeat_; // 同步为 tiles	186	+ nddmaLoopInfo_.loopSize[1] = nddmaRepeat_; // 同步为 tiles
181	nddmaLoopInfo_.loopSize[2] = td_->ubFactor;	187	nddmaLoopInfo_.loopSize[2] = td_->ubFactor;
182	}	188	}
183	}	189	}
	@@ -200,90 +206,78 @@ private:
200	__aicore__ inline void UpdateLayoutAxis0(uint64_t blockIdx)	206	__aicore__ inline void UpdateLayoutAxis0(uint64_t blockIdx)
201	{	207	{
202	int64_t outerStart = static_cast<int64_t>(td_->ubFactor) * static_cast<int64_t>(blockIdx);	208	int64_t outerStart = static_cast<int64_t>(td_->ubFactor) * static_cast<int64_t>(blockIdx);
203	- int64_t outerCount = min(	209	+ int64_t outerCount = min(static_cast<int64_t>(td_->ubFactor), td_->outShape[0] - outerStart);
204	- static_cast<int64_t>(td_->ubFactor),
205	- td_->outShape[0] - outerStart);
206		210
207	- layout_.axes[0].inStart = outerStart;	211	+ layout_.axes[0].inStart = outerStart;
208	layout_.axes[0].outStart = outerStart;	212	layout_.axes[0].outStart = outerStart;
209	- layout_.axes[0].length = outerCount;	213	+ layout_.axes[0].length = outerCount;
210	- layout_.axes[1].length = td_->outShape[1]; // tiles	214	+ layout_.axes[1].length = td_->outShape[1]; // tiles
211	- layout_.axes[2].length = td_->outShape[2]; // rowLength	215	+ layout_.axes[2].length = td_->outShape[2]; // rowLength
212		216
213	- layout_.inOffset = outerStart * layout_.axes[0].inStride;	217	+ layout_.inOffset = outerStart * layout_.axes[0].inStride;
214	- layout_.outOffset = outerStart * layout_.axes[0].outStride;	218	+ layout_.outOffset = outerStart * layout_.axes[0].outStride;
215		219
216	- ubLen_ = outerCount * rowOutLength_;	220	+ ubLen_ = outerCount * rowOutLength_;
217	- outGmBase_ = layout_.outOffset;	221	+ outGmBase_ = layout_.outOffset;
218	- inGmBase_ = layout_.inOffset;	222	+ inGmBase_ = layout_.inOffset;
219	}	223	}
220		224
221	// ---- 切 tiles (outShape=[outerDim,tiles,rowLength]) ----	225	// ---- 切 tiles (outShape=[outerDim,tiles,rowLength]) ----
222	__aicore__ inline void UpdateLayoutAxis1(uint64_t blockIdx)	226	__aicore__ inline void UpdateLayoutAxis1(uint64_t blockIdx)
223	{	227	{
224	- int64_t tileBlocksPerOuter = CeilDiv(	228	+ int64_t tileBlocksPerOuter = CeilDiv(td_->outShape[1], static_cast<int64_t>(td_->ubFactor));
225	- td_->outShape[1], static_cast<int64_t>(td_->ubFactor));	229	+ int64_t outerIdx = static_cast<int64_t>(blockIdx) / tileBlocksPerOuter;
226	- int64_t outerIdx = static_cast<int64_t>(blockIdx) / tileBlocksPerOuter;	230	+ int64_t tileStart = static_cast<int64_t>(td_->ubFactor) * (static_cast<int64_t>(blockIdx) % tileBlocksPerOuter);
227	- int64_t tileStart = static_cast<int64_t>(td_->ubFactor)	231	+ int64_t tileCount = min(static_cast<int64_t>(td_->ubFactor), td_->outShape[1] - tileStart);
228	- * (static_cast<int64_t>(blockIdx) % tileBlocksPerOuter);
229	- int64_t tileCount = min(
230	- static_cast<int64_t>(td_->ubFactor),
231	- td_->outShape[1] - tileStart);
232		232
233	- layout_.axes[0].inStart = outerIdx;	233	+ layout_.axes[0].inStart = outerIdx;
234	layout_.axes[0].outStart = outerIdx;	234	layout_.axes[0].outStart = outerIdx;
235	- layout_.axes[0].length = 1;	235	+ layout_.axes[0].length = 1;
236	layout_.axes[1].outStart = tileStart;	236	layout_.axes[1].outStart = tileStart;
237	- layout_.axes[1].length = tileCount;	237	+ layout_.axes[1].length = tileCount;
238	- layout_.axes[2].length = td_->outShape[2]; // rowLength	238	+ layout_.axes[2].length = td_->outShape[2]; // rowLength
239		239
240	- layout_.outOffset = outerIdx * layout_.axes[0].outStride	240	+ layout_.outOffset = outerIdx * layout_.axes[0].outStride + tileStart * layout_.axes[1].outStride;
241	- + tileStart * layout_.axes[1].outStride;
242		241
243	// 输入: 中间维恒为 1, tile 维 = 0	242	// 输入: 中间维恒为 1, tile 维 = 0
244	layout_.inOffset = outerIdx * layout_.axes[0].inStride;	243	layout_.inOffset = outerIdx * layout_.axes[0].inStride;
245		244
246	- ubLen_ = tileCount * td_->outShape[2];	245	+ ubLen_ = tileCount * td_->outShape[2];
247	- outGmBase_ = layout_.outOffset;	246	+ outGmBase_ = layout_.outOffset;
248	- inGmBase_ = layout_.inOffset;	247	+ inGmBase_ = layout_.inOffset;
249	}	248	}
250		249
251	// ---- 切 rowLength (outShape=[outerDim,tiles,rowLength])	250	// ---- 切 rowLength (outShape=[outerDim,tiles,rowLength])
252	// 每 block = 1 outer 行 x 1 tile x ubFactor 个连续元素 ----	251	// 每 block = 1 outer 行 x 1 tile x ubFactor 个连续元素 ----
253	__aicore__ inline void UpdateLayoutAxis2(uint64_t blockIdx)	252	__aicore__ inline void UpdateLayoutAxis2(uint64_t blockIdx)
254	{	253	{
255	- int64_t rowBlocksPerTile = CeilDiv(	254	+ int64_t rowBlocksPerTile = CeilDiv(td_->outShape[2], static_cast<int64_t>(td_->ubFactor));
256	- td_->outShape[2], static_cast<int64_t>(td_->ubFactor));	255	+ int64_t blocksPerOuter = td_->outShape[1] * rowBlocksPerTile;
257	- int64_t blocksPerOuter = td_->outShape[1] * rowBlocksPerTile;	256	+ int64_t outerIdx = static_cast<int64_t>(blockIdx) / blocksPerOuter;
258	- int64_t outerIdx = static_cast<int64_t>(blockIdx) / blocksPerOuter;	257	+ int64_t remainder = static_cast<int64_t>(blockIdx) % blocksPerOuter;
259	- int64_t remainder = static_cast<int64_t>(blockIdx) % blocksPerOuter;	258	+ int64_t tileIdx = remainder / rowBlocksPerTile;
260	- int64_t tileIdx = remainder / rowBlocksPerTile;	259	+ int64_t innerBlock = remainder % rowBlocksPerTile;
261	- int64_t innerBlock = remainder % rowBlocksPerTile;	260	+ int64_t innerStart = static_cast<int64_t>(td_->ubFactor) * innerBlock;
262	- int64_t innerStart = static_cast<int64_t>(td_->ubFactor) * innerBlock;	261	+ int64_t innerCount = min(static_cast<int64_t>(td_->ubFactor), td_->outShape[2] - innerStart);
263	- int64_t innerCount = min(
264	- static_cast<int64_t>(td_->ubFactor),
265	- td_->outShape[2] - innerStart);
266		262
267	// 输入: 中间维=1 (tile=0)	263	// 输入: 中间维=1 (tile=0)
268	- layout_.axes[0].inStart = outerIdx;	264	+ layout_.axes[0].inStart = outerIdx;
269	layout_.axes[0].outStart = outerIdx;	265	layout_.axes[0].outStart = outerIdx;
270	- layout_.axes[0].length = 1;	266	+ layout_.axes[0].length = 1;
271	layout_.axes[1].outStart = tileIdx;	267	layout_.axes[1].outStart = tileIdx;
272	- layout_.axes[1].inStart = 0; // 输入中间维=1	268	+ layout_.axes[1].inStart = 0; // 输入中间维=1
273	- layout_.axes[1].length = 1; // 仅 1 个 tile	269	+ layout_.axes[1].length = 1; // 仅 1 个 tile
274	- layout_.axes[2].inStart = innerStart;	270	+ layout_.axes[2].inStart = innerStart;
275	layout_.axes[2].outStart = innerStart;	271	layout_.axes[2].outStart = innerStart;
276	- layout_.axes[2].length = innerCount;	272	+ layout_.axes[2].length = innerCount;
277		273
278	- layout_.inOffset = outerIdx * layout_.axes[0].inStride	274	+ layout_.inOffset = outerIdx * layout_.axes[0].inStride + innerStart * layout_.axes[2].inStride;
279	- + innerStart * layout_.axes[2].inStride;	275	+ layout_.outOffset = outerIdx * layout_.axes[0].outStride + tileIdx * layout_.axes[1].outStride +
280	- layout_.outOffset = outerIdx * layout_.axes[0].outStride	276	+ innerStart * layout_.axes[2].outStride;
281	- + tileIdx * layout_.axes[1].outStride
282	- + innerStart * layout_.axes[2].outStride;
283		277
284	- ubLen_ = innerCount;	278	+ ubLen_ = innerCount;
285	- outGmBase_ = layout_.outOffset;	279	+ outGmBase_ = layout_.outOffset;
286	- inGmBase_ = layout_.inOffset;	280	+ inGmBase_ = layout_.inOffset;
287	}	281	}
288		282
289	// ================================================================	283	// ================================================================
	@@ -303,20 +297,15 @@ private:
303	}	297	}
304	NdDmaDci();	298	NdDmaDci();
305	AscendC::NdDmaParams<T, NDDMA_DIM> params = {loopInfo, 0};	299	AscendC::NdDmaParams<T, NDDMA_DIM> params = {loopInfo, 0};
306	- AscendC::DataCopy<T, NDDMA_DIM>(	300	+ AscendC::DataCopy<T, NDDMA_DIM>(ubLocal[bufOffset], xGm_[inGmBase_], params);
307	- ubLocal[bufOffset],
308	- xGm_[inGmBase_],
309	- params);
310	} else {	301	} else {
311	DataCopyExtParams copyIn;	302	DataCopyExtParams copyIn;
312	copyIn.blockCount = 1;	303	copyIn.blockCount = 1;
313	- copyIn.blockLen = static_cast<uint32_t>(ubLen_ * sizeof(T));	304	+ copyIn.blockLen = static_cast<uint32_t>(ubLen_ * sizeof(T));
314	- copyIn.srcStride = 0;	305	+ copyIn.srcStride = 0;
315	- copyIn.dstStride = 0;	306	+ copyIn.dstStride = 0;
316	DataCopyPadExtParams<T> padParams{false, 0, 0, 0};	307	DataCopyPadExtParams<T> padParams{false, 0, 0, 0};
317	- DataCopyPad(ubLocal[bufOffset],	308	+ DataCopyPad(ubLocal[bufOffset], xGm_[inGmBase_], copyIn, padParams);
318	- xGm_[inGmBase_],
319	- copyIn, padParams);
320	}	309	}
321	}	310	}
322		311
	@@ -326,12 +315,8 @@ private:
326	// ================================================================	315	// ================================================================
327	__aicore__ inline void CopyOut(const LocalTensor<T>& ubLocal, uint64_t bufOffset)	316	__aicore__ inline void CopyOut(const LocalTensor<T>& ubLocal, uint64_t bufOffset)
328	{	317	{
329	- DataCopyExtParams copyOut{	318	+ DataCopyExtParams copyOut{1, static_cast<uint32_t>(ubLen_ * sizeof(T)), 0, 0, 0};
330	- 1, static_cast<uint32_t>(ubLen_ * sizeof(T)), 0, 0, 0	319	+ DataCopyPad(yGm_[outGmBase_], ubLocal[bufOffset], copyOut);
331	- };
332	- DataCopyPad(yGm_[outGmBase_],
333	- ubLocal[bufOffset],
334	- copyOut);
335	}	320	}
336		321
337	// ================================================================	322	// ================================================================
	@@ -349,9 +334,9 @@ private:
349	// ================================================================	334	// ================================================================
350	// 每 Block 运行时变量	335	// 每 Block 运行时变量
351	// ================================================================	336	// ================================================================
352	- int64_t ubLen_; // 本 Block 输出元素数	337	+ int64_t ubLen_; // 本 Block 输出元素数
353	- int64_t outGmBase_; // 输出 GM 基址	338	+ int64_t outGmBase_; // 输出 GM 基址
354	- int64_t inGmBase_; // 输入 GM 基址	339	+ int64_t inGmBase_; // 输入 GM 基址
355	};	340	};
356		341
357	#endif // TILE_WITH_AXIS_H	342	#endif // TILE_WITH_AXIS_H

Mconversion/tile_with_axis/op_kernel/arch35/tile_with_axis_tiling_data.h+9-9

	@@ -28,17 +28,17 @@
28		28
29	struct TileWithAxisTilingData {	29	struct TileWithAxisTilingData {
30	// ==== A. LayoutTransform 单切分公共字段 ====	30	// ==== A. LayoutTransform 单切分公共字段 ====
31	- int64_t inShape[3]; // 输入展平 shape: [outerDim, 1, rowLength]	31	+ int64_t inShape[3]; // 输入展平 shape: [outerDim, 1, rowLength]
32	- int64_t outShape[3]; // 输出展平 shape: [outerDim, tiles, rowLength]	32	+ int64_t outShape[3]; // 输出展平 shape: [outerDim, tiles, rowLength]
33	- uint64_t totalCount; // UB 切分块总数（基于 outShape 计算）	33	+ uint64_t totalCount; // UB 切分块总数（基于 outShape 计算）
34	- uint64_t perCoreCount; // 每核处理的块数（totalCount 均分给 coreNum 个核）	34	+ uint64_t perCoreCount; // 每核处理的块数（totalCount 均分给 coreNum 个核）
35	- uint8_t ubAxis; // UB 切分轴（0=outerDim, 1=tiles, 2=rowLength）	35	+ uint8_t ubAxis; // UB 切分轴（0=outerDim, 1=tiles, 2=rowLength）
36	- uint32_t ubFactor; // 切分因子（沿 ubAxis 每次搬运该轴多少个元素）	36	+ uint32_t ubFactor; // 切分因子（沿 ubAxis 每次搬运该轴多少个元素）
37	- uint32_t bufferSize; // UB buffer 字节大小（ping-pong 单个 buffer 的字节数）	37	+ uint32_t bufferSize; // UB buffer 字节大小（ping-pong 单个 buffer 的字节数）
38		38
39	// ==== B. TileWithAxis 特有字段 ====	39	// ==== B. TileWithAxis 特有字段 ====
40	- int64_t tiles; // 复制次数	40	+ int64_t tiles; // 复制次数
41	- int64_t rowLength; // 折叠后每行长度 = axisDim * innerDim	41	+ int64_t rowLength; // 折叠后每行长度 = axisDim * innerDim
42	};	42	};
43		43
44	#endif // _TILE_WITH_AXIS_TILING_DATA_H_	44	#endif // _TILE_WITH_AXIS_TILING_DATA_H_

Mconversion/tile_with_axis/op_kernel/arch35/tile_with_axis_tiling_key.h+2-8

	@@ -22,14 +22,8 @@
22		22
23	#include "ascendc/host_api/tiling/template_argument.h"	23	#include "ascendc/host_api/tiling/template_argument.h"
24		24
25	-ASCENDC_TPL_ARGS_DECL(TileWithAxis,	25	+ASCENDC_TPL_ARGS_DECL(TileWithAxis, ASCENDC_TPL_UINT_DECL(UB_AXIS, 8, ASCENDC_TPL_UI_LIST, 0, 1, 2));
26	- ASCENDC_TPL_UINT_DECL(UB_AXIS, 8, ASCENDC_TPL_UI_LIST, 0, 1, 2)
27	-);
28		26
29	-ASCENDC_TPL_SEL(	27	+ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_UINT_SEL(UB_AXIS, ASCENDC_TPL_UI_LIST, 0, 1, 2)));
30	- ASCENDC_TPL_ARGS_SEL(
31	- ASCENDC_TPL_UINT_SEL(UB_AXIS, ASCENDC_TPL_UI_LIST, 0, 1, 2)
32	- )
33	-);
34		28
35	#endif // __TILE_WITH_AXIS_TILING_KEY_H__	29	#endif // __TILE_WITH_AXIS_TILING_KEY_H__

Mconversion/tile_with_axis/op_kernel/tile_with_axis_apt.cpp+2-4

	@@ -25,8 +25,7 @@
25	using namespace AscendC;	25	using namespace AscendC;
26		26
27	template <int UB_AXIS>	27	template <int UB_AXIS>
28	-__aicore__ inline void LaunchKernel(GM_ADDR x, GM_ADDR y,	28	+__aicore__ inline void LaunchKernel(GM_ADDR x, GM_ADDR y, const TileWithAxisTilingData* td)
29	- const TileWithAxisTilingData* td)
30	{	29	{
31	TileWithAxisKernel<DTYPE_X, UB_AXIS> op;	30	TileWithAxisKernel<DTYPE_X, UB_AXIS> op;
32	op.Init(x, y, td);	31	op.Init(x, y, td);
	@@ -34,8 +33,7 @@ __aicore__ inline void LaunchKernel(GM_ADDR x, GM_ADDR y,
34	}	33	}
35		34
36	template <int UB_AXIS>	35	template <int UB_AXIS>
37	-__global__ __aicore__ void tile_with_axis(	36	+__global__ __aicore__ void tile_with_axis(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
38	- GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
39	{	37	{
40	SetSysWorkspace(workspace);	38	SetSysWorkspace(workspace);
41	KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);	39	KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);

Aconversion/tile_with_axis/tests/assets/golden.py+149-0

@@ -0,0 +1,149 @@
	1	+#!/usr/bin/env python3
	2	+# -- coding: UTF-8 --
	3	+# ----------------------------------------------------------------------------
	4	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	5	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	6	+# CANN Open Software License Agreement Version 2.0 (the "License").
	7	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	8	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	9	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	10	+# See LICENSE in the root of the software repository for the full text of the License.
	11	+# ----------------------------------------------------------------------------
	12	+"""TileWithAxis 独立 Golden：Kernel/GEIR。
	13	+
	14	+沿指定 axis 将输入 x 复制 tiles 份并整块拼接（等价 concat([x]*tiles, dim=axis)），
	15	+仅 Caffe Tile 与之一一对应；PyTorch/TF 均无同名单算子 API
	16	+（torch.repeat / tf.tile 为全维度 multiples 形式，torch.repeat_interleave /
	17	+tf.repeat 为交错复制、语义不符）。
	18	+因此 golden 与三方标杆均以 torch 小算子拼接实现，并通过
	19	+torch_graph（torch.nn.Module）提供 TTK E2E torch.compile 模式的图封装。
	20	+
	21	+TTK kernel 模式消费：
	22	+ python3 -m ttk kernel -i xxx.csv --plugin tile_with_axis_golden_torch.py
	23	+
	24	+注册名 tile_with_axis 与 CSV 的 op_name 一致；
	25	+golden 入参为 numpy.ndarray（kernel 流程），返回 list（每个输出一个元素）。
	26	+"""
	27	+
	28	+import numpy as np
	29	+import torch
	30	+
	31	+__spec__ = {"tile_with_axis": "TileWithAxisTestSpec"}
	32	+
	33	+# 纯拷贝算子容差口径：
	34	+# - TTK tolerance 声明（本文件 _TOLERANCE）：整型 binary_equal；浮点同为 binary_equal
	35	+# （拷贝语义下 golden 与 kernel 输出应逐位一致，为最严声明）。
	36	+# - spec.yaml numerical_tolerance 与 CSV precision_tolerances 对浮点另声明 rtol/atol
	37	+# （fp16 1e-3 / bf16 4e-3 / fp32 1e-5，stat_rel_err 口径），为三方精度报告的宽容差上限；
	38	+# 实际 run 收据按 stat_rel_err 记录。两者宽严不同但均为拷贝语义下的合法声明，不构成冲突。
	39	+_TOLERANCE = {
	40	+ "float16": {"standard": "binary_equal"},
	41	+ "bfloat16": {"standard": "binary_equal"},
	42	+ "float32": {"standard": "binary_equal"},
	43	+ "int8": {"standard": "binary_equal"},
	44	+ "int16": {"standard": "binary_equal"},
	45	+ "int32": {"standard": "binary_equal"},
	46	+ "int64": {"standard": "binary_equal"},
	47	+ "uint8": {"standard": "binary_equal"},
	48	+ "uint16": {"standard": "binary_equal"},
	49	+ "uint32": {"standard": "binary_equal"},
	50	+ "uint64": {"standard": "binary_equal"},
	51	+}
	52	+
	53	+
	54	+def _tile_with_axis_compute(x, axis, tiles):
	55	+ """核心计算（torch 小算子拼接）：y = cat([x]*tiles, dim=axis)。
	56	+
	57	+ 整块复制拼接、非交错；负索引先规约到 [0, rank)。
	58	+ rank=0 标量输入（人工裁决 2026-10-20：支持）：输出 [tiles]，axis 不生效。
	59	+ """
	60	+ axis = int(axis)
	61	+ tiles = int(tiles)
	62	+ if tiles <= 0:
	63	+ raise ValueError(f"tiles must be > 0, got {tiles}")
	64	+ if x.dim() == 0:
	65	+ return x.repeat(tiles)
	66	+ if axis < 0:
	67	+ axis += x.dim()
	68	+ return torch.cat([x] * tiles, dim=axis)
	69	+
	70	+
	71	+def _numpy_to_torch(np_array):
	72	+ if np_array is None:
	73	+ return None
	74	+ if "bfloat16" in str(np_array.dtype):
	75	+ return torch.from_numpy(np_array.view(np.int16)).view(torch.bfloat16)
	76	+ return torch.from_numpy(np_array)
	77	+
	78	+
	79	+def _torch_to_numpy(torch_tensor, orig_dtype=None):
	80	+ if torch_tensor is None:
	81	+ return None
	82	+ if torch_tensor.dtype == torch.bfloat16:
	83	+ target = orig_dtype if orig_dtype is not None else np.dtype(np.int16)
	84	+ return torch_tensor.view(torch.int16).numpy().view(target)
	85	+ return torch_tensor.numpy()
	86	+
	87	+
	88	+def tile_with_axis_golden(x, , axis, tiles, *kwargs):
	89	+ """Golden function for tile_with_axis (kernel mode, numpy in / numpy out)。"""
	90	+ del kwargs
	91	+ orig_dtype = x.dtype
	92	+ x_torch = _numpy_to_torch(x)
	93	+ result = _tile_with_axis_compute(x_torch, axis, tiles)
	94	+ return [_torch_to_numpy(result, orig_dtype)]
	95	+
	96	+
	97	+class _TorchTileWithAxis:
	98	+ """竞品 third_party：torch 小算子拼接实现（无一一对应竞品算子）。
	99	+
	100	+ __init__ 吸收 CSV attributes（axis / tiles）；
	101	+ __call__ 直接使用输入张量计算，无多余 .to() 等数据转换，保证性能统计准确。
	102	+ 输入支持位置参数 torch.Tensor 或按名 kwargs（x / input）。
	103	+ """
	104	+
	105	+ def __init__(self, , axis=1, tiles=1, *kwargs):
	106	+ del kwargs
	107	+ self.axis = axis
	108	+ self.tiles = tiles
	109	+
	110	+ def __call__(self, args, *kwargs):
	111	+ x = next((a for a in args if isinstance(a, torch.Tensor)), None)
	112	+ if x is None:
	113	+ x = kwargs.get("x", kwargs.get("input"))
	114	+ if x is None:
	115	+ raise ValueError(
	116	+ "Missing tensor input; expected positional Tensor or one of: x, input"
	117	+ )
	118	+ return [_tile_with_axis_compute(x, self.axis, self.tiles)]
	119	+
	120	+
	121	+class TileWithAxisGraph(torch.nn.Module):
	122	+ """Graph module consumed by TTK E2E torch.compile modes.
	123	+
	124	+ 与 third_party 同一核心拼接实现，属性经 __init__ 固化为图常量。
	125	+ """
	126	+
	127	+ def __init__(self, , axis=1, tiles=1, *kwargs):
	128	+ super().__init__()
	129	+ del kwargs
	130	+ self.axis = axis
	131	+ self.tiles = tiles
	132	+
	133	+ def forward(self, x, **kwargs):
	134	+ del kwargs
	135	+ return _tile_with_axis_compute(x, self.axis, self.tiles)
	136	+
	137	+
	138	+class TileWithAxisTestSpec:
	139	+ """TileWithAxis operator golden spec (kernel flow: numpy in, numpy out).
	140	+
	141	+ x: input Tensor (numpy.ndarray)
	142	+ axis: int, the axis along which to tile
	143	+ tiles: int, number of copies to concatenate along axis
	144	+ """
	145	+
	146	+ golden = staticmethod(tile_with_axis_golden)
	147	+ third_party = {"torch": _TorchTileWithAxis}
	148	+ torch_graph = TileWithAxisGraph
	149	+ tolerance = _TOLERANCE

Mconversion/tile_with_axis/tests/ut/op_host/arch35/test_tile_with_axis_tiling.cpp+133-64

	@@ -36,43 +36,31 @@ static const std::string OP_NAME = "TileWithAxis";
36	// ============================================================================	36	// ============================================================================
37	// Helper: build TilingContextPara with attributes	37	// Helper: build TilingContextPara with attributes
38	// ============================================================================	38	// ============================================================================
39	-static TilingContextPara MakeTilingPara(	39	+static TilingContextPara MakeTilingPara(const string& caseName, initializer_list<int64_t> xShape, DataType xDtype,
40	- const string& caseName,	40	+ int64_t axis, int64_t tiles,
41	- initializer_list<int64_t> xShape,	41	+ uint64_t ubSize = 253952, // Ascend950 UB = 248KB
42	- DataType xDtype,	42	+ uint64_t coreNum = 64, uint64_t tilingDataSize = 4096, bool includeTiles = true)
43	- int64_t axis,
44	- int64_t tiles,
45	- uint64_t ubSize = 253952, // Ascend950 UB = 248KB
46	- uint64_t coreNum = 64,
47	- uint64_t tilingDataSize = 4096)
48	{	43	{
49	StorageShape shape = {xShape, xShape};	44	StorageShape shape = {xShape, xShape};
50	- vector<TilingContextPara::TensorDescription> inputs = {	45	+ vector<TilingContextPara::TensorDescription> inputs = {{shape, xDtype, FORMAT_ND}};
51	- {shape, xDtype, FORMAT_ND}
52	- };
53	// Output shape: not used by tiling function directly; infershape computes it	46	// Output shape: not used by tiling function directly; infershape computes it
54	- StorageShape outShape = {xShape, xShape}; // placeholder	47	+ StorageShape outShape = {xShape, xShape}; // placeholder
55	- vector<TilingContextPara::TensorDescription> outputs = {	48	+ vector<TilingContextPara::TensorDescription> outputs = {{outShape, xDtype, FORMAT_ND}};
56	- {outShape, xDtype, FORMAT_ND}
57	- };
58		49
59	- vector<TilingContextPara::OpAttr> attrs = {	50	+ vector<TilingContextPara::OpAttr> attrs = {{"axis", Ops::Math::AnyValue::CreateFrom<int64_t>(axis)}};
60	- {"axis", Ops::Math::AnyValue::CreateFrom<int64_t>(axis)},	51	+ if (includeTiles) {
61	- {"tiles", Ops::Math::AnyValue::CreateFrom<int64_t>(tiles)}	52	+ attrs.push_back({"tiles", Ops::Math::AnyValue::CreateFrom<int64_t>(tiles)});
62	- };	53	+ }
63		54
64	- struct TileWithAxisCompileInfo {} compileInfo;	55	+ struct TileWithAxisCompileInfo {
65	- return TilingContextPara(OP_NAME, inputs, outputs, attrs, &compileInfo,	56	+ } compileInfo;
66	- coreNum, ubSize, tilingDataSize);	57	+ return TilingContextPara(OP_NAME, inputs, outputs, attrs, &compileInfo, coreNum, ubSize, tilingDataSize);
67	}	58	}
68		59
69	// ============================================================================	60	// ============================================================================
70	// Helper: execute tiling and extract TilingData	61	// Helper: execute tiling and extract TilingData
71	// ============================================================================	62	// ============================================================================
72	-static bool GetTilingData(const TilingContextPara& para, TilingInfo& info)	63	+static bool GetTilingData(const TilingContextPara& para, TilingInfo& info) { return ExecuteTiling(para, info); }
73	-{
74	- return ExecuteTiling(para, info);
75	-}
76		64
77	// ============================================================================	65	// ============================================================================
78	// P0: ubAxis=0 分支 — 切 outerDim	66	// P0: ubAxis=0 分支 — 切 outerDim
	@@ -83,21 +71,20 @@ static bool GetTilingData(const TilingContextPara& para, TilingInfo& info)
83	// ============================================================================	71	// ============================================================================
84	TEST(TileWithAxisTiling, ubAxis0_cut_outerDim_fp16)	72	TEST(TileWithAxisTiling, ubAxis0_cut_outerDim_fp16)
85	{	73	{
86	- auto para = MakeTilingPara("ubAxis0_fp16",	74	+ auto para = MakeTilingPara("ubAxis0_fp16", {100, 10, 10}, DT_FLOAT16, 1, 3);
87	- {100, 10, 10}, DT_FLOAT16, 1, 3);
88	TilingInfo info;	75	TilingInfo info;
89	ASSERT_TRUE(GetTilingData(para, info));	76	ASSERT_TRUE(GetTilingData(para, info));
90		77
91	auto* td = reinterpret_cast<const TileWithAxisTilingData*>(info.tilingData.get());	78	auto* td = reinterpret_cast<const TileWithAxisTilingData*>(info.tilingData.get());
92	EXPECT_EQ(td->ubAxis, static_cast<uint8_t>(0));	79	EXPECT_EQ(td->ubAxis, static_cast<uint8_t>(0));
93	EXPECT_EQ(td->tiles, 3);	80	EXPECT_EQ(td->tiles, 3);
94	- EXPECT_EQ(td->rowLength, 100); // axisDim=10 * innerDim=10	81	+ EXPECT_EQ(td->rowLength, 100); // axisDim=10 * innerDim=10
95	- EXPECT_EQ(td->inShape[0], 100); // outerDim	82	+ EXPECT_EQ(td->inShape[0], 100); // outerDim
96	- EXPECT_EQ(td->inShape[1], 1); // 输入中间维=1	83	+ EXPECT_EQ(td->inShape[1], 1); // 输入中间维=1
97	- EXPECT_EQ(td->inShape[2], 100); // rowLength	84	+ EXPECT_EQ(td->inShape[2], 100); // rowLength
98	- EXPECT_EQ(td->outShape[0], 100); // outerDim	85	+ EXPECT_EQ(td->outShape[0], 100); // outerDim
99	- EXPECT_EQ(td->outShape[1], 3); // tiles	86	+ EXPECT_EQ(td->outShape[1], 3); // tiles
100	- EXPECT_EQ(td->outShape[2], 100); // rowLength	87	+ EXPECT_EQ(td->outShape[2], 100); // rowLength
101	// Step 2.5 核数感知: totalCount=1 < threshold=51, 降 ubFactor 增加 block	88	// Step 2.5 核数感知: totalCount=1 < threshold=51, 降 ubFactor 增加 block
102	// targetBlocks=min(51, 30000/4096=7, 100)=7, ubFactor=CeilDiv(100,7)=15	89	// targetBlocks=min(51, 30000/4096=7, 100)=7, ubFactor=CeilDiv(100,7)=15
103	EXPECT_EQ(td->totalCount, static_cast<uint64_t>(7));	90	EXPECT_EQ(td->totalCount, static_cast<uint64_t>(7));
	@@ -115,14 +102,13 @@ TEST(TileWithAxisTiling, ubAxis0_cut_outerDim_fp16)
115	// ============================================================================	102	// ============================================================================
116	TEST(TileWithAxisTiling, ubAxis0_cut_outerDim_fp32)	103	TEST(TileWithAxisTiling, ubAxis0_cut_outerDim_fp32)
117	{	104	{
118	- auto para = MakeTilingPara("ubAxis0_fp32",	105	+ auto para = MakeTilingPara("ubAxis0_fp32", {100, 6, 50}, DT_FLOAT, 1, 3);
119	- {100, 6, 50}, DT_FLOAT, 1, 3);
120	TilingInfo info;	106	TilingInfo info;
121	ASSERT_TRUE(GetTilingData(para, info));	107	ASSERT_TRUE(GetTilingData(para, info));
122		108
123	auto* td = reinterpret_cast<const TileWithAxisTilingData*>(info.tilingData.get());	109	auto* td = reinterpret_cast<const TileWithAxisTilingData*>(info.tilingData.get());
124	EXPECT_EQ(td->ubAxis, static_cast<uint8_t>(0));	110	EXPECT_EQ(td->ubAxis, static_cast<uint8_t>(0));
125	- EXPECT_EQ(td->rowLength, 300); // axisDim=6 * innerDim=50	111	+ EXPECT_EQ(td->rowLength, 300); // axisDim=6 * innerDim=50
126	EXPECT_EQ(td->outShape[0], 100);	112	EXPECT_EQ(td->outShape[0], 100);
127	EXPECT_EQ(td->outShape[1], 3);	113	EXPECT_EQ(td->outShape[1], 3);
128	EXPECT_EQ(td->outShape[2], 300);	114	EXPECT_EQ(td->outShape[2], 300);
	@@ -141,8 +127,7 @@ TEST(TileWithAxisTiling, ubAxis0_cut_outerDim_fp32)
141	// ============================================================================	127	// ============================================================================
142	TEST(TileWithAxisTiling, ubAxis1_cut_tiles_fp16)	128	TEST(TileWithAxisTiling, ubAxis1_cut_tiles_fp16)
143	{	129	{
144	- auto para = MakeTilingPara("ubAxis1_fp16",	130	+ auto para = MakeTilingPara("ubAxis1_fp16", {3, 1000, 60}, DT_FLOAT16, 1, 3);
145	- {3, 1000, 60}, DT_FLOAT16, 1, 3);
146	TilingInfo info;	131	TilingInfo info;
147	ASSERT_TRUE(GetTilingData(para, info));	132	ASSERT_TRUE(GetTilingData(para, info));
148		133
	@@ -151,7 +136,7 @@ TEST(TileWithAxisTiling, ubAxis1_cut_tiles_fp16)
151	// candTarget=51, bpot=CeilDiv(51,9)=6, ubFactor=CeilDiv(60000,6)=10000	136	// candTarget=51, bpot=CeilDiv(51,9)=6, ubFactor=CeilDiv(60000,6)=10000
152	EXPECT_EQ(td->ubAxis, static_cast<uint8_t>(2));	137	EXPECT_EQ(td->ubAxis, static_cast<uint8_t>(2));
153	EXPECT_EQ(td->ubFactor, static_cast<uint32_t>(10000));	138	EXPECT_EQ(td->ubFactor, static_cast<uint32_t>(10000));
154	- EXPECT_EQ(td->rowLength, 60000); // axisDim=1000 * innerDim=60	139	+ EXPECT_EQ(td->rowLength, 60000); // axisDim=1000 * innerDim=60
155	EXPECT_EQ(td->outShape[0], 3);	140	EXPECT_EQ(td->outShape[0], 3);
156	EXPECT_EQ(td->outShape[1], 3);	141	EXPECT_EQ(td->outShape[1], 3);
157	EXPECT_EQ(td->outShape[2], 60000);	142	EXPECT_EQ(td->outShape[2], 60000);
	@@ -167,8 +152,7 @@ TEST(TileWithAxisTiling, ubAxis1_cut_tiles_fp16)
167	// ============================================================================	152	// ============================================================================
168	TEST(TileWithAxisTiling, ubAxis1_cut_tiles_fp32)	153	TEST(TileWithAxisTiling, ubAxis1_cut_tiles_fp32)
169	{	154	{
170	- auto para = MakeTilingPara("ubAxis1_fp32",	155	+ auto para = MakeTilingPara("ubAxis1_fp32", {3, 500, 40}, DT_FLOAT, 1, 5);
171	- {3, 500, 40}, DT_FLOAT, 1, 5);
172	TilingInfo info;	156	TilingInfo info;
173	ASSERT_TRUE(GetTilingData(para, info));	157	ASSERT_TRUE(GetTilingData(para, info));
174		158
	@@ -192,8 +176,7 @@ TEST(TileWithAxisTiling, ubAxis1_cut_tiles_fp32)
192	// ============================================================================	176	// ============================================================================
193	TEST(TileWithAxisTiling, ubAxis2_cut_rowLength_fp16)	177	TEST(TileWithAxisTiling, ubAxis2_cut_rowLength_fp16)
194	{	178	{
195	- auto para = MakeTilingPara("ubAxis2_fp16",	179	+ auto para = MakeTilingPara("ubAxis2_fp16", {2, 65536, 16}, DT_FLOAT16, 1, 3);
196	- {2, 65536, 16}, DT_FLOAT16, 1, 3);
197	TilingInfo info;	180	TilingInfo info;
198	ASSERT_TRUE(GetTilingData(para, info));	181	ASSERT_TRUE(GetTilingData(para, info));
199		182
	@@ -217,8 +200,7 @@ TEST(TileWithAxisTiling, ubAxis2_cut_rowLength_fp16)
217	// ============================================================================	200	// ============================================================================
218	TEST(TileWithAxisTiling, ubAxis2_cut_rowLength_fp32)	201	TEST(TileWithAxisTiling, ubAxis2_cut_rowLength_fp32)
219	{	202	{
220	- auto para = MakeTilingPara("ubAxis2_fp32",	203	+ auto para = MakeTilingPara("ubAxis2_fp32", {2, 32768, 16}, DT_FLOAT, 1, 3);
221	- {2, 32768, 16}, DT_FLOAT, 1, 3);
222	TilingInfo info;	204	TilingInfo info;
223	ASSERT_TRUE(GetTilingData(para, info));	205	ASSERT_TRUE(GetTilingData(para, info));
224		206
	@@ -237,8 +219,7 @@ TEST(TileWithAxisTiling, ubAxis2_cut_rowLength_fp32)
237	// ============================================================================	219	// ============================================================================
238	TEST(TileWithAxisTiling, scalar_input)	220	TEST(TileWithAxisTiling, scalar_input)
239	{	221	{
240	- auto para = MakeTilingPara("scalar",	222	+ auto para = MakeTilingPara("scalar", {}, DT_FLOAT16, 0, 5);
241	- {}, DT_FLOAT16, 0, 5);
242	TilingInfo info;	223	TilingInfo info;
243	ASSERT_TRUE(GetTilingData(para, info));	224	ASSERT_TRUE(GetTilingData(para, info));
244		225
	@@ -265,8 +246,7 @@ TEST(TileWithAxisTiling, scalar_input)
265	// ============================================================================	246	// ============================================================================
266	TEST(TileWithAxisTiling, negative_axis)	247	TEST(TileWithAxisTiling, negative_axis)
267	{	248	{
268	- auto para = MakeTilingPara("negative_axis",	249	+ auto para = MakeTilingPara("negative_axis", {2, 3, 4}, DT_FLOAT16, -1, 3);
269	- {2, 3, 4}, DT_FLOAT16, -1, 3);
270	TilingInfo info;	250	TilingInfo info;
271	ASSERT_TRUE(GetTilingData(para, info));	251	ASSERT_TRUE(GetTilingData(para, info));
272		252
	@@ -275,8 +255,8 @@ TEST(TileWithAxisTiling, negative_axis)
275	// candTarget=12, bpot=CeilDiv(12,6)=2, ubFactor=CeilDiv(4,2)=2	255	// candTarget=12, bpot=CeilDiv(12,6)=2, ubFactor=CeilDiv(4,2)=2
276	EXPECT_EQ(td->ubAxis, static_cast<uint8_t>(2));	256	EXPECT_EQ(td->ubAxis, static_cast<uint8_t>(2));
277	EXPECT_EQ(td->tiles, 3);	257	EXPECT_EQ(td->tiles, 3);
278	- EXPECT_EQ(td->rowLength, 4); // axisDim=4, innerDim=1	258	+ EXPECT_EQ(td->rowLength, 4); // axisDim=4, innerDim=1
279	- EXPECT_EQ(td->inShape[0], 6); // outerDim = 2*3 = 6	259	+ EXPECT_EQ(td->inShape[0], 6); // outerDim = 2*3 = 6
280	EXPECT_EQ(td->outShape[0], 6);	260	EXPECT_EQ(td->outShape[0], 6);
281	EXPECT_EQ(td->outShape[1], 3);	261	EXPECT_EQ(td->outShape[1], 3);
282	EXPECT_EQ(td->outShape[2], 4);	262	EXPECT_EQ(td->outShape[2], 4);
	@@ -289,8 +269,7 @@ TEST(TileWithAxisTiling, negative_axis)
289	// ============================================================================	269	// ============================================================================
290	TEST(TileWithAxisTiling, tiles_equals_one)	270	TEST(TileWithAxisTiling, tiles_equals_one)
291	{	271	{
292	- auto para = MakeTilingPara("tiles_1",	272	+ auto para = MakeTilingPara("tiles_1", {2, 3, 4}, DT_FLOAT16, 1, 1);
293	- {2, 3, 4}, DT_FLOAT16, 1, 1);
294	TilingInfo info;	273	TilingInfo info;
295	ASSERT_TRUE(GetTilingData(para, info));	274	ASSERT_TRUE(GetTilingData(para, info));
296		275
	@@ -305,8 +284,7 @@ TEST(TileWithAxisTiling, tiles_equals_one)
305	// ============================================================================	284	// ============================================================================
306	TEST(TileWithAxisTiling, tiles_invalid_zero)	285	TEST(TileWithAxisTiling, tiles_invalid_zero)
307	{	286	{
308	- auto para = MakeTilingPara("tiles_0",	287	+ auto para = MakeTilingPara("tiles_0", {2, 3, 4}, DT_FLOAT16, 0, 0);
309	- {2, 3, 4}, DT_FLOAT16, 0, 0);
310	TilingInfo info;	288	TilingInfo info;
311	EXPECT_FALSE(GetTilingData(para, info));	289	EXPECT_FALSE(GetTilingData(para, info));
312	}	290	}
	@@ -317,8 +295,7 @@ TEST(TileWithAxisTiling, tiles_invalid_zero)
317	TEST(TileWithAxisTiling, axis_out_of_range)	295	TEST(TileWithAxisTiling, axis_out_of_range)
318	{	296	{
319	// rank=3, axis=3 → 越界	297	// rank=3, axis=3 → 越界
320	- auto para = MakeTilingPara("axis_oob",	298	+ auto para = MakeTilingPara("axis_oob", {2, 3, 4}, DT_FLOAT16, 3, 2);
321	- {2, 3, 4}, DT_FLOAT16, 3, 2);
322	TilingInfo info;	299	TilingInfo info;
323	EXPECT_FALSE(ExecuteTiling(para, info));	300	EXPECT_FALSE(ExecuteTiling(para, info));
324	}	301	}
	@@ -331,8 +308,7 @@ TEST(TileWithAxisTiling, axis_out_of_range)
331	// ============================================================================	308	// ============================================================================
332	TEST(TileWithAxisTiling, high_dim_4d_axis_2)	309	TEST(TileWithAxisTiling, high_dim_4d_axis_2)
333	{	310	{
334	- auto para = MakeTilingPara("4d_axis2",	311	+ auto para = MakeTilingPara("4d_axis2", {2, 4, 8, 16}, DT_FLOAT16, 2, 3);
335	- {2, 4, 8, 16}, DT_FLOAT16, 2, 3);
336	TilingInfo info;	312	TilingInfo info;
337	ASSERT_TRUE(GetTilingData(para, info));	313	ASSERT_TRUE(GetTilingData(para, info));
338		314
	@@ -340,10 +316,103 @@ TEST(TileWithAxisTiling, high_dim_4d_axis_2)
340	// Step 2.5 跨轴回退: ubAxis=0 totalCount=1 < threshold=51, 回退到 ubAxis=2	316	// Step 2.5 跨轴回退: ubAxis=0 totalCount=1 < threshold=51, 回退到 ubAxis=2
341	// candTarget=51, bpot=CeilDiv(51,24)=3, ubFactor=CeilDiv(128,3)=43	317	// candTarget=51, bpot=CeilDiv(51,24)=3, ubFactor=CeilDiv(128,3)=43
342	EXPECT_EQ(td->ubAxis, static_cast<uint8_t>(2));	318	EXPECT_EQ(td->ubAxis, static_cast<uint8_t>(2));
343	- EXPECT_EQ(td->rowLength, 128); // axisDim=8 * innerDim=16	319	+ EXPECT_EQ(td->rowLength, 128); // axisDim=8 * innerDim=16
344	- EXPECT_EQ(td->inShape[0], 8); // outerDim = 2*4	320	+ EXPECT_EQ(td->inShape[0], 8); // outerDim = 2*4
345	EXPECT_EQ(td->outShape[1], 3);	321	EXPECT_EQ(td->outShape[1], 3);
346	EXPECT_EQ(td->outShape[2], 128);	322	EXPECT_EQ(td->outShape[2], 128);
347	}	323	}
348		324
		325	+// ============================================================================
		326	+// 回归（SAFE.01/1 / CODE.09/1）: 标量输入 tiles 超过单缓冲容量时按 ubAxis=1 分块
		327	+// shape=[], fp16, tiles=200000, ubSize=253952
		328	+// bufferSize=126976B, bufferSizeElements=63488 < 200000
		329	+// → ubAxis=1, ubFactor=63488, totalCount=CeilDiv(200000,63488)=4
		330	+// perCoreCount=CeilDiv(4,64)=1, realCoreNum=CeilDiv(4,1)=4
		331	+// ============================================================================
		332	+TEST(TileWithAxisTiling, scalar_input_huge_tiles_chunked)
		333	+{
		334	+ auto para = MakeTilingPara("scalar_huge_tiles", {}, DT_FLOAT16, 0, 200000);
		335	+ TilingInfo info;
		336	+ ASSERT_TRUE(GetTilingData(para, info));
		337	+
		338	+ auto* td = reinterpret_cast<const TileWithAxisTilingData*>(info.tilingData.get());
		339	+ EXPECT_EQ(td->ubAxis, static_cast<uint8_t>(1));
		340	+ EXPECT_EQ(td->ubFactor, static_cast<uint32_t>(63488));
		341	+ EXPECT_EQ(td->totalCount, static_cast<uint64_t>(4));
		342	+ EXPECT_EQ(td->outShape[0], 1);
		343	+ EXPECT_EQ(td->outShape[1], 200000);
		344	+ EXPECT_EQ(td->outShape[2], 1);
		345	+ EXPECT_EQ(info.blockNum, static_cast<size_t>(4));
		346	+}
		347	+
		348	+// ============================================================================
		349	+// 回归（空张量兜底，REQUIREMENTS §3.5.2）: 某维为 0 时 totalCount=0
		350	+// SetBlockDim(1)（而非 0），kernel 空循环
		351	+// shape=[2, 0, 3], axis=1, tiles=2, fp16: rowLength=0 → totalCount=0
		352	+// ============================================================================
		353	+TEST(TileWithAxisTiling, empty_tensor_blockdim_one)
		354	+{
		355	+ auto para = MakeTilingPara("empty_tensor", {2, 0, 3}, DT_FLOAT16, 1, 2);
		356	+ TilingInfo info;
		357	+ ASSERT_TRUE(GetTilingData(para, info));
		358	+
		359	+ auto* td = reinterpret_cast<const TileWithAxisTilingData*>(info.tilingData.get());
		360	+ EXPECT_EQ(td->totalCount, static_cast<uint64_t>(0));
		361	+ EXPECT_EQ(td->perCoreCount, static_cast<uint64_t>(1));
		362	+ EXPECT_EQ(info.blockNum, static_cast<size_t>(1));
		363	+}
		364	+
		365	+// ============================================================================
		366	+// 回归（CODE.18/1 / CODE.06/1）: int64 维度乘积溢出时 tiling 干净拒绝
		367	+// rowLength = axisDim * innerDim = 4e18 * 2 = 8e18（int64 可表示），
		368	+// 但 rowOutElements = tiles * rowLength = 2 * 8e18 = 1.6e19 溢出 int64，
		369	+// 应返回 GRAPH_FAILED，而不是生成错误 TilingData。
		370	+// ============================================================================
		371	+TEST(TileWithAxisTiling, int64_product_overflow_rejected)
		372	+{
		373	+ auto para = MakeTilingPara("int64_overflow", {1, 4000000000000000000LL, 2}, DT_FLOAT16, 1, 2);
		374	+ TilingInfo info;
		375	+ EXPECT_FALSE(ExecuteTiling(para, info));
		376	+}
		377	+
		378	+// ============================================================================
		379	+// 回归（CODE.18/1 / CODE.06/1）: rowLength = axisDim * innerDim 溢出时干净拒绝
		380	+// rank=2, axis=0: axisDim=4e18, innerDim=3 → rowLength=1.2e19 溢出 int64
		381	+// ============================================================================
		382	+TEST(TileWithAxisTiling, rowLength_product_overflow_rejected)
		383	+{
		384	+ auto para = MakeTilingPara("rowLength_overflow", {4000000000000000000LL, 3}, DT_FLOAT16, 0, 2);
		385	+ TilingInfo info;
		386	+ EXPECT_FALSE(ExecuteTiling(para, info));
		387	+}
		388	+
		389	+// ============================================================================
		390	+// 回归（CODE.18/1 / CODE.06/1）: outerDim 连乘溢出时干净拒绝
		391	+// rank=3, axis=2: outerDim = 2e9 * 3e9 = 6e18 溢出 int64
		392	+// ============================================================================
		393	+TEST(TileWithAxisTiling, outerDim_product_overflow_rejected)
		394	+{
		395	+ auto para = MakeTilingPara("outerDim_overflow", {2000000000LL, 3000000000LL, 4}, DT_FLOAT16, 2, 2);
		396	+ TilingInfo info;
		397	+ EXPECT_FALSE(ExecuteTiling(para, info));
		398	+}
		399	+
		400	+// ============================================================================
		401	+// 回归（CODE.18/1 / CODE.06/1）: totalOutElements = outerDimtilesrowLength 溢出时干净拒绝
		402	+// outerDim=4e18, rowLength=3 → total = 4e1823 = 2.4e19 溢出 int64
		403	+// ============================================================================
		404	+TEST(TileWithAxisTiling, totalOutElements_overflow_rejected)
		405	+{
		406	+ auto para = MakeTilingPara("total_overflow", {4000000000000000000LL, 3}, DT_FLOAT16, 1, 2);
		407	+ TilingInfo info;
		408	+ EXPECT_FALSE(ExecuteTiling(para, info));
		409	+}
		410	+
		411	+TEST(TileWithAxisTiling, tiles_missing_rejected)
		412	+{
		413	+ auto para = MakeTilingPara("tiles_missing", {2, 3, 4}, DT_FLOAT16, 1, 0, 253952, 64, 4096, false);
		414	+ TilingInfo info;
		415	+ EXPECT_FALSE(ExecuteTiling(para, info));
		416	+}
		417	+
349	} // namespace TileWithAxisUT	418	} // namespace TileWithAxisUT

Mconversion/tile_with_axis/tests/ut/op_host/test_tile_with_axis_infershape.cpp+42-3

	@@ -34,7 +34,7 @@ static const string OP_NAME = "TileWithAxis";
34	// Helper: create InfershapeContextPara with axis/tiles attributes	34	// Helper: create InfershapeContextPara with axis/tiles attributes
35	// ============================================================================	35	// ============================================================================
36	static InfershapeContextPara MakeInfershapePara(const string& caseName, initializer_list<int64_t> xShape,	36	static InfershapeContextPara MakeInfershapePara(const string& caseName, initializer_list<int64_t> xShape,
37	- DataType dtype, int64_t axis, int64_t tiles)	37	+ DataType dtype, int64_t axis, int64_t tiles, bool includeTiles = true)
38	{	38	{
39	StorageShape inputShape = {xShape, xShape};	39	StorageShape inputShape = {xShape, xShape};
40	// Output placeholder: unknown shape	40	// Output placeholder: unknown shape
	@@ -42,8 +42,10 @@ static InfershapeContextPara MakeInfershapePara(const string& caseName, initiali
42		42
43	vector<InfershapeContextPara::TensorDescription> inputs = {{inputShape, dtype, FORMAT_ND}};	43	vector<InfershapeContextPara::TensorDescription> inputs = {{inputShape, dtype, FORMAT_ND}};
44	vector<InfershapeContextPara::TensorDescription> outputs = {{outputPlaceholder, dtype, FORMAT_ND}};	44	vector<InfershapeContextPara::TensorDescription> outputs = {{outputPlaceholder, dtype, FORMAT_ND}};
45	- vector<InfershapeContextPara::OpAttr> attrs = {{"axis", Ops::Math::AnyValue::CreateFrom<int64_t>(axis)},	45	+ vector<InfershapeContextPara::OpAttr> attrs = {{"axis", Ops::Math::AnyValue::CreateFrom<int64_t>(axis)}};
46	- {"tiles", Ops::Math::AnyValue::CreateFrom<int64_t>(tiles)}};	46	+ if (includeTiles) {
		47	+ attrs.push_back({"tiles", Ops::Math::AnyValue::CreateFrom<int64_t>(tiles)});
		48	+ }
47		49
48	return InfershapeContextPara(OP_NAME, inputs, outputs, attrs);	50	return InfershapeContextPara(OP_NAME, inputs, outputs, attrs);
49	}	51	}
	@@ -159,4 +161,41 @@ TEST(TileWithAxisInfershape, invalid_axis_negative_is_rejected)
159	ExecuteTestCase(para, GRAPH_FAILED);	161	ExecuteTestCase(para, GRAPH_FAILED);
160	}	162	}
161		163
		164	+// ============================================================================
		165	+// 回归: axis 维为未知维标记 -1 时原样传播（COMPAT.02/1 / GEIR.03/2）
		166	+// 输入 shape=[-1, 3, 4], axis=1, tiles=2 → 输出 shape=[-1, 6, 4]（非 -1*2=-2）
		167	+// ============================================================================
		168	+TEST(TileWithAxisInfershape, unknown_dim_axis_propagated)
		169	+{
		170	+ auto para = MakeInfershapePara("unknown_dim_axis", {-1, 3, 4}, DT_FLOAT16, 1, 2);
		171	+ vector<vector<int64_t>> expected = {{-1, 6, 4}};
		172	+ ExecuteTestCase(para, GRAPH_SUCCESS, expected);
		173	+}
		174	+
		175	+TEST(TileWithAxisInfershape, unknown_dim_negative_axis_propagated)
		176	+{
		177	+ auto para = MakeInfershapePara("unknown_dim_neg_axis", {2, -1, 4}, DT_FLOAT16, -1, 3);
		178	+ vector<vector<int64_t>> expected = {{2, -1, 12}};
		179	+ ExecuteTestCase(para, GRAPH_SUCCESS, expected);
		180	+}
		181	+
		182	+// ============================================================================
		183	+// 回归: axis 维乘法 int64 溢出时拒绝（GEIR.04/1）
		184	+// dim=5e9, tiles=2e9 → 乘积 1e19 > INT64_MAX，返回 GRAPH_FAILED
		185	+// ============================================================================
		186	+TEST(TileWithAxisInfershape, axis_dim_overflow_rejected)
		187	+{
		188	+ auto para = MakeInfershapePara("axis_dim_overflow", {5000000000LL, 3}, DT_FLOAT16, 0, 2000000000LL);
		189	+ ExecuteTestCase(para, GRAPH_FAILED);
		190	+}
		191	+
		192	+// ============================================================================
		193	+// 回归: tiles 属性缺失时拒绝而非静默回退 1（GEIR.07/1 / CODE.01/1）
		194	+// ============================================================================
		195	+TEST(TileWithAxisInfershape, tiles_missing_rejected)
		196	+{
		197	+ auto para = MakeInfershapePara("tiles_missing", {2, 3, 4}, DT_FLOAT16, 1, 0, false);
		198	+ ExecuteTestCase(para, GRAPH_FAILED);
		199	+}
		200	+
162	} // namespace TileWithAxisUT	201	} // namespace TileWithAxisUT