3、CUDA统一内存编程：cudaMallocManaged用法、页面迁移机制、内存预取与驻留控制

3.1 为什么需要统一内存？

做GPU编程的人，十有八九都被内存管理折磨过。

传统CUDA编程里，你得手动管两套内存：主机端用malloc，设备端用cudaMalloc。数据要搬来搬去，还得记着cudaMemcpy。代码写起来像这样：

float *h_data, *d_data;
h_data = (float*)malloc(N * sizeof(float));
cudaMalloc(&d_data, N * sizeof(float));
// 初始化h_data...
cudaMemcpy(d_data, h_data, N * sizeof(float), cudaMemcpyHostToDevice);
// 跑kernel...
cudaMemcpy(h_data, d_data, N * sizeof(float), cudaMemcpyDeviceToHost);
free(h_data);
cudaMalloc(d_data);

烦不烦？我当年刚入行时，就因为这破事debug了整整三天——最后发现是cudaMemcpy的方向写反了。

统一内存（Unified Memory）就是来解决这个痛点的。它让你只用一套指针，CPU和GPU都能访问。系统自动帮你搬数据。

3.2 cudaMallocManaged：一键分配，双端访问

核心API就一个：cudaMallocManaged。

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

// CPU端直接用
for (int i = 0; i < N; i++) data[i] = i * 1.0f;

// GPU端也直接用
my_kernel<<<grid, block>>>(data, N);

cudaDeviceSynchronize();
// CPU再读结果
printf("data[0] = %f\n", data[0]);

cudaFree(data);

你看，没有malloc，没有cudaMemcpy，没有两套指针。清爽多了。

但注意，cudaMallocManaged返回的指针，默认是“驻留在CPU端”的。GPU第一次访问时，才会触发迁移。

3.3 页面迁移机制：数据是怎么“跑”过去的？

统一内存背后，是CUDA驱动在管一个“页面迁移引擎”。

说白了，它把物理内存分成4KB大小的页面。每个页面有个“主人”标签——当前在CPU上，还是在GPU上。

当GPU要访问一个页面，发现它不在显存里，就会触发一个缺页中断。驱动把页面从CPU内存搬到显存，然后更新页表。

这个过程是透明的。但代价呢？

看到没？一次缺页中断，比一次普通内存访问慢了100倍。如果代码写得不好，频繁触发缺页，性能会崩得很难看。

3.4 内存预取：提前把数据准备好

既然缺页这么贵，那能不能提前把数据搬过去？

当然可以。CUDA提供了cudaMemPrefetchAsync。

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

// 初始化数据（在CPU上）
for (int i = 0; i < N; i++) data[i] = i * 1.0f;

// 预取到GPU，指定设备ID为0
cudaMemPrefetchAsync(data, N * sizeof(float), 0, NULL);

// 现在GPU访问data，不会触发缺页
my_kernel<<<grid, block>>>(data, N);

cudaDeviceSynchronize();

// 结果要读回CPU？再预取回来
cudaMemPrefetchAsync(data, N * sizeof(float), cudaCpuDeviceId, NULL);
printf("data[0] = %f\n", data[0]);

预取的本质，就是告诉驱动：“我马上要用这些数据了，你先搬过去。”

它是个异步操作。调用后立即返回，后台在PCIe上慢慢搬。你可以在搬数据的同时做其他计算，实现计算与数据传输重叠。

3.5 内存驻留控制：把数据“钉”在显存里

预取只是“建议”。驱动可能会因为显存不够，把一些页面踢回CPU。

如果你确定某个数据要长期留在GPU上，可以用cudaMemAdvise设置驻留提示。

cudaMemAdvise(data, N * sizeof(float), cudaMemAdviseSetPreferredLocation, 0);
cudaMemAdvise(data, N * sizeof(float), cudaMemAdviseSetAccessedBy, 0);

第一个参数：告诉驱动，这个数据“最好”放在设备0上。

