diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 02786427801e..ae636aa9db5c 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -11412,6 +11412,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">;
 def ext_sycl_2020_attr_spelling : ExtWarn<
   "use of attribute %0 is a SYCL 2020 extension">,
   InGroup<Sycl2017Compat>;
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index 563d4a8c5a86..d754e47a738f 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/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:
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index c2c11765a8bc..d4851208775b 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/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 &,
@@ -2419,15 +2372,6 @@ class SyclOptReportCreator : public SyclKernelFieldHandler {
   }
 
   bool handleSyclStreamType(FieldDecl *FD, QualType FieldTy) final {
-    // For the current implementation of stream class, the Visitor 'handles'
-    // stream argument and then visits each accessor field in stream. Therefore
-    // handleSpecialType in this case only adds a single argument for stream.
-    // The arguments corresponding to accessors in stream are handled in
-    // handleSyclAccessorType. The opt-report therefore does not diffrentiate
-    // between the accessors in streams and accessors captured by SYCL kernel.
-    // Once stream API is modified to use __init(), the visitor will no longer
-    // visit the stream object and opt-report output for stream class will be
-    // similar to that of other special types.
     return handleSpecialType(
         FD, FieldTy, KernelArgDescription(KernelArgDescription::Stream));
   }
@@ -2805,6 +2749,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 class.
+    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 +2849,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 +2926,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 +3239,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 +3271,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 +4072,7 @@ static const char *paramKind2Str(KernelParamKind K) {
     CASE(accessor);
     CASE(std_layout);
     CASE(sampler);
+    CASE(stream);
     CASE(specialization_constants_buffer);
     CASE(pointer);
   }
diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
index d27f3c9e1a67..f709e2ed1d5f 100644
--- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp
+++ b/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>
diff --git a/clang/test/CodeGenSYCL/stream.cpp b/clang/test/CodeGenSYCL/stream.cpp
index cfc354efa068..eb57aa207bea 100644
--- a/clang/test/CodeGenSYCL/stream.cpp
+++ b/clang/test/CodeGenSYCL/stream.cpp
@@ -1,10 +1,24 @@
 // 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_]+]])
+
+// Alloca and addrspace casts for kernel parameters
+// CHECK: [[ARG:%[a-zA-Z0-9_]+]].addr = alloca i8 addrspace(1)*, align 8
+// CHECK: [[ARG:%[a-zA-Z0-9_]+]].addr.ascast = addrspacecast i8 addrspace(1)** [[ARG]].addr to i8 addrspace(1)* addrspace(4)*
+// CHECK: [[ARG_LOAD:%[a-zA-Z0-9_]+]] = load i8 addrspace(1)*, i8 addrspace(1)* addrspace(4)* [[ARG]].addr.ascast, align 8,
+
+// Check __init and __finalize method calls
+// CHECK:  call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::stream" addrspace(4)* align 4 dereferenceable_or_null(16) %4, i8 addrspace(1)* [[ARG_LOAD]], %[[RANGE_TYPE]]* byval(%[[RANGE_TYPE]]) {{.*}}%{{.*}}
+// CHECK:  call spir_func void @_ZN2cl4sycl6stream10__finalizeEv(%{{.*}}cl::sycl::stream" addrspace(4)* align 4 dereferenceable_or_null(16) %{{[0-9]+}})
 
 #include "Inputs/sycl.hpp"
 
diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp
index 4782bafa74ce..88dc82a4f576 100644
--- a/clang/test/SemaSYCL/Inputs/sycl.hpp
+++ b/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 {
diff --git a/clang/test/SemaSYCL/decomposition.cpp b/clang/test/SemaSYCL/decomposition.cpp
index a5e3bebca0c2..f56fde742b11 100644
--- a/clang/test/SemaSYCL/decomposition.cpp
+++ b/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)'
   }
 
   {
diff --git a/clang/test/SemaSYCL/kernel-arg-opt-report.cpp b/clang/test/SemaSYCL/kernel-arg-opt-report.cpp
index 967e269dace2..cf6b28cb4f9a 100644
--- a/clang/test/SemaSYCL/kernel-arg-opt-report.cpp
+++ b/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
diff --git a/clang/test/SemaSYCL/streams.cpp b/clang/test/SemaSYCL/stream.cpp
similarity index 59%
rename from clang/test/SemaSYCL/streams.cpp
rename to clang/test/SemaSYCL/stream.cpp
index 933dabff649f..606f77925e13 100644
--- a/clang/test/SemaSYCL/streams.cpp
+++ b/clang/test/SemaSYCL/stream.cpp
@@ -52,60 +52,39 @@ int main() {
 
 // CHECK: InitListExpr {{.*}} '(lambda at
 // 'in_lambda'
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void (const sycl::stream &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const sycl::stream' lvalue <NoOp>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue ParmVar
-
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
 // 'in_lambda_array'
 // CHECK-NEXT: InitListExpr {{.*}} 'sycl::stream [2]'
 // element 0
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void (const sycl::stream &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const sycl::stream' lvalue <NoOp>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue ParmVar
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
 // element 1
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void (const sycl::stream &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const sycl::stream' lvalue <NoOp>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue ParmVar
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
 
 // 'in_lambda_mdarray'
 // CHECK-NEXT: InitListExpr {{.*}} 'sycl::stream [2][2]'
 // sub-array 0
 // CHECK-NEXT: InitListExpr {{.*}} 'sycl::stream [2]'
 // element 0
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void (const sycl::stream &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const sycl::stream' lvalue <NoOp>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue ParmVar
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
 // element 1
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void (const sycl::stream &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const sycl::stream' lvalue <NoOp>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue ParmVar
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
 // sub-array 1
 // CHECK-NEXT: InitListExpr {{.*}} 'sycl::stream [2]'
 // element 0
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void (const sycl::stream &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const sycl::stream' lvalue <NoOp>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue ParmVar
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
 // element 1
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void (const sycl::stream &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const sycl::stream' lvalue <NoOp>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue ParmVar
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
 
 // HasStreams struct
-// CHECK-NEXT: InitListExpr {{.*}} 'HasStreams'
+// CHECK: InitListExpr {{.*}} 'HasStreams'
 // HasStreams::s1
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void (const sycl::stream &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const sycl::stream' lvalue <NoOp>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue ParmVar
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
 // HasStreams::s_array
 // CHECK-NEXT: InitListExpr {{.*}} 'sycl::stream [2]'
 // element 0
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void (const sycl::stream &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const sycl::stream' lvalue <NoOp>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue ParmVar
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
 // element 1
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void (const sycl::stream &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const sycl::stream' lvalue <NoOp>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue ParmVar
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
 
 // HasArrayOfHasStreams
 // CHECK-NEXT: InitListExpr {{.*}} 'HasArrayOfHasStreams'
@@ -117,38 +96,27 @@ int main() {
 // HasStreams struct
 // CHECK-NEXT: InitListExpr {{.*}} 'HasStreams'
 // HasStreams::s1
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void (const sycl::stream &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const sycl::stream' lvalue <NoOp>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue ParmVar
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
 // HasStreams::s_array
 // CHECK-NEXT: InitListExpr {{.*}} 'sycl::stream [2]'
 // element 0
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void (const sycl::stream &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const sycl::stream' lvalue <NoOp>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue ParmVar
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
 // element 1
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void (const sycl::stream &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const sycl::stream' lvalue <NoOp>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue ParmVar
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
 // HasStreams struct
 // CHECK-NEXT: InitListExpr {{.*}} 'HasStreams'
 // HasStreams::s1
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void (const sycl::stream &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const sycl::stream' lvalue <NoOp>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue ParmVar
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
 // HasStreams::s_array
 // CHECK-NEXT: InitListExpr {{.*}} 'sycl::stream [2]'
 // element 0
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void (const sycl::stream &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const sycl::stream' lvalue <NoOp>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue ParmVar
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
 // element 1
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void (const sycl::stream &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const sycl::stream' lvalue <NoOp>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue ParmVar
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
 
 // HasArrayOfHasStreams Array
-// CHECK-NEXT: InitListExpr {{.*}} 'HasArrayOfHasStreams [2]'
+// CHECK: InitListExpr {{.*}} 'HasArrayOfHasStreams [2]'
+// // HasArrayOfHasStreams Struct
 // CHECK-NEXT: InitListExpr {{.*}} 'HasArrayOfHasStreams'
 // HasArrayOfHasStreams::i
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
@@ -158,35 +126,24 @@ int main() {
 // HasStreams struct
 // CHECK-NEXT: InitListExpr {{.*}} 'HasStreams'
 // HasStreams::s1
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void (const sycl::stream &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const sycl::stream' lvalue <NoOp>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue ParmVar
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
 // HasStreams::s_array
 // CHECK-NEXT: InitListExpr {{.*}} 'sycl::stream [2]'
 // element 0
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void (const sycl::stream &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const sycl::stream' lvalue <NoOp>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue ParmVar
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
 // element 1
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void (const sycl::stream &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const sycl::stream' lvalue <NoOp>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue ParmVar
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
 // HasStreams struct
 // CHECK-NEXT: InitListExpr {{.*}} 'HasStreams'
 // HasStreams::s1
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void (const sycl::stream &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const sycl::stream' lvalue <NoOp>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue ParmVar
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
 // HasStreams::s_array
 // CHECK-NEXT: InitListExpr {{.*}} 'sycl::stream [2]'
 // element 0
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void (const sycl::stream &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const sycl::stream' lvalue <NoOp>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue ParmVar
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
 // element 1
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void (const sycl::stream &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const sycl::stream' lvalue <NoOp>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue ParmVar
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
+// HasArrayOfHasStreams Struct
 // CHECK-NEXT: InitListExpr {{.*}} 'HasArrayOfHasStreams'
 // HasArrayOfHasStreams::i
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
@@ -196,102 +153,54 @@ int main() {
 // HasStreams struct
 // CHECK-NEXT: InitListExpr {{.*}} 'HasStreams'
 // HasStreams::s1
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void (const sycl::stream &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const sycl::stream' lvalue <NoOp>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue ParmVar
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
 // HasStreams::s_array
 // CHECK-NEXT: InitListExpr {{.*}} 'sycl::stream [2]'
 // element 0
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void (const sycl::stream &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const sycl::stream' lvalue <NoOp>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue ParmVar
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
 // element 1
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void (const sycl::stream &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const sycl::stream' lvalue <NoOp>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue ParmVar
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
 // HasStreams struct
 // CHECK-NEXT: InitListExpr {{.*}} 'HasStreams'
 // HasStreams::s1
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void (const sycl::stream &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const sycl::stream' lvalue <NoOp>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue ParmVar
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
 // HasStreams::s_array
 // CHECK-NEXT: InitListExpr {{.*}} 'sycl::stream [2]'
 // element 0
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void (const sycl::stream &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const sycl::stream' lvalue <NoOp>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue ParmVar
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
 // element 1
-// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void (const sycl::stream &) noexcept'
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const sycl::stream' lvalue <NoOp>
-// CHECK-NEXT: DeclRefExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue ParmVar
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept'
 
-// Calls to Init, note that the accessor in the stream comes first, since the
-// stream __init call depends on the accessor's call already having happened.
+// Calls to init
 // in_lambda __init
 // CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void (sycl::accessor{{.*}})' lvalue .__init
-// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc
-// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
-
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init
+// CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue .
 // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
 
-// _in_lambda_array
+// in_lambda_array
 // element 0
 // CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void (sycl::accessor{{.*}})' lvalue .__init
-// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc
+// CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue
-// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
-
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue
+// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .
 // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
 
 // element 1
 // CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void (sycl::accessor{{.*}})' lvalue .__init
-// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc
+// CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue
-// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue
+// CHECK-NEXT: MemberExpr  {{.*}} 'sycl::stream [2]' lvalue .
 // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
 
 // _in_lambda_mdarray
 // [0][0]
 // CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void (sycl::accessor{{.*}})' lvalue .__init
-// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream [2]' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream (*)[2]' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2][2]' lvalue
-// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init
+// CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream [2]' lvalue
@@ -302,18 +211,7 @@ int main() {
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
 // [0][1]
 // CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void (sycl::accessor{{.*}})' lvalue .__init
-// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream [2]' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream (*)[2]' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2][2]' lvalue
-// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init
+// CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream [2]' lvalue
@@ -324,18 +222,7 @@ int main() {
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
 // [1][0]
 // CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void (sycl::accessor{{.*}})' lvalue .__init
-// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream [2]' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream (*)[2]' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2][2]' lvalue
-// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init
+// CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream [2]' lvalue
@@ -346,18 +233,7 @@ int main() {
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
 // [1][1]
 // CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void (sycl::accessor{{.*}})' lvalue .__init
-// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream [2]' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream (*)[2]' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2][2]' lvalue
-// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init
+// CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream [2]' lvalue
@@ -369,28 +245,13 @@ int main() {
 
 // HasStreams
 // CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void (sycl::accessor{{.*}})' lvalue .__init
-// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc
-// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue .s1
-// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init
+// CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue .s1
 // CHECK-NEXT: MemberExpr {{.*}} 'HasStreams' lvalue .
 // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
 // array:
 // CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void (sycl::accessor{{.*}})' lvalue .__init
-// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array
-// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init
+// CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
 // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array
@@ -399,16 +260,7 @@ int main() {
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
 // element 1
 // CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void (sycl::accessor{{.*}})' lvalue .__init
-// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array
-// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init
+// CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
 // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array
@@ -419,17 +271,7 @@ int main() {
 // HasArrayOfHasStreams
 // First element
 // CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void (sycl::accessor{{.*}})' lvalue .__init
-// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc
-// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue .s1
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs
-// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init
+// CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue .s1
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' <ArrayToPointerDecay>
@@ -439,20 +281,7 @@ int main() {
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
 // array:
 // CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void (sycl::accessor{{.*}})' lvalue .__init
-// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs
-// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init
+// CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
 // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array
@@ -465,20 +294,7 @@ int main() {
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
 // element 1
 // CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void (sycl::accessor{{.*}})' lvalue .__init
-// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs
-// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init
+// CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
 // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array
@@ -491,17 +307,7 @@ int main() {
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
 // second element
 // CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void (sycl::accessor{{.*}})' lvalue .__init
-// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc
-// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue .s1
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs
-// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init
+// CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue .s1
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' <ArrayToPointerDecay>
@@ -511,20 +317,7 @@ int main() {
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
 // array:
 // CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void (sycl::accessor{{.*}})' lvalue .__init
-// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs
-// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init
+// CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
 // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array
@@ -537,20 +330,7 @@ int main() {
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
 // element 1
 // CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void (sycl::accessor{{.*}})' lvalue .__init
-// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs
-// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init
+// CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
 // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array
@@ -561,24 +341,10 @@ int main() {
 // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-
 // HasArrayOfHasStreams array
 // First element
 // CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void (sycl::accessor{{.*}})' lvalue .__init
-// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc
-// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue .s1
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs
-// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams [2]' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init
+// CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue .s1
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' <ArrayToPointerDecay>
@@ -591,23 +357,7 @@ int main() {
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
 // array:
 // CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void (sycl::accessor{{.*}})' lvalue .__init
-// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs
-// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams [2]' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init
+// CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
 // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array
@@ -623,23 +373,7 @@ int main() {
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
 // element 1
 // CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void (sycl::accessor{{.*}})' lvalue .__init
-// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs
-// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams [2]' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init
+// CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
 // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array
@@ -655,20 +389,7 @@ int main() {
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
 // second element
 // CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void (sycl::accessor{{.*}})' lvalue .__init
-// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc
-// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue .s1
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs
-// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams [2]' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init
+// CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue .s1
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' <ArrayToPointerDecay>
@@ -681,23 +402,7 @@ int main() {
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
 // array:
 // CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void (sycl::accessor{{.*}})' lvalue .__init
-// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs
-// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams [2]' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init
+// CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
 // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array
@@ -713,23 +418,7 @@ int main() {
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
 // element 1
 // CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void (sycl::accessor{{.*}})' lvalue .__init
-// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs
-// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams [2]' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init
+// CHECK-NEXT: MemberExpr {{.*}}  'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
 // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array
@@ -745,20 +434,7 @@ int main() {
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
 // second element
 // CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void (sycl::accessor{{.*}})' lvalue .__init
-// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc
-// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue .s1
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs
-// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams [2]' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init
+// CHECK-NEXT: MemberExpr {{.*}}  'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue .s1
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' <ArrayToPointerDecay>
@@ -771,23 +447,7 @@ int main() {
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
 // array:
 // CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void (sycl::accessor{{.*}})' lvalue .__init
-// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs
-// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams [2]' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init
+// CHECK-NEXT: MemberExpr {{.*}}  'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
 // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array
@@ -803,23 +463,7 @@ int main() {
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
 // element 1
 // CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void (sycl::accessor{{.*}})' lvalue .__init
-// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs
-// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams [2]' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init
+// CHECK-NEXT: MemberExpr {{.*}}  'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
 // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array
@@ -835,20 +479,7 @@ int main() {
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
 // second element
 // CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void (sycl::accessor{{.*}})' lvalue .__init
-// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc
-// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue .s1
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs
-// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams [2]' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init
+// CHECK-NEXT: MemberExpr {{.*}}  'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue .s1
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' <ArrayToPointerDecay>
@@ -861,23 +492,7 @@ int main() {
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
 // array:
 // CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void (sycl::accessor{{.*}})' lvalue .__init
-// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs
-// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams [2]' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init
+// CHECK-NEXT: MemberExpr {{.*}}  'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
 // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array
@@ -893,23 +508,7 @@ int main() {
 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0
 // element 1
 // CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void (sycl::accessor{{.*}})' lvalue .__init
-// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array
-// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs
-// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue
-// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' <ArrayToPointerDecay>
-// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams [2]' lvalue .
-// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1
-// CHECK: CXXMemberCallExpr {{.*}} 'void'
-// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init
+// CHECK-NEXT: MemberExpr {{.*}}  'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init
 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue
 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' <ArrayToPointerDecay>
 // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array
diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp
index 35a9d6e90880..b86c4bbe3471 100755
--- a/sycl/include/CL/sycl/accessor.hpp
+++ b/sycl/include/CL/sycl/accessor.hpp
@@ -203,7 +203,7 @@
 
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
-
+class stream;
 namespace ext {
 namespace intel {
 namespace experimental {
@@ -930,6 +930,7 @@ class accessor :
 #endif // __SYCL_DEVICE_ONLY__
 
 private:
+  friend class sycl::stream;
   friend class sycl::ext::intel::experimental::esimd::detail::
       AccessorPrivateProxy;
 
diff --git a/sycl/include/CL/sycl/detail/kernel_desc.hpp b/sycl/include/CL/sycl/detail/kernel_desc.hpp
index 7db5b75386e1..d3a4de78657b 100644
--- a/sycl/include/CL/sycl/detail/kernel_desc.hpp
+++ b/sycl/include/CL/sycl/detail/kernel_desc.hpp
@@ -35,6 +35,7 @@ enum class kernel_param_kind_t {
   kind_sampler = 2,
   kind_pointer = 3,
   kind_specialization_constants_buffer = 4,
+  kind_stream = 5,
 };
 
 // describes a kernel parameter
diff --git a/sycl/include/CL/sycl/stream.hpp b/sycl/include/CL/sycl/stream.hpp
index 678fc650668a..4bd46ddfdf6c 100644
--- a/sycl/include/CL/sycl/stream.hpp
+++ b/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,25 @@ 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) {
+    GlobalBuf.__init(GlobalBufPtr, GlobalBufAccRange, GlobalBufMemRange,
+                     GlobalBufId);
+    GlobalOffset.__init(GlobalOffsetPtr, GlobalOffsetAccRange,
+                        GlobalOffsetMemRange, GlobalOffsetId);
+    GlobalFlushBuf.__init(GlobalFlushPtr, GlobalFlushAccRange,
+                          GlobalFlushMemRange, GlobalFlushId);
+    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 +906,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 +1155,3 @@ template <> struct hash<cl::sycl::stream> {
   }
 };
 } // namespace std
-
diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp
index 5535e397b200..4ea9394f2034 100644
--- a/sycl/source/detail/scheduler/commands.cpp
+++ b/sycl/source/detail/scheduler/commands.cpp
@@ -1661,6 +1661,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);
diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp
index c7c1f6c06ab2..4c53b7eb5c1a 100644
--- a/sycl/source/handler.cpp
+++ b/sycl/source/handler.cpp
@@ -14,6 +14,7 @@
 #include <CL/sycl/event.hpp>
 #include <CL/sycl/handler.hpp>
 #include <CL/sycl/info/info_desc.hpp>
+#include <CL/sycl/stream.hpp>
 #include <detail/global_handler.hpp>
 #include <detail/kernel_bundle_impl.hpp>
 #include <detail/kernel_impl.hpp>
@@ -234,6 +235,41 @@ void handler::associateWithHandler(detail::AccessorBaseHost *AccBase,
                                    /*index*/ 0);
 }
 
+static void addArgsForGlobalAccessor(detail::Requirement *AccImpl, size_t Index,
+                                     size_t &IndexShift, int Size,
+                                     bool IsKernelCreatedFromSource,
+                                     size_t GlobalSize,
+                                     vector_class<detail::ArgDesc> &Args,
+                                     bool isESIMD) {
+  using detail::kernel_param_kind_t;
+  if (AccImpl->PerWI)
+    AccImpl->resize(GlobalSize);
+
+  Args.emplace_back(kernel_param_kind_t::kind_accessor, AccImpl, Size,
+                    Index + IndexShift);
+
+  // TODO ESIMD currently does not suport offset, memory and access ranges -
+  // accessor::init for ESIMD-mode accessor has a single field, translated
+  // to a single kernel argument set above.
+  if (!isESIMD && !IsKernelCreatedFromSource) {
+    // Dimensionality of the buffer is 1 when dimensionality of the
+    // accessor is 0.
+    const size_t SizeAccField =
+        sizeof(size_t) * (AccImpl->MDims == 0 ? 1 : AccImpl->MDims);
+    ++IndexShift;
+    Args.emplace_back(kernel_param_kind_t::kind_std_layout,
+                      &AccImpl->MAccessRange[0], SizeAccField,
+                      Index + IndexShift);
+    ++IndexShift;
+    Args.emplace_back(kernel_param_kind_t::kind_std_layout,
+                      &AccImpl->MMemoryRange[0], SizeAccField,
+                      Index + IndexShift);
+    ++IndexShift;
+    Args.emplace_back(kernel_param_kind_t::kind_std_layout,
+                      &AccImpl->MOffset[0], SizeAccField, Index + IndexShift);
+  }
+}
+
 // TODO remove this one once ABI breaking changes are allowed.
 void handler::processArg(void *Ptr, const detail::kernel_param_kind_t &Kind,
                          const int Size, const size_t Index, size_t &IndexShift,
@@ -253,6 +289,40 @@ void handler::processArg(void *Ptr, const detail::kernel_param_kind_t &Kind,
     MArgs.emplace_back(Kind, Ptr, Size, Index + IndexShift);
     break;
   }
+  case kernel_param_kind_t::kind_stream: {
+    // Stream contains several accessors inside.
+    stream *S = static_cast<stream *>(Ptr);
+
+    detail::AccessorBaseHost *GBufBase =
+        static_cast<detail::AccessorBaseHost *>(&S->GlobalBuf);
+    detail::AccessorImplPtr GBufImpl = detail::getSyclObjImpl(*GBufBase);
+    detail::Requirement *GBufReq = GBufImpl.get();
+    addArgsForGlobalAccessor(GBufReq, Index, IndexShift, Size,
+                             IsKernelCreatedFromSource,
+                             MNDRDesc.GlobalSize.size(), MArgs, IsESIMD);
+    ++IndexShift;
+    detail::AccessorBaseHost *GOffsetBase =
+        static_cast<detail::AccessorBaseHost *>(&S->GlobalOffset);
+    detail::AccessorImplPtr GOfssetImpl = detail::getSyclObjImpl(*GOffsetBase);
+    detail::Requirement *GOffsetReq = GOfssetImpl.get();
+    addArgsForGlobalAccessor(GOffsetReq, Index, IndexShift, Size,
+                             IsKernelCreatedFromSource,
+                             MNDRDesc.GlobalSize.size(), MArgs, IsESIMD);
+    ++IndexShift;
+    detail::AccessorBaseHost *GFlushBase =
+        static_cast<detail::AccessorBaseHost *>(&S->GlobalFlushBuf);
+    detail::AccessorImplPtr GFlushImpl = detail::getSyclObjImpl(*GFlushBase);
+    detail::Requirement *GFlushReq = GFlushImpl.get();
+    addArgsForGlobalAccessor(GFlushReq, Index, IndexShift, Size,
+                             IsKernelCreatedFromSource,
+                             MNDRDesc.GlobalSize.size(), MArgs, IsESIMD);
+    ++IndexShift;
+    MArgs.emplace_back(kernel_param_kind_t::kind_std_layout,
+                       &S->FlushBufferSize, sizeof(S->FlushBufferSize),
+                       Index + IndexShift);
+
+    break;
+  }
   case kernel_param_kind_t::kind_accessor: {
     // For args kind of accessor Size is information about accessor.
     // The first 11 bits of Size encodes the accessor target.
@@ -261,37 +331,9 @@ void handler::processArg(void *Ptr, const detail::kernel_param_kind_t &Kind,
     case access::target::global_buffer:
     case access::target::constant_buffer: {
       detail::Requirement *AccImpl = static_cast<detail::Requirement *>(Ptr);
-
-      // Stream implementation creates an accessor with initial size for
-      // work item. Number of work items is not available during
-      // stream construction, that is why size of the accessor is updated here
-      // using information about number of work items.
-      if (AccImpl->PerWI) {
-        AccImpl->resize(MNDRDesc.GlobalSize.size());
-      }
-      MArgs.emplace_back(Kind, AccImpl, Size, Index + IndexShift);
-
-      // TODO ESIMD currently does not suport offset, memory and access ranges -
-      // accessor::init for ESIMD-mode accessor has a single field, translated
-      // to a single kernel argument set above.
-      if (!IsKernelCreatedFromSource && !IsESIMD) {
-        // Dimensionality of the buffer is 1 when dimensionality of the
-        // accessor is 0.
-        const size_t SizeAccField =
-            sizeof(size_t) * (AccImpl->MDims == 0 ? 1 : AccImpl->MDims);
-        ++IndexShift;
-        MArgs.emplace_back(kernel_param_kind_t::kind_std_layout,
-                           &AccImpl->MAccessRange[0], SizeAccField,
-                           Index + IndexShift);
-        ++IndexShift;
-        MArgs.emplace_back(kernel_param_kind_t::kind_std_layout,
-                           &AccImpl->MMemoryRange[0], SizeAccField,
-                           Index + IndexShift);
-        ++IndexShift;
-        MArgs.emplace_back(kernel_param_kind_t::kind_std_layout,
-                           &AccImpl->MOffset[0], SizeAccField,
-                           Index + IndexShift);
-      }
+      addArgsForGlobalAccessor(AccImpl, Index, IndexShift, Size,
+                               IsKernelCreatedFromSource,
+                               MNDRDesc.GlobalSize.size(), MArgs, IsESIMD);
       break;
     }
     case access::target::local: {