Workgroup Occupancy Optimization for Spatial Kernels

A spatial compute kernel that hashes points into cells, tests primitives against a viewport, or accumulates a density grid is almost always memory-bound: each invocation issues a storage read, then waits hundreds of cycles for it to land. The only thing that hides that latency is having enough other invocations resident on the same compute unit to run while the first one stalls. That ratio — how many invocations are actually resident against how many the hardware could hold — is occupancy, and it is the single number that most often separates a kernel running at memory bandwidth from one running at a third of it. Yet occupancy is not something you request; it is an emergent consequence of three shader-level decisions: the @workgroup_size you compile in, how much var<workgroup> shared memory each workgroup reserves, and how many registers each invocation needs for its live values. Get any one wrong and the scheduler quietly launches fewer concurrent workgroups, collapsing the latency hiding that made the GPU worth using.

This area sits inside Performance Tuning & Profiling for WebGPU Spatial Pipelines. It is the counterpart to the memory-side work in VRAM budget management across tile zoom levels: occupancy governs how fast a resident kernel runs, while VRAM budgeting governs whether its buffers fit at all. The tuning here builds directly on the dispatch-level flags in optimization flags for compute dispatches, and it generalizes the empirical sizing sweep first shown for filtering in optimizing workgroup sizes for vector geometry filtering. Every claim below is verified by measurement, using the frame profiling with timestamp queries harness, because occupancy is a property of a specific kernel on specific silicon and cannot be reasoned about from source alone.

What occupancy actually measures

Formally, occupancy is the ratio of resident invocations to the hardware ceiling of invocations a single compute unit (an NVIDIA SM, an AMD CU, an Apple or Adreno equivalent) can hold at once. If $n_{\text{wg}}$ workgroups are co-resident and each holds $|\text{wg}|$ invocations, and the hardware ceiling is $I_{\max }$ invocations per compute unit, then

The number that matters is $n_{\text{wg}}$, and it is not something you set — it is the minimum of three independent ceilings the scheduler evaluates when it decides how many copies of your workgroup fit:

Here $W_{\max }$ is the hard cap on resident workgroup slots, $M$ is the shared memory available per compute unit and $m_{\text{wg}}$ is the var<workgroup> bytes your kernel declares, and $R$ is the register file size with $r$ the registers each invocation needs to hold its live temporaries. Whichever term is smallest is your limiter. A spatial kernel that declares a 16 KiB var<workgroup> scratch grid is almost always shared-memory limited; one that unrolls a Gaussian kernel into dozens of live accumulators is register limited; a tiny predicate kernel with neither is workgroup-slot limited and simply wants the largest legal size. The entire craft of occupancy tuning is identifying which term binds and relaxing it without breaking the algorithm.

Prerequisites

Before tuning occupancy for a spatial kernel, you should have:

1. A working WGSL compute kernel that is memory-bound, not compute-bound. Occupancy hides memory latency; it does nothing for a kernel already saturating the ALUs. Confirm the class with a timestamp sweep before tuning, as in optimizing workgroup sizes for vector geometry filtering.
2. The adapter limits in hand. maxComputeInvocationsPerWorkgroup, maxComputeWorkgroupSizeX/Y/Z, and maxComputeWorkgroupStorageSize are read from device.limits after negotiating them during initializing WebGPU devices for GIS workloads.
3. Timestamp queries wired up. Every occupancy claim is validated against a measured dispatch duration, using the frame profiling with timestamp queries harness — the browser exposes no direct occupancy counter, so wall-clock per-dispatch time is the proxy.
4. A Structure-of-Arrays input layout. Occupancy interacts with coalescing: contiguous array<vec2<f32>> positions let a resident workgroup issue one coalesced burst, whereas interleaved structs scatter the reads and waste the residency you fought for. The memory alignment for spatial data buffers reference covers the packing rules.
5. Knowledge of the target device classes. A size that maximizes occupancy on a 32-wide desktop NVIDIA part can under-fill a 64-wide AMD wavefront or a mobile Adreno; plan to sweep across the tiers you ship to.

Limits and occupancy reference

The limits below are the ceilings the scheduler divides against. The three maxCompute… values are guaranteed floors in the WebGPU specification; the physical per-compute-unit resources ($M$, $R$, $I_{\max }$) are not exposed to the browser and must be inferred by measurement.

These tables scroll horizontally on narrow viewports. For the normative wording on workgroup limits and the @workgroup_size attribute, consult the W3C WebGPU specification and the WGSL specification.

