Previews, as seen when this build job started (84bdbc7):
WebGPU | IDL
WGSL
Explainer

I realize I made a mistake when reading the Metal docs and unintentionally assumed more functionality than is guaranteed. threadgroup_barrier in MSL is effectively OpenCL C 1.2's barrier intrinsic which only has a memory scope of Workgroup. This means (in the context of #2229) that inter-workgroup communication cannot be support correctly on all underlying platforms.

Really? From the MSL manual v2.4, 6.9.1 "Threadgroup and SIMD-group synchronization functions"

The threadgroup_barrier (or simdgroup_barrier) function can also queue a memory fence
(reads and writes) to ensure the correct ordering of memory operations to threadgroup or
device memory.

And table 6.12 (mem fence flags): has

mem_device: Ensure correct ordering of memory operations to device
memory.

The OpenCL docs provide a bit better explanation, but there are a few things going on with barriers (and atomics).

For barrier and the work_group_barrier variant that does not take a memory scope, the scope is memory_scope_work_group.

That's from the unified OpenCL C 3.0 spec. From OpenCL's SPIR environment:

In an OpenCL 1.2 environment, for the Barrier Instructions OpControlBarrier and OpMemoryBarrier, the Scope for Memory must be Workgroup, and the memory-order constraint in Memory Semantics must be SequentiallyConsistent.

For barriers, I think SPIR breaks it down reasonably well, though the semantics operand is a bit muddled:

• Execution scope: which invocations participate in the control aspect
• Memory scope: the extent (in terms of invocations) of cache flushing
• Memory semantics: both memory ordering and the affected memory types

Metal's mem_fence_flags only affects the storage aspect of memory semantics, not the memory scope. That's the intent of this PR.

The important part is to update the mapping as that really describes the memory model aspect. We can do a follow up PR or two to try and clarify what these four (five?) things really are for barriers and atomics (no execution scope there of course).

Metal's mem_fence_flags only affects the storage aspect of memory semantics, not the memory scope. That's the intent of this PR.

Got it.

If I'm reading it right, it means Metal's threadgroup_barrier doesn't reliably support message passing between one threadgroup (workgroup) and another via device memory. Hence the first complaint in #2229.

A similar conclusion applies to vanilla barrier() in GLSL on Vulkan or OpenGL. You'd need to use a more capable barrier enabled by GL_KHR_memory_scope_semantics

I think your analysis of HLSL's barriers concluded that it does support what's needed. (And we made sure the SPIR-V backend on DXC translates appropriately when targeting Vulkan at least.)

Previews, as seen when this build job started (ee7b2ae):
WebGPU | IDL
WGSL
Explainer

To address the questions raised by @dneto0 about what's needed for message passing across workgroups (where the message can't be packed in a 32 bit word along with the flags), no, from what I can tell it can't be done in Metal at all (confirmed by testing and discussions with experts). On GLSL in compatibility mode, memoryBarrierBuffer should be sufficient, as it's defined in the spec to have queue family scope (which is fine, all I care about is within a dispatch; coordination between GPU and CPU and suchlike is a whole nother ball of wax). And on HLSL I am fairly certain DeviceMemoryBarrier has device scope as well.

On GLSL in compatibility mode, memoryBarrierBuffer should be sufficient, as it's defined in the spec to have queue family scope (which is fine, all I care about is within a dispatch; coordination between GPU and CPU and suchlike is a whole nother ball of wax).

