[SYCL] Add support for custom allocators in buffer

Added missed constructors and get_allocator method.

Signed-off-by: Mariya Podchishchaeva <[email protected]>

Author: Fznamznon

sycl/include/CL/sycl/accessor.hpp

@@ -150,9 +150,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<buffer_allocator<char>> *m_Buf = nullptr;
+  detail::buffer_impl<buffer_allocator> *m_Buf = nullptr;
 #else
-  char padding[sizeof(detail::buffer_impl<buffer_allocator<char>> *)];
+  char padding[sizeof(detail::buffer_impl<buffer_allocator> *)];
 #endif // __SYCL_DEVICE_ONLY__
 
   dataT *Data;
@@ -185,9 +185,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<buffer_allocator<char>> *m_Buf = nullptr;
+  detail::buffer_impl<buffer_allocator> *m_Buf = nullptr;
 #else
-  char padding[sizeof(detail::buffer_impl<buffer_allocator<char>> *)];
+  char padding[sizeof(detail::buffer_impl<buffer_allocator> *)];
 #endif // __SYCL_DEVICE_ONLY__
 
   dataT *Data;

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<char>>
+          typename AllocatorT = cl::sycl::buffer_allocator>
 class buffer {
 public:
   using value_type = T;
@@ -36,11 +36,11 @@ class buffer {
         get_count() * sizeof(T), propList);
   }
 
-  // buffer(const range<dimensions> &bufferRange, AllocatorT allocator,
-  // const property_list &propList = {}) {
-  //     impl = std::make_shared<detail::buffer_impl>(bufferRange, allocator,
-  //     propList);
-  // }
+  buffer(const range<dimensions> &bufferRange, AllocatorT allocator,
+         const property_list &propList = {}) {
+    impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
+        get_count() * sizeof(T), propList, allocator);
+  }
 
   buffer(T *hostData, const range<dimensions> &bufferRange,
          const property_list &propList = {})
@@ -49,11 +49,11 @@ class buffer {
         hostData, get_count() * sizeof(T), propList);
   }
 
-  // buffer(T *hostData, const range<dimensions> &bufferRange,
-  // AllocatorT allocator, const property_list &propList = {}) {
-  //     impl = std::make_shared<detail::buffer_impl>(hostData, bufferRange,
-  //     allocator, propList);
-  // }
+  buffer(T *hostData, const range<dimensions> &bufferRange,
+         AllocatorT allocator, const property_list &propList = {}) {
+    impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
+        hostData, get_count() * sizeof(T), propList, allocator);
+  }
 
   buffer(const T *hostData, const range<dimensions> &bufferRange,
          const property_list &propList = {})
@@ -62,18 +62,18 @@ class buffer {
         hostData, get_count() * sizeof(T), propList);
   }
 
-  // buffer(const T *hostData, const range<dimensions> &bufferRange,
-  // AllocatorT allocator, const property_list &propList = {}) {
-  //     impl = std::make_shared<detail::buffer_impl>(hostData, bufferRange,
-  //     allocator, propList);
-  // }
+  buffer(const T *hostData, const range<dimensions> &bufferRange,
+         AllocatorT allocator, const property_list &propList = {}) {
+    impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
+        hostData, get_count() * sizeof(T), propList, allocator);
+  }
 
-  // buffer(const shared_ptr_class<T> &hostData,
-  // const range<dimensions> &bufferRange, AllocatorT allocator,
-  // const property_list &propList = {}) {
-  //     impl = std::make_shared<detail::buffer_impl>(hostData, bufferRange,
-  //     allocator, propList);
-  // }
+  buffer(const shared_ptr_class<T> &hostData,
+         const range<dimensions> &bufferRange, AllocatorT allocator,
+         const property_list &propList = {}) {
+    impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
+        hostData, get_count() * sizeof(T), propList, allocator);
+  }
 
   buffer(const shared_ptr_class<T> &hostData,
          const range<dimensions> &bufferRange,
@@ -83,12 +83,13 @@ class buffer {
         hostData, get_count() * sizeof(T), propList);
   }
 
-  // template <class InputIterator>
-  // buffer<T, 1>(InputIterator first, InputIterator last, AllocatorT allocator,
-  // const property_list &propList = {}) {
-  //     impl = std::make_shared<detail::buffer_impl>(first, last, allocator,
-  //     propList);
-  // }
+  template <class InputIterator>
+  buffer(InputIterator first, InputIterator last, AllocatorT allocator,
+         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, allocator);
+  }
 
   template <class InputIterator, int N = dimensions,
             typename = std::enable_if<N == 1>>
