Home | History | Annotate | Download | only in NVPTX 
      1 //===- NVPTXIntrinsics.td - PTX Intrinsics Instructions -------*- tblgen -*-==//
      2 //
      3 //                     The LLVM Compiler Infrastructure
      4 //
      5 // This file is distributed under the University of Illinois Open Source
      6 // License. See LICENSE.TXT for details.
      7 //
      8 //===----------------------------------------------------------------------===//
      9 
     10 def immFloat0 : PatLeaf<(fpimm), [{
     11     float f = (float)N->getValueAPF().convertToFloat();
     12     return (f==0.0f);
     13 }]>;
     14 
     15 def immFloat1 : PatLeaf<(fpimm), [{
     16     float f = (float)N->getValueAPF().convertToFloat();
     17     return (f==1.0f);
     18 }]>;
     19 
     20 def immDouble0 : PatLeaf<(fpimm), [{
     21     double d = (double)N->getValueAPF().convertToDouble();
     22     return (d==0.0);
     23 }]>;
     24 
     25 def immDouble1 : PatLeaf<(fpimm), [{
     26     double d = (double)N->getValueAPF().convertToDouble();
     27     return (d==1.0);
     28 }]>;
     29 
     30 
     31 
     32 //-----------------------------------
     33 // Synchronization Functions
     34 //-----------------------------------
     35 def INT_CUDA_SYNCTHREADS : NVPTXInst<(outs), (ins),
     36                   "bar.sync \t0;",
     37       [(int_cuda_syncthreads)]>;
     38 def INT_BARRIER0 : NVPTXInst<(outs), (ins),
     39                   "bar.sync \t0;",
     40       [(int_nvvm_barrier0)]>;
     41 def INT_BARRIER0_POPC : NVPTXInst<(outs Int32Regs:$dst), (ins Int32Regs:$pred),
     42   !strconcat("{{ \n\t",
     43       !strconcat(".reg .pred \t%p1; \n\t",
     44       !strconcat("setp.ne.u32 \t%p1, $pred, 0; \n\t",
     45       !strconcat("bar.red.popc.u32 \t$dst, 0, %p1; \n\t",
     46         !strconcat("}}", ""))))),
     47       [(set Int32Regs:$dst, (int_nvvm_barrier0_popc Int32Regs:$pred))]>;
     48 def INT_BARRIER0_AND : NVPTXInst<(outs Int32Regs:$dst), (ins Int32Regs:$pred),
     49   !strconcat("{{ \n\t",
     50       !strconcat(".reg .pred \t%p1; \n\t",
     51       !strconcat(".reg .pred \t%p2; \n\t",
     52       !strconcat("setp.ne.u32 \t%p1, $pred, 0; \n\t",
     53       !strconcat("bar.red.and.pred \t%p2, 0, %p1; \n\t",
     54       !strconcat("selp.u32 \t$dst, 1, 0, %p2; \n\t",
     55         !strconcat("}}", ""))))))),
     56       [(set Int32Regs:$dst, (int_nvvm_barrier0_and Int32Regs:$pred))]>;
     57 def INT_BARRIER0_OR : NVPTXInst<(outs Int32Regs:$dst), (ins Int32Regs:$pred),
     58   !strconcat("{{ \n\t",
     59       !strconcat(".reg .pred \t%p1; \n\t",
     60       !strconcat(".reg .pred \t%p2; \n\t",
     61       !strconcat("setp.ne.u32 \t%p1, $pred, 0; \n\t",
     62       !strconcat("bar.red.or.pred \t%p2, 0, %p1; \n\t",
     63       !strconcat("selp.u32 \t$dst, 1, 0, %p2; \n\t",
     64         !strconcat("}}", ""))))))),
     65       [(set Int32Regs:$dst, (int_nvvm_barrier0_or Int32Regs:$pred))]>;
     66 
     67 
     68 //-----------------------------------
     69 // Explicit Memory Fence Functions
     70 //-----------------------------------
     71 class MEMBAR<string StrOp, Intrinsic IntOP> :
     72               NVPTXInst<(outs), (ins),
     73             StrOp, [(IntOP)]>;
     74 
     75 def INT_MEMBAR_CTA : MEMBAR<"membar.cta;", int_nvvm_membar_cta>;
     76 def INT_MEMBAR_GL  : MEMBAR<"membar.gl;",  int_nvvm_membar_gl>;
     77 def INT_MEMBAR_SYS : MEMBAR<"membar.sys;", int_nvvm_membar_sys>;
     78 
     79 
     80 //-----------------------------------
     81 // Math Functions
     82 //-----------------------------------
     83 
     84 // Map min(1.0, max(0.0, x)) to sat(x)
     85 multiclass SAT<NVPTXRegClass regclass, Operand fimm, Intrinsic IntMinOp,
     86   Intrinsic IntMaxOp, PatLeaf f0, PatLeaf f1, string OpStr> {
     87 
     88    // fmin(1.0, fmax(0.0, x)) => sat(x)
     89    def SAT11 : NVPTXInst<(outs regclass:$dst),
     90      (ins fimm:$srcf0, fimm:$srcf1, regclass:$src),
     91            OpStr,
     92      [(set regclass:$dst, (IntMinOp f1:$srcf0 ,
     93        (IntMaxOp f0:$srcf1, regclass:$src)))]>;
     94 
     95    // fmin(1.0, fmax(x, 0.0)) => sat(x)
     96    def SAT12 : NVPTXInst<(outs regclass:$dst),
     97      (ins fimm:$srcf0, fimm:$srcf1, regclass:$src),
     98            OpStr,
     99      [(set regclass:$dst, (IntMinOp f1:$srcf0 ,
    100        (IntMaxOp regclass:$src, f0:$srcf1)))]>;
    101 
    102    // fmin(fmax(0.0, x), 1.0) => sat(x)
    103    def SAT13 : NVPTXInst<(outs regclass:$dst),
    104      (ins fimm:$srcf0, fimm:$srcf1, regclass:$src),
    105            OpStr,
    106      [(set regclass:$dst, (IntMinOp
    107        (IntMaxOp f0:$srcf0, regclass:$src), f1:$srcf1))]>;
    108 
    109    // fmin(fmax(x, 0.0), 1.0) => sat(x)
    110    def SAT14 : NVPTXInst<(outs regclass:$dst),
    111      (ins fimm:$srcf0, fimm:$srcf1, regclass:$src),
    112          OpStr,
    113      [(set regclass:$dst, (IntMinOp
    114        (IntMaxOp regclass:$src, f0:$srcf0), f1:$srcf1))]>;
    115 
    116 }
    117 // Note that max(0.0, min(x, 1.0)) cannot be mapped to sat(x) because when x
    118 // is NaN
    119 // max(0.0, min(x, 1.0)) is 1.0 while sat(x) is 0.
    120 // Same story for fmax, fmin.
    121 
    122 defm SAT_fmin_fmax_f : SAT<Float32Regs, f32imm, int_nvvm_fmin_f,
    123   int_nvvm_fmax_f, immFloat0, immFloat1,
    124            "cvt.sat.f32.f32 \t$dst, $src; \n">;
    125 defm SAT_fmin_fmax_d : SAT<Float64Regs, f64imm, int_nvvm_fmin_d,
    126   int_nvvm_fmax_d, immDouble0, immDouble1,
    127            "cvt.sat.f64.f64 \t$dst, $src; \n">;
    128 
    129 
    130 // We need a full string for OpcStr here because we need to deal with case like
    131 // INT_PTX_RECIP.
    132 class F_MATH_1<string OpcStr, NVPTXRegClass target_regclass,
    133   NVPTXRegClass src_regclass, Intrinsic IntOP>
    134             : NVPTXInst<(outs target_regclass:$dst), (ins src_regclass:$src0),
    135             OpcStr,
    136         [(set target_regclass:$dst, (IntOP src_regclass:$src0))]>;
    137 
    138 // We need a full string for OpcStr here because we need to deal with the case
    139 // like INT_PTX_NATIVE_POWR_F.
    140 class F_MATH_2<string OpcStr, NVPTXRegClass t_regclass,
    141   NVPTXRegClass s0_regclass, NVPTXRegClass s1_regclass, Intrinsic IntOP>
    142             : NVPTXInst<(outs t_regclass:$dst),
    143               (ins s0_regclass:$src0, s1_regclass:$src1),
    144             OpcStr,
    145         [(set t_regclass:$dst, (IntOP s0_regclass:$src0, s1_regclass:$src1))]>;
    146 
    147 class F_MATH_3<string OpcStr, NVPTXRegClass t_regclass,
    148   NVPTXRegClass s0_regclass, NVPTXRegClass s1_regclass,
    149   NVPTXRegClass s2_regclass, Intrinsic IntOP>
    150             : NVPTXInst<(outs t_regclass:$dst),
    151               (ins s0_regclass:$src0, s1_regclass:$src1, s2_regclass:$src2),
    152             OpcStr,
    153         [(set t_regclass:$dst,
    154           (IntOP s0_regclass:$src0, s1_regclass:$src1, s2_regclass:$src2))]>;
    155 
    156 //
    157 // MISC
    158 //
    159 
    160 def INT_NVVM_CLZ_I : F_MATH_1<"clz.b32 \t$dst, $src0;", Int32Regs, Int32Regs,
    161   int_nvvm_clz_i>;
    162 def INT_NVVM_CLZ_LL : F_MATH_1<"clz.b64 \t$dst, $src0;", Int32Regs, Int64Regs,
    163   int_nvvm_clz_ll>;
    164 
    165 def INT_NVVM_POPC_I : F_MATH_1<"popc.b32 \t$dst, $src0;", Int32Regs, Int32Regs,
    166   int_nvvm_popc_i>;
    167 def INT_NVVM_POPC_LL : F_MATH_1<"popc.b64 \t$dst, $src0;", Int32Regs, Int64Regs,
    168   int_nvvm_popc_ll>;
    169 
    170 def INT_NVVM_PRMT : F_MATH_3<"prmt.b32 \t$dst, $src0, $src1, $src2;", Int32Regs,
    171   Int32Regs, Int32Regs, Int32Regs, int_nvvm_prmt>;
    172 
    173 //
    174 // Min Max
    175 //
    176 
    177 def INT_NVVM_MIN_I : F_MATH_2<"min.s32 \t$dst, $src0, $src1;", Int32Regs,
    178   Int32Regs, Int32Regs, int_nvvm_min_i>;
    179 def INT_NVVM_MIN_UI : F_MATH_2<"min.u32 \t$dst, $src0, $src1;", Int32Regs,
    180   Int32Regs, Int32Regs, int_nvvm_min_ui>;
    181 
    182 def INT_NVVM_MIN_LL : F_MATH_2<"min.s64 \t$dst, $src0, $src1;", Int64Regs,
    183   Int64Regs, Int64Regs, int_nvvm_min_ll>;
    184 def INT_NVVM_MIN_ULL : F_MATH_2<"min.u64 \t$dst, $src0, $src1;", Int64Regs,
    185   Int64Regs, Int64Regs, int_nvvm_min_ull>;
    186 
    187 def INT_NVVM_MAX_I : F_MATH_2<"max.s32 \t$dst, $src0, $src1;", Int32Regs,
    188   Int32Regs, Int32Regs, int_nvvm_max_i>;
    189 def INT_NVVM_MAX_UI : F_MATH_2<"max.u32 \t$dst, $src0, $src1;", Int32Regs,
    190   Int32Regs, Int32Regs, int_nvvm_max_ui>;
    191 
    192 def INT_NVVM_MAX_LL : F_MATH_2<"max.s64 \t$dst, $src0, $src1;", Int64Regs,
    193   Int64Regs, Int64Regs, int_nvvm_max_ll>;
    194 def INT_NVVM_MAX_ULL : F_MATH_2<"max.u64 \t$dst, $src0, $src1;", Int64Regs,
    195   Int64Regs, Int64Regs, int_nvvm_max_ull>;
    196 
    197 def INT_NVVM_FMIN_F : F_MATH_2<"min.f32 \t$dst, $src0, $src1;", Float32Regs,
    198   Float32Regs, Float32Regs, int_nvvm_fmin_f>;
    199 def INT_NVVM_FMIN_FTZ_F : F_MATH_2<"min.ftz.f32 \t$dst, $src0, $src1;",
    200   Float32Regs, Float32Regs, Float32Regs, int_nvvm_fmin_ftz_f>;
    201 
    202 def INT_NVVM_FMAX_F : F_MATH_2<"max.f32 \t$dst, $src0, $src1;", Float32Regs,
    203   Float32Regs, Float32Regs, int_nvvm_fmax_f>;
    204 def INT_NVVM_FMAX_FTZ_F : F_MATH_2<"max.ftz.f32 \t$dst, $src0, $src1;",
    205   Float32Regs, Float32Regs, Float32Regs, int_nvvm_fmax_ftz_f>;
    206 
    207 def INT_NVVM_FMIN_D : F_MATH_2<"min.f64 \t$dst, $src0, $src1;", Float64Regs,
    208   Float64Regs, Float64Regs, int_nvvm_fmin_d>;
    209 def INT_NVVM_FMAX_D : F_MATH_2<"max.f64 \t$dst, $src0, $src1;", Float64Regs,
    210   Float64Regs, Float64Regs, int_nvvm_fmax_d>;
    211 
    212 //
    213 // Multiplication
    214 //
    215 
    216 def INT_NVVM_MULHI_I : F_MATH_2<"mul.hi.s32 \t$dst, $src0, $src1;", Int32Regs,
    217   Int32Regs, Int32Regs, int_nvvm_mulhi_i>;
    218 def INT_NVVM_MULHI_UI : F_MATH_2<"mul.hi.u32 \t$dst, $src0, $src1;", Int32Regs,
    219   Int32Regs, Int32Regs, int_nvvm_mulhi_ui>;
    220 
    221 def INT_NVVM_MULHI_LL : F_MATH_2<"mul.hi.s64 \t$dst, $src0, $src1;", Int64Regs,
    222   Int64Regs, Int64Regs, int_nvvm_mulhi_ll>;
    223 def INT_NVVM_MULHI_ULL : F_MATH_2<"mul.hi.u64 \t$dst, $src0, $src1;", Int64Regs,
    224   Int64Regs, Int64Regs, int_nvvm_mulhi_ull>;
    225 
    226 def INT_NVVM_MUL_RN_FTZ_F : F_MATH_2<"mul.rn.ftz.f32 \t$dst, $src0, $src1;",
    227   Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rn_ftz_f>;
    228 def INT_NVVM_MUL_RN_F : F_MATH_2<"mul.rn.f32 \t$dst, $src0, $src1;",
    229   Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rn_f>;
    230 def INT_NVVM_MUL_RZ_FTZ_F : F_MATH_2<"mul.rz.ftz.f32 \t$dst, $src0, $src1;",
    231   Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rz_ftz_f>;
    232 def INT_NVVM_MUL_RZ_F : F_MATH_2<"mul.rz.f32 \t$dst, $src0, $src1;",
    233   Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rz_f>;
    234 def INT_NVVM_MUL_RM_FTZ_F : F_MATH_2<"mul.rm.ftz.f32 \t$dst, $src0, $src1;",
    235   Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rm_ftz_f>;
    236 def INT_NVVM_MUL_RM_F : F_MATH_2<"mul.rm.f32 \t$dst, $src0, $src1;",
    237   Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rm_f>;
    238 def INT_NVVM_MUL_RP_FTZ_F : F_MATH_2<"mul.rp.ftz.f32 \t$dst, $src0, $src1;",
    239   Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rp_ftz_f>;
    240 def INT_NVVM_MUL_RP_F : F_MATH_2<"mul.rp.f32 \t$dst, $src0, $src1;",
    241   Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rp_f>;
    242 
    243 def INT_NVVM_MUL_RN_D : F_MATH_2<"mul.rn.f64 \t$dst, $src0, $src1;",
    244   Float64Regs, Float64Regs, Float64Regs, int_nvvm_mul_rn_d>;
    245 def INT_NVVM_MUL_RZ_D : F_MATH_2<"mul.rz.f64 \t$dst, $src0, $src1;",
    246   Float64Regs, Float64Regs, Float64Regs, int_nvvm_mul_rz_d>;
    247 def INT_NVVM_MUL_RM_D : F_MATH_2<"mul.rm.f64 \t$dst, $src0, $src1;",
    248   Float64Regs, Float64Regs, Float64Regs, int_nvvm_mul_rm_d>;
    249 def INT_NVVM_MUL_RP_D : F_MATH_2<"mul.rp.f64 \t$dst, $src0, $src1;",
    250   Float64Regs, Float64Regs, Float64Regs, int_nvvm_mul_rp_d>;
    251 
    252 def INT_NVVM_MUL24_I : F_MATH_2<"mul24.lo.s32 \t$dst, $src0, $src1;",
    253   Int32Regs, Int32Regs, Int32Regs, int_nvvm_mul24_i>;
    254 def INT_NVVM_MUL24_UI : F_MATH_2<"mul24.lo.u32 \t$dst, $src0, $src1;",
    255   Int32Regs, Int32Regs, Int32Regs, int_nvvm_mul24_ui>;
    256 
    257 //
    258 // Div
    259 //
    260 
    261 def INT_NVVM_DIV_APPROX_FTZ_F
    262   : F_MATH_2<"div.approx.ftz.f32 \t$dst, $src0, $src1;", Float32Regs,
    263     Float32Regs, Float32Regs, int_nvvm_div_approx_ftz_f>;
    264 def INT_NVVM_DIV_APPROX_F : F_MATH_2<"div.approx.f32 \t$dst, $src0, $src1;",
    265   Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_approx_f>;
    266 
    267 def INT_NVVM_DIV_RN_FTZ_F : F_MATH_2<"div.rn.ftz.f32 \t$dst, $src0, $src1;",
    268   Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rn_ftz_f>;
    269 def INT_NVVM_DIV_RN_F     : F_MATH_2<"div.rn.f32 \t$dst, $src0, $src1;",
    270   Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rn_f>;
    271 def INT_NVVM_DIV_RZ_FTZ_F : F_MATH_2<"div.rz.ftz.f32 \t$dst, $src0, $src1;",
    272   Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rz_ftz_f>;
    273 def INT_NVVM_DIV_RZ_F     : F_MATH_2<"div.rz.f32 \t$dst, $src0, $src1;",
    274   Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rz_f>;
    275 def INT_NVVM_DIV_RM_FTZ_F : F_MATH_2<"div.rm.ftz.f32 \t$dst, $src0, $src1;",
    276   Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rm_ftz_f>;
    277 def INT_NVVM_DIV_RM_F     : F_MATH_2<"div.rm.f32 \t$dst, $src0, $src1;",
    278   Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rm_f>;
    279 def INT_NVVM_DIV_RP_FTZ_F : F_MATH_2<"div.rp.ftz.f32 \t$dst, $src0, $src1;",
    280   Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rp_ftz_f>;
    281 def INT_NVVM_DIV_RP_F     : F_MATH_2<"div.rp.f32 \t$dst, $src0, $src1;",
    282   Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rp_f>;
    283 
    284 def INT_NVVM_DIV_RN_D : F_MATH_2<"div.rn.f64 \t$dst, $src0, $src1;",
    285   Float64Regs, Float64Regs, Float64Regs, int_nvvm_div_rn_d>;
    286 def INT_NVVM_DIV_RZ_D : F_MATH_2<"div.rz.f64 \t$dst, $src0, $src1;",
    287   Float64Regs, Float64Regs, Float64Regs, int_nvvm_div_rz_d>;
    288 def INT_NVVM_DIV_RM_D : F_MATH_2<"div.rm.f64 \t$dst, $src0, $src1;",
    289   Float64Regs, Float64Regs, Float64Regs, int_nvvm_div_rm_d>;
    290 def INT_NVVM_DIV_RP_D : F_MATH_2<"div.rp.f64 \t$dst, $src0, $src1;",
    291   Float64Regs, Float64Regs, Float64Regs, int_nvvm_div_rp_d>;
    292 
    293 //
    294 // Brev
    295 //
    296 
    297 def INT_NVVM_BREV32 : F_MATH_1<"brev.b32 \t$dst, $src0;", Int32Regs, Int32Regs,
    298   int_nvvm_brev32>;
    299 def INT_NVVM_BREV64 : F_MATH_1<"brev.b64 \t$dst, $src0;", Int64Regs, Int64Regs,
    300   int_nvvm_brev64>;
    301 
    302 //
    303 // Sad
    304 //
    305 
    306 def INT_NVVM_SAD_I : F_MATH_3<"sad.s32 \t$dst, $src0, $src1, $src2;",
    307   Int32Regs, Int32Regs, Int32Regs, Int32Regs, int_nvvm_sad_i>;
    308 def INT_NVVM_SAD_UI : F_MATH_3<"sad.u32 \t$dst, $src0, $src1, $src2;",
    309   Int32Regs, Int32Regs, Int32Regs, Int32Regs, int_nvvm_sad_ui>;
    310 
    311 //
    312 // Floor  Ceil
    313 //
    314 
    315 def INT_NVVM_FLOOR_FTZ_F : F_MATH_1<"cvt.rmi.ftz.f32.f32 \t$dst, $src0;",
    316   Float32Regs, Float32Regs, int_nvvm_floor_ftz_f>;
    317 def INT_NVVM_FLOOR_F : F_MATH_1<"cvt.rmi.f32.f32 \t$dst, $src0;",
    318   Float32Regs, Float32Regs, int_nvvm_floor_f>;
    319 def INT_NVVM_FLOOR_D : F_MATH_1<"cvt.rmi.f64.f64 \t$dst, $src0;",
    320   Float64Regs, Float64Regs, int_nvvm_floor_d>;
    321 
    322 def INT_NVVM_CEIL_FTZ_F : F_MATH_1<"cvt.rpi.ftz.f32.f32 \t$dst, $src0;",
    323   Float32Regs, Float32Regs, int_nvvm_ceil_ftz_f>;
    324 def INT_NVVM_CEIL_F : F_MATH_1<"cvt.rpi.f32.f32 \t$dst, $src0;",
    325   Float32Regs, Float32Regs, int_nvvm_ceil_f>;
    326 def INT_NVVM_CEIL_D : F_MATH_1<"cvt.rpi.f64.f64 \t$dst, $src0;",
    327   Float64Regs, Float64Regs, int_nvvm_ceil_d>;
    328 
    329 //
    330 // Abs
    331 //
    332 
    333 def INT_NVVM_ABS_I : F_MATH_1<"abs.s32 \t$dst, $src0;", Int32Regs, Int32Regs,
    334   int_nvvm_abs_i>;
    335 def INT_NVVM_ABS_LL : F_MATH_1<"abs.s64 \t$dst, $src0;", Int64Regs, Int64Regs,
    336   int_nvvm_abs_ll>;
    337 
    338 def INT_NVVM_FABS_FTZ_F : F_MATH_1<"abs.ftz.f32 \t$dst, $src0;", Float32Regs,
    339   Float32Regs, int_nvvm_fabs_ftz_f>;
    340 def INT_NVVM_FABS_F : F_MATH_1<"abs.f32 \t$dst, $src0;", Float32Regs,
    341   Float32Regs, int_nvvm_fabs_f>;
    342 
    343 def INT_NVVM_FABS_D : F_MATH_1<"abs.f64 \t$dst, $src0;", Float64Regs,
    344   Float64Regs, int_nvvm_fabs_d>;
    345 
    346 //
    347 // Round
    348 //
    349 
    350 def INT_NVVM_ROUND_FTZ_F : F_MATH_1<"cvt.rni.ftz.f32.f32 \t$dst, $src0;",
    351   Float32Regs, Float32Regs, int_nvvm_round_ftz_f>;
    352 def INT_NVVM_ROUND_F : F_MATH_1<"cvt.rni.f32.f32 \t$dst, $src0;", Float32Regs,
    353   Float32Regs, int_nvvm_round_f>;
    354 
    355 def INT_NVVM_ROUND_D : F_MATH_1<"cvt.rni.f64.f64 \t$dst, $src0;", Float64Regs,
    356   Float64Regs, int_nvvm_round_d>;
    357 
    358 //
    359 // Trunc
    360 //
    361 
    362 def INT_NVVM_TRUNC_FTZ_F : F_MATH_1<"cvt.rzi.ftz.f32.f32 \t$dst, $src0;",
    363   Float32Regs, Float32Regs, int_nvvm_trunc_ftz_f>;
    364 def INT_NVVM_TRUNC_F : F_MATH_1<"cvt.rzi.f32.f32 \t$dst, $src0;", Float32Regs,
    365   Float32Regs, int_nvvm_trunc_f>;
    366 
    367 def INT_NVVM_TRUNC_D : F_MATH_1<"cvt.rzi.f64.f64 \t$dst, $src0;", Float64Regs,
    368   Float64Regs, int_nvvm_trunc_d>;
    369 
    370 //
    371 // Saturate
    372 //
    373 
    374 def INT_NVVM_SATURATE_FTZ_F : F_MATH_1<"cvt.sat.ftz.f32.f32 \t$dst, $src0;",
    375   Float32Regs, Float32Regs, int_nvvm_saturate_ftz_f>;
    376 def INT_NVVM_SATURATE_F : F_MATH_1<"cvt.sat.f32.f32 \t$dst, $src0;",
    377   Float32Regs, Float32Regs, int_nvvm_saturate_f>;
    378 
    379 def INT_NVVM_SATURATE_D : F_MATH_1<"cvt.sat.f64.f64 \t$dst, $src0;",
    380   Float64Regs, Float64Regs, int_nvvm_saturate_d>;
    381 
    382 //
    383 // Exp2  Log2
    384 //
    385 
    386 def INT_NVVM_EX2_APPROX_FTZ_F : F_MATH_1<"ex2.approx.ftz.f32 \t$dst, $src0;",
    387   Float32Regs, Float32Regs, int_nvvm_ex2_approx_ftz_f>;
    388 def INT_NVVM_EX2_APPROX_F : F_MATH_1<"ex2.approx.f32 \t$dst, $src0;",
    389   Float32Regs, Float32Regs, int_nvvm_ex2_approx_f>;
    390 def INT_NVVM_EX2_APPROX_D : F_MATH_1<"ex2.approx.f64 \t$dst, $src0;",
    391   Float64Regs, Float64Regs, int_nvvm_ex2_approx_d>;
    392 
    393 def INT_NVVM_LG2_APPROX_FTZ_F : F_MATH_1<"lg2.approx.ftz.f32 \t$dst, $src0;",
    394   Float32Regs, Float32Regs, int_nvvm_lg2_approx_ftz_f>;
    395 def INT_NVVM_LG2_APPROX_F : F_MATH_1<"lg2.approx.f32 \t$dst, $src0;",
    396   Float32Regs, Float32Regs, int_nvvm_lg2_approx_f>;
    397 def INT_NVVM_LG2_APPROX_D : F_MATH_1<"lg2.approx.f64 \t$dst, $src0;",
    398   Float64Regs, Float64Regs, int_nvvm_lg2_approx_d>;
    399 
    400 //
    401 // Sin  Cos
    402 //
    403 
    404 def INT_NVVM_SIN_APPROX_FTZ_F : F_MATH_1<"sin.approx.ftz.f32 \t$dst, $src0;",
    405   Float32Regs, Float32Regs, int_nvvm_sin_approx_ftz_f>;
    406 def INT_NVVM_SIN_APPROX_F : F_MATH_1<"sin.approx.f32 \t$dst, $src0;",
    407   Float32Regs, Float32Regs, int_nvvm_sin_approx_f>;
    408 
    409 def INT_NVVM_COS_APPROX_FTZ_F : F_MATH_1<"cos.approx.ftz.f32 \t$dst, $src0;",
    410   Float32Regs, Float32Regs, int_nvvm_cos_approx_ftz_f>;
    411 def INT_NVVM_COS_APPROX_F : F_MATH_1<"cos.approx.f32 \t$dst, $src0;",
    412   Float32Regs, Float32Regs, int_nvvm_cos_approx_f>;
    413 
    414 //
    415 // Fma
    416 //
    417 
    418 def INT_NVVM_FMA_RN_FTZ_F
    419   : F_MATH_3<"fma.rn.ftz.f32 \t$dst, $src0, $src1, $src2;", Float32Regs,
    420     Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rn_ftz_f>;
    421 def INT_NVVM_FMA_RN_F : F_MATH_3<"fma.rn.f32 \t$dst, $src0, $src1, $src2;",
    422   Float32Regs, Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rn_f>;
    423 def INT_NVVM_FMA_RZ_FTZ_F
    424   : F_MATH_3<"fma.rz.ftz.f32 \t$dst, $src0, $src1, $src2;", Float32Regs,
    425     Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rz_ftz_f>;
    426 def INT_NVVM_FMA_RZ_F : F_MATH_3<"fma.rz.f32 \t$dst, $src0, $src1, $src2;",
    427   Float32Regs, Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rz_f>;
    428 def INT_NVVM_FMA_RM_FTZ_F
    429   : F_MATH_3<"fma.rm.ftz.f32 \t$dst, $src0, $src1, $src2;", Float32Regs,
    430     Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rm_ftz_f>;
    431 def INT_NVVM_FMA_RM_F : F_MATH_3<"fma.rm.f32 \t$dst, $src0, $src1, $src2;",
    432   Float32Regs, Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rm_f>;
    433 def INT_NVVM_FMA_RP_FTZ_F
    434   : F_MATH_3<"fma.rp.ftz.f32 \t$dst, $src0, $src1, $src2;", Float32Regs,
    435     Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rp_ftz_f>;
    436 def INT_NVVM_FMA_RP_F : F_MATH_3<"fma.rp.f32 \t$dst, $src0, $src1, $src2;",
    437   Float32Regs, Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rp_f>;
    438 
    439 def INT_NVVM_FMA_RN_D : F_MATH_3<"fma.rn.f64 \t$dst, $src0, $src1, $src2;",
    440   Float64Regs, Float64Regs, Float64Regs, Float64Regs, int_nvvm_fma_rn_d>;
    441 def INT_NVVM_FMA_RZ_D : F_MATH_3<"fma.rz.f64 \t$dst, $src0, $src1, $src2;",
    442   Float64Regs, Float64Regs, Float64Regs, Float64Regs, int_nvvm_fma_rz_d>;
    443 def INT_NVVM_FMA_RM_D : F_MATH_3<"fma.rm.f64 \t$dst, $src0, $src1, $src2;",
    444   Float64Regs, Float64Regs, Float64Regs, Float64Regs, int_nvvm_fma_rm_d>;
    445 def INT_NVVM_FMA_RP_D : F_MATH_3<"fma.rp.f64 \t$dst, $src0, $src1, $src2;",
    446   Float64Regs, Float64Regs, Float64Regs, Float64Regs, int_nvvm_fma_rp_d>;
    447 
    448 //
    449 // Rcp
    450 //
    451 
    452 def INT_NVVM_RCP_RN_FTZ_F : F_MATH_1<"rcp.rn.ftz.f32 \t$dst, $src0;",
    453   Float32Regs, Float32Regs, int_nvvm_rcp_rn_ftz_f>;
    454 def INT_NVVM_RCP_RN_F : F_MATH_1<"rcp.rn.f32 \t$dst, $src0;",
    455   Float32Regs, Float32Regs, int_nvvm_rcp_rn_f>;
    456 def INT_NVVM_RCP_RZ_FTZ_F : F_MATH_1<"rcp.rz.ftz.f32 \t$dst, $src0;",
    457   Float32Regs, Float32Regs, int_nvvm_rcp_rz_ftz_f>;
    458 def INT_NVVM_RCP_RZ_F : F_MATH_1<"rcp.rz.f32 \t$dst, $src0;",
    459   Float32Regs, Float32Regs, int_nvvm_rcp_rz_f>;
    460 def INT_NVVM_RCP_RM_FTZ_F : F_MATH_1<"rcp.rm.ftz.f32 \t$dst, $src0;",
    461   Float32Regs, Float32Regs, int_nvvm_rcp_rm_ftz_f>;
    462 def INT_NVVM_RCP_RM_F : F_MATH_1<"rcp.rm.f32 \t$dst, $src0;",
    463   Float32Regs, Float32Regs, int_nvvm_rcp_rm_f>;
    464 def INT_NVVM_RCP_RP_FTZ_F : F_MATH_1<"rcp.rp.ftz.f32 \t$dst, $src0;",
    465   Float32Regs, Float32Regs, int_nvvm_rcp_rp_ftz_f>;
    466 def INT_NVVM_RCP_RP_F : F_MATH_1<"rcp.rp.f32 \t$dst, $src0;",
    467   Float32Regs, Float32Regs, int_nvvm_rcp_rp_f>;
    468 
    469 def INT_NVVM_RCP_RN_D : F_MATH_1<"rcp.rn.f64 \t$dst, $src0;", Float64Regs,
    470   Float64Regs, int_nvvm_rcp_rn_d>;
    471 def INT_NVVM_RCP_RZ_D : F_MATH_1<"rcp.rz.f64 \t$dst, $src0;", Float64Regs,
    472   Float64Regs, int_nvvm_rcp_rz_d>;
    473 def INT_NVVM_RCP_RM_D : F_MATH_1<"rcp.rm.f64 \t$dst, $src0;", Float64Regs,
    474   Float64Regs, int_nvvm_rcp_rm_d>;
    475 def INT_NVVM_RCP_RP_D : F_MATH_1<"rcp.rp.f64 \t$dst, $src0;", Float64Regs,
    476   Float64Regs, int_nvvm_rcp_rp_d>;
    477 
    478 def INT_NVVM_RCP_APPROX_FTZ_D : F_MATH_1<"rcp.approx.ftz.f64 \t$dst, $src0;",
    479   Float64Regs, Float64Regs, int_nvvm_rcp_approx_ftz_d>;
    480 
    481 //
    482 // Sqrt
    483 //
    484 
    485 def INT_NVVM_SQRT_RN_FTZ_F : F_MATH_1<"sqrt.rn.ftz.f32 \t$dst, $src0;",
    486   Float32Regs, Float32Regs, int_nvvm_sqrt_rn_ftz_f>;
    487 def INT_NVVM_SQRT_RN_F : F_MATH_1<"sqrt.rn.f32 \t$dst, $src0;", Float32Regs,
    488   Float32Regs, int_nvvm_sqrt_rn_f>;
    489 def INT_NVVM_SQRT_RZ_FTZ_F : F_MATH_1<"sqrt.rz.ftz.f32 \t$dst, $src0;",
    490   Float32Regs, Float32Regs, int_nvvm_sqrt_rz_ftz_f>;
    491 def INT_NVVM_SQRT_RZ_F : F_MATH_1<"sqrt.rz.f32 \t$dst, $src0;", Float32Regs,
    492   Float32Regs, int_nvvm_sqrt_rz_f>;
    493 def INT_NVVM_SQRT_RM_FTZ_F : F_MATH_1<"sqrt.rm.ftz.f32 \t$dst, $src0;",
    494   Float32Regs, Float32Regs, int_nvvm_sqrt_rm_ftz_f>;
    495 def INT_NVVM_SQRT_RM_F : F_MATH_1<"sqrt.rm.f32 \t$dst, $src0;", Float32Regs,
    496   Float32Regs, int_nvvm_sqrt_rm_f>;
    497 def INT_NVVM_SQRT_RP_FTZ_F : F_MATH_1<"sqrt.rp.ftz.f32 \t$dst, $src0;",
    498   Float32Regs, Float32Regs, int_nvvm_sqrt_rp_ftz_f>;
    499 def INT_NVVM_SQRT_RP_F : F_MATH_1<"sqrt.rp.f32 \t$dst, $src0;", Float32Regs,
    500   Float32Regs, int_nvvm_sqrt_rp_f>;
    501 def INT_NVVM_SQRT_APPROX_FTZ_F : F_MATH_1<"sqrt.approx.ftz.f32 \t$dst, $src0;",
    502   Float32Regs, Float32Regs, int_nvvm_sqrt_approx_ftz_f>;
    503 def INT_NVVM_SQRT_APPROX_F : F_MATH_1<"sqrt.approx.f32 \t$dst, $src0;",
    504   Float32Regs, Float32Regs, int_nvvm_sqrt_approx_f>;
    505 
    506 def INT_NVVM_SQRT_RN_D : F_MATH_1<"sqrt.rn.f64 \t$dst, $src0;", Float64Regs,
    507   Float64Regs, int_nvvm_sqrt_rn_d>;
    508 def INT_NVVM_SQRT_RZ_D : F_MATH_1<"sqrt.rz.f64 \t$dst, $src0;", Float64Regs,
    509   Float64Regs, int_nvvm_sqrt_rz_d>;
    510 def INT_NVVM_SQRT_RM_D : F_MATH_1<"sqrt.rm.f64 \t$dst, $src0;", Float64Regs,
    511   Float64Regs, int_nvvm_sqrt_rm_d>;
    512 def INT_NVVM_SQRT_RP_D : F_MATH_1<"sqrt.rp.f64 \t$dst, $src0;", Float64Regs,
    513   Float64Regs, int_nvvm_sqrt_rp_d>;
    514 
    515 //
    516 // Rsqrt
    517 //
    518 
    519 def INT_NVVM_RSQRT_APPROX_FTZ_F
    520   : F_MATH_1<"rsqrt.approx.ftz.f32 \t$dst, $src0;", Float32Regs, Float32Regs,
    521     int_nvvm_rsqrt_approx_ftz_f>;
    522 def INT_NVVM_RSQRT_APPROX_F : F_MATH_1<"rsqrt.approx.f32 \t$dst, $src0;",
    523   Float32Regs, Float32Regs, int_nvvm_rsqrt_approx_f>;
    524 def INT_NVVM_RSQRT_APPROX_D : F_MATH_1<"rsqrt.approx.f64 \t$dst, $src0;",
    525   Float64Regs, Float64Regs, int_nvvm_rsqrt_approx_d>;
    526 
    527 //
    528 // Add
    529 //
    530 
    531 def INT_NVVM_ADD_RN_FTZ_F : F_MATH_2<"add.rn.ftz.f32 \t$dst, $src0, $src1;",
    532   Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rn_ftz_f>;
    533 def INT_NVVM_ADD_RN_F : F_MATH_2<"add.rn.f32 \t$dst, $src0, $src1;",
    534   Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rn_f>;
    535 def INT_NVVM_ADD_RZ_FTZ_F : F_MATH_2<"add.rz.ftz.f32 \t$dst, $src0, $src1;",
    536   Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rz_ftz_f>;
    537 def INT_NVVM_ADD_RZ_F : F_MATH_2<"add.rz.f32 \t$dst, $src0, $src1;",
    538   Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rz_f>;
    539 def INT_NVVM_ADD_RM_FTZ_F : F_MATH_2<"add.rm.ftz.f32 \t$dst, $src0, $src1;",
    540   Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rm_ftz_f>;
    541 def INT_NVVM_ADD_RM_F : F_MATH_2<"add.rm.f32 \t$dst, $src0, $src1;",
    542   Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rm_f>;
    543 def INT_NVVM_ADD_RP_FTZ_F : F_MATH_2<"add.rp.ftz.f32 \t$dst, $src0, $src1;",
    544   Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rp_ftz_f>;
    545 def INT_NVVM_ADD_RP_F : F_MATH_2<"add.rp.f32 \t$dst, $src0, $src1;",
    546   Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rp_f>;
    547 
    548 def INT_NVVM_ADD_RN_D : F_MATH_2<"add.rn.f64 \t$dst, $src0, $src1;",
    549   Float64Regs, Float64Regs, Float64Regs, int_nvvm_add_rn_d>;
    550 def INT_NVVM_ADD_RZ_D : F_MATH_2<"add.rz.f64 \t$dst, $src0, $src1;",
    551   Float64Regs, Float64Regs, Float64Regs, int_nvvm_add_rz_d>;
    552 def INT_NVVM_ADD_RM_D : F_MATH_2<"add.rm.f64 \t$dst, $src0, $src1;",
    553   Float64Regs, Float64Regs, Float64Regs, int_nvvm_add_rm_d>;
    554 def INT_NVVM_ADD_RP_D : F_MATH_2<"add.rp.f64 \t$dst, $src0, $src1;",
    555   Float64Regs, Float64Regs, Float64Regs, int_nvvm_add_rp_d>;
    556 
    557 //
    558 // Convert
    559 //
    560 
    561 def INT_NVVM_D2F_RN_FTZ : F_MATH_1<"cvt.rn.ftz.f32.f64 \t$dst, $src0;",
    562   Float32Regs, Float64Regs, int_nvvm_d2f_rn_ftz>;
    563 def INT_NVVM_D2F_RN : F_MATH_1<"cvt.rn.f32.f64 \t$dst, $src0;",
    564   Float32Regs, Float64Regs, int_nvvm_d2f_rn>;
    565 def INT_NVVM_D2F_RZ_FTZ : F_MATH_1<"cvt.rz.ftz.f32.f64 \t$dst, $src0;",
    566   Float32Regs, Float64Regs, int_nvvm_d2f_rz_ftz>;
    567 def INT_NVVM_D2F_RZ : F_MATH_1<"cvt.rz.f32.f64 \t$dst, $src0;",
    568   Float32Regs, Float64Regs, int_nvvm_d2f_rz>;
    569 def INT_NVVM_D2F_RM_FTZ : F_MATH_1<"cvt.rm.ftz.f32.f64 \t$dst, $src0;",
    570   Float32Regs, Float64Regs, int_nvvm_d2f_rm_ftz>;
    571 def INT_NVVM_D2F_RM : F_MATH_1<"cvt.rm.f32.f64 \t$dst, $src0;",
    572   Float32Regs, Float64Regs, int_nvvm_d2f_rm>;
    573 def INT_NVVM_D2F_RP_FTZ : F_MATH_1<"cvt.rp.ftz.f32.f64 \t$dst, $src0;",
    574   Float32Regs, Float64Regs, int_nvvm_d2f_rp_ftz>;
    575 def INT_NVVM_D2F_RP : F_MATH_1<"cvt.rp.f32.f64 \t$dst, $src0;",
    576   Float32Regs, Float64Regs, int_nvvm_d2f_rp>;
    577 
    578 def INT_NVVM_D2I_RN : F_MATH_1<"cvt.rni.s32.f64 \t$dst, $src0;",
    579   Int32Regs, Float64Regs, int_nvvm_d2i_rn>;
    580 def INT_NVVM_D2I_RZ : F_MATH_1<"cvt.rzi.s32.f64 \t$dst, $src0;",
    581   Int32Regs, Float64Regs, int_nvvm_d2i_rz>;
    582 def INT_NVVM_D2I_RM : F_MATH_1<"cvt.rmi.s32.f64 \t$dst, $src0;",
    583   Int32Regs, Float64Regs, int_nvvm_d2i_rm>;
    584 def INT_NVVM_D2I_RP : F_MATH_1<"cvt.rpi.s32.f64 \t$dst, $src0;",
    585   Int32Regs, Float64Regs, int_nvvm_d2i_rp>;
    586 
    587 def INT_NVVM_D2UI_RN : F_MATH_1<"cvt.rni.u32.f64 \t$dst, $src0;",
    588   Int32Regs, Float64Regs, int_nvvm_d2ui_rn>;
    589 def INT_NVVM_D2UI_RZ : F_MATH_1<"cvt.rzi.u32.f64 \t$dst, $src0;",
    590   Int32Regs, Float64Regs, int_nvvm_d2ui_rz>;
    591 def INT_NVVM_D2UI_RM : F_MATH_1<"cvt.rmi.u32.f64 \t$dst, $src0;",
    592   Int32Regs, Float64Regs, int_nvvm_d2ui_rm>;
    593 def INT_NVVM_D2UI_RP : F_MATH_1<"cvt.rpi.u32.f64 \t$dst, $src0;",
    594   Int32Regs, Float64Regs, int_nvvm_d2ui_rp>;
    595 
    596 def INT_NVVM_I2D_RN : F_MATH_1<"cvt.rn.f64.s32 \t$dst, $src0;",
    597   Float64Regs, Int32Regs, int_nvvm_i2d_rn>;
    598 def INT_NVVM_I2D_RZ : F_MATH_1<"cvt.rz.f64.s32 \t$dst, $src0;",
    599   Float64Regs, Int32Regs, int_nvvm_i2d_rz>;
    600 def INT_NVVM_I2D_RM : F_MATH_1<"cvt.rm.f64.s32 \t$dst, $src0;",
    601   Float64Regs, Int32Regs, int_nvvm_i2d_rm>;
    602 def INT_NVVM_I2D_RP : F_MATH_1<"cvt.rp.f64.s32 \t$dst, $src0;",
    603   Float64Regs, Int32Regs, int_nvvm_i2d_rp>;
    604 
    605 def INT_NVVM_UI2D_RN : F_MATH_1<"cvt.rn.f64.u32 \t$dst, $src0;",
    606   Float64Regs, Int32Regs, int_nvvm_ui2d_rn>;
    607 def INT_NVVM_UI2D_RZ : F_MATH_1<"cvt.rz.f64.u32 \t$dst, $src0;",
    608   Float64Regs, Int32Regs, int_nvvm_ui2d_rz>;
    609 def INT_NVVM_UI2D_RM : F_MATH_1<"cvt.rm.f64.u32 \t$dst, $src0;",
    610   Float64Regs, Int32Regs, int_nvvm_ui2d_rm>;
    611 def INT_NVVM_UI2D_RP : F_MATH_1<"cvt.rp.f64.u32 \t$dst, $src0;",
    612   Float64Regs, Int32Regs, int_nvvm_ui2d_rp>;
    613 
    614 def INT_NVVM_F2I_RN_FTZ : F_MATH_1<"cvt.rni.ftz.s32.f32 \t$dst, $src0;",
    615   Int32Regs, Float32Regs, int_nvvm_f2i_rn_ftz>;
    616 def INT_NVVM_F2I_RN : F_MATH_1<"cvt.rni.s32.f32 \t$dst, $src0;", Int32Regs,
    617   Float32Regs, int_nvvm_f2i_rn>;
    618 def INT_NVVM_F2I_RZ_FTZ : F_MATH_1<"cvt.rzi.ftz.s32.f32 \t$dst, $src0;",
    619   Int32Regs, Float32Regs, int_nvvm_f2i_rz_ftz>;
    620 def INT_NVVM_F2I_RZ : F_MATH_1<"cvt.rzi.s32.f32 \t$dst, $src0;", Int32Regs,
    621   Float32Regs, int_nvvm_f2i_rz>;
    622 def INT_NVVM_F2I_RM_FTZ : F_MATH_1<"cvt.rmi.ftz.s32.f32 \t$dst, $src0;",
    623   Int32Regs, Float32Regs, int_nvvm_f2i_rm_ftz>;
    624 def INT_NVVM_F2I_RM : F_MATH_1<"cvt.rmi.s32.f32 \t$dst, $src0;", Int32Regs,
    625   Float32Regs, int_nvvm_f2i_rm>;
    626 def INT_NVVM_F2I_RP_FTZ : F_MATH_1<"cvt.rpi.ftz.s32.f32 \t$dst, $src0;",
    627   Int32Regs, Float32Regs, int_nvvm_f2i_rp_ftz>;
    628 def INT_NVVM_F2I_RP : F_MATH_1<"cvt.rpi.s32.f32 \t$dst, $src0;", Int32Regs,
    629   Float32Regs, int_nvvm_f2i_rp>;
    630 
    631 def INT_NVVM_F2UI_RN_FTZ : F_MATH_1<"cvt.rni.ftz.u32.f32 \t$dst, $src0;",
    632   Int32Regs, Float32Regs, int_nvvm_f2ui_rn_ftz>;
    633 def INT_NVVM_F2UI_RN : F_MATH_1<"cvt.rni.u32.f32 \t$dst, $src0;", Int32Regs,
    634   Float32Regs, int_nvvm_f2ui_rn>;
    635 def INT_NVVM_F2UI_RZ_FTZ : F_MATH_1<"cvt.rzi.ftz.u32.f32 \t$dst, $src0;",
    636   Int32Regs, Float32Regs, int_nvvm_f2ui_rz_ftz>;
    637 def INT_NVVM_F2UI_RZ : F_MATH_1<"cvt.rzi.u32.f32 \t$dst, $src0;", Int32Regs,
    638   Float32Regs, int_nvvm_f2ui_rz>;
    639 def INT_NVVM_F2UI_RM_FTZ : F_MATH_1<"cvt.rmi.ftz.u32.f32 \t$dst, $src0;",
    640   Int32Regs, Float32Regs, int_nvvm_f2ui_rm_ftz>;
    641 def INT_NVVM_F2UI_RM : F_MATH_1<"cvt.rmi.u32.f32 \t$dst, $src0;", Int32Regs,
    642   Float32Regs, int_nvvm_f2ui_rm>;
    643 def INT_NVVM_F2UI_RP_FTZ : F_MATH_1<"cvt.rpi.ftz.u32.f32 \t$dst, $src0;",
    644   Int32Regs, Float32Regs, int_nvvm_f2ui_rp_ftz>;
    645 def INT_NVVM_F2UI_RP : F_MATH_1<"cvt.rpi.u32.f32 \t$dst, $src0;", Int32Regs,
    646   Float32Regs, int_nvvm_f2ui_rp>;
    647 
    648 def INT_NVVM_I2F_RN : F_MATH_1<"cvt.rn.f32.s32 \t$dst, $src0;", Float32Regs,
    649   Int32Regs, int_nvvm_i2f_rn>;
    650 def INT_NVVM_I2F_RZ : F_MATH_1<"cvt.rz.f32.s32 \t$dst, $src0;", Float32Regs,
    651   Int32Regs, int_nvvm_i2f_rz>;
    652 def INT_NVVM_I2F_RM : F_MATH_1<"cvt.rm.f32.s32 \t$dst, $src0;", Float32Regs,
    653   Int32Regs, int_nvvm_i2f_rm>;
    654 def INT_NVVM_I2F_RP : F_MATH_1<"cvt.rp.f32.s32 \t$dst, $src0;", Float32Regs,
    655   Int32Regs, int_nvvm_i2f_rp>;
    656 
    657 def INT_NVVM_UI2F_RN : F_MATH_1<"cvt.rn.f32.u32 \t$dst, $src0;", Float32Regs,
    658   Int32Regs, int_nvvm_ui2f_rn>;
    659 def INT_NVVM_UI2F_RZ : F_MATH_1<"cvt.rz.f32.u32 \t$dst, $src0;", Float32Regs,
    660   Int32Regs, int_nvvm_ui2f_rz>;
    661 def INT_NVVM_UI2F_RM : F_MATH_1<"cvt.rm.f32.u32 \t$dst, $src0;", Float32Regs,
    662   Int32Regs, int_nvvm_ui2f_rm>;
    663 def INT_NVVM_UI2F_RP : F_MATH_1<"cvt.rp.f32.u32 \t$dst, $src0;", Float32Regs,
    664   Int32Regs, int_nvvm_ui2f_rp>;
    665 
    666 def INT_NVVM_LOHI_I2D : F_MATH_2<"mov.b64 \t$dst, {{$src0, $src1}};",
    667   Float64Regs, Int32Regs, Int32Regs, int_nvvm_lohi_i2d>;
    668 
    669 def INT_NVVM_D2I_LO : F_MATH_1<!strconcat("{{\n\t",
    670                        !strconcat(".reg .b32 %temp; \n\t",
    671              !strconcat("mov.b64 \t{$dst, %temp}, $src0;\n\t",
    672                "}}"))),
    673              Int32Regs, Float64Regs, int_nvvm_d2i_lo>;
    674 def INT_NVVM_D2I_HI : F_MATH_1<!strconcat("{{\n\t",
    675                        !strconcat(".reg .b32 %temp; \n\t",
    676                          !strconcat("mov.b64 \t{%temp, $dst}, $src0;\n\t",
    677                            "}}"))),
    678              Int32Regs, Float64Regs, int_nvvm_d2i_hi>;
    679 
    680 def INT_NVVM_F2LL_RN_FTZ : F_MATH_1<"cvt.rni.ftz.s64.f32 \t$dst, $src0;",
    681   Int64Regs, Float32Regs, int_nvvm_f2ll_rn_ftz>;
    682 def INT_NVVM_F2LL_RN : F_MATH_1<"cvt.rni.s64.f32 \t$dst, $src0;", Int64Regs,
    683   Float32Regs, int_nvvm_f2ll_rn>;
    684 def INT_NVVM_F2LL_RZ_FTZ : F_MATH_1<"cvt.rzi.ftz.s64.f32 \t$dst, $src0;",
    685   Int64Regs, Float32Regs, int_nvvm_f2ll_rz_ftz>;
    686 def INT_NVVM_F2LL_RZ : F_MATH_1<"cvt.rzi.s64.f32 \t$dst, $src0;", Int64Regs,
    687   Float32Regs, int_nvvm_f2ll_rz>;
    688 def INT_NVVM_F2LL_RM_FTZ : F_MATH_1<"cvt.rmi.ftz.s64.f32 \t$dst, $src0;",
    689   Int64Regs, Float32Regs, int_nvvm_f2ll_rm_ftz>;
    690 def INT_NVVM_F2LL_RM : F_MATH_1<"cvt.rmi.s64.f32 \t$dst, $src0;", Int64Regs,
    691   Float32Regs, int_nvvm_f2ll_rm>;
    692 def INT_NVVM_F2LL_RP_FTZ : F_MATH_1<"cvt.rpi.ftz.s64.f32 \t$dst, $src0;",
    693   Int64Regs, Float32Regs, int_nvvm_f2ll_rp_ftz>;
    694 def INT_NVVM_F2LL_RP : F_MATH_1<"cvt.rpi.s64.f32 \t$dst, $src0;", Int64Regs,
    695   Float32Regs, int_nvvm_f2ll_rp>;
    696 
    697 def INT_NVVM_F2ULL_RN_FTZ : F_MATH_1<"cvt.rni.ftz.u64.f32 \t$dst, $src0;",
    698   Int64Regs, Float32Regs, int_nvvm_f2ull_rn_ftz>;
    699 def INT_NVVM_F2ULL_RN : F_MATH_1<"cvt.rni.u64.f32 \t$dst, $src0;", Int64Regs,
    700   Float32Regs, int_nvvm_f2ull_rn>;
    701 def INT_NVVM_F2ULL_RZ_FTZ : F_MATH_1<"cvt.rzi.ftz.u64.f32 \t$dst, $src0;",
    702   Int64Regs, Float32Regs, int_nvvm_f2ull_rz_ftz>;
    703 def INT_NVVM_F2ULL_RZ : F_MATH_1<"cvt.rzi.u64.f32 \t$dst, $src0;", Int64Regs,
    704   Float32Regs, int_nvvm_f2ull_rz>;
    705 def INT_NVVM_F2ULL_RM_FTZ : F_MATH_1<"cvt.rmi.ftz.u64.f32 \t$dst, $src0;",
    706   Int64Regs, Float32Regs, int_nvvm_f2ull_rm_ftz>;
    707 def INT_NVVM_F2ULL_RM : F_MATH_1<"cvt.rmi.u64.f32 \t$dst, $src0;", Int64Regs,
    708   Float32Regs, int_nvvm_f2ull_rm>;
    709 def INT_NVVM_F2ULL_RP_FTZ : F_MATH_1<"cvt.rpi.ftz.u64.f32 \t$dst, $src0;",
    710   Int64Regs, Float32Regs, int_nvvm_f2ull_rp_ftz>;
    711 def INT_NVVM_F2ULL_RP : F_MATH_1<"cvt.rpi.u64.f32 \t$dst, $src0;", Int64Regs,
    712   Float32Regs, int_nvvm_f2ull_rp>;
    713 
    714 def INT_NVVM_D2LL_RN : F_MATH_1<"cvt.rni.s64.f64 \t$dst, $src0;", Int64Regs,
    715   Float64Regs, int_nvvm_d2ll_rn>;
    716 def INT_NVVM_D2LL_RZ : F_MATH_1<"cvt.rzi.s64.f64 \t$dst, $src0;", Int64Regs,
    717   Float64Regs, int_nvvm_d2ll_rz>;
    718 def INT_NVVM_D2LL_RM : F_MATH_1<"cvt.rmi.s64.f64 \t$dst, $src0;", Int64Regs,
    719   Float64Regs, int_nvvm_d2ll_rm>;
    720 def INT_NVVM_D2LL_RP : F_MATH_1<"cvt.rpi.s64.f64 \t$dst, $src0;", Int64Regs,
    721   Float64Regs, int_nvvm_d2ll_rp>;
    722 
    723 def INT_NVVM_D2ULL_RN : F_MATH_1<"cvt.rni.u64.f64 \t$dst, $src0;", Int64Regs,
    724   Float64Regs, int_nvvm_d2ull_rn>;
    725 def INT_NVVM_D2ULL_RZ : F_MATH_1<"cvt.rzi.u64.f64 \t$dst, $src0;", Int64Regs,
    726   Float64Regs, int_nvvm_d2ull_rz>;
    727 def INT_NVVM_D2ULL_RM : F_MATH_1<"cvt.rmi.u64.f64 \t$dst, $src0;", Int64Regs,
    728   Float64Regs, int_nvvm_d2ull_rm>;
    729 def INT_NVVM_D2ULL_RP : F_MATH_1<"cvt.rpi.u64.f64 \t$dst, $src0;", Int64Regs,
    730   Float64Regs, int_nvvm_d2ull_rp>;
    731 
    732 def INT_NVVM_LL2F_RN : F_MATH_1<"cvt.rn.f32.s64 \t$dst, $src0;", Float32Regs,
    733   Int64Regs, int_nvvm_ll2f_rn>;
    734 def INT_NVVM_LL2F_RZ : F_MATH_1<"cvt.rz.f32.s64 \t$dst, $src0;", Float32Regs,
    735   Int64Regs, int_nvvm_ll2f_rz>;
    736 def INT_NVVM_LL2F_RM : F_MATH_1<"cvt.rm.f32.s64 \t$dst, $src0;", Float32Regs,
    737   Int64Regs, int_nvvm_ll2f_rm>;
    738 def INT_NVVM_LL2F_RP : F_MATH_1<"cvt.rp.f32.s64 \t$dst, $src0;", Float32Regs,
    739   Int64Regs, int_nvvm_ll2f_rp>;
    740 def INT_NVVM_ULL2F_RN : F_MATH_1<"cvt.rn.f32.u64 \t$dst, $src0;", Float32Regs,
    741   Int64Regs, int_nvvm_ull2f_rn>;
    742 def INT_NVVM_ULL2F_RZ : F_MATH_1<"cvt.rz.f32.u64 \t$dst, $src0;", Float32Regs,
    743   Int64Regs, int_nvvm_ull2f_rz>;
    744 def INT_NVVM_ULL2F_RM : F_MATH_1<"cvt.rm.f32.u64 \t$dst, $src0;", Float32Regs,
    745   Int64Regs, int_nvvm_ull2f_rm>;
    746 def INT_NVVM_ULL2F_RP : F_MATH_1<"cvt.rp.f32.u64 \t$dst, $src0;", Float32Regs,
    747   Int64Regs, int_nvvm_ull2f_rp>;
    748 
    749 def INT_NVVM_LL2D_RN : F_MATH_1<"cvt.rn.f64.s64 \t$dst, $src0;", Float64Regs,
    750   Int64Regs, int_nvvm_ll2d_rn>;
    751 def INT_NVVM_LL2D_RZ : F_MATH_1<"cvt.rz.f64.s64 \t$dst, $src0;", Float64Regs,
    752   Int64Regs, int_nvvm_ll2d_rz>;
    753 def INT_NVVM_LL2D_RM : F_MATH_1<"cvt.rm.f64.s64 \t$dst, $src0;", Float64Regs,
    754   Int64Regs, int_nvvm_ll2d_rm>;
    755 def INT_NVVM_LL2D_RP : F_MATH_1<"cvt.rp.f64.s64 \t$dst, $src0;", Float64Regs,
    756   Int64Regs, int_nvvm_ll2d_rp>;
    757 def INT_NVVM_ULL2D_RN : F_MATH_1<"cvt.rn.f64.u64 \t$dst, $src0;", Float64Regs,
    758   Int64Regs, int_nvvm_ull2d_rn>;
    759 def INT_NVVM_ULL2D_RZ : F_MATH_1<"cvt.rz.f64.u64 \t$dst, $src0;", Float64Regs,
    760   Int64Regs, int_nvvm_ull2d_rz>;
    761 def INT_NVVM_ULL2D_RM : F_MATH_1<"cvt.rm.f64.u64 \t$dst, $src0;", Float64Regs,
    762   Int64Regs, int_nvvm_ull2d_rm>;
    763 def INT_NVVM_ULL2D_RP : F_MATH_1<"cvt.rp.f64.u64 \t$dst, $src0;", Float64Regs,
    764   Int64Regs, int_nvvm_ull2d_rp>;
    765 
    766 def INT_NVVM_F2H_RN_FTZ : F_MATH_1<!strconcat("{{\n\t",
    767                                    !strconcat(".reg .b16 %temp;\n\t",
    768            !strconcat("cvt.rn.ftz.f16.f32 \t%temp, $src0;\n\t",
    769            !strconcat("mov.b16 \t$dst, %temp;\n",
    770              "}}")))),
    771                                    Int16Regs, Float32Regs, int_nvvm_f2h_rn_ftz>;
    772 def INT_NVVM_F2H_RN : F_MATH_1<!strconcat("{{\n\t",
    773                                    !strconcat(".reg .b16 %temp;\n\t",
    774            !strconcat("cvt.rn.f16.f32 \t%temp, $src0;\n\t",
    775            !strconcat("mov.b16 \t$dst, %temp;\n",
    776              "}}")))),
    777            Int16Regs, Float32Regs, int_nvvm_f2h_rn>;
    778 
    779 def INT_NVVM_H2F : F_MATH_1<!strconcat("{{\n\t",
    780                             !strconcat(".reg .b16 %temp;\n\t",
    781           !strconcat("mov.b16 \t%temp, $src0;\n\t",
    782           !strconcat("cvt.f32.f16 \t$dst, %temp;\n\t",
    783             "}}")))),
    784           Float32Regs, Int16Regs, int_nvvm_h2f>;
    785 
    786 //
    787 // Bitcast
    788 //
    789 
    790 def INT_NVVM_BITCAST_F2I : F_MATH_1<"mov.b32 \t$dst, $src0;", Int32Regs,
    791   Float32Regs, int_nvvm_bitcast_f2i>;
    792 def INT_NVVM_BITCAST_I2F : F_MATH_1<"mov.b32 \t$dst, $src0;", Float32Regs,
    793   Int32Regs, int_nvvm_bitcast_i2f>;
    794 
    795 def INT_NVVM_BITCAST_LL2D : F_MATH_1<"mov.b64 \t$dst, $src0;", Float64Regs,
    796   Int64Regs, int_nvvm_bitcast_ll2d>;
    797 def INT_NVVM_BITCAST_D2LL : F_MATH_1<"mov.b64 \t$dst, $src0;", Int64Regs,
    798   Float64Regs, int_nvvm_bitcast_d2ll>;
    799 
    800 //-----------------------------------
    801 // Atomic Functions
    802 //-----------------------------------
    803 
    804 class ATOMIC_GLOBAL_CHK <dag ops, dag frag>
    805  : PatFrag<ops, frag, [{
    806    return ChkMemSDNodeAddressSpace(N, llvm::ADDRESS_SPACE_GLOBAL);
    807 }]>;
    808 class ATOMIC_SHARED_CHK <dag ops, dag frag>
    809  : PatFrag<ops, frag, [{
    810    return ChkMemSDNodeAddressSpace(N, llvm::ADDRESS_SPACE_SHARED);
    811 }]>;
    812 class ATOMIC_GENERIC_CHK <dag ops, dag frag>
    813  : PatFrag<ops, frag, [{
    814    return ChkMemSDNodeAddressSpace(N, llvm::ADDRESS_SPACE_GENERIC);
    815 }]>;
    816 
    817 multiclass F_ATOMIC_2_imp<NVPTXRegClass ptrclass, NVPTXRegClass regclass,
    818   string SpaceStr, string TypeStr, string OpcStr, PatFrag IntOp,
    819   Operand IMMType, SDNode IMM, Predicate Pred> {
    820   def reg : NVPTXInst<(outs regclass:$dst), (ins ptrclass:$addr, regclass:$b),
    821                !strconcat("atom",
    822          !strconcat(SpaceStr,
    823          !strconcat(OpcStr,
    824          !strconcat(TypeStr,
    825          !strconcat(" \t$dst, [$addr], $b;", ""))))),
    826          [(set regclass:$dst, (IntOp ptrclass:$addr, regclass:$b))]>,
    827   Requires<[Pred]>;
    828   def imm : NVPTXInst<(outs regclass:$dst), (ins ptrclass:$addr, IMMType:$b),
    829                !strconcat("atom",
    830          !strconcat(SpaceStr,
    831          !strconcat(OpcStr,
    832          !strconcat(TypeStr,
    833          !strconcat(" \t$dst, [$addr], $b;", ""))))),
    834          [(set regclass:$dst, (IntOp ptrclass:$addr, IMM:$b))]>,
    835   Requires<[Pred]>;
    836 }
    837 multiclass F_ATOMIC_2<NVPTXRegClass regclass, string SpaceStr, string TypeStr,
    838   string OpcStr, PatFrag IntOp, Operand IMMType, SDNode IMM, Predicate Pred> {
    839   defm p32 : F_ATOMIC_2_imp<Int32Regs, regclass, SpaceStr, TypeStr, OpcStr,
    840     IntOp, IMMType, IMM, Pred>;
    841   defm p64 : F_ATOMIC_2_imp<Int64Regs, regclass, SpaceStr, TypeStr, OpcStr,
    842     IntOp, IMMType, IMM, Pred>;
    843 }
    844 
    845 // has 2 operands, neg the second one
    846 multiclass F_ATOMIC_2_NEG_imp<NVPTXRegClass ptrclass, NVPTXRegClass regclass,
    847   string SpaceStr, string TypeStr, string OpcStr, PatFrag IntOp,
    848   Operand IMMType, Predicate Pred> {
    849   def reg : NVPTXInst<(outs regclass:$dst), (ins ptrclass:$addr, regclass:$b),
    850     !strconcat("{{ \n\t",
    851          !strconcat(".reg \t.s",
    852          !strconcat(TypeStr,
    853          !strconcat(" temp; \n\t",
    854          !strconcat("neg.s",
    855          !strconcat(TypeStr,
    856          !strconcat(" \ttemp, $b; \n\t",
    857                !strconcat("atom",
    858          !strconcat(SpaceStr,
    859          !strconcat(OpcStr,
    860          !strconcat(".u",
    861          !strconcat(TypeStr,
    862          !strconcat(" \t$dst, [$addr], temp; \n\t",
    863            !strconcat("}}", "")))))))))))))),
    864          [(set regclass:$dst, (IntOp ptrclass:$addr, regclass:$b))]>,
    865   Requires<[Pred]>;
    866 }
    867 multiclass F_ATOMIC_2_NEG<NVPTXRegClass regclass, string SpaceStr,
    868   string TypeStr, string OpcStr, PatFrag IntOp, Operand IMMType,
    869   Predicate Pred> {
    870  defm p32: F_ATOMIC_2_NEG_imp<Int32Regs, regclass, SpaceStr, TypeStr, OpcStr,
    871    IntOp, IMMType, Pred> ;
    872  defm p64: F_ATOMIC_2_NEG_imp<Int64Regs, regclass, SpaceStr, TypeStr, OpcStr,
    873    IntOp, IMMType, Pred> ;
    874 }
    875 
    876 // has 3 operands
    877 multiclass F_ATOMIC_3_imp<NVPTXRegClass ptrclass, NVPTXRegClass regclass,
    878   string SpaceStr, string TypeStr, string OpcStr, PatFrag IntOp,
    879   Operand IMMType, Predicate Pred> {
    880   def reg : NVPTXInst<(outs regclass:$dst),
    881     (ins ptrclass:$addr, regclass:$b, regclass:$c),
    882                !strconcat("atom",
    883          !strconcat(SpaceStr,
    884          !strconcat(OpcStr,
    885          !strconcat(TypeStr,
    886          !strconcat(" \t$dst, [$addr], $b, $c;", ""))))),
    887          [(set regclass:$dst,
    888            (IntOp ptrclass:$addr, regclass:$b, regclass:$c))]>,
    889          Requires<[Pred]>;
    890   def imm1 : NVPTXInst<(outs regclass:$dst),
    891     (ins ptrclass:$addr, IMMType:$b, regclass:$c),
    892                !strconcat("atom",
    893          !strconcat(SpaceStr,
    894          !strconcat(OpcStr,
    895          !strconcat(TypeStr,
    896          !strconcat(" \t$dst, [$addr], $b, $c;", ""))))),
    897          [(set regclass:$dst, (IntOp ptrclass:$addr, imm:$b, regclass:$c))]>,
    898   Requires<[Pred]>;
    899   def imm2 : NVPTXInst<(outs regclass:$dst),
    900     (ins ptrclass:$addr, regclass:$b, IMMType:$c),
    901                !strconcat("atom",
    902          !strconcat(SpaceStr,
    903          !strconcat(OpcStr,
    904          !strconcat(TypeStr,
    905          !strconcat(" \t$dst, [$addr], $b, $c;", ""))))),
    906          [(set regclass:$dst, (IntOp ptrclass:$addr, regclass:$b, imm:$c))]>,
    907   Requires<[Pred]>;
    908   def imm3 : NVPTXInst<(outs regclass:$dst),
    909     (ins ptrclass:$addr, IMMType:$b, IMMType:$c),
    910                !strconcat("atom",
    911          !strconcat(SpaceStr,
    912          !strconcat(OpcStr,
    913          !strconcat(TypeStr,
    914          !strconcat(" \t$dst, [$addr], $b, $c;", ""))))),
    915          [(set regclass:$dst, (IntOp ptrclass:$addr, imm:$b, imm:$c))]>,
    916   Requires<[Pred]>;
    917 }
    918 multiclass F_ATOMIC_3<NVPTXRegClass regclass, string SpaceStr, string TypeStr,
    919   string OpcStr, PatFrag IntOp, Operand IMMType, Predicate Pred> {
    920   defm p32 : F_ATOMIC_3_imp<Int32Regs, regclass, SpaceStr, TypeStr, OpcStr,
    921     IntOp, IMMType, Pred>;
    922   defm p64 : F_ATOMIC_3_imp<Int64Regs, regclass, SpaceStr, TypeStr, OpcStr,
    923     IntOp, IMMType, Pred>;
    924 }
    925 
    926 // atom_add
    927 
    928 def atomic_load_add_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
    929   (atomic_load_add_32 node:$a, node:$b)>;
    930 def atomic_load_add_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
    931   (atomic_load_add_32 node:$a, node:$b)>;
    932 def atomic_load_add_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
    933   (atomic_load_add_32 node:$a, node:$b)>;
    934 def atomic_load_add_64_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
    935   (atomic_load_add_64 node:$a, node:$b)>;
    936 def atomic_load_add_64_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
    937   (atomic_load_add_64 node:$a, node:$b)>;
    938 def atomic_load_add_64_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
    939   (atomic_load_add_64 node:$a, node:$b)>;
    940 def atomic_load_add_f32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
    941   (int_nvvm_atomic_load_add_f32 node:$a, node:$b)>;
    942 def atomic_load_add_f32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
    943   (int_nvvm_atomic_load_add_f32 node:$a, node:$b)>;
    944 def atomic_load_add_f32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
    945   (int_nvvm_atomic_load_add_f32 node:$a, node:$b)>;
    946 
    947 defm INT_PTX_ATOM_ADD_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".u32", ".add",
    948   atomic_load_add_32_g, i32imm, imm, hasAtomRedG32>;
    949 defm INT_PTX_ATOM_ADD_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".u32", ".add",
    950   atomic_load_add_32_s, i32imm, imm, hasAtomRedS32>;
    951 defm INT_PTX_ATOM_ADD_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".u32", ".add",
    952   atomic_load_add_32_gen, i32imm, imm, hasAtomRedGen32>;
    953 defm INT_PTX_ATOM_ADD_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".u32",
    954   ".add", atomic_load_add_32_gen, i32imm, imm, useAtomRedG32forGen32>;
    955 
    956 defm INT_PTX_ATOM_ADD_G_64 : F_ATOMIC_2<Int64Regs, ".global", ".u64", ".add",
    957   atomic_load_add_64_g, i64imm, imm, hasAtomRedG64>;
    958 defm INT_PTX_ATOM_ADD_S_64 : F_ATOMIC_2<Int64Regs, ".shared", ".u64", ".add",
    959   atomic_load_add_64_s, i64imm, imm, hasAtomRedS64>;
    960 defm INT_PTX_ATOM_ADD_GEN_64 : F_ATOMIC_2<Int64Regs, "", ".u64", ".add",
    961   atomic_load_add_64_gen, i64imm, imm, hasAtomRedGen64>;
    962 defm INT_PTX_ATOM_ADD_GEN_64_USE_G : F_ATOMIC_2<Int64Regs, ".global", ".u64",
    963   ".add", atomic_load_add_64_gen, i64imm, imm, useAtomRedG64forGen64>;
    964 
    965 defm INT_PTX_ATOM_ADD_G_F32 : F_ATOMIC_2<Float32Regs, ".global", ".f32", ".add",
    966   atomic_load_add_f32_g, f32imm, fpimm, hasAtomAddF32>;
    967 defm INT_PTX_ATOM_ADD_S_F32 : F_ATOMIC_2<Float32Regs, ".shared", ".f32", ".add",
    968   atomic_load_add_f32_s, f32imm, fpimm, hasAtomAddF32>;
    969 defm INT_PTX_ATOM_ADD_GEN_F32 : F_ATOMIC_2<Float32Regs, "", ".f32", ".add",
    970   atomic_load_add_f32_gen, f32imm, fpimm, hasAtomAddF32>;
    971 
    972 // atom_sub
    973 
    974 def atomic_load_sub_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
    975   (atomic_load_sub_32 node:$a, node:$b)>;
    976 def atomic_load_sub_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
    977   (atomic_load_sub_32 node:$a, node:$b)>;
    978 def atomic_load_sub_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
    979   (atomic_load_sub_32 node:$a, node:$b)>;
    980 def atomic_load_sub_64_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
    981   (atomic_load_sub_64 node:$a, node:$b)>;
    982 def atomic_load_sub_64_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
    983   (atomic_load_sub_64 node:$a, node:$b)>;
    984 def atomic_load_sub_64_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
    985   (atomic_load_sub_64 node:$a, node:$b)>;
    986 
    987 defm INT_PTX_ATOM_SUB_G_32 : F_ATOMIC_2_NEG<Int32Regs, ".global", "32", ".add",
    988   atomic_load_sub_32_g, i32imm, hasAtomRedG32>;
    989 defm INT_PTX_ATOM_SUB_G_64 : F_ATOMIC_2_NEG<Int64Regs, ".global", "64", ".add",
    990   atomic_load_sub_64_g, i64imm, hasAtomRedG64>;
    991 defm INT_PTX_ATOM_SUB_GEN_32 : F_ATOMIC_2_NEG<Int32Regs, "", "32", ".add",
    992   atomic_load_sub_32_gen, i32imm, hasAtomRedGen32>;
    993 defm INT_PTX_ATOM_SUB_GEN_32_USE_G : F_ATOMIC_2_NEG<Int32Regs, ".global", "32",
    994   ".add", atomic_load_sub_32_gen, i32imm, useAtomRedG32forGen32>;
    995 defm INT_PTX_ATOM_SUB_S_32 : F_ATOMIC_2_NEG<Int32Regs, ".shared", "32", ".add",
    996   atomic_load_sub_32_s, i32imm, hasAtomRedS32>;
    997 defm INT_PTX_ATOM_SUB_S_64 : F_ATOMIC_2_NEG<Int64Regs, ".shared", "64", ".add",
    998   atomic_load_sub_64_s, i64imm, hasAtomRedS64>;
    999 defm INT_PTX_ATOM_SUB_GEN_64 : F_ATOMIC_2_NEG<Int64Regs, "", "64", ".add",
   1000   atomic_load_sub_64_gen, i64imm, hasAtomRedGen64>;
   1001 defm INT_PTX_ATOM_SUB_GEN_64_USE_G : F_ATOMIC_2_NEG<Int64Regs, ".global", "64",
   1002   ".add", atomic_load_sub_64_gen, i64imm, useAtomRedG64forGen64>;
   1003 
   1004 // atom_swap
   1005 
   1006 def atomic_swap_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
   1007   (atomic_swap_32 node:$a, node:$b)>;
   1008 def atomic_swap_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
   1009   (atomic_swap_32 node:$a, node:$b)>;
   1010 def atomic_swap_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
   1011   (atomic_swap_32 node:$a, node:$b)>;
   1012 def atomic_swap_64_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
   1013   (atomic_swap_64 node:$a, node:$b)>;
   1014 def atomic_swap_64_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
   1015   (atomic_swap_64 node:$a, node:$b)>;
   1016 def atomic_swap_64_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
   1017   (atomic_swap_64 node:$a, node:$b)>;
   1018 
   1019 defm INT_PTX_ATOM_SWAP_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".b32", ".exch",
   1020   atomic_swap_32_g, i32imm, imm, hasAtomRedG32>;
   1021 defm INT_PTX_ATOM_SWAP_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".b32", ".exch",
   1022   atomic_swap_32_s, i32imm, imm, hasAtomRedS32>;
   1023 defm INT_PTX_ATOM_SWAP_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".b32", ".exch",
   1024   atomic_swap_32_gen, i32imm, imm, hasAtomRedGen32>;
   1025 defm INT_PTX_ATOM_SWAP_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".b32",
   1026   ".exch", atomic_swap_32_gen, i32imm, imm, useAtomRedG32forGen32>;
   1027 defm INT_PTX_ATOM_SWAP_G_64 : F_ATOMIC_2<Int64Regs, ".global", ".b64", ".exch",
   1028   atomic_swap_64_g, i64imm, imm, hasAtomRedG64>;
   1029 defm INT_PTX_ATOM_SWAP_S_64 : F_ATOMIC_2<Int64Regs, ".shared", ".b64", ".exch",
   1030   atomic_swap_64_s, i64imm, imm, hasAtomRedS64>;
   1031 defm INT_PTX_ATOM_SWAP_GEN_64 : F_ATOMIC_2<Int64Regs, "", ".b64", ".exch",
   1032   atomic_swap_64_gen, i64imm, imm, hasAtomRedGen64>;
   1033 defm INT_PTX_ATOM_SWAP_GEN_64_USE_G : F_ATOMIC_2<Int64Regs, ".global", ".b64",
   1034   ".exch", atomic_swap_64_gen, i64imm, imm, useAtomRedG64forGen64>;
   1035 
   1036 // atom_max
   1037 
   1038 def atomic_load_max_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b)
   1039   , (atomic_load_max_32 node:$a, node:$b)>;
   1040 def atomic_load_max_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
   1041   (atomic_load_max_32 node:$a, node:$b)>;
   1042 def atomic_load_max_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
   1043   (atomic_load_max_32 node:$a, node:$b)>;
   1044 def atomic_load_umax_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
   1045   (atomic_load_umax_32 node:$a, node:$b)>;
   1046 def atomic_load_umax_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
   1047   (atomic_load_umax_32 node:$a, node:$b)>;
   1048 def atomic_load_umax_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
   1049   (atomic_load_umax_32 node:$a, node:$b)>;
   1050 
   1051 defm INT_PTX_ATOM_LOAD_MAX_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".s32",
   1052   ".max", atomic_load_max_32_g, i32imm, imm, hasAtomRedG32>;
   1053 defm INT_PTX_ATOM_LOAD_MAX_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".s32",
   1054   ".max", atomic_load_max_32_s, i32imm, imm, hasAtomRedS32>;
   1055 defm INT_PTX_ATOM_LOAD_MAX_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".s32", ".max",
   1056   atomic_load_max_32_gen, i32imm, imm, hasAtomRedGen32>;
   1057 defm INT_PTX_ATOM_LOAD_MAX_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global",
   1058   ".s32", ".max", atomic_load_max_32_gen, i32imm, imm, useAtomRedG32forGen32>;
   1059 defm INT_PTX_ATOM_LOAD_UMAX_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".u32",
   1060   ".max", atomic_load_umax_32_g, i32imm, imm, hasAtomRedG32>;
   1061 defm INT_PTX_ATOM_LOAD_UMAX_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".u32",
   1062   ".max", atomic_load_umax_32_s, i32imm, imm, hasAtomRedS32>;
   1063 defm INT_PTX_ATOM_LOAD_UMAX_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".u32", ".max",
   1064   atomic_load_umax_32_gen, i32imm, imm, hasAtomRedGen32>;
   1065 defm INT_PTX_ATOM_LOAD_UMAX_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global",
   1066   ".u32", ".max", atomic_load_umax_32_gen, i32imm, imm, useAtomRedG32forGen32>;
   1067 
   1068 // atom_min
   1069 
   1070 def atomic_load_min_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
   1071   (atomic_load_min_32 node:$a, node:$b)>;
   1072 def atomic_load_min_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
   1073   (atomic_load_min_32 node:$a, node:$b)>;
   1074 def atomic_load_min_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
   1075   (atomic_load_min_32 node:$a, node:$b)>;
   1076 def atomic_load_umin_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
   1077   (atomic_load_umin_32 node:$a, node:$b)>;
   1078 def atomic_load_umin_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
   1079   (atomic_load_umin_32 node:$a, node:$b)>;
   1080 def atomic_load_umin_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
   1081   (atomic_load_umin_32 node:$a, node:$b)>;
   1082 
   1083 defm INT_PTX_ATOM_LOAD_MIN_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".s32",
   1084   ".min", atomic_load_min_32_g, i32imm, imm, hasAtomRedG32>;
   1085 defm INT_PTX_ATOM_LOAD_MIN_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".s32",
   1086   ".min", atomic_load_min_32_s, i32imm, imm, hasAtomRedS32>;
   1087 defm INT_PTX_ATOM_LOAD_MIN_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".s32", ".min",
   1088   atomic_load_min_32_gen, i32imm, imm, hasAtomRedGen32>;
   1089 defm INT_PTX_ATOM_LOAD_MIN_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global",
   1090   ".s32", ".min", atomic_load_min_32_gen, i32imm, imm, useAtomRedG32forGen32>;
   1091 defm INT_PTX_ATOM_LOAD_UMIN_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".u32",
   1092   ".min", atomic_load_umin_32_g, i32imm, imm, hasAtomRedG32>;
   1093 defm INT_PTX_ATOM_LOAD_UMIN_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".u32",
   1094   ".min", atomic_load_umin_32_s, i32imm, imm, hasAtomRedS32>;
   1095 defm INT_PTX_ATOM_LOAD_UMIN_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".u32", ".min",
   1096   atomic_load_umin_32_gen, i32imm, imm, hasAtomRedGen32>;
   1097 defm INT_PTX_ATOM_LOAD_UMIN_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global",
   1098   ".u32", ".min", atomic_load_umin_32_gen, i32imm, imm, useAtomRedG32forGen32>;
   1099 
   1100 // atom_inc  atom_dec
   1101 
   1102 def atomic_load_inc_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
   1103   (int_nvvm_atomic_load_inc_32 node:$a, node:$b)>;
   1104 def atomic_load_inc_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
   1105   (int_nvvm_atomic_load_inc_32 node:$a, node:$b)>;
   1106 def atomic_load_inc_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
   1107   (int_nvvm_atomic_load_inc_32 node:$a, node:$b)>;
   1108 def atomic_load_dec_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
   1109   (int_nvvm_atomic_load_dec_32 node:$a, node:$b)>;
   1110 def atomic_load_dec_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
   1111   (int_nvvm_atomic_load_dec_32 node:$a, node:$b)>;
   1112 def atomic_load_dec_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
   1113   (int_nvvm_atomic_load_dec_32 node:$a, node:$b)>;
   1114 
   1115 defm INT_PTX_ATOM_INC_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".u32", ".inc",
   1116   atomic_load_inc_32_g, i32imm, imm, hasAtomRedG32>;
   1117 defm INT_PTX_ATOM_INC_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".u32", ".inc",
   1118   atomic_load_inc_32_s, i32imm, imm, hasAtomRedS32>;
   1119 defm INT_PTX_ATOM_INC_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".u32", ".inc",
   1120   atomic_load_inc_32_gen, i32imm, imm, hasAtomRedGen32>;
   1121 defm INT_PTX_ATOM_INC_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".u32",
   1122   ".inc", atomic_load_inc_32_gen, i32imm, imm, useAtomRedG32forGen32>;
   1123 defm INT_PTX_ATOM_DEC_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".u32", ".dec",
   1124   atomic_load_dec_32_g, i32imm, imm, hasAtomRedG32>;
   1125 defm INT_PTX_ATOM_DEC_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".u32", ".dec",
   1126   atomic_load_dec_32_s, i32imm, imm, hasAtomRedS32>;
   1127 defm INT_PTX_ATOM_DEC_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".u32", ".dec",
   1128   atomic_load_dec_32_gen, i32imm, imm, hasAtomRedGen32>;
   1129 defm INT_PTX_ATOM_DEC_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".u32",
   1130   ".dec", atomic_load_dec_32_gen, i32imm, imm, useAtomRedG32forGen32>;
   1131 
   1132 // atom_and
   1133 
   1134 def atomic_load_and_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
   1135   (atomic_load_and_32 node:$a, node:$b)>;
   1136 def atomic_load_and_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
   1137   (atomic_load_and_32 node:$a, node:$b)>;
   1138 def atomic_load_and_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
   1139   (atomic_load_and_32 node:$a, node:$b)>;
   1140 
   1141 defm INT_PTX_ATOM_AND_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".b32", ".and",
   1142   atomic_load_and_32_g, i32imm, imm, hasAtomRedG32>;
   1143 defm INT_PTX_ATOM_AND_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".b32", ".and",
   1144   atomic_load_and_32_s, i32imm, imm, hasAtomRedS32>;
   1145 defm INT_PTX_ATOM_AND_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".b32", ".and",
   1146   atomic_load_and_32_gen, i32imm, imm, hasAtomRedGen32>;
   1147 defm INT_PTX_ATOM_AND_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".b32",
   1148   ".and", atomic_load_and_32_gen, i32imm, imm, useAtomRedG32forGen32>;
   1149 
   1150 // atom_or
   1151 
   1152 def atomic_load_or_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
   1153   (atomic_load_or_32 node:$a, node:$b)>;
   1154 def atomic_load_or_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
   1155   (atomic_load_or_32 node:$a, node:$b)>;
   1156 def atomic_load_or_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
   1157   (atomic_load_or_32 node:$a, node:$b)>;
   1158 
   1159 defm INT_PTX_ATOM_OR_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".b32", ".or",
   1160   atomic_load_or_32_g, i32imm, imm, hasAtomRedG32>;
   1161 defm INT_PTX_ATOM_OR_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".b32", ".or",
   1162   atomic_load_or_32_gen, i32imm, imm, hasAtomRedGen32>;
   1163 defm INT_PTX_ATOM_OR_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".b32",
   1164   ".or", atomic_load_or_32_gen, i32imm, imm, useAtomRedG32forGen32>;
   1165 defm INT_PTX_ATOM_OR_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".b32", ".or",
   1166   atomic_load_or_32_s, i32imm, imm, hasAtomRedS32>;
   1167 
   1168 // atom_xor
   1169 
   1170 def atomic_load_xor_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
   1171   (atomic_load_xor_32 node:$a, node:$b)>;
   1172 def atomic_load_xor_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
   1173   (atomic_load_xor_32 node:$a, node:$b)>;
   1174 def atomic_load_xor_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
   1175   (atomic_load_xor_32 node:$a, node:$b)>;
   1176 
   1177 defm INT_PTX_ATOM_XOR_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".b32", ".xor",
   1178   atomic_load_xor_32_g, i32imm, imm, hasAtomRedG32>;
   1179 defm INT_PTX_ATOM_XOR_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".b32", ".xor",
   1180   atomic_load_xor_32_s, i32imm, imm, hasAtomRedS32>;
   1181 defm INT_PTX_ATOM_XOR_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".b32", ".xor",
   1182   atomic_load_xor_32_gen, i32imm, imm, hasAtomRedGen32>;
   1183 defm INT_PTX_ATOM_XOR_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".b32",
   1184   ".xor", atomic_load_xor_32_gen, i32imm, imm, useAtomRedG32forGen32>;
   1185 
   1186 // atom_cas
   1187 
   1188 def atomic_cmp_swap_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b, node:$c),
   1189   (atomic_cmp_swap_32 node:$a, node:$b, node:$c)>;
   1190 def atomic_cmp_swap_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b, node:$c),
   1191   (atomic_cmp_swap_32 node:$a, node:$b, node:$c)>;
   1192 def atomic_cmp_swap_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b, node:$c),
   1193   (atomic_cmp_swap_32 node:$a, node:$b, node:$c)>;
   1194 def atomic_cmp_swap_64_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b, node:$c),
   1195   (atomic_cmp_swap_64 node:$a, node:$b, node:$c)>;
   1196 def atomic_cmp_swap_64_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b, node:$c),
   1197   (atomic_cmp_swap_64 node:$a, node:$b, node:$c)>;
   1198 def atomic_cmp_swap_64_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b, node:$c),
   1199   (atomic_cmp_swap_64 node:$a, node:$b, node:$c)>;
   1200 
   1201 defm INT_PTX_ATOM_CAS_G_32 : F_ATOMIC_3<Int32Regs, ".global", ".b32", ".cas",
   1202   atomic_cmp_swap_32_g, i32imm, hasAtomRedG32>;
   1203 defm INT_PTX_ATOM_CAS_S_32 : F_ATOMIC_3<Int32Regs, ".shared", ".b32", ".cas",
   1204   atomic_cmp_swap_32_s, i32imm, hasAtomRedS32>;
   1205 defm INT_PTX_ATOM_CAS_GEN_32 : F_ATOMIC_3<Int32Regs, "", ".b32", ".cas",
   1206   atomic_cmp_swap_32_gen, i32imm, hasAtomRedGen32>;
   1207 defm INT_PTX_ATOM_CAS_GEN_32_USE_G : F_ATOMIC_3<Int32Regs, ".global", ".b32",
   1208   ".cas", atomic_cmp_swap_32_gen, i32imm, useAtomRedG32forGen32>;
   1209 defm INT_PTX_ATOM_CAS_G_64 : F_ATOMIC_3<Int64Regs, ".global", ".b64", ".cas",
   1210   atomic_cmp_swap_64_g, i64imm, hasAtomRedG64>;
   1211 defm INT_PTX_ATOM_CAS_S_64 : F_ATOMIC_3<Int64Regs, ".shared", ".b64", ".cas",
   1212   atomic_cmp_swap_64_s, i64imm, hasAtomRedS64>;
   1213 defm INT_PTX_ATOM_CAS_GEN_64 : F_ATOMIC_3<Int64Regs, "", ".b64", ".cas",
   1214   atomic_cmp_swap_64_gen, i64imm, hasAtomRedGen64>;
   1215 defm INT_PTX_ATOM_CAS_GEN_64_USE_G : F_ATOMIC_3<Int64Regs, ".global", ".b64",
   1216   ".cas", atomic_cmp_swap_64_gen, i64imm, useAtomRedG64forGen64>;
   1217 
   1218 
   1219 //-----------------------------------
   1220 // Read Special Registers
   1221 //-----------------------------------
   1222 class F_SREG<string OpStr, NVPTXRegClass regclassOut, Intrinsic IntOp> :
   1223       NVPTXInst<(outs regclassOut:$dst), (ins),
   1224                OpStr,
   1225          [(set regclassOut:$dst, (IntOp))]>;
   1226 
   1227 def INT_PTX_SREG_TID_X : F_SREG<"mov.u32 \t$dst, %tid.x;", Int32Regs,
   1228   int_nvvm_read_ptx_sreg_tid_x>;
   1229 def INT_PTX_SREG_TID_Y : F_SREG<"mov.u32 \t$dst, %tid.y;", Int32Regs,
   1230   int_nvvm_read_ptx_sreg_tid_y>;
   1231 def INT_PTX_SREG_TID_Z : F_SREG<"mov.u32 \t$dst, %tid.z;", Int32Regs,
   1232   int_nvvm_read_ptx_sreg_tid_z>;
   1233 
   1234 def INT_PTX_SREG_NTID_X : F_SREG<"mov.u32 \t$dst, %ntid.x;", Int32Regs,
   1235   int_nvvm_read_ptx_sreg_ntid_x>;
   1236 def INT_PTX_SREG_NTID_Y : F_SREG<"mov.u32 \t$dst, %ntid.y;", Int32Regs,
   1237   int_nvvm_read_ptx_sreg_ntid_y>;
   1238 def INT_PTX_SREG_NTID_Z : F_SREG<"mov.u32 \t$dst, %ntid.z;", Int32Regs,
   1239   int_nvvm_read_ptx_sreg_ntid_z>;
   1240 
   1241 def INT_PTX_SREG_CTAID_X : F_SREG<"mov.u32 \t$dst, %ctaid.x;", Int32Regs,
   1242   int_nvvm_read_ptx_sreg_ctaid_x>;
   1243 def INT_PTX_SREG_CTAID_Y : F_SREG<"mov.u32 \t$dst, %ctaid.y;", Int32Regs,
   1244   int_nvvm_read_ptx_sreg_ctaid_y>;
   1245 def INT_PTX_SREG_CTAID_Z : F_SREG<"mov.u32 \t$dst, %ctaid.z;", Int32Regs,
   1246   int_nvvm_read_ptx_sreg_ctaid_z>;
   1247 
   1248 def INT_PTX_SREG_NCTAID_X : F_SREG<"mov.u32 \t$dst, %nctaid.x;", Int32Regs,
   1249   int_nvvm_read_ptx_sreg_nctaid_x>;
   1250 def INT_PTX_SREG_NCTAID_Y : F_SREG<"mov.u32 \t$dst, %nctaid.y;", Int32Regs,
   1251   int_nvvm_read_ptx_sreg_nctaid_y>;
   1252 def INT_PTX_SREG_NCTAID_Z : F_SREG<"mov.u32 \t$dst, %nctaid.z;", Int32Regs,
   1253   int_nvvm_read_ptx_sreg_nctaid_z>;
   1254 
   1255 def INT_PTX_SREG_WARPSIZE : F_SREG<"mov.u32 \t$dst, WARP_SZ;", Int32Regs,
   1256   int_nvvm_read_ptx_sreg_warpsize>;
   1257 
   1258 
   1259 //-----------------------------------
   1260 // Support for ldu on sm_20 or later
   1261 //-----------------------------------
   1262 
   1263 // Scalar
   1264 // @TODO: Revisit this, Changed imemAny to imem
   1265 multiclass LDU_G<string TyStr, NVPTXRegClass regclass, Intrinsic IntOp> {
   1266   def areg: NVPTXInst<(outs regclass:$result), (ins Int32Regs:$src),
   1267                !strconcat("ldu.global.", TyStr),
   1268          [(set regclass:$result, (IntOp Int32Regs:$src))]>, Requires<[hasLDU]>;
   1269   def areg64: NVPTXInst<(outs regclass:$result), (ins Int64Regs:$src),
   1270                !strconcat("ldu.global.", TyStr),
   1271          [(set regclass:$result, (IntOp Int64Regs:$src))]>, Requires<[hasLDU]>;
   1272  def avar:  NVPTXInst<(outs regclass:$result), (ins imem:$src),
   1273                !strconcat("ldu.global.", TyStr),
   1274                 [(set regclass:$result, (IntOp (Wrapper tglobaladdr:$src)))]>,
   1275                 Requires<[hasLDU]>;
   1276  def ari :  NVPTXInst<(outs regclass:$result), (ins MEMri:$src),
   1277                !strconcat("ldu.global.", TyStr),
   1278          [(set regclass:$result, (IntOp ADDRri:$src))]>, Requires<[hasLDU]>;
   1279  def ari64 :  NVPTXInst<(outs regclass:$result), (ins MEMri64:$src),
   1280                !strconcat("ldu.global.", TyStr),
   1281          [(set regclass:$result, (IntOp ADDRri64:$src))]>, Requires<[hasLDU]>;
   1282 }
   1283 
   1284 defm INT_PTX_LDU_GLOBAL_i8  : LDU_G<"u8 \t$result, [$src];",  Int8Regs,
   1285 int_nvvm_ldu_global_i>;
   1286 defm INT_PTX_LDU_GLOBAL_i16 : LDU_G<"u16 \t$result, [$src];", Int16Regs,
   1287 int_nvvm_ldu_global_i>;
   1288 defm INT_PTX_LDU_GLOBAL_i32 : LDU_G<"u32 \t$result, [$src];", Int32Regs,
   1289 int_nvvm_ldu_global_i>;
   1290 defm INT_PTX_LDU_GLOBAL_i64 : LDU_G<"u64 \t$result, [$src];", Int64Regs,
   1291 int_nvvm_ldu_global_i>;
   1292 defm INT_PTX_LDU_GLOBAL_f32 : LDU_G<"f32 \t$result, [$src];", Float32Regs,
   1293 int_nvvm_ldu_global_f>;
   1294 defm INT_PTX_LDU_GLOBAL_f64 : LDU_G<"f64 \t$result, [$src];", Float64Regs,
   1295 int_nvvm_ldu_global_f>;
   1296 defm INT_PTX_LDU_GLOBAL_p32 : LDU_G<"u32 \t$result, [$src];", Int32Regs,
   1297 int_nvvm_ldu_global_p>;
   1298 defm INT_PTX_LDU_GLOBAL_p64 : LDU_G<"u64 \t$result, [$src];", Int64Regs,
   1299 int_nvvm_ldu_global_p>;
   1300 
   1301 // vector
   1302 
   1303 // Elementized vector ldu
   1304 multiclass VLDU_G_ELE_V2<string TyStr, NVPTXRegClass regclass> {
   1305  def _32:     NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
   1306    (ins Int32Regs:$src),
   1307                      !strconcat("ldu.global.", TyStr), []>;
   1308  def _64:     NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
   1309    (ins Int64Regs:$src),
   1310                      !strconcat("ldu.global.", TyStr), []>;
   1311 }
   1312 
   1313 multiclass VLDU_G_ELE_V4<string TyStr, NVPTXRegClass regclass> {
   1314  def _32:    NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
   1315      regclass:$dst4), (ins Int32Regs:$src),
   1316                !strconcat("ldu.global.", TyStr), []>;
   1317  def _64:    NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
   1318      regclass:$dst4), (ins Int64Regs:$src),
   1319                !strconcat("ldu.global.", TyStr), []>;
   1320 }
   1321 
   1322 defm INT_PTX_LDU_G_v2i8_ELE
   1323   : VLDU_G_ELE_V2<"v2.u8 \t{{$dst1, $dst2}}, [$src];",  Int8Regs>;
   1324 defm INT_PTX_LDU_G_v2i16_ELE
   1325   : VLDU_G_ELE_V2<"v2.u16 \t{{$dst1, $dst2}}, [$src];", Int16Regs>;
   1326 defm INT_PTX_LDU_G_v2i32_ELE
   1327   : VLDU_G_ELE_V2<"v2.u32 \t{{$dst1, $dst2}}, [$src];", Int32Regs>;
   1328 defm INT_PTX_LDU_G_v2f32_ELE
   1329   : VLDU_G_ELE_V2<"v2.f32 \t{{$dst1, $dst2}}, [$src];", Float32Regs>;
   1330 defm INT_PTX_LDU_G_v2i64_ELE
   1331   : VLDU_G_ELE_V2<"v2.u64 \t{{$dst1, $dst2}}, [$src];", Int64Regs>;
   1332 defm INT_PTX_LDU_G_v2f64_ELE
   1333   : VLDU_G_ELE_V2<"v2.f64 \t{{$dst1, $dst2}}, [$src];", Float64Regs>;
   1334 defm INT_PTX_LDU_G_v4i8_ELE
   1335   : VLDU_G_ELE_V4<"v4.u8 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Int8Regs>;
   1336 defm INT_PTX_LDU_G_v4i16_ELE
   1337   : VLDU_G_ELE_V4<"v4.u16 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];",
   1338     Int16Regs>;
   1339 defm INT_PTX_LDU_G_v4i32_ELE
   1340   : VLDU_G_ELE_V4<"v4.u32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];",
   1341     Int32Regs>;
   1342 defm INT_PTX_LDU_G_v4f32_ELE
   1343   : VLDU_G_ELE_V4<"v4.f32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];",
   1344     Float32Regs>;
   1345 
   1346 
   1347 //-----------------------------------
   1348 // Support for ldg on sm_35 or later 
   1349 //-----------------------------------
   1350 
   1351 def ldg_i8 : PatFrag<(ops node:$ptr), (int_nvvm_ldg_global_i node:$ptr), [{
   1352   MemIntrinsicSDNode *M = cast<MemIntrinsicSDNode>(N);
   1353   return M->getMemoryVT() == MVT::i8;
   1354 }]>;
   1355 
   1356 multiclass LDG_G<string TyStr, NVPTXRegClass regclass, Intrinsic IntOp> {
   1357   def areg: NVPTXInst<(outs regclass:$result), (ins Int32Regs:$src),
   1358                !strconcat("ld.global.nc.", TyStr),
   1359          [(set regclass:$result, (IntOp Int32Regs:$src))]>, Requires<[hasLDG]>;
   1360   def areg64: NVPTXInst<(outs regclass:$result), (ins Int64Regs:$src),
   1361                !strconcat("ld.global.nc.", TyStr),
   1362          [(set regclass:$result, (IntOp Int64Regs:$src))]>, Requires<[hasLDG]>;
   1363  def avar:  NVPTXInst<(outs regclass:$result), (ins imem:$src),
   1364                !strconcat("ld.global.nc.", TyStr),
   1365          [(set regclass:$result, (IntOp (Wrapper tglobaladdr:$src)))]>,
   1366          Requires<[hasLDG]>;
   1367  def ari :  NVPTXInst<(outs regclass:$result), (ins MEMri:$src),
   1368                !strconcat("ld.global.nc.", TyStr),
   1369          [(set regclass:$result, (IntOp ADDRri:$src))]>, Requires<[hasLDG]>;
   1370  def ari64 :  NVPTXInst<(outs regclass:$result), (ins MEMri64:$src),
   1371                !strconcat("ld.global.nc.", TyStr),
   1372          [(set regclass:$result, (IntOp ADDRri64:$src))]>, Requires<[hasLDG]>;
   1373 }
   1374 
   1375 multiclass LDG_G_NOINTRIN<string TyStr, NVPTXRegClass regclass, PatFrag IntOp> {
   1376   def areg: NVPTXInst<(outs regclass:$result), (ins Int32Regs:$src),
   1377                !strconcat("ld.global.nc.", TyStr),
   1378          [(set regclass:$result, (IntOp Int32Regs:$src))]>, Requires<[hasLDG]>;
   1379   def areg64: NVPTXInst<(outs regclass:$result), (ins Int64Regs:$src),
   1380                !strconcat("ld.global.nc.", TyStr),
   1381          [(set regclass:$result, (IntOp Int64Regs:$src))]>, Requires<[hasLDG]>;
   1382  def avar:  NVPTXInst<(outs regclass:$result), (ins imem:$src),
   1383                !strconcat("ld.global.nc.", TyStr),
   1384          [(set regclass:$result, (IntOp (Wrapper tglobaladdr:$src)))]>,
   1385         Requires<[hasLDG]>;
   1386  def ari :  NVPTXInst<(outs regclass:$result), (ins MEMri:$src),
   1387                !strconcat("ld.global.nc.", TyStr),
   1388          [(set regclass:$result, (IntOp ADDRri:$src))]>, Requires<[hasLDG]>;
   1389  def ari64 :  NVPTXInst<(outs regclass:$result), (ins MEMri64:$src),
   1390                !strconcat("ld.global.nc.", TyStr),
   1391          [(set regclass:$result, (IntOp ADDRri64:$src))]>, Requires<[hasLDG]>;
   1392 }
   1393 
   1394 defm INT_PTX_LDG_GLOBAL_i8
   1395   : LDG_G_NOINTRIN<"u8 \t$result, [$src];",  Int16Regs, ldg_i8>;
   1396 defm INT_PTX_LDG_GLOBAL_i16
   1397   : LDG_G<"u16 \t$result, [$src];", Int16Regs,   int_nvvm_ldg_global_i>;
   1398 defm INT_PTX_LDG_GLOBAL_i32
   1399   : LDG_G<"u32 \t$result, [$src];", Int32Regs,   int_nvvm_ldg_global_i>;
   1400 defm INT_PTX_LDG_GLOBAL_i64
   1401   : LDG_G<"u64 \t$result, [$src];", Int64Regs,   int_nvvm_ldg_global_i>;
   1402 defm INT_PTX_LDG_GLOBAL_f32
   1403   : LDG_G<"f32 \t$result, [$src];", Float32Regs, int_nvvm_ldg_global_f>;
   1404 defm INT_PTX_LDG_GLOBAL_f64
   1405   : LDG_G<"f64 \t$result, [$src];", Float64Regs, int_nvvm_ldg_global_f>;
   1406 defm INT_PTX_LDG_GLOBAL_p32
   1407   : LDG_G<"u32 \t$result, [$src];", Int32Regs,   int_nvvm_ldg_global_p>;
   1408 defm INT_PTX_LDG_GLOBAL_p64
   1409   : LDG_G<"u64 \t$result, [$src];", Int64Regs,   int_nvvm_ldg_global_p>;
   1410 
   1411 // vector
   1412 
   1413 // Elementized vector ldg 
   1414 multiclass VLDG_G_ELE_V2<string TyStr, NVPTXRegClass regclass> {
   1415  def _32:     NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
   1416                      (ins Int32Regs:$src),
   1417                      !strconcat("ld.global.nc.", TyStr), []>;
   1418  def _64:     NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
   1419                      (ins Int64Regs:$src),
   1420                      !strconcat("ld.global.nc.", TyStr), []>;
   1421 }
   1422 
   1423 multiclass VLDG_G_ELE_V4<string TyStr, NVPTXRegClass regclass> { 
   1424  def _32:    NVPTXInst<(outs regclass:$dst1, regclass:$dst2,
   1425                         regclass:$dst3, regclass:$dst4), (ins Int32Regs:$src),
   1426                !strconcat("ld.global.nc.", TyStr), []>;
   1427  def _64:    NVPTXInst<(outs regclass:$dst1, regclass:$dst2,
   1428                         regclass:$dst3, regclass:$dst4), (ins Int64Regs:$src),
   1429                !strconcat("ld.global.nc.", TyStr), []>;
   1430 }
   1431 
   1432 // FIXME: 8-bit LDG should be fixed once LDG/LDU nodes are made into proper loads.
   1433 defm INT_PTX_LDG_G_v2i8_ELE
   1434   : VLDG_G_ELE_V2<"v2.u8 \t{{$dst1, $dst2}}, [$src];",  Int16Regs>;
   1435 defm INT_PTX_LDG_G_v2i16_ELE
   1436   : VLDG_G_ELE_V2<"v2.u16 \t{{$dst1, $dst2}}, [$src];", Int16Regs>;
   1437 defm INT_PTX_LDG_G_v2i32_ELE
   1438   : VLDG_G_ELE_V2<"v2.u32 \t{{$dst1, $dst2}}, [$src];", Int32Regs>;
   1439 defm INT_PTX_LDG_G_v2f32_ELE
   1440   : VLDG_G_ELE_V2<"v2.f32 \t{{$dst1, $dst2}}, [$src];", Float32Regs>;
   1441 defm INT_PTX_LDG_G_v2i64_ELE
   1442   : VLDG_G_ELE_V2<"v2.u64 \t{{$dst1, $dst2}}, [$src];", Int64Regs>;
   1443 defm INT_PTX_LDG_G_v2f64_ELE
   1444   : VLDG_G_ELE_V2<"v2.f64 \t{{$dst1, $dst2}}, [$src];", Float64Regs>;
   1445 defm INT_PTX_LDG_G_v4i8_ELE
   1446   : VLDG_G_ELE_V4<"v4.u8 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Int16Regs>;
   1447 defm INT_PTX_LDG_G_v4i16_ELE
   1448   : VLDG_G_ELE_V4<"v4.u16 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Int16Regs>;
   1449 defm INT_PTX_LDG_G_v4i32_ELE
   1450   : VLDG_G_ELE_V4<"v4.u32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Int32Regs>;
   1451 defm INT_PTX_LDG_G_v4f32_ELE
   1452   : VLDG_G_ELE_V4<"v4.f32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Float32Regs>;
   1453 
   1454 
   1455 multiclass NG_TO_G<string Str, Intrinsic Intrin> {
   1456    def _yes : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
   1457           !strconcat("cvta.", !strconcat(Str, ".u32 \t$result, $src;")),
   1458       [(set Int32Regs:$result, (Intrin Int32Regs:$src))]>,
   1459    Requires<[hasGenericLdSt]>;
   1460    def _yes_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
   1461           !strconcat("cvta.", !strconcat(Str, ".u64 \t$result, $src;")),
   1462       [(set Int64Regs:$result, (Intrin Int64Regs:$src))]>,
   1463    Requires<[hasGenericLdSt]>;
   1464 
   1465 // @TODO: Are these actually needed?  I believe global addresses will be copied
   1466 // to register values anyway.
   1467    /*def __addr_yes : NVPTXInst<(outs Int32Regs:$result), (ins imemAny:$src),
   1468           !strconcat("cvta.", !strconcat(Str, ".u32 \t$result, $src;")),
   1469       [(set Int32Regs:$result, (Intrin (Wrapper tglobaladdr:$src)))]>,
   1470       Requires<[hasGenericLdSt]>;
   1471    def __addr_yes_64 : NVPTXInst<(outs Int64Regs:$result), (ins imemAny:$src),
   1472           !strconcat("cvta.", !strconcat(Str, ".u64 \t$result, $src;")),
   1473       [(set Int64Regs:$result, (Intrin (Wrapper tglobaladdr:$src)))]>,
   1474       Requires<[hasGenericLdSt]>;*/
   1475 
   1476    def _no : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
   1477           "mov.u32 \t$result, $src;",
   1478       [(set Int32Regs:$result, (Intrin Int32Regs:$src))]>;
   1479    def _no_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
   1480           "mov.u64 \t$result, $src;",
   1481       [(set Int64Regs:$result, (Intrin Int64Regs:$src))]>;
   1482 
   1483 // @TODO: Are these actually needed?  I believe global addresses will be copied
   1484 // to register values anyway.
   1485    /*def _addr_no : NVPTXInst<(outs Int32Regs:$result), (ins imem:$src),
   1486           "mov.u32 \t$result, $src;",
   1487       [(set Int32Regs:$result, (Intrin (Wrapper tglobaladdr:$src)))]>;
   1488    def _addr_no_64 : NVPTXInst<(outs Int64Regs:$result), (ins imem:$src),
   1489           "mov.u64 \t$result, $src;",
   1490       [(set Int64Regs:$result, (Intrin (Wrapper tglobaladdr:$src)))]>;*/
   1491 }
   1492 
   1493 multiclass G_TO_NG<string Str, Intrinsic Intrin> {
   1494    def _yes : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
   1495           !strconcat("cvta.to.", !strconcat(Str, ".u32 \t$result, $src;")),
   1496       [(set Int32Regs:$result, (Intrin Int32Regs:$src))]>,
   1497    Requires<[hasGenericLdSt]>;
   1498    def _yes_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
   1499           !strconcat("cvta.to.", !strconcat(Str, ".u64 \t$result, $src;")),
   1500       [(set Int64Regs:$result, (Intrin Int64Regs:$src))]>,
   1501    Requires<[hasGenericLdSt]>;
   1502    def _no : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
   1503           "mov.u32 \t$result, $src;",
   1504       [(set Int32Regs:$result, (Intrin Int32Regs:$src))]>;
   1505    def _no_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
   1506           "mov.u64 \t$result, $src;",
   1507       [(set Int64Regs:$result, (Intrin Int64Regs:$src))]>;
   1508 }
   1509 
   1510 defm cvta_local  : NG_TO_G<"local", int_nvvm_ptr_local_to_gen>;
   1511 defm cvta_shared : NG_TO_G<"shared", int_nvvm_ptr_shared_to_gen>;
   1512 defm cvta_global : NG_TO_G<"global", int_nvvm_ptr_global_to_gen>;
   1513 
   1514 defm cvta_to_local   : G_TO_NG<"local", int_nvvm_ptr_gen_to_local>;
   1515 defm cvta_to_shared : G_TO_NG<"shared", int_nvvm_ptr_gen_to_shared>;
   1516 defm cvta_to_global : G_TO_NG<"global", int_nvvm_ptr_gen_to_global>;
   1517 
   1518 def cvta_const : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
   1519                "mov.u32 \t$result, $src;",
   1520      [(set Int32Regs:$result, (int_nvvm_ptr_constant_to_gen Int32Regs:$src))]>;
   1521 def cvta_const_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
   1522                "mov.u64 \t$result, $src;",
   1523      [(set Int64Regs:$result, (int_nvvm_ptr_constant_to_gen Int64Regs:$src))]>;
   1524 
   1525 
   1526 
   1527 // @TODO: Revisit this.  There is a type
   1528 // contradiction between iPTRAny and iPTR for the def.
   1529 /*def cvta_const_addr : NVPTXInst<(outs Int32Regs:$result), (ins imemAny:$src),
   1530                "mov.u32 \t$result, $src;",
   1531      [(set Int32Regs:$result, (int_nvvm_ptr_constant_to_gen
   1532      (Wrapper tglobaladdr:$src)))]>;
   1533 def cvta_const_addr_64 : NVPTXInst<(outs Int64Regs:$result), (ins imemAny:$src),
   1534                "mov.u64 \t$result, $src;",
   1535      [(set Int64Regs:$result, (int_nvvm_ptr_constant_to_gen
   1536      (Wrapper tglobaladdr:$src)))]>;*/
   1537 
   1538 
   1539 def cvta_to_const : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
   1540             "mov.u32 \t$result, $src;",
   1541      [(set Int32Regs:$result, (int_nvvm_ptr_gen_to_constant Int32Regs:$src))]>;
   1542 def cvta_to_const_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
   1543             "mov.u64 \t$result, $src;",
   1544      [(set Int64Regs:$result, (int_nvvm_ptr_gen_to_constant Int64Regs:$src))]>;
   1545 
   1546 
   1547 // nvvm.ptr.gen.to.param
   1548 def nvvm_ptr_gen_to_param : NVPTXInst<(outs Int32Regs:$result),
   1549   (ins Int32Regs:$src),
   1550                         "mov.u32 \t$result, $src;",
   1551                               [(set Int32Regs:$result,
   1552                                 (int_nvvm_ptr_gen_to_param Int32Regs:$src))]>;
   1553 def nvvm_ptr_gen_to_param_64 : NVPTXInst<(outs Int64Regs:$result),
   1554   (ins Int64Regs:$src),
   1555                         "mov.u64 \t$result, $src;",
   1556                               [(set Int64Regs:$result,
   1557                                 (int_nvvm_ptr_gen_to_param Int64Regs:$src))]>;
   1558 
   1559 
   1560 // nvvm.move intrinsicc
   1561 def nvvm_move_i8 : NVPTXInst<(outs Int8Regs:$r), (ins Int8Regs:$s),
   1562                              "mov.b16 \t$r, $s;",
   1563                              [(set Int8Regs:$r,
   1564                                (int_nvvm_move_i8 Int8Regs:$s))]>;
   1565 def nvvm_move_i16 : NVPTXInst<(outs Int16Regs:$r), (ins Int16Regs:$s),
   1566                              "mov.b16 \t$r, $s;",
   1567                              [(set Int16Regs:$r,
   1568                                (int_nvvm_move_i16 Int16Regs:$s))]>;
   1569 def nvvm_move_i32 : NVPTXInst<(outs Int32Regs:$r), (ins Int32Regs:$s),
   1570                              "mov.b32 \t$r, $s;",
   1571                              [(set Int32Regs:$r,
   1572                                (int_nvvm_move_i32 Int32Regs:$s))]>;
   1573 def nvvm_move_i64 : NVPTXInst<(outs Int64Regs:$r), (ins Int64Regs:$s),
   1574                              "mov.b64 \t$r, $s;",
   1575                              [(set Int64Regs:$r,
   1576                                (int_nvvm_move_i64 Int64Regs:$s))]>;
   1577 def nvvm_move_float : NVPTXInst<(outs Float32Regs:$r), (ins Float32Regs:$s),
   1578                              "mov.f32 \t$r, $s;",
   1579                              [(set Float32Regs:$r,
   1580                                (int_nvvm_move_float Float32Regs:$s))]>;
   1581 def nvvm_move_double : NVPTXInst<(outs Float64Regs:$r), (ins Float64Regs:$s),
   1582                              "mov.f64 \t$r, $s;",
   1583                              [(set Float64Regs:$r,
   1584                                (int_nvvm_move_double Float64Regs:$s))]>;
   1585 def nvvm_move_ptr32 : NVPTXInst<(outs Int32Regs:$r), (ins Int32Regs:$s),
   1586                              "mov.u32 \t$r, $s;",
   1587                              [(set Int32Regs:$r,
   1588                                (int_nvvm_move_ptr Int32Regs:$s))]>;
   1589 def nvvm_move_ptr64 : NVPTXInst<(outs Int64Regs:$r), (ins Int64Regs:$s),
   1590                              "mov.u64 \t$r, $s;",
   1591                              [(set Int64Regs:$r,
   1592                                (int_nvvm_move_ptr Int64Regs:$s))]>;
   1593 
   1594 // @TODO: Are these actually needed, or will we always just see symbols
   1595 // copied to registers first?
   1596 /*def nvvm_move_sym32 : NVPTXInst<(outs Int32Regs:$r), (ins imem:$s),
   1597                              "mov.u32 \t$r, $s;",
   1598                              [(set Int32Regs:$r,
   1599                              (int_nvvm_move_ptr texternalsym:$s))]>;
   1600 def nvvm_move_sym64 : NVPTXInst<(outs Int64Regs:$r), (ins imem:$s),
   1601                              "mov.u64 \t$r, $s;",
   1602                              [(set Int64Regs:$r,
   1603                              (int_nvvm_move_ptr texternalsym:$s))]>;*/
   1604 
   1605 
   1606 // MoveParam        %r1, param
   1607 // ptr_local_to_gen %r2, %r1
   1608 // ptr_gen_to_local %r3, %r2
   1609 // ->
   1610 // mov %r1, param
   1611 
   1612 // @TODO: Revisit this.  There is a type
   1613 // contradiction between iPTRAny and iPTR for the addr defs, so the move_sym
   1614 // instructions are not currently defined. However, we can use the ptr
   1615 // variants and the asm printer will do the right thing.
   1616 def : Pat<(i64 (int_nvvm_ptr_gen_to_local (int_nvvm_ptr_local_to_gen
   1617                 (MoveParam texternalsym:$src)))),
   1618                (nvvm_move_ptr64  texternalsym:$src)>;
   1619 def : Pat<(i32 (int_nvvm_ptr_gen_to_local (int_nvvm_ptr_local_to_gen
   1620                 (MoveParam texternalsym:$src)))),
   1621                (nvvm_move_ptr32  texternalsym:$src)>;
   1622 
   1623 
   1624 //-----------------------------------
   1625 // Compiler Error Warn
   1626 // - Just ignore them in codegen
   1627 //-----------------------------------
   1628 
   1629 def INT_NVVM_COMPILER_WARN_32 : NVPTXInst<(outs), (ins Int32Regs:$a),
   1630                 "// llvm.nvvm.compiler.warn()",
   1631                 [(int_nvvm_compiler_warn Int32Regs:$a)]>;
   1632 def INT_NVVM_COMPILER_WARN_64 : NVPTXInst<(outs), (ins Int64Regs:$a),
   1633                 "// llvm.nvvm.compiler.warn()",
   1634                 [(int_nvvm_compiler_warn Int64Regs:$a)]>;
   1635 def INT_NVVM_COMPILER_ERROR_32 : NVPTXInst<(outs), (ins Int32Regs:$a),
   1636                 "// llvm.nvvm.compiler.error()",
   1637                 [(int_nvvm_compiler_error Int32Regs:$a)]>;
   1638 def INT_NVVM_COMPILER_ERROR_64 : NVPTXInst<(outs), (ins Int64Regs:$a),
   1639                 "// llvm.nvvm.compiler.error()",
   1640                 [(int_nvvm_compiler_error Int64Regs:$a)]>;
   1641 
   1642 
   1643 
   1644 //===-- Old PTX Back-end Intrinsics ---------------------------------------===//
   1645 
   1646 // These intrinsics are handled to retain compatibility with the old backend.
   1647 
   1648 // PTX Special Purpose Register Accessor Intrinsics
   1649 
   1650 class PTX_READ_SPECIAL_REGISTER_R64<string regname, Intrinsic intop>
   1651   : NVPTXInst<(outs Int64Regs:$d), (ins),
   1652               !strconcat(!strconcat("mov.u64\t$d, %", regname), ";"),
   1653               [(set Int64Regs:$d, (intop))]>;
   1654 
   1655 class PTX_READ_SPECIAL_REGISTER_R32<string regname, Intrinsic intop>
   1656   : NVPTXInst<(outs Int32Regs:$d), (ins),
   1657               !strconcat(!strconcat("mov.u32\t$d, %", regname), ";"),
   1658               [(set Int32Regs:$d, (intop))]>;
   1659 
   1660 // TODO Add read vector-version of special registers
   1661 
   1662 def PTX_READ_TID_X   : PTX_READ_SPECIAL_REGISTER_R32<"tid.x",
   1663                                                      int_ptx_read_tid_x>;
   1664 def PTX_READ_TID_Y   : PTX_READ_SPECIAL_REGISTER_R32<"tid.y",
   1665                                                      int_ptx_read_tid_y>;
   1666 def PTX_READ_TID_Z   : PTX_READ_SPECIAL_REGISTER_R32<"tid.z",
   1667                                                      int_ptx_read_tid_z>;
   1668 def PTX_READ_TID_W   : PTX_READ_SPECIAL_REGISTER_R32<"tid.w",
   1669                                                      int_ptx_read_tid_w>;
   1670 
   1671 def PTX_READ_NTID_X   : PTX_READ_SPECIAL_REGISTER_R32<"ntid.x",
   1672                                                       int_ptx_read_ntid_x>;
   1673 def PTX_READ_NTID_Y   : PTX_READ_SPECIAL_REGISTER_R32<"ntid.y",
   1674                                                       int_ptx_read_ntid_y>;
   1675 def PTX_READ_NTID_Z   : PTX_READ_SPECIAL_REGISTER_R32<"ntid.z",
   1676                                                       int_ptx_read_ntid_z>;
   1677 def PTX_READ_NTID_W   : PTX_READ_SPECIAL_REGISTER_R32<"ntid.w",
   1678                                                       int_ptx_read_ntid_w>;
   1679 
   1680 def PTX_READ_LANEID  : PTX_READ_SPECIAL_REGISTER_R32<"laneid",
   1681                                                      int_ptx_read_laneid>;
   1682 def PTX_READ_WARPID  : PTX_READ_SPECIAL_REGISTER_R32<"warpid",
   1683                                                      int_ptx_read_warpid>;
   1684 def PTX_READ_NWARPID : PTX_READ_SPECIAL_REGISTER_R32<"nwarpid",
   1685                                                      int_ptx_read_nwarpid>;
   1686 
   1687 def PTX_READ_CTAID_X   : PTX_READ_SPECIAL_REGISTER_R32<"ctaid.x",
   1688                                                        int_ptx_read_ctaid_x>;
   1689 def PTX_READ_CTAID_Y   : PTX_READ_SPECIAL_REGISTER_R32<"ctaid.y",
   1690                                                        int_ptx_read_ctaid_y>;
   1691 def PTX_READ_CTAID_Z   : PTX_READ_SPECIAL_REGISTER_R32<"ctaid.z",
   1692                                                        int_ptx_read_ctaid_z>;
   1693 def PTX_READ_CTAID_W   : PTX_READ_SPECIAL_REGISTER_R32<"ctaid.w",
   1694                                                        int_ptx_read_ctaid_w>;
   1695 
   1696 def PTX_READ_NCTAID_X   : PTX_READ_SPECIAL_REGISTER_R32<"nctaid.x",
   1697                                                         int_ptx_read_nctaid_x>;
   1698 def PTX_READ_NCTAID_Y   : PTX_READ_SPECIAL_REGISTER_R32<"nctaid.y",
   1699                                                         int_ptx_read_nctaid_y>;
   1700 def PTX_READ_NCTAID_Z   : PTX_READ_SPECIAL_REGISTER_R32<"nctaid.z",
   1701                                                         int_ptx_read_nctaid_z>;
   1702 def PTX_READ_NCTAID_W   : PTX_READ_SPECIAL_REGISTER_R32<"nctaid.w",
   1703                                                         int_ptx_read_nctaid_w>;
   1704 
   1705 def PTX_READ_SMID  : PTX_READ_SPECIAL_REGISTER_R32<"smid",
   1706                                                    int_ptx_read_smid>;
   1707 def PTX_READ_NSMID  : PTX_READ_SPECIAL_REGISTER_R32<"nsmid",
   1708                                                     int_ptx_read_nsmid>;
   1709 def PTX_READ_GRIDID  : PTX_READ_SPECIAL_REGISTER_R32<"gridid",
   1710                                                      int_ptx_read_gridid>;
   1711 
   1712 def PTX_READ_LANEMASK_EQ
   1713   : PTX_READ_SPECIAL_REGISTER_R32<"lanemask_eq", int_ptx_read_lanemask_eq>;
   1714 def PTX_READ_LANEMASK_LE
   1715   : PTX_READ_SPECIAL_REGISTER_R32<"lanemask_le", int_ptx_read_lanemask_le>;
   1716 def PTX_READ_LANEMASK_LT
   1717   : PTX_READ_SPECIAL_REGISTER_R32<"lanemask_lt", int_ptx_read_lanemask_lt>;
   1718 def PTX_READ_LANEMASK_GE
   1719   : PTX_READ_SPECIAL_REGISTER_R32<"lanemask_ge", int_ptx_read_lanemask_ge>;
   1720 def PTX_READ_LANEMASK_GT
   1721   : PTX_READ_SPECIAL_REGISTER_R32<"lanemask_gt", int_ptx_read_lanemask_gt>;
   1722 
   1723 def PTX_READ_CLOCK
   1724   : PTX_READ_SPECIAL_REGISTER_R32<"clock", int_ptx_read_clock>;
   1725 def PTX_READ_CLOCK64
   1726   : PTX_READ_SPECIAL_REGISTER_R64<"clock64", int_ptx_read_clock64>;
   1727 
   1728 def PTX_READ_PM0 : PTX_READ_SPECIAL_REGISTER_R32<"pm0", int_ptx_read_pm0>;
   1729 def PTX_READ_PM1 : PTX_READ_SPECIAL_REGISTER_R32<"pm1", int_ptx_read_pm1>;
   1730 def PTX_READ_PM2 : PTX_READ_SPECIAL_REGISTER_R32<"pm2", int_ptx_read_pm2>;
   1731 def PTX_READ_PM3 : PTX_READ_SPECIAL_REGISTER_R32<"pm3", int_ptx_read_pm3>;
   1732 
   1733 // PTX Parallel Synchronization and Communication Intrinsics
   1734 
   1735 def PTX_BAR_SYNC : NVPTXInst<(outs), (ins i32imm:$i), "bar.sync\t$i;",
   1736                              [(int_ptx_bar_sync imm:$i)]>;
   1737