问题排查和常见问题解答¶

常见问题解答¶

问：以下运行时警告意味着什么？

WARN: IB HCA and GPU are not connected to a PCIe switch so InfiniBand
performance can be limited depending on the CPU generation

答：此警告与平台的 HCA 到 GPU 映射有关。有关更多信息，请参阅SHMEM_HCA_PE_MAPPING变量，位于环境变量。

问：以下运行时错误表明什么？

src/bootstrap/bootstrap_loader.cpp:46: NULL value Bootstrap unable to load 'nvshmem_bootstrap_mpi.so'
    nvshmem_bootstrap_mpi.so: cannot open shared object file: No such file or directory
src/bootstrap/bootstrap.cpp:26: non-zero status: -1 bootstrap_loader_init returned error
src/init/init.cpp:101: non-zero status: 7 bootstrap_init failed
src/init/init.cpp:767: non-zero status: 7 nvshmem_bootstrap failed
src/init/init.cpp:792: non-zero status: 7: Success, exiting... aborting due to error in nvshmemi_init_thread

答：NVSHMEM 为多个引导程序使用动态加载的引导模块，包括 MPI、OpenSHMEM 和 PMIx。上述错误表明无法加载 MPI 的引导模块。确保 NVSHMEM 库目录位于动态链接器的系统搜索路径中，或者 LD_LIBRARY_PATH 变量包含 NVSHMEM 库目录。或者，您可以设置 NSHMEM_BOOTSTRAP_PLUGIN 环境变量以帮助 NVSHMEM 找到插件。如果插件在 NVSHMEM 安装过程中未构建，或者您需要支持与安装期间使用的库（例如 MPI 库）不同的库，则可以从 share/nvshmem 目录下提供的源代码构建 NVSHMEM 引导插件。

问：以下运行时错误表明什么？

src/comm/transports/ibrc/ibrc.cpp:: NULL value mem registration failed.

答：如果平台上未启用 GPUDirect RDMA，则会发生这种情况，从而阻止使用 InfiniBand 驱动程序注册 cudaMalloc 内存。这通常表明缺少 nv_peer_mem 内核模块。安装 nv_peer_mem 后，lsmod 的输出类似于以下内容

~$ lsmod | grep nv_peer_mem
nv_peer_mem               20480 0
ib_core                   241664 11
rdma_cm,ib_cm,iw_cm,nv_peer_mem,mlx4_ib,mlx5_ib,ib_ucm,ib_umad,ib_uverbs,rdma_ucm,ib_ipoib
nvidia                  17596416 226
nv_peer_mem,gdrdrv,nvidia_modeset,nvidia_uvm

nv_peer_mem 可在此处获取：here。

问：以下运行时错误表明什么？

src/comm/transports/ibrc/ibrc.cpp: NULL value get_device_list failed

答：当系统上存在 ibverbs 库，但该库无法检测到系统上的任何 InfiniBand 设备时，会发生这种情况。确保 InfiniBand 设备可用和/或处于工作状态。

问：以下运行时错误表明什么？

src/comm/transports/libfabric/libfabric.cpp:482: non-zero status: -12 Error registering memory region: Cannot allocate memory
src/mem/mem.cpp:464: non-zero status: 7 transport get memhandle failed
src/comm/device/proxy_device.cu:695: NULL value failed allocating proxy_channel_g_bufchannel creation failed

答：当通过 CUDA VMM 接口分配的内存不受网络堆栈支持时，可能会发生这种情况。您可以尝试在环境中设置 NVSHMEM_DISABLE_CUDA_VMM=1，以在 NVSHMEM 分配内存时禁用 CUDA VMM 接口的使用。

问：什么会导致与 ibv_poll_cq 相关的运行时错误？

例如

src/comm/transports/ibrc/ibrc.cpp:962: non-zero status: 10 ibv_poll_cq failed, status: 10
src/comm/transports/ibrc/ibrc.cpp:1021: non-zero status: 7 progress_send failed, outstanding_count: 0
src/comm/transports/ibrc/ibrc.cpp:1294: non-zero status: 7 check_poll failed
src/comm/proxy/proxy.cu:progress_quiet:612: aborting due to error in progress_quiet

答：这些类型的错误发生在 NVSHMEM ibrc 远程传输中，以响应 infiniband 工作请求失败。工作请求可能因多种原因而失败，如 ibv_poll_cq 手册页中所述。在 NVSHMEM 中，我们在 “ibv_poll_cq failed, status: 10” 行中发布 ibv_poll_cq 错误的错误代码。该状态值可以在 ibv_poll_cq 文档中查找，以确定故障原因。

