KhronosGroup / KhronosGroup/Vulkan-Portability

Support for SPIR-V OpMemoryBarrier

Open
#6 0 comments 0 reactions 0 assignees View on GitHub
enhancement Metal
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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.