Thread blocks and grids

The Two-Level Hierarchy

GPUs execute millions of threads simultaneously, but hardware constraints force a structure:

Thread Block (Cooperative Thread Array):

A group of threads (typically 128–1024) that execute the same kernel code
All threads in a block run on the same Streaming Multiprocessor (SM)
Share fast on-chip shared memory (user-managed cache)
Can synchronize using __syncthreads() barrier
Maximum size: hardware-limited (e.g., 1024 threads/block on modern GPUs)

Grid:

A collection of thread blocks launched together
Blocks can be 1D, 2D, or 3D arrays (matching problem geometry)
Blocks execute independently and may run in any order
Spans all SMs on the GPU
No cross-block synchronization within a kernel (blocks are autonomous units)

Why This Design?

Why blocks exist:

Scalability: Same kernel runs on GPUs with 10 SMs or 100 SMs—runtime just schedules blocks differently
Data locality: Threads processing nearby data elements stay physically close, sharing fast memory
Synchronization: Only threads that need to coordinate (within a block) can synchronize—avoids global stalls

Why grids exist:

Problem decomposition: Natural mapping to data structures (2D grid for images, 3D for volumes)
Paralelism exposure: Programmer expresses maximum parallelism; hardware exploits what it can
Load balancing: Runtime distributes blocks to SMs dynamically, hiding idle time

Deriving Global Thread IDs

For a 1D problem (e.g., vector addition of length N):

Setup:

Grid has G blocks, each with B threads
Total threads: $T = G \times B$
Each thread processes one element

Why this indexing? Thread 0 in block 0 should map to element 0, thread 0 in block 1 should map to element B, etc.

Derivation:

Block 0: threads [0, B-1]      → elements [0, B-1]
Block 1: threads [0, B-1]      → elements [B, 2B-1]
Block k: threads [0, B-1]      → elements [kB, (k+1)B-1]

For thread t in block k: $\text{global\_id} = k \cdot B + t = \text{blockIdx.x} \times \text{blockDim.x} + \text{threadIdx.x}$

Why this step? The block index tells us how many complete blocks came before (each contributing B elements), and threadIdx gives our position within the current block.

Worked example Vector Addition (1D)

Add two arrays: C[i] = A[i] + B[i] for i = 0..N-1

// Host code: launch kernel
int N = 1000000;
int threadsPerBlock = 256;
int blocks = (N + threadsPerBlock - 1) / threadsPerBlock;  // ceiling division
vecAdd<<<blocks, threadsPerBlock>>(A, B, C, N);
// Kernel
__global__ void vecAdd(float *A, float *B, float *C, int N) {
   int i = blockIdx.x * blockDim.x + threadIdx.x;
   if (i < N) {  // Guard for oversubscription
       C[i] = A[i] + B[i];
   }
}

Why ceiling division? If N=1000 and block size=256, we need 4 blocks (not 3) to cover all elements. Formula: $\lceil N/B \rceil = (N + B - 1) / B$ integer math.

Why the if-guard? Last block may have extra threads (block 4 has 256 threads but only 232 elements left). Without the guard, threads access invalid memory.

Worked example Matrix Multiplication (2D)

Compute C = A × B where C[row][col] = sum(A[row][k] * B[k][col])

// Launch: 2D grid of2D blocks
dim3 threadsPerBlock(16, 16);  // 256 threads
dim3 blocks((N + 15)/16, (M + 15)/16);
matMul<<<blocks, threadsPerBlock>>>(A, B, C, M, N, K);
__global__ void matMul(float *A, float *B, float *C, int M, int N, int K) {
   int row = blockIdx.y * blockDim.y + threadIdx.y;
   int col = blockIdx.x * blockDim.x + threadIdx.x;
   if (row < M && col < N) {
       float sum = 0.0f;
       for (int k = 0; k < K; k++) {
           sum += A[row * K + k] * B[k * N + col];
       }
       C[row * N + col] = sum;
   }
}

Why 2D blocks? Each output element C[row][col] is independent—mapping threads to outputs directly is natural. 16×16 blocks hit sweet spot: enough threads to hide latency, small enough to fit shared memory (for optimized versions).

