DPC++ Runtime
Runtime libraries for oneAPI DPC++
stream_impl.cpp
Go to the documentation of this file.
1 //==----------------- stream_impl.cpp - SYCL standard header file ---------==//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8 
9 #include <detail/buffer_impl.hpp>
10 #include <detail/queue_impl.hpp>
12 #include <detail/stream_impl.hpp>
13 #include <sycl/queue.hpp>
14 
15 #include <cstdio>
16 
17 namespace sycl {
19 namespace detail {
20 
21 stream_impl::stream_impl(size_t BufferSize, size_t MaxStatementSize,
22  handler &CGH)
23  : stream_impl(BufferSize, MaxStatementSize, {}) {
24  (void)CGH;
25 }
26 
27 stream_impl::stream_impl(size_t BufferSize, size_t MaxStatementSize,
28  const property_list &PropList)
29  : BufferSize_(BufferSize), MaxStatementSize_(MaxStatementSize),
30  PropList_(PropList), Buf_(range<1>(BufferSize + OffsetSize + 1)),
31  FlushBuf_(range<1>(MaxStatementSize + FLUSH_BUF_OFFSET_SIZE)) {
32  // Additional place is allocated in the stream buffer for the offset variable
33  // and the end of line symbol. Buffers are created without host pointers so
34  // that they are released in a deferred manner. Disable copy back on buffer
35  // destruction. Copy is scheduled as a host task which fires up as soon as
36  // kernel has completed execution.
37  Buf_.set_write_back(false);
38  FlushBuf_.set_write_back(false);
39  // Initialize stream buffer with zeros, this is needed for two reasons:
40  // 1. We don't need to care about end of line when printing out
41  // streamed data.
42  // 2. Offset is properly initialized.
43  host_accessor Acc{Buf_};
44  char *Ptr = Acc.get_pointer();
45  std::memset(Ptr, 0, Buf_.size());
46 }
47 
48 // Method to provide an access to the global stream buffer
51  CGH, range<1>(BufferSize_), id<1>(OffsetSize));
52 }
53 
54 // Method to provide an accessor to the global flush buffer
57  CGH, range<1>(MaxStatementSize_ + FLUSH_BUF_OFFSET_SIZE), id<1>(0));
58 }
59 
60 // Method to provide an atomic access to the offset in the global stream
61 // buffer and offset in the flush buffer
63  auto OffsetSubBuf = buffer<char, 1>(Buf_, id<1>(0), range<1>(OffsetSize));
64  auto ReinterpretedBuf = OffsetSubBuf.reinterpret<unsigned, 1>(range<1>(2));
65  return ReinterpretedBuf.get_access<sycl::access::mode::atomic>(
66  CGH, range<1>(2), id<1>(0));
67 }
68 
69 size_t stream_impl::size() const noexcept { return BufferSize_; }
70 
72  return MaxStatementSize_;
73 }
74 
75 size_t stream_impl::get_size() const { return BufferSize_; }
76 
77 size_t stream_impl::get_max_statement_size() const { return MaxStatementSize_; }
78 
80  // Real size of full flush buffer is saved only in buffer_impl field of
81  // FlushBuf object.
82  size_t FlushBufSize = getSyclObjImpl(FlushBuf_)->size();
83 
84  auto Q = createSyclObjFromImpl<queue>(Queue);
85  Q.submit([&](handler &cgh) {
86  auto FlushBufAcc = FlushBuf_.get_access<access::mode::discard_write,
87  access::target::host_buffer>(
88  cgh, range<1>(1), id<1>(0));
89  cgh.host_task([=] {
90  char *FlushBufPtr = FlushBufAcc.get_pointer();
91  std::memset(FlushBufPtr, 0, FlushBufSize);
92  });
93  });
94 }
95 
96 void stream_impl::flush(const EventImplPtr &LeadEvent) {
97  // We don't want stream flushing to be blocking operation that is why submit a
98  // host task to print stream buffer. It will fire up as soon as the kernel
99  // finishes execution.
100  auto Q = detail::createSyclObjFromImpl<queue>(
101  sycl::detail::Scheduler::getInstance().getDefaultHostQueue());
102  event Event = Q.submit([&](handler &cgh) {
103  auto BufHostAcc =
104  Buf_.get_access<access::mode::read_write, access::target::host_buffer>(
105  cgh, range<1>(BufferSize_), id<1>(OffsetSize));
106  // Create accessor to the flush buffer even if not using it yet. Otherwise
107  // kernel will be a leaf for the flush buffer and scheduler will not be able
108  // to cleanup the kernel. TODO: get rid of finalize method by using host
109  // accessor to the flush buffer.
110  auto FlushBufHostAcc =
111  FlushBuf_
112  .get_access<access::mode::read_write, access::target::host_buffer>(
113  cgh);
114  cgh.host_task([=] {
115  printf("%s", &(BufHostAcc[0]));
116  fflush(stdout);
117  });
118  });
119  if (LeadEvent) {
120  LeadEvent->attachEventToComplete(detail::getSyclObjImpl(Event));
121  LeadEvent->getSubmittedQueue()->registerStreamServiceEvent(
122  detail::getSyclObjImpl(Event));
123  }
124 }
125 
126 void stream_impl::flush() { flush(nullptr); }
127 } // namespace detail
128 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
129 } // namespace sycl
sycl::_V1::buffer::size
size_t size() const noexcept
Definition: buffer.hpp:483
sycl::_V1::detail::stream_impl::stream_impl
stream_impl(size_t BufferSize, size_t MaxStatementSize, handler &CGH)
Definition: stream_impl.cpp:21
sycl::_V1::property_list
Objects of the property_list class are containers for the SYCL properties.
Definition: property_list.hpp:24
sycl::_V1::detail::stream_impl::accessGlobalFlushBuf
GlobalBufAccessorT accessGlobalFlushBuf(handler &CGH)
Definition: stream_impl.cpp:55
sycl::_V1::detail::stream_impl::get_size
size_t get_size() const
Definition: stream_impl.cpp:75
__SYCL_INLINE_VER_NAMESPACE
#define __SYCL_INLINE_VER_NAMESPACE(X)
Definition: defines_elementary.hpp:11
sycl::_V1::buffer< char, 1 >
sycl::_V1::detail::stream_impl::get_work_item_buffer_size
size_t get_work_item_buffer_size() const
Definition: stream_impl.cpp:71
sycl::_V1::host_accessor
Definition: accessor.hpp:3011
sycl::_V1::detail::stream_impl::get_max_statement_size
size_t get_max_statement_size() const
Definition: stream_impl.cpp:77
sycl
---— Error handling, matching OpenCL plugin semantics.
Definition: access.hpp:14
sycl::_V1::detail::stream_impl::flush
void flush()
Definition: stream_impl.cpp:126
queue_impl.hpp
sycl::_V1::id< 1 >
scheduler.hpp
sycl::_V1::range
Defines the iteration domain of either a single work-group in a parallel dispatch,...
Definition: buffer.hpp:28
sycl::_V1::detail::stream_impl
Definition: stream_impl.hpp:25
sycl::_V1::handler
Command group handler class.
Definition: handler.hpp:315
sycl::_V1::detail::stream_impl::size
size_t size() const noexcept
Definition: stream_impl.cpp:69
queue.hpp
sycl::_V1::read_write
constexpr mode_tag_t< access_mode::read_write > read_write
Definition: access.hpp:74
sycl::_V1::detail::EventImplPtr
std::shared_ptr< event_impl > EventImplPtr
Definition: cg.hpp:42
sycl::_V1::detail::stream_impl::initStreamHost
void initStreamHost(QueueImplPtr Queue)
Definition: stream_impl.cpp:79
sycl::_V1::detail::FLUSH_BUF_OFFSET_SIZE
constexpr unsigned FLUSH_BUF_OFFSET_SIZE
Definition: stream.hpp:60
sycl::_V1::accessor
Definition: accessor.hpp:225
sycl::_V1::detail::QueueImplPtr
std::shared_ptr< sycl::detail::queue_impl > QueueImplPtr
Definition: event_impl.hpp:32
sycl::_V1::detail::stream_impl::accessGlobalOffset
GlobalOffsetAccessorT accessGlobalOffset(handler &CGH)
Definition: stream_impl.cpp:62
sycl::_V1::buffer::get_access
accessor< T, dimensions, Mode, Target, access::placeholder::false_t, ext::oneapi::accessor_property_list<> > get_access(handler &CommandGroupHandler, const detail::code_location CodeLoc=detail::code_location::current())
Definition: buffer.hpp:498
sycl::_V1::detail::stream_impl::accessGlobalBuf
GlobalBufAccessorT accessGlobalBuf(handler &CGH)
Definition: stream_impl.cpp:49
sycl::_V1::buffer::set_write_back
void set_write_back(bool flag=true)
Definition: buffer.hpp:639
sycl::_V1::ext::oneapi::experimental::printf
int printf(const FormatT *__format, Args... args)
Definition: builtins.hpp:75
sycl::_V1::access::mode::read_write
@ read_write
stream_impl.hpp
sycl::_V1::access::mode::discard_write
@ discard_write
sycl::_V1::detail::getSyclObjImpl
decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject)
Definition: common.hpp:300
buffer_impl.hpp