Ch.8 ‐ Occupancy Tuning, Warp Efficiency, ILP - rFronteddu/general_wiki GitHub Wiki

if your kernel’s ALU utilization is low while its memory throughput is at 95% of peak, you are almost certainly memory-bandwidth bound.

if ALU utilization is near its maximum but memory throughput remains modest, the kernel is compute bound: gain speed only by increasing arithmetic throughput—typically by switching to lower-precision types (FP16, FP8, or FP4) and moving work onto the faster Tensor Cores of the GPUs

If both ALU utilization and memory throughput are low, your kernel may be experi‐ encing long-latency operations, synchronization overhead, or simply insufficient par‐ allel work. This could indicate low instruction-level parallelism or that you haven’t launched enough threads to fully utilize the GPU

You can frame this analysis using the Roofline model, which plots a kernel’s FLOPS against its arithmetic intensity (FLOPS per byte of memory accessed). The roofline defines a compute roof (the maximum FLOPS the GPU can sustain) and a memory roof (the maximum memory bandwidth).

If your kernel’s FLOP/byte ratio falls below the hardware’s compute-to-memory ratio, you are memory bound because you cannot supply data fast enough. If the ratio is high but actual FLOPS remains far below peak, the kernel may be latency bound or lacking sufficient ILP to saturate the compute units

In practice, always compare two key figures: your kernel’s memory throughput versus the hardware’s peak memory bandwidth—as well as your kernel’s compute through‐ put versus the hardware’s peak FLOPS. Those comparisons will tell you whether your next optimization should focus on memory access, compute work, or parallelism. Let’s look at each of these next.

You can use Nsight Compute’s memory chart to see how much traffic goes to L2 ver‐ sus DRAM. A high L2 hit rate could alleviate global memory bottlenecks, which means the kernel might actually be compute-limited despite performing heavy mem‐ ory accesses. The on-chip cache is servicing a lot of the memory accesses

On Blackwell, you can control L2 data persistence to keep critical working sets resident

Kernel Compute Throughput Versus Peak GPU FLOPS

You can use Nsight Compute’s Occupancy section and Source Counters to pinpoint these issues. Ensure the kernel launches enough threads to fill the GPU, up to the perSM resident warps limit for your device (e.g., 64 resident warps per SM.)

, examine instruction issue efficiency and throughput metrics. Blackwell scales to many SMs per device, so underutilization at the kernel level can translate to large aggregate losses.

If the achieved FLOPS are moderate to high but memory throughput is low, the ker‐ nel might be mostly compute-focused but limited by instruction dependencies. You can confirm by checking the “Exec Dependency” stalls metric in Nsight Compute— and any other stall reasons. If compute throughput is near peak, then you are truly compute bound. In this case, you have already optimized the kernel’s memory access patterns—or you are already using lower-precision Tensor Cores to reach such high FLOPS.

Prolonged near-peak compute utilization can sometimes invoke power management compute limiters on modern GPUs to keep them healthy. Be sure to take this into account when you’re profiling, benchmarking, and tuning. The easiest way to see if you’re being power-limited is to use the following to monitor the enforced power limit alongside any “HW Slowdown” flags in real time:

nvidia-smi \
--query-gpu=\
power.draw,clocks.current.sm,clocks.current.memory,\
clocks_event_reasons.active \
--format=csv -l 1

This command prints a new line every 1 second with current power draw, graphics/ memory clocks, and throttle reasons. With this information you can pinpoint when the GPU hits its power cap and downclocks.

You can also use NVIDIA’s Management Library (NVML) API, which provides pro‐ grammatic access to the CUDA C++ nvmlDeviceGetPowerUsage() and nvmlDevice GetEnforcedPowerLimit() APIs—as well as the equivalent NVML Python APIs. These are ideal for custom scripts or integration with monitoring systems.

Make sure you see high occupancy, high warp efficiency, and a balanced usage of execution units, including Tensor Cores

Iteratively Profiling and Determining the Kernel Bottleneck

GPUs can stall for four fundamentally different reasons: underutilization, latency bound, memory bound, and compute bound. Often, fixing one bottleneck will reveal another.

