Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

How to unwrap .dat files to examine assembly? #1759

Open
mysoreanoop opened this issue Mar 12, 2025 · 3 comments
Open

How to unwrap .dat files to examine assembly? #1759

mysoreanoop opened this issue Mar 12, 2025 · 3 comments

Comments

@mysoreanoop
Copy link

When running hipblaslt-bench, I am able to observe certain kernels being launched like named below. I want to examine the corresponding assembly code. When I grep for this string, I see it appearing in hipblaslt/build/release/Tensile/library/*.dat file. How can I unwrap the assembly from there? If not, is there any other way I can get the assembly execution trace?

Cijk_Ailk_Bljk_BBS_BH_Bias_HAS_SAV_UserArgs_MT160x160x64_MI16x16x1_SN_LDSB1_AFC1_AFEM1_AFEM1_ASEM1_CLR1_CADS0_DTVA0_DTVB0_EPS0_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMB_GLS0_ISA90a_K1_LBSPPA2560_LBSPPB128_LBSPPM0_LPA16_LPB16_LPM0_LRVW8_LWPMn1_MIAV0_MIWT5_5_MO40_NTn1_NTA0_NTB0_NTC0_NTD0_NTM0_NEPBS0_NLCA5_NLCB1_ONLL1_PGR2_PLR1_PKA1_SIA3_SS1_SPO0_SRVW0_SSO0_SVW1_SK0_SKXCCM0_TLDS1_ULSGRO0_USL1_UIOFGRO0_USFGROn1_VSn1_VWA1_VWB1_WSGRA0_WSGRB1_WG32_8_1

@mysoreanoop
Copy link
Author

I see the associated .co file with the same file name. Is there a way to disassemble it?

@harkgill-amd
Copy link

harkgill-amd commented Mar 14, 2025

Hi @mysoreanoop, you can use llvm-objdump to examine the assembly code associated with a .co file. For example,

/opt/rocm/llvm/bin/llvm-objdump --disassemble <.co file>

For more information on this, you can reference the documentation at llvm-objdump.

@mysoreanoop
Copy link
Author

mysoreanoop commented Mar 14, 2025

/opt/rocm/llvm/bin/llvm-objdump: error: '../../Tensile/library/TensileLibrary_BB_BB_HA_Bias_SAV_Type_BB_HPA_Contraction_l_Alik_Bljk_Cijk_Dijk_gfx90a.co': The file was not recognized as a valid object file

I tried it with various .co files. This is also what was selected when I invoked hipblaslt-bench.

ROCm v6.3, gfx90a

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Projects
None yet
Development

No branches or pull requests

2 participants