From 74d4c8ab92c1050865c421a64d9f5d3fc9a2da73 Mon Sep 17 00:00:00 2001
From: Yong He <yonghe@outlook.com>
Date: Wed, 3 Apr 2024 11:46:25 -0700
Subject: Update glsl intrinsic for GroupMemoryBarrierWithGroupSync (#3890)

* Update glsl intrinsic for `GroupMemoryBarrierWithGroupSync`,

* Add spirv tests for `GroupMemoryBarrierWithGroupSync`.
---
 source/slang/hlsl.meta.slang | 3 ++-
 1 file changed, 2 insertions(+), 1 deletion(-)

(limited to 'source')

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:
-- 
cgit v1.2.3