From 1884bd4bd8efa04d6afa5425a2ae286366d3cd62 Mon Sep 17 00:00:00 2001 From: Mikhail Lychkov Date: Tue, 25 Aug 2020 22:49:05 +0300 Subject: [PATCH 1/9] [SYCL][FPGA] Align clang with new spec of accessor_property_list Signed-off-by: Mikhail Lychkov --- .../clang/Basic/DiagnosticSemaKinds.td | 7 +- clang/lib/Sema/SemaSYCL.cpp | 86 +++++++++++++------ clang/test/CodeGenSYCL/Inputs/sycl.hpp | 40 +++++---- .../test/CodeGenSYCL/accessor_inheritance.cpp | 6 +- clang/test/CodeGenSYCL/buffer_location.cpp | 12 +-- clang/test/CodeGenSYCL/integration_header.cpp | 18 ++-- .../CodeGenSYCL/kernel-param-acc-array-ih.cpp | 2 +- .../kernel-param-member-acc-array-ih.cpp | 2 +- .../test/CodeGenSYCL/struct_kernel_param.cpp | 2 +- clang/test/SemaSYCL/Inputs/sycl.hpp | 13 ++- clang/test/SemaSYCL/accessor_inheritance.cpp | 4 +- clang/test/SemaSYCL/basic-kernel-wrapper.cpp | 2 +- clang/test/SemaSYCL/buffer_location.cpp | 32 +++---- clang/test/SemaSYCL/wrapped-accessor.cpp | 8 +- 14 files changed, 141 insertions(+), 93 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index c532b2272ad6e..5aff495d36f84 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -11033,9 +11033,10 @@ def err_sycl_compiletime_property_duplication : Error< def err_sycl_invalid_property_list_param_number : Error< "%0 must have exactly one template parameter">; def err_sycl_invalid_accessor_property_template_param : Error< - "Fifth template parameter of the accessor must be of a property_list type">; -def err_sycl_invalid_property_list_template_param : Error< - "%select{property_list|property_list pack argument|buffer_location}0 " + "Sixth template parameter of the accessor must be of accessor_property_list " + "or property_list type">; +def err_sycl_invalid_accessor_property_list_template_param : Error< + "%select{accessor_property_list|accessor_property_list pack argument|buffer_location}0 " "template parameter must be a " "%select{parameter pack|type|non-negative integer}1">; def warn_sycl_pass_by_value_deprecated diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index f5812b986e547..17bba2c16e9c7 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -85,6 +85,10 @@ class Util { /// property_list class. static bool isPropertyListType(const QualType &Ty); + /// Checks whether given clang type is a full specialization of the SYCL + /// accessor_property_list class. + static bool isAccessorPropertyListType(const QualType &Ty); + /// Checks whether given clang type is a full specialization of the SYCL /// buffer_location class. static bool isSyclBufferLocationType(const QualType &Ty); @@ -1194,29 +1198,33 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler { return; } QualType PropListTy = PropList.getAsType(); - if (!Util::isPropertyListType(PropListTy)) { + if (Util::isPropertyListType(PropListTy)) + return; + if (!Util::isAccessorPropertyListType(PropListTy)) { SemaRef.Diag(Loc, diag::err_sycl_invalid_accessor_property_template_param); return; } - const auto *PropListDecl = + const auto *AccPropListDecl = cast(PropListTy->getAsRecordDecl()); - if (PropListDecl->getTemplateArgs().size() != 1) { + if (AccPropListDecl->getTemplateArgs().size() != 1) { SemaRef.Diag(Loc, diag::err_sycl_invalid_property_list_param_number) - << "property_list"; + << "accessor_property_list"; return; } - const auto TemplArg = PropListDecl->getTemplateArgs()[0]; + const auto TemplArg = AccPropListDecl->getTemplateArgs()[0]; if (TemplArg.getKind() != TemplateArgument::ArgKind::Pack) { - SemaRef.Diag(Loc, diag::err_sycl_invalid_property_list_template_param) - << /*property_list*/ 0 << /*parameter pack*/ 0; + SemaRef.Diag(Loc, + diag::err_sycl_invalid_accessor_property_list_template_param) + << /*accessor_property_list*/ 0 << /*parameter pack*/ 0; return; } for (TemplateArgument::pack_iterator Prop = TemplArg.pack_begin(); Prop != TemplArg.pack_end(); ++Prop) { if (Prop->getKind() != TemplateArgument::ArgKind::Type) { - SemaRef.Diag(Loc, diag::err_sycl_invalid_property_list_template_param) - << /*property_list pack argument*/ 1 << /*type*/ 1; + SemaRef.Diag( + Loc, diag::err_sycl_invalid_accessor_property_list_template_param) + << /*accessor_property_list pack argument*/ 1 << /*type*/ 1; return; } QualType PropTy = Prop->getAsType(); @@ -1235,13 +1243,15 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler { } const auto BufferLoc = PropDecl->getTemplateArgs()[0]; if (BufferLoc.getKind() != TemplateArgument::ArgKind::Integral) { - SemaRef.Diag(Loc, diag::err_sycl_invalid_property_list_template_param) + SemaRef.Diag(Loc, + diag::err_sycl_invalid_accessor_property_list_template_param) << /*buffer_location*/ 2 << /*non-negative integer*/ 2; return; } int LocationID = static_cast(BufferLoc.getAsIntegral().getExtValue()); if (LocationID < 0) { - SemaRef.Diag(Loc, diag::err_sycl_invalid_property_list_template_param) + SemaRef.Diag(Loc, + diag::err_sycl_invalid_accessor_property_list_template_param) << /*buffer_location*/ 2 << /*non-negative integer*/ 2; return; } @@ -1251,17 +1261,29 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler { assert(Util::isSyclAccessorType(Ty) && "Should only be called on SYCL accessor types."); - const RecordDecl *RecD = Ty->getAsRecordDecl(); - if (const ClassTemplateSpecializationDecl *CTSD = - dyn_cast(RecD)) { + RecordDecl *RecD = Ty->getAsRecordDecl(); + if (auto *CTSD = dyn_cast(RecD)) { const TemplateArgumentList &TAL = CTSD->getTemplateArgs(); + TemplateArgument TA = TAL.get(0); const QualType TemplateArgTy = TA.getAsType(); - - if (TAL.size() > 5) - checkPropertyListType(TAL.get(5), Loc.getBegin()); llvm::DenseSet Visited; checkSYCLType(SemaRef, TemplateArgTy, Loc, Visited); + + if (TAL.size() < 6) { + // Not enough arguments for this parameter pack. + SemaRef.Diag(Loc.getBegin(), + diag::err_template_arg_list_different_arity) + << /*not enough args*/ 0 + << (int)SemaRef.getTemplateNameKindForDiagnostics( + TemplateName(CTSD->getSpecializedTemplate())) + << CTSD; + SemaRef.Diag(CTSD->getLocation(), diag::note_template_decl_here) + << CTSD->getSourceRange(); + return; + } + + checkPropertyListType(TAL.get(5), Loc.getBegin()); } } @@ -1402,19 +1424,21 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { } // Handle accessor properties. If any properties were found in - // the property_list - add the appropriate attributes to ParmVarDecl. + // the accessor_property_list - add the appropriate attributes to ParmVarDecl. void handleAccessorPropertyList(ParmVarDecl *Param, const CXXRecordDecl *RecordDecl, SourceLocation Loc) { const auto *AccTy = cast(RecordDecl); - // TODO: when SYCL headers' part is ready - replace this 'if' with an error if (AccTy->getTemplateArgs().size() < 6) return; const auto PropList = cast(AccTy->getTemplateArgs()[5]); QualType PropListTy = PropList.getAsType(); - const auto *PropListDecl = + // property_list contains runtime properties, it shouldn't be handled here. + if (Util::isPropertyListType(PropListTy)) + return; + const auto *AccPropListDecl = cast(PropListTy->getAsRecordDecl()); - const auto TemplArg = PropListDecl->getTemplateArgs()[0]; + const auto TemplArg = AccPropListDecl->getTemplateArgs()[0]; // Move through TemplateArgs list of a property list and search for // properties. If found - apply the appropriate attribute to ParmVarDecl. for (TemplateArgument::pack_iterator Prop = TemplArg.pack_begin(); @@ -3444,17 +3468,16 @@ bool Util::isSyclSpecConstantType(const QualType &Ty) { } bool Util::isPropertyListType(const QualType &Ty) { - return isSyclType(Ty, "property_list", true /*Tmpl*/); + return isSyclType(Ty, "property_list"); } bool Util::isSyclBufferLocationType(const QualType &Ty) { const StringRef &Name = "buffer_location"; - std::array Scopes = { + std::array Scopes = { Util::DeclContextDesc{clang::Decl::Kind::Namespace, "cl"}, Util::DeclContextDesc{clang::Decl::Kind::Namespace, "sycl"}, - // TODO: this doesn't belong to property namespace, instead it shall be - // in its own namespace. Change it, when the actual implementation in SYCL - // headers is ready + Util::DeclContextDesc{clang::Decl::Kind::Namespace, "ext"}, + Util::DeclContextDesc{clang::Decl::Kind::Namespace, "INTEL"}, Util::DeclContextDesc{clang::Decl::Kind::Namespace, "property"}, Util::DeclContextDesc{Decl::Kind::ClassTemplateSpecialization, Name}}; return matchQualifiedTypeName(Ty, Scopes); @@ -3470,6 +3493,17 @@ bool Util::isSyclType(const QualType &Ty, StringRef Name, bool Tmpl) { return matchQualifiedTypeName(Ty, Scopes); } +bool Util::isAccessorPropertyListType(const QualType &Ty) { + const StringRef &Name = "accessor_property_list"; + std::array Scopes = { + Util::DeclContextDesc{clang::Decl::Kind::Namespace, "cl"}, + Util::DeclContextDesc{clang::Decl::Kind::Namespace, "sycl"}, + Util::DeclContextDesc{clang::Decl::Kind::Namespace, "ext"}, + Util::DeclContextDesc{clang::Decl::Kind::Namespace, "ONEAPI"}, + Util::DeclContextDesc{Decl::Kind::ClassTemplateSpecialization, Name}}; + return matchQualifiedTypeName(Ty, Scopes); +} + bool Util::matchQualifiedTypeName(const QualType &Ty, ArrayRef Scopes) { // The idea: check the declaration context chain starting from the type diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index 0d77fd81dc7da..da1a78f4ac145 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -71,19 +71,11 @@ enum prop_type { base_prop }; -// Compile time known accessor property -// TODO: this doesn't belong to property namespace, instead it shall be in its -// own namespace. Change it, when the actual implementation in SYCL headers is -// ready -template -class buffer_location {}; - struct property_base { virtual prop_type type() const = 0; }; } // namespace property -template class property_list { public: template @@ -102,6 +94,21 @@ class property_list { bool operator!=(const property_list &rhs) const { return false; } }; +namespace ext { +namespace INTEL { +namespace property { +// Compile time known accessor property +template +class buffer_location {}; +} // namespace property +} // namespace INTEL + +namespace ONEAPI { +template +class accessor_property_list {}; +} // namespace ONEAPI +} // namespace ext + template struct id { template @@ -136,7 +143,7 @@ struct _ImplT { template > + typename propertyListT = property_list> class accessor { public: @@ -150,8 +157,6 @@ class accessor { private: void __init(__attribute__((opencl_global)) dataT *Ptr, range AccessRange, range MemRange, id Offset) {} - - propertyListT prop_list; }; template @@ -339,8 +344,7 @@ const stream& operator<<(const stream &S, T&&) { } template + typename AllocatorT = int /*fake type as AllocatorT is not used*/> class buffer { public: using value_type = T; @@ -352,13 +356,13 @@ class buffer { buffer(ParamTypes... args) {} // fake constructor buffer(const range &bufferRange, - const property_list &propList = {}) {} + const property_list &propList = {}) {} buffer(T *hostData, const range &bufferRange, - const property_list &propList = {}) {} + const property_list &propList = {}) {} buffer(const T *hostData, const range &bufferRange, - const property_list &propList = {}) {} + const property_list &propList = {}) {} buffer(const buffer &rhs) = default; @@ -426,12 +430,12 @@ enum class image_channel_type : unsigned int { fp32 }; -template +template class image { public: image(image_channel_order Order, image_channel_type Type, const range &Range, - const property_list &PropList = {}) {} + const property_list &PropList = {}) {} /* -- common interface members -- */ diff --git a/clang/test/CodeGenSYCL/accessor_inheritance.cpp b/clang/test/CodeGenSYCL/accessor_inheritance.cpp index b3857806a2bcd..b94c9fa7a9478 100644 --- a/clang/test/CodeGenSYCL/accessor_inheritance.cpp +++ b/clang/test/CodeGenSYCL/accessor_inheritance.cpp @@ -67,13 +67,13 @@ int main() { // CHECK: [[ACC_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base* [[BITCAST]], i32 0, i32 2 // CHECK: [[ACC1_AS_CAST:%[a-zA-Z0-9_]+]] = addrspacecast %"class{{.*}}cl::sycl::accessor"* [[ACC_FIELD]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)* // Default constructor call -// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_13property_listIJEEEEC1Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[ACC1_AS_CAST]]) +// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_13property_listEEC1Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[ACC1_AS_CAST]]) // CHECK: [[BITCAST1:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured* [[GEP]] to i8* -// CHECK: [[GEP1:%[a-zA-Z0-9_]+]] = getelementptr inbounds i8, i8* [[BITCAST1]], i64 24 +// CHECK: [[GEP1:%[a-zA-Z0-9_]+]] = getelementptr inbounds i8, i8* [[BITCAST1]], i64 20 // CHECK: [[BITCAST2:%[a-zA-Z0-9_]+]] = bitcast i8* [[GEP1]] to %"class{{.*}}cl::sycl::accessor"* // CHECK: [[ACC2_AS_CAST:%[a-zA-Z0-9_]+]] = addrspacecast %"class{{.*}}cl::sycl::accessor"* [[BITCAST2]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)* // Default constructor call -// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_13property_listIJEEEEC2Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[ACC2_AS_CAST]]) +// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_13property_listEEC2Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[ACC2_AS_CAST]]) // CHECK C field initialization // CHECK: [[FIELD_C:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Captured, %struct{{.*}}Captured* [[GEP]], i32 0, i32 2 diff --git a/clang/test/CodeGenSYCL/buffer_location.cpp b/clang/test/CodeGenSYCL/buffer_location.cpp index ae1b9088cc271..b18e8ecef8234 100644 --- a/clang/test/CodeGenSYCL/buffer_location.cpp +++ b/clang/test/CodeGenSYCL/buffer_location.cpp @@ -10,8 +10,8 @@ struct Base { cl::sycl::accessor>> + cl::sycl::ext::ONEAPI::accessor_property_list< + cl::sycl::ext::INTEL::property::buffer_location<2>>> AccField; }; @@ -19,8 +19,8 @@ struct Captured : Base, cl::sycl::accessor>> { + cl::sycl::ext::ONEAPI::accessor_property_list< + cl::sycl::ext::INTEL::property::buffer_location<2>>> { int C; }; @@ -29,8 +29,8 @@ int main() { cl::sycl::accessor>> + cl::sycl::ext::ONEAPI::accessor_property_list< + cl::sycl::ext::INTEL::property::buffer_location<3>>> accessorA; cl::sycl::kernel_single_task( [=]() { diff --git a/clang/test/CodeGenSYCL/integration_header.cpp b/clang/test/CodeGenSYCL/integration_header.cpp index 65c1e22be1889..a0859f4925c43 100644 --- a/clang/test/CodeGenSYCL/integration_header.cpp +++ b/clang/test/CodeGenSYCL/integration_header.cpp @@ -31,18 +31,18 @@ // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 1, 4 }, // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 8 }, // CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 12 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 6112, 28 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_sampler, 8, 48 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 6112, 24 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_sampler, 8, 40 }, // CHECK-EMPTY: // CHECK-NEXT: //--- _ZTSN16second_namespace13second_kernelIcEE // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 }, // CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 6112, 4 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_sampler, 8, 24 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_sampler, 8, 16 }, // CHECK-EMPTY: // CHECK-NEXT: //--- _ZTS12third_kernelILi1Ei5pointIZ4mainE1XEE // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 }, // CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 6112, 4 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_sampler, 8, 24 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_sampler, 8, 16 }, // CHECK-EMPTY: // CHECK-NEXT: //--- _ZTS13fourth_kernelIJN15template_arg_ns14namespaced_argILi1EEEEE // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 }, @@ -52,11 +52,11 @@ // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 }, // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 4 }, // CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 8 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 24 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 28 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 44 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 48 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 64 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 20 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 24 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 36 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 40 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 52 }, // CHECK-EMPTY: // CHECK-NEXT: }; // diff --git a/clang/test/CodeGenSYCL/kernel-param-acc-array-ih.cpp b/clang/test/CodeGenSYCL/kernel-param-acc-array-ih.cpp index 7bf277437a032..7f458efb36a57 100644 --- a/clang/test/CodeGenSYCL/kernel-param-acc-array-ih.cpp +++ b/clang/test/CodeGenSYCL/kernel-param-acc-array-ih.cpp @@ -21,7 +21,7 @@ // CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = { // CHECK-NEXT: //--- _ZTSZ4mainE8kernel_A // CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 0 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 16 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 12 }, // CHECK-EMPTY: // CHECK-NEXT: }; diff --git a/clang/test/CodeGenSYCL/kernel-param-member-acc-array-ih.cpp b/clang/test/CodeGenSYCL/kernel-param-member-acc-array-ih.cpp index 242d3c2f50697..db2ffcfb13fd5 100644 --- a/clang/test/CodeGenSYCL/kernel-param-member-acc-array-ih.cpp +++ b/clang/test/CodeGenSYCL/kernel-param-member-acc-array-ih.cpp @@ -21,7 +21,7 @@ // CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = { // CHECK-NEXT: //--- _ZTSZ4mainE8kernel_C // CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 0 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 16 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 12 }, // CHECK-EMPTY: // CHECK-NEXT: }; diff --git a/clang/test/CodeGenSYCL/struct_kernel_param.cpp b/clang/test/CodeGenSYCL/struct_kernel_param.cpp index b1382ec6c2b92..a00f147b0dee0 100644 --- a/clang/test/CodeGenSYCL/struct_kernel_param.cpp +++ b/clang/test/CodeGenSYCL/struct_kernel_param.cpp @@ -4,12 +4,12 @@ // CHECK: const kernel_param_desc_t kernel_signatures[] = { // CHECK-NEXT: //--- _ZTSZZ5test0vENK3$_0clERN2cl4sycl7handlerEE8MyKernel // CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 0 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 12 }, // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 16 }, // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 20 }, // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 24 }, // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 28 }, // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 32 }, -// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 36 }, // CHECK-EMPTY: // CHECK-NEXT:}; diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp index 65a77b01f165a..2e9bab2a2492a 100644 --- a/clang/test/SemaSYCL/Inputs/sycl.hpp +++ b/clang/test/SemaSYCL/Inputs/sycl.hpp @@ -37,13 +37,21 @@ enum class address_space : int { }; } // namespace access +class property_list {}; + +namespace ext { +namespace INTEL { namespace property { template class buffer_location {}; } // namespace property +} // namespace INTEL +namespace ONEAPI { template -class property_list {}; +class accessor_property_list {}; +} // namespace ONEAPI +} // namespace ext namespace detail { namespace half_impl { @@ -95,7 +103,7 @@ struct DeviceValueType { template > + typename propertyListT = property_list> class accessor { public: @@ -107,7 +115,6 @@ class accessor { using PtrType = typename DeviceValueType::type *; void __init(PtrType Ptr, range AccessRange, range MemRange, id Offset) {} - propertyListT prop_list; }; template diff --git a/clang/test/SemaSYCL/accessor_inheritance.cpp b/clang/test/SemaSYCL/accessor_inheritance.cpp index db07ee25a5457..f286706aeafd9 100644 --- a/clang/test/SemaSYCL/accessor_inheritance.cpp +++ b/clang/test/SemaSYCL/accessor_inheritance.cpp @@ -42,8 +42,8 @@ int main() { // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_A' 'int' // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_B' 'int' -// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::accessor':'cl::sycl::accessor>' 'void () noexcept' -// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::accessor':'cl::sycl::accessor>' 'void () noexcept' +// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::accessor':'cl::sycl::accessor' 'void () noexcept' +// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::accessor':'cl::sycl::accessor' 'void () noexcept' // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_C' 'int' diff --git a/clang/test/SemaSYCL/basic-kernel-wrapper.cpp b/clang/test/SemaSYCL/basic-kernel-wrapper.cpp index 9d32dd1445453..d953cbbb4aa19 100644 --- a/clang/test/SemaSYCL/basic-kernel-wrapper.cpp +++ b/clang/test/SemaSYCL/basic-kernel-wrapper.cpp @@ -43,7 +43,7 @@ int main() { // CHECK: CXXMemberCallExpr {{.*}} 'void' // 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: MemberExpr {{.*}} 'cl::sycl::accessor':'cl::sycl::accessor' lvalue . // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}basic-kernel-wrapper.cpp{{.*}})' lvalue Var // CHECK-NEXT: ImplicitCastExpr {{.*}} diff --git a/clang/test/SemaSYCL/buffer_location.cpp b/clang/test/SemaSYCL/buffer_location.cpp index 65742af0cbe1c..5cb02923a9751 100644 --- a/clang/test/SemaSYCL/buffer_location.cpp +++ b/clang/test/SemaSYCL/buffer_location.cpp @@ -14,8 +14,8 @@ struct Base { cl::sycl::accessor>> + cl::sycl::ext::ONEAPI::accessor_property_list< + cl::sycl::ext::INTEL::property::buffer_location<1>>> AccField; }; @@ -23,8 +23,8 @@ struct Captured : Base, cl::sycl::accessor>> { + cl::sycl::ext::ONEAPI::accessor_property_list< + cl::sycl::ext::INTEL::property::buffer_location<1>>> { int C; }; @@ -35,28 +35,30 @@ int main() { cl::sycl::accessor>> + cl::sycl::ext::ONEAPI::accessor_property_list< + cl::sycl::ext::INTEL::property::buffer_location<2>>> // CHECK: SYCLIntelBufferLocationAttr {{.*}} Implicit 2 accessorA; cl::sycl::accessor>> + cl::sycl::ext::ONEAPI::accessor_property_list< + another_property, + cl::sycl::ext::INTEL::property::buffer_location<3>>> // CHECK: SYCLIntelBufferLocationAttr {{.*}} Implicit 3 accessorB; cl::sycl::accessor> + cl::sycl::ext::ONEAPI::accessor_property_list< + another_property>> accessorC; #else cl::sycl::accessor>> + cl::sycl::ext::ONEAPI::accessor_property_list< + cl::sycl::ext::INTEL::property::buffer_location<-2>>> accessorD; cl::sycl::accessor, - cl::sycl::property::buffer_location<2>>> + cl::sycl::ext::ONEAPI::accessor_property_list< + cl::sycl::ext::INTEL::property::buffer_location<1>, + cl::sycl::ext::INTEL::property::buffer_location<2>>> accessorF; #endif cl::sycl::kernel_single_task( @@ -82,7 +84,7 @@ int main() { #else //expected-error@+1{{buffer_location template parameter must be a non-negative integer}} accessorD.use(); - //expected-error@+1{{Fifth template parameter of the accessor must be of a property_list type}} + //expected-error@+1{{Sixth template parameter of the accessor must be of accessor_property_list or property_list type}} accessorE.use(); //expected-error@+1{{Can't apply buffer_location property twice to the same accessor}} accessorF.use(); diff --git a/clang/test/SemaSYCL/wrapped-accessor.cpp b/clang/test/SemaSYCL/wrapped-accessor.cpp index c6ebc3c48730e..9347c0999e34b 100644 --- a/clang/test/SemaSYCL/wrapped-accessor.cpp +++ b/clang/test/SemaSYCL/wrapped-accessor.cpp @@ -35,14 +35,14 @@ int main() { // argument // CHECK: VarDecl {{.*}}'(lambda at {{.*}}wrapped-accessor.cpp{{.*}})' // CHECK-NEXT: InitListExpr {{.*}}'(lambda at {{.*}}wrapped-accessor.cpp{{.*}})' -// CHECK-NEXT: InitListExpr {{.*}}'AccWrapper>>' -// CHECK-NEXT: CXXConstructExpr {{.*}}'cl::sycl::accessor>':'cl::sycl::accessor>' 'void () noexcept' +// CHECK-NEXT: InitListExpr {{.*}}'AccWrapper>' +// CHECK-NEXT: CXXConstructExpr {{.*}}'cl::sycl::accessor':'cl::sycl::accessor' 'void () noexcept' // Check that accessor field of the wrapper object is initialized using __init method // CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void' // 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: 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 From d375f2a9896b2be1e5ae93663a3c690ae576214f Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Wed, 9 Sep 2020 12:47:53 +0300 Subject: [PATCH 2/9] Align with the latest revision and temporarily allow old accessors --- clang/lib/Sema/SemaSYCL.cpp | 34 +++++++------------- clang/test/CodeGenSYCL/Inputs/sycl.hpp | 7 ++-- clang/test/CodeGenSYCL/buffer_location.cpp | 12 +++---- clang/test/SemaSYCL/Inputs/sycl.hpp | 9 +++--- clang/test/SemaSYCL/accessor_inheritance.cpp | 4 +-- clang/test/SemaSYCL/basic-kernel-wrapper.cpp | 2 +- clang/test/SemaSYCL/buffer_location.cpp | 28 ++++++++-------- clang/test/SemaSYCL/wrapped-accessor.cpp | 8 ++--- 8 files changed, 45 insertions(+), 59 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 17bba2c16e9c7..87cc828c1a81e 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1270,20 +1270,8 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler { llvm::DenseSet Visited; checkSYCLType(SemaRef, TemplateArgTy, Loc, Visited); - if (TAL.size() < 6) { - // Not enough arguments for this parameter pack. - SemaRef.Diag(Loc.getBegin(), - diag::err_template_arg_list_different_arity) - << /*not enough args*/ 0 - << (int)SemaRef.getTemplateNameKindForDiagnostics( - TemplateName(CTSD->getSpecializedTemplate())) - << CTSD; - SemaRef.Diag(CTSD->getLocation(), diag::note_template_decl_here) - << CTSD->getSourceRange(); - return; - } - - checkPropertyListType(TAL.get(5), Loc.getBegin()); + if (TAL.size() > 5) + checkPropertyListType(TAL.get(5), Loc.getBegin()); } } @@ -3472,14 +3460,15 @@ bool Util::isPropertyListType(const QualType &Ty) { } bool Util::isSyclBufferLocationType(const QualType &Ty) { - const StringRef &Name = "buffer_location"; + const StringRef &PropertyName = "buffer_location"; + const StringRef &InstanceName = "instance"; std::array Scopes = { - Util::DeclContextDesc{clang::Decl::Kind::Namespace, "cl"}, - Util::DeclContextDesc{clang::Decl::Kind::Namespace, "sycl"}, - Util::DeclContextDesc{clang::Decl::Kind::Namespace, "ext"}, - Util::DeclContextDesc{clang::Decl::Kind::Namespace, "INTEL"}, - Util::DeclContextDesc{clang::Decl::Kind::Namespace, "property"}, - Util::DeclContextDesc{Decl::Kind::ClassTemplateSpecialization, Name}}; + Util::DeclContextDesc{Decl::Kind::Namespace, "cl"}, + Util::DeclContextDesc{Decl::Kind::Namespace, "sycl"}, + Util::DeclContextDesc{Decl::Kind::Namespace, "INTEL"}, + Util::DeclContextDesc{Decl::Kind::Namespace, "property"}, + Util::DeclContextDesc{Decl::Kind::CXXRecord, PropertyName}, + Util::DeclContextDesc{Decl::Kind::ClassTemplateSpecialization, InstanceName}}; return matchQualifiedTypeName(Ty, Scopes); } @@ -3495,10 +3484,9 @@ bool Util::isSyclType(const QualType &Ty, StringRef Name, bool Tmpl) { bool Util::isAccessorPropertyListType(const QualType &Ty) { const StringRef &Name = "accessor_property_list"; - std::array Scopes = { + std::array Scopes = { Util::DeclContextDesc{clang::Decl::Kind::Namespace, "cl"}, Util::DeclContextDesc{clang::Decl::Kind::Namespace, "sycl"}, - Util::DeclContextDesc{clang::Decl::Kind::Namespace, "ext"}, Util::DeclContextDesc{clang::Decl::Kind::Namespace, "ONEAPI"}, Util::DeclContextDesc{Decl::Kind::ClassTemplateSpecialization, Name}}; return matchQualifiedTypeName(Ty, Scopes); diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index da1a78f4ac145..92dc64d04cd08 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -94,12 +94,12 @@ class property_list { bool operator!=(const property_list &rhs) const { return false; } }; -namespace ext { namespace INTEL { namespace property { // Compile time known accessor property -template -class buffer_location {}; +struct buffer_location { + template class instance {}; +}; } // namespace property } // namespace INTEL @@ -107,7 +107,6 @@ namespace ONEAPI { template class accessor_property_list {}; } // namespace ONEAPI -} // namespace ext template struct id { diff --git a/clang/test/CodeGenSYCL/buffer_location.cpp b/clang/test/CodeGenSYCL/buffer_location.cpp index b18e8ecef8234..7e1fa2b459a56 100644 --- a/clang/test/CodeGenSYCL/buffer_location.cpp +++ b/clang/test/CodeGenSYCL/buffer_location.cpp @@ -10,8 +10,8 @@ struct Base { cl::sycl::accessor>> + cl::sycl::ONEAPI::accessor_property_list< + cl::sycl::INTEL::property::buffer_location::instance<2>>> AccField; }; @@ -19,8 +19,8 @@ struct Captured : Base, cl::sycl::accessor>> { + cl::sycl::ONEAPI::accessor_property_list< + cl::sycl::INTEL::property::buffer_location::instance<2>>> { int C; }; @@ -29,8 +29,8 @@ int main() { cl::sycl::accessor>> + cl::sycl::ONEAPI::accessor_property_list< + cl::sycl::INTEL::property::buffer_location::instance<3>>> accessorA; cl::sycl::kernel_single_task( [=]() { diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp index 2e9bab2a2492a..f2bfe5e357041 100644 --- a/clang/test/SemaSYCL/Inputs/sycl.hpp +++ b/clang/test/SemaSYCL/Inputs/sycl.hpp @@ -39,11 +39,11 @@ enum class address_space : int { class property_list {}; -namespace ext { namespace INTEL { namespace property { -template -class buffer_location {}; +struct buffer_location { + template class instance {}; +}; } // namespace property } // namespace INTEL @@ -51,7 +51,6 @@ namespace ONEAPI { template class accessor_property_list {}; } // namespace ONEAPI -} // namespace ext namespace detail { namespace half_impl { @@ -103,7 +102,7 @@ struct DeviceValueType { template + typename propertyListT = ONEAPI::accessor_property_list<>> class accessor { public: diff --git a/clang/test/SemaSYCL/accessor_inheritance.cpp b/clang/test/SemaSYCL/accessor_inheritance.cpp index f286706aeafd9..bc3598208ca11 100644 --- a/clang/test/SemaSYCL/accessor_inheritance.cpp +++ b/clang/test/SemaSYCL/accessor_inheritance.cpp @@ -42,8 +42,8 @@ int main() { // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_A' 'int' // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_B' 'int' -// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::accessor':'cl::sycl::accessor' 'void () noexcept' -// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::accessor':'cl::sycl::accessor' 'void () noexcept' +// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::accessor':'cl::sycl::accessor>' 'void () noexcept' +// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::accessor':'cl::sycl::accessor>' 'void () noexcept' // CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_C' 'int' diff --git a/clang/test/SemaSYCL/basic-kernel-wrapper.cpp b/clang/test/SemaSYCL/basic-kernel-wrapper.cpp index d953cbbb4aa19..bba39fdddf9ba 100644 --- a/clang/test/SemaSYCL/basic-kernel-wrapper.cpp +++ b/clang/test/SemaSYCL/basic-kernel-wrapper.cpp @@ -43,7 +43,7 @@ int main() { // CHECK: CXXMemberCallExpr {{.*}} 'void' // 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: MemberExpr {{.*}} 'cl::sycl::accessor':'cl::sycl::accessor>' lvalue . // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}basic-kernel-wrapper.cpp{{.*}})' lvalue Var // CHECK-NEXT: ImplicitCastExpr {{.*}} diff --git a/clang/test/SemaSYCL/buffer_location.cpp b/clang/test/SemaSYCL/buffer_location.cpp index 5cb02923a9751..ac6dc527b5ae2 100644 --- a/clang/test/SemaSYCL/buffer_location.cpp +++ b/clang/test/SemaSYCL/buffer_location.cpp @@ -14,8 +14,8 @@ struct Base { cl::sycl::accessor>> + cl::sycl::ONEAPI::accessor_property_list< + cl::sycl::INTEL::property::buffer_location::instance<1>>> AccField; }; @@ -23,8 +23,8 @@ struct Captured : Base, cl::sycl::accessor>> { + cl::sycl::ONEAPI::accessor_property_list< + cl::sycl::INTEL::property::buffer_location::instance<1>>> { int C; }; @@ -35,30 +35,30 @@ int main() { cl::sycl::accessor>> + cl::sycl::ONEAPI::accessor_property_list< + cl::sycl::INTEL::property::buffer_location::instance<2>>> // CHECK: SYCLIntelBufferLocationAttr {{.*}} Implicit 2 accessorA; cl::sycl::accessor>> + cl::sycl::INTEL::property::buffer_location::instance<3>>> // CHECK: SYCLIntelBufferLocationAttr {{.*}} Implicit 3 accessorB; cl::sycl::accessor> accessorC; #else cl::sycl::accessor>> + cl::sycl::ONEAPI::accessor_property_list< + cl::sycl::INTEL::property::buffer_location::instance<-2>>> accessorD; cl::sycl::accessor, - cl::sycl::ext::INTEL::property::buffer_location<2>>> + cl::sycl::ONEAPI::accessor_property_list< + cl::sycl::INTEL::property::buffer_location::instance<1>, + cl::sycl::INTEL::property::buffer_location::instance<2>>> accessorF; #endif cl::sycl::kernel_single_task( diff --git a/clang/test/SemaSYCL/wrapped-accessor.cpp b/clang/test/SemaSYCL/wrapped-accessor.cpp index 9347c0999e34b..c3dec6b9f3236 100644 --- a/clang/test/SemaSYCL/wrapped-accessor.cpp +++ b/clang/test/SemaSYCL/wrapped-accessor.cpp @@ -35,14 +35,14 @@ int main() { // argument // CHECK: VarDecl {{.*}}'(lambda at {{.*}}wrapped-accessor.cpp{{.*}})' // CHECK-NEXT: InitListExpr {{.*}}'(lambda at {{.*}}wrapped-accessor.cpp{{.*}})' -// CHECK-NEXT: InitListExpr {{.*}}'AccWrapper>' -// CHECK-NEXT: CXXConstructExpr {{.*}}'cl::sycl::accessor':'cl::sycl::accessor' 'void () noexcept' +// CHECK-NEXT: InitListExpr {{.*}}'AccWrapper>>' +// CHECK-NEXT: CXXConstructExpr {{.*}}'cl::sycl::accessor>':'cl::sycl::accessor>' 'void () noexcept' // Check that accessor field of the wrapper object is initialized using __init method // CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void' // 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: 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 From 1813143a1dbbb876eed12d99d0a42b4c1e49b421 Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Wed, 9 Sep 2020 12:58:31 +0300 Subject: [PATCH 3/9] Remove some unused code --- clang/lib/Sema/SemaSYCL.cpp | 13 ------------- 1 file changed, 13 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 87cc828c1a81e..7080352f48bc5 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -81,10 +81,6 @@ class Util { /// half class. static bool isSyclHalfType(const QualType &Ty); - /// Checks whether given clang type is a full specialization of the SYCL - /// property_list class. - static bool isPropertyListType(const QualType &Ty); - /// Checks whether given clang type is a full specialization of the SYCL /// accessor_property_list class. static bool isAccessorPropertyListType(const QualType &Ty); @@ -1198,8 +1194,6 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler { return; } QualType PropListTy = PropList.getAsType(); - if (Util::isPropertyListType(PropListTy)) - return; if (!Util::isAccessorPropertyListType(PropListTy)) { SemaRef.Diag(Loc, diag::err_sycl_invalid_accessor_property_template_param); @@ -1421,9 +1415,6 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { return; const auto PropList = cast(AccTy->getTemplateArgs()[5]); QualType PropListTy = PropList.getAsType(); - // property_list contains runtime properties, it shouldn't be handled here. - if (Util::isPropertyListType(PropListTy)) - return; const auto *AccPropListDecl = cast(PropListTy->getAsRecordDecl()); const auto TemplArg = AccPropListDecl->getTemplateArgs()[0]; @@ -3455,10 +3446,6 @@ bool Util::isSyclSpecConstantType(const QualType &Ty) { return matchQualifiedTypeName(Ty, Scopes); } -bool Util::isPropertyListType(const QualType &Ty) { - return isSyclType(Ty, "property_list"); -} - bool Util::isSyclBufferLocationType(const QualType &Ty) { const StringRef &PropertyName = "buffer_location"; const StringRef &InstanceName = "instance"; From e638f101910403e9f7eb3a69c6c9c02cc1008cbd Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Wed, 9 Sep 2020 12:59:14 +0300 Subject: [PATCH 4/9] Run clang-format --- clang/lib/Sema/SemaSYCL.cpp | 13 +++++++------ 1 file changed, 7 insertions(+), 6 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 7080352f48bc5..e5cda3bad514f 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -3450,12 +3450,13 @@ bool Util::isSyclBufferLocationType(const QualType &Ty) { const StringRef &PropertyName = "buffer_location"; const StringRef &InstanceName = "instance"; std::array Scopes = { - Util::DeclContextDesc{Decl::Kind::Namespace, "cl"}, - Util::DeclContextDesc{Decl::Kind::Namespace, "sycl"}, - Util::DeclContextDesc{Decl::Kind::Namespace, "INTEL"}, - Util::DeclContextDesc{Decl::Kind::Namespace, "property"}, - Util::DeclContextDesc{Decl::Kind::CXXRecord, PropertyName}, - Util::DeclContextDesc{Decl::Kind::ClassTemplateSpecialization, InstanceName}}; + Util::DeclContextDesc{Decl::Kind::Namespace, "cl"}, + Util::DeclContextDesc{Decl::Kind::Namespace, "sycl"}, + Util::DeclContextDesc{Decl::Kind::Namespace, "INTEL"}, + Util::DeclContextDesc{Decl::Kind::Namespace, "property"}, + Util::DeclContextDesc{Decl::Kind::CXXRecord, PropertyName}, + Util::DeclContextDesc{Decl::Kind::ClassTemplateSpecialization, + InstanceName}}; return matchQualifiedTypeName(Ty, Scopes); } From ef2cd88ad1edb83dc384e4a6d52cc22d00265ecd Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Wed, 9 Sep 2020 15:11:19 +0300 Subject: [PATCH 5/9] Fix tests --- clang/test/CodeGenSYCL/Inputs/sycl.hpp | 2 +- clang/test/CodeGenSYCL/accessor_inheritance.cpp | 4 ++-- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index 92dc64d04cd08..e266ae3d3bea5 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -142,7 +142,7 @@ struct _ImplT { template + typename propertyListT = ONEAPI::accessor_property_list<>> class accessor { public: diff --git a/clang/test/CodeGenSYCL/accessor_inheritance.cpp b/clang/test/CodeGenSYCL/accessor_inheritance.cpp index b94c9fa7a9478..2718819e95f91 100644 --- a/clang/test/CodeGenSYCL/accessor_inheritance.cpp +++ b/clang/test/CodeGenSYCL/accessor_inheritance.cpp @@ -67,13 +67,13 @@ int main() { // CHECK: [[ACC_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base* [[BITCAST]], i32 0, i32 2 // CHECK: [[ACC1_AS_CAST:%[a-zA-Z0-9_]+]] = addrspacecast %"class{{.*}}cl::sycl::accessor"* [[ACC_FIELD]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)* // Default constructor call -// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_13property_listEEC1Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[ACC1_AS_CAST]]) +// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEC1Ev(%"class._ZTSN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* [[ACC1_AS_CAST]]) // CHECK: [[BITCAST1:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured* [[GEP]] to i8* // CHECK: [[GEP1:%[a-zA-Z0-9_]+]] = getelementptr inbounds i8, i8* [[BITCAST1]], i64 20 // CHECK: [[BITCAST2:%[a-zA-Z0-9_]+]] = bitcast i8* [[GEP1]] to %"class{{.*}}cl::sycl::accessor"* // CHECK: [[ACC2_AS_CAST:%[a-zA-Z0-9_]+]] = addrspacecast %"class{{.*}}cl::sycl::accessor"* [[BITCAST2]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)* // Default constructor call -// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_13property_listEEC2Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[ACC2_AS_CAST]]) +// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEC2Ev(%"class._ZTSN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* [[ACC2_AS_CAST]]) // CHECK C field initialization // CHECK: [[FIELD_C:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Captured, %struct{{.*}}Captured* [[GEP]], i32 0, i32 2 From f6f1806ecdd024e7a35edaa41301fdd5c3da72a9 Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Thu, 10 Sep 2020 14:55:51 +0300 Subject: [PATCH 6/9] Address feedback --- clang/include/clang/Basic/DiagnosticSemaKinds.td | 4 ++-- clang/lib/Sema/SemaSYCL.cpp | 8 +++++--- clang/test/SemaSYCL/buffer_location.cpp | 2 +- 3 files changed, 8 insertions(+), 6 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 5aff495d36f84..6d78279aa7d16 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -11033,8 +11033,8 @@ def err_sycl_compiletime_property_duplication : Error< def err_sycl_invalid_property_list_param_number : Error< "%0 must have exactly one template parameter">; def err_sycl_invalid_accessor_property_template_param : Error< - "Sixth template parameter of the accessor must be of accessor_property_list " - "or property_list type">; + "sixth template parameter of the accessor must be of accessor_property_list " + "type">; def err_sycl_invalid_accessor_property_list_template_param : Error< "%select{accessor_property_list|accessor_property_list pack argument|buffer_location}0 " "template parameter must be a " diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index e5cda3bad514f..8f79b2f0806d9 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1256,16 +1256,18 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler { "Should only be called on SYCL accessor types."); RecordDecl *RecD = Ty->getAsRecordDecl(); - if (auto *CTSD = dyn_cast(RecD)) { + if (const ClassTemplateSpecializationDecl *CTSD = + dyn_cast(RecD)) { const TemplateArgumentList &TAL = CTSD->getTemplateArgs(); TemplateArgument TA = TAL.get(0); const QualType TemplateArgTy = TA.getAsType(); - llvm::DenseSet Visited; - checkSYCLType(SemaRef, TemplateArgTy, Loc, Visited); if (TAL.size() > 5) checkPropertyListType(TAL.get(5), Loc.getBegin()); + + llvm::DenseSet Visited; + checkSYCLType(SemaRef, TemplateArgTy, Loc, Visited); } } diff --git a/clang/test/SemaSYCL/buffer_location.cpp b/clang/test/SemaSYCL/buffer_location.cpp index ac6dc527b5ae2..f665549bf0b85 100644 --- a/clang/test/SemaSYCL/buffer_location.cpp +++ b/clang/test/SemaSYCL/buffer_location.cpp @@ -84,7 +84,7 @@ int main() { #else //expected-error@+1{{buffer_location template parameter must be a non-negative integer}} accessorD.use(); - //expected-error@+1{{Sixth template parameter of the accessor must be of accessor_property_list or property_list type}} + //expected-error@+1{{sixth template parameter of the accessor must be of accessor_property_list type}} accessorE.use(); //expected-error@+1{{Can't apply buffer_location property twice to the same accessor}} accessorF.use(); From 7b425629f7be4d4d0690602f785daaa0db395dd2 Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Fri, 11 Sep 2020 08:45:11 +0300 Subject: [PATCH 7/9] Address comments --- clang/lib/Sema/SemaSYCL.cpp | 4 +--- clang/test/CodeGenSYCL/accessor_inheritance.cpp | 4 ++-- 2 files changed, 3 insertions(+), 5 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 8f79b2f0806d9..deea2365fa084 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1255,17 +1255,15 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler { assert(Util::isSyclAccessorType(Ty) && "Should only be called on SYCL accessor types."); - RecordDecl *RecD = Ty->getAsRecordDecl(); + const RecordDecl *RecD = Ty->getAsRecordDecl(); if (const ClassTemplateSpecializationDecl *CTSD = dyn_cast(RecD)) { const TemplateArgumentList &TAL = CTSD->getTemplateArgs(); - TemplateArgument TA = TAL.get(0); const QualType TemplateArgTy = TA.getAsType(); if (TAL.size() > 5) checkPropertyListType(TAL.get(5), Loc.getBegin()); - llvm::DenseSet Visited; checkSYCLType(SemaRef, TemplateArgTy, Loc, Visited); } diff --git a/clang/test/CodeGenSYCL/accessor_inheritance.cpp b/clang/test/CodeGenSYCL/accessor_inheritance.cpp index 2718819e95f91..0edc4f445da24 100644 --- a/clang/test/CodeGenSYCL/accessor_inheritance.cpp +++ b/clang/test/CodeGenSYCL/accessor_inheritance.cpp @@ -67,13 +67,13 @@ int main() { // CHECK: [[ACC_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, %struct{{.*}}Base* [[BITCAST]], i32 0, i32 2 // CHECK: [[ACC1_AS_CAST:%[a-zA-Z0-9_]+]] = addrspacecast %"class{{.*}}cl::sycl::accessor"* [[ACC_FIELD]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)* // Default constructor call -// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEC1Ev(%"class._ZTSN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* [[ACC1_AS_CAST]]) +// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEC1Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[ACC1_AS_CAST]]) // CHECK: [[BITCAST1:%[a-zA-Z0-9_]+]] = bitcast %struct{{.*}}Captured* [[GEP]] to i8* // CHECK: [[GEP1:%[a-zA-Z0-9_]+]] = getelementptr inbounds i8, i8* [[BITCAST1]], i64 20 // CHECK: [[BITCAST2:%[a-zA-Z0-9_]+]] = bitcast i8* [[GEP1]] to %"class{{.*}}cl::sycl::accessor"* // CHECK: [[ACC2_AS_CAST:%[a-zA-Z0-9_]+]] = addrspacecast %"class{{.*}}cl::sycl::accessor"* [[BITCAST2]] to %"class{{.*}}cl::sycl::accessor" addrspace(4)* // Default constructor call -// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEC2Ev(%"class._ZTSN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEE.cl::sycl::accessor" addrspace(4)* [[ACC2_AS_CAST]]) +// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_6ONEAPI22accessor_property_listIJEEEEC2Ev(%"class{{.*}}cl::sycl::accessor" addrspace(4)* [[ACC2_AS_CAST]]) // CHECK C field initialization // CHECK: [[FIELD_C:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Captured, %struct{{.*}}Captured* [[GEP]], i32 0, i32 2 From 384152e7377ca1cd97b2e31721a15c7e608860cf Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Fri, 11 Sep 2020 09:06:04 +0300 Subject: [PATCH 8/9] Fix formatting --- clang/test/SemaSYCL/buffer_location.cpp | 30 ++++++++++++------------- 1 file changed, 15 insertions(+), 15 deletions(-) diff --git a/clang/test/SemaSYCL/buffer_location.cpp b/clang/test/SemaSYCL/buffer_location.cpp index f665549bf0b85..96a7094948344 100644 --- a/clang/test/SemaSYCL/buffer_location.cpp +++ b/clang/test/SemaSYCL/buffer_location.cpp @@ -9,22 +9,24 @@ template class another_property_list { }; +template +using buffer_location = cl::sycl::INTEL::property::buffer_location::instance; + struct Base { int A, B; cl::sycl::accessor>> + cl::sycl::ONEAPI::accessor_property_list>> AccField; }; -struct Captured : Base, - cl::sycl::accessor>> { +struct Captured + : Base, + cl::sycl::accessor>> { int C; }; @@ -35,8 +37,7 @@ int main() { cl::sycl::accessor>> + cl::sycl::ONEAPI::accessor_property_list>> // CHECK: SYCLIntelBufferLocationAttr {{.*}} Implicit 2 accessorA; cl::sycl::accessor>> + buffer_location<3>>> // CHECK: SYCLIntelBufferLocationAttr {{.*}} Implicit 3 accessorB; cl::sycl::accessor>> + cl::sycl::ONEAPI::accessor_property_list>> accessorD; cl::sycl::accessor, - cl::sycl::INTEL::property::buffer_location::instance<2>>> + buffer_location<1>, + buffer_location<2>>> accessorF; #endif cl::sycl::kernel_single_task( From 9201655cfc3bdda8b893199f5a3348d0904b06b9 Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Fri, 11 Sep 2020 09:10:16 +0300 Subject: [PATCH 9/9] Remove extra whitespace --- clang/test/SemaSYCL/buffer_location.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/test/SemaSYCL/buffer_location.cpp b/clang/test/SemaSYCL/buffer_location.cpp index 96a7094948344..766cbbb893a75 100644 --- a/clang/test/SemaSYCL/buffer_location.cpp +++ b/clang/test/SemaSYCL/buffer_location.cpp @@ -21,7 +21,7 @@ struct Base { AccField; }; -struct Captured +struct Captured : Base, cl::sycl::accessor