github/ROCm
1
1
Fork 0
mirror of https://github.com/ROCm/ROCm.git synced 2026-04-05 03:01:17 -04:00
Code Issues Packages Projects Releases Wiki Activity

ROCM IFU: Fix getValueLivenessRange

Browse Source
This commit is contained in:
Jason Furmanek
2023-09-28 21:35:43 +00:00
parent 28c571ea43
commit 92edee723b
1 changed files with 6 additions and 4 deletions
Download Patch File Download Diff File Expand all files Collapse all files

10
lib/Analysis/Allocation.cpp
View File

@@ -507,14 +507,16 @@ private:
// Analyze liveness of explicit buffers
Liveness liveness(operation);
auto getValueLivenessRange = [&](Value value) {
LivenessR ranges;
// Shared memory allocated by mbarrier cannot be reused
if (value.getDefiningOp() &&
isa<triton::nvidia_gpu::AllocMBarrierOp>(value.getDefiningOp()))
return Interval(std::numeric_limits<size_t>::min(),
std::numeric_limits<size_t>::max());
isa<triton::nvidia_gpu::AllocMBarrierOp>(value.getDefiningOp())) {
ranges.push_back(Interval(std::numeric_limits<size_t>::min(),
std::numeric_limits<size_t>::max()));
return ranges;
}
auto liveOperations = liveness.resolveLiveness(value);
LivenessR ranges;
std::for_each(liveOperations.begin(), liveOperations.end(),
[&](Operation *liveOp) {
ranges.add(operationId[liveOp]);