support fusion-compile ELF parsing flow in msobjdump-asc-tools-AtomGit

已合并

support fusion-compile ELF parsing flow in msobjdump #235

zhangyujia77创建于 4月9日从zhangyujia77/asc-tools： msobjdump 合入到 cann/asc-tools： master

support fusion-compile ELF parsing flow in msobjdump #235

已合并

zhangyujia77创建于 4月9日

从zhangyujia77/asc-tools msobjdump合入到cann/asc-tools master

共 10 个文件变更+1332-172

Mdocs/03_msobjdump.md+23-126

	@@ -1,145 +1,42 @@
1	-<!--声明：本文使用[Creative Commons License version 4.0](https://creativecommons.org/licenses/by/4.0/legalcode)许可协议，转载、引用或修改等操作请遵循此许可协议。-->	1	+# msobjdump
2	-# msobjdump
3	-
4		2
5	## 概述	3	## 概述
6	-本工具主要针对Kernel直调算子开发与工程化算子开发编译生成的算子ELF文件（Executable and Linkable Format）提供解析和解压功能，并将结果信息以可读形式呈现，方便开发者直观获得kernel文件信息。关于本工具的详细介绍请参考《[Ascend C算子开发](https://hiascend.com/document/redirect/CannCommunityOpdevAscendC)》中的“msobjdump工具”。	4	+本工具主要针对生成的算子ELF文件（Executable and Linkable Format）提供解析和解压功能，并将结果信息以可读形式呈现，方便开发者直观获得kernel文件信息。关于本工具的详细介绍请参考《[Ascend C算子开发](https://hiascend.com/document/redirect/CannCommunityOpdevAscendC)》中的“编程指南 > 附录 > msobjdump工具”。
7		5
8	-	6	+工具调用演示可参考[msobjdump样例](../examples/04_msobjdump/README.md)。
9	-## 工具安装
10	-
11	-本工具跟随CANN软件包发布，请参考[环境搭建](00_quick_start.md)进行使用工具前必要的环境准备。
12	-- 执行如下命令设置环境变量。
13	- ```
14	- source ${install_path}/latest/toolkit/bin/setenv.bash
15	- ```
16	-- 执行如下命令，若能正常显示--help或-h信息，则表示工具环境正常，功能可正常使用。
17	- ```
18	- msobjdump -h
19	- ```
20		7
21	## 命令格式	8	## 命令格式
22		9
23	- 解析ELF文件的命令	10	- 解析ELF文件的命令
24	- ```	11	+ ```bash
25	msobjdump --dump-elf <elf_file> [--verbose]	12	msobjdump --dump-elf <elf_file> [--verbose]
26	- ```	13	+ ```
27	--dump-elf <elf_file>为必选，表示待解析ELF文件路径。[--verbose]为可选，用于开启ELF文件中全量打印device信息功能。	14	--dump-elf <elf_file>为必选，表示待解析ELF文件路径。[--verbose]为可选，用于开启ELF文件中全量打印device信息功能。
28		15
29	- 解压ELF文件的命令	16	- 解压ELF文件的命令
30	- ```	17	+ ```bash
31	msobjdump --extract-elf <elf_file> [--out-dir <out_path>]	18	msobjdump --extract-elf <elf_file> [--out-dir <out_path>]
32	- ```	19	+ ```
33	--extract-elf <elf_file>为必选，表示待解析ELF文件路径。[--out-dir <out_path>]为可选，用于设置解压文件的落盘路径。	20	--extract-elf <elf_file>为必选，表示待解析ELF文件路径。[--out-dir <out_path>]为可选，用于设置解压文件的落盘路径。
34		21
35	- 获取ELF文件列表的命令	22	- 获取ELF文件列表的命令
36	- ```	23	+ ```bash
37	msobjdump --list-elf <elf_file>	24	msobjdump --list-elf <elf_file>
38	```	25	```
39	--list-elf <elf_file>为可选，获取ELF文件中包含的device信息文件列表，并打印显示。	26	--list-elf <elf_file>为可选，获取ELF文件中包含的device信息文件列表，并打印显示。
40		27
41	-## 使用样例（Kernel直调算子工程）	28	+ 下表为ELF文件中常见字段说明：
42	-以[matmul_kernellaunch](../examples/02_matmul_kernellaunch/README.md)算子为例（NPU模式），假设\${cmake_install_dir}为算子Cmake编译产物根目录，目录结构如下。	29	+\| 字段名 \| 含义 \| 是否必选 \| 打印说明 \|
43	-	30	+\| ---- \| ---- \| ---- \| ---- \|
44	-```	31	+\| `.ascend.meta. ${id}` \| 表示算子kernel函数名称，其中`${id}`表示meta信息的索引值。 \| 是 \| 不设置`--verbose`，默认打印。 \|
45	-out	32	+\| `VERSION` \| 表示版本号。 \| 是 \| 不设置`--verbose`，默认打印。 \|
46	-├── lib	33	+\| `DEBUG` \| 调试相关信息，包含如下两部分内容：<br>`debugBufSize`：调试信息需要的内存空间。<br>`debugOptions`：调试开关状态。取值如下：<br>`0`：调试开关关闭。<br>`1`：通过DumpTensor、printf打印进行调试。<br>`2`：通过assert断言进行调试。<br>`4`：通过时间戳打点功能进行调试。<br>`8`：通过内存越界检测进行调试。 \| 否 \| 不设置`--verbose`，默认打印。 \|
47	-│ ├── libascendc_kernels_npu.so	34	+\| `DYNAMIC_PARAM` \| 算子kernel函数是否启用动态参数。取值分别为：<br>`0`：关闭动态参数模式。<br>`1`：开启动态参数模式。 \| 否 \| 不设置`--verbose`，默认打印。 \|
48	-├── include	35	+\| `OPTIONAL_PARAM` \| 可选参数信息，包含如下两部分内容：<br>`optionalInputMode`：可选输入在算子kernel函数中是否需要占位。<br>`0`：可选输入不占位。<br>`1`：可选输入占位。<br>`optionalOutputMode`：可选输出在算子kernel函数中是否需要占位。<br>`0`：可选输出不占位。<br>`1`：可选输出占位。 \| 否 \| 不设置`--verbose`，默认打印。 \|
49	-│ ├── ascendc_kernels_npu	36	+\| `KERNEL_TYPE` \| 表示kernel函数运行时core类型。 \| 否 \| 不设置`--verbose`，默认打印。 \|
50	-│ ├── aclrtlaunch_matmul_custom.h	37	+\| `CROSS_CORE_SYNC` \| 表示硬同步syncall类型。<br>`USE_SYNC`：使用硬同步。<br>`NO_USE_SYNC`：不使用硬同步。 \| 否 \| 不设置`--verbose`，默认打印。 \|
51	-│ ├── aclrtlaunch_triple_chevrons_func.h	38	+\| `MIX_TASK_RATION` \| 表示kernel函数运行时的Cube核/Vector核占比分配类型。 \| 否 \| 不设置`--verbose`，默认打印。 \|
52	-├── bin	39	+\| `DETERMINISTIC_INFO` \| 表示算子是否为确定性计算。<br>`0`：不确定计算。<br>`1`：确定性计算。 \| 否 \| 不设置`--verbose`，默认打印。 \|
53	-│ ├── ascendc_kernels_bbit	40	+\| `BLOCK_NUM` \| 表示算子执行核数，该字段当前暂不支持，只打印默认值`0xFFFFFFFF`。 \| 否 \| 不设置`--verbose`，默认打印。 \|
54	-```	41	+\| `FUNCTION_ENTRY` \| 算子TilingKey的值。 \| 否 \| 不设置`--verbose`，默认打印。 \|
55	-	42	+\| `elf header infos` \| 包括ELF Header、Section Headers、Key to Flags、Program Headers、Symbol表等信息。 \| 否 \| 设置`--verbose`，开启全量打印。 \|
56	-工具对编译生成的库文件（如.so、.a等）进行解析和解压，功能实现命令样例如下：
57	-- 解析包含device信息的库文件
58	- 支持两种打印方式。
59	- - 简单打印
60	- ```
61	- msobjdump --dump-elf ${cmake_install_dir}/out/lib/libascendc_kernels_npu.so
62	- ```
63	- 执行上述命令，终端打印基础device信息，示例如下：
64	-
65	- ```
66	- ===========================
67	- [VERSION]: 1
68	- [TYPE COUNT]: 1
69	- ===========================
70	- [ELF FILE 0]: ascendxxxb1_ascendc_kernels_npu_0_mix.o
71	- [KERNEL TYPE]: mix
72	- [KERNEL LEN]: 511560
73	- [ASCEND META]: None
74	- ```
75	- - 全量打印
76	- ```
77	- msobjdump --dump-elf ${cmake_install_dir}/out/lib/libascendc_kernels_npu.so --verbose
78	- ```
79	- 执行上述命令，终端打印所有device信息，示例如下：
80	-
81	- ```
82	- ===========================
83	- [VERSION]: 1
84	- [TYPE COUNT]: 1
85	- ===========================
86	- [ELF FILE 0]: ascendxxxb1_ascendc_kernels_npu_0_mix.o
87	- [KERNEL TYPE]: mix
88	- [KERNEL LEN]: 511560
89	- [ASCEND META]: None
90	- ====== [elf heard infos] ======
91	- ELF Header:
92	- Magic: 7f 45 4c 46 02 01 01 00 00 00 00 00 00 00 00 00
93	- Class: ELF64
94	- Data: 2's complement, little endian
95	- Version: 1 (current)
96	- OS/ABI: UNIX - System V
97	- ABI Version: 0
98	- Type: EXEC (Executable file)
99	- Machine: <unknown>: 0x1029
100	- Version: 0x1
101	- Entry point address: 0x0
102	- Start of program headers: 64 (bytes into file)
103	- Start of section headers: 510280 (bytes into file)
104	- Flags: 0x940000
105	- Size of this header: 64 (bytes)
106	- Size of program headers: 56 (bytes)
107	- Number of program headers: 2
108	- Size of section headers: 64 (bytes)
109	- Number of section headers: 20
110	- Section header string table index: 18
111	-
112	- Section Headers:
113	- [Nr] Name Type Address Off Size ES Flg Lk Inf Al
114	- [ 0] NULL 0000000000000000 000000 000000 00 0 0 0
115	- [ 1] .text PROGBITS 0000000000000000 0000b0 010a08 00 AX 0 0 4
116	- .....................................................................................
117	- [19] .strtab STRTAB 0000000000000000 071278 00b6cb 00 0 0 1
118	- Key to Flags:
119	- W (write), A (alloc), X (execute), M (merge), S (strings), I (info),
120	- L (link order), O (extra OS processing required), G (group), T (TLS),
121	- C (compressed), x (unknown), o (OS specific), E (exclude),
122	- D (mbind), p (processor specific)
123	-
124	- There are no section groups in this file.
125	-
126	- Program Headers:
127	- Type Offset VirtAddr PhysAddr FileSiz MemSiz Flg Align
128	- LOAD 0x0000b0 0x0000000000000000 0x0000000000000000 0x010aa8 0x010aa8 R E 0x1000
129	- GNU_STACK 0x000000 0x0000000000000000 0x0000000000000000 0x000000 0x000000 RW 0
130	- ......
131	- ```
132	-- 解压包含device信息的库文件并落盘
133	- ```
134	- msobjdump --extract-elf ${cmake_install_dir}/out/lib/libascendc_kernels_npu.so
135	- ```
136	- 执行上述命令，默认在当前执行路径下落盘ascendxxxb1_ascendc_kernels_npu_0_mix.o文件。
137	-- 获取包含device信息的库文件列表
138	- ```
139	- msobjdump --list-elf ${cmake_install_dir}/out/lib/libascendc_kernels_npu.so
140	- ```
141	- 执行上述命令，终端会打印所有文件，屏显信息形如：
142	-
143	- ```
144	- ELF file 0: ascendxxxb1_ascendc_kernels_npu_0_mix.o
145	- ```

Aexamples/04_msobjdump/CMakeLists.txt+39-0

@@ -0,0 +1,39 @@
	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	+
	12	+cmake_minimum_required(VERSION 3.16)
	13	+
	14	+find_package(ASC REQUIRED)
	15	+
	16	+project(kernel_samples LANGUAGES ASC CXX)
	17	+
	18	+add_executable(demo
	19	+ matmul_leakyrelu.asc
	20	+)
	21	+
	22	+target_link_libraries(demo PRIVATE
	23	+ tiling_api
	24	+ register
	25	+ platform
	26	+ m
	27	+ dl
	28	+)
	29	+
	30	+# ======================================================================================
	31	+# NPU 编译选项配置
	32	+#
	33	+# 说明：
	34	+# - 需根据实际部署的 NPU 硬件架构选择对应的 `npu-arch` 参数。
	35	+# ======================================================================================
	36	+target_compile_options(demo PRIVATE
	37	+ $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-2201>
	38	+ # $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-3510>
	39	+)

Aexamples/04_msobjdump/README.md+204-0

@@ -0,0 +1,204 @@
	1	+# msobjdump样例
	2	+
	3	+## 概述
	4	+
	5	+本样例基于MatmulLeakyRelu算子，演示融合编译场景下`msobjdump`工具的使用方式。样例通过编译[matmul_leakyrelu.asc](./matmul_leakyrelu.asc)生成融合编译产物，再对生成的ELF文件执行解析。`msobjdump`工具的详细说明请参考[msobjdump工具](../../docs/03_msobjdump.md)。
	6	+
	7	+## 支持的产品
	8	+
	9	+- Ascend 950PR/Ascend 950DT
	10	+- Atlas A3 训练系列产品/Atlas A3 推理系列产品
	11	+- Atlas A2 训练系列产品/Atlas A2 推理系列产品
	12	+
	13	+## 目录结构介绍
	14	+
	15	+```
	16	+├── 04_msobjdump
	17	+│ ├── CMakeLists.txt // 编译工程文件
	18	+│ ├── data_utils.h // 数据读入写出函数
	19	+│ ├── matmul_leakyrelu.asc // Ascend C算子实现 & 调用样例
	20	+│ └── scripts
	21	+│ ├── gen_data.py // 输入数据和真值数据生成脚本文件
	22	+│ └── verify_result.py // 真值对比文件
	23	+```
	24	+
	25	+## 算子描述
	26	+
	27	+- 样例功能：
	28	+ MatmulLeakyRelu的计算公式为：
	29	+ ```
	30	+ C = A * B + Bias
	31	+ C = C > 0 ? C : C * 0.001
	32	+ ```
	33	+ 样例参数M = 1024，K = 256，N = 640，样例规格如下表所示：
	34	+ <table>
	35	+ <tr><td rowspan="1" align="center">样例类型(OpType)</td><td colspan="4" align="center">MatmulLeakyRelu</td></tr>
	36	+ <tr><td rowspan="4" align="center">样例输入</td><td align="center">name</td><td align="center">shape</td><td align="center">data type</td><td align="center">format</td></tr>
	37	+ <tr><td align="center">A</td><td align="center">[M, K]</td><td align="center">float16</td><td align="center">ND</td></tr>
	38	+ <tr><td align="center">B</td><td align="center">[K, N]</td><td align="center">float16</td><td align="center">ND</td></tr>
	39	+ <tr><td align="center">Bias</td><td align="center">[N]</td><td align="center">float</td><td align="center">ND</td></tr>
	40	+ <tr><td rowspan="1" align="center">样例输出</td><td align="center">C</td><td align="center">[M, N]</td><td align="center">float</td><td align="center">ND</td></tr>
	41	+ <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">matmul_leakyrelu_custom</td></tr>
	42	+ </table>
	43	+
	44	+- 样例实现：
	45	+ - 实现流程
	46	+ - 通过GenerateTiling实现host侧的Tiling计算
	47	+ - 通过CalcGMOffset完成分核计算
	48	+ - 通过Iterate接口完成矩阵乘计算
	49	+ - 通过LeakyRelu实现激活函数计算
	50	+
	51	+ - 调用实现
	52	+ 使用内核调用符<<<>>>调用核函数。
	53	+
	54	+## 编译运行
	55	+
	56	+在本样例根目录下执行如下步骤，编译并执行算子。
	57	+- 配置环境变量
	58	+ 请根据当前环境上CANN开发套件包的[安装方式](../../docs/00_quick_start.md#prepare&install)，选择对应配置环境变量的命令。
	59	+ - 默认路径，root用户安装CANN软件包
	60	+ ```bash
	61	+ source /usr/local/Ascend/cann/set_env.sh
	62	+ ```
	63	+
	64	+ - 默认路径，非root用户安装CANN软件包
	65	+ ```bash
	66	+ source $HOME/Ascend/cann/set_env.sh
	67	+ ```
	68	+
	69	+ - 指定路径install_path，安装CANN软件包
	70	+ ```bash
	71	+ source ${install_path}/cann/set_env.sh
	72	+ ```
	73	+
	74	+- 执行如下命令，若能正常显示-h信息，则表示工具环境正常，功能可正常使用。
	75	+ ```bash
	76	+ msobjdump -h
	77	+ ```
	78	+
	79	+- 样例执行
	80	+ ```bash
	81	+ mkdir -p build && cd build; # 创建并进入build目录
	82	+ cmake ..;make -j; # 编译工程
	83	+ python3 ../scripts/gen_data.py # 生成测试输入数据
	84	+ ./demo # 执行编译生成的可执行程序，执行样例
	85	+ python3 ../scripts/verify_result.py output/output.bin output/golden.bin # 验证输出结果是否正确，确认算法逻辑正确
	86	+ ```
	87	+ 执行结果如下，说明精度对比成功。
	88	+ ```bash
	89	+ test pass!
	90	+ ```
	91	+
	92	+- 调用msobjdump工具解析
	93	+ 样例中的`demo`为融合编译生成的ELF文件。若该ELF中包含`.aicore_binary`段，`msobjdump`会自动提取该段内容并继续解析，无需手工拆分中间文件。
	94	+
	95	+ - 解析融合编译产物
	96	+ ```bash
	97	+ msobjdump --dump-elf ./demo
	98	+ ```
	99	+
	100	+ 本样例实际输出如下：
	101	+
	102	+ ```Plain Text
	103	+ .ascend.meta META INFO
	104	+ RUNTIME_IMPLICIT_INFO: DOUBLE_PAGE_TABLE_ADDR
	105	+ RUNTIME_IMPLICIT_INFO: FFTS_ADDR
	106	+ VERSION: 1
	107	+ RUNTIME_IMPLICIT_INFO: DOUBLE_PAGE_TABLE_ADDR
	108	+ RUNTIME_IMPLICIT_INFO: FFTS_ADDR
	109	+ VERSION: 1
	110	+ RUNTIME_IMPLICIT_INFO: SIMD_TRACE_SPACE
	111	+ .ascend.meta. [0]: _Z23matmul_leakyrelu_customPhS_S_S_S_N7AscendC6tiling11TCubeTilingE_mix_aic
	112	+ KERNEL_TYPE: MIX_AIC_MAIN
	113	+ CROSS_CORE_SYNC: USE_SYNC
	114	+ MIX_TASK_RATION: [1:2]
	115	+ .ascend.meta. [0]: _Z23matmul_leakyrelu_customPhS_S_S_S_N7AscendC6tiling11TCubeTilingE_mix_aiv
	116	+ KERNEL_TYPE: MIX_AIC_MAIN
	117	+ CROSS_CORE_SYNC: USE_SYN
	118	+ ```
	119	+
	120	+ - 全量打印融合编译产物中的device信息
	121	+ ```bash
	122	+ msobjdump --dump-elf ./demo --verbose
	123	+ ```
	124	+ 本样例实际输出如下：
	125	+ ```Plain Text
	126	+ .ascend.meta META INFO
	127	+ RUNTIME_IMPLICIT_INFO: DOUBLE_PAGE_TABLE_ADDR
	128	+ RUNTIME_IMPLICIT_INFO: FFTS_ADDR
	129	+ VERSION: 1
	130	+ RUNTIME_IMPLICIT_INFO: DOUBLE_PAGE_TABLE_ADDR
	131	+ RUNTIME_IMPLICIT_INFO: FFTS_ADDR
	132	+ VERSION: 1
	133	+ RUNTIME_IMPLICIT_INFO: SIMD_TRACE_SPACE
	134	+ .ascend.meta. [0]: _Z23matmul_leakyrelu_customPhS_S_S_S_N7AscendC6tiling11TCubeTilingE_mix_aic
	135	+ KERNEL_TYPE: MIX_AIC_MAIN
	136	+ CROSS_CORE_SYNC: USE_SYNC
	137	+ MIX_TASK_RATION: [1:2]
	138	+ .ascend.meta. [0]: _Z23matmul_leakyrelu_customPhS_S_S_S_N7AscendC6tiling11TCubeTilingE_mix_aiv
	139	+ KERNEL_TYPE: MIX_AIC_MAIN
	140	+ CROSS_CORE_SYNC: USE_SYNC
	141	+ MIX_TASK_RATION: [1:2]
	142	+ ====== [elf heard infos] ======
	143	+ ELF Header:
	144	+ Magic: 7f 45 4c 46 02 01 01 00 00 00 00 00 00 00 00 00
	145	+ Class: ELF64
	146	+ Data: 2's complement, little endian
	147	+ Version: 1 (current)
	148	+ OS/ABI: UNIX - System V
	149	+ ABI Version: 0
	150	+ Type: EXEC (Executable file)
	151	+ Machine: <unknown>: 0x1029
	152	+ Version: 0x1
	153	+ Entry point address: 0x0
	154	+ Start of program headers: 64 (bytes into file)
	155	+ Start of section headers: 33504 (bytes into file)
	156	+ Flags: 0x940000
	157	+ Size of this header: 64 (bytes)
	158	+ Size of program headers: 56 (bytes)
	159	+ Number of program headers: 3
	160	+ Size of section headers: 64 (bytes)
	161	+ Number of section headers: 16
	162	+ Section header string table index: 14
	163	+
	164	+ Section Headers:
	165	+ [Nr] Name Type Address Off Size ES Flg Lk Inf Al
	166	+ [ 0] NULL 0000000000000000 000000 000000 00 0 0 0
	167	+ [ 1] .text PROGBITS 0000000000000000 0000e8 006c94 00 AX 0 0 4
	168	+ ......................................................................................
	169	+ [15] .strtab STRTAB 0000000000000000 007ce0 0005fc 00 0 0 1
	170	+ Key to Flags:
	171	+ W (write), A (alloc), X (execute), M (merge), S (strings), I (info),
	172	+ L (link order), O (extra OS processing required), G (group), T (TLS),
	173	+ C (compressed), x (unknown), o (OS specific), E (exclude),
	174	+ D (mbind), p (processor specific)
	175	+
	176	+ There are no section groups in this file.
	177	+
	178	+ Program Headers:
	179	+ Type Offset VirtAddr PhysAddr FileSiz MemSiz Flg Align
	180	+ LOAD 0x0000e8 0x0000000000000000 0x0000000000000000 0x006ca7 0x006ca7 R E 0x1000
	181	+ LOAD 0x0070e8 0x0000000000007000 0x0000000000007000 0x000210 0x000210 RW 0x1000
	182	+ GNU_STACK 0x000000 0x0000000000000000 0x0000000000000000 0x000000 0x000000 RW 0
	183	+
	184	+ ......
	185	+ ```
	186	+
	187	+ - 获取融合编译产物中的ELF文件列表
	188	+ ```bash
	189	+ msobjdump --list-elf ./demo
	190	+ ```
	191	+
	192	+ 对当前样例产物，终端提示如下：
	193	+
	194	+ ```Plain Text
	195	+ ELF file 0: demo.aicore.o
	196	+ ```
	197	+
	198	+ - 解压融合编译产物中的ELF文件
	199	+ ```bash
	200	+ mkdir -p objdump_out
	201	+ msobjdump --extract-elf ./demo
	202	+ ```
	203	+
	204	+ 执行上述命令，默认在当前执行路径下落盘`demo.aicore.o`文件，若需指定路径可通过--out-dir进行设置。

Aexamples/04_msobjdump/data_utils.h+94-0

@@ -0,0 +1,94 @@
	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	+
	12	+/* !
	13	+ * \file data_utils.h
	14	+ * \brief
	15	+ */
	16	+
	17	+#ifndef DATA_UTILS_H
	18	+#define DATA_UTILS_H
	19	+#include <fcntl.h>
	20	+#include <sys/stat.h>
	21	+#include <unistd.h>
	22	+#include <fstream>
	23	+
	24	+#define ERROR_LOG(fmt, args...) fprintf(stdout, "[ERROR] " fmt "\n", ##args)
	25	+
	26	+bool ReadFile(const std::string &filePath, size_t &fileSize, void *buffer, size_t bufferSize)
	27	+{
	28	+ struct stat sBuf;
	29	+ int fileStatus = stat(filePath.data(), &sBuf);
	30	+ if (fileStatus == -1) {
	31	+ ERROR_LOG("failed to get file");
	32	+ return false;
	33	+ }
	34	+ if (S_ISREG(sBuf.st_mode) == 0) {
	35	+ ERROR_LOG("%s is not a file, please enter a file", filePath.c_str());
	36	+ return false;
	37	+ }
	38	+
	39	+ std::ifstream file;
	40	+ file.open(filePath, std::ios::binary);
	41	+ if (!file.is_open()) {
	42	+ ERROR_LOG("Open file failed. path = %s", filePath.c_str());
	43	+ return false;
	44	+ }
	45	+
	46	+ std::filebuf *buf = file.rdbuf();
	47	+ size_t size = buf->pubseekoff(0, std::ios::end, std::ios::in);
	48	+ if (size == 0) {
	49	+ ERROR_LOG("file size is 0");
	50	+ file.close();
	51	+ return false;
	52	+ }
	53	+ if (size > bufferSize) {
	54	+ ERROR_LOG("file size is larger than buffer size");
	55	+ file.close();
	56	+ return false;
	57	+ }
	58	+ buf->pubseekpos(0, std::ios::in);
	59	+ buf->sgetn(static_cast<char *>(buffer), size);
	60	+ fileSize = size;
	61	+ file.close();
	62	+ return true;
	63	+}
	64	+
	65	+/**
	66	+ * @brief Write data to file
	67	+ * @param [in] filePath: file path
	68	+ * @param [in] buffer: data to write to file
	69	+ * @param [in] size: size to write
	70	+ * @return write result
	71	+ */
	72	+bool WriteFile(const std::string &filePath, const void *buffer, size_t size)
	73	+{
	74	+ if (buffer == nullptr) {
	75	+ ERROR_LOG("Write file failed. buffer is nullptr");
	76	+ return false;
	77	+ }
	78	+
	79	+ int fd = open(filePath.c_str(), O_RDWR \| O_CREAT \| O_TRUNC, S_IRUSR \| S_IWRITE);
	80	+ if (fd < 0) {
	81	+ ERROR_LOG("Open file failed. path = %s", filePath.c_str());
	82	+ return false;
	83	+ }
	84	+
	85	+ size_t writeSize = write(fd, buffer, size);
	86	+ (void)close(fd);
	87	+ if (writeSize != size) {
	88	+ ERROR_LOG("Write file Failed.");
	89	+ return false;
	90	+ }
	91	+
	92	+ return true;
	93	+}
	94	+#endif // DATA_UTILS_H

Aexamples/04_msobjdump/matmul_leakyrelu.asc+322-0

@@ -0,0 +1,322 @@
	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	+
	12	+/* !
	13	+ * \file matmul_leakyrelu.asc
	14	+ * \brief
	15	+ */
	16	+
	17	+#include "data_utils.h"
	18	+#include "kernel_tiling/kernel_tiling.h"
	19	+#include "tiling/platform/platform_ascendc.h"
	20	+#include "tiling/tiling_api.h"
	21	+#include "acl/acl.h"
	22	+#include "kernel_operator.h"
	23	+#include "lib/matmul_intf.h"
	24	+
	25	+__aicore__ inline uint32_t Ceiling(uint32_t a, uint32_t b)
	26	+{
	27	+ return (a + b - 1) / b;
	28	+}
	29	+
	30	+template <typename aType, typename bType, typename cType, typename biasType> class MatmulLeakyKernel {
	31	+public:
	32	+ __aicore__ inline MatmulLeakyKernel(){};
	33	+ __aicore__ inline void Init(GM_ADDR a, GM_ADDR b, GM_ADDR bias, GM_ADDR c, GM_ADDR workspace,
	34	+ const TCubeTiling &tiling, AscendC::TPipe *pipe);
	35	+ __aicore__ inline void Process(AscendC::TPipe *pipe);
	36	+
	37	+ __aicore__ inline void MatmulCompute();
	38	+ __aicore__ inline void LeakyReluCompute();
	39	+ __aicore__ inline void CopyOut(uint32_t count);
	40	+ __aicore__ inline void CalcOffset(int32_t blockIdx, const TCubeTiling &tiling, int32_t &offsetA, int32_t &offsetB,
	41	+ int32_t &offsetC, int32_t &offsetBias);
	42	+
	43	+ matmul::Matmul<matmul::MatmulType<AscendC::TPosition::GM, CubeFormat::ND, aType>,
	44	+ matmul::MatmulType<AscendC::TPosition::GM, CubeFormat::ND, bType>,
	45	+ matmul::MatmulType<AscendC::TPosition::VECIN, CubeFormat::ND, cType>,
	46	+ matmul::MatmulType<AscendC::TPosition::GM, CubeFormat::ND, biasType>>
	47	+ matmulObj;
	48	+
	49	+ AscendC::GlobalTensor<aType> aGlobal;
	50	+ AscendC::GlobalTensor<bType> bGlobal;
	51	+ AscendC::GlobalTensor<cType> cGlobal;
	52	+ AscendC::GlobalTensor<biasType> biasGlobal;
	53	+ AscendC::LocalTensor<cType> reluOutLocal;
	54	+ TCubeTiling tiling;
	55	+ AscendC::TQue<AscendC::TPosition::VECOUT, 1> reluOutQueue_;
	56	+};
	57	+
	58	+/**
	59	+ * @brief Set matmulLeaky input and output gm addr of current core.
	60	+ * @param a: A matrix gm addr.
	61	+ * @param b: B matrix gm addr.
	62	+ * @param bias: Bias gm addr.
	63	+ * @param c: C matrix gm addr.
	64	+ * @param workspace: Temporary gm space addr required by matmul calc.
	65	+ * @param tiling: matmul tiling data.
	66	+ * @param pipe: Global memory and sync management TPipe object.
	67	+ * @retval None
	68	+ */
	69	+template <typename aType, typename bType, typename cType, typename biasType>
	70	+__aicore__ inline void MatmulLeakyKernel<aType, bType, cType, biasType>::Init(GM_ADDR a, GM_ADDR b, GM_ADDR bias,
	71	+ GM_ADDR c, GM_ADDR workspace,
	72	+ const TCubeTiling &tiling, AscendC::TPipe *pipe)
	73	+{
	74	+ this->tiling = tiling;
	75	+ aGlobal.SetGlobalBuffer(reinterpret_cast<__gm__ aType >(a), tiling.M tiling.Ka);
	76	+ bGlobal.SetGlobalBuffer(reinterpret_cast<__gm__ bType >(b), tiling.Kb tiling.N);
	77	+ cGlobal.SetGlobalBuffer(reinterpret_cast<__gm__ cType >(c), tiling.M tiling.N);
	78	+ biasGlobal.SetGlobalBuffer(reinterpret_cast<__gm__ biasType *>(bias), tiling.N);
	79	+
	80	+ int32_t offsetA, offsetB, offsetC, offsetBias;
	81	+ CalcOffset(AscendC::GetBlockIdx(), tiling, offsetA, offsetB, offsetC, offsetBias); // Calculate the gm offset based on the blockidx.
	82	+ aGlobal = aGlobal[offsetA];
	83	+ bGlobal = bGlobal[offsetB];
	84	+ cGlobal = cGlobal[offsetC];
	85	+ biasGlobal = biasGlobal[offsetBias];
	86	+ pipe->InitBuffer(reluOutQueue_, 1, tiling.baseM * tiling.baseN * sizeof(cType)); // Init output buffer.
	87	+}
	88	+
	89	+/**
	90	+ * @brief Main process of matmul calculation
	91	+ * @param pipe: Global memory and sync management TPipe object.
	92	+ * @retval None
	93	+ */
	94	+template <typename aType, typename bType, typename cType, typename biasType>
	95	+__aicore__ inline void MatmulLeakyKernel<aType, bType, cType, biasType>::Process(AscendC::TPipe *pipe)
	96	+{
	97	+ uint32_t computeRound = 0;
	98	+
	99	+ AscendC::printf("-----------%d\n", computeRound);
	100	+ matmulObj.SetTensorA(aGlobal);
	101	+ matmulObj.SetTensorB(bGlobal);
	102	+ matmulObj.SetBias(biasGlobal);
	103	+ while (matmulObj.template Iterate<true>()) { // Once Iterate, compute baseM * baseN, sync is set true here.
	104	+ MatmulCompute(); // Get matmul compute result.
	105	+ LeakyReluCompute(); // Compute leakyRelu.
	106	+ CopyOut(computeRound); // Copy leakyRelu out result to GM.
	107	+ computeRound++;
	108	+ }
	109	+ matmulObj.End();
	110	+}
	111	+
	112	+template <typename aType, typename bType, typename cType, typename biasType>
	113	+__aicore__ inline void MatmulLeakyKernel<aType, bType, cType, biasType>::MatmulCompute()
	114	+{
	115	+ reluOutLocal = reluOutQueue_.AllocTensor<cType>();
	116	+ matmulObj.template GetTensorC<true>(reluOutLocal, false, true);
	117	+}
	118	+
	119	+template <typename aType, typename bType, typename cType, typename biasType>
	120	+__aicore__ inline void MatmulLeakyKernel<aType, bType, cType, biasType>::LeakyReluCompute()
	121	+{
	122	+ LeakyRelu(reluOutLocal, reluOutLocal, (cType)0.001, tiling.baseM * tiling.baseN);
	123	+ reluOutQueue_.EnQue(reluOutLocal);
	124	+}
	125	+
	126	+/**
	127	+ * @brief Copy leakyRelu out result to GM.
	128	+ * @param count: Iterate count(once Iterate, compute baseM * baseN).
	129	+ * @retval None
	130	+ */
	131	+template <typename aType, typename bType, typename cType, typename biasType>
	132	+__aicore__ inline void MatmulLeakyKernel<aType, bType, cType, biasType>::CopyOut(uint32_t count)
	133	+{
	134	+ reluOutQueue_.DeQue<cType>();
	135	+ const uint32_t roundM = tiling.singleCoreM / tiling.baseM;
	136	+ const uint32_t roundN = tiling.singleCoreN / tiling.baseN;
	137	+ uint32_t startOffset = (count % roundM * tiling.baseM * tiling.N + count / roundM * tiling.baseN);
	138	+ AscendC::DataCopyParams copyParam = {(uint16_t)tiling.baseM, (uint16_t)(tiling.baseN * sizeof(cType) / AscendC::DEFAULT_C0_SIZE), 0,
	139	+ (uint16_t)((tiling.N - tiling.baseN) * sizeof(cType) / AscendC::DEFAULT_C0_SIZE)};
	140	+ DataCopy(cGlobal[startOffset], reluOutLocal, copyParam);
	141	+ reluOutQueue_.FreeTensor(reluOutLocal);
	142	+}
	143	+
	144	+/**
	145	+ * @brief Calculate the gm offset based on the blockidx.
	146	+ * @param blockIdx: Current Core blockidx.
	147	+ * @param tiling: Matmul tiling data.
	148	+ * @param offsetA: Gm offset of A matrix.
	149	+ * @param offsetB: Gm offset of B matrix.
	150	+ * @param offsetC: Gm offset of C matrix.
	151	+ * @param offsetBias: Gm offset of Bias matrix.
	152	+ * @retval None
	153	+ */
	154	+template <typename aType, typename bType, typename cType, typename biasType>
	155	+__aicore__ inline void
	156	+MatmulLeakyKernel<aType, bType, cType, biasType>::CalcOffset(int32_t blockIdx, const TCubeTiling &tiling,
	157	+ int32_t &offsetA, int32_t &offsetB, int32_t &offsetC,
	158	+ int32_t &offsetBias)
	159	+{
	160	+ auto mSingleBlocks = Ceiling(tiling.M, tiling.singleCoreM);
	161	+ auto mCoreIndx = blockIdx % mSingleBlocks;
	162	+ auto nCoreIndx = blockIdx / mSingleBlocks;
	163	+
	164	+ offsetA = mCoreIndx * tiling.Ka * tiling.singleCoreM;
	165	+ offsetB = nCoreIndx * tiling.singleCoreN;
	166	+ offsetC = mCoreIndx * tiling.N * tiling.singleCoreM + nCoreIndx * tiling.singleCoreN;
	167	+ offsetBias = nCoreIndx * tiling.singleCoreN;
	168	+}
	169	+
	170	+/**
	171	+ * @brief matmul_leakyrelu kernel function entry
	172	+ * @param a: A matrix gm addr.
	173	+ * @param b: B matrix gm addr.
	174	+ * @param bias: Bias gm addr.
	175	+ * @param c: Out gm addr.
	176	+ * @param workspace: Temporary gm space addr required by matmul calc.
	177	+ * @param tiling: Tiling data.
	178	+ * @retval None
	179	+ */
	180	+__global__ __mix__(1, 2) void matmul_leakyrelu_custom(GM_ADDR a, GM_ADDR b, GM_ADDR bias, GM_ADDR c,
	181	+ __kfc_workspace__ GM_ADDR workspace, AscendC::tiling::TCubeTiling tiling)
	182	+{
	183	+ AscendC::TPipe pipe;
	184	+ MatmulLeakyKernel<half, half, float, float> matmulLeakyKernel;
	185	+ matmulLeakyKernel.Init(a, b, bias, c, workspace, tiling, &pipe);
	186	+ REGIST_MATMUL_OBJ(&pipe, GetSysWorkSpacePtr(), matmulLeakyKernel.matmulObj, &matmulLeakyKernel.tiling); // Initialize the matmul object.
	187	+ matmulLeakyKernel.Process(&pipe);
	188	+}
	189	+
	190	+/**
	191	+ * @brief Generate matmul tiling.
	192	+ * @param ascendcPlatform: platform info.
	193	+ */
	194	+AscendC::tiling::TCubeTiling GenerateTiling(platform_ascendc::PlatformAscendC* ascendcPlatform)
	195	+{
	196	+ using TPosition = matmul_tiling::TPosition;
	197	+ using CubeFormat = matmul_tiling::CubeFormat;
	198	+ using DataType = matmul_tiling::DataType;
	199	+ int M = 1024;
	200	+ int N = 640;
	201	+ int K = 256;
	202	+
	203	+ TPosition leftPosition = TPosition::GM;
	204	+ CubeFormat leftFormat = CubeFormat::ND;
	205	+ DataType leftDtype = DataType::DT_FLOAT16;
	206	+ bool isTransA = false;
	207	+
	208	+ TPosition rightPosition = TPosition::GM;
	209	+ CubeFormat rightFormat = CubeFormat::ND;
	210	+ DataType rightDtype = DataType::DT_FLOAT16;
	211	+ bool isTransB = false;
	212	+
	213	+ TPosition resultPosition = TPosition::GM;
	214	+ CubeFormat resultFormat = CubeFormat::ND;
	215	+ DataType resultDtype = DataType::DT_FLOAT;
	216	+
	217	+ TPosition biasPosition = TPosition::GM;
	218	+ CubeFormat biasFormat = CubeFormat::ND;
	219	+ DataType biasDtype = DataType::DT_FLOAT;
	220	+ bool isBias = true;
	221	+
	222	+ int usedCoreNum = 2;
	223	+ int baseM = 256;
	224	+ int baseN = 128;
	225	+
	226	+ matmul_tiling::MultiCoreMatmulTiling tilingApi(*ascendcPlatform);
	227	+
	228	+ tilingApi.SetDim(usedCoreNum); // Set the number of cores that participate in multi-core computaion is 2.
	229	+ tilingApi.SetAType(leftPosition, leftFormat, leftDtype, isTransA);
	230	+ tilingApi.SetBType(rightPosition, rightFormat, rightDtype, isTransB);
	231	+ tilingApi.SetCType(resultPosition, resultFormat, resultDtype);
	232	+ tilingApi.SetBiasType(biasPosition, biasFormat, biasDtype);
	233	+
	234	+ tilingApi.SetOrgShape(M, N, K);
	235	+ tilingApi.SetShape(M, N, K);
	236	+ tilingApi.SetBias(isBias);
	237	+ tilingApi.SetTraverse(matmul_tiling::MatrixTraverse::FIRSTM); // Set the matmul travse is FIRSTM.
	238	+ tilingApi.SetFixSplit(baseM, baseN, -1); // Set the fixed baseM=128, baseN=256.
	239	+ tilingApi.SetBufferSpace(-1, -1, -1);
	240	+
	241	+ AscendC::tiling::TCubeTiling tilingData;
	242	+ int64_t res = tilingApi.GetTiling(tilingData); // Get matmul tiling data.
	243	+ if (res == -1) {
	244	+ std::cout << "gen tiling failed" << std::endl;
	245	+ }
	246	+ tilingData.stepM = 1; // Set the matmul tiling stepM=1.
	247	+ tilingData.stepN = 1; // Set the matmul tiling stepN=1.
	248	+ return tilingData;
	249	+}
	250	+
	251	+int32_t main(int32_t argc, char *argv[])
	252	+{
	253	+ const char *socVersion = "Ascend910B1";
	254	+ auto ascendcPlatform = platform_ascendc::PlatformAscendCManager::GetInstance(socVersion);
	255	+ size_t aFileSize = 262144 * sizeof(int16_t);
	256	+ size_t bFileSize = 163840 * sizeof(int16_t);
	257	+ size_t cFileSize = 655360 * sizeof(float);
	258	+ size_t biasFileSize = 640 * sizeof(float);
	259	+ size_t tilingFileSize = sizeof(TCubeTiling);
	260	+ size_t userWorkspaceSize = 0;
	261	+ size_t systemWorkspaceSize = static_cast<size_t>(ascendcPlatform->GetLibApiWorkSpaceSize());
	262	+ size_t workspaceSize = userWorkspaceSize + systemWorkspaceSize;
	263	+ auto tiling = GenerateTiling(ascendcPlatform);
	264	+ uint32_t numBlocks = 1;
	265	+
	266	+ aclInit(nullptr);
	267	+ int32_t deviceId = 0;
	268	+ aclrtSetDevice(deviceId);
	269	+ aclrtStream stream = nullptr;
	270	+ aclrtCreateStream(&stream);
	271	+
	272	+ uint8_t *inputAHost;
	273	+ uint8_t *inputADevice;
	274	+ aclrtMallocHost((void **)(&inputAHost), aFileSize);
	275	+ aclrtMalloc((void **)&inputADevice, aFileSize, ACL_MEM_MALLOC_HUGE_FIRST);
	276	+ ReadFile("./input/x1_gm.bin", aFileSize, inputAHost, aFileSize);
	277	+ aclrtMemcpy(inputADevice, aFileSize, inputAHost, aFileSize, ACL_MEMCPY_HOST_TO_DEVICE);
	278	+
	279	+ uint8_t *inputBHost;
	280	+ uint8_t *inputBDevice;
	281	+ aclrtMallocHost((void **)(&inputBHost), bFileSize);
	282	+ aclrtMalloc((void **)&inputBDevice, bFileSize, ACL_MEM_MALLOC_HUGE_FIRST);
	283	+ ReadFile("./input/x2_gm.bin", bFileSize, inputBHost, bFileSize);
	284	+ aclrtMemcpy(inputBDevice, bFileSize, inputBHost, bFileSize, ACL_MEMCPY_HOST_TO_DEVICE);
	285	+
	286	+ uint8_t *outputCHost;
	287	+ uint8_t *outputCDevice;
	288	+ aclrtMallocHost((void **)(&outputCHost), cFileSize);
	289	+ aclrtMalloc((void **)&outputCDevice, cFileSize, ACL_MEM_MALLOC_HUGE_FIRST);
	290	+
	291	+ uint8_t *inputBiasHost;
	292	+ uint8_t *inputBiasDevice;
	293	+ aclrtMallocHost((void **)(&inputBiasHost), biasFileSize);
	294	+ aclrtMalloc((void **)&inputBiasDevice, biasFileSize, ACL_MEM_MALLOC_HUGE_FIRST);
	295	+ ReadFile("./input/bias.bin", biasFileSize, inputBiasHost, biasFileSize);
	296	+ aclrtMemcpy(inputBiasDevice, biasFileSize, inputBiasHost, biasFileSize, ACL_MEMCPY_HOST_TO_DEVICE);
	297	+
	298	+ uint8_t *workspaceDevice;
	299	+ aclrtMalloc((void **)&workspaceDevice, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
	300	+
	301	+ matmul_leakyrelu_custom<<<numBlocks, nullptr, stream>>>(inputADevice, inputBDevice, inputBiasDevice, outputCDevice,
	302	+ workspaceDevice, tiling);
	303	+
	304	+ aclrtSynchronizeStream(stream);
	305	+
	306	+ aclrtFree(inputADevice);
	307	+ aclrtFreeHost(inputAHost);
	308	+ aclrtFree(inputBDevice);
	309	+ aclrtFreeHost(inputBHost);
	310	+ aclrtMemcpy(outputCHost, cFileSize, outputCDevice, cFileSize, ACL_MEMCPY_DEVICE_TO_HOST);
	311	+ WriteFile("./output/output.bin", outputCHost, cFileSize);
	312	+ aclrtFree(outputCDevice);
	313	+ aclrtFreeHost(outputCHost);
	314	+ aclrtFree(inputBiasDevice);
	315	+ aclrtFreeHost(inputBiasHost);
	316	+ aclrtFree(workspaceDevice);
	317	+
	318	+ aclrtDestroyStream(stream);
	319	+ aclrtResetDevice(deviceId);
	320	+ aclFinalize();
	321	+ return 0;
	322	+}

Aexamples/04_msobjdump/scripts/gen_data.py+38-0

@@ -0,0 +1,38 @@
	1	+#!/usr/bin/python3
	2	+# coding=utf-8
	3	+
	4	+# ----------------------------------------------------------------------------------------------------------
	5	+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
	6	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	7	+# 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,
	10	+# INCLUDING 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	+
	15	+import os
	16	+import numpy as np
	17	+
	18	+
	19	+def gen_golden_data():
	20	+ m = 1024
	21	+ n = 640
	22	+ k = 256
	23	+ input_a = np.random.randint(1, 10, [m, k]).astype(np.float16)
	24	+ input_b = np.random.randint(1, 10, [k, n]).astype(np.float16)
	25	+ input_bias = np.random.randint(1, 10, [n]).astype(np.float32)
	26	+ alpha = 0.001
	27	+ golden = (np.matmul(input_a.astype(np.float32), input_b.astype(np.float32)) + input_bias).astype(np.float32)
	28	+ golden = np.where(golden >= 0, golden, golden * alpha)
	29	+ os.makedirs("input", exist_ok=True)
	30	+ os.makedirs("output", exist_ok=True)
	31	+ input_a.tofile("./input/x1_gm.bin")
	32	+ input_b.tofile("./input/x2_gm.bin")
	33	+ input_bias.tofile("./input/bias.bin")
	34	+ golden.tofile("./output/golden.bin")
	35	+
	36	+
	37	+if __name__ == "__main__":
	38	+ gen_golden_data()

Aexamples/04_msobjdump/scripts/verify_result.py+59-0

@@ -0,0 +1,59 @@
	1	+#!/usr/bin/python3
	2	+# coding=utf-8
	3	+
	4	+# ----------------------------------------------------------------------------------------------------------
	5	+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
	6	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	7	+# 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,
	10	+# INCLUDING 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	+
	15	+import sys
	16	+import numpy as np
	17	+
	18	+
	19	+RELATIVE_TOL = 1e-6
	20	+ABSOLUTE_TOL = 1e-9
	21	+ERROR_TOL = 1e-4
	22	+
	23	+
	24	+def verify_result(output, golden):
	25	+ output = np.fromfile(output, dtype=np.float32).reshape(-1)
	26	+ golden = np.fromfile(golden, dtype=np.float32).reshape(-1)
	27	+ different_element_results = np.isclose(output,
	28	+ golden,
	29	+ rtol=RELATIVE_TOL,
	30	+ atol=ABSOLUTE_TOL,
	31	+ equal_nan=True)
	32	+ different_element_indexes = np.where(different_element_results == False)[0]
	33	+ for index in range(len(different_element_indexes)):
	34	+ real_index = different_element_indexes[index]
	35	+ golden_data = golden[real_index]
	36	+ output_data = output[real_index]
	37	+ print(
	38	+ "data index: %06d, expected: %-.9f, actual: %-.9f, rdiff: %-.6f" %
	39	+ (real_index, golden_data, output_data,
	40	+ abs(output_data - golden_data) / golden_data))
	41	+ if index == 100:
	42	+ break
	43	+ print("golden_data : ", golden)
	44	+ print("output : ", output)
	45	+ error_ratio = float(different_element_indexes.size) / golden.size
	46	+ print("error ratio: %.4f, tolerance: %.4f" % (error_ratio, ERROR_TOL))
	47	+ return error_ratio <= ERROR_TOL
	48	+
	49	+
	50	+if __name__ == '__main__':
	51	+ try:
	52	+ res = verify_result(sys.argv[1], sys.argv[2])
	53	+ if not res:
	54	+ raise ValueError("[ERROR] result error")
	55	+ else:
	56	+ print("test pass!")
	57	+ except Exception as e:
	58	+ print(e)
	59	+ sys.exit(1)

Mtests/py_ut/testcase/msobjdump/test_msobjdump.py+414-1

	@@ -13,7 +13,7 @@ import os
13	import sys	13	import sys
14	import shutil	14	import shutil
15	import unittest	15	import unittest
16	-import subprocess	16	+import contextlib
17		17
18	from unittest.mock import MagicMock, patch	18	from unittest.mock import MagicMock, patch
19	from io import StringIO	19	from io import StringIO
	@@ -288,6 +288,313 @@ class TestMsObjdump(unittest.TestCase):
288	self.assertTrue(True)	288	self.assertTrue(True)
289	self._clean_out_dir(out_dir)	289	self._clean_out_dir(out_dir)
290		290
		291	+ @patch('msobjdump.msobjdump_main.ObjDump._get_segment_content')
		292	+ @patch('msobjdump.utils.extract_aicore_binary_from_elf')
		293	+ @patch('msobjdump.utils.get_all_section_symbols_in_file')
		294	+ @patch('msobjdump.utils.get_section_headers_in_file')
		295	+ @patch('msobjdump.utils.get_symbols_in_file')
		296	+ def test_dump_elf_fusion_compile(self, mock_symbol, mock_section, mock_all, mock_run, mock_elf_content):
		297	+ mock_symbol.side_effect = ['test section list \n', 'test section list \n']
		298	+ mock_all.return_value = 'test all section and symbols \n'
		299	+ mock_section.side_effect = [
		300	+ '[26] .aicore_binary PROGBITS 00000000000ab500 07b500 0072c0 00 WA 0 0 256\n',
		301	+ '[15] .ascend.meta.gen_FFN_2000_mix_aiv NOTE 0000000000000000 06e108 000010 00 0 0 4\n',
		302	+ '[15] .ascend.meta.gen_FFN_2000_mix_aiv NOTE 0000000000000000 06e108 000010 00 0 0 4\n'
		303	+ ]
		304	+ mock_elf_content.return_value = b'\x01\x00\x04\x00\x04\x00\x00\x00\x03\x00\x04\x00\x01\x00\x01\x00'
		305	+
		306	+ def _mock_objcopy(input_file, output_file):
		307	+ with open(output_file, 'wb') as f:
		308	+ f.write(b'fusion_meta')
		309	+ result = MagicMock()
		310	+ result.returncode = 0
		311	+ result.stderr = ''
		312	+ return result
		313	+
		314	+ mock_run.side_effect = _mock_objcopy
		315	+ parse_mock = MagicMock()
		316	+ out_dir = self._make_out_dir('test_dump_elf_fusion_compile')
		317	+ parse_mock.out_dir = out_dir
		318	+ elf_file = os.path.join(out_dir, 'test_dump_elf_fusion_compile')
		319	+ parse_mock.dump_elf = elf_file
		320	+ with open(elf_file, 'a+') as f:
		321	+ f.write('test')
		322	+ parse_mock.extr_elf = None
		323	+ parse_mock.list_elf = None
		324	+ parse_mock.verbose = True
		325	+ try:
		326	+ msobjdump_main.run_obj_dump(parse_mock)
		327	+ except RuntimeError:
		328	+ self.assertTrue(False)
		329	+ else:
		330	+ self.assertTrue(mock_run.called)
		331	+ self.assertTrue(mock_all.called)
		332	+ self._clean_out_dir(out_dir)
		333	+
		334	+ @patch('msobjdump.msobjdump_main.ObjDump._get_segment_content')
		335	+ @patch('msobjdump.utils.extract_aicore_binary_from_elf')
		336	+ @patch('msobjdump.utils.get_all_section_symbols_in_file')
		337	+ @patch('msobjdump.utils.get_section_headers_in_file')
		338	+ @patch('msobjdump.utils.get_symbols_in_file')
		339	+ def test_list_elf_fusion_compile(self, mock_symbol, mock_section, mock_all, mock_run, mock_elf_content):
		340	+ mock_symbol.side_effect = ['test section list \n', 'test section list \n']
		341	+ mock_all.return_value = 'test all section and symbols \n'
		342	+ kernel_section = \
		343	+ '[23] .ascend.kernel.ascend910b1.ascendc_kernels_npu PROGBITS 0000 0002 018 00 WA 0 0 8\n'
		344	+ mock_section.side_effect = [
		345	+ '[26] .aicore_binary PROGBITS 00000000000ab500 07b500 0072c0 00 WA 0 0 256\n',
		346	+ kernel_section,
		347	+ kernel_section
		348	+ ]
		349	+ mock_elf_content.return_value = b'\x00\x00\x00\x00\x01\x00\x00\x00\x00\x00\x00\x00\x04\x00\x00\x00\x01\x00\x00\x00\x00\x00\x00\x00'
		350	+
		351	+ def _mock_objcopy(input_file, output_file):
		352	+ with open(output_file, 'wb') as f:
		353	+ f.write(b'fusion_kernel')
		354	+ result = MagicMock()
		355	+ result.returncode = 0
		356	+ result.stderr = ''
		357	+ return result
		358	+
		359	+ mock_run.side_effect = _mock_objcopy
		360	+ parse_mock = MagicMock()
		361	+ out_dir = self._make_out_dir('test_list_elf_fusion_compile')
		362	+ parse_mock.out_dir = out_dir
		363	+ elf_file = os.path.join(out_dir, 'test_list_elf_fusion_compile')
		364	+ parse_mock.list_elf = elf_file
		365	+ with open(elf_file, 'a+') as f:
		366	+ f.write('test')
		367	+ parse_mock.extr_elf = None
		368	+ parse_mock.dump_elf = None
		369	+ try:
		370	+ msobjdump_main.run_obj_dump(parse_mock)
		371	+ except RuntimeError:
		372	+ self.assertTrue(False)
		373	+ else:
		374	+ self.assertTrue(mock_run.called)
		375	+ self._clean_out_dir(out_dir)
		376	+
		377	+ @patch('msobjdump.msobjdump_main.ObjDump._get_segment_content')
		378	+ @patch('msobjdump.utils.extract_aicore_binary_from_elf')
		379	+ @patch('msobjdump.utils.get_all_section_symbols_in_file')
		380	+ @patch('msobjdump.utils.get_section_headers_in_file')
		381	+ @patch('msobjdump.utils.get_symbols_in_file')
		382	+ def test_extract_elf_fusion_compile(self, mock_symbol, mock_section, mock_all, mock_run, mock_elf_content):
		383	+ mock_symbol.side_effect = ['test section list \n', 'test section list \n']
		384	+ mock_all.return_value = 'test all section and symbols \n'
		385	+ kernel_section = \
		386	+ '[23] .ascend.kernel.ascend910b1.ascendc_kernels_npu PROGBITS 0000 0002 018 00 WA 0 0 8\n'
		387	+ mock_section.side_effect = [
		388	+ '[26] .aicore_binary PROGBITS 00000000000ab500 07b500 0072c0 00 WA 0 0 256\n',
		389	+ kernel_section,
		390	+ kernel_section
		391	+ ]
		392	+ mock_elf_content.return_value = b'\x00\x00\x00\x00\x01\x00\x00\x00\x00\x00\x00\x00\x04\x00\x00\x00\x01\x00\x00\x00\x00\x00\x00\x00'
		393	+
		394	+ def _mock_objcopy(input_file, output_file):
		395	+ with open(output_file, 'wb') as f:
		396	+ f.write(b'fusion_kernel')
		397	+ result = MagicMock()
		398	+ result.returncode = 0
		399	+ result.stderr = ''
		400	+ return result
		401	+
		402	+ mock_run.side_effect = _mock_objcopy
		403	+ parse_mock = MagicMock()
		404	+ out_dir = self._make_out_dir('test_extract_elf_fusion_compile')
		405	+ parse_mock.out_dir = out_dir
		406	+ elf_file = os.path.join(out_dir, 'test_extract_elf_fusion_compile')
		407	+ parse_mock.extr_elf = elf_file
		408	+ with open(elf_file, 'a+') as f:
		409	+ f.write('test')
		410	+ parse_mock.list_elf = None
		411	+ parse_mock.dump_elf = None
		412	+ try:
		413	+ msobjdump_main.run_obj_dump(parse_mock)
		414	+ except RuntimeError:
		415	+ self.assertTrue(False)
		416	+ else:
		417	+ self.assertTrue(mock_run.called)
		418	+ self.assertTrue(os.path.exists(os.path.join(out_dir, 'test_extract_elf_fusion_compile.aicore.o')))
		419	+ self._clean_out_dir(out_dir)
		420	+
		421	+ @patch('msobjdump.msobjdump_main.ObjDump._get_segment_content')
		422	+ @patch('msobjdump.utils.extract_aicore_binary_from_elf')
		423	+ @patch('msobjdump.utils.get_section_headers_in_file')
		424	+ @patch('msobjdump.utils.get_symbols_in_file')
		425	+ def test_list_elf_fusion_compile_meta_lists_extracted_binary(self, mock_symbol, mock_section, mock_run, mock_elf_content):
		426	+ mock_symbol.side_effect = ['test section list \n', 'test section list \n']
		427	+ mock_section.side_effect = [
		428	+ '[26] .aicore_binary PROGBITS 00000000000ab500 07b500 0072c0 00 WA 0 0 256\n',
		429	+ '[15] .ascend.meta.gen_FFN_2000_mix_aiv NOTE 0000000000000000 06e108 000010 00 0 0 4\n'
		430	+ ]
		431	+ mock_elf_content.return_value = b'\x01\x00\x04\x00\x04\x00\x00\x00\x03\x00\x04\x00\x01\x00\x01\x00'
		432	+
		433	+ def _mock_objcopy(input_file, output_file):
		434	+ with open(output_file, 'wb') as f:
		435	+ f.write(b'fusion_meta')
		436	+ result = MagicMock()
		437	+ result.returncode = 0
		438	+ result.stderr = ''
		439	+ return result
		440	+
		441	+ mock_run.side_effect = _mock_objcopy
		442	+ parse_mock = MagicMock()
		443	+ out_dir = self._make_out_dir('test_list_elf_fusion_compile_meta')
		444	+ parse_mock.out_dir = out_dir
		445	+ elf_file = os.path.join(out_dir, 'test_list_elf_fusion_compile_meta')
		446	+ parse_mock.list_elf = elf_file
		447	+ with open(elf_file, 'a+') as f:
		448	+ f.write('test')
		449	+ parse_mock.extr_elf = None
		450	+ parse_mock.dump_elf = None
		451	+
		452	+ captured_output = StringIO()
		453	+ with contextlib.redirect_stdout(captured_output):
		454	+ msobjdump_main.run_obj_dump(parse_mock)
		455	+ self.assertIn('ELF file 0: test_list_elf_fusion_compile_meta.aicore.o', captured_output.getvalue())
		456	+ self._clean_out_dir(out_dir)
		457	+
		458	+ @patch('msobjdump.msobjdump_main.ObjDump._get_segment_content')
		459	+ @patch('msobjdump.utils.extract_aicore_binary_from_elf')
		460	+ @patch('msobjdump.utils.get_section_headers_in_file')
		461	+ @patch('msobjdump.utils.get_symbols_in_file')
		462	+ def test_extract_elf_fusion_compile_meta_outputs_extracted_binary(self, mock_symbol, mock_section, mock_run, mock_elf_content):
		463	+ mock_symbol.side_effect = ['test section list \n', 'test section list \n']
		464	+ mock_section.side_effect = [
		465	+ '[26] .aicore_binary PROGBITS 00000000000ab500 07b500 0072c0 00 WA 0 0 256\n',
		466	+ '[15] .ascend.meta.gen_FFN_2000_mix_aiv NOTE 0000000000000000 06e108 000010 00 0 0 4\n'
		467	+ ]
		468	+ mock_elf_content.return_value = b'\x01\x00\x04\x00\x04\x00\x00\x00\x03\x00\x04\x00\x01\x00\x01\x00'
		469	+
		470	+ def _mock_objcopy(input_file, output_file):
		471	+ with open(output_file, 'wb') as f:
		472	+ f.write(b'fusion_meta')
		473	+ result = MagicMock()
		474	+ result.returncode = 0
		475	+ result.stderr = ''
		476	+ return result
		477	+
		478	+ mock_run.side_effect = _mock_objcopy
		479	+ parse_mock = MagicMock()
		480	+ out_dir = self._make_out_dir('test_extract_elf_fusion_compile_meta')
		481	+ parse_mock.out_dir = out_dir
		482	+ elf_file = os.path.join(out_dir, 'test_extract_elf_fusion_compile_meta')
		483	+ parse_mock.extr_elf = elf_file
		484	+ with open(elf_file, 'a+') as f:
		485	+ f.write('test')
		486	+ parse_mock.list_elf = None
		487	+ parse_mock.dump_elf = None
		488	+
		489	+ msobjdump_main.run_obj_dump(parse_mock)
		490	+ self.assertTrue(os.path.exists(os.path.join(out_dir, 'test_extract_elf_fusion_compile_meta.aicore.o')))
		491	+ self._clean_out_dir(out_dir)
		492	+
		493	+ @patch('msobjdump.utils.extract_aicore_binary_from_elf')
		494	+ @patch('msobjdump.utils.get_section_headers_in_file')
		495	+ @patch('msobjdump.utils.get_symbols_in_file')
		496	+ def test_fusion_compile_objcopy_failed(self, mock_symbol, mock_section, mock_run):
		497	+ mock_symbol.return_value = 'test section list \n'
		498	+ mock_section.return_value = \
		499	+ '[26] .aicore_binary PROGBITS 00000000000ab500 07b500 0072c0 00 WA 0 0 256\n'
		500	+ result = MagicMock()
		501	+ result.returncode = 1
		502	+ result.stderr = 'mock objcopy failed'
		503	+ mock_run.return_value = result
		504	+
		505	+ parse_mock = MagicMock()
		506	+ out_dir = self._make_out_dir('test_fusion_compile_objcopy_failed')
		507	+ parse_mock.out_dir = out_dir
		508	+ elf_file = os.path.join(out_dir, 'test_fusion_compile_objcopy_failed')
		509	+ parse_mock.dump_elf = elf_file
		510	+ with open(elf_file, 'a+') as f:
		511	+ f.write('test')
		512	+ parse_mock.extr_elf = None
		513	+ parse_mock.list_elf = None
		514	+
		515	+ with self.assertRaises(RuntimeError):
		516	+ msobjdump_main.run_obj_dump(parse_mock)
		517	+ self._clean_out_dir(out_dir)
		518	+
		519	+ @patch('msobjdump.utils.extract_aicore_binary_from_elf')
		520	+ @patch('msobjdump.utils.get_section_headers_in_file')
		521	+ @patch('msobjdump.utils.get_symbols_in_file')
		522	+ def test_fusion_compile_unsupported_after_extract(self, mock_symbol, mock_section, mock_run):
		523	+ mock_symbol.side_effect = ['test section list \n', 'test section list \n']
		524	+ mock_section.side_effect = [
		525	+ '[26] .aicore_binary PROGBITS 00000000000ab500 07b500 0072c0 00 WA 0 0 256\n',
		526	+ '[26] .aicore_binary PROGBITS 00000000000ab500 07b500 0072c0 00 WA 0 0 256\n'
		527	+ ]
		528	+
		529	+ def _mock_objcopy(input_file, output_file):
		530	+ with open(output_file, 'wb') as f:
		531	+ f.write(b'fusion_unknown')
		532	+ result = MagicMock()
		533	+ result.returncode = 0
		534	+ result.stderr = ''
		535	+ return result
		536	+
		537	+ mock_run.side_effect = _mock_objcopy
		538	+ parse_mock = MagicMock()
		539	+ out_dir = self._make_out_dir('test_fusion_compile_unsupported_after_extract')
		540	+ parse_mock.out_dir = out_dir
		541	+ parse_mock.verbose = False
		542	+ elf_file = os.path.join(out_dir, 'test_fusion_compile_unsupported_after_extract')
		543	+ parse_mock.dump_elf = elf_file
		544	+ with open(elf_file, 'a+') as f:
		545	+ f.write('test')
		546	+ parse_mock.extr_elf = None
		547	+ parse_mock.list_elf = None
		548	+
		549	+ msobjdump_main.run_obj_dump(parse_mock)
		550	+ self._clean_out_dir(out_dir)
		551	+
		552	+ def test_extract_aicore_binary_edge_cases(self):
		553	+ with self.subTest("remove old tmp and extract success"):
		554	+ objdump = msobjdump_main.ObjDump.__new__(msobjdump_main.ObjDump)
		555	+ objdump.tmp_dir = self._make_out_dir('test_extract_aicore_binary_edge_success')
		556	+ objdump.obj = os.path.join(objdump.tmp_dir, 'input_bin')
		557	+ with patch('msobjdump.msobjdump_main.os.path.exists', side_effect=[True, True]) as mock_exists, \
		558	+ patch('msobjdump.msobjdump_main.os.path.getsize', return_value=8), \
		559	+ patch('msobjdump.msobjdump_main.os.remove') as mock_remove, \
		560	+ patch('msobjdump.msobjdump_main.utils.extract_aicore_binary_from_elf') as mock_run:
		561	+ result = MagicMock()
		562	+ result.returncode = 0
		563	+ result.stderr = ''
		564	+ mock_run.return_value = result
		565	+ tmp_file = objdump._extract_aicore_binary()
		566	+ self.assertTrue(mock_exists.called)
		567	+ self.assertTrue(mock_remove.called)
		568	+ self.assertEqual(tmp_file, os.path.join(objdump.tmp_dir, 'fusion_aicore_binary.aicore.o'))
		569	+ shutil.rmtree(objdump.tmp_dir, ignore_errors=True)
		570	+
		571	+ with self.subTest("llvm-objcopy not found"):
		572	+ objdump = msobjdump_main.ObjDump.__new__(msobjdump_main.ObjDump)
		573	+ objdump.tmp_dir = self._make_out_dir('test_extract_aicore_binary_edge_not_found')
		574	+ objdump.obj = os.path.join(objdump.tmp_dir, 'input_bin')
		575	+ with patch('msobjdump.msobjdump_main.os.path.exists', return_value=False), \
		576	+ patch('msobjdump.msobjdump_main.utils.extract_aicore_binary_from_elf',
		577	+ side_effect=FileNotFoundError('llvm-objcopy not found')):
		578	+ with self.assertRaises(RuntimeError) as ctx:
		579	+ objdump._extract_aicore_binary()
		580	+ self.assertIn('llvm-objcopy is not available', str(ctx.exception))
		581	+ shutil.rmtree(objdump.tmp_dir, ignore_errors=True)
		582	+
		583	+ with self.subTest("extract output missing or empty"):
		584	+ objdump = msobjdump_main.ObjDump.__new__(msobjdump_main.ObjDump)
		585	+ objdump.tmp_dir = self._make_out_dir('test_extract_aicore_binary_edge_empty')
		586	+ objdump.obj = os.path.join(objdump.tmp_dir, 'input_bin')
		587	+ with patch('msobjdump.msobjdump_main.os.path.exists', side_effect=[False, False]), \
		588	+ patch('msobjdump.msobjdump_main.utils.extract_aicore_binary_from_elf') as mock_run:
		589	+ result = MagicMock()
		590	+ result.returncode = 0
		591	+ result.stderr = ''
		592	+ mock_run.return_value = result
		593	+ with self.assertRaises(RuntimeError) as ctx:
		594	+ objdump._extract_aicore_binary()
		595	+ self.assertIn('file is empty or missing', str(ctx.exception))
		596	+ shutil.rmtree(objdump.tmp_dir, ignore_errors=True)
		597	+
291	def test_show_ascend_meta_tlv_block_num(self):	598	def test_show_ascend_meta_tlv_block_num(self):
292	content = b'\xff\xff\xff\xff'	599	content = b'\xff\xff\xff\xff'
293	t = 15	600	t = 15
	@@ -304,6 +611,112 @@ class TestMsObjdump(unittest.TestCase):
304		611
305	self.assertEqual(output, "BLOCK_NUM: 0xFFFFFFFF")	612	self.assertEqual(output, "BLOCK_NUM: 0xFFFFFFFF")
306		613
		614	+ def test_show_ascend_meta_op_tlv_dedup_for_aicore_binary(self):
		615	+ objdump = msobjdump_main.ObjDump.__new__(msobjdump_main.ObjDump)
		616	+ objdump.obj_type = msobjdump_main.ObjType.TYPE_AICORE_BINARY
		617	+ objdump._aicore_binary_meta_printed = set()
		618	+ version_content = b'\x01\x00\x00\x00'
		619	+ runtime_content = b'\x03\x00\x00\x00'
		620	+
		621	+ captured_output = StringIO()
		622	+ sys.stdout = captured_output
		623	+
		624	+ objdump._show_ascend_meta_op_tlv(
		625	+ runtime_content,
		626	+ msobjdump_main.B_TYPE_RUNTIME_IMPLICIT_INFO,
		627	+ 4,
		628	+ 0
		629	+ )
		630	+ objdump._show_ascend_meta_op_tlv(version_content, msobjdump_main.B_TYPE_VERSION, 4, 0)
		631	+ objdump._show_ascend_meta_op_tlv(version_content, msobjdump_main.B_TYPE_VERSION, 4, 0)
		632	+
		633	+ sys.stdout = sys.__stdout__
		634	+ output = captured_output.getvalue().strip().splitlines()
		635	+
		636	+ self.assertEqual(output, ["RUNTIME_IMPLICIT_INFO: L2Cache Hint Flag", "VERSION: 1"])
		637	+
		638	+ def test_show_ascend_meta_op_tlv_other_types(self):
		639	+ objdump = msobjdump_main.ObjDump.__new__(msobjdump_main.ObjDump)
		640	+ objdump.obj_type = msobjdump_main.ObjType.TYPE_ASCEND_META
		641	+ objdump._aicore_binary_meta_printed = set()
		642	+
		643	+ captured_output = StringIO()
		644	+ sys.stdout = captured_output
		645	+
		646	+ objdump._show_ascend_meta_op_tlv(
		647	+ b'\x08\x00\x00\x00\x02\x00\x00\x00',
		648	+ msobjdump_main.B_TYPE_DEBUG,
		649	+ 8,
		650	+ 0
		651	+ )
		652	+ objdump._show_ascend_meta_op_tlv(
		653	+ b'\x00\x00\x03\x00',
		654	+ msobjdump_main.B_TYPE_DYNAMIC_PARAM,
		655	+ 4,
		656	+ 0
		657	+ )
		658	+ objdump._show_ascend_meta_op_tlv(
		659	+ b'\x01\x00\x02\x00',
		660	+ msobjdump_main.B_TYPE_OPTIONAL_PARAM,
		661	+ 4,
		662	+ 0
		663	+ )
		664	+ objdump._show_ascend_meta_op_tlv(
		665	+ b'\x03\x00\x00\x00',
		666	+ msobjdump_main.B_TYPE_RUNTIME_IMPLICIT_INFO,
		667	+ 4,
		668	+ 0
		669	+ )
		670	+
		671	+ sys.stdout = sys.__stdout__
		672	+ output = captured_output.getvalue().strip().splitlines()
		673	+
		674	+ self.assertEqual(
		675	+ output,
		676	+ [
		677	+ "DEBUG: debugBufSize=8, debugOptions=2",
		678	+ "DYNAMIC_PARAM: dynamicParamMode=3",
		679	+ "OPTIONAL_PARAM: optionalInputMode=1, optionalOutputMode=2",
		680	+ "RUNTIME_IMPLICIT_INFO: L2Cache Hint Flag"
		681	+ ]
		682	+ )
		683	+
		684	+ def test_utils_helper_functions(self):
		685	+ self.assertEqual(utils.split_str_with_space(" foo\t bar baz "), ["foo", "bar", "baz"])
		686	+ self.assertEqual(utils.get_str_between("_binary_demo_o_start", "_binary_", "_start"), "demo_o")
		687	+ self.assertEqual(utils.get_str_between("abc", "_binary_", "_start"), "")
		688	+ self.assertTrue(utils.is_prefix_substring("Ascend910", ["config", "ascend"]))
		689	+ self.assertFalse(utils.is_prefix_substring("kernel", ["config", "ascend"]))
		690	+
		691	+ @patch('msobjdump.utils.subprocess.run')
		692	+ @patch('msobjdump.utils.shutil.copy')
		693	+ @patch('msobjdump.utils.os.makedirs')
		694	+ @patch('msobjdump.utils.os.path.exists')
		695	+ def test_utils_file_and_command_wrappers(self, mock_exists, mock_makedirs, mock_copy, mock_run):
		696	+ mock_exists.side_effect = [True, False]
		697	+ utils.copy_file_src_exist("/tmp/src.txt", "/tmp/out/dest.txt")
		698	+ mock_makedirs.assert_called_once_with("/tmp/out")
		699	+ mock_copy.assert_called_once_with("/tmp/src.txt", "/tmp/out/dest.txt")
		700	+
		701	+ mock_exists.reset_mock()
		702	+ mock_makedirs.reset_mock()
		703	+ mock_copy.reset_mock()
		704	+ mock_exists.side_effect = None
		705	+ mock_exists.return_value = False
		706	+ utils.copy_file_src_exist("/tmp/missing.txt", "/tmp/out/dest.txt")
		707	+ mock_makedirs.assert_not_called()
		708	+ mock_copy.assert_not_called()
		709	+
		710	+ mock_run.return_value = MagicMock(stdout="test output")
		711	+ self.assertEqual(utils.get_section_headers_in_file("test.o"), "test output")
		712	+ self.assertEqual(utils.get_symbols_in_file("test.o"), "test output")
		713	+ self.assertEqual(utils.get_all_section_symbols_in_file("test.o"), "test output")
		714	+
		715	+ mock_result = MagicMock()
		716	+ mock_run.return_value = mock_result
		717	+ self.assertIs(utils.get_o_file_from_a_file("libtest.a", "demo.o"), mock_result)
		718	+ self.assertIs(utils.extract_aicore_binary_from_elf("demo", "demo.aicore.o"), mock_result)
		719	+
307		720
308	if __name__ == "__main__":	721	if __name__ == "__main__":
309	unittest.main()	722	unittest.main()

Mutils/msobjdump/msobjdump/msobjdump_main.py+129-43

	@@ -29,6 +29,7 @@ from msobjdump import utils
29	KEY_ASCEND_META = '.ascend.meta.'	29	KEY_ASCEND_META = '.ascend.meta.'
30	KEY_ASCEND_META_OP = '.ascend.meta'	30	KEY_ASCEND_META_OP = '.ascend.meta'
31	KEY_ASCEND_KERNEL = '.ascend.kernel.'	31	KEY_ASCEND_KERNEL = '.ascend.kernel.'
		32	+KEY_AICORE_BINARY = '.aicore_binary'
32	KEY_O_JSON = '_binary_'	33	KEY_O_JSON = '_binary_'
33	KEY_A_FILE = '.a'	34	KEY_A_FILE = '.a'
34	KEY_O_FILE = '.cpp.o'	35	KEY_O_FILE = '.cpp.o'
	@@ -42,21 +43,36 @@ ACLNN_BINARY = ''
42	############### ascend kernel define ####################	43	############### ascend kernel define ####################
43	KERNEL_TYPE_MAP = {'0': 'mix', '1': 'aiv', '2': 'aic'}	44	KERNEL_TYPE_MAP = {'0': 'mix', '1': 'aiv', '2': 'aic'}
44		45
45	-############### ascend meta define ####################	46	+############### Function Meta Type ####################
46	-TYPE_KTYPE = 1	47	+F_TYPE_KTYPE = 1
47	-TYPE_CROSS_CORE_SYNC = 2	48	+F_TYPE_CROSS_CORE_SYNC = 2
48	-TYPE_MIX_TASK_RATION = 3	49	+F_TYPE_MIX_TASK_RATION = 3
49	-TYPE_DETERMINISTIC_INFO = 13	50	+F_TYPE_DETERMINISTIC_INFO = 13
50	-TYPE_FUNCTION_ENTRY = 14	51	+F_TYPE_FUNCTION_ENTRY = 14
51	-TYPE_BLOCK_NUM = 15	52	+F_TYPE_BLOCK_NUM = 15
52		53
53	F_TYPE_MAP = {	54	F_TYPE_MAP = {
54	- TYPE_KTYPE: "KERNEL_TYPE",	55	+ F_TYPE_KTYPE: "KERNEL_TYPE",
55	- TYPE_CROSS_CORE_SYNC: "CROSS_CORE_SYNC",	56	+ F_TYPE_CROSS_CORE_SYNC: "CROSS_CORE_SYNC",
56	- TYPE_MIX_TASK_RATION: "MIX_TASK_RATION",	57	+ F_TYPE_MIX_TASK_RATION: "MIX_TASK_RATION",
57	- TYPE_DETERMINISTIC_INFO: "DETERMINISTIC_INFO",	58	+ F_TYPE_DETERMINISTIC_INFO: "DETERMINISTIC_INFO",
58	- TYPE_FUNCTION_ENTRY: "FUNCTION_ENTRY",	59	+ F_TYPE_FUNCTION_ENTRY: "FUNCTION_ENTRY",
59	- TYPE_BLOCK_NUM: "BLOCK_NUM"	60	+ F_TYPE_BLOCK_NUM: "BLOCK_NUM"
		61	+}
		62	+
		63	+############### Binary Meta Type ####################
		64	+B_TYPE_VERSION = 0
		65	+B_TYPE_DEBUG = 1
		66	+B_TYPE_DYNAMIC_PARAM = 2
		67	+B_TYPE_OPTIONAL_PARAM = 3
		68	+B_TYPE_RUNTIME_IMPLICIT_INFO = 4
		69	+
		70	+B_TYPE_MAP = {
		71	+ B_TYPE_VERSION: "VERSION",
		72	+ B_TYPE_DEBUG: "DEBUG",
		73	+ B_TYPE_DYNAMIC_PARAM: "DYNAMIC_PARAM",
		74	+ B_TYPE_OPTIONAL_PARAM: "OPTIONAL_PARAM",
		75	+ B_TYPE_RUNTIME_IMPLICIT_INFO: "RUNTIME_IMPLICIT_INFO"
60	}	76	}
61		77
62	K_TYPE_MAP = {	78	K_TYPE_MAP = {
	@@ -69,12 +85,20 @@ K_TYPE_MAP = {
69	"7": "AIV_ROLLBACK"	85	"7": "AIV_ROLLBACK"
70	}	86	}
71	C_TYPE_MAP = {"0":"NO_USE_SYNC", "1": "USE_SYNC"}	87	C_TYPE_MAP = {"0":"NO_USE_SYNC", "1": "USE_SYNC"}
		88	+RUNTIME_IMPLICIT_INFO_MAP = {
		89	+ 1: "SIMD Printf Flag",
		90	+ 2: "Hardware Sync Flag",
		91	+ 3: "L2Cache Hint Flag",
		92	+ 4: "SIMT Printf Flag",
		93	+ 5: "SIMD Assert Flag"
		94	+}
72		95
73		96
74	class ObjType(Enum):	97	class ObjType(Enum):
75	TYPE_ASCEND_KERNEL = 0	98	TYPE_ASCEND_KERNEL = 0
76	TYPE_BINARY_O_JSON = 1	99	TYPE_BINARY_O_JSON = 1
77	TYPE_ASCEND_META = 2	100	TYPE_ASCEND_META = 2
		101	+ TYPE_AICORE_BINARY = 3
78		102
79		103
80	class ParseObjMode(Enum):	104	class ParseObjMode(Enum):
	@@ -107,53 +131,37 @@ class ObjDump:
107	'''	131	'''
108	def __init__(self, args):	132	def __init__(self, args):
109	self.obj = None # 待解析解压文件	133	self.obj = None # 待解析解压文件
		134	+ self.src_obj = None # 用户原始输入文件
110	self.parse_obj_mode = None # ParseObjMode 解析解压类型	135	self.parse_obj_mode = None # ParseObjMode 解析解压类型
111	self.obj_type = None # ObjType obj文件场景分类	136	self.obj_type = None # ObjType obj文件场景分类
112	self.tmp_dir = None # 存放临时文件结束后目录会删除	137	self.tmp_dir = None # 存放临时文件结束后目录会删除
113	self.out_dir = args.out_dir # 落盘文件目录，用户未设置时当前路径	138	self.out_dir = args.out_dir # 落盘文件目录，用户未设置时当前路径
		139	+ self._aicore_binary_meta_printed = set()
114	self._set_out_dir()	140	self._set_out_dir()
115	# preprocess args	141	# preprocess args
116	self._set_parse_obj_and_mode(args)	142	self._set_parse_obj_and_mode(args)
117		143
118	-
119	@staticmethod	144	@staticmethod
120	- def _show_ascend_meta_tlv(content: bytes, t: int, l: int, index: int):	145	+ def _show_ascend_meta_tlv(content: bytes, t: int, l: int, index: int):
121	- if t == TYPE_MIX_TASK_RATION:	146	+ if t == F_TYPE_MIX_TASK_RATION:
122	v1, v2 = struct.unpack("2H", content[index:index + 4])	147	v1, v2 = struct.unpack("2H", content[index:index + 4])
123	print(f"{F_TYPE_MAP.get(t)}: [{v1}:{v2}]")	148	print(f"{F_TYPE_MAP.get(t)}: [{v1}:{v2}]")
124	- elif t == TYPE_CROSS_CORE_SYNC :	149	+ elif t == F_TYPE_CROSS_CORE_SYNC:
125	v, = struct.unpack("I", content[index:index + 4])	150	v, = struct.unpack("I", content[index:index + 4])
126	print(f"{F_TYPE_MAP.get(t)}: {C_TYPE_MAP.get(str(v))}")	151	print(f"{F_TYPE_MAP.get(t)}: {C_TYPE_MAP.get(str(v))}")
127	- elif t == TYPE_KTYPE:	152	+ elif t == F_TYPE_KTYPE:
128	v, = struct.unpack("I", content[index:index + 4])	153	v, = struct.unpack("I", content[index:index + 4])
129	print(f"{F_TYPE_MAP.get(t)}: {K_TYPE_MAP.get(str(v))}")	154	print(f"{F_TYPE_MAP.get(t)}: {K_TYPE_MAP.get(str(v))}")
130	- elif t == TYPE_DETERMINISTIC_INFO:	155	+ elif t == F_TYPE_DETERMINISTIC_INFO:
131	v, = struct.unpack("I", content[index:index + 4])	156	v, = struct.unpack("I", content[index:index + 4])
132	print(f"{F_TYPE_MAP.get(t)}: {v}")	157	print(f"{F_TYPE_MAP.get(t)}: {v}")
133	- elif t == TYPE_FUNCTION_ENTRY:	158	+ elif t == F_TYPE_FUNCTION_ENTRY:
134	v, = struct.unpack("<Q", content[index + 4:index + 12])	159	v, = struct.unpack("<Q", content[index + 4:index + 12])
135	print(f"{F_TYPE_MAP.get(t)}: {v}")	160	print(f"{F_TYPE_MAP.get(t)}: {v}")
136	- elif t == TYPE_BLOCK_NUM:	161	+ elif t == F_TYPE_BLOCK_NUM:
137	v = "0xFFFFFFFF"	162	v = "0xFFFFFFFF"
138	print(f"{F_TYPE_MAP.get(t)}: {v}")	163	print(f"{F_TYPE_MAP.get(t)}: {v}")
139		164
140	-
141	- @staticmethod
142	- def _show_ascend_meta_op_tlv(content: bytes, t: int, l: int, index: int):
143	- if t == 0:
144	- v, = struct.unpack("I", content[index:index + 4])
145	- print(f"VERSION: {v}")
146	- if t == 1:
147	- debugBufSize, debugOptions = struct.unpack("II", content[index:index + 8])
148	- print(f"DEBUG: debugBufSize={debugBufSize}, debugOptions={debugOptions}")
149	- elif t == 2:
150	- dynamicParamMode, = struct.unpack("H", content[index + 2:index + 4])
151	- print(f"DYNAMIC_PARAM: dynamicParamMode={dynamicParamMode}")
152	- elif t == 3:
153	- optionalInputMode, optionalOutputMode = struct.unpack("HH", content[index:index + 4])
154	- print(f"OPTIONAL_PARAM: optionalInputMode={optionalInputMode}, optionalOutputMode={optionalOutputMode}")
155	-
156	-
157	@staticmethod	165	@staticmethod
158	def _unpack_buff_content_by_type(content: bytes, start_idx: int, read_len: int, type_str: str):	166	def _unpack_buff_content_by_type(content: bytes, start_idx: int, read_len: int, type_str: str):
159	try:	167	try:
	@@ -170,6 +178,38 @@ class ObjDump:
170	self._clean()	178	self._clean()
171	return	179	return
172		180
		181	+ def _show_ascend_meta_op_tlv(self, content: bytes, t: int, l: int, index: int):
		182	+ if t == B_TYPE_VERSION:
		183	+ v, = struct.unpack("I", content[index:index + 4])
		184	+ output = f"{B_TYPE_MAP.get(t)}: {v}"
		185	+ self._print_ascend_meta_op_tlv(output)
		186	+ if t == B_TYPE_DEBUG:
		187	+ debugBufSize, debugOptions = struct.unpack("II", content[index:index + 8])
		188	+ output = f"{B_TYPE_MAP.get(t)}: debugBufSize={debugBufSize}, debugOptions={debugOptions}"
		189	+ self._print_ascend_meta_op_tlv(output)
		190	+ elif t == B_TYPE_DYNAMIC_PARAM:
		191	+ dynamicParamMode, = struct.unpack("H", content[index + 2:index + 4])
		192	+ output = f"{B_TYPE_MAP.get(t)}: dynamicParamMode={dynamicParamMode}"
		193	+ self._print_ascend_meta_op_tlv(output)
		194	+ elif t == B_TYPE_OPTIONAL_PARAM:
		195	+ optionalInputMode, optionalOutputMode = struct.unpack("HH", content[index:index + 4])
		196	+ output = (
		197	+ f"{B_TYPE_MAP.get(t)}: optionalInputMode={optionalInputMode}, "
		198	+ f"optionalOutputMode={optionalOutputMode}"
		199	+ )
		200	+ self._print_ascend_meta_op_tlv(output)
		201	+ elif t == B_TYPE_RUNTIME_IMPLICIT_INFO:
		202	+ v, = struct.unpack("I", content[index:index + 4])
		203	+ output = f"{B_TYPE_MAP.get(t)}: {RUNTIME_IMPLICIT_INFO_MAP.get(v, v)}"
		204	+ self._print_ascend_meta_op_tlv(output)
		205	+
		206	+ def _print_ascend_meta_op_tlv(self, output: str):
		207	+ if self.obj_type == ObjType.TYPE_AICORE_BINARY:
		208	+ if output in self._aicore_binary_meta_printed:
		209	+ return
		210	+ self._aicore_binary_meta_printed.add(output)
		211	+ print(output)
		212	+
173	def _set_out_dir(self):	213	def _set_out_dir(self):
174	if not self.out_dir:	214	if not self.out_dir:
175	self.out_dir = os.getcwd()	215	self.out_dir = os.getcwd()
	@@ -206,7 +246,16 @@ class ObjDump:
206	self.obj = args.list_elf	246	self.obj = args.list_elf
207	self.parse_obj_mode = ParseObjMode.MODE_LIST_ELF	247	self.parse_obj_mode = ParseObjMode.MODE_LIST_ELF
208		248
209	- # parse obj type	249	+ self.src_obj = self.obj
		250	+ self._detect_obj_type()
		251	+ if self.obj_type is not None and self.obj_type != ObjType.TYPE_AICORE_BINARY:
		252	+ return
		253	+
		254	+ if self.obj_type == ObjType.TYPE_AICORE_BINARY:
		255	+ self.obj = self._extract_aicore_binary()
		256	+
		257	+ def _detect_obj_type(self):
		258	+ self.obj_type = None
210	output = utils.get_symbols_in_file(self.obj)	259	output = utils.get_symbols_in_file(self.obj)
211		260
212	if "_o_start" in output or "_json_start" in output:	261	if "_o_start" in output or "_json_start" in output:
	@@ -217,6 +266,27 @@ class ObjDump:
217	self.obj_type = ObjType.TYPE_ASCEND_META	266	self.obj_type = ObjType.TYPE_ASCEND_META
218	if KEY_ASCEND_KERNEL in output:	267	if KEY_ASCEND_KERNEL in output:
219	self.obj_type = ObjType.TYPE_ASCEND_KERNEL	268	self.obj_type = ObjType.TYPE_ASCEND_KERNEL
		269	+ if KEY_AICORE_BINARY in output:
		270	+ self.obj_type = ObjType.TYPE_AICORE_BINARY
		271	+
		272	+ def _extract_aicore_binary(self) -> str:
		273	+ src_obj = getattr(self, 'src_obj', None)
		274	+ src_name = os.path.basename(src_obj) if src_obj else 'fusion_aicore_binary'
		275	+ tmp_file = os.path.join(self.tmp_dir, f'{src_name}.aicore.o')
		276	+ if os.path.exists(tmp_file):
		277	+ os.remove(tmp_file)
		278	+ try:
		279	+ result = utils.extract_aicore_binary_from_elf(self.obj, tmp_file)
		280	+ except FileNotFoundError as e:
		281	+ raise RuntimeError('[ERROR]: llvm-objcopy is not available, cannot extract .aicore_binary.') from e
		282	+
		283	+ if result.returncode != 0:
		284	+ err_msg = result.stderr.strip() if result.stderr else 'unknown error'
		285	+ raise RuntimeError(f'[ERROR]: Extract .aicore_binary failed: {err_msg}')
		286	+
		287	+ if not os.path.exists(tmp_file) or os.path.getsize(tmp_file) == 0:
		288	+ raise RuntimeError('[ERROR]: Extracted .aicore_binary file is empty or missing.')
		289	+ return tmp_file
220		290
221	def _parse_process(self):	291	def _parse_process(self):
222	if self.parse_obj_mode == ParseObjMode.MODE_DUMP_ELF or self.parse_obj_mode == ParseObjMode.MODE_VERBOSE:	292	if self.parse_obj_mode == ParseObjMode.MODE_DUMP_ELF or self.parse_obj_mode == ParseObjMode.MODE_VERBOSE:
	@@ -246,6 +316,11 @@ class ObjDump:
246	self._show_elf_ascend_kernel_obj(kernel_obj_infos)	316	self._show_elf_ascend_kernel_obj(kernel_obj_infos)
247	elif self.obj_type == ObjType.TYPE_ASCEND_META:	317	elif self.obj_type == ObjType.TYPE_ASCEND_META:
248	self._show_elf_ascend_meta_obj()	318	self._show_elf_ascend_meta_obj()
		319	+ elif self.obj_type == ObjType.TYPE_AICORE_BINARY:
		320	+ self._show_elf_ascend_meta_obj()
		321	+ if self.parse_obj_mode == ParseObjMode.MODE_VERBOSE:
		322	+ print(f'====== [elf header infos] ======')
		323	+ print(utils.get_all_section_symbols_in_file(self.obj))
249	else:	324	else:
250	print(f'The kernel meta information cannot be found.')	325	print(f'The kernel meta information cannot be found.')
251		326
	@@ -260,6 +335,8 @@ class ObjDump:
260	self._move_file_to_outdir_ascend_kernel(kernel_obj_infos)	335	self._move_file_to_outdir_ascend_kernel(kernel_obj_infos)
261	elif self.obj_type == ObjType.TYPE_ASCEND_META:	336	elif self.obj_type == ObjType.TYPE_ASCEND_META:
262	print('[WARNING]: nothing to extra in single op elf file')	337	print('[WARNING]: nothing to extra in single op elf file')
		338	+ elif self.obj_type == ObjType.TYPE_AICORE_BINARY:
		339	+ self._move_extracted_aicore_binary_to_outdir()
263		340
264	def _list_elf(self):	341	def _list_elf(self):
265	if self.obj_type == ObjType.TYPE_BINARY_O_JSON:	342	if self.obj_type == ObjType.TYPE_BINARY_O_JSON:
	@@ -270,6 +347,14 @@ class ObjDump:
270	self._parse_elf_ascend_kernel_by_type(ascend_kernel_dict, 'list')	347	self._parse_elf_ascend_kernel_by_type(ascend_kernel_dict, 'list')
271	elif self.obj_type == ObjType.TYPE_ASCEND_META:	348	elif self.obj_type == ObjType.TYPE_ASCEND_META:
272	print('[WARNING]: nothing to list in single op elf file')	349	print('[WARNING]: nothing to list in single op elf file')
		350	+ elif self.obj_type == ObjType.TYPE_AICORE_BINARY:
		351	+ self._list_extracted_aicore_binary()
		352	+
		353	+ def _list_extracted_aicore_binary(self):
		354	+ print(f'ELF file 0: {os.path.basename(self.obj)}')
		355	+
		356	+ def _move_extracted_aicore_binary_to_outdir(self):
		357	+ utils.copy_file_src_exist(self.obj, os.path.join(self.out_dir, os.path.basename(self.obj)))
273		358
274	def _copy_file_o_json_to_out_dir(self, src_file: str, json_file: str, copy_files: list):	359	def _copy_file_o_json_to_out_dir(self, src_file: str, json_file: str, copy_files: list):
275	utils.copy_file_src_exist(src_file, os.path.join(self.out_dir, json_file))	360	utils.copy_file_src_exist(src_file, os.path.join(self.out_dir, json_file))
	@@ -356,7 +441,7 @@ class ObjDump:
356	if file_name.endswith('.json'):	441	if file_name.endswith('.json'):
357	continue	442	continue
358	output = utils.get_all_section_symbols_in_file(file_path)	443	output = utils.get_all_section_symbols_in_file(file_path)
359	- print(f'===== [elf heard infos] in {file_name} =====:')	444	+ print(f'===== [elf header infos] in {file_name} =====:')
360	for line in output.split('\n'):	445	for line in output.split('\n'):
361	print(line)	446	print(line)
362		447
	@@ -407,7 +492,7 @@ class ObjDump:
407	print(f'[KERNEL LEN]: {kernel_info.kernel_len}')	492	print(f'[KERNEL LEN]: {kernel_info.kernel_len}')
408	print(f'[ASCEND META]: {self._show_elf_ascend_meta_obj()}')	493	print(f'[ASCEND META]: {self._show_elf_ascend_meta_obj()}')
409	if self.parse_obj_mode == ParseObjMode.MODE_VERBOSE:	494	if self.parse_obj_mode == ParseObjMode.MODE_VERBOSE:
410	- print(f'====== [elf heard infos] ======')	495	+ print(f'====== [elf header infos] ======')
411	print(utils.get_all_section_symbols_in_file(kernel_info.kernel_file))	496	print(utils.get_all_section_symbols_in_file(kernel_info.kernel_file))
412		497
413		498
	@@ -431,7 +516,7 @@ class ObjDump:
431	index += 4	516	index += 4
432	if (index + l) <= len(content):	517	if (index + l) <= len(content):
433	self._show_ascend_meta_tlv(content, t, l, index)	518	self._show_ascend_meta_tlv(content, t, l, index)
434	- if t == TYPE_FUNCTION_ENTRY:	519	+ if t == F_TYPE_FUNCTION_ENTRY:
435	index = index + l + 3	520	index = index + l + 3
436	else:	521	else:
437	index += l	522	index += l
	@@ -439,12 +524,13 @@ class ObjDump:
439	def _get_elf_ascend_meta_op_tlv(self, meta_infos: dict):	524	def _get_elf_ascend_meta_op_tlv(self, meta_infos: dict):
440	meta_name = 'meta'	525	meta_name = 'meta'
441	if meta_name in meta_infos.keys():	526	if meta_name in meta_infos.keys():
		527	+ self._aicore_binary_meta_printed.clear()
442	meta_lists = meta_infos[meta_name]	528	meta_lists = meta_infos[meta_name]
443	content = self._get_segment_content(int(meta_lists[1], HEX_NUM), int(meta_lists[2], HEX_NUM))	529	content = self._get_segment_content(int(meta_lists[1], HEX_NUM), int(meta_lists[2], HEX_NUM))
444	idx = 0	530	idx = 0
445	index = 0	531	index = 0
446	print(f'{KEY_ASCEND_META_OP} META INFO')	532	print(f'{KEY_ASCEND_META_OP} META INFO')
447	- while index < len(content) and idx < 4:	533	+ while index < len(content) and (idx < 4 or self.obj_type == ObjType.TYPE_AICORE_BINARY):
448	if index + 4 > len(content):	534	if index + 4 > len(content):
449	break	535	break
450	t, l = struct.unpack("2H", content[index:index+4])	536	t, l = struct.unpack("2H", content[index:index+4])

Mutils/msobjdump/msobjdump/utils.py+10-2

	@@ -17,7 +17,7 @@ import subprocess
17		17
18		18
19	def split_str_with_space(input_str: str) -> list:	19	def split_str_with_space(input_str: str) -> list:
20	- result = re.split('\s+', input_str)	20	+ result = re.split(r'\s+', input_str)
21	return [element for element in result if element != '']	21	return [element for element in result if element != '']
22		22
23		23
	@@ -57,4 +57,12 @@ def get_all_section_symbols_in_file(file_name: str) -> str:
57		57
58		58
59	def get_o_file_from_a_file(file_path: str, file_name: str) -> str:	59	def get_o_file_from_a_file(file_path: str, file_name: str) -> str:
60	- return subprocess.run(['ar', 'x', file_path, file_name], capture_output=True)	60	+ return subprocess.run(['ar', 'x', file_path, file_name], capture_output=True)
		61	+
		62	+
		63	+def extract_aicore_binary_from_elf(input_file: str, output_file: str):
		64	+ return subprocess.run(
		65	+ ['llvm-objcopy', '-O', 'binary', '--only-section=.aicore_binary', input_file, output_file],
		66	+ capture_output=True,
		67	+ text=True
		68	+ )