4. 页锁定内存与零拷贝技术

在上一节中我们了解到，普通 malloc 或 new 分配的主机内存是“可分页”的：操作系统可以在物理内存与磁盘交换区之间移动这些页面。这带来一个关键问题——当 GPU 通过 PCIe 总线以 DMA 方式读取主机内存时，如果目标页面恰好被换出，传输就会失败或被迫降级。页锁定内存（Pinned Memory，又称锁页内存）正是为解决这一问题而设计的。

页锁定内存的分配与注册

页锁定内存的本质是向操作系统申请一块不会被换出、物理地址固定的主机内存。CUDA 提供两种获得方式：直接分配和注册已有内存。

直接分配使用 cudaMallocHost（或较新的 cudaHostAlloc），它一次性完成“分配 + 锁定”两步：

float *h_pinned;
size_t bytes = N * sizeof(float);
cudaMallocHost((void**)&h_pinned, bytes);   // 分配并锁定
// ... 使用 h_pinned ...
cudaFreeHost(h_pinned);                       // 释放并解锁

如果内存已经由普通方式分配好，则可用 cudaHostRegister 将其“就地注册”为页锁定内存，避免数据搬迁：

float *h_pageable = (float*)malloc(bytes);
cudaHostRegister(h_pageable, bytes, cudaHostRegisterDefault);
// ... 使用 ...
cudaHostUnregister(h_pageable);   // 取消注册
free(h_pageable);

注册方式适合已有缓冲区（如第三方库分配的内存、内存映射文件）的场景，但需注意注册的地址范围必须页对齐，且注册本身有开销，不适合频繁注册/注销的小块内存。页锁定内存的代价是：它占用的物理内存无法被换出，过量分配会拖慢甚至拖垮整个系统，因此应只对真正需要高频传输的缓冲区使用。

零拷贝内存映射

传统路径中，数据要经历“主机内存 → 设备显存 → 计算”的搬运。零拷贝（Zero-Copy）技术则利用页锁定内存的物理地址固定特性，让 GPU 通过 PCIe 总线直接访问主机内存，省去显式拷贝这一步。其实现依赖“映射”操作：

float *h_zero;
cudaHostAlloc((void**)&h_zero, bytes, cudaHostAllocMapped);

float *d_zero;
cudaHostGetDevicePointer((void**)&d_zero, h_zero, 0);
// 内核中直接使用 d_zero 指针访问主机内存
kernel<<<grid, block>>>(d_zero, N);

关键点是 cudaHostAllocMapped 标志与 cudaHostGetDevicePointer：前者声明该内存可被设备映射，后者取得设备端可用的指针。此后内核通过 d_zero 读写的就是主机内存，数据不再经过显存中转。在支持统一寻址（UVA）的 64 位平台上，主机指针与设备指针一致，可省略 cudaHostGetDevicePointer。

适用场景与限制

零拷贝最适合以下情形：数据量不大但传输频繁、需要与计算重叠；数据规模超过显存容量，只能“边算边取”；以及集成显卡（如 Jetson 系列）上主机与设备共享物理内存，零拷贝几乎无额外代价。反之，在独立显卡上对大规模数据做密集随机访问时，PCIe 带宽（PCIe 3.0 x16 约 12–16 GB/s，远低于显存的数百 GB/s）会成为瓶颈，此时传统“拷贝到显存再计算”反而更快。

下图对比了三条数据通路的差异，可直观看出零拷贝省去的是哪一段搬运：

实际传输性能测试与对比

要判断是否值得使用页锁定内存，最可靠的方式是实测。下面给出一个可复现的对比测试框架，分别测量可分页内存与页锁定内存的 H2D（主机到设备）拷贝带宽。示例参数为：数据量 256 MB，重复 20 次取平均，使用 cudaEvent 计时。

size_t bytes = 256UL << 20;   // 256 MB
float *h_page, *h_pin, *d_buf;
cudaMallocHost(&h_pin, bytes);
h_page = (float*)malloc(bytes);
cudaMalloc(&d_buf, bytes);

cudaEvent_t t0, t1;
cudaEventCreate(&t0); cudaEventCreate(&t1);

// 测可分页内存
cudaEventRecord(t0);
for (int i = 0; i < 20; i++)
    cudaMemcpy(d_buf, h_page, bytes, cudaMemcpyHostToDevice);
cudaEventRecord(t1);
cudaEventSynchronize(t1);
float ms_page; cudaEventElapsedTime(&ms_page, t0, t1);

// 测页锁定内存（同样循环）
// ... 用 h_pin 替换 h_page ...
float ms_pin; // 由同样方式测得

float bw_page = bytes * 20 / (ms_page * 1e6);  // GB/s
float bw_pin  = bytes * 20 / (ms_pin  * 1e6);

在典型 PCIe 3.0 x16 平台上，合理假设的示例结果大致为：可分页内存 H2D 带宽约 3–6 GB/s，页锁定内存可达 10–12 GB/s，提升约 2–3 倍。这一差距来自可分页内存需要驱动先做一次“暂存到内部锁页缓冲区”的额外拷贝，而页锁定内存可由 DMA 引擎直接搬运。若进一步使用 cudaMemcpyAsync 配合多流，页锁定内存还能实现拷贝与计算重叠，把有效吞吐再推高。

综合来看，选择策略可以归纳为：高频、小块、需重叠的传输优先用页锁定内存；数据超显存、需流式处理时考虑零拷贝映射；而一次性、大块、密集访问的数据，老老实实拷贝到显存往往仍是带宽最优解。理解这三条通路的代价差异，才能在真实工程中做出有依据的取舍。