diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 72a1551edd9d9..b3ecda585cf3e 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -617,286 +617,6 @@ static CXXRecordDecl *getKernelObjectType(FunctionDecl *Caller) { return (*Caller->param_begin())->getType()->getAsCXXRecordDecl(); } -// Creates body for new OpenCL kernel. This body contains initialization of SYCL -// kernel object fields with kernel parameters and a little bit transformed body -// of the kernel caller function. -static CompoundStmt *CreateOpenCLKernelBody(Sema &S, - FunctionDecl *KernelCallerFunc, - DeclContext *KernelDecl) { - using BodyStmtsT = llvm::SmallVector; - - BodyStmtsT BodyStmts; - BodyStmtsT FinalizeStmts; - CXXRecordDecl *LC = getKernelObjectType(KernelCallerFunc); - assert(LC && "Kernel object must be available"); - - if (getKernelInvocationKind(KernelCallerFunc) == InvokeParallelForWorkGroup) { - CXXRecordDecl *LambdaObjTy = - KernelCallerFunc->getParamDecl(0)->getType()->getAsCXXRecordDecl(); - assert(LambdaObjTy && - "unexpected kernel_parallel_for_work_group parameter type"); - FindPFWGLambdaFnVisitor V(LambdaObjTy); - V.TraverseStmt(KernelCallerFunc->getBody()); - CXXMethodDecl *WGLambdaFn = V.getLambdaFn(); - assert(WGLambdaFn && "PFWG lambda not found"); - // Mark the function that it "works" in a work group scope: - // NOTE: In case of parallel_for_work_item the marker call itself is marked - // with work item scope attribute, here the '()' operator of the - // object passed as parameter is marked. This is an optimization - - // there are a lot of locals created at parallel_for_work_group scope - // before calling the lambda - it is more efficient to have all of - // them in the private address space rather then sharing via the local - // AS. See parallel_for_work_group implementation in the SYCL headers. - if (!WGLambdaFn->hasAttr()) { - WGLambdaFn->addAttr(SYCLScopeAttr::CreateImplicit( - S.getASTContext(), SYCLScopeAttr::Level::WorkGroup)); - // Search and mark parallel_for_work_item calls: - MarkWIScopeFnVisitor MarkWIScope(S.getASTContext()); - MarkWIScope.TraverseDecl(WGLambdaFn); - // Now mark local variables declared in the PFWG lambda with work group - // scope attribute - addScopeAttrToLocalVars(*WGLambdaFn); - } - } - - TypeSourceInfo *TSInfo = LC->isLambda() ? LC->getLambdaTypeInfo() : nullptr; - - // Create a local kernel object (lambda or functor) assembled from the - // incoming formal parameters - auto KernelObjClone = VarDecl::Create( - S.Context, KernelDecl, SourceLocation(), SourceLocation(), - LC->getIdentifier(), QualType(LC->getTypeForDecl(), 0), TSInfo, SC_None); - Stmt *DS = new (S.Context) DeclStmt(DeclGroupRef(KernelObjClone), - SourceLocation(), SourceLocation()); - BodyStmts.push_back(DS); - auto KernelObjCloneRef = - DeclRefExpr::Create(S.Context, NestedNameSpecifierLoc(), SourceLocation(), - KernelObjClone, false, DeclarationNameInfo(), - QualType(LC->getTypeForDecl(), 0), VK_LValue); - - auto KernelFuncDecl = cast(KernelDecl); - auto KernelFuncParam = - KernelFuncDecl->param_begin(); // Iterator to ParamVarDecl (VarDecl) - if (KernelFuncParam) { - llvm::SmallVector InitExprs; - InitializedEntity VarEntity = - InitializedEntity::InitializeVariable(KernelObjClone); - for (auto Field : LC->fields()) { - // Creates Expression for special SYCL object: accessor or sampler. - // All special SYCL objects must have __init method, here we use it to - // initialize them. We create call of __init method and pass built kernel - // arguments as parameters to the __init method. - auto getExprForSpecialSYCLObj = [&](const QualType ¶mTy, - FieldDecl *Field, - const CXXRecordDecl *CRD, - Expr *Base, - const std::string &MethodName, - BodyStmtsT &Statements) { - CXXMethodDecl *Method = getMethodByName(CRD, MethodName); - assert(Method && - "The accessor/sampler/stream must have the __init method. Stream" - " must also have __finalize method"); - unsigned NumParams = Method->getNumParams(); - llvm::SmallVector ParamDREs(NumParams); - auto KFP = KernelFuncParam; - for (size_t I = 0; I < NumParams; ++KFP, ++I) { - QualType ParamType = (*KFP)->getOriginalType(); - ParamDREs[I] = DeclRefExpr::Create( - S.Context, NestedNameSpecifierLoc(), SourceLocation(), *KFP, - false, DeclarationNameInfo(), ParamType, VK_LValue); - } - - if (NumParams) - std::advance(KernelFuncParam, NumParams - 1); - - DeclAccessPair FieldDAP = DeclAccessPair::make(Field, AS_none); - // [kernel_obj or wrapper object].special_obj - auto SpecialObjME = MemberExpr::Create( - S.Context, Base, false, SourceLocation(), NestedNameSpecifierLoc(), - SourceLocation(), Field, FieldDAP, - DeclarationNameInfo(Field->getDeclName(), SourceLocation()), - nullptr, Field->getType(), VK_LValue, OK_Ordinary, NOUR_None); - - // [kernel_obj or wrapper object].special_obj.__init - DeclAccessPair MethodDAP = DeclAccessPair::make(Method, AS_none); - auto ME = MemberExpr::Create( - S.Context, SpecialObjME, false, SourceLocation(), - NestedNameSpecifierLoc(), SourceLocation(), Method, MethodDAP, - DeclarationNameInfo(Method->getDeclName(), SourceLocation()), - nullptr, Method->getType(), VK_LValue, OK_Ordinary, NOUR_None); - - // Not referenced -> not emitted - S.MarkFunctionReferenced(SourceLocation(), Method, true); - - QualType ResultTy = Method->getReturnType(); - ExprValueKind VK = Expr::getValueKindForType(ResultTy); - ResultTy = ResultTy.getNonLValueExprType(S.Context); - - llvm::SmallVector ParamStmts; - const auto *Proto = cast(Method->getType()); - S.GatherArgumentsForCall(SourceLocation(), Method, Proto, 0, - ParamDREs, ParamStmts); - // [kernel_obj or wrapper object].accessor.__init(_ValueType*, - // range, range, id) - CXXMemberCallExpr *Call = CXXMemberCallExpr::Create( - S.Context, ME, ParamStmts, ResultTy, VK, SourceLocation()); - Statements.push_back(Call); - }; - - // Recursively search for accessor fields to initialize them with kernel - // parameters - std::function - getExprForWrappedAccessorInit = - [&](const CXXRecordDecl *CRD, Expr *Base) { - for (auto *WrapperFld : CRD->fields()) { - QualType FldType = WrapperFld->getType(); - CXXRecordDecl *WrapperFldCRD = FldType->getAsCXXRecordDecl(); - if (FldType->isStructureOrClassType()) { - if (Util::isSyclAccessorType(FldType)) { - // Accessor field found - create expr to initialize this - // accessor object. Need to start from the next target - // function parameter, since current one is the wrapper - // object or parameter of the previous processed accessor - // object. - KernelFuncParam++; - getExprForSpecialSYCLObj(FldType, WrapperFld, - WrapperFldCRD, Base, - InitMethodName, BodyStmts); - } else if (Util::isSyclSpecConstantType(FldType)) { - // Specialization constants are "invisible" to the - // kernel argument creation and device-side SYCL object - // materialization infrastructure in this source. - // It is OK not to really materialize them on the kernel - // side, because their only use can be via - // 'spec_const_obj.get()' method, which is translated to - // an intrinsic and 'this' is really never used. - } else { - // Field is a structure or class so change the wrapper - // object and recursively search for accessor field. - DeclAccessPair WrapperFieldDAP = - DeclAccessPair::make(WrapperFld, AS_none); - auto NewBase = MemberExpr::Create( - S.Context, Base, false, SourceLocation(), - NestedNameSpecifierLoc(), SourceLocation(), - WrapperFld, WrapperFieldDAP, - DeclarationNameInfo(WrapperFld->getDeclName(), - SourceLocation()), - nullptr, WrapperFld->getType(), VK_LValue, - OK_Ordinary, NOUR_None); - getExprForWrappedAccessorInit(WrapperFldCRD, NewBase); - } - } - } - }; - - // Run through kernel object fields and add initialization for them using - // built kernel parameters. There are a several possible cases: - // - Kernel object field is a SYCL special object (SYCL accessor or SYCL - // sampler). These objects has a special initialization scheme - using - // __init method. - // - Kernel object field has a scalar type. In this case we should add - // simple initialization. - // - Kernel object field has a structure or class type. Same handling as - // a scalar but we should check if this structure/class contains - // accessors and add initialization for them properly. - QualType FieldType = Field->getType(); - CXXRecordDecl *CRD = FieldType->getAsCXXRecordDecl(); - InitializedEntity Entity = - InitializedEntity::InitializeMember(Field, &VarEntity); - if (Util::isSyclAccessorType(FieldType) || - Util::isSyclSamplerType(FieldType)) { - // Initialize with the default constructor. - InitializationKind InitKind = - InitializationKind::CreateDefault(SourceLocation()); - InitializationSequence InitSeq(S, Entity, InitKind, None); - ExprResult MemberInit = InitSeq.Perform(S, Entity, InitKind, None); - InitExprs.push_back(MemberInit.get()); - getExprForSpecialSYCLObj(FieldType, Field, CRD, KernelObjCloneRef, - InitMethodName, BodyStmts); - } else if (Util::isSyclSpecConstantType(FieldType)) { - // Just skip specialization constants - not part of signature. - } else if (CRD || FieldType->isScalarType()) { - // If field has built-in or a structure/class type just initialize - // this field with corresponding kernel argument using copy - // initialization. - QualType ParamType = (*KernelFuncParam)->getOriginalType(); - Expr *DRE = - DeclRefExpr::Create(S.Context, NestedNameSpecifierLoc(), - SourceLocation(), *KernelFuncParam, false, - DeclarationNameInfo(), ParamType, VK_LValue); - - if (FieldType->isPointerType() && - FieldType->getPointeeType().getAddressSpace() != - ParamType->getPointeeType().getAddressSpace()) - DRE = ImplicitCastExpr::Create(S.Context, FieldType, - CK_AddressSpaceConversion, DRE, - nullptr, VK_RValue); - InitializationKind InitKind = - InitializationKind::CreateCopy(SourceLocation(), SourceLocation()); - InitializationSequence InitSeq(S, Entity, InitKind, DRE); - - ExprResult MemberInit = InitSeq.Perform(S, Entity, InitKind, DRE); - InitExprs.push_back(MemberInit.get()); - - if (CRD) { - // If a structure/class type has accessor fields then we need to - // initialize these accessors in proper way by calling __init method - // of the accessor and passing corresponding kernel parameters. - DeclAccessPair FieldDAP = DeclAccessPair::make(Field, AS_none); - auto Lhs = MemberExpr::Create( - S.Context, KernelObjCloneRef, false, SourceLocation(), - NestedNameSpecifierLoc(), SourceLocation(), Field, FieldDAP, - DeclarationNameInfo(Field->getDeclName(), SourceLocation()), - nullptr, Field->getType(), VK_LValue, OK_Ordinary, NOUR_None); - getExprForWrappedAccessorInit(CRD, Lhs); - if (Util::isSyclStreamType(FieldType)) { - // Generate call to the __init method of the stream class after - // initializing accessors wrapped by this stream object - getExprForSpecialSYCLObj(FieldType, Field, CRD, KernelObjCloneRef, - InitMethodName, BodyStmts); - - // Generate call to the __finalize method of stream class. - // Will put it later to the end of function body. - getExprForSpecialSYCLObj(FieldType, Field, CRD, KernelObjCloneRef, - FinalizeMethodName, FinalizeStmts); - } - } - } else { - llvm_unreachable("Unsupported field type"); - } - KernelFuncParam++; - } - Expr *ILE = new (S.Context) - InitListExpr(S.Context, SourceLocation(), InitExprs, SourceLocation()); - ILE->setType(QualType(LC->getTypeForDecl(), 0)); - KernelObjClone->setInit(ILE); - } - - // In the kernel caller function kernel object is a function parameter, so we - // need to replace all refs to this kernel oject with refs to our clone - // declared inside kernel body. - Stmt *FunctionBody = KernelCallerFunc->getBody(); - ParmVarDecl *KernelObjParam = *(KernelCallerFunc->param_begin()); - - // DeclRefExpr with valid source location but with decl which is not marked - // as used is invalid. - KernelObjClone->setIsUsed(); - std::pair MappingPair; - MappingPair.first = KernelObjParam; - MappingPair.second = KernelObjClone; - - // Function scope might be empty, so we do push - S.PushFunctionScope(); - KernelBodyTransform KBT(MappingPair, S); - Stmt *NewBody = KBT.TransformStmt(FunctionBody).get(); - BodyStmts.push_back(NewBody); - - BodyStmts.insert(BodyStmts.end(), FinalizeStmts.begin(), FinalizeStmts.end()); - - return CompoundStmt::Create(S.Context, BodyStmts, SourceLocation(), - SourceLocation()); -} - /// Creates a kernel parameter descriptor /// \param Src field declaration to construct name from /// \param Ty the desired parameter type @@ -908,311 +628,777 @@ static ParamDesc makeParamDesc(const FieldDecl *Src, QualType Ty) { Ctx.getTrivialTypeSourceInfo(Ty)); } +static ParamDesc makeParamDesc(ASTContext &Ctx, const CXXBaseSpecifier &Src, + QualType Ty) { + // TODO: There is no name for the base available, but duplicate names are + // seemingly already possible, so we'll give them all the same name for now. + // This only happens with the accessor types. + std::string Name = "_arg__base"; + return std::make_tuple(Ty, &Ctx.Idents.get(Name), + Ctx.getTrivialTypeSourceInfo(Ty)); +} + /// \return the target of given SYCL accessor type static target getAccessTarget(const ClassTemplateSpecializationDecl *AccTy) { return static_cast( AccTy->getTemplateArgs()[3].getAsIntegral().getExtValue()); } -// Creates list of kernel parameters descriptors using KernelObj (kernel object) -// Fields of kernel object must be initialized with SYCL kernel arguments so -// in the following function we extract types of kernel object fields and add it -// to the array with kernel parameters descriptors. -// Returns true if all arguments are successfully built. -static bool buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj, - SmallVectorImpl &ParamDescs) { - auto CreateAndAddPrmDsc = [&](const FieldDecl *Fld, const QualType &ArgType) { - // Create a parameter descriptor and append it to the result - ParamDescs.push_back(makeParamDesc(Fld, ArgType)); - }; - - // Creates a parameter descriptor for SYCL special object - SYCL accessor or - // sampler. +// The first template argument to the kernel caller function is used to identify +// the kernel itself. +static QualType calculateKernelNameType(ASTContext &Ctx, + FunctionDecl *KernelCallerFunc) { + const TemplateArgumentList *TAL = + KernelCallerFunc->getTemplateSpecializationArgs(); + assert(TAL && "No template argument info"); + return TypeName::getFullyQualifiedType(TAL->get(0).getAsType(), Ctx, + /*WithGlobalNSPrefix=*/true); +} + +// Gets a name for the OpenCL kernel function, calculated from the first +// template argument of the kernel caller function. +static std::pair +constructKernelName(Sema &S, FunctionDecl *KernelCallerFunc, + MangleContext &MC) { + QualType KernelNameType = + calculateKernelNameType(S.getASTContext(), KernelCallerFunc); + + SmallString<256> Result; + llvm::raw_svector_ostream Out(Result); + + MC.mangleTypeName(KernelNameType, Out); + + return {std::string(Out.str()), + PredefinedExpr::ComputeName(S.getASTContext(), + PredefinedExpr::UniqueStableNameType, + KernelNameType)}; +} + +// anonymous namespace so these don't get linkage. +namespace { + +QualType getItemType(const FieldDecl *FD) { return FD->getType(); } +QualType getItemType(const CXXBaseSpecifier &BS) { return BS.getType(); } + +// Implements the 'for-each-visitor' pattern. +template +static void VisitAccessorWrapper(CXXRecordDecl *Owner, ParentTy &Parent, + CXXRecordDecl *Wrapper, + Handlers &... handlers); + +template +static void VisitAccessorWrapperHelper(CXXRecordDecl *Owner, RangeTy Range, + Handlers &... handlers) { + for (const auto &Item : Range) { + QualType ItemTy = getItemType(Item); + if (Util::isSyclAccessorType(ItemTy)) + (void)std::initializer_list{ + (handlers.handleSyclAccessorType(Item, ItemTy), 0)...}; + else if (ItemTy->isStructureOrClassType()) { + VisitAccessorWrapper(Owner, Item, ItemTy->getAsCXXRecordDecl(), + handlers...); + if (Util::isSyclStreamType(ItemTy)) + (void)std::initializer_list{ + (handlers.handleSyclStreamType(Item, ItemTy), 0)...}; + } + } +} + +// Parent contains the FieldDecl or CXXBaseSpecifier that was used to enter +// the Wrapper structure that we're currently visiting. Owner is the parent type +// (which doesn't exist in cases where it is a FieldDecl in the 'root'), and +// Wrapper is the current struct being unwrapped. +template +static void VisitAccessorWrapper(CXXRecordDecl *Owner, ParentTy &Parent, + CXXRecordDecl *Wrapper, + Handlers &... handlers) { + (void)std::initializer_list{(handlers.enterStruct(Owner, Parent), 0)...}; + VisitAccessorWrapperHelper(Wrapper, Wrapper->bases(), handlers...); + VisitAccessorWrapperHelper(Wrapper, Wrapper->fields(), handlers...); + (void)std::initializer_list{(handlers.leaveStruct(Owner, Parent), 0)...}; +} + +// A visitor function that dispatches to functions as defined in +// SyclKernelFieldHandler for the purposes of kernel generation. +template +static void VisitRecordFields(RecordDecl::field_range Fields, + Handlers &... handlers) { +#define KF_FOR_EACH(FUNC) \ + (void)std::initializer_list { (handlers.FUNC(Field, FieldTy), 0)... } + + for (const auto &Field : Fields) { + QualType FieldTy = Field->getType(); + + if (Util::isSyclAccessorType(FieldTy)) + KF_FOR_EACH(handleSyclAccessorType); + else if (Util::isSyclSamplerType(FieldTy)) + KF_FOR_EACH(handleSyclSamplerType); + else if (Util::isSyclSpecConstantType(FieldTy)) + KF_FOR_EACH(handleSyclSpecConstantType); + else if (Util::isSyclStreamType(FieldTy)) { + // Stream actually wraps accessors, so do recursion + CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl(); + VisitAccessorWrapper(nullptr, Field, RD, handlers...); + KF_FOR_EACH(handleSyclStreamType); + } else if (FieldTy->isStructureOrClassType()) { + KF_FOR_EACH(handleStructType); + CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl(); + VisitAccessorWrapper(nullptr, Field, RD, handlers...); + } else if (FieldTy->isReferenceType()) + KF_FOR_EACH(handleReferenceType); + else if (FieldTy->isPointerType()) + KF_FOR_EACH(handlePointerType); + else if (FieldTy->isArrayType()) + KF_FOR_EACH(handleArrayType); + else if (FieldTy->isScalarType()) + KF_FOR_EACH(handleScalarType); + else + KF_FOR_EACH(handleOtherType); + } +#undef KF_FOR_EACH +} + +// A base type that the SYCL OpenCL Kernel construction task uses to implement +// individual tasks. +template class SyclKernelFieldHandler { +protected: + Sema &SemaRef; + SyclKernelFieldHandler(Sema &S) : SemaRef(S) {} + +public: + // Mark these virutal so that we can use override in the implementer classes, + // despite virtual dispatch never being used. + + // Accessor can be a base class or a field decl, so both must be handled. + virtual void handleSyclAccessorType(const CXXBaseSpecifier &, QualType) {} + virtual void handleSyclAccessorType(FieldDecl *, QualType) {} + virtual void handleSyclSamplerType(FieldDecl *, QualType) {} + virtual void handleSyclSpecConstantType(FieldDecl *, QualType) {} + virtual void handleSyclStreamType(const CXXBaseSpecifier &, QualType) {} + virtual void handleSyclStreamType(FieldDecl *, QualType) {} + virtual void handleStructType(FieldDecl *, QualType) {} + virtual void handleReferenceType(FieldDecl *, QualType) {} + virtual void handlePointerType(FieldDecl *, QualType) {} + virtual void handleArrayType(FieldDecl *, QualType) {} + virtual void handleScalarType(FieldDecl *, QualType) {} + // Most handlers shouldn't be handling this, just the field checker. + virtual void handleOtherType(FieldDecl *, QualType) {} + + // The following are only used for keeping track of where we are in the base + // class/field graph. Int Headers use this to calculate offset, most others + // don't have a need for these. + + virtual void enterStruct(const CXXRecordDecl *, FieldDecl *) {} + virtual void leaveStruct(const CXXRecordDecl *, FieldDecl *) {} + virtual void enterStruct(const CXXRecordDecl *, const CXXBaseSpecifier &) {} + virtual void leaveStruct(const CXXRecordDecl *, const CXXBaseSpecifier &) {} +}; + +// A type to check the valididty of all of the argument types. +class SyclKernelFieldChecker + : public SyclKernelFieldHandler { + bool IsInvalid = false; + DiagnosticsEngine &Diag; + +public: + SyclKernelFieldChecker(Sema &S) + : SyclKernelFieldHandler(S), Diag(S.getASTContext().getDiagnostics()) {} + bool isValid() { return !IsInvalid; } + + void handleReferenceType(FieldDecl *FD, QualType FieldTy) final { + IsInvalid = Diag.Report(FD->getLocation(), diag::err_bad_kernel_param_type) + << FieldTy; + } + void handleStructType(FieldDecl *FD, QualType FieldTy) final { + if (SemaRef.getASTContext().getLangOpts().SYCLStdLayoutKernelParams && + !FieldTy->isStandardLayoutType()) + IsInvalid = + Diag.Report(FD->getLocation(), diag::err_sycl_non_std_layout_type) + << FieldTy; + else { + CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl(); + if (!RD->hasTrivialCopyConstructor()) + + IsInvalid = + Diag.Report(FD->getLocation(), + diag::err_sycl_non_trivially_copy_ctor_dtor_type) + << 0 << FieldTy; + else if (!RD->hasTrivialDestructor()) + IsInvalid = + Diag.Report(FD->getLocation(), + diag::err_sycl_non_trivially_copy_ctor_dtor_type) + << 1 << FieldTy; + } + } + + // We should be able to handle this, so we made it part of the visitor, but + // this is 'to be implemented'. + void handleArrayType(FieldDecl *FD, QualType FieldTy) final { + IsInvalid = Diag.Report(FD->getLocation(), diag::err_bad_kernel_param_type) + << FieldTy; + } + + void handleOtherType(FieldDecl *FD, QualType FieldTy) final { + IsInvalid = Diag.Report(FD->getLocation(), diag::err_bad_kernel_param_type) + << FieldTy; + } +}; + +// A type to Create and own the FunctionDecl for the kernel. +class SyclKernelDeclCreator + : public SyclKernelFieldHandler { + FunctionDecl *KernelDecl; + llvm::SmallVector Params; + SyclKernelFieldChecker &ArgChecker; + Sema::ContextRAII FuncContext; + // Holds the last handled field's first parameter. This doesn't store an + // iterator as push_back invalidates iterators. + size_t LastParamIndex = 0; + + void addParam(const FieldDecl *FD, QualType FieldTy) { + ParamDesc newParamDesc = makeParamDesc(FD, FieldTy); + addParam(newParamDesc, FieldTy); + } + + void addParam(const CXXBaseSpecifier &BS, QualType FieldTy) { + ParamDesc newParamDesc = + makeParamDesc(SemaRef.getASTContext(), BS, FieldTy); + addParam(newParamDesc, FieldTy); + } + + void addParam(ParamDesc newParamDesc, QualType FieldTy) { + // Create a new ParmVarDecl based on the new info. + auto *NewParam = ParmVarDecl::Create( + SemaRef.getASTContext(), KernelDecl, SourceLocation(), SourceLocation(), + std::get<1>(newParamDesc), std::get<0>(newParamDesc), + std::get<2>(newParamDesc), SC_None, /*DefArg*/ nullptr); + + NewParam->setScopeInfo(0, Params.size()); + NewParam->setIsUsed(); + + LastParamIndex = Params.size(); + Params.push_back(NewParam); + } + // All special SYCL objects must have __init method. We extract types for // kernel parameters from __init method parameters. We will use __init method // and kernel parameters which we build here to initialize special objects in // the kernel body. - auto createSpecialSYCLObjParamDesc = [&](const FieldDecl *Fld, - const QualType &ArgTy) { - const auto *RecordDecl = ArgTy->getAsCXXRecordDecl(); - assert(RecordDecl && "Special SYCL object must be of a record type"); + void handleSpecialType(FieldDecl *FD, QualType FieldTy) { + const auto *RecordDecl = FieldTy->getAsCXXRecordDecl(); + assert(RecordDecl && "The accessor/sampler must be a RecordDecl"); + CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, InitMethodName); + assert(InitMethod && "The accessor/sampler must have the __init method"); + + // Don't do -1 here because we count on this to be the first parameter added + // (if any). + size_t ParamIndex = Params.size(); + for (const ParmVarDecl *Param : InitMethod->parameters()) + addParam(FD, Param->getType().getCanonicalType()); + LastParamIndex = ParamIndex; + } + + static void setKernelImplicitAttrs(ASTContext &Context, FunctionDecl *FD, + StringRef Name) { + // Set implict attributes. + FD->addAttr(OpenCLKernelAttr::CreateImplicit(Context)); + FD->addAttr(AsmLabelAttr::CreateImplicit(Context, Name)); + FD->addAttr(ArtificialAttr::CreateImplicit(Context)); + } + + static FunctionDecl *createKernelDecl(ASTContext &Ctx, StringRef Name, + SourceLocation Loc, bool IsInline) { + // Create this with no prototype, and we can fix this up after we've seen + // all the params. + FunctionProtoType::ExtProtoInfo Info(CC_OpenCLKernel); + QualType FuncType = Ctx.getFunctionType(Ctx.VoidTy, {}, Info); + + FunctionDecl *FD = FunctionDecl::Create( + Ctx, Ctx.getTranslationUnitDecl(), Loc, Loc, &Ctx.Idents.get(Name), + FuncType, Ctx.getTrivialTypeSourceInfo(Ctx.VoidTy), SC_None); + FD->setImplicitlyInline(IsInline); + setKernelImplicitAttrs(Ctx, FD, Name); + + // Add kernel to translation unit to see it in AST-dump. + Ctx.getTranslationUnitDecl()->addDecl(FD); + return FD; + } + +public: + SyclKernelDeclCreator(Sema &S, SyclKernelFieldChecker &ArgChecker, + StringRef Name, SourceLocation Loc, bool IsInline) + : SyclKernelFieldHandler(S), + KernelDecl(createKernelDecl(S.getASTContext(), Name, Loc, IsInline)), + ArgChecker(ArgChecker), FuncContext(SemaRef, KernelDecl) {} + + ~SyclKernelDeclCreator() { + ASTContext &Ctx = SemaRef.getASTContext(); + FunctionProtoType::ExtProtoInfo Info(CC_OpenCLKernel); + + SmallVector ArgTys; + std::transform(std::begin(Params), std::end(Params), + std::back_inserter(ArgTys), + [](const ParmVarDecl *PVD) { return PVD->getType(); }); + + QualType FuncType = Ctx.getFunctionType(Ctx.VoidTy, ArgTys, Info); + KernelDecl->setType(FuncType); + KernelDecl->setParams(Params); + + if (ArgChecker.isValid()) + SemaRef.addSyclDeviceDecl(KernelDecl); + } + void handleSyclAccessorType(const CXXBaseSpecifier &BS, + QualType FieldTy) final { + const auto *RecordDecl = FieldTy->getAsCXXRecordDecl(); + assert(RecordDecl && "The accessor/sampler must be a RecordDecl"); CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, InitMethodName); assert(InitMethod && "The accessor/sampler must have the __init method"); - unsigned NumParams = InitMethod->getNumParams(); + + // Don't do -1 here because we count on this to be the first parameter added + // (if any). + size_t ParamIndex = Params.size(); + for (const ParmVarDecl *Param : InitMethod->parameters()) + addParam(BS, Param->getType().getCanonicalType()); + LastParamIndex = ParamIndex; + } + + void handleSyclAccessorType(FieldDecl *FD, QualType FieldTy) final { + handleSpecialType(FD, FieldTy); + } + + void handleSyclSamplerType(FieldDecl *FD, QualType FieldTy) final { + handleSpecialType(FD, FieldTy); + } + + void handlePointerType(FieldDecl *FD, QualType FieldTy) final { + // USM allows to use raw pointers instead of buffers/accessors, but these + // pointers point to the specially allocated memory. For pointer fields we + // add a kernel argument with the same type as field but global address + // space, because OpenCL requires it. + QualType PointeeTy = FieldTy->getPointeeType(); + Qualifiers Quals = PointeeTy.getQualifiers(); + Quals.setAddressSpace(LangAS::opencl_global); + PointeeTy = SemaRef.getASTContext().getQualifiedType( + PointeeTy.getUnqualifiedType(), Quals); + QualType ModTy = SemaRef.getASTContext().getPointerType(PointeeTy); + addParam(FD, ModTy); + } + + void handleScalarType(FieldDecl *FD, QualType FieldTy) final { + addParam(FD, FieldTy); + } + + void handleStructType(FieldDecl *FD, QualType FieldTy) final { + addParam(FD, FieldTy); + } + + void handleSyclStreamType(FieldDecl *FD, QualType FieldTy) final { + addParam(FD, FieldTy); + } + + void handleSyclStreamType(const CXXBaseSpecifier &, QualType FieldTy) final { + // FIXME SYCL stream should be usable as a base type + // See https://github.com/intel/llvm/issues/1552 + } + + void setBody(CompoundStmt *KB) { KernelDecl->setBody(KB); } + + FunctionDecl *getKernelDecl() { return KernelDecl; } + + llvm::ArrayRef getParamVarDeclsForCurrentField() { + return ArrayRef(std::begin(Params) + LastParamIndex, + std::end(Params)); + } +}; + +class SyclKernelBodyCreator + : public SyclKernelFieldHandler { + SyclKernelDeclCreator &DeclCreator; + llvm::SmallVector BodyStmts; + llvm::SmallVector FinalizeStmts; + llvm::SmallVector InitExprs; + VarDecl *KernelObjClone; + InitializedEntity VarEntity; + CXXRecordDecl *KernelObj; + llvm::SmallVector MemberExprBases; + FunctionDecl *KernelCallerFunc; + + // Using the statements/init expressions that we've created, this generates + // the kernel body compound stmt. CompoundStmt needs to know its number of + // statements in advance to allocate it, so we cannot do this as we go along. + CompoundStmt *createKernelBody() { + + Expr *ILE = new (SemaRef.getASTContext()) InitListExpr( + SemaRef.getASTContext(), SourceLocation(), InitExprs, SourceLocation()); + ILE->setType(QualType(KernelObj->getTypeForDecl(), 0)); + KernelObjClone->setInit(ILE); + Stmt *FunctionBody = KernelCallerFunc->getBody(); + + ParmVarDecl *KernelObjParam = *(KernelCallerFunc->param_begin()); + + // DeclRefExpr with valid source location but with decl which is not marked + // as used is invalid. + KernelObjClone->setIsUsed(); + std::pair MappingPair = + std::make_pair(KernelObjParam, KernelObjClone); + + // Push the Kernel function scope to ensure the scope isn't empty + SemaRef.PushFunctionScope(); + KernelBodyTransform KBT(MappingPair, SemaRef); + Stmt *NewBody = KBT.TransformStmt(FunctionBody).get(); + BodyStmts.push_back(NewBody); + + BodyStmts.insert(BodyStmts.end(), FinalizeStmts.begin(), + FinalizeStmts.end()); + return CompoundStmt::Create(SemaRef.getASTContext(), BodyStmts, {}, {}); + } + + void markParallelWorkItemCalls() { + if (getKernelInvocationKind(KernelCallerFunc) == + InvokeParallelForWorkGroup) { + FindPFWGLambdaFnVisitor V(KernelObj); + V.TraverseStmt(KernelCallerFunc->getBody()); + CXXMethodDecl *WGLambdaFn = V.getLambdaFn(); + assert(WGLambdaFn && "PFWG lambda not found"); + // Mark the function that it "works" in a work group scope: + // NOTE: In case of parallel_for_work_item the marker call itself is + // marked with work item scope attribute, here the '()' operator of the + // object passed as parameter is marked. This is an optimization - + // there are a lot of locals created at parallel_for_work_group + // scope before calling the lambda - it is more efficient to have + // all of them in the private address space rather then sharing via + // the local AS. See parallel_for_work_group implementation in the + // SYCL headers. + if (!WGLambdaFn->hasAttr()) { + WGLambdaFn->addAttr(SYCLScopeAttr::CreateImplicit( + SemaRef.getASTContext(), SYCLScopeAttr::Level::WorkGroup)); + // Search and mark parallel_for_work_item calls: + MarkWIScopeFnVisitor MarkWIScope(SemaRef.getASTContext()); + MarkWIScope.TraverseDecl(WGLambdaFn); + // Now mark local variables declared in the PFWG lambda with work group + // scope attribute + addScopeAttrToLocalVars(*WGLambdaFn); + } + } + } + + MemberExpr *BuildMemberExpr(Expr *Base, ValueDecl *Member) { + DeclAccessPair MemberDAP = DeclAccessPair::make(Member, AS_none); + MemberExpr *Result = SemaRef.BuildMemberExpr( + Base, /*IsArrow */ false, SourceLocation(), NestedNameSpecifierLoc(), + SourceLocation(), Member, MemberDAP, + /*HadMultipleCandidates*/ false, + DeclarationNameInfo(Member->getDeclName(), SourceLocation()), + Member->getType(), VK_LValue, OK_Ordinary); + return Result; + } + + void createExprForStructOrScalar(FieldDecl *FD) { + ParmVarDecl *KernelParameter = + DeclCreator.getParamVarDeclsForCurrentField()[0]; + InitializedEntity Entity = + InitializedEntity::InitializeMember(FD, &VarEntity); + QualType ParamType = KernelParameter->getOriginalType(); + Expr *DRE = SemaRef.BuildDeclRefExpr(KernelParameter, ParamType, VK_LValue, + SourceLocation()); + if (FD->getType()->isPointerType() && + FD->getType()->getPointeeType().getAddressSpace() != + ParamType->getPointeeType().getAddressSpace()) + DRE = ImplicitCastExpr::Create(SemaRef.Context, FD->getType(), + CK_AddressSpaceConversion, DRE, nullptr, + VK_RValue); + InitializationKind InitKind = + InitializationKind::CreateCopy(SourceLocation(), SourceLocation()); + InitializationSequence InitSeq(SemaRef, Entity, InitKind, DRE); + + ExprResult MemberInit = InitSeq.Perform(SemaRef, Entity, InitKind, DRE); + InitExprs.push_back(MemberInit.get()); + } + + void createSpecialMethodCall(const CXXRecordDecl *SpecialClass, Expr *Base, + const std::string &MethodName, + FieldDecl *Field) { + CXXMethodDecl *Method = getMethodByName(SpecialClass, MethodName); + assert(Method && + "The accessor/sampler/stream must have the __init method. Stream" + " must also have __finalize method"); + unsigned NumParams = Method->getNumParams(); + llvm::SmallVector ParamDREs(NumParams); + llvm::ArrayRef KernelParameters = + DeclCreator.getParamVarDeclsForCurrentField(); for (size_t I = 0; I < NumParams; ++I) { - ParmVarDecl *PD = InitMethod->getParamDecl(I); - CreateAndAddPrmDsc(Fld, PD->getType().getCanonicalType()); + QualType ParamType = KernelParameters[I]->getOriginalType(); + ParamDREs[I] = SemaRef.BuildDeclRefExpr(KernelParameters[I], ParamType, + VK_LValue, SourceLocation()); } - }; - - // Create parameter descriptor for accessor in case when it's wrapped with - // some class. - // TODO: Do we need support case when sampler is wrapped with some class or - // struct? - std::function - createParamDescForWrappedAccessors = - [&](const FieldDecl *Fld, const QualType &ArgTy) { - const auto *Wrapper = ArgTy->getAsCXXRecordDecl(); - for (const auto *WrapperFld : Wrapper->fields()) { - QualType FldType = WrapperFld->getType(); - if (FldType->isStructureOrClassType()) { - if (Util::isSyclAccessorType(FldType)) { - // Accessor field is found - create descriptor. - createSpecialSYCLObjParamDesc(WrapperFld, FldType); - } else if (Util::isSyclSpecConstantType(FldType)) { - // Don't try recursive search below. - } else { - // Field is some class or struct - recursively check for - // accessor fields. - createParamDescForWrappedAccessors(WrapperFld, FldType); - } - } - } - }; - - bool AllArgsAreValid = true; - // Run through kernel object fields and create corresponding kernel - // parameters descriptors. There are a several possible cases: - // - Kernel object field is a SYCL special object (SYCL accessor or SYCL - // sampler). These objects has a special initialization scheme - using - // __init method. - // - Kernel object field has a scalar type. In this case we should add - // kernel parameter with the same type. - // - Kernel object field has a structure or class type. Same handling as a - // scalar but we should check if this structure/class contains accessors - // and add parameter decriptor for them properly. - for (const auto *Fld : KernelObj->fields()) { - QualType ArgTy = Fld->getType(); - if (Util::isSyclAccessorType(ArgTy) || Util::isSyclSamplerType(ArgTy)) { - createSpecialSYCLObjParamDesc(Fld, ArgTy); - } else if (Util::isSyclSpecConstantType(ArgTy)) { - // Specialization constants are not added as arguments. - } else if (ArgTy->isStructureOrClassType()) { - if (Context.getLangOpts().SYCLStdLayoutKernelParams) { - if (!ArgTy->isStandardLayoutType()) { - Context.getDiagnostics().Report(Fld->getLocation(), - diag::err_sycl_non_std_layout_type) - << ArgTy; - AllArgsAreValid = false; - continue; - } - } - CXXRecordDecl *RD = - cast(ArgTy->getAs()->getDecl()); - if (!RD->hasTrivialCopyConstructor()) { - Context.getDiagnostics().Report( - Fld->getLocation(), - diag::err_sycl_non_trivially_copy_ctor_dtor_type) - << 0 << ArgTy; - AllArgsAreValid = false; - continue; - } - if (!RD->hasTrivialDestructor()) { - Context.getDiagnostics().Report( - Fld->getLocation(), - diag::err_sycl_non_trivially_copy_ctor_dtor_type) - << 1 << ArgTy; - AllArgsAreValid = false; - continue; - } + MemberExpr *SpecialObjME = BuildMemberExpr(Base, Field); + MemberExpr *MethodME = BuildMemberExpr(SpecialObjME, Method); + + QualType ResultTy = Method->getReturnType(); + ExprValueKind VK = Expr::getValueKindForType(ResultTy); + ResultTy = ResultTy.getNonLValueExprType(SemaRef.Context); + llvm::SmallVector ParamStmts; + const auto *Proto = cast(Method->getType()); + SemaRef.GatherArgumentsForCall(SourceLocation(), Method, Proto, 0, + ParamDREs, ParamStmts); + // [kernel_obj or wrapper object].accessor.__init(_ValueType*, + // range, range, id) + CXXMemberCallExpr *Call = CXXMemberCallExpr::Create( + SemaRef.Context, MethodME, ParamStmts, ResultTy, VK, SourceLocation()); + if (MethodName == FinalizeMethodName) + FinalizeStmts.push_back(Call); + else + BodyStmts.push_back(Call); + } - CreateAndAddPrmDsc(Fld, ArgTy); - - // Create descriptors for each accessor field in the class or struct - createParamDescForWrappedAccessors(Fld, ArgTy); - } else if (ArgTy->isReferenceType()) { - Context.getDiagnostics().Report( - Fld->getLocation(), diag::err_bad_kernel_param_type) << ArgTy; - AllArgsAreValid = false; - } else if (ArgTy->isPointerType()) { - // Pointer Arguments need to be in the global address space - QualType PointeeTy = ArgTy->getPointeeType(); - Qualifiers Quals = PointeeTy.getQualifiers(); - Quals.setAddressSpace(LangAS::opencl_global); - PointeeTy = - Context.getQualifiedType(PointeeTy.getUnqualifiedType(), Quals); - QualType ModTy = Context.getPointerType(PointeeTy); - - CreateAndAddPrmDsc(Fld, ModTy); - } else if (ArgTy->isScalarType()) { - CreateAndAddPrmDsc(Fld, ArgTy); - } else { - llvm_unreachable("Unsupported kernel parameter type"); + // FIXME Avoid creation of kernel obj clone. + // See https://github.com/intel/llvm/issues/1544 for details. + static VarDecl *createKernelObjClone(ASTContext &Ctx, DeclContext *DC, + CXXRecordDecl *KernelObj) { + TypeSourceInfo *TSInfo = + KernelObj->isLambda() ? KernelObj->getLambdaTypeInfo() : nullptr; + VarDecl *VD = VarDecl::Create( + Ctx, DC, SourceLocation(), SourceLocation(), KernelObj->getIdentifier(), + QualType(KernelObj->getTypeForDecl(), 0), TSInfo, SC_None); + + return VD; + } + + void handleSpecialType(FieldDecl *FD, QualType Ty) { + const auto *RecordDecl = Ty->getAsCXXRecordDecl(); + // Perform initialization only if it is field of kernel object + if (MemberExprBases.size() == 1) { + InitializedEntity Entity = + InitializedEntity::InitializeMember(FD, &VarEntity); + // Initialize with the default constructor. + InitializationKind InitKind = + InitializationKind::CreateDefault(SourceLocation()); + InitializationSequence InitSeq(SemaRef, Entity, InitKind, None); + ExprResult MemberInit = InitSeq.Perform(SemaRef, Entity, InitKind, None); + InitExprs.push_back(MemberInit.get()); } + createSpecialMethodCall(RecordDecl, MemberExprBases.back(), InitMethodName, + FD); } - return AllArgsAreValid; -} +public: + SyclKernelBodyCreator(Sema &S, SyclKernelDeclCreator &DC, + CXXRecordDecl *KernelObj, + FunctionDecl *KernelCallerFunc) + : SyclKernelFieldHandler(S), DeclCreator(DC), + KernelObjClone(createKernelObjClone(S.getASTContext(), + DC.getKernelDecl(), KernelObj)), + VarEntity(InitializedEntity::InitializeVariable(KernelObjClone)), + KernelObj(KernelObj), KernelCallerFunc(KernelCallerFunc) { + markParallelWorkItemCalls(); + + Stmt *DS = new (S.Context) DeclStmt(DeclGroupRef(KernelObjClone), + SourceLocation(), SourceLocation()); + BodyStmts.push_back(DS); + DeclRefExpr *KernelObjCloneRef = DeclRefExpr::Create( + S.Context, NestedNameSpecifierLoc(), SourceLocation(), KernelObjClone, + false, DeclarationNameInfo(), QualType(KernelObj->getTypeForDecl(), 0), + VK_LValue); + MemberExprBases.push_back(KernelObjCloneRef); + } + + ~SyclKernelBodyCreator() { + CompoundStmt *KernelBody = createKernelBody(); + DeclCreator.setBody(KernelBody); + } + + void handleSyclAccessorType(FieldDecl *FD, QualType Ty) final { + handleSpecialType(FD, Ty); + } + + void handleSyclAccessorType(const CXXBaseSpecifier &BS, QualType Ty) final { + // FIXME SYCL accessor should be usable as a base type + // See https://github.com/intel/llvm/issues/28. + } + + void handleSyclSamplerType(FieldDecl *FD, QualType Ty) final { + handleSpecialType(FD, Ty); + } + + void handleSyclStreamType(FieldDecl *FD, QualType Ty) final { + const auto *StreamDecl = Ty->getAsCXXRecordDecl(); + createExprForStructOrScalar(FD); + createSpecialMethodCall(StreamDecl, MemberExprBases.back(), InitMethodName, + FD); + createSpecialMethodCall(StreamDecl, MemberExprBases.back(), + FinalizeMethodName, FD); + } + + void handleSyclStreamType(const CXXBaseSpecifier &BS, QualType Ty) final { + // FIXME SYCL stream should be usable as a base type + // See https://github.com/intel/llvm/issues/1552 + } + + void handlePointerType(FieldDecl *FD, QualType FieldTy) final { + createExprForStructOrScalar(FD); + } + + void handleStructType(FieldDecl *FD, QualType FieldTy) final { + createExprForStructOrScalar(FD); + } + + void handleScalarType(FieldDecl *FD, QualType FieldTy) final { + createExprForStructOrScalar(FD); + } + + void enterStruct(const CXXRecordDecl *, FieldDecl *FD) final { + MemberExprBases.push_back(BuildMemberExpr(MemberExprBases.back(), FD)); + } + + void leaveStruct(const CXXRecordDecl *, FieldDecl *FD) final { + MemberExprBases.pop_back(); + } + + using SyclKernelFieldHandler::enterStruct; + using SyclKernelFieldHandler::leaveStruct; +}; -/// Adds necessary data describing given kernel to the integration header. -/// \param H the integration header object -/// \param Name kernel name -/// \param NameType type representing kernel name (first template argument -/// of single_task, parallel_for, etc) -/// \param KernelObjTy kernel object type -static void populateIntHeader(SYCLIntegrationHeader &H, const StringRef Name, - QualType NameType, CXXRecordDecl *KernelObjTy) { - - ASTContext &Ctx = KernelObjTy->getASTContext(); - const ASTRecordLayout &Layout = Ctx.getASTRecordLayout(KernelObjTy); - const std::string StableName = PredefinedExpr::ComputeName( - Ctx, PredefinedExpr::UniqueStableNameExpr, NameType); - H.startKernel(Name, NameType, StableName, KernelObjTy->getLocation()); - - auto populateHeaderForAccessor = [&](const QualType &ArgTy, uint64_t Offset) { - // The parameter is a SYCL accessor object. - // The Info field of the parameter descriptor for accessor contains - // two template parameters packed into an integer field: - // - target (e.g. global_buffer, constant_buffer, local); - // - dimension of the accessor. - const auto *AccTy = ArgTy->getAsCXXRecordDecl(); - assert(AccTy && "accessor must be of a record type"); - const auto *AccTmplTy = cast(AccTy); +class SyclKernelIntHeaderCreator + : public SyclKernelFieldHandler { + SYCLIntegrationHeader &Header; + const CXXRecordDecl *KernelLambda; + // Necessary to figure out the offset of the base class. + const CXXRecordDecl *CurStruct = nullptr; + int64_t CurOffset = 0; + + uint64_t getOffset(const CXXRecordDecl *RD) const { + assert(CurOffset && + "Cannot have a base class without setting the active struct"); + const ASTRecordLayout &Layout = + SemaRef.getASTContext().getASTRecordLayout(CurStruct); + return CurOffset + Layout.getBaseClassOffset(RD).getQuantity(); + } + uint64_t getOffset(const FieldDecl *FD) const { + return CurOffset + SemaRef.getASTContext().getFieldOffset(FD) / 8; + } + + void addParam(const FieldDecl *FD, QualType FieldTy, + SYCLIntegrationHeader::kernel_param_kind_t Kind) { + uint64_t Size = + SemaRef.getASTContext().getTypeSizeInChars(FieldTy).getQuantity(); + Header.addParamDesc(Kind, static_cast(Size), + static_cast(getOffset(FD))); + } + +public: + SyclKernelIntHeaderCreator(Sema &S, SYCLIntegrationHeader &H, + const CXXRecordDecl *KernelLambda, + QualType NameType, StringRef Name, + StringRef StableName) + : SyclKernelFieldHandler(S), Header(H), KernelLambda(KernelLambda) { + Header.startKernel(Name, NameType, StableName, KernelLambda->getLocation()); + } + + void handleSyclAccessorType(const CXXBaseSpecifier &BC, + QualType FieldTy) final { + const auto *AccTy = + cast(FieldTy->getAsRecordDecl()); + assert(AccTy->getTemplateArgs().size() >= 2 && + "Incorrect template args for Accessor Type"); int Dims = static_cast( - AccTmplTy->getTemplateArgs()[1].getAsIntegral().getExtValue()); - int Info = getAccessTarget(AccTmplTy) | (Dims << 11); - H.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info, Offset); - }; - - std::function - populateHeaderForWrappedAccessors = [&](const QualType &ArgTy, - uint64_t Offset) { - const auto *Wrapper = ArgTy->getAsCXXRecordDecl(); - for (const auto *WrapperFld : Wrapper->fields()) { - QualType FldType = WrapperFld->getType(); - if (FldType->isStructureOrClassType()) { - ASTContext &WrapperCtx = Wrapper->getASTContext(); - const ASTRecordLayout &WrapperLayout = - WrapperCtx.getASTRecordLayout(Wrapper); - // Get offset (in bytes) of the field in wrapper class or struct - uint64_t OffsetInWrapper = - WrapperLayout.getFieldOffset(WrapperFld->getFieldIndex()) / 8; - if (Util::isSyclAccessorType(FldType)) { - // This is an accesor - populate the header appropriately - populateHeaderForAccessor(FldType, Offset + OffsetInWrapper); - } else { - // This is an other class or struct - recursively search for an - // accessor field - populateHeaderForWrappedAccessors(FldType, - Offset + OffsetInWrapper); - } - } - } - }; - - for (const auto Fld : KernelObjTy->fields()) { - QualType ActualArgType; - QualType ArgTy = Fld->getType(); - - // Get offset in bytes - uint64_t Offset = Layout.getFieldOffset(Fld->getFieldIndex()) / 8; - - if (Util::isSyclAccessorType(ArgTy)) { - populateHeaderForAccessor(ArgTy, Offset); - } else if (Util::isSyclSamplerType(ArgTy)) { - // The parameter is a SYCL sampler object - const auto *SamplerTy = ArgTy->getAsCXXRecordDecl(); - assert(SamplerTy && "sampler must be of a record type"); - - CXXMethodDecl *InitMethod = getMethodByName(SamplerTy, InitMethodName); - assert(InitMethod && "sampler must have __init method"); - - // sampler __init method has only one argument - auto *FuncDecl = cast(InitMethod); - ParmVarDecl *SamplerArg = FuncDecl->getParamDecl(0); - assert(SamplerArg && "sampler __init method must have sampler parameter"); - uint64_t Sz = Ctx.getTypeSizeInChars(SamplerArg->getType()).getQuantity(); - H.addParamDesc(SYCLIntegrationHeader::kind_sampler, - static_cast(Sz), static_cast(Offset)); - } else if (ArgTy->isPointerType()) { - uint64_t Sz = Ctx.getTypeSizeInChars(Fld->getType()).getQuantity(); - H.addParamDesc(SYCLIntegrationHeader::kind_pointer, - static_cast(Sz), static_cast(Offset)); - } else if (Util::isSyclSpecConstantType(ArgTy)) { - // Add specialization constant ID to the header. - auto *TmplSpec = - cast(ArgTy->getAsCXXRecordDecl()); - const TemplateArgumentList *TemplateArgs = - &TmplSpec->getTemplateInstantiationArgs(); - // Get specialization constant ID type, which is the second template - // argument. - QualType SpecConstIDTy = TypeName::getFullyQualifiedType( - TemplateArgs->get(1).getAsType(), Ctx, true) - .getCanonicalType(); - const std::string SpecConstName = PredefinedExpr::ComputeName( - Ctx, PredefinedExpr::UniqueStableNameExpr, SpecConstIDTy); - H.addSpecConstant(SpecConstName, SpecConstIDTy); - // Spec constant lambda capture does not become a kernel argument. - } else if (ArgTy->isStructureOrClassType() || ArgTy->isScalarType()) { - // the parameter is an object of standard layout type or scalar; - // the check for standard layout is done elsewhere - uint64_t Sz = Ctx.getTypeSizeInChars(Fld->getType()).getQuantity(); - H.addParamDesc(SYCLIntegrationHeader::kind_std_layout, - static_cast(Sz), static_cast(Offset)); - - // check for accessor fields in structure or class and populate the - // integration header appropriately - if (ArgTy->isStructureOrClassType()) { - populateHeaderForWrappedAccessors(ArgTy, Offset); - } - } else { - llvm_unreachable("unsupported kernel parameter type"); - } + AccTy->getTemplateArgs()[1].getAsIntegral().getExtValue()); + int Info = getAccessTarget(AccTy) | (Dims << 11); + Header.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info, + getOffset(BC.getType()->getAsCXXRecordDecl())); } -} -// Creates a mangled kernel name for given kernel name type -static std::string constructKernelName(QualType KernelNameType, - MangleContext &MC) { - SmallString<256> Result; - llvm::raw_svector_ostream Out(Result); + void handleSyclAccessorType(FieldDecl *FD, QualType FieldTy) final { + const auto *AccTy = + cast(FieldTy->getAsRecordDecl()); + assert(AccTy->getTemplateArgs().size() >= 2 && + "Incorrect template args for Accessor Type"); + int Dims = static_cast( + AccTy->getTemplateArgs()[1].getAsIntegral().getExtValue()); + int Info = getAccessTarget(AccTy) | (Dims << 11); + Header.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info, + getOffset(FD)); + } - MC.mangleTypeName(KernelNameType, Out); - return std::string(Out.str()); -} + void handleSyclSamplerType(FieldDecl *FD, QualType FieldTy) final { + const auto *SamplerTy = FieldTy->getAsCXXRecordDecl(); + assert(SamplerTy && "Sampler type must be a C++ record type"); + CXXMethodDecl *InitMethod = getMethodByName(SamplerTy, InitMethodName); + assert(InitMethod && "sampler must have __init method"); -static FunctionDecl * -CreateOpenCLKernelDeclaration(ASTContext &Context, StringRef Name, - ArrayRef ParamDescs) { - - DeclContext *DC = Context.getTranslationUnitDecl(); - QualType RetTy = Context.VoidTy; - SmallVector ArgTys; - - // Extract argument types from the descriptor array: - std::transform( - ParamDescs.begin(), ParamDescs.end(), std::back_inserter(ArgTys), - [](const ParamDesc &PD) -> QualType { return std::get<0>(PD); }); - FunctionProtoType::ExtProtoInfo Info(CC_OpenCLKernel); - QualType FuncTy = Context.getFunctionType(RetTy, ArgTys, Info); - DeclarationName DN = DeclarationName(&Context.Idents.get(Name)); - - FunctionDecl *OpenCLKernel = FunctionDecl::Create( - Context, DC, SourceLocation(), SourceLocation(), DN, FuncTy, - Context.getTrivialTypeSourceInfo(RetTy), SC_None); - - llvm::SmallVector Params; - int i = 0; - for (const auto &PD : ParamDescs) { - auto P = ParmVarDecl::Create(Context, OpenCLKernel, SourceLocation(), - SourceLocation(), std::get<1>(PD), - std::get<0>(PD), std::get<2>(PD), SC_None, 0); - P->setScopeInfo(0, i++); - P->setIsUsed(); - Params.push_back(P); - } - OpenCLKernel->setParams(Params); - - OpenCLKernel->addAttr(OpenCLKernelAttr::CreateImplicit(Context)); - OpenCLKernel->addAttr(AsmLabelAttr::CreateImplicit(Context, Name)); - OpenCLKernel->addAttr(ArtificialAttr::CreateImplicit(Context)); - - // Add kernel to translation unit to see it in AST-dump - DC->addDecl(OpenCLKernel); - return OpenCLKernel; -} + // sampler __init method has only one argument + const ParmVarDecl *SamplerArg = InitMethod->getParamDecl(0); + assert(SamplerArg && "sampler __init method must have sampler parameter"); + + addParam(FD, SamplerArg->getType(), SYCLIntegrationHeader::kind_sampler); + } + + void handleSyclSpecConstantType(FieldDecl *FD, QualType FieldTy) final { + const TemplateArgumentList &TemplateArgs = + cast(FieldTy->getAsRecordDecl()) + ->getTemplateInstantiationArgs(); + assert(TemplateArgs.size() == 2 && + "Incorrect template args for Accessor Type"); + // Get specialization constant ID type, which is the second template + // argument. + QualType SpecConstIDTy = + TypeName::getFullyQualifiedType(TemplateArgs.get(1).getAsType(), + SemaRef.getASTContext(), true) + .getCanonicalType(); + const std::string SpecConstName = PredefinedExpr::ComputeName( + SemaRef.getASTContext(), PredefinedExpr::UniqueStableNameType, + SpecConstIDTy); + Header.addSpecConstant(SpecConstName, SpecConstIDTy); + } + + void handlePointerType(FieldDecl *FD, QualType FieldTy) final { + addParam(FD, FieldTy, SYCLIntegrationHeader::kind_pointer); + } + void handleStructType(FieldDecl *FD, QualType FieldTy) final { + addParam(FD, FieldTy, SYCLIntegrationHeader::kind_std_layout); + } + void handleScalarType(FieldDecl *FD, QualType FieldTy) final { + addParam(FD, FieldTy, SYCLIntegrationHeader::kind_std_layout); + } + + void handleSyclStreamType(FieldDecl *FD, QualType FieldTy) final { + addParam(FD, FieldTy, SYCLIntegrationHeader::kind_std_layout); + } + void handleSyclStreamType(const CXXBaseSpecifier &BC, + QualType FieldTy) final { + // FIXME SYCL stream should be usable as a base type + // See https://github.com/intel/llvm/issues/1552 + } + + // Keep track of the current struct offset. + void enterStruct(const CXXRecordDecl *RD, FieldDecl *FD) final { + CurStruct = FD->getType()->getAsCXXRecordDecl(); + CurOffset += SemaRef.getASTContext().getFieldOffset(FD) / 8; + } + + void leaveStruct(const CXXRecordDecl *RD, FieldDecl *FD) final { + CurStruct = RD; + CurOffset -= SemaRef.getASTContext().getFieldOffset(FD) / 8; + } + + void enterStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS) final { + CurStruct = BS.getType()->getAsCXXRecordDecl(); + const ASTRecordLayout &Layout = + SemaRef.getASTContext().getASTRecordLayout(RD); + CurOffset += Layout.getBaseClassOffset(BS.getType()->getAsCXXRecordDecl()) + .getQuantity(); + } + + void leaveStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS) final { + CurStruct = RD; + const ASTRecordLayout &Layout = + SemaRef.getASTContext().getASTRecordLayout(RD); + CurOffset -= Layout.getBaseClassOffset(BS.getType()->getAsCXXRecordDecl()) + .getQuantity(); + } +}; +} // namespace // Generates the OpenCL kernel using KernelCallerFunc (kernel caller // function) defined is SYCL headers. @@ -1238,52 +1424,32 @@ CreateOpenCLKernelDeclaration(ASTContext &Context, StringRef Name, // void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, MangleContext &MC) { - CXXRecordDecl *LE = getKernelObjectType(KernelCallerFunc); - assert(LE && "invalid kernel caller"); - - // Build list of kernel arguments - llvm::SmallVector ParamDescs; - if (!buildArgTys(getASTContext(), LE, ParamDescs)) - return; - - // Extract name from kernel caller parameters and mangle it. - const TemplateArgumentList *TemplateArgs = - KernelCallerFunc->getTemplateSpecializationArgs(); - assert(TemplateArgs && "No template argument info"); - QualType KernelNameType = TypeName::getFullyQualifiedType( - TemplateArgs->get(0).getAsType(), getASTContext(), true); - - std::string Name; - // TODO SYCLIntegrationHeader also computes a unique stable name. It should - // probably lose this responsibility and only use the name provided here. - if (getLangOpts().SYCLUnnamedLambda) - Name = PredefinedExpr::ComputeName( - getASTContext(), PredefinedExpr::UniqueStableNameExpr, KernelNameType); - else - Name = constructKernelName(KernelNameType, MC); - - // TODO Maybe don't emit integration header inside the Sema? - populateIntHeader(getSyclIntegrationHeader(), Name, KernelNameType, LE); - - FunctionDecl *OpenCLKernel = - CreateOpenCLKernelDeclaration(getASTContext(), Name, ParamDescs); - - ContextRAII FuncContext(*this, OpenCLKernel); - - // Let's copy source location of a functor/lambda to emit nicer diagnostics - OpenCLKernel->setLocation(LE->getLocation()); - - // If the source function is implicitly inline, the kernel should be marked - // such as well. This allows the kernel to be ODR'd if there are multiple uses - // in different translation units. - OpenCLKernel->setImplicitlyInline(KernelCallerFunc->isInlined()); + // The first argument to the KernelCallerFunc is the lambda object. + CXXRecordDecl *KernelLambda = getKernelObjectType(KernelCallerFunc); + assert(KernelLambda && "invalid kernel caller"); + + // Calculate both names, since Integration headers need both. + std::string CalculatedName, StableName; + std::tie(CalculatedName, StableName) = + constructKernelName(*this, KernelCallerFunc, MC); + StringRef KernelName(getLangOpts().SYCLUnnamedLambda ? StableName + : CalculatedName); + + SyclKernelFieldChecker checker(*this); + SyclKernelDeclCreator kernel_decl(*this, checker, KernelName, + KernelLambda->getLocation(), + KernelCallerFunc->isInlined()); + SyclKernelBodyCreator kernel_body(*this, kernel_decl, KernelLambda, + KernelCallerFunc); + SyclKernelIntHeaderCreator int_header( + *this, getSyclIntegrationHeader(), KernelLambda, + calculateKernelNameType(Context, KernelCallerFunc), KernelName, + StableName); ConstructingOpenCLKernel = true; - CompoundStmt *OpenCLKernelBody = - CreateOpenCLKernelBody(*this, KernelCallerFunc, OpenCLKernel); + VisitRecordFields(KernelLambda->fields(), checker, kernel_decl, kernel_body, + int_header); ConstructingOpenCLKernel = false; - OpenCLKernel->setBody(OpenCLKernelBody); - addSyclDeviceDecl(OpenCLKernel); } void Sema::MarkDevice(void) { diff --git a/clang/test/CodeGenSYCL/integration_header.cpp b/clang/test/CodeGenSYCL/integration_header.cpp index ad788c646c8e8..84b35578f48e6 100644 --- a/clang/test/CodeGenSYCL/integration_header.cpp +++ b/clang/test/CodeGenSYCL/integration_header.cpp @@ -1,4 +1,4 @@ -// RUN: %clang -I %S/Inputs -fsycl-device-only -Xclang -fsycl-int-header=%t.h %s -c -o %T/kernel.spv +// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s -fsyntax-only // RUN: FileCheck -input-file=%t.h %s // // CHECK: #include @@ -21,6 +21,7 @@ // CHECK-NEXT: "_ZTSN16second_namespace13second_kernelIcEE", // CHECK-NEXT: "_ZTS12third_kernelILi1Ei5pointIZ4mainE1XEE" // CHECK-NEXT: "_ZTS13fourth_kernelIJN15template_arg_ns14namespaced_argILi1EEEEE" +// CHECK-NEXT: "_ZTSZ4mainE16accessor_in_base" // CHECK-NEXT: }; // // CHECK: static constexpr @@ -45,6 +46,13 @@ // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 }, // CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 6112, 4 }, // CHECK-EMPTY: +// CHECK-NEXT: //--- _ZTSZ4mainE16accessor_in_base +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 64, 0 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 8 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 24 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 40 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 52 }, +// CHECK-EMPTY: // CHECK-NEXT: }; // // CHECK: template <> struct KernelInfo { @@ -77,6 +85,28 @@ struct namespaced_arg {}; template class fourth_kernel; +namespace accessor_in_base { +struct other_base { + int i; +}; +struct base { + int i, j; + cl::sycl::accessor acc; +}; + +struct base2 : other_base, + cl::sycl::accessor { + int i; + cl::sycl::accessor acc; +}; + +struct captured : base, base2 { + cl::sycl::accessor acc; + void use() const {} +}; + +}; // namespace accessor_in_base + int main() { cl::sycl::accessor acc1; @@ -121,5 +151,11 @@ int main() { } }); + // FIXME: We cannot use the member-capture because all the handlers except the + // integration header handler in SemaSYCL don't handle base types right. + accessor_in_base::captured c; + kernel_single_task([c]() { + }); + return 0; }