Kernel Occupancy

Each thread block requires a fixed number of registers per thread. The total register allocation for all active warps cannot exceed the streaming multiprocessor's (SM) physical register file capacity. High register usage per thread directly reduces the number of concurrent warps that can be resident.

• Example: An SM with 65,536 32-bit registers can support a maximum of 64 warps if each warp requires 1024 registers (32 threads * 32 registers/thread).
• Compiler Impact: Aggressive optimizations like loop unrolling can increase register pressure, potentially lowering occupancy. Techniques like register spilling to local memory are used when limits are exceeded, at a performance cost.

Shared memory is a software-managed, on-chip cache shared by all threads in a thread block. Its size is finite per SM (e.g., 64KB or 96KB). The total shared memory requested by all resident thread blocks must fit within this limit.

• Static vs. Dynamic: Allocation can be static (compile-time known) or dynamic (runtime variable). Dynamic allocation reduces flexibility for the scheduler.
• Trade-off: Using more shared memory per block for data reuse or communication reduces the number of blocks that can co-reside, potentially lowering occupancy but improving per-thread performance.

The dimensions of a thread block (blockDim) determine how warps are formed and impact several limits.

• Warp Count per Block: A block is composed of an integer number of warps. A block size not a multiple of the warp size (e.g., 33 threads) wastes resources, as the last warp is underpopulated but still consumes a full warp slot.
• Hardware Limits: Architectures impose a maximum number of threads per block (e.g., 1024) and a maximum number of blocks per SM. Small blocks may hit the block limit before other resources are exhausted, while very large blocks may be limited by registers or shared memory first.

The number of warp schedulers per SM and the total number of warps they can track is a fundamental hardware limit, often expressed as a 'Maximum Warps per SM' or 'Maximum Threads per SM' specification.

• Active Warps vs. Eligible Warps: Not all resident warps are ready to issue an instruction on every cycle. Higher occupancy increases the scheduler's pool of eligible warps, improving its ability to hide instruction and memory latency by keeping execution units busy.
• Architectural Variance: This limit varies significantly between GPU/NPU architectures (e.g., 64 warps/SM on some, 32 on others).

Additional finite resources per SM can become the limiting factor.

• Thread Slots: The maximum concurrent threads per SM is a hard limit (e.g., 2048). Occupancy cannot exceed (Active Threads) / (Max Threads).
• Barrier Count: The number of barrier synchronization primitives available per SM limits blocks with complex synchronization patterns.
• Constant Cache / L1 Pressure: While not a direct occupancy limit, high pressure on other cache partitions can indirectly affect performance, making a lower-occupancy, cache-friendly configuration optimal.

Register pressure is a primary limiter of occupancy. Each thread's declared variables consume registers, and exceeding the hardware limit forces threads to be serialized. Key strategies include:

• Use smaller data types (e.g., float16 instead of float32) where precision permits.
• Limit variable scope to reduce the number of live values requiring registers simultaneously.
• Employ compiler directives (e.g., __launch_bounds__) to hint at maximum thread block size, guiding the register allocator.
• Enable compiler optimizations like common subexpression elimination (CSE) and constant propagation to reduce redundant calculations and register demands. Excessive register usage can lead to register spilling, where values are moved to slower local memory, creating a double penalty of lower occupancy and higher latency.

Shared memory is a fast, software-managed cache on-chip, but its capacity is fixed per streaming multiprocessor (SM). Inefficient use directly caps the number of concurrent thread blocks. Techniques involve:

• Precisely size shared memory arrays to the minimum required for the tile (e.g., for a 32x32 tile, allocate float tile[32][32], not [64][64]).
• Implement bank conflict-free access patterns to avoid serialization of accesses within a warp.
• Dynamically allocate shared memory when possible, allowing the runtime to pack thread blocks more flexibly.
• Reuse shared memory buffers for multiple stages of a kernel's computation to reduce the total footprint. The goal is to fit more thread blocks per SM by minimizing the shared memory each block requires.

The dimensions of a thread block (blockDim) directly impact how resources are allocated and how many blocks can reside on an SM. Optimization involves:

• Choosing a multiple of the warp size (typically 32) for the x-dimension to ensure full warp utilization and efficient memory coalescing.
• Balancing block size; larger blocks use more registers and shared memory per block, potentially reducing the number of concurrent blocks. Smaller blocks may increase occupancy but can reduce per-thread work and increase global scheduling overhead.
• Experimenting with 1D, 2D, or 3D block shapes to best match the data access patterns of the algorithm. The optimal configuration is found through auto-tuning, empirically searching the parameter space (e.g., block sizes of 128, 256, 512) for the hardware target.

Kernel fusion merges multiple sequential kernels into one, which can dramatically improve effective occupancy by eliminating intermediate global memory transfers. Benefits include:

• Reduced global memory traffic: Intermediate results stay in registers or shared memory instead of being written to and read from DRAM.
• Fewer kernel launches: Launch latency and overhead are amortized over more work.
• Increased arithmetic intensity: The fused kernel performs more operations per byte fetched from global memory, moving the workload closer to being compute-bound. For example, fusing a convolution, bias addition, and ReLU activation into a single operation fusion kernel keeps all intermediate data on-chip, freeing up memory bandwidth and reducing pressure on occupancy limits from memory operations.

Modern compilers for accelerators provide directives and flags to influence resource allocation and occupancy.

• __launch_bounds__(maxThreadsPerBlock, minBlocksPerMultiprocessor): This CUDA directive informs the compiler of the intended launch configuration, allowing it to optimize register usage to hit the specified occupancy target.
• -maxrregcount flag: Sets a hard upper limit on the number of registers used per thread, forcing spilling but allowing more threads to be active concurrently—a trade-off to increase occupancy for memory-bound kernels.
• Loop unrolling and software pipelining: Controlled via pragmas (e.g., #pragma unroll), these transformations can reduce loop overhead and hide latency, sometimes allowing for a reduction in the number of live registers, thereby improving occupancy.

Systematic improvement requires measurement. Vendor tools provide essential occupancy metrics and bottleneck analysis.

• Nsight Compute (NVIDIA) & ROCProf (AMD): These profilers report achieved occupancy, the limiting factor (register, shared memory, or thread block limit), and detailed resource usage per kernel.
• Occupancy Calculator: Many SDKs provide spreadsheets or APIs to model occupancy based on kernel resource requirements (threads, registers, shared memory) and specific hardware specs.
• The Roofline Model: This analytical model helps classify the kernel as compute-bound or memory-bound. If memory-bound, increasing occupancy (to hide memory latency) is a high-priority optimization. If compute-bound, optimizing for instruction-level parallelism (ILP) or using tensor cores may be more beneficial than maximizing occupancy alone.

Warp divergence occurs in SIMT architectures when threads within a single warp (a group of 32 threads on NVIDIA GPUs) take different execution paths due to conditional branching (e.g., if/else statements). This forces the hardware to serialize the execution of all divergent paths, severely reducing the effective occupancy and parallel efficiency of the streaming multiprocessor. Minimizing warp divergence through data reorganization or algorithmic changes is essential for maintaining high occupancy and performance.

Register spilling is a critical compiler-level event that directly limits kernel occupancy. It occurs when the live variables in a kernel require more architectural registers than are physically available on the hardware. The compiler is forced to 'spill' these values to much slower memory (like local or global memory). This not only increases latency but also consumes shared memory or L1 cache resources, which are part of the occupancy limit calculation. High register usage is a primary bottleneck for achieving maximum occupancy.

Memory coalescing is an optimization where concurrent memory accesses from threads in a warp are combined into a single, wider transaction. While it primarily maximizes memory bandwidth, it has a secondary impact on occupancy. Efficient coalescing reduces memory transaction stalls, allowing warps to complete execution and free resources more quickly. This improves the throughput of warp scheduling, effectively allowing the hardware to sustain a higher level of active occupancy even if the theoretical limit is unchanged.

Kernel fusion is a compiler optimization that merges multiple separate computational kernels into a single, larger kernel. This directly impacts occupancy dynamics by:

• Eliminating intermediate kernel launch overhead and global memory writes/reads.
• Increasing the arithmetic intensity of the fused kernel, potentially shifting the performance bottleneck.
• Changing the resource profile (shared memory, registers) of the combined kernel, which must be recalculated against occupancy limits. Fusion is a key technique for moving from memory-bound to compute-bound kernels.

A thread block is a programmer-defined group of threads that can cooperate via shared memory and synchronization. The dimensions of the thread block (block size) are a primary lever for controlling kernel occupancy. Key considerations include:

• Block size must be a multiple of the warp size.
• Larger blocks may improve latency hiding but increase per-block register and shared memory usage, potentially lowering the number of concurrent blocks.
• The compiler and hardware scheduler determine how many blocks can reside on a streaming multiprocessor based on their resource consumption versus the hardware's limits.

The Roofline Model is an analytical performance model that provides context for the importance of occupancy. It plots attainable performance (GFLOPs/sec) against arithmetic intensity (ops/byte). Occupancy is a key factor in reaching the model's 'rooflines':

• Memory-Bound Region: Performance is limited by memory bandwidth. High occupancy is crucial here to hide the long latency of memory accesses by keeping the execution units busy with other warps.
• Compute-Bound Region: Performance is limited by peak compute throughput. While high occupancy is still beneficial, optimizing instruction-level parallelism and using Tensor Cores may become more critical than simply maximizing warp count.