4、统一内存与流：流关联、异步预取、流同步

好，咱们接着聊统一内存。前面几章我们把统一内存的基本概念、数据迁移、页错误这些机制都捋了一遍。这一章，我打算聊聊统一内存和 CUDA 流（Stream）怎么配合。

说实话，我刚接触统一内存时，觉得它就是个“自动搬运工”。后来发现，这个搬运工如果不管它，它可能会搬得不是时候、不是地方。这时候，流就派上用场了。

4.1 流关联：让数据跟着流走

先问个问题：你在 GPU 上启动了一个 kernel，它访问统一内存里的数据。CUDA 运行时怎么知道该把数据迁移到哪个设备？

答案是：它看的是当前默认流（或者你指定的流）。

默认情况下，所有统一内存的页错误处理、数据迁移，都发生在当前设备上。但如果你有多个流，每个流跑不同的任务，数据就可能被来回折腾。

流关联（Stream Association） 就是干这个的。它允许你把一段统一内存和一个特定的流绑定起来。绑了之后，这个流发起的访问，会优先触发数据迁移到该流所在的设备。

核心概念：流关联让 CUDA 运行时知道，“这个流要用这块数据，赶紧搬过来”。

我记得在做一个多流并行的项目时，两个流分别处理不同的数据块。一开始没做流关联，结果两个流互相抢数据，性能惨不忍睹。后来加上流关联，每个流只搬自己需要的数据，效率直接翻倍。

4.2 异步预取：提前把数据搬过来

统一内存的自动迁移，说白了是按需的。你访问了，它才搬。这有个问题：如果 kernel 启动后立刻访问大量数据，页错误会扎堆出现，性能瞬间崩掉。

怎么办？提前搬。

异步预取（Async Prefetching） 就是干这个的。它让你在 kernel 启动之前，主动告诉 CUDA：“把这块数据搬到 GPU 上，我马上要用。”

API 长这样：

cudaError_t cudaMemPrefetchAsync(
    const void *devPtr,
    size_t count,
    int dstDevice,
    cudaStream_t stream = 0
);

参数很简单：

  • devPtr：统一内存的起始地址
  • count：要预取的字节数
  • dstDevice：目标设备 ID（CPU 用 cudaCpuDeviceId）
  • stream：关联的流（可选）

举个例子：

float *d_data;
cudaMallocManaged(&d_data, N * sizeof(float));

// 把数据预取到 GPU 0
cudaMemPrefetchAsync(d_data, N * sizeof(float), 0, stream);

// 启动 kernel，此时数据已经在 GPU 上了
myKernel<<<grid, block, 0, stream>>>(d_data);

我个人习惯在 kernel 启动前，先做一次预取。尤其是数据量大的时候，这招特别管用。我曾经在一个图像处理项目里，预取把 kernel 启动延迟从 200ms 降到了 5ms。嗯，你没看错，就是这么大差距。

小技巧：预取是异步的，它不会阻塞 CPU。你可以一边预取，一边做其他准备工作。但要注意，预取操作本身也会占用 PCIe 带宽，别一次预取太多，否则其他传输会被堵住。

4.3 流同步：别让数据跑丢了

预取是异步的，那问题来了：你怎么知道数据已经搬完了？

答案是：用流同步。

CUDA 提供了几种同步方式：

  • cudaStreamSynchronize(stream)：等待流里的所有操作完成
  • cudaDeviceSynchronize()：等待设备上所有流完成
  • cudaEventSynchronize(event)：等待某个事件

但这里有个坑，我踩过好几次。

注意：统一内存的预取操作，虽然是在流里发起的，但它的完成时间点，并不一定和流里的其他操作完全对齐。说白了，预取可能在流同步之前就完成了，也可能在之后。这取决于硬件和驱动。

所以，我建议的做法是：

  1. 发起预取
  2. 在同一个流里启动 kernel
  3. 用 cudaStreamSynchronize 等待流完成

这样最保险。你想想看，预取和 kernel 在同一个流里，CUDA 会保证它们的顺序。预取没完成，kernel 不会启动。

但如果你在流 A 里预取，在流 B 里启动 kernel，那就得小心了。两个流之间没有依赖关系，数据可能还没到，kernel 就开始跑了。这时候，你需要用事件或者 cudaStreamWaitEvent 来手动同步。

我曾经在一个项目里，就是因为没注意流之间的依赖，导致 kernel 读到了旧数据。排查了半天，最后发现是预取和 kernel 不在同一个流里。从那以后，我养成了一个习惯：预取和 kernel 尽量放同一个流。

4.4 实战建议：流与统一内存的最佳实践

说了这么多，我总结几条实战经验：

场景 建议做法 原因
单流、单 GPU 预取 + kernel 放同一个流 简单、可靠、无需额外同步
多流、单 GPU 每个流预取自己的数据块 避免数据竞争和页错误风暴
多 GPU 每个 GPU 预取到本地显存 减少跨 GPU 访问延迟
CPU-GPU 混合 CPU 访问前预取到 CPU 端 避免 CPU 触发页错误

另外，别忘了 cudaMemAdvise 这个函数。它可以告诉 CUDA 运行时，某块数据在某个设备上的访问模式。比如：

// 告诉 CUDA，这块数据主要在 GPU 0 上被读取
cudaMemAdvise(d_data, N * sizeof(float), cudaMemAdviseSetReadMostly, 0);

这个建议配合预取使用，效果更好。CUDA 会根据你的建议，优化数据迁移策略。

一句话总结：统一内存 + 流 = 自动迁移 + 手动优化。预取让你掌控时机，流让你掌控顺序。两者结合，才能发挥统一内存的真正威力。

好，这一章就到这里。下一章我们聊聊统一内存的调试和性能分析工具。到时候我会分享一些我踩过的坑，以及怎么用工具快速定位问题。

← 上一章 📖 返回目录 下一章 →