Skip to content

[SYCL] Add noexcept and hidden friends to nd_item #22902

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

Merged
KornevNikita merged 2 commits into from
Aug 13, 2026
+180 −36
Merged

[SYCL] Add noexcept and hidden friends to nd_item #22902

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
68b2d3e
[SYCL] Add noexcept and hidden friends to nd_item
crystarm Aug 8, 2026
03c10c4
[SYCL] Complete nd_item interface updates
crystarm Aug 10, 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
77 changes: 42 additions & 35 deletions sycl/include/sycl/nd_item.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -47,21 +47,23 @@ template <int Dimensions = 1> class nd_item {
public:
static constexpr int dimensions = Dimensions;
Comment thread
KornevNikita marked this conversation as resolved.

id<Dimensions> get_global_id() const {
nd_item() = delete;

id<Dimensions> get_global_id() const noexcept {
#ifdef __SYCL_DEVICE_ONLY__
return __spirv::initBuiltInGlobalInvocationId<Dimensions, id<Dimensions>>();
#else
return {};
#endif
}

size_t __SYCL_ALWAYS_INLINE get_global_id(int Dimension) const {
size_t __SYCL_ALWAYS_INLINE get_global_id(int Dimension) const noexcept {
size_t Id = get_global_id()[Dimension];
__SYCL_ASSUME_ID_RANGE(Id);
return Id;
}

size_t __SYCL_ALWAYS_INLINE get_global_linear_id() const {
size_t __SYCL_ALWAYS_INLINE get_global_linear_id() const noexcept {
size_t LinId = 0;
id<Dimensions> Index = get_global_id();
range<Dimensions> Extent = get_global_range();
Expand All @@ -78,21 +80,21 @@ template <int Dimensions = 1> class nd_item {
return LinId;
}

id<Dimensions> get_local_id() const {
id<Dimensions> get_local_id() const noexcept {
#ifdef __SYCL_DEVICE_ONLY__
return __spirv::initBuiltInLocalInvocationId<Dimensions, id<Dimensions>>();
#else
return {};
#endif
}

size_t __SYCL_ALWAYS_INLINE get_local_id(int Dimension) const {
size_t __SYCL_ALWAYS_INLINE get_local_id(int Dimension) const noexcept {
size_t Id = get_local_id()[Dimension];
__SYCL_ASSUME_ID_RANGE(Id);
return Id;
}

size_t get_local_linear_id() const {
size_t get_local_linear_id() const noexcept {
size_t LinId = 0;
id<Dimensions> Index = get_local_id();
range<Dimensions> Extent = get_local_range();
Expand All @@ -108,23 +110,23 @@ template <int Dimensions = 1> class nd_item {
return LinId;
}

group<Dimensions> get_group() const {
group<Dimensions> get_group() const noexcept {
// TODO: ideally Group object should be stateless and have a contructor with
// no arguments.
return detail::Builder::createGroup(get_global_range(), get_local_range(),
get_group_range(), get_group_id());
}

// Out-of-class definition in sub_group.hpp
sub_group get_sub_group() const;
sub_group get_sub_group() const noexcept;

size_t __SYCL_ALWAYS_INLINE get_group(int Dimension) const {
size_t __SYCL_ALWAYS_INLINE get_group(int Dimension) const noexcept {
size_t Id = get_group_id()[Dimension];
__SYCL_ASSUME_ID_RANGE(Id);
return Id;
}

size_t __SYCL_ALWAYS_INLINE get_group_linear_id() const {
size_t __SYCL_ALWAYS_INLINE get_group_linear_id() const noexcept {
size_t LinId = 0;
id<Dimensions> Index = get_group_id();
range<Dimensions> Extent = get_group_range();
Expand All @@ -140,58 +142,58 @@ template <int Dimensions = 1> class nd_item {
return LinId;
}

range<Dimensions> get_group_range() const {
range<Dimensions> get_group_range() const noexcept {
#ifdef __SYCL_DEVICE_ONLY__
return __spirv::initBuiltInNumWorkgroups<Dimensions, range<Dimensions>>();
#else
return {};
#endif
}

size_t __SYCL_ALWAYS_INLINE get_group_range(int Dimension) const {
size_t __SYCL_ALWAYS_INLINE get_group_range(int Dimension) const noexcept {
size_t Range = get_group_range()[Dimension];
__SYCL_ASSUME_ID_RANGE(Range);
return Range;
}

range<Dimensions> get_global_range() const {
range<Dimensions> get_global_range() const noexcept {
#ifdef __SYCL_DEVICE_ONLY__
return __spirv::initBuiltInGlobalSize<Dimensions, range<Dimensions>>();
#else
return {};
#endif
}

size_t get_global_range(int Dimension) const {
size_t get_global_range(int Dimension) const noexcept {
size_t Val = get_global_range()[Dimension];
__SYCL_ASSUME_ID_RANGE(Val);
return Val;
}

range<Dimensions> get_local_range() const {
range<Dimensions> get_local_range() const noexcept {
#ifdef __SYCL_DEVICE_ONLY__
return __spirv::initBuiltInWorkgroupSize<Dimensions, range<Dimensions>>();
#else
return {};
#endif
}

size_t get_local_range(int Dimension) const {
size_t get_local_range(int Dimension) const noexcept {
size_t Id = get_local_range()[Dimension];
__SYCL_ASSUME_ID_RANGE(Id);
return Id;
}

__SYCL2020_DEPRECATED("offsets are deprecated in SYCL 2020")
id<Dimensions> get_offset() const {
id<Dimensions> get_offset() const noexcept {
#ifdef __SYCL_DEVICE_ONLY__
return __spirv::initBuiltInGlobalOffset<Dimensions, id<Dimensions>>();
#else
return {};
#endif
}

nd_range<Dimensions> get_nd_range() const {
nd_range<Dimensions> get_nd_range() const noexcept {
return nd_range<Dimensions>(get_global_range(), get_local_range(),
get_offset());
}
Expand Down Expand Up @@ -248,7 +250,7 @@ template <int Dimensions = 1> class nd_item {
src,
[[maybe_unused]] size_t numElements,
[[maybe_unused]] size_t srcStride)
const {
const noexcept {
#ifdef __SYCL_DEVICE_ONLY__
__ocl_event_t E = __spirv_GroupAsyncCopy(
__spv::Scope::Workgroup, detail::convertToOpenCLGroupAsyncCopyPtr(dest),
Expand All @@ -274,7 +276,7 @@ template <int Dimensions = 1> class nd_item {
[[maybe_unused]] local_ptr<dataT> src,
[[maybe_unused]] size_t numElements,
[[maybe_unused]] size_t destStride)
const {
const noexcept {
#ifdef __SYCL_DEVICE_ONLY__
__ocl_event_t E = __spirv_GroupAsyncCopy(
__spv::Scope::Workgroup, detail::convertToOpenCLGroupAsyncCopyPtr(dest),
Expand All @@ -299,7 +301,7 @@ template <int Dimensions = 1> class nd_item {
async_work_group_copy([[maybe_unused]] decorated_local_ptr<DestDataT> dest,
[[maybe_unused]] decorated_global_ptr<SrcDataT> src,
[[maybe_unused]] size_t numElements,
[[maybe_unused]] size_t srcStride) const {
[[maybe_unused]] size_t srcStride) const noexcept {
#ifdef __SYCL_DEVICE_ONLY__
__ocl_event_t E = __spirv_GroupAsyncCopy(
__spv::Scope::Workgroup, detail::convertToOpenCLGroupAsyncCopyPtr(dest),
Expand All @@ -324,7 +326,7 @@ template <int Dimensions = 1> class nd_item {
async_work_group_copy([[maybe_unused]] decorated_global_ptr<DestDataT> dest,
[[maybe_unused]] decorated_local_ptr<SrcDataT> src,
[[maybe_unused]] size_t numElements,
[[maybe_unused]] size_t destStride) const {
[[maybe_unused]] size_t destStride) const noexcept {
#ifdef __SYCL_DEVICE_ONLY__
__ocl_event_t E = __spirv_GroupAsyncCopy(
__spv::Scope::Workgroup, detail::convertToOpenCLGroupAsyncCopyPtr(dest),
Expand Down Expand Up @@ -352,7 +354,7 @@ template <int Dimensions = 1> class nd_item {
access::decorated::legacy>
Src,
size_t NumElements,
size_t Stride) const {
size_t Stride) const noexcept {
static_assert(sizeof(bool) == sizeof(uint8_t),
"Async copy to/from bool memory is not supported.");
auto DestP = multi_ptr<uint8_t, DestS, access::decorated::legacy>(
Expand All @@ -378,7 +380,7 @@ template <int Dimensions = 1> class nd_item {
access::decorated::legacy>
Src,
size_t NumElements,
size_t Stride) const {
size_t Stride) const noexcept {
static_assert(sizeof(bool) == sizeof(uint8_t),
"Async copy to/from bool memory is not supported.");
using VecT = detail::change_base_type_t<T, uint8_t>;
Expand All @@ -401,7 +403,7 @@ template <int Dimensions = 1> class nd_item {
device_event>
async_work_group_copy(multi_ptr<DestT, DestS, access::decorated::yes> Dest,
multi_ptr<SrcT, SrcS, access::decorated::yes> Src,
size_t NumElements, size_t Stride) const {
size_t NumElements, size_t Stride) const noexcept {
static_assert(sizeof(bool) == sizeof(uint8_t),
"Async copy to/from bool memory is not supported.");
using QualSrcT =
Expand Down Expand Up @@ -429,7 +431,7 @@ template <int Dimensions = 1> class nd_item {
device_event>
async_work_group_copy(multi_ptr<DestT, DestS, access::decorated::yes> Dest,
multi_ptr<SrcT, SrcS, access::decorated::yes> Src,
size_t NumElements, size_t Stride) const {
size_t NumElements, size_t Stride) const noexcept {
static_assert(sizeof(bool) == sizeof(uint8_t),
"Async copy to/from bool memory is not supported.");
using VecT = detail::change_base_type_t<DestT, uint8_t>;
Expand All @@ -455,7 +457,7 @@ template <int Dimensions = 1> class nd_item {
__SYCL2020_DEPRECATED("Use decorated multi_ptr arguments instead")
device_event
async_work_group_copy(local_ptr<dataT> dest, global_ptr<dataT> src,
size_t numElements) const {
size_t numElements) const noexcept {
return async_work_group_copy(dest, src, numElements, 1);
}

Expand All @@ -468,7 +470,7 @@ template <int Dimensions = 1> class nd_item {
__SYCL2020_DEPRECATED("Use decorated multi_ptr arguments instead")
device_event
async_work_group_copy(global_ptr<dataT> dest, local_ptr<dataT> src,
size_t numElements) const {
size_t numElements) const noexcept {
return async_work_group_copy(dest, src, numElements, 1);
}

Expand All @@ -483,7 +485,7 @@ template <int Dimensions = 1> class nd_item {
std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>, device_event>
async_work_group_copy(decorated_local_ptr<DestDataT> dest,
decorated_global_ptr<SrcDataT> src,
size_t numElements) const {
size_t numElements) const noexcept {
return async_work_group_copy(dest, src, numElements, 1);
}

Expand All @@ -498,11 +500,12 @@ template <int Dimensions = 1> class nd_item {
std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>, device_event>
async_work_group_copy(decorated_global_ptr<DestDataT> dest,
decorated_local_ptr<SrcDataT> src,
size_t numElements) const {
size_t numElements) const noexcept {
return async_work_group_copy(dest, src, numElements, 1);
}

template <typename... eventTN> void wait_for(eventTN... events) const {
template <typename... eventTN>
void wait_for(eventTN... events) const noexcept {
(events.wait(), ...);
}

Expand All @@ -517,16 +520,20 @@ template <int Dimensions = 1> class nd_item {
nd_item &operator=(const nd_item &rhs) = default;
nd_item &operator=(nd_item &&rhs) = default;

bool operator==(const nd_item &) const { return true; }
bool operator!=(const nd_item &rhs) const { return !((*this) == rhs); }
friend bool operator==(const nd_item &, const nd_item &) noexcept {
return true;
}

friend bool operator!=(const nd_item &lhs, const nd_item &rhs) noexcept {
return !(lhs == rhs);
}

protected:
friend class detail::Builder;
nd_item() {}
nd_item(const item<Dimensions, true> &, const item<Dimensions, false> &,
const group<Dimensions> &) {}

id<Dimensions> get_group_id() const {
id<Dimensions> get_group_id() const noexcept {
#ifdef __SYCL_DEVICE_ONLY__
return __spirv::initBuiltInWorkgroupId<Dimensions, id<Dimensions>>();
#else
Expand Down
3 changes: 2 additions & 1 deletion sycl/include/sycl/sub_group.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -197,7 +197,8 @@ struct sub_group {
sub_group() = default;
};

template <int Dimensions> sub_group nd_item<Dimensions>::get_sub_group() const {
template <int Dimensions>
sub_group nd_item<Dimensions>::get_sub_group() const noexcept {
return sub_group();
}

Expand Down
Loading
Loading