From 9ea493710f7f7f44e7027e9b18b5a54f35ead870 Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Fri, 21 Aug 2020 07:41:47 -0700 Subject: [PATCH 01/23] [NFCI][SYCL] WIP On Rework of how arrays function. This is a WIP patch, just for information of Mariya/Elizabeth to see how things are going. The intent of this patch is simply to clean up how arrays work in the SYCL visitors. This patch already changes the visitor in a way that I believe is optimal and has what is necessary to finish this. Integration header generation also seems to be correct. Work on this DID require quite a bit of a re-imagining of how the DeclBodyCreator works. There are a number of small refactors to fix things I found while working that I believe make it easier to maintain this in the future. However, the biggest change here is with the init-list-expr creation. This patch creates them in advance, rather than waiting until we are done with a struct/array/etc. This allows us to generate them as we go. The side effect of this is that the InitExprs array now ONLY contains collection InitListExprs, since it no longer has to assemble them on the other side. TODO List: There are a number of TODOs in the patch itself that should be worked out at one point or another. There are also currently 9 failing lit tests. However, I know that the following doesn't work: - Base Class initialization lists - Base Class MemberExprBases - Array MemberExprBases --- clang/lib/Sema/SemaSYCL.cpp | 726 ++++++++++++++++-------------------- 1 file changed, 324 insertions(+), 402 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index a207c91cef3c1..2fc6f9a7344bb 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -686,10 +686,10 @@ static void addScopeAttrToLocalVars(CXXMethodDecl &F) { /// Return method by name static CXXMethodDecl *getMethodByName(const CXXRecordDecl *CRD, - const std::string &MethodName) { + StringRef MethodName) { CXXMethodDecl *Method; auto It = std::find_if(CRD->methods().begin(), CRD->methods().end(), - [&MethodName](const CXXMethodDecl *Method) { + [MethodName](const CXXMethodDecl *Method) { return Method->getNameAsString() == MethodName; }); Method = (It != CRD->methods().end()) ? *It : nullptr; @@ -830,64 +830,6 @@ class KernelObjVisitor { // return handlers.f(FD, FDTy); \ // })...) - // Implements the 'for-each-visitor' pattern. - template - void VisitElementImpl(CXXRecordDecl *Owner, FieldDecl *ArrayField, - QualType ElementTy, Handlers &... handlers) { - if (Util::isSyclAccessorType(ElementTy)) - KF_FOR_EACH(handleSyclAccessorType, ArrayField, ElementTy); - else if (Util::isSyclStreamType(ElementTy)) - KF_FOR_EACH(handleSyclStreamType, ArrayField, ElementTy); - else if (Util::isSyclSamplerType(ElementTy)) - KF_FOR_EACH(handleSyclSamplerType, ArrayField, ElementTy); - else if (Util::isSyclHalfType(ElementTy)) - KF_FOR_EACH(handleSyclHalfType, ArrayField, ElementTy); - else if (ElementTy->isStructureOrClassType()) - VisitRecord(Owner, ArrayField, ElementTy->getAsCXXRecordDecl(), - handlers...); - else if (ElementTy->isUnionType()) - VisitUnion(Owner, ArrayField, ElementTy->getAsCXXRecordDecl(), - handlers...); - else if (ElementTy->isArrayType()) - VisitArrayElements(ArrayField, ElementTy, handlers...); - else if (ElementTy->isScalarType()) - KF_FOR_EACH(handleScalarType, ArrayField, ElementTy); - } - - template - void VisitFirstElement(CXXRecordDecl *Owner, FieldDecl *ArrayField, - QualType ElementTy, Handlers &... handlers) { - VisitElementImpl(Owner, ArrayField, ElementTy, handlers...); - } - - template - void VisitNthElement(CXXRecordDecl *Owner, FieldDecl *ArrayField, - QualType ElementTy, Handlers &... handlers); - - template - void VisitArrayElements(FieldDecl *FD, QualType FieldTy, - Handlers &... handlers) { - const ConstantArrayType *CAT = - SemaRef.getASTContext().getAsConstantArrayType(FieldTy); - assert(CAT && "Should only be called on constant-size array."); - QualType ET = CAT->getElementType(); - int64_t ElemCount = CAT->getSize().getSExtValue(); - std::initializer_list{(handlers.enterArray(), 0)...}; - - assert(ElemCount > 0 && "SYCL prohibits 0 sized arrays"); - VisitFirstElement(nullptr, FD, ET, handlers...); - (void)std::initializer_list{(handlers.nextElement(ET, 1), 0)...}; - - for (int64_t Count = 1; Count < ElemCount; Count++) { - VisitNthElement(nullptr, FD, ET, handlers...); - (void)std::initializer_list{ - (handlers.nextElement(ET, Count + 1), 0)...}; - } - - (void)std::initializer_list{ - (handlers.leaveArray(FD, ET, ElemCount), 0)...}; - } - // Parent contains the FieldDecl or CXXBaseSpecifier that was used to enter // the Wrapper structure that we're currently visiting. Owner is the parent // type (which doesn't exist in cases where it is a FieldDecl in the @@ -912,8 +854,6 @@ class KernelObjVisitor { clang::CXXRecordDecl::base_class_const_range Range, Handlers &... handlers) { for (const auto &Base : Range) { - (void)std::initializer_list{ - (handlers.enterField(Owner, Base), 0)...}; QualType BaseTy = Base.getType(); // Handle accessor class as base if (Util::isSyclAccessorType(BaseTy)) { @@ -926,8 +866,6 @@ class KernelObjVisitor { } else // For all other bases, visit the record VisitRecord(Owner, Base, BaseTy->getAsCXXRecordDecl(), handlers...); - (void)std::initializer_list{ - (handlers.leaveField(Owner, Base), 0)...}; } } @@ -946,16 +884,99 @@ class KernelObjVisitor { (handlers.enterStruct(Owner, Parent), 0)...}; for (const auto &Field : Wrapper->fields()) { QualType FieldTy = Field->getType(); - (void)std::initializer_list{ - (handlers.enterField(Wrapper, Field), 0)...}; // Required to initialize accessors inside streams. if (Util::isSyclAccessorType(FieldTy)) KF_FOR_EACH(handleSyclAccessorType, Field, FieldTy); - (void)std::initializer_list{ - (handlers.leaveField(Wrapper, Field), 0)...}; } + } + + template + void VisitArrayElementImpl(const CXXRecordDecl *Owner, FieldDecl *ArrayField, + QualType ElementTy, uint64_t Index, + Handlers &... handlers) { (void)std::initializer_list{ - (handlers.leaveStruct(Owner, Parent), 0)...}; + (handlers.nextElement(ElementTy, Index), 0)...}; + VisitField(Owner, ArrayField, ElementTy, handlers...); + } + + template + void VisitFirstArrayElement(const CXXRecordDecl *Owner, FieldDecl *ArrayField, + QualType ElementTy, Handlers &... handlers) { + VisitArrayElementImpl(Owner, ArrayField, ElementTy, 0, handlers...); + } + template + void VisitNthArrayElement(const CXXRecordDecl *Owner, FieldDecl *ArrayField, + QualType ElementTy, uint64_t Index, + Handlers &... handlers); + + template + void VisitArray(const CXXRecordDecl *Owner, FieldDecl *Field, + QualType FieldTy, Handlers &... handlers) { + // Array workflow is: + // handleArrayType + // enterArray + // nextElement + // VisitField (same as before, note that The FieldDecl is the of array + // itself, not the element) + // ... repeat per element, opt-out for duplicates. + // leaveArray + + if (!KF_FOR_EACH(handleArrayType, Field, FieldTy)) + return; + + const ConstantArrayType *CAT = + SemaRef.getASTContext().getAsConstantArrayType(FieldTy); + assert(CAT && "Should only be called on constant-size array."); + QualType ET = CAT->getElementType(); + uint64_t ElemCount = CAT->getSize().getZExtValue(); + assert(ElemCount > 0 && "SYCL prohibits 0 sized arrays"); + + (void)std::initializer_list{(handlers.enterArray(FieldTy, ET), 0)...}; + + VisitFirstArrayElement(Owner, Field, ET, handlers...); + for (uint64_t Index = 1; Index < ElemCount; ++Index) + VisitNthArrayElement(Owner, Field, ET, Index, handlers...); + + (void)std::initializer_list{ + (handlers.leaveArray(FieldTy, ET), 0)...}; + } + + template + void VisitField(const CXXRecordDecl *Owner, FieldDecl *Field, + QualType FieldTy, Handlers &... handlers) { + if (Util::isSyclAccessorType(FieldTy)) + KF_FOR_EACH(handleSyclAccessorType, Field, FieldTy); + else if (Util::isSyclSamplerType(FieldTy)) + KF_FOR_EACH(handleSyclSamplerType, Field, FieldTy); + else if (Util::isSyclHalfType(FieldTy)) + 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. + VisitStreamRecord(Owner, Field, RD, handlers...); + KF_FOR_EACH(handleSyclStreamType, Field, FieldTy); + } else if (FieldTy->isStructureOrClassType()) { + if (KF_FOR_EACH(handleStructType, Field, FieldTy)) { + CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl(); + VisitRecord(Owner, Field, RD, handlers...); + } + } else if (FieldTy->isUnionType()) { + if (KF_FOR_EACH(handleUnionType, Field, FieldTy)) { + CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl(); + VisitUnion(Owner, Field, RD, handlers...); + } + } else if (FieldTy->isReferenceType()) + KF_FOR_EACH(handleReferenceType, Field, FieldTy); + else if (FieldTy->isPointerType()) + KF_FOR_EACH(handlePointerType, Field, FieldTy); + else if (FieldTy->isArrayType()) + VisitArray(Owner, Field, FieldTy, handlers...); + else if (FieldTy->isScalarType() || FieldTy->isVectorType()) + KF_FOR_EACH(handleScalarType, Field, FieldTy); + else + KF_FOR_EACH(handleOtherType, Field, FieldTy); } public: @@ -971,49 +992,8 @@ class KernelObjVisitor { // SyclKernelFieldHandler for the purposes of kernel generation. template void VisitRecordFields(const CXXRecordDecl *Owner, Handlers &... handlers) { - - for (const auto Field : Owner->fields()) { - (void)std::initializer_list{ - (handlers.enterField(Owner, Field), 0)...}; - QualType FieldTy = Field->getType(); - - if (Util::isSyclAccessorType(FieldTy)) - KF_FOR_EACH(handleSyclAccessorType, Field, FieldTy); - else if (Util::isSyclSamplerType(FieldTy)) - KF_FOR_EACH(handleSyclSamplerType, Field, FieldTy); - else if (Util::isSyclHalfType(FieldTy)) - 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. - VisitStreamRecord(Owner, Field, RD, handlers...); - KF_FOR_EACH(handleSyclStreamType, Field, FieldTy); - } else if (FieldTy->isStructureOrClassType()) { - if (KF_FOR_EACH(handleStructType, Field, FieldTy)) { - CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl(); - VisitRecord(Owner, Field, RD, handlers...); - } - } else if (FieldTy->isUnionType()) { - if (KF_FOR_EACH(handleUnionType, Field, FieldTy)) { - CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl(); - VisitUnion(Owner, Field, RD, handlers...); - } - } else if (FieldTy->isReferenceType()) - KF_FOR_EACH(handleReferenceType, Field, FieldTy); - else if (FieldTy->isPointerType()) - KF_FOR_EACH(handlePointerType, Field, FieldTy); - else if (FieldTy->isArrayType()) { - if (KF_FOR_EACH(handleArrayType, Field, FieldTy)) - VisitArrayElements(Field, FieldTy, handlers...); - } else if (FieldTy->isScalarType() || FieldTy->isVectorType()) - KF_FOR_EACH(handleScalarType, Field, FieldTy); - else - KF_FOR_EACH(handleOtherType, Field, FieldTy); - (void)std::initializer_list{ - (handlers.leaveField(Owner, Field), 0)...}; - } + for (const auto Field : Owner->fields()) + VisitField(Owner, Field, Field->getType(), handlers...); } #undef KF_FOR_EACH }; @@ -1023,7 +1003,7 @@ class KernelObjVisitor { class SyclKernelFieldHandlerBase { public: static constexpr const bool VisitUnionBody = false; - static constexpr const bool VisitNthElement = true; + static constexpr const bool VisitNthArrayElement = true; // Mark these virtual so that we can use override in the implementer classes, // despite virtual dispatch never being used. @@ -1076,18 +1056,11 @@ class SyclKernelFieldHandlerBase { virtual bool leaveUnion(const CXXRecordDecl *, FieldDecl *) { return true; } // The following are used for stepping through array elements. + virtual bool enterArray(QualType ArrayTy, QualType ElementTy) { return true; } + virtual bool leaveArray(QualType ArrayTy, QualType ElementTy) { return true; } - virtual bool enterField(const CXXRecordDecl *, const CXXBaseSpecifier &) { - return true; - } - virtual bool leaveField(const CXXRecordDecl *, const CXXBaseSpecifier &) { - return true; - } - virtual bool enterField(const CXXRecordDecl *, FieldDecl *) { return true; } - virtual bool leaveField(const CXXRecordDecl *, FieldDecl *) { return true; } - virtual bool enterArray() { return true; } + // TODO: does this need the index? virtual bool nextElement(QualType, uint64_t) { return true; } - virtual bool leaveArray(FieldDecl *, QualType, int64_t) { return true; } virtual ~SyclKernelFieldHandlerBase() = default; }; @@ -1137,17 +1110,18 @@ void KernelObjVisitor::VisitUnion(const CXXRecordDecl *Owner, ParentTy &Parent, } template -void KernelObjVisitor::VisitNthElement(CXXRecordDecl *Owner, - FieldDecl *ArrayField, - QualType ElementTy, - Handlers &... handlers) { +void KernelObjVisitor::VisitNthArrayElement(const CXXRecordDecl *Owner, + FieldDecl *ArrayField, + QualType ElementTy, + uint64_t Index, + Handlers &... handlers) { // Don't continue descending if none of the handlers 'care'. This could be 'if // constexpr' starting in C++17. Until then, we have to count on the // optimizer to realize "if (false)" is a dead branch. - if (AnyTrue::Value) - VisitElementImpl( - Owner, ArrayField, ElementTy, - HandlerFilter(handlers) + if (AnyTrue::Value) + VisitArrayElementImpl( + Owner, ArrayField, ElementTy, Index, + HandlerFilter(handlers) .Handler...); } @@ -1273,7 +1247,7 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler { public: SyclKernelFieldChecker(Sema &S) : SyclKernelFieldHandler(S), Diag(S.getASTContext().getDiagnostics()) {} - static constexpr const bool VisitNthElement = false; + static constexpr const bool VisitNthArrayElement = false; bool isValid() { return !IsInvalid; } bool handleReferenceType(FieldDecl *FD, QualType FieldTy) final { @@ -1322,7 +1296,7 @@ class SyclKernelUnionChecker : public SyclKernelFieldHandler { : SyclKernelFieldHandler(S), Diag(S.getASTContext().getDiagnostics()) {} bool isValid() { return !IsInvalid; } static constexpr const bool VisitUnionBody = true; - static constexpr const bool VisitNthElement = false; + static constexpr const bool VisitNthArrayElement = false; bool checkType(SourceLocation Loc, QualType Ty) { if (UnionCount) { @@ -1382,10 +1356,6 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { int StructDepth = 0; void addParam(const FieldDecl *FD, QualType FieldTy) { - const ConstantArrayType *CAT = - SemaRef.getASTContext().getAsConstantArrayType(FieldTy); - if (CAT) - FieldTy = CAT->getElementType(); ParamDesc newParamDesc = makeParamDesc(FD, FieldTy); addParam(newParamDesc, FieldTy); } @@ -1657,26 +1627,30 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { using SyclKernelFieldHandler::leaveStruct; }; +// TODO: ERICH: const-correctness of all the functions? class SyclKernelBodyCreator : public SyclKernelFieldHandler { SyclKernelDeclCreator &DeclCreator; llvm::SmallVector BodyStmts; + llvm::SmallVector CollectionInitExprs; llvm::SmallVector FinalizeStmts; - llvm::SmallVector InitExprs; + llvm::SmallVector, 8> ArrayInfos; VarDecl *KernelObjClone; InitializedEntity VarEntity; const CXXRecordDecl *KernelObj; llvm::SmallVector MemberExprBases; FunctionDecl *KernelCallerFunc; + // Contains a count of how many containers we're in. This is used by the + // pointer-struct-wrapping code to ensure that we don't try to wrap + // non-top-level pointers. + uint64_t ContainerDepth = 0; // Using the statements/init expressions that we've created, this generates // the kernel body compound stmt. CompoundStmt needs to know its number of // statements in advance to allocate it, so we cannot do this as we go along. CompoundStmt *createKernelBody() { - - Expr *ILE = new (SemaRef.getASTContext()) InitListExpr( - SemaRef.getASTContext(), SourceLocation(), InitExprs, SourceLocation()); - ILE->setType(QualType(KernelObj->getTypeForDecl(), 0)); - KernelObjClone->setInit(ILE); + assert(CollectionInitExprs.size() == 1 && + "Should have been popped down to just the first one"); + KernelObjClone->setInit(CollectionInitExprs.back()); Stmt *FunctionBody = KernelCallerFunc->getBody(); ParmVarDecl *KernelObjParam = *(KernelCallerFunc->param_begin()); @@ -1727,147 +1701,113 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { } } - MemberExpr *BuildMemberExpr(Expr *Base, ValueDecl *Member) { - DeclAccessPair MemberDAP = DeclAccessPair::make(Member, AS_none); - MemberExpr *Result = SemaRef.BuildMemberExpr( - Base, /*IsArrow */ false, SourceLocation(), NestedNameSpecifierLoc(), - SourceLocation(), Member, MemberDAP, - /*HadMultipleCandidates*/ false, - DeclarationNameInfo(Member->getDeclName(), SourceLocation()), - Member->getType(), VK_LValue, OK_Ordinary); - return Result; + // Creates a DeclRefExpr to the ParmVar that represents the current field. + Expr* createParamReferenceExpr() { + ParmVarDecl *KernelParameter = + DeclCreator.getParamVarDeclsForCurrentField()[0]; + + QualType ParamType = KernelParameter->getOriginalType(); + Expr *DRE = SemaRef.BuildDeclRefExpr(KernelParameter, ParamType, VK_LValue, + SourceLocation()); + return DRE; } - Expr *createInitExpr(FieldDecl *FD) { + // Creates a DeclRefExpr to the ParmVar that represents the current pointer + // field. + Expr* createPointerParamReferenceExpr(QualType PointerTy, bool Wrapped) { ParmVarDecl *KernelParameter = DeclCreator.getParamVarDeclsForCurrentField()[0]; + QualType ParamType = KernelParameter->getOriginalType(); Expr *DRE = SemaRef.BuildDeclRefExpr(KernelParameter, ParamType, VK_LValue, SourceLocation()); - if (FD->getType()->isPointerType()) { - QualType ModifiedType = ParamType; - // Struct Type kernel arguments are decomposed. The pointer fields are - // then wrapped inside a compiler generated struct. Therefore when - // generating the initializers, we have to 'unwrap' the pointer. - if (MemberExprBases.size() > 2) { - CXXRecordDecl *WrapperStruct = ParamType->getAsCXXRecordDecl(); - // Pointer field wrapped inside __wrapper_class - FieldDecl *Pointer = *(WrapperStruct->field_begin()); - DRE = BuildMemberExpr(DRE, Pointer); - ModifiedType = Pointer->getType(); - } - - if (FD->getType()->getPointeeType().getAddressSpace() != - ModifiedType->getPointeeType().getAddressSpace()) - DRE = ImplicitCastExpr::Create(SemaRef.Context, FD->getType(), - CK_AddressSpaceConversion, DRE, nullptr, - VK_RValue); + // Struct Type kernel arguments are decomposed. The pointer fields are + // then wrapped inside a compiler generated struct. Therefore when + // generating the initializers, we have to 'unwrap' the pointer. + if (Wrapped) { + CXXRecordDecl *WrapperStruct = ParamType->getAsCXXRecordDecl(); + // Pointer field wrapped inside __wrapper_class + FieldDecl *Pointer = *(WrapperStruct->field_begin()); + DRE = BuildMemberExpr(DRE, Pointer); + ParamType = Pointer->getType(); } + + if (PointerTy->getPointeeType().getAddressSpace() != + ParamType->getPointeeType().getAddressSpace()) + DRE = ImplicitCastExpr::Create(SemaRef.Context, PointerTy, + CK_AddressSpaceConversion, DRE, nullptr, + VK_RValue); + return DRE; } - void createExprForStructOrScalar(FieldDecl *FD) { - InitializedEntity Entity = - InitializedEntity::InitializeMember(FD, &VarEntity); - InitializationKind InitKind = - InitializationKind::CreateCopy(SourceLocation(), SourceLocation()); - Expr *DRE = createInitExpr(FD); - InitializationSequence InitSeq(SemaRef, Entity, InitKind, DRE); - ExprResult MemberInit = InitSeq.Perform(SemaRef, Entity, InitKind, DRE); - InitExprs.push_back(MemberInit.get()); + // Returns 'true' if the thing we're visiting (Based on the FD/QualType pair) + // is an element of an array. This will determine whether we do + // MemberExprBases in some cases or not, AND determines how we initialize + // values. + bool IsArrayElement(FieldDecl *FD, QualType Ty) { + // TODO, better way to detect that we're in an array? + SemaRef.getASTContext().hasSameType(FD->getType(), Ty); + return FD->getType() != Ty; } - int getDims() { - int Dims = 0; - for (int i = MemberExprBases.size() - 1; i >= 0; --i) { - if (!isa(MemberExprBases[i])) - break; - ++Dims; - } - return Dims; - } - - int64_t getArrayIndex(int Idx) { - ArraySubscriptExpr *LastArrayRef = - cast(MemberExprBases[Idx]); - Expr *LastIdx = LastArrayRef->getIdx(); - llvm::APSInt Result; - SemaRef.VerifyIntegerConstantExpression(LastIdx, &Result); - return Result.getExtValue(); - } - - void createExprForScalarElement(FieldDecl *FD) { - llvm::SmallVector InitEntities; - - // For multi-dimensional arrays, an initialized entity needs to be - // generated for each 'dimension'. For example, the initialized entity - // for s.array[x][y][z] is constructed using initialized entities for - // s.array[x][y], s.array[x] and s.array. InitEntities is used to maintain - // this. - InitializedEntity Entity = - InitializedEntity::InitializeMember(FD, &VarEntity); - InitEntities.push_back(Entity); - - // Calculate dimension using ArraySubscriptExpressions in MemberExprBases. - // Each dimension has an ArraySubscriptExpression (maintains index) - // in MemberExprBases. For example, if we are currently handling element - // a[0][0][1], the top of stack entries are ArraySubscriptExpressions for - // indices 0,0 and 1, with 1 on top. - int Dims = getDims(); - - // MemberExprBasesIdx is used to get the index of each dimension, in correct - // order, from MemberExprBases. For example for a[0][0][1], getArrayIndex - // will return 0, 0 and then 1. - int MemberExprBasesIdx = MemberExprBases.size() - Dims; - for (int I = 0; I < Dims; ++I) { - InitializedEntity NewEntity = InitializedEntity::InitializeElement( - SemaRef.getASTContext(), getArrayIndex(MemberExprBasesIdx), - InitEntities.back()); - InitEntities.push_back(NewEntity); - ++MemberExprBasesIdx; - } + // Creates an initialized entity for a field/item. In the case where this is a + // field, returns a normal member initializer, if we're in a sub-array of a MD + // array, returns an element initializer. + InitializedEntity getFieldEntity(FieldDecl *FD, QualType Ty) { + if (IsArrayElement(FD, Ty)) + return InitializedEntity::InitializeElement(SemaRef.getASTContext(), + ArrayInfos.back().second, + ArrayInfos.back().first); + return InitializedEntity::InitializeMember(FD, &VarEntity); + } + // Should this take a ArrayRef instead, to avoid this foolish ParamRef ? + // ParamRef : None? + void addFieldInit(FieldDecl *FD, QualType Ty, MultiExprArg ParamRef) { + // TODO: Why is this by copy rather than 'forinit' or value init? InitializationKind InitKind = InitializationKind::CreateCopy(SourceLocation(), SourceLocation()); - Expr *DRE = createInitExpr(FD); - InitializationSequence InitSeq(SemaRef, InitEntities.back(), InitKind, DRE); - ExprResult MemberInit = - InitSeq.Perform(SemaRef, InitEntities.back(), InitKind, DRE); - InitExprs.push_back(MemberInit.get()); - } - - void addArrayInit(FieldDecl *FD, int64_t Count) { - llvm::SmallVector ArrayInitExprs; - for (int64_t I = 0; I < Count; I++) { - ArrayInitExprs.push_back(InitExprs.back()); - InitExprs.pop_back(); - } - std::reverse(ArrayInitExprs.begin(), ArrayInitExprs.end()); - Expr *ILE = new (SemaRef.getASTContext()) - InitListExpr(SemaRef.getASTContext(), SourceLocation(), ArrayInitExprs, - SourceLocation()); - - // We need to find the type of the element for which we are generating the - // InitListExpr. For example, for a multi-dimensional array say a[2][3][2], - // the types for InitListExpr of the array and its 'sub-arrays' are - - // int [2][3][2], int [3][2] and int [2]. This loop is used to obtain this - // information from MemberExprBases. MemberExprBases holds - // ArraySubscriptExprs and the top of stack shows how far we have descended - // down the array. getDims() calculates this depth. - QualType ILEType = FD->getType(); - for (int I = getDims(); I > 1; I--) { - const ConstantArrayType *CAT = - SemaRef.getASTContext().getAsConstantArrayType(ILEType); - assert(CAT && "Should only be called on constant-size array."); - ILEType = CAT->getElementType(); - } - ILE->setType(ILEType); - InitExprs.push_back(ILE); + addFieldInit(FD, Ty, ParamRef, InitKind); + } + + void addFieldInit(FieldDecl *FD, QualType Ty, MultiExprArg ParamRef, + InitializationKind InitKind) { + InitializedEntity Entity = getFieldEntity(FD, Ty); + + InitializationSequence InitSeq(SemaRef, Entity, InitKind, ParamRef); + ExprResult Init = + InitSeq.Perform(SemaRef, Entity, InitKind, ParamRef); + + InitListExpr *ParentILE = CollectionInitExprs.back(); + ParentILE->updateInit(SemaRef.getASTContext(), ParentILE->getNumInits(), + Init.get()); + } + + // Adds an initializer that handles a simple initialization of a field. + void addSimpleFieldInit(FieldDecl *FD, QualType Ty) { + Expr *ParamRef = createParamReferenceExpr(); + addFieldInit(FD, Ty, ParamRef); + } + + MemberExpr *BuildMemberExpr(Expr *Base, ValueDecl *Member) { + DeclAccessPair MemberDAP = DeclAccessPair::make(Member, AS_none); + MemberExpr *Result = SemaRef.BuildMemberExpr( + Base, /*IsArrow */ false, SourceLocation(), NestedNameSpecifierLoc(), + SourceLocation(), Member, MemberDAP, + /*HadMultipleCandidates*/ false, + DeclarationNameInfo(Member->getDeclName(), SourceLocation()), + Member->getType(), VK_LValue, OK_Ordinary); + return Result; } - CXXMemberCallExpr *createSpecialMethodCall(Expr *Base, CXXMethodDecl *Method, - FieldDecl *Field) { + void createSpecialMethodCall(const CXXRecordDecl *RD, StringRef MethodName, + SmallVectorImpl &AddTo) { + CXXMethodDecl *Method = getMethodByName(RD, MethodName); + if (!Method) + return; + unsigned NumParams = Method->getNumParams(); llvm::SmallVector ParamDREs(NumParams); llvm::ArrayRef KernelParameters = @@ -1877,7 +1817,8 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { ParamDREs[I] = SemaRef.BuildDeclRefExpr(KernelParameters[I], ParamType, VK_LValue, SourceLocation()); } - MemberExpr *MethodME = BuildMemberExpr(Base, Method); + + MemberExpr *MethodME = BuildMemberExpr(MemberExprBases.back(), Method); QualType ResultTy = Method->getReturnType(); ExprValueKind VK = Expr::getValueKindForType(ResultTy); @@ -1888,10 +1829,46 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { ParamDREs, ParamStmts); // [kernel_obj or wrapper object].accessor.__init(_ValueType*, // range, range, id) - CXXMemberCallExpr *Call = CXXMemberCallExpr::Create( + AddTo.push_back(CXXMemberCallExpr::Create( SemaRef.Context, MethodME, ParamStmts, ResultTy, VK, SourceLocation(), - FPOptionsOverride()); - return Call; + FPOptionsOverride())); + } + + // Creates an empty InitListExpr of the correct number of child-inits + // of this to append into. + void addCollectionInitListExpr(const CXXRecordDecl *RD) { + const ASTRecordLayout &Info = + SemaRef.getASTContext().getASTRecordLayout(RD); + uint64_t NumInitExprs = Info.getFieldCount() + RD->getNumBases(); + addCollectionInitListExpr(QualType(RD->getTypeForDecl(), 0), NumInitExprs); + } + + InitListExpr *CreateInitListExpr(const CXXRecordDecl *RD) { + const ASTRecordLayout &Info = + SemaRef.getASTContext().getASTRecordLayout(RD); + uint64_t NumInitExprs = Info.getFieldCount() + RD->getNumBases(); + return CreateInitListExpr(QualType(RD->getTypeForDecl(), 0), NumInitExprs); + } + + InitListExpr *CreateInitListExpr(QualType InitTy, uint64_t NumChildInits) { + InitListExpr *ILE = new (SemaRef.getASTContext()) InitListExpr( + SemaRef.getASTContext(), SourceLocation(), {}, SourceLocation()); + ILE->reserveInits(SemaRef.getASTContext(), NumChildInits); + ILE->setType(InitTy); + + return ILE; + } + + // Create an empty InitListExpr of the type/size for the rest of the visitor + // to append into. + void addCollectionInitListExpr(QualType InitTy, uint64_t NumChildInits) { + + InitListExpr *ILE = CreateInitListExpr(InitTy, NumChildInits); + InitListExpr *ParentILE = CollectionInitExprs.back(); + ParentILE->updateInit(SemaRef.getASTContext(), ParentILE->getNumInits(), + ILE); + + CollectionInitExprs.push_back(ILE); } // FIXME Avoid creation of kernel obj clone. @@ -1907,25 +1884,17 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { return VD; } + // Default inits the type, then calls the init-method in the body. bool handleSpecialType(FieldDecl *FD, QualType Ty) { + addFieldInit(FD, Ty, None, + InitializationKind::CreateDefault(SourceLocation())); + + // TODO ERICH: if this is in an array, we likely don't want this. + MemberExprBases.push_back(BuildMemberExpr(MemberExprBases.back(), FD)); const auto *RecordDecl = Ty->getAsCXXRecordDecl(); - // TODO: VarEntity is initialized entity for KernelObjClone, I guess we need - // to create new one when enter new struct. - InitializedEntity Entity = - InitializedEntity::InitializeMember(FD, &VarEntity); - // Initialize with the default constructor. - InitializationKind InitKind = - InitializationKind::CreateDefault(SourceLocation()); - InitializationSequence InitSeq(SemaRef, Entity, InitKind, None); - ExprResult MemberInit = InitSeq.Perform(SemaRef, Entity, InitKind, None); - InitExprs.push_back(MemberInit.get()); + createSpecialMethodCall(RecordDecl, InitMethodName, BodyStmts); + MemberExprBases.pop_back(); - CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, InitMethodName); - if (InitMethod) { - CXXMemberCallExpr *InitCall = - createSpecialMethodCall(MemberExprBases.back(), InitMethod, FD); - BodyStmts.push_back(InitCall); - } return true; } @@ -1940,14 +1909,9 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { InitializationKind::CreateDefault(SourceLocation()); InitializationSequence InitSeq(SemaRef, Entity, InitKind, None); ExprResult MemberInit = InitSeq.Perform(SemaRef, Entity, InitKind, None); - InitExprs.push_back(MemberInit.get()); + // TODO: ERICH: FIGURE OUT CollectionInitExprs.push_back(MemberInit.get()); - CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, InitMethodName); - if (InitMethod) { - CXXMemberCallExpr *InitCall = - createSpecialMethodCall(MemberExprBases.back(), InitMethod, nullptr); - BodyStmts.push_back(InitCall); - } + createSpecialMethodCall(RecordDecl, InitMethodName, BodyStmts); return true; } @@ -1960,6 +1924,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { DC.getKernelDecl(), KernelObj)), VarEntity(InitializedEntity::InitializeVariable(KernelObjClone)), KernelObj(KernelObj), KernelCallerFunc(KernelCallerFunc) { + CollectionInitExprs.push_back(CreateInitListExpr(KernelObj)); markParallelWorkItemCalls(); Stmt *DS = new (S.Context) DeclStmt(DeclGroupRef(KernelObjClone), @@ -1973,8 +1938,8 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { } ~SyclKernelBodyCreator() { - CompoundStmt *KernelBody = createKernelBody(); - DeclCreator.setBody(KernelBody); + CompoundStmt *KernelBody = createKernelBody(); + DeclCreator.setBody(KernelBody); } bool handleSyclAccessorType(FieldDecl *FD, QualType Ty) final { @@ -1996,21 +1961,18 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { bool handleSyclStreamType(FieldDecl *FD, QualType Ty) final { const auto *StreamDecl = Ty->getAsCXXRecordDecl(); - createExprForStructOrScalar(FD); - size_t NumBases = MemberExprBases.size(); - CXXMethodDecl *InitMethod = getMethodByName(StreamDecl, InitMethodName); - if (InitMethod) { - CXXMemberCallExpr *InitCall = - createSpecialMethodCall(MemberExprBases.back(), InitMethod, FD); - BodyStmts.push_back(InitCall); - } - CXXMethodDecl *FinalizeMethod = - getMethodByName(StreamDecl, FinalizeMethodName); - if (FinalizeMethod) { - CXXMemberCallExpr *FinalizeCall = createSpecialMethodCall( - MemberExprBases[NumBases - 2], FinalizeMethod, FD); - FinalizeStmts.push_back(FinalizeCall); - } + // Streams just get copied as a new init. + addSimpleFieldInit(FD, Ty); + + // Add a dummy init expression to catch the accessor initializers. + CollectionInitExprs.push_back(CreateInitListExpr(StreamDecl)); + + // Add init/finalize method calls. + // TODO ERICH: if this is in an array, we likely don't want this. + MemberExprBases.push_back(BuildMemberExpr(MemberExprBases.back(), FD)); + createSpecialMethodCall(StreamDecl, InitMethodName, BodyStmts); + createSpecialMethodCall(StreamDecl, FinalizeMethodName, FinalizeStmts); + MemberExprBases.pop_back(); return true; } @@ -2022,31 +1984,55 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { } bool handleSyclHalfType(FieldDecl *FD, QualType Ty) final { - createExprForStructOrScalar(FD); + addSimpleFieldInit(FD, Ty); return true; } bool handlePointerType(FieldDecl *FD, QualType FieldTy) final { - createExprForStructOrScalar(FD); + Expr *PointerRef = + createPointerParamReferenceExpr(FD->getType(), ContainerDepth != 0); + addFieldInit(FD, FieldTy, PointerRef); return true; } bool handleScalarType(FieldDecl *FD, QualType FieldTy) final { - if (dyn_cast(MemberExprBases.back())) - createExprForScalarElement(FD); - else - createExprForStructOrScalar(FD); + addSimpleFieldInit(FD, FieldTy); return true; } bool handleUnionType(FieldDecl *FD, QualType FieldTy) final { - return handleScalarType(FD, FieldTy); + addSimpleFieldInit(FD, FieldTy); + return true; + } + + bool enterStruct(const CXXRecordDecl *RD, FieldDecl *FD) final { + ++ContainerDepth; + // We handle adding a throw-away initializer in handleSyclStreamType since + // the 'default' init needs to stick around, but the accessors that are + // 'children' of it do not. + if (!Util::isSyclStreamType(FD->getType())) + addCollectionInitListExpr(FD->getType()->getAsCXXRecordDecl()); + + MemberExprBases.push_back(BuildMemberExpr(MemberExprBases.back(), FD)); + return true; + } + + bool leaveStruct(const CXXRecordDecl *, FieldDecl *FD) final { + --ContainerDepth; + // If this is a stream, this has popped the 'fake' one that was added in + // handleSyclStreamType, which hasn't been added as a child. + CollectionInitExprs.pop_back(); + MemberExprBases.pop_back(); + return true; } bool enterStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS) final { + ++ContainerDepth; + CXXCastPath BasePath; QualType DerivedTy(RD->getTypeForDecl(), 0); QualType BaseTy = BS.getType(); +// // TODO: Why is this here? Do we think this check could fail? SemaRef.CheckDerivedToBaseConversion(DerivedTy, BaseTy, SourceLocation(), SourceRange(), &BasePath, /*IgnoreBaseAccess*/ true); @@ -2054,106 +2040,47 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { SemaRef.Context, BaseTy, CK_DerivedToBase, MemberExprBases.back(), /* CXXCastPath=*/&BasePath, VK_LValue); MemberExprBases.push_back(Cast); - return true; - } - - void addStructInit(const CXXRecordDecl *RD) { - const ASTRecordLayout &Info = - SemaRef.getASTContext().getASTRecordLayout(RD); - int NumberOfFields = Info.getFieldCount(); - int popOut = NumberOfFields + RD->getNumBases(); - - llvm::SmallVector BaseInitExprs; - for (int I = 0; I < popOut; I++) { - BaseInitExprs.push_back(InitExprs.back()); - InitExprs.pop_back(); - } - std::reverse(BaseInitExprs.begin(), BaseInitExprs.end()); - - Expr *ILE = new (SemaRef.getASTContext()) - InitListExpr(SemaRef.getASTContext(), SourceLocation(), BaseInitExprs, - SourceLocation()); - ILE->setType(QualType(RD->getTypeForDecl(), 0)); - InitExprs.push_back(ILE); - } - bool leaveStruct(const CXXRecordDecl *, FieldDecl *FD) final { - // Handle struct when kernel object field is struct type or array of - // structs. - const CXXRecordDecl *RD = - FD->getType()->getBaseElementTypeUnsafe()->getAsCXXRecordDecl(); - - // Initializers for accessors inside stream not added. - if (!Util::isSyclStreamType(FD->getType())) - addStructInit(RD); - // Pop out unused initializers created in handleSyclAccesorType - // for accessors inside stream class. - else { - for (const auto &Field : RD->fields()) { - QualType FieldTy = Field->getType(); - if (Util::isSyclAccessorType(FieldTy)) - InitExprs.pop_back(); - } - } + addCollectionInitListExpr(BaseTy->getAsCXXRecordDecl()); return true; } bool leaveStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS) final { - const CXXRecordDecl *BaseClass = BS.getType()->getAsCXXRecordDecl(); - addStructInit(BaseClass); + --ContainerDepth; MemberExprBases.pop_back(); + CollectionInitExprs.pop_back(); return true; } - bool enterField(const CXXRecordDecl *RD, FieldDecl *FD) final { - if (!FD->getType()->isReferenceType()) - MemberExprBases.push_back(BuildMemberExpr(MemberExprBases.back(), FD)); + bool handleArrayType(FieldDecl *FD, QualType FieldTy) final { + // Add the 'array info' pair, which correctly initializes the initialized + // entity object. + ArrayInfos.emplace_back(getFieldEntity(FD, FieldTy), 0); return true; } - bool leaveField(const CXXRecordDecl *, FieldDecl *FD) final { - if (!FD->getType()->isReferenceType()) - MemberExprBases.pop_back(); + bool enterArray(QualType ArrayType, QualType ElementType) final { + uint64_t ArraySize = SemaRef.getASTContext() + .getAsConstantArrayType(ArrayType) + ->getSize() + .getZExtValue(); + addCollectionInitListExpr(ArrayType, ArraySize); return true; } - bool enterArray() final { - Expr *ArrayBase = MemberExprBases.back(); - ExprResult IndexExpr = SemaRef.ActOnIntegerConstant(SourceLocation(), 0); - ExprResult ElementBase = SemaRef.CreateBuiltinArraySubscriptExpr( - ArrayBase, SourceLocation(), IndexExpr.get(), SourceLocation()); - MemberExprBases.push_back(ElementBase.get()); + bool nextElement(QualType, uint64_t Index) final { + ArrayInfos.back().second = Index; return true; } - bool nextElement(QualType ET, uint64_t) final { - // Top of MemberExprBases holds ArraySubscriptExpression of element - // we just handled, or the Array base for the dimension we are - // currently visiting. - int64_t nextIndex = getArrayIndex(MemberExprBases.size() - 1) + 1; - MemberExprBases.pop_back(); - Expr *ArrayBase = MemberExprBases.back(); - ExprResult IndexExpr = - SemaRef.ActOnIntegerConstant(SourceLocation(), nextIndex); - ExprResult ElementBase = SemaRef.CreateBuiltinArraySubscriptExpr( - ArrayBase, SourceLocation(), IndexExpr.get(), SourceLocation()); - MemberExprBases.push_back(ElementBase.get()); + bool leaveArray(QualType ArrayType, QualType ElementType) final { + CollectionInitExprs.pop_back(); + ArrayInfos.pop_back(); return true; } - bool leaveArray(FieldDecl *FD, QualType, int64_t Count) final { - addArrayInit(FD, Count); - MemberExprBases.pop_back(); - return true; - } - - using SyclKernelFieldHandler::enterArray; - using SyclKernelFieldHandler::enterField; - using SyclKernelFieldHandler::enterStruct; using SyclKernelFieldHandler::handleSyclHalfType; using SyclKernelFieldHandler::handleSyclSamplerType; - using SyclKernelFieldHandler::leaveField; - using SyclKernelFieldHandler::leaveStruct; }; class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { @@ -2176,10 +2103,6 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { void addParam(const FieldDecl *FD, QualType ArgTy, SYCLIntegrationHeader::kernel_param_kind_t Kind) { uint64_t Size; - const ConstantArrayType *CAT = - SemaRef.getASTContext().getAsConstantArrayType(ArgTy); - if (CAT) - ArgTy = CAT->getElementType(); Size = SemaRef.getASTContext().getTypeSizeInChars(ArgTy).getQuantity(); Header.addParamDesc(Kind, static_cast(Size), static_cast(CurOffset + offsetOf(FD))); @@ -2307,20 +2230,19 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { return true; } - bool enterArray() final { + bool enterArray(QualType, QualType) final { ArrayBaseOffsets.push_back(CurOffset); return true; } bool nextElement(QualType ET, uint64_t Index) final { int64_t Size = SemaRef.getASTContext().getTypeSizeInChars(ET).getQuantity(); - CurOffset = ArrayBaseOffsets.back() + Size * (Index); + CurOffset = ArrayBaseOffsets.back() + Size * Index; return true; } - bool leaveArray(FieldDecl *, QualType ET, int64_t) final { - CurOffset = ArrayBaseOffsets.back(); - ArrayBaseOffsets.pop_back(); + bool leaveArray(QualType, QualType) final { + CurOffset = ArrayBaseOffsets.pop_back_val(); return true; } using SyclKernelFieldHandler::enterStruct; From 72b52d84932cc759e880b68611068df303707fc8 Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Tue, 25 Aug 2020 13:54:49 -0700 Subject: [PATCH 02/23] Add QualType to EnterStruct/LeaveStruct so that it properly knows how to create its base --- clang/lib/Sema/SemaSYCL.cpp | 65 +++++++++++++++++++++++-------------- 1 file changed, 40 insertions(+), 25 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 2fc6f9a7344bb..4ac139d07ce0e 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -836,13 +836,14 @@ class KernelObjVisitor { // 'root'), and Wrapper is the current struct being unwrapped. template void VisitRecord(const CXXRecordDecl *Owner, ParentTy &Parent, - const CXXRecordDecl *Wrapper, Handlers &... handlers) { + const CXXRecordDecl *Wrapper, QualType RecordTy, + Handlers &... handlers) { (void)std::initializer_list{ - (handlers.enterStruct(Owner, Parent), 0)...}; + (handlers.enterStruct(Owner, Parent, RecordTy), 0)...}; VisitRecordHelper(Wrapper, Wrapper->bases(), handlers...); VisitRecordHelper(Wrapper, Wrapper->fields(), handlers...); (void)std::initializer_list{ - (handlers.leaveStruct(Owner, Parent), 0)...}; + (handlers.leaveStruct(Owner, Parent, RecordTy), 0)...}; } template @@ -865,7 +866,8 @@ class KernelObjVisitor { (handlers.handleSyclStreamType(Owner, Base, BaseTy), 0)...}; } else // For all other bases, visit the record - VisitRecord(Owner, Base, BaseTy->getAsCXXRecordDecl(), handlers...); + VisitRecord(Owner, Base, BaseTy->getAsCXXRecordDecl(), BaseTy, + handlers...); } } @@ -879,15 +881,18 @@ class KernelObjVisitor { // FIXME: Can this be refactored/handled some other way? template void VisitStreamRecord(const CXXRecordDecl *Owner, ParentTy &Parent, - CXXRecordDecl *Wrapper, Handlers &... handlers) { + CXXRecordDecl *Wrapper, QualType RecordTy, + Handlers &... handlers) { (void)std::initializer_list{ - (handlers.enterStruct(Owner, Parent), 0)...}; + (handlers.enterStruct(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.leaveStruct(Owner, Parent, RecordTy), 0)...}; } template @@ -955,12 +960,12 @@ class KernelObjVisitor { else if (Util::isSyclStreamType(FieldTy)) { CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl(); // Handle accessors in stream class. - VisitStreamRecord(Owner, Field, RD, handlers...); + VisitStreamRecord(Owner, Field, RD, FieldTy, handlers...); KF_FOR_EACH(handleSyclStreamType, Field, FieldTy); } else if (FieldTy->isStructureOrClassType()) { if (KF_FOR_EACH(handleStructType, Field, FieldTy)) { CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl(); - VisitRecord(Owner, Field, RD, handlers...); + VisitRecord(Owner, Field, RD, FieldTy, handlers...); } } else if (FieldTy->isUnionType()) { if (KF_FOR_EACH(handleUnionType, Field, FieldTy)) { @@ -1044,14 +1049,21 @@ class SyclKernelFieldHandlerBase { // class/field graph. Int Headers use this to calculate offset, most others // don't have a need for these. - virtual bool enterStruct(const CXXRecordDecl *, FieldDecl *) { return true; } - virtual bool leaveStruct(const CXXRecordDecl *, FieldDecl *) { return true; } - virtual bool enterStruct(const CXXRecordDecl *, const CXXBaseSpecifier &) { + virtual bool enterStruct(const CXXRecordDecl *, FieldDecl *, QualType) { + return true; + } + virtual bool leaveStruct(const CXXRecordDecl *, FieldDecl *, QualType) { + return true; + } + virtual bool enterStruct(const CXXRecordDecl *, const CXXBaseSpecifier &, + QualType) { return true; } - virtual bool leaveStruct(const CXXRecordDecl *, const CXXBaseSpecifier &) { + virtual bool leaveStruct(const CXXRecordDecl *, const CXXBaseSpecifier &, + QualType) { return true; } + // TODO: Does enter-union need to be worried when it is in an array?! virtual bool enterUnion(const CXXRecordDecl *, FieldDecl *) { return true; } virtual bool leaveUnion(const CXXRecordDecl *, FieldDecl *) { return true; } @@ -1059,7 +1071,6 @@ class SyclKernelFieldHandlerBase { virtual bool enterArray(QualType ArrayTy, QualType ElementTy) { return true; } virtual bool leaveArray(QualType ArrayTy, QualType ElementTy) { return true; } - // TODO: does this need the index? virtual bool nextElement(QualType, uint64_t) { return true; } virtual ~SyclKernelFieldHandlerBase() = default; @@ -1502,12 +1513,12 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { SemaRef.addSyclDeviceDecl(KernelDecl); } - bool enterStruct(const CXXRecordDecl *, FieldDecl *) final { + bool enterStruct(const CXXRecordDecl *, FieldDecl *, QualType) final { ++StructDepth; return true; } - bool leaveStruct(const CXXRecordDecl *, FieldDecl *) final { + bool leaveStruct(const CXXRecordDecl *, FieldDecl *, QualType) final { --StructDepth; return true; } @@ -2005,19 +2016,19 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { return true; } - bool enterStruct(const CXXRecordDecl *RD, FieldDecl *FD) final { + bool enterStruct(const CXXRecordDecl *RD, FieldDecl *FD, QualType Ty) final { ++ContainerDepth; // We handle adding a throw-away initializer in handleSyclStreamType since // the 'default' init needs to stick around, but the accessors that are // 'children' of it do not. - if (!Util::isSyclStreamType(FD->getType())) - addCollectionInitListExpr(FD->getType()->getAsCXXRecordDecl()); + if (!Util::isSyclStreamType(Ty)) + addCollectionInitListExpr(Ty->getAsCXXRecordDecl()); MemberExprBases.push_back(BuildMemberExpr(MemberExprBases.back(), FD)); return true; } - bool leaveStruct(const CXXRecordDecl *, FieldDecl *FD) final { + bool leaveStruct(const CXXRecordDecl *, FieldDecl *FD, QualType) final { --ContainerDepth; // If this is a stream, this has popped the 'fake' one that was added in // handleSyclStreamType, which hasn't been added as a child. @@ -2026,7 +2037,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { return true; } - bool enterStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS) final { + bool enterStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS, QualType) final { ++ContainerDepth; CXXCastPath BasePath; @@ -2045,7 +2056,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { return true; } - bool leaveStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS) final { + bool leaveStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS, QualType) final { --ContainerDepth; MemberExprBases.pop_back(); CollectionInitExprs.pop_back(); @@ -2208,24 +2219,28 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { return true; } - bool enterStruct(const CXXRecordDecl *, FieldDecl *FD) final { + bool enterStruct(const CXXRecordDecl *, FieldDecl *FD, QualType) final { ++StructDepth; + // TODO: Is this right?! I think this only needs to be incremented when we + // aren't in an array, otherwise 'enterArray's base offsets should handle + // this right. Otherwise an array of structs is going to be in the middle + // of nowhere. CurOffset += offsetOf(FD); return true; } - bool leaveStruct(const CXXRecordDecl *, FieldDecl *FD) final { + bool leaveStruct(const CXXRecordDecl *, FieldDecl *FD, QualType) final { --StructDepth; CurOffset -= offsetOf(FD); return true; } - bool enterStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS) final { + bool enterStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS, QualType) final { CurOffset += offsetOf(RD, BS.getType()->getAsCXXRecordDecl()); return true; } - bool leaveStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS) final { + bool leaveStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS, QualType) final { CurOffset -= offsetOf(RD, BS.getType()->getAsCXXRecordDecl()); return true; } From eee0bffd37d838bf117774bf311f3b9899872f66 Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Tue, 25 Aug 2020 14:32:33 -0700 Subject: [PATCH 03/23] Get special base types to work --- clang/lib/Sema/SemaSYCL.cpp | 28 +++++++++++++++------------- 1 file changed, 15 insertions(+), 13 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 4ac139d07ce0e..e48b6ceff1381 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1774,8 +1774,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { return InitializedEntity::InitializeMember(FD, &VarEntity); } - // Should this take a ArrayRef instead, to avoid this foolish ParamRef ? - // ParamRef : None? + // TODO: This name seems outdated now :) void addFieldInit(FieldDecl *FD, QualType Ty, MultiExprArg ParamRef) { // TODO: Why is this by copy rather than 'forinit' or value init? InitializationKind InitKind = @@ -1796,6 +1795,18 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { Init.get()); } + void addBaseInit(const CXXBaseSpecifier &BS, QualType Ty, + InitializationKind InitKind) { + InitializedEntity Entity = InitializedEntity::InitializeBase( + SemaRef.Context, &BS, /*IsInheritedVirtualBase*/ false, &VarEntity); + InitializationSequence InitSeq(SemaRef, Entity, InitKind, None); + ExprResult Init = InitSeq.Perform(SemaRef, Entity, InitKind, None); + + InitListExpr *ParentILE = CollectionInitExprs.back(); + ParentILE->updateInit(SemaRef.getASTContext(), ParentILE->getNumInits(), + Init.get()); + } + // Adds an initializer that handles a simple initialization of a field. void addSimpleFieldInit(FieldDecl *FD, QualType Ty) { Expr *ParamRef = createParamReferenceExpr(); @@ -1911,17 +1922,8 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { bool handleSpecialType(const CXXBaseSpecifier &BS, QualType Ty) { const auto *RecordDecl = Ty->getAsCXXRecordDecl(); - // TODO: VarEntity is initialized entity for KernelObjClone, I guess we need - // to create new one when enter new struct. - InitializedEntity Entity = InitializedEntity::InitializeBase( - SemaRef.Context, &BS, /*IsInheritedVirtualBase*/ false, &VarEntity); - // Initialize with the default constructor. - InitializationKind InitKind = - InitializationKind::CreateDefault(SourceLocation()); - InitializationSequence InitSeq(SemaRef, Entity, InitKind, None); - ExprResult MemberInit = InitSeq.Perform(SemaRef, Entity, InitKind, None); - // TODO: ERICH: FIGURE OUT CollectionInitExprs.push_back(MemberInit.get()); - + addBaseInit(BS, Ty, InitializationKind::CreateDefault(SourceLocation())); + // TODO: Needs MemberExprBases entry? createSpecialMethodCall(RecordDecl, InitMethodName, BodyStmts); return true; } From 77b12c721f2b9ea57ebab1151d83b9aae76d9231 Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Tue, 25 Aug 2020 17:28:25 -0700 Subject: [PATCH 04/23] Fix Array Kernel Param test-- I think this test was just wrong, it seemed to be checking that we were constructing 2 arrays of accessors instead of accessors themselves. This patch changes those to be just the accessors themselves, which are inside an array of 2 accessors. --- clang/test/SemaSYCL/array-kernel-param.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/test/SemaSYCL/array-kernel-param.cpp b/clang/test/SemaSYCL/array-kernel-param.cpp index ab5df5329b890..02987fcfb4293 100644 --- a/clang/test/SemaSYCL/array-kernel-param.cpp +++ b/clang/test/SemaSYCL/array-kernel-param.cpp @@ -123,8 +123,8 @@ int main() { // CHECK-NEXT: InitListExpr {{.*}} '(lambda at {{.*}}array-kernel-param.cpp{{.*}})' // CHECK-NEXT: InitListExpr {{.*}} 'struct_acc_t' // CHECK-NEXT: InitListExpr {{.*}} 'Accessor [2]' -// CHECK-NEXT: CXXConstructExpr {{.*}} 'Accessor [2]' -// CHECK-NEXT: CXXConstructExpr {{.*}} 'Accessor [2]' +// CHECK-NEXT: CXXConstructExpr {{.*}} 'Accessor' +// CHECK-NEXT: CXXConstructExpr {{.*}} 'Accessor' // Check __init functions are called // CHECK: CXXMemberCallExpr {{.*}} 'void' From 25edc1ff4e12ac8000b0135127bac1f6509f8cb8 Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Wed, 26 Aug 2020 10:47:47 -0700 Subject: [PATCH 05/23] Fix streams but putting 'handle' first to match the rest of the collections --- clang/lib/Sema/SemaSYCL.cpp | 74 +++++++++++++++++++++++++++---------- 1 file changed, 55 insertions(+), 19 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index e48b6ceff1381..548f37d1b7ea7 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -916,7 +916,7 @@ class KernelObjVisitor { template void VisitArray(const CXXRecordDecl *Owner, FieldDecl *Field, - QualType FieldTy, Handlers &... handlers) { + QualType ArrayTy, Handlers &... handlers) { // Array workflow is: // handleArrayType // enterArray @@ -926,24 +926,25 @@ class KernelObjVisitor { // ... repeat per element, opt-out for duplicates. // leaveArray - if (!KF_FOR_EACH(handleArrayType, Field, FieldTy)) + if (!KF_FOR_EACH(handleArrayType, Field, ArrayTy)) return; const ConstantArrayType *CAT = - SemaRef.getASTContext().getAsConstantArrayType(FieldTy); + SemaRef.getASTContext().getAsConstantArrayType(ArrayTy); assert(CAT && "Should only be called on constant-size array."); QualType ET = CAT->getElementType(); uint64_t ElemCount = CAT->getSize().getZExtValue(); assert(ElemCount > 0 && "SYCL prohibits 0 sized arrays"); - (void)std::initializer_list{(handlers.enterArray(FieldTy, ET), 0)...}; + (void)std::initializer_list{ + (handlers.enterArray(Field, ArrayTy, ET), 0)...}; VisitFirstArrayElement(Owner, Field, ET, handlers...); for (uint64_t Index = 1; Index < ElemCount; ++Index) VisitNthArrayElement(Owner, Field, ET, Index, handlers...); (void)std::initializer_list{ - (handlers.leaveArray(FieldTy, ET), 0)...}; + (handlers.leaveArray(Field, ArrayTy, ET), 0)...}; } template @@ -960,8 +961,8 @@ class KernelObjVisitor { else if (Util::isSyclStreamType(FieldTy)) { CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl(); // Handle accessors in stream class. - VisitStreamRecord(Owner, Field, RD, FieldTy, handlers...); KF_FOR_EACH(handleSyclStreamType, Field, FieldTy); + VisitStreamRecord(Owner, Field, RD, FieldTy, handlers...); } else if (FieldTy->isStructureOrClassType()) { if (KF_FOR_EACH(handleStructType, Field, FieldTy)) { CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl(); @@ -1068,8 +1069,12 @@ class SyclKernelFieldHandlerBase { virtual bool leaveUnion(const CXXRecordDecl *, FieldDecl *) { return true; } // The following are used for stepping through array elements. - virtual bool enterArray(QualType ArrayTy, QualType ElementTy) { return true; } - virtual bool leaveArray(QualType ArrayTy, QualType ElementTy) { return true; } + virtual bool enterArray(FieldDecl *, QualType ArrayTy, QualType ElementTy) { + return true; + } + virtual bool leaveArray(FieldDecl *, QualType ArrayTy, QualType ElementTy) { + return true; + } virtual bool nextElement(QualType, uint64_t) { return true; } @@ -2065,30 +2070,61 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { return true; } - bool handleArrayType(FieldDecl *FD, QualType FieldTy) final { - // Add the 'array info' pair, which correctly initializes the initialized - // entity object. - ArrayInfos.emplace_back(getFieldEntity(FD, FieldTy), 0); - return true; - } - - bool enterArray(QualType ArrayType, QualType ElementType) final { + bool enterArray(FieldDecl *FD, QualType ArrayType, + QualType ElementType) final { uint64_t ArraySize = SemaRef.getASTContext() .getAsConstantArrayType(ArrayType) ->getSize() .getZExtValue(); addCollectionInitListExpr(ArrayType, ArraySize); + ArrayInfos.emplace_back(getFieldEntity(FD, ArrayType), 0); + + // If this is the top-level array, we need to make a MemberExpr in addition + // to an array subscript. + if (!IsArrayElement(FD, ArrayType)) + MemberExprBases.push_back(BuildMemberExpr(MemberExprBases.back(), FD)); return true; } bool nextElement(QualType, uint64_t Index) final { ArrayInfos.back().second = Index; + + // Pop off the last member expr base. + if (Index != 0) MemberExprBases.pop_back(); + + QualType SizeT = SemaRef.getASTContext().getSizeType(); + + llvm::APInt IndexVal{ + static_cast(SemaRef.getASTContext().getTypeSize(SizeT)), + Index, SizeT->isSignedIntegerType()}; + + auto IndexLiteral = IntegerLiteral::Create( + SemaRef.getASTContext(), IndexVal, SizeT, SourceLocation()); + + ExprResult IndexExpr = SemaRef.CreateBuiltinArraySubscriptExpr( + MemberExprBases.back(), SourceLocation{}, IndexLiteral, + SourceLocation{}); + + assert(!IndexExpr.isInvalid()); + MemberExprBases.push_back(IndexExpr.get()); return true; } - bool leaveArray(QualType ArrayType, QualType ElementType) final { + bool leaveArray(FieldDecl *FD, QualType ArrayType, + QualType ElementType) final { CollectionInitExprs.pop_back(); ArrayInfos.pop_back(); + + assert( + !SemaRef.getASTContext().getAsConstantArrayType(ArrayType)->getSize() == + 0 && + "Constant arrays must have at least 1 element"); + // Remove the IndexExpr. + MemberExprBases.pop_back(); + + // Remove the field access expr as well. + if (!IsArrayElement(FD, ArrayType)) + MemberExprBases.pop_back(); return true; } @@ -2247,7 +2283,7 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { return true; } - bool enterArray(QualType, QualType) final { + bool enterArray(FieldDecl*, QualType, QualType) final { ArrayBaseOffsets.push_back(CurOffset); return true; } @@ -2258,7 +2294,7 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { return true; } - bool leaveArray(QualType, QualType) final { + bool leaveArray(FieldDecl*, QualType, QualType) final { CurOffset = ArrayBaseOffsets.pop_back_val(); return true; } From 9b14462b839aafba825505a2aa568cb59aee3bca Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Wed, 26 Aug 2020 11:47:21 -0700 Subject: [PATCH 06/23] Fix init call for members inside an array, update the test to better reflect what is going on. Note that the test was wrong I think before? --- clang/lib/Sema/SemaSYCL.cpp | 18 ++++++---- .../CodeGenSYCL/kernel-param-acc-array.cpp | 33 ++++++++++--------- 2 files changed, 29 insertions(+), 22 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 548f37d1b7ea7..95c578dc2fd6e 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1916,11 +1916,14 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { addFieldInit(FD, Ty, None, InitializationKind::CreateDefault(SourceLocation())); - // TODO ERICH: if this is in an array, we likely don't want this. - MemberExprBases.push_back(BuildMemberExpr(MemberExprBases.back(), FD)); + if (!IsArrayElement(FD, Ty)) + MemberExprBases.push_back(BuildMemberExpr(MemberExprBases.back(), FD)); + const auto *RecordDecl = Ty->getAsCXXRecordDecl(); createSpecialMethodCall(RecordDecl, InitMethodName, BodyStmts); - MemberExprBases.pop_back(); + + if (!IsArrayElement(FD, Ty)) + MemberExprBases.pop_back(); return true; } @@ -1986,11 +1989,14 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { CollectionInitExprs.push_back(CreateInitListExpr(StreamDecl)); // Add init/finalize method calls. - // TODO ERICH: if this is in an array, we likely don't want this. - MemberExprBases.push_back(BuildMemberExpr(MemberExprBases.back(), FD)); + if (!IsArrayElement(FD, Ty)) + MemberExprBases.push_back(BuildMemberExpr(MemberExprBases.back(), FD)); + createSpecialMethodCall(StreamDecl, InitMethodName, BodyStmts); createSpecialMethodCall(StreamDecl, FinalizeMethodName, FinalizeStmts); - MemberExprBases.pop_back(); + + if (!IsArrayElement(FD, Ty)) + MemberExprBases.pop_back(); return true; } diff --git a/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp b/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp index cf17a9d7e3e83..9032dc7fe3f02 100644 --- a/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp +++ b/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp @@ -52,27 +52,28 @@ int main() { // CHECK accessor array default inits // CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0 // CHECK: [[BEGIN:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR:.*]]], [2 x [[ACCESSOR]]]* [[ACCESSOR_ARRAY1]], i64 0, i64 0 -// CHECK: [[END:%[a-zA-Z0-9._]*]] = getelementptr inbounds [[ACCESSOR:.*]], [[ACCESSOR]]* [[BEGIN]], i64 2 -// CHECK: [[NEXT0:%[a-zA-Z0-9._]*]] = getelementptr inbounds [[ACCESSOR]], [[ACCESSOR]]* {{.*}}, i64 1 -// CHECK: [[ELEMENT:%[a-zA-Z0-9._]*]] = getelementptr inbounds [[ACCESSOR]], [[ACCESSOR]]* {{.*}}, i64 1 -// CHECK: [[ELEMENT:%[a-zA-Z0-9._]*]] = getelementptr inbounds [[ACCESSOR]], [[ACCESSOR]]* {{.*}}, i64 2 -// CHECK: [[NEXT1:%[a-zA-Z0-9._]*]] = getelementptr inbounds [[ACCESSOR]], [[ACCESSOR]]* {{.*}}, i64 1 - -// CHECK: [[ACCESSOR_ARRAY2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0 -// CHECK: [[INDEX:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], [2 x [[ACCESSOR]]]* [[ACCESSOR_ARRAY2]], i64 0, i64 0 +// Clang takes advantage of element 1 having the same address as the array, so it doesn't do a GEP. +// CHECK: [[ELEM1_ASCAST:%[a-zA-Z0-9_.]+]] = addrspacecast [[ACCESSOR]]* [[BEGIN]] to [[ACCESSOR]] addrspace(4)* +// CTOR Call #1 +// CHECK: call spir_func void @{{.+}}([[ACCESSOR]] addrspace(4)* [[ELEM1_ASCAST]]) +// CHECK: [[ELEM2_GEP:%[a-zA-Z0-9_.]+]] = getelementptr inbounds [[ACCESSOR]], [[ACCESSOR]]* [[BEGIN]], i64 1 +// CHECK: [[ELEM2_ASCAST:%[a-zA-Z0-9_.]+]] = addrspacecast [[ACCESSOR]]* [[ELEM2_GEP]] to [[ACCESSOR]] addrspace(4)* +// CTOR Call #2 +// CHECK: call spir_func void @{{.+}}([[ACCESSOR]] addrspace(4)* [[ELEM2_ASCAST]]) +// CHECK acc[0] __init method call +// CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0 +// CHECK: [[INDEX1:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], [2 x [[ACCESSOR]]]* [[ACCESSOR_ARRAY1]], i64 0, i64 0 // CHECK load from kernel pointer argument alloca // CHECK: [[MEM_LOAD1:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG1]] - -// CHECK: [[ACC_CAST1:%[0-9]+]] = addrspacecast [[ACCESSOR]]* {{.*}} to [[ACCESSOR]] addrspace(4)* - -// CHECK acc[0] __init method call +// CHECK: [[ACC_CAST1:%[0-9]+]] = addrspacecast [[ACCESSOR]]* [[INDEX1]] to [[ACCESSOR]] addrspace(4)* // CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* [[ACC_CAST1]], i32 addrspace(1)* [[MEM_LOAD1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE1]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET1]]) -// CHECK load from kernel pointer argument alloca -// CHECK: [[MEM_LOAD2:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG2]] - -// CHECK: [[ACC_CAST2:%[0-9]+]] = addrspacecast [[ACCESSOR]]* {{.*}} to [[ACCESSOR]] addrspace(4)* // CHECK acc[1] __init method call +// CHECK: [[ACCESSOR_ARRAY2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0 +// CHECK: [[INDEX2:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], [2 x [[ACCESSOR]]]* [[ACCESSOR_ARRAY2]], i64 0, i64 1 +// CHECK load from kernel pointer argument alloca +// CHECK: [[MEM_LOAD2:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG2]] +// CHECK: [[ACC_CAST2:%[0-9]+]] = addrspacecast [[ACCESSOR]]* [[INDEX2]] to [[ACCESSOR]] addrspace(4)* // CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* [[ACC_CAST2]], i32 addrspace(1)* [[MEM_LOAD2]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE2]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE2]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET2]]) From 5ce171f4c6c4115a836fbf08f67cedc1c9c6217d Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Wed, 26 Aug 2020 11:58:02 -0700 Subject: [PATCH 07/23] Update test, now no longer loops through --- .../kernel-param-member-acc-array.cpp | 19 ++++++++++++------- 1 file changed, 12 insertions(+), 7 deletions(-) diff --git a/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp b/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp index b77eecfc85d68..87304dc9b5481 100644 --- a/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp +++ b/clang/test/CodeGenSYCL/kernel-param-member-acc-array.cpp @@ -52,13 +52,18 @@ int main() { // CHECK: [[MEM_RANGE2:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range" // CHECK: [[OFFSET2:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::id" -// Check loop which calls the default constructor for each element of accessor array is emitted. -// CHECK: [[GEP_LAMBDA:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0 -// CHECK: [[GEP_MEMBER_ACC:%[a-zA-Z_]+]] = getelementptr inbounds %struct.{{.*}}.struct_acc_t, %struct.{{.*}}.struct_acc_t* [[GEP_LAMBDA]], i32 0, i32 0 -// CHECK: [[ARRAY_BEGIN:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR:.*]]], [2 x [[ACCESSOR]]]* [[GEP_MEMBER_ACC]], i64 0, i64 0 -// CHECK: [[ARRAY_END:%[a-zA-Z0-9._]*]] = getelementptr inbounds [[ACCESSOR]], [[ACCESSOR]]* [[ARRAY_BEGIN]], i64 2 -// CHECK: br label %arrayctor.loop -// CHECK: arrayctor.loop: +// CHECK accessor array default inits +// CHECK: [[ACCESSOR_WRAPPER:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0 +// CHECK: [[ACCESSOR_ARRAY1:%[a-zA-Z0-9_.]+]] = getelementptr inbounds %struct.{{.*}}.struct_acc_t, %struct.{{.*}}.struct_acc_t* [[ACCESSOR_WRAPPER]], i32 0, i32 0 +// CHECK: [[BEGIN:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR:.*]]], [2 x [[ACCESSOR]]]* [[ACCESSOR_ARRAY1]], i64 0, i64 0 +// Clang takes advantage of element 1 having the same address as the array, so it doesn't do a GEP. +// CHECK: [[ELEM1_ASCAST:%[a-zA-Z0-9_.]+]] = addrspacecast [[ACCESSOR]]* [[BEGIN]] to [[ACCESSOR]] addrspace(4)* +// CTOR Call #1 +// CHECK: call spir_func void @{{.+}}([[ACCESSOR]] addrspace(4)* [[ELEM1_ASCAST]]) +// CHECK: [[ELEM2_GEP:%[a-zA-Z0-9_.]+]] = getelementptr inbounds [[ACCESSOR]], [[ACCESSOR]]* [[BEGIN]], i64 1 +// CHECK: [[ELEM2_ASCAST:%[a-zA-Z0-9_.]+]] = addrspacecast [[ACCESSOR]]* [[ELEM2_GEP]] to [[ACCESSOR]] addrspace(4)* +// CTOR Call #2 +// CHECK: call spir_func void @{{.+}}([[ACCESSOR]] addrspace(4)* [[ELEM2_ASCAST]]) // Check acc[0] __init method call // CHECK: [[GEP_LAMBDA1:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0 From 7e745eceb816d0e3e29cb6efeab8acec647611ff Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Wed, 26 Aug 2020 12:16:22 -0700 Subject: [PATCH 08/23] Clang-format --- clang/lib/Sema/SemaSYCL.cpp | 49 ++++++++++--------- .../CodeGenSYCL/kernel-param-acc-array.cpp | 1 - 2 files changed, 26 insertions(+), 24 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 95c578dc2fd6e..c599b8e2f4c9b 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -901,7 +901,7 @@ class KernelObjVisitor { Handlers &... handlers) { (void)std::initializer_list{ (handlers.nextElement(ElementTy, Index), 0)...}; - VisitField(Owner, ArrayField, ElementTy, handlers...); + VisitField(Owner, ArrayField, ElementTy, handlers...); } template @@ -1128,8 +1128,7 @@ void KernelObjVisitor::VisitUnion(const CXXRecordDecl *Owner, ParentTy &Parent, template void KernelObjVisitor::VisitNthArrayElement(const CXXRecordDecl *Owner, FieldDecl *ArrayField, - QualType ElementTy, - uint64_t Index, + QualType ElementTy, uint64_t Index, Handlers &... handlers) { // Don't continue descending if none of the handlers 'care'. This could be 'if // constexpr' starting in C++17. Until then, we have to count on the @@ -1718,7 +1717,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { } // Creates a DeclRefExpr to the ParmVar that represents the current field. - Expr* createParamReferenceExpr() { + Expr *createParamReferenceExpr() { ParmVarDecl *KernelParameter = DeclCreator.getParamVarDeclsForCurrentField()[0]; @@ -1730,7 +1729,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { // Creates a DeclRefExpr to the ParmVar that represents the current pointer // field. - Expr* createPointerParamReferenceExpr(QualType PointerTy, bool Wrapped) { + Expr *createPointerParamReferenceExpr(QualType PointerTy, bool Wrapped) { ParmVarDecl *KernelParameter = DeclCreator.getParamVarDeclsForCurrentField()[0]; @@ -1742,11 +1741,11 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { // then wrapped inside a compiler generated struct. Therefore when // generating the initializers, we have to 'unwrap' the pointer. if (Wrapped) { - CXXRecordDecl *WrapperStruct = ParamType->getAsCXXRecordDecl(); - // Pointer field wrapped inside __wrapper_class - FieldDecl *Pointer = *(WrapperStruct->field_begin()); - DRE = BuildMemberExpr(DRE, Pointer); - ParamType = Pointer->getType(); + CXXRecordDecl *WrapperStruct = ParamType->getAsCXXRecordDecl(); + // Pointer field wrapped inside __wrapper_class + FieldDecl *Pointer = *(WrapperStruct->field_begin()); + DRE = BuildMemberExpr(DRE, Pointer); + ParamType = Pointer->getType(); } if (PointerTy->getPointeeType().getAddressSpace() != @@ -1792,8 +1791,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { InitializedEntity Entity = getFieldEntity(FD, Ty); InitializationSequence InitSeq(SemaRef, Entity, InitKind, ParamRef); - ExprResult Init = - InitSeq.Perform(SemaRef, Entity, InitKind, ParamRef); + ExprResult Init = InitSeq.Perform(SemaRef, Entity, InitKind, ParamRef); InitListExpr *ParentILE = CollectionInitExprs.back(); ParentILE->updateInit(SemaRef.getASTContext(), ParentILE->getNumInits(), @@ -1879,7 +1877,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { InitListExpr *CreateInitListExpr(QualType InitTy, uint64_t NumChildInits) { InitListExpr *ILE = new (SemaRef.getASTContext()) InitListExpr( - SemaRef.getASTContext(), SourceLocation(), {}, SourceLocation()); + SemaRef.getASTContext(), SourceLocation(), {}, SourceLocation()); ILE->reserveInits(SemaRef.getASTContext(), NumChildInits); ILE->setType(InitTy); @@ -1959,8 +1957,8 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { } ~SyclKernelBodyCreator() { - CompoundStmt *KernelBody = createKernelBody(); - DeclCreator.setBody(KernelBody); + CompoundStmt *KernelBody = createKernelBody(); + DeclCreator.setBody(KernelBody); } bool handleSyclAccessorType(FieldDecl *FD, QualType Ty) final { @@ -2050,13 +2048,14 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { return true; } - bool enterStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS, QualType) final { + bool enterStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS, + QualType) final { ++ContainerDepth; CXXCastPath BasePath; QualType DerivedTy(RD->getTypeForDecl(), 0); QualType BaseTy = BS.getType(); -// // TODO: Why is this here? Do we think this check could fail? + // // TODO: Why is this here? Do we think this check could fail? SemaRef.CheckDerivedToBaseConversion(DerivedTy, BaseTy, SourceLocation(), SourceRange(), &BasePath, /*IgnoreBaseAccess*/ true); @@ -2069,7 +2068,8 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { return true; } - bool leaveStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS, QualType) final { + bool leaveStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS, + QualType) final { --ContainerDepth; MemberExprBases.pop_back(); CollectionInitExprs.pop_back(); @@ -2096,7 +2096,8 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { ArrayInfos.back().second = Index; // Pop off the last member expr base. - if (Index != 0) MemberExprBases.pop_back(); + if (Index != 0) + MemberExprBases.pop_back(); QualType SizeT = SemaRef.getASTContext().getSizeType(); @@ -2279,17 +2280,19 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { return true; } - bool enterStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS, QualType) final { + bool enterStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS, + QualType) final { CurOffset += offsetOf(RD, BS.getType()->getAsCXXRecordDecl()); return true; } - bool leaveStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS, QualType) final { + bool leaveStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS, + QualType) final { CurOffset -= offsetOf(RD, BS.getType()->getAsCXXRecordDecl()); return true; } - bool enterArray(FieldDecl*, QualType, QualType) final { + bool enterArray(FieldDecl *, QualType, QualType) final { ArrayBaseOffsets.push_back(CurOffset); return true; } @@ -2300,7 +2303,7 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { return true; } - bool leaveArray(FieldDecl*, QualType, QualType) final { + bool leaveArray(FieldDecl *, QualType, QualType) final { CurOffset = ArrayBaseOffsets.pop_back_val(); return true; } diff --git a/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp b/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp index 9032dc7fe3f02..618e850f514b5 100644 --- a/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp +++ b/clang/test/CodeGenSYCL/kernel-param-acc-array.cpp @@ -69,7 +69,6 @@ int main() { // CHECK: [[ACC_CAST1:%[0-9]+]] = addrspacecast [[ACCESSOR]]* [[INDEX1]] to [[ACCESSOR]] addrspace(4)* // CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor" addrspace(4)* [[ACC_CAST1]], i32 addrspace(1)* [[MEM_LOAD1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[ACC_RANGE1]], %"struct.{{.*}}.cl::sycl::range"* byval({{.*}}) align 4 [[MEM_RANGE1]], %"struct.{{.*}}.cl::sycl::id"* byval({{.*}}) align 4 [[OFFSET1]]) - // CHECK acc[1] __init method call // CHECK: [[ACCESSOR_ARRAY2:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0 // CHECK: [[INDEX2:%[a-zA-Z0-9._]*]] = getelementptr inbounds [2 x [[ACCESSOR]]], [2 x [[ACCESSOR]]]* [[ACCESSOR_ARRAY2]], i64 0, i64 1 From 03ad0d2990ca9bdcd570f33187e2c959f9a0d1a2 Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Thu, 27 Aug 2020 06:31:38 -0700 Subject: [PATCH 09/23] Add enter/leaveStream, since the handling for them is so different from the others. --- clang/lib/Sema/SemaSYCL.cpp | 92 ++++++++++++++++++++++++++----------- 1 file changed, 64 insertions(+), 28 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index c599b8e2f4c9b..d842279867c1e 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -884,7 +884,7 @@ class KernelObjVisitor { CXXRecordDecl *Wrapper, QualType RecordTy, Handlers &... handlers) { (void)std::initializer_list{ - (handlers.enterStruct(Owner, Parent, RecordTy), 0)...}; + (handlers.enterStream(Owner, Parent, RecordTy), 0)...}; for (const auto &Field : Wrapper->fields()) { QualType FieldTy = Field->getType(); // Required to initialize accessors inside streams. @@ -892,7 +892,7 @@ class KernelObjVisitor { KF_FOR_EACH(handleSyclAccessorType, Field, FieldTy); } (void)std::initializer_list{ - (handlers.leaveStruct(Owner, Parent, RecordTy), 0)...}; + (handlers.leaveStream(Owner, Parent, RecordTy), 0)...}; } template @@ -1056,6 +1056,12 @@ 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; @@ -1064,7 +1070,6 @@ class SyclKernelFieldHandlerBase { QualType) { return true; } - // TODO: Does enter-union need to be worried when it is in an array?! virtual bool enterUnion(const CXXRecordDecl *, FieldDecl *) { return true; } virtual bool leaveUnion(const CXXRecordDecl *, FieldDecl *) { return true; } @@ -1517,6 +1522,14 @@ 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; @@ -1778,9 +1791,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { return InitializedEntity::InitializeMember(FD, &VarEntity); } - // TODO: This name seems outdated now :) void addFieldInit(FieldDecl *FD, QualType Ty, MultiExprArg ParamRef) { - // TODO: Why is this by copy rather than 'forinit' or value init? InitializationKind InitKind = InitializationKind::CreateCopy(SourceLocation(), SourceLocation()); addFieldInit(FD, Ty, ParamRef, InitKind); @@ -1929,7 +1940,6 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { bool handleSpecialType(const CXXBaseSpecifier &BS, QualType Ty) { const auto *RecordDecl = Ty->getAsCXXRecordDecl(); addBaseInit(BS, Ty, InitializationKind::CreateDefault(SourceLocation())); - // TODO: Needs MemberExprBases entry? createSpecialMethodCall(RecordDecl, InitMethodName, BodyStmts); return true; } @@ -1979,22 +1989,8 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { } bool handleSyclStreamType(FieldDecl *FD, QualType Ty) final { - const auto *StreamDecl = Ty->getAsCXXRecordDecl(); // Streams just get copied as a new init. addSimpleFieldInit(FD, Ty); - - // Add a dummy init expression to catch the accessor initializers. - CollectionInitExprs.push_back(CreateInitListExpr(StreamDecl)); - - // Add init/finalize method calls. - if (!IsArrayElement(FD, Ty)) - MemberExprBases.push_back(BuildMemberExpr(MemberExprBases.back(), FD)); - - createSpecialMethodCall(StreamDecl, InitMethodName, BodyStmts); - createSpecialMethodCall(StreamDecl, FinalizeMethodName, FinalizeStmts); - - if (!IsArrayElement(FD, Ty)) - MemberExprBases.pop_back(); return true; } @@ -2027,23 +2023,47 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { return true; } + bool enterStream(const CXXRecordDecl *RD, FieldDecl *FD, QualType Ty) final { + ++ContainerDepth; + // Add a dummy init expression to catch the accessor initializers. + const auto *StreamDecl = Ty->getAsCXXRecordDecl(); + CollectionInitExprs.push_back(CreateInitListExpr(StreamDecl)); + + MemberExprBases.push_back(BuildMemberExpr(MemberExprBases.back(), FD)); + return true; + } + + bool leaveStream(const CXXRecordDecl *RD, FieldDecl *FD, QualType Ty) final { + --ContainerDepth; + // Stream requires that its 'init' calls happen after its accessors init + // calls, so add them here instead. + const auto *StreamDecl = Ty->getAsCXXRecordDecl(); + if (!IsArrayElement(FD, Ty)) + MemberExprBases.push_back(BuildMemberExpr(MemberExprBases.back(), FD)); + + createSpecialMethodCall(StreamDecl, InitMethodName, BodyStmts); + createSpecialMethodCall(StreamDecl, FinalizeMethodName, FinalizeStmts); + + if (!IsArrayElement(FD, Ty)) + MemberExprBases.pop_back(); + + CollectionInitExprs.pop_back(); + MemberExprBases.pop_back(); + return true; + } + bool enterStruct(const CXXRecordDecl *RD, FieldDecl *FD, QualType Ty) final { ++ContainerDepth; - // We handle adding a throw-away initializer in handleSyclStreamType since - // the 'default' init needs to stick around, but the accessors that are - // 'children' of it do not. - if (!Util::isSyclStreamType(Ty)) - addCollectionInitListExpr(Ty->getAsCXXRecordDecl()); + addCollectionInitListExpr(Ty->getAsCXXRecordDecl()); MemberExprBases.push_back(BuildMemberExpr(MemberExprBases.back(), FD)); return true; } - bool leaveStruct(const CXXRecordDecl *, FieldDecl *FD, QualType) final { + bool leaveStruct(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final { --ContainerDepth; - // If this is a stream, this has popped the 'fake' one that was added in - // handleSyclStreamType, which hasn't been added as a child. CollectionInitExprs.pop_back(); + MemberExprBases.pop_back(); return true; } @@ -2264,6 +2284,22 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { return true; } + bool enterStream(const CXXRecordDecl *, FieldDecl *FD, QualType) final { + ++StructDepth; + // TODO: Is this right?! I think this only needs to be incremented when we + // aren't in an array, otherwise 'enterArray's base offsets should handle + // this right. Otherwise an array of structs is going to be in the middle + // of nowhere. + CurOffset += offsetOf(FD); + return true; + } + + bool leaveStream(const CXXRecordDecl *, FieldDecl *FD, QualType) final { + --StructDepth; + CurOffset -= offsetOf(FD); + return true; + } + bool enterStruct(const CXXRecordDecl *, FieldDecl *FD, QualType) final { ++StructDepth; // TODO: Is this right?! I think this only needs to be incremented when we From 90ae330c86c2a5b75a602c31469d0cb323f64751 Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Thu, 27 Aug 2020 08:11:52 -0700 Subject: [PATCH 10/23] Correct behavior of integration header, now that enterField isn't here. It sorta worked before this, just by chance, but it would have broken in a couple of other cases I can think of. This makes sure that the integration-header code is array-aware. --- clang/lib/Sema/SemaSYCL.cpp | 69 ++++++++++++++++++++++++------------- 1 file changed, 45 insertions(+), 24 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index d842279867c1e..44aa588aa2fe0 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1774,8 +1774,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { // is an element of an array. This will determine whether we do // MemberExprBases in some cases or not, AND determines how we initialize // values. - bool IsArrayElement(FieldDecl *FD, QualType Ty) { - // TODO, better way to detect that we're in an array? + bool IsArrayElement(const FieldDecl *FD, QualType Ty) const { SemaRef.getASTContext().hasSameType(FD->getType(), Ty); return FD->getType() != Ty; } @@ -2075,7 +2074,6 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { CXXCastPath BasePath; QualType DerivedTy(RD->getTypeForDecl(), 0); QualType BaseTy = BS.getType(); - // // TODO: Why is this here? Do we think this check could fail? SemaRef.CheckDerivedToBaseConversion(DerivedTy, BaseTy, SourceLocation(), SourceRange(), &BasePath, /*IgnoreBaseAccess*/ true); @@ -2178,12 +2176,30 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { void addParam(const FieldDecl *FD, QualType ArgTy, SYCLIntegrationHeader::kernel_param_kind_t Kind) { + addParam(FD, ArgTy, Kind, IsArrayElement(FD, ArgTy)); + } + void addParam(const FieldDecl *FD, QualType ArgTy, + SYCLIntegrationHeader::kernel_param_kind_t Kind, + bool IsArrayElem) { uint64_t Size; Size = SemaRef.getASTContext().getTypeSizeInChars(ArgTy).getQuantity(); + uint64_t Offset = CurOffset; + if (!IsArrayElem) + Offset += offsetOf(FD); Header.addParamDesc(Kind, static_cast(Size), - static_cast(CurOffset + offsetOf(FD))); + static_cast(Offset)); + } + + // Returns 'true' if the thing we're visiting (Based on the FD/QualType pair) + // is an element of an array. This will determine whether we do + // MemberExprBases in some cases or not, AND determines how we initialize + // values. + bool IsArrayElement(const FieldDecl *FD, QualType Ty) const { + SemaRef.getASTContext().hasSameType(FD->getType(), Ty); + return FD->getType() != Ty; } + public: SyclKernelIntHeaderCreator(Sema &S, SYCLIntegrationHeader &H, const CXXRecordDecl *KernelObj, QualType NameType, @@ -2216,8 +2232,12 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { int Dims = static_cast( AccTy->getTemplateArgs()[1].getAsIntegral().getExtValue()); int Info = getAccessTarget(AccTy) | (Dims << 11); + + uint64_t Offset = CurOffset; + if (!IsArrayElement(FD, FieldTy)) + Offset += offsetOf(FD); Header.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info, - CurOffset + offsetOf(FD)); + Offset); return true; } @@ -2231,7 +2251,8 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { const ParmVarDecl *SamplerArg = InitMethod->getParamDecl(0); assert(SamplerArg && "sampler __init method must have sampler parameter"); - addParam(FD, SamplerArg->getType(), SYCLIntegrationHeader::kind_sampler); + addParam(FD, SamplerArg->getType(), SYCLIntegrationHeader::kind_sampler, + IsArrayElement(FD, FieldTy)); return true; } @@ -2284,35 +2305,31 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { return true; } - bool enterStream(const CXXRecordDecl *, FieldDecl *FD, QualType) final { + bool enterStream(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final { ++StructDepth; - // TODO: Is this right?! I think this only needs to be incremented when we - // aren't in an array, otherwise 'enterArray's base offsets should handle - // this right. Otherwise an array of structs is going to be in the middle - // of nowhere. - CurOffset += offsetOf(FD); + if (!IsArrayElement(FD, Ty)) + CurOffset += offsetOf(FD); return true; } - bool leaveStream(const CXXRecordDecl *, FieldDecl *FD, QualType) final { + bool leaveStream(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final { --StructDepth; - CurOffset -= offsetOf(FD); + if (!IsArrayElement(FD, Ty)) + CurOffset -= offsetOf(FD); return true; } - bool enterStruct(const CXXRecordDecl *, FieldDecl *FD, QualType) final { + bool enterStruct(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final { ++StructDepth; - // TODO: Is this right?! I think this only needs to be incremented when we - // aren't in an array, otherwise 'enterArray's base offsets should handle - // this right. Otherwise an array of structs is going to be in the middle - // of nowhere. - CurOffset += offsetOf(FD); + if (!IsArrayElement(FD, Ty)) + CurOffset += offsetOf(FD); return true; } - bool leaveStruct(const CXXRecordDecl *, FieldDecl *FD, QualType) final { + bool leaveStruct(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final { --StructDepth; - CurOffset -= offsetOf(FD); + if (!IsArrayElement(FD, Ty)) + CurOffset -= offsetOf(FD); return true; } @@ -2328,8 +2345,12 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { return true; } - bool enterArray(FieldDecl *, QualType, QualType) final { - ArrayBaseOffsets.push_back(CurOffset); + bool enterArray(FieldDecl *FD, QualType ArrayTy, QualType) final { + uint64_t Offset = CurOffset; + if (!IsArrayElement(FD, ArrayTy)) + Offset += offsetOf(FD); + + ArrayBaseOffsets.push_back(Offset); return true; } From 1a96e206b99b618ca8f920892a1497a47b10d410 Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Thu, 27 Aug 2020 08:46:49 -0700 Subject: [PATCH 11/23] Clang-format fixes --- clang/lib/Sema/SemaSYCL.cpp | 4 +--- 1 file changed, 1 insertion(+), 3 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 44aa588aa2fe0..5a90acf53c18f 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -2199,7 +2199,6 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { return FD->getType() != Ty; } - public: SyclKernelIntHeaderCreator(Sema &S, SYCLIntegrationHeader &H, const CXXRecordDecl *KernelObj, QualType NameType, @@ -2236,8 +2235,7 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { uint64_t Offset = CurOffset; if (!IsArrayElement(FD, FieldTy)) Offset += offsetOf(FD); - Header.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info, - Offset); + Header.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info, Offset); return true; } From 693ae2b25bff5bf757c9996504372d8dd8b880bb Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Thu, 27 Aug 2020 09:45:33 -0700 Subject: [PATCH 12/23] Revert "Correct behavior of integration header, now that enterField isn't here." Temporarily rever the int-header changes to see if this fixes the validation machine issues I see. This reverts commit 90ae330c86c2a5b75a602c31469d0cb323f64751. Conflicts: clang/lib/Sema/SemaSYCL.cpp --- clang/lib/Sema/SemaSYCL.cpp | 64 +++++++++++++------------------------ 1 file changed, 22 insertions(+), 42 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 5a90acf53c18f..155a9f2bc7ed4 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -2176,27 +2176,10 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { void addParam(const FieldDecl *FD, QualType ArgTy, SYCLIntegrationHeader::kernel_param_kind_t Kind) { - addParam(FD, ArgTy, Kind, IsArrayElement(FD, ArgTy)); - } - void addParam(const FieldDecl *FD, QualType ArgTy, - SYCLIntegrationHeader::kernel_param_kind_t Kind, - bool IsArrayElem) { uint64_t Size; Size = SemaRef.getASTContext().getTypeSizeInChars(ArgTy).getQuantity(); - uint64_t Offset = CurOffset; - if (!IsArrayElem) - Offset += offsetOf(FD); Header.addParamDesc(Kind, static_cast(Size), - static_cast(Offset)); - } - - // Returns 'true' if the thing we're visiting (Based on the FD/QualType pair) - // is an element of an array. This will determine whether we do - // MemberExprBases in some cases or not, AND determines how we initialize - // values. - bool IsArrayElement(const FieldDecl *FD, QualType Ty) const { - SemaRef.getASTContext().hasSameType(FD->getType(), Ty); - return FD->getType() != Ty; + static_cast(CurOffset + offsetOf(FD))); } public: @@ -2232,10 +2215,8 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { AccTy->getTemplateArgs()[1].getAsIntegral().getExtValue()); int Info = getAccessTarget(AccTy) | (Dims << 11); - uint64_t Offset = CurOffset; - if (!IsArrayElement(FD, FieldTy)) - Offset += offsetOf(FD); - Header.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info, Offset); + Header.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info, + CurOffset + offsetOf(FD)); return true; } @@ -2249,8 +2230,7 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { const ParmVarDecl *SamplerArg = InitMethod->getParamDecl(0); assert(SamplerArg && "sampler __init method must have sampler parameter"); - addParam(FD, SamplerArg->getType(), SYCLIntegrationHeader::kind_sampler, - IsArrayElement(FD, FieldTy)); + addParam(FD, SamplerArg->getType(), SYCLIntegrationHeader::kind_sampler); return true; } @@ -2303,31 +2283,35 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { return true; } - bool enterStream(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final { + bool enterStream(const CXXRecordDecl *, FieldDecl *FD, QualType) final { ++StructDepth; - if (!IsArrayElement(FD, Ty)) - CurOffset += offsetOf(FD); + // TODO: Is this right?! I think this only needs to be incremented when we + // aren't in an array, otherwise 'enterArray's base offsets should handle + // this right. Otherwise an array of structs is going to be in the middle + // of nowhere. + CurOffset += offsetOf(FD); return true; } - bool leaveStream(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final { + bool leaveStream(const CXXRecordDecl *, FieldDecl *FD, QualType) final { --StructDepth; - if (!IsArrayElement(FD, Ty)) - CurOffset -= offsetOf(FD); + CurOffset -= offsetOf(FD); return true; } - bool enterStruct(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final { + bool enterStruct(const CXXRecordDecl *, FieldDecl *FD, QualType) final { ++StructDepth; - if (!IsArrayElement(FD, Ty)) - CurOffset += offsetOf(FD); + // TODO: Is this right?! I think this only needs to be incremented when we + // aren't in an array, otherwise 'enterArray's base offsets should handle + // this right. Otherwise an array of structs is going to be in the middle + // of nowhere. + CurOffset += offsetOf(FD); return true; } - bool leaveStruct(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final { + bool leaveStruct(const CXXRecordDecl *, FieldDecl *FD, QualType) final { --StructDepth; - if (!IsArrayElement(FD, Ty)) - CurOffset -= offsetOf(FD); + CurOffset -= offsetOf(FD); return true; } @@ -2343,12 +2327,8 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { return true; } - bool enterArray(FieldDecl *FD, QualType ArrayTy, QualType) final { - uint64_t Offset = CurOffset; - if (!IsArrayElement(FD, ArrayTy)) - Offset += offsetOf(FD); - - ArrayBaseOffsets.push_back(Offset); + bool enterArray(FieldDecl *, QualType, QualType) final { + ArrayBaseOffsets.push_back(CurOffset); return true; } From a6ee42509b41bd7e149eba8758d0f38a3097699e Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Thu, 27 Aug 2020 10:06:14 -0700 Subject: [PATCH 13/23] Reapply Correct behavior of integration header, now that enterField isn't here. This reverts commit 693ae2b25bff5bf757c9996504372d8dd8b880bb. Fixed now :) --- clang/lib/Sema/SemaSYCL.cpp | 61 +++++++++++++++++++++---------------- 1 file changed, 35 insertions(+), 26 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 155a9f2bc7ed4..7a0a5d795f81a 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1775,8 +1775,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { // MemberExprBases in some cases or not, AND determines how we initialize // values. bool IsArrayElement(const FieldDecl *FD, QualType Ty) const { - SemaRef.getASTContext().hasSameType(FD->getType(), Ty); - return FD->getType() != Ty; + return !SemaRef.getASTContext().hasSameType(FD->getType(), Ty); } // Creates an initialized entity for a field/item. In the case where this is a @@ -2164,8 +2163,10 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { int StructDepth = 0; // A series of functions to calculate the change in offset based on the type. - int64_t offsetOf(const FieldDecl *FD) const { - return SemaRef.getASTContext().getFieldOffset(FD) / 8; + int64_t offsetOf(const FieldDecl *FD, QualType ArgTy) const { + return IsArrayElement(FD, ArgTy) + ? 0 + : SemaRef.getASTContext().getFieldOffset(FD) / 8; } int64_t offsetOf(const CXXRecordDecl *RD, const CXXRecordDecl *Base) const { @@ -2176,10 +2177,23 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { void addParam(const FieldDecl *FD, QualType ArgTy, SYCLIntegrationHeader::kernel_param_kind_t Kind) { + addParam(FD, ArgTy, Kind, offsetOf(FD, ArgTy)); + } + void addParam(const FieldDecl *FD, QualType ArgTy, + SYCLIntegrationHeader::kernel_param_kind_t Kind, + uint64_t OffsetAdj) { uint64_t Size; Size = SemaRef.getASTContext().getTypeSizeInChars(ArgTy).getQuantity(); Header.addParamDesc(Kind, static_cast(Size), - static_cast(CurOffset + offsetOf(FD))); + static_cast(CurOffset + OffsetAdj)); + } + + // Returns 'true' if the thing we're visiting (Based on the FD/QualType pair) + // is an element of an array. This will determine whether we do + // MemberExprBases in some cases or not, AND determines how we initialize + // values. + bool IsArrayElement(const FieldDecl *FD, QualType Ty) const { + return !SemaRef.getASTContext().hasSameType(FD->getType(), Ty); } public: @@ -2216,7 +2230,7 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { int Info = getAccessTarget(AccTy) | (Dims << 11); Header.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info, - CurOffset + offsetOf(FD)); + CurOffset + offsetOf(FD, FieldTy)); return true; } @@ -2230,7 +2244,8 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { const ParmVarDecl *SamplerArg = InitMethod->getParamDecl(0); assert(SamplerArg && "sampler __init method must have sampler parameter"); - addParam(FD, SamplerArg->getType(), SYCLIntegrationHeader::kind_sampler); + addParam(FD, SamplerArg->getType(), SYCLIntegrationHeader::kind_sampler, + offsetOf(FD, FieldTy)); return true; } @@ -2283,35 +2298,27 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { return true; } - bool enterStream(const CXXRecordDecl *, FieldDecl *FD, QualType) final { + bool enterStream(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final { ++StructDepth; - // TODO: Is this right?! I think this only needs to be incremented when we - // aren't in an array, otherwise 'enterArray's base offsets should handle - // this right. Otherwise an array of structs is going to be in the middle - // of nowhere. - CurOffset += offsetOf(FD); + CurOffset += offsetOf(FD, Ty); return true; } - bool leaveStream(const CXXRecordDecl *, FieldDecl *FD, QualType) final { + bool leaveStream(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final { --StructDepth; - CurOffset -= offsetOf(FD); + CurOffset -= offsetOf(FD, Ty); return true; } - bool enterStruct(const CXXRecordDecl *, FieldDecl *FD, QualType) final { + bool enterStruct(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final { ++StructDepth; - // TODO: Is this right?! I think this only needs to be incremented when we - // aren't in an array, otherwise 'enterArray's base offsets should handle - // this right. Otherwise an array of structs is going to be in the middle - // of nowhere. - CurOffset += offsetOf(FD); + CurOffset += offsetOf(FD, Ty); return true; } - bool leaveStruct(const CXXRecordDecl *, FieldDecl *FD, QualType) final { + bool leaveStruct(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final { --StructDepth; - CurOffset -= offsetOf(FD); + CurOffset -= offsetOf(FD, Ty); return true; } @@ -2327,8 +2334,8 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { return true; } - bool enterArray(FieldDecl *, QualType, QualType) final { - ArrayBaseOffsets.push_back(CurOffset); + bool enterArray(FieldDecl *FD, QualType ArrayTy, QualType) final { + ArrayBaseOffsets.push_back(CurOffset + offsetOf(FD, ArrayTy)); return true; } @@ -2338,10 +2345,12 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { return true; } - bool leaveArray(FieldDecl *, QualType, QualType) final { + bool leaveArray(FieldDecl *FD, QualType ArrayTy, QualType) final { CurOffset = ArrayBaseOffsets.pop_back_val(); + CurOffset -= offsetOf(FD, ArrayTy); return true; } + using SyclKernelFieldHandler::enterStruct; using SyclKernelFieldHandler::handleSyclHalfType; using SyclKernelFieldHandler::handleSyclSamplerType; From 7ed04d7065350d678584f68baaeeb9d87cafc467 Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Fri, 28 Aug 2020 09:00:42 -0700 Subject: [PATCH 14/23] Fix a pair of issues in stream/arrays of struct, add streams test that tests most array situations too. --- clang/lib/Sema/SemaSYCL.cpp | 12 +- clang/test/SemaSYCL/Inputs/sycl.hpp | 12 + clang/test/SemaSYCL/streams.cpp | 932 ++++++++++++++++++++++++++++ 3 files changed, 950 insertions(+), 6 deletions(-) create mode 100644 clang/test/SemaSYCL/streams.cpp diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 7a0a5d795f81a..94b38fc2bc8e2 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -2027,7 +2027,8 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { const auto *StreamDecl = Ty->getAsCXXRecordDecl(); CollectionInitExprs.push_back(CreateInitListExpr(StreamDecl)); - MemberExprBases.push_back(BuildMemberExpr(MemberExprBases.back(), FD)); + if (!IsArrayElement(FD, Ty)) + MemberExprBases.push_back(BuildMemberExpr(MemberExprBases.back(), FD)); return true; } @@ -2036,8 +2037,6 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { // Stream requires that its 'init' calls happen after its accessors init // calls, so add them here instead. const auto *StreamDecl = Ty->getAsCXXRecordDecl(); - if (!IsArrayElement(FD, Ty)) - MemberExprBases.push_back(BuildMemberExpr(MemberExprBases.back(), FD)); createSpecialMethodCall(StreamDecl, InitMethodName, BodyStmts); createSpecialMethodCall(StreamDecl, FinalizeMethodName, FinalizeStmts); @@ -2046,7 +2045,6 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { MemberExprBases.pop_back(); CollectionInitExprs.pop_back(); - MemberExprBases.pop_back(); return true; } @@ -2054,7 +2052,8 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { ++ContainerDepth; addCollectionInitListExpr(Ty->getAsCXXRecordDecl()); - MemberExprBases.push_back(BuildMemberExpr(MemberExprBases.back(), FD)); + if (!IsArrayElement(FD, Ty)) + MemberExprBases.push_back(BuildMemberExpr(MemberExprBases.back(), FD)); return true; } @@ -2062,7 +2061,8 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { --ContainerDepth; CollectionInitExprs.pop_back(); - MemberExprBases.pop_back(); + if (!IsArrayElement(FD, Ty)) + MemberExprBases.pop_back(); return true; } diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp index 2e01b5235a63a..f8178ad26e2da 100644 --- a/clang/test/SemaSYCL/Inputs/sycl.hpp +++ b/clang/test/SemaSYCL/Inputs/sycl.hpp @@ -218,6 +218,18 @@ class handler { } }; +class stream { +public: + stream(unsigned long BufferSize, unsigned long MaxStatementSize, + handler &CGH) {} + + void __init() {} + void use() const {} + + void __finalize() {} +}; + + namespace ONEAPI { namespace experimental { template diff --git a/clang/test/SemaSYCL/streams.cpp b/clang/test/SemaSYCL/streams.cpp new file mode 100644 index 0000000000000..dd3d93818398f --- /dev/null +++ b/clang/test/SemaSYCL/streams.cpp @@ -0,0 +1,932 @@ +// RUN: %clang_cc1 -S -I %S/Inputs -fsycl -fsycl-is-device -triple spir64 -ast-dump %s | FileCheck %s + +#include + +template +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { + kernelFunc(); +} + +using namespace cl::sycl; + +handler H; + +struct HasStreams { + stream s1{0,0,H}; + stream s_array[2] = {{0, 0, H}, {0, 0, H}}; +}; + +struct HasArrayOfHasStreams { + int i; + HasStreams hs[2]; +}; + +void use() { + stream in_lambda {0,0,H}; + stream in_lambda_array[2] = {{0, 0, H}, {0, 0, H}}; + stream in_lambda_mdarray[2][2] = {{{0, 0, H}, {0, 0, H}},{{0, 0, H}, {0, 0, H}}}; + + HasStreams Struct; + HasArrayOfHasStreams haohs; + HasArrayOfHasStreams haohs_array[2]; + + kernel([=]() { + in_lambda.use(); + in_lambda_array[1].use(); + in_lambda_mdarray[1][1].use(); + + Struct.s1.use(); + + haohs.hs[0].s1.use(); + haohs_array[0].hs[0].s1.use(); + }); + } + +// Function Declaration +// CHECK: FunctionDecl {{.*}}stream_test{{.*}} + +// Initializers: +// CHECK: InitListExpr {{.*}} '(lambda at +// 'in_lambda' +// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::stream' 'void (const cl::sycl::stream &) noexcept' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::stream' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::stream' lvalue ParmVar + +// 'in_lambda_array' +// CHECK-NEXT: InitListExpr {{.*}} 'cl::sycl::stream [2]' +// element 0 +// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::stream' 'void (const cl::sycl::stream &) noexcept' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::stream' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::stream' lvalue ParmVar +// element 1 +// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::stream' 'void (const cl::sycl::stream &) noexcept' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::stream' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::stream' lvalue ParmVar + +// 'in_lambda_mdarray' +// CHECK-NEXT: InitListExpr {{.*}} 'cl::sycl::stream [2][2]' +// sub-array 0 +// CHECK-NEXT: InitListExpr {{.*}} 'cl::sycl::stream [2]' +// element 0 +// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::stream' 'void (const cl::sycl::stream &) noexcept' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::stream' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::stream' lvalue ParmVar +// element 1 +// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::stream' 'void (const cl::sycl::stream &) noexcept' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::stream' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::stream' lvalue ParmVar +// sub-array 1 +// CHECK-NEXT: InitListExpr {{.*}} 'cl::sycl::stream [2]' +// element 0 +// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::stream' 'void (const cl::sycl::stream &) noexcept' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::stream' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::stream' lvalue ParmVar +// element 1 +// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::stream' 'void (const cl::sycl::stream &) noexcept' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::stream' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::stream' lvalue ParmVar + +// HasStreams struct +// CHECK-NEXT: InitListExpr {{.*}} 'HasStreams' +// HasStreams::s1 +// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::stream' 'void (const cl::sycl::stream &) noexcept' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::stream' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::stream' lvalue ParmVar +// HasStreams::s_array +// CHECK-NEXT: InitListExpr {{.*}} 'cl::sycl::stream [2]' +// element 0 +// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::stream' 'void (const cl::sycl::stream &) noexcept' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::stream' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::stream' lvalue ParmVar +// element 1 +// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::stream' 'void (const cl::sycl::stream &) noexcept' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::stream' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::stream' lvalue ParmVar + +// 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 {{.*}} 'cl::sycl::stream' 'void (const cl::sycl::stream &) noexcept' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::stream' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::stream' lvalue ParmVar +// HasStreams::s_array +// CHECK-NEXT: InitListExpr {{.*}} 'cl::sycl::stream [2]' +// element 0 +// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::stream' 'void (const cl::sycl::stream &) noexcept' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::stream' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::stream' lvalue ParmVar +// element 1 +// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::stream' 'void (const cl::sycl::stream &) noexcept' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::stream' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::stream' lvalue ParmVar +// HasStreams struct +// CHECK-NEXT: InitListExpr {{.*}} 'HasStreams' +// HasStreams::s1 +// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::stream' 'void (const cl::sycl::stream &) noexcept' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::stream' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::stream' lvalue ParmVar +// HasStreams::s_array +// CHECK-NEXT: InitListExpr {{.*}} 'cl::sycl::stream [2]' +// element 0 +// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::stream' 'void (const cl::sycl::stream &) noexcept' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::stream' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::stream' lvalue ParmVar +// element 1 +// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::stream' 'void (const cl::sycl::stream &) noexcept' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::stream' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::stream' lvalue ParmVar + +// HasArrayOfHasStreams Array +// CHECK-NEXT: 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 {{.*}} 'cl::sycl::stream' 'void (const cl::sycl::stream &) noexcept' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::stream' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::stream' lvalue ParmVar +// HasStreams::s_array +// CHECK-NEXT: InitListExpr {{.*}} 'cl::sycl::stream [2]' +// element 0 +// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::stream' 'void (const cl::sycl::stream &) noexcept' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::stream' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::stream' lvalue ParmVar +// element 1 +// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::stream' 'void (const cl::sycl::stream &) noexcept' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::stream' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::stream' lvalue ParmVar +// HasStreams struct +// CHECK-NEXT: InitListExpr {{.*}} 'HasStreams' +// HasStreams::s1 +// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::stream' 'void (const cl::sycl::stream &) noexcept' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::stream' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::stream' lvalue ParmVar +// HasStreams::s_array +// CHECK-NEXT: InitListExpr {{.*}} 'cl::sycl::stream [2]' +// element 0 +// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::stream' 'void (const cl::sycl::stream &) noexcept' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::stream' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::stream' lvalue ParmVar +// element 1 +// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::stream' 'void (const cl::sycl::stream &) noexcept' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::stream' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::stream' lvalue ParmVar +// 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 {{.*}} 'cl::sycl::stream' 'void (const cl::sycl::stream &) noexcept' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::stream' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::stream' lvalue ParmVar +// HasStreams::s_array +// CHECK-NEXT: InitListExpr {{.*}} 'cl::sycl::stream [2]' +// element 0 +// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::stream' 'void (const cl::sycl::stream &) noexcept' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::stream' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::stream' lvalue ParmVar +// element 1 +// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::stream' 'void (const cl::sycl::stream &) noexcept' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::stream' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::stream' lvalue ParmVar +// HasStreams struct +// CHECK-NEXT: InitListExpr {{.*}} 'HasStreams' +// HasStreams::s1 +// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::stream' 'void (const cl::sycl::stream &) noexcept' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::stream' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::stream' lvalue ParmVar +// HasStreams::s_array +// CHECK-NEXT: InitListExpr {{.*}} 'cl::sycl::stream [2]' +// element 0 +// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::stream' 'void (const cl::sycl::stream &) noexcept' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::stream' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::stream' lvalue ParmVar +// element 1 +// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::stream' 'void (const cl::sycl::stream &) noexcept' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::stream' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::stream' lvalue ParmVar + + +// Calls to Init +// in_lambda __init +// CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::stream' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at + +// _in_lambda_array +// element 0 +// CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::stream [2]' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// element 1 +// CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 .__init +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream [2]' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream (*)[2]' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 .__init +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream [2]' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream (*)[2]' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 .__init +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream [2]' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream (*)[2]' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 .__init +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream [2]' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream (*)[2]' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 .__init +// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::stream' lvalue .s1 +// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at +// array: +// CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 .__init +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 .__init +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 .__init +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 .__init +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 .__init +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 .__init +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 .__init +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 .__init +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 .__init +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 .__init +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::stream [2]' lvalue .s_array +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' +// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs +// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' +// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams [2]' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 +// second element +// CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 .__init +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 .__init +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::stream [2]' lvalue .s_array +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' +// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs +// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' +// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams [2]' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 +// second element +// CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 .__init +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::stream [2]' lvalue .s_array +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' +// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs +// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' +// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams [2]' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// element 1 +// CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 .__init +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 .__init +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 ()' lvalue .__init +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 + +// Finalize +// in_lambda __finalize +// CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream [2]' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream (*)[2]' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream [2]' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream (*)[2]' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream [2]' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream (*)[2]' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream [2]' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream (*)[2]' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 {{.*}} 'cl::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 {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 {{.*}} 'cl::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 {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 {{.*}} 'cl::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 {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 {{.*}} 'cl::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 {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::stream [2]' lvalue .s_array +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' +// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs +// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' +// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams [2]' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 +// second element +// CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::stream [2]' lvalue .s_array +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' +// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs +// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' +// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams [2]' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 +// second element +// CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::stream [2]' lvalue .s_array +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' +// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs +// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' +// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams [2]' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// element 1 +// CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__finalize +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 {{.*}} 'cl::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: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 ()' lvalue .__finalize +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 From 891454bd7aa74643284522e0300a17ba2979d6f8 Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Fri, 28 Aug 2020 09:06:32 -0700 Subject: [PATCH 15/23] Refactor to make memberexprbases less error prone --- clang/lib/Sema/SemaSYCL.cpp | 38 +++++++++++++++++++------------------ 1 file changed, 20 insertions(+), 18 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 94b38fc2bc8e2..c4a60e113a018 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1655,7 +1655,6 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { using SyclKernelFieldHandler::leaveStruct; }; -// TODO: ERICH: const-correctness of all the functions? class SyclKernelBodyCreator : public SyclKernelFieldHandler { SyclKernelDeclCreator &DeclCreator; llvm::SmallVector BodyStmts; @@ -1825,7 +1824,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { addFieldInit(FD, Ty, ParamRef); } - MemberExpr *BuildMemberExpr(Expr *Base, ValueDecl *Member) { + MemberExpr *BuildMemberExpr(Expr *Base, const ValueDecl *Member) { DeclAccessPair MemberDAP = DeclAccessPair::make(Member, AS_none); MemberExpr *Result = SemaRef.BuildMemberExpr( Base, /*IsArrow */ false, SourceLocation(), NestedNameSpecifierLoc(), @@ -1836,6 +1835,17 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { return Result; } + void AddFieldMemberExpr(const FieldDecl *FD, QualType Ty) { + if (!IsArrayType(FD, Ty)) + MemberExprBases.push_back(BuildMemberExpr(MemberExprBases.back(), FD)); + } + + void RemoveFieldMemberExpr(const FieldDecl *FD, QualType Ty) { + if (!IsArrayType(FD, Ty)) + MemberExprBases.pop_back(); + } + + void createSpecialMethodCall(const CXXRecordDecl *RD, StringRef MethodName, SmallVectorImpl &AddTo) { CXXMethodDecl *Method = getMethodByName(RD, MethodName); @@ -1923,14 +1933,12 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { addFieldInit(FD, Ty, None, InitializationKind::CreateDefault(SourceLocation())); - if (!IsArrayElement(FD, Ty)) - MemberExprBases.push_back(BuildMemberExpr(MemberExprBases.back(), FD)); + AddFieldMemberExpr(FD, Ty); const auto *RecordDecl = Ty->getAsCXXRecordDecl(); createSpecialMethodCall(RecordDecl, InitMethodName, BodyStmts); - if (!IsArrayElement(FD, Ty)) - MemberExprBases.pop_back(); + RemoveFieldMemberExpr(FD, Ty); return true; } @@ -2027,8 +2035,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { const auto *StreamDecl = Ty->getAsCXXRecordDecl(); CollectionInitExprs.push_back(CreateInitListExpr(StreamDecl)); - if (!IsArrayElement(FD, Ty)) - MemberExprBases.push_back(BuildMemberExpr(MemberExprBases.back(), FD)); + AddFieldMemberExpr(FD, Ty); return true; } @@ -2041,8 +2048,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { createSpecialMethodCall(StreamDecl, InitMethodName, BodyStmts); createSpecialMethodCall(StreamDecl, FinalizeMethodName, FinalizeStmts); - if (!IsArrayElement(FD, Ty)) - MemberExprBases.pop_back(); + RemoveFieldMemberExpr(FD, Ty); CollectionInitExprs.pop_back(); return true; @@ -2052,8 +2058,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { ++ContainerDepth; addCollectionInitListExpr(Ty->getAsCXXRecordDecl()); - if (!IsArrayElement(FD, Ty)) - MemberExprBases.push_back(BuildMemberExpr(MemberExprBases.back(), FD)); + AddFieldMemberExpr(FD, Ty); return true; } @@ -2061,8 +2066,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { --ContainerDepth; CollectionInitExprs.pop_back(); - if (!IsArrayElement(FD, Ty)) - MemberExprBases.pop_back(); + RemoveFieldMemberExpr(FD, Ty); return true; } @@ -2104,8 +2108,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { // If this is the top-level array, we need to make a MemberExpr in addition // to an array subscript. - if (!IsArrayElement(FD, ArrayType)) - MemberExprBases.push_back(BuildMemberExpr(MemberExprBases.back(), FD)); + AddFieldMemberExpr(FD, Ty); return true; } @@ -2147,8 +2150,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { MemberExprBases.pop_back(); // Remove the field access expr as well. - if (!IsArrayElement(FD, ArrayType)) - MemberExprBases.pop_back(); + AddFieldMemberExpr(FD, Ty); return true; } From 307b07a70fa29fe386ce80bcbfe03918c0ac3840 Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Fri, 28 Aug 2020 09:06:50 -0700 Subject: [PATCH 16/23] Clang-format --- clang/lib/Sema/SemaSYCL.cpp | 1 - clang/test/SemaSYCL/Inputs/sycl.hpp | 1 - clang/test/SemaSYCL/streams.cpp | 13 ++++--------- 3 files changed, 4 insertions(+), 11 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index c4a60e113a018..cca186023921d 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1845,7 +1845,6 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { MemberExprBases.pop_back(); } - void createSpecialMethodCall(const CXXRecordDecl *RD, StringRef MethodName, SmallVectorImpl &AddTo) { CXXMethodDecl *Method = getMethodByName(RD, MethodName); diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp index f8178ad26e2da..2d40dd806593b 100644 --- a/clang/test/SemaSYCL/Inputs/sycl.hpp +++ b/clang/test/SemaSYCL/Inputs/sycl.hpp @@ -229,7 +229,6 @@ class stream { void __finalize() {} }; - namespace ONEAPI { namespace experimental { template diff --git a/clang/test/SemaSYCL/streams.cpp b/clang/test/SemaSYCL/streams.cpp index dd3d93818398f..c703ad147b1d3 100644 --- a/clang/test/SemaSYCL/streams.cpp +++ b/clang/test/SemaSYCL/streams.cpp @@ -12,7 +12,7 @@ using namespace cl::sycl; handler H; struct HasStreams { - stream s1{0,0,H}; + stream s1{0, 0, H}; stream s_array[2] = {{0, 0, H}, {0, 0, H}}; }; @@ -22,9 +22,9 @@ struct HasArrayOfHasStreams { }; void use() { - stream in_lambda {0,0,H}; + stream in_lambda{0, 0, H}; stream in_lambda_array[2] = {{0, 0, H}, {0, 0, H}}; - stream in_lambda_mdarray[2][2] = {{{0, 0, H}, {0, 0, H}},{{0, 0, H}, {0, 0, H}}}; + stream in_lambda_mdarray[2][2] = {{{0, 0, H}, {0, 0, H}}, {{0, 0, H}, {0, 0, H}}}; HasStreams Struct; HasArrayOfHasStreams haohs; @@ -40,7 +40,7 @@ void use() { haohs.hs[0].s1.use(); haohs_array[0].hs[0].s1.use(); }); - } +} // Function Declaration // CHECK: FunctionDecl {{.*}}stream_test{{.*}} @@ -222,7 +222,6 @@ void use() { // CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::stream' lvalue // CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::stream' lvalue ParmVar - // Calls to Init // in_lambda __init // CHECK: CXXMemberCallExpr {{.*}} 'void' @@ -319,7 +318,6 @@ void use() { // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 - // HasArrayOfHasStreams // First element // CHECK: CXXMemberCallExpr {{.*}} 'void' @@ -394,7 +392,6 @@ void use() { // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 - // HasArrayOfHasStreams array // First element // CHECK: CXXMemberCallExpr {{.*}} 'void' @@ -673,7 +670,6 @@ void use() { // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 - // HasArrayOfHasStreams // First element // CHECK: CXXMemberCallExpr {{.*}} 'void' @@ -748,7 +744,6 @@ void use() { // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 - // HasArrayOfHasStreams array // First element // CHECK: CXXMemberCallExpr {{.*}} 'void' From 865d275e0ae7a5a3c59bcaf3879dbdff56cc4b56 Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Fri, 28 Aug 2020 09:09:39 -0700 Subject: [PATCH 17/23] Woops, fixed a build error, guess I committed too early --- clang/lib/Sema/SemaSYCL.cpp | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index cca186023921d..f8abc76690b28 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1824,7 +1824,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { addFieldInit(FD, Ty, ParamRef); } - MemberExpr *BuildMemberExpr(Expr *Base, const ValueDecl *Member) { + MemberExpr *BuildMemberExpr(Expr *Base, ValueDecl *Member) { DeclAccessPair MemberDAP = DeclAccessPair::make(Member, AS_none); MemberExpr *Result = SemaRef.BuildMemberExpr( Base, /*IsArrow */ false, SourceLocation(), NestedNameSpecifierLoc(), @@ -1835,13 +1835,13 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { return Result; } - void AddFieldMemberExpr(const FieldDecl *FD, QualType Ty) { - if (!IsArrayType(FD, Ty)) + void AddFieldMemberExpr(FieldDecl *FD, QualType Ty) { + if (!IsArrayElement(FD, Ty)) MemberExprBases.push_back(BuildMemberExpr(MemberExprBases.back(), FD)); } void RemoveFieldMemberExpr(const FieldDecl *FD, QualType Ty) { - if (!IsArrayType(FD, Ty)) + if (!IsArrayElement(FD, Ty)) MemberExprBases.pop_back(); } @@ -2107,7 +2107,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { // If this is the top-level array, we need to make a MemberExpr in addition // to an array subscript. - AddFieldMemberExpr(FD, Ty); + AddFieldMemberExpr(FD, ArrayType); return true; } @@ -2149,7 +2149,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { MemberExprBases.pop_back(); // Remove the field access expr as well. - AddFieldMemberExpr(FD, Ty); + RemoveFieldMemberExpr(FD, ArrayType); return true; } From 766887af2efca242bda5c24140d879ee240c60d1 Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Mon, 31 Aug 2020 06:43:29 -0700 Subject: [PATCH 18/23] Fix a couple of @Fznamznon comments --- clang/lib/Sema/SemaSYCL.cpp | 5 +++++ clang/test/SemaSYCL/streams.cpp | 2 +- 2 files changed, 6 insertions(+), 1 deletion(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index f8abc76690b28..9b172aa125f83 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1660,6 +1660,11 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { llvm::SmallVector BodyStmts; llvm::SmallVector CollectionInitExprs; llvm::SmallVector FinalizeStmts; + // This collection contains the information required to add/remove information + // about arrays as we enter them. The InitializedEntity component is + // necessary for initializing child members. uin64_t is the index of the + // current element being worked on, which is updated every time we visit + // nextElement. llvm::SmallVector, 8> ArrayInfos; VarDecl *KernelObjClone; InitializedEntity VarEntity; diff --git a/clang/test/SemaSYCL/streams.cpp b/clang/test/SemaSYCL/streams.cpp index c703ad147b1d3..2fb4f37775a7d 100644 --- a/clang/test/SemaSYCL/streams.cpp +++ b/clang/test/SemaSYCL/streams.cpp @@ -2,7 +2,7 @@ #include -template +template __attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { kernelFunc(); } From f208da77531a4768809ee6dd6df500be0fb1835c Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Mon, 31 Aug 2020 07:10:55 -0700 Subject: [PATCH 19/23] Update 'streams' test to have accessor in it --- clang/test/SemaSYCL/Inputs/sycl.hpp | 1 + clang/test/SemaSYCL/streams.cpp | 348 +++++++++++++++++++++++++++- 2 files changed, 348 insertions(+), 1 deletion(-) diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp index 2d40dd806593b..84c8ef83514d6 100644 --- a/clang/test/SemaSYCL/Inputs/sycl.hpp +++ b/clang/test/SemaSYCL/Inputs/sycl.hpp @@ -219,6 +219,7 @@ class handler { }; class stream { + accessor acc; public: stream(unsigned long BufferSize, unsigned long MaxStatementSize, handler &CGH) {} diff --git a/clang/test/SemaSYCL/streams.cpp b/clang/test/SemaSYCL/streams.cpp index 2fb4f37775a7d..0691acfcda95a 100644 --- a/clang/test/SemaSYCL/streams.cpp +++ b/clang/test/SemaSYCL/streams.cpp @@ -222,8 +222,15 @@ void use() { // CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::stream' lvalue // CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::stream' lvalue ParmVar -// Calls to Init +// 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: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void (cl::sycl::accessor{{.*}})' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc +// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::stream' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at + // CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init // CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::stream' lvalue . @@ -231,6 +238,15 @@ void use() { // _in_lambda_array // element 0 +// CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void (cl::sycl::accessor{{.*}})' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::stream [2]' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 + // CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue @@ -238,7 +254,17 @@ void use() { // CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::stream [2]' lvalue // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 + // element 1 +// CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void (cl::sycl::accessor{{.*}})' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::stream [2]' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 + // CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue @@ -250,6 +276,17 @@ void use() { // _in_lambda_mdarray // [0][0] // CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void (cl::sycl::accessor{{.*}})' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream [2]' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream (*)[2]' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' @@ -261,6 +298,17 @@ void use() { // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 // [0][1] // CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void (cl::sycl::accessor{{.*}})' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream [2]' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream (*)[2]' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' @@ -272,6 +320,17 @@ void use() { // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 // [1][0] // CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void (cl::sycl::accessor{{.*}})' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream [2]' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream (*)[2]' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' @@ -283,6 +342,17 @@ void use() { // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 // [1][1] // CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void (cl::sycl::accessor{{.*}})' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream [2]' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream (*)[2]' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' @@ -295,12 +365,27 @@ void use() { // HasStreams // CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void (cl::sycl::accessor{{.*}})' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc +// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::stream' lvalue .s1 +// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at +// CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init // CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::stream' lvalue .s1 // CHECK-NEXT: MemberExpr {{.*}} 'HasStreams' lvalue . // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at // array: // CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void (cl::sycl::accessor{{.*}})' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' @@ -310,6 +395,15 @@ void use() { // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 // element 1 // CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void (cl::sycl::accessor{{.*}})' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' @@ -321,6 +415,16 @@ void use() { // HasArrayOfHasStreams // First element // CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void (cl::sycl::accessor{{.*}})' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc +// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::stream' lvalue .s1 +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' +// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs +// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init // CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::stream' lvalue .s1 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue @@ -331,6 +435,19 @@ void use() { // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 // array: // CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void (cl::sycl::accessor{{.*}})' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::stream [2]' lvalue .s_array +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' +// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs +// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' @@ -344,6 +461,19 @@ void use() { // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 // element 1 // CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void (cl::sycl::accessor{{.*}})' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::stream [2]' lvalue .s_array +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' +// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs +// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 +// CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' @@ -357,6 +487,16 @@ void use() { // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 // second element // CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void (cl::sycl::accessor{{.*}})' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc +// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::stream' lvalue .s1 +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' +// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs +// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 +// CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init // CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::stream' lvalue .s1 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue @@ -367,6 +507,19 @@ void use() { // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 // array: // CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void (cl::sycl::accessor{{.*}})' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::stream [2]' lvalue .s_array +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' +// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs +// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' @@ -380,6 +533,19 @@ void use() { // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 // element 1 // CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void (cl::sycl::accessor{{.*}})' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 +// CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' @@ -395,6 +561,19 @@ void use() { // HasArrayOfHasStreams array // First element // CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void (cl::sycl::accessor{{.*}})' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc +// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::stream' lvalue .s1 +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' +// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs +// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' +// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams [2]' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init // CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::stream' lvalue .s1 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue @@ -408,6 +587,22 @@ void use() { // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 // array: // CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void (cl::sycl::accessor{{.*}})' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::stream [2]' lvalue .s_array +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' +// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs +// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' +// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams [2]' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' @@ -424,6 +619,22 @@ void use() { // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 // element 1 // CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void (cl::sycl::accessor{{.*}})' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::stream [2]' lvalue .s_array +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' +// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs +// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' +// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams [2]' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 +// CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' @@ -440,6 +651,19 @@ void use() { // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 // second element // CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void (cl::sycl::accessor{{.*}})' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc +// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::stream' lvalue .s1 +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' +// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs +// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' +// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams [2]' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 +// CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init // CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::stream' lvalue .s1 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue @@ -453,6 +677,22 @@ void use() { // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 // array: // CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void (cl::sycl::accessor{{.*}})' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::stream [2]' lvalue .s_array +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' +// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs +// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' +// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams [2]' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' @@ -469,6 +709,22 @@ void use() { // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 // element 1 // CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void (cl::sycl::accessor{{.*}})' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::stream [2]' lvalue .s_array +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' +// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs +// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' +// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams [2]' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 +// CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' @@ -485,6 +741,19 @@ void use() { // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 // second element // CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void (cl::sycl::accessor{{.*}})' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc +// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::stream' lvalue .s1 +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' +// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs +// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' +// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams [2]' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init // CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::stream' lvalue .s1 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue @@ -498,6 +767,22 @@ void use() { // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 // array: // CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void (cl::sycl::accessor{{.*}})' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::stream [2]' lvalue .s_array +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' +// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs +// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' +// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams [2]' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' @@ -514,6 +799,22 @@ void use() { // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 // element 1 // CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void (cl::sycl::accessor{{.*}})' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::stream [2]' lvalue .s_array +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasStreams *' +// CHECK-NEXT: MemberExpr {{.*}} 'HasStreams [2]' lvalue .hs +// CHECK-NEXT: ArraySubscriptExpr{{.*}} 'HasArrayOfHasStreams' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'HasArrayOfHasStreams *' +// CHECK-NEXT: MemberExpr {{.*}} 'HasArrayOfHasStreams [2]' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 +// CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 +// CHECK: CXXMemberCallExpr {{.*}} 'void' // CHECK-NEXT: MemberExpr {{.*}} 'void ()' lvalue .__init // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' @@ -530,6 +831,19 @@ void use() { // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 // second element // CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void (cl::sycl::accessor{{.*}})' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 {{.*}} 'cl::sycl::stream' lvalue .s1 // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'HasStreams' lvalue @@ -543,6 +857,22 @@ void use() { // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 1 // array: // CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void (cl::sycl::accessor{{.*}})' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 {{.*}} 'cl::sycl::stream' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' @@ -559,6 +889,22 @@ void use() { // CHECK-NEXT: IntegerLiteral {{.*}} '{{.*}}' 0 // element 1 // CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void (cl::sycl::accessor{{.*}})' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'accessor<{{.*}}' lvalue .acc +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'cl::sycl::stream' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' +// CHECK-NEXT: MemberExpr {{.*}} 'cl::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 {{.*}} 'cl::sycl::stream' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'cl::sycl::stream *' From 047ab16cfd249b9e54c4b9bba86c64afb5e77fe9 Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Mon, 31 Aug 2020 07:20:43 -0700 Subject: [PATCH 20/23] Code style fixes, function names start with a lower case --- clang/lib/Sema/SemaSYCL.cpp | 70 ++++++++++++++++++------------------- 1 file changed, 35 insertions(+), 35 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 9b172aa125f83..b26e4e7305bda 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -896,26 +896,26 @@ class KernelObjVisitor { } template - void VisitArrayElementImpl(const CXXRecordDecl *Owner, FieldDecl *ArrayField, + void visitArrayElementImpl(const CXXRecordDecl *Owner, FieldDecl *ArrayField, QualType ElementTy, uint64_t Index, Handlers &... handlers) { (void)std::initializer_list{ (handlers.nextElement(ElementTy, Index), 0)...}; - VisitField(Owner, ArrayField, ElementTy, handlers...); + visitField(Owner, ArrayField, ElementTy, handlers...); } template - void VisitFirstArrayElement(const CXXRecordDecl *Owner, FieldDecl *ArrayField, + void visitFirstArrayElement(const CXXRecordDecl *Owner, FieldDecl *ArrayField, QualType ElementTy, Handlers &... handlers) { - VisitArrayElementImpl(Owner, ArrayField, ElementTy, 0, handlers...); + visitArrayElementImpl(Owner, ArrayField, ElementTy, 0, handlers...); } template - void VisitNthArrayElement(const CXXRecordDecl *Owner, FieldDecl *ArrayField, + void visitNthArrayElement(const CXXRecordDecl *Owner, FieldDecl *ArrayField, QualType ElementTy, uint64_t Index, Handlers &... handlers); template - void VisitArray(const CXXRecordDecl *Owner, FieldDecl *Field, + void visitArray(const CXXRecordDecl *Owner, FieldDecl *Field, QualType ArrayTy, Handlers &... handlers) { // Array workflow is: // handleArrayType @@ -939,16 +939,16 @@ class KernelObjVisitor { (void)std::initializer_list{ (handlers.enterArray(Field, ArrayTy, ET), 0)...}; - VisitFirstArrayElement(Owner, Field, ET, handlers...); + visitFirstArrayElement(Owner, Field, ET, handlers...); for (uint64_t Index = 1; Index < ElemCount; ++Index) - VisitNthArrayElement(Owner, Field, ET, Index, handlers...); + visitNthArrayElement(Owner, Field, ET, Index, handlers...); (void)std::initializer_list{ (handlers.leaveArray(Field, ArrayTy, ET), 0)...}; } template - void VisitField(const CXXRecordDecl *Owner, FieldDecl *Field, + void visitField(const CXXRecordDecl *Owner, FieldDecl *Field, QualType FieldTy, Handlers &... handlers) { if (Util::isSyclAccessorType(FieldTy)) KF_FOR_EACH(handleSyclAccessorType, Field, FieldTy); @@ -978,7 +978,7 @@ class KernelObjVisitor { else if (FieldTy->isPointerType()) KF_FOR_EACH(handlePointerType, Field, FieldTy); else if (FieldTy->isArrayType()) - VisitArray(Owner, Field, FieldTy, handlers...); + visitArray(Owner, Field, FieldTy, handlers...); else if (FieldTy->isScalarType() || FieldTy->isVectorType()) KF_FOR_EACH(handleScalarType, Field, FieldTy); else @@ -999,7 +999,7 @@ class KernelObjVisitor { template void VisitRecordFields(const CXXRecordDecl *Owner, Handlers &... handlers) { for (const auto Field : Owner->fields()) - VisitField(Owner, Field, Field->getType(), handlers...); + visitField(Owner, Field, Field->getType(), handlers...); } #undef KF_FOR_EACH }; @@ -1131,7 +1131,7 @@ void KernelObjVisitor::VisitUnion(const CXXRecordDecl *Owner, ParentTy &Parent, } template -void KernelObjVisitor::VisitNthArrayElement(const CXXRecordDecl *Owner, +void KernelObjVisitor::visitNthArrayElement(const CXXRecordDecl *Owner, FieldDecl *ArrayField, QualType ElementTy, uint64_t Index, Handlers &... handlers) { @@ -1139,7 +1139,7 @@ void KernelObjVisitor::VisitNthArrayElement(const CXXRecordDecl *Owner, // constexpr' starting in C++17. Until then, we have to count on the // optimizer to realize "if (false)" is a dead branch. if (AnyTrue::Value) - VisitArrayElementImpl( + visitArrayElementImpl( Owner, ArrayField, ElementTy, Index, HandlerFilter(handlers) .Handler...); @@ -1778,7 +1778,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { // is an element of an array. This will determine whether we do // MemberExprBases in some cases or not, AND determines how we initialize // values. - bool IsArrayElement(const FieldDecl *FD, QualType Ty) const { + bool isArrayElement(const FieldDecl *FD, QualType Ty) const { return !SemaRef.getASTContext().hasSameType(FD->getType(), Ty); } @@ -1786,7 +1786,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { // field, returns a normal member initializer, if we're in a sub-array of a MD // array, returns an element initializer. InitializedEntity getFieldEntity(FieldDecl *FD, QualType Ty) { - if (IsArrayElement(FD, Ty)) + if (isArrayElement(FD, Ty)) return InitializedEntity::InitializeElement(SemaRef.getASTContext(), ArrayInfos.back().second, ArrayInfos.back().first); @@ -1840,13 +1840,13 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { return Result; } - void AddFieldMemberExpr(FieldDecl *FD, QualType Ty) { - if (!IsArrayElement(FD, Ty)) + void addFieldMemberExpr(FieldDecl *FD, QualType Ty) { + if (!isArrayElement(FD, Ty)) MemberExprBases.push_back(BuildMemberExpr(MemberExprBases.back(), FD)); } - void RemoveFieldMemberExpr(const FieldDecl *FD, QualType Ty) { - if (!IsArrayElement(FD, Ty)) + void removeFieldMemberExpr(const FieldDecl *FD, QualType Ty) { + if (!isArrayElement(FD, Ty)) MemberExprBases.pop_back(); } @@ -1891,14 +1891,14 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { addCollectionInitListExpr(QualType(RD->getTypeForDecl(), 0), NumInitExprs); } - InitListExpr *CreateInitListExpr(const CXXRecordDecl *RD) { + InitListExpr *createInitListExpr(const CXXRecordDecl *RD) { const ASTRecordLayout &Info = SemaRef.getASTContext().getASTRecordLayout(RD); uint64_t NumInitExprs = Info.getFieldCount() + RD->getNumBases(); - return CreateInitListExpr(QualType(RD->getTypeForDecl(), 0), NumInitExprs); + return createInitListExpr(QualType(RD->getTypeForDecl(), 0), NumInitExprs); } - InitListExpr *CreateInitListExpr(QualType InitTy, uint64_t NumChildInits) { + InitListExpr *createInitListExpr(QualType InitTy, uint64_t NumChildInits) { InitListExpr *ILE = new (SemaRef.getASTContext()) InitListExpr( SemaRef.getASTContext(), SourceLocation(), {}, SourceLocation()); ILE->reserveInits(SemaRef.getASTContext(), NumChildInits); @@ -1911,7 +1911,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { // to append into. void addCollectionInitListExpr(QualType InitTy, uint64_t NumChildInits) { - InitListExpr *ILE = CreateInitListExpr(InitTy, NumChildInits); + InitListExpr *ILE = createInitListExpr(InitTy, NumChildInits); InitListExpr *ParentILE = CollectionInitExprs.back(); ParentILE->updateInit(SemaRef.getASTContext(), ParentILE->getNumInits(), ILE); @@ -1937,12 +1937,12 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { addFieldInit(FD, Ty, None, InitializationKind::CreateDefault(SourceLocation())); - AddFieldMemberExpr(FD, Ty); + addFieldMemberExpr(FD, Ty); const auto *RecordDecl = Ty->getAsCXXRecordDecl(); createSpecialMethodCall(RecordDecl, InitMethodName, BodyStmts); - RemoveFieldMemberExpr(FD, Ty); + removeFieldMemberExpr(FD, Ty); return true; } @@ -1963,7 +1963,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { DC.getKernelDecl(), KernelObj)), VarEntity(InitializedEntity::InitializeVariable(KernelObjClone)), KernelObj(KernelObj), KernelCallerFunc(KernelCallerFunc) { - CollectionInitExprs.push_back(CreateInitListExpr(KernelObj)); + CollectionInitExprs.push_back(createInitListExpr(KernelObj)); markParallelWorkItemCalls(); Stmt *DS = new (S.Context) DeclStmt(DeclGroupRef(KernelObjClone), @@ -2037,9 +2037,9 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { ++ContainerDepth; // Add a dummy init expression to catch the accessor initializers. const auto *StreamDecl = Ty->getAsCXXRecordDecl(); - CollectionInitExprs.push_back(CreateInitListExpr(StreamDecl)); + CollectionInitExprs.push_back(createInitListExpr(StreamDecl)); - AddFieldMemberExpr(FD, Ty); + addFieldMemberExpr(FD, Ty); return true; } @@ -2052,7 +2052,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { createSpecialMethodCall(StreamDecl, InitMethodName, BodyStmts); createSpecialMethodCall(StreamDecl, FinalizeMethodName, FinalizeStmts); - RemoveFieldMemberExpr(FD, Ty); + removeFieldMemberExpr(FD, Ty); CollectionInitExprs.pop_back(); return true; @@ -2062,7 +2062,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { ++ContainerDepth; addCollectionInitListExpr(Ty->getAsCXXRecordDecl()); - AddFieldMemberExpr(FD, Ty); + addFieldMemberExpr(FD, Ty); return true; } @@ -2070,7 +2070,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { --ContainerDepth; CollectionInitExprs.pop_back(); - RemoveFieldMemberExpr(FD, Ty); + removeFieldMemberExpr(FD, Ty); return true; } @@ -2112,7 +2112,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { // If this is the top-level array, we need to make a MemberExpr in addition // to an array subscript. - AddFieldMemberExpr(FD, ArrayType); + addFieldMemberExpr(FD, ArrayType); return true; } @@ -2154,7 +2154,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { MemberExprBases.pop_back(); // Remove the field access expr as well. - RemoveFieldMemberExpr(FD, ArrayType); + removeFieldMemberExpr(FD, ArrayType); return true; } @@ -2170,7 +2170,7 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { // A series of functions to calculate the change in offset based on the type. int64_t offsetOf(const FieldDecl *FD, QualType ArgTy) const { - return IsArrayElement(FD, ArgTy) + return isArrayElement(FD, ArgTy) ? 0 : SemaRef.getASTContext().getFieldOffset(FD) / 8; } @@ -2198,7 +2198,7 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { // is an element of an array. This will determine whether we do // MemberExprBases in some cases or not, AND determines how we initialize // values. - bool IsArrayElement(const FieldDecl *FD, QualType Ty) const { + bool isArrayElement(const FieldDecl *FD, QualType Ty) const { return !SemaRef.getASTContext().hasSameType(FD->getType(), Ty); } From 33f92b03555e2ca4a0dad4989422d57fbb594614 Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Mon, 31 Aug 2020 07:21:16 -0700 Subject: [PATCH 21/23] Clang-format --- clang/test/SemaSYCL/Inputs/sycl.hpp | 1 + 1 file changed, 1 insertion(+) diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp index 84c8ef83514d6..65a77b01f165a 100644 --- a/clang/test/SemaSYCL/Inputs/sycl.hpp +++ b/clang/test/SemaSYCL/Inputs/sycl.hpp @@ -220,6 +220,7 @@ class handler { class stream { accessor acc; + public: stream(unsigned long BufferSize, unsigned long MaxStatementSize, handler &CGH) {} From 587cc77ee638a5a25529e56ffcd28d5dd7483328 Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Mon, 31 Aug 2020 07:45:11 -0700 Subject: [PATCH 22/23] More function/veriable name fixes --- clang/lib/Sema/SemaSYCL.cpp | 128 ++++++++++++++++++------------------ 1 file changed, 64 insertions(+), 64 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index b26e4e7305bda..e811e4926d481 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -791,17 +791,17 @@ template using bind_param_t = typename bind_param::type; class KernelObjVisitor { Sema &SemaRef; - template + template void VisitUnionImpl(const CXXRecordDecl *Owner, ParentTy &Parent, - const CXXRecordDecl *Wrapper, Handlers &... handlers) { + const CXXRecordDecl *Wrapper, HandlerTys &... Handlers) { (void)std::initializer_list{ - (handlers.enterUnion(Owner, Parent), 0)...}; - VisitRecordHelper(Wrapper, Wrapper->fields(), handlers...); + (Handlers.enterUnion(Owner, Parent), 0)...}; + VisitRecordHelper(Wrapper, Wrapper->fields(), Handlers...); (void)std::initializer_list{ - (handlers.leaveUnion(Owner, Parent), 0)...}; + (Handlers.leaveUnion(Owner, Parent), 0)...}; } - // These enable handler execution only when previous handlers succeed. + // These enable handler execution only when previous Handlers succeed. template bool handleField(FieldDecl *FD, QualType FDTy, Tn &&... tn) { bool result = true; @@ -819,72 +819,72 @@ class KernelObjVisitor { #define KF_FOR_EACH(FUNC, Item, Qt) \ handleField( \ Item, Qt, \ - std::bind(static_cast::*)( \ + std::bind(static_cast::*)( \ bind_param_t, QualType)>( \ - &std::decay_t::FUNC), \ - std::ref(handlers), _1, _2)...) + &std::decay_t::FUNC), \ + std::ref(Handlers), _1, _2)...) // The following simpler definition works with gcc 8.x and later. //#define KF_FOR_EACH(FUNC) \ // handleField(Field, FieldTy, ([&](FieldDecl *FD, QualType FDTy) { \ -// return handlers.f(FD, FDTy); \ +// return Handlers.f(FD, FDTy); \ // })...) // Parent contains the FieldDecl or CXXBaseSpecifier that was used to enter // the Wrapper structure that we're currently visiting. Owner is the parent // type (which doesn't exist in cases where it is a FieldDecl in the // 'root'), and Wrapper is the current struct being unwrapped. - template - void VisitRecord(const CXXRecordDecl *Owner, ParentTy &Parent, + template + void visitRecord(const CXXRecordDecl *Owner, ParentTy &Parent, const CXXRecordDecl *Wrapper, QualType RecordTy, - Handlers &... handlers) { + HandlerTys &... Handlers) { (void)std::initializer_list{ - (handlers.enterStruct(Owner, Parent, RecordTy), 0)...}; - VisitRecordHelper(Wrapper, Wrapper->bases(), handlers...); - VisitRecordHelper(Wrapper, Wrapper->fields(), handlers...); + (Handlers.enterStruct(Owner, Parent, RecordTy), 0)...}; + VisitRecordHelper(Wrapper, Wrapper->bases(), Handlers...); + VisitRecordHelper(Wrapper, Wrapper->fields(), Handlers...); (void)std::initializer_list{ - (handlers.leaveStruct(Owner, Parent, RecordTy), 0)...}; + (Handlers.leaveStruct(Owner, Parent, RecordTy), 0)...}; } - template + template void VisitUnion(const CXXRecordDecl *Owner, ParentTy &Parent, - const CXXRecordDecl *Wrapper, Handlers &... handlers); + const CXXRecordDecl *Wrapper, HandlerTys &... Handlers); - template + template void VisitRecordHelper(const CXXRecordDecl *Owner, clang::CXXRecordDecl::base_class_const_range Range, - Handlers &... handlers) { + HandlerTys &... Handlers) { for (const auto &Base : Range) { QualType BaseTy = Base.getType(); // Handle accessor class as base if (Util::isSyclAccessorType(BaseTy)) { (void)std::initializer_list{ - (handlers.handleSyclAccessorType(Owner, Base, BaseTy), 0)...}; + (Handlers.handleSyclAccessorType(Owner, Base, BaseTy), 0)...}; } else if (Util::isSyclStreamType(BaseTy)) { // Handle stream class as base (void)std::initializer_list{ - (handlers.handleSyclStreamType(Owner, Base, BaseTy), 0)...}; + (Handlers.handleSyclStreamType(Owner, Base, BaseTy), 0)...}; } else // For all other bases, visit the record - VisitRecord(Owner, Base, BaseTy->getAsCXXRecordDecl(), BaseTy, - handlers...); + visitRecord(Owner, Base, BaseTy->getAsCXXRecordDecl(), BaseTy, + Handlers...); } } - template + template void VisitRecordHelper(const CXXRecordDecl *Owner, RecordDecl::field_range Range, - Handlers &... handlers) { - VisitRecordFields(Owner, handlers...); + HandlerTys &... Handlers) { + VisitRecordFields(Owner, Handlers...); } // FIXME: Can this be refactored/handled some other way? - template - void VisitStreamRecord(const CXXRecordDecl *Owner, ParentTy &Parent, + template + void visitStreamRecord(const CXXRecordDecl *Owner, ParentTy &Parent, CXXRecordDecl *Wrapper, QualType RecordTy, - Handlers &... handlers) { + HandlerTys &... Handlers) { (void)std::initializer_list{ - (handlers.enterStream(Owner, Parent, RecordTy), 0)...}; + (Handlers.enterStream(Owner, Parent, RecordTy), 0)...}; for (const auto &Field : Wrapper->fields()) { QualType FieldTy = Field->getType(); // Required to initialize accessors inside streams. @@ -892,31 +892,31 @@ class KernelObjVisitor { KF_FOR_EACH(handleSyclAccessorType, Field, FieldTy); } (void)std::initializer_list{ - (handlers.leaveStream(Owner, Parent, RecordTy), 0)...}; + (Handlers.leaveStream(Owner, Parent, RecordTy), 0)...}; } - template + template void visitArrayElementImpl(const CXXRecordDecl *Owner, FieldDecl *ArrayField, QualType ElementTy, uint64_t Index, - Handlers &... handlers) { + HandlerTys &... Handlers) { (void)std::initializer_list{ - (handlers.nextElement(ElementTy, Index), 0)...}; - visitField(Owner, ArrayField, ElementTy, handlers...); + (Handlers.nextElement(ElementTy, Index), 0)...}; + visitField(Owner, ArrayField, ElementTy, Handlers...); } - template + template void visitFirstArrayElement(const CXXRecordDecl *Owner, FieldDecl *ArrayField, - QualType ElementTy, Handlers &... handlers) { - visitArrayElementImpl(Owner, ArrayField, ElementTy, 0, handlers...); + QualType ElementTy, HandlerTys &... Handlers) { + visitArrayElementImpl(Owner, ArrayField, ElementTy, 0, Handlers...); } - template + template void visitNthArrayElement(const CXXRecordDecl *Owner, FieldDecl *ArrayField, QualType ElementTy, uint64_t Index, - Handlers &... handlers); + HandlerTys &... Handlers); - template + template void visitArray(const CXXRecordDecl *Owner, FieldDecl *Field, - QualType ArrayTy, Handlers &... handlers) { + QualType ArrayTy, HandlerTys &... Handlers) { // Array workflow is: // handleArrayType // enterArray @@ -937,19 +937,19 @@ class KernelObjVisitor { assert(ElemCount > 0 && "SYCL prohibits 0 sized arrays"); (void)std::initializer_list{ - (handlers.enterArray(Field, ArrayTy, ET), 0)...}; + (Handlers.enterArray(Field, ArrayTy, ET), 0)...}; - visitFirstArrayElement(Owner, Field, ET, handlers...); + visitFirstArrayElement(Owner, Field, ET, Handlers...); for (uint64_t Index = 1; Index < ElemCount; ++Index) - visitNthArrayElement(Owner, Field, ET, Index, handlers...); + visitNthArrayElement(Owner, Field, ET, Index, Handlers...); (void)std::initializer_list{ - (handlers.leaveArray(Field, ArrayTy, ET), 0)...}; + (Handlers.leaveArray(Field, ArrayTy, ET), 0)...}; } - template + template void visitField(const CXXRecordDecl *Owner, FieldDecl *Field, - QualType FieldTy, Handlers &... handlers) { + QualType FieldTy, HandlerTys &... Handlers) { if (Util::isSyclAccessorType(FieldTy)) KF_FOR_EACH(handleSyclAccessorType, Field, FieldTy); else if (Util::isSyclSamplerType(FieldTy)) @@ -962,23 +962,23 @@ class KernelObjVisitor { CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl(); // Handle accessors in stream class. KF_FOR_EACH(handleSyclStreamType, Field, FieldTy); - VisitStreamRecord(Owner, Field, RD, FieldTy, handlers...); + visitStreamRecord(Owner, Field, RD, FieldTy, Handlers...); } else if (FieldTy->isStructureOrClassType()) { if (KF_FOR_EACH(handleStructType, Field, FieldTy)) { CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl(); - VisitRecord(Owner, Field, RD, FieldTy, handlers...); + visitRecord(Owner, Field, RD, FieldTy, Handlers...); } } else if (FieldTy->isUnionType()) { if (KF_FOR_EACH(handleUnionType, Field, FieldTy)) { CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl(); - VisitUnion(Owner, Field, RD, handlers...); + VisitUnion(Owner, Field, RD, Handlers...); } } else if (FieldTy->isReferenceType()) KF_FOR_EACH(handleReferenceType, Field, FieldTy); else if (FieldTy->isPointerType()) KF_FOR_EACH(handlePointerType, Field, FieldTy); else if (FieldTy->isArrayType()) - visitArray(Owner, Field, FieldTy, handlers...); + visitArray(Owner, Field, FieldTy, Handlers...); else if (FieldTy->isScalarType() || FieldTy->isVectorType()) KF_FOR_EACH(handleScalarType, Field, FieldTy); else @@ -988,18 +988,18 @@ class KernelObjVisitor { public: KernelObjVisitor(Sema &S) : SemaRef(S) {} - template + template void VisitRecordBases(const CXXRecordDecl *KernelFunctor, - Handlers &... handlers) { - VisitRecordHelper(KernelFunctor, KernelFunctor->bases(), handlers...); + HandlerTys &... Handlers) { + VisitRecordHelper(KernelFunctor, KernelFunctor->bases(), Handlers...); } // A visitor function that dispatches to functions as defined in // SyclKernelFieldHandler for the purposes of kernel generation. - template - void VisitRecordFields(const CXXRecordDecl *Owner, Handlers &... handlers) { + template + void VisitRecordFields(const CXXRecordDecl *Owner, HandlerTys &... Handlers) { for (const auto Field : Owner->fields()) - visitField(Owner, Field, Field->getType(), handlers...); + visitField(Owner, Field, Field->getType(), Handlers...); } #undef KF_FOR_EACH }; @@ -1761,7 +1761,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { CXXRecordDecl *WrapperStruct = ParamType->getAsCXXRecordDecl(); // Pointer field wrapped inside __wrapper_class FieldDecl *Pointer = *(WrapperStruct->field_begin()); - DRE = BuildMemberExpr(DRE, Pointer); + DRE = buildMemberExpr(DRE, Pointer); ParamType = Pointer->getType(); } @@ -1829,7 +1829,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { addFieldInit(FD, Ty, ParamRef); } - MemberExpr *BuildMemberExpr(Expr *Base, ValueDecl *Member) { + MemberExpr *buildMemberExpr(Expr *Base, ValueDecl *Member) { DeclAccessPair MemberDAP = DeclAccessPair::make(Member, AS_none); MemberExpr *Result = SemaRef.BuildMemberExpr( Base, /*IsArrow */ false, SourceLocation(), NestedNameSpecifierLoc(), @@ -1842,7 +1842,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { void addFieldMemberExpr(FieldDecl *FD, QualType Ty) { if (!isArrayElement(FD, Ty)) - MemberExprBases.push_back(BuildMemberExpr(MemberExprBases.back(), FD)); + MemberExprBases.push_back(buildMemberExpr(MemberExprBases.back(), FD)); } void removeFieldMemberExpr(const FieldDecl *FD, QualType Ty) { @@ -1866,7 +1866,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { VK_LValue, SourceLocation()); } - MemberExpr *MethodME = BuildMemberExpr(MemberExprBases.back(), Method); + MemberExpr *MethodME = buildMemberExpr(MemberExprBases.back(), Method); QualType ResultTy = Method->getReturnType(); ExprValueKind VK = Expr::getValueKindForType(ResultTy); From 5f7b8086ff13be576b03e72405bc4dd91b330f0c Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Mon, 31 Aug 2020 15:30:18 -0700 Subject: [PATCH 23/23] @elizabethandrews ' comments --- clang/lib/Sema/SemaSYCL.cpp | 21 +++++++++++++-------- 1 file changed, 13 insertions(+), 8 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index e811e4926d481..d04d3d240bbe7 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1674,7 +1674,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { // Contains a count of how many containers we're in. This is used by the // pointer-struct-wrapping code to ensure that we don't try to wrap // non-top-level pointers. - uint64_t ContainerDepth = 0; + uint64_t StructDepth = 0; // Using the statements/init expressions that we've created, this generates // the kernel body compound stmt. CompoundStmt needs to know its number of @@ -2018,7 +2018,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { bool handlePointerType(FieldDecl *FD, QualType FieldTy) final { Expr *PointerRef = - createPointerParamReferenceExpr(FD->getType(), ContainerDepth != 0); + createPointerParamReferenceExpr(FD->getType(), StructDepth != 0); addFieldInit(FD, FieldTy, PointerRef); return true; } @@ -2034,7 +2034,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { } bool enterStream(const CXXRecordDecl *RD, FieldDecl *FD, QualType Ty) final { - ++ContainerDepth; + ++StructDepth; // Add a dummy init expression to catch the accessor initializers. const auto *StreamDecl = Ty->getAsCXXRecordDecl(); CollectionInitExprs.push_back(createInitListExpr(StreamDecl)); @@ -2044,7 +2044,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { } bool leaveStream(const CXXRecordDecl *RD, FieldDecl *FD, QualType Ty) final { - --ContainerDepth; + --StructDepth; // Stream requires that its 'init' calls happen after its accessors init // calls, so add them here instead. const auto *StreamDecl = Ty->getAsCXXRecordDecl(); @@ -2059,7 +2059,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { } bool enterStruct(const CXXRecordDecl *RD, FieldDecl *FD, QualType Ty) final { - ++ContainerDepth; + ++StructDepth; addCollectionInitListExpr(Ty->getAsCXXRecordDecl()); addFieldMemberExpr(FD, Ty); @@ -2067,7 +2067,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { } bool leaveStruct(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final { - --ContainerDepth; + --StructDepth; CollectionInitExprs.pop_back(); removeFieldMemberExpr(FD, Ty); @@ -2076,7 +2076,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { bool enterStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS, QualType) final { - ++ContainerDepth; + ++StructDepth; CXXCastPath BasePath; QualType DerivedTy(RD->getTypeForDecl(), 0); @@ -2095,7 +2095,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { bool leaveStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS, QualType) final { - --ContainerDepth; + --StructDepth; MemberExprBases.pop_back(); CollectionInitExprs.pop_back(); return true; @@ -2250,6 +2250,11 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { const ParmVarDecl *SamplerArg = InitMethod->getParamDecl(0); assert(SamplerArg && "sampler __init method must have sampler parameter"); + // For samplers, we do some special work to ONLY initialize the first item + // to the InitMethod as a performance improvement presumably, so the normal + // offsetOf calculation wouldn't work correctly. Therefore, we need to call + // a version of addParam where we calculate the offset based on the true + // FieldDecl/FieldType pair, rather than the SampleArg type. addParam(FD, SamplerArg->getType(), SYCLIntegrationHeader::kind_sampler, offsetOf(FD, FieldTy)); return true;