//===- NVPTXIntrinsics.td - PTX Intrinsics Instructions -------*- tblgen -*-==//
//
//                     The LLVM Compiler Infrastructure
//
// This file is distributed under the University of Illinois Open Source
// License. See LICENSE.TXT for details.
//
//===----------------------------------------------------------------------===//

def immFloat0 : PatLeaf<(fpimm), [{
    float f = (float)N->getValueAPF().convertToFloat();
    return (f==0.0f);
}]>;

def immFloat1 : PatLeaf<(fpimm), [{
    float f = (float)N->getValueAPF().convertToFloat();
    return (f==1.0f);
}]>;

def immDouble0 : PatLeaf<(fpimm), [{
    double d = (double)N->getValueAPF().convertToDouble();
    return (d==0.0);
}]>;

def immDouble1 : PatLeaf<(fpimm), [{
    double d = (double)N->getValueAPF().convertToDouble();
    return (d==1.0);
}]>;



//-----------------------------------
// Synchronization Functions
//-----------------------------------
def INT_CUDA_SYNCTHREADS : NVPTXInst<(outs), (ins),
                  "bar.sync \t0;",
      [(int_cuda_syncthreads)]>;
def INT_BARRIER0 : NVPTXInst<(outs), (ins),
                  "bar.sync \t0;",
      [(int_nvvm_barrier0)]>;
def INT_BARRIER0_POPC : NVPTXInst<(outs Int32Regs:$dst), (ins Int32Regs:$pred),
  !strconcat("{{ \n\t",
      !strconcat(".reg .pred \t%p1; \n\t",
      !strconcat("setp.ne.u32 \t%p1, $pred, 0; \n\t",
      !strconcat("bar.red.popc.u32 \t$dst, 0, %p1; \n\t",
        !strconcat("}}", ""))))),
      [(set Int32Regs:$dst, (int_nvvm_barrier0_popc Int32Regs:$pred))]>;
def INT_BARRIER0_AND : NVPTXInst<(outs Int32Regs:$dst), (ins Int32Regs:$pred),
  !strconcat("{{ \n\t",
      !strconcat(".reg .pred \t%p1; \n\t",
      !strconcat(".reg .pred \t%p2; \n\t",
      !strconcat("setp.ne.u32 \t%p1, $pred, 0; \n\t",
      !strconcat("bar.red.and.pred \t%p2, 0, %p1; \n\t",
      !strconcat("selp.u32 \t$dst, 1, 0, %p2; \n\t",
        !strconcat("}}", ""))))))),
      [(set Int32Regs:$dst, (int_nvvm_barrier0_and Int32Regs:$pred))]>;
def INT_BARRIER0_OR : NVPTXInst<(outs Int32Regs:$dst), (ins Int32Regs:$pred),
  !strconcat("{{ \n\t",
      !strconcat(".reg .pred \t%p1; \n\t",
      !strconcat(".reg .pred \t%p2; \n\t",
      !strconcat("setp.ne.u32 \t%p1, $pred, 0; \n\t",
      !strconcat("bar.red.or.pred \t%p2, 0, %p1; \n\t",
      !strconcat("selp.u32 \t$dst, 1, 0, %p2; \n\t",
        !strconcat("}}", ""))))))),
      [(set Int32Regs:$dst, (int_nvvm_barrier0_or Int32Regs:$pred))]>;


//-----------------------------------
// Explicit Memory Fence Functions
//-----------------------------------
class MEMBAR<string StrOp, Intrinsic IntOP> :
              NVPTXInst<(outs), (ins),
            StrOp, [(IntOP)]>;

def INT_MEMBAR_CTA : MEMBAR<"membar.cta;", int_nvvm_membar_cta>;
def INT_MEMBAR_GL  : MEMBAR<"membar.gl;",  int_nvvm_membar_gl>;
def INT_MEMBAR_SYS : MEMBAR<"membar.sys;", int_nvvm_membar_sys>;


//-----------------------------------
// Math Functions
//-----------------------------------

// Map min(1.0, max(0.0, x)) to sat(x)
// Note that max(0.0, min(x, 1.0)) cannot be mapped to sat(x) because when x is
// NaN
// max(0.0, min(x, 1.0)) is 1.0 while sat(x) is 0.
// Same story for fmax, fmin.

def : Pat<(int_nvvm_fmin_f immFloat1,
            (int_nvvm_fmax_f immFloat0, Float32Regs:$a)),
          (CVT_f32_f32 Float32Regs:$a, CvtSAT)>;
def : Pat<(int_nvvm_fmin_f immFloat1,
            (int_nvvm_fmax_f Float32Regs:$a, immFloat0)),
          (CVT_f32_f32 Float32Regs:$a, CvtSAT)>;
def : Pat<(int_nvvm_fmin_f
            (int_nvvm_fmax_f immFloat0, Float32Regs:$a), immFloat1),
          (CVT_f32_f32 Float32Regs:$a, CvtSAT)>;
def : Pat<(int_nvvm_fmin_f
            (int_nvvm_fmax_f Float32Regs:$a, immFloat0), immFloat1),
          (CVT_f32_f32 Float32Regs:$a, CvtSAT)>;

def : Pat<(int_nvvm_fmin_d immDouble1,
            (int_nvvm_fmax_d immDouble0, Float64Regs:$a)),
          (CVT_f64_f64 Float64Regs:$a, CvtSAT)>;
def : Pat<(int_nvvm_fmin_d immDouble1,
            (int_nvvm_fmax_d Float64Regs:$a, immDouble0)),
          (CVT_f64_f64 Float64Regs:$a, CvtSAT)>;
def : Pat<(int_nvvm_fmin_d
            (int_nvvm_fmax_d immDouble0, Float64Regs:$a), immDouble1),
          (CVT_f64_f64 Float64Regs:$a, CvtSAT)>;
def : Pat<(int_nvvm_fmin_d
            (int_nvvm_fmax_d Float64Regs:$a, immDouble0), immDouble1),
          (CVT_f64_f64 Float64Regs:$a, CvtSAT)>;


// We need a full string for OpcStr here because we need to deal with case like
// INT_PTX_RECIP.
class F_MATH_1<string OpcStr, NVPTXRegClass target_regclass,
  NVPTXRegClass src_regclass, Intrinsic IntOP>
            : NVPTXInst<(outs target_regclass:$dst), (ins src_regclass:$src0),
            OpcStr,
        [(set target_regclass:$dst, (IntOP src_regclass:$src0))]>;

// We need a full string for OpcStr here because we need to deal with the case
// like INT_PTX_NATIVE_POWR_F.
class F_MATH_2<string OpcStr, NVPTXRegClass t_regclass,
  NVPTXRegClass s0_regclass, NVPTXRegClass s1_regclass, Intrinsic IntOP>
            : NVPTXInst<(outs t_regclass:$dst),
              (ins s0_regclass:$src0, s1_regclass:$src1),
            OpcStr,
        [(set t_regclass:$dst, (IntOP s0_regclass:$src0, s1_regclass:$src1))]>;

class F_MATH_3<string OpcStr, NVPTXRegClass t_regclass,
  NVPTXRegClass s0_regclass, NVPTXRegClass s1_regclass,
  NVPTXRegClass s2_regclass, Intrinsic IntOP>
            : NVPTXInst<(outs t_regclass:$dst),
              (ins s0_regclass:$src0, s1_regclass:$src1, s2_regclass:$src2),
            OpcStr,
        [(set t_regclass:$dst,
          (IntOP s0_regclass:$src0, s1_regclass:$src1, s2_regclass:$src2))]>;

//
// MISC
//

def INT_NVVM_CLZ_I : F_MATH_1<"clz.b32 \t$dst, $src0;", Int32Regs, Int32Regs,
  int_nvvm_clz_i>;
def INT_NVVM_CLZ_LL : F_MATH_1<"clz.b64 \t$dst, $src0;", Int32Regs, Int64Regs,
  int_nvvm_clz_ll>;

def INT_NVVM_POPC_I : F_MATH_1<"popc.b32 \t$dst, $src0;", Int32Regs, Int32Regs,
  int_nvvm_popc_i>;
def INT_NVVM_POPC_LL : F_MATH_1<"popc.b64 \t$dst, $src0;", Int32Regs, Int64Regs,
  int_nvvm_popc_ll>;

def INT_NVVM_PRMT : F_MATH_3<"prmt.b32 \t$dst, $src0, $src1, $src2;", Int32Regs,
  Int32Regs, Int32Regs, Int32Regs, int_nvvm_prmt>;

//
// Min Max
//

def INT_NVVM_MIN_I : F_MATH_2<"min.s32 \t$dst, $src0, $src1;", Int32Regs,
  Int32Regs, Int32Regs, int_nvvm_min_i>;
def INT_NVVM_MIN_UI : F_MATH_2<"min.u32 \t$dst, $src0, $src1;", Int32Regs,
  Int32Regs, Int32Regs, int_nvvm_min_ui>;

def INT_NVVM_MIN_LL : F_MATH_2<"min.s64 \t$dst, $src0, $src1;", Int64Regs,
  Int64Regs, Int64Regs, int_nvvm_min_ll>;
def INT_NVVM_MIN_ULL : F_MATH_2<"min.u64 \t$dst, $src0, $src1;", Int64Regs,
  Int64Regs, Int64Regs, int_nvvm_min_ull>;

def INT_NVVM_MAX_I : F_MATH_2<"max.s32 \t$dst, $src0, $src1;", Int32Regs,
  Int32Regs, Int32Regs, int_nvvm_max_i>;
def INT_NVVM_MAX_UI : F_MATH_2<"max.u32 \t$dst, $src0, $src1;", Int32Regs,
  Int32Regs, Int32Regs, int_nvvm_max_ui>;

def INT_NVVM_MAX_LL : F_MATH_2<"max.s64 \t$dst, $src0, $src1;", Int64Regs,
  Int64Regs, Int64Regs, int_nvvm_max_ll>;
def INT_NVVM_MAX_ULL : F_MATH_2<"max.u64 \t$dst, $src0, $src1;", Int64Regs,
  Int64Regs, Int64Regs, int_nvvm_max_ull>;

def INT_NVVM_FMIN_F : F_MATH_2<"min.f32 \t$dst, $src0, $src1;", Float32Regs,
  Float32Regs, Float32Regs, int_nvvm_fmin_f>;
def INT_NVVM_FMIN_FTZ_F : F_MATH_2<"min.ftz.f32 \t$dst, $src0, $src1;",
  Float32Regs, Float32Regs, Float32Regs, int_nvvm_fmin_ftz_f>;

def INT_NVVM_FMAX_F : F_MATH_2<"max.f32 \t$dst, $src0, $src1;", Float32Regs,
  Float32Regs, Float32Regs, int_nvvm_fmax_f>;
def INT_NVVM_FMAX_FTZ_F : F_MATH_2<"max.ftz.f32 \t$dst, $src0, $src1;",
  Float32Regs, Float32Regs, Float32Regs, int_nvvm_fmax_ftz_f>;

def INT_NVVM_FMIN_D : F_MATH_2<"min.f64 \t$dst, $src0, $src1;", Float64Regs,
  Float64Regs, Float64Regs, int_nvvm_fmin_d>;
def INT_NVVM_FMAX_D : F_MATH_2<"max.f64 \t$dst, $src0, $src1;", Float64Regs,
  Float64Regs, Float64Regs, int_nvvm_fmax_d>;

//
// Multiplication
//

def INT_NVVM_MULHI_I : F_MATH_2<"mul.hi.s32 \t$dst, $src0, $src1;", Int32Regs,
  Int32Regs, Int32Regs, int_nvvm_mulhi_i>;
def INT_NVVM_MULHI_UI : F_MATH_2<"mul.hi.u32 \t$dst, $src0, $src1;", Int32Regs,
  Int32Regs, Int32Regs, int_nvvm_mulhi_ui>;

def INT_NVVM_MULHI_LL : F_MATH_2<"mul.hi.s64 \t$dst, $src0, $src1;", Int64Regs,
  Int64Regs, Int64Regs, int_nvvm_mulhi_ll>;
def INT_NVVM_MULHI_ULL : F_MATH_2<"mul.hi.u64 \t$dst, $src0, $src1;", Int64Regs,
  Int64Regs, Int64Regs, int_nvvm_mulhi_ull>;

def INT_NVVM_MUL_RN_FTZ_F : F_MATH_2<"mul.rn.ftz.f32 \t$dst, $src0, $src1;",
  Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rn_ftz_f>;
def INT_NVVM_MUL_RN_F : F_MATH_2<"mul.rn.f32 \t$dst, $src0, $src1;",
  Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rn_f>;
def INT_NVVM_MUL_RZ_FTZ_F : F_MATH_2<"mul.rz.ftz.f32 \t$dst, $src0, $src1;",
  Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rz_ftz_f>;
def INT_NVVM_MUL_RZ_F : F_MATH_2<"mul.rz.f32 \t$dst, $src0, $src1;",
  Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rz_f>;
def INT_NVVM_MUL_RM_FTZ_F : F_MATH_2<"mul.rm.ftz.f32 \t$dst, $src0, $src1;",
  Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rm_ftz_f>;
def INT_NVVM_MUL_RM_F : F_MATH_2<"mul.rm.f32 \t$dst, $src0, $src1;",
  Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rm_f>;
def INT_NVVM_MUL_RP_FTZ_F : F_MATH_2<"mul.rp.ftz.f32 \t$dst, $src0, $src1;",
  Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rp_ftz_f>;
def INT_NVVM_MUL_RP_F : F_MATH_2<"mul.rp.f32 \t$dst, $src0, $src1;",
  Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rp_f>;

def INT_NVVM_MUL_RN_D : F_MATH_2<"mul.rn.f64 \t$dst, $src0, $src1;",
  Float64Regs, Float64Regs, Float64Regs, int_nvvm_mul_rn_d>;
def INT_NVVM_MUL_RZ_D : F_MATH_2<"mul.rz.f64 \t$dst, $src0, $src1;",
  Float64Regs, Float64Regs, Float64Regs, int_nvvm_mul_rz_d>;
def INT_NVVM_MUL_RM_D : F_MATH_2<"mul.rm.f64 \t$dst, $src0, $src1;",
  Float64Regs, Float64Regs, Float64Regs, int_nvvm_mul_rm_d>;
def INT_NVVM_MUL_RP_D : F_MATH_2<"mul.rp.f64 \t$dst, $src0, $src1;",
  Float64Regs, Float64Regs, Float64Regs, int_nvvm_mul_rp_d>;

def INT_NVVM_MUL24_I : F_MATH_2<"mul24.lo.s32 \t$dst, $src0, $src1;",
  Int32Regs, Int32Regs, Int32Regs, int_nvvm_mul24_i>;
def INT_NVVM_MUL24_UI : F_MATH_2<"mul24.lo.u32 \t$dst, $src0, $src1;",
  Int32Regs, Int32Regs, Int32Regs, int_nvvm_mul24_ui>;

//
// Div
//

def INT_NVVM_DIV_APPROX_FTZ_F
  : F_MATH_2<"div.approx.ftz.f32 \t$dst, $src0, $src1;", Float32Regs,
    Float32Regs, Float32Regs, int_nvvm_div_approx_ftz_f>;
def INT_NVVM_DIV_APPROX_F : F_MATH_2<"div.approx.f32 \t$dst, $src0, $src1;",
  Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_approx_f>;

def INT_NVVM_DIV_RN_FTZ_F : F_MATH_2<"div.rn.ftz.f32 \t$dst, $src0, $src1;",
  Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rn_ftz_f>;
def INT_NVVM_DIV_RN_F     : F_MATH_2<"div.rn.f32 \t$dst, $src0, $src1;",
  Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rn_f>;
def INT_NVVM_DIV_RZ_FTZ_F : F_MATH_2<"div.rz.ftz.f32 \t$dst, $src0, $src1;",
  Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rz_ftz_f>;
def INT_NVVM_DIV_RZ_F     : F_MATH_2<"div.rz.f32 \t$dst, $src0, $src1;",
  Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rz_f>;
def INT_NVVM_DIV_RM_FTZ_F : F_MATH_2<"div.rm.ftz.f32 \t$dst, $src0, $src1;",
  Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rm_ftz_f>;
def INT_NVVM_DIV_RM_F     : F_MATH_2<"div.rm.f32 \t$dst, $src0, $src1;",
  Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rm_f>;
def INT_NVVM_DIV_RP_FTZ_F : F_MATH_2<"div.rp.ftz.f32 \t$dst, $src0, $src1;",
  Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rp_ftz_f>;
def INT_NVVM_DIV_RP_F     : F_MATH_2<"div.rp.f32 \t$dst, $src0, $src1;",
  Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rp_f>;

def INT_NVVM_DIV_RN_D : F_MATH_2<"div.rn.f64 \t$dst, $src0, $src1;",
  Float64Regs, Float64Regs, Float64Regs, int_nvvm_div_rn_d>;
def INT_NVVM_DIV_RZ_D : F_MATH_2<"div.rz.f64 \t$dst, $src0, $src1;",
  Float64Regs, Float64Regs, Float64Regs, int_nvvm_div_rz_d>;
def INT_NVVM_DIV_RM_D : F_MATH_2<"div.rm.f64 \t$dst, $src0, $src1;",
  Float64Regs, Float64Regs, Float64Regs, int_nvvm_div_rm_d>;
def INT_NVVM_DIV_RP_D : F_MATH_2<"div.rp.f64 \t$dst, $src0, $src1;",
  Float64Regs, Float64Regs, Float64Regs, int_nvvm_div_rp_d>;

//
// Brev
//

def INT_NVVM_BREV32 : F_MATH_1<"brev.b32 \t$dst, $src0;", Int32Regs, Int32Regs,
  int_nvvm_brev32>;
def INT_NVVM_BREV64 : F_MATH_1<"brev.b64 \t$dst, $src0;", Int64Regs, Int64Regs,
  int_nvvm_brev64>;

//
// Sad
//

def INT_NVVM_SAD_I : F_MATH_3<"sad.s32 \t$dst, $src0, $src1, $src2;",
  Int32Regs, Int32Regs, Int32Regs, Int32Regs, int_nvvm_sad_i>;
def INT_NVVM_SAD_UI : F_MATH_3<"sad.u32 \t$dst, $src0, $src1, $src2;",
  Int32Regs, Int32Regs, Int32Regs, Int32Regs, int_nvvm_sad_ui>;

//
// Floor  Ceil
//

def : Pat<(int_nvvm_floor_ftz_f Float32Regs:$a),
          (CVT_f32_f32 Float32Regs:$a, CvtRMI_FTZ)>;
def : Pat<(int_nvvm_floor_f Float32Regs:$a),
          (CVT_f32_f32 Float32Regs:$a, CvtRMI)>;
def : Pat<(int_nvvm_floor_d Float64Regs:$a),
          (CVT_f64_f64 Float64Regs:$a, CvtRMI)>;

def : Pat<(int_nvvm_ceil_ftz_f Float32Regs:$a),
          (CVT_f32_f32 Float32Regs:$a, CvtRPI_FTZ)>;
def : Pat<(int_nvvm_ceil_f Float32Regs:$a),
          (CVT_f32_f32 Float32Regs:$a, CvtRPI)>;
def : Pat<(int_nvvm_ceil_d Float64Regs:$a),
          (CVT_f64_f64 Float64Regs:$a, CvtRPI)>;

//
// Abs
//

def INT_NVVM_ABS_I : F_MATH_1<"abs.s32 \t$dst, $src0;", Int32Regs, Int32Regs,
  int_nvvm_abs_i>;
def INT_NVVM_ABS_LL : F_MATH_1<"abs.s64 \t$dst, $src0;", Int64Regs, Int64Regs,
  int_nvvm_abs_ll>;

def INT_NVVM_FABS_FTZ_F : F_MATH_1<"abs.ftz.f32 \t$dst, $src0;", Float32Regs,
  Float32Regs, int_nvvm_fabs_ftz_f>;
def INT_NVVM_FABS_F : F_MATH_1<"abs.f32 \t$dst, $src0;", Float32Regs,
  Float32Regs, int_nvvm_fabs_f>;

def INT_NVVM_FABS_D : F_MATH_1<"abs.f64 \t$dst, $src0;", Float64Regs,
  Float64Regs, int_nvvm_fabs_d>;

//
// Round
//

def : Pat<(int_nvvm_round_ftz_f Float32Regs:$a),
          (CVT_f32_f32 Float32Regs:$a, CvtRNI_FTZ)>;
def : Pat<(int_nvvm_round_f Float32Regs:$a),
          (CVT_f32_f32 Float32Regs:$a, CvtRNI)>;
def : Pat<(int_nvvm_round_d Float64Regs:$a),
          (CVT_f64_f64 Float64Regs:$a, CvtRNI)>;

//
// Trunc
//

def : Pat<(int_nvvm_trunc_ftz_f Float32Regs:$a),
          (CVT_f32_f32 Float32Regs:$a, CvtRZI_FTZ)>;
def : Pat<(int_nvvm_trunc_f Float32Regs:$a),
          (CVT_f32_f32 Float32Regs:$a, CvtRZI)>;
def : Pat<(int_nvvm_trunc_d Float64Regs:$a),
          (CVT_f64_f64 Float64Regs:$a, CvtRZI)>;

//
// Saturate
//

def : Pat<(int_nvvm_saturate_ftz_f Float32Regs:$a),
          (CVT_f32_f32 Float32Regs:$a, CvtSAT_FTZ)>;
def : Pat<(int_nvvm_saturate_f Float32Regs:$a),
          (CVT_f32_f32 Float32Regs:$a, CvtSAT)>;
def : Pat<(int_nvvm_saturate_d Float64Regs:$a),
          (CVT_f64_f64 Float64Regs:$a, CvtSAT)>;

//
// Exp2  Log2
//

def INT_NVVM_EX2_APPROX_FTZ_F : F_MATH_1<"ex2.approx.ftz.f32 \t$dst, $src0;",
  Float32Regs, Float32Regs, int_nvvm_ex2_approx_ftz_f>;
def INT_NVVM_EX2_APPROX_F : F_MATH_1<"ex2.approx.f32 \t$dst, $src0;",
  Float32Regs, Float32Regs, int_nvvm_ex2_approx_f>;
def INT_NVVM_EX2_APPROX_D : F_MATH_1<"ex2.approx.f64 \t$dst, $src0;",
  Float64Regs, Float64Regs, int_nvvm_ex2_approx_d>;

