[libclc] 9d38926 - libclc: Fix logic for fence ordering based on mem flags (#185367)

[via cfe-commits]

Author: Matt Arsenault
Date: 2026-03-09T09:15:21+01:00
New Revision: 9d38926d51121c9820d22768bfb498cefc2857e0

URL: https://github.com/llvm/llvm-project/commit/9d38926d51121c9820d22768bfb498cefc2857e0
DIFF: https://github.com/llvm/llvm-project/commit/9d38926d51121c9820d22768bfb498cefc2857e0.diff

LOG: libclc: Fix logic for fence ordering based on mem flags (#185367)

Added: 
    

Modified: 
    libclc/clc/lib/amdgcn/synchronization/clc_work_group_barrier.cl

Removed: 
    


################################################################################
diff  --git a/libclc/clc/lib/amdgcn/synchronization/clc_work_group_barrier.cl b/libclc/clc/lib/amdgcn/synchronization/clc_work_group_barrier.cl
index 67b3d9b2f308b..0950af0aaf0d6 100644
--- a/libclc/clc/lib/amdgcn/synchronization/clc_work_group_barrier.cl
+++ b/libclc/clc/lib/amdgcn/synchronization/clc_work_group_barrier.cl
@@ -15,14 +15,13 @@ __clc_work_group_barrier(int memory_scope,
   if (memory_semantics == 0) {
     __builtin_amdgcn_s_barrier();
   } else {
-    int memory_order_before =
-        memory_semantics & (__CLC_MEMORY_GLOBAL | __CLC_MEMORY_LOCAL)
-            ? __ATOMIC_SEQ_CST
-            : __ATOMIC_RELEASE;
-    int memory_order_after =
-        memory_semantics & (__CLC_MEMORY_GLOBAL | __CLC_MEMORY_LOCAL)
-            ? __ATOMIC_SEQ_CST
-            : __ATOMIC_ACQUIRE;
+    uint seq_cst_mask = __CLC_MEMORY_GLOBAL | __CLC_MEMORY_LOCAL;
+    int memory_order_before = (memory_semantics & seq_cst_mask) == seq_cst_mask
+                                  ? __ATOMIC_SEQ_CST
+                                  : __ATOMIC_RELEASE;
+    int memory_order_after = (memory_semantics & seq_cst_mask) == seq_cst_mask
+                                 ? __ATOMIC_SEQ_CST
+                                 : __ATOMIC_ACQUIRE;
 
     __clc_mem_fence(memory_scope, memory_order_before, memory_semantics);
     __builtin_amdgcn_s_barrier();