[SYCL] Support *global_[device|host]_space in static_address_cast…

… (#15498) When these address spaces are used with regular `sycl::detail::spirv::GenericCastToPtr` they are turned into `unreachable`. We don't have a SPIR-V intrinsic yet (or maybe we shouldn't even have it, and will continue to rely on standard LLVM IR's `addrspacecast`), so use C-style cast and rely on the translator/backend to generate proper operation, similarly to `sycl::detail::cast_AS`.
oneapi-src · Sep 27, 2024 · 23fed07 · 23fed07
1 parent c956940
commit 23fed07
Show file tree

Hide file tree

Showing 3 changed files with 74 additions and 9 deletions.
diff --git a/sycl/include/sycl/ext/oneapi/experimental/address_cast.hpp b/sycl/include/sycl/ext/oneapi/experimental/address_cast.hpp
@@ -16,6 +16,9 @@ inline namespace _V1 {
 namespace ext {
 namespace oneapi {
 namespace experimental {
+namespace detail {
+using namespace sycl::detail;
+}
 // Shorthands for address space names
 constexpr inline access::address_space global_space = access::address_space::global_space;
 constexpr inline access::address_space local_space = access::address_space::local_space;
@@ -32,6 +35,18 @@ static_address_cast(ElementType *Ptr) {
   if constexpr (Space == generic_space) {
     // Undecorated raw pointer is in generic AS already, no extra casts needed.
     return ret_ty(Ptr);
+  } else if constexpr (Space == access::address_space::
+                                    ext_intel_global_device_space ||
+                       Space ==
+                           access::address_space::ext_intel_global_host_space) {
+#ifdef __ENABLE_USM_ADDR_SPACE__
+    // No SPIR-V intrinsic for this yet.
+    using raw_type = detail::DecoratedType<ElementType, Space>::type *;
+    auto CastPtr = (raw_type)(Ptr);
+#else
+    auto CastPtr = sycl::detail::spirv::GenericCastToPtr<global_space>(Ptr);
+#endif
+    return ret_ty(CastPtr);
   } else {
     auto CastPtr = sycl::detail::spirv::GenericCastToPtr<Space>(Ptr);
     return ret_ty(CastPtr);
@@ -60,6 +75,20 @@ dynamic_address_cast(ElementType *Ptr) {
                 "The extension expects undecorated raw pointers only!");
   if constexpr (Space == generic_space) {
     return ret_ty(Ptr);
+  } else if constexpr (Space == access::address_space::
+                                    ext_intel_global_device_space ||
+                       Space ==
+                           access::address_space::ext_intel_global_host_space) {
+#ifdef __ENABLE_USM_ADDR_SPACE__
+    static_assert(
+        Space != access::address_space::ext_intel_global_device_space &&
+            Space != access::address_space::ext_intel_global_host_space,
+        "Not supported yet!");
+    return ret_ty(nullptr);
+#else
+    auto CastPtr = sycl::detail::spirv::GenericCastToPtr<global_space>(Ptr);
+    return ret_ty(CastPtr);
+#endif
   } else {
     auto CastPtr = sycl::detail::spirv::GenericCastToPtrExplicit<Space>(Ptr);
     return ret_ty(CastPtr);

diff --git a/sycl/test/check_device_code/extensions/address_cast.cpp b/sycl/test/check_device_code/extensions/address_cast.cpp
@@ -1,5 +1,5 @@
 // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
-// RUN: %clangxx -O3 -fsycl -fsycl-device-only -fno-discard-value-names -S -emit-llvm -fno-sycl-instrument-device-code  -o - %s | FileCheck %s
+// RUN: %clangxx -D__ENABLE_USM_ADDR_SPACE__ -O3 -fsycl -fsycl-device-only -fno-discard-value-names -S -emit-llvm -fno-sycl-instrument-device-code  -o - %s | FileCheck %s
 
 // Linux/Windows have minor differences in the generated IR (e.g. TBAA
 // metadata). Having linux-only checks eases the maintenance without sacrifising
@@ -59,45 +59,67 @@ SYCL_EXTERNAL auto to_generic_decorated(decorated_generic_ptr<int> p) {
 SYCL_EXTERNAL auto to_generic_not_decorated(int *p) {
   return static_address_cast<access::address_space::generic_space>(p);
 }
+
+// CHECK-LABEL: define dso_local spir_func void @_ZN14static_as_cast16to_global_deviceEPi(
+// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.3") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR4]] !srcloc [[META35:![0-9]+]] !sycl_fixed_targets [[META7]] {
+// CHECK-NEXT:  [[ENTRY:.*:]]
+// CHECK-NEXT:    [[TMP0:%.*]] = addrspacecast ptr addrspace(4) [[P]] to ptr addrspace(5)
+// CHECK-NEXT:    store ptr addrspace(5) [[TMP0]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA36:![0-9]+]], !alias.scope [[META38:![0-9]+]]
+// CHECK-NEXT:    ret void
+//
+SYCL_EXTERNAL auto to_global_device(int *p) {
+  return static_address_cast<access::address_space::ext_intel_global_device_space>(p);
+}
+
+// CHECK-LABEL: define dso_local spir_func void @_ZN14static_as_cast14to_global_hostEPi(
+// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.4") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR4]] !srcloc [[META41:![0-9]+]] !sycl_fixed_targets [[META7]] {
+// CHECK-NEXT:  [[ENTRY:.*:]]
+// CHECK-NEXT:    [[TMP0:%.*]] = addrspacecast ptr addrspace(4) [[P]] to ptr addrspace(6)
+// CHECK-NEXT:    store ptr addrspace(6) [[TMP0]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA42:![0-9]+]], !alias.scope [[META44:![0-9]+]]
+// CHECK-NEXT:    ret void
+//
+SYCL_EXTERNAL auto to_global_host(int *p) {
+  return static_address_cast<access::address_space::ext_intel_global_host_space>(p);
+}
 } // namespace static_as_cast
 
 namespace dynamic_as_cast {
 // CHECK-LABEL: define dso_local spir_func void @_ZN15dynamic_as_cast19to_global_decoratedEN4sycl3_V19multi_ptrIiLNS1_6access13address_spaceE6ELNS3_9decoratedE1EEE(
-// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr") align 8 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::multi_ptr.0") align 8 [[P:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META35:![0-9]+]] !sycl_fixed_targets [[META7]] {
+// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr") align 8 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::multi_ptr.0") align 8 [[P:%.*]]) local_unnamed_addr #[[ATTR0]] !srcloc [[META47:![0-9]+]] !sycl_fixed_targets [[META7]] {
 // CHECK-NEXT:  [[ENTRY:.*:]]
 // CHECK-NEXT:    [[TMP0:%.*]] = load i64, ptr [[P]], align 8, !tbaa [[TBAA8]]
 // CHECK-NEXT:    [[TMP1:%.*]] = inttoptr i64 [[TMP0]] to ptr addrspace(4)
 // CHECK-NEXT:    [[CALL_I_I_I_I:%.*]] = tail call spir_func noundef ptr addrspace(1) @_Z41__spirv_GenericCastToPtrExplicit_ToGlobalPvi(ptr addrspace(4) noundef [[TMP1]], i32 noundef 5) #[[ATTR5]]
-// CHECK-NEXT:    store ptr addrspace(1) [[CALL_I_I_I_I]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA12]], !alias.scope [[META36:![0-9]+]]
+// CHECK-NEXT:    store ptr addrspace(1) [[CALL_I_I_I_I]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA12]], !alias.scope [[META48:![0-9]+]]
 // CHECK-NEXT:    ret void
 //
 SYCL_EXTERNAL auto to_global_decorated(decorated_generic_ptr<int> p) {
   return dynamic_address_cast<access::address_space::global_space>(p);
 }
 // CHECK-LABEL: define dso_local spir_func void @_ZN15dynamic_as_cast23to_global_not_decoratedEPi(
-// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.1") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR2]] !srcloc [[META41:![0-9]+]] !sycl_fixed_targets [[META7]] {
+// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.1") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR2]] !srcloc [[META53:![0-9]+]] !sycl_fixed_targets [[META7]] {
 // CHECK-NEXT:  [[ENTRY:.*:]]
 // CHECK-NEXT:    [[CALL_I_I_I:%.*]] = tail call spir_func noundef ptr addrspace(1) @_Z41__spirv_GenericCastToPtrExplicit_ToGlobalPvi(ptr addrspace(4) noundef [[P]], i32 noundef 5) #[[ATTR5]]
-// CHECK-NEXT:    store ptr addrspace(1) [[CALL_I_I_I]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA20]], !alias.scope [[META42:![0-9]+]]
+// CHECK-NEXT:    store ptr addrspace(1) [[CALL_I_I_I]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA20]], !alias.scope [[META54:![0-9]+]]
 // CHECK-NEXT:    ret void
 //
 SYCL_EXTERNAL auto to_global_not_decorated(int *p) {
   return dynamic_address_cast<access::address_space::global_space>(p);
 }
 // CHECK-LABEL: define dso_local spir_func void @_ZN15dynamic_as_cast20to_generic_decoratedEN4sycl3_V19multi_ptrIiLNS1_6access13address_spaceE6ELNS3_9decoratedE1EEE(
-// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.0") align 8 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::multi_ptr.0") align 8 [[P:%.*]]) local_unnamed_addr #[[ATTR3]] !srcloc [[META45:![0-9]+]] !sycl_fixed_targets [[META7]] {
+// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.0") align 8 [[AGG_RESULT:%.*]], ptr nocapture noundef readonly byval(%"class.sycl::_V1::multi_ptr.0") align 8 [[P:%.*]]) local_unnamed_addr #[[ATTR3]] !srcloc [[META57:![0-9]+]] !sycl_fixed_targets [[META7]] {
 // CHECK-NEXT:  [[ENTRY:.*:]]
 // CHECK-NEXT:    [[TMP0:%.*]] = load i64, ptr [[P]], align 8, !tbaa [[TBAA8]]
-// CHECK-NEXT:    store i64 [[TMP0]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA8]], !alias.scope [[META46:![0-9]+]]
+// CHECK-NEXT:    store i64 [[TMP0]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA8]], !alias.scope [[META58:![0-9]+]]
 // CHECK-NEXT:    ret void
 //
 SYCL_EXTERNAL auto to_generic_decorated(decorated_generic_ptr<int> p) {
   return dynamic_address_cast<access::address_space::generic_space>(p);
 }
 // CHECK-LABEL: define dso_local spir_func void @_ZN15dynamic_as_cast24to_generic_not_decoratedEPi(
-// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.2") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR4]] !srcloc [[META49:![0-9]+]] !sycl_fixed_targets [[META7]] {
+// CHECK-SAME: ptr addrspace(4) dead_on_unwind noalias nocapture writable writeonly sret(%"class.sycl::_V1::multi_ptr.2") align 8 [[AGG_RESULT:%.*]], ptr addrspace(4) noundef [[P:%.*]]) local_unnamed_addr #[[ATTR4]] !srcloc [[META61:![0-9]+]] !sycl_fixed_targets [[META7]] {
 // CHECK-NEXT:  [[ENTRY:.*:]]
-// CHECK-NEXT:    store ptr addrspace(4) [[P]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA30]], !alias.scope [[META50:![0-9]+]]
+// CHECK-NEXT:    store ptr addrspace(4) [[P]], ptr addrspace(4) [[AGG_RESULT]], align 8, !tbaa [[TBAA30]], !alias.scope [[META62:![0-9]+]]
 // CHECK-NEXT:    ret void
 //
 SYCL_EXTERNAL auto to_generic_not_decorated(int *p) {

diff --git a/sycl/test/extensions/address_cast_negative.cpp b/sycl/test/extensions/address_cast_negative.cpp
@@ -0,0 +1,14 @@
+// RUN: %clangxx -D__ENABLE_USM_ADDR_SPACE__ -fsycl -fsycl-device-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=warning,note %s
+
+#include <sycl/sycl.hpp>
+
+using namespace sycl::ext::oneapi::experimental;
+
+SYCL_EXTERNAL void test(int *p) {
+  // expected-error-re@sycl/ext/oneapi/experimental/address_cast.hpp:* {{{{.*}}Not supported yet!}}
+  std::ignore = dynamic_address_cast<
+      sycl::access::address_space::ext_intel_global_device_space>(p);
+  // expected-error-re@sycl/ext/oneapi/experimental/address_cast.hpp:* {{{{.*}}Not supported yet!}}
+  std::ignore = dynamic_address_cast<
+      sycl::access::address_space::ext_intel_global_host_space>(p);
+}