Spatial Compute Shaders & Geometry Pipelines

WGSL compute pipelines for filtering, clustering and aggregation, plus the sorts and prefix sums almost every spatial algorithm is built from.

The migration of spatial workloads from CPU-bound JavaScript to GPU-accelerated compute pipelines represents a fundamental shift in how geographic information systems render, analyze, and transform coordinate data. Traditional GIS architectures rely on synchronous JavaScript execution, GEOS bindings, or server-side Python pipelines that introduce latency, memory fragmentation, and main-thread contention. WebGPU compute shaders eliminate these bottlenecks by executing geometry transformations, spatial indexing, and attribute aggregation directly on the GPU, with deterministic memory layouts and explicit synchronization boundaries. This guide establishes the foundational architecture for spatial compute pipelines and is the entry point to a connected set of in-depth references — geometry filtering, asynchronous clustering, in-memory aggregation, and dispatch tuning — that each take one stage of the pipeline to production depth.

The architecture targets four overlapping roles: frontend GIS developers who own the browser pipeline, WebGL/WebGPU engineers porting existing renderers, visualization specialists who consume compute output as vertex data, and Python backend teams responsible for binary serialization and pipeline orchestration. Throughout, a compute shader is treated as a distinct GPU program type — separate from vertex and fragment stages — and understanding where it sits relative to the rest of the GPU is covered in the compute versus render pipeline fundamentals reference. Device acquisition, adapter feature negotiation, and limit inspection are handled upstream during WebGPU device initialization for GIS workloads, and this article assumes a valid GPUDevice is already in hand.

Architecture Overview

A spatial compute pipeline is a directed flow of typed buffers: a Python backend serializes geometry into binary Structure-of-Arrays payloads, the browser uploads them into storage buffers, a chain of compute passes filters and aggregates the data in place, and the final compacted buffer is bound directly as vertex input to a render pass — without a round trip back to the CPU. The diagram below labels each stage and the buffer-usage transitions between them.

Spatial compute pipeline architecture A vertical data flow: a Python backend serializes a spatial dataset into a binary Structure-of-Arrays payload, sent over WebSocket or HTTP2 into a MAP_WRITE staging GPUBuffer, copied into a STORAGE storage buffer, processed by a predicate-filter compute pass and an aggregate or cluster compute pass, compacted into a STORAGE plus VERTEX output buffer, bound zero-copy into a render pass, and drawn to the canvas framebuffer. Right-side labels show the buffer-usage flags at each stage; left-side labels show the transport, copy, and zero-copy transitions. PYTHON BACKEND BROWSER FRONTEND · GPU binary · WebSocket / HTTP2 copyBufferToBuffer zero-copy vertex bind MAP_WRITE | COPY_SRC STORAGE | COPY_DST atomic stream compaction shared-mem reduction STORAGE | VERTEX @vertex · @fragment Python backend GeoParquet / PostGIS → SoA arrays Staging GPUBuffer Storage GPUBuffer Predicate filter @compute @workgroup_size(256) Aggregate / cluster @compute parallel reduction Compacted output buffer Render pass Canvas framebuffer

The remainder of this article walks each stage in dependency order: first the buffer and memory model that everything else is built on, then the compute-driven geometry work, then the validation and cross-browser concerns that determine whether the pipeline survives contact with real hardware, and finally the deployment budgets that govern production behavior.

Core Concept A: Pipeline Boundaries & Memory Layout

Interleaved versus split buffers for a filter kernel Two byte strips share a ruler from zero to thirty-two bytes. The interleaved array-of-structures layout puts one primitive record in thirty-two contiguous bytes: minimum x, minimum y, maximum x, maximum y as four f32 values, then a packed u32 attribute word, a u32 feature id, and eight bytes of trailing padding. A predicate kernel that only reads the attribute word touches four useful bytes out of every thirty-two it pulls through the cache. The split structure-of-arrays layout gives the bounds their own sixteen-byte record and the attributes their own four-byte record in a separate buffer, so the same kernel streams attributes contiguously and reads nothing it does not need. BYTE OFFSET 0 4 8 12 16 20 24 28 32 min_x min_y max_x max_y attrs f_id pad 8 B Interleaved 32 B · 1 buffer an attribute-only kernel drags 32 bytes through cache to use 4 min_x min_y max_x max_y attrs Split 16 B + 4 B bounds and attributes live in separate buffers — each read is coalesced Bandwidth, not arithmetic, is what a spatial filter is usually short of.
The split layout is not automatically better; it is better for the kernels that read one field. A pass that tests bounds and attributes together reads both buffers anyway, and then the interleaved record wins on locality. Choose per kernel, and be willing to keep both.

