[SYCL] refactor and implement dynamic address space casts (intel#16604)

The patch provides an implementation for
dynamic address space cast for NVPTX and AMDGPU.

The patch also refactor the static cast by relying on a simple cast
rather than the SPIR-V builtin.

---------

Signed-off-by: Victor Lomuller <[email protected]>

Author: Victor Lomuller

libclc/libspirv/include/libspirv/conversion/GenericCastToPtrExplicit.h

@@ -0,0 +1,27 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#define GenericCastToPtrExplicit_To(ADDRSPACE, NAME)                           \
+  _CLC_DECL _CLC_OVERLOAD                                                      \
+      ADDRSPACE void *__spirv_GenericCastToPtrExplicit_To##NAME(               \
+          generic void *, int);                                                \
+  _CLC_DECL _CLC_OVERLOAD                                                      \
+      ADDRSPACE const void *__spirv_GenericCastToPtrExplicit_To##NAME(         \
+          generic const void *, int);                                          \
+  _CLC_DECL _CLC_OVERLOAD                                                      \
+      ADDRSPACE volatile void *__spirv_GenericCastToPtrExplicit_To##NAME(      \
+          generic volatile void *, int);                                       \
+  _CLC_DECL _CLC_OVERLOAD ADDRSPACE const volatile void *                      \
+      __spirv_GenericCastToPtrExplicit_To##NAME(generic const volatile void *, \
+                                                int)
+
+GenericCastToPtrExplicit_To(global, Global);
+GenericCastToPtrExplicit_To(local, Local);
+GenericCastToPtrExplicit_To(private, Private);
+
+#undef GenericCastToPtrExplicit_To

libclc/libspirv/include/libspirv/spirv.h

@@ -92,4 +92,7 @@
 #include <libspirv/image/image.h>
 #include <libspirv/image/image_defines.h>
 
+/* Pointer Conversion */
+#include <libspirv/conversion/GenericCastToPtrExplicit.h>
+
 #pragma OPENCL EXTENSION all : disable

libclc/libspirv/lib/amdgcn-amdhsa/SOURCES

@@ -14,6 +14,7 @@ atomic/atomic_min.cl
 atomic/atomic_max.cl
 atomic/atomic_sub.cl
 atomic/atomic_store.cl
+conversion/GenericCastToPtrExplicit.cl
 synchronization/barrier.cl
 math/acos.cl
 math/acosh.cl

libclc/libspirv/lib/amdgcn-amdhsa/conversion/GenericCastToPtrExplicit.cl

@@ -0,0 +1,51 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include <libspirv/spirv.h>
+
+
+_CLC_DEF static bool __clc_amdgcn_is_private(generic void *ptr) {
+  return __builtin_amdgcn_is_private(ptr);
+}
+_CLC_DEF static bool __clc_amdgcn_is_local(generic void *ptr) {
+  return __builtin_amdgcn_is_shared(ptr);  
+}
+_CLC_DEF static bool __clc_amdgcn_is_global(generic void *ptr) {
+  return !__clc_amdgcn_is_private(ptr) && !__clc_amdgcn_is_local(ptr);
+}
+
+#define GenericCastToPtrExplicit_To(ADDRSPACE, NAME)                           \
+  _CLC_DECL _CLC_OVERLOAD                                                      \
+      ADDRSPACE void *__spirv_GenericCastToPtrExplicit_To##NAME(               \
+          generic void *ptr, int unused) {                                     \
+    if (__clc_amdgcn_is_##ADDRSPACE(ptr))                                      \
+      return (ADDRSPACE void *)ptr;                                            \
+    return 0;                                                                  \
+  }                                                                            \
+  _CLC_DECL _CLC_OVERLOAD                                                      \
+      ADDRSPACE const void *__spirv_GenericCastToPtrExplicit_To##NAME(         \
+          generic const void *ptr, int unused) {                               \
+    return __spirv_GenericCastToPtrExplicit_To##NAME((generic void *)ptr,      \
+                                                     unused);                  \
+  }                                                                            \
+  _CLC_DECL _CLC_OVERLOAD                                                      \
+      ADDRSPACE volatile void *__spirv_GenericCastToPtrExplicit_To##NAME(      \
+          generic volatile void *ptr, int unused) {                            \
+    return __spirv_GenericCastToPtrExplicit_To##NAME((generic void *)ptr,      \
+                                                     unused);                  \
+  }                                                                            \
+  _CLC_DECL _CLC_OVERLOAD ADDRSPACE const volatile void                        \
+      *__spirv_GenericCastToPtrExplicit_To##NAME(                              \
+          generic const volatile void *ptr, int unused) {                      \
+    return __spirv_GenericCastToPtrExplicit_To##NAME((generic void *)ptr,      \
+                                                     unused);                  \
+  }
+
+GenericCastToPtrExplicit_To(global, Global)
+GenericCastToPtrExplicit_To(local, Local)
+GenericCastToPtrExplicit_To(private, Private)

