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 <[email protected]> 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(); _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
