[SYCL] Use a global flush buffer in stream

Change stream implementation to use a global buffer for all the
work-items to flush their outputs into. Previously a set of local flush
buffers were used for this purpose.

This change is made to avoid the issue of some devices running out of
local memory when available space is less than statement size times the
number of work-items.

Signed-off-by: Shahab Layeghi <[email protected]>

Author: Shahab Layeghi

sycl/include/CL/sycl/detail/accessor_impl.hpp

@@ -79,6 +79,8 @@ class __SYCL_EXPORT AccessorImplHost {
         MElemSize(Other.MElemSize), MOffsetInBytes(Other.MOffsetInBytes),
         MIsSubBuffer(Other.MIsSubBuffer) {}
 
+  void resize(size_t GlobalSize);
+
   id<3> MOffset;
   // The size of accessing region.
   range<3> MAccessRange;
@@ -96,6 +98,8 @@ class __SYCL_EXPORT AccessorImplHost {
   void *MData = nullptr;
 
   Command *MBlockedCmd = nullptr;
+
+  bool PerWI = false;
 };
 
 using AccessorImplPtr = shared_ptr_class<AccessorImplHost>;
@@ -139,23 +143,6 @@ class __SYCL_EXPORT LocalAccessorImplHost {
   int MDims;
   int MElemSize;
   std::vector<char> MMem;
-
-  bool PerWI = false;
-  size_t LocalMemSize;
-  size_t MaxWGSize;
-  void resize(size_t LocalSize, size_t GlobalSize) {
-    if (GlobalSize != 1 && LocalSize != 1) {
-      // If local size is not specified then work group size is chosen by
-      // runtime. That is why try to allocate based on max work group size or
-      // global size. In the worst case allocate 80% of local memory.
-      size_t MinEstWGSize = LocalSize ? LocalSize : GlobalSize;
-      MinEstWGSize = MinEstWGSize > MaxWGSize ? MaxWGSize : MinEstWGSize;
-      size_t NewSize = MinEstWGSize * MSize[0];
-      MSize[0] =
-          NewSize > 8 * LocalMemSize / 10 ? 8 * LocalMemSize / 10 : NewSize;
-      MMem.resize(NewSize * MElemSize);
-    }
-  }
 };
 
 using LocalAccessorImplPtr = shared_ptr_class<LocalAccessorImplHost>;

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

@@ -113,6 +113,8 @@ class __SYCL_EXPORT buffer_impl final : public SYCLMemObjT {
     } catch (...) {
     }
   }
+
+  void resize(size_t size) { BaseT::MSizeInBytes = size; }
 };
 
 } // namespace detail

sycl/include/CL/sycl/detail/stream_impl.hpp

@@ -89,16 +89,6 @@ class __SYCL_EXPORT stream_impl {
                cl::sycl::access::target::global_buffer,
                cl::sycl::access::placeholder::false_t>;
 
-  using FlushBufAccessorT =
-      accessor<char, 1, cl::sycl::access::mode::read_write,
-               cl::sycl::access::target::local,
-               cl::sycl::access::placeholder::false_t>;
-
-  using LocalOffsetAccessorT =
-      accessor<unsigned, 1, cl::sycl::access::mode::atomic,
-               cl::sycl::access::target::local,
-               cl::sycl::access::placeholder::false_t>;
-
   stream_impl(size_t BufferSize, size_t MaxStatementSize, handler &CGH);
 
   // Method to provide an access to the global stream buffer
@@ -107,6 +97,12 @@ class __SYCL_EXPORT stream_impl {
         CGH, range<1>(BufferSize_), id<1>(OffsetSize));
   }
 
+  // Method to provide an accessor to the global flush buffer
+  GlobalBufAccessorT accessGlobalFlushBuf(handler &CGH) {
+    return FlushBuf.get_access<cl::sycl::access::mode::read_write>(
+        CGH, range<1>(MaxStatementSize_), id<1>(0));
+  }
+
   // Method to provide an atomic access to the offset in the global stream
   // buffer
   GlobalOffsetAccessorT accessGlobalOffset(handler &CGH) {
@@ -116,12 +112,6 @@ class __SYCL_EXPORT stream_impl {
         CGH, range<1>(1), id<1>(0));
   }
 