Underutilization happens when you simply haven’t launched enough threads or work. In this case, both FLOPS and memory bandwidth stay low and the execution timeline has idle gaps. Once you increase parallelism, you will find your warps are now stalling and waiting on memory loads

Once you more fully utilize your GPU, you can now distinguish between latency bound and memory bound. A latency-bound kernel issues far fewer bytes/sec than the hardware can deliver because individual memory loads are stalling the warps. The fix is to increase memory-compute overlap by increasing occupancy, using more ILP, prefetching, and pipelining.

A memory-bound kernel, in contrast, is saturating DRAM bandwidth, but your ALUs are sitting idle, not because of stalls but because there’s simply no more data you can fetch per second due to the memory pipe saturation. In this case, you must raise arithmetic intensity with tiling, fusing, exploiting caches (L1/texture), or reducing precision to reduce memory traffic.

If neither of those fixes produces more speed, you’ve moved into compute-bound territory in which the GPU’s arithmetic pipes (e.g., ALUs and Tensor Cores) are the lim‐ iting factor. Here you can increase per-thread ILP by overlapping independent instructions with unrolling and software pipelining. On modern GPUs, unified cores cannot execute an INT32 and an FP32 instruction in the same clock.

As you optimize, you’ll often find yourself working through these regimes as follows:

underutilized → latency bound → memory bound → compute bound.

When optimizing GPU code, you should follow a structured approach as described here. First, you should profile. Next, you can identify the bottleneck. You can use Nsight Compute for kernel-level metrics (e.g., warp stalls, achieved occupancy, memory ver‐ sus compute utilization) and Nsight Systems for application-level timelines (e.g., con‐ currency, idle gaps).

Once you identify the bottlenecks, you can determine if the kernel is memory bound, compute bound, latency bound, or simply underutilizing the GPU. The GPU is underutilized when it is not issuing enough work.

A memory-bound kernel is one where performance is limited by memory throughput—specifically, if the GPU’s global memory is unable to feed data to the compute units fast enough. In this case, your kernel’s achieved FLOPS will sit near the roofline set by memory bandwidth. This happens if you have plenty of threads but can’t move data any faster. As such, you’re up against the memory-bandwidth ceiling.

Conversely, a compute-bound kernel saturates the GPU’s arithmetic units (ALU FP32 CUDA cores or reduced-precision Tensor Cores). This is noticeable if the kernel is approaching the peak FLOPS roofline for the cores. Profiling metrics like achieved occupancy, memory utilization, and execution dependency stalls can help confirm this classification

When a kernel is latency bound, on the other hand, each thread spends a large frac‐ tion of its time waiting on individual memory loads instead of doing useful work. In practical terms, this means that when a warp issues a global‐memory load to fetch A[idx], for instance, all 32 threads in that warp will stall until that fetch completes.

the code immediately issues another dependent load or computation, the warp simply sits idle for hundreds of cycles on each load. The GPU’s warp scheduler may switch to other warps when it is latency bound, but if every warp is structured the same way (e.g., one load → wait → compute → write), there is rarely any other work to fill in those idle cycles. As such, the kernel never has enough independent operations in flight to hide the latency of a long‐latency DRAM access. To break out of the latency‐bound situation, you need to give the GPU multiple oper‐ ations to overlap. One approach is to increase occupancy by launching more threads and warps. This way, when one warp stalls on its load, another warp is ready to run. Equally important, though, is increasing each warp’s ILP.

Optimizing the Kernel

In practice, performance tuning should follow a clear, step‐by‐step workflow. First, identify which regime your kernel occupies: memory bound, latency bound, compute bound, or underutilized. Next, apply the corresponding optimizations. If your kernel is memory bound, concentrate on reducing and hiding memory traffic. You can do this by improving coalescing, raising occupancy, increasing ILP, and introducing data reuse through caching or tiling.

