intel · Fznamznon · Aug 5, 2020
@@ -314,6 +314,7 @@ class SYCLIntegrationHeader {
     kind_accessor = kind_first,
     kind_std_layout,
     kind_sampler,
+    kind_stream,
     kind_pointer,
     kind_last = kind_pointer
   };

@@ -880,26 +880,6 @@ class KernelObjVisitor {
     VisitRecordFields(Owner, handlers...);
   }
 
-  // FIXME: Can this be refactored/handled some other way?
-  template <typename ParentTy, typename... Handlers>
-  void VisitStreamRecord(const CXXRecordDecl *Owner, ParentTy &Parent,
-                         CXXRecordDecl *Wrapper, Handlers &... handlers) {
-    (void)std::initializer_list<int>{
-        (handlers.enterStruct(Owner, Parent), 0)...};
-    for (const auto &Field : Wrapper->fields()) {
-      QualType FieldTy = Field->getType();
-      (void)std::initializer_list<int>{
-          (handlers.enterField(Wrapper, Field), 0)...};
-      // Required to initialize accessors inside streams.
-      if (Util::isSyclAccessorType(FieldTy))
-        KF_FOR_EACH(handleSyclAccessorType, Field, FieldTy);
-      (void)std::initializer_list<int>{
-          (handlers.leaveField(Wrapper, Field), 0)...};
-    }
-    (void)std::initializer_list<int>{
-        (handlers.leaveStruct(Owner, Parent), 0)...};
-  }
-
   template <typename... Handlers>
   void VisitRecordBases(const CXXRecordDecl *KernelFunctor,
                         Handlers &... handlers) {
@@ -924,12 +904,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.
-        VisitStreamRecord(Owner, Field, RD, handlers...);
+      else if (Util::isSyclStreamType(FieldTy))
         KF_FOR_EACH(handleSyclStreamType, Field, FieldTy);
-      } else if (FieldTy->isStructureOrClassType()) {
+      else if (FieldTy->isStructureOrClassType()) {
         if (KF_FOR_EACH(handleStructType, Field, FieldTy)) {
           CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl();
           VisitRecord(Owner, Field, RD, handlers...);
@@ -1297,8 +1274,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
   }
 
   bool handleSyclStreamType(FieldDecl *FD, QualType FieldTy) final {
-    addParam(FD, FieldTy);
-    return true;
+    return handleSpecialType(FD, FieldTy);
   }
 
   bool handleSyclStreamType(const CXXBaseSpecifier &, QualType FieldTy) final {
@@ -1515,6 +1491,13 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
           createSpecialMethodCall(MemberExprBases.back(), InitMethod, FD);
       BodyStmts.push_back(InitCall);
     }
+    CXXMethodDecl *FinalizeMethod =
+        getMethodByName(RecordDecl, FinalizeMethodName);
+    if (FinalizeMethod) {
+      CXXMemberCallExpr *FinalizeCall =
+          createSpecialMethodCall(MemberExprBases.back(), FinalizeMethod, FD);
+      FinalizeStmts.push_back(FinalizeCall);
+    }
     return true;
   }
 
@@ -1537,6 +1520,13 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
           createSpecialMethodCall(MemberExprBases.back(), InitMethod, nullptr);
       BodyStmts.push_back(InitCall);
     }
+    CXXMethodDecl *FinalizeMethod =
+        getMethodByName(RecordDecl, FinalizeMethodName);
+    if (FinalizeMethod) {
+      CXXMemberCallExpr *FinalizeCall = createSpecialMethodCall(
+          MemberExprBases.back(), FinalizeMethod, nullptr);
+      FinalizeStmts.push_back(FinalizeCall);
+    }
     return true;
   }
 
