diag_flat: arch35 split, config update, remove op

Mconversion/diag_flat/CMakeLists.txt+5-9

	@@ -9,12 +9,8 @@
9	# See LICENSE in the root of the software repository for the full text of the License.	9	# See LICENSE in the root of the software repository for the full text of the License.
10	# ----------------------------------------------------------------------------	10	# ----------------------------------------------------------------------------
11		11
12	-file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)	12	+set(SUPPORT_COMPUTE_UNIT "ascend310b" "ascend310p" "ascend910" "ascend910_93" "ascend910b" "kirin9030" "kirinx90" "ascend950")
13	-if(NOT ENABLE_TEST AND NOT BENCHMARK)	13	+set(SUPPORT_TILING_DIR "arch32" "arch32" "arch32" "arch32" "arch32" "arch32" "arch32" "arch35")
14	- list(REMOVE_ITEM CURRENT_DIRS tests)	14	+add_all_modules_sources(OPTYPE diag_flat ACLNNTYPE aclnn_exclude
15	-endif()	15	+ COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT}
16	-foreach(SUB_DIR ${CURRENT_DIRS})	16	+ TILING_DIR ${SUPPORT_TILING_DIR})
17	- if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
18	- add_subdirectory(${SUB_DIR})
19	- endif()
20	-endforeach()

Mconversion/diag_flat/examples/test_aclnn_diag_flat.cpp+2-2

	@@ -124,8 +124,8 @@ int main() {
124	}	124	}
125		125
126	// 6. 释放aclTensor，需要根据具体API的接口定义修改	126	// 6. 释放aclTensor，需要根据具体API的接口定义修改
127	- aclDestroyTensor(self);	127	+ (void)aclDestroyTensor(self);
128	- aclDestroyTensor(out);	128	+ (void)aclDestroyTensor(out);
129		129
130	// 7. 释放device资源	130	// 7. 释放device资源
131	aclrtFree(selfDeviceAddr);	131	aclrtFree(selfDeviceAddr);

Rconversion/diag_flat/op_host/op_api/aclnn_diag_flat.cpp→conversion/diag_flat/op_api/aclnn_diag_flat.cpp+0-0

文件重命名但无更改。

Rconversion/diag_flat/op_host/op_api/aclnn_diag_flat.h→conversion/diag_flat/op_api/aclnn_diag_flat.h+0-0

文件重命名但无更改。

Rconversion/diag_flat/op_host/op_api/diag_flat.cpp→conversion/diag_flat/op_api/diag_flat.cpp+0-0

文件重命名但无更改。

Rconversion/diag_flat/op_host/op_api/diag_flat.h→conversion/diag_flat/op_api/diag_flat.h+0-0

文件重命名但无更改。

Dconversion/diag_flat/op_host/CMakeLists.txt+0-12

	@@ -1,12 +0,0 @@
1	-# ----------------------------------------------------------------------------
2	-# This program is free software, you can redistribute it and/or modify it.
3	-# Copyright (c) 2025 Huawei Technologies Co., Ltd.
4	-# This file is a part of the CANN Open Software.
5	-# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6	-# Please refer to the License for details. You may not use this file except in compliance with the License.
7	-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8	-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9	-# See LICENSE in the root of the software repository for the full text of the License.
10	-# ----------------------------------------------------------------------------
11	-
12	-add_modules_sources(OPTYPE diag_flat ACLNNTYPE aclnn_exclude)

Rconversion/diag_flat/op_host/diag_flat_tiling.cpp→conversion/diag_flat/op_host/arch32/diag_flat_tiling.cpp+1-1

	@@ -13,7 +13,7 @@
13	* \brief	13	* \brief
14	*/	14	*/
15		15
16	-#include "../../diag_v2/op_host/diag_v2_tiling.h"	16	+#include "../../../diag_v2/op_host/arch32/diag_v2_tiling.h"
17	#include "util/math_util.h"	17	#include "util/math_util.h"
18	#include "log/log.h"	18	#include "log/log.h"
19	#include "platform/platform_ascendc.h"	19	#include "platform/platform_ascendc.h"

Aconversion/diag_flat/op_host/arch35/diag_flat_tiling.cpp+204-0

@@ -0,0 +1,204 @@
	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 diag_flat_tiling.cpp
	13	+ * \brief DiagFlat Host-side Tiling implementation (arch35, DAV_3510)
	14	+ *
	15	+ * Design: DESIGN.md v2.1 Sec 3.3
	16	+ *
	17	+ * TilingDiagFlatArch35() is non-static, exported via diag_flat_tiling.h.
	18	+ * Called by both DiagFlat (own tiling) and DiagV2 (rank==1 delegation).
	19	+ * One-way dependency: this file includes nothing from diag_v2.
	20	+ */
	21	+
	22	+#include "diag_flat_tiling.h"
	23	+#include "register/op_def_registry.h"
	24	+#include "op_common/log/log.h"
	25	+#include "op_common/op_host/util/math_util.h"
	26	+#include "op_common/op_host/util/platform_util.h"
	27	+#include "../../op_kernel/arch35/diag_flat_tiling_data.h"
	28	+#include "../../op_kernel/arch35/diag_flat_tiling_key.h"
	29	+#include <algorithm>
	30	+#include <cstring>
	31	+
	32	+namespace optiling {
	33	+
	34	+using Ops::Base::CeilDiv;
	35	+
	36	+constexpr int64_t MIN_WORK_PER_CORE = 256;
	37	+constexpr int64_t DIAG_FLAT_BUFFER_NUM = 2;
	38	+constexpr int64_t SIMT_DCACHE_SIZE = 64 * 1024;
	39	+
	40	+// ============================================================================
	41	+// Helper: Get platform info (ubSize, coreNum)
	42	+// ============================================================================
	43	+static ge::graphStatus GetPlatformInfo(gert::TilingContext* context,
	44	+ uint64_t* ubSize, int64_t* coreNum)
	45	+{
	46	+ fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo();
	47	+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr);
	48	+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
	49	+ *coreNum = ascendcPlatform.GetCoreNumAiv();
	50	+ OP_CHECK_IF(*coreNum == 0, OP_LOGE(context, "coreNum is 0"),
	51	+ return ge::GRAPH_FAILED);
	52	+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, *ubSize);
	53	+ OP_CHECK_IF(*ubSize == 0, OP_LOGE(context, "ubSize is 0"),
	54	+ return ge::GRAPH_FAILED);
	55	+ return ge::GRAPH_SUCCESS;
	56	+}
	57	+
	58	+// ============================================================================
	59	+// Helper: Get dtype size
	60	+// ============================================================================
	61	+static int64_t GetDtypeSize(ge::DataType dataType)
	zhanw_coding6月16日可以直接使用 ge::GetSizeByDataType
	62	+{
	63	+ switch (dataType) {
	64	+ case ge::DT_FLOAT16: case ge::DT_BF16:
	65	+ case ge::DT_INT16: case ge::DT_UINT16:
	66	+ return 2;
	67	+ case ge::DT_FLOAT: case ge::DT_INT32: case ge::DT_UINT32:
	68	+ return 4;
	69	+ case ge::DT_DOUBLE: case ge::DT_INT64: case ge::DT_UINT64:
	70	+ case ge::DT_COMPLEX64:
	71	+ return 8;
	72	+ case ge::DT_INT8: case ge::DT_UINT8:
	73	+ return 1;
	74	+ default:
	75	+ return 4;
	76	+ }
	77	+}
	78	+
	79	+// ============================================================================
	80	+// Exported: core 1D→2D tiling computation (one-way dependency: no diag_v2 includes)
	81	+// ============================================================================
	82	+ge::graphStatus TilingDiagFlatArch35(gert::TilingContext* context, DiagFlatTilingOutput* out)
	83	+{
	84	+ // 1. Get platform info
	85	+ uint64_t ubSize;
	86	+ int64_t hwCoreNum;
	87	+ OP_CHECK_IF(
	88	+ GetPlatformInfo(context, &ubSize, &hwCoreNum) != ge::GRAPH_SUCCESS,
	89	+ OP_LOGE(context, "GetPlatformInfo error"),
	90	+ return ge::GRAPH_FAILED);
	91	+
	92	+ // 2. Get input shape and numInput
	93	+ auto inputX = context->GetInputShape(0);
	94	+ OP_CHECK_NULL_WITH_CONTEXT(context, inputX);
	95	+ auto inputShape = inputX->GetStorageShape();
	96	+ int64_t numInput = inputShape.GetShapeSize();
	97	+
	98	+ // 3. Get diagonal attribute
	99	+ auto attrs = context->GetAttrs();
	100	+ OP_CHECK_NULL_WITH_CONTEXT(context, attrs);
	101	+ const int64_t* diagonalPtr = attrs->GetAttrPointer<int64_t>(0);
	102	+ int64_t diagonal = (diagonalPtr != nullptr) ? *diagonalPtr : 0;
	103	+
	104	+ // 4. Get dtype size
	105	+ auto inputDesc = context->GetInputDesc(0);
	106	+ OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc);
	107	+ auto dataType = inputDesc->GetDataType();
	108	+ int64_t dtypeSize = GetDtypeSize(dataType);
	109	+
	110	+ // 5. Compute output dimensions
	111	+ int64_t outWidth = numInput + (diagonal >= 0 ? diagonal : -diagonal);
	112	+ int64_t outTotal = outWidth * outWidth;
	113	+
	114	+ // 6. Multi-core split
	115	+ int64_t realCoreNum = std::min(hwCoreNum,
	116	+ std::max<int64_t>(1, (numInput + outTotal) / MIN_WORK_PER_CORE / 2));
	117	+
	118	+ // 7. outPerCore
	119	+ int64_t outPerCore = 0;
	120	+ if (outTotal > 0) {
	121	+ outPerCore = CeilDiv(outTotal, realCoreNum);
	122	+ }
	123	+
	124	+ // 8. tileLength
	125	+ OP_CHECK_IF(dtypeSize <= 0, OP_LOGE(context, "invalid dtypeSize %ld", dtypeSize),
	126	+ return ge::GRAPH_FAILED);
	127	+ int64_t ubPerBuf = (ubSize - SIMT_DCACHE_SIZE) / DIAG_FLAT_BUFFER_NUM;
	128	+ int64_t tileLength = std::min<int64_t>({ubPerBuf / dtypeSize, outPerCore});
	129	+
	130	+ // 9. Fill output struct
	131	+ out->numInput = numInput;
	132	+ out->diagonal = diagonal;
	133	+ out->outWidth = outWidth;
	134	+ out->outTotal = outTotal;
	135	+ out->outPerCore = outPerCore;
	136	+ out->tileLength = tileLength;
	137	+ out->realCoreNum = realCoreNum;
	138	+ out->localMemSize = static_cast<uint32_t>(ubSize - SIMT_DCACHE_SIZE);
	139	+
	140	+ // 10. Set context-level params (common to both callers)
	141	+ auto ret = context->SetLocalMemorySize(out->localMemSize);
	142	+ OP_CHECK_IF(ret != ge::GRAPH_SUCCESS,
	143	+ OP_LOGE(context, "set local memory size failed."), return ret);
	144	+
	145	+ context->SetBlockDim(realCoreNum);
	146	+
	147	+ size_t* currentWorkspace = context->GetWorkspaceSizes(1);
	148	+ OP_CHECK_NULL_WITH_CONTEXT(context, currentWorkspace);
	149	+ currentWorkspace[0] = 0;
	150	+
	151	+ return ge::GRAPH_SUCCESS;
	152	+}
	153	+
	154	+// ============================================================================
	155	+// DiagFlat's own tiling entry (static wrapper)
	156	+// ============================================================================
	157	+static ge::graphStatus DiagFlatTilingFunc(gert::TilingContext* context)
	158	+{
	159	+ DiagFlatTilingOutput out;
	160	+ OP_CHECK_IF(
	161	+ TilingDiagFlatArch35(context, &out) != ge::GRAPH_SUCCESS,
	162	+ OP_LOGE(context, "TilingDiagFlatArch35 error"),
	163	+ return ge::GRAPH_FAILED);
	164	+
	165	+ // Fill DiagFlat's own TilingData
	166	+ DiagFlatArch35TilingData* tiling = context->GetTilingData<DiagFlatArch35TilingData>();
	167	+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
	168	+ OP_CHECK_IF(
	169	+ memset_s(tiling, sizeof(DiagFlatArch35TilingData), 0, sizeof(DiagFlatArch35TilingData)) != EOK,
	170	+ OP_LOGE(context, "set tiling data error"),
	171	+ return ge::GRAPH_FAILED);
	172	+
	173	+ tiling->numInput = out.numInput;
	174	+ tiling->diagonal = out.diagonal;
	175	+ tiling->outWidth = out.outWidth;
	176	+ tiling->outTotal = out.outTotal;
	177	+ tiling->outPerCore = out.outPerCore;
	178	+ tiling->tileLength = out.tileLength;
	179	+ tiling->realCoreNum = out.realCoreNum;
	180	+
	181	+ // DiagFlat's own TilingKey
	182	+ ASCENDC_TPL_SEL_PARAM(context, static_cast<uint32_t>(3501));
	183	+
	184	+ return ge::GRAPH_SUCCESS;
	185	+}
	186	+
	187	+// ============================================================================
	188	+// TilingParse
	189	+// ============================================================================
	190	+static ge::graphStatus TilingParseForDiagFlat([[maybe_unused]] gert::TilingParseContext* context)
	191	+{
	192	+ return ge::GRAPH_SUCCESS;
	193	+}
	194	+
	195	+struct DiagFlatCompileInfo {};
	196	+
	197	+// ============================================================================
	198	+// Tiling registration
	199	+// ============================================================================
	200	+IMPL_OP_OPTILING(DiagFlat)
	201	+ .Tiling(DiagFlatTilingFunc)
	202	+ .TilingParse<DiagFlatCompileInfo>(TilingParseForDiagFlat);
	203	+
	204	+} // namespace optiling

