From 68e6c7fd2a05396e8b55020b9ef719094a723052 Mon Sep 17 00:00:00 2001 From: Grzegorz Mianowski Date: Thu, 6 Aug 2026 14:36:27 +0200 Subject: [PATCH 1/2] Add noexcept to group and sub_group methods. Equality operators changed to frienf functions --- sycl/include/sycl/group.hpp | 122 ++++++++++++++++++-------------- sycl/include/sycl/nd_item.hpp | 2 +- sycl/include/sycl/sub_group.hpp | 39 +++++----- 3 files changed, 89 insertions(+), 74 deletions(-) diff --git a/sycl/include/sycl/group.hpp b/sycl/include/sycl/group.hpp index b7ad9bd7a7bd9..eb44645327914 100644 --- a/sycl/include/sycl/group.hpp +++ b/sycl/include/sycl/group.hpp @@ -116,24 +116,24 @@ template class __SYCL_TYPE(group) group { group() = delete; __SYCL2020_DEPRECATED("use sycl::group::get_group_id() instead") - id get_id() const { return index; } + id get_id() const noexcept { return index; } __SYCL2020_DEPRECATED("use sycl::group::get_group_id() instead") - size_t get_id(int dimension) const { return index[dimension]; } + size_t get_id(int dimension) const noexcept { return index[dimension]; } - id get_group_id() const { return index; } + id get_group_id() const noexcept { return index; } - size_t get_group_id(int dimension) const { return index[dimension]; } + size_t get_group_id(int dimension) const noexcept { return index[dimension]; } __SYCL2020_DEPRECATED("calculate sycl::group::get_group_range() * " "sycl::group::get_max_local_range() instead") - range get_global_range() const { return globalRange; } + range get_global_range() const noexcept { return globalRange; } - size_t get_global_range(int dimension) const { + size_t get_global_range(int dimension) const noexcept { return globalRange[dimension]; } - id get_local_id() const { + id get_local_id() const noexcept { #ifdef __SYCL_DEVICE_ONLY__ return __spirv::initBuiltInLocalInvocationId>(); #else @@ -142,40 +142,48 @@ template class __SYCL_TYPE(group) group { #endif } - size_t get_local_id(int dimention) const { return get_local_id()[dimention]; } + size_t get_local_id(int dimention) const noexcept { + return get_local_id()[dimention]; + } - size_t get_local_linear_id() const { + size_t get_local_linear_id() const noexcept { return get_local_linear_id_impl(); } - range get_local_range() const { return localRange; } + range get_local_range() const noexcept { return localRange; } - size_t get_local_range(int dimension) const { return localRange[dimension]; } + size_t get_local_range(int dimension) const noexcept { + return localRange[dimension]; + } - size_t get_local_linear_range() const { + size_t get_local_linear_range() const noexcept { return get_local_linear_range_impl(); } - range get_group_range() const { return groupRange; } + range get_group_range() const noexcept { return groupRange; } - size_t get_group_range(int dimension) const { + size_t get_group_range(int dimension) const noexcept { return get_group_range()[dimension]; } - size_t get_group_linear_range() const { + size_t get_group_linear_range() const noexcept { return get_group_linear_range_impl(); } - range get_max_local_range() const { return get_local_range(); } + range get_max_local_range() const noexcept { + return get_local_range(); + } - size_t operator[](int dimension) const { return index[dimension]; } + size_t operator[](int dimension) const noexcept { return index[dimension]; } __SYCL2020_DEPRECATED("use sycl::group::get_group_linear_id() instead") - size_t get_linear_id() const { return get_group_linear_id(); } + size_t get_linear_id() const noexcept { return get_group_linear_id(); } - size_t get_group_linear_id() const { return get_group_linear_id_impl(); } + size_t get_group_linear_id() const noexcept { + return get_group_linear_id_impl(); + } - bool leader() const { return (get_local_linear_id() == 0); } + bool leader() const noexcept { return (get_local_linear_id() == 0); } // Note: These signatures for parallel_for_work_item are intentionally // non-conforming. The spec says this should take const WorkItemFunctionT &, @@ -188,7 +196,7 @@ template class __SYCL_TYPE(group) group { __attribute__((__libclc_call__)) #endif void - parallel_for_work_item(WorkItemFunctionT Func) const { + parallel_for_work_item(WorkItemFunctionT Func) const noexcept { // need barriers to enforce SYCL semantics for the work item loop - // compilers are expected to optimize when possible detail::workGroupBarrier(); @@ -244,7 +252,7 @@ template class __SYCL_TYPE(group) group { #endif void parallel_for_work_item(range flexibleRange, - WorkItemFunctionT Func) const { + WorkItemFunctionT Func) const noexcept { detail::workGroupBarrier(); #ifdef __SYCL_DEVICE_ONLY__ range GlobalSize{ @@ -308,7 +316,7 @@ template class __SYCL_TYPE(group) group { 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 { #ifdef __SYCL_DEVICE_ONLY__ uint32_t flags = detail::getSPIRVMemorySemanticsMask(accessSpace); // TODO: currently, there is no good way in SPIR-V to set the memory @@ -339,7 +347,7 @@ template class __SYCL_TYPE(group) group { 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), @@ -365,7 +373,7 @@ template class __SYCL_TYPE(group) group { [[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), @@ -390,7 +398,7 @@ template class __SYCL_TYPE(group) group { 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), @@ -415,7 +423,7 @@ template class __SYCL_TYPE(group) group { 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), @@ -443,7 +451,7 @@ template class __SYCL_TYPE(group) group { 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( @@ -469,7 +477,7 @@ template class __SYCL_TYPE(group) group { 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; @@ -492,7 +500,7 @@ template class __SYCL_TYPE(group) group { 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 = @@ -520,7 +528,7 @@ template class __SYCL_TYPE(group) group { 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; @@ -546,7 +554,7 @@ template class __SYCL_TYPE(group) group { __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); } @@ -559,7 +567,7 @@ template class __SYCL_TYPE(group) group { __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); } @@ -574,7 +582,7 @@ template class __SYCL_TYPE(group) group { 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); } @@ -589,24 +597,28 @@ template class __SYCL_TYPE(group) group { 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(), ...); } - bool operator==(const group &rhs) const { - bool Result = (rhs.globalRange == globalRange) && - (rhs.localRange == localRange) && (rhs.index == index); - __SYCL_ASSERT(rhs.groupRange == groupRange && + friend bool operator==(const group &lhs, + const group &rhs) noexcept { + bool Result = (rhs.globalRange == lhs.globalRange) && + (rhs.localRange == lhs.localRange) && + (rhs.index == lhs.index); + __SYCL_ASSERT(rhs.groupRange == lhs.groupRange && "inconsistent group class fields"); return Result; } - bool operator!=(const group &rhs) const { - return !((*this) == rhs); + friend bool operator!=(const group &lhs, + const group &rhs) noexcept { + return !(lhs == rhs); } private: @@ -617,21 +629,21 @@ template class __SYCL_TYPE(group) group { template typename std::enable_if_t<(dims == 1), size_t> - get_local_linear_id_impl() const { + get_local_linear_id_impl() const noexcept { id localId = get_local_id(); return localId[0]; } template typename std::enable_if_t<(dims == 2), size_t> - get_local_linear_id_impl() const { + get_local_linear_id_impl() const noexcept { id localId = get_local_id(); return localId[0] * localRange[1] + localId[1]; } template typename std::enable_if_t<(dims == 3), size_t> - get_local_linear_id_impl() const { + get_local_linear_id_impl() const noexcept { id localId = get_local_id(); return (localId[0] * localRange[1] * localRange[2]) + (localId[1] * localRange[2]) + localId[2]; @@ -639,55 +651,55 @@ template class __SYCL_TYPE(group) group { template typename std::enable_if_t<(dims == 1), size_t> - get_local_linear_range_impl() const { + get_local_linear_range_impl() const noexcept { auto localRange = get_local_range(); return localRange[0]; } template typename std::enable_if_t<(dims == 2), size_t> - get_local_linear_range_impl() const { + get_local_linear_range_impl() const noexcept { auto localRange = get_local_range(); return localRange[0] * localRange[1]; } template typename std::enable_if_t<(dims == 3), size_t> - get_local_linear_range_impl() const { + get_local_linear_range_impl() const noexcept { auto localRange = get_local_range(); return localRange[0] * localRange[1] * localRange[2]; } template typename std::enable_if_t<(dims == 1), size_t> - get_group_linear_range_impl() const { + get_group_linear_range_impl() const noexcept { auto groupRange = get_group_range(); return groupRange[0]; } template typename std::enable_if_t<(dims == 2), size_t> - get_group_linear_range_impl() const { + get_group_linear_range_impl() const noexcept { auto groupRange = get_group_range(); return groupRange[0] * groupRange[1]; } template typename std::enable_if_t<(dims == 3), size_t> - get_group_linear_range_impl() const { + get_group_linear_range_impl() const noexcept { auto groupRange = get_group_range(); return groupRange[0] * groupRange[1] * groupRange[2]; } template typename std::enable_if_t<(dims == 1), size_t> - get_group_linear_id_impl() const { + get_group_linear_id_impl() const noexcept { return index[0]; } template typename std::enable_if_t<(dims == 2), size_t> - get_group_linear_id_impl() const { + get_group_linear_id_impl() const noexcept { return index[0] * groupRange[1] + index[1]; } @@ -703,7 +715,7 @@ template class __SYCL_TYPE(group) group { // work-group id from a multi-dimensional index follows the equation 4.3. template typename std::enable_if_t<(dims == 3), size_t> - get_group_linear_id_impl() const { + get_group_linear_id_impl() const noexcept { return (index[0] * groupRange[1] * groupRange[2]) + (index[1] * groupRange[2]) + index[2]; } @@ -711,7 +723,7 @@ template class __SYCL_TYPE(group) group { protected: friend class detail::Builder; group(const range &G, const range &L, - const range GroupRange, const id &I) + const range GroupRange, const id &I) noexcept : globalRange(G), localRange(L), groupRange(GroupRange), index(I) {} }; } // namespace _V1 diff --git a/sycl/include/sycl/nd_item.hpp b/sycl/include/sycl/nd_item.hpp index 62eb7806e4517..707233ea97ab6 100644 --- a/sycl/include/sycl/nd_item.hpp +++ b/sycl/include/sycl/nd_item.hpp @@ -116,7 +116,7 @@ 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 Id = get_group_id()[Dimension]; diff --git a/sycl/include/sycl/sub_group.hpp b/sycl/include/sycl/sub_group.hpp index d25ea31a1b88a..677ec7e4b6858 100644 --- a/sycl/include/sycl/sub_group.hpp +++ b/sycl/include/sycl/sub_group.hpp @@ -10,11 +10,13 @@ #include #include + #ifndef __INTEL_PREVIEW_BREAKING_CHANGES #include // for __SYCL_DEPRECATED #endif // __INTEL_PREVIEW_BREAKING_CHANGES -#include // for id -#include // for memory_scope + +#include // for id +#include // for memory_scope #include #include // for range @@ -43,7 +45,7 @@ struct sub_group { /* --- common interface members --- */ - id_type get_local_id() const { + id_type get_local_id() const noexcept { #ifdef __SYCL_DEVICE_ONLY__ return __spirv_BuiltInSubgroupLocalInvocationId(); #else @@ -52,7 +54,7 @@ struct sub_group { #endif } - linear_id_type get_local_linear_id() const { + linear_id_type get_local_linear_id() const noexcept { #ifdef __SYCL_DEVICE_ONLY__ return static_cast(get_local_id()[0]); #else @@ -61,7 +63,7 @@ struct sub_group { #endif } - range_type get_local_range() const { + range_type get_local_range() const noexcept { #ifdef __SYCL_DEVICE_ONLY__ return __spirv_BuiltInSubgroupSize(); #else @@ -70,7 +72,7 @@ struct sub_group { #endif } - range_type get_max_local_range() const { + range_type get_max_local_range() const noexcept { #ifdef __SYCL_DEVICE_ONLY__ return __spirv_BuiltInSubgroupMaxSize(); #else @@ -79,7 +81,7 @@ struct sub_group { #endif } - id_type get_group_id() const { + id_type get_group_id() const noexcept { #ifdef __SYCL_DEVICE_ONLY__ return __spirv_BuiltInSubgroupId(); #else @@ -88,7 +90,7 @@ struct sub_group { #endif } - linear_id_type get_group_linear_id() const { + linear_id_type get_group_linear_id() const noexcept { #ifdef __SYCL_DEVICE_ONLY__ return static_cast(get_group_id()[0]); #else @@ -97,7 +99,7 @@ struct sub_group { #endif } - range_type get_group_range() const { + range_type get_group_range() const noexcept { #ifdef __SYCL_DEVICE_ONLY__ return __spirv_BuiltInNumSubgroups(); #else @@ -111,7 +113,7 @@ struct sub_group { __SYCL_DEPRECATED( "Sub-group barrier with no arguments is deprecated." "Use sycl::group_barrier with the sub-group as the argument instead.") - void barrier() const { + void barrier() const noexcept { #ifdef __SYCL_DEVICE_ONLY__ __spirv_ControlBarrier( __spv::Scope::Subgroup, __spv::Scope::Subgroup, @@ -128,7 +130,7 @@ struct sub_group { __SYCL_DEPRECATED( "Sub-group barrier accepting fence_space is deprecated." "Use sycl::group_barrier with the sub-group as the argument instead.") - void barrier(access::fence_space accessSpace) const { + void barrier(access::fence_space accessSpace) const noexcept { #ifdef __SYCL_DEVICE_ONLY__ int32_t flags = sycl::detail::getSPIRVMemorySemanticsMask(accessSpace); __spirv_ControlBarrier(__spv::Scope::Subgroup, __spv::Scope::Subgroup, @@ -141,7 +143,7 @@ struct sub_group { } #endif // __INTEL_PREVIEW_BREAKING_CHANGES - linear_id_type get_group_linear_range() const { + linear_id_type get_group_linear_range() const noexcept { #ifdef __SYCL_DEVICE_ONLY__ return static_cast(get_group_range()[0]); #else @@ -150,7 +152,7 @@ struct sub_group { #endif } - linear_id_type get_local_linear_range() const { + linear_id_type get_local_linear_range() const noexcept { #ifdef __SYCL_DEVICE_ONLY__ return static_cast(get_local_range()[0]); #else @@ -159,7 +161,7 @@ struct sub_group { #endif } - bool leader() const { + bool leader() const noexcept { #ifdef __SYCL_DEVICE_ONLY__ return get_local_linear_id() == 0; #else @@ -169,7 +171,7 @@ struct sub_group { } // Common member functions for by-value semantics - friend bool operator==(const sub_group &lhs, const sub_group &rhs) { + friend bool operator==(const sub_group &lhs, const sub_group &rhs) noexcept { #ifdef __SYCL_DEVICE_ONLY__ return lhs.get_group_id() == rhs.get_group_id(); #else @@ -180,7 +182,7 @@ struct sub_group { #endif } - friend bool operator!=(const sub_group &lhs, const sub_group &rhs) { + friend bool operator!=(const sub_group &lhs, const sub_group &rhs) noexcept { #ifdef __SYCL_DEVICE_ONLY__ return !(lhs == rhs); #else @@ -194,10 +196,11 @@ struct sub_group { protected: template friend class sycl::nd_item; friend sub_group ext::oneapi::this_work_item::get_sub_group(); - sub_group() = default; + sub_group() noexcept = default; }; -template sub_group nd_item::get_sub_group() const { +template +sub_group nd_item::get_sub_group() const noexcept { return sub_group(); } From f4188c3764dba2ce2645374f9fd3208f520ebdc5 Mon Sep 17 00:00:00 2001 From: Grzegorz Mianowski Date: Fri, 7 Aug 2026 10:02:57 +0200 Subject: [PATCH 2/2] Change throw to terminate --- sycl/include/sycl/group.hpp | 4 +-- sycl/include/sycl/sub_group.hpp | 44 +++++++++++---------------------- 2 files changed, 17 insertions(+), 31 deletions(-) diff --git a/sycl/include/sycl/group.hpp b/sycl/include/sycl/group.hpp index eb44645327914..33a1eb76720a7 100644 --- a/sycl/include/sycl/group.hpp +++ b/sycl/include/sycl/group.hpp @@ -25,6 +25,7 @@ #include // for range #ifndef __SYCL_DEVICE_ONLY__ +#include #include #include // for unique_ptr @@ -137,8 +138,7 @@ template class __SYCL_TYPE(group) group { #ifdef __SYCL_DEVICE_ONLY__ return __spirv::initBuiltInLocalInvocationId>(); #else - throw sycl::exception(make_error_code(errc::feature_not_supported), - "get_local_id() is not implemented on host"); + std::terminate(); #endif } diff --git a/sycl/include/sycl/sub_group.hpp b/sycl/include/sycl/sub_group.hpp index 677ec7e4b6858..16b9c7107a535 100644 --- a/sycl/include/sycl/sub_group.hpp +++ b/sycl/include/sycl/sub_group.hpp @@ -23,7 +23,7 @@ #include // for uint32_t #ifndef __SYCL_DEVICE_ONLY__ -#include // for exception, make_error... +#include // for terminate #endif namespace sycl { @@ -49,8 +49,7 @@ struct sub_group { #ifdef __SYCL_DEVICE_ONLY__ return __spirv_BuiltInSubgroupLocalInvocationId(); #else - throw sycl::exception(make_error_code(errc::feature_not_supported), - "Sub-groups are not supported on host."); + std::terminate(); #endif } @@ -58,8 +57,7 @@ struct sub_group { #ifdef __SYCL_DEVICE_ONLY__ return static_cast(get_local_id()[0]); #else - throw sycl::exception(make_error_code(errc::feature_not_supported), - "Sub-groups are not supported on host."); + std::terminate(); #endif } @@ -67,8 +65,7 @@ struct sub_group { #ifdef __SYCL_DEVICE_ONLY__ return __spirv_BuiltInSubgroupSize(); #else - throw sycl::exception(make_error_code(errc::feature_not_supported), - "Sub-groups are not supported on host."); + std::terminate(); #endif } @@ -76,8 +73,7 @@ struct sub_group { #ifdef __SYCL_DEVICE_ONLY__ return __spirv_BuiltInSubgroupMaxSize(); #else - throw sycl::exception(make_error_code(errc::feature_not_supported), - "Sub-groups are not supported on host."); + std::terminate(); #endif } @@ -85,8 +81,7 @@ struct sub_group { #ifdef __SYCL_DEVICE_ONLY__ return __spirv_BuiltInSubgroupId(); #else - throw sycl::exception(make_error_code(errc::feature_not_supported), - "Sub-groups are not supported on host."); + std::terminate(); #endif } @@ -94,8 +89,7 @@ struct sub_group { #ifdef __SYCL_DEVICE_ONLY__ return static_cast(get_group_id()[0]); #else - throw sycl::exception(make_error_code(errc::feature_not_supported), - "Sub-groups are not supported on host."); + std::terminate(); #endif } @@ -103,8 +97,7 @@ struct sub_group { #ifdef __SYCL_DEVICE_ONLY__ return __spirv_BuiltInNumSubgroups(); #else - throw sycl::exception(make_error_code(errc::feature_not_supported), - "Sub-groups are not supported on host."); + std::terminate(); #endif } @@ -122,8 +115,7 @@ struct sub_group { __spv::MemorySemanticsMask::WorkgroupMemory | __spv::MemorySemanticsMask::CrossWorkgroupMemory); #else - throw sycl::exception(make_error_code(errc::feature_not_supported), - "Sub-groups are not supported on host."); + std::terminate(); #endif } @@ -137,8 +129,7 @@ struct sub_group { flags); #else (void)accessSpace; - throw sycl::exception(make_error_code(errc::feature_not_supported), - "Sub-groups are not supported on host."); + std::terminate(); #endif } #endif // __INTEL_PREVIEW_BREAKING_CHANGES @@ -147,8 +138,7 @@ struct sub_group { #ifdef __SYCL_DEVICE_ONLY__ return static_cast(get_group_range()[0]); #else - throw sycl::exception(make_error_code(errc::feature_not_supported), - "Sub-groups are not supported on host."); + std::terminate(); #endif } @@ -156,8 +146,7 @@ struct sub_group { #ifdef __SYCL_DEVICE_ONLY__ return static_cast(get_local_range()[0]); #else - throw sycl::exception(make_error_code(errc::feature_not_supported), - "Sub-groups are not supported on host."); + std::terminate(); #endif } @@ -165,8 +154,7 @@ struct sub_group { #ifdef __SYCL_DEVICE_ONLY__ return get_local_linear_id() == 0; #else - throw sycl::exception(make_error_code(errc::feature_not_supported), - "Sub-groups are not supported on host."); + std::terminate(); #endif } @@ -177,8 +165,7 @@ struct sub_group { #else (void)lhs; (void)rhs; - throw sycl::exception(make_error_code(errc::feature_not_supported), - "Sub-groups are not supported on host."); + std::terminate(); #endif } @@ -188,8 +175,7 @@ struct sub_group { #else (void)lhs; (void)rhs; - throw sycl::exception(make_error_code(errc::feature_not_supported), - "Sub-groups are not supported on host."); + std::terminate(); #endif }