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

6.5 Concurrent Copying and Kernel Processing

Since CUDA applications must transfer data across the PCI Express bus for processing by the GPU, another performance opportunity presents itself in the form of performing those host↔︎device memory transfers concurrently with kernel processing. According to Amdahl’s Law, the maximum speedup achievable by using multiple processors is as follows:

\[Speedup = \frac{1}{r_{s} + \frac{r_{p}}{N}}\]

where \(r_{s} + r_{p} = 1\) and N is the number of processors. In the case of concurrent copying and kernel processing, the “number of processors” is the number of autonomous hardware units in the GPU: 1-2 copy engines, plus the SMs that execute the kernels. For N=2, Figure 6-6 shows the idealized speedup curve as \(r_{s}\) and \(r_{p}\) vary.

Figure 6-6. Idealized Amdahl’s Law Curve.

So in theory, a 2x performance improvement is possible on a GPU with one copy engine, but only if the program gets perfect overlap between the SMs and the copy engine, and only if the program spends equal time transferring and processing the data.

Before undertaking this endeavor, you should take a close look at whether it will benefit your application. Applications that are extremely transfer-bound (i.e., they spend most of their time transferring data to and from the GPU) or extremely compute-bound (i.e., they spend most of their time processing data on the GPU) will derive little benefit from overlapping transfer and compute.

6.5.1 concurrencyMemcpyKernel.cu

The program concurrencyMemcpyKernel.cu is designed to illustrate not only how to implement concurrent memcpy and kernel execution, but also how to determine whether it is worth doing at all.

Listing 6-3 gives AddKernel(), a “makework” kernel that has a parameter cycles to control how long it runs.

__global__ voidAddKernel( int *out, const int *in, size_t N, int addValue, int cycles ){    for ( size_t i = blockIdx.x*blockDim.x+threadIdx.x;                  i < N;                 i += blockDim.x*gridDim.x )    {        volatile int value = in[i];        for ( int j = 0; j < cycles; j++ ) {            value += addValue;        }        out[i] = value;    }}

AddKernel() streams an array of integers from in to out, looping over each input value cycles times. By varying the value of cycles, we can make the kernel range from a trivial streaming kernel that pushes the memory bandwidth limits of the machine, to a totally compute-bound kernel.

Two routines in the program measure the performance of AddKernel():

• TimeSequentialMemcpyKernel() copies the input data to the GPU, invokes AddKernel(), and copies the output back from the GPU in separate, sequential steps;

• TimeConcurrentMemcpyKernel() allocates a number of CUDA streams and performs the host→device memcpy’s, kernel processing, and device→host memcpy’s in parallel.

TimeSequentialMemcpyKernel(), given in Listing 6-4, uses four CUDA events to separately time the host→device memcpy, kernel processing, and device→host memcpy. It also reports back the total time, as measured by the CUDA events.

boolTimeSequentialMemcpyKernel(     float *timesHtoD,     float *timesKernel,     float *timesDtoH,     float *timesTotal,    size_t N,     const chShmooRange& cyclesRange,    int numBlocks ){    cudaError_t status;    bool ret = false;    int *hostIn = 0;    int *hostOut = 0;    int *deviceIn = 0;    int *deviceOut = 0;    const int numEvents = 4;    cudaEvent_t events[numEvents];     for ( int i = 0; i < numEvents; i++ ) {        events[i] = NULL;        cuda(EventCreate( &events[i] ) );    }    cudaMallocHost( &hostIn, N*sizeof(int) );    cudaMallocHost( &hostOut, N*sizeof(int) );    cudaMalloc( &deviceIn, N*sizeof(int) );    cudaMalloc( &deviceOut, N*sizeof(int) );     for ( size_t i = 0; i < N; i++ ) {        hostIn[i] = rand();    }     cudaDeviceSynchronize();     for ( chShmooIterator cycles(cyclesRange); cycles; cycles++ ) {         printf( "." ); fflush( stdout );         cudaEventRecord( events[0], NULL );        cudaMemcpyAsync( deviceIn, hostIn, N*sizeof(int),             cudaMemcpyHostToDevice, NULL );        cudaEventRecord( events[1], NULL );        AddKernel<<<numBlocks, 256>>>(             deviceOut, deviceIn, N, 0xcc, *cycles );        cudaEventRecord( events[2], NULL );        cudaMemcpyAsync( hostOut, deviceOut, N*sizeof(int),             cudaMemcpyDeviceToHost, NULL );        cudaEventRecord( events[3], NULL );         cudaDeviceSynchronize();         cudaEventElapsedTime( timesHtoD, events[0], events[1] );        cudaEventElapsedTime( timesKernel, events[1], events[2] );        cudaEventElapsedTime( timesDtoH, events[2], events[3] );        cudaEventElapsedTime( timesTotal, events[0], events[3] );         timesHtoD += 1;        timesKernel += 1;        timesDtoH += 1;        timesTotal += 1;    }     ret = true; Error:    for ( int i = 0; i < numEvents; i++ ) {        cudaEventDestroy( events[i] );    }    cudaFree( deviceIn );    cudaFree( deviceOut );    cudaFreeHost( hostOut );    cudaFreeHost( hostIn );    return ret;}

The cyclesRange parameter, which uses the “shmoo” functionality described in Section A.4, specifies the range of cycles values to use when invoking AddKernel(). On a g2.2xlarge instance in EC2, the times (in ms) for cycles values from 4..64 are as follows:

For values of *cycles around 48 (highlighted), where the kernel takes about the same amount of time as the memcpy operations, we presume there would be a benefit in performing the operations concurrently.

