1 ;; Machine description for NVPTX.
2 ;; Copyright (C) 2014-2023 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" [
35 UNSPEC_FPINT_NEARBYINT
55 (define_c_enum "unspecv" [
66 UNSPECV_UNIFORM_WARP_CHECK
85 (define_attr "subregs_ok" "false,true"
86 (const_string "false"))
88 (define_attr "atomic" "false,true"
89 (const_string "false"))
91 ;; The nvptx operand predicates, in general, don't permit subregs and
92 ;; only literal constants, which differ from the generic ones, which
93 ;; permit subregs and symbolc constants (as appropriate)
94 (define_predicate "nvptx_register_operand"
97 return register_operand (op, mode);
100 (define_predicate "nvptx_register_or_complex_di_df_register_operand"
101 (ior (match_code "reg")
102 (match_code "concat"))
104 if (GET_CODE (op) == CONCAT)
105 return ((GET_MODE (op) == DCmode || GET_MODE (op) == CDImode)
106 && nvptx_register_operand (XEXP (op, 0), mode)
107 && nvptx_register_operand (XEXP (op, 1), mode));
109 return nvptx_register_operand (op, mode);
112 (define_predicate "nvptx_nonimmediate_operand"
113 (match_code "mem,reg")
115 return (REG_P (op) ? register_operand (op, mode)
116 : memory_operand (op, mode));
119 (define_predicate "nvptx_nonmemory_operand"
120 (match_code "reg,const_int,const_double")
122 return (REG_P (op) ? register_operand (op, mode)
123 : immediate_operand (op, mode));
126 (define_predicate "const0_operand"
127 (and (match_code "const_int")
128 (match_test "op == const0_rtx")))
130 ;; True if this operator is valid for predication.
131 (define_predicate "predicate_operator"
132 (match_code "eq,ne"))
134 (define_predicate "ne_operator"
137 (define_predicate "nvptx_comparison_operator"
138 (match_code "eq,ne,le,ge,lt,gt,leu,geu,ltu,gtu"))
140 (define_predicate "nvptx_float_comparison_operator"
141 (match_code "eq,ne,le,ge,lt,gt,uneq,unle,unge,unlt,ungt,unordered,ordered"))
143 (define_predicate "nvptx_vector_index_operand"
144 (and (match_code "const_int")
145 (match_test "UINTVAL (op) < 4")))
147 ;; Test for a valid operand for a call instruction.
148 (define_predicate "call_insn_operand"
149 (match_code "symbol_ref,reg")
151 return REG_P (op) || SYMBOL_REF_FUNCTION_P (op);
154 ;; Return true if OP is a call with parallel USEs of the argument
156 (define_predicate "call_operation"
157 (match_code "parallel")
159 int arg_end = XVECLEN (op, 0);
161 for (int i = 1; i < arg_end; i++)
163 rtx elt = XVECEXP (op, 0, i);
165 if (GET_CODE (elt) != USE || !REG_P (XEXP (elt, 0)))
171 ;; Test for a function symbol ref operand
172 (define_predicate "symbol_ref_function_operand"
173 (match_code "symbol_ref")
175 return SYMBOL_REF_FUNCTION_P (op);
178 (define_attr "predicable" "no,yes"
179 (const_string "yes"))
182 [(match_operator 0 "predicate_operator"
183 [(match_operand:BI 1 "nvptx_register_operand" "")
184 (match_operand:BI 2 "const0_operand" "")])]
189 (define_constraint "P0"
190 "An integer with the value 0."
191 (and (match_code "const_int")
192 (match_test "ival == 0")))
194 (define_constraint "P1"
195 "An integer with the value 1."
196 (and (match_code "const_int")
197 (match_test "ival == 1")))
199 (define_constraint "Pn"
200 "An integer with the value -1."
201 (and (match_code "const_int")
202 (match_test "ival == -1")))
204 (define_constraint "R"
208 (define_constraint "Ia"
209 "Any integer constant."
210 (and (match_code "const_int") (match_test "true")))
212 (define_mode_iterator QHSDISDFM [QI HI SI DI SF DF])
213 (define_mode_iterator QHSDIM [QI HI SI DI])
214 (define_mode_iterator HSDIM [HI SI DI])
215 (define_mode_iterator BHSDIM [BI HI SI DI])
216 (define_mode_iterator SDIM [SI DI])
217 (define_mode_iterator SDISDFM [SI DI SF DF])
218 (define_mode_iterator QHIM [QI HI])
219 (define_mode_iterator QHSIM [QI HI SI])
220 (define_mode_iterator SDFM [SF DF])
221 (define_mode_iterator HSFM [HF SF])
222 (define_mode_iterator SDCM [SC DC])
223 (define_mode_iterator BITS [SI SF])
224 (define_mode_iterator BITD [DI DF])
225 (define_mode_iterator VECIM [V2SI V2DI])
227 ;; This mode iterator allows :P to be used for patterns that operate on
228 ;; pointer-sized quantities. Exactly one of the two alternatives will match.
229 (define_mode_iterator P [(SI "Pmode == SImode") (DI "Pmode == DImode")])
231 ;; Define element mode for each vector mode.
232 (define_mode_attr VECELEM [(V2SI "SI") (V2DI "DI")])
233 (define_mode_attr Vecelem [(V2SI "si") (V2DI "di")])
235 ;; We should get away with not defining memory alternatives, since we don't
236 ;; get variables in this mode and pseudos are never spilled.
238 [(set (match_operand:BI 0 "nvptx_register_operand" "=R,R,R")
239 (match_operand:BI 1 "nvptx_nonmemory_operand" "R,P0,P1"))]
242 %.\\tmov%t0\\t%0, %1;
243 %.\\tsetp.eq.u32\\t%0, 1, 0;
244 %.\\tsetp.eq.u32\\t%0, 1, 1;")
246 (define_insn "*mov<mode>_insn"
247 [(set (match_operand:VECIM 0 "nonimmediate_operand" "=R,R,m")
248 (match_operand:VECIM 1 "general_operand" "Ri,m,R"))]
249 "!MEM_P (operands[0]) || REG_P (operands[1])"
251 if (which_alternative == 1)
252 return "%.\\tld%A1%u1\\t%0, %1;";
253 if (which_alternative == 2)
254 return "%.\\tst%A0%u0\\t%0, %1;";
256 return nvptx_output_mov_insn (operands[0], operands[1]);
258 [(set_attr "subregs_ok" "true")])
260 (define_insn "*mov<mode>_insn"
261 [(set (match_operand:QHSDIM 0 "nonimmediate_operand" "=R,R,m")
262 (match_operand:QHSDIM 1 "general_operand" "Ri,m,R"))]
263 "!MEM_P (operands[0]) || REG_P (operands[1])"
265 if (which_alternative == 1)
266 return "%.\\tld%A1%u1\\t%0, %1;";
267 if (which_alternative == 2)
268 return "%.\\tst%A0%u0\\t%0, %1;";
270 return nvptx_output_mov_insn (operands[0], operands[1]);
272 [(set_attr "subregs_ok" "true")])
274 ;; ptxas segfaults on 'mov.u64 %r24,bar+4096', so break it up.
276 [(set (match_operand:DI 0 "nvptx_register_operand")
277 (const:DI (plus:DI (match_operand:DI 1 "symbol_ref_function_operand")
278 (match_operand 2 "const_int_operand"))))]
280 [(set (match_dup 0) (match_dup 1))
281 (set (match_dup 0) (plus:DI (match_dup 0) (match_dup 2)))
285 (define_insn "*mov<mode>_insn"
286 [(set (match_operand:SDFM 0 "nonimmediate_operand" "=R,R,m")
287 (match_operand:SDFM 1 "general_operand" "RF,m,R"))]
288 "!MEM_P (operands[0]) || REG_P (operands[1])"
290 if (which_alternative == 1)
291 return "%.\\tld%A1%u0\\t%0, %1;";
292 if (which_alternative == 2)
293 return "%.\\tst%A0%u1\\t%0, %1;";
295 return nvptx_output_mov_insn (operands[0], operands[1]);
297 [(set_attr "subregs_ok" "true")])
299 (define_insn "*movhf_insn"
300 [(set (match_operand:HF 0 "nonimmediate_operand" "=R,R,m")
301 (match_operand:HF 1 "nonimmediate_operand" "R,m,R"))]
302 "!MEM_P (operands[0]) || REG_P (operands[1])"
304 %.\\tmov.b16\\t%0, %1;
305 %.\\tld.b16\\t%0, %1;
306 %.\\tst.b16\\t%0, %1;"
307 [(set_attr "subregs_ok" "true")])
309 (define_expand "movhf"
310 [(set (match_operand:HF 0 "nonimmediate_operand" "")
311 (match_operand:HF 1 "nonimmediate_operand" ""))]
314 /* Load HFmode constants as SFmode with an explicit FLOAT_TRUNCATE. */
315 if (CONST_DOUBLE_P (operands[1]))
317 rtx tmp1 = gen_reg_rtx (SFmode);
318 REAL_VALUE_TYPE d = *CONST_DOUBLE_REAL_VALUE (operands[1]);
319 real_convert (&d, SFmode, &d);
320 emit_move_insn (tmp1, const_double_from_real_value (d, SFmode));
322 if (!REG_P (operands[0]))
324 rtx tmp2 = gen_reg_rtx (HFmode);
325 emit_insn (gen_truncsfhf2 (tmp2, tmp1));
326 emit_move_insn (operands[0], tmp2);
329 emit_insn (gen_truncsfhf2 (operands[0], tmp1));
333 if (MEM_P (operands[0]) && !REG_P (operands[1]))
335 rtx tmp = gen_reg_rtx (HFmode);
336 emit_move_insn (tmp, operands[1]);
337 emit_move_insn (operands[0], tmp);
342 (define_insn "load_arg_reg<mode>"
343 [(set (match_operand:QHIM 0 "nvptx_register_operand" "=R")
344 (unspec:QHIM [(match_operand 1 "const_int_operand" "n")]
347 "%.\\tcvt%t0.u32\\t%0, %%ar%1;")
349 (define_insn "load_arg_reg<mode>"
350 [(set (match_operand:SDISDFM 0 "nvptx_register_operand" "=R")
351 (unspec:SDISDFM [(match_operand 1 "const_int_operand" "n")]
354 "%.\\tmov%t0\\t%0, %%ar%1;")
356 (define_expand "mov<mode>"
357 [(set (match_operand:VECIM 0 "nonimmediate_operand" "")
358 (match_operand:VECIM 1 "general_operand" ""))]
361 if (MEM_P (operands[0]) && !REG_P (operands[1]))
363 rtx tmp = gen_reg_rtx (<MODE>mode);
364 emit_move_insn (tmp, operands[1]);
365 emit_move_insn (operands[0], tmp);
370 (define_expand "mov<mode>"
371 [(set (match_operand:QHSDISDFM 0 "nonimmediate_operand" "")
372 (match_operand:QHSDISDFM 1 "general_operand" ""))]
375 if (MEM_P (operands[0]) && !REG_P (operands[1]))
377 rtx tmp = gen_reg_rtx (<MODE>mode);
378 emit_move_insn (tmp, operands[1]);
379 emit_move_insn (operands[0], tmp);
383 if (GET_CODE (operands[1]) == LABEL_REF)
384 sorry ("target cannot support label values");
387 (define_insn "zero_extendqihi2"
388 [(set (match_operand:HI 0 "nvptx_register_operand" "=R,R")
389 (zero_extend:HI (match_operand:QI 1 "nvptx_nonimmediate_operand" "R,m")))]
392 %.\\tcvt.u16.u%T1\\t%0, %1;
393 %.\\tld%A1.u8\\t%0, %1;"
394 [(set_attr "subregs_ok" "true")])
396 (define_insn "zero_extend<mode>si2"
397 [(set (match_operand:SI 0 "nvptx_register_operand" "=R,R")
398 (zero_extend:SI (match_operand:QHIM 1 "nvptx_nonimmediate_operand" "R,m")))]
401 %.\\tcvt.u32.u%T1\\t%0, %1;
402 %.\\tld%A1.u%T1\\t%0, %1;"
403 [(set_attr "subregs_ok" "true")])
405 (define_insn "zero_extend<mode>di2"
406 [(set (match_operand:DI 0 "nvptx_register_operand" "=R,R")
407 (zero_extend:DI (match_operand:QHSIM 1 "nvptx_nonimmediate_operand" "R,m")))]
410 %.\\tcvt.u64.u%T1\\t%0, %1;
411 %.\\tld%A1%u1\\t%0, %1;"
412 [(set_attr "subregs_ok" "true")])
414 (define_insn "extendqihi2"
415 [(set (match_operand:HI 0 "nvptx_register_operand" "=R")
416 (sign_extend:HI (match_operand:QI 1 "nvptx_register_operand" "R")))]
418 "%.\\tcvt.s16.s8\\t%0, %1;"
419 [(set_attr "subregs_ok" "true")])
421 (define_insn "extend<mode>si2"
422 [(set (match_operand:SI 0 "nvptx_register_operand" "=R,R")
423 (sign_extend:SI (match_operand:QHIM 1 "nvptx_nonimmediate_operand" "R,m")))]
426 %.\\tcvt.s32.s%T1\\t%0, %1;
427 %.\\tld%A1.s%T1\\t%0, %1;"
428 [(set_attr "subregs_ok" "true")])
430 (define_insn "extend<mode>di2"
431 [(set (match_operand:DI 0 "nvptx_register_operand" "=R,R")
432 (sign_extend:DI (match_operand:QHSIM 1 "nvptx_nonimmediate_operand" "R,m")))]
435 %.\\tcvt.s64.s%T1\\t%0, %1;
436 %.\\tld%A1.s%T1\\t%0, %1;"
437 [(set_attr "subregs_ok" "true")])
439 (define_insn "trunchiqi2"
440 [(set (match_operand:QI 0 "nvptx_nonimmediate_operand" "=R,m")
441 (truncate:QI (match_operand:HI 1 "nvptx_register_operand" "R,R")))]
444 %.\\tcvt%t0.u16\\t%0, %1;
445 %.\\tst%A0.u8\\t%0, %1;"
446 [(set_attr "subregs_ok" "true")])
448 (define_insn "truncsi<mode>2"
449 [(set (match_operand:QHIM 0 "nvptx_nonimmediate_operand" "=R,m")
450 (truncate:QHIM (match_operand:SI 1 "nvptx_register_operand" "R,R")))]
453 if (which_alternative == 1)
454 return "%.\\tst%A0.u%T0\\t%0, %1;";
455 if (GET_MODE (operands[0]) == QImode)
456 return "%.\\tmov%t0\\t%0, %1;";
457 return "%.\\tcvt%t0.u32\\t%0, %1;";
459 [(set_attr "subregs_ok" "true")])
461 (define_insn "truncdi<mode>2"
462 [(set (match_operand:QHSIM 0 "nvptx_nonimmediate_operand" "=R,m")
463 (truncate:QHSIM (match_operand:DI 1 "nvptx_register_operand" "R,R")))]
466 %.\\tcvt%t0.u64\\t%0, %1;
467 %.\\tst%A0.u%T0\\t%0, %1;"
468 [(set_attr "subregs_ok" "true")])
470 ;; Sign-extensions of truncations
472 (define_insn "*extend_trunc_<mode>2_qi"
473 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
475 (truncate:QI (match_operand:HSDIM 1 "nvptx_register_operand" "R"))))]
477 "%.\\tcvt.s%T0.s8\\t%0, %1;"
478 [(set_attr "subregs_ok" "true")])
480 (define_insn "*extend_trunc_<mode>2_hi"
481 [(set (match_operand:SDIM 0 "nvptx_register_operand" "=R")
483 (truncate:HI (match_operand:SDIM 1 "nvptx_register_operand" "R"))))]
485 "%.\\tcvt.s%T0.s16\\t%0, %1;"
486 [(set_attr "subregs_ok" "true")])
488 (define_insn "*extend_trunc_di2_si"
489 [(set (match_operand:DI 0 "nvptx_register_operand" "=R")
491 (truncate:SI (match_operand:DI 1 "nvptx_register_operand" "R"))))]
493 "%.\\tcvt.s64.s32\\t%0, %1;"
494 [(set_attr "subregs_ok" "true")])
496 ;; Integer arithmetic
498 (define_insn "add<mode>3"
499 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
500 (plus:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")
501 (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
503 "%.\\tadd%t0\\t%0, %1, %2;")
505 (define_insn "*vadd_addsi4"
506 [(set (match_operand:SI 0 "nvptx_register_operand" "=R")
507 (plus:SI (plus:SI (match_operand:SI 1 "nvptx_register_operand" "R")
508 (match_operand:SI 2 "nvptx_register_operand" "R"))
509 (match_operand:SI 3 "nvptx_register_operand" "R")))]
511 "%.\\tvadd%t0%t1%t2.add\\t%0, %1, %2, %3;")
513 (define_insn "*vsub_addsi4"
514 [(set (match_operand:SI 0 "nvptx_register_operand" "=R")
515 (plus:SI (minus:SI (match_operand:SI 1 "nvptx_register_operand" "R")
516 (match_operand:SI 2 "nvptx_register_operand" "R"))
517 (match_operand:SI 3 "nvptx_register_operand" "R")))]
519 "%.\\tvsub%t0%t1%t2.add\\t%0, %1, %2, %3;")
521 (define_insn "sub<mode>3"
522 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
523 (minus:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")
524 (match_operand:HSDIM 2 "nvptx_register_operand" "R")))]
527 if (GET_MODE (operands[0]) == HImode)
528 /* Workaround https://developer.nvidia.com/nvidia_bug/3527713.
530 return "%.\\tsub.s16\\t%0, %1, %2;";
532 return "%.\\tsub%t0\\t%0, %1, %2;";
535 (define_insn "mul<mode>3"
536 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
537 (mult:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")
538 (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
540 "%.\\tmul.lo%t0\\t%0, %1, %2;")
542 (define_insn "*mad<mode>3"
543 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
544 (plus:HSDIM (mult:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")
545 (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri"))
546 (match_operand:HSDIM 3 "nvptx_nonmemory_operand" "Ri")))]
548 "%.\\tmad.lo%t0\\t%0, %1, %2, %3;")
550 (define_insn "div<mode>3"
551 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
552 (div:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")
553 (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
555 "%.\\tdiv.s%T0\\t%0, %1, %2;")
557 (define_insn "udiv<mode>3"
558 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
559 (udiv:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")
560 (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
562 "%.\\tdiv.u%T0\\t%0, %1, %2;")
564 (define_insn "mod<mode>3"
565 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
566 (mod:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "Ri")
567 (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
569 "%.\\trem.s%T0\\t%0, %1, %2;")
571 (define_insn "umod<mode>3"
572 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
573 (umod:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "Ri")
574 (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
576 "%.\\trem.u%T0\\t%0, %1, %2;")
578 (define_insn "smin<mode>3"
579 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
580 (smin:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")
581 (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
583 "%.\\tmin.s%T0\\t%0, %1, %2;")
585 (define_insn "umin<mode>3"
586 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
587 (umin:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")
588 (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
590 "%.\\tmin.u%T0\\t%0, %1, %2;")
592 (define_insn "smax<mode>3"
593 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
594 (smax:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")
595 (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
597 "%.\\tmax.s%T0\\t%0, %1, %2;")
599 (define_insn "umax<mode>3"
600 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
601 (umax:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")
602 (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
604 "%.\\tmax.u%T0\\t%0, %1, %2;")
606 (define_insn "abs<mode>2"
607 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
608 (abs:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")))]
610 "%.\\tabs.s%T0\\t%0, %1;")
612 (define_insn "neg<mode>2"
613 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
614 (neg:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")))]
616 "%.\\tneg.s%T0\\t%0, %1;")
618 (define_insn "one_cmpl<mode>2"
619 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
620 (not:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")))]
622 "%.\\tnot.b%T0\\t%0, %1;")
624 (define_insn "one_cmplbi2"
625 [(set (match_operand:BI 0 "nvptx_register_operand" "=R")
626 (not:BI (match_operand:BI 1 "nvptx_register_operand" "R")))]
628 "%.\\tnot.pred\\t%0, %1;")
630 (define_insn "*cnot<mode>2"
631 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
632 (eq:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")
635 "%.\\tcnot.b%T0\\t%0, %1;")
637 (define_insn "bitrev<mode>2"
638 [(set (match_operand:SDIM 0 "nvptx_register_operand" "=R")
639 (unspec:SDIM [(match_operand:SDIM 1 "nvptx_register_operand" "R")]
642 "%.\\tbrev.b%T0\\t%0, %1;")
644 (define_insn "clz<mode>2"
645 [(set (match_operand:SI 0 "nvptx_register_operand" "=R")
646 (clz:SI (match_operand:SDIM 1 "nvptx_register_operand" "R")))]
648 "%.\\tclz.b%T1\\t%0, %1;")
650 (define_expand "ctz<mode>2"
651 [(set (match_operand:SI 0 "nvptx_register_operand" "")
652 (ctz:SI (match_operand:SDIM 1 "nvptx_register_operand" "")))]
655 rtx tmpreg = gen_reg_rtx (<MODE>mode);
656 emit_insn (gen_bitrev<mode>2 (tmpreg, operands[1]));
657 emit_insn (gen_clz<mode>2 (operands[0], tmpreg));
661 (define_insn "popcount<mode>2"
662 [(set (match_operand:SI 0 "nvptx_register_operand" "=R")
663 (popcount:SI (match_operand:SDIM 1 "nvptx_register_operand" "R")))]
665 "%.\\tpopc.b%T1\\t%0, %1;")
667 ;; Multiplication variants
669 (define_insn "mulhisi3"
670 [(set (match_operand:SI 0 "nvptx_register_operand" "=R")
671 (mult:SI (sign_extend:SI
672 (match_operand:HI 1 "nvptx_register_operand" "R"))
674 (match_operand:HI 2 "nvptx_register_operand" "R"))))]
676 "%.\\tmul.wide.s16\\t%0, %1, %2;")
678 (define_insn "mulsidi3"
679 [(set (match_operand:DI 0 "nvptx_register_operand" "=R")
680 (mult:DI (sign_extend:DI
681 (match_operand:SI 1 "nvptx_register_operand" "R"))
683 (match_operand:SI 2 "nvptx_register_operand" "R"))))]
685 "%.\\tmul.wide.s32\\t%0, %1, %2;")
687 (define_insn "umulhisi3"
688 [(set (match_operand:SI 0 "nvptx_register_operand" "=R")
689 (mult:SI (zero_extend:SI
690 (match_operand:HI 1 "nvptx_register_operand" "R"))
692 (match_operand:HI 2 "nvptx_register_operand" "R"))))]
694 "%.\\tmul.wide.u16\\t%0, %1, %2;")
696 (define_insn "umulsidi3"
697 [(set (match_operand:DI 0 "nvptx_register_operand" "=R")
698 (mult:DI (zero_extend:DI
699 (match_operand:SI 1 "nvptx_register_operand" "R"))
701 (match_operand:SI 2 "nvptx_register_operand" "R"))))]
703 "%.\\tmul.wide.u32\\t%0, %1, %2;")
705 (define_expand "mulditi3"
706 [(set (match_operand:TI 0 "nvptx_register_operand")
707 (mult:TI (sign_extend:TI
708 (match_operand:DI 1 "nvptx_register_operand"))
710 (match_operand:DI 2 "nvptx_nonmemory_operand"))))]
713 rtx hi = gen_reg_rtx (DImode);
714 rtx lo = gen_reg_rtx (DImode);
715 emit_insn (gen_smuldi3_highpart (hi, operands[1], operands[2]));
716 emit_insn (gen_muldi3 (lo, operands[1], operands[2]));
717 emit_move_insn (gen_highpart (DImode, operands[0]), hi);
718 emit_move_insn (gen_lowpart (DImode, operands[0]), lo);
722 (define_expand "umulditi3"
723 [(set (match_operand:TI 0 "nvptx_register_operand")
724 (mult:TI (zero_extend:TI
725 (match_operand:DI 1 "nvptx_register_operand"))
727 (match_operand:DI 2 "nvptx_nonmemory_operand"))))]
730 rtx hi = gen_reg_rtx (DImode);
731 rtx lo = gen_reg_rtx (DImode);
732 emit_insn (gen_umuldi3_highpart (hi, operands[1], operands[2]));
733 emit_insn (gen_muldi3 (lo, operands[1], operands[2]));
734 emit_move_insn (gen_highpart (DImode, operands[0]), hi);
735 emit_move_insn (gen_lowpart (DImode, operands[0]), lo);
739 (define_insn "smul<mode>3_highpart"
740 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
742 (match_operand:HSDIM 1 "nvptx_register_operand" "R")
743 (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
745 "%.\\tmul.hi.s%T0\\t%0, %1, %2;")
747 (define_insn "umul<mode>3_highpart"
748 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
750 (match_operand:HSDIM 1 "nvptx_register_operand" "R")
751 (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
753 "%.\\tmul.hi.u%T0\\t%0, %1, %2;")
755 (define_insn "*smulhi3_highpart_2"
756 [(set (match_operand:HI 0 "nvptx_register_operand" "=R")
759 (mult:SI (sign_extend:SI
760 (match_operand:HI 1 "nvptx_register_operand" "R"))
762 (match_operand:HI 2 "nvptx_register_operand" "R")))
765 "%.\\tmul.hi.s16\\t%0, %1, %2;")
767 (define_insn "*smulsi3_highpart_2"
768 [(set (match_operand:SI 0 "nvptx_register_operand" "=R")
771 (mult:DI (sign_extend:DI
772 (match_operand:SI 1 "nvptx_register_operand" "R"))
774 (match_operand:SI 2 "nvptx_register_operand" "R")))
777 "%.\\tmul.hi.s32\\t%0, %1, %2;")
779 (define_insn "*umulhi3_highpart_2"
780 [(set (match_operand:HI 0 "nvptx_register_operand" "=R")
783 (mult:SI (zero_extend:SI
784 (match_operand:HI 1 "nvptx_register_operand" "R"))
786 (match_operand:HI 2 "nvptx_register_operand" "R")))
789 "%.\\tmul.hi.u16\\t%0, %1, %2;")
791 (define_insn "*umulsi3_highpart_2"
792 [(set (match_operand:SI 0 "nvptx_register_operand" "=R")
795 (mult:DI (zero_extend:DI
796 (match_operand:SI 1 "nvptx_register_operand" "R"))
798 (match_operand:SI 2 "nvptx_register_operand" "R")))
801 "%.\\tmul.hi.u32\\t%0, %1, %2;")
805 (define_insn "ashl<mode>3"
806 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
807 (ashift:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")
808 (match_operand:SI 2 "nvptx_nonmemory_operand" "Ri")))]
810 "%.\\tshl.b%T0\\t%0, %1, %2;")
812 (define_insn "ashr<mode>3"
813 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
814 (ashiftrt:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")
815 (match_operand:SI 2 "nvptx_nonmemory_operand" "Ri")))]
817 "%.\\tshr.s%T0\\t%0, %1, %2;")
819 (define_insn "lshr<mode>3"
820 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
821 (lshiftrt:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")
822 (match_operand:SI 2 "nvptx_nonmemory_operand" "Ri")))]
824 "%.\\tshr.u%T0\\t%0, %1, %2;")
826 (define_insn "rotlsi3"
827 [(set (match_operand:SI 0 "nvptx_register_operand" "=R")
828 (rotate:SI (match_operand:SI 1 "nvptx_register_operand" "R")
829 (and:SI (match_operand:SI 2 "nvptx_nonmemory_operand" "Ri")
832 "%.\\tshf.l.wrap.b32\\t%0, %1, %1, %2;")
834 (define_insn "rotrsi3"
835 [(set (match_operand:SI 0 "nvptx_register_operand" "=R")
836 (rotatert:SI (match_operand:SI 1 "nvptx_register_operand" "R")
837 (and:SI (match_operand:SI 2 "nvptx_nonmemory_operand" "Ri")
840 "%.\\tshf.r.wrap.b32\\t%0, %1, %1, %2;")
842 ;; Logical operations
844 (define_code_iterator any_logic [and ior xor])
845 (define_code_attr logic [(and "and") (ior "or") (xor "xor")])
846 (define_code_attr ilogic [(and "and") (ior "ior") (xor "xor")])
848 (define_insn "<ilogic><mode>3"
849 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
851 (match_operand:HSDIM 1 "nvptx_register_operand" "R")
852 (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
854 "%.\\t<logic>.b%T0\\t%0, %1, %2;")
856 (define_insn "<ilogic>bi3"
857 [(set (match_operand:BI 0 "nvptx_register_operand" "=R")
858 (any_logic:BI (match_operand:BI 1 "nvptx_register_operand" "R")
859 (match_operand:BI 2 "nvptx_register_operand" "R")))]
861 "%.\\t<logic>.pred\\t%0, %1, %2;")
864 [(set (match_operand:HSDIM 0 "nvptx_register_operand")
866 (ne:HSDIM (match_operand:BI 1 "nvptx_register_operand")
868 (ne:HSDIM (match_operand:BI 2 "nvptx_register_operand")
870 "can_create_pseudo_p ()"
871 [(set (match_dup 3) (any_logic:BI (match_dup 1) (match_dup 2)))
872 (set (match_dup 0) (ne:HSDIM (match_dup 3) (const_int 0)))]
874 operands[3] = gen_reg_rtx (BImode);
877 ;; Comparisons and branches
879 (define_insn "cmp<mode>"
880 [(set (match_operand:BI 0 "nvptx_register_operand" "=R")
881 (match_operator:BI 1 "nvptx_comparison_operator"
882 [(match_operand:HSDIM 2 "nvptx_register_operand" "R")
883 (match_operand:HSDIM 3 "nvptx_nonmemory_operand" "Ri")]))]
885 "%.\\tsetp%c1\\t%0, %2, %3;")
887 (define_insn "*cmp<mode>"
888 [(set (match_operand:BI 0 "nvptx_register_operand" "=R")
889 (match_operator:BI 1 "nvptx_float_comparison_operator"
890 [(match_operand:SDFM 2 "nvptx_register_operand" "R")
891 (match_operand:SDFM 3 "nvptx_nonmemory_operand" "RF")]))]
893 "%.\\tsetp%c1\\t%0, %2, %3;")
895 (define_insn "*cmphf"
896 [(set (match_operand:BI 0 "nvptx_register_operand" "=R")
897 (match_operator:BI 1 "nvptx_float_comparison_operator"
898 [(match_operand:HF 2 "nvptx_register_operand" "R")
899 (match_operand:HF 3 "nvptx_nonmemory_operand" "RF")]))]
901 "%.\\tsetp%c1\\t%0, %2, %3;")
905 (label_ref (match_operand 0 "" "")))]
909 (define_insn "br_true"
911 (if_then_else (ne (match_operand:BI 0 "nvptx_register_operand" "R")
913 (label_ref (match_operand 1 "" ""))
917 [(set_attr "predicable" "no")])
919 (define_insn "br_false"
921 (if_then_else (eq (match_operand:BI 0 "nvptx_register_operand" "R")
923 (label_ref (match_operand 1 "" ""))
927 [(set_attr "predicable" "no")])
929 ;; unified conditional branch
930 (define_insn "br_true_uni"
931 [(set (pc) (if_then_else
932 (ne (unspec:BI [(match_operand:BI 0 "nvptx_register_operand" "R")]
933 UNSPEC_BR_UNIFIED) (const_int 0))
934 (label_ref (match_operand 1 "" "")) (pc)))]
936 "%j0\\tbra.uni\\t%l1;"
937 [(set_attr "predicable" "no")])
939 (define_insn "br_false_uni"
940 [(set (pc) (if_then_else
941 (eq (unspec:BI [(match_operand:BI 0 "nvptx_register_operand" "R")]
942 UNSPEC_BR_UNIFIED) (const_int 0))
943 (label_ref (match_operand 1 "" "")) (pc)))]
945 "%J0\\tbra.uni\\t%l1;"
946 [(set_attr "predicable" "no")])
948 (define_expand "cbranch<mode>4"
950 (if_then_else (match_operator 0 "nvptx_comparison_operator"
951 [(match_operand:HSDIM 1 "nvptx_register_operand" "")
952 (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "")])
953 (label_ref (match_operand 3 "" ""))
957 rtx t = nvptx_expand_compare (operands[0]);
959 operands[1] = XEXP (t, 0);
960 operands[2] = XEXP (t, 1);
963 (define_expand "cbranch<mode>4"
965 (if_then_else (match_operator 0 "nvptx_float_comparison_operator"
966 [(match_operand:SDFM 1 "nvptx_register_operand" "")
967 (match_operand:SDFM 2 "nvptx_nonmemory_operand" "")])
968 (label_ref (match_operand 3 "" ""))
972 rtx t = nvptx_expand_compare (operands[0]);
974 operands[1] = XEXP (t, 0);
975 operands[2] = XEXP (t, 1);
978 (define_expand "cbranchbi4"
980 (if_then_else (match_operator 0 "predicate_operator"
981 [(match_operand:BI 1 "nvptx_register_operand" "")
982 (match_operand:BI 2 "const0_operand" "")])
983 (label_ref (match_operand 3 "" ""))
988 ;; Conditional stores
990 (define_insn "setcc<mode>_from_bi"
991 [(set (match_operand:QHSDIM 0 "nvptx_register_operand" "=R")
992 (ne:QHSDIM (match_operand:BI 1 "nvptx_register_operand" "R")
995 "%.\\tselp%t0\\t%0, 1, 0, %1;")
997 (define_insn "*setcc<mode>_from_not_bi"
998 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
999 (eq:HSDIM (match_operand:BI 1 "nvptx_register_operand" "R")
1002 "%.\\tselp%t0\\t%0, 0, 1, %1;")
1004 (define_insn "extendbi<mode>2"
1005 [(set (match_operand:QHSDIM 0 "nvptx_register_operand" "=R")
1007 (match_operand:BI 1 "nvptx_register_operand" "R")))]
1009 "%.\\tselp%t0\\t%0, -1, 0, %1;")
1011 (define_insn "zero_extendbi<mode>2"
1012 [(set (match_operand:QHSDIM 0 "nvptx_register_operand" "=R")
1014 (match_operand:BI 1 "nvptx_register_operand" "R")))]
1016 "%.\\tselp%t0\\t%0, 1, 0, %1;")
1018 (define_insn "sel_true<mode>"
1019 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
1021 (ne (match_operand:BI 1 "nvptx_register_operand" "R") (const_int 0))
1022 (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")
1023 (match_operand:HSDIM 3 "nvptx_nonmemory_operand" "Ri")))]
1025 "%.\\tselp%t0\\t%0, %2, %3, %1;")
1027 (define_insn "sel_true<mode>"
1028 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
1030 (ne (match_operand:BI 1 "nvptx_register_operand" "R") (const_int 0))
1031 (match_operand:SDFM 2 "nvptx_nonmemory_operand" "RF")
1032 (match_operand:SDFM 3 "nvptx_nonmemory_operand" "RF")))]
1034 "%.\\tselp%t0\\t%0, %2, %3, %1;")
1036 (define_insn "sel_false<mode>"
1037 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
1039 (eq (match_operand:BI 1 "nvptx_register_operand" "R") (const_int 0))
1040 (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")
1041 (match_operand:HSDIM 3 "nvptx_nonmemory_operand" "Ri")))]
1043 "%.\\tselp%t0\\t%0, %3, %2, %1;")
1045 (define_insn "sel_false<mode>"
1046 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
1048 (eq (match_operand:BI 1 "nvptx_register_operand" "R") (const_int 0))
1049 (match_operand:SDFM 2 "nvptx_nonmemory_operand" "RF")
1050 (match_operand:SDFM 3 "nvptx_nonmemory_operand" "RF")))]
1052 "%.\\tselp%t0\\t%0, %3, %2, %1;")
1054 (define_code_iterator eqne [eq ne])
1056 ;; Split negation of a predicate into a conditional move.
1057 (define_insn_and_split "*selp<mode>_neg_<code>"
1058 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
1059 (neg:HSDIM (eqne:HSDIM
1060 (match_operand:BI 1 "nvptx_register_operand" "R")
1067 (eqne (match_dup 1) (const_int 0))
1071 ;; Split bitwise not of a predicate into a conditional move.
1072 (define_insn_and_split "*selp<mode>_not_<code>"
1073 [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
1074 (not:HSDIM (eqne:HSDIM
1075 (match_operand:BI 1 "nvptx_register_operand" "R")
1082 (eqne (match_dup 1) (const_int 0))
1086 (define_insn "*setcc_int<mode>"
1087 [(set (match_operand:SI 0 "nvptx_register_operand" "=R")
1089 (match_operator:SI 1 "nvptx_comparison_operator"
1090 [(match_operand:HSDIM 2 "nvptx_register_operand" "R")
1091 (match_operand:HSDIM 3 "nvptx_nonmemory_operand" "Ri")])))]
1093 "%.\\tset%t0%c1\\t%0, %2, %3;")
1095 (define_insn "*setcc_int<mode>"
1096 [(set (match_operand:SI 0 "nvptx_register_operand" "=R")
1098 (match_operator:SI 1 "nvptx_float_comparison_operator"
1099 [(match_operand:SDFM 2 "nvptx_register_operand" "R")
1100 (match_operand:SDFM 3 "nvptx_nonmemory_operand" "RF")])))]
1102 "%.\\tset%t0%c1\\t%0, %2, %3;")
1104 (define_insn "setcc_float<mode>"
1105 [(set (match_operand:SF 0 "nvptx_register_operand" "=R")
1106 (match_operator:SF 1 "nvptx_comparison_operator"
1107 [(match_operand:HSDIM 2 "nvptx_register_operand" "R")
1108 (match_operand:HSDIM 3 "nvptx_nonmemory_operand" "Ri")]))]
1110 "%.\\tset%t0%c1\\t%0, %2, %3;")
1112 (define_insn "setcc_float<mode>"
1113 [(set (match_operand:SF 0 "nvptx_register_operand" "=R")
1114 (match_operator:SF 1 "nvptx_float_comparison_operator"
1115 [(match_operand:SDFM 2 "nvptx_register_operand" "R")
1116 (match_operand:SDFM 3 "nvptx_nonmemory_operand" "RF")]))]
1118 "%.\\tset%t0%c1\\t%0, %2, %3;")
1120 (define_expand "cstore<mode>4"
1121 [(set (match_operand:SI 0 "nvptx_register_operand")
1122 (match_operator:SI 1 "nvptx_comparison_operator"
1123 [(match_operand:HSDIM 2 "nvptx_register_operand")
1124 (match_operand:HSDIM 3 "nvptx_nonmemory_operand")]))]
1127 rtx reg = gen_reg_rtx (BImode);
1128 rtx cmp = gen_rtx_fmt_ee (GET_CODE (operands[1]), BImode,
1129 operands[2], operands[3]);
1130 emit_move_insn (reg, cmp);
1131 emit_insn (gen_setccsi_from_bi (operands[0], reg));
1135 (define_expand "cstore<mode>4"
1136 [(set (match_operand:SI 0 "nvptx_register_operand")
1137 (match_operator:SI 1 "nvptx_float_comparison_operator"
1138 [(match_operand:SDFM 2 "nvptx_register_operand")
1139 (match_operand:SDFM 3 "nvptx_nonmemory_operand")]))]
1142 rtx reg = gen_reg_rtx (BImode);
1143 rtx cmp = gen_rtx_fmt_ee (GET_CODE (operands[1]), BImode,
1144 operands[2], operands[3]);
1145 emit_move_insn (reg, cmp);
1146 emit_insn (gen_setccsi_from_bi (operands[0], reg));
1150 (define_expand "cstorehf4"
1151 [(set (match_operand:SI 0 "nvptx_register_operand")
1152 (match_operator:SI 1 "nvptx_float_comparison_operator"
1153 [(match_operand:HF 2 "nvptx_register_operand")
1154 (match_operand:HF 3 "nvptx_nonmemory_operand")]))]
1157 rtx reg = gen_reg_rtx (BImode);
1158 rtx cmp = gen_rtx_fmt_ee (GET_CODE (operands[1]), BImode,
1159 operands[2], operands[3]);
1160 emit_move_insn (reg, cmp);
1161 emit_insn (gen_setccsi_from_bi (operands[0], reg));
1167 (define_insn "call_insn_<mode>"
1168 [(match_parallel 2 "call_operation"
1169 [(call (mem:QI (match_operand:P 0 "call_insn_operand" "Rs"))
1170 (match_operand 1))])]
1173 return nvptx_output_call_insn (insn, NULL_RTX, operands[0]);
1176 (define_insn "call_value_insn_<mode>"
1177 [(match_parallel 3 "call_operation"
1178 [(set (match_operand 0 "nvptx_register_operand" "=R")
1179 (call (mem:QI (match_operand:P 1 "call_insn_operand" "Rs"))
1180 (match_operand 2)))])]
1183 return nvptx_output_call_insn (insn, operands[0], operands[1]);
1186 (define_expand "call"
1187 [(match_operand 0 "" "")]
1190 nvptx_expand_call (NULL_RTX, operands[0]);
1194 (define_expand "call_value"
1195 [(match_operand 0 "" "")
1196 (match_operand 1 "" "")]
1199 nvptx_expand_call (operands[0], operands[1]);
1203 ;; Floating point arithmetic.
1205 (define_insn "add<mode>3"
1206 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
1207 (plus:SDFM (match_operand:SDFM 1 "nvptx_register_operand" "R")
1208 (match_operand:SDFM 2 "nvptx_nonmemory_operand" "RF")))]
1210 "%.\\tadd%t0\\t%0, %1, %2;")
1212 (define_insn "sub<mode>3"
1213 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
1214 (minus:SDFM (match_operand:SDFM 1 "nvptx_register_operand" "R")
1215 (match_operand:SDFM 2 "nvptx_register_operand" "R")))]
1217 "%.\\tsub%t0\\t%0, %1, %2;")
1219 (define_insn "mul<mode>3"
1220 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
1221 (mult:SDFM (match_operand:SDFM 1 "nvptx_register_operand" "R")
1222 (match_operand:SDFM 2 "nvptx_nonmemory_operand" "RF")))]
1224 "%.\\tmul%t0\\t%0, %1, %2;")
1226 (define_insn "fma<mode>4"
1227 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
1228 (fma:SDFM (match_operand:SDFM 1 "nvptx_register_operand" "R")
1229 (match_operand:SDFM 2 "nvptx_nonmemory_operand" "RF")
1230 (match_operand:SDFM 3 "nvptx_nonmemory_operand" "RF")))]
1232 "%.\\tfma%#%t0\\t%0, %1, %2, %3;")
1234 (define_insn "*recip<mode>2"
1235 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
1237 (match_operand:SDFM 2 "const_double_operand" "F")
1238 (match_operand:SDFM 1 "nvptx_register_operand" "R")))]
1239 "CONST_DOUBLE_P (operands[2])
1240 && real_identical (CONST_DOUBLE_REAL_VALUE (operands[2]), &dconst1)"
1241 "%.\\trcp%#%t0\\t%0, %1;")
1243 (define_insn "div<mode>3"
1244 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
1245 (div:SDFM (match_operand:SDFM 1 "nvptx_register_operand" "R")
1246 (match_operand:SDFM 2 "nvptx_nonmemory_operand" "RF")))]
1248 "%.\\tdiv%#%t0\\t%0, %1, %2;")
1250 (define_insn "copysign<mode>3"
1251 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
1252 (unspec:SDFM [(match_operand:SDFM 1 "nvptx_nonmemory_operand" "RF")
1253 (match_operand:SDFM 2 "nvptx_nonmemory_operand" "RF")]
1256 "%.\\tcopysign%t0\\t%0, %2, %1;")
1258 (define_insn "smin<mode>3"
1259 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
1260 (smin:SDFM (match_operand:SDFM 1 "nvptx_register_operand" "R")
1261 (match_operand:SDFM 2 "nvptx_nonmemory_operand" "RF")))]
1263 "%.\\tmin%t0\\t%0, %1, %2;")
1265 (define_insn "smax<mode>3"
1266 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
1267 (smax:SDFM (match_operand:SDFM 1 "nvptx_register_operand" "R")
1268 (match_operand:SDFM 2 "nvptx_nonmemory_operand" "RF")))]
1270 "%.\\tmax%t0\\t%0, %1, %2;")
1272 (define_insn "abs<mode>2"
1273 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
1274 (abs:SDFM (match_operand:SDFM 1 "nvptx_register_operand" "R")))]
1276 "%.\\tabs%t0\\t%0, %1;")
1278 (define_insn "neg<mode>2"
1279 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
1280 (neg:SDFM (match_operand:SDFM 1 "nvptx_register_operand" "R")))]
1282 "%.\\tneg%t0\\t%0, %1;")
1284 (define_insn "sqrt<mode>2"
1285 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
1286 (sqrt:SDFM (match_operand:SDFM 1 "nvptx_register_operand" "R")))]
1288 "%.\\tsqrt%#%t0\\t%0, %1;")
1290 (define_expand "sincossf3"
1291 [(set (match_operand:SF 0 "nvptx_register_operand" "=R")
1292 (unspec:SF [(match_operand:SF 2 "nvptx_register_operand" "R")]
1294 (set (match_operand:SF 1 "nvptx_register_operand" "=R")
1295 (unspec:SF [(match_dup 2)] UNSPEC_SIN))]
1296 "flag_unsafe_math_optimizations"
1298 operands[2] = make_safe_from (operands[2], operands[0]);
1301 (define_insn "sinsf2"
1302 [(set (match_operand:SF 0 "nvptx_register_operand" "=R")
1303 (unspec:SF [(match_operand:SF 1 "nvptx_register_operand" "R")]
1305 "flag_unsafe_math_optimizations"
1306 "%.\\tsin.approx%t0\\t%0, %1;")
1308 (define_insn "cossf2"
1309 [(set (match_operand:SF 0 "nvptx_register_operand" "=R")
1310 (unspec:SF [(match_operand:SF 1 "nvptx_register_operand" "R")]
1312 "flag_unsafe_math_optimizations"
1313 "%.\\tcos.approx%t0\\t%0, %1;")
1315 (define_insn "log2sf2"
1316 [(set (match_operand:SF 0 "nvptx_register_operand" "=R")
1317 (unspec:SF [(match_operand:SF 1 "nvptx_register_operand" "R")]
1319 "flag_unsafe_math_optimizations"
1320 "%.\\tlg2.approx%t0\\t%0, %1;")
1322 (define_insn "exp2sf2"
1323 [(set (match_operand:SF 0 "nvptx_register_operand" "=R")
1324 (unspec:SF [(match_operand:SF 1 "nvptx_register_operand" "R")]
1326 "flag_unsafe_math_optimizations"
1327 "%.\\tex2.approx%t0\\t%0, %1;")
1329 (define_insn "setcc_isinf<mode>"
1330 [(set (match_operand:BI 0 "nvptx_register_operand" "=R")
1331 (unspec:BI [(match_operand:SDFM 1 "nvptx_register_operand" "R")]
1334 "%.\\ttestp.infinite%t1\\t%0, %1;")
1336 (define_expand "isinf<mode>2"
1337 [(set (match_operand:SI 0 "nvptx_register_operand" "=R")
1338 (unspec:SI [(match_operand:SDFM 1 "nvptx_register_operand" "R")]
1342 rtx pred = gen_reg_rtx (BImode);
1343 emit_insn (gen_setcc_isinf<mode> (pred, operands[1]));
1344 emit_insn (gen_setccsi_from_bi (operands[0], pred));
1348 ;; HFmode floating point arithmetic.
1350 (define_insn "addhf3"
1351 [(set (match_operand:HF 0 "nvptx_register_operand" "=R")
1352 (plus:HF (match_operand:HF 1 "nvptx_register_operand" "R")
1353 (match_operand:HF 2 "nvptx_register_operand" "R")))]
1355 "%.\\tadd.f16\\t%0, %1, %2;")
1357 (define_insn "subhf3"
1358 [(set (match_operand:HF 0 "nvptx_register_operand" "=R")
1359 (minus:HF (match_operand:HF 1 "nvptx_register_operand" "R")
1360 (match_operand:HF 2 "nvptx_register_operand" "R")))]
1362 "%.\\tsub.f16\\t%0, %1, %2;")
1364 (define_insn "mulhf3"
1365 [(set (match_operand:HF 0 "nvptx_register_operand" "=R")
1366 (mult:HF (match_operand:HF 1 "nvptx_register_operand" "R")
1367 (match_operand:HF 2 "nvptx_register_operand" "R")))]
1369 "%.\\tmul.f16\\t%0, %1, %2;")
1371 (define_insn "fmahf4"
1372 [(set (match_operand:HF 0 "nvptx_register_operand" "=R")
1373 (fma:HF (match_operand:HF 1 "nvptx_register_operand" "R")
1374 (match_operand:HF 2 "nvptx_nonmemory_operand" "RF")
1375 (match_operand:HF 3 "nvptx_nonmemory_operand" "RF")))]
1377 "%.\\tfma%#.f16\\t%0, %1, %2, %3;")
1379 (define_insn "neghf2"
1380 [(set (match_operand:HF 0 "nvptx_register_operand" "=R")
1381 (neg:HF (match_operand:HF 1 "nvptx_register_operand" "R")))]
1383 "%.\\txor.b16\\t%0, %1, -32768;")
1385 (define_insn "abshf2"
1386 [(set (match_operand:HF 0 "nvptx_register_operand" "=R")
1387 (abs:HF (match_operand:HF 1 "nvptx_register_operand" "R")))]
1389 "%.\\tand.b16\\t%0, %1, 32767;")
1391 (define_insn "exp2hf2"
1392 [(set (match_operand:HF 0 "nvptx_register_operand" "=R")
1393 (unspec:HF [(match_operand:HF 1 "nvptx_register_operand" "R")]
1395 "TARGET_SM75 && flag_unsafe_math_optimizations"
1396 "%.\\tex2.approx.f16\\t%0, %1;")
1398 (define_insn "tanh<mode>2"
1399 [(set (match_operand:HSFM 0 "nvptx_register_operand" "=R")
1400 (unspec:HSFM [(match_operand:HSFM 1 "nvptx_register_operand" "R")]
1402 "TARGET_SM75 && flag_unsafe_math_optimizations"
1403 "%.\\ttanh.approx%t0\\t%0, %1;")
1405 ;; HFmode floating point arithmetic.
1407 (define_insn "sminhf3"
1408 [(set (match_operand:HF 0 "nvptx_register_operand" "=R")
1409 (smin:HF (match_operand:HF 1 "nvptx_register_operand" "R")
1410 (match_operand:HF 2 "nvptx_register_operand" "R")))]
1412 "%.\\tmin.f16\\t%0, %1, %2;")
1414 (define_insn "smaxhf3"
1415 [(set (match_operand:HF 0 "nvptx_register_operand" "=R")
1416 (smax:HF (match_operand:HF 1 "nvptx_register_operand" "R")
1417 (match_operand:HF 2 "nvptx_register_operand" "R")))]
1419 "%.\\tmax.f16\\t%0, %1, %2;")
1421 ;; Conversions involving floating point
1423 (define_insn "extendsfdf2"
1424 [(set (match_operand:DF 0 "nvptx_register_operand" "=R")
1425 (float_extend:DF (match_operand:SF 1 "nvptx_register_operand" "R")))]
1427 "%.\\tcvt%t0%t1\\t%0, %1;")
1429 (define_insn "truncdfsf2"
1430 [(set (match_operand:SF 0 "nvptx_register_operand" "=R")
1431 (float_truncate:SF (match_operand:DF 1 "nvptx_register_operand" "R")))]
1433 "%.\\tcvt%#%t0%t1\\t%0, %1;")
1435 (define_insn "floatunssi<mode>2"
1436 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
1437 (unsigned_float:SDFM (match_operand:SI 1 "nvptx_register_operand" "R")))]
1439 "%.\\tcvt%#%t0.u%T1\\t%0, %1;")
1441 (define_insn "floatsi<mode>2"
1442 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
1443 (float:SDFM (match_operand:SI 1 "nvptx_register_operand" "R")))]
1445 "%.\\tcvt%#%t0.s%T1\\t%0, %1;")
1447 (define_insn "floatunsdi<mode>2"
1448 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
1449 (unsigned_float:SDFM (match_operand:DI 1 "nvptx_register_operand" "R")))]
1451 "%.\\tcvt%#%t0.u%T1\\t%0, %1;")
1453 (define_insn "floatdi<mode>2"
1454 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
1455 (float:SDFM (match_operand:DI 1 "nvptx_register_operand" "R")))]
1457 "%.\\tcvt%#%t0.s%T1\\t%0, %1;")
1459 (define_insn "fixuns_trunc<mode>si2"
1460 [(set (match_operand:SI 0 "nvptx_register_operand" "=R")
1461 (unsigned_fix:SI (match_operand:SDFM 1 "nvptx_register_operand" "R")))]
1463 "%.\\tcvt.rzi.u%T0%t1\\t%0, %1;")
1465 (define_insn "fix_trunc<mode>si2"
1466 [(set (match_operand:SI 0 "nvptx_register_operand" "=R")
1467 (fix:SI (match_operand:SDFM 1 "nvptx_register_operand" "R")))]
1469 "%.\\tcvt.rzi.s%T0%t1\\t%0, %1;")
1471 (define_insn "fixuns_trunc<mode>di2"
1472 [(set (match_operand:DI 0 "nvptx_register_operand" "=R")
1473 (unsigned_fix:DI (match_operand:SDFM 1 "nvptx_register_operand" "R")))]
1475 "%.\\tcvt.rzi.u%T0%t1\\t%0, %1;")
1477 (define_insn "fix_trunc<mode>di2"
1478 [(set (match_operand:DI 0 "nvptx_register_operand" "=R")
1479 (fix:DI (match_operand:SDFM 1 "nvptx_register_operand" "R")))]
1481 "%.\\tcvt.rzi.s%T0%t1\\t%0, %1;")
1483 (define_int_iterator FPINT [UNSPEC_FPINT_FLOOR UNSPEC_FPINT_BTRUNC
1484 UNSPEC_FPINT_CEIL UNSPEC_FPINT_NEARBYINT])
1485 (define_int_attr fpint_name [(UNSPEC_FPINT_FLOOR "floor")
1486 (UNSPEC_FPINT_BTRUNC "btrunc")
1487 (UNSPEC_FPINT_CEIL "ceil")
1488 (UNSPEC_FPINT_NEARBYINT "nearbyint")])
1489 (define_int_attr fpint_roundingmode [(UNSPEC_FPINT_FLOOR ".rmi")
1490 (UNSPEC_FPINT_BTRUNC ".rzi")
1491 (UNSPEC_FPINT_CEIL ".rpi")
1492 (UNSPEC_FPINT_NEARBYINT "%#i")])
1494 (define_insn "<FPINT:fpint_name><SDFM:mode>2"
1495 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
1496 (unspec:SDFM [(match_operand:SDFM 1 "nvptx_register_operand" "R")]
1499 "%.\\tcvt<FPINT:fpint_roundingmode>%t0%t1\\t%0, %1;")
1501 (define_int_iterator FPINT2 [UNSPEC_FPINT_FLOOR UNSPEC_FPINT_CEIL])
1502 (define_int_attr fpint2_name [(UNSPEC_FPINT_FLOOR "lfloor")
1503 (UNSPEC_FPINT_CEIL "lceil")])
1504 (define_int_attr fpint2_roundingmode [(UNSPEC_FPINT_FLOOR ".rmi")
1505 (UNSPEC_FPINT_CEIL ".rpi")])
1507 (define_insn "<FPINT2:fpint2_name><SDFM:mode><SDIM:mode>2"
1508 [(set (match_operand:SDIM 0 "nvptx_register_operand" "=R")
1509 (unspec:SDIM [(match_operand:SDFM 1 "nvptx_register_operand" "R")]
1512 "%.\\tcvt<FPINT2:fpint2_roundingmode>.s%T0%t1\\t%0, %1;")
1514 (define_insn "extendhf<mode>2"
1515 [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
1516 (float_extend:SDFM (match_operand:HF 1 "nvptx_register_operand" "R")))]
1518 "%.\\tcvt%t0%t1\\t%0, %1;")
1520 (define_insn "trunc<mode>hf2"
1521 [(set (match_operand:HF 0 "nvptx_register_operand" "=R")
1522 (float_truncate:HF (match_operand:SDFM 1 "nvptx_register_operand" "R")))]
1524 "%.\\tcvt%#%t0%t1\\t%0, %1;")
1526 ;; Vector operations
1528 (define_insn "*vec_set<mode>_0"
1529 [(set (match_operand:VECIM 0 "nvptx_register_operand" "=R")
1531 (vec_duplicate:VECIM
1532 (match_operand:<VECELEM> 1 "nvptx_register_operand" "R"))
1536 "%.\\tmov%t1\\t%0.x, %1;")
1538 (define_insn "*vec_set<mode>_1"
1539 [(set (match_operand:VECIM 0 "nvptx_register_operand" "=R")
1541 (vec_duplicate:VECIM
1542 (match_operand:<VECELEM> 1 "nvptx_register_operand" "R"))
1546 "%.\\tmov%t1\\t%0.y, %1;")
1548 (define_insn "*vec_set<mode>_2"
1549 [(set (match_operand:VECIM 0 "nvptx_register_operand" "=R")
1551 (vec_duplicate:VECIM
1552 (match_operand:<VECELEM> 1 "nvptx_register_operand" "R"))
1556 "%.\\tmov%t1\\t%0.z, %1;")
1558 (define_insn "*vec_set<mode>_3"
1559 [(set (match_operand:VECIM 0 "nvptx_register_operand" "=R")
1561 (vec_duplicate:VECIM
1562 (match_operand:<VECELEM> 1 "nvptx_register_operand" "R"))
1566 "%.\\tmov%t1\\t%0.w, %1;")
1568 (define_expand "vec_set<mode>"
1569 [(match_operand:VECIM 0 "nvptx_register_operand")
1570 (match_operand:<VECELEM> 1 "nvptx_register_operand")
1571 (match_operand:SI 2 "nvptx_vector_index_operand")]
1574 enum machine_mode mode = GET_MODE (operands[0]);
1575 int mask = 1 << INTVAL (operands[2]);
1576 rtx tmp = gen_rtx_VEC_DUPLICATE (mode, operands[1]);
1577 tmp = gen_rtx_VEC_MERGE (mode, tmp, operands[0], GEN_INT (mask));
1578 emit_insn (gen_rtx_SET (operands[0], tmp));
1582 (define_insn "vec_extract<mode><Vecelem>"
1583 [(set (match_operand:<VECELEM> 0 "nvptx_register_operand" "=R")
1584 (vec_select:<VECELEM>
1585 (match_operand:VECIM 1 "nvptx_register_operand" "R")
1586 (parallel [(match_operand:SI 2 "nvptx_vector_index_operand" "")])))]
1589 static const char *const asms[4] = {
1590 "%.\\tmov%t0\\t%0, %1.x;",
1591 "%.\\tmov%t0\\t%0, %1.y;",
1592 "%.\\tmov%t0\\t%0, %1.z;",
1593 "%.\\tmov%t0\\t%0, %1.w;"
1595 return asms[INTVAL (operands[2])];
1610 (define_insn "fake_nop"
1614 .reg .u32 %%nop_src;
1615 .reg .u32 %%nop_dst;
1616 mov.u32 %%nop_dst, %%nop_src;
1619 (define_insn "return"
1623 return nvptx_output_return ();
1625 [(set_attr "predicable" "no")])
1627 (define_expand "epilogue"
1628 [(clobber (const_int 0))]
1631 if (TARGET_SOFT_STACK)
1632 emit_insn (gen_set_softstack (Pmode, gen_rtx_REG (Pmode,
1633 SOFTSTACK_PREV_REGNUM)));
1634 emit_jump_insn (gen_return ());
1638 (define_expand "nonlocal_goto"
1639 [(match_operand 0 "" "")
1640 (match_operand 1 "" "")
1641 (match_operand 2 "" "")
1642 (match_operand 3 "" "")]
1645 sorry ("target cannot support nonlocal goto");
1646 emit_insn (gen_nop ());
1650 (define_expand "nonlocal_goto_receiver"
1654 sorry ("target cannot support nonlocal goto");
1657 (define_expand "allocate_stack"
1658 [(match_operand 0 "nvptx_register_operand")
1659 (match_operand 1 "nvptx_register_operand")]
1662 if (TARGET_SOFT_STACK)
1664 emit_move_insn (stack_pointer_rtx,
1665 gen_rtx_MINUS (Pmode, stack_pointer_rtx, operands[1]));
1666 emit_insn (gen_set_softstack (Pmode, stack_pointer_rtx));
1667 emit_move_insn (operands[0], virtual_stack_dynamic_rtx);
1670 /* The ptx documentation specifies an alloca intrinsic (for 32 bit
1671 only) but notes it is not implemented. The assembler emits a
1672 confused error message. Issue a blunt one now instead. */
1673 sorry ("target cannot support alloca");
1674 emit_insn (gen_nop ());
1678 (define_insn "@set_softstack_<mode>"
1679 [(unspec [(match_operand:P 0 "nvptx_register_operand" "R")]
1680 UNSPEC_SET_SOFTSTACK)]
1683 return nvptx_output_set_softstack (REGNO (operands[0]));
1686 (define_expand "restore_stack_block"
1687 [(match_operand 0 "register_operand" "")
1688 (match_operand 1 "register_operand" "")]
1691 if (TARGET_SOFT_STACK)
1693 emit_move_insn (operands[0], operands[1]);
1694 emit_insn (gen_set_softstack (Pmode, operands[0]));
1699 (define_expand "restore_stack_function"
1700 [(match_operand 0 "register_operand" "")
1701 (match_operand 1 "register_operand" "")]
1708 [(trap_if (const_int 1) (const_int 0))]
1712 (define_insn "trap_if_true"
1713 [(trap_if (ne (match_operand:BI 0 "nvptx_register_operand" "R")
1717 "%j0 trap; %j0 exit;"
1718 [(set_attr "predicable" "no")])
1720 (define_insn "trap_if_false"
1721 [(trap_if (eq (match_operand:BI 0 "nvptx_register_operand" "R")
1725 "%J0 trap; %J0 exit;"
1726 [(set_attr "predicable" "no")])
1728 (define_expand "ctrap<mode>4"
1729 [(trap_if (match_operator 0 "nvptx_comparison_operator"
1730 [(match_operand:SDIM 1 "nvptx_register_operand")
1731 (match_operand:SDIM 2 "nvptx_nonmemory_operand")])
1732 (match_operand 3 "const0_operand"))]
1735 rtx t = nvptx_expand_compare (operands[0]);
1736 emit_insn (gen_trap_if_true (t));
1740 (define_insn "oacc_dim_size"
1741 [(set (match_operand:SI 0 "nvptx_register_operand" "")
1742 (unspec:SI [(match_operand:SI 1 "const_int_operand" "")]
1746 static const char *const asms[] =
1747 { /* Must match oacc_loop_levels ordering. */
1748 "%.\\tmov.u32\\t%0, %%nctaid.x;", /* gang */
1749 "%.\\tmov.u32\\t%0, %%ntid.y;", /* worker */
1750 "%.\\tmov.u32\\t%0, %%ntid.x;", /* vector */
1752 return asms[INTVAL (operands[1])];
1755 (define_insn "oacc_dim_pos"
1756 [(set (match_operand:SI 0 "nvptx_register_operand" "")
1757 (unspec_volatile:SI [(match_operand:SI 1 "const_int_operand" "")]
1761 static const char *const asms[] =
1762 { /* Must match oacc_loop_levels ordering. */
1763 "%.\\tmov.u32\\t%0, %%ctaid.x;", /* gang */
1764 "%.\\tmov.u32\\t%0, %%tid.y;", /* worker */
1765 "%.\\tmov.u32\\t%0, %%tid.x;", /* vector */
1767 return asms[INTVAL (operands[1])];
1770 (define_insn "nvptx_fork"
1771 [(unspec_volatile:SI [(match_operand:SI 0 "const_int_operand" "")]
1775 [(set_attr "predicable" "no")])
1777 (define_insn "nvptx_forked"
1778 [(unspec_volatile:SI [(match_operand:SI 0 "const_int_operand" "")]
1782 [(set_attr "predicable" "no")])
1784 (define_insn "nvptx_joining"
1785 [(unspec_volatile:SI [(match_operand:SI 0 "const_int_operand" "")]
1789 [(set_attr "predicable" "no")])
1791 (define_insn "nvptx_join"
1792 [(unspec_volatile:SI [(match_operand:SI 0 "const_int_operand" "")]
1796 [(set_attr "predicable" "no")])
1798 (define_expand "oacc_fork"
1799 [(set (match_operand:SI 0 "nvptx_nonmemory_operand" "")
1800 (match_operand:SI 1 "general_operand" ""))
1801 (unspec_volatile:SI [(match_operand:SI 2 "const_int_operand" "")]
1805 if (operands[0] != const0_rtx)
1806 emit_move_insn (operands[0], operands[1]);
1807 nvptx_expand_oacc_fork (INTVAL (operands[2]));
1811 (define_expand "oacc_join"
1812 [(set (match_operand:SI 0 "nvptx_nonmemory_operand" "")
1813 (match_operand:SI 1 "general_operand" ""))
1814 (unspec_volatile:SI [(match_operand:SI 2 "const_int_operand" "")]
1818 if (operands[0] != const0_rtx)
1819 emit_move_insn (operands[0], operands[1]);
1820 nvptx_expand_oacc_join (INTVAL (operands[2]));
1824 ;; only 32-bit shuffles exist.
1825 (define_insn "nvptx_shuffle<mode>"
1826 [(set (match_operand:BITS 0 "nvptx_register_operand" "=R")
1828 [(match_operand:BITS 1 "nvptx_register_operand" "R")
1829 (match_operand:SI 2 "nvptx_nonmemory_operand" "Ri")
1830 (match_operand:SI 3 "const_int_operand" "n")]
1835 return "%.\\tshfl.sync%S3.b32\\t%0, %1, %2, 31, 0xffffffff;";
1837 return "%.\\tshfl%S3.b32\\t%0, %1, %2, 31;";
1840 (define_insn "nvptx_vote_ballot"
1841 [(set (match_operand:SI 0 "nvptx_register_operand" "=R")
1842 (unspec:SI [(match_operand:BI 1 "nvptx_register_operand" "R")]
1843 UNSPEC_VOTE_BALLOT))]
1847 return "%.\\tvote.sync.ballot.b32\\t%0, %1, 0xffffffff;";
1849 return "%.\\tvote.ballot.b32\\t%0, %1;";
1852 ;; Patterns for OpenMP SIMD-via-SIMT lowering
1854 (define_insn "@omp_simt_enter_<mode>"
1855 [(set (match_operand:P 0 "nvptx_register_operand" "=R")
1856 (unspec_volatile:P [(match_operand:P 1 "nvptx_nonmemory_operand" "Ri")
1857 (match_operand:P 2 "nvptx_nonmemory_operand" "Ri")]
1858 UNSPECV_SIMT_ENTER))]
1861 return nvptx_output_simt_enter (operands[0], operands[1], operands[2]);
1864 (define_expand "omp_simt_enter"
1865 [(match_operand 0 "nvptx_register_operand" "=R")
1866 (match_operand 1 "nvptx_nonmemory_operand" "Ri")
1867 (match_operand 2 "const_int_operand" "n")]
1870 if (!CONST_INT_P (operands[1]))
1871 cfun->machine->simt_stack_size = HOST_WIDE_INT_M1U;
1873 cfun->machine->simt_stack_size = MAX (UINTVAL (operands[1]),
1874 cfun->machine->simt_stack_size);
1875 cfun->machine->simt_stack_align = MAX (UINTVAL (operands[2]),
1876 cfun->machine->simt_stack_align);
1877 cfun->machine->has_simtreg = true;
1878 emit_insn (gen_omp_simt_enter (Pmode, operands[0], operands[1], operands[2]));
1882 (define_expand "omp_simt_exit"
1883 [(match_operand 0 "nvptx_register_operand" "R")]
1886 emit_insn (gen_omp_simt_exit (Pmode, operands[0]));
1888 emit_insn (gen_nvptx_warpsync ());
1890 emit_insn (gen_nvptx_uniform_warp_check ());
1894 (define_insn "@omp_simt_exit_<mode>"
1895 [(unspec_volatile [(match_operand:P 0 "nvptx_register_operand" "R")]
1899 return nvptx_output_simt_exit (operands[0]);
1902 ;; Implement IFN_GOMP_SIMT_LANE: set operand 0 to lane index
1903 (define_insn "omp_simt_lane"
1904 [(set (match_operand:SI 0 "nvptx_register_operand" "")
1905 (unspec:SI [(const_int 0)] UNSPEC_LANEID))]
1907 "%.\\tmov.u32\\t%0, %%laneid;")
1909 ;; Implement IFN_GOMP_SIMT_ORDERED: copy operand 1 to operand 0 and
1910 ;; place a compiler barrier to disallow unrolling/peeling the containing loop
1911 (define_expand "omp_simt_ordered"
1912 [(match_operand:SI 0 "nvptx_register_operand" "=R")
1913 (match_operand:SI 1 "nvptx_register_operand" "R")]
1916 emit_move_insn (operands[0], operands[1]);
1917 emit_insn (gen_nvptx_nounroll ());
1921 ;; Implement IFN_GOMP_SIMT_XCHG_BFLY: perform a "butterfly" exchange
1923 (define_expand "omp_simt_xchg_bfly"
1924 [(match_operand 0 "nvptx_register_or_complex_di_df_register_operand" "=R")
1925 (match_operand 1 "nvptx_register_or_complex_di_df_register_operand" "R")
1926 (match_operand:SI 2 "nvptx_nonmemory_operand" "Ri")]
1929 emit_insn (nvptx_gen_shuffle (operands[0], operands[1], operands[2],
1934 ;; Implement IFN_GOMP_SIMT_XCHG_IDX: broadcast value in operand 1
1935 ;; from lane given by index in operand 2 to operand 0 in all lanes
1936 (define_expand "omp_simt_xchg_idx"
1937 [(match_operand 0 "nvptx_register_or_complex_di_df_register_operand" "=R")
1938 (match_operand 1 "nvptx_register_or_complex_di_df_register_operand" "R")
1939 (match_operand:SI 2 "nvptx_nonmemory_operand" "Ri")]
1942 emit_insn (nvptx_gen_shuffle (operands[0], operands[1], operands[2],
1947 ;; Implement IFN_GOMP_SIMT_VOTE_ANY:
1948 ;; set operand 0 to zero iff all lanes supply zero in operand 1
1949 (define_expand "omp_simt_vote_any"
1950 [(match_operand:SI 0 "nvptx_register_operand" "=R")
1951 (match_operand:SI 1 "nvptx_register_operand" "R")]
1954 rtx pred = gen_reg_rtx (BImode);
1955 emit_move_insn (pred, gen_rtx_NE (BImode, operands[1], const0_rtx));
1956 emit_insn (gen_nvptx_vote_ballot (operands[0], pred));
1960 ;; Implement IFN_GOMP_SIMT_LAST_LANE:
1961 ;; set operand 0 to the lowest lane index that passed non-zero in operand 1
1962 (define_expand "omp_simt_last_lane"
1963 [(match_operand:SI 0 "nvptx_register_operand" "=R")
1964 (match_operand:SI 1 "nvptx_register_operand" "R")]
1967 rtx pred = gen_reg_rtx (BImode);
1968 rtx tmp = gen_reg_rtx (SImode);
1969 emit_move_insn (pred, gen_rtx_NE (BImode, operands[1], const0_rtx));
1970 emit_insn (gen_nvptx_vote_ballot (tmp, pred));
1971 emit_insn (gen_ctzsi2 (operands[0], tmp));
1975 ;; extract parts of a 64 bit object into 2 32-bit ints
1976 (define_insn "unpack<mode>si2"
1977 [(set (match_operand:SI 0 "nvptx_register_operand" "=R")
1978 (unspec:SI [(match_operand:BITD 2 "nvptx_register_operand" "R")
1979 (const_int 0)] UNSPEC_BIT_CONV))
1980 (set (match_operand:SI 1 "nvptx_register_operand" "=R")
1981 (unspec:SI [(match_dup 2) (const_int 1)] UNSPEC_BIT_CONV))]
1983 "%.\\tmov.b64\\t{%0,%1}, %2;")
1985 ;; pack 2 32-bit ints into a 64 bit object
1986 (define_insn "packsi<mode>2"
1987 [(set (match_operand:BITD 0 "nvptx_register_operand" "=R")
1988 (unspec:BITD [(match_operand:SI 1 "nvptx_register_operand" "R")
1989 (match_operand:SI 2 "nvptx_register_operand" "R")]
1992 "%.\\tmov.b64\\t%0, {%1,%2};")
1996 (define_expand "atomic_compare_and_swap<mode>"
1997 [(match_operand:SI 0 "nvptx_register_operand") ;; bool success output
1998 (match_operand:SDIM 1 "nvptx_register_operand") ;; oldval output
1999 (match_operand:SDIM 2 "memory_operand") ;; memory
2000 (match_operand:SDIM 3 "nvptx_register_operand") ;; expected input
2001 (match_operand:SDIM 4 "nvptx_register_operand") ;; newval input
2002 (match_operand:SI 5 "const_int_operand") ;; is_weak
2003 (match_operand:SI 6 "const_int_operand") ;; success model
2004 (match_operand:SI 7 "const_int_operand")] ;; failure model
2007 if (nvptx_mem_local_p (operands[2]))
2008 emit_insn (gen_atomic_compare_and_swap<mode>_1_local
2009 (operands[1], operands[2], operands[3], operands[4],
2012 emit_insn (gen_atomic_compare_and_swap<mode>_1
2013 (operands[1], operands[2], operands[3], operands[4],
2016 rtx cond = gen_reg_rtx (BImode);
2017 emit_move_insn (cond, gen_rtx_EQ (BImode, operands[1], operands[3]));
2018 emit_insn (gen_sel_truesi (operands[0], cond, GEN_INT (1), GEN_INT (0)));
2022 (define_insn "atomic_compare_and_swap<mode>_1_local"
2023 [(set (match_operand:SDIM 0 "nvptx_register_operand" "=R")
2024 (unspec_volatile:SDIM
2025 [(match_operand:SDIM 1 "memory_operand" "+m")
2026 (match_operand:SDIM 2 "nvptx_nonmemory_operand" "Ri")
2027 (match_operand:SDIM 3 "nvptx_nonmemory_operand" "Ri")
2028 (match_operand:SI 4 "const_int_operand")]
2031 (unspec_volatile:SDIM [(const_int 0)] UNSPECV_CAS_LOCAL))]
2034 output_asm_insn ("{", NULL);
2035 output_asm_insn ("\\t" ".reg.pred" "\\t" "%%eq_p;", NULL);
2036 output_asm_insn ("\\t" ".reg%t0" "\\t" "%%val;", operands);
2037 output_asm_insn ("\\t" "ld%A1%t0" "\\t" "%%val,%1;", operands);
2038 output_asm_insn ("\\t" "setp.eq%t0" "\\t" "%%eq_p, %%val, %2;",
2040 output_asm_insn ("@%%eq_p\\t" "st%A1%t0" "\\t" "%1,%3;", operands);
2041 output_asm_insn ("\\t" "mov%t0" "\\t" "%0,%%val;", operands);
2042 output_asm_insn ("}", NULL);
2045 [(set_attr "predicable" "no")])
2047 (define_insn "atomic_compare_and_swap<mode>_1"
2048 [(set (match_operand:SDIM 0 "nvptx_register_operand" "=R")
2049 (unspec_volatile:SDIM
2050 [(match_operand:SDIM 1 "memory_operand" "+m")
2051 (match_operand:SDIM 2 "nvptx_nonmemory_operand" "Ri")
2052 (match_operand:SDIM 3 "nvptx_nonmemory_operand" "Ri")
2053 (match_operand:SI 4 "const_int_operand")]
2056 (unspec_volatile:SDIM [(const_int 0)] UNSPECV_CAS))]
2060 = "%.\\tatom%A1.cas.b%T0\\t%x0, %1, %2, %3;";
2061 return nvptx_output_atomic_insn (t, operands, 1, 4);
2063 [(set_attr "atomic" "true")])
2065 (define_insn "atomic_exchange<mode>"
2066 [(set (match_operand:SDIM 0 "nvptx_register_operand" "=R") ;; output
2067 (unspec_volatile:SDIM
2068 [(match_operand:SDIM 1 "memory_operand" "+m") ;; memory
2069 (match_operand:SI 3 "const_int_operand")] ;; model
2072 (match_operand:SDIM 2 "nvptx_nonmemory_operand" "Ri"))] ;; input
2075 if (nvptx_mem_local_p (operands[1]))
2077 output_asm_insn ("{", NULL);
2078 output_asm_insn ("\\t" ".reg%t0" "\\t" "%%val;", operands);
2079 output_asm_insn ("%.\\t" "ld%A1%t0" "\\t" "%%val,%1;", operands);
2080 output_asm_insn ("%.\\t" "st%A1%t0" "\\t" "%1,%2;", operands);
2081 output_asm_insn ("%.\\t" "mov%t0" "\\t" "%0,%%val;", operands);
2082 output_asm_insn ("}", NULL);
2086 = "%.\tatom%A1.exch.b%T0\t%x0, %1, %2;";
2087 return nvptx_output_atomic_insn (t, operands, 1, 3);
2089 [(set_attr "atomic" "true")])
2091 (define_expand "atomic_store<mode>"
2092 [(match_operand:SDIM 0 "memory_operand" "=m") ;; memory
2093 (match_operand:SDIM 1 "nvptx_nonmemory_operand" "Ri") ;; input
2094 (match_operand:SI 2 "const_int_operand")] ;; model
2097 struct address_info info;
2098 decompose_mem_address (&info, operands[0]);
2099 if (info.base != NULL && REG_P (*info.base)
2100 && REGNO_PTR_FRAME_P (REGNO (*info.base)))
2102 emit_insn (gen_mov<mode> (operands[0], operands[1]));
2108 emit_insn (gen_nvptx_atomic_store_sm70<mode> (operands[0], operands[1],
2113 bool maybe_shared_p = nvptx_mem_maybe_shared_p (operands[0]);
2114 if (!maybe_shared_p)
2115 /* Fall back to expand_atomic_store. */
2118 emit_insn (gen_nvptx_atomic_store<mode> (operands[0], operands[1],
2123 (define_insn "nvptx_atomic_store_sm70<mode>"
2124 [(set (match_operand:SDIM 0 "memory_operand" "+m") ;; memory
2125 (unspec_volatile:SDIM
2126 [(match_operand:SDIM 1 "nvptx_nonmemory_operand" "Ri") ;; input
2127 (match_operand:SI 2 "const_int_operand")] ;; model
2132 = "%.\tst%A0.b%T0\t%0, %1;";
2133 return nvptx_output_atomic_insn (t, operands, 0, 2);
2135 [(set_attr "atomic" "false")]) ;; Note: st is not an atomic insn.
2137 (define_insn "nvptx_atomic_store<mode>"
2138 [(set (match_operand:SDIM 0 "memory_operand" "+m") ;; memory
2139 (unspec_volatile:SDIM
2140 [(match_operand:SDIM 1 "nvptx_nonmemory_operand" "Ri") ;; input
2141 (match_operand:SI 2 "const_int_operand")] ;; model
2146 = "%.\tatom%A0.exch.b%T0\t_, %0, %1;";
2147 return nvptx_output_atomic_insn (t, operands, 0, 2);
2149 [(set_attr "atomic" "true")])
2151 (define_insn "atomic_fetch_add<mode>"
2152 [(set (match_operand:SDIM 1 "memory_operand" "+m")
2153 (unspec_volatile:SDIM
2154 [(plus:SDIM (match_dup 1)
2155 (match_operand:SDIM 2 "nvptx_nonmemory_operand" "Ri"))
2156 (match_operand:SI 3 "const_int_operand")] ;; model
2158 (set (match_operand:SDIM 0 "nvptx_register_operand" "=R")
2162 if (nvptx_mem_local_p (operands[1]))
2164 output_asm_insn ("{", NULL);
2165 output_asm_insn ("\\t" ".reg%t0" "\\t" "%%val;", operands);
2166 output_asm_insn ("\\t" ".reg%t0" "\\t" "%%update;", operands);
2167 output_asm_insn ("%.\\t" "ld%A1%t0" "\\t" "%%val,%1;", operands);
2168 output_asm_insn ("%.\\t" "add%t0" "\\t" "%%update,%%val,%2;",
2170 output_asm_insn ("%.\\t" "st%A1%t0" "\\t" "%1,%%update;", operands);
2171 output_asm_insn ("%.\\t" "mov%t0" "\\t" "%0,%%val;", operands);
2172 output_asm_insn ("}", NULL);
2176 = "%.\\tatom%A1.add%t0\\t%x0, %1, %2;";
2177 return nvptx_output_atomic_insn (t, operands, 1, 3);
2179 [(set_attr "atomic" "true")])
2181 (define_insn "atomic_fetch_addsf"
2182 [(set (match_operand:SF 1 "memory_operand" "+m")
2184 [(plus:SF (match_dup 1)
2185 (match_operand:SF 2 "nvptx_nonmemory_operand" "RF"))
2186 (match_operand:SI 3 "const_int_operand")] ;; model
2188 (set (match_operand:SF 0 "nvptx_register_operand" "=R")
2192 if (nvptx_mem_local_p (operands[1]))
2194 output_asm_insn ("{", NULL);
2195 output_asm_insn ("\\t" ".reg%t0" "\\t" "%%val;", operands);
2196 output_asm_insn ("\\t" ".reg%t0" "\\t" "%%update;", operands);
2197 output_asm_insn ("%.\\t" "ld%A1%t0" "\\t" "%%val,%1;", operands);
2198 output_asm_insn ("%.\\t" "add%t0" "\\t" "%%update,%%val,%2;",
2200 output_asm_insn ("%.\\t" "st%A1%t0" "\\t" "%1,%%update;", operands);
2201 output_asm_insn ("%.\\t" "mov%t0" "\\t" "%0,%%val;", operands);
2202 output_asm_insn ("}", NULL);
2206 = "%.\\tatom%A1.add%t0\\t%x0, %1, %2;";
2207 return nvptx_output_atomic_insn (t, operands, 1, 3);
2209 [(set_attr "atomic" "true")])
2211 (define_insn "atomic_fetch_<logic><mode>"
2212 [(set (match_operand:SDIM 1 "memory_operand" "+m")
2213 (unspec_volatile:SDIM
2214 [(any_logic:SDIM (match_dup 1)
2215 (match_operand:SDIM 2 "nvptx_nonmemory_operand" "Ri"))
2216 (match_operand:SI 3 "const_int_operand")] ;; model
2218 (set (match_operand:SDIM 0 "nvptx_register_operand" "=R")
2220 "<MODE>mode == SImode || TARGET_SM35"
2222 if (nvptx_mem_local_p (operands[1]))
2224 output_asm_insn ("{", NULL);
2225 output_asm_insn ("\\t" ".reg.b%T0" "\\t" "%%val;", operands);
2226 output_asm_insn ("\\t" ".reg.b%T0" "\\t" "%%update;", operands);
2227 output_asm_insn ("%.\\t" "ld%A1%t0" "\\t" "%%val,%1;", operands);
2228 output_asm_insn ("%.\\t" "<logic>.b%T0" "\\t" "%%update,%%val,%2;",
2230 output_asm_insn ("%.\\t" "st%A1%t0" "\\t" "%1,%%update;", operands);
2231 output_asm_insn ("%.\\t" "mov%t0" "\\t" "%0,%%val;", operands);
2232 output_asm_insn ("}", NULL);
2236 = "%.\\tatom%A1.<logic>.b%T0\\t%x0, %1, %2;";
2237 return nvptx_output_atomic_insn (t, operands, 1, 3);
2240 [(set_attr "atomic" "true")])
2242 (define_expand "atomic_test_and_set"
2243 [(match_operand:SI 0 "nvptx_register_operand") ;; bool success output
2244 (match_operand:QI 1 "memory_operand") ;; memory
2245 (match_operand:SI 2 "const_int_operand")] ;; model
2250 libfunc = init_one_libfunc ("__atomic_test_and_set_1");
2251 addr = convert_memory_address (ptr_mode, XEXP (operands[1], 0));
2252 emit_library_call_value (libfunc, operands[0], LCT_NORMAL, SImode,
2254 operands[2], SImode);
2258 (define_insn "nvptx_barsync"
2259 [(unspec_volatile [(match_operand:SI 0 "nvptx_nonmemory_operand" "Ri")
2260 (match_operand:SI 1 "const_int_operand")]
2264 if (INTVAL (operands[1]) == 0)
2265 return (TARGET_PTX_6_0
2266 ? "\\tbarrier.sync.aligned\\t%0;"
2267 : "\\tbar.sync\\t%0;");
2269 return (TARGET_PTX_6_0
2270 ? "\\tbarrier.sync\\t%0, %1;"
2271 : "\\tbar.sync\\t%0, %1;");
2273 [(set_attr "predicable" "no")])
2275 (define_insn "nvptx_warpsync"
2276 [(unspec_volatile [(const_int 0)] UNSPECV_WARPSYNC)]
2278 "%.\\tbar.warp.sync\\t0xffffffff;")
2280 (define_int_iterator BARRED
2283 UNSPECV_BARRED_POPC])
2284 (define_int_attr barred_op
2285 [(UNSPECV_BARRED_AND "and")
2286 (UNSPECV_BARRED_OR "or")
2287 (UNSPECV_BARRED_POPC "popc")])
2288 (define_int_attr barred_mode
2289 [(UNSPECV_BARRED_AND "BI")
2290 (UNSPECV_BARRED_OR "BI")
2291 (UNSPECV_BARRED_POPC "SI")])
2292 (define_int_attr barred_ptxtype
2293 [(UNSPECV_BARRED_AND "pred")
2294 (UNSPECV_BARRED_OR "pred")
2295 (UNSPECV_BARRED_POPC "u32")])
2297 (define_insn "nvptx_barred_<barred_op>"
2298 [(set (match_operand:<barred_mode> 0 "nvptx_register_operand" "=R")
2300 [(match_operand:SI 1 "nvptx_nonmemory_operand" "Ri")
2301 (match_operand:SI 2 "nvptx_nonmemory_operand" "Ri")
2302 (match_operand:SI 3 "const_int_operand" "i")
2303 (match_operand:BI 4 "nvptx_register_operand" "R")]
2306 "\\tbar.red.<barred_op>.<barred_ptxtype> \\t%0, %1, %2, %p3%4;";"
2307 [(set_attr "predicable" "no")])
2309 (define_insn "nvptx_uniform_warp_check"
2310 [(unspec_volatile [(const_int 0)] UNSPECV_UNIFORM_WARP_CHECK)]
2313 const char *insns[] = {
2315 "\\t" ".reg.b32" "\\t" "%%r_act;",
2316 "%.\\t" "vote.ballot.b32" "\\t" "%%r_act,1;",
2317 "\\t" ".reg.pred" "\\t" "%%r_do_abort;",
2318 "\\t" "mov.pred" "\\t" "%%r_do_abort,0;",
2319 "%.\\t" "setp.ne.b32" "\\t" "%%r_do_abort,%%r_act,"
2321 "@ %%r_do_abort\\t" "trap;",
2322 "@ %%r_do_abort\\t" "exit;",
2326 for (const char **p = &insns[0]; *p != NULL; p++)
2327 output_asm_insn (*p, NULL);
2331 (define_expand "memory_barrier"
2333 (unspec_volatile:BLK [(match_dup 0)] UNSPECV_MEMBAR))]
2336 operands[0] = gen_rtx_MEM (BLKmode, gen_rtx_SCRATCH (Pmode));
2337 MEM_VOLATILE_P (operands[0]) = 1;
2340 ;; Ptx defines the memory barriers membar.cta, membar.gl and membar.sys
2341 ;; (corresponding to cuda functions threadfence_block, threadfence and
2342 ;; threadfence_system). For the insn memory_barrier we use membar.sys. This
2343 ;; may be overconservative, but before using membar.gl instead we'll need to
2344 ;; explain in detail why it's safe to use. For now, use membar.sys.
2345 (define_insn "*memory_barrier"
2346 [(set (match_operand:BLK 0 "" "")
2347 (unspec_volatile:BLK [(match_dup 0)] UNSPECV_MEMBAR))]
2350 [(set_attr "predicable" "no")])
2352 (define_expand "nvptx_membar_cta"
2354 (unspec_volatile:BLK [(match_dup 0)] UNSPECV_MEMBAR_CTA))]
2357 operands[0] = gen_rtx_MEM (BLKmode, gen_rtx_SCRATCH (Pmode));
2358 MEM_VOLATILE_P (operands[0]) = 1;
2361 (define_insn "*nvptx_membar_cta"
2362 [(set (match_operand:BLK 0 "" "")
2363 (unspec_volatile:BLK [(match_dup 0)] UNSPECV_MEMBAR_CTA))]
2366 [(set_attr "predicable" "no")])
2368 (define_expand "nvptx_membar_gl"
2370 (unspec_volatile:BLK [(match_dup 0)] UNSPECV_MEMBAR_GL))]
2373 operands[0] = gen_rtx_MEM (BLKmode, gen_rtx_SCRATCH (Pmode));
2374 MEM_VOLATILE_P (operands[0]) = 1;
2377 (define_insn "*nvptx_membar_gl"
2378 [(set (match_operand:BLK 0 "" "")
2379 (unspec_volatile:BLK [(match_dup 0)] UNSPECV_MEMBAR_GL))]
2382 [(set_attr "predicable" "no")])
2384 (define_insn "nvptx_nounroll"
2385 [(unspec_volatile [(const_int 0)] UNSPECV_NOUNROLL)]
2387 "\\t.pragma \\\"nounroll\\\";"
2388 [(set_attr "predicable" "no")])
2390 (define_insn "nvptx_red_partition"
2391 [(set (match_operand:DI 0 "nonimmediate_operand" "=R")
2392 (unspec_volatile:DI [(match_operand:DI 1 "const_int_operand")]
2396 return nvptx_output_red_partition (operands[0], operands[1]);
2398 [(set_attr "predicable" "no")])
2400 ;; Expand QI mode operations using SI mode instructions.
2401 (define_code_iterator any_sbinary [plus minus smin smax])
2402 (define_code_attr sbinary [(plus "add") (minus "sub") (smin "smin") (smax "smax")])
2404 (define_code_iterator any_ubinary [and ior xor umin umax])
2405 (define_code_attr ubinary [(and "and") (ior "ior") (xor "xor") (umin "umin")
2408 (define_code_iterator any_sunary [neg abs])
2409 (define_code_attr sunary [(neg "neg") (abs "abs")])
2411 (define_code_iterator any_uunary [not])
2412 (define_code_attr uunary [(not "one_cmpl")])
2414 (define_expand "<sbinary>qi3"
2415 [(set (match_operand:QI 0 "nvptx_register_operand")
2416 (any_sbinary:QI (match_operand:QI 1 "nvptx_nonmemory_operand")
2417 (match_operand:QI 2 "nvptx_nonmemory_operand")))]
2420 rtx reg = gen_reg_rtx (SImode);
2421 rtx op0 = convert_modes (SImode, QImode, operands[1], 0);
2422 rtx op1 = convert_modes (SImode, QImode, operands[2], 0);
2423 if (<CODE> == MINUS)
2424 op0 = force_reg (SImode, op0);
2425 emit_insn (gen_<sbinary>si3 (reg, op0, op1));
2426 emit_insn (gen_truncsiqi2 (operands[0], reg));
2430 (define_expand "<ubinary>qi3"
2431 [(set (match_operand:QI 0 "nvptx_register_operand")
2432 (any_ubinary:QI (match_operand:QI 1 "nvptx_nonmemory_operand")
2433 (match_operand:QI 2 "nvptx_nonmemory_operand")))]
2436 rtx reg = gen_reg_rtx (SImode);
2437 rtx op0 = convert_modes (SImode, QImode, operands[1], 1);
2438 rtx op1 = convert_modes (SImode, QImode, operands[2], 1);
2439 emit_insn (gen_<ubinary>si3 (reg, op0, op1));
2440 emit_insn (gen_truncsiqi2 (operands[0], reg));
2444 (define_expand "<sunary>qi2"
2445 [(set (match_operand:QI 0 "nvptx_register_operand")
2446 (any_sunary:QI (match_operand:QI 1 "nvptx_nonmemory_operand")))]
2449 rtx reg = gen_reg_rtx (SImode);
2450 rtx op0 = convert_modes (SImode, QImode, operands[1], 0);
2451 emit_insn (gen_<sunary>si2 (reg, op0));
2452 emit_insn (gen_truncsiqi2 (operands[0], reg));
2456 (define_expand "<uunary>qi2"
2457 [(set (match_operand:QI 0 "nvptx_register_operand")
2458 (any_uunary:QI (match_operand:QI 1 "nvptx_nonmemory_operand")))]
2461 rtx reg = gen_reg_rtx (SImode);
2462 rtx op0 = convert_modes (SImode, QImode, operands[1], 1);
2463 emit_insn (gen_<uunary>si2 (reg, op0));
2464 emit_insn (gen_truncsiqi2 (operands[0], reg));
2468 (define_expand "cstoreqi4"
2469 [(set (match_operand:SI 0 "nvptx_register_operand")
2470 (match_operator:SI 1 "nvptx_comparison_operator"
2471 [(match_operand:QI 2 "nvptx_nonmemory_operand")
2472 (match_operand:QI 3 "nvptx_nonmemory_operand")]))]
2475 rtx reg = gen_reg_rtx (BImode);
2476 enum rtx_code code = GET_CODE (operands[1]);
2477 int unsignedp = unsigned_condition_p (code);
2478 rtx op2 = convert_modes (SImode, QImode, operands[2], unsignedp);
2479 rtx op3 = convert_modes (SImode, QImode, operands[3], unsignedp);
2480 rtx cmp = gen_rtx_fmt_ee (code, SImode, op2, op3);
2481 emit_insn (gen_cmpsi (reg, cmp, op2, op3));
2482 emit_insn (gen_setccsi_from_bi (operands[0], reg));
2486 (define_insn "*ext_truncsi2_qi"
2487 [(set (match_operand:SI 0 "nvptx_register_operand" "=R")
2489 (truncate:QI (match_operand:SI 1 "nvptx_register_operand" "R"))))]
2491 "%.\\tcvt.s32.s8\\t%0, %1;")
2493 (define_insn "*zext_truncsi2_qi"
2494 [(set (match_operand:SI 0 "nvptx_register_operand" "=R")
2496 (truncate:QI (match_operand:SI 1 "nvptx_register_operand" "R"))))]
2498 "%.\\tcvt.u32.u8\\t%0, %1;")