WebGPU Compute Shaders: Dispatching Workgroups, Choosing Sizes, and Reading Back Results

Why compute shaders matter in the browser

When you need to move a large amount of data through a simple transformation—doubling a vector, applying a filter, or performing a reduction—CPU code can become a bottleneck. WebGPU’s compute pipeline lets you launch thousands of lightweight GPU threads that run in parallel, all within a web page, without any plug‑ins.

Dispatching workgroups with dispatchWorkgroups

The core of a compute operation is the GPUComputePassEncoder.dispatchWorkgroups(x, y, z) call. It tells the GPU how many workgroups to launch in each dimension. Each workgroup contains a fixed number of invocations (threads) defined by the @workgroup_size attribute in WGSL. The total number of invocations is x * y * z * workgroupSizeProduct. Because the dispatch is independent of any rendering pipeline, you can run a compute pass as often as you like, even inside a render loop.

Choosing a portable workgroup size

WGSL requires that @workgroup_size be a compile‑time constant. The product of the three dimensions must not exceed the device limit maxComputeInvocationsPerWorkgroup. You can query this limit after you have a GPUDevice:

const adapter = await navigator.gpu.requestAdapter();
const device = await adapter.requestDevice();
console.log(device.limits.maxComputeInvocationsPerWorkgroup); // e.g. 256

A safe, widely‑supported choice is @workgroup_size(64) or @workgroup_size(16, 16). If you need a larger grid, split your work into multiple dispatches or use a 3‑D layout. Remember that the workgroup size also affects how much shared memory (workgroup memory) you can allocate.

Storage buffers and bind group layout

Data that the shader reads or writes lives in GPUBuffer objects. You must declare the buffer usage flags when you create the buffer:

• GPUBufferUsage.STORAGE – the buffer is a storage resource that the compute shader can read or write.
• GPUBufferUsage.COPY_DST – allows the host to copy data into the buffer via queue.writeBuffer.
• GPUBufferUsage.COPY_SRC – allows the GPU to copy data out of the buffer (used for readback).
• GPUBufferUsage.UNIFORM – used for small constant data such as the array length.
• GPUBufferUsage.MAP_READ – required if you want to map the buffer for reading back data.

In WGSL you bind the buffers in a bind group that matches a bind group layout. The layout declares the buffer type (storage, uniform, etc.) and the access mode (read, write, read_write). Here is a minimal layout for a two‑buffer compute shader that doubles an array:

@group(0) @binding(0) var  input  : array;
@group(0) @binding(1) var output : array;
@group(0) @binding(2) var config : Config;

and the corresponding Config structure:

struct Config {
  length : u32;
};

Worked example: doubling a Float32Array

Below is a full, end‑to‑end example that doubles each element of a Float32Array using WebGPU compute. The code is written for modern browsers that support WebGPU (Chrome 121+, Edge 121+, Safari 17.4+). It uses async/await for readability.

WGSL compute shader

// double.wgsl
@group(0) @binding(0) var  input  : array;
@group(0) @binding(1) var output : array;
@group(0) @binding(2) var config : Config;

struct Config {
  length : u32;
};

@compute @workgroup_size(64)
fn main() {
  let i : u32 = global_invocation_id.x;
  if (i < config.length) {
    output[i] = input[i] * 2.0;
  }
}

JavaScript setup

