Skip to content

Commit 5e74e57

Browse files
AndreiZibrovAndreiZibrov
committed
SYCL] Fixed-size groups and partitions are renamed to "chunks"
1 parent fb1d33d commit 5e74e57

File tree

7 files changed

+32
-37
lines changed

7 files changed

+32
-37
lines changed

‎sycl/include/sycl/detail/spirv.hpp‎

Lines changed: 14 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -79,8 +79,8 @@ struct is_ballot_group<
7979
template <typename Group> struct is_chunk : std::false_type {};
8080

8181
template <size_t ChunkSize, typename ParentGroup>
82-
struct is_chunk<sycl::ext::oneapi::experimental::chunk<
83-
ChunkSize, ParentGroup>> : std::true_type {};
82+
struct is_chunk<sycl::ext::oneapi::experimental::chunk<ChunkSize, ParentGroup>>
83+
: std::true_type {};
8484

8585
template <typename Group> struct group_scope {};
8686

@@ -106,8 +106,8 @@ struct group_scope<sycl::ext::oneapi::experimental::ballot_group<ParentGroup>> {
106106
};
107107

108108
template <size_t ChunkSize, typename ParentGroup>
109-
struct group_scope<sycl::ext::oneapi::experimental::chunk<
110-
ChunkSize, ParentGroup>> {
109+
struct group_scope<
110+
sycl::ext::oneapi::experimental::chunk<ChunkSize, ParentGroup>> {
111111
static constexpr __spv::Scope::Flag value = group_scope<ParentGroup>::value;
112112
};
113113

@@ -175,9 +175,8 @@ bool GroupAll(ext::oneapi::experimental::ballot_group<ParentGroup> g,
175175
}
176176
}
177177
template <size_t ChunkSize, typename ParentGroup>
178-
bool GroupAll(
179-
ext::oneapi::experimental::chunk<ChunkSize, ParentGroup>,
180-
bool pred) {
178+
bool GroupAll(ext::oneapi::experimental::chunk<ChunkSize, ParentGroup>,
179+
bool pred) {
181180
// GroupNonUniformAll doesn't support cluster size, so use a reduction
182181
return __spirv_GroupNonUniformBitwiseAnd(
183182
group_scope<ParentGroup>::value,
@@ -211,9 +210,8 @@ bool GroupAny(ext::oneapi::experimental::ballot_group<ParentGroup> g,
211210
}
212211
}
213212
template <size_t ChunkSize, typename ParentGroup>
214-
bool GroupAny(
215-
ext::oneapi::experimental::chunk<ChunkSize, ParentGroup>,
216-
bool pred) {
213+
bool GroupAny(ext::oneapi::experimental::chunk<ChunkSize, ParentGroup>,
214+
bool pred) {
217215
// GroupNonUniformAny doesn't support cluster size, so use a reduction
218216
return __spirv_GroupNonUniformBitwiseOr(
219217
group_scope<ParentGroup>::value,
@@ -328,9 +326,9 @@ GroupBroadcast(sycl::ext::oneapi::experimental::ballot_group<ParentGroup> g,
328326
}
329327
}
330328
template <size_t ChunkSize, typename ParentGroup, typename T, typename IdT>
331-
EnableIfNativeBroadcast<T, IdT> GroupBroadcast(
332-
ext::oneapi::experimental::chunk<ChunkSize, ParentGroup> g,
333-
T x, IdT local_id) {
329+
EnableIfNativeBroadcast<T, IdT>
330+
GroupBroadcast(ext::oneapi::experimental::chunk<ChunkSize, ParentGroup> g, T x,
331+
IdT local_id) {
334332
// Remap local_id to its original numbering in ParentGroup
335333
auto LocalId = g.get_group_linear_id() * ChunkSize + local_id;
336334

@@ -1298,8 +1296,8 @@ ControlBarrier(Group g, memory_scope FenceScope, memory_order Order) {
12981296
} \
12991297
} \
13001298
\
1301-
template <__spv::GroupOperation Op, size_t ChunkSize, \
1302-
typename ParentGroup, typename T> \
1299+
template <__spv::GroupOperation Op, size_t ChunkSize, typename ParentGroup, \
1300+
typename T> \
13031301
inline T Group##Instruction( \
13041302
ext::oneapi::experimental::chunk<ChunkSize, ParentGroup> g, T x) { \
13051303
using ConvertedT = detail::ConvertToOpenCLType_t<T>; \
@@ -1319,7 +1317,7 @@ ControlBarrier(Group g, memory_scope FenceScope, memory_order Order) {
13191317
constexpr auto OpInt = \
13201318
static_cast<unsigned int>(__spv::GroupOperation::ClusteredReduce); \
13211319
return __spirv_GroupNonUniform##Instruction(Scope, OpInt, Arg, \
1322-
ChunkSize); \
1320+
ChunkSize); \
13231321
} else { \
13241322
T tmp; \
13251323
for (size_t Cluster = 0; Cluster < g.get_group_linear_range(); \

‎sycl/include/sycl/detail/type_traits.hpp‎

Lines changed: 2 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -22,8 +22,7 @@ inline namespace _V1 {
2222
namespace detail {
2323
template <class T> struct is_chunk : std::false_type {};
2424

25-
template <class T>
26-
inline constexpr bool is_chunk_v = is_chunk<T>::value;
25+
template <class T> inline constexpr bool is_chunk_v = is_chunk<T>::value;
2726

2827
template <typename VecT, typename OperationLeftT, typename OperationRightT,
2928
template <typename> class OperationCurrentT, int... Indexes>
@@ -157,7 +156,7 @@ template <typename T, int N> struct get_elem_type_unqual<vec<T, N>> {
157156
template <typename VecT, typename OperationLeftT, typename OperationRightT,
158157
template <typename> class OperationCurrentT, int... Indexes>
159158
struct get_elem_type_unqual<SwizzleOp<VecT, OperationLeftT, OperationRightT,
160-
OperationCurrentT, Indexes...>> {
159+
OperationCurrentT, Indexes...>> {
161160
using type = typename get_elem_type_unqual<std::remove_cv_t<VecT>>::type;
162161
};
163162

‎sycl/include/sycl/ext/oneapi/experimental/cuda/non_uniform_algorithms.hpp‎

Lines changed: 8 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -57,35 +57,35 @@ masked_reduction_cuda_sm80(Group g, T x, BinaryOperation binary_op,
5757
}
5858

5959
template <typename Group, typename T, class BinaryOperation>
60-
std::enable_if_t<(is_sugeninteger_v<T> ||
61-
is_sigeninteger_v<T>)&&IsPlus<T, BinaryOperation>::value,
60+
std::enable_if_t<(is_sugeninteger_v<T> || is_sigeninteger_v<T>) &&
61+
IsPlus<T, BinaryOperation>::value,
6262
T>
6363
masked_reduction_cuda_sm80(Group g, T x, BinaryOperation binary_op,
6464
const uint32_t MemberMask) {
6565
return __nvvm_redux_sync_add(x, MemberMask);
6666
}
6767

6868
template <typename Group, typename T, class BinaryOperation>
69-
std::enable_if_t<(is_sugeninteger_v<T> ||
70-
is_sigeninteger_v<T>)&&IsBitAND<T, BinaryOperation>::value,
69+
std::enable_if_t<(is_sugeninteger_v<T> || is_sigeninteger_v<T>) &&
70+
IsBitAND<T, BinaryOperation>::value,
7171
T>
7272
masked_reduction_cuda_sm80(Group g, T x, BinaryOperation binary_op,
7373
const uint32_t MemberMask) {
7474
return __nvvm_redux_sync_and(x, MemberMask);
7575
}
7676

7777
template <typename Group, typename T, class BinaryOperation>
78-
std::enable_if_t<(is_sugeninteger_v<T> ||
79-
is_sigeninteger_v<T>)&&IsBitOR<T, BinaryOperation>::value,
78+
std::enable_if_t<(is_sugeninteger_v<T> || is_sigeninteger_v<T>) &&
79+
IsBitOR<T, BinaryOperation>::value,
8080
T>
8181
masked_reduction_cuda_sm80(Group g, T x, BinaryOperation binary_op,
8282
const uint32_t MemberMask) {
8383
return __nvvm_redux_sync_or(x, MemberMask);
8484
}
8585

8686
template <typename Group, typename T, class BinaryOperation>
87-
std::enable_if_t<(is_sugeninteger_v<T> ||
88-
is_sigeninteger_v<T>)&&IsBitXOR<T, BinaryOperation>::value,
87+
std::enable_if_t<(is_sugeninteger_v<T> || is_sigeninteger_v<T>) &&
88+
IsBitXOR<T, BinaryOperation>::value,
8989
T>
9090
masked_reduction_cuda_sm80(Group g, T x, BinaryOperation binary_op,
9191
const uint32_t MemberMask) {

‎sycl/include/sycl/sycl.hpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -89,11 +89,11 @@
8989
#include <sycl/ext/oneapi/experimental/ballot_group.hpp>
9090
#include <sycl/ext/oneapi/experimental/bfloat16_math.hpp>
9191
#include <sycl/ext/oneapi/experimental/builtins.hpp>
92+
#include <sycl/ext/oneapi/experimental/chunk.hpp>
9293
#include <sycl/ext/oneapi/experimental/cluster_group_prop.hpp>
9394
#include <sycl/ext/oneapi/experimental/composite_device.hpp>
9495
#include <sycl/ext/oneapi/experimental/cuda/barrier.hpp>
9596
#include <sycl/ext/oneapi/experimental/enqueue_functions.hpp>
96-
#include <sycl/ext/oneapi/experimental/chunk.hpp>
9797
#include <sycl/ext/oneapi/experimental/forward_progress.hpp>
9898
#include <sycl/ext/oneapi/experimental/group_load_store.hpp>
9999
#include <sycl/ext/oneapi/experimental/group_sort.hpp>

‎sycl/test-e2e/NonUniformGroups/chunk.cpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -11,7 +11,7 @@
1111

1212
#include <vector>
1313

14-
//#ifdef __SYCL_DEVICE_ONLY__
14+
// #ifdef __SYCL_DEVICE_ONLY__
1515
//[[__sycl_detail__::__uses_aspects__(sycl::aspect::ext_oneapi_chunk)]]
1616

1717
#include <sycl/detail/core.hpp>
@@ -78,4 +78,4 @@ int main() {
7878
return 0;
7979
}
8080

81-
//# endif
81+
// #endif

‎sycl/test-e2e/NonUniformGroups/chunk_algorithms.cpp‎

Lines changed: 1 addition & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -14,7 +14,6 @@
1414
#include <sycl/group_barrier.hpp>
1515
#include <vector>
1616

17-
1817
namespace syclex = sycl::ext::oneapi::experimental;
1918

2019
template <size_t ChunkSize> class TestKernel;
@@ -76,8 +75,7 @@ template <size_t ChunkSize> void test() {
7675
uint32_t OriginalLID = SG.get_local_linear_id();
7776
uint32_t LID = Partition.get_local_linear_id();
7877

79-
uint32_t PartitionLeader =
80-
(OriginalLID / ChunkSize) * ChunkSize;
78+
uint32_t PartitionLeader = (OriginalLID / ChunkSize) * ChunkSize;
8179
uint32_t BroadcastResult =
8280
sycl::group_broadcast(Partition, OriginalLID, 0);
8381
BroadcastAcc[WI] = (BroadcastResult == PartitionLeader);

‎sycl/test-e2e/NonUniformGroups/is_user_constructed.cpp‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -9,10 +9,10 @@ namespace syclex = sycl::ext::oneapi::experimental;
99

1010
static_assert(
1111
syclex::is_user_constructed_group_v<syclex::ballot_group<sycl::sub_group>>);
12-
static_assert(syclex::is_user_constructed_group_v<
13-
syclex::chunk<1, sycl::sub_group>>);
14-
static_assert(syclex::is_user_constructed_group_v<
15-
syclex::chunk<2, sycl::sub_group>>);
12+
static_assert(
13+
syclex::is_user_constructed_group_v<syclex::chunk<1, sycl::sub_group>>);
14+
static_assert(
15+
syclex::is_user_constructed_group_v<syclex::chunk<2, sycl::sub_group>>);
1616
static_assert(
1717
syclex::is_user_constructed_group_v<syclex::tangle_group<sycl::sub_group>>);
1818
static_assert(syclex::is_user_constructed_group_v<syclex::opportunistic_group>);

0 commit comments

Comments
 (0)