Revert "[SYCL] Reuse user ptr in buffer & fix default memory allocati…

sycl/include/CL/sycl/accessor.hpp

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

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

@@ -10,20 +10,18 @@
 
 #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 {
-namespace detail {
-template <typename T> class aligned_allocator {
+template <typename T, size_t Alignment>
+class aligned_allocator {
 public:
   using value_type = T;
   using pointer = T*;
@@ -32,7 +30,10 @@ template <typename T> class aligned_allocator {
   using const_reference = const T&;
 
 public:
-  template <typename U> struct rebind { typedef aligned_allocator<U> other; };
+  template<typename U>
+  struct rebind {
+    typedef aligned_allocator<U, Alignment> other;
+  };
 
   // Construct an object
   void construct(pointer Ptr, const_reference Val) {
@@ -45,15 +46,11 @@ template <typename T> class aligned_allocator {
   pointer address(reference Val) const { return &Val; }
   const_pointer address(const_reference Val) { return &Val; }
 
-  // Allocate sufficiently aligned memory
+  // Allocate aligned (to Alignment) memory
   pointer allocate(size_t Size) {
-    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;
-
+    Size += Alignment - Size % Alignment;
     pointer Result = reinterpret_cast<pointer>(
-        detail::OSUtil::alignedAlloc(Alignment, NumBytes));
+      detail::OSUtil::alignedAlloc(Alignment, Size * sizeof(value_type)));
     if (!Result)
       throw std::bad_alloc();
     return Result;
@@ -68,6 +65,5 @@ template <typename T> 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,7 +22,6 @@
 #include <CL/sycl/stl.hpp>
 #include <CL/sycl/types.hpp>
 
-#include <cstdint>
 #include <functional>
 #include <memory>
 #include <type_traits>
@@ -36,7 +35,7 @@ class accessor;
 template <typename T, int Dimensions, typename AllocatorT> class buffer;
 class handler;
 
-using buffer_allocator = detail::aligned_allocator<char>;
+using buffer_allocator = aligned_allocator<char, /*Alignment*/64>;
 
 namespace detail {
 using EventImplPtr = std::shared_ptr<detail::event_impl>;
@@ -60,10 +59,7 @@ template <typename AllocatorT> class buffer_impl : public SYCLMemObjT {
       return;
 
     set_final_data(reinterpret_cast<char *>(HostData));
-    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>()) {
+    if (MProps.has_property<property::buffer::use_host_ptr>()) {
       MUserPtr = HostData;
       return;
     }

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

@@ -103,9 +103,6 @@ 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>;
+using image_allocator = aligned_allocator<byte, /*alignment*/ 64>;
 
 // utility function: Returns the Number of Channels for a given Order.
 uint8_t getImageNumberChannels(image_channel_order Order);

sycl/source/detail/common.cpp

@@ -237,17 +237,6 @@ 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