'memoryBarrierBuffer' was added in the vulkan memory model extension for GLSL (that's what you pointed at).

Prior to that extension, there is only barrier() and memoryBarrier().
The barrier() is not sufficient:

https://github.com/KhronosGroup/GLSL/blob/master/extensions/khr/GL_KHR_memory_scope_semantics.txt#L461

In compute shaders, barrier() is equivalent to controlBarrier() with
execution and memory scope equal to gl_ScopeWorkgroup, storage semantics
equal to gl_StorageSemanticsShared, and sem equal to
gl_SemanticsAcquireRelease. Informally, this means it synchronizes
accesses to shared memory between invocations in the same workgroup.

For memoryBarrier() I may stand corrected:

https://github.com/KhronosGroup/GLSL/blob/master/extensions/khr/GL_KHR_memory_scope_semantics.txt#L492

    void memoryBarrier()
    // equivalent to:
    // memoryBarrier(gl_ScopeQueueFamily,
    //               gl_StorageSemanticsBuffer |
    //               gl_StorageSemanticsShared |
    //               gl_StorageSemanticsImage,
    //               gl_SemanticsAcquireRelease)

That's a lot stronger than I expected. But then there's a caveat that we're relying on GLSL-before-Vulkan-memory-model chaining a barrier() with a memoryBarrier() will effectively form (essentially) a synchronizing chain. I'm not clear that we can guarantee that, but it's what we're shooting for in the WGSL memory model, and will test. We do guarantee that in the vulkan memory model case.

Ah, thanks. I thought memoryBarrierBuffer was in the compatibility set, I might switch it to memoryBarrier if I come across a device that has trouble with it.

I'm finding a lot of failures on actual GPUs as I exercise this stuff, as in I haven't yet come across a single card that executes all the tests in my suite without any issues. I think this is going to be a two-part effort, first spec'ing the memory model, then doing a considerable amount of validation work.

I'm finding a lot of failures on actual GPUs as I exercise this stuff, as in I haven't yet come across a single card that executes all the tests in my suite without any issues.

Thank you!

I think this is going to be a two-part effort, first spec'ing the memory model, then doing a considerable amount of validation work.

Definitely. We have to find the common set of functionality that actually works. The draft as landed, and amended here, is our best understanding at this point.

But Metal and D3D memory models are unspecified, and Vulkan is really only specified once we opt into "vulkan memory model". So we're relying on "lore", commonly-understood-but-never-really-well-specified. One way to look at it is that GPU-land is where CPU-land was prior to C++11.

As for testing, @alan-baker engaged with UCSC researchers. They have been writing memory model litmus tests for WebGPU. They've done a great job. See https://users.soe.ucsc.edu/~reeselevine/webgpu/ (requires Chrome canary with a flag). They plan to present an overview in an upcoming WGSL meeting (December 14).

There are lots of moving parts here:

• what's the actual behaviour of underlying platforms
• spec'ing that well
• testing that well
• ensuring browsers map correctly to underlying platforms
• ensuring platform compilers and drivers do what they promise (or don't promise a the case may be)

Thanks for pushing the envelope. It helps us all.

Previews, as seen when this build job started (ee7b2ae):
WebGPU | IDL
WGSL
Explainer

Minor update from the UCSC side, we now have an official subdomain for the memory model tests at https://gpuharbor.ucsc.edu/ (and it's in the origin trial so no need for Chrome Canary!).

I've been talking to @raphlinus about these issues as well, there are some ideas that I'm hoping to work into the tests there. Also, there's a version of his prefix sum code at https://gpuharbor.ucsc.edu/prefix-sum/, which we can hopefully use to reproduce some of the failures and look for more testing ideas!

I agree this represents the current state of understanding, and thanks. I would argue that the name storageBarrier is confusing for the new understanding (I'd much prefer the name be reserved for a barrier with device scope), and that the mismatch between scope and space will be difficult to document, but don't at this point have a better proposal.

A bit of a follow-up. I have indeed begun collaborating with @reeselevine on more rigorously testing the memory model, especially in light of the failures I've been seeing (which btw are not quite as bad in my last comment, some of the failures turned out to be a write-after-read hazard in my code). I should also clarify, the decoupled look-back test is now not expected to pass in WebGPU, but the program may still be useful in tracking bugs (it absolutely would show failures in current AMD 5700 XT Windows Vulkan if the barrier were upgraded to device scope), and of course in pointing the way to what might be added post-1.0 to make such workloads viable.

So far we haven't come up with a test that is valid WebGPU and unambiguously catches new failures, but we have one in the pipeline which may pan out, relating to coherence of relaxed atomic loads. We'll keep iterating on that and certainly file an issue if anything needs fixing in the spec.