CUDA NVSHMEM Interoperability

This section describes some key CUDA and NVSHMEM API interoperability considerations when developing applications using the NVSHMEM runtime.

Using CUDA Streams APIs

As recommended by the CUDA toolkit, users are encouraged to use the async APIs cudaMemcpyAsync and cudaMemsetAsync instead of the non-async versions of cudaMemcpy and cudaMemset because applications can be impacted by subtle synchronization behavior observed with the non-async versions of these APIs. For NVSHMEM applications, the following usage can lead to incorrect behavior:

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

Here is the correct usage:

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

NVSHMEM sets the CU_POINTER_ATTRIBUTE_SYNC_MEMOPS attribute, which automatically synchronizes the synchronous CUDA memory operations on the symmetric heap. As a result, the application does not need to call cudaDeviceSynchronize(). Starting with CUDA 11.3, NVSHMEM uses the CUDA VMM API for the symmetric heap. Support for synchronous memory operations was added for symmetric heaps created using the CUDA VMM API in CUDA 12.1 and NVSHMEM 2.10.1.

When NVSHMEM uses CUDA VMM and the CUDA version is earlier than 12.1, the application needs to explicitly use cudaDeviceSynchronize() to achieve behavioral parity with later CUDA releases that automatically synchronize CUDA memory operations. Additionally, users must be careful when using these async operations with GPUDirect async data transfers. If similar device synchronization or barriered operations are not used, this can lead to race conditions.