Cuda Programming 9: Advanced CUDA Optimization – Using CUDA Streams and Events for Concurrency

Introduction:

CUDA provides several ways to achieve concurrency and optimize the performance of GPU-accelerated applications. One of the most powerful features for improving parallelism is the use of CUDA streams and events. These tools allow for overlapping data transfer, computation, and kernel execution, which can lead to significant performance gains, especially in applications with complex workflows.

In this blog, we will cover how to use CUDA Streams and CUDA Events to maximize GPU utilization and improve the overall efficiency of your CUDA programs.

1. What is a CUDA Stream?

A CUDA stream is a sequence of operations (e.g., kernels, memory copies) that execute in order on the GPU. Streams allow for overlapping the execution of different operations. By default, all operations in CUDA run in stream 0 (the default stream). However, by creating additional streams, you can achieve concurrent execution of kernels and memory operations.

Why Use CUDA Streams?

• Concurrency: Stream operations can execute concurrently, allowing the CPU to initiate work on the GPU while the GPU works on a different task.
• Overlapping Computation and Communication: Streams enable overlapping of memory copies and kernel execution, significantly reducing idle times.

Example of Using CUDA Streams:

Let’s take a look at a simple example that demonstrates how to overlap kernel execution and memory transfers using streams.

cppCopyEdit#include <iostream>
#include <cuda_runtime.h>

#define N 1024

// Kernel to add two arrays
__global__ void add(int *A, int *B, int *C) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < N) {
        C[idx] = A[idx] + B[idx];
    }
}

int main() {
    int *h_A, *h_B, *h_C;
    int *d_A, *d_B, *d_C;
    size_t size = N * sizeof(int);
    
    // Allocate memory on host and device
    h_A = (int*)malloc(size);
    h_B = (int*)malloc(size);
    h_C = (int*)malloc(size);
    cudaMalloc(&d_A, size);
    cudaMalloc(&d_B, size);
    cudaMalloc(&d_C, size);

    // Initialize host arrays
    for (int i = 0; i < N; i++) {
        h_A[i] = i;
        h_B[i] = i * 2;
    }

    // Create streams
    cudaStream_t stream1, stream2;
    cudaStreamCreate(&stream1);
    cudaStreamCreate(&stream2);

    // Copy data to device in parallel using streams
    cudaMemcpyAsync(d_A, h_A, size, cudaMemcpyHostToDevice, stream1);
    cudaMemcpyAsync(d_B, h_B, size, cudaMemcpyHostToDevice, stream2);

    // Launch kernel in stream1
    add<<<(N + 255) / 256, 256, 0, stream1>>>(d_A, d_B, d_C);

    // Copy result back to host in parallel using streams
    cudaMemcpyAsync(h_C, d_C, size, cudaMemcpyDeviceToHost, stream1);

    // Wait for stream1 to finish before exiting
    cudaStreamSynchronize(stream1);

    // Clean up
    free(h_A);
    free(h_B);
    free(h_C);
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);
    cudaStreamDestroy(stream1);
    cudaStreamDestroy(stream2);

    return 0;
}

Explanation:

• Streams: Two streams (stream1 and stream2) are created using cudaStreamCreate(). These streams allow for concurrent execution of memory copy operations and kernel execution.
• cudaMemcpyAsync: Memory copies to and from the device are performed asynchronously, meaning they do not block the CPU while waiting for the GPU to finish the transfer.
• Kernel Execution: The kernel add is launched asynchronously in stream1, while data transfers happen in parallel.
• Synchronization: cudaStreamSynchronize() ensures that the program waits for stream1 to finish before the program exits, making sure all operations are completed.

By using streams, we can achieve better utilization of the GPU by overlapping memory transfers and kernel executions, which helps improve performance in high-performance applications.

2. What is a CUDA Event?

A CUDA event is a point in time that can be used to record when an operation starts or finishes. Events are typically used in conjunction with streams to track the progress of operations and synchronize different parts of the program.

Why Use CUDA Events?

• Synchronization: Events allow for precise synchronization between different streams and operations.
• Timing: Events can be used to measure the time taken for a particular operation, helping you to profile and optimize your code.

Example of Using CUDA Events:

Let’s look at how we can use CUDA events to measure the time taken for kernel execution and memory transfers.

cppCopyEdit#include <iostream>
#include <cuda_runtime.h>

#define N 1024

// Kernel to add two arrays
__global__ void add(int *A, int *B, int *C) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < N) {
        C[idx] = A[idx] + B[idx];
    }
}

int main() {
    int *h_A, *h_B, *h_C;
    int *d_A, *d_B, *d_C;
    size_t size = N * sizeof(int);

    // Allocate memory on host and device
    h_A = (int*)malloc(size);
    h_B = (int*)malloc(size);
    h_C = (int*)malloc(size);
    cudaMalloc(&d_A, size);
    cudaMalloc(&d_B, size);
    cudaMalloc(&d_C, size);

    // Initialize host arrays
    for (int i = 0; i < N; i++) {
        h_A[i] = i;
        h_B[i] = i * 2;
    }

    // Create CUDA events
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);

    // Record start event
    cudaEventRecord(start);

    // Perform memory copy and kernel execution
    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
    add<<<(N + 255) / 256, 256>>>(d_A, d_B, d_C);
    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

    // Record stop event
    cudaEventRecord(stop);

    // Synchronize the events
    cudaEventSynchronize(stop);

    // Measure the time taken
    float milliseconds = 0;
    cudaEventElapsedTime(&milliseconds, start, stop);
    std::cout << "Time taken: " << milliseconds << " ms" << std::endl;

    // Clean up
    free(h_A);
    free(h_B);
    free(h_C);
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);
    cudaEventDestroy(start);
    cudaEventDestroy(stop);

    return 0;
}

Explanation:

• Events for Timing: We create two events, start and stop, using cudaEventCreate(). These events are used to record when memory transfers and kernel executions begin and end.
• cudaEventRecord(): This function records the current timestamp to the event.
• cudaEventElapsedTime(): After synchronization, we use this function to measure the time between the start and stop events, giving us the elapsed time in milliseconds.
• Synchronization: We synchronize the stop event to ensure that the kernel and memory operations are complete before measuring the elapsed time.

By using CUDA events, we can accurately profile different sections of the program and optimize execution times.

3. Best Practices for Using Streams and Events

• Avoid Serialization: Make sure that independent operations are placed in different streams to allow for concurrent execution.
• Use cudaMemcpyAsync: For data transfers, always use asynchronous memory copy operations when working with streams to maximize performance.
• Timing and Profiling: Use CUDA events to profile your application and measure the time taken by different operations to help identify performance bottlenecks.

4. Conclusion

CUDA Streams and Events are powerful tools for achieving concurrency in GPU-accelerated applications. By allowing the GPU to perform multiple tasks in parallel, such as kernel execution and memory transfers, these features can lead to significant performance improvements, especially in more complex applications.

In the next blog, we’ll discuss CUDA Graphs, which provide a way to optimize complex execution dependencies, further improving the performance of your CUDA applications.

Call to Action:

• How have you used CUDA streams and events in your applications? Share your experiences in the comments below!
• Next up: Stay tuned for a deep dive into CUDA Graphs and their impact on GPU performance.