Skip to content

Commit 004d6d9

Browse files
authored
[SYCL][NATIVECPU] added __spir cast builtins to NativeCPU (#16676)
Added more support for __spir cast builtins. Needed for e2e etc
1 parent 8998b9b commit 004d6d9

File tree

2 files changed

+54
-10
lines changed

2 files changed

+54
-10
lines changed

libdevice/nativecpu_utils.cpp

Lines changed: 18 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -31,16 +31,7 @@ using __nativecpu_state = native_cpu::state;
3131

3232
#define OCL_LOCAL __attribute__((opencl_local))
3333
#define OCL_GLOBAL __attribute__((opencl_global))
34-
35-
DEVICE_EXTERNAL OCL_LOCAL void *
36-
__spirv_GenericCastToPtrExplicit_ToLocal(void *p, int) {
37-
return (OCL_LOCAL void *)p;
38-
}
39-
40-
DEVICE_EXTERNAL OCL_GLOBAL void *
41-
__spirv_GenericCastToPtrExplicit_ToGlobal(void *p, int) {
42-
return (OCL_GLOBAL void *)p;
43-
}
34+
#define OCL_PRIVATE __attribute__((opencl_private))
4435

4536
DEVICE_EXTERN_C void __mux_work_group_barrier(uint32_t id, uint32_t scope,
4637
uint32_t semantics);
@@ -61,6 +52,23 @@ __spirv_MemoryBarrier(uint32_t Memory, uint32_t Semantics) {
6152
// Turning clang format off here because it reorders macro invocations
6253
// making the following code very difficult to read.
6354
// clang-format off
55+
56+
#define DefGenericCastToPtrExplImpl(sfx, asp, cv)\
57+
DEVICE_EXTERNAL cv asp void *\
58+
__spirv_GenericCastToPtrExplicit_##sfx(cv void *p ,int) {\
59+
return (cv asp void *)p;\
60+
}
61+
62+
#define DefGenericCastToPtrExpl(sfx, asp)\
63+
DefGenericCastToPtrExplImpl(sfx, asp, )\
64+
DefGenericCastToPtrExplImpl(sfx, asp, const)\
65+
DefGenericCastToPtrExplImpl(sfx, asp, volatile)\
66+
DefGenericCastToPtrExplImpl(sfx, asp, const volatile)
67+
68+
DefGenericCastToPtrExpl(ToPrivate, OCL_PRIVATE)
69+
DefGenericCastToPtrExpl(ToLocal, OCL_LOCAL)
70+
DefGenericCastToPtrExpl(ToGlobal, OCL_GLOBAL)
71+
6472
#define DefSubgroupBlockINTEL1(Type, PType) \
6573
template <> \
6674
__SYCL_CONVERGENT__ DEVICE_EXTERNAL Type \
Lines changed: 36 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,36 @@
1+
// REQUIRES: native_cpu_ock
2+
// RUN: %clangxx -fsycl -fsycl-targets=native_cpu -Xclang -sycl-std=2020 -mllvm -sycl-opt -mllvm -inline-threshold=500 -O0 -mllvm -sycl-native-cpu-no-vecz -mllvm -sycl-native-dump-device-ir %s | FileCheck %s
3+
4+
// RUN: %clangxx -fsycl-device-only -fsycl-targets=native_cpu -fno-inline -Xclang -sycl-std=2020 -mllvm -sycl-opt -S -emit-llvm %s -o - | FileCheck %s --check-prefix=CHECK-DEV
5+
6+
// check that builtins are defined
7+
8+
// CHECK-NOT: {{.*}}__spirv_GenericCastToPtrExplicit
9+
// CHECK-DEV: {{.*}}__spirv_GenericCastToPtrExplicit
10+
11+
#include <sycl/sycl.hpp>
12+
13+
using namespace sycl;
14+
using namespace sycl::ext::oneapi::experimental;
15+
16+
#define DefTestCast(FName, Space, PType) \
17+
SYCL_EXTERNAL auto FName(PType p) { return dynamic_address_cast<Space>(p); }
18+
19+
// Turning clang format off here because it would change the indentations of
20+
// the macro invocations making the following code difficult to read.
21+
// clang-format off
22+
23+
#define DefTestCastForSpace(PType)\
24+
DefTestCast(to_local, access::address_space::local_space, PType)\
25+
DefTestCast(to_global, access::address_space::global_space, PType)\
26+
DefTestCast(to_private, access::address_space::private_space, PType)\
27+
DefTestCast(to_generic, access::address_space::generic_space, PType)
28+
29+
DefTestCastForSpace(int*)
30+
DefTestCastForSpace(const int*)
31+
DefTestCastForSpace(volatile int*)
32+
DefTestCastForSpace(const volatile int*)
33+
34+
int main(){}
35+
// check that the generated module has the is-native-cpu module flag set
36+
// CHECK: !{{[0-9]*}} = !{i32 1, !"is-native-cpu", i32 1}

0 commit comments

Comments
 (0)