KhronosGroup / KhronosGroup/Vulkan-Portability
Support for SPIR-V OpMemoryBarrier
- Dominant language
- No language data
- Stars
- 45
- Forks
- 7
- PR merge metrics
- No merged PRs in 30d
Description
MSL supports `OpControlBarrier` via `threadgroup_barrier()`, but MSL does not support the equivalent of `OpMemoryBarrier`, as distinct from this.
This becomes meaningful because `OpMemoryBarrier` does not have the same requirement for uniform flow (consistent `if` and `for` statement paths across threads) that govern `OpControlBarrier` and `threadgroup_barrier()`, and there are [use-cases](https://github.com/KhronosGroup/MoltenVK/issues/973) that depend on this.
@kvark commented:
> I double checked the specs and can confirm that Metal only supports uniform control flow barriers. A new flag for this might be needed, something like `nonUniformMemoryBarriers`.
Contributor guide
No contributing guide indexed for this repository
Research direction
Start by reviewing the existing OpControlBarrier handling through MSL’s threadgroup_barrier() and the distinction from OpMemoryBarrier. Consider the linked MoltenVK issue and whether a nonUniformMemoryBarriers flag is needed. Done means the portability layer supports the described non-uniform memory-barrier use cases without treating them as uniform control-flow barriers.
Written by the indexing model from the issue text.
Assessment
- Domain
- compilers, computer-graphics
- Issue type
- Feature
- Difficulty
- 5/5
- Estimated time
- Over a week
- Activity status
- Stale
- Clarity
- Mostly clear
- Newbie friendliness
- 25/100