[SYCL] Add missing generic cast builtin for non-const volatile types (#7505)

The headers use builtins for casting from generic pointers to other
address spaces. However, of these definitions it is missing variants for
non-const volatile. This commit adds these missing definitions.

Signed-off-by: Larsen, Steffen <[email protected]>

Author: steffenlarsen

sycl/include/CL/__spirv/spirv_ops.hpp

@@ -415,6 +415,14 @@ __SYCL_GenericCastToPtrExplicit_ToGlobal(const void *Ptr) noexcept {
           Ptr, __spv::StorageClass::CrossWorkgroup);
 }
 
+template <typename dataT>
+extern volatile __attribute__((opencl_global)) dataT *
+__SYCL_GenericCastToPtrExplicit_ToGlobal(volatile void *Ptr) noexcept {
+  return (volatile __attribute__((opencl_global)) dataT *)
+      __spirv_GenericCastToPtrExplicit_ToGlobal(
+          Ptr, __spv::StorageClass::CrossWorkgroup);
+}
+
 template <typename dataT>
 extern const volatile __attribute__((opencl_global)) dataT *
 __SYCL_GenericCastToPtrExplicit_ToGlobal(const volatile void *Ptr) noexcept {
@@ -439,6 +447,14 @@ __SYCL_GenericCastToPtrExplicit_ToLocal(const void *Ptr) noexcept {
                                                __spv::StorageClass::Workgroup);
 }
 
+template <typename dataT>
+extern volatile __attribute__((opencl_local)) dataT *
+__SYCL_GenericCastToPtrExplicit_ToLocal(volatile void *Ptr) noexcept {
+  return (volatile __attribute__((opencl_local)) dataT *)
+      __spirv_GenericCastToPtrExplicit_ToLocal(Ptr,
+                                               __spv::StorageClass::Workgroup);
+}
+
 template <typename dataT>
 extern const volatile __attribute__((opencl_local)) dataT *
 __SYCL_GenericCastToPtrExplicit_ToLocal(const volatile void *Ptr) noexcept {
@@ -463,6 +479,14 @@ __SYCL_GenericCastToPtrExplicit_ToPrivate(const void *Ptr) noexcept {
                                                  __spv::StorageClass::Function);
 }
 
+template <typename dataT>
+extern volatile __attribute__((opencl_private)) dataT *
+__SYCL_GenericCastToPtrExplicit_ToPrivate(volatile void *Ptr) noexcept {
+  return (volatile __attribute__((opencl_private)) dataT *)
+      __spirv_GenericCastToPtrExplicit_ToPrivate(Ptr,
+                                                 __spv::StorageClass::Function);
+}
+
 template <typename dataT>
 extern const volatile __attribute__((opencl_private)) dataT *
 __SYCL_GenericCastToPtrExplicit_ToPrivate(const volatile void *Ptr) noexcept {

sycl/test/regression/multi_ptr_gen_casting.cpp

@@ -0,0 +1,66 @@
+// RUN: %clangxx -fsycl %s -o %t.out
+//
+// Tests that casting multi_ptr to and from generic compiles for various
+// combinations of valid qualifiers.
+
+#include <sycl/sycl.hpp>
+
+using namespace sycl;
+
+template <typename T, access::address_space AddrSpace,
+          sycl::access::decorated IsDecorated>
+void test(queue &Q) {
+  T *GlobPtr = malloc_device<T>(1, Q);
+  Q.submit([&](handler &CGH) {
+     local_accessor<T> LocPtr{1, CGH};
+     CGH.single_task([=]() {
+       T X = 0;
+       T *InPtr;
+       if constexpr (AddrSpace == access::address_space::global_space)
+         InPtr = GlobPtr;
+       else if constexpr (AddrSpace == access::address_space::local_space)
+         InPtr = LocPtr.get_pointer();
+       else
+         InPtr = &X;
+
+       auto MPtr = address_space_cast<AddrSpace, IsDecorated>(InPtr);
+       multi_ptr<T, access::address_space::generic_space, IsDecorated> GenPtr;
+       GenPtr = MPtr;
+       MPtr = multi_ptr<T, AddrSpace, IsDecorated>{GenPtr};
+     });
+   }).wait();
+}
+
+template <typename T, access::address_space AddrSpace>
+void testAllDecos(queue &Q) {
+  test<T, AddrSpace, sycl::access::decorated::yes>(Q);
+  test<T, AddrSpace, sycl::access::decorated::no>(Q);
+}
+
+template <typename T> void testAllAddrSpace(queue &Q) {
+  testAllDecos<T, access::address_space::private_space>(Q);
+  testAllDecos<T, access::address_space::local_space>(Q);
+  testAllDecos<T, access::address_space::global_space>(Q);
+}
+
+template <typename T> void testAllQuals(queue &Q) {
+  using UnqualT = std::remove_cv_t<T>;
+  testAllAddrSpace<UnqualT>(Q);
+  testAllAddrSpace<std::add_const_t<UnqualT>>(Q);
+  testAllAddrSpace<std::add_volatile_t<UnqualT>>(Q);
+  testAllAddrSpace<std::add_cv_t<UnqualT>>(Q);
+}
+
+int main() {
+  queue Q;
+  testAllQuals<bool>(Q);
+  testAllQuals<char>(Q);
+  testAllQuals<short>(Q);
+  testAllQuals<int>(Q);
+  testAllQuals<long>(Q);
+  testAllQuals<long long>(Q);
+  testAllQuals<sycl::half>(Q);
+  testAllQuals<float>(Q);
+  testAllQuals<double>(Q);
+  return 0;
+}