; GCN-LABEL: {{^}}test_mfma_loop_zeroinit:
-; GCN-COUNT32: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
+; GCN-COUNT-32: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
; Check that we do not copy agprs to vgprs and back inside the loop.
; Final result should be read only once after the loop.
-; GCN-COUNT32: v_accvgpr_read_b32
+; GCN-COUNT-32: v_accvgpr_read_b32
define amdgpu_kernel void @test_mfma_loop_zeroinit(<32 x float> addrspace(1)* %arg) {
entry:
; Check that we do not use 32 temp sgprs as well.
; GCN: v_mov_b32_e32 [[TMP:v[0-9]+]], 0x42f60000
-; GCN-COUNT32: v_accvgpr_write_b32 a0, [[TMP]]
+; GCN-COUNT-32: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]]
; GCN: [[LOOP:BB[0-9_]+]]:
; GCN-NOT: v_accvgpr
; GCN-NOT: v_accvgpr
; GCN: s_cbranch_scc1 [[LOOP]]
-; GCN-COUNT32: v_accvgpr_read_b32
+; GCN-COUNT-32: v_accvgpr_read_b32
define amdgpu_kernel void @test_mfma_loop_unfoldable_splat(<32 x float> addrspace(1)* %arg) {
entry:
; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, 1.0{{$}}
-; GCN-COUNT30: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
+; GCN-COUNT-30: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
; GCN: [[LOOP:BB[0-9_]+]]:
; GCN-NOT: v_accvgpr
; GCN-NOT: v_accvgpr
; GCN: s_cbranch_scc1 [[LOOP]]
-; GCN-COUNT32: v_accvgpr_read_b32
+; GCN-COUNT-32: v_accvgpr_read_b32
define amdgpu_kernel void @test_mfma_loop_non_splat(<32 x float> addrspace(1)* %arg) {
entry:
; GCN: v_mov_b32_e32 [[TMP1:v[0-9]+]], 0x42f60000
; GCN: v_mov_b32_e32 [[TMP2:v[0-9]+]], 0x42f80000
; GCN: v_mov_b32_e32 [[TMP3:v[0-9]+]], 0x42fe0000
-; GCN-COUNT29: v_mov_b32_e32 v1, 0x4{{[0-9a-f]+}}
-; GCN-COUNT10: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]]
-; GCN-COUNT11: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]]
-; GCN-COUNT11: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]]
+; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]]
+; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]]
+; GCN: v_mov_b32_e32 [[TMP1]], 0x4{{[0-9a-f]+}}
+; GCN: v_mov_b32_e32 [[TMP2]], 0x4{{[0-9a-f]+}}
+; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]]
+; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]]
+; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]]
+; GCN: v_mov_b32_e32 [[TMP1]], 0x4{{[0-9a-f]+}}
+; GCN: v_mov_b32_e32 [[TMP2]], 0x4{{[0-9a-f]+}}
+; GCN: v_mov_b32_e32 [[TMP3]], 0x4{{[0-9a-f]+}}
+; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]]
+; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]]
+; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]]
+; GCN: v_mov_b32_e32 [[TMP1]], 0x4{{[0-9a-f]+}}
+; GCN: v_mov_b32_e32 [[TMP2]], 0x4{{[0-9a-f]+}}
+; GCN: v_mov_b32_e32 [[TMP3]], 0x4{{[0-9a-f]+}}
+; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]]
+; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]]
+; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]]
+; GCN: v_mov_b32_e32 [[TMP1]], 0x4{{[0-9a-f]+}}
+; GCN: v_mov_b32_e32 [[TMP2]], 0x4{{[0-9a-f]+}}
+; GCN: v_mov_b32_e32 [[TMP3]], 0x4{{[0-9a-f]+}}
+; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]]
+; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]]
+; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]]
+; GCN: v_mov_b32_e32 [[TMP1]], 0x4{{[0-9a-f]+}}
+; GCN: v_mov_b32_e32 [[TMP2]], 0x4{{[0-9a-f]+}}
+; GCN: v_mov_b32_e32 [[TMP3]], 0x4{{[0-9a-f]+}}
+; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]]
+; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]]
+; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]]
+; GCN: v_mov_b32_e32 [[TMP1]], 0x4{{[0-9a-f]+}}
+; GCN: v_mov_b32_e32 [[TMP2]], 0x4{{[0-9a-f]+}}
+; GCN: v_mov_b32_e32 [[TMP3]], 0x4{{[0-9a-f]+}}
+; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]]
+; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]]
+; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]]
+; GCN: v_mov_b32_e32 [[TMP1]], 0x4{{[0-9a-f]+}}
+; GCN: v_mov_b32_e32 [[TMP2]], 0x4{{[0-9a-f]+}}
+; GCN: v_mov_b32_e32 [[TMP3]], 0x4{{[0-9a-f]+}}
+; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]]
+; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]]
+; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]]
+; GCN: v_mov_b32_e32 [[TMP1]], 0x4{{[0-9a-f]+}}
+; GCN: v_mov_b32_e32 [[TMP2]], 0x4{{[0-9a-f]+}}
+; GCN: v_mov_b32_e32 [[TMP3]], 0x4{{[0-9a-f]+}}
+; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]]
+; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]]
+; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]]
+; GCN: v_mov_b32_e32 [[TMP1]], 0x4{{[0-9a-f]+}}
+; GCN: v_mov_b32_e32 [[TMP2]], 0x4{{[0-9a-f]+}}
+; GCN: v_mov_b32_e32 [[TMP3]], 0x4{{[0-9a-f]+}}
+; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]]
+; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]]
+; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]]
+; GCN: v_mov_b32_e32 [[TMP1]], 0x4{{[0-9a-f]+}}
+; GCN: v_mov_b32_e32 [[TMP2]], 0x4{{[0-9a-f]+}}
+; GCN: v_mov_b32_e32 [[TMP3]], 0x4{{[0-9a-f]+}}
+; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]]
+; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]]
+; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]]
; GCN: [[LOOP:BB[0-9_]+]]:
; GCN-NOT: v_accvgpr
; GCN-NOT: v_accvgpr
; GCN: s_cbranch_scc1 [[LOOP]]
-; GCN-COUNT32: v_accvgpr_read_b32
+; GCN-COUNT-32: v_accvgpr_read_b32
define amdgpu_kernel void @test_mfma_loop_unfoldable_seq(<32 x float> addrspace(1)* %arg) {
entry:
; GCN-LABEL: {{^}}test_mfma_loop_vgpr_init:
-; GCN-COUNT32: v_accvgpr_write_b32 a{{[0-9]+}}, v0{{$}}
+; GCN-COUNT-32: v_accvgpr_write_b32 a{{[0-9]+}}, v0{{$}}
; GCN: [[LOOP:BB[0-9_]+]]:
; GCN-NOT: v_accvgpr
; GCN-NOT: v_accvgpr
; GCN: s_cbranch_scc1 [[LOOP]]
-; GCN-COUNT32: v_accvgpr_read_b32
+; GCN-COUNT-32: v_accvgpr_read_b32
define amdgpu_kernel void @test_mfma_loop_vgpr_init(<32 x float> addrspace(1)* %arg) {
entry:
; GCN-LABEL: {{^}}test_mfma_loop_sgpr_init:
; GCN: v_mov_b32_e32 [[TMP:v[0-9]+]], s{{[0-9]+}}
-; GCN-COUNT32: v_accvgpr_write_b32 a0, [[TMP]]
+; GCN-COUNT-32: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]]
; GCN: [[LOOP:BB[0-9_]+]]:
; GCN-NOT: v_accvgpr
; GCN-NOT: v_accvgpr
; GCN: s_cbranch_scc1 [[LOOP]]
-; GCN-COUNT32: v_accvgpr_read_b32
+; GCN-COUNT-32: v_accvgpr_read_b32
define amdgpu_kernel void @test_mfma_loop_sgpr_init(<32 x float> addrspace(1)* %arg, float %init) {
entry:
; GCN-LABEL: {{^}}test_mfma_loop_mixed_init:
-; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, v0
-; GCN: v_mov_b32_e32 [[TMP:v[0-9]+]], s{{[0-9]+}}
-; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]]
-
-; GCN-COUNT30: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
+; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, v0
+; GCN-DAG: v_mov_b32_e32 [[TMP:v[0-9]+]], s{{[0-9]+}}
+; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]]
+; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
+; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
+; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
+; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
+; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
+; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
+; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
+; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
+; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
+; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
+; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
+; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
+; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
+; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
+; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
+; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
+; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
+; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
+; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
+; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
+; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
+; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
+; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
+; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
+; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
+; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
+; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
+; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
+; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
+; GCN-DAG: v_accvgpr_write_b32 a{{[0-9]+}}, 0{{$}}
; GCN: [[LOOP:BB[0-9_]+]]:
; GCN-NOT: v_accvgpr
; GCN-NOT: v_accvgpr
; GCN: s_cbranch_scc1 [[LOOP]]
-; GCN-COUNT32: v_accvgpr_read_b32
+; GCN-COUNT-32: v_accvgpr_read_b32
define amdgpu_kernel void @test_mfma_loop_mixed_init(<32 x float> addrspace(1)* %arg, float %x) {
entry:
; GCN-LABEL: {{^}}test_mfma_loop_mfma_forward_init:
-; GCN-COUNT32: v_accvgpr_write_b32 a{{[0-9]+}}, 0
+; GCN-COUNT-32: v_accvgpr_write_b32 a{{[0-9]+}}, 0
; GCN: v_mfma_f32_32x32x1f32
; GCN-NOT: v_accvgpr
; GCN-NOT: v_accvgpr
; GCN: s_cbranch_scc1 [[LOOP]]
-; GCN-COUNT32: v_accvgpr_read_b32
+; GCN-COUNT-32: v_accvgpr_read_b32
define amdgpu_kernel void @test_mfma_loop_mfma_forward_init(<32 x float> addrspace(1)* %arg) {
entry:
%mai.0 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> zeroinitializer, i32 0, i32 0, i32 0)
; GCN-LABEL: {{^}}test_mfma_loop_agpr_init:
-; GCN-COUNT32: v_accvgpr_write_b32 a{{[0-9]+}}, 0
+; GCN-COUNT-32: v_accvgpr_write_b32 a{{[0-9]+}}, 0
; GCN: v_mfma_f32_32x32x1f32
; Check that we are using only one tmp VGPR.
; GCN: v_accvgpr_read_b32 [[TMP:v[0-9]+]], a{{[0-9]+}}
-; GCN-COUNT32: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]]{{$}}
+; GCN-COUNT-31: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP]]{{$}}
; GCN: [[LOOP:BB[0-9_]+]]:
; GCN-NOT: v_accvgpr
; GCN-NOT: v_accvgpr
; GCN: s_cbranch_scc1 [[LOOP]]
-; GCN-COUNT32: v_accvgpr_read_b32
+; GCN-COUNT-32: v_accvgpr_read_b32
define amdgpu_kernel void @test_mfma_loop_agpr_init(<32 x float> addrspace(1)* %arg) {
entry:
%mai.0 = tail call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.0, float 2.0, <32 x float> zeroinitializer, i32 0, i32 0, i32 0)