环境搭建与编译：NVSHMEM的安装与配置、nvcc编译选项、链接NVSHMEM库、检查NVSHMEM版本与设备支持

好，咱们直接进入正题。环境搭建这事儿，说难不难，说简单吧，坑也不少。我刚开始接触NVSHMEM那会儿，光编译就折腾了大半天。后来发现，其实只要理清几个关键点，十分钟就能搞定。

NVSHMEM的安装与配置

先说说安装。NVSHMEM不是单独一个包，它依赖CUDA和MPI。我个人习惯用NVSHMEM官方提供的安装脚本，省心。

下载NVSHMEM源码后，解压，进入目录。我一般这么配：

./install.sh --prefix=/opt/nvshmem --with-cuda=/usr/local/cuda --with-mpi=/usr/local/mpi

这里要注意，--with-cuda和--with-mpi的路径得写对。我在项目中遇到过有人把CUDA路径写成了/usr/local/cuda-11.8，结果编译死活过不去。其实你只要确认nvcc --version能跑，路径就没错。

nvcc编译选项

编译NVSHMEM程序，跟普通CUDA程序不太一样。你得告诉编译器：我要用NVSHMEM的API，还要链接它的库。

最基本的编译命令长这样：

nvcc -arch=sm_80 -rdc=true -c my_kernel.cu -o my_kernel.o
nvcc -arch=sm_80 -rdc=true my_kernel.o -o my_program -lnvshmem -L/opt/nvshmem/lib -I/opt/nvshmem/include

嗯，这里有两个关键点：

• -rdc=true：这个必须加。NVSHMEM的设备端API需要可重定位设备代码。不加的话，链接阶段会报一堆看不懂的错误。
• -lnvshmem：链接NVSHMEM库。注意顺序，一般放在源文件后面。

为什么会这样？因为NVSHMEM的设备端函数是在GPU上执行的，编译器需要生成独立的设备代码段，然后由CUDA驱动在运行时加载。说白了，-rdc=true就是告诉编译器：别急着把所有东西都内联了，留点空间给后面的链接。

链接NVSHMEM库

链接这块，其实就两个事：告诉编译器头文件在哪，告诉链接器库文件在哪。

我习惯用环境变量来管理：

export NVSHMEM_HOME=/opt/nvshmem
export CUDA_HOME=/usr/local/cuda

nvcc -arch=sm_80 -rdc=true \
     -I$NVSHMEM_HOME/include \
     -L$NVSHMEM_HOME/lib \
     -lnvshmem \
     my_program.cu -o my_program

如果你用Makefile，可以这样写：

NVSHMEM_HOME ?= /opt/nvshmem
CUDA_HOME ?= /usr/local/cuda

NVCCFLAGS = -arch=sm_80 -rdc=true -I$(NVSHMEM_HOME)/include
LDFLAGS = -L$(NVSHMEM_HOME)/lib -lnvshmem

my_program: my_program.cu
	nvcc $(NVCCFLAGS) $< -o $@ $(LDFLAGS)

你想想看，如果每次编译都手动敲路径，多累。用环境变量或者Makefile，一劳永逸。

检查NVSHMEM版本与设备支持

环境搭好了，怎么确认能用？我一般写个简单的检查程序：

#include <stdio.h>
#include <nvshmem.h>

int main() {
    int my_pe = nvshmem_my_pe();
    int n_pes = nvshmem_n_pes();
    
    printf("Hello from PE %d of %d\n", my_pe, n_pes);
    printf("NVSHMEM version: %d.%d\n", 
           NVSHMEM_MAJOR_VERSION, NVSHMEM_MINOR_VERSION);
    
    // 检查设备支持
    int dev_count;
    cudaGetDeviceCount(&dev_count);
    printf("Number of CUDA devices: %d\n", dev_count);
    
    for (int i = 0; i < dev_count; i++) {
        cudaDeviceProp prop;
        cudaGetDeviceProperties(&prop, i);
        printf("Device %d: %s (Compute Capability %d.%d)\n", 
               i, prop.name, prop.major, prop.minor);
    }
    
    return 0;
}

编译运行：

mpirun -np 2 ./check_nvshmem

如果看到类似这样的输出，说明环境没问题：

Hello from PE 0 of 2
NVSHMEM version: 1.3
Number of CUDA devices: 2
Device 0: NVIDIA A100 (Compute Capability 8.0)
Device 1: NVIDIA A100 (Compute Capability 8.0)

另外，版本号也很重要。不同版本的NVSHMEM，API可能有细微差别。比如1.2版本引入了nvshmem_team_t，1.3版本增加了nvshmemx_扩展API。我建议至少用1.3以上版本，功能更全，bug也更少。

知识体系总览

为了让你对整个流程有个直观印象，我画了张图：

这张图把整个流程串起来了。从依赖安装开始，到最终验证通过，每一步都有坑，但每一步也都有解法。我个人觉得，最关键的还是-rdc=true和库路径这两个点。只要它们没错，基本就稳了。

---