Aconversion/diag_flat/op_host/arch35/diag_flat_tiling.h+48-0

@@ -0,0 +1,48 @@
	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 diag_flat_tiling.h
	13	+ * \brief DiagFlat arch35 tiling function export (called by diag_v2 for 1D→2D delegation)
	14	+ *
	15	+ * One-way dependency: this header is self-contained (no diag_v2 includes).
	16	+ * diag_v2 includes this header and calls TilingDiagFlatArch35() for rank==1 input.
	17	+ */
	18	+
	19	+#ifndef __DIAG_FLAT_ARCH35_TILING_H__
	20	+#define __DIAG_FLAT_ARCH35_TILING_H__
	21	+
	22	+#include <cstdint>
	23	+#include "register/op_impl_registry.h"
	24	+
	25	+namespace optiling {
	26	+
	27	+// Output struct returned by TilingDiagFlatArch35.
	28	+// Caller fills its own TilingData from these values.
	29	+struct DiagFlatTilingOutput {
	30	+ int64_t numInput;
	31	+ int64_t diagonal;
	32	+ int64_t outWidth;
	33	+ int64_t outTotal;
	34	+ int64_t outPerCore;
	35	+ int64_t tileLength;
	36	+ int64_t realCoreNum;
	37	+ uint32_t localMemSize; // ubSize - SIMT_DCACHE_SIZE, for SetLocalMemorySize
	38	+};
	39	+
	40	+// Core 1D→2D tiling computation. Does NOT call context->GetTilingData<>()
	41	+// — the caller is responsible for filling its own TilingData struct.
	42	+// Also sets BlockDim, LocalMemorySize, and Workspace on the context.
	43	+// TilingKey selection is the caller's responsibility.
	44	+ge::graphStatus TilingDiagFlatArch35(gert::TilingContext* context, DiagFlatTilingOutput* out);
	45	+
	46	+} // namespace optiling
	47	+
	48	+#endif // __DIAG_FLAT_ARCH35_TILING_H__

Mconversion/diag_flat/op_host/config/ascend950/diag_flat_binary.json+148-149

	@@ -1,150 +1,149 @@
1	{	1	{
2	- "op_type": "DiagFlat",	2	+ "op_type": "DiagFlat",
3	- "op_list": [	3	+ "op_list": [
4	- {	4	+ {
5	- "bin_filename": "DiagFlat_1980d2db1a80c523e7e59a032c95969880af8",	5	+ "bin_filename": "DiagFlat_1980d2db1a80c523e7e59a032c95969880af8",
6	- "inputs": [	6	+ "inputs": [
7	- {	7	+ {
8	- "name": "x",	8	+ "name": "x",
9	- "index": 0,	9	+ "index": 0,
10	- "dtype": "int8",	10	+ "dtype": "int8",
11	- "format": "ND",	11	+ "format": "ND",
12	- "paramType": "required",	12	+ "paramType": "required",
13	- "shape": [	13	+ "shape": [
14	- -2	14	+ -2
15	- ],	15	+ ],
16	- "dtype_match_mode": "DtypeByte"	16	+ "dtype_match_mode": "DtypeByte"
17	- }	17	+ }
18	- ],	18	+ ],
19	- "outputs": [	19	+ "outputs": [
20	- {	20	+ {
21	- "name": "y",	21	+ "name": "y",
22	- "index": 0,	22	+ "index": 0,
23	- "dtype": "int8",	23	+ "dtype": "int8",
24	- "format": "ND",	24	+ "format": "ND",
25	- "paramType": "required",	25	+ "paramType": "required",
26	- "shape": [	26	+ "shape": [
27	- -2	27	+ -2
28	- ],	28	+ ],
29	- "dtype_match_mode": "DtypeByte"	29	+ "dtype_match_mode": "DtypeByte"
30	- }	30	+ }
31	- ],	31	+ ],
32	- "attrs": [	32	+ "attrs": [
33	- {	33	+ {
34	- "name": "diagonal",	34	+ "name": "diagonal",
35	- "dtype": "int",	35	+ "dtype": "int",
36	- "value": null	36	+ "value": null
37	- }	37	+ }
38	- ]	38	+ ]
39	- },	39	+ },
40	- {	40	+ {
41	- "bin_filename": "DiagFlat_19808098b79498645e9c0da095babdeaf862i16",	41	+ "bin_filename": "DiagFlat_19808098b79498645e9c0da095babdeaf862i16",
42	- "inputs": [	42	+ "inputs": [
43	- {	43	+ {
44	- "name": "x",	44	+ "name": "x",
45	- "index": 0,	45	+ "index": 0,
46	- "dtype": "int16",	46	+ "dtype": "int16",
47	- "format": "ND",	47	+ "format": "ND",
48	- "paramType": "required",	48	+ "paramType": "required",
49	- "shape": [	49	+ "shape": [
50	- -2	50	+ -2
51	- ],	51	+ ],
52	- "dtype_match_mode": "DtypeByte"	52	+ "dtype_match_mode": "DtypeByte"
53	- }	53	+ }
54	- ],	54	+ ],
55	- "outputs": [	55	+ "outputs": [
56	- {	56	+ {
57	- "name": "y",	57	+ "name": "y",
58	- "index": 0,	58	+ "index": 0,
59	- "dtype": "int16",	59	+ "dtype": "int16",
60	- "format": "ND",	60	+ "format": "ND",
61	- "paramType": "required",	61	+ "paramType": "required",
62	- "shape": [	62	+ "shape": [
63	- -2	63	+ -2
64	- ],	64	+ ],
65	- "dtype_match_mode": "DtypeByte"	65	+ "dtype_match_mode": "DtypeByte"
66	- }	66	+ }
67	- ],	67	+ ],
68	- "attrs": [	68	+ "attrs": [
69	- {	69	+ {
70	- "name": "diagonal",	70	+ "name": "diagonal",
71	- "dtype": "int",	71	+ "dtype": "int",
72	- "value": null	72	+ "value": null
73	- }	73	+ }
74	- ]	74	+ ]
75	- },	75	+ },
76	- {	76	+ {
77	- "bin_filename": "DiagFlat_19805c5e58e043224cccf51a06108010ae13i32",	77	+ "bin_filename": "DiagFlat_19805c5e58e043224cccf51a06108010ae13i32",
78	- "inputs": [	78	+ "inputs": [
79	- {	79	+ {
80	- "name": "x",	80	+ "name": "x",
81	- "index": 0,	81	+ "index": 0,
82	- "dtype": "int32",	82	+ "dtype": "int32",
83	- "format": "ND",	83	+ "format": "ND",
84	- "paramType": "required",	84	+ "paramType": "required",
85	- "shape": [	85	+ "shape": [
86	- -2	86	+ -2
87	- ],	87	+ ],
88	- "dtype_match_mode": "DtypeByte"	88	+ "dtype_match_mode": "DtypeByte"
89	- }	89	+ }
90	- ],	90	+ ],
91	- "outputs": [	91	+ "outputs": [
92	- {	92	+ {
93	- "name": "y",	93	+ "name": "y",
94	- "index": 0,	94	+ "index": 0,
95	- "dtype": "int32",	95	+ "dtype": "int32",
96	- "format": "ND",	96	+ "format": "ND",
97	- "paramType": "required",	97	+ "paramType": "required",
98	- "shape": [	98	+ "shape": [
99	- -2	99	+ -2
100	- ],	100	+ ],
101	- "dtype_match_mode": "DtypeByte"	101	+ "dtype_match_mode": "DtypeByte"
102	- }	102	+ }
103	- ],	103	+ ],
104	- "attrs": [	104	+ "attrs": [
105	- {	105	+ {
106	- "name": "diagonal",	106	+ "name": "diagonal",
107	- "dtype": "int",	107	+ "dtype": "int",
108	- "value": null	108	+ "value": null
109	- }	109	+ }
110	- ]	110	+ ]
111	- },	111	+ },
112	- {	112	+ {
113	- "bin_filename": "DiagFlat_1980d2db1a80c523e7e59a032c95969880af64",	113	+ "bin_filename": "DiagFlat_1980d2db1a80c523e7e59a032c95969880af64",
114	- "inputs": [	114	+ "inputs": [
115	- {	115	+ {
116	- "name": "x",	116	+ "name": "x",
117	- "index": 0,	117	+ "index": 0,
118	- "dtype": "int64",	118	+ "dtype": "int64",
119	- "format": "ND",	119	+ "format": "ND",
120	- "paramType": "required",	120	+ "paramType": "required",
121	- "shape": [	121	+ "shape": [
122	- -2	122	+ -2
123	- ],	123	+ ],
124	- "dtype_match_mode": "DtypeByte"	124	+ "dtype_match_mode": "DtypeByte"
125	- }	125	+ }
126	- ],	126	+ ],
127	- "outputs": [	127	+ "outputs": [
128	- {	128	+ {
129	- "name": "y",	129	+ "name": "y",
130	- "index": 0,	130	+ "index": 0,
131	- "dtype": "int64",	131	+ "dtype": "int64",
132	- "format": "ND",	132	+ "format": "ND",
133	- "paramType": "required",	133	+ "paramType": "required",
134	- "shape": [	134	+ "shape": [
135	- -2	135	+ -2
136	- ],	136	+ ],
137	- "dtype_match_mode": "DtypeByte"	137	+ "dtype_match_mode": "DtypeByte"
138	- }	138	+ }
139	- ],	139	+ ],
140	- "attrs": [	140	+ "attrs": [
141	- {	141	+ {
142	- "name": "diagonal",	142	+ "name": "diagonal",
143	- "dtype": "int",	143	+ "dtype": "int",
144	- "value": null	144	+ "value": null
145	- }	145	+ }
146	- ]	146	+ ]
147	- }	147	+ }
148	- ]	148	+ ]
149	- }	149	+}
150	-

Mconversion/diag_flat/op_host/config/ascend950/diag_flat_simplified_key.ini+1-12

	@@ -1,13 +1,2 @@
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	[DiagFlat]	1	[DiagFlat]
13	-default=0	2	+default=0

Mconversion/diag_flat/op_host/diag_flat_def.cpp+9-1

	@@ -40,7 +40,15 @@ public:
40	this->Attr("diagonal").AttrType(OPTIONAL).Int(0);	40	this->Attr("diagonal").AttrType(OPTIONAL).Int(0);
41	this->AICore().AddConfig("ascend910b");	41	this->AICore().AddConfig("ascend910b");
42	this->AICore().AddConfig("ascend910_93");	42	this->AICore().AddConfig("ascend910_93");
43	- this->AICore().AddConfig("ascend950");	43	+
		44	+ OpAICoreConfig config950;
		45	+ config950.DynamicCompileStaticFlag(true)
		46	+ .DynamicRankSupportFlag(true)
		47	+ .DynamicShapeSupportFlag(true)
		48	+ .NeedCheckSupportFlag(false)
		49	+ .PrecisionReduceFlag(true)
		50	+ .ExtendCfgInfo("opFile.value", "diag_flat_apt");
		51	+ this->AICore().AddConfig("ascend950", config950);
44		52
45	OpAICoreConfig config_310p_910;	53	OpAICoreConfig config_310p_910;
46	config_310p_910.Input("x")	54	config_310p_910.Input("x")

