[libclc] f31e65f - libclc: Add atomic_work_item_fence (#184844)

[via cfe-commits]

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));
+}