Compute Shader vs CPU Spatial Indexing

A CPU spatial index is astonishingly good at what it was built for: flatbush answers a single bounding-box query against a million rectangles in microseconds, and rbush handles interactive hover-and-pick without ever troubling the GPU. The trouble starts when the query pattern inverts — hundreds of thousands of range queries per frame for viewport culling, or a nearest-neighbour lookup for every point in a streaming batch. At that point the per-query CPU cost, tiny as it is, multiplies into a frame-killer on a single thread, and a GPU compute shader that runs every query in parallel becomes the faster answer despite the cost of shipping the data across the bus. This reference gives you the break-even model that decides between them: the transfer and dispatch overhead the GPU pays up front, the per-query cost each side charges, and the batch size at which parallelism finally pays for the round-trip.

This is one area of Performance Tuning & Profiling for WebGPU Spatial Pipelines, and it is the query-side companion to the transform-side decision in WebGPU vs WebGL 2.0 for spatial workloads. The GPU-resident indexes it discusses — uniform grids and quadtrees built in WGSL — are the subject of WGSL spatial algorithms on the GPU, and the queries feed the same downstream stages as geometry filtering with WGSL compute shaders and spatial aggregation in GPU memory.

When to offload and when not to

The decision hinges on three quantities: the query batch size $Q$ (how many index queries you issue at once), the residency of the data (already in VRAM, or sitting in a CPU typed array), and whether you are optimizing latency (one query, answer now) or throughput (a flood of queries, answer them all soon). The CPU wins decisively on latency and small batches because it has zero transfer overhead and a mature, cache-friendly tree. The GPU wins on throughput once the batch is large enough to amortize the upload, dispatch, and readback across thousands of parallel lanes — and it wins earlier still when the points to query are already resident in VRAM from a prior pass, because then the largest cost, the transfer, is already paid.

Workload	Keep on CPU (rbush / flatbush)	Offload to GPU compute	Deciding factor
Single hover / click pick	Yes	No	Latency: one query, no batch to amortize transfer.
Interactive tooltip range query	Yes	No	Small $Q$ ; CPU tree answers in microseconds.
Viewport culling, millions of features/frame	No	Yes	Large $Q$ every frame; parallelism dominates.
Per-point nearest-neighbour over a streaming batch	No	Yes	$Q$ equals batch size; throughput-bound.
Points already produced by a prior compute pass	No	Yes	Data resident in VRAM; transfer cost already paid.
Index rebuilt every frame from tiny CPU data	Yes	No	Build cost dwarfs query savings at small $N$ .
Bulk join: every point vs every polygon	No	Yes	$Q \times$ candidates is huge; GPU absorbs it in parallel.
One-off query against a static CPU dataset	Yes	No	No reuse to amortize a device upload.
Latency-critical single query on hot path	Yes	No	Readback latency alone exceeds a CPU tree walk.

The GPU's fixed upload-and-readback overhead sets the intercept; the crossover Q* falls as that overhead shrinks — most sharply when the data is already resident in VRAM.

The chart formalizes into an inequality. Let $t_{c p u}$ be the per-query CPU cost, $t_{g p u}$ the effective per-query GPU cost after parallelism, and $t_{u p} + t_{d o w n}$ the fixed upload plus readback overhead. The GPU is faster exactly when

t_{u p} + t_{d o w n} + Q \cdot t_{g p u} < Q \cdot t_{c p u}

which rearranges to a break-even batch size

Q^{*} = \frac{t _{u p} + t _{d o w n}}{t _{c p u} - t _{g p u}}, GPU wins for Q > Q^{*} (assuming t_{c p u} > t_{g p u}) .

When the query points already live in VRAM, $t_{u p} \to 0$ and $Q^{*}$ collapses toward zero — offload almost always wins. The concrete measurement of $Q^{*}$ on real hardware is the subject of the guide on deciding when to offload R-tree queries to the GPU.

Prerequisites

Before wiring an offload path, you should have:

A working CPU baseline — an rbush or flatbush index you can time per query, so the break-even model has a real $t_{c p u}$ rather than a guess. Never offload without the number you are trying to beat.
A negotiated GPUDevice with maxStorageBufferBindingSize checked against your index and query buffers — device and limit negotiation is covered in initializing WebGPU devices for GIS workloads.
A GPU-friendly index structure. A uniform grid or a flattened quadtree maps to array<u32> cleanly; a pointer-chasing R-tree does not, which is why GPU offload usually rebuilds the index as a grid — see WGSL spatial algorithms on the GPU.
Coordinates in one projected linear space (Web Mercator metres or normalized tile space) so a grid cell index is a stable integer on both the CPU and the GPU.
A timing method — GPU timestamp queries for the dispatch and readback, and performance.now() around the CPU tree walk — so the terms of the break-even inequality are measured, not assumed.

