gpu programming model
**GPU Programming Model and Thread Hierarchy** is the **software abstraction that organizes millions of GPU threads into a hierarchical structure — grids of thread blocks (each containing hundreds of threads organized into warps of 32) — where the programmer expresses parallelism at the thread block level while the hardware scheduler dynamically maps blocks to Streaming Multiprocessors (SMs), enabling a single program to scale from a 10-SM laptop GPU to a 132-SM data center accelerator without code changes**.
**Thread Hierarchy**
```svg
```
- **Thread**: The finest granularity of execution. Each thread has its own registers and program counter (logically — physically, warps share a PC).
- **Warp (32 threads)**: The hardware scheduling unit. All 32 threads execute the same instruction simultaneously (SIMT). Divergent branches cause warp serialization.
- **Thread Block (32-1024 threads)**: The programmer-defined grouping. All threads in a block execute on the same SM, share shared memory (up to 228 KB on H100), and can synchronize with __syncthreads().
- **Grid**: All thread blocks in a kernel launch. Blocks execute independently in any order — the GPU hardware schedules them dynamically.
**Why This Hierarchy Works**
- **Scalability**: The programmer specifies blocks, not SM assignments. A grid of 1000 blocks runs on a 10-SM GPU with 100 blocks per SM (time-sliced) or a 100-SM GPU with 10 blocks per SM (all concurrent). The same kernel binary scales automatically.
- **Synchronization Scope**: Threads within a block can synchronize (barrier) and communicate (shared memory). Threads in different blocks cannot synchronize (no global barrier within a kernel) — this independence is what enables the scheduler's flexibility.
**Cooperative Groups (CUDA 9+)**
Extends the programming model beyond the block level:
- **Thread Block Tile**: Partition a block into fixed-size tiles (e.g., 32 threads = warp) with tile-level sync and collective operations.
- **Grid Group**: All blocks in a kernel can synchronize using cooperative launch (grid-wide barrier). Requires all blocks to be resident simultaneously — limits the number of blocks.
- **Multi-Grid Group**: Synchronization across multiple kernel launches.
**Occupancy and Scheduling**
The SM scheduler assigns as many blocks to each SM as resources allow (registers, shared memory, max threads per SM). For example, if each block uses 64 registers per thread × 256 threads = 16,384 registers per block, and the SM has 65,536 registers, then 4 blocks can be resident simultaneously. Higher occupancy (more warps in-flight) helps hide memory latency.
**Thread Indexing**
```
int gid = blockIdx.x * blockDim.x + threadIdx.x; // Global thread ID
int lid = threadIdx.x; // Local (block) ID
```
The global ID maps each thread to a unique data element. The local ID selects shared memory locations. Multi-dimensional indexing (3D grids and blocks) naturally maps to 2D/3D data structures.
The GPU Programming Model is **the abstraction that makes massively parallel hardware programmable** — hiding the complexity of warp scheduling, SM assignment, and hardware resource management behind a clean hierarchical model that lets programmers focus on the parallel algorithm rather than the machine architecture.