Another portion of applied remarks

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

Author: mlychkov

sycl/source/detail/scheduler/scheduler_helpers.hpp

@@ -16,22 +16,32 @@ 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.");
+  Scheduler::StreamBuffers *StrBufs{};
 
-  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();
+  {
+    std::lock_guard<std::mutex> lock(
+        Scheduler::getInstance().StreamBuffersPoolMutex);
+
+    auto StreamBuf =
+        Scheduler::getInstance().StreamBuffersPool.find(Stream.get());
+    assert((StreamBuf != Scheduler::getInstance().StreamBuffersPool.end()) &&
+           "Stream is unexpectedly not found in pool.");
+
+    StrBufs = StreamBuf->second;
+  }
+
+  assert(StrBufs && "No buffers for a stream.");
+
+  // Real size of full flush buffer is saved only in buffer_impl field of
+  // FlushBuf object.
+  size_t FlushBufSize = getSyclObjImpl(StrBufs->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));
+    auto FlushBufAcc =
+        StrBufs->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());

sycl/source/stream.cpp

@@ -6,7 +6,6 @@
 //
 //===----------------------------------------------------------------------===//
 
-//#include <CL/sycl/exception.hpp>
 #include <CL/sycl/stream.hpp>
 #include <detail/queue_impl.hpp>
 #include <detail/stream_impl.hpp>
@@ -27,10 +26,10 @@ stream::stream(size_t BufferSize, size_t MaxStatementSize, handler &CGH)
       GlobalFlushBuf(impl->accessGlobalFlushBuf(CGH)),
       FlushBufferSize(MaxStatementSize + detail::FLUSH_BUF_OFFSET_SIZE) {
   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);
+    throw sycl::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

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

@@ -0,0 +1,36 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: env SYCL_DEVICE_TYPE=HOST SYCL_PRINT_EXECUTION_GRAPH=after_addCG %t.out
+// RUN: cat graph_2after_addCG.dot | FileCheck %s
+// RUN: rm *.dot
+
+// This test is executed on host device only because flush buffer initialized
+// via a separate call group task for this device only.
+
+#include <CL/sycl.hpp>
+
+using namespace cl;
+
+int main() {
+  {
+    sycl::queue Queue;
+    Queue.submit([&](sycl::handler &cgh) {
+      sycl::stream Out(100, 100, cgh);
+      cgh.single_task<class test_flush_buf_init_dep>(
+          [=]() { Out << "Hello world!" << sycl::endl; });
+    });
+    Queue.wait();
+  }
+  // CHECK: [[MAIN_CG:"0x[0-9a-f]+"]] {{.*}}EXEC CG ON HOST{{.*}}test_flush_buf_init_dep
+  // First 3 lines show dependencies on global buffer related commands
+  // CHECK:   [[MAIN_CG]]
+  // CHECK:   [[MAIN_CG]]
+  // CHECK:   [[MAIN_CG]]
+  // CHECK:   [[MAIN_CG]] -> [[EMPTY_NODE:"0x[0-9a-f]+"]] [ label = "Access mode: read_write\nMemObj: [[FLUSHBUF_MEMOBJ:0x[0-9a-f]+]]
+  // CHECK: [[EMPTY_NODE]] {{.*}}EMPTY NODE
+  // CHECK:   [[EMPTY_NODE]] -> [[FILL_TASK:"0x[0-9a-f]+"]] [ label = "Access mode: discard_write\nMemObj: [[FLUSHBUF_MEMOBJ]]
+  // CHECK: [[FILL_TASK]] {{.*}}EXEC CG ON HOST\nCG type: host task
+  // CHECK:   [[FILL_TASK]] -> [[ALLOC_TASK:"0x[0-9a-f]+"]] [ label = "Access mode: discard_write\nMemObj: [[FLUSHBUF_MEMOBJ]]
+  // CHECK: [[ALLOC_TASK]] {{.*}}ALLOCA ON HOST\n MemObj : [[FLUSHBUF_MEMOBJ]]
+}
+  return 0;
+}