[SYCL] Limit decorations to init function for annotated_arg/ptr classes (#11881)

annotated_arg/annotated_ptr are special classes and arguments of the
__init() member function become kernel arguments. The classes have
decorated_global_ptr or other such type applied on the class members to
indicate that these must be in global address-space (or other address
space as indicated by the applied type). The goal of applying the type
on the member was to ensure that class member gets captured as a global
address space pointer when set as kernel argument.

However, the type was applied on class constructor arguments as well.
This PR restricts the type to only be present on the __init member
function argument, otherwise the presence of the type on constructors
causes functional errors on certain backends when objects are
constructed in the kernel body. The object constructed inside kernel
body also gets a global address space cast generated by the compiler
when the decorated_global_ptr type is applied on ctor arguments. This is
problematic for example when the object is constructed from a pointer to
some array defined inside the kernel.

Author: tiwaria1

sycl/include/sycl/ext/oneapi/experimental/annotated_arg/annotated_arg.hpp

@@ -66,20 +66,15 @@ class __SYCL_SPECIAL_CLASS
 __SYCL_TYPE(annotated_arg) annotated_arg<T *, detail::properties_t<Props...>> {
   using property_list_t = detail::properties_t<Props...>;
 
-#ifdef __SYCL_DEVICE_ONLY__
-  using global_pointer_t = typename decorated_global_ptr<T>::pointer;
-#else
-  using global_pointer_t = T *;
-#endif
-
-  global_pointer_t obj;
+  T *obj;
 
   template <typename T2, typename PropertyListT> friend class annotated_arg;
 
 #ifdef __SYCL_DEVICE_ONLY__
   void __init([[__sycl_detail__::add_ir_attributes_kernel_parameter(
       detail::PropertyMetaInfo<Props>::name...,
-      detail::PropertyMetaInfo<Props>::value...)]] global_pointer_t _obj) {
+      detail::PropertyMetaInfo<Props>::value...)]]
+              typename decorated_global_ptr<T>::pointer _obj) {
     obj = _obj;
   }
 #endif
@@ -91,7 +86,7 @@ __SYCL_TYPE(annotated_arg) annotated_arg<T *, detail::properties_t<Props...>> {
 
   annotated_arg(T *_ptr,
                 const property_list_t &PropList = properties{}) noexcept
-      : obj(global_pointer_t(_ptr)) {
+      : obj(_ptr) {
     (void)PropList;
   }
 
@@ -100,8 +95,7 @@ __SYCL_TYPE(annotated_arg) annotated_arg<T *, detail::properties_t<Props...>> {
   // variadic properties. The same property in `Props...` and
   // `PropertyValueTs...` must have the same property value.
   template <typename... PropertyValueTs>
-  annotated_arg(T *_ptr, const PropertyValueTs &...props) noexcept
-      : obj(global_pointer_t(_ptr)) {
+  annotated_arg(T *_ptr, const PropertyValueTs &...props) noexcept : obj(_ptr) {
     static constexpr bool has_same_properties = std::is_same<
         property_list_t,
         detail::merged_properties_t<property_list_t,

sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp

@@ -190,7 +190,6 @@ __SYCL_TYPE(annotated_ptr) annotated_ptr<T, detail::properties_t<Props...>> {
   using reference = sycl::ext::oneapi::experimental::annotated_ref<
       T, typename unpack<filtered_properties>::type>;
 
-#ifdef __SYCL_DEVICE_ONLY__
 #ifdef __ENABLE_USM_ADDR_SPACE__
   using global_pointer_t = std::conditional_t<
       detail::IsUsmKindDevice<property_list_t>::value,
@@ -202,11 +201,8 @@ __SYCL_TYPE(annotated_ptr) annotated_ptr<T, detail::properties_t<Props...>> {
 #else
   using global_pointer_t = typename decorated_global_ptr<T>::pointer;
 #endif // __ENABLE_USM_ADDR_SPACE__
-#else
-  using global_pointer_t = T *;
-#endif // __SYCL_DEVICE_ONLY__
 
-  global_pointer_t m_Ptr;
+  T *m_Ptr;
 
   template <typename T2, typename PropertyListT> friend class annotated_ptr;
 
@@ -225,15 +221,15 @@ __SYCL_TYPE(annotated_ptr) annotated_ptr<T, detail::properties_t<Props...>> {
 
   explicit annotated_ptr(T *Ptr,
                          const property_list_t & = properties{}) noexcept
-      : m_Ptr(global_pointer_t(Ptr)) {}
+      : m_Ptr(Ptr) {}
 
   // Constructs an annotated_ptr object from a raw pointer and variadic
   // properties. The new property set contains all properties of the input
   // variadic properties. The same property in `Props...` and
   // `PropertyValueTs...` must have the same property value.
   template <typename... PropertyValueTs>
   explicit annotated_ptr(T *Ptr, const PropertyValueTs &...props) noexcept
-      : m_Ptr(global_pointer_t(Ptr)) {
+      : m_Ptr(Ptr) {
     static constexpr bool has_same_properties = std::is_same<
         property_list_t,
         detail::merged_properties_t<property_list_t,