Chapter 06

Parallelism & Correctness

A compute dispatch is many invocations running the same code with different coordinates. Correctness comes from knowing which memory is private, which memory is shared, and where a before/after boundary is required.

Invocation coordinates define responsibility

global_invocation_id usually tells each lane which element to read or write. Guard the tail of a dispatch when the data length is not a clean multiple of the workgroup size.

override COUNT : u32 = 1024u;

@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) id : vec3u) {
  if (id.x >= COUNT) {
    return;
  }
  data[id.x] = data[id.x] * 2.0;
}

Barriers only order one workgroup

workgroupBarrier() lets lanes in the same workgroup agree that earlier shared-memory writes are visible before later reads. It does not synchronize separate workgroups.

Uniformity is a validation concern

Some operations require all relevant lanes to reach the same point in compatible control flow. When Studio flags uniformity, read it as a control-flow proof problem, not as a syntax typo.