libclc/libspirv/lib/ptx-nvidiacl/SOURCES

@@ -3,6 +3,7 @@ atomic/loadstore_helpers_release.ll
 atomic/loadstore_helpers_acquire.ll
 atomic/loadstore_helpers_seq_cst.ll
 cl_khr_int64_extended_atomics/minmax_helpers.ll
+conversion/GenericCastToPtrExplicit.cl
 integer/mul24.cl
 integer/mul_hi.cl
 math/acos.cl

libclc/libspirv/lib/ptx-nvidiacl/conversion/GenericCastToPtrExplicit.cl

@@ -0,0 +1,50 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include <libspirv/spirv.h>
+
+_CLC_DEF static bool __clc_nvvm_is_private(generic void *ptr) {
+  return __nvvm_isspacep_local(ptr);
+}
+_CLC_DEF static bool __clc_nvvm_is_local(generic void *ptr) {
+  return __nvvm_isspacep_shared(ptr);
+}
+_CLC_DEF static bool __clc_nvvm_is_global(generic void *ptr) {
+  return __nvvm_isspacep_global(ptr);
+}
+
+#define GenericCastToPtrExplicit_To(ADDRSPACE, NAME)                           \
+  _CLC_DECL _CLC_OVERLOAD                                                      \
+      ADDRSPACE void *__spirv_GenericCastToPtrExplicit_To##NAME(               \
+          generic void *ptr, int unused) {                                     \
+    if (__clc_nvvm_is_##ADDRSPACE(ptr))                                        \
+      return (ADDRSPACE void *)ptr;                                            \
+    return 0;                                                                  \
+  }                                                                            \
+  _CLC_DECL _CLC_OVERLOAD                                                      \
+      ADDRSPACE const void *__spirv_GenericCastToPtrExplicit_To##NAME(         \
+          generic const void *ptr, int unused) {                               \
+    return __spirv_GenericCastToPtrExplicit_To##NAME((generic void *)ptr,      \
+                                                     unused);                  \
+  }                                                                            \
+  _CLC_DECL _CLC_OVERLOAD                                                      \
+      ADDRSPACE volatile void *__spirv_GenericCastToPtrExplicit_To##NAME(      \
+          generic volatile void *ptr, int unused) {                            \
+    return __spirv_GenericCastToPtrExplicit_To##NAME((generic void *)ptr,      \
+                                                     unused);                  \
+  }                                                                            \
+  _CLC_DECL _CLC_OVERLOAD ADDRSPACE const volatile void                        \
+      *__spirv_GenericCastToPtrExplicit_To##NAME(                              \
+          generic const volatile void *ptr, int unused) {                      \
+    return __spirv_GenericCastToPtrExplicit_To##NAME((generic void *)ptr,      \
+                                                     unused);                  \
+  }
+
+GenericCastToPtrExplicit_To(global, Global)
+GenericCastToPtrExplicit_To(local, Local)
+GenericCastToPtrExplicit_To(private, Private)

sycl/include/sycl/access/access.hpp

