Skip to content

[SYCL] Consistent is_device_copyable for sycl::accessor #23250

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
inteltdelisle wants to merge 15 commits into
base: sycl
Choose a base branch
from
Open

[SYCL] Consistent is_device_copyable for sycl::accessor #23250

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
15 commits
Select commit Hold shift + click to select a range
54550f9
[SYCL] Consistent is_device_copyable for sycl::accessor
inteltdelisle Sep 23, 2026
4771df0
[SYCL] added user-defined destructors to accessors
inteltdelisle Sep 25, 2026
4bc0a59
Fixed formatting missed by git clang-format
inteltdelisle Sep 25, 2026
74cbd87
added exception for accessor to is_device_copyable
inteltdelisle Oct 1, 2026
3b87544
formatting
inteltdelisle Oct 1, 2026
7478b51
[SYCL] Re-write CheckDeviceCopyable to match spec more closely.
inteltdelisle Oct 2, 2026
8d933d1
[SYCL] update handling of sycl special types to finalize destructor
inteltdelisle Oct 5, 2026
faecfa9
[SYCL] Do not emit destructors for device work groups.
inteltdelisle Oct 6, 2026
37e133a
[SYCL] Add local accessor dtor to verified esimd functions.
inteltdelisle Oct 7, 2026
82c5e91
Merge branch 'sycl' into tdelisle/runtime/accessor_is_device_copyable
inteltdelisle Oct 7, 2026
47f7932
[SYCL] Add explicit copy and move ctor declarations.
inteltdelisle Oct 7, 2026
e966767
[SYCL] Relax deduction rules for sycl::handler::set_arg.
inteltdelisle Oct 8, 2026
f011bd9
[SYCL] Rework is_legal_kernel_parameter to be much stricter
inteltdelisle Oct 13, 2026
7b15c7c
Merge branch 'sycl' into tdelisle/runtime/accessor_is_device_copyable
inteltdelisle Oct 14, 2026
0adb720
Address review comments.
inteltdelisle Oct 19, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
14 changes: 14 additions & 0 deletions clang/lib/CodeGen/CGDeclCXX.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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

@elizabethandrews elizabethandrews Oct 10, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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 inteltdelisle Oct 13, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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 inteltdelisle Oct 13, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

#23461

// 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) {

@elizabethandrews elizabethandrews Oct 10, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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 inteltdelisle Oct 13, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I will submit a follow-up issue + PR.

