extract-asm-triton

name	extract-asm-triton
description	Extract GPU ISA from Triton kernels on XPU. Triton compiles through the same IGC backend as SYCL (Triton IR → SPIR-V → IGC → zebin). Extraction is identical to sycl-jit: IGC_ShaderDumpEnable=1 captures the zebin at runtime. Use when extracting ASM from torch.compile fusions (triton_per_fused, triton_poi, triton_red), standalone @triton.jit kernels, or Inductor-generated XPU kernels.

Extract ASM from Triton Kernels on XPU

Triton on XPU uses the same IGC backend as SYCL JIT. The path is: Triton IR → ttgir → SPIR-V → IGC → zebin. Extraction is mechanically identical to extract-asm-syclkernel-jit — both use IGC_ShaderDumpEnable=1 to capture the runtime-compiled zebin. The only difference is how the kernel enters the pipeline (Triton Python compiler vs DPC++ clang frontend).

How this fits the compilation stack

Triton path:
  @triton.jit Python → Triton IR (ttir) → Triton GPU IR (ttgir)
                                                    ↓
                                              SPIR-V (via triton-xpu backend)
                                                    ↓
                                              IGC (JIT) → zebin
                                                    ↓
                                         IGC_ShaderDumpEnable=1 → .asm

Compare with SYCL JIT:
  SYCL C++ → LLVM IR → SPIR-V → IGC (JIT) → zebin → .asm

The back half (SPIR-V → IGC → zebin) is IDENTICAL.

When to use

• Hot kernel is triton_per_fused_* / triton_poi_* / triton_red_* (from torch.compile / Inductor)
• Or: standalone @triton.jit kernel in a Python script
• You need the actual GPU ISA (not Triton IR)

When NOT to use

• Kernel is _ZTS… (SYCL) → use extract-asm-syclkernel-{aot,jit}
• Kernel is oneDNN ngen → use extract-asm-onednn (different stack)
• You only need Triton IR → read ${TRITON_CACHE_DIR:-~/.triton/cache}/ directly

Steps

1. Identify the kernel name.

   Inductor fusion (from torch.compile):

   TRITON_CACHE="${TRITON_CACHE_DIR:-$HOME/.triton/cache}"
   rm -rf "$TRITON_CACHE" /tmp/torchinductor_$USER
   TORCH_LOGS=output_code python <repro.py> 2>&1 \
     | grep -oE 'triton_(poi|per|red)_fused_[A-Za-z0-9_]+' | sort -u
   

   Standalone @triton.jit: use unitrace to pin the kernel:

   unitrace -d python <repro.py>
   # Or fallback:
   SYCL_UR_TRACE=-1 python <repro.py> 2>&1 \
     | grep -oP 'pKernelName = 0x[0-9a-f]+ \(\K[^)]+' | sort -u
   

2. Re-run with IGC dump (cold cache).

   TRITON_CACHE="${TRITON_CACHE_DIR:-$HOME/.triton/cache}"
   rm -rf "$TRITON_CACHE" /tmp/torchinductor_$USER
   OUT="<workdir>/triton_$(date +%Y%m%d_%H%M%S)"
   mkdir -p "$OUT/igc"
   
   IGC_ShaderDumpEnable=1 \
   IGC_DumpToCustomDir="$OUT/igc" \
   ONEAPI_DEVICE_SELECTOR=level_zero:0 \
     python <repro.py> 2>&1 | tee "$OUT/run.log"
   

3. Match kernel name to .asm.

   NAME="<fusion-or-kernel-name>"
   MATCH=$(grep -l "$NAME" "$OUT/igc"/OCL_asm*_simd*_entry_*.asm | head -1)
   echo "asm-file: $MATCH"
   

   Cross-check: .zeinfo reports simd_size — confirm it matches num_warps * 32.