async function run() {
  if (!navigator.gpu) { alert('WebGPU not supported'); return; }
  const adapter = await navigator.gpu.requestAdapter();
  if (!adapter) { alert('No suitable GPU adapter found'); return; }
  const device = await adapter.requestDevice();

  const size = 1024; // number of elements
  const byteSize = size * 4; // Float32Array

  // Create input buffer and copy data
  const inputBuffer = device.createBuffer({
    size: byteSize,
    usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
    mappedAtCreation: false,
  });
  const inputArray = new Float32Array(size);
  for (let i = 0; i < size; ++i) inputArray[i] = Math.random();
  device.queue.writeBuffer(inputBuffer, 0, inputArray.buffer, 0, byteSize);

  // Create output buffer (readable after compute)
  const outputBuffer = device.createBuffer({
    size: byteSize,
    usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC,
    mappedAtCreation: false,
  });

  // Uniform buffer for array length
  const configBuffer = device.createBuffer({
    size: 4,
    usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST,
    mappedAtCreation: false,
  });
  const lengthArray = new Uint32Array([size]);
  device.queue.writeBuffer(configBuffer, 0, lengthArray.buffer, 0, 4);

  // Bind group layout and bind group
  const bindGroupLayout = device.createBindGroupLayout({
    entries: [
      { binding: 0, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'read-only-storage' } },
      { binding: 1, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'storage' } },
      { binding: 2, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'uniform' } },
    ],
  });

  const bindGroup = device.createBindGroup({
    layout: bindGroupLayout,
    entries: [
      { binding: 0, resource: { buffer: inputBuffer } },
      { binding: 1, resource: { buffer: outputBuffer } },
      { binding: 2, resource: { buffer: configBuffer } },
    ],
  });

  // Compile shader module and create pipeline
  const shaderModule = device.createShaderModule({ code: await fetch('double.wgsl').then(r => r.text()) });
  const pipeline = device.createComputePipeline({
    layout: device.createPipelineLayout({ bindGroupLayouts: [bindGroupLayout] }),
    compute: { module: shaderModule, entryPoint: 'main' },
  });

  // Command encoder and compute pass
  const commandEncoder = device.createCommandEncoder();
  const passEncoder = commandEncoder.beginComputePass();
  passEncoder.setPipeline(pipeline);
  passEncoder.setBindGroup(0, bindGroup);
  const workgroupCount = Math.ceil(size / 64);
  passEncoder.dispatchWorkgroups(workgroupCount);
  passEncoder.endPass();

  // Copy result to a staging buffer that can be mapped for reading
  const readBuffer = device.createBuffer({
    size: byteSize,
    usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ,
    mappedAtCreation: false,
  });
  commandEncoder.copyBufferToBuffer(outputBuffer, 0, readBuffer, 0, byteSize);

  // Submit the command buffer
  device.queue.submit([commandEncoder.finish()]);

  // Read back the data
  await readBuffer.mapAsync(GPUMapMode.READ);
  const arrayBuffer = readBuffer.getMappedRange();
  const result = new Float32Array(arrayBuffer.slice(0));
  readBuffer.unmap();

  // Verify correctness (CPU fallback)
  const cpuResult = inputArray.map(v => v * 2);
  const ok = result.every((v, i) => Math.abs(v - cpuResult[i]) < 1e-6);
  console.log('GPU result matches CPU?', ok);
}
run();

Why this example works

• Each invocation processes a single array element via global_invocation_id.x.
• The if (i < config.length) guard prevents out‑of‑bounds memory access when the array size is not a multiple of 64.
• We use a uniform buffer to pass the array length because WGSL cannot query buffer sizes at runtime.
• All buffers are created with the correct usage flags; the staging buffer is mapped after the GPU finishes the copy.
• We query device.limits.maxComputeInvocationsPerWorkgroup implicitly by choosing 64, a value that is guaranteed to be supported on current desktop GPUs.

Performance considerations and the break‑even point

Compute shaders are most advantageous when the data stays on the GPU for the entire pipeline. In the example above we copy the input data from the CPU to the GPU and read the output back to the CPU. The time spent on queue.writeBuffer and mapAsync can dominate the total runtime for small arrays. A rough rule of thumb is that the GPU starts to beat a single‑threaded CPU loop when you process several thousand elements, but you should benchmark on your target hardware.

To find the break‑even point:

1. Wrap the compute code in a function that accepts an array length.
2. Measure the total time (including write, dispatch, copy, and readback) using performance.now().
3. Do the same on the CPU with Array.prototype.map or Float32Array.forEach.
4. Plot the times and look for the crossover.

Because mapAsync returns a Promise, you can use await to block until the readback completes. If you are running multiple dispatches per frame, reuse the buffers and bind groups to amortize the allocation cost.

Trade‑offs and limitations

• Asynchronous readback: The GPU may still be busy when you try to read the buffer. mapAsync guarantees that you will see the latest data, but the latency can be hundreds of microseconds.
• Workgroup synchronization: workgroupBarrier() synchronizes only within a single workgroup. There is no cross‑group barrier in WGSL, so reductions that span the entire dispatch must be implemented with multiple passes or atomic operations.
• Feature detection: Some browsers enable optional features such as GPUFeatureName.subgroups only on certain GPUs. Always check adapter.features before requesting a device with those features.
• Limits vary: The maximum number of workgroups per dimension, the maximum size of a storage buffer, and the maximum number of storage buffers per shader stage differ between devices. Query device.limits and design your code to adapt.
• Security context: WebGPU is only available on HTTPS or localhost. Running the example on an insecure page will fail at navigator.gpu.

Actionable next steps

• Open a secure page and run if (navigator.gpu) { … } to confirm WebGPU support.
• After obtaining a GPUDevice, log device.limits.maxComputeInvocationsPerWorkgroup and device.limits.maxComputeWorkgroupsPerDimension to guide your workgroup size.
• Write a simple WGSL compute shader that processes one element per invocation and includes a bounds check.
• Create the necessary storage and uniform buffers with the correct usage flags.
• Dispatch the compute pass and read back the result with mapAsync.
• Benchmark against a CPU implementation for increasing input sizes to find your break‑even point.
• Implement graceful fallback to a CPU path if navigator.gpu is missing or the adapter request fails.

Once you are comfortable with the basics, you can explore more advanced patterns: tiled reductions with workgroup memory, inter‑group communication via atomics, or pipeline parallelism across multiple compute passes. WebGPU’s explicit API gives you fine‑grained control, but it also requires you to manage data movement and synchronization yourself. Start small, validate your assumptions on the target device, and scale up when the performance benefits become clear.