Skip to content

Commit 17c8c1c

Browse files
bcahoonbcahoon
authored
[AMDGPU] Do not fold into v_accvpr_mov/write/read (#120475)
In SIFoldOperands, leave copies for moving between agpr and vgpr registers. The register coalescer is able to handle the copies more efficiently than v_accvgpr_mov, v_accvgpr_write, and v_accvgpr_read. Otherwise, the compiler generates unneccesary instructions such as v_accvgpr_mov a0, a0.
1 parent 8178d3c commit 17c8c1c

File tree

4 files changed

+537
-13
lines changed

4 files changed

+537
-13
lines changed

‎llvm/lib/Target/AMDGPU/SIFoldOperands.cpp

Lines changed: 0 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -1096,21 +1096,8 @@ void SIFoldOperandsImpl::foldOperand(
10961096
B.addImm(Defs[I].second);
10971097
}
10981098
LLVM_DEBUG(dbgs() << "Folded " << *UseMI);
1099-
return;
11001099
}
11011100

1102-
if (Size != 4)
1103-
return;
1104-
1105-
Register Reg0 = UseMI->getOperand(0).getReg();
1106-
Register Reg1 = UseMI->getOperand(1).getReg();
1107-
if (TRI->isAGPR(*MRI, Reg0) && TRI->isVGPR(*MRI, Reg1))
1108-
UseMI->setDesc(TII->get(AMDGPU::V_ACCVGPR_WRITE_B32_e64));
1109-
else if (TRI->isVGPR(*MRI, Reg0) && TRI->isAGPR(*MRI, Reg1))
1110-
UseMI->setDesc(TII->get(AMDGPU::V_ACCVGPR_READ_B32_e64));
1111-
else if (ST->hasGFX90AInsts() && TRI->isAGPR(*MRI, Reg0) &&
1112-
TRI->isAGPR(*MRI, Reg1))
1113-
UseMI->setDesc(TII->get(AMDGPU::V_ACCVGPR_MOV_B32));
11141101
return;
11151102
}
11161103

‎llvm/test/CodeGen/AMDGPU/no-fold-accvgpr-mov.ll