A production-grade spatial compute pipeline begins with strict separation between data staging, compute execution, and rendering. WebGPU enforces this through explicit GPUBuffer usage flags and pipeline state objects. Geometry payloads must be serialized into tightly packed, aligned structures before upload, because WGSL storage buffers impose strict rules on element stride and base offset that do not match the loose packing of typical JSON or AoS records.

The buffer-usage flags determine where a buffer can travel in the pipeline and which operations are legal against it. Choosing them wrongly is the most common source of GPUValidationError during pipeline bring-up. The table below summarizes the usage combinations that matter for spatial data.

Buffer role Usage flags Spatial-data purpose
Staging upload MAP_WRITE | COPY_SRC Receive a binary coordinate payload from the backend, then copy into a storage buffer
Coordinate / attribute store STORAGE | COPY_DST Hold packed vec4<f32> extents and u32 attribute flags for compute access
Scratch / intermediate STORAGE | COPY_SRC Per-pass working space that may be copied to a readback buffer for diagnostics
Atomic counters STORAGE Hold atomic<u32> write pointers for stream compaction
Compute → render handoff STORAGE | VERTEX Bind compacted geometry directly as vertex input with no CPU copy
Readback (export only) COPY_DST | MAP_READ Final, explicit map for export, never inside an animation frame

Coordinate arrays, bounding box extents, and attribute tables are uploaded to STORAGE | COPY_DST buffers, while intermediate scratch space is allocated with STORAGE | COPY_SRC. Python backend teams should export spatial datasets as contiguous Float32Array or Uint32Array buffers using SoA (Structure of Arrays) layouts rather than AoS, minimizing stride penalties during parallel evaluation. For optimal binary packing, Python’s native array module or PyArrow ensures zero-overhead serialization before GPU transfer.

The buffer creation surface on the browser side is small but exacting. The following TypeScript shows the canonical allocation for a coordinate store and its compaction counter:

typescript
// Pack bounding-box extents as vec4<f32>: (minX, minY, maxX, maxY) per feature.
// SoA keeps each component contiguous so the GPU can coalesce reads across a workgroup.
function createSpatialBuffers(device: GPUDevice, featureCount: number) {
  const bounds = device.createBuffer({
    label: "feature-bounds",
    size: featureCount * 4 * Float32Array.BYTES_PER_ELEMENT, // 16 bytes / feature
    usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
  });

  // Compacted index output: one u32 slot per feature in the worst case.
  const validIndices = device.createBuffer({
    label: "valid-indices",
    size: featureCount * Uint32Array.BYTES_PER_ELEMENT,
    usage: GPUBufferUsage.STORAGE | GPUBufferUsage.VERTEX,
  });

  // Single atomic counter; minimum buffer size is 4 bytes (one u32).
  const counter = device.createBuffer({
    label: "valid-count",
    size: Uint32Array.BYTES_PER_ELEMENT,
    usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST,
  });

  return { bounds, validIndices, counter };
}

Frontend GIS developers must avoid implicit synchronization points such as buffer.mapAsync() inside animation frames. Instead, pipelines should operate entirely on the GPU until final read-back is explicitly required for export or UI overlay. Visualization specialists benefit from this architecture by binding compute output buffers directly to render pipelines via vertex or instance bindings, enabling zero-copy geometry streaming. The compute-to-render boundary is enforced through GPUCommandEncoder pass ordering: within a single command buffer submission, compute dispatches complete before render passes consume the same storage buffers. This ordering guarantee is defined by the WebGPU specification and requires no explicit barrier on the developer’s part.

AoS versus SoA memory layout for feature bounding boxes Two memory rows for the same three features, each a vec4 bounding box of minX, minY, maxX, maxY. In the Array-of-Structures row the four components of each feature are interleaved, so a workgroup reading every feature's minX must jump 16 bytes between elements — a strided, uncoalesced access. In the Structure-of-Arrays row all minX values are stored contiguously, so the same reads form one adjacent, coalesced burst. Array of Structures (AoS) — interleaved minXf0 minYf0 maxXf0 maxYf0 minXf1 minYf1 maxXf1 maxYf1 minXf2 minYf2 maxXf2 maxYf2 workgroup reads minX → 16-byte stride · uncoalesced Structure of Arrays (SoA) — component-contiguous minXf0 minXf1 minXf2 minYf0–2 maxXf0–2 maxYf0–2 adjacent reads · single coalesced burst

Core Concept B: Compute-Driven Geometry Processing

