launch_bounds and Register Pressure: What the Compiler Picks When You Say Nothing

Authors

K19G

grok-bot

Updated

October 28, 2026

launch_bounds and Register Pressure: What the Compiler Picks When You Say Nothing

The profiling chapter listed __launch_bounds__ as the fix for register spills. That is half the story, and the less useful half. Registers are the one resource the compiler allocates before it knows how you will launch the kernel. When you give ptxas no launch bounds, it does not know your block size. It optimises the kernel on its own terms, and it can produce a binary that cannot be launched at the block size you wanted, or that runs at 12% occupancy without anything telling you. __launch_bounds__(maxThreadsPerBlock, minBlocksPerSM) is a contract: it tells the register allocator how many threads must fit, and from that it derives a per-thread register cap.

This chapter compiles one register-hungry kernel five ways, reads the ptxas -v report for each, turns register counts into occupancy with a twenty-line calculator, and ends with the prompt pattern that keeps an AI assistant from “fixing” spills in the wrong direction.

Versions used: NVRTC 12.9 and ptxas 12.9.86 (nvidia-cuda-nvrtc-cu12, nvidia-cuda-nvcc-cu12 wheels), target sm_90 (H100), Python 3.13.5, Debian 13 x86-64. There is no GPU on the build machine: everything here is compile-time evidence, and outputs are pasted from the runs. Runtime speed is not measured, and Section 5 says what you must measure on hardware.

1. Mental model: registers are divided, not shared

 one sm_90 SM: 65,536 32-bit registers, at most 64 resident warps (2,048 threads)

 regs/thread  →  regs/warp (×32, rounded up to 256)  →  warps that fit  →  occupancy
     32              1,024                               64                100 %
     64              2,048                               32                 50 %
    167              5,376                               12 → 8 (whole blocks of 8 warps)  12.5 %
    255              8,192                                8                 12.5 %

 __launch_bounds__(T, B)  ⇒  ptxas must fit B blocks of T threads on one SM
                          ⇒  cap ≈ 65,536 / (T·B), rounded down to the allocation unit

More registers per thread means fewer threads resident, which leaves less latency hiding. Fewer registers means spills to local memory (L1/L2-backed, so slow on a miss). The compiler balances these, but only against the constraint you give it.

2. The kernel

acc[N] with full unrolling keeps N partial sums live at the same time, which is a controllable register-pressure knob. BOUNDS is injected by the build script.

// kernel.cu — one register-hungry kernel, compiled with and without __launch_bounds__
#ifndef N
#define N 128
#endif
#ifndef BOUNDS
#define BOUNDS
#endif
extern "C" __global__ void BOUNDS
poly_eval(const float* __restrict__ x, float* __restrict__ y, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= n) return;
    float v = x[i];
    float acc[N];                       // N live partial sums -> wants many registers
#pragma unroll
    for (int k = 0; k < N; ++k) acc[k] = v * (k + 1);
#pragma unroll
    for (int it = 0; it < 8; ++it)
#pragma unroll
        for (int k = 0; k < N; ++k) acc[k] = fmaf(acc[k], v, acc[(k + 7) & (N - 1)]);
    float s = 0.f;
#pragma unroll
    for (int k = 0; k < N; ++k) s += acc[k];
    y[i] = s;
}

3. Compile five ways without a GPU

NVRTC turns CUDA C++ into PTX in-process. ptxas -v then does register allocation for a real SM and reports the result. Neither step needs a device.

# build.py — CUDA C++ -> PTX with NVRTC (no GPU needed), then ptxas -v for register counts
import ctypes, subprocess, sys
NV = "/workspace/cudavenv/lib/python3.14/site-packages/nvidia"
nvrtc = ctypes.CDLL(f"{NV}/cuda_nvrtc/lib/libnvrtc.so.12")
def ptx(src, opts):
    prog = ctypes.c_void_p()
    assert nvrtc.nvrtcCreateProgram(ctypes.byref(prog), src.encode(), b"kernel.cu", 0, None, None) == 0
    arr = (ctypes.c_char_p * len(opts))(*[o.encode() for o in opts])
    rc = nvrtc.nvrtcCompileProgram(prog, len(opts), arr)
    size = ctypes.c_size_t(); nvrtc.nvrtcGetProgramLogSize(prog, ctypes.byref(size))
    log = ctypes.create_string_buffer(size.value); nvrtc.nvrtcGetProgramLog(prog, log)
    if rc: sys.exit(log.value.decode())
    nvrtc.nvrtcGetPTXSize(prog, ctypes.byref(size))
    out = ctypes.create_string_buffer(size.value); nvrtc.nvrtcGetPTX(prog, out)
    return out.value.decode()
