┌──────┐ ┌──────┐ ┌──────┐ ┌──────┐ │ core │ │ core │ │ core │ │ core │ × thousands │ ## │ │ ## │ │ ## │ │ ## │ └──┬───┘ └──┬───┘ └──┬───┘ └──┬───┘ │ │ │ │ ═══╧════════╧════════╧════════╧═══ shared bus

CUDA Parallelism

Thousands of threads, one idea

01 The Big Picture

A CPU is a latency machine — a few powerful cores that rip through sequential tasks. A GPU is a throughput machine — thousands of modest cores that chew through data in parallel. CUDA is NVIDIA's programming model that lets you harness those cores from ordinary C/C++ code. You write a function once, and the hardware runs it across thousands of threads simultaneously.

The central abstraction is simple: instead of looping over data, you launch a grid of threads where every thread handles its own slice of the problem. The GPU scheduler maps those threads onto physical cores automatically. You think in parallel; the hardware executes in parallel.

02 Kernels — The Unit of Work

A kernel is a function marked with __global__ that runs on the GPU. When you call it, you specify how many threads to launch using the <<<blocks, threads>>> syntax. Every thread executes the same code but knows its own identity via built-in variables like threadIdx and blockIdx, so each thread operates on different data.

This pattern — one program, many data — is called SIMT (Single Instruction, Multiple Threads). It descends from the classical SIMD model but adds per-thread branching flexibility.

03 Thread Hierarchy

CUDA organizes threads into a two-level hierarchy:

Grid ┌─────────────────────────────────────────────┐ │ Block (0,0) Block (1,0) Block (2,0)│ │ ┌───────────┐ ┌───────────┐ ┌─────────┐ │ │ │ t0 t1 t2 │ │ t0 t1 t2 │ │ t0 t1 t2│ │ │ │ t3 t4 t5 │ │ t3 t4 t5 │ │ t3 t4 t5│ │ │ └───────────┘ └───────────┘ └─────────┘ │ │ │ │ Block (0,1) Block (1,1) Block (2,1)│ │ ┌───────────┐ ┌───────────┐ ┌─────────┐ │ │ │ t0 t1 t2 │ │ t0 t1 t2 │ │ t0 t1 t2│ │ │ │ t3 t4 t5 │ │ t3 t4 t5 │ │ t3 t4 t5│ │ │ └───────────┘ └───────────┘ └─────────┘ │ └─────────────────────────────────────────────┘

Blocks can be 1D, 2D, or 3D — a convenience that maps naturally to vectors, matrices, or volumes. The hardware doesn't care about dimensionality; it flattens everything internally.

· · · ──── ▸▸ ──── · · ·

04 Warps — Where the Magic Happens

The GPU doesn't schedule individual threads. It schedules warps — groups of 32 threads that execute in lockstep on a single Streaming Multiprocessor (SM). Every clock cycle, all 32 threads in a warp execute the same instruction.

When threads in a warp diverge at an if statement, the hardware serializes the paths: it runs one branch with some threads masked off, then the other. This is warp divergence, and it costs performance. The golden rule: keep threads within a warp on the same code path whenever possible.

05 Memory Hierarchy

Performance in CUDA is dominated not by compute but by memory access patterns. The hierarchy, from fastest to slowest:

Speed Level Scope Size ───────────────────────────────────────────────── █████ Registers per thread ~255 regs ████░ Shared Memory per block up to 228 KB ███░░ L1 / L2 Cache per SM / GPU varies █░░░░ Global Memory all threads up to 80 GB ▒░░░░ Host (CPU) RAM CPU side system RAM

Shared memory is the key tool. It lives on-chip, has roughly the same latency as L1 cache, and is explicitly managed by the programmer. Classic technique: threads in a block cooperatively load a tile of data from slow global memory into fast shared memory, synchronize with __syncthreads(), then compute on the tile at full speed.

06 Coalesced Access

Global memory is accessed in 32-byte or 128-byte transactions. When all 32 threads in a warp read consecutive addresses, the hardware merges — coalesces — those requests into a single wide transaction. Scattered reads waste bandwidth because the bus fetches full cache lines regardless. Aligning data structures and using Struct-of-Arrays instead of Array-of-

to be continued....