[llvm] [AMDGPU][NFC] Harden scheduler tests balanced on pressure boundaries (PR #228293)
Dhruva Chakrabarti via llvm-commits
llvm-commits at lists.llvm.org
Thu Oct 1 18:01:19 PDT 2026
https://github.com/dhruvachak created https://github.com/llvm/llvm-project/pull/228293
Several AMDGPU scheduler tests could stop exercising the intended code path for small changes in certain factors, e.g. register pressure, a VGPR block, etc. Move the affected tests away from those boundaries.
- sema-v-unsched-bundle.ll kept just enough values live across the GWS bundle to land on an appropriate number of VGPRs for the schedule to be reverted. Another live vector ensures that unscheduling of bundled instructions is covered with room to spare.
- sched-dynamic-vgpr-blocks.mir needed the region to grow from one 16-VGPR block to two. One more load/add pair ensures the growth.
- coexec-rewrite-mfma.ll needed ArchVGPR pressure above the addressable limit for RewriteMFMAFormStage to run at all. Two loop-carried phis and one more MFMA pair keep it there.
- rewrite-vgpr-mfma-to-agpr-spill-multi-store-mir relied on an accumulator set that only spilled at one particular pressure. Re-reduced the test to make it less dependent on the register pressure.
No functional change; all four tests pass before and after.
Assisted-by: Cursor/Claude Opus
>From d3a8ae8b43d5768f4c3f4b3f9b11e951826d2542 Mon Sep 17 00:00:00 2001
From: Dhruva Chakrabarti <Dhruva.Chakrabarti at amd.com>
Date: Fri, 25 Sep 2026 21:07:52 -0500
Subject: [PATCH] [AMDGPU][NFC] Harden scheduler tests balanced on pressure
boundaries
Several AMDGPU scheduler tests could stop exercising the intended code
path for small changes in certain factors, e.g. register pressure, a
VGPR block, etc. Move the affected tests away from those boundaries.
- sema-v-unsched-bundle.ll kept just enough values live across the GWS
bundle to land on an appropriate number of VGPRs for the schedule to
be reverted. Another live vector ensures that unscheduling of bundled
instructions is covered with room to spare.
- sched-dynamic-vgpr-blocks.mir needed the region to grow from one
16-VGPR block to two. One more load/add pair ensures the growth.
- coexec-rewrite-mfma.ll needed ArchVGPR pressure above the addressable
limit for RewriteMFMAFormStage to run at all. Two loop-carried phis and
one more MFMA pair keep it there.
- rewrite-vgpr-mfma-to-agpr-spill-multi-store-mir relied on an
accumulator set that only spilled at one particular pressure. Re-reduced
the test to make it less dependent on the register pressure.
No functional change; all four tests pass before and after.
Assisted-by: Cursor/Claude Opus
---
.../CodeGen/AMDGPU/coexec-rewrite-mfma.ll | 12 +-
...gpr-mfma-to-agpr-spill-multi-store-mir.mir | 881 +++++++++---------
.../AMDGPU/sched-dynamic-vgpr-blocks.mir | 8 +-
.../CodeGen/AMDGPU/sema-v-unsched-bundle.ll | 10 +-
4 files changed, 482 insertions(+), 429 deletions(-)
diff --git a/llvm/test/CodeGen/AMDGPU/coexec-rewrite-mfma.ll b/llvm/test/CodeGen/AMDGPU/coexec-rewrite-mfma.ll
index a632f5b3d161df0..f73edff901bbf80 100644
--- a/llvm/test/CodeGen/AMDGPU/coexec-rewrite-mfma.ll
+++ b/llvm/test/CodeGen/AMDGPU/coexec-rewrite-mfma.ll
@@ -3,12 +3,12 @@
; The coexec scheduler strategy pushes RewriteMFMAForm into its pipeline so MFMA
; chains can be rewritten to the AGPR form when ArchVGPR pressure is excessive.
;
-; This IR is reduced from a Triton-generated GEMM kernel with multiple (28 to be
+; This IR is reduced from a Triton-generated GEMM kernel with multiple (30 to be
; exact) MFMAs in a phi-carried loop with high VGPR pressure, every MFMA in the
; loop is rewritten to the AGPR-destination form (v_mfma_f32_16x16x32_f16 a*).
; CHECK-LABEL: v5_local_prefetch:
-; CHECK-COUNT-28: v_mfma_f32_16x16x32_f16 a
+; CHECK-COUNT-30: v_mfma_f32_16x16x32_f16 a
; CHECK-NOT: v_mfma_f32_16x16x32_f16 v
define amdgpu_kernel void @v5_local_prefetch() #0 {
@@ -58,6 +58,8 @@ define amdgpu_kernel void @v5_local_prefetch() #0 {
%40 = phi float [ 0.000000e+00, %.lr.ph ], [ %148, %3 ]
%41 = phi float [ 0.000000e+00, %.lr.ph ], [ %153, %3 ]
%42 = phi float [ 0.000000e+00, %.lr.ph ], [ %41, %3 ]
+ %p0 = phi float [ 0.000000e+00, %.lr.ph ], [ %e0, %3 ]
+ %p1 = phi float [ 0.000000e+00, %.lr.ph ], [ %e1, %3 ]
%43 = phi <2 x half> [ %1, %.lr.ph ], [ zeroinitializer, %3 ]
%44 = phi <2 x half> [ %2, %.lr.ph ], [ zeroinitializer, %3 ]
%45 = phi <8 x half> [ zeroinitializer, %.lr.ph ], [ <half 0.000000e+00, half 0.000000e+00, half poison, half poison, half poison, half poison, half poison, half poison>, %3 ]
@@ -169,6 +171,12 @@ define amdgpu_kernel void @v5_local_prefetch() #0 {
%151 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> zeroinitializer, <8 x half> zeroinitializer, <4 x float> %150, i32 0, i32 0, i32 0)
%152 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> zeroinitializer, <8 x half> zeroinitializer, <4 x float> %151, i32 0, i32 0, i32 0)
%153 = extractelement <4 x float> %152, i64 0
+ %x0 = insertelement <4 x float> zeroinitializer, float %p0, i64 0
+ %x1 = insertelement <4 x float> %x0, float %p1, i64 1
+ %x2 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> zeroinitializer, <8 x half> zeroinitializer, <4 x float> %x1, i32 0, i32 0, i32 0)
+ %x3 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> zeroinitializer, <8 x half> zeroinitializer, <4 x float> %x2, i32 0, i32 0, i32 0)
+ %e0 = extractelement <4 x float> %x3, i64 0
+ %e1 = extractelement <4 x float> %x3, i64 1
br label %3
}
diff --git a/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store-mir.mir b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store-mir.mir
index 111f86aff8b53ec..8ce7f38657aea49 100644
--- a/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store-mir.mir
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store-mir.mir
@@ -13,447 +13,482 @@
# CHECK: Split unspilled interval into {{[0-9]+}} components
--- |
- define amdgpu_kernel void @multi_store_spill_slot() #0 {
- entry:
- br label %do.body
-
- do.body:
- br label %DummyReturnBlock
-
- DummyReturnBlock:
- unreachable
- }
+ define amdgpu_kernel void @multi_store_spill_slot() #0 { ret void }
attributes #0 = { "amdgpu-flat-work-group-size"="1,256" "amdgpu-lds-size"="32768" }
...
---
name: multi_store_spill_slot
tracksRegLiveness: true
-noPhis: true
-isSSA: false
machineFunctionInfo:
- ldsSize: 32768
stackPtrOffsetReg: '$sgpr32'
+ occupancy: 2
body: |
- bb.0.entry:
+ bb.0:
successors: %bb.1(0x80000000)
-
- %39:vgpr_32 = AV_MOV_B32_IMM_PSEUDO 0, implicit $exec
- %41:sreg_32 = S_MOV_B32 0
- undef %42.sub0:sreg_64 = COPY %41
- %42.sub1:sreg_64 = COPY killed %41
- %44:av_64_align2 = COPY killed %42
- %45:vgpr_32 = V_MOV_B32_e32 2143289344, implicit $exec
- %50:vgpr_32 = AV_MOV_B32_IMM_PSEUDO 1065353216, implicit $exec
- %101:vgpr_32 = V_MOV_B32_e32 1006648320, implicit $exec
- undef %93.sub0:av_64_align2 = COPY %101
- %93.sub1:av_64_align2 = COPY killed %101
- %100:sreg_64 = S_AND_B64 $exec, -1, implicit-def dead $scc
- %102:av_32 = COPY %39
- %103:av_32 = COPY %39
- %104:av_32 = COPY %39
- %105:av_32 = COPY %39
- %106:av_32 = COPY %39
- %107:av_32 = COPY %39
- %108:av_32 = COPY %39
- %109:av_32 = COPY %39
- %110:av_32 = COPY %39
- %111:av_32 = COPY %39
- %112:av_32 = COPY %39
- %113:av_32 = COPY %39
- %114:av_32 = COPY %39
- %115:av_32 = COPY %39
- %116:av_32 = COPY %39
- %117:av_32 = COPY %39
- %118:av_32 = COPY %39
- %119:av_32 = COPY %39
- %120:av_32 = COPY %39
- %121:av_32 = COPY %39
- %122:av_32 = COPY %39
-
- bb.1.do.body:
+
+ %0:vgpr_32 = AV_MOV_B32_IMM_PSEUDO 0, implicit $exec
+ %1:vgpr_32 = AV_MOV_B32_IMM_PSEUDO 0, implicit $exec
+ %2:vgpr_32 = AV_MOV_B32_IMM_PSEUDO 0, implicit $exec
+ %3:vgpr_32 = AV_MOV_B32_IMM_PSEUDO 0, implicit $exec
+ %4:vgpr_32 = AV_MOV_B32_IMM_PSEUDO 0, implicit $exec
+ %5:vgpr_32 = AV_MOV_B32_IMM_PSEUDO 0, implicit $exec
+ %6:vgpr_32 = AV_MOV_B32_IMM_PSEUDO 0, implicit $exec
+ %7:vgpr_32 = AV_MOV_B32_IMM_PSEUDO 0, implicit $exec
+ %8:vgpr_32 = AV_MOV_B32_IMM_PSEUDO 0, implicit $exec
+ %9:vgpr_32 = AV_MOV_B32_IMM_PSEUDO 0, implicit $exec
+ %10:vgpr_32 = AV_MOV_B32_IMM_PSEUDO 0, implicit $exec
+ %11:vgpr_32 = AV_MOV_B32_IMM_PSEUDO 0, implicit $exec
+ %12:vgpr_32 = AV_MOV_B32_IMM_PSEUDO 0, implicit $exec
+ %13:vgpr_32 = AV_MOV_B32_IMM_PSEUDO 0, implicit $exec
+ %14:vgpr_32 = AV_MOV_B32_IMM_PSEUDO 0, implicit $exec
+ %15:vgpr_32 = AV_MOV_B32_IMM_PSEUDO 0, implicit $exec
+ %16:vgpr_32 = AV_MOV_B32_IMM_PSEUDO 0, implicit $exec
+ %17:vgpr_32 = AV_MOV_B32_IMM_PSEUDO 0, implicit $exec
+ %18:vgpr_32 = AV_MOV_B32_IMM_PSEUDO 0, implicit $exec
+ %19:vgpr_32 = AV_MOV_B32_IMM_PSEUDO 0, implicit $exec
+ %20:vgpr_32 = AV_MOV_B32_IMM_PSEUDO 0, implicit $exec
+ %21:vgpr_32 = AV_MOV_B32_IMM_PSEUDO 1065353216, implicit $exec
+ %22:sreg_32 = S_MOV_B32 0
+ undef %23.sub0:sreg_64 = COPY %22
+ %23.sub1:sreg_64 = COPY killed %22
+ %24:av_64_align2 = COPY killed %23
+ %25:vgpr_32 = V_MOV_B32_e32 2143289344, implicit $exec
+ %26:sreg_64 = S_AND_B64 $exec, -1, implicit-def dead $scc
+ %27:av_32 = COPY killed %1
+ %28:av_32 = COPY killed %2
+ %29:av_32 = COPY killed %3
+ %30:av_32 = COPY killed %4
+ %31:av_32 = COPY killed %5
+ %32:av_32 = COPY killed %6
+ %33:av_32 = COPY killed %7
+ %34:av_32 = COPY killed %8
+ %35:av_32 = COPY killed %9
+ %36:vgpr_32 = COPY %0
+ %37:av_32 = COPY killed %10
+ %38:av_32 = COPY killed %11
+ %39:av_32 = COPY killed %12
+ %40:av_32 = COPY killed %13
+ %41:av_32 = COPY killed %14
+ %42:av_32 = COPY killed %15
+ %43:av_32 = COPY killed %16
+ %44:av_32 = COPY killed %17
+ %45:av_32 = COPY killed %18
+ %46:av_32 = COPY killed %19
+ %47:av_32 = COPY killed %20
+
+ bb.1:
successors: %bb.2(0x04000000), %bb.1(0x7c000000)
-
- %19:av_32 = COPY killed %121
- %18:av_32 = COPY killed %120
- %17:av_32 = COPY killed %119
- %16:av_32 = COPY killed %118
- %15:av_32 = COPY killed %117
- %14:av_32 = COPY killed %116
- %13:av_32 = COPY killed %115
- %12:av_32 = COPY killed %114
- %11:av_32 = COPY killed %113
- %10:av_32 = COPY killed %112
- %9:av_32 = COPY killed %111
- %8:av_32 = COPY killed %110
- %7:av_32 = COPY killed %109
- %6:av_32 = COPY killed %108
- %5:av_32 = COPY killed %107
- %4:av_32 = COPY killed %106
- %3:av_32 = COPY killed %105
- %2:av_32 = COPY killed %104
- %1:av_32 = COPY killed %103
- %0:av_32 = COPY killed %102
- undef %40.sub0:vreg_512_align2 = COPY killed %0
- %40.sub1:vreg_512_align2 = COPY %39
- %40.sub2:vreg_512_align2 = COPY %39
- %40.sub3:vreg_512_align2 = COPY %39
- %40.sub4:vreg_512_align2 = COPY %39
- %40.sub5:vreg_512_align2 = COPY %39
- %40.sub6:vreg_512_align2 = COPY %39
- %40.sub7:vreg_512_align2 = COPY %39
- %40.sub8:vreg_512_align2 = COPY %39
- %40.sub9:vreg_512_align2 = COPY %39
- %40.sub10:vreg_512_align2 = COPY %39
- %40.sub11:vreg_512_align2 = COPY %39
- %40.sub12:vreg_512_align2 = COPY %39
- %40.sub13:vreg_512_align2 = COPY %39
- %40.sub14:vreg_512_align2 = COPY %39
- %40.sub15:vreg_512_align2 = COPY %39
- %43:vreg_512_align2 = COPY killed %40
- %43:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %43, 0, 0, 0, implicit $mode, implicit $exec
- undef %46.sub0:vreg_512_align2 = COPY %45
- %46.sub1:vreg_512_align2 = COPY %39
- %46.sub2:vreg_512_align2 = COPY %39
- %46.sub3:vreg_512_align2 = COPY %39
- %46.sub4:vreg_512_align2 = COPY %39
- %46.sub5:vreg_512_align2 = COPY %39
- %46.sub6:vreg_512_align2 = COPY %39
- %46.sub7:vreg_512_align2 = COPY %39
- %46.sub8:vreg_512_align2 = COPY %39
- %46.sub9:vreg_512_align2 = COPY %39
- %46.sub10:vreg_512_align2 = COPY %39
- %46.sub11:vreg_512_align2 = COPY %39
- %46.sub12:vreg_512_align2 = COPY %39
- %46.sub13:vreg_512_align2 = COPY killed %1
- %46.sub14:vreg_512_align2 = COPY killed %2
- %46.sub15:vreg_512_align2 = COPY %39
- %47:vreg_512_align2 = COPY killed %46
- %47:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %47, 0, 0, 0, implicit $mode, implicit $exec
- undef %48.sub0:vreg_512_align2 = COPY %39
- %48.sub1:vreg_512_align2 = COPY %39
- %48.sub2:vreg_512_align2 = COPY %39
- %48.sub3:vreg_512_align2 = COPY killed %5
- %48.sub4:vreg_512_align2 = COPY %39
- %48.sub5:vreg_512_align2 = COPY %39
- %48.sub6:vreg_512_align2 = COPY %39
- %48.sub7:vreg_512_align2 = COPY %39
- %48.sub8:vreg_512_align2 = COPY %39
- %48.sub9:vreg_512_align2 = COPY %39
- %48.sub10:vreg_512_align2 = COPY %39
- %48.sub11:vreg_512_align2 = COPY %39
- %48.sub12:vreg_512_align2 = COPY %39
- %48.sub13:vreg_512_align2 = COPY %39
- %48.sub14:vreg_512_align2 = COPY %39
- %48.sub15:vreg_512_align2 = COPY %39
- undef %49.sub0:vreg_512_align2 = COPY killed %6
- %49.sub1:vreg_512_align2 = COPY %39
- %49.sub2:vreg_512_align2 = COPY %39
- %49.sub3:vreg_512_align2 = COPY %39
- %49.sub4:vreg_512_align2 = COPY %39
- %49.sub5:vreg_512_align2 = COPY %39
- %49.sub6:vreg_512_align2 = COPY %39
- %49.sub7:vreg_512_align2 = COPY %39
- %49.sub8:vreg_512_align2 = COPY %39
- %49.sub9:vreg_512_align2 = COPY %39
- %49.sub10:vreg_512_align2 = COPY %39
- %49.sub11:vreg_512_align2 = COPY %39
- %49.sub12:vreg_512_align2 = COPY %39
- %49.sub13:vreg_512_align2 = COPY %39
- %49.sub14:vreg_512_align2 = COPY %39
- %49.sub15:vreg_512_align2 = COPY %39
- undef %51.sub0:vreg_512_align2 = COPY killed %8
- %51.sub1:vreg_512_align2 = COPY %50
- %51.sub2:vreg_512_align2 = COPY %50
- %51.sub3:vreg_512_align2 = COPY %50
- %51.sub4:vreg_512_align2 = COPY %50
- %51.sub5:vreg_512_align2 = COPY %50
- %51.sub6:vreg_512_align2 = COPY %50
- %51.sub7:vreg_512_align2 = COPY %50
- %51.sub8:vreg_512_align2 = COPY %50
- %51.sub9:vreg_512_align2 = COPY %50
- %51.sub10:vreg_512_align2 = COPY %50
- %51.sub11:vreg_512_align2 = COPY %50
- %51.sub12:vreg_512_align2 = COPY %50
- %51.sub13:vreg_512_align2 = COPY %50
- %51.sub14:vreg_512_align2 = COPY %50
- %51.sub15:vreg_512_align2 = COPY %50
- %52:vreg_512_align2 = COPY killed %51
- %52:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %52, 0, 0, 0, implicit $mode, implicit $exec
- undef %53.sub0:vreg_512_align2 = COPY killed %10
- %53.sub1:vreg_512_align2 = COPY killed %9
- %53.sub2:vreg_512_align2 = COPY %39
- %53.sub3:vreg_512_align2 = COPY %39
- %53.sub4:vreg_512_align2 = COPY %39
- %53.sub5:vreg_512_align2 = COPY %39
- %53.sub6:vreg_512_align2 = COPY %39
- %53.sub7:vreg_512_align2 = COPY %39
- %53.sub8:vreg_512_align2 = COPY %39
- %53.sub9:vreg_512_align2 = COPY %39
- %53.sub10:vreg_512_align2 = COPY %39
- %53.sub11:vreg_512_align2 = COPY %39
- %53.sub12:vreg_512_align2 = COPY %39
- %53.sub13:vreg_512_align2 = COPY %39
- %53.sub14:vreg_512_align2 = COPY %39
- %53.sub15:vreg_512_align2 = COPY %39
- undef %54.sub0:vreg_512_align2 = COPY killed %12
- %54.sub1:vreg_512_align2 = COPY %39
- %54.sub2:vreg_512_align2 = COPY %39
- %54.sub3:vreg_512_align2 = COPY %39
- %54.sub4:vreg_512_align2 = COPY %39
- %54.sub5:vreg_512_align2 = COPY %39
- %54.sub6:vreg_512_align2 = COPY %39
- %54.sub7:vreg_512_align2 = COPY %39
- %54.sub8:vreg_512_align2 = COPY %39
- %54.sub9:vreg_512_align2 = COPY %39
- %54.sub10:vreg_512_align2 = COPY %39
- %54.sub11:vreg_512_align2 = COPY %39
- %54.sub12:vreg_512_align2 = COPY %39
- %54.sub13:vreg_512_align2 = COPY %39
- %54.sub14:vreg_512_align2 = COPY %39
- %54.sub15:vreg_512_align2 = COPY %39
- undef %55.sub0:vreg_512_align2 = COPY %39
- %55.sub1:vreg_512_align2 = COPY %39
- %55.sub2:vreg_512_align2 = COPY %39
- %55.sub3:vreg_512_align2 = COPY %39
- %55.sub4:vreg_512_align2 = COPY %39
- %55.sub5:vreg_512_align2 = COPY %39
- %55.sub6:vreg_512_align2 = COPY %39
- %55.sub7:vreg_512_align2 = COPY %39
- %55.sub8:vreg_512_align2 = COPY %39
- %55.sub9:vreg_512_align2 = COPY %39
- %55.sub10:vreg_512_align2 = COPY %39
- %55.sub11:vreg_512_align2 = COPY %39
- %55.sub12:vreg_512_align2 = COPY %39
- %55.sub13:vreg_512_align2 = COPY %39
- %55.sub14:vreg_512_align2 = COPY %13
- %55.sub15:vreg_512_align2 = COPY killed %14
- %56:vreg_512_align2 = COPY killed %55
- %56:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %56, 0, 0, 0, implicit $mode, implicit $exec
- undef %57.sub0:vreg_512_align2 = COPY killed %15
- %57.sub1:vreg_512_align2 = COPY %39
- %57.sub2:vreg_512_align2 = COPY %39
- %57.sub3:vreg_512_align2 = COPY %39
- %57.sub4:vreg_512_align2 = COPY %39
- %57.sub5:vreg_512_align2 = COPY %39
- %57.sub6:vreg_512_align2 = COPY %39
- %57.sub7:vreg_512_align2 = COPY %39
- %57.sub8:vreg_512_align2 = COPY %39
- %57.sub9:vreg_512_align2 = COPY %39
- %57.sub10:vreg_512_align2 = COPY %39
- %57.sub11:vreg_512_align2 = COPY %39
- %57.sub12:vreg_512_align2 = COPY %39
- %57.sub13:vreg_512_align2 = COPY %39
- %57.sub14:vreg_512_align2 = COPY %39
- %57.sub15:vreg_512_align2 = COPY %39
- undef %58.sub0:vreg_512_align2 = COPY %17
- %58.sub1:vreg_512_align2 = COPY killed %16
- %58.sub2:vreg_512_align2 = COPY %39
- %58.sub3:vreg_512_align2 = COPY %39
- %58.sub4:vreg_512_align2 = COPY %39
- %58.sub5:vreg_512_align2 = COPY %39
- %58.sub6:vreg_512_align2 = COPY %39
- %58.sub7:vreg_512_align2 = COPY %39
- %58.sub8:vreg_512_align2 = COPY %39
- %58.sub9:vreg_512_align2 = COPY %39
- %58.sub10:vreg_512_align2 = COPY %39
- %58.sub11:vreg_512_align2 = COPY %39
- %58.sub12:vreg_512_align2 = COPY %39
- %58.sub13:vreg_512_align2 = COPY %39
- %58.sub14:vreg_512_align2 = COPY %39
- %58.sub15:vreg_512_align2 = COPY %39
- undef %59.sub0:vreg_512_align2 = COPY %39
- %59.sub1:vreg_512_align2 = COPY %39
- %59.sub2:vreg_512_align2 = COPY %39
- %59.sub3:vreg_512_align2 = COPY %39
- %59.sub4:vreg_512_align2 = COPY %39
- %59.sub5:vreg_512_align2 = COPY %39
- %59.sub6:vreg_512_align2 = COPY %39
- %59.sub7:vreg_512_align2 = COPY %39
- %59.sub8:vreg_512_align2 = COPY %39
- %59.sub9:vreg_512_align2 = COPY %39
- %59.sub10:vreg_512_align2 = COPY %39
- %59.sub11:vreg_512_align2 = COPY %39
- %59.sub12:vreg_512_align2 = COPY %39
- %59.sub13:vreg_512_align2 = COPY %39
- %59.sub14:vreg_512_align2 = COPY %39
- %59.sub15:vreg_512_align2 = COPY killed %18
- %20:av_32 = COPY killed %122
- undef %60.sub0:vreg_512_align2 = COPY killed %19
- %60.sub1:vreg_512_align2 = COPY killed %20
- %60.sub2:vreg_512_align2 = COPY %39
- %60.sub3:vreg_512_align2 = COPY %39
- %60.sub4:vreg_512_align2 = COPY %39
- %60.sub5:vreg_512_align2 = COPY %39
- %60.sub6:vreg_512_align2 = COPY %39
- %60.sub7:vreg_512_align2 = COPY %39
- %60.sub8:vreg_512_align2 = COPY %39
- %60.sub9:vreg_512_align2 = COPY %39
- %60.sub10:vreg_512_align2 = COPY %39
- %60.sub11:vreg_512_align2 = COPY %39
- %60.sub12:vreg_512_align2 = COPY %39
- %60.sub13:vreg_512_align2 = COPY %39
- %60.sub14:vreg_512_align2 = COPY %39
- %60.sub15:vreg_512_align2 = COPY %39
- %61:vreg_512_align2 = COPY killed %54
- %61:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %61, 0, 0, 0, implicit $mode, implicit $exec
- %62:vreg_512_align2 = COPY killed %61
- %62:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %62, 0, 0, 0, implicit $mode, implicit $exec
- %63:vreg_512_align2 = COPY killed %59
- %63:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %63, 0, 0, 0, implicit $mode, implicit $exec
- %64:vreg_512_align2 = COPY killed %63
- %64:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %64, 0, 0, 0, implicit $mode, implicit $exec
- %65:vreg_512_align2 = COPY killed %48
- %65:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %65, 0, 0, 0, implicit $mode, implicit $exec
- %66:vreg_512_align2 = COPY killed %43
- %66:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %66, 0, 0, 0, implicit $mode, implicit $exec
- %67:vreg_512_align2 = COPY killed %52
- %67:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %67, 0, 0, 0, implicit $mode, implicit $exec
- %68:vreg_512_align2 = COPY killed %53
- %68:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %68, 0, 0, 0, implicit $mode, implicit $exec
- %69:vreg_512_align2 = COPY killed %58
- %69:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %69, 0, 0, 0, implicit $mode, implicit $exec
- %70:vreg_512_align2 = COPY killed %60
- %70:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %70, 0, 0, 0, implicit $mode, implicit $exec
- %71:vreg_512_align2 = COPY killed %66
- %71:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %71, 0, 0, 0, implicit $mode, implicit $exec
- early-clobber %72:vreg_512_align2 = V_MFMA_F32_32X32X8F16_vgprcd_e64 %44, %44, %68, 0, 0, 0, implicit $mode, implicit $exec
- %73:vreg_512_align2 = COPY killed %56
- %73:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %73, 0, 0, 0, implicit $mode, implicit $exec
- %74:vreg_512_align2 = COPY killed %47
- %74:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %74, 0, 0, 0, implicit $mode, implicit $exec
- early-clobber %75:vreg_512_align2 = V_MFMA_F32_32X32X8F16_vgprcd_e64 %44, %44, %74, 0, 0, 0, implicit $mode, implicit $exec
- undef %76.sub0:vreg_512_align2 = COPY %3
- %76.sub1:vreg_512_align2 = COPY killed %4
- %76.sub2:vreg_512_align2 = COPY %39
- %76.sub3:vreg_512_align2 = COPY %39
- %76.sub4:vreg_512_align2 = COPY %39
- %76.sub5:vreg_512_align2 = COPY %39
- %76.sub6:vreg_512_align2 = COPY %39
- %76.sub7:vreg_512_align2 = COPY %39
- %76.sub8:vreg_512_align2 = COPY %39
- %76.sub9:vreg_512_align2 = COPY %39
- %76.sub10:vreg_512_align2 = COPY %39
- %76.sub11:vreg_512_align2 = COPY %39
- %76.sub12:vreg_512_align2 = COPY %39
- %76.sub13:vreg_512_align2 = COPY %39
- %76.sub14:vreg_512_align2 = COPY %39
- %76.sub15:vreg_512_align2 = COPY %39
+
+ %48:av_32 = COPY killed %46
+ %49:av_32 = COPY killed %45
+ %50:av_32 = COPY killed %44
+ %51:av_32 = COPY killed %43
+ %52:av_32 = COPY killed %42
+ %53:av_32 = COPY killed %41
+ %54:av_32 = COPY killed %40
+ %55:av_32 = COPY killed %39
+ %56:av_32 = COPY killed %38
+ %57:av_32 = COPY killed %37
+ %58:vgpr_32 = COPY killed %36
+ %59:av_32 = COPY killed %35
+ %60:av_32 = COPY killed %34
+ %61:av_32 = COPY killed %33
+ %62:av_32 = COPY killed %32
+ %63:av_32 = COPY killed %31
+ %64:av_32 = COPY killed %30
+ %65:av_32 = COPY killed %29
+ %66:av_32 = COPY killed %28
+ %67:av_32 = COPY killed %27
+ undef %68.sub0:vreg_512_align2 = COPY killed %67
+ %68.sub1:vreg_512_align2 = COPY %0
+ %68.sub2:vreg_512_align2 = COPY %0
+ %68.sub3:vreg_512_align2 = COPY %0
+ %68.sub4:vreg_512_align2 = COPY %0
+ %68.sub5:vreg_512_align2 = COPY %0
+ %68.sub6:vreg_512_align2 = COPY %0
+ %68.sub7:vreg_512_align2 = COPY %0
+ %68.sub8:vreg_512_align2 = COPY %0
+ %68.sub9:vreg_512_align2 = COPY %0
+ %68.sub10:vreg_512_align2 = COPY %0
+ %68.sub11:vreg_512_align2 = COPY %0
+ %68.sub12:vreg_512_align2 = COPY %0
+ %68.sub13:vreg_512_align2 = COPY %0
+ %68.sub14:vreg_512_align2 = COPY %0
+ %68.sub15:vreg_512_align2 = COPY %0
+ undef %69.sub0:vreg_512_align2 = COPY killed %66
+ %69.sub1:vreg_512_align2 = COPY %0
+ %69.sub2:vreg_512_align2 = COPY %0
+ %69.sub3:vreg_512_align2 = COPY %0
+ %69.sub4:vreg_512_align2 = COPY %0
+ %69.sub5:vreg_512_align2 = COPY %0
+ %69.sub6:vreg_512_align2 = COPY %0
+ %69.sub7:vreg_512_align2 = COPY %0
+ %69.sub8:vreg_512_align2 = COPY %0
+ %69.sub9:vreg_512_align2 = COPY %0
+ %69.sub10:vreg_512_align2 = COPY %0
+ %69.sub11:vreg_512_align2 = COPY %0
+ %69.sub12:vreg_512_align2 = COPY %0
+ %69.sub13:vreg_512_align2 = COPY %0
+ %69.sub14:vreg_512_align2 = COPY %0
+ %69.sub15:vreg_512_align2 = COPY %0
+ undef %70.sub0:vreg_512_align2 = COPY killed %63
+ %70.sub1:vreg_512_align2 = COPY %21
+ %70.sub2:vreg_512_align2 = COPY %21
+ %70.sub3:vreg_512_align2 = COPY %21
+ %70.sub4:vreg_512_align2 = COPY %21
+ %70.sub5:vreg_512_align2 = COPY %21
+ %70.sub6:vreg_512_align2 = COPY %21
+ %70.sub7:vreg_512_align2 = COPY %21
+ %70.sub8:vreg_512_align2 = COPY %21
+ %70.sub9:vreg_512_align2 = COPY %21
+ %70.sub10:vreg_512_align2 = COPY %21
+ %70.sub11:vreg_512_align2 = COPY %21
+ %70.sub12:vreg_512_align2 = COPY %21
+ %70.sub13:vreg_512_align2 = COPY %21
+ %70.sub14:vreg_512_align2 = COPY %21
+ %70.sub15:vreg_512_align2 = COPY %21
+ %71:vreg_512_align2 = COPY killed %70
+ %71:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %71, 0, 0, 0, implicit $mode, implicit $exec
+ undef %72.sub0:vreg_512_align2 = COPY %0
+ %72.sub1:vreg_512_align2 = COPY killed %55
+ %72.sub2:vreg_512_align2 = COPY %0
+ %72.sub3:vreg_512_align2 = COPY %0
+ %72.sub4:vreg_512_align2 = COPY %0
+ %72.sub5:vreg_512_align2 = COPY %0
+ %72.sub6:vreg_512_align2 = COPY %0
+ %72.sub7:vreg_512_align2 = COPY %0
+ %72.sub8:vreg_512_align2 = COPY %0
+ %72.sub9:vreg_512_align2 = COPY %0
+ %72.sub10:vreg_512_align2 = COPY %0
+ %72.sub11:vreg_512_align2 = COPY %0
+ %72.sub12:vreg_512_align2 = COPY %0
+ %72.sub13:vreg_512_align2 = COPY %0
+ %72.sub14:vreg_512_align2 = COPY %0
+ %72.sub15:vreg_512_align2 = COPY %0
+ undef %73.sub0:vreg_512_align2 = COPY killed %50
+ %73.sub1:vreg_512_align2 = COPY killed %51
+ %73.sub2:vreg_512_align2 = COPY %0
+ %73.sub3:vreg_512_align2 = COPY %0
+ %73.sub4:vreg_512_align2 = COPY %0
+ %73.sub5:vreg_512_align2 = COPY %0
+ %73.sub6:vreg_512_align2 = COPY %0
+ %73.sub7:vreg_512_align2 = COPY %0
+ %73.sub8:vreg_512_align2 = COPY %0
+ %73.sub9:vreg_512_align2 = COPY %0
+ %73.sub10:vreg_512_align2 = COPY %0
+ %73.sub11:vreg_512_align2 = COPY %0
+ %73.sub12:vreg_512_align2 = COPY %0
+ %73.sub13:vreg_512_align2 = COPY %0
+ %73.sub14:vreg_512_align2 = COPY %0
+ %73.sub15:vreg_512_align2 = COPY %0
+ undef %74.sub0:vreg_512_align2 = COPY %56
+ %74.sub1:vreg_512_align2 = COPY killed %57
+ %74.sub2:vreg_512_align2 = COPY %0
+ %74.sub3:vreg_512_align2 = COPY %0
+ %74.sub4:vreg_512_align2 = COPY %0
+ %74.sub5:vreg_512_align2 = COPY %0
+ %74.sub6:vreg_512_align2 = COPY %0
+ %74.sub7:vreg_512_align2 = COPY %0
+ %74.sub8:vreg_512_align2 = COPY %0
+ %74.sub9:vreg_512_align2 = COPY %0
+ %74.sub10:vreg_512_align2 = COPY %0
+ %74.sub11:vreg_512_align2 = COPY %0
+ %74.sub12:vreg_512_align2 = COPY %0
+ %74.sub13:vreg_512_align2 = COPY %0
+ %74.sub14:vreg_512_align2 = COPY %0
+ %74.sub15:vreg_512_align2 = COPY %0
+ undef %75.sub0:vreg_512_align2 = COPY killed %58
+ %75.sub1:vreg_512_align2 = COPY killed %59
+ %75.sub2:vreg_512_align2 = COPY %61
+ %75.sub3:vreg_512_align2 = COPY %0
+ %75.sub4:vreg_512_align2 = COPY %0
+ %75.sub5:vreg_512_align2 = COPY %0
+ %75.sub6:vreg_512_align2 = COPY %0
+ %75.sub7:vreg_512_align2 = COPY %0
+ %75.sub8:vreg_512_align2 = COPY %0
+ %75.sub9:vreg_512_align2 = COPY %0
+ %75.sub10:vreg_512_align2 = COPY %0
+ %75.sub11:vreg_512_align2 = COPY %0
+ %75.sub12:vreg_512_align2 = COPY %0
+ %75.sub13:vreg_512_align2 = COPY %0
+ %75.sub14:vreg_512_align2 = COPY %0
+ %75.sub15:vreg_512_align2 = COPY %0
+ undef %76.sub0:vreg_512_align2 = COPY killed %54
+ %76.sub1:vreg_512_align2 = COPY %25
+ %76.sub2:vreg_512_align2 = COPY %25
+ %76.sub3:vreg_512_align2 = COPY %25
+ %76.sub4:vreg_512_align2 = COPY %25
+ %76.sub5:vreg_512_align2 = COPY %25
+ %76.sub6:vreg_512_align2 = COPY %25
+ %76.sub7:vreg_512_align2 = COPY %25
+ %76.sub8:vreg_512_align2 = COPY %25
+ %76.sub9:vreg_512_align2 = COPY %25
+ %76.sub10:vreg_512_align2 = COPY %25
+ %76.sub11:vreg_512_align2 = COPY %25
+ %76.sub12:vreg_512_align2 = COPY %25
+ %76.sub13:vreg_512_align2 = COPY %25
+ %76.sub14:vreg_512_align2 = COPY %25
+ %76.sub15:vreg_512_align2 = COPY %25
%77:vreg_512_align2 = COPY killed %76
- %77:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %77, 0, 0, 0, implicit $mode, implicit $exec
- %78:vreg_512_align2 = COPY killed %65
- %78:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %78, 0, 0, 0, implicit $mode, implicit $exec
+ %77:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %77, 0, 0, 0, implicit $mode, implicit $exec
+ undef %78.sub0:vreg_512_align2 = COPY killed %53
+ %78.sub1:vreg_512_align2 = COPY %21
+ %78.sub2:vreg_512_align2 = COPY %21
+ %78.sub3:vreg_512_align2 = COPY %21
+ %78.sub4:vreg_512_align2 = COPY %21
+ %78.sub5:vreg_512_align2 = COPY %21
+ %78.sub6:vreg_512_align2 = COPY %21
+ %78.sub7:vreg_512_align2 = COPY %21
+ %78.sub8:vreg_512_align2 = COPY %21
+ %78.sub9:vreg_512_align2 = COPY %21
+ %78.sub10:vreg_512_align2 = COPY %21
+ %78.sub11:vreg_512_align2 = COPY %21
+ %78.sub12:vreg_512_align2 = COPY %21
+ %78.sub13:vreg_512_align2 = COPY %21
+ %78.sub14:vreg_512_align2 = COPY %21
+ %78.sub15:vreg_512_align2 = COPY %21
%79:vreg_512_align2 = COPY killed %78
- %79:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %79, 0, 0, 0, implicit $mode, implicit $exec
- %80:vreg_512_align2 = COPY killed %73
- %80:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %80, 0, 0, 0, implicit $mode, implicit $exec
- %81:vreg_512_align2 = COPY killed %69
- %81:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %81, 0, 0, 0, implicit $mode, implicit $exec
- %21:av_32 = COPY killed %71.sub0
- %82:vreg_512_align2 = COPY killed %75
- %82:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %82, 0, 0, 0, implicit $mode, implicit $exec
- %22:av_32 = COPY killed %82.sub0
- %23:av_32 = COPY killed %74.sub0
- %83:vreg_512_align2 = COPY killed %77
- %83:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %83, 0, 0, 0, implicit $mode, implicit $exec
- %24:av_32 = COPY killed %83.sub0
- %25:av_32 = COPY killed %79.sub0
- %84:vreg_512_align2 = COPY killed %49
- %84:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %84, 0, 0, 0, implicit $mode, implicit $exec
- %85:vreg_512_align2 = COPY killed %84
- %85:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %85, 0, 0, 0, implicit $mode, implicit $exec
- %26:av_32 = COPY killed %85.sub0
- undef %86.sub0:vreg_512_align2 = COPY killed %7
- %86.sub1:vreg_512_align2 = COPY %39
- %86.sub2:vreg_512_align2 = COPY %39
- %86.sub3:vreg_512_align2 = COPY %39
- %86.sub4:vreg_512_align2 = COPY %39
- %86.sub5:vreg_512_align2 = COPY %39
- %86.sub6:vreg_512_align2 = COPY %39
- %86.sub7:vreg_512_align2 = COPY %39
- %86.sub8:vreg_512_align2 = COPY %39
- %86.sub9:vreg_512_align2 = COPY %39
- %86.sub10:vreg_512_align2 = COPY %39
- %86.sub11:vreg_512_align2 = COPY %39
- %86.sub12:vreg_512_align2 = COPY %39
- %86.sub13:vreg_512_align2 = COPY %39
- %86.sub14:vreg_512_align2 = COPY %39
- %86.sub15:vreg_512_align2 = COPY %39
+ %79:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %79, 0, 0, 0, implicit $mode, implicit $exec
+ undef %80.sub0:vreg_512_align2 = COPY killed %65
+ %80.sub1:vreg_512_align2 = COPY %0
+ %80.sub2:vreg_512_align2 = COPY %0
+ %80.sub3:vreg_512_align2 = COPY %0
+ %80.sub4:vreg_512_align2 = COPY %0
+ %80.sub5:vreg_512_align2 = COPY %0
+ %80.sub6:vreg_512_align2 = COPY %0
+ %80.sub7:vreg_512_align2 = COPY %0
+ %80.sub8:vreg_512_align2 = COPY %0
+ %80.sub9:vreg_512_align2 = COPY %0
+ %80.sub10:vreg_512_align2 = COPY %0
+ %80.sub11:vreg_512_align2 = COPY %0
+ %80.sub12:vreg_512_align2 = COPY %0
+ %80.sub13:vreg_512_align2 = COPY %0
+ %80.sub14:vreg_512_align2 = COPY %0
+ %80.sub15:vreg_512_align2 = COPY %0
+ undef %81.sub0:vreg_512_align2 = COPY killed %64
+ %81.sub1:vreg_512_align2 = COPY %0
+ %81.sub2:vreg_512_align2 = COPY %0
+ %81.sub3:vreg_512_align2 = COPY %0
+ %81.sub4:vreg_512_align2 = COPY %0
+ %81.sub5:vreg_512_align2 = COPY %0
+ %81.sub6:vreg_512_align2 = COPY %0
+ %81.sub7:vreg_512_align2 = COPY %0
+ %81.sub8:vreg_512_align2 = COPY %0
+ %81.sub9:vreg_512_align2 = COPY %0
+ %81.sub10:vreg_512_align2 = COPY %0
+ %81.sub11:vreg_512_align2 = COPY %0
+ %81.sub12:vreg_512_align2 = COPY %0
+ %81.sub13:vreg_512_align2 = COPY %0
+ %81.sub14:vreg_512_align2 = COPY %0
+ %81.sub15:vreg_512_align2 = COPY %0
+ %82:vreg_512_align2 = COPY killed %81
+ %82:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %82, 0, 0, 0, implicit $mode, implicit $exec
+ %83:vreg_512_align2 = COPY killed %82
+ %83:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %83, 0, 0, 0, implicit $mode, implicit $exec
+ %84:vreg_512_align2 = COPY killed %74
+ %84:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %84, 0, 0, 0, implicit $mode, implicit $exec
+ %85:vreg_512_align2 = COPY killed %68
+ %85:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %85, 0, 0, 0, implicit $mode, implicit $exec
+ %86:vreg_512_align2 = COPY killed %69
+ %86:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %86, 0, 0, 0, implicit $mode, implicit $exec
%87:vreg_512_align2 = COPY killed %86
- %87:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %87, 0, 0, 0, implicit $mode, implicit $exec
+ %87:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %87, 0, 0, 0, implicit $mode, implicit $exec
%88:vreg_512_align2 = COPY killed %87
- %88:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %88, 0, 0, 0, implicit $mode, implicit $exec
- %27:av_32 = COPY killed %88.sub0
- %89:vreg_512_align2 = COPY killed %67
- %89:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %89, 0, 0, 0, implicit $mode, implicit $exec
- %28:av_32 = COPY killed %89.sub0
- %29:av_32 = COPY killed %68.sub0
- %90:vreg_512_align2 = COPY killed %72
- %90:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %90, 0, 0, 0, implicit $mode, implicit $exec
- %91:vreg_512_align2 = COPY killed %90
- %91:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %91, 0, 0, 0, implicit $mode, implicit $exec
+ %88:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %88, 0, 0, 0, implicit $mode, implicit $exec
+ %89:vreg_512_align2 = COPY killed %88
+ %89:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %89, 0, 0, 0, implicit $mode, implicit $exec
+ %90:vreg_512_align2 = COPY killed %89
+ %90:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %90, 0, 0, 0, implicit $mode, implicit $exec
+ %91:vreg_512_align2 = COPY killed %75
+ %91:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %91, 0, 0, 0, implicit $mode, implicit $exec
%92:vreg_512_align2 = COPY killed %91
- %92:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %93, %92, 0, 0, 0, implicit $mode, implicit $exec
- %30:av_32 = COPY killed %92.sub7
- undef %94.sub0:vreg_512_align2 = COPY killed %11
- %94.sub1:vreg_512_align2 = COPY %39
- %94.sub2:vreg_512_align2 = COPY %39
- %94.sub3:vreg_512_align2 = COPY %39
- %94.sub4:vreg_512_align2 = COPY %39
- %94.sub5:vreg_512_align2 = COPY %39
- %94.sub6:vreg_512_align2 = COPY %39
- %94.sub7:vreg_512_align2 = COPY %39
- %94.sub8:vreg_512_align2 = COPY %39
- %94.sub9:vreg_512_align2 = COPY %39
- %94.sub10:vreg_512_align2 = COPY %39
- %94.sub11:vreg_512_align2 = COPY %39
- %94.sub12:vreg_512_align2 = COPY %39
- %94.sub13:vreg_512_align2 = COPY %39
- %94.sub14:vreg_512_align2 = COPY %39
- %94.sub15:vreg_512_align2 = COPY %39
- %95:vreg_512_align2 = COPY killed %94
- %95:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %95, 0, 0, 0, implicit $mode, implicit $exec
+ %92:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %92, 0, 0, 0, implicit $mode, implicit $exec
+ %93:vreg_512_align2 = COPY killed %92
+ %93:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %93, 0, 0, 0, implicit $mode, implicit $exec
+ %94:vreg_512_align2 = COPY killed %93
+ %94:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %94, 0, 0, 0, implicit $mode, implicit $exec
+ %95:vreg_512_align2 = COPY killed %79
+ %95:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %95, 0, 0, 0, implicit $mode, implicit $exec
%96:vreg_512_align2 = COPY killed %95
- %96:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %96, 0, 0, 0, implicit $mode, implicit $exec
- %31:av_32 = COPY killed %96.sub0
- %32:av_32 = COPY killed %62.sub0
- %33:av_32 = COPY killed %80.sub0
- %97:vreg_512_align2 = COPY killed %57
- %97:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %97, 0, 0, 0, implicit $mode, implicit $exec
- %34:av_32 = COPY killed %97.sub0
- %98:vreg_512_align2 = COPY killed %81
- %98:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %98, 0, 0, 0, implicit $mode, implicit $exec
- %35:av_32 = COPY killed %98.sub0
- %36:av_32 = COPY killed %64.sub0
- %37:av_32 = COPY %70.sub0
- %99:vreg_512_align2 = COPY killed %70
- %99:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %44, %44, %99, 0, 0, 0, implicit $mode, implicit $exec
- %38:av_32 = COPY killed %99.sub0
- $vcc = COPY %100
- %102:av_32 = COPY killed %21
- %103:av_32 = COPY killed %22
- %104:av_32 = COPY killed %23
- %105:av_32 = COPY killed %24
- %106:av_32 = COPY killed %3
- %107:av_32 = COPY killed %25
- %108:av_32 = COPY killed %26
- %109:av_32 = COPY killed %27
- %110:av_32 = COPY killed %28
- %111:av_32 = COPY killed %29
- %112:av_32 = COPY killed %30
- %113:av_32 = COPY killed %31
- %114:av_32 = COPY killed %32
- %115:av_32 = COPY killed %33
- %116:av_32 = COPY killed %13
- %117:av_32 = COPY killed %34
- %118:av_32 = COPY killed %17
- %119:av_32 = COPY killed %35
- %120:av_32 = COPY killed %36
- %121:av_32 = COPY killed %37
- %122:av_32 = COPY killed %38
+ %96:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %96, 0, 0, 0, implicit $mode, implicit $exec
+ %97:vreg_512_align2 = COPY killed %96
+ %97:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %97, 0, 0, 0, implicit $mode, implicit $exec
+ %98:vreg_512_align2 = COPY killed %97
+ %98:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %98, 0, 0, 0, implicit $mode, implicit $exec
+ %99:vreg_512_align2 = COPY killed %73
+ %99:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %99, 0, 0, 0, implicit $mode, implicit $exec
+ early-clobber %100:vreg_512_align2 = V_MFMA_F32_32X32X8F16_vgprcd_e64 %24, %24, %99, 0, 0, 0, implicit $mode, implicit $exec
+ %101:vreg_512_align2 = COPY killed %100
+ %101:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %101, 0, 0, 0, implicit $mode, implicit $exec
+ %102:vreg_512_align2 = COPY killed %101
+ %102:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %102, 0, 0, 0, implicit $mode, implicit $exec
+ %103:vreg_512_align2 = COPY killed %102
+ %103:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %103, 0, 0, 0, implicit $mode, implicit $exec
+ %104:vreg_512_align2 = COPY killed %80
+ %104:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %104, 0, 0, 0, implicit $mode, implicit $exec
+ %105:vreg_512_align2 = COPY killed %104
+ %105:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %105, 0, 0, 0, implicit $mode, implicit $exec
+ undef %106.sub0:vreg_512_align2 = COPY killed %52
+ %106.sub1:vreg_512_align2 = COPY %0
+ %106.sub2:vreg_512_align2 = COPY %0
+ %106.sub3:vreg_512_align2 = COPY %0
+ %106.sub4:vreg_512_align2 = COPY %0
+ %106.sub5:vreg_512_align2 = COPY %0
+ %106.sub6:vreg_512_align2 = COPY %0
+ %106.sub7:vreg_512_align2 = COPY %0
+ %106.sub8:vreg_512_align2 = COPY %0
+ %106.sub9:vreg_512_align2 = COPY %0
+ %106.sub10:vreg_512_align2 = COPY %0
+ %106.sub11:vreg_512_align2 = COPY %0
+ %106.sub12:vreg_512_align2 = COPY %0
+ %106.sub13:vreg_512_align2 = COPY %0
+ %106.sub14:vreg_512_align2 = COPY %0
+ %106.sub15:vreg_512_align2 = COPY %0
+ %107:vreg_512_align2 = COPY killed %106
+ %107:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %107, 0, 0, 0, implicit $mode, implicit $exec
+ %108:vreg_512_align2 = COPY killed %85
+ %108:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %108, 0, 0, 0, implicit $mode, implicit $exec
+ %109:vreg_512_align2 = COPY killed %108
+ %109:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %109, 0, 0, 0, implicit $mode, implicit $exec
+ %110:vreg_512_align2 = COPY killed %90
+ %110:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %110, 0, 0, 0, implicit $mode, implicit $exec
+ %111:vreg_512_align2 = COPY killed %110
+ %111:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %111, 0, 0, 0, implicit $mode, implicit $exec
+ %112:vreg_512_align2 = COPY killed %94
+ %112:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %112, 0, 0, 0, implicit $mode, implicit $exec
+ %113:vreg_512_align2 = COPY killed %112
+ %113:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %113, 0, 0, 0, implicit $mode, implicit $exec
+ %114:vreg_512_align2 = COPY killed %77
+ %114:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %114, 0, 0, 0, implicit $mode, implicit $exec
+ %115:vreg_512_align2 = COPY killed %114
+ %115:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %115, 0, 0, 0, implicit $mode, implicit $exec
+ %116:vreg_512_align2 = COPY killed %98
+ %116:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %116, 0, 0, 0, implicit $mode, implicit $exec
+ %117:vreg_512_align2 = COPY killed %103
+ %117:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %117, 0, 0, 0, implicit $mode, implicit $exec
+ %118:av_32 = COPY killed %109.sub0
+ %119:vreg_512_align2 = COPY killed %111
+ %119:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %119, 0, 0, 0, implicit $mode, implicit $exec
+ %120:av_32 = COPY killed %119.sub1
+ %121:av_32 = COPY killed %105.sub0
+ %122:av_32 = COPY killed %83.sub0
+ %123:vreg_512_align2 = COPY killed %71
+ %123:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %123, 0, 0, 0, implicit $mode, implicit $exec
+ %124:av_32 = COPY killed %123.sub0
+ undef %125.sub0:vreg_512_align2 = COPY killed %62
+ %125.sub1:vreg_512_align2 = COPY %0
+ %125.sub2:vreg_512_align2 = COPY %0
+ %125.sub3:vreg_512_align2 = COPY %0
+ %125.sub4:vreg_512_align2 = COPY %0
+ %125.sub5:vreg_512_align2 = COPY %0
+ %125.sub6:vreg_512_align2 = COPY %0
+ %125.sub7:vreg_512_align2 = COPY %0
+ %125.sub8:vreg_512_align2 = COPY %0
+ %125.sub9:vreg_512_align2 = COPY %0
+ %125.sub10:vreg_512_align2 = COPY %0
+ %125.sub11:vreg_512_align2 = COPY %0
+ %125.sub12:vreg_512_align2 = COPY %0
+ %125.sub13:vreg_512_align2 = COPY %0
+ %125.sub14:vreg_512_align2 = COPY %0
+ %125.sub15:vreg_512_align2 = COPY %0
+ %126:vreg_512_align2 = COPY killed %125
+ %126:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %126, 0, 0, 0, implicit $mode, implicit $exec
+ %127:av_32 = COPY killed %126.sub1
+ %128:vreg_512_align2 = COPY killed %113
+ %128:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %128, 0, 0, 0, implicit $mode, implicit $exec
+ %129:vreg_512_align2 = COPY killed %128
+ %129:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %129, 0, 0, 0, implicit $mode, implicit $exec
+ %130:av_32 = COPY killed %129.sub0
+ %131:av_32 = COPY killed %84.sub0
+ %132:vreg_512_align2 = COPY killed %72
+ %132:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %132, 0, 0, 0, implicit $mode, implicit $exec
+ %133:vreg_512_align2 = COPY killed %132
+ %133:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %133, 0, 0, 0, implicit $mode, implicit $exec
+ %134:av_32 = COPY killed %133.sub0
+ %135:vreg_512_align2 = COPY killed %115
+ %135:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %135, 0, 0, 0, implicit $mode, implicit $exec
+ %136:av_32 = COPY killed %135.sub0
+ %137:vreg_512_align2 = COPY killed %116
+ %137:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %137, 0, 0, 0, implicit $mode, implicit $exec
+ %138:vreg_512_align2 = COPY killed %137
+ %138:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %138, 0, 0, 0, implicit $mode, implicit $exec
+ %139:av_32 = COPY killed %138.sub0
+ %140:vreg_512_align2 = COPY killed %107
+ %140:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %140, 0, 0, 0, implicit $mode, implicit $exec
+ %141:av_32 = COPY killed %140.sub0
+ %142:vreg_512_align2 = COPY killed %117
+ %142:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %142, 0, 0, 0, implicit $mode, implicit $exec
+ %143:vreg_512_align2 = COPY killed %142
+ %143:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %143, 0, 0, 0, implicit $mode, implicit $exec
+ %144:av_32 = COPY killed %143.sub0
+ %145:av_32 = COPY killed %99.sub0
+ undef %146.sub0:vreg_512_align2 = COPY %0
+ %146.sub1:vreg_512_align2 = COPY killed %49
+ %146.sub2:vreg_512_align2 = COPY %0
+ %146.sub3:vreg_512_align2 = COPY %0
+ %146.sub4:vreg_512_align2 = COPY %0
+ %146.sub5:vreg_512_align2 = COPY %0
+ %146.sub6:vreg_512_align2 = COPY %0
+ %146.sub7:vreg_512_align2 = COPY %0
+ %146.sub8:vreg_512_align2 = COPY %0
+ %146.sub9:vreg_512_align2 = COPY %0
+ %146.sub10:vreg_512_align2 = COPY %0
+ %146.sub11:vreg_512_align2 = COPY %0
+ %146.sub12:vreg_512_align2 = COPY %0
+ %146.sub13:vreg_512_align2 = COPY %0
+ %146.sub14:vreg_512_align2 = COPY %0
+ %146.sub15:vreg_512_align2 = COPY %0
+ %147:vreg_512_align2 = COPY killed %146
+ %147:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %147, 0, 0, 0, implicit $mode, implicit $exec
+ %148:vreg_512_align2 = COPY killed %147
+ %148:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %148, 0, 0, 0, implicit $mode, implicit $exec
+ %149:av_32 = COPY killed %148.sub0
+ %150:av_32 = COPY killed %47
+ undef %151.sub0:vreg_512_align2 = COPY killed %150
+ %151.sub1:vreg_512_align2 = COPY %0
+ %151.sub2:vreg_512_align2 = COPY %0
+ %151.sub3:vreg_512_align2 = COPY %0
+ %151.sub4:vreg_512_align2 = COPY %0
+ %151.sub5:vreg_512_align2 = COPY %0
+ %151.sub6:vreg_512_align2 = COPY %0
+ %151.sub7:vreg_512_align2 = COPY %0
+ %151.sub8:vreg_512_align2 = COPY %0
+ %151.sub9:vreg_512_align2 = COPY %0
+ %151.sub10:vreg_512_align2 = COPY %0
+ %151.sub11:vreg_512_align2 = COPY %0
+ %151.sub12:vreg_512_align2 = COPY %0
+ %151.sub13:vreg_512_align2 = COPY %0
+ %151.sub14:vreg_512_align2 = COPY %0
+ %151.sub15:vreg_512_align2 = COPY %0
+ %152:vreg_512_align2 = COPY killed %151
+ %152:vreg_512_align2 = V_MFMA_F32_32X32X8F16_mac_vgprcd_e64 %24, %24, %152, 0, 0, 0, implicit $mode, implicit $exec
+ %153:av_32 = COPY killed %152.sub0
+ $vcc = COPY %26
+ %27:av_32 = COPY killed %118
+ %28:av_32 = COPY killed %120
+ %29:av_32 = COPY killed %121
+ %30:av_32 = COPY killed %122
+ %31:av_32 = COPY killed %124
+ %32:av_32 = COPY killed %127
+ %33:av_32 = COPY killed %130
+ %34:av_32 = COPY killed %61
+ %35:av_32 = COPY killed %60
+ %36:vgpr_32 = COPY %21
+ %37:av_32 = COPY killed %56
+ %38:av_32 = COPY killed %131
+ %39:av_32 = COPY killed %134
+ %40:av_32 = COPY killed %136
+ %41:av_32 = COPY killed %139
+ %42:av_32 = COPY killed %141
+ %43:av_32 = COPY killed %144
+ %44:av_32 = COPY killed %145
+ %45:av_32 = COPY killed %149
+ %46:av_32 = COPY killed %153
+ %47:av_32 = COPY killed %48
S_CBRANCH_VCCNZ %bb.1, implicit killed $vcc
S_BRANCH %bb.2
-
- bb.2.DummyReturnBlock:
+
+ bb.2:
S_ENDPGM 0
...
diff --git a/llvm/test/CodeGen/AMDGPU/sched-dynamic-vgpr-blocks.mir b/llvm/test/CodeGen/AMDGPU/sched-dynamic-vgpr-blocks.mir
index 5b588783efb2015..a103a039370a1c8 100644
--- a/llvm/test/CodeGen/AMDGPU/sched-dynamic-vgpr-blocks.mir
+++ b/llvm/test/CodeGen/AMDGPU/sched-dynamic-vgpr-blocks.mir
@@ -4,7 +4,7 @@
# In dynamic VGPR mode we don't want to waste VGPR blocks, so we revert any
# schedule that increases the number of blocks we need to allocate. The
# scheduler will try to hoist the loads above the dependent adds, which will
-# increase the pressure in the region from 5 to 17 VGPRs, i.e. from one block of
+# increase the pressure in the region from 5 to 18 VGPRs, i.e. from one block of
# 16 to 2. This should trigger the code that performs the revert based on the
# number of VGPR blocks.
# The -amdgpu-stress-vgpr flag is only there to keep the test small (the
@@ -16,7 +16,7 @@
...
---
# DEBUG: Region register pressure: VGPRs: 5
-# DEBUG: Pressure after scheduling: VGPRs: 17
+# DEBUG: Pressure after scheduling: VGPRs: 18
# DEBUG-NEXT: Region: 0
# DEBUG-NEXT: Occupancy before scheduling: 16, after 16.
# DEBUG-NEXT: Reverting scheduling for region 0
@@ -68,5 +68,7 @@ body: |
%acc14:vgpr_32 = V_ADD_U32_e32 %acc13, %l13, implicit $exec
%l14:vgpr_32 = GLOBAL_LOAD_DWORD %addr, 56, 0, implicit $exec
%acc15:vgpr_32 = V_ADD_U32_e32 %acc14, %l14, implicit $exec
- S_ENDPGM 0, implicit %acc15
+ %l15:vgpr_32 = GLOBAL_LOAD_DWORD %addr, 60, 0, implicit $exec
+ %acc16:vgpr_32 = V_ADD_U32_e32 %acc15, %l15, implicit $exec
+ S_ENDPGM 0, implicit %acc16
...
diff --git a/llvm/test/CodeGen/AMDGPU/sema-v-unsched-bundle.ll b/llvm/test/CodeGen/AMDGPU/sema-v-unsched-bundle.ll
index d3788c65249b9ca..9bde7aaedb0e03e 100644
--- a/llvm/test/CodeGen/AMDGPU/sema-v-unsched-bundle.ll
+++ b/llvm/test/CodeGen/AMDGPU/sema-v-unsched-bundle.ll
@@ -3,15 +3,23 @@
; DEBUG: Reverting scheduling for region 0
+; Three vectors are kept live across the GWS bundle so that the region stays
+; well above the occupancy limit and scheduling is reverted, which is what
+; exercises the unscheduling of bundled instructions.
+
@G = global <32 x i8> splat (i8 1)
@G.1 = global <32 x i8> splat (i8 127)
+ at G.2 = global <32 x i8> splat (i8 63)
-define amdgpu_kernel void @gws_sema_v_offset0(i32 %val, <32 x i1>* %inp) {
+define amdgpu_kernel void @gws_sema_v_offset0(i32 %val, <32 x i1>* %inp, <32 x i1>* %inp2) {
%LGV1 = load <32 x i8>, ptr @G.1, align 32
%LGV = load <32 x i8>, ptr @G, align 32
+ %LGV2 = load <32 x i8>, ptr @G.2, align 32
call void @llvm.amdgcn.ds.gws.sema.v(i32 0)
%C = icmp ne <32 x i8> %LGV, %LGV1
+ %C2 = icmp ne <32 x i8> %LGV1, %LGV2
store <32 x i1> %C, ptr %inp, align 4
+ store <32 x i1> %C2, ptr %inp2, align 4
ret void
}
More information about the llvm-commits
mailing list