Paths
| Differential D116053
[MachineSink] Allow sinking of constant or ignorable physreg uses ClosedPublic Authored by vangthao on Dec 20 2021, 12:33 PM.
Details Summary For AMDGPU, any use of the physical register EXEC prevents sinking even if it is not a real physical register read. Add check to see if a physical Also perform same constant and ignorable physical register check when considering sinking in loops.
Diff Detail
Unit TestsFailed
Event TimelineHerald added subscribers: kerbowa, arphaman, hiraditya and 3 others. · View Herald TranscriptDec 20 2021, 12:33 PM Comment Actions I am not sure I can prove to myself this is legal. For example you are sinking a def into a loop with divergent condition and this def is used after the loop. Can this happen? If so a def might be done with an exec smaller than a use which creates an undef. Hoisting was OK because def was moved into a direction where exec is strictly not less than before. Did you run PSDB on it? Comment Actions
If SuccToSinkTo exists, all uses of the def must be dominated by block SuccToSinkTo. So I think it is ok to do the sinking. isConstantPhysReg() seems ok as any def to such register has no meaning. Better let an AMDGPU expert have a look especially at the code gen changes. Comment Actions
It's ok to sink a def with a smaller exec set as long as the use set is still a subset of the new def point exec Comment Actions
I believe the requirement for defs to dominate all uses prevents this from happening but I can add more test cases to check for this. This passed PSDB. Comment Actions
IR is essentially a single thread representation. The implicit exec use is our way to model mutithreaded divergence. Consider this transformation which shall now become legal: int lid = get_local_id(0); int lid = get_local_id(0);
int i = 0; int i = 0;
x = def(); do {
do { => x = def();
use1(x); use1(x);
} while(i++ < lid); } while(i++ < lid);
use2(x); use2(x);def dominates use2 in both cases, but in the second case not with every lane. All lanes except first will use an undef. Comment Actions
We will not sink into a loop if the def is outside of the loop. In the test case loop_sink_fmac, the def was already in a loop and was why it was able to be sinked. Comment Actions
If you are sure then it is OK. Note my example was with a loop which is guaranteed to execute at least one iteration. @arsenm ? This revision is now accepted and ready to land.Dec 22 2021, 11:28 AM Comment Actions This change seems to sink v_cmp instructions, which creates different results if the exec mask changed and that makes several Vulkan tests fail. Comment Actions
Does this help? diff --git a/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp b/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp index 88996f455227..0678ceeeea21 100644 --- a/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp +++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp @@ -130,10 +130,29 @@ bool SIInstrInfo::isReallyTriviallyReMaterializable(const MachineInstr &MI, return false; } +static bool readsExecAsData(const MachineInstr &MI) { + if (MI.isCompare()) + return true; + + switch (MI.getOpcode()) { + default: + break; + case AMDGPU::V_READFIRSTLANE_B32: + case AMDGPU::V_CNDMASK_B64_PSEUDO: + case AMDGPU::V_CNDMASK_B32_dpp: + case AMDGPU::V_CNDMASK_B32_e32: + case AMDGPU::V_CNDMASK_B32_e64: + case AMDGPU::V_CNDMASK_B32_sdwa: + return true; + } + + return false; +} + bool SIInstrInfo::isIgnorableUse(const MachineOperand &MO) const { // Any implicit use of exec by VALU is not a real register read. return MO.getReg() == AMDGPU::EXEC && MO.isImplicit() && - isVALU(*MO.getParent()); + isVALU(*MO.getParent()) && !readsExecAsData(*MO.getParent()); } bool SIInstrInfo::areLoadsFromSameBasePtr(SDNode *Load0, SDNode *Load1, Comment Actions Another, shorter, approach may be: diff --git a/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp b/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp index 88996f455227..f85a71941e66 100644 --- a/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp +++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp @@ -132,8 +132,10 @@ bool SIInstrInfo::isReallyTriviallyReMaterializable(const MachineInstr &MI, bool SIInstrInfo::isIgnorableUse(const MachineOperand &MO) const { // Any implicit use of exec by VALU is not a real register read. + // isRematerializable check excludes instructions reading EXEC as data, + // such as compares, v_cndmask_b32, and readfirstlane. return MO.getReg() == AMDGPU::EXEC && MO.isImplicit() && - isVALU(*MO.getParent()); + isVALU(*MO.getParent()) && MO.getParent()->isRematerializable(); } bool SIInstrInfo::areLoadsFromSameBasePtr(SDNode *Load0, SDNode *Load1,
Revision Contents
Diff 395511 llvm/include/llvm/CodeGen/TargetInstrInfo.h
llvm/lib/CodeGen/MachineSink.cpp
llvm/test/CodeGen/AMDGPU/GlobalISel/sdiv.i64.ll
llvm/test/CodeGen/AMDGPU/GlobalISel/srem.i64.ll
llvm/test/CodeGen/AMDGPU/GlobalISel/udiv.i64.ll
llvm/test/CodeGen/AMDGPU/GlobalISel/urem.i64.ll
llvm/test/CodeGen/AMDGPU/frame-index-elimination.ll
llvm/test/CodeGen/AMDGPU/machine-sink-ignorable-exec-use.mir
llvm/test/CodeGen/AMDGPU/mul24-pass-ordering.ll
llvm/test/CodeGen/AMDGPU/skip-if-dead.ll
llvm/test/CodeGen/AMDGPU/stack-pointer-offset-relative-frameindex.ll
| ||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||||