Extract XPU Kernel ASM (Dispatcher)

Classify which codegen path produced the kernel and delegate to the matching atomic skill to extract its Intel GPU ISA.

Compilation Stack

All XPU kernel compilation ultimately produces the same thing: GPU ISA bytes. The skills are separated not by compilation logic (which is largely identical), but by where the zebin lives:

┌─────────────────────────────────────────────────────────────────┐
│              Shared Compilation Stack                            │
│  SYCL/C++ → LLVM IR → SPIR-V → IGC → zebin (GPU ISA)          │
│                                                                 │
│  • -g / -gline-tables-only acts at frontend (SYCL → LLVM IR)   │
│  • Debug info flows: !dbg → OpLine → DebugLoc → .debug_line    │
│  • JIT vs AOT: same IR pipeline, different WHEN it runs         │
└─────────────────────────────────────────────────────────────────┘

┌───────────────┬──────────────────────────────────────────────────┐
│ Scenario      │ Where is the zebin? How to get it?               │
├───────────────┼──────────────────────────────────────────────────┤
│ sycl-aot      │ Embedded in host binary at build time            │
│               │ → DumpZEBin=1 NEOReadDebugKeys=1 → ocloc disasm  │
├───────────────┼──────────────────────────────────────────────────┤
│ sycl-jit      │ Generated at first launch, only in memory        │
│               │ → IGC_ShaderDumpEnable=1 → dump dir → .asm/.elf  │
├───────────────┼──────────────────────────────────────────────────┤
│ triton        │ Triton compiler → SPIR-V → IGC JIT at launch     │
│               │ → IGC_ShaderDumpEnable=1 (same as sycl-jit)      │
├───────────────┼──────────────────────────────────────────────────┤
│ onednn (ngen) │ BYPASSES the entire SPIR-V/IGC stack             │
│               │ Own JIT: ngen → raw ISA bytes (not zebin ELF)    │
│               │ → ONEDNN_JIT_DUMP=1 → IGA ctypes disassembly    │
└───────────────┴──────────────────────────────────────────────────┘

Key implications for downstream:

• sycl-aot, sycl-jit, triton all produce zebin ELF → same ocloc disasm / readelf / .debug_line workflow applies.
• onednn produces raw ISA bytes (no ELF wrapper, no .debug_line) → requires IGA ctypes, pattern-recognition only for source mapping.

When to use

• You have a hot kernel (from profiler / unitrace) and need its ISA.
• Entry points: PyTorch op, standalone @triton.jit, standalone DPC++.
• Hardware: Intel XPU (PVC / BMG-G31 / DG2 / LNL / ARL / MTL / PTL).

When NOT to use

• You only want wall-clock timing — use the profiler directly.
• You only want Triton IR (not ISA) — read ~/.triton/cache/.

Steps

Step 0: Locate tools

Key tools are NOT always on PATH. Probe before proceeding:

