From 9856031d130b218ba00931b5aceb802118d2a5af Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Mon, 14 Sep 2020 10:44:09 -0700 Subject: [PATCH 01/18] [SYCL] Implmenet no-decomposition for kernel types that don't need it. Kernel arguments don't need to be decomposed unless they contain a pointer or a special type, so we don't want to decompose structs/arrays containing these. This patch accomplishes that. First, we add a new attribute without a spelling that is added during the 'checking' stage, that the later vistiors can then check to see if decomposition is necessary. Next, we add a new checker to run during the checking stage that applies the attribute based on logic. Basically, a container doesn't need to be decomposed if all of its 'children' are acceptable, so we simply hold a stack of the containers to tell which need to be decomposed. This, of course, works recursively. Finally, we add some new calls to the visitor that handle the case of a 'simple array' and a 'simple struct', which are ones that don't require decomposition. --- clang/include/clang/Basic/Attr.td | 8 + clang/lib/Sema/SemaSYCL.cpp | 376 +++++++++++++++++++++++++++++- 2 files changed, 372 insertions(+), 12 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 6e1d15bed74e6..e859b30da71a2 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1198,6 +1198,14 @@ def SYCLIntelBufferLocation : InheritableAttr { let Documentation = [Undocumented]; } +def SYCLRequiresDecomposition : InheritableAttr { + // No spellings, as this is for internal use. + let Spellings = []; + let Subjects = SubjectList<[Named]>; + let LangOpts = [SYCLIsDevice, SYCLIsHost]; + let Documentation = [Undocumented]; +} + def SYCLIntelKernelArgsRestrict : InheritableAttr { let Spellings = [ CXX11<"intel", "kernel_args_restrict"> ]; let Subjects = SubjectList<[Function], ErrorDiag>; diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index fb54d682a1085..7bdebf5bf3abf 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -836,7 +836,7 @@ class KernelObjVisitor { // 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, + void visitComplexRecord(const CXXRecordDecl *Owner, ParentTy &Parent, const CXXRecordDecl *Wrapper, QualType RecordTy, HandlerTys &... Handlers) { (void)std::initializer_list{ @@ -847,6 +847,19 @@ class KernelObjVisitor { (Handlers.leaveStruct(Owner, Parent, RecordTy), 0)...}; } + template + void visitSimpleRecord(const CXXRecordDecl *Owner, ParentTy &Parent, + const CXXRecordDecl *Wrapper, QualType RecordTy, + HandlerTys &... Handlers) { + (void)std::initializer_list{ + (Handlers.handleNonDecompStruct(Owner, Parent, RecordTy), 0)...}; + } + + template + void visitRecord(const CXXRecordDecl *Owner, ParentTy &Parent, + const CXXRecordDecl *Wrapper, QualType RecordTy, + HandlerTys &... Handlers); + template void VisitUnion(const CXXRecordDecl *Owner, ParentTy &Parent, const CXXRecordDecl *Wrapper, HandlerTys &... Handlers); @@ -916,8 +929,15 @@ class KernelObjVisitor { HandlerTys &... Handlers); template - void visitArray(const CXXRecordDecl *Owner, FieldDecl *Field, - QualType ArrayTy, HandlerTys &... Handlers) { + void visitSimpleArray(const CXXRecordDecl *Owner, FieldDecl *Field, + QualType ArrayTy, HandlerTys &... Handlers) { + (void)std::initializer_list{ + (Handlers.handleSimpleArrayType(Field, ArrayTy), 0)...}; + } + + template + void visitComplexArray(const CXXRecordDecl *Owner, FieldDecl *Field, + QualType ArrayTy, HandlerTys &... Handlers) { // Array workflow is: // handleArrayType // enterArray @@ -948,6 +968,10 @@ class KernelObjVisitor { (Handlers.leaveArray(Field, ArrayTy, ET), 0)...}; } + template + void visitArray(const CXXRecordDecl *Owner, FieldDecl *Field, + QualType ArrayTy, HandlerTys &... Handlers); + template void visitField(const CXXRecordDecl *Owner, FieldDecl *Field, QualType FieldTy, HandlerTys &... Handlers) { @@ -1011,6 +1035,10 @@ class SyclKernelFieldHandlerBase { public: static constexpr const bool VisitUnionBody = false; static constexpr const bool VisitNthArrayElement = true; + // Opt-in based on whether we should visit inside Simple containers (structs, + // arrays). All of the 'check' types should likely be true, the int-header, + // and kernel decl creation types should not. + static constexpr const bool VisitInsideSimpleContainers = true; // Mark these virtual so that we can use override in the implementer classes, // despite virtual dispatch never being used. @@ -1047,6 +1075,28 @@ class SyclKernelFieldHandlerBase { // Most handlers shouldn't be handling this, just the field checker. virtual bool handleOtherType(FieldDecl *, QualType) { return true; } + // Handle a simple struct that doesn't need to be decomposed, only called on + // handlers with VisitInsideSimpleContainers as false. Replaces + // handleStructType, enterStruct, leaveStruct, and visiting of sub-elements. + virtual bool handleNonDecompStruct(const CXXRecordDecl *, FieldDecl *, + QualType) { + // TODO: Should this assert? Empty would have to do something about it. + return true; + } + virtual bool handleNonDecompStruct(const CXXRecordDecl *, + const CXXBaseSpecifier &, QualType) { + // TODO: Should this assert? Empty would have to do something about it. + return true; + } + + // Instead of handleArrayType, enterArray, leaveArray, and nextElement (plus + // descending down the elements), this function gets called in the event of an + // array containing simple elements (even in the case of an MD array). + virtual bool handleSimpleArrayType(FieldDecl *, QualType) { + // TODO: Should this assert? Empty would have to do something about it. + return true; + } + // The following are only used for keeping track of where we are in the base // class/field graph. Int Headers use this to calculate offset, most others // don't have a need for these. @@ -1118,6 +1168,14 @@ template struct AnyTrue { static constexpr bool Value = B || AnyTrue::Value; }; +template struct AllTrue; + +template struct AllTrue { static constexpr bool Value = B; }; + +template struct AllTrue { + static constexpr bool Value = B && AnyTrue::Value; +}; + template void KernelObjVisitor::VisitUnion(const CXXRecordDecl *Owner, ParentTy &Parent, const CXXRecordDecl *Wrapper, @@ -1146,6 +1204,53 @@ void KernelObjVisitor::visitNthArrayElement(const CXXRecordDecl *Owner, .Handler...); } +template +void KernelObjVisitor::visitRecord(const CXXRecordDecl *Owner, ParentTy &Parent, + const CXXRecordDecl *Wrapper, + QualType RecordTy, + HandlerTys &... Handlers) { + if (RecordTy->getAsRecordDecl()->hasAttr()) { + visitComplexRecord(Owner, Parent, Wrapper, RecordTy, Handlers...); + } else { + if (!AllTrue::Value) + visitSimpleRecord( + Owner, Parent, Wrapper, RecordTy, + HandlerFilter( + Handlers) + .Handler...); + + if (AnyTrue::Value) + visitComplexRecord( + Owner, Parent, Wrapper, RecordTy, + HandlerFilter( + Handlers) + .Handler...); + } +} + +template +void KernelObjVisitor::visitArray(const CXXRecordDecl *Owner, FieldDecl *Field, + QualType ArrayTy, HandlerTys &... Handlers) { + + if (Field->hasAttr()) { + visitComplexArray(Owner, Field, ArrayTy, Handlers...); + } else { + if (!AllTrue::Value) + visitSimpleArray( + Owner, Field, ArrayTy, + HandlerFilter( + Handlers) + .Handler...); + + if (AnyTrue::Value) + visitComplexArray( + Owner, Field, ArrayTy, + HandlerFilter( + Handlers) + .Handler...); + } +} + // A type to check the validity of all of the argument types. class SyclKernelFieldChecker : public SyclKernelFieldHandler { bool IsInvalid = false; @@ -1381,6 +1486,132 @@ class SyclKernelUnionChecker : public SyclKernelFieldHandler { } }; +// A type to mark whether a collection requires decomposition. +class SyclKernelDecompMarker : public SyclKernelFieldHandler { + llvm::SmallVector CollectionStack; + +public: + static constexpr const bool VisitUnionBody = false; + static constexpr const bool VisitNthArrayElement = false; + + SyclKernelDecompMarker(Sema &S) : SyclKernelFieldHandler(S) { + // In order to prevent checking this over and over, just add a dummy-base + // entry. + CollectionStack.push_back(true); + } + + bool handleSyclAccessorType(const CXXRecordDecl *, const CXXBaseSpecifier &, + QualType) final { + CollectionStack.back() = true; + return true; + } + bool handleSyclAccessorType(FieldDecl *, QualType) final { + CollectionStack.back() = true; + return true; + } + + bool handleSyclSamplerType(const CXXRecordDecl *, const CXXBaseSpecifier &, + QualType) final { + CollectionStack.back() = true; + return true; + } + bool handleSyclSamplerType(FieldDecl *, QualType) final { + CollectionStack.back() = true; + return true; + } + bool handleSyclSpecConstantType(FieldDecl *, QualType) final { + CollectionStack.back() = true; + return true; + } + bool handleSyclStreamType(const CXXRecordDecl *, const CXXBaseSpecifier &, + QualType) final { + CollectionStack.back() = true; + return true; + } + bool handleSyclStreamType(FieldDecl *, QualType) final { + CollectionStack.back() = true; + return true; + } + bool handleSyclHalfType(const CXXRecordDecl *, const CXXBaseSpecifier &, + QualType) final { + CollectionStack.back() = true; + return true; + } + bool handleSyclHalfType(FieldDecl *, QualType) final { + CollectionStack.back() = true; + return true; + } + + bool handlePointerType(FieldDecl *, QualType) final { + CollectionStack.back() = true; + return true; + } + + // Stream is always decomposed (and whether it gets decomposed is handled in + // handleSyclStreamType), but we need a CollectionStack entry to capture the + // accessors that get handled. + bool enterStream(const CXXRecordDecl *, FieldDecl *, QualType) final { + CollectionStack.push_back(false); + return true; + } + bool leaveStream(const CXXRecordDecl *, FieldDecl *, QualType Ty) final { + CollectionStack.pop_back(); + return true; + } + + bool enterStruct(const CXXRecordDecl *, FieldDecl *, QualType) final { + CollectionStack.push_back(false); + return true; + } + + bool leaveStruct(const CXXRecordDecl *, FieldDecl *, QualType Ty) final { + if (CollectionStack.pop_back_val()) { + RecordDecl *RD = Ty->getAsRecordDecl(); + if (!RD->hasAttr()) + RD->addAttr(SYCLRequiresDecompositionAttr::CreateImplicit( + SemaRef.getASTContext())); + CollectionStack.back() = true; + } + return true; + } + + bool enterStruct(const CXXRecordDecl *, const CXXBaseSpecifier &, + QualType) final { + CollectionStack.push_back(false); + return true; + } + + bool leaveStruct(const CXXRecordDecl *, const CXXBaseSpecifier &, + QualType Ty) final { + if (CollectionStack.pop_back_val()) { + RecordDecl *RD = Ty->getAsRecordDecl(); + if (!RD->hasAttr()) + RD->addAttr(SYCLRequiresDecompositionAttr::CreateImplicit( + SemaRef.getASTContext())); + CollectionStack.back() = true; + } + + return true; + } + + bool enterArray(FieldDecl *, QualType ArrayTy, QualType ElementTy) final { + CollectionStack.push_back(false); + return true; + } + + bool leaveArray(FieldDecl *FD, QualType ArrayTy, QualType ElementTy) final { + if (CollectionStack.pop_back_val()) { + // Cannot assert, since in MD arrays we'll end up marking them multiple + // times. + if (!FD->hasAttr()) + FD->addAttr(SYCLRequiresDecompositionAttr::CreateImplicit( + SemaRef.getASTContext())); + CollectionStack.back() = true; + } + return true; + } +}; + // A type to Create and own the FunctionDecl for the kernel. class SyclKernelDeclCreator : public SyclKernelFieldHandler { FunctionDecl *KernelDecl; @@ -1515,6 +1746,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { } public: + static constexpr const bool VisitInsideSimpleContainers = false; SyclKernelDeclCreator(Sema &S, StringRef Name, SourceLocation Loc, bool IsInline, bool IsSIMDKernel) : SyclKernelFieldHandler(S), @@ -1630,11 +1862,32 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { return true; } + bool handleSimpleArrayType(FieldDecl *FD, QualType FieldTy) final { + // Arrays are always wrapped in a struct since they cannot be passed + // directly. + RecordDecl *WrappedPointer = wrapField(FD, FieldTy); + QualType ModTy = SemaRef.getASTContext().getRecordType(WrappedPointer); + addParam(FD, ModTy); + return true; + } + bool handleScalarType(FieldDecl *FD, QualType FieldTy) final { addParam(FD, FieldTy); return true; } + bool handleNonDecompStruct(const CXXRecordDecl *, FieldDecl *FD, + QualType Ty) final { + addParam(FD, Ty); + return true; + } + + bool handleNonDecompStruct(const CXXRecordDecl *Base, + const CXXBaseSpecifier &BS, QualType Ty) final { + addParam(BS, Ty); + return true; + } + bool handleUnionType(FieldDecl *FD, QualType FieldTy) final { return handleScalarType(FD, FieldTy); } @@ -1691,6 +1944,7 @@ class SyclKernelArgsSizeChecker : public SyclKernelFieldHandler { } public: + static constexpr const bool VisitInsideSimpleContainers = false; SyclKernelArgsSizeChecker(Sema &S, SourceLocation Loc) : SyclKernelFieldHandler(S), KernelLoc(Loc) {} @@ -1730,6 +1984,23 @@ class SyclKernelArgsSizeChecker : public SyclKernelFieldHandler { return true; } + bool handleSimpleArrayType(FieldDecl *FD, QualType FieldTy) final { + addParam(FieldTy); + return true; + } + + bool handleNonDecompStruct(const CXXRecordDecl *, FieldDecl *FD, + QualType Ty) final { + addParam(Ty); + return true; + } + + bool handleNonDecompStruct(const CXXRecordDecl *Base, + const CXXBaseSpecifier &BS, QualType Ty) final { + addParam(Ty); + return true; + } + bool handleUnionType(FieldDecl *FD, QualType FieldTy) final { return handleScalarType(FD, FieldTy); } @@ -1875,6 +2146,23 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { return DRE; } + Expr *createSimpleArrayParamReferenceExpr(QualType ArrayTy) { + ParmVarDecl *KernelParameter = + DeclCreator.getParamVarDeclsForCurrentField()[0]; + QualType ParamType = KernelParameter->getOriginalType(); + Expr *DRE = SemaRef.BuildDeclRefExpr(KernelParameter, ParamType, VK_LValue, + KernelCallerSrcLoc); + + // Unwrapp the array. + CXXRecordDecl *WrapperStruct = ParamType->getAsCXXRecordDecl(); + FieldDecl *ArrayField = *(WrapperStruct->field_begin()); + DRE = buildMemberExpr(DRE, ArrayField); + + // TODO: do we need to do the L->R val conversion? I think this should + // happen automatically. + return DRE; + } + // 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 @@ -1902,8 +2190,11 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { void addFieldInit(FieldDecl *FD, QualType Ty, MultiExprArg ParamRef, InitializationKind InitKind) { - InitializedEntity Entity = getFieldEntity(FD, Ty); + addFieldInit(FD, Ty, ParamRef, InitKind, getFieldEntity(FD, Ty)); + } + void addFieldInit(FieldDecl *FD, QualType Ty, MultiExprArg ParamRef, + InitializationKind InitKind, InitializedEntity Entity) { InitializationSequence InitSeq(SemaRef, Entity, InitKind, ParamRef); ExprResult Init = InitSeq.Perform(SemaRef, Entity, InitKind, ParamRef); @@ -1924,6 +2215,22 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { Init.get()); } + void addSimpleBaseInit(const CXXBaseSpecifier &BS, QualType Ty) { + InitializationKind InitKind = + InitializationKind::CreateCopy(KernelCallerSrcLoc, KernelCallerSrcLoc); + + InitializedEntity Entity = InitializedEntity::InitializeBase( + SemaRef.Context, &BS, /*IsInheritedVirtualBase*/ false, &VarEntity); + + Expr *ParamRef = createParamReferenceExpr(); + 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(); @@ -2057,6 +2364,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { } public: + static constexpr const bool VisitInsideSimpleContainers = false; SyclKernelBodyCreator(Sema &S, SyclKernelDeclCreator &DC, const CXXRecordDecl *KernelObj, FunctionDecl *KernelCallerFunc) @@ -2126,6 +2434,30 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { return true; } + bool handleSimpleArrayType(FieldDecl *FD, QualType FieldTy) final { + Expr *ArrayRef = createSimpleArrayParamReferenceExpr(FieldTy); + InitializationKind InitKind = InitializationKind::CreateDirect({}, {}, {}); + + InitializedEntity Entity = + InitializedEntity::InitializeMember(FD, &VarEntity, /*Implicit*/ true); + + addFieldInit(FD, FieldTy, ArrayRef, InitKind, Entity); + return true; + } + + bool handleNonDecompStruct(const CXXRecordDecl *, FieldDecl *FD, + QualType Ty) final { + addSimpleFieldInit(FD, Ty); + return true; + } + + bool handleNonDecompStruct(const CXXRecordDecl *Base, + const CXXBaseSpecifier &BS, QualType Ty) final { + addSimpleBaseInit(BS, Ty); + return true; + } + + bool handleScalarType(FieldDecl *FD, QualType FieldTy) final { addSimpleFieldInit(FD, FieldTy); return true; @@ -2286,10 +2618,9 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { void addParam(const FieldDecl *FD, QualType ArgTy, SYCLIntegrationHeader::kernel_param_kind_t Kind) { - addParam(FD, ArgTy, Kind, offsetOf(FD, ArgTy)); + addParam(ArgTy, Kind, offsetOf(FD, ArgTy)); } - void addParam(const FieldDecl *FD, QualType ArgTy, - SYCLIntegrationHeader::kernel_param_kind_t Kind, + void addParam(QualType ArgTy, SYCLIntegrationHeader::kernel_param_kind_t Kind, uint64_t OffsetAdj) { uint64_t Size; Size = SemaRef.getASTContext().getTypeSizeInChars(ArgTy).getQuantity(); @@ -2306,6 +2637,7 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { } public: + static constexpr const bool VisitInsideSimpleContainers = false; SyclKernelIntHeaderCreator(Sema &S, SYCLIntegrationHeader &H, const CXXRecordDecl *KernelObj, QualType NameType, StringRef Name, StringRef StableName) @@ -2358,7 +2690,7 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { // 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, + addParam(SamplerArg->getType(), SYCLIntegrationHeader::kind_sampler, offsetOf(FD, FieldTy)); return true; } @@ -2391,6 +2723,25 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { return true; } + bool handleSimpleArrayType(FieldDecl *FD, QualType FieldTy) final { + // I think this is right, we need to always wrap arrays. + addParam(FD, FieldTy, SYCLIntegrationHeader::kind_std_layout); + return true; + } + + bool handleNonDecompStruct(const CXXRecordDecl *, FieldDecl *FD, + QualType Ty) final { + addParam(FD, Ty, SYCLIntegrationHeader::kind_std_layout); + return true; + } + + bool handleNonDecompStruct(const CXXRecordDecl *Base, + const CXXBaseSpecifier &, QualType Ty) final { + addParam(Ty, SYCLIntegrationHeader::kind_std_layout, + offsetOf(Base, Ty->getAsCXXRecordDecl())); + return true; + } + bool handleUnionType(FieldDecl *FD, QualType FieldTy) final { return handleScalarType(FD, FieldTy); } @@ -2491,8 +2842,6 @@ void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc, } } - SyclKernelFieldChecker FieldChecker(*this); - SyclKernelUnionChecker UnionChecker(*this); SyclKernelArgsSizeChecker ArgsSizeChecker(*this, Args[0]->getExprLoc()); // check that calling kernel conforms to spec QualType KernelParamTy = KernelFunc->getParamDecl(0)->getType(); @@ -2510,12 +2859,15 @@ void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc, if (KernelObj->isInvalidDecl()) return; + SyclKernelFieldChecker FieldChecker(*this); + SyclKernelUnionChecker UnionChecker(*this); + SyclKernelDecompMarker DecompMarker(*this); KernelObjVisitor Visitor{*this}; DiagnosingSYCLKernel = true; Visitor.VisitRecordBases(KernelObj, FieldChecker, UnionChecker, - ArgsSizeChecker); + ArgsSizeChecker, DecompMarker); Visitor.VisitRecordFields(KernelObj, FieldChecker, UnionChecker, - ArgsSizeChecker); + ArgsSizeChecker, DecompMarker); DiagnosingSYCLKernel = false; if (!FieldChecker.isValid() || !UnionChecker.isValid()) KernelFunc->setInvalidDecl(); From c4b2f2032beeda23da5d7a94a96ea36ed0cad19f Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Tue, 15 Sep 2020 11:12:12 -0700 Subject: [PATCH 02/18] Run clang-format --- clang/lib/Sema/SemaSYCL.cpp | 11 +++++------ 1 file changed, 5 insertions(+), 6 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 7bdebf5bf3abf..6fe5179d10efd 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -837,8 +837,8 @@ class KernelObjVisitor { // 'root'), and Wrapper is the current struct being unwrapped. template void visitComplexRecord(const CXXRecordDecl *Owner, ParentTy &Parent, - const CXXRecordDecl *Wrapper, QualType RecordTy, - HandlerTys &... Handlers) { + const CXXRecordDecl *Wrapper, QualType RecordTy, + HandlerTys &... Handlers) { (void)std::initializer_list{ (Handlers.enterStruct(Owner, Parent, RecordTy), 0)...}; VisitRecordHelper(Wrapper, Wrapper->bases(), Handlers...); @@ -849,8 +849,8 @@ class KernelObjVisitor { template void visitSimpleRecord(const CXXRecordDecl *Owner, ParentTy &Parent, - const CXXRecordDecl *Wrapper, QualType RecordTy, - HandlerTys &... Handlers) { + const CXXRecordDecl *Wrapper, QualType RecordTy, + HandlerTys &... Handlers) { (void)std::initializer_list{ (Handlers.handleNonDecompStruct(Owner, Parent, RecordTy), 0)...}; } @@ -1168,7 +1168,7 @@ template struct AnyTrue { static constexpr bool Value = B || AnyTrue::Value; }; -template struct AllTrue; +template struct AllTrue; template struct AllTrue { static constexpr bool Value = B; }; @@ -2457,7 +2457,6 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { return true; } - bool handleScalarType(FieldDecl *FD, QualType FieldTy) final { addSimpleFieldInit(FD, FieldTy); return true; From c5ee6407d86175ae7a8d2f28b944642bc4a7a5cd Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Tue, 15 Sep 2020 11:41:52 -0700 Subject: [PATCH 03/18] Fix bug that Prem found --- clang/lib/Sema/SemaSYCL.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 6fe5179d10efd..2888ef7832de4 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1173,7 +1173,7 @@ template struct AllTrue; template struct AllTrue { static constexpr bool Value = B; }; template struct AllTrue { - static constexpr bool Value = B && AnyTrue::Value; + static constexpr bool Value = B && AllTrue::Value; }; template From eb0896e48234203bdd54503b785c4225ab8e0a71 Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Wed, 16 Sep 2020 07:43:58 -0700 Subject: [PATCH 04/18] Fix union-kernel-param2.cpp by adding a pointer member so the struct gets decomped --- clang/test/SemaSYCL/union-kernel-param2.cpp | 10 ++++++---- 1 file changed, 6 insertions(+), 4 deletions(-) diff --git a/clang/test/SemaSYCL/union-kernel-param2.cpp b/clang/test/SemaSYCL/union-kernel-param2.cpp index 6fdd90f672993..0fd340b343b5e 100644 --- a/clang/test/SemaSYCL/union-kernel-param2.cpp +++ b/clang/test/SemaSYCL/union-kernel-param2.cpp @@ -25,7 +25,7 @@ int main() { float b; char c; } union_mem; - int d; + int *d; } struct_mem; a_kernel( @@ -53,9 +53,9 @@ int main() { // CHECK-NEXT: DeclRefExpr {{.*}} 'union MyUnion':'MyUnion' lvalue ParmVar {{.*}} '_arg_' 'union MyUnion':'MyUnion' // Check kernel_B parameters -// CHECK: FunctionDecl {{.*}}kernel_B{{.*}} 'void (union MyUnion, int)' +// CHECK: FunctionDecl {{.*}}kernel_B{{.*}} 'void (union MyUnion, __wrapper_class)' // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_union_mem 'union MyUnion':'MyStruct::MyUnion' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_d 'int' +// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_d '__wrapper_class' // Check kernel_B inits // CHECK-NEXT: CompoundStmt @@ -67,4 +67,6 @@ int main() { // CHECK-NEXT: ImplicitCastExpr {{.*}} 'const MyStruct::MyUnion' // CHECK-NEXT: DeclRefExpr {{.*}} 'union MyUnion':'MyStruct::MyUnion' lvalue ParmVar {{.*}} '_arg_union_mem' 'union MyUnion':'MyStruct::MyUnion' // CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_d' 'int' +// CHECK-NEXT: ImplicitCastExpr +// CHECK-NEXT: MemberExpr +// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_d' '__wrapper_class' From 911c5016faa6209429020b670e5c7160c768dbaf Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Wed, 16 Sep 2020 08:03:25 -0700 Subject: [PATCH 05/18] Move the decomp checker to its own invocation of the visitor This is necessary since the size-checker needs opt-in (so that it properly reflects the opencl kernel arguments). When they are in the same invocation, the size-checker is erronously called thinking that the first time we see a struct that it doesn't need to be decomposed. I opted to use the same visitor, since it doesn't have state. --- clang/lib/Sema/SemaSYCL.cpp | 7 +++++-- 1 file changed, 5 insertions(+), 2 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 2888ef7832de4..e17f85ea64561 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -2841,7 +2841,6 @@ void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc, } } - SyclKernelArgsSizeChecker ArgsSizeChecker(*this, Args[0]->getExprLoc()); // check that calling kernel conforms to spec QualType KernelParamTy = KernelFunc->getParamDecl(0)->getType(); if (KernelParamTy->isReferenceType()) { @@ -2858,10 +2857,14 @@ void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc, if (KernelObj->isInvalidDecl()) return; + SyclKernelDecompMarker DecompMarker(*this); SyclKernelFieldChecker FieldChecker(*this); SyclKernelUnionChecker UnionChecker(*this); - SyclKernelDecompMarker DecompMarker(*this); + SyclKernelArgsSizeChecker ArgsSizeChecker(*this, Args[0]->getExprLoc()); + KernelObjVisitor Visitor{*this}; + Visitor.VisitRecordBases(KernelObj, DecompMarker); + Visitor.VisitRecordFields(KernelObj, DecompMarker); DiagnosingSYCLKernel = true; Visitor.VisitRecordBases(KernelObj, FieldChecker, UnionChecker, ArgsSizeChecker, DecompMarker); From 23a608eafbcbad8d6af5cde59efc18a5ff91cc40 Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Wed, 16 Sep 2020 08:56:31 -0700 Subject: [PATCH 06/18] Change the call-order of the handlers, ArgsSizeChecker is what needs to be split out --- clang/lib/Sema/SemaSYCL.cpp | 15 ++++++++++----- 1 file changed, 10 insertions(+), 5 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index e17f85ea64561..9aa509fd496f5 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -2863,13 +2863,18 @@ void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc, SyclKernelArgsSizeChecker ArgsSizeChecker(*this, Args[0]->getExprLoc()); KernelObjVisitor Visitor{*this}; - Visitor.VisitRecordBases(KernelObj, DecompMarker); - Visitor.VisitRecordFields(KernelObj, DecompMarker); DiagnosingSYCLKernel = true; - Visitor.VisitRecordBases(KernelObj, FieldChecker, UnionChecker, - ArgsSizeChecker, DecompMarker); + Visitor.VisitRecordBases(KernelObj, FieldChecker, UnionChecker, DecompMarker); Visitor.VisitRecordFields(KernelObj, FieldChecker, UnionChecker, - ArgsSizeChecker, DecompMarker); + DecompMarker); + // ArgSizeChecker needs to happen after DecompMarker has completed, since it + // cares about the decomp attributes. DecompMarker cannot run before the + // others, since it counts on the FieldChecker to make sure it is visiting + // valid arrays/etc. Thus, ArgSizeChecker has its own visitation. + if (FieldChecker.isValid() && UnionChecker.isValid()) { + Visitor.VisitRecordBases(KernelObj, ArgsSizeChecker); + Visitor.VisitRecordFields(KernelObj, ArgsSizeChecker); + } DiagnosingSYCLKernel = false; if (!FieldChecker.isValid() || !UnionChecker.isValid()) KernelFunc->setInvalidDecl(); From 21c808767785cf9c4e09c1a6333c42f5cfd8d729 Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Wed, 16 Sep 2020 09:51:59 -0700 Subject: [PATCH 07/18] Fix fake-accessors, fake ones now passed as non-decomped structs --- clang/test/SemaSYCL/fake-accessors.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/clang/test/SemaSYCL/fake-accessors.cpp b/clang/test/SemaSYCL/fake-accessors.cpp index 3a7a44aa8dc68..d4dd2a0f60209 100644 --- a/clang/test/SemaSYCL/fake-accessors.cpp +++ b/clang/test/SemaSYCL/fake-accessors.cpp @@ -51,6 +51,6 @@ int main() { }); return 0; } -// CHECK: fake_accessors{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, int, int) -// CHECK: accessor_typedef{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, int, int) -// CHECK: accessor_alias{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, int, int) +// CHECK: fake_accessors{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor) +// CHECK: accessor_typedef{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor) +// CHECK: accessor_alias{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor) From 39c1561c9b65c5af160d2c9a92085ff75194ca9e Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Wed, 16 Sep 2020 10:53:21 -0700 Subject: [PATCH 08/18] Update SemaSYCL/array-kernel-param.cpp, add some additional array tests --- clang/test/SemaSYCL/array-kernel-param.cpp | 230 +++++++++++++-------- 1 file changed, 141 insertions(+), 89 deletions(-) diff --git a/clang/test/SemaSYCL/array-kernel-param.cpp b/clang/test/SemaSYCL/array-kernel-param.cpp index 945b7decd9f57..cc7611ed33839 100644 --- a/clang/test/SemaSYCL/array-kernel-param.cpp +++ b/clang/test/SemaSYCL/array-kernel-param.cpp @@ -24,6 +24,8 @@ int main() { Accessor acc[2]; int a[2]; + int *a_ptrs[2]; + struct struct_acc_t { Accessor member_acc[2]; } struct_acc; @@ -32,13 +34,13 @@ int main() { struct foo_inner { int foo_inner_x; int foo_inner_y; - int foo_inner_z[2]; + int *foo_inner_z[2]; }; struct foo { int foo_a; foo_inner foo_b[2]; - int foo_2D[2][1]; + int *foo_2D[2][1]; int foo_c; }; @@ -56,6 +58,12 @@ int main() { int local = a[1]; }); + a_kernel( + [=]() { + int local = *a_ptrs[1]; + }); + + a_kernel( [=]() { struct_acc.member_acc[2].use(); @@ -93,19 +101,40 @@ int main() { // CHECK-NEXT: MemberExpr {{.*}}__init // Check kernel_B parameters -// CHECK: FunctionDecl {{.*}}kernel_B{{.*}} 'void (int, int)' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'int' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'int' +// CHECK: FunctionDecl {{.*}}kernel_B{{.*}} 'void (__wrapper_class)' +// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ '__wrapper_class' // Check kernel_B inits // CHECK-NEXT: CompoundStmt // CHECK-NEXT: DeclStmt // CHECK-NEXT: VarDecl {{.*}} cinit // CHECK-NEXT: InitListExpr -// CHECK-NEXT: InitListExpr {{.*}} 'int [2]' -// CHECK: ImplicitCastExpr -// CHECK: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int' -// CHECK: ImplicitCastExpr -// CHECK: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int' +// CHECK-NEXT: ArrayInitLoopExpr {{.*}} 'int [2]' +// CHECK-NEXT: OpaqueValueExpr {{.*}} 'int [2]' lvalue +// CHECK-NEXT: MemberExpr {{.*}} 'int [2]' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_' '__wrapper_class' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'int' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' +// CHECK-NEXT: OpaqueValueExpr {{.*}} 'int [2]' lvalue +// CHECK-NEXT: MemberExpr {{.*}} 'int [2]' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_' '__wrapper_class' + +// Check kernel_B_ptrs parameters +// CHECK: FunctionDecl {{.*}}kernel_B_ptrs{{.*}} 'void (__global int *, __global int *)' +// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ '__global int *' +// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ '__global int *' +// Check kernel_B_ptrs inits +// CHECK-NEXT: CompoundStmt +// CHECK-NEXT: DeclStmt +// CHECK-NEXT: VarDecl {{.*}} cinit +// CHECK-NEXT: InitListExpr +// CHECK-NEXT: InitListExpr {{.*}} 'int *[2]' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' +// CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '_arg_' '__global int *' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' +// CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '_arg_' '__global int *' // Check kernel_C parameters // CHECK: FunctionDecl {{.*}}kernel_C{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' @@ -133,30 +162,30 @@ int main() { // CHECK-NEXT: MemberExpr {{.*}}__init // Check kernel_D parameters -// CHECK: FunctionDecl {{.*}}kernel_D{{.*}} 'void (int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int, int)' +// CHECK: FunctionDecl {{.*}}kernel_D{{.*}} 'void (int, int, int, __wrapper_class, __wrapper_class, int, int, __wrapper_class, __wrapper_class, __wrapper_class, __wrapper_class, int, int, int, int, __wrapper_class, __wrapper_class, int, int, __wrapper_class, __wrapper_class, __wrapper_class, __wrapper_class, int)' // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_a 'int' // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_x 'int' // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_y 'int' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z 'int' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z 'int' +// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z '__wrapper_class' +// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z '__wrapper_class' // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_x 'int' // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_y 'int' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z 'int' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z 'int' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_2D 'int' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_2D 'int' +// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z '__wrapper_class' +// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z '__wrapper_class' +// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_2D '__wrapper_class' +// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_2D '__wrapper_class' // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_c 'int' // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_a 'int' // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_x 'int' // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_y 'int' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z 'int' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z 'int' +// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z '__wrapper_class' +// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z '__wrapper_class' // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_x 'int' // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_y 'int' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z 'int' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z 'int' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_2D 'int' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_2D 'int' +// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z '__wrapper_class' +// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_inner_z '__wrapper_class' +// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_2D '__wrapper_class' +// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_2D '__wrapper_class' // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_foo_c 'int' // CHECK-NEXT: CompoundStmt // CHECK-NEXT: DeclStmt @@ -168,120 +197,143 @@ int main() { // Initializer for first element of struct_array // CHECK-NEXT: InitListExpr {{.*}} 'foo' -// CHECK-NEXT: ImplicitCastExpr +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_a' 'int' + // Initializer for struct array inside foo i.e. foo_inner foo_b[2] // CHECK-NEXT: InitListExpr {{.*}} 'foo_inner [2]' // Initializer for first element of inner struct array // CHECK-NEXT: InitListExpr {{.*}} 'foo_inner' -// CHECK-NEXT: ImplicitCastExpr +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_x' 'int' -// CHECK-NEXT: ImplicitCastExpr +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_y' 'int' -// CHECK-NEXT: InitListExpr {{.*}} 'int [2]' +// CHECK-NEXT: InitListExpr {{.*}} 'int *[2]' // CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int' +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' +// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_foo_inner_z' '__wrapper_class' // CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int' +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' +// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_foo_inner_z' '__wrapper_class' // Initializer for second element of inner struct array // CHECK-NEXT: InitListExpr {{.*}} 'foo_inner' // CHECK-NEXT: ImplicitCastExpr // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_x' 'int' // CHECK-NEXT: ImplicitCastExpr // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_y' 'int' -// CHECK-NEXT: InitListExpr {{.*}} 'int [2]' +// CHECK-NEXT: InitListExpr {{.*}} 'int *[2]' // CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int' +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' +// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_foo_inner_z' '__wrapper_class' // CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int' -// CHECK-NEXT: InitListExpr {{.*}} 'int [2][1]' -// CHECK-NEXT: InitListExpr {{.*}} 'int [1]' +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' +// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_foo_inner_z' '__wrapper_class' +// CHECK-NEXT: InitListExpr {{.*}} 'int *[2][1]' +// CHECK-NEXT: InitListExpr {{.*}} 'int *[1]' // CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_2D' 'int' -// CHECK-NEXT: InitListExpr {{.*}} 'int [1]' +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' +// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_foo_2D' '__wrapper_class' +// CHECK-NEXT: InitListExpr {{.*}} 'int *[1]' // CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_2D' 'int' +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' +// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_foo_2D' '__wrapper_class' // CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_c' 'int' +// CHECK-NEXT: DeclRefExpr{{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_c' 'int' -// Initializer for second element of struct_array +// Initializer for first element of struct_array // CHECK-NEXT: InitListExpr {{.*}} 'foo' -// CHECK-NEXT: ImplicitCastExpr +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_a' 'int' + +// Initializer for struct array inside foo i.e. foo_inner foo_b[2] // CHECK-NEXT: InitListExpr {{.*}} 'foo_inner [2]' +// Initializer for first element of inner struct array // CHECK-NEXT: InitListExpr {{.*}} 'foo_inner' -// CHECK-NEXT: ImplicitCastExpr +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_x' 'int' -// CHECK-NEXT: ImplicitCastExpr +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_y' 'int' -// CHECK-NEXT: InitListExpr {{.*}} 'int [2]' +// CHECK-NEXT: InitListExpr {{.*}} 'int *[2]' // CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int' +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' +// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_foo_inner_z' '__wrapper_class' // CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int' +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' +// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_foo_inner_z' '__wrapper_class' +// Initializer for second element of inner struct array // CHECK-NEXT: InitListExpr {{.*}} 'foo_inner' // CHECK-NEXT: ImplicitCastExpr // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_x' 'int' // CHECK-NEXT: ImplicitCastExpr // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_y' 'int' -// CHECK-NEXT: InitListExpr {{.*}} 'int [2]' +// CHECK-NEXT: InitListExpr {{.*}} 'int *[2]' // CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int' +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' +// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_foo_inner_z' '__wrapper_class' // CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_inner_z' 'int' -// CHECK-NEXT: InitListExpr {{.*}} 'int [2][1]' -// CHECK-NEXT: InitListExpr {{.*}} 'int [1]' +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' +// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_foo_inner_z' '__wrapper_class' +// CHECK-NEXT: InitListExpr {{.*}} 'int *[2][1]' +// CHECK-NEXT: InitListExpr {{.*}} 'int *[1]' // CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_2D' 'int' -// CHECK-NEXT: InitListExpr {{.*}} 'int [1]' +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' +// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_foo_2D' '__wrapper_class' +// CHECK-NEXT: InitListExpr {{.*}} 'int *[1]' // CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_2D' 'int' -// CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_c' 'int' +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' +// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_foo_2D' '__wrapper_class' // Check kernel_E parameters -// CHECK: FunctionDecl {{.*}}kernel_E{{.*}} 'void (int, int, int)' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_a 'int':'int' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_a 'int':'int' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_a 'int':'int' +// CHECK: FunctionDecl {{.*}}kernel_E{{.*}} 'void (S)' +// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'S':'S' // CHECK-NEXT: CompoundStmt // CHECK-NEXT: DeclStmt // CHECK-NEXT: VarDecl {{.*}} used '(lambda at {{.*}}array-kernel-param.cpp{{.*}})' cinit // CHECK-NEXT: InitListExpr {{.*}} '(lambda at {{.*}}array-kernel-param.cpp{{.*}})' -// CHECK-NEXT: InitListExpr {{.*}} 'S' -// CHECK-NEXT: InitListExpr {{.*}} 'int [3]' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int':'int' -// CHECK-NEXT: DeclRefExpr {{.*}} 'int':'int' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int':'int' -// CHECK-NEXT: DeclRefExpr {{.*}} 'int':'int' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int':'int' -// CHECK-NEXT: DeclRefExpr {{.*}} 'int':'int' +// CHECK-NEXT: CXXConstructExpr {{.*}} 'S':'S' 'void (const S &) noexcept' +// CHECK-NEXT: ImplicitCastExpr +// CHECK-NEXT: DeclRefExpr {{.*}} 'S':'S' lvalue ParmVar {{.*}} '_arg_' 'S':'S' // Check kernel_F parameters -// CHECK: FunctionDecl {{.*}}kernel_F{{.*}} 'void (int, int, int, int, int, int)' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'int' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'int' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'int' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'int' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'int' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ 'int' +// CHECK: FunctionDecl {{.*}}kernel_F{{.*}} 'void (__wrapper_class)' +// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ '__wrapper_class' // Check kernel_F inits // CHECK-NEXT: CompoundStmt // CHECK-NEXT: DeclStmt // CHECK-NEXT: VarDecl {{.*}} cinit // CHECK-NEXT: InitListExpr -// CHECK-NEXT: InitListExpr {{.*}} 'int [2][3]' -// CHECK-NEXT: InitListExpr {{.*}} 'int [3]' -// CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int' -// CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int' -// CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int' -// CHECK-NEXT: InitListExpr {{.*}} 'int [3]' -// CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int' -// CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int' -// CHECK-NEXT: ImplicitCastExpr -// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int' +// CHECK-NEXT: ArrayInitLoopExpr {{.*}} 'int [2][3]' +// CHECK-NEXT: OpaqueValueExpr {{.*}} 'int [2][3]' lvalue +// CHECK-NEXT: MemberExpr {{.*}} 'int [2][3]' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_' '__wrapper_class' +// CHECK-NEXT: ArrayInitLoopExpr {{.*}} 'int [3]' +// CHECK-NEXT: OpaqueValueExpr {{.*}} 'int [3]' lvalue +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'int [3]' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int (*)[3]' +// CHECK-NEXT: OpaqueValueExpr {{.*}} 'int [2][3]' lvalue +// CHECK-NEXT: MemberExpr {{.*}} 'int [2][3]' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_' '__wrapper_class' +// CHECK-NEXT: ArrayInitIndexExpr {{.*}} 'unsigned long' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'int' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' +// CHECK-NEXT: OpaqueValueExpr {{.*}} 'int [3]' lvalue +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'int [3]' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int (*)[3]' +// CHECK-NEXT: OpaqueValueExpr {{.*}} 'int [2][3]' lvalue +// CHECK-NEXT: MemberExpr {{.*}} 'int [2][3]' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_' '__wrapper_class' +// CHECK-NEXT: ArrayInitIndexExpr {{.*}} 'unsigned long' +// CHECK-NEXT: ArrayInitIndexExpr {{.*}} 'unsigned long' From dcddfa233ca257911be951ec8be8d64bb30ed790 Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Wed, 16 Sep 2020 12:27:43 -0700 Subject: [PATCH 09/18] Fixed an issue with pointer wrapping, fixed SemaSYCL/inheritence.cpp test --- clang/lib/Sema/SemaSYCL.cpp | 15 ++++++++++++--- clang/test/SemaSYCL/inheritance.cpp | 28 +++++++++++----------------- 2 files changed, 23 insertions(+), 20 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 9aa509fd496f5..a9c0b03a15a3e 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1788,6 +1788,18 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { return true; } + bool enterStruct(const CXXRecordDecl *, const CXXBaseSpecifier &BS, + QualType FieldTy) final { + ++StructDepth; + return true; + } + + bool leaveStruct(const CXXRecordDecl *, const CXXBaseSpecifier &BS, + QualType FieldTy) final { + --StructDepth; + return true; + } + bool handleSyclAccessorType(const CXXRecordDecl *, const CXXBaseSpecifier &BS, QualType FieldTy) final { const auto *RecordDecl = FieldTy->getAsCXXRecordDecl(); @@ -1919,9 +1931,6 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { } using SyclKernelFieldHandler::handleSyclHalfType; using SyclKernelFieldHandler::handleSyclSamplerType; - // Required to handle pointers inside structs - using SyclKernelFieldHandler::enterStruct; - using SyclKernelFieldHandler::leaveStruct; }; class SyclKernelArgsSizeChecker : public SyclKernelFieldHandler { diff --git a/clang/test/SemaSYCL/inheritance.cpp b/clang/test/SemaSYCL/inheritance.cpp index f8dcfaf84af94..ab1f505619fee 100644 --- a/clang/test/SemaSYCL/inheritance.cpp +++ b/clang/test/SemaSYCL/inheritance.cpp @@ -4,7 +4,7 @@ class second_base { public: - int e; + int *e; }; class InnerFieldBase { @@ -40,13 +40,11 @@ int main() { } // Check declaration of the kernel -// CHECK: derived{{.*}} 'void (int, int, int, int, int)' +// CHECK: derived{{.*}} 'void (base, __wrapper_class, int) // Check parameters of the kernel -// CHECK: ParmVarDecl {{.*}} used _arg_b 'int' -// CHECK: ParmVarDecl {{.*}} used _arg_d 'int' -// CHECK: ParmVarDecl {{.*}} used _arg_c 'int' -// CHECK: ParmVarDecl {{.*}} used _arg_e 'int' +// CHECK: ParmVarDecl {{.*}} used _arg__base 'base' +// CHECK: ParmVarDecl {{.*}} used _arg_e '__wrapper_class' // CHECK: ParmVarDecl {{.*}} used _arg_a 'int' // Check initializers for derived and base classes. @@ -54,17 +52,13 @@ int main() { // Base classes should be initialized first. // CHECK: VarDecl {{.*}} derived 'derived' cinit // CHECK-NEXT: InitListExpr {{.*}} 'derived' -// CHECK-NEXT: InitListExpr {{.*}} 'base' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' -// CHECK-NEXT: DeclRefExpr {{.*}} lvalue ParmVar {{.*}} '_arg_b' 'int' -// CHECK-NEXT: InitListExpr {{.*}} 'InnerField' -// CHECK-NEXT: InitListExpr {{.*}} 'InnerFieldBase' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' -// CHECK-NEXT: DeclRefExpr {{.*}} lvalue ParmVar {{.*}} '_arg_d' 'int' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' -// CHECK-NEXT: DeclRefExpr {{.*}} lvalue ParmVar {{.*}} '_arg_c' 'int' +// CHECK-NEXT: CXXConstructExpr {{.*}} 'base' 'void (const base &) noexcept' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const base' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} lvalue ParmVar {{.*}} '_arg__base' 'base' // CHECK-NEXT: InitListExpr {{.*}} 'second_base' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' -// CHECK-NEXT: DeclRefExpr {{.*}} lvalue ParmVar {{.*}} '_arg_e' 'int' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' +// CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} lvalue ParmVar {{.*}} '_arg_e' '__wrapper_class' // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' // CHECK-NEXT: DeclRefExpr {{.*}} lvalue ParmVar {{.*}} '_arg_a' 'int' From 08966133d0ba8a0b9b85560c73a2ee1d3fe1a916 Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Wed, 16 Sep 2020 12:31:59 -0700 Subject: [PATCH 10/18] Clang-format fix --- clang/test/SemaSYCL/array-kernel-param.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/clang/test/SemaSYCL/array-kernel-param.cpp b/clang/test/SemaSYCL/array-kernel-param.cpp index cc7611ed33839..aa0d69288292d 100644 --- a/clang/test/SemaSYCL/array-kernel-param.cpp +++ b/clang/test/SemaSYCL/array-kernel-param.cpp @@ -63,7 +63,6 @@ int main() { int local = *a_ptrs[1]; }); - a_kernel( [=]() { struct_acc.member_acc[2].use(); From eb0c8ce3e17ba960fb989e17b5c8046ce746eb5a Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Wed, 16 Sep 2020 13:12:38 -0700 Subject: [PATCH 11/18] update codegensycl/inheritence.cpp --- clang/test/CodeGenSYCL/inheritance.cpp | 60 +++++++++++--------------- 1 file changed, 24 insertions(+), 36 deletions(-) diff --git a/clang/test/CodeGenSYCL/inheritance.cpp b/clang/test/CodeGenSYCL/inheritance.cpp index 9cbea0ca8de48..b5007d1885447 100644 --- a/clang/test/CodeGenSYCL/inheritance.cpp +++ b/clang/test/CodeGenSYCL/inheritance.cpp @@ -4,7 +4,7 @@ class second_base { public: - int e; + int *e; }; class InnerFieldBase { @@ -40,45 +40,33 @@ int main() { } // Check kernel paramters -// CHECK: define spir_kernel void @{{.*}}derived(i32 %_arg_b, i32 %_arg_d, i32 %_arg_c, i32 %_arg_e, i32 %_arg_a) +// CHECK: define spir_kernel void @{{.*}}derived(%struct.{{.*}}.base* byval(%struct.{{.*}}.base) align 4 %_arg__base, %struct.{{.*}}.__wrapper_class* byval(%struct.{{.*}}.__wrapper_class) align 8 %_arg_e, i32 %_arg_a) // Check alloca for kernel paramters -// CHECK: %[[ARG_B:[a-zA-Z0-9_.]+]] = alloca i32, align 4 -// CHECK: %[[ARG_D:[a-zA-Z0-9_.]+]] = alloca i32, align 4 -// CHECK: %[[ARG_C:[a-zA-Z0-9_.]+]] = alloca i32, align 4 -// CHECK: %[[ARG_E:[a-zA-Z0-9_.]+]] = alloca i32, align 4 // CHECK: %[[ARG_A:[a-zA-Z0-9_.]+]] = alloca i32, align 4 - // Check alloca for local functor object -// CHECK: %[[LOCAL_OBJECT:[a-zA-Z0-9_.]+]] = alloca %struct.{{.*}}.derived, align 4 - -// Initialize field 'b' -// CHECK: %[[BITCAST1:[0-9]+]] = bitcast %struct.{{.*}}.derived* %[[LOCAL_OBJECT]] to %struct.{{.*}}.base* -// CHECK: %[[GEP_B:[a-zA-Z0-9]+]] = getelementptr inbounds %struct.{{.*}}.base, %struct.{{.*}}.base* %[[BITCAST1]], i32 0, i32 0 -// CHECK: %[[LOAD_B:[0-9]+]] = load i32, i32* %[[ARG_B]], align 4 -// CHECK: store i32 %[[LOAD_B]], i32* %[[GEP_B]], align 4 - -// Initialize field 'd' -// CHECK: %[[GEP_OBJ:[a-zA-Z0-9]+]] = getelementptr inbounds %struct.{{.*}}.base, %struct.{{.*}}.base* %[[BITCAST1]], i32 0, i32 1 -// CHECK: %[[BITCAST2:[0-9]+]] = bitcast %class.{{.*}}.InnerField* %[[GEP_OBJ]] to %class.{{.*}}.InnerFieldBase* -// CHECK: %[[GEP_D:[a-zA-Z0-9]+]] = getelementptr inbounds %class.{{.*}}.InnerFieldBase, %class.{{.*}}.InnerFieldBase* %[[BITCAST2]], i32 0, i32 0 -// CHECK: %[[LOAD_D:[0-9]+]] = load i32, i32* %[[ARG_D]], align 4 -// CHECK: store i32 %[[LOAD_D]], i32* %[[GEP_D]], align 4 - -// Initialize field 'c' -// CHECK: %[[GEP_C:[a-zA-Z0-9]+]] = getelementptr inbounds %class.{{.*}}.InnerField, %class.{{.*}}.InnerField* %[[GEP_OBJ]], i32 0, i32 1 -// CHECK: %[[LOAD_C:[0-9]+]] = load i32, i32* %[[ARG_C]], align 4 -// CHECK: store i32 %[[LOAD_C]], i32* %[[GEP_C]], align 4 - -// Initialize field 'e' -// CHECK: %[[BITCAST3:[0-9]+]] = bitcast %struct.{{.*}}.derived* %[[LOCAL_OBJECT]] to i8* -// CHECK: %[[GEP_DERIVED:[a-zA-Z0-9]+]] = getelementptr inbounds i8, i8* %[[BITCAST3]], i64 12 -// CHECK: %[[BITCAST4:[0-9]+]] = bitcast i8* %[[GEP_DERIVED]] to %class.{{.*}}.second_base* -// CHECK: %[[GEP_E:[a-zA-Z0-9]+]] = getelementptr inbounds %class.{{.*}}.second_base, %class.{{.*}}.second_base* %[[BITCAST4]], i32 0, i32 0 -// CHECK: %[[LOAD_E:[0-9]+]] = load i32, i32* %[[ARG_E]], align 4 -// CHECK: store i32 %[[LOAD_E]], i32* %[[GEP_E]], align 4 +// CHECK: %[[LOCAL_OBJECT:[a-zA-Z0-9_.]+]] = alloca %struct.{{.*}}.derived, align 8 +// CHECK: store i32 %_arg_a, i32* %[[ARG_A]], align 4 + +// Initialize 'base' subobject +// CHECK: %[[DERIVED_TO_BASE:.*]] = bitcast %struct.{{.*}}.derived* %[[LOCAL_OBJECT]] to %struct.{{.*}}.base* +// CHECK: %[[BASE_TO_PTR:.*]] = bitcast %struct.{{.*}}.base* %[[DERIVED_TO_BASE]] to i8* +// CHECK: %[[PARAM_TO_PTR:.*]] = bitcast %struct.{{.*}}.base* %_arg__base to i8* +// CHECK: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 %[[BASE_TO_PTR]], i8* align 4 %[[PARAM_TO_PTR]], i64 12, i1 false) + +// Initialize 'second_base' subobject +// First, derived-to-base cast with offset: +// CHECK: %[[DERIVED_PTR:.*]] = bitcast %struct.{{.*}}.derived* %[[LOCAL_OBJECT]] to i8* +// CHECK: %[[OFFSET_CALC:.*]] = getelementptr inbounds i8, i8* %[[DERIVED_PTR]], i64 16 +// CHECK: %[[TO_SECOND_BASE:.*]] = bitcast i8* %[[OFFSET_CALC]] to %class.{{.*}}.second_base* +// Initialize 'second_base::e' +// CHECK: %[[SECOND_BASE_PTR:.*]] = getelementptr inbounds %class.{{.*}}.second_base, %class.{{.*}}.second_base* %[[TO_SECOND_BASE]], i32 0, i32 0 +// CHECK: %[[PTR_TO_WRAPPER:.*]] = getelementptr inbounds %struct.{{.*}}.__wrapper_class, %struct.{{.*}}.__wrapper_class* %_arg_e, i32 0, i32 0 +// CHECK: %[[LOAD_PTR:.*]] = load i32 addrspace(1)*, i32 addrspace(1)** %[[PTR_TO_WRAPPER]] +// CHECK: %[[AS_CAST:.*]] = addrspacecast i32 addrspace(1)* %[[LOAD_PTR]] to i32 addrspace(4)* +// CHECK: store i32 addrspace(4)* %[[AS_CAST]], i32 addrspace(4)** %[[SECOND_BASE_PTR]] // Initialize field 'a' -// CHECK: %[[GEP_A:[a-zA-Z0-9]+]] = getelementptr inbounds %struct.{{.*}}.derived, %struct.{{.*}}.derived* %[[LOCAL_OBJECT]], i32 0, i32 2 +// CHECK: %[[GEP_A:[a-zA-Z0-9]+]] = getelementptr inbounds %struct.{{.*}}.derived, %struct.{{.*}}.derived* %[[LOCAL_OBJECT]], i32 0, i32 3 // CHECK: %[[LOAD_A:[0-9]+]] = load i32, i32* %[[ARG_A]], align 4 -// CHECK: store i32 %[[LOAD_A]], i32* %[[GEP_A]], align 4 +// CHECK: store i32 %[[LOAD_A]], i32* %[[GEP_A]] From 8f7fc77b43e13aa7b6b3caf89be69b052d0b7caf Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Wed, 16 Sep 2020 13:23:22 -0700 Subject: [PATCH 12/18] update CodeGenSYCL/struct_kernel_param.cpp --- clang/test/CodeGenSYCL/struct_kernel_param.cpp | 13 ++++++++----- 1 file changed, 8 insertions(+), 5 deletions(-) diff --git a/clang/test/CodeGenSYCL/struct_kernel_param.cpp b/clang/test/CodeGenSYCL/struct_kernel_param.cpp index e67915455539a..11c1526f41040 100644 --- a/clang/test/CodeGenSYCL/struct_kernel_param.cpp +++ b/clang/test/CodeGenSYCL/struct_kernel_param.cpp @@ -3,13 +3,16 @@ // CHECK: const kernel_param_desc_t kernel_signatures[] = { // CHECK-NEXT: //--- _ZTSZZ5test0vENK3$_0clERN2cl4sycl7handlerEE8MyKernel +// Accessor // CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 0 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 12 }, +// FldInt, offset to 16 because the float* causes the alignment of the structs +// to change. // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 16 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 20 }, +// FldArr // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 24 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 28 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 32 }, +// FldFloat +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 8, 32 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 12, 40 }, // CHECK-EMPTY: // CHECK-NEXT:}; @@ -21,7 +24,7 @@ using namespace cl::sycl; struct MyNestedStruct { int FldArr[1]; - float FldFloat; + float *FldFloat; }; struct MyStruct { From a681d1dd7bc6e84427a8280ceccc1c80311cdb86 Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Thu, 17 Sep 2020 07:18:30 -0700 Subject: [PATCH 13/18] Update kernel-param-pod-array test --- .../CodeGenSYCL/kernel-param-pod-array.cpp | 193 +++++++----------- 1 file changed, 72 insertions(+), 121 deletions(-) diff --git a/clang/test/CodeGenSYCL/kernel-param-pod-array.cpp b/clang/test/CodeGenSYCL/kernel-param-pod-array.cpp index 26b28a4a58cf6..c9f602261a8f7 100644 --- a/clang/test/CodeGenSYCL/kernel-param-pod-array.cpp +++ b/clang/test/CodeGenSYCL/kernel-param-pod-array.cpp @@ -46,136 +46,87 @@ int main() { // Check kernel_B parameters // CHECK: define spir_kernel void @{{.*}}kernel_B -// CHECK-SAME: i32 [[ELEM_ARG0:%[a-zA-Z0-9_]+]], -// CHECK-SAME: i32 [[ELEM_ARG1:%[a-zA-Z_]+_[0-9]+]]) +// CHECK-SAME:(%struct.{{.*}}.__wrapper_class* byval(%struct.{{.*}}.__wrapper_class) align 4 %[[ARR_ARG:.*]]) // Check local lambda object alloca -// CHECK: [[LOCAL_OBJECT:%[0-9]+]] = alloca %"class.{{.*}}.anon", align 4 - -// Check local variables created for parameters -// CHECK: store i32 [[ELEM_ARG0]], i32* [[ELEM_L0:%[a-zA-Z_]+.addr]], align 4 -// CHECK: store i32 [[ELEM_ARG1]], i32* [[ELEM_L1:%[a-zA-Z_]+.addr[0-9]*]], align 4 - -// Check init of local array -// CHECK: [[ARRAY:%[0-9]*]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[LOCAL_OBJECT]], i32 0, i32 0 -// CHECK: [[ARRAY_BEGIN:%[a-zA-Z_.]+]] = getelementptr inbounds [2 x i32], [2 x i32]* [[ARRAY]], i64 0, i64 0 -// CHECK: [[ARRAY0:%[0-9]*]] = load i32, i32* [[ELEM_L0]], align 4 -// CHECK: store i32 [[ARRAY0]], i32* [[ARRAY_BEGIN]], align 4 -// CHECK: [[ARRAY_ELEMENT:%[a-zA-Z_.]+]] = getelementptr inbounds i32, i32* %arrayinit.begin, i64 1 -// CHECK: [[ARRAY1:%[0-9]*]] = load i32, i32* [[ELEM_L1]], align 4 -// CHECK: store i32 [[ARRAY1]], i32* [[ARRAY_ELEMENT]], align 4 +// CHECK: %[[LOCAL_OBJECT:[0-9]+]] = alloca %"class.{{.*}}.anon", align 4 + +// Check for Array init loop +// CHECK: %[[LAMBDA_PTR:.+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* %[[LOCAL_OBJECT]], i32 0, i32 0 +// CHECK: %[[WRAPPER_PTR:.+]] = getelementptr inbounds %struct.{{.*}}.__wrapper_class, %struct.{{.*}}.__wrapper_class* %[[ARR_ARG]], i32 0, i32 0 +// CHECK: %[[ARRAY_BEGIN:.+]] = getelementptr inbounds [2 x i32], [2 x i32]* %[[LAMBDA_PTR]], i64 0, i64 0 +// CHECK: br label %[[ARRAYINITBODY:.+]] + +// The loop body itself +// CHECK: [[ARRAYINITBODY]]: +// CHECK: %[[ARRAYINDEX:.+]] = phi i64 [ 0, %{{.*}} ], [ %[[NEXTINDEX:.+]], %[[ARRAYINITBODY]] ] +// CHECK: %[[TARG_ARRAY_ELEM:.+]] = getelementptr inbounds i32, i32* %[[ARRAY_BEGIN]], i64 %[[ARRAYINDEX]] +// CHECK: %[[SRC_ELEM:.+]] = getelementptr inbounds [2 x i32], [2 x i32]* %[[WRAPPER_PTR]], i64 0, i64 %[[ARRAYINDEX]] +// CHECK: %[[SRC_VAL:.+]] = load i32, i32* %[[SRC_ELEM]] +// CHECK: store i32 %[[SRC_VAL]], i32* %[[TARG_ARRAY_ELEM]] +// CHECK: %[[NEXTINDEX]] = add nuw i64 %[[ARRAYINDEX]], 1 +// CHECK: %[[ISDONE:.+]] = icmp eq i64 %[[NEXTINDEX]], 2 +// CHECK: br i1 %[[ISDONE]], label %{{.*}}, label %[[ARRAYINITBODY]] // Check kernel_C parameters // CHECK: define spir_kernel void @{{.*}}kernel_C -// CHECK-SAME: i32 [[FOO1_A:%[a-zA-Z0-9_]+]], i32 [[FOO1_B1_X:%[a-zA-Z0-9_]+]], i32 [[FOO1_B1_Y:%[a-zA-Z0-9_]+]], i32 [[FOO1_B2_X:%[a-zA-Z0-9_]+]], i32 [[FOO1_B2_Y:%[a-zA-Z0-9_]+]], i32 [[FOO1_C:%[a-zA-Z0-9_]+]], -// CHECK-SAME: i32 [[FOO2_A:%[a-zA-Z0-9_]+]], i32 [[FOO2_B1_X:%[a-zA-Z0-9_]+]], i32 [[FOO2_B1_Y:%[a-zA-Z0-9_]+]], i32 [[FOO2_B2_X:%[a-zA-Z0-9_]+]], i32 [[FOO2_B2_Y:%[a-zA-Z0-9_]+]], i32 [[FOO2_C:%[a-zA-Z0-9_]+]] +// CHECK-SAME:(%struct.{{.*}}.__wrapper_class{{.*}}* byval(%struct.{{.*}}.__wrapper_class{{.*}}) align 4 %[[ARR_ARG:.*]]) // Check local lambda object alloca -// CHECK: [[KERNEL_OBJ:%[0-9]+]] = alloca %"class.{{.*}}.anon.0", align 4 - -// Check local stores -// CHECK: store i32 [[FOO1_A]], i32* [[FOO1_A_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4 -// CHECK: store i32 [[FOO1_B1_X]], i32* [[FOO1_B1_X_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4 -// CHECK: store i32 [[FOO1_B1_Y]], i32* [[FOO1_B1_Y_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4 -// CHECK: store i32 [[FOO1_B2_X]], i32* [[FOO1_B2_X_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4 -// CHECK: store i32 [[FOO1_B2_Y]], i32* [[FOO1_B2_Y_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4 -// CHECK: store i32 [[FOO1_C]], i32* [[FOO1_C_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4 -// CHECK: store i32 [[FOO2_A]], i32* [[FOO2_A_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4 -// CHECK: store i32 [[FOO2_B1_X]], i32* [[FOO2_B1_X_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4 -// CHECK: store i32 [[FOO2_B1_Y]], i32* [[FOO2_B1_Y_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4 -// CHECK: store i32 [[FOO2_B2_X]], i32* [[FOO2_B2_X_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4 -// CHECK: store i32 [[FOO2_B2_Y]], i32* [[FOO2_B2_Y_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4 -// CHECK: store i32 [[FOO2_C]], i32* [[FOO2_C_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4 - -// Check initialization of local array - -// Initialize struct_array[0].foo_a -// CHECK: [[GEP:%[a-zA-Z0-9_]+]] = getelementptr inbounds %"class.{{.*}}.anon.0", %"class.{{.*}}.anon.0"* [[KERNEL_OBJ]], i32 0, i32 0 -// CHECK: [[FOO_ARRAY_0:%[a-zA-Z_.]+]] = getelementptr inbounds [2 x %struct.{{.*}}.foo], [2 x %struct.{{.*}}.foo]* [[GEP]], i64 0, i64 0 -// CHECK: [[GEP_FOO1_A:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}.foo, %struct.{{.*}}.foo* [[FOO_ARRAY_0]], i32 0, i32 0 -// CHECK: [[LOAD_FOO1_A:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO1_A_LOCAL]], align 4 -// CHECK: store i32 [[LOAD_FOO1_A]], i32* [[GEP_FOO1_A]], align 4 - -// Initialize struct_array[0].foo_b[0].x -// CHECK: [[GEP_FOO1_B:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}.foo, %struct.{{.*}}.foo* [[FOO_ARRAY_0]], i32 0, i32 1 -// CHECK: [[B_ARRAY_0:%[a-zA-Z0-9_.]+]] = getelementptr inbounds [2 x %struct.{{.*}}foo_inner.foo_inner], [2 x %struct.{{.*}}foo_inner.foo_inner]* [[GEP_FOO1_B]], i64 0, i64 0 -// CHECK: [[GEP_FOO1_B1_X:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo_inner.foo_inner, %struct.{{.*}}foo_inner.foo_inner* [[B_ARRAY_0]], i32 0, i32 0 -// CHECK: [[LOAD_FOO1_B1_X:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO1_B1_X_LOCAL]], align 4 -// CHECK: store i32 [[LOAD_FOO1_B1_X]], i32* [[GEP_FOO1_B1_X]], align 4 - -// Initialize struct_array[0].foo_b[0].y -// CHECK: [[GEP_FOO1_B1_Y:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo_inner.foo_inner, %struct.{{.*}}foo_inner.foo_inner* [[B_ARRAY_0]], i32 0, i32 1 -// CHECK: [[LOAD_FOO1_B1_Y:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO1_B1_Y_LOCAL]], align 4 -// CHECK: store i32 [[LOAD_FOO1_B1_Y]], i32* [[GEP_FOO1_B1_Y]], align 4 - -// Initialize struct_array[0].foo_b[1].x -// CHECK: [[B_ARRAY_1:%[a-zA-Z0-9_.]+]] = getelementptr inbounds %struct.{{.*}}foo_inner.foo_inner, %struct.{{.*}}foo_inner.foo_inner* [[B_ARRAY_0]], i64 1 -// CHECK: [[GEP_FOO1_B2_X:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo_inner.foo_inner, %struct.{{.*}}foo_inner.foo_inner* [[B_ARRAY_1]], i32 0, i32 0 -// CHECK: [[LOAD_FOO1_B2_X:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO1_B2_X_LOCAL]], align 4 -// CHECK: store i32 [[LOAD_FOO1_B2_X]], i32* [[GEP_FOO1_B2_X]], align 4 - -// Initialize struct_array[0].foo_b[1].y -// CHECK: [[GEP_FOO1_B2_Y:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo_inner.foo_inner, %struct.{{.*}}foo_inner.foo_inner* [[B_ARRAY_1]], i32 0, i32 1 -// CHECK: [[LOAD_FOO1_B2_Y:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO1_B2_Y_LOCAL]], align 4 -// CHECK: store i32 [[LOAD_FOO1_B2_Y]], i32* [[GEP_FOO1_B2_Y]], align 4 - -// Initialize struct_array[0].foo_c -// CHECK: [[GEP_FOO1_C:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo.foo, %struct.{{.*}}foo.foo* [[FOO_ARRAY_0]], i32 0, i32 2 -// CHECK: [[LOAD_FOO1_C:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO1_C_LOCAL]], align 4 -// CHECK: store i32 [[LOAD_FOO1_C]], i32* [[GEP_FOO1_C]], align 4 - -// Initialize struct_array[1].foo_a -// CHECK: [[FOO_ARRAY_1:%[a-zA-Z0-9_.]+]] = getelementptr inbounds %struct._ZTS3foo.foo, %struct._ZTS3foo.foo* [[FOO_ARRAY_0]], i64 1 -// CHECK: [[GEP_FOO2_A:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo.foo, %struct.{{.*}}foo.foo* [[FOO_ARRAY_1]], i32 0, i32 0 -// CHECK: [[LOAD_FOO2_A:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO2_A_LOCAL]], align 4 -// CHECK: store i32 [[LOAD_FOO2_A]], i32* [[GEP_FOO2_A]], align 4 - -// Initialize struct_array[1].foo_b[0].x -// CHECK: [[GEP_FOO2_B:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}.foo, %struct.{{.*}}.foo* [[FOO_ARRAY_1]], i32 0, i32 1 -// CHECK: [[FOO2_B_ARRAY_0:%[a-zA-Z0-9_.]+]] = getelementptr inbounds [2 x %struct.{{.*}}foo_inner.foo_inner], [2 x %struct.{{.*}}foo_inner.foo_inner]* [[GEP_FOO2_B]], i64 0, i64 0 -// CHECK: [[GEP_FOO2_B1_X:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo_inner.foo_inner, %struct.{{.*}}foo_inner.foo_inner* [[FOO2_B_ARRAY_0]], i32 0, i32 0 -// CHECK: [[LOAD_FOO2_B1_X:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO2_B1_X_LOCAL]], align 4 -// CHECK: store i32 [[LOAD_FOO2_B1_X]], i32* [[GEP_FOO2_B1_X]] - -// Initialize struct_array[1].foo_b[0].y -// CHECK: [[GEP_FOO2_B1_Y:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo_inner.foo_inner, %struct.{{.*}}foo_inner.foo_inner* [[FOO2_B_ARRAY_0]], i32 0, i32 1 -// CHECK: [[LOAD_FOO2_B1_Y:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO2_B1_Y_LOCAL]], align 4 -// CHECK: store i32 [[LOAD_FOO2_B1_Y]], i32* [[GEP_FOO2_B1_Y]], align 4 - -// Initialize struct_array[1].foo_b[1].x -// CHECK: [[FOO2_B_ARRAY_1:%[a-zA-Z0-9_.]+]] = getelementptr inbounds %struct.{{.*}}foo_inner.foo_inner, %struct.{{.*}}foo_inner.foo_inner* [[FOO2_B_ARRAY_0]], i64 1 -// CHECK: [[GEP_FOO2_B2_X:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo_inner.foo_inner, %struct.{{.*}}foo_inner.foo_inner* [[FOO2_B_ARRAY_1]], i32 0, i32 0 -// CHECK: [[LOAD_FOO2_B2_X:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO2_B2_X_LOCAL]], align 4 -// CHECK: store i32 [[LOAD_FOO2_B2_X]], i32* [[GEP_FOO2_B2_X]], align 4 - -// Initialize struct_array[1].foo_b[1].y -// CHECK: [[GEP_FOO2_B2_Y:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo_inner.foo_inner, %struct.{{.*}}foo_inner.foo_inner* [[FOO2_B_ARRAY_1]], i32 0, i32 1 -// CHECK: [[LOAD_FOO2_B2_Y:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO2_B2_Y_LOCAL]], align 4 -// CHECK: store i32 [[LOAD_FOO2_B2_Y]], i32* [[GEP_FOO2_B2_Y]], align 4 - -// Initialize struct_array[1].foo_c -// CHECK: [[GEP_FOO2_C:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct.{{.*}}foo.foo, %struct.{{.*}}foo.foo* [[FOO_ARRAY_1]], i32 0, i32 2 -// CHECK: [[LOAD_FOO2_C:%[a-zA-Z0-9_]+]] = load i32, i32* [[FOO2_C_LOCAL]], align 4 -// CHECK: store i32 [[LOAD_FOO2_C]], i32* [[GEP_FOO2_C]], align 4 +// CHECK: %[[LOCAL_OBJECT:[0-9]+]] = alloca %"class.{{.*}}.anon{{.*}}", align 4 + +// Check for Array init loop +// CHECK: %[[LAMBDA_PTR:.+]] = getelementptr inbounds %"class.{{.*}}.anon{{.*}}", %"class.{{.*}}.anon{{.*}}"* %[[LOCAL_OBJECT]], i32 0, i32 0 +// CHECK: %[[WRAPPER_PTR:.+]] = getelementptr inbounds %struct.{{.*}}.__wrapper_class{{.*}}, %struct.{{.*}}.__wrapper_class{{.*}}* %[[ARR_ARG]], i32 0, i32 0 +// CHECK: %[[ARRAY_BEGIN:.+]] = getelementptr inbounds [2 x %struct.{{.*}}.foo], [2 x %struct.{{.*}}.foo]* %[[LAMBDA_PTR]], i64 0, i64 0 +// CHECK: br label %[[ARRAYINITBODY:.+]] + +// The loop body itself +// CHECK: [[ARRAYINITBODY]]: +// CHECK: %[[ARRAYINDEX:.+]] = phi i64 [ 0, %{{.*}} ], [ %[[NEXTINDEX:.+]], %[[ARRAYINITBODY]] ] +// CHECK: %[[TARG_ARRAY_ELEM:.+]] = getelementptr inbounds %struct.{{.*}}.foo, %struct.{{.*}}.foo* %[[ARRAY_BEGIN]], i64 %[[ARRAYINDEX]] +// CHECK: %[[SRC_ELEM:.+]] = getelementptr inbounds [2 x %struct.{{.*}}.foo], [2 x %struct.{{.*}}.foo]* %[[WRAPPER_PTR]], i64 0, i64 %[[ARRAYINDEX]] +// CHECK: %[[TARG_PTR:.+]] = bitcast %struct.{{.*}}.foo* %[[TARG_ARRAY_ELEM]] to i8* +// CHECK: %[[SRC_PTR:.+]] = bitcast %struct.{{.*}}.foo* %[[SRC_ELEM]] to i8* +// call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 %[[TARG_PTR]], i8* align %[[SRC_PTR]], i64 24, i1 false) +// CHECK: %[[NEXTINDEX]] = add nuw i64 %[[ARRAYINDEX]], 1 +// CHECK: %[[ISDONE:.+]] = icmp eq i64 %[[NEXTINDEX]], 2 +// CHECK: br i1 %[[ISDONE]], label %{{.*}}, label %[[ARRAYINITBODY]] // Check kernel_D parameters // CHECK: define spir_kernel void @{{.*}}kernel_D -// CHECK-SAME: i32 [[ARR_2D_1:%[a-zA-Z0-9_]+]], i32 [[ARR_2D_2:%[a-zA-Z0-9_]+]] +// CHECK-SAME:(%struct.{{.*}}.__wrapper_class{{.*}}* byval(%struct.{{.*}}.__wrapper_class{{.*}}) align 4 %[[ARR_ARG:.*]]) // Check local lambda object alloca -// CHECK: [[LAMBDA_OBJ:%[0-9]+]] = alloca %"class.{{.*}}.anon.1", align 4 - -// Check local stores -// CHECK: store i32 [[ARR_2D_1]], i32* [[ARR_2D_1_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4 -// CHECK: store i32 [[ARR_2D_2]], i32* [[ARR_2D_2_LOCAL:%[a-zA-Z_]+.addr[0-9]*]], align 4 - -// Check initialization of local array -// CHECK: [[GEP_ARR_2D:%[0-9]*]] = getelementptr inbounds %"class._ZTSZ4mainE3$_0.anon.1", %"class._ZTSZ4mainE3$_0.anon.1"* [[LAMBDA_OBJ]], i32 0, i32 0 -// CHECK: [[GEP_ARR_BEGIN1:%[a-zA-Z0-9_.]+]] = getelementptr inbounds [2 x [1 x i32]], [2 x [1 x i32]]* [[GEP_ARR_2D]], i64 0, i64 0 -// CHECK: [[GEP_ARR_ELEM0:%[a-zA-Z0-9_.]+]] = getelementptr inbounds [1 x i32], [1 x i32]* [[GEP_ARR_BEGIN1]], i64 0, i64 0 -// CHECK: [[ARR_2D_ELEM0:%[0-9]*]] = load i32, i32* [[ARR_2D_1_LOCAL]], align 4 -// CHECK: store i32 [[ARR_2D_ELEM0]], i32* [[GEP_ARR_ELEM0]], align 4 -// CHECK: [[GEP_ARR_BEGIN2:%[a-zA-Z_.]+]] = getelementptr inbounds [1 x i32], [1 x i32]* [[GEP_ARR_BEGIN1]], i64 1 -// CHECK: [[GEP_ARR_ELEM1:%[a-zA-Z0-9_.]+]] = getelementptr inbounds [1 x i32], [1 x i32]* [[GEP_ARR_BEGIN2]], i64 0, i64 0 -// CHECK: [[ARR_2D_ELEM1:%[0-9]*]] = load i32, i32* [[ARR_2D_2_LOCAL]], align 4 -// CHECK: store i32 [[ARR_2D_ELEM1]], i32* [[GEP_ARR_ELEM1]], align 4 +// CHECK: %[[LOCAL_OBJECT:[0-9]+]] = alloca %"class.{{.*}}.anon{{.*}}", align 4 + +// Check for Array init loop +// CHECK: %[[LAMBDA_PTR:.+]] = getelementptr inbounds %"class.{{.*}}.anon{{.*}}", %"class.{{.*}}.anon{{.*}}"* %[[LOCAL_OBJECT]], i32 0, i32 0 +// CHECK: %[[WRAPPER_PTR:.+]] = getelementptr inbounds %struct.{{.*}}.__wrapper_class{{.*}}, %struct.{{.*}}.__wrapper_class{{.*}}* %[[ARR_ARG]], i32 0, i32 0 +// CHECK: %[[ARRAY_BEGIN:.+]] = getelementptr inbounds [2 x [1 x i32]], [2 x [1 x i32]]* %[[LAMBDA_PTR]], i64 0, i64 0 +// CHECK: br label %[[ARRAYINITBODY:.+]] + +// Check Outer loop. +// CHECK: [[ARRAYINITBODY]]: +// CHECK: %[[ARRAYINDEX:.+]] = phi i64 [ 0, %{{.*}} ], [ %[[NEXTINDEX:.+]], %[[ARRAYINITEND:.+]] ] +// CHECK: %[[TARG_OUTER_ELEM:.+]] = getelementptr inbounds [1 x i32], [1 x i32]* %[[ARRAY_BEGIN]], i64 %[[ARRAYINDEX]] +// CHECK: %[[SRC_OUTER_ELEM:.+]] = getelementptr inbounds [2 x [1 x i32]], [2 x [1 x i32]]* %[[WRAPPER_PTR]], i64 0, i64 %[[ARRAYINDEX]] +// CHECK: %[[ARRAY_BEGIN_INNER:.+]] = getelementptr inbounds [1 x i32], [1 x i32]* %[[TARG_OUTER_ELEM]], i64 0, i64 0 +// CHECK: br label %[[ARRAYINITBODY_INNER:.+]] + +// Check Inner Loop +// CHECK: [[ARRAYINITBODY_INNER]]: +// CHECK: %[[ARRAYINDEX_INNER:.+]] = phi i64 [ 0, %{{.*}} ], [ %[[NEXTINDEX_INNER:.+]], %[[ARRAYINITBODY_INNER:.+]] ] +// CHECK: %[[TARG_INNER_ELEM:.+]] = getelementptr inbounds i32, i32* %[[ARRAY_BEGIN_INNER]], i64 %[[ARRAYINDEX_INNER]] +// CHECK: %[[SRC_INNER_ELEM:.+]] = getelementptr inbounds [1 x i32], [1 x i32]* %[[SRC_OUTER_ELEM]], i64 0, i64 %[[ARRAYINDEX_INNER]] +// CHECK: %[[SRC_LOAD:.+]] = load i32, i32* %[[SRC_INNER_ELEM]] +// CHECK: store i32 %[[SRC_LOAD]], i32* %[[TARG_INNER_ELEM]] +// CHECK: %[[NEXTINDEX_INNER]] = add nuw i64 %[[ARRAYINDEX_INNER]], 1 +// CHECK: %[[ISDONE_INNER:.+]] = icmp eq i64 %[[NEXTINDEX_INNER]], 1 +// CHECK: br i1 %[[ISDONE_INNER]], label %[[ARRAYINITEND]], label %[[ARRAYINITBODY_INNER]] + +// Check Inner loop 'end' +// CHECK: [[ARRAYINITEND]]: +// CHECK: %[[NEXTINDEX]] = add nuw i64 %[[ARRAYINDEX]], 1 +// CHECK: %[[ISDONE:.+]] = icmp eq i64 %[[NEXTINDEX]], 2 +// CHECK: br i1 %[[ISDONE]], label %{{.*}}, label %[[ARRAYINITBODY]] From 9d7f9aed23fa3e12be50525ea2d1b76d732fd4ed Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Thu, 17 Sep 2020 07:48:46 -0700 Subject: [PATCH 14/18] Finish integration header test --- clang/test/CodeGenSYCL/integration_header.cpp | 3 +- .../CodeGenSYCL/kernel-param-pod-array-ih.cpp | 30 ++++--------------- 2 files changed, 6 insertions(+), 27 deletions(-) diff --git a/clang/test/CodeGenSYCL/integration_header.cpp b/clang/test/CodeGenSYCL/integration_header.cpp index d5eca9624f3f2..5bc45080d4235 100644 --- a/clang/test/CodeGenSYCL/integration_header.cpp +++ b/clang/test/CodeGenSYCL/integration_header.cpp @@ -28,8 +28,7 @@ // CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = { // CHECK-NEXT: //--- _ZTSZ4mainE12first_kernel // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 1, 4 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 8 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 8, 4 }, // CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 12 }, // CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 6112, 24 }, // CHECK-NEXT: { kernel_param_kind_t::kind_sampler, 8, 40 }, diff --git a/clang/test/CodeGenSYCL/kernel-param-pod-array-ih.cpp b/clang/test/CodeGenSYCL/kernel-param-pod-array-ih.cpp index ad2b36524a78e..c5881940a3fa3 100644 --- a/clang/test/CodeGenSYCL/kernel-param-pod-array-ih.cpp +++ b/clang/test/CodeGenSYCL/kernel-param-pod-array-ih.cpp @@ -21,41 +21,21 @@ // CHECK: static constexpr // CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = { // CHECK-NEXT: //--- _ZTSZ4mainE8kernel_B -// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 4 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 8 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 12 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 16 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 20, 0 }, // CHECK-EMPTY: // CHECK-NEXT: //--- _ZTSZ4mainE8kernel_C -// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 4 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 8 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 12 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 16 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 20 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 24, 0 }, // CHECK-EMPTY: // CHECK-NEXT: //--- _ZTSZ4mainE8kernel_D -// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 4 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 8 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 12 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 16 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 20 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 24 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 28 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 32 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 36 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 40 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 44 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 48, 0 }, // CHECK-EMPTY: // CHECK-NEXT: }; // CHECK: static constexpr // CHECK-NEXT: const unsigned kernel_signature_start[] = { // CHECK-NEXT: 0, // _ZTSZ4mainE8kernel_B -// CHECK-NEXT: 6, // _ZTSZ4mainE8kernel_C -// CHECK-NEXT: 13 // _ZTSZ4mainE8kernel_D +// CHECK-NEXT: 2, // _ZTSZ4mainE8kernel_C +// CHECK-NEXT: 4 // _ZTSZ4mainE8kernel_D // CHECK-NEXT: }; // CHECK: template <> struct KernelInfo { From 6832565c6c6f416cc082fbbc1a053b6825529f78 Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Thu, 17 Sep 2020 09:19:34 -0700 Subject: [PATCH 15/18] Add test for decomposition situations --- clang/test/SemaSYCL/decomposition.cpp | 124 ++++++++++++++++++++++++++ 1 file changed, 124 insertions(+) create mode 100644 clang/test/SemaSYCL/decomposition.cpp diff --git a/clang/test/SemaSYCL/decomposition.cpp b/clang/test/SemaSYCL/decomposition.cpp new file mode 100644 index 0000000000000..6830c944f239c --- /dev/null +++ b/clang/test/SemaSYCL/decomposition.cpp @@ -0,0 +1,124 @@ +// RUN: %clang_cc1 -fsycl -fsycl-is-device -ast-dump %s | FileCheck %s + +#include "Inputs/sycl.hpp" + +using namespace cl::sycl; + +struct has_acc { + accessor acc; +}; + +struct acc_base : accessor { + int i; +}; + +struct has_sampler { + sampler sampl; +}; + +struct has_spec_const { + ONEAPI::experimental::spec_constant SC; +}; + +handler H; + +struct has_stream { + stream s1{0, 0, H}; +}; + +struct has_half { + half h; +}; + +struct non_decomposed { + int i; + float f; + double d; +}; + +struct use_non_decomposed : non_decomposed { + non_decomposed member; + float f; + double d; +}; + +template +struct Test1 { + T a; + T b[2]; + non_decomposed d; + int i; +}; + +template +struct Test2 : T { + non_decomposed d; + int i; +}; + +template +__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { + kernelFunc(); +} + +int main() { + + non_decomposed d; + non_decomposed ds[5]; + use_non_decomposed d2; + use_non_decomposed d2s[5]; + // Check to ensure that these are not decomposed. + kernel([=]() { return d.i + ds[0].i + d2.i + d2s[0].i; }); + // CHECK: FunctionDecl {{.*}}NonDecomp{{.*}} 'void (non_decomposed, __wrapper_class, use_non_decomposed, __wrapper_class)' + + { + Test1 t1; + kernel([=]() { return t1.i; }); + // CHECK: FunctionDecl {{.*}}Acc1{{.*}} 'void (__global char *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, __global char *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, __global char *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, non_decomposed, int)' + Test2 t2; + kernel([=]() { return t2.i; }); + // CHECK: FunctionDecl {{.*}}Acc2{{.*}} 'void (__global char *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, non_decomposed, int)' + Test1 t3; + kernel([=]() { return t3.i; }); + // CHECK: FunctionDecl {{.*}}Acc3{{.*}} 'void (__global char *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, int, __global char *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, int, __global char *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, int, non_decomposed, int)' + Test2 t4; + kernel([=]() { return t4.i; }); + // CHECK: FunctionDecl {{.*}}Acc4{{.*}} 'void (__global char *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, int, non_decomposed, int)' + } + + { + Test1 t1; + kernel([=]() { return t1.i; }); + // CHECK: FunctionDecl {{.*}}Sampl1{{.*}} 'void (sampler_t, sampler_t, sampler_t, non_decomposed, int)' + Test2 t2; + kernel([=]() { return t2.i; }); + // CHECK: FunctionDecl {{.*}}Sampl2{{.*}} 'void (sampler_t, non_decomposed, int)' + } + + { + Test1 t1; + kernel([=]() { return t1.i; }); + // CHECK: FunctionDecl {{.*}}SpecConst{{.*}} 'void (non_decomposed, int)' + Test2 t2; + kernel([=]() { return t2.i; }); + // CHECK: FunctionDecl {{.*}}SpecConst2{{.*}} 'void (non_decomposed, int)' + } + + { + Test1 t1; + kernel([=]() { return t1.i; }); + // CHECK: FunctionDecl {{.*}}Stream1{{.*}} 'void (cl::sycl::stream, __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, cl::sycl::stream, __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, cl::sycl::stream, __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, non_decomposed, int)' + Test2 t2; + kernel([=]() { return t2.i; }); + // CHECK: FunctionDecl {{.*}}Stream2{{.*}} 'void (cl::sycl::stream, __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, non_decomposed, int)' + } + + { + Test1 t1; + kernel([=]() { return t1.i; }); + // CHECK: FunctionDecl {{.*}}Half1{{.*}} 'void (cl::sycl::half, cl::sycl::half, cl::sycl::half, non_decomposed, int)' + Test2 t2; + kernel([=]() { return t2.i; }); + // CHECK: FunctionDecl {{.*}}Half2{{.*}} 'void (cl::sycl::half, non_decomposed, int)' + } +} From 0107920a30550ad17797e87f4d38a71fa0d8ae3f Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Fri, 18 Sep 2020 08:34:26 -0700 Subject: [PATCH 16/18] Fix check lines for windows, removing the quote allows the windows unsigned long long to work here --- clang/test/SemaSYCL/array-kernel-param.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/clang/test/SemaSYCL/array-kernel-param.cpp b/clang/test/SemaSYCL/array-kernel-param.cpp index aa0d69288292d..f2393c63c293a 100644 --- a/clang/test/SemaSYCL/array-kernel-param.cpp +++ b/clang/test/SemaSYCL/array-kernel-param.cpp @@ -324,7 +324,7 @@ int main() { // CHECK-NEXT: OpaqueValueExpr {{.*}} 'int [2][3]' lvalue // CHECK-NEXT: MemberExpr {{.*}} 'int [2][3]' lvalue . // CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_' '__wrapper_class' -// CHECK-NEXT: ArrayInitIndexExpr {{.*}} 'unsigned long' +// CHECK-NEXT: ArrayInitIndexExpr {{.*}} 'unsigned long // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' // CHECK-NEXT: ArraySubscriptExpr {{.*}} 'int' lvalue // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' @@ -334,5 +334,5 @@ int main() { // CHECK-NEXT: OpaqueValueExpr {{.*}} 'int [2][3]' lvalue // CHECK-NEXT: MemberExpr {{.*}} 'int [2][3]' lvalue . // CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_' '__wrapper_class' -// CHECK-NEXT: ArrayInitIndexExpr {{.*}} 'unsigned long' -// CHECK-NEXT: ArrayInitIndexExpr {{.*}} 'unsigned long' +// CHECK-NEXT: ArrayInitIndexExpr {{.*}} 'unsigned long +// CHECK-NEXT: ArrayInitIndexExpr {{.*}} 'unsigned long From 8e8519c5600c15e466ebc1a1ff22ef6430689f36 Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Tue, 22 Sep 2020 06:24:50 -0700 Subject: [PATCH 17/18] Fix things @elizabethandrews came up with --- clang/lib/Sema/SemaSYCL.cpp | 30 ++++++++++++++++-------------- 1 file changed, 16 insertions(+), 14 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index a9c0b03a15a3e..578586197f761 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1080,22 +1080,17 @@ class SyclKernelFieldHandlerBase { // handleStructType, enterStruct, leaveStruct, and visiting of sub-elements. virtual bool handleNonDecompStruct(const CXXRecordDecl *, FieldDecl *, QualType) { - // TODO: Should this assert? Empty would have to do something about it. return true; } virtual bool handleNonDecompStruct(const CXXRecordDecl *, const CXXBaseSpecifier &, QualType) { - // TODO: Should this assert? Empty would have to do something about it. return true; } // Instead of handleArrayType, enterArray, leaveArray, and nextElement (plus // descending down the elements), this function gets called in the event of an // array containing simple elements (even in the case of an MD array). - virtual bool handleSimpleArrayType(FieldDecl *, QualType) { - // TODO: Should this assert? Empty would have to do something about it. - return true; - } + virtual bool handleSimpleArrayType(FieldDecl *, QualType) { return true; } // The following are only used for keeping track of where we are in the base // class/field graph. Int Headers use this to calculate offset, most others @@ -1210,8 +1205,16 @@ void KernelObjVisitor::visitRecord(const CXXRecordDecl *Owner, ParentTy &Parent, QualType RecordTy, HandlerTys &... Handlers) { if (RecordTy->getAsRecordDecl()->hasAttr()) { + // If this container requires decomposition, we have to visit it as + // 'complex', so all handlers are called in this case with the 'complex' + // case. visitComplexRecord(Owner, Parent, Wrapper, RecordTy, Handlers...); } else { + // "Simple" Containers are those that do NOT need to be decomposed, + // "Complex" containers are those that DO. In the case where the container + // does NOT need to be decomposed, we can call VisitSimpleRecord on the + // handlers that have opted-out of VisitInsideSimpleContainers. The 'if' + // makes sure we only do that if at least 1 has opted out. if (!AllTrue::Value) visitSimpleRecord( Owner, Parent, Wrapper, RecordTy, @@ -1219,6 +1222,9 @@ void KernelObjVisitor::visitRecord(const CXXRecordDecl *Owner, ParentTy &Parent, Handlers) .Handler...); + // Even though this is a 'simple' container, some handlers (via + // VisitInsideSimpleContainers = true) need to treat it as if it needs + // decomposing, so we call VisitComplexRecord iif at least one has. if (AnyTrue::Value) visitComplexRecord( Owner, Parent, Wrapper, RecordTy, @@ -1877,8 +1883,8 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { bool handleSimpleArrayType(FieldDecl *FD, QualType FieldTy) final { // Arrays are always wrapped in a struct since they cannot be passed // directly. - RecordDecl *WrappedPointer = wrapField(FD, FieldTy); - QualType ModTy = SemaRef.getASTContext().getRecordType(WrappedPointer); + RecordDecl *WrappedArray = wrapField(FD, FieldTy); + QualType ModTy = SemaRef.getASTContext().getRecordType(WrappedArray); addParam(FD, ModTy); return true; } @@ -2162,14 +2168,10 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { Expr *DRE = SemaRef.BuildDeclRefExpr(KernelParameter, ParamType, VK_LValue, KernelCallerSrcLoc); - // Unwrapp the array. + // Unwrap the array. CXXRecordDecl *WrapperStruct = ParamType->getAsCXXRecordDecl(); FieldDecl *ArrayField = *(WrapperStruct->field_begin()); - DRE = buildMemberExpr(DRE, ArrayField); - - // TODO: do we need to do the L->R val conversion? I think this should - // happen automatically. - return DRE; + return buildMemberExpr(DRE, ArrayField); } // Returns 'true' if the thing we're visiting (Based on the FD/QualType pair) From f70a58a0bdca48510a1656f5368ce9504809e0ff Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Wed, 23 Sep 2020 06:10:00 -0700 Subject: [PATCH 18/18] Add test, fix nits from @elizabethandrews --- clang/lib/Sema/SemaSYCL.cpp | 5 +-- clang/test/SemaSYCL/array-kernel-param.cpp | 36 +++++++++++++++++++++- 2 files changed, 38 insertions(+), 3 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 578586197f761..9fb35412cac4a 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1035,7 +1035,7 @@ class SyclKernelFieldHandlerBase { public: static constexpr const bool VisitUnionBody = false; static constexpr const bool VisitNthArrayElement = true; - // Opt-in based on whether we should visit inside Simple containers (structs, + // Opt-in based on whether we should visit inside simple containers (structs, // arrays). All of the 'check' types should likely be true, the int-header, // and kernel decl creation types should not. static constexpr const bool VisitInsideSimpleContainers = true; @@ -2734,7 +2734,8 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { } bool handleSimpleArrayType(FieldDecl *FD, QualType FieldTy) final { - // I think this is right, we need to always wrap arrays. + // Arrays are always wrapped inside of structs, so just treat it as a simple + // struct. addParam(FD, FieldTy, SYCLIntegrationHeader::kind_std_layout); return true; } diff --git a/clang/test/SemaSYCL/array-kernel-param.cpp b/clang/test/SemaSYCL/array-kernel-param.cpp index f2393c63c293a..ea1d4ac4a899d 100644 --- a/clang/test/SemaSYCL/array-kernel-param.cpp +++ b/clang/test/SemaSYCL/array-kernel-param.cpp @@ -44,7 +44,15 @@ int main() { int foo_c; }; + // Not decomposed. + struct foo2 { + int foo_a; + int foo_2D[2][1]; + int foo_c; + }; + foo struct_array[2]; + foo2 struct_array2[2]; int array_2D[2][3]; @@ -82,6 +90,11 @@ int main() { [=]() { int local = array_2D[1][1]; }); + + a_kernel( + [=]() { + foo2 local = struct_array2[0]; + }); } // Check kernel_A parameters @@ -245,7 +258,7 @@ int main() { // CHECK-NEXT: ImplicitCastExpr // CHECK-NEXT: DeclRefExpr{{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_c' 'int' -// Initializer for first element of struct_array +// Initializer for second element of struct_array // CHECK-NEXT: InitListExpr {{.*}} 'foo' // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_foo_a' 'int' @@ -336,3 +349,24 @@ int main() { // CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_' '__wrapper_class' // CHECK-NEXT: ArrayInitIndexExpr {{.*}} 'unsigned long // CHECK-NEXT: ArrayInitIndexExpr {{.*}} 'unsigned long + +// Check kernel_G parameters. +// CHECK: FunctionDecl {{.*}}kernel_G{{.*}} 'void (__wrapper_class)' +// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ '__wrapper_class' +// Check kernel_G inits +// CHECK-NEXT: CompoundStmt +// CHECK-NEXT: DeclStmt +// CHECK-NEXT: VarDecl {{.*}} cinit +// CHECK-NEXT: InitListExpr +// CHECK-NEXT: ArrayInitLoopExpr {{.*}} 'foo2 [2]' +// CHECK-NEXT: OpaqueValueExpr {{.*}} 'foo2 [2]' lvalue +// CHECK-NEXT: MemberExpr {{.*}} 'foo2 [2]' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_' '__wrapper_class' +// CHECK-NEXT: CXXConstructExpr {{.*}} 'foo2' 'void (const foo2 &) noexcept' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const foo2' lvalue +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'foo2' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'foo2 *' +// CHECK-NEXT: OpaqueValueExpr {{.*}} 'foo2 [2]' lvalue +// CHECK-NEXT: MemberExpr {{.*}} 'foo2 [2]' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_' '__wrapper_class' +// CHECK-NEXT: ArrayInitIndexExpr {{.*}} 'unsigned long