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

[Matt Arsenault via cfe-commits]

https://github.com/arsenm created https://github.com/llvm/llvm-project/pull/185367

None

>From fef2aec12e1c4c5def72774f3cae50b955504c12 Mon Sep 17 00:00:00 2001
From: Matt Arsenault <Matthew.Arsenault at amd.com>
Date: Sun, 8 Mar 2026 18:22:33 +0100
Subject: [PATCH] libclc: Fix logic for fence ordering based on mem flags

---
 .../synchronization/clc_work_group_barrier.cl     | 15 +++++++--------
 1 file changed, 7 insertions(+), 8 deletions(-)

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();