Aconversion/diag_flat/op_kernel/arch35/diag_flat_simd.h+160-0

@@ -0,0 +1,160 @@
	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 diag_flat_simd.h
	13	+ * \brief DiagFlat SIMD outer shell + SIMT kernel (arch35, DAV_3510)
	14	+ *
	15	+ * Design: DESIGN.md v2.1 Sec 3.4
	16	+ *
	17	+ * Execution flow (per tile):
	18	+ * AllocTensor → asc_vf_call (SIMT: fill UB — diagonal=input, rest=0)
	19	+ * → EnQue → DeQue → DataCopyPad (UB → GM, continuous burst) → FreeTensor
	20	+ *
	21	+ * Note: SIMT kernel writes to EVERY UB position (even non-diagonal), matching diag_v2's
	22	+ * unconditional-write pattern, to ensure compiler emits full-vector stores with
	23	+ * 32B-aligned addresses. Conditional "scatter" writes would produce masked partial-vector
	24	+ * stores whose start address may not be 32B-aligned → VEC_ERROR 340 (ub addr misaligned).
	25	+ *
	26	+ * (ref: diag_v2/op_kernel/arch35/diag_v2.h)
	27	+ */
	28	+
	29	+#ifndef __DIAG_FLAT_ARCH35_SIMD_H__
	30	+#define __DIAG_FLAT_ARCH35_SIMD_H__
	31	+
	32	+#include "kernel_operator.h"
	33	+#include "simt_api/asc_simt.h"
	34	+#include "diag_flat_tiling_data.h" // kept for backward compat; fields now in DiagV2Arch35TilingData
	35	+
	36	+using namespace AscendC;
	37	+
	38	+constexpr int32_t DIAG_FLAT_BUFFER_NUM = 2;
	39	+constexpr uint32_t DIAG_FLAT_THREAD_NUM = 1024;
	40	+
	41	+// ================================================================
	42	+// SIMT kernel: fill UB buffer (diagonal ← xGm, non-diagonal ← 0)
	43	+// Every thread always writes — no divergent UB store.
	44	+// (ref: batch_to_space_nd_simt.h HALF_THREAD_NUM_LAUNCH_BOUND=1024 for arch35)
	45	+// ================================================================
	46	+
	47	+template <typename T>
	48	+__simt_vf__ __aicore__ LAUNCH_BOUND(DIAG_FLAT_THREAD_NUM) void SimtDiagFlat(
	49	+ __gm__ T* xGm, // input (1D contiguous)
	50	+ int64_t tileOffset, // global output element offset of this tile
	51	+ int64_t tileElems, // number of output elements in this tile
	52	+ int64_t outWidth, // W = N + \|k\|
	53	+ uint64_t outWidthMagic, // fast div magic for outWidth
	54	+ uint64_t outWidthShift, // fast div shift for outWidth
	55	+ int64_t diagonal, // k (signed)
	56	+ int64_t numInput, // N = total input elements
	57	+ __ubuf__ T* yUb) // UB output buffer (packed) — LAST
	58	+{
	59	+ for (int64_t pos = threadIdx.x; pos < tileElems; pos += blockDim.x) {
	60	+ int64_t globalPos = tileOffset + pos;
	61	+
	62	+ // Fast unsigned division replaces slow int64_t / and %
	63	+ // (ref: batch_to_space_nd_simt.h Simt::UintDiv pattern)
	64	+ uint64_t row = Simt::UintDiv(static_cast<uint64_t>(globalPos), outWidthMagic, outWidthShift);
	65	+ uint64_t col = static_cast<uint64_t>(globalPos) - row * static_cast<uint64_t>(outWidth);
	66	+
	67	+ int64_t inputIdx = (diagonal >= 0) ? static_cast<int64_t>(row) : static_cast<int64_t>(col);
	68	+ bool onDiag = (static_cast<int64_t>(col) - static_cast<int64_t>(row) == diagonal);
	69	+ bool inRange = (inputIdx >= 0) && (inputIdx < numInput);
	70	+ bool doRead = onDiag && inRange;
	71	+
	72	+ // Always read from GM (non-divergent), then branchless select
	73	+ T gmVal = xGm[doRead ? inputIdx : 0];
	74	+ yUb[pos] = doRead ? gmVal : static_cast<T>(0);
	75	+ }
	76	+}
	77	+
	78	+// ================================================================
	79	+// SIMD outer shell: pipe/buffer management + SIMT launch + DMA output
	80	+// ================================================================
	81	+
	82	+template <typename T>
	83	+class DiagFlatSimd {
	84	+public:
	85	+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y,
	86	+ const DiagFlatArch35TilingData* tilingData)
	87	+ {
	88	+ td_ = tilingData;
	89	+ xGm_.SetGlobalBuffer((__gm__ T*)x);
	90	+ yGm_.SetGlobalBuffer((__gm__ T*)y);
	91	+
	92	+ uint64_t blockIdx = GetBlockIdx();
	93	+ tileOffset_ = blockIdx * td_->outPerCore;
	94	+ remaining_ = min(td_->outPerCore, td_->outTotal - tileOffset_);
	95	+
	96	+ // InitBuffer 3rd arg = per-buffer bytes (ref: diag_v2, batch_to_space_nd).
	97	+ uint32_t bufBytes = static_cast<uint32_t>(
	98	+ ((td_->tileLength + DIAG_FLAT_THREAD_NUM - 1) / DIAG_FLAT_THREAD_NUM) *
	99	+ DIAG_FLAT_THREAD_NUM * sizeof(T));
	100	+ pipe_.InitBuffer(outQueue_, DIAG_FLAT_BUFFER_NUM, bufBytes);
	101	+
	102	+ // Pre-compute fast division magic for outWidth (ref: batch_to_space_nd pattern)
	103	+ GetUintDivMagicAndShift(outWidthMagic_, outWidthShift_,
	104	+ static_cast<uint64_t>(td_->outWidth));
	105	+ }
	106	+
	107	+ __aicore__ inline void Process()
	108	+ {
	109	+ if (remaining_ <= 0) return;
	110	+
	111	+ int64_t end = tileOffset_ + remaining_;
	112	+ int64_t curTileLen;
	113	+ for (int64_t curStart = tileOffset_; curStart < end; curStart += td_->tileLength) {
	114	+ curTileLen = min(td_->tileLength, end - curStart);
	115	+ // 1. Allocate UB buffer
	116	+ auto outBuf = outQueue_.AllocTensor<T>();
	117	+ auto yUb = reinterpret_cast<__ubuf__ T*>(outBuf.GetPhyAddr());
	118	+
	119	+
	120	+ // 2. SIMT: fill UB (diagonal ← input, non-diagonal ← 0)
	121	+ // Always use fixed threadNum, same as batch_to_space_nd pattern.
	122	+ // Loop condition `pos < tileElems` naturally filters idle threads.
	123	+ asc_vf_call<SimtDiagFlat<T>>(
	124	+ dim3(DIAG_FLAT_THREAD_NUM),
	125	+ (__gm__ T*)xGm_.GetPhyAddr(),
	126	+ curStart,
	127	+ curTileLen,
	128	+ td_->outWidth,
	129	+ outWidthMagic_,
	130	+ outWidthShift_,
	131	+ td_->diagonal,
	132	+ td_->numInput,
	133	+ yUb);
	134	+
	135	+ // 3. EnQue → DeQue (sync)
	136	+ outQueue_.EnQue(outBuf);
	137	+ LocalTensor<T> readyBuf = outQueue_.DeQue<T>();
	138	+
	139	+ // 4. DMA: continuous burst output
	140	+ DataCopyPad(yGm_[curStart], readyBuf,
	141	+ {1, static_cast<uint32_t>(curTileLen * sizeof(T)), 0, 0, 0});
	142	+
	143	+ // 5. Free UB buffer
	144	+ outQueue_.FreeTensor(outBuf);
	145	+ }
	146	+ }
	147	+
	148	+private:
	149	+ const DiagFlatArch35TilingData* td_;
	150	+ GlobalTensor<T> xGm_;
	151	+ GlobalTensor<T> yGm_;
	152	+ TPipe pipe_;
	153	+ TQue<QuePosition::VECOUT, DIAG_FLAT_BUFFER_NUM> outQueue_;
	154	+ int64_t tileOffset_;
	155	+ int64_t remaining_;
	156	+ uint64_t outWidthMagic_;
	157	+ uint64_t outWidthShift_;
	158	+};
	159	+
	160	+#endif // __DIAG_FLAT_ARCH35_SIMD_H__

Aconversion/diag_flat/op_kernel/arch35/diag_flat_tiling_data.h+35-0

@@ -0,0 +1,35 @@
	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 diag_flat_tiling_data.h
	13	+ * \brief DiagFlat TilingData struct definition (arch35, DAV_3510)
	14	+ *
	15	+ * Design basis: DESIGN.md v2.1 Sec 3.2
	16	+ * Standard C++ POD struct.
	17	+ * (ref: diag_v2/op_kernel/arch35/diag_v2_tiling_data.h)
	18	+ */
	19	+
	20	+#ifndef __DIAG_FLAT_ARCH35_TILING_DATA_H__
	21	+#define __DIAG_FLAT_ARCH35_TILING_DATA_H__
	22	+
	23	+#include <cstdint>
	24	+
	25	+struct DiagFlatArch35TilingData {
	26	+ int64_t numInput; // N = numel(x)
	27	+ int64_t diagonal; // k (signed)
	28	+ int64_t outWidth; // W = N + \|k\|
	29	+ int64_t outTotal; // W * W
	30	+ int64_t outPerCore; // ceil(outTotal / realCoreNum), elements per core
	31	+ int64_t tileLength; // max elements per tile
	32	+ int64_t realCoreNum; // actual number of cores used
	33	+};
	34	+
	35	+#endif // __DIAG_FLAT_ARCH35_TILING_DATA_H__

Aconversion/diag_flat/op_kernel/arch35/diag_flat_tiling_key.h+34-0

@@ -0,0 +1,34 @@
	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 diag_flat_tiling_key.h
	13	+ * \brief DiagFlat TilingKey template parameter definition (arch35, DAV_3510)
	14	+ *
	15	+ * Single TilingKey 3501: SIMD+SIMT hybrid, covers all 13 dtypes.
	16	+ * (ref: diag_v2/op_kernel/arch35/diag_v2_tiling_key.h)
	17	+ */
	18	+
	19	+#ifndef __DIAG_FLAT_ARCH35_TILING_KEY_H__
	20	+#define __DIAG_FLAT_ARCH35_TILING_KEY_H__
	21	+
	22	+#include "ascendc/host_api/tiling/template_argument.h"
	23	+
	24	+ASCENDC_TPL_ARGS_DECL(DiagFlat,
	25	+ ASCENDC_TPL_UINT_DECL(ARCH35_KEY, 32, ASCENDC_TPL_UI_LIST, 3501)
	26	+);
	27	+
	28	+ASCENDC_TPL_SEL(
	29	+ ASCENDC_TPL_ARGS_SEL(
	30	+ ASCENDC_TPL_UINT_SEL(ARCH35_KEY, ASCENDC_TPL_UI_LIST, 3501)
	31	+ )
	32	+);
	33	+
	34	+#endif // __DIAG_FLAT_ARCH35_TILING_KEY_H__

Mconversion/diag_flat/op_kernel/diag_flat.cpp+0-1

	@@ -12,7 +12,6 @@
12	* \file diag_flat.cpp	12	* \file diag_flat.cpp
13	* \brief	13	* \brief
14	*/	14	*/
15	-
16	#include "diag_flat_nd_to_2d.h"	15	#include "diag_flat_nd_to_2d.h"
17	#include "diag_flat_nd_to_2d_with_few.h"	16	#include "diag_flat_nd_to_2d_with_few.h"
18	#include "diag_flat_nd_to_2d_b16_more64.h"	17	#include "diag_flat_nd_to_2d_b16_more64.h"

Aconversion/diag_flat/op_kernel/diag_flat_apt.cpp+55-0

