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 // Note that max(0.0, min(x, 1.0)) cannot be mapped to sat(x) because when x is
  86 // NaN
  87 // max(0.0, min(x, 1.0)) is 1.0 while sat(x) is 0.
  88 // Same story for fmax, fmin.
  89
  90 def : Pat<(int_nvvm_fmin_f immFloat1,
  91             (int_nvvm_fmax_f immFloat0, Float32Regs:$a)),
  92           (CVT_f32_f32 Float32Regs:$a, CvtSAT)>;
  93 def : Pat<(int_nvvm_fmin_f immFloat1,
  94             (int_nvvm_fmax_f Float32Regs:$a, immFloat0)),
  95           (CVT_f32_f32 Float32Regs:$a, CvtSAT)>;
  96 def : Pat<(int_nvvm_fmin_f
  97             (int_nvvm_fmax_f immFloat0, Float32Regs:$a), immFloat1),
  98           (CVT_f32_f32 Float32Regs:$a, CvtSAT)>;
  99 def : Pat<(int_nvvm_fmin_f
 100             (int_nvvm_fmax_f Float32Regs:$a, immFloat0), immFloat1),
 101           (CVT_f32_f32 Float32Regs:$a, CvtSAT)>;
 102
 103 def : Pat<(int_nvvm_fmin_d immDouble1,
 104             (int_nvvm_fmax_d immDouble0, Float64Regs:$a)),
 105           (CVT_f64_f64 Float64Regs:$a, CvtSAT)>;
 106 def : Pat<(int_nvvm_fmin_d immDouble1,
 107             (int_nvvm_fmax_d Float64Regs:$a, immDouble0)),
 108           (CVT_f64_f64 Float64Regs:$a, CvtSAT)>;
 109 def : Pat<(int_nvvm_fmin_d
 110             (int_nvvm_fmax_d immDouble0, Float64Regs:$a), immDouble1),
 111           (CVT_f64_f64 Float64Regs:$a, CvtSAT)>;
 112 def : Pat<(int_nvvm_fmin_d
 113             (int_nvvm_fmax_d Float64Regs:$a, immDouble0), immDouble1),
 114           (CVT_f64_f64 Float64Regs:$a, CvtSAT)>;
 115
 116
 117 // We need a full string for OpcStr here because we need to deal with case like
 118 // INT_PTX_RECIP.
 119 class F_MATH_1<string OpcStr, NVPTXRegClass target_regclass,
 120   NVPTXRegClass src_regclass, Intrinsic IntOP>
 121             : NVPTXInst<(outs target_regclass:$dst), (ins src_regclass:$src0),
 122             OpcStr,
 123         [(set target_regclass:$dst, (IntOP src_regclass:$src0))]>;
 124
 125 // We need a full string for OpcStr here because we need to deal with the case
 126 // like INT_PTX_NATIVE_POWR_F.
 127 class F_MATH_2<string OpcStr, NVPTXRegClass t_regclass,
 128   NVPTXRegClass s0_regclass, NVPTXRegClass s1_regclass, Intrinsic IntOP>
 129             : NVPTXInst<(outs t_regclass:$dst),
 130               (ins s0_regclass:$src0, s1_regclass:$src1),
 131             OpcStr,
 132         [(set t_regclass:$dst, (IntOP s0_regclass:$src0, s1_regclass:$src1))]>;
 133
 134 class F_MATH_3<string OpcStr, NVPTXRegClass t_regclass,
 135   NVPTXRegClass s0_regclass, NVPTXRegClass s1_regclass,
 136   NVPTXRegClass s2_regclass, Intrinsic IntOP>
 137             : NVPTXInst<(outs t_regclass:$dst),
 138               (ins s0_regclass:$src0, s1_regclass:$src1, s2_regclass:$src2),
 139             OpcStr,
 140         [(set t_regclass:$dst,
 141           (IntOP s0_regclass:$src0, s1_regclass:$src1, s2_regclass:$src2))]>;
 142
 143 //
 144 // MISC
 145 //
 146
 147 def INT_NVVM_CLZ_I : F_MATH_1<"clz.b32 \t$dst, $src0;", Int32Regs, Int32Regs,
 148   int_nvvm_clz_i>;
 149 def INT_NVVM_CLZ_LL : F_MATH_1<"clz.b64 \t$dst, $src0;", Int32Regs, Int64Regs,
 150   int_nvvm_clz_ll>;
 151
 152 def INT_NVVM_POPC_I : F_MATH_1<"popc.b32 \t$dst, $src0;", Int32Regs, Int32Regs,
 153   int_nvvm_popc_i>;
 154 def INT_NVVM_POPC_LL : F_MATH_1<"popc.b64 \t$dst, $src0;", Int32Regs, Int64Regs,
 155   int_nvvm_popc_ll>;
 156
 157 def INT_NVVM_PRMT : F_MATH_3<"prmt.b32 \t$dst, $src0, $src1, $src2;", Int32Regs,
 158   Int32Regs, Int32Regs, Int32Regs, int_nvvm_prmt>;
 159
 160 //
 161 // Min Max
 162 //
 163
 164 def INT_NVVM_MIN_I : F_MATH_2<"min.s32 \t$dst, $src0, $src1;", Int32Regs,
 165   Int32Regs, Int32Regs, int_nvvm_min_i>;
 166 def INT_NVVM_MIN_UI : F_MATH_2<"min.u32 \t$dst, $src0, $src1;", Int32Regs,
 167   Int32Regs, Int32Regs, int_nvvm_min_ui>;
 168
 169 def INT_NVVM_MIN_LL : F_MATH_2<"min.s64 \t$dst, $src0, $src1;", Int64Regs,
 170   Int64Regs, Int64Regs, int_nvvm_min_ll>;
 171 def INT_NVVM_MIN_ULL : F_MATH_2<"min.u64 \t$dst, $src0, $src1;", Int64Regs,
 172   Int64Regs, Int64Regs, int_nvvm_min_ull>;
 173
 174 def INT_NVVM_MAX_I : F_MATH_2<"max.s32 \t$dst, $src0, $src1;", Int32Regs,
 175   Int32Regs, Int32Regs, int_nvvm_max_i>;
 176 def INT_NVVM_MAX_UI : F_MATH_2<"max.u32 \t$dst, $src0, $src1;", Int32Regs,
 177   Int32Regs, Int32Regs, int_nvvm_max_ui>;
 178
 179 def INT_NVVM_MAX_LL : F_MATH_2<"max.s64 \t$dst, $src0, $src1;", Int64Regs,
 180   Int64Regs, Int64Regs, int_nvvm_max_ll>;
 181 def INT_NVVM_MAX_ULL : F_MATH_2<"max.u64 \t$dst, $src0, $src1;", Int64Regs,
 182   Int64Regs, Int64Regs, int_nvvm_max_ull>;
 183
 184 def INT_NVVM_FMIN_F : F_MATH_2<"min.f32 \t$dst, $src0, $src1;", Float32Regs,
 185   Float32Regs, Float32Regs, int_nvvm_fmin_f>;
 186 def INT_NVVM_FMIN_FTZ_F : F_MATH_2<"min.ftz.f32 \t$dst, $src0, $src1;",
 187   Float32Regs, Float32Regs, Float32Regs, int_nvvm_fmin_ftz_f>;
 188
 189 def INT_NVVM_FMAX_F : F_MATH_2<"max.f32 \t$dst, $src0, $src1;", Float32Regs,
 190   Float32Regs, Float32Regs, int_nvvm_fmax_f>;
 191 def INT_NVVM_FMAX_FTZ_F : F_MATH_2<"max.ftz.f32 \t$dst, $src0, $src1;",
 192   Float32Regs, Float32Regs, Float32Regs, int_nvvm_fmax_ftz_f>;
 193
 194 def INT_NVVM_FMIN_D : F_MATH_2<"min.f64 \t$dst, $src0, $src1;", Float64Regs,
 195   Float64Regs, Float64Regs, int_nvvm_fmin_d>;
 196 def INT_NVVM_FMAX_D : F_MATH_2<"max.f64 \t$dst, $src0, $src1;", Float64Regs,
 197   Float64Regs, Float64Regs, int_nvvm_fmax_d>;
 198
 199 //
 200 // Multiplication
 201 //
 202
 203 def INT_NVVM_MULHI_I : F_MATH_2<"mul.hi.s32 \t$dst, $src0, $src1;", Int32Regs,
 204   Int32Regs, Int32Regs, int_nvvm_mulhi_i>;
 205 def INT_NVVM_MULHI_UI : F_MATH_2<"mul.hi.u32 \t$dst, $src0, $src1;", Int32Regs,
 206   Int32Regs, Int32Regs, int_nvvm_mulhi_ui>;
 207
 208 def INT_NVVM_MULHI_LL : F_MATH_2<"mul.hi.s64 \t$dst, $src0, $src1;", Int64Regs,
 209   Int64Regs, Int64Regs, int_nvvm_mulhi_ll>;
 210 def INT_NVVM_MULHI_ULL : F_MATH_2<"mul.hi.u64 \t$dst, $src0, $src1;", Int64Regs,
 211   Int64Regs, Int64Regs, int_nvvm_mulhi_ull>;
 212
 213 def INT_NVVM_MUL_RN_FTZ_F : F_MATH_2<"mul.rn.ftz.f32 \t$dst, $src0, $src1;",
 214   Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rn_ftz_f>;
 215 def INT_NVVM_MUL_RN_F : F_MATH_2<"mul.rn.f32 \t$dst, $src0, $src1;",
 216   Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rn_f>;
 217 def INT_NVVM_MUL_RZ_FTZ_F : F_MATH_2<"mul.rz.ftz.f32 \t$dst, $src0, $src1;",
 218   Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rz_ftz_f>;
 219 def INT_NVVM_MUL_RZ_F : F_MATH_2<"mul.rz.f32 \t$dst, $src0, $src1;",
 220   Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rz_f>;
 221 def INT_NVVM_MUL_RM_FTZ_F : F_MATH_2<"mul.rm.ftz.f32 \t$dst, $src0, $src1;",
 222   Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rm_ftz_f>;
 223 def INT_NVVM_MUL_RM_F : F_MATH_2<"mul.rm.f32 \t$dst, $src0, $src1;",
 224   Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rm_f>;
 225 def INT_NVVM_MUL_RP_FTZ_F : F_MATH_2<"mul.rp.ftz.f32 \t$dst, $src0, $src1;",
 226   Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rp_ftz_f>;
 227 def INT_NVVM_MUL_RP_F : F_MATH_2<"mul.rp.f32 \t$dst, $src0, $src1;",
 228   Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rp_f>;
 229
 230 def INT_NVVM_MUL_RN_D : F_MATH_2<"mul.rn.f64 \t$dst, $src0, $src1;",
 231   Float64Regs, Float64Regs, Float64Regs, int_nvvm_mul_rn_d>;
 232 def INT_NVVM_MUL_RZ_D : F_MATH_2<"mul.rz.f64 \t$dst, $src0, $src1;",
 233   Float64Regs, Float64Regs, Float64Regs, int_nvvm_mul_rz_d>;
 234 def INT_NVVM_MUL_RM_D : F_MATH_2<"mul.rm.f64 \t$dst, $src0, $src1;",
 235   Float64Regs, Float64Regs, Float64Regs, int_nvvm_mul_rm_d>;
 236 def INT_NVVM_MUL_RP_D : F_MATH_2<"mul.rp.f64 \t$dst, $src0, $src1;",
 237   Float64Regs, Float64Regs, Float64Regs, int_nvvm_mul_rp_d>;
 238
 239 def INT_NVVM_MUL24_I : F_MATH_2<"mul24.lo.s32 \t$dst, $src0, $src1;",
 240   Int32Regs, Int32Regs, Int32Regs, int_nvvm_mul24_i>;
 241 def INT_NVVM_MUL24_UI : F_MATH_2<"mul24.lo.u32 \t$dst, $src0, $src1;",
 242   Int32Regs, Int32Regs, Int32Regs, int_nvvm_mul24_ui>;
 243
 244 //
 245 // Div
 246 //
 247
 248 def INT_NVVM_DIV_APPROX_FTZ_F
 249   : F_MATH_2<"div.approx.ftz.f32 \t$dst, $src0, $src1;", Float32Regs,
 250     Float32Regs, Float32Regs, int_nvvm_div_approx_ftz_f>;
 251 def INT_NVVM_DIV_APPROX_F : F_MATH_2<"div.approx.f32 \t$dst, $src0, $src1;",
 252   Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_approx_f>;
 253
 254 def INT_NVVM_DIV_RN_FTZ_F : F_MATH_2<"div.rn.ftz.f32 \t$dst, $src0, $src1;",
 255   Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rn_ftz_f>;
 256 def INT_NVVM_DIV_RN_F     : F_MATH_2<"div.rn.f32 \t$dst, $src0, $src1;",
 257   Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rn_f>;
 258 def INT_NVVM_DIV_RZ_FTZ_F : F_MATH_2<"div.rz.ftz.f32 \t$dst, $src0, $src1;",
 259   Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rz_ftz_f>;
 260 def INT_NVVM_DIV_RZ_F     : F_MATH_2<"div.rz.f32 \t$dst, $src0, $src1;",
 261   Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rz_f>;
 262 def INT_NVVM_DIV_RM_FTZ_F : F_MATH_2<"div.rm.ftz.f32 \t$dst, $src0, $src1;",
 263   Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rm_ftz_f>;
 264 def INT_NVVM_DIV_RM_F     : F_MATH_2<"div.rm.f32 \t$dst, $src0, $src1;",
 265   Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rm_f>;
 266 def INT_NVVM_DIV_RP_FTZ_F : F_MATH_2<"div.rp.ftz.f32 \t$dst, $src0, $src1;",
 267   Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rp_ftz_f>;
 268 def INT_NVVM_DIV_RP_F     : F_MATH_2<"div.rp.f32 \t$dst, $src0, $src1;",
 269   Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rp_f>;
 270
 271 def INT_NVVM_DIV_RN_D : F_MATH_2<"div.rn.f64 \t$dst, $src0, $src1;",
 272   Float64Regs, Float64Regs, Float64Regs, int_nvvm_div_rn_d>;
 273 def INT_NVVM_DIV_RZ_D : F_MATH_2<"div.rz.f64 \t$dst, $src0, $src1;",
 274   Float64Regs, Float64Regs, Float64Regs, int_nvvm_div_rz_d>;
 275 def INT_NVVM_DIV_RM_D : F_MATH_2<"div.rm.f64 \t$dst, $src0, $src1;",
 276   Float64Regs, Float64Regs, Float64Regs, int_nvvm_div_rm_d>;
 277 def INT_NVVM_DIV_RP_D : F_MATH_2<"div.rp.f64 \t$dst, $src0, $src1;",
 278   Float64Regs, Float64Regs, Float64Regs, int_nvvm_div_rp_d>;
 279
 280 //
 281 // Brev
 282 //
 283
 284 def INT_NVVM_BREV32 : F_MATH_1<"brev.b32 \t$dst, $src0;", Int32Regs, Int32Regs,
 285   int_nvvm_brev32>;
 286 def INT_NVVM_BREV64 : F_MATH_1<"brev.b64 \t$dst, $src0;", Int64Regs, Int64Regs,
 287   int_nvvm_brev64>;
 288
 289 //
 290 // Sad
 291 //
 292
 293 def INT_NVVM_SAD_I : F_MATH_3<"sad.s32 \t$dst, $src0, $src1, $src2;",
 294   Int32Regs, Int32Regs, Int32Regs, Int32Regs, int_nvvm_sad_i>;
 295 def INT_NVVM_SAD_UI : F_MATH_3<"sad.u32 \t$dst, $src0, $src1, $src2;",
 296   Int32Regs, Int32Regs, Int32Regs, Int32Regs, int_nvvm_sad_ui>;
 297
 298 //
 299 // Floor  Ceil
 300 //
 301
 302 def : Pat<(int_nvvm_floor_ftz_f Float32Regs:$a),
 303           (CVT_f32_f32 Float32Regs:$a, CvtRMI_FTZ)>;
 304 def : Pat<(int_nvvm_floor_f Float32Regs:$a),
 305           (CVT_f32_f32 Float32Regs:$a, CvtRMI)>;
 306 def : Pat<(int_nvvm_floor_d Float64Regs:$a),
 307           (CVT_f64_f64 Float64Regs:$a, CvtRMI)>;
 308
 309 def : Pat<(int_nvvm_ceil_ftz_f Float32Regs:$a),
 310           (CVT_f32_f32 Float32Regs:$a, CvtRPI_FTZ)>;
 311 def : Pat<(int_nvvm_ceil_f Float32Regs:$a),
 312           (CVT_f32_f32 Float32Regs:$a, CvtRPI)>;
 313 def : Pat<(int_nvvm_ceil_d Float64Regs:$a),
 314           (CVT_f64_f64 Float64Regs:$a, CvtRPI)>;
 315
 316 //
 317 // Abs
 318 //
 319
 320 def INT_NVVM_ABS_I : F_MATH_1<"abs.s32 \t$dst, $src0;", Int32Regs, Int32Regs,
 321   int_nvvm_abs_i>;
 322 def INT_NVVM_ABS_LL : F_MATH_1<"abs.s64 \t$dst, $src0;", Int64Regs, Int64Regs,
 323   int_nvvm_abs_ll>;
 324
 325 def INT_NVVM_FABS_FTZ_F : F_MATH_1<"abs.ftz.f32 \t$dst, $src0;", Float32Regs,
 326   Float32Regs, int_nvvm_fabs_ftz_f>;
 327 def INT_NVVM_FABS_F : F_MATH_1<"abs.f32 \t$dst, $src0;", Float32Regs,
 328   Float32Regs, int_nvvm_fabs_f>;
 329
 330 def INT_NVVM_FABS_D : F_MATH_1<"abs.f64 \t$dst, $src0;", Float64Regs,
 331   Float64Regs, int_nvvm_fabs_d>;
 332
 333 //
 334 // Round
 335 //
 336
 337 def : Pat<(int_nvvm_round_ftz_f Float32Regs:$a),
 338           (CVT_f32_f32 Float32Regs:$a, CvtRNI_FTZ)>;
 339 def : Pat<(int_nvvm_round_f Float32Regs:$a),
 340           (CVT_f32_f32 Float32Regs:$a, CvtRNI)>;
 341 def : Pat<(int_nvvm_round_d Float64Regs:$a),
 342           (CVT_f64_f64 Float64Regs:$a, CvtRNI)>;
 343
 344 //
 345 // Trunc
 346 //
 347
 348 def : Pat<(int_nvvm_trunc_ftz_f Float32Regs:$a),
 349           (CVT_f32_f32 Float32Regs:$a, CvtRZI_FTZ)>;
 350 def : Pat<(int_nvvm_trunc_f Float32Regs:$a),
 351           (CVT_f32_f32 Float32Regs:$a, CvtRZI)>;
 352 def : Pat<(int_nvvm_trunc_d Float64Regs:$a),
 353           (CVT_f64_f64 Float64Regs:$a, CvtRZI)>;
 354
 355 //
 356 // Saturate
 357 //
 358
 359 def : Pat<(int_nvvm_saturate_ftz_f Float32Regs:$a),
 360           (CVT_f32_f32 Float32Regs:$a, CvtSAT_FTZ)>;
 361 def : Pat<(int_nvvm_saturate_f Float32Regs:$a),
 362           (CVT_f32_f32 Float32Regs:$a, CvtSAT)>;
 363 def : Pat<(int_nvvm_saturate_d Float64Regs:$a),
 364           (CVT_f64_f64 Float64Regs:$a, CvtSAT)>;
 365
 366 //
 367 // Exp2  Log2
 368 //
 369
 370 def INT_NVVM_EX2_APPROX_FTZ_F : F_MATH_1<"ex2.approx.ftz.f32 \t$dst, $src0;",
 371   Float32Regs, Float32Regs, int_nvvm_ex2_approx_ftz_f>;
 372 def INT_NVVM_EX2_APPROX_F : F_MATH_1<"ex2.approx.f32 \t$dst, $src0;",
 373   Float32Regs, Float32Regs, int_nvvm_ex2_approx_f>;
 374 def INT_NVVM_EX2_APPROX_D : F_MATH_1<"ex2.approx.f64 \t$dst, $src0;",
 375   Float64Regs, Float64Regs, int_nvvm_ex2_approx_d>;
 376
 377 def INT_NVVM_LG2_APPROX_FTZ_F : F_MATH_1<"lg2.approx.ftz.f32 \t$dst, $src0;",
 378   Float32Regs, Float32Regs, int_nvvm_lg2_approx_ftz_f>;
 379 def INT_NVVM_LG2_APPROX_F : F_MATH_1<"lg2.approx.f32 \t$dst, $src0;",
 380   Float32Regs, Float32Regs, int_nvvm_lg2_approx_f>;
 381 def INT_NVVM_LG2_APPROX_D : F_MATH_1<"lg2.approx.f64 \t$dst, $src0;",
 382   Float64Regs, Float64Regs, int_nvvm_lg2_approx_d>;
 383
 384 //
 385 // Sin  Cos
 386 //
 387
 388 def INT_NVVM_SIN_APPROX_FTZ_F : F_MATH_1<"sin.approx.ftz.f32 \t$dst, $src0;",
 389   Float32Regs, Float32Regs, int_nvvm_sin_approx_ftz_f>;
 390 def INT_NVVM_SIN_APPROX_F : F_MATH_1<"sin.approx.f32 \t$dst, $src0;",
 391   Float32Regs, Float32Regs, int_nvvm_sin_approx_f>;
 392
 393 def INT_NVVM_COS_APPROX_FTZ_F : F_MATH_1<"cos.approx.ftz.f32 \t$dst, $src0;",
 394   Float32Regs, Float32Regs, int_nvvm_cos_approx_ftz_f>;
 395 def INT_NVVM_COS_APPROX_F : F_MATH_1<"cos.approx.f32 \t$dst, $src0;",
 396   Float32Regs, Float32Regs, int_nvvm_cos_approx_f>;
 397
 398 //
 399 // Fma
 400 //
 401
 402 def INT_NVVM_FMA_RN_FTZ_F
 403   : F_MATH_3<"fma.rn.ftz.f32 \t$dst, $src0, $src1, $src2;", Float32Regs,
 404     Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rn_ftz_f>;
 405 def INT_NVVM_FMA_RN_F : F_MATH_3<"fma.rn.f32 \t$dst, $src0, $src1, $src2;",
 406   Float32Regs, Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rn_f>;
 407 def INT_NVVM_FMA_RZ_FTZ_F
 408   : F_MATH_3<"fma.rz.ftz.f32 \t$dst, $src0, $src1, $src2;", Float32Regs,
 409     Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rz_ftz_f>;
 410 def INT_NVVM_FMA_RZ_F : F_MATH_3<"fma.rz.f32 \t$dst, $src0, $src1, $src2;",
 411   Float32Regs, Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rz_f>;
 412 def INT_NVVM_FMA_RM_FTZ_F
 413   : F_MATH_3<"fma.rm.ftz.f32 \t$dst, $src0, $src1, $src2;", Float32Regs,
 414     Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rm_ftz_f>;
 415 def INT_NVVM_FMA_RM_F : F_MATH_3<"fma.rm.f32 \t$dst, $src0, $src1, $src2;",
 416   Float32Regs, Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rm_f>;
 417 def INT_NVVM_FMA_RP_FTZ_F
 418   : F_MATH_3<"fma.rp.ftz.f32 \t$dst, $src0, $src1, $src2;", Float32Regs,
 419     Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rp_ftz_f>;
 420 def INT_NVVM_FMA_RP_F : F_MATH_3<"fma.rp.f32 \t$dst, $src0, $src1, $src2;",
 421   Float32Regs, Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rp_f>;
 422
 423 def INT_NVVM_FMA_RN_D : F_MATH_3<"fma.rn.f64 \t$dst, $src0, $src1, $src2;",
 424   Float64Regs, Float64Regs, Float64Regs, Float64Regs, int_nvvm_fma_rn_d>;
 425 def INT_NVVM_FMA_RZ_D : F_MATH_3<"fma.rz.f64 \t$dst, $src0, $src1, $src2;",
 426   Float64Regs, Float64Regs, Float64Regs, Float64Regs, int_nvvm_fma_rz_d>;
 427 def INT_NVVM_FMA_RM_D : F_MATH_3<"fma.rm.f64 \t$dst, $src0, $src1, $src2;",
 428   Float64Regs, Float64Regs, Float64Regs, Float64Regs, int_nvvm_fma_rm_d>;
 429 def INT_NVVM_FMA_RP_D : F_MATH_3<"fma.rp.f64 \t$dst, $src0, $src1, $src2;",
 430   Float64Regs, Float64Regs, Float64Regs, Float64Regs, int_nvvm_fma_rp_d>;
 431
 432 //
 433 // Rcp
 434 //
 435
 436 def INT_NVVM_RCP_RN_FTZ_F : F_MATH_1<"rcp.rn.ftz.f32 \t$dst, $src0;",
 437   Float32Regs, Float32Regs, int_nvvm_rcp_rn_ftz_f>;
 438 def INT_NVVM_RCP_RN_F : F_MATH_1<"rcp.rn.f32 \t$dst, $src0;",
 439   Float32Regs, Float32Regs, int_nvvm_rcp_rn_f>;
 440 def INT_NVVM_RCP_RZ_FTZ_F : F_MATH_1<"rcp.rz.ftz.f32 \t$dst, $src0;",
 441   Float32Regs, Float32Regs, int_nvvm_rcp_rz_ftz_f>;
 442 def INT_NVVM_RCP_RZ_F : F_MATH_1<"rcp.rz.f32 \t$dst, $src0;",
 443   Float32Regs, Float32Regs, int_nvvm_rcp_rz_f>;
 444 def INT_NVVM_RCP_RM_FTZ_F : F_MATH_1<"rcp.rm.ftz.f32 \t$dst, $src0;",
 445   Float32Regs, Float32Regs, int_nvvm_rcp_rm_ftz_f>;
 446 def INT_NVVM_RCP_RM_F : F_MATH_1<"rcp.rm.f32 \t$dst, $src0;",
 447   Float32Regs, Float32Regs, int_nvvm_rcp_rm_f>;
 448 def INT_NVVM_RCP_RP_FTZ_F : F_MATH_1<"rcp.rp.ftz.f32 \t$dst, $src0;",
 449   Float32Regs, Float32Regs, int_nvvm_rcp_rp_ftz_f>;
 450 def INT_NVVM_RCP_RP_F : F_MATH_1<"rcp.rp.f32 \t$dst, $src0;",
 451   Float32Regs, Float32Regs, int_nvvm_rcp_rp_f>;
 452
 453 def INT_NVVM_RCP_RN_D : F_MATH_1<"rcp.rn.f64 \t$dst, $src0;", Float64Regs,
 454   Float64Regs, int_nvvm_rcp_rn_d>;
 455 def INT_NVVM_RCP_RZ_D : F_MATH_1<"rcp.rz.f64 \t$dst, $src0;", Float64Regs,
 456   Float64Regs, int_nvvm_rcp_rz_d>;
 457 def INT_NVVM_RCP_RM_D : F_MATH_1<"rcp.rm.f64 \t$dst, $src0;", Float64Regs,
 458   Float64Regs, int_nvvm_rcp_rm_d>;
 459 def INT_NVVM_RCP_RP_D : F_MATH_1<"rcp.rp.f64 \t$dst, $src0;", Float64Regs,
 460   Float64Regs, int_nvvm_rcp_rp_d>;
 461
 462 def INT_NVVM_RCP_APPROX_FTZ_D : F_MATH_1<"rcp.approx.ftz.f64 \t$dst, $src0;",
 463   Float64Regs, Float64Regs, int_nvvm_rcp_approx_ftz_d>;
 464
 465 //
 466 // Sqrt
 467 //
 468
 469 def INT_NVVM_SQRT_RN_FTZ_F : F_MATH_1<"sqrt.rn.ftz.f32 \t$dst, $src0;",
 470   Float32Regs, Float32Regs, int_nvvm_sqrt_rn_ftz_f>;
 471 def INT_NVVM_SQRT_RN_F : F_MATH_1<"sqrt.rn.f32 \t$dst, $src0;", Float32Regs,
 472   Float32Regs, int_nvvm_sqrt_rn_f>;
 473 def INT_NVVM_SQRT_RZ_FTZ_F : F_MATH_1<"sqrt.rz.ftz.f32 \t$dst, $src0;",
 474   Float32Regs, Float32Regs, int_nvvm_sqrt_rz_ftz_f>;
 475 def INT_NVVM_SQRT_RZ_F : F_MATH_1<"sqrt.rz.f32 \t$dst, $src0;", Float32Regs,
 476   Float32Regs, int_nvvm_sqrt_rz_f>;
 477 def INT_NVVM_SQRT_RM_FTZ_F : F_MATH_1<"sqrt.rm.ftz.f32 \t$dst, $src0;",
 478   Float32Regs, Float32Regs, int_nvvm_sqrt_rm_ftz_f>;
 479 def INT_NVVM_SQRT_RM_F : F_MATH_1<"sqrt.rm.f32 \t$dst, $src0;", Float32Regs,
 480   Float32Regs, int_nvvm_sqrt_rm_f>;
 481 def INT_NVVM_SQRT_RP_FTZ_F : F_MATH_1<"sqrt.rp.ftz.f32 \t$dst, $src0;",
 482   Float32Regs, Float32Regs, int_nvvm_sqrt_rp_ftz_f>;
 483 def INT_NVVM_SQRT_RP_F : F_MATH_1<"sqrt.rp.f32 \t$dst, $src0;", Float32Regs,
 484   Float32Regs, int_nvvm_sqrt_rp_f>;
 485 def INT_NVVM_SQRT_APPROX_FTZ_F : F_MATH_1<"sqrt.approx.ftz.f32 \t$dst, $src0;",
 486   Float32Regs, Float32Regs, int_nvvm_sqrt_approx_ftz_f>;
 487 def INT_NVVM_SQRT_APPROX_F : F_MATH_1<"sqrt.approx.f32 \t$dst, $src0;",
 488   Float32Regs, Float32Regs, int_nvvm_sqrt_approx_f>;
 489
 490 def INT_NVVM_SQRT_RN_D : F_MATH_1<"sqrt.rn.f64 \t$dst, $src0;", Float64Regs,
 491   Float64Regs, int_nvvm_sqrt_rn_d>;
 492 def INT_NVVM_SQRT_RZ_D : F_MATH_1<"sqrt.rz.f64 \t$dst, $src0;", Float64Regs,
 493   Float64Regs, int_nvvm_sqrt_rz_d>;
 494 def INT_NVVM_SQRT_RM_D : F_MATH_1<"sqrt.rm.f64 \t$dst, $src0;", Float64Regs,
 495   Float64Regs, int_nvvm_sqrt_rm_d>;
 496 def INT_NVVM_SQRT_RP_D : F_MATH_1<"sqrt.rp.f64 \t$dst, $src0;", Float64Regs,
 497   Float64Regs, int_nvvm_sqrt_rp_d>;
 498
 499 // nvvm_sqrt intrinsic
 500 def : Pat<(int_nvvm_sqrt_f Float32Regs:$a),
 501           (INT_NVVM_SQRT_RN_FTZ_F Float32Regs:$a)>, Requires<[doF32FTZ, do_SQRTF32_RN]>;
 502 def : Pat<(int_nvvm_sqrt_f Float32Regs:$a),
 503           (INT_NVVM_SQRT_RN_F Float32Regs:$a)>, Requires<[do_SQRTF32_RN]>;
 504 def : Pat<(int_nvvm_sqrt_f Float32Regs:$a),
 505           (INT_NVVM_SQRT_APPROX_FTZ_F Float32Regs:$a)>, Requires<[doF32FTZ]>;
 506 def : Pat<(int_nvvm_sqrt_f Float32Regs:$a),
 507           (INT_NVVM_SQRT_APPROX_F Float32Regs:$a)>;
 508
 509 //
 510 // Rsqrt
 511 //
 512
 513 def INT_NVVM_RSQRT_APPROX_FTZ_F
 514   : F_MATH_1<"rsqrt.approx.ftz.f32 \t$dst, $src0;", Float32Regs, Float32Regs,
 515     int_nvvm_rsqrt_approx_ftz_f>;
 516 def INT_NVVM_RSQRT_APPROX_F : F_MATH_1<"rsqrt.approx.f32 \t$dst, $src0;",
 517   Float32Regs, Float32Regs, int_nvvm_rsqrt_approx_f>;
 518 def INT_NVVM_RSQRT_APPROX_D : F_MATH_1<"rsqrt.approx.f64 \t$dst, $src0;",
 519   Float64Regs, Float64Regs, int_nvvm_rsqrt_approx_d>;
 520
 521 //
 522 // Add
 523 //
 524
 525 def INT_NVVM_ADD_RN_FTZ_F : F_MATH_2<"add.rn.ftz.f32 \t$dst, $src0, $src1;",
 526   Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rn_ftz_f>;
 527 def INT_NVVM_ADD_RN_F : F_MATH_2<"add.rn.f32 \t$dst, $src0, $src1;",
 528   Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rn_f>;
 529 def INT_NVVM_ADD_RZ_FTZ_F : F_MATH_2<"add.rz.ftz.f32 \t$dst, $src0, $src1;",
 530   Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rz_ftz_f>;
 531 def INT_NVVM_ADD_RZ_F : F_MATH_2<"add.rz.f32 \t$dst, $src0, $src1;",
 532   Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rz_f>;
 533 def INT_NVVM_ADD_RM_FTZ_F : F_MATH_2<"add.rm.ftz.f32 \t$dst, $src0, $src1;",
 534   Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rm_ftz_f>;
 535 def INT_NVVM_ADD_RM_F : F_MATH_2<"add.rm.f32 \t$dst, $src0, $src1;",
 536   Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rm_f>;
 537 def INT_NVVM_ADD_RP_FTZ_F : F_MATH_2<"add.rp.ftz.f32 \t$dst, $src0, $src1;",
 538   Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rp_ftz_f>;
 539 def INT_NVVM_ADD_RP_F : F_MATH_2<"add.rp.f32 \t$dst, $src0, $src1;",
 540   Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rp_f>;
 541
 542 def INT_NVVM_ADD_RN_D : F_MATH_2<"add.rn.f64 \t$dst, $src0, $src1;",
 543   Float64Regs, Float64Regs, Float64Regs, int_nvvm_add_rn_d>;
 544 def INT_NVVM_ADD_RZ_D : F_MATH_2<"add.rz.f64 \t$dst, $src0, $src1;",
 545   Float64Regs, Float64Regs, Float64Regs, int_nvvm_add_rz_d>;
 546 def INT_NVVM_ADD_RM_D : F_MATH_2<"add.rm.f64 \t$dst, $src0, $src1;",
 547   Float64Regs, Float64Regs, Float64Regs, int_nvvm_add_rm_d>;
 548 def INT_NVVM_ADD_RP_D : F_MATH_2<"add.rp.f64 \t$dst, $src0, $src1;",
 549   Float64Regs, Float64Regs, Float64Regs, int_nvvm_add_rp_d>;
 550
 551 //
 552 // Convert
 553 //
 554
 555 def : Pat<(int_nvvm_d2f_rn_ftz Float64Regs:$a),
 556           (CVT_f32_f64 Float64Regs:$a, CvtRN_FTZ)>;
 557 def : Pat<(int_nvvm_d2f_rn Float64Regs:$a),
 558           (CVT_f32_f64 Float64Regs:$a, CvtRN)>;
 559 def : Pat<(int_nvvm_d2f_rz_ftz Float64Regs:$a),
 560           (CVT_f32_f64 Float64Regs:$a, CvtRZ_FTZ)>;
 561 def : Pat<(int_nvvm_d2f_rz Float64Regs:$a),
 562           (CVT_f32_f64 Float64Regs:$a, CvtRZ)>;
 563 def : Pat<(int_nvvm_d2f_rm_ftz Float64Regs:$a),
 564           (CVT_f32_f64 Float64Regs:$a, CvtRM_FTZ)>;
 565 def : Pat<(int_nvvm_d2f_rm Float64Regs:$a),
 566           (CVT_f32_f64 Float64Regs:$a, CvtRM)>;
 567 def : Pat<(int_nvvm_d2f_rp_ftz Float64Regs:$a),
 568           (CVT_f32_f64 Float64Regs:$a, CvtRP_FTZ)>;
 569 def : Pat<(int_nvvm_d2f_rp Float64Regs:$a),
 570           (CVT_f32_f64 Float64Regs:$a, CvtRP)>;
 571
 572 def : Pat<(int_nvvm_d2i_rn Float64Regs:$a),
 573           (CVT_s32_f64 Float64Regs:$a, CvtRNI)>;
 574 def : Pat<(int_nvvm_d2i_rz Float64Regs:$a),
 575           (CVT_s32_f64 Float64Regs:$a, CvtRZI)>;
 576 def : Pat<(int_nvvm_d2i_rm Float64Regs:$a),
 577           (CVT_s32_f64 Float64Regs:$a, CvtRMI)>;
 578 def : Pat<(int_nvvm_d2i_rp Float64Regs:$a),
 579           (CVT_s32_f64 Float64Regs:$a, CvtRPI)>;
 580
 581 def : Pat<(int_nvvm_d2ui_rn Float64Regs:$a),
 582           (CVT_u32_f64 Float64Regs:$a, CvtRNI)>;
 583 def : Pat<(int_nvvm_d2ui_rz Float64Regs:$a),
 584           (CVT_u32_f64 Float64Regs:$a, CvtRZI)>;
 585 def : Pat<(int_nvvm_d2ui_rm Float64Regs:$a),
 586           (CVT_u32_f64 Float64Regs:$a, CvtRMI)>;
 587 def : Pat<(int_nvvm_d2ui_rp Float64Regs:$a),
 588           (CVT_u32_f64 Float64Regs:$a, CvtRPI)>;
 589
 590 def : Pat<(int_nvvm_i2d_rn Int32Regs:$a),
 591           (CVT_f64_s32 Int32Regs:$a, CvtRN)>;
 592 def : Pat<(int_nvvm_i2d_rz Int32Regs:$a),
 593           (CVT_f64_s32 Int32Regs:$a, CvtRZ)>;
 594 def : Pat<(int_nvvm_i2d_rm Int32Regs:$a),
 595           (CVT_f64_s32 Int32Regs:$a, CvtRM)>;
 596 def : Pat<(int_nvvm_i2d_rp Int32Regs:$a),
 597           (CVT_f64_s32 Int32Regs:$a, CvtRP)>;
 598
 599 def : Pat<(int_nvvm_ui2d_rn Int32Regs:$a),
 600           (CVT_f64_u32 Int32Regs:$a, CvtRN)>;
 601 def : Pat<(int_nvvm_ui2d_rz Int32Regs:$a),
 602           (CVT_f64_u32 Int32Regs:$a, CvtRZ)>;
 603 def : Pat<(int_nvvm_ui2d_rm Int32Regs:$a),
 604           (CVT_f64_u32 Int32Regs:$a, CvtRM)>;
 605 def : Pat<(int_nvvm_ui2d_rp Int32Regs:$a),
 606           (CVT_f64_u32 Int32Regs:$a, CvtRP)>;
 607
 608 def : Pat<(int_nvvm_f2i_rn_ftz Float32Regs:$a),
 609           (CVT_s32_f32 Float32Regs:$a, CvtRNI_FTZ)>;
 610 def : Pat<(int_nvvm_f2i_rn Float32Regs:$a),
 611           (CVT_s32_f32 Float32Regs:$a, CvtRNI)>;
 612 def : Pat<(int_nvvm_f2i_rz_ftz Float32Regs:$a),
 613           (CVT_s32_f32 Float32Regs:$a, CvtRZI_FTZ)>;
 614 def : Pat<(int_nvvm_f2i_rz Float32Regs:$a),
 615           (CVT_s32_f32 Float32Regs:$a, CvtRZI)>;
 616 def : Pat<(int_nvvm_f2i_rm_ftz Float32Regs:$a),
 617           (CVT_s32_f32 Float32Regs:$a, CvtRMI_FTZ)>;
 618 def : Pat<(int_nvvm_f2i_rm Float32Regs:$a),
 619           (CVT_s32_f32 Float32Regs:$a, CvtRMI)>;
 620 def : Pat<(int_nvvm_f2i_rp_ftz Float32Regs:$a),
 621           (CVT_s32_f32 Float32Regs:$a, CvtRPI_FTZ)>;
 622 def : Pat<(int_nvvm_f2i_rp Float32Regs:$a),
 623           (CVT_s32_f32 Float32Regs:$a, CvtRPI)>;
 624
 625 def : Pat<(int_nvvm_f2ui_rn_ftz Float32Regs:$a),
 626           (CVT_u32_f32 Float32Regs:$a, CvtRNI_FTZ)>;
 627 def : Pat<(int_nvvm_f2ui_rn Float32Regs:$a),
 628           (CVT_u32_f32 Float32Regs:$a, CvtRNI)>;
 629 def : Pat<(int_nvvm_f2ui_rz_ftz Float32Regs:$a),
 630           (CVT_u32_f32 Float32Regs:$a, CvtRZI_FTZ)>;
 631 def : Pat<(int_nvvm_f2ui_rz Float32Regs:$a),
 632           (CVT_u32_f32 Float32Regs:$a, CvtRZI)>;
 633 def : Pat<(int_nvvm_f2ui_rm_ftz Float32Regs:$a),
 634           (CVT_u32_f32 Float32Regs:$a, CvtRMI_FTZ)>;
 635 def : Pat<(int_nvvm_f2ui_rm Float32Regs:$a),
 636           (CVT_u32_f32 Float32Regs:$a, CvtRMI)>;
 637 def : Pat<(int_nvvm_f2ui_rp_ftz Float32Regs:$a),
 638           (CVT_u32_f32 Float32Regs:$a, CvtRPI_FTZ)>;
 639 def : Pat<(int_nvvm_f2ui_rp Float32Regs:$a),
 640           (CVT_u32_f32 Float32Regs:$a, CvtRPI)>;
 641
 642 def : Pat<(int_nvvm_i2f_rn Int32Regs:$a),
 643           (CVT_f32_s32 Int32Regs:$a, CvtRN)>;
 644 def : Pat<(int_nvvm_i2f_rz Int32Regs:$a),
 645           (CVT_f32_s32 Int32Regs:$a, CvtRZ)>;
 646 def : Pat<(int_nvvm_i2f_rm Int32Regs:$a),
 647           (CVT_f32_s32 Int32Regs:$a, CvtRM)>;
 648 def : Pat<(int_nvvm_i2f_rp Int32Regs:$a),
 649           (CVT_f32_s32 Int32Regs:$a, CvtRP)>;
 650
 651 def : Pat<(int_nvvm_ui2f_rn Int32Regs:$a),
 652           (CVT_f32_u32 Int32Regs:$a, CvtRN)>;
 653 def : Pat<(int_nvvm_ui2f_rz Int32Regs:$a),
 654           (CVT_f32_u32 Int32Regs:$a, CvtRZ)>;
 655 def : Pat<(int_nvvm_ui2f_rm Int32Regs:$a),
 656           (CVT_f32_u32 Int32Regs:$a, CvtRM)>;
 657 def : Pat<(int_nvvm_ui2f_rp Int32Regs:$a),
 658           (CVT_f32_u32 Int32Regs:$a, CvtRP)>;
 659
 660 def INT_NVVM_LOHI_I2D : F_MATH_2<"mov.b64 \t$dst, {{$src0, $src1}};",
 661   Float64Regs, Int32Regs, Int32Regs, int_nvvm_lohi_i2d>;
 662
 663 def INT_NVVM_D2I_LO : F_MATH_1<!strconcat("{{\n\t",
 664                        !strconcat(".reg .b32 %temp; \n\t",
 665              !strconcat("mov.b64 \t{$dst, %temp}, $src0;\n\t",
 666                "}}"))),
 667              Int32Regs, Float64Regs, int_nvvm_d2i_lo>;
 668 def INT_NVVM_D2I_HI : F_MATH_1<!strconcat("{{\n\t",
 669                        !strconcat(".reg .b32 %temp; \n\t",
 670                          !strconcat("mov.b64 \t{%temp, $dst}, $src0;\n\t",
 671                            "}}"))),
 672              Int32Regs, Float64Regs, int_nvvm_d2i_hi>;
 673
 674 def : Pat<(int_nvvm_f2ll_rn_ftz Float32Regs:$a),
 675           (CVT_s64_f32 Float32Regs:$a, CvtRNI_FTZ)>;
 676 def : Pat<(int_nvvm_f2ll_rn Float32Regs:$a),
 677           (CVT_s64_f32 Float32Regs:$a, CvtRNI)>;
 678 def : Pat<(int_nvvm_f2ll_rz_ftz Float32Regs:$a),
 679           (CVT_s64_f32 Float32Regs:$a, CvtRZI_FTZ)>;
 680 def : Pat<(int_nvvm_f2ll_rz Float32Regs:$a),
 681           (CVT_s64_f32 Float32Regs:$a, CvtRZI)>;
 682 def : Pat<(int_nvvm_f2ll_rm_ftz Float32Regs:$a),
 683           (CVT_s64_f32 Float32Regs:$a, CvtRMI_FTZ)>;
 684 def : Pat<(int_nvvm_f2ll_rm Float32Regs:$a),
 685           (CVT_s64_f32 Float32Regs:$a, CvtRMI)>;
 686 def : Pat<(int_nvvm_f2ll_rp_ftz Float32Regs:$a),
 687           (CVT_s64_f32 Float32Regs:$a, CvtRPI_FTZ)>;
 688 def : Pat<(int_nvvm_f2ll_rp Float32Regs:$a),
 689           (CVT_s64_f32 Float32Regs:$a, CvtRPI)>;
 690
 691 def : Pat<(int_nvvm_f2ull_rn_ftz Float32Regs:$a),
 692           (CVT_u64_f32 Float32Regs:$a, CvtRNI_FTZ)>;
 693 def : Pat<(int_nvvm_f2ull_rn Float32Regs:$a),
 694           (CVT_u64_f32 Float32Regs:$a, CvtRNI)>;
 695 def : Pat<(int_nvvm_f2ull_rz_ftz Float32Regs:$a),
 696           (CVT_u64_f32 Float32Regs:$a, CvtRZI_FTZ)>;
 697 def : Pat<(int_nvvm_f2ull_rz Float32Regs:$a),
 698           (CVT_u64_f32 Float32Regs:$a, CvtRZI)>;
 699 def : Pat<(int_nvvm_f2ull_rm_ftz Float32Regs:$a),
 700           (CVT_u64_f32 Float32Regs:$a, CvtRMI_FTZ)>;
 701 def : Pat<(int_nvvm_f2ull_rm Float32Regs:$a),
 702           (CVT_u64_f32 Float32Regs:$a, CvtRMI)>;
 703 def : Pat<(int_nvvm_f2ull_rp_ftz Float32Regs:$a),
 704           (CVT_u64_f32 Float32Regs:$a, CvtRPI_FTZ)>;
 705 def : Pat<(int_nvvm_f2ull_rp Float32Regs:$a),
 706           (CVT_u64_f32 Float32Regs:$a, CvtRPI)>;
 707
 708 def : Pat<(int_nvvm_d2ll_rn Float64Regs:$a),
 709           (CVT_s64_f64 Float64Regs:$a, CvtRNI)>;
 710 def : Pat<(int_nvvm_d2ll_rz Float64Regs:$a),
 711           (CVT_s64_f64 Float64Regs:$a, CvtRZI)>;
 712 def : Pat<(int_nvvm_d2ll_rm Float64Regs:$a),
 713           (CVT_s64_f64 Float64Regs:$a, CvtRMI)>;
 714 def : Pat<(int_nvvm_d2ll_rp Float64Regs:$a),
 715           (CVT_s64_f64 Float64Regs:$a, CvtRPI)>;
 716
 717 def : Pat<(int_nvvm_d2ull_rn Float64Regs:$a),
 718           (CVT_u64_f64 Float64Regs:$a, CvtRNI)>;
 719 def : Pat<(int_nvvm_d2ull_rz Float64Regs:$a),
 720           (CVT_u64_f64 Float64Regs:$a, CvtRZI)>;
 721 def : Pat<(int_nvvm_d2ull_rm Float64Regs:$a),
 722           (CVT_u64_f64 Float64Regs:$a, CvtRMI)>;
 723 def : Pat<(int_nvvm_d2ull_rp Float64Regs:$a),
 724           (CVT_u64_f64 Float64Regs:$a, CvtRPI)>;
 725
 726 def : Pat<(int_nvvm_ll2f_rn Int64Regs:$a),
 727           (CVT_f32_s64 Int64Regs:$a, CvtRN)>;
 728 def : Pat<(int_nvvm_ll2f_rz Int64Regs:$a),
 729           (CVT_f32_s64 Int64Regs:$a, CvtRZ)>;
 730 def : Pat<(int_nvvm_ll2f_rm Int64Regs:$a),
 731           (CVT_f32_s64 Int64Regs:$a, CvtRM)>;
 732 def : Pat<(int_nvvm_ll2f_rp Int64Regs:$a),
 733           (CVT_f32_s64 Int64Regs:$a, CvtRP)>;
 734
 735 def : Pat<(int_nvvm_ull2f_rn Int64Regs:$a),
 736           (CVT_f32_u64 Int64Regs:$a, CvtRN)>;
 737 def : Pat<(int_nvvm_ull2f_rz Int64Regs:$a),
 738           (CVT_f32_u64 Int64Regs:$a, CvtRZ)>;
 739 def : Pat<(int_nvvm_ull2f_rm Int64Regs:$a),
 740           (CVT_f32_u64 Int64Regs:$a, CvtRM)>;
 741 def : Pat<(int_nvvm_ull2f_rp Int64Regs:$a),
 742           (CVT_f32_u64 Int64Regs:$a, CvtRP)>;
 743
 744 def : Pat<(int_nvvm_ll2d_rn Int64Regs:$a),
 745           (CVT_f64_s64 Int64Regs:$a, CvtRN)>;
 746 def : Pat<(int_nvvm_ll2d_rz Int64Regs:$a),
 747           (CVT_f64_s64 Int64Regs:$a, CvtRZ)>;
 748 def : Pat<(int_nvvm_ll2d_rm Int64Regs:$a),
 749           (CVT_f64_s64 Int64Regs:$a, CvtRM)>;
 750 def : Pat<(int_nvvm_ll2d_rp Int64Regs:$a),
 751           (CVT_f64_s64 Int64Regs:$a, CvtRP)>;
 752
 753 def : Pat<(int_nvvm_ull2d_rn Int64Regs:$a),
 754           (CVT_f64_u64 Int64Regs:$a, CvtRN)>;
 755 def : Pat<(int_nvvm_ull2d_rz Int64Regs:$a),
 756           (CVT_f64_u64 Int64Regs:$a, CvtRZ)>;
 757 def : Pat<(int_nvvm_ull2d_rm Int64Regs:$a),
 758           (CVT_f64_u64 Int64Regs:$a, CvtRM)>;
 759 def : Pat<(int_nvvm_ull2d_rp Int64Regs:$a),
 760           (CVT_f64_u64 Int64Regs:$a, CvtRP)>;
 761
 762
 763 // FIXME: Ideally, we could use these patterns instead of the scope-creating
 764 // patterns, but ptxas does not like these since .s16 is not compatible with
 765 // .f16.  The solution is to use .bXX for all integer register types, but we
 766 // are not there yet.
 767 //def : Pat<(int_nvvm_f2h_rn_ftz Float32Regs:$a),
 768 //          (CVT_f16_f32 Float32Regs:$a, CvtRN_FTZ)>;
 769 //def : Pat<(int_nvvm_f2h_rn Float32Regs:$a),
 770 //          (CVT_f16_f32 Float32Regs:$a, CvtRN)>;
 771 //
 772 //def : Pat<(int_nvvm_h2f Int16Regs:$a),
 773 //          (CVT_f32_f16 Int16Regs:$a, CvtNONE)>;
 774
 775 def INT_NVVM_F2H_RN_FTZ : F_MATH_1<!strconcat("{{\n\t",
 776                                    !strconcat(".reg .b16 %temp;\n\t",
 777            !strconcat("cvt.rn.ftz.f16.f32 \t%temp, $src0;\n\t",
 778            !strconcat("mov.b16 \t$dst, %temp;\n",
 779              "}}")))),
 780                                    Int16Regs, Float32Regs, int_nvvm_f2h_rn_ftz>;
 781 def INT_NVVM_F2H_RN : F_MATH_1<!strconcat("{{\n\t",
 782                                    !strconcat(".reg .b16 %temp;\n\t",
 783            !strconcat("cvt.rn.f16.f32 \t%temp, $src0;\n\t",
 784            !strconcat("mov.b16 \t$dst, %temp;\n",
 785              "}}")))),
 786            Int16Regs, Float32Regs, int_nvvm_f2h_rn>;
 787
 788 def INT_NVVM_H2F : F_MATH_1<!strconcat("{{\n\t",
 789                             !strconcat(".reg .b16 %temp;\n\t",
 790           !strconcat("mov.b16 \t%temp, $src0;\n\t",
 791           !strconcat("cvt.f32.f16 \t$dst, %temp;\n\t",
 792             "}}")))),
 793           Float32Regs, Int16Regs, int_nvvm_h2f>;
 794
 795 def : Pat<(f32 (f16_to_f32 Int16Regs:$a)),
 796           (CVT_f32_f16 Int16Regs:$a, CvtNONE)>;
 797 def : Pat<(i16 (f32_to_f16 Float32Regs:$a)),
 798           (CVT_f16_f32 Float32Regs:$a, CvtRN_FTZ)>, Requires<[doF32FTZ]>;
 799 def : Pat<(i16 (f32_to_f16 Float32Regs:$a)),
 800           (CVT_f16_f32 Float32Regs:$a, CvtRN)>;
 801
 802 //
 803 // Bitcast
 804 //
 805
 806 def INT_NVVM_BITCAST_F2I : F_MATH_1<"mov.b32 \t$dst, $src0;", Int32Regs,
 807   Float32Regs, int_nvvm_bitcast_f2i>;
 808 def INT_NVVM_BITCAST_I2F : F_MATH_1<"mov.b32 \t$dst, $src0;", Float32Regs,
 809   Int32Regs, int_nvvm_bitcast_i2f>;
 810
 811 def INT_NVVM_BITCAST_LL2D : F_MATH_1<"mov.b64 \t$dst, $src0;", Float64Regs,
 812   Int64Regs, int_nvvm_bitcast_ll2d>;
 813 def INT_NVVM_BITCAST_D2LL : F_MATH_1<"mov.b64 \t$dst, $src0;", Int64Regs,
 814   Float64Regs, int_nvvm_bitcast_d2ll>;
 815
 816 //-----------------------------------
 817 // Atomic Functions
 818 //-----------------------------------
 819
 820 class ATOMIC_GLOBAL_CHK <dag ops, dag frag>
 821  : PatFrag<ops, frag, [{
 822    return ChkMemSDNodeAddressSpace(N, llvm::ADDRESS_SPACE_GLOBAL);
 823 }]>;
 824 class ATOMIC_SHARED_CHK <dag ops, dag frag>
 825  : PatFrag<ops, frag, [{
 826    return ChkMemSDNodeAddressSpace(N, llvm::ADDRESS_SPACE_SHARED);
 827 }]>;
 828 class ATOMIC_GENERIC_CHK <dag ops, dag frag>
 829  : PatFrag<ops, frag, [{
 830    return ChkMemSDNodeAddressSpace(N, llvm::ADDRESS_SPACE_GENERIC);
 831 }]>;
 832
 833 multiclass F_ATOMIC_2_imp<NVPTXRegClass ptrclass, NVPTXRegClass regclass,
 834   string SpaceStr, string TypeStr, string OpcStr, PatFrag IntOp,
 835   Operand IMMType, SDNode IMM, Predicate Pred> {
 836   def reg : NVPTXInst<(outs regclass:$dst), (ins ptrclass:$addr, regclass:$b),
 837                !strconcat("atom",
 838          !strconcat(SpaceStr,
 839          !strconcat(OpcStr,
 840          !strconcat(TypeStr,
 841          !strconcat(" \t$dst, [$addr], $b;", ""))))),
 842          [(set regclass:$dst, (IntOp ptrclass:$addr, regclass:$b))]>,
 843   Requires<[Pred]>;
 844   def imm : NVPTXInst<(outs regclass:$dst), (ins ptrclass:$addr, IMMType:$b),
 845                !strconcat("atom",
 846          !strconcat(SpaceStr,
 847          !strconcat(OpcStr,
 848          !strconcat(TypeStr,
 849          !strconcat(" \t$dst, [$addr], $b;", ""))))),
 850          [(set regclass:$dst, (IntOp ptrclass:$addr, IMM:$b))]>,
 851   Requires<[Pred]>;
 852 }
 853 multiclass F_ATOMIC_2<NVPTXRegClass regclass, string SpaceStr, string TypeStr,
 854   string OpcStr, PatFrag IntOp, Operand IMMType, SDNode IMM, Predicate Pred> {
 855   defm p32 : F_ATOMIC_2_imp<Int32Regs, regclass, SpaceStr, TypeStr, OpcStr,
 856     IntOp, IMMType, IMM, Pred>;
 857   defm p64 : F_ATOMIC_2_imp<Int64Regs, regclass, SpaceStr, TypeStr, OpcStr,
 858     IntOp, IMMType, IMM, Pred>;
 859 }
 860
 861 // has 2 operands, neg the second one
 862 multiclass F_ATOMIC_2_NEG_imp<NVPTXRegClass ptrclass, NVPTXRegClass regclass,
 863   string SpaceStr, string TypeStr, string OpcStr, PatFrag IntOp,
 864   Operand IMMType, Predicate Pred> {
 865   def reg : NVPTXInst<(outs regclass:$dst), (ins ptrclass:$addr, regclass:$b),
 866     !strconcat("{{ \n\t",
 867          !strconcat(".reg \t.s",
 868          !strconcat(TypeStr,
 869          !strconcat(" temp; \n\t",
 870          !strconcat("neg.s",
 871          !strconcat(TypeStr,
 872          !strconcat(" \ttemp, $b; \n\t",
 873                !strconcat("atom",
 874          !strconcat(SpaceStr,
 875          !strconcat(OpcStr,
 876          !strconcat(".u",
 877          !strconcat(TypeStr,
 878          !strconcat(" \t$dst, [$addr], temp; \n\t",
 879            !strconcat("}}", "")))))))))))))),
 880          [(set regclass:$dst, (IntOp ptrclass:$addr, regclass:$b))]>,
 881   Requires<[Pred]>;
 882 }
 883 multiclass F_ATOMIC_2_NEG<NVPTXRegClass regclass, string SpaceStr,
 884   string TypeStr, string OpcStr, PatFrag IntOp, Operand IMMType,
 885   Predicate Pred> {
 886  defm p32: F_ATOMIC_2_NEG_imp<Int32Regs, regclass, SpaceStr, TypeStr, OpcStr,
 887    IntOp, IMMType, Pred> ;
 888  defm p64: F_ATOMIC_2_NEG_imp<Int64Regs, regclass, SpaceStr, TypeStr, OpcStr,
 889    IntOp, IMMType, Pred> ;
 890 }
 891
 892 // has 3 operands
 893 multiclass F_ATOMIC_3_imp<NVPTXRegClass ptrclass, NVPTXRegClass regclass,
 894   string SpaceStr, string TypeStr, string OpcStr, PatFrag IntOp,
 895   Operand IMMType, Predicate Pred> {
 896   def reg : NVPTXInst<(outs regclass:$dst),
 897     (ins ptrclass:$addr, regclass:$b, regclass:$c),
 898                !strconcat("atom",
 899          !strconcat(SpaceStr,
 900          !strconcat(OpcStr,
 901          !strconcat(TypeStr,
 902          !strconcat(" \t$dst, [$addr], $b, $c;", ""))))),
 903          [(set regclass:$dst,
 904            (IntOp ptrclass:$addr, regclass:$b, regclass:$c))]>,
 905          Requires<[Pred]>;
 906   def imm1 : NVPTXInst<(outs regclass:$dst),
 907     (ins ptrclass:$addr, IMMType:$b, regclass:$c),
 908                !strconcat("atom",
 909          !strconcat(SpaceStr,
 910          !strconcat(OpcStr,
 911          !strconcat(TypeStr,
 912          !strconcat(" \t$dst, [$addr], $b, $c;", ""))))),
 913          [(set regclass:$dst, (IntOp ptrclass:$addr, imm:$b, regclass:$c))]>,
 914   Requires<[Pred]>;
 915   def imm2 : NVPTXInst<(outs regclass:$dst),
 916     (ins ptrclass:$addr, regclass:$b, IMMType:$c),
 917                !strconcat("atom",
 918          !strconcat(SpaceStr,
 919          !strconcat(OpcStr,
 920          !strconcat(TypeStr,
 921          !strconcat(" \t$dst, [$addr], $b, $c;", ""))))),
 922          [(set regclass:$dst, (IntOp ptrclass:$addr, regclass:$b, imm:$c))]>,
 923   Requires<[Pred]>;
 924   def imm3 : NVPTXInst<(outs regclass:$dst),
 925     (ins ptrclass:$addr, IMMType:$b, IMMType:$c),
 926                !strconcat("atom",
 927          !strconcat(SpaceStr,
 928          !strconcat(OpcStr,
 929          !strconcat(TypeStr,
 930          !strconcat(" \t$dst, [$addr], $b, $c;", ""))))),
 931          [(set regclass:$dst, (IntOp ptrclass:$addr, imm:$b, imm:$c))]>,
 932   Requires<[Pred]>;
 933 }
 934 multiclass F_ATOMIC_3<NVPTXRegClass regclass, string SpaceStr, string TypeStr,
 935   string OpcStr, PatFrag IntOp, Operand IMMType, Predicate Pred> {
 936   defm p32 : F_ATOMIC_3_imp<Int32Regs, regclass, SpaceStr, TypeStr, OpcStr,
 937     IntOp, IMMType, Pred>;
 938   defm p64 : F_ATOMIC_3_imp<Int64Regs, regclass, SpaceStr, TypeStr, OpcStr,
 939     IntOp, IMMType, Pred>;
 940 }
 941
 942 // atom_add
 943
 944 def atomic_load_add_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
 945   (atomic_load_add_32 node:$a, node:$b)>;
 946 def atomic_load_add_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
 947   (atomic_load_add_32 node:$a, node:$b)>;
 948 def atomic_load_add_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
 949   (atomic_load_add_32 node:$a, node:$b)>;
 950 def atomic_load_add_64_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
 951   (atomic_load_add_64 node:$a, node:$b)>;
 952 def atomic_load_add_64_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
 953   (atomic_load_add_64 node:$a, node:$b)>;
 954 def atomic_load_add_64_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
 955   (atomic_load_add_64 node:$a, node:$b)>;
 956 def atomic_load_add_f32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
 957   (int_nvvm_atomic_load_add_f32 node:$a, node:$b)>;
 958 def atomic_load_add_f32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
 959   (int_nvvm_atomic_load_add_f32 node:$a, node:$b)>;
 960 def atomic_load_add_f32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
 961   (int_nvvm_atomic_load_add_f32 node:$a, node:$b)>;
 962
 963 defm INT_PTX_ATOM_ADD_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".u32", ".add",
 964   atomic_load_add_32_g, i32imm, imm, hasAtomRedG32>;
 965 defm INT_PTX_ATOM_ADD_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".u32", ".add",
 966   atomic_load_add_32_s, i32imm, imm, hasAtomRedS32>;
 967 defm INT_PTX_ATOM_ADD_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".u32", ".add",
 968   atomic_load_add_32_gen, i32imm, imm, hasAtomRedGen32>;
 969 defm INT_PTX_ATOM_ADD_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".u32",
 970   ".add", atomic_load_add_32_gen, i32imm, imm, useAtomRedG32forGen32>;
 971
 972 defm INT_PTX_ATOM_ADD_G_64 : F_ATOMIC_2<Int64Regs, ".global", ".u64", ".add",
 973   atomic_load_add_64_g, i64imm, imm, hasAtomRedG64>;
 974 defm INT_PTX_ATOM_ADD_S_64 : F_ATOMIC_2<Int64Regs, ".shared", ".u64", ".add",
 975   atomic_load_add_64_s, i64imm, imm, hasAtomRedS64>;
 976 defm INT_PTX_ATOM_ADD_GEN_64 : F_ATOMIC_2<Int64Regs, "", ".u64", ".add",
 977   atomic_load_add_64_gen, i64imm, imm, hasAtomRedGen64>;
 978 defm INT_PTX_ATOM_ADD_GEN_64_USE_G : F_ATOMIC_2<Int64Regs, ".global", ".u64",
 979   ".add", atomic_load_add_64_gen, i64imm, imm, useAtomRedG64forGen64>;
 980
 981 defm INT_PTX_ATOM_ADD_G_F32 : F_ATOMIC_2<Float32Regs, ".global", ".f32", ".add",
 982   atomic_load_add_f32_g, f32imm, fpimm, hasAtomAddF32>;
 983 defm INT_PTX_ATOM_ADD_S_F32 : F_ATOMIC_2<Float32Regs, ".shared", ".f32", ".add",
 984   atomic_load_add_f32_s, f32imm, fpimm, hasAtomAddF32>;
 985 defm INT_PTX_ATOM_ADD_GEN_F32 : F_ATOMIC_2<Float32Regs, "", ".f32", ".add",
 986   atomic_load_add_f32_gen, f32imm, fpimm, hasAtomAddF32>;
 987
 988 // atom_sub
 989
 990 def atomic_load_sub_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
 991   (atomic_load_sub_32 node:$a, node:$b)>;
 992 def atomic_load_sub_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
 993   (atomic_load_sub_32 node:$a, node:$b)>;
 994 def atomic_load_sub_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
 995   (atomic_load_sub_32 node:$a, node:$b)>;
 996 def atomic_load_sub_64_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
 997   (atomic_load_sub_64 node:$a, node:$b)>;
 998 def atomic_load_sub_64_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
 999   (atomic_load_sub_64 node:$a, node:$b)>;
