22 #include <type_traits>
25 inline namespace _V1 {
26 namespace ext::oneapi::experimental {
28 template <
typename ParentGroup>
class ballot_group;
30 template <
typename Group>
31 #ifdef __SYCL_DEVICE_ONLY__
32 [[__sycl_detail__::__uses_aspects__(sycl::aspect::ext_oneapi_ballot_group)]]
34 inline std::enable_if_t<sycl::is_group_v<std::decay_t<Group>> &&
35 std::is_same_v<Group, sycl::sub_group>,
48 #ifdef __SYCL_DEVICE_ONLY__
51 throw runtime_error(
"Non-uniform groups are not supported on host device.",
52 PI_ERROR_INVALID_DEVICE);
57 #ifdef __SYCL_DEVICE_ONLY__
58 return sycl::detail::CallerPositionInMask(
Mask);
60 throw runtime_error(
"Non-uniform groups are not supported on host device.",
61 PI_ERROR_INVALID_DEVICE);
66 #ifdef __SYCL_DEVICE_ONLY__
69 throw runtime_error(
"Non-uniform groups are not supported on host device.",
70 PI_ERROR_INVALID_DEVICE);
75 #ifdef __SYCL_DEVICE_ONLY__
78 throw runtime_error(
"Non-uniform groups are not supported on host device.",
79 PI_ERROR_INVALID_DEVICE);
84 #ifdef __SYCL_DEVICE_ONLY__
87 throw runtime_error(
"Non-uniform groups are not supported on host device.",
88 PI_ERROR_INVALID_DEVICE);
93 #ifdef __SYCL_DEVICE_ONLY__
96 throw runtime_error(
"Non-uniform groups are not supported on host device.",
97 PI_ERROR_INVALID_DEVICE);
102 #ifdef __SYCL_DEVICE_ONLY__
105 throw runtime_error(
"Non-uniform groups are not supported on host device.",
106 PI_ERROR_INVALID_DEVICE);
111 #ifdef __SYCL_DEVICE_ONLY__
114 throw runtime_error(
"Non-uniform groups are not supported on host device.",
115 PI_ERROR_INVALID_DEVICE);
120 #ifdef __SYCL_DEVICE_ONLY__
121 uint32_t Lowest =
static_cast<uint32_t
>(
Mask.
find_low()[0]);
122 return __spirv_SubgroupLocalInvocationId() == Lowest;
124 throw runtime_error(
"Non-uniform groups are not supported on host device.",
125 PI_ERROR_INVALID_DEVICE);
136 get_ballot_group<ParentGroup>(ParentGroup g,
bool predicate);
142 template <
typename Group>
143 inline std::enable_if_t<sycl::is_group_v<std::decay_t<Group>> &&
144 std::is_same_v<Group, sycl::sub_group>,
148 #ifdef __SYCL_DEVICE_ONLY__
149 #if defined(__SPIR__) || defined(__NVPTX__)
167 throw runtime_error(
"Non-uniform groups are not supported on host device.",
168 PI_ERROR_INVALID_DEVICE);
172 template <
typename ParentGroup>
177 template <
typename ParentGroup>
178 struct is_group<ext::oneapi::experimental::ballot_group<ParentGroup>>
range_type get_group_range() const
id_type get_local_id() const
ballot_group(sub_group_mask m, bool p)
const sub_group_mask Mask
static constexpr sycl::memory_scope fence_scope
linear_id_type get_group_linear_id() const
range_type get_local_range() const
linear_id_type get_local_linear_range() const
id_type get_group_id() const
linear_id_type get_group_linear_range() const
linear_id_type get_local_linear_id() const
static constexpr int dimensions
typename ParentGroup::linear_id_type linear_id_type
fence_scope
The scope that fence() operation should apply to.
std::enable_if_t< sycl::is_group_v< std::decay_t< Group > > &&std::is_same_v< Group, sycl::sub_group >, ballot_group< Group > > get_ballot_group(Group group, bool predicate)
std::enable_if_t< std::is_same_v< std::decay_t< Group >, sub_group >||std::is_same_v< std::decay_t< Group >, sycl::sub_group >, sub_group_mask > group_ballot(Group g, bool predicate=true)
static constexpr size_t max_bits