From 82d874f49293d342eb2e2ffe2eb8174778d4ceab Mon Sep 17 00:00:00 2001
From: Victor Perez <victor.perez@codeplay.com>
Date: Mon, 11 Nov 2024 11:05:19 +0100
Subject: [PATCH] [SYCL] Add `(raw|decorated)_generic_ptr` aliases (#15389)

https://github.com/KhronosGroup/SYCL-Docs/pull/598 adds these aliases to
the spec. Prepare for spec change by adding them to the implementation
beforehand.

Extend tests so all `multi_ptr` aliases are checked.

---------

Signed-off-by: Victor Perez <victor.perez@codeplay.com>
---
 sycl/include/sycl/pointers.hpp                | 10 +++
 .../extensions/address_cast.cpp               |  6 --
 sycl/test/multi_ptr/aliases.cpp               | 87 +++++++++++++++++++
 3 files changed, 97 insertions(+), 6 deletions(-)
 create mode 100644 sycl/test/multi_ptr/aliases.cpp

diff --git a/sycl/include/sycl/pointers.hpp b/sycl/include/sycl/pointers.hpp
index c74a8e104ab82..cd082b9f79904 100644
--- a/sycl/include/sycl/pointers.hpp
+++ b/sycl/include/sycl/pointers.hpp
@@ -47,6 +47,11 @@ using private_ptr =
 // The interface exposes non-decorated pointer while keeping the
 // address space information internally.
 
+template <typename ElementType>
+using raw_generic_ptr =
+    multi_ptr<ElementType, access::address_space::generic_space,
+              access::decorated::no>;
+
 template <typename ElementType>
 using raw_global_ptr =
     multi_ptr<ElementType, access::address_space::global_space,
@@ -64,6 +69,11 @@ using raw_private_ptr =
 // Template specialization aliases for different pointer address spaces.
 // The interface exposes decorated pointer.
 
+template <typename ElementType>
+using decorated_generic_ptr =
+    multi_ptr<ElementType, access::address_space::generic_space,
+              access::decorated::yes>;
+
 template <typename ElementType>
 using decorated_global_ptr =
     multi_ptr<ElementType, access::address_space::global_space,
diff --git a/sycl/test/check_device_code/extensions/address_cast.cpp b/sycl/test/check_device_code/extensions/address_cast.cpp
index 36251cecbbc3b..c265a5e2fe74b 100644
--- a/sycl/test/check_device_code/extensions/address_cast.cpp
+++ b/sycl/test/check_device_code/extensions/address_cast.cpp
@@ -11,12 +11,6 @@
 using namespace sycl;
 using namespace sycl::ext::oneapi::experimental;
 
-// FIXME: should be removed when https://github.com/intel/llvm/pull/15389 is merged in.
-template <typename ElementType>
-using decorated_generic_ptr =
-    multi_ptr<ElementType, access::address_space::generic_space,
-              access::decorated::yes>;
-
 namespace static_as_cast {
 // CHECK-LABEL: define dso_local spir_func void @_ZN14static_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:[0-9]+]] !srcloc [[META6:![0-9]+]] !sycl_fixed_targets [[META7:![0-9]+]] {
diff --git a/sycl/test/multi_ptr/aliases.cpp b/sycl/test/multi_ptr/aliases.cpp
new file mode 100644
index 0000000000000..1911602e74386
--- /dev/null
+++ b/sycl/test/multi_ptr/aliases.cpp
@@ -0,0 +1,87 @@
+//==--------------- aliases.cpp - SYCL multi_ptr aliases test --------------==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s -Xclang -verify-ignore-unexpected=note,warning
+
+// expected-no-diagnostics
+
+#include <sycl/detail/core.hpp>
+
+#include <type_traits>
+
+template <typename T, sycl::access::decorated IsDecorated>
+void test_address_space_aliases() {
+  static_assert(std::is_same_v<
+                sycl::generic_ptr<T, IsDecorated>,
+                sycl::multi_ptr<T, sycl::access::address_space::generic_space,
+                                IsDecorated>>);
+  static_assert(std::is_same_v<
+                sycl::global_ptr<T, IsDecorated>,
+                sycl::multi_ptr<T, sycl::access ::address_space::global_space,
+                                IsDecorated>>);
+  static_assert(std::is_same_v<
+                sycl::local_ptr<T, IsDecorated>,
+                sycl::multi_ptr<T, sycl::access::address_space::local_space,
+                                IsDecorated>>);
+  static_assert(std::is_same_v<
+                sycl::private_ptr<T, IsDecorated>,
+                sycl::multi_ptr<T, sycl::access::address_space::private_space,
+                                IsDecorated>>);
+}
+
+template <typename T> void test_aliases() {
+  // Template specialization aliases for different pointer address spaces
+  test_address_space_aliases<T, sycl::access::decorated::yes>();
+  test_address_space_aliases<T, sycl::access::decorated::no>();
+  test_address_space_aliases<T, sycl::access::decorated::legacy>();
+
+  // Template specialization aliases for different pointer address spaces.
+  // The interface exposes non-decorated pointer while keeping the
+  // address space information internally.
+  static_assert(std::is_same_v<
+                sycl::raw_generic_ptr<T>,
+                sycl::multi_ptr<T, sycl::access::address_space::generic_space,
+                                sycl::access::decorated::no>>);
+  static_assert(std::is_same_v<
+                sycl::raw_global_ptr<T>,
+                sycl::multi_ptr<T, sycl::access::address_space::global_space,
+                                sycl::access::decorated::no>>);
+  static_assert(std::is_same_v<
+                sycl::raw_local_ptr<T>,
+                sycl::multi_ptr<T, sycl::access::address_space::local_space,
+                                sycl::access::decorated::no>>);
+  static_assert(std::is_same_v<
+                sycl::raw_private_ptr<T>,
+                sycl::multi_ptr<T, sycl::access::address_space::private_space,
+                                sycl::access::decorated::no>>);
+
+  // Template specialization aliases for different pointer address spaces.
+  // The interface exposes decorated pointer.
+  static_assert(std::is_same_v<
+                sycl::decorated_generic_ptr<T>,
+                sycl::multi_ptr<T, sycl::access::address_space::generic_space,
+                                sycl::access::decorated::yes>>);
+  static_assert(std::is_same_v<
+                sycl::decorated_global_ptr<T>,
+                sycl::multi_ptr<T, sycl::access::address_space::global_space,
+                                sycl::access::decorated::yes>>);
+  static_assert(std::is_same_v<
+                sycl::decorated_local_ptr<T>,
+                sycl::multi_ptr<T, sycl::access::address_space::local_space,
+                                sycl::access::decorated::yes>>);
+  static_assert(std::is_same_v<
+                sycl::decorated_private_ptr<T>,
+                sycl::multi_ptr<T, sycl::access::address_space::private_space,
+                                sycl::access::decorated::yes>>);
+}
+
+// Test "minimal set of types" in the CTS. As we are just testing aliases are
+// present in this test, this should work for any type.
+
+template void test_aliases<int>();
+template void test_aliases<float>();