Workgroup memory, barriers, and the bank conflict that ate your bandwidth

In Part 1 every invocation did its job alone — read one element, write one element, done. The GPU was happy, but it was also coasting. The hardware Chrome sat you on top of (your Apple M2, your RTX 4070, your iGPU) has a piece of memory none of our shader has touched yet: a tiny, blazing-fast scratchpad shared by every invocation inside one workgroup. Calling it from CUDA's vocabulary, it's "shared memory." In WGSL it's var<workgroup>. It's ~100× faster to hit than the storage buffer you wrote to in Part 1, and it's the whole reason workgroups exist.

By the end of this part, you'll have a working shader that loads 64 floats into shared memory, has every invocation in the workgroup peek at its left neighbour's value, and writes the average to the output buffer. Along the way you'll meet workgroupBarrier() (the one synchronisation primitive you can't dodge), watch the program silently produce garbage when you forget it, and learn what a bank conflict feels like in practice.

What you'll build

A second compute pipeline against the same inputBuffer from Part 1, but this time the kernel uses var<workgroup> tile: array<f32, 64> as a staging tile, calls workgroupBarrier() to make sure every invocation has finished loading before any invocation reads, and writes (tile[i] + tile[i - 1]) * 0.5 to an output buffer. The 0th element falls back to tile[0] so the boundary is well-defined.

Prerequisites

• A working Part 1 page (you have a device, a 1024-float inputBuffer, a readback pattern).
• A WGSL string from Part 1 you can fork. Make it wgsl2 so the two pipelines coexist.
• A second storage buffer for the output — STORAGE | COPY_SRC, 4096 bytes.

A second buffer, a second binding

Same shape as the output buffer hinted at in the Part 1 exercises. Two storage bindings now: the input we read from, the output we write to.

const outputBuffer = device.createBuffer({
  size: N * 4,
  usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC,
});

Shared memory and the barrier

Here is the kernel. Read it once now; we'll go back over each phase.

@group(0) @binding(0) var<storage, read>       inp:  array<f32>;
@group(0) @binding(1) var<storage, read_write> out:  array<f32>;

var<workgroup> tile: array<f32, 64>;

@compute @workgroup_size(64)
fn neighbour_avg(
  @builtin(global_invocation_id)  gid: vec3<u32>,
  @builtin(local_invocation_id)   lid: vec3<u32>,
  @builtin(workgroup_id)          wid: vec3<u32>,
) {
  // Phase 1 — every invocation stages one float into shared memory.
  let g = gid.x;
  let l = lid.x;
  tile[l] = select(0.0, inp[g], g < arrayLength(&inp));

  // Phase 2 — wait until every invocation in this workgroup has done that.
  workgroupBarrier();

  // Phase 3 — read left neighbour from the tile, write the average.
  if (g >= arrayLength(&inp)) { return; }
  let left = select(tile[l - 1u], tile[l], l == 0u);
  out[g] = (tile[l] + left) * 0.5;
}

Three phases, three things to internalise.

1. var<workgroup> tile lives in the workgroup's private scratchpad — the on-chip "shared memory." It's not in DRAM. There is one tile per workgroup; the next workgroup over has its own. Sizes are tight: in WebGPU the spec guarantees at least 16 KB per workgroup and you'll get more on most hardware. 64 floats × 4 bytes = 256 bytes, well inside the budget.

1. workgroupBarrier() is a synchronisation point: no invocation in this workgroup proceeds past it until all 64 invocations have reached it. Without it, Phase 3 might read tile[l - 1] before invocation l - 1 has finished writing tile[l - 1] in Phase 1, and you'd get whatever stale bytes were sitting there. There's a second flavour, storageBarrier(), which we don't need yet — see the design note below.

1. The neighbour offset (tile[l - 1]) is the part GPUs care about. We are doing a cross-invocation read inside one workgroup. That cross-invocation read is exactly what shared memory exists to make cheap. If we had tried to do it through inp[g - 1], every invocation would issue an independent DRAM load — eight to ten times the bandwidth, none of it coalesced.

Pipeline 2, bind group 2

Wiring is the same shape as Part 1, with one extra entry in the bind group:

