gfx-kernel-metadata

Reading AMD GPU kernel metadata from .hsaco

When tuning HIP kernels you often need to know what the compiler actually allocated: VGPRs, SGPRs, LDS bytes, and whether anything spilled. Those numbers gate occupancy and reveal whether the kernel is register-bound, LDS-bound, or neither.

What .hsaco actually is

A .hsaco is not a raw ELF — it's a __CLANG_OFFLOAD_BUNDLE__ container that wraps an amdgcn-amd-amdhsa--gfxNNN ELF (and possibly a host stub). You must unbundle it before any objdump / readelf tool will recognize it. First 16 bytes will look like __CLANG_OFFLOAD_, confirming the wrapper.

Step-by-step

# Pick your arch — auto-detect from the cache directory if hipfire's
# JIT cache is around, or set explicitly. This works for any RDNA or
# CDNA arch (gfx906/908/90a/942/1010/1030/1100/1101/1102/1150/1151/1200/1201).
ARCH="${ARCH:-$(basename "$(ls -1d .hipfire_kernels/gfx* 2>/dev/null | head -1)")}"
ARCH="${ARCH:-gfx1100}"  # fallback if no cache directory exists

# 1. List bundled targets (confirm the gfx target inside)
/opt/rocm/llvm/bin/clang-offload-bundler \
    --list --type=o \
    --input=path/to/kernel.hsaco
# → hipv4-amdgcn-amd-amdhsa--$ARCH
#   host-x86_64-unknown-linux-gnu-

# 2. Unbundle to a real ELF
/opt/rocm/llvm/bin/clang-offload-bundler \
    --type=o --unbundle \
    --input=path/to/kernel.hsaco \
    --output=/tmp/kernel.elf \
    --targets=hipv4-amdgcn-amd-amdhsa--$ARCH

# 3. Read the AMDGPU note section — this is the metadata
/opt/rocm/llvm/bin/llvm-readelf --notes /tmp/kernel.elf

The --notes output contains a YAML block under amdhsa.kernels: with the relevant fields:

.vgpr_count:                82       # VGPRs allocated per wave
.sgpr_count:                18       # SGPRs (scalar regs) per wave
.vgpr_spill_count:          0        # >0 means spill pressure
.sgpr_spill_count:          0
.group_segment_fixed_size:  256      # static LDS bytes per workgroup
.private_segment_fixed_size:0        # private (scratch) bytes per lane
.max_flat_workgroup_size:   32
.wavefront_size:            32       # 64 = CDNA, 32 = RDNA
.uses_dynamic_stack:        false

Spill canary: .vgpr_spill_count and .sgpr_spill_count may be elided from the YAML when zero (depends on toolchain version). The reliable proxy is .private_segment_fixed_size — if it's 0, the kernel uses no scratch memory and therefore has no register spills. Any non-zero value means VGPRs are spilling to private memory and you should investigate.

For a fast multi-kernel comparison:

for k in kernel1 kernel2 kernel3; do
  echo "=== $k ==="
  /opt/rocm/llvm/bin/llvm-readelf --notes "$k.elf" 2>&1 | \
    grep -E "\.(vgpr_count|sgpr_count|vgpr_spill|sgpr_spill|group_segment_fixed_size|private_segment_fixed_size|wavefront_size):"
done

Disassembly (when you need to see actual instructions)

/opt/rocm/llvm/bin/llvm-objdump --disassemble --mcpu=$ARCH /tmp/kernel.elf

Note: the --mcpu= flag must match the target the ELF was built for. Without it, you'll get incorrect or no decode. Useful for: spotting global_load_* patterns (memory parallelism), v_dot4_i32_i8 (dp4a), v_wmma_* (RDNA3+ matrix), v_mfma_* (CDNA matrix), ds_* (LDS) instructions, s_waitcnt and s_barrier placement.

Architecture cheat-sheet

Wave64 (CDNA) needs ~2× the VGPRs of wave32 RDNA for the same per- lane pressure: a 32-VGPR wave32 RDNA kernel ≈ 64-VGPR wave64. CDNA kernels frequently sit at 40–80 VGPR; RDNA kernels in the 30–90 range.

Interpreting the numbers — occupancy tables per arch family

Wave64 — gfx906/908 (256 VGPRs/SIMD)

Wave64 — gfx90a/942 (512 VGPRs/SIMD)

Theoretical waves/SIMD = floor(512 / vgpr_count) capped at 8. Common allocation granule = 8 VGPRs.

Wave32 — gfx1010+ RDNA (1024 VGPRs/SIMD; gfx1200+ has 1536)

Theoretical waves/SIMD = floor(VGPRs_per_SIMD / vgpr_count) capped at 16 (RDNA1: 20). Common allocation granule = 24 VGPRs (rounded up in groups of 8 for some toolchain versions).

For gfx1100 / gfx1151 (RDNA3, 1024 VGPRs/SIMD):