def INT_NVVM_LG2_APPROX_FTZ_F : F_MATH_1<"lg2.approx.ftz.f32 \t$dst, $src0;",
  Float32Regs, Float32Regs, int_nvvm_lg2_approx_ftz_f>;
def INT_NVVM_LG2_APPROX_F : F_MATH_1<"lg2.approx.f32 \t$dst, $src0;",
  Float32Regs, Float32Regs, int_nvvm_lg2_approx_f>;
def INT_NVVM_LG2_APPROX_D : F_MATH_1<"lg2.approx.f64 \t$dst, $src0;",
  Float64Regs, Float64Regs, int_nvvm_lg2_approx_d>;

//
// Sin  Cos
//

def INT_NVVM_SIN_APPROX_FTZ_F : F_MATH_1<"sin.approx.ftz.f32 \t$dst, $src0;",
  Float32Regs, Float32Regs, int_nvvm_sin_approx_ftz_f>;
def INT_NVVM_SIN_APPROX_F : F_MATH_1<"sin.approx.f32 \t$dst, $src0;",
  Float32Regs, Float32Regs, int_nvvm_sin_approx_f>;

def INT_NVVM_COS_APPROX_FTZ_F : F_MATH_1<"cos.approx.ftz.f32 \t$dst, $src0;",
  Float32Regs, Float32Regs, int_nvvm_cos_approx_ftz_f>;
def INT_NVVM_COS_APPROX_F : F_MATH_1<"cos.approx.f32 \t$dst, $src0;",
  Float32Regs, Float32Regs, int_nvvm_cos_approx_f>;

//
// Fma
//

def INT_NVVM_FMA_RN_FTZ_F
  : F_MATH_3<"fma.rn.ftz.f32 \t$dst, $src0, $src1, $src2;", Float32Regs,
    Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rn_ftz_f>;
def INT_NVVM_FMA_RN_F : F_MATH_3<"fma.rn.f32 \t$dst, $src0, $src1, $src2;",
  Float32Regs, Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rn_f>;
def INT_NVVM_FMA_RZ_FTZ_F
  : F_MATH_3<"fma.rz.ftz.f32 \t$dst, $src0, $src1, $src2;", Float32Regs,
    Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rz_ftz_f>;
def INT_NVVM_FMA_RZ_F : F_MATH_3<"fma.rz.f32 \t$dst, $src0, $src1, $src2;",
  Float32Regs, Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rz_f>;
def INT_NVVM_FMA_RM_FTZ_F
  : F_MATH_3<"fma.rm.ftz.f32 \t$dst, $src0, $src1, $src2;", Float32Regs,
    Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rm_ftz_f>;
def INT_NVVM_FMA_RM_F : F_MATH_3<"fma.rm.f32 \t$dst, $src0, $src1, $src2;",
  Float32Regs, Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rm_f>;
def INT_NVVM_FMA_RP_FTZ_F
  : F_MATH_3<"fma.rp.ftz.f32 \t$dst, $src0, $src1, $src2;", Float32Regs,
    Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rp_ftz_f>;
def INT_NVVM_FMA_RP_F : F_MATH_3<"fma.rp.f32 \t$dst, $src0, $src1, $src2;",
  Float32Regs, Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rp_f>;

def INT_NVVM_FMA_RN_D : F_MATH_3<"fma.rn.f64 \t$dst, $src0, $src1, $src2;",
  Float64Regs, Float64Regs, Float64Regs, Float64Regs, int_nvvm_fma_rn_d>;
def INT_NVVM_FMA_RZ_D : F_MATH_3<"fma.rz.f64 \t$dst, $src0, $src1, $src2;",
  Float64Regs, Float64Regs, Float64Regs, Float64Regs, int_nvvm_fma_rz_d>;
def INT_NVVM_FMA_RM_D : F_MATH_3<"fma.rm.f64 \t$dst, $src0, $src1, $src2;",
  Float64Regs, Float64Regs, Float64Regs, Float64Regs, int_nvvm_fma_rm_d>;
def INT_NVVM_FMA_RP_D : F_MATH_3<"fma.rp.f64 \t$dst, $src0, $src1, $src2;",
  Float64Regs, Float64Regs, Float64Regs, Float64Regs, int_nvvm_fma_rp_d>;

//
// Rcp
//

def INT_NVVM_RCP_RN_FTZ_F : F_MATH_1<"rcp.rn.ftz.f32 \t$dst, $src0;",
  Float32Regs, Float32Regs, int_nvvm_rcp_rn_ftz_f>;
def INT_NVVM_RCP_RN_F : F_MATH_1<"rcp.rn.f32 \t$dst, $src0;",
  Float32Regs, Float32Regs, int_nvvm_rcp_rn_f>;
def INT_NVVM_RCP_RZ_FTZ_F : F_MATH_1<"rcp.rz.ftz.f32 \t$dst, $src0;",
  Float32Regs, Float32Regs, int_nvvm_rcp_rz_ftz_f>;
def INT_NVVM_RCP_RZ_F : F_MATH_1<"rcp.rz.f32 \t$dst, $src0;",
  Float32Regs, Float32Regs, int_nvvm_rcp_rz_f>;
def INT_NVVM_RCP_RM_FTZ_F : F_MATH_1<"rcp.rm.ftz.f32 \t$dst, $src0;",
  Float32Regs, Float32Regs, int_nvvm_rcp_rm_ftz_f>;
def INT_NVVM_RCP_RM_F : F_MATH_1<"rcp.rm.f32 \t$dst, $src0;",
  Float32Regs, Float32Regs, int_nvvm_rcp_rm_f>;
def INT_NVVM_RCP_RP_FTZ_F : F_MATH_1<"rcp.rp.ftz.f32 \t$dst, $src0;",
  Float32Regs, Float32Regs, int_nvvm_rcp_rp_ftz_f>;
def INT_NVVM_RCP_RP_F : F_MATH_1<"rcp.rp.f32 \t$dst, $src0;",
  Float32Regs, Float32Regs, int_nvvm_rcp_rp_f>;

def INT_NVVM_RCP_RN_D : F_MATH_1<"rcp.rn.f64 \t$dst, $src0;", Float64Regs,
  Float64Regs, int_nvvm_rcp_rn_d>;
def INT_NVVM_RCP_RZ_D : F_MATH_1<"rcp.rz.f64 \t$dst, $src0;", Float64Regs,
  Float64Regs, int_nvvm_rcp_rz_d>;
def INT_NVVM_RCP_RM_D : F_MATH_1<"rcp.rm.f64 \t$dst, $src0;", Float64Regs,
  Float64Regs, int_nvvm_rcp_rm_d>;
def INT_NVVM_RCP_RP_D : F_MATH_1<"rcp.rp.f64 \t$dst, $src0;", Float64Regs,
  Float64Regs, int_nvvm_rcp_rp_d>;

def INT_NVVM_RCP_APPROX_FTZ_D : F_MATH_1<"rcp.approx.ftz.f64 \t$dst, $src0;",
  Float64Regs, Float64Regs, int_nvvm_rcp_approx_ftz_d>;

//
// Sqrt
//

def INT_NVVM_SQRT_RN_FTZ_F : F_MATH_1<"sqrt.rn.ftz.f32 \t$dst, $src0;",
  Float32Regs, Float32Regs, int_nvvm_sqrt_rn_ftz_f>;
def INT_NVVM_SQRT_RN_F : F_MATH_1<"sqrt.rn.f32 \t$dst, $src0;", Float32Regs,
  Float32Regs, int_nvvm_sqrt_rn_f>;
def INT_NVVM_SQRT_RZ_FTZ_F : F_MATH_1<"sqrt.rz.ftz.f32 \t$dst, $src0;",
  Float32Regs, Float32Regs, int_nvvm_sqrt_rz_ftz_f>;
def INT_NVVM_SQRT_RZ_F : F_MATH_1<"sqrt.rz.f32 \t$dst, $src0;", Float32Regs,
  Float32Regs, int_nvvm_sqrt_rz_f>;
def INT_NVVM_SQRT_RM_FTZ_F : F_MATH_1<"sqrt.rm.ftz.f32 \t$dst, $src0;",
  Float32Regs, Float32Regs, int_nvvm_sqrt_rm_ftz_f>;
def INT_NVVM_SQRT_RM_F : F_MATH_1<"sqrt.rm.f32 \t$dst, $src0;", Float32Regs,
  Float32Regs, int_nvvm_sqrt_rm_f>;
def INT_NVVM_SQRT_RP_FTZ_F : F_MATH_1<"sqrt.rp.ftz.f32 \t$dst, $src0;",
  Float32Regs, Float32Regs, int_nvvm_sqrt_rp_ftz_f>;
def INT_NVVM_SQRT_RP_F : F_MATH_1<"sqrt.rp.f32 \t$dst, $src0;", Float32Regs,
  Float32Regs, int_nvvm_sqrt_rp_f>;
def INT_NVVM_SQRT_APPROX_FTZ_F : F_MATH_1<"sqrt.approx.ftz.f32 \t$dst, $src0;",
  Float32Regs, Float32Regs, int_nvvm_sqrt_approx_ftz_f>;
def INT_NVVM_SQRT_APPROX_F : F_MATH_1<"sqrt.approx.f32 \t$dst, $src0;",
  Float32Regs, Float32Regs, int_nvvm_sqrt_approx_f>;

def INT_NVVM_SQRT_RN_D : F_MATH_1<"sqrt.rn.f64 \t$dst, $src0;", Float64Regs,
  Float64Regs, int_nvvm_sqrt_rn_d>;
def INT_NVVM_SQRT_RZ_D : F_MATH_1<"sqrt.rz.f64 \t$dst, $src0;", Float64Regs,
  Float64Regs, int_nvvm_sqrt_rz_d>;
def INT_NVVM_SQRT_RM_D : F_MATH_1<"sqrt.rm.f64 \t$dst, $src0;", Float64Regs,
  Float64Regs, int_nvvm_sqrt_rm_d>;
def INT_NVVM_SQRT_RP_D : F_MATH_1<"sqrt.rp.f64 \t$dst, $src0;", Float64Regs,
  Float64Regs, int_nvvm_sqrt_rp_d>;

// nvvm_sqrt intrinsic
def : Pat<(int_nvvm_sqrt_f Float32Regs:$a),
          (INT_NVVM_SQRT_RN_FTZ_F Float32Regs:$a)>, Requires<[doF32FTZ, do_SQRTF32_RN]>;
def : Pat<(int_nvvm_sqrt_f Float32Regs:$a),
          (INT_NVVM_SQRT_RN_F Float32Regs:$a)>, Requires<[do_SQRTF32_RN]>;
def : Pat<(int_nvvm_sqrt_f Float32Regs:$a),
          (INT_NVVM_SQRT_APPROX_FTZ_F Float32Regs:$a)>, Requires<[doF32FTZ]>;
def : Pat<(int_nvvm_sqrt_f Float32Regs:$a),
          (INT_NVVM_SQRT_APPROX_F Float32Regs:$a)>;

//
// Rsqrt
//

def INT_NVVM_RSQRT_APPROX_FTZ_F
  : F_MATH_1<"rsqrt.approx.ftz.f32 \t$dst, $src0;", Float32Regs, Float32Regs,
    int_nvvm_rsqrt_approx_ftz_f>;
def INT_NVVM_RSQRT_APPROX_F : F_MATH_1<"rsqrt.approx.f32 \t$dst, $src0;",
  Float32Regs, Float32Regs, int_nvvm_rsqrt_approx_f>;
def INT_NVVM_RSQRT_APPROX_D : F_MATH_1<"rsqrt.approx.f64 \t$dst, $src0;",
  Float64Regs, Float64Regs, int_nvvm_rsqrt_approx_d>;

//
// Add
//

def INT_NVVM_ADD_RN_FTZ_F : F_MATH_2<"add.rn.ftz.f32 \t$dst, $src0, $src1;",
  Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rn_ftz_f>;
def INT_NVVM_ADD_RN_F : F_MATH_2<"add.rn.f32 \t$dst, $src0, $src1;",
  Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rn_f>;
def INT_NVVM_ADD_RZ_FTZ_F : F_MATH_2<"add.rz.ftz.f32 \t$dst, $src0, $src1;",
  Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rz_ftz_f>;
def INT_NVVM_ADD_RZ_F : F_MATH_2<"add.rz.f32 \t$dst, $src0, $src1;",
  Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rz_f>;
def INT_NVVM_ADD_RM_FTZ_F : F_MATH_2<"add.rm.ftz.f32 \t$dst, $src0, $src1;",
  Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rm_ftz_f>;
def INT_NVVM_ADD_RM_F : F_MATH_2<"add.rm.f32 \t$dst, $src0, $src1;",
  Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rm_f>;
def INT_NVVM_ADD_RP_FTZ_F : F_MATH_2<"add.rp.ftz.f32 \t$dst, $src0, $src1;",
  Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rp_ftz_f>;
def INT_NVVM_ADD_RP_F : F_MATH_2<"add.rp.f32 \t$dst, $src0, $src1;",
  Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rp_f>;

def INT_NVVM_ADD_RN_D : F_MATH_2<"add.rn.f64 \t$dst, $src0, $src1;",
  Float64Regs, Float64Regs, Float64Regs, int_nvvm_add_rn_d>;
def INT_NVVM_ADD_RZ_D : F_MATH_2<"add.rz.f64 \t$dst, $src0, $src1;",
  Float64Regs, Float64Regs, Float64Regs, int_nvvm_add_rz_d>;
def INT_NVVM_ADD_RM_D : F_MATH_2<"add.rm.f64 \t$dst, $src0, $src1;",
  Float64Regs, Float64Regs, Float64Regs, int_nvvm_add_rm_d>;
def INT_NVVM_ADD_RP_D : F_MATH_2<"add.rp.f64 \t$dst, $src0, $src1;",
  Float64Regs, Float64Regs, Float64Regs, int_nvvm_add_rp_d>;

//
// Convert
//

def : Pat<(int_nvvm_d2f_rn_ftz Float64Regs:$a),
          (CVT_f32_f64 Float64Regs:$a, CvtRN_FTZ)>;
def : Pat<(int_nvvm_d2f_rn Float64Regs:$a),
          (CVT_f32_f64 Float64Regs:$a, CvtRN)>;
def : Pat<(int_nvvm_d2f_rz_ftz Float64Regs:$a),
          (CVT_f32_f64 Float64Regs:$a, CvtRZ_FTZ)>;
def : Pat<(int_nvvm_d2f_rz Float64Regs:$a),
          (CVT_f32_f64 Float64Regs:$a, CvtRZ)>;
def : Pat<(int_nvvm_d2f_rm_ftz Float64Regs:$a),
          (CVT_f32_f64 Float64Regs:$a, CvtRM_FTZ)>;
def : Pat<(int_nvvm_d2f_rm Float64Regs:$a),
          (CVT_f32_f64 Float64Regs:$a, CvtRM)>;
def : Pat<(int_nvvm_d2f_rp_ftz Float64Regs:$a),
          (CVT_f32_f64 Float64Regs:$a, CvtRP_FTZ)>;
def : Pat<(int_nvvm_d2f_rp Float64Regs:$a),
          (CVT_f32_f64 Float64Regs:$a, CvtRP)>;

def : Pat<(int_nvvm_d2i_rn Float64Regs:$a),
          (CVT_s32_f64 Float64Regs:$a, CvtRNI)>;
def : Pat<(int_nvvm_d2i_rz Float64Regs:$a),
          (CVT_s32_f64 Float64Regs:$a, CvtRZI)>;
def : Pat<(int_nvvm_d2i_rm Float64Regs:$a),
          (CVT_s32_f64 Float64Regs:$a, CvtRMI)>;
def : Pat<(int_nvvm_d2i_rp Float64Regs:$a),
          (CVT_s32_f64 Float64Regs:$a, CvtRPI)>;

def : Pat<(int_nvvm_d2ui_rn Float64Regs:$a),
          (CVT_u32_f64 Float64Regs:$a, CvtRNI)>;
def : Pat<(int_nvvm_d2ui_rz Float64Regs:$a),
          (CVT_u32_f64 Float64Regs:$a, CvtRZI)>;
def : Pat<(int_nvvm_d2ui_rm Float64Regs:$a),
          (CVT_u32_f64 Float64Regs:$a, CvtRMI)>;
def : Pat<(int_nvvm_d2ui_rp Float64Regs:$a),
          (CVT_u32_f64 Float64Regs:$a, CvtRPI)>;

def : Pat<(int_nvvm_i2d_rn Int32Regs:$a),
          (CVT_f64_s32 Int32Regs:$a, CvtRN)>;
def : Pat<(int_nvvm_i2d_rz Int32Regs:$a),
          (CVT_f64_s32 Int32Regs:$a, CvtRZ)>;
def : Pat<(int_nvvm_i2d_rm Int32Regs:$a),
          (CVT_f64_s32 Int32Regs:$a, CvtRM)>;
def : Pat<(int_nvvm_i2d_rp Int32Regs:$a),
          (CVT_f64_s32 Int32Regs:$a, CvtRP)>;

def : Pat<(int_nvvm_ui2d_rn Int32Regs:$a),
          (CVT_f64_u32 Int32Regs:$a, CvtRN)>;
def : Pat<(int_nvvm_ui2d_rz Int32Regs:$a),
          (CVT_f64_u32 Int32Regs:$a, CvtRZ)>;
def : Pat<(int_nvvm_ui2d_rm Int32Regs:$a),
          (CVT_f64_u32 Int32Regs:$a, CvtRM)>;
def : Pat<(int_nvvm_ui2d_rp Int32Regs:$a),
          (CVT_f64_u32 Int32Regs:$a, CvtRP)>;

def : Pat<(int_nvvm_f2i_rn_ftz Float32Regs:$a),
          (CVT_s32_f32 Float32Regs:$a, CvtRNI_FTZ)>;
def : Pat<(int_nvvm_f2i_rn Float32Regs:$a),
          (CVT_s32_f32 Float32Regs:$a, CvtRNI)>;
def : Pat<(int_nvvm_f2i_rz_ftz Float32Regs:$a),
          (CVT_s32_f32 Float32Regs:$a, CvtRZI_FTZ)>;
def : Pat<(int_nvvm_f2i_rz Float32Regs:$a),
          (CVT_s32_f32 Float32Regs:$a, CvtRZI)>;
def : Pat<(int_nvvm_f2i_rm_ftz Float32Regs:$a),
          (CVT_s32_f32 Float32Regs:$a, CvtRMI_FTZ)>;
def : Pat<(int_nvvm_f2i_rm Float32Regs:$a),
          (CVT_s32_f32 Float32Regs:$a, CvtRMI)>;
def : Pat<(int_nvvm_f2i_rp_ftz Float32Regs:$a),
          (CVT_s32_f32 Float32Regs:$a, CvtRPI_FTZ)>;
def : Pat<(int_nvvm_f2i_rp Float32Regs:$a),
          (CVT_s32_f32 Float32Regs:$a, CvtRPI)>;

def : Pat<(int_nvvm_f2ui_rn_ftz Float32Regs:$a),
          (CVT_u32_f32 Float32Regs:$a, CvtRNI_FTZ)>;
def : Pat<(int_nvvm_f2ui_rn Float32Regs:$a),
          (CVT_u32_f32 Float32Regs:$a, CvtRNI)>;
def : Pat<(int_nvvm_f2ui_rz_ftz Float32Regs:$a),
          (CVT_u32_f32 Float32Regs:$a, CvtRZI_FTZ)>;
def : Pat<(int_nvvm_f2ui_rz Float32Regs:$a),
          (CVT_u32_f32 Float32Regs:$a, CvtRZI)>;
def : Pat<(int_nvvm_f2ui_rm_ftz Float32Regs:$a),
          (CVT_u32_f32 Float32Regs:$a, CvtRMI_FTZ)>;
def : Pat<(int_nvvm_f2ui_rm Float32Regs:$a),
          (CVT_u32_f32 Float32Regs:$a, CvtRMI)>;
def : Pat<(int_nvvm_f2ui_rp_ftz Float32Regs:$a),
          (CVT_u32_f32 Float32Regs:$a, CvtRPI_FTZ)>;
def : Pat<(int_nvvm_f2ui_rp Float32Regs:$a),
          (CVT_u32_f32 Float32Regs:$a, CvtRPI)>;

def : Pat<(int_nvvm_i2f_rn Int32Regs:$a),
          (CVT_f32_s32 Int32Regs:$a, CvtRN)>;
def : Pat<(int_nvvm_i2f_rz Int32Regs:$a),
          (CVT_f32_s32 Int32Regs:$a, CvtRZ)>;
def : Pat<(int_nvvm_i2f_rm Int32Regs:$a),
          (CVT_f32_s32 Int32Regs:$a, CvtRM)>;
def : Pat<(int_nvvm_i2f_rp Int32Regs:$a),
          (CVT_f32_s32 Int32Regs:$a, CvtRP)>;

def : Pat<(int_nvvm_ui2f_rn Int32Regs:$a),
          (CVT_f32_u32 Int32Regs:$a, CvtRN)>;
def : Pat<(int_nvvm_ui2f_rz Int32Regs:$a),
          (CVT_f32_u32 Int32Regs:$a, CvtRZ)>;
def : Pat<(int_nvvm_ui2f_rm Int32Regs:$a),
          (CVT_f32_u32 Int32Regs:$a, CvtRM)>;
def : Pat<(int_nvvm_ui2f_rp Int32Regs:$a),
          (CVT_f32_u32 Int32Regs:$a, CvtRP)>;

def INT_NVVM_LOHI_I2D : F_MATH_2<"mov.b64 \t$dst, {{$src0, $src1}};",
  Float64Regs, Int32Regs, Int32Regs, int_nvvm_lohi_i2d>;

def INT_NVVM_D2I_LO : F_MATH_1<!strconcat("{{\n\t",
                       !strconcat(".reg .b32 %temp; \n\t",
             !strconcat("mov.b64 \t{$dst, %temp}, $src0;\n\t",
               "}}"))),
             Int32Regs, Float64Regs, int_nvvm_d2i_lo>;
def INT_NVVM_D2I_HI : F_MATH_1<!strconcat("{{\n\t",
                       !strconcat(".reg .b32 %temp; \n\t",
                         !strconcat("mov.b64 \t{%temp, $dst}, $src0;\n\t",
                           "}}"))),
             Int32Regs, Float64Regs, int_nvvm_d2i_hi>;

def : Pat<(int_nvvm_f2ll_rn_ftz Float32Regs:$a),
          (CVT_s64_f32 Float32Regs:$a, CvtRNI_FTZ)>;