在 NVSHMEM 中，IB 故障最常见的原因是状态：10 - 远程保护错误或状态：4 - 本地保护错误。当 RMA 或原子操作中使用不在 NVSHMEM 对称堆上的地址作为远程缓冲区时，会发生第一个错误。当本地缓冲区既不在对称堆中，也未在 NVSHMEM 中注册为本地缓冲区时，会发生第二个错误。如果在 NVSHMEM RMA/AMO 操作中使用了从 nvshmem_ptr 或 nvshmemx_mc_ptr 返回的地址，也可能发生这两种错误。

问：我的应用程序使用 CMake 构建系统。将 NVSHMEM 添加到构建系统会导致 CMake 版本低于 3.11 时中断。为什么？

答：设备链接支持在 3.11 版本中添加，NVSHMEM 需要此支持。

问：为什么我的 NVSHMEM 应用程序的 CMake 构建在版本 3.12 中失败，但在早期版本中没有失败？

答：新的 CMake 策略将 -pthread 添加到 nvcc 设备链接，从而导致链接失败。在 3.12 之前，默认策略未添加 -pthread。对于 3.12 及更高版本，请将 cmake_policy(SET CMP0074 OLD) 添加到 CMakeLists.txt。

问：构建 CUDA 或 NVSHMEM 应用程序需要哪些 CMake 设置？

答：将以下内容添加到 CMake 文件中，将目标 GPU 架构替换为 compute_70 和 sm_70。

string(APPEND CMAKE_CUDA_FLAGS "-gencode arch=compute_70,code=sm_70")

问：为什么我的 NVSHMEM Hydra 作业在 Summit 上变得无响应？

答：Summit 需要将附加选项 --launcher ssh 传递给命令行中的 nvshmrun。

问：多个 PE 可以使用 NVSHMEM 共享同一个 GPU 吗？

答：在 NVSHMEM 2.2.1 之前，NVSHMEM 假定 PE 与 GPU 之间存在 1:1 映射。不支持启动的 PE 数量多于可用 GPU 的作业。自 NVSHMEM 2.4.1 起，有限支持每个 GPU 多个进程 (GPU)。有关更多详细信息，请参阅多进程 GPU 支持。

问：使用 NVSHMEM 的 CUDA_VISIBLE_DEVICES 的正确方法是什么？

答：当将 CUDA_VISIBLE_DEVICES 与 NVSHMEM 一起使用时，应将相同的 CUDA_VISIBLE_DEVICES 值传递给所有 PE。请注意，我们可能会在未来的 NVSHMEM 版本中更改此设置。

先决条件常见问题解答¶

问：NVSHMEM 需要 CUDA 吗？

答：是的。必须安装 CUDA 才能使用 NVSHMEM。有关版本要求，请参阅安装指南。 NVSHMEM 是一个通信库，旨在用于两个或多个 GPU 之间的高效数据移动和同步。它目前不适用于不涉及 GPU 的数据移动。

问：NVSHMEM 需要 MPI 吗？

答：否。没有 MPI 依赖项的 NVSHMEM 应用程序可以使用 NVSHMEM，并可以使用 Hydra 启动器启动。可以使用 NVSHMEM 附带的安装脚本安装 Hydra。此脚本安装 Hydra 启动器二进制文件，名称为 nvshmrun 和 nvshmrun.hydra。也可以使用外部安装的 Hydra 启动器副本，它通常安装 Hydra 启动器二进制文件，名称为 mpiexec、mpirun 和 mpiexec.hydra。

问：我的 NVSHMEM 作业在 NVIDIA Volta GPU 上运行，但在 NVIDIA Kepler GPU 上挂起。为什么会这样？

答：CUDA 内核中的 NVSHMEM 同步 API 仅在 NVIDIA Volta 和更新的 GPU 上受支持。

运行 NVSHMEM 程序常见问题解答¶

问：当我尝试使用 PMIx 和 mpirun 的 OpenMPI 实现启动 NVSHMEM 应用程序时，出现缺少符号错误。如何解决此问题？

答：OpenMPI 中的内部 PMIx 实现与我们在 NVSHMEM 中使用的实现之间存在已知的不兼容性。尝试使用内部 PMIx 通过 OpenMPI 运行 nvshmem 应用程序很可能会导致以下错误

pmix_mca_base_component_repository_open: unable to open mca_gds_ds21:
perhaps a missing symbol, or compiled for a different version of OpenPMIx (ignored)

可以通过使用 --with-pmix={PMIX_PATH} 配置选项，使用 PMIx 的外部实现编译 OpenMPI 来解决此问题。