@@ -0,0 +1,55 @@
	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 diag_flat_apt.cpp
	13	+ * \brief DiagFlat kernel entry for arch35 (DAV_3510, __NPU_ARCH__=3101)
	14	+ *
	15	+ * Single TilingKey 3501: SIMD+SIMT hybrid, covers all 13 dtypes via DTYPE_X.
	16	+ * (ref: diag_v2/op_kernel/diag_v2_apt.cpp same pattern)
	17	+ */
	18	+
	19	+#include "arch35/diag_flat_tiling_key.h"
	20	+#include "arch35/diag_flat_simd.h"
	21	+
	22	+#include <type_traits>
	23	+#include "kernel_operator.h"
	24	+
	25	+using namespace AscendC;
	26	+
	27	+template <int ARCH35_KEY>
	28	+__aicore__ inline void LaunchKernel(GM_ADDR x, GM_ADDR y,
	29	+ const DiagFlatArch35TilingData* td)
	30	+{
	31	+ // 8-byte types (complex64/float64/int64/uint64): alias through int64_t
	32	+ // because SIMT kernel cannot directly dereference __gm__ Complex<float>* or
	33	+ // __gm__ double*, nor static_cast<Complex<float>>(0) / static_cast<double>(0) in SIMT.
	34	+ if constexpr (sizeof(DTYPE_X) == sizeof(int64_t)) {
	35	+ DiagFlatSimd<int64_t> op;
	36	+ op.Init(x, y, td);
	37	+ op.Process();
	38	+ } else {
	39	+ DiagFlatSimd<DTYPE_X> op;
	40	+ op.Init(x, y, td);
	41	+ op.Process();
	42	+ }
	43	+}
	44	+
	45	+template <int ARCH35_KEY>
	46	+__global__ __aicore__ void diag_flat(
	47	+ GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
	48	+{
	49	+ SetSysWorkspace(workspace);
	50	+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
	51	+ REGISTER_TILING_DEFAULT(DiagFlatArch35TilingData);
	52	+ GET_TILING_DATA_WITH_STRUCT(DiagFlatArch35TilingData, tilingData, tiling);
	53	+
	54	+ LaunchKernel<ARCH35_KEY>(x, y, &tilingData);
	55	+}

Dconversion/diag_flat/tests/CMakeLists.txt+0-16

	@@ -1,16 +0,0 @@
1	-# ----------------------------------------------------------------------------
2	-# Copyright (c) Huawei Technologies Co., Ltd. 2025. All rights reserved.
3	-# This file is a part of the CANN Open Software.
4	-# Licensed under 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	-file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12	-foreach(SUB_DIR ${CURRENT_DIRS})
13	- if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
14	- add_subdirectory(${SUB_DIR})
15	- endif()
16	-endforeach()

Dconversion/diag_flat/tests/ut/CMakeLists.txt+0-16

	@@ -1,16 +0,0 @@
1	-# ----------------------------------------------------------------------------
2	-# Copyright (c) Huawei Technologies Co., Ltd. 2025. All rights reserved.
3	-# This file is a part of the CANN Open Software.
4	-# Licensed under 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	-file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12	-foreach(SUB_DIR ${CURRENT_DIRS})
13	- if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
14	- add_subdirectory(${SUB_DIR})
15	- endif()
16	-endforeach()

Dconversion/diag_flat/tests/ut/op_host/CMakeLists.txt+0-21

	@@ -1,21 +0,0 @@
1	-# ----------------------------------------------------------------------------
2	-# Copyright (c) Huawei Technologies Co., Ltd. 2025. All rights reserved.
3	-# This file is a part of the CANN Open Software.
4	-# Licensed under 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	-if(UT_TEST_ALL OR OP_HOST_UT)
12	- add_modules_ut_sources(UT_NAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
13	- add_modules_ut_sources(UT_NAME ${OP_INFERSHAPE_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
14	-endif()
15	-
16	-file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
17	-foreach(SUB_DIR ${CURRENT_DIRS})
18	- if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
19	- add_subdirectory(${SUB_DIR})
20	- endif()
21	-endforeach()

Rconversion/diag_flat/tests/ut/op_host/test_diag_flat_tiling.cpp→conversion/diag_flat/tests/ut/op_host/arch32/test_diag_flat_tiling.cpp+1-1

	@@ -12,7 +12,7 @@
12	#include <fstream>	12	#include <fstream>
13	#include <vector>	13	#include <vector>
14	#include <gtest/gtest.h>	14	#include <gtest/gtest.h>
15	-#include "../../../../diag_v2/op_host/diag_v2_tiling.h"	15	+#include "../../../../../diag_v2/op_host/arch32/diag_v2_tiling.h"
16		16
17	#include "tiling_context_faker.h"	17	#include "tiling_context_faker.h"
18	#include "tiling_case_executor.h"	18	#include "tiling_case_executor.h"

Dconversion/diag_flat/tests/ut/op_kernel/CMakeLists.txt+0-27

	@@ -1,27 +0,0 @@
1	-# ----------------------------------------------------------------------------
2	-# Copyright (c) Huawei Technologies Co., Ltd. 2025. All rights reserved.
3	-# This file is a part of the CANN Open Software.
4	-# Licensed under 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	-if (UT_TEST_ALL OR OP_KERNEL_UT)
12	- # 需要将Tiling依赖的文件添加到CMakeLists.txt中
13	- # set(elewise_common_tiling_files
14	- # ${CANN_ROOT}/ops/built-in/op_tiling/runtime/elewise_tiling.cc
15	- # )
16	- # 算子自己的tiling文件路径
17	- set(diag_flat_tiling_files
18	- ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/diag_flat_tiling.cpp
19	- # ${elewise_common_tiling_files}
20	- )
21	- # 使用AddOpTestCase
22	- # param1：算子名称，以kernel方式命名
23	- # param2：soc版本，多个以分号分隔，例如："ascend950;ascend910b"
24	- # param3：自定义编译选项，一般填写测试的一种典型数据类型组合，不需要则传入空字符串，例如："-DDTYPE_X=float"，多个使用空格分隔，例如："-DDTYPE_X=float -DDTYPE_Y=float"
25	- # param4：该算子依赖的所有tiling源码文件
26	- AddOpTestCase(diag_flat "ascend910b" "-DDTYPE_X=float" "${diag_flat_tiling_files}")
27	-endif()

Dconversion/diag_flat/tests/ut/op_kernel/diag_flat_data/gen_data.py+0-137

	@@ -1,137 +0,0 @@
1	-#!/usr/bin/env python3
2	-# -- coding: utf-8 --
3	-# ----------------------------------------------------------------------------
4	-# This program is free software, you can redistribute it and/or modify it.
5	-# Copyright (c) 2025 Huawei Technologies Co., Ltd.
6	-# This file is a part of the CANN Open Software.
7	-# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
8	-# Please refer to the License for details. You may not use this file except in compliance with the License.
9	-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
10	-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
11	-# See LICENSE in the root of the software repository for the full text of the License.
12	-# ----------------------------------------------------------------------------
13	-
14	-import sys
15	-import os
16	-import numpy as np
17	-
18	-
19	-input1_int8 = np.arange(1, 255).astype(np.int8)
20	-input1_int16 = np.arange(1, 255).astype(np.int16)
21	-input1_int32 = np.arange(1, 255).astype(np.int32)
22	-input1_int64 = np.arange(1, 257).astype(np.int64)
23	-
24	-input2_int8 = np.arange(1, 260).astype(np.int8)
25	-input2_int16 = np.arange(1, 260).astype(np.int16)
26	-input2_int64 = np.arange(1, 260).astype(np.int64)
27	-
28	-
29	-input3_int8 = np.arange(1, 60).astype(np.int8)
30	-input3_int16 = np.arange(1, 60).astype(np.int16)
31	-input3_int64 = np.arange(1, 60).astype(np.int64)
32	-
33	-
34	-input_data1 = [(np.random.randint(1, 20, (1,)) +
35	- 1.j * np.random.randint(1, 20, (1,)))[0] for _ in range(1, 259 + 1)]
36	-input4_complex128 = np.asarray(input_data1).astype(np.complex128)
37	-
38	-input_data2 = [(np.random.randint(1, 20, (1,)) +
39	- 1.j * np.random.randint(1, 20, (1,)))[0] for _ in range(1, 59 + 1)]
40	-input5_complex128 = np.asarray(input_data2).astype(np.complex128)
41	-
42	-case0_params = {
43	- "input": input1_int8,
44	- "offset": -2,
45	- "d_type": np.int8,
46	-}
47	-
48	-case1_params = {
49	- "input": input1_int16,
50	- "offset": -2,
51	- "d_type": np.int16,
52	-}
53	-
54	-case2_params = {
55	- "input": input1_int32,
56	- "offset": -2,
57	- "d_type": np.int32,
58	-}
59	-
60	-case3_params = {
61	- "input": input1_int64,
62	- "offset": -2,
63	- "d_type": np.int64,
64	-}
65	-
66	-case4_params = {
67	- "input": input2_int8,
68	- "offset": 0,
69	- "d_type": np.int8,
70	-}
71	-
72	-case5_params = {
73	- "input": input2_int16,
74	- "offset": 0,
75	- "d_type": np.int16,
76	-}
77	-
78	-case6_params = {
79	- "input": input2_int64,
80	- "offset": 0,
81	- "d_type": np.int64,
82	-}
83	-
84	-case7_params = {
85	- "input": input3_int16,
86	- "offset": -2,
87	- "d_type": np.int16,
88	-}
89	-
90	-case8_params = {
91	- "input": input2_int64,
92	- "offset": -2,
93	- "d_type": np.int64,
94	-}
95	-
96	-case9_params = {
97	- "input": input4_complex128,
98	- "offset": 100,
99	- "d_type": np.complex128,
100	-}
101	-
102	-case10_params = {
103	- "input": input5_complex128,
104	- "offset": 0,
105	- "d_type": np.complex128,
106	-}
107	-
108	-
109	-test_cast = {
110	- 'case0': case0_params,
111	- 'case1': case1_params,
112	- 'case2': case2_params,
113	- 'case3': case3_params,
114	- 'case4': case4_params,
115	- 'case5': case5_params,
116	- 'case6': case6_params,
117	- 'case7': case7_params,
118	- 'case8': case8_params,
119	- 'case9': case9_params,
120	- 'case10':case10_params,
121	-}
122	-
123	-
124	-def gen_data_and_golden(case_num):
125	- case_params = test_cast[str(case_num)]
126	- input = case_params['input']
127	- d_type = case_params['d_type']
128	-
129	- input = input.astype(d_type)
130	- input.tofile('./input.bin')
131	- offset = case_params['offset']
132	- golden = np.diagflat(input, offset)
133	- golden.tofile('./golden.bin')
134	-
135	-
136	-if __name__ == "__main__":
137	- gen_data_and_golden(sys.argv[1])

Dconversion/diag_flat/tests/ut/op_kernel/diag_flat_data/gen_tiling.py+0-45

	@@ -1,45 +0,0 @@
1	-#!/usr/bin/env python3
2	-# -- coding: utf-8 --
3	-# ----------------------------------------------------------------------------
4	-# This program is free software, you can redistribute it and/or modify it.
5	-# Copyright (c) 2025 Huawei Technologies Co., Ltd.
6	-# This file is a part of the CANN Open Software.
7	-# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
8	-# Please refer to the License for details. You may not use this file except in compliance with the License.
9	-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
10	-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
11	-# See LICENSE in the root of the software repository for the full text of the License.
12	-# ----------------------------------------------------------------------------
13	-
14	-import numpy as np
15	-import sys
16	-
17	-caseflat_params = [326, 326, 0, 326, 3, 160, 6, 2101, 128, 59, 64, 64, 64, 64, 0, 0, 0, 0]
18	-case0_params = [326, 326, 0, 326, 3, 160, 6, 2101, 128, 0, 0, 0, 0, 0, 0, 0, 0, 0]
19	-case1_params = [326, 326, 0, 326, 3, 160, 6, 2101, 64, 0, 0, 0, 0, 0, 0, 0, 0, 0]
20	-case2_params = [64, 64, 0, 64, 2, 32, 32, 2405, 32, 0, 0, 0, 0, 0, 0, 0, 0, 0]
21	-case3_params = [128, 128, 0, 128, 2, 64, 64, 2104, 64, 0, 0, 0, 0, 0, 0, 0, 0, 0]
22	-case4_params = [326, 326, 0, 326, 3, 160, 6, 2101, 128, 0, 0, 0, 0, 0, 0, 0, 0, 0]
23	-case5_params = [64, 64, 0, 64, 2, 32, 32, 2405, 64, 0, 0, 0, 0, 0, 0, 0, 0, 0]
24	-
25	-params_info = {
26	- "case0": case0_params,
27	- "case1": case1_params,
28	- "case2": case2_params,
29	- "case3": case3_params,
30	- "case4": case4_params,
31	- "case5": case5_params,
32	- "caseflat": caseflat_params,
33	-}
34	-
35	-def main():
36	- params_list = params_info[sys.argv[1]] # python gen_tiling.py case0 sys.argv[1]="case0"
37	-
38	- base_params = np.array(params_list, dtype=np.int64)
39	-
40	- tiling_file = open("tiling.bin", "wb")
41	- base_params.tofile(tiling_file)
42	-
43	-
44	-if __name__ == '__main__':
45	- main()

Dconversion/diag_flat/tests/ut/op_kernel/test_diag_flat.cpp+0-45

	@@ -1,45 +0,0 @@
1	-/**
2	- * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3	- * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4	- * CANN Open Software License Agreement Version 2.0 (the "License").
5	- * Please refer to the License for details. You may not use this file except in compliance with the License.
6	- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7	- * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8	- * See LICENSE in the root of the software repository for the full text of the License.
9	- */
10	-
11	-#include <array>
12	-#include <vector>
13	-#include <iostream>
14	-#include <string>
15	-#include <fstream>
16	-#include <cstdio>
17	-#include <iomanip>
18	-#include <fcntl.h>
19	-#include <unistd.h>
20	-#include <sys/stat.h>
21	-#include "kernel_fp16.h"
22	-#include "gtest/gtest.h"
23	-#include "test_diag_flat_tiling.h"
24	-
25	-#ifdef __CCE_KT_TEST__
26	-#include <cstdint>
27	-#include "tikicpulib.h"
28	-#endif
29	-
30	-#define INFO_LOG(fmt, args...) fprintf(stdout, "[INFO] " fmt "\n", ##args)
31	-#define WARN_LOG(fmt, args...) fprintf(stdout, "[WARN] " fmt "\n", ##args)
32	-#define ERROR_LOG(fmt, args...) fprintf(stdout, "[ERROR] " fmt "\n", ##args)
33	-using namespace std;
34	-
35	-extern "C" __global__ __aicore__ void diag_flat(GM_ADDR input, GM_ADDR output, GM_ADDR tiling, GM_ADDR workspace);
36	-
37	-class diag_flat_test : public testing::Test {
38	-protected:
39	- static void SetUpTestCase() {
40	- std::cout << "========== diag_flat_test SetUp ==========\n" << std::endl;
41	- }
42	- static void TearDownTestCase() {
43	- std::cout << "========== diag_flat_test TearDown ==========\n" << std::endl;
44	- }
45	-};

Dconversion/diag_flat/tests/ut/op_kernel/test_diag_flat_tiling.h+0-68

	@@ -1,68 +0,0 @@
1	-/**
2	- * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3	- * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4	- * CANN Open Software License Agreement Version 2.0 (the "License").
5	- * Please refer to the License for details. You may not use this file except in compliance with the License.
6	- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7	- * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8	- * See LICENSE in the root of the software repository for the full text of the License.
9	- */
10	-
11	-#ifndef _DIAG_FLAT_TILING_H_
12	-#define _DIAG_FLAT_TILING_H_
13	-
14	-#include "kernel_tiling/kernel_tiling.h"
15	-
16	-struct DiagV2TilingDataTest {
17	- int64_t xWidth = 0;
18	- int64_t xHeight = 0;
19	- int64_t gmOffset = 0;
20	- int64_t numOut = 0;
21	- int64_t realCoreNum = 0;
22	- int64_t numPerCore = 0;
23	- int64_t tailNum = 0;
24	- int64_t tilingKey = 0;
25	- int64_t matrixRowLength = 0;
26	- int64_t inputNum = 0;
27	- int64_t usedCoreNum = 0;
28	- int64_t totalCoreNum = 0;
29	- int64_t normalCoreHandleNum = 0;
30	- int64_t lastCoreHandleNum = 0;
31	- int64_t diagonal = 0;
32	- int64_t align2 = 0;
33	- int64_t align3 = 0;
34	- int64_t align4 = 0;
35	-};
36	-
37	-#define DTYPE_X int64_t
38	-
39	-#pragma pack(1)
40	-
41	-#pragma pack()
42	-
43	-#define CONVERT_TILING_DATA(tilingStruct, tilingDataPointer, tilingPointer) \
44	- __ubuf__ tilingStruct* tilingDataPointer = \
45	- reinterpret_cast<__ubuf__ tilingStruct>((__ubuf__ uint8_t)(tilingPointer));
46	-
47	-#define INIT_TILING_DATA(tilingStruct, tilingDataPointer, tilingPointer) \
48	- CONVERT_TILING_DATA(tilingStruct, tilingDataPointer, tilingPointer);
49	-
50	-#define GET_TILING_DATA(tilingData, tilingPointer) \
51	- DiagV2TilingData tilingData; \
52	- INIT_TILING_DATA(DiagV2TilingData, tilingDataPointer, tilingPointer); \
53	- (tilingData).xWidth = tilingDataPointer->xWidth; \
54	- (tilingData).xHeight = tilingDataPointer->xHeight; \
55	- (tilingData).gmOffset = tilingDataPointer->gmOffset; \
56	- (tilingData).numOut = tilingDataPointer->numOut; \
57	- (tilingData).realCoreNum = tilingDataPointer->realCoreNum; \
58	- (tilingData).numPerCore = tilingDataPointer->numPerCore; \
59	- (tilingData).tailNum = tilingDataPointer->tailNum; \
60	- (tilingData).matrixRowLength = tilingDataPointer->matrixRowLength; \
61	- (tilingData).inputNum = tilingDataPointer->inputNum; \
62	- (tilingData).usedCoreNum = tilingDataPointer->usedCoreNum; \
63	- (tilingData).totalCoreNum = tilingDataPointer->totalCoreNum; \
64	- (tilingData).normalCoreHandleNum = tilingDataPointer->normalCoreHandleNum; \
65	- (tilingData).lastCoreHandleNum = tilingDataPointer->lastCoreHandleNum; \
66	- (tilingData).diagonal = tilingDataPointer->diagonal;
67	-
68	-#endif

Mconversion/diag_v2/CMakeLists.txt+6-1

	@@ -9,4 +9,9 @@
9	# See LICENSE in the root of the software repository for the full text of the License.	9	# See LICENSE in the root of the software repository for the full text of the License.
10	# ----------------------------------------------------------------------------	10	# ----------------------------------------------------------------------------
11		11
12	-add_all_modules_sources(OPTYPE diag_v2 ACLNNTYPE aclnn_exclude DEPENDENCIES diag_flat)	12	+set(SUPPORT_COMPUTE_UNIT "ascend310b" "ascend310p" "ascend910" "ascend910_93" "ascend910b" "kirin9030" "kirinx90" "ascend950")
		13	+set(SUPPORT_TILING_DIR "arch32" "arch32" "arch32" "arch32" "arch32" "arch32" "arch32" "arch35")
		14	+add_all_modules_sources(OPTYPE diag_v2 ACLNNTYPE aclnn_exclude
		15	+ COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT}
		16	+ TILING_DIR ${SUPPORT_TILING_DIR}
		17	+ DEPENDENCIES diag_flat)

Mconversion/diag_v2/op_api/aclnn_diag.cpp+1-1

	@@ -18,7 +18,7 @@
18	#include "aclnn_kernels/contiguous.h"	18	#include "aclnn_kernels/contiguous.h"
19	#include "diagv2.h"	19	#include "diagv2.h"
20	#include "conversion/fill/op_api/fill.h"	20	#include "conversion/fill/op_api/fill.h"
21	-#include "conversion/diag_flat/op_host/op_api/diag_flat.h"	21	+#include "conversion/diag_flat/op_api/diag_flat.h"
22	#include "aclnn_kernels/common/op_error_check.h"	22	#include "aclnn_kernels/common/op_error_check.h"
23	#include "opdev/common_types.h"	23	#include "opdev/common_types.h"
24	#include "opdev/data_type_utils.h"	24	#include "opdev/data_type_utils.h"

Dconversion/diag_v2/op_host/CMakeLists.txt+0-12

	@@ -1,12 +0,0 @@
1	-# ----------------------------------------------------------------------------
2	-# This program is free software, you can redistribute it and/or modify it.
3	-# Copyright (c) 2025 Huawei Technologies Co., Ltd.
4	-# This file is a part of the CANN Open Software.
5	-# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6	-# Please refer to the License for details. You may not use this file except in compliance with the License.
7	-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8	-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9	-# See LICENSE in the root of the software repository for the full text of the License.
10	-# ----------------------------------------------------------------------------
11	-
12	-add_modules_sources(OPTYPE diag_v2 ACLNNTYPE aclnn_exclude)

Rconversion/diag_v2/op_host/diag_v2_tiling.cpp→conversion/diag_v2/op_host/arch32/diag_v2_tiling.cpp+0-0

文件重命名但无更改。

Rconversion/diag_v2/op_host/diag_v2_tiling.h→conversion/diag_v2/op_host/arch32/diag_v2_tiling.h+0-0

文件重命名但无更改。

Aconversion/diag_v2/op_host/arch35/diag_v2_tiling.cpp+256-0

@@ -0,0 +1,256 @@
	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 diag_v2_tiling.cpp
	13	+ * \brief DiagV2 Host-side Tiling implementation (arch35, DAV_3510)
	14	+ *
	15	+ * Design: DESIGN.md v2.5
	16	+ *
	17	+ * One-way dependency: diag_v2 → diag_flat.
	18	+ * rank<=1 → call TilingDiagFlatArch35() (defined in diag_flat), fill own TilingData, IS_1D_INPUT=1
	19	+ * rank>=2 → local 2D→1D tiling, IS_1D_INPUT=0
	20	+ */
	21	+
	22	+#include "diag_v2_tiling.h"
	23	+#include "register/op_def_registry.h"
	24	+#include "op_common/log/log.h"
	25	+#include "op_common/op_host/util/math_util.h"
	26	+#include "op_common/op_host/util/platform_util.h"
	27	+#include "../../../diag_flat/op_host/arch35/diag_flat_tiling.h"
	28	+#include "../../op_kernel/arch35/diag_v2_tiling_data.h"
	29	+#include "../../op_kernel/arch35/diag_v2_tiling_key.h"
	30	+#include <algorithm>
	31	+#include <cstring>
	32	+
	33	+namespace optiling {
	34	+
	35	+using Ops::Base::CeilDiv;
	36	+
	37	+constexpr int64_t MIN_WORK_PER_CORE = 256;
	38	+constexpr int64_t TILE_LENGTH = 2048;
	39	+constexpr size_t ATTR_DIAGONAL_IDX = 0;
	40	+constexpr size_t WORKSPACE_NUM = 1;
	41	+constexpr uint32_t WS_SYS_SIZE = 0U;
	42	+
	43	+// ============================================================================
	44	+// Helper: Get platform info (ubSize, coreNum)
	45	+// ============================================================================
	46	+static ge::graphStatus GetPlatformInfo(gert::TilingContext* context,
	47	+ uint64_t* ubSize, int64_t* coreNum)
	48	+{
	49	+ fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo();
	50	+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr);
	51	+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
	52	+ *coreNum = ascendcPlatform.GetCoreNumAiv();
	53	+ OP_CHECK_IF(*coreNum == 0, OP_LOGE(context, "coreNum is 0"),
	54	+ return ge::GRAPH_FAILED);
	55	+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, *ubSize);
	56	+ OP_CHECK_IF(*ubSize == 0, OP_LOGE(context, "ubSize is 0"),
	57	+ return ge::GRAPH_FAILED);
	58	+ return ge::GRAPH_SUCCESS;
	59	+}
	60	+
	61	+// ============================================================================
	62	+// Helper: 32B alignment
	63	+// ============================================================================
	64	+static int64_t AlignUp(int64_t n, int64_t align)
	65	+{
	66	+ return ((n + align - 1) / align) * align;
	67	+}
	68	+
	69	+// ============================================================================
	70	+// Helper: Get dtype size
	71	+// ============================================================================
	72	+static int64_t GetDtypeSize(ge::DataType dataType)
	zhanw_coding6月16日建议直接使用 ge::GetSizeByDataType
	73	+{
	74	+ switch (dataType) {
	75	+ case ge::DT_FLOAT16: case ge::DT_BF16:
	76	+ case ge::DT_INT16: case ge::DT_UINT16:
	77	+ return 2;
	78	+ case ge::DT_FLOAT: case ge::DT_INT32: case ge::DT_UINT32:
	79	+ return 4;
	80	+ case ge::DT_DOUBLE: case ge::DT_INT64: case ge::DT_UINT64:
	81	+ case ge::DT_COMPLEX64:
	82	+ return 8;
	83	+ case ge::DT_INT8: case ge::DT_UINT8: case ge::DT_BOOL:
	84	+ return 1;
	85	+ default:
	86	+ return 4;
	87	+ }
	88	+}
	89	+
	90	+// ============================================================================
	91	+// 1D→2D via diag_flat (one-way: diag_v2 → diag_flat)
	92	+// ============================================================================
	93	+static inline ge::graphStatus ProcessDiagFlat(gert::TilingContext* context)
	94	+{
	95	+ DiagFlatTilingOutput out;
	96	+ OP_CHECK_IF(
	97	+ TilingDiagFlatArch35(context, &out) != ge::GRAPH_SUCCESS,
	98	+ OP_LOGE(context, "TilingDiagFlatArch35 error"),
	99	+ return ge::GRAPH_FAILED);
	100	+
	101	+ // Fill diag_v2's own TilingData from diag_flat's output
	102	+ DiagV2Arch35TilingData* tiling = context->GetTilingData<DiagV2Arch35TilingData>();
	103	+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
	104	+ OP_CHECK_IF(
	105	+ memset_s(tiling, sizeof(DiagV2Arch35TilingData), 0, sizeof(DiagV2Arch35TilingData)) != EOK,
	106	+ OP_LOGE(context, "set tiling data error"),
	107	+ return ge::GRAPH_FAILED);
	108	+
	109	+ tiling->diagonal = out.diagonal;
	110	+ tiling->realCoreNum = out.realCoreNum;
	111	+ tiling->tileLength = out.tileLength;
	112	+ tiling->numInput = out.numInput;
	113	+ tiling->outWidth = out.outWidth;
	114	+ tiling->outTotal = out.outTotal;
	115	+ tiling->outPerCore = out.outPerCore;
	116	+
	117	+ // Select diag_v2's own TilingKey: IS_1D_INPUT=1 → DiagFlatSimd kernel
	118	+ ASCENDC_TPL_SEL_PARAM(context, static_cast<uint32_t>(1));
	119	+
	120	+ size_t* currentWorkspace = context->GetWorkspaceSizes(WORKSPACE_NUM);
	121	+ OP_CHECK_NULL_WITH_CONTEXT(context, currentWorkspace);
	122	+ currentWorkspace[0] = WS_SYS_SIZE;
	123	+
	124	+ return ge::GRAPH_SUCCESS;
	125	+}
	126	+
	127	+// ============================================================================
	128	+// 2D→1D tiling (IS_1D_INPUT=0)
	129	+// ============================================================================
	130	+static ge::graphStatus Tiling2Dto1D(gert::TilingContext* context,
	131	+ int64_t hwCoreNum, int64_t dtypeSize,
	132	+ uint64_t ubSize)
	133	+{
	134	+ auto inputX = context->GetInputShape(0);
	135	+ OP_CHECK_NULL_WITH_CONTEXT(context, inputX);
	136	+ auto inputShape = inputX->GetStorageShape();
	137	+ int64_t xHeight = inputShape.GetDim(0);
	138	+ int64_t xWidth = inputShape.GetDim(1);
	139	+
	140	+ auto attrs = context->GetAttrs();
	141	+ OP_CHECK_NULL_WITH_CONTEXT(context, attrs);
	142	+ const int64_t* diagonalPtr = attrs->GetAttrPointer<int64_t>(ATTR_DIAGONAL_IDX);
	143	+ int64_t diagonal = (diagonalPtr != nullptr) ? *diagonalPtr : 0;
	144	+
	145	+ int64_t numOut = 0;
	146	+ int64_t gmOffset = 0;
	147	+
	148	+ if (xHeight == 0 \|\| xWidth == 0) {
	149	+ numOut = 0;
	150	+ } else if (diagonal >= 0) {
	151	+ if (diagonal >= xWidth) {
	152	+ numOut = 0;
	153	+ } else {
	154	+ numOut = std::min(xHeight, xWidth - diagonal);
	155	+ gmOffset = diagonal;
	156	+ }
	157	+ } else {
	158	+ if (-diagonal >= xHeight) {
	159	+ numOut = 0;
	160	+ } else {
	161	+ numOut = std::min(xHeight + diagonal, xWidth);
	162	+ gmOffset = -diagonal * xWidth;
	163	+ }
	164	+ }
	165	+
	166	+ int64_t realCoreNum = std::min(hwCoreNum, std::max<int64_t>(1, numOut / MIN_WORK_PER_CORE));
	167	+
	168	+ int64_t alignElems = std::max<int64_t>(1, 32 / dtypeSize);
	169	+ int64_t numPerCore = 0;
	170	+ if (numOut > 0) {
	171	+ numPerCore = AlignUp(CeilDiv(numOut, realCoreNum), alignElems);
	172	+ }
	173	+ int64_t tailNum = numOut - (realCoreNum - 1) * numPerCore;
	174	+
	175	+ int64_t tileLength = std::min(TILE_LENGTH, numPerCore);
	176	+ int64_t threadNum = std::min(TILE_LENGTH, numPerCore);
	177	+
	178	+ DiagV2Arch35TilingData* tiling = context->GetTilingData<DiagV2Arch35TilingData>();
	179	+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
	180	+ OP_CHECK_IF(
	181	+ memset_s(tiling, sizeof(DiagV2Arch35TilingData), 0, sizeof(DiagV2Arch35TilingData)) != EOK,
	182	+ OP_LOGE(context, "set tiling data error"),
	183	+ return ge::GRAPH_FAILED);
	184	+
	185	+ tiling->xWidth = xWidth;
	186	+ tiling->xHeight = xHeight;
	187	+ tiling->gmOffset = gmOffset;
	188	+ tiling->numOut = numOut;
	189	+ tiling->realCoreNum = realCoreNum;
	190	+ tiling->numPerCore = numPerCore;
	191	+ tiling->tailNum = tailNum;
	192	+ tiling->diagonal = diagonal;
	193	+ tiling->tileLength = tileLength;
	194	+ tiling->threadNum = threadNum;
	195	+
	196	+ context->SetLocalMemorySize(ubSize - 64 * 1024);
	zhanw_coding6月16日建议将这里的 64k 提取为常量，并说明其含义和作用
	197	+
	198	+ context->SetBlockDim(realCoreNum);
	199	+ return ge::GRAPH_SUCCESS;
	200	+}
	201	+
	202	+// ============================================================================
	203	+// Tiling entry
	204	+// ============================================================================
	205	+static ge::graphStatus DiagV2TilingFunc(gert::TilingContext* context)
	206	+{
	207	+ auto inputX = context->GetInputShape(0);
	208	+ OP_CHECK_NULL_WITH_CONTEXT(context, inputX);
	209	+ auto inputShape = inputX->GetStorageShape();
	210	+ int64_t rank = inputShape.GetDimNum();
	211	+
	212	+ // Route: rank<=1 → delegate to diag_flat, rank>=2 → 2D→1D
	213	+ if (rank <= 1) {
	214	+ return ProcessDiagFlat(context);
	215	+ }
	216	+
	217	+ // 2D→1D path
	218	+ uint64_t ubSize;
	219	+ int64_t hwCoreNum;
	220	+ OP_CHECK_IF(
	221	+ GetPlatformInfo(context, &ubSize, &hwCoreNum) != ge::GRAPH_SUCCESS,
	222	+ OP_LOGE(context, "GetPlatformInfo error"),
	223	+ return ge::GRAPH_FAILED);
	224	+
	225	+ auto inputDesc = context->GetInputDesc(0);
	226	+ OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc);
	227	+ int64_t dtypeSize = GetDtypeSize(inputDesc->GetDataType());
	228	+
	229	+ ge::graphStatus ret = Tiling2Dto1D(context, hwCoreNum, dtypeSize, ubSize);
	230	+ OP_CHECK_IF(ret != ge::GRAPH_SUCCESS, OP_LOGE(context, "Tiling2Dto1D error"), return ret);
	231	+
	232	+ size_t* currentWorkspace = context->GetWorkspaceSizes(WORKSPACE_NUM);
	233	+ OP_CHECK_NULL_WITH_CONTEXT(context, currentWorkspace);
	234	+ currentWorkspace[0] = WS_SYS_SIZE;
	235	+
	236	+ ASCENDC_TPL_SEL_PARAM(context, static_cast<uint32_t>(0)); // IS_1D_INPUT=0
	237	+
	238	+ return ge::GRAPH_SUCCESS;
	239	+}
	240	+
	241	+// ============================================================================
	242	+// TilingParse
	243	+// ============================================================================
	244	+static ge::graphStatus TilingParseForDiagV2([[maybe_unused]] gert::TilingParseContext* context)
	245	+{
	246	+ return ge::GRAPH_SUCCESS;
	247	+}
	248	+
	249	+// ============================================================================
	250	+// Tiling registration
	251	+// ============================================================================
	252	+IMPL_OP_OPTILING(DiagV2)
	253	+ .Tiling(DiagV2TilingFunc)
	254	+ .TilingParse<DiagV2CompileInfo>(TilingParseForDiagV2);
	255	+
	256	+} // namespace optiling

