Guess the constraints instead of using user-provided hints

Author: linuxrocks123

llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp

@@ -2392,13 +2392,56 @@ bool SchedGroup::canAddMI(const MachineInstr &MI) const {
     Result = false;
 
   else if (MI.isInlineAsm()) {
+    auto &TRI = TII->getRegisterInfo();
+    auto &MRI = MI.getParent()->getParent()->getRegInfo();
+    bool SGPR_used = false, VGPR_used = false, VMFMA_used = false,
+         MayLoad = MI.mayLoad(), MayStore = MI.mayStore();
+    for (const MachineOperand &Operand : MI.operands())
+      if (Operand.isReg()) {
+        auto &RegClass = *TRI.getRegClassForOperandReg(MRI, Operand);
+        if (TRI.isVGPRClass(&RegClass))
+          VGPR_used = true;
+        if (TRI.isAGPRClass(&RegClass) || TRI.getRegSizeInBits(RegClass) > 128)
+          VMFMA_used = true;
+        if (TRI.isSGPRClass(&RegClass))
+          SGPR_used = true;
+      }
+
+    unsigned long InlineAsmMask = 0;
+    if (VGPR_used && !VMFMA_used && !MayLoad && !MayStore)
+      InlineAsmMask |= (unsigned long)SchedGroupMask::VALU;
+    if (VMFMA_used)
+      InlineAsmMask |= (unsigned long)SchedGroupMask::MFMA;
+    if (VGPR_used && MayLoad)
+      InlineAsmMask |= (unsigned long)SchedGroupMask::VMEM_READ;
+    if (VGPR_used && MayStore)
+      InlineAsmMask |= (unsigned long)SchedGroupMask::VMEM_WRITE;
+    if (!VGPR_used && MayLoad)
+      InlineAsmMask |= (unsigned long)SchedGroupMask::DS_READ;
+    if (!VGPR_used && MayStore)
+      InlineAsmMask |= (unsigned long)SchedGroupMask::DS_WRITE;
+    if (InlineAsmMask & (unsigned long)SchedGroupMask::VALU ||
+        InlineAsmMask & (unsigned long)SchedGroupMask::SALU)
+      InlineAsmMask |= (unsigned long)SchedGroupMask::ALU;
+    if (InlineAsmMask & (unsigned long)SchedGroupMask::DS_READ ||
+        InlineAsmMask & (unsigned long)SchedGroupMask::DS_WRITE)
+      InlineAsmMask |= (unsigned long)SchedGroupMask::DS;
+    if (InlineAsmMask & (unsigned long)SchedGroupMask::VMEM_READ ||
+        InlineAsmMask & (unsigned long)SchedGroupMask::VMEM_WRITE)
+      InlineAsmMask |= (unsigned long)SchedGroupMask::VMEM;
+
+    Result = ((unsigned long)SGMask & InlineAsmMask) != 0;
+
+    // Original implementation
+#if 0
     StringRef Text = MI.getOperand(0).getSymbolName();
     if (Text.find("SGMASK:") != std::string::npos) {
       Text = Text.substr(Text.find("SGMASK:") + strlen("SGMASK:"));
       Text = Text.substr(0, Text.find_first_of(" \t\r\n"));
       unsigned long InlineAsmMask = std::stoul(Text.str(), nullptr, 0);
       Result = ((unsigned long)SGMask & InlineAsmMask) != 0;
     }
+#endif
   }
 
   else if (((SGMask & SchedGroupMask::ALU) != SchedGroupMask::NONE) &&

llvm/test/CodeGen/AMDGPU/inlineasm-sgmask.ll

@@ -28,20 +28,23 @@ if.then:                                          ; preds = %entry
   %mul3 = shl nsw i32 %add, 2
   %idx.ext4 = sext i32 %mul3 to i64
   %add.ptr5 = getelementptr inbounds float, ptr addrspace(1) %from.coerce, i64 %idx.ext4
-  %2 = load <4 x float>, ptr addrspace(1) %add.ptr5, align 16, !tbaa !0
-  %3 = extractelement <4 x float> %2, i64 3
-  %4 = extractelement <4 x float> %2, i64 0
-  %5 = tail call contract noundef float asm "v_add_f32_e32 $0, $1, $2 ; SGMASK:0x1", "=v,v,v"(float %3, float %4) #3, !srcloc !3
-  %6 = extractelement <4 x float> %2, i64 1
-  %7 = extractelement <4 x float> %2, i64 2
+  %2 = load <2 x float>, ptr addrspace(1) %add.ptr5, align 16, !tbaa !0
+  %a20 = add i64 %idx.ext4, 2
+  %a21 = getelementptr inbounds float, ptr addrspace(1) %from.coerce, i64 %a20
+  %a22 = load <2 x float>, ptr addrspace(1) %a21, align 16, !tbaa !0
+  %3 = extractelement <2 x float> %a22, i64 1
+  %4 = extractelement <2 x float> %2, i64 0
+  %5 = tail call contract noundef float asm "v_mfma_f64_4x4x4f64 $0, $1, $2, 0", "=a,v,v"(<2 x float> %2, <2 x float> %a22) #3, !srcloc !3
+  %6 = extractelement <2 x float> %2, i64 1
+  %7 = extractelement <2 x float> %a22, i64 0
   %add6 = fadd contract float %6, %7
   %add7 = fadd contract float %5, %add6
   store float %add7, ptr addrspace(1) %add.ptr, align 4, !tbaa !4
   tail call void @llvm.amdgcn.sched.group.barrier(i32 16, i32 1, i32 0)
   tail call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 5, i32 0)
   tail call void @llvm.amdgcn.sched.group.barrier(i32 16, i32 1, i32 0)
   tail call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 1, i32 0)
-  tail call void @llvm.amdgcn.sched.group.barrier(i32 1, i32 1, i32 0)
+  tail call void @llvm.amdgcn.sched.group.barrier(i32 8, i32 1, i32 0)
   tail call void @llvm.amdgcn.sched.group.barrier(i32 2, i32 1, i32 0)
   br label %if.end