NVSHMEM官方文档部分内容总结
Using NVSHMEM
1、NVSHMEM 简介
- 针对 CUDA 设备实现的基于 OpenSHMEM 的规范（PGAS 并行编程范式，partitioned global address space）
- OpenSHMEM 和 NVSHMEM 的区别
- NVSHMEM 通信过程中所有的 buffer arguments 必须是对称的
- Block 中线程进行 fetch 操作时只能保证数据的 weak ordering（如需强制顺序则要显式调用 nvshmem_fence，其会保证在 fence 之前的所有操作以保证内存一致性）
- 支持单边通信的方式：通过 Verbs API 或者 UCX 进行 IB 或 RoCE 传输
2、NVSHMEM 优点
- 支持多 GPU 间直接进行 data movement 或 synchronization，而无需考虑和 CPU 之间的互动之类的操作带来的 overhead
- 支持通过 composite kernel 实现 comp 和 comm 的 overlap
3、MPI 的问题
- 可能会引起 shared data 读写的 high locking/atomics overheads
- Message ordering 会引起 serialization overheads
- 接收方在发布描述符之前收到数据会引起 protocol overheads
4、NVSHMEM 对应的 addressing model（类似 PGAS）
- 一个 NVSHMEM job 中的所有 PE 必须同时被初始化，在 job 退出前也需要同时 finalize
- 通过 CPU-side NVSHMEM allocation API 可以在 GPU 上分配（必须是）大小相同的 shared symmetric memory（而其他方式开辟的内存则属于 private memory）
- 每个 PE 上的 symmetric memory 对其他 PE 均可见且可操作（通过 NVSHMEM API）
- 通过 NVSHMEM 分配 API（如 nvshmem_malloc ）返回的对称内存地址对于调用该 API 的 PE 对应的 GPU 来说，是一个可以通过 CUDA API 或者 Load/Store 操作直接访问的有效内存地址。因此如果仅需要对本地 PE 的对称内存进行操作那么可以可以利用 CUDA 原生操作和优化避免 NVSHMEM API 的额外开销
- 通过<symmetric_address, destination_PE>的组合可以获取对应的 symmetric objects，在 NVSHMEM runtime 中 symmetric_address 会被转换为真正的远端地址
5、simple shift 代码和编译
- 代码
#include <stdio.h>
#include <cuda.h>
#include <nvshmem.h>
#include <nvshmemx.h>
__global__ void simple_shift(int *destination) {
// 返回调用进程在全局NVSHMEM作业中的PE编号,不需要指定任何团队参数
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();
// 获取节点内PE编号,节点内编号通常从0开始因此适合用作设备ID
mype_node = nvshmem_team_my_pe(NVSHMEMX_TEAM_NODE);
cudaSetDevice(mype_node);
cudaStreamCreate(&stream);
// 每个PE上都会开辟一块叫做destination的symmetric memory
int *destination = (int *) nvshmem_malloc(sizeof(int));
simple_shift<<<1, 1, 0, stream>>>(destination);
// 在指定的CUDA流上执行全局屏障同步
// 确保所有PE在继续执行前都已到达屏障点即所有通信操作都已完成
// nvshmem_barrier_all(): 在当前CUDA上下文中同步所有PE但不与特定流关联
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;
}- 编译
6、Communication Model
- put/get：可以在 symmetric object 之间来回读写数据
- AMO：Atomic Memory Operations，特点有
- 原子性保证 ：操作在一个不可分割的步骤中完成，避免了竞态条件
- 远程执行 ：可以对远程 PE 的内存执行原子操作
- 多种操作类型 ：支持多种原子操作，如加、与、或、异或、交换
- 开销：原子操作通常比普通内存操作慢；远程原子操作比本地原子操作开销更大
- 与 CUDA 原子操作的区别：可以跨节点执行；提供了更丰富的原子操作类型；可以同时在设备和主机代码中被调用
- nvshmem_ptr：适用范围为全局 PE。用于获取远程 PE 上对称内存地址的本地映射指针
- 将远程 PE 上的对称内存地址转换为本地可访问的指针
- 使得本地 PE 可以直接通过指针对远程 PE 的 global memory 进行 load 或者 store 操作（而无需使用显式的通信操作）
- nvshmemx_mc_ptr：适用范围为单个 team，故需要需要指定一个 team 对象
- 调用后返回的是指向团队内所有 PE 的对称内存的多连接指针（multicast pointer）
- 适用于团队内的集体通信模式（支持 multimem load reduce and store broadcast to global multicast memory）
7、相较于 OpenSHMEM 的优势
- 所有通过 NVSHMEM API 分配的 symmetric memory 均为 pinned GPU device memory
- NVSHMEM 同时支持 GPU 侧发起或者 CPU 侧翻起的通信和同步，而 OpenSHMEM 中这类操作只能从 CPU 侧发起
8、NVSHMEM 应用必须保证以下几点
- 在第一次分配内存、同步、通信、启动 kernel API 等操作前必须保证 PE 已确定要使用的 GPU（例如通过 cudaSetDevice）
- 在使用任何设备端 NVSHMEM API 之前，必须在主机端完成 NVSHMEM 内存分配或同步操作，即只有在主机端分配对称内存后才能在设备端使用 NVSHMEM API
- 每个 PE 在整个 NVSHMEM 任务生命周期内只能使用一个固定的 GPU
- 一般一个 GPU 不能被多个 PE 共享使用
9、CUDA kernel 在使用 NVSHMEM 通信 API 时的两种启动方式及其限制条件
- Normal launch
- 启动方式：使用常规 <<<...>>> 语法或 cudaLaunchKernel 这类 API
- 特点：当 kernel 不使用 NVSHMEM synchronization API（如 nvshmem_quiet）或 collective API（如 nvshmem_broadcast）时可以使用该方式启动，且 kernel 中仍可以使用其他 NVSHMEM 设备端 API（如单边通信 API nvshmem_put_nbi）
- Collective launch
- 启动方式：使用例如 nvshmemx_collective_launch 等 NVSHMEM 特有的接口
- 特点：如果 kernel 使用了 NVSHMEM 的 synchronization/collective API 则必须使用 collective launch API 启动，否则会出现未定义行为
10、NVSHMEM 提供的多种对称内存访问方法
- PUT：将本地数据写入到远程 PE 的对称内存
- GET：从远程 PE 的对称内存读取数据到本地
- AMO：在远程内存上执行原子操作，如原子加、比较交换等，保证在并发环境下的数据一致性
- Signal：用于 PE 之间的同步和通知，该过程不涉及完整的数据传输
- Load/Store：通过 nvshmem_ptr 获取指向对称内存的指针后可以直接使用指针进行读写（类似普通内存操作），但仅限于同一节点内的访问
- Direct multimem load reduce and store broadcast：使用 nvshmemx_mc_ptr 返回的指针可以支持多内存操作，如归约和广播
- Collective functions：涉及多个 PE 的协同工作时需要用到的广播、归约等操作
- 等待和测试函数：允许一个 PE 等待本地对称内存中的值满足特定条件后再继续执行。这类函数在实现同步屏障、事件通知和条件执行时非常有用
11、nvshmem_fence 和 nvshmem_quiet 的区别
- nvshmem_fence：用于强制执行点对点通信的顺序。确保调用该函数之前的所有更新操作在调用之后的更新操作之前对特定目标 PE 可见
- nvshmem_quiet：用于全局同步。该函数确保调用之前 PE 发起的所有更新操作对所有 PE 可见
- 使用 nvshmem_fence 的场景：只需确保特定目标 PE 上的更新顺序，例如 PE0 先写入数据 A 再写入标志 B 到 PE1，需要确保 PE1 看到的顺序是先 A 后 B 即可
- 使用 nvshmem_quiet 的场景：确保更新对所有 PE 可见，例如 collective communication 就需要确保所有 PE 都看到最新的数据状态
12、NVSHMEM 内存更新
- 最终完成保证（Eventually Complete）：使用 NVSHMEM API（如 nvshmem_put 、 nvshmem_get 等）进行内存更新时，这些操作最终会完成，不需要源 PE 或目标 PE 执行额外的操作来确保完成
- 最终可见性（Eventually Visible）：一个 PE 对共享内存的更新最终会被其他 PE 通过 NVSHMEM API 看到，不需要显式的同步操作来确保可见性（当需要确保特定顺序或立即可见性时，可以使用 nvshmem_fence 或 nvshmem_quiet 等同步操作）
- 更新稳定性（Stability of Updates）：一旦更新对其他 PE 可见，它会保持可见直到被另一个更新替换
- 有限时间完成（Finite Time Completion）：nvshmem_fence 、 nvshmem_quiet 等同步操作会在有限时间内完成，不会出现无限等待或死锁情况
13、MPG（multiple PEs per GPU）模式：支持每个 GPU 绑定多个 PE
- CUDA MPS：允许多个 CPU 进程共享同一 GPU context，提高GPU利用率和占用率
- MPG without MPS：属于时间共享（time-sharing）模式，多个 PE 通过时间共享方式使用同一个 GPU
- 每个 PE 有自己的 CUDA 上下文
- GPU 需要在不同 PE 的 CUDA 任务之间切换上下文
- 多个 PE 不能同时在同一 GPU 上运行
- 不支持点对点同步 API 和集体通信 API
- MPG with MPS：允许多个 PE 各自的 CUDA 上下文同时在同一 GPU 上运行。当所有 PE 的 active thread 百分比总和小于 100% 时支持所有 NVSHMEM API
- MPG with MPS & oversubscription：当所有 PE 的 active thread 百分比总和超过 100% 时 CUDA 无法保证所有分配给同一 GPU 的 PE 能同时运行
评论
发表评论