-  // Method to provide an atomic access to the flush buffer size
-  GlobalOffsetAccessorT accessFlushBufferSize(handler &CGH) {
-    return FlushBufferSize.get_access<cl::sycl::access::mode::atomic>(
-        CGH, range<1>(1), id<1>(0));
-  }
-
   // Copy stream buffer to the host and print the contents
   void flush();
 
@@ -148,8 +138,8 @@ class __SYCL_EXPORT stream_impl {
   // Stream buffer
   buffer<char, 1> Buf;
 
-  //  Buffer for flush buffer size
-  buffer<unsigned, 1> FlushBufferSize;
+  // Global flush buffer
+  buffer<char, 1> FlushBuf;
 };
 
 template <typename T>
@@ -309,35 +299,34 @@ inline bool updateOffset(stream_impl::GlobalOffsetAccessorT &GlobalOffset,
 
 inline void flushBuffer(stream_impl::GlobalOffsetAccessorT &GlobalOffset,
                         stream_impl::GlobalBufAccessorT &GlobalBuf,
-                        stream_impl::FlushBufAccessorT &FlushBufs,
+                        stream_impl::GlobalBufAccessorT &GlobalFlushBuf,
                         unsigned &WIOffset, unsigned &Offset) {
-  // Copy data from flush buffer (local memory) to global buffer (global
-  // memory)
+
   unsigned Cur = 0;
   if (!updateOffset(GlobalOffset, GlobalBuf, Offset, Cur))
     return;
 
   for (unsigned I = WIOffset; I < WIOffset + Offset; I++) {
-    GlobalBuf[Cur++] = FlushBufs[I];
+    GlobalBuf[Cur++] = GlobalFlushBuf[I];
   }
   // Reset the offset in the flush buffer
   Offset = 0;
 }
 
-inline void write(stream_impl::FlushBufAccessorT &FlushBufs,
+inline void write(stream_impl::GlobalBufAccessorT &GlobalFlushBuf,
                   size_t FlushBufferSize, unsigned WIOffset, unsigned &Offset,
                   const char *Str, unsigned Len, unsigned Padding = 0) {
   if ((FlushBufferSize - Offset < Len + Padding) ||
-      (WIOffset + Offset + Len + Padding > FlushBufs.get_count()))
+      (WIOffset + Offset + Len + Padding > GlobalFlushBuf.get_count()))
     // TODO: flush here
     return;
 
   // Write padding
   for (size_t I = 0; I < Padding; ++I, ++Offset)
-    FlushBufs[WIOffset + Offset] = ' ';
+    GlobalFlushBuf[WIOffset + Offset] = ' ';
 
   for (size_t I = 0; I < Len; ++I, ++Offset) {
-    FlushBufs[WIOffset + Offset] = Str[I];
+    GlobalFlushBuf[WIOffset + Offset] = Str[I];
   }
 }
 
@@ -475,25 +464,25 @@ ScalarToStr(const T &Val, char *Buf, unsigned Flags, int, int Precision = -1) {
 
 template <typename T>
 inline typename std::enable_if<std::is_integral<T>::value>::type
-writeIntegral(stream_impl::FlushBufAccessorT &FlushBufs, size_t FlushBufferSize,
-              unsigned WIOffset, unsigned &Offset, unsigned Flags, int Width,
-              const T &Val) {
+writeIntegral(stream_impl::GlobalBufAccessorT &GlobalFlushBuf,
+              size_t FlushBufferSize, unsigned WIOffset, unsigned &Offset,
+              unsigned Flags, int Width, const T &Val) {
   char Digits[MAX_INTEGRAL_DIGITS] = {0};
   unsigned Len = ScalarToStr(Val, Digits, Flags, Width);
-  write(FlushBufs, FlushBufferSize, WIOffset, Offset, Digits, Len,
+  write(GlobalFlushBuf, FlushBufferSize, WIOffset, Offset, Digits, Len,
         (Width > 0 && static_cast<unsigned>(Width) > Len)
             ? static_cast<unsigned>(Width) - Len
             : 0);
 }
 
 template <typename T>
 inline EnableIfFP<T>
-writeFloatingPoint(stream_impl::FlushBufAccessorT &FlushBufs,
+writeFloatingPoint(stream_impl::GlobalBufAccessorT &GlobalFlushBuf,
                    size_t FlushBufferSize, unsigned WIOffset, unsigned &Offset,
                    unsigned Flags, int Width, int Precision, const T &Val) {
   char Digits[MAX_FLOATING_POINT_DIGITS] = {0};
   unsigned Len = ScalarToStr(Val, Digits, Flags, Width, Precision);
-  write(FlushBufs, FlushBufferSize, WIOffset, Offset, Digits, Len,
+  write(GlobalFlushBuf, FlushBufferSize, WIOffset, Offset, Digits, Len,
         (Width > 0 && static_cast<unsigned>(Width) > Len)
             ? static_cast<unsigned>(Width) - Len
             : 0);
@@ -531,7 +520,7 @@ VecToStr(const vec<T, VecLength> &Vec, char *VecStr, unsigned Flags, int Width,
 }
 
 template <typename T, int VecLength>
-inline void writeVec(stream_impl::FlushBufAccessorT &FlushBufs,
+inline void writeVec(stream_impl::GlobalBufAccessorT &GlobalFlushBuf,
                      size_t FlushBufferSize, unsigned WIOffset,
                      unsigned &Offset, unsigned Flags, int Width, int Precision,
                      const vec<T, VecLength> &Vec) {
@@ -540,7 +529,7 @@ inline void writeVec(stream_impl::FlushBufAccessorT &FlushBufs,
       MAX_FLOATING_POINT_DIGITS * VecLength + (VecLength - 1) * 2;
   char VecStr[MAX_VEC_SIZE] = {0};
   unsigned Len = VecToStr<T, VecLength>(Vec, VecStr, Flags, Width, Precision);
-  write(FlushBufs, FlushBufferSize, WIOffset, Offset, VecStr, Len,
+  write(GlobalFlushBuf, FlushBufferSize, WIOffset, Offset, VecStr, Len,
         (Width > 0 && Width > Len) ? Width - Len : 0);
 }
 
@@ -561,16 +550,16 @@ inline unsigned ArrayToStr(char *Buf, const array<ArrayLength> &Arr) {
 }
 
 template <int ArrayLength>
-inline void writeArray(stream_impl::FlushBufAccessorT &FlushBufs,
+inline void writeArray(stream_impl::GlobalBufAccessorT &GlobalFlushBuf,
                        size_t FlushBufferSize, unsigned WIOffset,
                        unsigned &Offset, const array<ArrayLength> &Arr) {
   char Buf[MAX_ARRAY_SIZE];
   unsigned Len = ArrayToStr(Buf, Arr);
-  write(FlushBufs, FlushBufferSize, WIOffset, Offset, Buf, Len);
+  write(GlobalFlushBuf, FlushBufferSize, WIOffset, Offset, Buf, Len);
 }
 
 template <int Dimensions>
-inline void writeItem(stream_impl::FlushBufAccessorT &FlushBufs,
+inline void writeItem(stream_impl::GlobalBufAccessorT &GlobalFlushBuf,
                       size_t FlushBufferSize, unsigned WIOffset,
                       unsigned &Offset, const item<Dimensions> &Item) {
   // Reserve space for 3 arrays and additional place (40 symbols) for printing
@@ -585,11 +574,11 @@ inline void writeItem(stream_impl::FlushBufAccessorT &FlushBufs,
   Len += append(Buf + Len, ", offset: ");
   Len += ArrayToStr(Buf + Len, Item.get_offset());
   Buf[Len++] = ')';
-  write(FlushBufs, FlushBufferSize, WIOffset, Offset, Buf, Len);
+  write(GlobalFlushBuf, FlushBufferSize, WIOffset, Offset, Buf, Len);
 }
 
 template <int Dimensions>
-inline void writeNDRange(stream_impl::FlushBufAccessorT &FlushBufs,
+inline void writeNDRange(stream_impl::GlobalBufAccessorT &GlobalFlushBuf,
                          size_t FlushBufferSize, unsigned WIOffset,
                          unsigned &Offset,
                          const nd_range<Dimensions> &ND_Range) {
@@ -605,11 +594,11 @@ inline void writeNDRange(stream_impl::FlushBufAccessorT &FlushBufs,
   Len += append(Buf + Len, ", offset: ");
   Len += ArrayToStr(Buf + Len, ND_Range.get_offset());
   Buf[Len++] = ')';
-  write(FlushBufs, FlushBufferSize, WIOffset, Offset, Buf, Len);
+  write(GlobalFlushBuf, FlushBufferSize, WIOffset, Offset, Buf, Len);
 }
 
 template <int Dimensions>
-inline void writeNDItem(stream_impl::FlushBufAccessorT &FlushBufs,
+inline void writeNDItem(stream_impl::GlobalBufAccessorT &GlobalFlushBuf,
                         size_t FlushBufferSize, unsigned WIOffset,
                         unsigned &Offset, const nd_item<Dimensions> &ND_Item) {
   // Reserve space for 2 arrays and additional place (40 symbols) for printing
@@ -622,11 +611,11 @@ inline void writeNDItem(stream_impl::FlushBufAccessorT &FlushBufs,
   Len += append(Buf + Len, ", local_id: ");
   Len += ArrayToStr(Buf + Len, ND_Item.get_local_id());
   Buf[Len++] = ')';
-  write(FlushBufs, FlushBufferSize, WIOffset, Offset, Buf, Len);
+  write(GlobalFlushBuf, FlushBufferSize, WIOffset, Offset, Buf, Len);
 }
 
 template <int Dimensions>
-inline void writeGroup(stream_impl::FlushBufAccessorT &FlushBufs,
+inline void writeGroup(stream_impl::GlobalBufAccessorT &GlobalFlushBuf,
                        size_t FlushBufferSize, unsigned WIOffset,
                        unsigned &Offset, const group<Dimensions> &Group) {
   // Reserve space for 4 arrays and additional place (60 symbols) for printing
@@ -643,7 +632,7 @@ inline void writeGroup(stream_impl::FlushBufAccessorT &FlushBufs,
   Len += append(Buf + Len, ", group_range: ");
   Len += ArrayToStr(Buf + Len, Group.get_group_range());
   Buf[Len++] = ')';
-  write(FlushBufs, FlushBufferSize, WIOffset, Offset, Buf, Len);
+  write(GlobalFlushBuf, FlushBufferSize, WIOffset, Offset, Buf, Len);
 }
 
 // Space for 2 arrays and additional place (20 symbols) for printing
@@ -663,7 +652,7 @@ inline unsigned ItemToStr(char *Buf, const item<Dimensions, false> &Item) {
 }
 
 template <int Dimensions>
-inline void writeHItem(stream_impl::FlushBufAccessorT &FlushBufs,
+inline void writeHItem(stream_impl::GlobalBufAccessorT &GlobalFlushBuf,
                        size_t FlushBufferSize, unsigned WIOffset,
                        unsigned &Offset, const h_item<Dimensions> &HItem) {
   // Reserve space for 3 items and additional place (60 symbols) for printing
@@ -680,7 +669,7 @@ inline void writeHItem(stream_impl::FlushBufAccessorT &FlushBufs,
                                                 : HItem.get_physical_local());
   }
   Len += append(Buf + Len, "\n)");
-  write(FlushBufs, FlushBufferSize, WIOffset, Offset, Buf, Len);
+  write(GlobalFlushBuf, FlushBufferSize, WIOffset, Offset, Buf, Len);
 }
 
 } // namespace detail