Magic-Mirror/llvm-capstone
1
0
Fork 0
mirror of https://github.com/capstone-engine/llvm-capstone.git synced 2025-02-03 16:03:21 +00:00
Code Issues Actions Packages Projects Releases Wiki Activity

[AMDGPU] Select AGPR in PHI operand legalization

Browse Source 
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

llvm-svn: 375446
This commit is contained in:
Stanislav Mekhanoshin 2019-10-21 19:25:27 +00:00
parent 7c15c4fb17
commit 33092194f2
2 changed files with 56 additions and 1 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
View File

@ -4576,6 +4576,10 @@ void SIInstrInfo::legalizeOperands(MachineInstr &MI,
VRC = RI.hasAGPRs(getOpRegClass(MI, 0))
? RI.getEquivalentAGPRClass(SRC)
: RI.getEquivalentVGPRClass(SRC);
} else {
VRC = RI.hasAGPRs(getOpRegClass(MI, 0))
? RI.getEquivalentAGPRClass(VRC)
: RI.getEquivalentVGPRClass(VRC);
}
RC = VRC;
} else {

53
llvm/test/CodeGen/AMDGPU/mfma-loop.ll
View File

@ -1,13 +1,64 @@
; RUN: llc -march=amdgcn -mcpu=gfx908 -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
; GCN-LABEL: {{^}}test_mfma_loop_zeroinit:
; GCN-COUNT32: v_accvgpr_write_b32
; Check that we do not use 32 temp vgprs, but rotate 3 vgprs only.
; 3 vgprs are needed to avoid wait states between writes.
; FIXME: We should not be using and temporary registers at all.
; At the moment we initialize an sgpr, then copy it via vgprs.
; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2:v[0-9]+]]
; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3:v[0-9]+]]
; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP1:v[0-9]+]]
; GCN: v_accvgpr_write_b32 a{{[0-9]+}}, [[TMP2]]
; 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_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_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_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_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_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_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_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_accvgpr_write_b32 a{{[0-9]+}}, [[TMP3]]
; Check that we do not copy agprs to vgprs and back inside the loop.
; GCN: [[LOOP:BB[0-9_]+]]:
; GCN-NOT: v_accvgpr
; GCN: v_mfma_f32_32x32x1f32
; GCN-NOT: v_accvgpr
; GCN: s_cbranch_scc1 [[LOOP]]
; Final result should be read only once after the loop.
; GCN-COUNT32: v_accvgpr_read_b32
define amdgpu_kernel void @test_mfma_loop_zeroinit(<32 x float> addrspace(1)* %arg) {
entry:
br label %for.cond.preheader