Made by AI, Directed by David Stein

Search
Creative Coding

Getting Started With WebGPU Compute Shaders

Device, buffers, a bind group, a WGSL kernel and a readback — the five pieces you need before any of the clever stuff. Plus the three mistakes that make a first compute pass silently return zeros.

6 min read

A fragment shader answers “what colour is this pixel.” A compute shader answers whatever you want, for as many items as you want, in parallel, with no geometry involved. Particle integration, FFTs, flocking, physics, image filters that need neighbours, sorting — all of it belongs in compute.

WebGPU is the first browser API that gives you compute properly. WebGL never had it. This is the minimum pipeline, and the parts that trip people up.

Before you start

You need a browser with WebGPU. Check for it rather than assuming:

if (!navigator.gpu) throw new Error('WebGPU not available');

On a laptop with two GPUs, the browser may give you the integrated one. That is usually fine for learning and not fine for benchmarking.

1. Adapter and device

Two steps, both async. The adapter is the physical GPU; the device is your handle on it.

const adapter = await navigator.gpu.requestAdapter();
const device  = await adapter.requestDevice();

Always attach an error handler before you do anything else, because WebGPU validation errors are informative and easy to miss:

device.addEventListener('uncapturederror', e => console.error(e.error.message));

2. The kernel

WGSL, WebGPU’s shader language. A trivial one that doubles every element:

@group(0) @binding(0) var<storage, read_write> data: array<f32>;

@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
  let i = gid.x;
  if (i >= arrayLength(&data)) { return; }
  data[i] = data[i] * 2.0;
}

Three things in there matter more than they look.

@workgroup_size(64) is how many invocations run as one tightly coupled group. 64 is a good default: it is a multiple of both 32 (NVIDIA warp) and 64 (AMD wavefront), so you waste no lanes on either. Use 256 for simple elementwise work, 8×8 for 2D image work.

global_invocation_id is this invocation’s index across the entire dispatch. There is also local_invocation_id (index within the workgroup) and workgroup_id. You will want all three eventually; start with the global one.

The bounds check is not optional. Dispatches are in whole workgroups, so if you have 100 elements and a workgroup size of 64 you dispatch 2 groups = 128 invocations, and 28 of them must do nothing. Omit the check and you corrupt memory past the end of the buffer.

3. Buffers

One storage buffer the shader writes, and one staging buffer you can read from the CPU. They have to be separate — this is the single most common beginner surprise.

const input = new Float32Array([1, 2, 3, 4, 5]);

const gpuBuffer = device.createBuffer({
  size: input.byteLength,
  usage: GPUBufferUsage.STORAGE
       | GPUBufferUsage.COPY_DST
       | GPUBufferUsage.COPY_SRC,
});
device.queue.writeBuffer(gpuBuffer, 0, input);

const readBuffer = device.createBuffer({
  size: input.byteLength,
  usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ,
});

MAP_READ cannot be combined with STORAGE. That restriction is the reason for the two-buffer dance: the GPU writes the storage buffer, you copy it into the mappable one, then you read that.

4. Pipeline and bind group

const module = device.createShaderModule({ code: wgsl });

const pipeline = device.createComputePipeline({
  layout: 'auto',
  compute: { module, entryPoint: 'main' },
});

const bindGroup = device.createBindGroup({
  layout: pipeline.getBindGroupLayout(0),
  entries: [{ binding: 0, resource: { buffer: gpuBuffer } }],
});

layout: 'auto' derives the bind group layout from the shader. It is the right choice while you are learning. You will outgrow it when you want to share one layout across several pipelines.

5. Encode, dispatch, submit

const encoder = device.createCommandEncoder();

const pass = encoder.beginComputePass();
pass.setPipeline(pipeline);
pass.setBindGroup(0, bindGroup);
pass.dispatchWorkgroups(Math.ceil(input.length / 64));
pass.end();

encoder.copyBufferToBuffer(gpuBuffer, 0, readBuffer, 0, input.byteLength);
device.queue.submit([encoder.finish()]);

dispatchWorkgroups takes a count of workgroups, not a count of items. This is mistake number two. Passing input.length there dispatches 5 × 64 = 320 invocations instead of 64, and only the bounds check saves you.

Note that the copy is encoded outside the compute pass but inside the same encoder. Ordering within one encoder is guaranteed, so the copy sees the finished compute work.

6. Read it back

await readBuffer.mapAsync(GPUMapMode.READ);
const out = new Float32Array(readBuffer.getMappedRange().slice(0));
readBuffer.unmap();
console.log(out);   // [2, 4, 6, 8, 10]

.slice(0) copies before you unmap. Mistake number three is keeping a view onto the mapped range, unmapping, and then finding your array has turned into zeros — the memory was borrowed, not given.

Where the actual speed comes from

Getting a correct result is step one. Making it fast is a different skill, and three ideas cover most of it.

Workgroup shared memory. var<workgroup> is fast scratch memory visible to all invocations in a group. Any algorithm where neighbours read the same data — blurs, reductions, FFT butterflies — loads a tile into shared memory once and then reads it many times. This is what web-ocean means when it says its radix-2 inverse FFTs work in workgroup memory.

Barriers. workgroupBarrier() makes every invocation in the group wait. You need it between writing shared memory and reading what your neighbours wrote. Every invocation in the group must reach the same barrier — putting one inside a divergent if is undefined behaviour.

Don’t read back. The readback in step 6 is for debugging. A per-frame mapAsync stalls the pipeline. In real use the compute output stays on the GPU and gets consumed by the next pass as a vertex, storage or texture binding.

Debugging, honestly

There is no printf. The techniques that work:

  • Write diagnostics into a buffer and read it once, offline. Reserve a few floats for “what did invocation 0 actually see.”
  • Run with one invocation. Dispatch a single workgroup of size 1 and check the arithmetic before you let 100,000 threads loose on it.
  • Check your bounds first. If the output is right for the first N elements and garbage after, it is the dispatch count or the bounds check.
  • Check for zeros vs. untouched. An all-zero result usually means the dispatch never ran or the bind group pointed at the wrong buffer. A partially-correct result means an indexing bug.

Next

Once this runs, port something you already have. A particle system is the natural first move: positions and velocities in two storage buffers, one compute pass to integrate, then render the position buffer directly as vertex data without a round trip through the CPU. That last part — never copying back — is the whole reason to be here.

Sources

Related Reading