Thread Block

A thread block is formally known as a Cooperative Thread Array (CTA) in hardware documentation. It is a logical grouping of threads (e.g., 32 to 1024 threads) that is mapped onto a single Stream Multiprocessor (SM) or NPU core for execution. This grouping is the smallest unit that can be independently scheduled by the hardware. The threads within a CTA can efficiently communicate via fast on-chip shared memory and synchronize using barrier primitives, enabling fine-grained data sharing and collaboration that is impossible between threads in different blocks.

The thread block is the granularity at which work is assigned to hardware execution resources. A Stream Multiprocessor (SM) or NPU core is allocated one or more thread blocks based on its resource limits (registers, shared memory). The hardware scheduler then manages the execution of these blocks. Threads within a block are further grouped into smaller units called warps (typically 32 threads) or wavefronts, which are the units of SIMT (Single Instruction, Multiple Threads) execution. This two-level hierarchy—blocks for scheduling and warps for instruction issue—is central to GPU/NPU architecture.

A defining feature of a thread block is its shared memory (also called scratchpad memory or local data share). This is a small, programmer-managed, high-bandwidth memory pool (typically 16-64 KB) that is private to the block. All threads within the block can read from and write to this memory with extremely low latency.

• Barrier Synchronization: Threads can call __syncthreads() (CUDA) or equivalent to ensure all threads in the block reach that point before any proceed.
• Scope Limitation: Threads in different blocks cannot directly synchronize or communicate via this shared memory, enforcing a clear boundary for data exchange and requiring coordination at a higher level (e.g., via global memory).

A thread block is defined with a 1D, 2D, or 3D grid of threads. The programmer specifies its dimensions (e.g., (256, 1, 1) or (16, 16, 1)). This logical structure aids in mapping computational patterns like matrix operations or image processing.

• Each thread has a unique thread index (threadIdx.x, .y, .z) within its block.
• Combined with the block's block index within the larger grid, each thread can compute a global index to determine which piece of data it operates on.
• The total number of threads per block is a key tuning parameter, balancing occupancy (hardware utilization) with available per-thread resources like registers and shared memory.

Thread blocks are designed for massive parallelism and scalability. A fundamental principle is that blocks execute independently and in any order. They can be scheduled on any available SM, and in any sequence (concurrently or serially). This independence guarantees that a program will execute correctly regardless of the number of SMs, making the model scalable across different hardware generations. Communication between blocks must occur via global memory and is typically orchestrated by completing one kernel launch (comprising all blocks) before launching another.

The number of thread blocks that can reside concurrently on an SM is limited by hardware resources, a key concept for performance optimization (occupancy). The main limiting resources are:

• Registers: Each thread consumes registers. A block's total register usage limits how many blocks can be active.
• Shared Memory: Each block allocates a portion of the SM's shared memory.
• Thread Slots: Each SM has a maximum number of threads (e.g., 2048) it can manage.

Optimizing a kernel involves choosing block dimensions and resource usage to maximize occupancy, keeping the hardware saturated with warps to hide memory and instruction latency.