An atomic is a read-modify-write no other invocation can interrupt. WGSL has atomicLoad, atomicStore, atomicAdd, atomicSub, atomicMax, atomicMin, bitwise atomicAnd/Or/Xor, atomicExchange and atomicCompareExchangeWeak, on atomic<u32> and atomic<i32> in storage or workgroup memory; all but the store return the old value, which is how GPU-Built Arguments appended to a list. atomic<f32> fails ("'atomic' only supports 'i32', 'u32' or 'vec2u' types"), so sum prices as integer cents. Atomics on one address serialize, so count in workgroup memory first and merge each workgroup's result with one storage atomic per bin. Here 16 workgroups build a star-rating histogram of 65,536 simulated reviews:
// out: 6 x u32, dispatch 16
@group(0) @binding(0) var<storage, read_write> out: array<atomic<u32>>;
var<workgroup> bins: array<atomic<u32>, 5>; // one histogram per workgroup
fn stars(review: u32) -> u32 { // a stand-in review score, 1 to 5
var h = review * 747796405u + 2891336453u; // PCG-style integer hash
h = ((h >> ((h >> 28u) + 4u)) ^ h) * 277803737u;
return min(5u, 1u + (h >> 22u) % 6u); // 1 to 6, so 5 stars twice as often
}
@compute @workgroup_size(256) fn main(@builtin(global_invocation_id) g: vec3u,
@builtin(local_invocation_index) i: u32) {
for (var k = 0u; k < 16; k++) { // 16 x 256 x 16 = 65,536 reviews
atomicAdd(&bins[stars(g.x * 16 + k) - 1], 1); // fast on-chip atomics
}
workgroupBarrier();
if (i < 5) {
let count = atomicLoad(&bins[i]);
atomicAdd(&out[i], count); // one storage atomic per bin
atomicMax(&out[5], count); // the largest per-workgroup bin
}
}Output
out: 10969, 10861, 11059, 10952, 21695, 1421 gpu: 0.00 ms (median of 5)
The bins match a JavaScript loop over the same hash; out[5] is the largest bin one workgroup merged. That took 80 storage atomics (16 workgroups x 5 bins) instead of 65,536. For updates no built-in expresses, loop on atomicCompareExchangeWeak(), which may fail spuriously: check exchanged and retry.
<!doctype html>
<style>
body { margin: 0; background: #f7f4ee; font: 14px system-ui, sans-serif; }
.stage { position: relative; width: 100%; max-width: 600px; }
.stage canvas { display: block; width: 100%; }
.stage canvas + canvas { position: absolute; inset: 0; pointer-events: none; }
</style>
<div class="stage">
<canvas id="view" width="600" height="330"></canvas>
<canvas id="labels" width="600" height="330"></canvas>
</div>
<script>
const canvas = document.getElementById('view');
const ink = document.getElementById('labels').getContext('2d');
function showMessage(text) { // 2D fallback when WebGPU is missing
const ctx = canvas.getContext('2d');
ctx.fillStyle = '#fbeaea'; ctx.fillRect(0, 0, canvas.width, canvas.height);
ctx.fillStyle = '#8a2b2b'; ctx.font = '18px system-ui, sans-serif'; ctx.textAlign = 'center';
ctx.fillText(text, canvas.width / 2, canvas.height / 2);
}
const code = /* wgsl */ `
@group(0) @binding(0) var<storage, read_write> out: array<atomic<u32>>;
var<workgroup> bins: array<atomic<u32>, 5>; // one histogram per workgroup
fn stars(review: u32) -> u32 { // a stand-in review score, 1 to 5
var h = review * 747796405u + 2891336453u; // PCG-style integer hash
h = ((h >> ((h >> 28u) + 4u)) ^ h) * 277803737u;
return min(5u, 1u + (h >> 22u) % 6u); // 1 to 6, so 5 stars twice as often
}
@compute @workgroup_size(256) fn main(@builtin(global_invocation_id) g: vec3u,
@builtin(local_invocation_index) i: u32) {
for (var k = 0u; k < 16; k++) { // 16 x 256 x 16 = 65,536 reviews
atomicAdd(&bins[stars(g.x * 16 + k) - 1], 1); // fast on-chip atomics
}
workgroupBarrier();
if (i < 5) {
let count = atomicLoad(&bins[i]);
atomicAdd(&out[i], count); // one storage atomic per bin
atomicMax(&out[5], count); // the largest per-workgroup bin
}
}`;
const view = /* wgsl */ `
@group(0) @binding(0) var<storage> out: array<u32>;
struct Out { @builtin(position) pos: vec4f, @location(0) uv: vec2f, @location(1) @interpolate(flat) bin: u32 }
@vertex fn vs(@builtin(vertex_index) v: u32, @builtin(instance_index) i: u32) -> Out {
let q = vec2f(f32(v & 1), f32(v >> 1));
let w = f32(out[i]) / 25000; // horizontal bars straight from the counts
return Out(vec4f(-0.62 + q.x * w * 1.5, 0.52 - f32(i) * 0.3 + q.y * 0.2, 0, 1), q, i);
}
@fragment fn fs(in: Out) -> @location(0) vec4f {
return vec4f(mix(vec3f(0.95, 0.72, 0.15), vec3f(0.85, 0.50, 0.10), in.uv.x), 1);
}`;
async function main() {
const adapter = await navigator.gpu?.requestAdapter();
if (!adapter) return showMessage('WebGPU is not available in this browser');
const device = await adapter.requestDevice();
const context = canvas.getContext('webgpu');
const format = navigator.gpu.getPreferredCanvasFormat();
context.configure({ device, format });
const B = GPUBufferUsage;
const out = device.createBuffer({ size: 24, usage: B.STORAGE | B.COPY_SRC });
const read = device.createBuffer({ size: 24, usage: B.COPY_DST | B.MAP_READ });
const compute = device.createComputePipeline({ layout: 'auto', compute: { module: device.createShaderModule({ code }) } });
const module = device.createShaderModule({ code: view });
const render = device.createRenderPipeline({ layout: 'auto', primitive: { topology: 'triangle-strip' },
vertex: { module }, fragment: { module, targets: [{ format }] } });
const encoder = device.createCommandEncoder();
const cp = encoder.beginComputePass();
cp.setPipeline(compute);
cp.setBindGroup(0, device.createBindGroup({ layout: compute.getBindGroupLayout(0), entries: [{ binding: 0, resource: { buffer: out } }] }));
cp.dispatchWorkgroups(16);
cp.end();
const pass = encoder.beginRenderPass({ colorAttachments: [{ view: context.getCurrentTexture().createView(),
clearValue: [0.97, 0.96, 0.93, 1], loadOp: 'clear', storeOp: 'store' }] });
pass.setPipeline(render);
pass.setBindGroup(0, device.createBindGroup({ layout: render.getBindGroupLayout(0), entries: [{ binding: 0, resource: { buffer: out } }] }));
pass.draw(4, 5);
pass.end();
encoder.copyBufferToBuffer(out, 0, read, 0, 24);
device.queue.submit([encoder.finish()]);
await read.mapAsync(GPUMapMode.READ);
const bins = new Uint32Array(read.getMappedRange().slice(0));
read.unmap();
// The same hash in JavaScript, to check the GPU's counts.
const js = [0, 0, 0, 0, 0];
for (let r = 0; r < 65536; r++) {
let h = (Math.imul(r, 747796405) + 2891336453) >>> 0;
h = Math.imul(((h >>> ((h >>> 28) + 4)) ^ h) >>> 0, 277803737) >>> 0;
js[Math.min(5, 1 + ((h >>> 22) % 6)) - 1]++;
}
ink.font = '13px system-ui, sans-serif';
for (let i = 0; i < 5; i++) {
const y = 67 + i * 49.5; // the middle of bar i
ink.fillStyle = '#b07a10'; ink.textAlign = 'right'; ink.fillText('★'.repeat(i + 1), 106, y);
ink.fillStyle = '#333'; ink.textAlign = 'left'; ink.fillText(bins[i].toLocaleString('en-US'), (0.38 + bins[i] / 25000 * 1.5) * 300 + 8, y);
}
ink.fillStyle = '#222'; ink.font = 'bold 13px system-ui, sans-serif';
ink.fillText('65,536 reviews, 16 workgroups of 256', 12, 24);
ink.font = '12px system-ui, sans-serif'; ink.fillStyle = '#444';
ink.fillText(`80 storage atomics (16 x 5) instead of 65,536; largest per-workgroup bin (atomicMax): ${bins[5]}`, 12, 44);
const same = js.every((c, i) => c === bins[i]);
ink.fillStyle = same ? '#1e6b3a' : '#8a2b2b';
ink.fillText(same ? 'A JavaScript loop over the same hash gives the same counts.' : `JavaScript counts: ${js.join(', ')}`, 12, 322);
}
main();
</script>