Aconversion/diag_v2/op_host/arch35/diag_v2_tiling.h+30-0

@@ -0,0 +1,30 @@
	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 diag_v2_tiling.h
	13	+ * \brief DiagV2 arch35 tiling declarations
	14	+ *
	15	+ * Design: DESIGN.md v2.5
	16	+ * One-way dependency: diag_v2 → diag_flat (includes diag_flat_tiling.h for 1D→2D delegation).
	17	+ */
	18	+
	19	+#ifndef __DIAG_V2_ARCH35_TILING_H__
	20	+#define __DIAG_V2_ARCH35_TILING_H__
	21	+
	22	+#include "register/op_impl_registry.h"
	23	+
	24	+namespace optiling {
	25	+
	26	+struct DiagV2CompileInfo {};
	27	+
	28	+} // namespace optiling
	29	+
	30	+#endif // __DIAG_V2_ARCH35_TILING_H__

Mconversion/diag_v2/op_host/diag_v2_def.cpp+29-1

	@@ -65,7 +65,35 @@ public:
65	ge::FORMAT_ND, ge::FORMAT_ND});	65	ge::FORMAT_ND, ge::FORMAT_ND});
66	this->AICore().AddConfig("ascend910b", config);	66	this->AICore().AddConfig("ascend910b", config);
67	this->AICore().AddConfig("ascend910_93", config);	67	this->AICore().AddConfig("ascend910_93", config);
68	- this->AICore().AddConfig("ascend950", config);	68	+
		69	+ OpAICoreConfig config950;
		70	+ config950.DynamicCompileStaticFlag(true)
		71	+ .DynamicRankSupportFlag(true)
		72	+ .DynamicShapeSupportFlag(true)
		73	+ .NeedCheckSupportFlag(false)
		74	+ .PrecisionReduceFlag(true)
		75	+ .ExtendCfgInfo("opFile.value", "diag_v2_apt");
		76	+ config950.Input("x")
		77	+ .ParamType(REQUIRED)
		78	+ .DataType(
		79	+ {ge::DT_INT8, ge::DT_UINT8, ge::DT_INT16, ge::DT_UINT16, ge::DT_INT32, ge::DT_UINT32, ge::DT_INT64,
		80	+ ge::DT_UINT64, ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16, ge::DT_DOUBLE, ge::DT_BOOL,
		81	+ ge::DT_COMPLEX64})
		82	+ .Format(
		83	+ {ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
		84	+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
		85	+ ge::FORMAT_ND, ge::FORMAT_ND});
		86	+ config950.Output("y")
		87	+ .ParamType(REQUIRED)
		88	+ .DataType(
		89	+ {ge::DT_INT8, ge::DT_UINT8, ge::DT_INT16, ge::DT_UINT16, ge::DT_INT32, ge::DT_UINT32, ge::DT_INT64,
		90	+ ge::DT_UINT64, ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16, ge::DT_DOUBLE, ge::DT_BOOL,
		91	+ ge::DT_COMPLEX64})
		92	+ .Format(
		93	+ {ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
		94	+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
		95	+ ge::FORMAT_ND, ge::FORMAT_ND});
		96	+ this->AICore().AddConfig("ascend950", config950);
		zhanw_coding6月16日这里的 config950 比原来的 config 少了 BF16
