From 8a95712bcc9a2a6d34ce66b6b5948c24204fa846 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Tue, 7 May 2019 16:47:08 -0700 Subject: [PATCH] [SYCL] Get sampler and accessor types without using named fields of kernel object Using the types of the parameter variables of __init method instead Signed-off-by: Sindhu Chittireddy --- clang/lib/Sema/SemaSYCL.cpp | 86 ++++++++++++++---------- clang/test/SemaSYCL/sampler.cpp | 4 +- clang/test/SemaSYCL/wrapped-accessor.cpp | 12 ++-- 3 files changed, 59 insertions(+), 43 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 454fbdb2da26..d63c8f7ab72b 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -412,6 +412,16 @@ static FunctionDecl *CreateSYCLKernelFunction(ASTContext &Context, DC->addDecl(SYCLKernel); return SYCLKernel; } +/// Return __init method +static CXXMethodDecl *getInitMethod(const CXXRecordDecl *CRD) { + CXXMethodDecl *InitMethod; + auto It = std::find_if(CRD->methods().begin(), CRD->methods().end(), + [](const CXXMethodDecl *Method) { + return Method->getNameAsString() == "__init"; + }); + InitMethod = (It != CRD->methods().end()) ? *It : nullptr; + return InitMethod; +} static CompoundStmt * CreateSYCLKernelBody(Sema &S, FunctionDecl *KernelCallerFunc, DeclContext *DC) { @@ -479,13 +489,7 @@ CreateSYCLKernelBody(Sema &S, FunctionDecl *KernelCallerFunc, DeclContext *DC) { DeclarationNameInfo(Field->getDeclName(), SourceLocation()), nullptr, Field->getType(), VK_LValue, OK_Ordinary); - CXXMethodDecl *InitMethod = nullptr; - for (auto Method : CRD->methods()) { - if (Method->getNameInfo().getName().getAsString() == "__init") { - InitMethod = Method; - break; - } - } + CXXMethodDecl *InitMethod = getInitMethod(CRD); assert(InitMethod && "The accessor must have the __init method"); // [kenrel_obj or wrapper object].accessor.__init @@ -581,13 +585,7 @@ CreateSYCLKernelBody(Sema &S, FunctionDecl *KernelCallerFunc, DeclContext *DC) { DeclarationNameInfo(Field->getDeclName(), SourceLocation()), nullptr, Field->getType(), VK_LValue, OK_Ordinary); - CXXMethodDecl *InitMethod = nullptr; - for (auto Method : CRD->methods()) { - if (Method->getNameInfo().getName().getAsString() == "__init") { - InitMethod = Method; - break; - } - } + CXXMethodDecl *InitMethod = getInitMethod(CRD); assert(InitMethod && "The sampler must have the __init method"); // kernel_obj.sampler.__init @@ -755,20 +753,27 @@ static void buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj, CreateAndAddPrmDsc(Fld, PointerType); - FieldDecl *AccessRangeFld = - getFieldDeclByName(RecordDecl, {"impl", "AccessRange"}); - assert(AccessRangeFld && - "The accessor.impl must contain the AccessRange field"); - CreateAndAddPrmDsc(AccessRangeFld, AccessRangeFld->getType()); + CXXMethodDecl *InitMethod = getInitMethod(RecordDecl); + assert(InitMethod && "accessor must have __init method"); - FieldDecl *MemRangeFld = - getFieldDeclByName(RecordDecl, {"impl", "MemRange"}); - assert(MemRangeFld && "The accessor.impl must contain the MemRange field"); - CreateAndAddPrmDsc(MemRangeFld, MemRangeFld->getType()); + // 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); - FieldDecl *OffsetFld = getFieldDeclByName(RecordDecl, {"impl", "Offset"}); - assert(OffsetFld && "The accessor.impl must contain the Offset field"); - CreateAndAddPrmDsc(OffsetFld, OffsetFld->getType()); + 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()); }; std::function @@ -799,10 +804,16 @@ static void buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj, const auto *RecordDecl = ArgTy->getAsCXXRecordDecl(); assert(RecordDecl && "sampler must be of a record type"); - FieldDecl *ImplFld = - getFieldDeclByName(RecordDecl, {"impl", "m_Sampler"}); - assert(ImplFld && "The sampler must contain impl field"); - CreateAndAddPrmDsc(ImplFld, ImplFld->getType()); + CXXMethodDecl *InitMethod = getInitMethod(RecordDecl); + assert(InitMethod && "sampler must have __init method"); + + // sampler __init method has only one parameter + // void __init(__spirv::OpTypeSampler *Sampler) + auto *FuncDecl = cast(InitMethod); + ParmVarDecl *SamplerArg = FuncDecl->getParamDecl(0); + assert(SamplerArg && "sampler __init method must have sampler parameter"); + + CreateAndAddPrmDsc(Fld, SamplerArg->getType()); } else if (Util::isSyclStreamType(ArgTy)) { // the parameter is a SYCL stream object llvm_unreachable("streams not supported yet"); @@ -893,13 +904,18 @@ static void populateIntHeader(SYCLIntegrationHeader &H, const StringRef Name, populateHeaderForAccessor(ArgTy, Offset); } else if (Util::isSyclSamplerType(ArgTy)) { // The parameter is a SYCL sampler object - // It has only one descriptor, "m_Sampler" const auto *SamplerTy = ArgTy->getAsCXXRecordDecl(); assert(SamplerTy && "sampler must be of a record type"); - FieldDecl *ImplFld = - getFieldDeclByName(SamplerTy, {"impl", "m_Sampler"}, &Offset); - uint64_t Sz = - Ctx.getTypeSizeInChars(ImplFld->getType()).getQuantity(); + + CXXMethodDecl *InitMethod = getInitMethod(SamplerTy); + assert(InitMethod && "sampler must have __init method"); + + // sampler __init method has only one argument + // void __init(__spirv::OpTypeSampler *Sampler) + auto *FuncDecl = cast(InitMethod); + ParmVarDecl *SamplerArg = FuncDecl->getParamDecl(0); + assert(SamplerArg && "sampler __init method must have sampler parameter"); + uint64_t Sz = Ctx.getTypeSizeInChars(SamplerArg->getType()).getQuantity(); H.addParamDesc(SYCLIntegrationHeader::kind_sampler, static_cast(Sz), static_cast(Offset)); } else if (Util::isSyclStreamType(ArgTy)) { diff --git a/clang/test/SemaSYCL/sampler.cpp b/clang/test/SemaSYCL/sampler.cpp index 8ac4102fe92b..9bf67860903c 100644 --- a/clang/test/SemaSYCL/sampler.cpp +++ b/clang/test/SemaSYCL/sampler.cpp @@ -19,7 +19,7 @@ int main() { // CHECK: FunctionDecl {{.*}}use_kernel_for_test 'void (__spirv::OpTypeSampler *)' // // Check parameters of the test kernel -// CHECK: ParmVarDecl {{.*}} used _arg_m_Sampler '__spirv::OpTypeSampler *' +// CHECK: ParmVarDecl {{.*}} used [[_arg_sampler:[0-9a-zA-Z_]+]] '__spirv::OpTypeSampler *' // // Check that sampler field of the test kernel object is initialized using __init method // CHECK: CXXMemberCallExpr {{.*}} 'void' @@ -30,4 +30,4 @@ int main() { // Check the parameters of __init method // CHECK-NEXT: ImplicitCastExpr {{.*}} '__spirv::OpTypeSampler *' // CHECK-NEXT: ImplicitCastExpr {{.*}} '__spirv::OpTypeSampler *' lvalue -// CHECK-NEXT: DeclRefExpr {{.*}} '__spirv::OpTypeSampler *' lvalue ParmVar {{.*}} '_arg_m_Sampler' '__spirv::OpTypeSampler +// CHECK-NEXT: DeclRefExpr {{.*}} '__spirv::OpTypeSampler *' lvalue ParmVar {{.*}} '[[_arg_sampler]]' '__spirv::OpTypeSampler diff --git a/clang/test/SemaSYCL/wrapped-accessor.cpp b/clang/test/SemaSYCL/wrapped-accessor.cpp index 15104bef2b0b..141dbe677ced 100644 --- a/clang/test/SemaSYCL/wrapped-accessor.cpp +++ b/clang/test/SemaSYCL/wrapped-accessor.cpp @@ -25,9 +25,9 @@ int main() { // Check parameters of the kernel // CHECK: ParmVarDecl {{.*}} used _arg_ 'AccWrapper >':'AccWrapper >' // CHECK: ParmVarDecl {{.*}} used _arg_accessor '__global int *' -// CHECK: ParmVarDecl {{.*}} used _arg_AccessRange 'range<1>':'cl::sycl::range<1>' -// CHECK: ParmVarDecl {{.*}} used _arg_MemRange 'range<1>':'cl::sycl::range<1>' -// CHECK: ParmVarDecl {{.*}} used _arg_Offset 'id<1>':'cl::sycl::id<1>' +// 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 that wrapper object itself is initialized with corresponding kernel argument using operator= // CHECK: BinaryOperator {{.*}} 'AccWrapper >':'AccWrapper >' lvalue '=' @@ -52,8 +52,8 @@ int main() { // CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' lvalue // 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: 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: 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: DeclRefExpr {{.*}} 'id<1>':'cl::sycl::id<1>' lvalue ParmVar {{.*}} '[[_arg_Offset]]' 'id<1>':'cl::sycl::id<1>'