The CUDA Handbook ▸ Chapter 9. Scaling and Data Movement

9.4 Inter-GPU Synchronization

CUDA events may be used for inter-GPU synchronization using cu(da)StreamWaitEvent(). If there is a producer/consumer relationship between two GPUs, the application can have the producer GPU record an event, then have the consumer GPU insert a stream-wait on that event into its command stream. When the consumer GPU encounters the stream wait, it will stop processing commands until the producer GPU has passed the point of execution where cu(da)EventRecord() was called.

The device runtime, described in Section 7.4, does not enable any inter-GPU synchronization whatsoever. That limitation may be relaxed in a future release.

Listing 9-1 gives chMemcpyPeerToPeer()¹, an implementation of peer-to-peer memcpy that uses portable memory and inter-GPU synchronization to implement the same type of memcpy that CUDA uses under the covers, if no direct mapping between the GPUs exists. The function works similarly to the chMemcpyHtoD() function of Listing 6-2 that performs host→device memcpy: a staging buffer is allocated in host memory and the memcpy begins by having the source GPU copy source data into the staging buffer and recording an event. But unlike the host→device memcpy, there is never any need for the CPU to synchronize, because all synchronization is done by the GPUs. Because both the memcpy and the event-record are asynchronous, immediately after kicking off the initial memcpy and event-record, the CPU can request that the destination GPU wait on that event and kick off a memcpy of the same buffer. Two staging buffers and two CUDA events are needed, so that the two GPUs can copy to and from staging buffers concurrently, much as the CPU and GPU concurrently operate on staging buffers during the host→device memcpy. The CPU loops over the input buffer and output buffers, issuing memcpy and event-record commands and ping-ponging between staging buffers, until it has requested copies for all bytes and all that’s left to do is wait for both GPUs to finish processing.

As with the implementations in the CUDA support provided by NVIDIA, our peer-to-peer memcpy is synchronous.

cudaError_tchMemcpyPeerToPeer(     void *_dst, int dstDevice,     const void *_src, int srcDevice,     size_t N ) {    cudaError_t status;    char *dst = (char *) _dst;    const char *src = (const char *) _src;    int stagingIndex = 0;    while ( N ) {        size_t thisCopySize = min( N, STAGING_BUFFER_SIZE );         cudaSetDevice( srcDevice );        cudaStreamWaitEvent( 0, g_events[dstDevice][stagingIndex], 0 );        cudaMemcpyAsync( g_hostBuffers[stagingIndex], src, thisCopySize,             cudaMemcpyDeviceToHost, NULL );        cudaEventRecord( g_events[srcDevice][stagingIndex] );         cudaSetDevice( dstDevice );        cudaStreamWaitEvent( 0, g_events[srcDevice][stagingIndex], 0 );        cudaMemcpyAsync( dst, g_hostBuffers[stagingIndex], thisCopySize,             cudaMemcpyHostToDevice, NULL );        cudaEventRecord( g_events[dstDevice][stagingIndex] );         dst += thisCopySize;        src += thisCopySize;        N -= thisCopySize;        stagingIndex = 1 - stagingIndex;    }    // Wait until both devices are done    cudaSetDevice( srcDevice );    cudaDeviceSynchronize();     cudaSetDevice( dstDevice );    cudaDeviceSynchronize(); Error:    return status;}