与 MPI 互操作性常见问题解答¶

问：NVSHMEM 可以用于 MPI 应用程序吗？

答：是的。 NVSHMEM 提供了一个初始化 API，它将 MPI 通信器作为属性。通信器中的每个 MPI 进程都成为 OpenSHMEM PE。目前，NVSHMEM 已在 OpenMPI 4.0.0 上进行了测试。原则上，预计其他 OpenMPI 衍生产品（例如 SpectrumMPI（在 Summit 和 Sierra 上可用））也能正常工作。

问：将 NVSHMEM 分配的缓冲区传递给 MPI 会导致错误。如何解决？

NVSHMEM 对称内存使用 CUDA IPC 映射。当这些缓冲区传递给 MPI 时，MPI 可能会尝试第二次为 CUDA IPC 映射这些缓冲区。在 CUDA 11.2 之前，这将导致错误。现在 CUDA 11.2 及更高版本中支持此功能。

从 NVSHMEM v2.2.1 开始，对称内存使用 CUDA VMM API 映射。这种映射对称堆的方法与某些 CUDA 感知的 MPI 库不兼容，从而导致如下崩溃

cma_ep.c:97   process_vm_readv(pid=434391 length=524288) returned -1: Bad address
==== backtrace (tid: 434392) ====
0x0000000000002558 uct_cma_ep_tx()  src/uct/sm/scopy/cma/cma_ep.c:95
0x000000000001a17c uct_scopy_ep_progress_tx()  src/uct/sm/scopy/base/scopy_ep.c:151
0x000000000004f7c1 ucs_arbiter_dispatch_nonempty()  src/ucs/datastruct/arbiter.c:321
0x0000000000019ce4 ucs_arbiter_dispatch()  src/ucs/datastruct/arbiter.h:386
0x000000000004fcfe ucs_callbackq_slow_proxy()  src/ucs/datastruct/callbackq.c:402
0x0000000000034eba ucs_callbackq_dispatch()  src/ucs/datastruct/callbackq.h:211
0x0000000000034eba uct_worker_progress()  src/uct/api/uct.h:2592
0x0000000000034eba ucp_worker_progress()  src/ucp/core/ucp_worker.c:2635
0x000000000003694c opal_progress()  ../opal/runtime/opal_progress.c:231
0x000000000004c983 ompi_request_default_test()  ../ompi/request/req_test.c:96
0x000000000004c983 ompi_request_default_test()  ../ompi/request/req_test.c:42
0x0000000000071d35 PMPI_Test()  ../ompi/mpi/c/profile/ptest.c:65
...

可以通过设置 NVSHMEM_DISABLE_CUDA_VMM=1 来避免这些问题。

与 OpenSHMEM 互操作性常见问题解答¶

问：NVSHMEM 可以用于 OpenSHMEM 应用程序吗？

答：是的。 NVSHMEM 提供了一个初始化 API，该 API 支持在 OpenMPI/OSHMEM 作业之上运行 NVSHMEM。每个 OSHMEM PE 都以 1:1 的方式映射到 NVSHMEM PE。 NVSHMEM 已在 OpenMPI 4.0.0/OSHMEM 和 OpenMPI3+/OSHMEM 上进行了测试，后者依赖于 UCX（NVSHMEM 已在 UCX 1.4.0 上进行了测试）。必须使用 --with-ucx 标志配置 OpenMPI-4.0.0 安装，以启用 OpenSHMEM + NVSHMEM 互操作性。

问：当使用 Open MPI SHMEM (OSHMEM) 时，如何初始化 NVSHMEM？

答：OSHMEM 存在一个已知问题，即在 shmem_init() 期间，它会错误地将 CUDA 上下文设置为设备 0。如果在调用 shmem_init() 后初始化 NVSHMEM，则 NVSHMEM 会检测到所有 PE 的设备都已设置为设备 0，因此会在多进程 GPU 共享模式下初始化自身，并将所有 PE 分配给 GPU 0。因此，当与 OSHMEM 一起运行时，用户必须在调用 NVSHMEM 初始化之前设置所需的 CUDA 上下文。

GPU-GPU 互连常见问题解答¶

问：我可以使用 NVSHMEM 在不同插槽上的 GPU 之间传输数据吗？

答：是的，如果存在 InfiniBand NIC，可以访问两个插槽上的 GPU。否则，NVSHMEM 要求所有 GPU 都是 P2P 可访问的。

问：我可以使用 NVSHMEM 在通过 PCIe 连接的 P2P 可访问 GPU 之间传输数据吗？

