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 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.

// ---- 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.

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

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

Failure modes

• 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.

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.

Related

• WGSL Spatial Algorithms on the GPU — the compaction skeleton this atomic-counter cull specializes.
• Point-Cloud LOD Compaction with Stream Compaction — the order-preserving prefix-sum alternative to this atomic-counter compaction.
• Geometry Filtering with WGSL Compute Shaders — predicate filtering that composes with the viewport test here.
• Finding Render Pass Bottlenecks with GPU Timestamps — confirming the cull, not the draw, is where a frame spends its time.

Up: WGSL Spatial Algorithms on the GPU