@@ -109,8 +109,9 @@ template <typename Group> bool GroupAll(Group g, bool pred) {
109109template <typename ParentGroup>
110110bool GroupAll (ext::oneapi::experimental::ballot_group<ParentGroup> g,
111111 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
114115 if (g.get_group_id () == 1 ) {
115116 return __spirv_GroupNonUniformAll (group_scope<ParentGroup>::value, pred);
116117 } else {
@@ -124,8 +125,9 @@ template <typename Group> bool GroupAny(Group g, bool pred) {
124125template <typename ParentGroup>
125126bool GroupAny (ext::oneapi::experimental::ballot_group<ParentGroup> g,
126127 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
129131 if (g.get_group_id () == 1 ) {
130132 return __spirv_GroupNonUniformAny (group_scope<ParentGroup>::value, pred);
131133 } else {
@@ -214,8 +216,9 @@ GroupBroadcast(sycl::ext::oneapi::experimental::ballot_group<ParentGroup> g,
214216 WidenedT OCLX = detail::convertDataToType<T, OCLT>(x);
215217 OCLIdT OCLId = detail::convertDataToType<GroupIdT, OCLIdT>(GroupLocalId);
216218
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
219222 if (g.get_group_id () == 1 ) {
220223 return __spirv_GroupNonUniformBroadcast (group_scope<ParentGroup>::value,
221224 OCLX, OCLId);
@@ -933,8 +936,9 @@ ControlBarrier(Group, memory_scope FenceScope, memory_order Order) {
933936 std::is_same<ConvertedT, cl_ushort>(), \
934937 cl_uint, ConvertedT>>; \
935938 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 */ \
938942 constexpr auto Scope = group_scope<ParentGroup>::value; \
939943 constexpr auto OpInt = static_cast <unsigned int >(Op); \
940944 if (g.get_group_id () == 1 ) { \
0 commit comments