Intrinsics1 min read

storageBarrier

Memory fence for storage-buffer writes within a workgroup.

Reading Time
1 min
Word Count
179
Sections
4
Try It Live

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 Playground

Live Demo Pending

This intrinsic does not yet have a self-contained interactive preview. Compute-only and external-resource intrinsics need a different demo path than the current fragment and 3D showcases.

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

bwsl
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:

bwsl
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

TargetOutput
SPIR-VOpMemoryBarrier (workgroup scope; uniform and workgroup memory)
GLSLgroupMemoryBarrier()
HLSLAllMemoryBarrier()
Metalthreadgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup)

See Also