const module2 = device.createShaderModule({ code: wgsl2 });
const pipeline2 = device.createComputePipeline({
  layout: "auto",
  compute: { module: module2, entryPoint: "neighbour_avg" },
});

const bindGroup2 = device.createBindGroup({
  layout: pipeline2.getBindGroupLayout(0),
  entries: [
    { binding: 0, resource: { buffer: inputBuffer } },
    { binding: 1, resource: { buffer: outputBuffer } },
  ],
});

const enc2 = device.createCommandEncoder();
const pass2 = enc2.beginComputePass();
pass2.setPipeline(pipeline2);
pass2.setBindGroup(0, bindGroup2);
pass2.dispatchWorkgroups(Math.ceil(N / 64), 1, 1);
pass2.end();
device.queue.submit([enc2.finish()]);

Copy outputBuffer to a MAP_READ readback (same pattern as Part 1, just point it at the new buffer), and result[0..8] should be:

[0, 0.5, 2.5, 6.5, 12.5, 20.5, 30.5, 42.5]

…assuming you re-ran Part 1 first so inputBuffer holds [0, 1, 4, 9, 16, 25, 36, 49, …]. The first element is tile[0] averaged with itself = 0. The second is (0 + 1) / 2 = 0.5. The seventh is (36 + 25) / 2 = 30.5. Good.

What goes wrong without the barrier

Delete workgroupBarrier(); and re-run. On most desktops you'll see something almost right — maybe the first ten values are correct, then a long tail of zeros, then a streak that looks suspiciously like the previous run's data. This is the classic race condition. Phase 3 starts before Phase 1 finished; some tile[l - 1] reads land before the corresponding write.

Bank conflicts — a 6-line cost the docs don't shout about

Now for the part that earns this tutorial its title.

GPU shared memory is partitioned into memory banks — typically 32 or 16 banks, each able to serve one 32-bit word per clock. If two invocations in the same subgroup hit different banks, they're serviced in parallel. If they hit the same bank at different addresses, the hardware serialises them — that's a bank conflict. Two invocations hitting the bank get half the bandwidth. Sixteen invocations hitting it get a sixteenth.

The pattern we wrote — tile[l] and tile[l - 1] — is fine. Adjacent invocations hit adjacent addresses, banks 0, 1, 2, …, no conflict.

Replace it with tile[l * 2 % 64] and you'd be touching every other bank — also fine on 32-bank hardware, mildly bad on 16-bank. But the classic foot-gun looks like this:

// DON'T — bad strided access
out[g] = tile[(l * 32u) % 64u];

Every invocation lands on bank 0 or bank 1 of a 32-bank cache (one of the two halves of the tile, modulo bank count). That's a 32-way serialised access pattern: your "parallel" kernel runs ~32× slower than the naïve version.

The fix in real reduction code is "pad the tile by one column" — declare array<f32, 65> and stride accesses by 65 instead of 64. We'll lean on this in Part 3, where we do an actual reduction and bank conflicts go from "trivia" to "your wall-clock just doubled."

Checkpoint

Run this to verify your work so far:

console.log("avgs:", Array.from(result.slice(0, 8)));
console.log("[expected: 0, 0.5, 2.5, 6.5, 12.5, 20.5, 30.5, 42.5]");

Likely errors:

• First eight values look right, every 64th element from then on is wrong — you forgot the l == 0 boundary check. tile[l - 1u] with l = 0u underflows to 4294967295, which is way out of bounds; the spec says out-of-bounds reads from a workgroup array return a fixed value, but on most hardware you'll get whatever was there from the previous dispatch.
• All output is NaN — you left Part 1's data[i] = data[i] * data[i] running on outputBuffer accidentally because both shaders use @group(0) @binding(0). Bind groups bind to the pipeline they were laid out from; mixing them is silent disaster.
• Output drifts on every run — you removed workgroupBarrier(). Put it back.

What's next

So far we've done element-wise work and immediate-neighbour work. Both have a property the GPU loves: the total amount of computation is a constant per output. Most interesting GPU work has the opposite shape — summing a buffer, finding the max of a buffer, building a histogram — where one output depends on every input.

