[SYCL] Refactor stream class handing implementation

clang/include/clang/Sema/Sema.h

@@ -316,7 +316,8 @@ class SYCLIntegrationHeader {
     kind_sampler,
     kind_pointer,
     kind_specialization_constants_buffer,
-    kind_last = kind_specialization_constants_buffer
+    kind_stream,
+    kind_last = kind_stream
   };
 
 public:

clang/lib/Sema/SemaSYCL.cpp

@@ -1025,23 +1025,6 @@ class KernelObjVisitor {
     VisitRecordFields(Owner, Handlers...);
   }
 
-  // FIXME: Can this be refactored/handled some other way?
-  template <typename ParentTy, typename... HandlerTys>
-  void visitStreamRecord(const CXXRecordDecl *Owner, ParentTy &Parent,
-                         CXXRecordDecl *Wrapper, QualType RecordTy,
-                         HandlerTys &... Handlers) {
-    (void)std::initializer_list<int>{
-        (Handlers.enterStream(Owner, Parent, RecordTy), 0)...};
-    for (const auto &Field : Wrapper->fields()) {
-      QualType FieldTy = Field->getType();
-      // Required to initialize accessors inside streams.
-      if (Util::isSyclAccessorType(FieldTy))
-        KF_FOR_EACH(handleSyclAccessorType, Field, FieldTy);
-    }
-    (void)std::initializer_list<int>{
-        (Handlers.leaveStream(Owner, Parent, RecordTy), 0)...};
-  }
-
   template <typename... HandlerTys>
   void visitArrayElementImpl(const CXXRecordDecl *Owner, FieldDecl *ArrayField,
                              QualType ElementTy, uint64_t Index,
@@ -1116,12 +1099,9 @@ class KernelObjVisitor {
       KF_FOR_EACH(handleSyclHalfType, Field, FieldTy);
     else if (Util::isSyclSpecConstantType(FieldTy))
       KF_FOR_EACH(handleSyclSpecConstantType, Field, FieldTy);
-    else if (Util::isSyclStreamType(FieldTy)) {
-      CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl();
-      // Handle accessors in stream class.
+    else if (Util::isSyclStreamType(FieldTy))
       KF_FOR_EACH(handleSyclStreamType, Field, FieldTy);
-      visitStreamRecord(Owner, Field, RD, FieldTy, Handlers...);
-    } else if (FieldTy->isStructureOrClassType()) {
+    else if (FieldTy->isStructureOrClassType()) {
       if (KF_FOR_EACH(handleStructType, Field, FieldTy)) {
         CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl();
         visitRecord(Owner, Field, RD, FieldTy, Handlers...);
@@ -1235,12 +1215,6 @@ class SyclKernelFieldHandlerBase {
   virtual bool leaveStruct(const CXXRecordDecl *, FieldDecl *, QualType) {
     return true;
   }
-  virtual bool enterStream(const CXXRecordDecl *, FieldDecl *, QualType) {
-    return true;
-  }
-  virtual bool leaveStream(const CXXRecordDecl *, FieldDecl *, QualType) {
-    return true;
-  }
   virtual bool enterStruct(const CXXRecordDecl *, const CXXBaseSpecifier &,
                            QualType) {
     return true;
@@ -1688,18 +1662,6 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler {
     return true;
   }
 
-  // Stream is always decomposed (and whether it gets decomposed is handled in
-  // handleSyclStreamType), but we need a CollectionStack entry to capture the
-  // accessors that get handled.
-  bool enterStream(const CXXRecordDecl *, FieldDecl *, QualType) final {
-    CollectionStack.push_back(false);
-    return true;
-  }
-  bool leaveStream(const CXXRecordDecl *, FieldDecl *, QualType Ty) final {
-    CollectionStack.pop_back();
-    return true;
-  }
-
   bool enterStruct(const CXXRecordDecl *, FieldDecl *, QualType) final {
     CollectionStack.push_back(false);
     return true;
@@ -1952,14 +1914,6 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
     SemaRef.addSyclDeviceDecl(KernelDecl);
   }
 
-  bool enterStream(const CXXRecordDecl *RD, FieldDecl *FD, QualType Ty) final {
-    return enterStruct(RD, FD, Ty);
-  }
-
-  bool leaveStream(const CXXRecordDecl *RD, FieldDecl *FD, QualType Ty) final {
-    return leaveStruct(RD, FD, Ty);
-  }
-
   bool enterStruct(const CXXRecordDecl *, FieldDecl *, QualType) final {
     ++StructDepth;
     return true;
@@ -2095,8 +2049,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
   }
 
   bool handleSyclStreamType(FieldDecl *FD, QualType FieldTy) final {
-    addParam(FD, FieldTy);
-    return true;
+    return handleSpecialType(FD, FieldTy);
   }
 
   bool handleSyclStreamType(const CXXRecordDecl *, const CXXBaseSpecifier &,
@@ -2591,6 +2544,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
 
     const auto *RecordDecl = Ty->getAsCXXRecordDecl();
     createSpecialMethodCall(RecordDecl, getInitMethodName(), BodyStmts);
+    createSpecialMethodCall(RecordDecl, FinalizeMethodName, BodyStmts);
 
     removeFieldMemberExpr(FD, Ty);
 
@@ -2684,9 +2638,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
   }
 
   bool handleSyclStreamType(FieldDecl *FD, QualType Ty) final {
-    // Streams just get copied as a new init.
-    addSimpleFieldInit(FD, Ty);
-    return true;
+    return handleSpecialType(FD, Ty);
   }
 
   bool handleSyclStreamType(const CXXRecordDecl *, const CXXBaseSpecifier &BS,
@@ -2763,31 +2715,6 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
       handleSpecialType(KernelHandlerArg->getType());
   }
 
-  bool enterStream(const CXXRecordDecl *RD, FieldDecl *FD, QualType Ty) final {
-    ++StructDepth;
-    // Add a dummy init expression to catch the accessor initializers.
-    const auto *StreamDecl = Ty->getAsCXXRecordDecl();
-    CollectionInitExprs.push_back(createInitListExpr(StreamDecl));
-
-    addFieldMemberExpr(FD, Ty);
-    return true;
-  }
-
-  bool leaveStream(const CXXRecordDecl *RD, FieldDecl *FD, QualType Ty) final {
-    --StructDepth;
-    // Stream requires that its 'init' calls happen after its accessors init
-    // calls, so add them here instead.
-    const auto *StreamDecl = Ty->getAsCXXRecordDecl();
-
-    createSpecialMethodCall(StreamDecl, getInitMethodName(), BodyStmts);
-    createSpecialMethodCall(StreamDecl, FinalizeMethodName, FinalizeStmts);
-
-    removeFieldMemberExpr(FD, Ty);
-
-    CollectionInitExprs.pop_back();
-    return true;
-  }
-
   bool enterStruct(const CXXRecordDecl *RD, FieldDecl *FD, QualType Ty) final {
     ++StructDepth;
     addCollectionInitListExpr(Ty->getAsCXXRecordDecl());
@@ -3101,7 +3028,7 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
   }
 
   bool handleSyclStreamType(FieldDecl *FD, QualType FieldTy) final {
-    addParam(FD, FieldTy, SYCLIntegrationHeader::kind_std_layout);
+    addParam(FD, FieldTy, SYCLIntegrationHeader::kind_stream);
     return true;
   }
 
@@ -3133,18 +3060,6 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
              SYCLIntegrationHeader::kind_specialization_constants_buffer, 0);
   }
 
-  bool enterStream(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final {
-    ++StructDepth;
-    CurOffset += offsetOf(FD, Ty);
-    return true;
-  }
-
-  bool leaveStream(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final {
-    --StructDepth;
-    CurOffset -= offsetOf(FD, Ty);
-    return true;
-  }
-
   bool enterStruct(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final {
     ++StructDepth;
     CurOffset += offsetOf(FD, Ty);
@@ -3948,6 +3863,7 @@ static const char *paramKind2Str(KernelParamKind K) {
     CASE(accessor);
     CASE(std_layout);
     CASE(sampler);
+    CASE(stream);
     CASE(specialization_constants_buffer);
     CASE(pointer);
   }

clang/test/CodeGenSYCL/Inputs/sycl.hpp

@@ -181,6 +181,7 @@ class accessor {
   void __init(__attribute__((opencl_global)) dataT *Ptr, range<dimensions> AccessRange,
               range<dimensions> MemRange, id<dimensions> Offset) {}
   void __init_esimd(__attribute__((opencl_global)) dataT *Ptr) {}
+  friend class stream;
 };
 
 template <int dimensions, access::mode accessmode, access::target accesstarget>
@@ -411,10 +412,22 @@ class stream {
 public:
   stream(unsigned long BufferSize, unsigned long MaxStatementSize,
          handler &CGH) {}
+#ifdef __SYCL_DEVICE_ONLY__
+  // Default constructor for objects later initialized with __init member.
+  stream() = default;
+#endif
 
-  void __init() {}
+  void __init(__attribute((opencl_global)) char *Ptr, range<1> AccessRange,
+              range<1> MemRange, id<1> Offset, int _FlushBufferSize) {
+    Acc.__init(Ptr, AccessRange, MemRange, Offset);
+    FlushBufferSize = _FlushBufferSize;
+  }
 
   void __finalize() {}
+
+private:
+  cl::sycl::accessor<char, 1, cl::sycl::access::mode::read_write> Acc;
+  int FlushBufferSize;
 };
 
 template <typename T>

clang/test/CodeGenSYCL/stream.cpp

@@ -1,10 +1,18 @@
 // RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o %t.ll
 // RUN: FileCheck < %t.ll --enable-var-scope %s
 //
-// CHECK: define {{.*}}spir_kernel void @"{{.*}}StreamTester"(%"{{.*}}cl::sycl::stream"* byval(%"{{.*}}cl::sycl::stream") {{.*}}){{.*}}
-// CHECK: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::stream{{.*}} addrspace(4)* {{[^,]*}} %{{[0-9]+}})
-// CHECK: call spir_func void @{{.*}}__finalize{{.*}}(%{{.*}}cl::sycl::stream{{.*}} addrspace(4)* {{[^,]*}} %{{[0-9]+}})
-//
+
+// CHECK: %[[RANGE_TYPE:"struct.*cl::sycl::range"]]
+// CHECK: %[[ID_TYPE:"struct.*cl::sycl::id"]]
+// CHECK: define dso_local spir_kernel void @{{.*}}StreamTester
+// CHECK-SAME: i8 addrspace(1)* [[ACC_DATA:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: %[[RANGE_TYPE]]* byval(%[[RANGE_TYPE]]) align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: %[[RANGE_TYPE]]* byval(%[[RANGE_TYPE]]) align 4 [[ACC_RANGE2:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: %[[ID_TYPE]]* byval(%[[ID_TYPE]]) align 4 [[ACC_ID:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: i32 [[ACC_INT:%[a-zA-Z0-9_]+]])
+
+// CHECK: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::stream{{.*}} addrspace(4)* dereferenceable_or_null(16) %{{[0-9]+}}, i8 addrspace(1)* %5, %[[RANGE_TYPE]]* byval(%[[RANGE_TYPE]]) {{.*}}%{{.*}}
+// CHECK: call spir_func void @{{.*}}__finalizeEv{{.*}}(%{{.*}}cl::sycl::stream{{.*}}" addrspace(4)* dereferenceable_or_null(16) %{{[0-9]+}})
 
 #include "Inputs/sycl.hpp"
 

sycl/include/CL/sycl/accessor.hpp

@@ -203,6 +203,7 @@
 
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
+class stream;
 namespace INTEL {
 namespace gpu {
 namespace detail {
@@ -929,6 +930,7 @@ class accessor :
 
 private:
   friend class sycl::INTEL::gpu::detail::AccessorPrivateProxy;
+  friend class sycl::stream;
 
 public:
   using value_type = DataT;

sycl/include/CL/sycl/detail/kernel_desc.hpp

@@ -33,8 +33,9 @@ enum class kernel_param_kind_t {
   kind_accessor = 0,
   kind_std_layout = 1, // standard layout object parameters
   kind_sampler = 2,
-  kind_pointer = 3,
-  kind_specialization_constants_buffer = 4,
+  kind_stream = 3,
+  kind_pointer = 4,
+  kind_specialization_constants_buffer = 5,
 };
 
 // describes a kernel parameter

sycl/include/CL/sycl/stream.hpp

@@ -68,11 +68,23 @@ using GlobalBufAccessorT = accessor<char, 1, cl::sycl::access::mode::read_write,
                                     cl::sycl::access::target::global_buffer,
                                     cl::sycl::access::placeholder::false_t>;
 
+constexpr static access::address_space GlobalBufAS =
+    TargetToAS<cl::sycl::access::target::global_buffer>::AS;
+using GlobalBufPtrType =
+    typename detail::DecoratedType<char, GlobalBufAS>::type *;
+constexpr static int GlobalBufDim = 1;
+
 using GlobalOffsetAccessorT =
     accessor<unsigned, 1, cl::sycl::access::mode::atomic,
              cl::sycl::access::target::global_buffer,
              cl::sycl::access::placeholder::false_t>;
 
+constexpr static access::address_space GlobalOffsetAS =
+    TargetToAS<cl::sycl::access::target::global_buffer>::AS;
+using GlobalOffsetPtrType =
+    typename detail::DecoratedType<unsigned, GlobalBufAS>::type *;
+constexpr static int GlobalOffsetDim = 1;
+
 // 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]);
@@ -729,6 +741,11 @@ inline __width_manipulator__ setw(int Width) {
 /// \ingroup sycl_api
 class __SYCL_EXPORT stream {
 public:
+#ifdef __SYCL_DEVICE_ONLY__
+  // Default constructor for objects later initialized with __init member.
+  stream() = default;
+#endif
+
   // Throws exception in case of invalid input parameters
   stream(size_t BufferSize, size_t MaxStatementSize, handler &CGH);
 
@@ -845,7 +862,27 @@ class __SYCL_EXPORT stream {
   }
 
 #ifdef __SYCL_DEVICE_ONLY__
-  void __init() {
+  void __init(detail::GlobalBufPtrType GlobalBufPtr,
+              range<detail::GlobalBufDim> GlobalBufAccRange,
+              range<detail::GlobalBufDim> GlobalBufMemRange,
+              id<detail::GlobalBufDim> GlobalBufId,
+              detail::GlobalOffsetPtrType GlobalOffsetPtr,
+              range<detail::GlobalOffsetDim> GlobalOffsetAccRange,
+              range<detail::GlobalOffsetDim> GlobalOffsetMemRange,
+              id<detail::GlobalOffsetDim> GlobalOffsetId,
+              detail::GlobalBufPtrType GlobalFlushPtr,
+              range<detail::GlobalBufDim> GlobalFlushAccRange,
+              range<detail::GlobalBufDim> GlobalFlushMemRange,
+              id<detail::GlobalBufDim> GlobalFlushId, size_t _FlushBufferSize) {
+#ifndef __SYCL_EXPLICIT_SIMD__
+    GlobalBuf.__init(GlobalBufPtr, GlobalBufAccRange, GlobalBufMemRange,
+                     GlobalBufId);
+    GlobalOffset.__init(GlobalOffsetPtr, GlobalOffsetAccRange,
+                        GlobalOffsetMemRange, GlobalOffsetId);
+    GlobalFlushBuf.__init(GlobalFlushPtr, GlobalFlushAccRange,
+                          GlobalFlushMemRange, GlobalFlushId);
+#endif
+    FlushBufferSize = _FlushBufferSize;
     // Calculate offset in the flush buffer for each work item in the global
     // work space. We need to avoid calling intrinsics to get global id because
     // when stream is used in a single_task kernel this could cause some
@@ -871,6 +908,8 @@ class __SYCL_EXPORT stream {
   }
 #endif
 
+  friend class handler;
+
   friend const stream &operator<<(const stream &, const char);
   friend const stream &operator<<(const stream &, const char *);
   template <typename ValueType>
@@ -1118,4 +1157,3 @@ template <> struct hash<cl::sycl::stream> {
   }
 };
 } // namespace std
-

sycl/source/detail/scheduler/commands.cpp

@@ -1660,6 +1660,8 @@ pi_result ExecCGCommand::SetKernelParamsAndLaunch(
     if (!EliminatedArgMask.empty() && EliminatedArgMask[Arg.MIndex])
       continue;
     switch (Arg.MType) {
+    case kernel_param_kind_t::kind_stream:
+      break;
     case kernel_param_kind_t::kind_accessor: {
       Requirement *Req = (Requirement *)(Arg.MPtr);
       AllocaCommandBase *AllocaCmd = getAllocaForReq(Req);