Performance: Warps, Memory Coalescing & the Ecosystem

Warps & Divergence

The hardware schedules and executes threads in groups of 32 called warps, in lockstep. If threads within a warp take different branches of a conditional (warp divergence), the hardware executes both paths serially for that warp, masking off inactive threads each pass — a real performance cost worth avoiding in hot loops.

Global vs. Shared Memory

__global__ void sumBlock(float* input, float* output) {
    __shared__ float cache[256]; // fast, on-chip, scoped to this block
    int tid = threadIdx.x;
    cache[tid] = input[blockIdx.x * 256 + tid];
    __syncthreads(); // barrier — wait for all threads to finish writing
    // ... reduce cache within the block ...
}

Global memory is large but relatively slow and visible to all threads; shared memory is much faster but small and scoped per block — commonly used to cache data reused by many threads, avoiding repeated slow global memory round-trips. __syncthreads() is the barrier that keeps cooperating threads from racing.

Coalesced Access & Minimizing Transfers

When threads in a warp access contiguous global memory addresses, the hardware combines them into fewer, wider transactions — scattered access patterns force many separate slow transactions instead. Similarly, transfers across the host-device PCIe/NVLink boundary are comparatively slow, so well-optimized code batches and minimizes them, overlapping transfer with computation via CUDA streams where possible.

Why CUDA Dominates AI Compute

Neural network training and inference are dominated by large parallel matrix/tensor operations, which map extremely well onto GPU architecture. CUDA's decade-plus head start — plus libraries like cuDNN (optimized deep learning primitives) and cuBLAS, deeply integrated into PyTorch/TensorFlow/JAX — is why most practitioners never write a raw kernel: framework-level tensor ops already sit on a mature, GPU-accelerated foundation.

CUDA vs. Vendor-Neutral Alternatives

CUDA only runs on NVIDIA GPUs. OpenCL (and increasingly SYCL) exist as open, cross-vendor standards for heterogeneous parallel computing — but CUDA's tooling and ecosystem maturity keep it dominant specifically for deep learning workloads.