[AMDGPU] Add llvm.amdgcn.wqm.demote intrinsic and live mask tracking

[critson]

Summary

Add intrinsic which demotes all active lanes to helper lanes.
This is used to implement demote to helper Vulkan extension.

In practice demoting a lane to helper simply means removing it
from the mask of live lanes used for WQM/WWM/Exact mode.
Where the shader does not use WQM, demotes just become kills.

To support demote live lanes must now be tracked through entire
shader. Move implementation of kill intrinsics to WQM pass
and use live lane tracking to enable early termination of
shaders are any point in control flow.

Additionally add llvm.amdgcn.live.mask intrinsic to complement
demote operations. In theory llvm.amdgcn.ps.live can be used
to detect helper lanes; however, ps.live can be moved by LICM.
The movement of ps.live cannot be remedied without changing
its type signature and such a change would require ps.live
users to update as well.

[arsenm]

The WQM pass should be moved to avoid recomputing these. Computing LiveIntervals twice is an unacceptable cost

[critson]

Update to move comprehensive implementation now fairly tested
for graphics.
Now preserves LiveInterval analysis.
This is dependent on D78417.

[arsenm]

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

1243

I assume this needs to be convergent

1458

Ditto, and nomem

llvm/lib/Target/AMDGPU/AMDGPUSearchableTables.td

101

Missing DivergenceAnalysis test

llvm/lib/Target/AMDGPU/SIWholeQuadMode.cpp

763–764

use Register

[critson]

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

1243

Could be, my understanding is that without flags the intrinsic is marked "has side effects", which is correct as then it will not moved by LICM or removed by CSE.

1458

Convergent maybe (as above), but not nomem.
If this is marked nomem then it will be eaten by early CSE.
Since this was modelled on kill, is there a reason we don't mark kill Convergent?

llvm/lib/Target/AMDGPU/SIWholeQuadMode.cpp

763–764

I assume you mean the variable name Src -> Register?

[critson]

Add DivergenceAnalysis test.
Mark intrinsics as convergent.
Rename a variable.

[nhaehnle]

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

1243

Rethinking this, convergent isn't correct here, because there is no implied cross-thread communication.

Rather, the semantics are that amdgcn.wqm.helper reads from some hidden memory that is written to by wqm.demote. So by that logic, this should arguably be ReadInaccessibleMemOnly.

1458

Following the logic above, this should not be convergent but WritesInaccessibleMemOnly. At least I *think* that captures the semantics most accurately.

[critson]

Update intrinsic parameters:
llvm.amdgcn.wqm.demote - IntrWriteMem, IntrInaccessibleMemOnly
llvm.amdgcn.wqm.helper - IntrReadMem, IntrInaccessibleMemOnly

[critson]

Ping

[critson]

Ping - I would like to get this moving again.

[nhaehnle]

In D67767#2103077, @critson wrote:

Ping - I would like to get this moving again.

Does it make sense to wait until the WQM pass is moved?

[critson]

In D67767#2108753, @nhaehnle wrote:

In D67767#2103077, @critson wrote:

Ping - I would like to get this moving again.

Does it make sense to wait until the WQM pass is moved?

Sure, I will revise this in light of D82190 and D82192.

[arsenm]

Can you add globalisel tests? I think you may need to add the operand class to the output patterns

[critson]

• Rebase and update for changes in other code.
• Rename intrinsic wqm.helper to wqm.live (it provides live lanes).
• Add initial GlobalISel support.
• Add GlobalISel tests (some ported by @mbrkusanin).

[foad]

llvm/lib/Target/AMDGPU/SIWholeQuadMode.cpp

763–764

I'm pretty sure he meant s/unsigned Src/Register Src/. clang-tidy thinks the same.

[nhaehnle]

Hi Carl, what's the status of this?

[critson]

In D67767#2303214, @nhaehnle wrote:

Hi Carl, what's the status of this?

As per your previous comment:

In D67767#2108753, @nhaehnle wrote:

Does it make sense to wait until the WQM pass is moved?

Pending review of D88081 after which this will need a little bit of reworking to fit into post-MI / non-SSA world.

[critson]

Heavily rewritten:

• Use/reuse single virtual register for live mask. This removes need for PHIs and live mask register tracking. Assumes WQM is running non-SSA. (Supporting non-SSA and SSA operation would bloat code.)
• Live mask tracks all kills and demotes.
• Live mask manipulations terminate shader if all lanes are killed, even in non-uniform control flow.
• Move all kill lowering to WQM pass, this simplifies later passes and avoids duplication when updating live mask. Removes the need for "clean up" operations.
• WQM pass always modifies shader if it has any kills or demotes, even if there is no WQM.

[foad]

llvm/lib/Target/AMDGPU/SIInstructions.td

344

EXEC,SCC

llvm/test/CodeGen/AMDGPU/atomic_optimizations_pixelshader.ll

210–211 ↗

