3. NVSHMEM 基础通信：点对点通信与同步机制
好，咱们正式开始聊 NVSHMEM 的核心通信原语。说实话，这部分内容我反复讲过很多次，但每次备课还是会想起自己第一次踩坑的经历。嗯，咱们先从最基础的点对点通信说起。
3.1 点对点通信：nvshmem_put 与 nvshmem_get
NVSHMEM 的点对点通信，说白了就是两个 PE（处理单元）之间直接传数据。我个人习惯把 nvshmem_put 理解为“主动推送”，把 nvshmem_get 理解为“主动拉取”。
3.1.1 nvshmem_put：数据推送
nvshmem_put 的作用是把本地数据写到远端 PE 的某个地址上。它的函数签名长这样：
void nvshmem_put(void *dest, const void *source, size_t nelems, int pe);
void nvshmem_putmem(void *dest, const void *source, size_t nelems, int pe);这里有个细节——nvshmem_put 是按类型大小传输的，而 nvshmem_putmem 是按字节数传输的。我在项目中遇到过有人把这两个搞混，结果数据量对不上，调试了一整天。
关键点：nvshmem_put 是单边通信。发送方不需要接收方配合，接收方甚至不知道自己被写了数据。这跟 MPI 的 Send/Recv 完全不一样。
3.1.2 nvshmem_get：数据拉取
反过来，nvshmem_get 是从远端 PE 读取数据到本地：
void nvshmem_get(void *dest, const void *source, size_t nelems, int pe);
void nvshmem_getmem(void *dest, const void *source, size_t nelems, int pe);你想想看，这个操作有多方便？你不需要通知远端 PE“我要读你的数据了”，直接读就行。但这也意味着——远端 PE 必须保证数据在读取时是有效的。我曾经在项目里犯过这个错：远端数据还没写完，我就开始 get，结果读到了一堆垃圾。
避坑指南：我曾经在做一个多 GPU 的流体模拟时，用 nvshmem_get 读取邻居的边界数据。因为没加同步，每次跑出来的结果都不一样。后来加了个 barrier，问题就解决了。记住：单边通信不代表不需要同步。
3.2 阻塞与非阻塞语义
NVSHMEM 提供了两种通信模式：阻塞和非阻塞。我个人建议，初学者先用阻塞版本，等理解了数据流再切到非阻塞。
3.2.1 阻塞通信
阻塞版本的 put/get 会在操作完成后才返回。比如：
// 阻塞 put
nvshmem_put(dest, src, 100, 1);
// 到这里时，数据已经保证被写到了 PE 1 的 dest 地址上
// 但注意：这并不保证 PE 1 能看到！嗯，这里要注意：阻塞只保证本地操作完成，不保证远端可见。这是个常见的误解。
3.2.2 非阻塞通信
非阻塞版本以 nvshmemx_ 开头，返回一个 nvshmemx_status_t 状态：
nvshmemx_status_t status;
status = nvshmemx_put(dest, src, 100, 1);
// 这里数据可能还没发出去
// 可以做点别的事
nvshmem_quiet(); // 等待所有未完成的 put/get 完成非阻塞的好处是能实现通信和计算的重叠。我习惯这样用：先发起非阻塞通信，然后做点本地计算，最后调用 nvshmem_quiet() 等待完成。
小技巧：非阻塞 put 之后，如果你确定远端不需要立即看到数据，可以延迟调用 quiet。这样计算和通信就能重叠起来。我在做 stencil 计算时，靠这个技巧把通信开销几乎隐藏掉了。
3.3 同步机制：nvshmem_barrier 与 nvshmem_sync
同步是分布式编程里绕不开的话题。NVSHMEM 提供了两种同步方式：全局同步和集合同步。
3.3.1 nvshmem_barrier：全局屏障
nvshmem_barrier 会让所有 PE 都到达同一个点，然后才能继续往下走：
// 所有 PE 都执行到这里
nvshmem_barrier_all();
// 然后才能继续说白了，这就是个“大家等一下”的机制。我刚开始用的时候觉得它很笨重，后来发现它其实是最安全的同步方式。
性能提示：nvshmem_barrier_all 的开销跟 PE 数量成正比。PE 越多，等待时间越长。所以能用局部同步就别用全局同步。
3.3.2 nvshmem_sync：集合同步
nvshmem_sync 是更灵活的同步方式。它只同步指定的 PE 集合：
// 只同步 PE 0, 1, 2
int pe_list[3] = {0, 1, 2};
nvshmem_sync(pe_list, 3);这个函数我在做子区域通信时特别喜欢用。比如一个 4x4 的网格，我只想同步某一行上的 PE，用 nvshmem_sync 就比全局 barrier 高效得多。
注意：nvshmem_sync 要求所有参与同步的 PE 都调用它。如果有 PE 没调用，那调用的 PE 就会一直等下去。我曾经在调试时遇到过这种死锁，排查了半天才发现是某个 PE 的同步调用被条件判断跳过了。
3.4 知识体系图
下面这张图总结了本章的核心内容，我建议你把它记在脑子里：
3.5 本章小结
咱们这一章聊了 NVSHMEM 最基础的通信原语。说白了就是三件事：怎么发数据（put/get）、怎么控制时机（阻塞/非阻塞）、怎么保证正确（同步）。
我个人觉得，理解单边通信的“不配合”特性是最关键的。你想想看，在 MPI 里你发数据必须有人收，但在 NVSHMEM 里你可以直接写别人的内存。这既是优势也是陷阱——优势是灵活，陷阱是容易忘记同步。
嗯，最后再啰嗦一句：刚开始用 NVSHMEM 时，先加 barrier 保证正确，再优化性能。这是我踩过无数坑之后总结出来的经验。