Skip to content

[SYCL] Align group and sub_group with SYCL 2020 #22888

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

Draft
mianowski wants to merge 2 commits into
base: sycl
Choose a base branch
from
+106 −105
Draft

[SYCL] Align group and sub_group with SYCL 2020 #22888

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
68e6c7f
Add noexcept to group and sub_group methods. Equality operators chang…
mianowski Aug 6, 2026
f4188c3
Change throw to terminate
mianowski Aug 7, 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
126 changes: 69 additions & 57 deletions sycl/include/sycl/group.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -25,6 +25,7 @@
#include <sycl/range.hpp> // for range

#ifndef __SYCL_DEVICE_ONLY__
#include <exception>
#include <sycl/exception.hpp>

#include <memory> // for unique_ptr
Expand Down Expand Up @@ -116,66 +117,73 @@ template <int Dimensions = 1> class __SYCL_TYPE(group) group {
group() = delete;

__SYCL2020_DEPRECATED("use sycl::group::get_group_id() instead")
id<Dimensions> get_id() const { return index; }
id<Dimensions> 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<Dimensions> get_group_id() const { return index; }
id<Dimensions> 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<Dimensions> get_global_range() const { return globalRange; }
range<Dimensions> 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<Dimensions> get_local_id() const {
id<Dimensions> get_local_id() const noexcept {
#ifdef __SYCL_DEVICE_ONLY__
return __spirv::initBuiltInLocalInvocationId<Dimensions, id<Dimensions>>();
#else
throw sycl::exception(make_error_code(errc::feature_not_supported),
"get_local_id() is not implemented on host");
std::terminate();
#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<Dimensions>();
}

range<Dimensions> get_local_range() const { return localRange; }
range<Dimensions> 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<Dimensions> get_group_range() const { return groupRange; }
range<Dimensions> 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<Dimensions> get_max_local_range() const { return get_local_range(); }
range<Dimensions> 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 &,
Expand All @@ -188,7 +196,7 @@ template <int Dimensions = 1> 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();
Expand Down Expand Up @@ -244,7 +252,7 @@ template <int Dimensions = 1> class __SYCL_TYPE(group) group {
#endif
void
parallel_for_work_item(range<Dimensions> flexibleRange,
WorkItemFunctionT Func) const {
WorkItemFunctionT Func) const noexcept {
detail::workGroupBarrier();
#ifdef __SYCL_DEVICE_ONLY__
range<Dimensions> GlobalSize{
Expand Down Expand Up @@ -308,7 +316,7 @@ template <int Dimensions = 1> 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
Expand Down Expand Up @@ -339,7 +347,7 @@ template <int Dimensions = 1> 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),
Expand All @@ -365,7 +373,7 @@ template <int Dimensions = 1> class __SYCL_TYPE(group) group {
[[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 @@ -390,7 +398,7 @@ template <int Dimensions = 1> class __SYCL_TYPE(group) group {
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 @@ -415,7 +423,7 @@ template <int Dimensions = 1> class __SYCL_TYPE(group) group {
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 @@ -443,7 +451,7 @@ template <int Dimensions = 1> 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<uint8_t, DestS, access::decorated::legacy>(
Expand All @@ -469,7 +477,7 @@ template <int Dimensions = 1> 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<T, uint8_t>;
Expand All @@ -492,7 +500,7 @@ template <int Dimensions = 1> class __SYCL_TYPE(group) group {
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 @@ -520,7 +528,7 @@ template <int Dimensions = 1> class __SYCL_TYPE(group) group {
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 @@ -546,7 +554,7 @@ template <int Dimensions = 1> class __SYCL_TYPE(group) group {
__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 @@ -559,7 +567,7 @@ template <int Dimensions = 1> class __SYCL_TYPE(group) group {
__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 @@ -574,7 +582,7 @@ template <int Dimensions = 1> class __SYCL_TYPE(group) group {
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 @@ -589,24 +597,28 @@ template <int Dimensions = 1> class __SYCL_TYPE(group) group {
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(), ...);
}

bool operator==(const group<Dimensions> &rhs) const {
bool Result = (rhs.globalRange == globalRange) &&
(rhs.localRange == localRange) && (rhs.index == index);
__SYCL_ASSERT(rhs.groupRange == groupRange &&
friend bool operator==(const group<Dimensions> &lhs,
const group<Dimensions> &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<Dimensions> &rhs) const {
return !((*this) == rhs);
friend bool operator!=(const group<Dimensions> &lhs,
const group<Dimensions> &rhs) noexcept {
return !(lhs == rhs);
}

private:
Expand All @@ -617,77 +629,77 @@ template <int Dimensions = 1> class __SYCL_TYPE(group) group {

template <int dims = Dimensions>
typename std::enable_if_t<(dims == 1), size_t>
get_local_linear_id_impl() const {
get_local_linear_id_impl() const noexcept {
id<Dimensions> localId = get_local_id();
return localId[0];
}

template <int dims = Dimensions>
typename std::enable_if_t<(dims == 2), size_t>
get_local_linear_id_impl() const {
get_local_linear_id_impl() const noexcept {
id<Dimensions> localId = get_local_id();
return localId[0] * localRange[1] + localId[1];
}

template <int dims = Dimensions>
typename std::enable_if_t<(dims == 3), size_t>
get_local_linear_id_impl() const {
get_local_linear_id_impl() const noexcept {
id<Dimensions> localId = get_local_id();
return (localId[0] * localRange[1] * localRange[2]) +
(localId[1] * localRange[2]) + localId[2];
}

template <int dims = Dimensions>
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 <int dims = Dimensions>
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 <int dims = Dimensions>
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 <int dims = Dimensions>
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 <int dims = Dimensions>
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 <int dims = Dimensions>
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 <int dims = Dimensions>
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 <int dims = Dimensions>
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];
}

Expand All @@ -703,15 +715,15 @@ template <int Dimensions = 1> class __SYCL_TYPE(group) group {
// work-group id from a multi-dimensional index follows the equation 4.3.
template <int dims = Dimensions>
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];
}

protected:
friend class detail::Builder;
group(const range<Dimensions> &G, const range<Dimensions> &L,
const range<Dimensions> GroupRange, const id<Dimensions> &I)
const range<Dimensions> GroupRange, const id<Dimensions> &I) noexcept
: globalRange(G), localRange(L), groupRange(GroupRange), index(I) {}
};
} // namespace _V1
Expand Down
2 changes: 1 addition & 1 deletion sycl/include/sycl/nd_item.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -116,7 +116,7 @@ template <int Dimensions = 1> 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];
Expand Down
Loading
Loading