CUDA NVSHMEM 互操作性#

本节介绍在使用 NVSHMEM 运行时开发应用程序时，一些关键的 CUDA 和 NVSHMEM API 互操作性注意事项。

使用 CUDA 流 API#

根据 CUDA 工具包的建议，鼓励用户使用异步 API cudaMemcpyAsync、cudaMemsetAsync，而不是非异步版本的 cudaMemcpy 和 cudaMemset，因为应用程序可能会受到这些 API 的非异步版本所观察到的细微同步行为的影响。 对于 NVSHMEM 应用程序，以下用法可能会导致不正确的行为

cudaMemcpy() // with target as device memory
app_kernel<<<>>>(); // Kernels Uses NVSHMEM API on target of previous cudaMemcpy() and can access stale data

以下是正确的用法

cudaMemcpyAsync(..., stream);
cudaStreamSynchronize(stream);
app_kernel<<<>>>();

NVSHMEM 设置了 CU_POINTER_ATTRIBUTE_SYNC_MEMOPS 属性，该属性自动同步对称堆上的同步 CUDA 内存操作。 因此，应用程序不需要调用 cudaDeviceSynchronize()。 从 CUDA 11.3 开始，NVSHMEM 将对对称堆使用 CUDA VMM API。 在 CUDA 12.1 和 NVSHMEM 2.10.1 中，为使用 CUDA VMM API 创建的对称堆添加了同步内存操作支持。

当 NVSHMEM 使用 CUDA VMM 且 CUDA 版本早于 12.1 时，为了与使用更高版本 CUDA（自动同步 CUDA 内存操作）的应用程序实现行为对等，应用程序需要显式使用 cudaDeviceSynchronize()。 此外，用户在使用这些异步操作和 GPUDirect 异步数据传输时必须小心。 如果不使用类似的设备同步或 barrier 操作，则此过程可能会导致竞争条件。

在本页中