Skip to content

[SYCL][Reduction] Limit reduction work non-uniform case #8458

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 7 commits into from
Mar 7, 2023
Merged

[SYCL][Reduction] Limit reduction work non-uniform case #8458

Changes from 1 commit
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
82359b3
[SYCL][Reduction] Limit reduction work non-uniform case
steffenlarsen Feb 22, 2023
47e9a5d
Fix build failure
steffenlarsen Feb 24, 2023
f377c82
Code-review for 8458 (#1)
aelovikov-intel Mar 1, 2023
61ff406
Add tree-reduction guarantee enum
steffenlarsen Mar 1, 2023
7c2f448
Specify tuple tree-reduction requirement
steffenlarsen Mar 1, 2023
5e0a904
Make enum an enum class to avoid value name leaks
steffenlarsen Mar 2, 2023
1eed6c7
Remove barrier argument
steffenlarsen Mar 2, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
26 changes: 17 additions & 9 deletions sycl/include/sycl/reduction.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1200,13 +1200,18 @@ template <> struct NDRangeReduction<reduction::strategy::range_basic> {
// Reduce each result separately
// TODO: Opportunity to parallelize across elements
for (int E = 0; E < NElements; ++E) {
auto LocalSum = Identity;
for (size_t I = LID; I < NWorkGroups; I += WGSize)
LocalSum = BOp(LocalSum, PartialSums[I * NElements + E]);
// The last work-group may not have enough work for all its items.
size_t RemainingWorkSize = std::min(NWorkGroups, WGSize);

LocalReds[LID] = LocalSum;
if (LID < RemainingWorkSize) {
auto LocalSum = PartialSums[LID * NElements + E];
for (size_t I = LID + WGSize; I < NWorkGroups; I += WGSize)
LocalSum = BOp(LocalSum, PartialSums[I * NElements + E]);

doTreeReduction(WGSize, LID, LocalReds, BOp,
LocalReds[LID] = LocalSum;
}

doTreeReduction(RemainingWorkSize, LID, LocalReds, BOp,
[&]() { workGroupBarrier(); });
if (LID == 0) {
auto V = LocalReds[0];
Expand Down Expand Up @@ -1571,13 +1576,16 @@ template <> struct NDRangeReduction<reduction::strategy::basic> {
size_t GID = NDIt.get_global_linear_id();

for (int E = 0; E < NElements; ++E) {
// The last work-group may not have enough work for all its items.
size_t RemainingWorkSize =
sycl::min(WGSize, NWorkItems - GrID * WGSize);

// Copy the element to local memory to prepare it for
// tree-reduction.
LocalReds[LID] = (UniformPow2WG || GID < NWorkItems)
? In[GID * NElements + E]
: ReduIdentity;
if (LID < RemainingWorkSize)
LocalReds[LID] = In[GID * NElements + E];

doTreeReduction(WGSize, LID, LocalReds, BOp,
doTreeReduction(RemainingWorkSize, LID, LocalReds, BOp,
[&]() { NDIt.barrier(); });

// Compute the partial sum/reduction for the work-group.
Expand Down