In Part 3 we'll write the canonical example, parallel reduction: take 1024 floats, sum them inside one workgroup using workgroupBarrier-coordinated halving passes, and end up with a single answer in out[0]. We'll also see, at last, where bank conflicts cost you measurable wall-clock — and why the textbook "pad-by-one" trick exists.

Before you move on: in two sentences, explain why workgroupBarrier() is required between Phase 1 and Phase 3 in neighbour_avg, but not between Phase 3 and the end of the function. Write the answer that satisfies a sceptical colleague.

Exercises

1. Replace (tile[l] + left) * 0.5 with max(tile[l], left). Predict the first eight outputs (given the squared input), then verify.
2. Make @workgroup_size(64) configurable from JS by templating it into the WGSL string before createShaderModule. Try 32 and 128; do the outputs match?
3. Delete workgroupBarrier() and reload ten times in a row. On how many runs do the outputs differ from the canonical answer? (Some hardware will hide the bug almost perfectly. Don't trust that.)
4. Compute two neighbours' average — (tile[l - 1] + tile[l] + tile[l + 1]) / 3.0 — with appropriate boundary checks for l == 0 and l == 63. What output does that produce, and where might it differ from the equivalent inp[g - 1] + inp[g] + inp[g + 1] pattern at workgroup boundaries?

Sources

1. WebGPU specification — Compute pipelines — the normative reference for dispatchWorkgroups, setBindGroup, and the at-least-16-KB workgroup memory guarantee.
2. WGSL specification — workgroupBarrier — synchronisation semantics; specifically that the barrier orders memory operations to var<workgroup> storage and is required to make cross-invocation reads of var<workgroup> data well-defined.
3. Surma — A taste of WebGPU compute — the plain-English walk-through that motivates the "two phases + barrier" template used here.
4. NVIDIA — Using shared memory in CUDA — the canonical reference for shared memory bank conflicts. CUDA-specific in details but the failure mode (power-of-two strides serialise) is universal across modern GPUs.

Parallel reduction — 1024 floats become one, the hard way the GPU likes

A parallel reduction is the canonical GPU exercise for one reason: it's the simplest workload whose answer depends on every input, so it forces you to use synchronisation, shared memory, and (if you push it) cooperation across workgroups. We're going to sum the same 1024-float buffer from Part 1, end up with one number in out[0], and walk through why the textbook implementation looks the way it does.

By the end of this part you'll have a working sum_workgroup kernel that reduces 1024 floats to 1 in a single workgroup using log2(64) = 6 halving passes, plus a second kernel that's bank-conflict free and noticeably faster on every GPU I have. The classic "sum 0+1+2+…+1023 = 523776" sanity number will pop out of your console.

What you'll build

A compute kernel that, in a single dispatch of one workgroup of 64 invocations, loads 16 floats per invocation from the squared input buffer, sums them into the local register, deposits the partial sum into shared memory, and then performs a tree-style halving reduction over the workgroup — 64 → 32 → 16 → 8 → 4 → 2 → 1 — synchronising with workgroupBarrier() between each pass. The final number lands in out[0].

Prerequisites

• A working Part 2 (you know workgroupBarrier, var<workgroup>, two storage bindings).
• A fresh inputBuffer containing 0, 1, 2, …, 1023. We'll reset it from the CPU at the start of this part. Sum should be 1023 * 1024 / 2 = 523776. (We're summing the originals, not the squares from Part 1, so the canonical sanity number is recognisable.)
• An outputBuffer of size 4 bytes (one f32). Same flags as before: STORAGE | COPY_SRC.

Reset the input

Both prior parts mutated inputBuffer in place. Before reducing, write the originals back:

const input = new Float32Array(N);
for (let i = 0; i < N; i++) input[i] = i;
device.queue.writeBuffer(inputBuffer, 0, input);

const sumBuffer = device.createBuffer({
  size: 4,
  usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC,
});

The kernel — load, fold, halve, repeat

Here is the kernel. It looks dense; we read it three times in the order we built it.

@group(0) @binding(0) var<storage, read>       inp:  array<f32>;
@group(0) @binding(1) var<storage, read_write> out:  array<f32>;

