安装

下载 NVSHMEM

要下载 NVSHMEM，请确保您已注册 NVIDIA 开发者计划。

1. 转到 NVIDIA NVSHMEM 主页。

2. 单击 开始使用。

3. 接受条款和条件。将出现一个包含最新 NVSHMEM 版本的按钮。

4. 单击 下载。

5. 要选择 NVSHMEM 的早期版本，请单击 NVSHMEM 存档。

6. 选择您要安装的 NVSHMEM 版本。

请参阅以下部分，根据您的 Linux 发行版选择正确的软件包。

安装 NVSHMEM

Ubuntu

在 Ubuntu 上安装 NVSHMEM 之前:

1. 向 APT 系统添加一个包含 NVSHMEM 软件包的存储库。

2. 使用 APT 通过以下存储库之一安装 NVSHMEM 软件包

• 本地存储库。

• 网络存储库。

为了在发布新版本时检索升级，我们建议您选择网络存储库。

1. 在以下命令中，将 <architecture> 替换为您的 CPU 架构 (x86_64 或 sbsa)，并将 <distro> 替换为 Ubuntu 版本，例如 ubuntu2004、 ubuntu2204 或 ubuntu2404。

2. 安装存储库。

   • 对于本地 NVSHMEM 存储库

     sudo dpkg -i nvshmem-repo-<version>.deb
     

     注意

     本地存储库安装将提示您安装它嵌入的本地密钥以及软件包的签名。请务必按照说明安装本地密钥，否则安装阶段稍后将失败。

   • 对于网络存储库

     wget https://developer.download.nvidia.com/compute/cuda/repos/<distro>/<architecture>/cuda-keyring_1.1-1_all.deb
     sudo dpkg -i cuda-keyring_1.1-1_all.deb
     

3. 更新 APT 数据库。

   sudo apt update
   

4. 使用 APT 安装 libnvshmem 软件包。此外，要使用 NVSHMEM 编译应用程序，您还可以安装 libnvshmem-dev 软件包。

   注意

   如果您使用的是网络存储库，则以下命令会将 NVSHMEM 升级到最新版本。

   sudo apt install libnvshmem3 libnvshmem3-dev
   

5. 要保留 NVSHMEM 和 CUDA 的早期版本，请指定特定版本，例如

   sudo apt install libnvshmem3=cuda-11_3.1.7-1 libnvshmem3-dev=cuda-11_3.1.7-1
   

   有关确切的软件包版本，请参阅下载页面。

RHEL/CentOS

在 RHEL 或 CentOS 上安装 NVSHMEM 之前:

1. 向 YUM 系统添加一个包含 NVSHMEM 软件包的存储库。

2. 通过 YUM 使用以下存储库之一安装 NVSHMEM 软件包

• 本地存储库。

• 网络存储库。

为了轻松检索发布的新版本升级，我们建议您使用网络存储库。

在以下命令中，<architecture> 应为您的 CPU 架构 (x86_64 或 sbsa)。

1. 安装存储库。

   • 对于本地 NVSHMEM 存储库

     sudo rpm -i nvshmem-repo-<version>.rpm
     

   • 对于网络存储库

     • RHEL 8

       sudo yum-config-manager --add-repo https://developer.download.nvidia.com/compute/cuda/repos/rhel8/<architecture>/cuda-rhel8.repo
       

     • RHEL 9

       sudo dnf config-manager --add-repo http://developer.download.nvidia.com/compute/cuda/repos/rhel9/<architecture>/cuda-rhel9.repo
       

2. 使用 YUM 安装 libnvshmem 软件包。要使用 NVSHMEM 编译应用程序，请安装 libnvshmem-devel 软件包；如果您计划在应用程序中静态链接 NVSHMEM，请安装 libnvshmem-static 软件包

   注意

   如果您使用的是网络存储库，则以下命令会将 NVSHMEM 升级到最新版本。

   sudo yum install libnvshmem3 libnvshmem3-devel libnvshmem3-static
   

   • 要保留 NVSHMEM 和 CUDA 的早期版本，请指定特定版本，例如

     sudo yum install libnvshmem3-cuda-11-3.1.7-1 libnvshmem3-devel-cuda-11-3.1.7-1 libnvshmem3-static-cuda-11-3.1.7-1
     

   有关确切的软件包版本，请参阅下载页面。

