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程序从启动到退出的完整流程。我建议你把它记在脑子里,后面所有程序都逃不出这个框架。

NVSHMEM 程序生命周期 nvshmem_init() nvshmem_my_pe() / nvshmem_n_pes() Put / Get / Barrier 等操作 nvshmem_finalize()

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操作。

专注资料整理