Flydsl Kernel Trace Analysis

Kernel Trace Analysis

Profile and analyze GPU kernel ATT traces to identify stall hotspots and produce an optimization plan.

All analysis is done by running the bundled hotspot_analyzer.py (ATT) and pmc_l2_analyzer.py (PMC CSVs); never recompute their metrics yourself. Do not use GUI tools.

If the question is only "did my change increase register pressure, spills, or LDS?", use /isa-resource-diff first — it is compile-only, needs no GPU or profiler run, and answers in seconds. Come here when you need to know why a kernel is slow rather than what resources it uses.

To pick an LLVM knob that changes what the compiler emits — a compile hint, a function attribute, an -mllvm flag — and to prove it took effect, use /llvm. That is also compile-only.

Arguments

Analyzer Scripts

• ${CLAUDE_SKILL_DIR}/scripts/hotspot_analyzer.py — reads a ui_output_agent_*_dispatch_* ATT directory; reports top-K stall hotspots, stall-type breakdown, and occupancy (combined-VGPR-pool model, reads accum/LDS/SGPR from out_kernel_trace.csv).
• ${CLAUDE_SKILL_DIR}/scripts/pmc_l2_analyzer.py — reads rocprofv3 PMC counter CSV(s); reports L2 hit rate, HBM 32B-partial fraction, and over-fetch ratio. Use when a kernel is memory-bound and you need to know why (ATT has no cache counters). See "L2 / HBM efficiency analysis" under Step 5.

Workflow

Mode A: Analyze existing dispatch directory

If the user provides --dir <path> or already has a ui_output_agent_*_dispatch_* directory:

bashCopy code
# Both analyzers ship with this skill -- do not reimplement them.
python ${CLAUDE_SKILL_DIR}/scripts/hotspot_analyzer.py <dispatch_dir> --kernel <kernel> --topk 15 --mode both --detail --context 4

Skip to Step 5: Interpret Results.

Mode B: Full collection workflow

Step 1: Kernel Discovery

bashCopy code
touch /tmp/trace_ts
rocprofv3 --stats --kernel-trace -f csv -- <CMD> 2>&1
find . -maxdepth 3 -name "*stats*" -newer /tmp/trace_ts -type f 2>/dev/null

Parse the stats CSV and present a kernel table:

Ask the user which kernel to trace if not obvious.

Prefer results.db if available — use sqlite3 for structured queries:

bashCopy code
sqlite3 results.db "
SELECT ks.KernelName, COUNT(*) calls,
       ROUND(AVG(kd.end-kd.start)/1000.0,1) avg_us
FROM rocpd_kernel_dispatch kd
JOIN rocpd_info_kernel_symbol ks ON kd.kernel_symbol_id=ks.id
GROUP BY ks.KernelName ORDER BY avg_us DESC LIMIT 20;"

Step 2: Configure input.yaml

Write /tmp/trace_input.yaml with the job below, setting kernel_include_regex to the kernel found in Step 1:

yamlCopy code
jobs:
   -
       kernel_include_regex: <KERNEL_NAME_PATTERN>
       kernel_iteration_range: "[1, [3-4]]"
       output_file: out
       output_directory: kernel_trace_output
       output_format: [csv]
       truncate_kernels: true
       sys_trace: true
       advanced_thread_trace: true
       att_target_cu: 1
       att_shader_engine_mask: "0xf"
       att_simd_select: "0xf"
       att_buffer_size: "0x6000000"

Key notes:

• kernel_iteration_range: "[1, [3-4]]" skips warmup, traces dispatches 3-4
• att_buffer_size: 96MB per SE; increase to "0xC000000" if truncated
• att_target_cu: 1: single CU keeps output manageable

Step 3: Collect ATT Trace

bashCopy code
FLYDSL_DEBUG_ENABLE_DEBUG_INFO=1 rocprofv3 -i /tmp/trace_input.yaml -- <CMD> 2>&1
find . -type d -name "ui_output_agent_*" -newer /tmp/trace_ts 2>/dev/null

If the rocprof-trace-decoder library is missing, install the release matching the ROCm version in /opt/rocm/.info/version. Set RTD_VERSION to that release before running the block below. build-rocm-image hardcodes its own RTD_VERSION for the image it builds and does not derive it either; if the two disagree, the ROCm version on the machine you are profiling wins:

