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 ++-
 tests/spirv/barrier.slang    | 14 ++++++++++++++
 2 files changed, 16 insertions(+), 1 deletion(-)
 create mode 100644 tests/spirv/barrier.slang

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:
diff --git a/tests/spirv/barrier.slang b/tests/spirv/barrier.slang
new file mode 100644
index 000000000..769bf9b7e
--- /dev/null
+++ b/tests/spirv/barrier.slang
@@ -0,0 +1,14 @@
+//TEST:SIMPLE(filecheck=CHECK):-target spirv -emit-spirv-directly -O0
+//TEST:SIMPLE(filecheck=CHECK):-target spirv -emit-spirv-via-glsl -entry main -stage compute -O0
+
+RWStructuredBuffer<float> output;
+
+[numthreads(1,1,1)]
+[shader("compute")]
+void main()
+{
+    output[0] = 0;
+    // CHECK: OpControlBarrier %uint_2 %uint_2 %uint_264
+    GroupMemoryBarrierWithGroupSync();
+    output[1] = 0;
+}
\ No newline at end of file
-- 
cgit v1.2.3