diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index c532b2272ad6e..6d78279aa7d16 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 " + "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..deea2365fa084 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -82,8 +82,8 @@ class Util { 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); + /// 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. @@ -1194,29 +1194,31 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler { return; } QualType PropListTy = PropList.getAsType(); - if (!Util::isPropertyListType(PropListTy)) { + 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 +1237,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; } @@ -1402,19 +1406,18 @@ 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 = + 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(); @@ -3443,20 +3446,17 @@ bool Util::isSyclSpecConstantType(const QualType &Ty) { return matchQualifiedTypeName(Ty, Scopes); } -bool Util::isPropertyListType(const QualType &Ty) { - return isSyclType(Ty, "property_list", true /*Tmpl*/); -} - bool Util::isSyclBufferLocationType(const QualType &Ty) { - const StringRef &Name = "buffer_location"; - 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, "property"}, - Util::DeclContextDesc{Decl::Kind::ClassTemplateSpecialization, Name}}; + 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}}; return matchQualifiedTypeName(Ty, Scopes); } @@ -3470,6 +3470,16 @@ 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, "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..e266ae3d3bea5 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,20 @@ class property_list { bool operator!=(const property_list &rhs) const { return false; } }; +namespace INTEL { +namespace property { +// Compile time known accessor property +struct buffer_location { + template class instance {}; +}; +} // namespace property +} // namespace INTEL + +namespace ONEAPI { +template +class accessor_property_list {}; +} // namespace ONEAPI + template struct id { template @@ -136,7 +142,7 @@ struct _ImplT { template > + typename propertyListT = ONEAPI::accessor_property_list<>> class accessor { public: @@ -150,8 +156,6 @@ class accessor { private: void __init(__attribute__((opencl_global)) dataT *Ptr, range AccessRange, range MemRange, id Offset) {} - - propertyListT prop_list; }; template @@ -339,8 +343,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 +355,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 +429,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..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_13property_listIJEEEEC1Ev(%"class{{.*}}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 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_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 diff --git a/clang/test/CodeGenSYCL/buffer_location.cpp b/clang/test/CodeGenSYCL/buffer_location.cpp index ae1b9088cc271..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/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..f2bfe5e357041 100644 --- a/clang/test/SemaSYCL/Inputs/sycl.hpp +++ b/clang/test/SemaSYCL/Inputs/sycl.hpp @@ -37,13 +37,20 @@ enum class address_space : int { }; } // namespace access +class property_list {}; + +namespace INTEL { namespace property { -template -class buffer_location {}; +struct buffer_location { + template class instance {}; +}; } // namespace property +} // namespace INTEL +namespace ONEAPI { template -class property_list {}; +class accessor_property_list {}; +} // namespace ONEAPI namespace detail { namespace half_impl { @@ -95,7 +102,7 @@ struct DeviceValueType { template > + typename propertyListT = ONEAPI::accessor_property_list<>> class accessor { public: @@ -107,7 +114,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..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 9d32dd1445453..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 65742af0cbe1c..766cbbb893a75 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,28 +37,28 @@ int main() { cl::sycl::accessor>> + cl::sycl::ONEAPI::accessor_property_list>> // CHECK: SYCLIntelBufferLocationAttr {{.*}} Implicit 2 accessorA; cl::sycl::accessor>> + cl::sycl::ONEAPI::accessor_property_list< + another_property, + buffer_location<3>>> // CHECK: SYCLIntelBufferLocationAttr {{.*}} Implicit 3 accessorB; cl::sycl::accessor> + cl::sycl::ONEAPI::accessor_property_list< + another_property>> accessorC; #else cl::sycl::accessor>> + cl::sycl::ONEAPI::accessor_property_list>> accessorD; cl::sycl::accessor, - cl::sycl::property::buffer_location<2>>> + cl::sycl::ONEAPI::accessor_property_list< + buffer_location<1>, + 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 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..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