[llvm] Reenable sched mfma rewrite (PR #180751)

[Matt Arsenault via llvm-commits]

================
@@ -0,0 +1,309 @@
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx942 -O1 -amdgpu-disable-rewrite-mfma-form-sched-stage=false -verify-machineinstrs < %s | FileCheck %s
+
+define amdgpu_kernel void @"fused_op_layout_cus_534d8e85a33330539c81dbeb009d324be60530f2_1x30x46x2016xbfloat16_1x30x46x2016xbfloat16_2016x1x1x2016xbfloat16$async_dispatch_0_matmul_like_2016x2016x1380_bf16xbf16xf32"(i32 %0, ptr addrspace(3) %1, ptr addrspace(3) %2, ptr addrspace(3) %3, ptr addrspace(3) %4, ptr addrspace(3) %5, ptr addrspace(3) %6, ptr addrspace(3) %7, ptr addrspace(3) %8, ptr addrspace(3) %9, ptr addrspace(3) %10, ptr addrspace(3) %11, ptr addrspace(3) %12, ptr addrspace(7) %invariant.gep, ptr addrspace(7) %gep, ptr addrspace(7) %gep106, ptr addrspace(7) %gep108, ptr addrspace(7) %gep112, ptr addrspace(7) %gep114, ptr addrspace(7) %gep116, ptr addrspace(7) %gep118, ptr addrspace(7) %gep120, ptr addrspace(7) %gep122, ptr addrspace(7) %gep124, ptr addrspace(3) %13, ptr addrspace(3) %14, ptr addrspace(3) %15, ptr addrspace(3) %16, ptr addrspace(3) %17, ptr addrspace(3) %18, ptr addrspace(3) %19, ptr addrspace(3) %20, ptr addrspace(3) %21, ptr addrspace(3) %22, ptr addrspace(3) %23, ptr addrspace(3) %24, ptr addrspace(3) %25, ptr addrspace(3) %26, ptr addrspace(3) %27, ptr addrspace(3) %28, ptr addrspace(3) %29, ptr addrspace(3) %30, ptr addrspace(3) %31, ptr addrspace(3) %32, ptr addrspace(3) %33, ptr addrspace(3) %34, ptr addrspace(3) %35, ptr addrspace(3) %36, ptr addrspace(3) %37, ptr addrspace(3) %38, ptr addrspace(3) %39, ptr addrspace(3) %40, ptr addrspace(3) %41, ptr addrspace(3) %42, ptr addrspace(3) %43, ptr addrspace(3) %44, ptr addrspace(3) %45, ptr addrspace(3) %46, ptr addrspace(3) %47, ptr addrspace(3) %48, ptr addrspace(3) %49, ptr addrspace(3) %50, ptr addrspace(3) %51, ptr addrspace(3) %52, ptr addrspace(3) %53, ptr addrspace(3) %54, ptr addrspace(3) %55, ptr addrspace(3) %56, ptr addrspace(3) %57, ptr addrspace(3) %58, ptr addrspace(3) %59, <4 x bfloat> %60, <4 x i16> %61, <4 x i16> %62, <4 x i16> %63, <4 x i16> %64, <4 x i16> %65, <8 x i1> %66, <4 x i16> %67, <4 x i16> %68, <4 x i16> %69, <4 x i16> %70, <8 x i1> %71) #0 {
+; This should be rejected for MFMA rewrite by the scheduler so we should see
+; no AGPR operands.
+; CHECK-NOT: v_mfma_f32_16x16x16_bf16 a[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}], a[{{[0-9:]+}}]
+  br label %73
+73:                                               ; preds = %73, %72
----------------
arsenm wrote:

Use named values in tests. Also this is pretty huge 

https://github.com/llvm/llvm-project/pull/180751