From 405174c0a454cd087f065d8f2643bd1d24a0d763 Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Tue, 11 Jun 2019 07:59:04 +0300 Subject: [PATCH 1/7] [SYCL] Unify kernel wrapper generation for sampler and accessor Signed-off-by: Mariya Podchishchaeva --- clang/lib/Sema/SemaSYCL.cpp | 244 +++++------------- .../test/CodeGenSYCL/basic-kernel-wrapper.cpp | 16 +- clang/test/CodeGenSYCL/sampler.cpp | 4 +- clang/test/SemaSYCL/Inputs/sycl.hpp | 21 +- clang/test/SemaSYCL/accessors-targets.cpp | 6 +- clang/test/SemaSYCL/basic-kernel-wrapper.cpp | 25 +- clang/test/SemaSYCL/fake-accessors.cpp | 6 +- clang/test/SemaSYCL/sampler.cpp | 7 +- clang/test/SemaSYCL/wrapped-accessor.cpp | 25 +- 9 files changed, 127 insertions(+), 227 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 28d8bf5119ca8..0c9f193624279 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -400,23 +400,26 @@ class KernelBodyTransform : public TreeTransform { Sema &SemaRef; }; -static FunctionDecl *CreateSYCLKernelFunction(ASTContext &Context, +static FunctionDecl *CreateSYCLKernelDeclaration(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( Context, DC, SourceLocation(), SourceLocation(), DN, FuncTy, Context.getTrivialTypeSourceInfo(RetTy), SC_None); + llvm::SmallVector Params; int i = 0; for (const auto &PD : ParamDescs) { @@ -434,7 +437,7 @@ static FunctionDecl *CreateSYCLKernelFunction(ASTContext &Context, SYCLKernel->addAttr(AsmLabelAttr::CreateImplicit(Context, Name)); SYCLKernel->addAttr(ArtificialAttr::CreateImplicit(Context)); - // To see kernel in AST-dump. + // Add kernel to translation unit to see it in AST-dump DC->addDecl(SYCLKernel); return SYCLKernel; } @@ -450,7 +453,7 @@ static CXXMethodDecl *getInitMethod(const CXXRecordDecl *CRD) { } static CompoundStmt * -CreateSYCLKernelBody(Sema &S, FunctionDecl *KernelCallerFunc, DeclContext *DC) { +CreateSYCLKernelBody(Sema &S, FunctionDecl *KernelCallerFunc, DeclContext *KernelDecl) { llvm::SmallVector BodyStmts; CXXRecordDecl *LC = getKernelObjectType(KernelCallerFunc); assert(LC && "Kernel object must be available"); @@ -458,67 +461,58 @@ CreateSYCLKernelBody(Sema &S, FunctionDecl *KernelCallerFunc, DeclContext *DC) { // 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(), + 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 = dyn_cast(KernelDecl); + assert(KernelFuncDecl && "No kernel function declaration?"); + 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, + // 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) { - // 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(); + // 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 + // [kenrel_obj or wrapper object].special_obj auto AccessorME = 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(), @@ -535,16 +529,13 @@ 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])); + llvm::SmallVector ParamStmts; + for (size_t I = 0; I < NumParams; ++I) { + ParamStmts.push_back(getExprForKernelParameter( + S, (*(ParamItr++))->getOriginalType(), ParamDREs[I])); + } // [kenrel_obj or wrapper object].accessor.__init(_ValueType*, // range, range, id) CXXMemberCallExpr *Call = CXXMemberCallExpr::Create( @@ -566,8 +557,8 @@ 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, + KernelFuncParam++; + getExprForSpecialSYCLObj(FldType, WrapperFld, WrapperFldCRD, Base); } else { // Field is a structure or class so change the wrapper object @@ -589,68 +580,22 @@ CreateSYCLKernelBody(Sema &S, FunctionDecl *KernelCallerFunc, DeclContext *DC) { 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 // 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); @@ -670,7 +615,7 @@ CreateSYCLKernelBody(Sema &S, FunctionDecl *KernelCallerFunc, DeclContext *DC) { } else { llvm_unreachable("unsupported field type"); } - TargetFuncParam++; + KernelFuncParam++; } } // In function from headers lambda is function parameter, we need @@ -719,61 +664,28 @@ static void buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj, ParamDescs.push_back(makeParamDesc(Fld, ArgType)); }; - auto createAccessorParamDesc = [&](const FieldDecl *Fld, + // Create 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) { - // the parameter is a SYCL accessor object 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()); + } }; + // 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 +695,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 @@ -795,22 +707,8 @@ static void buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj, 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()) { if (!ArgTy->isStandardLayoutType()) { const DeclaratorDecl *V = @@ -968,7 +866,7 @@ void Sema::ConstructSYCLKernel(FunctionDecl *KernelCallerFunc) { std::string Name = constructKernelName(KernelNameType, getASTContext()); populateIntHeader(getSyclIntegrationHeader(), Name, KernelNameType, LE); FunctionDecl *SYCLKernel = - CreateSYCLKernelFunction(getASTContext(), Name, ParamDescs); + CreateSYCLKernelDeclaration(getASTContext(), Name, ParamDescs); // Let's copy source location of a functor/lambda to emit nicer diagnostics SYCLKernel->setLocation(LE->getLocation()); diff --git a/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp b/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp index a0d68ca8334c9..e5d58b91f0c1a 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 258f3f8e12b96..8d75cc4508c7f 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 361c08630ecc5..6313ab843ea9e 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 958d216ea6762..de024e2891966 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 d46d77414b9ba..b8aa2e59ebedf 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 112f986d20083..acce120e49f68 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 cf3a1d4e4e015..0d7d97ad07c77 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 4e810e6b8120b..63b100ece9a50 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>' From 969129f9cd5b82d18c34753e645c17457510453b Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Tue, 11 Jun 2019 08:26:45 +0300 Subject: [PATCH 2/7] [SYCL] Refactor ThreeTransform in SemaSYCL.cpp We use this ThreeTransform to change only one object to another. No need to use DenseMap here, std::pair is enough. Signed-off-by: Mariya Podchishchaeva --- clang/lib/Sema/SemaSYCL.cpp | 42 ++++++++++++++++++------------------- 1 file changed, 20 insertions(+), 22 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 0c9f193624279..7cc40d8f79b0b 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,27 +374,25 @@ 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; }; @@ -618,21 +614,23 @@ CreateSYCLKernelBody(Sema &S, FunctionDecl *KernelCallerFunc, DeclContext *Kerne 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 kernel caller function lambda/functior is function parameter, we need + // to replace all refs to this lambda/functor with our kernel object 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(), From ce9fe95784bdf6cd681b76286b9dc59885faae80 Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Tue, 11 Jun 2019 12:17:50 +0300 Subject: [PATCH 3/7] [SYCL] Comment kernel wrapper generation code Signed-off-by: Mariya Podchishchaeva --- clang/lib/Sema/SemaSYCL.cpp | 145 ++++++++++++++++++++++++++---------- 1 file changed, 105 insertions(+), 40 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 7cc40d8f79b0b..d5d5e7ca3483e 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -396,9 +396,9 @@ class KernelBodyTransform : public TreeTransform { Sema &SemaRef; }; -static FunctionDecl *CreateSYCLKernelDeclaration(ASTContext &Context, - StringRef Name, - ArrayRef ParamDescs) { +static FunctionDecl * +CreateSYCLKernelDeclaration(ASTContext &Context, StringRef Name, + ArrayRef ParamDescs) { DeclContext *DC = Context.getTranslationUnitDecl(); QualType RetTy = Context.VoidTy; @@ -448,17 +448,22 @@ static CXXMethodDecl *getInitMethod(const CXXRecordDecl *CRD) { return InitMethod; } -static CompoundStmt * -CreateSYCLKernelBody(Sema &S, FunctionDecl *KernelCallerFunc, DeclContext *KernelDecl) { +// Creates body for new SYCL kernel. This body contains initialization of kernel +// object fields with kernel parameters and a little bit transformed body of the +// kernel caller function. +static CompoundStmt *CreateSYCLKernelBody(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, KernelDecl, 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); @@ -466,6 +471,7 @@ CreateSYCLKernelBody(Sema &S, FunctionDecl *KernelCallerFunc, DeclContext *Kerne DeclRefExpr::Create(S.Context, NestedNameSpecifierLoc(), SourceLocation(), KernelObjClone, false, DeclarationNameInfo(), QualType(LC->getTypeForDecl(), 0), VK_LValue); + auto KernelFuncDecl = dyn_cast(KernelDecl); assert(KernelFuncDecl && "No kernel function declaration?"); auto KernelFuncParam = @@ -484,11 +490,13 @@ CreateSYCLKernelBody(Sema &S, FunctionDecl *KernelCallerFunc, DeclContext *Kerne // 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) { + 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"); + assert(InitMethod && + "The accessor/sampler must have the __init method"); unsigned NumParams = InitMethod->getNumParams(); llvm::SmallVector ParamDREs(NumParams); auto KFP = KernelFuncParam; @@ -503,8 +511,8 @@ CreateSYCLKernelBody(Sema &S, FunctionDecl *KernelCallerFunc, DeclContext *Kerne DeclAccessPair FieldDAP = DeclAccessPair::make(Field, AS_none); // [kenrel_obj or wrapper object].special_obj auto AccessorME = MemberExpr::Create( - S.Context, Base, false, SourceLocation(), - NestedNameSpecifierLoc(), SourceLocation(), Field, FieldDAP, + S.Context, Base, false, SourceLocation(), NestedNameSpecifierLoc(), + SourceLocation(), Field, FieldDAP, DeclarationNameInfo(Field->getDeclName(), SourceLocation()), nullptr, Field->getType(), VK_LValue, OK_Ordinary); @@ -555,7 +563,7 @@ CreateSYCLKernelBody(Sema &S, FunctionDecl *KernelCallerFunc, DeclContext *Kerne // or parameter of the previous processed accessor object. KernelFuncParam++; getExprForSpecialSYCLObj(FldType, WrapperFld, WrapperFldCRD, - Base); + Base); } else { // Field is a structure or class so change the wrapper object // and recursively search for accessor field. @@ -574,13 +582,26 @@ CreateSYCLKernelBody(Sema &S, FunctionDecl *KernelCallerFunc, DeclContext *Kerne } }; + // 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) || Util::isSyclSamplerType(FieldType)) { getExprForSpecialSYCLObj(FieldType, Field, CRD, KernelObjCloneRef); + } else if (Util::isSyclStreamType(FieldType)) { + // TODO add support for streams + llvm_unreachable("Streams not supported yet"); } 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. @@ -609,14 +630,14 @@ CreateSYCLKernelBody(Sema &S, FunctionDecl *KernelCallerFunc, DeclContext *Kerne if (CRD) getExprForWrappedAccessorInit(CRD, Lhs); } else { - llvm_unreachable("unsupported field type"); + llvm_unreachable("Unsupported field type"); } KernelFuncParam++; } } - // In kernel caller function lambda/functior is function parameter, we need - // to replace all refs to this lambda/functor with our kernel object clone + // 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()); @@ -654,22 +675,26 @@ 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)); }; - // Create a parameter descriptor for SYCL special object - SYCL accessor or + // 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 QualType &ArgTy) { const auto *RecordDecl = ArgTy->getAsCXXRecordDecl(); assert(RecordDecl && "Special SYCL object must be of a record type"); @@ -682,6 +707,8 @@ static void buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj, } }; + // 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 @@ -703,11 +730,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) || 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()) @@ -715,15 +755,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"); } } } @@ -743,18 +782,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 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); }; std::function @@ -847,22 +886,48 @@ static std::string constructKernelName(QualType KernelNameType, return Out.str(); } +// Generates the "kernel wrapper" using KernelCallerFunc (kernel caller +// function) defined is SYCL headers. +// A "kernel wrapper" function 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 "kernel wrapper" 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(); +// } +// +// In the code below we call "kernel wrapper" SYCLKernel. +// void Sema::ConstructSYCLKernel(FunctionDecl *KernelCallerFunc) { - // TODO: Case when kernel is functor 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 = CreateSYCLKernelDeclaration(getASTContext(), Name, ParamDescs); From a5944b4415747d162ea73e87b4e4739cd0014c71 Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Mon, 17 Jun 2019 11:21:55 +0300 Subject: [PATCH 4/7] [SYCL] Minor changes Signed-off-by: Mariya Podchishchaeva --- clang/lib/Sema/SemaSYCL.cpp | 10 +++------- 1 file changed, 3 insertions(+), 7 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index d5d5e7ca3483e..10c99e1fac091 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -472,8 +472,7 @@ static CompoundStmt *CreateSYCLKernelBody(Sema &S, KernelObjClone, false, DeclarationNameInfo(), QualType(LC->getTypeForDecl(), 0), VK_LValue); - auto KernelFuncDecl = dyn_cast(KernelDecl); - assert(KernelFuncDecl && "No kernel function declaration?"); + auto KernelFuncDecl = cast(KernelDecl); auto KernelFuncParam = KernelFuncDecl->param_begin(); // Iterator to ParamVarDecl (VarDecl) if (KernelFuncParam) { @@ -510,7 +509,7 @@ static CompoundStmt *CreateSYCLKernelBody(Sema &S, DeclAccessPair FieldDAP = DeclAccessPair::make(Field, AS_none); // [kenrel_obj or wrapper object].special_obj - auto AccessorME = MemberExpr::Create( + auto SpecialObjME = MemberExpr::Create( S.Context, Base, false, SourceLocation(), NestedNameSpecifierLoc(), SourceLocation(), Field, FieldDAP, DeclarationNameInfo(Field->getDeclName(), SourceLocation()), @@ -519,7 +518,7 @@ static CompoundStmt *CreateSYCLKernelBody(Sema &S, // [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); @@ -597,9 +596,6 @@ static CompoundStmt *CreateSYCLKernelBody(Sema &S, if (Util::isSyclAccessorType(FieldType) || Util::isSyclSamplerType(FieldType)) { getExprForSpecialSYCLObj(FieldType, Field, CRD, KernelObjCloneRef); - } else if (Util::isSyclStreamType(FieldType)) { - // TODO add support for streams - llvm_unreachable("Streams not supported yet"); } else if (CRD || FieldType->isScalarType()) { // If field has built-in or a structure/class type just initialize // this field with corresponding kernel argument using '=' binary From 2d7c6b78ba458952e46d97379108345b6c74b8e7 Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Tue, 18 Jun 2019 11:47:47 +0300 Subject: [PATCH 5/7] [SYCL] Rename SYCLKernel and "kernel wrapper" with OpenCLKernel Signed-off-by: Mariya Podchishchaeva --- clang/include/clang/Sema/Sema.h | 2 +- clang/lib/Sema/SemaSYCL.cpp | 55 +++++++++---------- .../lib/Sema/SemaTemplateInstantiateDecl.cpp | 4 +- 3 files changed, 30 insertions(+), 31 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 6200e5560e49a..8d021f7e00d49 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 10c99e1fac091..3ecd2bd800c0e 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -397,7 +397,7 @@ class KernelBodyTransform : public TreeTransform { }; static FunctionDecl * -CreateSYCLKernelDeclaration(ASTContext &Context, StringRef Name, +CreateOpenCLKernelDeclaration(ASTContext &Context, StringRef Name, ArrayRef ParamDescs) { DeclContext *DC = Context.getTranslationUnitDecl(); @@ -412,30 +412,30 @@ CreateSYCLKernelDeclaration(ASTContext &Context, StringRef Name, 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)); // Add kernel to translation unit to see it in AST-dump - DC->addDecl(SYCLKernel); - return SYCLKernel; + DC->addDecl(OpenCLKernel); + return OpenCLKernel; } /// Return __init method static CXXMethodDecl *getInitMethod(const CXXRecordDecl *CRD) { @@ -448,12 +448,12 @@ static CXXMethodDecl *getInitMethod(const CXXRecordDecl *CRD) { return InitMethod; } -// Creates body for new SYCL kernel. This body contains initialization of kernel -// object fields with kernel parameters and a little bit transformed body of the -// kernel caller function. -static CompoundStmt *CreateSYCLKernelBody(Sema &S, - FunctionDecl *KernelCallerFunc, - DeclContext *KernelDecl) { +// 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"); @@ -882,13 +882,13 @@ static std::string constructKernelName(QualType KernelNameType, return Out.str(); } -// Generates the "kernel wrapper" using KernelCallerFunc (kernel caller +// Generates the OpenCL kernel using KernelCallerFunc (kernel caller // function) defined is SYCL headers. -// A "kernel wrapper" function contains the body of the kernel caller function, +// 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 "kernel wrapper" from KernelCallerFunc we put +// 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. @@ -903,9 +903,8 @@ static std::string constructKernelName(QualType KernelNameType, // KernelFuncObj(); // } // -// In the code below we call "kernel wrapper" SYCLKernel. // -void Sema::ConstructSYCLKernel(FunctionDecl *KernelCallerFunc) { +void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc) { CXXRecordDecl *LE = getKernelObjectType(KernelCallerFunc); assert(LE && "invalid kernel caller"); @@ -924,16 +923,16 @@ void Sema::ConstructSYCLKernel(FunctionDecl *KernelCallerFunc) { // TODO Maybe don't emit integration header inside the Sema? populateIntHeader(getSyclIntegrationHeader(), Name, KernelNameType, LE); - FunctionDecl *SYCLKernel = - CreateSYCLKernelDeclaration(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 ed5008c207592..032ba09b448c8 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); } From a6328b2edb7adcbfe12d46af8ecc6f5bfe48892d Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Wed, 19 Jun 2019 15:23:13 +0300 Subject: [PATCH 6/7] [SYCL] Apply clang-format Signed-off-by: Mariya Podchishchaeva --- clang/lib/Sema/SemaSYCL.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 3ecd2bd800c0e..6afecdd66b69e 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -398,7 +398,7 @@ class KernelBodyTransform : public TreeTransform { static FunctionDecl * CreateOpenCLKernelDeclaration(ASTContext &Context, StringRef Name, - ArrayRef ParamDescs) { + ArrayRef ParamDescs) { DeclContext *DC = Context.getTranslationUnitDecl(); QualType RetTy = Context.VoidTy; From 21f0aaad5c963b43e6a73bbe42769aecf845533c Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Wed, 19 Jun 2019 20:49:44 +0300 Subject: [PATCH 7/7] [SYCL] Fix typos in comments Signed-off-by: Mariya Podchishchaeva --- clang/lib/Sema/SemaSYCL.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 6afecdd66b69e..516d37dc2b05d 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -508,7 +508,7 @@ static CompoundStmt *CreateOpenCLKernelBody(Sema &S, std::advance(KernelFuncParam, NumParams - 1); DeclAccessPair FieldDAP = DeclAccessPair::make(Field, AS_none); - // [kenrel_obj or wrapper object].special_obj + // [kernel_obj or wrapper object].special_obj auto SpecialObjME = MemberExpr::Create( S.Context, Base, false, SourceLocation(), NestedNameSpecifierLoc(), SourceLocation(), Field, FieldDAP, @@ -539,7 +539,7 @@ static CompoundStmt *CreateOpenCLKernelBody(Sema &S, ParamStmts.push_back(getExprForKernelParameter( S, (*(ParamItr++))->getOriginalType(), ParamDREs[I])); } - // [kenrel_obj or wrapper object].accessor.__init(_ValueType*, + // [kernel_obj or wrapper object].accessor.__init(_ValueType*, // range, range, id) CXXMemberCallExpr *Call = CXXMemberCallExpr::Create( S.Context, ME, ParamStmts, ResultTy, VK, SourceLocation()); @@ -780,7 +780,7 @@ static void populateIntHeader(SYCLIntegrationHeader &H, const StringRef Name, 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: + // 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();