@@ -350,15 +350,6 @@ address_space_cast_is_possible(access::address_space Src,
 
 template <access::address_space Space, typename ElementType>
 auto static_address_cast(ElementType *Ptr) {
-  constexpr auto generic_space = access::address_space::generic_space;
-  constexpr auto global_space = access::address_space::global_space;
-  constexpr auto local_space = access::address_space::local_space;
-  constexpr auto private_space = access::address_space::private_space;
-  constexpr auto global_device =
-      access::address_space::ext_intel_global_device_space;
-  constexpr auto global_host =
-      access::address_space::ext_intel_global_host_space;
-
   constexpr auto SrcAS = deduce_AS<ElementType *>::value;
   static_assert(address_space_cast_is_possible(SrcAS, Space));
 
@@ -367,31 +358,7 @@ auto static_address_cast(ElementType *Ptr) {
 
   // Note: reinterpret_cast isn't enough for some of the casts between different
   // address spaces, use C-style cast instead.
-#if !defined(__SPIR__)
   return (dst_type)Ptr;
-#else
-  if constexpr (SrcAS != generic_space) {
-    return (dst_type)Ptr;
-  } else if constexpr (Space == global_space) {
-    return (dst_type)__spirv_GenericCastToPtr_ToGlobal(
-        Ptr, __spv::StorageClass::CrossWorkgroup);
-  } else if constexpr (Space == local_space) {
-    return (dst_type)__spirv_GenericCastToPtr_ToLocal(
-        Ptr, __spv::StorageClass::Workgroup);
-  } else if constexpr (Space == private_space) {
-    return (dst_type)__spirv_GenericCastToPtr_ToPrivate(
-        Ptr, __spv::StorageClass::Function);
-#if !defined(__ENABLE_USM_ADDR_SPACE__)
-  } else if constexpr (Space == global_device || Space == global_host) {
-    // If __ENABLE_USM_ADDR_SPACE__ isn't defined then both
-    // global_device/global_host are just aliases for global_space.
-    return (dst_type)__spirv_GenericCastToPtr_ToGlobal(
-        Ptr, __spv::StorageClass::CrossWorkgroup);
-#endif
-  } else {
-    return (dst_type)Ptr;
-  }
-#endif
 }
 
 // Previous implementation (`castAS`, used in `multi_ptr` ctors among other
@@ -427,14 +394,13 @@ auto dynamic_address_cast(ElementType *Ptr) {
 #if defined(__ENABLE_USM_ADDR_SPACE__)
     static_assert(SupressNotImplementedAssert || Space != Space,
                   "Not supported yet!");
-    return static_address_cast<Space>(Ptr);
+    return detail::static_address_cast<Space>(Ptr);
 #else
     // If __ENABLE_USM_ADDR_SPACE__ isn't defined then both
     // global_device/global_host are just aliases for global_space.
     static_assert(std::is_same_v<dst_type, ElementType *>);
     return (dst_type)Ptr;
 #endif
-#if defined(__SPIR__)
   } else if constexpr (Space == global_space) {
     return (dst_type)__spirv_GenericCastToPtrExplicit_ToGlobal(
         Ptr, __spv::StorageClass::CrossWorkgroup);
@@ -449,12 +415,11 @@ auto dynamic_address_cast(ElementType *Ptr) {
                        (Space == global_device || Space == global_host)) {
     return (dst_type)__spirv_GenericCastToPtrExplicit_ToGlobal(
         Ptr, __spv::StorageClass::CrossWorkgroup);
-#endif
 #endif
   } else {
     static_assert(SupressNotImplementedAssert || Space != Space,
                   "Not supported yet!");
-    return static_address_cast<Space>(Ptr);
+    return detail::static_address_cast<Space>(Ptr);
   }
 }
 #else  // __SYCL_DEVICE_ONLY__

sycl/test-e2e/AddressCast/dynamic_address_cast.cpp

@@ -7,9 +7,8 @@
 //===----------------------------------------------------------------------===//
 
 // Issue with OpenCL CPU runtime implementation of OpGenericCastToPtrExplicit
-// OpGenericCastToPtr* intrinsics not implemented on AMD or NVIDIA
 // FPGA emulator affected by same issue as OpenCL CPU runtime
-// UNSUPPORTED: cpu, hip, cuda, accelerator
+// UNSUPPORTED: cpu, accelerator
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 

sycl/test/check_device_code/atomic_ref.cpp

@@ -6,8 +6,8 @@
 // CHECK-LABEL: define dso_local spir_func noundef i32 @_Z17atomic_ref_globalRi(
 // CHECK-SAME: ptr addrspace(4) noundef align 4 dereferenceable(4) [[I:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] !srcloc [[META6:![0-9]+]] !sycl_fixed_targets [[META7:![0-9]+]] {
 // CHECK-NEXT:  [[ENTRY:.*:]]
-// CHECK-NEXT:    [[CALL_I_I_I_I_I_I:%.*]] = tail call spir_func noundef ptr addrspace(1) @_Z33__spirv_GenericCastToPtr_ToGlobalPvi(ptr addrspace(4) noundef align 4 dereferenceable(4) [[I]], i32 noundef 5) #[[ATTR3:[0-9]+]]
-// CHECK-NEXT:    [[CALL3_I_I:%.*]] = tail call spir_func noundef i32 @_Z18__spirv_AtomicLoadPU3AS1KiN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE(ptr addrspace(1) noundef [[CALL_I_I_I_I_I_I]], i32 noundef 1, i32 noundef 898) #[[ATTR4:[0-9]+]]
+// CHECK-NEXT:    [[TMP:%.*]] = addrspacecast ptr addrspace(4) [[I]] to ptr addrspace(1)
+// CHECK-NEXT:    [[CALL3_I_I:%.*]] = tail call spir_func noundef i32 @_Z18__spirv_AtomicLoadPU3AS1KiN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagE(ptr addrspace(1) noundef [[TMP]], i32 noundef 1, i32 noundef 898) #[[ATTR4:[0-9]+]]
 // CHECK-NEXT:    ret i32 [[CALL3_I_I]]
 //
 SYCL_EXTERNAL auto atomic_ref_global(int &i) {