@@ -109,8 +109,9 @@ template <typename Group> bool GroupAll(Group g, bool pred) {
109
109
template <typename ParentGroup>
110
110
bool GroupAll (ext::oneapi::experimental::ballot_group<ParentGroup> g,
111
111
bool pred) {
112
- // Each ballot_group implicitly represents two groups
113
- // We have to force each half down different control flow
112
+ // ballot_group partitions its parent into two groups (0 and 1)
113
+ // We have to force each group down different control flow
114
+ // Work-items in the "false" group (0) may still be active
114
115
if (g.get_group_id () == 1 ) {
115
116
return __spirv_GroupNonUniformAll (group_scope<ParentGroup>::value, pred);
116
117
} else {
@@ -124,8 +125,9 @@ template <typename Group> bool GroupAny(Group g, bool pred) {
124
125
template <typename ParentGroup>
125
126
bool GroupAny (ext::oneapi::experimental::ballot_group<ParentGroup> g,
126
127
bool pred) {
127
- // Each ballot_group implicitly represents two groups
128
- // We have to force each half down different control flow
128
+ // ballot_group partitions its parent into two groups (0 and 1)
129
+ // We have to force each group down different control flow
130
+ // Work-items in the "false" group (0) may still be active
129
131
if (g.get_group_id () == 1 ) {
130
132
return __spirv_GroupNonUniformAny (group_scope<ParentGroup>::value, pred);
131
133
} else {
@@ -214,8 +216,9 @@ GroupBroadcast(sycl::ext::oneapi::experimental::ballot_group<ParentGroup> g,
214
216
WidenedT OCLX = detail::convertDataToType<T, OCLT>(x);
215
217
OCLIdT OCLId = detail::convertDataToType<GroupIdT, OCLIdT>(GroupLocalId);
216
218
217
- // Each ballot_group implicitly represents two groups
218
- // We have to force each half down different control flow
219
+ // ballot_group partitions its parent into two groups (0 and 1)
220
+ // We have to force each group down different control flow
221
+ // Work-items in the "false" group (0) may still be active
219
222
if (g.get_group_id () == 1 ) {
220
223
return __spirv_GroupNonUniformBroadcast (group_scope<ParentGroup>::value,
221
224
OCLX, OCLId);
@@ -933,8 +936,9 @@ ControlBarrier(Group, memory_scope FenceScope, memory_order Order) {
933
936
std::is_same<ConvertedT, cl_ushort>(), \
934
937
cl_uint, ConvertedT>>; \
935
938
OCLT Arg = x; \
936
- /* Each ballot_group implicitly represents two groups */ \
937
- /* We have to force each half down different control flow */ \
939
+ /* ballot_group partitions its parent into two groups (0 and 1) */ \
940
+ /* We have to force each group down different control flow */ \
941
+ /* Work-items in the "false" group (0) may still be active */ \
938
942
constexpr auto Scope = group_scope<ParentGroup>::value; \
939
943
constexpr auto OpInt = static_cast <unsigned int >(Op); \
940
944
if (g.get_group_id () == 1 ) { \
0 commit comments