Add xDiT kernel source mapping support

[kyle-hoffmeyer]

Summary

This PR adds support for kernel source mapping for xDiT diffusion workloads. Each fix addresses a specific failure observed during source resolution of device kernels in these traces.

Change 1: Replace c++filt with itanium-demangler

File: source_resolver.py — _demangle_itanium() replacing _cxxfilt_base()

Problem: c++filt is the tool used to demangle kernel names. An example of a mangled kernel name:

_ZN2ck16tensor_operation6device12_GLOBAL__N_139kernel_grouped_conv_fwd_xdl_cshuffle_v3I...

These kernel names are very difficult to parse but can be demangled to more easily parse.
c++filt (GNU Binutils) fails on xDiT workloads due to deeply nested CK convolution kernel symbols (~20+ levels of template nesting, 1575+ characters). When demangling fails, c++filt returns the input string unchanged — an ambiguous failure signal that causes _non_patchable_kind() to misclassify CK kernels as patchable.

itanium-demangler handles deeper nesting (no hardcoded recursion limits) and returns "" on failure (unambiguous). This makes the existing if demangled: guard in _non_patchable_kind() work correctly — "" is falsy, so the mangled regex fallback fires and CK kernels are correctly classified as aiter_ck.

xDiT example: aten::miopen_convolution dispatches to CK grouped convolution:

_ZN2ck16tensor_operation6device12_GLOBAL__N_139kernel_grouped_conv_fwd_xdl_cshuffle_v3I... (1575 chars)

c++filt:            returns input unchanged (failure, no signal)
                    → if demangled: True → searches "ck::" in mangled string → no match
                    → mangled regex fallback never reached
                    ✗ CK kernel misclassified as patchable

itanium-demangler:  returns "" (clear failure signal)
                    → if demangled: False → falls through
                    → mangled regex fires → matches "ck"
                    ✓ CK kernel classified as non-patchable

14 CK convolution kernels across 4 xDiT traces are affected.

Comparison:

	c++filt (GNU Binutils)	itanium-demangler (Python)
Deeply nested CK templates (~20 levels)	Fails (returns input)	Handles them
Failure signal	Returns input unchanged (ambiguous)	Returns "" (clear)
Dependencies	Requires binutils installed in container	pip package, no system deps
Performance	Subprocess spawn per call (cached)	Pure Python, in-process (cached)
Imports required	shutil, subprocess	None (graceful ImportError fallback)

Change 2: Fix _base_from_demangled() for (anonymous namespace):: kernels

File: source_resolver.py — _base_from_demangled()

