contrib/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td

   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 // nvvm_sqrt intrinsic
 516 def : Pat<(int_nvvm_sqrt_f Float32Regs:$a),
 517           (INT_NVVM_SQRT_RN_FTZ_F Float32Regs:$a)>, Requires<[doF32FTZ, do_SQRTF32_RN]>;
 518 def : Pat<(int_nvvm_sqrt_f Float32Regs:$a),
 519           (INT_NVVM_SQRT_RN_F Float32Regs:$a)>, Requires<[do_SQRTF32_RN]>;
 520 def : Pat<(int_nvvm_sqrt_f Float32Regs:$a),
 521           (INT_NVVM_SQRT_APPROX_FTZ_F Float32Regs:$a)>, Requires<[doF32FTZ]>;
 522 def : Pat<(int_nvvm_sqrt_f Float32Regs:$a),
 523           (INT_NVVM_SQRT_APPROX_F Float32Regs:$a)>;
 524
 525 //
 526 // Rsqrt
 527 //
 528
 529 def INT_NVVM_RSQRT_APPROX_FTZ_F
 530   : F_MATH_1<"rsqrt.approx.ftz.f32 \t$dst, $src0;", Float32Regs, Float32Regs,
 531     int_nvvm_rsqrt_approx_ftz_f>;
 532 def INT_NVVM_RSQRT_APPROX_F : F_MATH_1<"rsqrt.approx.f32 \t$dst, $src0;",
 533   Float32Regs, Float32Regs, int_nvvm_rsqrt_approx_f>;
 534 def INT_NVVM_RSQRT_APPROX_D : F_MATH_1<"rsqrt.approx.f64 \t$dst, $src0;",
 535   Float64Regs, Float64Regs, int_nvvm_rsqrt_approx_d>;
 536
 537 //
 538 // Add
 539 //
 540
 541 def INT_NVVM_ADD_RN_FTZ_F : F_MATH_2<"add.rn.ftz.f32 \t$dst, $src0, $src1;",
 542   Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rn_ftz_f>;
 543 def INT_NVVM_ADD_RN_F : F_MATH_2<"add.rn.f32 \t$dst, $src0, $src1;",
 544   Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rn_f>;
 545 def INT_NVVM_ADD_RZ_FTZ_F : F_MATH_2<"add.rz.ftz.f32 \t$dst, $src0, $src1;",
 546   Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rz_ftz_f>;
 547 def INT_NVVM_ADD_RZ_F : F_MATH_2<"add.rz.f32 \t$dst, $src0, $src1;",
 548   Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rz_f>;
 549 def INT_NVVM_ADD_RM_FTZ_F : F_MATH_2<"add.rm.ftz.f32 \t$dst, $src0, $src1;",
 550   Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rm_ftz_f>;
 551 def INT_NVVM_ADD_RM_F : F_MATH_2<"add.rm.f32 \t$dst, $src0, $src1;",
 552   Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rm_f>;
 553 def INT_NVVM_ADD_RP_FTZ_F : F_MATH_2<"add.rp.ftz.f32 \t$dst, $src0, $src1;",
 554   Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rp_ftz_f>;
 555 def INT_NVVM_ADD_RP_F : F_MATH_2<"add.rp.f32 \t$dst, $src0, $src1;",
 556   Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rp_f>;
 557
 558 def INT_NVVM_ADD_RN_D : F_MATH_2<"add.rn.f64 \t$dst, $src0, $src1;",
 559   Float64Regs, Float64Regs, Float64Regs, int_nvvm_add_rn_d>;
 560 def INT_NVVM_ADD_RZ_D : F_MATH_2<"add.rz.f64 \t$dst, $src0, $src1;",
 561   Float64Regs, Float64Regs, Float64Regs, int_nvvm_add_rz_d>;
 562 def INT_NVVM_ADD_RM_D : F_MATH_2<"add.rm.f64 \t$dst, $src0, $src1;",
 563   Float64Regs, Float64Regs, Float64Regs, int_nvvm_add_rm_d>;
 564 def INT_NVVM_ADD_RP_D : F_MATH_2<"add.rp.f64 \t$dst, $src0, $src1;",
 565   Float64Regs, Float64Regs, Float64Regs, int_nvvm_add_rp_d>;
 566
 567 //
 568 // Convert
 569 //
 570
 571 def INT_NVVM_D2F_RN_FTZ : F_MATH_1<"cvt.rn.ftz.f32.f64 \t$dst, $src0;",
 572   Float32Regs, Float64Regs, int_nvvm_d2f_rn_ftz>;
 573 def INT_NVVM_D2F_RN : F_MATH_1<"cvt.rn.f32.f64 \t$dst, $src0;",
 574   Float32Regs, Float64Regs, int_nvvm_d2f_rn>;
 575 def INT_NVVM_D2F_RZ_FTZ : F_MATH_1<"cvt.rz.ftz.f32.f64 \t$dst, $src0;",
 576   Float32Regs, Float64Regs, int_nvvm_d2f_rz_ftz>;
 577 def INT_NVVM_D2F_RZ : F_MATH_1<"cvt.rz.f32.f64 \t$dst, $src0;",
 578   Float32Regs, Float64Regs, int_nvvm_d2f_rz>;
 579 def INT_NVVM_D2F_RM_FTZ : F_MATH_1<"cvt.rm.ftz.f32.f64 \t$dst, $src0;",
 580   Float32Regs, Float64Regs, int_nvvm_d2f_rm_ftz>;
 581 def INT_NVVM_D2F_RM : F_MATH_1<"cvt.rm.f32.f64 \t$dst, $src0;",
 582   Float32Regs, Float64Regs, int_nvvm_d2f_rm>;
 583 def INT_NVVM_D2F_RP_FTZ : F_MATH_1<"cvt.rp.ftz.f32.f64 \t$dst, $src0;",
 584   Float32Regs, Float64Regs, int_nvvm_d2f_rp_ftz>;
 585 def INT_NVVM_D2F_RP : F_MATH_1<"cvt.rp.f32.f64 \t$dst, $src0;",
 586   Float32Regs, Float64Regs, int_nvvm_d2f_rp>;
 587
 588 def INT_NVVM_D2I_RN : F_MATH_1<"cvt.rni.s32.f64 \t$dst, $src0;",
 589   Int32Regs, Float64Regs, int_nvvm_d2i_rn>;
 590 def INT_NVVM_D2I_RZ : F_MATH_1<"cvt.rzi.s32.f64 \t$dst, $src0;",
 591   Int32Regs, Float64Regs, int_nvvm_d2i_rz>;
 592 def INT_NVVM_D2I_RM : F_MATH_1<"cvt.rmi.s32.f64 \t$dst, $src0;",
 593   Int32Regs, Float64Regs, int_nvvm_d2i_rm>;
 594 def INT_NVVM_D2I_RP : F_MATH_1<"cvt.rpi.s32.f64 \t$dst, $src0;",
 595   Int32Regs, Float64Regs, int_nvvm_d2i_rp>;
 596
 597 def INT_NVVM_D2UI_RN : F_MATH_1<"cvt.rni.u32.f64 \t$dst, $src0;",
 598   Int32Regs, Float64Regs, int_nvvm_d2ui_rn>;
 599 def INT_NVVM_D2UI_RZ : F_MATH_1<"cvt.rzi.u32.f64 \t$dst, $src0;",
 600   Int32Regs, Float64Regs, int_nvvm_d2ui_rz>;
 601 def INT_NVVM_D2UI_RM : F_MATH_1<"cvt.rmi.u32.f64 \t$dst, $src0;",
 602   Int32Regs, Float64Regs, int_nvvm_d2ui_rm>;
 603 def INT_NVVM_D2UI_RP : F_MATH_1<"cvt.rpi.u32.f64 \t$dst, $src0;",
 604   Int32Regs, Float64Regs, int_nvvm_d2ui_rp>;
 605
 606 def INT_NVVM_I2D_RN : F_MATH_1<"cvt.rn.f64.s32 \t$dst, $src0;",
 607   Float64Regs, Int32Regs, int_nvvm_i2d_rn>;
 608 def INT_NVVM_I2D_RZ : F_MATH_1<"cvt.rz.f64.s32 \t$dst, $src0;",
 609   Float64Regs, Int32Regs, int_nvvm_i2d_rz>;
 610 def INT_NVVM_I2D_RM : F_MATH_1<"cvt.rm.f64.s32 \t$dst, $src0;",
 611   Float64Regs, Int32Regs, int_nvvm_i2d_rm>;
 612 def INT_NVVM_I2D_RP : F_MATH_1<"cvt.rp.f64.s32 \t$dst, $src0;",
 613   Float64Regs, Int32Regs, int_nvvm_i2d_rp>;
 614
 615 def INT_NVVM_UI2D_RN : F_MATH_1<"cvt.rn.f64.u32 \t$dst, $src0;",
 616   Float64Regs, Int32Regs, int_nvvm_ui2d_rn>;
 617 def INT_NVVM_UI2D_RZ : F_MATH_1<"cvt.rz.f64.u32 \t$dst, $src0;",
 618   Float64Regs, Int32Regs, int_nvvm_ui2d_rz>;
 619 def INT_NVVM_UI2D_RM : F_MATH_1<"cvt.rm.f64.u32 \t$dst, $src0;",
 620   Float64Regs, Int32Regs, int_nvvm_ui2d_rm>;
 621 def INT_NVVM_UI2D_RP : F_MATH_1<"cvt.rp.f64.u32 \t$dst, $src0;",
 622   Float64Regs, Int32Regs, int_nvvm_ui2d_rp>;
 623
 624 def INT_NVVM_F2I_RN_FTZ : F_MATH_1<"cvt.rni.ftz.s32.f32 \t$dst, $src0;",
 625   Int32Regs, Float32Regs, int_nvvm_f2i_rn_ftz>;
 626 def INT_NVVM_F2I_RN : F_MATH_1<"cvt.rni.s32.f32 \t$dst, $src0;", Int32Regs,
 627   Float32Regs, int_nvvm_f2i_rn>;
 628 def INT_NVVM_F2I_RZ_FTZ : F_MATH_1<"cvt.rzi.ftz.s32.f32 \t$dst, $src0;",
 629   Int32Regs, Float32Regs, int_nvvm_f2i_rz_ftz>;
 630 def INT_NVVM_F2I_RZ : F_MATH_1<"cvt.rzi.s32.f32 \t$dst, $src0;", Int32Regs,
 631   Float32Regs, int_nvvm_f2i_rz>;
 632 def INT_NVVM_F2I_RM_FTZ : F_MATH_1<"cvt.rmi.ftz.s32.f32 \t$dst, $src0;",
 633   Int32Regs, Float32Regs, int_nvvm_f2i_rm_ftz>;
 634 def INT_NVVM_F2I_RM : F_MATH_1<"cvt.rmi.s32.f32 \t$dst, $src0;", Int32Regs,
 635   Float32Regs, int_nvvm_f2i_rm>;
 636 def INT_NVVM_F2I_RP_FTZ : F_MATH_1<"cvt.rpi.ftz.s32.f32 \t$dst, $src0;",
 637   Int32Regs, Float32Regs, int_nvvm_f2i_rp_ftz>;
 638 def INT_NVVM_F2I_RP : F_MATH_1<"cvt.rpi.s32.f32 \t$dst, $src0;", Int32Regs,
 639   Float32Regs, int_nvvm_f2i_rp>;
 640
 641 def INT_NVVM_F2UI_RN_FTZ : F_MATH_1<"cvt.rni.ftz.u32.f32 \t$dst, $src0;",
 642   Int32Regs, Float32Regs, int_nvvm_f2ui_rn_ftz>;
 643 def INT_NVVM_F2UI_RN : F_MATH_1<"cvt.rni.u32.f32 \t$dst, $src0;", Int32Regs,
 644   Float32Regs, int_nvvm_f2ui_rn>;
 645 def INT_NVVM_F2UI_RZ_FTZ : F_MATH_1<"cvt.rzi.ftz.u32.f32 \t$dst, $src0;",
 646   Int32Regs, Float32Regs, int_nvvm_f2ui_rz_ftz>;
 647 def INT_NVVM_F2UI_RZ : F_MATH_1<"cvt.rzi.u32.f32 \t$dst, $src0;", Int32Regs,
 648   Float32Regs, int_nvvm_f2ui_rz>;
 649 def INT_NVVM_F2UI_RM_FTZ : F_MATH_1<"cvt.rmi.ftz.u32.f32 \t$dst, $src0;",
 650   Int32Regs, Float32Regs, int_nvvm_f2ui_rm_ftz>;
 651 def INT_NVVM_F2UI_RM : F_MATH_1<"cvt.rmi.u32.f32 \t$dst, $src0;", Int32Regs,
 652   Float32Regs, int_nvvm_f2ui_rm>;
 653 def INT_NVVM_F2UI_RP_FTZ : F_MATH_1<"cvt.rpi.ftz.u32.f32 \t$dst, $src0;",
 654   Int32Regs, Float32Regs, int_nvvm_f2ui_rp_ftz>;
 655 def INT_NVVM_F2UI_RP : F_MATH_1<"cvt.rpi.u32.f32 \t$dst, $src0;", Int32Regs,
 656   Float32Regs, int_nvvm_f2ui_rp>;
 657
 658 def INT_NVVM_I2F_RN : F_MATH_1<"cvt.rn.f32.s32 \t$dst, $src0;", Float32Regs,
 659   Int32Regs, int_nvvm_i2f_rn>;
 660 def INT_NVVM_I2F_RZ : F_MATH_1<"cvt.rz.f32.s32 \t$dst, $src0;", Float32Regs,
 661   Int32Regs, int_nvvm_i2f_rz>;
 662 def INT_NVVM_I2F_RM : F_MATH_1<"cvt.rm.f32.s32 \t$dst, $src0;", Float32Regs,
 663   Int32Regs, int_nvvm_i2f_rm>;
 664 def INT_NVVM_I2F_RP : F_MATH_1<"cvt.rp.f32.s32 \t$dst, $src0;", Float32Regs,
 665   Int32Regs, int_nvvm_i2f_rp>;
 666
 667 def INT_NVVM_UI2F_RN : F_MATH_1<"cvt.rn.f32.u32 \t$dst, $src0;", Float32Regs,
 668   Int32Regs, int_nvvm_ui2f_rn>;
 669 def INT_NVVM_UI2F_RZ : F_MATH_1<"cvt.rz.f32.u32 \t$dst, $src0;", Float32Regs,
 670   Int32Regs, int_nvvm_ui2f_rz>;
 671 def INT_NVVM_UI2F_RM : F_MATH_1<"cvt.rm.f32.u32 \t$dst, $src0;", Float32Regs,
 672   Int32Regs, int_nvvm_ui2f_rm>;
 673 def INT_NVVM_UI2F_RP : F_MATH_1<"cvt.rp.f32.u32 \t$dst, $src0;", Float32Regs,
 674   Int32Regs, int_nvvm_ui2f_rp>;
 675
 676 def INT_NVVM_LOHI_I2D : F_MATH_2<"mov.b64 \t$dst, {{$src0, $src1}};",
 677   Float64Regs, Int32Regs, Int32Regs, int_nvvm_lohi_i2d>;
 678
 679 def INT_NVVM_D2I_LO : F_MATH_1<!strconcat("{{\n\t",
 680                        !strconcat(".reg .b32 %temp; \n\t",
 681              !strconcat("mov.b64 \t{$dst, %temp}, $src0;\n\t",
 682                "}}"))),
 683              Int32Regs, Float64Regs, int_nvvm_d2i_lo>;
 684 def INT_NVVM_D2I_HI : F_MATH_1<!strconcat("{{\n\t",
 685                        !strconcat(".reg .b32 %temp; \n\t",
 686                          !strconcat("mov.b64 \t{%temp, $dst}, $src0;\n\t",
 687                            "}}"))),
 688              Int32Regs, Float64Regs, int_nvvm_d2i_hi>;
 689
 690 def INT_NVVM_F2LL_RN_FTZ : F_MATH_1<"cvt.rni.ftz.s64.f32 \t$dst, $src0;",
 691   Int64Regs, Float32Regs, int_nvvm_f2ll_rn_ftz>;
 692 def INT_NVVM_F2LL_RN : F_MATH_1<"cvt.rni.s64.f32 \t$dst, $src0;", Int64Regs,
 693   Float32Regs, int_nvvm_f2ll_rn>;
 694 def INT_NVVM_F2LL_RZ_FTZ : F_MATH_1<"cvt.rzi.ftz.s64.f32 \t$dst, $src0;",
 695   Int64Regs, Float32Regs, int_nvvm_f2ll_rz_ftz>;
 696 def INT_NVVM_F2LL_RZ : F_MATH_1<"cvt.rzi.s64.f32 \t$dst, $src0;", Int64Regs,
 697   Float32Regs, int_nvvm_f2ll_rz>;
 698 def INT_NVVM_F2LL_RM_FTZ : F_MATH_1<"cvt.rmi.ftz.s64.f32 \t$dst, $src0;",
 699   Int64Regs, Float32Regs, int_nvvm_f2ll_rm_ftz>;
 700 def INT_NVVM_F2LL_RM : F_MATH_1<"cvt.rmi.s64.f32 \t$dst, $src0;", Int64Regs,
 701   Float32Regs, int_nvvm_f2ll_rm>;
 702 def INT_NVVM_F2LL_RP_FTZ : F_MATH_1<"cvt.rpi.ftz.s64.f32 \t$dst, $src0;",
 703   Int64Regs, Float32Regs, int_nvvm_f2ll_rp_ftz>;
 704 def INT_NVVM_F2LL_RP : F_MATH_1<"cvt.rpi.s64.f32 \t$dst, $src0;", Int64Regs,
 705   Float32Regs, int_nvvm_f2ll_rp>;
 706
 707 def INT_NVVM_F2ULL_RN_FTZ : F_MATH_1<"cvt.rni.ftz.u64.f32 \t$dst, $src0;",
 708   Int64Regs, Float32Regs, int_nvvm_f2ull_rn_ftz>;
 709 def INT_NVVM_F2ULL_RN : F_MATH_1<"cvt.rni.u64.f32 \t$dst, $src0;", Int64Regs,
 710   Float32Regs, int_nvvm_f2ull_rn>;
 711 def INT_NVVM_F2ULL_RZ_FTZ : F_MATH_1<"cvt.rzi.ftz.u64.f32 \t$dst, $src0;",
 712   Int64Regs, Float32Regs, int_nvvm_f2ull_rz_ftz>;
 713 def INT_NVVM_F2ULL_RZ : F_MATH_1<"cvt.rzi.u64.f32 \t$dst, $src0;", Int64Regs,
 714   Float32Regs, int_nvvm_f2ull_rz>;
 715 def INT_NVVM_F2ULL_RM_FTZ : F_MATH_1<"cvt.rmi.ftz.u64.f32 \t$dst, $src0;",
 716   Int64Regs, Float32Regs, int_nvvm_f2ull_rm_ftz>;
 717 def INT_NVVM_F2ULL_RM : F_MATH_1<"cvt.rmi.u64.f32 \t$dst, $src0;", Int64Regs,
 718   Float32Regs, int_nvvm_f2ull_rm>;
 719 def INT_NVVM_F2ULL_RP_FTZ : F_MATH_1<"cvt.rpi.ftz.u64.f32 \t$dst, $src0;",
 720   Int64Regs, Float32Regs, int_nvvm_f2ull_rp_ftz>;
 721 def INT_NVVM_F2ULL_RP : F_MATH_1<"cvt.rpi.u64.f32 \t$dst, $src0;", Int64Regs,
 722   Float32Regs, int_nvvm_f2ull_rp>;
 723
 724 def INT_NVVM_D2LL_RN : F_MATH_1<"cvt.rni.s64.f64 \t$dst, $src0;", Int64Regs,
 725   Float64Regs, int_nvvm_d2ll_rn>;
 726 def INT_NVVM_D2LL_RZ : F_MATH_1<"cvt.rzi.s64.f64 \t$dst, $src0;", Int64Regs,
 727   Float64Regs, int_nvvm_d2ll_rz>;
 728 def INT_NVVM_D2LL_RM : F_MATH_1<"cvt.rmi.s64.f64 \t$dst, $src0;", Int64Regs,
 729   Float64Regs, int_nvvm_d2ll_rm>;
 730 def INT_NVVM_D2LL_RP : F_MATH_1<"cvt.rpi.s64.f64 \t$dst, $src0;", Int64Regs,
 731   Float64Regs, int_nvvm_d2ll_rp>;
 732
 733 def INT_NVVM_D2ULL_RN : F_MATH_1<"cvt.rni.u64.f64 \t$dst, $src0;", Int64Regs,
 734   Float64Regs, int_nvvm_d2ull_rn>;
 735 def INT_NVVM_D2ULL_RZ : F_MATH_1<"cvt.rzi.u64.f64 \t$dst, $src0;", Int64Regs,
 736   Float64Regs, int_nvvm_d2ull_rz>;
 737 def INT_NVVM_D2ULL_RM : F_MATH_1<"cvt.rmi.u64.f64 \t$dst, $src0;", Int64Regs,
 738   Float64Regs, int_nvvm_d2ull_rm>;
 739 def INT_NVVM_D2ULL_RP : F_MATH_1<"cvt.rpi.u64.f64 \t$dst, $src0;", Int64Regs,
 740   Float64Regs, int_nvvm_d2ull_rp>;
 741
 742 def INT_NVVM_LL2F_RN : F_MATH_1<"cvt.rn.f32.s64 \t$dst, $src0;", Float32Regs,
 743   Int64Regs, int_nvvm_ll2f_rn>;
 744 def INT_NVVM_LL2F_RZ : F_MATH_1<"cvt.rz.f32.s64 \t$dst, $src0;", Float32Regs,
 745   Int64Regs, int_nvvm_ll2f_rz>;
 746 def INT_NVVM_LL2F_RM : F_MATH_1<"cvt.rm.f32.s64 \t$dst, $src0;", Float32Regs,
 747   Int64Regs, int_nvvm_ll2f_rm>;
 748 def INT_NVVM_LL2F_RP : F_MATH_1<"cvt.rp.f32.s64 \t$dst, $src0;", Float32Regs,
 749   Int64Regs, int_nvvm_ll2f_rp>;
 750 def INT_NVVM_ULL2F_RN : F_MATH_1<"cvt.rn.f32.u64 \t$dst, $src0;", Float32Regs,
 751   Int64Regs, int_nvvm_ull2f_rn>;
 752 def INT_NVVM_ULL2F_RZ : F_MATH_1<"cvt.rz.f32.u64 \t$dst, $src0;", Float32Regs,
 753   Int64Regs, int_nvvm_ull2f_rz>;
 754 def INT_NVVM_ULL2F_RM : F_MATH_1<"cvt.rm.f32.u64 \t$dst, $src0;", Float32Regs,
 755   Int64Regs, int_nvvm_ull2f_rm>;
 756 def INT_NVVM_ULL2F_RP : F_MATH_1<"cvt.rp.f32.u64 \t$dst, $src0;", Float32Regs,
 757   Int64Regs, int_nvvm_ull2f_rp>;
 758
 759 def INT_NVVM_LL2D_RN : F_MATH_1<"cvt.rn.f64.s64 \t$dst, $src0;", Float64Regs,
 760   Int64Regs, int_nvvm_ll2d_rn>;
 761 def INT_NVVM_LL2D_RZ : F_MATH_1<"cvt.rz.f64.s64 \t$dst, $src0;", Float64Regs,
 762   Int64Regs, int_nvvm_ll2d_rz>;
 763 def INT_NVVM_LL2D_RM : F_MATH_1<"cvt.rm.f64.s64 \t$dst, $src0;", Float64Regs,
 764   Int64Regs, int_nvvm_ll2d_rm>;
 765 def INT_NVVM_LL2D_RP : F_MATH_1<"cvt.rp.f64.s64 \t$dst, $src0;", Float64Regs,
 766   Int64Regs, int_nvvm_ll2d_rp>;
 767 def INT_NVVM_ULL2D_RN : F_MATH_1<"cvt.rn.f64.u64 \t$dst, $src0;", Float64Regs,
 768   Int64Regs, int_nvvm_ull2d_rn>;
 769 def INT_NVVM_ULL2D_RZ : F_MATH_1<"cvt.rz.f64.u64 \t$dst, $src0;", Float64Regs,
 770   Int64Regs, int_nvvm_ull2d_rz>;
 771 def INT_NVVM_ULL2D_RM : F_MATH_1<"cvt.rm.f64.u64 \t$dst, $src0;", Float64Regs,
 772   Int64Regs, int_nvvm_ull2d_rm>;
 773 def INT_NVVM_ULL2D_RP : F_MATH_1<"cvt.rp.f64.u64 \t$dst, $src0;", Float64Regs,
 774   Int64Regs, int_nvvm_ull2d_rp>;
 775
 776 def INT_NVVM_F2H_RN_FTZ : F_MATH_1<!strconcat("{{\n\t",
 777                                    !strconcat(".reg .b16 %temp;\n\t",
 778            !strconcat("cvt.rn.ftz.f16.f32 \t%temp, $src0;\n\t",
 779            !strconcat("mov.b16 \t$dst, %temp;\n",
 780              "}}")))),
 781                                    Int16Regs, Float32Regs, int_nvvm_f2h_rn_ftz>;
 782 def INT_NVVM_F2H_RN : F_MATH_1<!strconcat("{{\n\t",
 783                                    !strconcat(".reg .b16 %temp;\n\t",
 784            !strconcat("cvt.rn.f16.f32 \t%temp, $src0;\n\t",
 785            !strconcat("mov.b16 \t$dst, %temp;\n",
 786              "}}")))),
 787            Int16Regs, Float32Regs, int_nvvm_f2h_rn>;
 788
 789 def INT_NVVM_H2F : F_MATH_1<!strconcat("{{\n\t",
 790                             !strconcat(".reg .b16 %temp;\n\t",
 791           !strconcat("mov.b16 \t%temp, $src0;\n\t",
 792           !strconcat("cvt.f32.f16 \t$dst, %temp;\n\t",
 793             "}}")))),
 794           Float32Regs, Int16Regs, int_nvvm_h2f>;
 795
 796 //
 797 // Bitcast
 798 //
 799
 800 def INT_NVVM_BITCAST_F2I : F_MATH_1<"mov.b32 \t$dst, $src0;", Int32Regs,
 801   Float32Regs, int_nvvm_bitcast_f2i>;
 802 def INT_NVVM_BITCAST_I2F : F_MATH_1<"mov.b32 \t$dst, $src0;", Float32Regs,
 803   Int32Regs, int_nvvm_bitcast_i2f>;
 804
 805 def INT_NVVM_BITCAST_LL2D : F_MATH_1<"mov.b64 \t$dst, $src0;", Float64Regs,
 806   Int64Regs, int_nvvm_bitcast_ll2d>;
 807 def INT_NVVM_BITCAST_D2LL : F_MATH_1<"mov.b64 \t$dst, $src0;", Int64Regs,
 808   Float64Regs, int_nvvm_bitcast_d2ll>;
 809
 810 //-----------------------------------
 811 // Atomic Functions
 812 //-----------------------------------
 813
 814 class ATOMIC_GLOBAL_CHK <dag ops, dag frag>
 815  : PatFrag<ops, frag, [{
 816    return ChkMemSDNodeAddressSpace(N, llvm::ADDRESS_SPACE_GLOBAL);
 817 }]>;
 818 class ATOMIC_SHARED_CHK <dag ops, dag frag>
 819  : PatFrag<ops, frag, [{
 820    return ChkMemSDNodeAddressSpace(N, llvm::ADDRESS_SPACE_SHARED);
 821 }]>;
 822 class ATOMIC_GENERIC_CHK <dag ops, dag frag>
 823  : PatFrag<ops, frag, [{
 824    return ChkMemSDNodeAddressSpace(N, llvm::ADDRESS_SPACE_GENERIC);
 825 }]>;
 826
 827 multiclass F_ATOMIC_2_imp<NVPTXRegClass ptrclass, NVPTXRegClass regclass,
 828   string SpaceStr, string TypeStr, string OpcStr, PatFrag IntOp,
 829   Operand IMMType, SDNode IMM, Predicate Pred> {
 830   def reg : NVPTXInst<(outs regclass:$dst), (ins ptrclass:$addr, regclass:$b),
 831                !strconcat("atom",
 832          !strconcat(SpaceStr,
 833          !strconcat(OpcStr,
 834          !strconcat(TypeStr,
 835          !strconcat(" \t$dst, [$addr], $b;", ""))))),
 836          [(set regclass:$dst, (IntOp ptrclass:$addr, regclass:$b))]>,
 837   Requires<[Pred]>;
 838   def imm : NVPTXInst<(outs regclass:$dst), (ins ptrclass:$addr, IMMType:$b),
 839                !strconcat("atom",
 840          !strconcat(SpaceStr,
 841          !strconcat(OpcStr,
 842          !strconcat(TypeStr,
 843          !strconcat(" \t$dst, [$addr], $b;", ""))))),
 844          [(set regclass:$dst, (IntOp ptrclass:$addr, IMM:$b))]>,
 845   Requires<[Pred]>;
 846 }
 847 multiclass F_ATOMIC_2<NVPTXRegClass regclass, string SpaceStr, string TypeStr,
 848   string OpcStr, PatFrag IntOp, Operand IMMType, SDNode IMM, Predicate Pred> {
 849   defm p32 : F_ATOMIC_2_imp<Int32Regs, regclass, SpaceStr, TypeStr, OpcStr,
 850     IntOp, IMMType, IMM, Pred>;
 851   defm p64 : F_ATOMIC_2_imp<Int64Regs, regclass, SpaceStr, TypeStr, OpcStr,
 852     IntOp, IMMType, IMM, Pred>;
 853 }
 854
 855 // has 2 operands, neg the second one
 856 multiclass F_ATOMIC_2_NEG_imp<NVPTXRegClass ptrclass, NVPTXRegClass regclass,
 857   string SpaceStr, string TypeStr, string OpcStr, PatFrag IntOp,
 858   Operand IMMType, Predicate Pred> {
 859   def reg : NVPTXInst<(outs regclass:$dst), (ins ptrclass:$addr, regclass:$b),
 860     !strconcat("{{ \n\t",
 861          !strconcat(".reg \t.s",
 862          !strconcat(TypeStr,
 863          !strconcat(" temp; \n\t",
 864          !strconcat("neg.s",
 865          !strconcat(TypeStr,
 866          !strconcat(" \ttemp, $b; \n\t",
 867                !strconcat("atom",
 868          !strconcat(SpaceStr,
 869          !strconcat(OpcStr,
 870          !strconcat(".u",
 871          !strconcat(TypeStr,
 872          !strconcat(" \t$dst, [$addr], temp; \n\t",
 873            !strconcat("}}", "")))))))))))))),
 874          [(set regclass:$dst, (IntOp ptrclass:$addr, regclass:$b))]>,
 875   Requires<[Pred]>;
 876 }
 877 multiclass F_ATOMIC_2_NEG<NVPTXRegClass regclass, string SpaceStr,
 878   string TypeStr, string OpcStr, PatFrag IntOp, Operand IMMType,
 879   Predicate Pred> {
 880  defm p32: F_ATOMIC_2_NEG_imp<Int32Regs, regclass, SpaceStr, TypeStr, OpcStr,
 881    IntOp, IMMType, Pred> ;
 882  defm p64: F_ATOMIC_2_NEG_imp<Int64Regs, regclass, SpaceStr, TypeStr, OpcStr,
 883    IntOp, IMMType, Pred> ;
 884 }
 885
 886 // has 3 operands
 887 multiclass F_ATOMIC_3_imp<NVPTXRegClass ptrclass, NVPTXRegClass regclass,
 888   string SpaceStr, string TypeStr, string OpcStr, PatFrag IntOp,
 889   Operand IMMType, Predicate Pred> {
 890   def reg : NVPTXInst<(outs regclass:$dst),
 891     (ins ptrclass:$addr, regclass:$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,
 898            (IntOp ptrclass:$addr, regclass:$b, regclass:$c))]>,
 899          Requires<[Pred]>;
 900   def imm1 : NVPTXInst<(outs regclass:$dst),
 901     (ins ptrclass:$addr, IMMType:$b, regclass:$c),
 902                !strconcat("atom",
 903          !strconcat(SpaceStr,
 904          !strconcat(OpcStr,
 905          !strconcat(TypeStr,
 906          !strconcat(" \t$dst, [$addr], $b, $c;", ""))))),
 907          [(set regclass:$dst, (IntOp ptrclass:$addr, imm:$b, regclass:$c))]>,
 908   Requires<[Pred]>;
 909   def imm2 : NVPTXInst<(outs regclass:$dst),
 910     (ins ptrclass:$addr, regclass:$b, IMMType:$c),
 911                !strconcat("atom",
 912          !strconcat(SpaceStr,
 913          !strconcat(OpcStr,
 914          !strconcat(TypeStr,
 915          !strconcat(" \t$dst, [$addr], $b, $c;", ""))))),
 916          [(set regclass:$dst, (IntOp ptrclass:$addr, regclass:$b, imm:$c))]>,
 917   Requires<[Pred]>;
 918   def imm3 : NVPTXInst<(outs regclass:$dst),
 919     (ins ptrclass:$addr, IMMType:$b, IMMType:$c),
 920                !strconcat("atom",
 921          !strconcat(SpaceStr,
 922          !strconcat(OpcStr,
 923          !strconcat(TypeStr,
 924          !strconcat(" \t$dst, [$addr], $b, $c;", ""))))),
 925          [(set regclass:$dst, (IntOp ptrclass:$addr, imm:$b, imm:$c))]>,
 926   Requires<[Pred]>;
 927 }
 928 multiclass F_ATOMIC_3<NVPTXRegClass regclass, string SpaceStr, string TypeStr,
 929   string OpcStr, PatFrag IntOp, Operand IMMType, Predicate Pred> {
 930   defm p32 : F_ATOMIC_3_imp<Int32Regs, regclass, SpaceStr, TypeStr, OpcStr,
 931     IntOp, IMMType, Pred>;
 932   defm p64 : F_ATOMIC_3_imp<Int64Regs, regclass, SpaceStr, TypeStr, OpcStr,
 933     IntOp, IMMType, Pred>;
 934 }
 935
 936 // atom_add
 937
 938 def atomic_load_add_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
 939   (atomic_load_add_32 node:$a, node:$b)>;
 940 def atomic_load_add_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
 941   (atomic_load_add_32 node:$a, node:$b)>;
 942 def atomic_load_add_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
 943   (atomic_load_add_32 node:$a, node:$b)>;
 944 def atomic_load_add_64_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
 945   (atomic_load_add_64 node:$a, node:$b)>;
 946 def atomic_load_add_64_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
 947   (atomic_load_add_64 node:$a, node:$b)>;
 948 def atomic_load_add_64_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
 949   (atomic_load_add_64 node:$a, node:$b)>;
 950 def atomic_load_add_f32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
 951   (int_nvvm_atomic_load_add_f32 node:$a, node:$b)>;
 952 def atomic_load_add_f32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
 953   (int_nvvm_atomic_load_add_f32 node:$a, node:$b)>;
 954 def atomic_load_add_f32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
 955   (int_nvvm_atomic_load_add_f32 node:$a, node:$b)>;
 956
 957 defm INT_PTX_ATOM_ADD_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".u32", ".add",
 958   atomic_load_add_32_g, i32imm, imm, hasAtomRedG32>;
 959 defm INT_PTX_ATOM_ADD_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".u32", ".add",
 960   atomic_load_add_32_s, i32imm, imm, hasAtomRedS32>;
 961 defm INT_PTX_ATOM_ADD_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".u32", ".add",
 962   atomic_load_add_32_gen, i32imm, imm, hasAtomRedGen32>;
 963 defm INT_PTX_ATOM_ADD_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".u32",
 964   ".add", atomic_load_add_32_gen, i32imm, imm, useAtomRedG32forGen32>;
 965
 966 defm INT_PTX_ATOM_ADD_G_64 : F_ATOMIC_2<Int64Regs, ".global", ".u64", ".add",
 967   atomic_load_add_64_g, i64imm, imm, hasAtomRedG64>;
 968 defm INT_PTX_ATOM_ADD_S_64 : F_ATOMIC_2<Int64Regs, ".shared", ".u64", ".add",
 969   atomic_load_add_64_s, i64imm, imm, hasAtomRedS64>;
 970 defm INT_PTX_ATOM_ADD_GEN_64 : F_ATOMIC_2<Int64Regs, "", ".u64", ".add",
 971   atomic_load_add_64_gen, i64imm, imm, hasAtomRedGen64>;
 972 defm INT_PTX_ATOM_ADD_GEN_64_USE_G : F_ATOMIC_2<Int64Regs, ".global", ".u64",
 973   ".add", atomic_load_add_64_gen, i64imm, imm, useAtomRedG64forGen64>;
 974
 975 defm INT_PTX_ATOM_ADD_G_F32 : F_ATOMIC_2<Float32Regs, ".global", ".f32", ".add",
 976   atomic_load_add_f32_g, f32imm, fpimm, hasAtomAddF32>;
 977 defm INT_PTX_ATOM_ADD_S_F32 : F_ATOMIC_2<Float32Regs, ".shared", ".f32", ".add",
 978   atomic_load_add_f32_s, f32imm, fpimm, hasAtomAddF32>;
 979 defm INT_PTX_ATOM_ADD_GEN_F32 : F_ATOMIC_2<Float32Regs, "", ".f32", ".add",
 980   atomic_load_add_f32_gen, f32imm, fpimm, hasAtomAddF32>;
 981
 982 // atom_sub
 983
 984 def atomic_load_sub_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
 985   (atomic_load_sub_32 node:$a, node:$b)>;
 986 def atomic_load_sub_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
 987   (atomic_load_sub_32 node:$a, node:$b)>;
 988 def atomic_load_sub_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
 989   (atomic_load_sub_32 node:$a, node:$b)>;
 990 def atomic_load_sub_64_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
 991   (atomic_load_sub_64 node:$a, node:$b)>;
 992 def atomic_load_sub_64_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
 993   (atomic_load_sub_64 node:$a, node:$b)>;
 994 def atomic_load_sub_64_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
 995   (atomic_load_sub_64 node:$a, node:$b)>;
 996
 997 defm INT_PTX_ATOM_SUB_G_32 : F_ATOMIC_2_NEG<Int32Regs, ".global", "32", ".add",
 998   atomic_load_sub_32_g, i32imm, hasAtomRedG32>;
 999 defm INT_PTX_ATOM_SUB_G_64 : F_ATOMIC_2_NEG<Int64Regs, ".global", "64", ".add",
