[SYCL] Implement cl::sycl::buffer::reinterpret

Main changes (required to implement this method):
Removed template parameters T and dimensions from buffer_impl class.
Added buffer range to accessor. Currently it done to save handler::copy
functionality but it also need to be used in multidimensional access.
Moved range from buffer_impl to buffer.
Used buffer_allocator as default allocator in buffer.

Signed-off-by: Mariya Podchishchaeva <mariya.podchishchaeva@intel.com>

Author: Fznamznon

sycl/include/CL/sycl/accessor.hpp

@@ -125,11 +125,12 @@ SYCL_ACCESSOR_IMPL(isTargetHostAccess(accessTarget) && dimensions == 0) {
 SYCL_ACCESSOR_IMPL(isTargetHostAccess(accessTarget) && dimensions > 0) {
   dataT *Data;
   range<dimensions> Range;
+  range<dimensions> BufRange;
   id<dimensions> Offset;
 
-  accessor_impl(dataT *Data, range<dimensions> Range,
-                id<dimensions> Offset = {})
-      : Data(Data), Range(Range), Offset(Offset) {}
+  accessor_impl(dataT * Data, range<dimensions> Range,
+                range<dimensions> BufRange, id<dimensions> Offset = {})
+      : Data(Data), Range(Range), BufRange(BufRange), Offset(Offset) {}
 
   // Returns the number of accessed elements.
   size_t get_count() const { return Range.size(); }
@@ -146,10 +147,9 @@ SYCL_ACCESSOR_IMPL(!isTargetHostAccess(accessTarget) &&
   // reinterpret casting while setting kernel arguments in order to get cl_mem
   // value from the buffer regardless of the accessor's dimensionality.
 #ifndef __SYCL_DEVICE_ONLY__
-  detail::buffer_impl<dataT, 1> *m_Buf = nullptr;
-
+  detail::buffer_impl<buffer_allocator<char>> *m_Buf = nullptr;
 #else
-  char padding[sizeof(detail::buffer_impl<dataT, dimensions> *)];
+  char padding[sizeof(detail::buffer_impl<buffer_allocator<char>> *)];
 #endif // __SYCL_DEVICE_ONLY__
 
   dataT *Data;
@@ -182,22 +182,23 @@ SYCL_ACCESSOR_IMPL(!isTargetHostAccess(accessTarget) &&
   // reinterpret casting while setting kernel arguments in order to get cl_mem
   // value from the buffer regardless of the accessor's dimensionality.
 #ifndef __SYCL_DEVICE_ONLY__
-  detail::buffer_impl<dataT, dimensions> *m_Buf = nullptr;
+  detail::buffer_impl<buffer_allocator<char>> *m_Buf = nullptr;
 #else
-  char padding[sizeof(detail::buffer_impl<dataT, dimensions> *)];
+  char padding[sizeof(detail::buffer_impl<buffer_allocator<char>> *)];
 #endif // __SYCL_DEVICE_ONLY__
 
   dataT *Data;
   range<dimensions> Range;
+  range<dimensions> BufRange;
   id<dimensions> Offset;
 
   // Device accessors must be associated with a command group handler.
   // The handler though can be nullptr at the creation point if the
   // accessor is a placeholder accessor.
-  accessor_impl(dataT *Data, range<dimensions> Range,
-                handler *Handler = nullptr, id<dimensions> Offset = {})
-      : Data(Data), Range(Range), Offset(Offset)
-  {}
+  accessor_impl(dataT * Data, range<dimensions> Range,
+                range<dimensions> BufRange, handler *Handler = nullptr,
+                id<dimensions> Offset = {})
+      : Data(Data), Range(Range), BufRange(BufRange), Offset(Offset) {}
 
   // Returns the number of accessed elements.
   size_t get_count() const { return Range.size(); }
@@ -633,8 +634,8 @@ class accessor
 #ifdef __SYCL_DEVICE_ONLY__
       ; // This ctor can't be used in device code, so no need to define it.
 #else // !__SYCL_DEVICE_ONLY__
-      : __impl(detail::getSyclObjImpl(bufferRef)->BufPtr,
-               detail::getSyclObjImpl(bufferRef)->Range,
+      : __impl((dataT *)detail::getSyclObjImpl(bufferRef)->BufPtr,
+               bufferRef.get_range(), bufferRef.get_range(),
                &commandGroupHandlerRef) {
     auto BufImpl = detail::getSyclObjImpl(bufferRef);
     if (BufImpl->OpenCLInterop && !BufImpl->isValidAccessToMem(accessMode)) {
@@ -669,8 +670,8 @@ class accessor
                AccessTarget == access::target::constant_buffer))) &&
             Dimensions > 0),
            buffer<DataT, Dimensions>>::type &bufferRef)
-      : __impl(detail::getSyclObjImpl(bufferRef)->BufPtr,
-               detail::getSyclObjImpl(bufferRef)->Range) {
+      : __impl((dataT *)detail::getSyclObjImpl(bufferRef)->BufPtr,
+               bufferRef.get_range(), bufferRef.get_range()) {
     auto BufImpl = detail::getSyclObjImpl(bufferRef);
     if (AccessTarget == access::target::host_buffer) {
       if (BufImpl->OpenCLInterop) {
@@ -701,17 +702,17 @@ class accessor
             access::target AccessTarget = accessTarget,
             access::placeholder IsPlaceholder = isPlaceholder>
   accessor(typename std::enable_if<
-           (IsPlaceholder == access::placeholder::false_t &&
-            (AccessTarget == access::target::global_buffer ||
-             AccessTarget == access::target::constant_buffer) &&
-            Dimensions > 0),
-           buffer<DataT, Dimensions>>::type &bufferRef,
+               (IsPlaceholder == access::placeholder::false_t &&
+                (AccessTarget == access::target::global_buffer ||
+                 AccessTarget == access::target::constant_buffer) &&
+                Dimensions > 0),
+               buffer<DataT, Dimensions>>::type &bufferRef,
            handler &commandGroupHandlerRef)
 #ifdef __SYCL_DEVICE_ONLY__
       ; // This ctor can't be used in device code, so no need to define it.
 #else
-      : __impl(detail::getSyclObjImpl(bufferRef)->BufPtr,
-               detail::getSyclObjImpl(bufferRef)->Range,
+      : __impl((dataT *)detail::getSyclObjImpl(bufferRef)->BufPtr,
+               bufferRef.get_range(), bufferRef.get_range(),
                &commandGroupHandlerRef) {
     auto BufImpl = detail::getSyclObjImpl(bufferRef);
     if (BufImpl->OpenCLInterop && !BufImpl->isValidAccessToMem(accessMode)) {
@@ -739,20 +740,19 @@ class accessor
             access::target AccessTarget = accessTarget,
             access::placeholder IsPlaceholder = isPlaceholder>
   accessor(typename std::enable_if<
-           ((IsPlaceholder == access::placeholder::false_t &&
-             AccessTarget == access::target::host_buffer) ||
-            (IsPlaceholder == access::placeholder::true_t &&
-             (AccessTarget == access::target::global_buffer ||
-              AccessTarget == access::target::constant_buffer) &&
-             Dimensions > 0)),
-           buffer<DataT, Dimensions>>::type &bufferRef,
-           range<Dimensions> Range,
-           id<Dimensions> Offset = {}
-          )
+               ((IsPlaceholder == access::placeholder::false_t &&
+                 AccessTarget == access::target::host_buffer) ||
+                (IsPlaceholder == access::placeholder::true_t &&
+                 (AccessTarget == access::target::global_buffer ||
+                  AccessTarget == access::target::constant_buffer) &&
+                 Dimensions > 0)),
+               buffer<DataT, Dimensions>>::type &bufferRef,
+           range<Dimensions> Range, id<Dimensions> Offset = {})
 #ifdef __SYCL_DEVICE_ONLY__
       ; // This ctor can't be used in device code, so no need to define it.
-#else // !__SYCL_DEVICE_ONLY__
-      : __impl(detail::getSyclObjImpl(bufferRef)->BufPtr, Range, Offset) {
+#else   // !__SYCL_DEVICE_ONLY__
+      : __impl(detail::getSyclObjImpl(bufferRef)->BufPtr, Range,
+               bufferRef.get_range(), Offset) {
     auto BufImpl = detail::getSyclObjImpl(bufferRef);
     if (AccessTarget == access::target::host_buffer) {
       if (BufImpl->OpenCLInterop) {
@@ -769,7 +769,7 @@ class accessor
           "interoperability buffer");
     }
   }
-#endif // !__SYCL_DEVICE_ONLY__
+#endif  // !__SYCL_DEVICE_ONLY__
 
   // buffer ctor #6:
   //   accessor(buffer &, handler &, range Range, id Offset);
@@ -784,20 +784,18 @@ class accessor
             access::target AccessTarget = accessTarget,
             access::placeholder IsPlaceholder = isPlaceholder>
   accessor(typename std::enable_if<
-           (IsPlaceholder == access::placeholder::false_t &&
-            (AccessTarget == access::target::global_buffer ||
-             AccessTarget == access::target::constant_buffer) &&
-            Dimensions > 0),
-           buffer<DataT, Dimensions>>::type &bufferRef,
-           handler &commandGroupHandlerRef,
-           range<Dimensions> Range,
-           id<Dimensions> Offset = {}
-          )
+               (IsPlaceholder == access::placeholder::false_t &&
+                (AccessTarget == access::target::global_buffer ||
+                 AccessTarget == access::target::constant_buffer) &&
+                Dimensions > 0),
+               buffer<DataT, Dimensions>>::type &bufferRef,
+           handler &commandGroupHandlerRef, range<Dimensions> Range,
+           id<Dimensions> Offset = {})
 #ifdef __SYCL_DEVICE_ONLY__
       ; // This ctor can't be used in device code, so no need to define it.
-#else // !__SYCL_DEVICE_ONLY__
-      : __impl(detail::getSyclObjImpl(bufferRef)->BufPtr, Range,
-               &commandGroupHandlerRef, Offset) {
+#else   // !__SYCL_DEVICE_ONLY__
+      : __impl((dataT *)detail::getSyclObjImpl(bufferRef)->BufPtr, Range,
+               bufferRef.get_range(), &commandGroupHandlerRef, Offset) {
     auto BufImpl = detail::getSyclObjImpl(bufferRef);
     if (BufImpl->OpenCLInterop && !BufImpl->isValidAccessToMem(accessMode)) {
       throw cl::sycl::runtime_error(
@@ -807,7 +805,7 @@ class accessor
     commandGroupHandlerRef.AddBufDep<AccessMode, AccessTarget>(*BufImpl);
     __impl.m_Buf = BufImpl.get();
   }
-#endif // !__SYCL_DEVICE_ONLY__
+#endif  // !__SYCL_DEVICE_ONLY__
 
   // TODO:
   // local accessor ctor #1

sycl/include/CL/sycl/buffer.hpp

@@ -21,7 +21,7 @@ class queue;
 template <int dimentions> class range;
 
 template <typename T, int dimensions = 1,
-          typename AllocatorT = cl::sycl::buffer_allocator<T>>
+          typename AllocatorT = cl::sycl::buffer_allocator<char>>
 class buffer {
 public:
   using value_type = T;
@@ -30,9 +30,10 @@ class buffer {
   using allocator_type = AllocatorT;
 
   buffer(const range<dimensions> &bufferRange,
-         const property_list &propList = {}) {
-    impl = std::make_shared<detail::buffer_impl<T, dimensions, AllocatorT>>(
-        bufferRange, propList);
+         const property_list &propList = {})
+      : Range(bufferRange) {
+    impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
+        get_count() * sizeof(T), propList);
   }
 
   // buffer(const range<dimensions> &bufferRange, AllocatorT allocator,
@@ -42,9 +43,10 @@ class buffer {
   // }
 
   buffer(T *hostData, const range<dimensions> &bufferRange,
-         const property_list &propList = {}) {
-    impl = std::make_shared<detail::buffer_impl<T, dimensions, AllocatorT>>(
-        hostData, bufferRange, propList);
+         const property_list &propList = {})
+      : Range(bufferRange) {
+    impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
+        hostData, get_count() * sizeof(T), propList);
   }
 
   // buffer(T *hostData, const range<dimensions> &bufferRange,
@@ -54,9 +56,10 @@ class buffer {
   // }
 
   buffer(const T *hostData, const range<dimensions> &bufferRange,
-         const property_list &propList = {}) {
-    impl = std::make_shared<detail::buffer_impl<T, dimensions, AllocatorT>>(
-        hostData, bufferRange, propList);
+         const property_list &propList = {})
+      : Range(bufferRange) {
+    impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
+        hostData, get_count() * sizeof(T), propList);
   }
 
   // buffer(const T *hostData, const range<dimensions> &bufferRange,
@@ -74,9 +77,10 @@ class buffer {
 
   buffer(const shared_ptr_class<T> &hostData,
          const range<dimensions> &bufferRange,
-         const property_list &propList = {}) {
-    impl = std::make_shared<detail::buffer_impl<T, dimensions, AllocatorT>>(
-        hostData, bufferRange, propList);
+         const property_list &propList = {})
+      : Range(bufferRange) {
+    impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
+        hostData, get_count() * sizeof(T), propList);
   }
 
   // template <class InputIterator>
@@ -89,9 +93,10 @@ class buffer {
   template <class InputIterator, int N = dimensions,
             typename = std::enable_if<N == 1>>
   buffer(InputIterator first, InputIterator last,
-         const property_list &propList = {}) {
-    impl = std::make_shared<detail::buffer_impl<T, dimensions, AllocatorT>>(
-        first, last, propList);
+         const property_list &propList = {})
+      : Range(range<1>(std::distance(first, last))) {
+    impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
+        first, last, get_count() * sizeof(T), propList);
   }
 
   // buffer(buffer<T, dimensions, AllocatorT> b, const id<dimensions>
@@ -102,7 +107,7 @@ class buffer {
   template <int N = dimensions, typename = std::enable_if<N == 1>>
   buffer(cl_mem MemObject, const context &SyclContext,
          event AvailableEvent = {}) {
-    impl = std::make_shared<detail::buffer_impl<T, dimensions, AllocatorT>>(
+    impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
         MemObject, SyclContext, AvailableEvent);
   }
 
@@ -124,26 +129,27 @@ class buffer {
 
   /* -- property interface members -- */
 
-  range<dimensions> get_range() const { return impl->get_range(); }
+  range<dimensions> get_range() const { return Range; }
 
-  size_t get_count() const { return impl->get_count(); }
+  size_t get_count() const { return Range.size(); }
 
   size_t get_size() const { return impl->get_size(); }
 
-  AllocatorT get_allocator() const { return impl->get_allocator(); }
+  // AllocatorT get_allocator() const { return impl->get_allocator(); }
 
   template <access::mode mode,
             access::target target = access::target::global_buffer>
   accessor<T, dimensions, mode, target, access::placeholder::false_t>
   get_access(handler &commandGroupHandler) {
-    return impl->template get_access<mode, target>(*this, commandGroupHandler);
+    return impl->template get_access<T, dimensions, mode, target>(
+        *this, commandGroupHandler);
   }
 
   template <access::mode mode>
   accessor<T, dimensions, mode, access::target::host_buffer,
            access::placeholder::false_t>
   get_access() {
-    return impl->template get_access<mode>(*this);
+    return impl->template get_access<T, dimensions, mode>(*this);
   }
 
   // template <access::mode mode, access::target target =
@@ -171,16 +177,29 @@ class buffer {
 
   // bool is_sub_buffer() const { return impl->is_sub_buffer(); }
 
-  // template <typename ReinterpretT, int ReinterpretDim>
-  // buffer<ReinterpretT, ReinterpretDim, AllocatorT>
-  // reinterpret(range<ReinterpretDim> reinterpretRange) const {
-  //     return impl->reinterpret((reinterpretRange));
-  // }
+  template <typename ReinterpretT, int ReinterpretDim>
+  buffer<ReinterpretT, ReinterpretDim, AllocatorT>
+  reinterpret(range<ReinterpretDim> reinterpretRange) const {
+    if (sizeof(ReinterpretT) * reinterpretRange.size() != get_size())
+      throw cl::sycl::invalid_object_error(
+          "Total size in bytes represented by the type and range of the "
+          "reinterpreted SYCL buffer does not equal the total size in bytes "
+          "represented by the type and range of this SYCL buffer");
+    return buffer<ReinterpretT, ReinterpretDim, AllocatorT>(impl,
+                                                            reinterpretRange);
+  }
 
 private:
-  shared_ptr_class<detail::buffer_impl<T, dimensions, AllocatorT>> impl;
+  shared_ptr_class<detail::buffer_impl<AllocatorT>> impl;
   template <class Obj>
   friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
+  template <typename A, int dims, typename C> friend class buffer;
+  range<dimensions> Range;
+
+  // Reinterpret contructor
+  buffer(shared_ptr_class<detail::buffer_impl<AllocatorT>> Impl,
+         range<dimensions> reinterpretRange)
+      : impl(Impl), Range(reinterpretRange){};
 };
 } // namespace sycl
 } // namespace cl
@@ -190,8 +209,7 @@ template <typename T, int dimensions, typename AllocatorT>
 struct hash<cl::sycl::buffer<T, dimensions, AllocatorT>> {
   size_t
   operator()(const cl::sycl::buffer<T, dimensions, AllocatorT> &b) const {
-    return hash<std::shared_ptr<
-        cl::sycl::detail::buffer_impl<T, dimensions, AllocatorT>>>()(
+    return hash<std::shared_ptr<cl::sycl::detail::buffer_impl<AllocatorT>>>()(
         cl::sycl::detail::getSyclObjImpl(b));
   }
 };