Index and query cost reference

The values below are the terms of the break-even model and the constraints that bound them. Per-query costs are order-of-magnitude, not device-exact; measure yours.

Term / rule	Typical value	Why it matters for the decision
$t_{c p u}$ (flatbush range query)	~0.1–1 µs	The number to beat; grows with result count $k$ , not just $lo g N$ .
$t_{g p u}$ (grid query, per lane)	~1–10 ns effective	Per-query cost after spreading across a workgroup; tiny once resident.
$t_{u p}$ (upload queries)	~0.1–2 ms	`writeBuffer` of the query buffer; dominates $Q^{*}$ for CPU-resident data.
$t_{d o w n}$ (readback results)	~0.2–1 ms	`mapAsync` round-trip latency; often the single largest fixed cost.
Grid buffer usage	`STORAGE \| COPY_DST`	The GPU index is a read-only storage buffer for the query kernel.
Result buffer usage	`STORAGE \| COPY_SRC`	Query outputs stage through `COPY_SRC` for the readback.
`@workgroup_size`	64–256	One invocation per query; 64 suits divergent grid walks, 256 uniform scans.
`maxStorageBufferBindingSize`	≥ 128 MiB	A grid over millions of features plus its cell lists can approach this.
Break-even $Q^{*}$ (CPU-resident)	~10k–100k queries	Below it, stay on CPU; above it, offload. Falls sharply when data is VRAM-resident.

For authoritative wording on storage buffers, dispatch, and readback, consult the W3C WebGPU specification and the WGSL specification.

Implementation walkthrough

The pipeline contrasts one range-query batch answered two ways: a CPU tree walk, and a GPU grid kernel with one invocation per query. Each step names the cost term it drives.

Step 1 — Build the index on each side

The CPU side builds a static flatbush R-tree; the GPU side builds a uniform grid — a flat array<u32> of cell start offsets plus a sorted feature list. The grid is chosen over an R-tree because it maps to storage buffers with no pointer chasing, which is what keeps $t_{g p u}$ low.

typescript

import Flatbush from "flatbush";

// CPU index: pack N bounding boxes into a static R-tree once.
function buildCpuIndex(boxes: Float32Array /* [minX,minY,maxX,maxY]*N */): Flatbush {
  const n = boxes.length / 4;
  const index = new Flatbush(n);
  for (let i = 0; i < n; i++) {
    const o = i * 4;
    index.add(boxes[o], boxes[o + 1], boxes[o + 2], boxes[o + 3]);
  }
  index.finish();                        // builds the packed Hilbert R-tree
  return index;
}

Building the equivalent GPU grid — bucket counts, a prefix sum, and a scatter of feature ids into cell lists — is a compute job in its own right, detailed in WGSL spatial algorithms on the GPU. Here we assume cellStart: array<u32> and cellItems: array<u32> already reside in VRAM.

Step 2 — Answer a query batch on the CPU

The CPU baseline loops the batch and walks the tree per query. This is the code path whose total time is $Q \cdot t_{c p u}$ — fast per query, linear in $Q$ , single-threaded.

typescript

// For each query box, collect the ids of features whose boxes intersect it.
function queryCpuBatch(index: Flatbush, queries: Float32Array /* [minX,minY,maxX,maxY]*Q */): number[][] {
  const results: number[][] = [];
  for (let q = 0; q * 4 < queries.length; q++) {
    const o = q * 4;
    // flatbush.search returns candidate indices; O(log N + k) per query.
    results.push(index.search(queries[o], queries[o + 1], queries[o + 2], queries[o + 3]));
  }
  return results;                        // total cost ~ Q * t_cpu, one thread
}

At a few hundred queries this finishes inside a frame with room to spare. At a few hundred thousand — one per visible feature for culling — the linear term dominates and the frame budget is gone.

Step 3 — Answer the same batch as a WGSL grid kernel

