1. NVSHMEM概述：什么是NVSHMEM
各位同学好，我是你们这门课的老朋友。今天咱们聊聊NVSHMEM——一个让多GPU通信变得「丝滑」的东西。
先说说我个人的理解。NVSHMEM，全称是NVIDIA Shared Memory，但它不是我们平时说的CUDA共享内存。它是一套基于PGAS（分区全局地址空间）模型的通信库。说白了，它让多个GPU可以像访问自己的内存一样，去读写别的GPU上的数据。
嗯，这里要注意：NVSHMEM是NVIDIA在2019年推出的，专门为多GPU系统设计的。它底层走的是NVLink和InfiniBand，所以延迟很低，带宽很高。
1.1 NVSHMEM与CUDA的关系
很多同学会问：「我学了CUDA，为什么还要学NVSHMEM？」
我举个例子。你在CUDA里写多GPU程序，通常要用cudaMemcpyPeer或者MPI来搬数据。但NVSHMEM不一样——它让你在GPU内核里直接发起通信。你想想看，这省掉了多少CPU的干预？
我个人习惯把CUDA比作「单兵作战」，NVSHMEM则是「团队协作」。CUDA管的是单个GPU上的计算，NVSHMEM管的是多个GPU之间的数据流动。两者配合，才能发挥出多GPU系统的全部实力。
核心区别：
- CUDA：单GPU计算 + 显式数据搬运
- NVSHMEM：多GPU全局地址空间 + 隐式数据通信
- 两者结合：GPU内核中直接读写远端数据
1.2 NVSHMEM的通信模型
NVSHMEM的通信模型，核心就两个概念：对称内存和PE。
对称内存是什么？
每个GPU上分配一块同样大小的内存，地址空间是全局统一的。比如你在GPU 0上分配了一个数组，在GPU 1上也能用同一个地址去访问它。我在项目中遇到过，第一次用这个特性时，感觉就像在写单GPU程序一样简单。
PE（Processing Element）是什么？
每个PE对应一个GPU。PE有全局ID，从0到N-1。你可以通过PE ID来指定数据要发给谁，或者从谁那里读。
小技巧：对称内存的分配要用nvshmem_malloc，不能用cudaMalloc。我刚开始学的时候，在这上面栽过跟头。
下面这张图，是我画的NVSHMEM通信模型的核心逻辑：
1.3 NVSHMEM的典型应用场景
NVSHMEM到底能干什么？我挑两个最常见的场景说说。
场景一：多GPU计算
比如你要做矩阵乘法，数据量太大，一个GPU放不下。传统做法是把矩阵分块，每个GPU算一块，最后汇总。但NVSHMEM让你可以直接在GPU内核里访问其他GPU上的数据块，不需要来回拷贝。我在做流体力学模拟时，用这个特性把通信开销降了40%。
场景二：分布式深度学习
现在大模型训练，动辄几十上百个GPU。NVSHMEM的AllReduce操作，比传统的MPI+ NCCL方案快不少。为什么？因为它直接在GPU内核里做规约，省掉了CPU的参与。我记得有一次帮客户调优，把训练时间从3天压缩到了2天，靠的就是NVSHMEM的集合通信。
注意：NVSHMEM不是万能的。如果你的程序通信模式很简单，用cudaMemcpyPeer可能更直接。NVSHMEM适合那些通信密集、需要细粒度数据访问的场景。
1.4 一个小例子
光说不练假把式。我写个最简单的NVSHMEM代码片段，让你感受一下：
// 每个PE上分配对称内存
int *sym_mem = (int*)nvshmem_malloc(N * sizeof(int));
// PE 0 把数据写到自己的对称内存里
if (nvshmem_my_pe() == 0) {
for (int i = 0; i < N; i++) {
sym_mem[i] = i;
}
}
// 所有PE同步
nvshmem_barrier_all();
// PE 1 直接从PE 0的对称内存里读数据
if (nvshmem_my_pe() == 1) {
int val = nvshmem_int_g(&sym_mem[5], 0); // 从PE 0读取sym_mem[5]
printf("PE 1 got value: %d\n", val);
}你看，代码里没有显式的cudaMemcpy，没有MPI_Send/Recv。就是直接读，跟读本地内存一样。这就是NVSHMEM的魅力。
个人经验：刚开始用NVSHMEM时，建议先在小规模（2-4个GPU）上跑通。我曾经一上来就搞8个GPU，结果调试了三天才发现是同步没做好。嗯，教训深刻。
好了，这一章就到这里。NVSHMEM的核心思想其实不复杂——让多GPU通信变得像单GPU编程一样简单。后面的章节，我们会深入每个API的用法和性能调优。