From 68b2d3e2ce447114130641c4ebdd806538adf568 Mon Sep 17 00:00:00 2001 From: Georgij Tsarin Date: Sat, 8 Aug 2026 13:23:02 +0300 Subject: [PATCH 1/2] [SYCL] Add noexcept and hidden friends to nd_item --- sycl/include/sycl/nd_item.hpp | 78 ++++++----- sycl/include/sycl/sub_group.hpp | 3 +- sycl/test/basic_tests/nd_item_interface.cpp | 139 ++++++++++++++++++++ 3 files changed, 183 insertions(+), 37 deletions(-) create mode 100644 sycl/test/basic_tests/nd_item_interface.cpp diff --git a/sycl/include/sycl/nd_item.hpp b/sycl/include/sycl/nd_item.hpp index 62eb7806e4517..27d6e3da27c48 100644 --- a/sycl/include/sycl/nd_item.hpp +++ b/sycl/include/sycl/nd_item.hpp @@ -47,7 +47,7 @@ template class nd_item { public: static constexpr int dimensions = Dimensions; - id get_global_id() const { + id get_global_id() const noexcept { #ifdef __SYCL_DEVICE_ONLY__ return __spirv::initBuiltInGlobalInvocationId>(); #else @@ -55,13 +55,13 @@ template class nd_item { #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 Index = get_global_id(); range Extent = get_global_range(); @@ -78,7 +78,7 @@ template class nd_item { return LinId; } - id get_local_id() const { + id get_local_id() const noexcept { #ifdef __SYCL_DEVICE_ONLY__ return __spirv::initBuiltInLocalInvocationId>(); #else @@ -86,13 +86,13 @@ template class nd_item { #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 Index = get_local_id(); range Extent = get_local_range(); @@ -108,7 +108,7 @@ template class nd_item { return LinId; } - group get_group() const { + group 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(), @@ -116,15 +116,15 @@ template class nd_item { } // 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 Index = get_group_id(); range Extent = get_group_range(); @@ -140,7 +140,7 @@ template class nd_item { return LinId; } - range get_group_range() const { + range get_group_range() const noexcept { #ifdef __SYCL_DEVICE_ONLY__ return __spirv::initBuiltInNumWorkgroups>(); #else @@ -148,13 +148,13 @@ template class nd_item { #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 get_global_range() const { + range get_global_range() const noexcept { #ifdef __SYCL_DEVICE_ONLY__ return __spirv::initBuiltInGlobalSize>(); #else @@ -162,13 +162,13 @@ template class nd_item { #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 get_local_range() const { + range get_local_range() const noexcept { #ifdef __SYCL_DEVICE_ONLY__ return __spirv::initBuiltInWorkgroupSize>(); #else @@ -176,14 +176,14 @@ template class nd_item { #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 get_offset() const { + id get_offset() const noexcept { #ifdef __SYCL_DEVICE_ONLY__ return __spirv::initBuiltInGlobalOffset>(); #else @@ -191,7 +191,7 @@ template class nd_item { #endif } - nd_range get_nd_range() const { + nd_range get_nd_range() const noexcept { return nd_range(get_global_range(), get_local_range(), get_offset()); } @@ -199,7 +199,7 @@ template class nd_item { #ifndef __INTEL_PREVIEW_BREAKING_CHANGES __SYCL2020_DEPRECATED("use sycl::group_barrier() free function instead") void barrier([[maybe_unused]] access::fence_space accessSpace = - access::fence_space::global_and_local) const { + access::fence_space::global_and_local) const noexcept { #ifdef __SYCL_DEVICE_ONLY__ uint32_t flags = _V1::detail::getSPIRVMemorySemanticsMask(accessSpace); __spirv_ControlBarrier(__spv::Scope::Workgroup, __spv::Scope::Workgroup, @@ -217,7 +217,7 @@ template class nd_item { accessMode == access::mode::write || accessMode == access::mode::read_write, access::fence_space> - accessSpace = access::fence_space::global_and_local) const { + accessSpace = access::fence_space::global_and_local) const noexcept { #if __SYCL_DEVICE_ONLY__ uint32_t flags = detail::getSPIRVMemorySemanticsMask(accessSpace); // TODO: currently, there is no good way in SPIR-V to set the memory @@ -248,7 +248,7 @@ template 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), @@ -274,7 +274,7 @@ template class nd_item { [[maybe_unused]] local_ptr 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), @@ -299,7 +299,7 @@ template class nd_item { async_work_group_copy([[maybe_unused]] decorated_local_ptr dest, [[maybe_unused]] decorated_global_ptr 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), @@ -324,7 +324,7 @@ template class nd_item { async_work_group_copy([[maybe_unused]] decorated_global_ptr dest, [[maybe_unused]] decorated_local_ptr 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), @@ -352,7 +352,7 @@ template 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( @@ -378,7 +378,7 @@ template 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; @@ -401,7 +401,7 @@ template class nd_item { device_event> async_work_group_copy(multi_ptr Dest, multi_ptr 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 = @@ -429,7 +429,7 @@ template class nd_item { device_event> async_work_group_copy(multi_ptr Dest, multi_ptr 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; @@ -455,7 +455,7 @@ template class nd_item { __SYCL2020_DEPRECATED("Use decorated multi_ptr arguments instead") device_event async_work_group_copy(local_ptr dest, global_ptr src, - size_t numElements) const { + size_t numElements) const noexcept { return async_work_group_copy(dest, src, numElements, 1); } @@ -468,7 +468,7 @@ template class nd_item { __SYCL2020_DEPRECATED("Use decorated multi_ptr arguments instead") device_event async_work_group_copy(global_ptr dest, local_ptr src, - size_t numElements) const { + size_t numElements) const noexcept { return async_work_group_copy(dest, src, numElements, 1); } @@ -483,7 +483,7 @@ template class nd_item { std::is_same_v>, device_event> async_work_group_copy(decorated_local_ptr dest, decorated_global_ptr src, - size_t numElements) const { + size_t numElements) const noexcept { return async_work_group_copy(dest, src, numElements, 1); } @@ -498,11 +498,12 @@ template class nd_item { std::is_same_v>, device_event> async_work_group_copy(decorated_global_ptr dest, decorated_local_ptr src, - size_t numElements) const { + size_t numElements) const noexcept { return async_work_group_copy(dest, src, numElements, 1); } - template void wait_for(eventTN... events) const { + template + void wait_for(eventTN... events) const noexcept { (events.wait(), ...); } @@ -517,8 +518,13 @@ template 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; @@ -526,7 +532,7 @@ template class nd_item { nd_item(const item &, const item &, const group &) {} - id get_group_id() const { + id get_group_id() const noexcept { #ifdef __SYCL_DEVICE_ONLY__ return __spirv::initBuiltInWorkgroupId>(); #else diff --git a/sycl/include/sycl/sub_group.hpp b/sycl/include/sycl/sub_group.hpp index d25ea31a1b88a..443c7c62c0c12 100644 --- a/sycl/include/sycl/sub_group.hpp +++ b/sycl/include/sycl/sub_group.hpp @@ -197,7 +197,8 @@ struct sub_group { sub_group() = default; }; -template sub_group nd_item::get_sub_group() const { +template +sub_group nd_item::get_sub_group() const noexcept { return sub_group(); } diff --git a/sycl/test/basic_tests/nd_item_interface.cpp b/sycl/test/basic_tests/nd_item_interface.cpp new file mode 100644 index 0000000000000..329ac442a4af0 --- /dev/null +++ b/sycl/test/basic_tests/nd_item_interface.cpp @@ -0,0 +1,139 @@ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -Wno-deprecated-declarations -fsyntax-only %s + +//==--------- nd_item_interface.cpp - nd_item interface test ---------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include +#include +#include +#include + +#include +#include +#include + +template +struct has_member_equal : std::false_type {}; + +template +struct has_member_equal< + T, std::void_t().operator==( + std::declval()))>> : std::true_type {}; + +template +struct has_member_not_equal : std::false_type {}; + +template +struct has_member_not_equal< + T, std::void_t().operator!=( + std::declval()))>> : std::true_type {}; + +using Item = sycl::nd_item<1>; + +static_assert(noexcept(std::declval().get_global_id())); +static_assert(noexcept(std::declval().get_global_id(0))); +static_assert(noexcept(std::declval().get_global_linear_id())); +static_assert(noexcept(std::declval().get_local_id())); +static_assert(noexcept(std::declval().get_local_id(0))); +static_assert(noexcept(std::declval().get_local_linear_id())); +static_assert(noexcept(std::declval().get_group())); +static_assert(noexcept(std::declval().get_sub_group())); +static_assert(noexcept(std::declval().get_group(0))); +static_assert(noexcept(std::declval().get_group_linear_id())); +static_assert(noexcept(std::declval().get_group_range())); +static_assert(noexcept(std::declval().get_group_range(0))); +static_assert(noexcept(std::declval().get_global_range())); +static_assert(noexcept(std::declval().get_global_range(0))); +static_assert(noexcept(std::declval().get_local_range())); +static_assert(noexcept(std::declval().get_local_range(0))); +static_assert(noexcept(std::declval().get_offset())); +static_assert(noexcept(std::declval().get_nd_range())); + +#ifndef __INTEL_PREVIEW_BREAKING_CHANGES +static_assert(noexcept(std::declval().barrier())); +static_assert(noexcept(std::declval().mem_fence())); +#endif + +static_assert(!has_member_equal::value); +static_assert(!has_member_not_equal::value); +static_assert(std::is_same_v() == + std::declval()), + bool>); +static_assert(std::is_same_v() != + std::declval()), + bool>); +static_assert(noexcept(std::declval() == + std::declval())); +static_assert(noexcept(std::declval() != + std::declval())); +static_assert(noexcept(operator==(std::declval(), + std::declval()))); +static_assert(noexcept(operator!=(std::declval(), + std::declval()))); + +using LegacyLocalIntPtr = sycl::local_ptr; +using LegacyGlobalIntPtr = sycl::global_ptr; +using DecoratedLocalIntPtr = sycl::decorated_local_ptr; +using DecoratedGlobalIntPtr = sycl::decorated_global_ptr; +using DecoratedLocalConstIntPtr = sycl::decorated_local_ptr; +using DecoratedGlobalConstIntPtr = sycl::decorated_global_ptr; + +static_assert(noexcept(std::declval().async_work_group_copy( + std::declval(), std::declval(), + std::size_t{}, std::size_t{}))); +static_assert(noexcept(std::declval().async_work_group_copy( + std::declval(), std::declval(), + std::size_t{}, std::size_t{}))); +static_assert(noexcept(std::declval().async_work_group_copy( + std::declval(), + std::declval(), std::size_t{}, std::size_t{}))); +static_assert(noexcept(std::declval().async_work_group_copy( + std::declval(), + std::declval(), std::size_t{}, std::size_t{}))); + +using BoolVector = sycl::vec; +using LegacyLocalBoolPtr = sycl::local_ptr; +using LegacyGlobalBoolPtr = sycl::global_ptr; +using LegacyLocalBoolVectorPtr = sycl::local_ptr; +using LegacyGlobalBoolVectorPtr = sycl::global_ptr; +using DecoratedLocalBoolPtr = sycl::decorated_local_ptr; +using DecoratedGlobalConstBoolPtr = sycl::decorated_global_ptr; +using DecoratedLocalBoolVectorPtr = sycl::decorated_local_ptr; +using DecoratedGlobalConstBoolVectorPtr = + sycl::decorated_global_ptr; + +static_assert(noexcept(std::declval().async_work_group_copy( + std::declval(), std::declval(), + std::size_t{}, std::size_t{}))); +static_assert(noexcept(std::declval().async_work_group_copy( + std::declval(), + std::declval(), std::size_t{}, std::size_t{}))); +static_assert(noexcept(std::declval().async_work_group_copy( + std::declval(), + std::declval(), std::size_t{}, + std::size_t{}))); +static_assert(noexcept(std::declval().async_work_group_copy( + std::declval(), + std::declval(), std::size_t{}, + std::size_t{}))); + +static_assert(noexcept(std::declval().async_work_group_copy( + std::declval(), std::declval(), + std::size_t{}))); +static_assert(noexcept(std::declval().async_work_group_copy( + std::declval(), std::declval(), + std::size_t{}))); +static_assert(noexcept(std::declval().async_work_group_copy( + std::declval(), + std::declval(), std::size_t{}))); +static_assert(noexcept(std::declval().async_work_group_copy( + std::declval(), + std::declval(), std::size_t{}))); + +static_assert(noexcept( + std::declval().wait_for(std::declval()))); From 03c10c43d7752791cbddb60e23b6b5572a85a9e2 Mon Sep 17 00:00:00 2001 From: Georgij Tsarin Date: Mon, 10 Aug 2026 16:37:08 +0300 Subject: [PATCH 2/2] [SYCL] Complete nd_item interface updates --- sycl/include/sycl/nd_item.hpp | 7 ++++--- sycl/test/basic_tests/nd_item_interface.cpp | 7 ++----- 2 files changed, 6 insertions(+), 8 deletions(-) diff --git a/sycl/include/sycl/nd_item.hpp b/sycl/include/sycl/nd_item.hpp index 27d6e3da27c48..fa1f95f8c5a69 100644 --- a/sycl/include/sycl/nd_item.hpp +++ b/sycl/include/sycl/nd_item.hpp @@ -47,6 +47,8 @@ template class nd_item { public: static constexpr int dimensions = Dimensions; + nd_item() = delete; + id get_global_id() const noexcept { #ifdef __SYCL_DEVICE_ONLY__ return __spirv::initBuiltInGlobalInvocationId>(); @@ -199,7 +201,7 @@ template class nd_item { #ifndef __INTEL_PREVIEW_BREAKING_CHANGES __SYCL2020_DEPRECATED("use sycl::group_barrier() free function instead") void barrier([[maybe_unused]] access::fence_space accessSpace = - access::fence_space::global_and_local) const noexcept { + access::fence_space::global_and_local) const { #ifdef __SYCL_DEVICE_ONLY__ uint32_t flags = _V1::detail::getSPIRVMemorySemanticsMask(accessSpace); __spirv_ControlBarrier(__spv::Scope::Workgroup, __spv::Scope::Workgroup, @@ -217,7 +219,7 @@ template class nd_item { accessMode == access::mode::write || accessMode == access::mode::read_write, access::fence_space> - accessSpace = access::fence_space::global_and_local) const noexcept { + accessSpace = access::fence_space::global_and_local) const { #if __SYCL_DEVICE_ONLY__ uint32_t flags = detail::getSPIRVMemorySemanticsMask(accessSpace); // TODO: currently, there is no good way in SPIR-V to set the memory @@ -528,7 +530,6 @@ template class nd_item { protected: friend class detail::Builder; - nd_item() {} nd_item(const item &, const item &, const group &) {} diff --git a/sycl/test/basic_tests/nd_item_interface.cpp b/sycl/test/basic_tests/nd_item_interface.cpp index 329ac442a4af0..7c0f1385767af 100644 --- a/sycl/test/basic_tests/nd_item_interface.cpp +++ b/sycl/test/basic_tests/nd_item_interface.cpp @@ -35,6 +35,8 @@ struct has_member_not_equal< using Item = sycl::nd_item<1>; +static_assert(!std::is_default_constructible_v); + static_assert(noexcept(std::declval().get_global_id())); static_assert(noexcept(std::declval().get_global_id(0))); static_assert(noexcept(std::declval().get_global_linear_id())); @@ -54,11 +56,6 @@ static_assert(noexcept(std::declval().get_local_range(0))); static_assert(noexcept(std::declval().get_offset())); static_assert(noexcept(std::declval().get_nd_range())); -#ifndef __INTEL_PREVIEW_BREAKING_CHANGES -static_assert(noexcept(std::declval().barrier())); -static_assert(noexcept(std::declval().mem_fence())); -#endif - static_assert(!has_member_equal::value); static_assert(!has_member_not_equal::value); static_assert(std::is_same_v() ==