[WGSL] Support Synchronization Built-in Functions - https://www.w3.org/TR/WGSL/#sync-builtin-functions First two are easy: [WGSL] => [MSL] mapping * storageBarrier => threadgroup_barrier(mem_flags::mem_device) * workgroupBarrier => threadgroup_barrier(mem_flags::mem_threadgroup) Last one expands to this I think * workgroupUniformLoad => uint metalWorkgroupUniformLoad(threadgroup T* const p) { threadgroup_barrier(mem_flags::mem_threadgroup); const T result = *p; threadgroup_barrier(mem_flags::mem_threadgroup); return result; }
<rdar://problem/115855666>