bashCopy code
RTD_VERSION=0.1.6   # <-- set to the release matching your ROCm version
RTD_INSTALLER="rocprof-trace-decoder-manylinux-2.28-${RTD_VERSION}-Linux.sh"
wget -q "https://github.com/ROCm/rocprof-trace-decoder/releases/download/${RTD_VERSION}/${RTD_INSTALLER}"
chmod +x "${RTD_INSTALLER}"
"./${RTD_INSTALLER}" --skip-license --prefix=/tmp/rtd-install
find /tmp/rtd-install -name '*.so*' -exec cp -a {} /opt/rocm/lib/ \;
ldconfig

Output structure:

Copy code
ui_output_agent_<PID>_dispatch_<N>/
├── code.json          ← PRIMARY: per-instruction stall/cycle data
├── snapshots.json     ← source file path mapping (virtual → local filename)
├── source_0_*.py      ← embedded source files
├── filenames.json     ← wave file index
├── occupancy.json     ← occupancy timeline
└── se*_sm*_sl*_wv*.json  ← per-wave raw traces

Step 4: Run hotspot_analyzer.py

The analyzer ships with this skill — do not reimplement it. Run:

bashCopy code
# Full report
python ${CLAUDE_SKILL_DIR}/scripts/hotspot_analyzer.py <dispatch_dir> --topk 15 --mode both

# Source-level with code context (best for optimization)
python ${CLAUDE_SKILL_DIR}/scripts/hotspot_analyzer.py <dispatch_dir> --topk 5 --mode src --detail --context 4

# ASM-only for instruction-level detail
python ${CLAUDE_SKILL_DIR}/scripts/hotspot_analyzer.py <dispatch_dir> --mode asm --topk 20

Step 5: Interpret Results

code.json field reference

Each row in code["code"] is:

Copy code
[asm, _, pc_index, source_loc, _, pc_addr, exec_count, total_cycles, stall_cycles, issue_cycles]
  0   1     2          3       4     5          6            7              8             9

• col[8] stall_cycles: cycles the instruction was blocked from issuing — primary hotspot metric
• col[7] total_cycles: total cycles charged to this instruction across all waves
• col[3] source_loc: "/path/to/file.py:LINE" — virtual path resolved via snapshots.json
• col[6] exec_count: number of wave-threads that executed this instruction

snapshots.json: resolving source paths

snapshots.json encodes a nested dict tree mapping virtual paths to local filenames:

jsonCopy code
{"/": {"FlyDSL": {"kernels": {"pa_decode_sw_fp8_ps.py": "source_0_pa_decode_sw_fp8_ps.py"}}}}

Flatten recursively: /FlyDSL/kernels/pa_decode_sw_fp8_ps.py → source_0_pa_decode_sw_fp8_ps.py

Stall type classification

Common hotspot patterns

The snippets below are schematic — they illustrate instruction scheduling, so loads are written in the abbreviated raw-intrinsic form. On the current surface these are fx.copy from a make_buffer_tensor view; the scheduling advice is unchanged either way.

Pattern 1: V/K loads inside MFMA loop → very high stall rate (80–95%)

pythonCopy code
# BAD: load and MFMA alternate — only 1 MFMA of hiding time
for k_step in range_constexpr(QKHELOOP * 2):
    if k_step % 2 == 0:
        v_data = buffer_ops.buffer_load(...)   # stall_rate ~92%
    acc = rocdl.mfma_f32_16x16x32_fp8_fp8(...)

# GOOD: batch all loads before the MFMA loop
for td in range_constexpr(TLOOP):
    v_prefetch[td] = [buffer_ops.buffer_load(...) for _ in range_constexpr(QKHELOOP)]

for td in range_constexpr(TLOOP):
    for k_step in range_constexpr(QKHELOOP * 2):
        acc = rocdl.mfma_f32_16x16x32_fp8_fp8(...)   # entire QK MFMA hides VMEM latency
    v_results[td] = v_prefetch[td]   # already in registers

Pattern 2: Sequential loads with no compute → VMEM queue saturation

pythonCopy code
# BAD: all loads back-to-back, no compute interleaved
for td in range_constexpr(TLOOP):
    for qkhe in range_constexpr(QKHELOOP):
        k4 = buffer_ops.buffer_load(k_rsrc, ka_dw, ...)   # queue fills up

# GOOD: prefetch next tile's K loads during current tile's MFMA computation

Pattern 3: LDS prob reads immediately before PV MFMA → lgkmcnt stall

