CUDA Python: Memory Transfers & Benchmarking

Transfers Dominate GPU Workloads (Theory)

In GPU applications, total time often looks like:

• Host↔Device transfers (PCIe/NVLink)
• Kernel execution
• CPU-side overhead (launch, preprocessing)
• Synchronization

Optimization priorities:

1. Move data once, do many computations on device
2. Fuse operations to reduce intermediate transfers
3. Use pinned memory for faster DMA
4. Avoid unnecessary copies (especially in tight loops)

Code Example 1 — Correct Benchmarking with Numba CUDA

import time
import numpy as np
from numba import cuda

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

def benchmark(N=1 << 24, iters=30):
    a = np.ones(N, dtype=np.float32)
    b = np.full(N, 2.0, dtype=np.float32)

    # Allocate device buffers once (important!)
    d_a = cuda.to_device(a)
    d_b = cuda.to_device(b)
    d_c = cuda.device_array_like(a)

    threads_per_block = 256
    blocks_per_grid = (N + threads_per_block - 1) // threads_per_block

    # Warm-up (JIT compilation + cache effects)
    vec_add[blocks_per_grid, threads_per_block](d_a, d_b, d_c)
    cuda.synchronize()

    # Timed section (kernel only)
    t0 = time.perf_counter()
    for _ in range(iters):
        vec_add[blocks_per_grid, threads_per_block](d_a, d_b, d_c)
    cuda.synchronize()  # ensure kernels finished
    t1 = time.perf_counter()

    ms_per_iter = (t1 - t0) * 1000 / iters
    print(f"Kernel time: {ms_per_iter:.3f} ms/iter")

    # Copy once for correctness check
    c = d_c.copy_to_host()
    print("c[0] =", c[0], "c[N-1] =", c[-1])

if __name__ == "__main__":
    benchmark()

Code Example 2 — Measuring End-to-End (Including Transfers)

If your real workload includes transfers each iteration, you must measure that too.

import time
import numpy as np
from numba import cuda

@cuda.jit
def scale_kernel(x, y, alpha):
    i = cuda.grid(1)
    if i < y.size:
        y[i] = alpha * x[i]

def benchmark_end_to_end(N=1 << 22, iters=50):
    threads = 256
    blocks  = (N + threads - 1) // threads

    alpha = 1.5
    x_host = np.ones(N, dtype=np.float32)

    # Warm-up kernel compile
    d_x = cuda.to_device(x_host)
    d_y = cuda.device_array_like(x_host)
    scale_kernel[blocks, threads](d_x, d_y, alpha)
    cuda.synchronize()

    t0 = time.perf_counter()
    for _ in range(iters):
        # Host->Device
        d_x = cuda.to_device(x_host)
        # Kernel
        scale_kernel[blocks, threads](d_x, d_y, alpha)
        cuda.synchronize()
        # Device->Host
        y_host = d_y.copy_to_host()
    t1 = time.perf_counter()

    ms_per_iter = (t1 - t0) * 1000 / iters
    print(f"End-to-end time: {ms_per_iter:.3f} ms/iter")

if __name__ == "__main__":
    benchmark_end_to_end()

Common Gotchas

• Timing without synchronization: kernels launch asynchronously; always cuda.synchronize() before reading timers.
• Copying device buffers repeatedly: allocate once if possible.
• Benchmarking after cold start: ignore first iteration (JIT + caching).
• Measuring only kernel time when transfers matter in production.

Quick Checklist

• Allocate device arrays once
• Warm up before timing
• Synchronize before stopping the timer
• Benchmark kernel-only and end-to-end if needed
• Minimize host↔device transfers