23 #include <type_traits>
26 inline namespace _V1 {
27 namespace ext::oneapi::experimental {
29 template <
size_t PartitionSize,
typename ParentGroup>
class fixed_size_group;
31 template <
size_t PartitionSize,
typename Group>
32 #ifdef __SYCL_DEVICE_ONLY__
33 [[__sycl_detail__::__uses_aspects__(sycl::aspect::ext_oneapi_fixed_size_group)]]
35 inline std::enable_if_t<sycl::is_group_v<std::decay_t<Group>> &&
36 std::is_same_v<Group, sycl::sub_group>,
37 fixed_size_group<PartitionSize, Group>>
49 #ifdef __SYCL_DEVICE_ONLY__
50 return __spirv_SubgroupLocalInvocationId() / PartitionSize;
53 "Non-uniform groups are not supported on host.");
58 #ifdef __SYCL_DEVICE_ONLY__
59 return __spirv_SubgroupLocalInvocationId() % PartitionSize;
62 "Non-uniform groups are not supported on host.");
67 #ifdef __SYCL_DEVICE_ONLY__
68 return __spirv_SubgroupSize() / PartitionSize;
71 "Non-uniform groups are not supported on host.");
76 #ifdef __SYCL_DEVICE_ONLY__
80 "Non-uniform groups are not supported on host.");
85 #ifdef __SYCL_DEVICE_ONLY__
89 "Non-uniform groups are not supported on host.");
94 #ifdef __SYCL_DEVICE_ONLY__
98 "Non-uniform groups are not supported on host.");
103 #ifdef __SYCL_DEVICE_ONLY__
107 "Non-uniform groups are not supported on host.");
112 #ifdef __SYCL_DEVICE_ONLY__
116 "Non-uniform groups are not supported on host.");
121 #ifdef __SYCL_DEVICE_ONLY__
125 "Non-uniform groups are not supported on host.");
130 #if defined(__SYCL_DEVICE_ONLY__) && defined(__NVPTX__)
134 #if defined(__SYCL_DEVICE_ONLY__) && defined(__NVPTX__)
141 get_fixed_size_group<PartitionSize, ParentGroup>(ParentGroup g);
144 sycl::detail::GetMask<fixed_size_group<PartitionSize, ParentGroup>>(
148 template <
size_t PartitionSize,
typename Group>
149 inline std::enable_if_t<sycl::is_group_v<std::decay_t<Group>> &&
150 std::is_same_v<Group, sycl::sub_group>,
154 #ifdef __SYCL_DEVICE_ONLY__
155 #if defined(__NVPTX__)
156 uint32_t loc_id =
group.get_local_linear_id();
157 uint32_t loc_size =
group.get_local_linear_range();
158 uint32_t bits = PartitionSize == 32
160 : ((1 << PartitionSize) - 1)
161 << ((loc_id / PartitionSize) * PartitionSize);
164 sycl::detail::Builder::createSubGroupMask<ext::oneapi::sub_group_mask>(
171 "Non-uniform groups are not supported on host.");
175 template <
size_t PartitionSize,
typename ParentGroup>
182 template <
size_t PartitionSize,
typename ParentGroup>
184 ext::oneapi::experimental::fixed_size_group<PartitionSize, ParentGroup>>
188 template <
size_t PartitionSize,
typename ParentGroup>
190 ext::oneapi::experimental::fixed_size_group<PartitionSize, ParentGroup>>
static constexpr sycl::memory_scope fence_scope
linear_id_type get_group_linear_range() const
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_local_id() const
typename ParentGroup::linear_id_type linear_id_type
id_type get_group_id() const
static constexpr int dimensions
linear_id_type get_local_linear_id() const
range_type get_group_range() const
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 >, fixed_size_group< PartitionSize, Group > > get_fixed_size_group(Group group)
std::error_code make_error_code(sycl::errc E) noexcept
Constructs an error code using e and sycl_category()