Barries are unsafe, but not clear why

[nazar-pc]

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.

[Firestar99]

I think we just made everything in spirv_std::arch unsafe since those are all intrinsics. But I could see us making many of them safe, #306 already started on subgroup. (We also have quite a few unsound things that need fixing, like shared memory being &mut isn't great when it rather should be a &[Cell<T>])

[nazar-pc]

Yes, shared memory (and global for that matter) are a little funny because they appear to provide &mut from multiple threads, which is at odds with Rust aliasing rules. It is like working with custom DSL: you see one thing, but imagine a different thing.

Since this is an explicitly multi-threaded situation, it can't be Cell, but rather UnsafeCell (or SyncUnsafeCell, though I guess no one will spawn more threads from existing thread, so it makes no difference), which unfortunately immediately causes developer experience to crater due to inherent complexity of dealing with UnsafeCell.

It may not be UB (or at very least not the same UB as in CPU code) to modify those from multiple threads, so might benefit from a custom container that is less painful to work with than UnsafeCell.

[schell]

A section about the safety would be a pretty easy addition, and we could mention that these are unsafe because they're intrinsics.

[Firestar99]

I'd rather not, if all intrinsics are unsafe people will stop reading the safety section as "it's just an intrinsic". I'd rather selectively mark functions as unsafe when they are actually unsafe.

[nazar-pc]

because they're intrinsics

This is not a correct justification. There are many intrinsics that are completely safe to use. In Rust standard library all standard atomic operations are safe, even most x86-64 SIMD intrinsics are actually safe in the latest stable (you have to use correct target_feature for it of course): https://doc.rust-lang.org/stable/core/arch/x86_64/fn._mm256_and_si256.html

So I agree with @Firestar99 that every intrinsic needs to be evaluated separately.

Let's take workgroup_memory_barrier_with_group_sync() intrinsic for example. When is it unsafe to call? "just an intrinsic" is a bad explanation because it doesn't really help the developer to use it correctly (and unsafe implies it can only be used correctly in certain ways). To me this particular intrinsic appears to be perfectly safe and there are probably many others like it.