diff options
| author | Yong He <yonghe@outlook.com> | 2024-04-03 11:46:25 -0700 |
|---|---|---|
| committer | GitHub <noreply@github.com> | 2024-04-03 11:46:25 -0700 |
| commit | 74d4c8ab92c1050865c421a64d9f5d3fc9a2da73 (patch) | |
| tree | c0173f5e59ca2c7f418ecd1568a26a95aab26ff6 /source | |
| parent | e0de98e9aabbe118f0eeca7821518c8fb4e1f6c4 (diff) | |
Update glsl intrinsic for GroupMemoryBarrierWithGroupSync (#3890)
* Update glsl intrinsic for `GroupMemoryBarrierWithGroupSync`,
* Add spirv tests for `GroupMemoryBarrierWithGroupSync`.
Diffstat (limited to 'source')
| -rw-r--r-- | source/slang/hlsl.meta.slang | 3 |
1 files changed, 2 insertions, 1 deletions
diff --git a/source/slang/hlsl.meta.slang b/source/slang/hlsl.meta.slang index 972702170..cd6de61b5 100644 --- a/source/slang/hlsl.meta.slang +++ b/source/slang/hlsl.meta.slang @@ -5105,11 +5105,12 @@ void __subgroupBarrier() } } +__glsl_extension(GL_KHR_memory_scope_semantics) void GroupMemoryBarrierWithGroupSync() { __target_switch { - case glsl: __intrinsic_asm "barrier"; + case glsl: __intrinsic_asm "controlBarrier(gl_ScopeWorkgroup, gl_ScopeWorkgroup, gl_StorageSemanticsShared, gl_SemanticsAcquireRelease)"; case hlsl: __intrinsic_asm "GroupMemoryBarrierWithGroupSync"; case cuda: __intrinsic_asm "__syncthreads()"; case spirv: |