def : Pat<(int_nvvm_f2ll_rn Float32Regs:$a),
          (CVT_s64_f32 Float32Regs:$a, CvtRNI)>;
def : Pat<(int_nvvm_f2ll_rz_ftz Float32Regs:$a),
          (CVT_s64_f32 Float32Regs:$a, CvtRZI_FTZ)>;
def : Pat<(int_nvvm_f2ll_rz Float32Regs:$a),
          (CVT_s64_f32 Float32Regs:$a, CvtRZI)>;
def : Pat<(int_nvvm_f2ll_rm_ftz Float32Regs:$a),
          (CVT_s64_f32 Float32Regs:$a, CvtRMI_FTZ)>;
def : Pat<(int_nvvm_f2ll_rm Float32Regs:$a),
          (CVT_s64_f32 Float32Regs:$a, CvtRMI)>;
def : Pat<(int_nvvm_f2ll_rp_ftz Float32Regs:$a),
          (CVT_s64_f32 Float32Regs:$a, CvtRPI_FTZ)>;
def : Pat<(int_nvvm_f2ll_rp Float32Regs:$a),
          (CVT_s64_f32 Float32Regs:$a, CvtRPI)>;

def : Pat<(int_nvvm_f2ull_rn_ftz Float32Regs:$a),
          (CVT_u64_f32 Float32Regs:$a, CvtRNI_FTZ)>;
def : Pat<(int_nvvm_f2ull_rn Float32Regs:$a),
          (CVT_u64_f32 Float32Regs:$a, CvtRNI)>;
def : Pat<(int_nvvm_f2ull_rz_ftz Float32Regs:$a),
          (CVT_u64_f32 Float32Regs:$a, CvtRZI_FTZ)>;
def : Pat<(int_nvvm_f2ull_rz Float32Regs:$a),
          (CVT_u64_f32 Float32Regs:$a, CvtRZI)>;
def : Pat<(int_nvvm_f2ull_rm_ftz Float32Regs:$a),
          (CVT_u64_f32 Float32Regs:$a, CvtRMI_FTZ)>;
def : Pat<(int_nvvm_f2ull_rm Float32Regs:$a),
          (CVT_u64_f32 Float32Regs:$a, CvtRMI)>;
def : Pat<(int_nvvm_f2ull_rp_ftz Float32Regs:$a),
          (CVT_u64_f32 Float32Regs:$a, CvtRPI_FTZ)>;
def : Pat<(int_nvvm_f2ull_rp Float32Regs:$a),
          (CVT_u64_f32 Float32Regs:$a, CvtRPI)>;

def : Pat<(int_nvvm_d2ll_rn Float64Regs:$a),
          (CVT_s64_f64 Float64Regs:$a, CvtRNI)>;
def : Pat<(int_nvvm_d2ll_rz Float64Regs:$a),
          (CVT_s64_f64 Float64Regs:$a, CvtRZI)>;
def : Pat<(int_nvvm_d2ll_rm Float64Regs:$a),
          (CVT_s64_f64 Float64Regs:$a, CvtRMI)>;
def : Pat<(int_nvvm_d2ll_rp Float64Regs:$a),
          (CVT_s64_f64 Float64Regs:$a, CvtRPI)>;

def : Pat<(int_nvvm_d2ull_rn Float64Regs:$a),
          (CVT_u64_f64 Float64Regs:$a, CvtRNI)>;
def : Pat<(int_nvvm_d2ull_rz Float64Regs:$a),
          (CVT_u64_f64 Float64Regs:$a, CvtRZI)>;
def : Pat<(int_nvvm_d2ull_rm Float64Regs:$a),
          (CVT_u64_f64 Float64Regs:$a, CvtRMI)>;
def : Pat<(int_nvvm_d2ull_rp Float64Regs:$a),
          (CVT_u64_f64 Float64Regs:$a, CvtRPI)>;

def : Pat<(int_nvvm_ll2f_rn Int64Regs:$a),
          (CVT_f32_s64 Int64Regs:$a, CvtRN)>;
def : Pat<(int_nvvm_ll2f_rz Int64Regs:$a),
          (CVT_f32_s64 Int64Regs:$a, CvtRZ)>;
def : Pat<(int_nvvm_ll2f_rm Int64Regs:$a),
          (CVT_f32_s64 Int64Regs:$a, CvtRM)>;
def : Pat<(int_nvvm_ll2f_rp Int64Regs:$a),
          (CVT_f32_s64 Int64Regs:$a, CvtRP)>;

def : Pat<(int_nvvm_ull2f_rn Int64Regs:$a),
          (CVT_f32_u64 Int64Regs:$a, CvtRN)>;
def : Pat<(int_nvvm_ull2f_rz Int64Regs:$a),
          (CVT_f32_u64 Int64Regs:$a, CvtRZ)>;
def : Pat<(int_nvvm_ull2f_rm Int64Regs:$a),
          (CVT_f32_u64 Int64Regs:$a, CvtRM)>;
def : Pat<(int_nvvm_ull2f_rp Int64Regs:$a),
          (CVT_f32_u64 Int64Regs:$a, CvtRP)>;

def : Pat<(int_nvvm_ll2d_rn Int64Regs:$a),
          (CVT_f64_s64 Int64Regs:$a, CvtRN)>;
def : Pat<(int_nvvm_ll2d_rz Int64Regs:$a),
          (CVT_f64_s64 Int64Regs:$a, CvtRZ)>;
def : Pat<(int_nvvm_ll2d_rm Int64Regs:$a),
          (CVT_f64_s64 Int64Regs:$a, CvtRM)>;
def : Pat<(int_nvvm_ll2d_rp Int64Regs:$a),
          (CVT_f64_s64 Int64Regs:$a, CvtRP)>;

def : Pat<(int_nvvm_ull2d_rn Int64Regs:$a),
          (CVT_f64_u64 Int64Regs:$a, CvtRN)>;
def : Pat<(int_nvvm_ull2d_rz Int64Regs:$a),
          (CVT_f64_u64 Int64Regs:$a, CvtRZ)>;
def : Pat<(int_nvvm_ull2d_rm Int64Regs:$a),
          (CVT_f64_u64 Int64Regs:$a, CvtRM)>;
def : Pat<(int_nvvm_ull2d_rp Int64Regs:$a),
          (CVT_f64_u64 Int64Regs:$a, CvtRP)>;


// FIXME: Ideally, we could use these patterns instead of the scope-creating
// patterns, but ptxas does not like these since .s16 is not compatible with
// .f16.  The solution is to use .bXX for all integer register types, but we
// are not there yet.
//def : Pat<(int_nvvm_f2h_rn_ftz Float32Regs:$a),
//          (CVT_f16_f32 Float32Regs:$a, CvtRN_FTZ)>;
//def : Pat<(int_nvvm_f2h_rn Float32Regs:$a),
//          (CVT_f16_f32 Float32Regs:$a, CvtRN)>;
//
//def : Pat<(int_nvvm_h2f Int16Regs:$a),
//          (CVT_f32_f16 Int16Regs:$a, CvtNONE)>;

def INT_NVVM_F2H_RN_FTZ : F_MATH_1<!strconcat("{{\n\t",
                                   !strconcat(".reg .b16 %temp;\n\t",
           !strconcat("cvt.rn.ftz.f16.f32 \t%temp, $src0;\n\t",
           !strconcat("mov.b16 \t$dst, %temp;\n",
             "}}")))),
                                   Int16Regs, Float32Regs, int_nvvm_f2h_rn_ftz>;
def INT_NVVM_F2H_RN : F_MATH_1<!strconcat("{{\n\t",
                                   !strconcat(".reg .b16 %temp;\n\t",
           !strconcat("cvt.rn.f16.f32 \t%temp, $src0;\n\t",
           !strconcat("mov.b16 \t$dst, %temp;\n",
             "}}")))),
           Int16Regs, Float32Regs, int_nvvm_f2h_rn>;

def INT_NVVM_H2F : F_MATH_1<!strconcat("{{\n\t",
                            !strconcat(".reg .b16 %temp;\n\t",
          !strconcat("mov.b16 \t%temp, $src0;\n\t",
          !strconcat("cvt.f32.f16 \t$dst, %temp;\n\t",
            "}}")))),
          Float32Regs, Int16Regs, int_nvvm_h2f>;

def : Pat<(f32 (f16_to_f32 Int16Regs:$a)),
          (CVT_f32_f16 Int16Regs:$a, CvtNONE)>;
def : Pat<(i16 (f32_to_f16 Float32Regs:$a)),
          (CVT_f16_f32 Float32Regs:$a, CvtRN_FTZ)>, Requires<[doF32FTZ]>;
def : Pat<(i16 (f32_to_f16 Float32Regs:$a)),
          (CVT_f16_f32 Float32Regs:$a, CvtRN)>;

//
// Bitcast
//

def INT_NVVM_BITCAST_F2I : F_MATH_1<"mov.b32 \t$dst, $src0;", Int32Regs,
  Float32Regs, int_nvvm_bitcast_f2i>;
def INT_NVVM_BITCAST_I2F : F_MATH_1<"mov.b32 \t$dst, $src0;", Float32Regs,
  Int32Regs, int_nvvm_bitcast_i2f>;

def INT_NVVM_BITCAST_LL2D : F_MATH_1<"mov.b64 \t$dst, $src0;", Float64Regs,
  Int64Regs, int_nvvm_bitcast_ll2d>;
def INT_NVVM_BITCAST_D2LL : F_MATH_1<"mov.b64 \t$dst, $src0;", Int64Regs,
  Float64Regs, int_nvvm_bitcast_d2ll>;

//-----------------------------------
// Atomic Functions
//-----------------------------------

class ATOMIC_GLOBAL_CHK <dag ops, dag frag>
 : PatFrag<ops, frag, [{
   return ChkMemSDNodeAddressSpace(N, llvm::ADDRESS_SPACE_GLOBAL);
}]>;
class ATOMIC_SHARED_CHK <dag ops, dag frag>
 : PatFrag<ops, frag, [{
   return ChkMemSDNodeAddressSpace(N, llvm::ADDRESS_SPACE_SHARED);
}]>;
class ATOMIC_GENERIC_CHK <dag ops, dag frag>
 : PatFrag<ops, frag, [{
   return ChkMemSDNodeAddressSpace(N, llvm::ADDRESS_SPACE_GENERIC);
}]>;

multiclass F_ATOMIC_2_imp<NVPTXRegClass ptrclass, NVPTXRegClass regclass,
  string SpaceStr, string TypeStr, string OpcStr, PatFrag IntOp,
  Operand IMMType, SDNode IMM, Predicate Pred> {
  def reg : NVPTXInst<(outs regclass:$dst), (ins ptrclass:$addr, regclass:$b),
               !strconcat("atom",
         !strconcat(SpaceStr,
         !strconcat(OpcStr,
         !strconcat(TypeStr,
         !strconcat(" \t$dst, [$addr], $b;", ""))))),
         [(set regclass:$dst, (IntOp ptrclass:$addr, regclass:$b))]>,
  Requires<[Pred]>;
  def imm : NVPTXInst<(outs regclass:$dst), (ins ptrclass:$addr, IMMType:$b),
               !strconcat("atom",
         !strconcat(SpaceStr,
         !strconcat(OpcStr,
         !strconcat(TypeStr,
         !strconcat(" \t$dst, [$addr], $b;", ""))))),
         [(set regclass:$dst, (IntOp ptrclass:$addr, IMM:$b))]>,
  Requires<[Pred]>;
}
multiclass F_ATOMIC_2<NVPTXRegClass regclass, string SpaceStr, string TypeStr,
  string OpcStr, PatFrag IntOp, Operand IMMType, SDNode IMM, Predicate Pred> {
  defm p32 : F_ATOMIC_2_imp<Int32Regs, regclass, SpaceStr, TypeStr, OpcStr,
    IntOp, IMMType, IMM, Pred>;
  defm p64 : F_ATOMIC_2_imp<Int64Regs, regclass, SpaceStr, TypeStr, OpcStr,
    IntOp, IMMType, IMM, Pred>;
}

// has 2 operands, neg the second one
multiclass F_ATOMIC_2_NEG_imp<NVPTXRegClass ptrclass, NVPTXRegClass regclass,
  string SpaceStr, string TypeStr, string OpcStr, PatFrag IntOp,
  Operand IMMType, Predicate Pred> {
  def reg : NVPTXInst<(outs regclass:$dst), (ins ptrclass:$addr, regclass:$b),
    !strconcat("{{ \n\t",
         !strconcat(".reg \t.s",
         !strconcat(TypeStr,
         !strconcat(" temp; \n\t",
         !strconcat("neg.s",
         !strconcat(TypeStr,
         !strconcat(" \ttemp, $b; \n\t",
               !strconcat("atom",
         !strconcat(SpaceStr,
         !strconcat(OpcStr,
         !strconcat(".u",
         !strconcat(TypeStr,
         !strconcat(" \t$dst, [$addr], temp; \n\t",
           !strconcat("}}", "")))))))))))))),
         [(set regclass:$dst, (IntOp ptrclass:$addr, regclass:$b))]>,
  Requires<[Pred]>;
}
multiclass F_ATOMIC_2_NEG<NVPTXRegClass regclass, string SpaceStr,
  string TypeStr, string OpcStr, PatFrag IntOp, Operand IMMType,
  Predicate Pred> {
 defm p32: F_ATOMIC_2_NEG_imp<Int32Regs, regclass, SpaceStr, TypeStr, OpcStr,
   IntOp, IMMType, Pred> ;
 defm p64: F_ATOMIC_2_NEG_imp<Int64Regs, regclass, SpaceStr, TypeStr, OpcStr,
   IntOp, IMMType, Pred> ;
}

// has 3 operands
multiclass F_ATOMIC_3_imp<NVPTXRegClass ptrclass, NVPTXRegClass regclass,
  string SpaceStr, string TypeStr, string OpcStr, PatFrag IntOp,
  Operand IMMType, Predicate Pred> {
  def reg : NVPTXInst<(outs regclass:$dst),
    (ins ptrclass:$addr, regclass:$b, regclass:$c),
               !strconcat("atom",
         !strconcat(SpaceStr,
         !strconcat(OpcStr,
         !strconcat(TypeStr,
         !strconcat(" \t$dst, [$addr], $b, $c;", ""))))),
         [(set regclass:$dst,
           (IntOp ptrclass:$addr, regclass:$b, regclass:$c))]>,
         Requires<[Pred]>;
  def imm1 : NVPTXInst<(outs regclass:$dst),
    (ins ptrclass:$addr, IMMType:$b, regclass:$c),
               !strconcat("atom",
         !strconcat(SpaceStr,
         !strconcat(OpcStr,
         !strconcat(TypeStr,
         !strconcat(" \t$dst, [$addr], $b, $c;", ""))))),
         [(set regclass:$dst, (IntOp ptrclass:$addr, imm:$b, regclass:$c))]>,
  Requires<[Pred]>;
  def imm2 : NVPTXInst<(outs regclass:$dst),
    (ins ptrclass:$addr, regclass:$b, IMMType:$c),
               !strconcat("atom",
         !strconcat(SpaceStr,
         !strconcat(OpcStr,
         !strconcat(TypeStr,
         !strconcat(" \t$dst, [$addr], $b, $c;", ""))))),
         [(set regclass:$dst, (IntOp ptrclass:$addr, regclass:$b, imm:$c))]>,
  Requires<[Pred]>;
  def imm3 : NVPTXInst<(outs regclass:$dst),
    (ins ptrclass:$addr, IMMType:$b, IMMType:$c),
               !strconcat("atom",
         !strconcat(SpaceStr,
         !strconcat(OpcStr,
         !strconcat(TypeStr,
         !strconcat(" \t$dst, [$addr], $b, $c;", ""))))),
         [(set regclass:$dst, (IntOp ptrclass:$addr, imm:$b, imm:$c))]>,
  Requires<[Pred]>;
}
multiclass F_ATOMIC_3<NVPTXRegClass regclass, string SpaceStr, string TypeStr,
  string OpcStr, PatFrag IntOp, Operand IMMType, Predicate Pred> {
  defm p32 : F_ATOMIC_3_imp<Int32Regs, regclass, SpaceStr, TypeStr, OpcStr,
    IntOp, IMMType, Pred>;
  defm p64 : F_ATOMIC_3_imp<Int64Regs, regclass, SpaceStr, TypeStr, OpcStr,
    IntOp, IMMType, Pred>;
}

// atom_add

def atomic_load_add_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
  (atomic_load_add_32 node:$a, node:$b)>;
def atomic_load_add_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
  (atomic_load_add_32 node:$a, node:$b)>;
def atomic_load_add_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
  (atomic_load_add_32 node:$a, node:$b)>;
def atomic_load_add_64_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
  (atomic_load_add_64 node:$a, node:$b)>;
def atomic_load_add_64_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
  (atomic_load_add_64 node:$a, node:$b)>;
def atomic_load_add_64_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
  (atomic_load_add_64 node:$a, node:$b)>;
def atomic_load_add_f32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
  (int_nvvm_atomic_load_add_f32 node:$a, node:$b)>;
def atomic_load_add_f32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
  (int_nvvm_atomic_load_add_f32 node:$a, node:$b)>;
def atomic_load_add_f32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
  (int_nvvm_atomic_load_add_f32 node:$a, node:$b)>;

defm INT_PTX_ATOM_ADD_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".u32", ".add",
  atomic_load_add_32_g, i32imm, imm, hasAtomRedG32>;
defm INT_PTX_ATOM_ADD_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".u32", ".add",
  atomic_load_add_32_s, i32imm, imm, hasAtomRedS32>;
defm INT_PTX_ATOM_ADD_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".u32", ".add",
  atomic_load_add_32_gen, i32imm, imm, hasAtomRedGen32>;
defm INT_PTX_ATOM_ADD_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".u32",
  ".add", atomic_load_add_32_gen, i32imm, imm, useAtomRedG32forGen32>;

defm INT_PTX_ATOM_ADD_G_64 : F_ATOMIC_2<Int64Regs, ".global", ".u64", ".add",
  atomic_load_add_64_g, i64imm, imm, hasAtomRedG64>;
defm INT_PTX_ATOM_ADD_S_64 : F_ATOMIC_2<Int64Regs, ".shared", ".u64", ".add",
  atomic_load_add_64_s, i64imm, imm, hasAtomRedS64>;
defm INT_PTX_ATOM_ADD_GEN_64 : F_ATOMIC_2<Int64Regs, "", ".u64", ".add",
  atomic_load_add_64_gen, i64imm, imm, hasAtomRedGen64>;
defm INT_PTX_ATOM_ADD_GEN_64_USE_G : F_ATOMIC_2<Int64Regs, ".global", ".u64",
  ".add", atomic_load_add_64_gen, i64imm, imm, useAtomRedG64forGen64>;

defm INT_PTX_ATOM_ADD_G_F32 : F_ATOMIC_2<Float32Regs, ".global", ".f32", ".add",
  atomic_load_add_f32_g, f32imm, fpimm, hasAtomAddF32>;
defm INT_PTX_ATOM_ADD_S_F32 : F_ATOMIC_2<Float32Regs, ".shared", ".f32", ".add",
  atomic_load_add_f32_s, f32imm, fpimm, hasAtomAddF32>;
defm INT_PTX_ATOM_ADD_GEN_F32 : F_ATOMIC_2<Float32Regs, "", ".f32", ".add",
  atomic_load_add_f32_gen, f32imm, fpimm, hasAtomAddF32>;

// atom_sub

def atomic_load_sub_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
  (atomic_load_sub_32 node:$a, node:$b)>;
def atomic_load_sub_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
  (atomic_load_sub_32 node:$a, node:$b)>;
def atomic_load_sub_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
  (atomic_load_sub_32 node:$a, node:$b)>;
def atomic_load_sub_64_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
  (atomic_load_sub_64 node:$a, node:$b)>;
def atomic_load_sub_64_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
  (atomic_load_sub_64 node:$a, node:$b)>;
def atomic_load_sub_64_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
  (atomic_load_sub_64 node:$a, node:$b)>;

defm INT_PTX_ATOM_SUB_G_32 : F_ATOMIC_2_NEG<Int32Regs, ".global", "32", ".add",
  atomic_load_sub_32_g, i32imm, hasAtomRedG32>;
defm INT_PTX_ATOM_SUB_G_64 : F_ATOMIC_2_NEG<Int64Regs, ".global", "64", ".add",
  atomic_load_sub_64_g, i64imm, hasAtomRedG64>;
defm INT_PTX_ATOM_SUB_GEN_32 : F_ATOMIC_2_NEG<Int32Regs, "", "32", ".add",
  atomic_load_sub_32_gen, i32imm, hasAtomRedGen32>;
defm INT_PTX_ATOM_SUB_GEN_32_USE_G : F_ATOMIC_2_NEG<Int32Regs, ".global", "32",
  ".add", atomic_load_sub_32_gen, i32imm, useAtomRedG32forGen32>;
defm INT_PTX_ATOM_SUB_S_32 : F_ATOMIC_2_NEG<Int32Regs, ".shared", "32", ".add",
  atomic_load_sub_32_s, i32imm, hasAtomRedS32>;
defm INT_PTX_ATOM_SUB_S_64 : F_ATOMIC_2_NEG<Int64Regs, ".shared", "64", ".add",
  atomic_load_sub_64_s, i64imm, hasAtomRedS64>;
defm INT_PTX_ATOM_SUB_GEN_64 : F_ATOMIC_2_NEG<Int64Regs, "", "64", ".add",
  atomic_load_sub_64_gen, i64imm, hasAtomRedGen64>;
defm INT_PTX_ATOM_SUB_GEN_64_USE_G : F_ATOMIC_2_NEG<Int64Regs, ".global", "64",
  ".add", atomic_load_sub_64_gen, i64imm, useAtomRedG64forGen64>;

// atom_swap

def atomic_swap_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
  (atomic_swap_32 node:$a, node:$b)>;
def atomic_swap_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
  (atomic_swap_32 node:$a, node:$b)>;
def atomic_swap_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
  (atomic_swap_32 node:$a, node:$b)>;
def atomic_swap_64_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
  (atomic_swap_64 node:$a, node:$b)>;
def atomic_swap_64_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
  (atomic_swap_64 node:$a, node:$b)>;
def atomic_swap_64_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
  (atomic_swap_64 node:$a, node:$b)>;

defm INT_PTX_ATOM_SWAP_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".b32", ".exch",
  atomic_swap_32_g, i32imm, imm, hasAtomRedG32>;