@@ -1583,23 +1573,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
   }
 
   bool handleSyclStreamType(FieldDecl *FD, QualType Ty) final {
-    const auto *StreamDecl = Ty->getAsCXXRecordDecl();
-    createExprForStructOrScalar(FD);
-    size_t NumBases = MemberExprBases.size();
-    CXXMethodDecl *InitMethod = getMethodByName(StreamDecl, InitMethodName);
-    if (InitMethod) {
-      CXXMemberCallExpr *InitCall =
-          createSpecialMethodCall(MemberExprBases.back(), InitMethod, FD);
-      BodyStmts.push_back(InitCall);
-    }
-    CXXMethodDecl *FinalizeMethod =
-        getMethodByName(StreamDecl, FinalizeMethodName);
-    if (FinalizeMethod) {
-      CXXMemberCallExpr *FinalizeCall = createSpecialMethodCall(
-          MemberExprBases[NumBases - 2], FinalizeMethod, FD);
-      FinalizeStmts.push_back(FinalizeCall);
-    }
-    return true;
+    return handleSpecialType(FD, Ty);
   }
 
   bool handleSyclStreamType(const CXXBaseSpecifier &BS, QualType Ty) final {
@@ -1666,18 +1640,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
     const CXXRecordDecl *RD =
         FD->getType()->getBaseElementTypeUnsafe()->getAsCXXRecordDecl();
 
-    // Initializers for accessors inside stream not added.
-    if (!Util::isSyclStreamType(FD->getType()))
-      addStructInit(RD);
-    // Pop out unused initializers created in handleSyclAccesorType
-    // for accessors inside stream class.
-    else {
-      for (const auto &Field : RD->fields()) {
-        QualType FieldTy = Field->getType();
-        if (Util::isSyclAccessorType(FieldTy))
-          InitExprs.pop_back();
-      }
-    }
+    addStructInit(RD);
     return true;
   }
 
@@ -1831,7 +1794,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;
   }
 
@@ -2211,6 +2174,7 @@ static const char *paramKind2Str(KernelParamKind K) {
     CASE(accessor);
     CASE(std_layout);
     CASE(sampler);
+    CASE(stream);
     CASE(pointer);
   default:
     return "<ERROR>";

@@ -141,6 +141,7 @@ class accessor {
 private:
   void __init(__attribute__((opencl_global)) dataT *Ptr, range<dimensions> AccessRange,
               range<dimensions> MemRange, id<dimensions> Offset) {}
+  friend class stream;
 };
 
 template <int dimensions, access::mode accessmode, access::target accesstarget>
@@ -314,10 +315,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>

@@ -1,8 +1,16 @@
 // RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -I %S/Inputs -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: %[[RANGE_TYPE:"struct.*cl::sycl::range"]]
+// CHECK: %[[ID_TYPE:"struct.*cl::sycl::id"]]
+// CHECK: define 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 [[ARG_INT:%[a-zA-Z0-9_]+]])
+
+// CHECK: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::stream{{.*}} addrspace(4)* %{{[0-9]+}}, i8 addrspace(1)* %5, %[[RANGE_TYPE]]* byval(%[[RANGE_TYPE]]) {{.*}} %{{.*}}
 // CHECK: call spir_func void @{{.*}}__finalize{{.*}}(%{{.*}}cl::sycl::stream{{.*}} addrspace(4)* %{{[0-9]+}})
 //
 

@@ -199,6 +199,7 @@
 
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
+class stream;
 namespace intel {
 namespace gpu {
 // Forward declare a "back-door" access class to support ESIMD.
@@ -886,6 +887,7 @@ class accessor :
 
 private:
   friend class sycl::intel::gpu::AccessorPrivateProxy;
+  friend class sycl::stream;
 
 public:
   using value_type = DataT;

@@ -27,6 +27,7 @@ enum class kernel_param_kind_t {
   kind_accessor,
   kind_std_layout, // standard layout object parameters
   kind_sampler,
+  kind_stream,
   kind_pointer
 };
 

@@ -63,11 +63,24 @@ 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::PtrValueType<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::PtrValueType<unsigned, GlobalOffsetAS>::type *;
+constexpr static int GlobalOffsetDim = 1;
+
+
 inline void write(GlobalBufAccessorT &GlobalFlushBuf, size_t FlushBufferSize,
                   unsigned WIOffset, unsigned &Offset, const char *Str,
                   unsigned Len, unsigned Padding = 0) {
@@ -697,6 +710,12 @@ 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
+
   stream(size_t BufferSize, size_t MaxStatementSize, handler &CGH);
 
   size_t get_size() const;
@@ -810,7 +829,28 @@ 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 work item's global id, this should be done once, that
     // is why this is done in _init method, call to __init method is generated
     // by frontend. As a result each work item will write to its own section
@@ -834,6 +874,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>

@@ -1652,6 +1652,8 @@ pi_result ExecCGCommand::SetKernelParamsAndLaunch(
   const detail::plugin &Plugin = MQueue->getPlugin();
   for (ArgDesc &Arg : ExecKernel->MArgs) {
     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);