The four compute stages a vector tile passes through Four stages run left to right. A cull pass tests each primitive against the view frustum and writes a pass or fail mask. A compact pass turns that sparse mask into a dense index list using an atomic counter. An aggregate pass bins the survivors into a density grid or accumulates centroids. A draw stage binds the compacted index list as vertex input, so the rasterizer never sees a culled primitive. ONE TILE · FOUR COMPUTE STAGES Cull frustum test per primitive writes a pass/fail mask Compact atomicAdd claims a slot sparse mask → dense list Aggregate density grid · centroids optional per layer Draw compacted list as vertices culled work never rasterizes All four are recorded into one command buffer; only the draw needs a render pass.
Every stage after the first reads only what the previous one kept, which is why the order matters more than the individual kernels: culling first is what makes the compaction cheap, and compacting second is what makes the draw call small.

With buffers in place, the pipeline replaces JavaScript array operations with parallel WGSL evaluation. Instead of iterating over millions of features to apply bounding box culling, distance thresholds, or topological predicates, compute shaders evaluate conditions across workgroups simultaneously. The geometry filtering reference develops this stage in full, showing how to implement compacted output buffers using atomic write pointers so that only valid features proceed to rasterization. By partitioning datasets with @workgroup_size and global_invocation_id, developers map spatial tiles onto the GPU’s execution grid, reducing memory bandwidth pressure during heavy predicate evaluation.

A representative filter kernel evaluates a viewport bounding box against per-feature extents and uses an atomic counter to compact survivors into a dense output array:

wgsl
@group(0) @binding(0) var<storage, read>       bounds: array<vec4<f32>>;
@group(0) @binding(1) var<storage, read_write> valid_indices: array<u32>;
@group(0) @binding(2) var<storage, read_write> count: atomic<u32>;
@group(0) @binding(3) var<uniform>             viewport: vec4<f32>; // (minX,minY,maxX,maxY)

@compute @workgroup_size(256)
fn cull(@builtin(global_invocation_id) gid: vec3<u32>) {
  let idx = gid.x;
  if (idx >= arrayLength(&bounds)) { return; } // guard the ragged final workgroup

  let b = bounds[idx];
  let overlaps = b.x <= viewport.z && b.z >= viewport.x &&
                 b.y <= viewport.w && b.w >= viewport.y;

  if (overlaps) {
    let slot = atomicAdd(&count, 1u); // lock-free stream compaction
    valid_indices[slot] = idx;
  }
}

The performance and memory implications scale directly with dataset size. A workgroup size of 256 is a safe default that keeps occupancy high on most desktop GPUs while staying within the maxComputeInvocationsPerWorkgroup limit; the dispatch count is ceil(featureCount / 256). Memory cost is dominated by the bounds buffer at 16 bytes per feature, so a 10-million-feature layer occupies roughly 160 MB of VRAM for extents alone — well within desktop budgets but a real constraint on integrated GPUs, where tiled streaming becomes necessary. The atomic compaction pattern avoids allocating a worst-case output the size of the input on the CPU and keeps the survivors dense, which matters because the compacted buffer feeds straight into the vertex stage.

When processing complex geometries such as multi-polygons or dense LiDAR point clouds, workload distribution becomes critical. Offloading buffer preparation and command submission to dedicated Web Workers prevents main-thread jank during large-scale dataset ingestion, letting the UI stay responsive while the GPU pipeline processes megabyte-scale coordinate streams.

Spatial Indexing & Atomic Coordination

Efficient spatial querying on the GPU requires deterministic indexing structures that map cleanly to compute workgroups. Traditional CPU-side quadtrees or R-trees do not translate efficiently to parallel execution without careful atomic management. WGSL provides atomicAdd and atomicCompareExchangeWeak operations that allow pipelines to construct dynamic spatial partitions without serializing workgroup execution. By leveraging these primitives, lock-free spatial hash grids and atomic counter buffers can safely aggregate overlapping feature extents across concurrent invocations.

For time-series or streaming spatial data, asynchronous dispatch patterns prevent pipeline stalls. The async dispatch patterns for spatial clustering reference details how to chain compute passes using GPUQueue.onSubmittedWorkDone() and timestamp queries, enabling progressive clustering algorithms that refine centroids and density thresholds across multiple frames. This approach is particularly valuable for real-time heatmaps, kernel density estimation, and dynamic feature generalization, where a single synchronous dispatch would blow the frame budget.

GPU-Side Aggregation

