hpc · advanced

The GPU / CUDA execution model

A GPU is not a faster CPU — it's a throughput device with thousands of simple cores running one kernel under the SIMT model (Single Instruction, Multiple Threads): you launch a grid of thread blocks and each thread processes one data element. Three things govern performance. First, the GPU has separate memory, so inputs and outputs cross PCIe via cudaMemcpy — often costlier than the compute, so you move data once and keep it resident. Second, threads run in lockstep 32-thread warps: branch divergence within a warp serializes both paths, and uncoalesced (scattered) memory access multiplies transactions, so consecutive threads should touch consecutive addresses. Third, the memory hierarchy (registers > per-block shared memory > global VRAM) must be used deliberately. GPUs win on massively data-parallel, high-arithmetic-intensity, regular-access workloads (dense linear algebra, deep learning, FFTs, Monte Carlo) and lose on branchy, pointer-chasing, or transfer-bound work.

The code

// CUDA kernel: ONE function, run by thousands of threads
__global__ void saxpy(int n, float a, float* x, float* y) {
    int i = blockIdx.x * blockDim.x + threadIdx.x; // global thread id
    if (i < n)
        y[i] = a * x[i] + y[i]; // each thread: 1 element
}

cudaMemcpy(d_x, x, bytes, cudaMemcpyHostToDevice);  // 1. copy in (PCIe)
saxpy<<<(n + 255) / 256, 256>>>(n, 2.0f, d_x, d_y); // 2. launch grid of blocks
cudaMemcpy(y, d_y, bytes, cudaMemcpyDeviceToHost);  // 3. copy results back

What this lesson walks through

1. 01SIMT — one kernel, thousands of threads
2. 02The PCIe tax — data movement dominates
3. 03Warps, divergence & coalescing

You write ONE kernel; the GPU runs it across a grid of blocks, each with hundreds of threads. Each thread computes one element via its global id blockIdx*blockDim + threadIdx.