gcc/config/nvptx/nvptx.md

   1 ;; Machine description for NVPTX.
   2 ;; Copyright (C) 2014-2019 Free Software Foundation, Inc.
   3 ;; Contributed by Bernd Schmidt <bernds@codesourcery.com>
   4 ;;
   5 ;; This file is part of GCC.
   6 ;;
   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)
  10 ;; any later version.
  11 ;;
  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.
  16 ;;
  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/>.
  20
  21 (define_c_enum "unspec" [
  22    UNSPEC_ARG_REG
  23
  24    UNSPEC_COPYSIGN
  25    UNSPEC_LOG2
  26    UNSPEC_EXP2
  27    UNSPEC_SIN
  28    UNSPEC_COS
  29
  30    UNSPEC_FPINT_FLOOR
  31    UNSPEC_FPINT_BTRUNC
  32    UNSPEC_FPINT_CEIL
  33    UNSPEC_FPINT_NEARBYINT
  34
  35    UNSPEC_BITREV
  36
  37    UNSPEC_ALLOCA
  38
  39    UNSPEC_SET_SOFTSTACK
  40
  41    UNSPEC_DIM_SIZE
  42
  43    UNSPEC_BIT_CONV
  44
  45    UNSPEC_VOTE_BALLOT
  46
  47    UNSPEC_LANEID
  48
  49    UNSPEC_SHUFFLE
  50    UNSPEC_BR_UNIFIED
  51 ])
  52
  53 (define_c_enum "unspecv" [
  54    UNSPECV_LOCK
  55    UNSPECV_CAS
  56    UNSPECV_XCHG
  57    UNSPECV_BARSYNC
  58    UNSPECV_MEMBAR
  59    UNSPECV_MEMBAR_CTA
  60    UNSPECV_DIM_POS
  61
  62    UNSPECV_FORK
  63    UNSPECV_FORKED
  64    UNSPECV_JOINING
  65    UNSPECV_JOIN
  66
  67    UNSPECV_NOUNROLL
  68
  69    UNSPECV_SIMT_ENTER
  70    UNSPECV_SIMT_EXIT
  71
  72    UNSPECV_RED_PART
  73 ])
  74
  75 (define_attr "subregs_ok" "false,true"
  76   (const_string "false"))
  77
  78 (define_attr "atomic" "false,true"
  79   (const_string "false"))
  80
  81 ;; The nvptx operand predicates, in general, don't permit subregs and
  82 ;; only literal constants, which differ from the generic ones, which
  83 ;; permit subregs and symbolc constants (as appropriate)
  84 (define_predicate "nvptx_register_operand"
  85   (match_code "reg")
  86 {
  87   return register_operand (op, mode);
  88 })
  89
  90 (define_predicate "nvptx_nonimmediate_operand"
  91   (match_code "mem,reg")
  92 {
  93   return (REG_P (op) ? register_operand (op, mode)
  94           : memory_operand (op, mode));
  95 })
  96
  97 (define_predicate "nvptx_nonmemory_operand"
  98   (match_code "reg,const_int,const_double")
  99 {
 100   return (REG_P (op) ? register_operand (op, mode)
 101           : immediate_operand (op, mode));
 102 })
 103
 104 (define_predicate "const0_operand"
 105   (and (match_code "const_int")
 106        (match_test "op == const0_rtx")))
 107
 108 ;; True if this operator is valid for predication.
 109 (define_predicate "predicate_operator"
 110   (match_code "eq,ne"))
 111
 112 (define_predicate "ne_operator"
 113   (match_code "ne"))
 114
 115 (define_predicate "nvptx_comparison_operator"
 116   (match_code "eq,ne,le,ge,lt,gt,leu,geu,ltu,gtu"))
 117
 118 (define_predicate "nvptx_float_comparison_operator"
 119   (match_code "eq,ne,le,ge,lt,gt,uneq,unle,unge,unlt,ungt,unordered,ordered"))
 120
 121 ;; Test for a valid operand for a call instruction.
 122 (define_predicate "call_insn_operand"
 123   (match_code "symbol_ref,reg")
 124 {
 125   return REG_P (op) || SYMBOL_REF_FUNCTION_P (op);
 126 })
 127
 128 ;; Return true if OP is a call with parallel USEs of the argument
 129 ;; pseudos.
 130 (define_predicate "call_operation"
 131   (match_code "parallel")
 132 {
 133   int arg_end = XVECLEN (op, 0);
 134
 135   for (int i = 1; i < arg_end; i++)
 136     {
 137       rtx elt = XVECEXP (op, 0, i);
 138
 139       if (GET_CODE (elt) != USE || !REG_P (XEXP (elt, 0)))
 140         return false;
 141     }
 142   return true;
 143 })
 144
 145 (define_attr "predicable" "false,true"
 146   (const_string "true"))
 147
 148 (define_cond_exec
 149   [(match_operator 0 "predicate_operator"
 150       [(match_operand:BI 1 "nvptx_register_operand" "")
 151        (match_operand:BI 2 "const0_operand" "")])]
 152   ""
 153   ""
 154   )
 155
 156 (define_constraint "P0"
 157   "An integer with the value 0."
 158   (and (match_code "const_int")
 159        (match_test "ival == 0")))
 160
 161 (define_constraint "P1"
 162   "An integer with the value 1."
 163   (and (match_code "const_int")
 164        (match_test "ival == 1")))
 165
 166 (define_constraint "Pn"
 167   "An integer with the value -1."
 168   (and (match_code "const_int")
 169        (match_test "ival == -1")))
 170
 171 (define_constraint "R"
 172   "A pseudo register."
 173   (match_code "reg"))
 174
 175 (define_constraint "Ia"
 176   "Any integer constant."
 177   (and (match_code "const_int") (match_test "true")))
 178
 179 (define_mode_iterator QHSDISDFM [QI HI SI DI SF DF])
 180 (define_mode_iterator QHSDIM [QI HI SI DI])
 181 (define_mode_iterator HSDIM [HI SI DI])
 182 (define_mode_iterator BHSDIM [BI HI SI DI])
 183 (define_mode_iterator SDIM [SI DI])
 184 (define_mode_iterator SDISDFM [SI DI SF DF])
 185 (define_mode_iterator QHIM [QI HI])
 186 (define_mode_iterator QHSIM [QI HI SI])
 187 (define_mode_iterator SDFM [SF DF])
 188 (define_mode_iterator SDCM [SC DC])
 189 (define_mode_iterator BITS [SI SF])
 190 (define_mode_iterator BITD [DI DF])
 191 (define_mode_iterator VECIM [V2SI V2DI])
 192
 193 ;; This mode iterator allows :P to be used for patterns that operate on
 194 ;; pointer-sized quantities.  Exactly one of the two alternatives will match.
 195 (define_mode_iterator P [(SI "Pmode == SImode") (DI "Pmode == DImode")])
 196
 197 ;; We should get away with not defining memory alternatives, since we don't
 198 ;; get variables in this mode and pseudos are never spilled.
 199 (define_insn "movbi"
 200   [(set (match_operand:BI 0 "nvptx_register_operand" "=R,R,R")
 201         (match_operand:BI 1 "nvptx_nonmemory_operand" "R,P0,Pn"))]
 202   ""
 203   "@
 204    %.\\tmov%t0\\t%0, %1;
 205    %.\\tsetp.eq.u32\\t%0, 1, 0;
 206    %.\\tsetp.eq.u32\\t%0, 1, 1;")
 207
 208 (define_insn "*mov<mode>_insn"
 209   [(set (match_operand:VECIM 0 "nonimmediate_operand" "=R,R,m")
 210         (match_operand:VECIM 1 "general_operand" "Ri,m,R"))]
 211   "!MEM_P (operands[0]) || REG_P (operands[1])"
 212 {
 213   if (which_alternative == 1)
 214     return "%.\\tld%A1%u1\\t%0, %1;";
 215   if (which_alternative == 2)
 216     return "%.\\tst%A0%u0\\t%0, %1;";
 217
 218   return nvptx_output_mov_insn (operands[0], operands[1]);
 219 }
 220   [(set_attr "subregs_ok" "true")])
 221
 222 (define_insn "*mov<mode>_insn"
 223   [(set (match_operand:QHSDIM 0 "nonimmediate_operand" "=R,R,m")
 224         (match_operand:QHSDIM 1 "general_operand" "Ri,m,R"))]
 225   "!MEM_P (operands[0]) || REG_P (operands[1])"
 226 {
 227   if (which_alternative == 1)
 228     return "%.\\tld%A1%u1\\t%0, %1;";
 229   if (which_alternative == 2)
 230     return "%.\\tst%A0%u0\\t%0, %1;";
 231
 232   return nvptx_output_mov_insn (operands[0], operands[1]);
 233 }
 234   [(set_attr "subregs_ok" "true")])
 235
 236 (define_insn "*mov<mode>_insn"
 237   [(set (match_operand:SDFM 0 "nonimmediate_operand" "=R,R,m")
 238         (match_operand:SDFM 1 "general_operand" "RF,m,R"))]
 239   "!MEM_P (operands[0]) || REG_P (operands[1])"
 240 {
 241   if (which_alternative == 1)
 242     return "%.\\tld%A1%u0\\t%0, %1;";
 243   if (which_alternative == 2)
 244     return "%.\\tst%A0%u1\\t%0, %1;";
 245
 246   return nvptx_output_mov_insn (operands[0], operands[1]);
 247 }
 248   [(set_attr "subregs_ok" "true")])
 249
 250 (define_insn "load_arg_reg<mode>"
 251   [(set (match_operand:QHIM 0 "nvptx_register_operand" "=R")
 252         (unspec:QHIM [(match_operand 1 "const_int_operand" "n")]
 253                      UNSPEC_ARG_REG))]
 254   ""
 255   "%.\\tcvt%t0.u32\\t%0, %%ar%1;")
 256
 257 (define_insn "load_arg_reg<mode>"
 258   [(set (match_operand:SDISDFM 0 "nvptx_register_operand" "=R")
 259         (unspec:SDISDFM [(match_operand 1 "const_int_operand" "n")]
 260                         UNSPEC_ARG_REG))]
 261   ""
 262   "%.\\tmov%t0\\t%0, %%ar%1;")
 263
 264  (define_expand "mov<mode>"
 265   [(set (match_operand:VECIM 0 "nonimmediate_operand" "")
 266         (match_operand:VECIM 1 "general_operand" ""))]
 267   ""
 268 {
 269   if (MEM_P (operands[0]) && !REG_P (operands[1]))
 270     {
 271       rtx tmp = gen_reg_rtx (<MODE>mode);
 272       emit_move_insn (tmp, operands[1]);
 273       emit_move_insn (operands[0], tmp);
 274       DONE;
 275     }
 276 })
 277
 278 (define_expand "mov<mode>"
 279   [(set (match_operand:QHSDISDFM 0 "nonimmediate_operand" "")
 280         (match_operand:QHSDISDFM 1 "general_operand" ""))]
 281   ""
 282 {
 283   if (MEM_P (operands[0]) && !REG_P (operands[1]))
 284     {
 285       rtx tmp = gen_reg_rtx (<MODE>mode);
 286       emit_move_insn (tmp, operands[1]);
 287       emit_move_insn (operands[0], tmp);
 288       DONE;
 289     }
 290
 291   if (GET_CODE (operands[1]) == LABEL_REF)
 292     sorry ("target cannot support label values");
 293 })
 294
 295 (define_insn "zero_extendqihi2"
 296   [(set (match_operand:HI 0 "nvptx_register_operand" "=R,R")
 297         (zero_extend:HI (match_operand:QI 1 "nvptx_nonimmediate_operand" "R,m")))]
 298   ""
 299   "@
 300    %.\\tcvt.u16.u%T1\\t%0, %1;
 301    %.\\tld%A1.u8\\t%0, %1;"
 302   [(set_attr "subregs_ok" "true")])
 303
 304 (define_insn "zero_extend<mode>si2"
 305   [(set (match_operand:SI 0 "nvptx_register_operand" "=R,R")
 306         (zero_extend:SI (match_operand:QHIM 1 "nvptx_nonimmediate_operand" "R,m")))]
 307   ""
 308   "@
 309    %.\\tcvt.u32.u%T1\\t%0, %1;
 310    %.\\tld%A1.u%T1\\t%0, %1;"
 311   [(set_attr "subregs_ok" "true")])
 312
 313 (define_insn "zero_extend<mode>di2"
 314   [(set (match_operand:DI 0 "nvptx_register_operand" "=R,R")
 315         (zero_extend:DI (match_operand:QHSIM 1 "nvptx_nonimmediate_operand" "R,m")))]
 316   ""
 317   "@
 318    %.\\tcvt.u64.u%T1\\t%0, %1;
 319    %.\\tld%A1%u1\\t%0, %1;"
 320   [(set_attr "subregs_ok" "true")])
 321
 322 (define_insn "extend<mode>si2"
 323   [(set (match_operand:SI 0 "nvptx_register_operand" "=R,R")
 324         (sign_extend:SI (match_operand:QHIM 1 "nvptx_nonimmediate_operand" "R,m")))]
 325   ""
 326   "@
 327    %.\\tcvt.s32.s%T1\\t%0, %1;
 328    %.\\tld%A1.s%T1\\t%0, %1;"
 329   [(set_attr "subregs_ok" "true")])
 330
 331 (define_insn "extend<mode>di2"
 332   [(set (match_operand:DI 0 "nvptx_register_operand" "=R,R")
 333         (sign_extend:DI (match_operand:QHSIM 1 "nvptx_nonimmediate_operand" "R,m")))]
 334   ""
 335   "@
 336    %.\\tcvt.s64.s%T1\\t%0, %1;
 337    %.\\tld%A1.s%T1\\t%0, %1;"
 338   [(set_attr "subregs_ok" "true")])
 339
 340 (define_insn "trunchiqi2"
 341   [(set (match_operand:QI 0 "nvptx_nonimmediate_operand" "=R,m")
 342         (truncate:QI (match_operand:HI 1 "nvptx_register_operand" "R,R")))]
 343   ""
 344   "@
 345    %.\\tcvt%t0.u16\\t%0, %1;
 346    %.\\tst%A0.u8\\t%0, %1;"
 347   [(set_attr "subregs_ok" "true")])
 348
 349 (define_insn "truncsi<mode>2"
 350   [(set (match_operand:QHIM 0 "nvptx_nonimmediate_operand" "=R,m")
 351         (truncate:QHIM (match_operand:SI 1 "nvptx_register_operand" "R,R")))]
 352   ""
 353   "@
 354    %.\\tcvt%t0.u32\\t%0, %1;
 355    %.\\tst%A0.u%T0\\t%0, %1;"
 356   [(set_attr "subregs_ok" "true")])
 357
 358 (define_insn "truncdi<mode>2"
 359   [(set (match_operand:QHSIM 0 "nvptx_nonimmediate_operand" "=R,m")
 360         (truncate:QHSIM (match_operand:DI 1 "nvptx_register_operand" "R,R")))]
 361   ""
 362   "@
 363    %.\\tcvt%t0.u64\\t%0, %1;
 364    %.\\tst%A0.u%T0\\t%0, %1;"
 365   [(set_attr "subregs_ok" "true")])
 366
 367 ;; Integer arithmetic
 368
 369 (define_insn "add<mode>3"
 370   [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
 371         (plus:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")
 372                     (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
 373   ""
 374   "%.\\tadd%t0\\t%0, %1, %2;")
 375
 376 (define_insn "sub<mode>3"
 377   [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
 378         (minus:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")
 379                      (match_operand:HSDIM 2 "nvptx_register_operand" "R")))]
 380   ""
 381   "%.\\tsub%t0\\t%0, %1, %2;")
 382
 383 (define_insn "mul<mode>3"
 384   [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
 385         (mult:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")
 386                     (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
 387   ""
 388   "%.\\tmul.lo%t0\\t%0, %1, %2;")
 389
 390 (define_insn "*mad<mode>3"
 391   [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
 392         (plus:HSDIM (mult:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")
 393                                 (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri"))
 394                     (match_operand:HSDIM 3 "nvptx_nonmemory_operand" "Ri")))]
 395   ""
 396   "%.\\tmad.lo%t0\\t%0, %1, %2, %3;")
 397
 398 (define_insn "div<mode>3"
 399   [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
 400         (div:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")
 401                    (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
 402   ""
 403   "%.\\tdiv.s%T0\\t%0, %1, %2;")
 404
 405 (define_insn "udiv<mode>3"
 406   [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
 407         (udiv:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")
 408                    (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
 409   ""
 410   "%.\\tdiv.u%T0\\t%0, %1, %2;")
 411
 412 (define_insn "mod<mode>3"
 413   [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
 414         (mod:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "Ri")
 415                    (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
 416   ""
 417   "%.\\trem.s%T0\\t%0, %1, %2;")
 418
 419 (define_insn "umod<mode>3"
 420   [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
 421         (umod:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "Ri")
 422                     (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
 423   ""
 424   "%.\\trem.u%T0\\t%0, %1, %2;")
 425
 426 (define_insn "smin<mode>3"
 427   [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
 428         (smin:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")
 429                     (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
 430   ""
 431   "%.\\tmin.s%T0\\t%0, %1, %2;")
 432
 433 (define_insn "umin<mode>3"
 434   [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
 435         (umin:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")
 436                     (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
 437   ""
 438   "%.\\tmin.u%T0\\t%0, %1, %2;")
 439
 440 (define_insn "smax<mode>3"
 441   [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
 442         (smax:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")
 443                     (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
 444   ""
 445   "%.\\tmax.s%T0\\t%0, %1, %2;")
 446
 447 (define_insn "umax<mode>3"
 448   [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
 449         (umax:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")
 450                     (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
 451   ""
 452   "%.\\tmax.u%T0\\t%0, %1, %2;")
 453
 454 (define_insn "abs<mode>2"
 455   [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
 456         (abs:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")))]
 457   ""
 458   "%.\\tabs.s%T0\\t%0, %1;")
 459
 460 (define_insn "neg<mode>2"
 461   [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
 462         (neg:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")))]
 463   ""
 464   "%.\\tneg.s%T0\\t%0, %1;")
 465
 466 (define_insn "one_cmpl<mode>2"
 467   [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
 468         (not:HSDIM (match_operand:HSDIM 1 "nvptx_register_operand" "R")))]
 469   ""
 470   "%.\\tnot.b%T0\\t%0, %1;")
 471
 472 (define_insn "bitrev<mode>2"
 473   [(set (match_operand:SDIM 0 "nvptx_register_operand" "=R")
 474         (unspec:SDIM [(match_operand:SDIM 1 "nvptx_register_operand" "R")]
 475                      UNSPEC_BITREV))]
 476   ""
 477   "%.\\tbrev.b%T0\\t%0, %1;")
 478
 479 (define_insn "clz<mode>2"
 480   [(set (match_operand:SI 0 "nvptx_register_operand" "=R")
 481         (clz:SI (match_operand:SDIM 1 "nvptx_register_operand" "R")))]
 482   ""
 483   "%.\\tclz.b%T1\\t%0, %1;")
 484
 485 (define_expand "ctz<mode>2"
 486   [(set (match_operand:SI 0 "nvptx_register_operand" "")
 487         (ctz:SI (match_operand:SDIM 1 "nvptx_register_operand" "")))]
 488   ""
 489 {
 490   rtx tmpreg = gen_reg_rtx (<MODE>mode);
 491   emit_insn (gen_bitrev<mode>2 (tmpreg, operands[1]));
 492   emit_insn (gen_clz<mode>2 (operands[0], tmpreg));
 493   DONE;
 494 })
 495
 496 ;; Shifts
 497
 498 (define_insn "ashl<mode>3"
 499   [(set (match_operand:SDIM 0 "nvptx_register_operand" "=R")
 500         (ashift:SDIM (match_operand:SDIM 1 "nvptx_register_operand" "R")
 501                      (match_operand:SI 2 "nvptx_nonmemory_operand" "Ri")))]
 502   ""
 503   "%.\\tshl.b%T0\\t%0, %1, %2;")
 504
 505 (define_insn "ashr<mode>3"
 506   [(set (match_operand:SDIM 0 "nvptx_register_operand" "=R")
 507         (ashiftrt:SDIM (match_operand:SDIM 1 "nvptx_register_operand" "R")
 508                        (match_operand:SI 2 "nvptx_nonmemory_operand" "Ri")))]
 509   ""
 510   "%.\\tshr.s%T0\\t%0, %1, %2;")
 511
 512 (define_insn "lshr<mode>3"
 513   [(set (match_operand:SDIM 0 "nvptx_register_operand" "=R")
 514         (lshiftrt:SDIM (match_operand:SDIM 1 "nvptx_register_operand" "R")
 515                        (match_operand:SI 2 "nvptx_nonmemory_operand" "Ri")))]
 516   ""
 517   "%.\\tshr.u%T0\\t%0, %1, %2;")
 518
 519 ;; Logical operations
 520
 521 (define_insn "and<mode>3"
 522   [(set (match_operand:BHSDIM 0 "nvptx_register_operand" "=R")
 523         (and:BHSDIM (match_operand:BHSDIM 1 "nvptx_register_operand" "R")
 524                     (match_operand:BHSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
 525   ""
 526   "%.\\tand.b%T0\\t%0, %1, %2;")
 527
 528 (define_insn "ior<mode>3"
 529   [(set (match_operand:BHSDIM 0 "nvptx_register_operand" "=R")
 530         (ior:BHSDIM (match_operand:BHSDIM 1 "nvptx_register_operand" "R")
 531                     (match_operand:BHSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
 532   ""
 533   "%.\\tor.b%T0\\t%0, %1, %2;")
 534
 535 (define_insn "xor<mode>3"
 536   [(set (match_operand:BHSDIM 0 "nvptx_register_operand" "=R")
 537         (xor:BHSDIM (match_operand:BHSDIM 1 "nvptx_register_operand" "R")
 538                     (match_operand:BHSDIM 2 "nvptx_nonmemory_operand" "Ri")))]
 539   ""
 540   "%.\\txor.b%T0\\t%0, %1, %2;")
 541
 542 ;; Comparisons and branches
 543
 544 (define_insn "*cmp<mode>"
 545   [(set (match_operand:BI 0 "nvptx_register_operand" "=R")
 546         (match_operator:BI 1 "nvptx_comparison_operator"
 547            [(match_operand:HSDIM 2 "nvptx_register_operand" "R")
 548             (match_operand:HSDIM 3 "nvptx_nonmemory_operand" "Ri")]))]
 549   ""
 550   "%.\\tsetp%c1\\t%0, %2, %3;")
 551
 552 (define_insn "*cmp<mode>"
 553   [(set (match_operand:BI 0 "nvptx_register_operand" "=R")
 554         (match_operator:BI 1 "nvptx_float_comparison_operator"
 555            [(match_operand:SDFM 2 "nvptx_register_operand" "R")
 556             (match_operand:SDFM 3 "nvptx_nonmemory_operand" "RF")]))]
 557   ""
 558   "%.\\tsetp%c1\\t%0, %2, %3;")
 559
 560 (define_insn "jump"
 561   [(set (pc)
 562         (label_ref (match_operand 0 "" "")))]
 563   ""
 564   "%.\\tbra\\t%l0;")
 565
 566 (define_insn "br_true"
 567   [(set (pc)
 568         (if_then_else (ne (match_operand:BI 0 "nvptx_register_operand" "R")
 569                           (const_int 0))
 570                       (label_ref (match_operand 1 "" ""))
 571                       (pc)))]
 572   ""
 573   "%j0\\tbra\\t%l1;"
 574   [(set_attr "predicable" "false")])
 575
 576 (define_insn "br_false"
 577   [(set (pc)
 578         (if_then_else (eq (match_operand:BI 0 "nvptx_register_operand" "R")
 579                           (const_int 0))
 580                       (label_ref (match_operand 1 "" ""))
 581                       (pc)))]
 582   ""
 583   "%J0\\tbra\\t%l1;"
 584   [(set_attr "predicable" "false")])
 585
 586 ;; unified conditional branch
 587 (define_insn "br_true_uni"
 588   [(set (pc) (if_then_else
 589         (ne (unspec:BI [(match_operand:BI 0 "nvptx_register_operand" "R")]
 590                        UNSPEC_BR_UNIFIED) (const_int 0))
 591         (label_ref (match_operand 1 "" "")) (pc)))]
 592   ""
 593   "%j0\\tbra.uni\\t%l1;"
 594   [(set_attr "predicable" "false")])
 595
 596 (define_insn "br_false_uni"
 597   [(set (pc) (if_then_else
 598         (eq (unspec:BI [(match_operand:BI 0 "nvptx_register_operand" "R")]
 599                        UNSPEC_BR_UNIFIED) (const_int 0))
 600         (label_ref (match_operand 1 "" "")) (pc)))]
 601   ""
 602   "%J0\\tbra.uni\\t%l1;"
 603   [(set_attr "predicable" "false")])
 604
 605 (define_expand "cbranch<mode>4"
 606   [(set (pc)
 607         (if_then_else (match_operator 0 "nvptx_comparison_operator"
 608                        [(match_operand:HSDIM 1 "nvptx_register_operand" "")
 609                         (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "")])
 610                       (label_ref (match_operand 3 "" ""))
 611                       (pc)))]
 612   ""
 613 {
 614   rtx t = nvptx_expand_compare (operands[0]);
 615   operands[0] = t;
 616   operands[1] = XEXP (t, 0);
 617   operands[2] = XEXP (t, 1);
 618 })
 619
 620 (define_expand "cbranch<mode>4"
 621   [(set (pc)
 622         (if_then_else (match_operator 0 "nvptx_float_comparison_operator"
 623                        [(match_operand:SDFM 1 "nvptx_register_operand" "")
 624                         (match_operand:SDFM 2 "nvptx_nonmemory_operand" "")])
 625                       (label_ref (match_operand 3 "" ""))
 626                       (pc)))]
 627   ""
 628 {
 629   rtx t = nvptx_expand_compare (operands[0]);
 630   operands[0] = t;
 631   operands[1] = XEXP (t, 0);
 632   operands[2] = XEXP (t, 1);
 633 })
 634
 635 (define_expand "cbranchbi4"
 636   [(set (pc)
 637         (if_then_else (match_operator 0 "predicate_operator"
 638                        [(match_operand:BI 1 "nvptx_register_operand" "")
 639                         (match_operand:BI 2 "const0_operand" "")])
 640                       (label_ref (match_operand 3 "" ""))
 641                       (pc)))]
 642   ""
 643   "")
 644
 645 ;; Conditional stores
 646
 647 (define_insn "setcc_from_bi"
 648   [(set (match_operand:SI 0 "nvptx_register_operand" "=R")
 649         (ne:SI (match_operand:BI 1 "nvptx_register_operand" "R")
 650                (const_int 0)))]
 651   ""
 652   "%.\\tselp%t0 %0,-1,0,%1;")
 653
 654 (define_insn "sel_true<mode>"
 655   [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
 656         (if_then_else:HSDIM
 657           (ne (match_operand:BI 1 "nvptx_register_operand" "R") (const_int 0))
 658           (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")
 659           (match_operand:HSDIM 3 "nvptx_nonmemory_operand" "Ri")))]
 660   ""
 661   "%.\\tselp%t0\\t%0, %2, %3, %1;")
 662
 663 (define_insn "sel_true<mode>"
 664   [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
 665         (if_then_else:SDFM
 666           (ne (match_operand:BI 1 "nvptx_register_operand" "R") (const_int 0))
 667           (match_operand:SDFM 2 "nvptx_nonmemory_operand" "RF")
 668           (match_operand:SDFM 3 "nvptx_nonmemory_operand" "RF")))]
 669   ""
 670   "%.\\tselp%t0\\t%0, %2, %3, %1;")
 671
 672 (define_insn "sel_false<mode>"
 673   [(set (match_operand:HSDIM 0 "nvptx_register_operand" "=R")
 674         (if_then_else:HSDIM
 675           (eq (match_operand:BI 1 "nvptx_register_operand" "R") (const_int 0))
 676           (match_operand:HSDIM 2 "nvptx_nonmemory_operand" "Ri")
 677           (match_operand:HSDIM 3 "nvptx_nonmemory_operand" "Ri")))]
 678   ""
 679   "%.\\tselp%t0\\t%0, %3, %2, %1;")
 680
 681 (define_insn "sel_false<mode>"
 682   [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
 683         (if_then_else:SDFM
 684           (eq (match_operand:BI 1 "nvptx_register_operand" "R") (const_int 0))
 685           (match_operand:SDFM 2 "nvptx_nonmemory_operand" "RF")
 686           (match_operand:SDFM 3 "nvptx_nonmemory_operand" "RF")))]
 687   ""
 688   "%.\\tselp%t0\\t%0, %3, %2, %1;")
 689
 690 (define_insn "setcc_int<mode>"
 691   [(set (match_operand:SI 0 "nvptx_register_operand" "=R")
 692         (match_operator:SI 1 "nvptx_comparison_operator"
 693           [(match_operand:HSDIM 2 "nvptx_register_operand" "R")
 694            (match_operand:HSDIM 3 "nvptx_nonmemory_operand" "Ri")]))]
 695   ""
 696   "%.\\tset%t0%c1\\t%0, %2, %3;")
 697
 698 (define_insn "setcc_int<mode>"
 699   [(set (match_operand:SI 0 "nvptx_register_operand" "=R")
 700         (match_operator:SI 1 "nvptx_float_comparison_operator"
 701            [(match_operand:SDFM 2 "nvptx_register_operand" "R")
 702             (match_operand:SDFM 3 "nvptx_nonmemory_operand" "RF")]))]
 703   ""
 704   "%.\\tset%t0%c1\\t%0, %2, %3;")
 705
 706 (define_insn "setcc_float<mode>"
 707   [(set (match_operand:SF 0 "nvptx_register_operand" "=R")
 708         (match_operator:SF 1 "nvptx_comparison_operator"
 709            [(match_operand:HSDIM 2 "nvptx_register_operand" "R")
 710             (match_operand:HSDIM 3 "nvptx_nonmemory_operand" "Ri")]))]
 711   ""
 712   "%.\\tset%t0%c1\\t%0, %2, %3;")
 713
 714 (define_insn "setcc_float<mode>"
 715   [(set (match_operand:SF 0 "nvptx_register_operand" "=R")
 716         (match_operator:SF 1 "nvptx_float_comparison_operator"
 717            [(match_operand:SDFM 2 "nvptx_register_operand" "R")
 718             (match_operand:SDFM 3 "nvptx_nonmemory_operand" "RF")]))]
 719   ""
 720   "%.\\tset%t0%c1\\t%0, %2, %3;")
 721
 722 (define_expand "cstorebi4"
 723   [(set (match_operand:SI 0 "nvptx_register_operand")
 724         (match_operator:SI 1 "ne_operator"
 725          [(match_operand:BI 2 "nvptx_register_operand")
 726           (match_operand:BI 3 "const0_operand")]))]
 727   ""
 728   "")
 729
 730 (define_expand "cstore<mode>4"
 731   [(set (match_operand:SI 0 "nvptx_register_operand")
 732         (match_operator:SI 1 "nvptx_comparison_operator"
 733          [(match_operand:HSDIM 2 "nvptx_register_operand")
 734           (match_operand:HSDIM 3 "nvptx_nonmemory_operand")]))]
 735   ""
 736   "")
 737
 738 (define_expand "cstore<mode>4"
 739   [(set (match_operand:SI 0 "nvptx_register_operand")
 740         (match_operator:SI 1 "nvptx_float_comparison_operator"
 741          [(match_operand:SDFM 2 "nvptx_register_operand")
 742           (match_operand:SDFM 3 "nvptx_nonmemory_operand")]))]
 743   ""
 744   "")
 745
 746 ;; Calls
 747
 748 (define_insn "call_insn_<mode>"
 749   [(match_parallel 2 "call_operation"
 750     [(call (mem:QI (match_operand:P 0 "call_insn_operand" "Rs"))
 751            (match_operand 1))])]
 752   ""
 753 {
 754   return nvptx_output_call_insn (insn, NULL_RTX, operands[0]);
 755 })
 756
 757 (define_insn "call_value_insn_<mode>"
 758   [(match_parallel 3 "call_operation"
 759     [(set (match_operand 0 "nvptx_register_operand" "=R")
 760           (call (mem:QI (match_operand:P 1 "call_insn_operand" "Rs"))
 761                 (match_operand 2)))])]
 762   ""
 763 {
 764   return nvptx_output_call_insn (insn, operands[0], operands[1]);
 765 })
 766
 767 (define_expand "call"
 768  [(match_operand 0 "" "")]
 769  ""
 770 {
 771   nvptx_expand_call (NULL_RTX, operands[0]);
 772   DONE;
 773 })
 774
 775 (define_expand "call_value"
 776   [(match_operand 0 "" "")
 777    (match_operand 1 "" "")]
 778  ""
 779 {
 780   nvptx_expand_call (operands[0], operands[1]);
 781   DONE;
 782 })
 783
 784 ;; Floating point arithmetic.
 785
 786 (define_insn "add<mode>3"
 787   [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
 788         (plus:SDFM (match_operand:SDFM 1 "nvptx_register_operand" "R")
 789                    (match_operand:SDFM 2 "nvptx_nonmemory_operand" "RF")))]
 790   ""
 791   "%.\\tadd%t0\\t%0, %1, %2;")
 792
 793 (define_insn "sub<mode>3"
 794   [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
 795         (minus:SDFM (match_operand:SDFM 1 "nvptx_register_operand" "R")
 796                     (match_operand:SDFM 2 "nvptx_register_operand" "R")))]
 797   ""
 798   "%.\\tsub%t0\\t%0, %1, %2;")
 799
 800 (define_insn "mul<mode>3"
 801   [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
 802         (mult:SDFM (match_operand:SDFM 1 "nvptx_register_operand" "R")
 803                    (match_operand:SDFM 2 "nvptx_nonmemory_operand" "RF")))]
 804   ""
 805   "%.\\tmul%t0\\t%0, %1, %2;")
 806
 807 (define_insn "fma<mode>4"
 808   [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
 809         (fma:SDFM (match_operand:SDFM 1 "nvptx_register_operand" "R")
 810                   (match_operand:SDFM 2 "nvptx_nonmemory_operand" "RF")
 811                   (match_operand:SDFM 3 "nvptx_nonmemory_operand" "RF")))]
 812   ""
 813   "%.\\tfma%#%t0\\t%0, %1, %2, %3;")
 814
 815 (define_insn "div<mode>3"
 816   [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
 817         (div:SDFM (match_operand:SDFM 1 "nvptx_register_operand" "R")
 818                   (match_operand:SDFM 2 "nvptx_nonmemory_operand" "RF")))]
 819   ""
 820   "%.\\tdiv%#%t0\\t%0, %1, %2;")
 821
 822 (define_insn "copysign<mode>3"
 823   [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
 824         (unspec:SDFM [(match_operand:SDFM 1 "nvptx_register_operand" "R")
 825                       (match_operand:SDFM 2 "nvptx_register_operand" "R")]
 826                       UNSPEC_COPYSIGN))]
 827   ""
 828   "%.\\tcopysign%t0\\t%0, %2, %1;")
 829
 830 (define_insn "smin<mode>3"
 831   [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
 832         (smin:SDFM (match_operand:SDFM 1 "nvptx_register_operand" "R")
 833                     (match_operand:SDFM 2 "nvptx_nonmemory_operand" "RF")))]
 834   ""
 835   "%.\\tmin%t0\\t%0, %1, %2;")
 836
 837 (define_insn "smax<mode>3"
 838   [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
 839         (smax:SDFM (match_operand:SDFM 1 "nvptx_register_operand" "R")
 840                     (match_operand:SDFM 2 "nvptx_nonmemory_operand" "RF")))]
 841   ""
 842   "%.\\tmax%t0\\t%0, %1, %2;")
 843
 844 (define_insn "abs<mode>2"
 845   [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
 846         (abs:SDFM (match_operand:SDFM 1 "nvptx_register_operand" "R")))]
 847   ""
 848   "%.\\tabs%t0\\t%0, %1;")
 849
 850 (define_insn "neg<mode>2"
 851   [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
 852         (neg:SDFM (match_operand:SDFM 1 "nvptx_register_operand" "R")))]
 853   ""
 854   "%.\\tneg%t0\\t%0, %1;")
 855
 856 (define_insn "sqrt<mode>2"
 857   [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
 858         (sqrt:SDFM (match_operand:SDFM 1 "nvptx_register_operand" "R")))]
 859   ""
 860   "%.\\tsqrt%#%t0\\t%0, %1;")
 861
 862 (define_expand "sincossf3"
 863   [(set (match_operand:SF 0 "nvptx_register_operand" "=R")
 864         (unspec:SF [(match_operand:SF 2 "nvptx_register_operand" "R")]
 865                    UNSPEC_COS))
 866    (set (match_operand:SF 1 "nvptx_register_operand" "=R")
 867         (unspec:SF [(match_dup 2)] UNSPEC_SIN))]
 868   "flag_unsafe_math_optimizations"
 869 {
 870   operands[2] = make_safe_from (operands[2], operands[0]);
 871 })
 872
 873 (define_insn "sinsf2"
 874   [(set (match_operand:SF 0 "nvptx_register_operand" "=R")
 875         (unspec:SF [(match_operand:SF 1 "nvptx_register_operand" "R")]
 876                    UNSPEC_SIN))]
 877   "flag_unsafe_math_optimizations"
 878   "%.\\tsin.approx%t0\\t%0, %1;")
 879
 880 (define_insn "cossf2"
 881   [(set (match_operand:SF 0 "nvptx_register_operand" "=R")
 882         (unspec:SF [(match_operand:SF 1 "nvptx_register_operand" "R")]
 883                    UNSPEC_COS))]
 884   "flag_unsafe_math_optimizations"
 885   "%.\\tcos.approx%t0\\t%0, %1;")
 886
 887 (define_insn "log2sf2"
 888   [(set (match_operand:SF 0 "nvptx_register_operand" "=R")
 889         (unspec:SF [(match_operand:SF 1 "nvptx_register_operand" "R")]
 890                    UNSPEC_LOG2))]
 891   "flag_unsafe_math_optimizations"
 892   "%.\\tlg2.approx%t0\\t%0, %1;")
 893
 894 (define_insn "exp2sf2"
 895   [(set (match_operand:SF 0 "nvptx_register_operand" "=R")
 896         (unspec:SF [(match_operand:SF 1 "nvptx_register_operand" "R")]
 897                    UNSPEC_EXP2))]
 898   "flag_unsafe_math_optimizations"
 899   "%.\\tex2.approx%t0\\t%0, %1;")
 900
 901 ;; Conversions involving floating point
 902
 903 (define_insn "extendsfdf2"
 904   [(set (match_operand:DF 0 "nvptx_register_operand" "=R")
 905         (float_extend:DF (match_operand:SF 1 "nvptx_register_operand" "R")))]
 906   ""
 907   "%.\\tcvt%t0%t1\\t%0, %1;")
 908
 909 (define_insn "truncdfsf2"
 910   [(set (match_operand:SF 0 "nvptx_register_operand" "=R")
 911         (float_truncate:SF (match_operand:DF 1 "nvptx_register_operand" "R")))]
 912   ""
 913   "%.\\tcvt%#%t0%t1\\t%0, %1;")
 914
 915 (define_insn "floatunssi<mode>2"
 916   [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
 917         (unsigned_float:SDFM (match_operand:SI 1 "nvptx_register_operand" "R")))]
 918   ""
 919   "%.\\tcvt%#%t0.u%T1\\t%0, %1;")
 920
 921 (define_insn "floatsi<mode>2"
 922   [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
 923         (float:SDFM (match_operand:SI 1 "nvptx_register_operand" "R")))]
 924   ""
 925   "%.\\tcvt%#%t0.s%T1\\t%0, %1;")
 926
 927 (define_insn "floatunsdi<mode>2"
 928   [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
 929         (unsigned_float:SDFM (match_operand:DI 1 "nvptx_register_operand" "R")))]
 930   ""
 931   "%.\\tcvt%#%t0.u%T1\\t%0, %1;")
 932
 933 (define_insn "floatdi<mode>2"
 934   [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
 935         (float:SDFM (match_operand:DI 1 "nvptx_register_operand" "R")))]
 936   ""
 937   "%.\\tcvt%#%t0.s%T1\\t%0, %1;")
 938
 939 (define_insn "fixuns_trunc<mode>si2"
 940   [(set (match_operand:SI 0 "nvptx_register_operand" "=R")
 941         (unsigned_fix:SI (match_operand:SDFM 1 "nvptx_register_operand" "R")))]
 942   ""
 943   "%.\\tcvt.rzi.u%T0%t1\\t%0, %1;")
 944
 945 (define_insn "fix_trunc<mode>si2"
 946   [(set (match_operand:SI 0 "nvptx_register_operand" "=R")
 947         (fix:SI (match_operand:SDFM 1 "nvptx_register_operand" "R")))]
 948   ""
 949   "%.\\tcvt.rzi.s%T0%t1\\t%0, %1;")
 950
 951 (define_insn "fixuns_trunc<mode>di2"
 952   [(set (match_operand:DI 0 "nvptx_register_operand" "=R")
 953         (unsigned_fix:DI (match_operand:SDFM 1 "nvptx_register_operand" "R")))]
 954   ""
 955   "%.\\tcvt.rzi.u%T0%t1\\t%0, %1;")
 956
 957 (define_insn "fix_trunc<mode>di2"
 958   [(set (match_operand:DI 0 "nvptx_register_operand" "=R")
 959         (fix:DI (match_operand:SDFM 1 "nvptx_register_operand" "R")))]
 960   ""
 961   "%.\\tcvt.rzi.s%T0%t1\\t%0, %1;")
 962
 963 (define_int_iterator FPINT [UNSPEC_FPINT_FLOOR UNSPEC_FPINT_BTRUNC
 964                             UNSPEC_FPINT_CEIL UNSPEC_FPINT_NEARBYINT])
 965 (define_int_attr fpint_name [(UNSPEC_FPINT_FLOOR "floor")
 966                              (UNSPEC_FPINT_BTRUNC "btrunc")
 967                              (UNSPEC_FPINT_CEIL "ceil")
 968                              (UNSPEC_FPINT_NEARBYINT "nearbyint")])
 969 (define_int_attr fpint_roundingmode [(UNSPEC_FPINT_FLOOR ".rmi")
 970                                      (UNSPEC_FPINT_BTRUNC ".rzi")
 971                                      (UNSPEC_FPINT_CEIL ".rpi")
 972                                      (UNSPEC_FPINT_NEARBYINT "%#i")])
 973
 974 (define_insn "<FPINT:fpint_name><SDFM:mode>2"
 975   [(set (match_operand:SDFM 0 "nvptx_register_operand" "=R")
 976         (unspec:SDFM [(match_operand:SDFM 1 "nvptx_register_operand" "R")]
 977                      FPINT))]
 978   ""
 979   "%.\\tcvt<FPINT:fpint_roundingmode>%t0%t1\\t%0, %1;")
 980
 981 (define_int_iterator FPINT2 [UNSPEC_FPINT_FLOOR UNSPEC_FPINT_CEIL])
 982 (define_int_attr fpint2_name [(UNSPEC_FPINT_FLOOR "lfloor")
 983                              (UNSPEC_FPINT_CEIL "lceil")])
 984 (define_int_attr fpint2_roundingmode [(UNSPEC_FPINT_FLOOR ".rmi")
 985                                      (UNSPEC_FPINT_CEIL ".rpi")])
 986
 987 (define_insn "<FPINT2:fpint2_name><SDFM:mode><SDIM:mode>2"
 988   [(set (match_operand:SDIM 0 "nvptx_register_operand" "=R")
 989         (unspec:SDIM [(match_operand:SDFM 1 "nvptx_register_operand" "R")]
 990                      FPINT2))]
 991   ""
 992   "%.\\tcvt<FPINT2:fpint2_roundingmode>.s%T0%t1\\t%0, %1;")
 993
 994 ;; Miscellaneous
 995
 996 (define_insn "nop"
 997   [(const_int 0)]
 998   ""
 999   "")
1000
1001 (define_insn "exit"
1002   [(const_int 1)]
1003   ""
1004   "exit;")
1005
1006 (define_insn "fake_nop"
1007   [(const_int 2)]
1008   ""
1009   "{
1010      .reg .u32 %%nop_src;
1011      .reg .u32 %%nop_dst;
1012      mov.u32 %%nop_dst, %%nop_src;
1013    }")
1014
1015 (define_insn "return"
1016   [(return)]
1017   ""
1018 {
1019   return nvptx_output_return ();
1020 }
1021   [(set_attr "predicable" "false")])
1022
1023 (define_expand "epilogue"
1024   [(clobber (const_int 0))]
1025   ""
1026 {
1027   if (TARGET_SOFT_STACK)
1028     emit_insn (gen_set_softstack_insn (gen_rtx_REG (Pmode,
1029                                                     SOFTSTACK_PREV_REGNUM)));
1030   emit_jump_insn (gen_return ());
1031   DONE;
1032 })
1033
1034 (define_expand "nonlocal_goto"
1035   [(match_operand 0 "" "")
1036    (match_operand 1 "" "")
1037    (match_operand 2 "" "")
1038    (match_operand 3 "" "")]
1039   ""
1040 {
1041   sorry ("target cannot support nonlocal goto.");
1042   emit_insn (gen_nop ());
1043   DONE;
1044 })
1045
1046 (define_expand "nonlocal_goto_receiver"
1047   [(const_int 0)]
1048   ""
1049 {
1050   sorry ("target cannot support nonlocal goto.");
1051 })
1052
1053 (define_expand "allocate_stack"
1054   [(match_operand 0 "nvptx_register_operand")
1055    (match_operand 1 "nvptx_register_operand")]
1056   ""
1057 {
1058   if (TARGET_SOFT_STACK)
1059     {
1060       emit_move_insn (stack_pointer_rtx,
1061                       gen_rtx_MINUS (Pmode, stack_pointer_rtx, operands[1]));
1062       emit_insn (gen_set_softstack_insn (stack_pointer_rtx));
1063       emit_move_insn (operands[0], virtual_stack_dynamic_rtx);
1064       DONE;
1065     }
1066   /* The ptx documentation specifies an alloca intrinsic (for 32 bit
1067      only)  but notes it is not implemented.  The assembler emits a
1068      confused error message.  Issue a blunt one now instead.  */
1069   sorry ("target cannot support alloca.");
1070   emit_insn (gen_nop ());
1071   DONE;
1072 })
1073
1074 (define_insn "set_softstack_<mode>"
1075   [(unspec [(match_operand:P 0 "nvptx_register_operand" "R")]
1076            UNSPEC_SET_SOFTSTACK)]
1077   "TARGET_SOFT_STACK"
1078 {
1079   return nvptx_output_set_softstack (REGNO (operands[0]));
1080 })
1081
1082 (define_expand "restore_stack_block"
1083   [(match_operand 0 "register_operand" "")
1084    (match_operand 1 "register_operand" "")]
1085   ""
1086 {
1087   if (TARGET_SOFT_STACK)
1088     {
1089       emit_move_insn (operands[0], operands[1]);
1090       emit_insn (gen_set_softstack_insn (operands[0]));
1091     }
1092   DONE;
1093 })
1094
1095 (define_expand "restore_stack_function"
1096   [(match_operand 0 "register_operand" "")
1097    (match_operand 1 "register_operand" "")]
1098   ""
1099 {
1100   DONE;
1101 })
1102
1103 (define_insn "trap"
1104   [(trap_if (const_int 1) (const_int 0))]
1105   ""
1106   "trap; exit;")
1107
1108 (define_insn "trap_if_true"
1109   [(trap_if (ne (match_operand:BI 0 "nvptx_register_operand" "R")
1110                 (const_int 0))
1111             (const_int 0))]
1112   ""
1113   "%j0 trap; %j0 exit;"
1114   [(set_attr "predicable" "false")])
1115
1116 (define_insn "trap_if_false"
1117   [(trap_if (eq (match_operand:BI 0 "nvptx_register_operand" "R")
1118                 (const_int 0))
1119             (const_int 0))]
1120   ""
1121   "%J0 trap; %J0 exit;"
1122   [(set_attr "predicable" "false")])
1123
1124 (define_expand "ctrap<mode>4"
1125   [(trap_if (match_operator 0 "nvptx_comparison_operator"
1126                             [(match_operand:SDIM 1 "nvptx_register_operand")
1127                              (match_operand:SDIM 2 "nvptx_nonmemory_operand")])
1128             (match_operand 3 "const0_operand"))]
1129   ""
1130 {
1131   rtx t = nvptx_expand_compare (operands[0]);
1132   emit_insn (gen_trap_if_true (t));
1133   DONE;
1134 })
1135
1136 (define_insn "oacc_dim_size"
1137   [(set (match_operand:SI 0 "nvptx_register_operand" "")
1138         (unspec:SI [(match_operand:SI 1 "const_int_operand" "")]
1139                    UNSPEC_DIM_SIZE))]
1140   ""
1141 {
1142   static const char *const asms[] =
1143 { /* Must match oacc_loop_levels ordering.  */
1144   "%.\\tmov.u32\\t%0, %%nctaid.x;",     /* gang */
1145   "%.\\tmov.u32\\t%0, %%ntid.y;",       /* worker */
1146   "%.\\tmov.u32\\t%0, %%ntid.x;",       /* vector */
1147 };
1148   return asms[INTVAL (operands[1])];
1149 })
1150
1151 (define_insn "oacc_dim_pos"
1152   [(set (match_operand:SI 0 "nvptx_register_operand" "")
1153         (unspec_volatile:SI [(match_operand:SI 1 "const_int_operand" "")]
1154                             UNSPECV_DIM_POS))]
1155   ""
1156 {
1157   static const char *const asms[] =
1158 { /* Must match oacc_loop_levels ordering.  */
1159   "%.\\tmov.u32\\t%0, %%ctaid.x;",      /* gang */
1160   "%.\\tmov.u32\\t%0, %%tid.y;",        /* worker */
1161   "%.\\tmov.u32\\t%0, %%tid.x;",        /* vector */
1162 };
1163   return asms[INTVAL (operands[1])];
1164 })
1165
1166 (define_insn "nvptx_fork"
1167   [(unspec_volatile:SI [(match_operand:SI 0 "const_int_operand" "")]
1168                        UNSPECV_FORK)]
1169   ""
1170   "// fork %0;"
1171   [(set_attr "predicable" "false")])
1172
1173 (define_insn "nvptx_forked"
1174   [(unspec_volatile:SI [(match_operand:SI 0 "const_int_operand" "")]
1175                        UNSPECV_FORKED)]
1176   ""
1177   "// forked %0;"
1178   [(set_attr "predicable" "false")])
1179
1180 (define_insn "nvptx_joining"
1181   [(unspec_volatile:SI [(match_operand:SI 0 "const_int_operand" "")]
1182                        UNSPECV_JOINING)]
1183   ""
1184   "// joining %0;"
1185   [(set_attr "predicable" "false")])
1186
1187 (define_insn "nvptx_join"
1188   [(unspec_volatile:SI [(match_operand:SI 0 "const_int_operand" "")]
1189                        UNSPECV_JOIN)]
1190   ""
1191   "// join %0;"
1192   [(set_attr "predicable" "false")])
1193
1194 (define_expand "oacc_fork"
1195   [(set (match_operand:SI 0 "nvptx_nonmemory_operand" "")
1196         (match_operand:SI 1 "general_operand" ""))
1197    (unspec_volatile:SI [(match_operand:SI 2 "const_int_operand" "")]
1198                         UNSPECV_FORKED)]
1199   ""
1200 {
1201   if (operands[0] != const0_rtx)
1202     emit_move_insn (operands[0], operands[1]);
1203   nvptx_expand_oacc_fork (INTVAL (operands[2]));
1204   DONE;
1205 })
1206
1207 (define_expand "oacc_join"
1208   [(set (match_operand:SI 0 "nvptx_nonmemory_operand" "")
1209         (match_operand:SI 1 "general_operand" ""))
1210    (unspec_volatile:SI [(match_operand:SI 2 "const_int_operand" "")]
1211                         UNSPECV_JOIN)]
1212   ""
1213 {
1214   if (operands[0] != const0_rtx)
1215     emit_move_insn (operands[0], operands[1]);
1216   nvptx_expand_oacc_join (INTVAL (operands[2]));
1217   DONE;
1218 })
1219
1220 ;; only 32-bit shuffles exist.
1221 (define_insn "nvptx_shuffle<mode>"
1222   [(set (match_operand:BITS 0 "nvptx_register_operand" "=R")
1223         (unspec:BITS
1224                 [(match_operand:BITS 1 "nvptx_register_operand" "R")
1225                  (match_operand:SI 2 "nvptx_nonmemory_operand" "Ri")
1226                  (match_operand:SI 3 "const_int_operand" "n")]
1227                   UNSPEC_SHUFFLE))]
1228   ""
1229   "%.\\tshfl%S3.b32\\t%0, %1, %2, 31;")
1230
1231 (define_insn "nvptx_vote_ballot"
1232   [(set (match_operand:SI 0 "nvptx_register_operand" "=R")
1233         (unspec:SI [(match_operand:BI 1 "nvptx_register_operand" "R")]
1234                    UNSPEC_VOTE_BALLOT))]
1235   ""
1236   "%.\\tvote.ballot.b32\\t%0, %1;")
1237
1238 ;; Patterns for OpenMP SIMD-via-SIMT lowering
1239
1240 (define_insn "omp_simt_enter_<mode>"
1241   [(set (match_operand:P 0 "nvptx_register_operand" "=R")
1242         (unspec_volatile:P [(match_operand:P 1 "nvptx_nonmemory_operand" "Ri")
1243                             (match_operand:P 2 "nvptx_nonmemory_operand" "Ri")]
1244                            UNSPECV_SIMT_ENTER))]
1245   ""
1246 {
1247   return nvptx_output_simt_enter (operands[0], operands[1], operands[2]);
1248 })
1249
1250 (define_expand "omp_simt_enter"
1251   [(match_operand 0 "nvptx_register_operand" "=R")
1252    (match_operand 1 "nvptx_nonmemory_operand" "Ri")
1253    (match_operand 2 "const_int_operand" "n")]
1254   ""
1255 {
1256   if (!CONST_INT_P (operands[1]))
1257     cfun->machine->simt_stack_size = HOST_WIDE_INT_M1U;
1258   else
1259     cfun->machine->simt_stack_size = MAX (UINTVAL (operands[1]),
1260                                           cfun->machine->simt_stack_size);
1261   cfun->machine->simt_stack_align = MAX (UINTVAL (operands[2]),
1262                                          cfun->machine->simt_stack_align);
1263   cfun->machine->has_simtreg = true;
1264   gcc_assert (GET_MODE (operands[0]) == Pmode);
1265   if (GET_MODE (operands[0]) == DImode)
1266     emit_insn (gen_omp_simt_enter_di (operands[0], operands[1], operands[2]));
1267   else if (GET_MODE (operands[0]) == SImode)
1268     emit_insn (gen_omp_simt_enter_si (operands[0], operands[1], operands[2]));
1269   else
1270     gcc_unreachable ();
1271   DONE;
1272 })
1273
1274 (define_expand "omp_simt_exit"
1275   [(match_operand 0 "nvptx_register_operand" "R")]
1276   ""
1277 {
1278   gcc_assert (GET_MODE (operands[0]) == Pmode);
1279   if (GET_MODE (operands[0]) == DImode)
1280     emit_insn (gen_omp_simt_exit_di (operands[0]));
1281   else if (GET_MODE (operands[0]) == SImode)
1282     emit_insn (gen_omp_simt_exit_si (operands[0]));
1283   else
1284     gcc_unreachable ();
1285   DONE;
1286 })
1287
1288 (define_insn "omp_simt_exit_<mode>"
1289   [(unspec_volatile [(match_operand:P 0 "nvptx_register_operand" "R")]
1290                     UNSPECV_SIMT_EXIT)]
1291   ""
1292 {
1293   return nvptx_output_simt_exit (operands[0]);
1294 })
1295
1296 ;; Implement IFN_GOMP_SIMT_LANE: set operand 0 to lane index
1297 (define_insn "omp_simt_lane"
1298   [(set (match_operand:SI 0 "nvptx_register_operand" "")
1299         (unspec:SI [(const_int 0)] UNSPEC_LANEID))]
1300   ""
1301   "%.\\tmov.u32\\t%0, %%laneid;")
1302
1303 ;; Implement IFN_GOMP_SIMT_ORDERED: copy operand 1 to operand 0 and
1304 ;; place a compiler barrier to disallow unrolling/peeling the containing loop
1305 (define_expand "omp_simt_ordered"
1306   [(match_operand:SI 0 "nvptx_register_operand" "=R")
1307    (match_operand:SI 1 "nvptx_register_operand" "R")]
1308   ""
1309 {
1310   emit_move_insn (operands[0], operands[1]);
1311   emit_insn (gen_nvptx_nounroll ());
1312   DONE;
1313 })
1314
1315 ;; Implement IFN_GOMP_SIMT_XCHG_BFLY: perform a "butterfly" exchange
1316 ;; across lanes
1317 (define_expand "omp_simt_xchg_bfly"
1318   [(match_operand 0 "nvptx_register_operand" "=R")
1319    (match_operand 1 "nvptx_register_operand" "R")
1320    (match_operand:SI 2 "nvptx_nonmemory_operand" "Ri")]
1321   ""
1322 {
1323   emit_insn (nvptx_gen_shuffle (operands[0], operands[1], operands[2],
1324                                 SHUFFLE_BFLY));
1325   DONE;
1326 })
1327
1328 ;; Implement IFN_GOMP_SIMT_XCHG_IDX: broadcast value in operand 1
1329 ;; from lane given by index in operand 2 to operand 0 in all lanes
1330 (define_expand "omp_simt_xchg_idx"
1331   [(match_operand 0 "nvptx_register_operand" "=R")
1332    (match_operand 1 "nvptx_register_operand" "R")
1333    (match_operand:SI 2 "nvptx_nonmemory_operand" "Ri")]
1334   ""
1335 {
1336   emit_insn (nvptx_gen_shuffle (operands[0], operands[1], operands[2],
1337                                 SHUFFLE_IDX));
1338   DONE;
1339 })
1340
1341 ;; Implement IFN_GOMP_SIMT_VOTE_ANY:
1342 ;; set operand 0 to zero iff all lanes supply zero in operand 1
1343 (define_expand "omp_simt_vote_any"
1344   [(match_operand:SI 0 "nvptx_register_operand" "=R")
1345    (match_operand:SI 1 "nvptx_register_operand" "R")]
1346   ""
1347 {
1348   rtx pred = gen_reg_rtx (BImode);
1349   emit_move_insn (pred, gen_rtx_NE (BImode, operands[1], const0_rtx));
1350   emit_insn (gen_nvptx_vote_ballot (operands[0], pred));
1351   DONE;
1352 })
1353
1354 ;; Implement IFN_GOMP_SIMT_LAST_LANE:
1355 ;; set operand 0 to the lowest lane index that passed non-zero in operand 1
1356 (define_expand "omp_simt_last_lane"
1357   [(match_operand:SI 0 "nvptx_register_operand" "=R")
1358    (match_operand:SI 1 "nvptx_register_operand" "R")]
1359   ""
1360 {
1361   rtx pred = gen_reg_rtx (BImode);
1362   rtx tmp = gen_reg_rtx (SImode);
1363   emit_move_insn (pred, gen_rtx_NE (BImode, operands[1], const0_rtx));
1364   emit_insn (gen_nvptx_vote_ballot (tmp, pred));
1365   emit_insn (gen_ctzsi2 (operands[0], tmp));
1366   DONE;
1367 })
1368
1369 ;; extract parts of a 64 bit object into 2 32-bit ints
1370 (define_insn "unpack<mode>si2"
1371   [(set (match_operand:SI 0 "nvptx_register_operand" "=R")
1372         (unspec:SI [(match_operand:BITD 2 "nvptx_register_operand" "R")
1373                     (const_int 0)] UNSPEC_BIT_CONV))
1374    (set (match_operand:SI 1 "nvptx_register_operand" "=R")
1375         (unspec:SI [(match_dup 2) (const_int 1)] UNSPEC_BIT_CONV))]
1376   ""
1377   "%.\\tmov.b64\\t{%0,%1}, %2;")
1378
1379 ;; pack 2 32-bit ints into a 64 bit object
1380 (define_insn "packsi<mode>2"
1381   [(set (match_operand:BITD 0 "nvptx_register_operand" "=R")
1382         (unspec:BITD [(match_operand:SI 1 "nvptx_register_operand" "R")
1383                       (match_operand:SI 2 "nvptx_register_operand" "R")]
1384                     UNSPEC_BIT_CONV))]
1385   ""
1386   "%.\\tmov.b64\\t%0, {%1,%2};")
1387
1388 ;; Atomic insns.
1389
1390 (define_expand "atomic_compare_and_swap<mode>"
1391   [(match_operand:SI 0 "nvptx_register_operand")        ;; bool success output
1392    (match_operand:SDIM 1 "nvptx_register_operand")      ;; oldval output
1393    (match_operand:SDIM 2 "memory_operand")              ;; memory
1394    (match_operand:SDIM 3 "nvptx_register_operand")      ;; expected input
1395    (match_operand:SDIM 4 "nvptx_register_operand")      ;; newval input
1396    (match_operand:SI 5 "const_int_operand")             ;; is_weak
1397    (match_operand:SI 6 "const_int_operand")             ;; success model
1398    (match_operand:SI 7 "const_int_operand")]            ;; failure model
1399   ""
1400 {
1401   emit_insn (gen_atomic_compare_and_swap<mode>_1
1402     (operands[1], operands[2], operands[3], operands[4], operands[6]));
1403
1404   rtx cond = gen_reg_rtx (BImode);
1405   emit_move_insn (cond, gen_rtx_EQ (BImode, operands[1], operands[3]));
1406   emit_insn (gen_sel_truesi (operands[0], cond, GEN_INT (1), GEN_INT (0)));
1407   DONE;
1408 })
1409
1410 (define_insn "atomic_compare_and_swap<mode>_1"
1411   [(set (match_operand:SDIM 0 "nvptx_register_operand" "=R")
1412         (unspec_volatile:SDIM
1413           [(match_operand:SDIM 1 "memory_operand" "+m")
1414            (match_operand:SDIM 2 "nvptx_nonmemory_operand" "Ri")
1415            (match_operand:SDIM 3 "nvptx_nonmemory_operand" "Ri")
1416            (match_operand:SI 4 "const_int_operand")]
1417           UNSPECV_CAS))
1418    (set (match_dup 1)
1419         (unspec_volatile:SDIM [(const_int 0)] UNSPECV_CAS))]
1420   ""
1421   "%.\\tatom%A1.cas.b%T0\\t%0, %1, %2, %3;"
1422   [(set_attr "atomic" "true")])
1423
1424 (define_insn "atomic_exchange<mode>"
1425   [(set (match_operand:SDIM 0 "nvptx_register_operand" "=R")    ;; output
1426         (unspec_volatile:SDIM
1427           [(match_operand:SDIM 1 "memory_operand" "+m")         ;; memory
1428            (match_operand:SI 3 "const_int_operand")]            ;; model
1429           UNSPECV_XCHG))
1430    (set (match_dup 1)
1431         (match_operand:SDIM 2 "nvptx_nonmemory_operand" "Ri"))] ;; input
1432   ""
1433   "%.\\tatom%A1.exch.b%T0\\t%0, %1, %2;"
1434   [(set_attr "atomic" "true")])
1435
1436 (define_insn "atomic_fetch_add<mode>"
1437   [(set (match_operand:SDIM 1 "memory_operand" "+m")
1438         (unspec_volatile:SDIM
1439           [(plus:SDIM (match_dup 1)
1440                       (match_operand:SDIM 2 "nvptx_nonmemory_operand" "Ri"))
1441            (match_operand:SI 3 "const_int_operand")]            ;; model
1442           UNSPECV_LOCK))
1443    (set (match_operand:SDIM 0 "nvptx_register_operand" "=R")
1444         (match_dup 1))]
1445   ""
1446   "%.\\tatom%A1.add%t0\\t%0, %1, %2;"
1447   [(set_attr "atomic" "true")])
1448
1449 (define_insn "atomic_fetch_addsf"
1450   [(set (match_operand:SF 1 "memory_operand" "+m")
1451         (unspec_volatile:SF
1452          [(plus:SF (match_dup 1)
1453                    (match_operand:SF 2 "nvptx_nonmemory_operand" "RF"))
1454            (match_operand:SI 3 "const_int_operand")]            ;; model
1455           UNSPECV_LOCK))
1456    (set (match_operand:SF 0 "nvptx_register_operand" "=R")
1457         (match_dup 1))]
1458   ""
1459   "%.\\tatom%A1.add%t0\\t%0, %1, %2;"
1460   [(set_attr "atomic" "true")])
1461
1462 (define_code_iterator any_logic [and ior xor])
1463 (define_code_attr logic [(and "and") (ior "or") (xor "xor")])
1464
1465 (define_insn "atomic_fetch_<logic><mode>"
1466   [(set (match_operand:SDIM 1 "memory_operand" "+m")
1467         (unspec_volatile:SDIM
1468           [(any_logic:SDIM (match_dup 1)
1469                            (match_operand:SDIM 2 "nvptx_nonmemory_operand" "Ri"))
1470            (match_operand:SI 3 "const_int_operand")]            ;; model
1471           UNSPECV_LOCK))
1472    (set (match_operand:SDIM 0 "nvptx_register_operand" "=R")
1473         (match_dup 1))]
1474   "<MODE>mode == SImode || TARGET_SM35"
1475   "%.\\tatom%A1.b%T0.<logic>\\t%0, %1, %2;"
1476   [(set_attr "atomic" "true")])
1477
1478 (define_insn "nvptx_barsync"
1479   [(unspec_volatile [(match_operand:SI 0 "nvptx_nonmemory_operand" "Ri")
1480                      (match_operand:SI 1 "const_int_operand")]
1481                     UNSPECV_BARSYNC)]
1482   ""
1483   {
1484     if (INTVAL (operands[1]) == 0)
1485       return "\\tbar.sync\\t%0;";
1486     else
1487       return "\\tbar.sync\\t%0, %1;";
1488   }
1489   [(set_attr "predicable" "false")])
1490
1491 (define_expand "memory_barrier"
1492   [(set (match_dup 0)
1493         (unspec_volatile:BLK [(match_dup 0)] UNSPECV_MEMBAR))]
1494   ""
1495 {
1496   operands[0] = gen_rtx_MEM (BLKmode, gen_rtx_SCRATCH (Pmode));
1497   MEM_VOLATILE_P (operands[0]) = 1;
1498 })
1499
1500 ;; Ptx defines the memory barriers membar.cta, membar.gl and membar.sys
1501 ;; (corresponding to cuda functions threadfence_block, threadfence and
1502 ;; threadfence_system).  For the insn memory_barrier we use membar.sys.  This
1503 ;; may be overconservative, but before using membar.gl instead we'll need to
1504 ;; explain in detail why it's safe to use.  For now, use membar.sys.
1505 (define_insn "*memory_barrier"
1506   [(set (match_operand:BLK 0 "" "")
1507         (unspec_volatile:BLK [(match_dup 0)] UNSPECV_MEMBAR))]
1508   ""
1509   "\\tmembar.sys;"
1510   [(set_attr "predicable" "false")])
1511
1512 (define_expand "nvptx_membar_cta"
1513   [(set (match_dup 0)
1514         (unspec_volatile:BLK [(match_dup 0)] UNSPECV_MEMBAR_CTA))]
1515   ""
1516 {
1517   operands[0] = gen_rtx_MEM (BLKmode, gen_rtx_SCRATCH (Pmode));
1518   MEM_VOLATILE_P (operands[0]) = 1;
1519 })
1520
1521 (define_insn "*nvptx_membar_cta"
1522   [(set (match_operand:BLK 0 "" "")
1523         (unspec_volatile:BLK [(match_dup 0)] UNSPECV_MEMBAR_CTA))]
1524   ""
1525   "\\tmembar.cta;"
1526   [(set_attr "predicable" "false")])
1527
1528 (define_insn "nvptx_nounroll"
1529   [(unspec_volatile [(const_int 0)] UNSPECV_NOUNROLL)]
1530   ""
1531   "\\t.pragma \\\"nounroll\\\";"
1532   [(set_attr "predicable" "false")])
1533
1534 (define_insn "nvptx_red_partition"
1535   [(set (match_operand:DI 0 "nonimmediate_operand" "=R")
1536         (unspec_volatile:DI [(match_operand:DI 1 "const_int_operand")]
1537          UNSPECV_RED_PART))]
1538   ""
1539   {
1540     return nvptx_output_red_partition (operands[0], operands[1]);
1541   }
1542   [(set_attr "predicable" "false")])