其他发行版

1. 通过选择 与操作系统无关的安装程序 下载 tar 文件包。有关更多信息，请参阅 下载 NVSHMEM。

2. 将 NVSHMEM 软件包解压到您的主目录，或者如果以 root 用户身份为所有用户安装，则解压到 /usr/local 中

   # cd /usr/local
   # tar xvf nvshmem-linux-<version>.txz
   

3. 编译应用程序时，请指定 NVSHMEM 的安装目录路径，例如 /usr/local/nvshmem-linux-<version>/。

4. 下载并解压缩存档后，解压后的目录包含以下文件和子目录

文件或目录	描述
src/	包含 NVSHMEM 源代码和标头。
perftest/	包含显示 NVSHMEM API 用法和性能报告的测试。
examples/	包含显示 NVSHMEM 一些常见用例的示例。
scripts/	包含帮助脚本，例如，用于下载、构建和安装 Hydra 的脚本。
changelog	存储库的更改历史记录。
License.txt	NVSHMEM 软件许可协议和版权信息。

1. 设置 CUDA_HOME 环境变量以指向 CUDA 工具包。

2. 设置 GDRCOPY_HOME 环境变量以指向 GDRCopy 安装。

   要在没有 GDRCopy 的情况下构建，请将环境变量设置为 NVSHMEM_USE_GDRCOPY=0。

   注意

   在没有 GDRCopy 的情况下，原子操作仅在 NVLink 连接之间受支持。

3. 如果需要 MPI 和/或 SHMEM 支持，请设置 NVSHMEM_MPI_SUPPORT=1 和/或 NVSHMEM_SHMEM_SUPPORT=1。

4. 设置 MPI_HOME 和 SHMEM_HOME 环境变量以分别指向 MPI 和 OpenSHMEM 安装。

5. 默认情况下，在 NVSHMEM 编译期间使用的 mpicc 位置设置为 $MPI_HOME/bin/mpicc。

   可以通过在环境中指定 MPICC=<path/to/mpicc> 来覆盖此位置。

   注意

   • 当使用 Open MPI 和 OSHMEM 时，路径是相同的。

   • 要使用 OSHMEM，需要使用 UCX 支持构建 Open MPI。

   • NVSHMEM 已经过 Open MPI 4.0.1 和 UCX 1.10 的测试。

   • 其他 MPI 和 OpenSHMEM 安装应该也可以工作。

   • 默认情况下，启用 MPI 支持，禁用 OpenSHMEM 支持。

6. 可选：要启用 UCX 支持，请设置 NVSHMEM_UCX_SUPPORT=1 和 UCX_HOME 到已安装的 UCX 目录。

7. 可选：要启用 libfabric 支持，请设置 NVSHMEM_LIBFABRIC_SUPPORT=1 和 LIBFABRIC_HOME 到已安装的 libfabric 目录。

8. 可选：要启用 NCCL 支持，请设置 NVSHMEM_USE_NCCL=1 和NCCL_HOME 到已安装的 NCCL 目录。

9. 可选：要启用 InfiniBand GPUDirect Async (IBGDA) 传输，请设置 NVSHMEM_IBGDA_SUPPORT=1

10. 可选：要启用 PMIx 支持，请设置 NVSHMEM_PMIX_SUPPORT=1 和PMIX_HOME 到已安装的 PMIx 目录。

11. 配置默认引导程序

    • 可以使用运行时 NVSHMEM_BOOTSTRAP_PMI 环境变量选择 PMI 引导方法。

      PMI-1（可与 Hydra 启动器一起使用）是 NVSHMEM 使用的默认 PMI 标准。

    • 要选择 PMIx 作为默认 PMI 接口，请设置 NVSHMEM_DEFAULT_PMIX=1。

    • 要选择 PMI-2 作为默认 PMI 接口，请设置 NVSHMEM_DEFAULT_PMI2=1。

12. 设置 NVSHMEM_PREFIX 以指定 NVSHMEM 的安装位置。

13. 要构建和安装库，请运行 cmake 和 make -j install。

使用 NVSHMEM cmake 构建系统

NVSHMEM 现在仅支持使用 cmake 版本 3.19 或更高版本进行构建

