Official agent skill

Asm Source Mapping

by intel in intel/torch-xpu-ops

Map Intel GPU ISA instruction addresses to precise SYCL/DPC++ source file:line numbers.

OfficialApache-2.0Auto-check passedDevelopment

Install Asm Source Mapping

skills CLI
$ npx skills add intel/torch-xpu-ops --skill asm-source-mapping -a claude-code

Project install by default; add -g for ~/.claude/skills/.

GitHub CLI
$ gh skill install intel/torch-xpu-ops asm-source-mapping --agent claude-code

Project scope by default; add --scope user for a personal install. Needs GitHub CLI 2.90.0 or later (public preview).

Manual copy
$ git clone --depth 1 https://github.com/intel/torch-xpu-ops.git skills-src && mkdir -p .claude/skills && cp -r skills-src/.claude/skills/asm-source-mapping .claude/skills/asm-source-mapping && rm -rf skills-src

Use ~/.claude/skills/ instead of .claude/skills for a personal install. The folder must contain SKILL.md.

Claude Code skills documentation · loads skills from .claude/skills/

Facts

Skill name
asm-source-mapping
GitHub stars
115
Token cost
~1.8k tokens
SKILL.md length
540 words
Files
1
Skills in repo
29
Repo updated
First seen
Licence
Apache-2.0

At a glance