Implementation walkthrough

The workflow is: read the limits, pick a size that is a subgroup multiple within them, account for shared memory and registers, measure, and iterate. Each step below carries the WGSL or TypeScript that makes it concrete for a spatial kernel.

Step 1 — Read the limits and derive the legal size window

Before choosing a size, clamp the search to what the device actually allows. The product of the three axes must not exceed maxComputeInvocationsPerWorkgroup, and a var<workgroup> declaration must fit maxComputeWorkgroupStorageSize. Derive both up front so no candidate size is illegal.

interface OccupancyBudget {
  maxInvocations: number;   // maxComputeInvocationsPerWorkgroup
  maxX: number;             // maxComputeWorkgroupSizeX
  maxSharedBytes: number;   // maxComputeWorkgroupStorageSize
  subgroupWidth: number;    // 32 desktop default; 64 on AMD; probe per device class
}

function readBudget(device: GPUDevice, subgroupWidth = 32): OccupancyBudget {
  const l = device.limits;
  return {
    maxInvocations: l.maxComputeInvocationsPerWorkgroup, // e.g. 256 floor
    maxX: l.maxComputeWorkgroupSizeX,                    // e.g. 256 floor
    maxSharedBytes: l.maxComputeWorkgroupStorageSize,    // e.g. 16384 floor
    subgroupWidth,
  };
}

// Candidate 1-D sizes: subgroup multiples that stay within the invocation cap.
function candidateSizes(b: OccupancyBudget): number[] {
  const out: number[] = [];
  for (let n = b.subgroupWidth; n <= Math.min(b.maxInvocations, b.maxX); n += b.subgroupWidth) {
    if ((n & (b.subgroupWidth - 1)) === 0) out.push(n); // whole subgroup multiples only
  }
  return out; // e.g. [32, 64, 96, 128, 160, 192, 224, 256] on a 32-wide part
}

Generating candidates as subgroup multiples rather than arbitrary integers is the spatial-kernel-specific discipline: a size of 100 leaves 28 lanes of the fourth subgroup permanently masked, so every point batch pays for silicon it never uses.

Step 2 — Choose the size as a multiple of subgroup width

The subgroup (warp on NVIDIA, wavefront on AMD) is the true unit of scheduling: the hardware runs all $s$ lanes of a subgroup in lockstep. A workgroup that is not a whole multiple of $s$ rounds up to one internally and masks the surplus lanes, so a spatial scatter kernel processing one point per invocation wastes exactly those lanes on every dispatch. Compile the size in as a WGSL override constant so one shader source can be specialized per device class.

// Overridable workgroup width — set at pipeline creation per device tier.
override WG_SIZE: u32 = 256u;

@group(0) @binding(0) var<storage, read>       positions: array<vec2<f32>>;
@group(0) @binding(1) var<storage, read_write> grid:      array<atomic<u32>>;
@group(0) @binding(2) var<uniform>             extent:    vec4<f32>;

const GRID_W: u32 = 1024u;
const GRID_H: u32 = 1024u;

@compute @workgroup_size(WG_SIZE)
fn scatter(@builtin(global_invocation_id) gid: vec3<u32>) {
    let i = gid.x;
    if (i >= arrayLength(&positions)) { return; } // ragged tail guard
    let p = positions[i];
    let n = (p - extent.xy) / (extent.zw - extent.xy);
    if (n.x < 0.0 || n.x >= 1.0 || n.y < 0.0 || n.y >= 1.0) { return; }
    let cell = u32(n.y * f32(GRID_H)) * GRID_W + u32(n.x * f32(GRID_W));
    atomicAdd(&grid[cell], 1u); // one coalesced-read, one scatter per invocation
}

Supplying WG_SIZE through the pipeline’s constants map means a single WGSL module compiles to 256 on desktop NVIDIA and 64 on a narrow mobile part without a source fork — the mechanism is detailed in choosing workgroup size for GPU occupancy.

Step 3 — Account for shared memory as the co-residency limiter

A tiled spatial kernel that accumulates into var<workgroup> scratch to cut atomic contention trades occupancy for locality: every byte of declared workgroup storage is a byte the scheduler subtracts from the pool it uses to co-resident workgroups. If physical shared memory per compute unit is $M$ and you declare $m_{\text{wg}}$ bytes, at most $\lfloor M/m_{\text{wg}}\rfloor$ workgroups fit. Size the scratch to the knee of that curve, not the maximum the limit allows.