For gfx1200 / gfx1201 (RDNA4, 1536 VGPRs/SIMD), divide 1536 by vgpr_count instead of 1024 — same table shifts by 1.5×.

LDS budget: 64 KB per CU on gfx9xx and gfx10xx-1101; 128 KB on gfx1200+. The kernel claims its share via .group_segment_fixed_size. LDS occupancy = floor(LDS_per_CU / group_segment_fixed_size) workgroups per CU; multiply by waves-per-WG to get waves/CU from the LDS side. Actual occupancy = min(VGPR-occupancy, LDS-occupancy, SGPR-occupancy, max-waves-per-CU).

SGPRs are rarely the binder — every modern AMD arch has 800+ SGPRs/CU.

What "high occupancy" does and doesn't tell you

• Spills (>0) = bug or budget breach. Always investigate. Common culprits: #pragma unroll on too-large loops, deep accumulator trees, large constant arrays. Measure spilled-vs-rolled both ways before committing.
• High theoretical occupancy + low VALUBusy = memory-bound. More occupancy won't help; you need more in-flight HBM transactions per wave (multi-quad interleave, half-wave splits, prefetch).
• Low occupancy (≤2 waves/SIMD) + high VALUBusy = look for register reuse / loop fusion to drop VGPR pressure.
• LDS = 0 + GEMV-shaped kernel is normal for memory-streaming kernels (decode GEMV). LDS only helps if you reuse data.
• WMMA / MFMA kernels run hot on VGPRs because the matrix accumulator (float8_t for WMMA-f32-16x16x16, float16_t etc.) occupies many VGPRs per wave. 80–120 VGPRs is common for fused GEMM kernels even without spills.

Reproducing — copy/paste recipe (parameterized)

# Defaults — override via environment.
ROCM="${ROCM:-/opt/rocm/llvm/bin}"
ARCH="${ARCH:-$(basename "$(ls -1d .hipfire_kernels/gfx* 2>/dev/null | head -1)")}"
ARCH="${ARCH:-gfx1100}"
KERNEL_DIR="${KERNEL_DIR:-.hipfire_kernels/$ARCH}"
TMP="${TMP:-/tmp/hsaco-extract-$ARCH}"
mkdir -p "$TMP"

# Inspect ALL kernels in the cache, or pass an explicit list as $@:
KERNELS=("$@")
if [ ${#KERNELS[@]} -eq 0 ]; then
  mapfile -t KERNELS < <(ls -1 "$KERNEL_DIR"/*.hsaco 2>/dev/null | xargs -n1 basename | sed 's/\.hsaco$//')
fi

printf "%-50s %4s %4s %5s %5s %6s\n" "kernel" "VGPR" "SGPR" "spill" "LDS" "wave"
for K in "${KERNELS[@]}"; do
  ELF="$TMP/$K.elf"
  $ROCM/clang-offload-bundler --type=o --unbundle \
      --input="$KERNEL_DIR/$K.hsaco" \
      --output="$ELF" \
      --targets=hipv4-amdgcn-amd-amdhsa--$ARCH 2>/dev/null
  notes=$($ROCM/llvm-readelf --notes "$ELF" 2>/dev/null)
  vgpr=$(echo "$notes" | grep -m1 vgpr_count | awk '{print $NF}')
  sgpr=$(echo "$notes" | grep -m1 sgpr_count | awk '{print $NF}')
  spill=$(echo "$notes" | grep -m1 private_segment_fixed_size | awk '{print $NF}')
  lds=$(echo "$notes" | grep -m1 group_segment_fixed_size | awk '{print $NF}')
  wave=$(echo "$notes" | grep -m1 wavefront_size | awk '{print $NF}')
  printf "%-50s %4s %4s %5s %5s %6s\n" "$K" "$vgpr" "$sgpr" "$spill" "$lds" "$wave"
done

(Hipfire's modern JIT-cache directory is .hipfire_kernels/<arch>/; older builds may use kernels/compiled/<arch>/. Adjust KERNEL_DIR accordingly.)

Gotchas

• The bundler --targets= string must match exactly. For gfx1100 use hipv4-amdgcn-amd-amdhsa--gfx1100. List with --list first if uncertain.
• llvm-objdump will fail with "not a valid object file" if you forget to unbundle. The error is unhelpful — always unbundle first.
• vgpr_count is post-allocation (granule-rounded). The actual in-use count may be lower; check disassembly's highest v<N> reference if you need the un-rounded value.
• code-object-v4 and v5 differ in metadata layout; --notes works for both but raw .amdhsa_* directive parsing does not.
• The spill fields (.vgpr_spill_count, .sgpr_spill_count) are sometimes omitted from the YAML when zero. .private_segment_fixed_size: 0 is the reliable "no spills" indicator — non-zero always means spills, regardless of which fields the toolchain emits.
• .max_flat_workgroup_size reports __launch_bounds__(N, ...) 's N value (not the M parameter). To find M (waves/SIMD limit), look at .amdhsa_* directives via raw disassembly.