From 5637abbd2b76a6c66195415dd2547fae0280bd89 Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Fri, 29 Jul 2022 09:11:12 -0700 Subject: [PATCH 01/18] Mark structs which contain pointers. Modify DeclCreator Signed-off-by: Elizabeth Andrews --- clang/include/clang/Basic/Attr.td | 8 +++++ clang/lib/Sema/SemaSYCL.cpp | 58 +++++++++++++++++++++++++++---- 2 files changed, 60 insertions(+), 6 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 5bcbb34b7b93e..4fb9e2d05328c 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1382,6 +1382,14 @@ def SYCLRequiresDecomposition : InheritableAttr { let Documentation = [Undocumented]; } +def SYCLGenerateNewType : 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 98f90d8c602da..0f767ea59d2a3 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1472,7 +1472,7 @@ void KernelObjVisitor::visitRecord(const CXXRecordDecl *Owner, ParentTy &Parent, HandlerTys &... Handlers) { RecordDecl *RD = RecordTy->getAsRecordDecl(); assert(RD && "should not be null."); - if (RD->hasAttr()) { + if (RD->hasAttr() || RD->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. @@ -1740,9 +1740,13 @@ class SyclKernelUnionChecker : public SyclKernelFieldHandler { } }; -// A type to mark whether a collection requires decomposition. +// A type to mark whether a collection requires decomposition +// or needs to be transformed to a new type. If a collection +// contains pointers, and is not decomposed, a new type must +// be generated with all pointers in global address space. class SyclKernelDecompMarker : public SyclKernelFieldHandler { llvm::SmallVector CollectionStack; + llvm::SmallVector PointerStack; public: static constexpr const bool VisitUnionBody = false; @@ -1752,6 +1756,7 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { // In order to prevent checking this over and over, just add a dummy-base // entry. CollectionStack.push_back(true); + PointerStack.push_back(true); } bool handleSyclSpecialType(const CXXRecordDecl *, const CXXBaseSpecifier &, @@ -1770,12 +1775,13 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { } bool handlePointerType(FieldDecl *, QualType) final { - CollectionStack.back() = true; + PointerStack.back() = true; return true; } bool enterStruct(const CXXRecordDecl *, FieldDecl *, QualType) final { CollectionStack.push_back(false); + PointerStack.push_back(false); return true; } @@ -1787,6 +1793,14 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { RD->addAttr(SYCLRequiresDecompositionAttr::CreateImplicit( SemaRef.getASTContext())); CollectionStack.back() = true; + PointerStack.pop_back(); + } else if (PointerStack.pop_back_val()) { + RecordDecl *RD = Ty->getAsRecordDecl(); + assert(RD && "should not be null."); + if (!RD->hasAttr()) + RD->addAttr( + SYCLGenerateNewTypeAttr::CreateImplicit(SemaRef.getASTContext())); + PointerStack.back() = true; } return true; } @@ -1794,6 +1808,7 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { bool enterStruct(const CXXRecordDecl *, const CXXBaseSpecifier &, QualType) final { CollectionStack.push_back(false); + PointerStack.push_back(false); return true; } @@ -1806,8 +1821,15 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { RD->addAttr(SYCLRequiresDecompositionAttr::CreateImplicit( SemaRef.getASTContext())); CollectionStack.back() = true; + PointerStack.pop_back(); + } else if (PointerStack.pop_back_val()) { + RecordDecl *RD = Ty->getAsRecordDecl(); + assert(RD && "should not be null."); + if (!RD->hasAttr()) + RD->addAttr( + SYCLGenerateNewTypeAttr::CreateImplicit(SemaRef.getASTContext())); + PointerStack.back() = true; } - return true; } @@ -2051,7 +2073,18 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { return true; } - bool leaveStruct(const CXXRecordDecl *, FieldDecl *, QualType) final { + bool leaveStruct(const CXXRecordDecl *RD, FieldDecl *, QualType) final { + if (RD->hasAttr()) { + auto NumFields = std::distance(RD->field_begin(), RD->field_end()); + for (unsigned I = 0; I < NumFields; I++) { + // Pop out from Params + auto PVD = Params.pop_back_val(); + // Remember this won't match order of fields in original class. So should store this is right order in temporary vector + // How do I create a new type here? Look at ASTContext::buildImplicitRecord, createPrivatesRecordDecl, addFieldToRecordDecl() + + } + } + --StructDepth; return true; } @@ -2062,8 +2095,21 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { return true; } - bool leaveStruct(const CXXRecordDecl *, const CXXBaseSpecifier &BS, + bool leaveStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS, QualType FieldTy) final { + if (RD->hasAttr()) { + auto NumFields = std::distance(RD->field_begin(), RD->field_end()); + for (unsigned I = 0; I < NumFields; I++) { + // Pop out from Params + auto PVD = Params.pop_back_val(); + auto test = PVD; + // Remember this won't match order of fields in original class. So should store this is right order in temporary vector + // Can I pop it out in right order by starting iteration at Params.size() - NumOfFields? + // How do I create a new global type here? - Look at how we create a kernel object + + } + } + --StructDepth; return true; } From 235a2b8ef226067584ca642d7253b1d177a64072 Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Fri, 29 Jul 2022 09:14:14 -0700 Subject: [PATCH 02/18] Remove DeclCreator changes Signed-off-by: Elizabeth Andrews --- clang/lib/Sema/SemaSYCL.cpp | 28 ++-------------------------- 1 file changed, 2 insertions(+), 26 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 0f767ea59d2a3..05639b19521b6 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -2073,18 +2073,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { return true; } - bool leaveStruct(const CXXRecordDecl *RD, FieldDecl *, QualType) final { - if (RD->hasAttr()) { - auto NumFields = std::distance(RD->field_begin(), RD->field_end()); - for (unsigned I = 0; I < NumFields; I++) { - // Pop out from Params - auto PVD = Params.pop_back_val(); - // Remember this won't match order of fields in original class. So should store this is right order in temporary vector - // How do I create a new type here? Look at ASTContext::buildImplicitRecord, createPrivatesRecordDecl, addFieldToRecordDecl() - - } - } - + bool leaveStruct(const CXXRecordDecl *, FieldDecl *, QualType) final { --StructDepth; return true; } @@ -2095,21 +2084,8 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { return true; } - bool leaveStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS, + bool leaveStruct(const CXXRecordDecl *, const CXXBaseSpecifier &BS, QualType FieldTy) final { - if (RD->hasAttr()) { - auto NumFields = std::distance(RD->field_begin(), RD->field_end()); - for (unsigned I = 0; I < NumFields; I++) { - // Pop out from Params - auto PVD = Params.pop_back_val(); - auto test = PVD; - // Remember this won't match order of fields in original class. So should store this is right order in temporary vector - // Can I pop it out in right order by starting iteration at Params.size() - NumOfFields? - // How do I create a new global type here? - Look at how we create a kernel object - - } - } - --StructDepth; return true; } From 88905fe98b3d40e09512168ee2352e4b7d96115f Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Wed, 7 Sep 2022 20:22:35 -0700 Subject: [PATCH 03/18] Broken first attempt. Missing BodyCreator Signed-off-by: Elizabeth Andrews --- clang/lib/Sema/SemaSYCL.cpp | 226 +++++++++++++++++++++++++++++++++++- 1 file changed, 221 insertions(+), 5 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 05639b19521b6..4eabbb6558d3b 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1323,6 +1323,7 @@ class SyclKernelFieldHandlerBase { // 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; + static constexpr const bool VisitInsideSimpleContainersWithPointer = false; // Mark these virtual so that we can use override in the implementer classes, // despite virtual dispatch never being used. @@ -1472,11 +1473,17 @@ void KernelObjVisitor::visitRecord(const CXXRecordDecl *Owner, ParentTy &Parent, HandlerTys &... Handlers) { RecordDecl *RD = RecordTy->getAsRecordDecl(); assert(RD && "should not be null."); - if (RD->hasAttr() || RD->hasAttr()) { + if (RD->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 if (AnyTrue:: + Value) { + if (RD->hasAttr()) + visitComplexRecord(Owner, Parent, Wrapper, RecordTy, Handlers...); + else + visitSimpleRecord(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 @@ -1833,6 +1840,7 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { return true; } + // Elizabeth - Look into handling arrays bool enterArray(FieldDecl *, QualType ArrayTy, QualType ElementTy) final { CollectionStack.push_back(false); return true; @@ -1851,6 +1859,175 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { } }; +static CXXRecordDecl *createNewType(ASTContext &Ctx, const CXXRecordDecl *RD) { + auto *ModifiedRD = CXXRecordDecl::Create( + Ctx, TTK_Struct, const_cast(RD->getDeclContext()), + SourceLocation(), SourceLocation(), RD->getIdentifier()); + ModifiedRD->startDefinition(); + return ModifiedRD; +} + +class SyclKernelPointerHandler : public SyclKernelFieldHandler { + llvm::SmallVector ModifiedRecords; + SmallVector ModifiedBases; + + void addField(const FieldDecl *FD, QualType FieldTy) { + assert(!ModifiedRecords.empty() && + "ModifiedRecords should have at least 1 record"); + ASTContext &Ctx = SemaRef.getASTContext(); + auto *Field = FieldDecl::Create( + Ctx, ModifiedRecords.back(), SourceLocation(), SourceLocation(), + FD->getIdentifier(), FieldTy, + Ctx.getTrivialTypeSourceInfo(FieldTy, SourceLocation()), /*BW=*/nullptr, + /*Mutable=*/false, ICIS_NoInit); + Field->setAccess(FD->getAccess()); + ModifiedRecords.back()->addDecl(Field); + } + + void createBaseSpecifier(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS) { + TypeSourceInfo *TInfo = SemaRef.getASTContext().getTrivialTypeSourceInfo( + QualType(RD->getTypeForDecl(), 0), SourceLocation()); + CXXBaseSpecifier *ModifiedBase = SemaRef.CheckBaseSpecifier( + const_cast(RD), SourceRange(), BS.isVirtual(), + BS.getAccessSpecifier(), TInfo, SourceLocation()); + ModifiedBases.push_back(ModifiedBase); + + } + +public: + static constexpr const bool VisitInsideSimpleContainersWithPointer = true; + SyclKernelPointerHandler(Sema &S, const CXXRecordDecl *RD) + : SyclKernelFieldHandler(S) { + CXXRecordDecl *ModifiedRD = createNewType(S.getASTContext(), RD); + ModifiedRecords.push_back(ModifiedRD); + } + + bool enterStruct(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final { + // Create the new record type. The fields (and base classes) of this + // record will be generated as the visitor traverses kernel object + // record fields. + // + // Elizabeth - Move this out to separate function + auto *RD = Ty->getAsCXXRecordDecl(); + auto *ModifiedRD = CXXRecordDecl::Create( + SemaRef.getASTContext(), TTK_Struct, + const_cast(RD->getDeclContext()), SourceLocation(), + SourceLocation(), RD->getIdentifier()); + ModifiedRD->startDefinition(); + ModifiedRecords.push_back(ModifiedRD); + return true; + } + + bool leaveStruct(const CXXRecordDecl *RD, FieldDecl *FD, QualType Ty) final { + // At this point, the new type has been created. Add this record as a field + // of it's parent record. + CXXRecordDecl *ModifiedRD = ModifiedRecords.pop_back_val(); + ModifiedRD->completeDefinition(); + if (!ModifiedRecords.empty()) + addField(FD, QualType(ModifiedRD->getTypeForDecl(), 0)); + return true; + } + + bool enterStruct(const CXXRecordDecl *, const CXXBaseSpecifier &BS, + QualType FieldTy) final { + // Create the new record type. The fields (and base classes) of this + // record will be generated as the visitor traverses kernel object + // record fields. + auto *RD = FieldTy->getAsCXXRecordDecl(); + auto *ModifiedRD = CXXRecordDecl::Create( + SemaRef.getASTContext(), TTK_Struct, + const_cast(RD->getDeclContext()), SourceLocation(), + SourceLocation(), RD->getIdentifier()); + ModifiedRD->startDefinition(); + ModifiedRecords.push_back(ModifiedRD); + return true; + } + + bool leaveStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS, + QualType FieldTy) final { + // Pop out generated class. + CXXRecordDecl *ModifiedRD = ModifiedRecords.pop_back_val(); + ModifiedRD->completeDefinition(); + + // If all bases classes for RD have been created, set the bases. + // Doesn't work for multi-level inheritance' + // Instead of handling parent decl, handle base classes for current struct itself? + //if (RD->getNumBases() == ModifiedBases.size()) { + //ModifiedRecords.back()->setBases(ModifiedBases.data(), RD->getNumBases()); + //ModifiedBases.clear(); + //} + + const auto *OldBaseDecl = FieldTy->getAsCXXRecordDecl(); + + if (OldBaseDecl->getNumBases() > 0) { + SmallVector BasesForGeneratedClass; + for (size_t I = 0; I < RD->getNumBases(); ++I) + BasesForGeneratedClass.push_back(ModifiedBases.pop_back_val()); + ModifiedRD->setBases(BasesForGeneratedClass.data(), RD->getNumBases()); + } + + // Create CXXBaseSpecifier for this generated class. + createBaseSpecifier(ModifiedRD, BS); + return true; + } + + bool handlePointerType(FieldDecl *FD, QualType FieldTy) final { + QualType PointeeTy = FieldTy->getPointeeType(); + Qualifiers Quals = PointeeTy.getQualifiers(); + auto AS = Quals.getAddressSpace(); + // Leave global_device and global_host address spaces as is to help FPGA + // device in memory allocations + if (!PointeeTy->isFunctionType() && AS != LangAS::sycl_global_device && + AS != LangAS::sycl_global_host) + Quals.setAddressSpace(LangAS::sycl_global); + PointeeTy = SemaRef.getASTContext().getQualifiedType( + PointeeTy.getUnqualifiedType(), Quals); + QualType ModTy = SemaRef.getASTContext().getPointerType(PointeeTy); + addField(FD, ModTy); + return true; + // Elizabeth - Do we care about pointer wrapping. We are already one level + // in at this point right? + } + + bool handleScalarType(FieldDecl *FD, QualType FieldTy) final { + addField(FD, FieldTy); + return true; + } + + bool handleUnionType(FieldDecl *FD, QualType FieldTy) final { + return handleScalarType(FD, FieldTy); + } + + bool handleNonDecompStruct(const CXXRecordDecl *RD, FieldDecl *FD, + QualType Ty) final { + addField(FD, Ty); + return true; + } + + bool handleNonDecompStruct(const CXXRecordDecl *Base, + const CXXBaseSpecifier &BS, QualType Ty) final { + // Elizabeth - Fix this + /*createBaseSpecifier(Base, BS); + + if (Base->getNumBases() > 0) { + SmallVector BasesForGeneratedClass; + for (size_t I = 0; I < Base->getNumBases(); ++I) + BasesForGeneratedClass.push_back(ModifiedBases.pop_back_val()); + ModifiedRecords.back()->setBases(BasesForGeneratedClass.data(), Base->getNumBases()); + }*/ + return true; + } + +public: + QualType getNewType() { + CXXRecordDecl *ModifiedRD = ModifiedRecords.pop_back_val(); + ModifiedRD->completeDefinition(); + return QualType(ModifiedRD->getTypeForDecl(), 0); + } + + // Elizabeth - Need to handle KernelHandler, array +}; + // A type to Create and own the FunctionDecl for the kernel. class SyclKernelDeclCreator : public SyclKernelFieldHandler { FunctionDecl *KernelDecl; @@ -2189,15 +2366,30 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { return true; } - bool handleNonDecompStruct(const CXXRecordDecl *, FieldDecl *FD, + bool handleNonDecompStruct(const CXXRecordDecl *RD, FieldDecl *FD, QualType Ty) final { - addParam(FD, Ty); + CXXRecordDecl *FieldRecordDecl = Ty->getAsCXXRecordDecl(); + SyclKernelPointerHandler PointerHandler(SemaRef, FieldRecordDecl); + KernelObjVisitor Visitor{SemaRef}; + Visitor.VisitRecordBases(FieldRecordDecl, PointerHandler); + Visitor.VisitRecordFields(FieldRecordDecl, PointerHandler); + addParam(FD, PointerHandler.getNewType()); return true; } bool handleNonDecompStruct(const CXXRecordDecl *Base, const CXXBaseSpecifier &BS, QualType Ty) final { - addParam(BS, Ty); + // We need to create a list of BaseSpecifiers for each class + // Actions.ActOnBaseSpecifiers(ClassDecl, BaseInfo); + // Or maybe just setBases() which requires us to create a bunch of + // CXXBaseSpecifiers + + CXXRecordDecl *BaseRecordDecl = Ty->getAsCXXRecordDecl(); + SyclKernelPointerHandler PointerHandler(SemaRef, BaseRecordDecl); + KernelObjVisitor Visitor{SemaRef}; + Visitor.VisitRecordBases(BaseRecordDecl, PointerHandler); + Visitor.VisitRecordFields(BaseRecordDecl, PointerHandler); + addParam(BS, PointerHandler.getNewType()); return true; } @@ -2694,6 +2886,25 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { addFieldInit(FD, Ty, ParamRef); } + // Call __builtin_memcpy to copy modified type with pointer + // to local clone + void generateBitwiseCopy(FieldDecl *FD, QualType Ty) { + // Need to change this to bitwise copy + addSimpleFieldInit(FD, Ty); + + // Compute the size of the memory buffer to be copied. + /*QualType SizeType = S.Context.getSizeType(); + llvm::APInt Size(SemaRef.Context.getTypeSize(SizeType), + SemaRef.Context.getTypeSizeInChars(Ty).getQuantity()); + + Expr *ParamRef = createParamReferenceExpr(); + From = UnaryOperator::Create(SemaRef.Context, ParamRef, UO_AddrOf, + SemaRef.Context.getPointerType(ParamRef->getType()), VK_PRValue, + OK_Ordinary, KernelCallerSrcLoc, false, SemaRef.CurFPFeatureOverrides()); + */ + // How do I get a reference to the field of local clone? + } + MemberExpr *buildMemberExpr(Expr *Base, ValueDecl *Member) { DeclAccessPair MemberDAP = DeclAccessPair::make(Member, AS_none); MemberExpr *Result = SemaRef.BuildMemberExpr( @@ -2923,7 +3134,12 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { bool handleNonDecompStruct(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final { - addSimpleFieldInit(FD, Ty); + CXXRecordDecl *RD = Ty->getAsCXXRecordDecl(); + assert(RD && "Type must be a C++ record type"); + if (RD->hasAttr()) + generateBitwiseCopy(FD, Ty); + else + addSimpleFieldInit(FD, Ty); return true; } From 59834111e7e376b9855b1d69dabb20fff7236eb8 Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Thu, 8 Sep 2022 18:46:59 -0700 Subject: [PATCH 04/18] Refactor code and fix base class handling Signed-off-by: Elizabeth Andrews --- clang/lib/Sema/SemaSYCL.cpp | 152 +++++++++++++++++------------------- 1 file changed, 72 insertions(+), 80 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 4eabbb6558d3b..90126a44c3525 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1859,115 +1859,109 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { } }; -static CXXRecordDecl *createNewType(ASTContext &Ctx, const CXXRecordDecl *RD) { - auto *ModifiedRD = CXXRecordDecl::Create( - Ctx, TTK_Struct, const_cast(RD->getDeclContext()), - SourceLocation(), SourceLocation(), RD->getIdentifier()); - ModifiedRD->startDefinition(); - return ModifiedRD; -} - class SyclKernelPointerHandler : public SyclKernelFieldHandler { llvm::SmallVector ModifiedRecords; SmallVector ModifiedBases; + IdentifierInfo *getModifiedName(IdentifierInfo *Id) { + std::string Name = (Id->getName() + Twine("_generated")).str(); + return &SemaRef.getASTContext().Idents.get(Name); + } + + // Create Decl for the new type we are generating. + // The fields (and base classes) of this record will be generated as + // the visitor traverses kernel object record fields. + void createNewType(const CXXRecordDecl *RD) { + auto *ModifiedRD = CXXRecordDecl::Create( + SemaRef.getASTContext(), TTK_Struct, + const_cast(RD->getDeclContext()), SourceLocation(), + SourceLocation(), getModifiedName(RD->getIdentifier())); + ModifiedRD->startDefinition(); + ModifiedRecords.push_back(ModifiedRD); + } + + // Create and add FieldDecl for FieldTy to generated record. void addField(const FieldDecl *FD, QualType FieldTy) { assert(!ModifiedRecords.empty() && "ModifiedRecords should have at least 1 record"); ASTContext &Ctx = SemaRef.getASTContext(); auto *Field = FieldDecl::Create( Ctx, ModifiedRecords.back(), SourceLocation(), SourceLocation(), - FD->getIdentifier(), FieldTy, + getModifiedName(FD->getIdentifier()), FieldTy, Ctx.getTrivialTypeSourceInfo(FieldTy, SourceLocation()), /*BW=*/nullptr, /*Mutable=*/false, ICIS_NoInit); Field->setAccess(FD->getAccess()); + // Add generated field to generated record. ModifiedRecords.back()->addDecl(Field); } - void createBaseSpecifier(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS) { + void createBaseSpecifier(const CXXRecordDecl *Parent, const CXXRecordDecl *RD, + const CXXBaseSpecifier &BS) { TypeSourceInfo *TInfo = SemaRef.getASTContext().getTrivialTypeSourceInfo( QualType(RD->getTypeForDecl(), 0), SourceLocation()); CXXBaseSpecifier *ModifiedBase = SemaRef.CheckBaseSpecifier( - const_cast(RD), SourceRange(), BS.isVirtual(), + const_cast(Parent), SourceRange(), BS.isVirtual(), BS.getAccessSpecifier(), TInfo, SourceLocation()); ModifiedBases.push_back(ModifiedBase); + } + CXXRecordDecl *getGeneratedNewRecord(const CXXRecordDecl *OldBaseDecl) { + // At this point we have finished generating fields for the new + // class corresponding to OldBaseDecl. Pop out the generated + // record. + CXXRecordDecl *ModifiedRD = ModifiedRecords.pop_back_val(); + ModifiedRD->completeDefinition(); + // Check the 'old' class for base classes. + // Set bases classes for newly generated class if it has any. + if (OldBaseDecl->getNumBases() > 0) { + SmallVector BasesForGeneratedClass; + for (size_t I = 0; I < OldBaseDecl->getNumBases(); ++I) + BasesForGeneratedClass.insert(BasesForGeneratedClass.begin(), + ModifiedBases.pop_back_val()); + ModifiedRD->setBases(BasesForGeneratedClass.data(), + OldBaseDecl->getNumBases()); + } + return ModifiedRD; } public: static constexpr const bool VisitInsideSimpleContainersWithPointer = true; SyclKernelPointerHandler(Sema &S, const CXXRecordDecl *RD) : SyclKernelFieldHandler(S) { - CXXRecordDecl *ModifiedRD = createNewType(S.getASTContext(), RD); - ModifiedRecords.push_back(ModifiedRD); + // Generate new type + createNewType(RD); } - bool enterStruct(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final { - // Create the new record type. The fields (and base classes) of this - // record will be generated as the visitor traverses kernel object - // record fields. - // - // Elizabeth - Move this out to separate function - auto *RD = Ty->getAsCXXRecordDecl(); - auto *ModifiedRD = CXXRecordDecl::Create( - SemaRef.getASTContext(), TTK_Struct, - const_cast(RD->getDeclContext()), SourceLocation(), - SourceLocation(), RD->getIdentifier()); - ModifiedRD->startDefinition(); - ModifiedRecords.push_back(ModifiedRD); + bool enterStruct(const CXXRecordDecl *, FieldDecl *, QualType Ty) final { + // Generate new type + createNewType(Ty->getAsCXXRecordDecl()); return true; } - bool leaveStruct(const CXXRecordDecl *RD, FieldDecl *FD, QualType Ty) final { - // At this point, the new type has been created. Add this record as a field - // of it's parent record. - CXXRecordDecl *ModifiedRD = ModifiedRecords.pop_back_val(); - ModifiedRD->completeDefinition(); + bool leaveStruct(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final { + // Get Decl of generated new type + CXXRecordDecl *ModifiedRD = getGeneratedNewRecord(Ty->getAsCXXRecordDecl()); + + // Add this record as a field of it's parent record. if (!ModifiedRecords.empty()) addField(FD, QualType(ModifiedRD->getTypeForDecl(), 0)); return true; } - bool enterStruct(const CXXRecordDecl *, const CXXBaseSpecifier &BS, - QualType FieldTy) final { - // Create the new record type. The fields (and base classes) of this - // record will be generated as the visitor traverses kernel object - // record fields. - auto *RD = FieldTy->getAsCXXRecordDecl(); - auto *ModifiedRD = CXXRecordDecl::Create( - SemaRef.getASTContext(), TTK_Struct, - const_cast(RD->getDeclContext()), SourceLocation(), - SourceLocation(), RD->getIdentifier()); - ModifiedRD->startDefinition(); - ModifiedRecords.push_back(ModifiedRD); + bool enterStruct(const CXXRecordDecl *, const CXXBaseSpecifier &, + QualType Ty) final { + // Generate new type + createNewType(Ty->getAsCXXRecordDecl()); return true; } - bool leaveStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS, - QualType FieldTy) final { - // Pop out generated class. - CXXRecordDecl *ModifiedRD = ModifiedRecords.pop_back_val(); - ModifiedRD->completeDefinition(); - - // If all bases classes for RD have been created, set the bases. - // Doesn't work for multi-level inheritance' - // Instead of handling parent decl, handle base classes for current struct itself? - //if (RD->getNumBases() == ModifiedBases.size()) { - //ModifiedRecords.back()->setBases(ModifiedBases.data(), RD->getNumBases()); - //ModifiedBases.clear(); - //} - - const auto *OldBaseDecl = FieldTy->getAsCXXRecordDecl(); - - if (OldBaseDecl->getNumBases() > 0) { - SmallVector BasesForGeneratedClass; - for (size_t I = 0; I < RD->getNumBases(); ++I) - BasesForGeneratedClass.push_back(ModifiedBases.pop_back_val()); - ModifiedRD->setBases(BasesForGeneratedClass.data(), RD->getNumBases()); - } + bool leaveStruct(const CXXRecordDecl *Parent, const CXXBaseSpecifier &BS, + QualType Ty) final { + // Get Decl of generated new type + CXXRecordDecl *ModifiedRD = getGeneratedNewRecord(Ty->getAsCXXRecordDecl()); // Create CXXBaseSpecifier for this generated class. - createBaseSpecifier(ModifiedRD, BS); + createBaseSpecifier(Parent, ModifiedRD, BS); return true; } @@ -1985,8 +1979,8 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { QualType ModTy = SemaRef.getASTContext().getPointerType(PointeeTy); addField(FD, ModTy); return true; - // Elizabeth - Do we care about pointer wrapping. We are already one level - // in at this point right? + // We do not need to wrap pointers since this is a pointer inside + // non-decomposed struct. } bool handleScalarType(FieldDecl *FD, QualType FieldTy) final { @@ -1998,23 +1992,15 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { return handleScalarType(FD, FieldTy); } - bool handleNonDecompStruct(const CXXRecordDecl *RD, FieldDecl *FD, - QualType Ty) final { + bool handleNonDecompStruct(const CXXRecordDecl *, FieldDecl *FD, + QualType Ty) final { addField(FD, Ty); return true; } - bool handleNonDecompStruct(const CXXRecordDecl *Base, + bool handleNonDecompStruct(const CXXRecordDecl *Parent, const CXXBaseSpecifier &BS, QualType Ty) final { - // Elizabeth - Fix this - /*createBaseSpecifier(Base, BS); - - if (Base->getNumBases() > 0) { - SmallVector BasesForGeneratedClass; - for (size_t I = 0; I < Base->getNumBases(); ++I) - BasesForGeneratedClass.push_back(ModifiedBases.pop_back_val()); - ModifiedRecords.back()->setBases(BasesForGeneratedClass.data(), Base->getNumBases()); - }*/ + createBaseSpecifier(Parent, Ty->getAsCXXRecordDecl(), BS); return true; } @@ -2022,6 +2008,12 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { QualType getNewType() { CXXRecordDecl *ModifiedRD = ModifiedRecords.pop_back_val(); ModifiedRD->completeDefinition(); + + // Elizabeth - Check if this will work if kernel functor has generated + // base class. + if (!ModifiedBases.empty()) + ModifiedRD->setBases(ModifiedBases.data(), ModifiedBases.size()); + return QualType(ModifiedRD->getTypeForDecl(), 0); } From f12a2d1a3ef7b7be189a7052e6863bb0698756c7 Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Sun, 11 Sep 2022 13:29:57 -0700 Subject: [PATCH 05/18] Add BodyCreator Statements (incomplete) Signed-off-by: Elizabeth Andrews --- clang/include/clang/Sema/Sema.h | 3 + clang/lib/Sema/SemaDeclCXX.cpp | 98 +++++++++++++++++---------------- clang/lib/Sema/SemaSYCL.cpp | 34 +++++++----- 3 files changed, 73 insertions(+), 62 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 6ba624c6b7c2b..b429f8fbf4a06 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -3098,6 +3098,9 @@ class Sema final { void SetParamDefaultArgument(ParmVarDecl *Param, Expr *DefaultArg, SourceLocation EqualLoc); + StmtResult BuildMemCpyCall(SourceLocation Loc, QualType T, Expr *From, + Expr *To); + // Contexts where using non-trivial C union types can be disallowed. This is // passed to err_non_trivial_c_union_in_invalid_context. enum NonTrivialCUnionContext { diff --git a/clang/lib/Sema/SemaDeclCXX.cpp b/clang/lib/Sema/SemaDeclCXX.cpp index 20ea8ccfce50b..046083ab6fcea 100644 --- a/clang/lib/Sema/SemaDeclCXX.cpp +++ b/clang/lib/Sema/SemaDeclCXX.cpp @@ -902,6 +902,55 @@ Sema::ActOnDecompositionDeclarator(Scope *S, Declarator &D, return New; } +StmtResult Sema::BuildMemCpyCall(SourceLocation Loc, QualType T, Expr *From, + Expr *To) { + // Compute the size of the memory buffer to be copied. + QualType SizeType = Context.getSizeType(); + llvm::APInt Size(Context.getTypeSize(SizeType), + Context.getTypeSizeInChars(T).getQuantity()); + + // Take the address of the field references for "from" and "to". We + // directly construct UnaryOperators here because semantic analysis + // does not permit us to take the address of an xvalue. + From = UnaryOperator::Create( + Context, From, UO_AddrOf, Context.getPointerType(From->getType()), + VK_PRValue, OK_Ordinary, Loc, false, CurFPFeatureOverrides()); + To = UnaryOperator::Create(Context, To, UO_AddrOf, + Context.getPointerType(To->getType()), VK_PRValue, + OK_Ordinary, Loc, false, CurFPFeatureOverrides()); + + const Type *E = T->getBaseElementTypeUnsafe(); + bool NeedsCollectableMemCpy = + !getLangOpts().SYCLIsDevice && E->isRecordType() && + E->castAs()->getDecl()->hasObjectMember(); + + // Create a reference to the __builtin_objc_memmove_collectable function + StringRef MemCpyName = NeedsCollectableMemCpy + ? "__builtin_objc_memmove_collectable" + : "__builtin_memcpy"; + LookupResult R(*this, &Context.Idents.get(MemCpyName), Loc, + Sema::LookupOrdinaryName); + LookupName(R, TUScope, true); + + FunctionDecl *MemCpy = R.getAsSingle(); + if (!MemCpy) + // Something went horribly wrong earlier, and we will have complained + // about it. + return StmtError(); + + ExprResult MemCpyRef = + BuildDeclRefExpr(MemCpy, Context.BuiltinFnTy, VK_PRValue, Loc, nullptr); + assert(MemCpyRef.isUsable() && "Builtin reference cannot fail"); + + Expr *CallArgs[] = {To, From, + IntegerLiteral::Create(Context, Size, SizeType, Loc)}; + ExprResult Call = + BuildCallExpr(/*Scope=*/nullptr, MemCpyRef.get(), Loc, CallArgs, Loc); + + assert(!Call.isInvalid() && "Call to __builtin_memcpy cannot fail!"); + return Call.getAs(); +} + static bool checkSimpleDecomposition( Sema &S, ArrayRef Bindings, ValueDecl *Src, QualType DecompType, const llvm::APSInt &NumElems, QualType ElemType, @@ -14050,54 +14099,7 @@ class SubscriptBuilder: public ExprBuilder { static StmtResult buildMemcpyForAssignmentOp(Sema &S, SourceLocation Loc, QualType T, const ExprBuilder &ToB, const ExprBuilder &FromB) { - // Compute the size of the memory buffer to be copied. - QualType SizeType = S.Context.getSizeType(); - llvm::APInt Size(S.Context.getTypeSize(SizeType), - S.Context.getTypeSizeInChars(T).getQuantity()); - - // Take the address of the field references for "from" and "to". We - // directly construct UnaryOperators here because semantic analysis - // does not permit us to take the address of an xvalue. - Expr *From = FromB.build(S, Loc); - From = UnaryOperator::Create( - S.Context, From, UO_AddrOf, S.Context.getPointerType(From->getType()), - VK_PRValue, OK_Ordinary, Loc, false, S.CurFPFeatureOverrides()); - Expr *To = ToB.build(S, Loc); - To = UnaryOperator::Create( - S.Context, To, UO_AddrOf, S.Context.getPointerType(To->getType()), - VK_PRValue, OK_Ordinary, Loc, false, S.CurFPFeatureOverrides()); - - const Type *E = T->getBaseElementTypeUnsafe(); - bool NeedsCollectableMemCpy = - E->isRecordType() && - E->castAs()->getDecl()->hasObjectMember(); - - // Create a reference to the __builtin_objc_memmove_collectable function - StringRef MemCpyName = NeedsCollectableMemCpy ? - "__builtin_objc_memmove_collectable" : - "__builtin_memcpy"; - LookupResult R(S, &S.Context.Idents.get(MemCpyName), Loc, - Sema::LookupOrdinaryName); - S.LookupName(R, S.TUScope, true); - - FunctionDecl *MemCpy = R.getAsSingle(); - if (!MemCpy) - // Something went horribly wrong earlier, and we will have complained - // about it. - return StmtError(); - - ExprResult MemCpyRef = S.BuildDeclRefExpr(MemCpy, S.Context.BuiltinFnTy, - VK_PRValue, Loc, nullptr); - assert(MemCpyRef.isUsable() && "Builtin reference cannot fail"); - - Expr *CallArgs[] = { - To, From, IntegerLiteral::Create(S.Context, Size, SizeType, Loc) - }; - ExprResult Call = S.BuildCallExpr(/*Scope=*/nullptr, MemCpyRef.get(), - Loc, CallArgs, Loc); - - assert(!Call.isInvalid() && "Call to __builtin_memcpy cannot fail!"); - return Call.getAs(); + return S.BuildMemCpyCall(Loc, T, FromB.build(S, Loc), ToB.build(S, Loc)); } /// Builds a statement that copies/moves the given entity from \p From to diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 90126a44c3525..649e782c01e06 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -2878,23 +2878,29 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { addFieldInit(FD, Ty, ParamRef); } + void createMemCpyCall(QualType Ty, SmallVectorImpl &AddTo) { + Expr *ParamRef = createParamReferenceExpr(); + Expr *FieldOfLocalClone = MemberExprBases.back(); + StmtResult Call = SemaRef.BuildMemCpyCall(KernelCallerSrcLoc, Ty, ParamRef, + FieldOfLocalClone); + + Expr *MemCpyCallExpr = Call.getAs(); + + AddTo.push_back(MemCpyCallExpr); + } + // Call __builtin_memcpy to copy modified type with pointer // to local clone - void generateBitwiseCopy(FieldDecl *FD, QualType Ty) { - // Need to change this to bitwise copy - addSimpleFieldInit(FD, Ty); + // Elizabeth - Refactor Sema once this works + void handleGeneratedType(FieldDecl *FD, QualType Ty) { + + addFieldInit(FD, Ty, None, + InitializationKind::CreateDefault(KernelCallerSrcLoc)); + addFieldMemberExpr(FD, Ty); - // Compute the size of the memory buffer to be copied. - /*QualType SizeType = S.Context.getSizeType(); - llvm::APInt Size(SemaRef.Context.getTypeSize(SizeType), - SemaRef.Context.getTypeSizeInChars(Ty).getQuantity()); + createMemCpyCall(Ty, BodyStmts); - Expr *ParamRef = createParamReferenceExpr(); - From = UnaryOperator::Create(SemaRef.Context, ParamRef, UO_AddrOf, - SemaRef.Context.getPointerType(ParamRef->getType()), VK_PRValue, - OK_Ordinary, KernelCallerSrcLoc, false, SemaRef.CurFPFeatureOverrides()); - */ - // How do I get a reference to the field of local clone? + removeFieldMemberExpr(FD, Ty); } MemberExpr *buildMemberExpr(Expr *Base, ValueDecl *Member) { @@ -3129,7 +3135,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { CXXRecordDecl *RD = Ty->getAsCXXRecordDecl(); assert(RD && "Type must be a C++ record type"); if (RD->hasAttr()) - generateBitwiseCopy(FD, Ty); + handleGeneratedType(FD, Ty); else addSimpleFieldInit(FD, Ty); return true; From 97f05952e59d508facc338cb693ebd990c89f95e Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Sun, 11 Sep 2022 16:30:58 -0700 Subject: [PATCH 06/18] Remove stale comments Signed-off-by: Elizabeth Andrews --- clang/lib/Sema/SemaSYCL.cpp | 14 +++----------- 1 file changed, 3 insertions(+), 11 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 649e782c01e06..b75a63deb66ed 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -2371,11 +2371,6 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { bool handleNonDecompStruct(const CXXRecordDecl *Base, const CXXBaseSpecifier &BS, QualType Ty) final { - // We need to create a list of BaseSpecifiers for each class - // Actions.ActOnBaseSpecifiers(ClassDecl, BaseInfo); - // Or maybe just setBases() which requires us to create a bunch of - // CXXBaseSpecifiers - CXXRecordDecl *BaseRecordDecl = Ty->getAsCXXRecordDecl(); SyclKernelPointerHandler PointerHandler(SemaRef, BaseRecordDecl); KernelObjVisitor Visitor{SemaRef}; @@ -2889,17 +2884,14 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { AddTo.push_back(MemCpyCallExpr); } - // Call __builtin_memcpy to copy modified type with pointer - // to local clone - // Elizabeth - Refactor Sema once this works + // Adds default initializer for generated type and creates + // a call to __builtin_memcpy to iniatilize local clone from + // kernel argument. void handleGeneratedType(FieldDecl *FD, QualType Ty) { - addFieldInit(FD, Ty, None, InitializationKind::CreateDefault(KernelCallerSrcLoc)); addFieldMemberExpr(FD, Ty); - createMemCpyCall(Ty, BodyStmts); - removeFieldMemberExpr(FD, Ty); } From 78eabb1e765150f9c79ab7a4c957edfdb67e74cd Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Mon, 12 Sep 2022 16:44:48 -0700 Subject: [PATCH 07/18] Fixed some bugs and added comments Signed-off-by: Elizabeth Andrews --- clang/lib/Sema/SemaSYCL.cpp | 96 +++++++++++++++++++++++++++++++------ 1 file changed, 81 insertions(+), 15 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index b75a63deb66ed..77ca4dba1bbe0 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1480,10 +1480,14 @@ void KernelObjVisitor::visitRecord(const CXXRecordDecl *Owner, ParentTy &Parent, visitComplexRecord(Owner, Parent, Wrapper, RecordTy, Handlers...); } else if (AnyTrue:: Value) { - if (RD->hasAttr()) + // We are currently in PointerHandler visitor. + if (RD->hasAttr()) { + // This is record containing pointers. visitComplexRecord(Owner, Parent, Wrapper, RecordTy, Handlers...); - else + } else { + // This is a record without pointers. visitSimpleRecord(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 @@ -1515,6 +1519,17 @@ void KernelObjVisitor::visitArray(const CXXRecordDecl *Owner, FieldDecl *Field, if (Field->hasAttr()) { visitComplexArray(Owner, Field, ArrayTy, Handlers...); + } else if (AnyTrue:: + Value) { + // We are currently in PointerHandler visitor. + if (Field->hasAttr()) { + // This is an array of pointers, or an array of a type containing + // pointers. + visitComplexArray(Owner, Field, ArrayTy, Handlers...); + } else { + // This is an array which does not contain pointers. + visitSimpleArray(Owner, Field, ArrayTy, Handlers...); + } } else { if (!AllTrue::Value) visitSimpleArray( @@ -1781,7 +1796,7 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { return true; } - bool handlePointerType(FieldDecl *, QualType) final { + bool handlePointerType(FieldDecl *FD, QualType Ty) final { PointerStack.back() = true; return true; } @@ -1793,6 +1808,10 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { } bool leaveStruct(const CXXRecordDecl *, FieldDecl *, QualType Ty) final { + // If a record needs to be decomposed, it is marked with + // SYCLRequiresDecompositionAttr. Else if a record contains + // a pointer, it is marked with SYCLGenerateNewTypeAttr. A record + // will never be marked with both attributes. if (CollectionStack.pop_back_val()) { RecordDecl *RD = Ty->getAsRecordDecl(); assert(RD && "should not be null."); @@ -1821,6 +1840,10 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { bool leaveStruct(const CXXRecordDecl *, const CXXBaseSpecifier &, QualType Ty) final { + // If a record needs to be decomposed, it is marked with + // SYCLRequiresDecompositionAttr. Else if a record contains + // a pointer, it is marked with SYCLGenerateNewTypeAttr. A record + // will never be marked with both attributes. if (CollectionStack.pop_back_val()) { RecordDecl *RD = Ty->getAsRecordDecl(); assert(RD && "should not be null."); @@ -1840,13 +1863,18 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { return true; } - // Elizabeth - Look into handling arrays bool enterArray(FieldDecl *, QualType ArrayTy, QualType ElementTy) final { CollectionStack.push_back(false); + PointerStack.push_back(false); return true; } bool leaveArray(FieldDecl *FD, QualType ArrayTy, QualType ElementTy) final { + // If an array needs to be decomposed, it is marked with + // SYCLRequiresDecompositionAttr. Else if the array is an array of pointers + // or an array of structs containing pointers, it is marked with + // SYCLGenerateNewTypeAttr. An array will never be marked with both + // attributes. if (CollectionStack.pop_back_val()) { // Cannot assert, since in MD arrays we'll end up marking them multiple // times. @@ -1854,11 +1882,19 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { FD->addAttr(SYCLRequiresDecompositionAttr::CreateImplicit( SemaRef.getASTContext())); CollectionStack.back() = true; + PointerStack.pop_back(); + } else if (PointerStack.pop_back_val()) { + if (!FD->hasAttr()) + FD->addAttr( + SYCLGenerateNewTypeAttr::CreateImplicit(SemaRef.getASTContext())); + PointerStack.back() = true; } return true; } }; +// This visitor is used to traverse a non-decomposed record/array to +// generate a new type corresponding to this record/array. class SyclKernelPointerHandler : public SyclKernelFieldHandler { llvm::SmallVector ModifiedRecords; SmallVector ModifiedBases; @@ -2004,6 +2040,14 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { return true; } + bool handleSimpleArrayType(FieldDecl *FD, QualType Ty) final { + addField(FD, Ty); + return true; + } + + // Elizabeth - We need to handle complex array by adding enterArray, etc to + // create new array type + public: QualType getNewType() { CXXRecordDecl *ModifiedRD = ModifiedRecords.pop_back_val(); @@ -2017,7 +2061,7 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { return QualType(ModifiedRD->getTypeForDecl(), 0); } - // Elizabeth - Need to handle KernelHandler, array + // Elizabeth - Need to handle KernelHandler, arrays }; // A type to Create and own the FunctionDecl for the kernel. @@ -2203,6 +2247,22 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { return FD; } + // If the record has been marked with SYCLGenerateNewTypeAttr, + // it implies that it contains a pointer within. This function + // defines a PointerHandler visitor which visits this record + // recursively and modifies the address spaces of any pointer + // found as required, thereby generating a new record with all + // pointers in 'right' address space. PointerHandler.getNewType() + // returns this generated type, which is then added an openCL + // kernel argument. + QualType GenerateNewType(const CXXRecordDecl *RD) { + SyclKernelPointerHandler PointerHandler(SemaRef, RD); + KernelObjVisitor Visitor{SemaRef}; + Visitor.VisitRecordBases(RD, PointerHandler); + Visitor.VisitRecordFields(RD, PointerHandler); + return PointerHandler.getNewType(); + } + public: static constexpr const bool VisitInsideSimpleContainers = false; SyclKernelDeclCreator(Sema &S, SourceLocation Loc, bool IsInline, @@ -2360,23 +2420,29 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { bool handleNonDecompStruct(const CXXRecordDecl *RD, FieldDecl *FD, QualType Ty) final { + // This is a field which should not be decomposed. CXXRecordDecl *FieldRecordDecl = Ty->getAsCXXRecordDecl(); - SyclKernelPointerHandler PointerHandler(SemaRef, FieldRecordDecl); - KernelObjVisitor Visitor{SemaRef}; - Visitor.VisitRecordBases(FieldRecordDecl, PointerHandler); - Visitor.VisitRecordFields(FieldRecordDecl, PointerHandler); - addParam(FD, PointerHandler.getNewType()); + assert(FieldRecordDecl && "Type must be a C++ record type"); + // Check if we need to generate a new type for this record, + // i.e. this record contains pointers. + if (FieldRecordDecl->hasAttr()) + addParam(FD, GenerateNewType(FieldRecordDecl)); + else + addParam(FD, Ty); return true; } bool handleNonDecompStruct(const CXXRecordDecl *Base, const CXXBaseSpecifier &BS, QualType Ty) final { + // This is a base class which should not be decomposed. CXXRecordDecl *BaseRecordDecl = Ty->getAsCXXRecordDecl(); - SyclKernelPointerHandler PointerHandler(SemaRef, BaseRecordDecl); - KernelObjVisitor Visitor{SemaRef}; - Visitor.VisitRecordBases(BaseRecordDecl, PointerHandler); - Visitor.VisitRecordFields(BaseRecordDecl, PointerHandler); - addParam(BS, PointerHandler.getNewType()); + assert(BaseRecordDecl && "Type must be a C++ record type"); + // Check if we need to generate a new type for this record, + // i.e. this record contains pointers. + if (BaseRecordDecl->hasAttr()) + addParam(BS, GenerateNewType(BaseRecordDecl)); + else + addParam(BS, Ty); return true; } From 4344fcbf440350479c7c64cda7413557f45d27d2 Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Wed, 14 Sep 2022 06:08:30 -0700 Subject: [PATCH 08/18] Remove code handling arrays. Will be done in followup PR Signed-off-by: Elizabeth Andrews --- clang/lib/Sema/SemaSYCL.cpp | 53 ++++++++++++++++++++++--------------- 1 file changed, 31 insertions(+), 22 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 77ca4dba1bbe0..ed5985a5bc9f0 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1521,15 +1521,14 @@ void KernelObjVisitor::visitArray(const CXXRecordDecl *Owner, FieldDecl *Field, visitComplexArray(Owner, Field, ArrayTy, Handlers...); } else if (AnyTrue:: Value) { - // We are currently in PointerHandler visitor. - if (Field->hasAttr()) { - // This is an array of pointers, or an array of a type containing - // pointers. - visitComplexArray(Owner, Field, ArrayTy, Handlers...); - } else { - // This is an array which does not contain pointers. - visitSimpleArray(Owner, Field, ArrayTy, Handlers...); - } + assert(!Field->hasAttr() && + "Arrays should trigger decomposition"); + // We are currently in PointerHandler visitor, which implies this is a + // 'simple' array i.e. one that does not include special types or pointers. + // Array of pointers/ array of type containing pointers will be handled in + // a follow-up PR. Currently, they continue to trigger decomposition, and + // will be handled in 'if' statment above. + visitSimpleArray(Owner, Field, ArrayTy, Handlers...); } else { if (!AllTrue::Value) visitSimpleArray( @@ -1770,6 +1769,14 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { llvm::SmallVector CollectionStack; llvm::SmallVector PointerStack; + // FIXME: Array of pointers/ array of type containing pointers + // will be handled in a follow up PR. Currently, they continue + // to trigger decomposition. + // TODO: Remove this method once arrays are handled correctly + bool isArrayElement(const FieldDecl *FD, QualType Ty) const { + return !SemaRef.getASTContext().hasSameType(FD->getType(), Ty); + } + public: static constexpr const bool VisitUnionBody = false; static constexpr const bool VisitNthArrayElement = false; @@ -1797,7 +1804,13 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { } bool handlePointerType(FieldDecl *FD, QualType Ty) final { - PointerStack.back() = true; + // FIXME: Array of pointers/ array of type containing pointers + // will be handled in a follow up PR. Currently, they continue + // to trigger decomposition. + if (isArrayElement(FD, Ty)) + CollectionStack.back() = true; + else + PointerStack.back() = true; return true; } @@ -1870,11 +1883,6 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { } bool leaveArray(FieldDecl *FD, QualType ArrayTy, QualType ElementTy) final { - // If an array needs to be decomposed, it is marked with - // SYCLRequiresDecompositionAttr. Else if the array is an array of pointers - // or an array of structs containing pointers, it is marked with - // SYCLGenerateNewTypeAttr. An array will never be marked with both - // attributes. if (CollectionStack.pop_back_val()) { // Cannot assert, since in MD arrays we'll end up marking them multiple // times. @@ -1884,10 +1892,10 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { CollectionStack.back() = true; PointerStack.pop_back(); } else if (PointerStack.pop_back_val()) { - if (!FD->hasAttr()) - FD->addAttr( - SYCLGenerateNewTypeAttr::CreateImplicit(SemaRef.getASTContext())); - PointerStack.back() = true; + // FIXME: Array of pointers/ array of type containing pointers + // will be handled in a follow up PR. Currently, they continue + // to trigger decomposition. + llvm_unreachable("PointerStack should not be true when handling arrays."); } return true; } @@ -2045,8 +2053,9 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { return true; } - // Elizabeth - We need to handle complex array by adding enterArray, etc to - // create new array type + // FIXME: Array of pointers/ array of types containing pointers + // will be handled in a follow-up PR. Currently they continue to + // trigger decomposition. public: QualType getNewType() { @@ -2061,7 +2070,7 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { return QualType(ModifiedRD->getTypeForDecl(), 0); } - // Elizabeth - Need to handle KernelHandler, arrays + // Elizabeth - Need to handle KernelHandler }; // A type to Create and own the FunctionDecl for the kernel. From a2dc51d10795a96f3c41e2be31e9252071cad3bf Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Wed, 14 Sep 2022 16:57:32 -0700 Subject: [PATCH 09/18] Remove code for array and fix base memcpy Signed-off-by: Elizabeth Andrews --- clang/include/clang/Sema/Sema.h | 3 - clang/lib/Sema/SemaDeclCXX.cpp | 98 ++++++++++++++--------------- clang/lib/Sema/SemaSYCL.cpp | 92 ++++++++++++++++++++++----- clang/test/SemaSYCL/inheritance.cpp | 60 ++++++++++++++++-- 4 files changed, 181 insertions(+), 72 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index b429f8fbf4a06..6ba624c6b7c2b 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -3098,9 +3098,6 @@ class Sema final { void SetParamDefaultArgument(ParmVarDecl *Param, Expr *DefaultArg, SourceLocation EqualLoc); - StmtResult BuildMemCpyCall(SourceLocation Loc, QualType T, Expr *From, - Expr *To); - // Contexts where using non-trivial C union types can be disallowed. This is // passed to err_non_trivial_c_union_in_invalid_context. enum NonTrivialCUnionContext { diff --git a/clang/lib/Sema/SemaDeclCXX.cpp b/clang/lib/Sema/SemaDeclCXX.cpp index 046083ab6fcea..20ea8ccfce50b 100644 --- a/clang/lib/Sema/SemaDeclCXX.cpp +++ b/clang/lib/Sema/SemaDeclCXX.cpp @@ -902,55 +902,6 @@ Sema::ActOnDecompositionDeclarator(Scope *S, Declarator &D, return New; } -StmtResult Sema::BuildMemCpyCall(SourceLocation Loc, QualType T, Expr *From, - Expr *To) { - // Compute the size of the memory buffer to be copied. - QualType SizeType = Context.getSizeType(); - llvm::APInt Size(Context.getTypeSize(SizeType), - Context.getTypeSizeInChars(T).getQuantity()); - - // Take the address of the field references for "from" and "to". We - // directly construct UnaryOperators here because semantic analysis - // does not permit us to take the address of an xvalue. - From = UnaryOperator::Create( - Context, From, UO_AddrOf, Context.getPointerType(From->getType()), - VK_PRValue, OK_Ordinary, Loc, false, CurFPFeatureOverrides()); - To = UnaryOperator::Create(Context, To, UO_AddrOf, - Context.getPointerType(To->getType()), VK_PRValue, - OK_Ordinary, Loc, false, CurFPFeatureOverrides()); - - const Type *E = T->getBaseElementTypeUnsafe(); - bool NeedsCollectableMemCpy = - !getLangOpts().SYCLIsDevice && E->isRecordType() && - E->castAs()->getDecl()->hasObjectMember(); - - // Create a reference to the __builtin_objc_memmove_collectable function - StringRef MemCpyName = NeedsCollectableMemCpy - ? "__builtin_objc_memmove_collectable" - : "__builtin_memcpy"; - LookupResult R(*this, &Context.Idents.get(MemCpyName), Loc, - Sema::LookupOrdinaryName); - LookupName(R, TUScope, true); - - FunctionDecl *MemCpy = R.getAsSingle(); - if (!MemCpy) - // Something went horribly wrong earlier, and we will have complained - // about it. - return StmtError(); - - ExprResult MemCpyRef = - BuildDeclRefExpr(MemCpy, Context.BuiltinFnTy, VK_PRValue, Loc, nullptr); - assert(MemCpyRef.isUsable() && "Builtin reference cannot fail"); - - Expr *CallArgs[] = {To, From, - IntegerLiteral::Create(Context, Size, SizeType, Loc)}; - ExprResult Call = - BuildCallExpr(/*Scope=*/nullptr, MemCpyRef.get(), Loc, CallArgs, Loc); - - assert(!Call.isInvalid() && "Call to __builtin_memcpy cannot fail!"); - return Call.getAs(); -} - static bool checkSimpleDecomposition( Sema &S, ArrayRef Bindings, ValueDecl *Src, QualType DecompType, const llvm::APSInt &NumElems, QualType ElemType, @@ -14099,7 +14050,54 @@ class SubscriptBuilder: public ExprBuilder { static StmtResult buildMemcpyForAssignmentOp(Sema &S, SourceLocation Loc, QualType T, const ExprBuilder &ToB, const ExprBuilder &FromB) { - return S.BuildMemCpyCall(Loc, T, FromB.build(S, Loc), ToB.build(S, Loc)); + // Compute the size of the memory buffer to be copied. + QualType SizeType = S.Context.getSizeType(); + llvm::APInt Size(S.Context.getTypeSize(SizeType), + S.Context.getTypeSizeInChars(T).getQuantity()); + + // Take the address of the field references for "from" and "to". We + // directly construct UnaryOperators here because semantic analysis + // does not permit us to take the address of an xvalue. + Expr *From = FromB.build(S, Loc); + From = UnaryOperator::Create( + S.Context, From, UO_AddrOf, S.Context.getPointerType(From->getType()), + VK_PRValue, OK_Ordinary, Loc, false, S.CurFPFeatureOverrides()); + Expr *To = ToB.build(S, Loc); + To = UnaryOperator::Create( + S.Context, To, UO_AddrOf, S.Context.getPointerType(To->getType()), + VK_PRValue, OK_Ordinary, Loc, false, S.CurFPFeatureOverrides()); + + const Type *E = T->getBaseElementTypeUnsafe(); + bool NeedsCollectableMemCpy = + E->isRecordType() && + E->castAs()->getDecl()->hasObjectMember(); + + // Create a reference to the __builtin_objc_memmove_collectable function + StringRef MemCpyName = NeedsCollectableMemCpy ? + "__builtin_objc_memmove_collectable" : + "__builtin_memcpy"; + LookupResult R(S, &S.Context.Idents.get(MemCpyName), Loc, + Sema::LookupOrdinaryName); + S.LookupName(R, S.TUScope, true); + + FunctionDecl *MemCpy = R.getAsSingle(); + if (!MemCpy) + // Something went horribly wrong earlier, and we will have complained + // about it. + return StmtError(); + + ExprResult MemCpyRef = S.BuildDeclRefExpr(MemCpy, S.Context.BuiltinFnTy, + VK_PRValue, Loc, nullptr); + assert(MemCpyRef.isUsable() && "Builtin reference cannot fail"); + + Expr *CallArgs[] = { + To, From, IntegerLiteral::Create(S.Context, Size, SizeType, Loc) + }; + ExprResult Call = S.BuildCallExpr(/*Scope=*/nullptr, MemCpyRef.get(), + Loc, CallArgs, Loc); + + assert(!Call.isInvalid() && "Call to __builtin_memcpy cannot fail!"); + return Call.getAs(); } /// Builds a statement that copies/moves the given entity from \p From to diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index ed5985a5bc9f0..cfa8a169a1193 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1908,7 +1908,8 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { SmallVector ModifiedBases; IdentifierInfo *getModifiedName(IdentifierInfo *Id) { - std::string Name = (Id->getName() + Twine("_generated")).str(); + std::string Name = + Id ? (Twine("_generated_") + Id->getName()).str() : "_generated_"; return &SemaRef.getASTContext().Idents.get(Name); } @@ -1917,7 +1918,7 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { // the visitor traverses kernel object record fields. void createNewType(const CXXRecordDecl *RD) { auto *ModifiedRD = CXXRecordDecl::Create( - SemaRef.getASTContext(), TTK_Struct, + SemaRef.getASTContext(), RD->getTagKind(), const_cast(RD->getDeclContext()), SourceLocation(), SourceLocation(), getModifiedName(RD->getIdentifier())); ModifiedRD->startDefinition(); @@ -2731,7 +2732,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { SourceLocation KernelCallerSrcLoc; // KernelCallerFunc source location. // Contains a count of how many containers we're in. This is used by the // pointer-struct-wrapping code to ensure that we don't try to wrap - // non-top-level pointers. + // top-level pointers. uint64_t StructDepth = 0; VarDecl *KernelHandlerClone = nullptr; @@ -2948,15 +2949,58 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { addFieldInit(FD, Ty, ParamRef); } - void createMemCpyCall(QualType Ty, SmallVectorImpl &AddTo) { - Expr *ParamRef = createParamReferenceExpr(); - Expr *FieldOfLocalClone = MemberExprBases.back(); - StmtResult Call = SemaRef.BuildMemCpyCall(KernelCallerSrcLoc, Ty, ParamRef, - FieldOfLocalClone); + Expr *addDerivedToBaseCastExpr(const CXXRecordDecl *RD, + const CXXBaseSpecifier &BS, + Expr *LocalCloneRef) { + CXXCastPath BasePath; + QualType DerivedTy(RD->getTypeForDecl(), 0); + QualType BaseTy = BS.getType(); + SemaRef.CheckDerivedToBaseConversion(DerivedTy, BaseTy, KernelCallerSrcLoc, + SourceRange(), &BasePath, + /*IgnoreBaseAccess*/ true); + auto Cast = ImplicitCastExpr::Create( + SemaRef.Context, SemaRef.Context.getPointerType(BaseTy), + CK_DerivedToBase, LocalCloneRef, + /* CXXCastPath=*/&BasePath, VK_LValue, FPOptionsOverride()); + return Cast; + } + + Expr *getAddressOf(Expr *E) { + return UnaryOperator::Create(SemaRef.Context, E, UO_AddrOf, + SemaRef.Context.getPointerType(E->getType()), + VK_PRValue, OK_Ordinary, KernelCallerSrcLoc, + false, SemaRef.CurFPFeatureOverrides()); + } - Expr *MemCpyCallExpr = Call.getAs(); + Expr *buildMemCpyCall(Expr *From, Expr *To, QualType T) { + // Compute the size of the memory buffer to be copied. + QualType SizeType = SemaRef.Context.getSizeType(); + llvm::APInt Size(SemaRef.Context.getTypeSize(SizeType), + SemaRef.Context.getTypeSizeInChars(T).getQuantity()); - AddTo.push_back(MemCpyCallExpr); + LookupResult R(SemaRef, &SemaRef.Context.Idents.get("__builtin_memcpy"), + KernelCallerSrcLoc, Sema::LookupOrdinaryName); + SemaRef.LookupName(R, SemaRef.TUScope, true); + + FunctionDecl *MemCpy = R.getAsSingle(); + + assert(MemCpy && "__builtin_memcpy should be found"); + + ExprResult MemCpyRef = + SemaRef.BuildDeclRefExpr(MemCpy, SemaRef.Context.BuiltinFnTy, + VK_PRValue, KernelCallerSrcLoc, nullptr); + + assert(MemCpyRef.isUsable() && "Builtin reference cannot fail"); + + Expr *CallArgs[] = {To, From, + IntegerLiteral::Create(SemaRef.Context, Size, SizeType, + KernelCallerSrcLoc)}; + ExprResult Call = + SemaRef.BuildCallExpr(/*Scope=*/nullptr, MemCpyRef.get(), + KernelCallerSrcLoc, CallArgs, KernelCallerSrcLoc); + + assert(!Call.isInvalid() && "Call to __builtin_memcpy cannot fail!"); + return Call.getAs(); } // Adds default initializer for generated type and creates @@ -2966,10 +3010,26 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { addFieldInit(FD, Ty, None, InitializationKind::CreateDefault(KernelCallerSrcLoc)); addFieldMemberExpr(FD, Ty); - createMemCpyCall(Ty, BodyStmts); + Expr *ParamRef = getAddressOf(createParamReferenceExpr()); + Expr *LocalCloneRef = getAddressOf(MemberExprBases.back()); + Expr *MemCpyCallExpr = buildMemCpyCall(ParamRef, LocalCloneRef, Ty); + BodyStmts.push_back(MemCpyCallExpr); removeFieldMemberExpr(FD, Ty); } + // Adds default initializer for generated base and creates + // a call to __builtin_memcpy to initialize the base of local clone + // from kernel argument. + void handleGeneratedType(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS, + QualType Ty) { + addBaseInit(BS, Ty, InitializationKind::CreateDefault(KernelCallerSrcLoc)); + Expr *ParamRef = getAddressOf(createParamReferenceExpr()); + Expr *LocalCloneRef = getAddressOf(MemberExprBases.back()); + LocalCloneRef = addDerivedToBaseCastExpr(RD, BS, LocalCloneRef); + Expr *MemCpyCallExpr = buildMemCpyCall(ParamRef, LocalCloneRef, Ty); + BodyStmts.push_back(MemCpyCallExpr); + } + MemberExpr *buildMemberExpr(Expr *Base, ValueDecl *Member) { DeclAccessPair MemberDAP = DeclAccessPair::make(Member, AS_none); MemberExpr *Result = SemaRef.BuildMemberExpr( @@ -3208,9 +3268,14 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { return true; } - bool handleNonDecompStruct(const CXXRecordDecl *Base, + bool handleNonDecompStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS, QualType Ty) final { - addSimpleBaseInit(BS, Ty); + CXXRecordDecl *BaseDecl = Ty->getAsCXXRecordDecl(); + assert(BaseDecl && "Type must be a C++ record type"); + if (BaseDecl->hasAttr()) + handleGeneratedType(RD, BS, Ty); + else + addSimpleBaseInit(BS, Ty); return true; } @@ -3276,7 +3341,6 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { SemaRef.Context, BaseTy, CK_DerivedToBase, MemberExprBases.back(), /* CXXCastPath=*/&BasePath, VK_LValue, FPOptionsOverride()); MemberExprBases.push_back(Cast); - addCollectionInitListExpr(BaseTy->getAsCXXRecordDecl()); return true; } diff --git a/clang/test/SemaSYCL/inheritance.cpp b/clang/test/SemaSYCL/inheritance.cpp index 684fa835e8599..8ca3ece92ac36 100644 --- a/clang/test/SemaSYCL/inheritance.cpp +++ b/clang/test/SemaSYCL/inheritance.cpp @@ -2,6 +2,12 @@ #include "Inputs/sycl.hpp" +class third_base { +public: + int *d; + sycl::accessor AccField; +}; + class second_base { public: int *e; @@ -21,7 +27,7 @@ struct base { InnerField obj; }; -struct derived : base, second_base { +struct derived : base, second_base, third_base{ int a; void operator()() const { @@ -40,11 +46,17 @@ int main() { } // Check declaration of the kernel -// CHECK: derived{{.*}} 'void (base, __wrapper_class, int) +// CHECK: derived{{.*}} 'void (base, _generated_second_base, __wrapper_class, +// CHECK-SAME: __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, int) // Check parameters of the kernel // CHECK: ParmVarDecl {{.*}} used _arg__base 'base' -// CHECK: ParmVarDecl {{.*}} used _arg_e '__wrapper_class' +// CHECK: ParmVarDecl {{.*}} used _arg__base '_generated_second_base' +// CHECK: ParmVarDecl {{.*}} used _arg_d '__wrapper_class' +// CHECK: ParmVarDecl {{.*}} used _arg_AccField '__global char *' +// CHECK: ParmVarDecl {{.*}} used _arg_AccField 'sycl::range<1>' +// CHECK: ParmVarDecl {{.*}} used _arg_AccField 'sycl::range<1>' +// CHECK: ParmVarDecl {{.*}} used _arg_AccField 'sycl::id<1>' // CHECK: ParmVarDecl {{.*}} used _arg_a 'int' // Check initializers for derived and base classes. @@ -52,13 +64,51 @@ int main() { // Base classes should be initialized first. // CHECK: VarDecl {{.*}} derived 'derived' cinit // CHECK-NEXT: InitListExpr {{.*}} 'derived' + +// base is a simple class with no corresponding generated type. Therefore +// copy from ParamVar // 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' + +// second_base contains pointers and therefore the ParamVar is a new generated +// type. Default construct this class and initialize second_base via memcpy in +// body statements. +// CHECK-NEXT: CXXConstructExpr {{.*}} 'second_base' 'void () noexcept' + +// third_base contains special type accessor. Therefore it is decomposed and it's +// date members are copied from corrsponding ParamVar +// CHECK-NEXT: InitListExpr {{.*}} 'third_base' // 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: DeclRefExpr {{.*}} lvalue ParmVar {{.*}} '_arg_d' '__wrapper_class' +// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::accessor' + +// Initialize fields of 'derived' // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' // CHECK-NEXT: DeclRefExpr {{.*}} lvalue ParmVar {{.*}} '_arg_a' 'int' + + +// Check kernel body for call to __builtin_memcpy to initialize second_base +// CHECK: CallExpr {{.*}} 'void *' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'void *(*)(void *, const void *, unsigned long) noexcept' +// CHECK-NEXT: DeclRefExpr {{.*}} Function {{.*}} '__builtin_memcpy' 'void *(void *, const void *, unsigned long) noexcept' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'void *' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'second_base *' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'second_base *' lvalue +// CHECK-NEXT: UnaryOperator {{.*}} 'derived *' prefix '&' cannot overflow +// CHECK-NEXT: DeclRefExpr {{.*}} 'derived' lvalue Var {{.*}} 'derived' 'derived' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const void *' +// CHECK-NEXT: UnaryOperator {{.*}} '_generated_second_base *' prefix '&' cannot overflow +// CHECK-NEXT: DeclRefExpr {{.*}} '_generated_second_base' lvalue ParmVar {{.*}} '_arg__base' '_generated_second_base' +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 8 + + + +// Check kernel body for call to __init function of accessor +// CHECK: CXXMemberCallExpr +// CHECK-NEXT: MemberExpr {{.*}} lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} lvalue .AccField +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'third_base' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} 'derived' lvalue Var {{.*}} 'derived' 'derived' From bcf0ab35a59ea4eaf91d316851770b7fe89c65a6 Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Thu, 15 Sep 2022 19:20:07 -0700 Subject: [PATCH 10/18] Exclude Lambdas and fix some tests Signed-off-by: Elizabeth Andrews --- clang/lib/Sema/SemaSYCL.cpp | 18 +++- .../no_opaque_pointers-in-structs.cpp | 4 +- .../test/CodeGenSYCL/struct_kernel_param.cpp | 8 +- clang/test/SemaSYCL/inheritance.cpp | 5 +- clang/test/SemaSYCL/union-kernel-param2.cpp | 90 ++++++++++++++----- 5 files changed, 91 insertions(+), 34 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index cfa8a169a1193..ce4a377ce764a 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1837,8 +1837,22 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { RecordDecl *RD = Ty->getAsRecordDecl(); assert(RD && "should not be null."); if (!RD->hasAttr()) - RD->addAttr( - SYCLGenerateNewTypeAttr::CreateImplicit(SemaRef.getASTContext())); + // Do not generate a new type if the record corresponds to a + // lambda. Currently the fields/bases of the local clone + // corresponding to these generated types are intialized using + // their default constructors(Actual initialization is done via + // memcpy in kernel body.) to maintain the integrity of the + // InitListExpr we generate for Kernel Object local clone. + // Records correspondng to lambdas which have captures do not have + // a default constructor and so current logic fails for lambdas. + // FIXME: Can/Should we stop triggering decomposition for lambdas + // with pointers? + if (!RD->isLambda()) + RD->addAttr( + SYCLGenerateNewTypeAttr::CreateImplicit(SemaRef.getASTContext())); + else + RD->addAttr(SYCLRequiresDecompositionAttr::CreateImplicit( + SemaRef.getASTContext())); PointerStack.back() = true; } return true; diff --git a/clang/test/CodeGenSYCL/no_opaque_pointers-in-structs.cpp b/clang/test/CodeGenSYCL/no_opaque_pointers-in-structs.cpp index 4ce18fe8fe797..e2815e3606596 100644 --- a/clang/test/CodeGenSYCL/no_opaque_pointers-in-structs.cpp +++ b/clang/test/CodeGenSYCL/no_opaque_pointers-in-structs.cpp @@ -35,14 +35,14 @@ int main() { // CHECK: %[[WRAPPER_F1:[a-zA-Z0-9_.]+]] = type { i32 addrspace(1)* } // CHECK: %[[WRAPPER_F2:[a-zA-Z0-9_.]+]] = type { float addrspace(1)* } -// CHECK: %[[WRAPPER_F:[a-zA-Z0-9_.]+]] = type { float addrspace(1)* } +// CHECK: %[[GENERATED_A:[a-zA-Z0-9_.]+]] = type { float addrspace(1)* } // CHECK: %[[WRAPPER_F4_1:[a-zA-Z0-9_.]+]] = type { i32 addrspace(1)* } // CHECK: %[[WRAPPER_F4_2:[a-zA-Z0-9_.]+]] = type { i32 addrspace(1)* } // CHECK: %[[WRAPPER_LAMBDA_PTR:[a-zA-Z0-9_.]+]] = type { float addrspace(1)* } // CHECK: define {{.*}}spir_kernel void @{{.*}}structs // CHECK-SAME: %[[WRAPPER_F1]]* noundef byval(%[[WRAPPER_F1]]) align 8 %_arg_F1, // CHECK-SAME: %[[WRAPPER_F2]]* noundef byval(%[[WRAPPER_F2]]) align 8 %_arg_F2, -// CHECK-SAME: %[[WRAPPER_F]]* noundef byval(%[[WRAPPER_F]]) align 8 %_arg_F, +// CHECK-SAME: %[[GENERATED_A]]* noundef byval(%[[GENERATED_A]]) align 8 %_arg_F3, // CHECK-SAME: %[[WRAPPER_F4_1]]* noundef byval(%[[WRAPPER_F4_1]]) align 8 %_arg_F4 // CHECK-SAME: %[[WRAPPER_F4_2]]* noundef byval(%[[WRAPPER_F4_2]]) align 8 %_arg_F41 // CHECK: define {{.*}}spir_kernel void @{{.*}}lambdas{{.*}}(%[[WRAPPER_LAMBDA_PTR]]* noundef byval(%[[WRAPPER_LAMBDA_PTR]]) align 8 %_arg_Ptr) diff --git a/clang/test/CodeGenSYCL/struct_kernel_param.cpp b/clang/test/CodeGenSYCL/struct_kernel_param.cpp index c5e5df2d7ef2c..16a81faefda49 100644 --- a/clang/test/CodeGenSYCL/struct_kernel_param.cpp +++ b/clang/test/CodeGenSYCL/struct_kernel_param.cpp @@ -7,12 +7,8 @@ // CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 0 }, // 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 }, -// FldArr -// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 24 }, -// FldFloat -// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 8, 32 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 12, 40 }, +// MyStruct is not decomposed since it does not contain special types. +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 40, 16 }, // CHECK-EMPTY: // CHECK-NEXT: { kernel_param_kind_t::kind_invalid, -987654321, -987654321 }, // CHECK-NEXT:}; diff --git a/clang/test/SemaSYCL/inheritance.cpp b/clang/test/SemaSYCL/inheritance.cpp index 8ca3ece92ac36..a1b8052210699 100644 --- a/clang/test/SemaSYCL/inheritance.cpp +++ b/clang/test/SemaSYCL/inheritance.cpp @@ -77,7 +77,7 @@ int main() { // CHECK-NEXT: CXXConstructExpr {{.*}} 'second_base' 'void () noexcept' // third_base contains special type accessor. Therefore it is decomposed and it's -// date members are copied from corrsponding ParamVar +// data members are copied from corresponding ParamVar // CHECK-NEXT: InitListExpr {{.*}} 'third_base' // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int *' // CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' @@ -89,7 +89,6 @@ int main() { // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' // CHECK-NEXT: DeclRefExpr {{.*}} lvalue ParmVar {{.*}} '_arg_a' 'int' - // Check kernel body for call to __builtin_memcpy to initialize second_base // CHECK: CallExpr {{.*}} 'void *' // CHECK-NEXT: ImplicitCastExpr {{.*}} 'void *(*)(void *, const void *, unsigned long) noexcept' @@ -104,8 +103,6 @@ int main() { // CHECK-NEXT: DeclRefExpr {{.*}} '_generated_second_base' lvalue ParmVar {{.*}} '_arg__base' '_generated_second_base' // CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 8 - - // Check kernel body for call to __init function of accessor // CHECK: CXXMemberCallExpr // CHECK-NEXT: MemberExpr {{.*}} lvalue .__init diff --git a/clang/test/SemaSYCL/union-kernel-param2.cpp b/clang/test/SemaSYCL/union-kernel-param2.cpp index 7d38aa751a59e..1bb0aab19d785 100644 --- a/clang/test/SemaSYCL/union-kernel-param2.cpp +++ b/clang/test/SemaSYCL/union-kernel-param2.cpp @@ -1,12 +1,11 @@ -// RUN: %clang_cc1 -fsycl-is-device -ast-dump %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -ast-dump %s | FileCheck %s // This test checks that compiler generates correct kernel arguments for // a struct-with-an-array-of-unions and a array-of-struct-with-a-union. -template -__attribute__((sycl_kernel)) void a_kernel(Func kernelFunc) { - kernelFunc(); -} +#include "sycl.hpp" + +sycl::queue myQueue; int main() { @@ -25,18 +24,38 @@ int main() { float b; char c; } union_mem; - int *d; + sycl::accessor AccField; } struct_mem; - a_kernel( - [=]() { - int local = union_mem.struct_mem.a[2]; - }); + struct MyStructWithPtr { + union MyUnion { + int a[3]; + float b; + char c; + } union_mem; + int *d; + } structWithPtr_mem; + + myQueue.submit([&](sycl::handler &h) { + h.single_task( + [=]() { + int local = union_mem.struct_mem.a[2]; + }); + }); - a_kernel( - [=]() { - int local = struct_mem.union_mem.a[2]; - }); + myQueue.submit([&](sycl::handler &h) { + h.single_task( + [=]() { + int local = struct_mem.union_mem.a[2]; + }); + }); + + myQueue.submit([&](sycl::handler &h) { + h.single_task( + [=]() { + int local = structWithPtr_mem.union_mem.a[2]; + }); + }); } // Check kernel_A parameters @@ -53,9 +72,12 @@ int main() { // CHECK-NEXT: DeclRefExpr {{.*}} 'union MyUnion':'MyUnion' lvalue ParmVar {{.*}} '_arg_union_mem' 'union MyUnion':'MyUnion' // Check kernel_B parameters -// CHECK: FunctionDecl {{.*}}kernel_B{{.*}} 'void (union MyUnion, __wrapper_class)' +// CHECK: FunctionDecl {{.*}}kernel_B{{.*}} 'void (union MyUnion, __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>)' // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_union_mem 'union MyUnion':'MyStruct::MyUnion' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_d '__wrapper_class' +// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_AccField '__global char *' +// CHECK: ParmVarDecl {{.*}} used _arg_AccField 'sycl::range<1>' +// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_AccField 'sycl::range<1>' +// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_AccField 'sycl::id<1>' // Check kernel_B inits // CHECK-NEXT: CompoundStmt @@ -66,7 +88,35 @@ int main() { // CHECK-NEXT: CXXConstructExpr {{.*}} 'union MyUnion':'MyStruct::MyUnion' 'void (const MyStruct::MyUnion &) noexcept' // 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: ImplicitCastExpr -// CHECK-NEXT: MemberExpr -// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_d' '__wrapper_class' +// CHECK-NEXT: CXXConstructExpr {{.*}} 'sycl::accessor' + +// Check call to __init to initialize AccField +// CHECK-NEXT: CXXMemberCallExpr +// CHECK-NEXT: MemberExpr {{.*}} lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} lvalue .AccField +// CHECK-NEXT: MemberExpr {{.*}} lvalue .struct_mem +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}union-kernel-param2.cpp:48:9)' lvalue Var {{.*}} '' '(lambda at {{.*}}union-kernel-param2.cpp:48:9)' + +// Check kernel_C parameters +// CHECK: FunctionDecl {{.*}}kernel_C{{.*}} 'void (_generated_MyStructWithPtr)' +// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_structWithPtr_mem '_generated_MyStructWithPtr' + +// Check kernel_C inits +// CHECK-NEXT: CompoundStmt +// CHECK-NEXT: DeclStmt +// CHECK-NEXT: VarDecl {{.*}} cinit +// CHECK-NEXT: InitListExpr +// CHECK-NEXT: CXXConstructExpr {{.*}} 'struct MyStructWithPtr':'MyStructWithPtr' 'void () noexcept' + +// Check call to __builtin_memcpy to initialize structWithPtr_mem +// CHECK-NEXT: CallExpr +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'void *(*)(void *, const void *, unsigned long) noexcept' +// CHECK-NEXT: DeclRefExpr {{.*}} Function {{.*}} '__builtin_memcpy' 'void *(void *, const void *, unsigned long) noexcept' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'void *' +// CHECK-NEXT: UnaryOperator {{.*}} 'struct MyStructWithPtr *' prefix '&' cannot overflow +// CHECK-NEXT: MemberExpr {{.*}} 'struct MyStructWithPtr':'MyStructWithPtr' lvalue .structWithPtr_mem +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}union-kernel-param2.cpp:55:9)' lvalue Var {{.*}} '' '(lambda at {{.*}}union-kernel-param2.cpp:55:9)' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const void *' +// CHECK-NEXT: UnaryOperator {{.*}} '_generated_MyStructWithPtr *' prefix '&' cannot overflow +// CHECK-NEXT: DeclRefExpr {{.*}} '_generated_MyStructWithPtr' lvalue ParmVar {{.*}} '_arg_structWithPtr_mem' '_generated_MyStructWithPtr' +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 24 From dc09e07d5004fbffdacd8c2a4b96e18d0ef896e7 Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Thu, 15 Sep 2022 20:14:04 -0700 Subject: [PATCH 11/18] Fix dangling else Signed-off-by: Elizabeth Andrews --- clang/lib/Sema/SemaSYCL.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 45efb33281c47..dccc2b2c22807 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1812,7 +1812,7 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { } else if (PointerStack.pop_back_val()) { RecordDecl *RD = Ty->getAsRecordDecl(); assert(RD && "should not be null."); - if (!RD->hasAttr()) + if (!RD->hasAttr()) { // Do not generate a new type if the record corresponds to a // lambda. Currently the fields/bases of the local clone // corresponding to these generated types are intialized using @@ -1829,6 +1829,7 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { else RD->addAttr(SYCLRequiresDecompositionAttr::CreateImplicit( SemaRef.getASTContext())); + } PointerStack.back() = true; } return true; From 6e612b3354bfeecf1e8babaed17e41f8e68de7b7 Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Fri, 16 Sep 2022 19:56:02 -0700 Subject: [PATCH 12/18] Fix tests Signed-off-by: Elizabeth Andrews --- clang/test/CodeGenSYCL/inheritance.cpp | 45 +++++++++-------- .../CodeGenSYCL/no_opaque_inheritance.cpp | 48 +++++++++++-------- .../test/CodeGenSYCL/pointers-in-structs.cpp | 4 +- clang/test/SemaSYCL/inheritance.cpp | 6 +-- clang/test/SemaSYCL/union-kernel-param2.cpp | 4 +- 5 files changed, 62 insertions(+), 45 deletions(-) diff --git a/clang/test/CodeGenSYCL/inheritance.cpp b/clang/test/CodeGenSYCL/inheritance.cpp index dc46231dac670..11c71756d9768 100644 --- a/clang/test/CodeGenSYCL/inheritance.cpp +++ b/clang/test/CodeGenSYCL/inheritance.cpp @@ -39,31 +39,38 @@ int main() { return 0; } +// CHECK: %struct.base = type { i32, %class.InnerField } +// CHECK: %class.InnerField = type { %class.InnerFieldBase, i32 } +// CHECK: %class.InnerFieldBase = type { i32 } +// CHECK: %class._generated_second_base = type { ptr addrspace(1) } +// CHECK: %struct.derived = type <{ %struct.base, [4 x i8], %class.second_base, i32, [4 x i8] }> +// CHECK: %class.second_base = type { ptr addrspace(4) } + // Check kernel paramters -// CHECK: define {{.*}}spir_kernel void @{{.*}}derived(ptr noundef byval(%struct.base) align 4 %_arg__base, ptr noundef byval(%struct.__wrapper_class) align 8 %_arg_e, i32 noundef %_arg_a) +// CHECK: define {{.*}}spir_kernel void @{{.*}}derived +// CHECK-SAME: ptr noundef byval(%struct.base) align 4 %_arg__base +// CHECK-SAME: ptr noundef byval(%class._generated_second_base) align 8 %_arg__base1 +// CHECK-SAME: i32 noundef %_arg_a -// Check alloca for kernel paramters -// CHECK: %[[ARG_AA:[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 8 -// CHECK: %[[ARG_A:[a-zA-Z0-9_.]+]] = addrspacecast ptr %[[ARG_AA]] to ptr addrspace(4) -// CHECK: %[[BASE_TO_PTR:[a-zA-Z0-9_.]+]] = addrspacecast ptr %[[LOCAL_OBJECT]] to ptr addrspace(4) +// Check allocas for kernel parameters and local functor object +// CHECK: %[[ARG_A_ALLOCA:[a-zA-Z0-9_.]+]] = alloca i32, align 4 +// CHECK: %[[LOCAL_OBJECT_ALLOCA:[a-zA-Z0-9_.]+]] = alloca %struct.derived, align 8 +// CHECK: %[[ARG_A:[a-zA-Z0-9_.]+]] = addrspacecast ptr %[[ARG_A_ALLOCA]] to ptr addrspace(4) +// CHECK: %[[LOCAL_OBJECT:[a-zA-Z0-9_.]+]] = addrspacecast ptr %[[LOCAL_OBJECT_ALLOCA]] to ptr addrspace(4) +// CHECK: %[[ARG_BASE:[a-zA-Z0-9_.]+]] = addrspacecast ptr %_arg__base to ptr addrspace(4) +// CHECK: %[[ARG_BASE1:[a-zA-Z0-9_.]+]] = addrspacecast ptr %_arg__base1 to ptr addrspace(4) // CHECK: store i32 %_arg_a, ptr addrspace(4) %[[ARG_A]], align 4 // Initialize 'base' subobject -// CHECK: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 8 %[[BASE_TO_PTR]], ptr addrspace(4) align 4 %_arg__base.ascast, i64 12, i1 false) - -// Initialize 'second_base' subobject -// First, derived-to-base cast with offset: -// CHECK: %[[OFFSET_CALC:.*]] = getelementptr inbounds i8, ptr addrspace(4) %[[LOCAL_OBJECT]].ascast, i64 16 -// Initialize 'second_base::e' -// CHECK: %[[SECOND_BASE_PTR:.*]] = getelementptr inbounds %class.second_base, ptr addrspace(4) %[[OFFSET_CALC]], i32 0, i32 0 -// CHECK: %[[PTR_TO_WRAPPER:.*]] = getelementptr inbounds %struct.__wrapper_class, ptr addrspace(4) %_arg_e.ascast, i32 0, i32 0 -// CHECK: %[[LOAD_PTR:.*]] = load ptr addrspace(1), ptr addrspace(4) %[[PTR_TO_WRAPPER]] -// CHECK: %[[AS_CAST:.*]] = addrspacecast ptr addrspace(1) %[[LOAD_PTR]] to ptr addrspace(4) -// CHECK: store ptr addrspace(4) %[[AS_CAST]], ptr addrspace(4) %[[SECOND_BASE_PTR]] +// CHECK: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 8 %[[LOCAL_OBJECT]], ptr addrspace(4) align 4 %[[ARG_BASE]], i64 12, i1 false) // Initialize field 'a' -// CHECK: %[[GEP_A:[a-zA-Z0-9]+]] = getelementptr inbounds %struct.derived, ptr addrspace(4) %[[LOCAL_OBJECT]].ascast, i32 0, i32 3 +// CHECK: %[[GEP_A:[a-zA-Z0-9]+]] = getelementptr inbounds %struct.derived, ptr addrspace(4) %[[LOCAL_OBJECT]], i32 0, i32 3 // CHECK: %[[LOAD_A:[0-9]+]] = load i32, ptr addrspace(4) %[[ARG_A]], align 4 // CHECK: store i32 %[[LOAD_A]], ptr addrspace(4) %[[GEP_A]] + +// Initialize 'second_base' subobject +// First, derived-to-base cast with offset: +// CHECK: %[[OFFSET_CALC:.*]] = getelementptr inbounds i8, ptr addrspace(4) %[[LOCAL_OBJECT]], i64 16 +// Initialize 'second_base' +// CHECK: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 8 %[[OFFSET_CALC]], ptr addrspace(4) align 8 %[[ARG_BASE1]], i64 8, i1 false) diff --git a/clang/test/CodeGenSYCL/no_opaque_inheritance.cpp b/clang/test/CodeGenSYCL/no_opaque_inheritance.cpp index a513da161b878..4c20f958db294 100644 --- a/clang/test/CodeGenSYCL/no_opaque_inheritance.cpp +++ b/clang/test/CodeGenSYCL/no_opaque_inheritance.cpp @@ -39,35 +39,45 @@ int main() { return 0; } +// CHECK: %struct.base = type { i32, %class.InnerField } +// CHECK: %class.InnerField = type { %class.InnerFieldBase, i32 } +// CHECK: %class.InnerFieldBase = type { i32 } +// CHECK: %class._generated_second_base = type { i32 addrspace(1)* } +// CHECK: %struct.derived = type <{ %struct.base, [4 x i8], %class.second_base, i32, [4 x i8] }> +// CHECK: %class.second_base = type { i32 addrspace(4)* } + // Check kernel paramters -// CHECK: define {{.*}}spir_kernel void @{{.*}}derived(%struct.base* noundef byval(%struct.base) align 4 %_arg__base, %struct.__wrapper_class* noundef byval(%struct.__wrapper_class) align 8 %_arg_e, i32 noundef %_arg_a) +// CHECK: define {{.*}}spir_kernel void @{{.*}}derived +// CHECK-SAME: %struct.base* noundef byval(%struct.base) align 4 %_arg__base +// CHECK-SAME: %class._generated_second_base* noundef byval(%class._generated_second_base) align 8 %_arg__base1 +// CHECK-SAME: i32 noundef %_arg_a -// Check alloca for kernel paramters -// CHECK: %[[ARG_AA:[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 8 -// CHECK: %[[ARG_A:[a-zA-Z0-9_.]+]] = addrspacecast i32* %[[ARG_AA]] to i32 addrspace(4)* +// Check allocas for kernel parameters and local functor object +// CHECK: %[[ARG_A_ALLOCA:[a-zA-Z0-9_.]+]] = alloca i32, align 4 +// CHECK: %[[LOCAL_OBJECT_ALLOCA:[a-zA-Z0-9_.]+]] = alloca %struct.derived, align 8 +// CHECK: %[[ARG_A:[a-zA-Z0-9_.]+]] = addrspacecast i32* %[[ARG_A_ALLOCA]] to i32 addrspace(4)* +// CHECK: %[[LOCAL_OBJECT:[a-zA-Z0-9_.]+]] = addrspacecast %struct.derived* %[[LOCAL_OBJECT_ALLOCA]] to %struct.derived addrspace(4)* +// CHECK: %[[ARG_BASE:[a-zA-Z0-9_.]+]] = addrspacecast %struct.base* %_arg__base to %struct.base addrspace(4)* +// CHECK: %[[ARG_BASE1:[a-zA-Z0-9_.]+]] = addrspacecast %class._generated_second_base* %_arg__base1 to %class._generated_second_base addrspace(4)* // CHECK: store i32 %_arg_a, i32 addrspace(4)* %[[ARG_A]], align 4 // Initialize 'base' subobject -// CHECK: %[[DERIVED_TO_BASE:.*]] = bitcast %struct.derived addrspace(4)* %[[LOCAL_OBJECT]].ascast to %struct.base addrspace(4)* +// CHECK: %[[DERIVED_TO_BASE:.*]] = bitcast %struct.derived addrspace(4)* %[[LOCAL_OBJECT]] to %struct.base addrspace(4)* // CHECK: %[[BASE_TO_PTR:.*]] = bitcast %struct.base addrspace(4)* %[[DERIVED_TO_BASE]] to i8 addrspace(4)* -// CHECK: %[[PARAM_TO_PTR:.*]] = bitcast %struct.base addrspace(4)* %_arg__base.ascast to i8 addrspace(4)* +// CHECK: %[[PARAM_TO_PTR:.*]] = bitcast %struct.base addrspace(4)* %[[ARG_BASE]] to i8 addrspace(4)* // CHECK: call void @llvm.memcpy.p4i8.p4i8.i64(i8 addrspace(4)* align 8 %[[BASE_TO_PTR]], i8 addrspace(4)* align 4 %[[PARAM_TO_PTR]], i64 12, i1 false) +// Initialize field 'a' +// CHECK: %[[GEP_A:[a-zA-Z0-9]+]] = getelementptr inbounds %struct.derived, %struct.derived addrspace(4)* %[[LOCAL_OBJECT]], i32 0, i32 3 +// CHECK: %[[LOAD_A:[0-9]+]] = load i32, i32 addrspace(4)* %[[ARG_A]], align 4 +// CHECK: store i32 %[[LOAD_A]], i32 addrspace(4)* %[[GEP_A]] + // Initialize 'second_base' subobject // First, derived-to-base cast with offset: -// CHECK: %[[DERIVED_PTR:.*]] = bitcast %struct.derived addrspace(4)* %[[LOCAL_OBJECT]].ascast to i8 addrspace(4)* +// CHECK: %[[DERIVED_PTR:.*]] = bitcast %struct.derived addrspace(4)* %[[LOCAL_OBJECT]] to i8 addrspace(4)* // CHECK: %[[OFFSET_CALC:.*]] = getelementptr inbounds i8, i8 addrspace(4)* %[[DERIVED_PTR]], i64 16 // CHECK: %[[TO_SECOND_BASE:.*]] = bitcast i8 addrspace(4)* %[[OFFSET_CALC]] to %class.second_base addrspace(4)* -// Initialize 'second_base::e' -// CHECK: %[[SECOND_BASE_PTR:.*]] = getelementptr inbounds %class.second_base, %class.second_base addrspace(4)* %[[TO_SECOND_BASE]], i32 0, i32 0 -// CHECK: %[[PTR_TO_WRAPPER:.*]] = getelementptr inbounds %struct.__wrapper_class, %struct.__wrapper_class addrspace(4)* %_arg_e.ascast, i32 0, i32 0 -// CHECK: %[[LOAD_PTR:.*]] = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(4)* %[[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)* addrspace(4)* %[[SECOND_BASE_PTR]] +// CHECK: %[[SECOND_BASE_TO_PTR:.*]] = bitcast %class.second_base addrspace(4)* %[[TO_SECOND_BASE]] to i8 addrspace(4)* +// CHECK: %[[SECOND_PARAM_TO_PTR:.*]] = bitcast %class._generated_second_base addrspace(4)* %[[ARG_BASE1]] to i8 addrspace(4)* +// CHECK: call void @llvm.memcpy.p4i8.p4i8.i64(i8 addrspace(4)* align 8 %[[SECOND_BASE_TO_PTR]], i8 addrspace(4)* align 8 %[[SECOND_PARAM_TO_PTR]], i64 8, i1 false) -// Initialize field 'a' -// CHECK: %[[GEP_A:[a-zA-Z0-9]+]] = getelementptr inbounds %struct.derived, %struct.derived addrspace(4)* %[[LOCAL_OBJECT]].ascast, i32 0, i32 3 -// CHECK: %[[LOAD_A:[0-9]+]] = load i32, i32 addrspace(4)* %[[ARG_A]], align 4 -// CHECK: store i32 %[[LOAD_A]], i32 addrspace(4)* %[[GEP_A]] diff --git a/clang/test/CodeGenSYCL/pointers-in-structs.cpp b/clang/test/CodeGenSYCL/pointers-in-structs.cpp index c786cb7919725..486149c8af470 100644 --- a/clang/test/CodeGenSYCL/pointers-in-structs.cpp +++ b/clang/test/CodeGenSYCL/pointers-in-structs.cpp @@ -35,14 +35,14 @@ int main() { // CHECK: %[[WRAPPER_F1:[a-zA-Z0-9_.]+]] = type { ptr addrspace(1) } // CHECK: %[[WRAPPER_F2:[a-zA-Z0-9_.]+]] = type { ptr addrspace(1) } -// CHECK: %[[WRAPPER_F:[a-zA-Z0-9_.]+]] = type { ptr addrspace(1) } +// CHECK: %[[GENERATED_A:[a-zA-Z0-9_.]+]] = type { ptr addrspace(1) } // CHECK: %[[WRAPPER_F4_1:[a-zA-Z0-9_.]+]] = type { ptr addrspace(1) } // CHECK: %[[WRAPPER_F4_2:[a-zA-Z0-9_.]+]] = type { ptr addrspace(1) } // CHECK: %[[WRAPPER_LAMBDA_PTR:[a-zA-Z0-9_.]+]] = type { ptr addrspace(1) } // CHECK: define {{.*}}spir_kernel void @{{.*}}structs // CHECK-SAME: ptr noundef byval(%[[WRAPPER_F1]]) align 8 %_arg_F1, // CHECK-SAME: ptr noundef byval(%[[WRAPPER_F2]]) align 8 %_arg_F2, -// CHECK-SAME: ptr noundef byval(%[[WRAPPER_F]]) align 8 %_arg_F, +// CHECK-SAME: ptr noundef byval(%[[GENERATED_A]]) align 8 %_arg_F3, // CHECK-SAME: ptr noundef byval(%[[WRAPPER_F4_1]]) align 8 %_arg_F4 // CHECK-SAME: ptr noundef byval(%[[WRAPPER_F4_2]]) align 8 %_arg_F41 // CHECK: define {{.*}}spir_kernel void @{{.*}}lambdas{{.*}}(ptr noundef byval(%[[WRAPPER_LAMBDA_PTR]]) align 8 %_arg_Ptr) diff --git a/clang/test/SemaSYCL/inheritance.cpp b/clang/test/SemaSYCL/inheritance.cpp index 5044c4170e83f..b6e067afd8b75 100644 --- a/clang/test/SemaSYCL/inheritance.cpp +++ b/clang/test/SemaSYCL/inheritance.cpp @@ -67,14 +67,14 @@ int main() { // base is a simple class with no corresponding generated type. Therefore // copy from ParamVar -// CHECK-NEXT: CXXConstructExpr {{.*}} 'base' 'void (const base &) noexcept' +// CHECK-NEXT: CXXConstructExpr {{.*}} 'base':'base' 'void (const base &) noexcept' // CHECK-NEXT: ImplicitCastExpr {{.*}} 'const base' lvalue // CHECK-NEXT: DeclRefExpr {{.*}} lvalue ParmVar {{.*}} '_arg__base' 'base' // second_base contains pointers and therefore the ParamVar is a new generated // type. Default construct this class and initialize second_base via memcpy in // body statements. -// CHECK-NEXT: CXXConstructExpr {{.*}} 'second_base' 'void () noexcept' +// CHECK-NEXT: CXXConstructExpr {{.*}} 'second_base':'second_base' 'void () noexcept' // third_base contains special type accessor. Therefore it is decomposed and it's // data members are copied from corresponding ParamVar @@ -107,5 +107,5 @@ int main() { // CHECK: CXXMemberCallExpr // CHECK-NEXT: MemberExpr {{.*}} lvalue .__init // CHECK-NEXT: MemberExpr {{.*}} lvalue .AccField -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'third_base' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'third_base':'third_base' lvalue // CHECK-NEXT: DeclRefExpr {{.*}} 'derived' lvalue Var {{.*}} 'derived' 'derived' diff --git a/clang/test/SemaSYCL/union-kernel-param2.cpp b/clang/test/SemaSYCL/union-kernel-param2.cpp index 1bb0aab19d785..3d3a8105fb2d2 100644 --- a/clang/test/SemaSYCL/union-kernel-param2.cpp +++ b/clang/test/SemaSYCL/union-kernel-param2.cpp @@ -95,7 +95,7 @@ int main() { // CHECK-NEXT: MemberExpr {{.*}} lvalue .__init // CHECK-NEXT: MemberExpr {{.*}} lvalue .AccField // CHECK-NEXT: MemberExpr {{.*}} lvalue .struct_mem -// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}union-kernel-param2.cpp:48:9)' lvalue Var {{.*}} '' '(lambda at {{.*}}union-kernel-param2.cpp:48:9)' +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}union-kernel-param2.cpp:48:9)' lvalue Var {{.*}} '__SYCLKernel' '(lambda at {{.*}}union-kernel-param2.cpp:48:9)' // Check kernel_C parameters // CHECK: FunctionDecl {{.*}}kernel_C{{.*}} 'void (_generated_MyStructWithPtr)' @@ -115,7 +115,7 @@ int main() { // CHECK-NEXT: ImplicitCastExpr {{.*}} 'void *' // CHECK-NEXT: UnaryOperator {{.*}} 'struct MyStructWithPtr *' prefix '&' cannot overflow // CHECK-NEXT: MemberExpr {{.*}} 'struct MyStructWithPtr':'MyStructWithPtr' lvalue .structWithPtr_mem -// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}union-kernel-param2.cpp:55:9)' lvalue Var {{.*}} '' '(lambda at {{.*}}union-kernel-param2.cpp:55:9)' +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}union-kernel-param2.cpp:55:9)' lvalue Var {{.*}} '__SYCLKernel' '(lambda at {{.*}}union-kernel-param2.cpp:55:9)' // CHECK-NEXT: ImplicitCastExpr {{.*}} 'const void *' // CHECK-NEXT: UnaryOperator {{.*}} '_generated_MyStructWithPtr *' prefix '&' cannot overflow // CHECK-NEXT: DeclRefExpr {{.*}} '_generated_MyStructWithPtr' lvalue ParmVar {{.*}} '_arg_structWithPtr_mem' '_generated_MyStructWithPtr' From dd907e49563c348e8f6750520e18a58bcd8b97c8 Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Sun, 18 Sep 2022 12:00:55 -0700 Subject: [PATCH 13/18] Do not generate new type for non-trivial types Signed-off-by: Elizabeth Andrews --- clang/lib/Sema/SemaSYCL.cpp | 39 ++++++++++++++++++++++--------------- 1 file changed, 23 insertions(+), 16 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index dccc2b2c22807..72641f1e2bd15 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1801,29 +1801,26 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { // SYCLRequiresDecompositionAttr. Else if a record contains // a pointer, it is marked with SYCLGenerateNewTypeAttr. A record // will never be marked with both attributes. + CXXRecordDecl *RD = Ty->getAsCXXRecordDecl(); + assert(RD && "should not be null."); if (CollectionStack.pop_back_val()) { - RecordDecl *RD = Ty->getAsRecordDecl(); - assert(RD && "should not be null."); if (!RD->hasAttr()) RD->addAttr(SYCLRequiresDecompositionAttr::CreateImplicit( SemaRef.getASTContext())); CollectionStack.back() = true; PointerStack.pop_back(); } else if (PointerStack.pop_back_val()) { - RecordDecl *RD = Ty->getAsRecordDecl(); - assert(RD && "should not be null."); if (!RD->hasAttr()) { - // Do not generate a new type if the record corresponds to a - // lambda. Currently the fields/bases of the local clone + // Do not generate a new type if the record is not default + // constructible. Currently the fields/bases of the local clone // corresponding to these generated types are intialized using // their default constructors(Actual initialization is done via // memcpy in kernel body.) to maintain the integrity of the // InitListExpr we generate for Kernel Object local clone. - // Records correspondng to lambdas which have captures do not have - // a default constructor and so current logic fails for lambdas. - // FIXME: Can/Should we stop triggering decomposition for lambdas - // with pointers? - if (!RD->isLambda()) + // So current logic fails for types without default constructors. + // FIXME: Stop triggering decomposition for non-trivial types with + // pointers + if (RD->isTrivial()) RD->addAttr( SYCLGenerateNewTypeAttr::CreateImplicit(SemaRef.getASTContext())); else @@ -1848,20 +1845,30 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { // SYCLRequiresDecompositionAttr. Else if a record contains // a pointer, it is marked with SYCLGenerateNewTypeAttr. A record // will never be marked with both attributes. + CXXRecordDecl *RD = Ty->getAsCXXRecordDecl(); + assert(RD && "should not be null."); if (CollectionStack.pop_back_val()) { - RecordDecl *RD = Ty->getAsRecordDecl(); - assert(RD && "should not be null."); if (!RD->hasAttr()) RD->addAttr(SYCLRequiresDecompositionAttr::CreateImplicit( SemaRef.getASTContext())); CollectionStack.back() = true; PointerStack.pop_back(); } else if (PointerStack.pop_back_val()) { - RecordDecl *RD = Ty->getAsRecordDecl(); - assert(RD && "should not be null."); - if (!RD->hasAttr()) + // Do not generate a new type if the record is not default + // constructible. Currently the fields/bases of the local clone + // corresponding to these generated types are intialized using + // their default constructors(Actual initialization is done via + // memcpy in kernel body.) to maintain the integrity of the + // InitListExpr we generate for Kernel Object local clone. + // So current logic fails for types without default constructors. + // FIXME: Stop triggering decomposition for non-trivial types with + // pointers + if (RD->isTrivial()) RD->addAttr( SYCLGenerateNewTypeAttr::CreateImplicit(SemaRef.getASTContext())); + else + RD->addAttr(SYCLRequiresDecompositionAttr::CreateImplicit( + SemaRef.getASTContext())); PointerStack.back() = true; } return true; From 324cc47db21c7020592bbb0febe3ff09d969ea8b Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Mon, 19 Sep 2022 14:02:53 -0700 Subject: [PATCH 14/18] Handle opt-report generation correctly and add a few tests Signed-off-by: Elizabeth Andrews --- clang/lib/Sema/SemaSYCL.cpp | 66 +++++++------ clang/test/SemaSYCL/decomposition.cpp | 26 ++++- clang/test/SemaSYCL/kernel-arg-opt-report.cpp | 99 ++++++++++++++----- 3 files changed, 133 insertions(+), 58 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 72641f1e2bd15..c8fdad4f7fd61 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -2061,15 +2061,11 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { CXXRecordDecl *ModifiedRD = ModifiedRecords.pop_back_val(); ModifiedRD->completeDefinition(); - // Elizabeth - Check if this will work if kernel functor has generated - // base class. if (!ModifiedBases.empty()) ModifiedRD->setBases(ModifiedBases.data(), ModifiedBases.size()); return QualType(ModifiedRD->getTypeForDecl(), 0); } - - // Elizabeth - Need to handle KernelHandler }; // A type to Create and own the FunctionDecl for the kernel. @@ -2553,7 +2549,7 @@ class SyclKernelArgsSizeChecker : public SyclKernelFieldHandler { }; std::string getKernelArgDesc(StringRef KernelArgDescription) { - if (KernelArgDescription == ":" || KernelArgDescription == "") + if (KernelArgDescription == "") return ""; return ("Compiler generated argument for " + KernelArgDescription + ",") .str(); @@ -2564,27 +2560,20 @@ class SyclOptReportCreator : public SyclKernelFieldHandler { SourceLocation KernelInvocationLoc; void addParam(const FieldDecl *KernelArg, QualType KernelArgType, - StringRef KernelArgDescription) { + StringRef KernelArgDescription, + bool IsCompilerGeneratedType = false) { StringRef NameToEmitInDescription = KernelArg->getName(); const RecordDecl *KernelArgParent = KernelArg->getParent(); if (KernelArgParent && KernelArgDescription == "decomposed struct/class") NameToEmitInDescription = KernelArgParent->getName(); - bool isWrappedField = KernelArgDescription == "WrappedPointer" || - KernelArgDescription == "WrappedArray"; - - KernelArgDescription = - (KernelArgDescription == "WrappedPointer" - ? "nested pointer" - : (KernelArgDescription == "WrappedArray" ? "array" - : KernelArgDescription)); - unsigned KernelArgSize = SemaRef.getASTContext().getTypeSizeInChars(KernelArgType).getQuantity(); SemaRef.getDiagnostics().getSYCLOptReport().AddKernelArgs( DC.getKernelDecl(), NameToEmitInDescription, - isWrappedField ? "Compiler generated" : KernelArgType.getAsString(), + IsCompilerGeneratedType ? "Compiler generated" + : KernelArgType.getAsString(), KernelInvocationLoc, KernelArgSize, getKernelArgDesc(KernelArgDescription), (KernelArgDescription == "decomposed struct/class") @@ -2593,10 +2582,8 @@ class SyclOptReportCreator : public SyclKernelFieldHandler { } void addParam(const FieldDecl *FD, QualType FieldTy) { - std::string KernelArgDescription = FieldTy.getAsString(); + std::string KernelArgDescription = ""; const RecordDecl *RD = FD->getParent(); - if (FieldTy->isScalarType()) - KernelArgDescription = ""; if (RD && RD->hasAttr()) KernelArgDescription = "decomposed struct/class"; @@ -2605,12 +2592,15 @@ class SyclOptReportCreator : public SyclKernelFieldHandler { // Handles base classes. void addParam(const CXXBaseSpecifier &, QualType KernelArgType, - StringRef KernelArgDescription) { + StringRef KernelArgDescription, + bool IsCompilerGeneratedType = false) { unsigned KernelArgSize = SemaRef.getASTContext().getTypeSizeInChars(KernelArgType).getQuantity(); SemaRef.getDiagnostics().getSYCLOptReport().AddKernelArgs( DC.getKernelDecl(), KernelArgType.getAsString(), - KernelArgType.getAsString(), KernelInvocationLoc, KernelArgSize, + IsCompilerGeneratedType ? "Compiler generated" + : KernelArgType.getAsString(), + KernelInvocationLoc, KernelArgSize, getKernelArgDesc(KernelArgDescription), ""); } @@ -2652,15 +2642,20 @@ class SyclOptReportCreator : public SyclKernelFieldHandler { } bool handlePointerType(FieldDecl *FD, QualType FieldTy) final { - std::string KernelArgDescription = ":"; + std::string KernelArgDescription = ""; + bool IsCompilerGeneratedType = false; ParmVarDecl *KernelParameter = DC.getParamVarDeclsForCurrentField()[0]; // Compiler generated openCL kernel argument for current pointer field // is not a pointer. This means we are processing a nested pointer and // the openCL kernel argument is of type __wrapper_class. - if (!KernelParameter->getType()->isPointerType()) - KernelArgDescription = "WrappedPointer"; + if (!KernelParameter->getType()->isPointerType()) { + KernelArgDescription = "nested pointer"; + IsCompilerGeneratedType = true; + } + for (const auto *Param : DC.getParamVarDeclsForCurrentField()) - addParam(FD, Param->getType(), KernelArgDescription); + addParam(FD, Param->getType(), KernelArgDescription, + /*IsCompilerGeneratedType*/ IsCompilerGeneratedType); return true; } @@ -2672,19 +2667,30 @@ class SyclOptReportCreator : public SyclKernelFieldHandler { bool handleSimpleArrayType(FieldDecl *FD, QualType FieldTy) final { // Simple arrays are always wrapped. for (const auto *Param : DC.getParamVarDeclsForCurrentField()) - addParam(FD, Param->getType(), "WrappedArray"); + addParam(FD, Param->getType(), "array", /*IsCompilerGeneratedType*/ true); return true; } bool handleNonDecompStruct(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final { - addParam(FD, Ty); + CXXRecordDecl *RD = Ty->getAsCXXRecordDecl(); + assert(RD && "Type must be a C++ record type"); + if (RD->hasAttr()) + addParam(FD, Ty, "object with pointer", /*IsCompilerGeneratedType*/ true); + else + addParam(FD, Ty); return true; } - bool handleNonDecompStruct(const CXXRecordDecl *Base, - const CXXBaseSpecifier &BS, QualType Ty) final { - addParam(BS, Ty, "base class"); + bool handleNonDecompStruct(const CXXRecordDecl *, const CXXBaseSpecifier &BS, + QualType Ty) final { + CXXRecordDecl *RD = Ty->getAsCXXRecordDecl(); + assert(RD && "Type must be a C++ record type"); + if (RD->hasAttr()) + addParam(BS, Ty, "base class with pointer", + /*IsCompilerGeneratedType*/ true); + else + addParam(BS, Ty, "base class"); return true; } diff --git a/clang/test/SemaSYCL/decomposition.cpp b/clang/test/SemaSYCL/decomposition.cpp index 2dd4599865e7f..7245f58347444 100644 --- a/clang/test/SemaSYCL/decomposition.cpp +++ b/clang/test/SemaSYCL/decomposition.cpp @@ -10,6 +10,7 @@ sycl::queue myQueue; struct StructWithAccessor { sycl::accessor acc; + int *ptr; }; struct StructInheritedAccessor : sycl::accessor { @@ -46,6 +47,12 @@ struct StructWithNonDecomposedStruct : StructNonDecomposed { double d; }; +struct StructWithPtr { + StructNonDecomposed member; + int *ptr; + int i; +}; + template struct StructWithArray { T a; @@ -66,6 +73,8 @@ int main() { StructNonDecomposed ArrayOfSimpleStruct[5]; StructWithNonDecomposedStruct NonDecompStruct; StructWithNonDecomposedStruct ArrayOfNonDecompStruct[5]; + StructWithPtr SimpleStructWithPtr; + // Check to ensure that these are not decomposed. myQueue.submit([&](sycl::handler &h) { h.single_task([=]() { return SimpleStruct.i + ArrayOfSimpleStruct[0].i + NonDecompStruct.i + ArrayOfNonDecompStruct[0].i; }); @@ -77,13 +86,13 @@ int main() { myQueue.submit([&](sycl::handler &h) { h.single_task([=]() { return t1.i; }); }); - // CHECK: FunctionDecl {{.*}}Acc1{{.*}} 'void (__global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, StructNonDecomposed, int)' + // CHECK: FunctionDecl {{.*}}Acc1{{.*}} 'void (__global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, __wrapper_class, __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, __wrapper_class, __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, __wrapper_class, StructNonDecomposed, int)' DerivedStruct t2; myQueue.submit([&](sycl::handler &h) { h.single_task([=]() { return t2.i; }); }); - // CHECK: FunctionDecl {{.*}}Acc2{{.*}} 'void (__global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, StructNonDecomposed, int)' + // CHECK: FunctionDecl {{.*}}Acc2{{.*}} 'void (__global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, __wrapper_class, StructNonDecomposed, int)' StructWithArray t3; myQueue.submit([&](sycl::handler &h) { @@ -152,4 +161,17 @@ int main() { }); // CHECK: FunctionDecl {{.*}}Half2{{.*}} 'void (DerivedStruct)' } + + { + myQueue.submit([&](sycl::handler &h) { + h.single_task([=]() { return SimpleStructWithPtr.i; }); + }); + // CHECK: FunctionDecl {{.*}}Pointer{{.*}} 'void (_generated_StructWithPtr)' + + DerivedStruct t1; + myQueue.submit([&](sycl::handler &h) { + h.single_task([=]() { return t1.i; }); + }); + // CHECK: FunctionDecl {{.*}}Pointer{{.*}} 'void (_generated_DerivedStruct)' + } } diff --git a/clang/test/SemaSYCL/kernel-arg-opt-report.cpp b/clang/test/SemaSYCL/kernel-arg-opt-report.cpp index df38453b5cb09..a9cca45099b62 100644 --- a/clang/test/SemaSYCL/kernel-arg-opt-report.cpp +++ b/clang/test/SemaSYCL/kernel-arg-opt-report.cpp @@ -25,11 +25,17 @@ struct NotDecomposedBase { int B; }; -struct KernelFunctor : NotDecomposedBase, DecomposedBase { +struct StructWithPointer { +public: + int *Ptr; +}; + +struct KernelFunctor : NotDecomposedBase, DecomposedBase, StructWithPointer { int A; int *Ptr; int Array[3]; sycl::sampler Sampl; + StructWithPointer Obj; void operator()() const { } }; @@ -63,7 +69,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 28, Column: 8 } +// SPIR-NEXT: Line: 33, Column: 8 } // SPIR-NEXT: Function: _ZTS13KernelFunctor // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -84,7 +90,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 28, Column: 8 } +// SPIR-NEXT: Line: 33, Column: 8 } // SPIR-NEXT: Function: _ZTS13KernelFunctor // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -105,7 +111,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 28, Column: 8 } +// SPIR-NEXT: Line: 33, Column: 8 } // SPIR-NEXT: Function: _ZTS13KernelFunctor // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -126,7 +132,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 28, Column: 8 } +// SPIR-NEXT: Line: 33, Column: 8 } // SPIR-NEXT: Function: _ZTS13KernelFunctor // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -147,7 +153,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 28, Column: 8 } +// SPIR-NEXT: Line: 33, Column: 8 } // SPIR-NEXT: Function: _ZTS13KernelFunctor // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -168,7 +174,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 28, Column: 8 } +// SPIR-NEXT: Line: 33, Column: 8 } // SPIR-NEXT: Function: _ZTS13KernelFunctor // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -189,7 +195,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 28, Column: 8 } +// SPIR-NEXT: Line: 33, Column: 8 } // SPIR-NEXT: Function: _ZTS13KernelFunctor // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -210,7 +216,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 28, Column: 8 } +// SPIR-NEXT: Line: 33, Column: 8 } // SPIR-NEXT: Function: _ZTS13KernelFunctor // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -231,7 +237,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 28, Column: 8 } +// SPIR-NEXT: Line: 33, Column: 8 } // SPIR-NEXT: Function: _ZTS13KernelFunctor // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -252,7 +258,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 28, Column: 8 } +// SPIR-NEXT: Line: 33, Column: 8 } // SPIR-NEXT: Function: _ZTS13KernelFunctor // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -273,7 +279,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 28, Column: 8 } +// SPIR-NEXT: Line: 33, Column: 8 } // SPIR-NEXT: Function: _ZTS13KernelFunctor // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -294,7 +300,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 28, Column: 8 } +// SPIR-NEXT: Line: 33, Column: 8 } // SPIR-NEXT: Function: _ZTS13KernelFunctor // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -315,12 +321,33 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 28, Column: 8 } +// SPIR-NEXT: Line: 33, Column: 8 } // SPIR-NEXT: Function: _ZTS13KernelFunctor // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' // SPIR-NEXT: Argument: '12' // SPIR-NEXT: String: ':' +// SPIR-NEXT: String: Compiler generated argument for base class with pointer, +// SPIR-NEXT: String: StructWithPointer +// SPIR-NEXT: String: ' (' +// SPIR-NEXT: String: '' +// SPIR-NEXT: String: 'Type:' +// SPIR-NEXT: String: Compiler generated +// SPIR-NEXT: String: ', ' +// SPIR-NEXT: String: 'Size: ' +// SPIR-NEXT: Argument: '8' +// SPIR-NEXT: String: ')' + +// SPIR: --- !Passed +// SPIR: Pass:{{.*}}sycl +// SPIR: Name:{{.*}}Region +// SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// SPIR-NEXT: Line: 33, Column: 8 } +// SPIR-NEXT: Function: _ZTS13KernelFunctor +// SPIR-NEXT: Args: +// SPIR-NEXT: String: 'Arg ' +// SPIR-NEXT: Argument: '13' +// SPIR-NEXT: String: ':' // SPIR-NEXT: String: '' // SPIR-NEXT: String: A // SPIR-NEXT: String: ' (' @@ -336,11 +363,11 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 28, Column: 8 } +// SPIR-NEXT: Line: 33, Column: 8 } // SPIR-NEXT: Function: _ZTS13KernelFunctor // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' -// SPIR-NEXT: Argument: '13' +// SPIR-NEXT: Argument: '14' // SPIR-NEXT: String: ':' // SPIR-NEXT: String: '' // SPIR-NEXT: String: Ptr @@ -357,11 +384,11 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 28, Column: 8 } +// SPIR-NEXT: Line: 33, Column: 8 } // SPIR-NEXT: Function: _ZTS13KernelFunctor // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' -// SPIR-NEXT: Argument: '14' +// SPIR-NEXT: Argument: '15' // SPIR-NEXT: String: ':' // SPIR-NEXT: String: Compiler generated argument for array, // SPIR-NEXT: String: Array @@ -378,11 +405,11 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 28, Column: 8 } +// SPIR-NEXT: Line: 33, Column: 8 } // SPIR-NEXT: Function: _ZTS13KernelFunctor // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' -// SPIR-NEXT: Argument: '15' +// SPIR-NEXT: Argument: '16' // SPIR-NEXT: String: ':' // SPIR-NEXT: String: 'Compiler generated argument for sycl::sampler,' // SPIR-NEXT: String: Sampl @@ -395,13 +422,33 @@ int main() { // SPIR-NEXT: Argument: '8' // SPIR-NEXT: String: ')' +// SPIR: --- !Passed +// SPIR: Pass:{{.*}}sycl +// SPIR: Name:{{.*}}Region +// SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// SPIR-NEXT: Line: 33, Column: 8 } +// SPIR-NEXT: Function: _ZTS13KernelFunctor +// SPIR-NEXT: Args: +// SPIR-NEXT: String: 'Arg ' +// SPIR-NEXT: Argument: '17' +// SPIR-NEXT: String: ':' +// SPIR-NEXT: String: Compiler generated argument for object with pointer, +// SPIR-NEXT: String: Obj +// SPIR-NEXT: String: ' (' +// SPIR-NEXT: String: '' +// SPIR-NEXT: String: 'Type:' +// SPIR-NEXT: String: Compiler generated +// SPIR-NEXT: String: ', ' +// SPIR-NEXT: String: 'Size: ' +// SPIR-NEXT: Argument: '8' +// SPIR-NEXT: String: ')' // Output for kernel XYZ // SPIR: --- !Passed // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 53, Column: 9 } +// SPIR-NEXT: Line: 59, Column: 9 } // SPIR-NEXT: Function: _ZTSZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_E3XYZ // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -422,7 +469,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 53, Column: 9 } +// SPIR-NEXT: Line: 59, Column: 9 } // SPIR-NEXT: Function: _ZTSZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_E3XYZ // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -443,7 +490,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 53, Column: 9 } +// SPIR-NEXT: Line: 59, Column: 9 } // SPIR-NEXT: Function: _ZTSZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_E3XYZ // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -464,7 +511,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 53, Column: 9 } +// SPIR-NEXT: Line: 59, Column: 9 } // SPIR-NEXT: Function: _ZTSZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_E3XYZ // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -485,7 +532,7 @@ int main() { // SPIR: Pass:{{.*}}sycl // SPIR: Name:{{.*}}Region // SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// SPIR-NEXT: Line: 53, Column: 9 } +// SPIR-NEXT: Line: 59, Column: 9 } // SPIR-NEXT: Function: _ZTSZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_E3XYZ // SPIR-NEXT: Args: // SPIR-NEXT: String: 'Arg ' @@ -506,7 +553,7 @@ int main() { // NVPTX: Pass:{{.*}}sycl // NVPTX: Name:{{.*}}Region // NVPTX: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// NVPTX: Line: 53, Column: 9 } +// NVPTX: Line: 59, Column: 9 } // NVPTX-NEXT: Function: _ZTSZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_E3XYZ // NVPTX-NEXT: Args: // NVPTX-NEXT: String: 'Arg ' From f03b49e164a931cbdc5fcd48265f21c2741cf33e Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Tue, 20 Sep 2022 09:01:56 -0700 Subject: [PATCH 15/18] Implement review comments Signed-off-by: Elizabeth Andrews --- clang/include/clang/Basic/Attr.td | 2 +- clang/lib/Sema/SemaSYCL.cpp | 21 ++++++++------------- clang/test/SemaSYCL/decomposition.cpp | 17 ++++++++++++++++- 3 files changed, 25 insertions(+), 15 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 9b7fb98b6d0a3..7ed985c795b48 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1434,7 +1434,7 @@ def SYCLGenerateNewType : InheritableAttr { let Spellings = []; let Subjects = SubjectList<[Named]>; let LangOpts = [SYCLIsDevice, SYCLIsHost]; - let Documentation = [Undocumented]; + let Documentation = [InternalOnly]; } def SYCLIntelKernelArgsRestrict : InheritableAttr { diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index c8fdad4f7fd61..96aedf22e48c7 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1971,18 +1971,15 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { static constexpr const bool VisitInsideSimpleContainersWithPointer = true; SyclKernelPointerHandler(Sema &S, const CXXRecordDecl *RD) : SyclKernelFieldHandler(S) { - // Generate new type createNewType(RD); } bool enterStruct(const CXXRecordDecl *, FieldDecl *, QualType Ty) final { - // Generate new type createNewType(Ty->getAsCXXRecordDecl()); return true; } bool leaveStruct(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final { - // Get Decl of generated new type CXXRecordDecl *ModifiedRD = getGeneratedNewRecord(Ty->getAsCXXRecordDecl()); // Add this record as a field of it's parent record. @@ -1993,14 +1990,12 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { bool enterStruct(const CXXRecordDecl *, const CXXBaseSpecifier &, QualType Ty) final { - // Generate new type createNewType(Ty->getAsCXXRecordDecl()); return true; } bool leaveStruct(const CXXRecordDecl *Parent, const CXXBaseSpecifier &BS, QualType Ty) final { - // Get Decl of generated new type CXXRecordDecl *ModifiedRD = getGeneratedNewRecord(Ty->getAsCXXRecordDecl()); // Create CXXBaseSpecifier for this generated class. @@ -2011,9 +2006,9 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { bool handlePointerType(FieldDecl *FD, QualType FieldTy) final { QualType PointeeTy = FieldTy->getPointeeType(); Qualifiers Quals = PointeeTy.getQualifiers(); - auto AS = Quals.getAddressSpace(); + LangAS AS = Quals.getAddressSpace(); // Leave global_device and global_host address spaces as is to help FPGA - // device in memory allocations + // device in memory allocations. if (!PointeeTy->isFunctionType() && AS != LangAS::sycl_global_device && AS != LangAS::sycl_global_host) Quals.setAddressSpace(LangAS::sycl_global); @@ -2655,7 +2650,7 @@ class SyclOptReportCreator : public SyclKernelFieldHandler { for (const auto *Param : DC.getParamVarDeclsForCurrentField()) addParam(FD, Param->getType(), KernelArgDescription, - /*IsCompilerGeneratedType*/ IsCompilerGeneratedType); + IsCompilerGeneratedType); return true; } @@ -2976,7 +2971,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { return Cast; } - Expr *getAddressOf(Expr *E) { + Expr *createGetAddressOf(Expr *E) { return UnaryOperator::Create(SemaRef.Context, E, UO_AddrOf, SemaRef.Context.getPointerType(E->getType()), VK_PRValue, OK_Ordinary, KernelCallerSrcLoc, @@ -3021,8 +3016,8 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { addFieldInit(FD, Ty, None, InitializationKind::CreateDefault(KernelCallerSrcLoc)); addFieldMemberExpr(FD, Ty); - Expr *ParamRef = getAddressOf(createParamReferenceExpr()); - Expr *LocalCloneRef = getAddressOf(MemberExprBases.back()); + Expr *ParamRef = createGetAddressOf(createParamReferenceExpr()); + Expr *LocalCloneRef = createGetAddressOf(MemberExprBases.back()); Expr *MemCpyCallExpr = buildMemCpyCall(ParamRef, LocalCloneRef, Ty); BodyStmts.push_back(MemCpyCallExpr); removeFieldMemberExpr(FD, Ty); @@ -3034,8 +3029,8 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { void handleGeneratedType(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS, QualType Ty) { addBaseInit(BS, Ty, InitializationKind::CreateDefault(KernelCallerSrcLoc)); - Expr *ParamRef = getAddressOf(createParamReferenceExpr()); - Expr *LocalCloneRef = getAddressOf(MemberExprBases.back()); + Expr *ParamRef = createGetAddressOf(createParamReferenceExpr()); + Expr *LocalCloneRef = createGetAddressOf(MemberExprBases.back()); LocalCloneRef = addDerivedToBaseCastExpr(RD, BS, LocalCloneRef); Expr *MemCpyCallExpr = buildMemCpyCall(ParamRef, LocalCloneRef, Ty); BodyStmts.push_back(MemCpyCallExpr); diff --git a/clang/test/SemaSYCL/decomposition.cpp b/clang/test/SemaSYCL/decomposition.cpp index 7245f58347444..8cd1fbf93f338 100644 --- a/clang/test/SemaSYCL/decomposition.cpp +++ b/clang/test/SemaSYCL/decomposition.cpp @@ -67,6 +67,12 @@ struct DerivedStruct : T { int i; }; +struct NonTrivialType { + int *Ptr; + int i; + NonTrivialType(int i){} +}; + int main() { StructNonDecomposed SimpleStruct; @@ -74,6 +80,7 @@ int main() { StructWithNonDecomposedStruct NonDecompStruct; StructWithNonDecomposedStruct ArrayOfNonDecompStruct[5]; StructWithPtr SimpleStructWithPtr; + NonTrivialType NonTrivialStructWithPtr(10); // Check to ensure that these are not decomposed. myQueue.submit([&](sycl::handler &h) { @@ -172,6 +179,14 @@ int main() { myQueue.submit([&](sycl::handler &h) { h.single_task([=]() { return t1.i; }); }); - // CHECK: FunctionDecl {{.*}}Pointer{{.*}} 'void (_generated_DerivedStruct)' + // CHECK: FunctionDecl {{.*}}PointerInBase{{.*}} 'void (_generated_DerivedStruct)' + } + + { + // FIXME: Stop decomposition for non-trivial types with pointers. + myQueue.submit([&](sycl::handler &h) { + h.single_task([=]() { return NonTrivialStructWithPtr.i;}); + }); + // CHECK: FunctionDecl {{.*}}NonTrivial{{.*}} 'void (__wrapper_class, int)' } } From 449deece8becb436156a2858cf78fe2e540b8aed Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Tue, 20 Sep 2022 12:02:19 -0700 Subject: [PATCH 16/18] Implement review comments Signed-off-by: Elizabeth Andrews --- clang/lib/Sema/SemaSYCL.cpp | 13 +++++-------- 1 file changed, 5 insertions(+), 8 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 96aedf22e48c7..54484740c41f6 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1458,7 +1458,7 @@ void KernelObjVisitor::visitRecord(const CXXRecordDecl *Owner, ParentTy &Parent, Value) { // We are currently in PointerHandler visitor. if (RD->hasAttr()) { - // This is record containing pointers. + // This is a record containing pointers. visitComplexRecord(Owner, Parent, Wrapper, RecordTy, Handlers...); } else { // This is a record without pointers. @@ -1503,7 +1503,7 @@ void KernelObjVisitor::visitArray(const CXXRecordDecl *Owner, FieldDecl *Field, // 'simple' array i.e. one that does not include special types or pointers. // Array of pointers/ array of type containing pointers will be handled in // a follow-up PR. Currently, they continue to trigger decomposition, and - // will be handled in 'if' statment above. + // will be handled in 'if' statement above. visitSimpleArray(Owner, Field, ArrayTy, Handlers...); } else { if (!AllTrue::Value) @@ -1745,10 +1745,7 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { llvm::SmallVector CollectionStack; llvm::SmallVector PointerStack; - // FIXME: Array of pointers/ array of type containing pointers - // will be handled in a follow up PR. Currently, they continue - // to trigger decomposition. - // TODO: Remove this method once arrays are handled correctly + // TODO: Remove this method once arrays are handled correctly. bool isArrayElement(const FieldDecl *FD, QualType Ty) const { return !SemaRef.getASTContext().hasSameType(FD->getType(), Ty); } @@ -1819,7 +1816,7 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { // InitListExpr we generate for Kernel Object local clone. // So current logic fails for types without default constructors. // FIXME: Stop triggering decomposition for non-trivial types with - // pointers + // pointers. if (RD->isTrivial()) RD->addAttr( SYCLGenerateNewTypeAttr::CreateImplicit(SemaRef.getASTContext())); @@ -1862,7 +1859,7 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { // InitListExpr we generate for Kernel Object local clone. // So current logic fails for types without default constructors. // FIXME: Stop triggering decomposition for non-trivial types with - // pointers + // pointers. if (RD->isTrivial()) RD->addAttr( SYCLGenerateNewTypeAttr::CreateImplicit(SemaRef.getASTContext())); From babd36cebb58b6c6a23190e30855a051522a1680 Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Wed, 21 Sep 2022 13:47:37 -0700 Subject: [PATCH 17/18] Fix crash on array of type containing pointers. Implement review comments Signed-off-by: Elizabeth Andrews --- clang/lib/Sema/SemaSYCL.cpp | 67 +++--- clang/test/CodeGenSYCL/inheritance.cpp | 4 +- .../CodeGenSYCL/no_opaque_inheritance.cpp | 8 +- clang/test/SemaSYCL/array-kernel-param.cpp | 194 +++++++++++++++++- clang/test/SemaSYCL/decomposition.cpp | 4 +- clang/test/SemaSYCL/inheritance.cpp | 8 +- clang/test/SemaSYCL/union-kernel-param2.cpp | 8 +- 7 files changed, 229 insertions(+), 64 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 54484740c41f6..d74d859edc393 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1777,13 +1777,7 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { } bool handlePointerType(FieldDecl *FD, QualType Ty) final { - // FIXME: Array of pointers/ array of type containing pointers - // will be handled in a follow up PR. Currently, they continue - // to trigger decomposition. - if (isArrayElement(FD, Ty)) - CollectionStack.back() = true; - else - PointerStack.back() = true; + PointerStack.back() = true; return true; } @@ -1793,7 +1787,7 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { return true; } - bool leaveStruct(const CXXRecordDecl *, FieldDecl *, QualType Ty) final { + bool leaveStruct(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final { // If a record needs to be decomposed, it is marked with // SYCLRequiresDecompositionAttr. Else if a record contains // a pointer, it is marked with SYCLGenerateNewTypeAttr. A record @@ -1807,24 +1801,20 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { CollectionStack.back() = true; PointerStack.pop_back(); } else if (PointerStack.pop_back_val()) { - if (!RD->hasAttr()) { - // Do not generate a new type if the record is not default - // constructible. Currently the fields/bases of the local clone - // corresponding to these generated types are intialized using - // their default constructors(Actual initialization is done via - // memcpy in kernel body.) to maintain the integrity of the - // InitListExpr we generate for Kernel Object local clone. - // So current logic fails for types without default constructors. - // FIXME: Stop triggering decomposition for non-trivial types with - // pointers. - if (RD->isTrivial()) + // FIXME: Stop triggering decomposition for non-trivial types with + // pointers + if (RD->isTrivial()) { + PointerStack.back() = true; + if (!RD->hasAttr()) RD->addAttr( SYCLGenerateNewTypeAttr::CreateImplicit(SemaRef.getASTContext())); - else + } else { + // We are visiting a non-trivial type with pointer. + CollectionStack.back() = true; + if (!RD->hasAttr()) RD->addAttr(SYCLRequiresDecompositionAttr::CreateImplicit( SemaRef.getASTContext())); } - PointerStack.back() = true; } return true; } @@ -1851,22 +1841,20 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { CollectionStack.back() = true; PointerStack.pop_back(); } else if (PointerStack.pop_back_val()) { - // Do not generate a new type if the record is not default - // constructible. Currently the fields/bases of the local clone - // corresponding to these generated types are intialized using - // their default constructors(Actual initialization is done via - // memcpy in kernel body.) to maintain the integrity of the - // InitListExpr we generate for Kernel Object local clone. - // So current logic fails for types without default constructors. // FIXME: Stop triggering decomposition for non-trivial types with - // pointers. - if (RD->isTrivial()) - RD->addAttr( - SYCLGenerateNewTypeAttr::CreateImplicit(SemaRef.getASTContext())); - else - RD->addAttr(SYCLRequiresDecompositionAttr::CreateImplicit( - SemaRef.getASTContext())); - PointerStack.back() = true; + // pointers + if (RD->isTrivial()) { + PointerStack.back() = true; + if (!RD->hasAttr()) + RD->addAttr( + SYCLGenerateNewTypeAttr::CreateImplicit(SemaRef.getASTContext())); + } else { + // We are visiting a non-trivial type with pointer. + CollectionStack.back() = true; + if (!RD->hasAttr()) + RD->addAttr(SYCLRequiresDecompositionAttr::CreateImplicit( + SemaRef.getASTContext())); + } } return true; } @@ -1890,7 +1878,10 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { // FIXME: Array of pointers/ array of type containing pointers // will be handled in a follow up PR. Currently, they continue // to trigger decomposition. - llvm_unreachable("PointerStack should not be true when handling arrays."); + if (!FD->hasAttr()) + FD->addAttr(SYCLRequiresDecompositionAttr::CreateImplicit( + SemaRef.getASTContext())); + CollectionStack.back() = true; } return true; } @@ -1904,7 +1895,7 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { IdentifierInfo *getModifiedName(IdentifierInfo *Id) { std::string Name = - Id ? (Twine("_generated_") + Id->getName()).str() : "_generated_"; + Id ? (Twine("__generated_") + Id->getName()).str() : "__generated_"; return &SemaRef.getASTContext().Idents.get(Name); } diff --git a/clang/test/CodeGenSYCL/inheritance.cpp b/clang/test/CodeGenSYCL/inheritance.cpp index 11c71756d9768..883ae6c0b8087 100644 --- a/clang/test/CodeGenSYCL/inheritance.cpp +++ b/clang/test/CodeGenSYCL/inheritance.cpp @@ -42,14 +42,14 @@ int main() { // CHECK: %struct.base = type { i32, %class.InnerField } // CHECK: %class.InnerField = type { %class.InnerFieldBase, i32 } // CHECK: %class.InnerFieldBase = type { i32 } -// CHECK: %class._generated_second_base = type { ptr addrspace(1) } +// CHECK: %class.__generated_second_base = type { ptr addrspace(1) } // CHECK: %struct.derived = type <{ %struct.base, [4 x i8], %class.second_base, i32, [4 x i8] }> // CHECK: %class.second_base = type { ptr addrspace(4) } // Check kernel paramters // CHECK: define {{.*}}spir_kernel void @{{.*}}derived // CHECK-SAME: ptr noundef byval(%struct.base) align 4 %_arg__base -// CHECK-SAME: ptr noundef byval(%class._generated_second_base) align 8 %_arg__base1 +// CHECK-SAME: ptr noundef byval(%class.__generated_second_base) align 8 %_arg__base1 // CHECK-SAME: i32 noundef %_arg_a // Check allocas for kernel parameters and local functor object diff --git a/clang/test/CodeGenSYCL/no_opaque_inheritance.cpp b/clang/test/CodeGenSYCL/no_opaque_inheritance.cpp index 4c20f958db294..b58390d30443f 100644 --- a/clang/test/CodeGenSYCL/no_opaque_inheritance.cpp +++ b/clang/test/CodeGenSYCL/no_opaque_inheritance.cpp @@ -42,14 +42,14 @@ int main() { // CHECK: %struct.base = type { i32, %class.InnerField } // CHECK: %class.InnerField = type { %class.InnerFieldBase, i32 } // CHECK: %class.InnerFieldBase = type { i32 } -// CHECK: %class._generated_second_base = type { i32 addrspace(1)* } +// CHECK: %class.__generated_second_base = type { i32 addrspace(1)* } // CHECK: %struct.derived = type <{ %struct.base, [4 x i8], %class.second_base, i32, [4 x i8] }> // CHECK: %class.second_base = type { i32 addrspace(4)* } // Check kernel paramters // CHECK: define {{.*}}spir_kernel void @{{.*}}derived // CHECK-SAME: %struct.base* noundef byval(%struct.base) align 4 %_arg__base -// CHECK-SAME: %class._generated_second_base* noundef byval(%class._generated_second_base) align 8 %_arg__base1 +// CHECK-SAME: %class.__generated_second_base* noundef byval(%class.__generated_second_base) align 8 %_arg__base1 // CHECK-SAME: i32 noundef %_arg_a // Check allocas for kernel parameters and local functor object @@ -58,7 +58,7 @@ int main() { // CHECK: %[[ARG_A:[a-zA-Z0-9_.]+]] = addrspacecast i32* %[[ARG_A_ALLOCA]] to i32 addrspace(4)* // CHECK: %[[LOCAL_OBJECT:[a-zA-Z0-9_.]+]] = addrspacecast %struct.derived* %[[LOCAL_OBJECT_ALLOCA]] to %struct.derived addrspace(4)* // CHECK: %[[ARG_BASE:[a-zA-Z0-9_.]+]] = addrspacecast %struct.base* %_arg__base to %struct.base addrspace(4)* -// CHECK: %[[ARG_BASE1:[a-zA-Z0-9_.]+]] = addrspacecast %class._generated_second_base* %_arg__base1 to %class._generated_second_base addrspace(4)* +// CHECK: %[[ARG_BASE1:[a-zA-Z0-9_.]+]] = addrspacecast %class.__generated_second_base* %_arg__base1 to %class.__generated_second_base addrspace(4)* // CHECK: store i32 %_arg_a, i32 addrspace(4)* %[[ARG_A]], align 4 // Initialize 'base' subobject @@ -78,6 +78,6 @@ int main() { // CHECK: %[[OFFSET_CALC:.*]] = getelementptr inbounds i8, i8 addrspace(4)* %[[DERIVED_PTR]], i64 16 // CHECK: %[[TO_SECOND_BASE:.*]] = bitcast i8 addrspace(4)* %[[OFFSET_CALC]] to %class.second_base addrspace(4)* // CHECK: %[[SECOND_BASE_TO_PTR:.*]] = bitcast %class.second_base addrspace(4)* %[[TO_SECOND_BASE]] to i8 addrspace(4)* -// CHECK: %[[SECOND_PARAM_TO_PTR:.*]] = bitcast %class._generated_second_base addrspace(4)* %[[ARG_BASE1]] to i8 addrspace(4)* +// CHECK: %[[SECOND_PARAM_TO_PTR:.*]] = bitcast %class.__generated_second_base addrspace(4)* %[[ARG_BASE1]] to i8 addrspace(4)* // CHECK: call void @llvm.memcpy.p4i8.p4i8.i64(i8 addrspace(4)* align 8 %[[SECOND_BASE_TO_PTR]], i8 addrspace(4)* align 8 %[[SECOND_PARAM_TO_PTR]], i64 8, i1 false) diff --git a/clang/test/SemaSYCL/array-kernel-param.cpp b/clang/test/SemaSYCL/array-kernel-param.cpp index 4069c7dd79d09..02f2d720b05ee 100644 --- a/clang/test/SemaSYCL/array-kernel-param.cpp +++ b/clang/test/SemaSYCL/array-kernel-param.cpp @@ -29,7 +29,7 @@ int main() { S s; - struct StructWithPointers { + struct StructWithArrayOfPointers { int x; int y; int *ArrayOfPtrs[2]; @@ -37,7 +37,7 @@ int main() { struct DecomposedStruct { int a; - StructWithPointers SWPtrsMem[2]; + StructWithArrayOfPointers SWPtrsMem[2]; int *Array_2D_Ptrs[2][1]; int c; }; @@ -49,8 +49,19 @@ int main() { int c; }; + struct StructWithSimplePointer { + int *Ptr; + int a; + }; + + struct StructWithNestedPointer { + StructWithSimplePointer SWPointer[2]; + }; + DecomposedStruct DecompStructArray[2]; NonDecomposedStruct NonDecompStructArray[2]; + StructWithSimplePointer StructWithSimplePointerArray[2]; + StructWithNestedPointer StructWithNestedPointerArray[2]; int array_2D[2][3]; @@ -109,6 +120,20 @@ int main() { NonDecomposedStruct local = NonDecompStructArray[0]; }); }); + + myQueue.submit([&](sycl::handler &h) { + h.single_task( + [=] { + StructWithSimplePointer local = StructWithSimplePointerArray[0]; + }); + }); + + myQueue.submit([&](sycl::handler &h) { + h.single_task( + [=] { + StructWithNestedPointer local = StructWithNestedPointerArray[0]; + }); + }); } // Check Kernel_Accessor parameters @@ -230,10 +255,10 @@ int main() { // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_a' 'int' -// Initializer for struct array inside DecomposedStruct i.e. StructWithPointers SWPtrsMem[2] -// CHECK-NEXT: InitListExpr {{.*}} 'StructWithPointers[2]' +// Initializer for struct array inside DecomposedStruct i.e. StructWithArrayOfPointers SWPtrsMem[2] +// CHECK-NEXT: InitListExpr {{.*}} 'StructWithArrayOfPointers[2]' // Initializer for first element of inner struct array -// CHECK-NEXT: InitListExpr {{.*}} 'StructWithPointers' +// CHECK-NEXT: InitListExpr {{.*}} 'StructWithArrayOfPointers' // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_x' 'int' // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' @@ -248,7 +273,7 @@ int main() { // CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . // CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPtrs' '__wrapper_class' // Initializer for second element of inner struct array -// CHECK-NEXT: InitListExpr {{.*}} 'StructWithPointers' +// CHECK-NEXT: InitListExpr {{.*}} 'StructWithArrayOfPointers' // CHECK-NEXT: ImplicitCastExpr // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_x' 'int' // CHECK-NEXT: ImplicitCastExpr @@ -281,10 +306,10 @@ int main() { // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_a' 'int' -// Initializer for struct array inside DecomposedStruct i.e. StructWithPointers SWPtrsMem[2] -// CHECK-NEXT: InitListExpr {{.*}} 'StructWithPointers[2]' +// Initializer for struct array inside DecomposedStruct i.e. StructWithArrayOfPointers SWPtrsMem[2] +// CHECK-NEXT: InitListExpr {{.*}} 'StructWithArrayOfPointers[2]' // Initializer for first element of inner struct array -// CHECK-NEXT: InitListExpr {{.*}} 'StructWithPointers' +// CHECK-NEXT: InitListExpr {{.*}} 'StructWithArrayOfPointers' // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_x' 'int' // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' @@ -299,7 +324,7 @@ int main() { // CHECK-NEXT: MemberExpr {{.*}} '__global int *' lvalue . // CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ArrayOfPtrs' '__wrapper_class' // Initializer for second element of inner struct array -// CHECK-NEXT: InitListExpr {{.*}} 'StructWithPointers' +// CHECK-NEXT: InitListExpr {{.*}} 'StructWithArrayOfPointers' // CHECK-NEXT: ImplicitCastExpr // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_x' 'int' // CHECK-NEXT: ImplicitCastExpr @@ -388,3 +413,152 @@ int main() { // CHECK-NEXT: MemberExpr {{.*}} 'NonDecomposedStruct[2]' lvalue . // CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_NonDecompStructArray' '__wrapper_class' // CHECK-NEXT: ArrayInitIndexExpr {{.*}} 'unsigned + +// Check Kernel_StructWithSimplePointer parameters. +// CHECK: FunctionDecl {{.*}}Kernel_StructWithSimplePointer{{.*}} 'void (__generated_StructWithSimplePointer, __generated_StructWithSimplePointer)' +// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_StructWithSimplePointerArray '__generated_StructWithSimplePointer' +// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_StructWithSimplePointerArray '__generated_StructWithSimplePointer' +// Check Kernel_StructWithSimplePointer inits +// CHECK-NEXT: CompoundStmt +// CHECK-NEXT: DeclStmt +// CHECK-NEXT: VarDecl {{.*}} cinit +// CHECK-NEXT: InitListExpr +// CHECK-NEXT: InitListExpr {{.*}} 'StructWithSimplePointer[2]' +// Default Initialize array elements +// CHECK-NEXT: CXXConstructExpr {{.*}} 'StructWithSimplePointer':'StructWithSimplePointer' 'void () noexcept' +// CHECK-NEXT: CXXConstructExpr {{.*}} 'StructWithSimplePointer':'StructWithSimplePointer' 'void () noexcept' + +// Memcopy first array element +// CHECK-NEXT: CallExpr +// CHECK-NEXT: ImplicitCastExpr {{.*}} +// CHECK-NEXT: DeclRefExpr {{.*}} '' {{.*}} '__builtin_memcpy' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'void *' +// CHECK-NEXT: UnaryOperator {{.*}} 'StructWithSimplePointer *' prefix '&' cannot overflow +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'StructWithSimplePointer':'StructWithSimplePointer' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'StructWithSimplePointer *' +// CHECK-NEXT: MemberExpr {{.*}} 'StructWithSimplePointer[2]' lvalue .StructWithSimplePointerArray +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}array-kernel-param.cpp:126:9)' lvalue Var +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 0 +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const void *' +// CHECK-NEXT: UnaryOperator {{.*}} '__generated_StructWithSimplePointer *' prefix '&' cannot overflow +// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_StructWithSimplePointer' lvalue ParmVar {{.*}} '_arg_StructWithSimplePointerArray' '__generated_StructWithSimplePointer' +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 16 + +// Memcopy second array element +// CHECK-NEXT: CallExpr +// CHECK-NEXT: ImplicitCastExpr {{.*}} +// CHECK-NEXT: DeclRefExpr {{.*}} '' {{.*}} '__builtin_memcpy' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'void *' +// CHECK-NEXT: UnaryOperator {{.*}} 'StructWithSimplePointer *' prefix '&' cannot overflow +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'StructWithSimplePointer':'StructWithSimplePointer' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'StructWithSimplePointer *' +// CHECK-NEXT: MemberExpr {{.*}} 'StructWithSimplePointer[2]' lvalue .StructWithSimplePointerArray +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}array-kernel-param.cpp:126:9)' lvalue Var +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 1 +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const void *' +// CHECK-NEXT: UnaryOperator {{.*}} '__generated_StructWithSimplePointer *' prefix '&' cannot overflow +// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_StructWithSimplePointer' lvalue ParmVar {{.*}} '_arg_StructWithSimplePointerArray' '__generated_StructWithSimplePointer' +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 16 + +// Check Kernel_StructWithNestedPointer parameters. +// CHECK: FunctionDecl {{.*}}Kernel_StructWithNestedPointer{{.*}} 'void (__generated_StructWithSimplePointer, __generated_StructWithSimplePointer, __generated_StructWithSimplePointer, __generated_StructWithSimplePointer)' +// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_SWPointer '__generated_StructWithSimplePointer' +// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_SWPointer '__generated_StructWithSimplePointer' +// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_SWPointer '__generated_StructWithSimplePointer' +// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_SWPointer '__generated_StructWithSimplePointer' +// Check Kernel_StructWithNestedPointer inits +// CHECK-NEXT: CompoundStmt +// CHECK-NEXT: DeclStmt +// CHECK-NEXT: VarDecl {{.*}} cinit +// CHECK-NEXT: InitListExpr +// CHECK-NEXT: InitListExpr {{.*}} 'StructWithNestedPointer[2]' +// InitListExpr for first array element of StructWithNestedPointerArray +// CHECK-NEXT: InitListExpr {{.*}} 'StructWithNestedPointer' +// CHECK-NEXT: InitListExpr {{.*}} 'StructWithSimplePointer[2]' +// CHECK-NEXT: CXXConstructExpr {{.*}} 'StructWithSimplePointer':'StructWithSimplePointer' 'void () noexcept' +// CHECK-NEXT: CXXConstructExpr {{.*}} 'StructWithSimplePointer':'StructWithSimplePointer' 'void () noexcept' +// InitListExpr for second array element of StructWithNestedPointerArray +// CHECK-NEXT: InitListExpr {{.*}} 'StructWithNestedPointer' +// CHECK-NEXT: InitListExpr {{.*}} 'StructWithSimplePointer[2]' +// CHECK-NEXT: CXXConstructExpr {{.*}} 'StructWithSimplePointer':'StructWithSimplePointer' 'void () noexcept' +// CHECK-NEXT: CXXConstructExpr {{.*}} 'StructWithSimplePointer':'StructWithSimplePointer' 'void () noexcept' + +// Memcpy StructWithNestedPointerArray[0].SWPointer[0] +// CHECK-NEXT: CallExpr {{.*}} 'void *' +// CHECK-NEXT: ImplicitCastExpr {{.*}} +// CHECK-NEXT: DeclRefExpr {{.*}} '' {{.*}} '__builtin_memcpy' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'void *' +// CHECK-NEXT: UnaryOperator {{.*}} 'StructWithSimplePointer *' prefix '&' cannot overflow +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'StructWithSimplePointer':'StructWithSimplePointer' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'StructWithSimplePointer *' +// CHECK-NEXT: MemberExpr {{.*}} 'StructWithSimplePointer[2]' lvalue .SWPointer +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'StructWithNestedPointer':'StructWithNestedPointer' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'StructWithNestedPointer *' +// CHECK-NEXT: MemberExpr {{.*}} 'StructWithNestedPointer[2]' lvalue .StructWithNestedPointerArray +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}array-kernel-param.cpp:133:9)' +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 0 +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 0 +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const void *' +// CHECK-NEXT: UnaryOperator {{.*}} '__generated_StructWithSimplePointer *' prefix '&' cannot overflow +// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_StructWithSimplePointer' lvalue ParmVar {{.*}} '_arg_SWPointer' '__generated_StructWithSimplePointer' +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 16 + +// Memcpy StructWithNestedPointerArray[0].SWPointer[1] +// CHECK-NEXT: CallExpr {{.*}} 'void *' +// CHECK-NEXT: ImplicitCastExpr {{.*}} +// CHECK-NEXT: DeclRefExpr {{.*}} '' {{.*}} '__builtin_memcpy' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'void *' +// CHECK-NEXT: UnaryOperator {{.*}} 'StructWithSimplePointer *' prefix '&' cannot overflow +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'StructWithSimplePointer':'StructWithSimplePointer' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'StructWithSimplePointer *' +// CHECK-NEXT: MemberExpr {{.*}} 'StructWithSimplePointer[2]' lvalue .SWPointer +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'StructWithNestedPointer':'StructWithNestedPointer' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'StructWithNestedPointer *' +// CHECK-NEXT: MemberExpr {{.*}} 'StructWithNestedPointer[2]' lvalue .StructWithNestedPointerArray +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}array-kernel-param.cpp:133:9)' +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 0 +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 1 +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const void *' +// CHECK-NEXT: UnaryOperator {{.*}} '__generated_StructWithSimplePointer *' prefix '&' cannot overflow +// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_StructWithSimplePointer' lvalue ParmVar {{.*}} '_arg_SWPointer' '__generated_StructWithSimplePointer' +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 16 + +// Memcpy StructWithNestedPointerArray[1].SWPointer[0] +// CHECK-NEXT: CallExpr {{.*}} 'void *' +// CHECK-NEXT: ImplicitCastExpr {{.*}} +// CHECK-NEXT: DeclRefExpr {{.*}} '' {{.*}} '__builtin_memcpy' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'void *' +// CHECK-NEXT: UnaryOperator {{.*}} 'StructWithSimplePointer *' prefix '&' cannot overflow +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'StructWithSimplePointer':'StructWithSimplePointer' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'StructWithSimplePointer *' +// CHECK-NEXT: MemberExpr {{.*}} 'StructWithSimplePointer[2]' lvalue .SWPointer +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'StructWithNestedPointer':'StructWithNestedPointer' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'StructWithNestedPointer *' +// CHECK-NEXT: MemberExpr {{.*}} 'StructWithNestedPointer[2]' lvalue .StructWithNestedPointerArray +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}array-kernel-param.cpp:133:9)' +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 1 +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 0 +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const void *' +// CHECK-NEXT: UnaryOperator {{.*}} '__generated_StructWithSimplePointer *' prefix '&' cannot overflow +// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_StructWithSimplePointer' lvalue ParmVar {{.*}} '_arg_SWPointer' '__generated_StructWithSimplePointer' +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 16 + +// Memcpy StructWithNestedPointerArray[1].SWPointer[1] +// CHECK-NEXT: CallExpr {{.*}} 'void *' +// CHECK-NEXT: ImplicitCastExpr {{.*}} +// CHECK-NEXT: DeclRefExpr {{.*}} '' {{.*}} '__builtin_memcpy' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'void *' +// CHECK-NEXT: UnaryOperator {{.*}} 'StructWithSimplePointer *' prefix '&' cannot overflow +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'StructWithSimplePointer':'StructWithSimplePointer' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'StructWithSimplePointer *' +// CHECK-NEXT: MemberExpr {{.*}} 'StructWithSimplePointer[2]' lvalue .SWPointer +// CHECK-NEXT: ArraySubscriptExpr {{.*}} 'StructWithNestedPointer':'StructWithNestedPointer' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'StructWithNestedPointer *' +// CHECK-NEXT: MemberExpr {{.*}} 'StructWithNestedPointer[2]' lvalue .StructWithNestedPointerArray +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}array-kernel-param.cpp:133:9)' +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 1 +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 0 +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const void *' +// CHECK-NEXT: UnaryOperator {{.*}} '__generated_StructWithSimplePointer *' prefix '&' cannot overflow +// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_StructWithSimplePointer' lvalue ParmVar {{.*}} '_arg_SWPointer' '__generated_StructWithSimplePointer' +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 16 diff --git a/clang/test/SemaSYCL/decomposition.cpp b/clang/test/SemaSYCL/decomposition.cpp index 8cd1fbf93f338..6ca217837d809 100644 --- a/clang/test/SemaSYCL/decomposition.cpp +++ b/clang/test/SemaSYCL/decomposition.cpp @@ -173,13 +173,13 @@ int main() { myQueue.submit([&](sycl::handler &h) { h.single_task([=]() { return SimpleStructWithPtr.i; }); }); - // CHECK: FunctionDecl {{.*}}Pointer{{.*}} 'void (_generated_StructWithPtr)' + // CHECK: FunctionDecl {{.*}}Pointer{{.*}} 'void (__generated_StructWithPtr)' DerivedStruct t1; myQueue.submit([&](sycl::handler &h) { h.single_task([=]() { return t1.i; }); }); - // CHECK: FunctionDecl {{.*}}PointerInBase{{.*}} 'void (_generated_DerivedStruct)' + // CHECK: FunctionDecl {{.*}}PointerInBase{{.*}} 'void (__generated_DerivedStruct)' } { diff --git a/clang/test/SemaSYCL/inheritance.cpp b/clang/test/SemaSYCL/inheritance.cpp index b6e067afd8b75..eedcd58e9d07b 100644 --- a/clang/test/SemaSYCL/inheritance.cpp +++ b/clang/test/SemaSYCL/inheritance.cpp @@ -46,12 +46,12 @@ int main() { } // Check declaration of the kernel -// CHECK: derived{{.*}} 'void (base, _generated_second_base, __wrapper_class, +// CHECK: derived{{.*}} 'void (base, __generated_second_base, __wrapper_class, // CHECK-SAME: __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, int) // Check parameters of the kernel // CHECK: ParmVarDecl {{.*}} used _arg__base 'base' -// CHECK: ParmVarDecl {{.*}} used _arg__base '_generated_second_base' +// CHECK: ParmVarDecl {{.*}} used _arg__base '__generated_second_base' // CHECK: ParmVarDecl {{.*}} used _arg_d '__wrapper_class' // CHECK: ParmVarDecl {{.*}} used _arg_AccField '__global char *' // CHECK: ParmVarDecl {{.*}} used _arg_AccField 'sycl::range<1>' @@ -99,8 +99,8 @@ int main() { // CHECK-NEXT: UnaryOperator {{.*}} 'derived *' prefix '&' cannot overflow // CHECK-NEXT: DeclRefExpr {{.*}} 'derived' lvalue Var {{.*}} 'derived' 'derived' // CHECK-NEXT: ImplicitCastExpr {{.*}} 'const void *' -// CHECK-NEXT: UnaryOperator {{.*}} '_generated_second_base *' prefix '&' cannot overflow -// CHECK-NEXT: DeclRefExpr {{.*}} '_generated_second_base' lvalue ParmVar {{.*}} '_arg__base' '_generated_second_base' +// CHECK-NEXT: UnaryOperator {{.*}} '__generated_second_base *' prefix '&' cannot overflow +// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_second_base' lvalue ParmVar {{.*}} '_arg__base' '__generated_second_base' // CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 8 // Check kernel body for call to __init function of accessor diff --git a/clang/test/SemaSYCL/union-kernel-param2.cpp b/clang/test/SemaSYCL/union-kernel-param2.cpp index 3d3a8105fb2d2..f55db4426e494 100644 --- a/clang/test/SemaSYCL/union-kernel-param2.cpp +++ b/clang/test/SemaSYCL/union-kernel-param2.cpp @@ -98,8 +98,8 @@ int main() { // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}union-kernel-param2.cpp:48:9)' lvalue Var {{.*}} '__SYCLKernel' '(lambda at {{.*}}union-kernel-param2.cpp:48:9)' // Check kernel_C parameters -// CHECK: FunctionDecl {{.*}}kernel_C{{.*}} 'void (_generated_MyStructWithPtr)' -// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_structWithPtr_mem '_generated_MyStructWithPtr' +// CHECK: FunctionDecl {{.*}}kernel_C{{.*}} 'void (__generated_MyStructWithPtr)' +// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_structWithPtr_mem '__generated_MyStructWithPtr' // Check kernel_C inits // CHECK-NEXT: CompoundStmt @@ -117,6 +117,6 @@ int main() { // CHECK-NEXT: MemberExpr {{.*}} 'struct MyStructWithPtr':'MyStructWithPtr' lvalue .structWithPtr_mem // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}union-kernel-param2.cpp:55:9)' lvalue Var {{.*}} '__SYCLKernel' '(lambda at {{.*}}union-kernel-param2.cpp:55:9)' // CHECK-NEXT: ImplicitCastExpr {{.*}} 'const void *' -// CHECK-NEXT: UnaryOperator {{.*}} '_generated_MyStructWithPtr *' prefix '&' cannot overflow -// CHECK-NEXT: DeclRefExpr {{.*}} '_generated_MyStructWithPtr' lvalue ParmVar {{.*}} '_arg_structWithPtr_mem' '_generated_MyStructWithPtr' +// CHECK-NEXT: UnaryOperator {{.*}} '__generated_MyStructWithPtr *' prefix '&' cannot overflow +// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_MyStructWithPtr' lvalue ParmVar {{.*}} '_arg_structWithPtr_mem' '__generated_MyStructWithPtr' // CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 24 From e8cea2ff164ae9d2229230c7e5dee4df19a9e292 Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Thu, 22 Sep 2022 14:15:44 -0700 Subject: [PATCH 18/18] Fix test. Minor NFC. Add test Signed-off-by: Elizabeth Andrews --- clang/lib/Sema/SemaSYCL.cpp | 11 ++---- clang/test/SemaSYCL/array-kernel-param.cpp | 2 +- clang/test/SemaSYCL/decomposition.cpp | 45 +++++++++++++++++----- 3 files changed, 39 insertions(+), 19 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index d74d859edc393..853fb568efcf8 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1745,11 +1745,6 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { llvm::SmallVector CollectionStack; llvm::SmallVector PointerStack; - // TODO: Remove this method once arrays are handled correctly. - bool isArrayElement(const FieldDecl *FD, QualType Ty) const { - return !SemaRef.getASTContext().hasSameType(FD->getType(), Ty); - } - public: static constexpr const bool VisitUnionBody = false; static constexpr const bool VisitNthArrayElement = false; @@ -1776,7 +1771,7 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { return true; } - bool handlePointerType(FieldDecl *FD, QualType Ty) final { + bool handlePointerType(FieldDecl *, QualType) final { PointerStack.back() = true; return true; } @@ -1787,7 +1782,7 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { return true; } - bool leaveStruct(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final { + bool leaveStruct(const CXXRecordDecl *, FieldDecl *, QualType Ty) final { // If a record needs to be decomposed, it is marked with // SYCLRequiresDecompositionAttr. Else if a record contains // a pointer, it is marked with SYCLGenerateNewTypeAttr. A record @@ -2998,7 +2993,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { } // Adds default initializer for generated type and creates - // a call to __builtin_memcpy to iniatilize local clone from + // a call to __builtin_memcpy to initialize local clone from // kernel argument. void handleGeneratedType(FieldDecl *FD, QualType Ty) { addFieldInit(FD, Ty, None, diff --git a/clang/test/SemaSYCL/array-kernel-param.cpp b/clang/test/SemaSYCL/array-kernel-param.cpp index 02f2d720b05ee..cd9a8e3355498 100644 --- a/clang/test/SemaSYCL/array-kernel-param.cpp +++ b/clang/test/SemaSYCL/array-kernel-param.cpp @@ -557,7 +557,7 @@ int main() { // CHECK-NEXT: MemberExpr {{.*}} 'StructWithNestedPointer[2]' lvalue .StructWithNestedPointerArray // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}array-kernel-param.cpp:133:9)' // CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 1 -// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 0 +// CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned long' 1 // CHECK-NEXT: ImplicitCastExpr {{.*}} 'const void *' // CHECK-NEXT: UnaryOperator {{.*}} '__generated_StructWithSimplePointer *' prefix '&' cannot overflow // CHECK-NEXT: DeclRefExpr {{.*}} '__generated_StructWithSimplePointer' lvalue ParmVar {{.*}} '_arg_SWPointer' '__generated_StructWithSimplePointer' diff --git a/clang/test/SemaSYCL/decomposition.cpp b/clang/test/SemaSYCL/decomposition.cpp index 6ca217837d809..8c79ab582ad9d 100644 --- a/clang/test/SemaSYCL/decomposition.cpp +++ b/clang/test/SemaSYCL/decomposition.cpp @@ -53,6 +53,17 @@ struct StructWithPtr { int i; }; +struct NonTrivialType { + int *Ptr; + int i; + NonTrivialType(int i){} +}; + +struct NonTrivialDerived : NonTrivialType { + int a; + NonTrivialDerived(int i) : NonTrivialType(i) {} +}; + template struct StructWithArray { T a; @@ -67,20 +78,12 @@ struct DerivedStruct : T { int i; }; -struct NonTrivialType { - int *Ptr; - int i; - NonTrivialType(int i){} -}; - int main() { StructNonDecomposed SimpleStruct; StructNonDecomposed ArrayOfSimpleStruct[5]; StructWithNonDecomposedStruct NonDecompStruct; StructWithNonDecomposedStruct ArrayOfNonDecompStruct[5]; - StructWithPtr SimpleStructWithPtr; - NonTrivialType NonTrivialStructWithPtr(10); // Check to ensure that these are not decomposed. myQueue.submit([&](sycl::handler &h) { @@ -170,23 +173,45 @@ int main() { } { + StructWithPtr SimpleStructWithPtr; myQueue.submit([&](sycl::handler &h) { h.single_task([=]() { return SimpleStructWithPtr.i; }); }); // CHECK: FunctionDecl {{.*}}Pointer{{.*}} 'void (__generated_StructWithPtr)' - DerivedStruct t1; + // FIXME: Stop decomposition of arrays with pointers + StructWithArray t1; + myQueue.submit([&](sycl::handler &h) { + h.single_task([=]() { return t1.i; }); + }); + // CHECK: FunctionDecl {{.*}}NestedArrayOfStructWithPointer{{.*}} 'void (__generated_StructWithPtr, __generated_StructWithPtr, __generated_StructWithPtr, StructNonDecomposed, int)' + + DerivedStruct t2; myQueue.submit([&](sycl::handler &h) { - h.single_task([=]() { return t1.i; }); + h.single_task([=]() { return t2.i; }); }); // CHECK: FunctionDecl {{.*}}PointerInBase{{.*}} 'void (__generated_DerivedStruct)' } { // FIXME: Stop decomposition for non-trivial types with pointers. + + NonTrivialType NonTrivialStructWithPtr(10); myQueue.submit([&](sycl::handler &h) { h.single_task([=]() { return NonTrivialStructWithPtr.i;}); }); // CHECK: FunctionDecl {{.*}}NonTrivial{{.*}} 'void (__wrapper_class, int)' + + NonTrivialType NonTrivialTypeArray[2]{0,0}; + myQueue.submit([&](sycl::handler &h) { + h.single_task([=]() { return NonTrivialTypeArray[0].i;}); + }); + // CHECK: FunctionDecl {{.*}}ArrayOfNonTrivialStruct{{.*}} 'void (__wrapper_class, int, __wrapper_class, int)' + + NonTrivialDerived NonTrivialDerivedStructWithPtr(10); + myQueue.submit([&](sycl::handler &h) { + h.single_task([=]() { return NonTrivialDerivedStructWithPtr.i;}); + }); + // CHECK: FunctionDecl {{.*}}NonTrivialStructInBase{{.*}} 'void (__wrapper_class, int, int)' } }