override WG_SIZE: u32 = 256u;

// A 32x32 u32 scratch tile = 4 KiB. At 4 KiB, an SM with 64 KiB shared memory
// holds up to 16 co-resident workgroups; a 16 KiB tile would hold only 4.
const TILE: u32 = 32u;
var<workgroup> scratch: array<atomic<u32>, TILE * TILE>;

@compute @workgroup_size(WG_SIZE)
fn scatter_tiled(@builtin(local_invocation_index) lid: u32,
                 @builtin(global_invocation_id) gid: vec3<u32>) {
    for (var s = lid; s < TILE * TILE; s += WG_SIZE) { atomicStore(&scratch[s], 0u); }
    workgroupBarrier();
    // ... hash the point into a tile-local cell and atomicAdd into scratch ...
    workgroupBarrier();
    // ... flush non-zero scratch cells to the global grid with one atomic each ...
}

Halving the scratch tile from 16 KiB to 4 KiB can quadruple co-residency, which for a contention-bound spatial scatter often beats the marginally better locality of the larger tile. This is the sole knob the reducing register pressure in WGSL spatial kernels guide leaves to a sibling, because shared memory and registers are two distinct terms in the same min.

Step 4 — Keep register pressure below the co-residency cliff

The third ceiling, $\lfloor R/(r\cdot |\text{wg}|)\rfloor$, is the subtlest because WGSL never names registers. Every live local variable, every element of a var<function> array that the compiler cannot prove dead, consumes register file that scales with the workgroup size. A spatial kernel that keeps a 5×5 Gaussian window in a local array holds 25 live f32 values across the whole loop; recomputing the weight per tap instead keeps two. Prefer recompute-over-store and narrow the live range of every temporary.

// Register-hungry: the whole neighbourhood is live at once in a local array.
fn smooth_heavy(cx: i32, cy: i32) -> f32 {
    var win: array<f32, 25>;                 // 25 live f32 — high register pressure
    var k = 0;
    for (var dy = -2; dy <= 2; dy++) {
        for (var dx = -2; dx <= 2; dx++) { win[k] = sample(cx + dx, cy + dy); k++; }
    }
    var acc = 0.0;
    for (var j = 0; j < 25; j++) { acc += win[j] * gauss(j); }
    return acc;
}

// Register-lean: one accumulator, weights recomputed — far more co-resident workgroups.
fn smooth_lean(cx: i32, cy: i32) -> f32 {
    var acc = 0.0;
    var norm = 0.0;
    for (var dy = -2; dy <= 2; dy++) {
        for (var dx = -2; dx <= 2; dx++) {
            let w = exp(-f32(dx * dx + dy * dy) * 0.25); // recompute, do not store
            acc += sample(cx + dx, cy + dy) * w;
            norm += w;
        }
    }
    return acc / norm;
}

The lean form holds a handful of live values instead of 25, so the register term stops being the limiter and co-residency jumps. Recognizing and rewriting these patterns is the whole subject of reducing register pressure in WGSL spatial kernels.

Step 5 — Measure occupancy indirectly with timestamp queries

WebGPU exposes no occupancy counter, so measure its effect: sweep the candidate sizes, time each dispatch through a timestamp query set, and read residency off the shape of the curve. A size where throughput plateaus is at peak occupancy; one where it falls off has crossed a co-residency cliff.

async function timeDispatch(
  device: GPUDevice,
  pipeline: GPUComputePipeline,
  bindGroup: GPUBindGroup,
  workgroups: number,
  querySet: GPUQuerySet,
  resolve: GPUBuffer,
  readback: GPUBuffer,
): Promise<number> {
  const enc = device.createCommandEncoder();
  const pass = enc.beginComputePass({
    timestampWrites: { querySet, beginningOfPassWriteIndex: 0, endOfPassWriteIndex: 1 },
  });
  pass.setPipeline(pipeline);
  pass.setBindGroup(0, bindGroup);
  pass.dispatchWorkgroups(workgroups);
  pass.end();
  enc.resolveQuerySet(querySet, 0, 2, resolve, 0);
  enc.copyBufferToBuffer(resolve, 0, readback, 0, 16);
  device.queue.submit([enc.finish()]);

  await readback.mapAsync(GPUMapMode.READ);
  const t = new BigUint64Array(readback.getMappedRange());
  const nanos = Number(t[1] - t[0]); // end minus begin, in nanoseconds
  readback.unmap();
  return nanos;
}