69		97
70	OpAICoreConfig config_kirin = GetKirinCoreConfig();	98	OpAICoreConfig config_kirin = GetKirinCoreConfig();
71	this->AICore().AddConfig("kirinx90", config_kirin);	99	this->AICore().AddConfig("kirinx90", config_kirin);

Aconversion/diag_v2/op_kernel/arch35/diag_v2.h+141-0

@@ -0,0 +1,141 @@
	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 diag_v2.h
	13	+ * \brief DiagV2 SIMD outer shell + SIMT kernel (arch35, DAV_3510)
	14	+ *
	15	+ * Design: DESIGN.md v2.3 Sec 3.4
	16	+ *
	17	+ * Execution flow (per tile):
	18	+ * AllocTensor → asc_vf_call (SIMT: GM gather → UB write)
	19	+ * → EnQue → DeQue → DataCopyPad (UB → GM) → FreeTensor
	20	+ *
	21	+ * (ref: tile_with_axis.h for SIMD class structure;
	22	+ * batch_to_space_nd_simt.h for __simt_vf__ pattern)
	23	+ */
	24	+
	25	+#ifndef __DIAG_V2_ARCH35_H__
	26	+#define __DIAG_V2_ARCH35_H__
	27	+
	28	+#include "kernel_operator.h"
	29	+#include "simt_api/asc_simt.h"
	30	+#include "diag_v2_tiling_data.h"
	31	+
	32	+using namespace AscendC;
	33	+
	34	+constexpr int32_t DIAG_V2_BUFFER_NUM = 2;
	35	+
	36	+// ================================================================
	37	+// SIMT kernel: diagonal element gather from GM → UB
	38	+// ================================================================
	39	+
	40	+template <typename T>
	41	+__simt_vf__ __aicore__ void SimtDiagV2(
	42	+ __gm__ T* xGm, // input matrix in GM
	43	+ __ubuf__ T* yUb, // output buffer in UB
	44	+ int64_t curTileStart, // start index (global output index)
	45	+ int64_t curTileLen, // number of elements in this tile
	46	+ int64_t xWidth, // input matrix width N
	47	+ int64_t diagonal) // diagonal offset k
	48	+{
	49	+ for (int64_t idx = threadIdx.x; idx < curTileLen; idx += blockDim.x) {
	50	+ int64_t globalIdx = curTileStart + idx;
	51	+
	52	+ // Compute 2D (row, col) from 1D output index
	53	+ int64_t row, col;
	54	+ if (diagonal >= 0) {
	55	+ row = globalIdx;
	56	+ col = globalIdx + diagonal;
	57	+ } else {
	58	+ row = globalIdx - diagonal;
	59	+ col = globalIdx;
	60	+ }
	61	+
	62	+ yUb[idx] = xGm[row * xWidth + col];
	63	+ }
	64	+}
	65	+
	66	+// ================================================================
	67	+// SIMD outer shell: pipe/buffer management + SIMT launch + DMA output
	68	+// ================================================================
	69	+
	70	+template <typename T>
	71	+class DiagV2Simd {
	72	+public:
	73	+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y,
	74	+ const DiagV2Arch35TilingData* tilingData)
	75	+ {
	76	+ td_ = tilingData;
	77	+ xGm_.SetGlobalBuffer((__gm__ T*)x);
	78	+ yGm_.SetGlobalBuffer((__gm__ T*)y);
	79	+ pipe_.InitBuffer(outQueue_, DIAG_V2_BUFFER_NUM,
	80	+ static_cast<uint32_t>(td_->tileLength * sizeof(T)));
	81	+
	82	+ uint64_t blockIdx = GetBlockIdx();
	83	+ startIdx_ = blockIdx * td_->numPerCore;
	84	+ endIdx_ = min(startIdx_ + td_->numPerCore, td_->numOut);
	85	+ }
	86	+
	87	+ __aicore__ inline void Process()
	88	+ {
	89	+ if (startIdx_ >= endIdx_) return;
	90	+
	91	+ int64_t curTileStart = startIdx_;
	92	+
	93	+ while (curTileStart < endIdx_) {
	94	+ int64_t curTileLen = min(td_->tileLength, endIdx_ - curTileStart);
	95	+
	96	+ // 1. Allocate UB buffer
	97	+ auto outBuf = outQueue_.AllocTensor<T>();
	98	+ __ubuf__ T* yUb = reinterpret_cast<__ubuf__ T*>(outBuf.GetPhyAddr());
	99	+
	100	+ // 2. Launch SIMT kernel: gather diagonal elements GM → UB
	101	+ asc_vf_call<SimtDiagV2<T>>(
	102	+ dim3(static_cast<uint32_t>(curTileLen)),
	103	+ (__gm__ T*)xGm_.GetPhyAddr(),
	104	+ yUb,
	105	+ curTileStart,
	106	+ curTileLen,
	107	+ td_->xWidth,
	108	+ td_->diagonal);
	109	+
	110	+ // 3. EnQue: mark buffer ready for consumer
	111	+ outQueue_.EnQue(outBuf);
	112	+
	113	+ // 4. DeQue: get ready buffer for DMA read
	114	+ LocalTensor<T> readyBuf = outQueue_.DeQue<T>();
	115	+
	116	+ // 5. DMA: Copy UB → GM output
	117	+ DataCopyParams copyParams;
	118	+ copyParams.blockCount = 1;
	119	+ copyParams.blockLen = static_cast<uint16_t>(curTileLen * sizeof(T));
	120	+ copyParams.srcStride = 0;
	121	+ copyParams.dstStride = 0;
	122	+ DataCopyPad(yGm_[curTileStart], readyBuf, copyParams);
	123	+
	124	+ // 6. Free UB buffer
	125	+ outQueue_.FreeTensor(outBuf);
	126	+
	127	+ curTileStart += curTileLen;
	128	+ }
	129	+ }
	130	+
	131	+private:
	132	+ const DiagV2Arch35TilingData* td_;
	133	+ GlobalTensor<T> xGm_;
	134	+ GlobalTensor<T> yGm_;
	135	+ TPipe pipe_;
	136	+ TQue<QuePosition::VECOUT, DIAG_V2_BUFFER_NUM> outQueue_;
	137	+ int64_t startIdx_;
	138	+ int64_t endIdx_;
	139	+};
	140	+
	141	+#endif // __DIAG_V2_ARCH35_H__

