https://github.com/arsenm updated https://github.com/llvm/llvm-project/pull/184844
>From e2d14bfd021313fc9e34a9e08efc968aacb71ff7 Mon Sep 17 00:00:00 2001 From: Matt Arsenault <[email protected]> Date: Thu, 5 Mar 2026 17:09:48 +0100 Subject: [PATCH 1/2] libclc: Add atomic_work_item_fence --- .../clc/lib/amdgcn/mem_fence/clc_mem_fence.cl | 2 ++ libclc/opencl/lib/generic/SOURCES | 1 + .../generic/atomic/atomic_work_item_fence.cl | 17 +++++++++++++++++ 3 files changed, 20 insertions(+) create mode 100644 libclc/opencl/lib/generic/atomic/atomic_work_item_fence.cl diff --git a/libclc/clc/lib/amdgcn/mem_fence/clc_mem_fence.cl b/libclc/clc/lib/amdgcn/mem_fence/clc_mem_fence.cl index 6d2a0962ba20d..f54c9e214752c 100644 --- a/libclc/clc/lib/amdgcn/mem_fence/clc_mem_fence.cl +++ b/libclc/clc/lib/amdgcn/mem_fence/clc_mem_fence.cl @@ -10,6 +10,8 @@ #define BUILTIN_FENCE_ORDER(memory_order, ...) \ switch (memory_order) { \ + case __ATOMIC_RELAXED: \ + break; \ case __ATOMIC_ACQUIRE: \ __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, __VA_ARGS__); \ break; \ diff --git a/libclc/opencl/lib/generic/SOURCES b/libclc/opencl/lib/generic/SOURCES index 893312f405fd0..a2c1122ca7634 100644 --- a/libclc/opencl/lib/generic/SOURCES +++ b/libclc/opencl/lib/generic/SOURCES @@ -41,6 +41,7 @@ atomic/atom_or.cl atomic/atom_sub.cl atomic/atom_xchg.cl atomic/atom_xor.cl +atomic/atomic_work_item_fence.cl common/degrees.cl common/mix.cl common/radians.cl diff --git a/libclc/opencl/lib/generic/atomic/atomic_work_item_fence.cl b/libclc/opencl/lib/generic/atomic/atomic_work_item_fence.cl new file mode 100644 index 0000000000000..19eb40b5447a2 --- /dev/null +++ b/libclc/opencl/lib/generic/atomic/atomic_work_item_fence.cl @@ -0,0 +1,17 @@ +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include "clc/mem_fence/clc_mem_fence.h" +#include "clc/opencl/synchronization/utils.h" + +_CLC_OVERLOAD _CLC_DEF void atomic_work_item_fence(cl_mem_fence_flags flags, + memory_order order, + memory_scope scope) { + __clc_mem_fence(__opencl_get_memory_scope(scope), order, + __opencl_get_memory_semantics(flags)); +} >From 0b5d5efe2ec6eb5cd5bd4ace29cb70b72ad2f41f Mon Sep 17 00:00:00 2001 From: Matt Arsenault <[email protected]> Date: Fri, 6 Mar 2026 00:05:50 +0100 Subject: [PATCH 2/2] __opencl_get_clang_memory_scope --- libclc/opencl/lib/generic/atomic/atomic_work_item_fence.cl | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/libclc/opencl/lib/generic/atomic/atomic_work_item_fence.cl b/libclc/opencl/lib/generic/atomic/atomic_work_item_fence.cl index 19eb40b5447a2..164b0ec2ff31c 100644 --- a/libclc/opencl/lib/generic/atomic/atomic_work_item_fence.cl +++ b/libclc/opencl/lib/generic/atomic/atomic_work_item_fence.cl @@ -8,10 +8,11 @@ #include "clc/mem_fence/clc_mem_fence.h" #include "clc/opencl/synchronization/utils.h" +#include "clc/opencl/utils.h" _CLC_OVERLOAD _CLC_DEF void atomic_work_item_fence(cl_mem_fence_flags flags, memory_order order, memory_scope scope) { - __clc_mem_fence(__opencl_get_memory_scope(scope), order, + __clc_mem_fence(__opencl_get_clang_memory_scope(scope), order, __opencl_get_memory_semantics(flags)); } _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
