CUDA Python: Streams, Overlap, and (Conceptual) Multi-GPU

Streams Enable Concurrency (Theory)

CUDA uses a command ordering model:

• A stream is a queue of operations that execute in issue order.
• Operations in the same stream are sequential.
• Operations in different streams may overlap, depending on hardware and resource availability.

Typical pipeline:

1. H2D transfer (host → device)
2. Kernel execution
3. D2H transfer (device → host)

Using streams, you can overlap these steps for different data chunks.

Code Example 1 — Overlap H2D copies and Kernel with Streams (Numba CUDA)

import numpy as np
from numba import cuda
import math

@cuda.jit
def add_kernel(a, b, out):
    i = cuda.grid(1)
    if i < out.size:
        out[i] = a[i] + b[i]

def run_overlap(N=1 << 22, chunk=1 << 20):
    # Host arrays (use pagelocked/pinned for best transfer performance)
    a = np.ones(N, dtype=np.float32)
    b = np.full(N, 2.0, dtype=np.float32)
    out = np.empty_like(a)

    # Prepare pinned buffers
    a_p = cuda.pinned_array_like(a)
    b_p = cuda.pinned_array_like(b)
    out_p = cuda.pinned_array_like(out)
    a_p[:] = a
    b_p[:] = b

    # Create two streams for ping-pong buffering
    stream1 = cuda.stream()
    stream2 = cuda.stream()

    threads = 256
    blocks = (chunk + threads - 1) // threads

    # Allocate device buffers for a chunk (reused each iteration)
    d_a = cuda.device_array(chunk, dtype=np.float32)
    d_b = cuda.device_array(chunk, dtype=np.float32)
    d_out = cuda.device_array(chunk, dtype=np.float32)

    num_chunks = math.ceil(N / chunk)

    for c in range(num_chunks):
        start = c * chunk
        end = min(N, (c + 1) * chunk)
        size = end - start

        # Select stream (alternate chunks)
        s = stream1 if (c % 2 == 0) else stream2

        # Copy chunk to device asynchronously on stream s
        # Note: Numba provides asynchronous copies when using streams.
        d_a[:size].copy_to_device(a_p[start:end], stream=s)
        d_b[:size].copy_to_device(b_p[start:end], stream=s)

        # Launch kernel in same stream
        add_kernel[blocks, threads, s](d_a, d_b, d_out)

        # Copy result back asynchronously in same stream
        d_out[:size].copy_to_host(out_p[start:end], stream=s)

    # Wait for both streams to finish before using results
    stream1.synchronize()
    stream2.synchronize()

    # out_p now contains the results
    print("out[0] =", float(out_p[0]), "out[N-1] =", float(out_p[N - 1]))

if __name__ == "__main__":
    run_overlap()

Why this works

• Each chunk’s H2D, kernel, and D2H are ordered within a stream.
• Different chunks use different streams, enabling overlap.

Code Example 2 — Synchronization Safety Pattern

A common bug: reading results before work is complete.

Rule of thumb:

• If you launch kernels/copies into stream s, then before you read out, do s.synchronize() (or use events).

# Pseudocode pattern
s = cuda.stream()
d_x.copy_to_device(x_host, stream=s)
kernel[blocks, threads, s](d_x, ...)
d_y.copy_to_host(y_host, stream=s)

# Safe consumption
s.synchronize()
print(y_host[0])

Common Gotchas

• Overlapping copies too much: too many concurrent operations can saturate memory bandwidth and reduce benefits.
• Non-pinned host memory: async transfers may degrade without pinned/page-locked buffers.
• Using wrong stream: ensure copies and kernel for a chunk use the same stream to preserve order.
• Multi-GPU requires explicit device management: cuda.select_device(i) (or equivalent) and separate buffers per device.

Quick Checklist

• Use streams to pipeline chunked work
• Use pinned host buffers for best async transfer performance
• Synchronize the correct streams before reading outputs
• Measure with real end-to-end workload timing