Asynchronous Commands in CUDA
As described by the CUDA C Programming Guide, asynchronous commands return control to the calling host thread before the device has finished the requested task (they are non-blocking). These commands are:
- Kernel launches;
- Memory copies between two addresses to the same device memory;
- Memory copies from host to device of a memory block of 64 KB or less;
- Memory copies performed by functions with the
Async
suffix; - Memory set function calls.
Specifying a stream for a kernel launch or host-device memory copy is optional; you can invoke CUDA commands without specifying a stream (or by setting the stream parameter to zero). The following two lines of code both launch a kernel on the default stream.
kernel<<< blocks, threads, bytes >>>(); // default stream
kernel<<< blocks, threads, bytes, 0 >>>(); // stream 0
The Default Stream
The default stream is useful where concurrency is not crucial to performance. Before CUDA 7, each device has a single default stream used for all host threads, which causes implicit synchronization. As the section “Implicit Synchronization” in the CUDA C Programming Guide explains, two commands from different streams cannot run concurrently if the host thread issues any CUDA command to the default stream between them.
CUDA 7 introduces a new option, the per-thread default stream, that has two effects. First, it gives each host thread its own default stream. This means that commands issued to the default stream by different host threads can run concurrently. Second, these default streams are regular streams. This means that commands in the default stream may run concurrently with commands in non-default streams.
To enable per-thread default streams in CUDA 7 and later, you can either compile with the nvcc
command-line option --default-stream per-thread
, or #define
theCUDA_API_PER_THREAD_DEFAULT_STREAM
preprocessor macro before including CUDA headers (cuda.h
or cuda_runtime.h
). It is important to note: you cannot use #define CUDA_API_PER_THREAD_DEFAULT_STREAM
to enable this behavior in a .cu file when the code is compiled by nvcc
because nvcc
implicitly includes cuda_runtime.h
at the top of the translation unit.
A Multi-Stream Example
Let’s look at a trivial example. The following code simply launches eight copies of a simple kernel on eight streams. We launch only a single thread block for each grid so there are plenty of resources to run multiple of them concurrently. As an example of how the legacy default stream causes serialization, we add dummy kernel launches on the default stream that do no work. Here’s the code.
const int N = 1 << 20;
__global__ void kernel(float *x, int n)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
for (int i = tid; i < n; i += blockDim.x * gridDim.x) {
x[i] = sqrt(pow(3.14159,i));
}
}
int main()
{
const int num_streams = 8;
cudaStream_t streams[num_streams];
float *data[num_streams];
for (int i = 0; i < num_streams; i++) {
cudaStreamCreate(&streams[i]);
cudaMalloc(&data[i], N * sizeof(float));
// launch one worker kernel per stream
kernel<<<1, 64, 0, streams[i]>>>(data[i], N);
// launch a dummy kernel on the default stream
kernel<<<1, 1>>>(0, 0);
}
cudaDeviceReset();
return 0;
}
First let’s check out the legacy behavior, by compiling with no options.
nvcc ./stream_test.cu -o stream_legacy
We can run the program in the NVIDIA Visual Profiler (nvvp
) to get a timeline showing all streams and kernel launches. Figure 1 shows the resulting kernel timeline on a Macbook Pro with an NVIDIA GeForce GT 750M (a Kepler GPU). You can see the very small bars for the dummy kernels on the default stream, and how they cause all of the other streams to serialize.
![Legacy Default Stream Serialization](https://devblogs.nvidia.com/wp-content/uploads/2015/01/stream_legacy-624x178.png)
Now let’s try the new per-thread default stream.
nvcc –default-stream per-thread ./stream_test.cu -o stream_per-thread
Figure 2 shows the results from nvvp
. Here you can see full concurrency between nine streams: the default stream, which in this case maps to Stream 14, and the eight other streams we created. Note that the dummy kernels run so quickly that it’s hard to see that there are eight calls on the default stream in this image.
![Multi-Stream Per-Thread Default Stream Behavior](https://devblogs.nvidia.com/wp-content/uploads/2015/01/stream_per_thread-624x168.png)
A Multi-threading Example
Let’s look at another example, designed to demonstrate how the new default stream behavior makes it easier to achieve execution concurrency in multi-threaded applications. The following example creates eight POSIX threads, and each thread calls our kernel on the default stream and then synchronizes the default stream. (We need the synchronization in this example to make sure the profiler gets the kernel start and end timestamps before the program exits.)
#include <pthread.h>
#include <stdio.h>
const int N = 1 << 20;
__global__ void kernel(float *x, int n)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
for (int i = tid; i < n; i += blockDim.x * gridDim.x) {
x[i] = sqrt(pow(3.14159,i));
}
}
void *launch_kernel(void *dummy)
{
float *data;
cudaMalloc(&data, N * sizeof(float));
kernel<<<1, 64>>>(data, N);
cudaStreamSynchronize(0);
return NULL;
}
int main()
{
const int num_threads = 8;
pthread_t threads[num_threads];
for (int i = 0; i < num_threads; i++) {
if (pthread_create(&threads[i], NULL, launch_kernel, 0)) {
fprintf(stderr, "Error creating threadn");
return 1;
}
}
for (int i = 0; i < num_threads; i++) {
if(pthread_join(threads[i], NULL)) {
fprintf(stderr, "Error joining threadn");
return 2;
}
}
cudaDeviceReset();
return 0;
}
First, let’s compile with no options to test the legacy default stream behavior.
nvcc ./pthread_test.cu -o pthreads_legacy
When we run this in nvvp
, we see a single stream, the default stream, with all kernel launches serialized, as Figure 3 shows.
![Legacy Default Stream Behavior in a Multi-threaded app.](https://devblogs.nvidia.com/wp-content/uploads/2015/01/pthreads_legacy-624x36.png)
Let’s compile it with the new per-thread default stream option.
nvcc –default-stream per-thread ./pthread_test.cu -o pthreads_per_thread
Figure 4 shows that with per-thread default streams, each thread creates a new stream automatically and they do not synchronize, so the kernels from all eight threads run concurrently.
![New Default Stream Behavior In a Multi-threaded App](https://devblogs.nvidia.com/wp-content/uploads/2015/01/pthreads_per_thread-624x167.png)
More Tips
Here are a few more things to keep in mind when programming for concurrency.
- Remember: With per-thread default streams, the default stream in each thread behaves the same as a regular stream, as far as synchronization and concurrency goes. This is not true with the legacy default stream.
- The
--default-stream
option is applied per compilation unit, so make sure to apply it to allnvcc
command lines that need it. cudaDeviceSynchronize()
continues to synchronize everything on the device, even with the new per-thread default stream option. If you want to only synchronize a single stream, usecudaStreamSynchronize(cudaStream_t stream)
, as in our second example.- Starting in CUDA 7 you can also explicitly access the per-thread default stream using the handle
cudaStreamPerThread
, and you can access the legacy default stream using the handlecudaStreamLegacy
. Note thatcudaStreamLegacy
still synchronizes implicitly with the per-thread default streams if you happen to mix them in a program. - You can create non-blocking streams which do not synchronize with the legacy default stream by passing the
cudaStreamNonBlocking
flag to cudaStreamCreate().