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