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 <CL/sycl/queue.hpp>
11 #include <detail/stream_impl.hpp>
12 
13 #include <cstdio>
14 
16 namespace sycl {
17 namespace detail {
18 
19 stream_impl::stream_impl(size_t BufferSize, size_t MaxStatementSize,
20  handler &CGH)
21  : stream_impl(BufferSize, MaxStatementSize, {}) {
22  (void)CGH;
23 }
24 
25 stream_impl::stream_impl(size_t BufferSize, size_t MaxStatementSize,
26  const property_list &PropList)
27  : BufferSize_(BufferSize), MaxStatementSize_(MaxStatementSize),
28  PropList_(PropList) {
29  // We need to store stream buffers in the scheduler because they need to be
30  // alive after submitting the kernel. They cannot be stored in the stream
31  // object because it causes loop dependency between objects and results in
32  // memory leak.
33  // Allocate additional place in the stream buffer for the offset variable and
34  // the end of line symbol.
36  this, BufferSize + OffsetSize + 1 /* size of the stream buffer */,
37  MaxStatementSize + FLUSH_BUF_OFFSET_SIZE /* size of the flush buffer */);
38 }
39 
40 // Method to provide an access to the global stream buffer
43  .StreamBuffersPool.find(this)
44  ->second->Buf.get_access<cl::sycl::access::mode::read_write>(
45  CGH, range<1>(BufferSize_), id<1>(OffsetSize));
46 }
47 
48 // Method to provide an accessor to the global flush buffer
51  .StreamBuffersPool.find(this)
52  ->second->FlushBuf.get_access<cl::sycl::access::mode::read_write>(
53  CGH, range<1>(MaxStatementSize_ + FLUSH_BUF_OFFSET_SIZE), id<1>(0));
54 }
55 
56 // Method to provide an atomic access to the offset in the global stream
57 // buffer and offset in the flush buffer
60  .StreamBuffersPool.find(this)
61  ->second->Buf,
62  id<1>(0), range<1>(OffsetSize));
63  auto ReinterpretedBuf = OffsetSubBuf.reinterpret<unsigned, 1>(range<1>(2));
64  return ReinterpretedBuf.get_access<cl::sycl::access::mode::atomic>(
65  CGH, range<1>(2), id<1>(0));
66 }
67 size_t stream_impl::get_size() const { return BufferSize_; }
68 
69 size_t stream_impl::get_max_statement_size() const { return MaxStatementSize_; }
70 
72  // We don't want stream flushing to be blocking operation that is why submit a
73  // host task to print stream buffer. It will fire up as soon as the kernel
74  // finishes execution.
75  auto Q = detail::createSyclObjFromImpl<queue>(
76  cl::sycl::detail::Scheduler::getInstance().getDefaultHostQueue());
77  Q.submit([&](handler &cgh) {
78  auto BufHostAcc =
80  .StreamBuffersPool.find(this)
81  ->second->Buf
83  cgh, range<1>(BufferSize_), id<1>(OffsetSize));
84  // Create accessor to the flush buffer even if not using it yet. Otherwise
85  // kernel will be a leaf for the flush buffer and scheduler will not be able
86  // to cleanup the kernel. TODO: get rid of finalize method by using host
87  // accessor to the flush buffer.
88  auto FlushBufHostAcc =
90  .StreamBuffersPool.find(this)
91  ->second->FlushBuf
93  cgh);
94  cgh.host_task([=] {
95  printf("%s", &(BufHostAcc[0]));
96  fflush(stdout);
97  });
98  });
99 }
100 } // namespace detail
101 } // namespace sycl
102 } // __SYCL_INLINE_NAMESPACE(cl)
103 
cl::sycl::detail::stream_impl
Definition: stream_impl.hpp:25
cl::sycl::detail::FLUSH_BUF_OFFSET_SIZE
constexpr unsigned FLUSH_BUF_OFFSET_SIZE
Definition: stream.hpp:58
cl::sycl::detail::stream_impl::accessGlobalFlushBuf
GlobalBufAccessorT accessGlobalFlushBuf(handler &CGH)
Definition: stream_impl.cpp:49
cl::sycl::id< 1 >
cl::sycl::detail::Scheduler::allocateStreamBuffers
void allocateStreamBuffers(stream_impl *, size_t, size_t)
Allocate buffers in the pool for a provided stream.
Definition: scheduler.cpp:379
cl::sycl::detail::stream_impl::accessGlobalBuf
GlobalBufAccessorT accessGlobalBuf(handler &CGH)
Definition: stream_impl.cpp:41
cl::sycl::detail::stream_impl::accessGlobalOffset
GlobalOffsetAccessorT accessGlobalOffset(handler &CGH)
Definition: stream_impl.cpp:58
cl::sycl::access::mode::atomic
@ atomic
sycl
Definition: invoke_simd.hpp:68
cl::sycl::property_list
Objects of the property_list class are containers for the SYCL properties.
Definition: property_list.hpp:26
cl::sycl::buffer< char, 1 >
scheduler.hpp
cl::sycl::range< 1 >
cl::sycl::detail::stream_impl::get_max_statement_size
size_t get_max_statement_size() const
Definition: stream_impl.cpp:69
cl::sycl::access::target::host_buffer
@ host_buffer
cl::sycl::detail::Scheduler::getInstance
static Scheduler & getInstance()
Definition: scheduler.cpp:215
cl::sycl::detail::stream_impl::flush
void flush()
Definition: stream_impl.cpp:71
cl::sycl::accessor< char, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t >
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
queue.hpp
cl::sycl::handler
Command group handler class.
Definition: handler.hpp:362
cl::sycl::detail::Scheduler::StreamBuffersPool
std::unordered_map< stream_impl *, StreamBuffers * > StreamBuffersPool
Definition: scheduler.hpp:824
cl::sycl::detail::stream_impl::stream_impl
stream_impl(size_t BufferSize, size_t MaxStatementSize, handler &CGH)
Definition: stream_impl.cpp:19
cl::sycl::access::mode::read_write
@ read_write
cl::sycl::detail::stream_impl::get_size
size_t get_size() const
Definition: stream_impl.cpp:67
cl::sycl::ext::oneapi::experimental::printf
int printf(const FormatT *__format, Args... args)
Definition: builtins.hpp:76
stream_impl.hpp
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12