Aconversion/diag_v2/op_kernel/arch35/diag_v2_tiling_data.h+49-0

@@ -0,0 +1,49 @@
	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 diag_v2_tiling_data.h
	13	+ * \brief DiagV2 TilingData struct definition (arch35)
	14	+ *
	15	+ * Design basis: DESIGN.md v2.4 Sec 3.2
	16	+ * Standard C++ struct (no BEGIN_TILING_DATA_DEF / TILING_KEY_IS macros).
	17	+ *
	18	+ * Common fields are always valid. 2D→1D fields only valid when IS_1D_INPUT=0.
	19	+ * 1D→2D fields only valid when IS_1D_INPUT=1.
	20	+ */
	21	+
	22	+#ifndef __DIAG_V2_ARCH35_TILING_DATA_H__
	23	+#define __DIAG_V2_ARCH35_TILING_DATA_H__
	24	+
	25	+#include <cstdint>
	26	+
	27	+struct DiagV2Arch35TilingData {
	28	+ // === Common fields ===
	29	+ int64_t diagonal; // Diagonal offset k
	30	+ int64_t realCoreNum; // Actual number of cores used
	31	+ int64_t tileLength; // Max elements per tile
	32	+
	33	+ // === 2D→1D fields (IS_1D_INPUT=0) ===
	34	+ int64_t xWidth; // Input matrix width N
	35	+ int64_t xHeight; // Input matrix height M
	36	+ int64_t gmOffset; // GM offset of first diagonal element (linear index)
	37	+ int64_t numOut; // Total number of output elements
	38	+ int64_t numPerCore; // Elements per core (32B-aligned)
	39	+ int64_t tailNum; // Tail element count for the last core
	40	+ int64_t threadNum; // SIMT thread count (≤ 2048)
	41	+
	42	+ // === 1D→2D fields (IS_1D_INPUT=1, mirror DiagFlatArch35TilingData) ===
	43	+ int64_t numInput; // N = numel(x)
	44	+ int64_t outWidth; // W = N + \|k\|
	45	+ int64_t outTotal; // W * W
	46	+ int64_t outPerCore; // ceil(outTotal / realCoreNum), elements per core
	47	+};
	48	+
	49	+#endif // __DIAG_V2_ARCH35_TILING_DATA_H__

Aconversion/diag_v2/op_kernel/arch35/diag_v2_tiling_key.h+36-0

@@ -0,0 +1,36 @@
	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 diag_v2_tiling_key.h
	13	+ * \brief DiagV2 TilingKey template parameter definition (arch35)
	14	+ *
	15	+ * Design: DESIGN.md v2.4 Sec 3.1.2
	16	+ *
	17	+ * IS_1D_INPUT = 0: 2D→1D diagonal extraction (DiagV2Simd)
	18	+ * IS_1D_INPUT = 1: 1D→2D diagonal matrix construction (DiagFlatSimd)
	19	+ */
	20	+
	21	+#ifndef __DIAG_V2_ARCH35_TILING_KEY_H__
	22	+#define __DIAG_V2_ARCH35_TILING_KEY_H__
	23	+
	24	+#include "ascendc/host_api/tiling/template_argument.h"
	25	+
	26	+ASCENDC_TPL_ARGS_DECL(DiagV2,
	27	+ ASCENDC_TPL_UINT_DECL(IS_1D_INPUT, 8, ASCENDC_TPL_UI_LIST, 0, 1)
	28	+);
	29	+
	30	+ASCENDC_TPL_SEL(
	31	+ ASCENDC_TPL_ARGS_SEL(
	32	+ ASCENDC_TPL_UINT_SEL(IS_1D_INPUT, ASCENDC_TPL_UI_LIST, 0, 1)
	33	+ )
	34	+);
	35	+
	36	+#endif // __DIAG_V2_ARCH35_TILING_KEY_H__

Aconversion/diag_v2/op_kernel/diag_v2_apt.cpp+93-0

