[SYCL] Consistent is_device_copyable for sycl::accessor

[inteltdelisle]

sycl::is_device_copyable_v (and std::is_trivially_copyable_v) evaluated differently in host and device compilation, because the accessor special members are implemented differently under SYCL_DEVICE_ONLY.

For example, this compiled for the device but not for the host:

    static_assert(sycl::is_device_copyable_v<sycl::accessor<int, 1, sycl::access_mode::read_write>>);
    static_assert(sycl::is_device_copyable_v<sycl::local_accessor<int, 1>>);

Traits that differ between compilation passes can make constexpr/SFINAE results diverge between host and device.

Fix (behind -fpreview-breaking-changes, since it changes trivial copyability and is ABI-breaking):

Accessor types (accessor, local_accessor/local_accessor_base, host_accessor, and the image accessors) get a user-provided empty destructor, so they are consistently not trivially copyable and, per SYCL 2020, not implicitly device copyable. The copy/move constructors and assignment operators are explicitly defaulted to avoid issues with implicit functions.

Kernel parameter legality is now separate from device copyability (Not ABI-breaking):

New sycl::detail::is_legal_kernel_parameter trait that follows the SYCL 2020 "legal kernel parameter" rules: device copyable types, device-target accessors, arrays of legal types, and structs whose fields and bases are all legal and that contain a special SYCL type. Host-target accessors are rejected.

CheckDeviceCopyable is replaced by legal_kernel_parameters::validate_functor, which checks kernel captures against this trait. handler::set_arg has also been updated for consistency, reusing the trait were possible.
The integration header now specializes detail::is_legal_kernel_parameter instead of is_device_copyable for free-function structs that contain special types.

Compiler changes needed because accessors now have a non-trivial destructor:

SemaSYCL: finalize the destructor of special-object clones (FinalizeVarWithDestructor). Without this change, no definition is created for the dtor, since this is an inlined template, and calls to the destructor may fail to link.

CGDeclCXX: don't register a destructor for SYCL work-group-scope variables (e.g. the parallel_for_work_group kernel object). SYCL requires these destructors to have no effect and platforms are not required to declared/implement __cxx_atexit.

ESIMDVerifier: allow the local_accessor and local_accessor_base destructors.

Other: the accessor property-list conversion constructor is now constrained with enable_if instead of a static_assert, so CTAD and same-type copies pick the copy constructor.

Known limitation:

A struct that contains only legal members but has an illegal user-defined copy constructor or destructor is not detected.

Follow-up on the work-group scope attribute handling: #23461.

Tests:

New/updated tests for accessor special members, implicitly device copyable SYCL types, is_device_copyable (positive and negative), and set_arg. Each has RUN lines with and without -fpreview-breaking-changes.

[koparasy]

A high level comment. What makes accessors change their behavior depending on host|device compilation?

[inteltdelisle]

A high level comment. What makes accessors change their behavior depending on host|device compilation?

By default is_device_copyable simply falls back on std::is_trivially_copyable, which in parts look at what the ctors are doing. sycl::accessor ctors are implemented differently depending on whether ___SYCL_DEVICE_ONLY__ is defined and that is what was causing the difference.

[slawekptak]

@inteltdelisle thanks for the fix. Could you please describe a bit of a background for this change - was some specific test failing? Thanks!
Adding @gmlueck

[inteltdelisle]

This fix was of result of a bug report submitted by @koparasy, who noted that the following code was inconsistent between device and host compilation:

#include <sycl/sycl.hpp>

using AccT = sycl::accessor<int, 1, sycl::access_mode::read_write>;
using LAccT = sycl::local_accessor<int, 1>;

static_assert(
   sycl::is_device_copyable_v<AccT>,
   "DEVICE pass: accessor IS device-copyable (trivially copyable here)");
static_assert(sycl::is_device_copyable_v<LAccT>,
             "DEVICE pass: local_accessor IS device-copyable here");
