AMDGPU: Add tests with forced selection of mfmas to VGPR form #150626

arsenm · 2025-07-25T14:52:36Z

Add some run lines to existing tests with VGPR selection enabled.

arsenm · 2025-07-25T14:52:53Z

This stack of pull requests is managed by Graphite. Learn more about stacking.

llvmbot · 2025-07-25T14:53:27Z

@llvm/pr-subscribers-backend-amdgpu

Author: Matt Arsenault (arsenm)

Changes

Add some run lines to existing tests with VGPR selection enabled.

Patch is 105.62 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/150626.diff

2 Files Affected:

(modified) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.gfx90a.ll (+909)
(modified) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.ll (+796)

diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.gfx90a.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.gfx90a.ll
index 780c7e93f80d0..adbc2df4b3474 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.gfx90a.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.gfx90a.ll
@@ -1,6 +1,8 @@
 ; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
 ; RUN: llc -mtriple=amdgcn -mcpu=gfx90a < %s | FileCheck -enable-var-scope --check-prefixes=GCN,GFX90A %s
 ; RUN: llc -mtriple=amdgcn -mcpu=gfx942 < %s | FileCheck -enable-var-scope --check-prefixes=GCN,GFX942 %s
+; RUN: llc -mtriple=amdgcn -mcpu=gfx90a -amdgpu-mfma-vgpr-form < %s | FileCheck -enable-var-scope --check-prefixes=VGPR,GFX90A-VGPR %s
+; RUN: llc -mtriple=amdgcn -mcpu=gfx942 -amdgpu-mfma-vgpr-form < %s | FileCheck -enable-var-scope --check-prefixes=VGPR,GFX942-VGPR %s
 
 declare <32 x float> @llvm.amdgcn.mfma.f32.32x32x4bf16.1k(<4 x i16>, <4 x i16>, <32 x float>, i32, i32, i32)
 declare <16 x float> @llvm.amdgcn.mfma.f32.16x16x4bf16.1k(<4 x i16>, <4 x i16>, <16 x float>, i32, i32, i32)
@@ -127,6 +129,122 @@ define amdgpu_kernel void @test_mfma_f32_32x32x4bf16_1k(ptr addrspace(1) %arg) #
 ; GFX942-NEXT:    global_store_dwordx4 v1, a[0:3], s[34:35]
 ; GFX942-NEXT:    global_store_dwordx4 v1, a[4:7], s[34:35] offset:16
 ; GFX942-NEXT:    s_endpgm
+;
+; GFX90A-VGPR-LABEL: test_mfma_f32_32x32x4bf16_1k:
+; GFX90A-VGPR:       ; %bb.0: ; %bb
+; GFX90A-VGPR-NEXT:    s_load_dwordx2 s[34:35], s[4:5], 0x24
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v33, 0
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v34, 1
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v35, v33
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v32, 2
+; GFX90A-VGPR-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX90A-VGPR-NEXT:    s_load_dwordx16 s[16:31], s[34:35], 0x0
+; GFX90A-VGPR-NEXT:    s_load_dwordx16 s[0:15], s[34:35], 0x40
+; GFX90A-VGPR-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v0, s16
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v1, s17
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v2, s18
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v3, s19
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v4, s20
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v5, s21
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v6, s22
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v7, s23
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v8, s24
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v9, s25
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v10, s26
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v11, s27
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v12, s28
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v13, s29
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v14, s30
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v15, s31
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v16, s0
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v17, s1
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v18, s2
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v19, s3
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v20, s4
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v21, s5
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v22, s6
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v23, s7
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v24, s8
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v25, s9
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v26, s10
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v27, s11
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v28, s12
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v29, s13
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v30, s14
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v31, s15
+; GFX90A-VGPR-NEXT:    s_nop 1
+; GFX90A-VGPR-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-VGPR-NEXT:    s_nop 7
+; GFX90A-VGPR-NEXT:    s_nop 7
+; GFX90A-VGPR-NEXT:    s_nop 2
+; GFX90A-VGPR-NEXT:    global_store_dwordx4 v33, v[24:27], s[34:35] offset:96
+; GFX90A-VGPR-NEXT:    global_store_dwordx4 v33, v[28:31], s[34:35] offset:112
+; GFX90A-VGPR-NEXT:    global_store_dwordx4 v33, v[16:19], s[34:35] offset:64
+; GFX90A-VGPR-NEXT:    global_store_dwordx4 v33, v[20:23], s[34:35] offset:80
+; GFX90A-VGPR-NEXT:    global_store_dwordx4 v33, v[8:11], s[34:35] offset:32
+; GFX90A-VGPR-NEXT:    global_store_dwordx4 v33, v[12:15], s[34:35] offset:48
+; GFX90A-VGPR-NEXT:    global_store_dwordx4 v33, v[0:3], s[34:35]
+; GFX90A-VGPR-NEXT:    global_store_dwordx4 v33, v[4:7], s[34:35] offset:16
+; GFX90A-VGPR-NEXT:    s_endpgm
+;
+; GFX942-VGPR-LABEL: test_mfma_f32_32x32x4bf16_1k:
+; GFX942-VGPR:       ; %bb.0: ; %bb
+; GFX942-VGPR-NEXT:    s_load_dwordx2 s[34:35], s[4:5], 0x24
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v33, 0
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v34, 1
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v35, v33
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v32, 2
+; GFX942-VGPR-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX942-VGPR-NEXT:    s_load_dwordx16 s[16:31], s[34:35], 0x0
+; GFX942-VGPR-NEXT:    s_load_dwordx16 s[0:15], s[34:35], 0x40
+; GFX942-VGPR-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v0, s16
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v1, s17
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v2, s18
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v3, s19
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v4, s20
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v5, s21
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v6, s22
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v7, s23
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v8, s24
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v9, s25
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v10, s26
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v11, s27
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v12, s28
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v13, s29
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v14, s30
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v15, s31
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v16, s0
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v17, s1
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v18, s2
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v19, s3
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v20, s4
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v21, s5
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v22, s6
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v23, s7
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v24, s8
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v25, s9
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v26, s10
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v27, s11
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v28, s12
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v29, s13
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v30, s14
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v31, s15
+; GFX942-VGPR-NEXT:    s_nop 1
+; GFX942-VGPR-NEXT:    v_mfma_f32_32x32x4_2b_bf16 v[0:31], v[34:35], v[32:33], v[0:31] cbsz:1 abid:2 blgp:3
+; GFX942-VGPR-NEXT:    s_nop 7
+; GFX942-VGPR-NEXT:    s_nop 7
+; GFX942-VGPR-NEXT:    s_nop 2
+; GFX942-VGPR-NEXT:    global_store_dwordx4 v33, v[24:27], s[34:35] offset:96
+; GFX942-VGPR-NEXT:    global_store_dwordx4 v33, v[28:31], s[34:35] offset:112
+; GFX942-VGPR-NEXT:    global_store_dwordx4 v33, v[16:19], s[34:35] offset:64
+; GFX942-VGPR-NEXT:    global_store_dwordx4 v33, v[20:23], s[34:35] offset:80
+; GFX942-VGPR-NEXT:    global_store_dwordx4 v33, v[8:11], s[34:35] offset:32
+; GFX942-VGPR-NEXT:    global_store_dwordx4 v33, v[12:15], s[34:35] offset:48
+; GFX942-VGPR-NEXT:    global_store_dwordx4 v33, v[0:3], s[34:35]
+; GFX942-VGPR-NEXT:    global_store_dwordx4 v33, v[4:7], s[34:35] offset:16
+; GFX942-VGPR-NEXT:    s_endpgm
 bb:
   %in.1 = load <32 x float>, ptr addrspace(1) %arg
   %a = bitcast i64 1 to <4 x i16>
