From ade97f5a3893a17bdcc457464bf7f5f6ee50a8e0 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Wed, 28 Apr 2021 12:58:32 -0700 Subject: [PATCH 01/20] Stream handing implementation - PR2268 Signed-off-by: Zahira Ammarguellat --- clang/include/clang/Sema/Sema.h | 1 + clang/lib/Sema/SemaSYCL.cpp | 34 ++----- clang/test/CodeGenSYCL/Inputs/sycl.hpp | 15 ++- clang/test/CodeGenSYCL/stream.cpp | 15 ++- sycl/include/CL/sycl/accessor.hpp | 2 + sycl/include/CL/sycl/detail/kernel_desc.hpp | 5 +- sycl/include/CL/sycl/stream.hpp | 43 +++++++- sycl/source/detail/scheduler/commands.cpp | 2 + sycl/source/handler.cpp | 104 ++++++++++++++------ 9 files changed, 154 insertions(+), 67 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index f1414c3af6e6a..7aae13dee5f59 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -314,6 +314,7 @@ class SYCLIntegrationHeader { kind_accessor = kind_first, kind_std_layout, kind_sampler, + kind_stream, kind_pointer, kind_specialization_constants_buffer, kind_last = kind_specialization_constants_buffer diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index b0459d97c01f0..9ceb53ad38b49 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1025,23 +1025,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, @@ -1116,12 +1099,9 @@ class KernelObjVisitor { KF_FOR_EACH(handleSyclHalfType, Field, FieldTy); else if (Util::isSyclSpecConstantType(FieldTy)) KF_FOR_EACH(handleSyclSpecConstantType, Field, FieldTy); - else if (Util::isSyclStreamType(FieldTy)) { - CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl(); - // Handle accessors in stream class. + else if (Util::isSyclStreamType(FieldTy)) KF_FOR_EACH(handleSyclStreamType, Field, FieldTy); - visitStreamRecord(Owner, Field, RD, FieldTy, Handlers...); - } else if (FieldTy->isStructureOrClassType()) { + else if (FieldTy->isStructureOrClassType()) { if (KF_FOR_EACH(handleStructType, Field, FieldTy)) { CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl(); visitRecord(Owner, Field, RD, FieldTy, Handlers...); @@ -2095,8 +2075,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 &, @@ -2684,9 +2663,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, @@ -3101,7 +3078,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; } @@ -3948,6 +3925,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..7f4092fa9ddc8 100644 --- a/clang/test/CodeGenSYCL/stream.cpp +++ b/clang/test/CodeGenSYCL/stream.cpp @@ -1,10 +1,17 @@ // 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]]) {{.*}}%{{.*}} #include "Inputs/sycl.hpp" diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp index 3245a68c668ad..494c325fc5d1b 100755 --- a/sycl/include/CL/sycl/accessor.hpp +++ b/sycl/include/CL/sycl/accessor.hpp @@ -203,6 +203,7 @@ __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { +class stream; namespace INTEL { namespace gpu { namespace detail { @@ -929,6 +930,7 @@ class accessor : private: friend class sycl::INTEL::gpu::detail::AccessorPrivateProxy; + friend class sycl::stream; public: using value_type = DataT; diff --git a/sycl/include/CL/sycl/detail/kernel_desc.hpp b/sycl/include/CL/sycl/detail/kernel_desc.hpp index 7db5b75386e17..413ee66a55cc2 100644 --- a/sycl/include/CL/sycl/detail/kernel_desc.hpp +++ b/sycl/include/CL/sycl/detail/kernel_desc.hpp @@ -33,8 +33,9 @@ enum class kernel_param_kind_t { kind_accessor = 0, kind_std_layout = 1, // standard layout object parameters kind_sampler = 2, - kind_pointer = 3, - kind_specialization_constants_buffer = 4, + kind_stream = 3, + kind_pointer = 4, + kind_specialization_constants_buffer = 5, }; // describes a kernel parameter diff --git a/sycl/include/CL/sycl/stream.hpp b/sycl/include/CL/sycl/stream.hpp index 678fc650668aa..77f1259b2f55d 100644 --- a/sycl/include/CL/sycl/stream.hpp +++ b/sycl/include/CL/sycl/stream.hpp @@ -68,11 +68,23 @@ using GlobalBufAccessorT = accessor; +constexpr static access::address_space GlobalBufAS = + TargetToAS::AS; +using GlobalBufPtrType = + typename detail::PtrValueType::type *; +constexpr static int GlobalBufDim = 1; + using GlobalOffsetAccessorT = accessor; +constexpr static access::address_space GlobalOffsetAS = + TargetToAS::AS; +using GlobalOffsetPtrType = + typename detail::PtrValueType::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(&GlobalFlushBuf[WIOffset]); @@ -729,6 +741,12 @@ inline __width_manipulator__ setw(int Width) { /// \ingroup sycl_api class __SYCL_EXPORT stream { public: + +#ifdef __SYCL_DEVICE_ONLY__ + // Default constructor for objects later initialized with __init member. + stream() = default; +#endif + // Throws exception in case of invalid input parameters stream(size_t BufferSize, size_t MaxStatementSize, handler &CGH); @@ -845,7 +863,28 @@ class __SYCL_EXPORT stream { } #ifdef __SYCL_DEVICE_ONLY__ - void __init() { + void __init(detail::GlobalBufPtrType GlobalBufPtr, + range GlobalBufAccRange, + range GlobalBufMemRange, + id GlobalBufId, + detail::GlobalOffsetPtrType GlobalOffsetPtr, + range GlobalOffsetAccRange, + range GlobalOffsetMemRange, + id GlobalOffsetId, + detail::GlobalBufPtrType GlobalFlushPtr, + range GlobalFlushAccRange, + range GlobalFlushMemRange, + id GlobalFlushId, + size_t _FlushBufferSize) { +#ifndef __SYCL_EXPLICIT_SIMD__ + GlobalBuf.__init(GlobalBufPtr, GlobalBufAccRange, GlobalBufMemRange, + GlobalBufId); + GlobalOffset.__init(GlobalOffsetPtr, GlobalOffsetAccRange, + GlobalOffsetMemRange, GlobalOffsetId); + GlobalFlushBuf.__init(GlobalFlushPtr, GlobalFlushAccRange, + GlobalFlushMemRange, GlobalFlushId); +#endif + FlushBufferSize = _FlushBufferSize; // Calculate offset in the flush buffer for each work item in the global // work space. We need to avoid calling intrinsics to get global id because // when stream is used in a single_task kernel this could cause some @@ -871,6 +910,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 diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp index 471c11e19e6b1..15a75c9871ff2 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/sycl/source/detail/scheduler/commands.cpp @@ -1660,6 +1660,8 @@ pi_result ExecCGCommand::SetKernelParamsAndLaunch( if (!EliminatedArgMask.empty() && EliminatedArgMask[Arg.MIndex]) continue; switch (Arg.MType) { + case kernel_param_kind_t::kind_stream: + break; case kernel_param_kind_t::kind_accessor: { Requirement *Req = (Requirement *)(Arg.MPtr); AllocaCommandBase *AllocaCmd = getAllocaForReq(Req); diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index 81c1bc6846fe2..924e5758263e6 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -14,6 +14,7 @@ #include #include #include +#include #include #include #include @@ -230,6 +231,41 @@ void handler::associateWithHandler(detail::AccessorBaseHost *AccBase, /*index*/ 0); } +static void addArgsForGlobalAccessor(detail::Requirement *AccImpl, + const size_t Index, size_t &IndexShift, + const int Size, + bool IsKernelCreatedFromSource, + size_t GlobalSize, + vector_class &Args) { + 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 (!AccImpl->MIsESIMDAcc && !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, @@ -249,6 +285,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(Ptr); + + detail::AccessorBaseHost *GBufBase = + (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); + ++IndexShift; + detail::AccessorBaseHost *GOffsetBase = + (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); + ++IndexShift; + detail::AccessorBaseHost *GFlushBase = + (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); + ++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. @@ -257,37 +327,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(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); break; } case access::target::local: { From 9a60e02d583a44add0c9eb185638ebfda179f4a0 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Wed, 28 Apr 2021 13:03:50 -0700 Subject: [PATCH 02/20] Stream handing implementation - PR2268 Signed-off-by: Zahira Ammarguellat --- sycl/include/CL/sycl/stream.hpp | 4 +--- 1 file changed, 1 insertion(+), 3 deletions(-) diff --git a/sycl/include/CL/sycl/stream.hpp b/sycl/include/CL/sycl/stream.hpp index 77f1259b2f55d..56310c9c5f5d2 100644 --- a/sycl/include/CL/sycl/stream.hpp +++ b/sycl/include/CL/sycl/stream.hpp @@ -741,7 +741,6 @@ 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; @@ -874,8 +873,7 @@ class __SYCL_EXPORT stream { detail::GlobalBufPtrType GlobalFlushPtr, range GlobalFlushAccRange, range GlobalFlushMemRange, - id GlobalFlushId, - size_t _FlushBufferSize) { + id GlobalFlushId, size_t _FlushBufferSize) { #ifndef __SYCL_EXPLICIT_SIMD__ GlobalBuf.__init(GlobalBufPtr, GlobalBufAccRange, GlobalBufMemRange, GlobalBufId); From eb0b35d1366d187c678020e8e080be5297898e90 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Thu, 29 Apr 2021 05:12:43 -0700 Subject: [PATCH 03/20] Stream handing implementation - code from draft PR2268 Signed-off-by: Zahira Ammarguellat --- sycl/include/CL/sycl/stream.hpp | 17 ++++++++--------- 1 file changed, 8 insertions(+), 9 deletions(-) diff --git a/sycl/include/CL/sycl/stream.hpp b/sycl/include/CL/sycl/stream.hpp index 56310c9c5f5d2..4c97d7f1a98e7 100644 --- a/sycl/include/CL/sycl/stream.hpp +++ b/sycl/include/CL/sycl/stream.hpp @@ -71,7 +71,7 @@ using GlobalBufAccessorT = accessor::AS; using GlobalBufPtrType = - typename detail::PtrValueType::type *; + typename detail::DecoratedType::type *; constexpr static int GlobalBufDim = 1; using GlobalOffsetAccessorT = @@ -82,7 +82,7 @@ using GlobalOffsetAccessorT = constexpr static access::address_space GlobalOffsetAS = TargetToAS::AS; using GlobalOffsetPtrType = - typename detail::PtrValueType::type *; + typename detail::DecoratedType::type *; constexpr static int GlobalOffsetDim = 1; // Read first 2 bytes of flush buffer to get buffer offset. @@ -630,12 +630,12 @@ inline void writeHItem(GlobalBufAccessorT &GlobalFlushBuf, unsigned Len = 0; Len += append(Buf, "h_item("); for (int I = 0; I < 3; ++I) { - Len += append(Buf + Len, I == 0 ? "\n global " - : I == 1 ? "\n logical local " - : "\n physical local "); - Len += ItemToStr(Buf + Len, I == 0 ? HItem.get_global() - : I == 1 ? HItem.get_logical_local() - : HItem.get_physical_local()); + Len += append(Buf + Len, I == 0 ? "\n global " + : I == 1 ? "\n logical local " + : "\n physical local "); + Len += ItemToStr(Buf + Len, I == 0 ? HItem.get_global() + : I == 1 ? HItem.get_logical_local() + : HItem.get_physical_local()); } Len += append(Buf + Len, "\n)"); write(GlobalFlushBuf, FlushBufferSize, WIOffset, Buf, Len); @@ -1157,4 +1157,3 @@ template <> struct hash { } }; } // namespace std - From d9f7e7906882eaad5b236e1588b5654dc31cd4b9 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Thu, 29 Apr 2021 05:18:05 -0700 Subject: [PATCH 04/20] Stream handing implementation - code from draft PR2268 Signed-off-by: Zahira Ammarguellat --- sycl/include/CL/sycl/stream.hpp | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/sycl/include/CL/sycl/stream.hpp b/sycl/include/CL/sycl/stream.hpp index 4c97d7f1a98e7..53349f31c3b2d 100644 --- a/sycl/include/CL/sycl/stream.hpp +++ b/sycl/include/CL/sycl/stream.hpp @@ -630,12 +630,12 @@ inline void writeHItem(GlobalBufAccessorT &GlobalFlushBuf, unsigned Len = 0; Len += append(Buf, "h_item("); for (int I = 0; I < 3; ++I) { - Len += append(Buf + Len, I == 0 ? "\n global " - : I == 1 ? "\n logical local " - : "\n physical local "); - Len += ItemToStr(Buf + Len, I == 0 ? HItem.get_global() - : I == 1 ? HItem.get_logical_local() - : HItem.get_physical_local()); + Len += append(Buf + Len, I == 0 ? "\n global " + : I == 1 ? "\n logical local " + : "\n physical local "); + Len += ItemToStr(Buf + Len, I == 0 ? HItem.get_global() + : I == 1 ? HItem.get_logical_local() + : HItem.get_physical_local()); } Len += append(Buf + Len, "\n)"); write(GlobalFlushBuf, FlushBufferSize, WIOffset, Buf, Len); From d50466285f8b7c3e9ff47283b715eb65cd5c34e7 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Thu, 29 Apr 2021 10:55:13 -0700 Subject: [PATCH 05/20] Adding finalize function Signed-off-by: Zahira Ammarguellat --- clang/include/clang/Sema/Sema.h | 4 +- clang/lib/Sema/SemaSYCL.cpp | 64 +------------------------------ clang/test/CodeGenSYCL/stream.cpp | 1 + 3 files changed, 4 insertions(+), 65 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index b028a7bc63c6d..aac9bf01d4603 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -314,10 +314,10 @@ class SYCLIntegrationHeader { kind_accessor = kind_first, kind_std_layout, kind_sampler, - kind_stream, 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 9ceb53ad38b49..8e254d09beea5 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1215,12 +1215,6 @@ class SyclKernelFieldHandlerBase { virtual bool leaveStruct(const CXXRecordDecl *, FieldDecl *, QualType) { return true; } - virtual bool enterStream(const CXXRecordDecl *, FieldDecl *, QualType) { - return true; - } - virtual bool leaveStream(const CXXRecordDecl *, FieldDecl *, QualType) { - return true; - } virtual bool enterStruct(const CXXRecordDecl *, const CXXBaseSpecifier &, QualType) { return true; @@ -1668,18 +1662,6 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { return true; } - // Stream is always decomposed (and whether it gets decomposed is handled in - // handleSyclStreamType), but we need a CollectionStack entry to capture the - // accessors that get handled. - bool enterStream(const CXXRecordDecl *, FieldDecl *, QualType) final { - CollectionStack.push_back(false); - return true; - } - bool leaveStream(const CXXRecordDecl *, FieldDecl *, QualType Ty) final { - CollectionStack.pop_back(); - return true; - } - bool enterStruct(const CXXRecordDecl *, FieldDecl *, QualType) final { CollectionStack.push_back(false); return true; @@ -1932,14 +1914,6 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { SemaRef.addSyclDeviceDecl(KernelDecl); } - bool enterStream(const CXXRecordDecl *RD, FieldDecl *FD, QualType Ty) final { - return enterStruct(RD, FD, Ty); - } - - bool leaveStream(const CXXRecordDecl *RD, FieldDecl *FD, QualType Ty) final { - return leaveStruct(RD, FD, Ty); - } - bool enterStruct(const CXXRecordDecl *, FieldDecl *, QualType) final { ++StructDepth; return true; @@ -2570,6 +2544,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { const auto *RecordDecl = Ty->getAsCXXRecordDecl(); createSpecialMethodCall(RecordDecl, getInitMethodName(), BodyStmts); + createSpecialMethodCall(RecordDecl, FinalizeMethodName, BodyStmts); removeFieldMemberExpr(FD, Ty); @@ -2740,31 +2715,6 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { handleSpecialType(KernelHandlerArg->getType()); } - bool enterStream(const CXXRecordDecl *RD, FieldDecl *FD, QualType Ty) final { - ++StructDepth; - // Add a dummy init expression to catch the accessor initializers. - const auto *StreamDecl = Ty->getAsCXXRecordDecl(); - CollectionInitExprs.push_back(createInitListExpr(StreamDecl)); - - addFieldMemberExpr(FD, Ty); - return true; - } - - bool leaveStream(const CXXRecordDecl *RD, FieldDecl *FD, QualType Ty) final { - --StructDepth; - // Stream requires that its 'init' calls happen after its accessors init - // calls, so add them here instead. - const auto *StreamDecl = Ty->getAsCXXRecordDecl(); - - createSpecialMethodCall(StreamDecl, getInitMethodName(), BodyStmts); - createSpecialMethodCall(StreamDecl, FinalizeMethodName, FinalizeStmts); - - removeFieldMemberExpr(FD, Ty); - - CollectionInitExprs.pop_back(); - return true; - } - bool enterStruct(const CXXRecordDecl *RD, FieldDecl *FD, QualType Ty) final { ++StructDepth; addCollectionInitListExpr(Ty->getAsCXXRecordDecl()); @@ -3110,18 +3060,6 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { SYCLIntegrationHeader::kind_specialization_constants_buffer, 0); } - bool enterStream(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final { - ++StructDepth; - CurOffset += offsetOf(FD, Ty); - return true; - } - - bool leaveStream(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final { - --StructDepth; - CurOffset -= offsetOf(FD, Ty); - return true; - } - bool enterStruct(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final { ++StructDepth; CurOffset += offsetOf(FD, Ty); diff --git a/clang/test/CodeGenSYCL/stream.cpp b/clang/test/CodeGenSYCL/stream.cpp index 7f4092fa9ddc8..179acbced4d7e 100644 --- a/clang/test/CodeGenSYCL/stream.cpp +++ b/clang/test/CodeGenSYCL/stream.cpp @@ -12,6 +12,7 @@ // 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" From cf327e59fa0ac4af955944c15d6c399d906bea67 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Thu, 29 Apr 2021 13:56:53 -0700 Subject: [PATCH 06/20] Adding finalize function Signed-off-by: Zahira Ammarguellat --- clang/lib/Sema/SemaSYCL.cpp | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 8e254d09beea5..b6a8883a8438e 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -2544,7 +2544,10 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { const auto *RecordDecl = Ty->getAsCXXRecordDecl(); createSpecialMethodCall(RecordDecl, getInitMethodName(), BodyStmts); - createSpecialMethodCall(RecordDecl, FinalizeMethodName, BodyStmts); + CXXMethodDecl *FinalizeMethod = + getMethodByName(RecordDecl, FinalizeMethodName); + if (FinalizeMethod) + createSpecialMethodCall(RecordDecl, FinalizeMethodName, BodyStmts); removeFieldMemberExpr(FD, Ty); From 75d0e57fb0bf85d3a298d2f2d67dc8e9fce9d02e Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Mon, 3 May 2021 05:03:02 -0700 Subject: [PATCH 07/20] Review comments fixes Signed-off-by: Zahira Ammarguellat --- clang/test/SemaSYCL/Inputs/sycl.hpp | 16 +- clang/test/SemaSYCL/decomposition.cpp | 5 +- clang/test/SemaSYCL/streams.cpp | 842 ++++++-------------------- 3 files changed, 199 insertions(+), 664 deletions(-) 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/streams.cpp b/clang/test/SemaSYCL/streams.cpp index 933dabff649fb..d65f21d937d5c 100644 --- a/clang/test/SemaSYCL/streams.cpp +++ b/clang/test/SemaSYCL/streams.cpp @@ -44,371 +44,247 @@ int main() { return 0; } - -// Function Declaration -// CHECK: FunctionDecl {{.*}}stream_test{{.*}} - -// Initializers: - +// CHECK: FunctionDecl {{.*}} main 'int ()' +// CHECK: FunctionDecl {{.*}}stream_test // 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 -// 'in_lambda_array' +// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept' // 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 -// 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 - -// 'in_lambda_mdarray' +// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept' +// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept' // 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 -// 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 -// sub-array 1 +// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept' +// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept' // 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 -// 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: 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 -// HasStreams::s_array +// CHECK: InitListExpr {{.*}} 'HasStreams' +// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept' // 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 -// 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 - -// HasArrayOfHasStreams +// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept' +// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept' // CHECK-NEXT: InitListExpr {{.*}} 'HasArrayOfHasStreams' -// HasArrayOfHasStreams::i // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar -// HasArrayOfHasStreams::hs // CHECK-NEXT: InitListExpr {{.*}} 'HasStreams [2]' -// 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 -// 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 -// 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 -// 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 -// 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 -// 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 - -// HasArrayOfHasStreams Array -// CHECK-NEXT: InitListExpr {{.*}} 'HasArrayOfHasStreams [2]' +// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept' + +// HasArrayOfHasStreams struct +// CHECK: InitListExpr {{.*}} 'HasArrayOfHasStreams [2]' // CHECK-NEXT: InitListExpr {{.*}} 'HasArrayOfHasStreams' -// HasArrayOfHasStreams::i // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar -// HasArrayOfHasStreams::hs + // CHECK-NEXT: InitListExpr {{.*}} 'HasStreams [2]' -// 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 -// HasStreams::s_array +// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept' // 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 -// 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 -// HasStreams struct +// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept' +// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept' + // 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 -// HasStreams::s_array +// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept' // 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 -// 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: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept' + // CHECK-NEXT: InitListExpr {{.*}} 'HasArrayOfHasStreams' -// HasArrayOfHasStreams::i // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar -// HasArrayOfHasStreams::hs + // CHECK-NEXT: InitListExpr {{.*}} 'HasStreams [2]' -// 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 -// HasStreams::s_array +// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept' // 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 -// 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 -// HasStreams struct +// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept' +// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept' + // 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 -// HasStreams::s_array +// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept' // 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 -// 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 - -// 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. -// in_lambda __init +// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept' +// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept' + +// Calls to init // 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: MemberExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue . // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at // CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue . // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at -// _in_lambda_array -// element 0 +//_in_lambda_array // CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void (sycl::accessor{{.*}})' lvalue .__init -// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc +// CHECK: 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: 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: MemberExpr {{.*}} 'void ()' lvalue .__finalize // 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: 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: MemberExpr {{.*}} 'void ()' lvalue .__finalize // 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: 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 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream (*)[2]' -// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2][2]' lvalue +// 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 ()' lvalue .__finalize // 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: MemberExpr {{.*}} 'sycl::stream [2][2]' lvalue . // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 -// 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: 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 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream (*)[2]' -// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2][2]' lvalue +// 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 ()' lvalue .__finalize // 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: MemberExpr {{.*}} 'sycl::stream [2][2]' lvalue . // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 // 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: CXXMemberCallExpr {{.*}} 'void' +// 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 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream (*)[2]' -// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2][2]' lvalue +// 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 ()' lvalue .__finalize // 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: MemberExpr {{.*}} 'sycl::stream [2][2]' lvalue . // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 -// [1][1] -// CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void (sycl::accessor{{.*}})' lvalue .__init -// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc + +// CHECK: CXXMemberCallExpr {{.*}} 'void' +// 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 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream (*)[2]' -// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2][2]' lvalue +// 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 ()' lvalue .__finalize // 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: MemberExpr {{.*}} 'sycl::stream [2][2]' lvalue . // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 -// HasStreams // 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: MemberExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue .s1 // CHECK-NEXT: MemberExpr {{.*}} 'HasStreams' lvalue . // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global char *' +// CHECK-NEXT: DeclRefExpr {{.*}} '__global char *' lvalue ParmVar {{.*}} '_arg_s1' '__global char *' + // CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize // 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: 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 // 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 ()' lvalue .__finalize // 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 -// 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 // 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 ()' lvalue .__finalize // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array @@ -416,11 +292,8 @@ int main() { // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 -// HasArrayOfHasStreams -// First element -// CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void (sycl::accessor{{.*}})' lvalue .__init -// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc +// CHECK: | |-CXXMemberCallExpr {{.*}} 'void' +// 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 *' @@ -428,8 +301,9 @@ int main() { // 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 ()' lvalue .__finalize // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue .s1 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' @@ -437,10 +311,9 @@ int main() { // CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams' lvalue . // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 -// array: + // 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 @@ -451,8 +324,9 @@ int main() { // 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 ()' lvalue .__finalize // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array @@ -463,10 +337,9 @@ int main() { // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 // 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 @@ -477,8 +350,9 @@ int main() { // 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 ()' lvalue .__finalize // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array @@ -489,10 +363,9 @@ int main() { // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 // 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 {{.*}} '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 *' @@ -500,8 +373,9 @@ int main() { // 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 ()' lvalue .__finalize // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue .s1 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' @@ -509,10 +383,9 @@ int main() { // CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams' lvalue . // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 -// array: + // 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 @@ -523,8 +396,9 @@ int main() { // 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 ()' lvalue .__finalize // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array @@ -535,10 +409,9 @@ int main() { // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 // 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,8 +422,9 @@ 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: MemberExpr {{.*}} 'void ()' lvalue .__finalize // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream *' // CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2]' lvalue .s_array @@ -562,643 +436,320 @@ int main() { // 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 {{.*}} '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 *' // CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs -// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue +// 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 ()' lvalue .__finalize // 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: 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 -// array: + // 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 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' // CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs -// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue +// 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 ()' lvalue .__finalize // 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: 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 -// 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 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' // CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs -// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue +// 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 ()' lvalue .__finalize // 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: 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 -// second element + // 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: 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: 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 ()' lvalue .__finalize // 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: 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 -// array: + // 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 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' // CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs -// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue +// 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 ()' lvalue .__finalize // 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: 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 -// 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 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' // CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs -// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue +// 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 ()' lvalue .__finalize // 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: 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 -// second element + // 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: 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: 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 ()' lvalue .__finalize // 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: 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 -// array: + // 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 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' // CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs -// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue +// 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 ()' lvalue .__finalize // 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: 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 -// 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 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' // CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs -// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue +// 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 ()' lvalue .__finalize // 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: 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 -// 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 *' // 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 -// 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: 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 -// 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: 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: 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 - -// Finalize -// in_lambda __finalize -// CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize -// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue . -// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at - -// _in_lambda_array -// element 0 -// CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize -// 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 -// element 1 -// CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize -// 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 -// _in_lambda_mdarray -// [0][0] -// CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize -// 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 -// [0][1] -// CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize -// 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 -// [1][0] -// CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize -// 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 -// [1][1] -// CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize -// 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 - -// HasStreams -// CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize -// 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 ()' lvalue .__finalize -// 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 -// element 1 -// CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize -// 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 - -// HasArrayOfHasStreams -// First element -// CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize -// 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 -// array: -// CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize -// 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 -// element 1 -// CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize -// 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 -// second element // CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize // 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 -// array: -// CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize -// 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 -// element 1 -// CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize -// 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 ()' lvalue .__finalize -// 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 -// array: -// CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize -// 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 -// element 1 -// CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize -// 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: 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 -// second element -// CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize -// 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 -// array: -// CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize -// 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 -// element 1 + // CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize +// 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 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' // CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs -// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue +// 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 -// second element -// CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize -// 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 -// array: + // CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue @@ -1207,59 +758,30 @@ int main() { // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' // CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs -// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue +// 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 -// element 1 -// CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize -// 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 -// second element -// CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize -// 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 -// array: + // CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize +// 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 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' // CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs -// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue +// 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 -// element 1 +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 + // CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue @@ -1268,7 +790,7 @@ int main() { // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' // CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs -// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasArrayOfHasStreams' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' // CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams [2]' lvalue . // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at From afdc281073721d129ac5f8c0041533e2f99a6c2c Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Tue, 4 May 2021 05:44:55 -0700 Subject: [PATCH 08/20] Review comments fixes Signed-off-by: Zahira Ammarguellat --- clang/lib/Sema/SemaSYCL.cpp | 2 +- clang/test/SemaSYCL/streams.cpp | 608 ++++++++++++++++++-------------- 2 files changed, 342 insertions(+), 268 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index b6a8883a8438e..1c82e891f08bd 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -2547,7 +2547,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { CXXMethodDecl *FinalizeMethod = getMethodByName(RecordDecl, FinalizeMethodName); if (FinalizeMethod) - createSpecialMethodCall(RecordDecl, FinalizeMethodName, BodyStmts); + createSpecialMethodCall(RecordDecl, FinalizeMethodName, FinalizeStmts); removeFieldMemberExpr(FD, Ty); diff --git a/clang/test/SemaSYCL/streams.cpp b/clang/test/SemaSYCL/streams.cpp index d65f21d937d5c..50c780fdefe5a 100644 --- a/clang/test/SemaSYCL/streams.cpp +++ b/clang/test/SemaSYCL/streams.cpp @@ -44,218 +44,210 @@ int main() { return 0; } -// CHECK: FunctionDecl {{.*}} main 'int ()' -// CHECK: FunctionDecl {{.*}}stream_test -// CHECK: InitListExpr {{.*}} '(lambda at +// Function Declaration +// CHECK: FunctionDecl {{.*}}stream_test{{.*}} + +// Initializers: + +// CHECK: InitListExpr {{.*}} '(lambda at +// 'in_lambda' // 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 () noexcept' +// element 1 // 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 () noexcept' +// element 1 // 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 () noexcept' +// element 1 // CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept' - // HasStreams struct // CHECK: InitListExpr {{.*}} 'HasStreams' +// HasStreams::s1 // 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 () noexcept' +// element 1 // CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept' + +// HasArrayOfHasStreams // CHECK-NEXT: InitListExpr {{.*}} 'HasArrayOfHasStreams' +// HasArrayOfHasStreams::i // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar +// HasArrayOfHasStreams::hs // CHECK-NEXT: InitListExpr {{.*}} 'HasStreams [2]' +// HasStreams struct +// CHECK-NEXT: InitListExpr {{.*}} 'HasStreams' +// HasStreams::s1 +// 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 () noexcept' +// element 1 +// 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 () noexcept' +// HasStreams::s_array +// CHECK-NEXT: InitListExpr {{.*}} 'sycl::stream [2]' +// element 0 +// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept' +// element 1 // CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept' // HasArrayOfHasStreams struct // CHECK: InitListExpr {{.*}} 'HasArrayOfHasStreams [2]' // CHECK-NEXT: InitListExpr {{.*}} 'HasArrayOfHasStreams' +// HasArrayOfHasStreams::i // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar - +// HasArrayOfHasStreams::hs // CHECK-NEXT: InitListExpr {{.*}} 'HasStreams [2]' +// HasStreams struct // CHECK-NEXT: InitListExpr {{.*}} 'HasStreams' +// HasStreams::s1 // 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 () noexcept' +// element 1 // 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 () noexcept' +// HasStreams::s_array // CHECK-NEXT: InitListExpr {{.*}} 'sycl::stream [2]' +// element 0 // CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept' +// element 1 // CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept' - // CHECK-NEXT: InitListExpr {{.*}} 'HasArrayOfHasStreams' +// HasArrayOfHasStreams::i // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar - +// HasArrayOfHasStreams::hs // CHECK-NEXT: InitListExpr {{.*}} 'HasStreams [2]' +// HasStreams struct // CHECK-NEXT: InitListExpr {{.*}} 'HasStreams' +// HasStreams::s1 // 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 () noexcept' +// element 1 // 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 () noexcept' +// HasStreams::s_array // CHECK-NEXT: InitListExpr {{.*}} 'sycl::stream [2]' +// element 0 // CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept' +// element 1 // CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept' // Calls to init +// in_lambda __init // CHECK: CXXMemberCallExpr {{.*}} 'void' // 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 +// element 0 // CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize -// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue . -// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at - -//_in_lambda_array -// CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK: 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 .__finalize +// 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 - +// element 1 // CHECK: CXXMemberCallExpr {{.*}} 'void' // 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 .__finalize -// 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 (__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 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream (*)[2]' -// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2][2]' lvalue . +// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2][2]' lvalue // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 - -// CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize -// 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 - - +// [0][1] // CHECK: CXXMemberCallExpr {{.*}} 'void' // 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 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream (*)[2]' -// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2][2]' lvalue . +// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2][2]' lvalue // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 - +// [1][0] // CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize -// 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-NEXT: CXXMemberCallExpr {{.*}} 'void' // 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 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'sycl::stream (*)[2]' -// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2][2]' lvalue . +// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream [2][2]' lvalue // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 - +// [1][1] // CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize -// 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 (__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 // 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 .__finalize -// 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: MemberExpr {{.*}} 'sycl::stream [2][2]' lvalue // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 +// HasStreams // CHECK: CXXMemberCallExpr {{.*}} 'void' // 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 -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global char *' -// CHECK-NEXT: DeclRefExpr {{.*}} '__global char *' lvalue ParmVar {{.*}} '_arg_s1' '__global char *' - -// CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize -// 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 (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue @@ -264,16 +256,7 @@ int main() { // CHECK-NEXT: MemberExpr {{.*}} 'HasStreams' lvalue . // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 - -// CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize -// 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 - +// element 1 // CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue @@ -283,16 +266,9 @@ int main() { // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 +// HasArrayOfHasStreams +// First element // CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize -// 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 (__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 @@ -301,17 +277,7 @@ int main() { // CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams' lvalue . // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 - -// CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize -// 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 - +// array: // CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue @@ -324,9 +290,9 @@ int main() { // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 - +// element 1 // CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize +// 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 @@ -336,23 +302,20 @@ int main() { // CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams' lvalue . // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 -// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 - +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 +// second element // CHECK: CXXMemberCallExpr {{.*}} 'void' // 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 +// 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-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 - +// array: // CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize +// 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 @@ -361,29 +324,36 @@ int main() { // 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-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// element 1 // CHECK: CXXMemberCallExpr {{.*}} 'void' // 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 {{.*}} '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 ()' lvalue .__finalize +// 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 *' // CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs -// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams' lvalue . +// 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 +// array: // CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue @@ -392,24 +362,43 @@ int main() { // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' // CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs -// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams' lvalue . +// 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-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// element 1 // CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize +// 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 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' // CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs -// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams' lvalue . +// 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 +// second element +// CHECK: CXXMemberCallExpr {{.*}} 'void' +// 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 *' +// 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 +// array: // CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue @@ -418,158 +407,240 @@ int main() { // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' // CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs -// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams' lvalue . +// 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-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// element 1 // CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize +// 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 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' // CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs -// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams' lvalue . +// 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 - +// second element // CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' 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 *' // CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs -// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasArrayOfHasStreams' lvalue +// 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 - +// array: // CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize -// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue .s1 +// 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 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' // CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs -// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasArrayOfHasStreams' lvalue +// 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 - +// element 1 // CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' 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 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' // CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs -// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasArrayOfHasStreams' lvalue +// 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-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 - - +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 +// second element // CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize -// 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 {{.*}} '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 *' // CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs -// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasArrayOfHasStreams' lvalue +// 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-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 +// array: // CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' 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 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' // CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs -// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasArrayOfHasStreams' lvalue +// 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-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// element 1 // CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize +// 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 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' // CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs -// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasArrayOfHasStreams' lvalue +// 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 + +// Finalize +// in_lambda __finalize +// CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize +// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at + +// _in_lambda_array +// element 0 +// CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize +// 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 +// element 1 +// CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize +// 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 +// _in_lambda_mdarray +// [0][0] +// CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize +// 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 +// [0][1] +// CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize +// 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 +// [1][0] +// CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize +// 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 +// [1][1] +// CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize +// 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 +// HasStreams // CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize // 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: MemberExpr {{.*}} 'HasStreams' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at +// array: +// CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize +// 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 +// element 1 +// CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize +// 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 +// HasArrayOfHasStreams +// First element // CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize // 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: MemberExpr {{.*}} 'HasArrayOfHasStreams' lvalue . // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 -// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 - +// array: // CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize // 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: MemberExpr {{.*}} 'HasArrayOfHasStreams' lvalue . // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 -// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 - +// element 1 // CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue @@ -578,30 +649,34 @@ int main() { // 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: MemberExpr {{.*}} 'HasArrayOfHasStreams' lvalue . // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 -// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 - +// second element // CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize +// 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 +// array: +// CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize // 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: MemberExpr {{.*}} 'HasArrayOfHasStreams' lvalue . // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at -// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 -// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 - +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// element 1 // CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue @@ -610,88 +685,87 @@ int main() { // 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: MemberExpr {{.*}} 'HasArrayOfHasStreams' lvalue . // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at -// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 +// HasArrayOfHasStreams array +// First element // CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize // 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: 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 +// array: // CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize -// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue .s1 +// 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: 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-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// element 1 // CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize // 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: 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-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 +// second element // CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize -// 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 {{.*}} '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: 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-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 +// array: // CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize // 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: 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-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// element 1 // CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue @@ -700,88 +774,88 @@ int main() { // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' // CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs -// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasArrayOfHasStreams' lvalue +// 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-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 +// second element // CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize // 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: 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 +// array: // CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize -// CHECK-NEXT: MemberExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue .s1 +// 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: 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-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// element 1 // CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize // 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: 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-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 +// second element // CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize -// 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 {{.*}} '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: 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 - +// array: // CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void (__global char *, range<1>, range<1>, id<1>, int)' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize // 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: 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-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// element 1 // CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'sycl::stream':'sycl::stream' lvalue @@ -790,7 +864,7 @@ int main() { // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' // CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs -// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasArrayOfHasStreams' lvalue +// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' // CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams [2]' lvalue . // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at From 872f494a5da918b4342edfd6f4430af5a531d5f1 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Thu, 6 May 2021 10:12:34 -0700 Subject: [PATCH 09/20] Review comments fixes Signed-off-by: Zahira Ammarguellat --- clang/lib/Sema/SemaSYCL.cpp | 6 +++--- sycl/include/CL/sycl/detail/kernel_desc.hpp | 6 +++--- sycl/source/handler.cpp | 6 +++--- 3 files changed, 9 insertions(+), 9 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 183c39548ccdc..fe70bbdb0b39e 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -2544,9 +2544,9 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { const auto *RecordDecl = Ty->getAsCXXRecordDecl(); createSpecialMethodCall(RecordDecl, getInitMethodName(), BodyStmts); - CXXMethodDecl *FinalizeMethod = - getMethodByName(RecordDecl, FinalizeMethodName); - if (FinalizeMethod) + // A finalize-method is expected for stream classes. + if (CXXMethodDecl *FinalizeMethod = + getMethodByName(RecordDecl, FinalizeMethodName)) createSpecialMethodCall(RecordDecl, FinalizeMethodName, FinalizeStmts); removeFieldMemberExpr(FD, Ty); diff --git a/sycl/include/CL/sycl/detail/kernel_desc.hpp b/sycl/include/CL/sycl/detail/kernel_desc.hpp index 413ee66a55cc2..d3a4de78657b3 100644 --- a/sycl/include/CL/sycl/detail/kernel_desc.hpp +++ b/sycl/include/CL/sycl/detail/kernel_desc.hpp @@ -33,9 +33,9 @@ enum class kernel_param_kind_t { kind_accessor = 0, kind_std_layout = 1, // standard layout object parameters kind_sampler = 2, - kind_stream = 3, - kind_pointer = 4, - kind_specialization_constants_buffer = 5, + kind_pointer = 3, + kind_specialization_constants_buffer = 4, + kind_stream = 5, }; // describes a kernel parameter diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index 4b47c65e60bae..560a726836da5 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -294,7 +294,7 @@ void handler::processArg(void *Ptr, const detail::kernel_param_kind_t &Kind, stream *S = static_cast(Ptr); detail::AccessorBaseHost *GBufBase = - (detail::AccessorBaseHost *)&S->GlobalBuf; + static_cast(&S->GlobalBuf); detail::AccessorImplPtr GBufImpl = detail::getSyclObjImpl(*GBufBase); detail::Requirement *GBufReq = GBufImpl.get(); addArgsForGlobalAccessor(GBufReq, Index, IndexShift, Size, @@ -302,7 +302,7 @@ void handler::processArg(void *Ptr, const detail::kernel_param_kind_t &Kind, MNDRDesc.GlobalSize.size(), MArgs); ++IndexShift; detail::AccessorBaseHost *GOffsetBase = - (detail::AccessorBaseHost *)&S->GlobalOffset; + static_cast(&S->GlobalOffset); detail::AccessorImplPtr GOfssetImpl = detail::getSyclObjImpl(*GOffsetBase); detail::Requirement *GOffsetReq = GOfssetImpl.get(); addArgsForGlobalAccessor(GOffsetReq, Index, IndexShift, Size, @@ -310,7 +310,7 @@ void handler::processArg(void *Ptr, const detail::kernel_param_kind_t &Kind, MNDRDesc.GlobalSize.size(), MArgs); ++IndexShift; detail::AccessorBaseHost *GFlushBase = - (detail::AccessorBaseHost *)&S->GlobalFlushBuf; + static_cast(&S->GlobalFlushBuf); detail::AccessorImplPtr GFlushImpl = detail::getSyclObjImpl(*GFlushBase); detail::Requirement *GFlushReq = GFlushImpl.get(); addArgsForGlobalAccessor(GFlushReq, Index, IndexShift, Size, From e979fe1cb36741ea2593fa56882eef969969ca2d Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Thu, 6 May 2021 10:39:52 -0700 Subject: [PATCH 10/20] Review comments fixes Signed-off-by: Zahira Ammarguellat --- sycl/include/CL/sycl/stream.hpp | 2 -- 1 file changed, 2 deletions(-) diff --git a/sycl/include/CL/sycl/stream.hpp b/sycl/include/CL/sycl/stream.hpp index 53349f31c3b2d..4bd46ddfdf6ce 100644 --- a/sycl/include/CL/sycl/stream.hpp +++ b/sycl/include/CL/sycl/stream.hpp @@ -874,14 +874,12 @@ class __SYCL_EXPORT stream { range GlobalFlushAccRange, range GlobalFlushMemRange, id GlobalFlushId, size_t _FlushBufferSize) { -#ifndef __SYCL_EXPLICIT_SIMD__ GlobalBuf.__init(GlobalBufPtr, GlobalBufAccRange, GlobalBufMemRange, GlobalBufId); GlobalOffset.__init(GlobalOffsetPtr, GlobalOffsetAccRange, GlobalOffsetMemRange, GlobalOffsetId); GlobalFlushBuf.__init(GlobalFlushPtr, GlobalFlushAccRange, GlobalFlushMemRange, GlobalFlushId); -#endif FlushBufferSize = _FlushBufferSize; // Calculate offset in the flush buffer for each work item in the global // work space. We need to avoid calling intrinsics to get global id because From c3e0dcc80c29fd03dcad4e1e4595d2cfd705960c Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Thu, 6 May 2021 12:49:40 -0700 Subject: [PATCH 11/20] Review comments fixes Signed-off-by: Zahira Ammarguellat --- clang/include/clang/Basic/DiagnosticSemaKinds.td | 2 ++ clang/lib/Sema/SemaSYCL.cpp | 7 +++++-- sycl/source/handler.cpp | 5 ++--- 3 files changed, 9 insertions(+), 5 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 693e58c1b2c08..a45160350cce9 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -11396,6 +11396,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 this stream class">; // errors of expect.with.probability def err_probability_not_constant_float : Error< diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index fe70bbdb0b39e..44cfefda8c800 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -2544,9 +2544,12 @@ 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 (CXXMethodDecl *FinalizeMethod = - getMethodByName(RecordDecl, FinalizeMethodName)) + if (!FinalizeMethod && Util::isSyclStreamType(Ty)) + SemaRef.Diag(FD->getLocation(), diag::err_sycl_expected_finalize_method); + else createSpecialMethodCall(RecordDecl, FinalizeMethodName, FinalizeStmts); removeFieldMemberExpr(FD, Ty); diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index 560a726836da5..ab6f69c6c6e8e 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -235,9 +235,8 @@ void handler::associateWithHandler(detail::AccessorBaseHost *AccBase, /*index*/ 0); } -static void addArgsForGlobalAccessor(detail::Requirement *AccImpl, - const size_t Index, size_t &IndexShift, - const int Size, +static void addArgsForGlobalAccessor(detail::Requirement *AccImpl, size_t Index, + size_t &IndexShift, int Size, bool IsKernelCreatedFromSource, size_t GlobalSize, vector_class &Args) { From 45cd12f17bc1882756275facdd4fdf36a95064ae Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Thu, 6 May 2021 13:01:44 -0700 Subject: [PATCH 12/20] Review comments fixes Signed-off-by: Zahira Ammarguellat --- clang/include/clang/Basic/DiagnosticSemaKinds.td | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index a45160350cce9..8e7fd49a8d9c4 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -11397,7 +11397,7 @@ def err_sycl_mismatch_group_size "kernel">; def note_sycl_kernel_declared_here : Note<"kernel declared here">; def err_sycl_expected_finalize_method : Error< - "Expected a finalize method for this stream class">; + "expected a 'finalize' method for the 'stream' class">; // errors of expect.with.probability def err_probability_not_constant_float : Error< From 8dad2229eb54fd69027b8957975288829487c711 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Fri, 7 May 2021 06:04:51 -0700 Subject: [PATCH 13/20] Completed resolving conflict Signed-off-by: Zahira Ammarguellat --- sycl/include/CL/sycl/accessor.hpp | 3 --- 1 file changed, 3 deletions(-) diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp index 6f9b56d5ffbcc..5ad393644deca 100755 --- a/sycl/include/CL/sycl/accessor.hpp +++ b/sycl/include/CL/sycl/accessor.hpp @@ -931,13 +931,10 @@ class accessor : #endif // __SYCL_DEVICE_ONLY__ private: -<<<<<<< HEAD friend class sycl::INTEL::gpu::detail::AccessorPrivateProxy; friend class sycl::stream; -======= friend class sycl::ext::intel::experimental::esimd::detail:: AccessorPrivateProxy; ->>>>>>> remote/sycl public: using value_type = DataT; From ea58203046f847f986eebf02d4ba4d977c3f3b71 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Fri, 7 May 2021 06:13:01 -0700 Subject: [PATCH 14/20] Completed resolving conflict Signed-off-by: Zahira Ammarguellat --- sycl/include/CL/sycl/accessor.hpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp index 5ad393644deca..6a7715d620ef7 100755 --- a/sycl/include/CL/sycl/accessor.hpp +++ b/sycl/include/CL/sycl/accessor.hpp @@ -2333,8 +2333,8 @@ host_accessor(buffer, Type1, Type2, Type3, Type4, #endif -} // namespace sycl -} // __SYCL_INLINE_NAMESPACE(cl) +} // namespace gpu +} // namespace INTEL namespace std { template Date: Fri, 7 May 2021 06:53:09 -0700 Subject: [PATCH 15/20] Fixed resolution conflict Signed-off-by: Zahira Ammarguellat --- sycl/include/CL/sycl/accessor.hpp | 7 ++----- 1 file changed, 2 insertions(+), 5 deletions(-) diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp index 6a7715d620ef7..4648c0015e1f7 100755 --- a/sycl/include/CL/sycl/accessor.hpp +++ b/sycl/include/CL/sycl/accessor.hpp @@ -204,8 +204,6 @@ __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { class stream; -namespace INTEL { -namespace gpu { namespace ext { namespace intel { namespace experimental { @@ -931,7 +929,6 @@ class accessor : #endif // __SYCL_DEVICE_ONLY__ private: - friend class sycl::INTEL::gpu::detail::AccessorPrivateProxy; friend class sycl::stream; friend class sycl::ext::intel::experimental::esimd::detail:: AccessorPrivateProxy; @@ -2333,8 +2330,8 @@ host_accessor(buffer, Type1, Type2, Type3, Type4, #endif -} // namespace gpu -} // namespace INTEL +} // namespace sycl +} // __SYCL_INLINE_NAMESPACE(cl) namespace std { template Date: Tue, 18 May 2021 04:07:59 -0700 Subject: [PATCH 16/20] Fixed LIT test Signed-off-by: Zahira Ammarguellat --- clang/test/SemaSYCL/kernel-arg-opt-report.cpp | 30 +++++++++---------- 1 file changed, 15 insertions(+), 15 deletions(-) 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 From 1dbf04cda200c0cc2ee2b5a03bf67edb0bae7a5d Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Tue, 25 May 2021 06:48:41 -0700 Subject: [PATCH 17/20] Fixing ESIMD test failures Signed-off-by: Zahira Ammarguellat --- sycl/source/handler.cpp | 13 +++++++------ 1 file changed, 7 insertions(+), 6 deletions(-) diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index ab6f69c6c6e8e..4c53b7eb5c1a6 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -239,7 +239,8 @@ static void addArgsForGlobalAccessor(detail::Requirement *AccImpl, size_t Index, size_t &IndexShift, int Size, bool IsKernelCreatedFromSource, size_t GlobalSize, - vector_class &Args) { + vector_class &Args, + bool isESIMD) { using detail::kernel_param_kind_t; if (AccImpl->PerWI) AccImpl->resize(GlobalSize); @@ -250,7 +251,7 @@ static void addArgsForGlobalAccessor(detail::Requirement *AccImpl, size_t Index, // 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 (!AccImpl->MIsESIMDAcc && !IsKernelCreatedFromSource) { + if (!isESIMD && !IsKernelCreatedFromSource) { // Dimensionality of the buffer is 1 when dimensionality of the // accessor is 0. const size_t SizeAccField = @@ -298,7 +299,7 @@ void handler::processArg(void *Ptr, const detail::kernel_param_kind_t &Kind, detail::Requirement *GBufReq = GBufImpl.get(); addArgsForGlobalAccessor(GBufReq, Index, IndexShift, Size, IsKernelCreatedFromSource, - MNDRDesc.GlobalSize.size(), MArgs); + MNDRDesc.GlobalSize.size(), MArgs, IsESIMD); ++IndexShift; detail::AccessorBaseHost *GOffsetBase = static_cast(&S->GlobalOffset); @@ -306,7 +307,7 @@ void handler::processArg(void *Ptr, const detail::kernel_param_kind_t &Kind, detail::Requirement *GOffsetReq = GOfssetImpl.get(); addArgsForGlobalAccessor(GOffsetReq, Index, IndexShift, Size, IsKernelCreatedFromSource, - MNDRDesc.GlobalSize.size(), MArgs); + MNDRDesc.GlobalSize.size(), MArgs, IsESIMD); ++IndexShift; detail::AccessorBaseHost *GFlushBase = static_cast(&S->GlobalFlushBuf); @@ -314,7 +315,7 @@ void handler::processArg(void *Ptr, const detail::kernel_param_kind_t &Kind, detail::Requirement *GFlushReq = GFlushImpl.get(); addArgsForGlobalAccessor(GFlushReq, Index, IndexShift, Size, IsKernelCreatedFromSource, - MNDRDesc.GlobalSize.size(), MArgs); + MNDRDesc.GlobalSize.size(), MArgs, IsESIMD); ++IndexShift; MArgs.emplace_back(kernel_param_kind_t::kind_std_layout, &S->FlushBufferSize, sizeof(S->FlushBufferSize), @@ -332,7 +333,7 @@ void handler::processArg(void *Ptr, const detail::kernel_param_kind_t &Kind, detail::Requirement *AccImpl = static_cast(Ptr); addArgsForGlobalAccessor(AccImpl, Index, IndexShift, Size, IsKernelCreatedFromSource, - MNDRDesc.GlobalSize.size(), MArgs); + MNDRDesc.GlobalSize.size(), MArgs, IsESIMD); break; } case access::target::local: { From d5f0fea566ae70a22ce81d9980f5905db63c5b6c Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Tue, 25 May 2021 08:54:33 -0700 Subject: [PATCH 18/20] Fixing LIT failure Signed-off-by: Zahira Ammarguellat --- clang/test/CodeGenSYCL/stream.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/test/CodeGenSYCL/stream.cpp b/clang/test/CodeGenSYCL/stream.cpp index 179acbced4d7e..cf68439ca383f 100644 --- a/clang/test/CodeGenSYCL/stream.cpp +++ b/clang/test/CodeGenSYCL/stream.cpp @@ -11,8 +11,8 @@ // 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]+}}) +// CHECK: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::stream{{.*}} addrspace(4)* align 4 dereferenceable_or_null(16) %4, i8 addrspace(1)* %5, %[[RANGE_TYPE]]* byval(%[[RANGE_TYPE]]) {{.*}}%{{.*}} +// CHECK: call spir_func void @{{.*}}__finalizeEv{{.*}}(%{{.*}}cl::sycl::stream{{.*}}" addrspace(4)* align 4 dereferenceable_or_null(16) %{{[0-9]+}}) #include "Inputs/sycl.hpp" From 84dc696b24fae202b0913d2d53c9ba6fc501c029 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Wed, 26 May 2021 13:27:08 -0700 Subject: [PATCH 19/20] Fix after review Signed-off-by: Zahira Ammarguellat --- clang/lib/Sema/SemaSYCL.cpp | 9 --------- clang/test/CodeGenSYCL/stream.cpp | 12 +++++++++--- clang/test/SemaSYCL/streams.cpp | 5 ++++- 3 files changed, 13 insertions(+), 13 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 6f4c0e41080eb..fac5f23b62d89 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -2372,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)); } diff --git a/clang/test/CodeGenSYCL/stream.cpp b/clang/test/CodeGenSYCL/stream.cpp index cf68439ca383f..eb57aa207bead 100644 --- a/clang/test/CodeGenSYCL/stream.cpp +++ b/clang/test/CodeGenSYCL/stream.cpp @@ -1,9 +1,9 @@ // 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: %[[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_]+]], @@ -11,8 +11,14 @@ // 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)* align 4 dereferenceable_or_null(16) %4, i8 addrspace(1)* %5, %[[RANGE_TYPE]]* byval(%[[RANGE_TYPE]]) {{.*}}%{{.*}} -// CHECK: call spir_func void @{{.*}}__finalizeEv{{.*}}(%{{.*}}cl::sycl::stream{{.*}}" addrspace(4)* align 4 dereferenceable_or_null(16) %{{[0-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/streams.cpp b/clang/test/SemaSYCL/streams.cpp index 50c780fdefe5a..606f77925e13f 100644 --- a/clang/test/SemaSYCL/streams.cpp +++ b/clang/test/SemaSYCL/streams.cpp @@ -114,8 +114,9 @@ int main() { // element 1 // CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept' -// HasArrayOfHasStreams struct +// HasArrayOfHasStreams Array // CHECK: InitListExpr {{.*}} 'HasArrayOfHasStreams [2]' +// // HasArrayOfHasStreams Struct // CHECK-NEXT: InitListExpr {{.*}} 'HasArrayOfHasStreams' // HasArrayOfHasStreams::i // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' @@ -142,6 +143,7 @@ int main() { // CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept' // element 1 // CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::stream':'sycl::stream' 'void () noexcept' +// HasArrayOfHasStreams Struct // CHECK-NEXT: InitListExpr {{.*}} 'HasArrayOfHasStreams' // HasArrayOfHasStreams::i // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' @@ -545,6 +547,7 @@ 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' From 59047b7fb86b415bf484191d0a92f895bc7a839d Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Thu, 27 May 2021 05:52:08 -0700 Subject: [PATCH 20/20] Fix after review and renamed SemaSYCL/streams.cpp to SemaSYCL/stream.cpp Signed-off-by: Zahira Ammarguellat --- clang/lib/Sema/SemaSYCL.cpp | 2 +- clang/test/SemaSYCL/{streams.cpp => stream.cpp} | 0 2 files changed, 1 insertion(+), 1 deletion(-) rename clang/test/SemaSYCL/{streams.cpp => stream.cpp} (100%) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index fac5f23b62d89..d4851208775be 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -2751,7 +2751,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { createSpecialMethodCall(RecordDecl, getInitMethodName(), BodyStmts); CXXMethodDecl *FinalizeMethod = getMethodByName(RecordDecl, FinalizeMethodName); - // A finalize-method is expected for stream classes. + // A finalize-method is expected for stream class. if (!FinalizeMethod && Util::isSyclStreamType(Ty)) SemaRef.Diag(FD->getLocation(), diag::err_sycl_expected_finalize_method); else diff --git a/clang/test/SemaSYCL/streams.cpp b/clang/test/SemaSYCL/stream.cpp similarity index 100% rename from clang/test/SemaSYCL/streams.cpp rename to clang/test/SemaSYCL/stream.cpp