[SYCL][NATIVECPU] Support atomic fence on Native CPU

libdevice/nativecpu_utils.cpp

@@ -52,6 +52,12 @@ __spirv_ControlBarrier(uint32_t Execution, uint32_t Memory,
     __mux_work_group_barrier(0, Execution, Semantics);
 }
 
+DEVICE_EXTERN_C void __mux_mem_barrier(uint32_t scope, uint32_t semantics);
+__SYCL_CONVERGENT__ DEVICE_EXTERNAL void
+__spirv_MemoryBarrier(uint32_t Memory, uint32_t Semantics) {
+  __mux_mem_barrier(Memory, Semantics);
+}
+
 // Turning clang format off here because it reorders macro invocations
 // making the following code very difficult to read.
 // clang-format off

sycl/cmake/modules/FetchUnifiedRuntime.cmake

@@ -117,13 +117,13 @@ if(SYCL_UR_USE_FETCH_CONTENT)
   endfunction()
 
   set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git")
-  # commit 4d19115165b5497b647ae1b2e110488f84d1806a
-  # Merge: fb6df497 3f128d09
-  # Author: Piotr Balcer <piotr.balcer@intel.com>
-  # Date:   Tue Sep 17 10:31:44 2024 +0200
-  #     Merge pull request #2087 from nrspruit/fix_driver_inorder_event_wait
-  #     [L0] Fix urEnqueueEventsWaitWithBarrier for driver in order lists
-  set(UNIFIED_RUNTIME_TAG 4d19115165b5497b647ae1b2e110488f84d1806a)
+  # commit cfecab08e6e6dbb694f614b4f6271a258a41fc10
+  # Merge: 10fd78c1 5bebef5d
+  # Author: Omar Ahmed <omar.ahmed@codeplay.com>
+  # Date:   Tue Sep 17 12:26:35 2024 +0100
+  #   Merge pull request #1874 from PietroGhg/pietro/membarrier
+  #   [NATIVECPU] Support atomic fence queries
+  set(UNIFIED_RUNTIME_TAG cfecab08e6e6dbb694f614b4f6271a258a41fc10)
 
   set(UMF_BUILD_EXAMPLES OFF CACHE INTERNAL "EXAMPLES")
   # Due to the use of dependentloadflag and no installer for UMF and hwloc we need

sycl/test/check_device_code/native_cpu/native_cpu_attrs.cpp

@@ -0,0 +1,27 @@
+// REQUIRES: native_cpu_ock
+// RUN: %clangxx -fsycl -fsycl-targets=native_cpu -Xclang -sycl-std=2020 -mllvm -sycl-native-dump-device-ir %s | FileCheck %s
+
+#include "sycl.hpp"
+using namespace sycl;
+
+class Test;
+int main() {
+  sycl::queue deviceQueue;
+  sycl::nd_range<2> r({1, 1}, {1,1});
+  deviceQueue.submit([&](handler &h) {
+    h.parallel_for<Test>(
+        r, [=](nd_item<2> it) { 
+          it.barrier(access::fence_space::local_space);
+          //CHECK-DAG: call void @__mux_work_group_barrier({{.*}})
+          atomic_fence(memory_order::acquire, memory_scope::work_group);
+          //CHECK-DAG: call void @__mux_mem_barrier({{.*}})
+        });
+  });
+
+}
+
+//CHECK-DAG: define{{.*}}@__mux_work_group_barrier{{.*}}#[[ATTR:[0-9]+]]
+//CHECK-DAG: [[ATTR]]{{.*}}convergent
+
+//CHECK-DAG: define{{.*}}@__mux_mem_barrier{{.*}}#[[ATTR_MEM:[0-9]+]]
+//CHECK-DAG: [[ATTR_MEM]]{{.*}}convergent