@@ -208,6 +326,62 @@ define amdgpu_kernel void @test_mfma_f32_16x16x4bf16_1k(ptr addrspace(1) %arg) #
 ; GFX942-NEXT:    global_store_dwordx4 v1, a[4:7], s[16:17] offset:16
 ; GFX942-NEXT:    global_store_dwordx4 v1, a[0:3], s[16:17]
 ; GFX942-NEXT:    s_endpgm
+;
+; GFX90A-VGPR-LABEL: test_mfma_f32_16x16x4bf16_1k:
+; GFX90A-VGPR:       ; %bb.0: ; %bb
+; GFX90A-VGPR-NEXT:    s_load_dwordx2 s[16:17], s[4:5], 0x24
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v17, 0
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v18, 1
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v19, v17
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v16, 2
+; GFX90A-VGPR-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX90A-VGPR-NEXT:    s_load_dwordx16 s[0:15], s[16:17], 0x0
+; GFX90A-VGPR-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX90A-VGPR-NEXT:    v_pk_mov_b32 v[0:1], s[0:1], s[0:1] op_sel:[0,1]
+; GFX90A-VGPR-NEXT:    v_pk_mov_b32 v[2:3], s[2:3], s[2:3] op_sel:[0,1]
+; GFX90A-VGPR-NEXT:    v_pk_mov_b32 v[4:5], s[4:5], s[4:5] op_sel:[0,1]
+; GFX90A-VGPR-NEXT:    v_pk_mov_b32 v[6:7], s[6:7], s[6:7] op_sel:[0,1]
+; GFX90A-VGPR-NEXT:    v_pk_mov_b32 v[8:9], s[8:9], s[8:9] op_sel:[0,1]
+; GFX90A-VGPR-NEXT:    v_pk_mov_b32 v[10:11], s[10:11], s[10:11] op_sel:[0,1]
+; GFX90A-VGPR-NEXT:    v_pk_mov_b32 v[12:13], s[12:13], s[12:13] op_sel:[0,1]
+; GFX90A-VGPR-NEXT:    v_pk_mov_b32 v[14:15], s[14:15], s[14:15] op_sel:[0,1]
+; GFX90A-VGPR-NEXT:    s_nop 1
+; GFX90A-VGPR-NEXT:    v_mfma_f32_16x16x4bf16_1k v[0:15], v[18:19], v[16:17], v[0:15] cbsz:1 abid:2 blgp:3
+; GFX90A-VGPR-NEXT:    s_nop 7
+; GFX90A-VGPR-NEXT:    s_nop 2
+; GFX90A-VGPR-NEXT:    global_store_dwordx4 v17, v[12:15], s[16:17] offset:48
+; GFX90A-VGPR-NEXT:    global_store_dwordx4 v17, v[8:11], s[16:17] offset:32
+; GFX90A-VGPR-NEXT:    global_store_dwordx4 v17, v[4:7], s[16:17] offset:16
+; GFX90A-VGPR-NEXT:    global_store_dwordx4 v17, v[0:3], s[16:17]
+; GFX90A-VGPR-NEXT:    s_endpgm
+;
+; GFX942-VGPR-LABEL: test_mfma_f32_16x16x4bf16_1k:
+; GFX942-VGPR:       ; %bb.0: ; %bb
+; GFX942-VGPR-NEXT:    s_load_dwordx2 s[16:17], s[4:5], 0x24
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v17, 0
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v18, 1
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v19, v17
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v16, 2
+; GFX942-VGPR-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX942-VGPR-NEXT:    s_load_dwordx16 s[0:15], s[16:17], 0x0
+; GFX942-VGPR-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX942-VGPR-NEXT:    v_mov_b64_e32 v[0:1], s[0:1]
+; GFX942-VGPR-NEXT:    v_mov_b64_e32 v[2:3], s[2:3]
+; GFX942-VGPR-NEXT:    v_mov_b64_e32 v[4:5], s[4:5]
+; GFX942-VGPR-NEXT:    v_mov_b64_e32 v[6:7], s[6:7]
+; GFX942-VGPR-NEXT:    v_mov_b64_e32 v[8:9], s[8:9]
+; GFX942-VGPR-NEXT:    v_mov_b64_e32 v[10:11], s[10:11]
+; GFX942-VGPR-NEXT:    v_mov_b64_e32 v[12:13], s[12:13]
+; GFX942-VGPR-NEXT:    v_mov_b64_e32 v[14:15], s[14:15]
+; GFX942-VGPR-NEXT:    s_nop 1
+; GFX942-VGPR-NEXT:    v_mfma_f32_16x16x4_4b_bf16 v[0:15], v[18:19], v[16:17], v[0:15] cbsz:1 abid:2 blgp:3
+; GFX942-VGPR-NEXT:    s_nop 7
+; GFX942-VGPR-NEXT:    s_nop 2
+; GFX942-VGPR-NEXT:    global_store_dwordx4 v17, v[12:15], s[16:17] offset:48
+; GFX942-VGPR-NEXT:    global_store_dwordx4 v17, v[8:11], s[16:17] offset:32
+; GFX942-VGPR-NEXT:    global_store_dwordx4 v17, v[4:7], s[16:17] offset:16
+; GFX942-VGPR-NEXT:    global_store_dwordx4 v17, v[0:3], s[16:17]
+; GFX942-VGPR-NEXT:    s_endpgm
 bb:
   %in.1 = load <16 x float>, ptr addrspace(1) %arg
   %a = bitcast i64 1 to <4 x i16>