defm INT_PTX_ATOM_SWAP_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".b32", ".exch",
  atomic_swap_32_s, i32imm, imm, hasAtomRedS32>;
defm INT_PTX_ATOM_SWAP_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".b32", ".exch",
  atomic_swap_32_gen, i32imm, imm, hasAtomRedGen32>;
defm INT_PTX_ATOM_SWAP_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".b32",
  ".exch", atomic_swap_32_gen, i32imm, imm, useAtomRedG32forGen32>;
defm INT_PTX_ATOM_SWAP_G_64 : F_ATOMIC_2<Int64Regs, ".global", ".b64", ".exch",
  atomic_swap_64_g, i64imm, imm, hasAtomRedG64>;
defm INT_PTX_ATOM_SWAP_S_64 : F_ATOMIC_2<Int64Regs, ".shared", ".b64", ".exch",
  atomic_swap_64_s, i64imm, imm, hasAtomRedS64>;
defm INT_PTX_ATOM_SWAP_GEN_64 : F_ATOMIC_2<Int64Regs, "", ".b64", ".exch",
  atomic_swap_64_gen, i64imm, imm, hasAtomRedGen64>;
defm INT_PTX_ATOM_SWAP_GEN_64_USE_G : F_ATOMIC_2<Int64Regs, ".global", ".b64",
  ".exch", atomic_swap_64_gen, i64imm, imm, useAtomRedG64forGen64>;

// atom_max

def atomic_load_max_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b)
  , (atomic_load_max_32 node:$a, node:$b)>;
def atomic_load_max_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
  (atomic_load_max_32 node:$a, node:$b)>;
def atomic_load_max_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
  (atomic_load_max_32 node:$a, node:$b)>;
def atomic_load_umax_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
  (atomic_load_umax_32 node:$a, node:$b)>;
def atomic_load_umax_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
  (atomic_load_umax_32 node:$a, node:$b)>;
def atomic_load_umax_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
  (atomic_load_umax_32 node:$a, node:$b)>;

defm INT_PTX_ATOM_LOAD_MAX_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".s32",
  ".max", atomic_load_max_32_g, i32imm, imm, hasAtomRedG32>;
defm INT_PTX_ATOM_LOAD_MAX_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".s32",
  ".max", atomic_load_max_32_s, i32imm, imm, hasAtomRedS32>;
defm INT_PTX_ATOM_LOAD_MAX_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".s32", ".max",
  atomic_load_max_32_gen, i32imm, imm, hasAtomRedGen32>;
defm INT_PTX_ATOM_LOAD_MAX_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global",
  ".s32", ".max", atomic_load_max_32_gen, i32imm, imm, useAtomRedG32forGen32>;
defm INT_PTX_ATOM_LOAD_UMAX_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".u32",
  ".max", atomic_load_umax_32_g, i32imm, imm, hasAtomRedG32>;
defm INT_PTX_ATOM_LOAD_UMAX_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".u32",
  ".max", atomic_load_umax_32_s, i32imm, imm, hasAtomRedS32>;
defm INT_PTX_ATOM_LOAD_UMAX_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".u32", ".max",
  atomic_load_umax_32_gen, i32imm, imm, hasAtomRedGen32>;
defm INT_PTX_ATOM_LOAD_UMAX_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global",
  ".u32", ".max", atomic_load_umax_32_gen, i32imm, imm, useAtomRedG32forGen32>;

// atom_min

def atomic_load_min_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
  (atomic_load_min_32 node:$a, node:$b)>;
def atomic_load_min_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
  (atomic_load_min_32 node:$a, node:$b)>;
def atomic_load_min_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
  (atomic_load_min_32 node:$a, node:$b)>;
def atomic_load_umin_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
  (atomic_load_umin_32 node:$a, node:$b)>;
def atomic_load_umin_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
  (atomic_load_umin_32 node:$a, node:$b)>;
def atomic_load_umin_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
  (atomic_load_umin_32 node:$a, node:$b)>;

defm INT_PTX_ATOM_LOAD_MIN_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".s32",
  ".min", atomic_load_min_32_g, i32imm, imm, hasAtomRedG32>;
defm INT_PTX_ATOM_LOAD_MIN_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".s32",
  ".min", atomic_load_min_32_s, i32imm, imm, hasAtomRedS32>;
defm INT_PTX_ATOM_LOAD_MIN_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".s32", ".min",
  atomic_load_min_32_gen, i32imm, imm, hasAtomRedGen32>;
defm INT_PTX_ATOM_LOAD_MIN_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global",
  ".s32", ".min", atomic_load_min_32_gen, i32imm, imm, useAtomRedG32forGen32>;
defm INT_PTX_ATOM_LOAD_UMIN_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".u32",
  ".min", atomic_load_umin_32_g, i32imm, imm, hasAtomRedG32>;
defm INT_PTX_ATOM_LOAD_UMIN_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".u32",
  ".min", atomic_load_umin_32_s, i32imm, imm, hasAtomRedS32>;
defm INT_PTX_ATOM_LOAD_UMIN_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".u32", ".min",
  atomic_load_umin_32_gen, i32imm, imm, hasAtomRedGen32>;
defm INT_PTX_ATOM_LOAD_UMIN_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global",
  ".u32", ".min", atomic_load_umin_32_gen, i32imm, imm, useAtomRedG32forGen32>;

// atom_inc  atom_dec

def atomic_load_inc_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
  (int_nvvm_atomic_load_inc_32 node:$a, node:$b)>;
def atomic_load_inc_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
  (int_nvvm_atomic_load_inc_32 node:$a, node:$b)>;
def atomic_load_inc_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
  (int_nvvm_atomic_load_inc_32 node:$a, node:$b)>;
def atomic_load_dec_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
  (int_nvvm_atomic_load_dec_32 node:$a, node:$b)>;
def atomic_load_dec_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
  (int_nvvm_atomic_load_dec_32 node:$a, node:$b)>;
def atomic_load_dec_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
  (int_nvvm_atomic_load_dec_32 node:$a, node:$b)>;

defm INT_PTX_ATOM_INC_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".u32", ".inc",
  atomic_load_inc_32_g, i32imm, imm, hasAtomRedG32>;
defm INT_PTX_ATOM_INC_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".u32", ".inc",
  atomic_load_inc_32_s, i32imm, imm, hasAtomRedS32>;
defm INT_PTX_ATOM_INC_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".u32", ".inc",
  atomic_load_inc_32_gen, i32imm, imm, hasAtomRedGen32>;
defm INT_PTX_ATOM_INC_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".u32",
  ".inc", atomic_load_inc_32_gen, i32imm, imm, useAtomRedG32forGen32>;
defm INT_PTX_ATOM_DEC_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".u32", ".dec",
  atomic_load_dec_32_g, i32imm, imm, hasAtomRedG32>;
defm INT_PTX_ATOM_DEC_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".u32", ".dec",
  atomic_load_dec_32_s, i32imm, imm, hasAtomRedS32>;
defm INT_PTX_ATOM_DEC_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".u32", ".dec",
  atomic_load_dec_32_gen, i32imm, imm, hasAtomRedGen32>;
defm INT_PTX_ATOM_DEC_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".u32",
  ".dec", atomic_load_dec_32_gen, i32imm, imm, useAtomRedG32forGen32>;

// atom_and

def atomic_load_and_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
  (atomic_load_and_32 node:$a, node:$b)>;
def atomic_load_and_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
  (atomic_load_and_32 node:$a, node:$b)>;
def atomic_load_and_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
  (atomic_load_and_32 node:$a, node:$b)>;

defm INT_PTX_ATOM_AND_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".b32", ".and",
  atomic_load_and_32_g, i32imm, imm, hasAtomRedG32>;
defm INT_PTX_ATOM_AND_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".b32", ".and",
  atomic_load_and_32_s, i32imm, imm, hasAtomRedS32>;
defm INT_PTX_ATOM_AND_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".b32", ".and",
  atomic_load_and_32_gen, i32imm, imm, hasAtomRedGen32>;
defm INT_PTX_ATOM_AND_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".b32",
  ".and", atomic_load_and_32_gen, i32imm, imm, useAtomRedG32forGen32>;

// atom_or

def atomic_load_or_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
  (atomic_load_or_32 node:$a, node:$b)>;
def atomic_load_or_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
  (atomic_load_or_32 node:$a, node:$b)>;
def atomic_load_or_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
  (atomic_load_or_32 node:$a, node:$b)>;

defm INT_PTX_ATOM_OR_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".b32", ".or",
  atomic_load_or_32_g, i32imm, imm, hasAtomRedG32>;
defm INT_PTX_ATOM_OR_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".b32", ".or",
  atomic_load_or_32_gen, i32imm, imm, hasAtomRedGen32>;
defm INT_PTX_ATOM_OR_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".b32",
  ".or", atomic_load_or_32_gen, i32imm, imm, useAtomRedG32forGen32>;
defm INT_PTX_ATOM_OR_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".b32", ".or",
  atomic_load_or_32_s, i32imm, imm, hasAtomRedS32>;

// atom_xor

def atomic_load_xor_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
  (atomic_load_xor_32 node:$a, node:$b)>;
def atomic_load_xor_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
  (atomic_load_xor_32 node:$a, node:$b)>;
def atomic_load_xor_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
  (atomic_load_xor_32 node:$a, node:$b)>;

defm INT_PTX_ATOM_XOR_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".b32", ".xor",
  atomic_load_xor_32_g, i32imm, imm, hasAtomRedG32>;
defm INT_PTX_ATOM_XOR_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".b32", ".xor",
  atomic_load_xor_32_s, i32imm, imm, hasAtomRedS32>;
defm INT_PTX_ATOM_XOR_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".b32", ".xor",
  atomic_load_xor_32_gen, i32imm, imm, hasAtomRedGen32>;
defm INT_PTX_ATOM_XOR_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".b32",
  ".xor", atomic_load_xor_32_gen, i32imm, imm, useAtomRedG32forGen32>;

// atom_cas

def atomic_cmp_swap_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b, node:$c),
  (atomic_cmp_swap_32 node:$a, node:$b, node:$c)>;
def atomic_cmp_swap_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b, node:$c),
  (atomic_cmp_swap_32 node:$a, node:$b, node:$c)>;
def atomic_cmp_swap_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b, node:$c),
  (atomic_cmp_swap_32 node:$a, node:$b, node:$c)>;
def atomic_cmp_swap_64_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b, node:$c),
  (atomic_cmp_swap_64 node:$a, node:$b, node:$c)>;
def atomic_cmp_swap_64_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b, node:$c),
  (atomic_cmp_swap_64 node:$a, node:$b, node:$c)>;
def atomic_cmp_swap_64_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b, node:$c),
  (atomic_cmp_swap_64 node:$a, node:$b, node:$c)>;

defm INT_PTX_ATOM_CAS_G_32 : F_ATOMIC_3<Int32Regs, ".global", ".b32", ".cas",
  atomic_cmp_swap_32_g, i32imm, hasAtomRedG32>;
defm INT_PTX_ATOM_CAS_S_32 : F_ATOMIC_3<Int32Regs, ".shared", ".b32", ".cas",
  atomic_cmp_swap_32_s, i32imm, hasAtomRedS32>;
defm INT_PTX_ATOM_CAS_GEN_32 : F_ATOMIC_3<Int32Regs, "", ".b32", ".cas",
  atomic_cmp_swap_32_gen, i32imm, hasAtomRedGen32>;
defm INT_PTX_ATOM_CAS_GEN_32_USE_G : F_ATOMIC_3<Int32Regs, ".global", ".b32",
  ".cas", atomic_cmp_swap_32_gen, i32imm, useAtomRedG32forGen32>;
defm INT_PTX_ATOM_CAS_G_64 : F_ATOMIC_3<Int64Regs, ".global", ".b64", ".cas",
  atomic_cmp_swap_64_g, i64imm, hasAtomRedG64>;
defm INT_PTX_ATOM_CAS_S_64 : F_ATOMIC_3<Int64Regs, ".shared", ".b64", ".cas",
  atomic_cmp_swap_64_s, i64imm, hasAtomRedS64>;
defm INT_PTX_ATOM_CAS_GEN_64 : F_ATOMIC_3<Int64Regs, "", ".b64", ".cas",
  atomic_cmp_swap_64_gen, i64imm, hasAtomRedGen64>;
defm INT_PTX_ATOM_CAS_GEN_64_USE_G : F_ATOMIC_3<Int64Regs, ".global", ".b64",
  ".cas", atomic_cmp_swap_64_gen, i64imm, useAtomRedG64forGen64>;


//-----------------------------------
// Read Special Registers
//-----------------------------------
class F_SREG<string OpStr, NVPTXRegClass regclassOut, Intrinsic IntOp> :
      NVPTXInst<(outs regclassOut:$dst), (ins),
               OpStr,
         [(set regclassOut:$dst, (IntOp))]>;

def INT_PTX_SREG_TID_X : F_SREG<"mov.u32 \t$dst, %tid.x;", Int32Regs,
  int_nvvm_read_ptx_sreg_tid_x>;
def INT_PTX_SREG_TID_Y : F_SREG<"mov.u32 \t$dst, %tid.y;", Int32Regs,
  int_nvvm_read_ptx_sreg_tid_y>;
def INT_PTX_SREG_TID_Z : F_SREG<"mov.u32 \t$dst, %tid.z;", Int32Regs,
  int_nvvm_read_ptx_sreg_tid_z>;

def INT_PTX_SREG_NTID_X : F_SREG<"mov.u32 \t$dst, %ntid.x;", Int32Regs,
  int_nvvm_read_ptx_sreg_ntid_x>;
def INT_PTX_SREG_NTID_Y : F_SREG<"mov.u32 \t$dst, %ntid.y;", Int32Regs,
  int_nvvm_read_ptx_sreg_ntid_y>;
def INT_PTX_SREG_NTID_Z : F_SREG<"mov.u32 \t$dst, %ntid.z;", Int32Regs,
  int_nvvm_read_ptx_sreg_ntid_z>;

def INT_PTX_SREG_CTAID_X : F_SREG<"mov.u32 \t$dst, %ctaid.x;", Int32Regs,
  int_nvvm_read_ptx_sreg_ctaid_x>;
def INT_PTX_SREG_CTAID_Y : F_SREG<"mov.u32 \t$dst, %ctaid.y;", Int32Regs,
  int_nvvm_read_ptx_sreg_ctaid_y>;
def INT_PTX_SREG_CTAID_Z : F_SREG<"mov.u32 \t$dst, %ctaid.z;", Int32Regs,
  int_nvvm_read_ptx_sreg_ctaid_z>;

def INT_PTX_SREG_NCTAID_X : F_SREG<"mov.u32 \t$dst, %nctaid.x;", Int32Regs,
  int_nvvm_read_ptx_sreg_nctaid_x>;
def INT_PTX_SREG_NCTAID_Y : F_SREG<"mov.u32 \t$dst, %nctaid.y;", Int32Regs,
  int_nvvm_read_ptx_sreg_nctaid_y>;
def INT_PTX_SREG_NCTAID_Z : F_SREG<"mov.u32 \t$dst, %nctaid.z;", Int32Regs,
  int_nvvm_read_ptx_sreg_nctaid_z>;

def INT_PTX_SREG_WARPSIZE : F_SREG<"mov.u32 \t$dst, WARP_SZ;", Int32Regs,
  int_nvvm_read_ptx_sreg_warpsize>;


//-----------------------------------
// Support for ldu on sm_20 or later
//-----------------------------------

def ldu_i8 : PatFrag<(ops node:$ptr), (int_nvvm_ldu_global_i node:$ptr), [{
  MemIntrinsicSDNode *M = cast<MemIntrinsicSDNode>(N);
  return M->getMemoryVT() == MVT::i8;
}]>;

// Scalar
// @TODO: Revisit this, Changed imemAny to imem
multiclass LDU_G<string TyStr, NVPTXRegClass regclass, Intrinsic IntOp> {
  def areg: NVPTXInst<(outs regclass:$result), (ins Int32Regs:$src),
               !strconcat("ldu.global.", TyStr),
         [(set regclass:$result, (IntOp Int32Regs:$src))]>, Requires<[hasLDU]>;
  def areg64: NVPTXInst<(outs regclass:$result), (ins Int64Regs:$src),
               !strconcat("ldu.global.", TyStr),
         [(set regclass:$result, (IntOp Int64Regs:$src))]>, Requires<[hasLDU]>;
 def avar:  NVPTXInst<(outs regclass:$result), (ins imem:$src),
               !strconcat("ldu.global.", TyStr),
                [(set regclass:$result, (IntOp (Wrapper tglobaladdr:$src)))]>,
                Requires<[hasLDU]>;
 def ari :  NVPTXInst<(outs regclass:$result), (ins MEMri:$src),
               !strconcat("ldu.global.", TyStr),
         [(set regclass:$result, (IntOp ADDRri:$src))]>, Requires<[hasLDU]>;
 def ari64 :  NVPTXInst<(outs regclass:$result), (ins MEMri64:$src),
               !strconcat("ldu.global.", TyStr),
         [(set regclass:$result, (IntOp ADDRri64:$src))]>, Requires<[hasLDU]>;
}

multiclass LDU_G_NOINTRIN<string TyStr, NVPTXRegClass regclass, PatFrag IntOp> {
  def areg: NVPTXInst<(outs regclass:$result), (ins Int32Regs:$src),
               !strconcat("ldu.global.", TyStr),
         [(set regclass:$result, (IntOp Int32Regs:$src))]>, Requires<[hasLDU]>;
  def areg64: NVPTXInst<(outs regclass:$result), (ins Int64Regs:$src),
               !strconcat("ldu.global.", TyStr),
         [(set regclass:$result, (IntOp Int64Regs:$src))]>, Requires<[hasLDU]>;
 def avar:  NVPTXInst<(outs regclass:$result), (ins imem:$src),
               !strconcat("ldu.global.", TyStr),
         [(set regclass:$result, (IntOp (Wrapper tglobaladdr:$src)))]>,
         Requires<[hasLDU]>;
 def ari :  NVPTXInst<(outs regclass:$result), (ins MEMri:$src),
               !strconcat("ldu.global.", TyStr),
         [(set regclass:$result, (IntOp ADDRri:$src))]>, Requires<[hasLDU]>;
 def ari64 :  NVPTXInst<(outs regclass:$result), (ins MEMri64:$src),
               !strconcat("ldu.global.", TyStr),
         [(set regclass:$result, (IntOp ADDRri64:$src))]>, Requires<[hasLDU]>;
}

defm INT_PTX_LDU_GLOBAL_i8  : LDU_G_NOINTRIN<"u8 \t$result, [$src];", Int16Regs,
                                             ldu_i8>;
defm INT_PTX_LDU_GLOBAL_i16 : LDU_G<"u16 \t$result, [$src];", Int16Regs,
int_nvvm_ldu_global_i>;
defm INT_PTX_LDU_GLOBAL_i32 : LDU_G<"u32 \t$result, [$src];", Int32Regs,
int_nvvm_ldu_global_i>;
defm INT_PTX_LDU_GLOBAL_i64 : LDU_G<"u64 \t$result, [$src];", Int64Regs,
int_nvvm_ldu_global_i>;
defm INT_PTX_LDU_GLOBAL_f32 : LDU_G<"f32 \t$result, [$src];", Float32Regs,
int_nvvm_ldu_global_f>;
defm INT_PTX_LDU_GLOBAL_f64 : LDU_G<"f64 \t$result, [$src];", Float64Regs,
int_nvvm_ldu_global_f>;
defm INT_PTX_LDU_GLOBAL_p32 : LDU_G<"u32 \t$result, [$src];", Int32Regs,
int_nvvm_ldu_global_p>;
defm INT_PTX_LDU_GLOBAL_p64 : LDU_G<"u64 \t$result, [$src];", Int64Regs,
int_nvvm_ldu_global_p>;

// vector

// Elementized vector ldu
multiclass VLDU_G_ELE_V2<string TyStr, NVPTXRegClass regclass> {
 def _areg32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
                     (ins Int32Regs:$src),
                     !strconcat("ldu.global.", TyStr), []>;
 def _areg64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
                     (ins Int64Regs:$src),
                     !strconcat("ldu.global.", TyStr), []>;
 def _ari32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
                     (ins MEMri:$src),
                     !strconcat("ldu.global.", TyStr), []>;
 def _ari64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
                     (ins MEMri64:$src),
                     !strconcat("ldu.global.", TyStr), []>;
 def _avar: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
                     (ins imemAny:$src),
                     !strconcat("ldu.global.", TyStr), []>;
}

multiclass VLDU_G_ELE_V4<string TyStr, NVPTXRegClass regclass> { 
 def _areg32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
                            regclass:$dst4), (ins Int32Regs:$src), 
               !strconcat("ldu.global.", TyStr), []>;
 def _areg64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
                            regclass:$dst4), (ins Int64Regs:$src), 
               !strconcat("ldu.global.", TyStr), []>;
 def _ari32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
                            regclass:$dst4), (ins MEMri:$src), 
               !strconcat("ldu.global.", TyStr), []>;
 def _ari64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
                            regclass:$dst4), (ins MEMri64:$src), 
               !strconcat("ldu.global.", TyStr), []>;
 def _avar: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
                            regclass:$dst4), (ins imemAny:$src), 
               !strconcat("ldu.global.", TyStr), []>;
}

defm INT_PTX_LDU_G_v2i8_ELE
  : VLDU_G_ELE_V2<"v2.u8 \t{{$dst1, $dst2}}, [$src];",  Int16Regs>;
defm INT_PTX_LDU_G_v2i16_ELE
  : VLDU_G_ELE_V2<"v2.u16 \t{{$dst1, $dst2}}, [$src];", Int16Regs>;
defm INT_PTX_LDU_G_v2i32_ELE
  : VLDU_G_ELE_V2<"v2.u32 \t{{$dst1, $dst2}}, [$src];", Int32Regs>;
defm INT_PTX_LDU_G_v2f32_ELE
  : VLDU_G_ELE_V2<"v2.f32 \t{{$dst1, $dst2}}, [$src];", Float32Regs>;
defm INT_PTX_LDU_G_v2i64_ELE
  : VLDU_G_ELE_V2<"v2.u64 \t{{$dst1, $dst2}}, [$src];", Int64Regs>;
defm INT_PTX_LDU_G_v2f64_ELE
  : VLDU_G_ELE_V2<"v2.f64 \t{{$dst1, $dst2}}, [$src];", Float64Regs>;
