From 1c93c84add1d00e7cf060553e105b8b19a1ffe61 Mon Sep 17 00:00:00 2001 From: Vladimir Lazarev Date: Sun, 30 Dec 2018 23:19:32 +0300 Subject: [PATCH] [SYCL] Support structures as kernel parameters (compiler part). Signed-off-by: Vladimir Lazarev --- .../clang/Basic/DiagnosticSemaKinds.td | 2 + clang/include/clang/Sema/Sema.h | 6 +- clang/lib/Sema/SemaSYCL.cpp | 74 +++++++++++-------- clang/test/CodeGenSYCL/integration_header.cpp | 14 ++-- 4 files changed, 53 insertions(+), 43 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 49d19382ad99d..65ff990d1bb28 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -9491,4 +9491,6 @@ def err_sycl_kernel_name_class_not_top_level : Error< def err_sycl_virtual_types : Error< "No class with a vtable can be used in a SYCL kernel or any code included in the kernel">; def note_sycl_used_here : Note<"used here">; +def err_sycl_non_std_layout_type : Error< + "kernel parameter has non-standard layout class/struct type">; } // end of sema component. diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 1fac6da7d698d..0fdebe207b555 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -299,11 +299,9 @@ class SYCLIntegrationHeader { enum kernel_param_kind_t { kind_first, kind_accessor = kind_first, - kind_scalar, - kind_struct, + kind_std_layout, kind_sampler, - kind_struct_padding, // can be added by the compiler to enforce alignment - kind_last = kind_struct_padding + kind_last = kind_sampler }; public: diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index f0458b5c020fb..4c85e000ec882 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -37,6 +37,24 @@ enum target { image_array }; +/// Various utilities. +class Util { +public: + // TODO SYCL use AST infrastructure instead of string matching + + /// Checks whether given clang type is a sycl accessor class. + static bool isSyclAccessorType(QualType Ty) { + std::string Name = Ty.getCanonicalType().getAsString(); + return Name.find("class cl::sycl::accessor") != std::string::npos; + } + + /// Checks whether given clang type is a sycl stream class. + static bool isSyclStreamType(QualType Ty) { + std::string Name = Ty.getCanonicalType().getAsString(); + return Name == "stream"; + } +}; + static CXXRecordDecl *getKernelCallerLambdaArg(FunctionDecl *FD) { auto FirstArg = (*FD->param_begin()); if (FirstArg) @@ -271,7 +289,7 @@ CreateSYCLKernelBody(Sema &S, FunctionDecl *KernelCallerFunc, DeclContext *DC) { QualType FieldType = Field->getType(); CXXRecordDecl *CRD = FieldType->getAsCXXRecordDecl(); - if (CRD) { + if (CRD && Util::isSyclAccessorType(FieldType)) { DeclAccessPair FieldDAP = DeclAccessPair::make(Field, AS_none); // lambda.accessor auto AccessorME = MemberExpr::Create( @@ -373,9 +391,11 @@ CreateSYCLKernelBody(Sema &S, FunctionDecl *KernelCallerFunc, DeclContext *DC) { "unsupported accessor and without initialized range"); } } - } else if (FieldType->isBuiltinType()) { - // If field have built-in type just initialize this field - // with corresponding kernel argument using '=' binary operator. + } else if (CRD || FieldType->isBuiltinType()) { + // 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. DeclAccessPair FieldDAP = DeclAccessPair::make(Field, AS_none); auto Lhs = MemberExpr::Create( S.Context, LambdaDRE, false, SourceLocation(), @@ -416,31 +436,13 @@ CreateSYCLKernelBody(Sema &S, FunctionDecl *KernelCallerFunc, DeclContext *DC) { SourceLocation()); } -/// Various utilities. -class Util { -public: - // TODO SYCL use AST infrastructure instead of string matching - - /// Checks whether given clang type is a sycl accessor class. - static bool isSyclAccessorType(QualType Ty) { - std::string Name = Ty.getCanonicalType().getAsString(); - return Name.find("class cl::sycl::accessor") != std::string::npos; - } - - /// Checks whether given clang type is a sycl stream class. - static bool isSyclStreamType(QualType Ty) { - std::string Name = Ty.getCanonicalType().getAsString(); - return Name == "stream"; - } -}; - /// Identifies context of kernel lambda capture visitor function /// invocation. enum VisitorContext { pre_visit, pre_visit_class_field, visit_accessor, - visit_scalar, + visit_std_layout, visit_stream, post_visit, }; @@ -508,9 +510,16 @@ static void visitKernelLambdaCaptures(const CXXRecordDecl *Lambda, // stream parameter context auto F = std::get(Vis); F(Cnt, V, *Fld); + } else if (ArgTy->isStructureOrClassType()) { + if (!ArgTy->isStandardLayoutType()) + Lambda->getASTContext().getDiagnostics().Report(V->getLocation(), + diag::err_sycl_non_std_layout_type); + // structure or class typed parameter - the same handling as a scalar + auto F = std::get(Vis); + F(Cnt, V, *Fld); } else if (ArgTy->isScalarType()) { // scalar typed parameter context - auto F = std::get(Vis); + auto F = std::get(Vis); F(Cnt, V, *Fld); } else { llvm_unreachable("unsupported kernel parameter type"); @@ -523,7 +532,7 @@ static void visitKernelLambdaCaptures(const CXXRecordDecl *Lambda, // pre-visit context the same like for accessor auto F1Range = std::get(Vis); F1Range(Cnt, V, *Fld, AccessorRangeField); - auto FRange = std::get(Vis); + auto FRange = std::get(Vis); FRange(Cnt, V, AccessorRangeField); // post-visit context auto F2Range = std::get(Vis); @@ -568,7 +577,7 @@ static void BuildArgTys(ASTContext &Context, CXXRecordDecl *Lambda, ActualArgType = Context.getQualifiedType(PointerType.getUnqualifiedType(), Quals); }, - // visit_scalar + // visit_std_layout [&](int CaptureN, VarDecl *CapturedVar, FieldDecl *CapturedVal) { ActualArgType = CapturedVal->getType(); }, @@ -643,9 +652,12 @@ static void populateIntHeader(SYCLIntegrationHeader &H, const StringRef Name, Knd = SYCLIntegrationHeader::kind_accessor; Info = static_cast(AccTrg); }, - // visit_scalar + // visit_std_layout [&](int CaptureN, VarDecl *CapturedVar, FieldDecl *CapturedVal) { - Knd = SYCLIntegrationHeader::kind_scalar; + // TODO this code (when used to handle a structure-typed scalar) relies + // on the host and device structure layouts and sizes to be the same. + // Need SYCL spec clarification on passing structures as parameters. + Knd = SYCLIntegrationHeader::kind_std_layout; Info = static_cast( Ctx.getTypeSizeInChars(CapturedVal->getType()).getQuantity()); }, @@ -740,10 +752,8 @@ static const char *paramKind2Str(KernelParamKind K) { return "kind_" #x switch (K) { CASE(accessor); - CASE(scalar); - CASE(struct); + CASE(std_layout); CASE(sampler); - CASE(struct_padding); default: return ""; } @@ -766,7 +776,7 @@ void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D) { cast(D)->getTemplatedDecl() : dyn_cast(D); if (TD && TD->isCompleteDefinition()) { - // defied class constituting the kernel name is not globally + // defined class constituting the kernel name is not globally // accessible - contradicts the spec Diag.Report(D->getSourceRange().getBegin(), diag::err_sycl_kernel_name_class_not_top_level); diff --git a/clang/test/CodeGenSYCL/integration_header.cpp b/clang/test/CodeGenSYCL/integration_header.cpp index 357fa08151567..a2540339d819c 100644 --- a/clang/test/CodeGenSYCL/integration_header.cpp +++ b/clang/test/CodeGenSYCL/integration_header.cpp @@ -21,21 +21,21 @@ // CHECK: static constexpr // CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = { // CHECK-NEXT: //--- first_kernel -// CHECK-NEXT: { kernel_param_kind_t::kind_scalar, 4, 0 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 }, // CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 2014, 4 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_scalar, 1, 4 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 1, 4 }, // CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 2016, 5 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_scalar, 1, 5 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 1, 5 }, // CHECK-EMPTY: // CHECK-NEXT: //--- ::second_namespace::second_kernel -// CHECK-NEXT: { kernel_param_kind_t::kind_scalar, 4, 0 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 }, // CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 2016, 4 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_scalar, 1, 4 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 1, 4 }, // CHECK-EMPTY: // CHECK-NEXT: //--- ::third_kernel<1, int, ::point > -// CHECK-NEXT: { kernel_param_kind_t::kind_scalar, 4, 0 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 }, // CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 2016, 4 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_scalar, 1, 4 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 1, 4 }, // CHECK-EMPTY: // CHECK-NEXT: }; //