On-GPU Viewport Culling for Vector Tiles

Panning a dense vector basemap re-tests visibility every frame: for each feature, is its bounding box inside the current viewport? Doing that on the CPU means walking the whole feature array on the main thread and rebuilding an index buffer before the draw, an O(N) scan that competes with rendering for the same thread and stalls the moment a tile carries a few hundred thousand features. The GPU alternative is one independent bounds test per feature in a compute pass, a single atomic counter that packs the survivors into a dense buffer, and an indirect draw that reads the survivor count straight from VRAM — so the CPU never learns how many features passed and never touches the geometry. This page gives the complete WGSL and TypeScript for that atomic-counter cull, the order-independent variant of the compaction skeleton in WGSL spatial algorithms on the GPU.

The atomic-counter approach differs from the prefix-sum compaction used for level-of-detail work: it does not preserve input order, but it needs no scan pass, so for culling — where draw order is irrelevant for opaque instanced features — it is one pass shorter and simpler. The order-preserving alternative, point-cloud LOD compaction with stream compaction, is the right choice only when the survivors must stay sorted for a following pass; for a flat instanced draw, the atomic cursor wins on simplicity and on avoiding the scan’s sequential dependency.

Two properties make this cheap at tile scale. First, the cull reads only each feature’s 16-byte bounds box, never its geometry, so the pass streams a fraction of the tile’s bytes and stays bandwidth-light. Second, because the survivor count lands in an indirect-argument buffer, the CPU issues the same drawIndirect call every frame regardless of how many features survive — the command buffer is identical whether one feature or a million pass, so it can be recorded once and replayed, and the driver never stalls waiting for a count. That is the structural payoff over a CPU cull, which must rebuild an index buffer and resize the draw each frame before it can submit.

Runnable reference implementation

The three outcomes of a frustum test across a tile grid A seven by five grid of vector tiles is classified against a viewport. Four tiles lie fully inside the frustum and are accepted whole, with no per-primitive test needed. Eight tiles straddle the frustum boundary and are marked partial, so their primitives must be tested individually. The remaining twenty-three tiles lie entirely outside and are rejected with a single bounding-box comparison each. FRUSTUM TEST · TILE CLASSIFICATION 4 tiles accepted whole 8 tiles need per-primitive tests 23 tiles rejected outright per-primitive work: 8 of 35 tiles fully inside — accept whole straddles — test primitives outside — one comparison Test at the tile level first — a two-level cull does an order of magnitude less work than a flat one.
The three-way classification is what makes culling cheap. Only the amber tiles pay for a per-primitive test; the green ones are accepted wholesale and the rest are dismissed by a single box comparison each.

The cull kernel tests each feature’s axis-aligned bounds against the viewport extent with a separating-axis test, and every survivor claims the next slot in the output buffer with atomicAdd on a shared cursor. That same cursor, after the pass, is the instance count — so a one-invocation finalize pass copies it into the indirect draw arguments. No prefix sum, no readback.

wgsl
// ---- Cull pass: bounds test + atomic compaction ----
struct Feature {
    bounds: vec4<f32>,   // min_x, min_y, max_x, max_y in projected space
    payload: vec2<u32>,  // e.g. feature id + style index
};

@group(0) @binding(0) var<storage, read>       features: array<Feature>;
@group(0) @binding(1) var<storage, read_write> visible:  array<Feature>;      // dense survivors
@group(0) @binding(2) var<storage, read_write> counter:  atomic<u32>;         // shared cursor
@group(0) @binding(3) var<uniform>             view:     vec4<f32>;           // viewport extent

@compute @workgroup_size(256)
fn cull(@builtin(global_invocation_id) gid: vec3<u32>) {
    let i = gid.x;
    if (i >= arrayLength(&features)) { return; }          // guard the ragged tail
    let b = features[i].bounds;

    // Separating-axis overlap: the feature AABB is culled only if it lies
    // entirely off one side of the viewport AABB. Overlap keeps it.
    let separated = b.z < view.x || b.x > view.z ||
                    b.w < view.y || b.y > view.w;
    if (separated) { return; }                            // outside — do not emit

    let slot = atomicAdd(&counter, 1u);                   // claim next dense slot
    if (slot < arrayLength(&visible)) {                   // guard the output capacity
        visible[slot] = features[i];                      // pack survivor, order not preserved
    }
}

