diff --git a/libdevice/nativecpu_utils.cpp b/libdevice/nativecpu_utils.cpp index 6eaf5bfaf9c27..4107663e6d1fd 100755 --- a/libdevice/nativecpu_utils.cpp +++ b/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 diff --git a/sycl/cmake/modules/FetchUnifiedRuntime.cmake b/sycl/cmake/modules/FetchUnifiedRuntime.cmake index 811b9ff0a4445..f29b09723ff66 100644 --- a/sycl/cmake/modules/FetchUnifiedRuntime.cmake +++ b/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 - # 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 + # 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 diff --git a/sycl/test/check_device_code/native_cpu/native_cpu_attrs.cpp b/sycl/test/check_device_code/native_cpu/native_cpu_attrs.cpp new file mode 100644 index 0000000000000..63cfa084eee77 --- /dev/null +++ b/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( + 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