From 3916d3b66a89806e3c814326eddb4b35c5e91f30 Mon Sep 17 00:00:00 2001 From: elizabethandrews Date: Tue, 27 Sep 2022 13:22:29 -0400 Subject: [PATCH] [SYCL] Redesign pointer handling for OpenCL kernel generation (#6728) Requirement - Do not decompose types with pointers when generating OpenCL kernel arguments. This PR adds logic to stop decomposing trivial types containing pointers. For every SYCL kernel argument which is a record type containing a pointer (or has a field or a base class with a pointer), we generate a new record type with all pointers in __global address space. This compiler generated type is the openCL kernel argument. In the kernel body, we initialize the local clone via memcpy. Limitations: 1. Array of pointers or array of types with pointers are still decomposed to it's elements. 2. Due to current implementation restrictions, types which are not default constructible, continue to trigger decomposition if they contain pointers. Both limitations above will hopefully be fixed in follow-up PRs. Signed-off-by: Elizabeth Andrews --- clang/include/clang/Basic/Attr.td | 8 + clang/lib/Sema/SemaSYCL.cpp | 455 ++++++++++++++++-- clang/test/CodeGenSYCL/inheritance.cpp | 45 +- .../CodeGenSYCL/no_opaque_inheritance.cpp | 48 +- .../no_opaque_pointers-in-structs.cpp | 4 +- .../test/CodeGenSYCL/pointers-in-structs.cpp | 4 +- .../test/CodeGenSYCL/struct_kernel_param.cpp | 8 +- clang/test/SemaSYCL/array-kernel-param.cpp | 194 +++++++- clang/test/SemaSYCL/decomposition.cpp | 66 ++- clang/test/SemaSYCL/inheritance.cpp | 61 ++- clang/test/SemaSYCL/kernel-arg-opt-report.cpp | 99 +++- clang/test/SemaSYCL/union-kernel-param2.cpp | 90 +++- 12 files changed, 928 insertions(+), 154 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index b3efdd61a0fb6..7ed985c795b48 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1429,6 +1429,14 @@ def SYCLRequiresDecomposition : InheritableAttr { let Documentation = [InternalOnly]; } +def SYCLGenerateNewType : InheritableAttr { + // No spellings, as this is for internal use. + let Spellings = []; + let Subjects = SubjectList<[Named]>; + let LangOpts = [SYCLIsDevice, SYCLIsHost]; + let Documentation = [InternalOnly]; +} + 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 1e5a7c5a4bc02..03c0b5db43cf4 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1299,6 +1299,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. @@ -1453,6 +1454,16 @@ void KernelObjVisitor::visitRecord(const CXXRecordDecl *Owner, ParentTy &Parent, // 'complex', so all handlers are called in this case with the 'complex' // case. visitComplexRecord(Owner, Parent, Wrapper, RecordTy, Handlers...); + } else if (AnyTrue:: + Value) { + // We are currently in PointerHandler visitor. + if (RD->hasAttr()) { + // This is a record containing pointers. + visitComplexRecord(Owner, Parent, Wrapper, RecordTy, Handlers...); + } 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 @@ -1484,6 +1495,16 @@ void KernelObjVisitor::visitArray(const CXXRecordDecl *Owner, FieldDecl *Field, if (Field->hasAttr()) { visitComplexArray(Owner, Field, ArrayTy, Handlers...); + } else if (AnyTrue:: + Value) { + 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' statement above. + visitSimpleArray(Owner, Field, ArrayTy, Handlers...); } else { if (!AllTrue::Value) visitSimpleArray( @@ -1716,9 +1737,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; @@ -1728,6 +1753,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 &, @@ -1746,23 +1772,44 @@ 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; } 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. + 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()) { + // 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 { + // We are visiting a non-trivial type with pointer. + CollectionStack.back() = true; + if (!RD->hasAttr()) + RD->addAttr(SYCLRequiresDecompositionAttr::CreateImplicit( + SemaRef.getASTContext())); + } } return true; } @@ -1770,25 +1817,46 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { bool enterStruct(const CXXRecordDecl *, const CXXBaseSpecifier &, QualType) final { CollectionStack.push_back(false); + PointerStack.push_back(false); return true; } 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. + 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()) { + // 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 { + // We are visiting a non-trivial type with pointer. + CollectionStack.back() = true; + if (!RD->hasAttr()) + RD->addAttr(SYCLRequiresDecompositionAttr::CreateImplicit( + SemaRef.getASTContext())); + } } - return true; } bool enterArray(FieldDecl *, QualType ArrayTy, QualType ElementTy) final { CollectionStack.push_back(false); + PointerStack.push_back(false); return true; } @@ -1800,9 +1868,182 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { FD->addAttr(SYCLRequiresDecompositionAttr::CreateImplicit( SemaRef.getASTContext())); CollectionStack.back() = true; + PointerStack.pop_back(); + } else if (PointerStack.pop_back_val()) { + // FIXME: Array of pointers/ array of type containing pointers + // will be handled in a follow up PR. Currently, they continue + // to trigger decomposition. + if (!FD->hasAttr()) + FD->addAttr(SYCLRequiresDecompositionAttr::CreateImplicit( + SemaRef.getASTContext())); + CollectionStack.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; + + IdentifierInfo *getModifiedName(IdentifierInfo *Id) { + std::string Name = + Id ? (Twine("__generated_") + Id->getName()).str() : "__generated_"; + 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(), RD->getTagKind(), + 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(), + 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 *Parent, const CXXRecordDecl *RD, + const CXXBaseSpecifier &BS) { + TypeSourceInfo *TInfo = SemaRef.getASTContext().getTrivialTypeSourceInfo( + QualType(RD->getTypeForDecl(), 0), SourceLocation()); + CXXBaseSpecifier *ModifiedBase = SemaRef.CheckBaseSpecifier( + 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) { + createNewType(RD); + } + + bool enterStruct(const CXXRecordDecl *, FieldDecl *, QualType Ty) final { + createNewType(Ty->getAsCXXRecordDecl()); + return true; + } + + bool leaveStruct(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final { + 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 &, + QualType Ty) final { + createNewType(Ty->getAsCXXRecordDecl()); + return true; + } + + bool leaveStruct(const CXXRecordDecl *Parent, const CXXBaseSpecifier &BS, + QualType Ty) final { + CXXRecordDecl *ModifiedRD = getGeneratedNewRecord(Ty->getAsCXXRecordDecl()); + + // Create CXXBaseSpecifier for this generated class. + createBaseSpecifier(Parent, ModifiedRD, BS); + return true; + } + + bool handlePointerType(FieldDecl *FD, QualType FieldTy) final { + QualType PointeeTy = FieldTy->getPointeeType(); + Qualifiers Quals = PointeeTy.getQualifiers(); + LangAS 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; + // We do not need to wrap pointers since this is a pointer inside + // non-decomposed struct. + } + + 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 *, FieldDecl *FD, + QualType Ty) final { + addField(FD, Ty); + return true; + } + + bool handleNonDecompStruct(const CXXRecordDecl *Parent, + const CXXBaseSpecifier &BS, QualType Ty) final { + createBaseSpecifier(Parent, Ty->getAsCXXRecordDecl(), BS); + return true; + } + + bool handleSimpleArrayType(FieldDecl *FD, QualType Ty) final { + addField(FD, Ty); return true; } + + // 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() { + CXXRecordDecl *ModifiedRD = ModifiedRecords.pop_back_val(); + ModifiedRD->completeDefinition(); + + if (!ModifiedBases.empty()) + ModifiedRD->setBases(ModifiedBases.data(), ModifiedBases.size()); + + return QualType(ModifiedRD->getTypeForDecl(), 0); + } }; // A type to Create and own the FunctionDecl for the kernel. @@ -1991,6 +2232,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, @@ -2144,15 +2401,31 @@ 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); + // This is a field which should not be decomposed. + CXXRecordDecl *FieldRecordDecl = Ty->getAsCXXRecordDecl(); + 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 { - addParam(BS, Ty); + // This is a base class which should not be decomposed. + CXXRecordDecl *BaseRecordDecl = Ty->getAsCXXRecordDecl(); + 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; } @@ -2254,7 +2527,7 @@ class SyclKernelArgsSizeChecker : public SyclKernelFieldHandler { }; std::string getKernelArgDesc(StringRef KernelArgDescription) { - if (KernelArgDescription == ":" || KernelArgDescription == "") + if (KernelArgDescription == "") return ""; return ("Compiler generated argument for " + KernelArgDescription + ",") .str(); @@ -2265,27 +2538,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") @@ -2294,10 +2560,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"; @@ -2306,12 +2570,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), ""); } @@ -2353,15 +2620,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); return true; } @@ -2373,19 +2645,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; } @@ -2431,7 +2714,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; @@ -2655,6 +2938,87 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { addFieldInit(FD, Ty, ParamRef); } + 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 *createGetAddressOf(Expr *E) { + return UnaryOperator::Create(SemaRef.Context, E, UO_AddrOf, + SemaRef.Context.getPointerType(E->getType()), + VK_PRValue, OK_Ordinary, KernelCallerSrcLoc, + false, SemaRef.CurFPFeatureOverrides()); + } + + 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()); + + 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 + // a call to __builtin_memcpy to initialize local clone from + // kernel argument. + void handleGeneratedType(FieldDecl *FD, QualType Ty) { + addFieldInit(FD, Ty, None, + InitializationKind::CreateDefault(KernelCallerSrcLoc)); + addFieldMemberExpr(FD, Ty); + Expr *ParamRef = createGetAddressOf(createParamReferenceExpr()); + Expr *LocalCloneRef = createGetAddressOf(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 = createGetAddressOf(createParamReferenceExpr()); + Expr *LocalCloneRef = createGetAddressOf(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( @@ -2886,13 +3250,23 @@ 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()) + handleGeneratedType(FD, Ty); + else + addSimpleFieldInit(FD, Ty); 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; } @@ -2958,7 +3332,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/CodeGenSYCL/inheritance.cpp b/clang/test/CodeGenSYCL/inheritance.cpp index dc46231dac670..883ae6c0b8087 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..b58390d30443f 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/no_opaque_pointers-in-structs.cpp b/clang/test/CodeGenSYCL/no_opaque_pointers-in-structs.cpp index 5ce576ce82184..12e631c65de99 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/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/CodeGenSYCL/struct_kernel_param.cpp b/clang/test/CodeGenSYCL/struct_kernel_param.cpp index f245ba0627b4d..1418a0c519a7a 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/array-kernel-param.cpp b/clang/test/SemaSYCL/array-kernel-param.cpp index 4069c7dd79d09..cd9a8e3355498 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' 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 diff --git a/clang/test/SemaSYCL/decomposition.cpp b/clang/test/SemaSYCL/decomposition.cpp index 2dd4599865e7f..8c79ab582ad9d 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,23 @@ struct StructWithNonDecomposedStruct : StructNonDecomposed { double d; }; +struct StructWithPtr { + StructNonDecomposed member; + int *ptr; + 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; @@ -66,6 +84,7 @@ int main() { StructNonDecomposed ArrayOfSimpleStruct[5]; StructWithNonDecomposedStruct NonDecompStruct; StructWithNonDecomposedStruct ArrayOfNonDecompStruct[5]; + // 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 +96,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 +171,47 @@ int main() { }); // CHECK: FunctionDecl {{.*}}Half2{{.*}} 'void (DerivedStruct)' } + + { + StructWithPtr SimpleStructWithPtr; + myQueue.submit([&](sycl::handler &h) { + h.single_task([=]() { return SimpleStructWithPtr.i; }); + }); + // CHECK: FunctionDecl {{.*}}Pointer{{.*}} 'void (__generated_StructWithPtr)' + + // 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 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)' + } } diff --git a/clang/test/SemaSYCL/inheritance.cpp b/clang/test/SemaSYCL/inheritance.cpp index 1b908dbeb3302..eedcd58e9d07b 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,48 @@ int main() { // Base classes should be initialized first. // CHECK: VarDecl {{.*}} used derived 'derived' cinit // CHECK-NEXT: InitListExpr {{.*}} 'derived' -// CHECK-NEXT: CXXConstructExpr {{.*}}'base' 'void (const base &) noexcept' -// CHECK-NEXT: ImplicitCastExpr {{.*}}'const base' lvalue + +// base is a simple class with no corresponding generated type. Therefore +// copy from ParamVar +// CHECK-NEXT: CXXConstructExpr {{.*}} 'base':'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':'second_base' 'void () noexcept' + +// third_base contains special type accessor. Therefore it is decomposed and it's +// data members are copied from corresponding 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':'third_base' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} 'derived' lvalue Var {{.*}} 'derived' 'derived' 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 ' diff --git a/clang/test/SemaSYCL/union-kernel-param2.cpp b/clang/test/SemaSYCL/union-kernel-param2.cpp index 7d38aa751a59e..f55db4426e494 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 {{.*}} '__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 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 {{.*}} '__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: IntegerLiteral {{.*}} 'unsigned long' 24