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:

Sequencing note: for the decode/latency tier, run inference-decode-step-budget 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, 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.

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:

   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:

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):

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:

--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):

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.

# 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

5. Export human-readable CSV

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:

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

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
  kubectl -n $NS cp $TARGET_POD:/profiling/ncu-${kernel}.ncu-rep $BUNDLE/ncu-profiles/ -c $SIDECAR
  kubectl -n $NS cp $TARGET_POD:/profiling/ncu-${kernel}-sol.csv $BUNDLE/ncu-profiles/ -c $SIDECAR
done

7. Populate kernel_internals field in inference_perfbench_v1.json

import json, csv, pathlib
bundle = pathlib.Path("$BUNDLE")
ipb = json.loads((bundle / "inference_perfbench_v1.json").read_text())
ipb["kernel_internals"] = {
    "captured_at": "<UTC>",
    "method": "ncu-sidecar-kernel-scoped",
    "sidecar_image": "ghcr.io/cfregly/nsys-sidecar:0.1.0@sha256:3146de96f6022a8cc36f86d1b8c0281cb940e51e2c3dc49c315646ad66ede43d",
    "vllm_pid": $VLLM_PID,
    "per_kernel": {
        # one entry per kernel captured:
        # "<kernel-name>": {
        #     "achieved_occupancy_pct": ...,
        #     "theoretical_occupancy_pct": ...,
        #     "dram_throughput_pct_peak": ...,
        #     "sm_throughput_pct_peak": ...,
        #     "regs_per_thread": ...,
        #     "smem_per_block_bytes": ...,
        #     "block_limit_factor": "registers" | "shared_mem" | "warps",
        # }
    },
}
(bundle / "inference_perfbench_v1.json").write_text(json.dumps(ipb, indent=2))

The inference_perfbench_v1 schema treats kernel_internals as optional - bundles without it are unaffected.

8. Cleanup

kubectl -n $NS exec $TARGET_POD -c $SIDECAR -- rm -rf /profiling/ncu-*
# The sidecar container will terminate after the 3600s sleep; pod
# auto-cleans the ephemeral container.

Cross-skill join

Roofline interpretation is most useful when joined with the absolute kernel-time from inference-kernel-profile (nsys) and the relative sample-share from zymtrace-anchored-query. Example join:

Kernel                                  zymtrace samples %  nsys total ms  ncu DRAM %  ncu SM %  Verdict
flashinfer trtllm_allreduce_fusion      14.2%               1,847          88%          12%      Memory-bound; comm-bound (good - kernel is at-the-roof)
fmhaSm100fKernel...Persistent           8.7%                1,144          47%          81%      Compute-bound (room for compute-side tuning)
multimem_all_reduce_kernel              5.3%                694            91%           9%      Memory-bound; comm-bound (at-the-roof)
bmm_E2m1_E2m1E2m1                       6.9%                901            61%          74%      Mixed (canonical NVFP4 BMM: balanced)

Without ncu, you have shares (zymtrace) + times (nsys) but no attribution to the dimension (memory vs compute vs occupancy). ncu closes that loop.

Cost + risk

• ~10-100x per-replayed-kernel slowdown during capture. Bounded by --launch-count (5 replays at 100x = 0.5s of kernel time = sub-second for typical kernels)
• ncu-rep files are typically 5-50 MB per kernel (manageable)
• Sidecar attach is non-disruptive (no vllm restart)
• Live-pod risk: see Safety section above

Pair this skill with evidence-bundle-init when starting a new investigation so the ncu-profiles/ subdir lands in the right place.

Full-context reporting (no bare numbers)

This skill is the preferred follow-up for the proper roofline scatter. Per the canon "Every performance number carries its full context (no bare numbers)" (docs/METHODOLOGY.md "Full-context reporting"): every number this skill emits MUST carry its full measurement-context descriptor, and every comparison MUST be matched on it. A bare number is a defect - it cannot set a default, ship a config, or appear in a report.