# Detect oneAPI root: check env vars first, then common install locations
ONEAPI=${ONEAPI_ROOT:-${CMPLR_ROOT:+${CMPLR_ROOT%/*}}}
if [ -z "$ONEAPI" ]; then
  for d in /opt/intel/oneapi ~/intel/oneapi /usr/local/oneapi; do
    [ -d "$d" ] && ONEAPI="$d" && break
  done
fi

# ocloc (for disassembling zebin ELFs)
command -v ocloc >/dev/null || echo "ocloc not found; source oneapi-vars.sh"

# libiga64.so (for oneDNN ngen disassembly only)
IGA_LIB=$(test -n "$ONEAPI" && find "$ONEAPI" -name 'libiga64.so' 2>/dev/null | head -1)

Step 1: Pin the actually-launched kernel

Identify which kernel was executed on the GPU. The kernel name determines which sub-skill to dispatch to (oneDNN vs Triton vs SYCL).

Preferred: unitrace (covers ALL code paths including native-handle):

unitrace -d <repro_cmd>

Parse the == L0 Backend == table — each data row has the kernel name as the first double-quoted field. Skip rows starting with ze* (those are API calls, not GPU kernels). Extract and deduplicate kernel names.

Probe unitrace location in order: $UNITRACE env var → command -v unitrace → $UNITRACE_HOME/unitrace → <pti-gpu-build>/tools/unitrace/build/unitrace.

Fallback: SYCL_UR_TRACE (zero-dep, but has a blind spot):

SYCL_UR_TRACE=-1 <repro_cmd> 2>&1 | grep -oP 'pKernelName = 0x[0-9a-f]+ \(\K[^)]+'

WARNING: SYCL_UR_TRACE is BLIND to kernels created via urKernelCreateWithNativeHandle (oneDNN ngen, SYCL-TLA, Triton-xpu ≥ 3.7.0). If the list is empty but the workload clearly ran GPU kernels, you MUST install unitrace before proceeding. Do NOT extract ASM blindly.

Step 2: Classify scenario

For _ZTS… kernels, determine AOT vs JIT:

A binary may contain __CLANG_OFFLOAD_BUNDLE but its AOT targets may not cover the current device (e.g. built for PVC, running on BMG → runtime falls back to JIT via SPIR-V). Test which path is actually active:

Method A (preferred): unitrace Kernel Properties

If unitrace was already used in Step 1, check the Compiled column:

• AOT → scenario = sycl-aot
• JIT → scenario = sycl-jit

Method B: IGC dump probe

IGC is only invoked at runtime for JIT compilation. If IGC dump files appear, the kernel was JIT-compiled:

# Quick test: if IGC produces dump files, JIT path is active
rm -rf /tmp/igc_probe && mkdir -p /tmp/igc_probe
IGC_ShaderDumpEnable=1 IGC_ShaderDumpPidDisable=1 IGC_DumpToCustomDir=/tmp/igc_probe <repro_cmd> 2>/dev/null
if ls /tmp/igc_probe/*.asm 2>/dev/null | grep -q .; then
  SCENARIO=sycl-jit
else
  SCENARIO=sycl-aot
fi
rm -rf /tmp/igc_probe

NOTE: Do NOT use DumpZEBin=1 for classification — it produces .elf files for both AOT and JIT scenarios (the runtime always has a zebin to submit, regardless of how it was produced).

If ambiguous → ask the user.

Step 3: Dispatch to atomic skill

onednn   → extract-asm-onednn
triton   → extract-asm-triton
sycl-aot → extract-asm-syclkernel-aot
sycl-jit → extract-asm-syclkernel-jit

Step 4: Emit normalized outputs

All fields REQUIRED:

Examples

Each example shows the dispatcher's classification — how to identify the scenario and what output to expect. The detailed extraction steps are in the respective sub-skill; these examples only demonstrate the classification signal and final result.

Example 1 — oneDNN ngen: BF16 GEMM via aten::matmul

Repro:

import torch
a = torch.randn(4096, 4096, dtype=torch.bfloat16, device='xpu')
b = torch.randn(4096, 4096, dtype=torch.bfloat16, device='xpu')
c = a @ b; torch.xpu.synchronize()

Classification signal: unitrace shows gemm_kernel → scenario = onednn → delegate to extract-asm-onednn.

Expected result:

scenario:    onednn
asm-file:    <workdir>/gemm.asm
kernel-name: gemm_kernel
validation:  grep -c dpas gemm.asm → non-zero (GEMM uses dpas instructions)

Example 2 — Triton: torch.compile(softmax)

Repro:

import torch
@torch.compile
def fn(x): return torch.softmax(x, dim=-1)
x = torch.randn(1024, 1024, device='xpu')
fn(x); torch.xpu.synchronize()

Classification signal: TORCH_LOGS=output_code shows a triton_per_fused_*softmax* kernel name → scenario = triton → delegate to extract-asm-triton.

Expected result:

scenario:    triton
asm-file:    <igc_dump>/OCL_asm*_simd*_entry_*.asm
kernel-name: triton_per_fused_*softmax* (exact name varies by PyTorch version)
validation:  grep 'libdevice.exp' <asm-file> confirms softmax exp computation

Example 3 — SYCL AOT: standalone DPC++ binary

Repro:

// vec_add.cpp
#include <sycl/sycl.hpp>
class VecAddKernel;
int main() {
    sycl::queue q;
    constexpr int N = 1 << 24;
    float *a = sycl::malloc_device<float>(N, q);
    float *c = sycl::malloc_device<float>(N, q);
    q.parallel_for<VecAddKernel>(N, [=](int i) { c[i] = a[i] + 1.0f; }).wait();
    sycl::free(a, q); sycl::free(c, q);
}

icpx -fsycl -O2 -fsycl-targets=spir64_gen -Xs "-device <dev>" vec_add.cpp -o vec_add

Classification signal: unitrace shows VecAddKernel with Compiled = AOT → scenario = sycl-aot → delegate to extract-asm-syclkernel-aot.

Expected result:

scenario:    sycl-aot
asm-file:    <workdir>/<name>_dump/.text._ZTS12VecAddKernel.asm
kernel-name: _ZTS12VecAddKernel
validation:  c++filt _ZTS12VecAddKernel → VecAddKernel

Example 4 — SYCL AOT: libtorch_xpu.so

Repro:

import torch
x = torch.randn(1024, dtype=torch.float, device='xpu')
y = x + 1.0; torch.xpu.synchronize()

Classification signal: unitrace shows the kernel with Compiled = AOT → scenario = sycl-aot → delegate to extract-asm-syclkernel-aot.

Expected result:

scenario:    sycl-aot
asm-file:    <workdir>/<name>_dump/.text._ZTSN2at6native3xpu...E.asm
kernel-name: matches the demangled kernel from c++filt

Example 5 — SYCL JIT: standalone DPC++ without AOT target

Repro:

// shift_reduce.cpp
#include <sycl/sycl.hpp>
class ShiftReduceKernel;
int main() {
    sycl::queue q;
    auto *buf = sycl::malloc_device<int>(1024, q);
    q.parallel_for<ShiftReduceKernel>(
        sycl::nd_range<1>(1024, 32), [=](sycl::nd_item<1> it) {
        int val = buf[it.get_global_id(0)];
        val += sycl::shift_group_left(it.get_sub_group(), val, 1);
        buf[it.get_global_id(0)] = val;
    }).wait();
    sycl::free(buf, q);
}

icpx -fsycl -O2 -g -fsycl-targets=spir64 shift_reduce.cpp -o shift_reduce

Classification signal: unitrace shows ShiftReduceKernel with Compiled = JIT → scenario = sycl-jit → delegate to extract-asm-syclkernel-jit.

Alternatively: IGC_ShaderDumpEnable=1 ./shift_reduce produces .asm files in the IGC dump directory (confirming IGC was invoked at runtime = JIT).

Expected result:

scenario:    sycl-jit
asm-file:    <igc_dump>/OCL_asm*_simd*_entry_*.asm
kernel-name: _ZTS17ShiftReduceKernel
validation:  compiled with -g → grep '// Line' shows source annotations

Example 6 — SYCL JIT fallback: AOT mismatch

Repro:

// vec_add.cpp (same as Example 3)

# Compile AOT for PVC, but run on BMG → runtime falls back to JIT via SPIR-V
icpx -fsycl -O2 -g -fsycl-targets=spir64_gen -Xs "-device pvc" vec_add.cpp -o vec_add_pvc

Classification signal: unitrace may crash on cross-device AOT binaries. Use IGC probe instead: IGC_ShaderDumpEnable=1 ./vec_add_pvc produces .asm files → IGC was called at runtime → JIT fallback confirmed.

Expected result:

scenario:    sycl-jit
asm-file:    <igc_dump>/OCL_asm*_simd*_entry_*.asm
kernel-name: _ZTSN4sycl3_V16detail19__pf_kernel_wrapperI12VecAddKernelEE
validation:  compiled with -g → grep '// Line' shows source annotations
             .platform shows XE2 (BMG), not PVC