现在使用 NVSHMEM 需要您使用 cmake 版本 3.19 或更高版本进行构建。有关已弃用和已移除的 make 构建系统之间的比较，请参阅以下部分。

cmake 构建系统向后兼容原始 Makefile 中使用的环境变量。也就是说，无论使用 make 还是 cmake 进行构建，相同的环境都将产生可比较的构建。

Cmake 本身支持一些环境变量和 cmake 变量，以方便发现 NVSHMEM 依赖项（例如 MPI 和 CUDA）。可以在 NVSHMEM 的上下文中使用这些本机设置，但描述它们超出了本文档的范围。

此外，除了 NVSHMEM_HOME（它已被安装前缀 NVSHMEM_PREFIX 取代）之外，所有以前的环境变量在作为 cmake 变量传递时都会受到尊重。

以下步骤概述了使用 cmake 时 NVSHMEM 的典型构建步骤

1. 如果通过环境设置构建配置，请按照上面第 3.2 节中列出的步骤 1-12 进行操作。

2. 要创建用于构建 NVSHMEM 的 makefile，请调用

   cmake [-D{VAR_NAME_1={VAR_VAL_1}} ...-D{VAR_NAME_N={VAR_VAL_N}} -S . -B {PATH_TO_BUILD_DIR}
   

3. 其中 VAR_NAME_X 可以是任何先前接受的环境变量或本机 cmake 变量。注意：如果变量已在环境中设置，则无需在命令行上再次设置它们。如果变量在命令行和环境中都设置了，则命令行将取代环境。要制作 NVSHMEM，请 cd {PATH_TO_BUILD_DIR} && make [-j] [install]

4. 注意：CMAKE 不遵守 NVCC_GENCODE。而是使用 cmake 变量 CUDA_ARCHITECTURE。例如，设置 cmake -DCUDA_ARCHITECTURES=”70;80” 以编译用于 Volta 和 Ampere GPU。

5. 从源软件包使用 CMAKE 构建时，可以通过将 CMAKE 变量 NVSHMEM_BUILD_TESTS 和 NVSHMEM_BUILD_EXAMPLES 分别设置为 0 来禁用性能测试和示例。

6. 二进制软件包可以从源软件包构建，并将 CMake 变量 NVSHMEM_BUILD_PACKAGES 设置为 1。

   用户还可以精细控制是否使用 CMake 变量 NVSHMEM_BUILD_DEB_PACKAGE 和 NVSHMEM_BUILD_RPM_PACKAGE 生成 RPM 或 DEB 文件。默认情况下，这些设置为 1。

在应用程序中使用 NVSHMEM

启动 NVSHMEM 程序

NVSHMEM 支持以下方法来启动您的应用程序

• 使用 PMI-1 兼容的启动器，例如 Hydra。

• 使用 PMI-2 兼容的启动器，例如 Slurm。

• 使用 PMIx 兼容的启动器，例如 Slurm 或 Open MPI mpirun。

• 作为现有 MPI 应用程序的一部分启动。

• 作为现有 OpenSHMEM 应用程序的一部分启动。

PMI-1 和 PMI-2 客户端位于 NVSHMEM 中，并在构建过程中自动构建。PMIx 客户端必须由用户通过安装 Open PMIx 或使用 Open MPI 或 Slurm 安装的 PMIx 客户端来提供。构建 Open MPI 时，请包含 --enable-install-libpmix configure 选项。构建 NVSHMEM 时，设置 NVSHMEM_PMIX_SUPPORT=1 和 PMIX_HOME=/path/to/openmpi。

要在运行时选择正确的 PMI 库，请将 NVSHMEM_BOOTSTRAP_PMI 设置为 PMI、 PMI-2 或 PMIx。要使用 MPI 或 OpenSHMEM 引导 NVSHMEM，请以典型方式启动应用程序，启动 MPI 或 OpenSHMEM，然后调用 nvshmemx_init_attr 函数以告知 NVSHMEM NVSHMEM 正在作为现有 MPI 或 OpenSHMEM 作业的一部分运行。

将 NVSHMEM 与每个 GPU 多个进程一起使用

从 2.5.0 版本开始，NVSHMEM 支持每个 GPU 多个进程 (MPG)，这不需要额外的配置，并且可以在启用或不启用 CUDA 多进程服务 (MPS) 的情况下运行。

但是，如果未启用 MPS，则仅支持以下 API

• 点对点 RMA

• nvshmem_barrier_all() 主机

• nvshmemx_barrier_all_on_stream()

• nvshmem_sync_all() 主机

• nvshmemx_sync_all_on_stream()

要启用完整的 NVSHMEM MPG 支持，必须在系统上安装并运行 NVIDIA MPS 服务器。要启用对完整 API 的支持，还必须配置 MPS 服务器以限制最大 GPU 利用率为 100%。

当 NVSHMEM 库在进程数多于 GPU 数的系统上运行时，它将自动检测并相应地分发进程。它还会自动检测 MPS 服务器守护程序和 GPU 利用率配置的存在，并相应地启用 API。如果在受限 MPG 运行中使用不支持的 API，则会打印错误消息，并且应用程序将退出。

将 NVSHMEM 与您的 C 或 C++ 程序一起使用

1. 包含来自 include/ 的 nvshmem.h 和 nvshmemx.h。

2. 指向 include/ 和 lib/ 路径。

3. NVSHMEM 用户：如果您的 C 或 C++ 程序仅使用 NVSHMEM，请使用 scripts/ 目录下的 install_hydra.sh bash 脚本安装 Hydra 进程管理器。

   1. 提供下载和安装位置作为参数，例如

      ./install_hydra.sh <download_path> <install_path>
      

   2. 要运行 NVSHMEM 作业，请使用 nvshmrun 启动器，该启动器位于 Hydra 安装路径的 bin/ 下。

将 NVSHMEM 与您的 MPI 或 OpenSHMEM 程序一起使用

注意

当前唯一经过测试的 MPI 库是 Open MPI，但任何标准兼容的 MPI 库都应该可以工作。

要运行混合 MPI + NVSHMEM 程序，请使用 MPI 安装中的 mpirun 启动器。

同样，NVSHMEM 可以从 OpenSHMEM 程序中使用，您必须为 OpenSHMEM 库使用相应的启动器。当前唯一经过测试的 OpenSHMEM 版本是 Open MPI 中的 OSHMEM。其他 OpenSHMEM 实现（例如 Sandia OpenSHMEM (SOS)）也应该可以工作，但这些实现尚未经过测试。要运行混合 OpenSHMEM/NVSHMEM 作业，请使用 OpenMPI 安装中的 oshrun 启动器，或遵循您的 OpenSHMEM 库的启动器规范。

NVSHMEM 依赖于用于引导的插件系统。默认情况下，为 NVSHMEM 构建 MPI 引导插件，并将其安装在 $(NVSHMEM_HOME)/lib 中。如果此目录不在您的动态链接器搜索路径中，您可能需要将其添加到 $LD_LIBRARY_PATH。如果使用 nvshmemx_init_attr 初始化函数请求 MPI 引导，或者如果设置了 NVSHMEM_BOOTSTRAP=”MPI”，则会在运行时自动选择此 MPI 插件。

MPI 引导插件的源代码安装在 $(NVSHMEM_HOME)/share/nvshmem/src/bootstrap-plugins 中，并且可以与 NVSHMEM 库分开构建（例如，为了支持其他 MPI 库）。自定义引导插件也是可能的，并且应实现 $(NVSHMEM_HOME)/include/nvshmem_bootstrap.h 中定义的接口。插件必须构建为可重定位的共享对象。

构建外部插件库后，可以通过指定 NVSHMEM_BOOTSTRAP=”plugin” 和 NVSHMEM_BOOTSTRAP_PLUGIN=”[name of plugin]” 在运行时将其指定给 NVSHMEM。例如，NVSHMEM_BOOTSTRAP=”MPI” 等于 NVSHMEM_BOOTSTRAP=”plugin” 和 NVSHMEM_BOOTSTRAP_PLUGIN=”nvshmem_bootstrap_mpi.so”。

运行性能测试

在您可以运行性能测试之前，您首先必须构建它们。

1. 如果 NVSHMEM 库是使用 NVSHMEM_MPI_SUPPORT=1 构建的，请设置 CUDA_HOME、 NVSHMEM_HOME 和 MPI_HOME 环境变量以构建 NVSHMEM 性能测试

   CUDA_HOME=<path to supported CUDA installation>
   NVSHMEM_HOME=<path to directory where NVSHMEM is installed>
   MPI_HOME=<path to MPI installation>
   

   如果您在构建 perftest/ 时使用 MPI 和 OpenSHMEM 支持（NVSHMEM_MPI_SUPPORT=1 和 NVSHMEM_SHMEM_SUPPORT=1）构建了 NVSHMEM，则必须启用 MPI 和 OpenSHMEM 支持。

   构建时不进行 SHMEM 互操作： 要构建不进行 SHMEM 互操作的 NVSHMEM 性能测试，请将环境变量 NVSHMEM_SHMEM_SUPPORT 设置为 0。默认情况下，性能测试安装在 perftest/perftest_install 下。要安装到不同的路径，请设置 NVSHMEM_PERFTEST_INSTALL 以指向正确的路径。

2. 更新 LD_LIBRARY_PATH 以指向 $CUDA_HOME/lib64、 $MPI_HOME/lib 和 $NVSHMEM_HOME/lib。

3. 假设 Hydra 安装在 HYDRA_HOME 下，使用以下命令将性能测试作为 NVSHMEM 作业、混合 MPI+NVSHMEM 作业或混合 OpenSHMEM+NVSHMEM 作业运行（以 perftest/device/pt-to-pt/put.cu 为例）

使用 Hydra (PMI-1) 的 NVSHMEM 作业

$HYDRA_HOME/bin/nvshmrun -n <up to number of P2P or InfiniBand
NIC accessible GPUs>
$NVSHMEM_PERFTEST_INSTALL/device/pt-to-pt/shmem_put_bw

使用 Slurm 的 NVSHMEM 作业

srun -n <up to number of P2P or InfiniBand NIC accessible GPUs>
$NVSHMEM_PERFTEST_INSTALL/device/pt-to-pt/shmem_put_bw

注意

当 Slurm 是使用与 NVSHMEM 默认值不匹配的 PMI 构建的时，例如，如果 Slurm 是使用 PMIx 支持构建的，并且在构建 NVSHMEM 时未设置 NVSHMEM_DEFAULT_PMIX=1，则可以使用 NVSHMEM_BOOTSTRAP_PMI 覆盖默认值。可能的值为 PMIX、PMI-2 和 PMI。Slurm --mpi= 选项 srun 可用于告知 Slurm 要使用哪个 PMI 接口。

混合 MPI/NVSHMEM 作业

$MPI_HOME/bin/mpirun -n <up to number of GPUs accessible by P2P
or InfiniBand NIC> -x NVSHMEMTEST_USE_MPI_LAUNCHER=1
$NVSHMEM_PERFTEST_INSTALL/device/pt-to-pt/shmem_put_bw

混合 OpenSHMEM/NVSHMEM 作业

$MPI_HOME/bin/oshrun -n <up to number of GPUs accessible by P2P
or InfiniBand NIC> -x USE_SHMEM_IN_TEST=1
$NVSHMEM_PERFTEST_INSTALL/device/pt-to-pt/shmem_put_bw

“Hello World” 示例

1. 将以下代码另存为 nvshmemHelloWorld.cu

   #include <stdio.h>
   #include <cuda.h>
   #include <nvshmem.h>
   #include <nvshmemx.h>
   
   __global__ void simple_shift(int *destination) {
       int mype = nvshmem_my_pe();
       int npes = nvshmem_n_pes();
       int peer = (mype + 1) % npes;
   
       nvshmem_int_p(destination, mype, peer);
   }
   
   int main(void) {
       int mype_node, msg;
       cudaStream_t stream;
   
       nvshmem_init();
       mype_node = nvshmem_team_my_pe(NVSHMEMX_TEAM_NODE);
       cudaSetDevice(mype_node);
       cudaStreamCreate(&stream);
   
       int *destination = (int *) nvshmem_malloc(sizeof(int));
   
       simple_shift<<<1, 1, 0, stream>>>(destination);
       nvshmemx_barrier_all_on_stream(stream);
       cudaMemcpyAsync(&msg, destination, sizeof(int), cudaMemcpyDeviceToHost, stream);
   
       cudaStreamSynchronize(stream);
       printf("%d: received message %d\n", nvshmem_my_pe(), msg);
   
       nvshmem_free(destination);
       nvshmem_finalize();
       return 0;
   }
   

2. 使用以下命令构建 nvshmemHelloWorld.cu

   当使用动态链接时

   nvcc -rdc=true -ccbin g++ -gencode=$NVCC_GENCODE -I
   $NVSHMEM_HOME/include nvshmemHelloWorld.cu -o
   nvshmemHelloWorld.out -L $NVSHMEM_HOME/lib -lnvshmem_host -lnvshmem_device
   

   当使用静态链接时

   nvcc -rdc=true -ccbin g++ -gencode=$NVCC_GENCODE -I
   $NVSHMEM_HOME/include nvshmemHelloWorld.cu -o
   nvshmemHelloWorld.out -L $NVSHMEM_HOME/lib -lnvshmem -lnvidia-ml -lcuda -lcudart
   

   其中 arch=compute_70,code=sm_70 是 V100 GPU 的 NVCC_GENCODE 的值。

3. 使用以下命令之一运行 nvshmemHelloWorld 示例

   当在一个主机上使用两个 GPU（通过 PCI-E、NVLink 或 Infiniband 连接）运行时

   $HYDRA_HOME/bin/nvshmrun -n 2 -ppn 2 ./nvshmemHelloWorld.out
   

   当在两个主机上运行时，每个主机有一个 GPU，并通过 InfiniBand 连接

   $HYDRA_HOME/bin/nvshmrun -n 2 -ppn 1 –-hosts hostname1,hostname2 ./nvshmemHelloWorld.out
   

基于 NVSHMEM LLVM Bitcode 库的 Hello World 示例

从 NVSHMEM 3.2 版本开始，Python pip wheels 和 conda 软件包附带一个 LLVM-IR 库，该库公开了设备端 API (libnvshmem_device.bc)。nvidia-nvshmem-cu12 Python 软件包附带一个位代码库，该库支持 Hopper 和 Blackwell 架构，而 nvidia-nvshmem-cu11 库支持 Volta 和 Ampere 架构。以下示例说明如何降低第 1 部分中创建的“hello world”应用程序通过 LLVM IR，将其链接到 libnvshmem_device.bc，并从链接的 LLVM-IR 编译 cubin，该 cubin 可以使用 nvshmemx_cumodule_init 导入到 NVSHMEM 应用程序中。

1. 使用 Clang 将应用程序的设备代码降低到 LLVM-IR。NVSHMEM_BITCODE_APPLICATION 变量指示编译器仅包含用于降低到 LLVM-IR 的适当标头。

   clang -c emit-llvm -std=c++11 -x cuda --cuda-path=${CUDA_HOME} --cuda-device-only --cuda-gpu-arch={desired architecture} -I$NVSHMEM_HOME/include -DNVSHMEM_BITCODE_APPLICATION nvshmemHelloWorld.cu nvshmemHelloWorldPrelink.bc
   

2. 使用 LLVM 实用程序 llvm-link 和 opt 链接和优化应用程序。only-needed 参数将显着减小生成的二进制文件的大小。

   llvm-link --only-needed nvshmemHelloWorldPrelink.bc $NVSHMEM_HOME/lib/libnvshmem_device.bc -o nvshmemHelloWorldPreopt.bc
   opt -O3 nvshmemHelloWorldPreopt.bc -o nvshmemHelloWorld.bc
   

3. 使用 llc LLVM 实用程序将代码降低到 PTX。

   llc -O3 --disable-tail-calls -mcpu={desired architecture}-mattr=ptx{YY} nvshmemHelloWorld.bc -o nvshmemHelloWorld.ptx
   

4. 使用 CUDA 工具包中的 ptxas 将 PTX 汇编为 SASS。这可以在一个或两个步骤中完成，但单独执行链接步骤已显示出更高的性能。

   注意

   链接步骤中没有单独的库。

   ptxas -c -arch={desired architecture} nvshmemHelloWorld.ptx -o nvshmemHelloWorld_prelink.cubin
   nvlink -arch={desired architecture} nvshmemHelloWorld_prelink.cubin -o nvshmemHelloWorld.cubin