From 54550f9e03723e646bd5c9d21b544a5e2e57abf5 Mon Sep 17 00:00:00 2001 From: "Delisle, Thierry" Date: Wed, 23 Sep 2026 14:10:12 -0400 Subject: [PATCH 01/13] [SYCL] Consistent is_device_copyable for sycl::accessor Previously, is_device_copyable_v evaluated to different values for device and host code. Added explicit instantiation to consistently return false. Also added tests for other basic sycl types. Signed-off-by: Delisle, Thierry --- sycl/include/sycl/accessor.hpp | 16 ++++++++++++++++ .../implicit_device_copyable_types.cpp | 17 +++++++++++++++++ 2 files changed, 33 insertions(+) diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp index 0d8ea29488948..c2e917d42119f 100644 --- a/sycl/include/sycl/accessor.hpp +++ b/sycl/include/sycl/accessor.hpp @@ -18,6 +18,7 @@ #include // for buffer (fwd) #include // for is_genint, Try... #include // for associateWithH... +#include // for is_device_copyable #include // for loop #include // for OwnerLessBase #include // for PropWithDataKind @@ -2826,6 +2827,21 @@ host_accessor(buffer, Type1, Type2, Type3, Type4, Type5) -> host_accessor()>; +template +struct is_device_copyable> + : std::false_type {}; + +template +struct is_device_copyable> : std::false_type { +}; + +template +struct is_device_copyable> + : std::false_type {}; + } // namespace _V1 } // namespace sycl diff --git a/sycl/test/basic_tests/implicit_device_copyable_types.cpp b/sycl/test/basic_tests/implicit_device_copyable_types.cpp index dd30dd47bc05c..3a373b0cec169 100644 --- a/sycl/test/basic_tests/implicit_device_copyable_types.cpp +++ b/sycl/test/basic_tests/implicit_device_copyable_types.cpp @@ -108,5 +108,22 @@ 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); + + static_assert(!sycl::is_device_copyable_v< + sycl::accessor>); + static_assert(!sycl::is_device_copyable_v>); + static_assert(!sycl::is_device_copyable_v>); + + static_assert(sycl::is_device_copyable_v); + return 0; } From 4771df0f66f9c9d163b16c227dd369790fe80e5f Mon Sep 17 00:00:00 2001 From: "Delisle, Thierry" Date: Fri, 25 Sep 2026 09:45:14 -0400 Subject: [PATCH 02/13] [SYCL] added user-defined destructors to accessors Accessors are no longer trivially copyable to match the sycl spec more closely. Change is behind PREVIEW_BREAKING_CHANGES. Signed-off-by: Delisle, Thierry --- sycl/include/sycl/accessor.hpp | 30 +++++++++---------- .../implicit_device_copyable_types.cpp | 2 ++ 2 files changed, 17 insertions(+), 15 deletions(-) diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp index c2e917d42119f..372bb8cfeec80 100644 --- a/sycl/include/sycl/accessor.hpp +++ b/sycl/include/sycl/accessor.hpp @@ -18,7 +18,6 @@ #include // for buffer (fwd) #include // for is_genint, Try... #include // for associateWithH... -#include // for is_device_copyable #include // for loop #include // for OwnerLessBase #include // for PropWithDataKind @@ -758,6 +757,10 @@ 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() {} +#endif // __INTEL_PREVIEW_BREAKING_CHANGES + #else accessor(const detail::AccessorImplPtr &Impl) : detail::AccessorBaseHost{Impl} {} @@ -822,6 +825,10 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor : /*IsPlaceH=*/false, /*OffsetInBytes=*/0, /*IsSubBuffer=*/false, /*PropertyList=*/{}){}; +#ifdef __INTEL_PREVIEW_BREAKING_CHANGES + ~accessor() {} +#endif // __INTEL_PREVIEW_BREAKING_CHANGES + template friend class host_accessor; #endif // __SYCL_DEVICE_ONLY__ @@ -2417,6 +2424,10 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor range>::template get<0>(); } +#ifdef __INTEL_PREVIEW_BREAKING_CHANGES + ~local_accessor() {} +#endif // __INTEL_PREVIEW_BREAKING_CHANGES + #else local_accessor(const detail::AccessorImplPtr &Impl) : local_acc{Impl} {} #endif @@ -2583,6 +2594,9 @@ class __SYCL_EBO host_accessor public: host_accessor() : AccessorT() {} +#ifdef __INTEL_PREVIEW_BREAKING_CHANGES + ~host_accessor() {} +#endif // __INTEL_PREVIEW_BREAKING_CHANGES // The list of host_accessor constructors with their arguments // -------+---------+-------+----+----------+-------------- @@ -2827,20 +2841,6 @@ host_accessor(buffer, Type1, Type2, Type3, Type4, Type5) -> host_accessor()>; -template -struct is_device_copyable> - : std::false_type {}; - -template -struct is_device_copyable> : std::false_type { -}; - -template -struct is_device_copyable> - : std::false_type {}; } // namespace _V1 } // namespace sycl diff --git a/sycl/test/basic_tests/implicit_device_copyable_types.cpp b/sycl/test/basic_tests/implicit_device_copyable_types.cpp index 3a373b0cec169..a676791cafb4d 100644 --- a/sycl/test/basic_tests/implicit_device_copyable_types.cpp +++ b/sycl/test/basic_tests/implicit_device_copyable_types.cpp @@ -118,10 +118,12 @@ int main() { }; static_assert(sycl::is_device_copyable_v); +#ifdef __INTEL_PREVIEW_BREAKING_CHANGES static_assert(!sycl::is_device_copyable_v< sycl::accessor>); static_assert(!sycl::is_device_copyable_v>); static_assert(!sycl::is_device_copyable_v>); +#endif // __INTEL_PREVIEW_BREAKING_CHANGES static_assert(sycl::is_device_copyable_v); From 4bc0a5957500af70bbcda0a3d5377b6f0a186b9a Mon Sep 17 00:00:00 2001 From: "Delisle, Thierry" Date: Fri, 25 Sep 2026 09:53:10 -0400 Subject: [PATCH 03/13] Fixed formatting missed by git clang-format --- sycl/include/sycl/accessor.hpp | 1 - 1 file changed, 1 deletion(-) diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp index 372bb8cfeec80..8a55195743cd2 100644 --- a/sycl/include/sycl/accessor.hpp +++ b/sycl/include/sycl/accessor.hpp @@ -2841,7 +2841,6 @@ host_accessor(buffer, Type1, Type2, Type3, Type4, Type5) -> host_accessor()>; - } // namespace _V1 } // namespace sycl From 74cbd8749e94586e67625f7a8a99dd3223e87c0b Mon Sep 17 00:00:00 2001 From: "Delisle, Thierry" Date: Thu, 1 Oct 2026 09:26:13 -0400 Subject: [PATCH 04/13] added exception for accessor to is_device_copyable --- sycl/include/sycl/accessor.hpp | 21 +++++++++++++++++++ .../sycl/detail/is_device_copyable.hpp | 12 ++++++++++- 2 files changed, 32 insertions(+), 1 deletion(-) diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp index 8a55195743cd2..5c43af60b78cb 100644 --- a/sycl/include/sycl/accessor.hpp +++ b/sycl/include/sycl/accessor.hpp @@ -2841,6 +2841,27 @@ 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 + struct is_sycl_accessor> : std::true_type {}; + + template + struct is_sycl_accessor> : std::true_type {}; + + template + struct is_sycl_accessor> : std::true_type {}; +} + } // 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..c42fcc807332c 100644 --- a/sycl/include/sycl/detail/is_device_copyable.hpp +++ b/sycl/include/sycl/detail/is_device_copyable.hpp @@ -109,11 +109,21 @@ inline constexpr bool check_if_device_copyable_v = template struct CheckFieldsAreDeviceCopyable; template struct CheckBasesAreDeviceCopyable; +template +struct is_sycl_accessor : std::false_type {}; + +template +inline constexpr bool is_sycl_accessor_v = is_sycl_accessor::value; + 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); + std::is_trivially_destructible_v) || + // Per the spec, sycl::accessor and related types are not considered + // device copyable, but can be captured in a kernel. Therefore, we + // add an explicit exception here + is_sycl_accessor_v; template struct CheckFieldsAreDeviceCopyable> { From 3b875444b0db60a4f3874ad0f9a5a245a29dba48 Mon Sep 17 00:00:00 2001 From: "Delisle, Thierry" Date: Thu, 1 Oct 2026 09:38:09 -0400 Subject: [PATCH 05/13] formatting --- sycl/include/sycl/accessor.hpp | 33 +++++++++---------- .../sycl/detail/is_device_copyable.hpp | 5 ++- 2 files changed, 18 insertions(+), 20 deletions(-) diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp index 5c43af60b78cb..7a0266d80b403 100644 --- a/sycl/include/sycl/accessor.hpp +++ b/sycl/include/sycl/accessor.hpp @@ -2842,26 +2842,25 @@ host_accessor(buffer, Type1, Type2, Type3, Type4, detail::deduceAccessMode()>; 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 - struct is_sycl_accessor> : std::true_type {}; +// 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 - struct is_sycl_accessor> : std::true_type {}; +template +struct is_sycl_accessor> + : std::true_type {}; - template - struct is_sycl_accessor> : std::true_type {}; -} +template +struct is_sycl_accessor> : std::true_type {}; +template +struct is_sycl_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 c42fcc807332c..e62620a8365b1 100644 --- a/sycl/include/sycl/detail/is_device_copyable.hpp +++ b/sycl/include/sycl/detail/is_device_copyable.hpp @@ -109,8 +109,7 @@ inline constexpr bool check_if_device_copyable_v = template struct CheckFieldsAreDeviceCopyable; template struct CheckBasesAreDeviceCopyable; -template -struct is_sycl_accessor : std::false_type {}; +template struct is_sycl_accessor : std::false_type {}; template inline constexpr bool is_sycl_accessor_v = is_sycl_accessor::value; @@ -121,7 +120,7 @@ inline constexpr bool is_deprecated_device_copyable_v = (std::is_trivially_copy_constructible_v && std::is_trivially_destructible_v) || // Per the spec, sycl::accessor and related types are not considered - // device copyable, but can be captured in a kernel. Therefore, we + // device copyable, but can be captured in a kernel. Therefore, we // add an explicit exception here is_sycl_accessor_v; From 7478b51663f9a1cbc7f043148898dfd34279f11f Mon Sep 17 00:00:00 2001 From: "Delisle, Thierry" Date: Fri, 2 Oct 2026 11:15:40 -0400 Subject: [PATCH 06/13] [SYCL] Re-write CheckDeviceCopyable to match spec more closely. Previous definition of CheckDeviceCopyable was relying on is_device_copyable as a shortcut for checking kernel parameter legality. Renamed to CheckKernelParametersAreLegal and re-wrote implementation to match the sycl 2020 spec more closely. Signed-off-by: Delisle, Thierry --- .../sycl/detail/is_device_copyable.hpp | 137 +++++++++++++----- .../sycl/detail/kernel_launch_helper.hpp | 2 +- .../implicit_device_copyable_types.cpp | 2 + 3 files changed, 105 insertions(+), 36 deletions(-) diff --git a/sycl/include/sycl/detail/is_device_copyable.hpp b/sycl/include/sycl/detail/is_device_copyable.hpp index e62620a8365b1..df4ddfbb55c22 100644 --- a/sycl/include/sycl/detail/is_device_copyable.hpp +++ b/sycl/include/sycl/detail/is_device_copyable.hpp @@ -106,26 +106,19 @@ inline constexpr bool check_if_device_copyable_v = check_if_device_copyable::value; #ifdef __SYCL_DEVICE_ONLY__ -template struct CheckFieldsAreDeviceCopyable; -template struct CheckBasesAreDeviceCopyable; - -template struct is_sycl_accessor : std::false_type {}; - -template -inline constexpr bool is_sycl_accessor_v = is_sycl_accessor::value; +#ifndef __INTEL_PREVIEW_BREAKING_CHANGES +template struct CheckFieldsAreLegalKernelParameters; +template struct CheckBasesAreLegalKernelParameters; 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) || - // Per the spec, sycl::accessor and related types are not considered - // device copyable, but can be captured in a kernel. Therefore, we - // add an explicit exception here - is_sycl_accessor_v; + std::is_trivially_destructible_v); template -struct CheckFieldsAreDeviceCopyable> { +struct CheckFieldsAreLegalKernelParameters> { static_assert(((is_deprecated_device_copyable_v< decltype(__builtin_field_type(T, FieldIds))> && ...)), @@ -133,47 +126,121 @@ struct CheckFieldsAreDeviceCopyable> { }; template -struct CheckBasesAreDeviceCopyable> { +struct CheckBasesAreLegalKernelParameters> { static_assert(((is_deprecated_device_copyable_v< decltype(__builtin_base_type(T, BaseIds))> && ...)), "The specified type is not device copyable"); }; +#else +template struct is_sycl_accessor : std::false_type {}; + +template +inline constexpr bool is_sycl_accessor_v = is_sycl_accessor::value; + +template struct IsLegalKernelParameterBase { + static constexpr bool value = + check_if_device_copyable_v || is_sycl_accessor_v; +}; + +template +struct IsLegalKernelParameter : IsLegalKernelParameterBase {}; + +template struct AreFieldsLegalKernelParameters; +template struct AreBasesLegalKernelParameters; + +template +struct AreFieldsLegalKernelParameters> + : std::bool_constant<(IsLegalKernelParameter::value && + ...)> {}; + +template +struct AreBasesLegalKernelParameters> + : std::bool_constant<(IsLegalKernelParameter::value && + ...)> {}; + +template +struct IsLegalKernelParameter< + T, std::void_t> + : std::bool_constant< + IsLegalKernelParameterBase::value || + (AreFieldsLegalKernelParameters< + T, std::make_index_sequence<__builtin_num_fields(T)>>::value && + AreBasesLegalKernelParameters< + T, std::make_index_sequence<__builtin_num_bases(T)>>::value)> {}; + +template +struct IsLegalKernelParameter : IsLegalKernelParameter {}; +#endif // __INTEL_PREVIEW_BREAKING_CHANGES -// 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. template -struct CheckDeviceCopyable - : CheckFieldsAreDeviceCopyable< +struct CheckKernelParametersAreLegal +#ifdef __INTEL_PREVIEW_BREAKING_CHANGES +{ + // All the captures of a lambda or functor of type FuncT passed to a kernel + // must be legal kernel parameters 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. + static_assert(IsLegalKernelParameter::value, + "The specified type is not a legal device parameter"); +}; +#else + // 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. + : CheckFieldsAreLegalKernelParameters< FuncT, std::make_index_sequence<__builtin_num_fields(FuncT)>>, - CheckBasesAreDeviceCopyable< - FuncT, std::make_index_sequence<__builtin_num_bases(FuncT)>> {}; + CheckBasesAreLegalKernelParameters< + FuncT, std::make_index_sequence<__builtin_num_bases(FuncT)>> { +}; +#endif // __INTEL_PREVIEW_BREAKING_CHANGES template class RoundedRangeKernel; template class RoundedRangeKernelWithKH; -// Below are two specializations for CheckDeviceCopyable when a kernel lambda -// is wrapped after range rounding optimization. +// Below are two specializations for CheckKernelParametersAreLegal when a kernel +// lambda is wrapped after range rounding optimization. template -struct CheckDeviceCopyable< +struct CheckKernelParametersAreLegal< RoundedRangeKernel> - : CheckDeviceCopyable {}; + : CheckKernelParametersAreLegal {}; template -struct CheckDeviceCopyable< +struct CheckKernelParametersAreLegal< RoundedRangeKernelWithKH> - : CheckDeviceCopyable {}; + : CheckKernelParametersAreLegal {}; #endif // __SYCL_DEVICE_ONLY__ } // namespace detail diff --git a/sycl/include/sycl/detail/kernel_launch_helper.hpp b/sycl/include/sycl/detail/kernel_launch_helper.hpp index a535b741ff830..ecb7d978bf475 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::CheckKernelParametersAreLegal(); #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/test/basic_tests/implicit_device_copyable_types.cpp b/sycl/test/basic_tests/implicit_device_copyable_types.cpp index a676791cafb4d..23bdafbce32c2 100644 --- a/sycl/test/basic_tests/implicit_device_copyable_types.cpp +++ b/sycl/test/basic_tests/implicit_device_copyable_types.cpp @@ -97,6 +97,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>); From 8d933d1444e4b51f35e9328cb43f767f28d85a99 Mon Sep 17 00:00:00 2001 From: "Delisle, Thierry" Date: Mon, 5 Oct 2026 14:55:32 -0400 Subject: [PATCH 07/13] [SYCL] update handling of sycl special types to finalize destructor Since sycl::accessor is not trivially copyable, special types require the normal destructor machinery. Signed-off-by: Delisle, Thierry --- clang/lib/Sema/SemaSYCL.cpp | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 11d2e5e001629..98f336a38d58b 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()) From faecfa9ed5d2e9b26215d55b5afe4ffb3307f7ab Mon Sep 17 00:00:00 2001 From: "Delisle, Thierry" Date: Tue, 6 Oct 2026 17:31:53 -0400 Subject: [PATCH 08/13] [SYCL] Do not emit destructors for device work groups. Signed-off-by: Delisle, Thierry --- clang/lib/CodeGen/CGDeclCXX.cpp | 14 ++++++++++++++ 1 file changed, 14 insertions(+) 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 From 37e133a3e5be3d61a808789c965ee77a51a57519 Mon Sep 17 00:00:00 2001 From: "Delisle, Thierry" Date: Wed, 7 Oct 2026 09:50:00 -0400 Subject: [PATCH 09/13] [SYCL] Add local accessor dtor to verified esimd functions. Signed-off-by: Delisle, Thierry --- llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp b/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp index 7e1712a4271f8..04578603658f4 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp @@ -39,6 +39,7 @@ 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", From 47f79326a3218501261b83a639c763c1454b2414 Mon Sep 17 00:00:00 2001 From: "Delisle, Thierry" Date: Wed, 7 Oct 2026 15:10:15 -0400 Subject: [PATCH 10/13] [SYCL] Add explicit copy and move ctor declarations. Signed-off-by: Delisle, Thierry --- sycl/include/sycl/accessor.hpp | 35 +++++++++++++++++++++++++++++++++- 1 file changed, 34 insertions(+), 1 deletion(-) diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp index 7a0266d80b403..f6eaad4c3e820 100644 --- a/sycl/include/sycl/accessor.hpp +++ b/sycl/include/sycl/accessor.hpp @@ -896,6 +896,14 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor : // -------+---------+-------+----+-----+-------------- public: +#if __INTEL_PREVIEW_BREAKING_CHANGES + // common reference semantics + accessor(const accessor &) noexcept = default; + accessor &operator=(const accessor &) noexcept = default; + accessor(accessor &&) noexcept = default; + accessor &operator=(accessor &&) noexcept = default; +#endif // __INTEL_PREVIEW_BREAKING_CHANGES + // implicit conversion between const / non-const types for read only accessors template + template < + typename... NewPropsT +#ifdef __INTEL_PREVIEW_BREAKING_CHANGES + , + // Same-type copies are handled by the copy constructor. Excluding them + // here also keeps CTAD from a prvalue accessor unambiguous. + typename = std::enable_if_t, + std::remove_cv_t>>> +#endif // __INTEL_PREVIEW_BREAKING_CHANGES + > accessor( const accessor> &Other, @@ -2432,6 +2450,15 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor local_accessor(const detail::AccessorImplPtr &Impl) : local_acc{Impl} {} #endif +public: +#ifdef __INTEL_PREVIEW_BREAKING_CHANGES + // 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; +#endif // __INTEL_PREVIEW_BREAKING_CHANGES + // implicit conversion between non-const read-write accessor to const // read-only accessor public: @@ -2596,6 +2623,12 @@ class __SYCL_EBO host_accessor host_accessor() : AccessorT() {} #ifdef __INTEL_PREVIEW_BREAKING_CHANGES ~host_accessor() {} + + // 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; #endif // __INTEL_PREVIEW_BREAKING_CHANGES // The list of host_accessor constructors with their arguments From e966767c15f49f4f84f2a39e8a37d6358db5425e Mon Sep 17 00:00:00 2001 From: "Delisle, Thierry" Date: Thu, 8 Oct 2026 10:15:12 -0400 Subject: [PATCH 11/13] [SYCL] Relax deduction rules for sycl::handler::set_arg. Signed-off-by: Delisle, Thierry --- sycl/include/sycl/accessor.hpp | 16 ++++++++++++---- sycl/include/sycl/handler.hpp | 12 ++++++++++++ 2 files changed, 24 insertions(+), 4 deletions(-) diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp index f6eaad4c3e820..27fbed45b310f 100644 --- a/sycl/include/sycl/accessor.hpp +++ b/sycl/include/sycl/accessor.hpp @@ -758,7 +758,9 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor : detail::InitializedVal::template get<0>()) {} #ifdef __INTEL_PREVIEW_BREAKING_CHANGES - ~accessor() {} + ~accessor() { + // Per the spec, the destructor of an accessor must not have any effect + } #endif // __INTEL_PREVIEW_BREAKING_CHANGES #else @@ -826,7 +828,9 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor : /*OffsetInBytes=*/0, /*IsSubBuffer=*/false, /*PropertyList=*/{}){}; #ifdef __INTEL_PREVIEW_BREAKING_CHANGES - ~accessor() {} + ~accessor() { + // Per the spec, the destructor of an accessor must not have any effect + } #endif // __INTEL_PREVIEW_BREAKING_CHANGES template friend class host_accessor; @@ -2443,7 +2447,9 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor } #ifdef __INTEL_PREVIEW_BREAKING_CHANGES - ~local_accessor() {} + ~local_accessor() { + // Per the spec, the destructor of an accessor must not have any effect + } #endif // __INTEL_PREVIEW_BREAKING_CHANGES #else @@ -2622,7 +2628,9 @@ class __SYCL_EBO host_accessor public: host_accessor() : AccessorT() {} #ifdef __INTEL_PREVIEW_BREAKING_CHANGES - ~host_accessor() {} + ~host_accessor() { + // Per the spec, the destructor of an accessor must not have any effect + } // common reference semantics host_accessor(const host_accessor &) noexcept = default; diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 56539906b52d3..1549ea5c40097 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -1176,7 +1176,19 @@ class __SYCL_EXPORT handler { std::is_pointer_v>) // USM || is_same_type::value // Interop || is_same_type::value // Stream +#ifndef __INTEL_PREVIEW_BREAKING_CHANGES || detail::check_if_device_copyable_v>; +#else + // Is device copyable is too narrow a definition because not all legal + // sycl types are considered device copyable, and structures that + // contain such types are still legal SYCL types. Once C++26 is + // available, we might have a more precise way to determine parameter + // legality. + || (!std::is_abstract_v> && + !std::is_polymorphic_v> && + std::is_destructible_v> && + std::is_copy_constructible_v>); +#endif // __INTEL_PREVIEW_BREAKING_CHANGES }; /// Sets argument for OpenCL interoperability kernels. From f011bd957fa9c871c238ecfdc67327adac4f4266 Mon Sep 17 00:00:00 2001 From: "Delisle, Thierry" Date: Tue, 13 Oct 2026 15:36:31 -0400 Subject: [PATCH 12/13] [SYCL] Rework is_legal_kernel_parameter to be much stricter Updated is_legal_kernel_parameter to no longer accept any types composed of legal types. Changes to is device copyable extended to accessors. Added tests to cover new changes. Signed-off-by: Delisle, Thierry --- clang/lib/Sema/SemaSYCL.cpp | 6 +- .../CodeGenSYCL/free_function_int_header.cpp | 10 +- llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp | 1 + sycl/include/sycl/accessor.hpp | 70 +++-- sycl/include/sycl/accessor_image.hpp | 99 ++++++- .../sycl/detail/is_device_copyable.hpp | 277 +++++++++++------- .../sycl/detail/kernel_launch_helper.hpp | 2 +- .../experimental/free_function_traits.hpp | 4 +- sycl/include/sycl/handler.hpp | 14 +- .../accessor/accessor_special_members.cpp | 116 ++++++++ .../device_copyable_checks_disabled.cpp | 4 +- .../implicit_device_copyable_types.cpp | 21 ++ sycl/test/basic_tests/is_device_copyable.cpp | 93 +++++- .../basic_tests/is_device_copyable_neg.cpp | 90 +++++- sycl/test/basic_tests/set_arg_error.cpp | 22 ++ .../free_function_special_types.cpp | 10 +- 16 files changed, 666 insertions(+), 173 deletions(-) create mode 100644 sycl/test/basic_tests/accessor/accessor_special_members.cpp diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 98f336a38d58b..b8584adb7ca83 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -7680,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 04578603658f4..4630c400e6b25 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp @@ -44,6 +44,7 @@ static const char *LegalSYCLFunctions[] = { "^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 27fbed45b310f..09b4c6d409576 100644 --- a/sycl/include/sycl/accessor.hpp +++ b/sycl/include/sycl/accessor.hpp @@ -759,7 +759,8 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor : #ifdef __INTEL_PREVIEW_BREAKING_CHANGES ~accessor() { - // Per the spec, the destructor of an accessor must not have any effect + // 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 @@ -829,7 +830,9 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor : #ifdef __INTEL_PREVIEW_BREAKING_CHANGES ~accessor() { - // Per the spec, the destructor of an accessor must not have any effect + // 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 @@ -900,13 +903,11 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor : // -------+---------+-------+----+-----+-------------- public: -#if __INTEL_PREVIEW_BREAKING_CHANGES // common reference semantics accessor(const accessor &) noexcept = default; accessor &operator=(const accessor &) noexcept = default; accessor(accessor &&) noexcept = default; accessor &operator=(accessor &&) noexcept = default; -#endif // __INTEL_PREVIEW_BREAKING_CHANGES // implicit conversion between const / non-const types for read only accessors template () #ifdef __INTEL_PREVIEW_BREAKING_CHANGES - , - // Same-type copies are handled by the copy constructor. Excluding them - // here also keeps CTAD from a prvalue accessor unambiguous. - typename = std::enable_if_t, - std::remove_cv_t>>> + // 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, @@ -1635,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, @@ -2271,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") @@ -2446,24 +2459,16 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor range>::template get<0>(); } -#ifdef __INTEL_PREVIEW_BREAKING_CHANGES - ~local_accessor() { - // Per the spec, the destructor of an accessor must not have any effect - } -#endif // __INTEL_PREVIEW_BREAKING_CHANGES - #else local_accessor(const detail::AccessorImplPtr &Impl) : local_acc{Impl} {} #endif public: -#ifdef __INTEL_PREVIEW_BREAKING_CHANGES // 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; -#endif // __INTEL_PREVIEW_BREAKING_CHANGES // implicit conversion between non-const read-write accessor to const // read-only accessor @@ -2629,15 +2634,15 @@ class __SYCL_EBO host_accessor host_accessor() : AccessorT() {} #ifdef __INTEL_PREVIEW_BREAKING_CHANGES ~host_accessor() { - // Per the spec, the destructor of an accessor must not have any effect + // 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; -#endif // __INTEL_PREVIEW_BREAKING_CHANGES // The list of host_accessor constructors with their arguments // -------+---------+-------+----+----------+-------------- @@ -2888,19 +2893,24 @@ namespace detail { // 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::true_type {}; + : std::bool_constant> {}; template struct is_sycl_accessor> : std::true_type {}; - -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 df4ddfbb55c22..f4aaa28fc4754 100644 --- a/sycl/include/sycl/detail/is_device_copyable.hpp +++ b/sycl/include/sycl/detail/is_device_copyable.hpp @@ -105,144 +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__ -#ifndef __INTEL_PREVIEW_BREAKING_CHANGES -template struct CheckFieldsAreLegalKernelParameters; -template struct CheckBasesAreLegalKernelParameters; +// 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 is_deprecated_device_copyable_v = - check_if_device_copyable_v || - (std::is_trivially_copy_constructible_v && - std::is_trivially_destructible_v); +inline constexpr bool contains_special_sycl_types_v = + contains_special_sycl_types::value; +template struct special_sycl_types_fields; template -struct CheckFieldsAreLegalKernelParameters> { - static_assert(((is_deprecated_device_copyable_v< - decltype(__builtin_field_type(T, FieldIds))> && - ...)), - "The specified type is not device copyable"); -}; +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 CheckBasesAreLegalKernelParameters> { - static_assert(((is_deprecated_device_copyable_v< - decltype(__builtin_base_type(T, BaseIds))> && - ...)), - "The specified type is not device copyable"); -}; -#else -template struct is_sycl_accessor : std::false_type {}; +struct special_sycl_types_bases> + : std::bool_constant<(contains_special_sycl_types_v< + decltype(__builtin_base_type(T, BaseIds))> || + ...)> {}; -template -inline constexpr bool is_sycl_accessor_v = is_sycl_accessor::value; +template +inline constexpr bool special_sycl_types_bases_v = + special_sycl_types_bases::value; -template struct IsLegalKernelParameterBase { - static constexpr bool value = - check_if_device_copyable_v || is_sycl_accessor_v; -}; +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 IsLegalKernelParameter : IsLegalKernelParameterBase {}; +struct check_fields_and_bases : std::false_type {}; -template struct AreFieldsLegalKernelParameters; -template struct AreBasesLegalKernelParameters; +template +inline constexpr bool check_fields_and_bases_v = + check_fields_and_bases::value; +template struct check_fields; template -struct AreFieldsLegalKernelParameters> - : std::bool_constant<(IsLegalKernelParameter::value && +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 AreBasesLegalKernelParameters> - : std::bool_constant<(IsLegalKernelParameter::value && +struct check_bases> + : std::bool_constant<((is_legal_kernel_parameter_v>) && ...)> {}; +template +inline constexpr bool check_bases_v = check_bases::value; + template -struct IsLegalKernelParameter< +struct check_fields_and_bases< T, std::void_t> : std::bool_constant< - IsLegalKernelParameterBase::value || - (AreFieldsLegalKernelParameters< - T, std::make_index_sequence<__builtin_num_fields(T)>>::value && - AreBasesLegalKernelParameters< - T, std::make_index_sequence<__builtin_num_bases(T)>>::value)> {}; + check_fields_v> && + check_bases_v>> { +}; template -struct IsLegalKernelParameter : IsLegalKernelParameter {}; -#endif // __INTEL_PREVIEW_BREAKING_CHANGES +struct check_fields_and_bases : check_fields_and_bases {}; -template -struct CheckKernelParametersAreLegal -#ifdef __INTEL_PREVIEW_BREAKING_CHANGES -{ - // All the captures of a lambda or functor of type FuncT passed to a kernel - // must be legal kernel parameters 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. - static_assert(IsLegalKernelParameter::value, - "The specified type is not a legal device parameter"); +template struct validate_fields; +template +struct validate_fields> { + static_assert( + ((is_legal_kernel_parameter_v< + std::remove_cv_t> && + ...)), + "The specified type is not a legal kernel parameter"); }; -#else - // 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. - : CheckFieldsAreLegalKernelParameters< - FuncT, std::make_index_sequence<__builtin_num_fields(FuncT)>>, - CheckBasesAreLegalKernelParameters< - FuncT, std::make_index_sequence<__builtin_num_bases(FuncT)>> { + +// 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 validate_bases> { + static_assert( + ((is_legal_kernel_parameter_v< + std::remove_cv_t> && + ...)), + "The specified type is not a legal kernel parameter"); }; -#endif // __INTEL_PREVIEW_BREAKING_CHANGES -template -class RoundedRangeKernel; -template -class RoundedRangeKernelWithKH; +// All the captures of a lambda or functor of type FuncT passed to a kernel +// must be legal kernel parameters. +template +struct validate_functor + : validate_fields>, + validate_bases> {}; -// Below are two specializations for CheckKernelParametersAreLegal when a kernel +// Below are two specializations for validate_functor when a kernel // lambda is wrapped after range rounding optimization. template -struct CheckKernelParametersAreLegal< +struct validate_functor< RoundedRangeKernel> - : CheckKernelParametersAreLegal {}; + : validate_functor {}; template -struct CheckKernelParametersAreLegal< +struct validate_functor< RoundedRangeKernelWithKH> - : CheckKernelParametersAreLegal {}; + : 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 ecb7d978bf475..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::CheckKernelParametersAreLegal(); + 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 1549ea5c40097..64a5479f30ec1 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -1176,19 +1176,7 @@ class __SYCL_EXPORT handler { std::is_pointer_v>) // USM || is_same_type::value // Interop || is_same_type::value // Stream -#ifndef __INTEL_PREVIEW_BREAKING_CHANGES - || detail::check_if_device_copyable_v>; -#else - // Is device copyable is too narrow a definition because not all legal - // sycl types are considered device copyable, and structures that - // contain such types are still legal SYCL types. Once C++26 is - // available, we might have a more precise way to determine parameter - // legality. - || (!std::is_abstract_v> && - !std::is_polymorphic_v> && - std::is_destructible_v> && - std::is_copy_constructible_v>); -#endif // __INTEL_PREVIEW_BREAKING_CHANGES + || detail::is_legal_kernel_parameter_v>; }; /// 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 23bdafbce32c2..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 @@ -123,8 +125,27 @@ int main() { #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); 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..52f0715d2d3cb 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,6 +24,19 @@ 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); @@ -41,6 +56,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); @@ -57,5 +74,10 @@ int main() { h.set_arg( // expected-error {{no matching member function for call to 'set_arg'}} 7, NonStdLayout{}); #endif + h.set_arg(8, 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'}} + 9, 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; From 0adb720367754e95aa0e00aaee4357cd74a83915 Mon Sep 17 00:00:00 2001 From: "Delisle, Thierry" Date: Mon, 19 Oct 2026 16:28:15 -0400 Subject: [PATCH 13/13] Address review comments. Signed-off-by: Delisle, Thierry --- sycl/include/sycl/handler.hpp | 27 ++++++++++++++++++++--- sycl/test/basic_tests/set_arg_error.cpp | 29 ++++++++++++++++++++----- 2 files changed, 48 insertions(+), 8 deletions(-) diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 64a5479f30ec1..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 +#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/set_arg_error.cpp b/sycl/test/basic_tests/set_arg_error.cpp index 52f0715d2d3cb..56fcd2a467e82 100644 --- a/sycl/test/basic_tests/set_arg_error.cpp +++ b/sycl/test/basic_tests/set_arg_error.cpp @@ -41,8 +41,21 @@ 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); @@ -66,18 +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(8, sycl121); + 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'}} - 9, accWrap); + 13, accWrap); }); }