[SYCL][NATIVECPU] added __spir cast builtins to NativeCPU

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                                     \

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 <sycl/sycl.hpp>
+
+using namespace sycl;
+using namespace sycl::ext::oneapi::experimental;
+
+#define DefTestCast(FName, Space, PType)                                       \
+  SYCL_EXTERNAL auto FName(PType p) { return dynamic_address_cast<Space>(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}