[SYCL] Reuse user ptr in buffer & fix default memory allocation (#243)

[SYCL] Reuse user ptr if buffer & fix default memory allocation

Reuse the pointer provided by the user in the buffer constructor (even
if use_host_ptr wasn't specified) if its alignment is sufficient.
Use the type-specific aligned_allocator<T> for buffer host memory
allocation instead of the non-template buffer_allocator (which is
aligned_allocator<char>).

Signed-off-by: Sergey Semenov <sergey.semenov@intel.com>

Author: sergey-semenov

sycl/include/CL/sycl/accessor.hpp

@@ -336,11 +336,11 @@ class accessor :
   using reference = DataT &;
   using const_reference = const DataT &;
 
-  template <int Dims = Dimensions>
+  template <typename AllocatorT, int Dims = Dimensions>
   accessor(
       enable_if_t<Dims == 0 && ((!IsPlaceH && IsHostBuf) ||
                                 (IsPlaceH && (IsGlobalBuf || IsConstantBuf))),
-                  buffer<DataT, 1>> &BufferRef)
+                  buffer<DataT, 1, AllocatorT>> &BufferRef)
 #ifdef __SYCL_DEVICE_ONLY__
       : impl(id<AdjustedDim>(), BufferRef.get_range(), BufferRef.MemRange) {
 #else
@@ -357,9 +357,9 @@ class accessor :
 #endif
   }
 
-  template <int Dims = Dimensions>
+  template <typename AllocatorT, int Dims = Dimensions>
   accessor(
-      buffer<DataT, 1> &BufferRef,
+      buffer<DataT, 1, AllocatorT> &BufferRef,
       enable_if_t<Dims == 0 && (!IsPlaceH && (IsGlobalBuf || IsConstantBuf)),
                   handler> &CommandGroupHandler)
 #ifdef __SYCL_DEVICE_ONLY__
@@ -376,11 +376,11 @@ class accessor :
   }
 #endif
 
-  template <int Dims = Dimensions,
+  template <typename AllocatorT, int Dims = Dimensions,
             typename = enable_if_t<
                 (Dims > 0) && ((!IsPlaceH && IsHostBuf) ||
                                (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>>
-  accessor(buffer<DataT, Dimensions> &BufferRef)
+  accessor(buffer<DataT, Dimensions, AllocatorT> &BufferRef)
 #ifdef __SYCL_DEVICE_ONLY__
       : impl(id<Dimensions>(), BufferRef.get_range(), BufferRef.MemRange) {
   }
@@ -398,10 +398,11 @@ class accessor :
   }
 #endif
 
-  template <int Dims = Dimensions,
+  template <typename AllocatorT, int Dims = Dimensions,
             typename = enable_if_t<
                 (Dims > 0) && (!IsPlaceH && (IsGlobalBuf || IsConstantBuf))>>
-  accessor(buffer<DataT, Dimensions> &BufferRef, handler &CommandGroupHandler)
+  accessor(buffer<DataT, Dimensions, AllocatorT> &BufferRef,
+           handler &CommandGroupHandler)
 #ifdef __SYCL_DEVICE_ONLY__
       : impl(id<AdjustedDim>(), BufferRef.get_range(), BufferRef.MemRange) {
   }
@@ -416,12 +417,12 @@ class accessor :
   }
 #endif
 
-  template <int Dims = Dimensions,
+  template <typename AllocatorT, int Dims = Dimensions,
             typename = enable_if_t<
                 (Dims > 0) && ((!IsPlaceH && IsHostBuf) ||
                                (IsPlaceH && (IsGlobalBuf || IsConstantBuf)))>>
-  accessor(buffer<DataT, Dimensions> &BufferRef, range<Dimensions> AccessRange,
-           id<Dimensions> AccessOffset = {})
+  accessor(buffer<DataT, Dimensions, AllocatorT> &BufferRef,
+           range<Dimensions> AccessRange, id<Dimensions> AccessOffset = {})
 #ifdef __SYCL_DEVICE_ONLY__
       : impl(AccessOffset, AccessRange, BufferRef.MemRange) {
   }
@@ -438,11 +439,12 @@ class accessor :
   }
 #endif
 
-  template <int Dims = Dimensions,
+  template <typename AllocatorT, int Dims = Dimensions,
             typename = enable_if_t<
                 (Dims > 0) && (!IsPlaceH && (IsGlobalBuf || IsConstantBuf))>>
-  accessor(buffer<DataT, Dimensions> &BufferRef, handler &CommandGroupHandler,
-           range<Dimensions> AccessRange, id<Dimensions> AccessOffset = {})
+  accessor(buffer<DataT, Dimensions, AllocatorT> &BufferRef,
+           handler &CommandGroupHandler, range<Dimensions> AccessRange,
+           id<Dimensions> AccessOffset = {})
 #ifdef __SYCL_DEVICE_ONLY__
       : impl(AccessOffset, AccessRange, BufferRef.MemRange) {
   }

sycl/include/CL/sycl/buffer.hpp

@@ -21,7 +21,7 @@ class queue;
 template <int dimensions> class range;
 
 template <typename T, int dimensions = 1,
