| //===- 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<llvm_i64_ty>; // (space)i64* |
| |
| // |
| // MISC |
| // |
| |
| // Helper class for construction of n-element list<LLVMtype> [t,t,...,t] |
| class RepLLVMType<int N, LLVMType T> { |
| list<LLVMType> ret = !if(N, !listconcat(RepLLVMType<!add(N,-1), T>.ret, [T]), []); |
| } |
| |
| // Helper class that represents a 'fragment' of an NVPTX *MMA instruction. |
| // Geom: m<M>n<N>k<K>. E.g. m8n32k16 |
| // Frag: [abcd] |
| // PtxEltType: PTX type for the element. |
| class WMMA_REGS<string Geom, string Frag, string PtxEltType> { |
| 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<LLVMType> regs = !cond( |
| // 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") : RepLLVMType<8, llvm_v2f16_ty>.ret, |
| !eq(ft,"b:f16") : RepLLVMType<8, llvm_v2f16_ty>.ret, |
| !eq(ft,"c:f16") : RepLLVMType<4, llvm_v2f16_ty>.ret, |
| !eq(ft,"d:f16") : RepLLVMType<4, llvm_v2f16_ty>.ret, |
| !eq(ft,"c:f32") : RepLLVMType<8, llvm_float_ty>.ret, |
| !eq(ft,"d:f32") : RepLLVMType<8, llvm_float_ty>.ret, |
| |
| // u8/s8 -> s32 @ m16n16k16/m8n32k16/m32n8k16 |
| !eq(gft,"m16n16k16:a:u8") : RepLLVMType<2, llvm_i32_ty>.ret, |
| !eq(gft,"m16n16k16:a:s8") : RepLLVMType<2, llvm_i32_ty>.ret, |
| !eq(gft,"m16n16k16:b:u8") : RepLLVMType<2, llvm_i32_ty>.ret, |
| !eq(gft,"m16n16k16:b:s8") : RepLLVMType<2, llvm_i32_ty>.ret, |
| !eq(gft,"m16n16k16:c:s32") : RepLLVMType<8, llvm_i32_ty>.ret, |
| !eq(gft,"m16n16k16:d:s32") : RepLLVMType<8, llvm_i32_ty>.ret, |
| |
| !eq(gft,"m8n32k16:a:u8") : [llvm_i32_ty], |
| !eq(gft,"m8n32k16:a:s8") : [llvm_i32_ty], |
| !eq(gft,"m8n32k16:b:u8") : RepLLVMType<4, llvm_i32_ty>.ret, |
| !eq(gft,"m8n32k16:b:s8") : RepLLVMType<4, llvm_i32_ty>.ret, |
| !eq(gft,"m8n32k16:c:s32") : RepLLVMType<8, llvm_i32_ty>.ret, |
| !eq(gft,"m8n32k16:d:s32") : RepLLVMType<8, llvm_i32_ty>.ret, |
| |
| !eq(gft,"m32n8k16:a:u8") : RepLLVMType<4, llvm_i32_ty>.ret, |
| !eq(gft,"m32n8k16:a:s8") : RepLLVMType<4, llvm_i32_ty>.ret, |
| !eq(gft,"m32n8k16:b:u8") : [llvm_i32_ty], |
| !eq(gft,"m32n8k16:b:s8") : [llvm_i32_ty], |
| !eq(gft,"m32n8k16:c:s32") : RepLLVMType<8, llvm_i32_ty>.ret, |
| !eq(gft,"m32n8k16:d:s32") : RepLLVMType<8, llvm_i32_ty>.ret, |
| |
| // 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") : RepLLVMType<2, llvm_i32_ty>.ret, |
| !eq(gft,"m8n8k128:d:s32") : RepLLVMType<2, llvm_i32_ty>.ret, |
| !eq(gft,"m8n8k32:c:s32") : RepLLVMType<2, llvm_i32_ty>.ret, |
| !eq(gft,"m8n8k32:d:s32") : RepLLVMType<2, llvm_i32_ty>.ret, |
| ); |
| } |
| |
| class WMMA_NAME_LDST<string Op, WMMA_REGS Frag, string Layout, int WithStride> { |
| 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<WMMA_REGS A, WMMA_REGS B, WMMA_REGS C, WMMA_REGS D> { |
| list<WMMA_REGS> 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. |
| 1: [D, C] |
| ); |
| string ret = !foldl("", id_frags, a, b, !strconcat(a, ".", b.ptx_elt_type)); |
| } |
| |
| class WMMA_NAME_MMA<string ALayout, string BLayout, int Satfinite, |
| WMMA_REGS A, WMMA_REGS B, WMMA_REGS C, WMMA_REGS D> { |
| string signature = MMA_SIGNATURE<A, B, C, D>.ret; |
| string llvm = "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<list<string> Geom, list<string> TypeA, list<string> TypeB, |
| list<string> TypeC, list<string> TypeD> { |
| list<list<WMMA_REGS>> ret = |
| !foldl([]<list<WMMA_REGS>>, Geom, t1, geom, !listconcat(t1, |
| !foldl([]<list<WMMA_REGS>>, TypeA, t2, type_a, !listconcat(t2, |
| !foldl([]<list<WMMA_REGS>>, !if(!size(TypeB), TypeB, [type_a]), t3, type_b, !listconcat(t3, |
| !foldl([]<list<WMMA_REGS>>, TypeC, t4, type_c, !listconcat(t4, |
| !foldl([]<list<WMMA_REGS>>, !if(!size(TypeC), TypeC, [type_c]), t5, type_d, !listconcat(t5, |
| [[WMMA_REGS<geom, "a", type_a>, |
| WMMA_REGS<geom, "b", type_b>, |
| WMMA_REGS<geom, "c", type_c>, |
| WMMA_REGS<geom, "d", type_d>]])))))))))); |
| // Debugging aid for readable representation of the list above. |
| list<list<string>> ops = !foreach(x, ret, [x[0].gft, x[1].gft, x[2].gft, x[3].gft]); |
| } |
| |
| class MMA_LDST_OPS<list<string> Geom, list<string> Frags, list<string> Types> { |
| list<WMMA_REGS> ret = |
| !foldl([]<WMMA_REGS>, Geom, t1, geom, !listconcat(t1, |
| !foldl([]<WMMA_REGS>, Frags, t2, frag, !listconcat(t2, |
| !foldl([]<WMMA_REGS>, Types, t3, type, !listconcat(t3, |
| [WMMA_REGS<geom, frag, type>])))))); |
| // Debugging aid for readable representation of the list above. |
| list<string> 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<int _ = 0> { |
| list<list<WMMA_REGS>> fp_mma_ops = MMA_OPS< |
| ["m16n16k16", "m32n8k16", "m8n32k16"], |
| ["f16"], [], ["f16", "f32"], ["f16", "f32"]>.ret; |
| list<list<WMMA_REGS>> int_mma_ops = MMA_OPS< |
| ["m16n16k16", "m32n8k16", "m8n32k16"], |
| ["s8", "u8"], [], ["s32"], []>.ret; |
| list<list<WMMA_REGS>> subint_mma_ops = MMA_OPS< |
| ["m8n8k32"], |
| ["s4", "u4"], [], ["s32"], []>.ret; |
| list<list<WMMA_REGS>> bit_mma_ops = MMA_OPS< |
| ["m8n8k128"], |
| ["b1"], [], ["s32"], []>.ret; |
| list<list<WMMA_REGS>> all_mma_ops = !listconcat(fp_mma_ops, int_mma_ops, |
| subint_mma_ops, bit_mma_ops); |
| |
| list<WMMA_REGS> ldst_ab_ops = MMA_LDST_OPS< |
| ["m16n16k16", "m32n8k16", "m8n32k16"], |
| ["a", "b"], ["f16", "u8", "s8"]>.ret; |
| list<WMMA_REGS> ldst_cd_ops = MMA_LDST_OPS< |
| ["m16n16k16", "m32n8k16", "m8n32k16"], |
| ["c", "d"], ["f16", "f32", "s32"]>.ret; |
| list<WMMA_REGS> ldst_subint_ab_ops = MMA_LDST_OPS< |
| ["m8n8k32"], ["a", "b"], ["s4","u4"]>.ret; |
| list<WMMA_REGS> ldst_bit_ab_ops = MMA_LDST_OPS< |
| ["m8n8k128"], ["a", "b"], ["b1"]>.ret; |
| list<WMMA_REGS> ldst_subint_cd_ops = MMA_LDST_OPS< |
| ["m8n8k32", "m8n8k128"], ["c", "d"], ["s32"]>.ret; |
| list<WMMA_REGS> 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<WMMA_REGS> all_ld_ops = !foldl([]<WMMA_REGS>, all_ldst_ops, a, b, |
| !listconcat(a, !if(!eq(b.frag,"d"), [],[b]))); |
| list<WMMA_REGS> all_st_ops = !foldl([]<WMMA_REGS>, all_ldst_ops, a, b, |
| !listconcat(a, !if(!eq(b.frag,"d"), [b],[]))); |
| } |
| |
| def NVVM_MMA_OPS : NVVM_MMA_OPS; |
| |
| // Returns [1] if this combination of layout/satf is supported, [] 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. |
| // foreach _ = NVVM_MMA_SUPPORTED<...>.ret in = |
| // def : FOO<>; // The record will only be defined for supported ops. |
| // |
| class NVVM_MMA_SUPPORTED<list<WMMA_REGS> 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; |
| list<int> ret = !cond( |
| // Sub-int MMA only supports fixed A/B layout. |
| // b1 does not support .satf. |
| !eq(mma#":"#satf, "b1:row:col:0") : [1], |
| !eq(mma, "s4:row:col") : [1], |
| !eq(mma, "u4:row:col") : [1], |
| !eq(mma, "s4:row:col") : [1], |
| !eq(mma, "u4:row:col") : [1], |
| // 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"))) : [1], |
| // All other sub-int ops are not supported. |
| !eq(t, "b1") : [], |
| !eq(t, "s4") : [], |
| !eq(t, "u4") : [], |
| // All other (non sub-int) are OK. |
| 1: [1] |
| ); |
| } |
| |
| 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_add_f32 : Intrinsic<[llvm_float_ty], |
| [LLVMAnyPointerType<llvm_float_ty>, llvm_float_ty], |
| [IntrArgMemOnly, NoCapture<0>]>; |
| // Atomic add of f64 requires sm_60. |
| def int_nvvm_atomic_load_add_f64 : Intrinsic<[llvm_double_ty], |
| [LLVMAnyPointerType<llvm_double_ty>, llvm_double_ty], |
| [IntrArgMemOnly, NoCapture<0>]>; |
| |
| def int_nvvm_atomic_load_inc_32 : Intrinsic<[llvm_i32_ty], |
| [LLVMAnyPointerType<llvm_i32_ty>, llvm_i32_ty], |
| [IntrArgMemOnly, NoCapture<0>]>; |
| def int_nvvm_atomic_load_dec_32 : Intrinsic<[llvm_i32_ty], |
| [LLVMAnyPointerType<llvm_i32_ty>, llvm_i32_ty], |
| [IntrArgMemOnly, NoCapture<0>]>; |
| |
| class SCOPED_ATOMIC2_impl<LLVMType elty> |
| : Intrinsic<[elty], |
| [LLVMAnyPointerType<LLVMMatchType<0>>, LLVMMatchType<0>], |
| [IntrArgMemOnly, NoCapture<0>]>; |
| class SCOPED_ATOMIC3_impl<LLVMType elty> |
| : Intrinsic<[elty], |
| [LLVMAnyPointerType<LLVMMatchType<0>>, LLVMMatchType<0>, |
| LLVMMatchType<0>], |
| [IntrArgMemOnly, NoCapture<0>]>; |
| |
| multiclass PTXAtomicWithScope2<LLVMType elty> { |
| def _cta : SCOPED_ATOMIC2_impl<elty>; |
| def _sys : SCOPED_ATOMIC2_impl<elty>; |
| } |
| multiclass PTXAtomicWithScope3<LLVMType elty> { |
| def _cta : SCOPED_ATOMIC3_impl<elty>; |
| def _sys : SCOPED_ATOMIC3_impl<elty>; |
| } |
| multiclass PTXAtomicWithScope2_fi { |
| defm _f: PTXAtomicWithScope2<llvm_anyfloat_ty>; |
| defm _i: PTXAtomicWithScope2<llvm_anyint_ty>; |
| } |
| defm int_nvvm_atomic_add_gen : PTXAtomicWithScope2_fi; |
| defm int_nvvm_atomic_inc_gen_i : PTXAtomicWithScope2<llvm_anyint_ty>; |
| defm int_nvvm_atomic_dec_gen_i : PTXAtomicWithScope2<llvm_anyint_ty>; |
| defm int_nvvm_atomic_exch_gen_i: PTXAtomicWithScope2<llvm_anyint_ty>; |
| defm int_nvvm_atomic_xor_gen_i : PTXAtomicWithScope2<llvm_anyint_ty>; |
| defm int_nvvm_atomic_max_gen_i : PTXAtomicWithScope2<llvm_anyint_ty>; |
| defm int_nvvm_atomic_min_gen_i : PTXAtomicWithScope2<llvm_anyint_ty>; |
| defm int_nvvm_atomic_or_gen_i : PTXAtomicWithScope2<llvm_anyint_ty>; |
| defm int_nvvm_atomic_and_gen_i : PTXAtomicWithScope2<llvm_anyint_ty>; |
| defm int_nvvm_atomic_cas_gen_i : PTXAtomicWithScope3<llvm_anyint_ty>; |
| |
| // 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<LLVMMatchType<0>>, llvm_i32_ty], |
| [IntrReadMem, IntrArgMemOnly, NoCapture<0>], |
| "llvm.nvvm.ldu.global.i">; |
| def int_nvvm_ldu_global_f : Intrinsic<[llvm_anyfloat_ty], |
| [LLVMAnyPointerType<LLVMMatchType<0>>, llvm_i32_ty], |
| [IntrReadMem, IntrArgMemOnly, NoCapture<0>], |
| "llvm.nvvm.ldu.global.f">; |
| def int_nvvm_ldu_global_p : Intrinsic<[llvm_anyptr_ty], |
| [LLVMAnyPointerType<LLVMMatchType<0>>, llvm_i32_ty], |
| [IntrReadMem, IntrArgMemOnly, NoCapture<0>], |
| "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<LLVMMatchType<0>>, llvm_i32_ty], |
| [IntrReadMem, IntrArgMemOnly, NoCapture<0>], |
| "llvm.nvvm.ldg.global.i">; |
| def int_nvvm_ldg_global_f : Intrinsic<[llvm_anyfloat_ty], |
| [LLVMAnyPointerType<LLVMMatchType<0>>, llvm_i32_ty], |
| [IntrReadMem, IntrArgMemOnly, NoCapture<0>], |
| "llvm.nvvm.ldg.global.f">; |
| def int_nvvm_ldg_global_p : Intrinsic<[llvm_anyptr_ty], |
| [LLVMAnyPointerType<LLVMMatchType<0>>, llvm_i32_ty], |
| [IntrReadMem, IntrArgMemOnly, NoCapture<0>], |
| "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<0>], "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], |
| |