[SYCL] Stream class

clang/lib/Sema/SemaSYCL.cpp

@@ -67,9 +67,6 @@ class Util {
   /// sampler class.
   static bool isSyclSamplerType(const QualType &Ty);
 
-  /// Checks whether given clang type is the SYCL stream class.
-  static bool isSyclStreamType(const QualType &Ty);
-
   /// Checks whether given clang type is declared in the given hierarchy of
   /// declaration contexts.
   /// \param Ty         the clang type being checked
@@ -814,9 +811,6 @@ static void buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj,
       assert(SamplerArg && "sampler __init method must have sampler parameter");
 
       CreateAndAddPrmDsc(Fld, SamplerArg->getType());
-    } else if (Util::isSyclStreamType(ArgTy)) {
-      // the parameter is a SYCL stream object
-      llvm_unreachable("streams not supported yet");
     } else if (ArgTy->isStructureOrClassType()) {
       if (!ArgTy->isStandardLayoutType()) {
         const DeclaratorDecl *V =
@@ -917,9 +911,6 @@ static void populateIntHeader(SYCLIntegrationHeader &H, const StringRef Name,
       uint64_t Sz = Ctx.getTypeSizeInChars(SamplerArg->getType()).getQuantity();
       H.addParamDesc(SYCLIntegrationHeader::kind_sampler,
                      static_cast<unsigned>(Sz), static_cast<unsigned>(Offset));
-    } else if (Util::isSyclStreamType(ArgTy)) {
-      // the parameter is a SYCL stream object
-      llvm_unreachable("streams not supported yet");
     } else if (ArgTy->isStructureOrClassType() || ArgTy->isScalarType()) {
       // the parameter is an object of standard layout type or scalar;
       // the check for standard layout is done elsewhere
@@ -1369,14 +1360,6 @@ bool Util::isSyclSamplerType(const QualType &Ty) {
   return matchQualifiedTypeName(Ty, Scopes);
 }
 
-bool Util::isSyclStreamType(const QualType &Ty) {
-  static std::array<DeclContextDesc, 3> Scopes = {
-      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "cl"},
-      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "sycl"},
-      Util::DeclContextDesc{clang::Decl::Kind::CXXRecord, "stream"}};
-  return matchQualifiedTypeName(Ty, Scopes);
-}
-
 bool Util::matchQualifiedTypeName(const QualType &Ty,
                                   ArrayRef<Util::DeclContextDesc> Scopes) {
   // The idea: check the declaration context chain starting from the type

sycl/CMakeLists.txt

@@ -142,6 +142,7 @@ add_library("${SYCLLibrary}" SHARED
   "${sourceRootPath}/detail/os_util.cpp"
   "${sourceRootPath}/detail/platform_util.cpp"
   "${sourceRootPath}/detail/sampler_impl.cpp"
+  "${sourceRootPath}/detail/stream_impl.cpp"
   "${sourceRootPath}/detail/scheduler/commands.cpp"
   "${sourceRootPath}/detail/scheduler/scheduler.cpp"
   "${sourceRootPath}/detail/scheduler/graph_processor.cpp"
@@ -157,6 +158,7 @@ add_library("${SYCLLibrary}" SHARED
   "${sourceRootPath}/platform.cpp"
   "${sourceRootPath}/queue.cpp"
   "${sourceRootPath}/sampler.cpp"
+  "${sourceRootPath}/stream.cpp"
   "${sourceRootPath}/spirv_ops.cpp"
 )
 

sycl/include/CL/sycl.hpp

@@ -34,5 +34,6 @@
 #include <CL/sycl/queue.hpp>
 #include <CL/sycl/range.hpp>
 #include <CL/sycl/sampler.hpp>
+#include <CL/sycl/stream.hpp>
 #include <CL/sycl/types.hpp>
 #include <CL/sycl/version.hpp>

sycl/include/CL/sycl/accessor.hpp

@@ -514,28 +514,28 @@ class accessor :
     return getQualifiedPtr()[LinearIndex];
   }
 
-  template <
-      int Dims = Dimensions,
-      typename = enable_if_t<AccessMode == access::mode::atomic && Dims == 0>>
-  operator atomic<DataT, AS>() const {
+  template <int Dims = Dimensions>
+  operator typename std::enable_if<AccessMode == access::mode::atomic &&
+                                       Dims == 0,
+                                   atomic<DataT, AS>>::type() const {
     const size_t LinearIndex = getLinearIndex(id<AdjustedDim>());
     return atomic<DataT, AS>(
         multi_ptr<DataT, AS>(getQualifiedPtr() + LinearIndex));
   }
 
-  template <
-      int Dims = Dimensions,
-      typename = enable_if_t<AccessMode == access::mode::atomic && (Dims > 0)>>
-  atomic<DataT, AS> operator[](id<Dimensions> Index) const {
+  template <int Dims = Dimensions>
+  typename std::enable_if<AccessMode == access::mode::atomic && (Dims > 0),
+                          atomic<DataT, AS>>::type
+  operator[](id<Dimensions> Index) const {
     const size_t LinearIndex = getLinearIndex(Index);
     return atomic<DataT, AS>(
         multi_ptr<DataT, AS>(getQualifiedPtr() + LinearIndex));
   }
 
-  template <
-      int Dims = Dimensions,
-      typename = enable_if_t<AccessMode == access::mode::atomic && Dims == 1>>
-  atomic<DataT, AS> operator[](size_t Index) const {
+  template <int Dims = Dimensions>
+  typename enable_if_t<AccessMode == access::mode::atomic && Dims == 1,
+                       atomic<DataT, AS>>::type
+  operator[](size_t Index) const {
     const size_t LinearIndex = getLinearIndex(id<AdjustedDim>(Index));
     return atomic<DataT, AS>(
         multi_ptr<DataT, AS>(getQualifiedPtr() + LinearIndex));

sycl/include/CL/sycl/buffer.hpp

@@ -158,7 +158,7 @@ class buffer {
 
   size_t get_count() const { return Range.size(); }
 
-  size_t get_size() const { return impl->get_size(); }
+  size_t get_size() const { return get_count() * sizeof(T); }
 
   AllocatorT get_allocator() const { return impl->get_allocator(); }
 

sycl/include/CL/sycl/detail/cg.hpp

@@ -220,6 +220,7 @@ class HostKernel : public HostKernelBase {
   ~HostKernel() = default;
 };
 
+class stream_impl;
 // The base class for all types of command groups.
 class CG {
 public:
@@ -272,6 +273,7 @@ class CGExecKernel : public CG {
   std::vector<ArgDesc> MArgs;
   std::string MKernelName;
   detail::OSModuleHandle MOSModuleHandle;
+  std::vector<std::shared_ptr<detail::stream_impl>> MStreams;
 
   CGExecKernel(NDRDescT NDRDesc, std::unique_ptr<HostKernelBase> HKernel,
                std::shared_ptr<detail::kernel_impl> SyclKernel,
@@ -280,15 +282,20 @@ class CGExecKernel : public CG {
                std::vector<std::shared_ptr<void>> SharedPtrStorage,
                std::vector<Requirement *> Requirements,
                std::vector<ArgDesc> Args, std::string KernelName,
-               detail::OSModuleHandle OSModuleHandle)
+               detail::OSModuleHandle OSModuleHandle,
+               std::vector<std::shared_ptr<detail::stream_impl>> Streams)
       : CG(KERNEL, std::move(ArgsStorage), std::move(AccStorage),
            std::move(SharedPtrStorage), std::move(Requirements)),
         MNDRDesc(std::move(NDRDesc)), MHostKernel(std::move(HKernel)),
         MSyclKernel(std::move(SyclKernel)), MArgs(std::move(Args)),
-        MKernelName(std::move(KernelName)), MOSModuleHandle(OSModuleHandle) {}
+        MKernelName(std::move(KernelName)), MOSModuleHandle(OSModuleHandle),
+        MStreams(std::move(Streams)) {}
 
   std::vector<ArgDesc> getArguments() const { return MArgs; }
   std::string getKernelName() const { return MKernelName; }
+  std::vector<std::shared_ptr<detail::stream_impl>> getStreams() const {
+    return MStreams;
+  }
 };
 
 // The class which represents "copy" command group.

sycl/include/CL/sycl/detail/scheduler/commands.hpp

@@ -229,6 +229,8 @@ class ExecCGCommand : public Command {
       : Command(CommandType::RUN_CG, std::move(Queue)),
         MCommandGroup(std::move(CommandGroup)) {}
 
+  void flushStreams();
+
 private:
   // Implementation of enqueueing of ExecCGCommand.
   cl_int enqueueImp() override;

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

@@ -0,0 +1,75 @@
+//==----------------- stream_impl.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/accessor.hpp>
+#include <CL/sycl/device_selector.hpp>
+#include <CL/sycl/queue.hpp>
+
+namespace cl {
+namespace sycl {
+
+namespace detail {
+class stream_impl {
+public:
+  using AccessorType = accessor<char, 1, cl::sycl::access::mode::read_write,
+                                cl::sycl::access::target::global_buffer,
+                                cl::sycl::access::placeholder::false_t>;
+
+  using OffsetAccessorType =
+      accessor<unsigned, 1, cl::sycl::access::mode::atomic,
+               cl::sycl::access::target::global_buffer,
+               cl::sycl::access::placeholder::false_t>;
+
+  stream_impl(size_t BufferSize, size_t MaxStatementSize, handler &CGH);
+
+  // Method to provide an access to the stream buffer
+  AccessorType getAccessor(handler &CGH) {
+    return Buf.get_access<cl::sycl::access::mode::read_write>(
+        CGH, range<1>(BufferSize_), id<1>(OffsetSize));
+  }
+
+  // Method to provide an atomic access to the offset in the stream buffer
+  OffsetAccessorType getOffsetAccessor(handler &CGH) {
+    auto OffsetSubBuf = buffer<char, 1>(Buf, id<1>(0), range<1>(OffsetSize));
+    auto ReinterpretedBuf = OffsetSubBuf.reinterpret<unsigned, 1>(range<1>(1));
+    return ReinterpretedBuf.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();
+
+  size_t get_size() const;
+
+  size_t get_max_statement_size() const;
+
+private:
+  // Size of the stream buffer
+  size_t BufferSize_;
+
+  // Maximum number of symbols which could be streamed from the beginning of a
+  // statement till the semicolon
+  size_t MaxStatementSize_;
+
+  // Size of the variable which is used as an offset in the stream buffer.
+  // Additinonal memory is allocated in the beginning of the stream buffer for
+  // this variable.
+  static const size_t OffsetSize = sizeof(unsigned);
+
+  // Vector on the host side which is used to initialize the stream buffer
+  std::vector<char> Data;
+
+  // Stream buffer
+  buffer<char, 1> Buf;
+};
+} // namespace detail
+} // namespace sycl
+} // namespace cl
+

sycl/include/CL/sycl/handler.hpp

@@ -109,6 +109,7 @@ DEFINE_INIT_SIZES(GlobalOffset)
 #endif //__SYCL_DEVICE_ONLY__
 
 class queue_impl;
+class stream_impl;
 template <typename RetType, typename Func, typename Arg>
 static Arg member_ptr_helper(RetType (Func::*)(Arg) const);
 
@@ -163,6 +164,7 @@ class handler {
   // we exit the method they are passed in.
   std::vector<std::vector<char>> MArgsStorage;
   std::vector<detail::AccessorImplPtr> MAccStorage;
+  std::vector<std::shared_ptr<detail::stream_impl>> MStreamStorage;
   std::vector<std::shared_ptr<void>> MSharedPtrStorage;
   // The list of arguments for the kernel.
   std::vector<detail::ArgDesc> MArgs;
@@ -360,8 +362,8 @@ class handler {
           std::move(MNDRDesc), std::move(MHostKernel), std::move(MSyclKernel),
           std::move(MArgsStorage), std::move(MAccStorage),
           std::move(MSharedPtrStorage), std::move(MRequirements),
-          std::move(MArgs), std::move(MKernelName),
-          std::move(MOSModuleHandle)));
+          std::move(MArgs), std::move(MKernelName), std::move(MOSModuleHandle),
+          std::move(MStreamStorage)));
       break;
     case detail::CG::COPY_ACC_TO_PTR:
     case detail::CG::COPY_PTR_TO_ACC:
@@ -393,6 +395,12 @@ class handler {
     return EventRet;
   }
 
+  // Save streams associated with this handler. Streams are then forwarded to
+  // command group and flushed in the scheduler.
+  void addStream(std::shared_ptr<detail::stream_impl> s) {
+    MStreamStorage.push_back(std::move(s));
+  }
+
   ~handler() = default;
 
   bool is_host() { return MIsHost; }
@@ -485,6 +493,8 @@ class handler {
   template <typename DataT, int Dims, access::mode AccMode,
             access::target AccTarget, access::placeholder isPlaceholder>
   friend class accessor;
+  // Make stream class friend to be able to keep the list of associated streams
+  friend class stream;
 
 public:
   handler(const handler &) = delete;