@@ -257,6 +431,42 @@ define amdgpu_kernel void @test_mfma_f32_4x4x4bf16_1k(ptr addrspace(1) %arg) #0
 ; GFX942-NEXT:    s_nop 4
 ; GFX942-NEXT:    global_store_dwordx4 v1, a[0:3], s[6:7]
 ; GFX942-NEXT:    s_endpgm
+;
+; GFX90A-VGPR-LABEL: test_mfma_f32_4x4x4bf16_1k:
+; GFX90A-VGPR:       ; %bb.0: ; %bb
+; GFX90A-VGPR-NEXT:    s_load_dwordx2 s[6:7], s[4:5], 0x24
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v5, 0
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v6, 1
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v7, v5
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v4, 2
+; GFX90A-VGPR-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX90A-VGPR-NEXT:    s_load_dwordx4 s[0:3], s[6:7], 0x0
+; GFX90A-VGPR-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX90A-VGPR-NEXT:    v_pk_mov_b32 v[0:1], s[0:1], s[0:1] op_sel:[0,1]
+; GFX90A-VGPR-NEXT:    v_pk_mov_b32 v[2:3], s[2:3], s[2:3] op_sel:[0,1]
+; GFX90A-VGPR-NEXT:    s_nop 1
+; GFX90A-VGPR-NEXT:    v_mfma_f32_4x4x4bf16_1k v[0:3], v[6:7], v[4:5], v[0:3] cbsz:1 abid:2 blgp:3
+; GFX90A-VGPR-NEXT:    s_nop 4
+; GFX90A-VGPR-NEXT:    global_store_dwordx4 v5, v[0:3], s[6:7]
+; GFX90A-VGPR-NEXT:    s_endpgm
+;
+; GFX942-VGPR-LABEL: test_mfma_f32_4x4x4bf16_1k:
+; GFX942-VGPR:       ; %bb.0: ; %bb
+; GFX942-VGPR-NEXT:    s_load_dwordx2 s[6:7], s[4:5], 0x24
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v5, 0
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v6, 1
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v7, v5
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v4, 2
+; GFX942-VGPR-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX942-VGPR-NEXT:    s_load_dwordx4 s[0:3], s[6:7], 0x0
+; GFX942-VGPR-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX942-VGPR-NEXT:    v_mov_b64_e32 v[0:1], s[0:1]
+; GFX942-VGPR-NEXT:    v_mov_b64_e32 v[2:3], s[2:3]
+; GFX942-VGPR-NEXT:    s_nop 1
+; GFX942-VGPR-NEXT:    v_mfma_f32_4x4x4_16b_bf16 v[0:3], v[6:7], v[4:5], v[0:3] cbsz:1 abid:2 blgp:3
+; GFX942-VGPR-NEXT:    s_nop 4
+; GFX942-VGPR-NEXT:    global_store_dwordx4 v5, v[0:3], s[6:7]
+; GFX942-VGPR-NEXT:    s_endpgm
 bb:
   %in.1 = load <4 x float>, ptr addrspace(1) %arg
   %a = bitcast i64 1 to <4 x i16>
@@ -339,6 +549,63 @@ define amdgpu_kernel void @test_mfma_f32_32x32x8bf16_1k(ptr addrspace(1) %arg) #
 ; GFX942-NEXT:    global_store_dwordx4 v1, a[4:7], s[16:17] offset:16
 ; GFX942-NEXT:    global_store_dwordx4 v1, a[0:3], s[16:17]
 ; GFX942-NEXT:    s_endpgm
