storageBarrier
Memory fence for storage-buffer writes within a workgroup.
Test storageBarrier in a live shader
Open the playground, start from a visual preset, and wire storageBarrier into the fragment stage to see how it behaves with real values.
Open PlaygroundLive Demo Pending
storageBarrier orders writes to storage buffers (and shared memory): writes the invocation made before the call become visible to other invocations in the workgroup before any write it makes after the call. Unlike barrier, it does not wait for other invocations.
Signature
storageBarrier :: () -> void {...}
storageBarrier is only valid in compute stages. Calling it in a vertex or fragment stage is a compile error.
Example
A fence on its own does not make another invocation wait. To publish data and then read what other invocations wrote, follow the fence with barrier:
compute "Publish" [64, 1, 1] {
uint i = input.global_id.x;
resources.results[i] = computeValue(i);
storageBarrier(); // Order this invocation's buffer writes
barrier(); // Wait for the rest of the workgroup
float neighbor = resources.results[i ^ 1u];
}
memoryBarrier and storageBarrier
The compiler currently emits the same fence for memoryBarrier and storageBarrier: a workgroup-scope memory barrier covering both workgroup (shared) memory and buffer memory. Use the name that states your intent; it keeps code correct if the two are ever split.
Compiled Output
| Target | Output |
|---|---|
| SPIR-V | OpMemoryBarrier (workgroup scope; uniform and workgroup memory) |
| GLSL | groupMemoryBarrier() |
| HLSL | AllMemoryBarrier() |
| Metal | threadgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup) |
See Also
- barrier - Execution barrier for the whole workgroup
- memoryBarrier - Currently the same fence
- Compute Shaders - Shared memory and workgroups