void f() {}

This code compiled for the device but not the host. The additional tests are simply there to verify the behavior of other basic sycl types.

Worth noting that I did not add a test to verify the behavior of is_device_copyable<buffer<...>>, I am not sure what the desired behavior is for that type.

[gmlueck]

I wonder if we should change the implementation of the accessor instead of specializing is_device_copyable. For example, adding an empty user-provided destructor will cause the types to be non-trivially copyable.

This seems better than specializing the trait because it also makes is_trivially_copyable consistent between host and device.

[inteltdelisle]

I'll confirm locally whether that works. I was hesitant do change the accessor type itself, since that will have a less targeted impact. On the other hand the spec does state "Any type that is trivially copyable (as defined by the C++ core language) is implicitly device copyable." so that solution match the spec more closely.

[inteltdelisle]

@gmlueck Adding the dtors instead works locally. I can push that change if it's a preferred solution.

[gmlueck]

Spec-wise, I think it's better. Maybe @intel/llvm-reviewers-runtime can weigh in from the implementation side.

[uditagarwal97]

@inteltdelisle Changing a type's trivial-copyability is an ABI breaking change. I am on board with @gmlueck's suggestion, but I think we should add the user-provided destructor under PREVIEW_BREAKING_CHANGES flag.

[koparasy]

@gmlueck Does SYCL require std::is_trivially_copyable_v<accessor> itself to evaluate consistently with sycl::is_device_copyable_v<accessor>? If not, wouldn't an explicit specialization of is_device_copyable be preferable to deliberately making the destructor non-trivial, since the latter changes additional C++ type properties unrelated to device copyability?

I am also wondering whether we actually want to break code that may rely on std::is_trivially_copyable_v<accessor>, even if the current behavior is somewhat obscure and differs between host and device compilation.

[gmlueck]

In general, I think it's bad for traits to evaluate differently in the host vs. the device compilation passes because this can lead to constexpr expressions having different values in the two passes. I think this could lead to quite confusing behavior (e.g. SFINAE constraints evaluating different in the two passes). Therefore, I think std::is_trivially_copyable_v<accessor> should evaluate the same in both passes.

TBH, I think the SYCL spec should be changed to guarantee that all standard traits and all SYCL traits evaluate the same in all compiler passes. This is a battle for another day, though.

