[llvm] [AMDGPU][NFC] Harden scheduler tests balanced on pressure boundaries (PR #228293)
via llvm-commits
llvm-commits at lists.llvm.org
Thu Oct 1 18:02:04 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-amdgpu
Author: Dhruva Chakrabarti (dhruvachak)
<details>
<summary>Changes</summary>
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
---
Patch is 48.19 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/228293.diff
4 Files Affected:
- (modified) llvm/test/CodeGen/AMDGPU/coexec-rewrite-mfma.ll (+10-2)
- (modified) llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-spill-multi-store-mir.mir (+458-423)
- (modified) llvm/test/CodeGen/AMDGPU/sched-dynamic-vgpr-blocks.mir (+5-3)
- (modified) llvm/test/CodeGen/AMDGPU/sema-v-unsched-bundle.ll (+9-1)
``````````diff
diff --git a/llvm/test/CodeGen/AMDGPU/coexec-rewrite-mfma.ll b/llvm/test/CodeGen/AMDGPU/coexec-rewrite-mfma.ll
index a632f5b3d161d..f73edff901bbf 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 111f86aff8b53..8ce7f38657aea 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
+ %6...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/228293
More information about the llvm-commits
mailing list