[SYCL] Implement and use aligned allocator

MrSidims · romanovvlad · commit f2070627a361 · 2019-04-19T16:29:28.000+03:00
Alignment is of 64 bytes by default.
Buffer is not longer zero-initialized.

Signed-off-by: Dmitry Sidorov &lt;dmitry.sidorov@intel.com&gt;
diff --git a/sycl/include/CL/sycl/detail/aligned_allocator.hpp b/sycl/include/CL/sycl/detail/aligned_allocator.hpp
@@ -0,0 +1,68 @@
+//==------------ aligned_allocator.hpp - SYCL standard header file ---------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <CL/cl.h>
+#include <CL/sycl/detail/cnri.h>
+#include <CL/sycl/range.hpp>
+
+#include <cstring>
+#include <cstdlib>
+#include <memory>
+#include <vector>
+
+namespace cl {
+namespace sycl {
+template <typename T, size_t Alignment>
+class aligned_allocator {
+public:
+  using value_type = T;
+  using pointer = T*;
+  using const_pointer = const T*;
+  using reference = T&;
+  using const_reference = const T&;
+
+public:
+  template<typename U>
+  struct rebind {
+    typedef aligned_allocator<U, Alignment> other;
+  };
+
+  // Construct an object
+  void construct(pointer Ptr, const_reference Val) {
+    new (Ptr) value_type(Val);
+  }
+
+  // Destroy an object
+  void destroy(pointer Ptr) { Ptr->~value_type(); }
+
+  pointer address(reference Val) const { return &Val; }
+  const_pointer address(const_reference Val) { return &Val; }
+
+  // Allocate aligned (to Alignment) memory
+  pointer allocate(size_t Size) {
+    Size += Alignment - Size % Alignment;
+    pointer Result = reinterpret_cast<pointer>(
+      aligned_alloc(Alignment, Size * sizeof(value_type)));
+    if (!Result)
+      throw std::bad_alloc();
+    return Result;
+  }
+
+  // Release allocated memory
+  void deallocate(pointer Ptr, size_t size) {
+    if (Ptr)
+      free(Ptr);
+  }
+
+  bool operator==(const aligned_allocator&) { return true; }
+  bool operator!=(const aligned_allocator& rhs) { return false; }
+};
+} // namespace sycl
+} // namespace cl
diff --git a/sycl/include/CL/sycl/detail/buffer_impl.hpp b/sycl/include/CL/sycl/detail/buffer_impl.hpp
@@ -12,6 +12,7 @@
 #include <CL/sycl/context.hpp>
 #include <CL/sycl/detail/common.hpp>
 #include <CL/sycl/detail/helpers.hpp>
+#include <CL/sycl/detail/aligned_allocator.hpp>
 #include <CL/sycl/detail/queue_impl.hpp>
 #include <CL/sycl/detail/scheduler/scheduler.h>
 #include <CL/sycl/handler.hpp>
@@ -37,7 +38,7 @@ class handler;
 class queue;
 template <int dimentions> class id;
 template <int dimentions> class range;
