[llvm-branch-commits] [llvm] AMDGPU: Add baseline test for vgpr fma with copied-from AGPR (PR #153020)
via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Mon Aug 11 07:40:44 PDT 2025
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-amdgpu
Author: Matt Arsenault (arsenm)
<details>
<summary>Changes</summary>
We currently handle the case where an MFMA is copied to an AGPR,
but not the case where the MFMA is copying from an AGPR.
---
Patch is 34.43 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/153020.diff
2 Files Affected:
- (added) llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-copy-from.mir (+261)
- (modified) llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr.ll (+387)
``````````diff
diff --git a/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-copy-from.mir b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-copy-from.mir
new file mode 100644
index 0000000000000..7fdc8c0d8019b
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-copy-from.mir
@@ -0,0 +1,261 @@
+# NOTE: Assertions have been autogenerated by utils/update_mir_test_checks.py UTC_ARGS: --version 5
+# RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx942 -run-pass=greedy,amdgpu-rewrite-agpr-copy-mfma -verify-machineinstrs -o - %s | FileCheck %s
+
+---
+name: test_rewrite_mfma_copy_from_agpr_physreg
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5, $agpr0_agpr1
+
+ ; CHECK-LABEL: name: test_rewrite_mfma_copy_from_agpr_physreg
+ ; CHECK: liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5, $agpr0_agpr1
+ ; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:vreg_64_align2 = COPY $vgpr4_vgpr5
+ ; CHECK-NEXT: [[COPY1:%[0-9]+]]:av_64_align2 = COPY $vgpr0_vgpr1
+ ; CHECK-NEXT: [[COPY2:%[0-9]+]]:av_64_align2 = COPY $vgpr2_vgpr3
+ ; CHECK-NEXT: [[COPY3:%[0-9]+]]:vreg_64_align2 = COPY $agpr0_agpr1
+ ; CHECK-NEXT: [[V_MFMA_F64_4X4X4F64_vgprcd_e64_:%[0-9]+]]:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]], [[COPY3]], 0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: GLOBAL_STORE_DWORDX2 [[COPY]], [[V_MFMA_F64_4X4X4F64_vgprcd_e64_]], 0, 0, implicit $exec :: (store (s64), addrspace 1)
+ ; CHECK-NEXT: SI_RETURN
+ %0:vreg_64_align2 = COPY $vgpr4_vgpr5
+ %1:av_64_align2 = COPY $vgpr0_vgpr1
+ %2:av_64_align2 = COPY $vgpr2_vgpr3
+ %3:vreg_64_align2 = COPY $agpr0_agpr1
+ %4:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, %3, 0, 0, 0, implicit $mode, implicit $exec
+ GLOBAL_STORE_DWORDX2 %0, %4, 0, 0, implicit $exec :: (store (s64), addrspace 1)
+ SI_RETURN
+...
+
+---
+name: test_rewrite_mfma_copy_from_agpr_unrewritable_use
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+
+ ; CHECK-LABEL: name: test_rewrite_mfma_copy_from_agpr_unrewritable_use
+ ; CHECK: liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+ ; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:vreg_64_align2 = COPY $vgpr4_vgpr5
+ ; CHECK-NEXT: [[COPY1:%[0-9]+]]:av_64_align2 = COPY $vgpr0_vgpr1
+ ; CHECK-NEXT: [[COPY2:%[0-9]+]]:av_64_align2 = COPY $vgpr2_vgpr3
+ ; CHECK-NEXT: [[GLOBAL_LOAD_DWORDX4_:%[0-9]+]]:areg_128_align2 = GLOBAL_LOAD_DWORDX4 [[COPY]], 0, 0, implicit $exec :: (load (s128), addrspace 1)
+ ; CHECK-NEXT: [[COPY3:%[0-9]+]]:vreg_128_align2 = COPY [[GLOBAL_LOAD_DWORDX4_]]
+ ; CHECK-NEXT: [[V_MFMA_F64_4X4X4F64_vgprcd_e64_:%[0-9]+]]:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]], [[COPY3]].sub0_sub1, 0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: INLINEASM &"; use $0", 1 /* sideeffect attdialect */, 3866633 /* reguse:VReg_64_Align2 */, [[V_MFMA_F64_4X4X4F64_vgprcd_e64_]]
+ ; CHECK-NEXT: SI_RETURN
+ %0:vreg_64_align2 = COPY $vgpr4_vgpr5
+ %1:av_64_align2 = COPY $vgpr0_vgpr1
+ %2:av_64_align2 = COPY $vgpr2_vgpr3
+ %3:areg_128_align2 = GLOBAL_LOAD_DWORDX4 %0, 0, 0, implicit $exec :: (load (s128), addrspace 1)
+ %4:vreg_128_align2 = COPY %3
+ %5:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, %4.sub0_sub1, 0, 0, 0, implicit $mode, implicit $exec
+ INLINEASM &"; use $0", 1 /* sideeffect attdialect */, 3866633 /* reguse:VReg_64_Align2 */, %5
+ SI_RETURN
+...
+
+---
+name: test_rewrite_mfma_copy_from_agpr_src2_subreg_use
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+
+ ; CHECK-LABEL: name: test_rewrite_mfma_copy_from_agpr_src2_subreg_use
+ ; CHECK: liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+ ; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:vreg_64_align2 = COPY $vgpr4_vgpr5
+ ; CHECK-NEXT: [[COPY1:%[0-9]+]]:av_64_align2 = COPY $vgpr0_vgpr1
+ ; CHECK-NEXT: [[COPY2:%[0-9]+]]:av_64_align2 = COPY $vgpr2_vgpr3
+ ; CHECK-NEXT: [[GLOBAL_LOAD_DWORDX4_:%[0-9]+]]:areg_128_align2 = GLOBAL_LOAD_DWORDX4 [[COPY]], 0, 0, implicit $exec :: (load (s128), addrspace 1)
+ ; CHECK-NEXT: [[COPY3:%[0-9]+]]:vreg_128_align2 = COPY [[GLOBAL_LOAD_DWORDX4_]]
+ ; CHECK-NEXT: [[V_MFMA_F64_4X4X4F64_vgprcd_e64_:%[0-9]+]]:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]], [[COPY3]].sub0_sub1, 0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: GLOBAL_STORE_DWORDX2 [[COPY]], [[V_MFMA_F64_4X4X4F64_vgprcd_e64_]], 0, 0, implicit $exec :: (store (s64), addrspace 1)
+ ; CHECK-NEXT: SI_RETURN
+ %0:vreg_64_align2 = COPY $vgpr4_vgpr5
+ %1:av_64_align2 = COPY $vgpr0_vgpr1
+ %2:av_64_align2 = COPY $vgpr2_vgpr3
+ %3:areg_128_align2 = GLOBAL_LOAD_DWORDX4 %0, 0, 0, implicit $exec :: (load (s128), addrspace 1)
+ %4:vreg_128_align2 = COPY %3
+ %5:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, %4.sub0_sub1, 0, 0, 0, implicit $mode, implicit $exec
+ GLOBAL_STORE_DWORDX2 %0, %5, 0, 0, implicit $exec :: (store (s64), addrspace 1)
+ SI_RETURN
+...
+
+---
+name: test_rewrite_mfma_copy_from_agpr_vdst_subreg_use
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+
+ ; CHECK-LABEL: name: test_rewrite_mfma_copy_from_agpr_vdst_subreg_use
+ ; CHECK: liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+ ; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:vreg_64_align2 = COPY $vgpr4_vgpr5
+ ; CHECK-NEXT: [[COPY1:%[0-9]+]]:av_64_align2 = COPY $vgpr0_vgpr1
+ ; CHECK-NEXT: [[COPY2:%[0-9]+]]:av_64_align2 = COPY $vgpr2_vgpr3
+ ; CHECK-NEXT: [[GLOBAL_LOAD_DWORDX4_:%[0-9]+]]:areg_128_align2 = GLOBAL_LOAD_DWORDX4 [[COPY]], 0, 0, implicit $exec :: (load (s128), addrspace 1)
+ ; CHECK-NEXT: [[COPY3:%[0-9]+]]:vreg_128_align2 = COPY [[GLOBAL_LOAD_DWORDX4_]]
+ ; CHECK-NEXT: [[COPY3:%[0-9]+]].sub0_sub1:vreg_128_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]], [[COPY3]].sub2_sub3, 0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: GLOBAL_STORE_DWORDX4 [[COPY]], [[COPY3]], 0, 0, implicit $exec :: (store (s128), addrspace 1)
+ ; CHECK-NEXT: SI_RETURN
+ %0:vreg_64_align2 = COPY $vgpr4_vgpr5
+ %1:av_64_align2 = COPY $vgpr0_vgpr1
+ %2:av_64_align2 = COPY $vgpr2_vgpr3
+ %3:areg_128_align2 = GLOBAL_LOAD_DWORDX4 %0, 0, 0, implicit $exec :: (load (s128), addrspace 1)
+ %4:vreg_128_align2 = COPY %3
+ %4.sub0_sub1:vreg_128_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, %4.sub2_sub3, 0, 0, 0, implicit $mode, implicit $exec
+ GLOBAL_STORE_DWORDX4 %0, %4, 0, 0, implicit $exec :: (store (s128), addrspace 1)
+ SI_RETURN
+...
+
+# A-to-V copy is performed subregister at a time instead.
+---
+name: test_rewrite_mfma_copy_from_agpr_split_copy
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+
+ ; CHECK-LABEL: name: test_rewrite_mfma_copy_from_agpr_split_copy
+ ; CHECK: liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+ ; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:vreg_64_align2 = COPY $vgpr4_vgpr5
+ ; CHECK-NEXT: [[COPY1:%[0-9]+]]:av_64_align2 = COPY $vgpr0_vgpr1
+ ; CHECK-NEXT: [[COPY2:%[0-9]+]]:av_64_align2 = COPY $vgpr2_vgpr3
+ ; CHECK-NEXT: [[GLOBAL_LOAD_DWORDX2_:%[0-9]+]]:areg_64_align2 = GLOBAL_LOAD_DWORDX2 [[COPY]], 0, 0, implicit $exec :: (load (s64), addrspace 1)
+ ; CHECK-NEXT: undef [[COPY3:%[0-9]+]].sub0:vreg_64_align2 = COPY [[GLOBAL_LOAD_DWORDX2_]].sub0
+ ; CHECK-NEXT: [[COPY3:%[0-9]+]].sub1:vreg_64_align2 = COPY [[GLOBAL_LOAD_DWORDX2_]].sub1
+ ; CHECK-NEXT: [[V_MFMA_F64_4X4X4F64_vgprcd_e64_:%[0-9]+]]:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]], [[COPY3]], 0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: GLOBAL_STORE_DWORDX2 [[COPY]], [[V_MFMA_F64_4X4X4F64_vgprcd_e64_]], 0, 0, implicit $exec :: (store (s64), addrspace 1)
+ ; CHECK-NEXT: SI_RETURN
+ %0:vreg_64_align2 = COPY $vgpr4_vgpr5
+ %1:av_64_align2 = COPY $vgpr0_vgpr1
+ %2:av_64_align2 = COPY $vgpr2_vgpr3
+ %3:areg_64_align2 = GLOBAL_LOAD_DWORDX2 %0, 0, 0, implicit $exec :: (load (s64), addrspace 1)
+ undef %4.sub0:vreg_64_align2 = COPY %3.sub0
+ %4.sub1:vreg_64_align2 = COPY %3.sub1
+ %5:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, %4, 0, 0, 0, implicit $mode, implicit $exec
+ GLOBAL_STORE_DWORDX2 %0, %5, 0, 0, implicit $exec :: (store (s64), addrspace 1)
+ SI_RETURN
+...
+
+---
+name: test_rewrite_mfma_copy_from_agpr_copyback
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+
+ ; CHECK-LABEL: name: test_rewrite_mfma_copy_from_agpr_copyback
+ ; CHECK: liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+ ; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:vreg_64_align2 = COPY $vgpr4_vgpr5
+ ; CHECK-NEXT: [[COPY1:%[0-9]+]]:av_64_align2 = COPY $vgpr0_vgpr1
+ ; CHECK-NEXT: [[COPY2:%[0-9]+]]:av_64_align2 = COPY $vgpr2_vgpr3
+ ; CHECK-NEXT: [[GLOBAL_LOAD_DWORDX2_:%[0-9]+]]:areg_64_align2 = GLOBAL_LOAD_DWORDX2 [[COPY]], 0, 0, implicit $exec :: (load (s64), addrspace 1)
+ ; CHECK-NEXT: [[COPY3:%[0-9]+]]:areg_64_align2 = COPY [[GLOBAL_LOAD_DWORDX2_]]
+ ; CHECK-NEXT: [[V_MFMA_F64_4X4X4F64_e64_:%[0-9]+]]:areg_64_align2 = V_MFMA_F64_4X4X4F64_e64 [[COPY1]], [[COPY2]], [[COPY3]], 0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[COPY4:%[0-9]+]]:areg_64_align2 = COPY [[V_MFMA_F64_4X4X4F64_e64_]]
+ ; CHECK-NEXT: GLOBAL_STORE_DWORDX2 [[COPY]], [[COPY4]], 0, 0, implicit $exec :: (store (s64), addrspace 1)
+ ; CHECK-NEXT: SI_RETURN
+ %0:vreg_64_align2 = COPY $vgpr4_vgpr5
+ %1:av_64_align2 = COPY $vgpr0_vgpr1
+ %2:av_64_align2 = COPY $vgpr2_vgpr3
+ %3:areg_64_align2 = GLOBAL_LOAD_DWORDX2 %0, 0, 0, implicit $exec :: (load (s64), addrspace 1)
+ %4:vreg_64_align2 = COPY %3
+ %5:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, %4, 0, 0, 0, implicit $mode, implicit $exec
+ %6:areg_64_align2 = COPY %5
+ GLOBAL_STORE_DWORDX2 %0, %6, 0, 0, implicit $exec :: (store (s64), addrspace 1)
+ SI_RETURN
+...
+
+# There is a read of the coy-from-agpr in the dst operand of the MFMA.
+---
+name: test_rewrite_mfma_copy_from_agpr_vdst_subreg_use_imm_src2
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+
+ ; CHECK-LABEL: name: test_rewrite_mfma_copy_from_agpr_vdst_subreg_use_imm_src2
+ ; CHECK: liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+ ; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:vreg_64_align2 = COPY $vgpr4_vgpr5
+ ; CHECK-NEXT: [[COPY1:%[0-9]+]]:av_64_align2 = COPY $vgpr0_vgpr1
+ ; CHECK-NEXT: [[COPY2:%[0-9]+]]:av_64_align2 = COPY $vgpr2_vgpr3
+ ; CHECK-NEXT: [[GLOBAL_LOAD_DWORDX4_:%[0-9]+]]:areg_128_align2 = GLOBAL_LOAD_DWORDX4 [[COPY]], 0, 0, implicit $exec :: (load (s128), addrspace 1)
+ ; CHECK-NEXT: [[COPY3:%[0-9]+]]:vreg_128_align2 = COPY [[GLOBAL_LOAD_DWORDX4_]]
+ ; CHECK-NEXT: [[COPY3:%[0-9]+]].sub0_sub1:vreg_128_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]], 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: GLOBAL_STORE_DWORDX4 [[COPY]], [[COPY3]], 0, 0, implicit $exec :: (store (s128), addrspace 1)
+ ; CHECK-NEXT: SI_RETURN
+ %0:vreg_64_align2 = COPY $vgpr4_vgpr5
+ %1:av_64_align2 = COPY $vgpr0_vgpr1
+ %2:av_64_align2 = COPY $vgpr2_vgpr3
+ %3:areg_128_align2 = GLOBAL_LOAD_DWORDX4 %0, 0, 0, implicit $exec :: (load (s128), addrspace 1)
+ %4:vreg_128_align2 = COPY %3
+ %4.sub0_sub1:vreg_128_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, 0, 0, 0, 0, implicit $mode, implicit $exec
+ GLOBAL_STORE_DWORDX4 %0, %4, 0, 0, implicit $exec :: (store (s128), addrspace 1)
+ SI_RETURN
+...
+
+# Degenerate case. Copy from AGPR to VGPR is dead undef subreg def
+---
+name: test_rewrite_mfma_copy_from_agpr_undef_vdst_subreg_use_imm_src2
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+
+ ; CHECK-LABEL: name: test_rewrite_mfma_copy_from_agpr_undef_vdst_subreg_use_imm_src2
+ ; CHECK: liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+ ; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:vreg_64_align2 = COPY $vgpr4_vgpr5
+ ; CHECK-NEXT: [[COPY1:%[0-9]+]]:av_64_align2 = COPY $vgpr0_vgpr1
+ ; CHECK-NEXT: [[COPY2:%[0-9]+]]:av_64_align2 = COPY $vgpr2_vgpr3
+ ; CHECK-NEXT: [[GLOBAL_LOAD_DWORDX4_:%[0-9]+]]:areg_128_align2 = GLOBAL_LOAD_DWORDX4 [[COPY]], 0, 0, implicit $exec :: (load (s128), addrspace 1)
+ ; CHECK-NEXT: dead [[COPY3:%[0-9]+]]:vreg_128_align2 = COPY [[GLOBAL_LOAD_DWORDX4_]]
+ ; CHECK-NEXT: undef [[V_MFMA_F64_4X4X4F64_vgprcd_e64_:%[0-9]+]].sub0_sub1:vreg_128_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]], 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: GLOBAL_STORE_DWORDX4 [[COPY]], [[V_MFMA_F64_4X4X4F64_vgprcd_e64_]], 0, 0, implicit $exec :: (store (s128), addrspace 1)
+ ; CHECK-NEXT: SI_RETURN
+ %0:vreg_64_align2 = COPY $vgpr4_vgpr5
+ %1:av_64_align2 = COPY $vgpr0_vgpr1
+ %2:av_64_align2 = COPY $vgpr2_vgpr3
+ %3:areg_128_align2 = GLOBAL_LOAD_DWORDX4 %0, 0, 0, implicit $exec :: (load (s128), addrspace 1)
+ %4:vreg_128_align2 = COPY %3
+ undef %4.sub0_sub1:vreg_128_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, 0, 0, 0, 0, implicit $mode, implicit $exec
+ GLOBAL_STORE_DWORDX4 %0, %4, 0, 0, implicit $exec :: (store (s128), addrspace 1)
+ SI_RETURN
+...
+
+# Degenerate case. Copy from AGPR to VGPR is dead, but same register
+# is redefined as whole register.
+---
+name: test_rewrite_mfma_copy_from_agpr_to_vdst_def_imm_src2
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+
+ ; CHECK-LABEL: name: test_rewrite_mfma_copy_from_agpr_to_vdst_def_imm_src2
+ ; CHECK: liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+ ; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:vreg_64_align2 = COPY $vgpr4_vgpr5
+ ; CHECK-NEXT: [[COPY1:%[0-9]+]]:av_64_align2 = COPY $vgpr0_vgpr1
+ ; CHECK-NEXT: [[COPY2:%[0-9]+]]:av_64_align2 = COPY $vgpr2_vgpr3
+ ; CHECK-NEXT: [[GLOBAL_LOAD_DWORDX2_:%[0-9]+]]:areg_64_align2 = GLOBAL_LOAD_DWORDX2 [[COPY]], 0, 0, implicit $exec :: (load (s64), addrspace 1)
+ ; CHECK-NEXT: dead [[COPY3:%[0-9]+]]:vreg_64_align2 = COPY [[GLOBAL_LOAD_DWORDX2_]]
+ ; CHECK-NEXT: [[V_MFMA_F64_4X4X4F64_vgprcd_e64_:%[0-9]+]]:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]], 0, 0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: GLOBAL_STORE_DWORDX2 [[COPY]], [[V_MFMA_F64_4X4X4F64_vgprcd_e64_]], 0, 0, implicit $exec :: (store (s64), addrspace 1)
+ ; CHECK-NEXT: SI_RETURN
+ %0:vreg_64_align2 = COPY $vgpr4_vgpr5
+ %1:av_64_align2 = COPY $vgpr0_vgpr1
+ %2:av_64_align2 = COPY $vgpr2_vgpr3
+ %3:areg_64_align2 = GLOBAL_LOAD_DWORDX2 %0, 0, 0, implicit $exec :: (load (s64), addrspace 1)
+ %4:vreg_64_align2 = COPY %3
+ %4:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, 0, 0, 0, 0, implicit $mode, implicit $exec
+ GLOBAL_STORE_DWORDX2 %0, %4, 0, 0, implicit $exec :: (store (s64), addrspace 1)
+ SI_RETURN
+...
diff --git a/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr.ll b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr.ll
index 58ec41e3928bd..81613f69c982b 100644
--- a/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr.ll
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr.ll
@@ -522,7 +522,394 @@ define void @test_rewrite_mfma_subreg_insert2(double %arg0, double %arg1, ptr ad
ret void
}
+define amdgpu_kernel void @test_rewrite_mfma_direct_copy_from_agpr_class(ptr addrspace(1) %arg0, ptr addrspace(1) %arg1) #0 {
+; CHECK-LABEL: test_rewrite_mfma_direct_copy_from_agpr_class:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: ;;#ASMSTART
+; CHECK-NEXT: ; def a[0:31]
+; CHECK-NEXT: ;;#ASMEND
+; CHECK-NEXT: v_accvgpr_write_b32 a32, v0
+; CHECK-NEXT: v_accvgpr_read_b32 v63, a31
+; CHECK-NEXT: v_accvgpr_read_b32 v62, a30
+; CHECK-NEXT: v_accvgpr_read_b32 v61, a29
+; CHECK-NEXT: v_accvgpr_read_b32 v60, a28
+; CHECK-NEXT: v_accvgpr_read_b32 v59, a27
+; CHECK-NEXT: v_accvgpr_read_b32 v58, a26
+; CHECK-NEXT: v_accvgpr_read_b32 v57, a25
+; CHECK-NEXT: v_accvgpr_read_b32 v56, a24
+; CHECK-NEXT: v_accvgpr_read_b32 v55, a23
+; CHECK-NEXT: v_accvgpr_read_b32 v54, a22
+; CHECK-NEXT: v_accvgpr_read_b32 v53, a21
+; CHECK-NEXT: v_accvgpr_read_b32 v52, a20
+; CHECK-NEXT: v_accvgpr_read_b32 v51, a19
+; CHECK-NEXT: v_accvgpr_read_b32 v50, a18
+; CHECK-NEXT: v_accvgpr_read_b32 v49, a17
+; CHECK-NEXT: v_accvgpr_read_b32 v48, a16
+; CHECK-NEXT: v_accvgpr_read_b32 v47, a15
+; CHECK-NEXT: v_accvgpr_read_b32 v46, a14
+; CHECK-NEXT: v_accvgpr_read_b32 v45, a13
+; CHECK-NEXT: v_accvgpr_read_b32 v44, a12
+; CHECK-NEXT: v_accvgpr_read_b32 v43, a11
+; CHECK-NEXT: v_accvgpr_read_b32 v42, a10
+; CHECK-NEXT: v_accvgpr_read_b32 v41, a9
+; CHECK-NEXT: v_accvgpr_read_b32 v40, a8
+; CHECK-NEXT: v_accvgpr_read_b32 v39, a7
+; CHECK-NEXT: v_accvgpr_read_b32 v38, a6
+; CHECK-NEXT: v_accvgpr_read_b32 v37, a5
+; CHECK-NEXT: v_accvgpr_read_b32 v36, a4
+; CHECK-NEXT: v_accvgpr_read_b32 v35, a3
+; CHECK-NEXT: v_accvgpr_read_b32 v34, a2
+; CHECK-NEXT: v_accvgpr_read_b32 v33, a1
+; CHECK-NEXT: v_accvgpr_read_b32 v32, a0
+; CHECK-NEXT: v_accvgpr_write_b32 a0, 2.0
+; CHECK-NEXT: v_accvgpr_write_b32 a1, 4.0
+; CHECK-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x0
+; CHECK-NEXT: s_nop 0
+; CHECK-NEXT: v_mfma_f32_32x32x1_2b_f32 v[0:31], a0, a1, v[32:63]
+; CHECK-NEXT: v_accvgpr_write_b32 a0, v32
+; CHECK-NEXT: v_accvgpr_write_b32 a1, v33
+; CHECK-NEXT: v_accvgpr_write_b32 a2, v34
+; CHECK-NEXT: v_accvgpr_write_b32 a3, v35
+; CHECK-NEXT: v_accvgpr_write_b32 a4, v36
+; CHECK-NEXT: v_accvgpr_write_b32 a5, v37
+; CHECK-NEXT: v_accvgpr_write_b32 a6, v38
+; CHECK-NEXT: v_accvgpr_write_b32 a7, v39
+; CHECK-NEXT: v_accvgpr_write_b32 a8, v40
+; CHECK-NEXT: v_accvgpr_write_b32 a9, v41
+; CHECK-NEXT: v_accvgpr_write_b32 a10, v42
+; CHECK-NEXT: v_accvgpr_write_b32 a11, v43
+; CHECK-NEXT: v_accvgpr_write_b32 a12, v44
+; CHECK-NEXT: v_accvgpr_write_b32 a13, v45
+; CHECK-NEXT: v_accvgpr_write_b32 a14, v46
+; CHECK-NEXT: v_accvgpr_write_b32 a15, v47
+; CHECK-NEXT: v_accvgpr_write_b32 a16, v48
+; CHECK-NEXT: v_accvgpr_write_b32 a17, v49
+; CHECK-NEXT: v_accvgpr_write_b32 a18, v50
+; CHECK-NEXT: v_accvgpr_write_b32 a19, v51
+; CHECK-NEXT: v_accvgpr_write_b32 a20, v52
+; CHECK-NEXT: v_accvgpr_write_b32 a21, v53
+; CHECK-NEXT: v_accvgpr_write_b32 a22, v54
+; CHECK-NEXT: v_accvgpr_write_b32 a23, v55
+; CHECK-NEXT: v_accvgpr_write_b32 a24, v56
+; CHECK-NEXT: v_accvgpr_write_b32 a25, v57
+; CHECK-NEXT: v_accvgpr_write_b32 a26, v58
+; CHECK-NEXT: v_accvgpr_write_b32 a27, v59
+; CHECK-NEXT: v_accvgpr_write_b32 a28, v60
+; CHECK-NEXT: v_accvgpr_write_b32 a29, v61
+; CHECK-NEXT: v_accvgpr_write_b32 a30, v62
+; CHECK-NEXT: v_accvgpr_write_b32 a31, v63
+; CHECK-NEXT: v_accvgpr_read_b32 v32, a32
+; CHECK-NEXT: v_mov_b32_e32 v33, 0x41000000
+; CHECK-NEXT: v_and_b32_e32 v32, 0x3ff, v32
+; CHECK-NEXT: v_lshlrev_b32_e32 v32, 7, v32
+; CHECK-NEXT: s_waitcnt lgkmcnt(0)
+; CHECK-NEXT: global_store_dwordx4 v32, v[28:31], s[0:1] offset:112
+; CHECK-NEXT: global_store_dwordx4 v32, v[24:27], s[0:1] offset:96
+; CHECK-NEXT: global_store_dwordx4 v32, v[20:23], s[0:1] offset:80
+; CHECK-NEXT: global_store_dwordx4 v32, v[16:19], s[0:1] offset:64
+; CHECK-NEXT: global_store_dwordx4 v32, v[12:15], s[0:1] offset:48
+; CHECK-NEXT: global_store_dwordx4 v32, v[8:11], s[0:1] offset:32
+; CHECK-NEXT: global_store_dwordx4 v32, v[4:7], s[0:1] offset:16
+; CHECK-NEXT: global_store_dwordx4 v32, v[0:3], s[0:1]
+; CHECK-NEXT: v_mov_b32_e32 v34, 0x41800000
+; CHECK-NEXT: s_nop 0
+; CHECK-NEXT: v_accvgpr_read_b32 v0, a0
+; CHECK-NEXT: v_accvgpr_read_b32 v1, a1
+; CHECK-NEXT: v_accvgpr_read_b3...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/153020
More information about the llvm-branch-commits
mailing list