The routine TimeConcurrentMemcpyKernel() divides the computation performed by AddKernel() evenly into segments of size streamIncrement, and uses a separate CUDA stream to compute each. The code fragment of Listing 6-5, from TimeConcurrentMemcpyKernel(), highlights the complexity of programming with streams.

    intsLeft = N;    for ( int stream = 0; stream < numStreams; stream++ ) {        size_t intsToDo = (intsLeft < intsPerStream) ?             intsLeft : intsPerStream;        cuda(MemcpyAsync(             deviceIn+stream*intsPerStream,             hostIn+stream*intsPerStream,             intsToDo*sizeof(int),             cudaMemcpyHostToDevice, streams[stream] ) );        intsLeft -= intsToDo;    }     intsLeft = N;    for ( int stream = 0; stream < numStreams; stream++ ) {        size_t intsToDo = (intsLeft < intsPerStream) ?             intsLeft : intsPerStream;        AddKernel<<<numBlocks, 256, 0, streams[stream]>>>(             deviceOut+stream*intsPerStream,             deviceIn+stream*intsPerStream,             intsToDo, 0xcc, *cycles );        intsLeft -= intsToDo;    }     intsLeft = N;    for ( int stream = 0; stream < numStreams; stream++ ) {        size_t intsToDo = (intsLeft < intsPerStream) ?             intsLeft : intsPerStream;        cuda(MemcpyAsync(             hostOut+stream*intsPerStream,             deviceOut+stream*intsPerStream,             intsToDo*sizeof(int),             cudaMemcpyDeviceToHost, streams[stream] ) );        intsLeft -= intsToDo;    }

Besides requiring the application to create and destroy CUDA streams, the streams must be looped over separately for each of the host→device memcpy, kernel processing, and device→host memcpy operations; without this “software-pipelining,” there would be no concurrent execution of the different streams’ work, as each streamed operation is preceded by an “interlock” operation that prevents the operation from proceeding until the previous operation in that stream has completed. The result would be not only a failure to get parallel execution between the engines, but an additional performance degradation due to the slight overhead of managing stream concurrency.

The computation cannot be made fully concurrent, since no kernel processing can be overlapped with the first or last memcpy’s; and there is some overhead in synchronizing between CUDA streams and, as we saw in the previous section, in invoking the memcpy and kernel operations themselves. As a result, the optimal number of streams depends on the application and should be determined empirically. In my experience, values between 4-8 work well for streaming computations. The concurrencyMemcpyKernel.cu program enables the number of streams to be specified on the command line using the --numStreams parameter.

6.5.2 Performance Results

The concurrencyMemcpyKernel.cu program generates a report on performance characteristics over a variety of cycles values, with a fixed buffer size and number of streams. On a g2.8xlarge instance in Amazon EC2, with a buffer size of 128M integers and 8 streams, the report is as follows for cycles values from 4..64:

The full graph for cycles values from 4..256 is given in Figure 6-7. Unfortunately, for these settings, the 50% speedup shown here falls well short of the 3x speedup that theoretically could be obtained.

Figure 6-7. Speedup Due To Memcpy/Kernel Concurrency (Tesla M2050, with community results from modern GPUs)

The benefit on a GeForce GTX 280, which contains only one copy engine, is more pronounced. Here, the results from varying cycles up to 512 are shown. The maximum speedup, shown in Figure 6-8, is much closer to the theoretical maximum of 2x.

Figure 6-8. Speedup Due To Memcpy/Kernel Concurrency (GeForce GTX 280, with community results from modern GPUs)

The community curves overlaid on Figures 6-7 and 6-8 show how the calculus has shifted. A GeForce RTX 3060, with two copy engines and PCI Express 4.0, transfers the data so quickly that the make-work kernel stays well under the memcpy time across the book’s cycle range; even so, overlapping the copies with compute yields a speedup that climbs from about 1.6× (pure streaming) to roughly 2.2× as the kernel is given more work.

As written, concurrencyMemcpyKernel.cu serves little more than an illustrative purpose, because AddKernel() is just make-work. But you can plug your own kernel(s) into this application to help determine whether the additional complexity of using streams is justified by the performance improvement. Note that unless concurrent kernel execution is desired (see next section), lines 9-14 of Listing 6-5 could include successive kernel invocations in the same stream, and the application will still get the desired concurrency.

As a note, the number of copy engines can be queried by calling cudaGetDeviceProperties() and examining cudaDeviceProp::asyncEngineCount, or calling cuDeviceGetAttribute() with CU_DEVICE_ATTRIBUTE_ASYNC_ENGINE_COUNT.

The Copy Engines accompanying early CUDA hardware could copy linear memory only, but modern Copy Engines offer full memcpy support, including 2D and 3D CUDA arrays.

6.5.3 Breaking Inter-Engine Concurrency

Using CUDA streams for concurrent memcpy and kernel execution introduces many more opportunities to “break concurrency.” In the previous section, CPU/GPU concurrency could be broken by unintentionally doing something that caused CUDA to perform a full CPU/GPU synchronization. Here, CPU/GPU concurrency can be broken by unintentionally performing an unstreamed CUDA operation. Recall that the NULL stream performs a “join” on all GPU engines; so even an asynchronous memcpy operation will stall inter-engine concurrency, if the NULL stream is specified.

Besides specifying the NULL stream explicitly, the main avenue for these unintentional “concurrency breaks” is calling functions that run in the NULL stream implicitly because they do not take a stream parameter. When streams were first introduced, functions such as cudaMemset() and cuMemcpyDtoD(), and the interfaces for libraries such as CUFFT and CUBLAS, did not have any way for applications to specify stream parameters.

The CUDA Visual Profiler will call out concurrency breaks in its reporting.