Optimizing Asynchronous Operations with CUDA Streams and Events

CUDA Streams: Pipelines for GPU Work

A CUDA stream is a sequence of operations that execute on the device in the order they are issued by the host. The fundamental rule is:

• Operations within a single stream are guaranteed to execute sequentially.
• Operations in different streams have no defined order and can execute concurrently if their resource requirements don't conflict.

By default, all CUDA operations are issued to a special default stream (also called the "null stream"). This stream has a unique, synchronizing behavior: an operation in the default stream will only begin after all preceding operations in any stream on the device have completed, and no subsequent operation in any stream will begin until the default stream operation finishes. This makes the default stream unsuitable for achieving concurrent execution between multiple GPU tasks.

To enable concurrency, we must explicitly create and use our own non-default streams.

Implementing with Streams

To use non-default streams, you first declare and create them in your host code. Then, you use special versions of CUDA API calls that accept a stream identifier.

The key functions are:

• cudaStream_t stream;: Declares a stream handle.
• cudaStreamCreate(&stream);: Creates a non-default stream.
• kernel_name<<<grid, block, 0, stream>>>(...);: Launches a kernel into the specified stream. The stream argument is the fourth execution configuration parameter.
• cudaMemcpyAsync(dst, src, count, kind, stream);: Initiates an asynchronous memory copy in the specified stream.
• cudaStreamDestroy(stream);: Destroys the stream once you are finished with it.

It is crucial to use cudaMemcpyAsync. The standard, synchronous cudaMemcpy cannot be used to achieve overlap as it blocks the host.

Let's consult the NVIDIA Developer blog for a detailed walkthrough of the API.

The Prerequisite: Pinned Memory

There's a critical subtlety to using cudaMemcpyAsync. For a host-to-device or device-to-host memory transfer to be truly asynchronous and overlap with other operations, the host-side memory must be pinned (or page-locked).

Standard host memory allocated with malloc or new is pageable, meaning the operating system can move it from physical RAM to a swap file on disk at any time. The GPU's DMA (Direct Memory Access) engine cannot access data that isn't in physical RAM. If you issue an async copy from pageable memory, the CUDA driver must first implicitly create a temporary pinned buffer, copy your data there, and then initiate the transfer to the GPU. This staging process effectively serializes the copy, defeating the purpose of asynchrony.

By allocating host memory as pinned, you guarantee to the OS that this memory will always reside in physical RAM, allowing the GPU's DMA engine to access it directly without CPU intervention.

We will cover the specific APIs for allocating pinned memory in our next lesson. For now, understand that it is a mandatory prerequisite for overlapping data transfers with kernel execution.

The Overlap Pattern: Pipelining Work

With streams and pinned memory, we can implement a powerful optimization pattern. Imagine you need to process a very large array. Instead of copying the entire array, processing it, and copying it back, you can break it into smaller chunks and process them in a pipeline.

The strategy is as follows:

1. Divide the large array into N chunks.
2. Create a small number of streams (e.g., 2 to 4).
3. Loop through the chunks, assigning each to a stream in a round-robin fashion. For each chunk i in stream s:
   • Issue an asynchronous copy of chunk i from host to device (H2D) in stream s.
   • Issue the kernel launch to process chunk i in stream s.
   • Issue an asynchronous copy of the results for chunk i from device to host (D2H) in stream s.

Because operations within a stream are ordered, the kernel for chunk i will correctly wait for its data to arrive. Because operations in different streams can run concurrently, the H2D copy for chunk i+1 (in another stream) can overlap with the kernel execution for chunk i.

This overlap is visualized perfectly in the following diagram, which shows four streams working in a staggered fashion to keep the H2D copy engine, the kernel execution engine (K), and the D2H copy engine all busy simultaneously.

The CUDA Best Practices Guide provides an excellent summary and a code skeleton for this pattern.

Fine-Grained Control with CUDA Events

Streams manage sequences, but what if you need to synchronize at specific points? For example, how do you make a kernel in Stream B wait for a data transfer in Stream A to finish? Or how do you accurately time a kernel's execution without the overhead of host-side timers? For this, we use CUDA events.

An event is essentially a marker that you can place into a stream. The GPU records a timestamp when it reaches the event.

The core Event API functions are:

• cudaEvent_t start, stop;: Declares event handles.
• cudaEventCreate(&start);: Creates an event.
• cudaEventRecord(start, stream);: Records the event in the specified stream. The event is marked as "occurred" only when the GPU has processed all preceding work in that stream.
• cudaStreamWaitEvent(stream_to_wait, event_to_wait_for, 0);: Makes all future work in stream_to_wait stall until event_to_wait_for has been recorded. This is the key to creating cross-stream dependencies.
• cudaEventSynchronize(stop);: Blocks the host CPU thread until the specified event has been recorded. This is a much more targeted way to synchronize than cudaDeviceSynchronize().
• cudaEventElapsedTime(&milliseconds, start, stop);: Calculates the wall-clock time in milliseconds between the recording of a start and stop event. This is the canonical way to accurately measure GPU performance.

Case Study: Overlap Strategies

Let's return to the How to Overlap Data Transfers blog post to see how these concepts are applied in a complete example. The author demonstrates two different ways to structure the loops for the chunking pattern and analyzes the performance implications, which depend on the GPU's hardware capabilities (specifically, the number of copy engines). Given your background in computing, this analysis of how software strategy interacts with hardware architecture will be particularly relevant.

Conclusion

In this lesson, you have moved beyond simple asynchronous launching to a sophisticated model of GPU concurrency. You now have the tools to orchestrate multiple streams of work to maximize hardware utilization.

Key Takeaways:

• CUDA Streams are sequences of commands that execute in order. Operations in different non-default streams can execute concurrently.
• The default stream is a synchronizing stream and should be avoided when trying to achieve overlap between GPU tasks.
• Pinned (Page-Locked) Memory is a prerequisite for cudaMemcpyAsync to be truly non-blocking and enable overlap between data transfers and computation.
• A common and effective pattern is to chunk a large problem and process the chunks in a pipeline across multiple streams.
• CUDA Events are markers used for fine-grained synchronization (e.g., creating dependencies between streams with cudaStreamWaitEvent) and for accurate, low-overhead performance measurement on the device (cudaEventElapsedTime).

In our next lesson, we will dive deeper into GPU memory management. We will look at the explicit APIs for allocating pinned memory and introduce Unified Memory, a higher-level abstraction that simplifies memory management by providing a single memory space for both the CPU and GPU.