barrier
Workgroup execution barrier that also makes shared-memory writes visible.
Test barrier in a live shader
Open the playground, start from a visual preset, and wire barrier into the fragment stage to see how it behaves with real values.
Open PlaygroundLive Demo Pending
barrier blocks each invocation until every invocation in the workgroup has reached it. It also orders shared memory: writes made before the barrier are visible to every invocation in the workgroup after it.
Signature
barrier :: () -> void {...}
barrier is only valid in compute stages. Calling it in a vertex or fragment stage is a compile error.
Example
compute "Blur" [64, 1, 1] {
shared float[64] tile;
uint i = input.local_index;
// Every invocation loads one value...
tile[i] = resources.input[input.global_id.x];
// ...and waits until the whole tile is loaded before reading neighbors.
barrier();
float left = tile[(i + 63u) % 64u];
float right = tile[(i + 1u) % 64u];
resources.output[input.global_id.x] = (left + tile[i] + right) / 3.0;
}
Uniform control flow
Every invocation in the workgroup must reach the same barrier() calls in the same order. Placing a barrier inside a branch or loop that only some invocations take is undefined behavior and can hang the GPU.
Compiled Output
| Target | Output |
|---|---|
| SPIR-V | OpControlBarrier (workgroup scope, workgroup memory) |
| GLSL | barrier() |
| HLSL | GroupMemoryBarrierWithGroupSync() |
| Metal | threadgroup_barrier(mem_flags::mem_threadgroup) |
See Also
- memoryBarrier - Memory ordering without waiting for other invocations
- storageBarrier - Memory ordering for storage buffers
- Compute Shaders - Shared memory and workgroups