From ade97f5a3893a17bdcc457464bf7f5f6ee50a8e0 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Wed, 28 Apr 2021 12:58:32 -0700 Subject: [PATCH 01/22] 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/22] 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/22] 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/22] 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/22] 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/22] 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/22] 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/22] 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/22] 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/22] 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/22] 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/22] 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/22] 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/22] 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/22] 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: Fri, 7 May 2021 13:58:11 -0700 Subject: [PATCH 16/22] stream changes FE only Signed-off-by: Zahira Ammarguellat --- sycl/include/CL/sycl/accessor.hpp | 2 - sycl/include/CL/sycl/detail/kernel_desc.hpp | 1 - sycl/include/CL/sycl/stream.hpp | 39 +------- sycl/source/detail/scheduler/commands.cpp | 2 - sycl/source/handler.cpp | 103 ++++++-------------- 5 files changed, 32 insertions(+), 115 deletions(-) diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp index 4648c0015e1f7..e998a2ef9a511 100755 --- a/sycl/include/CL/sycl/accessor.hpp +++ b/sycl/include/CL/sycl/accessor.hpp @@ -203,7 +203,6 @@ __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { -class stream; namespace ext { namespace intel { namespace experimental { @@ -929,7 +928,6 @@ class accessor : #endif // __SYCL_DEVICE_ONLY__ private: - friend class sycl::stream; friend class sycl::ext::intel::experimental::esimd::detail:: AccessorPrivateProxy; diff --git a/sycl/include/CL/sycl/detail/kernel_desc.hpp b/sycl/include/CL/sycl/detail/kernel_desc.hpp index d3a4de78657b3..7db5b75386e17 100644 --- a/sycl/include/CL/sycl/detail/kernel_desc.hpp +++ b/sycl/include/CL/sycl/detail/kernel_desc.hpp @@ -35,7 +35,6 @@ enum class kernel_param_kind_t { kind_sampler = 2, kind_pointer = 3, kind_specialization_constants_buffer = 4, - kind_stream = 5, }; // describes a kernel parameter diff --git a/sycl/include/CL/sycl/stream.hpp b/sycl/include/CL/sycl/stream.hpp index 4bd46ddfdf6ce..232e5ee0bae13 100644 --- a/sycl/include/CL/sycl/stream.hpp +++ b/sycl/include/CL/sycl/stream.hpp @@ -68,23 +68,11 @@ using GlobalBufAccessorT = accessor; -constexpr static access::address_space GlobalBufAS = - TargetToAS::AS; -using GlobalBufPtrType = - typename detail::DecoratedType::type *; -constexpr static int GlobalBufDim = 1; - using GlobalOffsetAccessorT = accessor; -constexpr static access::address_space GlobalOffsetAS = - TargetToAS::AS; -using GlobalOffsetPtrType = - typename detail::DecoratedType::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]); @@ -741,11 +729,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; -#endif - // Throws exception in case of invalid input parameters stream(size_t BufferSize, size_t MaxStatementSize, handler &CGH); @@ -862,25 +845,7 @@ class __SYCL_EXPORT stream { } #ifdef __SYCL_DEVICE_ONLY__ - 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) { - GlobalBuf.__init(GlobalBufPtr, GlobalBufAccRange, GlobalBufMemRange, - GlobalBufId); - GlobalOffset.__init(GlobalOffsetPtr, GlobalOffsetAccRange, - GlobalOffsetMemRange, GlobalOffsetId); - GlobalFlushBuf.__init(GlobalFlushPtr, GlobalFlushAccRange, - GlobalFlushMemRange, GlobalFlushId); - FlushBufferSize = _FlushBufferSize; + void __init() { // 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 @@ -906,8 +871,6 @@ 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 4ea9394f2034f..5535e397b2006 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/sycl/source/detail/scheduler/commands.cpp @@ -1661,8 +1661,6 @@ 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 ab6f69c6c6e8e..c7c1f6c06ab2c 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -14,7 +14,6 @@ #include #include #include -#include #include #include #include @@ -235,40 +234,6 @@ void handler::associateWithHandler(detail::AccessorBaseHost *AccBase, /*index*/ 0); } -static void addArgsForGlobalAccessor(detail::Requirement *AccImpl, size_t Index, - size_t &IndexShift, int Size, - bool IsKernelCreatedFromSource, - size_t GlobalSize, - vector_class &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, @@ -288,40 +253,6 @@ 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 = - static_cast(&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 = - static_cast(&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 = - static_cast(&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. @@ -330,9 +261,37 @@ 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); - addArgsForGlobalAccessor(AccImpl, Index, IndexShift, Size, - IsKernelCreatedFromSource, - MNDRDesc.GlobalSize.size(), MArgs); + + // 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); + } break; } case access::target::local: { From bc76ab414d3328475b1cb39987f48ba25e432c4c Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Mon, 17 May 2021 11:17:05 -0700 Subject: [PATCH 17/22] Fix 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 eb9d625ce3c51ab22a88932238b5ba1304958201 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Mon, 17 May 2021 12:37:06 -0700 Subject: [PATCH 18/22] Remove unrelated edit Signed-off-by: Zahira Ammarguellat --- sycl/include/CL/sycl/accessor.hpp | 1 + sycl/include/CL/sycl/stream.hpp | 1 + 2 files changed, 2 insertions(+) diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp index 26bbc242ef8c0..35a9d6e90880d 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 { + namespace ext { namespace intel { namespace experimental { diff --git a/sycl/include/CL/sycl/stream.hpp b/sycl/include/CL/sycl/stream.hpp index 232e5ee0bae13..678fc650668aa 100644 --- a/sycl/include/CL/sycl/stream.hpp +++ b/sycl/include/CL/sycl/stream.hpp @@ -1118,3 +1118,4 @@ template <> struct hash { } }; } // namespace std + From 0494c269ca13157dceef2c9a7277910eb34f5566 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Tue, 18 May 2021 07:25:41 -0700 Subject: [PATCH 19/22] Adding code to fix lit failures in llvm-test-suite Signed-off-by: Zahira Ammarguellat --- sycl/include/CL/sycl/stream.hpp | 40 +++++++++++++++++++++++++++++++-- sycl/source/handler.cpp | 1 + 2 files changed, 39 insertions(+), 2 deletions(-) diff --git a/sycl/include/CL/sycl/stream.hpp b/sycl/include/CL/sycl/stream.hpp index 678fc650668aa..095c497792958 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::DecoratedType::type *; +constexpr static int GlobalBufDim = 1; + using GlobalOffsetAccessorT = accessor; +constexpr static access::address_space GlobalOffsetAS = + TargetToAS::AS; +using GlobalOffsetPtrType = + typename detail::DecoratedType::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,11 @@ inline __width_manipulator__ setw(int Width) { /// \ingroup sycl_api class __SYCL_EXPORT stream { public: +#ifdef __SYCL_DEVICE_ONLY__ + // Default constructor for objects later initialized with __init member. + stream() = default; +#endif + // Throws exception in case of invalid input parameters stream(size_t BufferSize, size_t MaxStatementSize, handler &CGH); @@ -845,7 +862,25 @@ class __SYCL_EXPORT stream { } #ifdef __SYCL_DEVICE_ONLY__ - void __init() { +void __init(detail::GlobalBufPtrType GlobalBufPtr, + range 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) { + GlobalBuf.__init(GlobalBufPtr, GlobalBufAccRange, GlobalBufMemRange, + GlobalBufId); + GlobalOffset.__init(GlobalOffsetPtr, GlobalOffsetAccRange, + GlobalOffsetMemRange, GlobalOffsetId); + GlobalFlushBuf.__init(GlobalFlushPtr, GlobalFlushAccRange, + GlobalFlushMemRange, GlobalFlushId); + FlushBufferSize = _FlushBufferSize; // Calculate offset in the flush buffer for each work item in the global // work space. We need to avoid calling intrinsics to get global id because // when stream is used in a single_task kernel this could cause some @@ -871,6 +906,8 @@ class __SYCL_EXPORT stream { } #endif + friend class handler; + friend const stream &operator<<(const stream &, const char); friend const stream &operator<<(const stream &, const char *); template @@ -1118,4 +1155,3 @@ template <> struct hash { } }; } // namespace std - diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index c7c1f6c06ab2c..9af2278b83b62 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -14,6 +14,7 @@ #include #include #include +#include #include #include #include From 6205a7bc2a6a7452c624a4d4a5e7c5a6b4acf11a Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Tue, 18 May 2021 07:47:41 -0700 Subject: [PATCH 20/22] Indent Signed-off-by: Zahira Ammarguellat --- sycl/include/CL/sycl/stream.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/CL/sycl/stream.hpp b/sycl/include/CL/sycl/stream.hpp index 095c497792958..4bd46ddfdf6ce 100644 --- a/sycl/include/CL/sycl/stream.hpp +++ b/sycl/include/CL/sycl/stream.hpp @@ -862,7 +862,7 @@ class __SYCL_EXPORT stream { } #ifdef __SYCL_DEVICE_ONLY__ -void __init(detail::GlobalBufPtrType GlobalBufPtr, + void __init(detail::GlobalBufPtrType GlobalBufPtr, range GlobalBufAccRange, range GlobalBufMemRange, id GlobalBufId, From 76d02a3412b3ab2a54a246ab6bf4ae94e88c905c Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Tue, 18 May 2021 11:08:11 -0700 Subject: [PATCH 21/22] Remove changes Signed-off-by: Zahira Ammarguellat --- sycl/include/CL/sycl/stream.hpp | 40 ++------------------------------- sycl/source/handler.cpp | 1 - 2 files changed, 2 insertions(+), 39 deletions(-) diff --git a/sycl/include/CL/sycl/stream.hpp b/sycl/include/CL/sycl/stream.hpp index 4bd46ddfdf6ce..678fc650668aa 100644 --- a/sycl/include/CL/sycl/stream.hpp +++ b/sycl/include/CL/sycl/stream.hpp @@ -68,23 +68,11 @@ using GlobalBufAccessorT = accessor; -constexpr static access::address_space GlobalBufAS = - TargetToAS::AS; -using GlobalBufPtrType = - typename detail::DecoratedType::type *; -constexpr static int GlobalBufDim = 1; - using GlobalOffsetAccessorT = accessor; -constexpr static access::address_space GlobalOffsetAS = - TargetToAS::AS; -using GlobalOffsetPtrType = - typename detail::DecoratedType::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]); @@ -741,11 +729,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; -#endif - // Throws exception in case of invalid input parameters stream(size_t BufferSize, size_t MaxStatementSize, handler &CGH); @@ -862,25 +845,7 @@ class __SYCL_EXPORT stream { } #ifdef __SYCL_DEVICE_ONLY__ - 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) { - GlobalBuf.__init(GlobalBufPtr, GlobalBufAccRange, GlobalBufMemRange, - GlobalBufId); - GlobalOffset.__init(GlobalOffsetPtr, GlobalOffsetAccRange, - GlobalOffsetMemRange, GlobalOffsetId); - GlobalFlushBuf.__init(GlobalFlushPtr, GlobalFlushAccRange, - GlobalFlushMemRange, GlobalFlushId); - FlushBufferSize = _FlushBufferSize; + void __init() { // 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 @@ -906,8 +871,6 @@ 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 @@ -1155,3 +1118,4 @@ template <> struct hash { } }; } // namespace std + diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index 9af2278b83b62..c7c1f6c06ab2c 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -14,7 +14,6 @@ #include #include #include -#include #include #include #include From 2f7dbb044d3e12e9553adc21393bbfa1929c4ae2 Mon Sep 17 00:00:00 2001 From: Zahira Ammarguellat Date: Tue, 18 May 2021 13:14:33 -0700 Subject: [PATCH 22/22] Fixing lit tests failures Signed-off-by: Zahira Ammarguellat --- sycl/include/CL/sycl/accessor.hpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp index 35a9d6e90880d..b86c4bbe34714 100755 --- a/sycl/include/CL/sycl/accessor.hpp +++ b/sycl/include/CL/sycl/accessor.hpp @@ -203,7 +203,7 @@ __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { - +class stream; namespace ext { namespace intel { namespace experimental { @@ -930,6 +930,7 @@ class accessor : #endif // __SYCL_DEVICE_ONLY__ private: + friend class sycl::stream; friend class sycl::ext::intel::experimental::esimd::detail:: AccessorPrivateProxy;