atomic ops unsafe, but unclear why

[bcolloran]

this is basically a follow up to issue #382, where @nazar-pc said:

"""
Functions like workgroup_memory_barrier_with_group_sync() are unsafe, but it is not clear why (no "Safety" section in docs).

They should either not be unsafe or "Safety" section should be added to corresponding function docs explaining how to use them safely.
"""

Similar concern with the files in e.g. crates/spirv-std/src/arch/atomics.rs -- all of the atomic ops seem to be marked as unsafe. I don't understand the spir-v / Vulkan / whatevs memory model... like... at all... but I always thought atomic ops like these were specifically designed to provide a safe way to update memory that might be touched by multiple threads? If this is not the case, then as @nazar-pc suggested re barriers, a comment saying how to uphold any invariants required for safety would be a big help! Thanks!! :-)

[charles-r-earp]

The workgroup barrier must be in uniform control flow. That is, all the threads in the group must reach that instruction together, else it's UB.
This isn't checked by the compiler, so it's unsafe.
The atomics exposed in spirv_std are more akin to rust std atomic intrinsics which are all unsafe for similar reasons. These are then made safe by the Atomic* types.
It would be trivial to invalidate usage of atomic accesses by simultaneously accessing the same values via non atomic loads / stores. This is prevented by declaring a variable with an Atomic type, which then enforces that it can only be accessed atomically.
In both cases these are unfortunately low level interfaces that don't have high level counterparts.

[nazar-pc]

I think it is only unsafe in Rust because it is intrinsic, not necessarily because it is unsafe to call. The fact that you get a bad value when doing non-atomic loads/stores doesn't mean the code has a safety issue.

In fact a whole bunch of SIMD instructions in Rust are safe these days, here is a random one i have just opened: https://doc.rust-lang.org/stable/core/arch/x86_64/fn._mm256_mask_broadcast_i32x2.html

So yes, it is important to distinguish memory safety issues from garbage output and document things accordingly.