[SYCL2020] Implement sycl::atomic_fence

[nmnobre]

Hi Aksel,

Less than a year later, I'm trying to bring it full circle here. 🥳
This (properly) closes #767.

Let me know what you think so far.
I haven't implemented the SSCP path because I could only find existing implementations for barriers but not memory fences, and so I think I should wait for you to tell me where and how you want this. 🙂

Cheers,
-N

[illuhad]

Good idea to work on this! There's already https://github.com/OpenSYCL/OpenSYCL/blob/develop/include/hipSYCL/sycl/libkernel/detail/mem_fence.hpp
which does basically the same thing, except that it comes from the SYCL 1.2.1 days. Can we implement the old mem_fence on top of the newer one?

For SSCP I would suggest looking at barrier - every barrier is effectively a combination of a memfence and a blocking synchronization primitive. So probably we can factor the mem_fence out from there :)

[illuhad]

What is the benefit of these structs, when we do runtime if-else in atomic_fence() anyway?

[nmnobre]

Little perhaps.
It was mostly to keep it similar to mem_fence.hpp and perhaps bring more clarity to which of the order and scope combinations we actually support.

[nmnobre]

Completely rewritten, structs gone.

[nmnobre]

Can we implement the old mem_fence on top of the newer one?

Done.

For SSCP I would suggest looking at barrier - every barrier is effectively a combination of a memfence and a blocking synchronization primitive. So probably we can factor the mem_fence out from there :)

Done.

Can I get the declaration of __spirv_MemoryBarrier from spirv_ops.hpp in fence.cpp or better not? I did not, just in case we want a clear separation.

As written, the semantics for seq_cst + device for explicit multipass SPIR-V and SSCP SPIR-V is different as the memory semantics flags differ. I suspect the former includes MemorySemanticsMask::WorkgroupMemory because mem_fence could (and still can) use fence_space::global_and_local, but we were using it even with fence_space::global_space. In the barrier implementation we do not include MemorySemanticsMask::WorkgroupMemory. What should we do here?

I've completely rewritten my implementation so it reflects my understanding of the intended semantics.
There are now no differences for SPIR-V.
For HIP-like devices, the SSCP path might be more precise, in the sense that the IR fence primitives we end up calling expose more choices than just __threadfence_block() and __threadfence().
It may well be those IR primitives only implement sequential consistency at the workgroup and device scopes as well - I dunno - but we don't care about that... Indeed, if that were to change, SSCP compilations would automatically be taking advantage.