kayf.dev

03 / 03 GPU compute

WebGPU Compute Pipeline

A browser-based procedural generation engine. Perlin noise and Marching Cubes run entirely in WebGPU compute shaders, so infinite terrain streams in without geometry ever crossing back to the CPU.

Dark faceted terrain generated by Marching Cubes, with glowing chunk seams, one chunk mid-build as a wireframe and empty chunk volumes waiting beyond it
Domain
Procedural geometry
Stack
WebGPU · WGSL · C++ → WASM
Technique
Marching Cubes · fBm Perlin noise

01 Results

Generates 64×64×64 voxel chunks in < 2 ms, entirely on the GPU. Zero CPU–GPU memory stalling.

Voxels per chunk
64³
Per chunk, GPU only
< 2 ms
CPU–GPU memory stalls
0

02 Problem

Procedural terrain on the web usually meshes on the CPU — often in WASM — and uploads the result. Every chunk pays for the copy across the bus, blocks a thread while it builds, and caps how fast new ground can appear. Meanwhile the GPU sits idle, waiting for geometry it could have produced itself.

03 Architecture

The CPU decides what to build; the GPU samples, meshes and draws it. Nothing is ever read back.

  1. 01

    Chunk scheduler

    C++ in WASM streams chunks around the camera by distance and visibility.

  2. 02

    Density pass

    fBm Perlin noise evaluated per voxel into a 3D storage texture.

  3. 03

    Classify & compact

    A Marching Cubes case per cell; a prefix sum packs only the cells that emit triangles.

  4. 04

    Mesh pass

    Vertices and gradient normals written straight into a pooled GPU vertex buffer.

  5. 05

    Indirect draw

    Vertex counts written on the GPU feed drawIndirect. No readback, no stall.

04 Engineering

  1. 01

    Zero readback

    The CPU never learns how many triangles a chunk produced. The mesh pass writes the count into an indirect-arguments buffer and the render pass consumes it with drawIndirect — noise to pixels without a single mapAsync.

  2. 02

    Stream compaction, not atomics

    Most cells in a chunk are empty. A workgroup-level exclusive scan over per-cell vertex counts gives each non-empty cell its write offset, instead of every thread contending on one global atomic counter.

  3. 03

    Seamless by construction

    Each chunk samples a one-voxel apron of its neighbours’ density. Shared faces produce identical vertices and normals, so chunks never crack and never have to talk to each other.

  4. 04

    C++ where the logic branches

    Scheduling, LOD selection and culling are branchy and cache-sensitive, so they run as C++ compiled to WebAssembly and hand the GPU a compact list of chunk jobs each frame.

05 Excerpt

shaders/density.wgslWGSL
@group(0) @binding(0) var<uniform> chunk : Chunk;@group(0) @binding(1) var density : texture_storage_3d<r32float, write>; // One invocation per density sample: the chunk's cells plus a one-voxel// apron, so neighbouring chunks agree on every shared vertex.@compute @workgroup_size(4, 4, 4)fn sampleDensity(@builtin(global_invocation_id) id : vec3u) {  if (any(id >= chunk.samples)) { return; }  let p = chunk.origin + vec3f(id) * chunk.voxel;   var d = -p.y * chunk.ground;  // half-space: solid below the ground plane  var amplitude = 1.0;  var frequency = chunk.frequency;  for (var octave = 0u; octave < 6u; octave++) {    d += perlin3(p * frequency) * amplitude;  // fBm over 3D Perlin noise    amplitude *= 0.5;    frequency *= 2.0;  }  textureStore(density, id, vec4f(d, 0.0, 0.0, 0.0));}

The density texture never leaves the GPU: the classify, scan and mesh passes read it in place.

Close-up of a wireframe chunk of Marching Cubes triangles beside shaded terrain
A chunk mid-build: Marching Cubes triangles as the mesh pass emits them, before the chunk joins the shaded terrain.

06 Specification

API
WebGPU compute and render, one command encoder per frame
Chunks
64 × 64 × 64 voxels, generated in < 2 ms on the GPU
Kernels
WGSL, 4×4×4 workgroups
Meshing
Marching Cubes with prefix-sum stream compaction
Draw path
drawIndirect, zero CPU readback
Host
C++20 compiled to WebAssembly
Memory
Pooled vertex slots, no allocation while streaming