-using buffer_allocator = std::allocator<char>;
+using buffer_allocator = aligned_allocator<char, /*alignment*/ 64>;
 namespace detail {
 template <typename AllocatorT> class buffer_impl {
 public:
@@ -48,16 +49,15 @@ template <typename AllocatorT> class buffer_impl {
   buffer_impl(void *hostData, const size_t sizeInBytes,
               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 {
-      BufData.resize(get_size());
-      BufPtr = reinterpret_cast<void *>(BufData.data());
-      if (hostData != nullptr) {
-        auto HostPtr = reinterpret_cast<char *>(hostData);
-        set_final_data(HostPtr);
-        std::copy(HostPtr, HostPtr + SizeInBytes, BufData.data());
+      : SizeInBytes(sizeInBytes), Props(propList), BufPtr(hostData),
+        MAllocator(allocator) {
+    if (!Props.has_property<property::buffer::use_host_ptr>()) {
+      BufPtr = allocateHostMem();
+      if (hostData) {
+        set_final_data(reinterpret_cast<char *>(hostData));
+        std::copy(reinterpret_cast<char *>(hostData),
+                  reinterpret_cast<char *>(hostData) + SizeInBytes,
+                  reinterpret_cast<char *>(BufPtr));
       }
     }
   }
@@ -66,35 +66,33 @@ template <typename AllocatorT> class buffer_impl {
   buffer_impl(const void *hostData, const size_t sizeInBytes,
               const property_list &propList,
               AllocatorT allocator = AllocatorT())
-      : SizeInBytes(sizeInBytes), Props(propList), MAllocator(allocator) {
-    if (Props.has_property<property::buffer::use_host_ptr>()) {
+      : SizeInBytes(sizeInBytes), Props(propList),
+        BufPtr(const_cast<void *>(hostData)), MAllocator(allocator) {
+    if (!Props.has_property<property::buffer::use_host_ptr>()) {
       // TODO make this buffer read only
-      BufPtr = const_cast<void *>(hostData);
-    } else {
-      BufData.resize(get_size());
-      BufPtr = reinterpret_cast<void *>(BufData.data());
-      if (hostData != nullptr) {
-        std::copy((char *)hostData, (char *)hostData + SizeInBytes,
-                  BufData.data());
-      }
+      BufPtr = allocateHostMem();
+      if (hostData)
+        std::copy(const_cast<char *>(reinterpret_cast<const char *>(hostData)),
+                  const_cast<char *>(
+                    reinterpret_cast<const char *>(hostData)) + SizeInBytes,
+                  const_cast<char *>(reinterpret_cast<const char *>(BufPtr)));
     }
   }
 
   template <typename T>
   buffer_impl(const shared_ptr_class<T> &hostData, const size_t sizeInBytes,
               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 {
-      BufData.resize(get_size());
-      BufPtr = reinterpret_cast<void *>(BufData.data());
-      if (hostData.get() != nullptr) {
+      : SizeInBytes(sizeInBytes), Props(propList), BufPtr(hostData.get()),
+        MAllocator(allocator) {
+    if (!Props.has_property<property::buffer::use_host_ptr>()) {
+      BufPtr = allocateHostMem();
+      if (hostData.get()) {
         weak_ptr_class<T> hostDataWeak = hostData;
         set_final_data(hostDataWeak);
-        std::copy((char *)hostData.get(), (char *)hostData.get() + SizeInBytes,
-                  BufData.data());
+        std::copy(reinterpret_cast<char *>(hostData.get()),
+                  reinterpret_cast<char *>(hostData.get()) + SizeInBytes,
+                  reinterpret_cast<char *>(BufPtr));
       }
     }
   }
@@ -120,8 +118,7 @@ template <typename AllocatorT> class buffer_impl {
               const size_t sizeInBytes, const property_list &propList,
               AllocatorT allocator = AllocatorT())
       : SizeInBytes(sizeInBytes), Props(propList), MAllocator(allocator) {
-    BufData.resize(get_size());
-    BufPtr = reinterpret_cast<void *>(BufData.data());
+    BufPtr = allocateHostMem();
     // We need cast BufPtr to pointer to the iteration type to get correct
     // offset in std::copy when it will increment destination pointer.
     auto *Ptr =
@@ -137,8 +134,7 @@ template <typename AllocatorT> class buffer_impl {
               const size_t sizeInBytes, const property_list &propList,
               AllocatorT allocator = AllocatorT())
       : SizeInBytes(sizeInBytes), Props(propList), MAllocator(allocator) {
-    BufData.resize(get_size());
-    BufPtr = reinterpret_cast<void *>(BufData.data());
+    BufPtr = allocateHostMem();
     // We need cast BufPtr to pointer to the iteration type to get correct
     // offset in std::copy when it will increment destination pointer.
     typedef typename std::iterator_traits<InputIterator>::value_type value;
@@ -148,9 +144,10 @@ template <typename AllocatorT> class buffer_impl {
   }
 
   buffer_impl(cl_mem MemObject, const context &SyclContext,
-              const size_t sizeInBytes, event AvailableEvent = {})
+              const size_t sizeInBytes, event AvailableEvent = {},
+              AllocatorT allocator = AllocatorT())
       : OpenCLInterop(true), SizeInBytes(sizeInBytes),
-        AvailableEvent(AvailableEvent) {
+        AvailableEvent(AvailableEvent), MAllocator(allocator) {
     if (SyclContext.is_host())
       throw cl::sycl::invalid_parameter_error(
           "Creation of interoperability buffer using host context is not "
@@ -165,8 +162,7 @@ template <typename AllocatorT> class buffer_impl {
     OCLState.Mem = MemObject;
     CHECK_OCL_CODE(clRetainMemObject(MemObject));
 
-    BufData.resize(get_size());
-    BufPtr = reinterpret_cast<void *>(BufData.data());
+    BufPtr = allocateHostMem();
   }
 
   size_t get_size() const { return SizeInBytes; }
@@ -184,6 +180,12 @@ template <typename AllocatorT> class buffer_impl {
 
     if (OpenCLInterop)
       CHECK_OCL_CODE_NO_EXC(clReleaseMemObject(OCLState.Mem));
+
+    if (!Props.has_property<property::buffer::use_host_ptr>()) {
+      if (BufPtr)
+        MAllocator.deallocate(reinterpret_cast<
+            typename AllocatorT::pointer>(BufPtr), SizeInBytes);
+    }
   }
 
   void set_final_data(std::nullptr_t) { uploadData = nullptr; }
@@ -252,6 +254,13 @@ template <typename AllocatorT> class buffer_impl {
                                                   accessOffset);
   }
 
+  inline void *allocateHostMem() {
+    size_t AllocatorValueSize = sizeof(typename AllocatorT::value_type);
+    size_t AllocationSize = get_size() / AllocatorValueSize;
+    AllocationSize += (get_size() % AllocatorValueSize) ? 1 : 0;
+    return MAllocator.allocate(AllocationSize);
+  }
+
   template <typename propertyT> bool has_property() const {
     return Props.has_property<propertyT>();
   }
@@ -311,7 +320,6 @@ template <typename AllocatorT> class buffer_impl {
   event AvailableEvent;
   cl_context OpenCLContext = nullptr;
   void *BufPtr = nullptr;
-  vector_class<byte> BufData;
   // TODO: enable support of cl_mem objects from multiple contexts
   // TODO: at the current moment, using a buffer on multiple devices
   // or on a device and a host simultaneously is not supported (the
diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp
@@ -12,6 +12,7 @@
 #include <CL/sycl/access/access.hpp>
 #include <CL/sycl/context.hpp>
 #include <CL/sycl/detail/common.hpp>
+#include <CL/sycl/detail/aligned_allocator.hpp>
 #include <CL/sycl/detail/os_util.hpp>
 #include <CL/sycl/detail/scheduler/scheduler.h>
 #include <CL/sycl/event.hpp>
@@ -54,6 +55,7 @@ class queue;
 template <typename DataT, int Dimensions, access::mode AccessMode,
           access::target AccessTarget, access::placeholder IsPlaceholder>
 class accessor;
+using buffer_allocator = aligned_allocator<char, /*alignment*/ 64>;
 template <typename T, int Dimensions, typename AllocatorT> class buffer;
 namespace detail {
 
@@ -587,8 +589,7 @@ class handler {
     range<dim> Range =
         getAccessorRangeHelper<T_src, dim, mode, tgt,
                                isPlaceholder>::getAccessorRange(src);
-    // TODO use buffer_allocator when it is possible
-    buffer<T_src, dim, std::allocator<char>> Buffer(
+    buffer<T_src, dim, buffer_allocator> Buffer(
         (shared_ptr_class<T_src>)dest, Range,
         {property::buffer::use_host_ptr()});
     accessor<T_src, dim, access::mode::write, access::target::global_buffer,
@@ -608,8 +609,7 @@ class handler {
     range<dim> Range =
         getAccessorRangeHelper<T_dest, dim, mode, tgt,
                                isPlaceholder>::getAccessorRange(dest);
-    // TODO use buffer_allocator when it is possible
-    buffer<T_dest, dim, std::allocator<char>> Buffer(
+    buffer<T_dest, dim, buffer_allocator> Buffer(
         (shared_ptr_class<T_dest>)src, Range,
         {property::buffer::use_host_ptr()});
     accessor<T_dest, dim, access::mode::read, access::target::global_buffer,
@@ -628,8 +628,7 @@ class handler {
     range<dim> Range =
         getAccessorRangeHelper<T_src, dim, mode, tgt,
                                isPlaceholder>::getAccessorRange(src);
-    // TODO use buffer_allocator when it is possible
-    buffer<T_src, dim, std::allocator<char>> Buffer(
+    buffer<T_src, dim, buffer_allocator> Buffer(
         (T_src *)dest, Range, {property::buffer::use_host_ptr()});
     accessor<T_src, dim, access::mode::write, access::target::global_buffer,
              access::placeholder::false_t>
@@ -647,8 +646,7 @@ class handler {
     range<dim> Range =
         getAccessorRangeHelper<T_dest, dim, mode, tgt,
                                isPlaceholder>::getAccessorRange(dest);
-    // TODO use buffer_allocator when it is possible
-    buffer<T_dest, dim, std::allocator<char>> Buffer(
+    buffer<T_dest, dim, buffer_allocator> Buffer(
         (T_dest *)src, Range, {property::buffer::use_host_ptr()});
     accessor<T_dest, dim, access::mode::read, access::target::global_buffer,
              access::placeholder::false_t>
diff --git a/sycl/test/basic_tests/buffer/buffer.cpp b/sycl/test/basic_tests/buffer/buffer.cpp
@@ -328,6 +328,12 @@ int main() {
       buffer<int, 2> Buffer(range<2>(20, 20));
       Buffer.set_final_data((int *)result);
       queue myQueue;
+      myQueue.submit([&](handler &cgh) {
+        auto B = Buffer.get_access<access::mode::write>(cgh);
+        cgh.parallel_for<class bufferByRange2Init>(
+            range<2>{20, 20}, [=](id<2> index) { B[index] = 0; });
+      });
+
       myQueue.submit([&](handler &cgh) {
         auto B = Buffer.get_access<access::mode::write>(cgh);
         cgh.parallel_for<class bufferByRange2>(
@@ -362,6 +368,12 @@ int main() {
       buffer<int, 2> Buffer(range<2>(20, 20));
       Buffer.set_final_data((int *)result);
       queue myQueue;
+      myQueue.submit([&](handler &cgh) {
+        auto B = Buffer.get_access<access::mode::write>(cgh);
+        cgh.parallel_for<class bufferByRangeOffsetInit>(
+            range<2>{20, 20}, [=](id<2> index) { B[index] = 0; });
+      });
+
       myQueue.submit([&](handler &cgh) {
         accessor<int, 2, access::mode::write, access::target::global_buffer,
                access::placeholder::false_t>
diff --git a/sycl/test/sub_group/vote.cpp b/sycl/test/sub_group/vote.cpp
@@ -25,19 +25,29 @@ void check(queue Queue, const int G, const int L, const int D, const int R) {
     buffer<int, 1> sganybuf(G);
     buffer<int, 1> sgallbuf(G);
 
+    // Initialise buffer with zeros
+    Queue.submit([&](handler &cgh) {
+      auto sganyacc = sganybuf.get_access<access::mode::read_write>(cgh);
+      auto sgallacc = sgallbuf.get_access<access::mode::read_write>(cgh);
+      cgh.parallel_for<class init>(range<1>{(unsigned)G}, [=](id<1> index) {
+        sganyacc[index] = 0;
+        sgallacc[index] = 0;
+      });
+    });
+
     Queue.submit([&](handler &cgh) {
       auto sganyacc = sganybuf.get_access<access::mode::read_write>(cgh);
       auto sgallacc = sgallbuf.get_access<access::mode::read_write>(cgh);
       cgh.parallel_for<class subgr>(NdRange, [=](nd_item<1> NdItem) {
         intel::sub_group SG = NdItem.get_sub_group();
         /* Set to 1 if any local ID in subgroup devided by D has remainder R */
         if (SG.any(SG.get_local_id().get(0) % D == R)) {
-          sganyacc[NdItem.get_global_id()]++;
+          sganyacc[NdItem.get_global_id()] = 1;
         }
         /* Set to 1 if remainder of division of subgroup local ID by D is less
          * than R for all work items in subgroup */
         if (SG.all(SG.get_local_id().get(0) % D < R)) {
-          sgallacc[NdItem.get_global_id()]++;
+          sgallacc[NdItem.get_global_id()] = 1;
         }
       });
     });