I agree that changing the trait in the device compiler could potentially break an application (thought I think it's pretty unlikely). Therefore, I agree that it's better to make this change at the next major release.

[inteltdelisle]

@gmlueck I'm not super familiar with the release procedure yet, what does making this change at the next major release entail precisely? Do we merge this change in behind the PREVIEW_BREAKING_CHANGES?

[gmlueck]

Yes, that's right. You would merge the change as it is now in this PR behind the PREVIEW_BREAKING_CHANGES macro. At the next major release, we search for all these macros and enable the code they protect by default. Merging the changes now under the macro provides two benefits:

• It allows user to see what potential breaking changes might be coming in the next major release by compiling with -fpreview-breaking-changes.
• It lets us address issues like this early, rather than waiting until right before the next major release.

[koparasy]

Thnx!

[Copilot]

🟡 Changes recommended

Kernel legality predicates currently admit invalid host accessors, host-target accessors, and arbitrary non-device-copyable arguments.

4 open findings

• Reject host-target accessors as device kernel parameters · New
• Reject host_accessor as an illegal kernel parameter · New
• Restrict set_arg fallback to device-copyable kernel arguments · New
• Enable preview-mode RUN coverage for accessor trait assertions · New

What changed in this PR

Makes accessor device-copyability consistent while preserving legal kernel captures.

Changes:

• Adds non-trivial accessor destructors and explicit copy/move operations.
• Reworks kernel-parameter legality checks.
• Updates compiler lowering, ESIMD verification, and type-trait tests.

File	Description
sycl/​test/​basic_tests/​implicit_device_copyable_types.cpp	Adds device-copyability assertions.
sycl/​include/​sycl/​handler.hpp	Broadens set_arg eligibility.
sycl/​include/​sycl/​detail/​kernel_launch_helper.hpp	Invokes the renamed legality check.
sycl/​include/​sycl/​detail/​is_device_copyable.hpp	Implements recursive kernel-parameter legality checks.
sycl/​include/​sycl/​accessor.hpp	Adds accessor destructors, special members, and legality traits.
llvm/​lib/​SYCLLowerIR/​ESIMD/​ESIMDVerifier.cpp	Allows local-accessor destructors in ESIMD.
clang/​lib/​Sema/​SemaSYCL.cpp	Finalizes generated variables with destructors.
clang/​lib/​CodeGen/​CGDeclCXX.cpp	Skips work-group-scope destructor registration.

🧠 Review effort: Balanced

---

Give feedback about Copilot approvals in this survey to enter a drawing for a $150 gift card.

[elizabethandrews]

Please add a FE test.

[elizabethandrews]

We have the scope attribute lookup in at least 3 locations in CodeGen now and the checks are not consistent. Can you factor the attribute lookup into a shared helper? If you would prefer not doing it in this PR, can you file an issue?

[inteltdelisle]

I will submit a follow-up issue + PR.

[elizabethandrews]

We have a function in addScopeAttrToLocalVars which adds an implicit scope attribute to PFWG locals. This would affect those cases as well right? Is that intended?

[inteltdelisle]

My current understanding is that addScopeAttrToLocalVars can cause linking errors if the destructor is non-trivial. This replaces the linking error with silently dropping the destructor. Once I have a concise reproduction for that case, I'll submit a separate issue to address that.

[inteltdelisle]

#23461

[inteltdelisle]

@gmlueck @koparasy Now that the tests are passing, I believe this is the set of changes that are needed to get is_device_copyable to be consistent for accessors across host and device. I made accessors no-longer trivially copyable by making the destructors non-trivial. I did not investigate the alternative solution of making the copy/move constructors non trivial. (Claude Code assumes that alternative isn't simpler).

Any thoughts on this approach?

[koparasy]

@gmlueck @koparasy Now that the tests are passing, I believe this is the set of changes that are needed to get is_device_copyable to be consistent for accessors across host and device. I made accessors no-longer trivially copyable by making the destructors non-trivial. I did not investigate the alternative solution of making the copy/move constructors non trivial. (Claude Code assumes that alternative isn't simpler).

Any thoughts on this approach?

I will come back monday about this. Initially I imagined this will be a trivial direction. I got that wrong. I need to reason about this PR a little deeper.

[koparasy]

With is_legal_kernel_parameter_v in ShouldEnableSetArg, the generic set_arg(int, T&&) is now enabled on host for every type is_sycl_accessor matches. Accessors with the default property list still reach the accessor-specific setArgHelper, but an accessor with a non-empty accessor_property_list, and unsampled_image_accessor / sampled_image_accessor, have no matching overload. IIUC they fall into the generic setArgHelper(T&&), which copies the host object bytes as kind_std_layout. Previously these failed to compile; now they compile and pass a wrong argument. Could you verify how those are resolved with some tests, and act accordingly?

[inteltdelisle]

With is_legal_kernel_parameter_v in ShouldEnableSetArg, the generic set_arg(int, T&&) is now enabled on host for every type is_sycl_accessor matches. Accessors with the default property list still reach the accessor-specific setArgHelper, but an accessor with a non-empty accessor_property_list, and unsampled_image_accessor / sampled_image_accessor, have no matching overload. IIUC they fall into the generic setArgHelper(T&&), which copies the host object bytes as kind_std_layout. Previously these failed to compile; now they compile and pass a wrong argument. Could you verify how those are resolved with some tests, and act accordingly?

I refined the condition so that it bahaviour is unchanged and added tests to confirm this. I also added a static assert to make sure the std layour is not used incorrectly.