Lines changed: 120 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,120 @@
1+
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
2+
; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx942 %s -o - | FileCheck %s --check-prefixes=GFX942
3+
; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=gfx908 %s -o - | FileCheck %s --check-prefixes=GFX908
4+
5+
define amdgpu_kernel void @matmul_kernel(i32 %a0, i32 %a1) {
6+
; GFX942-LABEL: matmul_kernel:
7+
; GFX942: ; %bb.0: ; %entry
8+
; GFX942-NEXT: s_load_dwordx2 s[0:1], s[4:5], 0x0
9+
; GFX942-NEXT: v_mov_b32_e32 v1, 0
10+
; GFX942-NEXT: s_mov_b32 s2, 0
11+
; GFX942-NEXT: v_accvgpr_write_b32 a0, v1
12+
; GFX942-NEXT: s_mov_b32 s3, 0
13+
; GFX942-NEXT: s_waitcnt lgkmcnt(0)
14+
; GFX942-NEXT: s_cmp_lg_u32 s0, 0
15+
; GFX942-NEXT: s_cselect_b64 s[0:1], -1, 0
16+
; GFX942-NEXT: v_cndmask_b32_e64 v0, 0, 1, s[0:1]
17+
; GFX942-NEXT: v_cmp_ne_u32_e64 s[0:1], 1, v0
18+
; GFX942-NEXT: s_branch .LBB0_2
19+
; GFX942-NEXT: .LBB0_1: ; %bb2
20+
; GFX942-NEXT: ; in Loop: Header=BB0_2 Depth=1
21+
; GFX942-NEXT: s_or_b32 s4, s3, 1
22+
; GFX942-NEXT: s_ashr_i32 s5, s3, 31
23+
; GFX942-NEXT: s_mov_b32 s3, s2
24+
; GFX942-NEXT: v_mov_b64_e32 v[4:5], s[2:3]
25+
; GFX942-NEXT: v_accvgpr_read_b32 v0, a0
26+
; GFX942-NEXT: v_mov_b32_e32 v2, v1
27+
; GFX942-NEXT: v_mov_b32_e32 v3, v1
28+
; GFX942-NEXT: v_accvgpr_write_b32 a0, v0
29+
; GFX942-NEXT: v_accvgpr_write_b32 a1, v1
30+
; GFX942-NEXT: v_accvgpr_write_b32 a2, v2
31+
; GFX942-NEXT: v_accvgpr_write_b32 a3, v3
32+
; GFX942-NEXT: s_and_b32 s3, s5, s4
33+
; GFX942-NEXT: s_nop 0
34+
; GFX942-NEXT: v_mfma_f32_16x16x16_f16 a[0:3], v[4:5], v[4:5], a[0:3]
35+
; GFX942-NEXT: s_cbranch_execz .LBB0_4
36+
; GFX942-NEXT: .LBB0_2: ; %bb
37+
; GFX942-NEXT: ; =>This Inner Loop Header: Depth=1
38+
; GFX942-NEXT: s_and_b64 vcc, exec, s[0:1]
39+
; GFX942-NEXT: s_cbranch_vccz .LBB0_1
40+
; GFX942-NEXT: ; %bb.3:
41+
; GFX942-NEXT: ; implicit-def: $sgpr3
42+
; GFX942-NEXT: .LBB0_4: ; %common.ret
43+
; GFX942-NEXT: s_endpgm
44+
;
45+
; GFX908-LABEL: matmul_kernel:
46+
; GFX908: ; %bb.0: ; %entry
47+
; GFX908-NEXT: s_load_dwordx2 s[0:1], s[8:9], 0x0
48+
; GFX908-NEXT: v_mov_b32_e32 v1, 0
49+
; GFX908-NEXT: s_mov_b32 s2, 0
50+
; GFX908-NEXT: s_mov_b32 s3, 0
51+
; GFX908-NEXT: v_accvgpr_write_b32 a0, v1
52+
; GFX908-NEXT: s_waitcnt lgkmcnt(0)
53+
; GFX908-NEXT: s_cmp_lg_u32 s0, 0
54+
; GFX908-NEXT: s_cselect_b64 s[0:1], -1, 0
55+
; GFX908-NEXT: v_cndmask_b32_e64 v0, 0, 1, s[0:1]
56+
; GFX908-NEXT: v_cmp_ne_u32_e64 s[0:1], 1, v0
57+
; GFX908-NEXT: s_branch .LBB0_2
58+
; GFX908-NEXT: .LBB0_1: ; %bb2
59+
; GFX908-NEXT: ; in Loop: Header=BB0_2 Depth=1
60+
; GFX908-NEXT: s_or_b32 s4, s3, 1
61+
; GFX908-NEXT: s_ashr_i32 s5, s3, 31
62+
; GFX908-NEXT: s_mov_b32 s3, s2
63+
; GFX908-NEXT: s_nop 3
64+
; GFX908-NEXT: v_accvgpr_read_b32 v0, a0
65+
; GFX908-NEXT: v_mov_b32_e32 v5, s3
66+
; GFX908-NEXT: v_mov_b32_e32 v4, s2
67+
; GFX908-NEXT: v_mov_b32_e32 v2, v1
68+
; GFX908-NEXT: v_mov_b32_e32 v3, v1
69+
; GFX908-NEXT: v_accvgpr_write_b32 a0, v0
70+
; GFX908-NEXT: v_accvgpr_write_b32 a1, v1
71+
; GFX908-NEXT: v_accvgpr_write_b32 a2, v2
72+
; GFX908-NEXT: v_accvgpr_write_b32 a3, v3
73+
; GFX908-NEXT: s_and_b32 s3, s5, s4
74+
; GFX908-NEXT: v_mfma_f32_16x16x16f16 a[0:3], v[4:5], v[4:5], a[0:3]
75+
; GFX908-NEXT: s_cbranch_execz .LBB0_4
76+
; GFX908-NEXT: .LBB0_2: ; %bb
77+
; GFX908-NEXT: ; =>This Inner Loop Header: Depth=1
78+
; GFX908-NEXT: s_and_b64 vcc, exec, s[0:1]
79+
; GFX908-NEXT: s_cbranch_vccz .LBB0_1
80+
; GFX908-NEXT: ; %bb.3:
81+
; GFX908-NEXT: ; implicit-def: $sgpr3
82+
; GFX908-NEXT: .LBB0_4: ; %common.ret
83+
; GFX908-NEXT: s_endpgm
84+
entry:
85+
br label %bb
86+
87+
bb:
88+
%i = phi { float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float } [ %i10, %bb2 ], [ zeroinitializer, %entry ]
89+
%i1 = phi i32 [ %i5, %bb2 ], [ 0, %entry ]
90+
%c0 = icmp ne i32 %a0, 0
91+
br i1 %c0, label %bb2, label %bb11
92+
93+
bb2:
94+
%i3 = or i32 %i1, 1
95+
%i4 = icmp slt i32 %i1, 0
96+
%i5 = select i1 %i4, i32 %i3, i32 0
97+
%i6 = extractvalue { float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float } %i, 123
98+
%i7 = insertelement <4 x float> zeroinitializer, float %i6, i32 0
99+
%i8 = call <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half> zeroinitializer, <4 x half> zeroinitializer, <4 x float> %i7, i32 0, i32 0, i32 0)
100+
%i9 = extractelement <4 x float> %i8, i32 0
101+
%i10 = insertvalue { float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float } zeroinitializer, float %i9, 123
102+
br label %bb
103+
104+
bb11:
105+
%c1 = icmp ne i32 %a1, 0
106+
br i1 %c1, label %bb12, label %common.ret
107+
108+
common.ret:
109+
ret void
110+
111+
bb12:
112+
%i13 = extractvalue { float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float, float } %i, 0
113+
%i14 = insertelement <4 x float> zeroinitializer, float %i13, i32 0
114+
%i15 = insertelement <4 x float> %i14, float 0.000000e+00, i32 0
115+
%i16 = insertelement <4 x float> %i15, float 0.000000e+00, i32 0
116+
br label %common.ret
117+
}
118+
119+
; Function Attrs: convergent nocallback nofree nosync nounwind willreturn memory(none)
120+
declare <4 x float> @llvm.amdgcn.mfma.f32.16x16x16f16(<4 x half>, <4 x half>, <4 x float>, i32 immarg, i32 immarg, i32 immarg)

0 commit comments

Comments
 (0)