Compute & Particles

A vertex shader answers a question about a vertex and a fragment shader answers a question about a pixel. Neither chooses when it runs: one is called because something is about to be rasterised, the other because something is about to be blended into a framebuffer, and the answer is consumed by the stage that comes next. A compute shader is the same hardware with all of that taken away. It runs as many times as you ask it to, and its output is whatever it left behind in memory.

This lab runs a compute shader over a field of particles. Every position and velocity lives in a buffer the GPU owns; a compute pass steps all of them; the vertex stage then reads the same buffer back, six vertices at a time, and draws them. The CPU writes forty-eight bytes a frame and issues one dispatch and one draw, whether there are five thousand particles or a hundred thousand.

The GPU has a stage that draws nothing#

The entry point carries the attribute @compute @workgroup_size(64) and takes one argument, @builtin(global_invocation_id). That argument is an index, and it is the only input: there is no vertex to place and no pixel to colour. Invocation i loads particles[i], works out a force towards the attractor, advances a velocity and a position, and writes the particle back. The force is a softened inverse square with a tangential term added — (dir + tangent * swirl) * attraction / (dist * dist + 0.08) — where the tangent is what turns a collapse into an orbit and the + 0.08 is what stops the pull going to infinity when a particle crosses the centre.

WebGL has no such stage. Not a slow one, not a restricted one — the concept is absent from the API. The way this was done for a decade is to put the state where a fragment shader can write: encode positions into a floating-point texture, draw a full-screen quad so that one fragment lands on each texel, step the simulation there, and write into a second texture bound to a framebuffer. The two textures swap roles every frame, because a shader cannot read the texture it is writing to. The result is then sampled again by a vertex shader to get the positions back as geometry.

The workaround also decides where a result is allowed to go. A fragment shader writes to the pixel it was rasterised at and nowhere else, so every algorithm has to be phrased as a gather: each output asks which inputs it needs. A compute invocation writes wherever it likes in the buffer. This lab does not use that freedom — each invocation writes back to the slot it read from — but it is the freedom a sort, a spatial hash or a collision grid is built on, and it is why the stage exists at all rather than as a faster way of doing what the fragment stage already did.

The particles live in memory the GPU owns#

A Particle here is four floats — pos: vec2f and vel: vec2f — so sixteen bytes. The buffer holding them is created once, at the largest size the lab will ever need: 120,000 particles, 1,920,000 bytes. The count slider allocates nothing. It changes how much of that allocation is stepped and how much of it is drawn.

Two declarations point at the one buffer. var<storage, read_write> particles is what the compute stage sees; var<storage, read> readParticles is what the vertex stage sees. Same memory, two access modes, and the read-only one is not a courtesy: WebGPU does not permit a writable storage buffer in the vertex stage at all, so that binding has to be declared read-only-storage. A uniform buffer cannot stand in for either of them — uniforms are read-only to the shader and sized for a handful of values, which is exactly the job the lab’s other binding does.

Because the buffer outlives the frame, it also outlives what is on screen. Particles beyond the count are neither stepped nor drawn, and they keep whatever values they last held. Drag the particle slider up in the instrument below and they rejoin from wherever they were left; on a freshly seeded field that is a ring appearing at the outer edge, because the seed walks outward as the index rises. Nothing was recomputed to bring them back. They were in memory the whole time.

The one moment per-particle data crosses from the CPU is Reseed, which uploads the whole 1.92-megabyte seed array in a single write. It happens when you press a button, not sixty times a second.

A dispatch counts workgroups, not particles#

The call that starts the compute pass is dispatchWorkgroups(Math.ceil(count / 64)). It does not take a number of particles. It takes a number of workgroups, and the workgroup size — 64 here — is fixed in the shader at compile time, so the two have to be reconciled by rounding up. A hundred thousand particles divided by sixty-four is 1,562.5, which becomes 1,563 workgroups and 100,032 invocations: thirty-two more than there are particles.

The first line of the shader body is if (i >= u32(params.count)) { return; }. Every compute entry point written against a count that is not a multiple of the workgroup size has that line, or has a bug — here the extra thirty-two would step particles that are not being drawn, and at the top of the slider, where 120,000 divides by 64 exactly, none of them return at all.

Sixty-four is a choice, and this shader gives no algorithmic reason for it. The invocations of a workgroup are scheduled together and can share a var<workgroup> block of fast memory; nothing here talks to a neighbour, so the size is picked to sit well on the hardware rather than to fit the problem. The number of workgroups, by contrast, is a runtime argument with room to spare: no WebGPU implementation may report a limit below 65,535 per dimension, so a single dispatch of 64-wide groups covers four million particles. The largest dispatch this lab ever issues is 1,875.

The vertex stage reads the buffer the compute pass wrote#

The render pipeline declares no vertex buffers at all, and the vertex shader takes no attributes — its only input is @builtin(vertex_index). The draw call is draw(count * 6), which at a hundred thousand particles is 600,000 vertices and 200,000 triangles, not one of which is stored anywhere. vertex_index / 6u selects the particle out of the storage buffer and vertex_index % 6u selects a corner from a six-element array of offsets: two triangles making a square. The offset is scaled by the size control, and its x is divided by the canvas aspect so the sprite stays square on a wide canvas rather than stretching with it.

The square becomes a disc in the fragment shader, which carries the corner offset through as a uv and runs if (r > 1.0) { discard; } — the four corners are rasterised and thrown away. The colour is the particle’s speed mapped between a cool blue and a warm orange, which is the only reason the structure of the field is legible at all: velocity is otherwise invisible in a still frame.

Count what left the CPU while that happened. Forty-eight bytes: eleven floats — the attractor’s two coordinates, the six control values, the frame’s timestep, the canvas aspect and a flag for the theme — padded to a multiple of sixteen, because a uniform block has to be. One compute pass, one render pass, one dispatch, one draw. Move the particle slider from 5,000 to 100,000 and that list is unchanged; a single number inside the block is different. The GPU does twenty times the work and the CPU never notices, which is what it means to say the data lives on the GPU.

That is a claim about the CPU, not about the GPU, and it generalises past particles. The next lab takes the same measurement from the other direction — Draw Calls & Instancing draws one mesh ten thousand times and shows the cost sitting in the number of calls rather than the number of triangles.

Now run all of it at once#

Below is the whole thing: six sliders, a running toggle and a reseed button, with a readout that turns the particle count into workgroups and invocations as you drag it. Move the pointer over the canvas to take hold of the attractor; leave the canvas and it goes back to drifting on its own. The presets are the fastest way in — one of them switches the swirl off so the field collapses into a single dot, which is worth seeing once, because it is the tangential term rather than the pull that makes the shape.

If the browser has no WebGPU, the canvas will say so and stop. There is no fallback, and that is the honest position: the other labs here run on WebGL because their mathematics does not care which API draws it. This one is a stage WebGL does not have.