Skip to content

Commit 2744b41

Browse files
[SYCL][Reduction] Update a comment (#11143)
It wasn't obvious why we don't have a WG barrier before reporting WG has finished its work. Clarify that in the comment now.
1 parent 54cd00d commit 2744b41

File tree

1 file changed

+8
-2
lines changed

1 file changed

+8
-2
lines changed

sycl/include/sycl/reduction.hpp

Lines changed: 8 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1349,7 +1349,10 @@ struct NDRangeReduction<
13491349
// We're done.
13501350
return;
13511351

1352-
// Signal this work-group has finished after all values are reduced
1352+
// Signal this work-group has finished after all values are reduced. We
1353+
// had an implicit work-group barrier in reduce_over_group and all the
1354+
// work since has been done in (LID == 0) work-item, so no extra sync is
1355+
// needed.
13531356
if (LID == 0) {
13541357
auto NFinished =
13551358
sycl::atomic_ref<int, memory_order::acq_rel, memory_scope::device,
@@ -1561,7 +1564,10 @@ template <> struct NDRangeReduction<reduction::strategy::range_basic> {
15611564
}
15621565
}
15631566

1564-
// Signal this work-group has finished after all values are reduced
1567+
// Signal this work-group has finished after all values are reduced. We
1568+
// had an implicit work-group barrier in doTreeReduction and all the
1569+
// work since has been done in (LID == 0) work-item, so no extra sync is
1570+
// needed.
15651571
if (LID == 0) {
15661572
auto NFinished =
15671573
sycl::atomic_ref<int, memory_order::acq_rel, memory_scope::device,

0 commit comments

Comments
 (0)