When the kernel is latency bound, meaning individual load or instruction latencies dominate, make sure there is enough independent work in flight. You can do this by issuing multiple nondependent loads/operations per thread or increasing overall occupancy so the scheduler always has ready warps to run. If the kernel is compute‐bound (i.e., the ALUs or Tensor Cores are saturated while memory is idle), shifting to lower‐precision arithmetic (FP16/FP8/FP4), offloading work onto Tensor Cores, or fusing more operations together can raise arithmetic throughput. Finally, if the GPU appears underutilized with low occupancy and fre‐ quent idle cycles, simply launching more threads or blocks (so that all SMs have work) is often enough to get the hardware busy before applying deeper optimizations. Here is a list of high-level optimization techniques that help you treat GPU perfor‐ mance tuning as a scientific process.

Convert memory-bound workloads to compute bound If memory bound (low arithmetic intensity), increase data reuse and work per launch as follows: apply tiling to use fast shared memory and reduce redundant accesses, fuse kernels to avoid unnecessary memory round trips, ensure memory accesses are optimized (coalesced, avoiding bank conflicts as per Chapter 6), and consider using compression or lower precision to move less data.

Further optimize compute-bound workloads If compute bound (e.g., high utilization of ALUs but not hitting peak due to dependencies), increase effective instruction throughput as follows: use ILP tech‐ niques (e.g., unroll loops, multiple accumulators to overlap independent ops), check for branch divergence (see Chapter 6) and try to reorganize work to reduce it since divergence wastes ALU cycles, and move to Tensor Cores and lower pre‐ cision to raise the compute ceiling. If the kernel is at compute roof, see if reduc‐ ing precision (FP32 → FP16 → FP8 → FP4) can give further speedups.

Increase parallelism for latency-bound workloads If latency bound (e.g., frequent warp stalls and not enough parallelism to hide latency), increase concurrency at various levels as follows: launch more threads/ blocks if possible until latency is hidden, ensure registers/shared memory are not overly limiting occupancy (e.g., occupancy tuning and balancing resource usage), overlap memory and compute within each thread/warp using async copies (e.g., intra-kernel pipelining), and use multiple streams to overlap independent tasks or overlap copies with compute (e.g., inter-kernel concurrency described in Chapter 11). If launch overhead is an issue (e.g., many tiny kernels), consider merging them with cooperative groups or simply combining their code (e.g., ker‐ nel fusion).

Increase GPU utilization If the GPU is underutilized (e.g., low SM Active, low occupancy), ensure you launch enough work to use all SMs. Specifically, make sure your kernel is launch‐ ing with a grid size large enough to fully utilize the GPU. A common mistake is launching as many threads as elements but forgetting that each thread does only a little bit of work. Sometimes you need multiple passes or more threads per element.

Reduce synchronizations and host-side (CPU) stalls You should remove any unnecessary cudaDeviceSynchronize() or host-side waits that stall the GPU. If the workload is inherently small, consider batching it with other work or running multiple instances concurrently (e.g., CUDA streams). You can also use CUDA Graphs (see Chapter 12) to predefine and efficiently launch an execution graph of many small kernels.

Leverage specialized hardware and reduced/mixed precision Blackwell Tensor Cores expose fifth-generation MMA instructions in PTX as tcgen05.mma and associated loads/stores (e.g., tcgen05.ld and tcgen05.st). Tensor Cores accelerate microscaling formats including MXFP8, MXFP4 (OCP MX formats), and NVIDIA’s NVFP4 format. They also support block-scaled mat‐ muls (K-grouped) that libraries select automatically when scale metadata is present. TMEM and TMA underpin these high-throughput data paths. We’ll dis‐ cuss these techniques in more detail later in this chapter and in Chapter 9.

Verify and iterate After each optimization, reprofile. Confirm the targeted stall or metric improved (e.g., memory stalls reduced after tiling, achieved occupancy went up after tuning block size, “SM Throughput %” increased after using CUDA streams, etc.). Also watch total runtime improvement. Sometimes one bottleneck masks another; you might fix memory bandwidth only to become compute bound next (which is fine—then address that if needed).