defm INT_PTX_LDU_G_v4i8_ELE
  : VLDU_G_ELE_V4<"v4.u8 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Int16Regs>;
defm INT_PTX_LDU_G_v4i16_ELE
  : VLDU_G_ELE_V4<"v4.u16 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];",
    Int16Regs>;
defm INT_PTX_LDU_G_v4i32_ELE
  : VLDU_G_ELE_V4<"v4.u32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];",
    Int32Regs>;
defm INT_PTX_LDU_G_v4f32_ELE
  : VLDU_G_ELE_V4<"v4.f32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];",
    Float32Regs>;


//-----------------------------------
// Support for ldg on sm_35 or later 
//-----------------------------------

def ldg_i8 : PatFrag<(ops node:$ptr), (int_nvvm_ldg_global_i node:$ptr), [{
  MemIntrinsicSDNode *M = cast<MemIntrinsicSDNode>(N);
  return M->getMemoryVT() == MVT::i8;
}]>;

multiclass LDG_G<string TyStr, NVPTXRegClass regclass, Intrinsic IntOp> {
  def areg: NVPTXInst<(outs regclass:$result), (ins Int32Regs:$src),
               !strconcat("ld.global.nc.", TyStr),
         [(set regclass:$result, (IntOp Int32Regs:$src))]>, Requires<[hasLDG]>;
  def areg64: NVPTXInst<(outs regclass:$result), (ins Int64Regs:$src),
               !strconcat("ld.global.nc.", TyStr),
         [(set regclass:$result, (IntOp Int64Regs:$src))]>, Requires<[hasLDG]>;
 def avar:  NVPTXInst<(outs regclass:$result), (ins imem:$src),
               !strconcat("ld.global.nc.", TyStr),
         [(set regclass:$result, (IntOp (Wrapper tglobaladdr:$src)))]>,
         Requires<[hasLDG]>;
 def ari :  NVPTXInst<(outs regclass:$result), (ins MEMri:$src),
               !strconcat("ld.global.nc.", TyStr),
         [(set regclass:$result, (IntOp ADDRri:$src))]>, Requires<[hasLDG]>;
 def ari64 :  NVPTXInst<(outs regclass:$result), (ins MEMri64:$src),
               !strconcat("ld.global.nc.", TyStr),
         [(set regclass:$result, (IntOp ADDRri64:$src))]>, Requires<[hasLDG]>;
}

multiclass LDG_G_NOINTRIN<string TyStr, NVPTXRegClass regclass, PatFrag IntOp> {
  def areg: NVPTXInst<(outs regclass:$result), (ins Int32Regs:$src),
               !strconcat("ld.global.nc.", TyStr),
         [(set regclass:$result, (IntOp Int32Regs:$src))]>, Requires<[hasLDG]>;
  def areg64: NVPTXInst<(outs regclass:$result), (ins Int64Regs:$src),
               !strconcat("ld.global.nc.", TyStr),
         [(set regclass:$result, (IntOp Int64Regs:$src))]>, Requires<[hasLDG]>;
 def avar:  NVPTXInst<(outs regclass:$result), (ins imem:$src),
               !strconcat("ld.global.nc.", TyStr),
         [(set regclass:$result, (IntOp (Wrapper tglobaladdr:$src)))]>,
        Requires<[hasLDG]>;
 def ari :  NVPTXInst<(outs regclass:$result), (ins MEMri:$src),
               !strconcat("ld.global.nc.", TyStr),
         [(set regclass:$result, (IntOp ADDRri:$src))]>, Requires<[hasLDG]>;
 def ari64 :  NVPTXInst<(outs regclass:$result), (ins MEMri64:$src),
               !strconcat("ld.global.nc.", TyStr),
         [(set regclass:$result, (IntOp ADDRri64:$src))]>, Requires<[hasLDG]>;
}

defm INT_PTX_LDG_GLOBAL_i8
  : LDG_G_NOINTRIN<"u8 \t$result, [$src];",  Int16Regs, ldg_i8>;
defm INT_PTX_LDG_GLOBAL_i16
  : LDG_G<"u16 \t$result, [$src];", Int16Regs,   int_nvvm_ldg_global_i>;
defm INT_PTX_LDG_GLOBAL_i32
  : LDG_G<"u32 \t$result, [$src];", Int32Regs,   int_nvvm_ldg_global_i>;
defm INT_PTX_LDG_GLOBAL_i64
  : LDG_G<"u64 \t$result, [$src];", Int64Regs,   int_nvvm_ldg_global_i>;
defm INT_PTX_LDG_GLOBAL_f32
  : LDG_G<"f32 \t$result, [$src];", Float32Regs, int_nvvm_ldg_global_f>;
defm INT_PTX_LDG_GLOBAL_f64
  : LDG_G<"f64 \t$result, [$src];", Float64Regs, int_nvvm_ldg_global_f>;
defm INT_PTX_LDG_GLOBAL_p32
  : LDG_G<"u32 \t$result, [$src];", Int32Regs,   int_nvvm_ldg_global_p>;
defm INT_PTX_LDG_GLOBAL_p64
  : LDG_G<"u64 \t$result, [$src];", Int64Regs,   int_nvvm_ldg_global_p>;

// vector

// Elementized vector ldg 
multiclass VLDG_G_ELE_V2<string TyStr, NVPTXRegClass regclass> {
 def _areg32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
                     (ins Int32Regs:$src),
                     !strconcat("ld.global.nc.", TyStr), []>;
 def _areg64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
                     (ins Int64Regs:$src),
                     !strconcat("ld.global.nc.", TyStr), []>;
 def _ari32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
                     (ins MEMri:$src),
                     !strconcat("ld.global.nc.", TyStr), []>;
 def _ari64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
                     (ins MEMri64:$src),
                     !strconcat("ld.global.nc.", TyStr), []>;
 def _avar: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
                     (ins imemAny:$src),
                     !strconcat("ld.global.nc.", TyStr), []>;
}

multiclass VLDG_G_ELE_V4<string TyStr, NVPTXRegClass regclass> { 
  def _areg32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
                              regclass:$dst4), (ins Int32Regs:$src), 
               !strconcat("ld.global.nc.", TyStr), []>;
  def _areg64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
                               regclass:$dst4), (ins Int64Regs:$src), 
               !strconcat("ld.global.nc.", TyStr), []>;
  def _ari32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
                              regclass:$dst4), (ins MEMri:$src), 
               !strconcat("ld.global.nc.", TyStr), []>;
  def _ari64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
                              regclass:$dst4), (ins MEMri64:$src), 
               !strconcat("ld.global.nc.", TyStr), []>;
  def _avar: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
                             regclass:$dst4), (ins imemAny:$src), 
               !strconcat("ld.global.nc.", TyStr), []>;
}

// FIXME: 8-bit LDG should be fixed once LDG/LDU nodes are made into proper loads.
defm INT_PTX_LDG_G_v2i8_ELE
  : VLDG_G_ELE_V2<"v2.u8 \t{{$dst1, $dst2}}, [$src];",  Int16Regs>;
defm INT_PTX_LDG_G_v2i16_ELE
  : VLDG_G_ELE_V2<"v2.u16 \t{{$dst1, $dst2}}, [$src];", Int16Regs>;
defm INT_PTX_LDG_G_v2i32_ELE
  : VLDG_G_ELE_V2<"v2.u32 \t{{$dst1, $dst2}}, [$src];", Int32Regs>;
defm INT_PTX_LDG_G_v2f32_ELE
  : VLDG_G_ELE_V2<"v2.f32 \t{{$dst1, $dst2}}, [$src];", Float32Regs>;
defm INT_PTX_LDG_G_v2i64_ELE
  : VLDG_G_ELE_V2<"v2.u64 \t{{$dst1, $dst2}}, [$src];", Int64Regs>;
defm INT_PTX_LDG_G_v2f64_ELE
  : VLDG_G_ELE_V2<"v2.f64 \t{{$dst1, $dst2}}, [$src];", Float64Regs>;
defm INT_PTX_LDG_G_v4i8_ELE
  : VLDG_G_ELE_V4<"v4.u8 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Int16Regs>;
defm INT_PTX_LDG_G_v4i16_ELE
  : VLDG_G_ELE_V4<"v4.u16 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Int16Regs>;
defm INT_PTX_LDG_G_v4i32_ELE
  : VLDG_G_ELE_V4<"v4.u32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Int32Regs>;
defm INT_PTX_LDG_G_v4f32_ELE
  : VLDG_G_ELE_V4<"v4.f32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Float32Regs>;


multiclass NG_TO_G<string Str, Intrinsic Intrin> {
   def _yes : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
          !strconcat("cvta.", !strconcat(Str, ".u32 \t$result, $src;")),
      [(set Int32Regs:$result, (Intrin Int32Regs:$src))]>,
   Requires<[hasGenericLdSt]>;
   def _yes_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
          !strconcat("cvta.", !strconcat(Str, ".u64 \t$result, $src;")),
      [(set Int64Regs:$result, (Intrin Int64Regs:$src))]>,
   Requires<[hasGenericLdSt]>;

// @TODO: Are these actually needed?  I believe global addresses will be copied
// to register values anyway.
   /*def __addr_yes : NVPTXInst<(outs Int32Regs:$result), (ins imemAny:$src),
          !strconcat("cvta.", !strconcat(Str, ".u32 \t$result, $src;")),
      [(set Int32Regs:$result, (Intrin (Wrapper tglobaladdr:$src)))]>,
      Requires<[hasGenericLdSt]>;
   def __addr_yes_64 : NVPTXInst<(outs Int64Regs:$result), (ins imemAny:$src),
          !strconcat("cvta.", !strconcat(Str, ".u64 \t$result, $src;")),
      [(set Int64Regs:$result, (Intrin (Wrapper tglobaladdr:$src)))]>,
      Requires<[hasGenericLdSt]>;*/

   def _no : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
          "mov.u32 \t$result, $src;",
      [(set Int32Regs:$result, (Intrin Int32Regs:$src))]>;
   def _no_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
          "mov.u64 \t$result, $src;",
      [(set Int64Regs:$result, (Intrin Int64Regs:$src))]>;

// @TODO: Are these actually needed?  I believe global addresses will be copied
// to register values anyway.
   /*def _addr_no : NVPTXInst<(outs Int32Regs:$result), (ins imem:$src),
          "mov.u32 \t$result, $src;",
      [(set Int32Regs:$result, (Intrin (Wrapper tglobaladdr:$src)))]>;
   def _addr_no_64 : NVPTXInst<(outs Int64Regs:$result), (ins imem:$src),
          "mov.u64 \t$result, $src;",
      [(set Int64Regs:$result, (Intrin (Wrapper tglobaladdr:$src)))]>;*/
}

multiclass G_TO_NG<string Str, Intrinsic Intrin> {
   def _yes : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
          !strconcat("cvta.to.", !strconcat(Str, ".u32 \t$result, $src;")),
      [(set Int32Regs:$result, (Intrin Int32Regs:$src))]>,
   Requires<[hasGenericLdSt]>;
   def _yes_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
          !strconcat("cvta.to.", !strconcat(Str, ".u64 \t$result, $src;")),
      [(set Int64Regs:$result, (Intrin Int64Regs:$src))]>,
   Requires<[hasGenericLdSt]>;
   def _no : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
          "mov.u32 \t$result, $src;",
      [(set Int32Regs:$result, (Intrin Int32Regs:$src))]>;
   def _no_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
          "mov.u64 \t$result, $src;",
      [(set Int64Regs:$result, (Intrin Int64Regs:$src))]>;
}

defm cvta_local  : NG_TO_G<"local", int_nvvm_ptr_local_to_gen>;
defm cvta_shared : NG_TO_G<"shared", int_nvvm_ptr_shared_to_gen>;
defm cvta_global : NG_TO_G<"global", int_nvvm_ptr_global_to_gen>;
defm cvta_const  : NG_TO_G<"const", int_nvvm_ptr_constant_to_gen>;

defm cvta_to_local   : G_TO_NG<"local", int_nvvm_ptr_gen_to_local>;
defm cvta_to_shared : G_TO_NG<"shared", int_nvvm_ptr_gen_to_shared>;
defm cvta_to_global : G_TO_NG<"global", int_nvvm_ptr_gen_to_global>;
defm cvta_to_const  : G_TO_NG<"const", int_nvvm_ptr_gen_to_constant>;


// nvvm.ptr.gen.to.param
def nvvm_ptr_gen_to_param : NVPTXInst<(outs Int32Regs:$result),
  (ins Int32Regs:$src),
                        "mov.u32 \t$result, $src;",
                              [(set Int32Regs:$result,
                                (int_nvvm_ptr_gen_to_param Int32Regs:$src))]>;
def nvvm_ptr_gen_to_param_64 : NVPTXInst<(outs Int64Regs:$result),
  (ins Int64Regs:$src),
                        "mov.u64 \t$result, $src;",
                              [(set Int64Regs:$result,
                                (int_nvvm_ptr_gen_to_param Int64Regs:$src))]>;


// nvvm.move intrinsicc
def nvvm_move_i16 : NVPTXInst<(outs Int16Regs:$r), (ins Int16Regs:$s),
                             "mov.b16 \t$r, $s;",
                             [(set Int16Regs:$r,
                               (int_nvvm_move_i16 Int16Regs:$s))]>;
def nvvm_move_i32 : NVPTXInst<(outs Int32Regs:$r), (ins Int32Regs:$s),
                             "mov.b32 \t$r, $s;",
                             [(set Int32Regs:$r,
                               (int_nvvm_move_i32 Int32Regs:$s))]>;
def nvvm_move_i64 : NVPTXInst<(outs Int64Regs:$r), (ins Int64Regs:$s),
                             "mov.b64 \t$r, $s;",
                             [(set Int64Regs:$r,
                               (int_nvvm_move_i64 Int64Regs:$s))]>;
def nvvm_move_float : NVPTXInst<(outs Float32Regs:$r), (ins Float32Regs:$s),
                             "mov.f32 \t$r, $s;",
                             [(set Float32Regs:$r,
                               (int_nvvm_move_float Float32Regs:$s))]>;
def nvvm_move_double : NVPTXInst<(outs Float64Regs:$r), (ins Float64Regs:$s),
                             "mov.f64 \t$r, $s;",
                             [(set Float64Regs:$r,
                               (int_nvvm_move_double Float64Regs:$s))]>;
def nvvm_move_ptr32 : NVPTXInst<(outs Int32Regs:$r), (ins Int32Regs:$s),
                             "mov.u32 \t$r, $s;",
                             [(set Int32Regs:$r,
                               (int_nvvm_move_ptr Int32Regs:$s))]>;
def nvvm_move_ptr64 : NVPTXInst<(outs Int64Regs:$r), (ins Int64Regs:$s),
                             "mov.u64 \t$r, $s;",
                             [(set Int64Regs:$r,
                               (int_nvvm_move_ptr Int64Regs:$s))]>;

// @TODO: Are these actually needed, or will we always just see symbols
// copied to registers first?
/*def nvvm_move_sym32 : NVPTXInst<(outs Int32Regs:$r), (ins imem:$s),
                             "mov.u32 \t$r, $s;",
                             [(set Int32Regs:$r,
                             (int_nvvm_move_ptr texternalsym:$s))]>;
def nvvm_move_sym64 : NVPTXInst<(outs Int64Regs:$r), (ins imem:$s),
                             "mov.u64 \t$r, $s;",
                             [(set Int64Regs:$r,
                             (int_nvvm_move_ptr texternalsym:$s))]>;*/


// MoveParam        %r1, param
// ptr_local_to_gen %r2, %r1
// ptr_gen_to_local %r3, %r2
// ->
// mov %r1, param

// @TODO: Revisit this.  There is a type
// contradiction between iPTRAny and iPTR for the addr defs, so the move_sym
// instructions are not currently defined. However, we can use the ptr
// variants and the asm printer will do the right thing.
def : Pat<(i64 (int_nvvm_ptr_gen_to_local (int_nvvm_ptr_local_to_gen
                (MoveParam texternalsym:$src)))),
               (nvvm_move_ptr64  texternalsym:$src)>;
def : Pat<(i32 (int_nvvm_ptr_gen_to_local (int_nvvm_ptr_local_to_gen
                (MoveParam texternalsym:$src)))),
               (nvvm_move_ptr32  texternalsym:$src)>;

def texsurf_handles
  : NVPTXInst<(outs Int64Regs:$result), (ins imem:$src),
              "mov.u64 \t$result, $src;", []>;

//-----------------------------------
// Compiler Error Warn
// - Just ignore them in codegen
//-----------------------------------

def INT_NVVM_COMPILER_WARN_32 : NVPTXInst<(outs), (ins Int32Regs:$a),
                "// llvm.nvvm.compiler.warn()",
                [(int_nvvm_compiler_warn Int32Regs:$a)]>;
def INT_NVVM_COMPILER_WARN_64 : NVPTXInst<(outs), (ins Int64Regs:$a),
                "// llvm.nvvm.compiler.warn()",
                [(int_nvvm_compiler_warn Int64Regs:$a)]>;
def INT_NVVM_COMPILER_ERROR_32 : NVPTXInst<(outs), (ins Int32Regs:$a),
                "// llvm.nvvm.compiler.error()",
                [(int_nvvm_compiler_error Int32Regs:$a)]>;
def INT_NVVM_COMPILER_ERROR_64 : NVPTXInst<(outs), (ins Int64Regs:$a),
                "// llvm.nvvm.compiler.error()",
                [(int_nvvm_compiler_error Int64Regs:$a)]>;


//-----------------------------------
// Texture Intrinsics
//-----------------------------------

// NOTE: For Fermi support, any new texture/surface/sampler intrinsics must be
// also defined in NVPTXReplaceImageHandles.cpp


// Texture fetch instructions using handles
def TEX_1D_F32_I32
  : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
                    Float32Regs:$b, Float32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$x),
              "tex.1d.v4.f32.s32\t\\{$r, $g, $b, $a\\}, [$t, $s, \\{$x\\}];",
              []>;
def TEX_1D_F32_F32
  : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
                    Float32Regs:$b, Float32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x),
              "tex.1d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, [$t, $s, \\{$x\\}];",
              []>;
def TEX_1D_F32_F32_LEVEL
  : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
                    Float32Regs:$b, Float32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$lod),
              "tex.level.1d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
              "[$t, $s, \\{$x\\}], $lod;",
              []>;
def TEX_1D_F32_F32_GRAD
  : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
                    Float32Regs:$b, Float32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x,
                   Float32Regs:$gradx, Float32Regs:$grady),
              "tex.grad.1d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
              "[$t, $s, \\{$x\\}], \\{$gradx\\}, \\{$grady\\};",
              []>;
def TEX_1D_I32_I32
  : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
                    Int32Regs:$b, Int32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$x),
              "tex.1d.v4.s32.s32\t\\{$r, $g, $b, $a\\}, [$t, $s, \\{$x\\}];",
              []>;
def TEX_1D_I32_F32
  : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
                    Int32Regs:$b, Int32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x),
              "tex.1d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, [$t, $s, \\{$x\\}];",
              []>;
def TEX_1D_I32_F32_LEVEL
  : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
                    Int32Regs:$b, Int32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x,
                   Float32Regs:$lod),
              "tex.level.1d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
              "[$t, $s, \\{$x\\}], $lod;",
              []>;
def TEX_1D_I32_F32_GRAD
  : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
                    Int32Regs:$b, Int32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x,
                   Float32Regs:$gradx, Float32Regs:$grady),
              "tex.grad.1d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
              "[$t, $s, \\{$x\\}], \\{$gradx\\}, \\{$grady\\};",
              []>;

def TEX_1D_ARRAY_F32_I32
  : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
                    Float32Regs:$b, Float32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Int32Regs:$x),
              "tex.a1d.v4.f32.s32\t\\{$r, $g, $b, $a\\}, "
              "[$t, $s, \\{$l, $x\\}];",
              []>;
def TEX_1D_ARRAY_F32_F32
  : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
                    Float32Regs:$b, Float32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x),
              "tex.a1d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
              "[$t, $s, \\{$l, $x\\}];",
              []>;
def TEX_1D_ARRAY_F32_F32_LEVEL
  : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
                    Float32Regs:$b, Float32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
                   Float32Regs:$lod),
              "tex.level.a1d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
              "[$t, $s, \\{$l, $x\\}], $lod;",
              []>;
def TEX_1D_ARRAY_F32_F32_GRAD
  : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
                    Float32Regs:$b, Float32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
                   Float32Regs:$gradx, Float32Regs:$grady),
              "tex.grad.a1d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
              "[$t, $s, \\{$l, $x\\}], \\{$gradx\\}, \\{$grady\\};",
              []>;
def TEX_1D_ARRAY_I32_I32
  : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
                    Int32Regs:$b, Int32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Int32Regs:$x),
              "tex.a1d.v4.s32.s32\t\\{$r, $g, $b, $a\\}, "
              "[$t, $s, \\{$l, $x\\}];",
              []>;
def TEX_1D_ARRAY_I32_F32
  : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
                    Int32Regs:$b, Int32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x),
              "tex.a1d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
              "[$t, $s, \\{$l, $x\\}];",
              []>;
def TEX_1D_ARRAY_I32_F32_LEVEL
  : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
                    Int32Regs:$b, Int32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
                   Float32Regs:$lod),
              "tex.level.a1d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
              "[$t, $s, \\{$l, $x\\}], $lod;",
              []>;
def TEX_1D_ARRAY_I32_F32_GRAD
  : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
                    Int32Regs:$b, Int32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
                   Float32Regs:$gradx, Float32Regs:$grady),
              "tex.grad.a1d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
              "[$t, $s, \\{$l, $x\\}], \\{$gradx\\}, \\{$grady\\};",
              []>;

def TEX_2D_F32_I32
  : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
                    Float32Regs:$b, Float32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$x, Int32Regs:$y),
              "tex.2d.v4.f32.s32\t\\{$r, $g, $b, $a\\}, "
              "[$t, $s, \\{$x, $y\\}];",
              []>;
def TEX_2D_F32_F32
  : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
                    Float32Regs:$b, Float32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y),
              "tex.2d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
              "[$t, $s, \\{$x, $y\\}];",
              []>;
def TEX_2D_F32_F32_LEVEL
  : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
                    Float32Regs:$b, Float32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
                   Float32Regs:$lod),
              "tex.level.2d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
              "[$t, $s, \\{$x, $y\\}], $lod;",
              []>;