答：是的，NVSHMEM 支持 PCIe。但是，当使用 PCIe 进行 P2P 通信时，要么需要 InfiniBand 支持才能使用 NVSHMEM 原子内存操作 API，要么必须使用 NVSHMEM 的 UCX 传输（在 IB 不存在时，该传输将使用套接字进行原子操作）。

问：我可以使用 NVSHMEM 在通过 InfiniBand 连接的不同主机上的 GPU 之间传输数据吗？

答：是的。 NVSHMEM 支持 InfiniBand。 InfiniBand 不支持 Strided-RMA (shmem_iput/iget) 操作。

问：我可以在没有 InfiniBand NIC 的主机上运行 NVSHMEM 吗？

答：是的。 P2P 平台上的支持保持不变。

问：我可以在具有 InfiniBand NIC 的主机上运行 NVSHMEM 吗？在这些主机上，某些 NIC 已禁用或配置为非 InfiniBand 模式？

答：是的。请参阅“有用的环境变量”部分，了解如何显式指定 NIC 与 PE 的亲缘性。

NVSHMEM API 使用常见问题解答¶

问：例如，nvshmemx_putmem_on_stream 和 nvshmemx_putmem_nbi_on_stream 之间有什么区别？似乎两者都相对于主机线程是异步的，并且相对于给定的流是有序的。

答：函数 nvshmemx_putmem_nbi_on_stream 的实现方式更为延迟，它不会立即发出传输，而是使其等待流末尾的事件。如果同时（在另一个流上）有另一个传输正在进行，则可能会共享带宽。如果应用程序可以避免这种情况，则 nvshmemx_putmem_nbi_on_stream 提供了向 NVSHMEM 表达此意图的灵活性。但 NVSHMEM 目前不跟踪所有 CUDA 流上的活动。当前的实现记录用户提供的流上的事件，使 NVSHMEM 内部流等待该事件，然后在内部流上发出 put 操作。如果所有 nbi put 操作都落在同一个内部流上，则它们会被序列化，以便独占使用带宽。

问：我可以并发地在多个流上发出多个 nvshmemx_barrier_all_on_stream，然后在每个流上 cudaStreamSynchronize 吗？

答：多个并发 nvshmemx_barrier_all_on_stream / nvshmem_barrier_all 调用无效。在任何给定时间，同一组 PE 之间只能有一个屏障（或任何其他集合操作）处于飞行状态。要在部分重叠的团队之间使用并发屏障，可以使用 syncneighborhood_kernel 作为模板来实现自定义屏障。有关自定义屏障的示例，请参见以下内容 (multi-gpu-programming-models)。

问：假设存在正在进行的 nvshmem_putmem_on_stream 操作。 nvshmem_barrier_all() 是否确保完成流上挂起的 NVSHMEM 操作？

答：nvshmem_barrier_all() 操作不确保完成流上挂起的 NVSHMEM 操作。应在调用 nvshmem_barrier_all 之前调用 cudaStreamSynchronize 函数。

问：为什么 syncneighborhood_kernel 中需要 nvshmem_quiet？

答：nvshmem_barrier 语义要求这样做。正如 multi-gpu-programming-models 中所述，“nvshmem_barrier 确保团队中任何 PE 在默认上下文中完成的所有先前发出的存储和远程内存更新（包括 AMO 和 RMA 操作）在返回之前完成。”

问：如果内核使用 nvshmem_put_block 而不是 nvshmem_p，是否仍然需要 nvshmem_quiet？

答：这是 OpenSHMEM 对 put 语义的要求，该语义不保证将数据传递到远程 PE 上的目标数组。有关更多信息，请参阅 multi-gpu-programming-models。

问：我在我希望按顺序交付到目标端的同一 CUDA 流上使用主机端阻塞 API nvshmem_putmem_on_stream。即使没有非阻塞调用并且它们在单独的内核中发出，是否也需要 nvshmem_quiet？

答：在当前实现中，nvshmem_putmem_on_stream 包括 quiet。但是，根据 OpenSHMEM 规范，它仅需要释放本地缓冲区，而不一定需要交付到目标端。

问：在上述情况下，如果目标是同一个 PE，则使用 nvshmem_fence（而不是 nvshmem_quiet）是否足够？

答：在当前实现中，发送到同一 PE 的所有消息都按照 HCA 接收它们的顺序交付，这遵循流顺序。因此，甚至不需要 nvshmem_fence。但是，这些不是 OpenSHMEM 规范提供的语义。同一 CUDA 流上的 nvshmem_putmem_on_stream 函数仅确保传输的本地缓冲区将按相同顺序释放。

