A data race happens when two invocations touch the same memory, at least one writes, and nothing orders them; the result then changes from run to run. This shader, run as 1,024 workgroups of 64, has one race on purpose and one that a barrier prevents:
// out: 3 x u32, dispatch 1024
struct Counts { plain: u32, atomic: atomic<u32>, wrong: atomic<u32> }
@group(0) @binding(0) var<storage, read_write> out: Counts;
var<workgroup> tile: array<u32, 64>;
@compute @workgroup_size(64) fn main(@builtin(global_invocation_id) g: vec3u,
@builtin(local_invocation_index) i: u32) {
out.plain += 1; // load, add, store: other invocations interleave
atomicAdd(&out.atomic, 1); // one indivisible read-modify-write
tile[i] = g.x; // publish a value in shared memory
workgroupBarrier(); // delete this line and rerun
let mirrored = tile[63 - i]; // read the value another invocation wrote
if (mirrored != g.x - i + 63 - i) { atomicAdd(&out.wrong, 1); }
}Output
out: 5, 65536, 0 gpu: 0.00 ms (median of 5)
The plain counter reached 5 instead of 65,536 (three more runs: 7, 6 and 8): invocations loaded the same old value and stored over each other. The atomic counter is always exact. Without the barrier, 3,520, 3,808 and 3,616 invocations in three runs read a tile slot before its owner wrote it, about one in twenty, so a race can hide for a long time. The cures: one output slot per invocation, atomics for shared counters, a barrier between writes and dependent reads, and a new dispatch for steps that need every workgroup.
<!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="340"></canvas>
<canvas id="labels" width="600" height="340"></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 = (barrier) => /* wgsl */ `
struct Counts { plain: u32, atomic: atomic<u32>, wrong: atomic<u32> }
@group(0) @binding(0) var<storage, read_write> out: Counts;
var<workgroup> tile: array<u32, 64>;
@compute @workgroup_size(64) fn main(@builtin(global_invocation_id) g: vec3u,
@builtin(local_invocation_index) i: u32) {
out.plain += 1; // load, add, store: other invocations interleave
atomicAdd(&out.atomic, 1); // one indivisible read-modify-write
tile[i] = g.x; // publish a value in shared memory
${barrier ? 'workgroupBarrier(); // every slot written before any read' : '// (no barrier: some reads come too early)'}
let mirrored = tile[63 - i]; // read the value another invocation wrote
if (mirrored != g.x - i + 63 - i) { atomicAdd(&out.wrong, 1); }
}`;
// Bars on a log scale, straight from the result buffer: 6 runs x 3 counters, 256 bytes apart.
const view = /* wgsl */ `
@group(0) @binding(0) var<storage> results: array<u32>;
struct Out { @builtin(position) pos: vec4f, @location(0) color: vec3f }
@vertex fn vs(@builtin(vertex_index) v: u32, @builtin(instance_index) i: u32) -> Out {
let q = vec2f(f32(v & 1), f32(v >> 1));
let run = i / 3; let kind = i % 3;
let value = f32(results[run * 64 + kind]);
let w = log2(value + 1) / 16; // 65,536 = full width
let y = 0.66 - f32(run) * 0.27 - f32(kind) * 0.075;
let colors = array(vec3f(0.85, 0.35, 0.25), vec3f(0.16, 0.56, 0.30), vec3f(0.55, 0.30, 0.65));
return Out(vec4f(-0.42 + q.x * w * 1.0, y + q.y * 0.06, 0, 1), colors[kind]);
}
@fragment fn fs(in: Out) -> @location(0) vec4f { return vec4f(in.color, 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, RUNS = 6, SLOT = 256; // storage binding offsets must be multiples of 256
const results = device.createBuffer({ size: RUNS * SLOT, usage: B.STORAGE | B.COPY_SRC });
const read = device.createBuffer({ size: RUNS * SLOT, usage: B.COPY_DST | B.MAP_READ });
const pipelines = [false, true].map((barrier) => device.createComputePipeline({ layout: 'auto',
compute: { module: device.createShaderModule({ code: code(barrier) }) } }));
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();
for (let run = 0; run < RUNS; run++) { // runs 0-2 without the barrier, 3-5 with it
const pipeline = pipelines[run < 3 ? 0 : 1];
const pass = encoder.beginComputePass();
pass.setPipeline(pipeline);
pass.setBindGroup(0, device.createBindGroup({ layout: pipeline.getBindGroupLayout(0),
entries: [{ binding: 0, resource: { buffer: results, offset: run * SLOT, size: 12 } }] }));
pass.dispatchWorkgroups(1024); // 1,024 workgroups of 64 = 65,536 invocations
pass.end();
}
const draw = encoder.beginRenderPass({ colorAttachments: [{ view: context.getCurrentTexture().createView(),
clearValue: [0.97, 0.96, 0.93, 1], loadOp: 'clear', storeOp: 'store' }] });
draw.setPipeline(render);
draw.setBindGroup(0, device.createBindGroup({ layout: render.getBindGroupLayout(0), entries: [{ binding: 0, resource: { buffer: results } }] }));
draw.draw(4, RUNS * 3);
draw.end();
encoder.copyBufferToBuffer(results, 0, read, 0, RUNS * SLOT);
device.queue.submit([encoder.finish()]);
await read.mapAsync(GPUMapMode.READ);
const all = new Uint32Array(read.getMappedRange().slice(0));
read.unmap();
ink.font = 'bold 12.5px system-ui, sans-serif'; ink.fillStyle = '#222';
ink.fillText('log scale; 65,536 is the right answer for both counters', 12, 18);
ink.font = '11.5px system-ui, sans-serif';
for (let run = 0; run < RUNS; run++) {
const y = 170 * (1 - 0.66 + run * 0.27) - 4;
ink.fillStyle = '#222'; ink.textAlign = 'right';
ink.fillText(`${run < 3 ? 'no barrier' : 'barrier'}, run ${run % 3 + 1}`, 166, y);
ink.textAlign = 'left';
for (let kind = 0; kind < 3; kind++) { // each bar's value at its end
const value = all[run * 64 + kind];
const barEnd = 174 + Math.log2(value + 1) / 16 * 300;
ink.fillStyle = '#444'; ink.font = '10.5px system-ui, sans-serif';
ink.fillText(value.toLocaleString('en-US'), barEnd + 6, 170 * (1 - 0.66 + run * 0.27 + kind * 0.075) + 1);
}
ink.font = '11.5px system-ui, sans-serif';
}
ink.textAlign = 'left';
[['plain +=', '#d95940'], ['atomicAdd', '#2a8f4d'], ['read before written', '#8c4da6']].forEach(([t, c], k) => {
ink.fillStyle = c; ink.fillRect(12 + k * 150, 326, 12, 10); ink.fillStyle = '#333'; ink.fillText(t, 28 + k * 150, 335);
});
}
main();
</script>