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 viewBox="0 0 687 188" xmlns="http://www.w3.org/2000/svg" style="max-width:100%;height:auto" role="img"><rect x="0" y="0" width="687" height="188" rx="12" fill="#0d1117"/><g font-family="ui-monospace,SFMono-Regular,Menlo,Consolas,&quot;Liberation Mono&quot;,monospace" font-size="14"><text xml:space="preserve" x="20" y="31.7"><tspan fill="#c9d1d9">Grid (Kernel Launch)</tspan></text><text xml:space="preserve" x="20" y="50.7"><tspan fill="#6e7681">├──</tspan><tspan fill="#c9d1d9"> Block (0,0)          </tspan><tspan fill="#6e7681">←</tspan><tspan fill="#c9d1d9"> Thread Block: 32-1024 threads, scheduled on one SM</tspan></text><text xml:space="preserve" x="20" y="69.7"><tspan fill="#6e7681">│</tspan><tspan fill="#c9d1d9">   </tspan><tspan fill="#6e7681">├──</tspan><tspan fill="#c9d1d9"> Warp 0 (threads 0-31)   </tspan><tspan fill="#6e7681">←</tspan><tspan fill="#c9d1d9"> 32 threads executing in SIMT lockstep</tspan></text><text xml:space="preserve" x="20" y="88.7"><tspan fill="#6e7681">│</tspan><tspan fill="#c9d1d9">   </tspan><tspan fill="#6e7681">├──</tspan><tspan fill="#c9d1d9"> Warp 1 (threads 32-63)</tspan></text><text xml:space="preserve" x="20" y="107.7"><tspan fill="#6e7681">│</tspan><tspan fill="#c9d1d9">   </tspan><tspan fill="#6e7681">└──</tspan><tspan fill="#c9d1d9"> ...</tspan></text><text xml:space="preserve" x="20" y="126.7"><tspan fill="#6e7681">├──</tspan><tspan fill="#c9d1d9"> Block (0,1)</tspan></text><text xml:space="preserve" x="20" y="145.7"><tspan fill="#6e7681">├──</tspan><tspan fill="#c9d1d9"> Block (1,0)</tspan></text><text xml:space="preserve" x="20" y="164.7"><tspan fill="#6e7681">└──</tspan><tspan fill="#c9d1d9"> ... (up to 2^31 blocks)</tspan></text></g></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.