+;
+; GFX90A-VGPR-LABEL: test_mfma_f32_32x32x8bf16_1k:
+; GFX90A-VGPR:       ; %bb.0: ; %bb
+; GFX90A-VGPR-NEXT:    s_load_dwordx2 s[16:17], s[4:5], 0x24
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v17, 0
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v18, 1
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v19, v17
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v16, 2
+; GFX90A-VGPR-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX90A-VGPR-NEXT:    s_load_dwordx16 s[0:15], s[16:17], 0x0
+; GFX90A-VGPR-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX90A-VGPR-NEXT:    v_pk_mov_b32 v[0:1], s[0:1], s[0:1] op_sel:[0,1]
+; GFX90A-VGPR-NEXT:    v_pk_mov_b32 v[2:3], s[2:3], s[2:3] op_sel:[0,1]
+; GFX90A-VGPR-NEXT:    v_pk_mov_b32 v[4:5], s[4:5], s[4:5] op_sel:[0,1]
+; GFX90A-VGPR-NEXT:    v_pk_mov_b32 v[6:7], s[6:7], s[6:7] op_sel:[0,1]
+; GFX90A-VGPR-NEXT:    v_pk_mov_b32 v[8:9], s[8:9], s[8:9] op_sel:[0,1]
+; GFX90A-VGPR-NEXT:    v_pk_mov_b32 v[10:11], s[10:11], s[10:11] op_sel:[0,1]
+; GFX90A-VGPR-NEXT:    v_pk_mov_b32 v[12:13], s[12:13], s[12:13] op_sel:[0,1]
+; GFX90A-VGPR-NEXT:    v_pk_mov_b32 v[14:15], s[14:15], s[14:15] op_sel:[0,1]
+; GFX90A-VGPR-NEXT:    s_nop 1
+; GFX90A-VGPR-NEXT:    v_mfma_f32_32x32x8bf16_1k v[0:15], v[18:19], v[16:17], v[0:15] cbsz:1 abid:2 blgp:3
+; GFX90A-VGPR-NEXT:    s_nop 7
+; GFX90A-VGPR-NEXT:    s_nop 7
+; GFX90A-VGPR-NEXT:    s_nop 2
+; GFX90A-VGPR-NEXT:    global_store_dwordx4 v17, v[12:15], s[16:17] offset:48
+; GFX90A-VGPR-NEXT:    global_store_dwordx4 v17, v[8:11], s[16:17] offset:32
+; GFX90A-VGPR-NEXT:    global_store_dwordx4 v17, v[4:7], s[16:17] offset:16
+; GFX90A-VGPR-NEXT:    global_store_dwordx4 v17, v[0:3], s[16:17]
+; GFX90A-VGPR-NEXT:    s_endpgm
+;
+; GFX942-VGPR-LABEL: test_mfma_f32_32x32x8bf16_1k:
+; GFX942-VGPR:       ; %bb.0: ; %bb
+; GFX942-VGPR-NEXT:    s_load_dwordx2 s[16:17], s[4:5], 0x24
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v17, 0
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v18, 1
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v19, v17
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v16, 2
+; GFX942-VGPR-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX942-VGPR-NEXT:    s_load_dwordx16 s[0:15], s[16:17], 0x0
+; GFX942-VGPR-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX942-VGPR-NEXT:    v_mov_b64_e32 v[0:1], s[0:1]
+; GFX942-VGPR-NEXT:    v_mov_b64_e32 v[2:3], s[2:3]
+; GFX942-VGPR-NEXT:    v_mov_b64_e32 v[4:5], s[4:5]
+; GFX942-VGPR-NEXT:    v_mov_b64_e32 v[6:7], s[6:7]
+; GFX942-VGPR-NEXT:    v_mov_b64_e32 v[8:9], s[8:9]
+; GFX942-VGPR-NEXT:    v_mov_b64_e32 v[10:11], s[10:11]
+; GFX942-VGPR-NEXT:    v_mov_b64_e32 v[12:13], s[12:13]
+; GFX942-VGPR-NEXT:    v_mov_b64_e32 v[14:15], s[14:15]
+; GFX942-VGPR-NEXT:    s_nop 1
+; GFX942-VGPR-NEXT:    v_mfma_f32_32x32x8_bf16 v[0:15], v[18:19], v[16:17], v[0:15] cbsz:1 abid:2 blgp:3
+; GFX942-VGPR-NEXT:    s_nop 7
+; GFX942-VGPR-NEXT:    s_nop 2
+; GFX942-VGPR-NEXT:    global_store_dwordx4 v17, v[12:15], s[16:17] offset:48
+; GFX942-VGPR-NEXT:    global_store_dwordx4 v17, v[8:11], s[16:17] offset:32
+; GFX942-VGPR-NEXT:    global_store_dwordx4 v17, v[4:7], s[16:17] offset:16
+; GFX942-VGPR-NEXT:    global_store_dwordx4 v17, v[0:3], s[16:17]
+; GFX942-VGPR-NEXT:    s_endpgm
 bb:
   %in.1 = load <16 x float>, ptr addrspace(1) %arg
   %a = bitcast i64 1 to <4 x i16>
@@ -389,6 +656,43 @@ define amdgpu_kernel void @test_mfma_f32_16x16x16bf16_1k(ptr addrspace(1) %arg)
 ; GFX942-NEXT:    s_nop 6
 ; GFX942-NEXT:    global_store_dwordx4 v1, a[0:3], s[6:7]
 ; GFX942-NEXT:    s_endpgm
