第 8 章 · 性能分析与异步执行

学习目标

• 会用 Nsight Compute（kernel 内部）和 Nsight Systems（系统级 timeline）
• 读懂 roofline、内存吞吐、stall reason 等关键指标
• 用 CUDA streams 让 H2D / 计算 / D2H 三阶段并行
• 用 CUDA Graphs 把"成千上万次 launch"压成一次

前置知识

已完成 Ch05–07，能区分计算、访存和同步瓶颈，并会使用 GpuTimer。

核心概念

8.1 Roofline 模型：在哪个屋顶下？

性能分析的第一原则：先看你受限于什么。

graph LR
    A["算术强度 (FLOP/Byte)"] --> B{"比较"}
    B -->|"低"| M["Memory-bound
瓶颈在带宽"]
    B -->|"高"| C["Compute-bound
瓶颈在算力"]
    M --> M1["优化方向:
coalescing
shared mem 复用
量化降位宽"]
    C --> C1["优化方向:
Tensor Core
算子融合
增大 batch"]
    style M fill:#f3f1e8,stroke:#a86420
    style C fill:#f3f1e8,stroke:#2f5d3a

LLM 推理（小 batch）通常是 memory-bound——所有的工程套路（KV cache 优化、量化、PagedAttention）都为了减少访存。 LLM 训练 / 大 batch 推理偏 compute-bound——所以 Tensor Core / FP8 是关键。

8.2 Nsight 工具链

典型工作流

# 1. nsys 找 "kernel 占多少时间 / 哪里有 idle / H2D 重叠了吗"
nsys profile --stats=true --force-overwrite=true -o report.nsys-rep ./my_app
nsys stats report.nsys-rep     # 命令行表格

# 2. ncu 钻到具体某个 kernel 看为什么慢
ncu --set full --kernel-name matmul_tiled -o report.ncu-rep ./my_app
ncu-ui report.ncu-rep          # 图形界面看 roofline、source counter

8.3 关键指标速读

关键代码

8.4 CUDA Streams — H2D 与计算重叠

默认 stream（NULL stream）上所有操作串行。开多个 stream，CUDA 才能并发执行：

sequenceDiagram
    participant H as Host
    participant S1 as Stream_1
    participant S2 as Stream_2
    Note over H,S2: Async chunk 切片，每片走独立 stream
    H ->> S1: H2D chunk 0
    H ->> S2: H2D chunk 1
    S1 ->> S1: kernel chunk 0
    S2 ->> S2: kernel chunk 1
    S1 ->> H: D2H chunk 0
    S2 ->> H: D2H chunk 1

实现要点：

1. host buffer 必须是 pinned（cudaMallocHost），否则 cudaMemcpyAsync 退化成同步
2. 用 cudaMemcpyAsync，传 stream 参数
3. kernel 启动加第 4 个尖括号参数：kernel<<<g, b, 0, streamX>>>(...)

cudaStream_t streams[4];
for (int i = 0; i < 4; ++i) cudaStreamCreate(&streams[i]);

for (int s = 0; s < 4; ++s) {
    int off = s * chunk;
    cudaMemcpyAsync(d + off, h + off, chunk * 4, cudaMemcpyHostToDevice, streams[s]);
    my_kernel<<<g, b, 0, streams[s]>>>(d + off, chunk);
    cudaMemcpyAsync(h + off, d + off, chunk * 4, cudaMemcpyDeviceToHost, streams[s]);
}
for (int s = 0; s < 4; ++s) cudaStreamSynchronize(streams[s]);

实测见 multi_stream.cu。多 stream 是否加速取决于 H2D、kernel、D2H 三段时间是否接近，以及拷贝引擎和 SM 是否能同时工作；用 Nsight Systems 看 timeline 才能下结论。

8.5 CUDA Graphs — 干掉 launch overhead

kernel launch 有固定开销，具体数字随驱动、CPU、功耗状态和采集方式变化。LLM 推理一个 token 要触发很多小 kernel 时，这部分开销会在 timeline 上形成可见空洞，对小 model 和低 batch decode 尤其明显。

CUDA Graphs 思路：把一连串 launch "录制"成一个图对象，重放时只付一次 launch 开销。

cudaGraph_t graph;
cudaGraphExec_t exec;

// 1) 进入 capture 模式
cudaStreamBeginCapture(s, cudaStreamCaptureModeGlobal);
for (int i = 0; i < 1000; ++i) tiny_kernel<<<g, b, 0, s>>>(d, N);
cudaStreamEndCapture(s, &graph);

// 2) 编译为可执行图
cudaGraphInstantiate(&exec, graph, nullptr, nullptr, 0);

// 3) 反复执行 (每次只 1 次 launch overhead)
cudaGraphLaunch(exec, s);
cudaStreamSynchronize(s);

建议用 cuda_graph_demo.cu 采集本机结果：记录 graph 前后的 kernel 数、CPU launch gap、GPU idle gap 和端到端时间，而不是套用固定加速比。

运行结果

用 Nsight Systems 核对 H2D、kernel、D2H 是否真实重叠，并区分 eager launch 与 CUDA Graph replay。

自检清单

练习题

1. 用 nsys profile 跑 multi_stream.cu，截图 timeline，标出 H2D / kernel / D2H 的重叠区。
2. 把 cuda_graph_demo.cu 改成 capture 一个"矩阵乘 → softmax → 矩阵乘"序列（用前面章节的 kernel）。
3. 用 ncu 跑 Ch6 的 matmul_tiled，记录 Achieved Occupancy 和 L1 Hit Rate。
4. 给 vec_add（Ch3）改成 4 stream 版本，比对加速比。

8.8 工业实战：Nsight 工作流、流优先级、生产 profile 套路

8.8.1 Nsight Systems：第一步永远是看 timeline

定位性能问题的正确顺序是 "先大后小"：先用 nsys 看整体 timeline 是否有 GPU idle gap、CPU 阻塞、stream 是否 overlap，再用 ncu 钻具体 kernel。新手通常上来就用 ncu 调单 kernel，结果发现整体瓶颈不在那个 kernel。

# 1) 录制 timeline (限制时长避免文件过大)
nsys profile \
    --trace=cuda,nvtx,osrt,cudnn,cublas \
    --duration=10 \
    --force-overwrite=true \
    -o report.nsys-rep \
    ./my_llm_server

# 2) 命令行看 summary
nsys stats report.nsys-rep
# 关注 "CUDA Kernel Statistics" 表 — 哪个 kernel 占时间最多

# 3) GUI 看 timeline (本地 Mac/Win 用 Nsight Systems Desktop 打开)
# 关键视角:
#  - GPU 行: 看 kernel 之间有没有 gap (= idle)
#  - CUDA HW row: H2D/D2H 跟 kernel 是否 overlap
#  - CPU 行: 主线程是否被 cudaMemcpy 阻塞 (说明没用 pinned + async)

8.8.2 NVTX 注解 — 让 timeline 可读

不打 NVTX 标记，timeline 上一堆 mykernel_v2 / mykernel_v3 完全看不出哪是 attention、哪是 MLP。生产代码必须在每个逻辑段加 NVTX range：

#include <nvtx3/nvToolsExt.h>

void forward(...) {
    nvtxRangePushA("layer_0");
    {
        nvtxRangePushA("attention");
        nvtxRangePushA("qkv_proj");   qkv_proj_kernel(...);   nvtxRangePop();
        nvtxRangePushA("flash_attn"); flash_attn_kernel(...); nvtxRangePop();
        nvtxRangePushA("out_proj");   out_proj_kernel(...);   nvtxRangePop();
        nvtxRangePop();  // attention
        nvtxRangePushA("mlp"); /* ... */ nvtxRangePop();
    }
    nvtxRangePop();  // layer_0
}

PyTorch 已经把 NVTX 集成进 torch.profiler，会自动给每个 op 加标记。

8.8.3 Nsight Compute：钻进单 kernel

找到瓶颈 kernel 后用 ncu 看为什么慢。关键 set：

# 完整指标 (最慢, 但信息最全, 一次抓清)
ncu --set full --target-processes all -o report.ncu-rep ./app

# 只抓特定 kernel (按名字匹配)
ncu --set full --kernel-name regex:".*flash_attn.*" -o fa.ncu-rep ./app

# 仅看 roofline + memory
ncu --section MemoryWorkloadAnalysis --section SpeedOfLight ./app

# 一次只跑 N 次 (避免长 trace)
ncu --launch-count 1 --launch-skip 100 ./app

用 Nsight Compute UI 打开 .ncu-rep 文件，先看 "Speed Of Light" 节——两个百分比：

• SM % = 算力利用率。低 → compute-bound 失败 / 寄存器spill / divergence
• Memory % = HBM 带宽利用率。低 → 访存模式差 / cache miss / 算术强度低
• 两者都低（< 30%）→ kernel 在等同步、stream serialization、launch 开销过大

8.8.4 必看的 10 个 ncu 指标

8.8.5 PyTorch / Python 侧的 profile

不是所有人都写裸 CUDA。LLM 工程师常用 PyTorch profiler：

import torch
from torch.profiler import profile, schedule, tensorboard_trace_handler

with profile(
    activities=[torch.profiler.ProfilerActivity.CPU,
                torch.profiler.ProfilerActivity.CUDA],
    schedule=schedule(wait=1, warmup=2, active=3),
    on_trace_ready=tensorboard_trace_handler('./log'),
    record_shapes=True,
    with_stack=True,
) as prof:
    for step in range(10):
        model(inputs)
        prof.step()

# 跑完用 tensorboard 看, 或者 chrome://tracing 打开 .json

8.8.6 流优先级 — 让推理请求插队

多请求并发推理时，新来的低延迟请求应该比"大 batch background 任务"优先。CUDA stream 支持优先级：

int low, high;
cudaDeviceGetStreamPriorityRange(&low, &high);
// low = 0, high = -1 (典型, 越小越高优先级, 看你 GPU)

cudaStream_t s_high, s_low;
cudaStreamCreateWithPriority(&s_high, cudaStreamNonBlocking, high);
cudaStreamCreateWithPriority(&s_low,  cudaStreamNonBlocking, low);

// 新来的 prompt 走 s_high, 大 batch 走 s_low
// 硬件调度器在 SM 上抢占 (会 preempt 当前 block 边界)

典型应用：vLLM 实现"chat 优先于 batch"调度，TensorRT-LLM 把 prefill 与 decode 放不同优先级。

8.8.7 production profile 工作流

1. 定位：nsys 录 timeline → 看 stats 找耗时 top kernel
2. 分析：ncu --set full 抓 top kernel → 看 SoL 百分比定性是 compute / memory / latency bound
3. 对症：
   • compute-bound → 换 Tensor Core / 提高 ILP / 减 divergence
   • memory-bound → coalesce / tile / 量化降位宽
   • latency-bound (两者都低) → 检查 occupancy, launch 频率, sync 次数
4. 验证：再录 nsys 对比 timeline 是否变窄
5. 回归：把改动放进 CI 性能门禁（防止后续 commit 反向优化）

常见误区：

• 只看一个 size 的 benchmark → 改完发现其他 size 性能反退（CUTLASS autotune 就在防这个）
• 没 warmup 第一次 timing → 包含 CUDA context init、cuBLAS 加载、Tensor Core 预热
• 用 std::chrono 测 GPU 时间 → 没 sync 时返回的是 launch 时间不是 kernel 时间

8.9 研究前沿（2025-2026）：CUDA Graph 进化与新 profile 工具

8.9.1 CUDA Graph 条件节点（Conditional Nodes，CUDA 12.4+）

原始 CUDA Graph 是静态的：capture 时怎样，重放就怎样。LLM 推理里"循环到 EOS 才停"这种动态控制流没法 capture 到一张 graph 里——只能展开成 max_tokens 张 graph 或者分多次 launch。

CUDA 12.4 引入 conditional graph nodes：

cudaGraphConditionalHandle handle;
cudaGraphConditionalHandleCreate(&handle, graph, /*default=*/0, 0);

// 添加一个 "while (handle != 0)" 节点
cudaGraphNode_t while_node;
cudaGraphNodeParams params = {};
params.type = cudaGraphNodeTypeConditional;
params.conditional.handle = handle;
params.conditional.type   = cudaGraphCondTypeWhile;
params.conditional.size   = 1;
cudaGraphAddNode(&while_node, graph, deps, n_deps, &params);

// 在循环 body kernel 里更新 handle:
__device__ void check_eos(int token, cudaGraphConditionalHandle h) {
    if (token == EOS_ID) cudaGraphSetConditional(h, 0);   // 退出循环
    else                 cudaGraphSetConditional(h, 1);
}

意义：一次 launch 能跑完整个生成过程，包括动态停止判断。端到端收益取决于每 token 小 kernel 数量、batch bucket 稳定性和 host scheduler 开销，需要用 decode-only trace 验证。

8.9.2 Nsight Compute 2025 — fp4 / fp8 Roofline 支持

2025 版 Nsight Compute 关键更新：

• fp8 / fp4 算力作为独立 roofline 轴（之前只有 fp16/fp32）
• TMEM occupancy 指标（Blackwell）
• wgmma / TCGEN05 流水利用率细分
• Auto-tuning advisor：分析后给出"改 tile / 改 launch_bounds / 加 cp.async" 等具体建议
• SQLite report 格式（之前是 protobuf），方便用 Python 自动分析

8.9.3 PyTorch Profiler + Holistic Trace Analysis

Meta 开源 HTA (Holistic Trace Analysis) 把 PyTorch profiler trace 自动化分析：

from hta.trace_analysis import TraceAnalysis
ta = TraceAnalysis(trace_dir="./logs")
ta.get_temporal_breakdown()        # GPU compute / comm / idle 占比
ta.get_gpu_kernel_breakdown()      # 单 kernel top
ta.get_communication_comp_overlap()# NCCL 是否跟 compute overlap
ta.get_idle_time_breakdown()       # SM idle 是因为什么 stall

对大规模训练 / 推理服务来说，能省下手工看 timeline 几小时的工作。

8.9.4 持久化 kernel + CUDA Graph 混合模式

vLLM、TRT-LLM、SGLang 在 2024-2025 大量采用的模式：

graph LR
    A["请求队列"] --> B["Scheduler
(host)"]
    B --> C["Persistent kernel
常驻全部 SM"]
    B --> D["CUDA Graph cache
各 batch_size 一份"]
    C --> E["attention / fused 算子"]
    D --> F["batch=1, 2, 4, 8, 16, 32"]
    style C fill:#f3f1e8,stroke:#8b1538
    style D fill:#f3f1e8,stroke:#2f5d3a

• Persistent kernel 永远在线，从 host 的 work queue 拉任务（用于动态调度的 attention / sampling）
• CUDA Graph 预 capture 各 bucket 的"标准 decode 一步"
• scheduler 每步选 batch_size bucket → launch 对应 graph
• graph 内部参数（KV 位置、当前 token）用 cudaGraphExecKernelNodeSetParams 动态更新

端到端是否更快要用线上形状复现：CUDA Graph 主要减少 host launch gap，persistent kernel 主要减少调度往返，两者收益会被 batch、KV 读带宽、sampling 和网络通信稀释。

8.9.5 NVSHMEM / 多 GPU 异步通信

NCCL 之外的 2024-2026 新选择：NVSHMEM（NVIDIA 推的 PGAS 风格库），让 kernel 内直接 put / get 跨 GPU 内存，跟 NVLink 硬件深度绑定：

__global__ void cross_gpu_attn(...) {
    /* compute local */
    nvshmem_putmem_block(remote_addr, local_addr, bytes, peer_pe);
    nvshmem_barrier_all();
    /* compute global */
}

比 NCCL 集合通信粒度更细，对 Ring Attention / Stripe Attention / Expert Parallel 等"边算边通信"模式收益大。TensorRT-LLM 长 context 推理已在用。

8.9.6 端到端性能调优新范式

2025 的"性能调优"已经从"调单 kernel" 转向"端到端 trace 驱动"：

1. 跑 nsys + HTA 录 timeline 跨 GPU / CPU / 通信
2. 找关键路径（critical path）—— 通常是某 GPU 的 kernel + NCCL 阻塞
3. 优化关键路径上的 1-2 个 kernel（用 ncu 钻），其他保持现状
4. 验证 timeline 是否被压缩

关键 KPI：单 GPU 不是 throughput 而是"在关键路径上的占比"。这是大规模分布式训练 / 推理性能工程的工作模式。

常见坑

• 用 pageable host 内存做 async memcpy → 实际是同步，看 nsys 才发现重叠"消失"
• graph capture 时 kernel 参数指针不能改 → 改了需要 cudaGraphExecUpdate
• 多 stream 共享同一 device buffer 又不加事件同步 → race

到此阶段 B 结束。你已经具备了写高性能 kernel 的全部工具。下一章我们用这些工具一气呵成把 GEMM 推到接近硬件极限。

8.11 CUDA 官方手册精讲（CUDA Programming Guide 13.2（核验：2026-07-20））

PDL、Graph Lifecycle、Conditional Nodes、Stream-ordered Alloc、Green Context

Stream 优先级 + Programmatic Dependent Launch (PDL，sm_90+)

8.4 让 H2D 与计算并发；现在 Hopper 推进一步：同 stream 上的两个 kernel 也能"前一个还没完，后一个就先启动"。这叫 PDL（Programmatic Dependent Launch）。

动机：每个 kernel 都有"前奏"（preamble）—— 比如读常量、清 accumulator、做 epilogue 之前的 norm 准备。 这些前奏可以与上游 kernel 的结尾并行。

    gantt
        title PDL 前后对比
        dateFormat X
        axisFormat %s
        section 传统串行
        primary_kernel       :a1, 0, 60
        secondary_kernel     :a2, 60, 100
        section 启用 PDL
        primary_kernel       :b1, 0, 60
        secondary_preamble   :b2, 50, 60
        secondary_main       :b3, 60, 90
    

__global__ void primary_kernel(...) {
        // 真正的计算
        do_work();

        // 在这里告诉硬件："我已经发出全部数据写，你可以开始 secondary 的前奏了"
        cudaTriggerProgrammaticLaunchCompletion();

        // 这之后还可以做尾部工作 (与 secondary 的前奏并行)
        finalize();
    }

    __global__ void secondary_kernel(...) {
        // 与 primary 完全独立的前奏
        setup_constants();
        init_accumulator();

        // 真正需要 primary 输出之前: 等 primary 数据可见
        cudaGridDependencySynchronize();

        // 现在可以读 primary 写过的内存
        use_primary_output();
    }

    // Launch:
    primary_kernel<<<g, b, 0, stream>>>();

    cudaLaunchAttribute attr;
    attr.id  = cudaLaunchAttributeProgrammaticStreamSerialization;
    attr.val.programmaticStreamSerializationAllowed = 1;
    cudaLaunchConfig_t cfg{};
    cfg.gridDim = g; cfg.blockDim = b; cfg.stream = stream;
    cfg.attrs = &attr; cfg.numAttrs = 1;
    cudaLaunchKernelEx(&cfg, secondary_kernel);

CUDA Graphs 完整生命周期：Explicit API vs Stream Capture vs Update

8.5 给了 stream capture 的简单写法。生产代码常用混合三种构建方式，各有适用场景。

方式 ②：Explicit API（更灵活，LLM 推理常用）

cudaGraph_t g;
    cudaGraphCreate(&g, 0);

    // 创建一个 kernel 节点
    cudaGraphNode_t n_qkv, n_attn, n_proj;
    cudaKernelNodeParams kp_qkv{};
    kp_qkv.func = (void*)qkv_kernel;
    kp_qkv.gridDim = grid; kp_qkv.blockDim = blk;
    void* args_qkv[] = { ... };
    kp_qkv.kernelParams = args_qkv;
    cudaGraphAddKernelNode(&n_qkv, g, nullptr, 0, &kp_qkv);

    // attn 依赖 qkv
    cudaGraphAddKernelNode(&n_attn, g, &n_qkv, 1, &kp_attn);

    // proj 依赖 attn
    cudaGraphAddKernelNode(&n_proj, g, &n_attn, 1, &kp_proj);

    // 实例化为可执行图
    cudaGraphExec_t exec;
    cudaGraphInstantiate(&exec, g, nullptr, nullptr, 0);

方式 ③：每步只改参数，复用 graphExec（最省时间）

LLM decode 每步只改 (KV 位置 token_pos, 当前 token id)。如果每步重新 capture / instantiate，会把图构建成本放进热路径；用 update 可以复用已实例化的 graphExec。两者差距请在目标机器上测。

// 第一次：建好图 + instantiate
    // (省略 build...)
    cudaGraphInstantiate(&exec, g, nullptr, nullptr, 0);

    // 每步 decode：只改某个 kernel 节点的参数
    cudaKernelNodeParams new_params{};
    new_params.func = (void*)attn_kernel;
    new_params.gridDim = grid;
    void* new_args[] = { &KV_ptr, &token_pos_NEW, ... };
    new_params.kernelParams = new_args;

    // 关键 API：原地改 exec 里某个节点的参数
    cudaGraphExecKernelNodeSetParams(exec, n_attn, &new_params);

    // Replay after updating parameters.
    cudaGraphLaunch(exec, stream);

    graph LR
        A["第 1 次:
Build graph"] --> B["Instantiate
TODO(on GPU)"]
        B --> C["Launch
TODO(on GPU)"]
        C --> D["第 N 次:
SetKernelNodeParams
TODO(on GPU)"]
        D --> E["Re-Launch
TODO(on GPU)"]
        E --> D
        style B fill:#f3f1e8,stroke:#a86420
        style D fill:#f3f1e8,stroke:#2f5d3a
    

CUDA Graphs 条件节点（Conditional Nodes，CUDA 12.4+）

教程 8.9.1 给了 cudaGraphConditionalHandle 用法的轮廓。这里给四种条件节点的完整 API 表 + Device Graph Launch 关联。

简化版 LLM decode-loop graph 骨架：

cudaGraph_t graph;
    cudaGraphCreate(&graph, 0);

    // 1) 主输入节点：初始化 token_count = 0
    // (省略 memset 节点)

    // 2) 创建 conditional handle, 默认值 = 1 (进入循环)
    cudaGraphConditionalHandle keep_going;
    cudaGraphConditionalHandleCreate(&keep_going, graph, /*defaultLaunchValue=*/1,
                                     cudaGraphCondAssignDefault);

    // 3) 添加 While 节点：body 是子图
    cudaGraphNode_t while_node;
    cudaGraphNodeParams while_params{};
    while_params.type                  = cudaGraphNodeTypeConditional;
    while_params.conditional.handle    = keep_going;
    while_params.conditional.type      = cudaGraphCondTypeWhile;
    while_params.conditional.size      = 1;
    cudaGraphAddNode(&while_node, graph, /*deps=*/nullptr, 0, &while_params);

    // 4) 取出 while 的 body 子图, 往里加 kernel
    cudaGraph_t body = while_params.conditional.phGraph_out[0];
    add_attention_kernel(body, ...);
    add_mlp_kernel(body, ...);
    add_sample_kernel(body, ..., &new_token);
    // 5) 在 body 最后加一个 kernel: 读 new_token, 如果 == EOS 则把 handle 设为 0
    add_check_eos_kernel(body, &new_token, keep_going);

    // 6) Instantiate + Launch
    cudaGraphExec_t exec;
    cudaGraphInstantiate(&exec, graph, nullptr, nullptr, 0);
    cudaGraphLaunch(exec, stream);  // 一次 launch, 跑完整个 decode loop!

在 body 内的 kernel 用 device 函数更新 handle：

__device__ void cudaGraphSetConditional(cudaGraphConditionalHandle handle,
                                            unsigned int value);

    __global__ void check_eos(int* new_token, cudaGraphConditionalHandle h) {
        if (threadIdx.x == 0 && blockIdx.x == 0) {
            cudaGraphSetConditional(h, (*new_token == EOS_ID) ? 0 : 1);
        }
    }

Stream-Ordered Memory Allocator (cudaMallocAsync) — 流内分配释放

传统 cudaMalloc 与 cudaFree 是 全局同步 操作——会等所有 stream 上的 GPU 工作完成。在 LLM 推理里，每步 decode 都临时申请 KV / workspace 时，这会成为主要瓶颈。

解法：cudaMallocAsync / cudaFreeAsync —— 让分配和释放都挂到指定 stream 上，与 kernel 一样按 stream order 调度，不阻塞其他流：

cudaStream_t s;
    cudaStreamCreate(&s);

    void* d_workspace;
    cudaMallocAsync(&d_workspace, 64 * 1024 * 1024, s);  // 流序分配, 不阻塞

    my_kernel<<<g, b, 0, s>>>(d_workspace, ...);          // 用

    cudaFreeAsync(d_workspace, s);                          // 流序释放, 不阻塞

底层有一个 memory pool 复用机制：第二次 malloc 同样大小时，driver 直接拿回上次释放的 chunk，不进 OS。

调 release threshold 控制"留多少缓存"

cudaMemPool_t pool;
    cudaDeviceGetDefaultMempool(&pool, device);

    // 默认 0 → 每次 sync 都把所有空闲块还给 OS
    // 设大 (如 1 GB) → 池里保留 1 GB 不还, 下次再 alloc 就是池命中
    uint64_t threshold = 1ULL * 1024 * 1024 * 1024;
    cudaMemPoolSetAttribute(pool, cudaMemPoolAttrReleaseThreshold, &threshold);

    // 或者完全不还
    uint64_t never = UINT64_MAX;
    cudaMemPoolSetAttribute(pool, cudaMemPoolAttrReleaseThreshold, &never);

统计：池吃了多少显存

struct PoolStats { uint64_t reserved, used; };
    PoolStats st;
    cudaMemPoolGetAttribute(pool, cudaMemPoolAttrReservedMemCurrent, &st.reserved);
    cudaMemPoolGetAttribute(pool, cudaMemPoolAttrUsedMemCurrent,     &st.used);
    // reserved = 池占的物理内存; used = 已分配未释放的

Green Contexts：把 SM 切给"高优先级请求"用

8.8.6 的 stream 优先级是 "调度顺序" 提示，硬件并不保证立即抢占。Green Context（CUDA 12.4+，Hopper 起 driver API；CUDA 13.1+ runtime API）走得更彻底：直接把 GPU 物理切分，划一片专用 SM 给某个 stream，谁也抢不走。

#include <cuda.h>  // driver API

    // 1) 看 GPU 有多少 SM
    CUdevice dev; cuDeviceGet(&dev, 0);
    int n_sms;
    cuDeviceGetAttribute(&n_sms, CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT, dev);

    // 2) 创建 "资源描述" — 想要 16 SM
    CUdevResource sm_resource;
    cuDeviceGetDevResource(dev, &sm_resource, CU_DEV_RESOURCE_TYPE_SM);

    CUdevResource split[2]; unsigned int n_groups = 1;
    unsigned int min_sm_count = 16;
    cuDevSmResourceSplitByCount(split, &n_groups, &sm_resource,
                                /*remaining=*/&split[1], 0, min_sm_count);
    // split[0] = 16 SM (给高优 stream), split[1] = 剩下 SM

    // 3) 创建 Green Context
    CUgreenCtx gc_critical, gc_bulk;
    CUdevResourceDesc desc1; cuDevResourceGenerateDesc(&desc1, &split[0], 1);
    cuGreenCtxCreate(&gc_critical, desc1, dev, CU_GREEN_CTX_DEFAULT_STREAM);

    CUdevResourceDesc desc2; cuDevResourceGenerateDesc(&desc2, &split[1], 1);
    cuGreenCtxCreate(&gc_bulk, desc2, dev, CU_GREEN_CTX_DEFAULT_STREAM);

    // 4) 在 Green Context 上建 stream + launch
    CUstream s_critical, s_bulk;
    cuGreenCtxStreamCreate(&s_critical, gc_critical, CU_STREAM_NON_BLOCKING, 0);
    cuGreenCtxStreamCreate(&s_bulk,     gc_bulk,     CU_STREAM_NON_BLOCKING, 0);

    // 现在: 低延迟请求 → s_critical (永远有 16 SM 待命)
    //      大 batch → s_bulk (其他 SM)

验证 Green Context 时，记录 critical kernel 从提交到开始执行的等待时间、bulk kernel 的吞吐损失，以及 SM 分配比例。 目标是用 Nsight Systems 证明 critical stream 有独立 SM 资源，而不是只看平均延迟。

NVTX 高级用法：colored ranges、payloads、Domain 隔离

教程 8.8.2 用了 nvtxRangePushA / Pop。生产代码用得更多的是带颜色、带 payload、带 domain 的 v3 API，让 Nsight timeline 直接可读。

#include <nvtx3/nvToolsExtCuda.h>
    #include <nvtx3/nvToolsExt.h>

    // 1) Domain 隔离：把不同模块的 range 染色到不同行
    nvtxDomainHandle_t dom_attn = nvtxDomainCreateA("Attention");
    nvtxDomainHandle_t dom_mlp  = nvtxDomainCreateA("MLP");

    // 2) Event attributes：颜色 + payload (用于 timeline 上标 batch_size / seq_len)
    nvtxEventAttributes_t attr{};
    attr.version       = NVTX_VERSION;
    attr.size          = NVTX_EVENT_ATTRIB_STRUCT_SIZE;
    attr.colorType     = NVTX_COLOR_ARGB;
    attr.color         = 0xFFFF6B6B;          // 橙红
    attr.messageType   = NVTX_MESSAGE_TYPE_ASCII;
    attr.message.ascii = "FlashAttention";
    attr.payloadType   = NVTX_PAYLOAD_TYPE_UNSIGNED_INT64;
    attr.payload.ullValue = ((uint64_t)batch_size << 32) | seq_len;   // 编码 2 个值

    nvtxRangeId_t r = nvtxDomainRangeStartEx(dom_attn, &attr);
    flash_attention_kernel(...);
    nvtxDomainRangeEnd(dom_attn, r);

Pytorch 用户用 torch.cuda.nvtx：

import torch.cuda.nvtx as nvtx

    with nvtx.range("attention", color="red"):
        out = self.attn(x)

    # 或者带 payload (PyTorch 2.4+):
    nvtx.range_push(f"decode_step_{step_id}_bsz_{bsz}")
    forward(...)
    nvtx.range_pop()

Lazy Loading：减少程序启动时间的免费午餐

每个 CUDA 程序都要加载 .cubin / .ptx 模块到 GPU。大库（cuBLAS, cuDNN, Triton）可能包含很多 kernel。 Lazy Loading：等真正用到某个 kernel 时再加载，通常能降低冷启动成本，但第一次触发某个 kernel 会把加载成本推迟到请求路径上。

CUDA 12.2 起 Linux 默认开启，CUDA 12.3 起 Windows 默认开启。什么都不用改就生效—— 只要保证：

1. CUDA Toolkit ≥ 11.7（runtime 静态链入应用）
2. Driver ≥ 515（lazy loading 内核支持）

# 显式确认 / 切换
    export CUDA_MODULE_LOADING=LAZY   # 默认 (12.2+)
    export CUDA_MODULE_LOADING=EAGER  # 强制全部预加载

查看是否启用：

#include <cuda.h>
    CUmoduleLoadingMode mode;
    cuInit(0);
    cuModuleGetLoadingMode(&mode);
    printf("Mode: %s\n", mode == CU_MODULE_LAZY_LOADING ? "lazy" : "eager");

三个陷阱

for (int i = 0; i < 3; ++i) my_kernel<<<...>>>(...);  // warmup
    cudaDeviceSynchronize();
    // 真正开始计时

Memory Synchronization Domains：减少 NCCL × 计算的 fence 干扰

多卡训练里常见情况：本卡 GPU 做计算 kernel，同时 NCCL kernel 通过 NVLink 与远端通信。 两个 kernel 完成时各自要做 memory fence，fence 可能互相等对方的内存事务，导致计算 kernel 在通信重叠场景下出现额外等待。

Hopper（sm_90）起引入 Memory Synchronization Domain：每个 kernel launch 标个 domain ID，fence 只等同 domain 的写。

// 给 NCCL kernel 打 remote domain
    cudaLaunchAttribute attr;
    attr.id  = cudaLaunchAttributeMemSyncDomain;
    attr.val.memSyncDomain = cudaLaunchMemSyncDomainRemote;

    cudaLaunchConfig_t cfg{};
    cfg.gridDim = g; cfg.blockDim = b; cfg.stream = nccl_stream;
    cfg.attrs = &attr; cfg.numAttrs = 1;
    cudaLaunchKernelEx(&cfg, ncclAllReduce_kernel, ...);

    // 计算 kernel 用默认 domain — 不用改
    matmul_kernel<<<g, b, 0, compute_stream>>>(...);

或者直接给 stream 配 domain map（更省事）：

cudaLaunchAttributeValue map{};
    map.memSyncDomainMap.default_ = 0;       // 计算流走 phys domain 0
    map.memSyncDomainMap.remote   = 1;       // 通信流走 phys domain 1
    cudaStreamSetAttribute(nccl_stream, cudaLaunchAttributeMemSyncDomainMap, &map);

下一章导览

第 9 章进入 LLM 的主计算核心 GEMM，并把手写 kernel 与 NVIDIA 库基线放在同一证据链中。