asm-source-mapping

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.

   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.

   readelf -S "$ZEBIN" | grep -q .debug_line || { echo "NO debug_line"; exit 1; }
   

3. Decode the line table.

   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.

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:

// 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 pattern	Source construct
Dense dpas chain, 8×repeat	GEMM tile (matmul accumulate)
mul :f + broadcast src1	Scalar rescale (softmax denominator)
exp2 :f / log2 :f	Softmax numerics
mov :bf :f + store	BF16 epilogue / output write
send.slm	SLM load/store
send.slm store + fence.slm + send.slm load + add :f	Reduction tree (workgroup reduce)
send.ugm	Global memory load/store
send.gtwy + sync.bar	Barrier / synchronization
VxH indirect mov r[a0.0]	Sub-group shuffle (shift_group_left)

Output Format

JSON array, one entry per IP:

[
  {
    "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):

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
});

$ 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)

# 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