diff --git a/clang/lib/CodeGen/CGDeclCXX.cpp b/clang/lib/CodeGen/CGDeclCXX.cpp index 4719dc4b369d8..c846e4dead9b9 100644 --- a/clang/lib/CodeGen/CGDeclCXX.cpp +++ b/clang/lib/CodeGen/CGDeclCXX.cpp @@ -73,6 +73,20 @@ static void EmitDeclInit(CodeGenFunction &CGF, const VarDecl &D, /// static storage duration. static void EmitDeclDestroy(CodeGenFunction &CGF, const VarDecl &D, ConstantAddress Addr) { + // SYCL work-group scope variables (e.g. the kernel object of a + // parallel_for_work_group kernel) are local variables promoted to + // work-group local memory. There is no program exit on the device to run + // their destructor at, and SYCL 2020 requires the destructors of device + // copyable types to have no effect on the device, so don't register one. + if (CGF.getLangOpts().SYCLIsDevice) { + const SYCLScopeAttr *Scope = D.getAttr(); + if (!Scope) + if (const auto *RD = D.getType()->getAsCXXRecordDecl()) + Scope = RD->getAttr(); + if (Scope && Scope->isWorkGroup()) + return; + } + // Honor __attribute__((no_destroy)) and bail instead of attempting // to emit a reference to a possibly nonexistent destructor, which // in turn can cause a crash. This will result in a global constructor diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 11d2e5e001629..b8584adb7ca83 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -4547,7 +4547,7 @@ class FreeFunctionKernelBodyCreator : public SyclKernelFieldHandler { // wgm.__init(arg); // user_kernel(some arguments..., wgm, some arguments...); // } - const auto *RecordDecl = ParamTy->getAsCXXRecordDecl(); + auto *RecordDecl = ParamTy->getAsCXXRecordDecl(); AccessSpecifier DefaultConstructorAccess; auto DefaultConstructor = std::find_if(RecordDecl->ctor_begin(), RecordDecl->ctor_end(), @@ -4571,6 +4571,8 @@ class FreeFunctionKernelBodyCreator : public SyclKernelFieldHandler { SpecialObjectClone->setInit( SemaSYCLRef.SemaRef.MaybeCreateExprWithCleanups(Init.get())); SpecialObjectClone->setInitStyle(VarDecl::CallInit); + SemaSYCLRef.SemaRef.FinalizeVarWithDestructor(SpecialObjectClone, + RecordDecl); DefaultConstructor->setAccess(DefaultConstructorAccess); Stmt *DS = new (SemaSYCLRef.getASTContext()) @@ -7678,11 +7680,11 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { // contains information about the offset, size and parameter // kind of every field inside the struct at any nesting level // This facilitates setting the arguments in the runtime. - // We also define is_device_copyable trait to be true for this type to - // allow it being passed in device kernels. + // We also define detail::is_legal_kernel_parameter trait to be true for + // this type to allow it being passed in device kernels. O << "template <>\n"; O << "struct " - "sycl::is_device_copyable<"; + "sycl::detail::is_legal_kernel_parameter<"; Policy.SuppressTagKeyword = true; type.print(O, Policy); O << ">: std::true_type {};\n"; diff --git a/clang/test/CodeGenSYCL/free_function_int_header.cpp b/clang/test/CodeGenSYCL/free_function_int_header.cpp index 0e2851ac20cd4..930835bf1e3b1 100644 --- a/clang/test/CodeGenSYCL/free_function_int_header.cpp +++ b/clang/test/CodeGenSYCL/free_function_int_header.cpp @@ -1609,7 +1609,7 @@ void ff_28(TemplatedAccessorStruct arg1) { // CHECK: Forward declarations of kernel and its argument types: // CHECK: void ff_25(AccessorAndLocalAccessor arg1); // CHECK-NEXT: template <> -// CHECK-NEXT: struct sycl::is_device_copyable: std::true_type {}; +// CHECK-NEXT: struct sycl::detail::is_legal_kernel_parameter: std::true_type {}; // CHECK-NEXT: template <> // CHECK-NEXT: struct sycl::ext::oneapi::experimental::detail::is_struct_with_special_type { // CHECK-NEXT: inline static constexpr bool value = true; @@ -1637,7 +1637,7 @@ void ff_28(TemplatedAccessorStruct arg1) { // CHECK: Forward declarations of kernel and its argument types: // CHECK: void ff_26(AccessorAndLocalAccessor arg1, SecondLevelAccessor arg2); // CHECK-NEXT: template <> -// CHECK-NEXT: struct sycl::is_device_copyable: std::true_type {}; +// CHECK-NEXT: struct sycl::detail::is_legal_kernel_parameter: std::true_type {}; // CHECK-NEXT: template <> // CHECK-NEXT: struct sycl::ext::oneapi::experimental::detail::is_struct_with_special_type { // CHECK-NEXT: inline static constexpr bool value = true; @@ -1665,7 +1665,7 @@ void ff_28(TemplatedAccessorStruct arg1) { // CHECK: Forward declarations of kernel and its argument types: // CHECK: void ff_27(IntAndAccessor arg1, AccessorAndInt ); // CHECK-NEXT: template <> -// CHECK-NEXT: struct sycl::is_device_copyable: std::true_type {}; +// CHECK-NEXT: struct sycl::detail::is_legal_kernel_parameter: std::true_type {}; // CHECK-NEXT: template <> // CHECK-NEXT: struct sycl::ext::oneapi::experimental::detail::is_struct_with_special_type { // CHECK-NEXT: inline static constexpr bool value = true; @@ -1678,7 +1678,7 @@ void ff_28(TemplatedAccessorStruct arg1) { // CHECK-NEXT: }; // CHECK: template <> -// CHECK-NEXT: struct sycl::is_device_copyable: std::true_type {}; +// CHECK-NEXT: struct sycl::detail::is_legal_kernel_parameter: std::true_type {}; // CHECK-NEXT: template <> // CHECK-NEXT: struct sycl::ext::oneapi::experimental::detail::is_struct_with_special_type { // CHECK-NEXT: inline static constexpr bool value = true; @@ -1710,7 +1710,7 @@ void ff_28(TemplatedAccessorStruct arg1) { // CHECK: template struct TemplatedAccessorStruct; // CHECK: void ff_28(TemplatedAccessorStruct arg1); // CHECK-NEXT: template <> -// CHECK-NEXT: struct sycl::is_device_copyable>: std::true_type {}; +// CHECK-NEXT: struct sycl::detail::is_legal_kernel_parameter>: std::true_type {}; // CHECK-NEXT: template <> // CHECK-NEXT: struct sycl::ext::oneapi::experimental::detail::is_struct_with_special_type> { // CHECK-NEXT: inline static constexpr bool value = true; diff --git a/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp b/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp index 7e1712a4271f8..4630c400e6b25 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp @@ -39,10 +39,12 @@ static const char *LegalSYCLFunctions[] = { "^sycl::_V1::accessor<.+>::__init_esimd", "^sycl::_V1::address_space_cast", "^sycl::_V1::local_accessor<.+>::local_accessor", + "^sycl::_V1::local_accessor<.+>::~local_accessor", "^sycl::_V1::local_accessor<.+>::__init_esimd", "^sycl::_V1::local_accessor<.+>::get_pointer", "^sycl::_V1::local_accessor<.+>::get_multi_ptr", "^sycl::_V1::local_accessor_base<.+>::local_accessor_base", + "^sycl::_V1::local_accessor_base<.+>::~local_accessor_base", "^sycl::_V1::local_accessor_base<.+>::__init_esimd", "^sycl::_V1::local_accessor_base<.+>::getQualifiedPtr", "^sycl::_V1::local_accessor_base<.+>::getSize", diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp index 0d8ea29488948..09b4c6d409576 100644 --- a/sycl/include/sycl/accessor.hpp +++ b/sycl/include/sycl/accessor.hpp @@ -757,6 +757,13 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor : : impl({}, detail::InitializedVal::template get<0>(), detail::InitializedVal::template get<0>()) {} +#ifdef __INTEL_PREVIEW_BREAKING_CHANGES + ~accessor() { + // User-declared so the type is not trivially copyable; intentionally empty + // as the spec states the destructor of an accessor must not have any effect + } +#endif // __INTEL_PREVIEW_BREAKING_CHANGES + #else accessor(const detail::AccessorImplPtr &Impl) : detail::AccessorBaseHost{Impl} {} @@ -821,6 +828,14 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor : /*IsPlaceH=*/false, /*OffsetInBytes=*/0, /*IsSubBuffer=*/false, /*PropertyList=*/{}){}; +#ifdef __INTEL_PREVIEW_BREAKING_CHANGES + ~accessor() { + // User-declared so the type is not trivially copyable; intentionally empty + // as the spec states the destructor of an accessor must not have any + // effect + } +#endif // __INTEL_PREVIEW_BREAKING_CHANGES + template friend class host_accessor; #endif // __SYCL_DEVICE_ONLY__ @@ -888,6 +903,12 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor : // -------+---------+-------+----+-----+-------------- public: + // common reference semantics + accessor(const accessor &) noexcept = default; + accessor &operator=(const accessor &) noexcept = default; + accessor(accessor &&) noexcept = default; + accessor &operator=(accessor &&) noexcept = default; + // implicit conversion between const / non-const types for read only accessors template + template < + typename... NewPropsT, + // Only viable when the compile-time properties match, so that CTAD from + // an accessor with a different property list deduces the source type. + typename = std::enable_if_t< + PropertyListT::template areSameCompileTimeProperties() +#ifdef __INTEL_PREVIEW_BREAKING_CHANGES + // Same-type copies are handled by the copy constructor. + && !std::is_same_v, + ext::oneapi::accessor_property_list> +#endif // __INTEL_PREVIEW_BREAKING_CHANGES + >> accessor( const accessor> &Other, @@ -1605,9 +1637,6 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor : { static_assert(detail::IsCxPropertyList::value, "Conversion is only available for accessor_property_list"); - static_assert( - PropertyListT::template areSameCompileTimeProperties(), - "Compile-time-constant properties must be the same"); (void)CodeLoc; #ifndef __SYCL_DEVICE_ONLY__ detail::constructorNotification(getMemoryObject(), impl.get(), AccessTarget, @@ -2241,6 +2270,20 @@ class __SYCL_SPECIAL_CLASS local_accessor_base : } #endif +#ifdef __INTEL_PREVIEW_BREAKING_CHANGES + ~local_accessor_base() { + // User-declared so the type is not trivially copyable; intentionally empty + // as the spec states the destructor of an accessor must not have any effect + } +#else + ~local_accessor_base() = default; +#endif // __INTEL_PREVIEW_BREAKING_CHANGES + + local_accessor_base(const local_accessor_base &) = default; + local_accessor_base &operator=(const local_accessor_base &) = default; + local_accessor_base(local_accessor_base &&) = default; + local_accessor_base &operator=(local_accessor_base &&) = default; + size_t get_size() const { return getSize().size() * sizeof(DataT); } __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead") @@ -2420,6 +2463,13 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor local_accessor(const detail::AccessorImplPtr &Impl) : local_acc{Impl} {} #endif +public: + // common reference semantics + local_accessor(const local_accessor &) noexcept = default; + local_accessor &operator=(const local_accessor &) noexcept = default; + local_accessor(local_accessor &&) noexcept = default; + local_accessor &operator=(local_accessor &&) noexcept = default; + // implicit conversion between non-const read-write accessor to const // read-only accessor public: @@ -2582,6 +2632,17 @@ class __SYCL_EBO host_accessor public: host_accessor() : AccessorT() {} +#ifdef __INTEL_PREVIEW_BREAKING_CHANGES + ~host_accessor() { + // User-declared so the type is not trivially copyable; intentionally empty + } +#endif // __INTEL_PREVIEW_BREAKING_CHANGES + + // common reference semantics + host_accessor(const host_accessor &) noexcept = default; + host_accessor &operator=(const host_accessor &) noexcept = default; + host_accessor(host_accessor &&) noexcept = default; + host_accessor &operator=(host_accessor &&) noexcept = default; // The list of host_accessor constructors with their arguments // -------+---------+-------+----+----------+-------------- @@ -2826,6 +2887,31 @@ host_accessor(buffer, Type1, Type2, Type3, Type4, Type5) -> host_accessor()>; +namespace detail { +// Add specializations for SYCL accessor types to the is_sycl_accessor trait. +// This is needed so that is_device_copyable returns false, but +// they can still be captured in a kernel. +template struct is_sycl_accessor; + +template +inline constexpr bool is_valid_kernel_access_target_v = + AccessTarget == access::target::device || + AccessTarget == access::target::global_buffer || + AccessTarget == access::target::constant_buffer || + AccessTarget == access::target::local || + AccessTarget == access::target::image || + AccessTarget == access::target::image_array; + +template +struct is_sycl_accessor> + : std::bool_constant> {}; + +template +struct is_sycl_accessor> : std::true_type {}; +} // namespace detail } // namespace _V1 } // namespace sycl diff --git a/sycl/include/sycl/accessor_image.hpp b/sycl/include/sycl/accessor_image.hpp index be5eb3a697ae4..621158ff3ff7d 100644 --- a/sycl/include/sycl/accessor_image.hpp +++ b/sycl/include/sycl/accessor_image.hpp @@ -638,6 +638,23 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor< access::target::image); #endif } + +#ifdef __INTEL_PREVIEW_BREAKING_CHANGES + ~accessor() { + // User-declared so the type is not trivially copyable; intentionally + // empty as the spec states the destructor of an accessor must not have + // any effect + } +#else + ~accessor() = default; +#endif + + // common reference semantics + accessor(const accessor &) noexcept = default; + accessor &operator=(const accessor &) noexcept = default; + accessor(accessor &&) noexcept = default; + accessor &operator=(accessor &&) noexcept = default; + #ifdef __SYCL_DEVICE_ONLY__ private: using OCLImageTy = @@ -688,6 +705,21 @@ class __SYCL_EBO accessor operator[](size_t Index) const { return detail::__image_array_slice__ returns false, but +// they can still be captured in a kernel. +template struct is_sycl_accessor; + +template +struct is_sycl_accessor> + : std::true_type {}; + +template +struct is_sycl_accessor< + sampled_image_accessor> + : std::true_type {}; +} // namespace detail } // namespace _V1 } // namespace sycl diff --git a/sycl/include/sycl/detail/is_device_copyable.hpp b/sycl/include/sycl/detail/is_device_copyable.hpp index a1720120afb1f..f4aaa28fc4754 100644 --- a/sycl/include/sycl/detail/is_device_copyable.hpp +++ b/sycl/include/sycl/detail/is_device_copyable.hpp @@ -105,68 +105,219 @@ template inline constexpr bool check_if_device_copyable_v = check_if_device_copyable::value; +// Deprecated sycl 1.2.1 rules for now +template +inline constexpr bool deprecated_sycl121_device_copyable_v = + std::is_trivially_copy_constructible_v && + std::is_trivially_destructible_v; + +template struct is_legal_kernel_parameter; + +template +inline constexpr bool is_legal_kernel_parameter_v = + is_legal_kernel_parameter::value; + +template struct is_sycl_accessor : std::false_type {}; + +template +inline constexpr bool is_sycl_accessor_v = is_sycl_accessor::value; + +#ifdef __SYCL_DEVICE_ONLY__ +template +class RoundedRangeKernel; +template +class RoundedRangeKernelWithKH; +#endif + +namespace legal_kernel_parameters { +// Some types are considered special SYCL types, such as SYCL accessors. +// These types are accepted legal parameters even if they do not satisfy the +// usual device copyability requirements. +template +struct is_special_sycl_type + : std::bool_constant>> {}; + +template +inline constexpr bool is_special_sycl_type_v = is_special_sycl_type::value; + #ifdef __SYCL_DEVICE_ONLY__ -template struct CheckFieldsAreDeviceCopyable; -template struct CheckBasesAreDeviceCopyable; +// We need to be able to check whether a record contains special SYCL types. +template +struct contains_special_sycl_types : is_special_sycl_type {}; + +template +inline constexpr bool contains_special_sycl_types_v = + contains_special_sycl_types::value; + +template struct special_sycl_types_fields; +template +struct special_sycl_types_fields> + : std::bool_constant<(contains_special_sycl_types_v< + decltype(__builtin_field_type(T, FieldIds))> || + ...)> {}; + +template +inline constexpr bool special_sycl_types_fields_v = + special_sycl_types_fields::value; + +template struct special_sycl_types_bases; +template +struct special_sycl_types_bases> + : std::bool_constant<(contains_special_sycl_types_v< + decltype(__builtin_base_type(T, BaseIds))> || + ...)> {}; + +template +inline constexpr bool special_sycl_types_bases_v = + special_sycl_types_bases::value; + +template +struct contains_special_sycl_types< + T, std::void_t> + : std::bool_constant< + is_special_sycl_type_v || + special_sycl_types_fields_v< + T, std::make_index_sequence<__builtin_num_fields(T)>> || + special_sycl_types_bases_v< + T, std::make_index_sequence<__builtin_num_bases(T)>>> {}; + +template +struct contains_special_sycl_types : contains_special_sycl_types {}; + +// In the case where any fields or base is a special SYCL type, we need to check +// whether ALL fields and bases are legal kernel parameters. +// This catches cases where a record contains special SYCL types but not all of +// its fields and bases are legal kernel parameters, but does not catch cases +// where a record containing special SYCL types also contains illegal member +// functions. +template +struct check_fields_and_bases : std::false_type {}; template -inline constexpr bool is_deprecated_device_copyable_v = - check_if_device_copyable_v || - (std::is_trivially_copy_constructible_v && - std::is_trivially_destructible_v); +inline constexpr bool check_fields_and_bases_v = + check_fields_and_bases::value; + +template struct check_fields; +template +struct check_fields> + : std::bool_constant<((is_legal_kernel_parameter_v>) && + ...)> {}; + +template +inline constexpr bool check_fields_v = check_fields::value; + +template struct check_bases; +template +struct check_bases> + : std::bool_constant<((is_legal_kernel_parameter_v>) && + ...)> {}; + +template +inline constexpr bool check_bases_v = check_bases::value; + +template +struct check_fields_and_bases< + T, std::void_t> + : std::bool_constant< + check_fields_v> && + check_bases_v>> { +}; + +template +struct check_fields_and_bases : check_fields_and_bases {}; +template struct validate_fields; template -struct CheckFieldsAreDeviceCopyable> { - static_assert(((is_deprecated_device_copyable_v< - decltype(__builtin_field_type(T, FieldIds))> && - ...)), - "The specified type is not device copyable"); +struct validate_fields> { + static_assert( + ((is_legal_kernel_parameter_v< + std::remove_cv_t> && + ...)), + "The specified type is not a legal kernel parameter"); }; +// Insert static assertions to validate fields and bases of a type at compile +// time. Used by kernel_launch_helper.hpp when wrapping kernel functors +template struct validate_bases; template -struct CheckBasesAreDeviceCopyable> { - static_assert(((is_deprecated_device_copyable_v< - decltype(__builtin_base_type(T, BaseIds))> && - ...)), - "The specified type is not device copyable"); +struct validate_bases> { + static_assert( + ((is_legal_kernel_parameter_v< + std::remove_cv_t> && + ...)), + "The specified type is not a legal kernel parameter"); }; // All the captures of a lambda or functor of type FuncT passed to a kernel -// must be is_device_copyable, which extends to bases and fields of FuncT. -// Fields are captures of lambda/functors and bases are possible base classes -// of functors also allowed by SYCL. -// The SYCL-2020 implementation must check each of the fields & bases of the -// type FuncT, only one level deep, which is enough to see if they are all -// device copyable by using the result of is_device_copyable returned for them. -// At this moment though the check also allowes using types for which -// (is_trivially_copy_constructible && is_trivially_destructible) returns true -// and (is_device_copyable) returns false. That is the deprecated behavior and -// is currently/temporarily supported only to not break older SYCL programs. +// must be legal kernel parameters. template -struct CheckDeviceCopyable - : CheckFieldsAreDeviceCopyable< - FuncT, std::make_index_sequence<__builtin_num_fields(FuncT)>>, - CheckBasesAreDeviceCopyable< - FuncT, std::make_index_sequence<__builtin_num_bases(FuncT)>> {}; +struct validate_functor + : validate_fields>, + validate_bases> {}; +// Below are two specializations for validate_functor when a kernel +// lambda is wrapped after range rounding optimization. template -class RoundedRangeKernel; -template -class RoundedRangeKernelWithKH; - -// Below are two specializations for CheckDeviceCopyable when a kernel lambda -// is wrapped after range rounding optimization. -template -struct CheckDeviceCopyable< +struct validate_functor< RoundedRangeKernel> - : CheckDeviceCopyable {}; + : validate_functor {}; template -struct CheckDeviceCopyable< +struct validate_functor< RoundedRangeKernelWithKH> - : CheckDeviceCopyable {}; + : validate_functor {}; + +#endif +} // namespace legal_kernel_parameters +// Is legal kernel parameter as defined by the SYCL-2020 specification. +// (Copied here) +// - Any device copyable type is a legal parameter type. +// - The following SYCL types are legal parameter types: +// - accessor when templated with target::device; +// - accessor when templated with any of the deprecated parameters: +// target::global_buffer, target::constant_buffer, or target::local; +// - local_accessor; +// - unsampled_image_accessor when templated with image_target::device; +// - sampled_image_accessor when templated with image_target::device; +// - stream; +// - id; +// - range; +// - marray when T is device copyable; +// - vec. +// - An array of element types T is a legal parameter type if T is a legal +// parameter type. +// - A class type S with a non-static member variable of type T is a legal +// parameter type if T is a legal parameter type and if S would otherwise be +// a legal parameter type aside from this member variable. +// - A class type S with a non-virtual base class of type T is a legal +// parameter type if T is a legal parameter type and if S would otherwise be +// a legal parameter type aside from this base class. +template +struct is_legal_kernel_parameter +#ifdef __SYCL_DEVICE_ONLY__ + : std::bool_constant< + check_if_device_copyable_v || + legal_kernel_parameters::is_special_sycl_type_v || + (legal_kernel_parameters::contains_special_sycl_types_v && + legal_kernel_parameters::check_fields_and_bases_v) || + deprecated_sycl121_device_copyable_v> { +}; + +#else + : std::bool_constant || + legal_kernel_parameters::is_special_sycl_type_v || + deprecated_sycl121_device_copyable_v> { +}; #endif // __SYCL_DEVICE_ONLY__ + +template +struct is_legal_kernel_parameter : is_legal_kernel_parameter {}; } // namespace detail } // namespace _V1 } // namespace sycl diff --git a/sycl/include/sycl/detail/kernel_launch_helper.hpp b/sycl/include/sycl/detail/kernel_launch_helper.hpp index a535b741ff830..b7db4bcd5b641 100644 --- a/sycl/include/sycl/detail/kernel_launch_helper.hpp +++ b/sycl/include/sycl/detail/kernel_launch_helper.hpp @@ -213,7 +213,7 @@ struct KernelWrapper< static void wrap([[maybe_unused]] const KernelType &KernelFunc) { #ifdef __SYCL_DEVICE_ONLY__ - detail::CheckDeviceCopyable(); + detail::legal_kernel_parameters::validate_functor(); #endif // Note: the static_assert below need to be run on both the host and the // device ends to avoid test issues, so don't put it into the #ifdef diff --git a/sycl/include/sycl/ext/oneapi/experimental/free_function_traits.hpp b/sycl/include/sycl/ext/oneapi/experimental/free_function_traits.hpp index 81a776eee544a..a9214aa1da210 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/free_function_traits.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/free_function_traits.hpp @@ -69,7 +69,9 @@ template struct kernel_function_s {}; template inline constexpr kernel_function_s kernel_function; } // namespace ext::oneapi::experimental -template struct is_device_copyable; +namespace detail { +template struct is_legal_kernel_parameter; +} } // namespace _V1 } // namespace sycl diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 56539906b52d3..15955e4d24683 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -528,14 +528,26 @@ class __SYCL_EXPORT handler { template void setArgHelper(int ArgIndex, T &&Arg) { void *StoredArg = storePlainArg(Arg); - if (!std::is_same::value && std::is_pointer::value) { + if constexpr (!std::is_same::value && + std::is_pointer::value) { addArg(detail::kernel_param_kind_t::kind_pointer, StoredArg, sizeof(T), ArgIndex); - } else if (ext::oneapi::experimental::detail::is_struct_with_special_type< - remove_cv_ref_t>::value) { + } else if constexpr (ext::oneapi::experimental::detail:: + is_struct_with_special_type< + remove_cv_ref_t>::value) { addArg(detail::kernel_param_kind_t::kind_struct_with_special_type, StoredArg, sizeof(T), ArgIndex); } else { +#ifndef __SYCL_DEVICE_ONLY__ + // Extra paranoid check to validated that we don't pass + // non-device-copyable types to the kernel with a standard layout. + // The device never executes this code and does not have access to + // is_struct_with_special_type so there is no point in trying to + // check it on the device. + static_assert( + detail::check_if_device_copyable_v> || + detail::deprecated_sycl121_device_copyable_v>); +#endif addArg(detail::kernel_param_kind_t::kind_std_layout, StoredArg, sizeof(T), ArgIndex); } @@ -1176,7 +1188,16 @@ class __SYCL_EXPORT handler { std::is_pointer_v>) // USM || is_same_type::value // Interop || is_same_type::value // Stream - || detail::check_if_device_copyable_v>; +#ifdef __SYCL_DEVICE_ONLY__ + // On the device we don't have the more precise + // is_struct_with_special_type + || detail::is_legal_kernel_parameter_v>; +#else + || (ext::oneapi::experimental::detail::is_struct_with_special_type< + remove_cv_ref_t>::value) || + detail::check_if_device_copyable_v> || + detail::deprecated_sycl121_device_copyable_v>; +#endif }; /// Sets argument for OpenCL interoperability kernels. diff --git a/sycl/test/basic_tests/accessor/accessor_special_members.cpp b/sycl/test/basic_tests/accessor/accessor_special_members.cpp new file mode 100644 index 0000000000000..a7dd57f322ca0 --- /dev/null +++ b/sycl/test/basic_tests/accessor/accessor_special_members.cpp @@ -0,0 +1,116 @@ +// RUN: %clangxx -fsycl -fsyntax-only %s +// RUN: %clangxx -fsycl -fsyntax-only -fpreview-breaking-changes %s + +#include "sycl/access/access.hpp" +#include +#include + +using namespace sycl; + +static_assert(std::is_nothrow_copy_constructible_v>, + "Copy constructor must be noexcept"); +static_assert(std::is_nothrow_copy_assignable_v>, + "Copy assignment must be noexcept"); +static_assert(std::is_nothrow_move_constructible_v>, + "Move constructor must be noexcept"); +static_assert(std::is_nothrow_move_assignable_v>, + "Move assignment must be noexcept"); + +static_assert(std::is_nothrow_copy_constructible_v>, + "Copy constructor must be noexcept"); +static_assert(std::is_nothrow_copy_assignable_v>, + "Copy assignment must be noexcept"); +static_assert(std::is_nothrow_move_constructible_v>, + "Move constructor must be noexcept"); +static_assert(std::is_nothrow_move_assignable_v>, + "Move assignment must be noexcept"); + +static_assert(std::is_nothrow_copy_constructible_v>, + "Copy constructor must be noexcept"); +static_assert(std::is_nothrow_copy_assignable_v>, + "Copy assignment must be noexcept"); +static_assert(std::is_nothrow_move_constructible_v>, + "Move constructor must be noexcept"); +static_assert(std::is_nothrow_move_assignable_v>, + "Move assignment must be noexcept"); + +static_assert( + std::is_nothrow_copy_constructible_v< + accessor>, + "Copy constructor must be noexcept"); +static_assert( + std::is_nothrow_copy_assignable_v< + accessor>, + "Copy assignment must be noexcept"); +static_assert( + std::is_nothrow_move_constructible_v< + accessor>, + "Move constructor must be noexcept"); +static_assert( + std::is_nothrow_move_assignable_v< + accessor>, + "Move assignment must be noexcept"); + +static_assert(std::is_nothrow_copy_constructible_v< + unsampled_image_accessor>, + "Copy constructor must be noexcept"); +static_assert(std::is_nothrow_copy_assignable_v< + unsampled_image_accessor>, + "Copy assignment must be noexcept"); +static_assert(std::is_nothrow_move_constructible_v< + unsampled_image_accessor>, + "Move constructor must be noexcept"); +static_assert(std::is_nothrow_move_assignable_v< + unsampled_image_accessor>, + "Move assignment must be noexcept"); + +static_assert( + std::is_nothrow_copy_constructible_v>, + "Copy constructor must be noexcept"); +static_assert( + std::is_nothrow_copy_assignable_v>, + "Copy assignment must be noexcept"); +static_assert( + std::is_nothrow_move_constructible_v>, + "Move constructor must be noexcept"); +static_assert( + std::is_nothrow_move_assignable_v>, + "Move assignment must be noexcept"); + +int main() { + queue Q; + buffer Buf{range<1>{1}}; + + Q.submit([&](handler &CGH) { + accessor Acc{Buf, CGH}; + using AccT = accessor; + static_assert(std::is_same_v); + + // Copy, move, and prvalue deduction must all yield the source type. + accessor AccCopy(Acc); + static_assert(std::is_same_v); + accessor AccMove(std::move(AccCopy)); + static_assert(std::is_same_v); + accessor AccPrvalue{accessor{Buf, CGH}}; + static_assert(std::is_same_v); + + local_accessor LocalAcc{range<1>{1}, CGH}; + local_accessor LocalAccCopy(LocalAcc); + static_assert(std::is_same_v); + local_accessor LocalAccMove(std::move(LocalAccCopy)); + static_assert(std::is_same_v); + + accessor AccProps{ + Buf, CGH, ext::oneapi::accessor_property_list{ext::oneapi::no_alias}}; + accessor AccPropsCopy(AccProps); + static_assert(std::is_same_v); + }); + + host_accessor HostAcc{Buf}; + host_accessor HostAccCopy(HostAcc); + static_assert(std::is_same_v); + host_accessor HostAccMove(std::move(HostAccCopy)); + static_assert(std::is_same_v); + + return 0; +} \ No newline at end of file diff --git a/sycl/test/basic_tests/device_copyable_checks_disabled.cpp b/sycl/test/basic_tests/device_copyable_checks_disabled.cpp index 908858207510d..76344e80cf001 100644 --- a/sycl/test/basic_tests/device_copyable_checks_disabled.cpp +++ b/sycl/test/basic_tests/device_copyable_checks_disabled.cpp @@ -1,5 +1,7 @@ // RUN: %clangxx -fsycl -fsycl-device-only -fsyntax-only -Xclang -verify=checks-on -Xclang -verify-ignore-unexpected=warning,note %s // RUN: %clangxx -fsycl -fsycl-device-only -fsyntax-only -DSYCL_DISABLE_DEVICE_COPYABLE_CHECKS -Xclang -verify=checks-off -Xclang -verify-ignore-unexpected=warning,note %s +// RUN: %clangxx -fsycl -fsycl-device-only -fsyntax-only -fpreview-breaking-changes -Xclang -verify=checks-on -Xclang -verify-ignore-unexpected=warning,note %s +// RUN: %clangxx -fsycl -fsycl-device-only -fsyntax-only -fpreview-breaking-changes -DSYCL_DISABLE_DEVICE_COPYABLE_CHECKS -Xclang -verify=checks-off -Xclang -verify-ignore-unexpected=warning,note %s // checks-off-no-diagnostics @@ -17,7 +19,7 @@ static_assert(!sycl::is_device_copyable_v); int main() { NotDeviceCopyable Val; - // checks-on-error@*:* {{The specified type is not device copyable}} + // checks-on-error@*:* {{The specified type is not a legal kernel parameter}} sycl::queue{}.single_task([=] { (void)Val; }); // checks-on-error@*:* {{a buffer must be device copyable}} diff --git a/sycl/test/basic_tests/implicit_device_copyable_types.cpp b/sycl/test/basic_tests/implicit_device_copyable_types.cpp index dd30dd47bc05c..ed6d0645b7e5f 100644 --- a/sycl/test/basic_tests/implicit_device_copyable_types.cpp +++ b/sycl/test/basic_tests/implicit_device_copyable_types.cpp @@ -1,5 +1,7 @@ // RUN: %clangxx -fsycl -fsyntax-only %s // RUN: %clangxx -D__SYCL_USE_LIBSYCL8_VEC_IMPL=1 -fsycl -fsyntax-only %s +// RUN: %clangxx -fsycl -fsyntax-only -fpreview-breaking-changes %s +// RUN: %clangxx -D__SYCL_USE_LIBSYCL8_VEC_IMPL=1 -fsycl -fsyntax-only -fpreview-breaking-changes %s #include #include @@ -97,6 +99,8 @@ int main() { #endif static_assert(sycl::is_device_copyable_v>); + static_assert(sycl::is_device_copyable_v>); + // Extra checks static_assert(sycl::is_device_copyable_v>); static_assert(sycl::is_device_copyable_v>); @@ -108,5 +112,43 @@ int main() { }; static_assert(sycl::is_device_copyable_v); + static_assert(sycl::is_device_copyable_v>); + static_assert(sycl::is_device_copyable_v>); + static_assert(sycl::is_device_copyable_v< + sycl::marray>); + + struct S2 { + sycl::marray v; + }; + static_assert(sycl::is_device_copyable_v); + +#ifdef __INTEL_PREVIEW_BREAKING_CHANGES + static_assert(!sycl::is_device_copyable_v< + sycl::accessor>); + static_assert(!std::is_trivially_copyable_v>); + static_assert( + !std::is_trivially_destructible_v>); + static_assert(!sycl::is_device_copyable_v>); + static_assert(!sycl::is_device_copyable_v>); + static_assert(!sycl::is_device_copyable_v>); + static_assert( + !sycl::is_device_copyable_v>); + static_assert(!sycl::is_device_copyable_v< + sycl::accessor>); + static_assert(!sycl::is_device_copyable_v>); + static_assert(!sycl::is_device_copyable_v< + sycl::sampled_image_accessor>); + static_assert(!sycl::is_device_copyable_v>); + static_assert(!sycl::is_device_copyable_v< + sycl::host_sampled_image_accessor>); +#endif // __INTEL_PREVIEW_BREAKING_CHANGES + + static_assert(sycl::is_device_copyable_v); + return 0; } diff --git a/sycl/test/basic_tests/is_device_copyable.cpp b/sycl/test/basic_tests/is_device_copyable.cpp index 1c4199e954530..4aad1b672f576 100644 --- a/sycl/test/basic_tests/is_device_copyable.cpp +++ b/sycl/test/basic_tests/is_device_copyable.cpp @@ -1,5 +1,7 @@ // RUN: %clangxx -fsycl -fsycl-device-only -fsyntax-only -Xclang -verify \ // RUN: %s -I %sycl_include +// RUN: %clangxx -fsycl -fsycl-device-only -fsyntax-only -fpreview-breaking-changes -Xclang -verify \ +// RUN: %s -I %sycl_include // This test checks that compiler does not report any errors // for the device copyable types being passed from HOST to DEVICE. @@ -39,6 +41,61 @@ struct DCopyable { ~DCopyable(); }; +// Legal but not device copyable +struct E { + sycl::accessor acc; + sycl::unsampled_image_accessor + uiacc; + sycl::sampled_image_accessor siacc; + int i; + float f; +}; +static_assert(sycl::detail::is_legal_kernel_parameter_v); +E SomeE(); + +struct F { + sycl::accessor acc[2]; + int i; +}; +static_assert(sycl::detail::is_legal_kernel_parameter_v); + +struct G { + std::array, 2> acc; + int i; +}; +static_assert(sycl::detail::is_legal_kernel_parameter_v); + +struct H { + const sycl::accessor acc; + int i; +}; +static_assert(sycl::detail::is_legal_kernel_parameter_v); + +struct Derived : sycl::accessor { + int i; +}; +static_assert(sycl::detail::is_legal_kernel_parameter_v); + +struct Nested { + H h; +}; +static_assert(sycl::detail::is_legal_kernel_parameter_v); + +struct S { + sycl::stream Stream; + int i; +}; +static_assert(sycl::detail::is_legal_kernel_parameter_v); + +struct KnownLimitation { + sycl::accessor acc; + int i; + KnownLimitation(); + KnownLimitation(const KnownLimitation &other); +}; +// Not actually legal, but included to demonstrate known limitation. +static_assert(sycl::detail::is_legal_kernel_parameter_v); + struct FunctorA { FunctorA() : b(7) {} void operator()() const { @@ -63,26 +120,54 @@ struct FunctorB : public DCopyable { template <> struct is_device_copyable : std::true_type {}; template <> struct is_device_copyable : std::true_type {}; +static_assert(sycl::detail::is_legal_kernel_parameter_v); +static_assert(sycl::detail::is_legal_kernel_parameter_v); + +sycl::unsampled_image_accessor +SomeUnsampledImageAccessor(); +sycl::sampled_image_accessor SomeSampledImageAccessor(); + void test() { A IamGood; IamGood.i = 0; BCopyable IamBadButCopyable(1); C IamAlsoGood; DCopyable IamAlsoBadButCopyable{0}; + E IamGoodButSpecial = SomeE(); + F IamGoodButSpecial2; + G IamGoodButSpecial3; + H IamGoodButSpecial4; + Derived IamDerived; + Nested IamNested; marray MarrayForCopyableIsCopyable(0); range<2> Range{1,2}; id<3> Id{1,2,3}; vec VecForCopyableIsCopyable{3}; + auto UIAcc = SomeUnsampledImageAccessor(); + auto SIAcc = SomeSampledImageAccessor(); + S IamStream; + KnownLimitation IamKnownLimitation; + queue Q; Q.single_task([=] { int A = IamGood.i; int B = IamBadButCopyable.i; int C = IamAlsoBadButCopyable.i; int D = IamAlsoGood.i; - int E = MarrayForCopyableIsCopyable[0]; - int F = Range[1]; - int G = Id[2]; - int H = VecForCopyableIsCopyable[0]; + int E = IamGoodButSpecial.i; + int F = IamGoodButSpecial2.i; + int G = IamGoodButSpecial3.i; + int H = IamGoodButSpecial4.i; + int I = MarrayForCopyableIsCopyable[0]; + int J = Range[1]; + int K = Id[2]; + int L = VecForCopyableIsCopyable[0]; + int M = IamDerived.i; + sycl::float4 N = UIAcc.read(0); + sycl::float4 O = SIAcc.read(0.f); + int P = IamNested.h.i; + int Q = IamStream.i; + int R = IamKnownLimitation.i; }); Q.single_task(FunctorA{}); diff --git a/sycl/test/basic_tests/is_device_copyable_neg.cpp b/sycl/test/basic_tests/is_device_copyable_neg.cpp index 61f88d6369aad..1968f5f81f82e 100644 --- a/sycl/test/basic_tests/is_device_copyable_neg.cpp +++ b/sycl/test/basic_tests/is_device_copyable_neg.cpp @@ -1,9 +1,12 @@ // RUN: %clangxx -fsycl -fsycl-device-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=warning,note %s +// RUN: %clangxx -fsycl -fsycl-device-only -fsyntax-only -fpreview-breaking-changes -Xclang -verify -Xclang -verify-ignore-unexpected=warning,note %s // This test checks if compiler reports compilation error on an attempt to pass // a struct with type that is not device copyable as SYCL kernel parameter. +#include #include +#include using namespace sycl; @@ -30,6 +33,53 @@ struct D { ~D(); }; +// Edge case with special type and invalid type +struct E { + sycl::accessor acc; + std::vector vec; +}; + +static_assert(!sycl::detail::check_if_device_copyable_v); +static_assert(!sycl::detail::is_legal_kernel_parameter_v); +static_assert( + sycl::detail::legal_kernel_parameters::contains_special_sycl_types_v); +static_assert( + !sycl::detail::legal_kernel_parameters::check_fields_and_bases_v); + +struct NotDC { + int a; + ~NotDC(); // Needed because of sycl 1.2.1 deprecated rules +}; + +template <> struct sycl::is_device_copyable : std::false_type {}; +static_assert(!sycl::detail::check_if_device_copyable_v); +static_assert(!sycl::detail::is_legal_kernel_parameter_v); + +struct F { + sycl::accessor acc; + NotDC notDC; +}; +static_assert(!sycl::detail::check_if_device_copyable_v); +static_assert(!sycl::detail::is_legal_kernel_parameter_v); + +struct G : sycl::accessor { + std::string str; +}; +static_assert(!sycl::detail::check_if_device_copyable_v); +static_assert(!sycl::detail::is_legal_kernel_parameter_v); + +struct WouldBeGood { + sycl::accessor acc; + int i; +}; + +struct H { + WouldBeGood wbg; + std::vector vec; +}; +static_assert(!sycl::detail::check_if_device_copyable_v); +static_assert(!sycl::detail::is_legal_kernel_parameter_v); + struct FunctorA { FunctorA() {} void operator()() const { @@ -54,20 +104,40 @@ struct FunctorB : public D { void test() { A IamBad(1); B IamAlsoBad{0}; + E IamEdgeCase; + F IamNotCopyable; + G IamDerivedFromAccessor; + H IamContainingWouldBeGood; marray MarrayForNotCopyable; queue Q; - // expected-error@*:* {{static assertion failed due to requirement 'is_deprecated_device_copyable_v': The specified type is not device copyable}} - Q.single_task([=] { - int A = IamBad.i; - int B = IamAlsoBad.i; - int MB = MarrayForNotCopyable[0].i; - }); + // expected-error@*:* {{static assertion failed due to requirement 'is_legal_kernel_parameter_v': The specified type is not a legal kernel parameter}} + Q.single_task([=] { int A = IamBad.i; }); + + // expected-error@*:* {{static assertion failed due to requirement 'is_legal_kernel_parameter_v': The specified type is not a legal kernel parameter}} + Q.single_task([=] { int B = IamAlsoBad.i; }); + + // expected-error@*:* {{static assertion failed due to requirement 'is_legal_kernel_parameter_v>': The specified type is not a legal kernel parameter}} + Q.single_task([=] { int MB = MarrayForNotCopyable[0].i; }); + + // expected-error@*:* {{static assertion failed due to requirement 'is_legal_kernel_parameter_v': The specified type is not a legal kernel parameter}} + Q.single_task([=] { int ECAcc = IamEdgeCase.acc[0]; }); + + // expected-error@*:* {{static assertion failed due to requirement 'is_legal_kernel_parameter_v': The specified type is not a legal kernel parameter}} + Q.single_task([=] { int FAcc = IamNotCopyable.acc[0]; }); + + // expected-error@*:* {{static assertion failed due to requirement 'is_legal_kernel_parameter_v': The specified type is not a legal kernel parameter}} + Q.single_task( + [=] { int GStrSize = IamDerivedFromAccessor.str.size(); }); + + // expected-error@*:* {{static assertion failed due to requirement 'is_legal_kernel_parameter_v': The specified type is not a legal kernel parameter}} + Q.single_task( + [=] { int HVecSize = IamContainingWouldBeGood.vec.size(); }); FunctorA FA; - // expected-error@*:* {{static assertion failed due to requirement 'is_deprecated_device_copyable_v': The specified type is not device copyable}} - Q.single_task(FA); + // expected-error@*:* {{static assertion failed due to requirement 'is_legal_kernel_parameter_v': The specified type is not a legal kernel parameter}} + Q.single_task(FA); FunctorB FB; - // expected-error@*:* {{static assertion failed due to requirement 'is_deprecated_device_copyable_v': The specified type is not device copyable}} - Q.single_task(FB); + // expected-error@*:* {{static assertion failed due to requirement 'is_legal_kernel_parameter_v': The specified type is not a legal kernel parameter}} + Q.single_task(FB); } diff --git a/sycl/test/basic_tests/set_arg_error.cpp b/sycl/test/basic_tests/set_arg_error.cpp index 5febdede1533b..56fcd2a467e82 100644 --- a/sycl/test/basic_tests/set_arg_error.cpp +++ b/sycl/test/basic_tests/set_arg_error.cpp @@ -1,5 +1,7 @@ // RUN: %clangxx %fsycl-host-only -DUSE_DEPRECATED_LOCAL_ACC -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s // RUN: %clangxx %fsycl-host-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s +// RUN: %clangxx %fsycl-host-only -DUSE_DEPRECATED_LOCAL_ACC -fsyntax-only -fpreview-breaking-changes -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s +// RUN: %clangxx %fsycl-host-only -fsyntax-only -fpreview-breaking-changes -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s #include @@ -22,12 +24,38 @@ struct NonStdLayout { int b; }; +// Trivially copy-constructible but not trivially copyable +struct Sycl121rules { + int a; + Sycl121rules() = default; + ~Sycl121rules() = default; + Sycl121rules(const Sycl121rules &) = default; + Sycl121rules &operator=(const Sycl121rules &); +}; + +struct AccWrap { + sycl::accessor acc; +}; + int main() { constexpr size_t size = 1; sycl::buffer buf(size); sycl::queue q; + sycl::int4 data; + sycl::unsampled_image<1> uimg(sycl::image_format::r32b32g32a32_sint, size); + sycl::sampled_image<1> simg( + &data, sycl::image_format::r32b32g32a32_sint, + {sycl::addressing_mode::none, + sycl::coordinate_normalization_mode::unnormalized, + sycl::filtering_mode::nearest}, + size); q.submit([&](sycl::handler &h) { auto global_acc = buf.get_access(h); + sycl::accessor prop_acc( + buf, h, + sycl::ext::oneapi::accessor_property_list{sycl::ext::oneapi::no_alias}); + auto uimg_acc = uimg.get_access(h); + auto simg_acc = simg.get_access(h); sycl::sampler samp(sycl::coordinate_normalization_mode::normalized, sycl::addressing_mode::clamp, sycl::filtering_mode::nearest); @@ -41,6 +69,8 @@ int main() { TriviallyCopyable tc{1, 2}; NonTriviallyCopyable ntc; + Sycl121rules sycl121; + AccWrap accWrap; h.set_arg(0, local_acc); h.set_arg(1, global_acc); h.set_arg(2, samp); @@ -49,13 +79,24 @@ int main() { h.set_arg( // expected-error {{no matching member function for call to 'set_arg'}} 5, ntc); h.set_arg( // expected-error {{no matching member function for call to 'set_arg'}} - 4, NonTriviallyCopyable{}); + 6, NonTriviallyCopyable{}); + h.set_arg( // expected-error {{no matching member function for call to 'set_arg'}} + 7, prop_acc); + h.set_arg( // expected-error {{no matching member function for call to 'set_arg'}} + 8, uimg_acc); + h.set_arg( // expected-error {{no matching member function for call to 'set_arg'}} + 9, simg_acc); #if SYCL_LANGUAGE_VERSION && SYCL_LANGUAGE_VERSION <= 201707 NonStdLayout nstd; h.set_arg( // expected-error {{no matching member function for call to 'set_arg'}} - 6, nstd); + 10, nstd); h.set_arg( // expected-error {{no matching member function for call to 'set_arg'}} - 7, NonStdLayout{}); + 11, NonStdLayout{}); #endif + h.set_arg(12, sycl121); + // No free-function kernel takes AccWrap, so the integration header doesn't + // mark it as a legal kernel parameter and the host-side set_arg rejects it. + h.set_arg( // expected-error {{no matching member function for call to 'set_arg'}} + 13, accWrap); }); } diff --git a/sycl/test/extensions/free_function_special_types.cpp b/sycl/test/extensions/free_function_special_types.cpp index 57363625455b1..6071aaea6f87a 100644 --- a/sycl/test/extensions/free_function_special_types.cpp +++ b/sycl/test/extensions/free_function_special_types.cpp @@ -1,4 +1,5 @@ // RUN: %clangxx -fsyntax-only -fsycl %s +// RUN: %clangxx -fsyntax-only -fsycl -fpreview-breaking-changes %s // Verify that we can pass top-level special type parameters to free function // kernels. @@ -7,10 +8,14 @@ using namespace sycl; +struct AccWrap { + accessor acc; +}; + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY( (ext::oneapi::experimental::nd_range_kernel<1>)) void foo(accessor acc, local_accessor lacc, sampler S, - ext::oneapi::experimental::annotated_ptr ptr) {} + ext::oneapi::experimental::annotated_ptr ptr, AccWrap aw) {} int main() { queue Q; @@ -26,7 +31,8 @@ int main() { sycl::filtering_mode::nearest); ext::oneapi::experimental::annotated_ptr ptr; sycl::stream str(8192, 1024, h); - h.set_args(acc, lacc, S, str, ptr); + AccWrap aw; + h.set_args(acc, lacc, S, str, ptr, aw); h.parallel_for(nd_range{{1}, {1}}, Kernel); }); return 0;