Occupancy

Occupancy quantifies the utilization of a processor's parallel execution resources. It is calculated as:

Occupancy = (Active Warps per SM) / (Maximum Warps per SM)

• Active Warps: The number of warps currently executing or ready to execute on a Streaming Multiprocessor (SM).
• Maximum Warps: The theoretical limit of warps an SM can hold, determined by its hardware resources (registers, shared memory).

High occupancy ensures the hardware scheduler has ample work to hide instruction and memory latency, but it is not the sole determinant of peak performance.

Occupancy is constrained by the scarcest of three primary hardware resources on an SM:

• Thread Block Size: The number of threads per block. Larger blocks require more registers and shared memory, potentially reducing the number of concurrent blocks.
• Register Usage per Thread: Each thread allocates a set of registers. High register usage limits the number of active threads, as the SM's total register file is fixed.
• Shared Memory Usage per Block: Each block allocates shared memory. High usage limits the number of concurrent blocks an SM can host.

The occupancy calculator, often provided in vendor SDKs, models these constraints to predict the occupancy for a given kernel configuration.

The primary purpose of high occupancy is to enable latency hiding. When one warp stalls (e.g., waiting for a global memory access), the hardware scheduler can immediately switch to another ready-to-execute warp, keeping the compute units busy.

• Low Occupancy: Few active warps. Stalls are poorly hidden, leading to underutilization of ALUs and reduced throughput.
• High Occupancy: Many active warps. The scheduler has a large pool of warps to choose from, effectively masking stalls and improving hardware utilization.

This is why occupancy is often described as providing 'thread-level parallelism' to hide latency.

High occupancy does not guarantee high performance. It is a necessary but not sufficient condition. Other critical factors include:

• Memory Coalescing: Poor access patterns can make a kernel memory-bound, stalling warps regardless of occupancy.
• Instruction-Level Parallelism (ILP): Kernels with dependent instructions have low ILP, causing execution dependencies that occupancy cannot hide.
• Divergent Execution: Thread divergence within a warp causes serialization, reducing effective throughput.

An optimal kernel balances occupancy with compute throughput and memory bandwidth efficiency. Auto-tuners often treat occupancy as one of several optimization objectives.

Occupancy is measured using hardware performance counters and visualized in kernel profiler tools (e.g., NVIDIA Nsight Compute, AMD ROCprofiler). Key profiling steps:

1. Collect Metrics: Tools report achieved_occupancy (average active warps / max warps) and theoretical_occupancy (limit based on resource usage).
2. Analyze Limits: The profiler identifies the primary limiting factor (registers, shared memory, or block size).
3. Cross-Reference: Correlate occupancy with other metrics like compute throughput (FLOPS) and memory bandwidth utilization to identify the true performance bottleneck.

This data drives auto-tuning to adjust kernel parameters for a better balance.

To improve occupancy, engineers adjust kernel launch parameters and resource usage:

• Adjust Workgroup/Thread Block Size: Experiment with sizes (e.g., 128, 256, 512 threads) to better fit the SM's resource limits.
• Limit Register Usage: Use compiler directives (e.g., __launch_bounds__, maxrregcount) to cap per-thread register allocation, allowing more concurrent threads.
• Optimize Shared Memory: Reduce shared memory footprint per block or design kernels to use configurable amounts, allowing more blocks per SM.
• Employ Concurrent Kernels: If a single kernel cannot fully utilize the chip, launch independent kernels to occupy unused SMs.

These strategies are often automated via auto-tuning frameworks that search the configuration space of these parameters.

A state where a kernel's execution time is limited by the speed of the processor's arithmetic logic units (ALUs), not by memory or I/O speeds. When a workload is compute-bound, the ALUs are saturated, and performance improvements come from increasing compute throughput (e.g., FLOPS). This is the ideal state for maximizing NPU efficiency, as it implies data is readily available for processing. Kernels with high arithmetic intensity (ratio of operations to bytes accessed) are typically compute-bound.

A state where a kernel's execution time is limited by the speed or bandwidth of the memory subsystem, causing the compute units to idle while waiting for data. This is a common bottleneck in NPU workloads. Key indicators include:

• Low occupancy due to threads waiting for memory requests.
• High latency for memory accesses.
• Performance that scales with improved memory bandwidth. Optimizations focus on improving data locality through memory coalescing, increasing cache hit rates, and using shared memory to reduce trips to global memory.

A condition in parallel computing where threads within the same warp or wavefront follow different execution paths due to conditional branching (e.g., if/else statements). This forces the hardware to serialize execution for each divergent path, drastically reducing effective occupancy and compute throughput. On NPUs and GPUs, all threads in a warp typically execute the same instruction. Minimizing divergence is critical for performance, often requiring algorithmic redesigns to ensure threads follow uniform control flow.

The number of threads grouped together for cooperative execution and synchronization within a single compute unit (e.g., a Streaming Multiprocessor on an NVIDIA GPU). This is a primary tunable parameter in auto-tuning. The optimal workgroup size balances:

• Occupancy: Maximizing resident warps to hide latency.
• Resource Limits: Staying within shared memory and register file constraints.
• Hardware Scheduling: Aligning with the warp/wavefront size of the target NPU. Selecting the wrong size can lead to underutilization and resource contention.

The automated process of searching a configuration space (defined by parameters like workgroup size, tile size, and unroll factor) to find the optimal setup for a given kernel on specific hardware. It directly addresses the challenge of maximizing metrics like occupancy and throughput. Methods include:

• Exhaustive Search: Testing all combinations (feasible only for small spaces).
• Bayesian Optimization: Using a probabilistic model to guide the search efficiently.
• Performance Model-guided search: Using analytical models to prune the search space. Tools like a Kernel Tuner automate this process.

A specialized hardware register that counts low-level microarchitectural events within a processor. These counters provide the raw data needed to calculate occupancy and diagnose performance issues. Common counters for occupancy analysis include:

• Active Warps/Cycles: The number of warps actively executing per cycle.
• Eligible Warps/Cycles: Warps that are ready to execute but may be stalled.
• Memory Transactions: Counts of loads/stores, helping identify memory-bound kernels.
• Stall Reasons: Counts of cycles stalled due to memory, execution dependency, or synchronization. Profilers aggregate these counters to create a performance profile.