Merge pull request #432 from AerialMantis/SYCL-2020/multi-ptr-errata

We have discovered that some of the changes in SYCL 2020 for multi_ptr have introduced breaking changes for SYCL 1.2.1 applications moving to SYCL 2020. At the time of these changes it was thought to be necessary but since then the SYCL working group has adopted a philosophy of deprecating interfaces before removing them outright.

This pull request aims to address these changes, such that SYCL 1.2.1 code will still be compatible with SYCL 2020 implementations, avoiding errors in moving from one version to the other and in portability between implementations.
1. Default access::decorated value for multi_ptr

The first issue that was identified was that the SYCL 2020 wording introduces the access::decorated template parameter in multi_ptr, however, does not provide a default. This means that any code which declares the multi_ptr without this template argument, as was the case in SYCL 1.2.1 would no longer compiler.

For example:

multi_ptr<float, access::address_space::global_space> multiPtr;

To solve this issue, this PR has the multi_ptr class template default to access::decorated::legacy. It also specifies in the access::decorated enum class that access::decorated::legacy is deprecated.
2. accessor returning a raw pointer

The second issue that was identified was that the SYCL 2020 wording changes the return type of accessor::get_pointer from a multi_ptr to a raw pointer. This means that any code which expects the return value of accessor::get_pointer to be a multi_ptr will fail to compile.

For example:

multi_ptr<float, access::address_space::global_space> multiPtr = acc.get_pointer();

or

float f4;
f4.load(offset, acc.get_pointer());

To solve this issue, this PR reverts back to accessor::get_pointer returning a multi_ptr with access::decorated::legacy. It also introduces wording for accessor::get_pointer which was previously missing and specifies that it is deprecated.
3. Missing interfaces for multi_ptr with access::decorated::legacy

The third issue that was identified was that the SYCL 2020 wording introduces new interfaces for the multi_ptr class; specifically the multi_ptr::get_raw and multi_ptr::get_decorated member functions, but not for the access::decorated::legacy specialization. This means any new code which uses these member functions will not be compatible with old code which uses the access::decorated::legacy specialization of multi_ptr and will therefore have to multi-version for the two specializations.

For example:

template <typename T, access::address_space AddressSpace, access::decorated IsDecorated>
void foo(multi_ptr<T, AddressSpace,IsDecorated> multiPtr) {
  T *ptr = multiPtr.get_decorated();
  // do something with ptr
}

To solve this issue, this PR adds the multi_ptr::get_raw and multi_ptr::get_decorated member functions to the access::decorated::legacy specialization of multi_ptr as well.
4. async_work_group_copy requiring access::decorated::yes

The fourth issue that was identified was that the SYCL 2020 wording changes the nd_item::async_work_group_copy and group::async_work_group_copy member functions such that they now take a multi_ptr with access::decorated::yes explicitly. This means that any code which uses the access::decorated::legacy specialization will not be able to call these functions.

For example:

multi_ptr<float, access::address_space::global_space> srcPtr = srcAcc.get_pointer();
multi_ptr<float, access::address_space::global_space> destPtr= destAcc.get_pointer();
async_work_group_copy(destPtr, srcPtr, numElements);

To solve this issue, this PR re-introduces the original SYCL 1.2.1 overloads of  async_work_group_copy which take a multi_ptr with access::decorated::legacy, and specifies them as deprecated.

