[llvm] AMDGPU: Select VGPR MFMAs by default (PR #159493)

[Jeffrey Byrnes via llvm-commits]

================
@@ -17,115 +17,115 @@ define amdgpu_kernel void @test_mfma_f32_32x32x4bf16_1k(ptr addrspace(1) %arg) #
 ; GFX90A-LABEL: test_mfma_f32_32x32x4bf16_1k:
 ; GFX90A:       ; %bb.0: ; %bb
 ; GFX90A-NEXT:    s_load_dwordx2 s[34:35], s[4:5], 0x24
-; GFX90A-NEXT:    v_mov_b32_e32 v1, 0
-; GFX90A-NEXT:    v_mov_b32_e32 v2, 1
-; GFX90A-NEXT:    v_mov_b32_e32 v3, v1
-; GFX90A-NEXT:    v_mov_b32_e32 v0, 2
+; GFX90A-NEXT:    v_mov_b32_e32 v33, 0
+; GFX90A-NEXT:    v_mov_b32_e32 v34, 1
+; GFX90A-NEXT:    v_mov_b32_e32 v35, v33
+; GFX90A-NEXT:    v_mov_b32_e32 v32, 2
 ; GFX90A-NEXT:    s_waitcnt lgkmcnt(0)
 ; GFX90A-NEXT:    s_load_dwordx16 s[16:31], s[34:35], 0x0
 ; GFX90A-NEXT:    s_load_dwordx16 s[0:15], s[34:35], 0x40
 ; GFX90A-NEXT:    s_waitcnt lgkmcnt(0)
-; GFX90A-NEXT:    v_accvgpr_write_b32 a0, s16
-; GFX90A-NEXT:    v_accvgpr_write_b32 a1, s17
-; GFX90A-NEXT:    v_accvgpr_write_b32 a2, s18
-; GFX90A-NEXT:    v_accvgpr_write_b32 a3, s19
-; GFX90A-NEXT:    v_accvgpr_write_b32 a4, s20
-; GFX90A-NEXT:    v_accvgpr_write_b32 a5, s21
-; GFX90A-NEXT:    v_accvgpr_write_b32 a6, s22
-; GFX90A-NEXT:    v_accvgpr_write_b32 a7, s23
-; GFX90A-NEXT:    v_accvgpr_write_b32 a8, s24
-; GFX90A-NEXT:    v_accvgpr_write_b32 a9, s25
-; GFX90A-NEXT:    v_accvgpr_write_b32 a10, s26
-; GFX90A-NEXT:    v_accvgpr_write_b32 a11, s27
-; GFX90A-NEXT:    v_accvgpr_write_b32 a12, s28
-; GFX90A-NEXT:    v_accvgpr_write_b32 a13, s29
-; GFX90A-NEXT:    v_accvgpr_write_b32 a14, s30
-; GFX90A-NEXT:    v_accvgpr_write_b32 a15, s31
-; GFX90A-NEXT:    v_accvgpr_write_b32 a16, s0
-; GFX90A-NEXT:    v_accvgpr_write_b32 a17, s1
-; GFX90A-NEXT:    v_accvgpr_write_b32 a18, s2
-; GFX90A-NEXT:    v_accvgpr_write_b32 a19, s3
-; GFX90A-NEXT:    v_accvgpr_write_b32 a20, s4
-; GFX90A-NEXT:    v_accvgpr_write_b32 a21, s5
-; GFX90A-NEXT:    v_accvgpr_write_b32 a22, s6
-; GFX90A-NEXT:    v_accvgpr_write_b32 a23, s7
-; GFX90A-NEXT:    v_accvgpr_write_b32 a24, s8
-; GFX90A-NEXT:    v_accvgpr_write_b32 a25, s9
-; GFX90A-NEXT:    v_accvgpr_write_b32 a26, s10
-; GFX90A-NEXT:    v_accvgpr_write_b32 a27, s11
-; GFX90A-NEXT:    v_accvgpr_write_b32 a28, s12
-; GFX90A-NEXT:    v_accvgpr_write_b32 a29, s13
-; GFX90A-NEXT:    v_accvgpr_write_b32 a30, s14
-; GFX90A-NEXT:    v_accvgpr_write_b32 a31, s15
+; GFX90A-NEXT:    v_mov_b32_e32 v0, s16
+; GFX90A-NEXT:    v_mov_b32_e32 v1, s17
+; GFX90A-NEXT:    v_mov_b32_e32 v2, s18
+; GFX90A-NEXT:    v_mov_b32_e32 v3, s19
+; GFX90A-NEXT:    v_mov_b32_e32 v4, s20
+; GFX90A-NEXT:    v_mov_b32_e32 v5, s21
+; GFX90A-NEXT:    v_mov_b32_e32 v6, s22
+; GFX90A-NEXT:    v_mov_b32_e32 v7, s23
+; GFX90A-NEXT:    v_mov_b32_e32 v8, s24
+; GFX90A-NEXT:    v_mov_b32_e32 v9, s25
+; GFX90A-NEXT:    v_mov_b32_e32 v10, s26
+; GFX90A-NEXT:    v_mov_b32_e32 v11, s27
+; GFX90A-NEXT:    v_mov_b32_e32 v12, s28
+; GFX90A-NEXT:    v_mov_b32_e32 v13, s29
+; GFX90A-NEXT:    v_mov_b32_e32 v14, s30
+; GFX90A-NEXT:    v_mov_b32_e32 v15, s31
+; GFX90A-NEXT:    v_mov_b32_e32 v16, s0
+; GFX90A-NEXT:    v_mov_b32_e32 v17, s1
+; GFX90A-NEXT:    v_mov_b32_e32 v18, s2
+; GFX90A-NEXT:    v_mov_b32_e32 v19, s3
+; GFX90A-NEXT:    v_mov_b32_e32 v20, s4
+; GFX90A-NEXT:    v_mov_b32_e32 v21, s5
+; GFX90A-NEXT:    v_mov_b32_e32 v22, s6
+; GFX90A-NEXT:    v_mov_b32_e32 v23, s7
+; GFX90A-NEXT:    v_mov_b32_e32 v24, s8
+; GFX90A-NEXT:    v_mov_b32_e32 v25, s9
+; GFX90A-NEXT:    v_mov_b32_e32 v26, s10
+; GFX90A-NEXT:    v_mov_b32_e32 v27, s11
+; GFX90A-NEXT:    v_mov_b32_e32 v28, s12
+; GFX90A-NEXT:    v_mov_b32_e32 v29, s13
+; GFX90A-NEXT:    v_mov_b32_e32 v30, s14
+; GFX90A-NEXT:    v_mov_b32_e32 v31, s15
 ; GFX90A-NEXT:    s_nop 1
-; GFX90A-NEXT:    v_mfma_f32_32x32x4bf16_1k a[0:31], v[2:3], v[0:1], a[0:31] cbsz:1 abid:2 blgp:3
+; GFX90A-NEXT:    v_mfma_f32_32x32x4bf16_1k v[0:31], v[34:35], v[32:33], v[0:31] cbsz:1 abid:2 blgp:3
 ; GFX90A-NEXT:    s_nop 15
 ; GFX90A-NEXT:    s_nop 2
-; GFX90A-NEXT:    global_store_dwordx4 v1, a[24:27], s[34:35] offset:96
-; GFX90A-NEXT:    global_store_dwordx4 v1, a[28:31], s[34:35] offset:112
-; GFX90A-NEXT:    global_store_dwordx4 v1, a[16:19], s[34:35] offset:64
-; GFX90A-NEXT:    global_store_dwordx4 v1, a[20:23], s[34:35] offset:80
-; GFX90A-NEXT:    global_store_dwordx4 v1, a[8:11], s[34:35] offset:32
-; GFX90A-NEXT:    global_store_dwordx4 v1, a[12:15], s[34:35] offset:48
-; GFX90A-NEXT:    global_store_dwordx4 v1, a[0:3], s[34:35]
-; GFX90A-NEXT:    global_store_dwordx4 v1, a[4:7], s[34:35] offset:16
+; GFX90A-NEXT:    global_store_dwordx4 v33, v[24:27], s[34:35] offset:96
----------------
jrbyrnes wrote:

Not suggesting we need to change anything here, but these changes are misleading I think since the test file does not use amdgpu-agpr-alloc (which the attributor will attach)

https://github.com/llvm/llvm-project/pull/159493