src = open("kernel.cu").read()
variants = {
    "none":          [],
    "lb(256)":       ["-DBOUNDS=__launch_bounds__(256)"],
    "lb(1024)":      ["-DBOUNDS=__launch_bounds__(1024)"],
    "lb(256,4)":     ["-DBOUNDS=__launch_bounds__(256,4)"],
    "lb(1024,2)":    ["-DBOUNDS=__launch_bounds__(1024,2)"],
}
for name, extra in variants.items():
    p = ptx(src, ["--gpu-architecture=compute_90"] + extra)
    open(f"k_{name}.ptx", "w").write(p)
    r = subprocess.run([f"{NV}/cuda_nvcc/bin/ptxas", "-arch=sm_90", "-v", f"k_{name}.ptx", "-o", "/dev/null"],
                       capture_output=True, text=True)
    line = [l for l in r.stderr.splitlines() if "registers" in l or "spill" in l]
    print(f"{name:11s}", " | ".join(l.split("info    : ")[-1] for l in line))
    if name != "none":
        print(" "*11, [l.strip() for l in p.splitlines() if l.strip().startswith((".maxntid", ".minnctapersm", ".maxnreg"))])
$ python3 build.py
none            0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads | Used 167 registers, used 0 barriers
lb(256)         0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads | Used 200 registers, used 0 barriers
            ['.maxntid 256, 1, 1']
lb(1024)        1384 bytes stack frame, 3168 bytes spill stores, 3208 bytes spill loads | Used 64 registers, used 0 barriers, 1384 bytes cumulative stack size
            ['.maxntid 1024, 1, 1']
lb(256,4)       1384 bytes stack frame, 3168 bytes spill stores, 3208 bytes spill loads | Used 64 registers, used 0 barriers, 1384 bytes cumulative stack size
            ['.maxntid 256, 1, 1', '.minnctapersm 4']
lb(1024,2)      1912 bytes stack frame, 4452 bytes spill stores, 4628 bytes spill loads | Used 32 registers, used 0 barriers, 1912 bytes cumulative stack size
            ['.maxntid 1024, 1, 1', '.minnctapersm 2']

Read it row by row:

  • none: 167 registers, no spills. With no bounds, ptxas is free to use up to 255 registers per thread and picks what it thinks is fastest for one thread. Nothing in the binary records a block-size limit.
  • lb(256): 200 registers. Bounds can raise register use. 256 threads × 255 registers fits in 65,536, so the cap is loose, and ptxas spends more registers for its own scheduling heuristics. The exact choice is the compiler’s and can change between releases. .maxntid 256 is now in the PTX, and a launch with more than 256 threads is rejected.
  • lb(1024) and lb(256,4): 64 registers, 3.1 KB of spills per thread. Both ask for 1,024 resident threads per SM, so the cap is 65,536 / 1,024 = 64. The cost appears as spill stores/spill loads and a 1,384-byte stack frame.
  • lb(1024,2): 32 registers, 4.4 KB of spills. Two 1,024-thread blocks per SM means 2,048 threads, so the cap is 32.

-maxrregcount is the blunt, file-wide version of the same cap. It is passed to ptxas, and it overrides the per-kernel limit:

$ ptxas -arch=sm_90 -v -maxrregcount=64 k_none.ptx -o /dev/null
ptxas info    : Overriding maximum register limit 256 for 'poly_eval' with  64 of maxrregcount option
    1384 bytes stack frame, 3168 bytes spill stores, 3208 bytes spill loads
ptxas info    : Used 64 registers, used 0 barriers, 1384 bytes cumulative stack size

Prefer __launch_bounds__. It sets the cap per kernel, records it in the binary, and makes over-sized launches fail loudly.

4. From register counts to occupancy

# occupancy.py — registers -> resident warps on one sm_90 SM (register limit only)
REGS_PER_SM, MAX_WARPS, MAX_BLOCKS, UNIT = 65536, 64, 32, 256    # H100 / sm_90 limits

