第 5 章 · 内存层级

学习目标

• 背下 GPU 的内存层级金字塔（register / shared / L1 / L2 / global / host）
• 掌握合并访问规则：什么样的访问模式能让 warp 一次取够 128B
• 会用 __shared__、__constant__
• 知道 cudaMallocManaged（unified memory）什么时候能用、什么时候坑

前置知识

已完成 Ch04，理解 warp、occupancy，以及 host/device 指针和显式数据搬运。

核心概念

5.1 内存金字塔

graph TD
    R["Registers
~1 cycle, thread 私有
256 KB / SM"]
    S["Shared Memory / L1
~30 cycles, block 内共享
192 KB / SM (A100)"]
    L2["L2 Cache
~200 cycles, 全 GPU 共享
40 MB (A100)"]
    G["Global / HBM
~500 cycles
40-80 GB"]
    H["Host RAM
~50000 cycles via PCIe"]
    R --> S --> L2 --> G --> H
    style R fill:#f3f1e8,stroke:#2f5d3a
    style S fill:#f3f1e8,stroke:#8b1538
    style G fill:#f3f1e8,stroke:#a86420

一句话总结："越靠上越快越小"。性能优化 = 让数据尽量待在上层。

关键代码

5.2 合并访问 (Memory Coalescing)

GPU 一次最少读 128B（= 32 个 float = 一个 warp 的份）。如果 warp 内 32 个 lane 恰好访问相邻 32 个 float，硬件合并为一次内存事务。如果它们访问跳跃的位置，就需要 32 次独立事务。

合并 vs 跳访

// ✅ COALESCED — warp 内 lane k 访问 in[base + k]
out[i] = in[i];

// ❌ STRIDED  — warp 内 lane k 访问 in[base + k*STRIDE]
out[i*STRIDE] = in[i*STRIDE];

// ❌ TRANSPOSED — 经典 row-major 数据按列访问就是 strided
out[col*M + row] = in[row*N + col];

跑 coalesce_vs_strided.cu，在你的目标 GPU 上填这张表：

5.3 Shared Memory: 片上 SRAM

__shared__ 关键字让你在 SM 的片上 SRAM 上开一块给整个 block 共享的数组。读写延迟接近寄存器，是把 global memory 数据驻留在芯片上重复使用的关键。

__global__ void block_sum_shared(const float* x, float* partial, int n) {
    __shared__ float sdata[256];          // 静态 shared，编译期定大小

    int tid = threadIdx.x;
    int gid = blockIdx.x * blockDim.x + threadIdx.x;

    sdata[tid] = (gid < n) ? x[gid] : 0;   // 1) 拉数据
    __syncthreads();                       // 2) 等齐

    for (int s = blockDim.x / 2; s > 0; s >>= 1) {  // 3) tree-reduce
        if (tid < s) sdata[tid] += sdata[tid + s];
        __syncthreads();
    }
    if (tid == 0) partial[blockIdx.x] = sdata[0];
}

• 静态大小：__shared__ float sdata[256];（编译期常量）
• 动态大小：extern __shared__ float sdata[];，然后 launch 时 kernel<<<g, b, N*sizeof(float)>>>
• 必须 __syncthreads()：保证全 block 都把数据填进 sdata 后才开始读，否则 race

5.4 Constant Memory

__constant__ 是 64 KB 大小的全局只读区，配有专门的 constant cache。最适合的访问模式：warp 内 32 lane 同时访问同一地址——硬件一次广播给 32 lane，0 额外开销。

__constant__ float c_coeff[16];   // 文件作用域声明

// host:
cudaMemcpyToSymbol(c_coeff, h_coeff, sizeof h_coeff);

// kernel: 所有 lane 读 c_coeff[k] → 广播
for (int k = 0; k < 16; ++k) acc += c_coeff[k] * x;

但如果 warp 内不同 lane 读不同的 constant 地址 → 序列化，反而比 global 慢。所以适用场景窄：小型 LUT、模型超参数、kernel 内不变的因子。

5.5 三种分配方式对比

跑 mem_modes.cu 实测三种差距。结论：

• 生产代码用 cudaMalloc。
• 需要 H2D/D2H 异步重叠时用 pinned（cudaMallocHost）。
• unified / zero-copy 仅用于原型，不要在性能敏感路径用。

运行结果

运行本章四个示例，先验证输出，再记录不同 stride、内存模式和 shared-memory 路径的实际带宽。

性能数据

自检清单

练习题

1. 修 01_coalesce_starter.cu：调换 row/col indexing 让 copy 变 coalesced。
2. 用 shared_demo.cu 改成把 256 个 float 求最大值而不是和。
3. 把 constant_demo.cu 里 c_coeff 改成 1024（超 4 KB），看是否还有加速？为啥？（提示：constant cache 只有 8 KB / SM。）
4. 给自己 GPU 测实际 vs 理论带宽比，记下结果——后面所有 kernel 的 "好不好" 都拿这个比。

5.8 工业实战：内存池、cp.async、显存碎片、NCCL

5.8.1 cudaMallocAsync — 解决 LLM 服务的显存碎片

问题：LLM 推理服务里，每个请求要分配几十 MB 的 KV cache，请求完后释放。频繁 cudaMalloc/cudaFree 在 hot path 上消耗几百微秒，且显存逐渐碎片化——明明 free 显存还有 20 GB 却分不出一块连续 5 GB。

CUDA 11.2+ 提供异步内存池（基于 stream-ordered memory allocator）：

// 创建 / 配置 memory pool
cudaMemPool_t pool;
cudaDeviceGetDefaultMemPool(&pool, 0);

// 设置 release threshold: 池中保留多少字节不归还 OS
size_t threshold = size_t(20) << 30;   // 20 GB
cudaMemPoolSetAttribute(pool, cudaMemPoolAttrReleaseThreshold, &threshold);

// 异步 malloc / free, 跟 stream 绑定
void* ptr;
cudaMallocAsync(&ptr, bytes, stream);     // 几乎 0 开销 (复用池)
// ... use ptr on stream ...
cudaFreeAsync(ptr, stream);                // 也几乎 0 开销

vLLM / TensorRT-LLM 都用这个。配合 PyTorch 用 torch.cuda.set_per_process_memory_fraction 和 PYTORCH_CUDA_ALLOC_CONF：

# PyTorch 推荐配置（避免 OOM + 碎片）:
export PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True,max_split_size_mb:512

5.8.2 cp.async — Ampere+ 的异步 shared 加载

Ch6 已经提过，这里看 PTX 级细节。普通 shared 加载：

// global -> register -> shared (两次访存事务)
LDG.E.SYS R0, [R2]      ; 从 global 读到 register
STS [R5], R0            ; 从 register 写到 shared

Ampere 的 cp.async：

// global -> shared 一步到位, 不经过 register, 异步
LDGSTS.E.BYPASS.LTC256B [R5], [R2]
// 之后 kernel 可继续做别的, 真要用 shared 时:
LDGSTS.WAIT

C++ 侧两种 API：

// 1) High-level: cuda::pipeline
#include <cuda/pipeline>
auto pipeline = cuda::make_pipeline();
pipeline.producer_acquire();
cuda::memcpy_async(&smem[0], &gmem[0], cuda::aligned_size_t<16>(64), pipeline);
pipeline.producer_commit();
pipeline.consumer_wait();
// ... use smem ...
pipeline.consumer_release();

// 2) Low-level: __pipeline_memcpy_async
__pipeline_memcpy_async(&smem[i], &gmem[i], sizeof(float4));
__pipeline_commit();
__pipeline_wait_prior(0);

收益要用 Nsight 验证：在 GEMM mainloop 里配 double-buffer 后，重点看 HBM 加载是否与 compute 重叠、Long Scoreboard 是否下降、Tensor Core 是否更连续地工作。不同 tile、寄存器压力和硬件代际下，端到端提升差异很大。

5.8.3 cudaMemPrefetchAsync — Unified Memory 的正确姿势

5.5 提到 unified memory（cudaMallocManaged）在生产中很少用，因为 page fault 慢。但如果非用不可（例如显存不够，要 host-managed swap），cudaMemPrefetchAsync 是关键：

float* m; cudaMallocManaged(&m, bytes);
// 提前把数据从 host 拉到 device, 避免 kernel 内 page fault
cudaMemPrefetchAsync(m, bytes, /*device=*/0, stream);
my_kernel<<<..., stream>>>(m);
// 用完拉回 host (eg. checkpoint)
cudaMemPrefetchAsync(m, bytes, cudaCpuDeviceId, stream);

NVIDIA Grace Hopper Superchip (GH200) 的 NVLink-C2C 让 unified memory 性能逼近显存，CPU+GPU 共享 480 GB——是为超大模型设计的，开始改变 unified 在生产中的地位。

5.8.4 跨 GPU：NCCL 集合通信

多 GPU 训练 / 张量并行推理必用 NCCL (NVIDIA Collective Communications Library)：

#include <nccl.h>
ncclComm_t comm;
ncclCommInitAll(&comm, n_gpus, devs);

// All-Reduce: 每张卡有 local result, 跨卡求和后所有卡都有 sum
ncclAllReduce(send_buf, recv_buf, count, ncclFloat, ncclSum, comm, stream);

// All-Gather: 每张卡有 1/N 数据, gather 完每张卡都有完整数据
ncclAllGather(send_buf, recv_buf, count, ncclFloat, comm, stream);

关键 pattern: 把 NCCL 调用放在独立 stream 上，跟计算 stream 并行：

cudaStream_t compute_s, comm_s;
cudaStreamCreate(&compute_s); cudaStreamCreate(&comm_s);

// 前向第 N 层算完, 异步把梯度 reduce, 同时算第 N+1 层
backward_layer_N<<<..., compute_s>>>(grad);
ncclAllReduce(grad, grad, ..., comm_s);    // overlap!
backward_layer_N1<<<..., compute_s>>>(...);

NCCL 拓扑感知（自动用 NVLink、忽略 PCIe 慢路径）。但你必须在同一个 process初始化所有 GPU 或用 NCCL_P2P_LEVEL 调优。详细见 NCCL 文档。

5.8.5 显存占用 debug 工具

# 看哪些进程占用 GPU 显存
nvidia-smi --query-compute-apps=pid,process_name,used_memory --format=csv

# 进入 Python 看 PyTorch 的分配
python -c "import torch; torch.cuda.memory._dump_snapshot('mem.pickle')"
# 然后用 https://pytorch.org/memory_viz 在线可视化

# 在 PyTorch 训练里加打印
print(torch.cuda.memory_summary())
print(f"allocated: {torch.cuda.memory_allocated()/1e9:.2f} GB")
print(f"reserved : {torch.cuda.memory_reserved()/1e9:.2f} GB")

典型显存泄漏症状：训练 N 步后 OOM，重启就好。9 成原因是 hold 住了 computation graph（debug 时 forget .detach()），或者把 tensor 加进 list 不释放。

5.8.6 LLM 推理的显存分布

7B 模型 fp16 推理，典型显存占用：

结论：A100-80G 跑 7B 服务约能容 32-64 并发；70B fp16 + KV 直接超过 80G，必须 TP 切到多卡或者 W4A16 量化。

5.9 研究前沿（2025-2026）：Blackwell TMEM 与 KV cache 革命

5.9.1 Tensor Memory (TMEM) — 内存层级新成员

Blackwell 把所有 SM 加了第六层内存。新的金字塔：

graph TD
    R["Registers
~1 cycle, thread 私有
256 KB / SM"]
    T["Tensor Memory (TMEM)
~5 cycles, SM 内 MMA 专用
256 KB / SM, sm_100+ 新增"]
    S["Shared Memory / L1
~30 cycles, block / cluster
228 KB / SM"]
    L2["L2 Cache
~200 cycles
50-80 MB"]
    G["Global / HBM
~500 cycles
B200 180 GB; GB200/GPU 186 GB; B300 288 GB"]
    H["Host RAM
~50000 cycles via PCIe / NVLink-C2C"]
    R --> T --> S --> L2 --> G --> H
    style T fill:#f3f1e8,stroke:#a86420

TMEM 的特殊性：

• 只能由 MMA 指令写入（普通 store 不行）
• 不参与一般运算，必须先 tcgen05.ld 读回 register 才能 ALU 运算
• 专为 fp8 / fp4 MMA accumulator 设计：fp4 GEMM 的 fp32 acc 占空间多，从 register 挪到 TMEM 释放了寄存器

普通 CUDA 开发者：通过 CUTLASS 4.6.1 / CuTe DSL（2026-07 快照） 自动使用，不需要直接写。

5.9.2 KV cache 压缩前沿（2024-2026）

LLM 推理的最大显存怪兽是 KV cache。2024-2026 出现了五种主流压缩思路：

实战组合：vLLM 默认用 PagedAttention + 可选 fp8/int4 KV 量化；SGLang 加上 RadixAttention；DeepSeek-V3 / V2 用 MLA + paged + fp8。多种技术叠加后，KV 显存从“按最大长度预留”转为“按真实 token 与共享前缀付费”，实际节省幅度必须按模型结构、上下文长度和请求分布测算。

5.9.3 PagedAttention 与显存碎片化的"解药"

vLLM 论文（SOSP 2023）后两年的演进：

• vAttention (Microsoft, 2024)：直接用 CUDA virtual memory API (cuMemAddressReserve + cuMemMap) 让 KV cache 逻辑连续但物理分页，不再需要 block table indirection，kernel 写得跟原生 attention 一样
• Pinned memory KV cache（Mooncake）：长 prompt 历史 KV 卸载到 CPU，需要时 NVLink-C2C / PCIe 重传
• KV cache 共享（LMCache, 2024）：跨进程 / 跨节点 KV 共享池，用 RDMA 传递

5.9.4 NVLink-C2C 与 unified memory 复活

过去 unified memory 的 page migration 往往不适合延迟敏感路径。Grace Hopper 与 Grace Blackwell 的 NVLink-C2C 改变了 CPU/GPU 共享内存的访问边界，但是否适合仍需按驻留位置、访问模式和目标系统实测：

// On a coherent Grace Hopper system, test managed placement explicitly.
float* w; cudaMallocManaged(&w, 600 * (1u << 30));   // 600 GB 模型权重
// w 落在 Grace LPDDR5X (480 GB), 由 NVLink-C2C 让 GPU 访问
// 比起 fp4 量化, 这种"显存外溢"对精度敏感模型更友好

典型用例是超大模型的混合内存推理：热权重和当前层激活尽量驻留 GPU HBM，冷权重、长尾 KV 或可重取数据放在 Grace 内存侧，按需通过 C2C 访问。是否划算取决于访问频率和重算成本。

5.9.5 显存优化的 2026 工业组合

一个完整 LLM 推理服务的显存优化技术栈：

1. fp8 / fp4 量化权重     → 权重显存减半到 1/4
2. fp8 KV cache           → KV 显存减半
3. PagedAttention         → 降低按 T_max 预分配造成的内部碎片
4. RadixAttention 前缀共享  → 多请求公共部分一份 KV
5. MLA 或 GQA             → KV head 维度压缩 (模型架构层面)
6. Stream-K + Stream-P    → batch 间显存峰值降低
7. cudaMallocAsync 内存池  → 减少碎片
8. NVLink-C2C 外溢 (GH200) → 大模型放得下

效果: 能否单卡服务 70B 级模型取决于量化格式、KV 长度、batch、并发调度和权重常驻策略

5.9.6 CUDA 12.6+ 内存子系统改进

• cudaMallocAsync 默认：从 12.6 起 PyTorch 默认走 stream-ordered allocator，碎片化大幅缓解
• cuMemMap-based KV cache：vAttention 风格的虚拟内存 KV 进 vLLM/TRT-LLM
• HMM (Heterogeneous Memory Management)：Linux 5.x+ 内核支持，让 host-device 内存接近统一
• Tensor Memory Accelerator (TMA) 在 Blackwell 上更强：单指令 multicast 到 cluster 内多个 block 的 shared

常见坑

• shared mem 数组开太大 → kernel launch 失败 too many resources requested
• 少写 __syncthreads → kernel 输出"几乎对"但偶尔错，race condition
• __constant__ 数组太大 → 编译报错 maximum constant size exceeded
• kernel 内 printf + 大量数据 → 缓冲区溢出, 输出截断

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

L1/Shared Carveout、DSMEM、Hints API、Pinned Memory、内存池

跨代 GPU 内存层级数字一览

选 kernel 优化方向之前先看你的硬件落在哪一格：

典型延迟（周期数，跨代基本不变）：

L1 / Shared 切片配置 — 解锁 48 KB 以上 shared

从 Volta 起，L1 和 shared memory 共用同一块物理 SRAM（"unified data cache"）。CUDA 让你 per-kernel 调整两者比例：

API 调用模式

__global__ void my_big_shared_kernel(float* x) {
        extern __shared__ float buf[];   // 动态 shared, 大小在 launch 时给
        // ...
    }

    int main() {
        // 1) 申请超过默认的 shared 大小 (例: A100 上想用 96 KB)
        int maxbytes = 96 * 1024;
        cudaFuncSetAttribute(my_big_shared_kernel,
                             cudaFuncAttributeMaxDynamicSharedMemorySize,
                             maxbytes);

        // 2) 偏好 shared 多一点 / L1 多一点 (driver 把它当 hint)
        cudaFuncSetAttribute(my_big_shared_kernel,
                             cudaFuncAttributePreferredSharedMemoryCarveout,
                             cudaSharedmemCarveoutMaxShared);   // 或一个 0-100 整数百分比

        // 3) launch 时第三个参数 = 动态 shared 字节数
        my_big_shared_kernel<<<grid, block, maxbytes>>>(x);
    }
    

三个 carveout 常量：

• cudaSharedmemCarveoutDefault — driver 自己挑
• cudaSharedmemCarveoutMaxL1 — 给 L1 越多越好
• cudaSharedmemCarveoutMaxShared — 给 shared 越多越好

FlashAttention 在 A100 上的典型配置：用 MaxDynamicSharedMemorySize = 100 KB + CarveoutMaxShared，把整个 KQ block + V tile 都放 shared，几乎不依赖 L1。Hopper 上进一步用 228 KB 全 shared 跑更大的 head_dim。

Constant cache 内部行为与 texture 简介

Constant cache 真正的"广播"

Constant memory 总容量 64 KB（全 GPU 共享），每 SM 上有 ~8 KB 的 constant cache。它的硬件特性：

所以 constant memory 只在"全 warp 读相同标量"这一个场景下有优势。RoPE 系数、attention scale 因子、batchnorm 的 affine 系数都是好用例；如果 LUT 大到要按 lane 取不同元素，立刻退化到比 global 慢。

Texture / Surface（仅简介）

texture memory 是为图形渲染设计的，对 CUDA 程序员意味着：

• Spatial locality cache：2D / 3D 块化缓存，对"邻近像素一起读"友好
• 硬件线性插值：tex2D<float>(t, 1.5f, 2.5f) 直接返回 4 个像素的双线性插值
• 边界处理：超出边界自动 clamp / wrap / mirror，无需写 if
• 归一化坐标：用 0-1 而不是 0-N，渲染常用

LLM 推理几乎用不到 texture——transformer 全是 1D 向量和 2D 矩阵的"按规则读"，spatial locality 帮不上。但你看 stable diffusion 的 VAE decoder、ControlNet 的 image warping 仍然有 cudaTextureObject_t，知道存在即可。

// 用 cuda runtime 创建 texture object (现代 API, CUDA 5.0+)
    cudaResourceDesc resDesc = {};
    resDesc.resType = cudaResourceTypeLinear;
    resDesc.res.linear.devPtr = d_data;
    resDesc.res.linear.desc = cudaCreateChannelDesc<float>();
    resDesc.res.linear.sizeInBytes = N * sizeof(float);

    cudaTextureDesc texDesc = {};
    texDesc.readMode = cudaReadModeElementType;

    cudaTextureObject_t tex;
    cudaCreateTextureObject(&tex, &resDesc, &texDesc, nullptr);

    // kernel 内: tex1Dfetch<float>(tex, i);
    

Distributed Shared Memory (DSMEM) — Hopper 新增内存空间

5.5.1.1 Thread Block Cluster 概念

Hopper (sm_90) 在 grid → block → warp → thread 之间插了一层：cluster。一个 cluster = 2~16 个 thread block，它们保证被同时调度到同一个 GPC（GPU Processing Cluster，含若干物理相邻的 SM）上。同一 cluster 内的所有 block 可以：

1. 互相访问对方的 shared memory（distributed shared memory, DSMEM）
2. 共享同一个硬件 barrier（cluster-scope barrier）
3. 通过 TMA multicast 一次性把 global memory 数据写进 cluster 内多个 block 的 shared

    graph TB
        subgraph GPC["GPC (GPU Processing Cluster, 含 8-16 个相邻 SM)"]
            subgraph C["Thread Block Cluster (sm_90+)"]
                B0["Block 0\nshared mem"]
                B1["Block 1\nshared mem"]
                B2["Block 2\nshared mem"]
                B3["Block 3\nshared mem"]
                B0 <--> B1
                B0 <--> B2
                B0 <--> B3
            end
        end
        style C fill:#f3f1e8,stroke:#8b1538
    

5.5.1.2 启动一个带 cluster 的 kernel

// kernel 声明: 用 __cluster_dims__ 编译期固定 cluster 形状
    __global__ void __cluster_dims__(2, 2, 1) cluster_kernel(/*...*/) {
        namespace cg = cooperative_groups;
        auto cluster = cg::this_cluster();
        auto block   = cg::this_thread_block();

        extern __shared__ float smem[];      // 本 block 的 shared
        smem[threadIdx.x] = some_value;
        cluster.sync();                       // 等 cluster 内所有 block 都填好

        // 拿到 cluster 内 block 1 的 shared mem 指针
        float* peer_smem = cluster.map_shared_rank(smem, /*peer_block_rank=*/1);
        float v = peer_smem[threadIdx.x];     // 直接读对面 block 的 shared!
    }

    // host: launch 时也要带 cluster 配置 (runtime API)
    cudaLaunchConfig_t config = {0};
    config.gridDim = dim3(8, 1, 1);
    config.blockDim = dim3(128, 1, 1);
    config.dynamicSmemBytes = 16384;

    cudaLaunchAttribute attrs[1];
    attrs[0].id = cudaLaunchAttributeClusterDimension;
    attrs[0].val.clusterDim = {2, 2, 1};      // 4-block cluster
    config.attrs = attrs;
    config.numAttrs = 1;

    cudaLaunchKernelEx(&config, cluster_kernel, /*args*/);
    

5.5.1.3 用什么场景？

• FlashAttention v3：在 cluster 内多个 block 间共享 K / V tile，跨 block 复用可减少重复 HBM 读取；实际比例用 NCU memory workload 统计
• GEMM split-K：cluster 内每个 block 算 K 维度的一部分，最后通过 DSMEM accumulate
• TMA multicast：sm_90a 上 1 条 TMA 指令把 global tile 同时写到 cluster 内所有 block 的 shared，省 N-1 次 HBM 读

5.5.1.4 查询当前 kernel 可支持的最大 cluster 大小

// 不同 SM 占用情况下能装下的最大 cluster
    int max_cluster_size;
    cudaOccupancyMaxPotentialClusterSize(
        &max_cluster_size, cluster_kernel, &config);

    // 同时能并发跑多少个该尺寸的 cluster
    int max_active_clusters;
    cudaOccupancyMaxActiveClusters(
        &max_active_clusters, cluster_kernel, &config);
    

cudaMemAdvise — Unified memory 性能调优三件套

5.8.3 的 prefetch 只是"把数据搬过去"。真正控制 unified memory 行为的是 cudaMemAdvise 提供的三种持久 hints，配合 prefetch 一起用：

5.8.7.1 三种 hints 的实际效果

// 场景: 共享只读权重, CPU 偶尔更新, GPU 0/1 都要高速读
    float* w;
    size_t bytes = 4ULL * (1u << 30);   // 4 GB
    cudaMallocManaged(&w, bytes);
    init_on_host(w, bytes);

    // (1) 标记 "ReadMostly" → driver 会在每个 device 上各放一份只读副本
    cudaMemLocation gpu0 = {.type = cudaMemLocationTypeDevice, .id = 0};
    cudaMemAdvise(w, bytes, cudaMemAdviseSetReadMostly, gpu0);

    // (2) preferred location = GPU 0
    cudaMemAdvise(w, bytes, cudaMemAdviseSetPreferredLocation, gpu0);

    // (3) 告诉系统 GPU 1 也会读 → 预建映射
    cudaMemLocation gpu1 = {.type = cudaMemLocationTypeDevice, .id = 1};
    cudaMemAdvise(w, bytes, cudaMemAdviseSetAccessedBy, gpu1);

    // (4) 真正物理 prefetch 到 GPU 0
    cudaMemPrefetchAsync(w, bytes, gpu0, /*flags=*/0, stream);

    // 现在两张 GPU 都能高速读 w, 无 page fault
    my_kernel<<<..., stream>>>(w);
    

5.8.7.2 重要陷阱：避免 host 频繁写 GPU-resident 内存

PDF §4.1.1.2.8 强调一个反直觉：

如果 unified memory 物理上在 GPU 上，CPU 写它时会触发整个 cache line 先从 GPU 读到 CPU cache，CPU 写完再传回 GPU。频繁小量写效率极低。

正确做法：把要 host 频繁写的小变量 preferred location 设为 host，让 GPU 跨 NVLink-C2C 去读：

int* control_flag;       // 主机准备频繁更新的小标志
    cudaMallocManaged(&control_flag, sizeof(int));

    cudaMemLocation host = {.type = cudaMemLocationTypeHost};
    cudaMemAdvise(control_flag, sizeof(int),
                  cudaMemAdviseSetPreferredLocation, host);
    cudaMemAdvise(control_flag, sizeof(int),
                  cudaMemAdviseSetAccessedBy, host);

    for (int i = 0; i < N; ++i) {
        *control_flag = i;                                  // 直接写 CPU 内存, 快
        poll_kernel<<<...>>>(control_flag);                 // GPU 跨 C2C 读
        cudaDeviceSynchronize();
    }
    

5.8.7.3 查询当前内存范围的属性

// 查这块内存现在 preferred location 是谁
    cudaMemLocationType locType;
    cudaMemRangeGetAttribute(&locType, sizeof(locType),
                              cudaMemRangeAttributePreferredLocationType,
                              w, bytes);

    // 查上一次 prefetch 到的位置
    int last_loc;
    cudaMemRangeGetAttribute(&last_loc, sizeof(last_loc),
                              cudaMemRangeAttributeLastPrefetchLocationId,
                              w, bytes);
    

Pinned / mapped memory 完整 API

5.5 节给了三种分配的速查表。这里展开 host 端 page-locked memory 的细节——这是 H2D / D2H 异步 overlap、CPU-GPU 双工管道的基石。

5.8.8.1 三种创建 pinned memory 的 API

5.8.8.2 cudaHostAlloc 的 flags

void* p;

    // (1) 默认: 仅 host 可访问, 适合作 H2D async copy 的源
    cudaHostAlloc(&p, n, cudaHostAllocDefault);

    // (2) portable: 多 GPU context 都能看到 (默认只 alloc 时 current context 能看)
    cudaHostAlloc(&p, n, cudaHostAllocPortable);

    // (3) mapped: GPU 可以直接通过 PCIe 访问（zero-copy）
    //    适合: 单次小数据传输, 或 GPU 偶尔 poll host 标志
    //    Avoid on a bandwidth-critical path; accesses traverse the host link.
    cudaHostAlloc(&p, n, cudaHostAllocMapped);
    void* d_p;
    cudaHostGetDevicePointer(&d_p, p, 0);  // 拿到对应的设备端指针

    // (4) write-combined: 给 CPU 写、GPU 读的 buffer
    //    CPU 写性能高 (不走 cache, 直接 burst)
    //    但 CPU 读会奇慢 (cache miss + 走总线)
    cudaHostAlloc(&p, n, cudaHostAllocWriteCombined);

    // 组合: 异步源 + 多 device 可见 + write-combined
    cudaHostAlloc(&p, n,
        cudaHostAllocPortable | cudaHostAllocWriteCombined);
    

5.8.8.3 为什么 pinned 才能异步？

普通 malloc 出来的内存可能被 OS 换出到磁盘 (swap)。CUDA 做 H2D 的时候必须保证物理地址稳定，所以：

• 非 pinned：driver 先把数据 staging 到内部 pinned buffer（一次额外 memcpy），然后才 DMA 到 device。cudaMemcpyAsync 在非 pinned 上行为退化为同步
• pinned：driver 直接发 DMA，CPU 不参与，可以与 kernel 并行

throughput 差别请在目标机器上测，尤其要记录 PCIe/NVLink 拓扑、NUMA 绑定、传输大小和是否与 kernel overlap：

5.8.8.4 pinned memory 的代价

5.8.8.5 LLM 推理的典型 pinned 用法

// vLLM / TRT-LLM 风格: 启动时一次性分配 pinned staging buffer
    size_t staging_size = 256 * 1024 * 1024;   // 256 MB
    void* h_staging;
    cudaHostAlloc(&h_staging, staging_size, cudaHostAllocDefault);

    // 请求来时, 用户 prompt 先拷到 staging, 再异步推到 GPU
    memcpy(h_staging, user_prompt_tokens, n * sizeof(int));
    cudaMemcpyAsync(d_input, h_staging, n * sizeof(int),
                    cudaMemcpyHostToDevice, stream);
    prefill_kernel<<<..., stream>>>(d_input);

    // 同时另一个 stream 把上一轮的 output token 拷回 host
    cudaMemcpyAsync(h_staging_out, d_output, m * sizeof(int),
                    cudaMemcpyDeviceToHost, stream_out);
    

stream-ordered memory pool 内部机制

5.8.1 给了最朴素的 cudaMallocAsync。其底层是 memory pool，是真正决定显存碎片率的对象。深入掌握三件事：

1. 每个 device 有一个 default memory pool，可以自己 create 显式 pool
2. Pool 有 release threshold：低于此阈值的空闲内存不归还 OS，下次 malloc 直接复用
3. Pool 有 reuse policy：决定 stream A free 的内存 stream B 能否立即用

5.8.9.1 设置 release threshold 减少 OS 调用

cudaMemPool_t pool;
    cudaDeviceGetDefaultMemPool(&pool, /*device=*/0);

    // 池中保留 20 GB 不归还 OS——下次 malloc 几乎 0 开销
    uint64_t threshold = 20ULL * (1u << 30);
    cudaMemPoolSetAttribute(pool,
        cudaMemPoolAttrReleaseThreshold, &threshold);

    // 实际用 cudaMallocAsync 时自动走 pool
    void* ptr;
    cudaMallocAsync(&ptr, bytes, stream);    // 复用池, 不调 cuMemMap
    cudaFreeAsync(ptr, stream);               // 还给池
    

5.8.9.2 三种 reuse policy

int enable = 1;
    cudaMemPoolSetAttribute(pool,
        cudaMemPoolReuseAllowOpportunistic, &enable);
    

5.8.9.3 跨进程共享 pool（IPC）

vLLM 多进程架构里，每个 worker 一个 process，但共享同一显存池可省 N-1 份 reserve：

// 进程 A: 创建 IPC-enabled pool
    cudaMemPoolProps props = {};
    props.allocType = cudaMemAllocationTypePinned;
    props.location.type = cudaMemLocationTypeDevice;
    props.location.id = 0;
    props.handleTypes = cudaMemHandleTypePosixFileDescriptor;

    cudaMemPool_t pool;
    cudaMemPoolCreate(&pool, &props);

    int fd;
    cudaMemPoolExportToShareableHandle(&fd, pool,
        cudaMemHandleTypePosixFileDescriptor, 0);
    // 通过 UNIX domain socket 把 fd 发给进程 B...

    // 进程 B: 导入
    cudaMemPool_t imp_pool;
    cudaMemPoolImportFromShareableHandle(&imp_pool, (void*)&fd,
        cudaMemHandleTypePosixFileDescriptor, 0);
    

5.8.9.4 监控 pool 占用

uint64_t reserved, used, reserved_high, used_high;
    cudaMemPoolGetAttribute(pool, cudaMemPoolAttrReservedMemCurrent, &reserved);
    cudaMemPoolGetAttribute(pool, cudaMemPoolAttrUsedMemCurrent,     &used);
    cudaMemPoolGetAttribute(pool, cudaMemPoolAttrReservedMemHigh,    &reserved_high);
    cudaMemPoolGetAttribute(pool, cudaMemPoolAttrUsedMemHigh,        &used_high);

    printf("pool reserved=%.1f GB, in-use=%.1f GB, peak=%.1f GB\n",
           reserved / 1e9, used / 1e9, reserved_high / 1e9);

    // reset 高水位计
    uint64_t zero = 0;
    cudaMemPoolSetAttribute(pool, cudaMemPoolAttrReservedMemHigh, &zero);
    

下一章导览

第 6 章把 shared memory 组织成 tile，用数据复用把朴素矩阵乘推向高吞吐实现。