# M3.6 — Nsight trace analysis runbook

> Companion to [`M3.6-decode-throughput-no-retraction-tax.md`](M3.6-decode-throughput-no-retraction-tax.md).
> When a Phase 1/2 trace lands, follow this runbook to read it the same way
> across Claude and Codex sessions — keeps interpretation consistent and
> stops one trace from generating five different conclusions.

## Pre-flight

| Check | Command | Pass condition |
|---|---|---|
| nsight present | `which nsys ncu` | both `/usr/bin/...` |
| ARLE NVTX scaffolding present | `rg "step_mixed_kernel_launch\|step_total" infer/src/scheduler/cuda` | both ranges present (`998bfee` or later) |
| wrapper present | `ls scripts/profile_nsys_guidellm.sh` | exists |
| bench output dir empty for label | `ls bench-output/<date>-<label>/ 2>/dev/null` | doesn't exist (don't overwrite) |

If any FAILS, stop and fix before capturing — partial traces waste hours.

## Capture (M3.6 Phase 1)

ARLE server boot (separate terminal — does NOT need `nsys` wrapping; nsys
attaches by PID):

```bash
RUST_LOG=info NVCC_CCBIN=/usr/bin/g++-14 \
INFER_TILELANG_PYTHON=/home/ckl/projects/arle/.venv/bin/python \
TORCH_CUDA_ARCH_LIST=8.9 \
target/release/infer \
  --model-path /home/ckl/projects/arle/infer/models/Qwen3-4B \
  --port 8000 \
  --max-seq-len 2048 \
  --num-slots 48
```

Wait for `[INFO] HTTP listening on :8000` and run a smoke `curl
http://localhost:8000/v1/completions` first to confirm warmup is past
(autotune/CUDA graph capture can dominate a short trace if you start
during it).

Trace capture (in a third terminal):

```bash
scripts/profile_nsys_guidellm.sh m3.6-arle-s48-highconc \
  --target http://localhost:8000 \
  --bench-profile sweep \
  --concurrencies 64 \
  --max-seconds 30 \
  --warmup 5
```

Output lands in `bench-output/<date>-m3.6-arle-s48-highconc/`:
- `*.nsys-rep` — raw capture (gitignored, 100s MB)
- `*.sqlite` — queryable (also gitignored)
- `nsys_stats.txt` — pre-computed summary (commit-eligible if < 500 KB)
- `summary.md` — wrapper-generated mini-report (commit it)

### vLLM control (Phase 2)

Same wrapper, target the vLLM server (codex's spawn command in
`m6-cuda-vllm-gap-followups.md` §"Phase 1 verify"). Match shape exactly:
1024 in / 256 out, c=64, --max-seconds 30 --warmup 5.

## Read the trace — top-down

Open `bench-output/<...>/nsys_stats.txt` first; only crack the GUI for
unresolved questions.

### Step 1 — overall time distribution (`nsys stats --report cuda_gpu_kern_sum`)

Look for the top 5 kernels by `Total Time`:

| Pattern | Hypothesis it implies |
|---|---|
| Top kernel is `tilelang_decode_attn_*` consuming ≥ 70% | Decode kernel dominant; matches expectation. ARLE-vs-vLLM throughput gap likely lives in this kernel. |
| Top kernel is `*gemv*` or `*gemm*` | Decode attn is fast, projection is slow. Different M_b candidate. |
| Top is a memcpy / memset | Buffer reuse path is broken; investigate before any kernel work. |
| No single kernel > 30% | Many small kernels — launch overhead candidate (Step 3). |

### Step 2 — NVTX phase breakdown (`nsys stats --report nvtx_sum`)

Look for the wall-clock share of M3.6's named ranges:

| Range | Acceptable share at high-conc steady state | Red flag if |
|---|---|---|
| `step_total` | 100% by definition | n/a |
| `step_admission` | < 5% | > 15% — admission CPU-bound |
| `step_plan` | < 3% | > 10% — plan selection logic too heavy |
| `step_mixed_launch_retract` | < 2% (idle most ticks) | > 10% — F2 candidate strengthens |
| `step_decode_kernel_launch` | 60-80% (the bulk) | < 40% — GPU under-utilized |
| `step_mixed_kernel_launch` | workload-dependent; expected non-trivial in high-conc mixed steps | absent or 0% while `/v1/stats` reports many `mixed` plans — NVTX placement is wrong |
| `step_prefill_kernel_launch` | 10-20% | > 30% at c=64 — prefill share too high |
| `step_dispatch_emits` | < 5% | > 15% — emit path leaking time |

Cross-check: `step_decode_kernel_launch` + `step_mixed_kernel_launch` +
`step_prefill_kernel_launch` should explain the GPU-launch-heavy share implied
by `/v1/stats` plan labels during the same window. If they diverge by > 20
percentage points, the NVTX placement is wrong — fix before drawing
conclusions.

