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, and mt_k, are now surfaced in impl_param columns 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.

setup.py

Added itanium-demangler dependency

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.
• Added additional cutlass, ck, and nvjet parsing logic.
• Added sources for every kernel pattern / library
• Updated PR description

[ajassani]

Tile size is an implementation detail, not a workload identity. compare_perf_reports_pytorch.py uses every param: column as a merge key, so adding mt_* to param_details breaks cross-hardware comparison. Rather than adding a one-off exclusion for these fields, I suggest introducing an impl_param: prefix: keep logical workload identity under param: and surface kernel-selection details such as mt_m, mt_n, and mt_k under impl_param:. The comparator can continue merging only on param: columns while preserving implementation details as per-report columns. This also scales to future fields such as kernel library, pipeline stages, wave tile, split-K, and algorithm ID.

[ajassani]

Six unrelated SGLang references (InferenceAttention_fwd.csv, RMSNorm_fwd.csv, and UnaryElementwise.csv under both decode and prefill/decode) only change None/nan serialization to empty cells. These changes are unrelated to GEMM kernel parsing and add regression-reference noise. Could they be reverted from this PR?

[ajassani]

The GRIDWISE_CONFIG comment references docs/ck_kernel_tile_indices.md, but that file is not present anywhere in this PR revision. Could you either add the referenced parameter table or remove/update the comment to point to an existing source?

[ajassani]

itanium-demangler is added to setup.py under install_requires, so the implementation is consistent as a mandatory dependency and normal installation will provide it. I am fine with making it required because it is a small pure-Python library with no runtime dependencies. The PR description currently says the dependency is optional and that demangling is skipped when it is unavailable; that is not how the code behaves because it is imported unconditionally. Please update the description to present it as a required dependency.

[kyle-hoffmeyer]

@ajassani, added impl_param columns, reverted unrelated changes to CSVs, fixed docstrings, and updated PR description.