Problem: _base_from_demangled() is the function used to parse a demangled kernel name. _base_from_demangled() used re.split(r"[(<]", ...) to extract the kernel name. For signatures containing (anonymous namespace), the ( in the namespace qualifier was matched first, truncating the name to "void".

Fix: Rewrite _base_from_demangled() to match the logic in _bypass_source_resolver._demangle_kernel_name(), which already handles this case correctly. Both functions now use the same approach for extracting the bare function identifier: strip void prefix, strip (anonymous namespace)::, then remove template/function args. Ideally this shared logic would live in a common utility function, but that belongs to a broader cleanup PR.

Change 3: Add Tensile and MIOpen non-patchable classification

File: source_resolver.py — _non_patchable_kind()

Problem: Tensile GEMM kernels (Cijk_*) and MIOpen convolution kernels were reported as "unresolved" with no explanation. These kernels are written in GPU assembly and shipped pre-compiled — no source code exists.

Fix: _non_patchable_kind() now accepts an op_name parameter. Tensile is identified by device kernel name prefix (Cijk_*). MIOpen is identified by op name ("miopen" in op_name) rather than device kernel name patterns, which vary across MIOpen versions (e.g., igemm_*, batched_transpose_*, Im2d2Col, Im3d2Col, SubTensorOpWithScalar*). Using the op name catches all MIOpen kernels regardless of the device kernel naming convention.

xDiT examples:

Tensile GEMMs (~37% of xDiT device kernels):

aten::mm       → Cijk_Ailk_Bljk_SBB_BFP16_BFP16_BFP16_...   → tensile_precompiled
aten::addmm    → Cijk_Ailk_Bljk_SB_BFP16_BFP16_BFP16_...    → tensile_precompiled
aten::bmm      → Cijk_Alik_Bljk_SB_BFP16_BFP16_BFP16_...    → tensile_precompiled

MIOpen convolutions (all caught via op name aten::miopen_convolution):

aten::miopen_convolution → igemm_fwd_gtcx3_nchw_fp16_...     → miopen_precompiled
aten::miopen_convolution → batched_transpose_b_n_c_h_w_...   → miopen_precompiled
aten::miopen_convolution → Im3d2Col                           → miopen_precompiled

These are now reported as "non_patchable" with a reason string instead of bare "unresolved".

Change 4: Mark Inductor-generated Triton kernels as non-patchable

File: _bypass_source_resolver.py — resolve_triton_py()
File: _bypass_report.py

Problem: torch.compile Inductor-generated Triton kernels are not patchable as the source for them is generated at compile time within Inductor. These kernels were previously reported as "unresolved" from resolve_triton_py. Tier 1 correctly identified them as non-editable (via torchinductor in the kernel_file path), but returned "unresolved" instead of "non_patchable", causing tiers 2 and 3 to run unnecessarily.

Fix: When resolve_triton_py() has a kernel_file path but editable_trace_source() rejects it, return "non_patchable" instead of "unresolved". The framework scan/repo scan gates in _bypass_report.py now use source_method == "unresolved", so only genuinely unresolved kernels fall through to tiers 2 and 3.

xDiT example: Inductor-generated Triton kernels:

triton_poi_fused_addmm_cat_gelu_slice_view_25
  kernel_file: /tmp/torchinductor_root/4o/c4odgozyt6ydnxpdsf...
  Before: "unresolved" (framework + repo scan ran, found nothing)
  After:  "non_patchable" (framework + repo scan skipped)

Verification

Tested via the full pipeline (analyze_trace(top_k=0) → build_candidates(top_k=0)) in tracelens-xdit:v26.7 (TraceLens patched xDiT image) against 4 xDiT diffusion workloads. Traces can be found in xDiT folders in https://github.com/AMD-AGI/TraceLens/tree/main/tests/traces/inference.

Kernel type	Before	After
Tensile GEMM (Cijk_*)	unresolved	non_patchable (tensile_precompiled)
MIOpen IGEMM (igemm_*)	unresolved	non_patchable (miopen_precompiled)
MIOpen helpers (batched_transpose_*)	unresolved	non_patchable (miopen_precompiled)
CK conv (_ZN2ck...)	unresolved (c++filt failed silently)	non_patchable (aiter_ck)
Inductor Triton (triton_poi_*)	unresolved (framework + repo scan wasted)	non_patchable (framework + repo scan skipped)
AITER groupnorm	unresolved (_base_from_demangled returned "void")	resolved → groupnorm.cu
AITER/Triton attention	resolved	resolved (unchanged)

New dependency

itanium-demangler (pip package, pure Python, no system deps). Added to pyproject.toml runtime optional dependencies. When not installed, a warning is logged ("itanium-demangler is not installed. Kernel classification may be degraded.") and demangling falls back to the mangled regex path.

[github-actions]

CI E2E report — ✅ Succeeded

item	value
result	✅ Succeeded
model	Qwen/Qwen3-0.6B (dense)
resources	1× GPU, TP=1
PR branch	khoffmey/feat/op_source_mapping_enhancements
commit	d2857d49aba8a8bcc1b0c1aa133ce68d1b84207e
session_id	2b3b4188-bdf0-46a5-ae0c-393488a88481
queue → dispatch	0s
run time	178m 50s
total	178m 50s

details

[kyle-hoffmeyer]

I'm not sure if I added the itanium-demangler dependency correctly. Please take a look.

[Ahmedhasssan-aig]

Can you please try to pass all the CI tests?

[Ahmedhasssan-aig]

Can we keep the _cxxfilt_base (to use it as fallback) if itanium is not installed or missing for any reason?