llvm.org GIT mirror llvm / d0edb1f
AMDGPU: Add s_dcache_* instructions git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@248533 91177308-0d34-0410-b5e6-96231b3b80d8 Matt Arsenault 5 years ago
12 changed file(s) with 220 addition(s) and 16 deletion(s). Raw diff Collapse all Expand all
9999 GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1">,
100100 Intrinsic<[], [], []>;
101101
102 def int_amdgcn_s_dcache_inv :
103 GCCBuiltin<"__builtin_amdgcn_s_dcache_inv">,
104 Intrinsic<[], [], []>;
105
106 // CI+
107 def int_amdgcn_s_dcache_inv_vol :
108 GCCBuiltin<"__builtin_amdgcn_s_dcache_inv_vol">,
109 Intrinsic<[], [], []>;
110
111 // VI
112 def int_amdgcn_s_dcache_wb :
113 GCCBuiltin<"__builtin_amdgcn_s_dcache_wb">,
114 Intrinsic<[], [], []>;
115
116 // VI
117 def int_amdgcn_s_dcache_wb_vol :
118 GCCBuiltin<"__builtin_amdgcn_s_dcache_wb_vol">,
119 Intrinsic<[], [], []>;
120
102121 }
88 // Instruction definitions for CI and newer.
99 //===----------------------------------------------------------------------===//
1010 // Remaining instructions:
11 // FLAT_*
1211 // S_CBRANCH_CDBGUSER
1312 // S_CBRANCH_CDBGSYS
1413 // S_CBRANCH_CDBGSYS_OR_USER
1514 // S_CBRANCH_CDBGSYS_AND_USER
16 // S_DCACHE_INV_VOL
1715 // DS_NOP
1816 // DS_GWS_SEMA_RELEASE_ALL
1917 // DS_WRAP_RTN_B32
9795
9896 // DS_CONDXCHG32_RTN_B64
9997 // DS_CONDXCHG32_RTN_B128
98
99 //===----------------------------------------------------------------------===//
100 // SMRD Instructions
101 //===----------------------------------------------------------------------===//
102
103 defm S_DCACHE_INV_VOL : SMRD_Inval ,
104 "s_dcache_inv_vol", int_amdgcn_s_dcache_inv_vol>;
100105
101106 //===----------------------------------------------------------------------===//
102107 // MUBUF Instructions
139139 Counters SIInsertWaits::getHwCounts(MachineInstr &MI) {
140140
141141 uint64_t TSFlags = TII->get(MI.getOpcode()).TSFlags;
142 Counters Result;
142 Counters Result = { { 0, 0, 0 } };
143143
144144 Result.Named.VM = !!(TSFlags & SIInstrFlags::VM_CNT);
145145
152152
153153 if (TII->isSMRD(MI.getOpcode())) {
154154
155 MachineOperand &Op = MI.getOperand(0);
156 assert(Op.isReg() && "First LGKM operand must be a register!");
157
158 unsigned Reg = Op.getReg();
159 unsigned Size = TRI->getMinimalPhysRegClass(Reg)->getSize();
160 Result.Named.LGKM = Size > 4 ? 2 : 1;
161
155 if (MI.getNumOperands() != 0) {
156 MachineOperand &Op = MI.getOperand(0);
157 assert(Op.isReg() && "First LGKM operand must be a register!");
158
159 unsigned Reg = Op.getReg();
160
161 // XXX - What if this is a write into a super register?
162 unsigned Size = TRI->getMinimalPhysRegClass(Reg)->getSize();
163 Result.Named.LGKM = Size > 4 ? 2 : 1;
164 } else {
165 // s_dcache_inv etc. do not have a a destination register. Assume we
166 // want a wait on these.
167 // XXX - What is the right value?
168 Result.Named.LGKM = 1;
169 }
162170 } else {
163171 // DS
164172 Result.Named.LGKM = 1;
7272 }
7373
7474 // Specify an SMRD opcode for SI and SMEM opcode for VI
75 class smrd si, bits<5> vi = si> {
76 field bits<5> SI = si;
77 field bits<8> VI = { 0, 0, 0, vi };
75
76 // FIXME: This should really be bits<5> si, Tablegen crashes if
77 // parameter default value is other parameter with different bit size
78 class smrd si, bits<8> vi = si> {
79 field bits<5> SI = si{4-0};
80 field bits<8> VI = vi;
7881 }
7982
8083 // Execpt for the NONE field, this must be kept in sync with the SISubtarget enum
898901 }
899902
900903 class SMRD_Real_vi op, string opName, bit imm, dag outs, dag ins,
901 string asm> :
902 SMRD ,
904 string asm, list pattern = []> :
905 SMRD ,
903906 SMEMe_vi ,
904907 SIMCInstr {
905908 let AssemblerPredicates = [isVI];
917920 let glc = 0 in {
918921 def _vi : SMRD_Real_vi ;
919922 }
923 }
924
925 multiclass SMRD_Inval
926 SDPatternOperator node> {
927 let hasSideEffects = 1, mayStore = 1 in {
928 def "" : SMRD_Pseudo ;
929
930 let sbase = 0, offset = 0 in {
931 let sdst = 0 in {
932 def _si : SMRD_Real_si ;
933 }
934
935 let glc = 0, sdata = 0 in {
936 def _vi : SMRD_Real_vi ;
937 }
938 }
939 }
940 }
941
942 class SMEM_Inval op, string opName, SDPatternOperator node> :
943 SMRD_Real_vi {
944 let hasSideEffects = 1;
945 let mayStore = 1;
946 let sbase = 0;
947 let sdata = 0;
948 let glc = 0;
949 let offset = 0;
920950 }
921951
922952 multiclass SMRD_Helper
9292 } // mayLoad = 1
9393
9494 //def S_MEMTIME : SMRD_ <0x0000001e, "s_memtime", []>;
95 //def S_DCACHE_INV : SMRD_ <0x0000001f, "s_dcache_inv", []>;
95
96 defm S_DCACHE_INV : SMRD_Inval , "s_dcache_inv",
97 int_amdgcn_s_dcache_inv>;
9698
9799 //===----------------------------------------------------------------------===//
98100 // SOP1 Instructions
8888 def : SI2_VI3Alias <"v_cvt_pknorm_u16_f32", V_CVT_PKNORM_U16_F32_e64_vi>;
8989 def : SI2_VI3Alias <"v_cvt_pkrtz_f16_f32", V_CVT_PKRTZ_F16_F32_e64_vi>;
9090
91 //===----------------------------------------------------------------------===//
92 // SMEM Instructions
93 //===----------------------------------------------------------------------===//
94
95 def S_DCACHE_WB : SMEM_Inval <0x21,
96 "s_dcache_wb", int_amdgcn_s_dcache_wb>;
97
98 def S_DCACHE_WB_VOL : SMEM_Inval <0x23,
99 "s_dcache_wb_vol", int_amdgcn_s_dcache_wb_vol>;
100
91101 } // End SIAssemblerPredicate = DisableInst, SubtargetPredicate = isVI
92102
93103 //===----------------------------------------------------------------------===//
0 ; RUN: llc -march=amdgcn -mcpu=tahiti -show-mc-encoding < %s | FileCheck -check-prefix=GCN -check-prefix=SI %s
1 ; RUN: llc -march=amdgcn -mcpu=fiji -show-mc-encoding < %s | FileCheck -check-prefix=GCN -check-prefix=VI %s
2
3 declare void @llvm.amdgcn.s.dcache.inv() #0
4
5 ; GCN-LABEL: {{^}}test_s_dcache_inv:
6 ; GCN-NEXT: ; BB#0:
7 ; SI-NEXT: s_dcache_inv ; encoding: [0x00,0x00,0xc0,0xc7]
8 ; VI-NEXT: s_dcache_inv ; encoding: [0x00,0x00,0x80,0xc0,0x00,0x00,0x00,0x00]
9 ; GCN-NEXT: s_endpgm
10 define void @test_s_dcache_inv() #0 {
11 call void @llvm.amdgcn.s.dcache.inv()
12 ret void
13 }
14
15 ; GCN-LABEL: {{^}}test_s_dcache_inv_insert_wait:
16 ; GCN-NEXT: ; BB#0:
17 ; GCN-NEXT: s_dcache_inv
18 ; GCN-NEXT: s_waitcnt lgkmcnt(0) ; encoding
19 define void @test_s_dcache_inv_insert_wait() #0 {
20 call void @llvm.amdgcn.s.dcache.inv()
21 br label %end
22
23 end:
24 store volatile i32 3, i32 addrspace(1)* undef
25 ret void
26 }
27
28 attributes #0 = { nounwind }
0 ; RUN: llc -march=amdgcn -mcpu=bonaire -show-mc-encoding < %s | FileCheck -check-prefix=GCN -check-prefix=CI %s
1 ; RUN: llc -march=amdgcn -mcpu=tonga -show-mc-encoding < %s | FileCheck -check-prefix=GCN -check-prefix=VI %s
2
3 declare void @llvm.amdgcn.s.dcache.inv.vol() #0
4
5 ; GCN-LABEL: {{^}}test_s_dcache_inv_vol:
6 ; GCN-NEXT: ; BB#0:
7 ; CI-NEXT: s_dcache_inv_vol ; encoding: [0x00,0x00,0x40,0xc7]
8 ; VI-NEXT: s_dcache_inv_vol ; encoding: [0x00,0x00,0x88,0xc0,0x00,0x00,0x00,0x00]
9 ; GCN-NEXT: s_endpgm
10 define void @test_s_dcache_inv_vol() #0 {
11 call void @llvm.amdgcn.s.dcache.inv.vol()
12 ret void
13 }
14
15 ; GCN-LABEL: {{^}}test_s_dcache_inv_vol_insert_wait:
16 ; GCN-NEXT: ; BB#0:
17 ; GCN-NEXT: s_dcache_inv_vol
18 ; GCN-NEXT: s_waitcnt lgkmcnt(0) ; encoding
19 define void @test_s_dcache_inv_vol_insert_wait() #0 {
20 call void @llvm.amdgcn.s.dcache.inv.vol()
21 br label %end
22
23 end:
24 store volatile i32 3, i32 addrspace(1)* undef
25 ret void
26 }
27
28 attributes #0 = { nounwind }
0 ; RUN: llc -march=amdgcn -mcpu=fiji -show-mc-encoding < %s | FileCheck -check-prefix=VI %s
1
2 declare void @llvm.amdgcn.s.dcache.wb() #0
3
4 ; VI-LABEL: {{^}}test_s_dcache_wb:
5 ; VI-NEXT: ; BB#0:
6 ; VI-NEXT: s_dcache_wb ; encoding: [0x00,0x00,0x84,0xc0,0x00,0x00,0x00,0x00]
7 ; VI-NEXT: s_endpgm
8 define void @test_s_dcache_wb() #0 {
9 call void @llvm.amdgcn.s.dcache.wb()
10 ret void
11 }
12
13 ; VI-LABEL: {{^}}test_s_dcache_wb_insert_wait:
14 ; VI-NEXT: ; BB#0:
15 ; VI-NEXT: s_dcache_wb
16 ; VI-NEXT: s_waitcnt lgkmcnt(0) ; encoding
17 define void @test_s_dcache_wb_insert_wait() #0 {
18 call void @llvm.amdgcn.s.dcache.wb()
19 br label %end
20
21 end:
22 store volatile i32 3, i32 addrspace(1)* undef
23 ret void
24 }
25
26 attributes #0 = { nounwind }
0 ; RUN: llc -march=amdgcn -mcpu=fiji -show-mc-encoding < %s | FileCheck -check-prefix=VI %s
1
2 declare void @llvm.amdgcn.s.dcache.wb.vol() #0
3
4 ; VI-LABEL: {{^}}test_s_dcache_wb_vol:
5 ; VI-NEXT: ; BB#0:
6 ; VI-NEXT: s_dcache_wb_vol ; encoding: [0x00,0x00,0x8c,0xc0,0x00,0x00,0x00,0x00]
7 ; VI-NEXT: s_endpgm
8 define void @test_s_dcache_wb_vol() #0 {
9 call void @llvm.amdgcn.s.dcache.wb.vol()
10 ret void
11 }
12
13 ; VI-LABEL: {{^}}test_s_dcache_wb_vol_insert_wait:
14 ; VI-NEXT: ; BB#0:
15 ; VI-NEXT: s_dcache_wb_vol
16 ; VI-NEXT: s_waitcnt lgkmcnt(0) ; encoding
17 define void @test_s_dcache_wb_vol_insert_wait() #0 {
18 call void @llvm.amdgcn.s.dcache.wb.vol()
19 br label %end
20
21 end:
22 store volatile i32 3, i32 addrspace(1)* undef
23 ret void
24 }
25
26 attributes #0 = { nounwind }
0 // RUN: llvm-mc -arch=amdgcn -mcpu=tonga -show-encoding %s | FileCheck -check-prefix=GCN -check-prefix=VI %s
1 // RUN: not llvm-mc -arch=amdgcn -mcpu=tahiti %s 2>&1 | FileCheck -check-prefix=NOSI %s
2 // RUN: not llvm-mc -arch=amdgcn -mcpu=bonaire %s 2>&1 | FileCheck -check-prefix=NOSI %s
3
4 s_dcache_wb
5 ; VI: s_dcache_wb ; encoding: [0x00,0x00,0x84,0xc0,0x00,0x00,0x00,0x00]
6 ; NOSI: error: instruction not supported on this GPU
7
8 s_dcache_wb_vol
9 ; VI: s_dcache_wb_vol ; encoding: [0x00,0x00,0x8c,0xc0,0x00,0x00,0x00,0x00]
10 ; NOSI: error: instruction not supported on this GPU
5050
5151 s_load_dwordx16 s[16:31], s[2:3], s4
5252 // GCN: s_load_dwordx16 s[16:31], s[2:3], s4 ; encoding: [0x04,0x02,0x08,0xc1]
53
54 s_dcache_inv
55 // GCN: s_dcache_inv ; encoding: [0x00,0x00,0xc0,0xc7]
56
57 s_dcache_inv_vol
58 // CI: s_dcache_inv_vol ; encoding: [0x00,0x00,0x40,0xc7]
59 // NOSI: error: instruction not supported on this GPU