def TEX_2D_F32_F32_GRAD
  : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
                    Float32Regs:$b, Float32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
                   Float32Regs:$gradx0, Float32Regs:$gradx1,
                   Float32Regs:$grady0, Float32Regs:$grady1),
              "tex.grad.2d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
              "[$t, $s, \\{$x, $y\\}], \\{$gradx0, $gradx1\\}, "
              "\\{$grady0, $grady1\\};",
              []>;
def TEX_2D_I32_I32
  : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
                    Int32Regs:$b, Int32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$x, Int32Regs:$y),
              "tex.2d.v4.s32.s32\t\\{$r, $g, $b, $a\\}, "
              "[$t, $s, \\{$x, $y\\}];",
              []>;
def TEX_2D_I32_F32
  : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
                    Int32Regs:$b, Int32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y),
              "tex.2d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
              "[$t, $s, \\{$x, $y\\}];",
              []>;
def TEX_2D_I32_F32_LEVEL
  : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
                    Int32Regs:$b, Int32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
                   Float32Regs:$lod),
              "tex.level.2d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
              "[$t, $s, \\{$x, $y\\}], $lod;",
              []>;
def TEX_2D_I32_F32_GRAD
  : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
                    Int32Regs:$b, Int32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
                   Float32Regs:$gradx0, Float32Regs:$gradx1,
                   Float32Regs:$grady0, Float32Regs:$grady1),
              "tex.grad.2d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
              "[$t, $s, \\{$x, $y\\}], \\{$gradx0, $gradx1\\}, "
              "\\{$grady0, $grady1\\};",
              []>;

def TEX_2D_ARRAY_F32_I32
  : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
                    Float32Regs:$b, Float32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
                   Int32Regs:$y),
              "tex.a2d.v4.f32.s32\t\\{$r, $g, $b, $a\\}, "
              "[$t, $s, \\{$l, $x, $y, $y\\}];",
              []>;
def TEX_2D_ARRAY_F32_F32
  : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
                    Float32Regs:$b, Float32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
                   Float32Regs:$y),
              "tex.a2d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
              "[$t, $s, \\{$l, $x, $y, $y\\}];",
              []>;
def TEX_2D_ARRAY_F32_F32_LEVEL
  : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
                    Float32Regs:$b, Float32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
                   Float32Regs:$y, Float32Regs:$lod),
              "tex.level.a2d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
              "[$t, $s, \\{$l, $x, $y, $y\\}], $lod;",
              []>;
def TEX_2D_ARRAY_F32_F32_GRAD
  : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
                    Float32Regs:$b, Float32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
                   Float32Regs:$y, Float32Regs:$gradx0, Float32Regs:$gradx1,
                   Float32Regs:$grady0, Float32Regs:$grady1),
              "tex.grad.a2d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
              "[$t, $s, \\{$l, $x, $y, $y\\}], \\{$gradx0, $gradx1\\}, "
              "\\{$grady0, $grady1\\};",
              []>;
def TEX_2D_ARRAY_I32_I32
  : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
                    Int32Regs:$b, Int32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
                   Int32Regs:$y),
              "tex.a2d.v4.s32.s32\t\\{$r, $g, $b, $a\\}, "
              "[$t, $s, \\{$l, $x, $y, $y\\}];",
              []>;
def TEX_2D_ARRAY_I32_F32
  : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
                    Int32Regs:$b, Int32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
                   Float32Regs:$y),
              "tex.a2d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
              "[$t, $s, \\{$l, $x, $y, $y\\}];",
              []>;
def TEX_2D_ARRAY_I32_F32_LEVEL
  : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
                    Int32Regs:$b, Int32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
                   Float32Regs:$y, Float32Regs:$lod),
              "tex.level.a2d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
              "[$t, $s, \\{$l, $x, $y, $y\\}], $lod;",
              []>;
def TEX_2D_ARRAY_I32_F32_GRAD
  : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
                    Int32Regs:$b, Int32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
                   Float32Regs:$y,
                   Float32Regs:$gradx0, Float32Regs:$gradx1,
                   Float32Regs:$grady0, Float32Regs:$grady1),
              "tex.grad.a2d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
              "[$t, $s, \\{$l, $x, $y, $y\\}], \\{$gradx0, $gradx1\\}, "
              "\\{$grady0, $grady1\\};",
              []>;

def TEX_3D_F32_I32
  : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
                    Float32Regs:$b, Float32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
                   Int32Regs:$z),
              "tex.3d.v4.f32.s32\t\\{$r, $g, $b, $a\\}, "
              "[$t, $s, \\{$x, $y, $z, $z\\}];",
              []>;
def TEX_3D_F32_F32
  : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
                    Float32Regs:$b, Float32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
                   Float32Regs:$z),
              "tex.3d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
              "[$t, $s, \\{$x, $y, $z, $z\\}];",
              []>;
def TEX_3D_F32_F32_LEVEL
  : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
                    Float32Regs:$b, Float32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
                   Float32Regs:$z, Float32Regs:$lod),
              "tex.level.3d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
              "[$t, $s, \\{$x, $y, $z, $z\\}], $lod;",
              []>;
def TEX_3D_F32_F32_GRAD
  : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
                    Float32Regs:$b, Float32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
                   Float32Regs:$z,
                   Float32Regs:$gradx0, Float32Regs:$gradx1,
                   Float32Regs:$gradx2, Float32Regs:$grady0,
                   Float32Regs:$grady1, Float32Regs:$grady2),
              "tex.grad.3d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
              "[$t, $s, \\{$x, $y, $z, $z\\}], "
              "\\{$gradx0, $gradx1, $gradx2, $gradx2\\}, "
              "\\{$grady0, $grady1, $grady2, $grady2\\};",
              []>;
def TEX_3D_I32_I32
  : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
                    Int32Regs:$b, Int32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
                   Int32Regs:$z),
              "tex.3d.v4.s32.s32\t\\{$r, $g, $b, $a\\}, "
              "[$t, $s, \\{$x, $y, $z, $z\\}];",
              []>;
def TEX_3D_I32_F32
  : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
                    Int32Regs:$b, Int32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
                   Float32Regs:$z),
              "tex.3d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
              "[$t, $s, \\{$x, $y, $z, $z\\}];",
              []>;
def TEX_3D_I32_F32_LEVEL
  : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
                    Int32Regs:$b, Int32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
                   Float32Regs:$z, Float32Regs:$lod),
              "tex.level.3d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
              "[$t, $s, \\{$x, $y, $z, $z\\}], $lod;",
              []>;
def TEX_3D_I32_F32_GRAD
  : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
                    Int32Regs:$b, Int32Regs:$a),
              (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
                   Float32Regs:$z,
                   Float32Regs:$gradx0, Float32Regs:$gradx1,
                   Float32Regs:$gradx2, Float32Regs:$grady0,
                   Float32Regs:$grady1, Float32Regs:$grady2),
              "tex.grad.3d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
              "[$t, $s, \\{$x, $y, $z, $z\\}], "
              "\\{$gradx0, $gradx1, $gradx2, $gradx2\\}, "
              "\\{$grady0, $grady1, $grady2, $grady2\\};",
              []>;


// Surface load instructions
def SULD_1D_I8_TRAP
  : NVPTXInst<(outs Int16Regs:$r),
              (ins Int64Regs:$s, Int32Regs:$x),
              "suld.b.1d.b8.trap \\{$r\\}, [$s, \\{$x\\}];",
              []>;
def SULD_1D_I16_TRAP
  : NVPTXInst<(outs Int16Regs:$r),
              (ins Int64Regs:$s, Int32Regs:$x),
              "suld.b.1d.b16.trap \\{$r\\}, [$s, \\{$x\\}];",
              []>;
def SULD_1D_I32_TRAP
  : NVPTXInst<(outs Int32Regs:$r),
              (ins Int64Regs:$s, Int32Regs:$x),
              "suld.b.1d.b32.trap \\{$r\\}, [$s, \\{$x\\}];",
              []>;
def SULD_1D_V2I8_TRAP
  : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g),
              (ins Int64Regs:$s, Int32Regs:$x),
              "suld.b.1d.v2.b8.trap \\{$r, $g\\}, [$s, \\{$x\\}];",
              []>;
def SULD_1D_V2I16_TRAP
  : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g),
              (ins Int64Regs:$s, Int32Regs:$x),
              "suld.b.1d.v2.b16.trap \\{$r, $g\\}, [$s, \\{$x\\}];",
              []>;
def SULD_1D_V2I32_TRAP
  : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g),
              (ins Int64Regs:$s, Int32Regs:$x),
              "suld.b.1d.v2.b32.trap \\{$r, $g\\}, [$s, \\{$x\\}];",
              []>;
def SULD_1D_V4I8_TRAP
  : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
              (ins Int64Regs:$s, Int32Regs:$x),
              "suld.b.1d.v4.b8.trap \\{$r, $g, $b, $a\\}, [$s, \\{$x\\}];",
              []>;
def SULD_1D_V4I16_TRAP
  : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
              (ins Int64Regs:$s, Int32Regs:$x),
              "suld.b.1d.v4.b16.trap \\{$r, $g, $b, $a\\}, [$s, \\{$x\\}];",
              []>;
def SULD_1D_V4I32_TRAP
  : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
              (ins Int64Regs:$s, Int32Regs:$x),
              "suld.b.1d.v4.b32.trap \\{$r, $g, $b, $a\\}, [$s, \\{$x\\}];",
              []>;

def SULD_1D_ARRAY_I8_TRAP
  : NVPTXInst<(outs Int16Regs:$r),
              (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x),
              "suld.b.a1d.b8.trap \\{$r\\}, [$s, \\{$l, $x\\}];",
              []>;
def SULD_1D_ARRAY_I16_TRAP
  : NVPTXInst<(outs Int16Regs:$r),
              (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x),
              "suld.b.a1d.b16.trap \\{$r\\}, [$s, \\{$l, $x\\}];",
              []>;
def SULD_1D_ARRAY_I32_TRAP
  : NVPTXInst<(outs Int32Regs:$r),
              (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x),
              "suld.b.a1d.b32.trap \\{$r\\}, [$s, \\{$l, $x\\}];",
              []>;
def SULD_1D_ARRAY_V2I8_TRAP
  : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g),
              (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x),
              "suld.b.a1d.v2.b8.trap \\{$r, $g\\}, [$s, \\{$l, $x\\}];",
              []>;
def SULD_1D_ARRAY_V2I16_TRAP
  : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g),
              (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x),
              "suld.b.a1d.v2.b16.trap \\{$r, $g\\}, [$s, \\{$l, $x\\}];",
              []>;
def SULD_1D_ARRAY_V2I32_TRAP
  : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g),
              (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x),
              "suld.b.a1d.v2.b32.trap \\{$r, $g\\}, [$s, \\{$l, $x\\}];",
              []>;
def SULD_1D_ARRAY_V4I8_TRAP
  : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
              (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x),
              "suld.b.a1d.v4.b8.trap \\{$r, $g, $b, $a\\}, "
              "[$s, \\{$l, $x\\}];",
              []>;
def SULD_1D_ARRAY_V4I16_TRAP
  : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
              (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x),
              "suld.b.a1d.v4.b16.trap \\{$r, $g, $b, $a\\}, "
              "[$s, \\{$l, $x\\}];",
              []>;
def SULD_1D_ARRAY_V4I32_TRAP
  : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
              (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x),
              "suld.b.a1d.v4.b32.trap \\{$r, $g, $b, $a\\}, "
              "[$s, \\{$l, $x\\}];",
              []>;

def SULD_2D_I8_TRAP
  : NVPTXInst<(outs Int16Regs:$r),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y),
              "suld.b.2d.b8.trap \\{$r\\}, [$s, \\{$x, $y\\}];",
              []>;
def SULD_2D_I16_TRAP
  : NVPTXInst<(outs Int16Regs:$r),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y),
              "suld.b.2d.b16.trap \\{$r\\}, [$s, \\{$x, $y\\}];",
              []>;
def SULD_2D_I32_TRAP
  : NVPTXInst<(outs Int32Regs:$r),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y),
              "suld.b.2d.b32.trap \\{$r\\}, [$s, \\{$x, $y\\}];",
              []>;
def SULD_2D_V2I8_TRAP
  : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y),
              "suld.b.2d.v2.b8.trap \\{$r, $g\\}, [$s, \\{$x, $y\\}];",
              []>;
def SULD_2D_V2I16_TRAP
  : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y),
              "suld.b.2d.v2.b16.trap \\{$r, $g\\}, [$s, \\{$x, $y\\}];",
              []>;
def SULD_2D_V2I32_TRAP
  : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y),
              "suld.b.2d.v2.b32.trap \\{$r, $g\\}, [$s, \\{$x, $y\\}];",
              []>;
def SULD_2D_V4I8_TRAP
  : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y),
              "suld.b.2d.v4.b8.trap \\{$r, $g, $b, $a\\}, [$s, \\{$x, $y\\}];",
              []>;
def SULD_2D_V4I16_TRAP
  : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y),
              "suld.b.2d.v4.b16.trap \\{$r, $g, $b, $a\\}, [$s, \\{$x, $y\\}];",
              []>;
def SULD_2D_V4I32_TRAP
  : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y),
              "suld.b.2d.v4.b32.trap \\{$r, $g, $b, $a\\}, [$s, \\{$x, $y\\}];",
              []>;

def SULD_2D_ARRAY_I8_TRAP
  : NVPTXInst<(outs Int16Regs:$r),
              (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y),
              "suld.b.a2d.b8.trap \\{$r\\}, [$s, \\{$l, $x, $y, $y\\}];",
              []>;
def SULD_2D_ARRAY_I16_TRAP
  : NVPTXInst<(outs Int16Regs:$r),
              (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y),
              "suld.b.a2d.b16.trap \\{$r\\}, [$s, \\{$l, $x, $y, $y\\}];",
              []>;
def SULD_2D_ARRAY_I32_TRAP
  : NVPTXInst<(outs Int32Regs:$r),
              (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y),
              "suld.b.a2d.b32.trap \\{$r\\}, [$s, \\{$l, $x, $y, $y\\}];",
              []>;
def SULD_2D_ARRAY_V2I8_TRAP
  : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g),
              (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y),
              "suld.b.a2d.v2.b8.trap \\{$r, $g\\}, "
              "[$s, \\{$l, $x, $y, $y\\}];",
              []>;
def SULD_2D_ARRAY_V2I16_TRAP
  : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g),
              (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y),
              "suld.b.a2d.v2.b16.trap \\{$r, $g\\}, "
              "[$s, \\{$l, $x, $y, $y\\}];",
              []>;
def SULD_2D_ARRAY_V2I32_TRAP
  : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g),
              (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y),
              "suld.b.a2d.v2.b32.trap \\{$r, $g\\}, "
              "[$s, \\{$l, $x, $y, $y\\}];",
              []>;
def SULD_2D_ARRAY_V4I8_TRAP
  : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
              (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y),
              "suld.b.a2d.v4.b8.trap \\{$r, $g, $b, $a\\}, "
              "[$s, \\{$l, $x, $y, $y\\}];",
              []>;
def SULD_2D_ARRAY_V4I16_TRAP
  : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
              (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y),
              "suld.b.a2d.v4.b16.trap \\{$r, $g, $b, $a\\}, "
              "[$s, \\{$l, $x, $y, $y\\}];",
              []>;
def SULD_2D_ARRAY_V4I32_TRAP
  : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
              (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y),
              "suld.b.a2d.v4.b32.trap \\{$r, $g, $b, $a\\}, "
              "[$s, \\{$l, $x, $y, $y\\}];",
              []>;

def SULD_3D_I8_TRAP
  : NVPTXInst<(outs Int16Regs:$r),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z),
              "suld.b.3d.b8.trap \\{$r\\}, [$s, \\{$x, $y, $z, $z\\}];",
              []>;
def SULD_3D_I16_TRAP
  : NVPTXInst<(outs Int16Regs:$r),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z),
              "suld.b.3d.b16.trap \\{$r\\}, [$s, \\{$x, $y, $z, $z\\}];",
              []>;
def SULD_3D_I32_TRAP
  : NVPTXInst<(outs Int32Regs:$r),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z),
              "suld.b.3d.b32.trap \\{$r\\}, [$s, \\{$x, $y, $z, $z\\}];",
              []>;
def SULD_3D_V2I8_TRAP
  : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z),
              "suld.b.3d.v2.b8.trap \\{$r, $g\\}, [$s, \\{$x, $y, $z, $z\\}];",
              []>;
def SULD_3D_V2I16_TRAP
  : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z),
              "suld.b.3d.v2.b16.trap \\{$r, $g\\}, [$s, \\{$x, $y, $z, $z\\}];",
              []>;
def SULD_3D_V2I32_TRAP
  : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z),
              "suld.b.3d.v2.b32.trap \\{$r, $g\\}, [$s, \\{$x, $y, $z, $z\\}];",
              []>;
def SULD_3D_V4I8_TRAP
  : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z),
              "suld.b.3d.v4.b8.trap \\{$r, $g, $b, $a\\}, "
              "[$s, \\{$x, $y, $z, $z\\}];",
              []>;
def SULD_3D_V4I16_TRAP
  : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z),
              "suld.b.3d.v4.b16.trap \\{$r, $g, $b, $a\\}, "
              "[$s, \\{$x, $y, $z, $z\\}];",
              []>;
def SULD_3D_V4I32_TRAP
  : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z),
              "suld.b.3d.v4.b32.trap \\{$r, $g, $b, $a\\}, "
              "[$s, \\{$x, $y, $z, $z\\}];",
              []>;


//-----------------------------------
// Texture Query Intrinsics
//-----------------------------------
def TXQ_CHANNEL_ORDER
  : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
              "txq.channel_order.b32 \t$d, [$a];",
              []>;
def TXQ_CHANNEL_DATA_TYPE
  : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
              "txq.channel_data_type.b32 \t$d, [$a];",
              []>;
def TXQ_WIDTH
  : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
              "txq.width.b32 \t$d, [$a];",
              []>;
def TXQ_HEIGHT
  : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
              "txq.height.b32 \t$d, [$a];",
              []>;
def TXQ_DEPTH
  : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
              "txq.depth.b32 \t$d, [$a];",
              []>;
def TXQ_ARRAY_SIZE
  : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
              "txq.array_size.b32 \t$d, [$a];",
              []>;
def TXQ_NUM_SAMPLES
  : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
              "txq.num_samples.b32 \t$d, [$a];",
              []>;
def TXQ_NUM_MIPMAP_LEVELS
  : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
              "txq.num_mipmap_levels.b32 \t$d, [$a];",
              []>;

def : Pat<(int_nvvm_txq_channel_order Int64Regs:$a),
          (TXQ_CHANNEL_ORDER Int64Regs:$a)>;
def : Pat<(int_nvvm_txq_channel_data_type Int64Regs:$a),
          (TXQ_CHANNEL_DATA_TYPE Int64Regs:$a)>;
def : Pat<(int_nvvm_txq_width Int64Regs:$a),
          (TXQ_WIDTH Int64Regs:$a)>;
def : Pat<(int_nvvm_txq_height Int64Regs:$a),
          (TXQ_HEIGHT Int64Regs:$a)>;
def : Pat<(int_nvvm_txq_depth Int64Regs:$a),
          (TXQ_DEPTH Int64Regs:$a)>;
def : Pat<(int_nvvm_txq_array_size Int64Regs:$a),
          (TXQ_ARRAY_SIZE Int64Regs:$a)>;
def : Pat<(int_nvvm_txq_num_samples Int64Regs:$a),
          (TXQ_NUM_SAMPLES Int64Regs:$a)>;
def : Pat<(int_nvvm_txq_num_mipmap_levels Int64Regs:$a),
          (TXQ_NUM_MIPMAP_LEVELS Int64Regs:$a)>;


//-----------------------------------
// Surface Query Intrinsics
//-----------------------------------
def SUQ_CHANNEL_ORDER
  : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
              "suq.channel_order.b32 \t$d, [$a];",
              []>;
def SUQ_CHANNEL_DATA_TYPE
  : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
              "suq.channel_data_type.b32 \t$d, [$a];",
              []>;
def SUQ_WIDTH
  : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
              "suq.width.b32 \t$d, [$a];",
              []>;
def SUQ_HEIGHT
  : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
              "suq.height.b32 \t$d, [$a];",
              []>;
def SUQ_DEPTH
  : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
              "suq.depth.b32 \t$d, [$a];",
              []>;
def SUQ_ARRAY_SIZE
  : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
              "suq.array_size.b32 \t$d, [$a];",
              []>;

def : Pat<(int_nvvm_suq_channel_order Int64Regs:$a),
          (SUQ_CHANNEL_ORDER Int64Regs:$a)>;
def : Pat<(int_nvvm_suq_channel_data_type Int64Regs:$a),
          (SUQ_CHANNEL_DATA_TYPE Int64Regs:$a)>;
def : Pat<(int_nvvm_suq_width Int64Regs:$a),
          (SUQ_WIDTH Int64Regs:$a)>;
def : Pat<(int_nvvm_suq_height Int64Regs:$a),
          (SUQ_HEIGHT Int64Regs:$a)>;
def : Pat<(int_nvvm_suq_depth Int64Regs:$a),
          (SUQ_DEPTH Int64Regs:$a)>;
def : Pat<(int_nvvm_suq_array_size Int64Regs:$a),
          (SUQ_ARRAY_SIZE Int64Regs:$a)>;


//===- Handle Query -------------------------------------------------------===//

// TODO: These intrinsics are not yet finalized, pending PTX ISA design work
def ISTYPEP_SAMPLER
  : NVPTXInst<(outs Int1Regs:$d), (ins Int64Regs:$a),
              "istypep.samplerref \t$d, $a;",
              [(set Int1Regs:$d, (int_nvvm_istypep_sampler Int64Regs:$a))]>;
def ISTYPEP_SURFACE
  : NVPTXInst<(outs Int1Regs:$d), (ins Int64Regs:$a),
              "istypep.surfref \t$d, $a;",
              [(set Int1Regs:$d, (int_nvvm_istypep_surface Int64Regs:$a))]>;
def ISTYPEP_TEXTURE
  : NVPTXInst<(outs Int1Regs:$d), (ins Int64Regs:$a),
              "istypep.texref \t$d, $a;",
              [(set Int1Regs:$d, (int_nvvm_istypep_texture Int64Regs:$a))]>;

//===- Surface Stores -----------------------------------------------------===//

// Unformatted

