📚 课程目录第4章统一内存与零拷贝工程权衡

4. 统一内存与零拷贝工程权衡

统一内存（Unified Memory, UM）把 CPU 与 GPU 的地址空间抽象成同一份指针，让开发者不必再手工维护两份数据副本。但“统一”只是编程模型层面的统一，物理上显存与主机内存依然分离，数据仍要在两者之间流动。理解这一点，是判断何时该用 cudaMallocManaged、何时该退回显式拷贝的前提。

按需页迁移：统一内存的真实代价

用 cudaMallocManaged 分配的托管内存，初始并不驻留在任何物理位置。当 CPU 或 GPU 首次访问某个页时，会触发一次页错误（page fault），驱动把该页迁移到访问方所在的物理内存，并建立映射。这个过程称为按需页迁移（on-demand page migration）。

关键点在于：迁移以页为粒度，且是同步阻塞的。如果 GPU 内核在循环中反复访问一块尚未驻留显存的托管数据，第一次访问会触发大量页错误，驱动逐页搬运，内核被迫等待。更糟的是，若 CPU 与 GPU 交替访问同一块数据，页面会在两端来回“弹跳”（thrashing），每次弹跳都伴随一次迁移开销。因此，托管内存的性能瓶颈往往不在带宽，而在页错误的延迟与迁移次数。

下面这张图展示了同一块托管内存在 CPU 与 GPU 之间按需迁移的路径，以及 cudaMemAdvise 与 cudaMemPrefetchAsync 分别作用在哪个环节。

两种调优手段的分工

cudaMemAdvise 是“声明意图”，它不搬运数据，只告诉驱动这块内存的访问模式，让驱动在后续迁移决策中少犯错。常用的三个 advice：

cudaMemAdviseSetReadMostly：声明只读。驱动可在 CPU 与 GPU 各保留一份副本，读操作直接命中本地副本，避免迁移。
cudaMemAdviseSetPreferredLocation：指定首选驻留设备。例如把一块主要给 GPU 用的数据首选位置设为 GPU，减少回迁。
cudaMemAdviseSetAccessedBy：声明某设备会访问该内存，驱动可提前建立映射，降低页错误成本。

cudaMemPrefetchAsync 则是“主动搬运”，在访问发生之前把页迁到目标设备。它把原本分散在首次访问时的页错误，集中成一次可预期的异步迁移，从而与计算重叠。典型用法是在内核启动前预取到 GPU，在内核结束后预取回 CPU：

int dev = 0;
cudaMemAdvise(data, N*sizeof(float),
              cudaMemAdviseSetReadMostly, dev);
cudaMemPrefetchAsync(data, N*sizeof(float), dev, stream);
kernel<<<grid, block, 0, stream>>>(data, N);
cudaMemPrefetchAsync(data, N*sizeof(float),
                     cudaCpuDeviceId, stream);
cudaStreamSynchronize(stream);

注意 cudaCpuDeviceId 是预取回主机端的特殊设备号。若省略回迁，CPU 后续访问会再次触发页错误，把预取省下的时间又还回去。

零拷贝映射内存的适用边界

零拷贝（zero-copy）走的是另一条路：用 cudaHostAlloc 配合 cudaHostAllocMapped 分配锁页主机内存，再通过 cudaHostGetDevicePointer 拿到设备侧指针，内核直接经 PCIe/NVLink 访问主机内存，不做显式拷贝。

它的代价是：每次访问都要跨总线，延迟远高于显存，且带宽受限于互连。因此零拷贝只适合访问次数少、数据量小、或访问具有强局部性的场景，例如内核只读取少量控制参数、稀疏地取几个元素。若内核要反复遍历大数组，零拷贝会把 PCIe 带宽变成瓶颈，性能通常远低于“先拷贝到显存再算”。

与显式拷贝的性能对比

三种方式的核心差异可以归纳为下表（数值为量级示意，非实测）：

方式	数据搬运时机	典型瓶颈	适用场景
显式拷贝 cudaMemcpy	开发者手动、成块	拷贝本身带宽	数据量大、访问模式清晰
托管内存 + 预取	预取时集中搬运	页错误与迁移次数	数据结构复杂、指针多
零拷贝映射	访问时逐次跨总线	互连延迟与带宽	小数据、稀疏访问

判断顺序建议：先问数据是否会被 GPU 反复遍历。是，则优先显式拷贝或托管内存加预取；否，且访问稀疏，才考虑零拷贝。托管内存不是“免费的统一”，预取与 advice 才是让它跑得快的关键。

常见误区：只调用 cudaMallocManaged 而不做任何预取或 advice，就期望它接近显式拷贝的性能。在数据量大、CPU/GPU 交替访问的场景下，页弹跳会让性能显著劣化，此时显式拷贝反而更稳。

工程上更稳妥的做法是：对生命周期长、访问模式稳定的数据，用显式拷贝把控制权握在自己手里；对指针结构复杂、难以手工管理的对象图，用托管内存并配合 cudaMemAdvise 声明意图、cudaMemPrefetchAsync 提前搬运；只有在数据量小且访问稀疏时，才启用零拷贝映射。三者不是替代关系，而是按数据特征分工的工具。