Chapter 07

Atomics & Memory Synchronization

Atomics are for the places where many invocations touch the same memory location. Barriers are for the places where a workgroup needs a clear before/after boundary. Both are small tools, but they sit exactly where GPU code stops being "many independent lanes" and starts being shared-state programming.

Atomic types

atomic<T> is only useful when the same value can be accessed by multiple invocations. In WGSL the component type is i32 or u32, and atomics belong in shared address spaces such as workgroup or storage.

var<workgroup> local_count : atomic<u32>;

@group(0) @binding(0)
var<storage, read_write> global_count : atomic<u32>;

Use workgroup atomics when lanes inside one workgroup cooperate. Use storage atomics when many workgroups need to accumulate into a shared buffer visible to the host or to another pass.

Atomic operations

Every atomic operation gives you a single, indivisible update for one memory location. The important habit is to remember which operations return the old value and which operations simply publish a value.

A histogram is the classic first example

Many invocations read independent input values, but they may land in the same bin. Without atomics, two lanes can read the same old bin value and lose one increment. atomicAdd turns the increment into one indivisible update.

@group(0) @binding(0)
var<storage, read> input : array<u32>;

@group(0) @binding(1)
var<storage, read_write> bins : array<atomic<u32>>;

@compute @workgroup_size(256)
fn histogram(@builtin(global_invocation_id) gid : vec3u) {
  let i = gid.x;
  if (i >= arrayLength(&input)) {
    return;
  }

  let bin = input[i] & 255u;
  atomicAdd(&bins[bin], 1u);
}

Barriers draw a workgroup boundary

A barrier is a promise made by every invocation in a workgroup: everyone reaches this point before anyone continues past it. That is why barrier calls must be reached uniformly. They are not a general "make all GPU work done" button; they synchronize lanes in the scope described by the operation.

Workgroup scratch memory

The most common barrier pattern is: load into workgroup memory, wait, reduce or reuse that memory, wait again if another phase needs the result. Each invocation owns one lane of the algorithm, but the array is shared by the workgroup.

var<workgroup> scratch : array<u32, 256>;

@group(0) @binding(0)
var<storage, read> input : array<u32>;

@group(0) @binding(1)
var<storage, read_write> partials : array<u32>;

@compute @workgroup_size(256)
fn sum_block(
  @builtin(local_invocation_id) lid : vec3u,
  @builtin(global_invocation_id) gid : vec3u,
  @builtin(workgroup_id) wid : vec3u
) {
  let local_index = lid.x;
  scratch[local_index] = select(0u, input[gid.x], gid.x < arrayLength(&input));
  workgroupBarrier();

  var stride = 128u;
  loop {
    if (local_index < stride) {
      scratch[local_index] = scratch[local_index] + scratch[local_index + stride];
    }
    workgroupBarrier();

    if (stride == 1u) {
      break;
    }
    stride = stride / 2u;
  }

  if (local_index == 0u) {
    partials[wid.x] = scratch[0];
  }
}

Memory layout still matters

Atomics and barriers solve synchronization problems. They do not remove layout rules. When a buffer contains structs, the host and shader still need the same alignment and stride story.

Study loop

1. Name the shared location. Is it workgroup memory, storage memory, or both?
2. Find every write. Decide whether two invocations can hit the same address.
3. If the address can collide, choose the smallest atomic operation that matches the update.
4. If later code reads shared workgroup memory, place a barrier where every lane reaches it.
5. Use Studio diagnostics for legality, reflection for bindings, and readback for the final data shape.