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`](../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. Per the [`Dean-Ghemawat performance-hints adaptation`](../../server/docs/performance-hints.md), bound the candidate win by the kernel's measured end-to-end share before capture. Hardware-counter improvements cannot recover time outside that share. ## 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@sha256:3146de96f6022a8cc36f86d1b8c0281cb940e51e2c3dc49c315646ad66ede43d`** (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`](../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. ### Operator-owned deployment and capture flow This repository does not ship model deployment manifests. Keep the sister pod and Helm values in the workload repository that owns the deployment. A safe operator flow must wait for readiness, warm the model, attach the sidecar, capture the selected kernels, extract the results, and remove the temporary release. Record the launch count, section set, replay mode, campaign, cell ID, deployment revision, and cleanup receipt in the evidence bundle. For the TP=8 NVFP4 `ContextSaveFailed` kernel-replay blocker, use application replay with the minimum roofline metrics and the acknowledgement required by your cluster workflow. ### 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 `DRAM Throughput % = 92%` is memory-bound. One with `Compute SM Throughput % = 88%` is compute-bound. One with `Achieved Occupancy = 0.31` and `Block Limit Registers = limited` is suffering register pressure. ### 6. Extract artifacts ```bash BUNDLE=experiments/artifacts/inference-perf-bench/<bundle>/ mkdir -p $BUNDLE/ncu-profiles for kernel in multimem_all_reduce_kernel flashinfer_trtllm_allreduce_fusion fmhaSm100fKernel concat_and_cache_mla_kernel bmm_E2m1; do