• Home
  • Line#
  • Scopes#
  • Navigate#
  • Raw
  • Download
1; RUN: llc < %s -march=ptx32 -mattr=+ptx20 | FileCheck %s
2
3define ptx_device i32 @test_tid_x() {
4; CHECK: mov.u32 %ret0, %tid.x;
5; CHECK: ret;
6	%x = call i32 @llvm.ptx.read.tid.x()
7	ret i32 %x
8}
9
10define ptx_device i32 @test_tid_y() {
11; CHECK: mov.u32 %ret0, %tid.y;
12; CHECK: ret;
13	%x = call i32 @llvm.ptx.read.tid.y()
14	ret i32 %x
15}
16
17define ptx_device i32 @test_tid_z() {
18; CHECK: mov.u32 %ret0, %tid.z;
19; CHECK: ret;
20	%x = call i32 @llvm.ptx.read.tid.z()
21	ret i32 %x
22}
23
24define ptx_device i32 @test_tid_w() {
25; CHECK: mov.u32 %ret0, %tid.w;
26; CHECK: ret;
27	%x = call i32 @llvm.ptx.read.tid.w()
28	ret i32 %x
29}
30
31define ptx_device i32 @test_ntid_x() {
32; CHECK: mov.u32 %ret0, %ntid.x;
33; CHECK: ret;
34	%x = call i32 @llvm.ptx.read.ntid.x()
35	ret i32 %x
36}
37
38define ptx_device i32 @test_ntid_y() {
39; CHECK: mov.u32 %ret0, %ntid.y;
40; CHECK: ret;
41	%x = call i32 @llvm.ptx.read.ntid.y()
42	ret i32 %x
43}
44
45define ptx_device i32 @test_ntid_z() {
46; CHECK: mov.u32 %ret0, %ntid.z;
47; CHECK: ret;
48	%x = call i32 @llvm.ptx.read.ntid.z()
49	ret i32 %x
50}
51
52define ptx_device i32 @test_ntid_w() {
53; CHECK: mov.u32 %ret0, %ntid.w;
54; CHECK: ret;
55	%x = call i32 @llvm.ptx.read.ntid.w()
56	ret i32 %x
57}
58
59define ptx_device i32 @test_laneid() {
60; CHECK: mov.u32 %ret0, %laneid;
61; CHECK: ret;
62	%x = call i32 @llvm.ptx.read.laneid()
63	ret i32 %x
64}
65
66define ptx_device i32 @test_warpid() {
67; CHECK: mov.u32 %ret0, %warpid;
68; CHECK: ret;
69	%x = call i32 @llvm.ptx.read.warpid()
70	ret i32 %x
71}
72
73define ptx_device i32 @test_nwarpid() {
74; CHECK: mov.u32 %ret0, %nwarpid;
75; CHECK: ret;
76	%x = call i32 @llvm.ptx.read.nwarpid()
77	ret i32 %x
78}
79
80define ptx_device i32 @test_ctaid_x() {
81; CHECK: mov.u32 %ret0, %ctaid.x;
82; CHECK: ret;
83	%x = call i32 @llvm.ptx.read.ctaid.x()
84	ret i32 %x
85}
86
87define ptx_device i32 @test_ctaid_y() {
88; CHECK: mov.u32 %ret0, %ctaid.y;
89; CHECK: ret;
90	%x = call i32 @llvm.ptx.read.ctaid.y()
91	ret i32 %x
92}
93
94define ptx_device i32 @test_ctaid_z() {
95; CHECK: mov.u32 %ret0, %ctaid.z;
96; CHECK: ret;
97	%x = call i32 @llvm.ptx.read.ctaid.z()
98	ret i32 %x
99}
100
101define ptx_device i32 @test_ctaid_w() {
102; CHECK: mov.u32 %ret0, %ctaid.w;
103; CHECK: ret;
104	%x = call i32 @llvm.ptx.read.ctaid.w()
105	ret i32 %x
106}
107
108define ptx_device i32 @test_nctaid_x() {
109; CHECK: mov.u32 %ret0, %nctaid.x;
110; CHECK: ret;
111	%x = call i32 @llvm.ptx.read.nctaid.x()
112	ret i32 %x
113}
114
115define ptx_device i32 @test_nctaid_y() {
116; CHECK: mov.u32 %ret0, %nctaid.y;
117; CHECK: ret;
118	%x = call i32 @llvm.ptx.read.nctaid.y()
119	ret i32 %x
120}
121
122define ptx_device i32 @test_nctaid_z() {
123; CHECK: mov.u32 %ret0, %nctaid.z;
124; CHECK: ret;
125	%x = call i32 @llvm.ptx.read.nctaid.z()
126	ret i32 %x
127}
128
129define ptx_device i32 @test_nctaid_w() {
130; CHECK: mov.u32 %ret0, %nctaid.w;
131; CHECK: ret;
132	%x = call i32 @llvm.ptx.read.nctaid.w()
133	ret i32 %x
134}
135
136define ptx_device i32 @test_smid() {
137; CHECK: mov.u32 %ret0, %smid;
138; CHECK: ret;
139	%x = call i32 @llvm.ptx.read.smid()
140	ret i32 %x
141}
142
143define ptx_device i32 @test_nsmid() {
144; CHECK: mov.u32 %ret0, %nsmid;
145; CHECK: ret;
146	%x = call i32 @llvm.ptx.read.nsmid()
147	ret i32 %x
148}
149
150define ptx_device i32 @test_gridid() {
151; CHECK: mov.u32 %ret0, %gridid;
152; CHECK: ret;
153	%x = call i32 @llvm.ptx.read.gridid()
154	ret i32 %x
155}
156
157define ptx_device i32 @test_lanemask_eq() {
158; CHECK: mov.u32 %ret0, %lanemask_eq;
159; CHECK: ret;
160	%x = call i32 @llvm.ptx.read.lanemask.eq()
161	ret i32 %x
162}
163
164define ptx_device i32 @test_lanemask_le() {
165; CHECK: mov.u32 %ret0, %lanemask_le;
166; CHECK: ret;
167	%x = call i32 @llvm.ptx.read.lanemask.le()
168	ret i32 %x
169}
170
171define ptx_device i32 @test_lanemask_lt() {
172; CHECK: mov.u32 %ret0, %lanemask_lt;
173; CHECK: ret;
174	%x = call i32 @llvm.ptx.read.lanemask.lt()
175	ret i32 %x
176}
177
178define ptx_device i32 @test_lanemask_ge() {
179; CHECK: mov.u32 %ret0, %lanemask_ge;
180; CHECK: ret;
181	%x = call i32 @llvm.ptx.read.lanemask.ge()
182	ret i32 %x
183}
184
185define ptx_device i32 @test_lanemask_gt() {
186; CHECK: mov.u32 %ret0, %lanemask_gt;
187; CHECK: ret;
188	%x = call i32 @llvm.ptx.read.lanemask.gt()
189	ret i32 %x
190}
191
192define ptx_device i32 @test_clock() {
193; CHECK: mov.u32 %ret0, %clock;
194; CHECK: ret;
195	%x = call i32 @llvm.ptx.read.clock()
196	ret i32 %x
197}
198
199define ptx_device i64 @test_clock64() {
200; CHECK: mov.u64 %ret0, %clock64;
201; CHECK: ret;
202	%x = call i64 @llvm.ptx.read.clock64()
203	ret i64 %x
204}
205
206define ptx_device i32 @test_pm0() {
207; CHECK: mov.u32 %ret0, %pm0;
208; CHECK: ret;
209	%x = call i32 @llvm.ptx.read.pm0()
210	ret i32 %x
211}
212
213define ptx_device i32 @test_pm1() {
214; CHECK: mov.u32 %ret0, %pm1;
215; CHECK: ret;
216	%x = call i32 @llvm.ptx.read.pm1()
217	ret i32 %x
218}
219
220define ptx_device i32 @test_pm2() {
221; CHECK: mov.u32 %ret0, %pm2;
222; CHECK: ret;
223	%x = call i32 @llvm.ptx.read.pm2()
224	ret i32 %x
225}
226
227define ptx_device i32 @test_pm3() {
228; CHECK: mov.u32 %ret0, %pm3;
229; CHECK: ret;
230	%x = call i32 @llvm.ptx.read.pm3()
231	ret i32 %x
232}
233
234define ptx_device void @test_bar_sync() {
235; CHECK: bar.sync 0
236; CHECK: ret;
237	call void @llvm.ptx.bar.sync(i32 0)
238	ret void
239}
240
241declare i32 @llvm.ptx.read.tid.x()
242declare i32 @llvm.ptx.read.tid.y()
243declare i32 @llvm.ptx.read.tid.z()
244declare i32 @llvm.ptx.read.tid.w()
245declare i32 @llvm.ptx.read.ntid.x()
246declare i32 @llvm.ptx.read.ntid.y()
247declare i32 @llvm.ptx.read.ntid.z()
248declare i32 @llvm.ptx.read.ntid.w()
249
250declare i32 @llvm.ptx.read.laneid()
251declare i32 @llvm.ptx.read.warpid()
252declare i32 @llvm.ptx.read.nwarpid()
253
254declare i32 @llvm.ptx.read.ctaid.x()
255declare i32 @llvm.ptx.read.ctaid.y()
256declare i32 @llvm.ptx.read.ctaid.z()
257declare i32 @llvm.ptx.read.ctaid.w()
258declare i32 @llvm.ptx.read.nctaid.x()
259declare i32 @llvm.ptx.read.nctaid.y()
260declare i32 @llvm.ptx.read.nctaid.z()
261declare i32 @llvm.ptx.read.nctaid.w()
262
263declare i32 @llvm.ptx.read.smid()
264declare i32 @llvm.ptx.read.nsmid()
265declare i32 @llvm.ptx.read.gridid()
266
267declare i32 @llvm.ptx.read.lanemask.eq()
268declare i32 @llvm.ptx.read.lanemask.le()
269declare i32 @llvm.ptx.read.lanemask.lt()
270declare i32 @llvm.ptx.read.lanemask.ge()
271declare i32 @llvm.ptx.read.lanemask.gt()
272
273declare i32 @llvm.ptx.read.clock()
274declare i64 @llvm.ptx.read.clock64()
275
276declare i32 @llvm.ptx.read.pm0()
277declare i32 @llvm.ptx.read.pm1()
278declare i32 @llvm.ptx.read.pm2()
279declare i32 @llvm.ptx.read.pm3()
280
281declare void @llvm.ptx.bar.sync(i32 %i)
282