@@ -0,0 +1,93 @@
	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 diag_v2_apt.cpp
	13	+ * \brief DiagV2 kernel entry for arch35 (DAV_3510, __NPU_ARCH__=3101)
	14	+ *
	15	+ * Design: DESIGN.md v2.4 Sec 3.4.1
	16	+ *
	17	+ * Dual dispatch on IS_1D_INPUT:
	18	+ * IS_1D_INPUT=0 → 2D→1D diagonal extraction (DiagV2Simd kernel)
	19	+ * IS_1D_INPUT=1 → 1D→2D diagonal matrix construction (DiagFlatSimd kernel, shared with diag_flat)
	20	+ */
	21	+
	22	+#include "arch35/diag_v2_tiling_key.h"
	23	+#include "arch35/diag_v2.h"
	24	+#include "../diag_flat/arch35/diag_flat_simd.h"
	25	+
	26	+#include <type_traits>
	27	+#include "kernel_operator.h"
	28	+
	29	+using namespace AscendC;
	30	+
	31	+template <int IS_1D_INPUT>
	32	+__aicore__ inline void LaunchKernel(GM_ADDR x, GM_ADDR y,
	33	+ const DiagV2Arch35TilingData* td)
	34	+{
	35	+ if constexpr (IS_1D_INPUT == 0) {
	36	+ // 2D→1D: diagonal extraction
	37	+ // 8-byte types (complex64/double/int64/uint64): alias through int64_t
	38	+ // because SIMT cannot dereference __gm__ double* / __gm__ Complex<float>*,
	39	+ // and DataCopyPad does not support complex64.
	40	+ // bool (1 byte): alias through uint8_t because DataCopyPad does not support bool.
	41	+ if constexpr (sizeof(DTYPE_X) == sizeof(int64_t)) {
	42	+ DiagV2Simd<int64_t> op;
	43	+ op.Init(x, y, td);
	44	+ op.Process();
	45	+ } else if constexpr (std::is_same_v<DTYPE_X, bool>) {
	46	+ DiagV2Simd<uint8_t> op;
	47	+ op.Init(x, y, td);
	48	+ op.Process();
	49	+ } else {
	50	+ DiagV2Simd<DTYPE_X> op;
	51	+ op.Init(x, y, td);
	52	+ op.Process();
	53	+ }
	54	+ } else {
	55	+ // 1D→2D: construct diagonal matrix, delegate to diag_flat kernel
	56	+ DiagFlatArch35TilingData flatTd;
	57	+ flatTd.numInput = td->numInput;
	58	+ flatTd.diagonal = td->diagonal;
	59	+ flatTd.outWidth = td->outWidth;
	60	+ flatTd.outTotal = td->outTotal;
	61	+ flatTd.outPerCore = td->outPerCore;
	62	+ flatTd.tileLength = td->tileLength;
	63	+ flatTd.realCoreNum = td->realCoreNum;
	64	+
	65	+ // 8-byte types (complex64/double/int64/uint64): alias through int64_t
	66	+ // bool (1 byte): alias through uint8_t (DataCopyPad does not support bool)
	67	+ if constexpr (sizeof(DTYPE_X) == sizeof(int64_t)) {
	68	+ DiagFlatSimd<int64_t> op;
	69	+ op.Init(x, y, &flatTd);
	70	+ op.Process();
	71	+ } else if constexpr (std::is_same_v<DTYPE_X, bool>) {
	72	+ DiagFlatSimd<uint8_t> op;
	73	+ op.Init(x, y, &flatTd);
	74	+ op.Process();
	75	+ } else {
	76	+ DiagFlatSimd<DTYPE_X> op;
	77	+ op.Init(x, y, &flatTd);
	78	+ op.Process();
	79	+ }
	80	+ }
	81	+}
	82	+
	83	+template <int IS_1D_INPUT>
	84	+__global__ __aicore__ void diag_v2(
	85	+ GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
	86	+{
	87	+ SetSysWorkspace(workspace);
	88	+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
	89	+ REGISTER_TILING_DEFAULT(DiagV2Arch35TilingData);
	90	+ GET_TILING_DATA_WITH_STRUCT(DiagV2Arch35TilingData, tilingData, tiling);
	91	+
	92	+ LaunchKernel<IS_1D_INPUT>(x, y, &tilingData);
	93	+}

Dconversion/diag_v2/tests/ut/op_host/test_diag_v2_tiling.cpp+0-49

	@@ -1,49 +0,0 @@
1	-/**
2	- * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3	- * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4	- * CANN Open Software License Agreement Version 2.0 (the "License").
5	- * Please refer to the License for details. You may not use this file except in compliance with the License.
6	- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7	- * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8	- * See LICENSE in the root of the software repository for the full text of the License.
9	- */
10	-
11	-#include <iostream>
12	-#include <gtest/gtest.h>
13	-#include "tiling_context_faker.h"
14	-#include "tiling_case_executor.h"
15	-
16	-#include "../../../op_host/diag_v2_tiling.h"
17	-
18	-using namespace std;
19	-using namespace ge;
20	-
21	-class DiagV2Tiling : public testing::Test {
22	- protected:
23	- static void SetUpTestCase() {
24	- std::cout << "DiagV2Tiling SetUp" << std::endl;
25	- }
26	-
27	- static void TearDownTestCase() {
28	- std::cout << "DiagV2Tiling TearDown" << std::endl;
29	- }
30	-};
31	-
32	-TEST_F(DiagV2Tiling, ascend910B1_test_tiling__001)
33	-{
34	- optiling::DiagV2CompileInfo compileInfo = {48, 196608};
35	- gert::TilingContextPara tilingContextPara(
36	- "DiagV2",
37	- {
38	- {{{8, 8}, {8, 8}}, ge::DT_FLOAT16, ge::FORMAT_ND},
39	- },
40	- {
41	- {{{8}, {8}}, ge::DT_FLOAT16, ge::FORMAT_ND},
42	- },
43	- {gert::TilingContextPara::OpAttr("diagonal", Ops::Math::AnyValue::CreateFrom<int64_t>(0))},
44	- &compileInfo);
45	- uint64_t expectTilingKey = 2102;
46	- string expectTilingData = "8 8 0 8 1 16 8 2102 128 0 0 0 0 0 0 0 0 0 ";
47	- std::vector<size_t> expectWorkspaces = {16777216};
48	- ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
49	-}

Dconversion/diag_v2/tests/ut/op_kernel/CMakeLists.txt+0-28

	@@ -1,28 +0,0 @@
1	-# ----------------------------------------------------------------------------
2	-# This program is free software, you can redistribute it and/or modify it.
3	-# Copyright (c) 2025 Huawei Technologies Co., Ltd.
4	-# This file is a part of the CANN Open Software.
5	-# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6	-# Please refer to the License for details. You may not use this file except in compliance with the License.
7	-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8	-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9	-# See LICENSE in the root of the software repository for the full text of the License.
10	-# ----------------------------------------------------------------------------
11	-
12	-if (UT_TEST_ALL OR OP_KERNEL_UT)
13	- # 需要将Tiling依赖的文件添加到CMakeLists.txt中
14	- # set(elewise_common_tiling_files
15	- # ${CANN_ROOT}/ops/built-in/op_tiling/runtime/elewise_tiling.cc
16	- # )
17	- # 算子自己的tiling文件路径
18	- set(diag_v2_tiling_files
19	- ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/diag_v2_tiling.cpp
20	- ${CMAKE_CURRENT_SOURCE_DIR}/../../../../../conversion/diag_flat/op_host/diag_flat_tiling.cpp
21	- )
22	- # 使用AddOpTestCase
23	- # param1：算子名称，以kernel方式命名
24	- # param2：soc版本，多个以分号分隔，例如："ascend950;ascend910b"
25	- # param3：自定义编译选项，一般填写测试的一种典型数据类型组合，不需要则传入空字符串，例如："-DDTYPE_X=float"，多个使用空格分隔，例如："-DDTYPE_X=float -DDTYPE_Y=float"
26	- # param4：该算子依赖的所有tiling源码文件
27	- AddOpTestCase(diag_v2 "ascend910b" "-DDTYPE_X=int32_t -D__CCE_UT_TEST__" "${diag_v2_tiling_files}")
28	-endif()

Dconversion/diag_v2/tests/ut/op_kernel/diag_v2_data/gen_data.py+0-40

	@@ -1,40 +0,0 @@
1	-#!/usr/bin/env python3
2	-# -- coding: utf-8 --
3	-# ----------------------------------------------------------------------------
4	-# This program is free software, you can redistribute it and/or modify it.
5	-# Copyright (c) 2025 Huawei Technologies Co., Ltd.
6	-# This file is a part of the CANN Open Software.
7	-# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
8	-# Please refer to the License for details. You may not use this file except in compliance with the License.
9	-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
10	-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
11	-# See LICENSE in the root of the software repository for the full text of the License.
12	-# ----------------------------------------------------------------------------
13	-
14	-import sys
15	-import numpy as np
16	-import torch
17	-
18	-
19	-def gen_golden_data(params, data_dir):
20	- input_x = np.random.uniform(-100, 100, params[0].shape).astype(params[0].np_dtype)
21	- inputX = torch.as_tensor(input_x)
22	- outputY = torch.diag(inputX)
23	- golden = outputY.numpy().astype(params[1].np_dtype)
24	-
25	- input_x.tofile(str(data_dir / params[0].data_path))
26	- golden.tofile(str(data_dir / params[1].golden_path))
27	-
28	-
29	-def gen_golden_data_simple(x, y, dtype):
30	- input_x = np.random.uniform(-1, 100, [int(x), int(y)]).astype(dtype)
31	- inputX = torch.as_tensor(input_x)
32	- outputY = torch.diag(inputX)
33	- golden = outputY.numpy().astype(dtype)
34	-
35	- input_x.tofile("./input_x.bin")
36	- golden.tofile("./golden.bin")
37	-
38	-
39	-if __name__ == "__main__":
40	- gen_golden_data_simple(sys.argv[1], sys.argv[2], sys.argv[3])

Dconversion/diag_v2/tests/ut/op_kernel/diag_v2_data/gen_tiling.py+0-44

	@@ -1,44 +0,0 @@
1	-#!/usr/bin/env python3
2	-# -- coding: utf-8 --
3	-# ----------------------------------------------------------------------------
4	-# This program is free software, you can redistribute it and/or modify it.
5	-# Copyright (c) 2025 Huawei Technologies Co., Ltd.
6	-# This file is a part of the CANN Open Software.
7	-# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
8	-# Please refer to the License for details. You may not use this file except in compliance with the License.
9	-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
10	-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
11	-# See LICENSE in the root of the software repository for the full text of the License.
12	-# ----------------------------------------------------------------------------
13	-
14	-import numpy as np
15	-import sys
16	-
17	-
18	-case0_params = [326, 326, 0, 326, 3, 160, 6, 2101, 128, 0, 0, 0, 0, 0, 0, 0, 0, 0]
19	-case1_params = [326, 326, 0, 326, 3, 160, 6, 2101, 64, 0, 0, 0, 0, 0, 0, 0, 0, 0]
20	-case2_params = [64, 64, 0, 64, 2, 32, 32, 2405, 32, 0, 0, 0, 0, 0, 0, 0, 0, 0]
21	-case3_params = [128, 128, 0, 128, 2, 64, 64, 2104, 64, 0, 0, 0, 0, 0, 0, 0, 0, 0]
22	-case4_params = [326, 326, 0, 326, 3, 160, 6, 2101, 128, 0, 0, 0, 0, 0, 0, 0, 0, 0]
23	-case5_params = [64, 64, 0, 64, 2, 32, 32, 2405, 64, 0, 0, 0, 0, 0, 0, 0, 0, 0]
24	-
25	-params_info = {
26	- "case0": case0_params,
27	- "case1": case1_params,
28	- "case2": case2_params,
29	- "case3": case3_params,
30	- "case4": case4_params,
31	- "case5": case5_params,
32	-}
33	-
34	-def main():
35	- params_list = params_info[sys.argv[1]] # python gen_tiling.py case0 sys.argv[1]="case0"
36	-
37	- base_params = np.array(params_list, dtype=np.int64)
38	-
39	- tiling_file = open("tiling.bin", "wb")
40	- base_params.tofile(tiling_file)
41	-
42	-
43	-if __name__ == '__main__':
44	- main()

Dconversion/diag_v2/tests/ut/op_kernel/test_diag_v2.cpp+0-76

	@@ -1,76 +0,0 @@
1	-/**
2	- * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3	- * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4	- * CANN Open Software License Agreement Version 2.0 (the "License").
5	- * Please refer to the License for details. You may not use this file except in compliance with the License.
6	- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7	- * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8	- * See LICENSE in the root of the software repository for the full text of the License.
9	- */
10	-
11	-#include <array>
12	-#include <vector>
13	-#include "gtest/gtest.h"
14	-#include "../../../op_host/diag_v2_tiling.h"
15	-
16	-#ifdef __CCE_KT_TEST__
17	-#include "tikicpulib.h"
18	-#include "data_utils.h"
19	-#include "string.h"
20	-#include <iostream>
21	-#include <string>
22	-#endif
23	-
24	-#include <cstdint>
25	-
26	-using namespace std;
27	-//using namespace AscendC;
28	-
29	-extern "C" __global__ __aicore__ void diag_v2(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling);
30	-
31	-
32	-class diag_v2_test : public testing::Test {
33	- protected:
34	- static void SetUpTestCase() {
35	- cout << "diag_v2_test SetUp\n " << endl;
36	- }
37	- static void TearDownTestCase() {
38	- cout << "diag_v2_test TearDown\n" << endl;
39	- }
40	-};
41	-
42	-TEST_F(diag_v2_test, test_case_0) {
43	- // x
44	- size_t inputByteSize = 326 * 326 * sizeof(int8_t);
45	- // y
46	- size_t outputByteSize = 326 * sizeof(int8_t);
47	-
48	- size_t tiling_data_size = sizeof(DiagV2TilingData);
49	-
50	- uint8_t x = (uint8_t )AscendC::GmAlloc(inputByteSize);
51	- uint8_t y = (uint8_t )AscendC::GmAlloc(outputByteSize);
52	- uint8_t workspace = (uint8_t )AscendC::GmAlloc(4096 * 16);
53	- uint8_t tiling = (uint8_t )AscendC::GmAlloc(tiling_data_size);
54	- uint32_t numBlocks = 3;
55	- system("cp -r ../../../../math/diag_v2/tests/ut/op_kernel/diag_v2_data ./");
56	- system("chmod -R 755 ./diag_v2_data/");
57	- system("cd ./diag_v2_data/ && rm -rf ./*bin");
58	- system("cd ./diag_v2_data/ && python3 gen_data.py 326 326 int8");
59	- system("cd ./diag_v2_data/ && python3 gen_tiling.py case0");
60	-
61	- char * path_ = get_current_dir_name();
62	- string path(path_);
63	- ReadFile(path + "/diag_v2_data/input_x.bin", inputByteSize, x, inputByteSize);
64	- ReadFile(path + "/diag_v2_data/tiling.bin", tiling_data_size, tiling, tiling_data_size);
65	-
66	- ICPU_SET_TILING_KEY(2101);
67	- AscendC::SetKernelMode(KernelMode::AIV_MODE);
68	- ICPU_RUN_KF(diag_v2, numBlocks, x, y, workspace, tiling);
69	-
70	- AscendC::GmFree(x);
71	- AscendC::GmFree(y);
72	- AscendC::GmFree(workspace);
73	- AscendC::GmFree(tiling);
74	- free(path_);
75	-}
76	-