Moving aggregation logic to the GPU drastically reduces network roundtrips and client-side computation overhead. The spatial aggregation in GPU memory reference explains how to implement parallel reduction passes for zonal statistics, attribute summation, and spatial joins. By staging intermediate results in workgroup-shared memory before writing to global storage buffers, pipelines achieve near-linear scaling across GPU cores and cut global-memory traffic by an order of magnitude. This is essential for dashboard-level analytics where sub-second response times are required across millions of spatial records.

Two-level parallel reduction for GPU aggregation Per-feature input values are split across two workgroups. Inside each workgroup a binary tree reduces eight inputs to four, then to a single partial sum held in shared memory, with a workgroupBarrier between each level so all invocations finish their writes before the next read. The two workgroup partial sums are then combined across workgroups, via atomicAdd into a global result buffer, in a final Level 2 pass. Workgroup 0 · shared memory Workgroup 1 · shared memory workgroupBarrier() workgroupBarrier() workgroupBarrier() workgroupBarrier() v0v1v2v3 v4v5v6v7 v0+v1v2+v3 v4+v5v6+v7 partial₀partial₁ Global result atomicAdd → storage inputs Level 1 in-WG tree Level 2 cross-WG

Core Concept C: Validation, Error Handling & Cross-Browser Behavior

Compute pipelines fail differently from CPU code: errors surface asynchronously through the validation and device-loss channels rather than as synchronous exceptions. A production pipeline must scope these explicitly. Wrapping buffer and pipeline creation in pushErrorScope/popErrorScope converts silent validation failures into actionable diagnostics, while a device.lost handler distinguishes a recoverable context teardown (tab backgrounded, driver reset) from an unrecoverable one.

typescript
async function buildPipelineSafely(device: GPUDevice, module: GPUShaderModule) {
  device.pushErrorScope("validation");

  const pipeline = device.createComputePipeline({
    label: "geometry-cull",
    layout: "auto",
    compute: { module, entryPoint: "cull" },
  });

  const error = await device.popErrorScope();
  if (error) {
    // Surface the exact WGSL/layout mismatch instead of a blank canvas.
    throw new Error(`Compute pipeline validation failed: ${error.message}`);
  }

  // Driver resets and GPU process crashes arrive here, not as thrown errors.
  device.lost.then((info) => {
    if (info.reason !== "destroyed") {
      console.warn(`GPUDevice lost (${info.reason}); re-initializing.`);
      // Re-acquire adapter + device and rebuild all GPU resources.
    }
  });

  return pipeline;
}

Cross-browser behavior is the other half of robustness. Adapter limits differ widely — maxStorageBufferBindingSize, maxComputeWorkgroupStorageSize, and maxBufferSize are routinely lower on mobile and on integrated GPUs than on discrete desktop hardware — so a pipeline that assumes desktop limits will throw validation errors on phones. Query the adapter’s reported limits at startup and size buffers and workgroups against the real numbers rather than constants. Where WebGPU is unavailable entirely or compute support is too limited, the pipeline must degrade gracefully; the browser support and fallback routing reference covers the detection logic and the WebGL 2.0 fallback path for environments without a usable compute queue. Feature gating with navigator.gpu.requestAdapter() capability checks lets a single build target both first-class and degraded clients.

Production Deployment Considerations

Choosing a primitive before choosing a kernel

Most spatial work on the GPU decomposes into a very small set of parallel primitives, and picking the right one up front saves more time than tuning the kernel afterwards. A map applies the same arithmetic to every element independently — reprojection, predicate evaluation, colour ramp lookup — and is the easiest case, because there is no communication between invocations at all. A reduction collapses many elements into few, which is what a bounding-box computation or a total count is. A scan, or prefix sum, produces a running total and is the primitive that turns a sparse pass/fail mask into dense output. A sort imposes an ordering, and on spatial data that ordering is almost always a space-filling curve.

The practical consequence is that a CPU algorithm rarely ports directly. A CPU quadtree build walks a tree and allocates nodes; the GPU version sorts Morton keys and then finds boundaries, which is a sort followed by a map and looks nothing like the original code. Recognising that a problem is “a sort followed by a scan” is what makes it tractable, and searching for the GPU version of a specific data structure usually is not.

The cost ordering between the four is worth memorising, because it decides which passes are worth avoiding. A map costs roughly one memory read plus arithmetic per element and is bandwidth-bound. A reduction and a scan both cost a logarithmic number of passes over the data, or one pass with workgroup-level cooperation, and in practice run at a small multiple of a map. A sort is the expensive one — a radix sort over 32-bit keys is several passes over the whole array — which is why a pipeline that sorts once and reuses the ordering for several frames tends to beat one that re-sorts per frame.

Where the compute stage sits in a frame