+;
+; GFX90A-VGPR-LABEL: test_mfma_f32_16x16x16bf16_1k:
+; GFX90A-VGPR:       ; %bb.0: ; %bb
+; GFX90A-VGPR-NEXT:    s_load_dwordx2 s[6:7], s[4:5], 0x24
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v5, 0
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v6, 1
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v7, v5
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v4, 2
+; GFX90A-VGPR-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX90A-VGPR-NEXT:    s_load_dwordx4 s[0:3], s[6:7], 0x0
+; GFX90A-VGPR-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX90A-VGPR-NEXT:    v_pk_mov_b32 v[0:1], s[0:1], s[0:1] op_sel:[0,1]
+; GFX90A-VGPR-NEXT:    v_pk_mov_b32 v[2:3], s[2:3], s[2:3] op_sel:[0,1]
+; GFX90A-VGPR-NEXT:    s_nop 1
+; GFX90A-VGPR-NEXT:    v_mfma_f32_16x16x16bf16_1k v[0:3], v[6:7], v[4:5], v[0:3] cbsz:1 abid:2 blgp:3
+; GFX90A-VGPR-NEXT:    s_nop 7
+; GFX90A-VGPR-NEXT:    s_nop 2
+; GFX90A-VGPR-NEXT:    global_store_dwordx4 v5, v[0:3], s[6:7]
+; GFX90A-VGPR-NEXT:    s_endpgm
+;
+; GFX942-VGPR-LABEL: test_mfma_f32_16x16x16bf16_1k:
+; GFX942-VGPR:       ; %bb.0: ; %bb
+; GFX942-VGPR-NEXT:    s_load_dwordx2 s[6:7], s[4:5], 0x24
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v5, 0
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v6, 1
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v7, v5
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v4, 2
+; GFX942-VGPR-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX942-VGPR-NEXT:    s_load_dwordx4 s[0:3], s[6:7], 0x0
+; GFX942-VGPR-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX942-VGPR-NEXT:    v_mov_b64_e32 v[0:1], s[0:1]
+; GFX942-VGPR-NEXT:    v_mov_b64_e32 v[2:3], s[2:3]
+; GFX942-VGPR-NEXT:    s_nop 1
+; GFX942-VGPR-NEXT:    v_mfma_f32_16x16x16_bf16 v[0:3], v[6:7], v[4:5], v[0:3] cbsz:1 abid:2 blgp:3
+; GFX942-VGPR-NEXT:    s_nop 6
+; GFX942-VGPR-NEXT:    global_store_dwordx4 v5, v[0:3], s[6:7]
+; GFX942-VGPR-NEXT:    s_endpgm
 bb:
   %in.1 = load <4 x float>, ptr addrspace(1) %arg
   %a = bitcast i64 1 to <4 x i16>
