Skip to content
Fragment Colorv0.12.7
Blog

A particle system

Somewhere past a million, triangles stop paying for themselves. Shrinking one saves little, because at four pixels wide most of its cost goes into preparing it, before any pixel gets a color.

Here, eight million particles move every frame, and none of them is a triangle.

Compute shaders and storage buffers

Section titled “Compute shaders and storage buffers”

In this program, a compute shader moves every particle and counts it in a grid of 1280 × 720 cells. One full-screen triangle then colors each pixel by the count in its cell: the more particles a cell holds, the brighter it glows.

A compute shader runs a function for every item, in workgroups of 256 here. It reads and writes storage buffers: arrays on the GPU that keep their contents from frame to frame. One of them stores our particles, each as a position and a velocity. A second holds a count per grid cell. FragmentColor keeps one storage buffer per name and binds it to every shader that declares it. The clear, simulate and draw shaders all reach the same cells through counts.

Three passes per frame

Section titled “Three passes per frame”
fn setup() -> Result<State, FragmentColorError> {
    // All three shaders take their sizes from these WGSL constants.
    let sizes = format!(
        "const COUNT: u32 = {COUNT}u;
         const GRID = vec2<u32>({WIDTH}u, {HEIGHT}u);
         const CELLS: u32 = {WIDTH}u * {HEIGHT}u;"
    );
    let clear = Shader::new([sizes.as_str(), CLEAR])?;
    let simulate = Shader::new([sizes.as_str(), SIMULATE])?;
    let draw = Shader::new([sizes.as_str(), DRAW])?;


    // One workgroup handles 256 items, as `@workgroup_size(256)` declares.
    let clear_pass = Pass::from_shader("clear", &clear);
    clear_pass.set_compute_dispatch((WIDTH * HEIGHT).div_ceil(256), 1, 1);


    let simulate_pass = Pass::from_shader("simulate", &simulate);
    simulate_pass.set_compute_dispatch(COUNT.div_ceil(256), 1, 1);
    simulate_pass.require(&clear_pass)?;


    let output = Pass::from_shader("draw", &draw);
    output.require(&simulate_pass)?;


    Ok(State {
        simulate,
        draw,
        output,
    })
}

The sizes string holds three WGSL constants, formatted from Rust numbers. Each Shader list puts them ahead of the shader’s own source. COUNT makes the particle buffer eight million entries of 16 bytes, 128 MB in all. That’s just under 128 MiB (134,217,728 bytes), and every WebGPU device must allow a storage buffer binding up to that size.

A Pass made from a compute shader dispatches workgroups instead of drawing. set_compute_dispatch says how many, in up to three dimensions: enough groups of 256 along the first to cover every item, and 1 along the other two. The require calls chain the passes in order, and frame hands only the last one to render.

clearsimulatedrawcomputecomputerenderparticles8,000,000 × 16 bytescounts1280 × 720 cellsatomicStoreread, writeatomicAddread
One frame. The passes run left to right, and each one reaches the buffers it declares by name.

Clearing the grid

Section titled “Clearing the grid”

The first pass sets every count back to zero:

@group(0) @binding(0) var<storage, read_write> counts: array<atomic<u32>, CELLS>;


@compute @workgroup_size(256)
fn cs_main(@builtin(global_invocation_id) id: vec3<u32>) {
    if (id.x < CELLS) {
        atomicStore(&counts[id.x], 0u);
    }
}

Moving the particles

Section titled “Moving the particles”
struct Particle {
    position: vec2<f32>,
    velocity: vec2<f32>,
};


@group(0) @binding(0) var<storage, read_write> particles: array<Particle, COUNT>;
@group(0) @binding(1) var<storage, read_write> counts: array<atomic<u32>, CELLS>;
@group(0) @binding(2) var<uniform> resolution: vec2<f32>;
@group(0) @binding(3) var<uniform> time: f32;


fn hash(value: u32) -> u32 {
    let state = value * 747796405u + 2891336453u;
    let word = ((state >> ((state >> 28u) + 4u)) ^ state) * 277803737u;
    return (word >> 22u) ^ word;
}


fn random(seed: u32) -> f32 {
    return f32(hash(seed)) / 4294967295.0;
}


@compute @workgroup_size(256)
fn cs_main(@builtin(global_invocation_id) id: vec3<u32>) {
    let index = id.x;
    if (index >= COUNT) {
        return;
    }
    var particle = particles[index];


    // A new buffer holds zeros, so the first frame starts every particle
    // circling the center. After that, about one in 300 starts over on each frame.
    // Mixing in the bit pattern of time gives each frame its own dice.
    let seed = hash(index ^ bitcast<u32>(time));
    let fresh = all(particle.velocity == vec2<f32>(0.0));
    if (fresh || random(seed) < 0.003) {
        let radius = 0.2 + 0.8 * sqrt(random(seed + 1u));
        let angle = random(seed + 2u) * 6.283;
        let spot = vec2<f32>(cos(angle), sin(angle));
        particle.position = spot * radius;
        particle.velocity = vec2<f32>(-spot.y, spot.x) * 0.6;
    }


    // An attractor wanders through the middle and pulls every particle.
    // The pull weakens with distance; 0.05 keeps it finite right at the attractor.
    let attractor = vec2<f32>(cos(time * 0.37), sin(time * 0.61)) * 0.4;
    let offset = attractor - particle.position;
    let pull = offset / (dot(offset, offset) + 0.05);


    // Each frame advances the simulation by 1/60 of a second, with a little drag.
    let step = 1.0 / 60.0;
    particle.velocity = (particle.velocity + pull * step) * 0.995;
    particle.position += particle.velocity * step;
    particles[index] = particle;


    // Count the particle in the grid cell under it.
    let fit = min(resolution.x, resolution.y) / resolution;
    let uv = particle.position * fit * 0.5 + 0.5;
    if (all(uv >= vec2<f32>(0.0)) && all(uv < vec2<f32>(1.0))) {
        let cell = vec2<u32>(uv * vec2<f32>(GRID));
        atomicAdd(&counts[cell.y * GRID.x + cell.x], 1u);
    }
}

