llvm.org GIT mirror llvm / 6e6c5e7
[AMDGPU] Select AGPR in PHI operand legalization If a PHI defines AGPR legalize its operands to AGPR. At the moment we can get an AGPR PHI with VGPR operands. I am not aware of any problems as it seems to be handled gracefully in RA, but this is not right anyway. It also slightly decreases VGPR pressure in some cases because we do not have to a copy via VGPR. Differential Revision: https://reviews.llvm.org/D69206 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@375446 91177308-0d34-0410-b5e6-96231b3b80d8 Stanislav Mekhanoshin 1 year, 1 month ago
2 changed file(s) with 56 addition(s) and 1 deletion(s). Raw diff Collapse all Expand all
45754575 VRC = RI.hasAGPRs(getOpRegClass(MI, 0))
45764576 ? RI.getEquivalentAGPRClass(SRC)
45774577 : RI.getEquivalentVGPRClass(SRC);
4578 } else {
4579 VRC = RI.hasAGPRs(getOpRegClass(MI, 0))
4580 ? RI.getEquivalentAGPRClass(VRC)
4581 : RI.getEquivalentVGPRClass(VRC);
45784582 }
45794583 RC = VRC;
45804584 } else {
0 ; RUN: llc -march=amdgcn -mcpu=gfx908 -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
11
22 ; GCN-LABEL: {{^}}test_mfma_loop_zeroinit:
3 ; GCN-COUNT32: v_accvgpr_write_b32
3
4 ; Check that we do not use 32 temp vgprs, but rotate 3 vgprs only.
5 ; 3 vgprs are needed to avoid wait states between writes.
6
7 ; FIXME: We should not be using and temporary registers at all.
8 ; At the moment we initialize an sgpr, then copy it via vgprs.
9
10 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2:v[0-9]+]]
11 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3:v[0-9]+]]
12
13 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1:v[0-9]+]]
14 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]]
15 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]]
16
17 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]]
18 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]]
19 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]]
20
21 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]]
22 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]]
23 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]]
24
25 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]]
26 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]]
27 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]]
28
29 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]]
30 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]]
31 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]]
32
33 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]]
34 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]]
35 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]]
36
37 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]]
38 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]]
39 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]]
40
41 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]]
42 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]]
43 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]]
44
45 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1]]
46 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]]
47 ; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]]
48
49 ; Check that we do not copy agprs to vgprs and back inside the loop.
50
451 ; GCN: [[LOOP:BB[0-9_]+]]:
552 ; GCN-NOT: v_accvgpr
653 ; GCN: v_mfma_f32_32x32x1f32
754 ; GCN-NOT: v_accvgpr
855 ; GCN: s_cbranch_scc1 [[LOOP]]
56
57 ; Final result should be read only once after the loop.
58
959 ; GCN-COUNT32: v_accvgpr_read_b32
60
1061 define amdgpu_kernel void @test_mfma_loop_zeroinit(<32 x float> addrspace(1)* %arg) {
1162 entry:
1263 br label %for.cond.preheader