//===- IntrinsicsNVVM.td - Defines NVVM intrinsics ---------*- tablegen -*-===// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// // // This file defines all of the NVVM-specific intrinsics for use with NVPTX. // //===----------------------------------------------------------------------===// // The following intrinsics were once defined here, but are now auto-upgraded // to target-generic LLVM intrinsics. // // * llvm.nvvm.brev32 --> llvm.bitreverse.i32 // * llvm.nvvm.brev64 --> llvm.bitreverse.i64 // * llvm.nvvm.clz.i --> llvm.ctlz.i32 // * llvm.nvvm.clz.ll --> trunc i64 llvm.ctlz.i64(x) to i32 // * llvm.nvvm.popc.i --> llvm.ctpop.i32 // * llvm.nvvm.popc.ll --> trunc i64 llvm.ctpop.i64 to i32 // * llvm.nvvm.abs.i --> select(x >= -x, x, -x) // * llvm.nvvm.abs.ll --> ibid. // * llvm.nvvm.max.i --> select(x sge y, x, y) // * llvm.nvvm.max.ll --> ibid. // * llvm.nvvm.max.ui --> select(x uge y, x, y) // * llvm.nvvm.max.ull --> ibid. // * llvm.nvvm.max.i --> select(x sle y, x, y) // * llvm.nvvm.max.ll --> ibid. // * llvm.nvvm.max.ui --> select(x ule y, x, y) // * llvm.nvvm.max.ull --> ibid. // * llvm.nvvm.h2f --> llvm.convert.to.fp16.f32 def llvm_anyi64ptr_ty : LLVMAnyPointerType; // (space)i64* // // MISC // // Helper class that represents a 'fragment' of an NVPTX *MMA instruction. // Geom: mnk. E.g. m8n32k16 // Frag: [abcd] // PtxEltType: PTX type for the element. class WMMA_REGS { string geom = Geom; string frag = Frag; string ptx_elt_type = PtxEltType; string gft = Geom#":"#Frag#":"#ptx_elt_type; string ft = frag#":"#ptx_elt_type; list regs = !cond( // mma.sync.m8n8k4 uses smaller a/b fragments than wmma fp ops !eq(gft,"m8n8k4:a:f16") : !listsplat(llvm_v2f16_ty, 2), !eq(gft,"m8n8k4:b:f16") : !listsplat(llvm_v2f16_ty, 2), // fp16 -> fp16/fp32 @ m16n16k16/m8n32k16/m32n8k16 // All currently supported geometries use the same fragment format, // so we only need to consider {fragment, type}. !eq(ft,"a:f16") : !listsplat(llvm_v2f16_ty, 8), !eq(ft,"b:f16") : !listsplat(llvm_v2f16_ty, 8), !eq(ft,"c:f16") : !listsplat(llvm_v2f16_ty, 4), !eq(ft,"d:f16") : !listsplat(llvm_v2f16_ty, 4), !eq(ft,"c:f32") : !listsplat(llvm_float_ty, 8), !eq(ft,"d:f32") : !listsplat(llvm_float_ty, 8), // u8/s8 -> s32 @ m16n16k16/m8n32k16/m32n8k16 !eq(gft,"m16n16k16:a:u8") : !listsplat(llvm_i32_ty, 2), !eq(gft,"m16n16k16:a:s8") : !listsplat(llvm_i32_ty, 2), !eq(gft,"m16n16k16:b:u8") : !listsplat(llvm_i32_ty, 2), !eq(gft,"m16n16k16:b:s8") : !listsplat(llvm_i32_ty, 2), !eq(gft,"m16n16k16:c:s32") : !listsplat(llvm_i32_ty, 8), !eq(gft,"m16n16k16:d:s32") : !listsplat(llvm_i32_ty, 8), !eq(gft,"m8n32k16:a:u8") : [llvm_i32_ty], !eq(gft,"m8n32k16:a:s8") : [llvm_i32_ty], !eq(gft,"m8n32k16:b:u8") : !listsplat(llvm_i32_ty, 4), !eq(gft,"m8n32k16:b:s8") : !listsplat(llvm_i32_ty, 4), !eq(gft,"m8n32k16:c:s32") : !listsplat(llvm_i32_ty, 8), !eq(gft,"m8n32k16:d:s32") : !listsplat(llvm_i32_ty, 8), !eq(gft,"m32n8k16:a:u8") : !listsplat(llvm_i32_ty, 4), !eq(gft,"m32n8k16:a:s8") : !listsplat(llvm_i32_ty, 4), !eq(gft,"m32n8k16:b:u8") : [llvm_i32_ty], !eq(gft,"m32n8k16:b:s8") : [llvm_i32_ty], !eq(gft,"m32n8k16:c:s32") : !listsplat(llvm_i32_ty, 8), !eq(gft,"m32n8k16:d:s32") : !listsplat(llvm_i32_ty, 8), // u4/s4/b1 -> s32 @ m8n8k32 (u4/s4), m8n8k128(b1) !eq(gft,"m8n8k128:a:b1") : [llvm_i32_ty], !eq(gft,"m8n8k32:a:u4") : [llvm_i32_ty], !eq(gft,"m8n8k32:a:s4") : [llvm_i32_ty], !eq(gft,"m8n8k128:b:b1") : [llvm_i32_ty], !eq(gft,"m8n8k32:b:u4") : [llvm_i32_ty], !eq(gft,"m8n8k32:b:s4") : [llvm_i32_ty], !eq(gft,"m8n8k128:c:s32") : !listsplat(llvm_i32_ty, 2), !eq(gft,"m8n8k128:d:s32") : !listsplat(llvm_i32_ty, 2), !eq(gft,"m8n8k32:c:s32") : !listsplat(llvm_i32_ty, 2), !eq(gft,"m8n8k32:d:s32") : !listsplat(llvm_i32_ty, 2), ); } class WMMA_NAME_LDST { string intr = "llvm.nvvm.wmma." # Frag.geom # "." # Op # "." # Frag.frag # "." # Layout # !if(WithStride, ".stride", "") # "." # Frag.ptx_elt_type ; // TODO(tra): record name should ideally use the same field order as the intrinsic. // E.g. string record = !subst("llvm", "int", // !subst(".", "_", llvm)); string record = "int_nvvm_wmma_" # Frag.geom # "_" # Op # "_" # Frag.frag # "_" # Frag.ptx_elt_type # "_" # Layout # !if(WithStride, "_stride", ""); } class MMA_SIGNATURE { list id_frags = !cond( // int and sub-int ops are identified by input type. !eq(A.ptx_elt_type, "s8") : [A], !eq(A.ptx_elt_type, "u8") : [A], !eq(A.ptx_elt_type, "s4") : [A], !eq(A.ptx_elt_type, "u4") : [A], !eq(A.ptx_elt_type, "b1") : [A], // the rest are FP ops identified by accumulator & result type. true: [D, C] ); string ret = !foldl("", id_frags, a, b, !strconcat(a, ".", b.ptx_elt_type)); } class WMMA_NAME_MMA { string signature = MMA_SIGNATURE.ret; string llvm = !if( !eq(A.geom, "m8n8k4"), "llvm.nvvm.mma.m8n8k4" # "." # ALayout # "." # BLayout # signature, "llvm.nvvm.wmma." # A.geom # ".mma" # "." # ALayout # "." # BLayout # signature # !if(Satfinite, ".satfinite", "")); string record = !subst(".", "_", !subst("llvm.", "int_", llvm)); } // Generates list of 4-tuples of WMMA_REGS representing a valid MMA op. // Geom: list of supported geometries. // TypeN: PTX type of the corresponding fragment's element. // TypeB and TypeD may be empty if it must match that of TypeA or TypeC. class MMA_OPS Geom, list TypeA, list TypeB, list TypeC, list TypeD> { list> ret = !foldl([]>, Geom, t1, geom, !listconcat(t1, !foldl([]>, TypeA, t2, type_a, !listconcat(t2, !foldl([]>, !if(!size(TypeB), TypeB, [type_a]), t3, type_b, !listconcat(t3, !foldl([]>, TypeC, t4, type_c, !listconcat(t4, !foldl([]>, !if(!size(TypeD), TypeD, [type_c]), t5, type_d, !listconcat(t5, [[WMMA_REGS, WMMA_REGS, WMMA_REGS, WMMA_REGS]])))))))))); // Debugging aid for readable representation of the list above. list> ops = !foreach(x, ret, [x[0].gft, x[1].gft, x[2].gft, x[3].gft]); } class MMA_LDST_OPS Geom, list Frags, list Types> { list ret = !foldl([], Geom, t1, geom, !listconcat(t1, !foldl([], Frags, t2, frag, !listconcat(t2, !foldl([], Types, t3, type, !listconcat(t3, [WMMA_REGS])))))); // Debugging aid for readable representation of the list above. list ops = !foreach(x, ret, x.gft); } // Creates list of valid combinations of fragments. This is the master list that // drives generation of corresponding intrinsics and instructions. class NVVM_MMA_OPS { list> fp_mma_ops = MMA_OPS< ["m8n8k4"], ["f16"], [], ["f16", "f32"], ["f16", "f32"]>.ret; list> fp_wmma_ops = MMA_OPS< ["m16n16k16", "m32n8k16", "m8n32k16"], ["f16"], [], ["f16", "f32"], ["f16", "f32"]>.ret; list> int_wmma_ops = MMA_OPS< ["m16n16k16", "m32n8k16", "m8n32k16"], ["s8", "u8"], [], ["s32"], []>.ret; list> subint_wmma_ops = MMA_OPS< ["m8n8k32"], ["s4", "u4"], [], ["s32"], []>.ret; list> bit_wmma_ops = MMA_OPS< ["m8n8k128"], ["b1"], [], ["s32"], []>.ret; list> all_mma_ops = !listconcat( fp_mma_ops, fp_wmma_ops, int_wmma_ops, subint_wmma_ops, bit_wmma_ops); list ldst_ab_ops = MMA_LDST_OPS< ["m16n16k16", "m32n8k16", "m8n32k16"], ["a", "b"], ["f16", "u8", "s8"]>.ret; list ldst_cd_ops = MMA_LDST_OPS< ["m16n16k16", "m32n8k16", "m8n32k16"], ["c", "d"], ["f16", "f32", "s32"]>.ret; list ldst_subint_ab_ops = MMA_LDST_OPS< ["m8n8k32"], ["a", "b"], ["s4","u4"]>.ret; list ldst_bit_ab_ops = MMA_LDST_OPS< ["m8n8k128"], ["a", "b"], ["b1"]>.ret; list ldst_subint_cd_ops = MMA_LDST_OPS< ["m8n8k32", "m8n8k128"], ["c", "d"], ["s32"]>.ret; list all_ldst_ops = !listconcat(ldst_ab_ops, ldst_cd_ops, ldst_subint_ab_ops, ldst_bit_ab_ops, ldst_subint_cd_ops); // Separate A/B/C fragments (loads) from D (stores). list all_ld_ops = !filter(op, all_ldst_ops, !ne(op.frag, "d")); list all_st_ops = !filter(op, all_ldst_ops, !eq(op.frag, "d")); } def NVVM_MMA_OPS : NVVM_MMA_OPS; // Returns true if this combination of layout/satf is supported; false otherwise. // MMA ops must provide all parameters. Loads and stores -- only frags and layout_a. // The class is used to prevent generation of records for the unsupported variants. // E.g. // if NVVM_MMA_SUPPORTED<...>.ret then // def : FOO<>; // The record will only be defined for supported ops. // class NVVM_MMA_SUPPORTED frags, string layout_a, string layout_b="-", int satf=-1> { // MMA ops check both layouts. string mma = frags[0].ptx_elt_type # ":" # layout_a # ":" # layout_b; // Load ops only need type/fragment/layout. string ld = frags[0].ptx_elt_type # ":" # frags[0].frag # ":" # layout_a ; string ldf = frags[0].ptx_elt_type # ":" # frags[0].frag ; string t = frags[0].ptx_elt_type; // gcd is a shortcut used to identify instructions that depend on // geom+frag_c+frag_d. Not all instances of this class have all fragments // specified. If there are not enough fragments, the tail evaluates to '?'. string gcd = frags[0].geom # ":" # !if(!eq(!size(frags), 4), frags[2].ptx_elt_type # frags[3].ptx_elt_type, "?"); bit ret = !cond( // Sub-int MMA only supports fixed A/B layout. // b1 does not support .satf. !eq(mma#":"#satf, "b1:row:col:0") : true, // mma.m8n8k4 has no .satf modifier. !and(!eq(frags[0].geom, "m8n8k4"), !ne(satf, 0)): false, // mma.m8n8k4 has no C=f32 D=f16 variant. !eq(gcd, "m8n8k4:f32f16"): false, !eq(mma, "s4:row:col") : true, !eq(mma, "u4:row:col") : true, !eq(mma, "s4:row:col") : true, !eq(mma, "u4:row:col") : true, // Sub-int load/stores have fixed layout for A and B. !and(!eq(layout_b, "-"), // It's a Load or Store op !or(!eq(ld, "b1:a:row"), !eq(ld, "b1:b:col"), !eq(ldf, "b1:c"), !eq(ldf, "b1:d"), !eq(ld, "s4:a:row"), !eq(ld, "s4:b:col"), !eq(ldf, "s4:c"), !eq(ldf, "s4:d"), !eq(ld, "u4:a:row"), !eq(ld, "u4:b:col"), !eq(ldf, "u4:c"), !eq(ldf, "u4:d"))) : true, // All other sub-int ops are not supported. !eq(t, "b1") : false, !eq(t, "s4") : false, !eq(t, "u4") : false, // All other (non sub-int) are OK. true: true ); } class SHFL_INFO { string Suffix = !if(sync, "sync_", "") # mode # "_" # type # !if(return_pred, "p", ""); string Name = "int_nvvm_shfl_" # Suffix; string Builtin = "__nvvm_shfl_" # Suffix; string IntrName = "llvm.nvvm.shfl." # !subst("_",".", Suffix); bit withGccBuiltin = !not(return_pred); bit withoutGccBuiltin = return_pred; LLVMType OpType = !cond( !eq(type,"i32"): llvm_i32_ty, !eq(type,"f32"): llvm_float_ty); list RetTy = !if(return_pred, [OpType, llvm_i1_ty], [OpType]); list ArgsTy = !if(sync, [llvm_i32_ty, OpType, llvm_i32_ty, llvm_i32_ty], [OpType, llvm_i32_ty, llvm_i32_ty]); } let TargetPrefix = "nvvm" in { def int_nvvm_prmt : GCCBuiltin<"__nvvm_prmt">, Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem, Commutative]>; // // Min Max // def int_nvvm_fmin_f : GCCBuiltin<"__nvvm_fmin_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_fmin_ftz_f : GCCBuiltin<"__nvvm_fmin_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_fmax_f : GCCBuiltin<"__nvvm_fmax_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty] , [IntrNoMem, Commutative]>; def int_nvvm_fmax_ftz_f : GCCBuiltin<"__nvvm_fmax_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_fmin_d : GCCBuiltin<"__nvvm_fmin_d">, Intrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty], [IntrNoMem, Commutative]>; def int_nvvm_fmax_d : GCCBuiltin<"__nvvm_fmax_d">, Intrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty], [IntrNoMem, Commutative]>; // // Multiplication // def int_nvvm_mulhi_i : GCCBuiltin<"__nvvm_mulhi_i">, Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, Commutative]>; def int_nvvm_mulhi_ui : GCCBuiltin<"__nvvm_mulhi_ui">, Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, Commutative]>; def int_nvvm_mulhi_ll : GCCBuiltin<"__nvvm_mulhi_ll">, Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i64_ty], [IntrNoMem, Commutative]>; def int_nvvm_mulhi_ull : GCCBuiltin<"__nvvm_mulhi_ull">, Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i64_ty], [IntrNoMem, Commutative]>; def int_nvvm_mul_rn_ftz_f : GCCBuiltin<"__nvvm_mul_rn_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_mul_rn_f : GCCBuiltin<"__nvvm_mul_rn_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_mul_rz_ftz_f : GCCBuiltin<"__nvvm_mul_rz_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_mul_rz_f : GCCBuiltin<"__nvvm_mul_rz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_mul_rm_ftz_f : GCCBuiltin<"__nvvm_mul_rm_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_mul_rm_f : GCCBuiltin<"__nvvm_mul_rm_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_mul_rp_ftz_f : GCCBuiltin<"__nvvm_mul_rp_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_mul_rp_f : GCCBuiltin<"__nvvm_mul_rp_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_mul_rn_d : GCCBuiltin<"__nvvm_mul_rn_d">, Intrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty], [IntrNoMem, Commutative]>; def int_nvvm_mul_rz_d : GCCBuiltin<"__nvvm_mul_rz_d">, Intrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty], [IntrNoMem, Commutative]>; def int_nvvm_mul_rm_d : GCCBuiltin<"__nvvm_mul_rm_d">, Intrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty], [IntrNoMem, Commutative]>; def int_nvvm_mul_rp_d : GCCBuiltin<"__nvvm_mul_rp_d">, Intrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty], [IntrNoMem, Commutative]>; def int_nvvm_mul24_i : GCCBuiltin<"__nvvm_mul24_i">, Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, Commutative]>; def int_nvvm_mul24_ui : GCCBuiltin<"__nvvm_mul24_ui">, Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, Commutative]>; // // Div // def int_nvvm_div_approx_ftz_f : GCCBuiltin<"__nvvm_div_approx_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_div_approx_f : GCCBuiltin<"__nvvm_div_approx_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_div_rn_ftz_f : GCCBuiltin<"__nvvm_div_rn_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_div_rn_f : GCCBuiltin<"__nvvm_div_rn_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_div_rz_ftz_f : GCCBuiltin<"__nvvm_div_rz_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_div_rz_f : GCCBuiltin<"__nvvm_div_rz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_div_rm_ftz_f : GCCBuiltin<"__nvvm_div_rm_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_div_rm_f : GCCBuiltin<"__nvvm_div_rm_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_div_rp_ftz_f : GCCBuiltin<"__nvvm_div_rp_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_div_rp_f : GCCBuiltin<"__nvvm_div_rp_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_div_rn_d : GCCBuiltin<"__nvvm_div_rn_d">, Intrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty], [IntrNoMem, Commutative]>; def int_nvvm_div_rz_d : GCCBuiltin<"__nvvm_div_rz_d">, Intrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty], [IntrNoMem, Commutative]>; def int_nvvm_div_rm_d : GCCBuiltin<"__nvvm_div_rm_d">, Intrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty], [IntrNoMem, Commutative]>; def int_nvvm_div_rp_d : GCCBuiltin<"__nvvm_div_rp_d">, Intrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty], [IntrNoMem, Commutative]>; // // Sad // def int_nvvm_sad_i : GCCBuiltin<"__nvvm_sad_i">, Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem, Commutative]>; def int_nvvm_sad_ui : GCCBuiltin<"__nvvm_sad_ui">, Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem, Commutative]>; // // Floor Ceil // def int_nvvm_floor_ftz_f : GCCBuiltin<"__nvvm_floor_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_floor_f : GCCBuiltin<"__nvvm_floor_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_floor_d : GCCBuiltin<"__nvvm_floor_d">, Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem]>; def int_nvvm_ceil_ftz_f : GCCBuiltin<"__nvvm_ceil_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_ceil_f : GCCBuiltin<"__nvvm_ceil_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_ceil_d : GCCBuiltin<"__nvvm_ceil_d">, Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem]>; // // Abs // def int_nvvm_fabs_ftz_f : GCCBuiltin<"__nvvm_fabs_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_fabs_f : GCCBuiltin<"__nvvm_fabs_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_fabs_d : GCCBuiltin<"__nvvm_fabs_d">, Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem]>; // // Round // def int_nvvm_round_ftz_f : GCCBuiltin<"__nvvm_round_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_round_f : GCCBuiltin<"__nvvm_round_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_round_d : GCCBuiltin<"__nvvm_round_d">, Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem]>; // // Trunc // def int_nvvm_trunc_ftz_f : GCCBuiltin<"__nvvm_trunc_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_trunc_f : GCCBuiltin<"__nvvm_trunc_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_trunc_d : GCCBuiltin<"__nvvm_trunc_d">, Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem]>; // // Saturate // def int_nvvm_saturate_ftz_f : GCCBuiltin<"__nvvm_saturate_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_saturate_f : GCCBuiltin<"__nvvm_saturate_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_saturate_d : GCCBuiltin<"__nvvm_saturate_d">, Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem]>; // // Exp2 Log2 // def int_nvvm_ex2_approx_ftz_f : GCCBuiltin<"__nvvm_ex2_approx_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_ex2_approx_f : GCCBuiltin<"__nvvm_ex2_approx_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_ex2_approx_d : GCCBuiltin<"__nvvm_ex2_approx_d">, Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem]>; def int_nvvm_lg2_approx_ftz_f : GCCBuiltin<"__nvvm_lg2_approx_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_lg2_approx_f : GCCBuiltin<"__nvvm_lg2_approx_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_lg2_approx_d : GCCBuiltin<"__nvvm_lg2_approx_d">, Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem]>; // // Sin Cos // def int_nvvm_sin_approx_ftz_f : GCCBuiltin<"__nvvm_sin_approx_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_sin_approx_f : GCCBuiltin<"__nvvm_sin_approx_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_cos_approx_ftz_f : GCCBuiltin<"__nvvm_cos_approx_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_cos_approx_f : GCCBuiltin<"__nvvm_cos_approx_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; // // Fma // def int_nvvm_fma_rn_ftz_f : GCCBuiltin<"__nvvm_fma_rn_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_fma_rn_f : GCCBuiltin<"__nvvm_fma_rn_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_fma_rz_ftz_f : GCCBuiltin<"__nvvm_fma_rz_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_fma_rz_f : GCCBuiltin<"__nvvm_fma_rz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_fma_rm_ftz_f : GCCBuiltin<"__nvvm_fma_rm_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_fma_rm_f : GCCBuiltin<"__nvvm_fma_rm_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_fma_rp_ftz_f : GCCBuiltin<"__nvvm_fma_rp_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_fma_rp_f : GCCBuiltin<"__nvvm_fma_rp_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_fma_rn_d : GCCBuiltin<"__nvvm_fma_rn_d">, Intrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty, llvm_double_ty], [IntrNoMem, Commutative]>; def int_nvvm_fma_rz_d : GCCBuiltin<"__nvvm_fma_rz_d">, Intrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty, llvm_double_ty], [IntrNoMem, Commutative]>; def int_nvvm_fma_rm_d : GCCBuiltin<"__nvvm_fma_rm_d">, Intrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty, llvm_double_ty], [IntrNoMem, Commutative]>; def int_nvvm_fma_rp_d : GCCBuiltin<"__nvvm_fma_rp_d">, Intrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty, llvm_double_ty], [IntrNoMem, Commutative]>; // // Rcp // def int_nvvm_rcp_rn_ftz_f : GCCBuiltin<"__nvvm_rcp_rn_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_rcp_rn_f : GCCBuiltin<"__nvvm_rcp_rn_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_rcp_rz_ftz_f : GCCBuiltin<"__nvvm_rcp_rz_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_rcp_rz_f : GCCBuiltin<"__nvvm_rcp_rz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_rcp_rm_ftz_f : GCCBuiltin<"__nvvm_rcp_rm_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_rcp_rm_f : GCCBuiltin<"__nvvm_rcp_rm_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_rcp_rp_ftz_f : GCCBuiltin<"__nvvm_rcp_rp_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_rcp_rp_f : GCCBuiltin<"__nvvm_rcp_rp_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_rcp_rn_d : GCCBuiltin<"__nvvm_rcp_rn_d">, Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem]>; def int_nvvm_rcp_rz_d : GCCBuiltin<"__nvvm_rcp_rz_d">, Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem]>; def int_nvvm_rcp_rm_d : GCCBuiltin<"__nvvm_rcp_rm_d">, Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem]>; def int_nvvm_rcp_rp_d : GCCBuiltin<"__nvvm_rcp_rp_d">, Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem]>; def int_nvvm_rcp_approx_ftz_d : GCCBuiltin<"__nvvm_rcp_approx_ftz_d">, Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem]>; // // Sqrt // def int_nvvm_sqrt_f : GCCBuiltin<"__nvvm_sqrt_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_sqrt_rn_ftz_f : GCCBuiltin<"__nvvm_sqrt_rn_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_sqrt_rn_f : GCCBuiltin<"__nvvm_sqrt_rn_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_sqrt_rz_ftz_f : GCCBuiltin<"__nvvm_sqrt_rz_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_sqrt_rz_f : GCCBuiltin<"__nvvm_sqrt_rz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_sqrt_rm_ftz_f : GCCBuiltin<"__nvvm_sqrt_rm_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_sqrt_rm_f : GCCBuiltin<"__nvvm_sqrt_rm_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_sqrt_rp_ftz_f : GCCBuiltin<"__nvvm_sqrt_rp_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_sqrt_rp_f : GCCBuiltin<"__nvvm_sqrt_rp_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_sqrt_approx_ftz_f : GCCBuiltin<"__nvvm_sqrt_approx_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_sqrt_approx_f : GCCBuiltin<"__nvvm_sqrt_approx_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_sqrt_rn_d : GCCBuiltin<"__nvvm_sqrt_rn_d">, Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem]>; def int_nvvm_sqrt_rz_d : GCCBuiltin<"__nvvm_sqrt_rz_d">, Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem]>; def int_nvvm_sqrt_rm_d : GCCBuiltin<"__nvvm_sqrt_rm_d">, Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem]>; def int_nvvm_sqrt_rp_d : GCCBuiltin<"__nvvm_sqrt_rp_d">, Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem]>; // // Rsqrt // def int_nvvm_rsqrt_approx_ftz_f : GCCBuiltin<"__nvvm_rsqrt_approx_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_rsqrt_approx_f : GCCBuiltin<"__nvvm_rsqrt_approx_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_rsqrt_approx_d : GCCBuiltin<"__nvvm_rsqrt_approx_d">, Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem]>; // // Add // def int_nvvm_add_rn_ftz_f : GCCBuiltin<"__nvvm_add_rn_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_add_rn_f : GCCBuiltin<"__nvvm_add_rn_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_add_rz_ftz_f : GCCBuiltin<"__nvvm_add_rz_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_add_rz_f : GCCBuiltin<"__nvvm_add_rz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_add_rm_ftz_f : GCCBuiltin<"__nvvm_add_rm_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_add_rm_f : GCCBuiltin<"__nvvm_add_rm_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_add_rp_ftz_f : GCCBuiltin<"__nvvm_add_rp_ftz_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_add_rp_f : GCCBuiltin<"__nvvm_add_rp_f">, Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty], [IntrNoMem, Commutative]>; def int_nvvm_add_rn_d : GCCBuiltin<"__nvvm_add_rn_d">, Intrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty], [IntrNoMem, Commutative]>; def int_nvvm_add_rz_d : GCCBuiltin<"__nvvm_add_rz_d">, Intrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty], [IntrNoMem, Commutative]>; def int_nvvm_add_rm_d : GCCBuiltin<"__nvvm_add_rm_d">, Intrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty], [IntrNoMem, Commutative]>; def int_nvvm_add_rp_d : GCCBuiltin<"__nvvm_add_rp_d">, Intrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty], [IntrNoMem, Commutative]>; // // Convert // def int_nvvm_d2f_rn_ftz : GCCBuiltin<"__nvvm_d2f_rn_ftz">, Intrinsic<[llvm_float_ty], [llvm_double_ty], [IntrNoMem]>; def int_nvvm_d2f_rn : GCCBuiltin<"__nvvm_d2f_rn">, Intrinsic<[llvm_float_ty], [llvm_double_ty], [IntrNoMem]>; def int_nvvm_d2f_rz_ftz : GCCBuiltin<"__nvvm_d2f_rz_ftz">, Intrinsic<[llvm_float_ty], [llvm_double_ty], [IntrNoMem]>; def int_nvvm_d2f_rz : GCCBuiltin<"__nvvm_d2f_rz">, Intrinsic<[llvm_float_ty], [llvm_double_ty], [IntrNoMem]>; def int_nvvm_d2f_rm_ftz : GCCBuiltin<"__nvvm_d2f_rm_ftz">, Intrinsic<[llvm_float_ty], [llvm_double_ty], [IntrNoMem]>; def int_nvvm_d2f_rm : GCCBuiltin<"__nvvm_d2f_rm">, Intrinsic<[llvm_float_ty], [llvm_double_ty], [IntrNoMem]>; def int_nvvm_d2f_rp_ftz : GCCBuiltin<"__nvvm_d2f_rp_ftz">, Intrinsic<[llvm_float_ty], [llvm_double_ty], [IntrNoMem]>; def int_nvvm_d2f_rp : GCCBuiltin<"__nvvm_d2f_rp">, Intrinsic<[llvm_float_ty], [llvm_double_ty], [IntrNoMem]>; def int_nvvm_d2i_rn : GCCBuiltin<"__nvvm_d2i_rn">, Intrinsic<[llvm_i32_ty], [llvm_double_ty], [IntrNoMem]>; def int_nvvm_d2i_rz : GCCBuiltin<"__nvvm_d2i_rz">, Intrinsic<[llvm_i32_ty], [llvm_double_ty], [IntrNoMem]>; def int_nvvm_d2i_rm : GCCBuiltin<"__nvvm_d2i_rm">, Intrinsic<[llvm_i32_ty], [llvm_double_ty], [IntrNoMem]>; def int_nvvm_d2i_rp : GCCBuiltin<"__nvvm_d2i_rp">, Intrinsic<[llvm_i32_ty], [llvm_double_ty], [IntrNoMem]>; def int_nvvm_d2ui_rn : GCCBuiltin<"__nvvm_d2ui_rn">, Intrinsic<[llvm_i32_ty], [llvm_double_ty], [IntrNoMem]>; def int_nvvm_d2ui_rz : GCCBuiltin<"__nvvm_d2ui_rz">, Intrinsic<[llvm_i32_ty], [llvm_double_ty], [IntrNoMem]>; def int_nvvm_d2ui_rm : GCCBuiltin<"__nvvm_d2ui_rm">, Intrinsic<[llvm_i32_ty], [llvm_double_ty], [IntrNoMem]>; def int_nvvm_d2ui_rp : GCCBuiltin<"__nvvm_d2ui_rp">, Intrinsic<[llvm_i32_ty], [llvm_double_ty], [IntrNoMem]>; def int_nvvm_i2d_rn : GCCBuiltin<"__nvvm_i2d_rn">, Intrinsic<[llvm_double_ty], [llvm_i32_ty], [IntrNoMem]>; def int_nvvm_i2d_rz : GCCBuiltin<"__nvvm_i2d_rz">, Intrinsic<[llvm_double_ty], [llvm_i32_ty], [IntrNoMem]>; def int_nvvm_i2d_rm : GCCBuiltin<"__nvvm_i2d_rm">, Intrinsic<[llvm_double_ty], [llvm_i32_ty], [IntrNoMem]>; def int_nvvm_i2d_rp : GCCBuiltin<"__nvvm_i2d_rp">, Intrinsic<[llvm_double_ty], [llvm_i32_ty], [IntrNoMem]>; def int_nvvm_ui2d_rn : GCCBuiltin<"__nvvm_ui2d_rn">, Intrinsic<[llvm_double_ty], [llvm_i32_ty], [IntrNoMem]>; def int_nvvm_ui2d_rz : GCCBuiltin<"__nvvm_ui2d_rz">, Intrinsic<[llvm_double_ty], [llvm_i32_ty], [IntrNoMem]>; def int_nvvm_ui2d_rm : GCCBuiltin<"__nvvm_ui2d_rm">, Intrinsic<[llvm_double_ty], [llvm_i32_ty], [IntrNoMem]>; def int_nvvm_ui2d_rp : GCCBuiltin<"__nvvm_ui2d_rp">, Intrinsic<[llvm_double_ty], [llvm_i32_ty], [IntrNoMem]>; def int_nvvm_f2i_rn_ftz : GCCBuiltin<"__nvvm_f2i_rn_ftz">, Intrinsic<[llvm_i32_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_f2i_rn : GCCBuiltin<"__nvvm_f2i_rn">, Intrinsic<[llvm_i32_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_f2i_rz_ftz : GCCBuiltin<"__nvvm_f2i_rz_ftz">, Intrinsic<[llvm_i32_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_f2i_rz : GCCBuiltin<"__nvvm_f2i_rz">, Intrinsic<[llvm_i32_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_f2i_rm_ftz : GCCBuiltin<"__nvvm_f2i_rm_ftz">, Intrinsic<[llvm_i32_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_f2i_rm : GCCBuiltin<"__nvvm_f2i_rm">, Intrinsic<[llvm_i32_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_f2i_rp_ftz : GCCBuiltin<"__nvvm_f2i_rp_ftz">, Intrinsic<[llvm_i32_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_f2i_rp : GCCBuiltin<"__nvvm_f2i_rp">, Intrinsic<[llvm_i32_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_f2ui_rn_ftz : GCCBuiltin<"__nvvm_f2ui_rn_ftz">, Intrinsic<[llvm_i32_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_f2ui_rn : GCCBuiltin<"__nvvm_f2ui_rn">, Intrinsic<[llvm_i32_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_f2ui_rz_ftz : GCCBuiltin<"__nvvm_f2ui_rz_ftz">, Intrinsic<[llvm_i32_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_f2ui_rz : GCCBuiltin<"__nvvm_f2ui_rz">, Intrinsic<[llvm_i32_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_f2ui_rm_ftz : GCCBuiltin<"__nvvm_f2ui_rm_ftz">, Intrinsic<[llvm_i32_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_f2ui_rm : GCCBuiltin<"__nvvm_f2ui_rm">, Intrinsic<[llvm_i32_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_f2ui_rp_ftz : GCCBuiltin<"__nvvm_f2ui_rp_ftz">, Intrinsic<[llvm_i32_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_f2ui_rp : GCCBuiltin<"__nvvm_f2ui_rp">, Intrinsic<[llvm_i32_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_i2f_rn : GCCBuiltin<"__nvvm_i2f_rn">, Intrinsic<[llvm_float_ty], [llvm_i32_ty], [IntrNoMem]>; def int_nvvm_i2f_rz : GCCBuiltin<"__nvvm_i2f_rz">, Intrinsic<[llvm_float_ty], [llvm_i32_ty], [IntrNoMem]>; def int_nvvm_i2f_rm : GCCBuiltin<"__nvvm_i2f_rm">, Intrinsic<[llvm_float_ty], [llvm_i32_ty], [IntrNoMem]>; def int_nvvm_i2f_rp : GCCBuiltin<"__nvvm_i2f_rp">, Intrinsic<[llvm_float_ty], [llvm_i32_ty], [IntrNoMem]>; def int_nvvm_ui2f_rn : GCCBuiltin<"__nvvm_ui2f_rn">, Intrinsic<[llvm_float_ty], [llvm_i32_ty], [IntrNoMem]>; def int_nvvm_ui2f_rz : GCCBuiltin<"__nvvm_ui2f_rz">, Intrinsic<[llvm_float_ty], [llvm_i32_ty], [IntrNoMem]>; def int_nvvm_ui2f_rm : GCCBuiltin<"__nvvm_ui2f_rm">, Intrinsic<[llvm_float_ty], [llvm_i32_ty], [IntrNoMem]>; def int_nvvm_ui2f_rp : GCCBuiltin<"__nvvm_ui2f_rp">, Intrinsic<[llvm_float_ty], [llvm_i32_ty], [IntrNoMem]>; def int_nvvm_lohi_i2d : GCCBuiltin<"__nvvm_lohi_i2d">, Intrinsic<[llvm_double_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, Commutative]>; def int_nvvm_d2i_lo : GCCBuiltin<"__nvvm_d2i_lo">, Intrinsic<[llvm_i32_ty], [llvm_double_ty], [IntrNoMem]>; def int_nvvm_d2i_hi : GCCBuiltin<"__nvvm_d2i_hi">, Intrinsic<[llvm_i32_ty], [llvm_double_ty], [IntrNoMem]>; def int_nvvm_f2ll_rn_ftz : GCCBuiltin<"__nvvm_f2ll_rn_ftz">, Intrinsic<[llvm_i64_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_f2ll_rn : GCCBuiltin<"__nvvm_f2ll_rn">, Intrinsic<[llvm_i64_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_f2ll_rz_ftz : GCCBuiltin<"__nvvm_f2ll_rz_ftz">, Intrinsic<[llvm_i64_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_f2ll_rz : GCCBuiltin<"__nvvm_f2ll_rz">, Intrinsic<[llvm_i64_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_f2ll_rm_ftz : GCCBuiltin<"__nvvm_f2ll_rm_ftz">, Intrinsic<[llvm_i64_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_f2ll_rm : GCCBuiltin<"__nvvm_f2ll_rm">, Intrinsic<[llvm_i64_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_f2ll_rp_ftz : GCCBuiltin<"__nvvm_f2ll_rp_ftz">, Intrinsic<[llvm_i64_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_f2ll_rp : GCCBuiltin<"__nvvm_f2ll_rp">, Intrinsic<[llvm_i64_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_f2ull_rn_ftz : GCCBuiltin<"__nvvm_f2ull_rn_ftz">, Intrinsic<[llvm_i64_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_f2ull_rn : GCCBuiltin<"__nvvm_f2ull_rn">, Intrinsic<[llvm_i64_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_f2ull_rz_ftz : GCCBuiltin<"__nvvm_f2ull_rz_ftz">, Intrinsic<[llvm_i64_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_f2ull_rz : GCCBuiltin<"__nvvm_f2ull_rz">, Intrinsic<[llvm_i64_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_f2ull_rm_ftz : GCCBuiltin<"__nvvm_f2ull_rm_ftz">, Intrinsic<[llvm_i64_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_f2ull_rm : GCCBuiltin<"__nvvm_f2ull_rm">, Intrinsic<[llvm_i64_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_f2ull_rp_ftz : GCCBuiltin<"__nvvm_f2ull_rp_ftz">, Intrinsic<[llvm_i64_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_f2ull_rp : GCCBuiltin<"__nvvm_f2ull_rp">, Intrinsic<[llvm_i64_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_d2ll_rn : GCCBuiltin<"__nvvm_d2ll_rn">, Intrinsic<[llvm_i64_ty], [llvm_double_ty], [IntrNoMem]>; def int_nvvm_d2ll_rz : GCCBuiltin<"__nvvm_d2ll_rz">, Intrinsic<[llvm_i64_ty], [llvm_double_ty], [IntrNoMem]>; def int_nvvm_d2ll_rm : GCCBuiltin<"__nvvm_d2ll_rm">, Intrinsic<[llvm_i64_ty], [llvm_double_ty], [IntrNoMem]>; def int_nvvm_d2ll_rp : GCCBuiltin<"__nvvm_d2ll_rp">, Intrinsic<[llvm_i64_ty], [llvm_double_ty], [IntrNoMem]>; def int_nvvm_d2ull_rn : GCCBuiltin<"__nvvm_d2ull_rn">, Intrinsic<[llvm_i64_ty], [llvm_double_ty], [IntrNoMem]>; def int_nvvm_d2ull_rz : GCCBuiltin<"__nvvm_d2ull_rz">, Intrinsic<[llvm_i64_ty], [llvm_double_ty], [IntrNoMem]>; def int_nvvm_d2ull_rm : GCCBuiltin<"__nvvm_d2ull_rm">, Intrinsic<[llvm_i64_ty], [llvm_double_ty], [IntrNoMem]>; def int_nvvm_d2ull_rp : GCCBuiltin<"__nvvm_d2ull_rp">, Intrinsic<[llvm_i64_ty], [llvm_double_ty], [IntrNoMem]>; def int_nvvm_ll2f_rn : GCCBuiltin<"__nvvm_ll2f_rn">, Intrinsic<[llvm_float_ty], [llvm_i64_ty], [IntrNoMem]>; def int_nvvm_ll2f_rz : GCCBuiltin<"__nvvm_ll2f_rz">, Intrinsic<[llvm_float_ty], [llvm_i64_ty], [IntrNoMem]>; def int_nvvm_ll2f_rm : GCCBuiltin<"__nvvm_ll2f_rm">, Intrinsic<[llvm_float_ty], [llvm_i64_ty], [IntrNoMem]>; def int_nvvm_ll2f_rp : GCCBuiltin<"__nvvm_ll2f_rp">, Intrinsic<[llvm_float_ty], [llvm_i64_ty], [IntrNoMem]>; def int_nvvm_ull2f_rn : GCCBuiltin<"__nvvm_ull2f_rn">, Intrinsic<[llvm_float_ty], [llvm_i64_ty], [IntrNoMem]>; def int_nvvm_ull2f_rz : GCCBuiltin<"__nvvm_ull2f_rz">, Intrinsic<[llvm_float_ty], [llvm_i64_ty], [IntrNoMem]>; def int_nvvm_ull2f_rm : GCCBuiltin<"__nvvm_ull2f_rm">, Intrinsic<[llvm_float_ty], [llvm_i64_ty], [IntrNoMem]>; def int_nvvm_ull2f_rp : GCCBuiltin<"__nvvm_ull2f_rp">, Intrinsic<[llvm_float_ty], [llvm_i64_ty], [IntrNoMem]>; def int_nvvm_ll2d_rn : GCCBuiltin<"__nvvm_ll2d_rn">, Intrinsic<[llvm_double_ty], [llvm_i64_ty], [IntrNoMem]>; def int_nvvm_ll2d_rz : GCCBuiltin<"__nvvm_ll2d_rz">, Intrinsic<[llvm_double_ty], [llvm_i64_ty], [IntrNoMem]>; def int_nvvm_ll2d_rm : GCCBuiltin<"__nvvm_ll2d_rm">, Intrinsic<[llvm_double_ty], [llvm_i64_ty], [IntrNoMem]>; def int_nvvm_ll2d_rp : GCCBuiltin<"__nvvm_ll2d_rp">, Intrinsic<[llvm_double_ty], [llvm_i64_ty], [IntrNoMem]>; def int_nvvm_ull2d_rn : GCCBuiltin<"__nvvm_ull2d_rn">, Intrinsic<[llvm_double_ty], [llvm_i64_ty], [IntrNoMem]>; def int_nvvm_ull2d_rz : GCCBuiltin<"__nvvm_ull2d_rz">, Intrinsic<[llvm_double_ty], [llvm_i64_ty], [IntrNoMem]>; def int_nvvm_ull2d_rm : GCCBuiltin<"__nvvm_ull2d_rm">, Intrinsic<[llvm_double_ty], [llvm_i64_ty], [IntrNoMem]>; def int_nvvm_ull2d_rp : GCCBuiltin<"__nvvm_ull2d_rp">, Intrinsic<[llvm_double_ty], [llvm_i64_ty], [IntrNoMem]>; def int_nvvm_f2h_rn_ftz : GCCBuiltin<"__nvvm_f2h_rn_ftz">, Intrinsic<[llvm_i16_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_f2h_rn : GCCBuiltin<"__nvvm_f2h_rn">, Intrinsic<[llvm_i16_ty], [llvm_float_ty], [IntrNoMem]>; // // Bitcast // def int_nvvm_bitcast_f2i : GCCBuiltin<"__nvvm_bitcast_f2i">, Intrinsic<[llvm_i32_ty], [llvm_float_ty], [IntrNoMem]>; def int_nvvm_bitcast_i2f : GCCBuiltin<"__nvvm_bitcast_i2f">, Intrinsic<[llvm_float_ty], [llvm_i32_ty], [IntrNoMem]>; def int_nvvm_bitcast_ll2d : GCCBuiltin<"__nvvm_bitcast_ll2d">, Intrinsic<[llvm_double_ty], [llvm_i64_ty], [IntrNoMem]>; def int_nvvm_bitcast_d2ll : GCCBuiltin<"__nvvm_bitcast_d2ll">, Intrinsic<[llvm_i64_ty], [llvm_double_ty], [IntrNoMem]>; // FNS def int_nvvm_fns : GCCBuiltin<"__nvvm_fns">, Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>; // Atomics not available as llvm intrinsics. def int_nvvm_atomic_load_inc_32 : Intrinsic<[llvm_i32_ty], [LLVMAnyPointerType, llvm_i32_ty], [IntrArgMemOnly, NoCapture>]>; def int_nvvm_atomic_load_dec_32 : Intrinsic<[llvm_i32_ty], [LLVMAnyPointerType, llvm_i32_ty], [IntrArgMemOnly, NoCapture>]>; class SCOPED_ATOMIC2_impl : Intrinsic<[elty], [LLVMAnyPointerType>, LLVMMatchType<0>], [IntrArgMemOnly, NoCapture>]>; class SCOPED_ATOMIC3_impl : Intrinsic<[elty], [LLVMAnyPointerType>, LLVMMatchType<0>, LLVMMatchType<0>], [IntrArgMemOnly, NoCapture>]>; multiclass PTXAtomicWithScope2 { def _cta : SCOPED_ATOMIC2_impl; def _sys : SCOPED_ATOMIC2_impl; } multiclass PTXAtomicWithScope3 { def _cta : SCOPED_ATOMIC3_impl; def _sys : SCOPED_ATOMIC3_impl; } multiclass PTXAtomicWithScope2_fi { defm _f: PTXAtomicWithScope2; defm _i: PTXAtomicWithScope2; } defm int_nvvm_atomic_add_gen : PTXAtomicWithScope2_fi; defm int_nvvm_atomic_inc_gen_i : PTXAtomicWithScope2; defm int_nvvm_atomic_dec_gen_i : PTXAtomicWithScope2; defm int_nvvm_atomic_exch_gen_i: PTXAtomicWithScope2; defm int_nvvm_atomic_xor_gen_i : PTXAtomicWithScope2; defm int_nvvm_atomic_max_gen_i : PTXAtomicWithScope2; defm int_nvvm_atomic_min_gen_i : PTXAtomicWithScope2; defm int_nvvm_atomic_or_gen_i : PTXAtomicWithScope2; defm int_nvvm_atomic_and_gen_i : PTXAtomicWithScope2; defm int_nvvm_atomic_cas_gen_i : PTXAtomicWithScope3; // Bar.Sync // The builtin for "bar.sync 0" is called __syncthreads. Unlike most of the // intrinsics in this file, this one is a user-facing API. def int_nvvm_barrier0 : GCCBuiltin<"__syncthreads">, Intrinsic<[], [], [IntrConvergent]>; // Synchronize all threads in the CTA at barrier 'n'. def int_nvvm_barrier_n : GCCBuiltin<"__nvvm_bar_n">, Intrinsic<[], [llvm_i32_ty], [IntrConvergent]>; // Synchronize 'm', a multiple of warp size, (arg 2) threads in // the CTA at barrier 'n' (arg 1). def int_nvvm_barrier : GCCBuiltin<"__nvvm_bar">, Intrinsic<[], [llvm_i32_ty, llvm_i32_ty], [IntrConvergent]>; def int_nvvm_barrier0_popc : GCCBuiltin<"__nvvm_bar0_popc">, Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrConvergent]>; def int_nvvm_barrier0_and : GCCBuiltin<"__nvvm_bar0_and">, Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrConvergent]>; def int_nvvm_barrier0_or : GCCBuiltin<"__nvvm_bar0_or">, Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrConvergent]>; def int_nvvm_bar_sync : Intrinsic<[], [llvm_i32_ty], [IntrConvergent]>, GCCBuiltin<"__nvvm_bar_sync">; def int_nvvm_bar_warp_sync : Intrinsic<[], [llvm_i32_ty], [IntrConvergent]>, GCCBuiltin<"__nvvm_bar_warp_sync">; // barrier.sync id[, cnt] def int_nvvm_barrier_sync : Intrinsic<[], [llvm_i32_ty], [IntrConvergent]>, GCCBuiltin<"__nvvm_barrier_sync">; def int_nvvm_barrier_sync_cnt : Intrinsic<[], [llvm_i32_ty, llvm_i32_ty], [IntrConvergent]>, GCCBuiltin<"__nvvm_barrier_sync_cnt">; // Membar def int_nvvm_membar_cta : GCCBuiltin<"__nvvm_membar_cta">, Intrinsic<[], [], []>; def int_nvvm_membar_gl : GCCBuiltin<"__nvvm_membar_gl">, Intrinsic<[], [], []>; def int_nvvm_membar_sys : GCCBuiltin<"__nvvm_membar_sys">, Intrinsic<[], [], []>; // Generated within nvvm. Use for ldu on sm_20 or later. Second arg is the // pointer's alignment. def int_nvvm_ldu_global_i : Intrinsic<[llvm_anyint_ty], [LLVMAnyPointerType>, llvm_i32_ty], [IntrReadMem, IntrArgMemOnly, NoCapture>], "llvm.nvvm.ldu.global.i">; def int_nvvm_ldu_global_f : Intrinsic<[llvm_anyfloat_ty], [LLVMAnyPointerType>, llvm_i32_ty], [IntrReadMem, IntrArgMemOnly, NoCapture>], "llvm.nvvm.ldu.global.f">; def int_nvvm_ldu_global_p : Intrinsic<[llvm_anyptr_ty], [LLVMAnyPointerType>, llvm_i32_ty], [IntrReadMem, IntrArgMemOnly, NoCapture>], "llvm.nvvm.ldu.global.p">; // Generated within nvvm. Use for ldg on sm_35 or later. Second arg is the // pointer's alignment. def int_nvvm_ldg_global_i : Intrinsic<[llvm_anyint_ty], [LLVMAnyPointerType>, llvm_i32_ty], [IntrReadMem, IntrArgMemOnly, NoCapture>], "llvm.nvvm.ldg.global.i">; def int_nvvm_ldg_global_f : Intrinsic<[llvm_anyfloat_ty], [LLVMAnyPointerType>, llvm_i32_ty], [IntrReadMem, IntrArgMemOnly, NoCapture>], "llvm.nvvm.ldg.global.f">; def int_nvvm_ldg_global_p : Intrinsic<[llvm_anyptr_ty], [LLVMAnyPointerType>, llvm_i32_ty], [IntrReadMem, IntrArgMemOnly, NoCapture>], "llvm.nvvm.ldg.global.p">; // Use for generic pointers // - These intrinsics are used to convert address spaces. // - The input pointer and output pointer must have the same type, except for // the address-space. (This restriction is not enforced here as there is // currently no way to describe it). // - This complements the llvm bitcast, which can be used to cast one type // of pointer to another type of pointer, while the address space remains // the same. def int_nvvm_ptr_local_to_gen: Intrinsic<[llvm_anyptr_ty], [llvm_anyptr_ty], [IntrNoMem], "llvm.nvvm.ptr.local.to.gen">; def int_nvvm_ptr_shared_to_gen: Intrinsic<[llvm_anyptr_ty], [llvm_anyptr_ty], [IntrNoMem], "llvm.nvvm.ptr.shared.to.gen">; def int_nvvm_ptr_global_to_gen: Intrinsic<[llvm_anyptr_ty], [llvm_anyptr_ty], [IntrNoMem], "llvm.nvvm.ptr.global.to.gen">; def int_nvvm_ptr_constant_to_gen: Intrinsic<[llvm_anyptr_ty], [llvm_anyptr_ty], [IntrNoMem], "llvm.nvvm.ptr.constant.to.gen">; def int_nvvm_ptr_gen_to_global: Intrinsic<[llvm_anyptr_ty], [llvm_anyptr_ty], [IntrNoMem], "llvm.nvvm.ptr.gen.to.global">; def int_nvvm_ptr_gen_to_shared: Intrinsic<[llvm_anyptr_ty], [llvm_anyptr_ty], [IntrNoMem], "llvm.nvvm.ptr.gen.to.shared">; def int_nvvm_ptr_gen_to_local: Intrinsic<[llvm_anyptr_ty], [llvm_anyptr_ty], [IntrNoMem], "llvm.nvvm.ptr.gen.to.local">; def int_nvvm_ptr_gen_to_constant: Intrinsic<[llvm_anyptr_ty], [llvm_anyptr_ty], [IntrNoMem], "llvm.nvvm.ptr.gen.to.constant">; // Used in nvvm internally to help address space opt and ptx code generation // This is for params that are passed to kernel functions by pointer by-val. def int_nvvm_ptr_gen_to_param: Intrinsic<[llvm_anyptr_ty], [llvm_anyptr_ty], [IntrNoMem], "llvm.nvvm.ptr.gen.to.param">; // Move intrinsics, used in nvvm internally def int_nvvm_move_i16 : Intrinsic<[llvm_i16_ty], [llvm_i16_ty], [IntrNoMem], "llvm.nvvm.move.i16">; def int_nvvm_move_i32 : Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem], "llvm.nvvm.move.i32">; def int_nvvm_move_i64 : Intrinsic<[llvm_i64_ty], [llvm_i64_ty], [IntrNoMem], "llvm.nvvm.move.i64">; def int_nvvm_move_float : Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem], "llvm.nvvm.move.float">; def int_nvvm_move_double : Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem], "llvm.nvvm.move.double">; def int_nvvm_move_ptr : Intrinsic<[llvm_anyptr_ty], [llvm_anyptr_ty], [IntrNoMem, NoCapture>], "llvm.nvvm.move.ptr">; // For getting the handle from a texture or surface variable def int_nvvm_texsurf_handle : Intrinsic<[llvm_i64_ty], [llvm_metadata_ty, llvm_anyi64ptr_ty], [IntrNoMem], "llvm.nvvm.texsurf.handle">; def int_nvvm_texsurf_handle_internal : Intrinsic<[llvm_i64_ty], [llvm_anyptr_ty], [IntrNoMem], "llvm.nvvm.texsurf.handle.internal">; /// Error / Warn def int_nvvm_compiler_error : Intrinsic<[], [llvm_anyptr_ty], [], "llvm.nvvm.compiler.error">; def int_nvvm_compiler_warn : Intrinsic<[], [llvm_anyptr_ty], [], "llvm.nvvm.compiler.warn">; def int_nvvm_reflect : Intrinsic<[llvm_i32_ty], [llvm_anyptr_ty], [IntrNoMem], "llvm.nvvm.reflect">; // isspacep.{const, global, local, shared} def int_nvvm_isspacep_const : Intrinsic<[llvm_i1_ty], [llvm_ptr_ty], [IntrNoMem], "llvm.nvvm.isspacep.const">, GCCBuiltin<"__nvvm_isspacep_const">; def int_nvvm_isspacep_global : Intrinsic<[llvm_i1_ty], [llvm_ptr_ty], [IntrNoMem], "llvm.nvvm.isspacep.global">, GCCBuiltin<"__nvvm_isspacep_global">; def int_nvvm_isspacep_local : Intrinsic<[llvm_i1_ty], [llvm_ptr_ty], [IntrNoMem], "llvm.nvvm.isspacep.local">, GCCBuiltin<"__nvvm_isspacep_local">; def int_nvvm_isspacep_shared : Intrinsic<[llvm_i1_ty], [llvm_ptr_ty], [IntrNoMem], "llvm.nvvm.isspacep.shared">, GCCBuiltin<"__nvvm_isspacep_shared">; // Environment register read def int_nvvm_read_ptx_sreg_envreg0 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem], "llvm.nvvm.read.ptx.sreg.envreg0">, GCCBuiltin<"__nvvm_read_ptx_sreg_envreg0">; def int_nvvm_read_ptx_sreg_envreg1 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem], "llvm.nvvm.read.ptx.sreg.envreg1">, GCCBuiltin<"__nvvm_read_ptx_sreg_envreg1">; def int_nvvm_read_ptx_sreg_envreg2 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem], "llvm.nvvm.read.ptx.sreg.envreg2">, GCCBuiltin<"__nvvm_read_ptx_sreg_envreg2">; def int_nvvm_read_ptx_sreg_envreg3 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem], "llvm.nvvm.read.ptx.sreg.envreg3">, GCCBuiltin<"__nvvm_read_ptx_sreg_envreg3">; def int_nvvm_read_ptx_sreg_envreg4 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem], "llvm.nvvm.read.ptx.sreg.envreg4">, GCCBuiltin<"__nvvm_read_ptx_sreg_envreg4">; def int_nvvm_read_ptx_sreg_envreg5 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem], "llvm.nvvm.read.ptx.sreg.envreg5">, GCCBuiltin<"__nvvm_read_ptx_sreg_envreg5">; def int_nvvm_read_ptx_sreg_envreg6 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem], "llvm.nvvm.read.ptx.sreg.envreg6">, GCCBuiltin<"__nvvm_read_ptx_sreg_envreg6">; def int_nvvm_read_ptx_sreg_envreg7 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem], "llvm.nvvm.read.ptx.sreg.envreg7">, GCCBuiltin<"__nvvm_read_ptx_sreg_envreg7">; def int_nvvm_read_ptx_sreg_envreg8 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem], "llvm.nvvm.read.ptx.sreg.envreg8">, GCCBuiltin<"__nvvm_read_ptx_sreg_envreg8">; def int_nvvm_read_ptx_sreg_envreg9 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem], "llvm.nvvm.read.ptx.sreg.envreg9">, GCCBuiltin<"__nvvm_read_ptx_sreg_envreg9">; def int_nvvm_read_ptx_sreg_envreg10 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem], "llvm.nvvm.read.ptx.sreg.envreg10">, GCCBuiltin<"__nvvm_read_ptx_sreg_envreg10">; def int_nvvm_read_ptx_sreg_envreg11 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem], "llvm.nvvm.read.ptx.sreg.envreg11">, GCCBuiltin<"__nvvm_read_ptx_sreg_envreg11">; def int_nvvm_read_ptx_sreg_envreg12 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem], "llvm.nvvm.read.ptx.sreg.envreg12">, GCCBuiltin<"__nvvm_read_ptx_sreg_envreg12">; def int_nvvm_read_ptx_sreg_envreg13 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem], "llvm.nvvm.read.ptx.sreg.envreg13">, GCCBuiltin<"__nvvm_read_ptx_sreg_envreg13">; def int_nvvm_read_ptx_sreg_envreg14 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem], "llvm.nvvm.read.ptx.sreg.envreg14">, GCCBuiltin<"__nvvm_read_ptx_sreg_envreg14">; def int_nvvm_read_ptx_sreg_envreg15 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem], "llvm.nvvm.read.ptx.sreg.envreg15">, GCCBuiltin<"__nvvm_read_ptx_sreg_envreg15">; def int_nvvm_read_ptx_sreg_envreg16 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem], "llvm.nvvm.read.ptx.sreg.envreg16">, GCCBuiltin<"__nvvm_read_ptx_sreg_envreg16">; def int_nvvm_read_ptx_sreg_envreg17 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem], "llvm.nvvm.read.ptx.sreg.envreg17">, GCCBuiltin<"__nvvm_read_ptx_sreg_envreg17">; def int_nvvm_read_ptx_sreg_envreg18 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem], "llvm.nvvm.read.ptx.sreg.envreg18">, GCCBuiltin<"__nvvm_read_ptx_sreg_envreg18">; def int_nvvm_read_ptx_sreg_envreg19 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem], "llvm.nvvm.read.ptx.sreg.envreg19">, GCCBuiltin<"__nvvm_read_ptx_sreg_envreg19">; def int_nvvm_read_ptx_sreg_envreg20 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem], "llvm.nvvm.read.ptx.sreg.envreg20">, GCCBuiltin<"__nvvm_read_ptx_sreg_envreg20">; def int_nvvm_read_ptx_sreg_envreg21 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem], "llvm.nvvm.read.ptx.sreg.envreg21">, GCCBuiltin<"__nvvm_read_ptx_sreg_envreg21">; def int_nvvm_read_ptx_sreg_envreg22 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem], "llvm.nvvm.read.ptx.sreg.envreg22">, GCCBuiltin<"__nvvm_read_ptx_sreg_envreg22">; def int_nvvm_read_ptx_sreg_envreg23 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem], "llvm.nvvm.read.ptx.sreg.envreg23">, GCCBuiltin<"__nvvm_read_ptx_sreg_envreg23">; def int_nvvm_read_ptx_sreg_envreg24 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem], "llvm.nvvm.read.ptx.sreg.envreg24">, GCCBuiltin<"__nvvm_read_ptx_sreg_envreg24">; def int_nvvm_read_ptx_sreg_envreg25 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem], "llvm.nvvm.read.ptx.sreg.envreg25">, GCCBuiltin<"__nvvm_read_ptx_sreg_envreg25">; def int_nvvm_read_ptx_sreg_envreg26 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem], "llvm.nvvm.read.ptx.sreg.envreg26">, GCCBuiltin<"__nvvm_read_ptx_sreg_envreg26">; def int_nvvm_read_ptx_sreg_envreg27 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem], "llvm.nvvm.read.ptx.sreg.envreg27">, GCCBuiltin<"__nvvm_read_ptx_sreg_envreg27">; def int_nvvm_read_ptx_sreg_envreg28 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem], "llvm.nvvm.read.ptx.sreg.envreg28">, GCCBuiltin<"__nvvm_read_ptx_sreg_envreg28">; def int_nvvm_read_ptx_sreg_envreg29 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem], "llvm.nvvm.read.ptx.sreg.envreg29">, GCCBuiltin<"__nvvm_read_ptx_sreg_envreg29">; def int_nvvm_read_ptx_sreg_envreg30 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem], "llvm.nvvm.read.ptx.sreg.envreg30">, GCCBuiltin<"__nvvm_read_ptx_sreg_envreg30">; def int_nvvm_read_ptx_sreg_envreg31 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem], "llvm.nvvm.read.ptx.sreg.envreg31">, GCCBuiltin<"__nvvm_read_ptx_sreg_envreg31">; // Texture Fetch // texmode_independent def int_nvvm_tex_1d_v4f32_s32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.tex.1d.v4f32.s32">; def int_nvvm_tex_1d_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty], [], "llvm.nvvm.tex.1d.v4f32.f32">; def int_nvvm_tex_1d_level_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.1d.level.v4f32.f32">; def int_nvvm_tex_1d_grad_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.1d.grad.v4f32.f32">; def int_nvvm_tex_1d_v4s32_s32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.tex.1d.v4s32.s32">; def int_nvvm_tex_1d_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty], [], "llvm.nvvm.tex.1d.v4s32.f32">; def int_nvvm_tex_1d_level_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.1d.level.v4s32.f32">; def int_nvvm_tex_1d_grad_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.1d.grad.v4s32.f32">; def int_nvvm_tex_1d_v4u32_s32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.tex.1d.v4u32.s32">; def int_nvvm_tex_1d_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty], [], "llvm.nvvm.tex.1d.v4u32.f32">; def int_nvvm_tex_1d_level_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.1d.level.v4u32.f32">; def int_nvvm_tex_1d_grad_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.1d.grad.v4u32.f32">; def int_nvvm_tex_1d_array_v4f32_s32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.tex.1d.array.v4f32.s32">; def int_nvvm_tex_1d_array_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty], [], "llvm.nvvm.tex.1d.array.v4f32.f32">; def int_nvvm_tex_1d_array_level_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.1d.array.level.v4f32.f32">; def int_nvvm_tex_1d_array_grad_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.1d.array.grad.v4f32.f32">; def int_nvvm_tex_1d_array_v4s32_s32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.tex.1d.array.v4s32.s32">; def int_nvvm_tex_1d_array_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty], [], "llvm.nvvm.tex.1d.array.v4s32.f32">; def int_nvvm_tex_1d_array_level_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.1d.array.level.v4s32.f32">; def int_nvvm_tex_1d_array_grad_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.1d.array.grad.v4s32.f32">; def int_nvvm_tex_1d_array_v4u32_s32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.tex.1d.array.v4u32.s32">; def int_nvvm_tex_1d_array_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty], [], "llvm.nvvm.tex.1d.array.v4u32.f32">; def int_nvvm_tex_1d_array_level_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.1d.array.level.v4u32.f32">; def int_nvvm_tex_1d_array_grad_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.1d.array.grad.v4u32.f32">; def int_nvvm_tex_2d_v4f32_s32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.tex.2d.v4f32.s32">; def int_nvvm_tex_2d_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.2d.v4f32.f32">; def int_nvvm_tex_2d_level_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.2d.level.v4f32.f32">; def int_nvvm_tex_2d_grad_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.2d.grad.v4f32.f32">; def int_nvvm_tex_2d_v4s32_s32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.tex.2d.v4s32.s32">; def int_nvvm_tex_2d_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.2d.v4s32.f32">; def int_nvvm_tex_2d_level_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.2d.level.v4s32.f32">; def int_nvvm_tex_2d_grad_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.2d.grad.v4s32.f32">; def int_nvvm_tex_2d_v4u32_s32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.tex.2d.v4u32.s32">; def int_nvvm_tex_2d_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.2d.v4u32.f32">; def int_nvvm_tex_2d_level_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.2d.level.v4u32.f32">; def int_nvvm_tex_2d_grad_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.2d.grad.v4u32.f32">; def int_nvvm_tex_2d_array_v4f32_s32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.tex.2d.array.v4f32.s32">; def int_nvvm_tex_2d_array_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.2d.array.v4f32.f32">; def int_nvvm_tex_2d_array_level_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.2d.array.level.v4f32.f32">; def int_nvvm_tex_2d_array_grad_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.2d.array.grad.v4f32.f32">; def int_nvvm_tex_2d_array_v4s32_s32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.tex.2d.array.v4s32.s32">; def int_nvvm_tex_2d_array_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.2d.array.v4s32.f32">; def int_nvvm_tex_2d_array_level_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.2d.array.level.v4s32.f32">; def int_nvvm_tex_2d_array_grad_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.2d.array.grad.v4s32.f32">; def int_nvvm_tex_2d_array_v4u32_s32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.tex.2d.array.v4u32.s32">; def int_nvvm_tex_2d_array_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.2d.array.v4u32.f32">; def int_nvvm_tex_2d_array_level_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.2d.array.level.v4u32.f32">; def int_nvvm_tex_2d_array_grad_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.2d.array.grad.v4u32.f32">; def int_nvvm_tex_3d_v4f32_s32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.tex.3d.v4f32.s32">; def int_nvvm_tex_3d_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.3d.v4f32.f32">; def int_nvvm_tex_3d_level_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.3d.level.v4f32.f32">; def int_nvvm_tex_3d_grad_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.3d.grad.v4f32.f32">; def int_nvvm_tex_3d_v4s32_s32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.tex.3d.v4s32.s32">; def int_nvvm_tex_3d_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.3d.v4s32.f32">; def int_nvvm_tex_3d_level_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.3d.level.v4s32.f32">; def int_nvvm_tex_3d_grad_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.3d.grad.v4s32.f32">; def int_nvvm_tex_3d_v4u32_s32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.tex.3d.v4u32.s32">; def int_nvvm_tex_3d_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.3d.v4u32.f32">; def int_nvvm_tex_3d_level_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.3d.level.v4u32.f32">; def int_nvvm_tex_3d_grad_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.3d.grad.v4u32.f32">; def int_nvvm_tex_cube_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.cube.v4f32.f32">; def int_nvvm_tex_cube_level_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.cube.level.v4f32.f32">; def int_nvvm_tex_cube_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.cube.v4s32.f32">; def int_nvvm_tex_cube_level_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.cube.level.v4s32.f32">; def int_nvvm_tex_cube_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.cube.v4u32.f32">; def int_nvvm_tex_cube_level_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.cube.level.v4u32.f32">; def int_nvvm_tex_cube_array_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.cube.array.v4f32.f32">; def int_nvvm_tex_cube_array_level_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.cube.array.level.v4f32.f32">; def int_nvvm_tex_cube_array_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.cube.array.v4s32.f32">; def int_nvvm_tex_cube_array_level_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.cube.array.level.v4s32.f32">; def int_nvvm_tex_cube_array_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.cube.array.v4u32.f32">; def int_nvvm_tex_cube_array_level_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.cube.array.level.v4u32.f32">; def int_nvvm_tld4_r_2d_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tld4.r.2d.v4f32.f32">; def int_nvvm_tld4_g_2d_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tld4.g.2d.v4f32.f32">; def int_nvvm_tld4_b_2d_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tld4.b.2d.v4f32.f32">; def int_nvvm_tld4_a_2d_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tld4.a.2d.v4f32.f32">; def int_nvvm_tld4_r_2d_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tld4.r.2d.v4s32.f32">; def int_nvvm_tld4_g_2d_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tld4.g.2d.v4s32.f32">; def int_nvvm_tld4_b_2d_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tld4.b.2d.v4s32.f32">; def int_nvvm_tld4_a_2d_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tld4.a.2d.v4s32.f32">; def int_nvvm_tld4_r_2d_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tld4.r.2d.v4u32.f32">; def int_nvvm_tld4_g_2d_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tld4.g.2d.v4u32.f32">; def int_nvvm_tld4_b_2d_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tld4.b.2d.v4u32.f32">; def int_nvvm_tld4_a_2d_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tld4.a.2d.v4u32.f32">; // texmode_unified def int_nvvm_tex_unified_1d_v4f32_s32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.tex.unified.1d.v4f32.s32">; def int_nvvm_tex_unified_1d_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.1d.v4f32.f32">; def int_nvvm_tex_unified_1d_level_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.1d.level.v4f32.f32">; def int_nvvm_tex_unified_1d_grad_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.1d.grad.v4f32.f32">; def int_nvvm_tex_unified_1d_v4s32_s32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.tex.unified.1d.v4s32.s32">; def int_nvvm_tex_unified_1d_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.1d.v4s32.f32">; def int_nvvm_tex_unified_1d_level_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.1d.level.v4s32.f32">; def int_nvvm_tex_unified_1d_grad_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.1d.grad.v4s32.f32">; def int_nvvm_tex_unified_1d_v4u32_s32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.tex.unified.1d.v4u32.s32">; def int_nvvm_tex_unified_1d_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.1d.v4u32.f32">; def int_nvvm_tex_unified_1d_level_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.1d.level.v4u32.f32">; def int_nvvm_tex_unified_1d_grad_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.1d.grad.v4u32.f32">; def int_nvvm_tex_unified_1d_array_v4f32_s32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.tex.unified.1d.array.v4f32.s32">; def int_nvvm_tex_unified_1d_array_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i32_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.1d.array.v4f32.f32">; def int_nvvm_tex_unified_1d_array_level_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.1d.array.level.v4f32.f32">; def int_nvvm_tex_unified_1d_array_grad_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.1d.array.grad.v4f32.f32">; def int_nvvm_tex_unified_1d_array_v4s32_s32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.tex.unified.1d.array.v4s32.s32">; def int_nvvm_tex_unified_1d_array_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.1d.array.v4s32.f32">; def int_nvvm_tex_unified_1d_array_level_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.1d.array.level.v4s32.f32">; def int_nvvm_tex_unified_1d_array_grad_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.1d.array.grad.v4s32.f32">; def int_nvvm_tex_unified_1d_array_v4u32_s32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.tex.unified.1d.array.v4u32.s32">; def int_nvvm_tex_unified_1d_array_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.1d.array.v4u32.f32">; def int_nvvm_tex_unified_1d_array_level_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.1d.array.level.v4u32.f32">; def int_nvvm_tex_unified_1d_array_grad_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.1d.array.grad.v4u32.f32">; def int_nvvm_tex_unified_2d_v4f32_s32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.tex.unified.2d.v4f32.s32">; def int_nvvm_tex_unified_2d_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.2d.v4f32.f32">; def int_nvvm_tex_unified_2d_level_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.2d.level.v4f32.f32">; def int_nvvm_tex_unified_2d_grad_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.2d.grad.v4f32.f32">; def int_nvvm_tex_unified_2d_v4s32_s32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.tex.unified.2d.v4s32.s32">; def int_nvvm_tex_unified_2d_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.2d.v4s32.f32">; def int_nvvm_tex_unified_2d_level_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.2d.level.v4s32.f32">; def int_nvvm_tex_unified_2d_grad_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.2d.grad.v4s32.f32">; def int_nvvm_tex_unified_2d_v4u32_s32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.tex.unified.2d.v4u32.s32">; def int_nvvm_tex_unified_2d_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.2d.v4u32.f32">; def int_nvvm_tex_unified_2d_level_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.2d.level.v4u32.f32">; def int_nvvm_tex_unified_2d_grad_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.2d.grad.v4u32.f32">; def int_nvvm_tex_unified_2d_array_v4f32_s32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.tex.unified.2d.array.v4f32.s32">; def int_nvvm_tex_unified_2d_array_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.2d.array.v4f32.f32">; def int_nvvm_tex_unified_2d_array_level_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.2d.array.level.v4f32.f32">; def int_nvvm_tex_unified_2d_array_grad_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.2d.array.grad.v4f32.f32">; def int_nvvm_tex_unified_2d_array_v4s32_s32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.tex.unified.2d.array.v4s32.s32">; def int_nvvm_tex_unified_2d_array_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.2d.array.v4s32.f32">; def int_nvvm_tex_unified_2d_array_level_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.2d.array.level.v4s32.f32">; def int_nvvm_tex_unified_2d_array_grad_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.2d.array.grad.v4s32.f32">; def int_nvvm_tex_unified_2d_array_v4u32_s32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.tex.unified.2d.array.v4u32.s32">; def int_nvvm_tex_unified_2d_array_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.2d.array.v4u32.f32">; def int_nvvm_tex_unified_2d_array_level_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.2d.array.level.v4u32.f32">; def int_nvvm_tex_unified_2d_array_grad_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.2d.array.grad.v4u32.f32">; def int_nvvm_tex_unified_3d_v4f32_s32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.tex.unified.3d.v4f32.s32">; def int_nvvm_tex_unified_3d_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.3d.v4f32.f32">; def int_nvvm_tex_unified_3d_level_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.3d.level.v4f32.f32">; def int_nvvm_tex_unified_3d_grad_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.3d.grad.v4f32.f32">; def int_nvvm_tex_unified_3d_v4s32_s32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.tex.unified.3d.v4s32.s32">; def int_nvvm_tex_unified_3d_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.3d.v4s32.f32">; def int_nvvm_tex_unified_3d_level_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.3d.level.v4s32.f32">; def int_nvvm_tex_unified_3d_grad_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.3d.grad.v4s32.f32">; def int_nvvm_tex_unified_3d_v4u32_s32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.tex.unified.3d.v4u32.s32">; def int_nvvm_tex_unified_3d_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.3d.v4u32.f32">; def int_nvvm_tex_unified_3d_level_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.3d.level.v4u32.f32">; def int_nvvm_tex_unified_3d_grad_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.3d.grad.v4u32.f32">; def int_nvvm_tex_unified_cube_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.cube.v4f32.f32">; def int_nvvm_tex_unified_cube_level_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.cube.level.v4f32.f32">; def int_nvvm_tex_unified_cube_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.cube.v4s32.f32">; def int_nvvm_tex_unified_cube_level_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.cube.level.v4s32.f32">; def int_nvvm_tex_unified_cube_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.cube.v4u32.f32">; def int_nvvm_tex_unified_cube_level_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.cube.level.v4u32.f32">; def int_nvvm_tex_unified_cube_array_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.cube.array.v4f32.f32">; def int_nvvm_tex_unified_cube_array_level_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.cube.array.level.v4f32.f32">; def int_nvvm_tex_unified_cube_array_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.cube.array.v4s32.f32">; def int_nvvm_tex_unified_cube_array_level_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.cube.array.level.v4s32.f32">; def int_nvvm_tex_unified_cube_array_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.cube.array.v4u32.f32">; def int_nvvm_tex_unified_cube_array_level_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tex.unified.cube.array.level.v4u32.f32">; def int_nvvm_tld4_unified_r_2d_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tld4.unified.r.2d.v4f32.f32">; def int_nvvm_tld4_unified_g_2d_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tld4.unified.g.2d.v4f32.f32">; def int_nvvm_tld4_unified_b_2d_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tld4.unified.b.2d.v4f32.f32">; def int_nvvm_tld4_unified_a_2d_v4f32_f32 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tld4.unified.a.2d.v4f32.f32">; def int_nvvm_tld4_unified_r_2d_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tld4.unified.r.2d.v4s32.f32">; def int_nvvm_tld4_unified_g_2d_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tld4.unified.g.2d.v4s32.f32">; def int_nvvm_tld4_unified_b_2d_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tld4.unified.b.2d.v4s32.f32">; def int_nvvm_tld4_unified_a_2d_v4s32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tld4.unified.a.2d.v4s32.f32">; def int_nvvm_tld4_unified_r_2d_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tld4.unified.r.2d.v4u32.f32">; def int_nvvm_tld4_unified_g_2d_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tld4.unified.g.2d.v4u32.f32">; def int_nvvm_tld4_unified_b_2d_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tld4.unified.b.2d.v4u32.f32">; def int_nvvm_tld4_unified_a_2d_v4u32_f32 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_float_ty, llvm_float_ty], [], "llvm.nvvm.tld4.unified.a.2d.v4u32.f32">; //=== Surface Load // .clamp variants def int_nvvm_suld_1d_i8_clamp : Intrinsic<[llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.i8.clamp">; def int_nvvm_suld_1d_i16_clamp : Intrinsic<[llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.i16.clamp">; def int_nvvm_suld_1d_i32_clamp : Intrinsic<[llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.i32.clamp">; def int_nvvm_suld_1d_i64_clamp : Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.i64.clamp">; def int_nvvm_suld_1d_v2i8_clamp : Intrinsic<[llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.v2i8.clamp">; def int_nvvm_suld_1d_v2i16_clamp : Intrinsic<[llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.v2i16.clamp">; def int_nvvm_suld_1d_v2i32_clamp : Intrinsic<[llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.v2i32.clamp">; def int_nvvm_suld_1d_v2i64_clamp : Intrinsic<[llvm_i64_ty, llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.v2i64.clamp">; def int_nvvm_suld_1d_v4i8_clamp : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.v4i8.clamp">; def int_nvvm_suld_1d_v4i16_clamp : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.v4i16.clamp">; def int_nvvm_suld_1d_v4i32_clamp : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.v4i32.clamp">; def int_nvvm_suld_1d_array_i8_clamp : Intrinsic<[llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.array.i8.clamp">; def int_nvvm_suld_1d_array_i16_clamp : Intrinsic<[llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.array.i16.clamp">; def int_nvvm_suld_1d_array_i32_clamp : Intrinsic<[llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.array.i32.clamp">; def int_nvvm_suld_1d_array_i64_clamp : Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.array.i64.clamp">; def int_nvvm_suld_1d_array_v2i8_clamp : Intrinsic<[llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.array.v2i8.clamp">; def int_nvvm_suld_1d_array_v2i16_clamp : Intrinsic<[llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.array.v2i16.clamp">; def int_nvvm_suld_1d_array_v2i32_clamp : Intrinsic<[llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.array.v2i32.clamp">; def int_nvvm_suld_1d_array_v2i64_clamp : Intrinsic<[llvm_i64_ty, llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.array.v2i64.clamp">; def int_nvvm_suld_1d_array_v4i8_clamp : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.array.v4i8.clamp">; def int_nvvm_suld_1d_array_v4i16_clamp : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.array.v4i16.clamp">; def int_nvvm_suld_1d_array_v4i32_clamp : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.array.v4i32.clamp">; def int_nvvm_suld_2d_i8_clamp : Intrinsic<[llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.i8.clamp">; def int_nvvm_suld_2d_i16_clamp : Intrinsic<[llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.i16.clamp">; def int_nvvm_suld_2d_i32_clamp : Intrinsic<[llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.i32.clamp">; def int_nvvm_suld_2d_i64_clamp : Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.i64.clamp">; def int_nvvm_suld_2d_v2i8_clamp : Intrinsic<[llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.v2i8.clamp">; def int_nvvm_suld_2d_v2i16_clamp : Intrinsic<[llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.v2i16.clamp">; def int_nvvm_suld_2d_v2i32_clamp : Intrinsic<[llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.v2i32.clamp">; def int_nvvm_suld_2d_v2i64_clamp : Intrinsic<[llvm_i64_ty, llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.v2i64.clamp">; def int_nvvm_suld_2d_v4i8_clamp : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.v4i8.clamp">; def int_nvvm_suld_2d_v4i16_clamp : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.v4i16.clamp">; def int_nvvm_suld_2d_v4i32_clamp : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.v4i32.clamp">; def int_nvvm_suld_2d_array_i8_clamp : Intrinsic<[llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.array.i8.clamp">; def int_nvvm_suld_2d_array_i16_clamp : Intrinsic<[llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.array.i16.clamp">; def int_nvvm_suld_2d_array_i32_clamp : Intrinsic<[llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.array.i32.clamp">; def int_nvvm_suld_2d_array_i64_clamp : Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.array.i64.clamp">; def int_nvvm_suld_2d_array_v2i8_clamp : Intrinsic<[llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.array.v2i8.clamp">; def int_nvvm_suld_2d_array_v2i16_clamp : Intrinsic<[llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.array.v2i16.clamp">; def int_nvvm_suld_2d_array_v2i32_clamp : Intrinsic<[llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.array.v2i32.clamp">; def int_nvvm_suld_2d_array_v2i64_clamp : Intrinsic<[llvm_i64_ty, llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.array.v2i64.clamp">; def int_nvvm_suld_2d_array_v4i8_clamp : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.array.v4i8.clamp">; def int_nvvm_suld_2d_array_v4i16_clamp : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.array.v4i16.clamp">; def int_nvvm_suld_2d_array_v4i32_clamp : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.array.v4i32.clamp">; def int_nvvm_suld_3d_i8_clamp : Intrinsic<[llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.3d.i8.clamp">; def int_nvvm_suld_3d_i16_clamp : Intrinsic<[llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.3d.i16.clamp">; def int_nvvm_suld_3d_i32_clamp : Intrinsic<[llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.3d.i32.clamp">; def int_nvvm_suld_3d_i64_clamp : Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.3d.i64.clamp">; def int_nvvm_suld_3d_v2i8_clamp : Intrinsic<[llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.3d.v2i8.clamp">; def int_nvvm_suld_3d_v2i16_clamp : Intrinsic<[llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.3d.v2i16.clamp">; def int_nvvm_suld_3d_v2i32_clamp : Intrinsic<[llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.3d.v2i32.clamp">; def int_nvvm_suld_3d_v2i64_clamp : Intrinsic<[llvm_i64_ty, llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.3d.v2i64.clamp">; def int_nvvm_suld_3d_v4i8_clamp : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.3d.v4i8.clamp">; def int_nvvm_suld_3d_v4i16_clamp : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.3d.v4i16.clamp">; def int_nvvm_suld_3d_v4i32_clamp : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.3d.v4i32.clamp">; // .trap variants def int_nvvm_suld_1d_i8_trap : Intrinsic<[llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.i8.trap">; def int_nvvm_suld_1d_i16_trap : Intrinsic<[llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.i16.trap">; def int_nvvm_suld_1d_i32_trap : Intrinsic<[llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.i32.trap">; def int_nvvm_suld_1d_i64_trap : Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.i64.trap">; def int_nvvm_suld_1d_v2i8_trap : Intrinsic<[llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.v2i8.trap">; def int_nvvm_suld_1d_v2i16_trap : Intrinsic<[llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.v2i16.trap">; def int_nvvm_suld_1d_v2i32_trap : Intrinsic<[llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.v2i32.trap">; def int_nvvm_suld_1d_v2i64_trap : Intrinsic<[llvm_i64_ty, llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.v2i64.trap">; def int_nvvm_suld_1d_v4i8_trap : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.v4i8.trap">; def int_nvvm_suld_1d_v4i16_trap : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.v4i16.trap">; def int_nvvm_suld_1d_v4i32_trap : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.v4i32.trap">; def int_nvvm_suld_1d_array_i8_trap : Intrinsic<[llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.array.i8.trap">; def int_nvvm_suld_1d_array_i16_trap : Intrinsic<[llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.array.i16.trap">; def int_nvvm_suld_1d_array_i32_trap : Intrinsic<[llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.array.i32.trap">; def int_nvvm_suld_1d_array_i64_trap : Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.array.i64.trap">; def int_nvvm_suld_1d_array_v2i8_trap : Intrinsic<[llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.array.v2i8.trap">; def int_nvvm_suld_1d_array_v2i16_trap : Intrinsic<[llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.array.v2i16.trap">; def int_nvvm_suld_1d_array_v2i32_trap : Intrinsic<[llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.array.v2i32.trap">; def int_nvvm_suld_1d_array_v2i64_trap : Intrinsic<[llvm_i64_ty, llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.array.v2i64.trap">; def int_nvvm_suld_1d_array_v4i8_trap : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.array.v4i8.trap">; def int_nvvm_suld_1d_array_v4i16_trap : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.array.v4i16.trap">; def int_nvvm_suld_1d_array_v4i32_trap : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.array.v4i32.trap">; def int_nvvm_suld_2d_i8_trap : Intrinsic<[llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.i8.trap">; def int_nvvm_suld_2d_i16_trap : Intrinsic<[llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.i16.trap">; def int_nvvm_suld_2d_i32_trap : Intrinsic<[llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.i32.trap">; def int_nvvm_suld_2d_i64_trap : Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.i64.trap">; def int_nvvm_suld_2d_v2i8_trap : Intrinsic<[llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.v2i8.trap">; def int_nvvm_suld_2d_v2i16_trap : Intrinsic<[llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.v2i16.trap">; def int_nvvm_suld_2d_v2i32_trap : Intrinsic<[llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.v2i32.trap">; def int_nvvm_suld_2d_v2i64_trap : Intrinsic<[llvm_i64_ty, llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.v2i64.trap">; def int_nvvm_suld_2d_v4i8_trap : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.v4i8.trap">; def int_nvvm_suld_2d_v4i16_trap : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.v4i16.trap">; def int_nvvm_suld_2d_v4i32_trap : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.v4i32.trap">; def int_nvvm_suld_2d_array_i8_trap : Intrinsic<[llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.array.i8.trap">; def int_nvvm_suld_2d_array_i16_trap : Intrinsic<[llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.array.i16.trap">; def int_nvvm_suld_2d_array_i32_trap : Intrinsic<[llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.array.i32.trap">; def int_nvvm_suld_2d_array_i64_trap : Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.array.i64.trap">; def int_nvvm_suld_2d_array_v2i8_trap : Intrinsic<[llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.array.v2i8.trap">; def int_nvvm_suld_2d_array_v2i16_trap : Intrinsic<[llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.array.v2i16.trap">; def int_nvvm_suld_2d_array_v2i32_trap : Intrinsic<[llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.array.v2i32.trap">; def int_nvvm_suld_2d_array_v2i64_trap : Intrinsic<[llvm_i64_ty, llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.array.v2i64.trap">; def int_nvvm_suld_2d_array_v4i8_trap : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.array.v4i8.trap">; def int_nvvm_suld_2d_array_v4i16_trap : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.array.v4i16.trap">; def int_nvvm_suld_2d_array_v4i32_trap : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.array.v4i32.trap">; def int_nvvm_suld_3d_i8_trap : Intrinsic<[llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.3d.i8.trap">; def int_nvvm_suld_3d_i16_trap : Intrinsic<[llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.3d.i16.trap">; def int_nvvm_suld_3d_i32_trap : Intrinsic<[llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.3d.i32.trap">; def int_nvvm_suld_3d_i64_trap : Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.3d.i64.trap">; def int_nvvm_suld_3d_v2i8_trap : Intrinsic<[llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.3d.v2i8.trap">; def int_nvvm_suld_3d_v2i16_trap : Intrinsic<[llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.3d.v2i16.trap">; def int_nvvm_suld_3d_v2i32_trap : Intrinsic<[llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.3d.v2i32.trap">; def int_nvvm_suld_3d_v2i64_trap : Intrinsic<[llvm_i64_ty, llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.3d.v2i64.trap">; def int_nvvm_suld_3d_v4i8_trap : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.3d.v4i8.trap">; def int_nvvm_suld_3d_v4i16_trap : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.3d.v4i16.trap">; def int_nvvm_suld_3d_v4i32_trap : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.3d.v4i32.trap">; // .zero variants def int_nvvm_suld_1d_i8_zero : Intrinsic<[llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.i8.zero">; def int_nvvm_suld_1d_i16_zero : Intrinsic<[llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.i16.zero">; def int_nvvm_suld_1d_i32_zero : Intrinsic<[llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.i32.zero">; def int_nvvm_suld_1d_i64_zero : Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.i64.zero">; def int_nvvm_suld_1d_v2i8_zero : Intrinsic<[llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.v2i8.zero">; def int_nvvm_suld_1d_v2i16_zero : Intrinsic<[llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.v2i16.zero">; def int_nvvm_suld_1d_v2i32_zero : Intrinsic<[llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.v2i32.zero">; def int_nvvm_suld_1d_v2i64_zero : Intrinsic<[llvm_i64_ty, llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.v2i64.zero">; def int_nvvm_suld_1d_v4i8_zero : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.v4i8.zero">; def int_nvvm_suld_1d_v4i16_zero : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.v4i16.zero">; def int_nvvm_suld_1d_v4i32_zero : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.v4i32.zero">; def int_nvvm_suld_1d_array_i8_zero : Intrinsic<[llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.array.i8.zero">; def int_nvvm_suld_1d_array_i16_zero : Intrinsic<[llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.array.i16.zero">; def int_nvvm_suld_1d_array_i32_zero : Intrinsic<[llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.array.i32.zero">; def int_nvvm_suld_1d_array_i64_zero : Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.array.i64.zero">; def int_nvvm_suld_1d_array_v2i8_zero : Intrinsic<[llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.array.v2i8.zero">; def int_nvvm_suld_1d_array_v2i16_zero : Intrinsic<[llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.array.v2i16.zero">; def int_nvvm_suld_1d_array_v2i32_zero : Intrinsic<[llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.array.v2i32.zero">; def int_nvvm_suld_1d_array_v2i64_zero : Intrinsic<[llvm_i64_ty, llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.array.v2i64.zero">; def int_nvvm_suld_1d_array_v4i8_zero : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.array.v4i8.zero">; def int_nvvm_suld_1d_array_v4i16_zero : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.array.v4i16.zero">; def int_nvvm_suld_1d_array_v4i32_zero : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.1d.array.v4i32.zero">; def int_nvvm_suld_2d_i8_zero : Intrinsic<[llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.i8.zero">; def int_nvvm_suld_2d_i16_zero : Intrinsic<[llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.i16.zero">; def int_nvvm_suld_2d_i32_zero : Intrinsic<[llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.i32.zero">; def int_nvvm_suld_2d_i64_zero : Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.i64.zero">; def int_nvvm_suld_2d_v2i8_zero : Intrinsic<[llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.v2i8.zero">; def int_nvvm_suld_2d_v2i16_zero : Intrinsic<[llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.v2i16.zero">; def int_nvvm_suld_2d_v2i32_zero : Intrinsic<[llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.v2i32.zero">; def int_nvvm_suld_2d_v2i64_zero : Intrinsic<[llvm_i64_ty, llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.v2i64.zero">; def int_nvvm_suld_2d_v4i8_zero : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.v4i8.zero">; def int_nvvm_suld_2d_v4i16_zero : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.v4i16.zero">; def int_nvvm_suld_2d_v4i32_zero : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.v4i32.zero">; def int_nvvm_suld_2d_array_i8_zero : Intrinsic<[llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.array.i8.zero">; def int_nvvm_suld_2d_array_i16_zero : Intrinsic<[llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.array.i16.zero">; def int_nvvm_suld_2d_array_i32_zero : Intrinsic<[llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.array.i32.zero">; def int_nvvm_suld_2d_array_i64_zero : Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.array.i64.zero">; def int_nvvm_suld_2d_array_v2i8_zero : Intrinsic<[llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.array.v2i8.zero">; def int_nvvm_suld_2d_array_v2i16_zero : Intrinsic<[llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.array.v2i16.zero">; def int_nvvm_suld_2d_array_v2i32_zero : Intrinsic<[llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.array.v2i32.zero">; def int_nvvm_suld_2d_array_v2i64_zero : Intrinsic<[llvm_i64_ty, llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.array.v2i64.zero">; def int_nvvm_suld_2d_array_v4i8_zero : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.array.v4i8.zero">; def int_nvvm_suld_2d_array_v4i16_zero : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.array.v4i16.zero">; def int_nvvm_suld_2d_array_v4i32_zero : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.2d.array.v4i32.zero">; def int_nvvm_suld_3d_i8_zero : Intrinsic<[llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.3d.i8.zero">; def int_nvvm_suld_3d_i16_zero : Intrinsic<[llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.3d.i16.zero">; def int_nvvm_suld_3d_i32_zero : Intrinsic<[llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.3d.i32.zero">; def int_nvvm_suld_3d_i64_zero : Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.3d.i64.zero">; def int_nvvm_suld_3d_v2i8_zero : Intrinsic<[llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.3d.v2i8.zero">; def int_nvvm_suld_3d_v2i16_zero : Intrinsic<[llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.3d.v2i16.zero">; def int_nvvm_suld_3d_v2i32_zero : Intrinsic<[llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.3d.v2i32.zero">; def int_nvvm_suld_3d_v2i64_zero : Intrinsic<[llvm_i64_ty, llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.3d.v2i64.zero">; def int_nvvm_suld_3d_v4i8_zero : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.3d.v4i8.zero">; def int_nvvm_suld_3d_v4i16_zero : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.3d.v4i16.zero">; def int_nvvm_suld_3d_v4i32_zero : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.suld.3d.v4i32.zero">; //===- Texture Query ------------------------------------------------------===// def int_nvvm_txq_channel_order : Intrinsic<[llvm_i32_ty], [llvm_i64_ty], [IntrNoMem], "llvm.nvvm.txq.channel.order">, GCCBuiltin<"__nvvm_txq_channel_order">; def int_nvvm_txq_channel_data_type : Intrinsic<[llvm_i32_ty], [llvm_i64_ty], [IntrNoMem], "llvm.nvvm.txq.channel.data.type">, GCCBuiltin<"__nvvm_txq_channel_data_type">; def int_nvvm_txq_width : Intrinsic<[llvm_i32_ty], [llvm_i64_ty], [IntrNoMem], "llvm.nvvm.txq.width">, GCCBuiltin<"__nvvm_txq_width">; def int_nvvm_txq_height : Intrinsic<[llvm_i32_ty], [llvm_i64_ty], [IntrNoMem], "llvm.nvvm.txq.height">, GCCBuiltin<"__nvvm_txq_height">; def int_nvvm_txq_depth : Intrinsic<[llvm_i32_ty], [llvm_i64_ty], [IntrNoMem], "llvm.nvvm.txq.depth">, GCCBuiltin<"__nvvm_txq_depth">; def int_nvvm_txq_array_size : Intrinsic<[llvm_i32_ty], [llvm_i64_ty], [IntrNoMem], "llvm.nvvm.txq.array.size">, GCCBuiltin<"__nvvm_txq_array_size">; def int_nvvm_txq_num_samples : Intrinsic<[llvm_i32_ty], [llvm_i64_ty], [IntrNoMem], "llvm.nvvm.txq.num.samples">, GCCBuiltin<"__nvvm_txq_num_samples">; def int_nvvm_txq_num_mipmap_levels : Intrinsic<[llvm_i32_ty], [llvm_i64_ty], [IntrNoMem], "llvm.nvvm.txq.num.mipmap.levels">, GCCBuiltin<"__nvvm_txq_num_mipmap_levels">; //===- Surface Query ------------------------------------------------------===// def int_nvvm_suq_channel_order : Intrinsic<[llvm_i32_ty], [llvm_i64_ty], [IntrNoMem], "llvm.nvvm.suq.channel.order">, GCCBuiltin<"__nvvm_suq_channel_order">; def int_nvvm_suq_channel_data_type : Intrinsic<[llvm_i32_ty], [llvm_i64_ty], [IntrNoMem], "llvm.nvvm.suq.channel.data.type">, GCCBuiltin<"__nvvm_suq_channel_data_type">; def int_nvvm_suq_width : Intrinsic<[llvm_i32_ty], [llvm_i64_ty], [IntrNoMem], "llvm.nvvm.suq.width">, GCCBuiltin<"__nvvm_suq_width">; def int_nvvm_suq_height : Intrinsic<[llvm_i32_ty], [llvm_i64_ty], [IntrNoMem], "llvm.nvvm.suq.height">, GCCBuiltin<"__nvvm_suq_height">; def int_nvvm_suq_depth : Intrinsic<[llvm_i32_ty], [llvm_i64_ty], [IntrNoMem], "llvm.nvvm.suq.depth">, GCCBuiltin<"__nvvm_suq_depth">; def int_nvvm_suq_array_size : Intrinsic<[llvm_i32_ty], [llvm_i64_ty], [IntrNoMem], "llvm.nvvm.suq.array.size">, GCCBuiltin<"__nvvm_suq_array_size">; //===- Handle Query -------------------------------------------------------===// def int_nvvm_istypep_sampler : Intrinsic<[llvm_i1_ty], [llvm_i64_ty], [IntrNoMem], "llvm.nvvm.istypep.sampler">, GCCBuiltin<"__nvvm_istypep_sampler">; def int_nvvm_istypep_surface : Intrinsic<[llvm_i1_ty], [llvm_i64_ty], [IntrNoMem], "llvm.nvvm.istypep.surface">, GCCBuiltin<"__nvvm_istypep_surface">; def int_nvvm_istypep_texture : Intrinsic<[llvm_i1_ty], [llvm_i64_ty], [IntrNoMem], "llvm.nvvm.istypep.texture">, GCCBuiltin<"__nvvm_istypep_texture">; //===- Surface Stores -----------------------------------------------------===// // Unformatted // .clamp variant def int_nvvm_sust_b_1d_i8_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.1d.i8.clamp">, GCCBuiltin<"__nvvm_sust_b_1d_i8_clamp">; def int_nvvm_sust_b_1d_i16_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.1d.i16.clamp">, GCCBuiltin<"__nvvm_sust_b_1d_i16_clamp">; def int_nvvm_sust_b_1d_i32_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.1d.i32.clamp">, GCCBuiltin<"__nvvm_sust_b_1d_i32_clamp">; def int_nvvm_sust_b_1d_i64_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i64_ty], [], "llvm.nvvm.sust.b.1d.i64.clamp">, GCCBuiltin<"__nvvm_sust_b_1d_i64_clamp">; def int_nvvm_sust_b_1d_v2i8_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.1d.v2i8.clamp">, GCCBuiltin<"__nvvm_sust_b_1d_v2i8_clamp">; def int_nvvm_sust_b_1d_v2i16_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.1d.v2i16.clamp">, GCCBuiltin<"__nvvm_sust_b_1d_v2i16_clamp">; def int_nvvm_sust_b_1d_v2i32_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.1d.v2i32.clamp">, GCCBuiltin<"__nvvm_sust_b_1d_v2i32_clamp">; def int_nvvm_sust_b_1d_v2i64_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i64_ty, llvm_i64_ty], [], "llvm.nvvm.sust.b.1d.v2i64.clamp">, GCCBuiltin<"__nvvm_sust_b_1d_v2i64_clamp">; def int_nvvm_sust_b_1d_v4i8_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.1d.v4i8.clamp">, GCCBuiltin<"__nvvm_sust_b_1d_v4i8_clamp">; def int_nvvm_sust_b_1d_v4i16_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.1d.v4i16.clamp">, GCCBuiltin<"__nvvm_sust_b_1d_v4i16_clamp">; def int_nvvm_sust_b_1d_v4i32_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.1d.v4i32.clamp">, GCCBuiltin<"__nvvm_sust_b_1d_v4i32_clamp">; def int_nvvm_sust_b_1d_array_i8_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.1d.array.i8.clamp">, GCCBuiltin<"__nvvm_sust_b_1d_array_i8_clamp">; def int_nvvm_sust_b_1d_array_i16_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.1d.array.i16.clamp">, GCCBuiltin<"__nvvm_sust_b_1d_array_i16_clamp">; def int_nvvm_sust_b_1d_array_i32_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.1d.array.i32.clamp">, GCCBuiltin<"__nvvm_sust_b_1d_array_i32_clamp">; def int_nvvm_sust_b_1d_array_i64_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i64_ty], [], "llvm.nvvm.sust.b.1d.array.i64.clamp">, GCCBuiltin<"__nvvm_sust_b_1d_array_i64_clamp">; def int_nvvm_sust_b_1d_array_v2i8_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.1d.array.v2i8.clamp">, GCCBuiltin<"__nvvm_sust_b_1d_array_v2i8_clamp">; def int_nvvm_sust_b_1d_array_v2i16_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.1d.array.v2i16.clamp">, GCCBuiltin<"__nvvm_sust_b_1d_array_v2i16_clamp">; def int_nvvm_sust_b_1d_array_v2i32_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.1d.array.v2i32.clamp">, GCCBuiltin<"__nvvm_sust_b_1d_array_v2i32_clamp">; def int_nvvm_sust_b_1d_array_v2i64_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i64_ty, llvm_i64_ty], [], "llvm.nvvm.sust.b.1d.array.v2i64.clamp">, GCCBuiltin<"__nvvm_sust_b_1d_array_v2i64_clamp">; def int_nvvm_sust_b_1d_array_v4i8_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.1d.array.v4i8.clamp">, GCCBuiltin<"__nvvm_sust_b_1d_array_v4i8_clamp">; def int_nvvm_sust_b_1d_array_v4i16_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.1d.array.v4i16.clamp">, GCCBuiltin<"__nvvm_sust_b_1d_array_v4i16_clamp">; def int_nvvm_sust_b_1d_array_v4i32_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.1d.array.v4i32.clamp">, GCCBuiltin<"__nvvm_sust_b_1d_array_v4i32_clamp">; def int_nvvm_sust_b_2d_i8_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.2d.i8.clamp">, GCCBuiltin<"__nvvm_sust_b_2d_i8_clamp">; def int_nvvm_sust_b_2d_i16_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.2d.i16.clamp">, GCCBuiltin<"__nvvm_sust_b_2d_i16_clamp">; def int_nvvm_sust_b_2d_i32_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.2d.i32.clamp">, GCCBuiltin<"__nvvm_sust_b_2d_i32_clamp">; def int_nvvm_sust_b_2d_i64_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i64_ty], [], "llvm.nvvm.sust.b.2d.i64.clamp">, GCCBuiltin<"__nvvm_sust_b_2d_i64_clamp">; def int_nvvm_sust_b_2d_v2i8_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.2d.v2i8.clamp">, GCCBuiltin<"__nvvm_sust_b_2d_v2i8_clamp">; def int_nvvm_sust_b_2d_v2i16_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.2d.v2i16.clamp">, GCCBuiltin<"__nvvm_sust_b_2d_v2i16_clamp">; def int_nvvm_sust_b_2d_v2i32_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.2d.v2i32.clamp">, GCCBuiltin<"__nvvm_sust_b_2d_v2i32_clamp">; def int_nvvm_sust_b_2d_v2i64_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i64_ty, llvm_i64_ty], [], "llvm.nvvm.sust.b.2d.v2i64.clamp">, GCCBuiltin<"__nvvm_sust_b_2d_v2i64_clamp">; def int_nvvm_sust_b_2d_v4i8_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.2d.v4i8.clamp">, GCCBuiltin<"__nvvm_sust_b_2d_v4i8_clamp">; def int_nvvm_sust_b_2d_v4i16_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.2d.v4i16.clamp">, GCCBuiltin<"__nvvm_sust_b_2d_v4i16_clamp">; def int_nvvm_sust_b_2d_v4i32_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.2d.v4i32.clamp">, GCCBuiltin<"__nvvm_sust_b_2d_v4i32_clamp">; def int_nvvm_sust_b_2d_array_i8_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.2d.array.i8.clamp">, GCCBuiltin<"__nvvm_sust_b_2d_array_i8_clamp">; def int_nvvm_sust_b_2d_array_i16_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.2d.array.i16.clamp">, GCCBuiltin<"__nvvm_sust_b_2d_array_i16_clamp">; def int_nvvm_sust_b_2d_array_i32_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.2d.array.i32.clamp">, GCCBuiltin<"__nvvm_sust_b_2d_array_i32_clamp">; def int_nvvm_sust_b_2d_array_i64_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i64_ty], [], "llvm.nvvm.sust.b.2d.array.i64.clamp">, GCCBuiltin<"__nvvm_sust_b_2d_array_i64_clamp">; def int_nvvm_sust_b_2d_array_v2i8_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.2d.array.v2i8.clamp">, GCCBuiltin<"__nvvm_sust_b_2d_array_v2i8_clamp">; def int_nvvm_sust_b_2d_array_v2i16_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.2d.array.v2i16.clamp">, GCCBuiltin<"__nvvm_sust_b_2d_array_v2i16_clamp">; def int_nvvm_sust_b_2d_array_v2i32_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.2d.array.v2i32.clamp">, GCCBuiltin<"__nvvm_sust_b_2d_array_v2i32_clamp">; def int_nvvm_sust_b_2d_array_v2i64_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i64_ty, llvm_i64_ty], [], "llvm.nvvm.sust.b.2d.array.v2i64.clamp">, GCCBuiltin<"__nvvm_sust_b_2d_array_v2i64_clamp">; def int_nvvm_sust_b_2d_array_v4i8_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.2d.array.v4i8.clamp">, GCCBuiltin<"__nvvm_sust_b_2d_array_v4i8_clamp">; def int_nvvm_sust_b_2d_array_v4i16_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.2d.array.v4i16.clamp">, GCCBuiltin<"__nvvm_sust_b_2d_array_v4i16_clamp">; def int_nvvm_sust_b_2d_array_v4i32_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.2d.array.v4i32.clamp">, GCCBuiltin<"__nvvm_sust_b_2d_array_v4i32_clamp">; def int_nvvm_sust_b_3d_i8_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.3d.i8.clamp">, GCCBuiltin<"__nvvm_sust_b_3d_i8_clamp">; def int_nvvm_sust_b_3d_i16_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.3d.i16.clamp">, GCCBuiltin<"__nvvm_sust_b_3d_i16_clamp">; def int_nvvm_sust_b_3d_i32_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.3d.i32.clamp">, GCCBuiltin<"__nvvm_sust_b_3d_i32_clamp">; def int_nvvm_sust_b_3d_i64_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i64_ty], [], "llvm.nvvm.sust.b.3d.i64.clamp">, GCCBuiltin<"__nvvm_sust_b_3d_i64_clamp">; def int_nvvm_sust_b_3d_v2i8_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.3d.v2i8.clamp">, GCCBuiltin<"__nvvm_sust_b_3d_v2i8_clamp">; def int_nvvm_sust_b_3d_v2i16_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.3d.v2i16.clamp">, GCCBuiltin<"__nvvm_sust_b_3d_v2i16_clamp">; def int_nvvm_sust_b_3d_v2i32_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.3d.v2i32.clamp">, GCCBuiltin<"__nvvm_sust_b_3d_v2i32_clamp">; def int_nvvm_sust_b_3d_v2i64_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i64_ty, llvm_i64_ty], [], "llvm.nvvm.sust.b.3d.v2i64.clamp">, GCCBuiltin<"__nvvm_sust_b_3d_v2i64_clamp">; def int_nvvm_sust_b_3d_v4i8_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.3d.v4i8.clamp">, GCCBuiltin<"__nvvm_sust_b_3d_v4i8_clamp">; def int_nvvm_sust_b_3d_v4i16_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.3d.v4i16.clamp">, GCCBuiltin<"__nvvm_sust_b_3d_v4i16_clamp">; def int_nvvm_sust_b_3d_v4i32_clamp : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.3d.v4i32.clamp">, GCCBuiltin<"__nvvm_sust_b_3d_v4i32_clamp">; // .trap variant def int_nvvm_sust_b_1d_i8_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.1d.i8.trap">, GCCBuiltin<"__nvvm_sust_b_1d_i8_trap">; def int_nvvm_sust_b_1d_i16_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.1d.i16.trap">, GCCBuiltin<"__nvvm_sust_b_1d_i16_trap">; def int_nvvm_sust_b_1d_i32_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.1d.i32.trap">, GCCBuiltin<"__nvvm_sust_b_1d_i32_trap">; def int_nvvm_sust_b_1d_i64_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i64_ty], [], "llvm.nvvm.sust.b.1d.i64.trap">, GCCBuiltin<"__nvvm_sust_b_1d_i64_trap">; def int_nvvm_sust_b_1d_v2i8_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.1d.v2i8.trap">, GCCBuiltin<"__nvvm_sust_b_1d_v2i8_trap">; def int_nvvm_sust_b_1d_v2i16_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.1d.v2i16.trap">, GCCBuiltin<"__nvvm_sust_b_1d_v2i16_trap">; def int_nvvm_sust_b_1d_v2i32_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.1d.v2i32.trap">, GCCBuiltin<"__nvvm_sust_b_1d_v2i32_trap">; def int_nvvm_sust_b_1d_v2i64_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i64_ty, llvm_i64_ty], [], "llvm.nvvm.sust.b.1d.v2i64.trap">, GCCBuiltin<"__nvvm_sust_b_1d_v2i64_trap">; def int_nvvm_sust_b_1d_v4i8_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.1d.v4i8.trap">, GCCBuiltin<"__nvvm_sust_b_1d_v4i8_trap">; def int_nvvm_sust_b_1d_v4i16_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.1d.v4i16.trap">, GCCBuiltin<"__nvvm_sust_b_1d_v4i16_trap">; def int_nvvm_sust_b_1d_v4i32_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.1d.v4i32.trap">, GCCBuiltin<"__nvvm_sust_b_1d_v4i32_trap">; def int_nvvm_sust_b_1d_array_i8_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.1d.array.i8.trap">, GCCBuiltin<"__nvvm_sust_b_1d_array_i8_trap">; def int_nvvm_sust_b_1d_array_i16_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.1d.array.i16.trap">, GCCBuiltin<"__nvvm_sust_b_1d_array_i16_trap">; def int_nvvm_sust_b_1d_array_i32_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.1d.array.i32.trap">, GCCBuiltin<"__nvvm_sust_b_1d_array_i32_trap">; def int_nvvm_sust_b_1d_array_i64_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i64_ty], [], "llvm.nvvm.sust.b.1d.array.i64.trap">, GCCBuiltin<"__nvvm_sust_b_1d_array_i64_trap">; def int_nvvm_sust_b_1d_array_v2i8_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.1d.array.v2i8.trap">, GCCBuiltin<"__nvvm_sust_b_1d_array_v2i8_trap">; def int_nvvm_sust_b_1d_array_v2i16_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.1d.array.v2i16.trap">, GCCBuiltin<"__nvvm_sust_b_1d_array_v2i16_trap">; def int_nvvm_sust_b_1d_array_v2i32_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.1d.array.v2i32.trap">, GCCBuiltin<"__nvvm_sust_b_1d_array_v2i32_trap">; def int_nvvm_sust_b_1d_array_v2i64_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i64_ty, llvm_i64_ty], [], "llvm.nvvm.sust.b.1d.array.v2i64.trap">, GCCBuiltin<"__nvvm_sust_b_1d_array_v2i64_trap">; def int_nvvm_sust_b_1d_array_v4i8_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.1d.array.v4i8.trap">, GCCBuiltin<"__nvvm_sust_b_1d_array_v4i8_trap">; def int_nvvm_sust_b_1d_array_v4i16_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.1d.array.v4i16.trap">, GCCBuiltin<"__nvvm_sust_b_1d_array_v4i16_trap">; def int_nvvm_sust_b_1d_array_v4i32_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.1d.array.v4i32.trap">, GCCBuiltin<"__nvvm_sust_b_1d_array_v4i32_trap">; def int_nvvm_sust_b_2d_i8_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.2d.i8.trap">, GCCBuiltin<"__nvvm_sust_b_2d_i8_trap">; def int_nvvm_sust_b_2d_i16_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.2d.i16.trap">, GCCBuiltin<"__nvvm_sust_b_2d_i16_trap">; def int_nvvm_sust_b_2d_i32_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.2d.i32.trap">, GCCBuiltin<"__nvvm_sust_b_2d_i32_trap">; def int_nvvm_sust_b_2d_i64_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i64_ty], [], "llvm.nvvm.sust.b.2d.i64.trap">, GCCBuiltin<"__nvvm_sust_b_2d_i64_trap">; def int_nvvm_sust_b_2d_v2i8_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.2d.v2i8.trap">, GCCBuiltin<"__nvvm_sust_b_2d_v2i8_trap">; def int_nvvm_sust_b_2d_v2i16_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.2d.v2i16.trap">, GCCBuiltin<"__nvvm_sust_b_2d_v2i16_trap">; def int_nvvm_sust_b_2d_v2i32_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.2d.v2i32.trap">, GCCBuiltin<"__nvvm_sust_b_2d_v2i32_trap">; def int_nvvm_sust_b_2d_v2i64_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i64_ty, llvm_i64_ty], [], "llvm.nvvm.sust.b.2d.v2i64.trap">, GCCBuiltin<"__nvvm_sust_b_2d_v2i64_trap">; def int_nvvm_sust_b_2d_v4i8_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.2d.v4i8.trap">, GCCBuiltin<"__nvvm_sust_b_2d_v4i8_trap">; def int_nvvm_sust_b_2d_v4i16_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.2d.v4i16.trap">, GCCBuiltin<"__nvvm_sust_b_2d_v4i16_trap">; def int_nvvm_sust_b_2d_v4i32_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.2d.v4i32.trap">, GCCBuiltin<"__nvvm_sust_b_2d_v4i32_trap">; def int_nvvm_sust_b_2d_array_i8_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.2d.array.i8.trap">, GCCBuiltin<"__nvvm_sust_b_2d_array_i8_trap">; def int_nvvm_sust_b_2d_array_i16_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.2d.array.i16.trap">, GCCBuiltin<"__nvvm_sust_b_2d_array_i16_trap">; def int_nvvm_sust_b_2d_array_i32_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.2d.array.i32.trap">, GCCBuiltin<"__nvvm_sust_b_2d_array_i32_trap">; def int_nvvm_sust_b_2d_array_i64_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i64_ty], [], "llvm.nvvm.sust.b.2d.array.i64.trap">, GCCBuiltin<"__nvvm_sust_b_2d_array_i64_trap">; def int_nvvm_sust_b_2d_array_v2i8_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.2d.array.v2i8.trap">, GCCBuiltin<"__nvvm_sust_b_2d_array_v2i8_trap">; def int_nvvm_sust_b_2d_array_v2i16_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.2d.array.v2i16.trap">, GCCBuiltin<"__nvvm_sust_b_2d_array_v2i16_trap">; def int_nvvm_sust_b_2d_array_v2i32_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.2d.array.v2i32.trap">, GCCBuiltin<"__nvvm_sust_b_2d_array_v2i32_trap">; def int_nvvm_sust_b_2d_array_v2i64_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i64_ty, llvm_i64_ty], [], "llvm.nvvm.sust.b.2d.array.v2i64.trap">, GCCBuiltin<"__nvvm_sust_b_2d_array_v2i64_trap">; def int_nvvm_sust_b_2d_array_v4i8_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.2d.array.v4i8.trap">, GCCBuiltin<"__nvvm_sust_b_2d_array_v4i8_trap">; def int_nvvm_sust_b_2d_array_v4i16_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.2d.array.v4i16.trap">, GCCBuiltin<"__nvvm_sust_b_2d_array_v4i16_trap">; def int_nvvm_sust_b_2d_array_v4i32_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.2d.array.v4i32.trap">, GCCBuiltin<"__nvvm_sust_b_2d_array_v4i32_trap">; def int_nvvm_sust_b_3d_i8_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.3d.i8.trap">, GCCBuiltin<"__nvvm_sust_b_3d_i8_trap">; def int_nvvm_sust_b_3d_i16_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.3d.i16.trap">, GCCBuiltin<"__nvvm_sust_b_3d_i16_trap">; def int_nvvm_sust_b_3d_i32_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.3d.i32.trap">, GCCBuiltin<"__nvvm_sust_b_3d_i32_trap">; def int_nvvm_sust_b_3d_i64_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i64_ty], [], "llvm.nvvm.sust.b.3d.i64.trap">, GCCBuiltin<"__nvvm_sust_b_3d_i64_trap">; def int_nvvm_sust_b_3d_v2i8_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.3d.v2i8.trap">, GCCBuiltin<"__nvvm_sust_b_3d_v2i8_trap">; def int_nvvm_sust_b_3d_v2i16_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.3d.v2i16.trap">, GCCBuiltin<"__nvvm_sust_b_3d_v2i16_trap">; def int_nvvm_sust_b_3d_v2i32_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.3d.v2i32.trap">, GCCBuiltin<"__nvvm_sust_b_3d_v2i32_trap">; def int_nvvm_sust_b_3d_v2i64_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i64_ty, llvm_i64_ty], [], "llvm.nvvm.sust.b.3d.v2i64.trap">, GCCBuiltin<"__nvvm_sust_b_3d_v2i64_trap">; def int_nvvm_sust_b_3d_v4i8_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.3d.v4i8.trap">, GCCBuiltin<"__nvvm_sust_b_3d_v4i8_trap">; def int_nvvm_sust_b_3d_v4i16_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.3d.v4i16.trap">, GCCBuiltin<"__nvvm_sust_b_3d_v4i16_trap">; def int_nvvm_sust_b_3d_v4i32_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.3d.v4i32.trap">, GCCBuiltin<"__nvvm_sust_b_3d_v4i32_trap">; // .zero variant def int_nvvm_sust_b_1d_i8_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.1d.i8.zero">, GCCBuiltin<"__nvvm_sust_b_1d_i8_zero">; def int_nvvm_sust_b_1d_i16_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.1d.i16.zero">, GCCBuiltin<"__nvvm_sust_b_1d_i16_zero">; def int_nvvm_sust_b_1d_i32_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.1d.i32.zero">, GCCBuiltin<"__nvvm_sust_b_1d_i32_zero">; def int_nvvm_sust_b_1d_i64_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i64_ty], [], "llvm.nvvm.sust.b.1d.i64.zero">, GCCBuiltin<"__nvvm_sust_b_1d_i64_zero">; def int_nvvm_sust_b_1d_v2i8_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.1d.v2i8.zero">, GCCBuiltin<"__nvvm_sust_b_1d_v2i8_zero">; def int_nvvm_sust_b_1d_v2i16_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.1d.v2i16.zero">, GCCBuiltin<"__nvvm_sust_b_1d_v2i16_zero">; def int_nvvm_sust_b_1d_v2i32_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.1d.v2i32.zero">, GCCBuiltin<"__nvvm_sust_b_1d_v2i32_zero">; def int_nvvm_sust_b_1d_v2i64_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i64_ty, llvm_i64_ty], [], "llvm.nvvm.sust.b.1d.v2i64.zero">, GCCBuiltin<"__nvvm_sust_b_1d_v2i64_zero">; def int_nvvm_sust_b_1d_v4i8_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.1d.v4i8.zero">, GCCBuiltin<"__nvvm_sust_b_1d_v4i8_zero">; def int_nvvm_sust_b_1d_v4i16_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.1d.v4i16.zero">, GCCBuiltin<"__nvvm_sust_b_1d_v4i16_zero">; def int_nvvm_sust_b_1d_v4i32_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.1d.v4i32.zero">, GCCBuiltin<"__nvvm_sust_b_1d_v4i32_zero">; def int_nvvm_sust_b_1d_array_i8_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.1d.array.i8.zero">, GCCBuiltin<"__nvvm_sust_b_1d_array_i8_zero">; def int_nvvm_sust_b_1d_array_i16_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.1d.array.i16.zero">, GCCBuiltin<"__nvvm_sust_b_1d_array_i16_zero">; def int_nvvm_sust_b_1d_array_i32_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.1d.array.i32.zero">, GCCBuiltin<"__nvvm_sust_b_1d_array_i32_zero">; def int_nvvm_sust_b_1d_array_i64_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i64_ty], [], "llvm.nvvm.sust.b.1d.array.i64.zero">, GCCBuiltin<"__nvvm_sust_b_1d_array_i64_zero">; def int_nvvm_sust_b_1d_array_v2i8_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.1d.array.v2i8.zero">, GCCBuiltin<"__nvvm_sust_b_1d_array_v2i8_zero">; def int_nvvm_sust_b_1d_array_v2i16_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.1d.array.v2i16.zero">, GCCBuiltin<"__nvvm_sust_b_1d_array_v2i16_zero">; def int_nvvm_sust_b_1d_array_v2i32_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.1d.array.v2i32.zero">, GCCBuiltin<"__nvvm_sust_b_1d_array_v2i32_zero">; def int_nvvm_sust_b_1d_array_v2i64_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i64_ty, llvm_i64_ty], [], "llvm.nvvm.sust.b.1d.array.v2i64.zero">, GCCBuiltin<"__nvvm_sust_b_1d_array_v2i64_zero">; def int_nvvm_sust_b_1d_array_v4i8_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.1d.array.v4i8.zero">, GCCBuiltin<"__nvvm_sust_b_1d_array_v4i8_zero">; def int_nvvm_sust_b_1d_array_v4i16_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.1d.array.v4i16.zero">, GCCBuiltin<"__nvvm_sust_b_1d_array_v4i16_zero">; def int_nvvm_sust_b_1d_array_v4i32_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.1d.array.v4i32.zero">, GCCBuiltin<"__nvvm_sust_b_1d_array_v4i32_zero">; def int_nvvm_sust_b_2d_i8_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.2d.i8.zero">, GCCBuiltin<"__nvvm_sust_b_2d_i8_zero">; def int_nvvm_sust_b_2d_i16_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.2d.i16.zero">, GCCBuiltin<"__nvvm_sust_b_2d_i16_zero">; def int_nvvm_sust_b_2d_i32_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.2d.i32.zero">, GCCBuiltin<"__nvvm_sust_b_2d_i32_zero">; def int_nvvm_sust_b_2d_i64_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i64_ty], [], "llvm.nvvm.sust.b.2d.i64.zero">, GCCBuiltin<"__nvvm_sust_b_2d_i64_zero">; def int_nvvm_sust_b_2d_v2i8_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.2d.v2i8.zero">, GCCBuiltin<"__nvvm_sust_b_2d_v2i8_zero">; def int_nvvm_sust_b_2d_v2i16_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.2d.v2i16.zero">, GCCBuiltin<"__nvvm_sust_b_2d_v2i16_zero">; def int_nvvm_sust_b_2d_v2i32_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.2d.v2i32.zero">, GCCBuiltin<"__nvvm_sust_b_2d_v2i32_zero">; def int_nvvm_sust_b_2d_v2i64_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i64_ty, llvm_i64_ty], [], "llvm.nvvm.sust.b.2d.v2i64.zero">, GCCBuiltin<"__nvvm_sust_b_2d_v2i64_zero">; def int_nvvm_sust_b_2d_v4i8_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.2d.v4i8.zero">, GCCBuiltin<"__nvvm_sust_b_2d_v4i8_zero">; def int_nvvm_sust_b_2d_v4i16_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.2d.v4i16.zero">, GCCBuiltin<"__nvvm_sust_b_2d_v4i16_zero">; def int_nvvm_sust_b_2d_v4i32_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.2d.v4i32.zero">, GCCBuiltin<"__nvvm_sust_b_2d_v4i32_zero">; def int_nvvm_sust_b_2d_array_i8_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.2d.array.i8.zero">, GCCBuiltin<"__nvvm_sust_b_2d_array_i8_zero">; def int_nvvm_sust_b_2d_array_i16_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.2d.array.i16.zero">, GCCBuiltin<"__nvvm_sust_b_2d_array_i16_zero">; def int_nvvm_sust_b_2d_array_i32_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.2d.array.i32.zero">, GCCBuiltin<"__nvvm_sust_b_2d_array_i32_zero">; def int_nvvm_sust_b_2d_array_i64_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i64_ty], [], "llvm.nvvm.sust.b.2d.array.i64.zero">, GCCBuiltin<"__nvvm_sust_b_2d_array_i64_zero">; def int_nvvm_sust_b_2d_array_v2i8_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.2d.array.v2i8.zero">, GCCBuiltin<"__nvvm_sust_b_2d_array_v2i8_zero">; def int_nvvm_sust_b_2d_array_v2i16_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.2d.array.v2i16.zero">, GCCBuiltin<"__nvvm_sust_b_2d_array_v2i16_zero">; def int_nvvm_sust_b_2d_array_v2i32_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.2d.array.v2i32.zero">, GCCBuiltin<"__nvvm_sust_b_2d_array_v2i32_zero">; def int_nvvm_sust_b_2d_array_v2i64_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i64_ty, llvm_i64_ty], [], "llvm.nvvm.sust.b.2d.array.v2i64.zero">, GCCBuiltin<"__nvvm_sust_b_2d_array_v2i64_zero">; def int_nvvm_sust_b_2d_array_v4i8_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.2d.array.v4i8.zero">, GCCBuiltin<"__nvvm_sust_b_2d_array_v4i8_zero">; def int_nvvm_sust_b_2d_array_v4i16_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.2d.array.v4i16.zero">, GCCBuiltin<"__nvvm_sust_b_2d_array_v4i16_zero">; def int_nvvm_sust_b_2d_array_v4i32_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.2d.array.v4i32.zero">, GCCBuiltin<"__nvvm_sust_b_2d_array_v4i32_zero">; def int_nvvm_sust_b_3d_i8_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.3d.i8.zero">, GCCBuiltin<"__nvvm_sust_b_3d_i8_zero">; def int_nvvm_sust_b_3d_i16_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.3d.i16.zero">, GCCBuiltin<"__nvvm_sust_b_3d_i16_zero">; def int_nvvm_sust_b_3d_i32_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.3d.i32.zero">, GCCBuiltin<"__nvvm_sust_b_3d_i32_zero">; def int_nvvm_sust_b_3d_i64_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i64_ty], [], "llvm.nvvm.sust.b.3d.i64.zero">, GCCBuiltin<"__nvvm_sust_b_3d_i64_zero">; def int_nvvm_sust_b_3d_v2i8_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.3d.v2i8.zero">, GCCBuiltin<"__nvvm_sust_b_3d_v2i8_zero">; def int_nvvm_sust_b_3d_v2i16_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.3d.v2i16.zero">, GCCBuiltin<"__nvvm_sust_b_3d_v2i16_zero">; def int_nvvm_sust_b_3d_v2i32_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.3d.v2i32.zero">, GCCBuiltin<"__nvvm_sust_b_3d_v2i32_zero">; def int_nvvm_sust_b_3d_v2i64_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i64_ty, llvm_i64_ty], [], "llvm.nvvm.sust.b.3d.v2i64.zero">, GCCBuiltin<"__nvvm_sust_b_3d_v2i64_zero">; def int_nvvm_sust_b_3d_v4i8_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.3d.v4i8.zero">, GCCBuiltin<"__nvvm_sust_b_3d_v4i8_zero">; def int_nvvm_sust_b_3d_v4i16_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.b.3d.v4i16.zero">, GCCBuiltin<"__nvvm_sust_b_3d_v4i16_zero">; def int_nvvm_sust_b_3d_v4i32_zero : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.b.3d.v4i32.zero">, GCCBuiltin<"__nvvm_sust_b_3d_v4i32_zero">; // Formatted def int_nvvm_sust_p_1d_i8_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.p.1d.i8.trap">, GCCBuiltin<"__nvvm_sust_p_1d_i8_trap">; def int_nvvm_sust_p_1d_i16_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.p.1d.i16.trap">, GCCBuiltin<"__nvvm_sust_p_1d_i16_trap">; def int_nvvm_sust_p_1d_i32_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.p.1d.i32.trap">, GCCBuiltin<"__nvvm_sust_p_1d_i32_trap">; def int_nvvm_sust_p_1d_v2i8_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.p.1d.v2i8.trap">, GCCBuiltin<"__nvvm_sust_p_1d_v2i8_trap">; def int_nvvm_sust_p_1d_v2i16_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.p.1d.v2i16.trap">, GCCBuiltin<"__nvvm_sust_p_1d_v2i16_trap">; def int_nvvm_sust_p_1d_v2i32_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.p.1d.v2i32.trap">, GCCBuiltin<"__nvvm_sust_p_1d_v2i32_trap">; def int_nvvm_sust_p_1d_v4i8_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.p.1d.v4i8.trap">, GCCBuiltin<"__nvvm_sust_p_1d_v4i8_trap">; def int_nvvm_sust_p_1d_v4i16_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.p.1d.v4i16.trap">, GCCBuiltin<"__nvvm_sust_p_1d_v4i16_trap">; def int_nvvm_sust_p_1d_v4i32_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.p.1d.v4i32.trap">, GCCBuiltin<"__nvvm_sust_p_1d_v4i32_trap">; def int_nvvm_sust_p_1d_array_i8_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.p.1d.array.i8.trap">, GCCBuiltin<"__nvvm_sust_p_1d_array_i8_trap">; def int_nvvm_sust_p_1d_array_i16_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.p.1d.array.i16.trap">, GCCBuiltin<"__nvvm_sust_p_1d_array_i16_trap">; def int_nvvm_sust_p_1d_array_i32_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.p.1d.array.i32.trap">, GCCBuiltin<"__nvvm_sust_p_1d_array_i32_trap">; def int_nvvm_sust_p_1d_array_v2i8_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.p.1d.array.v2i8.trap">, GCCBuiltin<"__nvvm_sust_p_1d_array_v2i8_trap">; def int_nvvm_sust_p_1d_array_v2i16_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.p.1d.array.v2i16.trap">, GCCBuiltin<"__nvvm_sust_p_1d_array_v2i16_trap">; def int_nvvm_sust_p_1d_array_v2i32_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.p.1d.array.v2i32.trap">, GCCBuiltin<"__nvvm_sust_p_1d_array_v2i32_trap">; def int_nvvm_sust_p_1d_array_v4i8_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.p.1d.array.v4i8.trap">, GCCBuiltin<"__nvvm_sust_p_1d_array_v4i8_trap">; def int_nvvm_sust_p_1d_array_v4i16_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.p.1d.array.v4i16.trap">, GCCBuiltin<"__nvvm_sust_p_1d_array_v4i16_trap">; def int_nvvm_sust_p_1d_array_v4i32_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.p.1d.array.v4i32.trap">, GCCBuiltin<"__nvvm_sust_p_1d_array_v4i32_trap">; def int_nvvm_sust_p_2d_i8_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.p.2d.i8.trap">, GCCBuiltin<"__nvvm_sust_p_2d_i8_trap">; def int_nvvm_sust_p_2d_i16_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.p.2d.i16.trap">, GCCBuiltin<"__nvvm_sust_p_2d_i16_trap">; def int_nvvm_sust_p_2d_i32_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.p.2d.i32.trap">, GCCBuiltin<"__nvvm_sust_p_2d_i32_trap">; def int_nvvm_sust_p_2d_v2i8_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.p.2d.v2i8.trap">, GCCBuiltin<"__nvvm_sust_p_2d_v2i8_trap">; def int_nvvm_sust_p_2d_v2i16_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.p.2d.v2i16.trap">, GCCBuiltin<"__nvvm_sust_p_2d_v2i16_trap">; def int_nvvm_sust_p_2d_v2i32_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.p.2d.v2i32.trap">, GCCBuiltin<"__nvvm_sust_p_2d_v2i32_trap">; def int_nvvm_sust_p_2d_v4i8_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.p.2d.v4i8.trap">, GCCBuiltin<"__nvvm_sust_p_2d_v4i8_trap">; def int_nvvm_sust_p_2d_v4i16_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.p.2d.v4i16.trap">, GCCBuiltin<"__nvvm_sust_p_2d_v4i16_trap">; def int_nvvm_sust_p_2d_v4i32_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.p.2d.v4i32.trap">, GCCBuiltin<"__nvvm_sust_p_2d_v4i32_trap">; def int_nvvm_sust_p_2d_array_i8_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.p.2d.array.i8.trap">, GCCBuiltin<"__nvvm_sust_p_2d_array_i8_trap">; def int_nvvm_sust_p_2d_array_i16_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.p.2d.array.i16.trap">, GCCBuiltin<"__nvvm_sust_p_2d_array_i16_trap">; def int_nvvm_sust_p_2d_array_i32_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.p.2d.array.i32.trap">, GCCBuiltin<"__nvvm_sust_p_2d_array_i32_trap">; def int_nvvm_sust_p_2d_array_v2i8_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.p.2d.array.v2i8.trap">, GCCBuiltin<"__nvvm_sust_p_2d_array_v2i8_trap">; def int_nvvm_sust_p_2d_array_v2i16_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.p.2d.array.v2i16.trap">, GCCBuiltin<"__nvvm_sust_p_2d_array_v2i16_trap">; def int_nvvm_sust_p_2d_array_v2i32_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.p.2d.array.v2i32.trap">, GCCBuiltin<"__nvvm_sust_p_2d_array_v2i32_trap">; def int_nvvm_sust_p_2d_array_v4i8_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.p.2d.array.v4i8.trap">, GCCBuiltin<"__nvvm_sust_p_2d_array_v4i8_trap">; def int_nvvm_sust_p_2d_array_v4i16_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.p.2d.array.v4i16.trap">, GCCBuiltin<"__nvvm_sust_p_2d_array_v4i16_trap">; def int_nvvm_sust_p_2d_array_v4i32_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.p.2d.array.v4i32.trap">, GCCBuiltin<"__nvvm_sust_p_2d_array_v4i32_trap">; def int_nvvm_sust_p_3d_i8_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.p.3d.i8.trap">, GCCBuiltin<"__nvvm_sust_p_3d_i8_trap">; def int_nvvm_sust_p_3d_i16_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [], "llvm.nvvm.sust.p.3d.i16.trap">, GCCBuiltin<"__nvvm_sust_p_3d_i16_trap">; def int_nvvm_sust_p_3d_i32_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.p.3d.i32.trap">, GCCBuiltin<"__nvvm_sust_p_3d_i32_trap">; def int_nvvm_sust_p_3d_v2i8_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.p.3d.v2i8.trap">, GCCBuiltin<"__nvvm_sust_p_3d_v2i8_trap">; def int_nvvm_sust_p_3d_v2i16_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.p.3d.v2i16.trap">, GCCBuiltin<"__nvvm_sust_p_3d_v2i16_trap">; def int_nvvm_sust_p_3d_v2i32_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.p.3d.v2i32.trap">, GCCBuiltin<"__nvvm_sust_p_3d_v2i32_trap">; def int_nvvm_sust_p_3d_v4i8_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.p.3d.v4i8.trap">, GCCBuiltin<"__nvvm_sust_p_3d_v4i8_trap">; def int_nvvm_sust_p_3d_v4i16_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [], "llvm.nvvm.sust.p.3d.v4i16.trap">, GCCBuiltin<"__nvvm_sust_p_3d_v4i16_trap">; def int_nvvm_sust_p_3d_v4i32_trap : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [], "llvm.nvvm.sust.p.3d.v4i32.trap">, GCCBuiltin<"__nvvm_sust_p_3d_v4i32_trap">; def int_nvvm_rotate_b32 : Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem], "llvm.nvvm.rotate.b32">, GCCBuiltin<"__nvvm_rotate_b32">; def int_nvvm_rotate_b64 :Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty], [IntrNoMem], "llvm.nvvm.rotate.b64">, GCCBuiltin<"__nvvm_rotate_b64">; def int_nvvm_rotate_right_b64 : Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty], [IntrNoMem], "llvm.nvvm.rotate.right.b64">, GCCBuiltin<"__nvvm_rotate_right_b64">; def int_nvvm_swap_lo_hi_b64 : Intrinsic<[llvm_i64_ty], [llvm_i64_ty], [IntrNoMem], "llvm.nvvm.swap.lo.hi.b64">, GCCBuiltin<"__nvvm_swap_lo_hi_b64">; // Accessing special registers. multiclass PTXReadSRegIntrinsic_v4i32 { // FIXME: Do we need the 128-bit integer type version? // def _r64 : Intrinsic<[llvm_i128_ty], [], [IntrNoMem]>; // FIXME: Enable this once v4i32 support is enabled in back-end. // def _v4i16 : Intrinsic<[llvm_v4i32_ty], [], [IntrNoMem]>; def _x : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>, GCCBuiltin<"__nvvm_read_ptx_sreg_" # regname # "_x">; def _y : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>, GCCBuiltin<"__nvvm_read_ptx_sreg_" # regname # "_y">; def _z : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>, GCCBuiltin<"__nvvm_read_ptx_sreg_" # regname # "_z">; def _w : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>, GCCBuiltin<"__nvvm_read_ptx_sreg_" # regname # "_w">; } class PTXReadSRegIntrinsic_r32 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>, GCCBuiltin<"__nvvm_read_ptx_sreg_" # name>; class PTXReadSRegIntrinsic_r64 : Intrinsic<[llvm_i64_ty], [], [IntrNoMem]>, GCCBuiltin<"__nvvm_read_ptx_sreg_" # name>; // Intrinsics to read registers with non-constant values. E.g. the values that // do change over the kernel lifetime. Such reads should not be CSE'd. class PTXReadNCSRegIntrinsic_r32 : Intrinsic<[llvm_i32_ty], [], [IntrInaccessibleMemOnly]>, GCCBuiltin<"__nvvm_read_ptx_sreg_" # name>; class PTXReadNCSRegIntrinsic_r64 : Intrinsic<[llvm_i64_ty], [], [IntrInaccessibleMemOnly]>, GCCBuiltin<"__nvvm_read_ptx_sreg_" # name>; defm int_nvvm_read_ptx_sreg_tid : PTXReadSRegIntrinsic_v4i32<"tid">; defm int_nvvm_read_ptx_sreg_ntid : PTXReadSRegIntrinsic_v4i32<"ntid">; def int_nvvm_read_ptx_sreg_laneid : PTXReadSRegIntrinsic_r32<"laneid">; def int_nvvm_read_ptx_sreg_warpid : PTXReadSRegIntrinsic_r32<"warpid">; def int_nvvm_read_ptx_sreg_nwarpid : PTXReadSRegIntrinsic_r32<"nwarpid">; defm int_nvvm_read_ptx_sreg_ctaid : PTXReadSRegIntrinsic_v4i32<"ctaid">; defm int_nvvm_read_ptx_sreg_nctaid : PTXReadSRegIntrinsic_v4i32<"nctaid">; def int_nvvm_read_ptx_sreg_smid : PTXReadSRegIntrinsic_r32<"smid">; def int_nvvm_read_ptx_sreg_nsmid : PTXReadSRegIntrinsic_r32<"nsmid">; def int_nvvm_read_ptx_sreg_gridid : PTXReadSRegIntrinsic_r32<"gridid">; def int_nvvm_read_ptx_sreg_lanemask_eq : PTXReadSRegIntrinsic_r32<"lanemask_eq">; def int_nvvm_read_ptx_sreg_lanemask_le : PTXReadSRegIntrinsic_r32<"lanemask_le">; def int_nvvm_read_ptx_sreg_lanemask_lt : PTXReadSRegIntrinsic_r32<"lanemask_lt">; def int_nvvm_read_ptx_sreg_lanemask_ge : PTXReadSRegIntrinsic_r32<"lanemask_ge">; def int_nvvm_read_ptx_sreg_lanemask_gt : PTXReadSRegIntrinsic_r32<"lanemask_gt">; def int_nvvm_read_ptx_sreg_clock : PTXReadNCSRegIntrinsic_r32<"clock">; def int_nvvm_read_ptx_sreg_clock64 : PTXReadNCSRegIntrinsic_r64<"clock64">; def int_nvvm_read_ptx_sreg_pm0 : PTXReadNCSRegIntrinsic_r32<"pm0">; def int_nvvm_read_ptx_sreg_pm1 : PTXReadNCSRegIntrinsic_r32<"pm1">; def int_nvvm_read_ptx_sreg_pm2 : PTXReadNCSRegIntrinsic_r32<"pm2">; def int_nvvm_read_ptx_sreg_pm3 : PTXReadNCSRegIntrinsic_r32<"pm3">; def int_nvvm_read_ptx_sreg_warpsize : PTXReadSRegIntrinsic_r32<"warpsize">; // // SHUFFLE // // Generate intrinsics for all variants of shfl instruction. foreach sync = [false, true] in { foreach mode = ["up", "down", "bfly", "idx"] in { foreach type = ["i32", "f32"] in { foreach return_pred = [false, true] in { foreach i = [SHFL_INFO] in { if i.withGccBuiltin then { def i.Name : GCCBuiltin, Intrinsic; } if i.withoutGccBuiltin then { def i.Name : Intrinsic; } } } } } } // // VOTE // // vote.all pred def int_nvvm_vote_all : Intrinsic<[llvm_i1_ty], [llvm_i1_ty], [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.vote.all">, GCCBuiltin<"__nvvm_vote_all">; // vote.any pred def int_nvvm_vote_any : Intrinsic<[llvm_i1_ty], [llvm_i1_ty], [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.vote.any">, GCCBuiltin<"__nvvm_vote_any">; // vote.uni pred def int_nvvm_vote_uni : Intrinsic<[llvm_i1_ty], [llvm_i1_ty], [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.vote.uni">, GCCBuiltin<"__nvvm_vote_uni">; // vote.ballot pred def int_nvvm_vote_ballot : Intrinsic<[llvm_i32_ty], [llvm_i1_ty], [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.vote.ballot">, GCCBuiltin<"__nvvm_vote_ballot">; // // VOTE.SYNC // // vote.sync.all mask, pred def int_nvvm_vote_all_sync : Intrinsic<[llvm_i1_ty], [llvm_i32_ty, llvm_i1_ty], [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.vote.all.sync">, GCCBuiltin<"__nvvm_vote_all_sync">; // vote.sync.any mask, pred def int_nvvm_vote_any_sync : Intrinsic<[llvm_i1_ty], [llvm_i32_ty, llvm_i1_ty], [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.vote.any.sync">, GCCBuiltin<"__nvvm_vote_any_sync">; // vote.sync.uni mask, pred def int_nvvm_vote_uni_sync : Intrinsic<[llvm_i1_ty], [llvm_i32_ty, llvm_i1_ty], [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.vote.uni.sync">, GCCBuiltin<"__nvvm_vote_uni_sync">; // vote.sync.ballot mask, pred def int_nvvm_vote_ballot_sync : Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i1_ty], [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.vote.ballot.sync">, GCCBuiltin<"__nvvm_vote_ballot_sync">; // // MATCH.SYNC // // match.any.sync.b32 mask, value def int_nvvm_match_any_sync_i32 : Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.match.any.sync.i32">, GCCBuiltin<"__nvvm_match_any_sync_i32">; // match.any.sync.b64 mask, value def int_nvvm_match_any_sync_i64 : Intrinsic<[llvm_i64_ty], [llvm_i32_ty, llvm_i64_ty], [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.match.any.sync.i64">, GCCBuiltin<"__nvvm_match_any_sync_i64">; // match.all instruction have two variants -- one returns a single value, another // returns a pair {value, predicate}. We currently only implement the latter as // that's the variant exposed by CUDA API. // match.all.sync.b32p mask, value def int_nvvm_match_all_sync_i32p : Intrinsic<[llvm_i32_ty, llvm_i1_ty], [llvm_i32_ty, llvm_i32_ty], [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.match.all.sync.i32p">; // match.all.sync.b64p mask, value def int_nvvm_match_all_sync_i64p : Intrinsic<[llvm_i64_ty, llvm_i1_ty], [llvm_i32_ty, llvm_i64_ty], [IntrInaccessibleMemOnly, IntrConvergent], "llvm.nvvm.match.all.sync.i64p">; // // WMMA instructions // // WMMA.LOAD class NVVM_WMMA_LD : Intrinsic>, NoCapture>], WMMA_NAME_LDST<"load", Frag, Layout, WithStride>.intr>; // WMMA.STORE.D class NVVM_WMMA_ST : Intrinsic<[], !listconcat( [llvm_anyptr_ty], Frag.regs, !if(WithStride, [llvm_i32_ty], [])), [IntrWriteMem, IntrArgMemOnly, WriteOnly>, NoCapture>], WMMA_NAME_LDST<"store", Frag, Layout, WithStride>.intr>; // Create all load/store variants foreach layout = ["row", "col"] in { foreach stride = [0, 1] in { foreach frag = NVVM_MMA_OPS.all_ld_ops in if NVVM_MMA_SUPPORTED<[frag], layout>.ret then def WMMA_NAME_LDST<"load", frag, layout, stride>.record : NVVM_WMMA_LD; foreach frag = NVVM_MMA_OPS.all_st_ops in if NVVM_MMA_SUPPORTED<[frag], layout>.ret then def WMMA_NAME_LDST<"store", frag, layout, stride>.record : NVVM_WMMA_ST; } } // WMMA.MMA class NVVM_WMMA_MMA : Intrinsic.llvm>; foreach layout_a = ["row", "col"] in { foreach layout_b = ["row", "col"] in { foreach satf = [0, 1] in { foreach op = NVVM_MMA_OPS.all_mma_ops in { if NVVM_MMA_SUPPORTED.ret then { def WMMA_NAME_MMA.record : NVVM_WMMA_MMA; } } } // satf } // layout_b } // layout_a } // let TargetPrefix = "nvvm"