Skip to content

Commit

Permalink
[SYCL] Get sampler and accessor types without using named fields of
Browse files Browse the repository at this point in the history 
kernel object

Using the types of the parameter variables of __init method instead

Signed-off-by: Sindhu Chittireddy <sindhu.chittireddy@intel.com>
  • Loading branch information
@schittir @bader
schittir authored and bader committed May 23, 2019
1 parent a5d763b commit 8a95712
Show file tree
Hide file tree
Showing 3 changed files with 59 additions and 43 deletions.
86 changes: 51 additions & 35 deletions clang/lib/Sema/SemaSYCL.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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) {
Expand Down Expand Up @@ -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
Expand Down Expand Up @@ -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
Expand Down Expand Up @@ -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<dimensions> AccessRange,
// range<dimensions> MemRange, id<dimensions> Offset)
auto *FuncDecl = cast<FunctionDecl>(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<void(const FieldDecl *, const QualType &ArgTy)>
Expand Down Expand Up @@ -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<FunctionDecl>(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");
Expand Down Expand Up @@ -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<FunctionDecl>(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<unsigned>(Sz), static_cast<unsigned>(Offset));
} else if (Util::isSyclStreamType(ArgTy)) {
Expand Down
4 changes: 2 additions & 2 deletions clang/test/SemaSYCL/sampler.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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'
Expand All @@ -30,4 +30,4 @@ int main() {
// Check the parameters of __init method
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__spirv::OpTypeSampler *' <LValueToRValue>
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__spirv::OpTypeSampler *' lvalue <NoOp>
// CHECK-NEXT: DeclRefExpr {{.*}} '__spirv::OpTypeSampler *' lvalue ParmVar {{.*}} '_arg_m_Sampler' '__spirv::OpTypeSampler
// CHECK-NEXT: DeclRefExpr {{.*}} '__spirv::OpTypeSampler *' lvalue ParmVar {{.*}} '[[_arg_sampler]]' '__spirv::OpTypeSampler
12 changes: 6 additions & 6 deletions clang/test/SemaSYCL/wrapped-accessor.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -25,9 +25,9 @@ int main() {
// Check parameters of the kernel
// CHECK: ParmVarDecl {{.*}} used _arg_ 'AccWrapper<cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t> >':'AccWrapper<cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t> >'
// 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<cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t> >':'AccWrapper<cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t> >' lvalue '='
Expand All @@ -52,8 +52,8 @@ int main() {
// CHECK-NEXT: ImplicitCastExpr {{.*}} '__global int *' lvalue <NoOp>
// CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '_arg_accessor' '__global int *'
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'range<1>':'cl::sycl::range<1>' <NoOp>
// 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>' <NoOp>
// 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>' <NoOp>
// 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>'

0 comments on commit 8a95712

Please sign in to comment.