[libclc] libclc: Add atomic_work_item_fence (PR #184844)

[Matt Arsenault via cfe-commits]

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 <Matthew.Arsenault at amd.com>
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 <Matthew.Arsenault at amd.com>
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));
 }