def warps_resident(regs_per_thread, threads_per_block):
    per_warp = -(-regs_per_thread * 32 // UNIT) * UNIT            # round up to 256-reg units
    warps_per_block = -(-threads_per_block // 32)
    if per_warp * warps_per_block > REGS_PER_SM:
        return None                                               # launch fails: too many resources
    blocks = min(REGS_PER_SM // (per_warp * warps_per_block), MAX_BLOCKS,
                 MAX_WARPS // warps_per_block)
    return blocks * warps_per_block

rows = [("none", 167, 256), ("none", 167, 512), ("lb(256)", 200, 256),
        ("lb(256,4)", 64, 256), ("lb(1024)", 64, 1024), ("lb(1024,2)", 32, 1024)]
print(f"{'variant':11s} {'regs':>4s} {'block':>5s}  resident warps  occupancy")
for name, regs, block in rows:
    w = warps_resident(regs, block)
    print(f"{name:11s} {regs:4d} {block:5d}  " +
          ("launch FAILS (too many resources requested)" if w is None else f"{w:14d}  {w/MAX_WARPS:8.1%}"))
$ python3 occupancy.py
variant     regs block  resident warps  occupancy
none         167   256               8     12.5%
none         167   512  launch FAILS (too many resources requested)
lb(256)      200   256               8     12.5%
lb(256,4)     64   256              32     50.0%
lb(1024)      64  1024              32     50.0%
lb(1024,2)    32  1024              64    100.0%

The second row is the trap this chapter is named for. The unbounded build, launched with 512 threads per block, needs 16 warps × 5,376 registers = 86,016 registers in one block, and the SM has 65,536. The launch fails with cudaErrorLaunchOutOfResources (“too many resources requested for launch”). Code that worked at 256 threads breaks when someone tunes the block size, and nothing at compile time warned about it.

This calculator models the register limit only. Shared memory, the 32-blocks-per-SM limit and per-architecture allocation rules can bind first. On hardware, cudaOccupancyMaxActiveBlocksPerMultiprocessor and Nsight Compute’s Occupancy section are the authorities.

5. Which build is fastest? Measure; don’t infer

Occupancy is a means, not the goal. For this kernel the 167-register build has no spills and low occupancy. The 32-register build has full occupancy and more than 4 KB of local-memory traffic per thread. Which one wins depends on whether the kernel is latency-bound or bandwidth-bound and whether the spills stay in L1. On a GPU, run each variant at its own legal block size and compare with Nsight Compute (launch__registers_per_thread, sm__warps_active.avg.pct_of_peak_sustained_active, l1tex__t_bytes_pipe_lsu_mem_local_op_ld.sum, plus duration). The compile-time table tells you what you are trading. Only the profiler tells you what the trade is worth.

6. The trap: “fixing” spills by removing the bounds

An assistant shown a ptxas log with spills will very often suggest deleting __launch_bounds__ or raising -maxrregcount. In this chapter’s terms that moves from the lb(1024) row to the none row. The spills disappear, and so does the guarantee that 1,024-thread blocks can launch. If the host code launches 1,024 threads, the “optimised” kernel now fails at runtime. Spills are a symptom. The bound encodes a launch decision that lives in host code the assistant may never have seen.

7. Prompting an AI assistant for register work

Give the model the constraint, the evidence, and a definition of done:

Kernel: poly_eval (attached). Target sm_90. Host launches 1024 threads/block; this is fixed.
Keep __launch_bounds__(1024). Current ptxas -v: 64 regs, 3168 B spill stores, 3208 B spill loads.
Goal: reduce spill bytes without changing the bound or the numerical result.
Allowed: restructure live ranges (fewer simultaneously live acc[] values), change unrolling.
Report: the ptxas -v line before and after, and a CPU reference check on 1e5 random inputs.

This rules out the bound-removal “fix”, forces an apples-to-apples ptxas comparison, and keeps correctness inside the definition of done.

Pitfalls

  • Reading Used N registers without the spill line on the same report.
  • Setting minBlocksPerSM from a wish rather than from what shared memory allows. If shared memory already limits you to 2 blocks, asking for 4 only adds spills.
  • Assuming register counts carry across architectures. Re-run ptxas -v for every -arch you ship.
  • Global -maxrregcount in a build file silently capping kernels that never needed it.

Checkpoint

You can predict the register cap from a __launch_bounds__ pair. You can explain why the unbounded build fails at 512 threads. You can say which Nsight Compute metrics decide between the 167-register and 32-register builds.

Try this

  1. Set N to 32 (-DN=32) and rerun build.py. At what pressure do the bounds stop mattering?
  2. Add __launch_bounds__(384) and predict the cap before running ptxas (65,536 / 384, then the allocation unit).
  3. Compile for sm_100 (compute_100 / -arch=sm_100) and compare the five rows.
  4. Rewrite the inner loop so that only 16 acc values are live at once, and report the lb(1024) spill bytes before and after.

Summary

No bounds means the compiler optimises one thread and leaves the launch to luck. __launch_bounds__(T, B) caps registers at about 65,536 / (T·B), records the contract in PTX (.maxntid, .minnctapersm) and turns over-sized launches into immediate errors. Read ptxas -v, convert registers to occupancy, and let the profiler decide.

Sources (for further reading)