memoryBarrier
Memory fence that orders shared- and storage-memory writes within a workgroup.
Test memoryBarrier in a live shader
Open the playground, start from a visual preset, and wire memoryBarrier into the fragment stage to see how it behaves with real values.
Open PlaygroundLive Demo Pending
memoryBarrier orders memory accesses: writes the invocation made before the call to shared memory and to storage buffers 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
memoryBarrier :: () -> void {...}
memoryBarrier 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);
memoryBarrier(); // 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
- storageBarrier - Currently the same fence
- Compute Shaders - Shared memory and workgroups