3. Hello World:第一个NVSHMEM程序、初始化与终止
好,咱们直接上手写代码。
学任何并行编程框架,第一个程序都是“Hello World”。这不仅是仪式感,更是验证环境、理解初始化流程的关键一步。NVSHMEM也不例外。
我个人习惯,拿到一个新库,先跑通最小示例。这样后面遇到问题,至少知道是环境没配好,还是代码逻辑有坑。
3.1 程序骨架:init 与 finalize
NVSHMEM程序的生命周期,说白了就三件事:启动、干活、收工。
启动靠 nvshmem_init(),收工靠 nvshmem_finalize()。所有PE(Processing Element,你可以理解为GPU进程)必须同步调用这两个函数。
注意: 我曾经见过有人把
nvshmem_init() 放在条件分支里,结果某些PE没调用,整个程序直接挂掉。记住,所有PE必须一致地调用init和finalize。
3.2 获取PE数量和ID
每个PE都有一个唯一的ID,范围从0到 npes-1。这两个信息是后续通信的基础。
nvshmem_n_pes():返回总共有多少个PE。nvshmem_my_pe():返回当前PE的ID。
你想想看,如果没有这两个函数,你怎么知道“我该跟谁说话”?
3.3 完整代码示例
下面这个例子,每个PE打印自己的ID和总PE数。代码很简单,但包含了NVSHMEM程序的所有核心要素。
#include <stdio.h>
#include <nvshmem.h>
int main(void) {
int mype, npes;
// 初始化NVSHMEM环境
nvshmem_init();
// 获取当前PE的ID和总PE数
mype = nvshmem_my_pe();
npes = nvshmem_n_pes();
// 打印Hello World
printf("Hello World from PE %d of %d\n", mype, npes);
// 确保所有PE都打印完毕
nvshmem_barrier_all();
// 终止NVSHMEM环境
nvshmem_finalize();
return 0;
}
3.4 编译与运行
编译时,你需要链接NVSHMEM库。我建议用 nvshmem_run 工具来启动,它会自动处理PE的分配。
# 编译
nvcc -o hello_nvshmem hello_nvshmem.cu -lnvshmem
# 运行,启动4个PE
nvshmem_run -np 4 ./hello_nvshmem
输出大概是这样:
Hello World from PE 0 of 4
Hello World from PE 1 of 4
Hello World from PE 2 of 4
Hello World from PE 3 of 4
注意,PE的打印顺序不一定是0,1,2,3。因为每个PE是独立运行的,谁先抢到输出谁就先打印。这很正常。
小技巧: 如果你想让输出有序,可以在
printf 前后加 nvshmem_barrier_all()。但说实话,调试阶段没必要,反而掩盖了并行执行的本质。
3.5 核心逻辑流程图
下面这张图,展示了NVSHMEM程序从启动到退出的完整流程。我建议你把它记在脑子里,后面所有程序都逃不出这个框架。
3.6 避坑指南
嗯,这里有几个我踩过的坑,你注意一下:
- 忘记调用
nvshmem_init():直接调用nvshmem_my_pe()会返回0,但后续通信全乱套。我曾经调试了一下午,才发现是init没写。 - PE数量不匹配:如果你用
nvshmem_run -np 4启动,但代码里假设只有2个PE,那nvshmem_n_pes()返回4,你的逻辑可能就崩了。 - 忘记同步:虽然这个例子不需要同步,但后续的Put/Get操作,如果没有
nvshmem_barrier_all()或nvshmem_quiet(),数据可能还没到达对方PE,你就开始读了。
核心要点: NVSHMEM程序必须遵循“init → 干活 → finalize”的流程。所有PE必须同步调用init和finalize。获取PE数量和ID是后续通信的基础。
好了,第一个程序跑通了。后面咱们就开始玩真的——单边通信的Put和Get操作。