https://github.com/arsenm created
https://github.com/llvm/llvm-project/pull/153018
Currently only cases rooted at a full copy of an MFMA result are handled.
Prepare to relax that by testing more intricate subregister usage.
Currently only full copies are handled, add some tests to help work
towards handling subregisters.
>From 361949e60111ec0bdd99761dc2f59d5c3a07434c Mon Sep 17 00:00:00 2001
From: Matt Arsenault <matthew.arsena...@amd.com>
Date: Fri, 8 Aug 2025 23:56:03 +0900
Subject: [PATCH] AMDGPU: Add some baseline test for mfma rewrite with
subregister copies
Currently only cases rooted at a full copy of an MFMA result are handled.
Prepare to relax that by testing more intricate subregister usage.
Currently only full copies are handled, add some tests to help work
towards handling subregisters.
---
...gpr-mfma-to-agpr-subreg-insert-extract.mir | 218 ++++++++++++
...te-vgpr-mfma-to-agpr-subreg-src2-chain.mir | 327 ++++++++++++++++++
.../AMDGPU/rewrite-vgpr-mfma-to-agpr.ll | 69 ++++
3 files changed, 614 insertions(+)
create mode 100644
llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-subreg-insert-extract.mir
create mode 100644
llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-subreg-src2-chain.mir
diff --git
a/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-subreg-insert-extract.mir
b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-subreg-insert-extract.mir
new file mode 100644
index 0000000000000..73784fb405ce5
--- /dev/null
+++
b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-subreg-insert-extract.mir
@@ -0,0 +1,218 @@
+# 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 -o - %s | FileCheck %s
+
+# V-to-A copy is a subregister insert
+---
+name: test_rewrite_mfma_copy_subreg_insert
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+
+ ; CHECK-LABEL: name: test_rewrite_mfma_copy_subreg_insert
+ ; 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]+]]:vreg_64_align2 =
GLOBAL_LOAD_DWORDX2 [[COPY]], 0, 0, implicit $exec :: (load (s64), addrspace 1)
+ ; CHECK-NEXT: [[V_MFMA_F64_4X4X4F64_vgprcd_e64_:%[0-9]+]]:vreg_64_align2 =
V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]], [[GLOBAL_LOAD_DWORDX2_]],
0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: undef [[COPY3:%[0-9]+]].sub0_sub1:areg_128_align2 = COPY
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_]]
+ ; CHECK-NEXT: INLINEASM &"; use $0", 1 /* sideeffect attdialect */,
6291465 /* reguse:AReg_128_Align2 */, [[COPY3]]
+ ; CHECK-NEXT: GLOBAL_STORE_DWORDX4 [[COPY]], [[COPY3]], 0, 0, implicit
$exec :: (store (s128), addrspace 1)
+ ; CHECK-NEXT: GLOBAL_STORE_DWORDX2 [[COPY]], [[COPY3]].sub2_sub3, 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:vreg_64_align2 = GLOBAL_LOAD_DWORDX2 %0, 0, 0, implicit $exec :: (load
(s64), addrspace 1)
+ %4:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, %3, 0, 0, 0,
implicit $mode, implicit $exec
+ undef %5.sub0_sub1:areg_128_align2 = COPY %4
+ INLINEASM &"; use $0", 1 /* sideeffect attdialect */, 6291465 /*
reguse:AReg_128_Align2 */, %5
+ GLOBAL_STORE_DWORDX4 %0, %5, 0, 0, implicit $exec :: (store (s128),
addrspace 1)
+ GLOBAL_STORE_DWORDX2 %0, %5.sub2_sub3, 0, 0, implicit $exec :: (store
(s128), addrspace 1)
+ SI_RETURN
+...
+
+# V-to-A copy is a subregister extract
+---
+name: test_rewrite_mfma_copy_subreg_extract
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+
+ ; CHECK-LABEL: name: test_rewrite_mfma_copy_subreg_extract
+ ; 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]+]]:vreg_64_align2 =
GLOBAL_LOAD_DWORDX2 [[COPY]], 0, 0, implicit $exec :: (load (s64), addrspace 1)
+ ; CHECK-NEXT: [[V_MFMA_F64_4X4X4F64_vgprcd_e64_:%[0-9]+]]:vreg_64_align2 =
V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]], [[GLOBAL_LOAD_DWORDX2_]],
0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[COPY3:%[0-9]+]]:agpr_32 = COPY
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_]].sub0
+ ; CHECK-NEXT: GLOBAL_STORE_DWORD [[COPY]], [[COPY3]], 0, 0, implicit $exec
:: (store (s32), 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 = GLOBAL_LOAD_DWORDX2 %0, 0, 0, implicit $exec :: (load
(s64), addrspace 1)
+ %4:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, %3, 0, 0, 0,
implicit $mode, implicit $exec
+ %5:agpr_32 = COPY %4.sub0
+ GLOBAL_STORE_DWORD %0, %5, 0, 0, implicit $exec :: (store (s32), addrspace
1)
+ SI_RETURN
+...
+
+# V-to-A copy is a subregister-to-subregister copy
+---
+name: test_rewrite_mfma_copy_subreg_insert_extract_same_subreg
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+
+ ; CHECK-LABEL: name:
test_rewrite_mfma_copy_subreg_insert_extract_same_subreg
+ ; 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]+]]:vreg_64_align2 =
GLOBAL_LOAD_DWORDX2 [[COPY]], 0, 0, implicit $exec :: (load (s64), addrspace 1)
+ ; CHECK-NEXT: [[V_MFMA_F64_4X4X4F64_vgprcd_e64_:%[0-9]+]]:vreg_64_align2 =
V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]], [[GLOBAL_LOAD_DWORDX2_]],
0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: undef [[COPY3:%[0-9]+]].sub0:areg_64_align2 = COPY
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_]].sub0
+ ; CHECK-NEXT: GLOBAL_STORE_DWORDX2 [[COPY]], [[COPY3]], 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 = GLOBAL_LOAD_DWORDX2 %0, 0, 0, implicit $exec :: (load
(s64), addrspace 1)
+ %4:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, %3, 0, 0, 0,
implicit $mode, implicit $exec
+ undef %5.sub0:areg_64_align2 = COPY %4.sub0
+ GLOBAL_STORE_DWORDX2 %0, %5, 0, 0, implicit $exec :: (store (s64),
addrspace 1)
+ SI_RETURN
+...
+
+# V-to-A copy is a subregister-to-subregister copy
+---
+name: test_rewrite_mfma_copy_subreg_insert_extract_different_subreg
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+
+ ; CHECK-LABEL: name:
test_rewrite_mfma_copy_subreg_insert_extract_different_subreg
+ ; 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]+]]:vreg_64_align2 =
GLOBAL_LOAD_DWORDX2 [[COPY]], 0, 0, implicit $exec :: (load (s64), addrspace 1)
+ ; CHECK-NEXT: [[V_MFMA_F64_4X4X4F64_vgprcd_e64_:%[0-9]+]]:vreg_64_align2 =
V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]], [[GLOBAL_LOAD_DWORDX2_]],
0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: undef [[COPY3:%[0-9]+]].sub0:areg_64_align2 = COPY
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_]].sub1
+ ; CHECK-NEXT: GLOBAL_STORE_DWORDX2 [[COPY]], [[COPY3]], 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 = GLOBAL_LOAD_DWORDX2 %0, 0, 0, implicit $exec :: (load
(s64), addrspace 1)
+ %4:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, %3, 0, 0, 0,
implicit $mode, implicit $exec
+ undef %5.sub0:areg_64_align2 = COPY %4.sub1
+ GLOBAL_STORE_DWORDX2 %0, %5, 0, 0, implicit $exec :: (store (s64),
addrspace 1)
+ SI_RETURN
+...
+
+# V-to-A copy is a subregister extract, from a subregister def
+---
+name: test_rewrite_mfma_copy_subreg_extract_from_subreg_def
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+
+ ; CHECK-LABEL: name: test_rewrite_mfma_copy_subreg_extract_from_subreg_def
+ ; 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]+]]:vreg_64_align2 =
GLOBAL_LOAD_DWORDX2 [[COPY]], 0, 0, implicit $exec :: (load (s64), addrspace 1)
+ ; CHECK-NEXT: undef
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_:%[0-9]+]].sub2_sub3:vreg_128_align2 =
V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]], [[GLOBAL_LOAD_DWORDX2_]],
0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[COPY3:%[0-9]+]]:agpr_32 = COPY
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_]].sub0
+ ; CHECK-NEXT: GLOBAL_STORE_DWORD [[COPY]], [[COPY3]], 0, 0, implicit $exec
:: (store (s32), 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 = GLOBAL_LOAD_DWORDX2 %0, 0, 0, implicit $exec :: (load
(s64), addrspace 1)
+ undef %4.sub2_sub3:vreg_128_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1,
%2, %3, 0, 0, 0, implicit $mode, implicit $exec
+ %5:agpr_32 = COPY %4.sub0
+ GLOBAL_STORE_DWORD %0, %5, 0, 0, implicit $exec :: (store (s32), addrspace
1)
+ SI_RETURN
+...
+
+# V-to-A copy is a subregister insert from a subregister_def
+---
+name: test_rewrite_mfma_copy_subreg_insert_from_subreg_def_tuple
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+
+ ; CHECK-LABEL: name:
test_rewrite_mfma_copy_subreg_insert_from_subreg_def_tuple
+ ; 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]+]]:vreg_64_align2 =
GLOBAL_LOAD_DWORDX2 [[COPY]], 0, 0, implicit $exec :: (load (s64), addrspace 1)
+ ; CHECK-NEXT: undef
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_:%[0-9]+]].sub2_sub3:vreg_128_align2 =
V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]], [[GLOBAL_LOAD_DWORDX2_]],
0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: undef [[COPY3:%[0-9]+]].sub0_sub1:areg_128_align2 = COPY
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_]].sub2_sub3
+ ; CHECK-NEXT: INLINEASM &"; use $0", 1 /* sideeffect attdialect */,
6291465 /* reguse:AReg_128_Align2 */, [[COPY3]]
+ ; CHECK-NEXT: GLOBAL_STORE_DWORDX4 [[COPY]], [[COPY3]], 0, 0, implicit
$exec :: (store (s128), addrspace 1)
+ ; CHECK-NEXT: GLOBAL_STORE_DWORDX2 [[COPY]], [[COPY3]].sub2_sub3, 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:vreg_64_align2 = GLOBAL_LOAD_DWORDX2 %0, 0, 0, implicit $exec :: (load
(s64), addrspace 1)
+ undef %4.sub2_sub3:vreg_128_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1,
%2, %3, 0, 0, 0, implicit $mode, implicit $exec
+ undef %5.sub0_sub1:areg_128_align2 = COPY %4.sub2_sub3
+ INLINEASM &"; use $0", 1 /* sideeffect attdialect */, 6291465 /*
reguse:AReg_128_Align2 */, %5
+ GLOBAL_STORE_DWORDX4 %0, %5, 0, 0, implicit $exec :: (store (s128),
addrspace 1)
+ GLOBAL_STORE_DWORDX2 %0, %5.sub2_sub3, 0, 0, implicit $exec :: (store
(s128), addrspace 1)
+ SI_RETURN
+...
+
+# V-to-A copy is a subregister insert of a subregister from a
+# subregister_def
+---
+name: test_rewrite_mfma_copy_subreg_insert_from_subreg_def_subreg
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+
+ ; CHECK-LABEL: name:
test_rewrite_mfma_copy_subreg_insert_from_subreg_def_subreg
+ ; 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]+]]:vreg_64_align2 =
GLOBAL_LOAD_DWORDX2 [[COPY]], 0, 0, implicit $exec :: (load (s64), addrspace 1)
+ ; CHECK-NEXT: undef
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_:%[0-9]+]].sub2_sub3:vreg_128_align2 =
V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]], [[GLOBAL_LOAD_DWORDX2_]],
0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: undef [[COPY3:%[0-9]+]].sub1:areg_128_align2 = COPY
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_]].sub2
+ ; CHECK-NEXT: INLINEASM &"; use $0", 1 /* sideeffect attdialect */,
6291465 /* reguse:AReg_128_Align2 */, [[COPY3]]
+ ; CHECK-NEXT: GLOBAL_STORE_DWORDX4 [[COPY]], [[COPY3]], 0, 0, implicit
$exec :: (store (s128), addrspace 1)
+ ; CHECK-NEXT: GLOBAL_STORE_DWORDX2 [[COPY]], [[COPY3]].sub2_sub3, 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:vreg_64_align2 = GLOBAL_LOAD_DWORDX2 %0, 0, 0, implicit $exec :: (load
(s64), addrspace 1)
+ undef %4.sub2_sub3:vreg_128_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1,
%2, %3, 0, 0, 0, implicit $mode, implicit $exec
+ undef %5.sub1:areg_128_align2 = COPY %4.sub2
+ INLINEASM &"; use $0", 1 /* sideeffect attdialect */, 6291465 /*
reguse:AReg_128_Align2 */, %5
+ GLOBAL_STORE_DWORDX4 %0, %5, 0, 0, implicit $exec :: (store (s128),
addrspace 1)
+ GLOBAL_STORE_DWORDX2 %0, %5.sub2_sub3, 0, 0, implicit $exec :: (store
(s128), addrspace 1)
+ SI_RETURN
+...
diff --git
a/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-subreg-src2-chain.mir
b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-subreg-src2-chain.mir
new file mode 100644
index 0000000000000..38b13052246c1
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-subreg-src2-chain.mir
@@ -0,0 +1,327 @@
+# 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 -o - %s | FileCheck %s
+
+---
+name: test_rewrite_mfma_src2_is_subreg_0
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+
+ ; CHECK-LABEL: name: test_rewrite_mfma_src2_is_subreg_0
+ ; 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]+]]:vreg_128_align2 =
GLOBAL_LOAD_DWORDX4 [[COPY]], 0, 0, implicit $exec :: (load (s128), addrspace 1)
+ ; CHECK-NEXT: [[V_MFMA_F64_4X4X4F64_vgprcd_e64_:%[0-9]+]]:vreg_64_align2 =
V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]],
[[GLOBAL_LOAD_DWORDX4_]].sub0_sub1, 0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: undef [[COPY3:%[0-9]+]].sub0_sub1:areg_128_align2 = COPY
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_]]
+ ; CHECK-NEXT: INLINEASM &"; use $0", 1 /* sideeffect attdialect */,
6291465 /* reguse:AReg_128_Align2 */, [[COPY3]]
+ ; 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:vreg_128_align2 = GLOBAL_LOAD_DWORDX4 %0, 0, 0, implicit $exec :: (load
(s128), addrspace 1)
+ %4:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, %3.sub0_sub1,
0, 0, 0, implicit $mode, implicit $exec
+ undef %5.sub0_sub1:areg_128_align2 = COPY %4
+ INLINEASM &"; use $0", 1 /* sideeffect attdialect */, 6291465 /*
reguse:AReg_128_Align2 */, %5
+ GLOBAL_STORE_DWORDX4 %0, %5, 0, 0, implicit $exec :: (store (s128),
addrspace 1)
+ SI_RETURN
+...
+
+---
+name: test_rewrite_mfma_src2_is_subreg_1
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+
+ ; CHECK-LABEL: name: test_rewrite_mfma_src2_is_subreg_1
+ ; 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]+]]:vreg_128_align2 =
GLOBAL_LOAD_DWORDX4 [[COPY]], 0, 0, implicit $exec :: (load (s128), addrspace 1)
+ ; CHECK-NEXT: [[V_MFMA_F64_4X4X4F64_vgprcd_e64_:%[0-9]+]]:vreg_64_align2 =
V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]],
[[GLOBAL_LOAD_DWORDX4_]].sub2_sub3, 0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: undef [[COPY3:%[0-9]+]].sub0_sub1:areg_128_align2 = COPY
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_]]
+ ; CHECK-NEXT: INLINEASM &"; use $0", 1 /* sideeffect attdialect */,
6291465 /* reguse:AReg_128_Align2 */, [[COPY3]]
+ ; 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:vreg_128_align2 = GLOBAL_LOAD_DWORDX4 %0, 0, 0, implicit $exec :: (load
(s128), addrspace 1)
+ %4:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, %3.sub2_sub3,
0, 0, 0, implicit $mode, implicit $exec
+ undef %5.sub0_sub1:areg_128_align2 = COPY %4
+ INLINEASM &"; use $0", 1 /* sideeffect attdialect */, 6291465 /*
reguse:AReg_128_Align2 */, %5
+ GLOBAL_STORE_DWORDX4 %0, %5, 0, 0, implicit $exec :: (store (s128),
addrspace 1)
+ SI_RETURN
+...
+
+---
+name: test_rewrite_mfma_src2_is_subreg_chain_mfma_full_def
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+
+ ; CHECK-LABEL: name: test_rewrite_mfma_src2_is_subreg_chain_mfma_full_def
+ ; 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]+]]:vreg_64_align2 =
GLOBAL_LOAD_DWORDX2 [[COPY]], 0, 0, implicit $exec :: (load (s64), addrspace 1)
+ ; CHECK-NEXT: undef
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_:%[0-9]+]].sub0_sub1:vreg_128_align2 =
V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]], [[GLOBAL_LOAD_DWORDX2_]],
0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: dead %other_use:vreg_64_align2 = COPY
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_]].sub0_sub1
+ ; CHECK-NEXT: [[V_MFMA_F64_4X4X4F64_vgprcd_e64_1:%[0-9]+]]:vreg_64_align2
= V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]],
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_]].sub0_sub1, 0, 0, 0, implicit $mode,
implicit $exec
+ ; CHECK-NEXT: [[COPY3:%[0-9]+]]:areg_64_align2 = COPY
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_1]]
+ ; CHECK-NEXT: INLINEASM &"; use $0", 1 /* sideeffect attdialect */,
3735561 /* reguse:AReg_64_Align2 */, [[COPY3]]
+ ; CHECK-NEXT: GLOBAL_STORE_DWORDX2 [[COPY]], [[COPY3]], 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 = GLOBAL_LOAD_DWORDX2 %0, 0, 0, implicit $exec :: (load
(s64), addrspace 1)
+ undef %4.sub0_sub1:vreg_128_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1,
%2, %3, 0, 0, 0, implicit $mode, implicit $exec
+ %other_use:vreg_64_align2 = COPY %4.sub0_sub1
+ %5:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, %4.sub0_sub1,
0, 0, 0, implicit $mode, implicit $exec
+ %6:areg_64_align2 = COPY %5
+ INLINEASM &"; use $0", 1 /* sideeffect attdialect */, 3735561 /*
reguse:AReg_64_Align2 */, %6:areg_64_align2
+ GLOBAL_STORE_DWORDX2 %0, %6, 0, 0, implicit $exec :: (store (s64),
addrspace 1)
+ SI_RETURN
+...
+
+---
+name: test_rewrite_mfma_src2_is_subreg_chain_mfma_sub_def
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+
+ ; CHECK-LABEL: name: test_rewrite_mfma_src2_is_subreg_chain_mfma_sub_def
+ ; 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: [[V_MFMA_F64_4X4X4F64_e64_:%[0-9]+]]:areg_64_align2 =
V_MFMA_F64_4X4X4F64_e64 [[COPY1]], [[COPY2]], [[GLOBAL_LOAD_DWORDX2_]], 0, 0,
0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: dead %other_use:vreg_64_align2 = COPY
[[V_MFMA_F64_4X4X4F64_e64_]]
+ ; CHECK-NEXT: undef
[[V_MFMA_F64_4X4X4F64_e64_1:%[0-9]+]].sub0_sub1:areg_128_align2 =
V_MFMA_F64_4X4X4F64_e64 [[COPY1]], [[COPY2]], [[V_MFMA_F64_4X4X4F64_e64_]], 0,
0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[V_MFMA_F64_4X4X4F64_e64_2:%[0-9]+]]:areg_64_align2 =
V_MFMA_F64_4X4X4F64_e64 [[COPY1]], [[COPY2]],
[[V_MFMA_F64_4X4X4F64_e64_1]].sub0_sub1, 0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: [[COPY3:%[0-9]+]]:areg_64_align2 = COPY
[[V_MFMA_F64_4X4X4F64_e64_2]]
+ ; CHECK-NEXT: INLINEASM &"; use $0", 1 /* sideeffect attdialect */,
3735561 /* reguse:AReg_64_Align2 */, [[COPY3]]
+ ; CHECK-NEXT: GLOBAL_STORE_DWORDX2 [[COPY]], [[COPY3]], 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 = GLOBAL_LOAD_DWORDX2 %0, 0, 0, implicit $exec :: (load
(s64), addrspace 1)
+ %4:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, %3, 0, 0, 0,
implicit $mode, implicit $exec
+ %other_use:vreg_64_align2 = COPY %4
+ undef %5.sub0_sub1:vreg_128_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1,
%2, %4, 0, 0, 0, implicit $mode, implicit $exec
+ %6:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, %5.sub0_sub1,
0, 0, 0, implicit $mode, implicit $exec
+ %7:areg_64_align2 = COPY %6
+ INLINEASM &"; use $0", 1 /* sideeffect attdialect */, 3735561 /*
reguse:AReg_64_Align2 */, %7
+ GLOBAL_STORE_DWORDX2 %0, %7, 0, 0, implicit $exec :: (store (s64),
addrspace 1)
+ SI_RETURN
+
+...
+
+---
+name: test_rewrite_mfma_src2_is_subreg_chain_mfma_sub_def_copy_subreg_insert
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+
+ ; CHECK-LABEL: name:
test_rewrite_mfma_src2_is_subreg_chain_mfma_sub_def_copy_subreg_insert
+ ; 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]+]]:vreg_64_align2 =
GLOBAL_LOAD_DWORDX2 [[COPY]], 0, 0, implicit $exec :: (load (s64), addrspace 1)
+ ; CHECK-NEXT: [[V_MFMA_F64_4X4X4F64_vgprcd_e64_:%[0-9]+]]:vreg_64_align2 =
V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]], [[GLOBAL_LOAD_DWORDX2_]],
0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: undef
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_1:%[0-9]+]].sub0_sub1:vreg_128_align2 =
V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]],
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_]], 0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: dead %other_use:vreg_64_align2 = COPY
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_1]].sub0_sub1
+ ; CHECK-NEXT: [[V_MFMA_F64_4X4X4F64_vgprcd_e64_2:%[0-9]+]]:vreg_64_align2
= V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]],
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_1]].sub0_sub1, 0, 0, 0, implicit $mode,
implicit $exec
+ ; CHECK-NEXT: undef [[COPY3:%[0-9]+]].sub0_sub1:areg_128_align2 = COPY
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_2]]
+ ; CHECK-NEXT: INLINEASM &"; use $0", 1 /* sideeffect attdialect */,
6291465 /* reguse:AReg_128_Align2 */, [[COPY3]]
+ ; 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:vreg_64_align2 = GLOBAL_LOAD_DWORDX2 %0, 0, 0, implicit $exec :: (load
(s64), addrspace 1)
+ %4:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, %3, 0, 0, 0,
implicit $mode, implicit $exec
+ undef %5.sub0_sub1:vreg_128_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1,
%2, %4, 0, 0, 0, implicit $mode, implicit $exec
+ %other_use:vreg_64_align2 = COPY %5.sub0_sub1
+ %6:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, %5.sub0_sub1,
0, 0, 0, implicit $mode, implicit $exec
+ undef %8.sub0_sub1:areg_128_align2 = COPY %6
+ INLINEASM &"; use $0", 1 /* sideeffect attdialect */, 6291465 /*
reguse:AReg_128_Align2 */, %8:areg_128_align2
+ GLOBAL_STORE_DWORDX4 %0, %8, 0, 0, implicit $exec :: (store (s128),
addrspace 1)
+ SI_RETURN
+
+...
+
+---
+name: test_rewrite_mfma_src2_is_subreg_chain_mfma_sub_def_copy_subreg_extract
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+
+ ; CHECK-LABEL: name:
test_rewrite_mfma_src2_is_subreg_chain_mfma_sub_def_copy_subreg_extract
+ ; 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]+]]:vreg_64_align2 =
GLOBAL_LOAD_DWORDX2 [[COPY]], 0, 0, implicit $exec :: (load (s64), addrspace 1)
+ ; CHECK-NEXT: [[V_MFMA_F64_4X4X4F64_vgprcd_e64_:%[0-9]+]]:vreg_64_align2 =
V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]], [[GLOBAL_LOAD_DWORDX2_]],
0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: dead %other_use0:vreg_64_align2 = COPY
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_]].sub0
+ ; CHECK-NEXT: undef
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_1:%[0-9]+]].sub0_sub1:vreg_128_align2 =
V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]],
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_]], 0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: dead %other_use1:vreg_64_align2 = COPY
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_1]].sub0_sub1
+ ; CHECK-NEXT: [[V_MFMA_F64_4X4X4F64_vgprcd_e64_2:%[0-9]+]]:vreg_64_align2
= V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]],
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_1]].sub0_sub1, 0, 0, 0, implicit $mode,
implicit $exec
+ ; CHECK-NEXT: [[COPY3:%[0-9]+]]:agpr_32 = COPY
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_2]]
+ ; CHECK-NEXT: INLINEASM &"; use $0", 1 /* sideeffect attdialect */,
2162697 /* reguse:AGPR_32 */, [[COPY3]]
+ ; CHECK-NEXT: GLOBAL_STORE_DWORD [[COPY]], [[COPY3]], 0, 0, implicit $exec
:: (store (s32), 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 = GLOBAL_LOAD_DWORDX2 %0, 0, 0, implicit $exec :: (load
(s64), addrspace 1)
+ %4:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, %3, 0, 0, 0,
implicit $mode, implicit $exec
+ %other_use0:vreg_64_align2 = COPY %4.sub0
+ undef %5.sub0_sub1:vreg_128_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1,
%2, %4, 0, 0, 0, implicit $mode, implicit $exec
+ %other_use1:vreg_64_align2 = COPY %5.sub0_sub1
+ %6:vreg_64_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2, %5.sub0_sub1,
0, 0, 0, implicit $mode, implicit $exec
+ %8:agpr_32 = COPY %6
+ INLINEASM &"; use $0", 1 /* sideeffect attdialect */, 2162697 /*
reguse:AGPR_32 */, %8:agpr_32
+ GLOBAL_STORE_DWORD %0, %8, 0, 0, implicit $exec :: (store (s32), addrspace
1)
+ SI_RETURN
+
+...
+
+# Multiple MFMAs in the chain use the same register with different
+# subregisters.
+---
+name: test_rewrite_mfma_src2_chain_different_subregs_same_reg_full_copy_use
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+
+ ; CHECK-LABEL: name:
test_rewrite_mfma_src2_chain_different_subregs_same_reg_full_copy_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: undef
[[V_MFMA_F64_4X4X4F64_e64_:%[0-9]+]].sub2_sub3:areg_128_align2 =
V_MFMA_F64_4X4X4F64_e64 [[COPY1]], [[COPY2]],
[[GLOBAL_LOAD_DWORDX4_]].sub0_sub1, 0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: dead %other_use0:vreg_64_align2 = COPY
[[V_MFMA_F64_4X4X4F64_e64_]].sub0_sub1
+ ; CHECK-NEXT:
[[V_MFMA_F64_4X4X4F64_e64_:%[0-9]+]].sub0_sub1:areg_128_align2 =
V_MFMA_F64_4X4X4F64_e64 [[COPY1]], [[COPY2]],
[[GLOBAL_LOAD_DWORDX4_]].sub2_sub3, 0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: dead %other_use1:vreg_64_align2 = COPY
[[V_MFMA_F64_4X4X4F64_e64_]].sub2_sub3
+ ; CHECK-NEXT: dead %other_use2:vreg_64 = COPY
[[V_MFMA_F64_4X4X4F64_e64_]].sub1_sub2
+ ; CHECK-NEXT: [[COPY3:%[0-9]+]]:areg_128_align2 = COPY
[[V_MFMA_F64_4X4X4F64_e64_]]
+ ; CHECK-NEXT: INLINEASM &"; use $0", 1 /* sideeffect attdialect */,
6291465 /* reguse:AReg_128_Align2 */, [[COPY3]]
+ ; 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:vreg_128_align2 = GLOBAL_LOAD_DWORDX4 %0, 0, 0, implicit $exec :: (load
(s128), addrspace 1)
+ undef %4.sub2_sub3:vreg_128_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1,
%2, %3.sub0_sub1, 0, 0, 0, implicit $mode, implicit $exec
+ %other_use0:vreg_64_align2 = COPY %4.sub0_sub1
+ %4.sub0_sub1:vreg_128_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2,
%3.sub2_sub3, 0, 0, 0, implicit $mode, implicit $exec
+ %other_use1:vreg_64_align2 = COPY %4.sub2_sub3
+ %other_use2:vreg_64 = COPY %4.sub1_sub2
+ %6:areg_128_align2 = COPY %4
+ INLINEASM &"; use $0", 1 /* sideeffect attdialect */, 6291465 /*
reguse:AReg_128_Align2 */, %6:areg_128_align2
+ GLOBAL_STORE_DWORDX4 %0, %6, 0, 0, implicit $exec :: (store (s128),
addrspace 1)
+ SI_RETURN
+...
+
+# Multiple MFMAs in the chain use the same register with different
+# subregisters.
+---
+name: test_rewrite_mfma_src2_chain_different_subregs_same_reg_extract_use_0
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+
+ ; CHECK-LABEL: name:
test_rewrite_mfma_src2_chain_different_subregs_same_reg_extract_use_0
+ ; 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]+]]:vreg_128_align2 =
GLOBAL_LOAD_DWORDX4 [[COPY]], 0, 0, implicit $exec :: (load (s128), addrspace 1)
+ ; CHECK-NEXT: undef
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_:%[0-9]+]].sub2_sub3:vreg_128_align2 =
V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]],
[[GLOBAL_LOAD_DWORDX4_]].sub0_sub1, 0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: dead %other_use0:vreg_64_align2 = COPY
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_]].sub0_sub1
+ ; CHECK-NEXT:
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_:%[0-9]+]].sub0_sub1:vreg_128_align2 =
V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]],
[[GLOBAL_LOAD_DWORDX4_]].sub2_sub3, 0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: %other_use1:vreg_64_align2 = COPY
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_]].sub2_sub3
+ ; CHECK-NEXT: dead %other_use2:vreg_64 = COPY
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_]].sub1_sub2
+ ; CHECK-NEXT: [[COPY3:%[0-9]+]]:areg_64 = COPY
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_]].sub1_sub2
+ ; CHECK-NEXT: INLINEASM &"; use $0", 1 /* sideeffect attdialect */,
3473417 /* reguse:AReg_64 */, [[COPY3]]
+ ; CHECK-NEXT: GLOBAL_STORE_DWORDX2 [[COPY]], %other_use1, 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_128_align2 = GLOBAL_LOAD_DWORDX4 %0, 0, 0, implicit $exec :: (load
(s128), addrspace 1)
+ undef %4.sub2_sub3:vreg_128_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1,
%2, %3.sub0_sub1, 0, 0, 0, implicit $mode, implicit $exec
+ %other_use0:vreg_64_align2 = COPY %4.sub0_sub1
+ %4.sub0_sub1:vreg_128_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2,
%3.sub2_sub3, 0, 0, 0, implicit $mode, implicit $exec
+ %other_use1:vreg_64_align2 = COPY %4.sub2_sub3
+ %other_use2:vreg_64 = COPY %4.sub1_sub2
+ %6:areg_64 = COPY %4.sub1_sub2
+ INLINEASM &"; use $0", 1 /* sideeffect attdialect */, 3473417 /*
reguse:AReg_64 */, %6:areg_64
+ GLOBAL_STORE_DWORDX2 %0, %other_use1, 0, 0, implicit $exec :: (store
(s64), addrspace 1)
+ SI_RETURN
+...
+
+---
+name: test_rewrite_mfma_src2_chain_different_subregs_same_reg_extract_use_1
+tracksRegLiveness: true
+body: |
+ bb.0:
+ liveins: $vgpr0, $vgpr1, $vgpr2, $vgpr3, $vgpr4, $vgpr5
+
+ ; CHECK-LABEL: name:
test_rewrite_mfma_src2_chain_different_subregs_same_reg_extract_use_1
+ ; 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]+]]:vreg_128_align2 =
GLOBAL_LOAD_DWORDX4 [[COPY]], 0, 0, implicit $exec :: (load (s128), addrspace 1)
+ ; CHECK-NEXT: undef
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_:%[0-9]+]].sub2_sub3:vreg_128_align2 =
V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]],
[[GLOBAL_LOAD_DWORDX4_]].sub0_sub1, 0, 0, 0, implicit $mode, implicit $exec
+ ; CHECK-NEXT: dead %other_use0:vreg_64_align2 = COPY
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_]].sub0_sub1
+ ; CHECK-NEXT: dead
[[GLOBAL_LOAD_DWORDX4_:%[0-9]+]].sub2_sub3:vreg_128_align2 =
V_MFMA_F64_4X4X4F64_vgprcd_e64 [[COPY1]], [[COPY2]],
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_]].sub2_sub3, 0, 0, 0, implicit $mode,
implicit $exec
+ ; CHECK-NEXT: %other_use1:vreg_64_align2 = COPY
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_]].sub2_sub3
+ ; CHECK-NEXT: dead %other_use2:vreg_64 = COPY
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_]].sub1_sub2
+ ; CHECK-NEXT: [[COPY3:%[0-9]+]]:areg_64 = COPY
[[V_MFMA_F64_4X4X4F64_vgprcd_e64_]].sub1_sub2
+ ; CHECK-NEXT: INLINEASM &"; use $0", 1 /* sideeffect attdialect */,
3473417 /* reguse:AReg_64 */, [[COPY3]]
+ ; CHECK-NEXT: GLOBAL_STORE_DWORDX2 [[COPY]], %other_use1, 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_128_align2 = GLOBAL_LOAD_DWORDX4 %0, 0, 0, implicit $exec :: (load
(s128), addrspace 1)
+ undef %4.sub2_sub3:vreg_128_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1,
%2, %3.sub0_sub1, 0, 0, 0, implicit $mode, implicit $exec
+ %other_use0:vreg_64_align2 = COPY %4.sub0_sub1
+ %3.sub2_sub3:vreg_128_align2 = V_MFMA_F64_4X4X4F64_vgprcd_e64 %1, %2,
%4.sub2_sub3, 0, 0, 0, implicit $mode, implicit $exec
+ %other_use1:vreg_64_align2 = COPY %4.sub2_sub3
+ %other_use2:vreg_64 = COPY %4.sub1_sub2
+ %6:areg_64 = COPY %4.sub1_sub2
+ INLINEASM &"; use $0", 1 /* sideeffect attdialect */, 3473417 /*
reguse:AReg_64 */, %6:areg_64
+ GLOBAL_STORE_DWORDX2 %0, %other_use1, 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 4bc0270e61a7a..701a83421524d 100644
--- a/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr.ll
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr.ll
@@ -479,6 +479,75 @@ entry:
ret void
}
+define void @test_rewrite_mfma_subreg_insert0(float %arg0, float %arg1, ptr
addrspace(1) %ptr) #0 {
+; CHECK-LABEL: test_rewrite_mfma_subreg_insert0:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: global_load_dwordx4 v[2:5], v[2:3], off
+; CHECK-NEXT: s_waitcnt vmcnt(0)
+; CHECK-NEXT: v_mfma_f32_4x4x1_16b_f32 v[0:3], v0, v1, v[2:5]
+; CHECK-NEXT: s_nop 3
+; CHECK-NEXT: v_accvgpr_write_b32 a0, v0
+; CHECK-NEXT: v_accvgpr_write_b32 a1, v1
+; CHECK-NEXT: v_accvgpr_write_b32 a2, v2
+; CHECK-NEXT: v_accvgpr_write_b32 a3, v3
+; CHECK-NEXT: ;;#ASMSTART
+; CHECK-NEXT: ; use a[0:7]
+; CHECK-NEXT: ;;#ASMEND
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %src2 = load <4 x float>, ptr addrspace(1) %ptr
+ %mai = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float %arg0, float
%arg1, <4 x float> %src2, i32 0, i32 0, i32 0)
+ %insert.sub0 = shufflevector <4 x float> %mai, <4 x float> poison, <8 x i32>
<i32 0, i32 1, i32 2, i32 3, i32 poison, i32 poison, i32 poison, i32 poison>
+ call void asm sideeffect "; use $0", "a"(<8 x float> %insert.sub0)
+ ret void
+}
+
+; odd offset
+define void @test_rewrite_mfma_subreg_insert1(float %arg0, float %arg1, ptr
addrspace(1) %ptr) #0 {
+; CHECK-LABEL: test_rewrite_mfma_subreg_insert1:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: global_load_dwordx4 v[2:5], v[2:3], off
+; CHECK-NEXT: s_waitcnt vmcnt(0)
+; CHECK-NEXT: v_mfma_f32_4x4x1_16b_f32 v[0:3], v0, v1, v[2:5]
+; CHECK-NEXT: s_nop 3
+; CHECK-NEXT: v_pk_mov_b32 v[0:1], v[0:1], v[2:3] op_sel:[1,0]
+; CHECK-NEXT: s_nop 0
+; CHECK-NEXT: v_accvgpr_write_b32 a0, v0
+; CHECK-NEXT: v_accvgpr_write_b32 a1, v1
+; CHECK-NEXT: v_accvgpr_write_b32 a2, v3
+; CHECK-NEXT: ;;#ASMSTART
+; CHECK-NEXT: ; use a[0:7]
+; CHECK-NEXT: ;;#ASMEND
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %src2 = load <4 x float>, ptr addrspace(1) %ptr
+ %mai = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float %arg0, float
%arg1, <4 x float> %src2, i32 0, i32 0, i32 0)
+ %insert.sub0 = shufflevector <4 x float> %mai, <4 x float> poison, <8 x i32>
<i32 1, i32 2, i32 3, i32 4, i32 poison, i32 poison, i32 poison, i32 poison>
+ call void asm sideeffect "; use $0", "a"(<8 x float> %insert.sub0)
+ ret void
+}
+
+define void @test_rewrite_mfma_subreg_insert2(double %arg0, double %arg1, ptr
addrspace(1) %ptr) #0 {
+; CHECK-LABEL: test_rewrite_mfma_subreg_insert2:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; CHECK-NEXT: global_load_dwordx2 v[4:5], v[4:5], off
+; CHECK-NEXT: s_waitcnt vmcnt(0)
+; CHECK-NEXT: v_mfma_f64_4x4x4_4b_f64 v[0:1], v[0:1], v[2:3], v[4:5]
+; CHECK-NEXT: s_nop 5
+; CHECK-NEXT: v_accvgpr_write_b32 a0, v0
+; CHECK-NEXT: v_accvgpr_write_b32 a1, v1
+; CHECK-NEXT: ;;#ASMSTART
+; CHECK-NEXT: ; use a[0:3]
+; CHECK-NEXT: ;;#ASMEND
+; CHECK-NEXT: s_setpc_b64 s[30:31]
+ %src2 = load double, ptr addrspace(1) %ptr
+ %mai = call double @llvm.amdgcn.mfma.f64.4x4x4f64(double %arg0, double
%arg1, double %src2, i32 0, i32 0, i32 0)
+ %insert.sub0 = insertelement <2 x double> poison, double %mai, i32 0
+ call void asm sideeffect "; use $0", "a"(<2 x double> %insert.sub0)
+ ret void
+}
+
declare <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half>, <4 x half>,
<4 x float>, i32 immarg, i32 immarg, i32 immarg) #2
declare <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float, float, <32 x
float>, i32 immarg, i32 immarg, i32 immarg) #2
declare noundef range(i32 0, 1024) i32 @llvm.amdgcn.workitem.id.x() #3
_______________________________________________
llvm-branch-commits mailing list
llvm-branch-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
