Skip to content

[SYCL] Consistently using unsigned loop counter for reduction #10795

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
steffenlarsen merged 2 commits into from
Aug 21, 2023
Merged

[SYCL] Consistently using unsigned loop counter for reduction #10795

Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
57eb53a
Consistently use unsigned counter in loop over elements
reble Aug 11, 2023
7d6fcbb
add missing cases
reble Aug 11, 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: 13 additions & 13 deletions sycl/include/sycl/reduction.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -990,7 +990,7 @@ class reduction_impl_algo {
// instantiated for the code below.
auto DoIt = [&](auto &Out) {
auto RWReduVal = std::make_shared<std::array<T, num_elements>>();
for (int i = 0; i < num_elements; ++i) {
for (size_t i = 0; i < num_elements; ++i) {
(*RWReduVal)[i] = decltype(MIdentityContainer)::getIdentity();
}
auto Buf = std::make_shared<buffer<T, 1>>(RWReduVal.get()->data(),
Expand Down Expand Up @@ -1021,7 +1021,7 @@ class reduction_impl_algo {
size_t NElements = num_elements;

CopyHandler.single_task<__sycl_init_mem_for<KernelName>>([=] {
for (int i = 0; i < NElements; ++i) {
for (size_t i = 0; i < NElements; ++i) {
if (IsUpdateOfUserVar)
Out[i] = BOp(Out[i], Mem[i]);
else
Expand Down Expand Up @@ -1205,7 +1205,7 @@ void reduSaveFinalResultToUserMem(handler &CGH, Reduction &Redu) {
bool IsUpdateOfUserVar = !Redu.initializeToIdentity();
auto BOp = Redu.getBinaryOperation();
CGH.single_task<KernelName>([=] {
for (int i = 0; i < NElements; ++i) {
for (size_t i = 0; i < NElements; ++i) {
auto Elem = InAcc[i];
if (IsUpdateOfUserVar)
UserVarPtr[i] = BOp(UserVarPtr[i], *Elem);
Expand Down Expand Up @@ -1328,7 +1328,7 @@ struct NDRangeReduction<
// If there are multiple values, reduce each separately
// reduce_over_group is only defined for each T, not for span<T, ...>
size_t LID = NDId.get_local_id(0);
for (int E = 0; E < NElements; ++E) {
for (size_t E = 0; E < NElements; ++E) {
auto &RedElem = *getReducerAccess(Reducer).getElement(E);
RedElem = reduce_over_group(Group, RedElem, BOp);
if (LID == 0) {
Expand Down Expand Up @@ -1362,7 +1362,7 @@ struct NDRangeReduction<
if (DoReducePartialSumsInLastWG[0]) {
// Reduce each result separately
// TODO: Opportunity to parallelize across elements.
for (int E = 0; E < NElements; ++E) {
for (size_t E = 0; E < NElements; ++E) {
auto LocalSum = getReducerAccess(Reducer).getIdentity();
for (size_t I = LID; I < NWorkGroups; I += WGSize)
LocalSum = BOp(LocalSum, PartialSums[I * NElements + E]);
Expand Down Expand Up @@ -1538,7 +1538,7 @@ template <> struct NDRangeReduction<reduction::strategy::range_basic> {
// If there are multiple values, reduce each separately
// This prevents local memory from scaling with elements
size_t LID = NDId.get_local_linear_id();
for (int E = 0; E < NElements; ++E) {
for (size_t E = 0; E < NElements; ++E) {

doTreeReduction<WorkSizeGuarantees::Equal>(
WGSize, NDId, LocalReds, ElementCombiner,
Expand Down Expand Up @@ -1574,7 +1574,7 @@ template <> struct NDRangeReduction<reduction::strategy::range_basic> {
if (DoReducePartialSumsInLastWG[0]) {
// Reduce each result separately
// TODO: Opportunity to parallelize across elements
for (int E = 0; E < NElements; ++E) {
for (size_t E = 0; E < NElements; ++E) {
doTreeReduction<WorkSizeGuarantees::None>(
NWorkGroups, NDId, LocalReds, ElementCombiner,
[&](size_t I) { return PartialSums[I * NElements + E]; });
Expand Down Expand Up @@ -1614,7 +1614,7 @@ struct NDRangeReduction<reduction::strategy::group_reduce_and_atomic_cross_wg> {
KernelFunc(NDIt, Reducer);

typename Reduction::binary_operation BOp;
for (int E = 0; E < NElements; ++E) {
for (size_t E = 0; E < NElements; ++E) {
auto &ReducerElem = getReducerAccess(Reducer).getElement(E);
*ReducerElem = reduce_over_group(NDIt.get_group(), *ReducerElem, BOp);
}
Expand Down Expand Up @@ -1663,7 +1663,7 @@ struct NDRangeReduction<

// If there are multiple values, reduce each separately
// This prevents local memory from scaling with elements
for (int E = 0; E < NElements; ++E) {
for (size_t E = 0; E < NElements; ++E) {

doTreeReduction<WorkSizeGuarantees::Equal>(
WGSize, NDIt, LocalReds, ElementCombiner,
Expand Down Expand Up @@ -1738,7 +1738,7 @@ struct NDRangeReduction<
// Compute the partial sum/reduction for the work-group.
size_t WGID = NDIt.get_group_linear_id();
typename Reduction::binary_operation BOp;
for (int E = 0; E < NElements; ++E) {
for (size_t E = 0; E < NElements; ++E) {
typename Reduction::result_type PSum;
PSum = *getReducerAccess(Reducer).getElement(E);
PSum = reduce_over_group(NDIt.get_group(), PSum, BOp);
Expand Down Expand Up @@ -1800,7 +1800,7 @@ struct NDRangeReduction<
size_t WGID = NDIt.get_group_linear_id();
size_t GID = NDIt.get_global_linear_id();

for (int E = 0; E < NElements; ++E) {
for (size_t E = 0; E < NElements; ++E) {
typename Reduction::result_type PSum =
(HasUniformWG || (GID < NWorkItems))
? *In[GID * NElements + E]
Expand Down Expand Up @@ -1897,7 +1897,7 @@ template <> struct NDRangeReduction<reduction::strategy::basic> {

// If there are multiple values, reduce each separately
// This prevents local memory from scaling with elements
for (int E = 0; E < NElements; ++E) {
for (size_t E = 0; E < NElements; ++E) {

doTreeReduction<WorkSizeGuarantees::Equal>(
WGSize, NDIt, LocalReds, ElementCombiner,
Expand Down Expand Up @@ -2004,7 +2004,7 @@ template <> struct NDRangeReduction<reduction::strategy::basic> {
return LHS.combine(BOp, RHS);
};

for (int E = 0; E < NElements; ++E) {
for (size_t E = 0; E < NElements; ++E) {

doTreeReduction<WorkSizeGuarantees::LessOrEqual>(
RemainingWorkSize, NDIt, LocalReds, ElementCombiner,
Expand Down