Improve fix by applying new review remarks

Signed-off-by: Mikhail Lychkov <[email protected]>

Author: mlychkov

sycl/include/CL/sycl/stream.hpp

@@ -73,16 +73,26 @@ using GlobalOffsetAccessorT =
              cl::sycl::access::target::global_buffer,
              cl::sycl::access::placeholder::false_t>;
 
+// Read first 2 bytes of flush buffer to get buffer offset.
+// TODO: Should be optimized to the following:
+//   return *reinterpret_cast<uint16_t *>(&GlobalFlushBuf[WIOffset]);
+// when an issue with device code compilation using this optimization is fixed.
 inline unsigned GetFlushBufOffset(const GlobalBufAccessorT &GlobalFlushBuf,
                                   unsigned WIOffset) {
-  return (((uint8_t)GlobalFlushBuf[WIOffset]) << 8) +
-         (uint8_t)GlobalFlushBuf[WIOffset + 1];
+  return ((static_cast<unsigned>(static_cast<uint8_t>(GlobalFlushBuf[WIOffset]))
+           << 8) +
+          static_cast<uint8_t>(GlobalFlushBuf[WIOffset + 1]));
 }
 
+// Write flush buffer's offset into first 2 bytes of that buffer.
+// TODO: Should be optimized to the following:
+//   *reinterpret_cast<uint16_t *>(&GlobalFlushBuf[WIOffset]) =
+//       static_cast<uint16_t>(Offset);
+// when an issue with device code compilation using this optimization is fixed.
 inline void SetFlushBufOffset(GlobalBufAccessorT &GlobalFlushBuf,
                               unsigned WIOffset, unsigned Offset) {
-  GlobalFlushBuf[WIOffset] = (Offset >> 8) & 0xff;
-  GlobalFlushBuf[WIOffset + 1] = Offset & 0xff;
+  GlobalFlushBuf[WIOffset] = static_cast<char>((Offset >> 8) & 0xff);
+  GlobalFlushBuf[WIOffset + 1] = static_cast<char>(Offset & 0xff);
 }
 
 inline void write(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize,
@@ -719,6 +729,7 @@ inline __width_manipulator__ setw(int Width) {
 /// \ingroup sycl_api
 class __SYCL_EXPORT stream {
 public:
+  // Throws exception in case of invalid input parameters
   stream(size_t BufferSize, size_t MaxStatementSize, handler &CGH);
 
   size_t get_size() const;
@@ -841,6 +852,10 @@ class __SYCL_EXPORT stream {
     // overhead on FPGA target. That is why use global atomic variable to
     // calculate offsets.
     WIOffset = GlobalOffset[1].fetch_add(FlushBufferSize);
+
+    // Initialize flush subbuffer's offset for each work item on device.
+    // Initialization on host device is performed via submition of additional
+    // host task.
     SetFlushBufOffset(GlobalFlushBuf, WIOffset, 0);
   }
 

sycl/source/detail/scheduler/scheduler.cpp

@@ -11,6 +11,7 @@
 #include <detail/global_handler.hpp>
 #include <detail/queue_impl.hpp>
 #include <detail/scheduler/scheduler.hpp>
+#include <detail/scheduler/scheduler_helpers.hpp>
 #include <detail/stream_impl.hpp>
 
 #include <chrono>
@@ -75,10 +76,12 @@ EventImplPtr Scheduler::addCG(std::unique_ptr<detail::CG> CommandGroup,
 
   if (IsKernel) {
     Streams = ((CGExecKernel *)CommandGroup.get())->getStreams();
+    // Stream's flush buffer memeory is mainly initialized in stream's __init
+    // method. However, this method is not available on host device.
+    // Initializing stream's flush buffer on the host side in a separate task.
     if (Queue->is_host()) {
-      // Initializing stream's flush buffer on the host side.
-      for (auto StreamImplPtr : Streams) {
-        StreamImplPtr->fill(Queue);
+      for (const StreamImplPtr &Stream : Streams) {
+        initStream(Stream, Queue);
       }
     }
   }

sycl/source/detail/scheduler/scheduler.hpp

@@ -764,6 +764,7 @@ class Scheduler {
   };
 
   friend class stream_impl;
+  friend void initStream(StreamImplPtr, QueueImplPtr);
 
   // Protects stream buffers pool
   std::recursive_mutex StreamBuffersPoolMutex;

sycl/source/detail/scheduler/scheduler_helpers.hpp

@@ -0,0 +1,44 @@
+//==---------- scheduler_helpers.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/sycl/queue.hpp>
+#include <detail/scheduler/scheduler.hpp>
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace detail {
+
+void initStream(StreamImplPtr Stream, QueueImplPtr Queue) {
+  auto StreamBuf =
+      Scheduler::getInstance().StreamBuffersPool.find(Stream.get());
+  assert((StreamBuf != Scheduler::getInstance().StreamBuffersPool.end()) &&
+         "Stream is unexpectedly not found in pool.");
+
+  auto &FlushBuf = StreamBuf->second->FlushBuf;
+  // Only size of buffer_impl object has been resized.
+  // Value of Range field of FlushBuf instance is still equal to
+  // MaxStatementSize only.
+  size_t FlushBufSize = getSyclObjImpl(FlushBuf)->get_count();
+
+  auto Q = createSyclObjFromImpl<queue>(Queue);
+  Q.submit([&](handler &cgh) {
+    auto FlushBufAcc = FlushBuf.get_access<access::mode::discard_write,
+                                           access::target::host_buffer>(
+        cgh, range<1>(FlushBufSize), id<1>(0));
+    cgh.codeplay_host_task([=] {
+      char *FlushBufPtr = FlushBufAcc.get_pointer();
+      std::memset(FlushBufPtr, 0, FlushBufAcc.get_size());
+    });
+  });
+}
+
+} // namespace detail
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)

sycl/source/detail/stream_impl.cpp

@@ -61,29 +61,6 @@ size_t stream_impl::get_size() const { return BufferSize_; }
 
 size_t stream_impl::get_max_statement_size() const { return MaxStatementSize_; }
 
-void stream_impl::fill(QueueImplPtr Queue) {
-  auto Q = detail::createSyclObjFromImpl<queue>(Queue);
-  Q.submit([&](handler &cgh) {
-    auto StreamBuf =
-        detail::Scheduler::getInstance().StreamBuffersPool.find(this);
-    assert((StreamBuf !=
-            detail::Scheduler::getInstance().StreamBuffersPool.end()) &&
-           "Stream is unexpectedly not found in pool");
-    auto &FlushBuf = StreamBuf->second->FlushBuf;
-    // Only size of buffer_impl object has been resized.
-    // Value of Range field of FlushBuf instance is still equal to
-    // MaxStatementSize only.
-    size_t FlushBufSize = detail::getSyclObjImpl(FlushBuf)->get_count();
-    auto FlushBufAcc = FlushBuf.get_access<access::mode::read_write,
-                                           access::target::global_buffer>(
-        cgh, range<1>(FlushBufSize), id<1>(0));
-    cgh.codeplay_host_task([=] {
-      char *FlushBufPtr = FlushBufAcc.get_pointer();
-      std::memset(FlushBufPtr, 0, FlushBufAcc.get_size());
-    });
-  });
-}
-
 void stream_impl::flush() {
   // We don't want stream flushing to be blocking operation that is why submit a
   // host task to print stream buffer. It will fire up as soon as the kernel

sycl/source/detail/stream_impl.hpp

@@ -37,8 +37,6 @@ class __SYCL_EXPORT stream_impl {
   // buffer and offset in the flush buffer
   GlobalOffsetAccessorT accessGlobalOffset(handler &CGH);
 
-  void fill(QueueImplPtr Queue);
-
   // Enqueue task to copy stream buffer to the host and print the contents
   void flush();
 

sycl/source/stream.cpp

@@ -6,6 +6,7 @@
 //
 //===----------------------------------------------------------------------===//
 
+//#include <CL/sycl/exception.hpp>
 #include <CL/sycl/stream.hpp>
 #include <detail/queue_impl.hpp>
 #include <detail/stream_impl.hpp>
@@ -25,8 +26,12 @@ stream::stream(size_t BufferSize, size_t MaxStatementSize, handler &CGH)
       // Allocate the flush buffer, which contains space for each work item
       GlobalFlushBuf(impl->accessGlobalFlushBuf(CGH)),
       FlushBufferSize(MaxStatementSize + detail::FLUSH_BUF_OFFSET_SIZE) {
-  assert((MaxStatementSize <= MAX_STATEMENT_SIZE) &&
-         "Maximum statement size too large.");
+  if (MaxStatementSize > MAX_STATEMENT_SIZE) {
+    throw invalid_parameter_error("Maximum statement size exceeds limit of " +
+                                      std::to_string(MAX_STATEMENT_SIZE) +
+                                      " bytes.",
+                                  PI_INVALID_VALUE);
+  }
 
   // Save stream implementation in the handler so that stream will be alive
   // during kernel execution

sycl/test/abi/sycl_symbols_linux.dump

@@ -3725,7 +3725,6 @@ _ZN2cl4sycl6detail11buffer_impl11allocateMemESt10shared_ptrINS1_12context_implEE
 _ZN2cl4sycl6detail11stream_impl15accessGlobalBufERNS0_7handlerE
 _ZN2cl4sycl6detail11stream_impl18accessGlobalOffsetERNS0_7handlerE
 _ZN2cl4sycl6detail11stream_impl20accessGlobalFlushBufERNS0_7handlerE
-_ZN2cl4sycl6detail11stream_impl4fillESt10shared_ptrINS1_10queue_implEE
 _ZN2cl4sycl6detail11stream_impl5flushEv
 _ZN2cl4sycl6detail11stream_implC1EmmRNS0_7handlerE
 _ZN2cl4sycl6detail11stream_implC2EmmRNS0_7handlerE

sycl/test/on-device/basic_tests/stream/stream_copies_buffer_sync.cpp

@@ -4,14 +4,6 @@
 // RUN: %GPU_RUN_ON_LINUX_PLACEHOLDER %t.out %GPU_CHECK_ON_LINUX_PLACEHOLDER
 // RUN: %ACC_RUN_PLACEHOLDER %t.out %ACC_CHECK_PLACEHOLDER
 
-//==-- stream_copies_buffer_sync.cpp - SYCL stream flush buffer sync test --==//
-//
-// 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
-//
-//===----------------------------------------------------------------------===//
-
 #include <CL/sycl.hpp>
 
 using namespace cl::sycl;

sycl/test/on-device/basic_tests/stream/stream_max_stmt_exceed.cpp

@@ -0,0 +1,27 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: env SYCL_DEVICE_TYPE=HOST %t.out | FileCheck %s
+// RUN: %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER
+// RUN: %GPU_RUN_ON_LINUX_PLACEHOLDER %t.out %GPU_CHECK_ON_LINUX_PLACEHOLDER
+// RUN: %ACC_RUN_PLACEHOLDER %t.out %ACC_CHECK_PLACEHOLDER
+
+#include <CL/sycl.hpp>
+
+#include <cassert>
+
+using namespace cl;
+
+int main() {
+  sycl::queue Queue;
+  try {
+    Queue.submit([&](sycl::handler &cgh) {
+      sycl::stream Out(100, 65536, cgh);
+      cgh.single_task<class test_max_stmt_exceed>(
+          [=]() { Out << "Hello world!" << sycl::endl; });
+    });
+    Queue.wait();
+  } catch (sycl::exception &ExpectedException) {
+    // CHECK: Maximum statement size exceeds limit of 65535 bytes
+    std::cout << ExpectedException.what() << std::endl;
+  }
+  return 0;
+}