?  search:
summary | log | graphiclog1 | graphiclog2 | commit | commitdiff | tree | refs | edit | fork
blame | history | raw | HEAD
gcc/
[official-gcc.git] / libgomp / oacc-ptx.h
blob2419a460c72f3da619bea184378712694cb47c0b
1 /* Copyright (C) 2014-2015 Free Software Foundation, Inc.
2
3 Contributed by Mentor Embedded.
4
5 This file is part of the GNU Offloading and Multi Processing Library
6 (libgomp).
7
8 Libgomp is free software; you can redistribute it and/or modify it
9 under the terms of the GNU General Public License as published by
10 the Free Software Foundation; either version 3, or (at your option)
11 any later version.
12
13 Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
14 WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
15 FOR A PARTICULAR PURPOSE. See the GNU General Public License for
16 more details.
17
18 Under Section 7 of GPL version 3, you are granted additional
19 permissions described in the GCC Runtime Library Exception, version
20 3.1, as published by the Free Software Foundation.
21
22 You should have received a copy of the GNU General Public License and
23 a copy of the GCC Runtime Library Exception along with this program;
24 see the files COPYING3 and COPYING.RUNTIME respectively. If not, see
25 <http://www.gnu.org/licenses/>. */
26
27 #define ABORT_PTX \
28 ".version 3.1\n" \
29 ".target sm_30\n" \
30 ".address_size 64\n" \
31 ".visible .func abort;\n" \
32 ".visible .func abort\n" \
33 "{\n" \
34 "trap;\n" \
35 "ret;\n" \
36 "}\n" \
37 ".visible .func _gfortran_abort;\n" \
38 ".visible .func _gfortran_abort\n" \
39 "{\n" \
40 "trap;\n" \
41 "ret;\n" \
42 "}\n" \
43
44 /* Generated with:
45
46 $ echo 'int acc_on_device(int d) { return __builtin_acc_on_device(d); } int acc_on_device_h_(int *d) { return acc_on_device(*d); }' | accel-gcc/xgcc -Baccel-gcc -x c - -o - -S -m64 -O3 -fno-builtin-acc_on_device -fno-inline
47 */
48 #define ACC_ON_DEVICE_PTX \
49 " .version 3.1\n" \
50 " .target sm_30\n" \
51 " .address_size 64\n" \
52 ".visible .func (.param.u32 %out_retval)acc_on_device(.param.u32 %in_ar1);\n" \
53 ".visible .func (.param.u32 %out_retval)acc_on_device(.param.u32 %in_ar1)\n" \
54 "{\n" \
55 " .reg.u32 %ar1;\n" \
56 ".reg.u32 %retval;\n" \
57 " .reg.u64 %hr10;\n" \
58 " .reg.u32 %r24;\n" \
59 " .reg.u32 %r25;\n" \
60 " .reg.pred %r27;\n" \
61 " .reg.u32 %r30;\n" \
62 " ld.param.u32 %ar1, [%in_ar1];\n" \
63 " mov.u32 %r24, %ar1;\n" \
64 " setp.ne.u32 %r27,%r24,4;\n" \
65 " set.u32.eq.u32 %r30,%r24,5;\n" \
66 " neg.s32 %r25, %r30;\n" \
67 " @%r27 bra $L3;\n" \
68 " mov.u32 %r25, 1;\n" \
69 "$L3:\n" \
70 " mov.u32 %retval, %r25;\n" \
71 " st.param.u32 [%out_retval], %retval;\n" \
72 " ret;\n" \
73 " }\n" \
74 ".visible .func (.param.u32 %out_retval)acc_on_device_h_(.param.u64 %in_ar1);\n" \
75 ".visible .func (.param.u32 %out_retval)acc_on_device_h_(.param.u64 %in_ar1)\n" \
76 "{\n" \
77 " .reg.u64 %ar1;\n" \
78 ".reg.u32 %retval;\n" \
79 " .reg.u64 %hr10;\n" \
80 " .reg.u64 %r25;\n" \
81 " .reg.u32 %r26;\n" \
82 " .reg.u32 %r27;\n" \
83 " ld.param.u64 %ar1, [%in_ar1];\n" \
84 " mov.u64 %r25, %ar1;\n" \
85 " ld.u32 %r26, [%r25];\n" \
86 " {\n" \
87 " .param.u32 %retval_in;\n" \
88 " {\n" \
89 " .param.u32 %out_arg0;\n" \
90 " st.param.u32 [%out_arg0], %r26;\n" \
91 " call (%retval_in), acc_on_device, (%out_arg0);\n" \
92 " }\n" \
93 " ld.param.u32 %r27, [%retval_in];\n" \
94 "}\n" \
95 " mov.u32 %retval, %r27;\n" \
96 " st.param.u32 [%out_retval], %retval;\n" \
97 " ret;\n" \
98 " }"
99
100 #define GOACC_INTERNAL_PTX \
101 ".version 3.1\n" \
102 ".target sm_30\n" \
103 ".address_size 64\n" \
104 ".visible .func (.param .u32 %out_retval) GOACC_tid (.param .u32 %in_ar1);\n" \
105 ".visible .func (.param .u32 %out_retval) GOACC_ntid (.param .u32 %in_ar1);\n" \
106 ".visible .func (.param .u32 %out_retval) GOACC_ctaid (.param .u32 %in_ar1);\n" \
107 ".visible .func (.param .u32 %out_retval) GOACC_nctaid (.param .u32 %in_ar1);\n" \
108 ".visible .func (.param .u32 %out_retval) GOACC_get_num_threads;\n" \
109 ".visible .func (.param .u32 %out_retval) GOACC_get_thread_num;\n" \
110 ".extern .func abort;\n" \
111 ".visible .func (.param .u32 %out_retval) GOACC_tid (.param .u32 %in_ar1)\n" \
112 "{\n" \
113 ".reg .u32 %ar1;\n" \
114 ".reg .u32 %retval;\n" \
115 ".reg .u64 %hr10;\n" \
116 ".reg .u32 %r22;\n" \
117 ".reg .u32 %r23;\n" \
118 ".reg .u32 %r24;\n" \
119 ".reg .u32 %r25;\n" \
120 ".reg .u32 %r26;\n" \
121 ".reg .u32 %r27;\n" \
122 ".reg .u32 %r28;\n" \
123 ".reg .u32 %r29;\n" \
124 ".reg .pred %r30;\n" \
125 ".reg .u32 %r31;\n" \
126 ".reg .pred %r32;\n" \
127 ".reg .u32 %r33;\n" \
128 ".reg .pred %r34;\n" \
129 ".local .align 8 .b8 %frame[4];\n" \
130 "ld.param.u32 %ar1,[%in_ar1];\n" \
131 "mov.u32 %r27,%ar1;\n" \
132 "st.local.u32 [%frame],%r27;\n" \
133 "ld.local.u32 %r28,[%frame];\n" \
134 "mov.u32 %r29,1;\n" \
135 "setp.eq.u32 %r30,%r28,%r29;\n" \
136 "@%r30 bra $L4;\n" \
137 "mov.u32 %r31,2;\n" \
138 "setp.eq.u32 %r32,%r28,%r31;\n" \
139 "@%r32 bra $L5;\n" \
140 "mov.u32 %r33,0;\n" \
141 "setp.eq.u32 %r34,%r28,%r33;\n" \
142 "@!%r34 bra $L8;\n" \
143 "mov.u32 %r23,%tid.x;\n" \
144 "mov.u32 %r22,%r23;\n" \
145 "bra $L7;\n" \
146 "$L4:\n" \
147 "mov.u32 %r24,%tid.y;\n" \
148 "mov.u32 %r22,%r24;\n" \
149 "bra $L7;\n" \
150 "$L5:\n" \
151 "mov.u32 %r25,%tid.z;\n" \
152 "mov.u32 %r22,%r25;\n" \
153 "bra $L7;\n" \
154 "$L8:\n" \
155 "{\n" \
156 "{\n" \
157 "call abort;\n" \
158 "}\n" \
159 "}\n" \
160 "$L7:\n" \
161 "mov.u32 %r26,%r22;\n" \
162 "mov.u32 %retval,%r26;\n" \
163 "st.param.u32 [%out_retval],%retval;\n" \
164 "ret;\n" \
165 "}\n" \
166 ".visible .func (.param .u32 %out_retval) GOACC_ntid (.param .u32 %in_ar1)\n" \
167 "{\n" \
168 ".reg .u32 %ar1;\n" \
169 ".reg .u32 %retval;\n" \
170 ".reg .u64 %hr10;\n" \
171 ".reg .u32 %r22;\n" \
172 ".reg .u32 %r23;\n" \
173 ".reg .u32 %r24;\n" \
174 ".reg .u32 %r25;\n" \
175 ".reg .u32 %r26;\n" \
176 ".reg .u32 %r27;\n" \
177 ".reg .u32 %r28;\n" \
178 ".reg .u32 %r29;\n" \
179 ".reg .pred %r30;\n" \
180 ".reg .u32 %r31;\n" \
181 ".reg .pred %r32;\n" \
182 ".reg .u32 %r33;\n" \
183 ".reg .pred %r34;\n" \
184 ".local .align 8 .b8 %frame[4];\n" \
185 "ld.param.u32 %ar1,[%in_ar1];\n" \
186 "mov.u32 %r27,%ar1;\n" \
187 "st.local.u32 [%frame],%r27;\n" \
188 "ld.local.u32 %r28,[%frame];\n" \
189 "mov.u32 %r29,1;\n" \
190 "setp.eq.u32 %r30,%r28,%r29;\n" \
191 "@%r30 bra $L11;\n" \
192 "mov.u32 %r31,2;\n" \
193 "setp.eq.u32 %r32,%r28,%r31;\n" \
194 "@%r32 bra $L12;\n" \
195 "mov.u32 %r33,0;\n" \
196 "setp.eq.u32 %r34,%r28,%r33;\n" \
197 "@!%r34 bra $L15;\n" \
198 "mov.u32 %r23,%ntid.x;\n" \
199 "mov.u32 %r22,%r23;\n" \
200 "bra $L14;\n" \
201 "$L11:\n" \
202 "mov.u32 %r24,%ntid.y;\n" \
203 "mov.u32 %r22,%r24;\n" \
204 "bra $L14;\n" \
205 "$L12:\n" \
206 "mov.u32 %r25,%ntid.z;\n" \
207 "mov.u32 %r22,%r25;\n" \
208 "bra $L14;\n" \
209 "$L15:\n" \
210 "{\n" \
211 "{\n" \
212 "call abort;\n" \
213 "}\n" \
214 "}\n" \
215 "$L14:\n" \
216 "mov.u32 %r26,%r22;\n" \
217 "mov.u32 %retval,%r26;\n" \
218 "st.param.u32 [%out_retval],%retval;\n" \
219 "ret;\n" \
220 "}\n" \
221 ".visible .func (.param .u32 %out_retval) GOACC_ctaid (.param .u32 %in_ar1)\n" \
222 "{\n" \
223 ".reg .u32 %ar1;\n" \
224 ".reg .u32 %retval;\n" \
225 ".reg .u64 %hr10;\n" \
226 ".reg .u32 %r22;\n" \
227 ".reg .u32 %r23;\n" \
228 ".reg .u32 %r24;\n" \
229 ".reg .u32 %r25;\n" \
230 ".reg .u32 %r26;\n" \
231 ".reg .u32 %r27;\n" \
232 ".reg .u32 %r28;\n" \
233 ".reg .u32 %r29;\n" \
234 ".reg .pred %r30;\n" \
235 ".reg .u32 %r31;\n" \
236 ".reg .pred %r32;\n" \
237 ".reg .u32 %r33;\n" \
238 ".reg .pred %r34;\n" \
239 ".local .align 8 .b8 %frame[4];\n" \
240 "ld.param.u32 %ar1,[%in_ar1];\n" \
241 "mov.u32 %r27,%ar1;\n" \
242 "st.local.u32 [%frame],%r27;\n" \
243 "ld.local.u32 %r28,[%frame];\n" \
244 "mov.u32 %r29,1;\n" \
245 "setp.eq.u32 %r30,%r28,%r29;\n" \
246 "@%r30 bra $L18;\n" \
247 "mov.u32 %r31,2;\n" \
248 "setp.eq.u32 %r32,%r28,%r31;\n" \
249 "@%r32 bra $L19;\n" \
250 "mov.u32 %r33,0;\n" \
251 "setp.eq.u32 %r34,%r28,%r33;\n" \
252 "@!%r34 bra $L22;\n" \
253 "mov.u32 %r23,%ctaid.x;\n" \
254 "mov.u32 %r22,%r23;\n" \
255 "bra $L21;\n" \
256 "$L18:\n" \
257 "mov.u32 %r24,%ctaid.y;\n" \
258 "mov.u32 %r22,%r24;\n" \
259 "bra $L21;\n" \
260 "$L19:\n" \
261 "mov.u32 %r25,%ctaid.z;\n" \
262 "mov.u32 %r22,%r25;\n" \
263 "bra $L21;\n" \
264 "$L22:\n" \
265 "{\n" \
266 "{\n" \
267 "call abort;\n" \
268 "}\n" \
269 "}\n" \
270 "$L21:\n" \
271 "mov.u32 %r26,%r22;\n" \
272 "mov.u32 %retval,%r26;\n" \
273 "st.param.u32 [%out_retval],%retval;\n" \
274 "ret;\n" \
275 "}\n" \
276 ".visible .func (.param .u32 %out_retval) GOACC_nctaid (.param .u32 %in_ar1)\n" \
277 "{\n" \
278 ".reg .u32 %ar1;\n" \
279 ".reg .u32 %retval;\n" \
280 ".reg .u64 %hr10;\n" \
281 ".reg .u32 %r22;\n" \
282 ".reg .u32 %r23;\n" \
283 ".reg .u32 %r24;\n" \
284 ".reg .u32 %r25;\n" \
285 ".reg .u32 %r26;\n" \
286 ".reg .u32 %r27;\n" \
287 ".reg .u32 %r28;\n" \
288 ".reg .u32 %r29;\n" \
289 ".reg .pred %r30;\n" \
290 ".reg .u32 %r31;\n" \
291 ".reg .pred %r32;\n" \
292 ".reg .u32 %r33;\n" \
293 ".reg .pred %r34;\n" \
294 ".local .align 8 .b8 %frame[4];\n" \
295 "ld.param.u32 %ar1,[%in_ar1];\n" \
296 "mov.u32 %r27,%ar1;\n" \
297 "st.local.u32 [%frame],%r27;\n" \
298 "ld.local.u32 %r28,[%frame];\n" \
299 "mov.u32 %r29,1;\n" \
300 "setp.eq.u32 %r30,%r28,%r29;\n" \
301 "@%r30 bra $L25;\n" \
302 "mov.u32 %r31,2;\n" \
303 "setp.eq.u32 %r32,%r28,%r31;\n" \
304 "@%r32 bra $L26;\n" \
305 "mov.u32 %r33,0;\n" \
306 "setp.eq.u32 %r34,%r28,%r33;\n" \
307 "@!%r34 bra $L29;\n" \
308 "mov.u32 %r23,%nctaid.x;\n" \
309 "mov.u32 %r22,%r23;\n" \
310 "bra $L28;\n" \
311 "$L25:\n" \
312 "mov.u32 %r24,%nctaid.y;\n" \
313 "mov.u32 %r22,%r24;\n" \
314 "bra $L28;\n" \
315 "$L26:\n" \
316 "mov.u32 %r25,%nctaid.z;\n" \
317 "mov.u32 %r22,%r25;\n" \
318 "bra $L28;\n" \
319 "$L29:\n" \
320 "{\n" \
321 "{\n" \
322 "call abort;\n" \
323 "}\n" \
324 "}\n" \
325 "$L28:\n" \
326 "mov.u32 %r26,%r22;\n" \
327 "mov.u32 %retval,%r26;\n" \
328 "st.param.u32 [%out_retval],%retval;\n" \
329 "ret;\n" \
330 "}\n" \
331 ".visible .func (.param .u32 %out_retval) GOACC_get_num_threads\n" \
332 "{\n" \
333 ".reg .u32 %retval;\n" \
334 ".reg .u64 %hr10;\n" \
335 ".reg .u32 %r22;\n" \
336 ".reg .u32 %r23;\n" \
337 ".reg .u32 %r24;\n" \
338 ".reg .u32 %r25;\n" \
339 ".reg .u32 %r26;\n" \
340 ".reg .u32 %r27;\n" \
341 ".reg .u32 %r28;\n" \
342 ".reg .u32 %r29;\n" \
343 "mov.u32 %r26,0;\n" \
344 "{\n" \
345 ".param .u32 %retval_in;\n" \
346 "{\n" \
347 ".param .u32 %out_arg0;\n" \
348 "st.param.u32 [%out_arg0],%r26;\n" \
349 "call (%retval_in),GOACC_ntid,(%out_arg0);\n" \
350 "}\n" \
351 "ld.param.u32 %r27,[%retval_in];\n" \
352 "}\n" \
353 "mov.u32 %r22,%r27;\n" \
354 "mov.u32 %r28,0;\n" \
355 "{\n" \
356 ".param .u32 %retval_in;\n" \
357 "{\n" \
358 ".param .u32 %out_arg0;\n" \
359 "st.param.u32 [%out_arg0],%r28;\n" \
360 "call (%retval_in),GOACC_nctaid,(%out_arg0);\n" \
361 "}\n" \
362 "ld.param.u32 %r29,[%retval_in];\n" \
363 "}\n" \
364 "mov.u32 %r23,%r29;\n" \
365 "mul.lo.u32 %r24,%r22,%r23;\n" \
366 "mov.u32 %r25,%r24;\n" \
367 "mov.u32 %retval,%r25;\n" \
368 "st.param.u32 [%out_retval],%retval;\n" \
369 "ret;\n" \
370 "}\n" \
371 ".visible .func (.param .u32 %out_retval) GOACC_get_thread_num\n" \
372 "{\n" \
373 ".reg .u32 %retval;\n" \
374 ".reg .u64 %hr10;\n" \
375 ".reg .u32 %r22;\n" \
376 ".reg .u32 %r23;\n" \
377 ".reg .u32 %r24;\n" \
378 ".reg .u32 %r25;\n" \
379 ".reg .u32 %r26;\n" \
380 ".reg .u32 %r27;\n" \
381 ".reg .u32 %r28;\n" \
382 ".reg .u32 %r29;\n" \
383 ".reg .u32 %r30;\n" \
384 ".reg .u32 %r31;\n" \
385 ".reg .u32 %r32;\n" \
386 ".reg .u32 %r33;\n" \
387 "mov.u32 %r28,0;\n" \
388 "{\n" \
389 ".param .u32 %retval_in;\n" \
390 "{\n" \
391 ".param .u32 %out_arg0;\n" \
392 "st.param.u32 [%out_arg0],%r28;\n" \
393 "call (%retval_in),GOACC_ntid,(%out_arg0);\n" \
394 "}\n" \
395 "ld.param.u32 %r29,[%retval_in];\n" \
396 "}\n" \
397 "mov.u32 %r22,%r29;\n" \
398 "mov.u32 %r30,0;\n" \
399 "{\n" \
400 ".param .u32 %retval_in;\n" \
401 "{\n" \
402 ".param .u32 %out_arg0;\n" \
403 "st.param.u32 [%out_arg0],%r30;\n" \
404 "call (%retval_in),GOACC_ctaid,(%out_arg0);\n" \
405 "}\n" \
406 "ld.param.u32 %r31,[%retval_in];\n" \
407 "}\n" \
408 "mov.u32 %r23,%r31;\n" \
409 "mul.lo.u32 %r24,%r22,%r23;\n" \
410 "mov.u32 %r32,0;\n" \
411 "{\n" \
412 ".param .u32 %retval_in;\n" \
413 "{\n" \
414 ".param .u32 %out_arg0;\n" \
415 "st.param.u32 [%out_arg0],%r32;\n" \
416 "call (%retval_in),GOACC_tid,(%out_arg0);\n" \
417 "}\n" \
418 "ld.param.u32 %r33,[%retval_in];\n" \
419 "}\n" \
420 "mov.u32 %r25,%r33;\n" \
421 "add.u32 %r26,%r24,%r25;\n" \
422 "mov.u32 %r27,%r26;\n" \
423 "mov.u32 %retval,%r27;\n" \
424 "st.param.u32 [%out_retval],%retval;\n" \
425 "ret;\n" \
426 "}\n"
Mirror of the offical gcc git repository hosted at gcc.gnu.org