def SUST_B_1D_B8_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r),
              "sust.b.1d.b8.trap \t[$s, \\{$x\\}], \\{$r\\};",
              []>;
def SUST_B_1D_B16_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r),
              "sust.b.1d.b16.trap \t[$s, \\{$x\\}], \\{$r\\};",
              []>;
def SUST_B_1D_B32_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$r),
              "sust.b.1d.b32.trap \t[$s, \\{$x\\}], \\{$r\\};",
              []>;
def SUST_B_1D_V2B8_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g),
              "sust.b.1d.v2.b8.trap \t[$s, \\{$x\\}], \\{$r, $g\\};",
              []>;
def SUST_B_1D_V2B16_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g),
              "sust.b.1d.v2.b16.trap \t[$s, \\{$x\\}], \\{$r, $g\\};",
              []>;
def SUST_B_1D_V2B32_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$r, Int32Regs:$g),
              "sust.b.1d.v2.b32.trap \t[$s, \\{$x\\}], \\{$r, $g\\};",
              []>;
def SUST_B_1D_V4B8_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g,
                   Int16Regs:$b, Int16Regs:$a),
              "sust.b.1d.v4.b8.trap \t[$s, \\{$x\\}], \\{$r, $g, $b, $a\\};",
              []>;
def SUST_B_1D_V4B16_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g,
                   Int16Regs:$b, Int16Regs:$a),
              "sust.b.1d.v4.b16.trap \t[$s, \\{$x\\}], \\{$r, $g, $b, $a\\};",
              []>;
def SUST_B_1D_V4B32_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$r, Int32Regs:$g,
                   Int32Regs:$b, Int32Regs:$a),
              "sust.b.1d.v4.b32.trap \t[$s, \\{$x\\}], \\{$r, $g, $b, $a\\};",
              []>;


def SUST_B_1D_ARRAY_B8_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int16Regs:$r),
              "sust.b.a1d.b8.trap \t[$s, \\{$idx, $x\\}], \\{$r\\};",
              []>;
def SUST_B_1D_ARRAY_B16_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int16Regs:$r),
              "sust.b.a1d.b16.trap \t[$s, \\{$idx, $x\\}], \\{$r\\};",
              []>;
def SUST_B_1D_ARRAY_B32_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$r),
              "sust.b.a1d.b32.trap \t[$s, \\{$idx, $x\\}], \\{$r\\};",
              []>;
def SUST_B_1D_ARRAY_V2B8_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int16Regs:$r,
                   Int16Regs:$g),
              "sust.b.a1d.v2.b8.trap \t[$s, \\{$idx, $x\\}], \\{$r, $g\\};",
              []>;
def SUST_B_1D_ARRAY_V2B16_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int16Regs:$r,
                   Int16Regs:$g),
              "sust.b.a1d.v2.b16.trap \t[$s, \\{$idx, $x\\}], \\{$r, $g\\};",
              []>;
def SUST_B_1D_ARRAY_V2B32_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$r,
                   Int32Regs:$g),
              "sust.b.a1d.v2.b32.trap \t[$s, \\{$idx, $x\\}], \\{$r, $g\\};",
              []>;
def SUST_B_1D_ARRAY_V4B8_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int16Regs:$r,
                   Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
              "sust.b.a1d.v4.b8.trap \t[$s, \\{$idx, $x\\}], "
              "\\{$r, $g, $b, $a\\};",
              []>;
def SUST_B_1D_ARRAY_V4B16_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int16Regs:$r,
                   Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
             "sust.b.a1d.v4.b16.trap \t[$s, \\{$idx, $x\\}], "
             "\\{$r, $g, $b, $a\\};",
              []>;
def SUST_B_1D_ARRAY_V4B32_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$r,
                   Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
             "sust.b.a1d.v4.b32.trap \t[$s, \\{$idx, $x\\}], "
             "\\{$r, $g, $b, $a\\};",
              []>;


def SUST_B_2D_B8_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r),
              "sust.b.2d.b8.trap \t[$s, \\{$x, $y\\}], \\{$r\\};",
              []>;
def SUST_B_2D_B16_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r),
              "sust.b.2d.b16.trap \t[$s, \\{$x, $y\\}], \\{$r\\};",
              []>;
def SUST_B_2D_B32_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r),
              "sust.b.2d.b32.trap \t[$s, \\{$x, $y\\}], \\{$r\\};",
              []>;
def SUST_B_2D_V2B8_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r,
                   Int16Regs:$g),
              "sust.b.2d.v2.b8.trap \t[$s, \\{$x, $y\\}], \\{$r, $g\\};",
              []>;
def SUST_B_2D_V2B16_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r,
                   Int16Regs:$g),
              "sust.b.2d.v2.b16.trap \t[$s, \\{$x, $y\\}], \\{$r, $g\\};",
              []>;
def SUST_B_2D_V2B32_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r,
                   Int32Regs:$g),
              "sust.b.2d.v2.b32.trap \t[$s, \\{$x, $y\\}], \\{$r, $g\\};",
              []>;
def SUST_B_2D_V4B8_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r,
                   Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
              "sust.b.2d.v4.b8.trap \t[$s, \\{$x, $y\\}], "
              "\\{$r, $g, $b, $a\\};",
              []>;
def SUST_B_2D_V4B16_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r,
                   Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
             "sust.b.2d.v4.b16.trap \t[$s, \\{$x, $y\\}], "
             "\\{$r, $g, $b, $a\\};",
              []>;
def SUST_B_2D_V4B32_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r,
                   Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
             "sust.b.2d.v4.b32.trap \t[$s, \\{$x, $y\\}], "
             "\\{$r, $g, $b, $a\\};",
              []>;


def SUST_B_2D_ARRAY_B8_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
                   Int16Regs:$r),
              "sust.b.a2d.b8.trap \t[$s, \\{$idx, $x, $y, $y\\}], \\{$r\\};",
              []>;
def SUST_B_2D_ARRAY_B16_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
                   Int16Regs:$r),
              "sust.b.a2d.b16.trap \t[$s, \\{$idx, $x, $y, $y\\}], \\{$r\\};",
              []>;
def SUST_B_2D_ARRAY_B32_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
                   Int32Regs:$r),
              "sust.b.a2d.b32.trap \t[$s, \\{$idx, $x, $y, $y\\}], \\{$r\\};",
              []>;
def SUST_B_2D_ARRAY_V2B8_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
                   Int16Regs:$r, Int16Regs:$g),
              "sust.b.a2d.v2.b8.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
              "\\{$r, $g\\};",
              []>;
def SUST_B_2D_ARRAY_V2B16_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
                   Int16Regs:$r, Int16Regs:$g),
             "sust.b.a2d.v2.b16.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
             "\\{$r, $g\\};",
              []>;
def SUST_B_2D_ARRAY_V2B32_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
                   Int32Regs:$r, Int32Regs:$g),
             "sust.b.a2d.v2.b32.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
             "\\{$r, $g\\};",
              []>;
def SUST_B_2D_ARRAY_V4B8_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
                   Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
      "sust.b.a2d.v4.b8.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
      "\\{$r, $g, $b, $a\\};",
              []>;
def SUST_B_2D_ARRAY_V4B16_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
                   Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
     "sust.b.a2d.v4.b16.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
     "\\{$r, $g, $b, $a\\};",
              []>;
def SUST_B_2D_ARRAY_V4B32_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
                   Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
     "sust.b.a2d.v4.b32.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
     "\\{$r, $g, $b, $a\\};",
              []>;


def SUST_B_3D_B8_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
                   Int16Regs:$r),
              "sust.b.3d.b8.trap \t[$s, \\{$x, $y, $z, $z\\}], \\{$r\\};",
              []>;
def SUST_B_3D_B16_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
                   Int16Regs:$r),
              "sust.b.3d.b16.trap \t[$s, \\{$x, $y, $z, $z\\}], \\{$r\\};",
              []>;
def SUST_B_3D_B32_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
                   Int32Regs:$r),
              "sust.b.3d.b32.trap \t[$s, \\{$x, $y, $z, $z\\}], \\{$r\\};",
              []>;
def SUST_B_3D_V2B8_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
                   Int16Regs:$r, Int16Regs:$g),
              "sust.b.3d.v2.b8.trap \t[$s, \\{$x, $y, $z, $z\\}], "
              "\\{$r, $g\\};",
              []>;
def SUST_B_3D_V2B16_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
                   Int16Regs:$r, Int16Regs:$g),
              "sust.b.3d.v2.b16.trap \t[$s, \\{$x, $y, $z, $z\\}], "
              "\\{$r, $g\\};",
              []>;
def SUST_B_3D_V2B32_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
                   Int32Regs:$r, Int32Regs:$g),
              "sust.b.3d.v2.b32.trap \t[$s, \\{$x, $y, $z, $z\\}], "
              "\\{$r, $g\\};",
              []>;
def SUST_B_3D_V4B8_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
                   Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
         "sust.b.3d.v4.b8.trap \t[$s, \\{$x, $y, $z, $z\\}], "
         "\\{$r, $g, $b, $a\\};",
              []>;
def SUST_B_3D_V4B16_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
                   Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
        "sust.b.3d.v4.b16.trap \t[$s, \\{$x, $y, $z, $z\\}], "
        "\\{$r, $g, $b, $a\\};",
              []>;
def SUST_B_3D_V4B32_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
                   Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
        "sust.b.3d.v4.b32.trap \t[$s, \\{$x, $y, $z, $z\\}], "
        "\\{$r, $g, $b, $a\\};",
              []>;

// Formatted

def SUST_P_1D_B8_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r),
              "sust.p.1d.b8.trap \t[$s, \\{$x\\}], \\{$r\\};",
              []>;
def SUST_P_1D_B16_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r),
              "sust.p.1d.b16.trap \t[$s, \\{$x\\}], \\{$r\\};",
              []>;
def SUST_P_1D_B32_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$r),
              "sust.p.1d.b32.trap \t[$s, \\{$x\\}], \\{$r\\};",
              []>;
def SUST_P_1D_V2B8_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g),
              "sust.p.1d.v2.b8.trap \t[$s, \\{$x\\}], \\{$r, $g\\};",
              []>;
def SUST_P_1D_V2B16_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g),
              "sust.p.1d.v2.b16.trap \t[$s, \\{$x\\}], \\{$r, $g\\};",
              []>;
def SUST_P_1D_V2B32_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$r, Int32Regs:$g),
              "sust.p.1d.v2.b32.trap \t[$s, \\{$x\\}], \\{$r, $g\\};",
              []>;
def SUST_P_1D_V4B8_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g,
                   Int16Regs:$b, Int16Regs:$a),
              "sust.p.1d.v4.b8.trap \t[$s, \\{$x\\}], \\{$r, $g, $b, $a\\};",
              []>;
def SUST_P_1D_V4B16_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g,
                   Int16Regs:$b, Int16Regs:$a),
              "sust.p.1d.v4.b16.trap \t[$s, \\{$x\\}], \\{$r, $g, $b, $a\\};",
              []>;
def SUST_P_1D_V4B32_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$r, Int32Regs:$g,
                   Int32Regs:$b, Int32Regs:$a),
              "sust.p.1d.v4.b32.trap \t[$s, \\{$x\\}], \\{$r, $g, $b, $a\\};",
              []>;


def SUST_P_1D_ARRAY_B8_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int16Regs:$r),
              "sust.p.a1d.b8.trap \t[$s, \\{$idx, $x\\}], \\{$r\\};",
              []>;
def SUST_P_1D_ARRAY_B16_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int16Regs:$r),
              "sust.p.a1d.b16.trap \t[$s, \\{$idx, $x\\}], \\{$r\\};",
              []>;
def SUST_P_1D_ARRAY_B32_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$r),
              "sust.p.a1d.b32.trap \t[$s, \\{$idx, $x\\}], \\{$r\\};",
              []>;
def SUST_P_1D_ARRAY_V2B8_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int16Regs:$r,
                   Int16Regs:$g),
              "sust.p.a1d.v2.b8.trap \t[$s, \\{$idx, $x\\}], \\{$r, $g\\};",
              []>;
def SUST_P_1D_ARRAY_V2B16_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int16Regs:$r,
                   Int16Regs:$g),
              "sust.p.a1d.v2.b16.trap \t[$s, \\{$idx, $x\\}], \\{$r, $g\\};",
              []>;
def SUST_P_1D_ARRAY_V2B32_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$r,
                   Int32Regs:$g),
              "sust.p.a1d.v2.b32.trap \t[$s, \\{$idx, $x\\}], \\{$r, $g\\};",
              []>;
def SUST_P_1D_ARRAY_V4B8_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int16Regs:$r,
                   Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
              "sust.p.a1d.v4.b8.trap \t[$s, \\{$idx, $x\\}], "
              "\\{$r, $g, $b, $a\\};",
              []>;
def SUST_P_1D_ARRAY_V4B16_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int16Regs:$r,
                   Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
             "sust.p.a1d.v4.b16.trap \t[$s, \\{$idx, $x\\}], "
             "\\{$r, $g, $b, $a\\};",
              []>;
def SUST_P_1D_ARRAY_V4B32_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$r,
                   Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
             "sust.p.a1d.v4.b32.trap \t[$s, \\{$idx, $x\\}], "
             "\\{$r, $g, $b, $a\\};",
              []>;


def SUST_P_2D_B8_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r),
              "sust.p.2d.b8.trap \t[$s, \\{$x, $y\\}], \\{$r\\};",
              []>;
def SUST_P_2D_B16_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r),
              "sust.p.2d.b16.trap \t[$s, \\{$x, $y\\}], \\{$r\\};",
              []>;
def SUST_P_2D_B32_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r),
              "sust.p.2d.b32.trap \t[$s, \\{$x, $y\\}], \\{$r\\};",
              []>;
def SUST_P_2D_V2B8_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r,
                   Int16Regs:$g),
              "sust.p.2d.v2.b8.trap \t[$s, \\{$x, $y\\}], \\{$r, $g\\};",
              []>;
def SUST_P_2D_V2B16_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r,
                   Int16Regs:$g),
              "sust.p.2d.v2.b16.trap \t[$s, \\{$x, $y\\}], \\{$r, $g\\};",
              []>;
def SUST_P_2D_V2B32_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r,
                   Int32Regs:$g),
              "sust.p.2d.v2.b32.trap \t[$s, \\{$x, $y\\}], \\{$r, $g\\};",
              []>;
def SUST_P_2D_V4B8_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r,
                   Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
              "sust.p.2d.v4.b8.trap \t[$s, \\{$x, $y\\}], "
              "\\{$r, $g, $b, $a\\};",
              []>;
def SUST_P_2D_V4B16_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r,
                   Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
             "sust.p.2d.v4.b16.trap \t[$s, \\{$x, $y\\}], "
             "\\{$r, $g, $b, $a\\};",
              []>;
def SUST_P_2D_V4B32_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r,
                   Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
             "sust.p.2d.v4.b32.trap \t[$s, \\{$x, $y\\}], "
             "\\{$r, $g, $b, $a\\};",
              []>;


def SUST_P_2D_ARRAY_B8_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
                   Int16Regs:$r),
              "sust.p.a2d.b8.trap \t[$s, \\{$idx, $x, $y, $y\\}], \\{$r\\};",
              []>;
def SUST_P_2D_ARRAY_B16_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
                   Int16Regs:$r),
              "sust.p.a2d.b16.trap \t[$s, \\{$idx, $x, $y, $y\\}], \\{$r\\};",
              []>;
def SUST_P_2D_ARRAY_B32_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
                   Int32Regs:$r),
              "sust.p.a2d.b32.trap \t[$s, \\{$idx, $x, $y, $y\\}], \\{$r\\};",
              []>;
def SUST_P_2D_ARRAY_V2B8_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
                   Int16Regs:$r, Int16Regs:$g),
              "sust.p.a2d.v2.b8.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
              "\\{$r, $g\\};",
              []>;
def SUST_P_2D_ARRAY_V2B16_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
                   Int16Regs:$r, Int16Regs:$g),
             "sust.p.a2d.v2.b16.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
             "\\{$r, $g\\};",
              []>;
def SUST_P_2D_ARRAY_V2B32_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
                   Int32Regs:$r, Int32Regs:$g),
             "sust.p.a2d.v2.b32.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
             "\\{$r, $g\\};",
              []>;
def SUST_P_2D_ARRAY_V4B8_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
                   Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
      "sust.p.a2d.v4.b8.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
      "\\{$r, $g, $b, $a\\};",
              []>;
def SUST_P_2D_ARRAY_V4B16_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
                   Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
     "sust.p.a2d.v4.b16.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
     "\\{$r, $g, $b, $a\\};",
              []>;
def SUST_P_2D_ARRAY_V4B32_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
                   Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
     "sust.p.a2d.v4.b32.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
     "\\{$r, $g, $b, $a\\};",
              []>;


def SUST_P_3D_B8_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
                   Int16Regs:$r),
              "sust.p.3d.b8.trap \t[$s, \\{$x, $y, $z, $z\\}], \\{$r\\};",
              []>;
def SUST_P_3D_B16_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
                   Int16Regs:$r),
              "sust.p.3d.b16.trap \t[$s, \\{$x, $y, $z, $z\\}], \\{$r\\};",
              []>;
def SUST_P_3D_B32_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
                   Int32Regs:$r),
              "sust.p.3d.b32.trap \t[$s, \\{$x, $y, $z, $z\\}], \\{$r\\};",
              []>;
def SUST_P_3D_V2B8_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
                   Int16Regs:$r, Int16Regs:$g),
              "sust.p.3d.v2.b8.trap \t[$s, \\{$x, $y, $z, $z\\}], "
              "\\{$r, $g\\};",
              []>;
def SUST_P_3D_V2B16_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
                   Int16Regs:$r, Int16Regs:$g),
              "sust.p.3d.v2.b16.trap \t[$s, \\{$x, $y, $z, $z\\}], "
              "\\{$r, $g\\};",
              []>;
def SUST_P_3D_V2B32_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
                   Int32Regs:$r, Int32Regs:$g),
              "sust.p.3d.v2.b32.trap \t[$s, \\{$x, $y, $z, $z\\}], "
              "\\{$r, $g\\};",
              []>;
def SUST_P_3D_V4B8_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
                   Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
         "sust.p.3d.v4.b8.trap \t[$s, \\{$x, $y, $z, $z\\}], "
         "\\{$r, $g, $b, $a\\};",
              []>;
def SUST_P_3D_V4B16_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
                   Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
        "sust.p.3d.v4.b16.trap \t[$s, \\{$x, $y, $z, $z\\}], "
        "\\{$r, $g, $b, $a\\};",
              []>;
