[llvm] [SPIRV] Preserve Workgroup execution scope for OpenCL barriers (PR #228961)

Sun Oct 4 19:12:18 PDT 2026

llvmorg-github-actions[bot] wrote:




@llvm/pr-subscribers-backend-spir-v

Author: Alexander Shaposhnikov (alexander-shaposhnikov)

<details>
<summary>Changes</summary>

Consider a work-group of 64 work-items split into two subgroups of 32:

  work_group_barrier(CLK_GLOBAL_MEM_FENCE, memory_scope_sub_group);

All 64 work-items must reach the barrier before any can continue. The memory-scope argument only controls memory ordering.

The lowering starts with Scope = Workgroup and maps the argument to MemScope = Subgroup. However, the old condition includes !(MemFlags & CLK_LOCAL_MEM_FENCE). This is true for a global-only fence, so Scope = MemScope overwrites the execution scope with Subgroup. The resulting OpControlBarrier lets one subgroup continue before the other has arrived.

Restrict that assignment to Opcode == OpMemoryBarrier. work_group_barrier emits OpControlBarrier, so its execution scope now stays Workgroup while its memory scope remains Subgroup. OpMemoryBarrier has no execution-scope operand: its first operand is the memory scope, so it still needs the assignment.

P.S.
The local-memory fence is still handled earlier:
  if (MemFlags & SPIRV::CLK_LOCAL_MEM_FENCE)
    MemSemantics |= SPIRV::MemorySemantics::WorkgroupMemory;


---
Full diff: https://github.com/llvm/llvm-project/pull/228961.diff


3 Files Affected:

- (modified) llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp (+1-2) 
- (modified) llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_split_barrier/split_work_group_barrier_20.ll (+6) 
- (modified) llvm/test/CodeGen/SPIRV/transcoding/OpenCL/work_group_barrier.ll (+14) 


``````````diff

diff --git a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
index 84196697060e3..c73bb4fab5930 100644
--- a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp
@@ -980,8 +980,7 @@ static bool buildBarrierInst(const SPIRV::IncomingCall *Call, unsigned Opcode,
     SPIRV::CLMemoryScope CLScope =
         static_cast<SPIRV::CLMemoryScope>(getIConstVal(ScopeArg, MRI));
     MemScope = getSPIRVScope(CLScope);
-    if (!(MemFlags & SPIRV::CLK_LOCAL_MEM_FENCE) ||
-        (Opcode == SPIRV::OpMemoryBarrier))
+    if (Opcode == SPIRV::OpMemoryBarrier)
       Scope = MemScope;
     if (CLScope == static_cast<unsigned>(Scope))
       ScopeReg = Call->Arguments[1];
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_split_barrier/split_work_group_barrier_20.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_split_barrier/split_work_group_barrier_20.ll
index 2d94e1f35094b..fe0a59a077a1c 100644
--- a/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_split_barrier/split_work_group_barrier_20.ll
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_split_barrier/split_work_group_barrier_20.ll
@@ -24,6 +24,8 @@
 ;;     intel_work_group_barrier_wait(CLK_LOCAL_MEM_FENCE, memory_scope_all_svm_devices);
 ;;     intel_work_group_barrier_arrive(CLK_LOCAL_MEM_FENCE, memory_scope_sub_group);
 ;;     intel_work_group_barrier_wait(CLK_LOCAL_MEM_FENCE, memory_scope_sub_group);
+;;     intel_work_group_barrier_arrive(CLK_GLOBAL_MEM_FENCE, memory_scope_device);
+;;     intel_work_group_barrier_wait(CLK_GLOBAL_MEM_FENCE, memory_scope_device);
 ;;}
 
 ; RUN: not llc -O0 -mtriple=spirv64-unknown-unknown %s -o %t.spvt 2>&1 | FileCheck %s --check-prefix=CHECK-ERROR
@@ -88,6 +90,8 @@
 ; CHECK: OpControlBarrierWaitINTEL %[[#SCOPE_WORK_GROUP]] %[[#SCOPE_CROSS_DEVICE]] %[[#ACQUIRE_LOCAL]]
 ; CHECK: OpControlBarrierArriveINTEL %[[#SCOPE_WORK_GROUP]] %[[#SCOPE_SUBGROUP]] %[[#RELEASE_LOCAL]]
 ; CHECK: OpControlBarrierWaitINTEL %[[#SCOPE_WORK_GROUP]] %[[#SCOPE_SUBGROUP]] %[[#ACQUIRE_LOCAL]]
+; CHECK: OpControlBarrierArriveINTEL %[[#SCOPE_WORK_GROUP]] %[[#SCOPE_DEVICE]] %[[#RELEASE_GLOBAL]]
+; CHECK: OpControlBarrierWaitINTEL %[[#SCOPE_WORK_GROUP]] %[[#SCOPE_DEVICE]] %[[#ACQUIRE_GLOBAL]]
 
 ; Function Attrs: convergent norecurse nounwind
 define dso_local spir_kernel void @test(ptr addrspace(1) nocapture noundef readnone align 4 %0) local_unnamed_addr #0 !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !6 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 {
@@ -111,6 +115,8 @@ define dso_local spir_kernel void @test(ptr addrspace(1) nocapture noundef readn
   tail call spir_func void @_Z29intel_work_group_barrier_waitj12memory_scope(i32 noundef 1, i32 noundef 3) #2
   tail call spir_func void @_Z31intel_work_group_barrier_arrivej12memory_scope(i32 noundef 1, i32 noundef 4) #2
   tail call spir_func void @_Z29intel_work_group_barrier_waitj12memory_scope(i32 noundef 1, i32 noundef 4) #2
+  tail call spir_func void @_Z31intel_work_group_barrier_arrivej12memory_scope(i32 noundef 2, i32 noundef 2) #2
+  tail call spir_func void @_Z29intel_work_group_barrier_waitj12memory_scope(i32 noundef 2, i32 noundef 2) #2
   ret void
 }
 
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/OpenCL/work_group_barrier.ll b/llvm/test/CodeGen/SPIRV/transcoding/OpenCL/work_group_barrier.ll
index a3c79d554a588..dd0ede4b106bd 100644
--- a/llvm/test/CodeGen/SPIRV/transcoding/OpenCL/work_group_barrier.ll
+++ b/llvm/test/CodeGen/SPIRV/transcoding/OpenCL/work_group_barrier.ll
@@ -1,4 +1,5 @@
 ; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
 ; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %}
 
 ;; This test checks that the backend is capable to correctly translate
@@ -22,6 +23,10 @@
 ;;
   ;; barrier should also work (preserved for backward compatibility)
 ;;   barrier(CLK_GLOBAL_MEM_FENCE);
+;;   work_group_barrier(CLK_GLOBAL_MEM_FENCE, memory_scope_work_group);
+;;   work_group_barrier(CLK_GLOBAL_MEM_FENCE, memory_scope_device);
+;;   work_group_barrier(CLK_GLOBAL_MEM_FENCE, memory_scope_all_svm_devices);
+;;   work_group_barrier(CLK_GLOBAL_MEM_FENCE, memory_scope_sub_group);
 ;; }
 ;;
 ;; __kernel void test_barrier_non_const_flags(cl_mem_fence_flags flags, memory_scope scope) {
@@ -77,6 +82,11 @@
 ; CHECK-SPIRV: OpControlBarrier %[[#SCOPE_WORK_GROUP]] %[[#SCOPE_CROSS_DEVICE]] %[[#LOCAL]]
 ; CHECK-SPIRV: OpControlBarrier %[[#SCOPE_WORK_GROUP]] %[[#SCOPE_SUBGROUP]] %[[#LOCAL]]
 ; CHECK-SPIRV: OpControlBarrier %[[#SCOPE_WORK_GROUP]] %[[#SCOPE_WORK_GROUP]] %[[#GLOBAL]]
+; Explicit memory scope must not change the Workgroup execution scope.
+; CHECK-SPIRV: OpControlBarrier %[[#SCOPE_WORK_GROUP]] %[[#SCOPE_WORK_GROUP]] %[[#GLOBAL]]
+; CHECK-SPIRV: OpControlBarrier %[[#SCOPE_WORK_GROUP]] %[[#SCOPE_DEVICE]] %[[#GLOBAL]]
+; CHECK-SPIRV: OpControlBarrier %[[#SCOPE_WORK_GROUP]] %[[#SCOPE_CROSS_DEVICE]] %[[#GLOBAL]]
+; CHECK-SPIRV: OpControlBarrier %[[#SCOPE_WORK_GROUP]] %[[#SCOPE_SUBGROUP]] %[[#GLOBAL]]
 
 define dso_local spir_kernel void @test_barrier_const_flags() local_unnamed_addr {
 entry:
@@ -92,6 +102,10 @@ entry:
   tail call spir_func void @_Z18work_group_barrierj12memory_scope(i32 noundef 1, i32 noundef 3)
   tail call spir_func void @_Z18work_group_barrierj12memory_scope(i32 noundef 1, i32 noundef 4)
   tail call spir_func void @_Z7barrierj(i32 noundef 2)
+  tail call spir_func void @_Z18work_group_barrierj12memory_scope(i32 noundef 2, i32 noundef 1)
+  tail call spir_func void @_Z18work_group_barrierj12memory_scope(i32 noundef 2, i32 noundef 2)
+  tail call spir_func void @_Z18work_group_barrierj12memory_scope(i32 noundef 2, i32 noundef 3)
+  tail call spir_func void @_Z18work_group_barrierj12memory_scope(i32 noundef 2, i32 noundef 4)
   ret void
 }
 

``````````

</details>


https://github.com/llvm/llvm-project/pull/228961