• Identity: model (+HF path), hardware (exact ceiling token GB300/B200), quant, kv-cache dtype.
• Parallelism: TP, DP (replicas), PP, EP, parallel_strategy.
• Serving cfg: max-num-seqs, max-num-batched-tokens, gpu-memory-utilization, max-model-len, cudagraph_mode/enforce_eager, async_scheduling, prefix-caching.
• Workload: dataset, ISL/OSL (or mean in/out tokens), concurrency, num-prompts.
• Regime: warm vs cold. Latency vs throughput tier.
• Stack: image/vllm commit, bench backend, serving engine.
• Grounding: %SoL (+ ceiling key from configs/sol-ceilings.yaml - never inline a peak), sol_rigor (L1-L4), trials n (mean±std), same-node, baseline named.
• Per-number exact shape (no smoothing): when reporting more than one number, keep EACH with its own exact shape (ISL/OSL, concurrency, dataset, regime) - never normalize a set to one uniform descriptor that hides per-point variation (e.g. c=1 @ ISL1024/OSL256 + c=64 @ ISL4096/OSL512, NOT one shared "random").

This skill is the preferred path to a real per-kernel arithmetic-intensity-vs-roofline scatter plot per docs/METHODOLOGY.md "Speed-of-light framing". The other two siblings (zymtrace-anchored-query, inference-kernel-profile) deliver time-share + wall-clock duration but cannot compute %SoL per kernel because they lack the FLOPS + bytes counters this skill captures (sm__sass_thread_inst_executed_op_*_pred_on.sum, dram__bytes.sum).

Workflow when proper %SoL is needed:

1. Identify the dominant hot kernel via zymtrace-anchored-query + inference-kernel-profile.
2. Run THIS skill against that kernel name with the roofline-min counter set (see Reference below).
3. Compute arithmetic intensity = FLOPS / bytes. Place the point on the B200 / GB300 / H100 roofline using peaks from configs/sol-ceilings.yaml (b200_sm100.nvfp4_dense_pflops for compute ceiling, b200_sm100.hbm3e_tbps for bandwidth ceiling). Cite by key path, never inline magic numbers.
4. Record the resulting %SoL in the per-run sol-summary.md doc for the campaign / cluster probe this kernel came from.

Next lever / BREAKTHROUGH (Grind Mandate)

If this skill emits a measured result, its output MUST end by naming the next perf lever, its expected unlock (direction + rough magnitude), and the gate that proves/refutes it, per docs/METHODOLOGY.md "Always be grinding (next-lever framing)". A measured win is the new floor, not the finish -- so do everything we can to find the next BREAKTHROUGH: the highest-EV unlock toward Speed-of-Light (a new champion / kernel / router / quant / parallelism / spec-decode win, or an unblocked stack), not just the next micro-lever. Rank the candidate breakthrough levers by value x cost (the GRIND FRONTIER, perftunereport value_view), pursue the top, bank the rest with evidence. Record WHY a refuted lever loses, update the standing frontier in the active bundle's HANDOFF.md. Never conclude "exhausted/optimal/done" without an explicit next-lever frontier (an empty frontier AND a documented SoL wall only). Delete this section ONLY if the skill produces no measurements.

Wiring ncu into the perf-report page-5 roofline scatter

The capture's ncu-profiles/*-sol.csv + *-raw.csv are ingested into a perf-report campaign cell as cells/<cell-id>/ncu_kernels.json (the page-5 input) via the import_ncu verb:

perftunereport import_ncu --campaign <slug> --cell-id <cell-id> --bundle <bundle-dir>
perftunereport report_render --campaign <slug>   # -> page 5 populated

Run these checked-in import commands explicitly after the capture. Notes:

• Two CSV shapes are auto-handled. ncu's --page raw is wide (one row per kernel, metric columns). Ncu-2026's --page details --section SpeedOfLight is long/melted (Metric Name / Metric Value rows). The importer detects and pivots the long shape, so either export works.
• Counter set determines what page 5 can show. --set=full or --roofline-min collect dram__bytes.sum + the sm__sass_thread_inst_executed_op_*_pred_on.sum FLOP op-counts, so the importer computes arithmetic intensity + achieved TFLOPS and page 5 plots a real roofline point. --set=basic collects neither, so AI is null and page 5 renders a hollow "%SoL only / AI unmeasured" marker placed at SM% x the kernel family's compute ceiling (category_ceiling_map in configs/sol-ceilings.yaml) - the y is honest, the x is a placeholder. AI is never fabricated.
• SM-busy% is NOT %-of-FLOP-SoL - never conclude "at-roof" from --set basic. The --set basic sm__throughput.avg.pct_of_peak ("SM busy") can read 88-92% while the kernel is at <15% of its FLOP-roofline - a persistent / split-K / spinning kernel keeps SMs busy doing little tensor work. To decide at-roof vs headroom you MUST run --set full and compare achieved TFLOPS to the FLOP ceiling (fp8_dense_pflops / nvfp4_dense_pflops / bf16_dense_pflops). WORKED EXAMPLE: the DeepSeek-V4-Flash FP8 GEMM (sm100_fp8_fp4_gemm_1d1d, ~70% of decode compute) read 88-92% SM (--set basic, which mislabeled it "at-roof") but --set full showed only 1-14.5% of the FP8 FLOP-SoL (M=1→256) - the kernel is concurrency-starved, not exhausted. The throughput tier was the real headroom, not a kernel rewrite.
• TP=8 NVFP4 blocker. Multi-pass kernel-replay fails at TP=8 NVFP4 (ContextSaveFailed). See the replay-mode-application runbook in the sister-deploy scaffolding above for the --replay-mode application path that gets to a full AI-grounded point.

