The CUDA Handbook ▸ Chapter 6. Streams, Events, and Graphs

6.3 CUDA Events: CPU/GPU Synchronization

CUDA events enable “partial” CPU/GPU synchronization: instead of full CPU/GPU synchronization where the CPU waits until the GPU is idle, introducing a bubble into the GPU’s work pipeline, CUDA events may be recorded into the asynchronous stream of GPU commands. The CPU then can wait until the work preceding the event has been done. The GPU may continue doing whatever work was submitted after the cuEventRecord()/cudaEventRecord().

As an example of CPU/GPU concurrency, let’s implement a memcpy routine for pageable memory. The code for this program implements the algorithm described in Figure 6-3, and is located in pageableMemcpyHtoD.cu. It uses two pinned memory buffers, stored in global variables declared as follows:

void *g_hostBuffers[2];

and two CUDA events, declared as:

cudaEvent_t g_events[2];

voidchMemcpyHtoD( void *device, const void *host, size_t N ) {    cudaError_t status;    char *dst = (char *) device;    const char *src = (const char *) host;    int stagingIndex = 0;    while ( N ) {        size_t thisCopySize = min( N, STAGING_BUFFER_SIZE );         cuda(EventSynchronize( g_events[stagingIndex] ) );        memcpy( g_hostBuffers[stagingIndex], src, thisCopySize );         cuda(MemcpyAsync( dst, g_hostBuffers[stagingIndex],            thisCopySize, cudaMemcpyHostToDevice, NULL ) );        cuda(EventRecord( g_events[1-stagingIndex], NULL ) );        dst += thisCopySize;        src += thisCopySize;        N -= thisCopySize;        stagingIndex = 1 - stagingIndex;    }Error:    return;}

chMemcpyHtoD() is designed to maximize CPU/GPU concurrency, by “ping-ponging” between the two host buffers: the CPU copies into one buffer while the GPU pulls from the other. There is some “overhang” where no CPU/GPU concurrency is possible at the beginning and end of the operation, when the CPU is copying the first and last buffers, respectively.

In this program, the only synchronization needed – the cudaEventSynchronize() of Line 11 – ensures that the GPU has finished with a buffer before starting to copy into it. cudaMemcpyAsync() returns as soon as the GPU commands have been submitted: it does not wait until the operation is complete. The cudaEventRecord() is also asynchronous – it causes the event to be signaled when the just-requested asynchronous memcpy has been completed.

The CUDA events are recorded immediately after creation, so that the first cudaEventSynchronize() calls of line 11 work correctly:

    cuda(EventCreate( &g_events[0] ) );
    cuda(EventCreate( &g_events[1] ) );
    // record events so they are signaled on first synchronize
    cuda(EventRecord( g_events[0], 0 ) );
    cuda(EventRecord( g_events[1], 0 ) );

If you run pageableMemcpyHtoD.cu, it will report a bandwidth number much smaller than the pageable memcpy bandwidth delivered by the CUDA driver. That’s because the C runtime’s memcpy() implementation is not optimized to move memory as fast as the CPU can. For best performance, the memory must be copied using SSE instructions that can move data 16 bytes at a time. Writing a general-purpose memcpy using these instructions is complicated by their alignment restrictions, but a simple version that requires the source, destination, and byte count to be 16-byte aligned is not difficult⁴:

#include <xmmintrin.h>
bool
memcpy16( void *_dst, const void *_src, size_t N )
{
    if ( N & 0xf ) {
        return false;
    }
    float *dst = (float *) _dst;
    const float *src = (const float *) _src;
    while ( N ) {
        _mm_store_ps( dst, _mm_load_ps( src ) );
        src += 4;
        dst += 4;
        N -= 16;
    }
    return true;
}

When the C runtime memcpy() is replaced by this one, performance on an Amazon EC2 cg1.4xlarge instance increases from 2155MB/s to 3267MB/s. More-complicated memcpy routines can deal with relaxed alignment constraints, and slightly higher performance is possible by unrolling the inner loop. On g2.2xlarge, the CUDA driver’s more-optimized SSE memcpy achieves about 100MB/s higher performance than pageableMemcpyHtoD16.cu. On a modern PCI Express 4.0 system (a GeForce RTX 3060 hosted by a Ryzen 7 7700X), the same staged approach sustains around 14GB/s – roughly half the pinned-memory bandwidth, since staging through a pinned buffer still costs one CPU-side copy.

How important is the CPU/GPU concurrency for performance of pageable memcpy? If we move the event synchronization, we can make the host→device memcpy synchronous:

    while ( N ) {
        size_t thisCopySize = min( N, STAGING_BUFFER_SIZE );
< cuda(EventSynchronize( g_events[stagingIndex] ) );
        memcpy( g_hostBuffers[stagingIndex], src, thisCopySize );
        cuda(MemcpyAsync( dst, g_hostBuffers[stagingIndex],
            thisCopySize, cudaMemcpyHostToDevice, NULL ) );
        cuda(EventRecord( g_events[1-stagingIndex], NULL ) );
> cuda(EventSynchronize( g_events[1-stagingIndex] ) );
        dst += thisCopySize;
        src += thisCopySize;
        N -= thisCopySize;
        stagingIndex = 1 - stagingIndex;
    }

This code is available in pageableMemcpyHtoD16Synchronous.cu, and is about 70% as fast (2334MB/s instead of 3267MB/s) on the same cg1.4xlarge instance.

6.3.1 Blocking Events

CUDA events also optionally can be made “blocking,” in which they use an interrupt-based mechanism for CPU synchronization. The CUDA driver then implements cu(da)EventSynchronize() calls using thread synchronization primitives that suspend the CPU thread instead of polling the event’s 32-bit tracking value.

For latency-sensitive applications, blocking events may impose a performance penalty. In the case of our pageable memcpy routine, using blocking events causes a slight slowdown (about 100MB/s) on our cg1.4xlarge instance. But for more GPU-intensive applications, or for applications with “mixed workloads” that need significant amounts of processing from both CPU and GPU, the benefits of having the CPU thread idle outweigh the costs of handling the interrupt that occurs when the wait is over. An example of a mixed workload is video transcoding, which features divergent bit-twiddling suitable for the CPU, and signal and pixel processing suitable for the GPU.

6.3.2 Queries

Both CUDA streams and CUDA events may be queried with cu(da)StreamQuery() and cu(da)EventQuery(), respectively. If cu(da)StreamQuery() returns success, all of the operations pending in a given stream have been completed; if cu(da)EventQuery() returns success, the event has been recorded.

Although these queries are intended to be lightweight, if ECC is enabled, they do perform kernel thunks to check the current error status of the GPU. Additionally, on Windows, any pending commands will be submitted to the GPU, which also requires a kernel thunk.