CPU/GPU concurrency refers to the ability of the CPU to continue processing after having sent some request to the GPU. Arguably, the most important use of CPU/GPU concurrency is to hide the overhead of requesting work from the GPU.
Kernel launches have always been asynchronous: a series of kernel launches, with no intervening CUDA operations in between, causes the CPU to submit the kernel launch to the GPU and return control to the caller before the GPU has finished processing.
We can measure the driver overhead by bracketing a series of NULL
kernel launches with timing operations. Listing 6-1 shows
nullKernelAsync.cu, a small program that measures the amount of time
needed to perform a kernel launch.
#include <stdio.h>#include__global__voidNullKernel(){}intmain( int argc, char *argv[] ){ const int cIterations = 1000000; printf( "Launches... " ); fflush( stdout ); std::chrono::steady_clock::time_point start, stop; start = std::chrono::steady_clock::now(); for ( int i = 0; i < cIterations; i++ ) { NullKernel<<<1,1>>>(); } cudaDeviceSynchronize(); stop = std::chrono::steady_clock::now(); double us = 1e6*std::chrono::duration(stop - start).count(); double usPerLaunch = us / (float) cIterations; printf( "%.2f us\n", usPerLaunch ); return 0;}nullKernelAsync.cuThe std::chrono::steady_clock calls, described in
Appendix A, use the host operating system’s high-resolution timing
facilities. The cudaDeviceSynchronize() call of line 18 is needed for
accurate timing: without it, the GPU would still be processing the last
kernel invocations when the end time is recorded with this function
call:
stop = std::chrono::steady_clock::now();
If you run this program, you will see that invoking a kernel, even a kernel that does nothing, costs on the order of a microsecond or two on current hardware (and anywhere from 2.0 to 8.0 microseconds on the systems of the book’s first edition).
Most of that time is spent in the driver. The CPU/GPU concurrency enabled by kernel launches only helps if the kernel runs for longer than it takes the driver to invoke it!
To underscore the importance of CPU/GPU concurrency for small kernel
launches, let’s move the cudaDeviceSynchronize() call into the inner
loop1:
start = std::chrono::steady_clock::now();
for ( int i = 0; i < cIterations; i++ ) {
NullKernel<<<1,1>>>();
cudaDeviceSynchronize();
}
stop = std::chrono::steady_clock::now();
The only difference here is that the CPU is waiting until the GPU has finished processing each NULL kernel launch before undertaking to launch the next kernel, as shown in Figure 6-1.
Figure 6-1. CPU/GPU concurrency
Note that cIterations is set to 10,000 instead of
1,000,000 because it takes so much longer to run! As an example, on a
GeForce RTX 3060, nullKernelAsync reports about 1.2 μs per
launch and nullKernelSync about 3.9 μs. (On the
ECC-disabled Amazon EC2 instance used for the first edition, the
contrast was far sharper: 3.4 μs versus 100 μs.) Either way, besides
giving up CPU/GPU concurrency, the synchronization itself is worth
avoiding.
Even without synchronizations, if the kernel doesn’t run for longer
than the amount of time it took to launch the kernel (about 1.2 μs on
the RTX 3060), the GPU may go idle before the CPU has submitted more
work. To explore just how much work a kernel might need to do, in order
to make the launch worth doing, let’s switch to a kernel that busy-waits
until a certain number of clock cycles (according to the clock()
intrinsic) has completed:
__device__ int deviceTime;
__global__
void
WaitKernel( int cycles, bool bWrite )
{
int start = clock();
int stop;
do {
stop = clock();
} while ( stop - start < cycles );
if ( bWrite && threadIdx.x==0 && blockIdx.x==0 ) {
deviceTime = stop - start;
}
}
By conditionally writing the result to deviceTime, this
kernel prevents the compiler from optimizing out the busy wait – the
compiler does not know that we are just going to pass false as the
second parameter2. The code in our main() function
then checks the launch time for various values of cycles, from 0 to
2500.
for ( int cycles = 0; cycles < 2500; cycles += 100 ) {
printf( "Cycles: %d - ", cycles ); fflush( stdout );
start = std::chrono::steady_clock::now();
for ( int i = 0; i < cIterations; i++ ) {
WaitKernel<<<1,1>>>( cycles, false );
}
cudaDeviceSynchronize();
stop = std::chrono::steady_clock::now();
double us = 1e6*std::chrono::duration(stop - start).count();
double usPerLaunch = us / (float) cIterations;
printf( "%.2f us\n", usPerLaunch );
}
This program may be found in breakevenKernelAsync.cu
(waitKernelAsync.cu in the first edition). Plotting its output on the
book’s EC2 instance gives Figure 6-2:
Figure 6-2. Microseconds/cycles plot for
waitKernelAsync.cu
On this host platform, the breakeven mark where the kernel launch time crosses over 2× that of a NULL kernel launch (4.90 μs) is at 4500 GPU clock cycles. Both numbers shrink dramatically on modern hardware: on a GeForce RTX 3060, the NULL launch is about 1.2 μs, so breakeven (around 2.4 μs) arrives at roughly 2600 clock cycles.
These performance characteristics can vary widely and depend on many factors, including:
Performance of the host CPU,
Host operating system,
Driver version,
Whether ECC is enabled on the GPU.
But the common underlying theme is that for most CUDA applications, developers should do their best to avoid breaking CPU/GPU concurrency. Only applications that are very compute-intensive and only perform large data transfers can afford to ignore this overhead.
To take advantage of CPU/GPU concurrency when performing memory copies as well as kernel launches, developers must use asynchronous memcpy.
This program is in the source code as
nullKernelSync.cu, not reproduced here because it is almost
identical to Listing 6-1.↩︎
The compiler could still invalidate our timing results
by branching around the loop if bWrite is false. If the timing results
looked suspicious, we could see if this is happening by looking at the
microcode with cuobjdump.↩︎