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 programmatically via `hotspot_analyzer.py` + `code.json`. 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*. ## Arguments | Argument | Description | |----------|-------------| | `<CMD>` | Command to profile. Example: `python bench_pa.py --batch 32` | | `--dir <path>` | Skip collection; analyze existing `ui_output_agent_*_dispatch_*` directory | | `--topk N` | Show top-N hotspots (default: 15) | --- ## 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: ```bash # Both analyzers ship with this skill -- do not reimplement them. python ${CLAUDE_SKILL_DIR}/scripts/hotspot_analyzer.py <dispatch_dir> --topk 15 --mode both python ${CLAUDE_SKILL_DIR}/scripts/hotspot_analyzer.py <dispatch_dir> --topk 5 --mode src --detail --context 4 ``` Skip to **Step 5: Interpret Results**. --- ### Mode B: Full collection workflow #### Step 1: Kernel Discovery ```bash 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: | Rank | Kernel Name | Calls | Total (us) | Avg (us) | % GPU Time | |------|-------------|-------|------------|----------|------------| Ask the user which kernel to trace if not obvious. **Prefer `results.db`** if available — use sqlite3 for structured queries: ```bash 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: ```yaml 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 ```bash 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: ```bash 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:** ``` 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: ```bash # 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: ``` [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: ```json {"/": {"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 | Type | Instructions | Root Cause | |------|-------------|------------| | `VMEM-load` | `buffer_load_*`, `global_load_*` | Load itself stalled (VMEM queue full or back-pressure from no compute to hide behind) | | `VMEM-wait` | `s_waitcnt vmcnt(N)` | Waiting for outstanding VMEM loads to complete | | `LDS/SMEM-wait` | `s_waitcnt lgkmcnt(N)` | Waiting for LDS or SMEM ops | | `barrier` | `s_barrier` | Cross-wave sync — slowest wave dominates | | `MFMA/FMA` | `v_mfma_*` | MFMA dependency chain (RAW hazard) | | `LDS` | `ds_read_*`, `ds_write_*` | LDS access latency | ### 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%) ```python # 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 ```python # 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 ```python # 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 ```python # 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**: ``` 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.) | Property | CDNA3 (gfx942) | CDNA4 (gfx950) | |---|---|---| | VGPR pool | 512 combined (256 arch + 256 accum, unified budget) | 512 combined (same) | | Occupancy formula (VGPR) | `512 / (arch_alloc + accum_alloc)` | `512 / (arch_alloc + accum_alloc)` | | Alloc granularity | 8 VGPRs | 8 VGPRs | | LDS size | 64 KB | 160 KB | | LDS alloc block | 256 bytes | 1280 bytes | | VMCNT width | 6 bits (max 63 in-flight) | 6 bits (max 63 in-flight) | | LGKMCNT width | 4 bits (max 15 in-flight) | 4 bits (max 15 in-flight) | 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 /