@@ -430,6 +734,38 @@ define amdgpu_kernel void @test_mfma_f64_4x4x4f64(ptr addrspace(1) %arg, double
 ; GFX942-NEXT:    s_nop 7
 ; GFX942-NEXT:    global_store_dwordx2 v0, a[0:1], s[0:1]
 ; GFX942-NEXT:    s_endpgm
+;
+; GFX90A-VGPR-LABEL: test_mfma_f64_4x4x4f64:
+; GFX90A-VGPR:       ; %bb.0: ; %bb
+; GFX90A-VGPR-NEXT:    s_load_dwordx4 s[0:3], s[4:5], 0x24
+; GFX90A-VGPR-NEXT:    s_load_dwordx2 s[6:7], s[4:5], 0x34
+; GFX90A-VGPR-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX90A-VGPR-NEXT:    v_pk_mov_b32 v[0:1], s[2:3], s[2:3] op_sel:[0,1]
+; GFX90A-VGPR-NEXT:    v_pk_mov_b32 v[2:3], s[6:7], s[6:7] op_sel:[0,1]
+; GFX90A-VGPR-NEXT:    s_nop 1
+; GFX90A-VGPR-NEXT:    v_mfma_f64_4x4x4f64 v[4:5], v[0:1], v[2:3], 0
+; GFX90A-VGPR-NEXT:    s_nop 3
+; GFX90A-VGPR-NEXT:    v_mfma_f64_4x4x4f64 v[0:1], v[0:1], v[2:3], v[4:5] cbsz:1 abid:2 blgp:3
+; GFX90A-VGPR-NEXT:    v_mov_b32_e32 v2, 0
+; GFX90A-VGPR-NEXT:    s_nop 7
+; GFX90A-VGPR-NEXT:    global_store_dwordx2 v2, v[0:1], s[0:1]
+; GFX90A-VGPR-NEXT:    s_endpgm
+;
+; GFX942-VGPR-LABEL: test_mfma_f64_4x4x4f64:
+; GFX942-VGPR:       ; %bb.0: ; %bb
+; GFX942-VGPR-NEXT:    s_load_dwordx4 s[0:3], s[4:5], 0x24
+; GFX942-VGPR-NEXT:    s_load_dwordx2 s[6:7], s[4:5], 0x34
+; GFX942-VGPR-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX942-VGPR-NEXT:    v_mov_b64_e32 v[0:1], s[2:3]
+; GFX942-VGPR-NEXT:    v_mov_b64_e32 v[2:3], s[6:7]
+; GFX942-VGPR-NEXT:    s_nop 1
+; GFX942-VGPR-NEXT:    v_mfma_f64_4x4x4_4b_f64 v[4:5], v[0:1], v[2:3], 0
+; GFX942-VGPR-NEXT:    s_nop 3
+; GFX942-VGPR-NEXT:    v_mfma_f64_4x4x4_4b_f64 v[0:1], v[0:1], v[2:3], v[4:5] cbsz:1 abid:2 neg:[1,1,0]
+; GFX942-VGPR-NEXT:    v_mov_b32_e32 v2, 0
+; GFX942-VGPR-NEXT:    s_nop 7
+; GFX942-VGPR-NEXT:    global_store_dwordx2 v2, v[0:1], s[0:1]
+; GFX942-VGPR-NEXT:    s_endpgm
 bb:
   %mai.1 = tail call double @llvm.amdgcn.mfma.f64.4x4x4f64(double %a, double %b, double 0.0, i32 0, i32 0, i32 0)
   %mai.2 = tail call double @llvm.amdgcn.mfma.f64.4x4x4f64(double %a, double %b, double %mai.1, i32 1, i32 2, i32 3)
@@ -493,6 +829,54 @@ define amdgpu_kernel void @test_mfma_f64_16x16x4f64(ptr addrspace(1) %arg, doubl
 ; GFX942-NEX...
[truncated]

arsenm · 2025-07-25T16:01:17Z

Merge activity

Jul 25, 4:01 PM UTC: A user started a stack merge that includes this pull request via Graphite.
Jul 25, 4:03 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 4:09 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 4:13 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 4:19 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 4:24 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 4:27 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 4:34 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 4:44 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 4:50 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 4:53 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 5:18 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 5:25 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 5:29 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 5:32 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 5:35 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 5:38 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 5:42 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 5:46 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 5:58 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 6:09 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 6:17 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 6:26 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 6:29 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 6:47 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 6:53 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 7:02 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 7:05 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 7:08 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 7:13 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 7:16 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 7:24 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 7:27 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 7:33 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 7:36 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 7:46 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 7:49 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 7:53 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 8:05 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 8:08 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 8:11 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 8:18 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 8:25 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 8:43 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 8:49 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 8:55 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 9:16 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 9:28 PM UTC: Graphite rebased this pull request as part of a merge.
Jul 25, 9:43 PM UTC: Graphite couldn't merge this PR because it was not satisfying all requirements (Failed CI: 'Build and Test Linux').
Jul 26, 12:18 AM UTC: Graphite rebased this pull request as part of a merge.
Jul 26, 12:46 AM UTC: @arsenm merged this pull request with Graphite.

shiltian · 2025-07-25T20:28:25Z

It seems like graphite can no longer directly bypass the checks after the confirmation checkbox is added. It keeps rebasing this branch.

Add some run lines to existing tests with VGPR selection enabled.

llvm-ci · 2025-07-26T00:51:49Z

LLVM Buildbot has detected a new failure on builder lldb-x86_64-debian running on lldb-x86_64-debian while building llvm at step 6 "test".

Full details are available at: https://lab.llvm.org/buildbot/#/builders/162/builds/27616

Here is the relevant piece of the build log for the reference

Step 6 (test) failure: build (failure)
...
UNSUPPORTED: lldb-shell :: Process/Optimization.test (3089 of 3101)
UNSUPPORTED: lldb-shell :: ScriptInterpreter/Python/Crashlog/interactive_crashlog_arm64_register.test (3090 of 3101)
UNSUPPORTED: lldb-shell :: RPC/Generator/Tests/Server/CheckConstCharPointer.test (3091 of 3101)
UNSUPPORTED: lldb-shell :: ScriptInterpreter/Python/Crashlog/no-args.test (3092 of 3101)
PASS: lldb-unit :: Utility/./UtilityTests/115/129 (3093 of 3101)
UNSUPPORTED: lldb-shell :: ScriptInterpreter/Lua/quit.test (3094 of 3101)
UNSUPPORTED: lldb-shell :: SymbolFile/NativePDB/stack_unwinding01.cpp (3095 of 3101)
UNSUPPORTED: lldb-shell :: ScriptInterpreter/Python/Crashlog/text.test (3096 of 3101)
PASS: lldb-api :: terminal/TestEditlineCompletions.py (3097 of 3101)
UNRESOLVED: lldb-api :: tools/lldb-dap/launch/TestDAP_launch.py (3098 of 3101)
******************** TEST 'lldb-api :: tools/lldb-dap/launch/TestDAP_launch.py' FAILED ********************
Script:
--
/usr/bin/python3 /home/worker/2.0.1/lldb-x86_64-debian/llvm-project/lldb/test/API/dotest.py -u CXXFLAGS -u CFLAGS --env LLVM_LIBS_DIR=/home/worker/2.0.1/lldb-x86_64-debian/build/./lib --env LLVM_INCLUDE_DIR=/home/worker/2.0.1/lldb-x86_64-debian/build/include --env LLVM_TOOLS_DIR=/home/worker/2.0.1/lldb-x86_64-debian/build/./bin --arch x86_64 --build-dir /home/worker/2.0.1/lldb-x86_64-debian/build/lldb-test-build.noindex --lldb-module-cache-dir /home/worker/2.0.1/lldb-x86_64-debian/build/lldb-test-build.noindex/module-cache-lldb/lldb-api --clang-module-cache-dir /home/worker/2.0.1/lldb-x86_64-debian/build/lldb-test-build.noindex/module-cache-clang/lldb-api --executable /home/worker/2.0.1/lldb-x86_64-debian/build/./bin/lldb --compiler /home/worker/2.0.1/lldb-x86_64-debian/build/./bin/clang --dsymutil /home/worker/2.0.1/lldb-x86_64-debian/build/./bin/dsymutil --make /usr/bin/gmake --llvm-tools-dir /home/worker/2.0.1/lldb-x86_64-debian/build/./bin --lldb-obj-root /home/worker/2.0.1/lldb-x86_64-debian/build/tools/lldb --lldb-libs-dir /home/worker/2.0.1/lldb-x86_64-debian/build/./lib --cmake-build-type Release -t /home/worker/2.0.1/lldb-x86_64-debian/llvm-project/lldb/test/API/tools/lldb-dap/launch -p TestDAP_launch.py
--
Exit Code: 1

Command Output (stdout):
--
lldb version 22.0.0git (https://github.com/llvm/llvm-project.git revision e4d5969d7181e39d549e344fcf2ec2b61a77947e)
  clang revision e4d5969d7181e39d549e344fcf2ec2b61a77947e
  llvm revision e4d5969d7181e39d549e344fcf2ec2b61a77947e
Skipping the following test categories: ['libc++', 'msvcstl', 'dsym', 'gmodules', 'debugserver', 'objc']

--
Command Output (stderr):
--
Change dir to: /home/worker/2.0.1/lldb-x86_64-debian/llvm-project/lldb/test/API/tools/lldb-dap/launch
runCmd: settings clear --all

output: 

runCmd: settings set symbols.enable-external-lookup false

output: 

runCmd: settings set target.inherit-tcc true

output: 

runCmd: settings set target.disable-aslr false

output: 

runCmd: settings set target.detach-on-error false

output: 

runCmd: settings set target.auto-apply-fixits false

llvm-ci · 2025-07-26T00:56:29Z

LLVM Buildbot has detected a new failure on builder clang-aarch64-quick running on linaro-clang-aarch64-quick while building llvm at step 5 "ninja check 1".

Full details are available at: https://lab.llvm.org/buildbot/#/builders/65/builds/20214

Here is the relevant piece of the build log for the reference

Step 5 (ninja check 1) failure: stage 1 checked (failure)
******************** TEST 'Clangd Unit Tests :: ./ClangdTests/89/162' FAILED ********************
Script(shard):
--
GTEST_OUTPUT=json:/home/tcwg-buildbot/worker/clang-aarch64-quick/stage1/tools/clang/tools/extra/clangd/unittests/./ClangdTests-Clangd Unit Tests-1955129-89-162.json GTEST_SHUFFLE=0 GTEST_TOTAL_SHARDS=162 GTEST_SHARD_INDEX=89 /home/tcwg-buildbot/worker/clang-aarch64-quick/stage1/tools/clang/tools/extra/clangd/unittests/./ClangdTests
--

Note: This is test shard 90 of 162.
[==========] Running 8 tests from 8 test suites.
[----------] Global test environment set-up.
[----------] 1 test from CompletionTest
[ RUN      ] CompletionTest.Snippets
Built preamble of size 707284 for file /clangd-test/foo.cpp version null in 0.03 seconds
Code complete: sema context DotMemberAccess, query scopes [] (AnyScope=false), expected type <none>
Code complete: 2 results from Sema, 0 from Index, 0 matched, 0 from identifiers, 2 returned.
[       OK ] CompletionTest.Snippets (40 ms)
[----------] 1 test from CompletionTest (40 ms total)

[----------] 1 test from CompletionStringTest
[ RUN      ] CompletionStringTest.ObjectiveCMethodOneArgument
[       OK ] CompletionStringTest.ObjectiveCMethodOneArgument (0 ms)
[----------] 1 test from CompletionStringTest (0 ms total)

[----------] 1 test from IncludeFixerTest
[ RUN      ] IncludeFixerTest.CImplicitFunctionDecl
Built preamble of size 705972 for file /clangd-test/test.c version null in 0.02 seconds
Trying to fix unresolved name "foo" in scopes: []
Built preamble of size 705836 for file /clangd-test/test.c version null in 0.02 seconds
Trying to fix unresolved name "foo" in scopes: []
[       OK ] IncludeFixerTest.CImplicitFunctionDecl (63 ms)
[----------] 1 test from IncludeFixerTest (63 ms total)

[----------] 1 test from FuzzyMatch
[ RUN      ] FuzzyMatch.Segmentation
[       OK ] FuzzyMatch.Segmentation (0 ms)
[----------] 1 test from FuzzyMatch (0 ms total)

[----------] 1 test from ParameterHints
[ RUN      ] ParameterHints.VariadicNoReferenceHintForwardingRefStdForward
Built preamble of size 707136 for file /clangd-test/TestTU.cpp version null in 0.02 seconds
[       OK ] ParameterHints.VariadicNoReferenceHintForwardingRefStdForward (33 ms)
[----------] 1 test from ParameterHints (33 ms total)

[----------] 1 test from CrossFileRenameTests
[ RUN      ] CrossFileRenameTests.WithUpToDateIndex
ASTWorker building file /clangd-test/foo.h version null with command 
[/clangd-test]
clang -xobjective-c++ /clangd-test/foo.h
Driver produced command: cc1 -cc1 -triple aarch64-unknown-linux-gnu -fsyntax-only -disable-free -clear-ast-before-backend -main-file-name foo.h -mrelocation-model pic -pic-level 2 -pic-is-pie -mframe-pointer=non-leaf -fmath-errno -ffp-contract=on -fno-rounding-math -mconstructor-aliases -funwind-tables=2 -enable-tlsdesc -target-cpu generic -target-feature +v8a -target-feature +fp-armv8 -target-feature +neon -target-abi aapcs -debugger-tuning=gdb -fdebug-compilation-dir=/clangd-test -fcoverage-compilation-dir=/clangd-test -resource-dir lib/clang/22 -internal-isystem lib/clang/22/include -internal-isystem /usr/local/include -internal-externc-isystem /include -internal-externc-isystem /usr/include -fdeprecated-macro -ferror-limit 19 -fno-signed-char -fgnuc-version=4.2.1 -fskip-odr-check-in-gmf -fobjc-runtime=gcc -fobjc-encode-cxx-class-template-spec -fobjc-exceptions -fcxx-exceptions -fexceptions -no-round-trip-args -target-feature -fmv -faddrsig -D__GCC_HAVE_DWARF2_CFI_ASM=1 -x objective-c++ /clangd-test/foo.h
Building first preamble for /clangd-test/foo.h version null
Built preamble of size 820588 for file /clangd-test/foo.h version null in 0.03 seconds
...

llvm-ci · 2025-07-26T17:05:35Z

LLVM Buildbot has detected a new failure on builder llvm-clang-x86_64-expensive-checks-debian running on gribozavr4 while building llvm at step 6 "test-build-unified-tree-check-all".

Full details are available at: https://lab.llvm.org/buildbot/#/builders/16/builds/23303

Here is the relevant piece of the build log for the reference

Step 6 (test-build-unified-tree-check-all) failure: test (failure)
******************** TEST 'LLVM :: TableGen/RuntimeLibcallEmitter.td' FAILED ********************
Exit Code: 1

Command Output (stderr):
--
/b/1/llvm-clang-x86_64-expensive-checks-debian/build/bin/llvm-tblgen -gen-runtime-libcalls -I /b/1/llvm-clang-x86_64-expensive-checks-debian/llvm-project/llvm/test/TableGen/../../include /b/1/llvm-clang-x86_64-expensive-checks-debian/llvm-project/llvm/test/TableGen/RuntimeLibcallEmitter.td | /b/1/llvm-clang-x86_64-expensive-checks-debian/build/bin/FileCheck /b/1/llvm-clang-x86_64-expensive-checks-debian/llvm-project/llvm/test/TableGen/RuntimeLibcallEmitter.td # RUN: at line 1
+ /b/1/llvm-clang-x86_64-expensive-checks-debian/build/bin/llvm-tblgen -gen-runtime-libcalls -I /b/1/llvm-clang-x86_64-expensive-checks-debian/llvm-project/llvm/test/TableGen/../../include /b/1/llvm-clang-x86_64-expensive-checks-debian/llvm-project/llvm/test/TableGen/RuntimeLibcallEmitter.td
+ /b/1/llvm-clang-x86_64-expensive-checks-debian/build/bin/FileCheck /b/1/llvm-clang-x86_64-expensive-checks-debian/llvm-project/llvm/test/TableGen/RuntimeLibcallEmitter.td
/b/1/llvm-clang-x86_64-expensive-checks-debian/llvm-project/llvm/test/TableGen/RuntimeLibcallEmitter.td:98:16: error: CHECK-NEXT: expected string not found in input
// CHECK-NEXT: sqrtl_f80 = 7, // sqrtl
               ^
<stdin>:32:23: note: scanning from here
 calloc = 6, // calloc
                      ^
<stdin>:34:2: note: possible intended match here
 sqrtl_f80 = 8, // sqrtl
 ^

Input file: <stdin>
Check file: /b/1/llvm-clang-x86_64-expensive-checks-debian/llvm-project/llvm/test/TableGen/RuntimeLibcallEmitter.td

-dump-input=help explains the following input dump.

Input was:
<<<<<<
           .
           .
           .
          27:  ___memcpy = 1, // ___memcpy 
          28:  ___memset = 2, // ___memset 
          29:  __ashlsi3 = 3, // __ashlsi3 
          30:  __lshrdi3 = 4, // __lshrdi3 
          31:  bzero = 5, // bzero 
          32:  calloc = 6, // calloc 
next:98'0                           X error: no match found
          33:  sqrtl_f128 = 7, // sqrtl 
next:98'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~
          34:  sqrtl_f80 = 8, // sqrtl 
next:98'0     ~~~~~~~~~~~~~~~~~~~~~~~~~
next:98'1      ?                        possible intended match
          35:  NumLibcallImpls = 9 
next:98'0     ~~~~~~~~~~~~~~~~~~~~~
          36: }; 
next:98'0     ~~~
          37: } // End namespace RTLIB 
next:98'0     ~~~~~~~~~~~~~~~~~~~~~~~~~
          38: } // End namespace llvm 