The GPU kernel runs one invocation per query. Each invocation maps its query box to the grid cells it overlaps and tests the features in those cells, writing a hit count (or a compacted id list) to an output buffer. Every query runs in parallel, so the effective per-query cost is $t_{g p u}$ , far below $t_{c p u}$ .

wgsl

@group(0) @binding(0) var<storage, read>       queries:   array<vec4<f32>>;  // Q query boxes
@group(0) @binding(1) var<storage, read>       featBoxes: array<vec4<f32>>;  // feature boxes
@group(0) @binding(2) var<storage, read>       cellStart: array<u32>;        // grid offsets
@group(0) @binding(3) var<storage, read>       cellItems: array<u32>;        // feature ids per cell
@group(0) @binding(4) var<storage, read_write> hits:      array<atomic<u32>>;// per-query hit count
@group(0) @binding(5) var<uniform>             gridInfo:  vec4<f32>;          // originX,originY,cellSize,gridW

@compute @workgroup_size(64)
fn query(@builtin(global_invocation_id) gid: vec3<u32>) {
    let q = gid.x;
    if (q >= arrayLength(&queries)) { return; }          // guard the ragged tail
    let box = queries[q];
    let cs = gridInfo.z;
    let gw = u32(gridInfo.w);
    // Grid cell range the query box overlaps.
    let x0 = u32(max(0.0, floor((box.x - gridInfo.x) / cs)));
    let y0 = u32(max(0.0, floor((box.y - gridInfo.y) / cs)));
    let x1 = u32(max(0.0, floor((box.z - gridInfo.x) / cs)));
    let y1 = u32(max(0.0, floor((box.w - gridInfo.y) / cs)));

    var count: u32 = 0u;
    for (var cy = y0; cy <= y1; cy++) {
        for (var cx = x0; cx <= x1; cx++) {
            let cell = cy * gw + cx;
            let start = cellStart[cell];
            let end = cellStart[cell + 1u];              // sentinel-terminated
            for (var k = start; k < end; k++) {
                let f = featBoxes[cellItems[k]];
                // Axis-aligned box overlap test.
                if (f.x <= box.z && f.z >= box.x && f.y <= box.w && f.w >= box.y) {
                    count += 1u;
                }
            }
        }
    }
    atomicStore(&hits[q], count);                        // one result per query lane
}

The double cellStart[cell] / cellStart[cell + 1u] read is the grid’s answer to R-tree pointer chasing: a contiguous slice per cell, no tree descent, which is why the kernel stays coalesced and $t_{g p u}$ stays flat as $Q$ grows. Storing only a count keeps the output small; a full id list needs a prefix-sum compaction pass like the ones in geometry filtering with WGSL compute shaders.

Step 4 — Account for transfer and readback

The offload is only worth it if the fixed cost is amortized. The upload of the query batch and the readback of results are the $t_{u p}$ and $t_{d o w n}$ terms; measure them, because for CPU-resident data they set $Q^{*}$ almost single-handedly.

typescript

async function queryGpuBatch(
  device: GPUDevice, pipeline: GPUComputePipeline, bind: GPUBindGroup,
  queryBuf: GPUBuffer, queries: Float32Array, hitsBuf: GPUBuffer, readback: GPUBuffer, Q: number,
): Promise<Uint32Array> {
  device.queue.writeBuffer(queryBuf, 0, queries);        // t_up: upload the batch
  const enc = device.createCommandEncoder();
  const pass = enc.beginComputePass();
  pass.setPipeline(pipeline);
  pass.setBindGroup(0, bind);
  pass.dispatchWorkgroups(Math.ceil(Q / 64));            // one lane per query
  pass.end();
  enc.copyBufferToBuffer(hitsBuf, 0, readback, 0, Q * 4);
  device.queue.submit([enc.finish()]);

  await readback.mapAsync(GPUMapMode.READ);              // t_down: readback latency
  const out = new Uint32Array(readback.getMappedRange().slice(0));
  readback.unmap();
  return out;
}

When the queries are generated on the GPU — for instance, the culling of features a prior pass produced — the writeBuffer disappears and the result never leaves VRAM if the next stage consumes it directly. That is the residency case where offload wins at almost any batch size.

Step 5 — Choose the path per frame

With the terms measured, the runtime picks a side from the break-even batch size, not a hunch. Route small batches to the CPU and large ones to the GPU, and treat VRAM-resident queries as always-offload.

typescript

