[llvm] 1fc2b37 - [AMDGPU][GISel] RegBankLegalize rules for BF16 variants of gfx950 MFMA intrinsics (#214042)
via llvm-commits
llvm-commits at lists.llvm.org
Tue Aug 4 15:48:34 PDT 2026
Author: Chinmay Deshpande
Date: 2026-08-04T15:48:30-07:00
New Revision: 1fc2b37c6892b68962b7a916f9648d5204549085
URL: https://github.com/llvm/llvm-project/commit/1fc2b37c6892b68962b7a916f9648d5204549085
DIFF: https://github.com/llvm/llvm-project/commit/1fc2b37c6892b68962b7a916f9648d5204549085.diff
LOG: [AMDGPU][GISel] RegBankLegalize rules for BF16 variants of gfx950 MFMA intrinsics (#214042)
Added:
Modified:
llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.gfx950.bf16.ll
llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.gfx950.ll
llvm/test/CodeGen/AMDGPU/llvm.amdgcn.smfmac.gfx950.ll
Removed:
################################################################################
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
index dfb80f53eb70d..4bdbce7c2aa8e 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
@@ -2691,35 +2691,34 @@ RegBankLegalizeRules::RegBankLegalizeRules(const GCNSubtarget &_ST,
HasGFX90AInsts);
// gfx90a+ only MFMAs
- addRulesForIOpcs(
- {
- amdgcn_mfma_f32_32x32x4bf16_1k,
- amdgcn_mfma_f32_16x16x4bf16_1k,
- amdgcn_mfma_f32_4x4x4bf16_1k,
- amdgcn_mfma_f32_32x32x8bf16_1k,
- amdgcn_mfma_f32_16x16x16bf16_1k,
- amdgcn_mfma_f64_16x16x4f64,
- amdgcn_mfma_f64_4x4x4f64,
- amdgcn_mfma_i32_16x16x32_i8,
- amdgcn_mfma_i32_32x32x16_i8,
- amdgcn_mfma_f32_16x16x8_xf32,
- amdgcn_mfma_f32_32x32x4_xf32,
- amdgcn_mfma_f32_16x16x32_bf8_bf8,
- amdgcn_mfma_f32_16x16x32_bf8_fp8,
- amdgcn_mfma_f32_16x16x32_fp8_bf8,
- amdgcn_mfma_f32_16x16x32_fp8_fp8,
- amdgcn_mfma_f32_32x32x16_bf8_bf8,
- amdgcn_mfma_f32_32x32x16_bf8_fp8,
- amdgcn_mfma_f32_32x32x16_fp8_bf8,
- amdgcn_mfma_f32_32x32x16_fp8_fp8,
- // gfx950
- amdgcn_mfma_f32_16x16x32_f16,
- amdgcn_mfma_f32_32x32x16_f16,
- amdgcn_mfma_i32_16x16x64_i8,
- amdgcn_mfma_i32_32x32x32_i8,
- // TODO: bf16 variants fail in IRTranslator.
- // amdgcn_mfma_f32_16x16x32_bf16, amdgcn_mfma_f32_32x32x16_bf16,
- })
+ addRulesForIOpcs({
+ amdgcn_mfma_f32_32x32x4bf16_1k,
+ amdgcn_mfma_f32_16x16x4bf16_1k,
+ amdgcn_mfma_f32_4x4x4bf16_1k,
+ amdgcn_mfma_f32_32x32x8bf16_1k,
+ amdgcn_mfma_f32_16x16x16bf16_1k,
+ amdgcn_mfma_f64_16x16x4f64,
+ amdgcn_mfma_f64_4x4x4f64,
+ amdgcn_mfma_i32_16x16x32_i8,
+ amdgcn_mfma_i32_32x32x16_i8,
+ amdgcn_mfma_f32_16x16x8_xf32,
+ amdgcn_mfma_f32_32x32x4_xf32,
+ amdgcn_mfma_f32_16x16x32_bf8_bf8,
+ amdgcn_mfma_f32_16x16x32_bf8_fp8,
+ amdgcn_mfma_f32_16x16x32_fp8_bf8,
+ amdgcn_mfma_f32_16x16x32_fp8_fp8,
+ amdgcn_mfma_f32_32x32x16_bf8_bf8,
+ amdgcn_mfma_f32_32x32x16_bf8_fp8,
+ amdgcn_mfma_f32_32x32x16_fp8_bf8,
+ amdgcn_mfma_f32_32x32x16_fp8_fp8,
+ // gfx950
+ amdgcn_mfma_f32_16x16x32_f16,
+ amdgcn_mfma_f32_32x32x16_f16,
+ amdgcn_mfma_i32_16x16x64_i8,
+ amdgcn_mfma_i32_32x32x32_i8,
+ amdgcn_mfma_f32_16x16x32_bf16,
+ amdgcn_mfma_f32_32x32x16_bf16,
+ })
.Any({{DivAnyTy},
{{VgprOrAgprAnyTy},
{IntrId, VgprAnyTy, VgprAnyTy, VgprOrAgprAnyTy}}});
@@ -2735,6 +2734,7 @@ RegBankLegalizeRules::RegBankLegalizeRules(const GCNSubtarget &_ST,
amdgcn_smfmac_f32_32x32x32_fp8_bf8, amdgcn_smfmac_f32_32x32x32_fp8_fp8,
// gfx950+
amdgcn_smfmac_f32_16x16x64_f16, amdgcn_smfmac_f32_32x32x32_f16,
+ amdgcn_smfmac_f32_16x16x64_bf16, amdgcn_smfmac_f32_32x32x32_bf16,
amdgcn_smfmac_i32_16x16x128_i8, amdgcn_smfmac_i32_32x32x64_i8,
amdgcn_smfmac_f32_16x16x128_bf8_bf8, amdgcn_smfmac_f32_16x16x128_bf8_fp8,
amdgcn_smfmac_f32_16x16x128_fp8_bf8, amdgcn_smfmac_f32_16x16x128_fp8_fp8,
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.gfx950.bf16.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.gfx950.bf16.ll
index c28536adace87..ce660cd4bb6c3 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.gfx950.bf16.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.gfx950.bf16.ll
@@ -1,9 +1,8 @@
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 4
-; RUN: llc -mtriple=amdgpu9.50 -amdgpu-mfma-vgpr-form=0 < %s | FileCheck -enable-var-scope --check-prefix=AGPR %s
-; RUN: llc -mtriple=amdgpu9.50 < %s | FileCheck -enable-var-scope --check-prefix=VGPR %s
-
-; FIXME: bfloat vector arguments are broken in globalisel.
-; https://github.com/llvm/llvm-project/issues/77055
+; RUN: llc -mtriple=amdgpu9.50 -global-isel=0 -amdgpu-mfma-vgpr-form=0 < %s | FileCheck -enable-var-scope --check-prefixes=AGPR,SDAG-AGPR %s
+; RUN: llc -mtriple=amdgpu9.50 -global-isel=1 -amdgpu-mfma-vgpr-form=0 < %s | FileCheck -enable-var-scope --check-prefixes=AGPR,GISEL-AGPR %s
+; RUN: llc -mtriple=amdgpu9.50 -global-isel=0 < %s | FileCheck -enable-var-scope --check-prefixes=VGPR,SDAG-VGPR %s
+; RUN: llc -mtriple=amdgpu9.50 -global-isel=1 < %s | FileCheck -enable-var-scope --check-prefixes=VGPR,GISEL-VGPR %s
declare <16 x float> @llvm.amdgcn.mfma.f32.32x32x16.bf16(<8 x bfloat>, <8 x bfloat>, <16 x float>, i32 immarg, i32 immarg, i32 immarg)
@@ -78,130 +77,256 @@ define amdgpu_kernel void @test_mfma_f32_32x32x16_bf16(<8 x bfloat> %arg0, <8 x
; GCN-NEXT: s_waitcnt vmcnt(0)
; GCN-NEXT: s_endpgm
;
-; AGPR-LABEL: test_mfma_f32_32x32x16_bf16:
-; AGPR: ; %bb.0:
-; AGPR-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x24
-; AGPR-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x64
-; AGPR-NEXT: v_mov_b64_e32 v[0:1], 48
-; AGPR-NEXT: v_mov_b64_e32 v[2:3], 32
-; AGPR-NEXT: v_mov_b64_e32 v[4:5], 16
-; AGPR-NEXT: s_waitcnt lgkmcnt(0)
-; AGPR-NEXT: v_mov_b64_e32 v[8:9], s[24:25]
-; AGPR-NEXT: v_mov_b64_e32 v[10:11], s[26:27]
-; AGPR-NEXT: v_mov_b64_e32 v[12:13], s[28:29]
-; AGPR-NEXT: v_mov_b64_e32 v[14:15], s[30:31]
-; AGPR-NEXT: v_accvgpr_write_b32 a0, s8
-; AGPR-NEXT: v_accvgpr_write_b32 a1, s9
-; AGPR-NEXT: v_accvgpr_write_b32 a2, s10
-; AGPR-NEXT: v_accvgpr_write_b32 a3, s11
-; AGPR-NEXT: v_accvgpr_write_b32 a4, s12
-; AGPR-NEXT: v_accvgpr_write_b32 a5, s13
-; AGPR-NEXT: v_accvgpr_write_b32 a6, s14
-; AGPR-NEXT: v_accvgpr_write_b32 a7, s15
-; AGPR-NEXT: v_accvgpr_write_b32 a8, s16
-; AGPR-NEXT: v_accvgpr_write_b32 a9, s17
-; AGPR-NEXT: v_accvgpr_write_b32 a10, s18
-; AGPR-NEXT: v_accvgpr_write_b32 a11, s19
-; AGPR-NEXT: v_accvgpr_write_b32 a12, s20
-; AGPR-NEXT: v_accvgpr_write_b32 a13, s21
-; AGPR-NEXT: v_accvgpr_write_b32 a14, s22
-; AGPR-NEXT: v_accvgpr_write_b32 a15, s23
-; AGPR-NEXT: v_mov_b64_e32 v[6:7], 0
-; AGPR-NEXT: v_mov_b32_e32 v16, s16
-; AGPR-NEXT: v_mfma_f32_32x32x16_bf16 a[16:31], v[8:11], v[12:15], a[0:15]
-; AGPR-NEXT: v_mov_b32_e32 v8, s20
-; AGPR-NEXT: v_mov_b32_e32 v9, s21
-; AGPR-NEXT: v_mov_b32_e32 v10, s22
-; AGPR-NEXT: v_mov_b32_e32 v11, s23
-; AGPR-NEXT: v_mov_b32_e32 v17, s17
-; AGPR-NEXT: v_mov_b32_e32 v18, s18
-; AGPR-NEXT: v_mov_b32_e32 v19, s19
-; AGPR-NEXT: s_nop 4
-; AGPR-NEXT: global_store_dwordx4 v[0:1], a[28:31], off sc0 sc1
-; AGPR-NEXT: s_waitcnt vmcnt(0)
-; AGPR-NEXT: global_store_dwordx4 v[2:3], a[24:27], off sc0 sc1
-; AGPR-NEXT: s_waitcnt vmcnt(0)
-; AGPR-NEXT: global_store_dwordx4 v[4:5], a[20:23], off sc0 sc1
-; AGPR-NEXT: s_waitcnt vmcnt(0)
-; AGPR-NEXT: global_store_dwordx4 v[6:7], a[16:19], off sc0 sc1
-; AGPR-NEXT: s_waitcnt vmcnt(0)
-; AGPR-NEXT: global_store_dwordx4 v[2:3], v[16:19], off sc0 sc1
-; AGPR-NEXT: s_waitcnt vmcnt(0)
-; AGPR-NEXT: global_store_dwordx4 v[0:1], v[8:11], off sc0 sc1
-; AGPR-NEXT: s_waitcnt vmcnt(0)
-; AGPR-NEXT: v_mov_b32_e32 v0, s8
-; AGPR-NEXT: v_mov_b32_e32 v1, s9
-; AGPR-NEXT: v_mov_b32_e32 v2, s10
-; AGPR-NEXT: v_mov_b32_e32 v3, s11
-; AGPR-NEXT: global_store_dwordx4 v[6:7], v[0:3], off sc0 sc1
-; AGPR-NEXT: s_waitcnt vmcnt(0)
-; AGPR-NEXT: s_nop 0
-; AGPR-NEXT: v_mov_b32_e32 v0, s12
-; AGPR-NEXT: v_mov_b32_e32 v1, s13
-; AGPR-NEXT: v_mov_b32_e32 v2, s14
-; AGPR-NEXT: v_mov_b32_e32 v3, s15
-; AGPR-NEXT: global_store_dwordx4 v[4:5], v[0:3], off sc0 sc1
-; AGPR-NEXT: s_waitcnt vmcnt(0)
-; AGPR-NEXT: s_endpgm
+; SDAG-AGPR-LABEL: test_mfma_f32_32x32x16_bf16:
+; SDAG-AGPR: ; %bb.0:
+; SDAG-AGPR-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x24
+; SDAG-AGPR-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x64
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[0:1], 48
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[2:3], 32
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[4:5], 16
+; SDAG-AGPR-NEXT: s_waitcnt lgkmcnt(0)
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[8:9], s[24:25]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[10:11], s[26:27]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[12:13], s[28:29]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[14:15], s[30:31]
+; SDAG-AGPR-NEXT: v_accvgpr_write_b32 a0, s8
+; SDAG-AGPR-NEXT: v_accvgpr_write_b32 a1, s9
+; SDAG-AGPR-NEXT: v_accvgpr_write_b32 a2, s10
+; SDAG-AGPR-NEXT: v_accvgpr_write_b32 a3, s11
+; SDAG-AGPR-NEXT: v_accvgpr_write_b32 a4, s12
+; SDAG-AGPR-NEXT: v_accvgpr_write_b32 a5, s13
+; SDAG-AGPR-NEXT: v_accvgpr_write_b32 a6, s14
+; SDAG-AGPR-NEXT: v_accvgpr_write_b32 a7, s15
+; SDAG-AGPR-NEXT: v_accvgpr_write_b32 a8, s16
+; SDAG-AGPR-NEXT: v_accvgpr_write_b32 a9, s17
+; SDAG-AGPR-NEXT: v_accvgpr_write_b32 a10, s18
+; SDAG-AGPR-NEXT: v_accvgpr_write_b32 a11, s19
+; SDAG-AGPR-NEXT: v_accvgpr_write_b32 a12, s20
+; SDAG-AGPR-NEXT: v_accvgpr_write_b32 a13, s21
+; SDAG-AGPR-NEXT: v_accvgpr_write_b32 a14, s22
+; SDAG-AGPR-NEXT: v_accvgpr_write_b32 a15, s23
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[6:7], 0
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v16, s16
+; SDAG-AGPR-NEXT: v_mfma_f32_32x32x16_bf16 a[16:31], v[8:11], v[12:15], a[0:15]
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v8, s20
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v9, s21
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v10, s22
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v11, s23
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v17, s17
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v18, s18
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v19, s19
+; SDAG-AGPR-NEXT: s_nop 4
+; SDAG-AGPR-NEXT: global_store_dwordx4 v[0:1], a[28:31], off sc0 sc1
+; SDAG-AGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-AGPR-NEXT: global_store_dwordx4 v[2:3], a[24:27], off sc0 sc1
+; SDAG-AGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-AGPR-NEXT: global_store_dwordx4 v[4:5], a[20:23], off sc0 sc1
+; SDAG-AGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-AGPR-NEXT: global_store_dwordx4 v[6:7], a[16:19], off sc0 sc1
+; SDAG-AGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-AGPR-NEXT: global_store_dwordx4 v[2:3], v[16:19], off sc0 sc1
+; SDAG-AGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-AGPR-NEXT: global_store_dwordx4 v[0:1], v[8:11], off sc0 sc1
+; SDAG-AGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v0, s8
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v1, s9
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v2, s10
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v3, s11
+; SDAG-AGPR-NEXT: global_store_dwordx4 v[6:7], v[0:3], off sc0 sc1
+; SDAG-AGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-AGPR-NEXT: s_nop 0
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v0, s12
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v1, s13
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v2, s14
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v3, s15
+; SDAG-AGPR-NEXT: global_store_dwordx4 v[4:5], v[0:3], off sc0 sc1
+; SDAG-AGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-AGPR-NEXT: s_endpgm
;
-; VGPR-LABEL: test_mfma_f32_32x32x16_bf16:
-; VGPR: ; %bb.0:
-; VGPR-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x24
-; VGPR-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x64
-; VGPR-NEXT: v_mov_b64_e32 v[32:33], 48
-; VGPR-NEXT: v_mov_b64_e32 v[34:35], 32
-; VGPR-NEXT: v_mov_b64_e32 v[36:37], 16
-; VGPR-NEXT: s_waitcnt lgkmcnt(0)
-; VGPR-NEXT: v_mov_b64_e32 v[42:43], s[26:27]
-; VGPR-NEXT: v_mov_b64_e32 v[40:41], s[24:25]
-; VGPR-NEXT: v_mov_b64_e32 v[46:47], s[30:31]
-; VGPR-NEXT: v_mov_b64_e32 v[44:45], s[28:29]
-; VGPR-NEXT: v_mov_b64_e32 v[0:1], s[8:9]
-; VGPR-NEXT: v_mov_b64_e32 v[2:3], s[10:11]
-; VGPR-NEXT: v_mov_b64_e32 v[4:5], s[12:13]
-; VGPR-NEXT: v_mov_b64_e32 v[6:7], s[14:15]
-; VGPR-NEXT: v_mov_b64_e32 v[8:9], s[16:17]
-; VGPR-NEXT: v_mov_b64_e32 v[10:11], s[18:19]
-; VGPR-NEXT: v_mov_b64_e32 v[12:13], s[20:21]
-; VGPR-NEXT: v_mov_b64_e32 v[14:15], s[22:23]
-; VGPR-NEXT: v_mov_b64_e32 v[38:39], 0
-; VGPR-NEXT: v_mov_b32_e32 v48, s16
-; VGPR-NEXT: v_mfma_f32_32x32x16_bf16 v[16:31], v[40:43], v[44:47], v[0:15]
-; VGPR-NEXT: v_mov_b32_e32 v49, s17
-; VGPR-NEXT: v_mov_b32_e32 v50, s18
-; VGPR-NEXT: v_mov_b32_e32 v51, s19
-; VGPR-NEXT: s_nop 8
-; VGPR-NEXT: global_store_dwordx4 v[32:33], v[28:31], off sc0 sc1
-; VGPR-NEXT: s_waitcnt vmcnt(0)
-; VGPR-NEXT: global_store_dwordx4 v[34:35], v[24:27], off sc0 sc1
-; VGPR-NEXT: s_waitcnt vmcnt(0)
-; VGPR-NEXT: global_store_dwordx4 v[36:37], v[20:23], off sc0 sc1
-; VGPR-NEXT: s_waitcnt vmcnt(0)
-; VGPR-NEXT: global_store_dwordx4 v[38:39], v[16:19], off sc0 sc1
-; VGPR-NEXT: s_waitcnt vmcnt(0)
-; VGPR-NEXT: global_store_dwordx4 v[34:35], v[48:51], off sc0 sc1
-; VGPR-NEXT: s_waitcnt vmcnt(0)
-; VGPR-NEXT: v_mov_b32_e32 v0, s20
-; VGPR-NEXT: v_mov_b32_e32 v1, s21
-; VGPR-NEXT: v_mov_b32_e32 v2, s22
-; VGPR-NEXT: v_mov_b32_e32 v3, s23
-; VGPR-NEXT: global_store_dwordx4 v[32:33], v[0:3], off sc0 sc1
-; VGPR-NEXT: s_waitcnt vmcnt(0)
-; VGPR-NEXT: s_nop 0
-; VGPR-NEXT: v_mov_b32_e32 v0, s8
-; VGPR-NEXT: v_mov_b32_e32 v1, s9
-; VGPR-NEXT: v_mov_b32_e32 v2, s10
-; VGPR-NEXT: v_mov_b32_e32 v3, s11
-; VGPR-NEXT: global_store_dwordx4 v[38:39], v[0:3], off sc0 sc1
-; VGPR-NEXT: s_waitcnt vmcnt(0)
-; VGPR-NEXT: s_nop 0
-; VGPR-NEXT: v_mov_b32_e32 v0, s12
-; VGPR-NEXT: v_mov_b32_e32 v1, s13
-; VGPR-NEXT: v_mov_b32_e32 v2, s14
-; VGPR-NEXT: v_mov_b32_e32 v3, s15
-; VGPR-NEXT: global_store_dwordx4 v[36:37], v[0:3], off sc0 sc1
-; VGPR-NEXT: s_waitcnt vmcnt(0)
-; VGPR-NEXT: s_endpgm
+; GISEL-AGPR-LABEL: test_mfma_f32_32x32x16_bf16:
+; GISEL-AGPR: ; %bb.0:
+; GISEL-AGPR-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x24
+; GISEL-AGPR-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x64
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[20:21], 0
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[22:23], 16
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[24:25], 32
+; GISEL-AGPR-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[0:1], s[24:25]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[2:3], s[26:27]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[4:5], s[28:29]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[6:7], s[30:31]
+; GISEL-AGPR-NEXT: v_accvgpr_write_b32 a0, s8
+; GISEL-AGPR-NEXT: v_accvgpr_write_b32 a1, s9
+; GISEL-AGPR-NEXT: v_accvgpr_write_b32 a2, s10
+; GISEL-AGPR-NEXT: v_accvgpr_write_b32 a3, s11
+; GISEL-AGPR-NEXT: v_accvgpr_write_b32 a4, s12
+; GISEL-AGPR-NEXT: v_accvgpr_write_b32 a5, s13
+; GISEL-AGPR-NEXT: v_accvgpr_write_b32 a6, s14
+; GISEL-AGPR-NEXT: v_accvgpr_write_b32 a7, s15
+; GISEL-AGPR-NEXT: v_accvgpr_write_b32 a8, s16
+; GISEL-AGPR-NEXT: v_accvgpr_write_b32 a9, s17
+; GISEL-AGPR-NEXT: v_accvgpr_write_b32 a10, s18
+; GISEL-AGPR-NEXT: v_accvgpr_write_b32 a11, s19
+; GISEL-AGPR-NEXT: v_accvgpr_write_b32 a12, s20
+; GISEL-AGPR-NEXT: v_accvgpr_write_b32 a13, s21
+; GISEL-AGPR-NEXT: v_accvgpr_write_b32 a14, s22
+; GISEL-AGPR-NEXT: v_accvgpr_write_b32 a15, s23
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[26:27], 48
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[18:19], s[10:11]
+; GISEL-AGPR-NEXT: v_mfma_f32_32x32x16_bf16 a[16:31], v[0:3], v[4:7], a[0:15]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[16:17], s[8:9]
+; GISEL-AGPR-NEXT: s_nop 10
+; GISEL-AGPR-NEXT: v_accvgpr_read_b32 v0, a16
+; GISEL-AGPR-NEXT: v_accvgpr_read_b32 v1, a17
+; GISEL-AGPR-NEXT: v_accvgpr_read_b32 v2, a18
+; GISEL-AGPR-NEXT: v_accvgpr_read_b32 v3, a19
+; GISEL-AGPR-NEXT: v_accvgpr_read_b32 v4, a20
+; GISEL-AGPR-NEXT: v_accvgpr_read_b32 v5, a21
+; GISEL-AGPR-NEXT: v_accvgpr_read_b32 v6, a22
+; GISEL-AGPR-NEXT: v_accvgpr_read_b32 v7, a23
+; GISEL-AGPR-NEXT: v_accvgpr_read_b32 v8, a24
+; GISEL-AGPR-NEXT: v_accvgpr_read_b32 v9, a25
+; GISEL-AGPR-NEXT: v_accvgpr_read_b32 v10, a26
+; GISEL-AGPR-NEXT: v_accvgpr_read_b32 v11, a27
+; GISEL-AGPR-NEXT: v_accvgpr_read_b32 v12, a28
+; GISEL-AGPR-NEXT: v_accvgpr_read_b32 v13, a29
+; GISEL-AGPR-NEXT: v_accvgpr_read_b32 v14, a30
+; GISEL-AGPR-NEXT: v_accvgpr_read_b32 v15, a31
+; GISEL-AGPR-NEXT: global_store_dwordx4 v[20:21], v[0:3], off sc0 sc1
+; GISEL-AGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-AGPR-NEXT: global_store_dwordx4 v[22:23], v[4:7], off sc0 sc1
+; GISEL-AGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-AGPR-NEXT: global_store_dwordx4 v[24:25], v[8:11], off sc0 sc1
+; GISEL-AGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-AGPR-NEXT: global_store_dwordx4 v[26:27], v[12:15], off sc0 sc1
+; GISEL-AGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-AGPR-NEXT: global_store_dwordx4 v[20:21], v[16:19], off sc0 sc1
+; GISEL-AGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[0:1], s[12:13]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[2:3], s[14:15]
+; GISEL-AGPR-NEXT: global_store_dwordx4 v[22:23], v[0:3], off sc0 sc1
+; GISEL-AGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-AGPR-NEXT: s_nop 0
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[0:1], s[16:17]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[2:3], s[18:19]
+; GISEL-AGPR-NEXT: global_store_dwordx4 v[24:25], v[0:3], off sc0 sc1
+; GISEL-AGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-AGPR-NEXT: s_nop 0
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[0:1], s[20:21]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[2:3], s[22:23]
+; GISEL-AGPR-NEXT: global_store_dwordx4 v[26:27], v[0:3], off sc0 sc1
+; GISEL-AGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-AGPR-NEXT: s_endpgm
+;
+; SDAG-VGPR-LABEL: test_mfma_f32_32x32x16_bf16:
+; SDAG-VGPR: ; %bb.0:
+; SDAG-VGPR-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x24
+; SDAG-VGPR-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x64
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[32:33], 48
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[34:35], 32
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[36:37], 16
+; SDAG-VGPR-NEXT: s_waitcnt lgkmcnt(0)
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[42:43], s[26:27]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[40:41], s[24:25]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[46:47], s[30:31]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[44:45], s[28:29]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[0:1], s[8:9]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[2:3], s[10:11]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[4:5], s[12:13]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[6:7], s[14:15]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[8:9], s[16:17]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[10:11], s[18:19]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[12:13], s[20:21]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[14:15], s[22:23]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[38:39], 0
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v48, s16
+; SDAG-VGPR-NEXT: v_mfma_f32_32x32x16_bf16 v[16:31], v[40:43], v[44:47], v[0:15]
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v49, s17
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v50, s18
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v51, s19
+; SDAG-VGPR-NEXT: s_nop 8
+; SDAG-VGPR-NEXT: global_store_dwordx4 v[32:33], v[28:31], off sc0 sc1
+; SDAG-VGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-VGPR-NEXT: global_store_dwordx4 v[34:35], v[24:27], off sc0 sc1
+; SDAG-VGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-VGPR-NEXT: global_store_dwordx4 v[36:37], v[20:23], off sc0 sc1
+; SDAG-VGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-VGPR-NEXT: global_store_dwordx4 v[38:39], v[16:19], off sc0 sc1
+; SDAG-VGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-VGPR-NEXT: global_store_dwordx4 v[34:35], v[48:51], off sc0 sc1
+; SDAG-VGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v0, s20
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v1, s21
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v2, s22
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v3, s23
+; SDAG-VGPR-NEXT: global_store_dwordx4 v[32:33], v[0:3], off sc0 sc1
+; SDAG-VGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-VGPR-NEXT: s_nop 0
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v0, s8
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v1, s9
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v2, s10
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v3, s11
+; SDAG-VGPR-NEXT: global_store_dwordx4 v[38:39], v[0:3], off sc0 sc1
+; SDAG-VGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-VGPR-NEXT: s_nop 0
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v0, s12
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v1, s13
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v2, s14
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v3, s15
+; SDAG-VGPR-NEXT: global_store_dwordx4 v[36:37], v[0:3], off sc0 sc1
+; SDAG-VGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-VGPR-NEXT: s_endpgm
+;
+; GISEL-VGPR-LABEL: test_mfma_f32_32x32x16_bf16:
+; GISEL-VGPR: ; %bb.0:
+; GISEL-VGPR-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x24
+; GISEL-VGPR-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x64
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[44:45], 0
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[46:47], 16
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[48:49], 32
+; GISEL-VGPR-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[34:35], s[26:27]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[32:33], s[24:25]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[38:39], s[30:31]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[36:37], s[28:29]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[0:1], s[8:9]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[2:3], s[10:11]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[4:5], s[12:13]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[6:7], s[14:15]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[8:9], s[16:17]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[10:11], s[18:19]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[12:13], s[20:21]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[14:15], s[22:23]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[50:51], 48
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[42:43], s[10:11]
+; GISEL-VGPR-NEXT: v_mfma_f32_32x32x16_bf16 v[16:31], v[32:35], v[36:39], v[0:15]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[40:41], s[8:9]
+; GISEL-VGPR-NEXT: s_nop 10
+; GISEL-VGPR-NEXT: global_store_dwordx4 v[44:45], v[16:19], off sc0 sc1
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-VGPR-NEXT: global_store_dwordx4 v[46:47], v[20:23], off sc0 sc1
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-VGPR-NEXT: global_store_dwordx4 v[48:49], v[24:27], off sc0 sc1
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-VGPR-NEXT: global_store_dwordx4 v[50:51], v[28:31], off sc0 sc1
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-VGPR-NEXT: global_store_dwordx4 v[44:45], v[40:43], off sc0 sc1
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[0:1], s[12:13]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[2:3], s[14:15]
+; GISEL-VGPR-NEXT: global_store_dwordx4 v[46:47], v[0:3], off sc0 sc1
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-VGPR-NEXT: s_nop 0
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[0:1], s[16:17]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[2:3], s[18:19]
+; GISEL-VGPR-NEXT: global_store_dwordx4 v[48:49], v[0:3], off sc0 sc1
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-VGPR-NEXT: s_nop 0
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[0:1], s[20:21]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[2:3], s[22:23]
+; GISEL-VGPR-NEXT: global_store_dwordx4 v[50:51], v[0:3], off sc0 sc1
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-VGPR-NEXT: s_endpgm
%result = call <16 x float> @llvm.amdgcn.mfma.f32.32x32x16.bf16(<8 x bfloat> %arg0, <8 x bfloat> %arg1, <16 x float> %arg2, i32 0, i32 0, i32 0)
store volatile <16 x float> %result, ptr addrspace(1) null
store volatile <16 x float> %arg2, ptr addrspace(1) null
@@ -275,130 +400,256 @@ define amdgpu_kernel void @test_mfma_f32_32x32x16_bf16__flags(<8 x bfloat> %arg0
; GCN-NEXT: s_waitcnt vmcnt(0)
; GCN-NEXT: s_endpgm
;
-; AGPR-LABEL: test_mfma_f32_32x32x16_bf16__flags:
-; AGPR: ; %bb.0:
-; AGPR-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x24
-; AGPR-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x64
-; AGPR-NEXT: v_mov_b64_e32 v[0:1], 48
-; AGPR-NEXT: v_mov_b64_e32 v[2:3], 32
-; AGPR-NEXT: v_mov_b64_e32 v[4:5], 16
-; AGPR-NEXT: s_waitcnt lgkmcnt(0)
-; AGPR-NEXT: v_mov_b64_e32 v[8:9], s[24:25]
-; AGPR-NEXT: v_mov_b64_e32 v[10:11], s[26:27]
-; AGPR-NEXT: v_mov_b64_e32 v[12:13], s[28:29]
-; AGPR-NEXT: v_mov_b64_e32 v[14:15], s[30:31]
-; AGPR-NEXT: v_accvgpr_write_b32 a0, s8
-; AGPR-NEXT: v_accvgpr_write_b32 a1, s9
-; AGPR-NEXT: v_accvgpr_write_b32 a2, s10
-; AGPR-NEXT: v_accvgpr_write_b32 a3, s11
-; AGPR-NEXT: v_accvgpr_write_b32 a4, s12
-; AGPR-NEXT: v_accvgpr_write_b32 a5, s13
-; AGPR-NEXT: v_accvgpr_write_b32 a6, s14
-; AGPR-NEXT: v_accvgpr_write_b32 a7, s15
-; AGPR-NEXT: v_accvgpr_write_b32 a8, s16
-; AGPR-NEXT: v_accvgpr_write_b32 a9, s17
-; AGPR-NEXT: v_accvgpr_write_b32 a10, s18
-; AGPR-NEXT: v_accvgpr_write_b32 a11, s19
-; AGPR-NEXT: v_accvgpr_write_b32 a12, s20
-; AGPR-NEXT: v_accvgpr_write_b32 a13, s21
-; AGPR-NEXT: v_accvgpr_write_b32 a14, s22
-; AGPR-NEXT: v_accvgpr_write_b32 a15, s23
-; AGPR-NEXT: v_mov_b64_e32 v[6:7], 0
-; AGPR-NEXT: v_mov_b32_e32 v16, s16
-; AGPR-NEXT: v_mfma_f32_32x32x16_bf16 a[16:31], v[8:11], v[12:15], a[0:15] cbsz:2 abid:3 blgp:1
-; AGPR-NEXT: v_mov_b32_e32 v8, s20
-; AGPR-NEXT: v_mov_b32_e32 v9, s21
-; AGPR-NEXT: v_mov_b32_e32 v10, s22
-; AGPR-NEXT: v_mov_b32_e32 v11, s23
-; AGPR-NEXT: v_mov_b32_e32 v17, s17
-; AGPR-NEXT: v_mov_b32_e32 v18, s18
-; AGPR-NEXT: v_mov_b32_e32 v19, s19
-; AGPR-NEXT: s_nop 4
-; AGPR-NEXT: global_store_dwordx4 v[0:1], a[28:31], off sc0 sc1
-; AGPR-NEXT: s_waitcnt vmcnt(0)
-; AGPR-NEXT: global_store_dwordx4 v[2:3], a[24:27], off sc0 sc1
-; AGPR-NEXT: s_waitcnt vmcnt(0)
-; AGPR-NEXT: global_store_dwordx4 v[4:5], a[20:23], off sc0 sc1
-; AGPR-NEXT: s_waitcnt vmcnt(0)
-; AGPR-NEXT: global_store_dwordx4 v[6:7], a[16:19], off sc0 sc1
-; AGPR-NEXT: s_waitcnt vmcnt(0)
-; AGPR-NEXT: global_store_dwordx4 v[2:3], v[16:19], off sc0 sc1
-; AGPR-NEXT: s_waitcnt vmcnt(0)
-; AGPR-NEXT: global_store_dwordx4 v[0:1], v[8:11], off sc0 sc1
-; AGPR-NEXT: s_waitcnt vmcnt(0)
-; AGPR-NEXT: v_mov_b32_e32 v0, s8
-; AGPR-NEXT: v_mov_b32_e32 v1, s9
-; AGPR-NEXT: v_mov_b32_e32 v2, s10
-; AGPR-NEXT: v_mov_b32_e32 v3, s11
-; AGPR-NEXT: global_store_dwordx4 v[6:7], v[0:3], off sc0 sc1
-; AGPR-NEXT: s_waitcnt vmcnt(0)
-; AGPR-NEXT: s_nop 0
-; AGPR-NEXT: v_mov_b32_e32 v0, s12
-; AGPR-NEXT: v_mov_b32_e32 v1, s13
-; AGPR-NEXT: v_mov_b32_e32 v2, s14
-; AGPR-NEXT: v_mov_b32_e32 v3, s15
-; AGPR-NEXT: global_store_dwordx4 v[4:5], v[0:3], off sc0 sc1
-; AGPR-NEXT: s_waitcnt vmcnt(0)
-; AGPR-NEXT: s_endpgm
+; SDAG-AGPR-LABEL: test_mfma_f32_32x32x16_bf16__flags:
+; SDAG-AGPR: ; %bb.0:
+; SDAG-AGPR-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x24
+; SDAG-AGPR-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x64
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[0:1], 48
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[2:3], 32
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[4:5], 16
+; SDAG-AGPR-NEXT: s_waitcnt lgkmcnt(0)
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[8:9], s[24:25]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[10:11], s[26:27]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[12:13], s[28:29]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[14:15], s[30:31]
+; SDAG-AGPR-NEXT: v_accvgpr_write_b32 a0, s8
+; SDAG-AGPR-NEXT: v_accvgpr_write_b32 a1, s9
+; SDAG-AGPR-NEXT: v_accvgpr_write_b32 a2, s10
+; SDAG-AGPR-NEXT: v_accvgpr_write_b32 a3, s11
+; SDAG-AGPR-NEXT: v_accvgpr_write_b32 a4, s12
+; SDAG-AGPR-NEXT: v_accvgpr_write_b32 a5, s13
+; SDAG-AGPR-NEXT: v_accvgpr_write_b32 a6, s14
+; SDAG-AGPR-NEXT: v_accvgpr_write_b32 a7, s15
+; SDAG-AGPR-NEXT: v_accvgpr_write_b32 a8, s16
+; SDAG-AGPR-NEXT: v_accvgpr_write_b32 a9, s17
+; SDAG-AGPR-NEXT: v_accvgpr_write_b32 a10, s18
+; SDAG-AGPR-NEXT: v_accvgpr_write_b32 a11, s19
+; SDAG-AGPR-NEXT: v_accvgpr_write_b32 a12, s20
+; SDAG-AGPR-NEXT: v_accvgpr_write_b32 a13, s21
+; SDAG-AGPR-NEXT: v_accvgpr_write_b32 a14, s22
+; SDAG-AGPR-NEXT: v_accvgpr_write_b32 a15, s23
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[6:7], 0
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v16, s16
+; SDAG-AGPR-NEXT: v_mfma_f32_32x32x16_bf16 a[16:31], v[8:11], v[12:15], a[0:15] cbsz:2 abid:3 blgp:1
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v8, s20
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v9, s21
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v10, s22
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v11, s23
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v17, s17
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v18, s18
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v19, s19
+; SDAG-AGPR-NEXT: s_nop 4
+; SDAG-AGPR-NEXT: global_store_dwordx4 v[0:1], a[28:31], off sc0 sc1
+; SDAG-AGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-AGPR-NEXT: global_store_dwordx4 v[2:3], a[24:27], off sc0 sc1
+; SDAG-AGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-AGPR-NEXT: global_store_dwordx4 v[4:5], a[20:23], off sc0 sc1
+; SDAG-AGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-AGPR-NEXT: global_store_dwordx4 v[6:7], a[16:19], off sc0 sc1
+; SDAG-AGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-AGPR-NEXT: global_store_dwordx4 v[2:3], v[16:19], off sc0 sc1
+; SDAG-AGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-AGPR-NEXT: global_store_dwordx4 v[0:1], v[8:11], off sc0 sc1
+; SDAG-AGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v0, s8
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v1, s9
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v2, s10
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v3, s11
+; SDAG-AGPR-NEXT: global_store_dwordx4 v[6:7], v[0:3], off sc0 sc1
+; SDAG-AGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-AGPR-NEXT: s_nop 0
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v0, s12
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v1, s13
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v2, s14
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v3, s15
+; SDAG-AGPR-NEXT: global_store_dwordx4 v[4:5], v[0:3], off sc0 sc1
+; SDAG-AGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-AGPR-NEXT: s_endpgm
;
-; VGPR-LABEL: test_mfma_f32_32x32x16_bf16__flags:
-; VGPR: ; %bb.0:
-; VGPR-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x24
-; VGPR-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x64
-; VGPR-NEXT: v_mov_b64_e32 v[32:33], 48
-; VGPR-NEXT: v_mov_b64_e32 v[34:35], 32
-; VGPR-NEXT: v_mov_b64_e32 v[36:37], 16
-; VGPR-NEXT: s_waitcnt lgkmcnt(0)
-; VGPR-NEXT: v_mov_b64_e32 v[42:43], s[26:27]
-; VGPR-NEXT: v_mov_b64_e32 v[40:41], s[24:25]
-; VGPR-NEXT: v_mov_b64_e32 v[46:47], s[30:31]
-; VGPR-NEXT: v_mov_b64_e32 v[44:45], s[28:29]
-; VGPR-NEXT: v_mov_b64_e32 v[0:1], s[8:9]
-; VGPR-NEXT: v_mov_b64_e32 v[2:3], s[10:11]
-; VGPR-NEXT: v_mov_b64_e32 v[4:5], s[12:13]
-; VGPR-NEXT: v_mov_b64_e32 v[6:7], s[14:15]
-; VGPR-NEXT: v_mov_b64_e32 v[8:9], s[16:17]
-; VGPR-NEXT: v_mov_b64_e32 v[10:11], s[18:19]
-; VGPR-NEXT: v_mov_b64_e32 v[12:13], s[20:21]
-; VGPR-NEXT: v_mov_b64_e32 v[14:15], s[22:23]
-; VGPR-NEXT: v_mov_b64_e32 v[38:39], 0
-; VGPR-NEXT: v_mov_b32_e32 v48, s16
-; VGPR-NEXT: v_mfma_f32_32x32x16_bf16 v[16:31], v[40:43], v[44:47], v[0:15] cbsz:2 abid:3 blgp:1
-; VGPR-NEXT: v_mov_b32_e32 v49, s17
-; VGPR-NEXT: v_mov_b32_e32 v50, s18
-; VGPR-NEXT: v_mov_b32_e32 v51, s19
-; VGPR-NEXT: s_nop 8
-; VGPR-NEXT: global_store_dwordx4 v[32:33], v[28:31], off sc0 sc1
-; VGPR-NEXT: s_waitcnt vmcnt(0)
-; VGPR-NEXT: global_store_dwordx4 v[34:35], v[24:27], off sc0 sc1
-; VGPR-NEXT: s_waitcnt vmcnt(0)
-; VGPR-NEXT: global_store_dwordx4 v[36:37], v[20:23], off sc0 sc1
-; VGPR-NEXT: s_waitcnt vmcnt(0)
-; VGPR-NEXT: global_store_dwordx4 v[38:39], v[16:19], off sc0 sc1
-; VGPR-NEXT: s_waitcnt vmcnt(0)
-; VGPR-NEXT: global_store_dwordx4 v[34:35], v[48:51], off sc0 sc1
-; VGPR-NEXT: s_waitcnt vmcnt(0)
-; VGPR-NEXT: v_mov_b32_e32 v0, s20
-; VGPR-NEXT: v_mov_b32_e32 v1, s21
-; VGPR-NEXT: v_mov_b32_e32 v2, s22
-; VGPR-NEXT: v_mov_b32_e32 v3, s23
-; VGPR-NEXT: global_store_dwordx4 v[32:33], v[0:3], off sc0 sc1
-; VGPR-NEXT: s_waitcnt vmcnt(0)
-; VGPR-NEXT: s_nop 0
-; VGPR-NEXT: v_mov_b32_e32 v0, s8
-; VGPR-NEXT: v_mov_b32_e32 v1, s9
-; VGPR-NEXT: v_mov_b32_e32 v2, s10
-; VGPR-NEXT: v_mov_b32_e32 v3, s11
-; VGPR-NEXT: global_store_dwordx4 v[38:39], v[0:3], off sc0 sc1
-; VGPR-NEXT: s_waitcnt vmcnt(0)
-; VGPR-NEXT: s_nop 0
-; VGPR-NEXT: v_mov_b32_e32 v0, s12
-; VGPR-NEXT: v_mov_b32_e32 v1, s13
-; VGPR-NEXT: v_mov_b32_e32 v2, s14
-; VGPR-NEXT: v_mov_b32_e32 v3, s15
-; VGPR-NEXT: global_store_dwordx4 v[36:37], v[0:3], off sc0 sc1
-; VGPR-NEXT: s_waitcnt vmcnt(0)
-; VGPR-NEXT: s_endpgm
+; GISEL-AGPR-LABEL: test_mfma_f32_32x32x16_bf16__flags:
+; GISEL-AGPR: ; %bb.0:
+; GISEL-AGPR-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x24
+; GISEL-AGPR-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x64
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[20:21], 0
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[22:23], 16
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[24:25], 32
+; GISEL-AGPR-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[0:1], s[24:25]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[2:3], s[26:27]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[4:5], s[28:29]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[6:7], s[30:31]
+; GISEL-AGPR-NEXT: v_accvgpr_write_b32 a0, s8
+; GISEL-AGPR-NEXT: v_accvgpr_write_b32 a1, s9
+; GISEL-AGPR-NEXT: v_accvgpr_write_b32 a2, s10
+; GISEL-AGPR-NEXT: v_accvgpr_write_b32 a3, s11
+; GISEL-AGPR-NEXT: v_accvgpr_write_b32 a4, s12
+; GISEL-AGPR-NEXT: v_accvgpr_write_b32 a5, s13
+; GISEL-AGPR-NEXT: v_accvgpr_write_b32 a6, s14
+; GISEL-AGPR-NEXT: v_accvgpr_write_b32 a7, s15
+; GISEL-AGPR-NEXT: v_accvgpr_write_b32 a8, s16
+; GISEL-AGPR-NEXT: v_accvgpr_write_b32 a9, s17
+; GISEL-AGPR-NEXT: v_accvgpr_write_b32 a10, s18
+; GISEL-AGPR-NEXT: v_accvgpr_write_b32 a11, s19
+; GISEL-AGPR-NEXT: v_accvgpr_write_b32 a12, s20
+; GISEL-AGPR-NEXT: v_accvgpr_write_b32 a13, s21
+; GISEL-AGPR-NEXT: v_accvgpr_write_b32 a14, s22
+; GISEL-AGPR-NEXT: v_accvgpr_write_b32 a15, s23
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[26:27], 48
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[18:19], s[10:11]
+; GISEL-AGPR-NEXT: v_mfma_f32_32x32x16_bf16 a[16:31], v[0:3], v[4:7], a[0:15] cbsz:2 abid:3 blgp:1
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[16:17], s[8:9]
+; GISEL-AGPR-NEXT: s_nop 10
+; GISEL-AGPR-NEXT: v_accvgpr_read_b32 v0, a16
+; GISEL-AGPR-NEXT: v_accvgpr_read_b32 v1, a17
+; GISEL-AGPR-NEXT: v_accvgpr_read_b32 v2, a18
+; GISEL-AGPR-NEXT: v_accvgpr_read_b32 v3, a19
+; GISEL-AGPR-NEXT: v_accvgpr_read_b32 v4, a20
+; GISEL-AGPR-NEXT: v_accvgpr_read_b32 v5, a21
+; GISEL-AGPR-NEXT: v_accvgpr_read_b32 v6, a22
+; GISEL-AGPR-NEXT: v_accvgpr_read_b32 v7, a23
+; GISEL-AGPR-NEXT: v_accvgpr_read_b32 v8, a24
+; GISEL-AGPR-NEXT: v_accvgpr_read_b32 v9, a25
+; GISEL-AGPR-NEXT: v_accvgpr_read_b32 v10, a26
+; GISEL-AGPR-NEXT: v_accvgpr_read_b32 v11, a27
+; GISEL-AGPR-NEXT: v_accvgpr_read_b32 v12, a28
+; GISEL-AGPR-NEXT: v_accvgpr_read_b32 v13, a29
+; GISEL-AGPR-NEXT: v_accvgpr_read_b32 v14, a30
+; GISEL-AGPR-NEXT: v_accvgpr_read_b32 v15, a31
+; GISEL-AGPR-NEXT: global_store_dwordx4 v[20:21], v[0:3], off sc0 sc1
+; GISEL-AGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-AGPR-NEXT: global_store_dwordx4 v[22:23], v[4:7], off sc0 sc1
+; GISEL-AGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-AGPR-NEXT: global_store_dwordx4 v[24:25], v[8:11], off sc0 sc1
+; GISEL-AGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-AGPR-NEXT: global_store_dwordx4 v[26:27], v[12:15], off sc0 sc1
+; GISEL-AGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-AGPR-NEXT: global_store_dwordx4 v[20:21], v[16:19], off sc0 sc1
+; GISEL-AGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[0:1], s[12:13]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[2:3], s[14:15]
+; GISEL-AGPR-NEXT: global_store_dwordx4 v[22:23], v[0:3], off sc0 sc1
+; GISEL-AGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-AGPR-NEXT: s_nop 0
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[0:1], s[16:17]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[2:3], s[18:19]
+; GISEL-AGPR-NEXT: global_store_dwordx4 v[24:25], v[0:3], off sc0 sc1
+; GISEL-AGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-AGPR-NEXT: s_nop 0
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[0:1], s[20:21]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[2:3], s[22:23]
+; GISEL-AGPR-NEXT: global_store_dwordx4 v[26:27], v[0:3], off sc0 sc1
+; GISEL-AGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-AGPR-NEXT: s_endpgm
+;
+; SDAG-VGPR-LABEL: test_mfma_f32_32x32x16_bf16__flags:
+; SDAG-VGPR: ; %bb.0:
+; SDAG-VGPR-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x24
+; SDAG-VGPR-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x64
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[32:33], 48
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[34:35], 32
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[36:37], 16
+; SDAG-VGPR-NEXT: s_waitcnt lgkmcnt(0)
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[42:43], s[26:27]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[40:41], s[24:25]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[46:47], s[30:31]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[44:45], s[28:29]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[0:1], s[8:9]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[2:3], s[10:11]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[4:5], s[12:13]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[6:7], s[14:15]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[8:9], s[16:17]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[10:11], s[18:19]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[12:13], s[20:21]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[14:15], s[22:23]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[38:39], 0
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v48, s16
+; SDAG-VGPR-NEXT: v_mfma_f32_32x32x16_bf16 v[16:31], v[40:43], v[44:47], v[0:15] cbsz:2 abid:3 blgp:1
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v49, s17
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v50, s18
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v51, s19
+; SDAG-VGPR-NEXT: s_nop 8
+; SDAG-VGPR-NEXT: global_store_dwordx4 v[32:33], v[28:31], off sc0 sc1
+; SDAG-VGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-VGPR-NEXT: global_store_dwordx4 v[34:35], v[24:27], off sc0 sc1
+; SDAG-VGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-VGPR-NEXT: global_store_dwordx4 v[36:37], v[20:23], off sc0 sc1
+; SDAG-VGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-VGPR-NEXT: global_store_dwordx4 v[38:39], v[16:19], off sc0 sc1
+; SDAG-VGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-VGPR-NEXT: global_store_dwordx4 v[34:35], v[48:51], off sc0 sc1
+; SDAG-VGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v0, s20
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v1, s21
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v2, s22
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v3, s23
+; SDAG-VGPR-NEXT: global_store_dwordx4 v[32:33], v[0:3], off sc0 sc1
+; SDAG-VGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-VGPR-NEXT: s_nop 0
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v0, s8
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v1, s9
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v2, s10
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v3, s11
+; SDAG-VGPR-NEXT: global_store_dwordx4 v[38:39], v[0:3], off sc0 sc1
+; SDAG-VGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-VGPR-NEXT: s_nop 0
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v0, s12
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v1, s13
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v2, s14
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v3, s15
+; SDAG-VGPR-NEXT: global_store_dwordx4 v[36:37], v[0:3], off sc0 sc1
+; SDAG-VGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-VGPR-NEXT: s_endpgm
+;
+; GISEL-VGPR-LABEL: test_mfma_f32_32x32x16_bf16__flags:
+; GISEL-VGPR: ; %bb.0:
+; GISEL-VGPR-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x24
+; GISEL-VGPR-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x64
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[44:45], 0
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[46:47], 16
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[48:49], 32
+; GISEL-VGPR-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[34:35], s[26:27]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[32:33], s[24:25]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[38:39], s[30:31]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[36:37], s[28:29]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[0:1], s[8:9]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[2:3], s[10:11]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[4:5], s[12:13]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[6:7], s[14:15]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[8:9], s[16:17]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[10:11], s[18:19]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[12:13], s[20:21]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[14:15], s[22:23]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[50:51], 48
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[42:43], s[10:11]
+; GISEL-VGPR-NEXT: v_mfma_f32_32x32x16_bf16 v[16:31], v[32:35], v[36:39], v[0:15] cbsz:2 abid:3 blgp:1
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[40:41], s[8:9]
+; GISEL-VGPR-NEXT: s_nop 10
+; GISEL-VGPR-NEXT: global_store_dwordx4 v[44:45], v[16:19], off sc0 sc1
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-VGPR-NEXT: global_store_dwordx4 v[46:47], v[20:23], off sc0 sc1
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-VGPR-NEXT: global_store_dwordx4 v[48:49], v[24:27], off sc0 sc1
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-VGPR-NEXT: global_store_dwordx4 v[50:51], v[28:31], off sc0 sc1
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-VGPR-NEXT: global_store_dwordx4 v[44:45], v[40:43], off sc0 sc1
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[0:1], s[12:13]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[2:3], s[14:15]
+; GISEL-VGPR-NEXT: global_store_dwordx4 v[46:47], v[0:3], off sc0 sc1
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-VGPR-NEXT: s_nop 0
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[0:1], s[16:17]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[2:3], s[18:19]
+; GISEL-VGPR-NEXT: global_store_dwordx4 v[48:49], v[0:3], off sc0 sc1
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-VGPR-NEXT: s_nop 0
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[0:1], s[20:21]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[2:3], s[22:23]
+; GISEL-VGPR-NEXT: global_store_dwordx4 v[50:51], v[0:3], off sc0 sc1
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-VGPR-NEXT: s_endpgm
%result = call <16 x float> @llvm.amdgcn.mfma.f32.32x32x16.bf16(<8 x bfloat> %arg0, <8 x bfloat> %arg1, <16 x float> %arg2, i32 2, i32 3, i32 1)
store volatile <16 x float> %result, ptr addrspace(1) null
store volatile <16 x float> %arg2, ptr addrspace(1) null
@@ -677,119 +928,211 @@ define amdgpu_kernel void @test_mfma_f32_32x32x16_bf16__vgprcd(<8 x bfloat> %arg
; GCN-NEXT: s_waitcnt vmcnt(0)
; GCN-NEXT: s_endpgm
;
-; AGPR-LABEL: test_mfma_f32_32x32x16_bf16__vgprcd:
-; AGPR: ; %bb.0:
-; AGPR-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x24
-; AGPR-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x64
-; AGPR-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0xa4
-; AGPR-NEXT: v_mov_b32_e32 v32, 0
-; AGPR-NEXT: s_waitcnt lgkmcnt(0)
-; AGPR-NEXT: v_mov_b64_e32 v[36:37], s[26:27]
-; AGPR-NEXT: v_mov_b64_e32 v[34:35], s[24:25]
-; AGPR-NEXT: v_mov_b64_e32 v[40:41], s[30:31]
-; AGPR-NEXT: v_mov_b64_e32 v[38:39], s[28:29]
-; AGPR-NEXT: v_mov_b64_e32 v[30:31], s[22:23]
-; AGPR-NEXT: v_mov_b64_e32 v[28:29], s[20:21]
-; AGPR-NEXT: v_mov_b64_e32 v[26:27], s[18:19]
-; AGPR-NEXT: v_mov_b64_e32 v[24:25], s[16:17]
-; AGPR-NEXT: v_mov_b64_e32 v[22:23], s[14:15]
-; AGPR-NEXT: v_mov_b64_e32 v[20:21], s[12:13]
-; AGPR-NEXT: v_mov_b64_e32 v[18:19], s[10:11]
-; AGPR-NEXT: v_mov_b64_e32 v[16:17], s[8:9]
-; AGPR-NEXT: v_mov_b32_e32 v42, s20
-; AGPR-NEXT: v_mov_b32_e32 v43, s21
-; AGPR-NEXT: v_mfma_f32_32x32x16_bf16 v[0:15], v[34:37], v[38:41], v[16:31]
-; AGPR-NEXT: v_mov_b32_e32 v44, s22
-; AGPR-NEXT: v_mov_b32_e32 v45, s23
-; AGPR-NEXT: global_store_dwordx4 v32, v[42:45], s[0:1] offset:48 sc0 sc1
-; AGPR-NEXT: s_waitcnt vmcnt(0)
-; AGPR-NEXT: s_nop 2
-; AGPR-NEXT: v_mov_b32_e32 v16, s16
-; AGPR-NEXT: v_mov_b32_e32 v17, s17
-; AGPR-NEXT: v_mov_b32_e32 v18, s18
-; AGPR-NEXT: v_mov_b32_e32 v19, s19
-; AGPR-NEXT: global_store_dwordx4 v32, v[16:19], s[0:1] offset:32 sc0 sc1
-; AGPR-NEXT: s_waitcnt vmcnt(0)
-; AGPR-NEXT: s_nop 0
-; AGPR-NEXT: v_mov_b32_e32 v16, s12
-; AGPR-NEXT: v_mov_b32_e32 v17, s13
-; AGPR-NEXT: v_mov_b32_e32 v18, s14
-; AGPR-NEXT: v_mov_b32_e32 v19, s15
-; AGPR-NEXT: global_store_dwordx4 v32, v[16:19], s[0:1] offset:16 sc0 sc1
-; AGPR-NEXT: s_waitcnt vmcnt(0)
-; AGPR-NEXT: s_nop 0
-; AGPR-NEXT: v_mov_b32_e32 v16, s8
-; AGPR-NEXT: v_mov_b32_e32 v17, s9
-; AGPR-NEXT: v_mov_b32_e32 v18, s10
-; AGPR-NEXT: v_mov_b32_e32 v19, s11
-; AGPR-NEXT: global_store_dwordx4 v32, v[16:19], s[0:1] sc0 sc1
-; AGPR-NEXT: s_waitcnt vmcnt(0)
-; AGPR-NEXT: global_store_dwordx4 v32, v[8:11], s[0:1] offset:32 sc0 sc1
-; AGPR-NEXT: s_waitcnt vmcnt(0)
-; AGPR-NEXT: global_store_dwordx4 v32, v[12:15], s[0:1] offset:48 sc0 sc1
-; AGPR-NEXT: s_waitcnt vmcnt(0)
-; AGPR-NEXT: global_store_dwordx4 v32, v[0:3], s[0:1] sc0 sc1
-; AGPR-NEXT: s_waitcnt vmcnt(0)
-; AGPR-NEXT: global_store_dwordx4 v32, v[4:7], s[0:1] offset:16 sc0 sc1
-; AGPR-NEXT: s_waitcnt vmcnt(0)
-; AGPR-NEXT: s_endpgm
+; SDAG-AGPR-LABEL: test_mfma_f32_32x32x16_bf16__vgprcd:
+; SDAG-AGPR: ; %bb.0:
+; SDAG-AGPR-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x24
+; SDAG-AGPR-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x64
+; SDAG-AGPR-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0xa4
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v32, 0
+; SDAG-AGPR-NEXT: s_waitcnt lgkmcnt(0)
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[36:37], s[26:27]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[34:35], s[24:25]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[40:41], s[30:31]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[38:39], s[28:29]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[30:31], s[22:23]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[28:29], s[20:21]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[26:27], s[18:19]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[24:25], s[16:17]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[22:23], s[14:15]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[20:21], s[12:13]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[18:19], s[10:11]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[16:17], s[8:9]
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v42, s20
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v43, s21
+; SDAG-AGPR-NEXT: v_mfma_f32_32x32x16_bf16 v[0:15], v[34:37], v[38:41], v[16:31]
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v44, s22
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v45, s23
+; SDAG-AGPR-NEXT: global_store_dwordx4 v32, v[42:45], s[0:1] offset:48 sc0 sc1
+; SDAG-AGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-AGPR-NEXT: s_nop 2
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v16, s16
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v17, s17
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v18, s18
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v19, s19
+; SDAG-AGPR-NEXT: global_store_dwordx4 v32, v[16:19], s[0:1] offset:32 sc0 sc1
+; SDAG-AGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-AGPR-NEXT: s_nop 0
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v16, s12
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v17, s13
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v18, s14
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v19, s15
+; SDAG-AGPR-NEXT: global_store_dwordx4 v32, v[16:19], s[0:1] offset:16 sc0 sc1
+; SDAG-AGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-AGPR-NEXT: s_nop 0
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v16, s8
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v17, s9
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v18, s10
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v19, s11
+; SDAG-AGPR-NEXT: global_store_dwordx4 v32, v[16:19], s[0:1] sc0 sc1
+; SDAG-AGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-AGPR-NEXT: global_store_dwordx4 v32, v[8:11], s[0:1] offset:32 sc0 sc1
+; SDAG-AGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-AGPR-NEXT: global_store_dwordx4 v32, v[12:15], s[0:1] offset:48 sc0 sc1
+; SDAG-AGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-AGPR-NEXT: global_store_dwordx4 v32, v[0:3], s[0:1] sc0 sc1
+; SDAG-AGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-AGPR-NEXT: global_store_dwordx4 v32, v[4:7], s[0:1] offset:16 sc0 sc1
+; SDAG-AGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-AGPR-NEXT: s_endpgm
;
-; VGPR-LABEL: test_mfma_f32_32x32x16_bf16__vgprcd:
-; VGPR: ; %bb.0:
-; VGPR-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x24
-; VGPR-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x64
-; VGPR-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0xa4
-; VGPR-NEXT: v_mov_b32_e32 v32, 0
-; VGPR-NEXT: s_waitcnt lgkmcnt(0)
-; VGPR-NEXT: v_mov_b64_e32 v[36:37], s[26:27]
-; VGPR-NEXT: v_mov_b64_e32 v[34:35], s[24:25]
-; VGPR-NEXT: v_mov_b64_e32 v[40:41], s[30:31]
-; VGPR-NEXT: v_mov_b64_e32 v[38:39], s[28:29]
-; VGPR-NEXT: v_mov_b64_e32 v[30:31], s[22:23]
-; VGPR-NEXT: v_mov_b64_e32 v[28:29], s[20:21]
-; VGPR-NEXT: v_mov_b64_e32 v[26:27], s[18:19]
-; VGPR-NEXT: v_mov_b64_e32 v[24:25], s[16:17]
-; VGPR-NEXT: v_mov_b64_e32 v[22:23], s[14:15]
-; VGPR-NEXT: v_mov_b64_e32 v[20:21], s[12:13]
-; VGPR-NEXT: v_mov_b64_e32 v[18:19], s[10:11]
-; VGPR-NEXT: v_mov_b64_e32 v[16:17], s[8:9]
-; VGPR-NEXT: v_mov_b32_e32 v42, s20
-; VGPR-NEXT: v_mov_b32_e32 v43, s21
-; VGPR-NEXT: v_mfma_f32_32x32x16_bf16 v[0:15], v[34:37], v[38:41], v[16:31]
-; VGPR-NEXT: v_mov_b32_e32 v44, s22
-; VGPR-NEXT: v_mov_b32_e32 v45, s23
-; VGPR-NEXT: global_store_dwordx4 v32, v[42:45], s[0:1] offset:48 sc0 sc1
-; VGPR-NEXT: s_waitcnt vmcnt(0)
-; VGPR-NEXT: s_nop 2
-; VGPR-NEXT: v_mov_b32_e32 v16, s16
-; VGPR-NEXT: v_mov_b32_e32 v17, s17
-; VGPR-NEXT: v_mov_b32_e32 v18, s18
-; VGPR-NEXT: v_mov_b32_e32 v19, s19
-; VGPR-NEXT: global_store_dwordx4 v32, v[16:19], s[0:1] offset:32 sc0 sc1
-; VGPR-NEXT: s_waitcnt vmcnt(0)
-; VGPR-NEXT: s_nop 0
-; VGPR-NEXT: v_mov_b32_e32 v16, s12
-; VGPR-NEXT: v_mov_b32_e32 v17, s13
-; VGPR-NEXT: v_mov_b32_e32 v18, s14
-; VGPR-NEXT: v_mov_b32_e32 v19, s15
-; VGPR-NEXT: global_store_dwordx4 v32, v[16:19], s[0:1] offset:16 sc0 sc1
-; VGPR-NEXT: s_waitcnt vmcnt(0)
-; VGPR-NEXT: s_nop 0
-; VGPR-NEXT: v_mov_b32_e32 v16, s8
-; VGPR-NEXT: v_mov_b32_e32 v17, s9
-; VGPR-NEXT: v_mov_b32_e32 v18, s10
-; VGPR-NEXT: v_mov_b32_e32 v19, s11
-; VGPR-NEXT: global_store_dwordx4 v32, v[16:19], s[0:1] sc0 sc1
-; VGPR-NEXT: s_waitcnt vmcnt(0)
-; VGPR-NEXT: global_store_dwordx4 v32, v[8:11], s[0:1] offset:32 sc0 sc1
-; VGPR-NEXT: s_waitcnt vmcnt(0)
-; VGPR-NEXT: global_store_dwordx4 v32, v[12:15], s[0:1] offset:48 sc0 sc1
-; VGPR-NEXT: s_waitcnt vmcnt(0)
-; VGPR-NEXT: global_store_dwordx4 v32, v[0:3], s[0:1] sc0 sc1
-; VGPR-NEXT: s_waitcnt vmcnt(0)
-; VGPR-NEXT: global_store_dwordx4 v32, v[4:7], s[0:1] offset:16 sc0 sc1
-; VGPR-NEXT: s_waitcnt vmcnt(0)
-; VGPR-NEXT: s_endpgm
+; GISEL-AGPR-LABEL: test_mfma_f32_32x32x16_bf16__vgprcd:
+; GISEL-AGPR: ; %bb.0:
+; GISEL-AGPR-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x24
+; GISEL-AGPR-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x64
+; GISEL-AGPR-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0xa4
+; GISEL-AGPR-NEXT: v_mov_b32_e32 v56, 0
+; GISEL-AGPR-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[34:35], s[26:27]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[32:33], s[24:25]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[38:39], s[30:31]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[36:37], s[28:29]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[0:1], s[8:9]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[2:3], s[10:11]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[4:5], s[12:13]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[6:7], s[14:15]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[8:9], s[16:17]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[10:11], s[18:19]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[12:13], s[20:21]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[14:15], s[22:23]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[42:43], s[10:11]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[40:41], s[8:9]
+; GISEL-AGPR-NEXT: v_mfma_f32_32x32x16_bf16 v[16:31], v[32:35], v[36:39], v[0:15]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[46:47], s[14:15]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[44:45], s[12:13]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[50:51], s[18:19]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[48:49], s[16:17]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[54:55], s[22:23]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[52:53], s[20:21]
+; GISEL-AGPR-NEXT: global_store_dwordx4 v56, v[40:43], s[0:1] sc0 sc1
+; GISEL-AGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-AGPR-NEXT: global_store_dwordx4 v56, v[44:47], s[0:1] offset:16 sc0 sc1
+; GISEL-AGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-AGPR-NEXT: global_store_dwordx4 v56, v[48:51], s[0:1] offset:32 sc0 sc1
+; GISEL-AGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-AGPR-NEXT: global_store_dwordx4 v56, v[52:55], s[0:1] offset:48 sc0 sc1
+; GISEL-AGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-AGPR-NEXT: global_store_dwordx4 v56, v[16:19], s[0:1] sc0 sc1
+; GISEL-AGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-AGPR-NEXT: global_store_dwordx4 v56, v[20:23], s[0:1] offset:16 sc0 sc1
+; GISEL-AGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-AGPR-NEXT: global_store_dwordx4 v56, v[24:27], s[0:1] offset:32 sc0 sc1
+; GISEL-AGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-AGPR-NEXT: global_store_dwordx4 v56, v[28:31], s[0:1] offset:48 sc0 sc1
+; GISEL-AGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-AGPR-NEXT: s_endpgm
+;
+; SDAG-VGPR-LABEL: test_mfma_f32_32x32x16_bf16__vgprcd:
+; SDAG-VGPR: ; %bb.0:
+; SDAG-VGPR-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x24
+; SDAG-VGPR-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x64
+; SDAG-VGPR-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0xa4
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v32, 0
+; SDAG-VGPR-NEXT: s_waitcnt lgkmcnt(0)
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[36:37], s[26:27]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[34:35], s[24:25]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[40:41], s[30:31]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[38:39], s[28:29]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[30:31], s[22:23]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[28:29], s[20:21]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[26:27], s[18:19]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[24:25], s[16:17]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[22:23], s[14:15]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[20:21], s[12:13]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[18:19], s[10:11]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[16:17], s[8:9]
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v42, s20
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v43, s21
+; SDAG-VGPR-NEXT: v_mfma_f32_32x32x16_bf16 v[0:15], v[34:37], v[38:41], v[16:31]
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v44, s22
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v45, s23
+; SDAG-VGPR-NEXT: global_store_dwordx4 v32, v[42:45], s[0:1] offset:48 sc0 sc1
+; SDAG-VGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-VGPR-NEXT: s_nop 2
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v16, s16
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v17, s17
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v18, s18
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v19, s19
+; SDAG-VGPR-NEXT: global_store_dwordx4 v32, v[16:19], s[0:1] offset:32 sc0 sc1
+; SDAG-VGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-VGPR-NEXT: s_nop 0
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v16, s12
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v17, s13
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v18, s14
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v19, s15
+; SDAG-VGPR-NEXT: global_store_dwordx4 v32, v[16:19], s[0:1] offset:16 sc0 sc1
+; SDAG-VGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-VGPR-NEXT: s_nop 0
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v16, s8
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v17, s9
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v18, s10
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v19, s11
+; SDAG-VGPR-NEXT: global_store_dwordx4 v32, v[16:19], s[0:1] sc0 sc1
+; SDAG-VGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-VGPR-NEXT: global_store_dwordx4 v32, v[8:11], s[0:1] offset:32 sc0 sc1
+; SDAG-VGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-VGPR-NEXT: global_store_dwordx4 v32, v[12:15], s[0:1] offset:48 sc0 sc1
+; SDAG-VGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-VGPR-NEXT: global_store_dwordx4 v32, v[0:3], s[0:1] sc0 sc1
+; SDAG-VGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-VGPR-NEXT: global_store_dwordx4 v32, v[4:7], s[0:1] offset:16 sc0 sc1
+; SDAG-VGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-VGPR-NEXT: s_endpgm
+;
+; GISEL-VGPR-LABEL: test_mfma_f32_32x32x16_bf16__vgprcd:
+; GISEL-VGPR: ; %bb.0:
+; GISEL-VGPR-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x24
+; GISEL-VGPR-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x64
+; GISEL-VGPR-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0xa4
+; GISEL-VGPR-NEXT: v_mov_b32_e32 v56, 0
+; GISEL-VGPR-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[34:35], s[26:27]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[32:33], s[24:25]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[38:39], s[30:31]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[36:37], s[28:29]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[0:1], s[8:9]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[2:3], s[10:11]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[4:5], s[12:13]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[6:7], s[14:15]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[8:9], s[16:17]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[10:11], s[18:19]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[12:13], s[20:21]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[14:15], s[22:23]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[42:43], s[10:11]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[40:41], s[8:9]
+; GISEL-VGPR-NEXT: v_mfma_f32_32x32x16_bf16 v[16:31], v[32:35], v[36:39], v[0:15]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[46:47], s[14:15]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[44:45], s[12:13]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[50:51], s[18:19]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[48:49], s[16:17]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[54:55], s[22:23]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[52:53], s[20:21]
+; GISEL-VGPR-NEXT: global_store_dwordx4 v56, v[40:43], s[0:1] sc0 sc1
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-VGPR-NEXT: global_store_dwordx4 v56, v[44:47], s[0:1] offset:16 sc0 sc1
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-VGPR-NEXT: global_store_dwordx4 v56, v[48:51], s[0:1] offset:32 sc0 sc1
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-VGPR-NEXT: global_store_dwordx4 v56, v[52:55], s[0:1] offset:48 sc0 sc1
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-VGPR-NEXT: global_store_dwordx4 v56, v[16:19], s[0:1] sc0 sc1
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-VGPR-NEXT: global_store_dwordx4 v56, v[20:23], s[0:1] offset:16 sc0 sc1
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-VGPR-NEXT: global_store_dwordx4 v56, v[24:27], s[0:1] offset:32 sc0 sc1
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-VGPR-NEXT: global_store_dwordx4 v56, v[28:31], s[0:1] offset:48 sc0 sc1
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-VGPR-NEXT: s_endpgm
%result = call <16 x float> @llvm.amdgcn.mfma.f32.32x32x16.bf16(<8 x bfloat> %arg0, <8 x bfloat> %arg1, <16 x float> %arg2, i32 0, i32 0, i32 0)
store volatile <16 x float> %arg2, ptr addrspace(1) %out
store volatile <16 x float> %result, ptr addrspace(1) %out
@@ -854,119 +1197,211 @@ define amdgpu_kernel void @test_mfma_f32_32x32x16_bf16__vgprcd__flags(<8 x bfloa
; GCN-NEXT: s_waitcnt vmcnt(0)
; GCN-NEXT: s_endpgm
;
-; AGPR-LABEL: test_mfma_f32_32x32x16_bf16__vgprcd__flags:
-; AGPR: ; %bb.0:
-; AGPR-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x24
-; AGPR-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x64
-; AGPR-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0xa4
-; AGPR-NEXT: v_mov_b32_e32 v32, 0
-; AGPR-NEXT: s_waitcnt lgkmcnt(0)
-; AGPR-NEXT: v_mov_b64_e32 v[36:37], s[26:27]
-; AGPR-NEXT: v_mov_b64_e32 v[34:35], s[24:25]
-; AGPR-NEXT: v_mov_b64_e32 v[40:41], s[30:31]
-; AGPR-NEXT: v_mov_b64_e32 v[38:39], s[28:29]
-; AGPR-NEXT: v_mov_b64_e32 v[30:31], s[22:23]
-; AGPR-NEXT: v_mov_b64_e32 v[28:29], s[20:21]
-; AGPR-NEXT: v_mov_b64_e32 v[26:27], s[18:19]
-; AGPR-NEXT: v_mov_b64_e32 v[24:25], s[16:17]
-; AGPR-NEXT: v_mov_b64_e32 v[22:23], s[14:15]
-; AGPR-NEXT: v_mov_b64_e32 v[20:21], s[12:13]
-; AGPR-NEXT: v_mov_b64_e32 v[18:19], s[10:11]
-; AGPR-NEXT: v_mov_b64_e32 v[16:17], s[8:9]
-; AGPR-NEXT: v_mov_b32_e32 v42, s20
-; AGPR-NEXT: v_mov_b32_e32 v43, s21
-; AGPR-NEXT: v_mfma_f32_32x32x16_bf16 v[0:15], v[34:37], v[38:41], v[16:31] cbsz:1 abid:2 blgp:3
-; AGPR-NEXT: v_mov_b32_e32 v44, s22
-; AGPR-NEXT: v_mov_b32_e32 v45, s23
-; AGPR-NEXT: global_store_dwordx4 v32, v[42:45], s[0:1] offset:48 sc0 sc1
-; AGPR-NEXT: s_waitcnt vmcnt(0)
-; AGPR-NEXT: s_nop 2
-; AGPR-NEXT: v_mov_b32_e32 v16, s16
-; AGPR-NEXT: v_mov_b32_e32 v17, s17
-; AGPR-NEXT: v_mov_b32_e32 v18, s18
-; AGPR-NEXT: v_mov_b32_e32 v19, s19
-; AGPR-NEXT: global_store_dwordx4 v32, v[16:19], s[0:1] offset:32 sc0 sc1
-; AGPR-NEXT: s_waitcnt vmcnt(0)
-; AGPR-NEXT: s_nop 0
-; AGPR-NEXT: v_mov_b32_e32 v16, s12
-; AGPR-NEXT: v_mov_b32_e32 v17, s13
-; AGPR-NEXT: v_mov_b32_e32 v18, s14
-; AGPR-NEXT: v_mov_b32_e32 v19, s15
-; AGPR-NEXT: global_store_dwordx4 v32, v[16:19], s[0:1] offset:16 sc0 sc1
-; AGPR-NEXT: s_waitcnt vmcnt(0)
-; AGPR-NEXT: s_nop 0
-; AGPR-NEXT: v_mov_b32_e32 v16, s8
-; AGPR-NEXT: v_mov_b32_e32 v17, s9
-; AGPR-NEXT: v_mov_b32_e32 v18, s10
-; AGPR-NEXT: v_mov_b32_e32 v19, s11
-; AGPR-NEXT: global_store_dwordx4 v32, v[16:19], s[0:1] sc0 sc1
-; AGPR-NEXT: s_waitcnt vmcnt(0)
-; AGPR-NEXT: global_store_dwordx4 v32, v[8:11], s[0:1] offset:32 sc0 sc1
-; AGPR-NEXT: s_waitcnt vmcnt(0)
-; AGPR-NEXT: global_store_dwordx4 v32, v[12:15], s[0:1] offset:48 sc0 sc1
-; AGPR-NEXT: s_waitcnt vmcnt(0)
-; AGPR-NEXT: global_store_dwordx4 v32, v[0:3], s[0:1] sc0 sc1
-; AGPR-NEXT: s_waitcnt vmcnt(0)
-; AGPR-NEXT: global_store_dwordx4 v32, v[4:7], s[0:1] offset:16 sc0 sc1
-; AGPR-NEXT: s_waitcnt vmcnt(0)
-; AGPR-NEXT: s_endpgm
+; SDAG-AGPR-LABEL: test_mfma_f32_32x32x16_bf16__vgprcd__flags:
+; SDAG-AGPR: ; %bb.0:
+; SDAG-AGPR-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x24
+; SDAG-AGPR-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x64
+; SDAG-AGPR-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0xa4
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v32, 0
+; SDAG-AGPR-NEXT: s_waitcnt lgkmcnt(0)
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[36:37], s[26:27]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[34:35], s[24:25]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[40:41], s[30:31]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[38:39], s[28:29]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[30:31], s[22:23]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[28:29], s[20:21]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[26:27], s[18:19]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[24:25], s[16:17]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[22:23], s[14:15]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[20:21], s[12:13]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[18:19], s[10:11]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[16:17], s[8:9]
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v42, s20
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v43, s21
+; SDAG-AGPR-NEXT: v_mfma_f32_32x32x16_bf16 v[0:15], v[34:37], v[38:41], v[16:31] cbsz:1 abid:2 blgp:3
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v44, s22
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v45, s23
+; SDAG-AGPR-NEXT: global_store_dwordx4 v32, v[42:45], s[0:1] offset:48 sc0 sc1
+; SDAG-AGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-AGPR-NEXT: s_nop 2
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v16, s16
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v17, s17
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v18, s18
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v19, s19
+; SDAG-AGPR-NEXT: global_store_dwordx4 v32, v[16:19], s[0:1] offset:32 sc0 sc1
+; SDAG-AGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-AGPR-NEXT: s_nop 0
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v16, s12
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v17, s13
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v18, s14
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v19, s15
+; SDAG-AGPR-NEXT: global_store_dwordx4 v32, v[16:19], s[0:1] offset:16 sc0 sc1
+; SDAG-AGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-AGPR-NEXT: s_nop 0
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v16, s8
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v17, s9
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v18, s10
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v19, s11
+; SDAG-AGPR-NEXT: global_store_dwordx4 v32, v[16:19], s[0:1] sc0 sc1
+; SDAG-AGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-AGPR-NEXT: global_store_dwordx4 v32, v[8:11], s[0:1] offset:32 sc0 sc1
+; SDAG-AGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-AGPR-NEXT: global_store_dwordx4 v32, v[12:15], s[0:1] offset:48 sc0 sc1
+; SDAG-AGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-AGPR-NEXT: global_store_dwordx4 v32, v[0:3], s[0:1] sc0 sc1
+; SDAG-AGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-AGPR-NEXT: global_store_dwordx4 v32, v[4:7], s[0:1] offset:16 sc0 sc1
+; SDAG-AGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-AGPR-NEXT: s_endpgm
;
-; VGPR-LABEL: test_mfma_f32_32x32x16_bf16__vgprcd__flags:
-; VGPR: ; %bb.0:
-; VGPR-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x24
-; VGPR-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x64
-; VGPR-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0xa4
-; VGPR-NEXT: v_mov_b32_e32 v32, 0
-; VGPR-NEXT: s_waitcnt lgkmcnt(0)
-; VGPR-NEXT: v_mov_b64_e32 v[36:37], s[26:27]
-; VGPR-NEXT: v_mov_b64_e32 v[34:35], s[24:25]
-; VGPR-NEXT: v_mov_b64_e32 v[40:41], s[30:31]
-; VGPR-NEXT: v_mov_b64_e32 v[38:39], s[28:29]
-; VGPR-NEXT: v_mov_b64_e32 v[30:31], s[22:23]
-; VGPR-NEXT: v_mov_b64_e32 v[28:29], s[20:21]
-; VGPR-NEXT: v_mov_b64_e32 v[26:27], s[18:19]
-; VGPR-NEXT: v_mov_b64_e32 v[24:25], s[16:17]
-; VGPR-NEXT: v_mov_b64_e32 v[22:23], s[14:15]
-; VGPR-NEXT: v_mov_b64_e32 v[20:21], s[12:13]
-; VGPR-NEXT: v_mov_b64_e32 v[18:19], s[10:11]
-; VGPR-NEXT: v_mov_b64_e32 v[16:17], s[8:9]
-; VGPR-NEXT: v_mov_b32_e32 v42, s20
-; VGPR-NEXT: v_mov_b32_e32 v43, s21
-; VGPR-NEXT: v_mfma_f32_32x32x16_bf16 v[0:15], v[34:37], v[38:41], v[16:31] cbsz:1 abid:2 blgp:3
-; VGPR-NEXT: v_mov_b32_e32 v44, s22
-; VGPR-NEXT: v_mov_b32_e32 v45, s23
-; VGPR-NEXT: global_store_dwordx4 v32, v[42:45], s[0:1] offset:48 sc0 sc1
-; VGPR-NEXT: s_waitcnt vmcnt(0)
-; VGPR-NEXT: s_nop 2
-; VGPR-NEXT: v_mov_b32_e32 v16, s16
-; VGPR-NEXT: v_mov_b32_e32 v17, s17
-; VGPR-NEXT: v_mov_b32_e32 v18, s18
-; VGPR-NEXT: v_mov_b32_e32 v19, s19
-; VGPR-NEXT: global_store_dwordx4 v32, v[16:19], s[0:1] offset:32 sc0 sc1
-; VGPR-NEXT: s_waitcnt vmcnt(0)
-; VGPR-NEXT: s_nop 0
-; VGPR-NEXT: v_mov_b32_e32 v16, s12
-; VGPR-NEXT: v_mov_b32_e32 v17, s13
-; VGPR-NEXT: v_mov_b32_e32 v18, s14
-; VGPR-NEXT: v_mov_b32_e32 v19, s15
-; VGPR-NEXT: global_store_dwordx4 v32, v[16:19], s[0:1] offset:16 sc0 sc1
-; VGPR-NEXT: s_waitcnt vmcnt(0)
-; VGPR-NEXT: s_nop 0
-; VGPR-NEXT: v_mov_b32_e32 v16, s8
-; VGPR-NEXT: v_mov_b32_e32 v17, s9
-; VGPR-NEXT: v_mov_b32_e32 v18, s10
-; VGPR-NEXT: v_mov_b32_e32 v19, s11
-; VGPR-NEXT: global_store_dwordx4 v32, v[16:19], s[0:1] sc0 sc1
-; VGPR-NEXT: s_waitcnt vmcnt(0)
-; VGPR-NEXT: global_store_dwordx4 v32, v[8:11], s[0:1] offset:32 sc0 sc1
-; VGPR-NEXT: s_waitcnt vmcnt(0)
-; VGPR-NEXT: global_store_dwordx4 v32, v[12:15], s[0:1] offset:48 sc0 sc1
-; VGPR-NEXT: s_waitcnt vmcnt(0)
-; VGPR-NEXT: global_store_dwordx4 v32, v[0:3], s[0:1] sc0 sc1
-; VGPR-NEXT: s_waitcnt vmcnt(0)
-; VGPR-NEXT: global_store_dwordx4 v32, v[4:7], s[0:1] offset:16 sc0 sc1
-; VGPR-NEXT: s_waitcnt vmcnt(0)
-; VGPR-NEXT: s_endpgm
+; GISEL-AGPR-LABEL: test_mfma_f32_32x32x16_bf16__vgprcd__flags:
+; GISEL-AGPR: ; %bb.0:
+; GISEL-AGPR-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x24
+; GISEL-AGPR-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x64
+; GISEL-AGPR-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0xa4
+; GISEL-AGPR-NEXT: v_mov_b32_e32 v56, 0
+; GISEL-AGPR-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[34:35], s[26:27]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[32:33], s[24:25]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[38:39], s[30:31]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[36:37], s[28:29]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[0:1], s[8:9]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[2:3], s[10:11]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[4:5], s[12:13]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[6:7], s[14:15]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[8:9], s[16:17]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[10:11], s[18:19]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[12:13], s[20:21]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[14:15], s[22:23]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[42:43], s[10:11]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[40:41], s[8:9]
+; GISEL-AGPR-NEXT: v_mfma_f32_32x32x16_bf16 v[16:31], v[32:35], v[36:39], v[0:15] cbsz:1 abid:2 blgp:3
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[46:47], s[14:15]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[44:45], s[12:13]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[50:51], s[18:19]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[48:49], s[16:17]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[54:55], s[22:23]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[52:53], s[20:21]
+; GISEL-AGPR-NEXT: global_store_dwordx4 v56, v[40:43], s[0:1] sc0 sc1
+; GISEL-AGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-AGPR-NEXT: global_store_dwordx4 v56, v[44:47], s[0:1] offset:16 sc0 sc1
+; GISEL-AGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-AGPR-NEXT: global_store_dwordx4 v56, v[48:51], s[0:1] offset:32 sc0 sc1
+; GISEL-AGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-AGPR-NEXT: global_store_dwordx4 v56, v[52:55], s[0:1] offset:48 sc0 sc1
+; GISEL-AGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-AGPR-NEXT: global_store_dwordx4 v56, v[16:19], s[0:1] sc0 sc1
+; GISEL-AGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-AGPR-NEXT: global_store_dwordx4 v56, v[20:23], s[0:1] offset:16 sc0 sc1
+; GISEL-AGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-AGPR-NEXT: global_store_dwordx4 v56, v[24:27], s[0:1] offset:32 sc0 sc1
+; GISEL-AGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-AGPR-NEXT: global_store_dwordx4 v56, v[28:31], s[0:1] offset:48 sc0 sc1
+; GISEL-AGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-AGPR-NEXT: s_endpgm
+;
+; SDAG-VGPR-LABEL: test_mfma_f32_32x32x16_bf16__vgprcd__flags:
+; SDAG-VGPR: ; %bb.0:
+; SDAG-VGPR-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x24
+; SDAG-VGPR-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x64
+; SDAG-VGPR-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0xa4
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v32, 0
+; SDAG-VGPR-NEXT: s_waitcnt lgkmcnt(0)
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[36:37], s[26:27]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[34:35], s[24:25]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[40:41], s[30:31]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[38:39], s[28:29]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[30:31], s[22:23]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[28:29], s[20:21]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[26:27], s[18:19]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[24:25], s[16:17]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[22:23], s[14:15]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[20:21], s[12:13]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[18:19], s[10:11]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[16:17], s[8:9]
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v42, s20
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v43, s21
+; SDAG-VGPR-NEXT: v_mfma_f32_32x32x16_bf16 v[0:15], v[34:37], v[38:41], v[16:31] cbsz:1 abid:2 blgp:3
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v44, s22
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v45, s23
+; SDAG-VGPR-NEXT: global_store_dwordx4 v32, v[42:45], s[0:1] offset:48 sc0 sc1
+; SDAG-VGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-VGPR-NEXT: s_nop 2
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v16, s16
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v17, s17
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v18, s18
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v19, s19
+; SDAG-VGPR-NEXT: global_store_dwordx4 v32, v[16:19], s[0:1] offset:32 sc0 sc1
+; SDAG-VGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-VGPR-NEXT: s_nop 0
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v16, s12
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v17, s13
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v18, s14
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v19, s15
+; SDAG-VGPR-NEXT: global_store_dwordx4 v32, v[16:19], s[0:1] offset:16 sc0 sc1
+; SDAG-VGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-VGPR-NEXT: s_nop 0
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v16, s8
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v17, s9
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v18, s10
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v19, s11
+; SDAG-VGPR-NEXT: global_store_dwordx4 v32, v[16:19], s[0:1] sc0 sc1
+; SDAG-VGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-VGPR-NEXT: global_store_dwordx4 v32, v[8:11], s[0:1] offset:32 sc0 sc1
+; SDAG-VGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-VGPR-NEXT: global_store_dwordx4 v32, v[12:15], s[0:1] offset:48 sc0 sc1
+; SDAG-VGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-VGPR-NEXT: global_store_dwordx4 v32, v[0:3], s[0:1] sc0 sc1
+; SDAG-VGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-VGPR-NEXT: global_store_dwordx4 v32, v[4:7], s[0:1] offset:16 sc0 sc1
+; SDAG-VGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-VGPR-NEXT: s_endpgm
+;
+; GISEL-VGPR-LABEL: test_mfma_f32_32x32x16_bf16__vgprcd__flags:
+; GISEL-VGPR: ; %bb.0:
+; GISEL-VGPR-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x24
+; GISEL-VGPR-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x64
+; GISEL-VGPR-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0xa4
+; GISEL-VGPR-NEXT: v_mov_b32_e32 v56, 0
+; GISEL-VGPR-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[34:35], s[26:27]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[32:33], s[24:25]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[38:39], s[30:31]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[36:37], s[28:29]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[0:1], s[8:9]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[2:3], s[10:11]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[4:5], s[12:13]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[6:7], s[14:15]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[8:9], s[16:17]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[10:11], s[18:19]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[12:13], s[20:21]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[14:15], s[22:23]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[42:43], s[10:11]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[40:41], s[8:9]
+; GISEL-VGPR-NEXT: v_mfma_f32_32x32x16_bf16 v[16:31], v[32:35], v[36:39], v[0:15] cbsz:1 abid:2 blgp:3
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[46:47], s[14:15]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[44:45], s[12:13]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[50:51], s[18:19]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[48:49], s[16:17]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[54:55], s[22:23]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[52:53], s[20:21]
+; GISEL-VGPR-NEXT: global_store_dwordx4 v56, v[40:43], s[0:1] sc0 sc1
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-VGPR-NEXT: global_store_dwordx4 v56, v[44:47], s[0:1] offset:16 sc0 sc1
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-VGPR-NEXT: global_store_dwordx4 v56, v[48:51], s[0:1] offset:32 sc0 sc1
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-VGPR-NEXT: global_store_dwordx4 v56, v[52:55], s[0:1] offset:48 sc0 sc1
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-VGPR-NEXT: global_store_dwordx4 v56, v[16:19], s[0:1] sc0 sc1
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-VGPR-NEXT: global_store_dwordx4 v56, v[20:23], s[0:1] offset:16 sc0 sc1
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-VGPR-NEXT: global_store_dwordx4 v56, v[24:27], s[0:1] offset:32 sc0 sc1
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-VGPR-NEXT: global_store_dwordx4 v56, v[28:31], s[0:1] offset:48 sc0 sc1
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-VGPR-NEXT: s_endpgm
%result = call <16 x float> @llvm.amdgcn.mfma.f32.32x32x16.bf16(<8 x bfloat> %arg0, <8 x bfloat> %arg1, <16 x float> %arg2, i32 1, i32 2, i32 3)
store volatile <16 x float> %arg2, ptr addrspace(1) %out
store volatile <16 x float> %result, ptr addrspace(1) %out
@@ -1002,61 +1437,117 @@ define amdgpu_kernel void @test_mfma_f32_32x32x16_bf16__vgprcd_mac(<8 x bfloat>
; GCN-NEXT: global_store_dwordx4 v16, v[0:3], s[0:1]
; GCN-NEXT: s_endpgm
;
-; AGPR-LABEL: test_mfma_f32_32x32x16_bf16__vgprcd_mac:
-; AGPR: ; %bb.0:
-; AGPR-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x24
-; AGPR-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x64
-; AGPR-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0xa4
-; AGPR-NEXT: s_waitcnt lgkmcnt(0)
-; AGPR-NEXT: v_mov_b64_e32 v[16:17], s[24:25]
-; AGPR-NEXT: v_mov_b64_e32 v[18:19], s[26:27]
-; AGPR-NEXT: v_mov_b64_e32 v[20:21], s[28:29]
-; AGPR-NEXT: v_mov_b64_e32 v[22:23], s[30:31]
-; AGPR-NEXT: v_mov_b64_e32 v[0:1], s[8:9]
-; AGPR-NEXT: v_mov_b64_e32 v[2:3], s[10:11]
-; AGPR-NEXT: v_mov_b64_e32 v[4:5], s[12:13]
-; AGPR-NEXT: v_mov_b64_e32 v[6:7], s[14:15]
-; AGPR-NEXT: v_mov_b64_e32 v[8:9], s[16:17]
-; AGPR-NEXT: v_mov_b64_e32 v[10:11], s[18:19]
-; AGPR-NEXT: v_mov_b64_e32 v[12:13], s[20:21]
-; AGPR-NEXT: v_mov_b64_e32 v[14:15], s[22:23]
-; AGPR-NEXT: s_nop 1
-; AGPR-NEXT: v_mfma_f32_32x32x16_bf16 v[0:15], v[16:19], v[20:23], v[0:15]
-; AGPR-NEXT: v_mov_b32_e32 v16, 0
-; AGPR-NEXT: s_nop 10
-; AGPR-NEXT: global_store_dwordx4 v16, v[12:15], s[0:1] offset:48
-; AGPR-NEXT: global_store_dwordx4 v16, v[8:11], s[0:1] offset:32
-; AGPR-NEXT: global_store_dwordx4 v16, v[4:7], s[0:1] offset:16
-; AGPR-NEXT: global_store_dwordx4 v16, v[0:3], s[0:1]
-; AGPR-NEXT: s_endpgm
+; SDAG-AGPR-LABEL: test_mfma_f32_32x32x16_bf16__vgprcd_mac:
+; SDAG-AGPR: ; %bb.0:
+; SDAG-AGPR-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x24
+; SDAG-AGPR-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x64
+; SDAG-AGPR-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0xa4
+; SDAG-AGPR-NEXT: s_waitcnt lgkmcnt(0)
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[16:17], s[24:25]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[18:19], s[26:27]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[20:21], s[28:29]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[22:23], s[30:31]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[0:1], s[8:9]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[2:3], s[10:11]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[4:5], s[12:13]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[6:7], s[14:15]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[8:9], s[16:17]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[10:11], s[18:19]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[12:13], s[20:21]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[14:15], s[22:23]
+; SDAG-AGPR-NEXT: s_nop 1
+; SDAG-AGPR-NEXT: v_mfma_f32_32x32x16_bf16 v[0:15], v[16:19], v[20:23], v[0:15]
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v16, 0
+; SDAG-AGPR-NEXT: s_nop 10
+; SDAG-AGPR-NEXT: global_store_dwordx4 v16, v[12:15], s[0:1] offset:48
+; SDAG-AGPR-NEXT: global_store_dwordx4 v16, v[8:11], s[0:1] offset:32
+; SDAG-AGPR-NEXT: global_store_dwordx4 v16, v[4:7], s[0:1] offset:16
+; SDAG-AGPR-NEXT: global_store_dwordx4 v16, v[0:3], s[0:1]
+; SDAG-AGPR-NEXT: s_endpgm
;
-; VGPR-LABEL: test_mfma_f32_32x32x16_bf16__vgprcd_mac:
-; VGPR: ; %bb.0:
-; VGPR-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x24
-; VGPR-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x64
-; VGPR-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0xa4
-; VGPR-NEXT: s_waitcnt lgkmcnt(0)
-; VGPR-NEXT: v_mov_b64_e32 v[16:17], s[24:25]
-; VGPR-NEXT: v_mov_b64_e32 v[18:19], s[26:27]
-; VGPR-NEXT: v_mov_b64_e32 v[20:21], s[28:29]
-; VGPR-NEXT: v_mov_b64_e32 v[22:23], s[30:31]
-; VGPR-NEXT: v_mov_b64_e32 v[0:1], s[8:9]
-; VGPR-NEXT: v_mov_b64_e32 v[2:3], s[10:11]
-; VGPR-NEXT: v_mov_b64_e32 v[4:5], s[12:13]
-; VGPR-NEXT: v_mov_b64_e32 v[6:7], s[14:15]
-; VGPR-NEXT: v_mov_b64_e32 v[8:9], s[16:17]
-; VGPR-NEXT: v_mov_b64_e32 v[10:11], s[18:19]
-; VGPR-NEXT: v_mov_b64_e32 v[12:13], s[20:21]
-; VGPR-NEXT: v_mov_b64_e32 v[14:15], s[22:23]
-; VGPR-NEXT: s_nop 1
-; VGPR-NEXT: v_mfma_f32_32x32x16_bf16 v[0:15], v[16:19], v[20:23], v[0:15]
-; VGPR-NEXT: v_mov_b32_e32 v16, 0
-; VGPR-NEXT: s_nop 10
-; VGPR-NEXT: global_store_dwordx4 v16, v[12:15], s[0:1] offset:48
-; VGPR-NEXT: global_store_dwordx4 v16, v[8:11], s[0:1] offset:32
-; VGPR-NEXT: global_store_dwordx4 v16, v[4:7], s[0:1] offset:16
-; VGPR-NEXT: global_store_dwordx4 v16, v[0:3], s[0:1]
-; VGPR-NEXT: s_endpgm
+; GISEL-AGPR-LABEL: test_mfma_f32_32x32x16_bf16__vgprcd_mac:
+; GISEL-AGPR: ; %bb.0:
+; GISEL-AGPR-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x24
+; GISEL-AGPR-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x64
+; GISEL-AGPR-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0xa4
+; GISEL-AGPR-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[16:17], s[24:25]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[18:19], s[26:27]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[20:21], s[28:29]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[22:23], s[30:31]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[0:1], s[8:9]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[2:3], s[10:11]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[4:5], s[12:13]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[6:7], s[14:15]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[8:9], s[16:17]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[10:11], s[18:19]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[12:13], s[20:21]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[14:15], s[22:23]
+; GISEL-AGPR-NEXT: s_nop 1
+; GISEL-AGPR-NEXT: v_mfma_f32_32x32x16_bf16 v[0:15], v[16:19], v[20:23], v[0:15]
+; GISEL-AGPR-NEXT: v_mov_b32_e32 v16, 0
+; GISEL-AGPR-NEXT: s_nop 10
+; GISEL-AGPR-NEXT: global_store_dwordx4 v16, v[0:3], s[0:1]
+; GISEL-AGPR-NEXT: global_store_dwordx4 v16, v[4:7], s[0:1] offset:16
+; GISEL-AGPR-NEXT: global_store_dwordx4 v16, v[8:11], s[0:1] offset:32
+; GISEL-AGPR-NEXT: global_store_dwordx4 v16, v[12:15], s[0:1] offset:48
+; GISEL-AGPR-NEXT: s_endpgm
+;
+; SDAG-VGPR-LABEL: test_mfma_f32_32x32x16_bf16__vgprcd_mac:
+; SDAG-VGPR: ; %bb.0:
+; SDAG-VGPR-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x24
+; SDAG-VGPR-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x64
+; SDAG-VGPR-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0xa4
+; SDAG-VGPR-NEXT: s_waitcnt lgkmcnt(0)
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[16:17], s[24:25]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[18:19], s[26:27]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[20:21], s[28:29]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[22:23], s[30:31]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[0:1], s[8:9]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[2:3], s[10:11]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[4:5], s[12:13]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[6:7], s[14:15]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[8:9], s[16:17]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[10:11], s[18:19]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[12:13], s[20:21]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[14:15], s[22:23]
+; SDAG-VGPR-NEXT: s_nop 1
+; SDAG-VGPR-NEXT: v_mfma_f32_32x32x16_bf16 v[0:15], v[16:19], v[20:23], v[0:15]
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v16, 0
+; SDAG-VGPR-NEXT: s_nop 10
+; SDAG-VGPR-NEXT: global_store_dwordx4 v16, v[12:15], s[0:1] offset:48
+; SDAG-VGPR-NEXT: global_store_dwordx4 v16, v[8:11], s[0:1] offset:32
+; SDAG-VGPR-NEXT: global_store_dwordx4 v16, v[4:7], s[0:1] offset:16
+; SDAG-VGPR-NEXT: global_store_dwordx4 v16, v[0:3], s[0:1]
+; SDAG-VGPR-NEXT: s_endpgm
+;
+; GISEL-VGPR-LABEL: test_mfma_f32_32x32x16_bf16__vgprcd_mac:
+; GISEL-VGPR: ; %bb.0:
+; GISEL-VGPR-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x24
+; GISEL-VGPR-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x64
+; GISEL-VGPR-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0xa4
+; GISEL-VGPR-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[16:17], s[24:25]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[18:19], s[26:27]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[20:21], s[28:29]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[22:23], s[30:31]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[0:1], s[8:9]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[2:3], s[10:11]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[4:5], s[12:13]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[6:7], s[14:15]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[8:9], s[16:17]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[10:11], s[18:19]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[12:13], s[20:21]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[14:15], s[22:23]
+; GISEL-VGPR-NEXT: s_nop 1
+; GISEL-VGPR-NEXT: v_mfma_f32_32x32x16_bf16 v[0:15], v[16:19], v[20:23], v[0:15]
+; GISEL-VGPR-NEXT: v_mov_b32_e32 v16, 0
+; GISEL-VGPR-NEXT: s_nop 10
+; GISEL-VGPR-NEXT: global_store_dwordx4 v16, v[0:3], s[0:1]
+; GISEL-VGPR-NEXT: global_store_dwordx4 v16, v[4:7], s[0:1] offset:16
+; GISEL-VGPR-NEXT: global_store_dwordx4 v16, v[8:11], s[0:1] offset:32
+; GISEL-VGPR-NEXT: global_store_dwordx4 v16, v[12:15], s[0:1] offset:48
+; GISEL-VGPR-NEXT: s_endpgm
%result = call <16 x float> @llvm.amdgcn.mfma.f32.32x32x16.bf16(<8 x bfloat> %arg0, <8 x bfloat> %arg1, <16 x float> %arg2, i32 0, i32 0, i32 0)
store <16 x float> %result, ptr addrspace(1) %out
ret void
@@ -1091,61 +1582,117 @@ define amdgpu_kernel void @test_mfma_f32_32x32x16_bf16__vgprcd_mac_flags(<8 x bf
; GCN-NEXT: global_store_dwordx4 v16, v[0:3], s[0:1]
; GCN-NEXT: s_endpgm
;
-; AGPR-LABEL: test_mfma_f32_32x32x16_bf16__vgprcd_mac_flags:
-; AGPR: ; %bb.0:
-; AGPR-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x24
-; AGPR-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x64
-; AGPR-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0xa4
-; AGPR-NEXT: s_waitcnt lgkmcnt(0)
-; AGPR-NEXT: v_mov_b64_e32 v[16:17], s[24:25]
-; AGPR-NEXT: v_mov_b64_e32 v[18:19], s[26:27]
-; AGPR-NEXT: v_mov_b64_e32 v[20:21], s[28:29]
-; AGPR-NEXT: v_mov_b64_e32 v[22:23], s[30:31]
-; AGPR-NEXT: v_mov_b64_e32 v[0:1], s[8:9]
-; AGPR-NEXT: v_mov_b64_e32 v[2:3], s[10:11]
-; AGPR-NEXT: v_mov_b64_e32 v[4:5], s[12:13]
-; AGPR-NEXT: v_mov_b64_e32 v[6:7], s[14:15]
-; AGPR-NEXT: v_mov_b64_e32 v[8:9], s[16:17]
-; AGPR-NEXT: v_mov_b64_e32 v[10:11], s[18:19]
-; AGPR-NEXT: v_mov_b64_e32 v[12:13], s[20:21]
-; AGPR-NEXT: v_mov_b64_e32 v[14:15], s[22:23]
-; AGPR-NEXT: s_nop 1
-; AGPR-NEXT: v_mfma_f32_32x32x16_bf16 v[0:15], v[16:19], v[20:23], v[0:15] cbsz:3 abid:2 blgp:1
-; AGPR-NEXT: v_mov_b32_e32 v16, 0
-; AGPR-NEXT: s_nop 10
-; AGPR-NEXT: global_store_dwordx4 v16, v[12:15], s[0:1] offset:48
-; AGPR-NEXT: global_store_dwordx4 v16, v[8:11], s[0:1] offset:32
-; AGPR-NEXT: global_store_dwordx4 v16, v[4:7], s[0:1] offset:16
-; AGPR-NEXT: global_store_dwordx4 v16, v[0:3], s[0:1]
-; AGPR-NEXT: s_endpgm
+; SDAG-AGPR-LABEL: test_mfma_f32_32x32x16_bf16__vgprcd_mac_flags:
+; SDAG-AGPR: ; %bb.0:
+; SDAG-AGPR-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x24
+; SDAG-AGPR-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x64
+; SDAG-AGPR-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0xa4
+; SDAG-AGPR-NEXT: s_waitcnt lgkmcnt(0)
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[16:17], s[24:25]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[18:19], s[26:27]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[20:21], s[28:29]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[22:23], s[30:31]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[0:1], s[8:9]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[2:3], s[10:11]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[4:5], s[12:13]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[6:7], s[14:15]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[8:9], s[16:17]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[10:11], s[18:19]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[12:13], s[20:21]
+; SDAG-AGPR-NEXT: v_mov_b64_e32 v[14:15], s[22:23]
+; SDAG-AGPR-NEXT: s_nop 1
+; SDAG-AGPR-NEXT: v_mfma_f32_32x32x16_bf16 v[0:15], v[16:19], v[20:23], v[0:15] cbsz:3 abid:2 blgp:1
+; SDAG-AGPR-NEXT: v_mov_b32_e32 v16, 0
+; SDAG-AGPR-NEXT: s_nop 10
+; SDAG-AGPR-NEXT: global_store_dwordx4 v16, v[12:15], s[0:1] offset:48
+; SDAG-AGPR-NEXT: global_store_dwordx4 v16, v[8:11], s[0:1] offset:32
+; SDAG-AGPR-NEXT: global_store_dwordx4 v16, v[4:7], s[0:1] offset:16
+; SDAG-AGPR-NEXT: global_store_dwordx4 v16, v[0:3], s[0:1]
+; SDAG-AGPR-NEXT: s_endpgm
;
-; VGPR-LABEL: test_mfma_f32_32x32x16_bf16__vgprcd_mac_flags:
-; VGPR: ; %bb.0:
-; VGPR-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x24
-; VGPR-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x64
-; VGPR-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0xa4
-; VGPR-NEXT: s_waitcnt lgkmcnt(0)
-; VGPR-NEXT: v_mov_b64_e32 v[16:17], s[24:25]
-; VGPR-NEXT: v_mov_b64_e32 v[18:19], s[26:27]
-; VGPR-NEXT: v_mov_b64_e32 v[20:21], s[28:29]
-; VGPR-NEXT: v_mov_b64_e32 v[22:23], s[30:31]
-; VGPR-NEXT: v_mov_b64_e32 v[0:1], s[8:9]
-; VGPR-NEXT: v_mov_b64_e32 v[2:3], s[10:11]
-; VGPR-NEXT: v_mov_b64_e32 v[4:5], s[12:13]
-; VGPR-NEXT: v_mov_b64_e32 v[6:7], s[14:15]
-; VGPR-NEXT: v_mov_b64_e32 v[8:9], s[16:17]
-; VGPR-NEXT: v_mov_b64_e32 v[10:11], s[18:19]
-; VGPR-NEXT: v_mov_b64_e32 v[12:13], s[20:21]
-; VGPR-NEXT: v_mov_b64_e32 v[14:15], s[22:23]
-; VGPR-NEXT: s_nop 1
-; VGPR-NEXT: v_mfma_f32_32x32x16_bf16 v[0:15], v[16:19], v[20:23], v[0:15] cbsz:3 abid:2 blgp:1
-; VGPR-NEXT: v_mov_b32_e32 v16, 0
-; VGPR-NEXT: s_nop 10
-; VGPR-NEXT: global_store_dwordx4 v16, v[12:15], s[0:1] offset:48
-; VGPR-NEXT: global_store_dwordx4 v16, v[8:11], s[0:1] offset:32
-; VGPR-NEXT: global_store_dwordx4 v16, v[4:7], s[0:1] offset:16
-; VGPR-NEXT: global_store_dwordx4 v16, v[0:3], s[0:1]
-; VGPR-NEXT: s_endpgm
+; GISEL-AGPR-LABEL: test_mfma_f32_32x32x16_bf16__vgprcd_mac_flags:
+; GISEL-AGPR: ; %bb.0:
+; GISEL-AGPR-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x24
+; GISEL-AGPR-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x64
+; GISEL-AGPR-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0xa4
+; GISEL-AGPR-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[16:17], s[24:25]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[18:19], s[26:27]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[20:21], s[28:29]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[22:23], s[30:31]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[0:1], s[8:9]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[2:3], s[10:11]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[4:5], s[12:13]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[6:7], s[14:15]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[8:9], s[16:17]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[10:11], s[18:19]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[12:13], s[20:21]
+; GISEL-AGPR-NEXT: v_mov_b64_e32 v[14:15], s[22:23]
+; GISEL-AGPR-NEXT: s_nop 1
+; GISEL-AGPR-NEXT: v_mfma_f32_32x32x16_bf16 v[0:15], v[16:19], v[20:23], v[0:15] cbsz:3 abid:2 blgp:1
+; GISEL-AGPR-NEXT: v_mov_b32_e32 v16, 0
+; GISEL-AGPR-NEXT: s_nop 10
+; GISEL-AGPR-NEXT: global_store_dwordx4 v16, v[0:3], s[0:1]
+; GISEL-AGPR-NEXT: global_store_dwordx4 v16, v[4:7], s[0:1] offset:16
+; GISEL-AGPR-NEXT: global_store_dwordx4 v16, v[8:11], s[0:1] offset:32
+; GISEL-AGPR-NEXT: global_store_dwordx4 v16, v[12:15], s[0:1] offset:48
+; GISEL-AGPR-NEXT: s_endpgm
+;
+; SDAG-VGPR-LABEL: test_mfma_f32_32x32x16_bf16__vgprcd_mac_flags:
+; SDAG-VGPR: ; %bb.0:
+; SDAG-VGPR-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x24
+; SDAG-VGPR-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x64
+; SDAG-VGPR-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0xa4
+; SDAG-VGPR-NEXT: s_waitcnt lgkmcnt(0)
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[16:17], s[24:25]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[18:19], s[26:27]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[20:21], s[28:29]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[22:23], s[30:31]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[0:1], s[8:9]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[2:3], s[10:11]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[4:5], s[12:13]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[6:7], s[14:15]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[8:9], s[16:17]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[10:11], s[18:19]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[12:13], s[20:21]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[14:15], s[22:23]
+; SDAG-VGPR-NEXT: s_nop 1
+; SDAG-VGPR-NEXT: v_mfma_f32_32x32x16_bf16 v[0:15], v[16:19], v[20:23], v[0:15] cbsz:3 abid:2 blgp:1
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v16, 0
+; SDAG-VGPR-NEXT: s_nop 10
+; SDAG-VGPR-NEXT: global_store_dwordx4 v16, v[12:15], s[0:1] offset:48
+; SDAG-VGPR-NEXT: global_store_dwordx4 v16, v[8:11], s[0:1] offset:32
+; SDAG-VGPR-NEXT: global_store_dwordx4 v16, v[4:7], s[0:1] offset:16
+; SDAG-VGPR-NEXT: global_store_dwordx4 v16, v[0:3], s[0:1]
+; SDAG-VGPR-NEXT: s_endpgm
+;
+; GISEL-VGPR-LABEL: test_mfma_f32_32x32x16_bf16__vgprcd_mac_flags:
+; GISEL-VGPR: ; %bb.0:
+; GISEL-VGPR-NEXT: s_load_dwordx8 s[24:31], s[4:5], 0x24
+; GISEL-VGPR-NEXT: s_load_dwordx16 s[8:23], s[4:5], 0x64
+; GISEL-VGPR-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0xa4
+; GISEL-VGPR-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[16:17], s[24:25]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[18:19], s[26:27]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[20:21], s[28:29]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[22:23], s[30:31]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[0:1], s[8:9]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[2:3], s[10:11]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[4:5], s[12:13]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[6:7], s[14:15]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[8:9], s[16:17]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[10:11], s[18:19]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[12:13], s[20:21]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[14:15], s[22:23]
+; GISEL-VGPR-NEXT: s_nop 1
+; GISEL-VGPR-NEXT: v_mfma_f32_32x32x16_bf16 v[0:15], v[16:19], v[20:23], v[0:15] cbsz:3 abid:2 blgp:1
+; GISEL-VGPR-NEXT: v_mov_b32_e32 v16, 0
+; GISEL-VGPR-NEXT: s_nop 10
+; GISEL-VGPR-NEXT: global_store_dwordx4 v16, v[0:3], s[0:1]
+; GISEL-VGPR-NEXT: global_store_dwordx4 v16, v[4:7], s[0:1] offset:16
+; GISEL-VGPR-NEXT: global_store_dwordx4 v16, v[8:11], s[0:1] offset:32
+; GISEL-VGPR-NEXT: global_store_dwordx4 v16, v[12:15], s[0:1] offset:48
+; GISEL-VGPR-NEXT: s_endpgm
%result = call <16 x float> @llvm.amdgcn.mfma.f32.32x32x16.bf16(<8 x bfloat> %arg0, <8 x bfloat> %arg1, <16 x float> %arg2, i32 3, i32 2, i32 1)
store <16 x float> %result, ptr addrspace(1) %out
ret void
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.gfx950.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.gfx950.ll
index cee17039bdf59..42cd877c42261 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.gfx950.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.gfx950.ll
@@ -1,6 +1,6 @@
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 4
; RUN: llc -mtriple=amdgpu9.50 -global-isel=0 < %s | FileCheck -enable-var-scope --check-prefixes=GCN,SDAG %s
-; RUN: llc -mtriple=amdgpu9.50 -global-isel=1 -global-isel-abort=2 < %s | FileCheck -enable-var-scope --check-prefixes=GCN,GISEL %s
+; RUN: llc -mtriple=amdgpu9.50 -global-isel=1 < %s | FileCheck -enable-var-scope --check-prefixes=GCN,GISEL %s
; RUN: llc -mtriple=amdgpu9.50 -global-isel=0 --amdgpu-mfma-vgpr-form=0 < %s | FileCheck -enable-var-scope --check-prefixes=HEURRC %s
; RUN: llc -mtriple=amdgpu9.50 -global-isel=0 --amdgpu-mfma-vgpr-form=1 < %s | FileCheck -enable-var-scope --check-prefixes=VGPRRC %s
@@ -5224,24 +5224,43 @@ define <4 x float> @test_mfma_f32_16x16x32_bf16__flags(<8 x bfloat> %arg0, <8 x
}
define amdgpu_kernel void @test_mfma_f32_16x16x32_bf16_no_agpr__vgprcd(ptr addrspace(1) %out, <8 x bfloat> %arg0, <8 x bfloat> %arg1, <4 x float> %arg2) #0 {
-; GCN-LABEL: test_mfma_f32_16x16x32_bf16_no_agpr__vgprcd:
-; GCN: ; %bb.0:
-; GCN-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x34
-; GCN-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x54
-; GCN-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24
-; GCN-NEXT: v_mov_b32_e32 v4, 0
-; GCN-NEXT: s_waitcnt lgkmcnt(0)
-; GCN-NEXT: v_mov_b64_e32 v[6:7], s[8:9]
-; GCN-NEXT: v_mov_b64_e32 v[8:9], s[10:11]
-; GCN-NEXT: v_mov_b64_e32 v[10:11], s[12:13]
-; GCN-NEXT: v_mov_b64_e32 v[12:13], s[14:15]
-; GCN-NEXT: v_mov_b64_e32 v[0:1], s[0:1]
-; GCN-NEXT: v_mov_b64_e32 v[2:3], s[2:3]
-; GCN-NEXT: s_nop 1
-; GCN-NEXT: v_mfma_f32_16x16x32_bf16 v[0:3], v[6:9], v[10:13], v[0:3]
-; GCN-NEXT: s_nop 7
-; GCN-NEXT: global_store_dwordx4 v4, v[0:3], s[6:7]
-; GCN-NEXT: s_endpgm
+; SDAG-LABEL: test_mfma_f32_16x16x32_bf16_no_agpr__vgprcd:
+; SDAG: ; %bb.0:
+; SDAG-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x34
+; SDAG-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x54
+; SDAG-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24
+; SDAG-NEXT: v_mov_b32_e32 v4, 0
+; SDAG-NEXT: s_waitcnt lgkmcnt(0)
+; SDAG-NEXT: v_mov_b64_e32 v[6:7], s[8:9]
+; SDAG-NEXT: v_mov_b64_e32 v[8:9], s[10:11]
+; SDAG-NEXT: v_mov_b64_e32 v[10:11], s[12:13]
+; SDAG-NEXT: v_mov_b64_e32 v[12:13], s[14:15]
+; SDAG-NEXT: v_mov_b64_e32 v[0:1], s[0:1]
+; SDAG-NEXT: v_mov_b64_e32 v[2:3], s[2:3]
+; SDAG-NEXT: s_nop 1
+; SDAG-NEXT: v_mfma_f32_16x16x32_bf16 v[0:3], v[6:9], v[10:13], v[0:3]
+; SDAG-NEXT: s_nop 7
+; SDAG-NEXT: global_store_dwordx4 v4, v[0:3], s[6:7]
+; SDAG-NEXT: s_endpgm
+;
+; GISEL-LABEL: test_mfma_f32_16x16x32_bf16_no_agpr__vgprcd:
+; GISEL: ; %bb.0:
+; GISEL-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x34
+; GISEL-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x54
+; GISEL-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24
+; GISEL-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-NEXT: v_mov_b64_e32 v[0:1], s[8:9]
+; GISEL-NEXT: v_mov_b64_e32 v[2:3], s[10:11]
+; GISEL-NEXT: v_mov_b64_e32 v[4:5], s[12:13]
+; GISEL-NEXT: v_mov_b64_e32 v[6:7], s[14:15]
+; GISEL-NEXT: v_mov_b64_e32 v[10:11], s[2:3]
+; GISEL-NEXT: v_mov_b64_e32 v[8:9], s[0:1]
+; GISEL-NEXT: s_nop 1
+; GISEL-NEXT: v_mfma_f32_16x16x32_bf16 v[0:3], v[0:3], v[4:7], v[8:11]
+; GISEL-NEXT: v_mov_b32_e32 v4, 0
+; GISEL-NEXT: s_nop 6
+; GISEL-NEXT: global_store_dwordx4 v4, v[0:3], s[6:7]
+; GISEL-NEXT: s_endpgm
;
; HEURRC-LABEL: test_mfma_f32_16x16x32_bf16_no_agpr__vgprcd:
; HEURRC: ; %bb.0:
@@ -5324,24 +5343,43 @@ define amdgpu_kernel void @test_mfma_f32_16x16x32_bf16_no_agpr__vgprcd(ptr addrs
}
define amdgpu_kernel void @test_mfma_f32_16x16x32_bf16_no_agpr__vgprcd__flags(ptr addrspace(1) %out, <8 x bfloat> %arg0, <8 x bfloat> %arg1, <4 x float> %arg2) #0 {
-; GCN-LABEL: test_mfma_f32_16x16x32_bf16_no_agpr__vgprcd__flags:
-; GCN: ; %bb.0:
-; GCN-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x34
-; GCN-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x54
-; GCN-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24
-; GCN-NEXT: v_mov_b32_e32 v4, 0
-; GCN-NEXT: s_waitcnt lgkmcnt(0)
-; GCN-NEXT: v_mov_b64_e32 v[6:7], s[8:9]
-; GCN-NEXT: v_mov_b64_e32 v[8:9], s[10:11]
-; GCN-NEXT: v_mov_b64_e32 v[10:11], s[12:13]
-; GCN-NEXT: v_mov_b64_e32 v[12:13], s[14:15]
-; GCN-NEXT: v_mov_b64_e32 v[0:1], s[0:1]
-; GCN-NEXT: v_mov_b64_e32 v[2:3], s[2:3]
-; GCN-NEXT: s_nop 1
-; GCN-NEXT: v_mfma_f32_16x16x32_bf16 v[0:3], v[6:9], v[10:13], v[0:3] cbsz:3 abid:2 blgp:1
-; GCN-NEXT: s_nop 7
-; GCN-NEXT: global_store_dwordx4 v4, v[0:3], s[6:7]
-; GCN-NEXT: s_endpgm
+; SDAG-LABEL: test_mfma_f32_16x16x32_bf16_no_agpr__vgprcd__flags:
+; SDAG: ; %bb.0:
+; SDAG-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x34
+; SDAG-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x54
+; SDAG-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24
+; SDAG-NEXT: v_mov_b32_e32 v4, 0
+; SDAG-NEXT: s_waitcnt lgkmcnt(0)
+; SDAG-NEXT: v_mov_b64_e32 v[6:7], s[8:9]
+; SDAG-NEXT: v_mov_b64_e32 v[8:9], s[10:11]
+; SDAG-NEXT: v_mov_b64_e32 v[10:11], s[12:13]
+; SDAG-NEXT: v_mov_b64_e32 v[12:13], s[14:15]
+; SDAG-NEXT: v_mov_b64_e32 v[0:1], s[0:1]
+; SDAG-NEXT: v_mov_b64_e32 v[2:3], s[2:3]
+; SDAG-NEXT: s_nop 1
+; SDAG-NEXT: v_mfma_f32_16x16x32_bf16 v[0:3], v[6:9], v[10:13], v[0:3] cbsz:3 abid:2 blgp:1
+; SDAG-NEXT: s_nop 7
+; SDAG-NEXT: global_store_dwordx4 v4, v[0:3], s[6:7]
+; SDAG-NEXT: s_endpgm
+;
+; GISEL-LABEL: test_mfma_f32_16x16x32_bf16_no_agpr__vgprcd__flags:
+; GISEL: ; %bb.0:
+; GISEL-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x34
+; GISEL-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x54
+; GISEL-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24
+; GISEL-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-NEXT: v_mov_b64_e32 v[0:1], s[8:9]
+; GISEL-NEXT: v_mov_b64_e32 v[2:3], s[10:11]
+; GISEL-NEXT: v_mov_b64_e32 v[4:5], s[12:13]
+; GISEL-NEXT: v_mov_b64_e32 v[6:7], s[14:15]
+; GISEL-NEXT: v_mov_b64_e32 v[10:11], s[2:3]
+; GISEL-NEXT: v_mov_b64_e32 v[8:9], s[0:1]
+; GISEL-NEXT: s_nop 1
+; GISEL-NEXT: v_mfma_f32_16x16x32_bf16 v[0:3], v[0:3], v[4:7], v[8:11] cbsz:3 abid:2 blgp:1
+; GISEL-NEXT: v_mov_b32_e32 v4, 0
+; GISEL-NEXT: s_nop 6
+; GISEL-NEXT: global_store_dwordx4 v4, v[0:3], s[6:7]
+; GISEL-NEXT: s_endpgm
;
; HEURRC-LABEL: test_mfma_f32_16x16x32_bf16_no_agpr__vgprcd__flags:
; HEURRC: ; %bb.0:
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.smfmac.gfx950.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.smfmac.gfx950.ll
index 76b2d9ab4dc40..0f8c5ca34c501 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.smfmac.gfx950.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.smfmac.gfx950.ll
@@ -1,9 +1,9 @@
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 4
-; RUN: llc -mtriple=amdgpu9.50 -global-isel=0 -amdgpu-mfma-vgpr-form=0 < %s | FileCheck -enable-var-scope --check-prefixes=GCN,SDAG %s
-; RUN: llc -mtriple=amdgpu9.50 -global-isel=1 -global-isel-abort=2 -amdgpu-mfma-vgpr-form=0 < %s | FileCheck -enable-var-scope --check-prefixes=GCN,GISEL %s
+; RUN: llc -mtriple=amdgpu9.50 -global-isel=0 -amdgpu-mfma-vgpr-form=0 < %s | FileCheck -enable-var-scope --check-prefix=SDAG %s
+; RUN: llc -mtriple=amdgpu9.50 -global-isel=1 -amdgpu-mfma-vgpr-form=0 < %s | FileCheck -enable-var-scope --check-prefix=GISEL %s
; RUN: llc -mtriple=amdgpu9.50 -global-isel=0 < %s | FileCheck -enable-var-scope --check-prefixes=GCN-VGPR,SDAG-VGPR %s
-; RUN: llc -mtriple=amdgpu9.50 -global-isel=1 -global-isel-abort=2 < %s | FileCheck -enable-var-scope --check-prefixes=GCN-VGPR,GISEL-VGPR %s
+; RUN: llc -mtriple=amdgpu9.50 -global-isel=1 < %s | FileCheck -enable-var-scope --check-prefixes=GCN-VGPR,GISEL-VGPR %s
declare i32 @llvm.amdgcn.workitem.id.x()
@@ -857,53 +857,103 @@ define <16 x float> @test_smfmac_f32_32x32x32_f16__sgpr(<8 x half> inreg %arg0,
declare <4 x float> @llvm.amdgcn.smfmac.f32.16x16x64.bf16(<8 x bfloat>, <16 x bfloat>, <4 x float>, i32, i32 immarg, i32 immarg)
define amdgpu_kernel void @test_smfmac_f32_16x16x64_bf16__vgpr(ptr addrspace(1) %arg, <8 x bfloat> %a, <16 x bfloat> %b, i32 %idx) #0 {
-; GCN-LABEL: test_smfmac_f32_16x16x64_bf16__vgpr:
-; GCN: ; %bb.0: ; %bb
-; GCN-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24
-; GCN-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x34
-; GCN-NEXT: v_and_b32_e32 v0, 0x3ff, v0
-; GCN-NEXT: v_lshlrev_b32_e32 v0, 4, v0
-; GCN-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x44
-; GCN-NEXT: s_load_dword s16, s[4:5], 0x64
-; GCN-NEXT: s_waitcnt lgkmcnt(0)
-; GCN-NEXT: global_load_dwordx4 v[10:13], v0, s[6:7]
-; GCN-NEXT: v_mov_b64_e32 v[16:17], s[2:3]
-; GCN-NEXT: v_mov_b64_e32 v[14:15], s[0:1]
-; GCN-NEXT: v_mov_b64_e32 v[2:3], s[8:9]
-; GCN-NEXT: v_mov_b64_e32 v[4:5], s[10:11]
-; GCN-NEXT: v_mov_b64_e32 v[6:7], s[12:13]
-; GCN-NEXT: v_mov_b64_e32 v[8:9], s[14:15]
-; GCN-NEXT: v_mov_b32_e32 v1, s16
-; GCN-NEXT: v_mov_b32_e32 v0, 0
-; GCN-NEXT: s_waitcnt vmcnt(0)
-; GCN-NEXT: v_smfmac_f32_16x16x64_bf16 v[10:13], v[14:17], v[2:9], v1 cbsz:1 abid:2
-; GCN-NEXT: s_nop 7
-; GCN-NEXT: global_store_dwordx4 v0, v[10:13], s[6:7]
-; GCN-NEXT: s_endpgm
-;
-; GCN-VGPR-LABEL: test_smfmac_f32_16x16x64_bf16__vgpr:
-; GCN-VGPR: ; %bb.0: ; %bb
-; GCN-VGPR-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24
-; GCN-VGPR-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x34
-; GCN-VGPR-NEXT: v_and_b32_e32 v0, 0x3ff, v0
-; GCN-VGPR-NEXT: v_lshlrev_b32_e32 v0, 4, v0
-; GCN-VGPR-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x44
-; GCN-VGPR-NEXT: s_load_dword s16, s[4:5], 0x64
-; GCN-VGPR-NEXT: s_waitcnt lgkmcnt(0)
-; GCN-VGPR-NEXT: global_load_dwordx4 v[10:13], v0, s[6:7]
-; GCN-VGPR-NEXT: v_mov_b64_e32 v[16:17], s[2:3]
-; GCN-VGPR-NEXT: v_mov_b64_e32 v[14:15], s[0:1]
-; GCN-VGPR-NEXT: v_mov_b64_e32 v[2:3], s[8:9]
-; GCN-VGPR-NEXT: v_mov_b64_e32 v[4:5], s[10:11]
-; GCN-VGPR-NEXT: v_mov_b64_e32 v[6:7], s[12:13]
-; GCN-VGPR-NEXT: v_mov_b64_e32 v[8:9], s[14:15]
-; GCN-VGPR-NEXT: v_mov_b32_e32 v1, s16
-; GCN-VGPR-NEXT: v_mov_b32_e32 v0, 0
-; GCN-VGPR-NEXT: s_waitcnt vmcnt(0)
-; GCN-VGPR-NEXT: v_smfmac_f32_16x16x64_bf16 v[10:13], v[14:17], v[2:9], v1 cbsz:1 abid:2
-; GCN-VGPR-NEXT: s_nop 7
-; GCN-VGPR-NEXT: global_store_dwordx4 v0, v[10:13], s[6:7]
-; GCN-VGPR-NEXT: s_endpgm
+; SDAG-LABEL: test_smfmac_f32_16x16x64_bf16__vgpr:
+; SDAG: ; %bb.0: ; %bb
+; SDAG-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24
+; SDAG-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x34
+; SDAG-NEXT: v_and_b32_e32 v0, 0x3ff, v0
+; SDAG-NEXT: v_lshlrev_b32_e32 v0, 4, v0
+; SDAG-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x44
+; SDAG-NEXT: s_load_dword s16, s[4:5], 0x64
+; SDAG-NEXT: s_waitcnt lgkmcnt(0)
+; SDAG-NEXT: global_load_dwordx4 v[10:13], v0, s[6:7]
+; SDAG-NEXT: v_mov_b64_e32 v[16:17], s[2:3]
+; SDAG-NEXT: v_mov_b64_e32 v[14:15], s[0:1]
+; SDAG-NEXT: v_mov_b64_e32 v[2:3], s[8:9]
+; SDAG-NEXT: v_mov_b64_e32 v[4:5], s[10:11]
+; SDAG-NEXT: v_mov_b64_e32 v[6:7], s[12:13]
+; SDAG-NEXT: v_mov_b64_e32 v[8:9], s[14:15]
+; SDAG-NEXT: v_mov_b32_e32 v1, s16
+; SDAG-NEXT: v_mov_b32_e32 v0, 0
+; SDAG-NEXT: s_waitcnt vmcnt(0)
+; SDAG-NEXT: v_smfmac_f32_16x16x64_bf16 v[10:13], v[14:17], v[2:9], v1 cbsz:1 abid:2
+; SDAG-NEXT: s_nop 7
+; SDAG-NEXT: global_store_dwordx4 v0, v[10:13], s[6:7]
+; SDAG-NEXT: s_endpgm
+;
+; GISEL-LABEL: test_smfmac_f32_16x16x64_bf16__vgpr:
+; GISEL: ; %bb.0: ; %bb
+; GISEL-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24
+; GISEL-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x34
+; GISEL-NEXT: v_and_b32_e32 v0, 0x3ff, v0
+; GISEL-NEXT: v_lshlrev_b32_e32 v0, 4, v0
+; GISEL-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x44
+; GISEL-NEXT: s_load_dword s16, s[4:5], 0x64
+; GISEL-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-NEXT: global_load_dwordx4 v[14:17], v0, s[6:7]
+; GISEL-NEXT: v_mov_b64_e32 v[10:11], s[2:3]
+; GISEL-NEXT: v_mov_b64_e32 v[8:9], s[0:1]
+; GISEL-NEXT: v_mov_b64_e32 v[0:1], s[8:9]
+; GISEL-NEXT: v_mov_b64_e32 v[2:3], s[10:11]
+; GISEL-NEXT: v_mov_b64_e32 v[4:5], s[12:13]
+; GISEL-NEXT: v_mov_b64_e32 v[6:7], s[14:15]
+; GISEL-NEXT: v_mov_b32_e32 v12, s16
+; GISEL-NEXT: s_waitcnt vmcnt(0)
+; GISEL-NEXT: s_nop 0
+; GISEL-NEXT: v_smfmac_f32_16x16x64_bf16 v[14:17], v[8:11], v[0:7], v12 cbsz:1 abid:2
+; GISEL-NEXT: v_mov_b32_e32 v0, 0
+; GISEL-NEXT: s_nop 6
+; GISEL-NEXT: global_store_dwordx4 v0, v[14:17], s[6:7]
+; GISEL-NEXT: s_endpgm
+;
+; SDAG-VGPR-LABEL: test_smfmac_f32_16x16x64_bf16__vgpr:
+; SDAG-VGPR: ; %bb.0: ; %bb
+; SDAG-VGPR-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24
+; SDAG-VGPR-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x34
+; SDAG-VGPR-NEXT: v_and_b32_e32 v0, 0x3ff, v0
+; SDAG-VGPR-NEXT: v_lshlrev_b32_e32 v0, 4, v0
+; SDAG-VGPR-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x44
+; SDAG-VGPR-NEXT: s_load_dword s16, s[4:5], 0x64
+; SDAG-VGPR-NEXT: s_waitcnt lgkmcnt(0)
+; SDAG-VGPR-NEXT: global_load_dwordx4 v[10:13], v0, s[6:7]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[16:17], s[2:3]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[14:15], s[0:1]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[2:3], s[8:9]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[4:5], s[10:11]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[6:7], s[12:13]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[8:9], s[14:15]
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v1, s16
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v0, 0
+; SDAG-VGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-VGPR-NEXT: v_smfmac_f32_16x16x64_bf16 v[10:13], v[14:17], v[2:9], v1 cbsz:1 abid:2
+; SDAG-VGPR-NEXT: s_nop 7
+; SDAG-VGPR-NEXT: global_store_dwordx4 v0, v[10:13], s[6:7]
+; SDAG-VGPR-NEXT: s_endpgm
+;
+; GISEL-VGPR-LABEL: test_smfmac_f32_16x16x64_bf16__vgpr:
+; GISEL-VGPR: ; %bb.0: ; %bb
+; GISEL-VGPR-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24
+; GISEL-VGPR-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x34
+; GISEL-VGPR-NEXT: v_and_b32_e32 v0, 0x3ff, v0
+; GISEL-VGPR-NEXT: v_lshlrev_b32_e32 v0, 4, v0
+; GISEL-VGPR-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x44
+; GISEL-VGPR-NEXT: s_load_dword s16, s[4:5], 0x64
+; GISEL-VGPR-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-VGPR-NEXT: global_load_dwordx4 v[14:17], v0, s[6:7]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[10:11], s[2:3]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[8:9], s[0:1]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[0:1], s[8:9]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[2:3], s[10:11]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[4:5], s[12:13]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[6:7], s[14:15]
+; GISEL-VGPR-NEXT: v_mov_b32_e32 v12, s16
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-VGPR-NEXT: s_nop 0
+; GISEL-VGPR-NEXT: v_smfmac_f32_16x16x64_bf16 v[14:17], v[8:11], v[0:7], v12 cbsz:1 abid:2
+; GISEL-VGPR-NEXT: v_mov_b32_e32 v0, 0
+; GISEL-VGPR-NEXT: s_nop 6
+; GISEL-VGPR-NEXT: global_store_dwordx4 v0, v[14:17], s[6:7]
+; GISEL-VGPR-NEXT: s_endpgm
bb:
%id = call i32 @llvm.amdgcn.workitem.id.x()
%gep = getelementptr <4 x float>, ptr addrspace(1) %arg, i32 %id
@@ -914,16 +964,25 @@ bb:
}
define <4 x float> @test_smfmac_f32_16x16x64_bf16(<8 x bfloat> %arg0, <16 x bfloat> %arg1, <4 x float> %arg2, i32 %arg3) #1 {
-; GCN-LABEL: test_smfmac_f32_16x16x64_bf16:
-; GCN: ; %bb.0:
-; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GCN-NEXT: v_smfmac_f32_16x16x64_bf16 v[12:15], v[0:3], v[4:11], v16
-; GCN-NEXT: s_nop 7
-; GCN-NEXT: v_mov_b32_e32 v0, v12
-; GCN-NEXT: v_mov_b32_e32 v1, v13
-; GCN-NEXT: v_mov_b32_e32 v2, v14
-; GCN-NEXT: v_mov_b32_e32 v3, v15
-; GCN-NEXT: s_setpc_b64 s[30:31]
+; SDAG-LABEL: test_smfmac_f32_16x16x64_bf16:
+; SDAG: ; %bb.0:
+; SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; SDAG-NEXT: v_smfmac_f32_16x16x64_bf16 v[12:15], v[0:3], v[4:11], v16
+; SDAG-NEXT: s_nop 7
+; SDAG-NEXT: v_mov_b32_e32 v0, v12
+; SDAG-NEXT: v_mov_b32_e32 v1, v13
+; SDAG-NEXT: v_mov_b32_e32 v2, v14
+; SDAG-NEXT: v_mov_b32_e32 v3, v15
+; SDAG-NEXT: s_setpc_b64 s[30:31]
+;
+; GISEL-LABEL: test_smfmac_f32_16x16x64_bf16:
+; GISEL: ; %bb.0:
+; GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GISEL-NEXT: v_smfmac_f32_16x16x64_bf16 v[12:15], v[0:3], v[4:11], v16
+; GISEL-NEXT: s_nop 7
+; GISEL-NEXT: v_mov_b64_e32 v[0:1], v[12:13]
+; GISEL-NEXT: v_mov_b64_e32 v[2:3], v[14:15]
+; GISEL-NEXT: s_setpc_b64 s[30:31]
;
; GCN-VGPR-LABEL: test_smfmac_f32_16x16x64_bf16:
; GCN-VGPR: ; %bb.0:
@@ -940,16 +999,25 @@ define <4 x float> @test_smfmac_f32_16x16x64_bf16(<8 x bfloat> %arg0, <16 x bflo
}
define <4 x float> @test_smfmac_f32_16x16x64_bf16__flags0(<8 x bfloat> %arg0, <16 x bfloat> %arg1, <4 x float> %arg2, i32 %arg3) #1 {
-; GCN-LABEL: test_smfmac_f32_16x16x64_bf16__flags0:
-; GCN: ; %bb.0:
-; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GCN-NEXT: v_smfmac_f32_16x16x64_bf16 v[12:15], v[0:3], v[4:11], v16 cbsz:1 abid:3
-; GCN-NEXT: s_nop 7
-; GCN-NEXT: v_mov_b32_e32 v0, v12
-; GCN-NEXT: v_mov_b32_e32 v1, v13
-; GCN-NEXT: v_mov_b32_e32 v2, v14
-; GCN-NEXT: v_mov_b32_e32 v3, v15
-; GCN-NEXT: s_setpc_b64 s[30:31]
+; SDAG-LABEL: test_smfmac_f32_16x16x64_bf16__flags0:
+; SDAG: ; %bb.0:
+; SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; SDAG-NEXT: v_smfmac_f32_16x16x64_bf16 v[12:15], v[0:3], v[4:11], v16 cbsz:1 abid:3
+; SDAG-NEXT: s_nop 7
+; SDAG-NEXT: v_mov_b32_e32 v0, v12
+; SDAG-NEXT: v_mov_b32_e32 v1, v13
+; SDAG-NEXT: v_mov_b32_e32 v2, v14
+; SDAG-NEXT: v_mov_b32_e32 v3, v15
+; SDAG-NEXT: s_setpc_b64 s[30:31]
+;
+; GISEL-LABEL: test_smfmac_f32_16x16x64_bf16__flags0:
+; GISEL: ; %bb.0:
+; GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GISEL-NEXT: v_smfmac_f32_16x16x64_bf16 v[12:15], v[0:3], v[4:11], v16 cbsz:1 abid:3
+; GISEL-NEXT: s_nop 7
+; GISEL-NEXT: v_mov_b64_e32 v[0:1], v[12:13]
+; GISEL-NEXT: v_mov_b64_e32 v[2:3], v[14:15]
+; GISEL-NEXT: s_setpc_b64 s[30:31]
;
; GCN-VGPR-LABEL: test_smfmac_f32_16x16x64_bf16__flags0:
; GCN-VGPR: ; %bb.0:
@@ -966,16 +1034,25 @@ define <4 x float> @test_smfmac_f32_16x16x64_bf16__flags0(<8 x bfloat> %arg0, <1
}
define <4 x float> @test_smfmac_f32_16x16x64_bf16__flags1(<8 x bfloat> %arg0, <16 x bfloat> %arg1, <4 x float> %arg2, i32 %arg3) #1 {
-; GCN-LABEL: test_smfmac_f32_16x16x64_bf16__flags1:
-; GCN: ; %bb.0:
-; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GCN-NEXT: v_smfmac_f32_16x16x64_bf16 v[12:15], v[0:3], v[4:11], v16 cbsz:3 abid:1
-; GCN-NEXT: s_nop 7
-; GCN-NEXT: v_mov_b32_e32 v0, v12
-; GCN-NEXT: v_mov_b32_e32 v1, v13
-; GCN-NEXT: v_mov_b32_e32 v2, v14
-; GCN-NEXT: v_mov_b32_e32 v3, v15
-; GCN-NEXT: s_setpc_b64 s[30:31]
+; SDAG-LABEL: test_smfmac_f32_16x16x64_bf16__flags1:
+; SDAG: ; %bb.0:
+; SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; SDAG-NEXT: v_smfmac_f32_16x16x64_bf16 v[12:15], v[0:3], v[4:11], v16 cbsz:3 abid:1
+; SDAG-NEXT: s_nop 7
+; SDAG-NEXT: v_mov_b32_e32 v0, v12
+; SDAG-NEXT: v_mov_b32_e32 v1, v13
+; SDAG-NEXT: v_mov_b32_e32 v2, v14
+; SDAG-NEXT: v_mov_b32_e32 v3, v15
+; SDAG-NEXT: s_setpc_b64 s[30:31]
+;
+; GISEL-LABEL: test_smfmac_f32_16x16x64_bf16__flags1:
+; GISEL: ; %bb.0:
+; GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GISEL-NEXT: v_smfmac_f32_16x16x64_bf16 v[12:15], v[0:3], v[4:11], v16 cbsz:3 abid:1
+; GISEL-NEXT: s_nop 7
+; GISEL-NEXT: v_mov_b64_e32 v[0:1], v[12:13]
+; GISEL-NEXT: v_mov_b64_e32 v[2:3], v[14:15]
+; GISEL-NEXT: s_setpc_b64 s[30:31]
;
; GCN-VGPR-LABEL: test_smfmac_f32_16x16x64_bf16__flags1:
; GCN-VGPR: ; %bb.0:
@@ -992,53 +1069,90 @@ define <4 x float> @test_smfmac_f32_16x16x64_bf16__flags1(<8 x bfloat> %arg0, <1
}
define <4 x float> @test_smfmac_f32_16x16x64_bf16__sgpr(<8 x bfloat> inreg %arg0, <16 x bfloat> inreg %arg1, <4 x float> inreg %arg2, i32 inreg %arg3) #1 {
-; GCN-LABEL: test_smfmac_f32_16x16x64_bf16__sgpr:
-; GCN: ; %bb.0:
-; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GCN-NEXT: v_mov_b32_e32 v14, s0
-; GCN-NEXT: v_mov_b32_e32 v15, s1
-; GCN-NEXT: v_mov_b32_e32 v16, s2
-; GCN-NEXT: v_mov_b32_e32 v17, s3
-; GCN-NEXT: v_mov_b32_e32 v6, s16
-; GCN-NEXT: v_mov_b32_e32 v7, s17
-; GCN-NEXT: v_mov_b32_e32 v8, s18
-; GCN-NEXT: v_mov_b32_e32 v9, s19
-; GCN-NEXT: v_mov_b32_e32 v10, s20
-; GCN-NEXT: v_mov_b32_e32 v11, s21
-; GCN-NEXT: v_mov_b32_e32 v12, s22
-; GCN-NEXT: v_mov_b32_e32 v13, s23
-; GCN-NEXT: v_mov_b32_e32 v0, s24
-; GCN-NEXT: v_mov_b32_e32 v1, s25
-; GCN-NEXT: v_mov_b32_e32 v2, s26
-; GCN-NEXT: v_mov_b32_e32 v3, s27
-; GCN-NEXT: v_mov_b32_e32 v4, s28
-; GCN-NEXT: s_nop 1
-; GCN-NEXT: v_smfmac_f32_16x16x64_bf16 v[0:3], v[14:17], v[6:13], v4
-; GCN-NEXT: s_setpc_b64 s[30:31]
-;
-; GCN-VGPR-LABEL: test_smfmac_f32_16x16x64_bf16__sgpr:
-; GCN-VGPR: ; %bb.0:
-; GCN-VGPR-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GCN-VGPR-NEXT: v_mov_b32_e32 v14, s0
-; GCN-VGPR-NEXT: v_mov_b32_e32 v15, s1
-; GCN-VGPR-NEXT: v_mov_b32_e32 v16, s2
-; GCN-VGPR-NEXT: v_mov_b32_e32 v17, s3
-; GCN-VGPR-NEXT: v_mov_b32_e32 v6, s16
-; GCN-VGPR-NEXT: v_mov_b32_e32 v7, s17
-; GCN-VGPR-NEXT: v_mov_b32_e32 v8, s18
-; GCN-VGPR-NEXT: v_mov_b32_e32 v9, s19
-; GCN-VGPR-NEXT: v_mov_b32_e32 v10, s20
-; GCN-VGPR-NEXT: v_mov_b32_e32 v11, s21
-; GCN-VGPR-NEXT: v_mov_b32_e32 v12, s22
-; GCN-VGPR-NEXT: v_mov_b32_e32 v13, s23
-; GCN-VGPR-NEXT: v_mov_b32_e32 v0, s24
-; GCN-VGPR-NEXT: v_mov_b32_e32 v1, s25
-; GCN-VGPR-NEXT: v_mov_b32_e32 v2, s26
-; GCN-VGPR-NEXT: v_mov_b32_e32 v3, s27
-; GCN-VGPR-NEXT: v_mov_b32_e32 v4, s28
-; GCN-VGPR-NEXT: s_nop 1
-; GCN-VGPR-NEXT: v_smfmac_f32_16x16x64_bf16 v[0:3], v[14:17], v[6:13], v4
-; GCN-VGPR-NEXT: s_setpc_b64 s[30:31]
+; SDAG-LABEL: test_smfmac_f32_16x16x64_bf16__sgpr:
+; SDAG: ; %bb.0:
+; SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; SDAG-NEXT: v_mov_b32_e32 v14, s0
+; SDAG-NEXT: v_mov_b32_e32 v15, s1
+; SDAG-NEXT: v_mov_b32_e32 v16, s2
+; SDAG-NEXT: v_mov_b32_e32 v17, s3
+; SDAG-NEXT: v_mov_b32_e32 v6, s16
+; SDAG-NEXT: v_mov_b32_e32 v7, s17
+; SDAG-NEXT: v_mov_b32_e32 v8, s18
+; SDAG-NEXT: v_mov_b32_e32 v9, s19
+; SDAG-NEXT: v_mov_b32_e32 v10, s20
+; SDAG-NEXT: v_mov_b32_e32 v11, s21
+; SDAG-NEXT: v_mov_b32_e32 v12, s22
+; SDAG-NEXT: v_mov_b32_e32 v13, s23
+; SDAG-NEXT: v_mov_b32_e32 v0, s24
+; SDAG-NEXT: v_mov_b32_e32 v1, s25
+; SDAG-NEXT: v_mov_b32_e32 v2, s26
+; SDAG-NEXT: v_mov_b32_e32 v3, s27
+; SDAG-NEXT: v_mov_b32_e32 v4, s28
+; SDAG-NEXT: s_nop 1
+; SDAG-NEXT: v_smfmac_f32_16x16x64_bf16 v[0:3], v[14:17], v[6:13], v4
+; SDAG-NEXT: s_setpc_b64 s[30:31]
+;
+; GISEL-LABEL: test_smfmac_f32_16x16x64_bf16__sgpr:
+; GISEL: ; %bb.0:
+; GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GISEL-NEXT: v_mov_b64_e32 v[10:11], s[2:3]
+; GISEL-NEXT: v_mov_b64_e32 v[8:9], s[0:1]
+; GISEL-NEXT: v_mov_b64_e32 v[0:1], s[16:17]
+; GISEL-NEXT: v_mov_b64_e32 v[2:3], s[18:19]
+; GISEL-NEXT: v_mov_b64_e32 v[4:5], s[20:21]
+; GISEL-NEXT: v_mov_b64_e32 v[6:7], s[22:23]
+; GISEL-NEXT: v_mov_b32_e32 v14, s24
+; GISEL-NEXT: v_mov_b32_e32 v15, s25
+; GISEL-NEXT: v_mov_b32_e32 v16, s26
+; GISEL-NEXT: v_mov_b32_e32 v17, s27
+; GISEL-NEXT: v_mov_b32_e32 v12, s28
+; GISEL-NEXT: s_nop 1
+; GISEL-NEXT: v_smfmac_f32_16x16x64_bf16 v[14:17], v[8:11], v[0:7], v12
+; GISEL-NEXT: s_nop 7
+; GISEL-NEXT: v_mov_b64_e32 v[0:1], v[14:15]
+; GISEL-NEXT: v_mov_b64_e32 v[2:3], v[16:17]
+; GISEL-NEXT: s_setpc_b64 s[30:31]
+;
+; SDAG-VGPR-LABEL: test_smfmac_f32_16x16x64_bf16__sgpr:
+; SDAG-VGPR: ; %bb.0:
+; SDAG-VGPR-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v14, s0
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v15, s1
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v16, s2
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v17, s3
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v6, s16
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v7, s17
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v8, s18
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v9, s19
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v10, s20
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v11, s21
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v12, s22
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v13, s23
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v0, s24
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v1, s25
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v2, s26
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v3, s27
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v4, s28
+; SDAG-VGPR-NEXT: s_nop 1
+; SDAG-VGPR-NEXT: v_smfmac_f32_16x16x64_bf16 v[0:3], v[14:17], v[6:13], v4
+; SDAG-VGPR-NEXT: s_setpc_b64 s[30:31]
+;
+; GISEL-VGPR-LABEL: test_smfmac_f32_16x16x64_bf16__sgpr:
+; GISEL-VGPR: ; %bb.0:
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[14:15], s[2:3]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[12:13], s[0:1]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[4:5], s[16:17]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[6:7], s[18:19]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[8:9], s[20:21]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[10:11], s[22:23]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[0:1], s[24:25]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[2:3], s[26:27]
+; GISEL-VGPR-NEXT: v_mov_b32_e32 v16, s28
+; GISEL-VGPR-NEXT: s_nop 1
+; GISEL-VGPR-NEXT: v_smfmac_f32_16x16x64_bf16 v[0:3], v[12:15], v[4:11], v16
+; GISEL-VGPR-NEXT: s_setpc_b64 s[30:31]
%result = call <4 x float> @llvm.amdgcn.smfmac.f32.16x16x64.bf16(<8 x bfloat> %arg0, <16 x bfloat> %arg1, <4 x float> %arg2, i32 %arg3, i32 immarg 0, i32 immarg 0)
ret <4 x float> %result
}
@@ -1050,112 +1164,189 @@ define <4 x float> @test_smfmac_f32_16x16x64_bf16__sgpr(<8 x bfloat> inreg %arg0
declare <16 x float> @llvm.amdgcn.smfmac.f32.32x32x32.bf16(<8 x bfloat>, <16 x bfloat>, <16 x float>, i32, i32 immarg, i32 immarg)
define amdgpu_kernel void @test_smfmac_f32_32x32x32_bf16__vgpr(ptr addrspace(1) %arg, <8 x bfloat> %a, <16 x bfloat> %b, i32 %idx) #0 {
-; GCN-LABEL: test_smfmac_f32_32x32x32_bf16__vgpr:
-; GCN: ; %bb.0: ; %bb
-; GCN-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24
-; GCN-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x34
-; GCN-NEXT: v_and_b32_e32 v0, 0x3ff, v0
-; GCN-NEXT: v_lshlrev_b32_e32 v16, 6, v0
-; GCN-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x44
-; GCN-NEXT: s_load_dword s16, s[4:5], 0x64
-; GCN-NEXT: s_waitcnt lgkmcnt(0)
-; GCN-NEXT: global_load_dwordx4 v[12:15], v16, s[6:7] offset:48
-; GCN-NEXT: global_load_dwordx4 v[8:11], v16, s[6:7] offset:32
-; GCN-NEXT: global_load_dwordx4 v[4:7], v16, s[6:7] offset:16
-; GCN-NEXT: global_load_dwordx4 v[0:3], v16, s[6:7]
-; GCN-NEXT: v_mov_b64_e32 v[28:29], s[2:3]
-; GCN-NEXT: v_mov_b64_e32 v[26:27], s[0:1]
-; GCN-NEXT: v_mov_b64_e32 v[24:25], s[14:15]
-; GCN-NEXT: v_mov_b64_e32 v[22:23], s[12:13]
-; GCN-NEXT: v_mov_b64_e32 v[20:21], s[10:11]
-; GCN-NEXT: v_mov_b64_e32 v[18:19], s[8:9]
-; GCN-NEXT: v_mov_b32_e32 v16, s16
-; GCN-NEXT: s_waitcnt vmcnt(0)
-; GCN-NEXT: s_nop 0
-; GCN-NEXT: v_smfmac_f32_32x32x32_bf16 v[0:15], v[26:29], v[18:25], v16 cbsz:1 abid:2
-; GCN-NEXT: v_mov_b32_e32 v16, 0
-; GCN-NEXT: s_nop 10
-; GCN-NEXT: global_store_dwordx4 v16, v[8:11], s[6:7] offset:32
-; GCN-NEXT: global_store_dwordx4 v16, v[12:15], s[6:7] offset:48
-; GCN-NEXT: global_store_dwordx4 v16, v[0:3], s[6:7]
-; GCN-NEXT: global_store_dwordx4 v16, v[4:7], s[6:7] offset:16
-; GCN-NEXT: s_endpgm
-;
-; GCN-VGPR-LABEL: test_smfmac_f32_32x32x32_bf16__vgpr:
-; GCN-VGPR: ; %bb.0: ; %bb
-; GCN-VGPR-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24
-; GCN-VGPR-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x34
-; GCN-VGPR-NEXT: v_and_b32_e32 v0, 0x3ff, v0
-; GCN-VGPR-NEXT: v_lshlrev_b32_e32 v16, 6, v0
-; GCN-VGPR-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x44
-; GCN-VGPR-NEXT: s_load_dword s16, s[4:5], 0x64
-; GCN-VGPR-NEXT: s_waitcnt lgkmcnt(0)
-; GCN-VGPR-NEXT: global_load_dwordx4 v[12:15], v16, s[6:7] offset:48
-; GCN-VGPR-NEXT: global_load_dwordx4 v[8:11], v16, s[6:7] offset:32
-; GCN-VGPR-NEXT: global_load_dwordx4 v[4:7], v16, s[6:7] offset:16
-; GCN-VGPR-NEXT: global_load_dwordx4 v[0:3], v16, s[6:7]
-; GCN-VGPR-NEXT: v_mov_b64_e32 v[28:29], s[2:3]
-; GCN-VGPR-NEXT: v_mov_b64_e32 v[26:27], s[0:1]
-; GCN-VGPR-NEXT: v_mov_b64_e32 v[24:25], s[14:15]
-; GCN-VGPR-NEXT: v_mov_b64_e32 v[22:23], s[12:13]
-; GCN-VGPR-NEXT: v_mov_b64_e32 v[20:21], s[10:11]
-; GCN-VGPR-NEXT: v_mov_b64_e32 v[18:19], s[8:9]
-; GCN-VGPR-NEXT: v_mov_b32_e32 v16, s16
-; GCN-VGPR-NEXT: s_waitcnt vmcnt(0)
-; GCN-VGPR-NEXT: s_nop 0
-; GCN-VGPR-NEXT: v_smfmac_f32_32x32x32_bf16 v[0:15], v[26:29], v[18:25], v16 cbsz:1 abid:2
-; GCN-VGPR-NEXT: v_mov_b32_e32 v16, 0
-; GCN-VGPR-NEXT: s_nop 10
-; GCN-VGPR-NEXT: global_store_dwordx4 v16, v[8:11], s[6:7] offset:32
-; GCN-VGPR-NEXT: global_store_dwordx4 v16, v[12:15], s[6:7] offset:48
-; GCN-VGPR-NEXT: global_store_dwordx4 v16, v[0:3], s[6:7]
-; GCN-VGPR-NEXT: global_store_dwordx4 v16, v[4:7], s[6:7] offset:16
-; GCN-VGPR-NEXT: s_endpgm
-bb:
- %id = call i32 @llvm.amdgcn.workitem.id.x()
- %gep = getelementptr <16 x float>, ptr addrspace(1) %arg, i32 %id
- %in.1 = load <16 x float>, ptr addrspace(1) %gep
- %mai.1 = tail call <16 x float> @llvm.amdgcn.smfmac.f32.32x32x32.bf16(<8 x bfloat> %a, <16 x bfloat> %b, <16 x float> %in.1, i32 %idx, i32 1, i32 2)
- store <16 x float> %mai.1, ptr addrspace(1) %arg
- ret void
-}
-
-define <16 x float> @test_smfmac_f32_32x32x32_bf16(<8 x bfloat> %arg0, <16 x bfloat> %arg1, <16 x float> %arg2, i32 %arg3) #1 {
-; GCN-LABEL: test_smfmac_f32_32x32x32_bf16:
-; GCN: ; %bb.0:
-; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GCN-NEXT: v_smfmac_f32_32x32x32_bf16 v[12:27], v[0:3], v[4:11], v28
-; GCN-NEXT: s_nop 11
-; GCN-NEXT: v_mov_b32_e32 v0, v12
-; GCN-NEXT: v_mov_b32_e32 v1, v13
-; GCN-NEXT: v_mov_b32_e32 v2, v14
-; GCN-NEXT: v_mov_b32_e32 v3, v15
-; GCN-NEXT: v_mov_b32_e32 v4, v16
-; GCN-NEXT: v_mov_b32_e32 v5, v17
-; GCN-NEXT: v_mov_b32_e32 v6, v18
-; GCN-NEXT: v_mov_b32_e32 v7, v19
-; GCN-NEXT: v_mov_b32_e32 v8, v20
-; GCN-NEXT: v_mov_b32_e32 v9, v21
-; GCN-NEXT: v_mov_b32_e32 v10, v22
-; GCN-NEXT: v_mov_b32_e32 v11, v23
-; GCN-NEXT: v_mov_b32_e32 v12, v24
-; GCN-NEXT: v_mov_b32_e32 v13, v25
-; GCN-NEXT: v_mov_b32_e32 v14, v26
-; GCN-NEXT: v_mov_b32_e32 v15, v27
-; GCN-NEXT: s_setpc_b64 s[30:31]
+; SDAG-LABEL: test_smfmac_f32_32x32x32_bf16__vgpr:
+; SDAG: ; %bb.0: ; %bb
+; SDAG-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24
+; SDAG-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x34
+; SDAG-NEXT: v_and_b32_e32 v0, 0x3ff, v0
+; SDAG-NEXT: v_lshlrev_b32_e32 v16, 6, v0
+; SDAG-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x44
+; SDAG-NEXT: s_load_dword s16, s[4:5], 0x64
+; SDAG-NEXT: s_waitcnt lgkmcnt(0)
+; SDAG-NEXT: global_load_dwordx4 v[12:15], v16, s[6:7] offset:48
+; SDAG-NEXT: global_load_dwordx4 v[8:11], v16, s[6:7] offset:32
+; SDAG-NEXT: global_load_dwordx4 v[4:7], v16, s[6:7] offset:16
+; SDAG-NEXT: global_load_dwordx4 v[0:3], v16, s[6:7]
+; SDAG-NEXT: v_mov_b64_e32 v[28:29], s[2:3]
+; SDAG-NEXT: v_mov_b64_e32 v[26:27], s[0:1]
+; SDAG-NEXT: v_mov_b64_e32 v[24:25], s[14:15]
+; SDAG-NEXT: v_mov_b64_e32 v[22:23], s[12:13]
+; SDAG-NEXT: v_mov_b64_e32 v[20:21], s[10:11]
+; SDAG-NEXT: v_mov_b64_e32 v[18:19], s[8:9]
+; SDAG-NEXT: v_mov_b32_e32 v16, s16
+; SDAG-NEXT: s_waitcnt vmcnt(0)
+; SDAG-NEXT: s_nop 0
+; SDAG-NEXT: v_smfmac_f32_32x32x32_bf16 v[0:15], v[26:29], v[18:25], v16 cbsz:1 abid:2
+; SDAG-NEXT: v_mov_b32_e32 v16, 0
+; SDAG-NEXT: s_nop 10
+; SDAG-NEXT: global_store_dwordx4 v16, v[8:11], s[6:7] offset:32
+; SDAG-NEXT: global_store_dwordx4 v16, v[12:15], s[6:7] offset:48
+; SDAG-NEXT: global_store_dwordx4 v16, v[0:3], s[6:7]
+; SDAG-NEXT: global_store_dwordx4 v16, v[4:7], s[6:7] offset:16
+; SDAG-NEXT: s_endpgm
;
-; GCN-VGPR-LABEL: test_smfmac_f32_32x32x32_bf16:
-; GCN-VGPR: ; %bb.0:
-; GCN-VGPR-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GCN-VGPR-NEXT: v_smfmac_f32_32x32x32_bf16 v[12:27], v[0:3], v[4:11], v28
-; GCN-VGPR-NEXT: s_nop 11
-; GCN-VGPR-NEXT: v_mov_b32_e32 v0, v12
-; GCN-VGPR-NEXT: v_mov_b32_e32 v1, v13
-; GCN-VGPR-NEXT: v_mov_b32_e32 v2, v14
-; GCN-VGPR-NEXT: v_mov_b32_e32 v3, v15
-; GCN-VGPR-NEXT: v_mov_b32_e32 v4, v16
-; GCN-VGPR-NEXT: v_mov_b32_e32 v5, v17
-; GCN-VGPR-NEXT: v_mov_b32_e32 v6, v18
+; GISEL-LABEL: test_smfmac_f32_32x32x32_bf16__vgpr:
+; GISEL: ; %bb.0: ; %bb
+; GISEL-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24
+; GISEL-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x34
+; GISEL-NEXT: v_and_b32_e32 v0, 0x3ff, v0
+; GISEL-NEXT: v_lshlrev_b32_e32 v16, 6, v0
+; GISEL-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x44
+; GISEL-NEXT: s_load_dword s16, s[4:5], 0x64
+; GISEL-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-NEXT: global_load_dwordx4 v[0:3], v16, s[6:7]
+; GISEL-NEXT: global_load_dwordx4 v[4:7], v16, s[6:7] offset:16
+; GISEL-NEXT: global_load_dwordx4 v[8:11], v16, s[6:7] offset:32
+; GISEL-NEXT: global_load_dwordx4 v[12:15], v16, s[6:7] offset:48
+; GISEL-NEXT: v_mov_b64_e32 v[26:27], s[2:3]
+; GISEL-NEXT: v_mov_b64_e32 v[24:25], s[0:1]
+; GISEL-NEXT: v_mov_b64_e32 v[22:23], s[14:15]
+; GISEL-NEXT: v_mov_b64_e32 v[20:21], s[12:13]
+; GISEL-NEXT: v_mov_b64_e32 v[18:19], s[10:11]
+; GISEL-NEXT: v_mov_b64_e32 v[16:17], s[8:9]
+; GISEL-NEXT: v_mov_b32_e32 v28, s16
+; GISEL-NEXT: s_waitcnt vmcnt(0)
+; GISEL-NEXT: s_nop 0
+; GISEL-NEXT: v_smfmac_f32_32x32x32_bf16 v[0:15], v[24:27], v[16:23], v28 cbsz:1 abid:2
+; GISEL-NEXT: v_mov_b32_e32 v16, 0
+; GISEL-NEXT: s_nop 10
+; GISEL-NEXT: global_store_dwordx4 v16, v[0:3], s[6:7]
+; GISEL-NEXT: global_store_dwordx4 v16, v[4:7], s[6:7] offset:16
+; GISEL-NEXT: global_store_dwordx4 v16, v[8:11], s[6:7] offset:32
+; GISEL-NEXT: global_store_dwordx4 v16, v[12:15], s[6:7] offset:48
+; GISEL-NEXT: s_endpgm
+;
+; SDAG-VGPR-LABEL: test_smfmac_f32_32x32x32_bf16__vgpr:
+; SDAG-VGPR: ; %bb.0: ; %bb
+; SDAG-VGPR-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24
+; SDAG-VGPR-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x34
+; SDAG-VGPR-NEXT: v_and_b32_e32 v0, 0x3ff, v0
+; SDAG-VGPR-NEXT: v_lshlrev_b32_e32 v16, 6, v0
+; SDAG-VGPR-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x44
+; SDAG-VGPR-NEXT: s_load_dword s16, s[4:5], 0x64
+; SDAG-VGPR-NEXT: s_waitcnt lgkmcnt(0)
+; SDAG-VGPR-NEXT: global_load_dwordx4 v[12:15], v16, s[6:7] offset:48
+; SDAG-VGPR-NEXT: global_load_dwordx4 v[8:11], v16, s[6:7] offset:32
+; SDAG-VGPR-NEXT: global_load_dwordx4 v[4:7], v16, s[6:7] offset:16
+; SDAG-VGPR-NEXT: global_load_dwordx4 v[0:3], v16, s[6:7]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[28:29], s[2:3]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[26:27], s[0:1]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[24:25], s[14:15]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[22:23], s[12:13]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[20:21], s[10:11]
+; SDAG-VGPR-NEXT: v_mov_b64_e32 v[18:19], s[8:9]
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v16, s16
+; SDAG-VGPR-NEXT: s_waitcnt vmcnt(0)
+; SDAG-VGPR-NEXT: s_nop 0
+; SDAG-VGPR-NEXT: v_smfmac_f32_32x32x32_bf16 v[0:15], v[26:29], v[18:25], v16 cbsz:1 abid:2
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v16, 0
+; SDAG-VGPR-NEXT: s_nop 10
+; SDAG-VGPR-NEXT: global_store_dwordx4 v16, v[8:11], s[6:7] offset:32
+; SDAG-VGPR-NEXT: global_store_dwordx4 v16, v[12:15], s[6:7] offset:48
+; SDAG-VGPR-NEXT: global_store_dwordx4 v16, v[0:3], s[6:7]
+; SDAG-VGPR-NEXT: global_store_dwordx4 v16, v[4:7], s[6:7] offset:16
+; SDAG-VGPR-NEXT: s_endpgm
+;
+; GISEL-VGPR-LABEL: test_smfmac_f32_32x32x32_bf16__vgpr:
+; GISEL-VGPR: ; %bb.0: ; %bb
+; GISEL-VGPR-NEXT: s_load_dwordx2 s[6:7], s[4:5], 0x24
+; GISEL-VGPR-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x34
+; GISEL-VGPR-NEXT: v_and_b32_e32 v0, 0x3ff, v0
+; GISEL-VGPR-NEXT: v_lshlrev_b32_e32 v16, 6, v0
+; GISEL-VGPR-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x44
+; GISEL-VGPR-NEXT: s_load_dword s16, s[4:5], 0x64
+; GISEL-VGPR-NEXT: s_waitcnt lgkmcnt(0)
+; GISEL-VGPR-NEXT: global_load_dwordx4 v[0:3], v16, s[6:7]
+; GISEL-VGPR-NEXT: global_load_dwordx4 v[4:7], v16, s[6:7] offset:16
+; GISEL-VGPR-NEXT: global_load_dwordx4 v[8:11], v16, s[6:7] offset:32
+; GISEL-VGPR-NEXT: global_load_dwordx4 v[12:15], v16, s[6:7] offset:48
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[26:27], s[2:3]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[24:25], s[0:1]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[22:23], s[14:15]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[20:21], s[12:13]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[18:19], s[10:11]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[16:17], s[8:9]
+; GISEL-VGPR-NEXT: v_mov_b32_e32 v28, s16
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-VGPR-NEXT: s_nop 0
+; GISEL-VGPR-NEXT: v_smfmac_f32_32x32x32_bf16 v[0:15], v[24:27], v[16:23], v28 cbsz:1 abid:2
+; GISEL-VGPR-NEXT: v_mov_b32_e32 v16, 0
+; GISEL-VGPR-NEXT: s_nop 10
+; GISEL-VGPR-NEXT: global_store_dwordx4 v16, v[0:3], s[6:7]
+; GISEL-VGPR-NEXT: global_store_dwordx4 v16, v[4:7], s[6:7] offset:16
+; GISEL-VGPR-NEXT: global_store_dwordx4 v16, v[8:11], s[6:7] offset:32
+; GISEL-VGPR-NEXT: global_store_dwordx4 v16, v[12:15], s[6:7] offset:48
+; GISEL-VGPR-NEXT: s_endpgm
+bb:
+ %id = call i32 @llvm.amdgcn.workitem.id.x()
+ %gep = getelementptr <16 x float>, ptr addrspace(1) %arg, i32 %id
+ %in.1 = load <16 x float>, ptr addrspace(1) %gep
+ %mai.1 = tail call <16 x float> @llvm.amdgcn.smfmac.f32.32x32x32.bf16(<8 x bfloat> %a, <16 x bfloat> %b, <16 x float> %in.1, i32 %idx, i32 1, i32 2)
+ store <16 x float> %mai.1, ptr addrspace(1) %arg
+ ret void
+}
+
+define <16 x float> @test_smfmac_f32_32x32x32_bf16(<8 x bfloat> %arg0, <16 x bfloat> %arg1, <16 x float> %arg2, i32 %arg3) #1 {
+; SDAG-LABEL: test_smfmac_f32_32x32x32_bf16:
+; SDAG: ; %bb.0:
+; SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; SDAG-NEXT: v_smfmac_f32_32x32x32_bf16 v[12:27], v[0:3], v[4:11], v28
+; SDAG-NEXT: s_nop 11
+; SDAG-NEXT: v_mov_b32_e32 v0, v12
+; SDAG-NEXT: v_mov_b32_e32 v1, v13
+; SDAG-NEXT: v_mov_b32_e32 v2, v14
+; SDAG-NEXT: v_mov_b32_e32 v3, v15
+; SDAG-NEXT: v_mov_b32_e32 v4, v16
+; SDAG-NEXT: v_mov_b32_e32 v5, v17
+; SDAG-NEXT: v_mov_b32_e32 v6, v18
+; SDAG-NEXT: v_mov_b32_e32 v7, v19
+; SDAG-NEXT: v_mov_b32_e32 v8, v20
+; SDAG-NEXT: v_mov_b32_e32 v9, v21
+; SDAG-NEXT: v_mov_b32_e32 v10, v22
+; SDAG-NEXT: v_mov_b32_e32 v11, v23
+; SDAG-NEXT: v_mov_b32_e32 v12, v24
+; SDAG-NEXT: v_mov_b32_e32 v13, v25
+; SDAG-NEXT: v_mov_b32_e32 v14, v26
+; SDAG-NEXT: v_mov_b32_e32 v15, v27
+; SDAG-NEXT: s_setpc_b64 s[30:31]
+;
+; GISEL-LABEL: test_smfmac_f32_32x32x32_bf16:
+; GISEL: ; %bb.0:
+; GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GISEL-NEXT: v_smfmac_f32_32x32x32_bf16 v[12:27], v[0:3], v[4:11], v28
+; GISEL-NEXT: s_nop 11
+; GISEL-NEXT: v_mov_b64_e32 v[0:1], v[12:13]
+; GISEL-NEXT: v_mov_b64_e32 v[2:3], v[14:15]
+; GISEL-NEXT: v_mov_b64_e32 v[4:5], v[16:17]
+; GISEL-NEXT: v_mov_b64_e32 v[6:7], v[18:19]
+; GISEL-NEXT: v_mov_b64_e32 v[8:9], v[20:21]
+; GISEL-NEXT: v_mov_b64_e32 v[10:11], v[22:23]
+; GISEL-NEXT: v_mov_b64_e32 v[12:13], v[24:25]
+; GISEL-NEXT: v_mov_b64_e32 v[14:15], v[26:27]
+; GISEL-NEXT: s_setpc_b64 s[30:31]
+;
+; GCN-VGPR-LABEL: test_smfmac_f32_32x32x32_bf16:
+; GCN-VGPR: ; %bb.0:
+; GCN-VGPR-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GCN-VGPR-NEXT: v_smfmac_f32_32x32x32_bf16 v[12:27], v[0:3], v[4:11], v28
+; GCN-VGPR-NEXT: s_nop 11
+; GCN-VGPR-NEXT: v_mov_b32_e32 v0, v12
+; GCN-VGPR-NEXT: v_mov_b32_e32 v1, v13
+; GCN-VGPR-NEXT: v_mov_b32_e32 v2, v14
+; GCN-VGPR-NEXT: v_mov_b32_e32 v3, v15
+; GCN-VGPR-NEXT: v_mov_b32_e32 v4, v16
+; GCN-VGPR-NEXT: v_mov_b32_e32 v5, v17
+; GCN-VGPR-NEXT: v_mov_b32_e32 v6, v18
; GCN-VGPR-NEXT: v_mov_b32_e32 v7, v19
; GCN-VGPR-NEXT: v_mov_b32_e32 v8, v20
; GCN-VGPR-NEXT: v_mov_b32_e32 v9, v21
@@ -1171,28 +1362,43 @@ define <16 x float> @test_smfmac_f32_32x32x32_bf16(<8 x bfloat> %arg0, <16 x bfl
}
define <16 x float> @test_smfmac_f32_32x32x32_bf16__flags0(<8 x bfloat> %arg0, <16 x bfloat> %arg1, <16 x float> %arg2, i32 %arg3) #1 {
-; GCN-LABEL: test_smfmac_f32_32x32x32_bf16__flags0:
-; GCN: ; %bb.0:
-; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GCN-NEXT: v_smfmac_f32_32x32x32_bf16 v[12:27], v[0:3], v[4:11], v28 cbsz:1 abid:3
-; GCN-NEXT: s_nop 11
-; GCN-NEXT: v_mov_b32_e32 v0, v12
-; GCN-NEXT: v_mov_b32_e32 v1, v13
-; GCN-NEXT: v_mov_b32_e32 v2, v14
-; GCN-NEXT: v_mov_b32_e32 v3, v15
-; GCN-NEXT: v_mov_b32_e32 v4, v16
-; GCN-NEXT: v_mov_b32_e32 v5, v17
-; GCN-NEXT: v_mov_b32_e32 v6, v18
-; GCN-NEXT: v_mov_b32_e32 v7, v19
-; GCN-NEXT: v_mov_b32_e32 v8, v20
-; GCN-NEXT: v_mov_b32_e32 v9, v21
-; GCN-NEXT: v_mov_b32_e32 v10, v22
-; GCN-NEXT: v_mov_b32_e32 v11, v23
-; GCN-NEXT: v_mov_b32_e32 v12, v24
-; GCN-NEXT: v_mov_b32_e32 v13, v25
-; GCN-NEXT: v_mov_b32_e32 v14, v26
-; GCN-NEXT: v_mov_b32_e32 v15, v27
-; GCN-NEXT: s_setpc_b64 s[30:31]
+; SDAG-LABEL: test_smfmac_f32_32x32x32_bf16__flags0:
+; SDAG: ; %bb.0:
+; SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; SDAG-NEXT: v_smfmac_f32_32x32x32_bf16 v[12:27], v[0:3], v[4:11], v28 cbsz:1 abid:3
+; SDAG-NEXT: s_nop 11
+; SDAG-NEXT: v_mov_b32_e32 v0, v12
+; SDAG-NEXT: v_mov_b32_e32 v1, v13
+; SDAG-NEXT: v_mov_b32_e32 v2, v14
+; SDAG-NEXT: v_mov_b32_e32 v3, v15
+; SDAG-NEXT: v_mov_b32_e32 v4, v16
+; SDAG-NEXT: v_mov_b32_e32 v5, v17
+; SDAG-NEXT: v_mov_b32_e32 v6, v18
+; SDAG-NEXT: v_mov_b32_e32 v7, v19
+; SDAG-NEXT: v_mov_b32_e32 v8, v20
+; SDAG-NEXT: v_mov_b32_e32 v9, v21
+; SDAG-NEXT: v_mov_b32_e32 v10, v22
+; SDAG-NEXT: v_mov_b32_e32 v11, v23
+; SDAG-NEXT: v_mov_b32_e32 v12, v24
+; SDAG-NEXT: v_mov_b32_e32 v13, v25
+; SDAG-NEXT: v_mov_b32_e32 v14, v26
+; SDAG-NEXT: v_mov_b32_e32 v15, v27
+; SDAG-NEXT: s_setpc_b64 s[30:31]
+;
+; GISEL-LABEL: test_smfmac_f32_32x32x32_bf16__flags0:
+; GISEL: ; %bb.0:
+; GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GISEL-NEXT: v_smfmac_f32_32x32x32_bf16 v[12:27], v[0:3], v[4:11], v28 cbsz:1 abid:3
+; GISEL-NEXT: s_nop 11
+; GISEL-NEXT: v_mov_b64_e32 v[0:1], v[12:13]
+; GISEL-NEXT: v_mov_b64_e32 v[2:3], v[14:15]
+; GISEL-NEXT: v_mov_b64_e32 v[4:5], v[16:17]
+; GISEL-NEXT: v_mov_b64_e32 v[6:7], v[18:19]
+; GISEL-NEXT: v_mov_b64_e32 v[8:9], v[20:21]
+; GISEL-NEXT: v_mov_b64_e32 v[10:11], v[22:23]
+; GISEL-NEXT: v_mov_b64_e32 v[12:13], v[24:25]
+; GISEL-NEXT: v_mov_b64_e32 v[14:15], v[26:27]
+; GISEL-NEXT: s_setpc_b64 s[30:31]
;
; GCN-VGPR-LABEL: test_smfmac_f32_32x32x32_bf16__flags0:
; GCN-VGPR: ; %bb.0:
@@ -1221,28 +1427,43 @@ define <16 x float> @test_smfmac_f32_32x32x32_bf16__flags0(<8 x bfloat> %arg0, <
}
define <16 x float> @test_smfmac_f32_32x32x32_bf16__flags1(<8 x bfloat> %arg0, <16 x bfloat> %arg1, <16 x float> %arg2, i32 %arg3) #1 {
-; GCN-LABEL: test_smfmac_f32_32x32x32_bf16__flags1:
-; GCN: ; %bb.0:
-; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GCN-NEXT: v_smfmac_f32_32x32x32_bf16 v[12:27], v[0:3], v[4:11], v28 cbsz:3 abid:1
-; GCN-NEXT: s_nop 11
-; GCN-NEXT: v_mov_b32_e32 v0, v12
-; GCN-NEXT: v_mov_b32_e32 v1, v13
-; GCN-NEXT: v_mov_b32_e32 v2, v14
-; GCN-NEXT: v_mov_b32_e32 v3, v15
-; GCN-NEXT: v_mov_b32_e32 v4, v16
-; GCN-NEXT: v_mov_b32_e32 v5, v17
-; GCN-NEXT: v_mov_b32_e32 v6, v18
-; GCN-NEXT: v_mov_b32_e32 v7, v19
-; GCN-NEXT: v_mov_b32_e32 v8, v20
-; GCN-NEXT: v_mov_b32_e32 v9, v21
-; GCN-NEXT: v_mov_b32_e32 v10, v22
-; GCN-NEXT: v_mov_b32_e32 v11, v23
-; GCN-NEXT: v_mov_b32_e32 v12, v24
-; GCN-NEXT: v_mov_b32_e32 v13, v25
-; GCN-NEXT: v_mov_b32_e32 v14, v26
-; GCN-NEXT: v_mov_b32_e32 v15, v27
-; GCN-NEXT: s_setpc_b64 s[30:31]
+; SDAG-LABEL: test_smfmac_f32_32x32x32_bf16__flags1:
+; SDAG: ; %bb.0:
+; SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; SDAG-NEXT: v_smfmac_f32_32x32x32_bf16 v[12:27], v[0:3], v[4:11], v28 cbsz:3 abid:1
+; SDAG-NEXT: s_nop 11
+; SDAG-NEXT: v_mov_b32_e32 v0, v12
+; SDAG-NEXT: v_mov_b32_e32 v1, v13
+; SDAG-NEXT: v_mov_b32_e32 v2, v14
+; SDAG-NEXT: v_mov_b32_e32 v3, v15
+; SDAG-NEXT: v_mov_b32_e32 v4, v16
+; SDAG-NEXT: v_mov_b32_e32 v5, v17
+; SDAG-NEXT: v_mov_b32_e32 v6, v18
+; SDAG-NEXT: v_mov_b32_e32 v7, v19
+; SDAG-NEXT: v_mov_b32_e32 v8, v20
+; SDAG-NEXT: v_mov_b32_e32 v9, v21
+; SDAG-NEXT: v_mov_b32_e32 v10, v22
+; SDAG-NEXT: v_mov_b32_e32 v11, v23
+; SDAG-NEXT: v_mov_b32_e32 v12, v24
+; SDAG-NEXT: v_mov_b32_e32 v13, v25
+; SDAG-NEXT: v_mov_b32_e32 v14, v26
+; SDAG-NEXT: v_mov_b32_e32 v15, v27
+; SDAG-NEXT: s_setpc_b64 s[30:31]
+;
+; GISEL-LABEL: test_smfmac_f32_32x32x32_bf16__flags1:
+; GISEL: ; %bb.0:
+; GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GISEL-NEXT: v_smfmac_f32_32x32x32_bf16 v[12:27], v[0:3], v[4:11], v28 cbsz:3 abid:1
+; GISEL-NEXT: s_nop 11
+; GISEL-NEXT: v_mov_b64_e32 v[0:1], v[12:13]
+; GISEL-NEXT: v_mov_b64_e32 v[2:3], v[14:15]
+; GISEL-NEXT: v_mov_b64_e32 v[4:5], v[16:17]
+; GISEL-NEXT: v_mov_b64_e32 v[6:7], v[18:19]
+; GISEL-NEXT: v_mov_b64_e32 v[8:9], v[20:21]
+; GISEL-NEXT: v_mov_b64_e32 v[10:11], v[22:23]
+; GISEL-NEXT: v_mov_b64_e32 v[12:13], v[24:25]
+; GISEL-NEXT: v_mov_b64_e32 v[14:15], v[26:27]
+; GISEL-NEXT: s_setpc_b64 s[30:31]
;
; GCN-VGPR-LABEL: test_smfmac_f32_32x32x32_bf16__flags1:
; GCN-VGPR: ; %bb.0:
@@ -1271,101 +1492,203 @@ define <16 x float> @test_smfmac_f32_32x32x32_bf16__flags1(<8 x bfloat> %arg0, <
}
define <16 x float> @test_smfmac_f32_32x32x32_bf16__sgpr(<8 x bfloat> inreg %arg0, <16 x bfloat> inreg %arg1, <16 x float> inreg %arg2, i32 inreg %arg3) #1 {
-; GCN-LABEL: test_smfmac_f32_32x32x32_bf16__sgpr:
-; GCN: ; %bb.0:
-; GCN-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GCN-NEXT: v_mov_b32_e32 v26, s0
-; GCN-NEXT: v_mov_b32_e32 v27, s1
-; GCN-NEXT: v_mov_b32_e32 v28, s2
-; GCN-NEXT: v_mov_b32_e32 v29, s3
-; GCN-NEXT: v_mov_b32_e32 v16, v10
-; GCN-NEXT: v_readfirstlane_b32 s4, v9
-; GCN-NEXT: v_readfirstlane_b32 s5, v8
-; GCN-NEXT: v_readfirstlane_b32 s6, v7
-; GCN-NEXT: v_readfirstlane_b32 s7, v6
-; GCN-NEXT: v_readfirstlane_b32 s8, v5
-; GCN-NEXT: v_readfirstlane_b32 s9, v4
-; GCN-NEXT: v_readfirstlane_b32 s10, v3
-; GCN-NEXT: v_readfirstlane_b32 s11, v2
-; GCN-NEXT: v_readfirstlane_b32 s12, v1
-; GCN-NEXT: v_readfirstlane_b32 s13, v0
-; GCN-NEXT: v_readfirstlane_b32 s0, v16
-; GCN-NEXT: v_mov_b32_e32 v0, s24
-; GCN-NEXT: v_mov_b32_e32 v1, s25
-; GCN-NEXT: v_mov_b32_e32 v2, s26
-; GCN-NEXT: v_mov_b32_e32 v3, s27
-; GCN-NEXT: v_mov_b32_e32 v4, s28
-; GCN-NEXT: v_mov_b32_e32 v5, s29
-; GCN-NEXT: v_mov_b32_e32 v6, s13
-; GCN-NEXT: v_mov_b32_e32 v7, s12
-; GCN-NEXT: v_mov_b32_e32 v8, s11
-; GCN-NEXT: v_mov_b32_e32 v9, s10
-; GCN-NEXT: v_mov_b32_e32 v10, s9
-; GCN-NEXT: v_mov_b32_e32 v11, s8
-; GCN-NEXT: v_mov_b32_e32 v12, s7
-; GCN-NEXT: v_mov_b32_e32 v13, s6
-; GCN-NEXT: v_mov_b32_e32 v14, s5
-; GCN-NEXT: v_mov_b32_e32 v15, s4
-; GCN-NEXT: v_mov_b32_e32 v18, s16
-; GCN-NEXT: v_mov_b32_e32 v19, s17
-; GCN-NEXT: v_mov_b32_e32 v20, s18
-; GCN-NEXT: v_mov_b32_e32 v21, s19
-; GCN-NEXT: v_mov_b32_e32 v22, s20
-; GCN-NEXT: v_mov_b32_e32 v23, s21
-; GCN-NEXT: v_mov_b32_e32 v24, s22
-; GCN-NEXT: v_mov_b32_e32 v25, s23
-; GCN-NEXT: v_mov_b32_e32 v16, s0
-; GCN-NEXT: s_nop 1
-; GCN-NEXT: v_smfmac_f32_32x32x32_bf16 v[0:15], v[26:29], v[18:25], v16
-; GCN-NEXT: s_setpc_b64 s[30:31]
-;
-; GCN-VGPR-LABEL: test_smfmac_f32_32x32x32_bf16__sgpr:
-; GCN-VGPR: ; %bb.0:
-; GCN-VGPR-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GCN-VGPR-NEXT: v_mov_b32_e32 v26, s0
-; GCN-VGPR-NEXT: v_mov_b32_e32 v27, s1
-; GCN-VGPR-NEXT: v_mov_b32_e32 v28, s2
-; GCN-VGPR-NEXT: v_mov_b32_e32 v29, s3
-; GCN-VGPR-NEXT: v_mov_b32_e32 v16, v10
-; GCN-VGPR-NEXT: v_readfirstlane_b32 s4, v9
-; GCN-VGPR-NEXT: v_readfirstlane_b32 s5, v8
-; GCN-VGPR-NEXT: v_readfirstlane_b32 s6, v7
-; GCN-VGPR-NEXT: v_readfirstlane_b32 s7, v6
-; GCN-VGPR-NEXT: v_readfirstlane_b32 s8, v5
-; GCN-VGPR-NEXT: v_readfirstlane_b32 s9, v4
-; GCN-VGPR-NEXT: v_readfirstlane_b32 s10, v3
-; GCN-VGPR-NEXT: v_readfirstlane_b32 s11, v2
-; GCN-VGPR-NEXT: v_readfirstlane_b32 s12, v1
-; GCN-VGPR-NEXT: v_readfirstlane_b32 s13, v0
-; GCN-VGPR-NEXT: v_readfirstlane_b32 s0, v16
-; GCN-VGPR-NEXT: v_mov_b32_e32 v0, s24
-; GCN-VGPR-NEXT: v_mov_b32_e32 v1, s25
-; GCN-VGPR-NEXT: v_mov_b32_e32 v2, s26
-; GCN-VGPR-NEXT: v_mov_b32_e32 v3, s27
-; GCN-VGPR-NEXT: v_mov_b32_e32 v4, s28
-; GCN-VGPR-NEXT: v_mov_b32_e32 v5, s29
-; GCN-VGPR-NEXT: v_mov_b32_e32 v6, s13
-; GCN-VGPR-NEXT: v_mov_b32_e32 v7, s12
-; GCN-VGPR-NEXT: v_mov_b32_e32 v8, s11
-; GCN-VGPR-NEXT: v_mov_b32_e32 v9, s10
-; GCN-VGPR-NEXT: v_mov_b32_e32 v10, s9
-; GCN-VGPR-NEXT: v_mov_b32_e32 v11, s8
-; GCN-VGPR-NEXT: v_mov_b32_e32 v12, s7
-; GCN-VGPR-NEXT: v_mov_b32_e32 v13, s6
-; GCN-VGPR-NEXT: v_mov_b32_e32 v14, s5
-; GCN-VGPR-NEXT: v_mov_b32_e32 v15, s4
-; GCN-VGPR-NEXT: v_mov_b32_e32 v18, s16
-; GCN-VGPR-NEXT: v_mov_b32_e32 v19, s17
-; GCN-VGPR-NEXT: v_mov_b32_e32 v20, s18
-; GCN-VGPR-NEXT: v_mov_b32_e32 v21, s19
-; GCN-VGPR-NEXT: v_mov_b32_e32 v22, s20
-; GCN-VGPR-NEXT: v_mov_b32_e32 v23, s21
-; GCN-VGPR-NEXT: v_mov_b32_e32 v24, s22
-; GCN-VGPR-NEXT: v_mov_b32_e32 v25, s23
-; GCN-VGPR-NEXT: v_mov_b32_e32 v16, s0
-; GCN-VGPR-NEXT: s_nop 1
-; GCN-VGPR-NEXT: v_smfmac_f32_32x32x32_bf16 v[0:15], v[26:29], v[18:25], v16
-; GCN-VGPR-NEXT: s_setpc_b64 s[30:31]
+; SDAG-LABEL: test_smfmac_f32_32x32x32_bf16__sgpr:
+; SDAG: ; %bb.0:
+; SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; SDAG-NEXT: v_mov_b32_e32 v26, s0
+; SDAG-NEXT: v_mov_b32_e32 v27, s1
+; SDAG-NEXT: v_mov_b32_e32 v28, s2
+; SDAG-NEXT: v_mov_b32_e32 v29, s3
+; SDAG-NEXT: v_mov_b32_e32 v16, v10
+; SDAG-NEXT: v_readfirstlane_b32 s4, v9
+; SDAG-NEXT: v_readfirstlane_b32 s5, v8
+; SDAG-NEXT: v_readfirstlane_b32 s6, v7
+; SDAG-NEXT: v_readfirstlane_b32 s7, v6
+; SDAG-NEXT: v_readfirstlane_b32 s8, v5
+; SDAG-NEXT: v_readfirstlane_b32 s9, v4
+; SDAG-NEXT: v_readfirstlane_b32 s10, v3
+; SDAG-NEXT: v_readfirstlane_b32 s11, v2
+; SDAG-NEXT: v_readfirstlane_b32 s12, v1
+; SDAG-NEXT: v_readfirstlane_b32 s13, v0
+; SDAG-NEXT: v_readfirstlane_b32 s0, v16
+; SDAG-NEXT: v_mov_b32_e32 v0, s24
+; SDAG-NEXT: v_mov_b32_e32 v1, s25
+; SDAG-NEXT: v_mov_b32_e32 v2, s26
+; SDAG-NEXT: v_mov_b32_e32 v3, s27
+; SDAG-NEXT: v_mov_b32_e32 v4, s28
+; SDAG-NEXT: v_mov_b32_e32 v5, s29
+; SDAG-NEXT: v_mov_b32_e32 v6, s13
+; SDAG-NEXT: v_mov_b32_e32 v7, s12
+; SDAG-NEXT: v_mov_b32_e32 v8, s11
+; SDAG-NEXT: v_mov_b32_e32 v9, s10
+; SDAG-NEXT: v_mov_b32_e32 v10, s9
+; SDAG-NEXT: v_mov_b32_e32 v11, s8
+; SDAG-NEXT: v_mov_b32_e32 v12, s7
+; SDAG-NEXT: v_mov_b32_e32 v13, s6
+; SDAG-NEXT: v_mov_b32_e32 v14, s5
+; SDAG-NEXT: v_mov_b32_e32 v15, s4
+; SDAG-NEXT: v_mov_b32_e32 v18, s16
+; SDAG-NEXT: v_mov_b32_e32 v19, s17
+; SDAG-NEXT: v_mov_b32_e32 v20, s18
+; SDAG-NEXT: v_mov_b32_e32 v21, s19
+; SDAG-NEXT: v_mov_b32_e32 v22, s20
+; SDAG-NEXT: v_mov_b32_e32 v23, s21
+; SDAG-NEXT: v_mov_b32_e32 v24, s22
+; SDAG-NEXT: v_mov_b32_e32 v25, s23
+; SDAG-NEXT: v_mov_b32_e32 v16, s0
+; SDAG-NEXT: s_nop 1
+; SDAG-NEXT: v_smfmac_f32_32x32x32_bf16 v[0:15], v[26:29], v[18:25], v16
+; SDAG-NEXT: s_setpc_b64 s[30:31]
+;
+; GISEL-LABEL: test_smfmac_f32_32x32x32_bf16__sgpr:
+; GISEL: ; %bb.0:
+; GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GISEL-NEXT: v_mov_b64_e32 v[26:27], s[2:3]
+; GISEL-NEXT: v_mov_b64_e32 v[24:25], s[0:1]
+; GISEL-NEXT: v_readfirstlane_b32 s4, v0
+; GISEL-NEXT: v_readfirstlane_b32 s5, v1
+; GISEL-NEXT: v_readfirstlane_b32 s6, v2
+; GISEL-NEXT: v_readfirstlane_b32 s7, v3
+; GISEL-NEXT: v_readfirstlane_b32 s8, v4
+; GISEL-NEXT: v_readfirstlane_b32 s9, v5
+; GISEL-NEXT: v_readfirstlane_b32 s10, v6
+; GISEL-NEXT: v_readfirstlane_b32 s11, v7
+; GISEL-NEXT: v_readfirstlane_b32 s12, v8
+; GISEL-NEXT: v_readfirstlane_b32 s13, v9
+; GISEL-NEXT: v_readfirstlane_b32 s14, v10
+; GISEL-NEXT: v_mov_b64_e32 v[16:17], s[16:17]
+; GISEL-NEXT: v_mov_b64_e32 v[18:19], s[18:19]
+; GISEL-NEXT: v_mov_b64_e32 v[20:21], s[20:21]
+; GISEL-NEXT: v_mov_b64_e32 v[22:23], s[22:23]
+; GISEL-NEXT: v_mov_b32_e32 v0, s24
+; GISEL-NEXT: v_mov_b32_e32 v1, s25
+; GISEL-NEXT: v_mov_b32_e32 v2, s26
+; GISEL-NEXT: v_mov_b32_e32 v3, s27
+; GISEL-NEXT: v_mov_b32_e32 v4, s28
+; GISEL-NEXT: v_mov_b32_e32 v5, s29
+; GISEL-NEXT: v_mov_b32_e32 v6, s4
+; GISEL-NEXT: v_mov_b32_e32 v7, s5
+; GISEL-NEXT: v_mov_b32_e32 v8, s6
+; GISEL-NEXT: v_mov_b32_e32 v9, s7
+; GISEL-NEXT: v_mov_b32_e32 v10, s8
+; GISEL-NEXT: v_mov_b32_e32 v11, s9
+; GISEL-NEXT: v_mov_b32_e32 v12, s10
+; GISEL-NEXT: v_mov_b32_e32 v13, s11
+; GISEL-NEXT: v_mov_b32_e32 v14, s12
+; GISEL-NEXT: v_mov_b32_e32 v15, s13
+; GISEL-NEXT: v_mov_b32_e32 v28, s14
+; GISEL-NEXT: s_nop 1
+; GISEL-NEXT: v_smfmac_f32_32x32x32_bf16 v[0:15], v[24:27], v[16:23], v28
+; GISEL-NEXT: s_setpc_b64 s[30:31]
+;
+; SDAG-VGPR-LABEL: test_smfmac_f32_32x32x32_bf16__sgpr:
+; SDAG-VGPR: ; %bb.0:
+; SDAG-VGPR-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v26, s0
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v27, s1
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v28, s2
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v29, s3
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v16, v10
+; SDAG-VGPR-NEXT: v_readfirstlane_b32 s4, v9
+; SDAG-VGPR-NEXT: v_readfirstlane_b32 s5, v8
+; SDAG-VGPR-NEXT: v_readfirstlane_b32 s6, v7
+; SDAG-VGPR-NEXT: v_readfirstlane_b32 s7, v6
+; SDAG-VGPR-NEXT: v_readfirstlane_b32 s8, v5
+; SDAG-VGPR-NEXT: v_readfirstlane_b32 s9, v4
+; SDAG-VGPR-NEXT: v_readfirstlane_b32 s10, v3
+; SDAG-VGPR-NEXT: v_readfirstlane_b32 s11, v2
+; SDAG-VGPR-NEXT: v_readfirstlane_b32 s12, v1
+; SDAG-VGPR-NEXT: v_readfirstlane_b32 s13, v0
+; SDAG-VGPR-NEXT: v_readfirstlane_b32 s0, v16
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v0, s24
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v1, s25
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v2, s26
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v3, s27
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v4, s28
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v5, s29
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v6, s13
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v7, s12
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v8, s11
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v9, s10
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v10, s9
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v11, s8
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v12, s7
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v13, s6
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v14, s5
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v15, s4
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v18, s16
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v19, s17
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v20, s18
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v21, s19
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v22, s20
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v23, s21
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v24, s22
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v25, s23
+; SDAG-VGPR-NEXT: v_mov_b32_e32 v16, s0
+; SDAG-VGPR-NEXT: s_nop 1
+; SDAG-VGPR-NEXT: v_smfmac_f32_32x32x32_bf16 v[0:15], v[26:29], v[18:25], v16
+; SDAG-VGPR-NEXT: s_setpc_b64 s[30:31]
+;
+; GISEL-VGPR-LABEL: test_smfmac_f32_32x32x32_bf16__sgpr:
+; GISEL-VGPR: ; %bb.0:
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GISEL-VGPR-NEXT: s_xor_saveexec_b64 s[4:5], -1
+; GISEL-VGPR-NEXT: scratch_store_dword off, v29, s32 ; 4-byte Folded Spill
+; GISEL-VGPR-NEXT: s_mov_b64 exec, s[4:5]
+; GISEL-VGPR-NEXT: v_writelane_b32 v29, s36, 0
+; GISEL-VGPR-NEXT: v_writelane_b32 v29, s37, 1
+; GISEL-VGPR-NEXT: v_writelane_b32 v29, s38, 2
+; GISEL-VGPR-NEXT: v_writelane_b32 v29, s39, 3
+; GISEL-VGPR-NEXT: v_writelane_b32 v29, s48, 4
+; GISEL-VGPR-NEXT: v_writelane_b32 v29, s49, 5
+; GISEL-VGPR-NEXT: v_writelane_b32 v29, s50, 6
+; GISEL-VGPR-NEXT: v_writelane_b32 v29, s51, 7
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[26:27], s[2:3]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[24:25], s[0:1]
+; GISEL-VGPR-NEXT: s_mov_b32 s36, s24
+; GISEL-VGPR-NEXT: s_mov_b32 s37, s25
+; GISEL-VGPR-NEXT: s_mov_b32 s38, s26
+; GISEL-VGPR-NEXT: s_mov_b32 s39, s27
+; GISEL-VGPR-NEXT: s_mov_b32 s40, s28
+; GISEL-VGPR-NEXT: s_mov_b32 s41, s29
+; GISEL-VGPR-NEXT: v_readfirstlane_b32 s42, v0
+; GISEL-VGPR-NEXT: v_readfirstlane_b32 s43, v1
+; GISEL-VGPR-NEXT: v_readfirstlane_b32 s44, v2
+; GISEL-VGPR-NEXT: v_readfirstlane_b32 s45, v3
+; GISEL-VGPR-NEXT: v_readfirstlane_b32 s46, v4
+; GISEL-VGPR-NEXT: v_readfirstlane_b32 s47, v5
+; GISEL-VGPR-NEXT: v_readfirstlane_b32 s48, v6
+; GISEL-VGPR-NEXT: v_readfirstlane_b32 s49, v7
+; GISEL-VGPR-NEXT: v_readfirstlane_b32 s50, v8
+; GISEL-VGPR-NEXT: v_readfirstlane_b32 s51, v9
+; GISEL-VGPR-NEXT: v_readfirstlane_b32 s0, v10
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[16:17], s[16:17]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[18:19], s[18:19]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[20:21], s[20:21]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[22:23], s[22:23]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[0:1], s[36:37]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[2:3], s[38:39]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[4:5], s[40:41]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[6:7], s[42:43]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[8:9], s[44:45]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[10:11], s[46:47]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[12:13], s[48:49]
+; GISEL-VGPR-NEXT: v_mov_b64_e32 v[14:15], s[50:51]
+; GISEL-VGPR-NEXT: v_mov_b32_e32 v28, s0
+; GISEL-VGPR-NEXT: v_readlane_b32 s51, v29, 7
+; GISEL-VGPR-NEXT: v_readlane_b32 s50, v29, 6
+; GISEL-VGPR-NEXT: v_smfmac_f32_32x32x32_bf16 v[0:15], v[24:27], v[16:23], v28
+; GISEL-VGPR-NEXT: v_readlane_b32 s49, v29, 5
+; GISEL-VGPR-NEXT: v_readlane_b32 s48, v29, 4
+; GISEL-VGPR-NEXT: v_readlane_b32 s39, v29, 3
+; GISEL-VGPR-NEXT: v_readlane_b32 s38, v29, 2
+; GISEL-VGPR-NEXT: v_readlane_b32 s37, v29, 1
+; GISEL-VGPR-NEXT: v_readlane_b32 s36, v29, 0
+; GISEL-VGPR-NEXT: s_xor_saveexec_b64 s[0:1], -1
+; GISEL-VGPR-NEXT: scratch_load_dword v29, off, s32 ; 4-byte Folded Reload
+; GISEL-VGPR-NEXT: s_mov_b64 exec, s[0:1]
+; GISEL-VGPR-NEXT: s_waitcnt vmcnt(0)
+; GISEL-VGPR-NEXT: s_setpc_b64 s[30:31]
%result = call <16 x float> @llvm.amdgcn.smfmac.f32.32x32x32.bf16(<8 x bfloat> %arg0, <16 x bfloat> %arg1, <16 x float> %arg2, i32 %arg3, i32 immarg 0, i32 immarg 0)
ret <16 x float> %result
}
More information about the llvm-commits
mailing list