next:98'0     ~~~~~~~~~~~~~~~~~~~~~~~~
          39: #endif 
next:98'0     ~~~~~~~
...

…50626) Add some run lines to existing tests with VGPR selection enabled.

This was referenced Jul 25, 2025

AMDGPU: Add a few mfma test with immediate splat src2 #150627

Merged

AMDGPU: Fix not folding splat immediate into VGPR MFMA src2 #150628

Merged

arsenm added the backend:AMDGPU label Jul 25, 2025 — with Graphite App

arsenm requested review from jrbyrnes, rampitec, shiltian and srpande July 25, 2025 14:53

arsenm marked this pull request as ready for review July 25, 2025 14:53

shiltian approved these changes Jul 25, 2025

View reviewed changes

arsenm force-pushed the users/arsenm/amdgpu/add-tests-mfma-forced-vgpr-selection branch 17 times, most recently from 41c6bc3 to 4328b06 Compare July 25, 2025 17:41

rampitec approved these changes Jul 25, 2025

View reviewed changes

arsenm force-pushed the users/arsenm/amdgpu/add-tests-mfma-forced-vgpr-selection branch 15 times, most recently from 29195b5 to dd31101 Compare July 25, 2025 20:24

arsenm force-pushed the users/arsenm/amdgpu/add-tests-mfma-forced-vgpr-selection branch 6 times, most recently from b81d061 to 86c6961 Compare July 25, 2025 23:48