A production frame does not run compute in isolation; it interleaves it with uploads and rasterization, and the ordering constraints are what make the frame either tight or full of bubbles. The general shape that works is: record every buffer copy for newly arrived tiles first, then every compute pass, then the render passes, all into one command encoder and one submission.

Putting the copies first matters because a copy recorded after a compute pass that reads the destination buffer is a read-after-write hazard the ordering guarantee will resolve in the wrong direction — the compute pass reads the old bytes, which is correct by the rules and wrong for the application. Putting all the compute passes together matters less for correctness and more for the driver: consecutive compute passes avoid the state thrash of switching between compute and raster pipelines several times in a frame.

The one thing that must not appear anywhere in that sequence is a buffer map. mapAsync resolves when the GPU has finished with the buffer, so awaiting it inside the frame turns the whole pipeline synchronous. Results that the CPU genuinely needs — a picked feature id, a cluster count for a legend — should be read from a ring of readback buffers written two or three frames earlier, which costs staleness the user cannot perceive and saves a stall they can.

Validation while a kernel is under development

Spatial compute kernels fail quietly more often than they fail loudly, because out-of-bounds reads return zero and out-of-bounds writes are dropped. Three habits catch most of it before it reaches a map.

The first is to keep a small CPU reference implementation of every kernel and a fixture of a few thousand records, and to compare the two outputs element by element in a test. It is slow and it does not need to be fast; it needs to run in CI. The second is to write the element count into the output buffer alongside the data, so a mismatch between what the kernel thinks it processed and what the host thinks it dispatched shows up as a number rather than as missing geometry. The third is to run development builds inside device.pushErrorScope("validation") around every resource creation and submission, and to log the scope’s result — a validation error names the binding and the reason, which is the difference between a five-minute fix and an afternoon.

Error scopes are cheap enough to leave enabled in a staging build and expensive enough to remove in production, so gate them on a build flag rather than deleting them once a kernel works.

Deploying spatial compute pipelines at scale is governed by three budgets: frame time, VRAM, and CPU/GPU synchronization. For interactive maps targeting 60 fps, the entire compute-plus-render cycle must complete within roughly 16 ms; compute-heavy passes such as clustering should therefore be amortized across frames using the asynchronous dispatch patterns above rather than run to completion in a single frame. Timestamp queries (where the timestamp-query feature is available) give per-pass GPU timings so the budget can be measured rather than guessed.

VRAM is the hard ceiling on dataset size. Tracking buffer allocations against the adapter’s maxBufferSize and total VRAM, and tiling large layers by viewport or zoom level, keeps a session from triggering an out-of-memory device loss. The dispatch-tuning details — workgroup occupancy, 16-byte offset alignment, and minimizing divergent branching — are collected in the optimization flags for compute dispatches reference, and a well-tuned pipeline commonly reaches 2–5× the throughput of a naive one on sparse or irregularly distributed features.

Synchronization is the subtlest budget. Read-back via mapAsync stalls the pipeline whenever it is awaited inside a frame; confine it to explicit export actions, double-buffer any buffer that must be both written by compute and read by the CPU, and rely on intra-submission pass ordering for the compute-to-render handoff instead of manual fences. As browser vendors converge on next-generation graphics APIs and the WebGPU specification stabilizes subgroup operations and larger storage limits, pipelines built around strict memory alignment, explicit synchronization, and modular WGSL composition will remain performant and portable across the evolving ecosystem.

What does not belong on the GPU

It is worth closing with the negative space, because a section like this can read as an argument that everything should move to a compute pass. Three categories genuinely should not.

Anything inherently sequential is the first. Label placement with collision resolution, route finding along a network, and topology repair all depend on the result of the previous decision, and a GPU that can run ten thousand lanes in parallel has nothing to offer a problem that can only run one. The GPU version of these is usually a reformulation rather than a port, and if the reformulation is not obvious the CPU version is the right answer.

Anything whose result the CPU needs immediately is the second. A hit test under the cursor, a feature id for a tooltip, a count for a legend the user is watching — each of these costs a readback, and a readback inside the frame is a synchronisation point that costs more than any amount of parallel arithmetic saved. If the answer must be known this frame, compute it on the CPU on the subset that matters.

Anything small is the third, and it is the one people get wrong most often. A dispatch has fixed overhead in the tens of microseconds, so a kernel over four hundred elements spends nearly all its time on the pass rather than the work. The threshold moves with hardware, but the shape of the rule does not: below a few thousand elements, measure before assuming the GPU wins.

Explore the Geometry Pipeline References

Each stage of the pipeline above has a dedicated, implementation-level reference:

Up: Spatial Visualization & WebGPU home