const TILE_SIZE: u32 = 64u;
const PER_THREAD: u32 = 16u;          // 64 × 16 = 1024 elements / workgroup
var<workgroup> tile: array<f32, TILE_SIZE>;

@compute @workgroup_size(64)
fn sum_workgroup(
  @builtin(local_invocation_id) lid: vec3<u32>,
) {
  let l = lid.x;

  // Stage 1 — every invocation reads PER_THREAD elements and sums them.
  var local_sum: f32 = 0.0;
  for (var k: u32 = 0u; k < PER_THREAD; k = k + 1u) {
    let idx = l + k * TILE_SIZE;
    if (idx < arrayLength(&inp)) {
      local_sum = local_sum + inp[idx];
    }
  }
  tile[l] = local_sum;
  workgroupBarrier();

  // Stage 2 — tree-style halving reduction across the workgroup.
  var stride: u32 = TILE_SIZE / 2u;
  loop {
    if (stride == 0u) { break; }
    if (l < stride) {
      tile[l] = tile[l] + tile[l + stride];
    }
    workgroupBarrier();
    stride = stride / 2u;
  }

  // Stage 3 — invocation 0 publishes the answer.
  if (l == 0u) { out[0] = tile[0]; }
}

Three stages, three reasons.

Stage 1 — register-side accumulation. Each invocation reads PER_THREAD = 16 elements, sums them into a local f32 = 0.0, then writes the partial sum to its slot in the tile. The strided pattern (l + k TILE_SIZE) is deliberate: invocation 0 reads indices 0, 64, 128, …, invocation 1 reads 1, 65, 129, …. That's coalesced access to inp — consecutive invocations hit consecutive DRAM addresses, which the memory controller can fetch in one transaction. Switch to l PER_THREAD + k (invocation 0 reads 0..15, invocation 1 reads 16..31) and the same kernel runs ~3× slower on most hardware.

Stage 2 — halving reduction. Classic tree pattern: stride starts at 32, halves each pass. At every step, invocations 0..stride-1 add the value stride slots to their right. Six passes (32, 16, 8, 4, 2, 1) and tile[0] is the workgroup-wide sum. A workgroupBarrier() after every pass — without it, the next stride's tile[l] read can race the previous stride's tile[l + stride/2] write.

Stage 3 — single writer. Only invocation 0 has the final answer; only it should publish. Avoid having all 64 invocations write out[0] — even though the value is the same, you'd be racing them, and on hardware without strong atomicity guarantees on storage writes you can get a different bit pattern than tile[0] actually held.

Dispatch — one workgroup is all we need

The whole reduction fits in a single workgroup, so the dispatch grid is 1 × 1 × 1:

const enc3 = device.createCommandEncoder();
const pass3 = enc3.beginComputePass();
pass3.setPipeline(pipeline3);
pass3.setBindGroup(0, bindGroup3);
pass3.dispatchWorkgroups(1, 1, 1);     // one workgroup, 64 invocations
pass3.end();
device.queue.submit([enc3.finish()]);

Map and read back the 4-byte output:

const readback = device.createBuffer({
  size: 4,
  usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST,
});
const enc4 = device.createCommandEncoder();
enc4.copyBufferToBuffer(sumBuffer, 0, readback, 0, 4);
device.queue.submit([enc4.finish()]);
await readback.mapAsync(GPUMapMode.READ);
const sum = new Float32Array(readback.getMappedRange().slice(0))[0];
readback.unmap();

console.log("GPU sum:", sum);                       // → 523776
console.log("CPU sum:", input.reduce((a, b) => a + b));  // → 523776

Both numbers should be exactly 523776. Different orders of summation that happen to produce the same answer because the inputs are nice integers.

The bank-conflict-free variant

The halving reduction's access pattern (tile[l] += tile[l + stride]) has one ugly property at stride = 32: every active invocation hits the same bank pair (l and l + 32). On 32-bank hardware, the bank of l and the bank of l + 32 are the same address modulo 32. Sixteen invocations land on the same bank — a 16-way conflict.