Map Intel GPU ISA instruction addresses to precise SYCL/DPC++ source file:line numbers.

  • Works in 5 steps: Locate the zebin ELF. → Verify .debug_line exists. → Decode the line table. → …
  • Mapping ASM to source code
  • SKILL.md covers Background, When to use, When NOT to use and Method Priority (strict…, plus 5 more sections
  • Instructions only: no scripts, shell commands, URLs or credentials in SKILL.md

What it does

Asm Source Mapping is an agent skill from intel/torch-xpu-ops, published by the product's own GitHub organization. Map Intel GPU ISA instruction addresses to precise SYCL/DPC++ source file:line numbers. Primary method reads the DWARF .debugline section from the GPU zebin ELF. Fallback uses structural pattern recognition by opcode mix. Use when mapping ASM to source code, finding which source line a GPU instruction comes from, or doing DWARF line table analysis on GPU binaries.

Its SKILL.md is about 1.8k tokens, which your agent loads only when the skill is triggered. It is a single SKILL.md file with no bundled scripts.

It sits in Development. It works with C++. The licence is Apache-2.0.

When your agent uses it

  • Mapping ASM to source code
  • Finding which source line a GPU instruction comes from
  • Doing DWARF line table analysis on GPU binaries

Example prompts

  • “/asm-source-mapping”

Workflow steps

5 steps, taken from the first numbered list in SKILL.md.

  1. Locate the zebin ELF.
  2. Verify .debug_line exists.
  3. Decode the line table.
  4. Build a lookup table. Parse into sorted (offset, file, line) tuples.
  5. Correlate with ASM file. The .asm file uses labels L where N

What it can do on your machine

Read from SKILL.md and the folder at commit 0187b3b. It shows what the files ask for, not the result of running them.

  • Tool permissions

    Pre-approves nothing: there is no allowed-tools line, so your agent's usual permission prompts apply.

    From allowed-tools in the SKILL.md frontmatter.

  • Runs code

    No scripts in the folder and no shell commands in SKILL.md (its code samples are bash, asm, json and cpp).

    From the folder's file list and the shell code blocks in SKILL.md.

  • Network

    No URLs in SKILL.md.

    From URLs in SKILL.md, links to its own repository left out.

  • Credentials

    Names no API keys, tokens, secrets or passwords.

    From names ending in _API_KEY, _TOKEN, _SECRET, _KEY or _PASSWORD in SKILL.md.

Context cost

Asm Source Mapping loads about 1.8k tokens when it runs. Until then it costs about 97 tokens; SKILL.md has 540 words of instructions outside code blocks.

Always · name and description, kept in context so the agent knows when to use it
~97
When it runs · the whole SKILL.md, loaded when a task matches
~1.8k

Estimates: characters ÷ 4, the usual rule of thumb; real counts depend on the model's tokenizer. Scripts and assets cost tokens only if the agent reads them.

Safety

Auto-check passed

The automated check found no risky patterns in SKILL.md.

Automated static check — not a guarantee. Review scripts before installing. It scans the text of SKILL.md for risky patterns (piping downloads into a shell, reading credential files, hidden Unicode, destructive commands); files beside SKILL.md are not scanned.

SKILL.md

The full file from intel/torch-xpu-ops at commit 0187b3b, republished under its Apache-2.0 licence (© intel). 540 words, ~1,787 tokens.

Download SKILL.mdSave it as .claude/skills/asm-source-mapping/SKILL.md (or your agent's skills folder).
name
asm-source-mapping
description
Map Intel GPU ISA instruction addresses to precise SYCL/DPC++ source file:line numbers. Primary method reads the DWARF .debug_line section from the GPU zebin ELF. Fallback uses structural pattern recognition by opcode mix. Use when mapping ASM to source code, finding which source line a GPU instruction comes from, or doing DWARF line table analysis on GPU binaries.

ASM → SYCL Source Mapping

Given one or more GPU instruction addresses (IPs) and an ASM directory, resolve each IP to the original SYCL/DPC++ source file:line and an optional structural construct label.

Background

Debug line info flows: C++ source → LLVM IR !dbg → SPIR-V OpLine → IGC vISA DebugLoc → zebin .debug_line (byte offset → file:line). The zebin's .debug_line section maps GPU ISA offsets to source, and is the primary data source for this skill. Requires -g at compile time (-gline-tables-only is not supported for spir64_gen AOT targets and is silently ignored by the compiler now and may fix it in the future version). oneDNN ngen has no .debug_line — pattern recognition is the only option.

When to use

  • You have one or more GPU instruction addresses (IPs) to map to source.
  • An ASM directory is available.
  • You need source attribution for performance analysis or debugging.

When NOT to use

  • ASM is not available (blocked binary, closed-source kernel).
  • The kernel is a tiny helper (< 1% of total GPU time) — mapping is unlikely useful.

Method Priority (strict waterfall)

Step 1: Try zebin .debug_line          → method = "debug-line"
         ↓ (not found)
Step 2: Try IGC ASM inline comments    → method = "debug-line"
         ↓ (not found)
Step 3: Fallback: pattern recognition  → method = "pattern-recognition"

Rule: if Method 1 succeeds, STOP. Do NOT run pattern-recognition for line attribution. Only add an optional construct label as enrichment.

Method 1 — Zebin .debug_line (primary)

The zebin ELF (e.g. dumped_zebin_module_0.elf) must contain a .debug_line section (present when compiled with -g or -gline-tables-only).

  1. Locate the zebin ELF.

    The zebin is in the extraction working directory (where DumpZEBin=1 NEOReadDebugKeys=1 dumped the .elf file). Typical locations:

    • AOT: <workdir>/dumped_zebin_module_N.elf (from DumpZEBin runtime dump)
    • JIT: zebin is in memory only — use Method 2 instead.
    bash
    ZEBIN=<path to zebin ELF from extraction step>
    file "$ZEBIN"  # must show: ELF 64-bit LSB relocatable, *unknown arch 0xcd*
  2. Verify .debug_line exists.

    bash
    readelf -S "$ZEBIN" | grep -q .debug_line || { echo "NO debug_line"; exit 1; }
  3. Decode the line table.

    bash
    readelf --debug-dump=decodedline "$ZEBIN" 2>/dev/null

    Note: readelf may warn about unknown reloc types for e_machine=205 (Intel GPU). Ignore it — the decoded output is still correct.

    Output format:

    shift_reduce.cpp                  77                   0               x
    shift_reduce.cpp                  91                0xe8               x
    shift_reduce.cpp                  93               0x238               x
  4. Build a lookup table. Parse into sorted (offset, file, line) tuples. For a given IP offset, find the entry with the largest offset ≤ IP (floor lookup).

  5. Correlate with ASM file. The .asm file uses labels L<N> where N is the decimal byte offset in .text. Map IP offset → label → ASM line.

Show full SKILL.md (204 more words)Show less

Method 2 — IGC ASM Inline Comments (secondary)

JIT-compiled kernels (IGC_ShaderDumpEnable=1) produce .asm files with inline source comments. Two formats exist depending on IGC version:

asm
// Format A (newer IGC, requires -g):
// Line 8:  int val = buf[it.get_global_id(0)];
(W) send.ugm (32|M0) r16 r14 ...
// Line 9:  val += sycl::shift_group_left(...);
(W) mov (16|M0) r20.0<1>:d r16.0<1;1,0>:d

// Format B (older IGC):
(W) add (M1, 16) r14.0<1>:d r14.0<1;1,0>:d 0x40:w  // shift_reduce.cpp:93

Scan for // Line <N>: or // <file>:<line> patterns. For each IP, walk backward to the nearest preceding annotated line.

Method 3 — Pattern Recognition (fallback only)

Use ONLY when Methods 1 and 2 both fail (no -g at compile time, or oneDNN ngen JIT which has no DWARF).

Classify each IP by surrounding instruction mix:

Instruction patternSource construct
Dense dpas chain, 8×repeatGEMM tile (matmul accumulate)
mul :f + broadcast src1Scalar rescale (softmax denominator)
exp2 :f / log2 :fSoftmax numerics
mov :bf :f + storeBF16 epilogue / output write
send.slmSLM load/store
send.slm store + fence.slm + send.slm load + add :fReduction tree (workgroup reduce)
send.ugmGlobal memory load/store
send.gtwy + sync.barBarrier / synchronization
VxH indirect mov r[a0.0]Sub-group shuffle (shift_group_left)

Output Format

JSON array, one entry per IP:

json
[
  {
    "ip": "0x238",
    "asm_offset": 568,
    "asm_label": "L568",
    "sycl_file": "shift_reduce.cpp",
    "sycl_line": 93,
    "source_construct": "ALU compute (value += other)",
    "method": "debug-line"
  }
]

Note: asm_label = L<decimal byte offset> (matching .asm file labels). asm_offset is the same value as an integer. ip is the hex form.

Examples

Example 1 — AOT binary with .debug_line (Method 1)

Source (compiled with icpx -fsycl -g -fsycl-targets=spir64_gen):

cpp
class ShiftReduceKernel;
q.parallel_for<ShiftReduceKernel>(nd_range<1>(1024, 32), [=](nd_item<1> it) {
    int val = buf[it.get_global_id(0)];                    // line 8
    val += sycl::shift_group_left(it.get_sub_group(), val, 1); // line 9
    buf[it.get_global_id(0)] = val;                        // line 10
});
bash
$ readelf --debug-dump=decodedline dumped_zebin_module_0.elf
# Output (offsets and lines vary by compiler version):
# <source_file>         <line>    <offset>
# shift_reduce.cpp        8        0xe8
# shift_reduce.cpp        9        0x238
# ...

# Given IP: 0x238
# Floor lookup: largest offset ≤ 0x238 → line 9
# Result:
#   ip=0x238 → shift_reduce.cpp:9
#   construct="ALU compute (shift_group_left)"
#   method=debug-line
Example 2 — JIT with inline comments (Method 2)
asm
# From IGC_ShaderDumpEnable=1 dump (compiled with -g):
// Line 8:  int val = buf[it.get_global_id(0)];
(W)     send.ugm (32|M0)  r16  r14  ...
// Line 9:  val += sycl::shift_group_left(...);
(W)     mov (16|M0)  r20.0<1>:d  r16.0<1;1,0>:d

# IP is at the mov instruction → walk backward → "// Line 9"
# Result:
#   sycl_line=9  method=debug-line
Example 3 — No debug info, oneDNN ngen (Method 3 fallback)
# No .debug_line, no // comments (oneDNN ngen JIT kernel)
# IP surrounded by dense dpas instructions:
#   dpas (8) r40.0<1>:f  r36.0<8;8,1>:hf  r20.0<8;4,1>:hf
#   dpas (8) r48.0<1>:f  r44.0<8;8,1>:hf  r20.0<8;4,1>:hf
# Result:
#   sycl_file=null  construct="GEMM tile/dpas"  method=pattern-recognition

© intel, Apache-2.0. Rendered from Markdown: HTML in the file is shown as text, images as links, and headings moved down two levels. Raw file

Files

Just SKILL.md in .claude/skills/asm-source-mapping of intel/torch-xpu-ops.

Open the folder on GitHubat commit 0187b3b

Compare with similar skills

Asm Source Mapping next to the 5 skills that share the most tags, products or categories with it. Stars are the repository's; “used in” counts other GitHub owners with a copy.

Asm Source Mapping compared with similar skills
SkillStarsUsed inTokensAuto-checkLicenceRepo updated
Asm Source Mapping this skillintel/torch-xpu-ops115—~1.8kAutomated safety check: PassApache-2.0
Qt C++ Code Reviewx-tools-author/x-tools1.1k2 repos~4.3kAutomated safety check: PassBSD-3-Clause
YugabyteDB ASH Instrumentationyugabyte/yugabyte-db11k—~4.5kAutomated safety check: PassCustom licence
pybind11 Release Preparationpybind/pybind1118k—~1.7kAutomated safety check: PassCustom licence
Qt Cpp ReviewSerial-Studio/Serial-Studio7.2k—~4.3kAutomated safety check: PassCustom licence
Paddle Eager GraphPaddlePaddle/Paddle24k—~562Automated safety check: PassApache-2.0

Similar skills

  • Qt C++ Code Review

    x-tools-author/x-tools

    Read-only review of Qt6 C++ code that combines a deterministic lint script with six parallel analysis agents and reports only high-confidence issues.

    1.1k GitHub starsUsed in 2 repos~4.3k tokens
    DevelopmentAuto-check passed
  • YugabyteDB ASH Instrumentation

    yugabyte/yugabyte-db

    Procedure for adding or changing YugabyteDB Active Session History wait states in TServer and DocDB C++ code, including the macro to use for sync and async paths.

    11k GitHub stars~4.5k tokensUpdated today
    DevelopmentAuto-check passed
  • Opens the pybind11 release-preparation pull request: picking the release base, bumping the version in common.h and integrating the changelog, following docs/release.rst.

    18k GitHub stars~1.7k tokensUpdated today
    DevelopmentAuto-check passed
  • Qt Cpp Review

    Serial-Studio/Serial-Studio

    Qt6/C++ deep code review for Serial Studio. An agent skill from Serial-Studio/Serial-Studio.

    7.2k GitHub stars~4.3k tokensUpdated yesterday
    DevelopmentAuto-check passed
  • Paddle Eager Graph

    PaddlePaddle/Paddle

    A skill your agent uses when navigating Paddle eager-mode (dynamic graph) source code, tracing forward/backward execution, debugging autograd issues, understanding PyLayer, or investigating…

    24k GitHub stars~562 tokensUpdated 7 days ago
    DevelopmentAuto-check passed
  • ONNX Runtime Source Build

    microsoft/onnxruntime

    Official

    Builds ONNX Runtime from source with its build scripts, explaining the update, build and test phases, key flags and where the build output lands.

    22k GitHub stars~1.4k tokensUpdated today
    DevelopmentAuto-check passed

More from intel/torch-xpu-ops

All 29 skills in this repo
  • Intel GPU Device Selection

    intel/torch-xpu-ops

    Official

    Select the Intel GPU device to use when a system has multiple Intel GPU devices.

    115 GitHub stars~508 tokensUpdated yesterday
    Auto-check passed
  • Xpu CI Health Check

    intel/torch-xpu-ops

    Official

    Check PyTorch ciflow/xpu (xpu.yml) on the main branch, collect the failing XPU test cases from the most recent completed run(s), analyze the ROOT CAUSE of each failure with AI, and produce a list…

    115 GitHub stars~1.5k tokensUpdated yesterday
    Auto-check passed
  • At Dispatch V2

    intel/torch-xpu-ops

    Official

    Convert PyTorch ATDISPATCH macros to ATDISPATCHV2 format in ATen C++ code.

    115 GitHub starsUsed in 3 repos~2.2k tokens
    Auto-check passed
  • PR Review

    intel/torch-xpu-ops

    Official

    Review pull requests for XPU operator or backend code. An agent skill from intel/torch-xpu-ops.

    115 GitHub stars~4.2k tokensUpdated yesterday
    Auto-check passed
  • Skill Writer

    intel/torch-xpu-ops

    Official

    Guide users through creating Agent Skills for Claude Code. An agent skill from intel/torch-xpu-ops.

    115 GitHub starsUsed in 3 repos~2.4k tokens
    Auto-check passed
  • Ut Issue Authoring

    intel/torch-xpu-ops

    Official

    Read the evidence a nightly UT run produced, decide which failures share a root cause and which are machine breakage rather than product bugs, and write one issue draft per root cause to drafts.json.

    115 GitHub stars~2k tokensUpdated yesterday
    Auto-check passed

Works with

Categories

Questions about Asm Source Mapping

What does Asm Source Mapping do?

Map Intel GPU ISA instruction addresses to precise SYCL/DPC++ source file:line numbers. Asm Source Mapping is an agent skill from intel/torch-xpu-ops, published by the product's own GitHub organization. Map Intel GPU ISA instruction addresses to precise SYCL/DPC++ source file:line numbers.

When should I use Asm Source Mapping?

Asm Source Mapping fits situations like: mapping ASM to source code; finding which source line a GPU instruction comes from; doing DWARF line table analysis on GPU binaries.

How do I install Asm Source Mapping in Claude Code?

Run `npx skills add intel/torch-xpu-ops --skill asm-source-mapping -a claude-code`. Or copy the skill folder (.claude/skills/asm-source-mapping in intel/torch-xpu-ops) into .claude/skills/asm-source-mapping in your project. Claude Code loads it when a task matches its description.

How do I install Asm Source Mapping in Codex?

Run `npx skills add intel/torch-xpu-ops --skill asm-source-mapping -a codex`. Or copy the skill folder (.claude/skills/asm-source-mapping in intel/torch-xpu-ops) into .agents/skills/asm-source-mapping in your project. Codex loads it when a task matches its description.

Can I use Asm Source Mapping in Cursor, Gemini CLI or GitHub Copilot?

Cursor, Gemini CLI, GitHub Copilot and OpenCode also load SKILL.md folders. With the skills CLI, run `npx skills add intel/torch-xpu-ops --skill asm-source-mapping -a cursor` (or -a gemini-cli, github-copilot or opencode for the others). To copy it by hand, put the folder in .cursor/skills/asm-source-mapping, .gemini/skills/asm-source-mapping, .github/skills/asm-source-mapping and .opencode/skills/asm-source-mapping in your project.

What does Asm Source Mapping need to run?

SKILL.md names no scripts, command-line tools or credentials: Asm Source Mapping is instructions for the agent only.

Does Asm Source Mapping access the network?

SKILL.md contains no URLs. Any network use would come from the scripts or tools the agent runs. This is read from the text; nothing was executed.

Is Asm Source Mapping safe to install?

Our automated static check of SKILL.md found no risky patterns, such as piping downloads into a shell, reading credential files or hidden Unicode. It is not a guarantee. Review the folder before installing.

What licence does Asm Source Mapping use?

Asm Source Mapping is published under the Apache-2.0 licence (the repository's licence). It allows redistribution, so the full SKILL.md is shown on this page.

How many tokens does Asm Source Mapping use?

About 1.8k tokens (SKILL.md is roughly 7.1k characters). Agents keep only the skill's name and description in context until a task matches; then they load SKILL.md in full.

What are the alternatives to Asm Source Mapping?

Skills that share tags, products or a category with Asm Source Mapping: Qt C++ Code Review (x-tools-author/x-tools, 1.1k stars), YugabyteDB ASH Instrumentation (yugabyte/yugabyte-db, 11k stars), pybind11 Release Preparation (pybind/pybind11, 18k stars) and Qt Cpp Review (Serial-Studio/Serial-Studio, 7.2k stars). The comparison table on this page puts their stars, adoption, token cost, safety result and licence side by side.

Who maintains Asm Source Mapping?

intel (a GitHub organization, an official publisher) maintains it in intel/torch-xpu-ops, which has 115 GitHub stars. The repository holds 29 skills in this directory. The repository was last updated on October 6, 2026.

Source: intel/torch-xpu-ops on GitHub. Facts on this page come from the repository at the commit we read; the author's words are quoted as theirs.