diff options
| -rw-r--r-- | source/slang/hlsl.meta.slang | 3 | ||||
| -rw-r--r-- | tests/spirv/barrier.slang | 14 |
2 files changed, 16 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: 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 |