1、共享内存基础：CPU-GPU异构计算概述、统一内存模型(UM)概念、共享内存的物理与逻辑视图
1.1 CPU-GPU异构计算：为什么我们需要共享内存？
说实话，我刚入行那会儿，CPU和GPU是各玩各的。CPU管逻辑控制，GPU管并行计算。数据得先从CPU拷贝到GPU显存，算完再拷回来。这个过程，我称之为「数据搬运工」模式。
你想想看，一个深度学习训练任务，数据在CPU和GPU之间来回倒腾，光PCIe带宽就占了大半。我曾在一次图像处理项目中，发现数据搬运占了总耗时的40%以上。这太浪费了。
异构计算的核心思想很简单：让合适的硬件干合适的活。CPU擅长复杂逻辑和分支预测，GPU擅长大规模并行计算。但问题来了——它们怎么高效地共享数据？
这就引出了我们今天的主角：共享内存。
核心观点：共享内存不是简单的「内存共用」，而是一种软硬件协同的机制。它让CPU和GPU能直接访问同一块物理内存，省去了显式数据拷贝的开销。
1.2 统一内存模型(UM)：一个美丽的抽象
统一内存模型（Unified Memory，简称UM），说白了就是让程序员写代码时，不用再操心数据在CPU还是GPU上。你只管分配内存，系统自动帮你迁移。
我记得第一次接触CUDA 6.0的UM时，心里直犯嘀咕：「这玩意儿靠谱吗？」后来在项目中实测，对于某些访存模式，UM确实能省掉大量手动cudaMemcpy的代码。但要注意，它不是银弹。
UM的工作原理可以概括为三点：
- 按需迁移：数据在CPU和GPU之间按需移动，不是一次性全搬
- 页粒度管理：以4KB或64KB的页为单位，由驱动和硬件协同管理
- 透明访问：CPU和GPU都能用同一个指针访问数据
来看一个简单的例子：
// 传统方式：手动拷贝
float *d_data;
cudaMalloc(&d_data, N * sizeof(float));
cudaMemcpy(d_data, h_data, N * sizeof(float), cudaMemcpyHostToDevice);
kernel<<<grid, block>>>(d_data);
cudaMemcpy(h_data, d_data, N * sizeof(float), cudaMemcpyDeviceToHost);
// UM方式：统一内存
float *data;
cudaMallocManaged(&data, N * sizeof(float));
// 初始化在CPU上
for (int i = 0; i < N; i++) data[i] = i;
// GPU直接访问，系统自动迁移
kernel<<<grid, block>>>(data);
// CPU直接读取结果
cudaDeviceSynchronize();
我的经验：UM在数据访问模式简单时效果很好。但如果你的数据频繁在CPU和GPU之间来回切换，反而会触发大量页迁移，性能还不如手动拷贝。我曾经在一个迭代求解器中踩过这个坑，后来改成了手动管理。
1.3 共享内存的物理视图：硬件到底长什么样？
嗯，这里要注意。共享内存在物理上并不是一块「新」的内存。它其实是利用了现有的硬件资源。
从物理角度看，有两种实现方式：
| 实现方式 | 物理位置 | 延迟 | 带宽 | 典型场景 |
|---|---|---|---|---|
| 统一虚拟地址（UVA） | CPU内存 + GPU显存 | 高（跨PCIe） | 受限于PCIe | 数据量小、访问不频繁 |
| 硬件共享内存（如HBM） | 同一物理芯片 | 低 | 高 | 高性能计算、实时处理 |
我个人的习惯是：先搞清楚你的硬件平台支持哪种。比如NVIDIA的Grace Hopper超级芯片，CPU和GPU通过NVLink-C2C互联，共享同一块HBM内存。这种物理共享，延迟只有传统PCIe方案的十分之一。
避坑指南：我曾经在一个项目中，想当然地以为所有平台都支持硬件级共享内存。结果部署到一台老服务器上，发现UM实际上走的是PCIe，性能直接崩了。所以，一定要查硬件规格书。
1.4 共享内存的逻辑视图：程序员看到的世界
逻辑视图就简单多了。从程序员的角度看，共享内存就是一个统一的地址空间。CPU和GPU都能用同一个指针读写。
但这里有个关键点：一致性。CPU和GPU同时访问同一块内存，会不会出问题？
答案是：会。而且很常见。
我举个例子：
// CPU写，GPU读
data[0] = 42; // CPU写入
kernel<<<1, 1>>>(data); // GPU读取
// 问题：GPU读到的是42吗？不一定！
为什么会这样？因为CPU和GPU都有自己的缓存。CPU写入了数据，但可能还在自己的缓存里，没写回主存。GPU去读，读到的可能是旧值。
解决方案有两种：
- 显式同步：调用cudaDeviceSynchronize()或cudaStreamSynchronize()
- 原子操作：使用atomicAdd等原子指令保证一致性
我的建议：在UM模式下，尽量让数据「归属」明确。要么大部分时间在GPU上，要么在CPU上。频繁切换归属，会触发大量缺页中断和页迁移，性能惨不忍睹。
1.5 小结：共享内存的本质
说了这么多，其实共享内存的本质就一句话：让CPU和GPU能高效地「对话」。
物理上，它可能是同一块芯片上的HBM，也可能是通过高速互联连接的独立内存。逻辑上，它给程序员提供了一个统一的地址空间，省去了手动数据搬运的麻烦。
但要注意，共享内存不是免费的午餐。它带来了性能提升的潜力，也引入了一致性和迁移开销的问题。我个人的经验是：先理解你的数据访问模式，再选择合适的内存模型。
下一章，我们会深入讨论共享内存的编程模型和优化技巧。到时候我会分享一些实际项目中的调优案例，保证让你有收获。