The CUDA Handbook ▸ Chapter 11. Streaming Workloads

11.3 Streams

For workloads that benefit from concurrent memcpy and kernel execution (GPU/GPU overlap), CUDA streams can be used to coordinate execution. The stream3Streams.cu application splits the input and output arrays into k streams, then invokes k host→device memcpy’s, kernels, and device→host memcpy’s each in their own stream. Associating the transfers and computations with different streams lets CUDA know that the computations are completely independent, and CUDA will exploit whatever parallelism opportunities the hardware can support. On GPUs with multiple copy engines, the GPU may be transferring data both to and from device memory while processing other data with the SMs.

Listing 11-6 shows an excerpt from stream3Streams.cu, the same portion of the application as shown in Listing 11-5. On the test system, the output from this application reads as follows:

Measuring times with 128M floats
Testing with default max of 8 streams (set with --maxStreams <count>)

Streams Time (ms)   MB/s
1   65.31 ms    23035.75
2   54.14 ms    27373.68
3   50.60 ms    29106.94
4   48.66 ms    30099.54
5   47.85 ms    30608.95
6   47.12 ms    31053.01
7   46.59 ms    31372.53
8   46.32 ms    31537.79

The RTX 3060 has two copy engines, so it can perform host→device and device→host transfers simultaneously, and splitting the arrays across more streams steadily improves performance—from 65 ms with a single stream to 46 ms with eight. The gains taper off with each additional stream because the first kernel launch cannot begin processing until the first host→device memcpy is done, and the final device→host memcpy cannot begin until the last kernel launch is done. If the kernel processing took more time, this “overhang” would be more pronounced. Pushing beyond the eight streams shown here confirms where the benefit bottoms out: the wall clock time inches down to about 45 ms around two dozen streams, then creeps back up as the per-stream overhead of the additional memcpy and launch calls overtakes the shrinking overhang.

For our application, the best wall clock time of 46 ms shows that most of the kernel processing (4.88 ms) has been hidden.

Note that in this formulation, partly due to hardware limitations, we are not trying to insert any cudaEventRecord() calls between operations, as we did in Listing 11-5. On most CUDA hardware, trying to record events between the streamed operations of Listing 11-6 would break concurrency and reduce performance. Instead, we bracket the operations with one cudaEventRecord() before and one cudaEventRecord() after.

for ( int iStream = 0; iStream < nStreams; iStream++ ) {    cuda(MemcpyAsync(                       dptrX+iStream*streamStep,                       hptrX+iStream*streamStep,                       streamStep*sizeof(float),                       cudaMemcpyHostToDevice,                       streams[iStream] ) );    cuda(MemcpyAsync(                       dptrY+iStream*streamStep,                       hptrY+iStream*streamStep,                       streamStep*sizeof(float),                       cudaMemcpyHostToDevice,                       streams[iStream] ) );} for ( int iStream = 0; iStream < nStreams; iStream++ ) {    saxpyGPU<<<nBlocks, nThreads, 0, streams[iStream]>>>(         dptrOut+iStream*streamStep,         dptrX+iStream*streamStep,         dptrY+iStream*streamStep,         streamStep,         alpha );} for ( int iStream = 0; iStream < nStreams; iStream++ ) {    cuda(MemcpyAsync(                       hptrOut+iStream*streamStep,                       dptrOut+iStream*streamStep,                       streamStep*sizeof(float),                       cudaMemcpyDeviceToHost,                       streams[iStream] ) );}