diff --git a/libdevice/nativecpu_utils.cpp b/libdevice/nativecpu_utils.cpp index c3e8bb61657a7..eb5c3ff2ebdf4 100644 --- a/libdevice/nativecpu_utils.cpp +++ b/libdevice/nativecpu_utils.cpp @@ -31,16 +31,7 @@ using __nativecpu_state = native_cpu::state; #define OCL_LOCAL __attribute__((opencl_local)) #define OCL_GLOBAL __attribute__((opencl_global)) - -DEVICE_EXTERNAL OCL_LOCAL void * -__spirv_GenericCastToPtrExplicit_ToLocal(void *p, int) { - return (OCL_LOCAL void *)p; -} - -DEVICE_EXTERNAL OCL_GLOBAL void * -__spirv_GenericCastToPtrExplicit_ToGlobal(void *p, int) { - return (OCL_GLOBAL void *)p; -} +#define OCL_PRIVATE __attribute__((opencl_private)) DEVICE_EXTERN_C void __mux_work_group_barrier(uint32_t id, uint32_t scope, uint32_t semantics); @@ -61,6 +52,23 @@ __spirv_MemoryBarrier(uint32_t Memory, uint32_t Semantics) { // Turning clang format off here because it reorders macro invocations // making the following code very difficult to read. // clang-format off + +#define DefGenericCastToPtrExplImpl(sfx, asp, cv)\ +DEVICE_EXTERNAL cv asp void *\ +__spirv_GenericCastToPtrExplicit_##sfx(cv void *p ,int) {\ + return (cv asp void *)p;\ +} + +#define DefGenericCastToPtrExpl(sfx, asp)\ + DefGenericCastToPtrExplImpl(sfx, asp, )\ + DefGenericCastToPtrExplImpl(sfx, asp, const)\ + DefGenericCastToPtrExplImpl(sfx, asp, volatile)\ + DefGenericCastToPtrExplImpl(sfx, asp, const volatile) + +DefGenericCastToPtrExpl(ToPrivate, OCL_PRIVATE) +DefGenericCastToPtrExpl(ToLocal, OCL_LOCAL) +DefGenericCastToPtrExpl(ToGlobal, OCL_GLOBAL) + #define DefSubgroupBlockINTEL1(Type, PType) \ template <> \ __SYCL_CONVERGENT__ DEVICE_EXTERNAL Type \ diff --git a/sycl/test/check_device_code/native_cpu/device_builtins.cpp b/sycl/test/check_device_code/native_cpu/device_builtins.cpp new file mode 100644 index 0000000000000..cd2a392ea4b5e --- /dev/null +++ b/sycl/test/check_device_code/native_cpu/device_builtins.cpp @@ -0,0 +1,36 @@ +// REQUIRES: native_cpu_ock +// 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 + +// 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 + +// check that builtins are defined + +// CHECK-NOT: {{.*}}__spirv_GenericCastToPtrExplicit +// CHECK-DEV: {{.*}}__spirv_GenericCastToPtrExplicit + +#include + +using namespace sycl; +using namespace sycl::ext::oneapi::experimental; + +#define DefTestCast(FName, Space, PType) \ + SYCL_EXTERNAL auto FName(PType p) { return dynamic_address_cast(p); } + +// Turning clang format off here because it would change the indentations of +// the macro invocations making the following code difficult to read. +// clang-format off + +#define DefTestCastForSpace(PType)\ + DefTestCast(to_local, access::address_space::local_space, PType)\ + DefTestCast(to_global, access::address_space::global_space, PType)\ + DefTestCast(to_private, access::address_space::private_space, PType)\ + DefTestCast(to_generic, access::address_space::generic_space, PType) + +DefTestCastForSpace(int*) +DefTestCastForSpace(const int*) +DefTestCastForSpace(volatile int*) +DefTestCastForSpace(const volatile int*) + +int main(){} +// check that the generated module has the is-native-cpu module flag set +// CHECK: !{{[0-9]*}} = !{i32 1, !"is-native-cpu", i32 1}