const SYCLScopeAttr *Scope = D.getAttr<SYCLScopeAttr>();
if (!Scope)
if (const auto *RD = D.getType()->getAsCXXRecordDecl())
Scope = RD->getAttr<SYCLScopeAttr>();
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
Expand Down
10 changes: 6 additions & 4 deletions clang/lib/Sema/SemaSYCL.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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(),
Expand All @@ -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())
Expand Down Expand Up @@ -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";
Expand Down
10 changes: 5 additions & 5 deletions clang/test/CodeGenSYCL/free_function_int_header.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1609,7 +1609,7 @@ void ff_28(TemplatedAccessorStruct<int> 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<AccessorAndLocalAccessor>: std::true_type {};
// CHECK-NEXT: struct sycl::detail::is_legal_kernel_parameter<AccessorAndLocalAccessor>: std::true_type {};
// CHECK-NEXT: template <>
// CHECK-NEXT: struct sycl::ext::oneapi::experimental::detail::is_struct_with_special_type<AccessorAndLocalAccessor> {
// CHECK-NEXT: inline static constexpr bool value = true;
Expand Down Expand Up @@ -1637,7 +1637,7 @@ void ff_28(TemplatedAccessorStruct<int> 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<SecondLevelAccessor>: std::true_type {};
// CHECK-NEXT: struct sycl::detail::is_legal_kernel_parameter<SecondLevelAccessor>: std::true_type {};
// CHECK-NEXT: template <>
// CHECK-NEXT: struct sycl::ext::oneapi::experimental::detail::is_struct_with_special_type<SecondLevelAccessor> {
// CHECK-NEXT: inline static constexpr bool value = true;
Expand Down Expand Up @@ -1665,7 +1665,7 @@ void ff_28(TemplatedAccessorStruct<int> 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<IntAndAccessor>: std::true_type {};
// CHECK-NEXT: struct sycl::detail::is_legal_kernel_parameter<IntAndAccessor>: std::true_type {};
// CHECK-NEXT: template <>
// CHECK-NEXT: struct sycl::ext::oneapi::experimental::detail::is_struct_with_special_type<IntAndAccessor> {
// CHECK-NEXT: inline static constexpr bool value = true;
Expand All @@ -1678,7 +1678,7 @@ void ff_28(TemplatedAccessorStruct<int> arg1) {
// CHECK-NEXT: };

// CHECK: template <>
// CHECK-NEXT: struct sycl::is_device_copyable<AccessorAndInt>: std::true_type {};
// CHECK-NEXT: struct sycl::detail::is_legal_kernel_parameter<AccessorAndInt>: std::true_type {};
// CHECK-NEXT: template <>
// CHECK-NEXT: struct sycl::ext::oneapi::experimental::detail::is_struct_with_special_type<AccessorAndInt> {
// CHECK-NEXT: inline static constexpr bool value = true;
Expand Down Expand Up @@ -1710,7 +1710,7 @@ void ff_28(TemplatedAccessorStruct<int> arg1) {
// CHECK: template <typename T> struct TemplatedAccessorStruct;
// CHECK: void ff_28(TemplatedAccessorStruct<int> arg1);
// CHECK-NEXT: template <>
// CHECK-NEXT: struct sycl::is_device_copyable<TemplatedAccessorStruct<int>>: std::true_type {};
// CHECK-NEXT: struct sycl::detail::is_legal_kernel_parameter<TemplatedAccessorStruct<int>>: std::true_type {};
// CHECK-NEXT: template <>
// CHECK-NEXT: struct sycl::ext::oneapi::experimental::detail::is_struct_with_special_type<TemplatedAccessorStruct<int>> {
// CHECK-NEXT: inline static constexpr bool value = true;
Expand Down
2 changes: 2 additions & 0 deletions llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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",
Expand Down
94 changes: 90 additions & 4 deletions sycl/include/sycl/accessor.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -757,6 +757,13 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
: impl({}, detail::InitializedVal<AdjustedDim, range>::template get<0>(),
detail::InitializedVal<AdjustedDim, range>::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} {}
Expand Down Expand Up @@ -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 <typename, int, access_mode> friend class host_accessor;

#endif // __SYCL_DEVICE_ONLY__
Expand Down Expand Up @@ -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 <typename DataT_,
typename = std::enable_if_t<
Expand Down Expand Up @@ -1592,7 +1613,18 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
: accessor(BufferRef, CommandGroupHandler, AccessRange, AccessOffset,
PropertyList, CodeLoc) {}

template <typename... NewPropsT>
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<NewPropsT...>()
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
// Same-type copies are handled by the copy constructor.
&& !std::is_same_v<std::remove_cv_t<PropertyListT>,
ext::oneapi::accessor_property_list<NewPropsT...>>
#endif // __INTEL_PREVIEW_BREAKING_CHANGES
>>
accessor(
const accessor<DataT, Dimensions, AccessMode, AccessTarget, IsPlaceholder,
ext::oneapi::accessor_property_list<NewPropsT...>> &Other,
Expand All @@ -1605,9 +1637,6 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
{
static_assert(detail::IsCxPropertyList<PropertyListT>::value,
"Conversion is only available for accessor_property_list");
static_assert(
PropertyListT::template areSameCompileTimeProperties<NewPropsT...>(),
"Compile-time-constant properties must be the same");
(void)CodeLoc;
#ifndef __SYCL_DEVICE_ONLY__
detail::constructorNotification(getMemoryObject(), impl.get(), AccessTarget,
Expand Down Expand Up @@ -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")
Expand Down Expand Up @@ -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:
Expand Down Expand Up @@ -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
// -------+---------+-------+----+----------+--------------
Expand Down Expand Up @@ -2826,6 +2887,31 @@ host_accessor(buffer<DataT, Dimensions, AllocatorT>, Type1, Type2, Type3, Type4,
Type5) -> host_accessor<DataT, Dimensions,
detail::deduceAccessMode<Type4, Type5>()>;

namespace detail {
// Add specializations for SYCL accessor types to the is_sycl_accessor trait.
// This is needed so that is_device_copyable<accessor> returns false, but
// they can still be captured in a kernel.
template <typename T> struct is_sycl_accessor;

template <access::target AccessTarget>
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 <typename DataT, int Dimensions, access::mode AccessMode,
access::target AccessTarget, access::placeholder IsPlaceholder,
typename PropertyListT>
struct is_sycl_accessor<accessor<DataT, Dimensions, AccessMode, AccessTarget,
IsPlaceholder, PropertyListT>>
: std::bool_constant<is_valid_kernel_access_target_v<AccessTarget>> {};

template <typename DataT, int Dimensions>
struct is_sycl_accessor<local_accessor<DataT, Dimensions>> : std::true_type {};
} // namespace detail
} // namespace _V1
} // namespace sycl

Expand Down
Loading
Loading