第 15 章 · GPU Profiling 与 LLM 推理优化实战

前 14 章已经把 CUDA、GEMM、softmax、attention、KV cache 和一个 mini GPT-2 推理引擎串起来。 这一章补上工业环境里真正每天会用的能力：拿到一个慢 kernel 或一条慢的 decode path，如何用 GPU 性能分析 (GPU Profiling) 证明瓶颈、提出改动、验证收益，并把结论沉淀成可复现 benchmark。

学习目标

• 建立 LLM 推理 profiling 的固定工作流：先 timeline，再 kernel，再 roofline，再代码改动。
• 读懂 Nsight Systems / Nsight Compute 的核心指标：launch gap、SM throughput、DRAM throughput、occupancy、stall reason、register spilling。
• 会用 NVTX 与 cudaProfilerStart/Stop 把 profile 范围缩到 prefill、decode、sampler 或某个 layer。
• 掌握 decode GEMV、KV cache layout、paged cache、量化权重读取这些 LLM infra 热点的 profiling 方法。
• 能分析 prefill 与 decode 混跑时的 head-of-line blocking、chunked prefill 和 stream priority 边界。
• 能分析 continuous batching 里的 ragged batch、活跃请求压缩和 padding 浪费。
• 能分析 MoE expert routing 的 dispatch/gather、capacity padding 和 expert load imbalance。
• 能分析 host/device staging 里的 pageable/pinned buffer、H2D/D2H copy 和 copy/compute overlap。
• 能分析 residual、RMSNorm、activation quant 这类小算子是否应该做 kernel fusion。
• 能按主流 NVIDIA GPU 的显存、带宽、Tensor Core 精度和功耗选择推理部署形态。

前置知识

• 第 5 章：global memory coalescing、cache、HBM 带宽。
• 第 7 章：warp reduce / block reduce。
• 第 8 章：Nsight Systems / Compute 基础、streams、CUDA Graph。
• 第 13 章：RoPE、SwiGLU、KV cache、sampling。
• 第 14 章：prefill vs decode、mini-LLM forward loop。

核心概念

15.1 工业 profiling 工作流

不要一上来就改 kernel。生产里最可靠的顺序是：