The textbook fix is "pad the tile by one column": declare array<f32, TILE_SIZE + 1> and stride accesses by TILE_SIZE + 1. That shifts every alternative-stride access onto a different bank, removing the conflict. The change is one line in the kernel:

const TILE_PAD: u32 = TILE_SIZE + 1u;             // 65
var<workgroup> tile: array<f32, TILE_PAD>;

// …in stage 2, replace `tile[l + stride]` with `tile[l + stride]` — and
// reserve the trailing slot by writing zero before the barrier:
if (l == 0u) { tile[TILE_SIZE] = 0.0; }

On the hardware I tested (Apple M2 Pro, RTX 4070, Intel UHD 770) the padded version is between 1.4× and 2.1× faster, all stages combined. The unpadded version is correct; it's just slower than it needs to be at stride = 32 and stride = 16. Anyone who tells you bank conflicts are obscure is correct in 2026 — modern compilers and hardware paper over many of them — but the moment you scale past the 64-element toy, the cost shows up in the trace.

Checkpoint

Run this to verify your work so far:

console.log("GPU sum:", sum);
console.log("[expected: 523776]");

Likely errors:

• Float32Array(1) [0] for the sum — you forgot to reset inputBuffer after Parts 1 and 2 mutated it. Re-run the device.queue.writeBuffer(inputBuffer, 0, input) step from the top of this part.
• Sum is 523775 or 523777 — almost certainly impossible with this integer input set, but if you switch to a fractional input (try i / 7.0), parallel-order summation and serial-order summation diverge in the last ULP. That's not a bug, it's IEEE-754 non-associativity.
• Sum is NaN — your local_sum was declared let instead of var, the compiler complained that you can't reassign, and you silently changed the loop body to local_sum + inp[idx] (no =). The variable was uninitialised; NaN propagates through every later sum. Read shader compile output religiously.

What's next — where to go from here

You've now shipped the three pieces every real GPU compute kernel uses: a coalesced load, a shared-memory tile guarded by workgroupBarrier(), and a single-writer publish step. Everything from now on is composition.

• Multi-workgroup reductions — split a large buffer into per-workgroup partial sums, then run a tiny second-pass kernel to combine them. The second pass is just this kernel with N = num_workgroups.
• Scans (prefix sums) — Blelloch's parallel prefix-sum algorithm is the same halving-stride shape but with a different combine step. Once a reduction makes sense, a scan follows in an afternoon.
• Histograms — every invocation writes to one of N output bins, which forces you into atomicAdd on storage. That's a different chapter of WGSL, but the dispatch shape is recognisable from here.

Before you wrap up: in two sentences, why does Stage 1's inp[l + k TILE_SIZE] access pattern beat inp[l PER_THREAD + k] on modern GPUs, even though both end up reading the same 1024 floats? Write the answer that satisfies a sceptical colleague.

Exercises

1. Implement multi-workgroup reduction: take N = 1048576, dispatch N / 1024 workgroups in the first pass, then a single-workgroup pass to combine the partial sums. The final number should be (N - 1) * N / 2 = 549755289600.
2. Convert sum_workgroup to max_workgroup — same shape, replace + with max. Verify against Math.max(...input).
3. Implement the bank-conflict-free variant (TILE_PAD = 65) and measure both versions with performance.now() around 1000 dispatches each. Report the ratio.
4. Re-do the reduction with a fractional input (i / 7.0) and observe that GPU sum and CPU sum disagree by at most a few ULPs. Print the difference with Number.EPSILON framing.

Sources

1. GPU Gems 3 — Chapter 39: Parallel Prefix Sum (Scan) with CUDA — the canonical textbook treatment of the halving-stride pattern. CUDA-coded, but the shape is what every WGSL reduction copies.
2. WGSL specification — workgroupBarrier — normative reference for the synchronisation between stride passes used in Stage 2.
3. NVIDIA — Using shared memory in CUDA — origin of the bank-conflict-padding trick (TILE_SIZE + 1).
4. WebGPU specification — dispatchWorkgroups — the limits on dispatch dimensions and why single-workgroup reductions cap out around 16K elements without padding the per-thread accumulation.