// ---- Finalize: turn the survivor count into indirect draw arguments ----
@group(0) @binding(0) var<storage, read>       counter:  u32;                 // final survivor total
@group(0) @binding(1) var<storage, read_write> drawArgs: array<u32>;          // 4 x u32

const VERTS_PER_FEATURE: u32 = 6u;   // two triangles per instanced quad

@compute @workgroup_size(1)
fn finalize() {
    drawArgs[0] = VERTS_PER_FEATURE;                      // vertexCount
    drawArgs[1] = min(counter, arrayLength(&drawArgs));   // instanceCount (clamped to capacity)
    drawArgs[2] = 0u;                                     // firstVertex
    drawArgs[3] = 0u;                                     // firstInstance
}

The TypeScript driver clears the atomic counter to zero each frame (a stale cursor overcounts), records the cull and finalize passes, then draws with drawIndirect against the compacted buffer. The counter and the draw-arguments buffer are separate: the counter is a plain atomic<u32> written during the cull, and finalize copies it into the INDIRECT-usage argument buffer the render pass consumes.

typescript
function cullAndDraw(
  device: GPUDevice,
  ctx: GPUCanvasContext,
  view: Float32Array,               // [min_x, min_y, max_x, max_y]
  buffers: { counter: GPUBuffer; viewUbo: GPUBuffer; drawArgs: GPUBuffer },
  cull: { pipeline: GPUComputePipeline; bind: GPUBindGroup },
  finalize: { pipeline: GPUComputePipeline; bind: GPUBindGroup },
  render: { pipeline: GPURenderPipeline; bind: GPUBindGroup },
  featureCount: number,
): void {
  device.queue.writeBuffer(buffers.viewUbo, 0, view);

  const enc = device.createCommandEncoder();
  enc.clearBuffer(buffers.counter);                       // reset cursor; stale count overdraws

  const c = enc.beginComputePass();
  c.setPipeline(cull.pipeline);
  c.setBindGroup(0, cull.bind);
  c.dispatchWorkgroups(Math.ceil(featureCount / 256));
  c.end();

  const f = enc.beginComputePass();
  f.setPipeline(finalize.pipeline);
  f.setBindGroup(0, finalize.bind);
  f.dispatchWorkgroups(1);
  f.end();

  const r = enc.beginRenderPass({
    colorAttachments: [{
      view: ctx.getCurrentTexture().createView(),
      loadOp: "clear", storeOp: "store", clearValue: { r: 0, g: 0, b: 0, a: 1 },
    }],
  });
  r.setPipeline(render.pipeline);
  r.setBindGroup(0, render.bind);                         // binds `visible` as instance data
  r.drawIndirect(buffers.drawArgs, 0);                    // instance count read from VRAM
  r.end();

  device.queue.submit([enc.finish()]);                   // one submission: cull → finalize → draw
}

At the scale of a whole basemap — thousands of tiles, tens of millions of features — testing every feature every frame is still wasteful when most tiles are entirely off screen. The natural extension is coarse-to-fine: cull tile bounding boxes first, then run the per-feature cull only over the features of tiles that survived. The node ranges produced by generating a quadtree on the GPU with WGSL give exactly that hierarchy for free: a node’s Morton prefix bounds a spatial cell, so the same separating-axis test applied to node cells prunes whole subtrees before a single feature is touched. The compute pass shown here becomes the leaf level of that traversal, and the atomic counter still gathers the final survivors into one dense buffer for a single indirect draw.

Parameter reference

What two-level culling removes from the primitive test A bar chart of primitive tests performed per frame for a scene of 35 tiles holding 12000 primitives each. A flat cull tests all 420000 primitives against the frustum. A two-level cull performs 35 tile-level box tests, then tests only the 96000 primitives inside the eight straddling tiles, because the four fully-inside tiles are accepted whole and the twenty-three outside tiles are rejected whole. PRIMITIVE TESTS PER FRAME · THOUSANDS Flat cull 420 K tests Two-level cull 96 K tests + 35 box tests 0 150 300 450 K per-primitive tests after tile classification The four fully-inside tiles are the real saving — accepting a tile whole skips its primitives entirely.
Thirty-five cheap box tests remove seventy-seven percent of the expensive ones. The ratio improves as the viewport gets smaller relative to the dataset, which is exactly the case that matters when a user zooms in.

