Skip to content

Commit da480db

Browse files
authored
[SYCL][NATIVECPU] Support atomic fence on Native CPU (#14619)
Adds `__spirv_MemoryBarrier` to the Native CPU device library. UR PR: oneapi-src/unified-runtime#1874
1 parent e89da32 commit da480db

File tree

3 files changed

+40
-7
lines changed

3 files changed

+40
-7
lines changed

libdevice/nativecpu_utils.cpp

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -52,6 +52,12 @@ __spirv_ControlBarrier(uint32_t Execution, uint32_t Memory,
5252
__mux_work_group_barrier(0, Execution, Semantics);
5353
}
5454

55+
DEVICE_EXTERN_C void __mux_mem_barrier(uint32_t scope, uint32_t semantics);
56+
__SYCL_CONVERGENT__ DEVICE_EXTERNAL void
57+
__spirv_MemoryBarrier(uint32_t Memory, uint32_t Semantics) {
58+
__mux_mem_barrier(Memory, Semantics);
59+
}
60+
5561
// Turning clang format off here because it reorders macro invocations
5662
// making the following code very difficult to read.
5763
// clang-format off

sycl/cmake/modules/FetchUnifiedRuntime.cmake

Lines changed: 7 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -117,13 +117,13 @@ if(SYCL_UR_USE_FETCH_CONTENT)
117117
endfunction()
118118

119119
set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git")
120-
# commit 4d19115165b5497b647ae1b2e110488f84d1806a
121-
# Merge: fb6df497 3f128d09
122-
# Author: Piotr Balcer <piotr.balcer@intel.com>
123-
# Date: Tue Sep 17 10:31:44 2024 +0200
124-
# Merge pull request #2087 from nrspruit/fix_driver_inorder_event_wait
125-
# [L0] Fix urEnqueueEventsWaitWithBarrier for driver in order lists
126-
set(UNIFIED_RUNTIME_TAG 4d19115165b5497b647ae1b2e110488f84d1806a)
120+
# commit cfecab08e6e6dbb694f614b4f6271a258a41fc10
121+
# Merge: 10fd78c1 5bebef5d
122+
# Author: Omar Ahmed <omar.ahmed@codeplay.com>
123+
# Date: Tue Sep 17 12:26:35 2024 +0100
124+
# Merge pull request #1874 from PietroGhg/pietro/membarrier
125+
# [NATIVECPU] Support atomic fence queries
126+
set(UNIFIED_RUNTIME_TAG cfecab08e6e6dbb694f614b4f6271a258a41fc10)
127127

128128
set(UMF_BUILD_EXAMPLES OFF CACHE INTERNAL "EXAMPLES")
129129
# Due to the use of dependentloadflag and no installer for UMF and hwloc we need
Lines changed: 27 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,27 @@
1+
// REQUIRES: native_cpu_ock
2+
// RUN: %clangxx -fsycl -fsycl-targets=native_cpu -Xclang -sycl-std=2020 -mllvm -sycl-native-dump-device-ir %s | FileCheck %s
3+
4+
#include "sycl.hpp"
5+
using namespace sycl;
6+
7+
class Test;
8+
int main() {
9+
sycl::queue deviceQueue;
10+
sycl::nd_range<2> r({1, 1}, {1,1});
11+
deviceQueue.submit([&](handler &h) {
12+
h.parallel_for<Test>(
13+
r, [=](nd_item<2> it) {
14+
it.barrier(access::fence_space::local_space);
15+
//CHECK-DAG: call void @__mux_work_group_barrier({{.*}})
16+
atomic_fence(memory_order::acquire, memory_scope::work_group);
17+
//CHECK-DAG: call void @__mux_mem_barrier({{.*}})
18+
});
19+
});
20+
21+
}
22+
23+
//CHECK-DAG: define{{.*}}@__mux_work_group_barrier{{.*}}#[[ATTR:[0-9]+]]
24+
//CHECK-DAG: [[ATTR]]{{.*}}convergent
25+
26+
//CHECK-DAG: define{{.*}}@__mux_mem_barrier{{.*}}#[[ATTR_MEM:[0-9]+]]
27+
//CHECK-DAG: [[ATTR_MEM]]{{.*}}convergent

0 commit comments

Comments
 (0)