Why blockIdx.y for rows? Convention: x-axis is horizontal (columns), y-axis is vertical (rows matching mathematical matrix notation.

Block Scheduling

Warp-based execution within blocks:

Threads in are grouped into warps (32 threads)
All threads in a warp execute the same instruction simultaneously (SIMT model)
Block size should be a multiple of 32 for efficiency

Example: 256-thread block = 8 warps

If threads diverge (if-else), some warp threads are masked—still consume cycles
Synchronization via __syncthreads() ensures all warps in a block reach the barrier before any proceeds

Resource limits per SM:

Maximum threads (e.g., 2048 on Ampere)
Registers per thread
Shared memory per block
Maximum blocks (typically 16-32)

Why block size matters:

Too small (e.g., 32): underutilizes SM, can't hide memory latency
Too large (e.g., 1024): fewer blocks fit on SM, reduces occupancy if one block stalls
Sweet spot: 128–512 threads, chosen empirically or via occupancy calculator

Common mistake "Blocks Synchronize Automatically"

Wrong idea: Assume block0 finishes before block 1 starts, or use global memory to "pass messages" between blocks within the same kernel.

Why it feels right: Blocks are numbered sequentially (blockIdx.x = 0, 1, 2...), suggesting order.

The fix: Blocks run in arbitrary order and may overlap in time. The hardware can schedule block 5 before block 1 if SM2 is idle. If you need cross-block coordination:

Option 1: Split into multiple kernel launches (implicit global barrier)
Option 2: Use grid-level synchronization (CUDA 9+, cooperative groups)
Option 3: Atomic operations on global memory (expensive)

Example: Don't do this:

// Block 0 writes flag=1, other blocks wait—DEADLOCK!
if (blockIdx.x == 0) {
   work();
   flag[0] = 1;
} else {
   while (flag[0] == 0);  // ❌ May run before block 0!
}

Recall Explain to a 12-Year-Old

Imagine you're organizing a giant school art project: 10,000 students painting a massive mural. You can't just let everyone swarm the wall—chaos! Instead:

Split students into teams of 30 (thread blocks). Each team gets a section of the wall and shares paint buckets (shared memory) at their table.
Teams can talk to each other ("Hey, is everyone done with the sky?") but only within their own table—shouting across the gym is too slow.
The gym has 10 workstations (SMs), so 10 teams work simultaneously. As teams finish, the teacher assigns new sections to those workstations.
The grid is the entire mural divided into sections, one per team.

GPUs work the same way: millions of tiny workers (threads) grouped into cooperative teams (blocks) that paint different parts of the canvas (grid) without neding to coordinate with distant teams. This is how your GPU renders a game frame with billions of pixels in milliseconds!

Connections

CUDA Programming Model – syntax and API for launching grids/blocks
Streaming Multiprocessors – hardware that executes thread blocks
Warp Execution – how threads within a block run in SIMT groups
Shared Memory – fast on-chip storage for intra-block communication
Thread Synchronization – __syncthreads() and cooperative groups
GPU Occupancy – maximizing active threads per SM
Memory Coalescing – how thread indexing affects memory bandwidth
Parallel Algorithm Design – decomposing problems into grids of blocks

#flashcards/hardware

What are the two main levels in the GPU execution hierarchy? :: Thread blocks and grids. Threads are grouped into blocks that share resources, and blocks are organized into a grid covering the problem space.

Why can't threads in different blocks synchronize within a kernel?

Blocks execute independently and may run in any order on any SM. This independence enables scalability across GPUs with different SM counts.

What is the formula for a thread's global ID in1D?

globalID = blockIdx.x * blockDim.x + threadIdx.x. This accounts for all threads in previous blocks plus position in current block.

Why do we use ceiling division when calculating grid size?

To ensure enough blocks cover all data elements. Formula (N + B - 1) / B rounds up, so the last block handles remaining elements even if N isn't divisible by block size.

What are the built-in CUDA variables for thread identification?

threadIdx (position in block), blockIdx (block position in grid), blockDim (block size), gridDim (grid size). Each has x, y, z components.

Why is a 32-multiple block size recommended?

Threads execute in warps of 32. Non-multiples waste warp slots and reduce efficiency due to the SIMT execution model.

What limits how many blocks can run on an SM simultaneously?

Hardware limits on total threads per SM, registers per thread, shared memory per block, and maximum concurrent blocks. The tightest constraint determines occupancy.

Why can't you assume block 0 executes before block 1?

GPU runtime schedules blocks independently to any available SM in arbitrary order for load balancing and scalability.

What happens to extra threads when N isn't divisible by block size? :: The last block has excess threads. Use an if-guard if (i < N) to prevent out-of-bounds memory access.

Why use 2D blocks for matrix operations?

Natural mapping: each thread computes one output element at (row, col). Matches problem geometry and simplifies indexing logic.

What is GPU occupancy?

Ratio of active warps to maximum possible warps on an SM. Higher occupancy helps hide memory latency through paralelism.

Why are thread blocks called "cooperative thread arrays"?

Threads in a block can cooperate through shared memory and synchronization barriers (__syncthreads()), unlike threads in different blocks.

Concept Map

Hinglish (regional understanding)

Intuition Hinglish mein samjho

GPU programming mein ek fundamental concept hai thread blocks aur grids. Socho ki tumhe ek bahut bada kaam karna hai—jaise 1 million numbers ko add karna.Ek single thread se yeh kaam bahut slow hoga, toh GPU hume lakhs threads ek sath chalane ki power deta hai.

Lekin itne threads ko manage karna mushkil hai, toh GPU ek two-level hierarchy use karta hai. Pehle level pe hum threads ko blocks mein group karte hain (typically 128-256 threads). Ek block ke andar ke threads ek dusre se fast shared memory ke through baat kar sakte hain aur __syncthreads() se synchronize ho sakte hain—matlab sab ek sath checkpoint pe pohochte hain. Yeh threads physically bhi pas hote hain, same Streaming Multiprocessor (SM) pe run karte hain.

Dusre level pe hum in blocks ko grid mein arrange karte hain. Grid1D, 2D, ya 3D ho sakta hai—tumhare problem ki geometry ke hisaab se. Jaise image processing ke liye 2D grid natural lagta hai. Har block independently run hota hai aur GPU scheduler inhe dynamically different SMs pe assign karta hai. Yeh design scalable hai—same code10 SMs wale laptop GPU pe bhi chalega aur 100 SMs wale datacenter GPU pe bhi.

Thread indexing ka formula simple hai: globalID = blockIdx.x * blockDim.x + threadIdx.x. Matlab pehle dekho kitne complete blocks pehle aye (har block mein blockDim threads), phir current block ke andar apna position add karo. Yeh unique ID har thread ko apna kaam assign karne mein help karti hai—like array ke kaunse element pe kaam karna hai. Common mistake yeh hai ki log blocks ko sequential samajhte hain (block 0 pehle, phir block 1), lekin GPU inhe any order mein run kar sakta hai. Isliye blocks ke bech coordination nahi ho sakta ek kernel andar—cross-block synchronization ke liye separate kernel launch chahiye.

The Two-Level Hierarchy

Why This Design?

Deriving Global Thread IDs

Block Scheduling

Connections

Concept Map

Hinglish (regional understanding)

Go deeper — visual, from zero

Test yourself — GPU Architecture

Connections