-          typename AllocatorT = cl::sycl::buffer_allocator>
+          typename AllocatorT = cl::sycl::detail::aligned_allocator<T>>
 class buffer {
 public:
   using value_type = T;

sycl/include/CL/sycl/detail/aligned_allocator.hpp

@@ -10,18 +10,20 @@
 
 #include <CL/cl.h>
 #include <CL/sycl/detail/cnri.h>
+#include <CL/sycl/detail/common.hpp>
 #include <CL/sycl/detail/os_util.hpp>
 #include <CL/sycl/range.hpp>
 
+#include <algorithm>
 #include <cstring>
 #include <cstdlib>
 #include <memory>
 #include <vector>
 
 namespace cl {
 namespace sycl {
-template <typename T, size_t Alignment>
-class aligned_allocator {
+namespace detail {
+template <typename T> class aligned_allocator {
 public:
   using value_type = T;
   using pointer = T*;
@@ -30,10 +32,7 @@ class aligned_allocator {
   using const_reference = const T&;
 
 public:
-  template<typename U>
-  struct rebind {
-    typedef aligned_allocator<U, Alignment> other;
-  };
+  template <typename U> struct rebind { typedef aligned_allocator<U> other; };
 
   // Construct an object
   void construct(pointer Ptr, const_reference Val) {
@@ -46,11 +45,15 @@ class aligned_allocator {
   pointer address(reference Val) const { return &Val; }
   const_pointer address(const_reference Val) { return &Val; }
 
-  // Allocate aligned (to Alignment) memory
+  // Allocate sufficiently aligned memory
   pointer allocate(size_t Size) {
-    Size += Alignment - Size % Alignment;
+    size_t NumBytes = Size * sizeof(value_type);
+    const size_t Alignment =
+        std::max<size_t>(getNextPowerOfTwo(sizeof(value_type)), 64);
+    NumBytes = ((NumBytes - 1) | (Alignment - 1)) + 1;
+
     pointer Result = reinterpret_cast<pointer>(
-      detail::OSUtil::alignedAlloc(Alignment, Size * sizeof(value_type)));
+        detail::OSUtil::alignedAlloc(Alignment, NumBytes));
     if (!Result)
       throw std::bad_alloc();
     return Result;
@@ -65,5 +68,6 @@ class aligned_allocator {
   bool operator==(const aligned_allocator&) { return true; }
   bool operator!=(const aligned_allocator& rhs) { return false; }
 };
+} // namespace detail
 } // namespace sycl
 } // namespace cl

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

@@ -22,6 +22,7 @@
 #include <CL/sycl/stl.hpp>
 #include <CL/sycl/types.hpp>
 
+#include <cstdint>
 #include <functional>
 #include <memory>
 #include <type_traits>
@@ -35,7 +36,7 @@ class accessor;
 template <typename T, int Dimensions, typename AllocatorT> class buffer;
 class handler;
 
-using buffer_allocator = aligned_allocator<char, /*Alignment*/64>;
+using buffer_allocator = detail::aligned_allocator<char>;
 
 namespace detail {
 using EventImplPtr = std::shared_ptr<detail::event_impl>;
@@ -59,7 +60,10 @@ template <typename AllocatorT> class buffer_impl : public SYCLMemObjT {
       return;
 
     set_final_data(reinterpret_cast<char *>(HostData));
-    if (MProps.has_property<property::buffer::use_host_ptr>()) {
+    size_t RequiredAlignment =
+        getNextPowerOfTwo(sizeof(typename AllocatorT::value_type));
+    if (reinterpret_cast<std::uintptr_t>(HostData) % RequiredAlignment == 0 ||
+        MProps.has_property<property::buffer::use_host_ptr>()) {
       MUserPtr = HostData;
       return;
     }

sycl/include/CL/sycl/detail/common.hpp

@@ -103,6 +103,9 @@ template <class T> T createSyclObjFromImpl(decltype(T::impl) ImplObj) {
   return T(ImplObj);
 }
 
+// Returns the smallest power of two not less than Var
+size_t getNextPowerOfTwo(size_t Var);
+
 } // namespace detail
 } // namespace sycl
 } // namespace cl

sycl/include/CL/sycl/detail/image_impl.hpp

@@ -26,7 +26,7 @@ enum class image_channel_type : unsigned int;
 namespace detail {
 
 // utility functions and typedefs for image_impl
-using image_allocator = aligned_allocator<byte, /*alignment*/ 64>;
+using image_allocator = aligned_allocator<byte>;
 
 // utility function: Returns the Number of Channels for a given Order.
 uint8_t getImageNumberChannels(image_channel_order Order);

sycl/source/detail/common.cpp

@@ -237,6 +237,17 @@ vector_class<string_class> split_string(const string_class &str,
   return result;
 }
 
+size_t getNextPowerOfTwo(size_t Var) {
+  --Var;
+  Var |= Var >> 1;
+  Var |= Var >> 2;
+  Var |= Var >> 4;
+  Var |= Var >> 8;
+  Var |= Var >> 16;
+  Var |= Var >> 32;
+  return ++Var;
+}
+
 } // namespace detail
 } // namespace sycl
 } // namespace cl