3、统一内存编程基础：cudaMallocManaged、cudaMemPrefetchAsync、cudaMemAdvise
好，我们直接进入正题。统一内存（Unified Memory）这个概念，说白了就是让 CPU 和 GPU 共享同一份地址空间。你不需要再手动来回搬运数据，也不用操心指针是设备端还是主机端。听起来很美好对吧？但实际用起来，有几个关键 API 你必须吃透。
我个人习惯把统一内存编程拆成三个层次：分配、迁移、优化。正好对应三个核心函数：cudaMallocManaged、cudaMemPrefetchAsync、cudaMemAdvise。咱们一个一个来。
3.1 cudaMallocManaged：一键分配，双端可见
这个函数是统一内存的入口。你用它分配一块内存，CPU 和 GPU 都能直接访问。不需要 cudaMalloc 再 cudaMemcpy，省事不少。
float *data;
cudaMallocManaged(&data, N * sizeof(float));
// CPU 可以直接读写
for (int i = 0; i < N; i++) data[i] = i * 1.0f;
// GPU 也能直接读写
kernel<<<grid, block>>>(data, N);
cudaDeviceSynchronize();嗯，这里要注意：cudaMallocManaged 分配的内存，默认是“惰性迁移”的。什么意思？就是数据一开始可能放在 CPU 上，等你 GPU 要用了，它才悄悄搬过去。这个机制叫 page fault 驱动迁移。
所以，cudaMallocManaged 适合写原型，或者数据访问模式比较简单的场景。但如果你追求极致性能，就得手动干预迁移策略。
3.2 cudaMemPrefetchAsync：提前搬家，避免缺页
这个函数就是用来解决上面那个问题的。你可以提前告诉 CUDA 运行时：“这块数据马上要在 GPU 上用，你先把它搬过去。” 这样 GPU 执行时就不会触发 page fault，延迟大幅降低。
cudaMemPrefetchAsync(data, N * sizeof(float), deviceId, stream);参数很简单：数据指针、大小、目标设备 ID、流。注意最后一个参数是流，你可以把它放到异步操作里，不阻塞主线程。
我个人习惯在 kernel launch 之前，先 prefetch 一下。尤其是处理大数组时，效果立竿见影。
cudaCpuDeviceId 把数据 prefetch 回 CPU。比如你 CPU 要读结果，提前 prefetch 回来，避免读的时候再触发缺页。
你想想看，prefetch 的本质就是“预判”。你预判得越准，性能越好。我曾经在一个图像处理项目里，把 prefetch 和 stream 结合使用，吞吐量提升了将近 30%。
3.3 cudaMemAdvise：给运行时提建议
这个函数更高级一点。它不是直接搬数据，而是告诉 CUDA 运行时：“我接下来会怎么用这块内存，你帮我优化一下。” 说白了，就是给运行时提建议。
cudaMemAdvise(data, N * sizeof(float), cudaMemAdviseSetPreferredLocation, deviceId);常用的建议类型有几种：
| 建议类型 | 含义 | 适用场景 |
|---|---|---|
cudaMemAdviseSetPreferredLocation |
设置首选位置 | 数据主要被某个设备访问 |
cudaMemAdviseSetAccessedBy |
标记会被哪个设备访问 | 多 GPU 共享数据 |
cudaMemAdviseSetReadMostly |
数据只读，可以复制副本 | 只读查找表、常量数据 |
举个例子。如果你有一张只读的查找表，所有 GPU 线程都要读它。你可以用 cudaMemAdviseSetReadMostly，这样每个 GPU 都会保留一份本地副本，不用跨 PCIe 去读。
cudaMemAdvise(lut, size, cudaMemAdviseSetReadMostly, deviceId);我记得有一次做分子动力学模拟，有一张很大的力场参数表，所有 GPU 都要频繁读取。用了 SetReadMostly 之后，带宽压力直接减半。
3.4 三者配合：一个完整示例
光说不练假把式。咱们看一个完整的例子，把三个 API 串起来。
float *data;
int N = 1 << 20;
int deviceId;
cudaGetDevice(&deviceId);
// 1. 分配统一内存
cudaMallocManaged(&data, N * sizeof(float));
// 2. 初始化数据（CPU 端）
for (int i = 0; i < N; i++) data[i] = i * 0.5f;
// 3. 建议：数据主要被 GPU 访问
cudaMemAdvise(data, N * sizeof(float),
cudaMemAdviseSetPreferredLocation, deviceId);
// 4. 预取到 GPU
cudaMemPrefetchAsync(data, N * sizeof(float), deviceId, 0);
// 5. 启动 kernel
kernel<<<grid, block>>>(data, N);
cudaDeviceSynchronize();
// 6. 结果 prefetch 回 CPU
cudaMemPrefetchAsync(data, N * sizeof(float), cudaCpuDeviceId, 0);
// 7. CPU 读取结果
float sum = 0;
for (int i = 0; i < N; i++) sum += data[i];
cudaFree(data);你看，整个流程非常清晰。没有手动 cudaMemcpy，但数据迁移的时机完全可控。
cudaDeviceSynchronize 或者用 stream 同步。我曾经因为忘了同步，在 CPU 端读到一半的数据还是旧的，排查了半天。
3.5 避坑指南与个人经验
最后，分享几个我实际项目中踩过的坑：
- 不要滥用 managed memory。 如果你的数据访问模式非常固定，手动
cudaMemcpy反而更快。managed 的开销在于 page fault 处理。 - 多 GPU 场景要小心。 每个 GPU 都有自己的内存，managed 虽然能自动迁移，但频繁跨 GPU 访问会触发大量缺页。建议用
cudaMemAdviseSetAccessedBy提前声明。 - prefetch 不是万能的。 如果你 prefetch 的数据量太大，超过了 GPU 显存，它会自动回退到缺页模式。所以还是要控制好数据规模。
- 调试时多用
cuda-memcheck。 managed 内存的越界访问很难排查，这个工具能帮你定位。
嗯，统一内存编程基础就讲到这里。这三个 API 是统一内存的基石，理解透了，后面的高级话题才能接得住。下一章我们会深入 cudaStream 和异步并发，到时候你会看到 prefetch 和 stream 配合起来有多强大。