第二个参数：告诉驱动，设备0会频繁访问这个数据。

这两个提示组合使用，驱动会尽量把数据留在显存里，不会轻易换出。

还有一个更狠的：cudaMemAdviseSetReadMostly。如果数据是只读的（比如查找表），加上这个提示，驱动会在多个设备间复制一份，避免反复迁移。

// 只读数据，每个设备都保留一份副本
cudaMemAdvise(lookup_table, size, cudaMemAdviseSetReadMostly, 0);

3.6 实战：一个完整的预取+驻留示例

把上面这些串起来，写一个完整的例子：

#include <cuda_runtime.h>
#include <stdio.h>

__global__ void vector_add(float *a, float *b, float *c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) c[idx] = a[idx] + b[idx];
}

int main() {
    int n = 1 << 20;  // 约100万个元素
    size_t size = n * sizeof(float);

    float *a, *b, *c;
    cudaMallocManaged(&a, size);
    cudaMallocManaged(&b, size);
    cudaMallocManaged(&c, size);

    // 初始化（在CPU上）
    for (int i = 0; i < n; i++) {
        a[i] = i * 1.0f;
        b[i] = (n - i) * 1.0f;
    }

    // 设置驻留提示：a和b是只读的，c是输出
    cudaMemAdvise(a, size, cudaMemAdviseSetReadMostly, 0);
    cudaMemAdvise(b, size, cudaMemAdviseSetReadMostly, 0);
    cudaMemAdvise(c, size, cudaMemAdviseSetPreferredLocation, 0);

    // 预取到GPU
    cudaMemPrefetchAsync(a, size, 0, NULL);
    cudaMemPrefetchAsync(b, size, 0, NULL);
    cudaMemPrefetchAsync(c, size, 0, NULL);

    // 跑kernel
    int threads = 256;
    int blocks = (n + threads - 1) / threads;
    vector_add<<<blocks, threads>>>(a, b, c, n);

    cudaDeviceSynchronize();

    // 结果预取回CPU
    cudaMemPrefetchAsync(c, size, cudaCpuDeviceId, NULL);
    cudaDeviceSynchronize();

    // 验证
    for (int i = 0; i < 10; i++) {
        printf("c[%d] = %f (expected %f)\n", i, c[i], (float)n);
    }

    cudaFree(a);
    cudaFree(b);
    cudaFree(c);
    return 0;
}

这个例子，把预取和驻留都用上了。a和b是只读的，每个GPU都保留副本。c是输出，固定在设备0上。kernel跑的时候，几乎不会触发缺页。

3.7 避坑指南

统一内存虽好，但有几个坑要小心：

• 不要在小数据上滥用：如果数据只有几百字节，预取的开销比直接缺页还大。小数据就让驱动自己处理。
• 注意同步点：cudaMemPrefetchAsync是异步的。如果你紧接着就启动kernel，可能数据还没搬完。建议在预取和kernel之间加一个cudaStreamSynchronize，或者用事件同步。
• 多GPU场景要小心：每个GPU有自己的显存。如果数据在多个GPU间频繁切换，迁移开销会爆炸。尽量让数据固定在一个设备上。
• 不要和手动cudaMemcpy混用：统一内存的页面由驱动管理。如果你手动cudaMemcpy一部分数据，可能会破坏页表一致性。要么全自动，要么全手动。

3.8 小结

统一内存，说白了就是让CPU和GPU共享一套地址空间。cudaMallocManaged帮你分配，驱动帮你搬数据。

但“自动”不等于“免费”。缺页中断很贵，预取和驻留提示是性能的关键。

我个人习惯是：原型阶段无脑用cudaMallocManaged，快速验证。优化阶段再根据profile结果，加上预取和驻留控制。这样既保证了开发效率，又不会牺牲最终性能。

下一章，我们会深入统一内存的底层实现——看看页表是怎么管理的，以及如何用UVMA（Unified Virtual Memory Architecture）做更精细的控制。