AMDGPU: Treat asm as a hazard for all register read-after-write hazards

[arsenm]

Summary

We have to assume there is a hazardous instruction if there's a def
of the appropriate register. We're missing test coverage for some of
these.

[rampitec]

I don't like this idea. Inserting 19 nops just in case is an overkill.

[arsenm]

In D135733#3853303, @rampitec wrote:

I don't like this idea. Inserting 19 nops just in case is an overkill.

The scheduler will still try to avoid the hazard, and only for the relevant registers, so it isn't blindly using the worst case. I don't really care about the performance when using inline asm. It's better to get this correct than continue to waste time debugging these kinds of problems. If people see the nops maybe they'll finally get the message to not use it

[arsenm]

For some of the other hazards (like setreg), I'm thinking of basing it on whether asm is marked sideeffect or not

[b-sumner]

We're in a lose-lose situation here. No matter what we do, someone will complain. But we're are definitely on more solid footing by going for correctness.

[foad]

Could you detect empty inline asms and do this only for the non-empty ones? In graphics we use empty inline asms as some kind of scheduling barrier. (Granted this is just a workaround for problems elsewhere, but I don't really want it broken or penalised just now.)

[arsenm]

Special case empty inline asm. Doesn't work for the device libraries since those are non-empty comments

[foad]

Thanks.

[foad]

In D135733#3855154, @foad wrote:

Could you detect empty inline asms and do this only for the non-empty ones? In graphics we use empty inline asms as some kind of scheduling barrier. (Granted this is just a workaround for problems elsewhere, but I don't really want it broken or penalised just now.)

Sorry I was wrong about the empty strings. LLPC typically uses "; %1" as the string - although we could change that.

[kerbowa]

In D135733#3855154, @foad wrote:

Could you detect empty inline asms and do this only for the non-empty ones? In graphics we use empty inline asms as some kind of scheduling barrier. (Granted this is just a workaround for problems elsewhere, but I don't really want it broken or penalised just now.)

Could you use llvm.amdgcn.sched.barrier instead?

[arsenm]

In D135733#3862393, @kerbowa wrote:

In D135733#3855154, @foad wrote:

Could you detect empty inline asms and do this only for the non-empty ones? In graphics we use empty inline asms as some kind of scheduling barrier. (Granted this is just a workaround for problems elsewhere, but I don't really want it broken or penalised just now.)

Could you use llvm.amdgcn.sched.barrier instead?

No. It's not really a scheduling barrier, it's a workaround for not having convergence tokens. The real problem is hoisting cross lane operations

llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp

for (MachineInstr *MI : EmittedInstrs) {
     ++WaitStates;
 
     if (WaitStates >= Limit)
       break;
   }
   return std::numeric_limits<int>::max();
 }
 
+// Check if this is inline asm, but ignoring asm with empty content. These are
+// currently used as hacks to work around the lack of convergence tokens.
+static bool isNonEmptyInlineAsm(const MachineInstr &MI) {
+  if (!MI.isInlineAsm())
+    return false;
+
+  // TODO: Device libraries using asm hacks with non-empty strings.
+  const char *AsmStr = MI.getOperand(0).getSymbolName();
+  return AsmStr && AsmStr[0];
+}
+
 int GCNHazardRecognizer::getWaitStatesSinceDef(unsigned Reg,
                                                IsHazardFn IsHazardDef,
                                                int Limit) {
   const SIRegisterInfo *TRI = ST.getRegisterInfo();
 
   auto IsHazardFn = [IsHazardDef, TRI, Reg](const MachineInstr &MI) {
-    return IsHazardDef(MI) && MI.modifiesRegister(Reg, TRI);
+    return (IsHazardDef(MI) || isNonEmptyInlineAsm(MI)) &&
+           MI.modifiesRegister(Reg, TRI);
   };
 
   return getWaitStatesSince(IsHazardFn, Limit);
 }
 
 int GCNHazardRecognizer::getWaitStatesSinceSetReg(IsHazardFn IsHazard,
                                                   int Limit) {
   auto IsHazardFn = [IsHazard](const MachineInstr &MI) {
int GCNHazardRecognizer::checkMFMAPadding(MachineInstr *MI) {
   return std::max(0, NeighborMFMAPaddingNeeded);
 }
 
 int GCNHazardRecognizer::checkMAIHazards908(MachineInstr *MI) {
   int WaitStatesNeeded = 0;
   unsigned Opc = MI->getOpcode();
 
   auto IsVALUFn = [](const MachineInstr &MI) {
-    return SIInstrInfo::isVALU(MI) || MI.isInlineAsm();
+    return SIInstrInfo::isVALU(MI);
   };
 
   if (Opc != AMDGPU::V_ACCVGPR_READ_B32_e64) { // MFMA or v_accvgpr_write
     const int LegacyVALUWritesVGPRWaitStates = 2;
     const int VALUWritesExecWaitStates = 4;
     const int MaxWaitStates = 4;
 
     int WaitStatesNeededForUse = VALUWritesExecWaitStates -
for (const MachineOperand &Use : MI->explicit_uses()) {
     const int DMFMA16x16WritesVGPROverlappedMFMASrcABWaitStates = 11;
     const int DMFMA4x4WritesVGPRFullSrcCWaitStates = 4;
     const int GFX940_SMFMA4x4WritesVGPRFullSrcCWaitStates = 2;
     const int MaxWaitStates = 19;
 
     if (!Use.isReg())
       continue;
     Register Reg = Use.getReg();
-    bool FullReg;
+    bool FullReg = false;
     const MachineInstr *MI1;
 
     auto IsOverlappedMFMAFn = [Reg, &FullReg, &MI1,
                                this](const MachineInstr &MI) {
-      if (!SIInstrInfo::isMFMA(MI))
-        return false;
+      if (SIInstrInfo::isMFMA(MI)) {
+        MI1 = &MI;
         Register DstReg = MI.getOperand(0).getReg();
         FullReg = (DstReg == Reg);
-      MI1 = &MI;
         return TRI.regsOverlap(DstReg, Reg);
+      }
+
+      if (MI.isInlineAsm()) {
+        MI1 = &MI;
+        for (const MachineOperand &MO : MI.operands()) {
+          if (!MO.isReg() || !MO.isDef())
+            continue;
+
+          Register DstReg = MO.getReg();
+          if (TRI.regsOverlap(DstReg, Reg))
+            return true;
+        }
+      }
+
+      return false;
     };
 
     WaitStatesNeededForUse = LegacyVALUNotDotWritesVGPRWaitStates -
       getWaitStatesSinceDef(Reg, IsLegacyVALUNotDotFn, MaxWaitStates);
     WaitStatesNeeded = std::max(WaitStatesNeeded, WaitStatesNeededForUse);
 
     int NumWaitStates =
         getWaitStatesSinceDef(Reg, IsOverlappedMFMAFn, MaxWaitStates);
if (OpNo == SrcCIdx) {
           if (!isXDL(ST, *MI))
             NeedWaitStates = DMFMA16x16WritesVGPROverlappedSrcCWaitStates;
           break;
         case AMDGPU::V_MFMA_F64_4X4X4F64_e64:
         case AMDGPU::V_MFMA_F64_4X4X4F64_vgprcd_e64:
           if (!isXDL(ST, *MI))
             NeedWaitStates = DMFMA4x4WritesVGPROverlappedSrcCWaitStates;
           break;
+        case AMDGPU::INLINEASM:
+        case AMDGPU::INLINEASM_BR:
+          NeedWaitStates = MaxWaitStates;
+          break;
         default:
           if (ST.hasGFX940Insts() && isXDL(ST, *MI) && !isXDL(ST, *MI1))
             break;
           switch (TSchedModel.computeInstrLatency(MI1)) {
           case 2:
             NeedWaitStates = ST.hasGFX940Insts()
               ? isXDL(ST, *MI1)
                 ? GFX940_XDL2PassWritesVGPROverlappedSMFMASrcCWaitStates
if (OpNo == SrcCIdx) {
       case AMDGPU::V_MFMA_F64_16X16X4F64_mac_e64:
       case AMDGPU::V_MFMA_F64_16X16X4F64_mac_vgprcd_e64:
         NeedWaitStates = DMFMA16x16WritesVGPROverlappedMFMASrcABWaitStates;
         break;
       case AMDGPU::V_MFMA_F64_4X4X4F64_e64:
       case AMDGPU::V_MFMA_F64_4X4X4F64_vgprcd_e64:
         NeedWaitStates = DMFMA4x4WritesVGPROverlappedMFMASrcABWaitStates;
         break;
+      case AMDGPU::INLINEASM:
+      case AMDGPU::INLINEASM_BR:
+        NeedWaitStates = MaxWaitStates;
+        break;
       default:
         switch (TSchedModel.computeInstrLatency(MI1)) {
         case 2:
           NeedWaitStates = ST.hasGFX940Insts()
             ? isXDL(ST, *MI1)
               ? GFX940_XDL2PassWritesVGPROverlappedSrcABWaitStates
               : GFX940_SMFMA2PassWritesVGPROverlappedSrcABWaitStates
             : SMFMA4x4WritesVGPROverlappedSrcABWaitStates;

llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.writelane.ll

 define amdgpu_ps float @test_writelane_m0_s_v(i32 inreg %lane, i32 %vdst.in) #0 {
 ; GFX7-LABEL: test_writelane_m0_s_v:
 ; GFX7:       ; %bb.0:
 ; GFX7-NEXT:    ;;#ASMSTART
 ; GFX7-NEXT:    s_mov_b32 m0, -1
 ; GFX7-NEXT:    ;;#ASMEND
 ; GFX7-NEXT:    s_mov_b32 s0, m0
 ; GFX7-NEXT:    s_mov_b32 m0, s2
+; GFX7-NEXT:    s_nop 1
 ; GFX7-NEXT:    v_writelane_b32 v0, s0, m0
 ; GFX7-NEXT:    ; return to shader part epilog
 ;
 ; GFX8-LABEL: test_writelane_m0_s_v:
 ; GFX8:       ; %bb.0:
 ; GFX8-NEXT:    ;;#ASMSTART
 ; GFX8-NEXT:    s_mov_b32 m0, -1
 ; GFX8-NEXT:    ;;#ASMEND
 ; GFX8-NEXT:    s_mov_b32 s0, m0
 ; GFX8-NEXT:    s_mov_b32 m0, s2
+; GFX8-NEXT:    s_nop 1
 ; GFX8-NEXT:    v_writelane_b32 v0, s0, m0
 ; GFX8-NEXT:    ; return to shader part epilog
 ;
 ; GFX10-LABEL: test_writelane_m0_s_v:
 ; GFX10:       ; %bb.0:
 ; GFX10-NEXT:    ;;#ASMSTART
 ; GFX10-NEXT:    s_mov_b32 m0, -1
 ; GFX10-NEXT:    ;;#ASMEND
 ; GFX10-NEXT:    v_writelane_b32 v0, m0, s2
 ; GFX10-NEXT:    ; return to shader part epilog
   %m0 = call i32 asm "s_mov_b32 m0, -1", "={m0}"()
   %writelane = call i32 @llvm.amdgcn.writelane(i32 %m0, i32 %lane, i32 %vdst.in)
   %writelane.cast = bitcast i32 %writelane to float
   ret float %writelane.cast
 }
 
 define amdgpu_ps float @test_writelane_s_m0_v(i32 inreg %data, i32 %vdst.in) #0 {
 ; GFX7-LABEL: test_writelane_s_m0_v:
 ; GFX7:       ; %bb.0:
 ; GFX7-NEXT:    ;;#ASMSTART
 ; GFX7-NEXT:    s_mov_b32 m0, -1
 ; GFX7-NEXT:    ;;#ASMEND
+; GFX7-NEXT:    s_nop 3
 ; GFX7-NEXT:    v_writelane_b32 v0, s2, m0
 ; GFX7-NEXT:    ; return to shader part epilog
 ;
 ; GFX8-LABEL: test_writelane_s_m0_v:
 ; GFX8:       ; %bb.0:
 ; GFX8-NEXT:    ;;#ASMSTART
 ; GFX8-NEXT:    s_mov_b32 m0, -1
 ; GFX8-NEXT:    ;;#ASMEND
+; GFX8-NEXT:    s_nop 3
 ; GFX8-NEXT:    v_writelane_b32 v0, s2, m0
 ; GFX8-NEXT:    ; return to shader part epilog
 ;
 ; GFX10-LABEL: test_writelane_s_m0_v:
 ; GFX10:       ; %bb.0:
 ; GFX10-NEXT:    ;;#ASMSTART
 ; GFX10-NEXT:    s_mov_b32 m0, -1
 ; GFX10-NEXT:    ;;#ASMEND

llvm/test/CodeGen/AMDGPU/agpr-copy-no-free-registers.ll

 ; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx908 < %s | FileCheck -check-prefix=GFX908 %s
 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a < %s | FileCheck -check-prefix=GFX90A %s
 
 ; This testcase would fail on GFX908 due to not having a free VGPR available to
 ; copy between AGPRs.
 define void @no_free_vgprs_at_agpr_to_agpr_copy(float %v0, float %v1) #0 {
 ; GFX908-LABEL: no_free_vgprs_at_agpr_to_agpr_copy:
 ; GFX908:       ; %bb.0:
 ; GFX908-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
 ; GFX908-NEXT:    v_mov_b32_e32 v32, v1
 ; GFX908-NEXT:    v_mov_b32_e32 v33, v0
 ; GFX908-NEXT:    ;;#ASMSTART
 ; GFX908-NEXT:    ; def v[0:31] a[0:15]
 ; GFX908-NEXT:    ;;#ASMEND
+; GFX908-NEXT:    s_nop 7
+; GFX908-NEXT:    s_nop 7
+; GFX908-NEXT:    s_nop 1
 ; GFX908-NEXT:    v_accvgpr_read_b32 v34, a15
 ; GFX908-NEXT:    s_nop 1
 ; GFX908-NEXT:    v_accvgpr_write_b32 a31, v34
 ; GFX908-NEXT:    v_accvgpr_read_b32 v34, a14
 ; GFX908-NEXT:    s_nop 1
 ; GFX908-NEXT:    v_accvgpr_write_b32 a30, v34
 ; GFX908-NEXT:    v_accvgpr_read_b32 v34, a13
 ; GFX908-NEXT:    s_nop 1

 ; GFX908-NEXT:    buffer_store_dword v34, off, s[0:3], s32 offset:36 ; 4-byte Folded Spill
 ; GFX908-NEXT:    v_accvgpr_read_b32 v34, a10 ; Reload Reuse
 ; GFX908-NEXT:    s_nop 1
 ; GFX908-NEXT:    buffer_store_dword v34, off, s[0:3], s32 offset:40 ; 4-byte Folded Spill
 ; GFX908-NEXT:    v_accvgpr_read_b32 v34, a15 ; Reload Reuse
 ; GFX908-NEXT:    ;;#ASMSTART
 ; GFX908-NEXT:    ; copy
 ; GFX908-NEXT:    ;;#ASMEND
+; GFX908-NEXT:    s_nop 7
+; GFX908-NEXT:    s_nop 7
+; GFX908-NEXT:    s_nop 1
 ; GFX908-NEXT:    v_accvgpr_read_b32 v34, a1
 ; GFX908-NEXT:    s_nop 1
 ; GFX908-NEXT:    v_accvgpr_write_b32 a16, v34
 ; GFX908-NEXT:    buffer_load_dword v34, off, s[0:3], s32 ; 4-byte Folded Reload
 ; GFX908-NEXT:    s_waitcnt vmcnt(0)
 ; GFX908-NEXT:    v_accvgpr_write_b32 a0, v34 ; Reload Reuse
 ; GFX908-NEXT:    buffer_load_dword v34, off, s[0:3], s32 offset:4 ; 4-byte Folded Reload
 ; GFX908-NEXT:    s_waitcnt vmcnt(0)

 ; GFX908-NEXT:    v_accvgpr_write_b32 a11, v38 ; Reload Reuse
 ; GFX908-NEXT:    v_accvgpr_write_b32 a12, v37 ; Reload Reuse
 ; GFX908-NEXT:    v_accvgpr_write_b32 a13, v36 ; Reload Reuse
 ; GFX908-NEXT:    v_accvgpr_write_b32 a14, v35 ; Reload Reuse
 ; GFX908-NEXT:    v_accvgpr_write_b32 a15, v34 ; Reload Reuse
 ; GFX908-NEXT:    ;;#ASMSTART
 ; GFX908-NEXT:    ; copy
 ; GFX908-NEXT:    ;;#ASMEND
+; GFX908-NEXT:    s_nop 7
+; GFX908-NEXT:    s_nop 7
+; GFX908-NEXT:    s_nop 1
 ; GFX908-NEXT:    v_accvgpr_read_b32 v32, a2
 ; GFX908-NEXT:    s_nop 1
 ; GFX908-NEXT:    v_accvgpr_write_b32 a3, v32
 ; GFX908-NEXT:    ;;#ASMSTART
 ; GFX908-NEXT:    ; use a3 v[0:31]
 ; GFX908-NEXT:    ;;#ASMEND
 ; GFX908-NEXT:    s_setpc_b64 s[30:31]
 ;

 ; GFX908-LABEL: v32_asm_def_use:
 ; GFX908:       ; %bb.0:
 ; GFX908-NEXT:    s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
 ; GFX908-NEXT:    v_mov_b32_e32 v33, v1
 ; GFX908-NEXT:    v_mov_b32_e32 v34, v0
 ; GFX908-NEXT:    ;;#ASMSTART
 ; GFX908-NEXT:    ; def v[0:31] a[0:15]
 ; GFX908-NEXT:    ;;#ASMEND
-; GFX908-NEXT:    v_accvgpr_read_b32 v35, a15
 ; GFX908-NEXT:    ;;#ASMSTART
 ; GFX908-NEXT:    ; def v32
 ; GFX908-NEXT:    ;;#ASMEND
+; GFX908-NEXT:    s_nop 7
+; GFX908-NEXT:    s_nop 7
+; GFX908-NEXT:    s_nop 1
+; GFX908-NEXT:    v_accvgpr_read_b32 v35, a15
 ; GFX908-NEXT:    s_nop 1
 ; GFX908-NEXT:    v_accvgpr_write_b32 a31, v35
 ; GFX908-NEXT:    v_accvgpr_read_b32 v35, a14
 ; GFX908-NEXT:    s_nop 1
 ; GFX908-NEXT:    v_accvgpr_write_b32 a30, v35
 ; GFX908-NEXT:    v_accvgpr_read_b32 v35, a13
 ; GFX908-NEXT:    s_nop 1
 ; GFX908-NEXT:    v_accvgpr_write_b32 a29, v35

 ; GFX908-NEXT:    s_nop 1
 ; GFX908-NEXT:    v_accvgpr_write_b32 a17, v35
 ; GFX908-NEXT:    v_accvgpr_read_b32 v35, a0
 ; GFX908-NEXT:    s_nop 1
 ; GFX908-NEXT:    v_accvgpr_write_b32 a16, v35
 ; GFX908-NEXT:    ;;#ASMSTART
 ; GFX908-NEXT:    ; copy
 ; GFX908-NEXT:    ;;#ASMEND
+; GFX908-NEXT:    s_nop 7
+; GFX908-NEXT:    s_nop 7
+; GFX908-NEXT:    s_nop 1
 ; GFX908-NEXT:    v_accvgpr_read_b32 v35, a1
 ; GFX908-NEXT:    v_mfma_f32_16x16x1f32 a[0:15], v34, v33, a[16:31]
 ; GFX908-NEXT:    s_nop 0
 ; GFX908-NEXT:    v_accvgpr_write_b32 a32, v35
 ; GFX908-NEXT:    ;;#ASMSTART
 ; GFX908-NEXT:    ; copy
 ; GFX908-NEXT:    ;;#ASMEND
 ; GFX908-NEXT:    s_nop 7
+; GFX908-NEXT:    s_nop 7
+; GFX908-NEXT:    s_nop 1
 ; GFX908-NEXT:    v_accvgpr_read_b32 v33, a2
 ; GFX908-NEXT:    s_nop 1
 ; GFX908-NEXT:    v_accvgpr_write_b32 a3, v33
 ; GFX908-NEXT:    ;;#ASMSTART
 ; GFX908-NEXT:    ; use a3 v[0:31]
 ; GFX908-NEXT:    ;;#ASMEND
 ; GFX908-NEXT:    ;;#ASMSTART
 ; GFX908-NEXT:    ; use v32

 ; GFX908-NEXT:    buffer_store_dword v34, off, s[0:3], s32 offset:36 ; 4-byte Folded Spill
 ; GFX908-NEXT:    v_accvgpr_read_b32 v34, a10 ; Reload Reuse
 ; GFX908-NEXT:    s_nop 1
 ; GFX908-NEXT:    buffer_store_dword v34, off, s[0:3], s32 offset:40 ; 4-byte Folded Spill
 ; GFX908-NEXT:    v_accvgpr_read_b32 v34, a15 ; Reload Reuse
 ; GFX908-NEXT:    ;;#ASMSTART
 ; GFX908-NEXT:    ; copy
 ; GFX908-NEXT:    ;;#ASMEND
+; GFX908-NEXT:    s_nop 7
+; GFX908-NEXT:    s_nop 7
+; GFX908-NEXT:    s_nop 1
 ; GFX908-NEXT:    v_accvgpr_read_b32 v34, a1
 ; GFX908-NEXT:    s_nop 1
 ; GFX908-NEXT:    v_accvgpr_write_b32 a32, v34
 ; GFX908-NEXT:    buffer_load_dword v34, off, s[0:3], s32 ; 4-byte Folded Reload
 ; GFX908-NEXT:    s_waitcnt vmcnt(0)
 ; GFX908-NEXT:    v_accvgpr_write_b32 a0, v34 ; Reload Reuse
 ; GFX908-NEXT:    buffer_load_dword v34, off, s[0:3], s32 offset:4 ; 4-byte Folded Reload
 ; GFX908-NEXT:    s_waitcnt vmcnt(0)

 ; GFX908-NEXT:    v_accvgpr_write_b32 a11, v38 ; Reload Reuse
 ; GFX908-NEXT:    v_accvgpr_write_b32 a12, v37 ; Reload Reuse
 ; GFX908-NEXT:    v_accvgpr_write_b32 a13, v36 ; Reload Reuse
 ; GFX908-NEXT:    v_accvgpr_write_b32 a14, v35 ; Reload Reuse
 ; GFX908-NEXT:    v_accvgpr_write_b32 a15, v34 ; Reload Reuse
 ; GFX908-NEXT:    ;;#ASMSTART
 ; GFX908-NEXT:    ; copy
 ; GFX908-NEXT:    ;;#ASMEND
+; GFX908-NEXT:    s_nop 7
+; GFX908-NEXT:    s_nop 7
+; GFX908-NEXT:    s_nop 1
 ; GFX908-NEXT:    v_accvgpr_read_b32 v32, a2
 ; GFX908-NEXT:    s_nop 1
 ; GFX908-NEXT:    v_accvgpr_write_b32 a3, v32
 ; GFX908-NEXT:    ;;#ASMSTART
 ; GFX908-NEXT:    ; use a3 v[0:31]
 ; GFX908-NEXT:    ;;#ASMEND
 ; GFX908-NEXT:    s_setpc_b64 s[30:31]
 ;

llvm/test/CodeGen/AMDGPU/inserted-wait-states.mir

-# RUN: llc -march=amdgcn -mcpu=tahiti -run-pass post-RA-hazard-rec  %s -o - | FileCheck %s -check-prefixes=GCN,SICI
+# RUN: llc -march=amdgcn -mcpu=tahiti -run-pass post-RA-hazard-rec  %s -o - | FileCheck %s -check-prefixes=GCN,SICI,SI
 # RUN: llc -march=amdgcn -mcpu=hawaii -run-pass post-RA-hazard-rec  %s -o - | FileCheck %s -check-prefixes=GCN,CIVI,SICI
-# RUN: llc -march=amdgcn -mcpu=fiji -run-pass post-RA-hazard-rec  %s -o - | FileCheck %s -check-prefixes=GCN,CIVI,VI
-# RUN: llc -march=amdgcn -mcpu=gfx900 -run-pass post-RA-hazard-rec  %s -o - | FileCheck %s -check-prefixes=GCN,CIVI,VI,GFX9
+# RUN: llc -march=amdgcn -mcpu=fiji -run-pass post-RA-hazard-rec  %s -o - | FileCheck %s -check-prefixes=GCN,CIVI,VI,GFX89
+# RUN: llc -march=amdgcn -mcpu=gfx900 -run-pass post-RA-hazard-rec  %s -o - | FileCheck %s -check-prefixes=GCN,CIVI,VI,GFX9,GFX89
 
 --- |
   define amdgpu_kernel void @div_fmas() { ret void }
   define amdgpu_kernel void @s_getreg() { ret void }
   define amdgpu_kernel void @s_setreg() { ret void }
   define amdgpu_kernel void @vmem_gt_8dw_store() { ret void }
   define amdgpu_kernel void @readwrite_lane() { ret void }
   define amdgpu_kernel void @rfe() { ret void }
   define amdgpu_kernel void @s_movrel() { ret void }
   define amdgpu_kernel void @v_interp() { ret void }
   define amdgpu_kernel void @dpp() { ret void }
+  define amdgpu_kernel void @salu_write_buffer_smrd() { ret void }
+  define amdgpu_kernel void  @vmem_read_sgpr_valu_def() { ret void }
 ...
 ---
 # GCN-LABEL: name: div_fmas
 
 # GCN-LABEL: bb.0:
 # GCN: S_MOV_B64
 # GCN-NOT: S_NOP
 # GCN: V_DIV_FMAS
 
 # GCN-LABEL: bb.1:
 # GCN: V_CMP_EQ_I32
 # GCN: S_NOP 3
 # GCN: V_DIV_FMAS_F32
 
 # GCN-LABEL: bb.2:
 # GCN: V_CMP_EQ_I32
 # GCN: S_NOP 3
 # GCN: V_DIV_FMAS_F32
 
 # GCN-LABEL: bb.3:
 # GCN: V_DIV_SCALE_F32
 # GCN: S_NOP 3
 # GCN: V_DIV_FMAS_F32
+
+# GCN-LABEL: bb.4:
+# GCN: INLINEASM
+# GCN-NEXT: S_NOP 3
+# GCN-NEXT: V_DIV_FMAS_F32
+
+# GCN-LABEL: bb.5:
+# GCN: INLINEASM
+# GCN-NEXT: V_DIV_FMAS_F32
+
 name: div_fmas
 
 body: |
   bb.0:
     $vcc = S_MOV_B64 0
     $vgpr0 = V_DIV_FMAS_F32_e64 0, $vgpr1, 0, $vgpr2, 0, $vgpr3, 0, 0, implicit $mode, implicit $vcc, implicit $exec
     S_BRANCH %bb.1
 
   bb.1:
     implicit $vcc = V_CMP_EQ_I32_e32 $vgpr1, $vgpr2, implicit $exec
     $vgpr0 = V_DIV_FMAS_F32_e64 0, $vgpr1, 0, $vgpr2, 0, $vgpr3, 0, 0, implicit $mode, implicit $vcc, implicit $exec
     S_BRANCH %bb.2
 
   bb.2:
     $vcc = V_CMP_EQ_I32_e64 $vgpr1, $vgpr2, implicit $exec
     $vgpr0 = V_DIV_FMAS_F32_e64 0, $vgpr1, 0, $vgpr2, 0, $vgpr3, 0, 0, implicit $mode, implicit $vcc, implicit $exec
     S_BRANCH %bb.3
 
   bb.3:
     $vgpr4, $vcc = V_DIV_SCALE_F32_e64 0, $vgpr1, 0, $vgpr1, 0, $vgpr3, 0, 0, implicit $mode, implicit $exec
     $vgpr0 = V_DIV_FMAS_F32_e64 0, $vgpr1, 0, $vgpr2, 0, $vgpr3, 0, 0, implicit $mode, implicit $vcc, implicit $exec
-    S_ENDPGM 0
 
-...
+  bb.4:
+    INLINEASM &"; non-empy asm", 1, 851978, def $vcc
+    $vgpr0 = V_DIV_FMAS_F32_e64 0, $vgpr1, 0, $vgpr2, 0, $vgpr3, 0, 0, implicit $mode, implicit $vcc, implicit $exec
 
+    ; No hazard here because the asm is empty
+  bb.5:
+    INLINEASM &"", 1, 851978, def $vcc
+    $vgpr0 = V_DIV_FMAS_F32_e64 0, $vgpr1, 0, $vgpr2, 0, $vgpr3, 0, 0, implicit $mode, implicit $vcc, implicit $exec
+    S_ENDPGM 0
 ...
 ---
 # GCN-LABEL: name: s_getreg
 
 # GCN-LABEL: bb.0:
 # GCN: S_SETREG
 # GCN: S_NOP 1
 # GCN: S_GETREG
bb.0:
     $vgpr1 = V_MOV_B32_dpp $vgpr1, $vgpr0, 0, 15, 15, 0, implicit $exec
     S_BRANCH %bb.1
 
   bb.1:
     implicit $exec, implicit $vcc = V_CMPX_EQ_I32_e32 $vgpr0, $vgpr1, implicit $exec
     $vgpr3 = V_MOV_B32_dpp $vgpr3, $vgpr0, 0, 15, 15, 0, implicit $exec
     S_ENDPGM 0
 ...
+---
+
+# GCN-LABEL: name: salu_write_buffer_smrd
+# GCN-LABEL: bb.0:
+# GCN: $sgpr15 = S_MOV_B32 123
+# SI-NEXT:  S_NOP 3
+# GCN-NEXT: S_BUFFER_LOAD_DWORD_SGPR
+
+# GCN-LABEL: bb.1:
+# GCN: INLINEASM
+# SI-NEXT:  S_NOP 3
+# GCN-NEXT: S_BUFFER_LOAD_DWORD_SGPR
+
+name: salu_write_buffer_smrd
+body: |
+  bb.0:
+    liveins: $sgpr12_sgpr13_sgpr14_sgpr15
+    $sgpr15 = S_MOV_B32 123
+    $sgpr17 = S_BUFFER_LOAD_DWORD_SGPR $sgpr12_sgpr13_sgpr14_sgpr15, $sgpr16, 0
+
+  bb.1:
+    liveins: $sgpr17
+    INLINEASM &"; foo", 1, 851978, def $sgpr14
+    $sgpr18 = S_BUFFER_LOAD_DWORD_SGPR $sgpr12_sgpr13_sgpr14_sgpr15, $sgpr16, 0
+    S_ENDPGM 0, implicit $sgpr17, implicit $sgpr18
+...
+
+---
+
+# GCN-LABEL: name: vmem_read_sgpr_valu_def
+# GCN-LABEL: bb.0:
+# GCN: V_READLANE_B32
+# GFX89-NEXT: S_NOP 4
+# GCN-NEXT: BUFFER_STORE_DWORD_OFFSET
+
+# GCN-LABEL: bb.1:
+# GCN: INLINEASM
+# GFX89-NEXT: S_NOP 4
+# GCN-NEXT: BUFFER_STORE_DWORD_OFFSET
+
+name: vmem_read_sgpr_valu_def
+body: |
+  bb.0:
+    liveins: $sgpr12_sgpr13_sgpr14_sgpr15, $vgpr0
+
+    $sgpr6 = V_READLANE_B32 $vgpr0, 0
+    BUFFER_STORE_DWORD_OFFSET $vgpr3, $sgpr0_sgpr1_sgpr2_sgpr3, $sgpr6, 0, 0, 0, 0, implicit $exec
+
+  bb.1:
+    INLINEASM &"; foo", 1, 851978, def $sgpr7
+    BUFFER_STORE_DWORD_OFFSET $vgpr3, $sgpr0_sgpr1_sgpr2_sgpr3, $sgpr7, 0, 0, 0, 0, implicit $exec
+    S_ENDPGM 0
+...

llvm/test/CodeGen/AMDGPU/mai-hazards-gfx90a.mir

bb.0:
     $vgpr0 = V_MOV_B32_e32 0, implicit $exec
 
   bb.1:
 
   bb.2:
     $agpr0_agpr1 = V_MFMA_F64_4X4X4F64_e64 $agpr0_agpr1, $agpr0_agpr1, $agpr0_agpr1, 0, 0, 0, implicit $mode, implicit $exec
     BUFFER_STORE_DWORDX2_OFFEN_exact $vgpr2_vgpr3, $vgpr0, $sgpr0_sgpr1_sgpr2_sgpr3, 0, 0, 0, 0, 0, implicit $exec
 ...
+
+---
+
+# GCN-LABEL: name: mfma_write_agpr_mfma_read_overlap_asm
+# GCN:      INLINEASM
+# GCN-NEXT: S_NOP 7
+# GCN-NEXT: S_NOP 7
+# GCN-NEXT: S_NOP 2
+# GCN-NEXT: V_MFMA
+name:            mfma_write_agpr_mfma_read_overlap_asm
+body:             |
+  bb.0:
+    INLINEASM &"; def $0", 1 /* sideeffect attdialect */, 1966090 /* regdef:AReg_128 */, def $agpr1_agpr2_agpr3_agpr4
+    $agpr0_agpr1_agpr2_agpr3 = V_MFMA_F32_4X4X1F32_e64 killed $vgpr1, killed $vgpr0, killed $agpr0_agpr1_agpr2_agpr3, 0, 0, 0, implicit $mode, implicit $exec
+...
+
+---
+
+# GCN-LABEL: name: mfma_write_agpr_mfma_read_overlap_asm_second_def
+# GCN:      INLINEASM
+# GCN-NEXT: S_NOP 7
+# GCN-NEXT: S_NOP 7
+# GCN-NEXT: S_NOP 2
+# GCN-NEXT: V_MFMA
+name:            mfma_write_agpr_mfma_read_overlap_asm_second_def
+body:             |
+  bb.0:
+    INLINEASM &"", 1 /* sideeffect attdialect */, 3211274 /* regdef:VReg_64 */, def $vgpr0_vgpr1, 3145738 /* regdef:AReg_64 */, def $agpr3_agpr4
+    $agpr0_agpr1_agpr2_agpr3 = V_MFMA_F32_4X4X1F32_e64 killed $vgpr1, killed $vgpr0, killed $agpr0_agpr1_agpr2_agpr3, 0, 0, 0, implicit $mode, implicit $exec
+...
+
+...
+# GCN-LABEL: name: sgemm4x4_mfma_write_agpr_mfma_read_overlap_asm
+# GCN:      INLINEASM
+# GCN-NEXT: S_NOP 7
+# GCN-NEXT: S_NOP 7
+# GCN-NEXT: S_NOP 2
+# GCN-NEXT: V_MFMA
+name:            sgemm4x4_mfma_write_agpr_mfma_read_overlap_asm
+body:             |
+  bb.0:
+    INLINEASM &"", 1 /* sideeffect attdialect */, 5832714 /* regdef:VReg_128 */, def $vgpr0_vgpr1_vgpr2_vgpr3, 5767178 /* regdef:AReg_128 */, def $agpr2_agpr3_agpr4_agpr5
+  $agpr0_agpr1_agpr2_agpr3 = V_MFMA_F32_4X4X1F32_e64 $vgpr1, $vgpr0, $agpr0_agpr1_agpr2_agpr3, 0, 0, 0, implicit $mode, implicit $exec
+...

llvm/test/CodeGen/AMDGPU/mai-hazards.mir

 name:            mfma_write_agpr_mfma_read_overlap
 body:             |
   bb.0:
     $agpr1_agpr2_agpr3_agpr4 = V_MFMA_F32_4X4X1F32_e64 killed $vgpr1, killed $vgpr0, killed $agpr0_agpr1_agpr2_agpr3, 0, 0, 0, implicit $mode, implicit $exec
     $agpr0_agpr1_agpr2_agpr3 = V_MFMA_F32_4X4X1F32_e64 killed $vgpr1, killed $vgpr0, killed $agpr0_agpr1_agpr2_agpr3, 0, 0, 0, implicit $mode, implicit $exec
 ...
 ---
 
+# GCN-LABEL: name: mfma_write_agpr_mfma_read_overlap_asm
+# GCN:      INLINEASM
+# GCN-NEXT: S_NOP 1
+# GCN-NEXT: V_MFMA
+name:            mfma_write_agpr_mfma_read_overlap_asm
+body:             |
+  bb.0:
+    INLINEASM &"; def $0", 1 /* sideeffect attdialect */, 1966090 /* regdef:AReg_128 */, def $agpr1_agpr2_agpr3_agpr4
+    $agpr0_agpr1_agpr2_agpr3 = V_MFMA_F32_4X4X1F32_e64 killed $vgpr1, killed $vgpr0, killed $agpr0_agpr1_agpr2_agpr3, 0, 0, 0, implicit $mode, implicit $exec
+...
+---
+
 # GCN-LABEL: name: mfma_write_agpr_mfma_read_partial
 # GCN:      V_MFMA
 # GCN-NEXT: S_NOP 1
 # GCN-NEXT: V_MFMA
 name:            mfma_write_agpr_mfma_read_partial
 body:             |
   bb.0:
     $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15 = V_MFMA_F32_16X16X1F32_e64 killed $vgpr1, killed $vgpr0, killed $agpr0_agpr1_agpr2_agpr3_agpr4_agpr5_agpr6_agpr7_agpr8_agpr9_agpr10_agpr11_agpr12_agpr13_agpr14_agpr15, 0, 0, 0, implicit $mode, implicit $exec

 # GCN-NEXT: FLAT_LOAD_DWORD
 name:            valu_write_vgpr_accvgpr_read_load_2_and_3_depend
 body:             |
   bb.0:
     $vgpr0 = V_MOV_B32_e32 1, implicit $exec
     $vgpr2 = V_ACCVGPR_READ_B32_e64 killed $agpr0, implicit $exec
     $vgpr4 = FLAT_LOAD_DWORD $vgpr2_vgpr3, 0, 0, implicit $exec, implicit $flat_scr
 ...
----