AMDGPU: Add tests with forced selection of mfmas to VGPR form

5810f83

Add some run lines to existing tests with VGPR selection enabled.

arsenm force-pushed the users/arsenm/amdgpu/add-tests-mfma-forced-vgpr-selection branch from 86c6961 to 5810f83 Compare July 26, 2025 00:18

arsenm merged commit e4d5969 into main Jul 26, 2025
9 checks passed

arsenm deleted the users/arsenm/amdgpu/add-tests-mfma-forced-vgpr-selection branch July 26, 2025 00:46

mahesh-attarde pushed a commit to mahesh-attarde/llvm-project that referenced this pull request Jul 28, 2025

AMDGPU: Add tests with forced selection of mfmas to VGPR form (llvm#1…

aad17b0

…50626) Add some run lines to existing tests with VGPR selection enabled.

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

AMDGPU: Add tests with forced selection of mfmas to VGPR form #150626

AMDGPU: Add tests with forced selection of mfmas to VGPR form #150626

arsenm commented Jul 25, 2025

Uh oh!

arsenm commented Jul 25, 2025 •

edited

Loading

Uh oh!

llvmbot commented Jul 25, 2025

Uh oh!

arsenm commented Jul 25, 2025 •

edited

Loading

Uh oh!

shiltian commented Jul 25, 2025

Uh oh!

Uh oh!

llvm-ci commented Jul 26, 2025

Uh oh!

llvm-ci commented Jul 26, 2025

Uh oh!

llvm-ci commented Jul 26, 2025

Uh oh!

Uh oh!

AMDGPU: Add tests with forced selection of mfmas to VGPR form #150626

AMDGPU: Add tests with forced selection of mfmas to VGPR form #150626

Conversation

arsenm commented Jul 25, 2025

Uh oh!

arsenm commented Jul 25, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

llvmbot commented Jul 25, 2025

Uh oh!

arsenm commented Jul 25, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Merge activity

Uh oh!

shiltian commented Jul 25, 2025

Uh oh!

Uh oh!

llvm-ci commented Jul 26, 2025

Uh oh!

llvm-ci commented Jul 26, 2025

Uh oh!

llvm-ci commented Jul 26, 2025

Uh oh!

Uh oh!

arsenm commented Jul 25, 2025 •

edited

Loading

arsenm commented Jul 25, 2025 •

edited

Loading