Kernel rubric (K/R/H/P/A) - this skill is the H + P enforcement point

When this capture backs a custom-kernel comparison (a candidate kernel vs a baseline), apply the kernel rubric (see docs/METHODOLOGY.md "Kernel-work classification"). ncu is the only profiler in the sibling set that can prove the two axes a kernel win actually turns on:

• H (hardware specialization) - the Compute (SM) Throughput [%] and tensor-pipe active % (from the SpeedOfLight section) reveal whether a kernel engages the frontier path. On SM100 the production libraries are H4 (sm100f tensor cores / NVFP4 tensor cores via bmm_E2m1, nvjet_tst_* in the canonical 5-kernel set). A candidate kernel that shows near-zero tensor-pipe activity / SM% dominated by FMA is H1 - and an H1 candidate cannot beat an H3-H4 baseline on a K3-K4 op no matter how it schedules. Read tensor-core engagement off ncu BEFORE believing any win.
• P (performance target) - the roofline (%SoL via the FLOPS + dram__bytes.sum counters, "Speed-of-light reporting" above) is the P-axis: P4 means at/above the best library. A kernel far below its family ceiling that still "wins" an e2e A/B is winning on something other than GPU work (re-check methodology).

The gate, operationally: record (K,R,H,P,A) for the candidate AND the named baseline in the bundle's SOURCE.md/summary.md. A speedup over a baseline at strictly lower H or R (e.g. a tensor-core candidate vs a generic-Triton baseline when production runs the sm100f library) is a DRAFT, never a VERDICT - it fails the "Fair baseline" clause. The canonical worked failure: warp-decode (K4/R2/H1) beat generic Triton in microbench but ncu/zymtrace confirmed it never engaged tensor cores, and it lost 1.51x to FlashInfer-TRTLLM (K4/R1/H4) on real GPU time.

Emit it as a structured krhpa: block (not just prose). When the campaign renders an L4 ncu roofline (page 5, sol_rigor=L4), perftunereport publish_to_lake FAILS CLOSED unless the campaign config.yaml carries a krhpa: block classifying both arms - prose in SOURCE.md/summary.md alone does not satisfy the gate for an L4 campaign. Add to the campaign config:

krhpa:
  candidate: {K: 4, R: 2, H: 1, P: 2, A: 1, name: "warp-decode (Triton FMA)"}
  baseline:  {K: 4, R: 1, H: 4, P: 4, A: 1, name: "FlashInfer-TRTLLM bmm_*_sm100f"}

Each axis is an int 1..4 (= L1..L4), baseline.name must name the production-representative kernel/library. The gate (lake_writer.krhpa_problems) refuses a missing or malformed block under --strict and records + warns otherwise.

Reference: full hardware-counter set names

ncu's --set full enables ~60 metric groups. For roofline analysis the minimum set is:

sm__sass_thread_inst_executed_op_*_pred_on.sum  (per-op work counts)
dram__bytes.sum                                  (DRAM bytes read+written)
gpc__cycles_elapsed.max                          (peak cycles)
smsp__warps_active.avg.pct_of_peak_sustained_active  (occupancy)
launch__registers_per_thread                     (regs/thread)
launch__shared_mem_per_block_static              (smem/block)

These are also available under the named section group --section SpeedOfLight_HierarchicalTensorRooflineChart for direct roofline plotting.