pythonCopy code
# BAD: LDS reads and MFMA in same loop
for vhe in ...:
    for vt in ...:
        p_i64 = lds_read(...)    # issued here
        tmp = mfma(v_i64, p_i64, ...)   # immediately consumed → lgkmcnt stall

# GOOD: batch all LDS reads first, then all MFMAs
for vhe in ...:
    for vt in ...:
        p_i64s.append(lds_read(...))    # all LDS reads issued first

for vhe in ...:
    for vt in ...:
        tmp = mfma(v_i64s[...], p_i64s[...], ...)   # LDS data already ready

Pattern 4: Scale loads too close to usage

pythonCopy code
# BAD: scale load and usage separated by only TLOOP MFMAs
for td in range_constexpr(TLOOP):
    k_scale = buffer_ops.buffer_load(ks_rsrc, ...)   # issued here
# ... small compute gap ...
    result = acc * k_scale   # used too soon → stall

# GOOD: issue scale loads at the very beginning of the block,
# before K loads, to maximise latency hiding distance

Pattern 5: Hotspot attributed to kernel entry line

When @flyc.kernel / kernel decorator line appears as the top hotspot with a mix of VMEM-wait + barrier stall types — this is a debug info aggregation artifact. MLIR/compiler-generated instructions (address arithmetic, cndmask, prologue setup) map to the outermost scope line. Ignore this line; focus on lines with explicit user ops.

Register pressure check (architecture-aware)

hotspot_analyzer.py auto-detects the GPU architecture from ISA instruction patterns and computes occupancy (waves/SIMD) as the minimum across every resource limiter:

Copy code
occupancy = min(vgpr_limit, lds_limit, sgpr_limit, hw_max=8)
  vgpr_limit = 512 // (arch_vgpr_alloc + accum_vgpr_alloc)              # per SIMD
  lds_limit  = (LDS_total // lds_per_wg) * waves_per_wg // 4_SIMDs      # per SIMD
  sgpr_limit = 800 // sgpr_alloc                                        # per SIMD

VGPR is a combined 512-entry pool on BOTH gfx942 and gfx950. CDNA2 (gfx90a) unified the arch (256) and accum (256) VGPR files into one 512 budget per SIMD, and gfx942/gfx950 inherit that. Occupancy from VGPR is 512 / (arch + accum) on both — NOT 256 / max(arch, accum). (The separate-pool 256/max model only applied to gfx908 / CDNA1, where accum VGPRs were a distinct file accessible only by MFMA.)

What actually changed in CDNA4 vs CDNA3 is the LDS size (64KB→160KB) and the LDS alloc granularity — not the VGPR pooling model.

Reading the real counts. code.json only holds the (often single-CU, often vgpr-form) disassembly, so it cannot reveal accum_vgpr / LDS / SGPR / workgroup size — an AGPR-form-blind ISA scan reports accum=0 and gets occupancy badly wrong. The analyzer reads out_kernel_trace.csv (staged next to the dispatch dir) for the authoritative Accum_VGPR_Count / LDS_Block_Size / SGPR_Count / Workgroup_Size_*. arch_vgpr is taken as max(ISA_scan, CSV) so a bogus-low CSV VGPR_Count field can't under-report. If no CSV is found it falls back to ISA-only and prints a warning.

Auto-detection: gfx950-specific instructions (v_mfma_scale_f32_*, v_mfma_f32_16x16x128_*, v_mfma_f32_32x32x64_*) indicate CDNA4. Absence indicates CDNA3.

bashCopy code
sqlite3 results.db "
SELECT ks.KernelName, ki.arch_vgpr_count, ki.accum_vgpr_count, ki.lds_size
FROM rocpd_kernel_dispatch kd
JOIN rocpd_info_kernel_symbol ks ON kd.kernel_symbol_id=ks.id
JOIN rocpd_info_kernel ki ON kd.kernel_id=ki.id LIMIT 5;"

Worked example (PA decode, gfx942): arch 144 + accum 136 = 280 combined → 512//280 = 1 wave/SIMD, VGPR-bound (LDS allows 5, SGPR allows 7). Reaching 2 waves needs combined ≤ 256, e.g. freeing ~24 VGPRs.

Note: maxnreg has been removed (it never reached LLVM). Forcing accum_vgpr=0 with it doubled occupancy but caused MFMA spills through arch_vgpr — measured 4.5x GPU slowdown. Use waves_per_eu instead.

L2 / HBM efficiency analysis (PMC, not ATT)

When the ATT hotspots are dominated by VMEM-load at high stall rate (e.g. 40-50% of stall, ~94% per-load), the kernel is memory-bound and the next question is why — and ATT cannot answer it (it has no cache counters). Capture PMC counters (see capture-kernel-trace "PMC Mode") and analyze with ${CLAUDE_SKILL_DIR}/scripts/pmc_l2_analyzer.py:

bashCopy code
python ${CLAUDE_SKILL_DIR}/scripts/pmc_l2_analyzer.py \
    /tmp/pmc_out/pass_1/pmc_l2_counter_collection.csv \
    /tmp/pmc_ea_out/pass_1/pmc_ea_counter_collection.csv \
    --kernel <kernel> --ideal-gb <bytes_per_dispatch_GB> --ea-channels 2

Three metrics, three decisions:

Decision tree for a memory-bound decode kernel:

1. L2 hit rate < 5% → pure streaming, no reuse. This is expected and correct for decode with independent per-sequence paged KV — each KV byte is read once; the GQA (×heads) and MTP (×seq) reuse is captured in registers/LDS, never re-reads L2. "Improving L2 hit rate" is a non-goal. The only thing that raises it is real KV reuse = shared-prefix serving (a workload/scheduling property, not a kernel change).
2. 32B fraction ≈ 0% → full 64B cache lines, no spatial-locality waste. Nothing to fix at the line level. (High 32B% would point to scattered/ misaligned access worth restructuring.)
3. over-fetch ≈ 1.0x → the kernel reads exactly the data it needs. The achieved bandwidth (compute as ideal_bytes / kernel_time) is then the real ceiling for this access pattern. 50-60% of theoretical HBM peak is normal even for clean streaming; paged-gather decode living at ~54% with 0% partial + ~1.0x over-fetch is healthy, not a defect.

Worked example (PA decode, gfx942, bs=16, ctx=131072, batch=256): L2 hit 1.7%, 32B 0%, over-fetch 1.04x, 2.85 TB/s = 54% peak. Conclusion: the memory subsystem is clean; there is no KV-load optimization left — verified by also testing block_size 16→64 (regressed +7.8%) and confirming dwordx8 doesn't exist on CDNA3 (dwordx4 / 16B is the max single vector load).

Counter-capture caveat: keep each PMC job to ≤ ~4 TCC counters (single hardware pass). Multi-pass collection has triggered a GPU Hang on gfx942 — see capture-kernel-trace.

MFMA latency reference (cycles = pipeline depth)

MFMA dependency NOPs (CDNA4, from ISA reference Table 38)

These are the minimum independent instructions (or s_nop counts) required between MFMA result production and consumption. The values vary by MFMA variant:

Wait counts for "5, 8, 12, or 20" depend on MFMA variant:

• 5 waits: 16x16 4-block variants (8 cycle MFMAs)
• 8 waits: 16x16x16 F16/BF16 etc. (16 cycle MFMAs)
• 12 waits: 32x32x8, 16x16x4 F32, etc. (32 cycle MFMAs)
• 20 waits: 32x32x4 F32, 32x32x2 F32 (64 cycle MFMAs)

Step 6: Optimization Plan

After running hotspot_analyzer.py --detail, produce a prioritized plan:

Copy code
## Stall Summary
- Total stalls: X cycles (Y% of kernel)
- Top stall type: VMEM-load (Z%)

## Hotspot Analysis

### #1 :LINE  stall=XK (N%)  VMEM-load  stall_rate=92%
Root cause: buffer_load inside QK MFMA loop — only 1 MFMA of hiding time.
Fix: Move all V loads before the QK MFMA loop.
Estimated gain: ~20% kernel cycle reduction (hypothesis; verify correctness and re-time).

### #2 :LINE  stall=XK (N%)  VMEM-load  stall_rate=80%
Root cause: K loads sequential with no compute interleaved.
Fix: Prefetch next tile's K during current tile's MFMA (double-buffer pattern).
See /prefetch-data-load skill.

### #3 ...

## Priority Order
1. [HIGH]  Fix V-load position (24% of all stalls, easy refactor)
2. [HIGH]  K-load cross-tile prefetch (8% of stalls, needs _process_block restructure)
3. [MED]   Move scale loads earlier (8% of stalls, trivial move)
4. [LOW]   Batch LDS reads before PV MFMA (4% of stalls, loop split)

Error Handling