### Step 3 — GPU bubbles (`nsys stats --report cuda_api_gpu_sum`)

Look at `Idle Time` per CUDA stream. At kv_util 95% with active=47, the
default decode stream should idle < 5% of wall clock. If it idles more:

- 5-15% idle → CPU dispatch is borderline — `step_admission` /
  `step_dispatch_emits` worth investigating.
- > 15% idle → CPU dispatch is the bottleneck (matches WukLab's
  finding for vLLM pre-v0.4). Compare `osrt` thread time against the
  decode thread to confirm.

### Step 4 — CUDA graph replay sanity

`nsys stats --report cuda_graph_sum` should show ARLE replaying captured
CUDA graphs every step. Each replay should have stable duration (< 5%
variance) — wide variance = capture/replay shape mismatch (a known
regression class, 2026-04 incident).

### Step 5 — per-kernel duration vs vLLM

Cross-reference Step 1 of ARLE trace against Step 1 of vLLM trace at the
same effective batch width. For each kernel ARLE has that vLLM also has
(decode attention, gemv, layernorm, rms_norm), compare `Avg Time`.
Differences > 1.5× are ncu-deep-dive candidates → Phase 4.

## Decision matrix — what does the trace license?

| Trace finding | Recommended fix tier | Don't do |
|---|---|---|
| ARLE decode kernel `Avg Time` > vLLM decode kernel `Avg Time` × 1.5 | Phase 4: `ncu --set full` on the hot kernel; hand off to M3.5 / M_b | Don't change scheduler — kernel work belongs in a different milestone |
| `step_mixed_launch_retract` > 10% AND `cuda_api_gpu_sum` shows wide GPU bubble during retract NVTX | F2 (decode-protected kv_util threshold, see M3.6 Candidate fixes) | Don't go straight to F3 — F2 is the smaller change |
| `step_admission` > 15% | scheduler CPU overhead (likely sched-loop collapse — see M3.5 plan) | Don't add NVTX-based mitigation — fix the algorithm |
| Decode stream idle > 15% AND no NVTX range explains it | OS-level investigation: `osrt` thread time, `cudaStreamSynchronize` waits | Don't blame the kernel |
| `cuda_graph_sum` variance > 5% | CUDA graph capture/replay shape regression | Don't recapture without isolating the offending step |

## Common false positives

- **First 5s of trace** is warmup-dominated (autotune, graph capture).
  Always set `--warmup 5` and analyze ONLY post-warmup window.
- **NVTX ranges from helper threads** can pollute totals. Filter by
  thread ID matching the scheduler's main thread.
- **`cudaLaunchKernel` API time != GPU kernel time**. Confusion between
  API time (CPU side) and kernel time (GPU side) is the most common
  misread. NVTX `step_*_kernel_launch` covers API time; for GPU time
  use `cuda_gpu_kern_sum`.
- **Process spawn overhead** can show up as a 100ms+ block at trace
  start if `nsys` was launched via `nsys profile <cmd>` instead of
  attaching by PID. The wrapper attaches by PID, so this should not
  occur — flag it if you see one.

## Artifact policy

- Raw `.nsys-rep` / `.sqlite` stay in `bench-output/<date>-<label>/`
  (gitignored). Cite the path + sha256 in the wins entry per
  [`docs/bench-and-trace-spec.md`](../bench-and-trace-spec.md) §§8-9.
- `nsys_stats.txt` and the wrapper's `summary.md` MAY be committed if
  < 500 KB.
- Annotated screenshot PNGs (< 500 KB each) MAY be committed alongside
  the wins entry as `bench-output/<date>-<label>/highlight-*.png`.
  Prefer them for findings that are easier shown than described
  (e.g. "see the GPU bubble at t=12.3s").

## Cross-link

- Discovery + falsified F1 evidence:
  [`m6-cuda-vllm-gap-followups.md`](m6-cuda-vllm-gap-followups.md).
- Parent plan with candidate fixes:
  [`M3.6-decode-throughput-no-retraction-tax.md`](M3.6-decode-throughput-no-retraction-tax.md).
- Existing wrapper:
  [`scripts/profile_nsys_guidellm.sh`](../../scripts/profile_nsys_guidellm.sh).
- Spec for trace artifact policy:
  [`docs/bench-and-trace-spec.md`](../bench-and-trace-spec.md).
- WukLab MLSys reference (scheduling overhead can dominate):
  https://mlsys.wuklab.io/posts/scheduling_overhead/.
- vLLM official profiling guide (canonical nsys flags):
  https://docs.vllm.ai/en/stable/contributing/profiling/.
