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