CUDA is best-known for enabling fine-grained concurrency, with hardware facilities that enable thousands or millions of threads to run in parallel; but it also has hardware and software facilities that enable more coarse-grained concurrency:
CPU/GPU concurrency: since they are separate devices, the CPU and GPU can operate independently of one another;
Memcpy/kernel processing concurrency: for GPUs that have one or more Copy Engines, host↔︎device memcpy can be performed while the SMs are processing kernels,
Kernel concurrency: kernels can run in parallel.
Multi-GPU concurrency: For problems with enough computational density, multiple GPUs can operate in parallel. (Chapter 9 is dedicated to multi-GPU programming.)
Within a stream, operations are performed in sequential order, but operations in different streams may be performed in parallel.
CUDA events complement CUDA streams by providing the synchronization mechanisms needed to coordinate the parallel execution enabled by streams. CUDA events may be asynchronously “recorded” into a stream, and the CUDA event becomes signaled when the operations preceding the CUDA event have been completed.
CUDA events may be used for CPU/GPU synchronization, for synchronization between the engines on the GPU, and for synchronization between GPUs.
CUDA events also provide a GPU-based timing mechanism, which cannot be perturbed by system events such as page faults or interrupts from disk or network controllers. Wall clock timers are best for overall timing, but CUDA events are useful for optimizing kernels or figuring out which of a series of pipelined GPU operations is taking the longest.
This chapter also covers CUDA graphs, which can encapsulate a series of CUDA operations—from memcpys to kernel launches, to memsets or host-side callbacks or even child graphs—into a form that can be instantiated and played back later to reduce CPU overhead.
cudaStreamWaitEvent()