(On Diff #304383)

Why the extra move here?

[critson]

llvm/lib/Target/AMDGPU/SIInstructions.td

344

Thanks!

llvm/test/CodeGen/AMDGPU/atomic_optimizations_pixelshader.ll

210–211 ↗

(On Diff #304383)

This is unfortunate cruft from moving WQM pass later in compiler where there is less COPY elimination.
I have been thinking about putting some additional COPY clean up optimisations in another pass.
I had originally put some in WQM pass, but agreed with Nicolai this was probably not appropriate as the pass is already getting quite complicated.
In this specific case, first copy is saving live lane mask at begin of shader, second copy is PS_LIVE which copies live lane mask.
In practice these are very unlikely to be next to each other, and determining that they are the trivial copies of each other will need to consider control flow.

[critson]

Rebase

[critson]

Split into D94746, D94747 and D94748.

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

 
 // Pixel shaders only: whether the current pixel is live (i.e. not a helper
 // invocation for derivative computation).
 def int_amdgcn_ps_live : Intrinsic <
   [llvm_i1_ty],
   [],
   [IntrNoMem]>;
 
+// Like ps.live, but cannot be moved by LICM.
+def int_amdgcn_wqm_helper : Intrinsic <[llvm_i1_ty], [], []>;
+
 def int_amdgcn_mbcnt_lo :
   GCCBuiltin<"__builtin_amdgcn_mbcnt_lo">,
   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
 
 def int_amdgcn_mbcnt_hi :
   GCCBuiltin<"__builtin_amdgcn_mbcnt_hi">,
   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
 

 // the function.
 def int_amdgcn_wqm_vote : Intrinsic<[llvm_i1_ty],
   [llvm_i1_ty], [IntrNoMem, IntrConvergent]
 >;
 
 // If false, set EXEC=0 for the current thread until the end of program.
 def int_amdgcn_kill : Intrinsic<[], [llvm_i1_ty], []>;
 
+// If false, mark all active lanes as helper lanes until the end of program.
+def int_amdgcn_wqm_demote : Intrinsic<[], [llvm_i1_ty], []>;
+
 // Copies the active channels of the source value to the destination value,
 // with the guarantee that the source value is computed as if the entire
 // program were executed in Whole Wavefront Mode, i.e. with all channels
 // enabled, with a few exceptions: - Phi nodes with require WWM return an
 // undefined value.
 def int_amdgcn_wwm : Intrinsic<[llvm_any_ty],
   [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable, IntrConvergent]
 >;

llvm/lib/Target/AMDGPU/AMDGPUSearchableTables.td

 def : SourceOfDivergence<int_amdgcn_struct_buffer_atomic_umax>;
 def : SourceOfDivergence<int_amdgcn_struct_buffer_atomic_and>;
 def : SourceOfDivergence<int_amdgcn_struct_buffer_atomic_or>;
 def : SourceOfDivergence<int_amdgcn_struct_buffer_atomic_xor>;
 def : SourceOfDivergence<int_amdgcn_struct_buffer_atomic_inc>;
 def : SourceOfDivergence<int_amdgcn_struct_buffer_atomic_dec>;
 def : SourceOfDivergence<int_amdgcn_struct_buffer_atomic_cmpswap>;
 def : SourceOfDivergence<int_amdgcn_ps_live>;
+def : SourceOfDivergence<int_amdgcn_wqm_helper>;
 def : SourceOfDivergence<int_amdgcn_ds_swizzle>;
 def : SourceOfDivergence<int_amdgcn_ds_ordered_add>;
 def : SourceOfDivergence<int_amdgcn_ds_ordered_swap>;
 def : SourceOfDivergence<int_amdgcn_permlane16>;
 def : SourceOfDivergence<int_amdgcn_permlanex16>;
 def : SourceOfDivergence<int_amdgcn_mov_dpp>;
 def : SourceOfDivergence<int_amdgcn_mov_dpp8>;
 def : SourceOfDivergence<int_amdgcn_update_dpp>;

llvm/lib/Target/AMDGPU/SIInsertSkips.cpp

if (TRI->isVGPR(MBB.getParent()->getRegInfo(),
         .add(MI.getOperand(1))
         .addImm(0)  // src1 modifiers
         .add(MI.getOperand(0));
 
       I.addImm(0);  // omod
     }
     break;
   }
+  case AMDGPU::SI_DEMOTE_I1_TERMINATOR:
   case AMDGPU::SI_KILL_I1_TERMINATOR: {
     const MachineFunction *MF = MI.getParent()->getParent();
     const GCNSubtarget &ST = MF->getSubtarget<GCNSubtarget>();
     unsigned Exec = ST.isWave32() ? AMDGPU::EXEC_LO : AMDGPU::EXEC;
     const MachineOperand &Op = MI.getOperand(0);
     int64_t KillVal = MI.getOperand(1).getImm();
     assert(KillVal == 0 || KillVal == -1);
 
for (I = MBB.begin(); I != MBB.end(); I = Next) {
           // inserted after the current one and let skip the two instructions
           // performing the kill if the exec mask is non-zero.
           MI.eraseFromParent();
         }
         break;
 
       case AMDGPU::SI_KILL_F32_COND_IMM_TERMINATOR:
       case AMDGPU::SI_KILL_I1_TERMINATOR:
+      case AMDGPU::SI_DEMOTE_I1_TERMINATOR:
         MadeChange = true;
         kill(MI);
 
-        if (ExecBranchStack.empty()) {
+        if (ExecBranchStack.empty() &&
+            MI.getOpcode() != AMDGPU::SI_DEMOTE_I1_TERMINATOR) {
           if (NextBB != BE && skipIfDead(MI, *NextBB)) {
             HaveSkipBlock = true;
             NextBB = std::next(BI);
             BE = MF.end();
           }
         } else {
           HaveKill = true;
         }

llvm/lib/Target/AMDGPU/SIInstrInfo.cpp

case AMDGPU::S_ANDN2_B64_term:
     break;
 
   case AMDGPU::S_ANDN2_B32_term:
     // This is only a terminator to get the correct spill code placement during
     // register allocation.
     MI.setDesc(get(AMDGPU::S_ANDN2_B32));
     break;
 
+  case AMDGPU::S_AND_B64_term:
+    // This is only a terminator to get the correct spill code placement during
+    // register allocation.
+    MI.setDesc(get(AMDGPU::S_AND_B64));
+    break;
+
+  case AMDGPU::S_AND_B32_term:
+    // This is only a terminator to get the correct spill code placement during
+    // register allocation.
+    MI.setDesc(get(AMDGPU::S_AND_B32));
+    break;
+
   case AMDGPU::V_MOV_B64_PSEUDO: {
     Register Dst = MI.getOperand(0).getReg();
     Register DstLo = RI.getSubReg(Dst, AMDGPU::sub0);
     Register DstHi = RI.getSubReg(Dst, AMDGPU::sub1);
 
     const MachineOperand &SrcOp = MI.getOperand(1);
     // FIXME: Will this work for 64-bit floating point immediates?
     assert(!SrcOp.isFPImm());
bool SIInstrInfo::analyzeBranch(MachineBasicBlock &MBB, MachineBasicBlock *&TBB,
   // exec management.
   while (I != E && !I->isBranch() && !I->isReturn() &&
          I->getOpcode() != AMDGPU::SI_MASK_BRANCH) {
     switch (I->getOpcode()) {
     case AMDGPU::SI_MASK_BRANCH:
     case AMDGPU::S_MOV_B64_term:
     case AMDGPU::S_XOR_B64_term:
     case AMDGPU::S_ANDN2_B64_term:
+    case AMDGPU::S_AND_B64_term:
     case AMDGPU::S_MOV_B32_term:
     case AMDGPU::S_XOR_B32_term:
     case AMDGPU::S_OR_B32_term:
     case AMDGPU::S_ANDN2_B32_term:
+    case AMDGPU::S_AND_B32_term:
       break;
     case AMDGPU::SI_IF:
     case AMDGPU::SI_ELSE:
     case AMDGPU::SI_KILL_I1_TERMINATOR:
     case AMDGPU::SI_KILL_F32_COND_IMM_TERMINATOR:
+    case AMDGPU::SI_DEMOTE_I1_TERMINATOR:
       // FIXME: It's messy that these need to be considered here at all.
       return true;
     default:
       llvm_unreachable("unexpected non-branch terminator inst");
     }
 
     ++I;
   }
MachineInstrBuilder SIInstrInfo::getAddNoCarry(MachineBasicBlock &MBB,
   return BuildMI(MBB, I, DL, get(AMDGPU::V_ADD_I32_e64), DestReg)
            .addReg(UnusedCarry, RegState::Define | RegState::Dead);
 }
 
 bool SIInstrInfo::isKillTerminator(unsigned Opcode) {
   switch (Opcode) {
   case AMDGPU::SI_KILL_F32_COND_IMM_TERMINATOR:
   case AMDGPU::SI_KILL_I1_TERMINATOR:
+  case AMDGPU::SI_DEMOTE_I1_TERMINATOR:
     return true;
   default:
     return false;
   }
 }
 
 const MCInstrDesc &SIInstrInfo::getKillTerminatorFromPseudo(unsigned Opcode) const {
   switch (Opcode) {

llvm/lib/Target/AMDGPU/SIInstructions.td

class WrapTerminatorInst<SOP_Pseudo base_inst> : SPseudoInstSI<
   let UseNamedOperandTable = base_inst.UseNamedOperandTable;
   let CodeSize = base_inst.CodeSize;
 }
 
 let WaveSizePredicate = isWave64 in {
 def S_MOV_B64_term : WrapTerminatorInst<S_MOV_B64>;
 def S_XOR_B64_term : WrapTerminatorInst<S_XOR_B64>;
 def S_ANDN2_B64_term : WrapTerminatorInst<S_ANDN2_B64>;
+def S_AND_B64_term : WrapTerminatorInst<S_AND_B64>;
 }
 
 let WaveSizePredicate = isWave32 in {
 def S_MOV_B32_term : WrapTerminatorInst<S_MOV_B32>;
 def S_XOR_B32_term : WrapTerminatorInst<S_XOR_B32>;
 def S_OR_B32_term : WrapTerminatorInst<S_OR_B32>;
 def S_ANDN2_B32_term : WrapTerminatorInst<S_ANDN2_B32>;
+def S_AND_B32_term : WrapTerminatorInst<S_AND_B32>;
 }
 
 def WAVE_BARRIER : SPseudoInstSI<(outs), (ins),
   [(int_amdgcn_wave_barrier)]> {
   let SchedRW = [];
   let hasNoSchedulingInfo = 1;
   let hasSideEffects = 1;
   let mayLoad = 1;

 }
 
 def SI_PS_LIVE : PseudoInstSI <
   (outs SReg_1:$dst), (ins),
   [(set i1:$dst, (int_amdgcn_ps_live))]> {
   let SALU = 1;
 }
 
+let Uses = [EXEC] in {
+def SI_WQM_HELPER : PseudoInstSI <
+  (outs SReg_1:$dst), (ins),
+  [(set i1:$dst, (int_amdgcn_wqm_helper))]> {
+  let SALU = 1;
+}
+
+let Defs = [EXEC] in {
+def SI_DEMOTE_I1 : SPseudoInstSI <(outs), (ins SCSrc_i1:$src, i1imm:$killvalue)> {
+}
+def SI_DEMOTE_I1_TERMINATOR : SPseudoInstSI <(outs), (ins SCSrc_i1:$src, i1imm:$killvalue)> {
+  let isTerminator = 1;
+}
+} // End Defs = [EXEC]
+
+} // End Uses = [EXEC]
+
 def SI_MASKED_UNREACHABLE : SPseudoInstSI <(outs), (ins),
   [(int_amdgcn_unreachable)],
   "; divergent unreachable"> {
   let Size = 0;
   let hasNoSchedulingInfo = 1;
   let FixedSize = 1;
 }
 

 >;
 
 def : Pat <
   (int_amdgcn_kill (i1 (not i1:$src))),
   (SI_KILL_I1_PSEUDO $src, -1)
 >;
 
 def : Pat <
+  (int_amdgcn_wqm_demote i1:$src),
+  (SI_DEMOTE_I1 $src, 0)
+>;
+
+def : Pat <
+  (int_amdgcn_wqm_demote (i1 (not i1:$src))),
+  (SI_DEMOTE_I1 $src, -1)
+>;
+
+def : Pat <
   (AMDGPUkill i32:$src),
   (SI_KILL_F32_COND_IMM_PSEUDO $src, 0, 3) // 3 means SETOGE
 >;
 
 def : Pat <
   (int_amdgcn_kill (i1 (setcc f32:$src, InlineFPImm<f32>:$imm, cond:$cond))),
   (SI_KILL_F32_COND_IMM_PSEUDO $src, (bitcast_fpimm_to_i32 $imm), (cond_as_i32imm $cond))
 >;

llvm/lib/Target/AMDGPU/SIOptimizeExecMasking.cpp

case AMDGPU::S_ANDN2_B64_term: {
     return true;
   }
   case AMDGPU::S_ANDN2_B32_term: {
     // This is only a terminator to get the correct spill code placement during
     // register allocation.
     MI.setDesc(TII.get(AMDGPU::S_ANDN2_B32));
     return true;
   }
+  case AMDGPU::S_AND_B64_term: {
+    // This is only a terminator to get the correct spill code placement during
+    // register allocation.
+    MI.setDesc(TII.get(AMDGPU::S_AND_B64));
+    return true;
+  }
+  case AMDGPU::S_AND_B32_term: {
+    // This is only a terminator to get the correct spill code placement during
+    // register allocation.
+    MI.setDesc(TII.get(AMDGPU::S_AND_B32));
+    return true;
+  }
   default:
     return false;
   }
 }
 
 static MachineBasicBlock::reverse_iterator fixTerminators(
   const SIInstrInfo &TII,
   MachineBasicBlock &MBB) {

llvm/lib/Target/AMDGPU/SIWholeQuadMode.cpp

struct InstrInfo {
   char Disabled = 0;
   char OutNeeds = 0;
 };
 
 struct BlockInfo {
   char Needs = 0;
   char InNeeds = 0;
   char OutNeeds = 0;
+  char InitialState = 0;
+  unsigned LiveMaskIn = 0; // Initial live mask register
+  unsigned LiveMaskOut = 0; // Outgoing live mask register
 };
 
 struct WorkItem {
   MachineBasicBlock *MBB = nullptr;
   MachineInstr *MI = nullptr;
 
   WorkItem() = default;
   WorkItem(MachineBasicBlock *MBB) : MBB(MBB) {}
   WorkItem(MachineInstr *MI) : MI(MI) {}
 };
 
 class SIWholeQuadMode : public MachineFunctionPass {
 private:
   CallingConv::ID CallingConv;
   const SIInstrInfo *TII;
   const SIRegisterInfo *TRI;
   const GCNSubtarget *ST;
   MachineRegisterInfo *MRI;
   LiveIntervals *LIS;
 
   DenseMap<const MachineInstr *, InstrInfo> Instructions;
   DenseMap<MachineBasicBlock *, BlockInfo> Blocks;
-  SmallVector<MachineInstr *, 1> LiveMaskQueries;
+
+  // Tracks live mask output of instructions
+  DenseMap<const MachineInstr *, unsigned> LiveMaskRegs;
+  // Tracks state (WQM/WWM/Exact) after a given instruction
+  DenseMap<const MachineInstr *, char> StateTransition;
+
+  SmallVector<MachineInstr *, 2> LiveMaskQueries;
   SmallVector<MachineInstr *, 4> LowerToCopyInstrs;
+  SmallVector<MachineInstr *, 4> DemoteInstrs;
 
   void printInfo();
 
   void markInstruction(MachineInstr &MI, char Flag,
                        std::vector<WorkItem> &Worklist);
   void markInstructionUses(const MachineInstr &MI, char Flag,
                            std::vector<WorkItem> &Worklist);
   char scanInstructions(MachineFunction &MF, std::vector<WorkItem> &Worklist);
   void propagateInstruction(MachineInstr &MI, std::vector<WorkItem> &Worklist);
   void propagateBlock(MachineBasicBlock &MBB, std::vector<WorkItem> &Worklist);
   char analyzeFunction(MachineFunction &MF);
 
+  void scanLiveLanes(MachineBasicBlock &MBB,
+                     std::vector<MachineBasicBlock *> &Worklist);
+  void analyzeLiveLanes(MachineFunction &MF);
+
   bool requiresCorrectState(const MachineInstr &MI) const;
 
   MachineBasicBlock::iterator saveSCC(MachineBasicBlock &MBB,
                                       MachineBasicBlock::iterator Before);
   MachineBasicBlock::iterator
   prepareInsertion(MachineBasicBlock &MBB, MachineBasicBlock::iterator First,
                    MachineBasicBlock::iterator Last, bool PreferLast,
-                   bool SaveSCC);
+                   bool SaveSCC, bool CheckPhys);
   void toExact(MachineBasicBlock &MBB, MachineBasicBlock::iterator Before,
                unsigned SaveWQM, unsigned LiveMaskReg);
   void toWQM(MachineBasicBlock &MBB, MachineBasicBlock::iterator Before,
              unsigned SavedWQM);
   void toWWM(MachineBasicBlock &MBB, MachineBasicBlock::iterator Before,
              unsigned SaveOrig);
   void fromWWM(MachineBasicBlock &MBB, MachineBasicBlock::iterator Before,
-               unsigned SavedOrig);
-  void processBlock(MachineBasicBlock &MBB, unsigned LiveMaskReg, bool isEntry);
+               unsigned SavedOrig, char NonWWMState);
 
-  void lowerLiveMaskQueries(unsigned LiveMaskReg);
+  bool canSplitBlockAt(MachineBasicBlock *BB, MachineInstr *MI);
+  MachineBasicBlock *splitBlock(MachineBasicBlock *BB,
+                                MachineInstr *TermMI);
+  void lowerBlock(MachineBasicBlock &MBB);
+
+  unsigned findLiveMaskReg(MachineBasicBlock &MBB, BlockInfo &BI,
+                           MachineBasicBlock::iterator &Before);
+  void processBlock(MachineBasicBlock &MBB, bool isEntry);
+
+  bool lowerLiveMaskQueries(unsigned LiveMaskReg);
   void lowerCopyInstrs();
+  bool lowerDemoteInstrs();
+
+  void lowerLiveMaskQuery(MachineBasicBlock &MBB,
+                          MachineInstr &MI,
+                          unsigned LiveMaskReg,
+                          bool isWQM);
+  bool lowerDemote(MachineBasicBlock &MBB, MachineInstr &MI,
+                   unsigned LiveMaskIn, unsigned LiveMaskOut,
+                   bool isWQM);
 
 public:
   static char ID;
 
   SIWholeQuadMode() :
     MachineFunctionPass(ID) { }
 
   bool runOnMachineFunction(MachineFunction &MF) override;
 
   StringRef getPassName() const override { return "SI Whole Quad Mode"; }
 
   void getAnalysisUsage(AnalysisUsage &AU) const override {
     AU.addRequired<LiveIntervals>();
-    AU.addPreserved<SlotIndexes>();
-    AU.addPreserved<LiveIntervals>();
-    AU.setPreservesCFG();
     MachineFunctionPass::getAnalysisUsage(AU);
   }
 };
 
 } // end anonymous namespace
 
 char SIWholeQuadMode::ID = 0;
 
for (auto II = MBB.begin(), IE = MBB.end(); II != IE; ++II) {
         if (!(BBI.InNeeds & StateExact)) {
           BBI.InNeeds |= StateExact;
           Worklist.push_back(&MBB);
         }
         GlobalFlags |= StateExact;
         III.Disabled = StateWQM | StateWWM;
         continue;
       } else {
-        if (Opcode == AMDGPU::SI_PS_LIVE) {
+        if (Opcode == AMDGPU::SI_PS_LIVE || Opcode == AMDGPU::SI_WQM_HELPER) {
           LiveMaskQueries.push_back(&MI);
+        } else if (Opcode == AMDGPU::SI_DEMOTE_I1) {
+          DemoteInstrs.push_back(&MI);
         } else if (WQMOutputs) {
           // The function is in machine SSA form, which means that physical
           // VGPRs correspond to shader inputs and outputs. Inputs are
           // only used, outputs are only defined.
           for (const MachineOperand &MO : MI.defs()) {
             if (!MO.isReg())
               continue;
 
if (WI.MI)
       propagateInstruction(*WI.MI, Worklist);
     else
       propagateBlock(*WI.MBB, Worklist);
   }
 
   return GlobalFlags;
 }
 
+// Trace live mask manipulate through block, creating new virtual registers.
+// Additionally insert PHI nodes when block has multiple predecessors
+// which manipulated the mask.
+void SIWholeQuadMode::scanLiveLanes(MachineBasicBlock &MBB,
+                                    std::vector<MachineBasicBlock *> &Worklist) {
+  BlockInfo &BI = Blocks[&MBB];
+
+  if (BI.LiveMaskIn && BI.LiveMaskOut)
+    return; // Block has been fully traced already.
+
+  if (!BI.LiveMaskIn) {
+    // Find the incoming live mask, or insert PHI if there are multiple.
+    unsigned LastPredReg = 0;
+    unsigned Count = 0;
+    bool Valid = true;
+
+    // Find predecessor live masks while performing basic deduplication.
+    for (MachineBasicBlock *Pred : MBB.predecessors()) {
+      BlockInfo &PredBI = Blocks[Pred];
+      if (!PredBI.LiveMaskOut) {
+        Valid = false;
+        break;
+      }
+      if (PredBI.LiveMaskOut != LastPredReg) {
+        LastPredReg = PredBI.LiveMaskOut;
+        Count++;
+      }
+    }
+
+    if (Valid) {
+      // All predecessors have live mask outputs.
+      if (Count > 1) {
+        BI.LiveMaskIn = MRI->createVirtualRegister(TRI->getBoolRC());
+        MachineInstrBuilder PHI = BuildMI(MBB, MBB.begin(), DebugLoc(),
+                                    TII->get(TargetOpcode::PHI),
+                                    BI.LiveMaskIn);
+        for (MachineBasicBlock *Pred : MBB.predecessors()) {
+          BlockInfo &PredBI = Blocks[Pred];
+          PHI.addReg(PredBI.LiveMaskOut);
+          PHI.addMBB(Pred);
+        }
+        LIS->InsertMachineInstrInMaps(*PHI);
+      } else {
+        BI.LiveMaskIn = LastPredReg;
+      }
+    } else {
+      // Not all predecessor blocks have live mask outputs,
+      // so this block will need to be revisited.
+
+      if (!BI.LiveMaskOut) {
+        // Give this block a live mask output to ensure forward progress.
+        BI.LiveMaskOut = MRI->createVirtualRegister(TRI->getBoolRC());
+      }
+
+      // Queue this block to be revisited and visit predecessors.
+      Worklist.push_back(&MBB);
+      for (MachineBasicBlock *Pred : MBB.predecessors()) {
+        BlockInfo &PredBI = Blocks[Pred];
+        if (!PredBI.LiveMaskOut)
+          Worklist.push_back(Pred);
+      }
+      return;
+    }
+  }
+
+  assert(BI.LiveMaskIn);
+
+  // Now that the initial live mask register is known the block can
+  // be traced and intermediate live mask registers assigned for instructions
+  // which manipulate the mask.
+  unsigned CurrentLive = BI.LiveMaskIn;
+  auto II = MBB.getFirstNonPHI(), IE = MBB.end();
+  while (II != IE) {
+    MachineInstr &MI = *II;
+    if (MI.getOpcode() == AMDGPU::SI_DEMOTE_I1) {
+      unsigned NewLive = MRI->createVirtualRegister(TRI->getBoolRC());
+      LiveMaskRegs[&MI] = NewLive;
+      CurrentLive = NewLive;
+    }
+    II++;
+  }
+
+  // If an output register was assigned to guarantee forward progress
+  // then it is possible the current live register will not become the output
+  // live mask register.  This will be resolved during block lowering.
+  if (!BI.LiveMaskOut) {
+    BI.LiveMaskOut = CurrentLive;
+  }
+}
+
+// Scan blocks for live mask manipulation operations in reverse post order
+// to minimise rescans: a block will have to be rescanned if it's
+// predecessors live mask output is not defined.
+void SIWholeQuadMode::analyzeLiveLanes(MachineFunction &MF) {
+  std::vector<MachineBasicBlock *> Worklist;
+
+  ReversePostOrderTraversal<MachineFunction *> RPOT(&MF);
+  for (auto BI = RPOT.begin(), BE = RPOT.end(); BI != BE; ++BI) {
+    MachineBasicBlock &MBB = **BI;
+    scanLiveLanes(MBB, Worklist);
+  }
+
+  while (!Worklist.empty()) {
+    MachineBasicBlock *MBB = Worklist.back();
+    Worklist.pop_back();
+    scanLiveLanes(*MBB, Worklist);
+  }
+}
+
 /// Whether \p MI really requires the exec state computed during analysis.
 ///
 /// Scalar instructions must occasionally be marked WQM for correct propagation
 /// (e.g. thread masks leading up to branches), but when it comes to actual
 /// execution, they don't care about EXEC.
 bool SIWholeQuadMode::requiresCorrectState(const MachineInstr &MI) const {
   if (MI.isTerminator())
     return true;
SIWholeQuadMode::saveSCC(MachineBasicBlock &MBB,
   return Restore;
 }
 
 // Return an iterator in the (inclusive) range [First, Last] at which
 // instructions can be safely inserted, keeping in mind that some of the
 // instructions we want to add necessarily clobber SCC.
 MachineBasicBlock::iterator SIWholeQuadMode::prepareInsertion(
     MachineBasicBlock &MBB, MachineBasicBlock::iterator First,
-    MachineBasicBlock::iterator Last, bool PreferLast, bool SaveSCC) {
+    MachineBasicBlock::iterator Last, bool PreferLast, bool SaveSCC,
+    bool CheckPhys) {
   if (!SaveSCC)
     return PreferLast ? Last : First;
 
   LiveRange &LR = LIS->getRegUnit(*MCRegUnitIterator(AMDGPU::SCC, TRI));
   auto MBBE = MBB.end();
   SlotIndex FirstIdx = First != MBBE ? LIS->getInstructionIndex(*First)
                                      : LIS->getMBBEndIdx(&MBB);
   SlotIndex LastIdx =
if (PreferLast) {
       if (Next > LastIdx)
         break;
       Idx = Next;
     }
   }
 
   MachineBasicBlock::iterator MBBI;
 
-  if (MachineInstr *MI = LIS->getInstructionFromIndex(Idx))
+  if (MachineInstr *MI = LIS->getInstructionFromIndex(Idx)) {
     MBBI = MI;
-  else {
+    if (CheckPhys) {
+      // Make sure insertion point is after any COPY instructions
+      // accessing physical live in registers.  This is ensures that
+      // block splitting does not occur before all live ins have been copied.
+      while (MBBI != Last) {
+        if (MBBI->getOpcode() != AMDGPU::COPY)
+          break;
+        unsigned Src = MBBI->getOperand(1).getReg();
+        if (!Register::isVirtualRegister(Src) && MBB.isLiveIn(Src)) {
+          MBBI++;
+        } else {
+          break;
+        }
+      }
+    }
+  } else {
     assert(Idx == LIS->getMBBEndIdx(&MBB));
     MBBI = MBB.end();
   }
 
   if (S)
     MBBI = saveSCC(MBB, MBBI);
 
   return MBBI;
if (SaveWQM) {
     MI = BuildMI(MBB, Before, DebugLoc(), TII->get(ST->isWave32() ?
                    AMDGPU::S_AND_B32 : AMDGPU::S_AND_B64),
                  Exec)
              .addReg(Exec)
              .addReg(LiveMaskReg);
   }
 
   LIS->InsertMachineInstrInMaps(*MI);
+  StateTransition[MI] = StateExact;
 }
 
 void SIWholeQuadMode::toWQM(MachineBasicBlock &MBB,
                             MachineBasicBlock::iterator Before,
                             unsigned SavedWQM) {
   MachineInstr *MI;
 
   unsigned Exec = ST->isWave32() ? AMDGPU::EXEC_LO : AMDGPU::EXEC;
   if (SavedWQM) {
     MI = BuildMI(MBB, Before, DebugLoc(), TII->get(AMDGPU::COPY), Exec)
              .addReg(SavedWQM);
   } else {
     MI = BuildMI(MBB, Before, DebugLoc(), TII->get(ST->isWave32() ?
                    AMDGPU::S_WQM_B32 : AMDGPU::S_WQM_B64),
                  Exec)
              .addReg(Exec);
   }
 
   LIS->InsertMachineInstrInMaps(*MI);
+  StateTransition[MI] = StateWQM;
 }
 
 void SIWholeQuadMode::toWWM(MachineBasicBlock &MBB,
                             MachineBasicBlock::iterator Before,
                             unsigned SaveOrig) {
   MachineInstr *MI;
 
   assert(SaveOrig);
   MI = BuildMI(MBB, Before, DebugLoc(), TII->get(AMDGPU::ENTER_WWM), SaveOrig)
            .addImm(-1);
   LIS->InsertMachineInstrInMaps(*MI);
+  StateTransition[MI] = StateWWM;
 }
 
 void SIWholeQuadMode::fromWWM(MachineBasicBlock &MBB,
                               MachineBasicBlock::iterator Before,
-                              unsigned SavedOrig) {
+                              unsigned SavedOrig,
+                              char NonWWMState) {
   MachineInstr *MI;
 
   assert(SavedOrig);
   MI = BuildMI(MBB, Before, DebugLoc(), TII->get(AMDGPU::EXIT_WWM),
                ST->isWave32() ? AMDGPU::EXEC_LO : AMDGPU::EXEC)
            .addReg(SavedOrig);
   LIS->InsertMachineInstrInMaps(*MI);
+  StateTransition[MI] = NonWWMState;
 }
 
-void SIWholeQuadMode::processBlock(MachineBasicBlock &MBB, unsigned LiveMaskReg,
-                                   bool isEntry) {
+void SIWholeQuadMode::lowerLiveMaskQuery(MachineBasicBlock &MBB,
+                                         MachineInstr &MI,
+                                         unsigned LiveMaskReg,
+                                         bool isWQM) {
+  const DebugLoc &DL = MI.getDebugLoc();
+  unsigned Dest = MI.getOperand(0).getReg();
+  MachineInstr *Copy =
+      BuildMI(MBB, MI, DL, TII->get(AMDGPU::COPY), Dest)
+          .addReg(LiveMaskReg);
+  LIS->ReplaceMachineInstrInMaps(MI, *Copy);
+  MBB.remove(&MI);
+}
+
+// Lower an instruction which demotes lanes to helpers by adding
+// appropriate live mask manipulation.  Note this is also applied to kills.
+bool SIWholeQuadMode::lowerDemote(MachineBasicBlock &MBB, MachineInstr &MI,
+                                  unsigned LiveMaskIn, unsigned LiveMaskOut,
+                                  bool isWQM) {
+  const unsigned Exec = ST->isWave32() ? AMDGPU::EXEC_LO : AMDGPU::EXEC;
+  const unsigned AndN2 =
+    ST->isWave32() ? AMDGPU::S_ANDN2_B32 : AMDGPU::S_ANDN2_B64;
+  const unsigned And =
+    ST->isWave32() ? AMDGPU::S_AND_B32 : AMDGPU::S_AND_B64;
+
+  const DebugLoc &DL = MI.getDebugLoc();
+  MachineInstr *NewMI = nullptr;
+  bool NeedSplit = false;
+
+  const MachineOperand &Op = MI.getOperand(0);
+  int64_t KillVal = MI.getOperand(1).getImm();
+  if (Op.isImm()) {
+    int64_t Imm = Op.getImm();
+    if (Imm == KillVal) {
+      NewMI = BuildMI(MBB, MI, DL,
+                      TII->get(AndN2),
+                      LiveMaskOut)
+        .addReg(LiveMaskIn)
+        .addReg(Exec);
+    }
+  } else {
+    unsigned Opcode = KillVal ? AndN2 : And;
+    NewMI = BuildMI(MBB, MI, DL,
+                    TII->get(Opcode),
+                    LiveMaskOut)
+      .addReg(LiveMaskIn)
+      .add(Op);
+  }
+
+  if (MI.getOpcode() == AMDGPU::SI_DEMOTE_I1) {
+    if (isWQM) {
+      // Inside WQM demotes are replaced with live mask manipulation
+      LIS->RemoveMachineInstrFromMaps(MI);
+      MBB.remove(&MI);
+    } else {
+      // Outside WQM demotes become kills terminating the block
+      NeedSplit = true;
+    }
+  }
+
+  if (NewMI) {
+    LIS->InsertMachineInstrInMaps(*NewMI);
+  }
+
+  return NeedSplit;
+}
+
+bool SIWholeQuadMode::canSplitBlockAt(MachineBasicBlock *BB,
+                                      MachineInstr *MI) {
+  // Cannot split immediately before the epilog
+  // because there are values in physical registers
+  if (MI->getOpcode() == AMDGPU::SI_RETURN_TO_EPILOG) {
+    return false;
+  }
+
+  return true;
+}
+
+MachineBasicBlock *SIWholeQuadMode::splitBlock(MachineBasicBlock *BB,
+                                               MachineInstr *TermMI) {
+  MachineBasicBlock::iterator SplitPoint(TermMI);
+  SplitPoint++;
+
+  LLVM_DEBUG(dbgs() << "Split block " << printMBBReference(*BB)
+                    << " @ " << *TermMI << "\n");
+
+  MachineBasicBlock *SplitBB = nullptr;
+
+  // Only split the block if the split point is not
+  // already the end of the block.
+  if (SplitPoint != BB->getFirstTerminator()) {
+    MachineFunction *MF = BB->getParent();
+    SplitBB = MF->CreateMachineBasicBlock(BB->getBasicBlock());
+
+    MachineFunction::iterator MBBI(BB);
+    ++MBBI;
+    MF->insert(MBBI, SplitBB);
+
+    SplitBB->splice(SplitBB->begin(), BB, SplitPoint, BB->end());
+    SplitBB->transferSuccessorsAndUpdatePHIs(BB);
+    BB->addSuccessor(SplitBB);
+  }
+
+  // Convert last instruction in to a terminator.
+  // Note: this only covers the expected patterns
+  switch (TermMI->getOpcode()) {
+  case AMDGPU::S_AND_B32:
+    TermMI->setDesc(TII->get(AMDGPU::S_AND_B32_term));
+    break;
+  case AMDGPU::S_AND_B64:
+    TermMI->setDesc(TII->get(AMDGPU::S_AND_B64_term));
+    break;
+  case AMDGPU::SI_DEMOTE_I1:
+    TermMI->setDesc(TII->get(AMDGPU::SI_DEMOTE_I1_TERMINATOR));
+    break;
+  case AMDGPU::SI_DEMOTE_I1_TERMINATOR:
+    break;
+  default:
+    if (BB->getFirstTerminator() == BB->end()) {
+      assert(SplitBB != nullptr);
+      BuildMI(*BB, BB->end(), DebugLoc(), TII->get(AMDGPU::S_BRANCH))
+        .addMBB(SplitBB);
+    }
+    break;
+  }
+
+  return SplitBB;
+}
+
+// Replace (or supplement) instructions accessing live mask.
+// This can only happen once all the live mask registers have been created
+// and the execute state (WQM/WWM/Exact) of instructions is known.
+void SIWholeQuadMode::lowerBlock(MachineBasicBlock &MBB) {
   auto BII = Blocks.find(&MBB);
   if (BII == Blocks.end())
     return;
 
+  LLVM_DEBUG(dbgs() << "\nLowering block " << printMBBReference(MBB)
+                    << ":\n");
+
   const BlockInfo &BI = BII->second;
 
+  SmallVector<MachineInstr *, 4> SplitPoints;
+  unsigned LiveMaskReg = BI.LiveMaskIn;
+  char State = BI.InitialState;
+
+  auto II = MBB.getFirstNonPHI(), IE = MBB.end();
+  while (II != IE) {
+    auto Next = std::next(II);
+    MachineInstr &MI = *II;
+
+    if (StateTransition.count(&MI)) {
+      // Mark transitions to Exact mode as split points so they become
+      // block terminators.
+      if (State != StateTransition[&MI] && StateTransition[&MI] == StateExact) {
+        if (State != StateWWM && canSplitBlockAt(&MBB, &MI))
+          SplitPoints.push_back(&MI);
+      }
+      State = StateTransition[&MI];
+    }
+
+    switch (MI.getOpcode()) {
+    case AMDGPU::SI_PS_LIVE:
+    case AMDGPU::SI_WQM_HELPER:
+      lowerLiveMaskQuery(MBB, MI, LiveMaskReg, State == StateWQM);
+      break;
+    case AMDGPU::SI_DEMOTE_I1: {
+      bool NeedSplit = lowerDemote(MBB, MI, LiveMaskReg,
+                                   LiveMaskRegs[&MI],
+                                   State == StateWQM);
+      if (NeedSplit)
+        SplitPoints.push_back(&MI);
+      break;
+    }
+    default:
+      break;
+    }
+
+    if (LiveMaskRegs.count(&MI))
+      LiveMaskReg = LiveMaskRegs[&MI];
+
+    II = Next;
+  }
+
+  if (BI.LiveMaskOut != LiveMaskReg) {
+    // If the final live mask register does not match the expected
+    // register of successor blocks then insert a copy.
+    MachineBasicBlock::instr_iterator Terminator =
+      MBB.getFirstInstrTerminator();
+    MachineInstr *MI = BuildMI(MBB, Terminator, DebugLoc(),
+                               TII->get(AMDGPU::COPY), BI.LiveMaskOut)
+      .addReg(LiveMaskReg);
+    LIS->InsertMachineInstrInMaps(*MI);
+  }
+
+  // Perform splitting after instruction scan to simplify iteration.
+  if (!SplitPoints.empty()) {
+    MachineBasicBlock *BB = &MBB;
+    for (MachineInstr *MI : SplitPoints) {
+      BB = splitBlock(BB, MI);
+    }
+  }
+}
+
+unsigned SIWholeQuadMode::findLiveMaskReg(MachineBasicBlock &MBB, BlockInfo &BI,
+                                          MachineBasicBlock::iterator &Before) {
+  assert(BI.LiveMaskIn);
+  if (BI.LiveMaskIn == BI.LiveMaskOut)
+    return BI.LiveMaskIn;
+
+  // FIXME: make this more efficient than scanning all instructions in a block
+  unsigned LiveMaskReg = BI.LiveMaskIn;
+  auto II = MBB.getFirstNonPHI(), IE = MBB.end();
+
+  while ((II != IE) && (II != Before)) {
+    MachineInstr *I = &*II;
+    if (LiveMaskRegs.count(I))
+      LiveMaskReg = LiveMaskRegs[I];
+    II++;
+  }
+
+  assert(LiveMaskReg);
+  return LiveMaskReg;
+}
+
+void SIWholeQuadMode::processBlock(MachineBasicBlock &MBB, bool isEntry) {
+  auto BII = Blocks.find(&MBB);
+  if (BII == Blocks.end())
+    return;
+
+  BlockInfo &BI = BII->second;
+
   // This is a non-entry block that is WQM throughout, so no need to do
   // anything.
-  if (!isEntry && BI.Needs == StateWQM && BI.OutNeeds != StateExact)
+  if (!isEntry && BI.Needs == StateWQM && BI.OutNeeds != StateExact) {
+    BI.InitialState = StateWQM;
     return;
+  }
 
   LLVM_DEBUG(dbgs() << "\nProcessing block " << printMBBReference(MBB)
                     << ":\n");
 
   unsigned SavedWQMReg = 0;
   unsigned SavedNonWWMReg = 0;
   bool WQMFromExec = isEntry;
   char State = (isEntry || !(BI.InNeeds & StateWQM)) ? StateExact : StateWQM;
   char NonWWMState = 0;
   const TargetRegisterClass *BoolRC = TRI->getBoolRC();
 
   auto II = MBB.getFirstNonPHI(), IE = MBB.end();
   if (isEntry)
     ++II; // Skip the instruction that saves LiveMask
 
   // This stores the first instruction where it's safe to switch from WQM to
   // Exact or vice versa.
   MachineBasicBlock::iterator FirstWQM = IE;
 
   // This stores the first instruction where it's safe to switch from WWM to
   // Exact/WQM or to switch to WWM. It must always be the same as, or after,
   // FirstWQM since if it's safe to switch to/from WWM, it must be safe to
   // switch to/from WQM as well.
   MachineBasicBlock::iterator FirstWWM = IE;
+
+  // Record initial state is block information.
+  BI.InitialState = State;
+
   for (;;) {
     MachineBasicBlock::iterator Next = II;
     char Needs = StateExact | StateWQM; // WWM is disabled by default
     char OutNeeds = 0;
 
     if (FirstWQM == IE)
       FirstWQM = II;
 
if (!(Needs & State)) {
         First = FirstWWM;
       } else {
         // We only need to switch to/from WQM, so we can use FirstWQM
         First = FirstWQM;
       }
 
       MachineBasicBlock::iterator Before =
           prepareInsertion(MBB, First, II, Needs == StateWQM,
-                           Needs == StateExact || WQMFromExec);
+                           Needs == StateExact || WQMFromExec,
+                           Needs == StateExact && isEntry);
 
       if (State == StateWWM) {
         assert(SavedNonWWMReg);
-        fromWWM(MBB, Before, SavedNonWWMReg);
+        fromWWM(MBB, Before, SavedNonWWMReg, NonWWMState);
         State = NonWWMState;
       }
 
       if (Needs == StateWWM) {
         NonWWMState = State;
         SavedNonWWMReg = MRI->createVirtualRegister(BoolRC);
         toWWM(MBB, Before, SavedNonWWMReg);
         State = StateWWM;
       } else {
         if (State == StateWQM && (Needs & StateExact) && !(Needs & StateWQM)) {
           if (!WQMFromExec && (OutNeeds & StateWQM))
             SavedWQMReg = MRI->createVirtualRegister(BoolRC);
 
-          toExact(MBB, Before, SavedWQMReg, LiveMaskReg);
+          toExact(MBB, Before, SavedWQMReg, findLiveMaskReg(MBB, BI, Before));
           State = StateExact;
         } else if (State == StateExact && (Needs & StateWQM) &&
                    !(Needs & StateExact)) {
           assert(WQMFromExec == (SavedWQMReg == 0));
 
           toWQM(MBB, Before, SavedWQMReg);
 
           if (SavedWQMReg) {
for (;;) {
     if (Needs != (StateExact | StateWQM | StateWWM)) {
       if (Needs != (StateExact | StateWQM))
         FirstWQM = IE;
       FirstWWM = IE;
     }
 
     if (II == IE)
       break;
+
     II = Next;
   }
 }
 
-void SIWholeQuadMode::lowerLiveMaskQueries(unsigned LiveMaskReg) {
+bool SIWholeQuadMode::lowerLiveMaskQueries(unsigned LiveMaskReg) {
+  bool Changed = false;
   for (MachineInstr *MI : LiveMaskQueries) {
     const DebugLoc &DL = MI->getDebugLoc();
     Register Dest = MI->getOperand(0).getReg();
     MachineInstr *Copy =
         BuildMI(*MI->getParent(), MI, DL, TII->get(AMDGPU::COPY), Dest)
             .addReg(LiveMaskReg);
 
     LIS->ReplaceMachineInstrInMaps(*MI, *Copy);
     MI->eraseFromParent();
+    Changed = true;
+  }
+  return Changed;
+}
+
+bool SIWholeQuadMode::lowerDemoteInstrs() {
+  bool Changed = false;
+  for (MachineInstr *MI : DemoteInstrs) {
+    MachineBasicBlock *MBB = MI->getParent();
+    splitBlock(MBB, MI);
+    Changed = true;
   }
+  return Changed;
 }
 
 void SIWholeQuadMode::lowerCopyInstrs() {
   for (MachineInstr *MI : LowerToCopyInstrs) {
     for (unsigned i = MI->getNumExplicitOperands() - 1; i > 1; i--)
       MI->RemoveOperand(i);
 
     const Register Reg = MI->getOperand(0).getReg();
void SIWholeQuadMode::lowerCopyInstrs() {
   }
 }
 
 bool SIWholeQuadMode::runOnMachineFunction(MachineFunction &MF) {
   Instructions.clear();
   Blocks.clear();
   LiveMaskQueries.clear();
   LowerToCopyInstrs.clear();
+  DemoteInstrs.clear();
+  LiveMaskRegs.clear();
+  StateTransition.clear();
+
   CallingConv = MF.getFunction().getCallingConv();
 
   ST = &MF.getSubtarget<GCNSubtarget>();
 
   TII = ST->getInstrInfo();
   TRI = &TII->getRegisterInfo();
   MRI = &MF.getRegInfo();
   LIS = &getAnalysis<LiveIntervals>();
 
-  char GlobalFlags = analyzeFunction(MF);
-  unsigned LiveMaskReg = 0;
-  unsigned Exec = ST->isWave32() ? AMDGPU::EXEC_LO : AMDGPU::EXEC;
-  if (!(GlobalFlags & StateWQM)) {
-    lowerLiveMaskQueries(Exec);
-    if (!(GlobalFlags & StateWWM) && LowerToCopyInstrs.empty())
-      return !LiveMaskQueries.empty();
-  } else {
-    // Store a copy of the original live mask when required
-    MachineBasicBlock &Entry = MF.front();
-    MachineBasicBlock::iterator EntryMI = Entry.getFirstNonPHI();
+  const char GlobalFlags = analyzeFunction(MF);
+  const bool NeedsLiveMask =
+    !(DemoteInstrs.empty() && LiveMaskQueries.empty());
+  const unsigned Exec = ST->isWave32() ? AMDGPU::EXEC_LO : AMDGPU::EXEC;
+  unsigned LiveMaskReg = Exec;
+
+  if (!(GlobalFlags & (StateWQM | StateWWM)) && LowerToCopyInstrs.empty()) {
+    // Shader only needs Exact mode
+    const bool LoweredQueries = lowerLiveMaskQueries(LiveMaskReg);
+    const bool LoweredDemotes = lowerDemoteInstrs();
+    return LoweredQueries || LoweredDemotes;
+  }
+
 
-    if (GlobalFlags & StateExact || !LiveMaskQueries.empty()) {
-      LiveMaskReg = MRI->createVirtualRegister(TRI->getBoolRC());
-      MachineInstr *MI = BuildMI(Entry, EntryMI, DebugLoc(),
-                                 TII->get(AMDGPU::COPY), LiveMaskReg)
-                             .addReg(Exec);
-      LIS->InsertMachineInstrInMaps(*MI);
-    }
+  // Store a copy of the original live mask when required
+  if (NeedsLiveMask || (GlobalFlags & StateWQM)) {
 
-    lowerLiveMaskQueries(LiveMaskReg);
-
-    if (GlobalFlags == StateWQM) {
-      // For a shader that needs only WQM, we can just set it once.
-      BuildMI(Entry, EntryMI, DebugLoc(), TII->get(ST->isWave32() ?
-                AMDGPU::S_WQM_B32 : AMDGPU::S_WQM_B64),
-              Exec)
-          .addReg(Exec);
+  if ((GlobalFlags == StateWQM) && DemoteInstrs.empty()) {
+    // Shader only needs WQM
 
-      lowerCopyInstrs();
-      // EntryMI may become invalid here
-      return true;
-    }
+    lowerLiveMaskQueries(LiveMaskReg);
+
+  if (NeedsLiveMask && (GlobalFlags & StateWQM)) {
+    BlockInfo &BI = Blocks[&Entry];
+    BI.LiveMaskIn = LiveMaskReg;
+    analyzeLiveLanes(MF);
+  } else {
+    for (auto BII : Blocks) {
+      BlockInfo &BI = Blocks[&*BII.first];
+      BI.LiveMaskIn = LiveMaskReg;
+      BI.LiveMaskOut = LiveMaskReg;
+    }
   }
 
   LLVM_DEBUG(printInfo());
 
   lowerCopyInstrs();
 
-  // Handle the general case
-  for (auto BII : Blocks)
-    processBlock(*BII.first, LiveMaskReg, BII.first == &*MF.begin());
-
-  // Physical registers like SCC aren't tracked by default anyway, so just
-  // removing the ranges we computed is the simplest option for maintaining
-  // the analysis results.
-  LIS->removeRegUnit(*MCRegUnitIterator(AMDGPU::SCC, TRI));
+  for (auto BII : Blocks) {
+    processBlock(*BII.first, BII.first == &Entry);
+  }
+
+  if (NeedsLiveMask && (GlobalFlags & StateWQM)) {
+    // Lowering blocks causes block splitting.
+    // Hence live ranges and slot indexes cease to be valid here.
+    for (auto BII : Blocks) {
+      lowerBlock(*BII.first);
+    }
+  } else {
+    lowerLiveMaskQueries(LiveMaskReg);
+    lowerDemoteInstrs();
+  }
 
   return true;
 }

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.wqm.demote.ll

+; RUN: llc -march=amdgcn -mcpu=tonga -verify-machineinstrs < %s | FileCheck -check-prefixes=GCN,SI,GCN-64 %s
+; RUN: llc -march=amdgcn -mcpu=gfx900 -verify-machineinstrs < %s | FileCheck -check-prefixes=GCN,GFX9,GCN-64 %s
+; RUN: llc -march=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize32,-wavefrontsize64 -verify-machineinstrs < %s | FileCheck -check-prefixes=GCN,GFX10,GCN-32 %s
+; RUN: llc -march=amdgcn -mcpu=gfx1010 -mattr=-wavefrontsize32,+wavefrontsize64 -verify-machineinstrs < %s | FileCheck -check-prefixes=GCN,GFX10,GCN-64 %s
+
+; GCN-LABEL: {{^}}static_exact:
+; GCN-32: v_cmp_gt_f32_e32 [[CMP:vcc_lo]], 0, v0
+; GCN-64: v_cmp_gt_f32_e32 [[CMP:vcc]], 0, v0
+; GCN-32: s_mov_b32 exec_lo, 0
+; GCN-64: s_mov_b64 exec, 0
+; GCN: v_cndmask_b32_e64 v{{[0-9]+}}, 0, 1.0, [[CMP]]
+; GCN: exp mrt1 v0, v0, v0, v0 done vm
+define amdgpu_ps void @static_exact(float %arg0, float %arg1) {
+.entry:
+  %c0 = fcmp olt float %arg0, 0.000000e+00
+  %c1 = fcmp oge float %arg1, 0.0
+  call void @llvm.amdgcn.wqm.demote(i1 false)
+  %tmp1 = select i1 %c0, float 1.000000e+00, float 0.000000e+00
+  call void @llvm.amdgcn.exp.f32(i32 1, i32 15, float %tmp1, float %tmp1, float %tmp1, float %tmp1, i1 true, i1 true) #0
+  ret void
+}
+
+; GCN-LABEL: {{^}}dynamic_exact:
+; GCN-32: v_cmp_le_f32_e64 [[CND:s[0-9]+]], 0, v1
+; GCN-64: v_cmp_le_f32_e64 [[CND:s\[[0-9]+:[0-9]+\]]], 0, v1
+; GCN-32: v_cmp_gt_f32_e32 [[CMP:vcc_lo]], 0, v0
+; GCN-64: v_cmp_gt_f32_e32 [[CMP:vcc]], 0, v0
+; GCN-32: s_and_b32 exec_lo, exec_lo, [[CND]]
+; GCN-64: s_and_b64 exec, exec, [[CND]]
+; GCN: v_cndmask_b32_e64 v{{[0-9]+}}, 0, 1.0, [[CMP]]
+; GCN: exp mrt1 v0, v0, v0, v0 done vm
+define amdgpu_ps void @dynamic_exact(float %arg0, float %arg1) {
+.entry:
+  %c0 = fcmp olt float %arg0, 0.000000e+00
+  %c1 = fcmp oge float %arg1, 0.0
+  call void @llvm.amdgcn.wqm.demote(i1 %c1)
+  %tmp1 = select i1 %c0, float 1.000000e+00, float 0.000000e+00
+  call void @llvm.amdgcn.exp.f32(i32 1, i32 15, float %tmp1, float %tmp1, float %tmp1, float %tmp1, i1 true, i1 true) #0
+  ret void
+}
+
+; GCN-LABEL: {{^}}branch:
+; GCN-32: s_and_saveexec_b32 s1, s0
+; GCN-64: s_and_saveexec_b64 s[2:3], s[0:1]
+; GCN-32: s_xor_b32 s0, exec_lo, s1
+; GCN-64: s_xor_b64 s[0:1], exec, s[2:3]
+; GCN-32: s_mov_b32 exec_lo, 0
+; GCN-64: s_mov_b64 exec, 0
+; GCN-32: s_or_b32 exec_lo, exec_lo, s0
+; GCN-64: s_or_b64 exec, exec, s[0:1]
+; GCN: v_cndmask_b32_e64 v0, 0, 1.0, vcc
+; GCN: exp mrt1 v0, v0, v0, v0 done vm
+define amdgpu_ps void @branch(float %arg0, float %arg1) {
+.entry:
+  %i0 = fptosi float %arg0 to i32
+  %i1 = fptosi float %arg1 to i32
+  %c0 = or i32 %i0, %i1
+  %c1 = and i32 %c0, 1
+  %c2 = icmp eq i32 %c1, 0
+  br i1 %c2, label %.continue, label %.demote
+
+.demote:
+  call void @llvm.amdgcn.wqm.demote(i1 false)
+  br label %.continue
+
+.continue:
+  %tmp1 = select i1 %c2, float 1.000000e+00, float 0.000000e+00
+  call void @llvm.amdgcn.exp.f32(i32 1, i32 15, float %tmp1, float %tmp1, float %tmp1, float %tmp1, i1 true, i1 true) #0
+  ret void
+}
+
+
+; GCN-LABEL: {{^}}wqm_demote_1:
+; GCN-NEXT: ; %.entry
+; GCN-32: s_mov_b32 [[ORIG:s[0-9]+]], exec_lo
+; GCN-64: s_mov_b64 [[ORIG:s\[[0-9]+:[0-9]+\]]], exec
+; GCN-32: s_wqm_b32 exec_lo, exec_lo
+; GCN-64: s_wqm_b64 exec, exec
+; GCN: ; %.demote
+; GCN-32-NEXT: s_andn2_b32 [[LIVE:s[0-9]+]], [[ORIG]], exec
+; GCN-64-NEXT: s_andn2_b64 [[LIVE:s\[[0-9]+:[0-9]+\]]], [[ORIG]], exec
+; GCN: ; %.continue
+; GCN: image_sample
+; GCN: v_add_f32_e32
+; GCN-32: s_and_b32 exec_lo, exec_lo, [[LIVE]]
+; GCN-64: s_and_b64 exec, exec, [[LIVE]]
+; GCN: image_sample
+define amdgpu_ps <4 x float> @wqm_demote_1(<8 x i32> inreg %rsrc, <4 x i32> inreg %sampler, i32 %idx, float %data, float %coord, float %coord2, float %z) {
+.entry:
+  %z.cmp = fcmp olt float %z, 0.0
+  br i1 %z.cmp, label %.continue, label %.demote
+
+.demote:
+  call void @llvm.amdgcn.wqm.demote(i1 false)
+  br label %.continue
+
+.continue:
+  %tex = call <4 x float> @llvm.amdgcn.image.sample.1d.v4f32.f32(i32 15, float %coord, <8 x i32> %rsrc, <4 x i32> %sampler, i1 0, i32 0, i32 0) #0
+  %tex0 = extractelement <4 x float> %tex, i32 0
+  %tex1 = extractelement <4 x float> %tex, i32 0
+  %coord1 = fadd float %tex0, %tex1
+  %rtex = call <4 x float> @llvm.amdgcn.image.sample.1d.v4f32.f32(i32 15, float %coord1, <8 x i32> %rsrc, <4 x i32> %sampler, i1 0, i32 0, i32 0) #0
+
+  ret <4 x float> %rtex
+}
+
+; GCN-LABEL: {{^}}wqm_demote_2:
+; GCN-NEXT: ; %.entry
+; GCN-32: s_mov_b32 [[ORIG:s[0-9]+]], exec_lo
+; GCN-64: s_mov_b64 [[ORIG:s\[[0-9]+:[0-9]+\]]], exec
+; GCN-32: s_wqm_b32 exec_lo, exec_lo
+; GCN-64: s_wqm_b64 exec, exec
+; GCN: image_sample
+; GCN: ; %.demote
+; GCN-32-NEXT: s_andn2_b32 [[LIVE:s[0-9]+]], [[ORIG]], exec
+; GCN-64-NEXT: s_andn2_b64 [[LIVE:s\[[0-9]+:[0-9]+\]]], [[ORIG]], exec
+; GCN: ; %.continue
+; GCN: v_add_f32_e32
+; GCN-32: s_and_b32 exec_lo, exec_lo, [[LIVE]]
+; GCN-64: s_and_b64 exec, exec, [[LIVE]]
+; GCN: image_sample
+define amdgpu_ps <4 x float> @wqm_demote_2(<8 x i32> inreg %rsrc, <4 x i32> inreg %sampler, i32 %idx, float %data, float %coord, float %coord2, float %z) {
+.entry:
+  %tex = call <4 x float> @llvm.amdgcn.image.sample.1d.v4f32.f32(i32 15, float %coord, <8 x i32> %rsrc, <4 x i32> %sampler, i1 0, i32 0, i32 0) #0
+  %tex0 = extractelement <4 x float> %tex, i32 0
+  %tex1 = extractelement <4 x float> %tex, i32 0
+  %z.cmp = fcmp olt float %tex0, 0.0
+  br i1 %z.cmp, label %.continue, label %.demote
+
+.demote:
+  call void @llvm.amdgcn.wqm.demote(i1 false)
+  br label %.continue
+
+.continue:
+  %coord1 = fadd float %tex0, %tex1
+  %rtex = call <4 x float> @llvm.amdgcn.image.sample.1d.v4f32.f32(i32 15, float %coord1, <8 x i32> %rsrc, <4 x i32> %sampler, i1 0, i32 0, i32 0) #0
+
+  ret <4 x float> %rtex
+}
+
+
+; GCN-LABEL: {{^}}wqm_deriv:
+; GCN-NEXT: ; %.entry
+; GCN-32: s_mov_b32 [[ORIG:s[0-9]+]], exec_lo
+; GCN-64: s_mov_b64 [[ORIG:s\[[0-9]+:[0-9]+\]]], exec
+; GCN-32: s_wqm_b32 exec_lo, exec_lo
+; GCN-64: s_wqm_b64 exec, exec
+; GCN: ; %.demote0
+; GCN-32-NEXT: s_andn2_b32 [[LIVE:s[0-9]+]], [[ORIG]], exec
+; GCN-64-NEXT: s_andn2_b64 [[LIVE:s\[[0-9]+:[0-9]+\]]], [[ORIG]], exec
+; GCN: ; %.continue0
+; GCN: v_cndmask_b32_e64 [[DST:v[0-9]+]], 1.0, 0, [[LIVE]]
+; GCN-32: s_and_b32 exec_lo, exec_lo, [[LIVE]]
+; GCN-64: s_and_b64 exec, exec, [[LIVE]]
+; GCN: ; %.demote1
+; GCN-32-NEXT: s_mov_b32 exec_lo, 0
+; GCN-64-NEXT: s_mov_b64 exec, 0
+; GCN: ; %.continue1
+; GCN: exp mrt0
+define amdgpu_ps void @wqm_deriv(<2 x float> %input, float %arg, i32 %index) {
+.entry:
+  %p0 = extractelement <2 x float> %input, i32 0
+  %p1 = extractelement <2 x float> %input, i32 1
+  %x0 = call float @llvm.amdgcn.interp.p1(float %p0, i32 immarg 0, i32 immarg 0, i32 %index) #2
+  %x1 = call float @llvm.amdgcn.interp.p2(float %x0, float %p1, i32 immarg 0, i32 immarg 0, i32 %index) #2
+  %argi = fptosi float %arg to i32
+  %cond0 = icmp eq i32 %argi, 0
+  br i1 %cond0, label %.continue0, label %.demote0
+
+.demote0:
+  call void @llvm.amdgcn.wqm.demote(i1 false)
+  br label %.continue0
+
+.continue0:
+  %live = call i1 @llvm.amdgcn.wqm.helper()
+  %live.cond = select i1 %live, i32 0, i32 1065353216
+  %live.v0 = call i32 @llvm.amdgcn.mov.dpp.i32(i32 %live.cond, i32 85, i32 15, i32 15, i1 true)
+  %live.v0f = bitcast i32 %live.v0 to float
+  %live.v1 = call i32 @llvm.amdgcn.mov.dpp.i32(i32 %live.cond, i32 0, i32 15, i32 15, i1 true)
+  %live.v1f = bitcast i32 %live.v1 to float
+  %v0 = fsub float %live.v0f, %live.v1f
+  %v0.wqm = call float @llvm.amdgcn.wqm.f32(float %v0)
+  %cond1 = fcmp oeq float %v0.wqm, 0.000000e+00
+  %cond2 = and i1 %live, %cond1
+  br i1 %cond2, label %.continue1, label %.demote1
+
+.demote1:
+  call void @llvm.amdgcn.wqm.demote(i1 false)
+  br label %.continue1
+
+.continue1:
+  call void @llvm.amdgcn.exp.compr.v2f16(i32 immarg 0, i32 immarg 15, <2 x half> <half 0xH3C00, half 0xH0000>, <2 x half> <half 0xH0000, half 0xH3C00>, i1 immarg true, i1 immarg true) #3
+  ret void
+}
+
+; GCN-LABEL: {{^}}wqm_deriv_loop:
+; GCN-NEXT: ; %.entry
+; GCN-32: s_mov_b32 [[ORIG:s[0-9]+]], exec_lo
+; GCN-64: s_mov_b64 [[ORIG:s\[[0-9]+:[0-9]+\]]], exec
+; GCN-32: s_wqm_b32 exec_lo, exec_lo
+; GCN-64: s_wqm_b64 exec, exec
+; GCN: ; %.demote0
+; GCN-32-NEXT: s_andn2_b32 [[LIVE:s[0-9]+]], [[ORIG]], exec
+; GCN-64-NEXT: s_andn2_b64 [[LIVE:s\[[0-9]+:[0-9]+\]]], [[ORIG]], exec
+; GCN: ; %.continue0.preheader
+; GCN: ; %.continue0
+; GCN: v_cndmask_b32_e64 [[DST:v[0-9]+]], [[SRC:v[0-9]+]], 0, [[LIVE]]
+; GCN: ; %.demote1
+; GCN-32: s_andn2_b32 [[LIVE]], [[LIVE]], exec
+; GCN-64: s_andn2_b64 [[LIVE]], [[LIVE]], exec
+; GCN: ; %.return
+; GCN-32: s_and_b32 exec_lo, exec_lo, [[LIVE]]
+; GCN-64: s_and_b64 exec, exec, [[LIVE]]
+; GCN: exp mrt0
+define amdgpu_ps void @wqm_deriv_loop(<2 x float> %input, float %arg, i32 %index, i32 %limit) {
+.entry:
+  %p0 = extractelement <2 x float> %input, i32 0
+  %p1 = extractelement <2 x float> %input, i32 1
+  %x0 = call float @llvm.amdgcn.interp.p1(float %p0, i32 immarg 0, i32 immarg 0, i32 %index) #2
+  %x1 = call float @llvm.amdgcn.interp.p2(float %x0, float %p1, i32 immarg 0, i32 immarg 0, i32 %index) #2
+  %argi = fptosi float %arg to i32
+  %cond0 = icmp eq i32 %argi, 0
+  br i1 %cond0, label %.continue0, label %.demote0
+
+.demote0:
+  call void @llvm.amdgcn.wqm.demote(i1 false)
+  br label %.continue0
+
+.continue0:
+  %count = phi i32 [ 0, %.entry ], [ 0, %.demote0 ], [ %next, %.continue1 ]
+  %live = call i1 @llvm.amdgcn.wqm.helper()
+  %live.cond = select i1 %live, i32 0, i32 %count
+  %live.v0 = call i32 @llvm.amdgcn.mov.dpp.i32(i32 %live.cond, i32 85, i32 15, i32 15, i1 true)
+  %live.v0f = bitcast i32 %live.v0 to float
+  %live.v1 = call i32 @llvm.amdgcn.mov.dpp.i32(i32 %live.cond, i32 0, i32 15, i32 15, i1 true)
+  %live.v1f = bitcast i32 %live.v1 to float
+  %v0 = fsub float %live.v0f, %live.v1f
+  %v0.wqm = call float @llvm.amdgcn.wqm.f32(float %v0)
+  %cond1 = fcmp oeq float %v0.wqm, 0.000000e+00
+  %cond2 = and i1 %live, %cond1
+  br i1 %cond2, label %.continue1, label %.demote1
+
+.demote1:
+  call void @llvm.amdgcn.wqm.demote(i1 false)
+  br label %.continue1
+
+.continue1:
+  %next = add i32 %count, 1
+  %loop.cond = icmp slt i32 %next, %limit
+  br i1 %loop.cond, label %.continue0, label %.return
+
+.return:
+  call void @llvm.amdgcn.exp.compr.v2f16(i32 immarg 0, i32 immarg 15, <2 x half> <half 0xH3C00, half 0xH0000>, <2 x half> <half 0xH0000, half 0xH3C00>, i1 immarg true, i1 immarg true) #3
+  ret void
+}
+
+declare void @llvm.amdgcn.wqm.demote(i1) #0
+declare i1 @llvm.amdgcn.wqm.helper() #0
+declare void @llvm.amdgcn.exp.f32(i32, i32, float, float, float, float, i1, i1) #0
+declare <4 x float> @llvm.amdgcn.image.sample.1d.v4f32.f32(i32, float, <8 x i32>, <4 x i32>, i1, i32, i32) #1
+declare float @llvm.amdgcn.wqm.f32(float) #1
+declare float @llvm.amdgcn.interp.p1(float, i32 immarg, i32 immarg, i32) #2
+declare float @llvm.amdgcn.interp.p2(float, float, i32 immarg, i32 immarg, i32) #2
+declare void @llvm.amdgcn.exp.compr.v2f16(i32 immarg, i32 immarg, <2 x half>, <2 x half>, i1 immarg, i1 immarg) #3
+declare i32 @llvm.amdgcn.mov.dpp.i32(i32, i32 immarg, i32 immarg, i32 immarg, i1 immarg) #4
+
+attributes #0 = { nounwind }
+attributes #1 = { nounwind readnone }
+attributes #2 = { nounwind readnone speculatable }
+attributes #3 = { inaccessiblememonly nounwind }
+attributes #4 = { convergent nounwind readnone }