diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 619c208597763..ac072908599ea 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/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< diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index e2786808c040b..df1e1ba4b56e1 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 d1876a67bea6c..5f2273b6c8993 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 - void visitStreamRecord(const CXXRecordDecl *Owner, ParentTy &Parent, - CXXRecordDecl *Wrapper, QualType RecordTy, - HandlerTys &... Handlers) { - (void)std::initializer_list{ - (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{ - (Handlers.leaveStream(Owner, Parent, RecordTy), 0)...}; - } - template 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); } diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index d27f3c9e1a67d..f709e2ed1d5f0 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 AccessRange, range MemRange, id Offset) {} void __init_esimd(__attribute__((opencl_global)) dataT *Ptr) {} + friend class stream; }; template @@ -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 Acc; + int FlushBufferSize; }; template diff --git a/clang/test/CodeGenSYCL/stream.cpp b/clang/test/CodeGenSYCL/stream.cpp index cfc354efa0686..179acbced4d7e 100644 --- a/clang/test/CodeGenSYCL/stream.cpp +++ b/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" diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp index 4782bafa74ce5..88dc82a4f5764 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::type *; void __init(PtrType Ptr, range AccessRange, range MemRange, id Offset) {} + friend class stream; }; template @@ -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 Acc; + int FlushBufferSize; }; namespace ONEAPI { diff --git a/clang/test/SemaSYCL/decomposition.cpp b/clang/test/SemaSYCL/decomposition.cpp index a5e3bebca0c27..f56fde742b11f 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([=]() { 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 t2; myQueue.submit([&](sycl::handler &h) { h.single_task([=]() { 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 967e269dace2d..cf6b28cb4f9ac 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/streams.cpp index 933dabff649fb..50c780fdefe5a 100644 --- a/clang/test/SemaSYCL/streams.cpp +++ b/clang/test/SemaSYCL/streams.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 -// 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 -// 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 -// 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 -// 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 -// 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 -// 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 -// 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 -// 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 -// 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 -// 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,26 @@ 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 -// 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 -// 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 -// 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 -// 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 -// 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 -// 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]' +// HasArrayOfHasStreams struct +// CHECK: InitListExpr {{.*}} 'HasArrayOfHasStreams [2]' // CHECK-NEXT: InitListExpr {{.*}} 'HasArrayOfHasStreams' // HasArrayOfHasStreams::i // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' @@ -158,35 +125,23 @@ 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 -// 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 -// 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 -// 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 -// 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 -// 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 -// CHECK-NEXT: DeclRefExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue ParmVar +// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept' // CHECK-NEXT: InitListExpr {{.*}} 'HasArrayOfHasStreams' // HasArrayOfHasStreams::i // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' @@ -196,102 +151,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 -// 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 -// 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 -// 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 -// 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 -// 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 -// 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 *' -// 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 *' -// 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 *' -// 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 *' -// 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 *' -// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream [2]' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream (*)[2]' -// 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 *' // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream [2]' lvalue @@ -302,18 +209,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 *' -// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream [2]' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream (*)[2]' -// 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 *' // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream [2]' lvalue @@ -324,18 +220,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 *' -// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream [2]' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream (*)[2]' -// 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 *' // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream [2]' lvalue @@ -346,18 +231,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 *' -// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream [2]' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream (*)[2]' -// 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 *' // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream [2]' lvalue @@ -369,28 +243,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 *' -// 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 *' // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array @@ -399,16 +258,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 *' -// 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 *' // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array @@ -419,17 +269,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 *' -// 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 *' @@ -439,20 +279,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 *' -// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array -// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' -// 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 *' // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array @@ -465,20 +292,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 *' -// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array -// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' -// 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 *' // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array @@ -491,17 +305,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 *' -// 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 *' @@ -511,20 +315,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 *' -// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array -// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' -// 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 *' // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array @@ -537,8 +328,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: 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 *' // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array @@ -549,36 +339,10 @@ int main() { // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 // 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 *' -// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array -// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' -// 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 - // 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 *' -// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs -// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' -// 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 *' @@ -591,23 +355,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 *' -// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array -// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' -// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs -// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' -// 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 *' // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array @@ -623,23 +371,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 *' -// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array -// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' -// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs -// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' -// 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 *' // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array @@ -655,20 +387,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 *' -// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs -// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' -// 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 *' @@ -681,23 +400,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 *' -// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array -// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' -// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs -// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' -// 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 *' // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array @@ -713,23 +416,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 *' -// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array -// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' -// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs -// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' -// 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 *' // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array @@ -745,20 +432,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 *' -// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs -// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' -// 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 *' @@ -771,23 +445,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 *' -// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array -// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' -// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs -// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' -// 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 *' // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array @@ -803,23 +461,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 *' -// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array -// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' -// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs -// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' -// 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 *' // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array @@ -835,20 +477,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 *' -// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs -// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' -// 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 *' @@ -861,23 +490,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 *' -// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array -// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' -// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs -// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' -// 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 *' // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array @@ -893,23 +506,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 *' -// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array -// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' -// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs -// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' -// 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 *' // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array @@ -948,7 +545,6 @@ int main() { // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 - // _in_lambda_mdarray // [0][0] // CHECK: CXXMemberCallExpr {{.*}} 'void' diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp index 35a9d6e90880d..b86c4bbe34714 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;