diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 6200e5560e49..8d021f7e00d4 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -11266,7 +11266,7 @@ class Sema { return *SyclIntHeader.get(); } - void ConstructSYCLKernel(FunctionDecl *KernelCallerFunc); + void ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc); void MarkDevice(void); }; diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 28d8bf5119ca..516d37dc2b05 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -26,8 +26,6 @@ using namespace clang; -typedef llvm::DenseMap DeclMap; - using KernelParamKind = SYCLIntegrationHeader::kernel_param_kind_t; enum target { @@ -376,67 +374,68 @@ class MarkDeviceFunction : public RecursiveASTVisitor { class KernelBodyTransform : public TreeTransform { public: - KernelBodyTransform(llvm::DenseMap &Map, + KernelBodyTransform(std::pair &MPair, Sema &S) - : TreeTransform(S), DMap(Map), SemaRef(S) {} + : TreeTransform(S), MappingPair(MPair), SemaRef(S) {} bool AlwaysRebuild() { return true; } ExprResult TransformDeclRefExpr(DeclRefExpr *DRE) { auto Ref = dyn_cast(DRE->getDecl()); - if (Ref) { - auto NewDecl = DMap[Ref]; - if (NewDecl) { - return DeclRefExpr::Create( - SemaRef.getASTContext(), DRE->getQualifierLoc(), - DRE->getTemplateKeywordLoc(), NewDecl, false, DRE->getNameInfo(), - NewDecl->getType(), DRE->getValueKind()); - } + if (Ref && Ref == MappingPair.first) { + auto NewDecl = MappingPair.second; + return DeclRefExpr::Create( + SemaRef.getASTContext(), DRE->getQualifierLoc(), + DRE->getTemplateKeywordLoc(), NewDecl, false, DRE->getNameInfo(), + NewDecl->getType(), DRE->getValueKind()); } return DRE; } private: - DeclMap DMap; + std::pair MappingPair; Sema &SemaRef; }; -static FunctionDecl *CreateSYCLKernelFunction(ASTContext &Context, - StringRef Name, - ArrayRef ParamDescs) { +static FunctionDecl * +CreateOpenCLKernelDeclaration(ASTContext &Context, StringRef Name, + ArrayRef ParamDescs) { DeclContext *DC = Context.getTranslationUnitDecl(); - FunctionProtoType::ExtProtoInfo Info(CC_OpenCLKernel); QualType RetTy = Context.VoidTy; SmallVector ArgTys; - // extract argument types from the descriptor array: + + // 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 *SYCLKernel = FunctionDecl::Create( + + 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, SYCLKernel, SourceLocation(), + 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); } - SYCLKernel->setParams(Params); + OpenCLKernel->setParams(Params); - SYCLKernel->addAttr(SYCLDeviceAttr::CreateImplicit(Context)); - SYCLKernel->addAttr(OpenCLKernelAttr::CreateImplicit(Context)); - SYCLKernel->addAttr(AsmLabelAttr::CreateImplicit(Context, Name)); - SYCLKernel->addAttr(ArtificialAttr::CreateImplicit(Context)); + OpenCLKernel->addAttr(SYCLDeviceAttr::CreateImplicit(Context)); + OpenCLKernel->addAttr(OpenCLKernelAttr::CreateImplicit(Context)); + OpenCLKernel->addAttr(AsmLabelAttr::CreateImplicit(Context, Name)); + OpenCLKernel->addAttr(ArtificialAttr::CreateImplicit(Context)); - // To see kernel in AST-dump. - DC->addDecl(SYCLKernel); - return SYCLKernel; + // Add kernel to translation unit to see it in AST-dump + DC->addDecl(OpenCLKernel); + return OpenCLKernel; } /// Return __init method static CXXMethodDecl *getInitMethod(const CXXRecordDecl *CRD) { @@ -449,79 +448,77 @@ static CXXMethodDecl *getInitMethod(const CXXRecordDecl *CRD) { return InitMethod; } -static CompoundStmt * -CreateSYCLKernelBody(Sema &S, FunctionDecl *KernelCallerFunc, DeclContext *DC) { +// 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) { llvm::SmallVector BodyStmts; CXXRecordDecl *LC = getKernelObjectType(KernelCallerFunc); assert(LC && "Kernel object must be available"); 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, DC, SourceLocation(), SourceLocation(), LC->getIdentifier(), - QualType(LC->getTypeForDecl(), 0), TSInfo, SC_None); + 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 CloneRef = + auto KernelObjCloneRef = DeclRefExpr::Create(S.Context, NestedNameSpecifierLoc(), SourceLocation(), KernelObjClone, false, DeclarationNameInfo(), QualType(LC->getTypeForDecl(), 0), VK_LValue); - auto TargetFunc = dyn_cast(DC); - assert(TargetFunc && "Not FunctionDecl"); - auto TargetFuncParam = - TargetFunc->param_begin(); // Iterator to ParamVarDecl (VarDecl) - if (TargetFuncParam) { + + auto KernelFuncDecl = cast(KernelDecl); + auto KernelFuncParam = + KernelFuncDecl->param_begin(); // Iterator to ParamVarDecl (VarDecl) + if (KernelFuncParam) { for (auto Field : LC->fields()) { - auto getExprForPointer = [](Sema &S, const QualType ¶mTy, - DeclRefExpr *DRE) { - // C++ address space attribute != OpenCL address space attribute - Expr *qualifiersCast = ImplicitCastExpr::Create( - S.Context, paramTy, CK_NoOp, DRE, nullptr, VK_LValue); - Expr *Res = - ImplicitCastExpr::Create(S.Context, paramTy, CK_LValueToRValue, - qualifiersCast, nullptr, VK_RValue); - return Res; - }; - auto getExprForRangeOrOffset = [](Sema &S, const QualType ¶mTy, - DeclRefExpr *DRE) { - Expr *Res = ImplicitCastExpr::Create(S.Context, paramTy, CK_NoOp, DRE, - nullptr, VK_RValue); + auto getExprForKernelParameter = [](Sema &S, const QualType ¶mTy, + DeclRefExpr *DRE) { + Expr *Res = ImplicitCastExpr::Create( + S.Context, paramTy, CK_LValueToRValue, DRE, nullptr, VK_RValue); return Res; }; - auto getExprForAccessorInit = [&](const QualType ¶mTy, - FieldDecl *Field, - const CXXRecordDecl *CRD, Expr *Base) { - // Since this is an accessor next 4 TargetFuncParams including current - // should be set in __init method: _ValueType*, range, range, - // id - const size_t NumParams = 4; - llvm::SmallVector ParamDREs(NumParams); - auto TFP = TargetFuncParam; - for (size_t I = 0; I < NumParams; ++TFP, ++I) { - QualType ParamType = (*TFP)->getOriginalType(); + // 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) { + // All special SYCL objects must have __init method + CXXMethodDecl *InitMethod = getInitMethod(CRD); + assert(InitMethod && + "The accessor/sampler must have the __init method"); + unsigned NumParams = InitMethod->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(), *TFP, + S.Context, NestedNameSpecifierLoc(), SourceLocation(), *KFP, false, DeclarationNameInfo(), ParamType, VK_LValue); } - std::advance(TargetFuncParam, NumParams - 1); + std::advance(KernelFuncParam, NumParams - 1); DeclAccessPair FieldDAP = DeclAccessPair::make(Field, AS_none); - // [kenrel_obj or wrapper object].accessor - auto AccessorME = MemberExpr::Create( - S.Context, Base, false, SourceLocation(), - NestedNameSpecifierLoc(), SourceLocation(), Field, FieldDAP, + // [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); - CXXMethodDecl *InitMethod = getInitMethod(CRD); - assert(InitMethod && "The accessor must have the __init method"); - - // [kenrel_obj or wrapper object].accessor.__init + // [kernel_obj or wrapper object].special_obj.__init DeclAccessPair MethodDAP = DeclAccessPair::make(InitMethod, AS_none); auto ME = MemberExpr::Create( - S.Context, AccessorME, false, SourceLocation(), + S.Context, SpecialObjME, false, SourceLocation(), NestedNameSpecifierLoc(), SourceLocation(), InitMethod, MethodDAP, InitMethod->getNameInfo(), nullptr, InitMethod->getType(), VK_LValue, OK_Ordinary); @@ -535,17 +532,14 @@ CreateSYCLKernelBody(Sema &S, FunctionDecl *KernelCallerFunc, DeclContext *DC) { // __init needs four parameter auto ParamItr = InitMethod->param_begin(); + // kernel_parameters - llvm::SmallVector ParamStmts; - ParamStmts.push_back(getExprForPointer( - S, (*(ParamItr++))->getOriginalType(), ParamDREs[0])); - ParamStmts.push_back(getExprForRangeOrOffset( - S, ((*ParamItr++))->getOriginalType(), ParamDREs[1])); - ParamStmts.push_back(getExprForRangeOrOffset( - S, ((*ParamItr++))->getOriginalType(), ParamDREs[2])); - ParamStmts.push_back(getExprForRangeOrOffset( - S, ((*ParamItr++))->getOriginalType(), ParamDREs[3])); - // [kenrel_obj or wrapper object].accessor.__init(_ValueType*, + llvm::SmallVector ParamStmts; + for (size_t I = 0; I < NumParams; ++I) { + ParamStmts.push_back(getExprForKernelParameter( + S, (*(ParamItr++))->getOriginalType(), ParamDREs[I])); + } + // [kernel_obj or wrapper object].accessor.__init(_ValueType*, // range, range, id) CXXMemberCallExpr *Call = CXXMemberCallExpr::Create( S.Context, ME, ParamStmts, ResultTy, VK, SourceLocation()); @@ -566,9 +560,9 @@ CreateSYCLKernelBody(Sema &S, FunctionDecl *KernelCallerFunc, DeclContext *DC) { // 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. - TargetFuncParam++; - getExprForAccessorInit(FldType, WrapperFld, WrapperFldCRD, - Base); + KernelFuncParam++; + getExprForSpecialSYCLObj(FldType, WrapperFld, WrapperFldCRD, + Base); } else { // Field is a structure or class so change the wrapper object // and recursively search for accessor field. @@ -587,70 +581,34 @@ CreateSYCLKernelBody(Sema &S, FunctionDecl *KernelCallerFunc, DeclContext *DC) { } }; + // 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 using binary '=' operator. + // - 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(); - if (Util::isSyclAccessorType(FieldType)) { - getExprForAccessorInit(FieldType, Field, CRD, CloneRef); - } else if (CRD && Util::isSyclSamplerType(FieldType)) { - - // Sampler has only one TargetFuncParam, which should be set in - // __init method: _ValueType - const size_t NumParams = 1; - llvm::SmallVector ParamDREs(NumParams); - auto TFP = TargetFuncParam; - QualType ParamType = (*TFP)->getOriginalType(); - ParamDREs[0] = DeclRefExpr::Create( - S.Context, NestedNameSpecifierLoc(), SourceLocation(), *TFP, - false, DeclarationNameInfo(), ParamType, VK_LValue); - DeclAccessPair FieldDAP = DeclAccessPair::make(Field, AS_none); - - // kernel_obj.sampler - auto SamplerME = MemberExpr::Create( - S.Context, CloneRef, false, SourceLocation(), - NestedNameSpecifierLoc(), SourceLocation(), Field, FieldDAP, - DeclarationNameInfo(Field->getDeclName(), SourceLocation()), - nullptr, Field->getType(), VK_LValue, OK_Ordinary); - - CXXMethodDecl *InitMethod = getInitMethod(CRD); - assert(InitMethod && "The sampler must have the __init method"); - - // kernel_obj.sampler.__init - DeclAccessPair MethodDAP = DeclAccessPair::make(InitMethod, AS_none); - auto ME = MemberExpr::Create( - S.Context, SamplerME, false, SourceLocation(), - NestedNameSpecifierLoc(), SourceLocation(), InitMethod, MethodDAP, - InitMethod->getNameInfo(), nullptr, InitMethod->getType(), - VK_LValue, OK_Ordinary); - - // Not referenced -> not emitted - S.MarkFunctionReferenced(SourceLocation(), InitMethod, true); - - QualType ResultTy = InitMethod->getReturnType(); - ExprValueKind VK = Expr::getValueKindForType(ResultTy); - ResultTy = ResultTy.getNonLValueExprType(S.Context); - - // __init needs one parameter - auto ParamItr = InitMethod->param_begin(); - // kernel_parameters - llvm::SmallVector ParamStmts; - ParamStmts.push_back(getExprForPointer( - S, (*ParamItr)->getOriginalType(), ParamDREs[0])); - CXXMemberCallExpr *Call = CXXMemberCallExpr::Create( - S.Context, ME, ParamStmts, ResultTy, VK, SourceLocation()); - BodyStmts.push_back(Call); + if (Util::isSyclAccessorType(FieldType) || + Util::isSyclSamplerType(FieldType)) { + getExprForSpecialSYCLObj(FieldType, Field, CRD, KernelObjCloneRef); } else if (CRD || FieldType->isScalarType()) { - // If field have built-in or a structure/class type just initialize + // If field has built-in or a structure/class type just initialize // this field with corresponding kernel argument using '=' binary // operator. The structure/class type must be copy assignable - this // holds because SYCL kernel lambdas capture arguments by copy. - QualType ParamType = (*TargetFuncParam)->getOriginalType(); + QualType ParamType = (*KernelFuncParam)->getOriginalType(); auto DRE = DeclRefExpr::Create(S.Context, NestedNameSpecifierLoc(), - SourceLocation(), *TargetFuncParam, false, + SourceLocation(), *KernelFuncParam, false, DeclarationNameInfo(), ParamType, VK_LValue); DeclAccessPair FieldDAP = DeclAccessPair::make(Field, AS_none); auto Lhs = MemberExpr::Create( - S.Context, CloneRef, false, SourceLocation(), + S.Context, KernelObjCloneRef, false, SourceLocation(), NestedNameSpecifierLoc(), SourceLocation(), Field, FieldDAP, DeclarationNameInfo(Field->getDeclName(), SourceLocation()), nullptr, Field->getType(), VK_LValue, OK_Ordinary); @@ -668,26 +626,28 @@ CreateSYCLKernelBody(Sema &S, FunctionDecl *KernelCallerFunc, DeclContext *DC) { if (CRD) getExprForWrappedAccessorInit(CRD, Lhs); } else { - llvm_unreachable("unsupported field type"); + llvm_unreachable("Unsupported field type"); } - TargetFuncParam++; + KernelFuncParam++; } } - // In function from headers lambda is function parameter, we need - // to replace all refs to this lambda with our vardecl. - // I used TreeTransform here, but I'm not sure that it is good solution - // Also I used map and I'm not sure about it too. - // TODO SYCL review the above design concerns + + // 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(); - DeclMap DMap; ParmVarDecl *KernelObjParam = *(KernelCallerFunc->param_begin()); + // DeclRefExpr with valid source location but with decl which is not marked // as used is invalid. KernelObjClone->setIsUsed(); - DMap[KernelObjParam] = KernelObjClone; - // Without PushFunctionScope I had segfault. Maybe we also need to do pop. + std::pair MappingPair; + MappingPair.first = KernelObjParam; + MappingPair.second = KernelObjClone; + + // Function scope might be empty, so we do push S.PushFunctionScope(); - KernelBodyTransform KBT(DMap, S); + KernelBodyTransform KBT(MappingPair, S); Stmt *NewBody = KBT.TransformStmt(FunctionBody).get(); BodyStmts.push_back(NewBody); return CompoundStmt::Create(S.Context, BodyStmts, SourceLocation(), @@ -711,69 +671,42 @@ static target getAccessTarget(const ClassTemplateSpecializationDecl *AccTy) { 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. static void buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj, SmallVectorImpl &ParamDescs) { const LambdaCapture *Cpt = KernelObj->captures_begin(); auto CreateAndAddPrmDsc = [&](const FieldDecl *Fld, const QualType &ArgType) { - // create a parameter descriptor and append it to the result + // Create a parameter descriptor and append it to the result ParamDescs.push_back(makeParamDesc(Fld, ArgType)); }; - auto createAccessorParamDesc = [&](const FieldDecl *Fld, - const QualType &ArgTy) { - // the parameter is a SYCL accessor object + // Creates a parameter descriptor for SYCL special object - SYCL accessor or + // sampler. + // 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 && "accessor must be of a record type"); - const auto *TemplateDecl = - cast(RecordDecl); - // First accessor template parameter - data type - QualType PointeeType = TemplateDecl->getTemplateArgs()[0].getAsType(); - // Fourth parameter - access target - target AccessTarget = getAccessTarget(TemplateDecl); - Qualifiers Quals = PointeeType.getQualifiers(); - // TODO: Support all access targets - switch (AccessTarget) { - case target::global_buffer: - Quals.setAddressSpace(LangAS::opencl_global); - break; - case target::constant_buffer: - Quals.setAddressSpace(LangAS::opencl_constant); - break; - case target::local: - Quals.setAddressSpace(LangAS::opencl_local); - break; - default: - llvm_unreachable("Unsupported access target"); - } - PointeeType = - Context.getQualifiedType(PointeeType.getUnqualifiedType(), Quals); - QualType PointerType = Context.getPointerType(PointeeType); - - CreateAndAddPrmDsc(Fld, PointerType); + assert(RecordDecl && "Special SYCL object must be of a record type"); CXXMethodDecl *InitMethod = getInitMethod(RecordDecl); - assert(InitMethod && "accessor must have __init method"); - - // Expected accessor __init method has four parameters - // void __init(_ValueType *Ptr, range AccessRange, - // range MemRange, id Offset) - auto *FuncDecl = cast(InitMethod); - ParmVarDecl *AccessRangeFld = FuncDecl->getParamDecl(1); - ParmVarDecl *MemRangeFld = FuncDecl->getParamDecl(2); - ParmVarDecl *OffsetFld = FuncDecl->getParamDecl(3); - - assert(AccessRangeFld && - "The accessor __init method must contain the AccessRange parameter"); - assert(MemRangeFld && - "The accessor __init method must contain the MemRange parameter"); - assert(OffsetFld && - "The accessor __init method must contain the Offset parameter"); - - CreateAndAddPrmDsc(Fld, AccessRangeFld->getType()); - CreateAndAddPrmDsc(Fld, MemRangeFld->getType()); - CreateAndAddPrmDsc(Fld, OffsetFld->getType()); + assert(InitMethod && "The accessor/sampler must have the __init method"); + unsigned NumParams = InitMethod->getNumParams(); + for (size_t I = 0; I < NumParams; ++I) { + ParmVarDecl *PD = InitMethod->getParamDecl(I); + CreateAndAddPrmDsc(Fld, PD->getType().getCanonicalType()); + } }; + // 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) { @@ -783,7 +716,7 @@ static void buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj, if (FldType->isStructureOrClassType()) { if (Util::isSyclAccessorType(FldType)) { // accessor field is found - create descriptor - createAccessorParamDesc(WrapperFld, FldType); + createSpecialSYCLObjParamDesc(WrapperFld, FldType); } else { // field is some class or struct - recursively check for // accessor fields @@ -793,25 +726,24 @@ static void buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj, } }; + // 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)) { - createAccessorParamDesc(Fld, ArgTy); - } else if (Util::isSyclSamplerType(ArgTy)) { - // the parameter is a SYCL sampler object - const auto *RecordDecl = ArgTy->getAsCXXRecordDecl(); - assert(RecordDecl && "sampler must be of a record type"); - - CXXMethodDecl *InitMethod = getInitMethod(RecordDecl); - assert(InitMethod && "sampler must have __init method"); - - // sampler __init method has only one parameter - auto *FuncDecl = cast(InitMethod); - ParmVarDecl *SamplerArg = FuncDecl->getParamDecl(0); - assert(SamplerArg && "sampler __init method must have sampler parameter"); - - CreateAndAddPrmDsc(Fld, SamplerArg->getType()); + if (Util::isSyclAccessorType(ArgTy) || Util::isSyclSamplerType(ArgTy)) { + createSpecialSYCLObjParamDesc(Fld, ArgTy); } else if (ArgTy->isStructureOrClassType()) { + // SYCL v1.2.1 s4.8.10 p5: + // C++ non-standard layout values must not be passed as arguments to a + // kernel that is compiled for a device. if (!ArgTy->isStandardLayoutType()) { const DeclaratorDecl *V = Cpt ? cast(Cpt->getCapturedVar()) @@ -819,15 +751,14 @@ static void buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj, KernelObj->getASTContext().getDiagnostics().Report( V->getLocation(), diag::err_sycl_non_std_layout_type); } - // structure or class typed parameter - the same handling as a scalar CreateAndAddPrmDsc(Fld, ArgTy); - // create descriptors for each accessor field in the class or struct + + // Create descriptors for each accessor field in the class or struct createParamDescForWrappedAccessors(Fld, ArgTy); } else if (ArgTy->isScalarType()) { - // scalar typed parameter CreateAndAddPrmDsc(Fld, ArgTy); } else { - llvm_unreachable("unsupported kernel parameter type"); + llvm_unreachable("Unsupported kernel parameter type"); } } } @@ -847,18 +778,18 @@ static void populateIntHeader(SYCLIntegrationHeader &H, const StringRef Name, H.startKernel(Name, NameType); 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 thid 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); - int Dims = static_cast( - AccTmplTy->getTemplateArgs()[1].getAsIntegral().getExtValue()); - int Info = getAccessTarget(AccTmplTy) | (Dims << 11); - H.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info, 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); + int Dims = static_cast( + AccTmplTy->getTemplateArgs()[1].getAsIntegral().getExtValue()); + int Info = getAccessTarget(AccTmplTy) | (Dims << 11); + H.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info, Offset); }; std::function @@ -951,32 +882,57 @@ static std::string constructKernelName(QualType KernelNameType, return Out.str(); } -void Sema::ConstructSYCLKernel(FunctionDecl *KernelCallerFunc) { - // TODO: Case when kernel is functor +// Generates the OpenCL kernel using KernelCallerFunc (kernel caller +// function) defined is SYCL headers. +// Generated OpenCL kernel contains the body of the kernel caller function, +// receives OpenCL like parameters and additionally does some manipulation to +// initialize captured lambda/functor fields with these parameters. +// SYCL runtime marks kernel caller function with sycl_kernel attribute. +// To be able to generate OpenCL kernel from KernelCallerFunc we put +// the following requirements to the function which SYCL runtime can mark with +// sycl_kernel attribute: +// - Must be template function with at least two template parameters. +// First parameter must represent "unique kernel name" +// Second parameter must be the function object type +// - Must have only one function parameter - function object. +// +// Example of kernel caller function: +// template +// __attribute__((sycl_kernel)) void kernel_caller_function(KernelType +// KernelFuncObj) { +// KernelFuncObj(); +// } +// +// +void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc) { CXXRecordDecl *LE = getKernelObjectType(KernelCallerFunc); assert(LE && "invalid kernel caller"); + + // Build list of kernel arguments llvm::SmallVector ParamDescs; buildArgTys(getASTContext(), LE, ParamDescs); - // Get Name for our kernel. + + // Extract name from kernel caller parameters and mangle it. const TemplateArgumentList *TemplateArgs = KernelCallerFunc->getTemplateSpecializationArgs(); assert(TemplateArgs && "No template argument info"); - // The first template argument always describes the kernel name - whether - // it is lambda or functor. QualType KernelNameType = TypeName::getFullyQualifiedType( TemplateArgs->get(0).getAsType(), getASTContext(), true); std::string Name = constructKernelName(KernelNameType, getASTContext()); + + // TODO Maybe don't emit integration header inside the Sema? populateIntHeader(getSyclIntegrationHeader(), Name, KernelNameType, LE); - FunctionDecl *SYCLKernel = - CreateSYCLKernelFunction(getASTContext(), Name, ParamDescs); + + FunctionDecl *OpenCLKernel = + CreateOpenCLKernelDeclaration(getASTContext(), Name, ParamDescs); // Let's copy source location of a functor/lambda to emit nicer diagnostics - SYCLKernel->setLocation(LE->getLocation()); + OpenCLKernel->setLocation(LE->getLocation()); - CompoundStmt *SYCLKernelBody = - CreateSYCLKernelBody(*this, KernelCallerFunc, SYCLKernel); - SYCLKernel->setBody(SYCLKernelBody); - AddSyclKernel(SYCLKernel); + CompoundStmt *OpenCLKernelBody = + CreateOpenCLKernelBody(*this, KernelCallerFunc, OpenCLKernel); + OpenCLKernel->setBody(OpenCLKernelBody); + AddSyclKernel(OpenCLKernel); } void Sema::MarkDevice(void) { diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp index ed5008c20759..032ba09b448c 100644 --- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -5523,7 +5523,7 @@ void Sema::PerformPendingInstantiations(bool LocalOnly) { // so we are checking for SYCL kernel attribute after instantination. if (getLangOpts().SYCLIsDevice && CurFD->hasAttr()) { - ConstructSYCLKernel(CurFD); + ConstructOpenCLKernel(CurFD); } CurFD->setInstantiationIsPending(false); } @@ -5537,7 +5537,7 @@ void Sema::PerformPendingInstantiations(bool LocalOnly) { // so we are checking for SYCL kernel attribute after instantination. if (getLangOpts().SYCLIsDevice && Function->hasAttr()) { - ConstructSYCLKernel(Function); + ConstructOpenCLKernel(Function); } Function->setInstantiationIsPending(false); } diff --git a/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp b/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp index a0d68ca8334c..e5d58b91f0c1 100644 --- a/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp +++ b/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp @@ -26,9 +26,6 @@ int main() { // Check lambda object alloca // CHECK: [[ANON:%[0-9]+]] = alloca %"class.{{.*}}.anon" // Check allocas for ranges -// CHECK: [[ACC_RANGE_COPY:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range" -// CHECK: [[MEM_RANGE_COPY:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::range" -// CHECK: [[OFFSET_COPY:%[a-zA-Z0-9_.]+]] = alloca %"struct.{{.*}}.cl::sycl::id" // // Check store of kernel pointer argument to alloca // CHECK: store i32 addrspace(1)* [[MEM_ARG]], i32 addrspace(1)** [[MEM_ARG]].addr, align 8 @@ -39,19 +36,8 @@ int main() { // Check load from kernel pointer argument alloca // CHECK: [[MEM_LOAD:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG]].addr -// Check that ranges/offsets arguments are copied to allocas -// CHECK: [[BITCAST1:%[0-9]+]] = bitcast %"struct.{{.*}}.cl::sycl::range"* [[ACC_RANGE_COPY]] to i8* -// CHECK: [[BITCAST2:%[0-9]+]] = bitcast %"struct.{{.*}}.cl::sycl::range"* [[ACC_RANGE]] to i8* -// CHECK: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 [[BITCAST1]], i8* align 4 [[BITCAST2]], i64 4, i1 false) -// CHECK: [[BITCAST3:%[0-9]+]] = bitcast %"struct.{{.*}}.cl::sycl::range"* [[MEM_RANGE_COPY]] to i8* -// CHECK: [[BITCAST4:%[0-9]+]] = bitcast %"struct.{{.*}}.cl::sycl::range"* [[MEM_RANGE]] to i8* -// CHECK: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 [[BITCAST3]], i8* align 4 [[BITCAST4]], i64 4, i1 false) -// CHECK: [[BITCAST5:%[0-9]+]] = bitcast %"struct.{{.*}}.cl::sycl::id"* [[OFFSET_COPY]] to i8* -// CHECK: [[BITCAST6:%[0-9]+]] = bitcast %"struct.{{.*}}.cl::sycl::id"* [[OFFSET]] to i8* -// CHECK: call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 4 [[BITCAST5]], i8* align 4 [[BITCAST6]], i64 4, i1 false) - // Check accessor __init method call -// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor"* [[ACCESSOR]], i32 addrspace(1)* [[MEM_LOAD]], %"struct.{{.*}}.cl::sycl::range"* byval align 4 [[ACC_RANGE_COPY]], %"struct.{{.*}}.cl::sycl::range"* byval align 4 [[MEM_RANGE_COPY]], %"struct.{{.*}}.cl::sycl::id"* byval align 4 [[OFFSET_COPY]]) +// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.{{.*}}.cl::sycl::accessor"* [[ACCESSOR]], i32 addrspace(1)* [[MEM_LOAD]], %"struct.{{.*}}.cl::sycl::range"* byval align 4 [[ACC_RANGE]], %"struct.{{.*}}.cl::sycl::range"* byval align 4 [[MEM_RANGE]], %"struct.{{.*}}.cl::sycl::id"* byval align 4 [[OFFSET]]) // Check lambda "()" operator call // CHECK: call spir_func void @{{.*}}(%"class.{{.*}}.anon"* [[ANON]]) diff --git a/clang/test/CodeGenSYCL/sampler.cpp b/clang/test/CodeGenSYCL/sampler.cpp index 258f3f8e12b9..8d75cc4508c7 100644 --- a/clang/test/CodeGenSYCL/sampler.cpp +++ b/clang/test/CodeGenSYCL/sampler.cpp @@ -3,11 +3,11 @@ // CHECK-NEXT: entry: // CHECK-NEXT: [[SAMPLER_ARG]].addr = alloca %opencl.sampler_t addrspace(2)*, align 8 // CHECK-NEXT: [[ANON:%[0-9]+]] = alloca %"class.{{.*}}.anon", align 8 -// CHECK-NEXT: store %opencl.sampler_t addrspace(2)* [[SAMPLER_ARG]], %opencl.sampler_t addrspace(2)** [[SAMPLER_ARG]].addr, align 8, !tbaa !9 +// CHECK-NEXT: store %opencl.sampler_t addrspace(2)* [[SAMPLER_ARG]], %opencl.sampler_t addrspace(2)** [[SAMPLER_ARG]].addr, align 8 // CHECK-NEXT: [[BITCAST:%[0-9]+]] = bitcast %"class.{{.*}}.anon"* [[ANON]] to i8* // CHECK-NEXT: call void @llvm.lifetime.start.p0i8(i64 8, i8* [[BITCAST]]) #4 // CHECK-NEXT: [[GEP:%[0-9]+]] = getelementptr inbounds %"class.{{.*}}.anon", %"class.{{.*}}.anon"* [[ANON]], i32 0, i32 0 -// CHECK-NEXT: [[LOAD_SAMPLER_ARG:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)** [[SAMPLER_ARG]].addr, align 8, !tbaa !9 +// CHECK-NEXT: [[LOAD_SAMPLER_ARG:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)** [[SAMPLER_ARG]].addr, align 8 // CHECK-NEXT: call spir_func void @{{[a-zA-Z0-9_]+}}(%"class.{{.*}}.cl::sycl::sampler"* [[GEP]], %opencl.sampler_t addrspace(2)* [[LOAD_SAMPLER_ARG]]) // #include "sycl.hpp" diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp index 361c08630ecc..6313ab843ea9 100644 --- a/clang/test/SemaSYCL/Inputs/sycl.hpp +++ b/clang/test/SemaSYCL/Inputs/sycl.hpp @@ -56,6 +56,24 @@ struct _ImplT { id Offset; }; +template +struct DeviceValueType; + +template +struct DeviceValueType { + using type = __global dataT; +}; + +template +struct DeviceValueType { + using type = __constant dataT; +}; + +template +struct DeviceValueType { + using type = __local dataT; +}; + template @@ -67,7 +85,8 @@ class accessor { _ImplT impl; private: - void __init(__global dataT *Ptr, range AccessRange, + using PtrType = typename DeviceValueType::type *; + void __init(PtrType Ptr, range AccessRange, range MemRange, id Offset) {} }; diff --git a/clang/test/SemaSYCL/accessors-targets.cpp b/clang/test/SemaSYCL/accessors-targets.cpp index 958d216ea676..de024e289196 100644 --- a/clang/test/SemaSYCL/accessors-targets.cpp +++ b/clang/test/SemaSYCL/accessors-targets.cpp @@ -36,6 +36,6 @@ int main() { constant_acc.use(); }); } -// CHECK: {{.*}}use_local 'void (__local int *, range<1>, range<1>, id<1>)' -// CHECK: {{.*}}use_global 'void (__global int *, range<1>, range<1>, id<1>)' -// CHECK: {{.*}}use_constant 'void (__constant int *, range<1>, range<1>, id<1>)' +// CHECK: {{.*}}use_local 'void (__local int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' +// CHECK: {{.*}}use_global 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' +// CHECK: {{.*}}use_constant 'void (__constant int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' diff --git a/clang/test/SemaSYCL/basic-kernel-wrapper.cpp b/clang/test/SemaSYCL/basic-kernel-wrapper.cpp index d46d77414b9b..b8aa2e59ebed 100644 --- a/clang/test/SemaSYCL/basic-kernel-wrapper.cpp +++ b/clang/test/SemaSYCL/basic-kernel-wrapper.cpp @@ -23,14 +23,14 @@ int main() { // Check declaration of the kernel -// CHECK: FunctionDecl {{.*}}kernel_wrapper 'void (__global int *, range<1>, range<1>, id<1>)' +// CHECK: FunctionDecl {{.*}}kernel_wrapper 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' // Check parameters of the kernel // CHECK: ParmVarDecl {{.*}} used [[_arg_Mem:[0-9a-zA-Z_]+]] '__global int *' -// CHECK: ParmVarDecl {{.*}} used [[_arg_AccessRange:[0-9a-zA-Z_]+]] 'range<1>':'cl::sycl::range<1>' -// CHECK: ParmVarDecl {{.*}} used [[_arg_MemRange:[0-9a-zA-Z_]+]] 'range<1>':'cl::sycl::range<1>' -// CHECK: ParmVarDecl {{.*}} used [[_arg_Offset:[0-9a-zA-Z_]+]] 'id<1>':'cl::sycl::id<1>' +// CHECK: ParmVarDecl {{.*}} used [[_arg_AccessRange:[0-9a-zA-Z_]+]] 'cl::sycl::range<1>' +// CHECK: ParmVarDecl {{.*}} used [[_arg_MemRange:[0-9a-zA-Z_]+]] 'cl::sycl::range<1>' +// CHECK: ParmVarDecl {{.*}} used [[_arg_Offset:[0-9a-zA-Z_]+]] 'cl::sycl::id<1>' // Check body of the kernel @@ -42,18 +42,17 @@ int main() { // Check accessor initialization // CHECK: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void (__global int *, range<1>, range<1>, id<1>)' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'void ({{.*}}PtrType, range<1>, range<1>, id<1>)' lvalue .__init // CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::accessor':'cl::sycl::accessor' lvalue . // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}basic-kernel-wrapper.cpp{{.*}})' lvalue Var -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} // CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '[[_arg_Mem]]' '__global int *' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'range<1>':'cl::sycl::range<1>' -// CHECK-NEXT: DeclRefExpr {{.*}} 'range<1>':'cl::sycl::range<1>' lvalue ParmVar {{.*}} '[[_arg_AccessRange]]' 'range<1>':'cl::sycl::range<1>' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'range<1>':'cl::sycl::range<1>' -// CHECK-NEXT: DeclRefExpr {{.*}} 'range<1>':'cl::sycl::range<1>' lvalue ParmVar {{.*}} '[[_arg_MemRange]]' 'range<1>':'cl::sycl::range<1>' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'id<1>':'cl::sycl::id<1>' -// CHECK-NEXT: DeclRefExpr {{.*}} 'id<1>':'cl::sycl::id<1>' lvalue ParmVar {{.*}} '[[_arg_Offset]]' 'id<1>':'cl::sycl::id<1>' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'range<1>':'cl::sycl::range<1>' +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::range<1>' lvalue ParmVar {{.*}} '[[_arg_AccessRange]]' 'cl::sycl::range<1>' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'range<1>':'cl::sycl::range<1>' +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::range<1>' lvalue ParmVar {{.*}} '[[_arg_MemRange]]' 'cl::sycl::range<1>' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'id<1>':'cl::sycl::id<1>' +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::id<1>' lvalue ParmVar {{.*}} '[[_arg_Offset]]' 'cl::sycl::id<1>' // Check that body of the kernel caller function is included into kernel diff --git a/clang/test/SemaSYCL/fake-accessors.cpp b/clang/test/SemaSYCL/fake-accessors.cpp index 112f986d2008..acce120e49f6 100644 --- a/clang/test/SemaSYCL/fake-accessors.cpp +++ b/clang/test/SemaSYCL/fake-accessors.cpp @@ -51,6 +51,6 @@ int main() { }); return 0; } -// CHECK: fake_accessors 'void (__global int *, range<1>, range<1>, id<1>, foo::cl::sycl::accessor, accessor) -// CHECK: accessor_typedef 'void (__global int *, range<1>, range<1>, id<1>, foo::cl::sycl::accessor, accessor) -// CHECK: accessor_alias 'void (__global int *, range<1>, range<1>, id<1>, foo::cl::sycl::accessor, accessor) +// CHECK: fake_accessors 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor) +// CHECK: accessor_typedef 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor) +// CHECK: accessor_alias 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor) diff --git a/clang/test/SemaSYCL/sampler.cpp b/clang/test/SemaSYCL/sampler.cpp index cf3a1d4e4e01..0d7d97ad07c7 100644 --- a/clang/test/SemaSYCL/sampler.cpp +++ b/clang/test/SemaSYCL/sampler.cpp @@ -16,10 +16,10 @@ int main() { } // Check declaration of the test kernel -// CHECK: FunctionDecl {{.*}}use_kernel_for_test 'void (__ocl_sampler_t)' +// CHECK: FunctionDecl {{.*}}use_kernel_for_test 'void (sampler_t)' // // Check parameters of the test kernel -// CHECK: ParmVarDecl {{.*}} used [[_arg_sampler:[0-9a-zA-Z_]+]] '__ocl_sampler_t' +// CHECK: ParmVarDecl {{.*}} used [[_arg_sampler:[0-9a-zA-Z_]+]] 'sampler_t' // // Check that sampler field of the test kernel object is initialized using __init method // CHECK: CXXMemberCallExpr {{.*}} 'void' @@ -29,5 +29,4 @@ int main() { // // Check the parameters of __init method // CHECK-NEXT: ImplicitCastExpr {{.*}} '__ocl_sampler_t':'sampler_t' -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__ocl_sampler_t':'sampler_t' lvalue -// CHECK-NEXT: DeclRefExpr {{.*}} '__ocl_sampler_t':'sampler_t' lvalue ParmVar {{.*}} '[[_arg_sampler]]' '__ocl_sampler_t':'sampler_t' +// CHECK-NEXT: DeclRefExpr {{.*}} 'sampler_t' lvalue ParmVar {{.*}} '[[_arg_sampler]]' 'sampler_t' diff --git a/clang/test/SemaSYCL/wrapped-accessor.cpp b/clang/test/SemaSYCL/wrapped-accessor.cpp index 4e810e6b8120..63b100ece9a5 100644 --- a/clang/test/SemaSYCL/wrapped-accessor.cpp +++ b/clang/test/SemaSYCL/wrapped-accessor.cpp @@ -23,14 +23,14 @@ int main() { } // Check declaration of the kernel -// CHECK: wrapped_access 'void (AccWrapper >, __global int *, range<1>, range<1>, id<1>)' +// CHECK: wrapped_access 'void (AccWrapper >, __global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' // Check parameters of the kernel // CHECK: ParmVarDecl {{.*}} used _arg_ 'AccWrapper >':'AccWrapper >' // CHECK: ParmVarDecl {{.*}} used _arg_accessor '__global int *' -// CHECK: ParmVarDecl {{.*}} used [[_arg_AccessRange:[0-9a-zA-Z_]+]] 'range<1>':'cl::sycl::range<1>' -// CHECK: ParmVarDecl {{.*}} used [[_arg_MemRange:[0-9a-zA-Z_]+]] 'range<1>':'cl::sycl::range<1>' -// CHECK: ParmVarDecl {{.*}} used [[_arg_Offset:[0-9a-zA-Z_]+]] 'id<1>':'cl::sycl::id<1>' +// CHECK: ParmVarDecl {{.*}} used [[_arg_AccessRange:[0-9a-zA-Z_]+]] 'cl::sycl::range<1>' +// CHECK: ParmVarDecl {{.*}} used [[_arg_MemRange:[0-9a-zA-Z_]+]] 'cl::sycl::range<1>' +// CHECK: ParmVarDecl {{.*}} used [[_arg_Offset:[0-9a-zA-Z_]+]] 'cl::sycl::id<1>' // Check that wrapper object itself is initialized with corresponding kernel argument using operator= // CHECK: BinaryOperator {{.*}} 'AccWrapper >':'AccWrapper >' lvalue '=' @@ -45,18 +45,17 @@ int main() { // Check that accessor field of the wrapper object is initialized using __init method // CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void' -// CHECK-NEXT: MemberExpr {{.*}} 'void (__global int *, range<1>, range<1>, id<1>)' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'void ({{.*}}PtrType, range<1>, range<1>, id<1>)' lvalue .__init // CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::accessor':'cl::sycl::accessor' lvalue .accessor {{.*}} // CHECK-NEXT: MemberExpr {{.*}} 'AccWrapper >':'AccWrapper >' lvalue . // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}wrapped-accessor.cpp{{.*}})' lvalue Var {{.*}} '(lambda at {{.*}}wrapped-accessor.cpp{{.*}})' // Parameters of the _init method -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' -// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' lvalue +// CHECK-NEXT: ImplicitCastExpr {{.*}} // CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '_arg_accessor' '__global int *' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'range<1>':'cl::sycl::range<1>' -// CHECK-NEXT: DeclRefExpr {{.*}} 'range<1>':'cl::sycl::range<1>' lvalue ParmVar {{.*}} '[[_arg_AccessRange]]' 'range<1>':'cl::sycl::range<1>' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'range<1>':'cl::sycl::range<1>' -// CHECK-NEXT: DeclRefExpr {{.*}} 'range<1>':'cl::sycl::range<1>' lvalue ParmVar {{.*}} '[[_arg_MemRange]]' 'range<1>':'cl::sycl::range<1>' -// CHECK-NEXT: ImplicitCastExpr {{.*}} 'id<1>':'cl::sycl::id<1>' -// CHECK-NEXT: DeclRefExpr {{.*}} 'id<1>':'cl::sycl::id<1>' lvalue ParmVar {{.*}} '[[_arg_Offset]]' 'id<1>':'cl::sycl::id<1>' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'range<1>':'cl::sycl::range<1>' +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::range<1>' lvalue ParmVar {{.*}} '[[_arg_AccessRange]]' 'cl::sycl::range<1>' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'range<1>':'cl::sycl::range<1>' +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::range<1>' lvalue ParmVar {{.*}} '[[_arg_MemRange]]' 'cl::sycl::range<1>' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'id<1>':'cl::sycl::id<1>' +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::id<1>' lvalue ParmVar {{.*}} '[[_arg_Offset]]' 'cl::sycl::id<1>'