Skip to content
This repository was archived by the owner on Apr 28, 2023. It is now read-only.

Commit 04a4377

Browse files
author
Sven Verdoolaege
committed
MappedScop::mapToThreads: fix depth of threadIdxX in case of reduction
Commit 88764d3 (split out reduction band after thread mapping to remove nested mapping, Mon Apr 9 17:43:18 2018 +0200) removed the condition that the reduction member always appears first in its band at the point of MappedScop::mapToThreads. The actual mapping was updated to take into account the depth inside the band, but the threadIdxXScheduleDepthState was not updated accordingly.
1 parent 1371777 commit 04a4377

File tree

1 file changed

+2
-2
lines changed

1 file changed

+2
-2
lines changed

tc/core/polyhedral/cuda/mapped_scop.cc

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -334,10 +334,10 @@ size_t MappedScop::mapToThreads(detail::ScheduleTree* band) {
334334
if (reductionBandUpdates_.count(band) == 1) {
335335
// A reduction is assumed to get mapped to threadIdx.x
336336
CHECK(reductionBandUpdates_.at(band).separated);
337+
auto reductionDim = reductionBandUpdates_.at(band).reductionDim;
337338
threadIdxXScheduleDepthState.emplace_back(std::make_pair(
338339
activeDomainPoints(schedule(), band),
339-
band->scheduleDepth(schedule()) + 0));
340-
auto reductionDim = reductionBandUpdates_.at(band).reductionDim;
340+
band->scheduleDepth(schedule()) + reductionDim));
341341
band = map(band, reductionDim, mapping::ThreadId::x());
342342
nMappedReductionThreads = 1;
343343
}

0 commit comments

Comments
 (0)