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 70b2e08c833a0..853fb568efcf8 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