Author: Matt Arsenault Date: 2026-03-06T07:57:04+01:00 New Revision: f31e65feaa50878f36c04ec46d0efc0085488fe5
URL: https://github.com/llvm/llvm-project/commit/f31e65feaa50878f36c04ec46d0efc0085488fe5 DIFF: https://github.com/llvm/llvm-project/commit/f31e65feaa50878f36c04ec46d0efc0085488fe5.diff LOG: libclc: Add atomic_work_item_fence (#184844) Added: libclc/opencl/lib/generic/atomic/atomic_work_item_fence.cl Modified: libclc/clc/lib/amdgcn/mem_fence/clc_mem_fence.cl libclc/opencl/lib/generic/SOURCES Removed: ################################################################################ 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 be94a34e9af08..43f70cf37a377 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..164b0ec2ff31c --- /dev/null +++ b/libclc/opencl/lib/generic/atomic/atomic_work_item_fence.cl @@ -0,0 +1,18 @@ +//===----------------------------------------------------------------------===// +// +// 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" +#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_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
