1 ;; Machine description for NVPTX.
2 ;; Copyright (C) 2014-2018 Free Software Foundation, Inc.
3 ;; Contributed by Bernd Schmidt <bernds@codesourcery.com>
5 ;; This file is part of GCC.
7 ;; GCC is free software; you can redistribute it and/or modify
8 ;; it under the terms of the GNU General Public License as published by
9 ;; the Free Software Foundation; either version 3, or (at your option)
12 ;; GCC is distributed in the hope that it will be useful,
13 ;; but WITHOUT ANY WARRANTY; without even the implied warranty of
14 ;; MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
15 ;; GNU General Public License for more details.
17 ;; You should have received a copy of the GNU General Public License
18 ;; along with GCC; see the file COPYING3. If not see
19 ;; <http://www.gnu.org/licenses/>.
21 (define_c_enum "unspec" [
33 UNSPEC_FPINT_NEARBYINT
53 (define_c_enum "unspecv" [
71 (define_attr "subregs_ok" "false,true"
72 (const_string "false"))
74 (define_attr "atomic" "false,true"
75 (const_string "false"))
77 ;; The nvptx operand predicates, in general, don't permit subregs and
78 ;; only literal constants, which differ from the generic ones, which
79 ;; permit subregs and symbolc constants (as appropriate)
80 (define_predicate "nvptx_register_operand"
83 return register_operand (op, mode);
86 (define_predicate "nvptx_nonimmediate_operand"
87 (match_code "mem,reg")
89 return (REG_P (op) ? register_operand (op, mode)
90 : memory_operand (op, mode));
93 (define_predicate "nvptx_nonmemory_operand"
94 (match_code "reg,const_int,const_double")
96 return (REG_P (op) ? register_operand (op, mode)
97 : immediate_operand (op, mode));
100 (define_predicate "const0_operand"
101 (and (match_code "const_int")
102 (match_test "op == const0_rtx")))
104 ;; True if this operator is valid for predication.
105 (define_predicate "predicate_operator"
106 (match_code "eq,ne"))
108 (define_predicate "ne_operator"
111 (define_predicate "nvptx_comparison_operator"
112 (match_code "eq,ne,le,ge,lt,gt,leu,geu,ltu,gtu"))
114 (define_predicate "nvptx_float_comparison_operator"
115 (match_code "eq,ne,le,ge,lt,gt,uneq,unle,unge,unlt,ungt,unordered,ordered"))
117 ;; Test for a valid operand for a call instruction.
118 (define_predicate "call_insn_operand"
119 (match_code "symbol_ref,reg")
121 return REG_P (op) || SYMBOL_REF_FUNCTION_P (op);
124 ;; Return true if OP is a call with parallel USEs of the argument
126 (define_predicate "call_operation"
127 (match_code "parallel")
129 int arg_end = XVECLEN (op, 0);
131 for (int i = 1; i < arg_end; i++)
133 rtx elt = XVECEXP (op, 0, i);
135 if (GET_CODE (elt) != USE || !REG_P (XEXP (elt, 0)))
141 (define_attr "predicable" "false,true"
142 (const_string "true"))
145 [(match_operator 0 "predicate_operator"
146 [(match_operand:BI 1 "nvptx_register_operand" "")
147 (match_operand:BI 2 "const0_operand" "")])]
152 (define_constraint "P0"
153 "An integer with the value 0."
154 (and (match_code "const_int")
155 (match_test "ival == 0")))
157 (define_constraint "P1"
158 "An integer with the value 1."
159 (and (match_code "const_int")
160 (match_test "ival == 1")))
162 (define_constraint "Pn"
163 "An integer with the value -1."
164 (and (match_code "const_int")
165 (match_test "ival == -1")))
167 (define_constraint "R"
171 (define_constraint "Ia"
172 "Any integer constant."
173 (and (match_code "const_int") (match_test "true")))
175 (define_mode_iterator QHSDISDFM [QI HI SI DI SF DF])
176 (define_mode_iterator QHSDIM [QI HI SI DI])
177 (define_mode_iterator HSDIM [HI SI DI])
178 (define_mode_iterator BHSDIM [BI HI SI DI])
179 (define_mode_iterator SDIM [SI DI])
180 (define_mode_iterator SDISDFM [SI DI SF DF])
181 (define_mode_iterator QHIM [QI HI])
182 (define_mode_iterator QHSIM [QI HI SI])
183 (define_mode_iterator SDFM [SF DF])
184 (define_mode_iterator SDCM [SC DC])
185 (define_mode_iterator BITS [SI SF])
186 (define_mode_iterator BITD [DI DF])
187 (define_mode_iterator VECIM [V2SI V2DI])
189 ;; This mode iterator allows :P to be used for patterns that operate on
190 ;; pointer-sized quantities. Exactly one of the two alternatives will match.
191 (define_mode_iterator P [(SI "Pmode == SImode") (DI "Pmode == DImode")])
193 ;; We should get away with not defining memory alternatives, since we don't
194 ;; get variables in this mode and pseudos are never spilled.
196 [(set (match_operand:BI 0 "nvptx_register_operand" "=R,R,R")
197 (match_operand:BI 1 "nvptx_nonmemory_operand" "R,P0,Pn"))]
200 %.\\tmov%t0\\t%0, %1;
201 %.\\tsetp.eq.u32\\t%0, 1, 0;
202 %.\\tsetp.eq.u32\\t%0, 1, 1;")
204 (define_insn "*mov<mode>_insn"
205 [(set (match_operand:VECIM 0 "nonimmediate_operand" "=R,R,m")
206 (match_operand:VECIM 1 "general_operand" "Ri,m,R"))]
207 "!MEM_P (operands[0]) || REG_P (operands[1])"
209 if (which_alternative == 1)
210 return "%.\\tld%A1%u1\\t%0, %1;";
211 if (which_alternative == 2)
212 return "%.\\tst%A0%u0\\t%0, %1;";
214 return nvptx_output_mov_insn (operands[0], operands[1]);
216 [(set_attr "subregs_ok" "true")])
218 (define_insn "*mov<mode>_insn"
219 [(set (match_operand:QHSDIM 0 "nonimmediate_operand" "=R,R,m")
220 (match_operand:QHSDIM 1 "general_operand" "Ri,m,R"))]
221 "!MEM_P (operands[0]) || REG_P (operands[1])"
223 if (which_alternative == 1)
224 return "%.\\tld%A1%u1\\t%0, %1;";
225 if (which_alternative == 2)
226 return "%.\\tst%A0%u0\\t%0, %1;";
228 return nvptx_output_mov_insn (operands[0], operands[1]);
230 [(set_attr "subregs_ok" "true")])
232 (define_insn "*mov<mode>_insn"
233 [(set (match_operand:SDFM 0 "nonimmediate_operand" "=R,R,m")
234 (match_operand:SDFM 1 "general_operand" "RF,m,R"))]
235 "!MEM_P (operands[0]) || REG_P (operands[1])"
237 if (which_alternative == 1)
238 return "%.\\tld%A1%u0\\t%0, %1;";
239 if (which_alternative == 2)
240 return "%.\\tst%A0%u1\\t%0, %1;";
242 return nvptx_output_mov_insn (operands[0], operands[1]);
244 [(set_attr "subregs_ok" "true")])
246 (define_insn "load_arg_reg<mode>"
247 [(set (match_operand:QHIM 0 "nvptx_register_operand" "=R")
248 (unspec:QHIM [(match_operand 1 "const_int_operand" "n")]
251 "%.\\tcvt%t0.u32\\t%0, %%ar%1;")
253 (define_insn "load_arg_reg<mode>"
254 [(set (match_operand:SDISDFM 0 "nvptx_register_operand" "=R")
255 (unspec:SDISDFM [(match_operand 1 "const_int_operand" "n")]
258 "%.\\tmov%t0\\t%0, %%ar%1;")
260 (define_expand "mov<mode>"
261 [(set (match_operand:VECIM 0 "nonimmediate_operand" "")
262 (match_operand:VECIM 1 "general_operand" ""))]
265 if (MEM_P (operands[0]) && !REG_P (operands[1]))
267 rtx tmp = gen_reg_rtx (<MODE>mode);
268 emit_move_insn (tmp, operands[1]);
269 emit_move_insn (operands[0], tmp);
274 (define_expand "mov<mode>"
275 [(set (match_operand:QHSDISDFM 0 "nonimmediate_operand" "")
276 (match_operand:QHSDISDFM 1 "general_operand" ""))]
279 if (MEM_P (operands[0]) && !REG_P (operands[1]))
281 rtx tmp = gen_reg_rtx (<MODE>mode);
282 emit_move_insn (tmp, operands[1]);
283 emit_move_insn (operands[0], tmp);
287 if (GET_CODE (operands[1]) == LABEL_REF)
288 sorry ("target cannot support label values");
291 (define_insn "zero_extendqihi2"
292 [(set (match_operand:HI 0 "nvptx_register_operand" "=R,R")
293 (zero_extend:HI (match_operand:QI 1 "nvptx_nonimmediate_operand" "R,m")))]
296 %.\\tcvt.u16.u%T1\\t%0, %1;
297 %.\\tld%A1.u8\\t%0, %1;"
298 [(set_attr "subregs_ok" "true")])
300 (define_insn "zero_extend<mode>si2"
301 [(set (match_operand:SI 0 "nvptx_register_operand" "=R,R")
302 (zero_extend:SI (match_operand:QHIM 1 "nvptx_nonimmediate_operand" "R,m")))]
305 %.\\tcvt.u32.u%T1\\t%0, %1;
306 %.\\tld%A1.u%T1\\t%0, %1;"
307 [(set_attr "subregs_ok" "true")])
309 (define_insn "zero_extend<mode>di2"
310 [(set (match_operand:DI 0 "nvptx_register_operand" "=R,R")
311 (zero_extend:DI (match_operand:QHSIM 1 "nvptx_nonimmediate_operand" "R,m")))]
314 %.\\tcvt.u64.u%T1\\t%0, %1;
315 %.\\tld%A1%u1\\t%0, %1;"
316 [(set_attr "subregs_ok" "true")])
318 (define_insn "extend<mode>si2"
319 [(set (match_operand:SI 0 "nvptx_register_operand" "=R,R")
320 (sign_extend:SI (match_operand:QHIM 1 "nvptx_nonimmediate_operand" "R,m")))]
323 %.\\tcvt.s32.s%T1\\t%0, %1;
324 %.\\tld%A1.s%T1\\t%0, %1;"
325 [(set_attr "subregs_ok" "true")])
327 (define_insn "extend<mode>di2"
328 [(set (match_operand:DI 0 "nvptx_register_operand" "=R,R")
329 (sign_extend:DI (match_operand:QHSIM 1 "nvptx_nonimmediate_operand" "R,m")))]
332 %.\\tcvt.s64.s%T1\\t%0, %1;
333 %.\\tld%A1.s%T1\\t%0, %1;"
334 [(set_attr "subregs_ok" "true")])
336 (define_insn "trunchiqi2"
337 [(set (match_operand:QI 0 "nvptx_nonimmediate_operand" "=R,m")
338 (truncate:QI (match_operand:HI 1 "nvptx_register_operand" "R,R")))]
341 %.\\tcvt%t0.u16\\t%0, %1;
342 %.\\tst%A0.u8\\t%0, %1;"
343 [(set_attr "subregs_ok" "true")])
345 (define_insn "truncsi<mode>2"
346 [(set (match_operand:QHIM 0 "nvptx_nonimmediate_operand" "=R,m")
347 (truncate:QHIM (match_operand:SI 1 "nvptx_register_operand" "R,R")))]
350 %.\\tcvt%t0.u32\\t%0, %1;
351 %.\\tst%A0.u%T0\\t%0, %1;"
352 [(set_attr "subregs_ok" "true")])
354 (define_insn "truncdi<mode>2"
355 [(set (match_operand:QHSIM 0 "nvptx_nonimmediate_operand" "=R,m")
356 (truncate:QHSIM (match_operand:DI 1 "nvptx_register_operand" "R,R")))]
359 %.\\tcvt%t0.u64\\t%0, %1;
360 %.\\tst%A0.u%T0\\t%0, %1;"
361 [(set_attr "subregs_ok" "true")])
363 ;; Integer arithmetic
365 (define_insn "add<mode>3"
366 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
367 (plus:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")
368 (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
370 "%.\\tadd%t0\\t%0, %1, %2;")
372 (define_insn "sub<mode>3"
373 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
374 (minus:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")
375 (match_operand:HSDIM 2 "nvptx_register_operand" "R")))]
377 "%.\\tsub%t0\\t%0, %1, %2;")
379 (define_insn "mul<mode>3"
380 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
381 (mult:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")
382 (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
384 "%.\\tmul.lo%t0\\t%0, %1, %2;")
386 (define_insn "*mad<mode>3"
387 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
388 (plus:HSDIM (mult:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")
389 (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri"))
390 (match_operand:HSDIM 3 "nvptx_nonmemory_operand" "Ri")))]
392 "%.\\tmad.lo%t0\\t%0, %1, %2, %3;")
394 (define_insn "div<mode>3"
395 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
396 (div:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")
397 (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
399 "%.\\tdiv.s%T0\\t%0, %1, %2;")
401 (define_insn "udiv<mode>3"
402 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
403 (udiv:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")
404 (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
406 "%.\\tdiv.u%T0\\t%0, %1, %2;")
408 (define_insn "mod<mode>3"
409 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
410 (mod:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "Ri")
411 (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
413 "%.\\trem.s%T0\\t%0, %1, %2;")
415 (define_insn "umod<mode>3"
416 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
417 (umod:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "Ri")
418 (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
420 "%.\\trem.u%T0\\t%0, %1, %2;")
422 (define_insn "smin<mode>3"
423 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
424 (smin:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")
425 (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
427 "%.\\tmin.s%T0\\t%0, %1, %2;")
429 (define_insn "umin<mode>3"
430 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
431 (umin:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")
432 (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
434 "%.\\tmin.u%T0\\t%0, %1, %2;")
436 (define_insn "smax<mode>3"
437 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
438 (smax:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")
439 (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
441 "%.\\tmax.s%T0\\t%0, %1, %2;")
443 (define_insn "umax<mode>3"
444 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
445 (umax:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")
446 (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
448 "%.\\tmax.u%T0\\t%0, %1, %2;")
450 (define_insn "abs<mode>2"
451 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
452 (abs:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")))]
454 "%.\\tabs.s%T0\\t%0, %1;")
456 (define_insn "neg<mode>2"
457 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
458 (neg:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")))]
460 "%.\\tneg.s%T0\\t%0, %1;")
462 (define_insn "one_cmpl<mode>2"
463 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
464 (not:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")))]
466 "%.\\tnot.b%T0\\t%0, %1;")
468 (define_insn "bitrev<mode>2"
469 [(set (match_operand:SDIM 0 "nvptx_register_operand" "=R")
470 (unspec:SDIM [(match_operand:SDIM 1 "nvptx_register_operand" "R")]
473 "%.\\tbrev.b%T0\\t%0, %1;")
475 (define_insn "clz<mode>2"
476 [(set (match_operand:SI 0 "nvptx_register_operand" "=R")
477 (clz:SI (match_operand:SDIM 1 "nvptx_register_operand" "R")))]
479 "%.\\tclz.b%T1\\t%0, %1;")
481 (define_expand "ctz<mode>2"
482 [(set (match_operand:SI 0 "nvptx_register_operand" "")
483 (ctz:SI (match_operand:SDIM 1 "nvptx_register_operand" "")))]
486 rtx tmpreg = gen_reg_rtx (<MODE>mode);
487 emit_insn (gen_bitrev<mode>2 (tmpreg, operands[1]));
488 emit_insn (gen_clz<mode>2 (operands[0], tmpreg));
494 (define_insn "ashl<mode>3"
495 [(set (match_operand:SDIM 0 "nvptx_register_operand" "=R")
496 (ashift:SDIM (match_operand:SDIM 1 "nvptx_register_operand" "R")
497 (match_operand:SI 2 "nvptx_nonmemory_operand" "Ri")))]
499 "%.\\tshl.b%T0\\t%0, %1, %2;")
501 (define_insn "ashr<mode>3"
502 [(set (match_operand:SDIM 0 "nvptx_register_operand" "=R")
503 (ashiftrt:SDIM (match_operand:SDIM 1 "nvptx_register_operand" "R")
504 (match_operand:SI 2 "nvptx_nonmemory_operand" "Ri")))]
506 "%.\\tshr.s%T0\\t%0, %1, %2;")
508 (define_insn "lshr<mode>3"
509 [(set (match_operand:SDIM 0 "nvptx_register_operand" "=R")
510 (lshiftrt:SDIM (match_operand:SDIM 1 "nvptx_register_operand" "R")
511 (match_operand:SI 2 "nvptx_nonmemory_operand" "Ri")))]
513 "%.\\tshr.u%T0\\t%0, %1, %2;")
515 ;; Logical operations
517 (define_insn "and<mode>3"
518 [(set (match_operand:BHSDIM 0 "nvptx_register_operand" "=R")
519 (and:BHSDIM (match_operand:BHSDIM 1 "nvptx_register_operand" "R")
520 (match_operand:BHSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
522 "%.\\tand.b%T0\\t%0, %1, %2;")
524 (define_insn "ior<mode>3"
525 [(set (match_operand:BHSDIM 0 "nvptx_register_operand" "=R")
526 (ior:BHSDIM (match_operand:BHSDIM 1 "nvptx_register_operand" "R")
527 (match_operand:BHSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
529 "%.\\tor.b%T0\\t%0, %1, %2;")
531 (define_insn "xor<mode>3"
532 [(set (match_operand:BHSDIM 0 "nvptx_register_operand" "=R")
533 (xor:BHSDIM (match_operand:BHSDIM 1 "nvptx_register_operand" "R")
534 (match_operand:BHSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
536 "%.\\txor.b%T0\\t%0, %1, %2;")
538 ;; Comparisons and branches
540 (define_insn "*cmp<mode>"
541 [(set (match_operand:BI 0 "nvptx_register_operand" "=R")
542 (match_operator:BI 1 "nvptx_comparison_operator"
543 [(match_operand:HSDIM 2 "nvptx_register_operand" "R")
544 (match_operand:HSDIM 3 "nvptx_nonmemory_operand" "Ri")]))]
546 "%.\\tsetp%c1\\t%0, %2, %3;")
548 (define_insn "*cmp<mode>"
549 [(set (match_operand:BI 0 "nvptx_register_operand" "=R")
550 (match_operator:BI 1 "nvptx_float_comparison_operator"
551 [(match_operand:SDFM 2 "nvptx_register_operand" "R")
552 (match_operand:SDFM 3 "nvptx_nonmemory_operand" "RF")]))]
554 "%.\\tsetp%c1\\t%0, %2, %3;")
558 (label_ref (match_operand 0 "" "")))]
562 (define_insn "br_true"
564 (if_then_else (ne (match_operand:BI 0 "nvptx_register_operand" "R")
566 (label_ref (match_operand 1 "" ""))
570 [(set_attr "predicable" "false")])
572 (define_insn "br_false"
574 (if_then_else (eq (match_operand:BI 0 "nvptx_register_operand" "R")
576 (label_ref (match_operand 1 "" ""))
580 [(set_attr "predicable" "false")])
582 ;; unified conditional branch
583 (define_insn "br_true_uni"
584 [(set (pc) (if_then_else
585 (ne (unspec:BI [(match_operand:BI 0 "nvptx_register_operand" "R")]
586 UNSPEC_BR_UNIFIED) (const_int 0))
587 (label_ref (match_operand 1 "" "")) (pc)))]
589 "%j0\\tbra.uni\\t%l1;"
590 [(set_attr "predicable" "false")])
592 (define_insn "br_false_uni"
593 [(set (pc) (if_then_else
594 (eq (unspec:BI [(match_operand:BI 0 "nvptx_register_operand" "R")]
595 UNSPEC_BR_UNIFIED) (const_int 0))
596 (label_ref (match_operand 1 "" "")) (pc)))]
598 "%J0\\tbra.uni\\t%l1;"
599 [(set_attr "predicable" "false")])
601 (define_expand "cbranch<mode>4"
603 (if_then_else (match_operator 0 "nvptx_comparison_operator"
604 [(match_operand:HSDIM 1 "nvptx_register_operand" "")
605 (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "")])
606 (label_ref (match_operand 3 "" ""))
610 rtx t = nvptx_expand_compare (operands[0]);
612 operands[1] = XEXP (t, 0);
613 operands[2] = XEXP (t, 1);
616 (define_expand "cbranch<mode>4"
618 (if_then_else (match_operator 0 "nvptx_float_comparison_operator"
619 [(match_operand:SDFM 1 "nvptx_register_operand" "")
620 (match_operand:SDFM 2 "nvptx_nonmemory_operand" "")])
621 (label_ref (match_operand 3 "" ""))
625 rtx t = nvptx_expand_compare (operands[0]);
627 operands[1] = XEXP (t, 0);
628 operands[2] = XEXP (t, 1);
631 (define_expand "cbranchbi4"
633 (if_then_else (match_operator 0 "predicate_operator"
634 [(match_operand:BI 1 "nvptx_register_operand" "")
635 (match_operand:BI 2 "const0_operand" "")])
636 (label_ref (match_operand 3 "" ""))
641 ;; Conditional stores
643 (define_insn "setcc_from_bi"
644 [(set (match_operand:SI 0 "nvptx_register_operand" "=R")
645 (ne:SI (match_operand:BI 1 "nvptx_register_operand" "R")
648 "%.\\tselp%t0 %0,-1,0,%1;")
650 (define_insn "sel_true<mode>"
651 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
653 (ne (match_operand:BI 1 "nvptx_register_operand" "R") (const_int 0))
654 (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")
655 (match_operand:HSDIM 3 "nvptx_nonmemory_operand" "Ri")))]
657 "%.\\tselp%t0\\t%0, %2, %3, %1;")
659 (define_insn "sel_true<mode>"
660 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
662 (ne (match_operand:BI 1 "nvptx_register_operand" "R") (const_int 0))
663 (match_operand:SDFM 2 "nvptx_nonmemory_operand" "RF")
664 (match_operand:SDFM 3 "nvptx_nonmemory_operand" "RF")))]
666 "%.\\tselp%t0\\t%0, %2, %3, %1;")
668 (define_insn "sel_false<mode>"
669 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
671 (eq (match_operand:BI 1 "nvptx_register_operand" "R") (const_int 0))
672 (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")
673 (match_operand:HSDIM 3 "nvptx_nonmemory_operand" "Ri")))]
675 "%.\\tselp%t0\\t%0, %3, %2, %1;")
677 (define_insn "sel_false<mode>"
678 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
680 (eq (match_operand:BI 1 "nvptx_register_operand" "R") (const_int 0))
681 (match_operand:SDFM 2 "nvptx_nonmemory_operand" "RF")
682 (match_operand:SDFM 3 "nvptx_nonmemory_operand" "RF")))]
684 "%.\\tselp%t0\\t%0, %3, %2, %1;")
686 (define_insn "setcc_int<mode>"
687 [(set (match_operand:SI 0 "nvptx_register_operand" "=R")
688 (match_operator:SI 1 "nvptx_comparison_operator"
689 [(match_operand:HSDIM 2 "nvptx_register_operand" "R")
690 (match_operand:HSDIM 3 "nvptx_nonmemory_operand" "Ri")]))]
692 "%.\\tset%t0%c1\\t%0, %2, %3;")
694 (define_insn "setcc_int<mode>"
695 [(set (match_operand:SI 0 "nvptx_register_operand" "=R")
696 (match_operator:SI 1 "nvptx_float_comparison_operator"
697 [(match_operand:SDFM 2 "nvptx_register_operand" "R")
698 (match_operand:SDFM 3 "nvptx_nonmemory_operand" "RF")]))]
700 "%.\\tset%t0%c1\\t%0, %2, %3;")
702 (define_insn "setcc_float<mode>"
703 [(set (match_operand:SF 0 "nvptx_register_operand" "=R")
704 (match_operator:SF 1 "nvptx_comparison_operator"
705 [(match_operand:HSDIM 2 "nvptx_register_operand" "R")
706 (match_operand:HSDIM 3 "nvptx_nonmemory_operand" "Ri")]))]
708 "%.\\tset%t0%c1\\t%0, %2, %3;")
710 (define_insn "setcc_float<mode>"
711 [(set (match_operand:SF 0 "nvptx_register_operand" "=R")
712 (match_operator:SF 1 "nvptx_float_comparison_operator"
713 [(match_operand:SDFM 2 "nvptx_register_operand" "R")
714 (match_operand:SDFM 3 "nvptx_nonmemory_operand" "RF")]))]
716 "%.\\tset%t0%c1\\t%0, %2, %3;")
718 (define_expand "cstorebi4"
719 [(set (match_operand:SI 0 "nvptx_register_operand")
720 (match_operator:SI 1 "ne_operator"
721 [(match_operand:BI 2 "nvptx_register_operand")
722 (match_operand:BI 3 "const0_operand")]))]
726 (define_expand "cstore<mode>4"
727 [(set (match_operand:SI 0 "nvptx_register_operand")
728 (match_operator:SI 1 "nvptx_comparison_operator"
729 [(match_operand:HSDIM 2 "nvptx_register_operand")
730 (match_operand:HSDIM 3 "nvptx_nonmemory_operand")]))]
734 (define_expand "cstore<mode>4"
735 [(set (match_operand:SI 0 "nvptx_register_operand")
736 (match_operator:SI 1 "nvptx_float_comparison_operator"
737 [(match_operand:SDFM 2 "nvptx_register_operand")
738 (match_operand:SDFM 3 "nvptx_nonmemory_operand")]))]
744 (define_insn "call_insn"
745 [(match_parallel 2 "call_operation"
746 [(call (mem:QI (match_operand 0 "call_insn_operand" "Rs"))
747 (match_operand 1))])]
750 return nvptx_output_call_insn (insn, NULL_RTX, operands[0]);
753 (define_insn "call_value_insn"
754 [(match_parallel 3 "call_operation"
755 [(set (match_operand 0 "nvptx_register_operand" "=R")
756 (call (mem:QI (match_operand 1 "call_insn_operand" "Rs"))
757 (match_operand 2)))])]
760 return nvptx_output_call_insn (insn, operands[0], operands[1]);
763 (define_expand "call"
764 [(match_operand 0 "" "")]
767 nvptx_expand_call (NULL_RTX, operands[0]);
771 (define_expand "call_value"
772 [(match_operand 0 "" "")
773 (match_operand 1 "" "")]
776 nvptx_expand_call (operands[0], operands[1]);
780 ;; Floating point arithmetic.
782 (define_insn "add<mode>3"
783 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
784 (plus:SDFM (match_operand:SDFM 1 "nvptx_register_operand" "R")
785 (match_operand:SDFM 2 "nvptx_nonmemory_operand" "RF")))]
787 "%.\\tadd%t0\\t%0, %1, %2;")
789 (define_insn "sub<mode>3"
790 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
791 (minus:SDFM (match_operand:SDFM 1 "nvptx_register_operand" "R")
792 (match_operand:SDFM 2 "nvptx_register_operand" "R")))]
794 "%.\\tsub%t0\\t%0, %1, %2;")
796 (define_insn "mul<mode>3"
797 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
798 (mult:SDFM (match_operand:SDFM 1 "nvptx_register_operand" "R")
799 (match_operand:SDFM 2 "nvptx_nonmemory_operand" "RF")))]
801 "%.\\tmul%t0\\t%0, %1, %2;")
803 (define_insn "fma<mode>4"
804 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
805 (fma:SDFM (match_operand:SDFM 1 "nvptx_register_operand" "R")
806 (match_operand:SDFM 2 "nvptx_nonmemory_operand" "RF")
807 (match_operand:SDFM 3 "nvptx_nonmemory_operand" "RF")))]
809 "%.\\tfma%#%t0\\t%0, %1, %2, %3;")
811 (define_insn "div<mode>3"
812 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
813 (div:SDFM (match_operand:SDFM 1 "nvptx_register_operand" "R")
814 (match_operand:SDFM 2 "nvptx_nonmemory_operand" "RF")))]
816 "%.\\tdiv%#%t0\\t%0, %1, %2;")
818 (define_insn "copysign<mode>3"
819 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
820 (unspec:SDFM [(match_operand:SDFM 1 "nvptx_register_operand" "R")
821 (match_operand:SDFM 2 "nvptx_register_operand" "R")]
824 "%.\\tcopysign%t0\\t%0, %2, %1;")
826 (define_insn "smin<mode>3"
827 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
828 (smin:SDFM (match_operand:SDFM 1 "nvptx_register_operand" "R")
829 (match_operand:SDFM 2 "nvptx_nonmemory_operand" "RF")))]
831 "%.\\tmin%t0\\t%0, %1, %2;")
833 (define_insn "smax<mode>3"
834 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
835 (smax:SDFM (match_operand:SDFM 1 "nvptx_register_operand" "R")
836 (match_operand:SDFM 2 "nvptx_nonmemory_operand" "RF")))]
838 "%.\\tmax%t0\\t%0, %1, %2;")
840 (define_insn "abs<mode>2"
841 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
842 (abs:SDFM (match_operand:SDFM 1 "nvptx_register_operand" "R")))]
844 "%.\\tabs%t0\\t%0, %1;")
846 (define_insn "neg<mode>2"
847 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
848 (neg:SDFM (match_operand:SDFM 1 "nvptx_register_operand" "R")))]
850 "%.\\tneg%t0\\t%0, %1;")
852 (define_insn "sqrt<mode>2"
853 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
854 (sqrt:SDFM (match_operand:SDFM 1 "nvptx_register_operand" "R")))]
856 "%.\\tsqrt%#%t0\\t%0, %1;")
858 (define_expand "sincossf3"
859 [(set (match_operand:SF 0 "nvptx_register_operand" "=R")
860 (unspec:SF [(match_operand:SF 2 "nvptx_register_operand" "R")]
862 (set (match_operand:SF 1 "nvptx_register_operand" "=R")
863 (unspec:SF [(match_dup 2)] UNSPEC_SIN))]
864 "flag_unsafe_math_optimizations"
866 operands[2] = make_safe_from (operands[2], operands[0]);
869 (define_insn "sinsf2"
870 [(set (match_operand:SF 0 "nvptx_register_operand" "=R")
871 (unspec:SF [(match_operand:SF 1 "nvptx_register_operand" "R")]
873 "flag_unsafe_math_optimizations"
874 "%.\\tsin.approx%t0\\t%0, %1;")
876 (define_insn "cossf2"
877 [(set (match_operand:SF 0 "nvptx_register_operand" "=R")
878 (unspec:SF [(match_operand:SF 1 "nvptx_register_operand" "R")]
880 "flag_unsafe_math_optimizations"
881 "%.\\tcos.approx%t0\\t%0, %1;")
883 (define_insn "log2sf2"
884 [(set (match_operand:SF 0 "nvptx_register_operand" "=R")
885 (unspec:SF [(match_operand:SF 1 "nvptx_register_operand" "R")]
887 "flag_unsafe_math_optimizations"
888 "%.\\tlg2.approx%t0\\t%0, %1;")
890 (define_insn "exp2sf2"
891 [(set (match_operand:SF 0 "nvptx_register_operand" "=R")
892 (unspec:SF [(match_operand:SF 1 "nvptx_register_operand" "R")]
894 "flag_unsafe_math_optimizations"
895 "%.\\tex2.approx%t0\\t%0, %1;")
897 ;; Conversions involving floating point
899 (define_insn "extendsfdf2"
900 [(set (match_operand:DF 0 "nvptx_register_operand" "=R")
901 (float_extend:DF (match_operand:SF 1 "nvptx_register_operand" "R")))]
903 "%.\\tcvt%t0%t1\\t%0, %1;")
905 (define_insn "truncdfsf2"
906 [(set (match_operand:SF 0 "nvptx_register_operand" "=R")
907 (float_truncate:SF (match_operand:DF 1 "nvptx_register_operand" "R")))]
909 "%.\\tcvt%#%t0%t1\\t%0, %1;")
911 (define_insn "floatunssi<mode>2"
912 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
913 (unsigned_float:SDFM (match_operand:SI 1 "nvptx_register_operand" "R")))]
915 "%.\\tcvt%#%t0.u%T1\\t%0, %1;")
917 (define_insn "floatsi<mode>2"
918 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
919 (float:SDFM (match_operand:SI 1 "nvptx_register_operand" "R")))]
921 "%.\\tcvt%#%t0.s%T1\\t%0, %1;")
923 (define_insn "floatunsdi<mode>2"
924 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
925 (unsigned_float:SDFM (match_operand:DI 1 "nvptx_register_operand" "R")))]
927 "%.\\tcvt%#%t0.u%T1\\t%0, %1;")
929 (define_insn "floatdi<mode>2"
930 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
931 (float:SDFM (match_operand:DI 1 "nvptx_register_operand" "R")))]
933 "%.\\tcvt%#%t0.s%T1\\t%0, %1;")
935 (define_insn "fixuns_trunc<mode>si2"
936 [(set (match_operand:SI 0 "nvptx_register_operand" "=R")
937 (unsigned_fix:SI (match_operand:SDFM 1 "nvptx_register_operand" "R")))]
939 "%.\\tcvt.rzi.u%T0%t1\\t%0, %1;")
941 (define_insn "fix_trunc<mode>si2"
942 [(set (match_operand:SI 0 "nvptx_register_operand" "=R")
943 (fix:SI (match_operand:SDFM 1 "nvptx_register_operand" "R")))]
945 "%.\\tcvt.rzi.s%T0%t1\\t%0, %1;")
947 (define_insn "fixuns_trunc<mode>di2"
948 [(set (match_operand:DI 0 "nvptx_register_operand" "=R")
949 (unsigned_fix:DI (match_operand:SDFM 1 "nvptx_register_operand" "R")))]
951 "%.\\tcvt.rzi.u%T0%t1\\t%0, %1;")
953 (define_insn "fix_trunc<mode>di2"
954 [(set (match_operand:DI 0 "nvptx_register_operand" "=R")
955 (fix:DI (match_operand:SDFM 1 "nvptx_register_operand" "R")))]
957 "%.\\tcvt.rzi.s%T0%t1\\t%0, %1;")
959 (define_int_iterator FPINT [UNSPEC_FPINT_FLOOR UNSPEC_FPINT_BTRUNC
960 UNSPEC_FPINT_CEIL UNSPEC_FPINT_NEARBYINT])
961 (define_int_attr fpint_name [(UNSPEC_FPINT_FLOOR "floor")
962 (UNSPEC_FPINT_BTRUNC "btrunc")
963 (UNSPEC_FPINT_CEIL "ceil")
964 (UNSPEC_FPINT_NEARBYINT "nearbyint")])
965 (define_int_attr fpint_roundingmode [(UNSPEC_FPINT_FLOOR ".rmi")
966 (UNSPEC_FPINT_BTRUNC ".rzi")
967 (UNSPEC_FPINT_CEIL ".rpi")
968 (UNSPEC_FPINT_NEARBYINT "%#i")])
970 (define_insn "<FPINT:fpint_name><SDFM:mode>2"
971 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
972 (unspec:SDFM [(match_operand:SDFM 1 "nvptx_register_operand" "R")]
975 "%.\\tcvt<FPINT:fpint_roundingmode>%t0%t1\\t%0, %1;")
977 (define_int_iterator FPINT2 [UNSPEC_FPINT_FLOOR UNSPEC_FPINT_CEIL])
978 (define_int_attr fpint2_name [(UNSPEC_FPINT_FLOOR "lfloor")
979 (UNSPEC_FPINT_CEIL "lceil")])
980 (define_int_attr fpint2_roundingmode [(UNSPEC_FPINT_FLOOR ".rmi")
981 (UNSPEC_FPINT_CEIL ".rpi")])
983 (define_insn "<FPINT2:fpint2_name><SDFM:mode><SDIM:mode>2"
984 [(set (match_operand:SDIM 0 "nvptx_register_operand" "=R")
985 (unspec:SDIM [(match_operand:SDFM 1 "nvptx_register_operand" "R")]
988 "%.\\tcvt<FPINT2:fpint2_roundingmode>.s%T0%t1\\t%0, %1;")
1002 (define_insn "fake_nop"
1006 .reg .u32 %%nop_src;
1007 .reg .u32 %%nop_dst;
1008 mov.u32 %%nop_dst, %%nop_src;
1011 (define_insn "return"
1015 return nvptx_output_return ();
1017 [(set_attr "predicable" "false")])
1019 (define_expand "epilogue"
1020 [(clobber (const_int 0))]
1023 if (TARGET_SOFT_STACK)
1024 emit_insn (gen_set_softstack_insn (gen_rtx_REG (Pmode,
1025 SOFTSTACK_PREV_REGNUM)));
1026 emit_jump_insn (gen_return ());
1030 (define_expand "nonlocal_goto"
1031 [(match_operand 0 "" "")
1032 (match_operand 1 "" "")
1033 (match_operand 2 "" "")
1034 (match_operand 3 "" "")]
1037 sorry ("target cannot support nonlocal goto.");
1038 emit_insn (gen_nop ());
1042 (define_expand "nonlocal_goto_receiver"
1046 sorry ("target cannot support nonlocal goto.");
1049 (define_expand "allocate_stack"
1050 [(match_operand 0 "nvptx_register_operand")
1051 (match_operand 1 "nvptx_register_operand")]
1054 if (TARGET_SOFT_STACK)
1056 emit_move_insn (stack_pointer_rtx,
1057 gen_rtx_MINUS (Pmode, stack_pointer_rtx, operands[1]));
1058 emit_insn (gen_set_softstack_insn (stack_pointer_rtx));
1059 emit_move_insn (operands[0], virtual_stack_dynamic_rtx);
1062 /* The ptx documentation specifies an alloca intrinsic (for 32 bit
1063 only) but notes it is not implemented. The assembler emits a
1064 confused error message. Issue a blunt one now instead. */
1065 sorry ("target cannot support alloca.");
1066 emit_insn (gen_nop ());
1070 (define_insn "set_softstack_insn"
1071 [(unspec [(match_operand 0 "nvptx_register_operand" "R")]
1072 UNSPEC_SET_SOFTSTACK)]
1075 return nvptx_output_set_softstack (REGNO (operands[0]));
1078 (define_expand "restore_stack_block"
1079 [(match_operand 0 "register_operand" "")
1080 (match_operand 1 "register_operand" "")]
1083 if (TARGET_SOFT_STACK)
1085 emit_move_insn (operands[0], operands[1]);
1086 emit_insn (gen_set_softstack_insn (operands[0]));
1091 (define_expand "restore_stack_function"
1092 [(match_operand 0 "register_operand" "")
1093 (match_operand 1 "register_operand" "")]
1100 [(trap_if (const_int 1) (const_int 0))]
1104 (define_insn "trap_if_true"
1105 [(trap_if (ne (match_operand:BI 0 "nvptx_register_operand" "R")
1110 [(set_attr "predicable" "false")])
1112 (define_insn "trap_if_false"
1113 [(trap_if (eq (match_operand:BI 0 "nvptx_register_operand" "R")
1118 [(set_attr "predicable" "false")])
1120 (define_expand "ctrap<mode>4"
1121 [(trap_if (match_operator 0 "nvptx_comparison_operator"
1122 [(match_operand:SDIM 1 "nvptx_register_operand")
1123 (match_operand:SDIM 2 "nvptx_nonmemory_operand")])
1124 (match_operand 3 "const0_operand"))]
1127 rtx t = nvptx_expand_compare (operands[0]);
1128 emit_insn (gen_trap_if_true (t));
1132 (define_insn "oacc_dim_size"
1133 [(set (match_operand:SI 0 "nvptx_register_operand" "")
1134 (unspec:SI [(match_operand:SI 1 "const_int_operand" "")]
1138 static const char *const asms[] =
1139 { /* Must match oacc_loop_levels ordering. */
1140 "%.\\tmov.u32\\t%0, %%nctaid.x;", /* gang */
1141 "%.\\tmov.u32\\t%0, %%ntid.y;", /* worker */
1142 "%.\\tmov.u32\\t%0, %%ntid.x;", /* vector */
1144 return asms[INTVAL (operands[1])];
1147 (define_insn "oacc_dim_pos"
1148 [(set (match_operand:SI 0 "nvptx_register_operand" "")
1149 (unspec_volatile:SI [(match_operand:SI 1 "const_int_operand" "")]
1153 static const char *const asms[] =
1154 { /* Must match oacc_loop_levels ordering. */
1155 "%.\\tmov.u32\\t%0, %%ctaid.x;", /* gang */
1156 "%.\\tmov.u32\\t%0, %%tid.y;", /* worker */
1157 "%.\\tmov.u32\\t%0, %%tid.x;", /* vector */
1159 return asms[INTVAL (operands[1])];
1162 (define_insn "nvptx_fork"
1163 [(unspec_volatile:SI [(match_operand:SI 0 "const_int_operand" "")]
1167 [(set_attr "predicable" "false")])
1169 (define_insn "nvptx_forked"
1170 [(unspec_volatile:SI [(match_operand:SI 0 "const_int_operand" "")]
1174 [(set_attr "predicable" "false")])
1176 (define_insn "nvptx_joining"
1177 [(unspec_volatile:SI [(match_operand:SI 0 "const_int_operand" "")]
1181 [(set_attr "predicable" "false")])
1183 (define_insn "nvptx_join"
1184 [(unspec_volatile:SI [(match_operand:SI 0 "const_int_operand" "")]
1188 [(set_attr "predicable" "false")])
1190 (define_expand "oacc_fork"
1191 [(set (match_operand:SI 0 "nvptx_nonmemory_operand" "")
1192 (match_operand:SI 1 "general_operand" ""))
1193 (unspec_volatile:SI [(match_operand:SI 2 "const_int_operand" "")]
1197 if (operands[0] != const0_rtx)
1198 emit_move_insn (operands[0], operands[1]);
1199 nvptx_expand_oacc_fork (INTVAL (operands[2]));
1203 (define_expand "oacc_join"
1204 [(set (match_operand:SI 0 "nvptx_nonmemory_operand" "")
1205 (match_operand:SI 1 "general_operand" ""))
1206 (unspec_volatile:SI [(match_operand:SI 2 "const_int_operand" "")]
1210 if (operands[0] != const0_rtx)
1211 emit_move_insn (operands[0], operands[1]);
1212 nvptx_expand_oacc_join (INTVAL (operands[2]));
1216 ;; only 32-bit shuffles exist.
1217 (define_insn "nvptx_shuffle<mode>"
1218 [(set (match_operand:BITS 0 "nvptx_register_operand" "=R")
1220 [(match_operand:BITS 1 "nvptx_register_operand" "R")
1221 (match_operand:SI 2 "nvptx_nonmemory_operand" "Ri")
1222 (match_operand:SI 3 "const_int_operand" "n")]
1225 "%.\\tshfl%S3.b32\\t%0, %1, %2, 31;")
1227 (define_insn "nvptx_vote_ballot"
1228 [(set (match_operand:SI 0 "nvptx_register_operand" "=R")
1229 (unspec:SI [(match_operand:BI 1 "nvptx_register_operand" "R")]
1230 UNSPEC_VOTE_BALLOT))]
1232 "%.\\tvote.ballot.b32\\t%0, %1;")
1234 ;; Patterns for OpenMP SIMD-via-SIMT lowering
1236 (define_insn "omp_simt_enter_insn"
1237 [(set (match_operand 0 "nvptx_register_operand" "=R")
1238 (unspec_volatile [(match_operand 1 "nvptx_nonmemory_operand" "Ri")
1239 (match_operand 2 "nvptx_nonmemory_operand" "Ri")]
1240 UNSPECV_SIMT_ENTER))]
1243 return nvptx_output_simt_enter (operands[0], operands[1], operands[2]);
1246 (define_expand "omp_simt_enter"
1247 [(match_operand 0 "nvptx_register_operand" "=R")
1248 (match_operand 1 "nvptx_nonmemory_operand" "Ri")
1249 (match_operand 2 "const_int_operand" "n")]
1252 if (!CONST_INT_P (operands[1]))
1253 cfun->machine->simt_stack_size = HOST_WIDE_INT_M1U;
1255 cfun->machine->simt_stack_size = MAX (UINTVAL (operands[1]),
1256 cfun->machine->simt_stack_size);
1257 cfun->machine->simt_stack_align = MAX (UINTVAL (operands[2]),
1258 cfun->machine->simt_stack_align);
1259 cfun->machine->has_simtreg = true;
1260 emit_insn (gen_omp_simt_enter_insn (operands[0], operands[1], operands[2]));
1264 (define_insn "omp_simt_exit"
1265 [(unspec_volatile [(match_operand 0 "nvptx_register_operand" "R")]
1269 return nvptx_output_simt_exit (operands[0]);
1272 ;; Implement IFN_GOMP_SIMT_LANE: set operand 0 to lane index
1273 (define_insn "omp_simt_lane"
1274 [(set (match_operand:SI 0 "nvptx_register_operand" "")
1275 (unspec:SI [(const_int 0)] UNSPEC_LANEID))]
1277 "%.\\tmov.u32\\t%0, %%laneid;")
1279 ;; Implement IFN_GOMP_SIMT_ORDERED: copy operand 1 to operand 0 and
1280 ;; place a compiler barrier to disallow unrolling/peeling the containing loop
1281 (define_expand "omp_simt_ordered"
1282 [(match_operand:SI 0 "nvptx_register_operand" "=R")
1283 (match_operand:SI 1 "nvptx_register_operand" "R")]
1286 emit_move_insn (operands[0], operands[1]);
1287 emit_insn (gen_nvptx_nounroll ());
1291 ;; Implement IFN_GOMP_SIMT_XCHG_BFLY: perform a "butterfly" exchange
1293 (define_expand "omp_simt_xchg_bfly"
1294 [(match_operand 0 "nvptx_register_operand" "=R")
1295 (match_operand 1 "nvptx_register_operand" "R")
1296 (match_operand:SI 2 "nvptx_nonmemory_operand" "Ri")]
1299 emit_insn (nvptx_gen_shuffle (operands[0], operands[1], operands[2],
1304 ;; Implement IFN_GOMP_SIMT_XCHG_IDX: broadcast value in operand 1
1305 ;; from lane given by index in operand 2 to operand 0 in all lanes
1306 (define_expand "omp_simt_xchg_idx"
1307 [(match_operand 0 "nvptx_register_operand" "=R")
1308 (match_operand 1 "nvptx_register_operand" "R")
1309 (match_operand:SI 2 "nvptx_nonmemory_operand" "Ri")]
1312 emit_insn (nvptx_gen_shuffle (operands[0], operands[1], operands[2],
1317 ;; Implement IFN_GOMP_SIMT_VOTE_ANY:
1318 ;; set operand 0 to zero iff all lanes supply zero in operand 1
1319 (define_expand "omp_simt_vote_any"
1320 [(match_operand:SI 0 "nvptx_register_operand" "=R")
1321 (match_operand:SI 1 "nvptx_register_operand" "R")]
1324 rtx pred = gen_reg_rtx (BImode);
1325 emit_move_insn (pred, gen_rtx_NE (BImode, operands[1], const0_rtx));
1326 emit_insn (gen_nvptx_vote_ballot (operands[0], pred));
1330 ;; Implement IFN_GOMP_SIMT_LAST_LANE:
1331 ;; set operand 0 to the lowest lane index that passed non-zero in operand 1
1332 (define_expand "omp_simt_last_lane"
1333 [(match_operand:SI 0 "nvptx_register_operand" "=R")
1334 (match_operand:SI 1 "nvptx_register_operand" "R")]
1337 rtx pred = gen_reg_rtx (BImode);
1338 rtx tmp = gen_reg_rtx (SImode);
1339 emit_move_insn (pred, gen_rtx_NE (BImode, operands[1], const0_rtx));
1340 emit_insn (gen_nvptx_vote_ballot (tmp, pred));
1341 emit_insn (gen_ctzsi2 (operands[0], tmp));
1345 ;; extract parts of a 64 bit object into 2 32-bit ints
1346 (define_insn "unpack<mode>si2"
1347 [(set (match_operand:SI 0 "nvptx_register_operand" "=R")
1348 (unspec:SI [(match_operand:BITD 2 "nvptx_register_operand" "R")
1349 (const_int 0)] UNSPEC_BIT_CONV))
1350 (set (match_operand:SI 1 "nvptx_register_operand" "=R")
1351 (unspec:SI [(match_dup 2) (const_int 1)] UNSPEC_BIT_CONV))]
1353 "%.\\tmov.b64\\t{%0,%1}, %2;")
1355 ;; pack 2 32-bit ints into a 64 bit object
1356 (define_insn "packsi<mode>2"
1357 [(set (match_operand:BITD 0 "nvptx_register_operand" "=R")
1358 (unspec:BITD [(match_operand:SI 1 "nvptx_register_operand" "R")
1359 (match_operand:SI 2 "nvptx_register_operand" "R")]
1362 "%.\\tmov.b64\\t%0, {%1,%2};")
1366 (define_expand "atomic_compare_and_swap<mode>"
1367 [(match_operand:SI 0 "nvptx_register_operand") ;; bool success output
1368 (match_operand:SDIM 1 "nvptx_register_operand") ;; oldval output
1369 (match_operand:SDIM 2 "memory_operand") ;; memory
1370 (match_operand:SDIM 3 "nvptx_register_operand") ;; expected input
1371 (match_operand:SDIM 4 "nvptx_register_operand") ;; newval input
1372 (match_operand:SI 5 "const_int_operand") ;; is_weak
1373 (match_operand:SI 6 "const_int_operand") ;; success model
1374 (match_operand:SI 7 "const_int_operand")] ;; failure model
1377 emit_insn (gen_atomic_compare_and_swap<mode>_1
1378 (operands[1], operands[2], operands[3], operands[4], operands[6]));
1380 rtx cond = gen_reg_rtx (BImode);
1381 emit_move_insn (cond, gen_rtx_EQ (BImode, operands[1], operands[3]));
1382 emit_insn (gen_sel_truesi (operands[0], cond, GEN_INT (1), GEN_INT (0)));
1386 (define_insn "atomic_compare_and_swap<mode>_1"
1387 [(set (match_operand:SDIM 0 "nvptx_register_operand" "=R")
1388 (unspec_volatile:SDIM
1389 [(match_operand:SDIM 1 "memory_operand" "+m")
1390 (match_operand:SDIM 2 "nvptx_nonmemory_operand" "Ri")
1391 (match_operand:SDIM 3 "nvptx_nonmemory_operand" "Ri")
1392 (match_operand:SI 4 "const_int_operand")]
1395 (unspec_volatile:SDIM [(const_int 0)] UNSPECV_CAS))]
1397 "%.\\tatom%A1.cas.b%T0\\t%0, %1, %2, %3;"
1398 [(set_attr "atomic" "true")])
1400 (define_insn "atomic_exchange<mode>"
1401 [(set (match_operand:SDIM 0 "nvptx_register_operand" "=R") ;; output
1402 (unspec_volatile:SDIM
1403 [(match_operand:SDIM 1 "memory_operand" "+m") ;; memory
1404 (match_operand:SI 3 "const_int_operand")] ;; model
1407 (match_operand:SDIM 2 "nvptx_nonmemory_operand" "Ri"))] ;; input
1409 "%.\\tatom%A1.exch.b%T0\\t%0, %1, %2;"
1410 [(set_attr "atomic" "true")])
1412 (define_insn "atomic_fetch_add<mode>"
1413 [(set (match_operand:SDIM 1 "memory_operand" "+m")
1414 (unspec_volatile:SDIM
1415 [(plus:SDIM (match_dup 1)
1416 (match_operand:SDIM 2 "nvptx_nonmemory_operand" "Ri"))
1417 (match_operand:SI 3 "const_int_operand")] ;; model
1419 (set (match_operand:SDIM 0 "nvptx_register_operand" "=R")
1422 "%.\\tatom%A1.add%t0\\t%0, %1, %2;"
1423 [(set_attr "atomic" "true")])
1425 (define_insn "atomic_fetch_addsf"
1426 [(set (match_operand:SF 1 "memory_operand" "+m")
1428 [(plus:SF (match_dup 1)
1429 (match_operand:SF 2 "nvptx_nonmemory_operand" "RF"))
1430 (match_operand:SI 3 "const_int_operand")] ;; model
1432 (set (match_operand:SF 0 "nvptx_register_operand" "=R")
1435 "%.\\tatom%A1.add%t0\\t%0, %1, %2;"
1436 [(set_attr "atomic" "true")])
1438 (define_code_iterator any_logic [and ior xor])
1439 (define_code_attr logic [(and "and") (ior "or") (xor "xor")])
1441 ;; Currently disabled until we add better subtarget support - requires sm_32.
1442 (define_insn "atomic_fetch_<logic><mode>"
1443 [(set (match_operand:SDIM 1 "memory_operand" "+m")
1444 (unspec_volatile:SDIM
1445 [(any_logic:SDIM (match_dup 1)
1446 (match_operand:SDIM 2 "nvptx_nonmemory_operand" "Ri"))
1447 (match_operand:SI 3 "const_int_operand")] ;; model
1449 (set (match_operand:SDIM 0 "nvptx_register_operand" "=R")
1452 "%.\\tatom%A1.b%T0.<logic>\\t%0, %1, %2;"
1453 [(set_attr "atomic" "true")])
1455 (define_insn "nvptx_barsync"
1456 [(unspec_volatile [(match_operand:SI 0 "const_int_operand" "")]
1460 [(set_attr "predicable" "false")])
1462 (define_insn "nvptx_nounroll"
1463 [(unspec_volatile [(const_int 0)] UNSPECV_NOUNROLL)]
1465 "\\t.pragma \\\"nounroll\\\";"
1466 [(set_attr "predicable" "false")])