1000 def atomic_load_sub_64_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1001   (atomic_load_sub_64 node:$a, node:$b)>;
1002
1003 defm INT_PTX_ATOM_SUB_G_32 : F_ATOMIC_2_NEG<Int32Regs, ".global", "32", ".add",
1004   atomic_load_sub_32_g, i32imm, hasAtomRedG32>;
1005 defm INT_PTX_ATOM_SUB_G_64 : F_ATOMIC_2_NEG<Int64Regs, ".global", "64", ".add",
1006   atomic_load_sub_64_g, i64imm, hasAtomRedG64>;
1007 defm INT_PTX_ATOM_SUB_GEN_32 : F_ATOMIC_2_NEG<Int32Regs, "", "32", ".add",
1008   atomic_load_sub_32_gen, i32imm, hasAtomRedGen32>;
1009 defm INT_PTX_ATOM_SUB_GEN_32_USE_G : F_ATOMIC_2_NEG<Int32Regs, ".global", "32",
1010   ".add", atomic_load_sub_32_gen, i32imm, useAtomRedG32forGen32>;
1011 defm INT_PTX_ATOM_SUB_S_32 : F_ATOMIC_2_NEG<Int32Regs, ".shared", "32", ".add",
1012   atomic_load_sub_32_s, i32imm, hasAtomRedS32>;
1013 defm INT_PTX_ATOM_SUB_S_64 : F_ATOMIC_2_NEG<Int64Regs, ".shared", "64", ".add",
1014   atomic_load_sub_64_s, i64imm, hasAtomRedS64>;
1015 defm INT_PTX_ATOM_SUB_GEN_64 : F_ATOMIC_2_NEG<Int64Regs, "", "64", ".add",
1016   atomic_load_sub_64_gen, i64imm, hasAtomRedGen64>;
1017 defm INT_PTX_ATOM_SUB_GEN_64_USE_G : F_ATOMIC_2_NEG<Int64Regs, ".global", "64",
1018   ".add", atomic_load_sub_64_gen, i64imm, useAtomRedG64forGen64>;
1019
1020 // atom_swap
1021
1022 def atomic_swap_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1023   (atomic_swap_32 node:$a, node:$b)>;
1024 def atomic_swap_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1025   (atomic_swap_32 node:$a, node:$b)>;
1026 def atomic_swap_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1027   (atomic_swap_32 node:$a, node:$b)>;
1028 def atomic_swap_64_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1029   (atomic_swap_64 node:$a, node:$b)>;
1030 def atomic_swap_64_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1031   (atomic_swap_64 node:$a, node:$b)>;
1032 def atomic_swap_64_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1033   (atomic_swap_64 node:$a, node:$b)>;
1034
1035 defm INT_PTX_ATOM_SWAP_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".b32", ".exch",
1036   atomic_swap_32_g, i32imm, imm, hasAtomRedG32>;
1037 defm INT_PTX_ATOM_SWAP_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".b32", ".exch",
1038   atomic_swap_32_s, i32imm, imm, hasAtomRedS32>;
1039 defm INT_PTX_ATOM_SWAP_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".b32", ".exch",
1040   atomic_swap_32_gen, i32imm, imm, hasAtomRedGen32>;
1041 defm INT_PTX_ATOM_SWAP_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".b32",
1042   ".exch", atomic_swap_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1043 defm INT_PTX_ATOM_SWAP_G_64 : F_ATOMIC_2<Int64Regs, ".global", ".b64", ".exch",
1044   atomic_swap_64_g, i64imm, imm, hasAtomRedG64>;
1045 defm INT_PTX_ATOM_SWAP_S_64 : F_ATOMIC_2<Int64Regs, ".shared", ".b64", ".exch",
1046   atomic_swap_64_s, i64imm, imm, hasAtomRedS64>;
1047 defm INT_PTX_ATOM_SWAP_GEN_64 : F_ATOMIC_2<Int64Regs, "", ".b64", ".exch",
1048   atomic_swap_64_gen, i64imm, imm, hasAtomRedGen64>;
1049 defm INT_PTX_ATOM_SWAP_GEN_64_USE_G : F_ATOMIC_2<Int64Regs, ".global", ".b64",
1050   ".exch", atomic_swap_64_gen, i64imm, imm, useAtomRedG64forGen64>;
1051
1052 // atom_max
1053
1054 def atomic_load_max_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b)
1055   , (atomic_load_max_32 node:$a, node:$b)>;
1056 def atomic_load_max_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1057   (atomic_load_max_32 node:$a, node:$b)>;
1058 def atomic_load_max_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1059   (atomic_load_max_32 node:$a, node:$b)>;
1060 def atomic_load_umax_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1061   (atomic_load_umax_32 node:$a, node:$b)>;
1062 def atomic_load_umax_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1063   (atomic_load_umax_32 node:$a, node:$b)>;
1064 def atomic_load_umax_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1065   (atomic_load_umax_32 node:$a, node:$b)>;
1066
1067 defm INT_PTX_ATOM_LOAD_MAX_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".s32",
1068   ".max", atomic_load_max_32_g, i32imm, imm, hasAtomRedG32>;
1069 defm INT_PTX_ATOM_LOAD_MAX_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".s32",
1070   ".max", atomic_load_max_32_s, i32imm, imm, hasAtomRedS32>;
1071 defm INT_PTX_ATOM_LOAD_MAX_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".s32", ".max",
1072   atomic_load_max_32_gen, i32imm, imm, hasAtomRedGen32>;
1073 defm INT_PTX_ATOM_LOAD_MAX_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global",
1074   ".s32", ".max", atomic_load_max_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1075 defm INT_PTX_ATOM_LOAD_UMAX_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".u32",
1076   ".max", atomic_load_umax_32_g, i32imm, imm, hasAtomRedG32>;
1077 defm INT_PTX_ATOM_LOAD_UMAX_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".u32",
1078   ".max", atomic_load_umax_32_s, i32imm, imm, hasAtomRedS32>;
1079 defm INT_PTX_ATOM_LOAD_UMAX_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".u32", ".max",
1080   atomic_load_umax_32_gen, i32imm, imm, hasAtomRedGen32>;
1081 defm INT_PTX_ATOM_LOAD_UMAX_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global",
1082   ".u32", ".max", atomic_load_umax_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1083
1084 // atom_min
1085
1086 def atomic_load_min_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1087   (atomic_load_min_32 node:$a, node:$b)>;
1088 def atomic_load_min_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1089   (atomic_load_min_32 node:$a, node:$b)>;
1090 def atomic_load_min_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1091   (atomic_load_min_32 node:$a, node:$b)>;
1092 def atomic_load_umin_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1093   (atomic_load_umin_32 node:$a, node:$b)>;
1094 def atomic_load_umin_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1095   (atomic_load_umin_32 node:$a, node:$b)>;
1096 def atomic_load_umin_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1097   (atomic_load_umin_32 node:$a, node:$b)>;
1098
1099 defm INT_PTX_ATOM_LOAD_MIN_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".s32",
1100   ".min", atomic_load_min_32_g, i32imm, imm, hasAtomRedG32>;
1101 defm INT_PTX_ATOM_LOAD_MIN_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".s32",
1102   ".min", atomic_load_min_32_s, i32imm, imm, hasAtomRedS32>;
1103 defm INT_PTX_ATOM_LOAD_MIN_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".s32", ".min",
1104   atomic_load_min_32_gen, i32imm, imm, hasAtomRedGen32>;
1105 defm INT_PTX_ATOM_LOAD_MIN_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global",
1106   ".s32", ".min", atomic_load_min_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1107 defm INT_PTX_ATOM_LOAD_UMIN_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".u32",
1108   ".min", atomic_load_umin_32_g, i32imm, imm, hasAtomRedG32>;
1109 defm INT_PTX_ATOM_LOAD_UMIN_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".u32",
1110   ".min", atomic_load_umin_32_s, i32imm, imm, hasAtomRedS32>;
1111 defm INT_PTX_ATOM_LOAD_UMIN_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".u32", ".min",
1112   atomic_load_umin_32_gen, i32imm, imm, hasAtomRedGen32>;
1113 defm INT_PTX_ATOM_LOAD_UMIN_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global",
1114   ".u32", ".min", atomic_load_umin_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1115
1116 // atom_inc  atom_dec
1117
1118 def atomic_load_inc_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1119   (int_nvvm_atomic_load_inc_32 node:$a, node:$b)>;
1120 def atomic_load_inc_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1121   (int_nvvm_atomic_load_inc_32 node:$a, node:$b)>;
1122 def atomic_load_inc_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1123   (int_nvvm_atomic_load_inc_32 node:$a, node:$b)>;
1124 def atomic_load_dec_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1125   (int_nvvm_atomic_load_dec_32 node:$a, node:$b)>;
1126 def atomic_load_dec_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1127   (int_nvvm_atomic_load_dec_32 node:$a, node:$b)>;
1128 def atomic_load_dec_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1129   (int_nvvm_atomic_load_dec_32 node:$a, node:$b)>;
1130
1131 defm INT_PTX_ATOM_INC_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".u32", ".inc",
1132   atomic_load_inc_32_g, i32imm, imm, hasAtomRedG32>;
1133 defm INT_PTX_ATOM_INC_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".u32", ".inc",
1134   atomic_load_inc_32_s, i32imm, imm, hasAtomRedS32>;
1135 defm INT_PTX_ATOM_INC_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".u32", ".inc",
1136   atomic_load_inc_32_gen, i32imm, imm, hasAtomRedGen32>;
1137 defm INT_PTX_ATOM_INC_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".u32",
1138   ".inc", atomic_load_inc_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1139 defm INT_PTX_ATOM_DEC_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".u32", ".dec",
1140   atomic_load_dec_32_g, i32imm, imm, hasAtomRedG32>;
1141 defm INT_PTX_ATOM_DEC_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".u32", ".dec",
1142   atomic_load_dec_32_s, i32imm, imm, hasAtomRedS32>;
1143 defm INT_PTX_ATOM_DEC_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".u32", ".dec",
1144   atomic_load_dec_32_gen, i32imm, imm, hasAtomRedGen32>;
1145 defm INT_PTX_ATOM_DEC_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".u32",
1146   ".dec", atomic_load_dec_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1147
1148 // atom_and
1149
1150 def atomic_load_and_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1151   (atomic_load_and_32 node:$a, node:$b)>;
1152 def atomic_load_and_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1153   (atomic_load_and_32 node:$a, node:$b)>;
1154 def atomic_load_and_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1155   (atomic_load_and_32 node:$a, node:$b)>;
1156
1157 defm INT_PTX_ATOM_AND_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".b32", ".and",
1158   atomic_load_and_32_g, i32imm, imm, hasAtomRedG32>;
1159 defm INT_PTX_ATOM_AND_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".b32", ".and",
1160   atomic_load_and_32_s, i32imm, imm, hasAtomRedS32>;
1161 defm INT_PTX_ATOM_AND_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".b32", ".and",
1162   atomic_load_and_32_gen, i32imm, imm, hasAtomRedGen32>;
1163 defm INT_PTX_ATOM_AND_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".b32",
1164   ".and", atomic_load_and_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1165
1166 // atom_or
1167
1168 def atomic_load_or_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1169   (atomic_load_or_32 node:$a, node:$b)>;
1170 def atomic_load_or_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1171   (atomic_load_or_32 node:$a, node:$b)>;
1172 def atomic_load_or_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1173   (atomic_load_or_32 node:$a, node:$b)>;
1174
1175 defm INT_PTX_ATOM_OR_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".b32", ".or",
1176   atomic_load_or_32_g, i32imm, imm, hasAtomRedG32>;
1177 defm INT_PTX_ATOM_OR_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".b32", ".or",
1178   atomic_load_or_32_gen, i32imm, imm, hasAtomRedGen32>;
1179 defm INT_PTX_ATOM_OR_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".b32",
1180   ".or", atomic_load_or_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1181 defm INT_PTX_ATOM_OR_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".b32", ".or",
1182   atomic_load_or_32_s, i32imm, imm, hasAtomRedS32>;
1183
1184 // atom_xor
1185
1186 def atomic_load_xor_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1187   (atomic_load_xor_32 node:$a, node:$b)>;
1188 def atomic_load_xor_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1189   (atomic_load_xor_32 node:$a, node:$b)>;
1190 def atomic_load_xor_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1191   (atomic_load_xor_32 node:$a, node:$b)>;
1192
1193 defm INT_PTX_ATOM_XOR_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".b32", ".xor",
1194   atomic_load_xor_32_g, i32imm, imm, hasAtomRedG32>;
1195 defm INT_PTX_ATOM_XOR_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".b32", ".xor",
1196   atomic_load_xor_32_s, i32imm, imm, hasAtomRedS32>;
1197 defm INT_PTX_ATOM_XOR_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".b32", ".xor",
1198   atomic_load_xor_32_gen, i32imm, imm, hasAtomRedGen32>;
1199 defm INT_PTX_ATOM_XOR_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".b32",
1200   ".xor", atomic_load_xor_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1201
1202 // atom_cas
1203
1204 def atomic_cmp_swap_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b, node:$c),
1205   (atomic_cmp_swap_32 node:$a, node:$b, node:$c)>;
1206 def atomic_cmp_swap_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b, node:$c),
1207   (atomic_cmp_swap_32 node:$a, node:$b, node:$c)>;
1208 def atomic_cmp_swap_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b, node:$c),
1209   (atomic_cmp_swap_32 node:$a, node:$b, node:$c)>;
1210 def atomic_cmp_swap_64_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b, node:$c),
1211   (atomic_cmp_swap_64 node:$a, node:$b, node:$c)>;
1212 def atomic_cmp_swap_64_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b, node:$c),
1213   (atomic_cmp_swap_64 node:$a, node:$b, node:$c)>;
1214 def atomic_cmp_swap_64_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b, node:$c),
1215   (atomic_cmp_swap_64 node:$a, node:$b, node:$c)>;
1216
1217 defm INT_PTX_ATOM_CAS_G_32 : F_ATOMIC_3<Int32Regs, ".global", ".b32", ".cas",
1218   atomic_cmp_swap_32_g, i32imm, hasAtomRedG32>;
1219 defm INT_PTX_ATOM_CAS_S_32 : F_ATOMIC_3<Int32Regs, ".shared", ".b32", ".cas",
1220   atomic_cmp_swap_32_s, i32imm, hasAtomRedS32>;
1221 defm INT_PTX_ATOM_CAS_GEN_32 : F_ATOMIC_3<Int32Regs, "", ".b32", ".cas",
1222   atomic_cmp_swap_32_gen, i32imm, hasAtomRedGen32>;
1223 defm INT_PTX_ATOM_CAS_GEN_32_USE_G : F_ATOMIC_3<Int32Regs, ".global", ".b32",
1224   ".cas", atomic_cmp_swap_32_gen, i32imm, useAtomRedG32forGen32>;
1225 defm INT_PTX_ATOM_CAS_G_64 : F_ATOMIC_3<Int64Regs, ".global", ".b64", ".cas",
1226   atomic_cmp_swap_64_g, i64imm, hasAtomRedG64>;
1227 defm INT_PTX_ATOM_CAS_S_64 : F_ATOMIC_3<Int64Regs, ".shared", ".b64", ".cas",
1228   atomic_cmp_swap_64_s, i64imm, hasAtomRedS64>;
1229 defm INT_PTX_ATOM_CAS_GEN_64 : F_ATOMIC_3<Int64Regs, "", ".b64", ".cas",
1230   atomic_cmp_swap_64_gen, i64imm, hasAtomRedGen64>;
1231 defm INT_PTX_ATOM_CAS_GEN_64_USE_G : F_ATOMIC_3<Int64Regs, ".global", ".b64",
1232   ".cas", atomic_cmp_swap_64_gen, i64imm, useAtomRedG64forGen64>;
1233
1234
1235 //-----------------------------------
1236 // Read Special Registers
1237 //-----------------------------------
1238 class F_SREG<string OpStr, NVPTXRegClass regclassOut, Intrinsic IntOp> :
1239       NVPTXInst<(outs regclassOut:$dst), (ins),
1240                OpStr,
1241          [(set regclassOut:$dst, (IntOp))]>;
1242
1243 def INT_PTX_SREG_TID_X : F_SREG<"mov.u32 \t$dst, %tid.x;", Int32Regs,
1244   int_nvvm_read_ptx_sreg_tid_x>;
1245 def INT_PTX_SREG_TID_Y : F_SREG<"mov.u32 \t$dst, %tid.y;", Int32Regs,
1246   int_nvvm_read_ptx_sreg_tid_y>;
1247 def INT_PTX_SREG_TID_Z : F_SREG<"mov.u32 \t$dst, %tid.z;", Int32Regs,
1248   int_nvvm_read_ptx_sreg_tid_z>;
1249
1250 def INT_PTX_SREG_NTID_X : F_SREG<"mov.u32 \t$dst, %ntid.x;", Int32Regs,
1251   int_nvvm_read_ptx_sreg_ntid_x>;
1252 def INT_PTX_SREG_NTID_Y : F_SREG<"mov.u32 \t$dst, %ntid.y;", Int32Regs,
1253   int_nvvm_read_ptx_sreg_ntid_y>;
1254 def INT_PTX_SREG_NTID_Z : F_SREG<"mov.u32 \t$dst, %ntid.z;", Int32Regs,
1255   int_nvvm_read_ptx_sreg_ntid_z>;
1256
1257 def INT_PTX_SREG_CTAID_X : F_SREG<"mov.u32 \t$dst, %ctaid.x;", Int32Regs,
1258   int_nvvm_read_ptx_sreg_ctaid_x>;
1259 def INT_PTX_SREG_CTAID_Y : F_SREG<"mov.u32 \t$dst, %ctaid.y;", Int32Regs,
1260   int_nvvm_read_ptx_sreg_ctaid_y>;
1261 def INT_PTX_SREG_CTAID_Z : F_SREG<"mov.u32 \t$dst, %ctaid.z;", Int32Regs,
1262   int_nvvm_read_ptx_sreg_ctaid_z>;
1263
1264 def INT_PTX_SREG_NCTAID_X : F_SREG<"mov.u32 \t$dst, %nctaid.x;", Int32Regs,
1265   int_nvvm_read_ptx_sreg_nctaid_x>;
1266 def INT_PTX_SREG_NCTAID_Y : F_SREG<"mov.u32 \t$dst, %nctaid.y;", Int32Regs,
1267   int_nvvm_read_ptx_sreg_nctaid_y>;
1268 def INT_PTX_SREG_NCTAID_Z : F_SREG<"mov.u32 \t$dst, %nctaid.z;", Int32Regs,
1269   int_nvvm_read_ptx_sreg_nctaid_z>;
1270
1271 def INT_PTX_SREG_WARPSIZE : F_SREG<"mov.u32 \t$dst, WARP_SZ;", Int32Regs,
1272   int_nvvm_read_ptx_sreg_warpsize>;
1273
1274
1275 //-----------------------------------
1276 // Support for ldu on sm_20 or later
1277 //-----------------------------------
1278
1279 def ldu_i8 : PatFrag<(ops node:$ptr), (int_nvvm_ldu_global_i node:$ptr), [{
1280   MemIntrinsicSDNode *M = cast<MemIntrinsicSDNode>(N);
1281   return M->getMemoryVT() == MVT::i8;
1282 }]>;
1283
1284 // Scalar
1285 // @TODO: Revisit this, Changed imemAny to imem
1286 multiclass LDU_G<string TyStr, NVPTXRegClass regclass, Intrinsic IntOp> {
1287   def areg: NVPTXInst<(outs regclass:$result), (ins Int32Regs:$src),
1288                !strconcat("ldu.global.", TyStr),
1289          [(set regclass:$result, (IntOp Int32Regs:$src))]>, Requires<[hasLDU]>;
1290   def areg64: NVPTXInst<(outs regclass:$result), (ins Int64Regs:$src),
1291                !strconcat("ldu.global.", TyStr),
1292          [(set regclass:$result, (IntOp Int64Regs:$src))]>, Requires<[hasLDU]>;
1293  def avar:  NVPTXInst<(outs regclass:$result), (ins imem:$src),
1294                !strconcat("ldu.global.", TyStr),
1295                 [(set regclass:$result, (IntOp (Wrapper tglobaladdr:$src)))]>,
1296                 Requires<[hasLDU]>;
1297  def ari :  NVPTXInst<(outs regclass:$result), (ins MEMri:$src),
1298                !strconcat("ldu.global.", TyStr),
1299          [(set regclass:$result, (IntOp ADDRri:$src))]>, Requires<[hasLDU]>;
1300  def ari64 :  NVPTXInst<(outs regclass:$result), (ins MEMri64:$src),
1301                !strconcat("ldu.global.", TyStr),
1302          [(set regclass:$result, (IntOp ADDRri64:$src))]>, Requires<[hasLDU]>;
1303 }
1304
1305 multiclass LDU_G_NOINTRIN<string TyStr, NVPTXRegClass regclass, PatFrag IntOp> {
1306   def areg: NVPTXInst<(outs regclass:$result), (ins Int32Regs:$src),
1307                !strconcat("ldu.global.", TyStr),
1308          [(set regclass:$result, (IntOp Int32Regs:$src))]>, Requires<[hasLDU]>;
1309   def areg64: NVPTXInst<(outs regclass:$result), (ins Int64Regs:$src),
1310                !strconcat("ldu.global.", TyStr),
1311          [(set regclass:$result, (IntOp Int64Regs:$src))]>, Requires<[hasLDU]>;
1312  def avar:  NVPTXInst<(outs regclass:$result), (ins imem:$src),
1313                !strconcat("ldu.global.", TyStr),
1314          [(set regclass:$result, (IntOp (Wrapper tglobaladdr:$src)))]>,
1315          Requires<[hasLDU]>;
1316  def ari :  NVPTXInst<(outs regclass:$result), (ins MEMri:$src),
1317                !strconcat("ldu.global.", TyStr),
1318          [(set regclass:$result, (IntOp ADDRri:$src))]>, Requires<[hasLDU]>;
1319  def ari64 :  NVPTXInst<(outs regclass:$result), (ins MEMri64:$src),
1320                !strconcat("ldu.global.", TyStr),
1321          [(set regclass:$result, (IntOp ADDRri64:$src))]>, Requires<[hasLDU]>;
1322 }
1323
1324 defm INT_PTX_LDU_GLOBAL_i8  : LDU_G_NOINTRIN<"u8 \t$result, [$src];", Int16Regs,
1325                                              ldu_i8>;
1326 defm INT_PTX_LDU_GLOBAL_i16 : LDU_G<"u16 \t$result, [$src];", Int16Regs,
1327 int_nvvm_ldu_global_i>;
1328 defm INT_PTX_LDU_GLOBAL_i32 : LDU_G<"u32 \t$result, [$src];", Int32Regs,
1329 int_nvvm_ldu_global_i>;
1330 defm INT_PTX_LDU_GLOBAL_i64 : LDU_G<"u64 \t$result, [$src];", Int64Regs,
1331 int_nvvm_ldu_global_i>;
1332 defm INT_PTX_LDU_GLOBAL_f32 : LDU_G<"f32 \t$result, [$src];", Float32Regs,
1333 int_nvvm_ldu_global_f>;
1334 defm INT_PTX_LDU_GLOBAL_f64 : LDU_G<"f64 \t$result, [$src];", Float64Regs,
1335 int_nvvm_ldu_global_f>;
1336 defm INT_PTX_LDU_GLOBAL_p32 : LDU_G<"u32 \t$result, [$src];", Int32Regs,
1337 int_nvvm_ldu_global_p>;
1338 defm INT_PTX_LDU_GLOBAL_p64 : LDU_G<"u64 \t$result, [$src];", Int64Regs,
1339 int_nvvm_ldu_global_p>;
1340
1341 // vector
1342
1343 // Elementized vector ldu
1344 multiclass VLDU_G_ELE_V2<string TyStr, NVPTXRegClass regclass> {
1345  def _areg32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
1346                      (ins Int32Regs:$src),
1347                      !strconcat("ldu.global.", TyStr), []>;
1348  def _areg64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
1349                      (ins Int64Regs:$src),
1350                      !strconcat("ldu.global.", TyStr), []>;
1351  def _ari32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
1352                      (ins MEMri:$src),
1353                      !strconcat("ldu.global.", TyStr), []>;
1354  def _ari64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
1355                      (ins MEMri64:$src),
1356                      !strconcat("ldu.global.", TyStr), []>;
1357  def _avar: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
1358                      (ins imemAny:$src),
1359                      !strconcat("ldu.global.", TyStr), []>;
1360 }
1361
1362 multiclass VLDU_G_ELE_V4<string TyStr, NVPTXRegClass regclass> {
1363  def _areg32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
1364                             regclass:$dst4), (ins Int32Regs:$src),
1365                !strconcat("ldu.global.", TyStr), []>;
1366  def _areg64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
1367                             regclass:$dst4), (ins Int64Regs:$src),
1368                !strconcat("ldu.global.", TyStr), []>;
1369  def _ari32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
1370                             regclass:$dst4), (ins MEMri:$src),
1371                !strconcat("ldu.global.", TyStr), []>;
1372  def _ari64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
1373                             regclass:$dst4), (ins MEMri64:$src),
1374                !strconcat("ldu.global.", TyStr), []>;
1375  def _avar: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
1376                             regclass:$dst4), (ins imemAny:$src),
1377                !strconcat("ldu.global.", TyStr), []>;
1378 }
1379
1380 defm INT_PTX_LDU_G_v2i8_ELE
1381   : VLDU_G_ELE_V2<"v2.u8 \t{{$dst1, $dst2}}, [$src];",  Int16Regs>;
1382 defm INT_PTX_LDU_G_v2i16_ELE
1383   : VLDU_G_ELE_V2<"v2.u16 \t{{$dst1, $dst2}}, [$src];", Int16Regs>;
1384 defm INT_PTX_LDU_G_v2i32_ELE
1385   : VLDU_G_ELE_V2<"v2.u32 \t{{$dst1, $dst2}}, [$src];", Int32Regs>;
1386 defm INT_PTX_LDU_G_v2f32_ELE
1387   : VLDU_G_ELE_V2<"v2.f32 \t{{$dst1, $dst2}}, [$src];", Float32Regs>;
1388 defm INT_PTX_LDU_G_v2i64_ELE
1389   : VLDU_G_ELE_V2<"v2.u64 \t{{$dst1, $dst2}}, [$src];", Int64Regs>;
1390 defm INT_PTX_LDU_G_v2f64_ELE
1391   : VLDU_G_ELE_V2<"v2.f64 \t{{$dst1, $dst2}}, [$src];", Float64Regs>;
1392 defm INT_PTX_LDU_G_v4i8_ELE
1393   : VLDU_G_ELE_V4<"v4.u8 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Int16Regs>;
1394 defm INT_PTX_LDU_G_v4i16_ELE
1395   : VLDU_G_ELE_V4<"v4.u16 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];",
1396     Int16Regs>;
1397 defm INT_PTX_LDU_G_v4i32_ELE
1398   : VLDU_G_ELE_V4<"v4.u32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];",
1399     Int32Regs>;
1400 defm INT_PTX_LDU_G_v4f32_ELE
1401   : VLDU_G_ELE_V4<"v4.f32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];",
1402     Float32Regs>;
1403
1404
1405 //-----------------------------------
1406 // Support for ldg on sm_35 or later
1407 //-----------------------------------
1408
1409 def ldg_i8 : PatFrag<(ops node:$ptr), (int_nvvm_ldg_global_i node:$ptr), [{
1410   MemIntrinsicSDNode *M = cast<MemIntrinsicSDNode>(N);
1411   return M->getMemoryVT() == MVT::i8;
1412 }]>;
1413
1414 multiclass LDG_G<string TyStr, NVPTXRegClass regclass, Intrinsic IntOp> {
1415   def areg: NVPTXInst<(outs regclass:$result), (ins Int32Regs:$src),
1416                !strconcat("ld.global.nc.", TyStr),
1417          [(set regclass:$result, (IntOp Int32Regs:$src))]>, Requires<[hasLDG]>;
1418   def areg64: NVPTXInst<(outs regclass:$result), (ins Int64Regs:$src),
1419                !strconcat("ld.global.nc.", TyStr),
1420          [(set regclass:$result, (IntOp Int64Regs:$src))]>, Requires<[hasLDG]>;
1421  def avar:  NVPTXInst<(outs regclass:$result), (ins imem:$src),
1422                !strconcat("ld.global.nc.", TyStr),
1423          [(set regclass:$result, (IntOp (Wrapper tglobaladdr:$src)))]>,
1424          Requires<[hasLDG]>;
1425  def ari :  NVPTXInst<(outs regclass:$result), (ins MEMri:$src),
1426                !strconcat("ld.global.nc.", TyStr),
1427          [(set regclass:$result, (IntOp ADDRri:$src))]>, Requires<[hasLDG]>;
1428  def ari64 :  NVPTXInst<(outs regclass:$result), (ins MEMri64:$src),
1429                !strconcat("ld.global.nc.", TyStr),
1430          [(set regclass:$result, (IntOp ADDRri64:$src))]>, Requires<[hasLDG]>;
1431 }
1432
1433 multiclass LDG_G_NOINTRIN<string TyStr, NVPTXRegClass regclass, PatFrag IntOp> {
1434   def areg: NVPTXInst<(outs regclass:$result), (ins Int32Regs:$src),
1435                !strconcat("ld.global.nc.", TyStr),
1436          [(set regclass:$result, (IntOp Int32Regs:$src))]>, Requires<[hasLDG]>;
1437   def areg64: NVPTXInst<(outs regclass:$result), (ins Int64Regs:$src),
1438                !strconcat("ld.global.nc.", TyStr),
1439          [(set regclass:$result, (IntOp Int64Regs:$src))]>, Requires<[hasLDG]>;
1440  def avar:  NVPTXInst<(outs regclass:$result), (ins imem:$src),
1441                !strconcat("ld.global.nc.", TyStr),
1442          [(set regclass:$result, (IntOp (Wrapper tglobaladdr:$src)))]>,
1443         Requires<[hasLDG]>;
1444  def ari :  NVPTXInst<(outs regclass:$result), (ins MEMri:$src),
1445                !strconcat("ld.global.nc.", TyStr),
1446          [(set regclass:$result, (IntOp ADDRri:$src))]>, Requires<[hasLDG]>;
1447  def ari64 :  NVPTXInst<(outs regclass:$result), (ins MEMri64:$src),
1448                !strconcat("ld.global.nc.", TyStr),
1449          [(set regclass:$result, (IntOp ADDRri64:$src))]>, Requires<[hasLDG]>;
1450 }
1451
1452 defm INT_PTX_LDG_GLOBAL_i8
1453   : LDG_G_NOINTRIN<"u8 \t$result, [$src];",  Int16Regs, ldg_i8>;
1454 defm INT_PTX_LDG_GLOBAL_i16
1455   : LDG_G<"u16 \t$result, [$src];", Int16Regs,   int_nvvm_ldg_global_i>;
1456 defm INT_PTX_LDG_GLOBAL_i32
1457   : LDG_G<"u32 \t$result, [$src];", Int32Regs,   int_nvvm_ldg_global_i>;
1458 defm INT_PTX_LDG_GLOBAL_i64
1459   : LDG_G<"u64 \t$result, [$src];", Int64Regs,   int_nvvm_ldg_global_i>;
1460 defm INT_PTX_LDG_GLOBAL_f32
1461   : LDG_G<"f32 \t$result, [$src];", Float32Regs, int_nvvm_ldg_global_f>;
1462 defm INT_PTX_LDG_GLOBAL_f64
1463   : LDG_G<"f64 \t$result, [$src];", Float64Regs, int_nvvm_ldg_global_f>;
1464 defm INT_PTX_LDG_GLOBAL_p32
1465   : LDG_G<"u32 \t$result, [$src];", Int32Regs,   int_nvvm_ldg_global_p>;
1466 defm INT_PTX_LDG_GLOBAL_p64
1467   : LDG_G<"u64 \t$result, [$src];", Int64Regs,   int_nvvm_ldg_global_p>;
1468
1469 // vector
1470
1471 // Elementized vector ldg
1472 multiclass VLDG_G_ELE_V2<string TyStr, NVPTXRegClass regclass> {
1473  def _areg32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
1474                      (ins Int32Regs:$src),
1475                      !strconcat("ld.global.nc.", TyStr), []>;
1476  def _areg64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
1477                      (ins Int64Regs:$src),
1478                      !strconcat("ld.global.nc.", TyStr), []>;
1479  def _ari32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
1480                      (ins MEMri:$src),
1481                      !strconcat("ld.global.nc.", TyStr), []>;
1482  def _ari64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
1483                      (ins MEMri64:$src),
1484                      !strconcat("ld.global.nc.", TyStr), []>;
1485  def _avar: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
1486                      (ins imemAny:$src),
1487                      !strconcat("ld.global.nc.", TyStr), []>;
1488 }
1489
1490 multiclass VLDG_G_ELE_V4<string TyStr, NVPTXRegClass regclass> {
1491   def _areg32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
1492                               regclass:$dst4), (ins Int32Regs:$src),
1493                !strconcat("ld.global.nc.", TyStr), []>;
1494   def _areg64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
1495                                regclass:$dst4), (ins Int64Regs:$src),
1496                !strconcat("ld.global.nc.", TyStr), []>;
1497   def _ari32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
1498                               regclass:$dst4), (ins MEMri:$src),
1499                !strconcat("ld.global.nc.", TyStr), []>;
1500   def _ari64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
1501                               regclass:$dst4), (ins MEMri64:$src),
1502                !strconcat("ld.global.nc.", TyStr), []>;
1503   def _avar: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
1504                              regclass:$dst4), (ins imemAny:$src),
1505                !strconcat("ld.global.nc.", TyStr), []>;
1506 }
1507
1508 // FIXME: 8-bit LDG should be fixed once LDG/LDU nodes are made into proper loads.
1509 defm INT_PTX_LDG_G_v2i8_ELE
1510   : VLDG_G_ELE_V2<"v2.u8 \t{{$dst1, $dst2}}, [$src];",  Int16Regs>;
1511 defm INT_PTX_LDG_G_v2i16_ELE
1512   : VLDG_G_ELE_V2<"v2.u16 \t{{$dst1, $dst2}}, [$src];", Int16Regs>;
1513 defm INT_PTX_LDG_G_v2i32_ELE
1514   : VLDG_G_ELE_V2<"v2.u32 \t{{$dst1, $dst2}}, [$src];", Int32Regs>;
1515 defm INT_PTX_LDG_G_v2f32_ELE
1516   : VLDG_G_ELE_V2<"v2.f32 \t{{$dst1, $dst2}}, [$src];", Float32Regs>;
1517 defm INT_PTX_LDG_G_v2i64_ELE
1518   : VLDG_G_ELE_V2<"v2.u64 \t{{$dst1, $dst2}}, [$src];", Int64Regs>;
1519 defm INT_PTX_LDG_G_v2f64_ELE
1520   : VLDG_G_ELE_V2<"v2.f64 \t{{$dst1, $dst2}}, [$src];", Float64Regs>;
1521 defm INT_PTX_LDG_G_v4i8_ELE
1522   : VLDG_G_ELE_V4<"v4.u8 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Int16Regs>;
1523 defm INT_PTX_LDG_G_v4i16_ELE
1524   : VLDG_G_ELE_V4<"v4.u16 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Int16Regs>;
1525 defm INT_PTX_LDG_G_v4i32_ELE
1526   : VLDG_G_ELE_V4<"v4.u32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Int32Regs>;
1527 defm INT_PTX_LDG_G_v4f32_ELE
1528   : VLDG_G_ELE_V4<"v4.f32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Float32Regs>;
1529
1530
1531 multiclass NG_TO_G<string Str, Intrinsic Intrin> {
1532    def _yes : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
1533           !strconcat("cvta.", !strconcat(Str, ".u32 \t$result, $src;")),
1534       [(set Int32Regs:$result, (Intrin Int32Regs:$src))]>,
1535    Requires<[hasGenericLdSt]>;
1536    def _yes_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
1537           !strconcat("cvta.", !strconcat(Str, ".u64 \t$result, $src;")),
1538       [(set Int64Regs:$result, (Intrin Int64Regs:$src))]>,
1539    Requires<[hasGenericLdSt]>;
1540
1541 // @TODO: Are these actually needed?  I believe global addresses will be copied
1542 // to register values anyway.
1543    /*def __addr_yes : NVPTXInst<(outs Int32Regs:$result), (ins imemAny:$src),
1544           !strconcat("cvta.", !strconcat(Str, ".u32 \t$result, $src;")),
1545       [(set Int32Regs:$result, (Intrin (Wrapper tglobaladdr:$src)))]>,
1546       Requires<[hasGenericLdSt]>;
1547    def __addr_yes_64 : NVPTXInst<(outs Int64Regs:$result), (ins imemAny:$src),
1548           !strconcat("cvta.", !strconcat(Str, ".u64 \t$result, $src;")),
1549       [(set Int64Regs:$result, (Intrin (Wrapper tglobaladdr:$src)))]>,
1550       Requires<[hasGenericLdSt]>;*/
1551
1552    def _no : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
1553           "mov.u32 \t$result, $src;",
1554       [(set Int32Regs:$result, (Intrin Int32Regs:$src))]>;
1555    def _no_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
1556           "mov.u64 \t$result, $src;",
1557       [(set Int64Regs:$result, (Intrin Int64Regs:$src))]>;
1558
1559 // @TODO: Are these actually needed?  I believe global addresses will be copied
1560 // to register values anyway.
1561    /*def _addr_no : NVPTXInst<(outs Int32Regs:$result), (ins imem:$src),
1562           "mov.u32 \t$result, $src;",
1563       [(set Int32Regs:$result, (Intrin (Wrapper tglobaladdr:$src)))]>;
1564    def _addr_no_64 : NVPTXInst<(outs Int64Regs:$result), (ins imem:$src),
1565           "mov.u64 \t$result, $src;",
1566       [(set Int64Regs:$result, (Intrin (Wrapper tglobaladdr:$src)))]>;*/
1567 }
1568
1569 multiclass G_TO_NG<string Str, Intrinsic Intrin> {
1570    def _yes : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
1571           !strconcat("cvta.to.", !strconcat(Str, ".u32 \t$result, $src;")),
1572       [(set Int32Regs:$result, (Intrin Int32Regs:$src))]>,
1573    Requires<[hasGenericLdSt]>;
1574    def _yes_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
1575           !strconcat("cvta.to.", !strconcat(Str, ".u64 \t$result, $src;")),
1576       [(set Int64Regs:$result, (Intrin Int64Regs:$src))]>,
1577    Requires<[hasGenericLdSt]>;
1578    def _no : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
1579           "mov.u32 \t$result, $src;",
1580       [(set Int32Regs:$result, (Intrin Int32Regs:$src))]>;
1581    def _no_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
1582           "mov.u64 \t$result, $src;",
1583       [(set Int64Regs:$result, (Intrin Int64Regs:$src))]>;
1584 }
1585
1586 defm cvta_local  : NG_TO_G<"local", int_nvvm_ptr_local_to_gen>;
1587 defm cvta_shared : NG_TO_G<"shared", int_nvvm_ptr_shared_to_gen>;
1588 defm cvta_global : NG_TO_G<"global", int_nvvm_ptr_global_to_gen>;
1589 defm cvta_const  : NG_TO_G<"const", int_nvvm_ptr_constant_to_gen>;
1590
1591 defm cvta_to_local   : G_TO_NG<"local", int_nvvm_ptr_gen_to_local>;
1592 defm cvta_to_shared : G_TO_NG<"shared", int_nvvm_ptr_gen_to_shared>;
1593 defm cvta_to_global : G_TO_NG<"global", int_nvvm_ptr_gen_to_global>;
1594 defm cvta_to_const  : G_TO_NG<"const", int_nvvm_ptr_gen_to_constant>;
1595
1596
1597 // nvvm.ptr.gen.to.param
1598 def nvvm_ptr_gen_to_param : NVPTXInst<(outs Int32Regs:$result),
1599   (ins Int32Regs:$src),
1600                         "mov.u32 \t$result, $src;",
1601                               [(set Int32Regs:$result,
1602                                 (int_nvvm_ptr_gen_to_param Int32Regs:$src))]>;
1603 def nvvm_ptr_gen_to_param_64 : NVPTXInst<(outs Int64Regs:$result),
1604   (ins Int64Regs:$src),
1605                         "mov.u64 \t$result, $src;",
1606                               [(set Int64Regs:$result,
1607                                 (int_nvvm_ptr_gen_to_param Int64Regs:$src))]>;
1608
1609
1610 // nvvm.move intrinsicc
1611 def nvvm_move_i16 : NVPTXInst<(outs Int16Regs:$r), (ins Int16Regs:$s),
1612                              "mov.b16 \t$r, $s;",
1613                              [(set Int16Regs:$r,
1614                                (int_nvvm_move_i16 Int16Regs:$s))]>;
1615 def nvvm_move_i32 : NVPTXInst<(outs Int32Regs:$r), (ins Int32Regs:$s),
1616                              "mov.b32 \t$r, $s;",
1617                              [(set Int32Regs:$r,
1618                                (int_nvvm_move_i32 Int32Regs:$s))]>;
1619 def nvvm_move_i64 : NVPTXInst<(outs Int64Regs:$r), (ins Int64Regs:$s),
1620                              "mov.b64 \t$r, $s;",
1621                              [(set Int64Regs:$r,
1622                                (int_nvvm_move_i64 Int64Regs:$s))]>;
1623 def nvvm_move_float : NVPTXInst<(outs Float32Regs:$r), (ins Float32Regs:$s),
1624                              "mov.f32 \t$r, $s;",
1625                              [(set Float32Regs:$r,
1626                                (int_nvvm_move_float Float32Regs:$s))]>;
1627 def nvvm_move_double : NVPTXInst<(outs Float64Regs:$r), (ins Float64Regs:$s),
1628                              "mov.f64 \t$r, $s;",
1629                              [(set Float64Regs:$r,
1630                                (int_nvvm_move_double Float64Regs:$s))]>;
1631 def nvvm_move_ptr32 : NVPTXInst<(outs Int32Regs:$r), (ins Int32Regs:$s),
1632                              "mov.u32 \t$r, $s;",
1633                              [(set Int32Regs:$r,
1634                                (int_nvvm_move_ptr Int32Regs:$s))]>;
1635 def nvvm_move_ptr64 : NVPTXInst<(outs Int64Regs:$r), (ins Int64Regs:$s),
1636                              "mov.u64 \t$r, $s;",
1637                              [(set Int64Regs:$r,
1638                                (int_nvvm_move_ptr Int64Regs:$s))]>;
1639
1640 // @TODO: Are these actually needed, or will we always just see symbols
1641 // copied to registers first?
1642 /*def nvvm_move_sym32 : NVPTXInst<(outs Int32Regs:$r), (ins imem:$s),
1643                              "mov.u32 \t$r, $s;",
1644                              [(set Int32Regs:$r,
1645                              (int_nvvm_move_ptr texternalsym:$s))]>;
1646 def nvvm_move_sym64 : NVPTXInst<(outs Int64Regs:$r), (ins imem:$s),
1647                              "mov.u64 \t$r, $s;",
1648                              [(set Int64Regs:$r,
1649                              (int_nvvm_move_ptr texternalsym:$s))]>;*/
1650
1651
1652 // MoveParam        %r1, param
1653 // ptr_local_to_gen %r2, %r1
1654 // ptr_gen_to_local %r3, %r2
1655 // ->
1656 // mov %r1, param
1657
1658 // @TODO: Revisit this.  There is a type
1659 // contradiction between iPTRAny and iPTR for the addr defs, so the move_sym
1660 // instructions are not currently defined. However, we can use the ptr
1661 // variants and the asm printer will do the right thing.
1662 def : Pat<(i64 (int_nvvm_ptr_gen_to_local (int_nvvm_ptr_local_to_gen
1663                 (MoveParam texternalsym:$src)))),
1664                (nvvm_move_ptr64  texternalsym:$src)>;
1665 def : Pat<(i32 (int_nvvm_ptr_gen_to_local (int_nvvm_ptr_local_to_gen
1666                 (MoveParam texternalsym:$src)))),
1667                (nvvm_move_ptr32  texternalsym:$src)>;
1668
1669 def texsurf_handles
1670   : NVPTXInst<(outs Int64Regs:$result), (ins imem:$src),
1671               "mov.u64 \t$result, $src;", []>;
1672
1673 //-----------------------------------
1674 // Compiler Error Warn
1675 // - Just ignore them in codegen
1676 //-----------------------------------
1677
1678 def INT_NVVM_COMPILER_WARN_32 : NVPTXInst<(outs), (ins Int32Regs:$a),
1679                 "// llvm.nvvm.compiler.warn()",
1680                 [(int_nvvm_compiler_warn Int32Regs:$a)]>;
1681 def INT_NVVM_COMPILER_WARN_64 : NVPTXInst<(outs), (ins Int64Regs:$a),
1682                 "// llvm.nvvm.compiler.warn()",
1683                 [(int_nvvm_compiler_warn Int64Regs:$a)]>;
1684 def INT_NVVM_COMPILER_ERROR_32 : NVPTXInst<(outs), (ins Int32Regs:$a),
1685                 "// llvm.nvvm.compiler.error()",
1686                 [(int_nvvm_compiler_error Int32Regs:$a)]>;
1687 def INT_NVVM_COMPILER_ERROR_64 : NVPTXInst<(outs), (ins Int64Regs:$a),
1688                 "// llvm.nvvm.compiler.error()",
1689                 [(int_nvvm_compiler_error Int64Regs:$a)]>;
1690
1691
1692 // isspacep
1693
1694 def ISSPACEP_CONST_32
1695   : NVPTXInst<(outs Int1Regs:$d), (ins Int32Regs:$a),
1696               "isspacep.const \t$d, $a;",
1697               [(set Int1Regs:$d, (int_nvvm_isspacep_const Int32Regs:$a))]>,
1698     Requires<[hasPTX31]>;
1699 def ISSPACEP_CONST_64
1700   : NVPTXInst<(outs Int1Regs:$d), (ins Int64Regs:$a),
1701               "isspacep.const \t$d, $a;",
1702               [(set Int1Regs:$d, (int_nvvm_isspacep_const Int64Regs:$a))]>,
1703     Requires<[hasPTX31]>;
1704 def ISSPACEP_GLOBAL_32
1705   : NVPTXInst<(outs Int1Regs:$d), (ins Int32Regs:$a),
1706               "isspacep.global \t$d, $a;",
1707               [(set Int1Regs:$d, (int_nvvm_isspacep_global Int32Regs:$a))]>;
1708 def ISSPACEP_GLOBAL_64
1709   : NVPTXInst<(outs Int1Regs:$d), (ins Int64Regs:$a),
1710               "isspacep.global \t$d, $a;",
1711               [(set Int1Regs:$d, (int_nvvm_isspacep_global Int64Regs:$a))]>;
1712 def ISSPACEP_LOCAL_32
1713   : NVPTXInst<(outs Int1Regs:$d), (ins Int32Regs:$a),
1714               "isspacep.local \t$d, $a;",
1715               [(set Int1Regs:$d, (int_nvvm_isspacep_local Int32Regs:$a))]>;
1716 def ISSPACEP_LOCAL_64
1717   : NVPTXInst<(outs Int1Regs:$d), (ins Int64Regs:$a),
1718               "isspacep.local \t$d, $a;",
1719               [(set Int1Regs:$d, (int_nvvm_isspacep_local Int64Regs:$a))]>;
1720 def ISSPACEP_SHARED_32
1721   : NVPTXInst<(outs Int1Regs:$d), (ins Int32Regs:$a),
1722               "isspacep.shared \t$d, $a;",
1723               [(set Int1Regs:$d, (int_nvvm_isspacep_shared Int32Regs:$a))]>;
1724 def ISSPACEP_SHARED_64
1725   : NVPTXInst<(outs Int1Regs:$d), (ins Int64Regs:$a),
1726               "isspacep.shared \t$d, $a;",
1727               [(set Int1Regs:$d, (int_nvvm_isspacep_shared Int64Regs:$a))]>;
1728
1729
1730 // Special register reads
1731 def MOV_SPECIAL : NVPTXInst<(outs Int32Regs:$d),
1732                             (ins SpecialRegs:$r),
1733                             "mov.b32\t$d, $r;", []>;
1734
1735 def : Pat<(int_nvvm_read_ptx_sreg_envreg0), (MOV_SPECIAL ENVREG0)>;
1736 def : Pat<(int_nvvm_read_ptx_sreg_envreg1), (MOV_SPECIAL ENVREG1)>;
1737 def : Pat<(int_nvvm_read_ptx_sreg_envreg2), (MOV_SPECIAL ENVREG2)>;
1738 def : Pat<(int_nvvm_read_ptx_sreg_envreg3), (MOV_SPECIAL ENVREG3)>;
1739 def : Pat<(int_nvvm_read_ptx_sreg_envreg4), (MOV_SPECIAL ENVREG4)>;
1740 def : Pat<(int_nvvm_read_ptx_sreg_envreg5), (MOV_SPECIAL ENVREG5)>;
1741 def : Pat<(int_nvvm_read_ptx_sreg_envreg6), (MOV_SPECIAL ENVREG6)>;
1742 def : Pat<(int_nvvm_read_ptx_sreg_envreg7), (MOV_SPECIAL ENVREG7)>;
1743 def : Pat<(int_nvvm_read_ptx_sreg_envreg8), (MOV_SPECIAL ENVREG8)>;
1744 def : Pat<(int_nvvm_read_ptx_sreg_envreg9), (MOV_SPECIAL ENVREG9)>;
1745 def : Pat<(int_nvvm_read_ptx_sreg_envreg10), (MOV_SPECIAL ENVREG10)>;
1746 def : Pat<(int_nvvm_read_ptx_sreg_envreg11), (MOV_SPECIAL ENVREG11)>;
1747 def : Pat<(int_nvvm_read_ptx_sreg_envreg12), (MOV_SPECIAL ENVREG12)>;
1748 def : Pat<(int_nvvm_read_ptx_sreg_envreg13), (MOV_SPECIAL ENVREG13)>;
1749 def : Pat<(int_nvvm_read_ptx_sreg_envreg14), (MOV_SPECIAL ENVREG14)>;
1750 def : Pat<(int_nvvm_read_ptx_sreg_envreg15), (MOV_SPECIAL ENVREG15)>;
1751 def : Pat<(int_nvvm_read_ptx_sreg_envreg16), (MOV_SPECIAL ENVREG16)>;
1752 def : Pat<(int_nvvm_read_ptx_sreg_envreg17), (MOV_SPECIAL ENVREG17)>;
1753 def : Pat<(int_nvvm_read_ptx_sreg_envreg18), (MOV_SPECIAL ENVREG18)>;
1754 def : Pat<(int_nvvm_read_ptx_sreg_envreg19), (MOV_SPECIAL ENVREG19)>;
1755 def : Pat<(int_nvvm_read_ptx_sreg_envreg20), (MOV_SPECIAL ENVREG20)>;
1756 def : Pat<(int_nvvm_read_ptx_sreg_envreg21), (MOV_SPECIAL ENVREG21)>;
1757 def : Pat<(int_nvvm_read_ptx_sreg_envreg22), (MOV_SPECIAL ENVREG22)>;
1758 def : Pat<(int_nvvm_read_ptx_sreg_envreg23), (MOV_SPECIAL ENVREG23)>;
1759 def : Pat<(int_nvvm_read_ptx_sreg_envreg24), (MOV_SPECIAL ENVREG24)>;
1760 def : Pat<(int_nvvm_read_ptx_sreg_envreg25), (MOV_SPECIAL ENVREG25)>;
1761 def : Pat<(int_nvvm_read_ptx_sreg_envreg26), (MOV_SPECIAL ENVREG26)>;
1762 def : Pat<(int_nvvm_read_ptx_sreg_envreg27), (MOV_SPECIAL ENVREG27)>;
1763 def : Pat<(int_nvvm_read_ptx_sreg_envreg28), (MOV_SPECIAL ENVREG28)>;
1764 def : Pat<(int_nvvm_read_ptx_sreg_envreg29), (MOV_SPECIAL ENVREG29)>;
1765 def : Pat<(int_nvvm_read_ptx_sreg_envreg30), (MOV_SPECIAL ENVREG30)>;
1766 def : Pat<(int_nvvm_read_ptx_sreg_envreg31), (MOV_SPECIAL ENVREG31)>;
1767
1768
1769 //-----------------------------------
1770 // Texture Intrinsics
1771 //-----------------------------------
1772
1773 // NOTE: For Fermi support, any new texture/surface/sampler intrinsics must be
1774 // also defined in NVPTXReplaceImageHandles.cpp
1775
1776
1777 // Texture fetch instructions using handles
1778 def TEX_1D_F32_I32
1779   : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
1780                     Float32Regs:$b, Float32Regs:$a),
1781               (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$x),
1782               "tex.1d.v4.f32.s32\t\\{$r, $g, $b, $a\\}, [$t, $s, \\{$x\\}];",
1783               []>;
1784 def TEX_1D_F32_F32
1785   : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
1786                     Float32Regs:$b, Float32Regs:$a),
1787               (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x),
1788               "tex.1d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, [$t, $s, \\{$x\\}];",
1789               []>;
1790 def TEX_1D_F32_F32_LEVEL
1791   : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
1792                     Float32Regs:$b, Float32Regs:$a),
1793               (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$lod),
1794               "tex.level.1d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
1795               "[$t, $s, \\{$x\\}], $lod;",
1796               []>;
1797 def TEX_1D_F32_F32_GRAD
1798   : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
1799                     Float32Regs:$b, Float32Regs:$a),
1800               (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x,
1801                    Float32Regs:$gradx, Float32Regs:$grady),
1802               "tex.grad.1d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
1803               "[$t, $s, \\{$x\\}], \\{$gradx\\}, \\{$grady\\};",
1804               []>;
1805 def TEX_1D_I32_I32
1806   : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
1807                     Int32Regs:$b, Int32Regs:$a),
1808               (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$x),
1809               "tex.1d.v4.s32.s32\t\\{$r, $g, $b, $a\\}, [$t, $s, \\{$x\\}];",
1810               []>;
1811 def TEX_1D_I32_F32
1812   : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
1813                     Int32Regs:$b, Int32Regs:$a),
1814               (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x),
1815               "tex.1d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, [$t, $s, \\{$x\\}];",
1816               []>;
1817 def TEX_1D_I32_F32_LEVEL
1818   : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
1819                     Int32Regs:$b, Int32Regs:$a),
1820               (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x,
1821                    Float32Regs:$lod),
1822               "tex.level.1d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
1823               "[$t, $s, \\{$x\\}], $lod;",
1824               []>;
1825 def TEX_1D_I32_F32_GRAD
1826   : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
1827                     Int32Regs:$b, Int32Regs:$a),
1828               (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x,
1829                    Float32Regs:$gradx, Float32Regs:$grady),
1830               "tex.grad.1d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
1831               "[$t, $s, \\{$x\\}], \\{$gradx\\}, \\{$grady\\};",
1832               []>;
1833
1834 def TEX_1D_ARRAY_F32_I32
1835   : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
1836                     Float32Regs:$b, Float32Regs:$a),
1837               (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Int32Regs:$x),
1838               "tex.a1d.v4.f32.s32\t\\{$r, $g, $b, $a\\}, "
1839               "[$t, $s, \\{$l, $x\\}];",
1840               []>;
1841 def TEX_1D_ARRAY_F32_F32
1842   : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
1843                     Float32Regs:$b, Float32Regs:$a),
1844               (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x),
1845               "tex.a1d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
1846               "[$t, $s, \\{$l, $x\\}];",
1847               []>;
1848 def TEX_1D_ARRAY_F32_F32_LEVEL
1849   : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
1850                     Float32Regs:$b, Float32Regs:$a),
1851               (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
1852                    Float32Regs:$lod),
1853               "tex.level.a1d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
1854               "[$t, $s, \\{$l, $x\\}], $lod;",
1855               []>;
1856 def TEX_1D_ARRAY_F32_F32_GRAD
1857   : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
1858                     Float32Regs:$b, Float32Regs:$a),
1859               (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
1860                    Float32Regs:$gradx, Float32Regs:$grady),
1861               "tex.grad.a1d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
1862               "[$t, $s, \\{$l, $x\\}], \\{$gradx\\}, \\{$grady\\};",
1863               []>;
1864 def TEX_1D_ARRAY_I32_I32
1865   : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
1866                     Int32Regs:$b, Int32Regs:$a),
1867               (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Int32Regs:$x),
1868               "tex.a1d.v4.s32.s32\t\\{$r, $g, $b, $a\\}, "
1869               "[$t, $s, \\{$l, $x\\}];",
1870               []>;
1871 def TEX_1D_ARRAY_I32_F32
1872   : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
1873                     Int32Regs:$b, Int32Regs:$a),
1874               (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x),
1875               "tex.a1d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
1876               "[$t, $s, \\{$l, $x\\}];",
1877               []>;
1878 def TEX_1D_ARRAY_I32_F32_LEVEL
1879   : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
1880                     Int32Regs:$b, Int32Regs:$a),
1881               (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
1882                    Float32Regs:$lod),
1883               "tex.level.a1d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
1884               "[$t, $s, \\{$l, $x\\}], $lod;",
1885               []>;
1886 def TEX_1D_ARRAY_I32_F32_GRAD
1887   : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
1888                     Int32Regs:$b, Int32Regs:$a),
1889               (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
1890                    Float32Regs:$gradx, Float32Regs:$grady),
1891               "tex.grad.a1d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
1892               "[$t, $s, \\{$l, $x\\}], \\{$gradx\\}, \\{$grady\\};",
1893               []>;
1894
1895 def TEX_2D_F32_I32
1896   : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
1897                     Float32Regs:$b, Float32Regs:$a),
1898               (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$x, Int32Regs:$y),
1899               "tex.2d.v4.f32.s32\t\\{$r, $g, $b, $a\\}, "
1900               "[$t, $s, \\{$x, $y\\}];",
1901               []>;
1902 def TEX_2D_F32_F32
1903   : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
1904                     Float32Regs:$b, Float32Regs:$a),
1905               (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y),
1906               "tex.2d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
1907               "[$t, $s, \\{$x, $y\\}];",
1908               []>;
1909 def TEX_2D_F32_F32_LEVEL
1910   : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
1911                     Float32Regs:$b, Float32Regs:$a),
1912               (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
1913                    Float32Regs:$lod),
1914               "tex.level.2d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
1915               "[$t, $s, \\{$x, $y\\}], $lod;",
1916               []>;
1917 def TEX_2D_F32_F32_GRAD
1918   : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
1919                     Float32Regs:$b, Float32Regs:$a),
1920               (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
1921                    Float32Regs:$gradx0, Float32Regs:$gradx1,
1922                    Float32Regs:$grady0, Float32Regs:$grady1),
1923               "tex.grad.2d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
1924               "[$t, $s, \\{$x, $y\\}], \\{$gradx0, $gradx1\\}, "
1925               "\\{$grady0, $grady1\\};",
1926               []>;
1927 def TEX_2D_I32_I32
1928   : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
1929                     Int32Regs:$b, Int32Regs:$a),
1930               (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$x, Int32Regs:$y),
1931               "tex.2d.v4.s32.s32\t\\{$r, $g, $b, $a\\}, "
1932               "[$t, $s, \\{$x, $y\\}];",
1933               []>;
1934 def TEX_2D_I32_F32
1935   : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
1936                     Int32Regs:$b, Int32Regs:$a),
1937               (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y),
1938               "tex.2d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
1939               "[$t, $s, \\{$x, $y\\}];",
1940               []>;
1941 def TEX_2D_I32_F32_LEVEL
1942   : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
1943                     Int32Regs:$b, Int32Regs:$a),
1944               (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
1945                    Float32Regs:$lod),
1946               "tex.level.2d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
1947               "[$t, $s, \\{$x, $y\\}], $lod;",
1948               []>;
1949 def TEX_2D_I32_F32_GRAD
1950   : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
1951                     Int32Regs:$b, Int32Regs:$a),
1952               (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
1953                    Float32Regs:$gradx0, Float32Regs:$gradx1,
1954                    Float32Regs:$grady0, Float32Regs:$grady1),
1955               "tex.grad.2d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
1956               "[$t, $s, \\{$x, $y\\}], \\{$gradx0, $gradx1\\}, "
1957               "\\{$grady0, $grady1\\};",
1958               []>;
1959
1960 def TEX_2D_ARRAY_F32_I32
1961   : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
1962                     Float32Regs:$b, Float32Regs:$a),
1963               (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
1964                    Int32Regs:$y),
1965               "tex.a2d.v4.f32.s32\t\\{$r, $g, $b, $a\\}, "
1966               "[$t, $s, \\{$l, $x, $y, $y\\}];",
1967               []>;
1968 def TEX_2D_ARRAY_F32_F32
1969   : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
1970                     Float32Regs:$b, Float32Regs:$a),
1971               (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
1972                    Float32Regs:$y),
1973               "tex.a2d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
1974               "[$t, $s, \\{$l, $x, $y, $y\\}];",
1975               []>;
1976 def TEX_2D_ARRAY_F32_F32_LEVEL
1977   : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
1978                     Float32Regs:$b, Float32Regs:$a),
1979               (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
1980                    Float32Regs:$y, Float32Regs:$lod),
1981               "tex.level.a2d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
1982               "[$t, $s, \\{$l, $x, $y, $y\\}], $lod;",
1983               []>;
1984 def TEX_2D_ARRAY_F32_F32_GRAD
1985   : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
1986                     Float32Regs:$b, Float32Regs:$a),
1987               (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
1988                    Float32Regs:$y, Float32Regs:$gradx0, Float32Regs:$gradx1,
1989                    Float32Regs:$grady0, Float32Regs:$grady1),
1990               "tex.grad.a2d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
1991               "[$t, $s, \\{$l, $x, $y, $y\\}], \\{$gradx0, $gradx1\\}, "
1992               "\\{$grady0, $grady1\\};",
1993               []>;
1994 def TEX_2D_ARRAY_I32_I32
1995   : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
1996                     Int32Regs:$b, Int32Regs:$a),
1997               (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
1998                    Int32Regs:$y),
1999               "tex.a2d.v4.s32.s32\t\\{$r, $g, $b, $a\\}, "
2000               "[$t, $s, \\{$l, $x, $y, $y\\}];",
2001               []>;
2002 def TEX_2D_ARRAY_I32_F32
2003   : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
2004                     Int32Regs:$b, Int32Regs:$a),
2005               (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
2006                    Float32Regs:$y),
2007               "tex.a2d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
2008               "[$t, $s, \\{$l, $x, $y, $y\\}];",
2009               []>;
2010 def TEX_2D_ARRAY_I32_F32_LEVEL
2011   : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
2012                     Int32Regs:$b, Int32Regs:$a),
2013               (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
2014                    Float32Regs:$y, Float32Regs:$lod),
2015               "tex.level.a2d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
2016               "[$t, $s, \\{$l, $x, $y, $y\\}], $lod;",
2017               []>;
2018 def TEX_2D_ARRAY_I32_F32_GRAD
2019   : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
2020                     Int32Regs:$b, Int32Regs:$a),
2021               (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
2022                    Float32Regs:$y,
2023                    Float32Regs:$gradx0, Float32Regs:$gradx1,
2024                    Float32Regs:$grady0, Float32Regs:$grady1),
2025               "tex.grad.a2d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
2026               "[$t, $s, \\{$l, $x, $y, $y\\}], \\{$gradx0, $gradx1\\}, "
2027               "\\{$grady0, $grady1\\};",
2028               []>;
2029
2030 def TEX_3D_F32_I32
2031   : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
2032                     Float32Regs:$b, Float32Regs:$a),
2033               (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
2034                    Int32Regs:$z),
2035               "tex.3d.v4.f32.s32\t\\{$r, $g, $b, $a\\}, "
2036               "[$t, $s, \\{$x, $y, $z, $z\\}];",
2037               []>;
2038 def TEX_3D_F32_F32
2039   : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
2040                     Float32Regs:$b, Float32Regs:$a),
2041               (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
2042                    Float32Regs:$z),
2043               "tex.3d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
2044               "[$t, $s, \\{$x, $y, $z, $z\\}];",
2045               []>;
2046 def TEX_3D_F32_F32_LEVEL
2047   : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
2048                     Float32Regs:$b, Float32Regs:$a),
2049               (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
2050                    Float32Regs:$z, Float32Regs:$lod),
2051               "tex.level.3d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
2052               "[$t, $s, \\{$x, $y, $z, $z\\}], $lod;",
2053               []>;
2054 def TEX_3D_F32_F32_GRAD
2055   : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
2056                     Float32Regs:$b, Float32Regs:$a),
2057               (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
2058                    Float32Regs:$z,
2059                    Float32Regs:$gradx0, Float32Regs:$gradx1,
2060                    Float32Regs:$gradx2, Float32Regs:$grady0,
2061                    Float32Regs:$grady1, Float32Regs:$grady2),
2062               "tex.grad.3d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
2063               "[$t, $s, \\{$x, $y, $z, $z\\}], "
2064               "\\{$gradx0, $gradx1, $gradx2, $gradx2\\}, "
2065               "\\{$grady0, $grady1, $grady2, $grady2\\};",
2066               []>;
2067 def TEX_3D_I32_I32
2068   : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
2069                     Int32Regs:$b, Int32Regs:$a),
2070               (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
2071                    Int32Regs:$z),
2072               "tex.3d.v4.s32.s32\t\\{$r, $g, $b, $a\\}, "
2073               "[$t, $s, \\{$x, $y, $z, $z\\}];",
2074               []>;
2075 def TEX_3D_I32_F32
2076   : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
2077                     Int32Regs:$b, Int32Regs:$a),
2078               (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
2079                    Float32Regs:$z),
2080               "tex.3d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
2081               "[$t, $s, \\{$x, $y, $z, $z\\}];",
2082               []>;
2083 def TEX_3D_I32_F32_LEVEL
2084   : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
2085                     Int32Regs:$b, Int32Regs:$a),
2086               (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
2087                    Float32Regs:$z, Float32Regs:$lod),
2088               "tex.level.3d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
2089               "[$t, $s, \\{$x, $y, $z, $z\\}], $lod;",
2090               []>;
2091 def TEX_3D_I32_F32_GRAD
2092   : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
2093                     Int32Regs:$b, Int32Regs:$a),
2094               (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
2095                    Float32Regs:$z,
2096                    Float32Regs:$gradx0, Float32Regs:$gradx1,
2097                    Float32Regs:$gradx2, Float32Regs:$grady0,
2098                    Float32Regs:$grady1, Float32Regs:$grady2),
2099               "tex.grad.3d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
2100               "[$t, $s, \\{$x, $y, $z, $z\\}], "
2101               "\\{$gradx0, $gradx1, $gradx2, $gradx2\\}, "
2102               "\\{$grady0, $grady1, $grady2, $grady2\\};",
2103               []>;
2104
2105
2106 // Surface load instructions
2107 def SULD_1D_I8_TRAP
2108   : NVPTXInst<(outs Int16Regs:$r),
2109               (ins Int64Regs:$s, Int32Regs:$x),
2110               "suld.b.1d.b8.trap \\{$r\\}, [$s, \\{$x\\}];",
2111               []>;
2112 def SULD_1D_I16_TRAP
2113   : NVPTXInst<(outs Int16Regs:$r),
2114               (ins Int64Regs:$s, Int32Regs:$x),
2115               "suld.b.1d.b16.trap \\{$r\\}, [$s, \\{$x\\}];",
2116               []>;
2117 def SULD_1D_I32_TRAP
2118   : NVPTXInst<(outs Int32Regs:$r),
2119               (ins Int64Regs:$s, Int32Regs:$x),
2120               "suld.b.1d.b32.trap \\{$r\\}, [$s, \\{$x\\}];",
2121               []>;
2122 def SULD_1D_V2I8_TRAP
2123   : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g),
2124               (ins Int64Regs:$s, Int32Regs:$x),
2125               "suld.b.1d.v2.b8.trap \\{$r, $g\\}, [$s, \\{$x\\}];",
2126               []>;
2127 def SULD_1D_V2I16_TRAP
2128   : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g),
2129               (ins Int64Regs:$s, Int32Regs:$x),
2130               "suld.b.1d.v2.b16.trap \\{$r, $g\\}, [$s, \\{$x\\}];",
2131               []>;
2132 def SULD_1D_V2I32_TRAP
2133   : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g),
2134               (ins Int64Regs:$s, Int32Regs:$x),
2135               "suld.b.1d.v2.b32.trap \\{$r, $g\\}, [$s, \\{$x\\}];",
2136               []>;
2137 def SULD_1D_V4I8_TRAP
2138   : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2139               (ins Int64Regs:$s, Int32Regs:$x),
2140               "suld.b.1d.v4.b8.trap \\{$r, $g, $b, $a\\}, [$s, \\{$x\\}];",
2141               []>;
2142 def SULD_1D_V4I16_TRAP
2143   : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2144               (ins Int64Regs:$s, Int32Regs:$x),
2145               "suld.b.1d.v4.b16.trap \\{$r, $g, $b, $a\\}, [$s, \\{$x\\}];",
2146               []>;
2147 def SULD_1D_V4I32_TRAP
2148   : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
2149               (ins Int64Regs:$s, Int32Regs:$x),
2150               "suld.b.1d.v4.b32.trap \\{$r, $g, $b, $a\\}, [$s, \\{$x\\}];",
2151               []>;
2152
2153 def SULD_1D_ARRAY_I8_TRAP
2154   : NVPTXInst<(outs Int16Regs:$r),
2155               (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x),
2156               "suld.b.a1d.b8.trap \\{$r\\}, [$s, \\{$l, $x\\}];",
2157               []>;
2158 def SULD_1D_ARRAY_I16_TRAP
2159   : NVPTXInst<(outs Int16Regs:$r),
2160               (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x),
2161               "suld.b.a1d.b16.trap \\{$r\\}, [$s, \\{$l, $x\\}];",
2162               []>;
2163 def SULD_1D_ARRAY_I32_TRAP
2164   : NVPTXInst<(outs Int32Regs:$r),
2165               (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x),
2166               "suld.b.a1d.b32.trap \\{$r\\}, [$s, \\{$l, $x\\}];",
2167               []>;
2168 def SULD_1D_ARRAY_V2I8_TRAP
2169   : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g),
2170               (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x),
2171               "suld.b.a1d.v2.b8.trap \\{$r, $g\\}, [$s, \\{$l, $x\\}];",
2172               []>;
2173 def SULD_1D_ARRAY_V2I16_TRAP
2174   : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g),
2175               (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x),
2176               "suld.b.a1d.v2.b16.trap \\{$r, $g\\}, [$s, \\{$l, $x\\}];",
2177               []>;
2178 def SULD_1D_ARRAY_V2I32_TRAP
2179   : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g),
2180               (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x),
2181               "suld.b.a1d.v2.b32.trap \\{$r, $g\\}, [$s, \\{$l, $x\\}];",
2182               []>;
2183 def SULD_1D_ARRAY_V4I8_TRAP
2184   : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2185               (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x),
2186               "suld.b.a1d.v4.b8.trap \\{$r, $g, $b, $a\\}, "
2187               "[$s, \\{$l, $x\\}];",
2188               []>;
2189 def SULD_1D_ARRAY_V4I16_TRAP
2190   : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2191               (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x),
2192               "suld.b.a1d.v4.b16.trap \\{$r, $g, $b, $a\\}, "
2193               "[$s, \\{$l, $x\\}];",
2194               []>;
2195 def SULD_1D_ARRAY_V4I32_TRAP
2196   : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
2197               (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x),
2198               "suld.b.a1d.v4.b32.trap \\{$r, $g, $b, $a\\}, "
2199               "[$s, \\{$l, $x\\}];",
2200               []>;
2201
2202 def SULD_2D_I8_TRAP
2203   : NVPTXInst<(outs Int16Regs:$r),
2204               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y),
2205               "suld.b.2d.b8.trap \\{$r\\}, [$s, \\{$x, $y\\}];",
2206               []>;
2207 def SULD_2D_I16_TRAP
2208   : NVPTXInst<(outs Int16Regs:$r),
2209               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y),
2210               "suld.b.2d.b16.trap \\{$r\\}, [$s, \\{$x, $y\\}];",
2211               []>;
2212 def SULD_2D_I32_TRAP
2213   : NVPTXInst<(outs Int32Regs:$r),
2214               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y),
2215               "suld.b.2d.b32.trap \\{$r\\}, [$s, \\{$x, $y\\}];",
2216               []>;
2217 def SULD_2D_V2I8_TRAP
2218   : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g),
2219               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y),
2220               "suld.b.2d.v2.b8.trap \\{$r, $g\\}, [$s, \\{$x, $y\\}];",
2221               []>;
2222 def SULD_2D_V2I16_TRAP
2223   : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g),
2224               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y),
2225               "suld.b.2d.v2.b16.trap \\{$r, $g\\}, [$s, \\{$x, $y\\}];",
2226               []>;
2227 def SULD_2D_V2I32_TRAP
2228   : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g),
2229               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y),
2230               "suld.b.2d.v2.b32.trap \\{$r, $g\\}, [$s, \\{$x, $y\\}];",
2231               []>;
2232 def SULD_2D_V4I8_TRAP
2233   : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2234               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y),
2235               "suld.b.2d.v4.b8.trap \\{$r, $g, $b, $a\\}, [$s, \\{$x, $y\\}];",
2236               []>;
2237 def SULD_2D_V4I16_TRAP
2238   : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2239               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y),
2240               "suld.b.2d.v4.b16.trap \\{$r, $g, $b, $a\\}, [$s, \\{$x, $y\\}];",
2241               []>;
2242 def SULD_2D_V4I32_TRAP
2243   : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
2244               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y),
2245               "suld.b.2d.v4.b32.trap \\{$r, $g, $b, $a\\}, [$s, \\{$x, $y\\}];",
2246               []>;
2247
2248 def SULD_2D_ARRAY_I8_TRAP
2249   : NVPTXInst<(outs Int16Regs:$r),
2250               (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y),
2251               "suld.b.a2d.b8.trap \\{$r\\}, [$s, \\{$l, $x, $y, $y\\}];",
2252               []>;
2253 def SULD_2D_ARRAY_I16_TRAP
2254   : NVPTXInst<(outs Int16Regs:$r),
2255               (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y),
2256               "suld.b.a2d.b16.trap \\{$r\\}, [$s, \\{$l, $x, $y, $y\\}];",
2257               []>;
2258 def SULD_2D_ARRAY_I32_TRAP
2259   : NVPTXInst<(outs Int32Regs:$r),
2260               (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y),
2261               "suld.b.a2d.b32.trap \\{$r\\}, [$s, \\{$l, $x, $y, $y\\}];",
2262               []>;
2263 def SULD_2D_ARRAY_V2I8_TRAP
2264   : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g),
2265               (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y),
2266               "suld.b.a2d.v2.b8.trap \\{$r, $g\\}, "
2267               "[$s, \\{$l, $x, $y, $y\\}];",
2268               []>;
2269 def SULD_2D_ARRAY_V2I16_TRAP
2270   : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g),
2271               (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y),
2272               "suld.b.a2d.v2.b16.trap \\{$r, $g\\}, "
2273               "[$s, \\{$l, $x, $y, $y\\}];",
2274               []>;
2275 def SULD_2D_ARRAY_V2I32_TRAP
2276   : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g),
2277               (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y),
2278               "suld.b.a2d.v2.b32.trap \\{$r, $g\\}, "
2279               "[$s, \\{$l, $x, $y, $y\\}];",
2280               []>;
2281 def SULD_2D_ARRAY_V4I8_TRAP
2282   : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2283               (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y),
2284               "suld.b.a2d.v4.b8.trap \\{$r, $g, $b, $a\\}, "
2285               "[$s, \\{$l, $x, $y, $y\\}];",
2286               []>;
2287 def SULD_2D_ARRAY_V4I16_TRAP
2288   : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2289               (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y),
2290               "suld.b.a2d.v4.b16.trap \\{$r, $g, $b, $a\\}, "
2291               "[$s, \\{$l, $x, $y, $y\\}];",
2292               []>;
2293 def SULD_2D_ARRAY_V4I32_TRAP
2294   : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
2295               (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y),
2296               "suld.b.a2d.v4.b32.trap \\{$r, $g, $b, $a\\}, "
2297               "[$s, \\{$l, $x, $y, $y\\}];",
2298               []>;
2299
2300 def SULD_3D_I8_TRAP
2301   : NVPTXInst<(outs Int16Regs:$r),
2302               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z),
2303               "suld.b.3d.b8.trap \\{$r\\}, [$s, \\{$x, $y, $z, $z\\}];",
2304               []>;
2305 def SULD_3D_I16_TRAP
2306   : NVPTXInst<(outs Int16Regs:$r),
2307               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z),
2308               "suld.b.3d.b16.trap \\{$r\\}, [$s, \\{$x, $y, $z, $z\\}];",
2309               []>;
2310 def SULD_3D_I32_TRAP
2311   : NVPTXInst<(outs Int32Regs:$r),
2312               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z),
2313               "suld.b.3d.b32.trap \\{$r\\}, [$s, \\{$x, $y, $z, $z\\}];",
2314               []>;
2315 def SULD_3D_V2I8_TRAP
2316   : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g),
2317               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z),
2318               "suld.b.3d.v2.b8.trap \\{$r, $g\\}, [$s, \\{$x, $y, $z, $z\\}];",
2319               []>;
2320 def SULD_3D_V2I16_TRAP
2321   : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g),
2322               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z),
2323               "suld.b.3d.v2.b16.trap \\{$r, $g\\}, [$s, \\{$x, $y, $z, $z\\}];",
2324               []>;
2325 def SULD_3D_V2I32_TRAP
2326   : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g),
2327               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z),
2328               "suld.b.3d.v2.b32.trap \\{$r, $g\\}, [$s, \\{$x, $y, $z, $z\\}];",
2329               []>;
2330 def SULD_3D_V4I8_TRAP
2331   : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2332               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z),
2333               "suld.b.3d.v4.b8.trap \\{$r, $g, $b, $a\\}, "
2334               "[$s, \\{$x, $y, $z, $z\\}];",
2335               []>;
2336 def SULD_3D_V4I16_TRAP
2337   : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2338               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z),
2339               "suld.b.3d.v4.b16.trap \\{$r, $g, $b, $a\\}, "
2340               "[$s, \\{$x, $y, $z, $z\\}];",
2341               []>;
2342 def SULD_3D_V4I32_TRAP
2343   : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
2344               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z),
2345               "suld.b.3d.v4.b32.trap \\{$r, $g, $b, $a\\}, "
2346               "[$s, \\{$x, $y, $z, $z\\}];",
2347               []>;
2348
2349
2350 //-----------------------------------
2351 // Texture Query Intrinsics
2352 //-----------------------------------
2353 def TXQ_CHANNEL_ORDER
2354   : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
2355               "txq.channel_order.b32 \t$d, [$a];",
2356               []>;
2357 def TXQ_CHANNEL_DATA_TYPE
2358   : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
2359               "txq.channel_data_type.b32 \t$d, [$a];",
2360               []>;
2361 def TXQ_WIDTH
2362   : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
2363               "txq.width.b32 \t$d, [$a];",
2364               []>;
2365 def TXQ_HEIGHT
2366   : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
2367               "txq.height.b32 \t$d, [$a];",
2368               []>;
2369 def TXQ_DEPTH
2370   : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
2371               "txq.depth.b32 \t$d, [$a];",
2372               []>;
2373 def TXQ_ARRAY_SIZE
2374   : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
2375               "txq.array_size.b32 \t$d, [$a];",
2376               []>;
2377 def TXQ_NUM_SAMPLES
2378   : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
2379               "txq.num_samples.b32 \t$d, [$a];",
2380               []>;
2381 def TXQ_NUM_MIPMAP_LEVELS
2382   : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
2383               "txq.num_mipmap_levels.b32 \t$d, [$a];",
2384               []>;
2385
2386 def : Pat<(int_nvvm_txq_channel_order Int64Regs:$a),
2387           (TXQ_CHANNEL_ORDER Int64Regs:$a)>;
2388 def : Pat<(int_nvvm_txq_channel_data_type Int64Regs:$a),
2389           (TXQ_CHANNEL_DATA_TYPE Int64Regs:$a)>;
2390 def : Pat<(int_nvvm_txq_width Int64Regs:$a),
2391           (TXQ_WIDTH Int64Regs:$a)>;
2392 def : Pat<(int_nvvm_txq_height Int64Regs:$a),
2393           (TXQ_HEIGHT Int64Regs:$a)>;
2394 def : Pat<(int_nvvm_txq_depth Int64Regs:$a),
2395           (TXQ_DEPTH Int64Regs:$a)>;
2396 def : Pat<(int_nvvm_txq_array_size Int64Regs:$a),
2397           (TXQ_ARRAY_SIZE Int64Regs:$a)>;
2398 def : Pat<(int_nvvm_txq_num_samples Int64Regs:$a),
2399           (TXQ_NUM_SAMPLES Int64Regs:$a)>;
2400 def : Pat<(int_nvvm_txq_num_mipmap_levels Int64Regs:$a),
2401           (TXQ_NUM_MIPMAP_LEVELS Int64Regs:$a)>;
2402
2403
2404 //-----------------------------------
2405 // Surface Query Intrinsics
2406 //-----------------------------------
2407 def SUQ_CHANNEL_ORDER
2408   : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
2409               "suq.channel_order.b32 \t$d, [$a];",
2410               []>;
2411 def SUQ_CHANNEL_DATA_TYPE
2412   : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
2413               "suq.channel_data_type.b32 \t$d, [$a];",
2414               []>;
2415 def SUQ_WIDTH
2416   : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
2417               "suq.width.b32 \t$d, [$a];",
2418               []>;
2419 def SUQ_HEIGHT
2420   : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
2421               "suq.height.b32 \t$d, [$a];",
2422               []>;
2423 def SUQ_DEPTH
2424   : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
2425               "suq.depth.b32 \t$d, [$a];",
2426               []>;
2427 def SUQ_ARRAY_SIZE
2428   : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
2429               "suq.array_size.b32 \t$d, [$a];",
2430               []>;
2431
2432 def : Pat<(int_nvvm_suq_channel_order Int64Regs:$a),
2433           (SUQ_CHANNEL_ORDER Int64Regs:$a)>;
2434 def : Pat<(int_nvvm_suq_channel_data_type Int64Regs:$a),
2435           (SUQ_CHANNEL_DATA_TYPE Int64Regs:$a)>;
2436 def : Pat<(int_nvvm_suq_width Int64Regs:$a),
2437           (SUQ_WIDTH Int64Regs:$a)>;
2438 def : Pat<(int_nvvm_suq_height Int64Regs:$a),
2439           (SUQ_HEIGHT Int64Regs:$a)>;
2440 def : Pat<(int_nvvm_suq_depth Int64Regs:$a),
2441           (SUQ_DEPTH Int64Regs:$a)>;
2442 def : Pat<(int_nvvm_suq_array_size Int64Regs:$a),
2443           (SUQ_ARRAY_SIZE Int64Regs:$a)>;
2444
2445
2446 //===- Handle Query -------------------------------------------------------===//
2447
2448 // TODO: These intrinsics are not yet finalized, pending PTX ISA design work
2449 def ISTYPEP_SAMPLER
2450   : NVPTXInst<(outs Int1Regs:$d), (ins Int64Regs:$a),
2451               "istypep.samplerref \t$d, $a;",
2452               [(set Int1Regs:$d, (int_nvvm_istypep_sampler Int64Regs:$a))]>;
2453 def ISTYPEP_SURFACE
2454   : NVPTXInst<(outs Int1Regs:$d), (ins Int64Regs:$a),
2455               "istypep.surfref \t$d, $a;",
2456               [(set Int1Regs:$d, (int_nvvm_istypep_surface Int64Regs:$a))]>;
2457 def ISTYPEP_TEXTURE
2458   : NVPTXInst<(outs Int1Regs:$d), (ins Int64Regs:$a),
2459               "istypep.texref \t$d, $a;",
2460               [(set Int1Regs:$d, (int_nvvm_istypep_texture Int64Regs:$a))]>;
2461
2462 //===- Surface Stores -----------------------------------------------------===//
2463
2464 // Unformatted
2465
2466 def SUST_B_1D_B8_TRAP
2467   : NVPTXInst<(outs),
2468               (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r),
2469               "sust.b.1d.b8.trap \t[$s, \\{$x\\}], \\{$r\\};",
2470               []>;
2471 def SUST_B_1D_B16_TRAP
2472   : NVPTXInst<(outs),
2473               (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r),
2474               "sust.b.1d.b16.trap \t[$s, \\{$x\\}], \\{$r\\};",
2475               []>;
2476 def SUST_B_1D_B32_TRAP
2477   : NVPTXInst<(outs),
2478               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$r),
2479               "sust.b.1d.b32.trap \t[$s, \\{$x\\}], \\{$r\\};",
2480               []>;
2481 def SUST_B_1D_V2B8_TRAP
2482   : NVPTXInst<(outs),
2483               (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g),
2484               "sust.b.1d.v2.b8.trap \t[$s, \\{$x\\}], \\{$r, $g\\};",
2485               []>;
2486 def SUST_B_1D_V2B16_TRAP
2487   : NVPTXInst<(outs),
2488               (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g),
2489               "sust.b.1d.v2.b16.trap \t[$s, \\{$x\\}], \\{$r, $g\\};",
2490               []>;
2491 def SUST_B_1D_V2B32_TRAP
2492   : NVPTXInst<(outs),
2493               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$r, Int32Regs:$g),
2494               "sust.b.1d.v2.b32.trap \t[$s, \\{$x\\}], \\{$r, $g\\};",
2495               []>;
2496 def SUST_B_1D_V4B8_TRAP
2497   : NVPTXInst<(outs),
2498               (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g,
2499                    Int16Regs:$b, Int16Regs:$a),
2500               "sust.b.1d.v4.b8.trap \t[$s, \\{$x\\}], \\{$r, $g, $b, $a\\};",
2501               []>;
2502 def SUST_B_1D_V4B16_TRAP
2503   : NVPTXInst<(outs),
2504               (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g,
2505                    Int16Regs:$b, Int16Regs:$a),
2506               "sust.b.1d.v4.b16.trap \t[$s, \\{$x\\}], \\{$r, $g, $b, $a\\};",
2507               []>;
2508 def SUST_B_1D_V4B32_TRAP
2509   : NVPTXInst<(outs),
2510               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$r, Int32Regs:$g,
2511                    Int32Regs:$b, Int32Regs:$a),
2512               "sust.b.1d.v4.b32.trap \t[$s, \\{$x\\}], \\{$r, $g, $b, $a\\};",
2513               []>;
2514
2515
2516 def SUST_B_1D_ARRAY_B8_TRAP
2517   : NVPTXInst<(outs),
2518               (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int16Regs:$r),
2519               "sust.b.a1d.b8.trap \t[$s, \\{$idx, $x\\}], \\{$r\\};",
2520               []>;
2521 def SUST_B_1D_ARRAY_B16_TRAP
2522   : NVPTXInst<(outs),
2523               (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int16Regs:$r),
2524               "sust.b.a1d.b16.trap \t[$s, \\{$idx, $x\\}], \\{$r\\};",
2525               []>;
2526 def SUST_B_1D_ARRAY_B32_TRAP
2527   : NVPTXInst<(outs),
2528               (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$r),
2529               "sust.b.a1d.b32.trap \t[$s, \\{$idx, $x\\}], \\{$r\\};",
2530               []>;
2531 def SUST_B_1D_ARRAY_V2B8_TRAP
2532   : NVPTXInst<(outs),
2533               (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int16Regs:$r,
2534                    Int16Regs:$g),
2535               "sust.b.a1d.v2.b8.trap \t[$s, \\{$idx, $x\\}], \\{$r, $g\\};",
2536               []>;
2537 def SUST_B_1D_ARRAY_V2B16_TRAP
2538   : NVPTXInst<(outs),
2539               (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int16Regs:$r,
2540                    Int16Regs:$g),
2541               "sust.b.a1d.v2.b16.trap \t[$s, \\{$idx, $x\\}], \\{$r, $g\\};",
2542               []>;
2543 def SUST_B_1D_ARRAY_V2B32_TRAP
2544   : NVPTXInst<(outs),
2545               (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$r,
2546                    Int32Regs:$g),
2547               "sust.b.a1d.v2.b32.trap \t[$s, \\{$idx, $x\\}], \\{$r, $g\\};",
2548               []>;
2549 def SUST_B_1D_ARRAY_V4B8_TRAP
2550   : NVPTXInst<(outs),
2551               (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int16Regs:$r,
2552                    Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2553               "sust.b.a1d.v4.b8.trap \t[$s, \\{$idx, $x\\}], "
2554               "\\{$r, $g, $b, $a\\};",
2555               []>;
2556 def SUST_B_1D_ARRAY_V4B16_TRAP
2557   : NVPTXInst<(outs),
2558               (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int16Regs:$r,
2559                    Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2560              "sust.b.a1d.v4.b16.trap \t[$s, \\{$idx, $x\\}], "
2561              "\\{$r, $g, $b, $a\\};",
2562               []>;
2563 def SUST_B_1D_ARRAY_V4B32_TRAP
2564   : NVPTXInst<(outs),
2565               (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$r,
2566                    Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
2567              "sust.b.a1d.v4.b32.trap \t[$s, \\{$idx, $x\\}], "
2568              "\\{$r, $g, $b, $a\\};",
2569               []>;
2570
2571
2572 def SUST_B_2D_B8_TRAP
2573   : NVPTXInst<(outs),
2574               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r),
2575               "sust.b.2d.b8.trap \t[$s, \\{$x, $y\\}], \\{$r\\};",
2576               []>;
2577 def SUST_B_2D_B16_TRAP
2578   : NVPTXInst<(outs),
2579               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r),
2580               "sust.b.2d.b16.trap \t[$s, \\{$x, $y\\}], \\{$r\\};",
2581               []>;
2582 def SUST_B_2D_B32_TRAP
2583   : NVPTXInst<(outs),
2584               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r),
2585               "sust.b.2d.b32.trap \t[$s, \\{$x, $y\\}], \\{$r\\};",
2586               []>;
2587 def SUST_B_2D_V2B8_TRAP
2588   : NVPTXInst<(outs),
2589               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r,
2590                    Int16Regs:$g),
2591               "sust.b.2d.v2.b8.trap \t[$s, \\{$x, $y\\}], \\{$r, $g\\};",
2592               []>;
2593 def SUST_B_2D_V2B16_TRAP
2594   : NVPTXInst<(outs),
2595               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r,
2596                    Int16Regs:$g),
2597               "sust.b.2d.v2.b16.trap \t[$s, \\{$x, $y\\}], \\{$r, $g\\};",
2598               []>;
2599 def SUST_B_2D_V2B32_TRAP
2600   : NVPTXInst<(outs),
2601               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r,
2602                    Int32Regs:$g),
2603               "sust.b.2d.v2.b32.trap \t[$s, \\{$x, $y\\}], \\{$r, $g\\};",
2604               []>;
2605 def SUST_B_2D_V4B8_TRAP
2606   : NVPTXInst<(outs),
2607               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r,
2608                    Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2609               "sust.b.2d.v4.b8.trap \t[$s, \\{$x, $y\\}], "
2610               "\\{$r, $g, $b, $a\\};",
2611               []>;
2612 def SUST_B_2D_V4B16_TRAP
2613   : NVPTXInst<(outs),
2614               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r,
2615                    Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2616              "sust.b.2d.v4.b16.trap \t[$s, \\{$x, $y\\}], "
2617              "\\{$r, $g, $b, $a\\};",
2618               []>;
2619 def SUST_B_2D_V4B32_TRAP
2620   : NVPTXInst<(outs),
2621               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r,
2622                    Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
2623              "sust.b.2d.v4.b32.trap \t[$s, \\{$x, $y\\}], "
2624              "\\{$r, $g, $b, $a\\};",
2625               []>;
2626
2627
2628 def SUST_B_2D_ARRAY_B8_TRAP
2629   : NVPTXInst<(outs),
2630               (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
2631                    Int16Regs:$r),
2632               "sust.b.a2d.b8.trap \t[$s, \\{$idx, $x, $y, $y\\}], \\{$r\\};",
2633               []>;
2634 def SUST_B_2D_ARRAY_B16_TRAP
2635   : NVPTXInst<(outs),
2636               (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
2637                    Int16Regs:$r),
2638               "sust.b.a2d.b16.trap \t[$s, \\{$idx, $x, $y, $y\\}], \\{$r\\};",
2639               []>;
2640 def SUST_B_2D_ARRAY_B32_TRAP
2641   : NVPTXInst<(outs),
2642               (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
2643                    Int32Regs:$r),
2644               "sust.b.a2d.b32.trap \t[$s, \\{$idx, $x, $y, $y\\}], \\{$r\\};",
2645               []>;
2646 def SUST_B_2D_ARRAY_V2B8_TRAP
2647   : NVPTXInst<(outs),
2648               (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
2649                    Int16Regs:$r, Int16Regs:$g),
2650               "sust.b.a2d.v2.b8.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
2651               "\\{$r, $g\\};",
2652               []>;
2653 def SUST_B_2D_ARRAY_V2B16_TRAP
2654   : NVPTXInst<(outs),
2655               (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
2656                    Int16Regs:$r, Int16Regs:$g),
2657              "sust.b.a2d.v2.b16.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
2658              "\\{$r, $g\\};",
2659               []>;
2660 def SUST_B_2D_ARRAY_V2B32_TRAP
2661   : NVPTXInst<(outs),
2662               (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
2663                    Int32Regs:$r, Int32Regs:$g),
2664              "sust.b.a2d.v2.b32.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
2665              "\\{$r, $g\\};",
2666               []>;
2667 def SUST_B_2D_ARRAY_V4B8_TRAP
2668   : NVPTXInst<(outs),
2669               (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
2670                    Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2671       "sust.b.a2d.v4.b8.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
2672       "\\{$r, $g, $b, $a\\};",
2673               []>;
2674 def SUST_B_2D_ARRAY_V4B16_TRAP
2675   : NVPTXInst<(outs),
2676               (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
2677                    Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2678      "sust.b.a2d.v4.b16.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
2679      "\\{$r, $g, $b, $a\\};",
2680               []>;
2681 def SUST_B_2D_ARRAY_V4B32_TRAP
2682   : NVPTXInst<(outs),
2683               (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
2684                    Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
2685      "sust.b.a2d.v4.b32.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
2686      "\\{$r, $g, $b, $a\\};",
2687               []>;
2688
2689
2690 def SUST_B_3D_B8_TRAP
2691   : NVPTXInst<(outs),
2692               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
2693                    Int16Regs:$r),
2694               "sust.b.3d.b8.trap \t[$s, \\{$x, $y, $z, $z\\}], \\{$r\\};",
2695               []>;
2696 def SUST_B_3D_B16_TRAP
2697   : NVPTXInst<(outs),
2698               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
2699                    Int16Regs:$r),
2700               "sust.b.3d.b16.trap \t[$s, \\{$x, $y, $z, $z\\}], \\{$r\\};",
2701               []>;
2702 def SUST_B_3D_B32_TRAP
2703   : NVPTXInst<(outs),
2704               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
2705                    Int32Regs:$r),
2706               "sust.b.3d.b32.trap \t[$s, \\{$x, $y, $z, $z\\}], \\{$r\\};",
2707               []>;
2708 def SUST_B_3D_V2B8_TRAP
2709   : NVPTXInst<(outs),
2710               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
2711                    Int16Regs:$r, Int16Regs:$g),
2712               "sust.b.3d.v2.b8.trap \t[$s, \\{$x, $y, $z, $z\\}], "
2713               "\\{$r, $g\\};",
2714               []>;
2715 def SUST_B_3D_V2B16_TRAP
2716   : NVPTXInst<(outs),
2717               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
2718                    Int16Regs:$r, Int16Regs:$g),
2719               "sust.b.3d.v2.b16.trap \t[$s, \\{$x, $y, $z, $z\\}], "
2720               "\\{$r, $g\\};",
2721               []>;
2722 def SUST_B_3D_V2B32_TRAP
2723   : NVPTXInst<(outs),
2724               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
2725                    Int32Regs:$r, Int32Regs:$g),
2726               "sust.b.3d.v2.b32.trap \t[$s, \\{$x, $y, $z, $z\\}], "
2727               "\\{$r, $g\\};",
2728               []>;
2729 def SUST_B_3D_V4B8_TRAP
2730   : NVPTXInst<(outs),
2731               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
2732                    Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2733          "sust.b.3d.v4.b8.trap \t[$s, \\{$x, $y, $z, $z\\}], "
2734          "\\{$r, $g, $b, $a\\};",
2735               []>;
2736 def SUST_B_3D_V4B16_TRAP
2737   : NVPTXInst<(outs),
2738               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
2739                    Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2740         "sust.b.3d.v4.b16.trap \t[$s, \\{$x, $y, $z, $z\\}], "
2741         "\\{$r, $g, $b, $a\\};",
2742               []>;
2743 def SUST_B_3D_V4B32_TRAP
2744   : NVPTXInst<(outs),
2745               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
2746                    Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
2747         "sust.b.3d.v4.b32.trap \t[$s, \\{$x, $y, $z, $z\\}], "
2748         "\\{$r, $g, $b, $a\\};",
2749               []>;
2750
2751 // Formatted
2752
2753 def SUST_P_1D_B8_TRAP
2754   : NVPTXInst<(outs),
2755               (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r),
2756               "sust.p.1d.b8.trap \t[$s, \\{$x\\}], \\{$r\\};",
2757               []>;
2758 def SUST_P_1D_B16_TRAP
2759   : NVPTXInst<(outs),
2760               (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r),
2761               "sust.p.1d.b16.trap \t[$s, \\{$x\\}], \\{$r\\};",
2762               []>;
2763 def SUST_P_1D_B32_TRAP
2764   : NVPTXInst<(outs),
2765               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$r),
2766               "sust.p.1d.b32.trap \t[$s, \\{$x\\}], \\{$r\\};",
2767               []>;
2768 def SUST_P_1D_V2B8_TRAP
2769   : NVPTXInst<(outs),
2770               (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g),
2771               "sust.p.1d.v2.b8.trap \t[$s, \\{$x\\}], \\{$r, $g\\};",
2772               []>;
2773 def SUST_P_1D_V2B16_TRAP
2774   : NVPTXInst<(outs),
2775               (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g),
2776               "sust.p.1d.v2.b16.trap \t[$s, \\{$x\\}], \\{$r, $g\\};",
2777               []>;
2778 def SUST_P_1D_V2B32_TRAP
2779   : NVPTXInst<(outs),
2780               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$r, Int32Regs:$g),
2781               "sust.p.1d.v2.b32.trap \t[$s, \\{$x\\}], \\{$r, $g\\};",
2782               []>;
2783 def SUST_P_1D_V4B8_TRAP
2784   : NVPTXInst<(outs),
2785               (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g,
2786                    Int16Regs:$b, Int16Regs:$a),
2787               "sust.p.1d.v4.b8.trap \t[$s, \\{$x\\}], \\{$r, $g, $b, $a\\};",
2788               []>;
2789 def SUST_P_1D_V4B16_TRAP
2790   : NVPTXInst<(outs),
2791               (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g,
2792                    Int16Regs:$b, Int16Regs:$a),
2793               "sust.p.1d.v4.b16.trap \t[$s, \\{$x\\}], \\{$r, $g, $b, $a\\};",
2794               []>;
2795 def SUST_P_1D_V4B32_TRAP
2796   : NVPTXInst<(outs),
2797               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$r, Int32Regs:$g,
2798                    Int32Regs:$b, Int32Regs:$a),
2799               "sust.p.1d.v4.b32.trap \t[$s, \\{$x\\}], \\{$r, $g, $b, $a\\};",
2800               []>;
2801
2802
2803 def SUST_P_1D_ARRAY_B8_TRAP
2804   : NVPTXInst<(outs),
2805               (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int16Regs:$r),
2806               "sust.p.a1d.b8.trap \t[$s, \\{$idx, $x\\}], \\{$r\\};",
2807               []>;
2808 def SUST_P_1D_ARRAY_B16_TRAP
2809   : NVPTXInst<(outs),
2810               (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int16Regs:$r),
2811               "sust.p.a1d.b16.trap \t[$s, \\{$idx, $x\\}], \\{$r\\};",
2812               []>;
2813 def SUST_P_1D_ARRAY_B32_TRAP
2814   : NVPTXInst<(outs),
2815               (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$r),
2816               "sust.p.a1d.b32.trap \t[$s, \\{$idx, $x\\}], \\{$r\\};",
2817               []>;
2818 def SUST_P_1D_ARRAY_V2B8_TRAP
2819   : NVPTXInst<(outs),
2820               (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int16Regs:$r,
2821                    Int16Regs:$g),
2822               "sust.p.a1d.v2.b8.trap \t[$s, \\{$idx, $x\\}], \\{$r, $g\\};",
2823               []>;
2824 def SUST_P_1D_ARRAY_V2B16_TRAP
2825   : NVPTXInst<(outs),
2826               (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int16Regs:$r,
2827                    Int16Regs:$g),
2828               "sust.p.a1d.v2.b16.trap \t[$s, \\{$idx, $x\\}], \\{$r, $g\\};",
2829               []>;
2830 def SUST_P_1D_ARRAY_V2B32_TRAP
2831   : NVPTXInst<(outs),
2832               (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$r,
2833                    Int32Regs:$g),
2834               "sust.p.a1d.v2.b32.trap \t[$s, \\{$idx, $x\\}], \\{$r, $g\\};",
2835               []>;
2836 def SUST_P_1D_ARRAY_V4B8_TRAP
2837   : NVPTXInst<(outs),
2838               (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int16Regs:$r,
2839                    Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2840               "sust.p.a1d.v4.b8.trap \t[$s, \\{$idx, $x\\}], "
2841               "\\{$r, $g, $b, $a\\};",
2842               []>;
2843 def SUST_P_1D_ARRAY_V4B16_TRAP
2844   : NVPTXInst<(outs),
2845               (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int16Regs:$r,
2846                    Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2847              "sust.p.a1d.v4.b16.trap \t[$s, \\{$idx, $x\\}], "
2848              "\\{$r, $g, $b, $a\\};",
2849               []>;
2850 def SUST_P_1D_ARRAY_V4B32_TRAP
2851   : NVPTXInst<(outs),
2852               (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$r,
2853                    Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
2854              "sust.p.a1d.v4.b32.trap \t[$s, \\{$idx, $x\\}], "
2855              "\\{$r, $g, $b, $a\\};",
2856               []>;
2857
2858
2859 def SUST_P_2D_B8_TRAP
2860   : NVPTXInst<(outs),
2861               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r),
2862               "sust.p.2d.b8.trap \t[$s, \\{$x, $y\\}], \\{$r\\};",
2863               []>;
2864 def SUST_P_2D_B16_TRAP
2865   : NVPTXInst<(outs),
2866               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r),
2867               "sust.p.2d.b16.trap \t[$s, \\{$x, $y\\}], \\{$r\\};",
2868               []>;
2869 def SUST_P_2D_B32_TRAP
2870   : NVPTXInst<(outs),
2871               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r),
2872               "sust.p.2d.b32.trap \t[$s, \\{$x, $y\\}], \\{$r\\};",
2873               []>;
2874 def SUST_P_2D_V2B8_TRAP
2875   : NVPTXInst<(outs),
2876               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r,
2877                    Int16Regs:$g),
2878               "sust.p.2d.v2.b8.trap \t[$s, \\{$x, $y\\}], \\{$r, $g\\};",
2879               []>;
2880 def SUST_P_2D_V2B16_TRAP
2881   : NVPTXInst<(outs),
2882               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r,
2883                    Int16Regs:$g),
2884               "sust.p.2d.v2.b16.trap \t[$s, \\{$x, $y\\}], \\{$r, $g\\};",
2885               []>;
2886 def SUST_P_2D_V2B32_TRAP
2887   : NVPTXInst<(outs),
2888               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r,
2889                    Int32Regs:$g),
2890               "sust.p.2d.v2.b32.trap \t[$s, \\{$x, $y\\}], \\{$r, $g\\};",
2891               []>;
2892 def SUST_P_2D_V4B8_TRAP
2893   : NVPTXInst<(outs),
2894               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r,
2895                    Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2896               "sust.p.2d.v4.b8.trap \t[$s, \\{$x, $y\\}], "
2897               "\\{$r, $g, $b, $a\\};",
2898               []>;
2899 def SUST_P_2D_V4B16_TRAP
2900   : NVPTXInst<(outs),
2901               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r,
2902                    Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2903              "sust.p.2d.v4.b16.trap \t[$s, \\{$x, $y\\}], "
2904              "\\{$r, $g, $b, $a\\};",
2905               []>;
2906 def SUST_P_2D_V4B32_TRAP
2907   : NVPTXInst<(outs),
2908               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r,
2909                    Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
2910              "sust.p.2d.v4.b32.trap \t[$s, \\{$x, $y\\}], "
2911              "\\{$r, $g, $b, $a\\};",
2912               []>;
2913
2914
2915 def SUST_P_2D_ARRAY_B8_TRAP
2916   : NVPTXInst<(outs),
2917               (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
2918                    Int16Regs:$r),
2919               "sust.p.a2d.b8.trap \t[$s, \\{$idx, $x, $y, $y\\}], \\{$r\\};",
2920               []>;
2921 def SUST_P_2D_ARRAY_B16_TRAP
2922   : NVPTXInst<(outs),
2923               (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
2924                    Int16Regs:$r),
2925               "sust.p.a2d.b16.trap \t[$s, \\{$idx, $x, $y, $y\\}], \\{$r\\};",
2926               []>;
2927 def SUST_P_2D_ARRAY_B32_TRAP
2928   : NVPTXInst<(outs),
2929               (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
2930                    Int32Regs:$r),
2931               "sust.p.a2d.b32.trap \t[$s, \\{$idx, $x, $y, $y\\}], \\{$r\\};",
2932               []>;
2933 def SUST_P_2D_ARRAY_V2B8_TRAP
2934   : NVPTXInst<(outs),
2935               (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
2936                    Int16Regs:$r, Int16Regs:$g),
2937               "sust.p.a2d.v2.b8.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
2938               "\\{$r, $g\\};",
2939               []>;
2940 def SUST_P_2D_ARRAY_V2B16_TRAP
2941   : NVPTXInst<(outs),
2942               (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
2943                    Int16Regs:$r, Int16Regs:$g),
2944              "sust.p.a2d.v2.b16.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
2945              "\\{$r, $g\\};",
2946               []>;
2947 def SUST_P_2D_ARRAY_V2B32_TRAP
2948   : NVPTXInst<(outs),
2949               (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
2950                    Int32Regs:$r, Int32Regs:$g),
2951              "sust.p.a2d.v2.b32.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
2952              "\\{$r, $g\\};",
2953               []>;
2954 def SUST_P_2D_ARRAY_V4B8_TRAP
2955   : NVPTXInst<(outs),
2956               (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
2957                    Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2958       "sust.p.a2d.v4.b8.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
2959       "\\{$r, $g, $b, $a\\};",
2960               []>;
2961 def SUST_P_2D_ARRAY_V4B16_TRAP
2962   : NVPTXInst<(outs),
2963               (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
2964                    Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2965      "sust.p.a2d.v4.b16.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
2966      "\\{$r, $g, $b, $a\\};",
2967               []>;
2968 def SUST_P_2D_ARRAY_V4B32_TRAP
2969   : NVPTXInst<(outs),
2970               (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
2971                    Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
2972      "sust.p.a2d.v4.b32.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
2973      "\\{$r, $g, $b, $a\\};",
2974               []>;
2975
2976
2977 def SUST_P_3D_B8_TRAP
2978   : NVPTXInst<(outs),
2979               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
2980                    Int16Regs:$r),
2981               "sust.p.3d.b8.trap \t[$s, \\{$x, $y, $z, $z\\}], \\{$r\\};",
2982               []>;
2983 def SUST_P_3D_B16_TRAP
2984   : NVPTXInst<(outs),
2985               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
2986                    Int16Regs:$r),
2987               "sust.p.3d.b16.trap \t[$s, \\{$x, $y, $z, $z\\}], \\{$r\\};",
2988               []>;
2989 def SUST_P_3D_B32_TRAP
2990   : NVPTXInst<(outs),
2991               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
2992                    Int32Regs:$r),
2993               "sust.p.3d.b32.trap \t[$s, \\{$x, $y, $z, $z\\}], \\{$r\\};",
2994               []>;
2995 def SUST_P_3D_V2B8_TRAP
2996   : NVPTXInst<(outs),
2997               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
2998                    Int16Regs:$r, Int16Regs:$g),
2999               "sust.p.3d.v2.b8.trap \t[$s, \\{$x, $y, $z, $z\\}], "
3000               "\\{$r, $g\\};",
3001               []>;
3002 def SUST_P_3D_V2B16_TRAP
3003   : NVPTXInst<(outs),
3004               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3005                    Int16Regs:$r, Int16Regs:$g),
3006               "sust.p.3d.v2.b16.trap \t[$s, \\{$x, $y, $z, $z\\}], "
3007               "\\{$r, $g\\};",
3008               []>;
3009 def SUST_P_3D_V2B32_TRAP
3010   : NVPTXInst<(outs),
3011               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3012                    Int32Regs:$r, Int32Regs:$g),
3013               "sust.p.3d.v2.b32.trap \t[$s, \\{$x, $y, $z, $z\\}], "
3014               "\\{$r, $g\\};",
3015               []>;
3016 def SUST_P_3D_V4B8_TRAP
3017   : NVPTXInst<(outs),
3018               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3019                    Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3020          "sust.p.3d.v4.b8.trap \t[$s, \\{$x, $y, $z, $z\\}], "
3021          "\\{$r, $g, $b, $a\\};",
3022               []>;
3023 def SUST_P_3D_V4B16_TRAP
3024   : NVPTXInst<(outs),
3025               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3026                    Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3027         "sust.p.3d.v4.b16.trap \t[$s, \\{$x, $y, $z, $z\\}], "
3028         "\\{$r, $g, $b, $a\\};",
3029               []>;
3030 def SUST_P_3D_V4B32_TRAP
3031   : NVPTXInst<(outs),
3032               (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3033                    Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
3034         "sust.p.3d.v4.b32.trap \t[$s, \\{$x, $y, $z, $z\\}], "
3035         "\\{$r, $g, $b, $a\\};",
3036               []>;
3037
3038
3039 // Surface store instruction patterns
3040 // I'm not sure why we can't just include these in the instruction definitions,
3041 // but TableGen complains of type errors :(
3042
3043 def : Pat<(int_nvvm_sust_b_1d_i8_trap
3044            Int64Regs:$s, Int32Regs:$x, Int16Regs:$r),
3045           (SUST_B_1D_B8_TRAP Int64Regs:$s, Int32Regs:$x, Int16Regs:$r)>;
3046
3047 def : Pat<(int_nvvm_sust_b_1d_i16_trap
3048            Int64Regs:$s, Int32Regs:$x, Int16Regs:$r),
3049           (SUST_B_1D_B16_TRAP Int64Regs:$s, Int32Regs:$x, Int16Regs:$r)>;
3050
3051 def : Pat<(int_nvvm_sust_b_1d_i32_trap
3052            Int64Regs:$s, Int32Regs:$x, Int32Regs:$r),
3053           (SUST_B_1D_B32_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$r)>;
3054
3055 def : Pat<(int_nvvm_sust_b_1d_v2i8_trap
3056            Int64Regs:$s, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g),
3057           (SUST_B_1D_V2B8_TRAP Int64Regs:$s, Int32Regs:$x,
3058            Int16Regs:$r, Int16Regs:$g)>;
3059
3060 def : Pat<(int_nvvm_sust_b_1d_v2i16_trap
3061            Int64Regs:$s, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g),
3062           (SUST_B_1D_V2B16_TRAP Int64Regs:$s, Int32Regs:$x,
3063            Int16Regs:$r, Int16Regs:$g)>;
3064
3065 def : Pat<(int_nvvm_sust_b_1d_v2i32_trap
3066            Int64Regs:$s, Int32Regs:$x, Int32Regs:$r, Int32Regs:$g),
3067           (SUST_B_1D_V2B32_TRAP Int64Regs:$s, Int32Regs:$x,
3068            Int32Regs:$r, Int32Regs:$g)>;
3069
3070 def : Pat<(int_nvvm_sust_b_1d_v4i8_trap
3071            Int64Regs:$s, Int32Regs:$x,
3072            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3073           (SUST_B_1D_V4B8_TRAP Int64Regs:$s, Int32Regs:$x,
3074            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3075
3076 def : Pat<(int_nvvm_sust_b_1d_v4i16_trap
3077            Int64Regs:$s, Int32Regs:$x,
3078            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3079           (SUST_B_1D_V4B16_TRAP Int64Regs:$s, Int32Regs:$x,
3080            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3081
3082 def : Pat<(int_nvvm_sust_b_1d_v4i32_trap
3083            Int64Regs:$s, Int32Regs:$x,
3084            Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
3085           (SUST_B_1D_V4B32_TRAP Int64Regs:$s, Int32Regs:$x,
3086            Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a)>;
3087
3088
3089
3090 def : Pat<(int_nvvm_sust_b_1d_array_i8_trap
3091            Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int16Regs:$r),
3092           (SUST_B_1D_ARRAY_B8_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3093            Int16Regs:$r)>;
3094
3095 def : Pat<(int_nvvm_sust_b_1d_array_i16_trap
3096            Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int16Regs:$r),
3097           (SUST_B_1D_ARRAY_B16_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3098            Int16Regs:$r)>;
3099
3100 def : Pat<(int_nvvm_sust_b_1d_array_i32_trap
3101            Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$r),
3102           (SUST_B_1D_ARRAY_B32_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3103            Int32Regs:$r)>;
3104
3105 def : Pat<(int_nvvm_sust_b_1d_array_v2i8_trap
3106           Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g),
3107           (SUST_B_1D_ARRAY_V2B8_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3108            Int16Regs:$r, Int16Regs:$g)>;
3109
3110 def : Pat<(int_nvvm_sust_b_1d_array_v2i16_trap
3111           Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g),
3112           (SUST_B_1D_ARRAY_V2B16_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3113            Int16Regs:$r, Int16Regs:$g)>;
3114
3115 def : Pat<(int_nvvm_sust_b_1d_array_v2i32_trap
3116           Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$r, Int32Regs:$g),
3117           (SUST_B_1D_ARRAY_V2B32_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3118            Int32Regs:$r, Int32Regs:$g)>;
3119
3120 def : Pat<(int_nvvm_sust_b_1d_array_v4i8_trap
3121            Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3122            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3123           (SUST_B_1D_ARRAY_V4B8_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3124            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3125
3126 def : Pat<(int_nvvm_sust_b_1d_array_v4i16_trap
3127            Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3128            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3129           (SUST_B_1D_ARRAY_V4B16_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3130            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3131
3132 def : Pat<(int_nvvm_sust_b_1d_array_v4i32_trap
3133            Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3134            Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
3135           (SUST_B_1D_ARRAY_V4B32_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3136            Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a)>;
3137
3138
3139
3140 def : Pat<(int_nvvm_sust_b_2d_i8_trap
3141            Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r),
3142           (SUST_B_2D_B8_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3143            Int16Regs:$r)>;
3144
3145 def : Pat<(int_nvvm_sust_b_2d_i16_trap
3146            Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r),
3147           (SUST_B_2D_B16_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3148            Int16Regs:$r)>;
3149
3150 def : Pat<(int_nvvm_sust_b_2d_i32_trap
3151            Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r),
3152           (SUST_B_2D_B32_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3153            Int32Regs:$r)>;
3154
3155 def : Pat<(int_nvvm_sust_b_2d_v2i8_trap
3156           Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r, Int16Regs:$g),
3157           (SUST_B_2D_V2B8_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3158            Int16Regs:$r, Int16Regs:$g)>;
3159
3160 def : Pat<(int_nvvm_sust_b_2d_v2i16_trap
3161           Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r, Int16Regs:$g),
3162           (SUST_B_2D_V2B16_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3163            Int16Regs:$r, Int16Regs:$g)>;
3164
3165 def : Pat<(int_nvvm_sust_b_2d_v2i32_trap
3166           Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r, Int32Regs:$g),
3167           (SUST_B_2D_V2B32_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3168            Int32Regs:$r, Int32Regs:$g)>;
3169
3170 def : Pat<(int_nvvm_sust_b_2d_v4i8_trap
3171            Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3172            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3173           (SUST_B_2D_V4B8_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3174            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3175
3176 def : Pat<(int_nvvm_sust_b_2d_v4i16_trap
3177            Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3178            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3179           (SUST_B_2D_V4B16_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3180            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3181
3182 def : Pat<(int_nvvm_sust_b_2d_v4i32_trap
3183            Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3184            Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
3185           (SUST_B_2D_V4B32_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3186            Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a)>;
3187
3188
3189
3190 def : Pat<(int_nvvm_sust_b_2d_array_i8_trap
3191           Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r),
3192           (SUST_B_2D_ARRAY_B8_TRAP Int64Regs:$s,
3193            Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3194            Int16Regs:$r)>;
3195
3196 def : Pat<(int_nvvm_sust_b_2d_array_i16_trap
3197           Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r),
3198           (SUST_B_2D_ARRAY_B16_TRAP Int64Regs:$s,
3199            Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3200            Int16Regs:$r)>;
3201
3202 def : Pat<(int_nvvm_sust_b_2d_array_i32_trap
3203           Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r),
3204           (SUST_B_2D_ARRAY_B32_TRAP Int64Regs:$s,
3205            Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3206            Int32Regs:$r)>;
3207
3208 def : Pat<(int_nvvm_sust_b_2d_array_v2i8_trap
3209            Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3210            Int16Regs:$r, Int16Regs:$g),
3211           (SUST_B_2D_ARRAY_V2B8_TRAP Int64Regs:$s, Int32Regs:$l,
3212            Int32Regs:$x, Int32Regs:$y,
3213            Int16Regs:$r, Int16Regs:$g)>;
3214
3215 def : Pat<(int_nvvm_sust_b_2d_array_v2i16_trap
3216            Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3217            Int16Regs:$r, Int16Regs:$g),
3218           (SUST_B_2D_ARRAY_V2B16_TRAP Int64Regs:$s, Int32Regs:$l,
3219            Int32Regs:$x, Int32Regs:$y,
3220            Int16Regs:$r, Int16Regs:$g)>;
3221
3222 def : Pat<(int_nvvm_sust_b_2d_array_v2i32_trap
3223            Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r,
3224            Int32Regs:$g),
3225           (SUST_B_2D_ARRAY_V2B32_TRAP Int64Regs:$s, Int32Regs:$l,
3226            Int32Regs:$x, Int32Regs:$y, Int32Regs:$r, Int32Regs:$g)>;
3227
3228 def : Pat<(int_nvvm_sust_b_2d_array_v4i8_trap
3229            Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3230            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3231           (SUST_B_2D_ARRAY_V4B8_TRAP Int64Regs:$s,
3232            Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3233            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3234
3235 def : Pat<(int_nvvm_sust_b_2d_array_v4i16_trap
3236            Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3237            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3238           (SUST_B_2D_ARRAY_V4B16_TRAP Int64Regs:$s,
3239            Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3240            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3241
3242 def : Pat<(int_nvvm_sust_b_2d_array_v4i32_trap
3243            Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3244            Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
3245           (SUST_B_2D_ARRAY_V4B32_TRAP Int64Regs:$s, Int32Regs:$l,
3246            Int32Regs:$x, Int32Regs:$y,
3247            Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a)>;
3248
3249
3250
3251 def : Pat<(int_nvvm_sust_b_3d_i8_trap
3252            Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3253            Int16Regs:$r),
3254           (SUST_B_3D_B8_TRAP Int64Regs:$s,
3255            Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3256            Int16Regs:$r)>;
3257
3258 def : Pat<(int_nvvm_sust_b_3d_i16_trap
3259            Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3260            Int16Regs:$r),
3261           (SUST_B_3D_B16_TRAP Int64Regs:$s,
3262            Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3263            Int16Regs:$r)>;
3264
3265 def : Pat<(int_nvvm_sust_b_3d_i32_trap
3266            Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3267            Int32Regs:$r),
3268           (SUST_B_3D_B32_TRAP Int64Regs:$s,
3269            Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3270            Int32Regs:$r)>;
3271
3272 def : Pat<(int_nvvm_sust_b_3d_v2i8_trap
3273            Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3274            Int16Regs:$r, Int16Regs:$g),
3275           (SUST_B_3D_V2B8_TRAP Int64Regs:$s,
3276            Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3277            Int16Regs:$r, Int16Regs:$g)>;
3278
3279 def : Pat<(int_nvvm_sust_b_3d_v2i16_trap
3280            Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3281            Int16Regs:$r, Int16Regs:$g),
3282           (SUST_B_3D_V2B16_TRAP Int64Regs:$s,
3283            Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3284            Int16Regs:$r, Int16Regs:$g)>;
3285
3286 def : Pat<(int_nvvm_sust_b_3d_v2i32_trap
3287            Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3288            Int32Regs:$r, Int32Regs:$g),
3289           (SUST_B_3D_V2B32_TRAP Int64Regs:$s,
3290            Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3291            Int32Regs:$r, Int32Regs:$g)>;
3292
3293 def : Pat<(int_nvvm_sust_b_3d_v4i8_trap
3294            Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3295            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3296           (SUST_B_3D_V4B8_TRAP Int64Regs:$s,
3297            Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3298            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3299
3300 def : Pat<(int_nvvm_sust_b_3d_v4i16_trap
3301            Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3302            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3303           (SUST_B_3D_V4B16_TRAP Int64Regs:$s,
3304            Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3305            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3306
3307 def : Pat<(int_nvvm_sust_b_3d_v4i32_trap
3308            Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3309            Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
3310           (SUST_B_3D_V4B32_TRAP Int64Regs:$s,
3311            Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3312            Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a)>;
3313
3314
3315
3316
3317 def : Pat<(int_nvvm_sust_p_1d_i8_trap
3318            Int64Regs:$s, Int32Regs:$x, Int16Regs:$r),
3319           (SUST_P_1D_B8_TRAP Int64Regs:$s, Int32Regs:$x, Int16Regs:$r)>;
3320
3321 def : Pat<(int_nvvm_sust_p_1d_i16_trap
3322            Int64Regs:$s, Int32Regs:$x, Int16Regs:$r),
3323           (SUST_P_1D_B16_TRAP Int64Regs:$s, Int32Regs:$x, Int16Regs:$r)>;
3324
3325 def : Pat<(int_nvvm_sust_p_1d_i32_trap
3326            Int64Regs:$s, Int32Regs:$x, Int32Regs:$r),
3327           (SUST_P_1D_B32_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$r)>;
3328
3329 def : Pat<(int_nvvm_sust_p_1d_v2i8_trap
3330            Int64Regs:$s, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g),
3331           (SUST_P_1D_V2B8_TRAP Int64Regs:$s, Int32Regs:$x,
3332            Int16Regs:$r, Int16Regs:$g)>;
3333
3334 def : Pat<(int_nvvm_sust_p_1d_v2i16_trap
3335            Int64Regs:$s, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g),
3336           (SUST_P_1D_V2B16_TRAP Int64Regs:$s, Int32Regs:$x,
3337            Int16Regs:$r, Int16Regs:$g)>;
3338
3339 def : Pat<(int_nvvm_sust_p_1d_v2i32_trap
3340            Int64Regs:$s, Int32Regs:$x, Int32Regs:$r, Int32Regs:$g),
3341           (SUST_P_1D_V2B32_TRAP Int64Regs:$s, Int32Regs:$x,
3342            Int32Regs:$r, Int32Regs:$g)>;
3343
3344 def : Pat<(int_nvvm_sust_p_1d_v4i8_trap
3345            Int64Regs:$s, Int32Regs:$x,
3346            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3347           (SUST_P_1D_V4B8_TRAP Int64Regs:$s, Int32Regs:$x,
3348            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3349
3350 def : Pat<(int_nvvm_sust_p_1d_v4i16_trap
3351            Int64Regs:$s, Int32Regs:$x,
3352            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3353           (SUST_P_1D_V4B16_TRAP Int64Regs:$s, Int32Regs:$x,
3354            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3355
3356 def : Pat<(int_nvvm_sust_p_1d_v4i32_trap
3357            Int64Regs:$s, Int32Regs:$x,
3358            Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
3359           (SUST_P_1D_V4B32_TRAP Int64Regs:$s, Int32Regs:$x,
3360            Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a)>;
3361
3362
3363
3364 def : Pat<(int_nvvm_sust_p_1d_array_i8_trap
3365            Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int16Regs:$r),
3366           (SUST_P_1D_ARRAY_B8_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3367            Int16Regs:$r)>;
3368
3369 def : Pat<(int_nvvm_sust_p_1d_array_i16_trap
3370            Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int16Regs:$r),
3371           (SUST_P_1D_ARRAY_B16_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3372            Int16Regs:$r)>;
3373
3374 def : Pat<(int_nvvm_sust_p_1d_array_i32_trap
3375            Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$r),
3376           (SUST_P_1D_ARRAY_B32_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3377            Int32Regs:$r)>;
3378
3379 def : Pat<(int_nvvm_sust_p_1d_array_v2i8_trap
3380           Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g),
3381           (SUST_P_1D_ARRAY_V2B8_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3382            Int16Regs:$r, Int16Regs:$g)>;
3383
3384 def : Pat<(int_nvvm_sust_p_1d_array_v2i16_trap
3385           Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g),
3386           (SUST_P_1D_ARRAY_V2B16_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3387            Int16Regs:$r, Int16Regs:$g)>;
3388
3389 def : Pat<(int_nvvm_sust_p_1d_array_v2i32_trap
3390           Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$r, Int32Regs:$g),
3391           (SUST_P_1D_ARRAY_V2B32_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3392            Int32Regs:$r, Int32Regs:$g)>;
3393
3394 def : Pat<(int_nvvm_sust_p_1d_array_v4i8_trap
3395            Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3396            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3397           (SUST_P_1D_ARRAY_V4B8_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3398            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3399
3400 def : Pat<(int_nvvm_sust_p_1d_array_v4i16_trap
3401            Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3402            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3403           (SUST_P_1D_ARRAY_V4B16_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3404            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3405
3406 def : Pat<(int_nvvm_sust_p_1d_array_v4i32_trap
3407            Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3408            Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
3409           (SUST_P_1D_ARRAY_V4B32_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3410            Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a)>;
3411
3412
3413
3414 def : Pat<(int_nvvm_sust_p_2d_i8_trap
3415            Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r),
3416           (SUST_P_2D_B8_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3417            Int16Regs:$r)>;
3418
3419 def : Pat<(int_nvvm_sust_p_2d_i16_trap
3420            Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r),
3421           (SUST_P_2D_B16_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3422            Int16Regs:$r)>;
3423
3424 def : Pat<(int_nvvm_sust_p_2d_i32_trap
3425            Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r),
3426           (SUST_P_2D_B32_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3427            Int32Regs:$r)>;
3428
3429 def : Pat<(int_nvvm_sust_p_2d_v2i8_trap
3430           Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r, Int16Regs:$g),
3431           (SUST_P_2D_V2B8_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3432            Int16Regs:$r, Int16Regs:$g)>;
3433
3434 def : Pat<(int_nvvm_sust_p_2d_v2i16_trap
3435           Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r, Int16Regs:$g),
3436           (SUST_P_2D_V2B16_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3437            Int16Regs:$r, Int16Regs:$g)>;
3438
3439 def : Pat<(int_nvvm_sust_p_2d_v2i32_trap
3440           Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r, Int32Regs:$g),
3441           (SUST_P_2D_V2B32_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3442            Int32Regs:$r, Int32Regs:$g)>;
3443
3444 def : Pat<(int_nvvm_sust_p_2d_v4i8_trap
3445            Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3446            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3447           (SUST_P_2D_V4B8_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3448            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3449
3450 def : Pat<(int_nvvm_sust_p_2d_v4i16_trap
3451            Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3452            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3453           (SUST_P_2D_V4B16_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3454            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3455
3456 def : Pat<(int_nvvm_sust_p_2d_v4i32_trap
3457            Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3458            Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
3459           (SUST_P_2D_V4B32_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3460            Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a)>;
3461
3462
3463
3464 def : Pat<(int_nvvm_sust_p_2d_array_i8_trap
3465           Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r),
3466           (SUST_P_2D_ARRAY_B8_TRAP Int64Regs:$s,
3467            Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3468            Int16Regs:$r)>;
3469
3470 def : Pat<(int_nvvm_sust_p_2d_array_i16_trap
3471           Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r),
3472           (SUST_P_2D_ARRAY_B16_TRAP Int64Regs:$s,
3473            Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3474            Int16Regs:$r)>;
3475
3476 def : Pat<(int_nvvm_sust_p_2d_array_i32_trap
3477           Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r),
3478           (SUST_P_2D_ARRAY_B32_TRAP Int64Regs:$s,
3479            Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3480            Int32Regs:$r)>;
3481
3482 def : Pat<(int_nvvm_sust_p_2d_array_v2i8_trap
3483            Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3484            Int16Regs:$r, Int16Regs:$g),
3485           (SUST_P_2D_ARRAY_V2B8_TRAP Int64Regs:$s, Int32Regs:$l,
3486            Int32Regs:$x, Int32Regs:$y,
3487            Int16Regs:$r, Int16Regs:$g)>;
3488
3489 def : Pat<(int_nvvm_sust_p_2d_array_v2i16_trap
3490            Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3491            Int16Regs:$r, Int16Regs:$g),
3492           (SUST_P_2D_ARRAY_V2B16_TRAP Int64Regs:$s, Int32Regs:$l,
3493            Int32Regs:$x, Int32Regs:$y,
3494            Int16Regs:$r, Int16Regs:$g)>;
3495
3496 def : Pat<(int_nvvm_sust_p_2d_array_v2i32_trap
3497            Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r,
3498            Int32Regs:$g),
3499           (SUST_P_2D_ARRAY_V2B32_TRAP Int64Regs:$s, Int32Regs:$l,
3500            Int32Regs:$x, Int32Regs:$y, Int32Regs:$r, Int32Regs:$g)>;
3501
3502 def : Pat<(int_nvvm_sust_p_2d_array_v4i8_trap
3503            Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3504            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3505           (SUST_P_2D_ARRAY_V4B8_TRAP Int64Regs:$s,
3506            Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3507            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3508
3509 def : Pat<(int_nvvm_sust_p_2d_array_v4i16_trap
3510            Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3511            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3512           (SUST_P_2D_ARRAY_V4B16_TRAP Int64Regs:$s,
3513            Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3514            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3515
3516 def : Pat<(int_nvvm_sust_p_2d_array_v4i32_trap
3517            Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3518            Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
3519           (SUST_P_2D_ARRAY_V4B32_TRAP Int64Regs:$s, Int32Regs:$l,
3520            Int32Regs:$x, Int32Regs:$y,
3521            Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a)>;
3522
3523
3524
3525 def : Pat<(int_nvvm_sust_p_3d_i8_trap
3526            Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3527            Int16Regs:$r),
3528           (SUST_P_3D_B8_TRAP Int64Regs:$s,
3529            Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3530            Int16Regs:$r)>;
3531
3532 def : Pat<(int_nvvm_sust_p_3d_i16_trap
3533            Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3534            Int16Regs:$r),
3535           (SUST_P_3D_B16_TRAP Int64Regs:$s,
3536            Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3537            Int16Regs:$r)>;
3538
3539 def : Pat<(int_nvvm_sust_p_3d_i32_trap
3540            Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3541            Int32Regs:$r),
3542           (SUST_P_3D_B32_TRAP Int64Regs:$s,
3543            Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3544            Int32Regs:$r)>;
3545
3546 def : Pat<(int_nvvm_sust_p_3d_v2i8_trap
3547            Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3548            Int16Regs:$r, Int16Regs:$g),
3549           (SUST_P_3D_V2B8_TRAP Int64Regs:$s,
3550            Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3551            Int16Regs:$r, Int16Regs:$g)>;
3552
3553 def : Pat<(int_nvvm_sust_p_3d_v2i16_trap
3554            Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3555            Int16Regs:$r, Int16Regs:$g),
3556           (SUST_P_3D_V2B16_TRAP Int64Regs:$s,
3557            Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3558            Int16Regs:$r, Int16Regs:$g)>;
3559
3560 def : Pat<(int_nvvm_sust_p_3d_v2i32_trap
3561            Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3562            Int32Regs:$r, Int32Regs:$g),
3563           (SUST_P_3D_V2B32_TRAP Int64Regs:$s,
3564            Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3565            Int32Regs:$r, Int32Regs:$g)>;
3566
3567 def : Pat<(int_nvvm_sust_p_3d_v4i8_trap
3568            Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3569            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3570           (SUST_P_3D_V4B8_TRAP Int64Regs:$s,
3571            Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3572            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3573
3574 def : Pat<(int_nvvm_sust_p_3d_v4i16_trap
3575            Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3576            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3577           (SUST_P_3D_V4B16_TRAP Int64Regs:$s,
3578            Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3579            Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3580
3581 def : Pat<(int_nvvm_sust_p_3d_v4i32_trap
3582            Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3583            Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
3584           (SUST_P_3D_V4B32_TRAP Int64Regs:$s,
3585            Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3586            Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a)>;
3587
3588
3589
3590 //===-- Old PTX Back-end Intrinsics ---------------------------------------===//
3591
3592 // These intrinsics are handled to retain compatibility with the old backend.
3593
3594 // PTX Special Purpose Register Accessor Intrinsics
3595
3596 class PTX_READ_SPECIAL_REGISTER_R64<string regname, Intrinsic intop>
3597   : NVPTXInst<(outs Int64Regs:$d), (ins),
3598               !strconcat(!strconcat("mov.u64\t$d, %", regname), ";"),
3599               [(set Int64Regs:$d, (intop))]>;
3600
3601 class PTX_READ_SPECIAL_REGISTER_R32<string regname, Intrinsic intop>
3602   : NVPTXInst<(outs Int32Regs:$d), (ins),
3603               !strconcat(!strconcat("mov.u32\t$d, %", regname), ";"),
3604               [(set Int32Regs:$d, (intop))]>;
3605
3606 // TODO Add read vector-version of special registers
3607
3608 def PTX_READ_TID_X   : PTX_READ_SPECIAL_REGISTER_R32<"tid.x",
3609                                                      int_ptx_read_tid_x>;
3610 def PTX_READ_TID_Y   : PTX_READ_SPECIAL_REGISTER_R32<"tid.y",
3611                                                      int_ptx_read_tid_y>;
3612 def PTX_READ_TID_Z   : PTX_READ_SPECIAL_REGISTER_R32<"tid.z",
3613                                                      int_ptx_read_tid_z>;
3614 def PTX_READ_TID_W   : PTX_READ_SPECIAL_REGISTER_R32<"tid.w",
3615                                                      int_ptx_read_tid_w>;
3616
3617 def PTX_READ_NTID_X   : PTX_READ_SPECIAL_REGISTER_R32<"ntid.x",
3618                                                       int_ptx_read_ntid_x>;
3619 def PTX_READ_NTID_Y   : PTX_READ_SPECIAL_REGISTER_R32<"ntid.y",
3620                                                       int_ptx_read_ntid_y>;
3621 def PTX_READ_NTID_Z   : PTX_READ_SPECIAL_REGISTER_R32<"ntid.z",
3622                                                       int_ptx_read_ntid_z>;
3623 def PTX_READ_NTID_W   : PTX_READ_SPECIAL_REGISTER_R32<"ntid.w",
3624                                                       int_ptx_read_ntid_w>;
3625
3626 def PTX_READ_LANEID  : PTX_READ_SPECIAL_REGISTER_R32<"laneid",
3627                                                      int_ptx_read_laneid>;
3628 def PTX_READ_WARPID  : PTX_READ_SPECIAL_REGISTER_R32<"warpid",
3629                                                      int_ptx_read_warpid>;
3630 def PTX_READ_NWARPID : PTX_READ_SPECIAL_REGISTER_R32<"nwarpid",
3631                                                      int_ptx_read_nwarpid>;
3632
3633 def PTX_READ_CTAID_X   : PTX_READ_SPECIAL_REGISTER_R32<"ctaid.x",
3634                                                        int_ptx_read_ctaid_x>;
3635 def PTX_READ_CTAID_Y   : PTX_READ_SPECIAL_REGISTER_R32<"ctaid.y",
3636                                                        int_ptx_read_ctaid_y>;
3637 def PTX_READ_CTAID_Z   : PTX_READ_SPECIAL_REGISTER_R32<"ctaid.z",
3638                                                        int_ptx_read_ctaid_z>;
3639 def PTX_READ_CTAID_W   : PTX_READ_SPECIAL_REGISTER_R32<"ctaid.w",
3640                                                        int_ptx_read_ctaid_w>;
3641
3642 def PTX_READ_NCTAID_X   : PTX_READ_SPECIAL_REGISTER_R32<"nctaid.x",
3643                                                         int_ptx_read_nctaid_x>;
3644 def PTX_READ_NCTAID_Y   : PTX_READ_SPECIAL_REGISTER_R32<"nctaid.y",
3645                                                         int_ptx_read_nctaid_y>;
3646 def PTX_READ_NCTAID_Z   : PTX_READ_SPECIAL_REGISTER_R32<"nctaid.z",
3647                                                         int_ptx_read_nctaid_z>;
3648 def PTX_READ_NCTAID_W   : PTX_READ_SPECIAL_REGISTER_R32<"nctaid.w",
3649                                                         int_ptx_read_nctaid_w>;
3650
3651 def PTX_READ_SMID  : PTX_READ_SPECIAL_REGISTER_R32<"smid",
3652                                                    int_ptx_read_smid>;
3653 def PTX_READ_NSMID  : PTX_READ_SPECIAL_REGISTER_R32<"nsmid",
3654                                                     int_ptx_read_nsmid>;
3655 def PTX_READ_GRIDID  : PTX_READ_SPECIAL_REGISTER_R32<"gridid",
3656                                                      int_ptx_read_gridid>;
3657
3658 def PTX_READ_LANEMASK_EQ
3659   : PTX_READ_SPECIAL_REGISTER_R32<"lanemask_eq", int_ptx_read_lanemask_eq>;
3660 def PTX_READ_LANEMASK_LE
3661   : PTX_READ_SPECIAL_REGISTER_R32<"lanemask_le", int_ptx_read_lanemask_le>;
3662 def PTX_READ_LANEMASK_LT
3663   : PTX_READ_SPECIAL_REGISTER_R32<"lanemask_lt", int_ptx_read_lanemask_lt>;
3664 def PTX_READ_LANEMASK_GE
3665   : PTX_READ_SPECIAL_REGISTER_R32<"lanemask_ge", int_ptx_read_lanemask_ge>;
3666 def PTX_READ_LANEMASK_GT
3667   : PTX_READ_SPECIAL_REGISTER_R32<"lanemask_gt", int_ptx_read_lanemask_gt>;
3668
3669 def PTX_READ_CLOCK
3670   : PTX_READ_SPECIAL_REGISTER_R32<"clock", int_ptx_read_clock>;
3671 def PTX_READ_CLOCK64
3672   : PTX_READ_SPECIAL_REGISTER_R64<"clock64", int_ptx_read_clock64>;
3673
3674 def PTX_READ_PM0 : PTX_READ_SPECIAL_REGISTER_R32<"pm0", int_ptx_read_pm0>;
3675 def PTX_READ_PM1 : PTX_READ_SPECIAL_REGISTER_R32<"pm1", int_ptx_read_pm1>;
3676 def PTX_READ_PM2 : PTX_READ_SPECIAL_REGISTER_R32<"pm2", int_ptx_read_pm2>;
3677 def PTX_READ_PM3 : PTX_READ_SPECIAL_REGISTER_R32<"pm3", int_ptx_read_pm3>;
3678
3679 // PTX Parallel Synchronization and Communication Intrinsics
3680
3681 def PTX_BAR_SYNC : NVPTXInst<(outs), (ins i32imm:$i), "bar.sync\t$i;",
3682                              [(int_ptx_bar_sync imm:$i)]>;