菜单
Extract GPU ISA from AOT-compiled SYCL binaries using NEO driver runtime dump. Use when the kernel is a SYCL symbol (_ZTS...) from an AOT binary such as libtorch_xpu.so, SYCL-TLA FMHA, or any DPC++ binary compiled with -fsycl-targets=spir64_gen -Xs "-device <gpu>".
用 Codex 或 Claude 帮你安装 复制这段 Prompt，粘贴到 Codex、Claude 或其他助手里，让它检查 Skill 页面并帮你完成安装。
基于 SOC 职业分类
Analyze UT (unit test) results for a torch-xpu-ops PR. Use when asked to check test results, analyze CI failures, or evaluate test coverage for a PR. Produces a structured report of new failures, failure relevance, and new test coverage.
Rules for automatically determining which disable_* labels to apply to a PR based on the file paths changed. Used by the auto-label workflow.
Extract GPU ISA from oneDNN ngen-JIT kernels. This is the ONLY codegen path that bypasses the standard SYCL/SPIR-V/IGC stack. oneDNN uses its own native code generator (ngen) that directly emits GPU ISA bytes — no SPIR-V, no IGC, no zebin ELF, no .debug_line. Use when extracting ASM from oneDNN kernels (gemm_kernel, gen_conv_kernel), matmul, linear, conv, or SDPA-graph ops dispatched via mkldnn.
Extract GPU ISA from JIT-compiled SYCL kernels. The binary only contains SPIR-V; IGC compiles it to native ISA at first launch. The zebin only exists in memory at runtime, so IGC_ShaderDumpEnable=1 is required to capture it. Use when extracting ASM from SYCL JIT kernels, torch-xpu-ops built with TORCH_XPU_ARCH_LIST=none, or standalone DPC++ without AOT flags.
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 Intel GPU ISA (assembly) from any XPU kernel. Classifies the codegen path (SYCL AOT, SYCL JIT, Triton, or oneDNN ngen) and delegates to the matching extraction skill. Use when asked to extract ASM, disassemble XPU kernels, get GPU ISA for an aten op, dump shader for PyTorch XPU, or disassemble a standalone DPC++/Triton binary.
| name | extract-asm-syclkernel-aot |
| description | Extract GPU ISA from AOT-compiled SYCL binaries using NEO driver runtime dump. Use when the kernel is a SYCL symbol (_ZTS...) from an AOT binary such as libtorch_xpu.so, SYCL-TLA FMHA, or any DPC++ binary compiled with -fsycl-targets=spir64_gen -Xs "-device <gpu>". |
IGC_ShaderDumpEnable does NOT work — IGC is never invoked at runtime.ocloc disasm to produce readable ASM._ZTS…) from an AOT-compiled binary whose
AOT target matches the current device:
libtorch_xpu.so / libtorch_xpu_ops.so (default PyTorch XPU build).so-fsycl-targets=spir64_genextract-asm-syclkernel-jitextract-asm-syclkernel-jitgemm_kernel, gen_conv_kernel) → use extract-asm-onednntriton_*) → use extract-asm-tritoncommand -v ocloc >/dev/null || { echo "ocloc not found; source oneapi-vars.sh"; exit 1; }
Run the workload with NEO debug keys. Only kernels actually dispatched to the GPU will be dumped — no unused kernels appear.
OUT="<workdir>/aot_$(date +%Y%m%d_%H%M%S)"
mkdir -p "$OUT" && cd "$OUT"
DumpZEBin=1 NEOReadDebugKeys=1 <repro_cmd>
ls *.elf
# Detect device name (bmg, dg2, pvc, etc.)
DEVICE=$(sycl-ls 2>/dev/null | grep -oP '(?<=\[)[^]]+' | head -1 | awk '{print tolower($NF)}')
DEVICE=${DEVICE:-bmg} # fallback
for elf in *.elf; do
name=$(basename "$elf" .elf)
ocloc disasm -file "$elf" -dump "${name}_dump" -device "$DEVICE"
done
ASM files are at <name>_dump/.text._ZTSxxxx.asm. Demangle to find yours:
for f in *_dump/.text._ZTS*.asm; do
mangled=$(basename "$f" .asm | sed 's/^\.text\.//')
echo "$f → $(c++filt "$mangled")"
done
Match the demangled name to the kernel you are analyzing.