Maintain correctness and acceptable accuracy When using lower-precision or new parallel strategies, test with assertions or comparisons to reference results. Ensure the speedup doesn’t come at the cost of accuracy unless that’s acceptable for the application. Typically, techniques like FP16 or even FP8 are carefully validated to have negligible accuracy impact on AI models.

Tuning Occupancy

occupancy is the ratio of active warps on the SM to the maximum number of warps that could be active on the SM. Low occupancy (< 50%) means that, on average, half of the possible warps were active. This might indicate that your kernel is limited by resources such as registers or shared memory per thread block—rather than just available parallelism. Occupancy is the measure of how many threads, or warps, are active on an SM rela‐ tive to the hardware’s maximum capacity. Higher occupancy, or more warps in flight, allows the GPU to better hide latency. his is because when one warp stalls waiting on a memory load, for instance, the scheduler can quickly switch to run another warp. More warps per SM generally means the GPU’s pipelines stay busier, and fewer cycles are wasted waiting on memory. Occupancy tuning is the practice of adjusting your kernel launch parameters and resource usage

Remember that the goal of occupancy tuning is to keep enough warps active to fully utilize the SM’s pipelines and hide long‐latency operations. In an ideal case, you would achieve 100% occupancy, filling all available warp slots.

This per-SM limit of 64 warps has remained the same for modern datacenter GPU architectures like Ampere, Hopper, and Blackwell— even as the overall GPU core counts have increased. The hardwareperformance improvements come from more SMs, larger caches, multidie, etc. You can use Nsight Compute’s Occupancy analysis to confirm the exact limits on your target device.

a compute-bound kernel might achieve peak performance with only 50% occupancy because each warp is doing lots of work without waiting. In contrast, a memory-bound kernel often benefits from high occupancy since some warps can run on the SM since other warps are stalled waiting for memory transfers.

Find the Right Occupancy for Your Workload

In practice, effective occupancy tuning often produces diminishing returns after a certain point. If a kernel is severely memory bound, for instance, going from 10% to 50% achieved occupancy might give a huge boost because now you have enough warps to cover latency. But going from 50% to 100% might give only a small further gain since other factors start to dominate, such as cache misses, memory bandwidth saturation, etc. Profiling helps determine the optimal occupancy. For example, you can evaluate profile metrics such as eligible warps per cycle and active warps per scheduler. These give insight into how many warps are ready to issue versus how many the hardware could handle.

Eligible warps per cycle reports the average number of warps that are in a “ready-torun” state each cycle and have no outstanding data or dependency stalls. Active warps per scheduler, often equal to the number of schedulers on the SM, is the maximum number of warps that could issue an instruction per cycle. If eligible warps per cycle is lower than active warps per scheduler, the GPU often runs out of ready warps. In this case, when one warp stalls on memory or a longlatency instruction, there isn’t another ready warp to switch to. This indicates you need more concurrency in the form of higher occupancy or ILP to hide the latenc

In contrast, if eligible warps per cycle meets or exceeds the scheduler limit but your kernel still runs slowly, it means you have enough warps ready, but they cannot issue because of other stalls such as memory-bandwidth saturation or execution dependencies. In this case, it’s best to focus on hiding memory latency through better coalescing and asynchronous copies—or increasing ILP by unrolling independent work. This is a better approach than simply adding more threads

And if you raise occupancy and see the “Stall: Not Selected” or idle percentages drop, and memory pipes are busy more often, you’ve successfully improved occupancy. If occupancy is high but Long Scoreboard is still dominant, you might need other tech‐ niques like improving memory access patterns or overlapping computation. In practice, once you reach a moderate occupancy (e.g., 60%–70%), the returns will start to diminish. As such, it’s often more effective to pursue better memory locality, higher ILP, and the use of on-chip memory like shared memory and registers to cache data. While maximizing occupancy ensures that many warps are available to run, those warps might still be idle if waiting on memory or if executing divergent code. There‐ fore, after achieving a reasonable occupancy (e.g., 50%–70%), focus on warp effi‐ ciency and latency-hiding rather than obsessing over 100% occupancy.

Occupancy is a means to an end (hiding latency), not the end goal itself. Once you have enough warps to keep the GPU busy, other optimizations will give better returns.