inference-kernel-ncu-profile

# inference-kernel-ncu-profile ## Purpose Capture per-kernel **hardware-counter** data from a **live vLLM inference pod** without rebuilding the production image. Uses `kubectl debug --share-processes` to attach an `ncu`-enabled sidecar to an existing pod. The sidecar replays a small number of launches of each named kernel with the hardware-counter set enabled and writes `.ncu-rep` files into a sidecar emptyDir that you extract via `kubectl cp`. The three sibling skills cover complementary kernel-profile dimensions: | Question | Tool | Skill | |---|---|---| | What kernels run + relative sample-share? | zymtrace | `zymtrace-anchored-query` | | Absolute per-kernel duration (ns), CUDA graph timeline, NVTX ranges | nsys | `inference-kernel-profile` | | **Per-kernel occupancy / regs / smem / DRAM-BW / arithmetic intensity / warp stalls** | **ncu** | **this skill** | | c=1 decode-step budget (GPU-busy vs host-gap) + kernel-vs-host-vs-comm verdict | vLLM profiler endpoint / nsys | `inference-decode-step-budget` | **Sequencing note:** for the decode/latency tier, run [`inference-decode-step-budget`](/plugins/profile-and-optimize/skills/inference-decode-step-budget/SKILL.md) FIRST. If at c=1 it returns "host-bound" (GPU idle >> busy per step, as GLM-5.1 does), per-kernel SoL is **moot** - the kernels are a small fraction of TPOT and are launch-latency-scale at one token. Only escalate to this skill when the budget says kernel-bound, or for the prefill/throughput tier. ## When to use - Investigating WHY a specific kernel is hot. The `zymtrace-anchored-query` + `inference-kernel-profile` skills tell you WHICH kernels are hot and HOW LONG they take. Ncu tells you whether they're memory-bound vs compute-bound, whether occupancy is the ceiling, whether warps are stalling on memory or sync, etc. - Roofline analysis to prove a kernel is at-the-roof vs has-room. - Pre-PR-submission justification: when proposing a custom kernel or a kernel-tuning change, ncu evidence (e.g. "this kernel achieves only 47% theoretical occupancy due to register pressure") is the kind of argument upstream reviewers expect. Do **not** use this skill for: - Continuous always-on capture - kernel replay is 10-100x slow per replayed kernel. Ncu is a per-shot tool. - Untargeted "give me everything" sweeps - without `--kernel-name` scoping, ncu will replay EVERY kernel, which on a busy vLLM pod is a heavy operation that risks Kubelet liveness-probe failures. - Multi-pod fleet capture - this skill targets ONE pod at a time. ## Sidecar image Same image as `inference-kernel-profile`: **`ghcr.io/cfregly/nsys-sidecar:0.1.0`** (publicly readable). Image includes nsys 2025.6.3 + **ncu 2026.1.1** + py-spy 0.4.1 + python3 + jq. Image is 11.87 GB. First-pull on a new GPU node is ~3-7 min from ghcr.io. Subsequent attaches on the same node are instant. ## Safety: live-serving impact ncu's `--set full` capture flag enables every hardware-counter group, which forces kernel replay (the kernel is launched, ncu reads the counters, the kernel is launched AGAIN, ncu reads the next group, ...). Replay multiplies the kernel's wall-clock by 10-100x. **On a single-replica deploy**, the replay slowdown can spike pod CPU to the point of Kubelet liveness-probe failure. Mitigate by: - Using `--launch-count 3` not 5 - Using `--set basic` instead of `--set full` for the first capture (basic captures only the metric groups required for roofline, ~3x slowdown instead of 30x) - Scoping to ONE kernel per capture, not 5 **On a multi-replica deploy**, a temporary per-replica capacity dip during ncu replay is usually acceptable, so `--launch-count 5` + `--set full` is safe - the slow replica is just 1 of N and the Service load-balancer distributes around it. **Fallback**: if even the conservative path is too risky, spin a sister deploy `<deploy>-ncu` at 0 service routing (a separate Helm release serving the same model, NOT in the Service selector), and ncu against THAT pod. Adds ~30 min for the sister helm install, but zero risk to production serving. **Memory caveat:** at `gpu-mem-util >= 0.85` there is little headroom for a profiler's host/device buffers. A long capture window (observed with a 40s nsys window on a GLM-5.1 canary) can OOM-crash vllm mid-capture and truncate the report. Keep windows short, scope ncu to one kernel + low `--launch-count`, and prefer the endpoint-triggered bounded windows in [`inference-decode-step-budget`](/plugins/profile-and-optimize/skills/inference-decode-step-budget/SKILL.md). ## Gate 0 (precondition): CUPTI must even initialize (CUDA image-vs-driver skew) ncu is a CUPTI client, so on GB300 nodes it can record **0 kernels** for the same reason nsys + the vLLM torch profiler do: a CUDA 12.9 serve image against the node's CUDA 13.0 driver makes CUPTI fail to init (`CUPTI_ERROR_INVALID_DEVICE`). This is NOT permission (`RmProfilingAdminOnly:0`), NOT a missing lib (`libcupti.so.12` present + linked), and NOT a wrong-process attach - it is a CUDA major-version toolkit-vs-driver skew. If ncu returns 0 kernels, grep the launch log for `CUDA versions. CUPTI/Runtime/Driver`: a 12.x-toolkit / 13.x-driver split needs a CUDA-13-aligned image or zymtrace (non-CUPTI), not more attach tuning. **SOLVED PATH (GB300 default): pull a Blackwell-capable ncu from NGC** instead of using the in-image ncu: copy ncu **2026.x** from `nvcr.io/nvidia/pytorch:26.05-py3` (arm64/sbsa) into the serve container via an initContainer (no prod rebuild. Needs an NGC image-pull secret in the namespace). Use this NGC-tools path on GB300 for both nsys and ncu, the legacy in-image / apt paths are superseded there. Verified end-to-end on a GB300 MoE deploy (per-kernel Compute(SM)/Occupancy resolved). ## ncu capture-hygiene (empty-rep / non-deterministic gates) - empty != tooling limit Beyond Gate 0 (CUPTI), three ncu traps on a vLLM stack each yield "No kernels were profiled" or a non-deterministic-count failure. NONE is a tooling limit (each reproduced + fixed on a DeepSeek-V4-Flash GB300 deploy). Canon: `docs/METHODOLOGY.md` "Capture hygiene". 1. **`cudaProfilerStart` breaks under any vLLM `direct_register_custom_op` import.** If your harness imports a module that runs `direct_register_custom_op` at load (`mhc.tilelang`, `fp8_utils`, …), ncu `--profile-from-start off` + an in-script `torch.cuda.cudart().cudaProfilerStart()` captures **0 kernels** (the custom-op/triton registration perturbs the CUPTI profiler-start hook). FIX: do NOT rely on cudaProfilerStart - profile from start and exclude warmup with `--launch-skip <warmup> --launch-count N` + `--kernel-name` scoping: ```bash M=256 N=4096 K=8192 ncu --kernel-name 'regex:sm100_fp8_fp4_gemm' \ --launch-skip 3 --launch-count 1 --set full --target-processes application-only \ --export out -- python3 one_op.py # one_op.py warms the op 3x then loops; ncu skips the 3 warmups ``` 2. **`-- env VAR=v python3` defeats `--target-processes application-only`** - ncu profiles the `env` process, not its python child → 0 kernels. FIX: set the env on the ncu process itself (`VAR=v ncu … -- python3 …`), never via an `env` wrapper after `--`. 3. **TP=N MoE+sparse in-engine `--replay-mode application` is non-deterministic.** A large MoE + sparse-attention model relaunched per metric-pass varies its kernel-launch count → `==ERROR== Unexpected number of profiled kernels`, `--app-replay-mode relaxed` + no-op `profile_run` + prefix-caching-off + eager + fixed-blocks do NOT fix it. The canonical L4 path is an **isolated single-op kernel-replay harness**: call the production library op directly (e.g. `vllm.utils.deep_gemm.fp8_gemm_nt` / `tf32_hc_prenorm_gemm`), TP=1, **NO `mhc.tilelang` import**, synthetic fixed-shape inputs (AI/%SoL is shape-structural). Mirror the production caller's quant + SF layout (e.g. `per_token_group_quant_fp8` + `per_block_cast_to_fp8` + `deepgemm_post_process_fp8_weight_block`) so the exact kernel template fires. ## Two blockers on prod-direct attach An older `ncu --attach <PID>` recipe attached ncu to a running vllm process directly. That recipe is **broken** on current production deploys for two independent reasons. Read this section before attempting any live capture. ### Blocker 1: zymtrace already owns `CUDA_INJECTION64_PATH` Production vllm deploys instrumented with zymtrace run with: ``` CUDA_INJECTION64_PATH=/var/lib/zymtrace/profiler/libzymtracecudaprofiler.so ``` This is the zymtrace profiler agent injecting itself as the CUDA driver-level injection library. The CUDA driver only honors **one** injection library per process. Ncu cannot stack on top. There's no way to add ncu's `libcompute.so` to a running zymtrace-instrumented process. **Exception: latency/canary deploys often have NO zymtrace injection.** A canary that runs with no `CUDA_INJECTION64_PATH` in its env escapes Blocker 1 - ncu launch-mode (or the preload+TCP attach below) works directly on it without a separate sister-deploy. Check the target's env (`kubectl get deploy <name> -o jsonpath='{...env...}'`) before assuming a sister is required. The sister-deploy is only mandatory for the zymtrace-instrumented production releases. ### Blocker 2: ncu 2026.1.1 dropped `--attach <PID>` The deprecated `--attach <PID>` flag was removed in ncu 2026.x: ``` ==ERROR== unrecognised option '--attach'. Use --help for further details. ``` The modern ncu interaction modes (per `ncu --help`) are: - `--mode launch-and-attach` (default) - wrap the application at process startup - `--mode launch` - launch the application and suspend for a later attach - `--mode attach` - connect to a process pre-instrumented with ncu's injection library, via TCP (`--hostname` + `--port`) The only path that works against an **already-running** vllm process is `--mode attach`, and it requires `NV_NCU_INJECTION64_PATH` to have been set when vllm started so the process exposes the ncu profiling endpoint. Production vllm (instrumented with zymtrace instead) does not expose this endpoint. ## Recipe: sister-deploy with preload + TCP attach Both blockers are dodged by spinning up a parallel **sister-deploy** helm release running the same vllm image but with: - zymtrace `CUDA_INJECTION64_PATH` unset - `NV_NCU_INJECTION64_PATH` pointing at ncu's `libcompute.so` (mounted from the nsys-sidecar image via an initContainer + emptyDir volume) - `replicas: 1` - a Service selector / label that **excludes** prod traffic (the prod Service matches `app: basic-inference`, the sister uses `app: basic-inference-ncu`, so prod LB routes around it) The vllm process on the sister pod exposes ncu's TCP profiling endpoint on `127.0.0.1:49152`. From a `kubectl debug` sidecar with `--share-processes`, run: ```bash ncu --mode attach --hostname 127.0.0.1 --port 49152 \ --kernel-name '<KERNEL_REGEX>' \ --launch-count 5 \ --set full \ --target-processes all \ --export /profiling/ncu-$K.ncu-rep ``` Multiple captures against the same long-lived sister pod are supported - vllm stays warm. The sidecar runs ncu N times sequentially. ### Reference helm + capture scaffolding A reusable sister-deploy scaffold (one directory per model) contains: - `values-ncu-sister.yaml` - helm values overrides - `install.sh` - `helm upgrade --install` wrapper - `capture.sh` - full capture flow (wait-ready, warmup, sidecar attach, N-kernel loop, extract, helm uninstall). Flags: `--launch-count N`, `--set basic|full`, `--roofline-min` (AI-roofline `--metrics` set only), `--replay-mode kernel|application|range`, and `--campaign <slug> --cell-id <cell>` to auto-import the result into a perf-report campaign cell. - `README.md` - runbook + canonical 5-kernel list - a replay-mode-application runbook - the ack-gated path past the TP=8 NVFP4 `ContextSaveFailed` kernel-replay blocker (use `--replay-mode application --roofline-min`). ### Identifying hot kernels ncu `--kernel-name` accepts substring or regex match against the demangled or mangled kernel name. Demangle a mangled symbol via `c++filt` (available in the nsys-sidecar image): ```bash echo "_ZN66_GLOBAL__N__51f43292_25_CUDASymmetricMemoryOps_cu_c81adf72_277485226multimem_all_reduce_kernelIN3c108BFloat16ELi16EEEvPT_mPPjmm" | c++filt # -> (anonymous namespace)::multimem_all_reduce_kernel<c10::BFloat16, 16>(...) ``` Either form is fine for `--kernel-name`: ```bash --kernel-name multimem_all_reduce_kernel --kernel-name "regex:multimem_all_reduce.*BFloat16" ``` Canonical 5-kernel set for B200 NVFP4 inference (Kimi K2.6, GLM-5.1, DeepSeek): | Kernel substring | Category | Expected bottleneck | | --- | --- | --- | | `multimem_all_reduce_kernel` | NCCL | NVLink5 bandwidth | | `bmm_E2m1` | BMM-NVFP4 | NVFP4 Tensor Core compute | | `fmhaSm100fKernel` | FMHA | HBM3e bandwidth | | `nvjet_tst_.*_2cta` | MoE / GEMM | NVFP4 Tensor Core compute | | `triton_red_fused` | Triton-fused | Mixed | ## Recipe (deprecated - DO NOT USE on prod-direct) An earlier recipe used `ncu --attach $VLLM_PID` against the zymtrace-instrumented production pod. That recipe is preserved below ONLY as historical reference. It WILL FAIL on every current inference deploy. Use the sister-deploy section above. <details> <summary>Historical recipe (broken. Click to expand)</summary> ```bash # THIS WILL FAIL: # ERROR: unrecognised option '--attach' (ncu 2026.1.1) # ERROR: CUDA injection slot taken by zymtrace ncu --kernel-name '$KERNEL' --launch-count 5 --set full \ --export /profiling/ncu-$KERNEL.ncu-rep \ --attach $VLLM_PID ``` </details> ### 5. Export human-readable CSV ```bash kubectl -n $NS exec $TARGET_POD -c $SIDECAR -- bash -c " ncu --import /profiling/ncu-${KERNEL}.ncu-rep \ --csv --page raw > /profiling/ncu-${KERNEL}-raw.csv ncu --import /profiling/ncu-${KERNEL}.ncu-rep \ --csv --page details --section SpeedOfLight \ > /profiling/ncu-${KERNEL}-sol.csv " ``` Key columns from `--page details --section SpeedOfLight`: | Column | Meaning | |---|---| | `Memory Throughput [%]` | DRAM throughput vs theoretical peak | | `DRAM Throughput [%]` | Same, restricted to DRAM (vs L2) | | `Compute (SM) Throughput [%]` | SM utilization | | `Achieved Occupancy` | Actual warps-active / theoretical-max-warps-active | | `Theoretical Occupancy` | What the kernel COULD achieve given regs+smem | | `Achieved Active Warps Per SM` | Average warps active per SM | | `Block Limit Registers` | Limit imposed by registers/thread | | `Block Limit Shared Mem` | Limit imposed by smem/block | | `Block Limit Warps` | Limit imposed by max warps per SM | These are the roofline + occupancy diagnostics. A kernel with