global_invocation_id numbers the runs, as instance_index numbers the copies of a mesh.

The ID is a vec3<u32> to match the three dispatch dimensions, and a flat list of particles needs only id.x. Rounding the dispatch up can start a few runs past the end. The if at the top of each compute shader skips them.

Every new storage buffer holds zeros. A particle standing perfectly still takes that as its cue to pick an orbit. hash scrambles its index and the time into well-mixed bits, and random turns those into numbers from 0 to 1: one for the orbit’s radius, one for its angle. After that, every frame adds the pull of the attractor, a moving point near the middle, to the velocity. The particle then moves by that velocity for 1/60 of a second. About one particle in 300 also starts over on each frame, which keeps the attractor from swallowing them all.

centerattractorvelocitypull
One particle's next two seconds. Its velocity carries it around the center while the pull bends its path toward the attractor. The attractor itself wanders along the dashed curve.

The grid is a flat array, one row after another: the cell in column x and row y is element y * GRID.x + x.

More than a thousand particles can share a cell during one frame. atomicAdd makes sure each of them counts, where an ordinary += from several runs at once could lose some. Like atomicStore in the first pass, it takes a pointer to the cell, written with &.

4141113212221111012345670123xycounts[2 * 8 + 5]
A grid with GRID.x set to 8. Every particle adds 1 to the cell under it. The marked cell is counts[21].

Drawing the counts

Section titled “Drawing the counts”
@group(0) @binding(0) var<storage, read> counts: array<u32, CELLS>;
@group(0) @binding(1) var<uniform> resolution: vec2<f32>;


@vertex
fn vs_main(@builtin(vertex_index) i: u32) -> @builtin(position) vec4<f32> {
    // One triangle large enough to cover the whole target.
    var corners = array<vec2<f32>, 3>(
        vec2<f32>(-1.0, -1.0), vec2<f32>(3.0, -1.0), vec2<f32>(-1.0, 3.0),
    );
    return vec4<f32>(corners[i], 0.0, 1.0);
}


@fragment
fn fs_main(@builtin(position) pixel: vec4<f32>) -> @location(0) vec4<f32> {
    // Pixel rows count down from the top, and grid rows count up from the bottom.
    let uv = vec2<f32>(pixel.x, resolution.y - pixel.y) / resolution;
    let cell = min(vec2<u32>(uv * vec2<f32>(GRID)), GRID - 1u);
    let count = f32(counts[cell.y * GRID.x + cell.x]);


    // A logarithm keeps a lone particle visible next to a crowd of a thousand.
    let heat = clamp(log2(1.0 + count) / 10.0, 0.0, 1.0);
    let night = vec3<f32>(0.02, 0.02, 0.06);
    let ember = vec3<f32>(0.95, 0.45, 0.25);
    return vec4<f32>(mix(night, ember, heat) + heat * heat * 0.5, 1.0);
}

The draw shader just reads, so it declares counts as plain u32 with read access: the same bytes without the atomic wrapper. In a fragment stage, @builtin(position) holds the pixel’s coordinates in the target, counted from the top-left corner, which is why fs_main flips y before it looks up a cell. The grid stretches to fill the canvas, while fit in the simulate shader keeps the swirl round. Only at 1280 × 720 does a cell match exactly one pixel.

pixelcellxy(0, 0)xy(0, 0)(384.5, 180.5)(384, 539)
One pixel on a 1280 by 720 canvas. @builtin(position) holds its center and counts rows from the top. The grid counts them from the bottom.

On every frame, the program sets the time and the canvas size and renders output:

fn frame(
    state: &State,
    renderer: &Renderer,
    target: &impl Target,
    time: f32,
) -> Result<(), FragmentColorError> {
    let size = target.size();
    let resolution = [size.width as f32, size.height as f32];
    state.simulate.set("resolution", resolution)?;
    state.simulate.set("time", time)?;
    state.draw.set("resolution", resolution)?;
    renderer.render(&state.output, target)?;
    Ok(())
}

Most demos on this site run on WebGPU, or on WebGL2 where a browser lacks it. This one needs WebGPU, since WebGL2 has no compute shaders, and shows a note instead where it’s missing.

Where the time goes

Section titled “Where the time goes”

A particle here is one run of the compute shader and one atomicAdd, with no triangle for the GPU to prepare. Item for item, it costs about a sixteenth as much as an instanced triangle. This program moves eight times as many particles as the last step of Hello, million triangles! draws triangles, in about half the time.

Your own numbers depend on your GPU. The frame budget section of Hello, million triangles! shows how to measure them, and FC_TUTORIAL_FRAMES=120 ./fc example tutorial_particle_system_splat prints this program’s average frame time.

What just happened

Section titled “What just happened”
  • A compute shader runs a function per item and needs no geometry.
  • Storage buffers keep their data on the GPU between frames, and every shader that declares one by name shares it.
  • atomicAdd lets many runs count into the same cell without losing any.
  • One full-screen draw turns the counts into a picture.