function chooseIndexPath(Q: number, dataInVram: boolean, qStar: number): "cpu" | "gpu" {
  if (dataInVram) return "gpu";          // t_up ~ 0, so Q* collapses toward zero
  return Q > qStar ? "gpu" : "cpu";      // above break-even, parallelism pays
}

Recomputing $Q^{*}$ periodically from live timings keeps the router honest across device classes, since a phone’s transfer cost and a workstation’s are not the same number.

Memory and performance implications

The break-even model makes the trade explicit, but two second-order effects decide whether a real pipeline lands where the model predicts. The first is build amortization: a GPU grid must be built before it can be queried, and if the underlying features change every frame the build cost joins the fixed overhead and pushes $Q^{*}$ up. Offload pays off when the index is built once and queried many times — static basemap features culled against a moving viewport are the ideal case, because the grid is built once and the query batch is enormous and repeated. The second is result size: storing a per-query count is cheap, but materializing full candidate lists means a compaction pass and a larger readback, inflating $t_{d o w n}$ . Keep the GPU output as small as the downstream stage allows, and prefer to consume results on the GPU — feeding them straight into spatial aggregation in GPU memory — so the readback never happens.

On the CPU side, the cost that surprises teams is not the tree walk but garbage: flatbush.search returning fresh arrays per query allocates under a large batch and triggers GC pauses that the per-query microbenchmark never shows. A GPU kernel writing into a pre-allocated buffer has no such tail. And on the latency axis, the readback is a hard floor: mapAsync cannot resolve faster than a submit-to-map round-trip, so for a single interactive query the CPU answers before the GPU has even finished uploading — which is why the decision matrix keeps every single-query, latency-bound workload on the CPU regardless of how fast the kernel is. Workgroup sizing for the query kernel follows the same occupancy logic as any spatial dispatch; the tuning levers are covered under the broader performance work.

Failure modes and diagnostics

Offloading below break-even. Moving a small query batch to the GPU pays $t_{u p} + t_{d o w n}$ for savings smaller than the overhead, so the “optimization” is slower. Detection: measured GPU total exceeds the CPU baseline for the batch. Fix: compute $Q^{*}$ and route batches under it to the CPU; only offload above it or when data is VRAM-resident.
Readback stall on the hot path. Awaiting mapAsync blocks the result until a full round-trip completes, adding milliseconds of latency to an interactive query. Detection: input-to-response lag spikes even though the kernel is fast. Fix: keep single, latency-critical queries on the CPU; for batches, consume results on the GPU without a readback.
GPUValidationError — result buffer too small. Sizing the hits buffer for the wrong $Q$ , or forgetting the cellStart sentinel, overruns the binding. Detection: synchronous validation error at submit, or a kernel reading cellStart[cell + 1u] out of bounds. Fix: allocate hits for the maximum $Q$ and append a trailing sentinel entry to cellStart.
Grid degeneracy under skew. A uniform grid with a cell size tuned for average density piles a dense city block into one cell, so a few query lanes scan thousands of features while the rest idle. Detection: kernel time dominated by a handful of workgroups; timestamp variance is huge. Fix: size cells to the dense region, or switch to a GPU quadtree from WGSL spatial algorithms on the GPU.
Stale index after data change. Querying a GPU grid built from a previous frame’s features returns wrong hits when the features moved. Detection: query results lag the visible data by a frame. Fix: rebuild or incrementally update the grid before querying, and fold the build cost into the break-even calculation so offload is only chosen when it still wins.
Device lost on an oversized query dispatch. A batch large enough to run the kernel past the driver’s TDR watchdog invalidates every buffer. Detection: device.lost resolves mid-frame. Fix: chunk the batch across dispatches and fall back to the CPU path through the browser support fallback routing strategies.

Continue in this section

Deciding When to Offload R-tree Queries to the GPU — a runnable model that measures the break-even batch size on your own hardware and picks the path from live timings.

Performance Tuning & Profiling for WebGPU Spatial Pipelines — the profiling practices that turn the break-even terms into measured numbers.
WebGPU vs WebGL 2.0 for Spatial Workloads — the transform-side API decision this query-side one runs parallel to.
WGSL Spatial Algorithms on the GPU — how to build the GPU-resident grids and quadtrees this offload queries.
Geometry Filtering with WGSL Compute Shaders — the compaction pattern for turning per-query counts into candidate id lists.
Spatial Aggregation in GPU Memory — a downstream stage that consumes query results without a readback.

Up: Performance Tuning & Profiling for WebGPU Spatial Pipelines