问：当在设备内核中使用 nvshmem_quiet 时，quiet 操作是否限定在内核运行的流中？换句话说，它是否确保完成所有操作，还是仅确保完成发送到同一流的操作？

答：它确保完成所有 GPU 发起的操作。设备上的 nvshmem_quiet 调用不会 quiet 来自主机的正在进行的操作。

调试常见问题解答¶

注意：确保您遵循 CUDA 最佳实践指南以简化 CUDA 程序的调试。例如，阅读错误处理。

问：有什么诊断挂起的提示吗？

答：检查应用程序中是否存在流 0 阻塞 CUDA 调用，例如 cudaDeviceSynchronize 或 cudaMemcpy，尤其是在应用程序的迭代阶段。初始化和最终化阶段中的流 0 阻塞调用通常是安全的。检查用于 NVSHMEM _on_stream 调用的用户流的优先级是否使用 cudaStreamCreateWithPriority 显式设置。检查挂起的确定性是否随单节点（仅通过 NVLink 或 PCI-E 连接的所有 GPU 对）与单节点（通过 InfiniBand 回环连接的不同插槽上的 GPU）或多节点（通过 InfiniBand 连接的 GPU）相比而变化。

问：如何转储调试信息？

答：请参阅运行时环境变量：NVSHMEM_INFO、NVSHMEM_DEBUG 和 NVSHMEM_DEBUG_FILE，位于环境变量。

问：即使在使用标志同步后，为什么接收缓冲区仍未用远程数据更新？

答：对于使用标志的同步，应用程序必须使用 nvshmem_wait_until 或 nvshmem_test API。仅使用 while 循环或 if 条件来检查标志值是不够的。 NVSHMEM 需要执行一致性操作，以确保在使用标志值同步后，数据对 GPU 可见。

其他常见问题解答¶

问：指针算术是否适用于 shmem 指针？例如，

int* outmsg = (int *) shmem_malloc(2* sizeof(int));
shmem_int_p(target + 1, mype, peer);

答：是的。

问：我可以避免使用 cudaDeviceSynchronize + MPI_Barrier 来跨多个 GPU 同步吗？

答：是的，可以从主机线程调用带有 cudaStreamSynchronize 的 nvshmem_barrier_all_on_stream。如果在与主机线程同步之前可能发生多个屏障同步事件，则可以提高性能。如果在屏障同步事件之后，同一个 CUDA 内核可以执行其他操作，则可以从 CUDA 内核内部调用 nvshmem_barrier_all 以进行集体同步。为了同步某些 PE 对而不是全部，可以使用发起者的成对 nvshmem_atomic_set 调用和目标的 nvshmem_wait_until 或 nvshmem_test 调用。

问：我应该如何为 NVSHMEM 分配内存？

答：有关分配对称内存的信息，请参阅内存管理。请注意，对于与通过 InfiniBand 连接的远程对等方进行通信，NVSHMEM 要求本地和远程指针都必须是对称的。如果远程对等方是 P2P 可访问的（PCI-E 或 NVLink），则可以使用 cudaMalloc 获取本地指针，并且不需要来自对称堆。

问：是否有使用 NVSHMEM 编写的微型应用程序示例？

答：是的。 multi-GPU programming models GitHub 存储库包含一个使用 NVSHMEM 编写的 Jacobi 微型应用程序示例。

问：在使用 NCCL 支持运行 NVSHMEM 时，我观察到性能下降。如何修复性能？

答：NVSHMEM 和 NCCL 都启动 CPU 代理线程以通过 IB 进行通信。这些线程可能会相互干扰，导致上下文切换，从而导致性能不佳。这些线程应具有专用的硬件线程（或核心）以获得更好的性能。仅当 NVSHMEM 使用 NCCL 进行集合通信时，NCCL 代理线程才处于活动状态。尝试绑定选项（如 –bind-to numa 或 –bind-to core，–bind-to none）以提高性能。

问：如何检查在 NVSHMEM 构建和应用程序运行时使用了哪些环境/配置变量？

答：NVSHMEM 构建包含一个二进制 nvshmem-info 实用程序，可以运行 -a 来提供此信息。也可以运行此二进制文件 -h 选项，以提供有关如何运行此二进制文件的帮助或用法。

./build/src/bin/nvshmem-info -h
Print information about NVSHMEM
Usage: nvshmem-info [options]
Options:
    -h  This help message
    -a  Print all output
    -n  Print version number
    -b  Print build information
    -e  Print environment variables
    -d  Include hidden environment variables in output
    -r  RST format environment variable output