def SUST_P_3D_V4B32_TRAP
  : NVPTXInst<(outs),
              (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
                   Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
        "sust.p.3d.v4.b32.trap \t[$s, \\{$x, $y, $z, $z\\}], "
        "\\{$r, $g, $b, $a\\};",
              []>;


// Surface store instruction patterns
// I'm not sure why we can't just include these in the instruction definitions,
// but TableGen complains of type errors :(

def : Pat<(int_nvvm_sust_b_1d_i8_trap
           Int64Regs:$s, Int32Regs:$x, Int16Regs:$r),
          (SUST_B_1D_B8_TRAP Int64Regs:$s, Int32Regs:$x, Int16Regs:$r)>;

def : Pat<(int_nvvm_sust_b_1d_i16_trap
           Int64Regs:$s, Int32Regs:$x, Int16Regs:$r),
          (SUST_B_1D_B16_TRAP Int64Regs:$s, Int32Regs:$x, Int16Regs:$r)>;

def : Pat<(int_nvvm_sust_b_1d_i32_trap
           Int64Regs:$s, Int32Regs:$x, Int32Regs:$r),
          (SUST_B_1D_B32_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$r)>;

def : Pat<(int_nvvm_sust_b_1d_v2i8_trap
           Int64Regs:$s, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g),
          (SUST_B_1D_V2B8_TRAP Int64Regs:$s, Int32Regs:$x,
           Int16Regs:$r, Int16Regs:$g)>;

def : Pat<(int_nvvm_sust_b_1d_v2i16_trap
           Int64Regs:$s, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g),
          (SUST_B_1D_V2B16_TRAP Int64Regs:$s, Int32Regs:$x,
           Int16Regs:$r, Int16Regs:$g)>;

def : Pat<(int_nvvm_sust_b_1d_v2i32_trap
           Int64Regs:$s, Int32Regs:$x, Int32Regs:$r, Int32Regs:$g),
          (SUST_B_1D_V2B32_TRAP Int64Regs:$s, Int32Regs:$x,
           Int32Regs:$r, Int32Regs:$g)>;

def : Pat<(int_nvvm_sust_b_1d_v4i8_trap
           Int64Regs:$s, Int32Regs:$x,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
          (SUST_B_1D_V4B8_TRAP Int64Regs:$s, Int32Regs:$x,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;

def : Pat<(int_nvvm_sust_b_1d_v4i16_trap
           Int64Regs:$s, Int32Regs:$x,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
          (SUST_B_1D_V4B16_TRAP Int64Regs:$s, Int32Regs:$x,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;

def : Pat<(int_nvvm_sust_b_1d_v4i32_trap
           Int64Regs:$s, Int32Regs:$x,
           Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
          (SUST_B_1D_V4B32_TRAP Int64Regs:$s, Int32Regs:$x,
           Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a)>;



def : Pat<(int_nvvm_sust_b_1d_array_i8_trap
           Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int16Regs:$r),
          (SUST_B_1D_ARRAY_B8_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
           Int16Regs:$r)>;

def : Pat<(int_nvvm_sust_b_1d_array_i16_trap
           Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int16Regs:$r),
          (SUST_B_1D_ARRAY_B16_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
           Int16Regs:$r)>;

def : Pat<(int_nvvm_sust_b_1d_array_i32_trap
           Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$r),
          (SUST_B_1D_ARRAY_B32_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
           Int32Regs:$r)>;

def : Pat<(int_nvvm_sust_b_1d_array_v2i8_trap
          Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g),
          (SUST_B_1D_ARRAY_V2B8_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
           Int16Regs:$r, Int16Regs:$g)>;

def : Pat<(int_nvvm_sust_b_1d_array_v2i16_trap
          Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g),
          (SUST_B_1D_ARRAY_V2B16_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
           Int16Regs:$r, Int16Regs:$g)>;

def : Pat<(int_nvvm_sust_b_1d_array_v2i32_trap
          Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$r, Int32Regs:$g),
          (SUST_B_1D_ARRAY_V2B32_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
           Int32Regs:$r, Int32Regs:$g)>;

def : Pat<(int_nvvm_sust_b_1d_array_v4i8_trap
           Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
          (SUST_B_1D_ARRAY_V4B8_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;

def : Pat<(int_nvvm_sust_b_1d_array_v4i16_trap
           Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
          (SUST_B_1D_ARRAY_V4B16_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;

def : Pat<(int_nvvm_sust_b_1d_array_v4i32_trap
           Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
           Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
          (SUST_B_1D_ARRAY_V4B32_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
           Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a)>;



def : Pat<(int_nvvm_sust_b_2d_i8_trap
           Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r),
          (SUST_B_2D_B8_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
           Int16Regs:$r)>;

def : Pat<(int_nvvm_sust_b_2d_i16_trap
           Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r),
          (SUST_B_2D_B16_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
           Int16Regs:$r)>;

def : Pat<(int_nvvm_sust_b_2d_i32_trap
           Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r),
          (SUST_B_2D_B32_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
           Int32Regs:$r)>;

def : Pat<(int_nvvm_sust_b_2d_v2i8_trap
          Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r, Int16Regs:$g),
          (SUST_B_2D_V2B8_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
           Int16Regs:$r, Int16Regs:$g)>;

def : Pat<(int_nvvm_sust_b_2d_v2i16_trap
          Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r, Int16Regs:$g),
          (SUST_B_2D_V2B16_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
           Int16Regs:$r, Int16Regs:$g)>;

def : Pat<(int_nvvm_sust_b_2d_v2i32_trap
          Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r, Int32Regs:$g),
          (SUST_B_2D_V2B32_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
           Int32Regs:$r, Int32Regs:$g)>;

def : Pat<(int_nvvm_sust_b_2d_v4i8_trap
           Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
          (SUST_B_2D_V4B8_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;

def : Pat<(int_nvvm_sust_b_2d_v4i16_trap
           Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
          (SUST_B_2D_V4B16_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;

def : Pat<(int_nvvm_sust_b_2d_v4i32_trap
           Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
           Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
          (SUST_B_2D_V4B32_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
           Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a)>;



def : Pat<(int_nvvm_sust_b_2d_array_i8_trap
          Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r),
          (SUST_B_2D_ARRAY_B8_TRAP Int64Regs:$s,
           Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
           Int16Regs:$r)>;

def : Pat<(int_nvvm_sust_b_2d_array_i16_trap
          Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r),
          (SUST_B_2D_ARRAY_B16_TRAP Int64Regs:$s,
           Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
           Int16Regs:$r)>;

def : Pat<(int_nvvm_sust_b_2d_array_i32_trap
          Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r),
          (SUST_B_2D_ARRAY_B32_TRAP Int64Regs:$s,
           Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
           Int32Regs:$r)>;

def : Pat<(int_nvvm_sust_b_2d_array_v2i8_trap
           Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
           Int16Regs:$r, Int16Regs:$g),
          (SUST_B_2D_ARRAY_V2B8_TRAP Int64Regs:$s, Int32Regs:$l,
           Int32Regs:$x, Int32Regs:$y,
           Int16Regs:$r, Int16Regs:$g)>;

def : Pat<(int_nvvm_sust_b_2d_array_v2i16_trap
           Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
           Int16Regs:$r, Int16Regs:$g),
          (SUST_B_2D_ARRAY_V2B16_TRAP Int64Regs:$s, Int32Regs:$l,
           Int32Regs:$x, Int32Regs:$y,
           Int16Regs:$r, Int16Regs:$g)>;

def : Pat<(int_nvvm_sust_b_2d_array_v2i32_trap
           Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r,
           Int32Regs:$g),
          (SUST_B_2D_ARRAY_V2B32_TRAP Int64Regs:$s, Int32Regs:$l,
           Int32Regs:$x, Int32Regs:$y, Int32Regs:$r, Int32Regs:$g)>;

def : Pat<(int_nvvm_sust_b_2d_array_v4i8_trap
           Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
          (SUST_B_2D_ARRAY_V4B8_TRAP Int64Regs:$s,
           Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;

def : Pat<(int_nvvm_sust_b_2d_array_v4i16_trap
           Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
          (SUST_B_2D_ARRAY_V4B16_TRAP Int64Regs:$s,
           Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;

def : Pat<(int_nvvm_sust_b_2d_array_v4i32_trap
           Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
           Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
          (SUST_B_2D_ARRAY_V4B32_TRAP Int64Regs:$s, Int32Regs:$l,
           Int32Regs:$x, Int32Regs:$y,
           Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a)>;



def : Pat<(int_nvvm_sust_b_3d_i8_trap
           Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
           Int16Regs:$r),
          (SUST_B_3D_B8_TRAP Int64Regs:$s,
           Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
           Int16Regs:$r)>;

def : Pat<(int_nvvm_sust_b_3d_i16_trap
           Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
           Int16Regs:$r),
          (SUST_B_3D_B16_TRAP Int64Regs:$s,
           Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
           Int16Regs:$r)>;

def : Pat<(int_nvvm_sust_b_3d_i32_trap
           Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
           Int32Regs:$r),
          (SUST_B_3D_B32_TRAP Int64Regs:$s,
           Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
           Int32Regs:$r)>;

def : Pat<(int_nvvm_sust_b_3d_v2i8_trap
           Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
           Int16Regs:$r, Int16Regs:$g),
          (SUST_B_3D_V2B8_TRAP Int64Regs:$s,
           Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
           Int16Regs:$r, Int16Regs:$g)>;

def : Pat<(int_nvvm_sust_b_3d_v2i16_trap
           Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
           Int16Regs:$r, Int16Regs:$g),
          (SUST_B_3D_V2B16_TRAP Int64Regs:$s,
           Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
           Int16Regs:$r, Int16Regs:$g)>;

def : Pat<(int_nvvm_sust_b_3d_v2i32_trap
           Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
           Int32Regs:$r, Int32Regs:$g),
          (SUST_B_3D_V2B32_TRAP Int64Regs:$s,
           Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
           Int32Regs:$r, Int32Regs:$g)>;

def : Pat<(int_nvvm_sust_b_3d_v4i8_trap
           Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
          (SUST_B_3D_V4B8_TRAP Int64Regs:$s,
           Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;

def : Pat<(int_nvvm_sust_b_3d_v4i16_trap
           Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
          (SUST_B_3D_V4B16_TRAP Int64Regs:$s,
           Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;

def : Pat<(int_nvvm_sust_b_3d_v4i32_trap
           Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
           Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
          (SUST_B_3D_V4B32_TRAP Int64Regs:$s,
           Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
           Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a)>;




def : Pat<(int_nvvm_sust_p_1d_i8_trap
           Int64Regs:$s, Int32Regs:$x, Int16Regs:$r),
          (SUST_P_1D_B8_TRAP Int64Regs:$s, Int32Regs:$x, Int16Regs:$r)>;

def : Pat<(int_nvvm_sust_p_1d_i16_trap
           Int64Regs:$s, Int32Regs:$x, Int16Regs:$r),
          (SUST_P_1D_B16_TRAP Int64Regs:$s, Int32Regs:$x, Int16Regs:$r)>;

def : Pat<(int_nvvm_sust_p_1d_i32_trap
           Int64Regs:$s, Int32Regs:$x, Int32Regs:$r),
          (SUST_P_1D_B32_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$r)>;

def : Pat<(int_nvvm_sust_p_1d_v2i8_trap
           Int64Regs:$s, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g),
          (SUST_P_1D_V2B8_TRAP Int64Regs:$s, Int32Regs:$x,
           Int16Regs:$r, Int16Regs:$g)>;

def : Pat<(int_nvvm_sust_p_1d_v2i16_trap
           Int64Regs:$s, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g),
          (SUST_P_1D_V2B16_TRAP Int64Regs:$s, Int32Regs:$x,
           Int16Regs:$r, Int16Regs:$g)>;

def : Pat<(int_nvvm_sust_p_1d_v2i32_trap
           Int64Regs:$s, Int32Regs:$x, Int32Regs:$r, Int32Regs:$g),
          (SUST_P_1D_V2B32_TRAP Int64Regs:$s, Int32Regs:$x,
           Int32Regs:$r, Int32Regs:$g)>;

def : Pat<(int_nvvm_sust_p_1d_v4i8_trap
           Int64Regs:$s, Int32Regs:$x,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
          (SUST_P_1D_V4B8_TRAP Int64Regs:$s, Int32Regs:$x,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;

def : Pat<(int_nvvm_sust_p_1d_v4i16_trap
           Int64Regs:$s, Int32Regs:$x,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
          (SUST_P_1D_V4B16_TRAP Int64Regs:$s, Int32Regs:$x,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;

def : Pat<(int_nvvm_sust_p_1d_v4i32_trap
           Int64Regs:$s, Int32Regs:$x,
           Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
          (SUST_P_1D_V4B32_TRAP Int64Regs:$s, Int32Regs:$x,
           Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a)>;



def : Pat<(int_nvvm_sust_p_1d_array_i8_trap
           Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int16Regs:$r),
          (SUST_P_1D_ARRAY_B8_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
           Int16Regs:$r)>;

def : Pat<(int_nvvm_sust_p_1d_array_i16_trap
           Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int16Regs:$r),
          (SUST_P_1D_ARRAY_B16_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
           Int16Regs:$r)>;

def : Pat<(int_nvvm_sust_p_1d_array_i32_trap
           Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$r),
          (SUST_P_1D_ARRAY_B32_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
           Int32Regs:$r)>;

def : Pat<(int_nvvm_sust_p_1d_array_v2i8_trap
          Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g),
          (SUST_P_1D_ARRAY_V2B8_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
           Int16Regs:$r, Int16Regs:$g)>;

def : Pat<(int_nvvm_sust_p_1d_array_v2i16_trap
          Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g),
          (SUST_P_1D_ARRAY_V2B16_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
           Int16Regs:$r, Int16Regs:$g)>;

def : Pat<(int_nvvm_sust_p_1d_array_v2i32_trap
          Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$r, Int32Regs:$g),
          (SUST_P_1D_ARRAY_V2B32_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
           Int32Regs:$r, Int32Regs:$g)>;

def : Pat<(int_nvvm_sust_p_1d_array_v4i8_trap
           Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
          (SUST_P_1D_ARRAY_V4B8_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;

def : Pat<(int_nvvm_sust_p_1d_array_v4i16_trap
           Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
          (SUST_P_1D_ARRAY_V4B16_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;

def : Pat<(int_nvvm_sust_p_1d_array_v4i32_trap
           Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
           Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
          (SUST_P_1D_ARRAY_V4B32_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
           Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a)>;



def : Pat<(int_nvvm_sust_p_2d_i8_trap
           Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r),
          (SUST_P_2D_B8_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
           Int16Regs:$r)>;

def : Pat<(int_nvvm_sust_p_2d_i16_trap
           Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r),
          (SUST_P_2D_B16_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
           Int16Regs:$r)>;

def : Pat<(int_nvvm_sust_p_2d_i32_trap
           Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r),
          (SUST_P_2D_B32_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
           Int32Regs:$r)>;

def : Pat<(int_nvvm_sust_p_2d_v2i8_trap
          Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r, Int16Regs:$g),
          (SUST_P_2D_V2B8_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
           Int16Regs:$r, Int16Regs:$g)>;

def : Pat<(int_nvvm_sust_p_2d_v2i16_trap
          Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r, Int16Regs:$g),
          (SUST_P_2D_V2B16_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
           Int16Regs:$r, Int16Regs:$g)>;

def : Pat<(int_nvvm_sust_p_2d_v2i32_trap
          Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r, Int32Regs:$g),
          (SUST_P_2D_V2B32_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
           Int32Regs:$r, Int32Regs:$g)>;

def : Pat<(int_nvvm_sust_p_2d_v4i8_trap
           Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
          (SUST_P_2D_V4B8_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;

def : Pat<(int_nvvm_sust_p_2d_v4i16_trap
           Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
          (SUST_P_2D_V4B16_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;

def : Pat<(int_nvvm_sust_p_2d_v4i32_trap
           Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
           Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
          (SUST_P_2D_V4B32_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
           Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a)>;



def : Pat<(int_nvvm_sust_p_2d_array_i8_trap
          Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r),
          (SUST_P_2D_ARRAY_B8_TRAP Int64Regs:$s,
           Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
           Int16Regs:$r)>;

def : Pat<(int_nvvm_sust_p_2d_array_i16_trap
          Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r),
          (SUST_P_2D_ARRAY_B16_TRAP Int64Regs:$s,
           Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
           Int16Regs:$r)>;

def : Pat<(int_nvvm_sust_p_2d_array_i32_trap
          Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r),
          (SUST_P_2D_ARRAY_B32_TRAP Int64Regs:$s,
           Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
           Int32Regs:$r)>;

def : Pat<(int_nvvm_sust_p_2d_array_v2i8_trap
           Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
           Int16Regs:$r, Int16Regs:$g),
          (SUST_P_2D_ARRAY_V2B8_TRAP Int64Regs:$s, Int32Regs:$l,
           Int32Regs:$x, Int32Regs:$y,
           Int16Regs:$r, Int16Regs:$g)>;

def : Pat<(int_nvvm_sust_p_2d_array_v2i16_trap
           Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
           Int16Regs:$r, Int16Regs:$g),
          (SUST_P_2D_ARRAY_V2B16_TRAP Int64Regs:$s, Int32Regs:$l,
           Int32Regs:$x, Int32Regs:$y,
           Int16Regs:$r, Int16Regs:$g)>;

def : Pat<(int_nvvm_sust_p_2d_array_v2i32_trap
           Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r,
           Int32Regs:$g),
          (SUST_P_2D_ARRAY_V2B32_TRAP Int64Regs:$s, Int32Regs:$l,
           Int32Regs:$x, Int32Regs:$y, Int32Regs:$r, Int32Regs:$g)>;

def : Pat<(int_nvvm_sust_p_2d_array_v4i8_trap
           Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
          (SUST_P_2D_ARRAY_V4B8_TRAP Int64Regs:$s,
           Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;

def : Pat<(int_nvvm_sust_p_2d_array_v4i16_trap
           Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
          (SUST_P_2D_ARRAY_V4B16_TRAP Int64Regs:$s,
           Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;

def : Pat<(int_nvvm_sust_p_2d_array_v4i32_trap
           Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
           Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
          (SUST_P_2D_ARRAY_V4B32_TRAP Int64Regs:$s, Int32Regs:$l,
           Int32Regs:$x, Int32Regs:$y,
           Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a)>;



def : Pat<(int_nvvm_sust_p_3d_i8_trap
           Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
           Int16Regs:$r),
          (SUST_P_3D_B8_TRAP Int64Regs:$s,
           Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
           Int16Regs:$r)>;

def : Pat<(int_nvvm_sust_p_3d_i16_trap
           Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
           Int16Regs:$r),
          (SUST_P_3D_B16_TRAP Int64Regs:$s,
           Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
           Int16Regs:$r)>;

def : Pat<(int_nvvm_sust_p_3d_i32_trap
           Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
           Int32Regs:$r),
          (SUST_P_3D_B32_TRAP Int64Regs:$s,
           Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
           Int32Regs:$r)>;

def : Pat<(int_nvvm_sust_p_3d_v2i8_trap
           Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
           Int16Regs:$r, Int16Regs:$g),
          (SUST_P_3D_V2B8_TRAP Int64Regs:$s,
           Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
           Int16Regs:$r, Int16Regs:$g)>;

def : Pat<(int_nvvm_sust_p_3d_v2i16_trap
           Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
           Int16Regs:$r, Int16Regs:$g),
          (SUST_P_3D_V2B16_TRAP Int64Regs:$s,
           Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
           Int16Regs:$r, Int16Regs:$g)>;

def : Pat<(int_nvvm_sust_p_3d_v2i32_trap
           Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
           Int32Regs:$r, Int32Regs:$g),
          (SUST_P_3D_V2B32_TRAP Int64Regs:$s,
           Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
           Int32Regs:$r, Int32Regs:$g)>;

def : Pat<(int_nvvm_sust_p_3d_v4i8_trap
           Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
          (SUST_P_3D_V4B8_TRAP Int64Regs:$s,
           Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;

def : Pat<(int_nvvm_sust_p_3d_v4i16_trap
           Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
          (SUST_P_3D_V4B16_TRAP Int64Regs:$s,
           Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
           Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;

def : Pat<(int_nvvm_sust_p_3d_v4i32_trap
           Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
           Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
          (SUST_P_3D_V4B32_TRAP Int64Regs:$s,
           Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
           Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a)>;



//===-- Old PTX Back-end Intrinsics ---------------------------------------===//

// These intrinsics are handled to retain compatibility with the old backend.

// PTX Special Purpose Register Accessor Intrinsics

class PTX_READ_SPECIAL_REGISTER_R64<string regname, Intrinsic intop>
  : NVPTXInst<(outs Int64Regs:$d), (ins),
              !strconcat(!strconcat("mov.u64\t$d, %", regname), ";"),
              [(set Int64Regs:$d, (intop))]>;

class PTX_READ_SPECIAL_REGISTER_R32<string regname, Intrinsic intop>
  : NVPTXInst<(outs Int32Regs:$d), (ins),
              !strconcat(!strconcat("mov.u32\t$d, %", regname), ";"),
              [(set Int32Regs:$d, (intop))]>;

// TODO Add read vector-version of special registers

def PTX_READ_TID_X   : PTX_READ_SPECIAL_REGISTER_R32<"tid.x",
                                                     int_ptx_read_tid_x>;
def PTX_READ_TID_Y   : PTX_READ_SPECIAL_REGISTER_R32<"tid.y",
                                                     int_ptx_read_tid_y>;
def PTX_READ_TID_Z   : PTX_READ_SPECIAL_REGISTER_R32<"tid.z",
                                                     int_ptx_read_tid_z>;
def PTX_READ_TID_W   : PTX_READ_SPECIAL_REGISTER_R32<"tid.w",
                                                     int_ptx_read_tid_w>;

def PTX_READ_NTID_X   : PTX_READ_SPECIAL_REGISTER_R32<"ntid.x",
                                                      int_ptx_read_ntid_x>;
def PTX_READ_NTID_Y   : PTX_READ_SPECIAL_REGISTER_R32<"ntid.y",
                                                      int_ptx_read_ntid_y>;
def PTX_READ_NTID_Z   : PTX_READ_SPECIAL_REGISTER_R32<"ntid.z",
                                                      int_ptx_read_ntid_z>;
def PTX_READ_NTID_W   : PTX_READ_SPECIAL_REGISTER_R32<"ntid.w",
                                                      int_ptx_read_ntid_w>;

def PTX_READ_LANEID  : PTX_READ_SPECIAL_REGISTER_R32<"laneid",
                                                     int_ptx_read_laneid>;
def PTX_READ_WARPID  : PTX_READ_SPECIAL_REGISTER_R32<"warpid",
                                                     int_ptx_read_warpid>;
def PTX_READ_NWARPID : PTX_READ_SPECIAL_REGISTER_R32<"nwarpid",
                                                     int_ptx_read_nwarpid>;

def PTX_READ_CTAID_X   : PTX_READ_SPECIAL_REGISTER_R32<"ctaid.x",
                                                       int_ptx_read_ctaid_x>;
def PTX_READ_CTAID_Y   : PTX_READ_SPECIAL_REGISTER_R32<"ctaid.y",
                                                       int_ptx_read_ctaid_y>;
def PTX_READ_CTAID_Z   : PTX_READ_SPECIAL_REGISTER_R32<"ctaid.z",
                                                       int_ptx_read_ctaid_z>;
def PTX_READ_CTAID_W   : PTX_READ_SPECIAL_REGISTER_R32<"ctaid.w",
                                                       int_ptx_read_ctaid_w>;

def PTX_READ_NCTAID_X   : PTX_READ_SPECIAL_REGISTER_R32<"nctaid.x",
                                                        int_ptx_read_nctaid_x>;
def PTX_READ_NCTAID_Y   : PTX_READ_SPECIAL_REGISTER_R32<"nctaid.y",
                                                        int_ptx_read_nctaid_y>;
def PTX_READ_NCTAID_Z   : PTX_READ_SPECIAL_REGISTER_R32<"nctaid.z",
                                                        int_ptx_read_nctaid_z>;
def PTX_READ_NCTAID_W   : PTX_READ_SPECIAL_REGISTER_R32<"nctaid.w",
                                                        int_ptx_read_nctaid_w>;

def PTX_READ_SMID  : PTX_READ_SPECIAL_REGISTER_R32<"smid",
                                                   int_ptx_read_smid>;
def PTX_READ_NSMID  : PTX_READ_SPECIAL_REGISTER_R32<"nsmid",
                                                    int_ptx_read_nsmid>;
def PTX_READ_GRIDID  : PTX_READ_SPECIAL_REGISTER_R32<"gridid",
                                                     int_ptx_read_gridid>;

def PTX_READ_LANEMASK_EQ
  : PTX_READ_SPECIAL_REGISTER_R32<"lanemask_eq", int_ptx_read_lanemask_eq>;
def PTX_READ_LANEMASK_LE
  : PTX_READ_SPECIAL_REGISTER_R32<"lanemask_le", int_ptx_read_lanemask_le>;
def PTX_READ_LANEMASK_LT
  : PTX_READ_SPECIAL_REGISTER_R32<"lanemask_lt", int_ptx_read_lanemask_lt>;
def PTX_READ_LANEMASK_GE
  : PTX_READ_SPECIAL_REGISTER_R32<"lanemask_ge", int_ptx_read_lanemask_ge>;
def PTX_READ_LANEMASK_GT
  : PTX_READ_SPECIAL_REGISTER_R32<"lanemask_gt", int_ptx_read_lanemask_gt>;

def PTX_READ_CLOCK
  : PTX_READ_SPECIAL_REGISTER_R32<"clock", int_ptx_read_clock>;
def PTX_READ_CLOCK64
  : PTX_READ_SPECIAL_REGISTER_R64<"clock64", int_ptx_read_clock64>;

def PTX_READ_PM0 : PTX_READ_SPECIAL_REGISTER_R32<"pm0", int_ptx_read_pm0>;
def PTX_READ_PM1 : PTX_READ_SPECIAL_REGISTER_R32<"pm1", int_ptx_read_pm1>;
def PTX_READ_PM2 : PTX_READ_SPECIAL_REGISTER_R32<"pm2", int_ptx_read_pm2>;
def PTX_READ_PM3 : PTX_READ_SPECIAL_REGISTER_R32<"pm3", int_ptx_read_pm3>;

// PTX Parallel Synchronization and Communication Intrinsics

def PTX_BAR_SYNC : NVPTXInst<(outs), (ins i32imm:$i), "bar.sync\t$i;",
                             [(int_ptx_bar_sync imm:$i)]>;