graph TD
    A["定义问题
p50 / p99 / tokens/s / GPU util"] --> B["Nsight Systems
看 timeline 与空洞"]
    B --> C{"GPU 是否忙?"}
    C -->|"不忙"| D["查 CPU 调度、launch gap、同步、allocator、网络"]
    C -->|"忙"| E["Nsight Compute
钻热 kernel"]
    E --> F{"roofline 落点"}
    F -->|"memory-bound"| G["降字节: quant / cache layout / fusion / coalescing"]
    F -->|"compute-bound"| H["升算力: Tensor Core / tile / occupancy / pipeline"]
    G --> I["微基准复现
固定 shape + seed + clock"]
    H --> I
    I --> J["端到端回归
FTL / TPOT / TTFT / p99"]
    style B fill:#f3f1e8,stroke:#1a4d80
    style E fill:#f3f1e8,stroke:#8b1538
    style J fill:#ecf3eb,stroke:#2f5d3a
    

LLM 服务里，常见的目标指标不是“某 kernel 快了多少”，而是： TTFT（time to first token，首 token 延迟）、 TPOT（time per output token，decode 单 token 延迟）、 tokens/s/GPU、p95/p99 latency、HBM 占用和功耗。 kernel 优化必须最终回到这些指标上，否则局部 kernel 的明显变化可能只带来很小的端到端收益。

15.2 工具分层：每个工具回答一个问题

Nsight Compute 官方 Profiling Guide 说明了 section set 的成本差异： --set full 会采集很多 section，可能触发 kernel replay；生产复现时先用 SpeedOfLight、MemoryWorkloadAnalysis、Occupancy 这几个定向 section。 参考：Nsight Compute Profiling Guide。

15.3 LLM kernel 指标速查

15.4 Prefill 与 Decode 要分开 profile

同一个模型 forward，prefill 和 decode 的热点完全不同：

实战中把 profile 文件命名成 model_gpu_phase_shape_commit，例如 qwen2_7b_h100_decode_b16_t1_20260715.nsys-rep。同事拿到文件后能立刻知道它代表什么， 也方便后续做性能回归。

关键代码

15.5 NVTX + cudaProfilerStart/Stop

profiling_ranges.cu 用两个 kernel 模拟 LLM 推理里的两类热点：memory-bound 的向量读写、compute-bound 的 FMA 链。 关键点是把 profiler 采集范围缩小到热区。

CUDA_CHECK(cudaProfilerStart());

{
    NvtxRange range("decode/memory_bound_saxpy");
    for (int r = 0; r < repeats; ++r) {
        saxpy_kernel<<<grid, block>>>(a, d_x.ptr, d_y.ptr, n);
    }
    CUDA_CHECK(cudaDeviceSynchronize());
}

{
    NvtxRange range("decode/compute_bound_fma");
    fma_chain_kernel<<<grid, block>>>(d_x.ptr, d_y.ptr, n, iters);
    CUDA_CHECK(cudaDeviceSynchronize());
}

CUDA_CHECK(cudaProfilerStop());

Nsight Systems 的 --capture-range=cudaProfilerApi 会等到 cudaProfilerStart() 才开始采集；Nsight Compute range replay 也支持 profiler API 和 NVTX range。 CUDA Runtime API 文档明确说明 cudaProfilerStart/Stop 用来控制采集粒度： Profiler Control。

15.6 Chunked prefill：别让长 prefill 挡住 decode

chunked_prefill_profile.cu 模拟线上服务里最常见的调度冲突：一个长 prefill 请求占住 GPU，旁边已有请求的 decode token step 被迫等待， 形成 队头阻塞 (head-of-line blocking)。 示例对比三种路径：monolithic prefill 跑完再 decode、把 prefill 切成 chunk 后在 chunk 间插入 decode、 以及低优先级 prefill stream + 高优先级 decode stream。它不试图复刻完整 scheduler，而是把 分块 prefill (chunked prefill)、first-token proxy、 流优先级 (stream priority) 和 kernel boundary 的关系拆开观察。

// Bad for latency: a long prefill kernel sits before the first decode step.
prefill_work_kernel<<<prefill_grid, 256, 0, stream>>>(prefill_in, prefill_out, n, iters);
decode_step_kernel<<<decode_grid, 256, 0, stream>>>(decode_state, tokens, n, 0, iters);

// Better for serving: create scheduling boundaries, then let decode run between chunks.
for (int c = 0; c < chunks; ++c) {
    prefill_work_kernel<<<chunk_grid, 256, 0, stream>>>(chunk_in, chunk_out, chunk_n, iters);
    decode_step_kernel<<<decode_grid, 256, 0, stream>>>(decode_state, tokens, n, step++, iters);
}

工业排查时不要只报告平均 tokens/s。要把长 prompt prefill、短 prompt prefill、decode-only、spec verify 分桶采集，分别记录 TTFT、TPOT、p95/p99 和 active request 数。Nsight Systems 里重点看 NVTX range： 如果 prefill 大 kernel 或 FlashAttention batch 长时间占满 GPU，decode stream 即使优先级高也只能等 kernel 边界； 如果 chunk 太小，又会引入 launch gap、降低 Tensor Core 利用率和吞吐。调参目标通常不是“chunk 越小越好”，而是在 prefill throughput 与 decode tail latency 之间找服务级 Pareto 点。

15.7 Decode GEMV roofline

decode_roofline.cu 把 decode 阶段的线性层抽象成 y[N] = W[N,K] * x[K]。 静态算术强度约为：

FLOP = 2 * N * K
Bytes ≈ (N*K + K + N) * sizeof(dtype)
Intensity ≈ 2 FLOP / 4 Byte = 0.5 FLOP/Byte   # fp32
Intensity ≈ 1.0 FLOP/Byte                      # fp16
Intensity 更低于 H100/H200/B200 的 Tensor Core 屋顶很多，因此 decode GEMV 常被 HBM 带宽限制。

这就是 W4A16 / W8A8 量化对 decode 特别有效的原因：不是因为 Tensor Core 算不动，而是因为每 token 都要从 HBM 读一遍权重。把权重从 fp16 降到 int4，理论上权重读字节减少 4×；实际收益会被 dequant、 scale 读取、group size、cache reuse、batch 大小稀释。

15.8 量化 decode GEMV：W8A16 / W4A16

quant_decode_profile.cu 对比三种 decode 线性层读取路径：fp32 权重、group-wise int8 权重、packed int4 权重。 这里没有调用 Tensor Core，因为目标是专门观察 decode 的 M=1 GEMV：权重每 token 从 HBM 扫一遍， 量化首先是在减字节，不是提高矩阵乘峰值。 示例里的 activation 用 fp32 保存，方便复用前面章节的校验代码；工业 W8A16/W4A16 通常会把 activation 放在 fp16/bf16，同时保留同样的“读低比特权重 → 读 scale → dequant → accumulate”profiling 问题。

fp32 bytes ≈ N*K*4 + K*4 + N*4
w8a16 bytes ≈ N*K*1 + N*groups*4 + K*4 + N*4
w4a16 bytes ≈ N*ceil(K/2)*1 + N*groups*4 + K*4 + N*4
groups = ceil(K / group_size)

profile 时要同时看三件事：第一，dram__throughput 是否下降到更少字节但更高有效 tokens/s； 第二，scale 读取是否因为 group 太小变成新的带宽项；第三，dequant 指令是否让 kernel 从 memory-bound 变成 instruction/latency-bound。工业里常见 group size 是 32/64/128/256，需要按模型、GPU、batch 和 kernel 实现扫。

15.9 KV cache layout 微基准

kv_cache_profile.cu 对比三种 cache 写入布局：

profile 时重点不是只看 kernel 时间，还要看 L2 hit、DRAM throughput、全局 load/store transaction 是否合并。 Paged cache 的页大小通常是 16 或 32 token；页太小会增加页表查询，页太大又会增加内部碎片。

15.10 Prefix cache：命中、复用与错误的复制路径

prefix_cache_profile.cu 模拟线上推理常见的 前缀缓存 (prefix cache)： 多个请求共享相同系统提示词、RAG 模板或长上下文前缀时，scheduler 可以直接复用已经写好的 paged KV block。 这个示例把三条路径拆开 profile：

// Fast path: hit pages alias existing physical blocks; only misses allocate/write.
alias_hit_pages_kernel<<<page_grid, 256>>>(block_table, page_probe, cached_probe, ...);
fill_miss_pages_kernel<<<miss_grid, 256>>>(request_cache, page_probe, block_table, ...);

// Slow fallback: copying hit KV blocks materializes large HBM traffic.
copy_hit_pages_kernel<<<hit_grid, 256>>>(cached_pages, request_cache, page_probe, block_table, ...);

真实系统里 prefix cache 不能只报告“命中率”。需要同时报告命中 token/page 数、alias 是否成功、miss 分配耗时、 block table 更新耗时、cache eviction、跨 GPU / 跨实例迁移，以及命中后端到端 TTFT 是否下降。 如果命中后仍在复制大量 KV block，Nsight Compute 会看到 copy_hit_pages_kernel 的 DRAM throughput 很高； Nsight Systems 会看到这些 copy kernel 或 memcpy 占据 prefill 前后的 GPU 时间，说明 cache 复用路径没有真正省掉 HBM。

15.11 KV cache quantization：少读 K/V，别忽略 scale

kv_quant_profile.cu 对比 fp32 KV cache scan 与 int8 KV cache scan。它模拟的是长上下文 decode 的核心矛盾： 每生成一个 token 都要扫描历史 K/V，KV cache 越长，HBM 读取越重。 KV cache 量化 (KV cache quantization) 把 K/V payload 从 fp16/fp32 压到 FP8/INT8/INT4 等格式，收益来自少读历史 K/V，而代价是 scale 元数据读取、dequant 指令和质量风险。

// fp32 path: read q, K, V and reduce.
kv_scan_fp32_kernel<<<B * H, 256, smem>>>(q, k_cache, v_cache, out, B, T, H, D);

// q8 path: read q, int8 K/V, per-token scales, then dequant inside the scan.
kv_scan_q8_kernel<<<B * H, 256, smem>>>(
    q, k_q8, v_q8, k_scale, v_scale, out, B, T, H, D);

实战里不要把 KV quant 和 weight quant 混在一起下结论。W4A16 decode GEMV 解决的是“每 token 读全量权重”的带宽； KV quant 解决的是“每 token 读历史 K/V”的带宽。它们可以叠加，但 profiler 要分开看： 如果 kv_scan_q8_kernel 的 DRAM 降了、但 SM 指令或 scale 读取变成瓶颈，端到端 TPOT 不一定等比例下降。 验收时至少同时报告 TPOT、最大上下文长度、KV 显存峰值、质量/困惑度回归，以及 scale granularity。

15.12 MQA/GQA：减少 decode KV scan 字节

gqa_decode_profile.cu 对比 多头注意力 (MHA)、 分组查询注意力 (GQA) 和 多查询注意力 (MQA) 的 decode KV scan。 在长上下文 decode 中，每个输出 token 都要读历史 K/V；减少 KV head 数能直接降低 KV cache 显存占用，并减少 HBM 读字节。

// One block handles one KV head. MHA: group=1, GQA: group=QH/KVH, MQA: group=QH.
grouped_kv_decode_scan_kernel<<<B * KVH, 256, group * smem>>>(
    q, k_cache, v_cache, out, B, T, QH, KVH, D);

这个示例不是 FlashAttention decode kernel，而是一个 profiling 夹具：它把 KV 读字节、query head fanout、 shared memory reduction 和 L2 复用关系单独暴露出来。真实服务里还要叠加 paged KV、KV quant、RoPE layout、 split-K 和 batch scheduler；但如果这个微基准已经显示 KV scan memory-bound，优化方向通常先是 GQA/MQA、 KV quant、page size 和 length bucket，而不是继续堆 Tensor Core 峰值。

15.13 Continuous batching：ragged decode 与 active compaction

ragged_batch_profile.cu 模拟在线推理服务里的 连续批处理 (continuous batching)： scheduler 维护 Bmax 个 slot，请求完成后留下空洞，新请求插入，活跃请求的上下文长度也不一样。 如果 kernel 仍按 Bmax*Tmax 扫 KV cache，就会为已结束 slot 和短序列 padding 消耗大量 HBM 带宽。

// Worst case: every slot scans Tmax, padding data must be zero for correctness.
decode_kv_padded_full_kernel<<<Bmax, 256, smem>>>(q, k_cache, out, Bmax, Tmax, D);

// Better: each slot scans its real length, but inactive slots still get launched.
decode_kv_masked_ragged_kernel<<<Bmax, 256, smem>>>(q, k_cache, seq_lens, out, Bmax, Tmax, D);

// Serving-style: scheduler compacts active request ids, kernel launches only active blocks.
decode_kv_compact_active_kernel<<<active, 256, smem>>>(
    q, k_cache, seq_lens, active_ids, out, active, Tmax, D);

真实引擎通常还会叠加 length bucket、paged KV block table、prefix cache、spec decode verify batch。 这个微基准的价值是先把浪费账算清：padded_waste = 1 - sum(seq_len)/(Bmax*Tmax)。 当 padding waste 在当前 trace 中占主导时，单 kernel 的 Nsight Compute 指标再漂亮也不代表服务端吞吐好；需要回到 scheduler、 active list compaction 和 batch policy 上优化。

15.14 MoE routing：expert padding 与 token compaction

moe_routing_profile.cu 模拟 混合专家 (MoE) 层里常见的 专家路由 (expert routing) 路径： 把 token 按 expert id dispatch，执行每 expert 计算，再 gather 回原 token 顺序。 示例里的 expert 计算只是 per-expert affine，用来隔离 routing 成本；真实模型里这里通常是 grouped GEMM、 expert MLP、或者跨 GPU expert parallel 的 all-to-all。

// Fixed-capacity path: simple but burns work on empty expert capacity slots.
padded_dispatch_kernel<<<grid, 256>>>(x, expert_ids, token_ranks, routed, T, H, capacity);
padded_expert_affine_kernel<<<padded_grid, 256>>>(routed, scale, bias, expert_out, E, capacity, H);
padded_gather_kernel<<<grid, 256>>>(expert_out, expert_ids, token_ranks, y, T, H, capacity);

// Compact path: prefix offsets create a dense [T,H] routed buffer.
compact_dispatch_kernel<<<grid, 256>>>(x, compact_pos, routed, T, H);
compact_expert_affine_kernel<<<grid, 256>>>(routed, compact_expert_ids, scale, bias, expert_out, T, H);
compact_gather_kernel<<<grid, 256>>>(expert_out, compact_pos, y, T, H);

MoE profiling 不要只看专家 MLP 的 GEMM。线上慢点经常来自 token 分布偏斜：某几个 expert 热、其它 expert 空， fixed capacity 为了防溢出预留了很多空位，dispatch/gather 又是低算术强度的全局内存搬运。 因此每份 MoE profile 都应该记录 expert counts、capacity、drop rate、top-k、是否跨 GPU expert parallel。

15.15 Logits / sampling 微基准

sampler_profile.cu 把 decode tail 拆成三个常见小 kernel：全 vocab argmax、softmax 归一化统计、top-4 风格扫描。 这些算子 FLOP 不大，但每个 output token 都可能触发，而且 vocabulary 通常有 32K 到 128K 项； 如果 batch 很小，launch gap 和 host-device 同步会比 kernel 本身更显眼。

sampler_argmax_kernel<<<B, 256>>>(logits, max_values, max_indices, B, V);
sampler_softmax_sum_kernel<<<B, 256>>>(logits, max_values, sum_exp, B, V, inv_temperature);
sampler_top4_kernel<<<B, 256>>>(logits, top_values, top_indices, B, V);

这不是生产级 top-p 实现，而是一个 profiling 夹具：它让你在 Nsight Systems 里看到 sampling 的短 kernel 是否被 CPU launch 间隙淹没，也能在 Nsight Compute 里确认全 vocab 扫描更像 memory-bound reduction。 真实引擎通常会把 lm_head、top-k/top-p、随机数和 token 选择融合，或者至少做 batch sampler 来摊薄 launch 开销。

15.16 Logits processor：temperature、penalty、mask 与 top-p

logits_processor_profile.cu 模拟线上采样前常见的 Logits 处理器 (logits processor)： temperature scaling、repetition penalty、bad-word/forbidden token mask，以及 核采样 (top-p) cutoff。 这些逻辑通常不是模型主干的一部分，但每个输出 token 都会执行；在低 batch、短 kernel decode 中，它们很容易和 sampler 一起成为尾部瓶颈。

// Separate path: simple but creates several tiny launches around the vocab buffer.
temperature_kernel<<<grid, 256>>>(logits, processed, B, V, inv_temperature);
repetition_penalty_scatter_kernel<<<grid_recent, 256>>>(processed, recent_ids, ...);
bad_words_mask_scatter_kernel<<<grid_bad, 256>>>(processed, bad_ids, ...);

// Fused path: one full-vocab pass, but each token checks processor lists.
fused_logits_processor_kernel<<<grid, 256>>>(
    logits, processed, recent_ids, bad_ids, B, V, recent, bad, inv_temperature, penalty);

top_p_cutoff_sorted_kernel<<<B, 256>>>(processed, cutoff, mass, B, V, top_p);

实战里 logits processor 的优化不要只看单个 kernel。需要结合请求分布看： recent token 列表有多长、bad-word mask 是否按请求变化、top-p/top-k 是否每个请求不同、采样是否在 CPU 上做随机数和 token 选择。 如果 Nsight Systems 显示 sampler tail 被多个小 kernel 撕碎，优先考虑 batch sampler、processor fusion、device-side random/token select； 如果 Nsight Compute 显示 fused processor 因列表扫描变慢，则应该把 processor 数据结构做成 bitset、sorted small list、或者把低频规则留给 scatter kernel。

15.17 CPU/GPU 同步：token 回读与 decode 空洞

sync_profile.cu 对比三种 decode loop：每 step 后 cudaDeviceSynchronize、每 step 阻塞拷贝一个 token 到 CPU、 以及把 token 写在 device buffer 里最后一次性 async 回读。它模拟的是线上服务常见现象： GPU kernel 很短，但 CPU 每个 token 都要等一个 scalar 结果，导致 GPU row 出现周期性空洞。

// Blocking path: each token forces host to wait for the GPU.
decode_step_kernel<<<grid, 256>>>(state, tokens, n, step);
cudaMemcpy(&host_token, tokens + step, sizeof(int), cudaMemcpyDeviceToHost);

// Better for batched decisions: keep tokens on device and copy once.
for (int step = 0; step < steps; ++step) {
    decode_step_kernel<<<grid, 256, 0, stream>>>(state, tokens, n, step);
}
cudaMemcpyAsync(pinned_tokens, tokens, steps * sizeof(int), cudaMemcpyDeviceToHost, stream);

并不是所有 token 回读都能消除：流式输出最终要把 token 交给 CPU、网络或上层 scheduler。 但你应该知道它在 timeline 里花了多少，并尽量把 CPU 决策批量化、把采样留在 GPU、用 pinned host buffer 做 async copy、 或者让 scheduler 只在必要边界同步。Nsight Systems 里重点看 CUDA API row 的 D2H copy、cudaDeviceSynchronize 和 GPU row 的空洞是否成对出现。

15.18 小算子融合：residual + RMSNorm + quant

fusion_profile.cu 用 小算子融合 (kernel fusion) 对比两条 decode tail 路径：第一条把 residual add、RMSNorm、row-wise absmax、int8 quant 拆成 4 个 kernel；第二条把它们融合成 1 个 kernel。它对应工业推理里的常见 epilogue： attention/MLP 输出先和 residual 相加，再做 RMSNorm，然后为下一层 W8/W4 kernel 生成量化 activation。

// Separate path: easy to write, but it creates four launches and extra global traffic.
residual_add_kernel<<<grid, 256>>>(x, residual, added, n);
rmsnorm_kernel<<<B, 256, smem>>>(added, gamma, norm, B, H);
row_absmax_scale_kernel<<<B, 256, smem>>>(norm, scales, B, H);
quantize_kernel<<<grid, 256>>>(norm, scales, quant, B, H);

// Fused path: one launch, no materialized added tensor, scale and quant stay local to the row.
fused_residual_rmsnorm_quant_kernel<<<B, 256, smem>>>(
    x, residual, gamma, norm, scales, quant, B, H);

这个示例故意保留 norm 输出，便于和 CPU reference 做 allclose；真实服务如果下一层直接吃量化 activation， 可以不写 fp32/fp16 norm，只写 quant 和 scale。profile 时分两层看：Nsight Systems 先证明 4 个小 kernel 是否被 launch gap 放大；Nsight Compute 再确认 fused kernel 是否真的减少了 global load/store，且没有因为寄存器过多、 shared memory 同步或 SFU 指令把收益吃掉。

15.19 RoPE + layout transform：SFU、LUT 与融合写出

rope_layout_profile.cu 对比三种 RoPE 路径： 先做 RoPE 再把 BSHD 转成 BHSD、直接 fused RoPE+layout、以及用预计算 sin/cos LUT 的 fused 路径。 这对应真实 LLM 里的 Q/K 准备阶段：Q 通常给 attention kernel 直接消费，K 要写入 KV cache，layout 选择会影响后续 decode scan。

// Separate path: two launches and one materialized temporary tensor.
rope_sincos_bshd_kernel<<<pair_grid, 256>>>(x, tmp, B, T, H, D);
bshd_to_bhsd_kernel<<<elem_grid, 256>>>(tmp, out, B, T, H, D);

// Fused path: rotate pair and directly store the attention/cache-friendly layout.
fused_rope_lut_bhsd_kernel<<<pair_grid, 256>>>(
    x, out, cos_lut, sin_lut, B, T, H, D);

工业实现通常不会孤立优化 RoPE：它会和 Q/K projection epilogue、Q/K layout transform、KV cache append、 quant/dequant 或 attention kernel prologue 融合。profile 时先用 Nsight Systems 看 RoPE 是否是独立小 kernel； 如果是，再用 Nsight Compute 看 fused_rope_lut_bhsd_kernel 是 SFU-bound、memory-bound， 还是被非合并写出拖慢。LUT 不总是更快：短上下文、小 batch、L2 足够时通常划算；超长上下文或跨请求乱序访问时，要验证 LUT hit rate。

15.20 CUDA Graph：把 decode 小 kernel 串变成 replay

cuda_graph_decode.cu 模拟一轮 decode 的短 kernel 链：RMSNorm、RoPE-like pair rotation、residual mix、argmax。 eager 模式每个 token step 发起 4 次 kernel launch；CUDA Graph 模式先捕获这 4 个 kernel，再每个 token step 只 replay 一次 graph。

cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);
rmsnorm_small_kernel<<<B, 256, smem, stream>>>(...);
rope_pair_kernel<<<pair_grid, 256, 0, stream>>>(...);
residual_mix_kernel<<<elem_grid, 256, 0, stream>>>(...);
argmax_hidden_kernel<<<B, 256, 0, stream>>>(...);
cudaStreamEndCapture(stream, &graph);
cudaGraphInstantiate(&graph_exec, graph, nullptr, nullptr, 0);

for (int s = 0; s < steps; ++s) {
    cudaGraphLaunch(graph_exec, stream);
}

graph 不是“免费加速开关”。它最适合 shape 稳定、kernel 拓扑稳定、buffer 地址稳定的 decode bucket： 例如 B=16, T=1, H=4096 这类固定批次。真实服务通常要按 batch、beam、是否 speculative verify、 KV page size、max sequence bucket 捕获多份 graph；shape 一直变化时，capture/instantiate 成本会吞掉收益。 Nsight Systems 里应该能看到 eager loop 的 CUDA API row 有密集 launch，而 graph replay loop 的 CPU launch 事件显著减少。

15.21 Speculative decoding：verify batch 与接受率

spec_decode_profile.cu 对比 target-only decode 和 推测解码 (speculative decoding) verify batch。target-only 路径每接受一个 token 就发一次 target kernel；spec 路径每个 window 先验证 draft model 提议的多个 token，再 compact 被接受的 token。它模拟的是线上常见权衡：launch 少了，但被拒绝 token 的 verify work 是额外成本。

// Target-only: many tiny launches, no rejected-token work.
for (int w = 0; w < windows; ++w) {
    for (int j = 0; j < accept; ++j) {
        target_only_token_kernel<<<grid, 256>>>(x, accepted, n, out_idx, proposal_id);
    }
}

// Speculative verify: fewer launches, but verifies draft tokens that may be rejected.
for (int w = 0; w < windows; ++w) {
    spec_verify_window_kernel<<<verify_grid, 256>>>(x, proposals, n, start, draft);
    compact_accepted_window_kernel<<<compact_grid, 256>>>(proposals, accepted, n, start, out, accept);
}

实战里 spec decode profile 至少要拆三段：draft model、target verify、accepted token compact / sampler。 Nsight Systems 用来看 verify batch 是否形成更粗的 GPU work；Nsight Compute 用来看 verify kernel 是 Tensor Core 受限、 memory-bound，还是被 ragged verify shape 拖低 occupancy。最终不要只报告“verify kernel 更快”，而要报告 accepted tokens/s、TPOT、accept rate、最大并发和质量指标。

15.22 Allocator：hot path 分配与 workspace pool

allocator_profile.cu 对比三种临时 workspace 管理方式：每 step cudaMalloc/cudaFree、预分配 ring pool、 以及 CUDA 11.2+ 的 cudaMallocAsync/cudaFreeAsync。LLM 服务里的 KV block、logits buffer、 attention workspace、采样临时数组如果在 hot path 反复分配，会在 Nsight Systems 的 CUDA API row 上形成明显空洞。

// Bad in decode hot path: allocation can synchronize and serialize CPU/GPU work.
cudaMalloc(&tmp, bytes);
workspace_kernel<<<grid, 256, 0, stream>>>(tmp, out, n, seed);
cudaFree(tmp);

// Better: reuse stable addresses from a preallocated pool.
float* tmp = pool + (step % slots) * elems;
workspace_kernel<<<grid, 256, 0, stream>>>(tmp, out, n, seed);

真实工程里优先做预分配：按最大 batch、最大上下文、并发上限和 workspace 类型建立 pool，把 alloc/free 从 decode loop 移到初始化或请求入队阶段。cudaMallocAsync 是有用工具，但仍要设置并复用 memory pool， 并确认 Nsight Systems 上没有每 token 的 allocator API gap。对服务端来说，稳定地址还有一个额外好处： CUDA Graph capture/replay 更容易复用。

15.23 多 GPU P2P：NVLink / PCIe 通信与 overlap

multi_gpu_p2p_profile.cu 用纯 CUDA Runtime 测量两张 GPU 之间的 cudaMemcpyPeerAsync 带宽，并同时在目标 GPU 上跑一个 compute stream，观察通信和计算能否 overlap。它不依赖 NCCL，因此适合作为读 tensor parallel / pipeline parallel profiler 之前的最小夹具。

cudaDeviceCanAccessPeer(&can, src, dst);
cudaDeviceEnablePeerAccess(dst, 0);

cudaMemcpyPeerAsync(dst_ptr, dst, src_ptr, src, bytes, copy_stream);
fma_stress_kernel<<<grid, 256, 0, compute_stream>>>(work_ptr, n, iters);

真实 LLM 多卡推理里，通信热点通常来自 tensor parallel 的 all-reduce / all-gather、expert parallel 的 token routing、 pipeline parallel 的 activation 传递，以及 KV / prefix cache 迁移。这个示例不替代 NCCL profile，但能先回答两个底层问题： 这两张卡之间是否有 direct P2P？走的是 NVLink 还是 PCIe？copy engine 上的传输是否挡住了 compute stream？ 在 Nsight Systems 里重点看 CUDA API row、Memcpy row、NVLink/PCIe row 和 compute kernel 是否真正重叠。

15.24 Tensor Parallel all-reduce：通信是否挡住 decode

tp_allreduce_profile.cu 用纯 CUDA Runtime 模拟 2-GPU 张量并行 (tensor parallel) all-reduce：两个 rank 各持有一段 activation shard，互相 P2P 拷贝 peer shard，然后本地做 vector sum。 它不是 NCCL 替代品，而是一个能在没有 NCCL 依赖时观察通信形态的 profiling 夹具。

// Sequential path: rank0 copy+reduce, then rank1 copy+reduce. Easy to understand, poor overlap.
cudaMemcpyPeerAsync(rank0_peer, gpu0, rank1_local, gpu1, bytes, rank0_stream);
vector_sum_kernel<<<grid, 256, 0, rank0_stream>>>(rank0_local, rank0_peer, rank0_out, n);

// Bidirectional path: both ranks exchange peer shards and reduce on their own streams.
cudaMemcpyPeerAsync(rank0_peer, gpu0, rank1_local, gpu1, bytes, rank0_stream);
cudaMemcpyPeerAsync(rank1_peer, gpu1, rank0_local, gpu0, bytes, rank1_stream);

真实 tensor parallel profile 仍应以 NCCL trace 为准：看 all-reduce/all-gather 是否在每层 MLP/attention 后挡住下一步， TP size 是否过大，rank placement 是否跨 PCIe root complex，NVLink 是否全互联，是否能用 reduce-scatter/all-gather 或 compute/comm overlap 缩短 critical path。这个 toy benchmark 的价值是先把“通信字节、拓扑、copy engine、 本地 reduce kernel、compute overlap”拆开，避免一上来就把所有问题都归因于 NCCL。

15.25 Host/device staging：pageable、pinned 与 copy engine

host_transfer_profile.cu 对比三条服务端数据通路：普通 可分页主机内存 (pageable host memory) 的同步 H2D/D2H roundtrip、页锁定内存 (pinned host memory) 的 async roundtrip、 以及两个 stream 上的 pinned chunk pipeline。它模拟真实推理里的 request staging、logits/token 回传、LoRA adapter 热切换、KV offload/prefetch 等路径。很多线上 p99 抖动不是 kernel 算得慢，而是 CPU 线程在 pageable 拷贝、隐式同步、 或 拷贝引擎 (copy engine) 排队上消耗了时间。

// Blocking path: pageable host memory makes the host wait around the copy.
cudaMemcpy(d_in, pageable_in, bytes, cudaMemcpyHostToDevice);
transfer_compute_kernel<<<grid, 256>>>(d_in, d_out, n, compute_iters);
cudaMemcpy(pageable_out, d_out, bytes, cudaMemcpyDeviceToHost);

// Better path: pinned memory makes H2D/D2H visible as async work on a stream.
cudaHostAlloc(&pinned_in, bytes, cudaHostAllocDefault);
cudaMemcpyAsync(d_in, pinned_in, bytes, cudaMemcpyHostToDevice, stream);
transfer_compute_kernel<<<grid, 256, 0, stream>>>(d_in, d_out, n, compute_iters);
cudaMemcpyAsync(pinned_out, d_out, bytes, cudaMemcpyDeviceToHost, stream);

profile 时先看 cudaDeviceProp.asyncEngineCount、Nsight Systems 的 Memcpy row、CUDA API row 和 GPU kernel row。 如果 H2D/D2H 明明用了 cudaMemcpyAsync 却没有 overlap，常见原因是 host buffer 不是 pinned、所有操作都进了默认 stream、 D2H 结果马上被 CPU 阻塞读取、或者 copy engine 已被 P2P/NCCL/KV offload 占满。工业服务通常用 pinned ring buffer 绑定 scheduler：输入 token、采样结果、beam 状态和少量 metadata 走小而稳定的 async copy；大块权重/KV 迁移则单独标 NVTX range， 避免和 decode critical path 混在一起。

15.26 Paged KV decode scan：页表、page size 与碎片化

paged_kv_decode_profile.cu 对比三条长上下文 decode KV scan 路径：连续 BHSD cache、物理页按逻辑顺序排列的 paged cache、 以及物理页被打散的 fragmented paged cache。它补齐了 kv_cache_profile.cu 只测写入布局的空白，专门观察 decode 读历史 K/V 时，块表 (block table) 间接寻址和物理页局部性如何影响 L2/DRAM。

// Contiguous path: address is arithmetic only.
base = ((((b * H + h) * T + t) * D + d) * 2);
acc += combine_kv(kv[base + 0], kv[base + 1]);

// Paged path: load block_table, then read the physical page.
phys_page = block_table[b * max_pages + logical_page];
base = ((((phys_page * page_size + offset) * H + h) * D + d) * 2);
acc += combine_kv(pool[base + 0], pool[base + 1]);

真正的 PagedAttention kernel 会把 QK、softmax、V 累加和 block table 访问融合在一个 attention kernel 内； 这个微基准故意只保留 KV payload scan 和 block table 读取，便于用 Nsight Compute 单独看 memory workload。 实战里扫 page=8/16/32/64，同时记录内部碎片、block table bytes、L2 hit、DRAM throughput 和质量无关的 TPOT。 如果 fragmented 路径明显慢，优先检查 KV allocator、prefix cache eviction、请求迁移和跨 GPU KV page placement。

运行结果

15.27 编译与普通运行

cd code/ch15_profiling
make ARCH=sm_80          # A100/A800
make ARCH=sm_89          # RTX 4090 / L40S / L4
make ARCH=sm_90a         # H100/H200
make NVTX=0              # 如果环境缺少 libnvToolsExt

make run
make bench

示例输出格式如下，具体数值请在 GPU 上实测：

[saxpy] PASS  max_abs=...  max_rel=...  l2=...
[fma_chain] PASS  max_abs=...  max_rel=...  l2=...
memory-bound saxpy: TODO(on GPU)
compute-bound fma : TODO(on GPU)

[gemv_row] PASS  max_abs=...  max_rel=...  l2=...
[gemv_splitk] PASS  max_abs=...  max_rel=...  l2=...
roofline static model: ... FLOP/Byte

[monolithic_prefill] PASS ...
[monolithic_decode] PASS ...
[chunked_prefill] PASS ...
[chunked_decode] PASS ...
[priority_prefill] PASS ...
[priority_decode] PASS ...
chunked prefill shape prefillMB=... chunks=... decode_steps=...

[fp32_gemv] PASS ...
[q8_gemv] PASS ...
[q4_gemv] PASS ...
quantization drift vs fp32: ...
quant decode shape N=... K=... group=...

[BSHD] PASS ...
[BHSD] PASS ...
[Paged] PASS ...

[contiguous_scan] PASS ...
[paged_ordered_scan] PASS ...
[paged_fragmented_scan] PASS ...
paged KV decode shape B=... T=... H=... D=... page=...

[no_cache] PASS ...
[prefix_alias_cache] PASS ...
[prefix_alias_probe] PASS ...
[prefix_alias_table] PASS ...
[prefix_copy] PASS ...
prefix cache shape B=... prefix_pages=... hit_pages=... miss_pages=...

[fp32_kv_scan] PASS ...
[q8_kv_scan] PASS ...
KV quant shape B=... T=... H=... D=... repeats=...
quantization drift vs fp32: ...

[mha_scan] PASS ...
[gqa_scan] PASS ...
[mqa_scan] PASS ...
gqa decode shape B=... T=... QH=... GQA_KVH=... D=...

[padded_full] PASS ...
[masked_ragged] PASS ...
[compact_active] PASS ...
ragged batch shape Bmax=... active=... Tmax=... D=...

[padded_moe] PASS ...
[compact_moe] PASS ...
moe routing shape T=... H=... E=... capacity=...

[argmax_value] PASS ...
[argmax_index] PASS
[softmax_sum] PASS ...
[top4_value] PASS ...
[top4_index] PASS
sampler shape B=... V=... temperature=... repeats=...

[separate_processor] PASS ...
[fused_processor] PASS ...
[top_p_cutoff] PASS ...
[top_p_mass] PASS ...
logits processor shape B=... V=... recent=... bad=...

[sync_each_step_state] PASS ...
[blocking_readback_state] PASS ...
[batched_async_state] PASS ...
sync profile shape N=... steps=...

[pageable_roundtrip] PASS ...
[pinned_async_roundtrip] PASS ...
[pinned_overlap_pipeline] PASS ...
host transfer shape MB=... chunks=... iters=... compute=...

[separate_norm] PASS ...
[fused_norm] PASS ...
[separate_scale] PASS ...
[fused_scale] PASS ...
[separate_quant] PASS ...
[fused_quant] PASS ...
fusion profile shape B=... H=... repeats=...

[separate_sincos] PASS ...
[fused_sincos] PASS ...
[fused_lut] PASS ...
rope layout shape B=... T=... H=... D=... repeats=...

[eager_state] PASS ...
[graph_state] PASS ...
[eager_max] PASS ...
[graph_max] PASS ...
[eager_token] PASS
[graph_token] PASS
decode graph shape B=... H=... steps=..., kernels_per_step=4

[target_only] PASS ...
[spec_verify] PASS ...
spec decode shape N=... windows=... draft=... accept=...

[malloc_free] PASS ...
[prealloc_pool] PASS ...
[malloc_async] PASS ... 或 SKIP
allocator shape MB=... iters=... slots=...

[p2p_copy] PASS ...
[overlap_copy] PASS ...
P2P direct access src->dst=... dst->src=...

[sequential_allreduce_gpu0] PASS ...
[bidirectional_allreduce_gpu0] PASS ...
[overlap_allreduce_gpu0] PASS ...
TP all-reduce P2P direct src->dst=... dst->src=...
或 [tp_allreduce] SKIP ...

15.28 Nsight Systems：先看 timeline

nsys profile --stats=true \
  --trace=cuda,nvtx,osrt \
  --capture-range=cudaProfilerApi \
  --capture-range-end=stop \
  --force-overwrite=true \
  -o profiling_ranges \
  ./profiling_ranges --N=16777216 --iters=4096 --repeats=20

nsys stats profiling_ranges.nsys-rep

看图顺序：

1. CUDA API row：有没有大量 cudaMalloc、cudaFree、cudaMemcpy、cudaDeviceSynchronize。
2. GPU row：kernel 之间是否有空洞，空洞对应 CPU 哪个线程在跑。
3. NVTX row：prefill / decode / sampling / layer range 是否清晰。
4. 每个 stream：是否所有工作挤在默认 stream，H2D/D2H/P2P 是否真的 overlap。

Nsight Systems 支持 --capture-range=cudaProfilerApi、--capture-range=nvtx 等方式限制采集范围，详见 Nsight Systems User Guide。

15.28.1 推荐：用脚本保存可复现证据

真实排查不要只留一个 .nsys-rep 文件。用 scripts/profile_llm.sh 同时保存 GPU 健康、拓扑、环境变量、git 状态和实际命令。

# Ch14 端到端 decode timeline
./scripts/profile_llm.sh --name mini_llm_decode --mode nsys -- \
  code/ch14_mini_llm/mini_llm \
  --weights=data/gpt2-small.bin \
  --tokens=15496,11,612,318 \
  --max_new=4 \
  --profile

# Ch15 KV layout 单 kernel memory analysis
./scripts/profile_llm.sh --name kv_layout --mode ncu \
  --kernel 'fill_.*' --ncu-set memory -- \
  code/ch15_profiling/kv_cache_profile \
  --B=8 --T=512 --H=32 --D=128 --page=16

# Ch15 paged KV decode scan and fragmented page table
./scripts/profile_llm.sh --name paged_kv_decode --mode both \
  --kernel '(contiguous_scan|paged_scan).*' --ncu-set memory -- \
  code/ch15_profiling/paged_kv_decode_profile \
  --B=2 --T=2048 --H=16 --D=128 --page=16 --repeats=20

# Ch15 chunked prefill vs decode-priority scheduling
./scripts/profile_llm.sh --name chunked_prefill --mode nsys -- \
  code/ch15_profiling/chunked_prefill_profile \
  --prefillMB=32 --chunks=8 --decodeN=65536 --decode_steps=32 --prefill_compute=64 --decode_compute=16

# Ch15 prefix cache alias vs copy-hit path
./scripts/profile_llm.sh --name prefix_cache --mode both \
  --kernel '(fill_no_cache|alias_hit|fill_miss|copy_hit).*' --ncu-set memory -- \
  code/ch15_profiling/prefix_cache_profile \
  --B=8 --prefix=128 --hit=75 --share=4 --page=16 --H=16 --D=128 --repeats=20

# Ch15 quantized KV cache decode scan
./scripts/profile_llm.sh --name kv_quant --mode both \
  --kernel 'kv_scan_(fp32|q8).*' --ncu-set memory -- \
  code/ch15_profiling/kv_quant_profile \
  --B=4 --T=2048 --H=16 --D=128 --repeats=20

# Ch15 continuous batching ragged decode
./scripts/profile_llm.sh --name ragged_decode --mode both \
  --kernel 'decode_kv_.*' --ncu-set memory -- \
  code/ch15_profiling/ragged_batch_profile \
  --Bmax=64 --active=37 --Tmax=2048 --D=128 --repeats=20

# Ch15 MHA/GQA/MQA decode KV scan
./scripts/profile_llm.sh --name gqa_decode --mode both \
  --kernel 'grouped_kv_decode_scan.*' --ncu-set memory -- \
  code/ch15_profiling/gqa_decode_profile \
  --B=2 --T=2048 --QH=32 --GQA_KVH=8 --D=128 --repeats=20

# Ch15 MoE expert routing and padding waste
./scripts/profile_llm.sh --name moe_routing --mode both \
  --kernel '(padded|compact)_.*' --ncu-set memory -- \
  code/ch15_profiling/moe_routing_profile \
  --T=4096 --H=1024 --E=8 --repeats=20

# Ch15 quantized decode GEMV memory analysis
./scripts/profile_llm.sh --name quant_decode --mode ncu \
  --kernel 'gemv_(fp32|q8|q4).*' --ncu-set memory -- \
  code/ch15_profiling/quant_decode_profile \
  --N=4096 --K=4096 --group=128 --repeats=20

# Ch15 sampler vocab scan
./scripts/profile_llm.sh --name sampler_tail --mode both \
  --kernel 'sampler_.*' --ncu-set memory -- \
  code/ch15_profiling/sampler_profile \
  --B=16 --V=50257 --repeats=200

# Ch15 logits processor and top-p tail
./scripts/profile_llm.sh --name logits_processor --mode both \
  --kernel '(temperature|repetition|bad_words|fused_logits|top_p).*' --ncu-set memory -- \
  code/ch15_profiling/logits_processor_profile \
  --B=16 --V=50257 --recent=16 --bad=32 --temperature=0.8 --penalty=1.1 --top_p=0.9 --repeats=100

# Ch15 decode synchronization and token readback
./scripts/profile_llm.sh --name sync_gap --mode nsys -- \
  code/ch15_profiling/sync_profile \
  --N=1048576 --steps=256

# Ch15 host/device staging and copy engine overlap
./scripts/profile_llm.sh --name host_transfer --mode nsys -- \
  code/ch15_profiling/host_transfer_profile \
  --MB=16 --chunks=8 --iters=20 --compute=64

# Ch15 residual + RMSNorm + quant fusion
./scripts/profile_llm.sh --name fusion_norm_quant --mode both \
  --kernel '(residual_add|rmsnorm|row_absmax|quantize|fused_residual).*' \
  --ncu-set memory -- \
  code/ch15_profiling/fusion_profile \
  --B=16 --H=4096 --repeats=100

# Ch15 RoPE + Q/K layout transform
./scripts/profile_llm.sh --name rope_layout --mode both \
  --kernel '(rope_sincos|bshd_to_bhsd|fused_rope).*' --ncu-set memory -- \
  code/ch15_profiling/rope_layout_profile \
  --B=4 --T=1024 --H=32 --D=128 --repeats=20

# Ch15 CUDA Graph decode launch gap
./scripts/profile_llm.sh --name graph_decode --mode nsys -- \
  code/ch15_profiling/cuda_graph_decode \
  --B=16 --H=4096 --steps=256

# Ch15 speculative decoding verify batch
./scripts/profile_llm.sh --name spec_decode --mode both \
  --kernel '(target_only|spec_verify|compact_accepted).*' --ncu-set memory -- \
  code/ch15_profiling/spec_decode_profile \
  --N=4096 --windows=256 --draft=4 --accept=3

# Ch15 allocator hot-path API gap
./scripts/profile_llm.sh --name allocator_gap --mode nsys -- \
  code/ch15_profiling/allocator_profile \
  --MB=64 --iters=100 --slots=8

# Ch15 multi-GPU P2P copy and overlap
./scripts/profile_llm.sh --name p2p_overlap --mode nsys -- \
  code/ch15_profiling/multi_gpu_p2p_profile \
  --src=0 --dst=1 --MB=256 --iters=20 --compute=256

# Ch15 tensor-parallel toy all-reduce
./scripts/profile_llm.sh --name tp_allreduce --mode nsys -- \
  code/ch15_profiling/tp_allreduce_profile \
  --src=0 --dst=1 --MB=64 --iters=20 --compute=128

15.29 Nsight Compute：再钻热 kernel

ncu --section SpeedOfLight \
    --section MemoryWorkloadAnalysis \
    --section Occupancy \
    --kernel-name 'regex:gemv_.*' \
    --force-overwrite \
    -o decode_roofline \
    ./decode_roofline --N=4096 --K=4096 --splits=8

ncu --set full \
    --kernel-name 'regex:fill_.*' \
    --force-overwrite \
    -o kv_cache_profile \
    ./kv_cache_profile --B=8 --T=512 --H=32 --D=128 --page=16

ncu --section SpeedOfLight \
    --section MemoryWorkloadAnalysis \
    --section Occupancy \
    --kernel-name 'regex:(contiguous_scan|paged_scan).*' \
    --force-overwrite \
    -o paged_kv_decode_profile \
    ./paged_kv_decode_profile --B=2 --T=2048 --H=16 --D=128 --page=16 --repeats=20

ncu --section SpeedOfLight \
    --section MemoryWorkloadAnalysis \
    --section Occupancy \
    --kernel-name 'regex:(fill_no_cache|alias_hit|fill_miss|copy_hit).*' \
    --force-overwrite \
    -o prefix_cache_profile \
    ./prefix_cache_profile --B=8 --prefix=128 --hit=75 --share=4 --page=16 --H=16 --D=128 --repeats=20

ncu --section SpeedOfLight \
    --section MemoryWorkloadAnalysis \
    --section Occupancy \
    --kernel-name 'regex:kv_scan_(fp32|q8).*' \
    --force-overwrite \
    -o kv_quant_profile \
    ./kv_quant_profile --B=4 --T=2048 --H=16 --D=128 --repeats=20

ncu --section SpeedOfLight \
    --section MemoryWorkloadAnalysis \
    --section Occupancy \
    --kernel-name 'regex:decode_kv_.*' \
    --force-overwrite \
    -o ragged_batch_profile \
    ./ragged_batch_profile --Bmax=64 --active=37 --Tmax=2048 --D=128 --repeats=20

ncu --section SpeedOfLight \
    --section MemoryWorkloadAnalysis \
    --section Occupancy \
    --kernel-name 'regex:grouped_kv_decode_scan.*' \
    --force-overwrite \
    -o gqa_decode_profile \
    ./gqa_decode_profile --B=2 --T=2048 --QH=32 --GQA_KVH=8 --D=128 --repeats=20

ncu --section SpeedOfLight \
    --section MemoryWorkloadAnalysis \
    --section Occupancy \
    --kernel-name 'regex:(padded|compact)_.*' \
    --force-overwrite \
    -o moe_routing_profile \
    ./moe_routing_profile --T=4096 --H=1024 --E=8 --repeats=20

ncu --section SpeedOfLight \
    --section MemoryWorkloadAnalysis \
    --section Occupancy \
    --kernel-name 'regex:gemv_(fp32|q8|q4).*' \
    --force-overwrite \
    -o quant_decode_profile \
    ./quant_decode_profile --N=4096 --K=4096 --group=128 --repeats=20

ncu --section SpeedOfLight \
    --section MemoryWorkloadAnalysis \
    --section Occupancy \
    --kernel-name 'regex:sampler_.*' \
    --force-overwrite \
    -o sampler_profile \
    ./sampler_profile --B=16 --V=50257 --repeats=200

ncu --section SpeedOfLight \
    --section MemoryWorkloadAnalysis \
    --section Occupancy \
    --kernel-name 'regex:(temperature|repetition|bad_words|fused_logits|top_p).*' \
    --force-overwrite \
    -o logits_processor_profile \
    ./logits_processor_profile --B=16 --V=50257 --recent=16 --bad=32 --temperature=0.8 --penalty=1.1 --top_p=0.9 --repeats=100

ncu --section SpeedOfLight \
    --section MemoryWorkloadAnalysis \
    --section Occupancy \
    --kernel-name 'regex:(residual_add|rmsnorm|row_absmax|quantize|fused_residual).*' \
    --force-overwrite \
    -o fusion_profile \
    ./fusion_profile --B=16 --H=4096 --repeats=100

ncu --section SpeedOfLight \
    --section MemoryWorkloadAnalysis \
    --section Occupancy \
    --kernel-name 'regex:(rope_sincos|bshd_to_bhsd|fused_rope).*' \
    --force-overwrite \
    -o rope_layout_profile \
    ./rope_layout_profile --B=4 --T=1024 --H=32 --D=128 --repeats=20

ncu --section SpeedOfLight \
    --section MemoryWorkloadAnalysis \
    --section Occupancy \
    --kernel-name 'regex:(target_only|spec_verify|compact_accepted).*' \
    --force-overwrite \
    -o spec_decode_profile \
    ./spec_decode_profile --N=4096 --windows=256 --draft=4 --accept=3

性能数据

15.30 主流 GPU 推理选型速查（截至 2026-07）

下表优先列 NVIDIA 官方页面直接给出的公开规格中对 LLM 推理最关键的维度。带 * 的 Tensor Core 数字是 NVIDIA 官方表中的 sparse 口径； dense 口径通常是其一半。消费卡或工作站卡页面没有列出完整显存带宽时，用“以厂商规格为准”标记。实际 tokens/s 取决于 batch、序列长度、量化、kernel、调度器与功耗墙。

15.31 LLM infra 热点与优化动作

自检清单

1. 能解释为什么 prefill 和 decode 不能混在一个 profile 里下结论。
2. 能用 Nsight Systems 判断长 prefill 是否挡住 decode，并解释 chunked prefill 的吞吐/延迟权衡。
3. 能用 nsys 判断慢是 CPU 调度/launch 问题还是 GPU kernel 问题。
4. 能用 ncu 判断一个 kernel 是 memory-bound、compute-bound、occupancy-bound 还是同步/指令受限。
5. 能给一个 decode GEMV 算出静态 FLOP、bytes、arithmetic intensity。
6. 能解释 W8A16/W4A16 decode GEMV 的权重字节、scale 字节和 dequant 指令成本分别怎么进入 profile。
7. 能说明 BSHD、BHSD、Paged KV cache 各自的 profiler 预期。
8. 能解释 paged KV decode scan 中 block table、page size、内部碎片和物理页局部性分别如何影响 TPOT。
9. 能说明 prefix cache 命中后为什么应优先复用 block table，而不是复制命中 KV block。
10. 能区分 weight quant 和 KV cache quant，并能解释 KV scale 读取和 dequant 指令如何影响 TPOT。
11. 能计算 MHA/GQA/MQA 在同一 B,T,QH,D 下的 KV cache 字节差异。
12. 能解释 continuous batching 中 padded scan、masked ragged scan、active compaction 的差异。
13. 能解释 MoE routing 中 expert load imbalance、capacity padding、dispatch/gather 对 profile 的影响。
14. 能解释为什么 sampling 侧小 kernel 在低 batch decode 里会放大 launch gap。
15. 能判断 logits processor 是应该融合、保留 scatter 小 kernel，还是改成 device-side batch sampler。
16. 能识别每 token 阻塞 D2H 回读和 cudaDeviceSynchronize 造成的 decode 空洞。
17. 能用 Nsight Systems 判断 pageable/pinned host buffer、H2D/D2H copy 和 copy/compute overlap 是否拖慢服务路径。
18. 能判断 residual + RMSNorm + quant 是否值得融合，并能说明 fusion 过度导致 register spill 的风险。
19. 能判断 RoPE 是 SFU-bound、LUT memory-bound，还是被 Q/K layout transform 的 global traffic 拖慢。
20. 能判断一个 decode path 是否适合 CUDA Graph，并说清需要按哪些 shape 分桶。
21. 能说明 speculative decoding 的 accept rate、verify waste、accepted tokens/s 为什么必须一起报告。
22. 能在 Nsight Systems 里识别 allocator hot path，并说明为什么预分配 pool 通常优先于每 token 分配。
23. 能用 nvidia-smi topo -m 和 Nsight Systems 判断两张 GPU 间通信走 NVLink、PCIe 还是被 CPU 调度挡住。
24. 能解释 tensor parallel all-reduce 为什么需要同时看通信字节、rank placement、TP size 和 compute overlap。
25. 能根据模型大小、上下文长度、吞吐目标和预算选择 T4/L4/L40S/A100/H100/H200/B200 这类 GPU。
26. 能写出一个带 NVTX range、CPU reference、allclose、可用 nsys/ncu 复现的微基准。

练习题

1. 把 decode_roofline.cu 改成 fp16 输入、fp32 累加，比较 fp32 与 fp16 的 arithmetic intensity 和实际吞吐。
2. 用 chunked_prefill_profile.cu 扫 chunks=1/2/4/8/16，记录 first_decode、total、TPOT proxy 和 Nsight Systems 中 decode 是否插进 prefill chunk 间隙。
3. 用 quant_decode_profile.cu 扫 group=32/64/128/256，记录 W8/W4 的 DRAM throughput、scale 读取和 quantization drift。
4. 给 kv_cache_profile.cu 增加 decode 单 token append 模式，比较 T=1 与 prefill bulk fill 的差异。
5. 用 ncu --section MemoryWorkloadAnalysis profile BSHD/BHSD/Paged 三个 kernel，记录 DRAM throughput、L2 hit rate、global store transactions。
6. 用 paged_kv_decode_profile.cu 扫 page=8/16/32/64 和 T=512/2048/8192，比较 contiguous、paged ordered、paged fragmented 的 L2 hit、DRAM throughput、table bytes 和内部碎片。
7. 用 prefix_cache_profile.cu 扫 hit=0/25/50/75/100 和 share=1/2/4/8，记录 alias 与 copy-hit 的 HBM 字节和 TTFT 影响。
8. 用 kv_quant_profile.cu 扫 T=512/2048/8192 和 D=64/128，记录 q8/fp32 bytes ratio、scale L2 hit、dequant 指令和 quantization drift。
9. 用 gqa_decode_profile.cu 扫 GQA_KVH=1/2/4/8/16/32，记录 KV cache bytes、L2 hit 和 DRAM throughput。
10. 用 ragged_batch_profile.cu 扫 active=8/16/32/64 和不同长度分布，记录 padding waste、DRAM throughput、tail effect。
11. 用 moe_routing_profile.cu 扫 E=4/8/16 和 capacity，记录 expert counts、padding waste、dispatch/gather 带宽。
12. 用 sampler_profile.cu 比较 B=1/4/16/64 时 argmax、softmax sum、top-4 的 launch gap 和 DRAM throughput。
13. 用 logits_processor_profile.cu 扫 recent=0/8/32/128 和 bad=0/32/256，比较 separate、fused、top-p cutoff 的 launch gap、branch divergence 和 vocab traffic。
14. 用 sync_profile.cu 对比逐 step 同步、阻塞 token 回读、batched async 回读的 TPOT 和 CUDA API row。
15. 用 host_transfer_profile.cu 扫 compute=0/32/128 与 chunks=1/2/4/8，比较 pageable、pinned async、pinned pipeline 的 Memcpy row 和 copy/compute overlap。
16. 用 fusion_profile.cu 比较拆分 kernel 与 fused kernel 的 launch 数、DRAM throughput、occupancy 和 register spill。
17. 把 fusion_profile.cu 改成“不写 norm，只写 quant + scale”，观察 memory traffic 和 allclose 校验应如何调整。
18. 用 rope_layout_profile.cu 扫 T=128/512/2048/8192，比较 separate、fused sin/cos、fused LUT 的 SFU、L2 hit 和 DRAM throughput。
19. 用 cuda_graph_decode.cu 比较 eager loop 与 CUDA Graph replay 的 CUDA API row、GPU gap、us/token-step。
20. 用 spec_decode_profile.cu 扫 draft=2/4/8 和 accept=1..draft，记录 launch reduction、verify waste、us/accepted-token。
21. 用 allocator_profile.cu 对比 cudaMalloc/cudaFree、预分配 pool、cudaMallocAsync 在 CUDA API row 上的差异。
22. 在双卡机器上运行 multi_gpu_p2p_profile.cu，对比 NVLink 与 PCIe 拓扑下的 P2P GB/s 和 copy/compute overlap。
23. 在双卡机器上运行 tp_allreduce_profile.cu，比较 sequential、bidirectional、overlap_compute 三段在 Nsight Systems 中的 critical path。
24. 用第 14 章 mini_llm.cu --profile 生成一份只采集 decode 的 .nsys-rep，展开 layer_N 对比 attention / MLP / logits / sampler 的占比。
25. 在 H100/H200 上把 CUDA_DEVICE_MAX_CONNECTIONS 分别设成 8/16/32，观察多 stream decode 的 timeline 是否变化。
26. 为一个真实模型写 profile 记录表：GPU、driver、CUDA、batch、ISL、OSL、dtype、quant、TTFT、TPOT、tokens/s、profile 文件名。

常见坑

• 把 profiler 开销当成真实性能：ncu --set full 会 replay kernel，端到端耗时不能直接拿来做 p99 结论。
• 只看 GPU util：GPU util 接近满值不代表高效，可能大部分时间都在低效访存；要结合 SM/DRAM throughput。
• 不固定频率：不同温度、功耗墙、MIG、共享机器会导致结果抖动。严谨 benchmark 要记录 clocks 和 power limit。
• profile 了 warmup：第一次 launch 可能包含 lazy module loading、JIT、cuBLAS handle 初始化；先 warmup 再采集。
• shape 不真实：只测 B=1,T=128 无法代表线上 B=32,T=4k，prefill/decode/spec decode 要分桶。
• 以为 stream priority 能抢占长 kernel：CUDA stream priority 通常只能影响后续 work 的调度顺序，已经运行的大 prefill kernel 不会因为 decode 来了就被切开；低延迟需要 chunked prefill 或更细粒度 kernel 边界。
• 把量化字节收益当端到端收益：W4 权重读少了不代表 TPOT 直接 4×；scale 读取、unpack/dequant、sampler、KV scan、调度都会稀释收益。
• 只看 attention FLOP：decode attention 常先受 KV cache 字节限制；MQA/GQA/KV quant 的收益来自少读 HBM，不是 Tensor Core 峰值。
• Paged KV 只调 page size：页大减少 block table 读取但增加内部碎片，页小更灵活但 table/L2 压力更高；还要看物理页是否碎片化、prefix alias 是否跨 GPU 迁移。
• 把 KV quant 当成无损压缩：KV cache 量化会改变 attention score/value 累加，必须报告质量回归、scale granularity 和长上下文稳定性。
• 只报 prefix cache hit rate：命中率高不代表 TTFT 降低；如果命中页仍被复制或跨 GPU 迁移，HBM/PCIe/NVLink 流量会吞掉收益。
• 忽略 ragged batch 浪费：continuous batching 里平均活跃长度远小于 Tmax 时，padded scan 会把 HBM 带宽花在零 padding 和 inactive slot 上。
• 只调 MoE GEMM：expert MLP 很快也可能被 token dispatch/gather、capacity padding、跨 GPU expert all-to-all 拖住端到端。
• 把采样尾部当免费：temperature、repetition penalty、mask、top-p 和随机数如果拆成多个小 kernel，再加每 token 回读，会在低 batch decode 中制造明显 p99 抖动。
• 每 token 阻塞回读：CPU 为了拿一个 token 调 cudaMemcpy 或 cudaDeviceSynchronize，会把 GPU decode 变成“跑一下、等一下”。
• 以为 cudaMemcpyAsync 天然异步：host buffer 不是 pinned、马上读取 D2H 结果、或者全进默认 stream 时，timeline 仍会表现成阻塞拷贝。
• 盲目融合小算子：fusion 能减少 launch 和 global traffic，但也可能增加寄存器、降低 occupancy、引入 spill；要用 ncu 验证。
• 默认 LUT 一定更快：RoPE LUT 能减少 SFU，但如果上下文很长、访问乱序或 L2 放不下，LUT 读取也可能变成新的 memory-bound 热点。
• Graph 捕获粒度不对：每个请求都重新 capture/instantiate 通常会变慢；应按稳定 shape 分桶复用 graph，并保证 buffer 地址生命周期稳定。
• 只报 speculative speedup：spec decode 必须同时报告 accept rate、accepted tokens/s、质量回归和 rejected-token verify waste。
• hot path 分配显存：cudaMalloc/cudaFree 可能触发同步和 allocator 锁；decode loop 里反复分配 workspace 会直接制造 p99 抖动。
• 忽略 GPU 拓扑：同样是 8 卡，NVLink 全互联、双 socket PCIe、跨节点 IB 的通信瓶颈完全不同；profile 文件必须附 nvidia-smi topo -m。
• 把 TP 通信只看成带宽问题：all-reduce 慢可能来自跨 root complex、rank 放置、TP size 过大、同步点太早，或者没有和下一层 compute overlap。
• NVTX range 太粗：只标 forward 没意义；至少标 phase、layer、attention/MLP/sampling。
• 把理论峰值当承诺：官方 Tensor Core 峰值常有 sparse/dense 口径差异，真实 kernel 还受内存、指令、同步、调度影响。
• 忽略线上约束：最优单 kernel 不一定最优服务；continuous batching、KV pool、NCCL、功耗和尾延迟都可能改变结论。

下一章导览

到这里，教程主线已经从 CUDA 入门走到 mini LLM，再补上了工业 profiling 方法。 后续继续扩展时，最自然的方向是两个：一是把第 14 章 capstone 改造成带 KV cache、fp16/WMMA、CUDA Graph 的 decode engine； 二是新增多 GPU / NCCL / tensor parallel 的推理章节，把单卡 kernel 优化连接到集群级 token throughput。