Every tunable value in the implementation above. These tables scroll horizontally on narrow viewports.

Parameter Typical value Guidance
visible capacity worst-case feature count The atomic cursor is guarded against overflow; size the output to the densest tile so a full-viewport frame never drops survivors.
@workgroup_size 256 Multiple of warp/wavefront width. The cull is bandwidth-bound, so larger groups rarely help.
Feature stride 32 bytes vec4<f32> bounds (16 B) + vec2<u32> payload (8 B) + 8 B pad to 16-byte alignment. Match the TypeScript layout exactly.
view extent frame uniform The current viewport AABB in the same projected space as the bounds. Add a margin to pre-load features about to scroll in.
drawArgs usage STORAGE | INDIRECT | COPY_DST Missing INDIRECT fails drawIndirect validation.
counter usage STORAGE | COPY_DST COPY_DST lets clearBuffer zero it each frame.
VERTS_PER_FEATURE 6 Two triangles per instanced quad; change with your instancing geometry.

Failure modes

Culling defects that only appear at the edges A table of three culling defects with cause and fix. Primitives popping in at the screen edge during a pan mean the frustum was tested without a margin, and the fix is to expand the test rectangle by the widest stroke or symbol on the layer. Large polygons vanishing when the camera is inside them mean only the vertices were tested, not the polygon extent, and the fix is to test the bounding box against the frustum rather than the points. A tile flickering between accepted and rejected means the tile bounds were computed in a different coordinate space from the frustum, and the fix is to compute both from the same view matrix. CULLING DEFECT · CAUSE · FIX Cause Fix popping at the edge no margin on the test expand by stroke width big polygons vanish vertices tested, not extent test the bounding box a tile flickers mismatched coordinate space one view matrix for both The margin has to cover half the stroke width plus any symbol anchored off-centre.
All three are invisible on a static screenshot and obvious in motion, which is why culling changes need to be reviewed by panning and zooming rather than by comparing frames.
  • Overdraw from a stale counter. Skipping clearBuffer on the atomic cursor leaves last frame’s total, so the cursor starts high and the instance count grows every frame until it saturates. Detection: instance count climbs monotonically on a static view. Fix: zero the counter before every cull pass.
  • Dropped survivors past capacity. If a full-viewport frame has more survivors than the visible buffer holds, the slot < arrayLength(&visible) guard silently discards the overflow. Detection: features vanish only when the whole tile is on screen. Fix: size visible to the worst-case feature count, not the average.
  • GPUValidationError — indirect buffer usage. Binding drawArgs to drawIndirect without GPUBufferUsage.INDIRECT fails validation at draw time. Detection: synchronous error at drawIndirect. Fix: create drawArgs with STORAGE | INDIRECT | COPY_DST.
  • Features cull one frame late. Writing the view uniform after recording the cull pass, or reusing a previous frame’s extent, lags visibility by a frame and pops features at the viewport edge. Detection: a strip of missing features trails the pan direction. Fix: writeBuffer the current extent before beginComputePass, and add an extent margin to pre-load.

It is worth stating the invariant this pass has to preserve: a primitive that would have been visible must never be culled, while a primitive that is culled and should not have been is merely wasted work. Because the test is asymmetric in cost, every margin and every rounding decision in the code below errs towards keeping geometry rather than dropping it.

Backend / Python interop note

Culling only reads bounds, so the server can ship a compact vec4<f32> bounds array plus a payload table, deferring the full geometry until a feature survives — the cull runs over 16 bytes per feature instead of the whole polygon. Precompute each feature’s axis-aligned bounds server-side in the same projected space the client uses (Web Mercator metres or normalized tile space), because a bounds box computed in degrees will not overlap a viewport expressed in metres and every feature culls. The indirect-draw hand-off that consumes the compacted visible buffer in a deck.gl layer is covered in indirect draw calls for deck.gl instanced layers.

Up: WGSL Spatial Algorithms on the GPU