@@ -135,7 +136,7 @@ class buffer {
 
   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>

sycl/include/CL/sycl/detail/buffer_impl.hpp

@@ -37,16 +37,18 @@ class handler;
 class queue;
 template <int dimentions> class id;
 template <int dimentions> class range;
-template <class T> using buffer_allocator = std::allocator<T>;
+using buffer_allocator = std::allocator<char>;
 namespace detail {
 template <typename AllocatorT> class buffer_impl {
 public:
-  buffer_impl(const size_t sizeInBytes, const property_list &propList)
-      : buffer_impl((void *)nullptr, sizeInBytes, propList) {}
+  buffer_impl(const size_t sizeInBytes, const property_list &propList,
+              AllocatorT allocator = AllocatorT())
+      : buffer_impl((void *)nullptr, sizeInBytes, propList, allocator) {}
 
   buffer_impl(void *hostData, const size_t sizeInBytes,
-              const property_list &propList)
-      : SizeInBytes(sizeInBytes), Props(propList) {
+              const property_list &propList,
+              AllocatorT allocator = AllocatorT())
+      : SizeInBytes(sizeInBytes), Props(propList), MAllocator(allocator) {
     if (Props.has_property<property::buffer::use_host_ptr>()) {
       BufPtr = hostData;
     } else {
@@ -62,8 +64,9 @@ template <typename AllocatorT> class buffer_impl {
 
   // TODO temporary solution for allowing initialisation with const data
   buffer_impl(const void *hostData, const size_t sizeInBytes,
-              const property_list &propList)
-      : SizeInBytes(sizeInBytes), Props(propList) {
+              const property_list &propList,
+              AllocatorT allocator = AllocatorT())
+      : SizeInBytes(sizeInBytes), Props(propList), MAllocator(allocator) {
     if (Props.has_property<property::buffer::use_host_ptr>()) {
       // TODO make this buffer read only
       BufPtr = const_cast<void *>(hostData);
@@ -79,8 +82,9 @@ template <typename AllocatorT> class buffer_impl {
 
   template <typename T>
   buffer_impl(const shared_ptr_class<T> &hostData, const size_t sizeInBytes,
-              const property_list &propList)
-      : SizeInBytes(sizeInBytes), Props(propList) {
+              const property_list &propList,
+              AllocatorT allocator = AllocatorT())
+      : SizeInBytes(sizeInBytes), Props(propList), MAllocator(allocator) {
     if (Props.has_property<property::buffer::use_host_ptr>()) {
       BufPtr = hostData.get();
     } else {
@@ -97,8 +101,9 @@ template <typename AllocatorT> class buffer_impl {
 
   template <class InputIterator>
   buffer_impl(InputIterator first, InputIterator last, const size_t sizeInBytes,
-              const property_list &propList)
-      : SizeInBytes(sizeInBytes), Props(propList) {
+              const property_list &propList,
+              AllocatorT allocator = AllocatorT())
+      : SizeInBytes(sizeInBytes), Props(propList), MAllocator(allocator) {
     if (Props.has_property<property::buffer::use_host_ptr>()) {
       // TODO next line looks unsafe
       BufPtr = &*first;
@@ -170,7 +175,7 @@ template <typename AllocatorT> class buffer_impl {
       throw cl::sycl::runtime_error(
           "set_final_data could not be used with interoperability buffer");
     static_assert(!std::is_const<Destination>::value,
-                  "Сan not write in a constant Destination. Destination should "
+                  "Can not write in a constant Destination. Destination should "
                   "not be const.");
     uploadData = [this, final_data]() mutable {
       auto *Ptr =
@@ -182,6 +187,8 @@ template <typename AllocatorT> class buffer_impl {
     };
   }
 
+  AllocatorT get_allocator() const { return MAllocator; }
+
   template <typename T, int dimensions, access::mode mode,
             access::target target = access::target::global_buffer>
   accessor<T, dimensions, mode, target, access::placeholder::false_t>
@@ -243,6 +250,7 @@ template <typename AllocatorT> class buffer_impl {
   // This field must be the first to guarantee that it's safe to use
   // reinterpret casting while setting kernel arguments in order to get cl_mem
   // value from the buffer regardless of its dimensionality.
+  AllocatorT MAllocator;
   OpenCLMemState OCLState;
   bool OpenCLInterop = false;
   event AvailableEvent;

sycl/include/CL/sycl/detail/scheduler/scheduler.cpp

@@ -61,8 +61,7 @@ template <typename dataT, int dimensions, access::mode accessMode,
 void Node::addAccRequirement(
     accessor<dataT, dimensions, accessMode, accessTarget, isPlaceholder> &&Acc,
     int argIndex) {
-  detail::buffer_impl<buffer_allocator<char>> *buf =
-      Acc.__get_impl()->m_Buf;
+  detail::buffer_impl<buffer_allocator> *buf = Acc.__get_impl()->m_Buf;
   addBufRequirement<accessMode, accessTarget>(*buf);
   addInteropArg(nullptr, buf->get_size(), argIndex,
                 getReqForBuffer(m_Bufs, *buf));
@@ -134,7 +133,7 @@ void Node::addExplicitMemOp(
   auto *DestBase = Dest.__get_impl();
   assert(DestBase != nullptr &&
          "Accessor should have an initialized accessor_base");
-  detail::buffer_impl<buffer_allocator<char>> *Buf = DestBase->m_Buf;
+  detail::buffer_impl<buffer_allocator> *Buf = DestBase->m_Buf;
 
   range<Dimensions> Range = DestBase->AccessRange;
   id<Dimensions> Offset = DestBase->Offset;
@@ -162,10 +161,10 @@ void Node::addExplicitMemOp(
   assert(DestBase != nullptr &&
          "Accessor should have an initialized accessor_base");
 
-  detail::buffer_impl<buffer_allocator<char>> *SrcBuf = SrcBase->m_Buf;
+  detail::buffer_impl<buffer_allocator> *SrcBuf = SrcBase->m_Buf;
   assert(SrcBuf != nullptr &&
          "Accessor should have an initialized buffer_impl");
-  detail::buffer_impl<buffer_allocator<char>> *DestBuf = DestBase->m_Buf;
+  detail::buffer_impl<buffer_allocator> *DestBuf = DestBase->m_Buf;
   assert(DestBuf != nullptr &&
          "Accessor should have an initialized buffer_impl");
 
@@ -195,7 +194,7 @@ void Scheduler::updateHost(
   auto *AccBase = Acc.__get_impl();
   assert(AccBase != nullptr &&
          "Accessor should have an initialized accessor_base");
-  detail::buffer_impl<buffer_allocator<char>> *Buf = AccBase->m_Buf;
+  detail::buffer_impl<buffer_allocator> *Buf = AccBase->m_Buf;
 
   updateHost<mode, tgt>(*Buf, Event);
 }