Stream class fe

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -11398,6 +11398,8 @@ def err_sycl_mismatch_group_size
             "have a sub group size that matches the size specified for the "
             "kernel">;
 def note_sycl_kernel_declared_here : Note<"kernel declared here">;
+def err_sycl_expected_finalize_method : Error<
+  "expected a 'finalize' method for the 'stream' class">;
 
 // errors of expect.with.probability
 def err_probability_not_constant_float : Error<

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

@@ -1034,23 +1034,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,
@@ -1125,12 +1108,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...);
@@ -1244,12 +1224,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;
@@ -1697,18 +1671,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;
@@ -1956,14 +1918,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;
@@ -2099,8 +2053,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 &,
@@ -2805,6 +2758,13 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
 
     const auto *RecordDecl = Ty->getAsCXXRecordDecl();
     createSpecialMethodCall(RecordDecl, getInitMethodName(), BodyStmts);
+    CXXMethodDecl *FinalizeMethod =
+        getMethodByName(RecordDecl, FinalizeMethodName);
+    // A finalize-method is expected for stream classes.
+    if (!FinalizeMethod && Util::isSyclStreamType(Ty))
+      SemaRef.Diag(FD->getLocation(), diag::err_sycl_expected_finalize_method);
+    else
+      createSpecialMethodCall(RecordDecl, FinalizeMethodName, FinalizeStmts);
 
     removeFieldMemberExpr(FD, Ty);
 
@@ -2898,9 +2858,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,
@@ -2977,31 +2935,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());
@@ -3315,7 +3248,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;
   }
 
@@ -3347,18 +3280,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);
@@ -4160,6 +4081,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"
 

clang/test/SemaSYCL/Inputs/sycl.hpp

@@ -114,6 +114,7 @@ class accessor {
   using PtrType = typename DeviceValueType<dataT, accessTarget>::type *;
   void __init(PtrType Ptr, range<dimensions> AccessRange,
               range<dimensions> MemRange, id<dimensions> Offset) {}
+  friend class stream;
 };
 
 template <int dimensions, access::mode accessmode, access::target accesstarget>
@@ -291,11 +292,24 @@ 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(__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 __init() {}
   void use() const {}
 
   void __finalize() {}
+
+private:
+  cl::sycl::accessor<char, 1, cl::sycl::access::mode::read_write> Acc;
+  int FlushBufferSize;
 };
 
 namespace ONEAPI {

clang/test/SemaSYCL/decomposition.cpp

@@ -131,13 +131,12 @@ int main() {
     myQueue.submit([&](sycl::handler &h) {
       h.single_task<class Stream1>([=]() { return t1.i; });
     });
-    // CHECK: FunctionDecl {{.*}}Stream1{{.*}} 'void (sycl::stream, __global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>, sycl::stream, __global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>, sycl::stream, __global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>, StructNonDecomposed, int)'
-
+    // CHECK: FunctionDecl {{.*}}Stream1{{.*}} 'void (__global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, int, __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, int, __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, int, StructNonDecomposed, int)'
     DerivedStruct<StructWithStream> t2;
     myQueue.submit([&](sycl::handler &h) {
       h.single_task<class Stream2>([=]() { return t2.i; });
     });
-    // CHECK: FunctionDecl {{.*}}Stream2{{.*}} 'void (sycl::stream, __global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>, StructNonDecomposed, int)'
+    // CHECK: FunctionDecl {{.*}}Stream2{{.*}} 'void (__global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, int, StructNonDecomposed, int)'
   }
 
   {

clang/test/SemaSYCL/kernel-arg-opt-report.cpp

@@ -221,10 +221,10 @@ int main() {
 // SPIR-NEXT: String:          '  ('
 // SPIR-NEXT: String:          ''
 // SPIR-NEXT: String:          'Type:'
-// SPIR-NEXT: String:          'sycl::stream'
+// SPIR-NEXT: String:          '__global char *'
 // SPIR-NEXT: String:          ', '
 // SPIR-NEXT: String:          'Size: '
-// SPIR-NEXT: Argument:        '3'
+// SPIR-NEXT: Argument:        '8'
 // SPIR-NEXT: String:          ')'
 
 // SPIR: --- !Passed
@@ -237,15 +237,15 @@ int main() {
 // SPIR-NEXT: String:          'Arg '
 // SPIR-NEXT: Argument:        '8'
 // SPIR-NEXT: String:          ':'
-// SPIR-NEXT: String:          Compiler generated argument for accessor,
-// SPIR-NEXT: String:          acc
+// SPIR-NEXT: String:          Compiler generated argument for stream,
+// SPIR-NEXT: String:          DecompStream
 // SPIR-NEXT: String:          '  ('
 // SPIR-NEXT: String:          ''
 // SPIR-NEXT: String:          'Type:'
-// SPIR-NEXT: String:          '__global int *'
+// SPIR-NEXT: String:          'struct sycl::range<1>'
 // SPIR-NEXT: String:          ', '
 // SPIR-NEXT: String:          'Size: '
-// SPIR-NEXT: Argument:        '8'
+// SPIR-NEXT: Argument:        '1'
 // SPIR-NEXT: String:          ')'
 
 // SPIR: --- !Passed
@@ -258,8 +258,8 @@ int main() {
 // SPIR-NEXT: String:          'Arg '
 // SPIR-NEXT: Argument:        '9'
 // SPIR-NEXT: String:          ':'
-// SPIR-NEXT: String:          Compiler generated argument for accessor,
-// SPIR-NEXT: String:          acc
+// SPIR-NEXT: String:          Compiler generated argument for stream,
+// SPIR-NEXT: String:          DecompStream
 // SPIR-NEXT: String:          '  ('
 // SPIR-NEXT: String:          ''
 // SPIR-NEXT: String:          'Type:'
@@ -279,12 +279,12 @@ int main() {
 // SPIR-NEXT: String:          'Arg '
 // SPIR-NEXT: Argument:        '10'
 // SPIR-NEXT: String:          ':'
-// SPIR-NEXT: String:          Compiler generated argument for accessor,
-// SPIR-NEXT: String:          acc
+// SPIR-NEXT: String:          Compiler generated argument for stream,
+// SPIR-NEXT: String:          DecompStream
 // SPIR-NEXT: String:          '  ('
 // SPIR-NEXT: String:          ''
 // SPIR-NEXT: String:          'Type:'
-// SPIR-NEXT: String:          'struct sycl::range<1>'
+// SPIR-NEXT: String:          'struct sycl::id<1>'
 // SPIR-NEXT: String:          ', '
 // SPIR-NEXT: String:          'Size: '
 // SPIR-NEXT: Argument:        '1'
@@ -300,15 +300,15 @@ int main() {
 // SPIR-NEXT: String:          'Arg '
 // SPIR-NEXT: Argument:        '11'
 // SPIR-NEXT: String:          ':'
-// SPIR-NEXT: String:          Compiler generated argument for accessor,
-// SPIR-NEXT: String:          acc
+// SPIR-NEXT: String:          Compiler generated argument for stream,
+// SPIR-NEXT: String:          DecompStream
 // SPIR-NEXT: String:          '  ('
 // SPIR-NEXT: String:          ''
 // SPIR-NEXT: String:          'Type:'
-// SPIR-NEXT: String:          'struct sycl::id<1>'
+// SPIR-NEXT: String:          int
 // SPIR-NEXT: String:          ', '
 // SPIR-NEXT: String:          'Size: '
-// SPIR-NEXT: Argument:        '1'
+// SPIR-NEXT: Argument:        '4'
 // SPIR-NEXT: String:          ')'
 
 // SPIR: --- !Passed