Note, the examples for these are based on the previous changes being applied, however, those issues are all reproducible even if they weren't using other SYCL 2020 interfaces.
Edits

    b84963d: Have local_accessor::get_pointer() also return multi_ptr` for consistency and address comments.
    Rebased onto gmlueck/SYCL-Docs@gmlueck/gentype-funcs.
    a101c3b: Removed constraints on the Space and IsDecorated template parameters for builtin functions taking a multi_ptr.
    323de84: Maintained builtin functions constraint on mutli_ptr parameters with constant_space and removed const qualifiers on global_ptr parameters to legacy async_work_group_copy overloads.
    99c0b3b: Added fixes based on feedback; fix for async_work_group_copy definition, removal of accessor_ptr for get_pointer() and additional get_pointer() declaration for the host_task case.

adoc/chapters/what_changed.adoc

@@ -119,10 +119,6 @@ Changes to buffers, images and accessors:
     [code]#accessor# have been changed to be [code]#const# types for read-only
     accessors.
 
-  * The [code]#accessor# member function [code]#get_pointer()# now returns
-    a raw pointer.  The [code]#get_multi_ptr()# member function was introduced
-    which returns the [code]#multi_ptr# class to the appropriate space.
-
   * The [code]#accessor# class now meets the {cpp} requirement of
     [code]#ReversibleContainer#.  This includes (but is not limited to)
     returning [code]#begin# and [code]#end# iterators, specifying a default

adoc/headers/accessorBuffer.h

@@ -184,6 +184,11 @@ class accessor {
   cl::sycl::atomic<DataT, access::address_space::global_space>
   operator[](id<Dimensions> index) const;
 
+  /* Deprecated in SYCL 2020
+  Available only when: (AccessTarget == target::device) */
+  global_ptr<DataT> get_pointer() const noexcept;
+
+  /* Available only when (AccessTarget == target::host_task) */
   std::add_pointer_t<value_type> get_pointer() const noexcept;
 
   template <access::decorated IsDecorated>

adoc/headers/accessorLocal.h

@@ -62,7 +62,8 @@ template <typename DataT, int Dimensions = 1> class local_accessor {
   /* Available only when: (Dimensions == 1) */
   reference operator[](size_t index) const;
 
-  std::add_pointer_t<value_type> get_pointer() const noexcept;
+  /* Deprecated in SYCL 2020 */
+  local_ptr<DataT> get_pointer() const noexcept;
 
   template <access::decorated IsDecorated>
   accessor_ptr<IsDecorated> get_multi_ptr() const noexcept;

adoc/headers/group.h

@@ -49,25 +49,59 @@ template <int Dimensions = 1> class group {
   void parallel_for_work_item(range<Dimensions> logicalRange,
                               const WorkItemFunctionT& func) const;
 
+  // Deprecated in SYCL 2020. 
   template <typename DataT>
-  device_event async_work_group_copy(decorated_local_ptr<DataT> dest,
-                                     decorated_global_ptr<const DataT> src,
+  device_event async_work_group_copy(local_ptr<DataT> dest,
+                                     global_ptr<DataT> src,
                                      size_t numElements) const;
 
+  // Deprecated in SYCL 2020.
   template <typename DataT>
-  device_event async_work_group_copy(decorated_global_ptr<DataT> dest,
-                                     decorated_local_ptr<const DataT> src,
+  device_event async_work_group_copy(global_ptr<DataT> dest,
+                                     local_ptr<DataT> src,
                                      size_t numElements) const;
 
+  // Deprecated in SYCL 2020.
   template <typename DataT>
-  device_event async_work_group_copy(decorated_local_ptr<DataT> dest,
-                                     decorated_global_ptr<const DataT> src,
+  device_event async_work_group_copy(local_ptr<DataT> dest,
+                                     global_ptr<DataT> src,
                                      size_t numElements,
                                      size_t srcStride) const;
 
+  // Deprecated in SYCL 2020.
   template <typename DataT>
-  device_event async_work_group_copy(decorated_global_ptr<DataT> dest,
-                                     decorated_local_ptr<const DataT> src,
+  device_event async_work_group_copy(global_ptr<DataT> dest,
+                                     local_ptr<DataT> src,
+                                     size_t numElements,
+                                     size_t destStride) const;
+
+  /* Available only when: (std::is_same_v<DestDataT,
+       std::remove_const_t<SrcDataT>> == true) */
+  template <typename DestDataT, typename SrcDataT>
+  device_event async_work_group_copy(decorated_local_ptr<DestDataT> dest,
+                                     decorated_global_ptr<SrcDataT> src,
+                                     size_t numElements) const;
+
+  /* Available only when: (std::is_same_v<DestDataT,
+       std::remove_const_t<SrcDataT>> == true) */
+  template <typename DestDataT, typename SrcDataT>
+  device_event async_work_group_copy(decorated_global_ptr<DestDataT> dest,
+                                     decorated_local_ptr<SrcDataT> src,
+                                     size_t numElements) const;
+
+  /* Available only when: (std::is_same_v<DestDataT,
+       std::remove_const_t<SrcDataT>> == true) */
+  template <typename DestDataT, typename SrcDataT>
+  device_event async_work_group_copy(decorated_local_ptr<DestDataT> dest,
+                                     decorated_global_ptr<SrcDataT> src,
+                                     size_t numElements,
+                                     size_t srcStride) const;
+
+  /* Available only when: (std::is_same_v<DestDataT,
+       std::remove_const_t<SrcDataT>> == true) */
+  template <typename DestDataT, typename SrcDataT>
+  device_event async_work_group_copy(decorated_global_ptr<DestDataT> dest,
+                                     decorated_local_ptr<SrcDataT> src,
                                      size_t numElements,
                                      size_t destStride) const;
 

adoc/headers/multipointer.h

@@ -12,7 +12,11 @@ enum class address_space : /* unspecified */ {
   generic_space
 };
 
-enum class decorated : /* unspecified */ { no, yes, legacy };
+enum class decorated : /* unspecified */ {
+  no,
+  yes,
+  legacy // Deprecated in SYCL 2020
+};
 
 } // namespace access
 
@@ -23,7 +27,7 @@ template <typename T> struct remove_decoration {
 template <typename T> using remove_decoration_t = remove_decoration<T>::type;
 
 template <typename ElementType, access::address_space Space,
-          access::decorated DecorateAddress>
+          access::decorated DecorateAddress = access::decorated::legacy>
 class multi_ptr {
  public:
   static constexpr bool is_decorated =

adoc/headers/multipointerlegacy.h

@@ -61,6 +61,10 @@ class [[deprecated]] multi_ptr<ElementType, Space, access::decorated::legacy> {
   // Returns the underlying OpenCL C pointer
   pointer_t get() const;
 
+  std::add_pointer_t<value_type> get_raw() const;
+
+  pointer_t get_decorated() const;
+
   // Implicit conversion to the underlying pointer type
   operator ElementType*() const;
 

adoc/headers/nditem.h

@@ -47,25 +47,59 @@ template <int Dimensions = 1> class nd_item {
 
   nd_range<Dimensions> get_nd_range() const;
 
+  // Deprecated in SYCL 2020. 
   template <typename DataT>
-  device_event async_work_group_copy(decorated_local_ptr<DataT> dest,
-                                     decorated_global_ptr<const DataT> src,
+  device_event async_work_group_copy(local_ptr<DataT> dest,
+                                     global_ptr<DataT> src,
                                      size_t numElements) const;
 
+  // Deprecated in SYCL 2020.
   template <typename DataT>
-  device_event async_work_group_copy(decorated_global_ptr<DataT> dest,
-                                     decorated_local_ptr<const DataT> src,
+  device_event async_work_group_copy(global_ptr<DataT> dest,
+                                     local_ptr<DataT> src,
                                      size_t numElements) const;
 
+  // Deprecated in SYCL 2020.
   template <typename DataT>
-  device_event async_work_group_copy(decorated_local_ptr<DataT> dest,
-                                     decorated_global_ptr<const DataT> src,
+  device_event async_work_group_copy(local_ptr<DataT> dest,
+                                     global_ptr<DataT> src,
                                      size_t numElements,
                                      size_t srcStride) const;
 
+  // Deprecated in SYCL 2020.
   template <typename DataT>
-  device_event async_work_group_copy(decorated_global_ptr<DataT> dest,
-                                     decorated_local_ptr<const DataT> src,
+  device_event async_work_group_copy(global_ptr<DataT> dest,
+                                     local_ptr<DataT> src,
+                                     size_t numElements,
+                                     size_t destStride) const;
+
+  /* Available only when: (std::is_same_v<DestDataT,
+       std::remove_const_t<SrcDataT>> == true) */
+  template <typename DestDataT, typename SrcDataT>
+  device_event async_work_group_copy(decorated_local_ptr<DestDataT> dest,
+                                     decorated_global_ptr<SrcDataT> src,
+                                     size_t numElements) const;
+
+  /* Available only when: (std::is_same_v<DestDataT,
+       std::remove_const_t<SrcDataT>> == true) */
+  template <typename DestDataT, typename SrcDataT>
+  device_event async_work_group_copy(decorated_global_ptr<DestDataT> dest,
+                                     decorated_local_ptr<SrcDataT> src,
+                                     size_t numElements) const;
+
+  /* Available only when: (std::is_same_v<DestDataT,
+       std::remove_const_t<SrcDataT>> == true) */
+  template <typename DestDataT, typename SrcDataT>
+  device_event async_work_group_copy(decorated_local_ptr<DestDataT> dest,
+                                     decorated_global_ptr<SrcDataT> src,
+                                     size_t numElements,
+                                     size_t srcStride) const;
+
+  /* Available only when: (std::is_same_v<DestDataT,
+       std::remove_const_t<SrcDataT>> == true) */
+  template <typename DestDataT, typename SrcDataT>
+  device_event async_work_group_copy(decorated_global_ptr<DestDataT> dest,
+                                     decorated_local_ptr<SrcDataT> src,
                                      size_t numElements,
                                      size_t destStride) const;