Improve GEMM kernel parsing

[kyle-hoffmeyer]

NOTE: Reference reports regenerated. Difference in reference reports is due to addition of new fields in param_details in perf reports. Change in param: transpose column also affected the rows that are matched via compare_perf_report.

Summary

Extends kernel name parsing to extract tile sizes (mt_m, mt_n, mt_k) and transpose flags from GEMM GPU kernel names across all major kernel libraries observed in MI300 and H100 traces. Previously only Tensile and nvjet kernels were parsed. Tile info is now surfaced in param_details alongside the existing transpose field.

Changes

General

Added six new parser pairs covering all GEMM kernel libraries observed in traces:

Library	Detection	Tile pattern	Transpose pattern	Example	Source
Tensile Cijk_*	is_tensile_gemm (renamed)	MT(\d+)x(\d+)x(\d+)	axis letters: Alik/Ailk, Bljk/Bjlk	Cijk_Alik_Bljk_..._MT64x32x256_MI16x16x1_...	https://rocm.docs.amd.com/projects/Tensile/en/docs-7.0.2/src/reference/nomenclature.html, https://github.com/ROCm/Tensile/wiki/Kernel-Parameters#kernel-names
Triton GEMM	is_triton_gemm	BLOCK_SIZE_[MNK]_(\d+)	not extractable	_gemm_a8w8_blockscale_kernel_..._BLOCK_SIZE_M_128_BLOCK_SIZE_N_128_BLOCK_SIZE_K_128_...	https://rocm.docs.amd.com/projects/ai-ecosystem/en/latest/optimization/workload-optimization.html
CK demangled	is_ck_gemm	parameter index. varies across kernels	first two type params: RowMajor=not transposed, ColumnMajor=transposed	Varies	https://rocm.docs.amd.com/projects/composable_kernel/en/docs-7.1.0/doxygen/html/. Full list of sources is in GRIDWISE_CONFIG in parse_ck_gemm
iGEMM	is_igemm	bt(\d+)x(\d+)x(\d+)	not applicable — iGEMM is always NHWC convolution; fwd/bwd/wrw encodes the pass, not matrix transpose	igemm_fwd_gtcx3_nhwc_bf16_bx0_ex1_bt128x128x32_wt32x32x8_...	https://github.com/ROCm/rocm-libraries/tree/develop/projects/miopen/src/kernels/dynamic_igemm/igemm_gtc_xdlops_nhwc. Contains kernel names with config underneath.
cuBLAS XMMA	is_cublas_xmma_gemm	tilesize(\d+)x(\d+)x(\d+)	_([tn])([tn])_	sm90_xmma_gemm_bf16bf16_bf16f32_f32_tn_n_tilesize128x128x64_warpgroupsize1x1x1_..._cublas	Can't find, but parameters can be inferred.
CUTLASS	is_cutlass_gemm	(\d+)x(\d+)_(\d+)x(\d+) or (\d+)x(\d+)x(\d+)	_([tn]{2})_	void cutlass::Kernel2<cutlass_80_tensorop_bf16_s16816gemm_bf16_256x128_64x3_tn_align2>(...)	https://docs.nvidia.com/cutlass/latest/media/docs/cpp/profiler.html
nvjet	is_nvjet_gemm (renamed)	nvjet_\w+_(\d+)x(\d+)_(\d+)x\d+	last 3 chars: T/N per position (A, B, C)	nvjet_tst_144x128_64x6_2x1_v_bz_bias_TNN	Can't find, but parameters can be inferred.

Also renamed is_rocm_gemm→is_tensile_gemm, parse_rocm_gemm→parse_tensile_gemm, is_cuda_gemm→is_nvjet_gemm, parse_cuda_gemm→parse_nvjet_gemm for clarity.

Transpose is not in Triton kernel names — layout is fixed by the kernel function (e.g., a8w8_blockscale is always activations × weights = row × column) and compiled into the JIT binary rather than parameterized in the name. For iGEMM, the concept of transpose doesn't apply — these are always NHWC convolution kernels where fwd/bwd/wrw indicates the convolution pass direction, not matrix orientation.

CK Kernels

The profiler records CK kernels in two forms: demangled names are human-readable C++ (e.g., void ck::kernel_gemm_xdl_cshuffle_v3<ck::GridwiseGemmMultiD_ABScale_xdl_cshuffle_v3<...>>), produced when the profiler resolves symbol names; mangled names are the raw encoded form the C++ compiler emits (e.g., _ZN2ck16tensor_operation6device...), where integers appear as ELi128E rather than 128.

Using python package itanium-demangler, the mangled kernel name can be demangled and parsed in the same way demangled kernels are parsed.

TraceLens/PerfModel/kernel_name_parser.py

Extended kernel parsing logic to cover different libraries

TraceLens/PerfModel/perf_model.py

mt_m, mt_n, mt_k, and transpose is now surfaced in param_details for all parsed kernel libraries. Fields default to "Unknown" if the kernel name doesn't match any parser or the field is not extractable for that library.

The K-tile field was previously named depth_u, which was Tensile specific. It has been renamed to mt_k to be consistent with mt_m/mt_n.

tests/test_kernel_name_parser.py (new)

Integration tests using verbatim kernel names extracted from trace files.
5 tests per library (35 total).

[gabeweisz]

Looks good - fix the failing test and I will approve

[ajassani]

@kyle-hoffmeyer are you sure of all the patterns here? Can we include the sources?

[kyle-hoffmeyer]

@gabeweisz, @ajassani,

• Added de-mangling logic for mangled CK kernels using itanium-demangler rather than trying to directly parse mangled kernel names which was unreliable. Added this package in setup.py. This new package isn't blocking. If a user doesn't have it installed, it will skip kernel de-mangling and won't error.
• Added additional cutlass, ck, and nvjet parsing logic.
• Added sources for every kernel pattern / library
• Updated PR description