1000   atomic_load_sub_64_g, i64imm, hasAtomRedG64>;
1001 defm INT_PTX_ATOM_SUB_GEN_32 : F_ATOMIC_2_NEG<Int32Regs, "", "32", ".add",
1002   atomic_load_sub_32_gen, i32imm, hasAtomRedGen32>;
1003 defm INT_PTX_ATOM_SUB_GEN_32_USE_G : F_ATOMIC_2_NEG<Int32Regs, ".global", "32",
1004   ".add", atomic_load_sub_32_gen, i32imm, useAtomRedG32forGen32>;
1005 defm INT_PTX_ATOM_SUB_S_32 : F_ATOMIC_2_NEG<Int32Regs, ".shared", "32", ".add",
1006   atomic_load_sub_32_s, i32imm, hasAtomRedS32>;
1007 defm INT_PTX_ATOM_SUB_S_64 : F_ATOMIC_2_NEG<Int64Regs, ".shared", "64", ".add",
1008   atomic_load_sub_64_s, i64imm, hasAtomRedS64>;
1009 defm INT_PTX_ATOM_SUB_GEN_64 : F_ATOMIC_2_NEG<Int64Regs, "", "64", ".add",
1010   atomic_load_sub_64_gen, i64imm, hasAtomRedGen64>;
1011 defm INT_PTX_ATOM_SUB_GEN_64_USE_G : F_ATOMIC_2_NEG<Int64Regs, ".global", "64",
1012   ".add", atomic_load_sub_64_gen, i64imm, useAtomRedG64forGen64>;
1013
1014 // atom_swap
1015
1016 def atomic_swap_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1017   (atomic_swap_32 node:$a, node:$b)>;
1018 def atomic_swap_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1019   (atomic_swap_32 node:$a, node:$b)>;
1020 def atomic_swap_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1021   (atomic_swap_32 node:$a, node:$b)>;
1022 def atomic_swap_64_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1023   (atomic_swap_64 node:$a, node:$b)>;
1024 def atomic_swap_64_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1025   (atomic_swap_64 node:$a, node:$b)>;
1026 def atomic_swap_64_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1027   (atomic_swap_64 node:$a, node:$b)>;
1028
1029 defm INT_PTX_ATOM_SWAP_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".b32", ".exch",
1030   atomic_swap_32_g, i32imm, imm, hasAtomRedG32>;
1031 defm INT_PTX_ATOM_SWAP_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".b32", ".exch",
1032   atomic_swap_32_s, i32imm, imm, hasAtomRedS32>;
1033 defm INT_PTX_ATOM_SWAP_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".b32", ".exch",
1034   atomic_swap_32_gen, i32imm, imm, hasAtomRedGen32>;
1035 defm INT_PTX_ATOM_SWAP_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".b32",
1036   ".exch", atomic_swap_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1037 defm INT_PTX_ATOM_SWAP_G_64 : F_ATOMIC_2<Int64Regs, ".global", ".b64", ".exch",
1038   atomic_swap_64_g, i64imm, imm, hasAtomRedG64>;
1039 defm INT_PTX_ATOM_SWAP_S_64 : F_ATOMIC_2<Int64Regs, ".shared", ".b64", ".exch",
1040   atomic_swap_64_s, i64imm, imm, hasAtomRedS64>;
1041 defm INT_PTX_ATOM_SWAP_GEN_64 : F_ATOMIC_2<Int64Regs, "", ".b64", ".exch",
1042   atomic_swap_64_gen, i64imm, imm, hasAtomRedGen64>;
1043 defm INT_PTX_ATOM_SWAP_GEN_64_USE_G : F_ATOMIC_2<Int64Regs, ".global", ".b64",
1044   ".exch", atomic_swap_64_gen, i64imm, imm, useAtomRedG64forGen64>;
1045
1046 // atom_max
1047
1048 def atomic_load_max_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b)
1049   , (atomic_load_max_32 node:$a, node:$b)>;
1050 def atomic_load_max_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1051   (atomic_load_max_32 node:$a, node:$b)>;
1052 def atomic_load_max_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1053   (atomic_load_max_32 node:$a, node:$b)>;
1054 def atomic_load_umax_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1055   (atomic_load_umax_32 node:$a, node:$b)>;
1056 def atomic_load_umax_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1057   (atomic_load_umax_32 node:$a, node:$b)>;
1058 def atomic_load_umax_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1059   (atomic_load_umax_32 node:$a, node:$b)>;
1060
1061 defm INT_PTX_ATOM_LOAD_MAX_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".s32",
1062   ".max", atomic_load_max_32_g, i32imm, imm, hasAtomRedG32>;
1063 defm INT_PTX_ATOM_LOAD_MAX_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".s32",
1064   ".max", atomic_load_max_32_s, i32imm, imm, hasAtomRedS32>;
1065 defm INT_PTX_ATOM_LOAD_MAX_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".s32", ".max",
1066   atomic_load_max_32_gen, i32imm, imm, hasAtomRedGen32>;
1067 defm INT_PTX_ATOM_LOAD_MAX_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global",
1068   ".s32", ".max", atomic_load_max_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1069 defm INT_PTX_ATOM_LOAD_UMAX_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".u32",
1070   ".max", atomic_load_umax_32_g, i32imm, imm, hasAtomRedG32>;
1071 defm INT_PTX_ATOM_LOAD_UMAX_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".u32",
1072   ".max", atomic_load_umax_32_s, i32imm, imm, hasAtomRedS32>;
1073 defm INT_PTX_ATOM_LOAD_UMAX_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".u32", ".max",
1074   atomic_load_umax_32_gen, i32imm, imm, hasAtomRedGen32>;
1075 defm INT_PTX_ATOM_LOAD_UMAX_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global",
1076   ".u32", ".max", atomic_load_umax_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1077
1078 // atom_min
1079
1080 def atomic_load_min_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1081   (atomic_load_min_32 node:$a, node:$b)>;
1082 def atomic_load_min_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1083   (atomic_load_min_32 node:$a, node:$b)>;
1084 def atomic_load_min_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1085   (atomic_load_min_32 node:$a, node:$b)>;
1086 def atomic_load_umin_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1087   (atomic_load_umin_32 node:$a, node:$b)>;
1088 def atomic_load_umin_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1089   (atomic_load_umin_32 node:$a, node:$b)>;
1090 def atomic_load_umin_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1091   (atomic_load_umin_32 node:$a, node:$b)>;
1092
1093 defm INT_PTX_ATOM_LOAD_MIN_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".s32",
1094   ".min", atomic_load_min_32_g, i32imm, imm, hasAtomRedG32>;
1095 defm INT_PTX_ATOM_LOAD_MIN_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".s32",
1096   ".min", atomic_load_min_32_s, i32imm, imm, hasAtomRedS32>;
1097 defm INT_PTX_ATOM_LOAD_MIN_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".s32", ".min",
1098   atomic_load_min_32_gen, i32imm, imm, hasAtomRedGen32>;
1099 defm INT_PTX_ATOM_LOAD_MIN_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global",
1100   ".s32", ".min", atomic_load_min_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1101 defm INT_PTX_ATOM_LOAD_UMIN_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".u32",
1102   ".min", atomic_load_umin_32_g, i32imm, imm, hasAtomRedG32>;
1103 defm INT_PTX_ATOM_LOAD_UMIN_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".u32",
1104   ".min", atomic_load_umin_32_s, i32imm, imm, hasAtomRedS32>;
1105 defm INT_PTX_ATOM_LOAD_UMIN_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".u32", ".min",
1106   atomic_load_umin_32_gen, i32imm, imm, hasAtomRedGen32>;
1107 defm INT_PTX_ATOM_LOAD_UMIN_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global",
1108   ".u32", ".min", atomic_load_umin_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1109
1110 // atom_inc  atom_dec
1111
1112 def atomic_load_inc_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1113   (int_nvvm_atomic_load_inc_32 node:$a, node:$b)>;
1114 def atomic_load_inc_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1115   (int_nvvm_atomic_load_inc_32 node:$a, node:$b)>;
1116 def atomic_load_inc_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1117   (int_nvvm_atomic_load_inc_32 node:$a, node:$b)>;
1118 def atomic_load_dec_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1119   (int_nvvm_atomic_load_dec_32 node:$a, node:$b)>;
1120 def atomic_load_dec_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1121   (int_nvvm_atomic_load_dec_32 node:$a, node:$b)>;
1122 def atomic_load_dec_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1123   (int_nvvm_atomic_load_dec_32 node:$a, node:$b)>;
1124
1125 defm INT_PTX_ATOM_INC_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".u32", ".inc",
1126   atomic_load_inc_32_g, i32imm, imm, hasAtomRedG32>;
1127 defm INT_PTX_ATOM_INC_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".u32", ".inc",
1128   atomic_load_inc_32_s, i32imm, imm, hasAtomRedS32>;
1129 defm INT_PTX_ATOM_INC_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".u32", ".inc",
1130   atomic_load_inc_32_gen, i32imm, imm, hasAtomRedGen32>;
1131 defm INT_PTX_ATOM_INC_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".u32",
1132   ".inc", atomic_load_inc_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1133 defm INT_PTX_ATOM_DEC_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".u32", ".dec",
1134   atomic_load_dec_32_g, i32imm, imm, hasAtomRedG32>;
1135 defm INT_PTX_ATOM_DEC_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".u32", ".dec",
1136   atomic_load_dec_32_s, i32imm, imm, hasAtomRedS32>;
1137 defm INT_PTX_ATOM_DEC_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".u32", ".dec",
1138   atomic_load_dec_32_gen, i32imm, imm, hasAtomRedGen32>;
1139 defm INT_PTX_ATOM_DEC_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".u32",
1140   ".dec", atomic_load_dec_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1141
1142 // atom_and
1143
1144 def atomic_load_and_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1145   (atomic_load_and_32 node:$a, node:$b)>;
1146 def atomic_load_and_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1147   (atomic_load_and_32 node:$a, node:$b)>;
1148 def atomic_load_and_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1149   (atomic_load_and_32 node:$a, node:$b)>;
1150
1151 defm INT_PTX_ATOM_AND_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".b32", ".and",
1152   atomic_load_and_32_g, i32imm, imm, hasAtomRedG32>;
1153 defm INT_PTX_ATOM_AND_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".b32", ".and",
1154   atomic_load_and_32_s, i32imm, imm, hasAtomRedS32>;
1155 defm INT_PTX_ATOM_AND_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".b32", ".and",
1156   atomic_load_and_32_gen, i32imm, imm, hasAtomRedGen32>;
1157 defm INT_PTX_ATOM_AND_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".b32",
1158   ".and", atomic_load_and_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1159
1160 // atom_or
1161
1162 def atomic_load_or_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1163   (atomic_load_or_32 node:$a, node:$b)>;
1164 def atomic_load_or_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1165   (atomic_load_or_32 node:$a, node:$b)>;
1166 def atomic_load_or_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1167   (atomic_load_or_32 node:$a, node:$b)>;
1168
1169 defm INT_PTX_ATOM_OR_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".b32", ".or",
1170   atomic_load_or_32_g, i32imm, imm, hasAtomRedG32>;
1171 defm INT_PTX_ATOM_OR_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".b32", ".or",
1172   atomic_load_or_32_gen, i32imm, imm, hasAtomRedGen32>;
1173 defm INT_PTX_ATOM_OR_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".b32",
1174   ".or", atomic_load_or_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1175 defm INT_PTX_ATOM_OR_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".b32", ".or",
1176   atomic_load_or_32_s, i32imm, imm, hasAtomRedS32>;
1177
1178 // atom_xor
1179
1180 def atomic_load_xor_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1181   (atomic_load_xor_32 node:$a, node:$b)>;
1182 def atomic_load_xor_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1183   (atomic_load_xor_32 node:$a, node:$b)>;
1184 def atomic_load_xor_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1185   (atomic_load_xor_32 node:$a, node:$b)>;
1186
1187 defm INT_PTX_ATOM_XOR_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".b32", ".xor",
1188   atomic_load_xor_32_g, i32imm, imm, hasAtomRedG32>;
1189 defm INT_PTX_ATOM_XOR_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".b32", ".xor",
1190   atomic_load_xor_32_s, i32imm, imm, hasAtomRedS32>;
1191 defm INT_PTX_ATOM_XOR_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".b32", ".xor",
1192   atomic_load_xor_32_gen, i32imm, imm, hasAtomRedGen32>;
1193 defm INT_PTX_ATOM_XOR_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".b32",
1194   ".xor", atomic_load_xor_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1195
1196 // atom_cas
1197
1198 def atomic_cmp_swap_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b, node:$c),
1199   (atomic_cmp_swap_32 node:$a, node:$b, node:$c)>;
1200 def atomic_cmp_swap_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b, node:$c),
1201   (atomic_cmp_swap_32 node:$a, node:$b, node:$c)>;
1202 def atomic_cmp_swap_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b, node:$c),
1203   (atomic_cmp_swap_32 node:$a, node:$b, node:$c)>;
1204 def atomic_cmp_swap_64_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b, node:$c),
1205   (atomic_cmp_swap_64 node:$a, node:$b, node:$c)>;
1206 def atomic_cmp_swap_64_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b, node:$c),
1207   (atomic_cmp_swap_64 node:$a, node:$b, node:$c)>;
1208 def atomic_cmp_swap_64_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b, node:$c),
1209   (atomic_cmp_swap_64 node:$a, node:$b, node:$c)>;
1210
1211 defm INT_PTX_ATOM_CAS_G_32 : F_ATOMIC_3<Int32Regs, ".global", ".b32", ".cas",
1212   atomic_cmp_swap_32_g, i32imm, hasAtomRedG32>;
1213 defm INT_PTX_ATOM_CAS_S_32 : F_ATOMIC_3<Int32Regs, ".shared", ".b32", ".cas",
1214   atomic_cmp_swap_32_s, i32imm, hasAtomRedS32>;
1215 defm INT_PTX_ATOM_CAS_GEN_32 : F_ATOMIC_3<Int32Regs, "", ".b32", ".cas",
1216   atomic_cmp_swap_32_gen, i32imm, hasAtomRedGen32>;
1217 defm INT_PTX_ATOM_CAS_GEN_32_USE_G : F_ATOMIC_3<Int32Regs, ".global", ".b32",
1218   ".cas", atomic_cmp_swap_32_gen, i32imm, useAtomRedG32forGen32>;
1219 defm INT_PTX_ATOM_CAS_G_64 : F_ATOMIC_3<Int64Regs, ".global", ".b64", ".cas",
1220   atomic_cmp_swap_64_g, i64imm, hasAtomRedG64>;
1221 defm INT_PTX_ATOM_CAS_S_64 : F_ATOMIC_3<Int64Regs, ".shared", ".b64", ".cas",
1222   atomic_cmp_swap_64_s, i64imm, hasAtomRedS64>;
1223 defm INT_PTX_ATOM_CAS_GEN_64 : F_ATOMIC_3<Int64Regs, "", ".b64", ".cas",
1224   atomic_cmp_swap_64_gen, i64imm, hasAtomRedGen64>;
1225 defm INT_PTX_ATOM_CAS_GEN_64_USE_G : F_ATOMIC_3<Int64Regs, ".global", ".b64",
1226   ".cas", atomic_cmp_swap_64_gen, i64imm, useAtomRedG64forGen64>;
1227
1228
1229 //-----------------------------------
1230 // Read Special Registers
1231 //-----------------------------------
1232 class F_SREG<string OpStr, NVPTXRegClass regclassOut, Intrinsic IntOp> :
1233       NVPTXInst<(outs regclassOut:$dst), (ins),
1234                OpStr,
1235          [(set regclassOut:$dst, (IntOp))]>;
1236
1237 def INT_PTX_SREG_TID_X : F_SREG<"mov.u32 \t$dst, %tid.x;", Int32Regs,
1238   int_nvvm_read_ptx_sreg_tid_x>;
1239 def INT_PTX_SREG_TID_Y : F_SREG<"mov.u32 \t$dst, %tid.y;", Int32Regs,
1240   int_nvvm_read_ptx_sreg_tid_y>;
1241 def INT_PTX_SREG_TID_Z : F_SREG<"mov.u32 \t$dst, %tid.z;", Int32Regs,
1242   int_nvvm_read_ptx_sreg_tid_z>;
1243
1244 def INT_PTX_SREG_NTID_X : F_SREG<"mov.u32 \t$dst, %ntid.x;", Int32Regs,
1245   int_nvvm_read_ptx_sreg_ntid_x>;
1246 def INT_PTX_SREG_NTID_Y : F_SREG<"mov.u32 \t$dst, %ntid.y;", Int32Regs,
1247   int_nvvm_read_ptx_sreg_ntid_y>;
1248 def INT_PTX_SREG_NTID_Z : F_SREG<"mov.u32 \t$dst, %ntid.z;", Int32Regs,
1249   int_nvvm_read_ptx_sreg_ntid_z>;
1250
1251 def INT_PTX_SREG_CTAID_X : F_SREG<"mov.u32 \t$dst, %ctaid.x;", Int32Regs,
1252   int_nvvm_read_ptx_sreg_ctaid_x>;
1253 def INT_PTX_SREG_CTAID_Y : F_SREG<"mov.u32 \t$dst, %ctaid.y;", Int32Regs,
1254   int_nvvm_read_ptx_sreg_ctaid_y>;
1255 def INT_PTX_SREG_CTAID_Z : F_SREG<"mov.u32 \t$dst, %ctaid.z;", Int32Regs,
1256   int_nvvm_read_ptx_sreg_ctaid_z>;
1257
1258 def INT_PTX_SREG_NCTAID_X : F_SREG<"mov.u32 \t$dst, %nctaid.x;", Int32Regs,
1259   int_nvvm_read_ptx_sreg_nctaid_x>;
1260 def INT_PTX_SREG_NCTAID_Y : F_SREG<"mov.u32 \t$dst, %nctaid.y;", Int32Regs,
1261   int_nvvm_read_ptx_sreg_nctaid_y>;
1262 def INT_PTX_SREG_NCTAID_Z : F_SREG<"mov.u32 \t$dst, %nctaid.z;", Int32Regs,
1263   int_nvvm_read_ptx_sreg_nctaid_z>;
1264
1265 def INT_PTX_SREG_WARPSIZE : F_SREG<"mov.u32 \t$dst, WARP_SZ;", Int32Regs,
1266   int_nvvm_read_ptx_sreg_warpsize>;
1267
1268
1269 //-----------------------------------
1270 // Support for ldu on sm_20 or later
1271 //-----------------------------------
1272
1273 // Scalar
1274 // @TODO: Revisit this, Changed imemAny to imem
1275 multiclass LDU_G<string TyStr, NVPTXRegClass regclass, Intrinsic IntOp> {
1276   def areg: NVPTXInst<(outs regclass:$result), (ins Int32Regs:$src),
1277                !strconcat("ldu.global.", TyStr),
1278          [(set regclass:$result, (IntOp Int32Regs:$src))]>, Requires<[hasLDU]>;
1279   def areg64: NVPTXInst<(outs regclass:$result), (ins Int64Regs:$src),
1280                !strconcat("ldu.global.", TyStr),
1281          [(set regclass:$result, (IntOp Int64Regs:$src))]>, Requires<[hasLDU]>;
1282  def avar:  NVPTXInst<(outs regclass:$result), (ins imem:$src),
1283                !strconcat("ldu.global.", TyStr),
1284                 [(set regclass:$result, (IntOp (Wrapper tglobaladdr:$src)))]>,
1285                 Requires<[hasLDU]>;
1286  def ari :  NVPTXInst<(outs regclass:$result), (ins MEMri:$src),
1287                !strconcat("ldu.global.", TyStr),
1288          [(set regclass:$result, (IntOp ADDRri:$src))]>, Requires<[hasLDU]>;
1289  def ari64 :  NVPTXInst<(outs regclass:$result), (ins MEMri64:$src),
1290                !strconcat("ldu.global.", TyStr),
1291          [(set regclass:$result, (IntOp ADDRri64:$src))]>, Requires<[hasLDU]>;
1292 }
1293
1294 defm INT_PTX_LDU_GLOBAL_i8  : LDU_G<"u8 \t$result, [$src];",  Int8Regs,
1295 int_nvvm_ldu_global_i>;
1296 defm INT_PTX_LDU_GLOBAL_i16 : LDU_G<"u16 \t$result, [$src];", Int16Regs,
1297 int_nvvm_ldu_global_i>;
1298 defm INT_PTX_LDU_GLOBAL_i32 : LDU_G<"u32 \t$result, [$src];", Int32Regs,
1299 int_nvvm_ldu_global_i>;
1300 defm INT_PTX_LDU_GLOBAL_i64 : LDU_G<"u64 \t$result, [$src];", Int64Regs,
1301 int_nvvm_ldu_global_i>;
1302 defm INT_PTX_LDU_GLOBAL_f32 : LDU_G<"f32 \t$result, [$src];", Float32Regs,
1303 int_nvvm_ldu_global_f>;
1304 defm INT_PTX_LDU_GLOBAL_f64 : LDU_G<"f64 \t$result, [$src];", Float64Regs,
1305 int_nvvm_ldu_global_f>;
1306 defm INT_PTX_LDU_GLOBAL_p32 : LDU_G<"u32 \t$result, [$src];", Int32Regs,
1307 int_nvvm_ldu_global_p>;
1308 defm INT_PTX_LDU_GLOBAL_p64 : LDU_G<"u64 \t$result, [$src];", Int64Regs,
1309 int_nvvm_ldu_global_p>;
1310
1311 // vector
1312
1313 // Elementized vector ldu
1314 multiclass VLDU_G_ELE_V2<string TyStr, NVPTXRegClass regclass> {
1315  def _32:     NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
1316    (ins Int32Regs:$src),
1317                      !strconcat("ldu.global.", TyStr), []>;
1318  def _64:     NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
1319    (ins Int64Regs:$src),
1320                      !strconcat("ldu.global.", TyStr), []>;
1321 }
1322
1323 multiclass VLDU_G_ELE_V4<string TyStr, NVPTXRegClass regclass> {
1324  def _32:    NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
1325      regclass:$dst4), (ins Int32Regs:$src),
1326                !strconcat("ldu.global.", TyStr), []>;
1327  def _64:    NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
1328      regclass:$dst4), (ins Int64Regs:$src),
1329                !strconcat("ldu.global.", TyStr), []>;
1330 }
1331
1332 defm INT_PTX_LDU_G_v2i8_ELE
1333   : VLDU_G_ELE_V2<"v2.u8 \t{{$dst1, $dst2}}, [$src];",  Int8Regs>;
1334 defm INT_PTX_LDU_G_v2i16_ELE
1335   : VLDU_G_ELE_V2<"v2.u16 \t{{$dst1, $dst2}}, [$src];", Int16Regs>;
1336 defm INT_PTX_LDU_G_v2i32_ELE
1337   : VLDU_G_ELE_V2<"v2.u32 \t{{$dst1, $dst2}}, [$src];", Int32Regs>;
1338 defm INT_PTX_LDU_G_v2f32_ELE
1339   : VLDU_G_ELE_V2<"v2.f32 \t{{$dst1, $dst2}}, [$src];", Float32Regs>;
1340 defm INT_PTX_LDU_G_v2i64_ELE
1341   : VLDU_G_ELE_V2<"v2.u64 \t{{$dst1, $dst2}}, [$src];", Int64Regs>;
1342 defm INT_PTX_LDU_G_v2f64_ELE
1343   : VLDU_G_ELE_V2<"v2.f64 \t{{$dst1, $dst2}}, [$src];", Float64Regs>;
1344 defm INT_PTX_LDU_G_v4i8_ELE
1345   : VLDU_G_ELE_V4<"v4.u8 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Int8Regs>;
1346 defm INT_PTX_LDU_G_v4i16_ELE
1347   : VLDU_G_ELE_V4<"v4.u16 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];",
1348     Int16Regs>;
1349 defm INT_PTX_LDU_G_v4i32_ELE
1350   : VLDU_G_ELE_V4<"v4.u32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];",
1351     Int32Regs>;
1352 defm INT_PTX_LDU_G_v4f32_ELE
1353   : VLDU_G_ELE_V4<"v4.f32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];",
1354     Float32Regs>;
1355
1356
1357 //-----------------------------------
1358 // Support for ldg on sm_35 or later
1359 //-----------------------------------
1360
1361 def ldg_i8 : PatFrag<(ops node:$ptr), (int_nvvm_ldg_global_i node:$ptr), [{
1362   MemIntrinsicSDNode *M = cast<MemIntrinsicSDNode>(N);
1363   return M->getMemoryVT() == MVT::i8;
1364 }]>;
1365
1366 multiclass LDG_G<string TyStr, NVPTXRegClass regclass, Intrinsic IntOp> {
1367   def areg: NVPTXInst<(outs regclass:$result), (ins Int32Regs:$src),
1368                !strconcat("ld.global.nc.", TyStr),
1369          [(set regclass:$result, (IntOp Int32Regs:$src))]>, Requires<[hasLDG]>;
1370   def areg64: NVPTXInst<(outs regclass:$result), (ins Int64Regs:$src),
1371                !strconcat("ld.global.nc.", TyStr),
1372          [(set regclass:$result, (IntOp Int64Regs:$src))]>, Requires<[hasLDG]>;
1373  def avar:  NVPTXInst<(outs regclass:$result), (ins imem:$src),
1374                !strconcat("ld.global.nc.", TyStr),
1375          [(set regclass:$result, (IntOp (Wrapper tglobaladdr:$src)))]>,
1376          Requires<[hasLDG]>;
1377  def ari :  NVPTXInst<(outs regclass:$result), (ins MEMri:$src),
1378                !strconcat("ld.global.nc.", TyStr),
1379          [(set regclass:$result, (IntOp ADDRri:$src))]>, Requires<[hasLDG]>;
1380  def ari64 :  NVPTXInst<(outs regclass:$result), (ins MEMri64:$src),
1381                !strconcat("ld.global.nc.", TyStr),
1382          [(set regclass:$result, (IntOp ADDRri64:$src))]>, Requires<[hasLDG]>;
1383 }
1384
1385 multiclass LDG_G_NOINTRIN<string TyStr, NVPTXRegClass regclass, PatFrag IntOp> {
1386   def areg: NVPTXInst<(outs regclass:$result), (ins Int32Regs:$src),
1387                !strconcat("ld.global.nc.", TyStr),
1388          [(set regclass:$result, (IntOp Int32Regs:$src))]>, Requires<[hasLDG]>;
1389   def areg64: NVPTXInst<(outs regclass:$result), (ins Int64Regs:$src),
1390                !strconcat("ld.global.nc.", TyStr),
1391          [(set regclass:$result, (IntOp Int64Regs:$src))]>, Requires<[hasLDG]>;
1392  def avar:  NVPTXInst<(outs regclass:$result), (ins imem:$src),
1393                !strconcat("ld.global.nc.", TyStr),
1394          [(set regclass:$result, (IntOp (Wrapper tglobaladdr:$src)))]>,
1395         Requires<[hasLDG]>;
1396  def ari :  NVPTXInst<(outs regclass:$result), (ins MEMri:$src),
1397                !strconcat("ld.global.nc.", TyStr),
1398          [(set regclass:$result, (IntOp ADDRri:$src))]>, Requires<[hasLDG]>;
1399  def ari64 :  NVPTXInst<(outs regclass:$result), (ins MEMri64:$src),
1400                !strconcat("ld.global.nc.", TyStr),
1401          [(set regclass:$result, (IntOp ADDRri64:$src))]>, Requires<[hasLDG]>;
1402 }
1403
1404 defm INT_PTX_LDG_GLOBAL_i8
1405   : LDG_G_NOINTRIN<"u8 \t$result, [$src];",  Int16Regs, ldg_i8>;
1406 defm INT_PTX_LDG_GLOBAL_i16
1407   : LDG_G<"u16 \t$result, [$src];", Int16Regs,   int_nvvm_ldg_global_i>;
1408 defm INT_PTX_LDG_GLOBAL_i32
1409   : LDG_G<"u32 \t$result, [$src];", Int32Regs,   int_nvvm_ldg_global_i>;
1410 defm INT_PTX_LDG_GLOBAL_i64
1411   : LDG_G<"u64 \t$result, [$src];", Int64Regs,   int_nvvm_ldg_global_i>;
1412 defm INT_PTX_LDG_GLOBAL_f32
1413   : LDG_G<"f32 \t$result, [$src];", Float32Regs, int_nvvm_ldg_global_f>;
1414 defm INT_PTX_LDG_GLOBAL_f64
1415   : LDG_G<"f64 \t$result, [$src];", Float64Regs, int_nvvm_ldg_global_f>;
1416 defm INT_PTX_LDG_GLOBAL_p32
1417   : LDG_G<"u32 \t$result, [$src];", Int32Regs,   int_nvvm_ldg_global_p>;
1418 defm INT_PTX_LDG_GLOBAL_p64
1419   : LDG_G<"u64 \t$result, [$src];", Int64Regs,   int_nvvm_ldg_global_p>;
1420
1421 // vector
1422
1423 // Elementized vector ldg
1424 multiclass VLDG_G_ELE_V2<string TyStr, NVPTXRegClass regclass> {
1425  def _32:     NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
1426                      (ins Int32Regs:$src),
1427                      !strconcat("ld.global.nc.", TyStr), []>;
1428  def _64:     NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
1429                      (ins Int64Regs:$src),
1430                      !strconcat("ld.global.nc.", TyStr), []>;
1431 }
1432
1433 multiclass VLDG_G_ELE_V4<string TyStr, NVPTXRegClass regclass> {
1434  def _32:    NVPTXInst<(outs regclass:$dst1, regclass:$dst2,
1435                         regclass:$dst3, regclass:$dst4), (ins Int32Regs:$src),
1436                !strconcat("ld.global.nc.", TyStr), []>;
1437  def _64:    NVPTXInst<(outs regclass:$dst1, regclass:$dst2,
1438                         regclass:$dst3, regclass:$dst4), (ins Int64Regs:$src),
1439                !strconcat("ld.global.nc.", TyStr), []>;
1440 }
1441
1442 // FIXME: 8-bit LDG should be fixed once LDG/LDU nodes are made into proper loads.
1443 defm INT_PTX_LDG_G_v2i8_ELE
1444   : VLDG_G_ELE_V2<"v2.u8 \t{{$dst1, $dst2}}, [$src];",  Int16Regs>;
1445 defm INT_PTX_LDG_G_v2i16_ELE
1446   : VLDG_G_ELE_V2<"v2.u16 \t{{$dst1, $dst2}}, [$src];", Int16Regs>;
1447 defm INT_PTX_LDG_G_v2i32_ELE
1448   : VLDG_G_ELE_V2<"v2.u32 \t{{$dst1, $dst2}}, [$src];", Int32Regs>;
1449 defm INT_PTX_LDG_G_v2f32_ELE
1450   : VLDG_G_ELE_V2<"v2.f32 \t{{$dst1, $dst2}}, [$src];", Float32Regs>;
1451 defm INT_PTX_LDG_G_v2i64_ELE
1452   : VLDG_G_ELE_V2<"v2.u64 \t{{$dst1, $dst2}}, [$src];", Int64Regs>;
1453 defm INT_PTX_LDG_G_v2f64_ELE
1454   : VLDG_G_ELE_V2<"v2.f64 \t{{$dst1, $dst2}}, [$src];", Float64Regs>;
1455 defm INT_PTX_LDG_G_v4i8_ELE
1456   : VLDG_G_ELE_V4<"v4.u8 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Int16Regs>;
1457 defm INT_PTX_LDG_G_v4i16_ELE
1458   : VLDG_G_ELE_V4<"v4.u16 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Int16Regs>;
1459 defm INT_PTX_LDG_G_v4i32_ELE
1460   : VLDG_G_ELE_V4<"v4.u32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Int32Regs>;
1461 defm INT_PTX_LDG_G_v4f32_ELE
1462   : VLDG_G_ELE_V4<"v4.f32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Float32Regs>;
1463
1464
1465 multiclass NG_TO_G<string Str, Intrinsic Intrin> {
1466    def _yes : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
1467           !strconcat("cvta.", !strconcat(Str, ".u32 \t$result, $src;")),
1468       [(set Int32Regs:$result, (Intrin Int32Regs:$src))]>,
1469    Requires<[hasGenericLdSt]>;
1470    def _yes_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
1471           !strconcat("cvta.", !strconcat(Str, ".u64 \t$result, $src;")),
1472       [(set Int64Regs:$result, (Intrin Int64Regs:$src))]>,
1473    Requires<[hasGenericLdSt]>;
1474
1475 // @TODO: Are these actually needed?  I believe global addresses will be copied
1476 // to register values anyway.
1477    /*def __addr_yes : NVPTXInst<(outs Int32Regs:$result), (ins imemAny:$src),
1478           !strconcat("cvta.", !strconcat(Str, ".u32 \t$result, $src;")),
1479       [(set Int32Regs:$result, (Intrin (Wrapper tglobaladdr:$src)))]>,
1480       Requires<[hasGenericLdSt]>;
1481    def __addr_yes_64 : NVPTXInst<(outs Int64Regs:$result), (ins imemAny:$src),
1482           !strconcat("cvta.", !strconcat(Str, ".u64 \t$result, $src;")),
1483       [(set Int64Regs:$result, (Intrin (Wrapper tglobaladdr:$src)))]>,
1484       Requires<[hasGenericLdSt]>;*/
1485
1486    def _no : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
1487           "mov.u32 \t$result, $src;",
1488       [(set Int32Regs:$result, (Intrin Int32Regs:$src))]>;
1489    def _no_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
1490           "mov.u64 \t$result, $src;",
1491       [(set Int64Regs:$result, (Intrin Int64Regs:$src))]>;
1492
1493 // @TODO: Are these actually needed?  I believe global addresses will be copied
1494 // to register values anyway.
1495    /*def _addr_no : NVPTXInst<(outs Int32Regs:$result), (ins imem:$src),
1496           "mov.u32 \t$result, $src;",
1497       [(set Int32Regs:$result, (Intrin (Wrapper tglobaladdr:$src)))]>;
1498    def _addr_no_64 : NVPTXInst<(outs Int64Regs:$result), (ins imem:$src),
1499           "mov.u64 \t$result, $src;",
1500       [(set Int64Regs:$result, (Intrin (Wrapper tglobaladdr:$src)))]>;*/
1501 }
1502
1503 multiclass G_TO_NG<string Str, Intrinsic Intrin> {
1504    def _yes : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
1505           !strconcat("cvta.to.", !strconcat(Str, ".u32 \t$result, $src;")),
1506       [(set Int32Regs:$result, (Intrin Int32Regs:$src))]>,
1507    Requires<[hasGenericLdSt]>;
1508    def _yes_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
1509           !strconcat("cvta.to.", !strconcat(Str, ".u64 \t$result, $src;")),
1510       [(set Int64Regs:$result, (Intrin Int64Regs:$src))]>,
1511    Requires<[hasGenericLdSt]>;
1512    def _no : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
1513           "mov.u32 \t$result, $src;",
1514       [(set Int32Regs:$result, (Intrin Int32Regs:$src))]>;
1515    def _no_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
1516           "mov.u64 \t$result, $src;",
1517       [(set Int64Regs:$result, (Intrin Int64Regs:$src))]>;
1518 }
1519
1520 defm cvta_local  : NG_TO_G<"local", int_nvvm_ptr_local_to_gen>;
1521 defm cvta_shared : NG_TO_G<"shared", int_nvvm_ptr_shared_to_gen>;
1522 defm cvta_global : NG_TO_G<"global", int_nvvm_ptr_global_to_gen>;
1523 defm cvta_const  : NG_TO_G<"const", int_nvvm_ptr_constant_to_gen>;
1524
1525 defm cvta_to_local   : G_TO_NG<"local", int_nvvm_ptr_gen_to_local>;
1526 defm cvta_to_shared : G_TO_NG<"shared", int_nvvm_ptr_gen_to_shared>;
1527 defm cvta_to_global : G_TO_NG<"global", int_nvvm_ptr_gen_to_global>;
1528 defm cvta_to_const  : G_TO_NG<"const", int_nvvm_ptr_gen_to_constant>;
1529
1530
1531 // nvvm.ptr.gen.to.param
1532 def nvvm_ptr_gen_to_param : NVPTXInst<(outs Int32Regs:$result),
1533   (ins Int32Regs:$src),
1534                         "mov.u32 \t$result, $src;",
1535                               [(set Int32Regs:$result,
1536                                 (int_nvvm_ptr_gen_to_param Int32Regs:$src))]>;
1537 def nvvm_ptr_gen_to_param_64 : NVPTXInst<(outs Int64Regs:$result),
1538   (ins Int64Regs:$src),
1539                         "mov.u64 \t$result, $src;",
1540                               [(set Int64Regs:$result,
1541                                 (int_nvvm_ptr_gen_to_param Int64Regs:$src))]>;
1542
1543
1544 // nvvm.move intrinsicc
1545 def nvvm_move_i8 : NVPTXInst<(outs Int8Regs:$r), (ins Int8Regs:$s),
1546                              "mov.b16 \t$r, $s;",
1547                              [(set Int8Regs:$r,
1548                                (int_nvvm_move_i8 Int8Regs:$s))]>;
1549 def nvvm_move_i16 : NVPTXInst<(outs Int16Regs:$r), (ins Int16Regs:$s),
1550                              "mov.b16 \t$r, $s;",
1551                              [(set Int16Regs:$r,
1552                                (int_nvvm_move_i16 Int16Regs:$s))]>;
1553 def nvvm_move_i32 : NVPTXInst<(outs Int32Regs:$r), (ins Int32Regs:$s),
1554                              "mov.b32 \t$r, $s;",
1555                              [(set Int32Regs:$r,
1556                                (int_nvvm_move_i32 Int32Regs:$s))]>;
1557 def nvvm_move_i64 : NVPTXInst<(outs Int64Regs:$r), (ins Int64Regs:$s),
1558                              "mov.b64 \t$r, $s;",
1559                              [(set Int64Regs:$r,
1560                                (int_nvvm_move_i64 Int64Regs:$s))]>;
1561 def nvvm_move_float : NVPTXInst<(outs Float32Regs:$r), (ins Float32Regs:$s),
1562                              "mov.f32 \t$r, $s;",
1563                              [(set Float32Regs:$r,
1564                                (int_nvvm_move_float Float32Regs:$s))]>;
1565 def nvvm_move_double : NVPTXInst<(outs Float64Regs:$r), (ins Float64Regs:$s),
1566                              "mov.f64 \t$r, $s;",
1567                              [(set Float64Regs:$r,
1568                                (int_nvvm_move_double Float64Regs:$s))]>;
1569 def nvvm_move_ptr32 : NVPTXInst<(outs Int32Regs:$r), (ins Int32Regs:$s),
1570                              "mov.u32 \t$r, $s;",
1571                              [(set Int32Regs:$r,
1572                                (int_nvvm_move_ptr Int32Regs:$s))]>;
1573 def nvvm_move_ptr64 : NVPTXInst<(outs Int64Regs:$r), (ins Int64Regs:$s),
1574                              "mov.u64 \t$r, $s;",
1575                              [(set Int64Regs:$r,
1576                                (int_nvvm_move_ptr Int64Regs:$s))]>;
1577
1578 // @TODO: Are these actually needed, or will we always just see symbols
1579 // copied to registers first?
1580 /*def nvvm_move_sym32 : NVPTXInst<(outs Int32Regs:$r), (ins imem:$s),
1581                              "mov.u32 \t$r, $s;",
1582                              [(set Int32Regs:$r,
1583                              (int_nvvm_move_ptr texternalsym:$s))]>;
1584 def nvvm_move_sym64 : NVPTXInst<(outs Int64Regs:$r), (ins imem:$s),
1585                              "mov.u64 \t$r, $s;",
1586                              [(set Int64Regs:$r,
1587                              (int_nvvm_move_ptr texternalsym:$s))]>;*/
1588
1589
1590 // MoveParam        %r1, param
1591 // ptr_local_to_gen %r2, %r1
1592 // ptr_gen_to_local %r3, %r2
1593 // ->
1594 // mov %r1, param
1595
1596 // @TODO: Revisit this.  There is a type
1597 // contradiction between iPTRAny and iPTR for the addr defs, so the move_sym
1598 // instructions are not currently defined. However, we can use the ptr
1599 // variants and the asm printer will do the right thing.
1600 def : Pat<(i64 (int_nvvm_ptr_gen_to_local (int_nvvm_ptr_local_to_gen
1601                 (MoveParam texternalsym:$src)))),
1602                (nvvm_move_ptr64  texternalsym:$src)>;
1603 def : Pat<(i32 (int_nvvm_ptr_gen_to_local (int_nvvm_ptr_local_to_gen
1604                 (MoveParam texternalsym:$src)))),
1605                (nvvm_move_ptr32  texternalsym:$src)>;
1606
1607
1608 //-----------------------------------
1609 // Compiler Error Warn
1610 // - Just ignore them in codegen
1611 //-----------------------------------
1612
1613 def INT_NVVM_COMPILER_WARN_32 : NVPTXInst<(outs), (ins Int32Regs:$a),
1614                 "// llvm.nvvm.compiler.warn()",
1615                 [(int_nvvm_compiler_warn Int32Regs:$a)]>;
1616 def INT_NVVM_COMPILER_WARN_64 : NVPTXInst<(outs), (ins Int64Regs:$a),
1617                 "// llvm.nvvm.compiler.warn()",
1618                 [(int_nvvm_compiler_warn Int64Regs:$a)]>;
1619 def INT_NVVM_COMPILER_ERROR_32 : NVPTXInst<(outs), (ins Int32Regs:$a),
1620                 "// llvm.nvvm.compiler.error()",
1621                 [(int_nvvm_compiler_error Int32Regs:$a)]>;
1622 def INT_NVVM_COMPILER_ERROR_64 : NVPTXInst<(outs), (ins Int64Regs:$a),
1623                 "// llvm.nvvm.compiler.error()",
1624                 [(int_nvvm_compiler_error Int64Regs:$a)]>;
1625
1626
1627
1628 //===-- Old PTX Back-end Intrinsics ---------------------------------------===//
1629
1630 // These intrinsics are handled to retain compatibility with the old backend.
1631
1632 // PTX Special Purpose Register Accessor Intrinsics
1633
1634 class PTX_READ_SPECIAL_REGISTER_R64<string regname, Intrinsic intop>
1635   : NVPTXInst<(outs Int64Regs:$d), (ins),
1636               !strconcat(!strconcat("mov.u64\t$d, %", regname), ";"),
1637               [(set Int64Regs:$d, (intop))]>;
1638
1639 class PTX_READ_SPECIAL_REGISTER_R32<string regname, Intrinsic intop>
1640   : NVPTXInst<(outs Int32Regs:$d), (ins),
1641               !strconcat(!strconcat("mov.u32\t$d, %", regname), ";"),
1642               [(set Int32Regs:$d, (intop))]>;
1643
1644 // TODO Add read vector-version of special registers
1645
1646 def PTX_READ_TID_X   : PTX_READ_SPECIAL_REGISTER_R32<"tid.x",
1647                                                      int_ptx_read_tid_x>;
1648 def PTX_READ_TID_Y   : PTX_READ_SPECIAL_REGISTER_R32<"tid.y",
1649                                                      int_ptx_read_tid_y>;
1650 def PTX_READ_TID_Z   : PTX_READ_SPECIAL_REGISTER_R32<"tid.z",
1651                                                      int_ptx_read_tid_z>;
1652 def PTX_READ_TID_W   : PTX_READ_SPECIAL_REGISTER_R32<"tid.w",
1653                                                      int_ptx_read_tid_w>;
1654
1655 def PTX_READ_NTID_X   : PTX_READ_SPECIAL_REGISTER_R32<"ntid.x",
1656                                                       int_ptx_read_ntid_x>;
1657 def PTX_READ_NTID_Y   : PTX_READ_SPECIAL_REGISTER_R32<"ntid.y",
1658                                                       int_ptx_read_ntid_y>;
1659 def PTX_READ_NTID_Z   : PTX_READ_SPECIAL_REGISTER_R32<"ntid.z",
1660                                                       int_ptx_read_ntid_z>;
1661 def PTX_READ_NTID_W   : PTX_READ_SPECIAL_REGISTER_R32<"ntid.w",
1662                                                       int_ptx_read_ntid_w>;
1663
1664 def PTX_READ_LANEID  : PTX_READ_SPECIAL_REGISTER_R32<"laneid",
1665                                                      int_ptx_read_laneid>;
1666 def PTX_READ_WARPID  : PTX_READ_SPECIAL_REGISTER_R32<"warpid",
1667                                                      int_ptx_read_warpid>;
1668 def PTX_READ_NWARPID : PTX_READ_SPECIAL_REGISTER_R32<"nwarpid",
1669                                                      int_ptx_read_nwarpid>;
1670
1671 def PTX_READ_CTAID_X   : PTX_READ_SPECIAL_REGISTER_R32<"ctaid.x",
1672                                                        int_ptx_read_ctaid_x>;
1673 def PTX_READ_CTAID_Y   : PTX_READ_SPECIAL_REGISTER_R32<"ctaid.y",
1674                                                        int_ptx_read_ctaid_y>;
1675 def PTX_READ_CTAID_Z   : PTX_READ_SPECIAL_REGISTER_R32<"ctaid.z",
1676                                                        int_ptx_read_ctaid_z>;
1677 def PTX_READ_CTAID_W   : PTX_READ_SPECIAL_REGISTER_R32<"ctaid.w",
1678                                                        int_ptx_read_ctaid_w>;
1679
1680 def PTX_READ_NCTAID_X   : PTX_READ_SPECIAL_REGISTER_R32<"nctaid.x",
1681                                                         int_ptx_read_nctaid_x>;
1682 def PTX_READ_NCTAID_Y   : PTX_READ_SPECIAL_REGISTER_R32<"nctaid.y",
1683                                                         int_ptx_read_nctaid_y>;
1684 def PTX_READ_NCTAID_Z   : PTX_READ_SPECIAL_REGISTER_R32<"nctaid.z",
1685                                                         int_ptx_read_nctaid_z>;
1686 def PTX_READ_NCTAID_W   : PTX_READ_SPECIAL_REGISTER_R32<"nctaid.w",
1687                                                         int_ptx_read_nctaid_w>;
1688
1689 def PTX_READ_SMID  : PTX_READ_SPECIAL_REGISTER_R32<"smid",
1690                                                    int_ptx_read_smid>;
1691 def PTX_READ_NSMID  : PTX_READ_SPECIAL_REGISTER_R32<"nsmid",
1692                                                     int_ptx_read_nsmid>;
1693 def PTX_READ_GRIDID  : PTX_READ_SPECIAL_REGISTER_R32<"gridid",
1694                                                      int_ptx_read_gridid>;
1695
1696 def PTX_READ_LANEMASK_EQ
1697   : PTX_READ_SPECIAL_REGISTER_R32<"lanemask_eq", int_ptx_read_lanemask_eq>;
1698 def PTX_READ_LANEMASK_LE
1699   : PTX_READ_SPECIAL_REGISTER_R32<"lanemask_le", int_ptx_read_lanemask_le>;
1700 def PTX_READ_LANEMASK_LT
1701   : PTX_READ_SPECIAL_REGISTER_R32<"lanemask_lt", int_ptx_read_lanemask_lt>;
1702 def PTX_READ_LANEMASK_GE
1703   : PTX_READ_SPECIAL_REGISTER_R32<"lanemask_ge", int_ptx_read_lanemask_ge>;
1704 def PTX_READ_LANEMASK_GT
1705   : PTX_READ_SPECIAL_REGISTER_R32<"lanemask_gt", int_ptx_read_lanemask_gt>;
1706
1707 def PTX_READ_CLOCK
1708   : PTX_READ_SPECIAL_REGISTER_R32<"clock", int_ptx_read_clock>;
1709 def PTX_READ_CLOCK64
1710   : PTX_READ_SPECIAL_REGISTER_R64<"clock64", int_ptx_read_clock64>;
1711
1712 def PTX_READ_PM0 : PTX_READ_SPECIAL_REGISTER_R32<"pm0", int_ptx_read_pm0>;
1713 def PTX_READ_PM1 : PTX_READ_SPECIAL_REGISTER_R32<"pm1", int_ptx_read_pm1>;
1714 def PTX_READ_PM2 : PTX_READ_SPECIAL_REGISTER_R32<"pm2", int_ptx_read_pm2>;
1715 def PTX_READ_PM3 : PTX_READ_SPECIAL_REGISTER_R32<"pm3", int_ptx_read_pm3>;
1716
1717 // PTX Parallel Synchronization and Communication Intrinsics
1718
1719 def PTX_BAR_SYNC : NVPTXInst<(outs), (ins i32imm:$i), "bar.sync\t$i;",
1720                              [(int_ptx_bar_sync imm:$i)]>;