Run this warmed and take the median across several submissions, exactly as the frame profiling with timestamp queries reference prescribes — a single sample is dominated by scheduling jitter and will mislead the sizing decision.

Memory and performance implications

Occupancy is a latency-hiding mechanism, so its payoff is proportional to how memory-bound the kernel already is. For a spatial scatter that reads one vec2<f32> and does a single atomicAdd, the arithmetic intensity is near zero and residency is almost everything: going from 25% to 90% occupancy on such a kernel can nearly triple throughput. For a projection kernel doing a dozen transcendental operations per point, the ALUs are busy enough to hide their own reads, and pushing occupancy past the point where latency is covered buys nothing — the extra residency competes for the same execution ports. The practical rule is to raise occupancy only until the timestamp curve flattens, then stop; beyond the knee you are trading register file and shared memory for throughput you cannot use.

There is a direct tension with the two other resident resources. Every kilobyte of var<workgroup> scratch you add to cut atomic contention subtracts from the shared-memory pool that bounds co-residency, and every live temporary you keep to avoid recomputation subtracts from the register file that bounds it. The best configuration for a contention-heavy density scatter is frequently a smaller workgroup with a smaller scratch tile running at high occupancy, not the 16 KiB tile the limit would permit — the reference implementation in spatial aggregation in GPU memory is sized with exactly this trade in mind. Because occupancy is per compute unit and dispatch count is global, the two compose: pick the size for occupancy first, then dispatch ceil(points / size) workgroups so the whole device stays saturated, a division handled by the dispatch-side flags in optimization flags for compute dispatches.

Failure modes and diagnostics

• Pipeline creation fails — workgroup size over the invocation cap. A @workgroup_size(32, 32) totals 1024, over the default maxComputeInvocationsPerWorkgroup of 256. Detection: a GPUValidationError at createComputePipeline, naming the exceeded limit. Fix: keep the product x*y*z within the limit, or request a higher one at device creation if the adapter reports it.
• var<workgroup> declaration exceeds storage limit. A scratch array larger than maxComputeWorkgroupStorageSize (16 KiB floor) is rejected. Detection: GPUValidationError at pipeline creation referencing workgroup storage. Fix: shrink the tile (Step 3), or move part of the accumulation to global memory.
• Occupancy collapse from register spilling. A kernel with large local arrays compiles and runs but times far slower than a leaner variant. Detection: the timestamp sweep shows throughput falling as size grows past a threshold where a lean kernel keeps climbing. Fix: apply the recompute-over-store rewrites in reducing register pressure in WGSL spatial kernels.
• Masked subgroup lanes from a non-multiple size. A @workgroup_size(100) runs correctly but wastes 28 of every 128 lanes on a 32-wide part. Detection: throughput is a step below the nearest subgroup-multiple size in the sweep, for no algorithmic reason. Fix: round the size to a whole multiple of the subgroup width (Step 1).
• Peak occupancy, no speedup — the kernel is compute-bound. Raising occupancy does nothing because the ALUs, not memory, are the bottleneck. Detection: the timestamp curve is flat across all sizes. Fix: stop tuning occupancy and reduce arithmetic per invocation, or accept the ALU ceiling.
• A size tuned on desktop halves throughput on mobile. A 256-wide size chosen on a 32-wide NVIDIA part over-subscribes a narrow Adreno register file. Detection: the median dispatch time regresses only on the mobile tier. Fix: sweep and specialize the override constant per device class, and route unsupported devices through the browser support fallback strategies.

Continue in this section

• Choosing Workgroup Size for GPU Occupancy — a runnable sweep that picks and benchmarks @workgroup_size across desktop, mobile, and integrated device classes with an override constant.
• Reducing Register Pressure in WGSL Spatial Kernels — before-and-after kernels that trim live temporaries and large local arrays so the register term stops bounding co-residency.

Related

• Optimization Flags for Compute Dispatches — the dispatch-level flags and workgroup-count math that compound with the occupancy chosen here.
• Optimizing Workgroup Sizes for Vector Geometry Filtering — the timestamp sweep applied to a filter pass, the concrete origin of this generalized method.
• Frame Profiling with Timestamp Queries — the measurement harness every occupancy claim above is validated against.
• Spatial Aggregation in GPU Memory — the scatter-and-smooth kernels whose scratch sizing trades directly against occupancy.
• Memory Alignment for Spatial Data Buffers — the packing rules that keep a resident workgroup’s reads coalesced.

Up: Performance Tuning & Profiling for WebGPU Spatial Pipelines