Lines Matching refs:Intrinsic
1255 auto MakeIntrinsicCall = [&](Intrinsic::ID IID) { in getSqrtEstimate()
1266 return MakeIntrinsicCall(Ftz ? Intrinsic::nvvm_rsqrt_approx_ftz_f in getSqrtEstimate()
1267 : Intrinsic::nvvm_rsqrt_approx_f); in getSqrtEstimate()
1269 return MakeIntrinsicCall(Intrinsic::nvvm_rsqrt_approx_d); in getSqrtEstimate()
1274 return MakeIntrinsicCall(Ftz ? Intrinsic::nvvm_sqrt_approx_ftz_f in getSqrtEstimate()
1275 : Intrinsic::nvvm_sqrt_approx_f); in getSqrtEstimate()
1283 DAG.getConstant(Intrinsic::nvvm_rcp_approx_ftz_d, DL, MVT::i32), in getSqrtEstimate()
1284 MakeIntrinsicCall(Intrinsic::nvvm_rsqrt_approx_d)); in getSqrtEstimate()
2783 static unsigned getOpcForTextureInstr(unsigned Intrinsic) { in getOpcForTextureInstr() argument
2784 switch (Intrinsic) { in getOpcForTextureInstr()
2788 case Intrinsic::nvvm_tex_1d_v4f32_s32: in getOpcForTextureInstr()
2790 case Intrinsic::nvvm_tex_1d_v4f32_f32: in getOpcForTextureInstr()
2792 case Intrinsic::nvvm_tex_1d_level_v4f32_f32: in getOpcForTextureInstr()
2794 case Intrinsic::nvvm_tex_1d_grad_v4f32_f32: in getOpcForTextureInstr()
2796 case Intrinsic::nvvm_tex_1d_v4s32_s32: in getOpcForTextureInstr()
2798 case Intrinsic::nvvm_tex_1d_v4s32_f32: in getOpcForTextureInstr()
2800 case Intrinsic::nvvm_tex_1d_level_v4s32_f32: in getOpcForTextureInstr()
2802 case Intrinsic::nvvm_tex_1d_grad_v4s32_f32: in getOpcForTextureInstr()
2804 case Intrinsic::nvvm_tex_1d_v4u32_s32: in getOpcForTextureInstr()
2806 case Intrinsic::nvvm_tex_1d_v4u32_f32: in getOpcForTextureInstr()
2808 case Intrinsic::nvvm_tex_1d_level_v4u32_f32: in getOpcForTextureInstr()
2810 case Intrinsic::nvvm_tex_1d_grad_v4u32_f32: in getOpcForTextureInstr()
2813 case Intrinsic::nvvm_tex_1d_array_v4f32_s32: in getOpcForTextureInstr()
2815 case Intrinsic::nvvm_tex_1d_array_v4f32_f32: in getOpcForTextureInstr()
2817 case Intrinsic::nvvm_tex_1d_array_level_v4f32_f32: in getOpcForTextureInstr()
2819 case Intrinsic::nvvm_tex_1d_array_grad_v4f32_f32: in getOpcForTextureInstr()
2821 case Intrinsic::nvvm_tex_1d_array_v4s32_s32: in getOpcForTextureInstr()
2823 case Intrinsic::nvvm_tex_1d_array_v4s32_f32: in getOpcForTextureInstr()
2825 case Intrinsic::nvvm_tex_1d_array_level_v4s32_f32: in getOpcForTextureInstr()
2827 case Intrinsic::nvvm_tex_1d_array_grad_v4s32_f32: in getOpcForTextureInstr()
2829 case Intrinsic::nvvm_tex_1d_array_v4u32_s32: in getOpcForTextureInstr()
2831 case Intrinsic::nvvm_tex_1d_array_v4u32_f32: in getOpcForTextureInstr()
2833 case Intrinsic::nvvm_tex_1d_array_level_v4u32_f32: in getOpcForTextureInstr()
2835 case Intrinsic::nvvm_tex_1d_array_grad_v4u32_f32: in getOpcForTextureInstr()
2838 case Intrinsic::nvvm_tex_2d_v4f32_s32: in getOpcForTextureInstr()
2840 case Intrinsic::nvvm_tex_2d_v4f32_f32: in getOpcForTextureInstr()
2842 case Intrinsic::nvvm_tex_2d_level_v4f32_f32: in getOpcForTextureInstr()
2844 case Intrinsic::nvvm_tex_2d_grad_v4f32_f32: in getOpcForTextureInstr()
2846 case Intrinsic::nvvm_tex_2d_v4s32_s32: in getOpcForTextureInstr()
2848 case Intrinsic::nvvm_tex_2d_v4s32_f32: in getOpcForTextureInstr()
2850 case Intrinsic::nvvm_tex_2d_level_v4s32_f32: in getOpcForTextureInstr()
2852 case Intrinsic::nvvm_tex_2d_grad_v4s32_f32: in getOpcForTextureInstr()
2854 case Intrinsic::nvvm_tex_2d_v4u32_s32: in getOpcForTextureInstr()
2856 case Intrinsic::nvvm_tex_2d_v4u32_f32: in getOpcForTextureInstr()
2858 case Intrinsic::nvvm_tex_2d_level_v4u32_f32: in getOpcForTextureInstr()
2860 case Intrinsic::nvvm_tex_2d_grad_v4u32_f32: in getOpcForTextureInstr()
2863 case Intrinsic::nvvm_tex_2d_array_v4f32_s32: in getOpcForTextureInstr()
2865 case Intrinsic::nvvm_tex_2d_array_v4f32_f32: in getOpcForTextureInstr()
2867 case Intrinsic::nvvm_tex_2d_array_level_v4f32_f32: in getOpcForTextureInstr()
2869 case Intrinsic::nvvm_tex_2d_array_grad_v4f32_f32: in getOpcForTextureInstr()
2871 case Intrinsic::nvvm_tex_2d_array_v4s32_s32: in getOpcForTextureInstr()
2873 case Intrinsic::nvvm_tex_2d_array_v4s32_f32: in getOpcForTextureInstr()
2875 case Intrinsic::nvvm_tex_2d_array_level_v4s32_f32: in getOpcForTextureInstr()
2877 case Intrinsic::nvvm_tex_2d_array_grad_v4s32_f32: in getOpcForTextureInstr()
2879 case Intrinsic::nvvm_tex_2d_array_v4u32_s32: in getOpcForTextureInstr()
2881 case Intrinsic::nvvm_tex_2d_array_v4u32_f32: in getOpcForTextureInstr()
2883 case Intrinsic::nvvm_tex_2d_array_level_v4u32_f32: in getOpcForTextureInstr()
2885 case Intrinsic::nvvm_tex_2d_array_grad_v4u32_f32: in getOpcForTextureInstr()
2888 case Intrinsic::nvvm_tex_3d_v4f32_s32: in getOpcForTextureInstr()
2890 case Intrinsic::nvvm_tex_3d_v4f32_f32: in getOpcForTextureInstr()
2892 case Intrinsic::nvvm_tex_3d_level_v4f32_f32: in getOpcForTextureInstr()
2894 case Intrinsic::nvvm_tex_3d_grad_v4f32_f32: in getOpcForTextureInstr()
2896 case Intrinsic::nvvm_tex_3d_v4s32_s32: in getOpcForTextureInstr()
2898 case Intrinsic::nvvm_tex_3d_v4s32_f32: in getOpcForTextureInstr()
2900 case Intrinsic::nvvm_tex_3d_level_v4s32_f32: in getOpcForTextureInstr()
2902 case Intrinsic::nvvm_tex_3d_grad_v4s32_f32: in getOpcForTextureInstr()
2904 case Intrinsic::nvvm_tex_3d_v4u32_s32: in getOpcForTextureInstr()
2906 case Intrinsic::nvvm_tex_3d_v4u32_f32: in getOpcForTextureInstr()
2908 case Intrinsic::nvvm_tex_3d_level_v4u32_f32: in getOpcForTextureInstr()
2910 case Intrinsic::nvvm_tex_3d_grad_v4u32_f32: in getOpcForTextureInstr()
2913 case Intrinsic::nvvm_tex_cube_v4f32_f32: in getOpcForTextureInstr()
2915 case Intrinsic::nvvm_tex_cube_level_v4f32_f32: in getOpcForTextureInstr()
2917 case Intrinsic::nvvm_tex_cube_v4s32_f32: in getOpcForTextureInstr()
2919 case Intrinsic::nvvm_tex_cube_level_v4s32_f32: in getOpcForTextureInstr()
2921 case Intrinsic::nvvm_tex_cube_v4u32_f32: in getOpcForTextureInstr()
2923 case Intrinsic::nvvm_tex_cube_level_v4u32_f32: in getOpcForTextureInstr()
2926 case Intrinsic::nvvm_tex_cube_array_v4f32_f32: in getOpcForTextureInstr()
2928 case Intrinsic::nvvm_tex_cube_array_level_v4f32_f32: in getOpcForTextureInstr()
2930 case Intrinsic::nvvm_tex_cube_array_v4s32_f32: in getOpcForTextureInstr()
2932 case Intrinsic::nvvm_tex_cube_array_level_v4s32_f32: in getOpcForTextureInstr()
2934 case Intrinsic::nvvm_tex_cube_array_v4u32_f32: in getOpcForTextureInstr()
2936 case Intrinsic::nvvm_tex_cube_array_level_v4u32_f32: in getOpcForTextureInstr()
2939 case Intrinsic::nvvm_tld4_r_2d_v4f32_f32: in getOpcForTextureInstr()
2941 case Intrinsic::nvvm_tld4_g_2d_v4f32_f32: in getOpcForTextureInstr()
2943 case Intrinsic::nvvm_tld4_b_2d_v4f32_f32: in getOpcForTextureInstr()
2945 case Intrinsic::nvvm_tld4_a_2d_v4f32_f32: in getOpcForTextureInstr()
2947 case Intrinsic::nvvm_tld4_r_2d_v4s32_f32: in getOpcForTextureInstr()
2949 case Intrinsic::nvvm_tld4_g_2d_v4s32_f32: in getOpcForTextureInstr()
2951 case Intrinsic::nvvm_tld4_b_2d_v4s32_f32: in getOpcForTextureInstr()
2953 case Intrinsic::nvvm_tld4_a_2d_v4s32_f32: in getOpcForTextureInstr()
2955 case Intrinsic::nvvm_tld4_r_2d_v4u32_f32: in getOpcForTextureInstr()
2957 case Intrinsic::nvvm_tld4_g_2d_v4u32_f32: in getOpcForTextureInstr()
2959 case Intrinsic::nvvm_tld4_b_2d_v4u32_f32: in getOpcForTextureInstr()
2961 case Intrinsic::nvvm_tld4_a_2d_v4u32_f32: in getOpcForTextureInstr()
2964 case Intrinsic::nvvm_tex_unified_1d_v4f32_s32: in getOpcForTextureInstr()
2966 case Intrinsic::nvvm_tex_unified_1d_v4f32_f32: in getOpcForTextureInstr()
2968 case Intrinsic::nvvm_tex_unified_1d_level_v4f32_f32: in getOpcForTextureInstr()
2970 case Intrinsic::nvvm_tex_unified_1d_grad_v4f32_f32: in getOpcForTextureInstr()
2972 case Intrinsic::nvvm_tex_unified_1d_v4s32_s32: in getOpcForTextureInstr()
2974 case Intrinsic::nvvm_tex_unified_1d_v4s32_f32: in getOpcForTextureInstr()
2976 case Intrinsic::nvvm_tex_unified_1d_level_v4s32_f32: in getOpcForTextureInstr()
2978 case Intrinsic::nvvm_tex_unified_1d_grad_v4s32_f32: in getOpcForTextureInstr()
2980 case Intrinsic::nvvm_tex_unified_1d_v4u32_s32: in getOpcForTextureInstr()
2982 case Intrinsic::nvvm_tex_unified_1d_v4u32_f32: in getOpcForTextureInstr()
2984 case Intrinsic::nvvm_tex_unified_1d_level_v4u32_f32: in getOpcForTextureInstr()
2986 case Intrinsic::nvvm_tex_unified_1d_grad_v4u32_f32: in getOpcForTextureInstr()
2989 case Intrinsic::nvvm_tex_unified_1d_array_v4f32_s32: in getOpcForTextureInstr()
2991 case Intrinsic::nvvm_tex_unified_1d_array_v4f32_f32: in getOpcForTextureInstr()
2993 case Intrinsic::nvvm_tex_unified_1d_array_level_v4f32_f32: in getOpcForTextureInstr()
2995 case Intrinsic::nvvm_tex_unified_1d_array_grad_v4f32_f32: in getOpcForTextureInstr()
2997 case Intrinsic::nvvm_tex_unified_1d_array_v4s32_s32: in getOpcForTextureInstr()
2999 case Intrinsic::nvvm_tex_unified_1d_array_v4s32_f32: in getOpcForTextureInstr()
3001 case Intrinsic::nvvm_tex_unified_1d_array_level_v4s32_f32: in getOpcForTextureInstr()
3003 case Intrinsic::nvvm_tex_unified_1d_array_grad_v4s32_f32: in getOpcForTextureInstr()
3005 case Intrinsic::nvvm_tex_unified_1d_array_v4u32_s32: in getOpcForTextureInstr()
3007 case Intrinsic::nvvm_tex_unified_1d_array_v4u32_f32: in getOpcForTextureInstr()
3009 case Intrinsic::nvvm_tex_unified_1d_array_level_v4u32_f32: in getOpcForTextureInstr()
3011 case Intrinsic::nvvm_tex_unified_1d_array_grad_v4u32_f32: in getOpcForTextureInstr()
3014 case Intrinsic::nvvm_tex_unified_2d_v4f32_s32: in getOpcForTextureInstr()
3016 case Intrinsic::nvvm_tex_unified_2d_v4f32_f32: in getOpcForTextureInstr()
3018 case Intrinsic::nvvm_tex_unified_2d_level_v4f32_f32: in getOpcForTextureInstr()
3020 case Intrinsic::nvvm_tex_unified_2d_grad_v4f32_f32: in getOpcForTextureInstr()
3022 case Intrinsic::nvvm_tex_unified_2d_v4s32_s32: in getOpcForTextureInstr()
3024 case Intrinsic::nvvm_tex_unified_2d_v4s32_f32: in getOpcForTextureInstr()
3026 case Intrinsic::nvvm_tex_unified_2d_level_v4s32_f32: in getOpcForTextureInstr()
3028 case Intrinsic::nvvm_tex_unified_2d_grad_v4s32_f32: in getOpcForTextureInstr()
3030 case Intrinsic::nvvm_tex_unified_2d_v4u32_s32: in getOpcForTextureInstr()
3032 case Intrinsic::nvvm_tex_unified_2d_v4u32_f32: in getOpcForTextureInstr()
3034 case Intrinsic::nvvm_tex_unified_2d_level_v4u32_f32: in getOpcForTextureInstr()
3036 case Intrinsic::nvvm_tex_unified_2d_grad_v4u32_f32: in getOpcForTextureInstr()
3039 case Intrinsic::nvvm_tex_unified_2d_array_v4f32_s32: in getOpcForTextureInstr()
3041 case Intrinsic::nvvm_tex_unified_2d_array_v4f32_f32: in getOpcForTextureInstr()
3043 case Intrinsic::nvvm_tex_unified_2d_array_level_v4f32_f32: in getOpcForTextureInstr()
3045 case Intrinsic::nvvm_tex_unified_2d_array_grad_v4f32_f32: in getOpcForTextureInstr()
3047 case Intrinsic::nvvm_tex_unified_2d_array_v4s32_s32: in getOpcForTextureInstr()
3049 case Intrinsic::nvvm_tex_unified_2d_array_v4s32_f32: in getOpcForTextureInstr()
3051 case Intrinsic::nvvm_tex_unified_2d_array_level_v4s32_f32: in getOpcForTextureInstr()
3053 case Intrinsic::nvvm_tex_unified_2d_array_grad_v4s32_f32: in getOpcForTextureInstr()
3055 case Intrinsic::nvvm_tex_unified_2d_array_v4u32_s32: in getOpcForTextureInstr()
3057 case Intrinsic::nvvm_tex_unified_2d_array_v4u32_f32: in getOpcForTextureInstr()
3059 case Intrinsic::nvvm_tex_unified_2d_array_level_v4u32_f32: in getOpcForTextureInstr()
3061 case Intrinsic::nvvm_tex_unified_2d_array_grad_v4u32_f32: in getOpcForTextureInstr()
3064 case Intrinsic::nvvm_tex_unified_3d_v4f32_s32: in getOpcForTextureInstr()
3066 case Intrinsic::nvvm_tex_unified_3d_v4f32_f32: in getOpcForTextureInstr()
3068 case Intrinsic::nvvm_tex_unified_3d_level_v4f32_f32: in getOpcForTextureInstr()
3070 case Intrinsic::nvvm_tex_unified_3d_grad_v4f32_f32: in getOpcForTextureInstr()
3072 case Intrinsic::nvvm_tex_unified_3d_v4s32_s32: in getOpcForTextureInstr()
3074 case Intrinsic::nvvm_tex_unified_3d_v4s32_f32: in getOpcForTextureInstr()
3076 case Intrinsic::nvvm_tex_unified_3d_level_v4s32_f32: in getOpcForTextureInstr()
3078 case Intrinsic::nvvm_tex_unified_3d_grad_v4s32_f32: in getOpcForTextureInstr()
3080 case Intrinsic::nvvm_tex_unified_3d_v4u32_s32: in getOpcForTextureInstr()
3082 case Intrinsic::nvvm_tex_unified_3d_v4u32_f32: in getOpcForTextureInstr()
3084 case Intrinsic::nvvm_tex_unified_3d_level_v4u32_f32: in getOpcForTextureInstr()
3086 case Intrinsic::nvvm_tex_unified_3d_grad_v4u32_f32: in getOpcForTextureInstr()
3089 case Intrinsic::nvvm_tex_unified_cube_v4f32_f32: in getOpcForTextureInstr()
3091 case Intrinsic::nvvm_tex_unified_cube_level_v4f32_f32: in getOpcForTextureInstr()
3093 case Intrinsic::nvvm_tex_unified_cube_v4s32_f32: in getOpcForTextureInstr()
3095 case Intrinsic::nvvm_tex_unified_cube_level_v4s32_f32: in getOpcForTextureInstr()
3097 case Intrinsic::nvvm_tex_unified_cube_v4u32_f32: in getOpcForTextureInstr()
3099 case Intrinsic::nvvm_tex_unified_cube_level_v4u32_f32: in getOpcForTextureInstr()
3102 case Intrinsic::nvvm_tex_unified_cube_array_v4f32_f32: in getOpcForTextureInstr()
3104 case Intrinsic::nvvm_tex_unified_cube_array_level_v4f32_f32: in getOpcForTextureInstr()
3106 case Intrinsic::nvvm_tex_unified_cube_array_v4s32_f32: in getOpcForTextureInstr()
3108 case Intrinsic::nvvm_tex_unified_cube_array_level_v4s32_f32: in getOpcForTextureInstr()
3110 case Intrinsic::nvvm_tex_unified_cube_array_v4u32_f32: in getOpcForTextureInstr()
3112 case Intrinsic::nvvm_tex_unified_cube_array_level_v4u32_f32: in getOpcForTextureInstr()
3115 case Intrinsic::nvvm_tld4_unified_r_2d_v4f32_f32: in getOpcForTextureInstr()
3117 case Intrinsic::nvvm_tld4_unified_g_2d_v4f32_f32: in getOpcForTextureInstr()
3119 case Intrinsic::nvvm_tld4_unified_b_2d_v4f32_f32: in getOpcForTextureInstr()
3121 case Intrinsic::nvvm_tld4_unified_a_2d_v4f32_f32: in getOpcForTextureInstr()
3123 case Intrinsic::nvvm_tld4_unified_r_2d_v4s32_f32: in getOpcForTextureInstr()
3125 case Intrinsic::nvvm_tld4_unified_g_2d_v4s32_f32: in getOpcForTextureInstr()
3127 case Intrinsic::nvvm_tld4_unified_b_2d_v4s32_f32: in getOpcForTextureInstr()
3129 case Intrinsic::nvvm_tld4_unified_a_2d_v4s32_f32: in getOpcForTextureInstr()
3131 case Intrinsic::nvvm_tld4_unified_r_2d_v4u32_f32: in getOpcForTextureInstr()
3133 case Intrinsic::nvvm_tld4_unified_g_2d_v4u32_f32: in getOpcForTextureInstr()
3135 case Intrinsic::nvvm_tld4_unified_b_2d_v4u32_f32: in getOpcForTextureInstr()
3137 case Intrinsic::nvvm_tld4_unified_a_2d_v4u32_f32: in getOpcForTextureInstr()
3142 static unsigned getOpcForSurfaceInstr(unsigned Intrinsic) { in getOpcForSurfaceInstr() argument
3143 switch (Intrinsic) { in getOpcForSurfaceInstr()
3146 case Intrinsic::nvvm_suld_1d_i8_clamp: in getOpcForSurfaceInstr()
3148 case Intrinsic::nvvm_suld_1d_i16_clamp: in getOpcForSurfaceInstr()
3150 case Intrinsic::nvvm_suld_1d_i32_clamp: in getOpcForSurfaceInstr()
3152 case Intrinsic::nvvm_suld_1d_i64_clamp: in getOpcForSurfaceInstr()
3154 case Intrinsic::nvvm_suld_1d_v2i8_clamp: in getOpcForSurfaceInstr()
3156 case Intrinsic::nvvm_suld_1d_v2i16_clamp: in getOpcForSurfaceInstr()
3158 case Intrinsic::nvvm_suld_1d_v2i32_clamp: in getOpcForSurfaceInstr()
3160 case Intrinsic::nvvm_suld_1d_v2i64_clamp: in getOpcForSurfaceInstr()
3162 case Intrinsic::nvvm_suld_1d_v4i8_clamp: in getOpcForSurfaceInstr()
3164 case Intrinsic::nvvm_suld_1d_v4i16_clamp: in getOpcForSurfaceInstr()
3166 case Intrinsic::nvvm_suld_1d_v4i32_clamp: in getOpcForSurfaceInstr()
3168 case Intrinsic::nvvm_suld_1d_array_i8_clamp: in getOpcForSurfaceInstr()
3170 case Intrinsic::nvvm_suld_1d_array_i16_clamp: in getOpcForSurfaceInstr()
3172 case Intrinsic::nvvm_suld_1d_array_i32_clamp: in getOpcForSurfaceInstr()
3174 case Intrinsic::nvvm_suld_1d_array_i64_clamp: in getOpcForSurfaceInstr()
3176 case Intrinsic::nvvm_suld_1d_array_v2i8_clamp: in getOpcForSurfaceInstr()
3178 case Intrinsic::nvvm_suld_1d_array_v2i16_clamp: in getOpcForSurfaceInstr()
3180 case Intrinsic::nvvm_suld_1d_array_v2i32_clamp: in getOpcForSurfaceInstr()
3182 case Intrinsic::nvvm_suld_1d_array_v2i64_clamp: in getOpcForSurfaceInstr()
3184 case Intrinsic::nvvm_suld_1d_array_v4i8_clamp: in getOpcForSurfaceInstr()
3186 case Intrinsic::nvvm_suld_1d_array_v4i16_clamp: in getOpcForSurfaceInstr()
3188 case Intrinsic::nvvm_suld_1d_array_v4i32_clamp: in getOpcForSurfaceInstr()
3190 case Intrinsic::nvvm_suld_2d_i8_clamp: in getOpcForSurfaceInstr()
3192 case Intrinsic::nvvm_suld_2d_i16_clamp: in getOpcForSurfaceInstr()
3194 case Intrinsic::nvvm_suld_2d_i32_clamp: in getOpcForSurfaceInstr()
3196 case Intrinsic::nvvm_suld_2d_i64_clamp: in getOpcForSurfaceInstr()
3198 case Intrinsic::nvvm_suld_2d_v2i8_clamp: in getOpcForSurfaceInstr()
3200 case Intrinsic::nvvm_suld_2d_v2i16_clamp: in getOpcForSurfaceInstr()
3202 case Intrinsic::nvvm_suld_2d_v2i32_clamp: in getOpcForSurfaceInstr()
3204 case Intrinsic::nvvm_suld_2d_v2i64_clamp: in getOpcForSurfaceInstr()
3206 case Intrinsic::nvvm_suld_2d_v4i8_clamp: in getOpcForSurfaceInstr()
3208 case Intrinsic::nvvm_suld_2d_v4i16_clamp: in getOpcForSurfaceInstr()
3210 case Intrinsic::nvvm_suld_2d_v4i32_clamp: in getOpcForSurfaceInstr()
3212 case Intrinsic::nvvm_suld_2d_array_i8_clamp: in getOpcForSurfaceInstr()
3214 case Intrinsic::nvvm_suld_2d_array_i16_clamp: in getOpcForSurfaceInstr()
3216 case Intrinsic::nvvm_suld_2d_array_i32_clamp: in getOpcForSurfaceInstr()
3218 case Intrinsic::nvvm_suld_2d_array_i64_clamp: in getOpcForSurfaceInstr()
3220 case Intrinsic::nvvm_suld_2d_array_v2i8_clamp: in getOpcForSurfaceInstr()
3222 case Intrinsic::nvvm_suld_2d_array_v2i16_clamp: in getOpcForSurfaceInstr()
3224 case Intrinsic::nvvm_suld_2d_array_v2i32_clamp: in getOpcForSurfaceInstr()
3226 case Intrinsic::nvvm_suld_2d_array_v2i64_clamp: in getOpcForSurfaceInstr()
3228 case Intrinsic::nvvm_suld_2d_array_v4i8_clamp: in getOpcForSurfaceInstr()
3230 case Intrinsic::nvvm_suld_2d_array_v4i16_clamp: in getOpcForSurfaceInstr()
3232 case Intrinsic::nvvm_suld_2d_array_v4i32_clamp: in getOpcForSurfaceInstr()
3234 case Intrinsic::nvvm_suld_3d_i8_clamp: in getOpcForSurfaceInstr()
3236 case Intrinsic::nvvm_suld_3d_i16_clamp: in getOpcForSurfaceInstr()
3238 case Intrinsic::nvvm_suld_3d_i32_clamp: in getOpcForSurfaceInstr()
3240 case Intrinsic::nvvm_suld_3d_i64_clamp: in getOpcForSurfaceInstr()
3242 case Intrinsic::nvvm_suld_3d_v2i8_clamp: in getOpcForSurfaceInstr()
3244 case Intrinsic::nvvm_suld_3d_v2i16_clamp: in getOpcForSurfaceInstr()
3246 case Intrinsic::nvvm_suld_3d_v2i32_clamp: in getOpcForSurfaceInstr()
3248 case Intrinsic::nvvm_suld_3d_v2i64_clamp: in getOpcForSurfaceInstr()
3250 case Intrinsic::nvvm_suld_3d_v4i8_clamp: in getOpcForSurfaceInstr()
3252 case Intrinsic::nvvm_suld_3d_v4i16_clamp: in getOpcForSurfaceInstr()
3254 case Intrinsic::nvvm_suld_3d_v4i32_clamp: in getOpcForSurfaceInstr()
3256 case Intrinsic::nvvm_suld_1d_i8_trap: in getOpcForSurfaceInstr()
3258 case Intrinsic::nvvm_suld_1d_i16_trap: in getOpcForSurfaceInstr()
3260 case Intrinsic::nvvm_suld_1d_i32_trap: in getOpcForSurfaceInstr()
3262 case Intrinsic::nvvm_suld_1d_i64_trap: in getOpcForSurfaceInstr()
3264 case Intrinsic::nvvm_suld_1d_v2i8_trap: in getOpcForSurfaceInstr()
3266 case Intrinsic::nvvm_suld_1d_v2i16_trap: in getOpcForSurfaceInstr()
3268 case Intrinsic::nvvm_suld_1d_v2i32_trap: in getOpcForSurfaceInstr()
3270 case Intrinsic::nvvm_suld_1d_v2i64_trap: in getOpcForSurfaceInstr()
3272 case Intrinsic::nvvm_suld_1d_v4i8_trap: in getOpcForSurfaceInstr()
3274 case Intrinsic::nvvm_suld_1d_v4i16_trap: in getOpcForSurfaceInstr()
3276 case Intrinsic::nvvm_suld_1d_v4i32_trap: in getOpcForSurfaceInstr()
3278 case Intrinsic::nvvm_suld_1d_array_i8_trap: in getOpcForSurfaceInstr()
3280 case Intrinsic::nvvm_suld_1d_array_i16_trap: in getOpcForSurfaceInstr()
3282 case Intrinsic::nvvm_suld_1d_array_i32_trap: in getOpcForSurfaceInstr()
3284 case Intrinsic::nvvm_suld_1d_array_i64_trap: in getOpcForSurfaceInstr()
3286 case Intrinsic::nvvm_suld_1d_array_v2i8_trap: in getOpcForSurfaceInstr()
3288 case Intrinsic::nvvm_suld_1d_array_v2i16_trap: in getOpcForSurfaceInstr()
3290 case Intrinsic::nvvm_suld_1d_array_v2i32_trap: in getOpcForSurfaceInstr()
3292 case Intrinsic::nvvm_suld_1d_array_v2i64_trap: in getOpcForSurfaceInstr()
3294 case Intrinsic::nvvm_suld_1d_array_v4i8_trap: in getOpcForSurfaceInstr()
3296 case Intrinsic::nvvm_suld_1d_array_v4i16_trap: in getOpcForSurfaceInstr()
3298 case Intrinsic::nvvm_suld_1d_array_v4i32_trap: in getOpcForSurfaceInstr()
3300 case Intrinsic::nvvm_suld_2d_i8_trap: in getOpcForSurfaceInstr()
3302 case Intrinsic::nvvm_suld_2d_i16_trap: in getOpcForSurfaceInstr()
3304 case Intrinsic::nvvm_suld_2d_i32_trap: in getOpcForSurfaceInstr()
3306 case Intrinsic::nvvm_suld_2d_i64_trap: in getOpcForSurfaceInstr()
3308 case Intrinsic::nvvm_suld_2d_v2i8_trap: in getOpcForSurfaceInstr()
3310 case Intrinsic::nvvm_suld_2d_v2i16_trap: in getOpcForSurfaceInstr()
3312 case Intrinsic::nvvm_suld_2d_v2i32_trap: in getOpcForSurfaceInstr()
3314 case Intrinsic::nvvm_suld_2d_v2i64_trap: in getOpcForSurfaceInstr()
3316 case Intrinsic::nvvm_suld_2d_v4i8_trap: in getOpcForSurfaceInstr()
3318 case Intrinsic::nvvm_suld_2d_v4i16_trap: in getOpcForSurfaceInstr()
3320 case Intrinsic::nvvm_suld_2d_v4i32_trap: in getOpcForSurfaceInstr()
3322 case Intrinsic::nvvm_suld_2d_array_i8_trap: in getOpcForSurfaceInstr()
3324 case Intrinsic::nvvm_suld_2d_array_i16_trap: in getOpcForSurfaceInstr()
3326 case Intrinsic::nvvm_suld_2d_array_i32_trap: in getOpcForSurfaceInstr()
3328 case Intrinsic::nvvm_suld_2d_array_i64_trap: in getOpcForSurfaceInstr()
3330 case Intrinsic::nvvm_suld_2d_array_v2i8_trap: in getOpcForSurfaceInstr()
3332 case Intrinsic::nvvm_suld_2d_array_v2i16_trap: in getOpcForSurfaceInstr()
3334 case Intrinsic::nvvm_suld_2d_array_v2i32_trap: in getOpcForSurfaceInstr()
3336 case Intrinsic::nvvm_suld_2d_array_v2i64_trap: in getOpcForSurfaceInstr()
3338 case Intrinsic::nvvm_suld_2d_array_v4i8_trap: in getOpcForSurfaceInstr()
3340 case Intrinsic::nvvm_suld_2d_array_v4i16_trap: in getOpcForSurfaceInstr()
3342 case Intrinsic::nvvm_suld_2d_array_v4i32_trap: in getOpcForSurfaceInstr()
3344 case Intrinsic::nvvm_suld_3d_i8_trap: in getOpcForSurfaceInstr()
3346 case Intrinsic::nvvm_suld_3d_i16_trap: in getOpcForSurfaceInstr()
3348 case Intrinsic::nvvm_suld_3d_i32_trap: in getOpcForSurfaceInstr()
3350 case Intrinsic::nvvm_suld_3d_i64_trap: in getOpcForSurfaceInstr()
3352 case Intrinsic::nvvm_suld_3d_v2i8_trap: in getOpcForSurfaceInstr()
3354 case Intrinsic::nvvm_suld_3d_v2i16_trap: in getOpcForSurfaceInstr()
3356 case Intrinsic::nvvm_suld_3d_v2i32_trap: in getOpcForSurfaceInstr()
3358 case Intrinsic::nvvm_suld_3d_v2i64_trap: in getOpcForSurfaceInstr()
3360 case Intrinsic::nvvm_suld_3d_v4i8_trap: in getOpcForSurfaceInstr()
3362 case Intrinsic::nvvm_suld_3d_v4i16_trap: in getOpcForSurfaceInstr()
3364 case Intrinsic::nvvm_suld_3d_v4i32_trap: in getOpcForSurfaceInstr()
3366 case Intrinsic::nvvm_suld_1d_i8_zero: in getOpcForSurfaceInstr()
3368 case Intrinsic::nvvm_suld_1d_i16_zero: in getOpcForSurfaceInstr()
3370 case Intrinsic::nvvm_suld_1d_i32_zero: in getOpcForSurfaceInstr()
3372 case Intrinsic::nvvm_suld_1d_i64_zero: in getOpcForSurfaceInstr()
3374 case Intrinsic::nvvm_suld_1d_v2i8_zero: in getOpcForSurfaceInstr()
3376 case Intrinsic::nvvm_suld_1d_v2i16_zero: in getOpcForSurfaceInstr()
3378 case Intrinsic::nvvm_suld_1d_v2i32_zero: in getOpcForSurfaceInstr()
3380 case Intrinsic::nvvm_suld_1d_v2i64_zero: in getOpcForSurfaceInstr()
3382 case Intrinsic::nvvm_suld_1d_v4i8_zero: in getOpcForSurfaceInstr()
3384 case Intrinsic::nvvm_suld_1d_v4i16_zero: in getOpcForSurfaceInstr()
3386 case Intrinsic::nvvm_suld_1d_v4i32_zero: in getOpcForSurfaceInstr()
3388 case Intrinsic::nvvm_suld_1d_array_i8_zero: in getOpcForSurfaceInstr()
3390 case Intrinsic::nvvm_suld_1d_array_i16_zero: in getOpcForSurfaceInstr()
3392 case Intrinsic::nvvm_suld_1d_array_i32_zero: in getOpcForSurfaceInstr()
3394 case Intrinsic::nvvm_suld_1d_array_i64_zero: in getOpcForSurfaceInstr()
3396 case Intrinsic::nvvm_suld_1d_array_v2i8_zero: in getOpcForSurfaceInstr()
3398 case Intrinsic::nvvm_suld_1d_array_v2i16_zero: in getOpcForSurfaceInstr()
3400 case Intrinsic::nvvm_suld_1d_array_v2i32_zero: in getOpcForSurfaceInstr()
3402 case Intrinsic::nvvm_suld_1d_array_v2i64_zero: in getOpcForSurfaceInstr()
3404 case Intrinsic::nvvm_suld_1d_array_v4i8_zero: in getOpcForSurfaceInstr()
3406 case Intrinsic::nvvm_suld_1d_array_v4i16_zero: in getOpcForSurfaceInstr()
3408 case Intrinsic::nvvm_suld_1d_array_v4i32_zero: in getOpcForSurfaceInstr()
3410 case Intrinsic::nvvm_suld_2d_i8_zero: in getOpcForSurfaceInstr()
3412 case Intrinsic::nvvm_suld_2d_i16_zero: in getOpcForSurfaceInstr()
3414 case Intrinsic::nvvm_suld_2d_i32_zero: in getOpcForSurfaceInstr()
3416 case Intrinsic::nvvm_suld_2d_i64_zero: in getOpcForSurfaceInstr()
3418 case Intrinsic::nvvm_suld_2d_v2i8_zero: in getOpcForSurfaceInstr()
3420 case Intrinsic::nvvm_suld_2d_v2i16_zero: in getOpcForSurfaceInstr()
3422 case Intrinsic::nvvm_suld_2d_v2i32_zero: in getOpcForSurfaceInstr()
3424 case Intrinsic::nvvm_suld_2d_v2i64_zero: in getOpcForSurfaceInstr()
3426 case Intrinsic::nvvm_suld_2d_v4i8_zero: in getOpcForSurfaceInstr()
3428 case Intrinsic::nvvm_suld_2d_v4i16_zero: in getOpcForSurfaceInstr()
3430 case Intrinsic::nvvm_suld_2d_v4i32_zero: in getOpcForSurfaceInstr()
3432 case Intrinsic::nvvm_suld_2d_array_i8_zero: in getOpcForSurfaceInstr()
3434 case Intrinsic::nvvm_suld_2d_array_i16_zero: in getOpcForSurfaceInstr()
3436 case Intrinsic::nvvm_suld_2d_array_i32_zero: in getOpcForSurfaceInstr()
3438 case Intrinsic::nvvm_suld_2d_array_i64_zero: in getOpcForSurfaceInstr()
3440 case Intrinsic::nvvm_suld_2d_array_v2i8_zero: in getOpcForSurfaceInstr()
3442 case Intrinsic::nvvm_suld_2d_array_v2i16_zero: in getOpcForSurfaceInstr()
3444 case Intrinsic::nvvm_suld_2d_array_v2i32_zero: in getOpcForSurfaceInstr()
3446 case Intrinsic::nvvm_suld_2d_array_v2i64_zero: in getOpcForSurfaceInstr()
3448 case Intrinsic::nvvm_suld_2d_array_v4i8_zero: in getOpcForSurfaceInstr()
3450 case Intrinsic::nvvm_suld_2d_array_v4i16_zero: in getOpcForSurfaceInstr()
3452 case Intrinsic::nvvm_suld_2d_array_v4i32_zero: in getOpcForSurfaceInstr()
3454 case Intrinsic::nvvm_suld_3d_i8_zero: in getOpcForSurfaceInstr()
3456 case Intrinsic::nvvm_suld_3d_i16_zero: in getOpcForSurfaceInstr()
3458 case Intrinsic::nvvm_suld_3d_i32_zero: in getOpcForSurfaceInstr()
3460 case Intrinsic::nvvm_suld_3d_i64_zero: in getOpcForSurfaceInstr()
3462 case Intrinsic::nvvm_suld_3d_v2i8_zero: in getOpcForSurfaceInstr()
3464 case Intrinsic::nvvm_suld_3d_v2i16_zero: in getOpcForSurfaceInstr()
3466 case Intrinsic::nvvm_suld_3d_v2i32_zero: in getOpcForSurfaceInstr()
3468 case Intrinsic::nvvm_suld_3d_v2i64_zero: in getOpcForSurfaceInstr()
3470 case Intrinsic::nvvm_suld_3d_v4i8_zero: in getOpcForSurfaceInstr()
3472 case Intrinsic::nvvm_suld_3d_v4i16_zero: in getOpcForSurfaceInstr()
3474 case Intrinsic::nvvm_suld_3d_v4i32_zero: in getOpcForSurfaceInstr()
3486 MachineFunction &MF, unsigned Intrinsic) const { in getTgtMemIntrinsic()
3487 switch (Intrinsic) { in getTgtMemIntrinsic()
3490 case Intrinsic::nvvm_match_all_sync_i32p: in getTgtMemIntrinsic()
3491 case Intrinsic::nvvm_match_all_sync_i64p: in getTgtMemIntrinsic()
3501 case Intrinsic::nvvm_wmma_m16n16k16_load_a_f16_col: in getTgtMemIntrinsic()
3502 case Intrinsic::nvvm_wmma_m16n16k16_load_a_f16_row: in getTgtMemIntrinsic()
3503 case Intrinsic::nvvm_wmma_m16n16k16_load_a_f16_col_stride: in getTgtMemIntrinsic()
3504 case Intrinsic::nvvm_wmma_m16n16k16_load_a_f16_row_stride: in getTgtMemIntrinsic()
3505 case Intrinsic::nvvm_wmma_m16n16k16_load_b_f16_col: in getTgtMemIntrinsic()
3506 case Intrinsic::nvvm_wmma_m16n16k16_load_b_f16_row: in getTgtMemIntrinsic()
3507 case Intrinsic::nvvm_wmma_m16n16k16_load_b_f16_col_stride: in getTgtMemIntrinsic()
3508 case Intrinsic::nvvm_wmma_m16n16k16_load_b_f16_row_stride: in getTgtMemIntrinsic()
3509 case Intrinsic::nvvm_wmma_m32n8k16_load_a_f16_col: in getTgtMemIntrinsic()
3510 case Intrinsic::nvvm_wmma_m32n8k16_load_a_f16_row: in getTgtMemIntrinsic()
3511 case Intrinsic::nvvm_wmma_m32n8k16_load_a_f16_col_stride: in getTgtMemIntrinsic()
3512 case Intrinsic::nvvm_wmma_m32n8k16_load_a_f16_row_stride: in getTgtMemIntrinsic()
3513 case Intrinsic::nvvm_wmma_m32n8k16_load_b_f16_col: in getTgtMemIntrinsic()
3514 case Intrinsic::nvvm_wmma_m32n8k16_load_b_f16_row: in getTgtMemIntrinsic()
3515 case Intrinsic::nvvm_wmma_m32n8k16_load_b_f16_col_stride: in getTgtMemIntrinsic()
3516 case Intrinsic::nvvm_wmma_m32n8k16_load_b_f16_row_stride: in getTgtMemIntrinsic()
3517 case Intrinsic::nvvm_wmma_m8n32k16_load_a_f16_col: in getTgtMemIntrinsic()
3518 case Intrinsic::nvvm_wmma_m8n32k16_load_a_f16_row: in getTgtMemIntrinsic()
3519 case Intrinsic::nvvm_wmma_m8n32k16_load_a_f16_col_stride: in getTgtMemIntrinsic()
3520 case Intrinsic::nvvm_wmma_m8n32k16_load_a_f16_row_stride: in getTgtMemIntrinsic()
3521 case Intrinsic::nvvm_wmma_m8n32k16_load_b_f16_col: in getTgtMemIntrinsic()
3522 case Intrinsic::nvvm_wmma_m8n32k16_load_b_f16_row: in getTgtMemIntrinsic()
3523 case Intrinsic::nvvm_wmma_m8n32k16_load_b_f16_col_stride: in getTgtMemIntrinsic()
3524 case Intrinsic::nvvm_wmma_m8n32k16_load_b_f16_row_stride: { in getTgtMemIntrinsic()
3533 case Intrinsic::nvvm_wmma_m16n16k16_load_a_s8_col: in getTgtMemIntrinsic()
3534 case Intrinsic::nvvm_wmma_m16n16k16_load_a_s8_col_stride: in getTgtMemIntrinsic()
3535 case Intrinsic::nvvm_wmma_m16n16k16_load_a_u8_col_stride: in getTgtMemIntrinsic()
3536 case Intrinsic::nvvm_wmma_m16n16k16_load_a_u8_col: in getTgtMemIntrinsic()
3537 case Intrinsic::nvvm_wmma_m16n16k16_load_a_s8_row: in getTgtMemIntrinsic()
3538 case Intrinsic::nvvm_wmma_m16n16k16_load_a_s8_row_stride: in getTgtMemIntrinsic()
3539 case Intrinsic::nvvm_wmma_m16n16k16_load_a_u8_row_stride: in getTgtMemIntrinsic()
3540 case Intrinsic::nvvm_wmma_m16n16k16_load_a_u8_row: in getTgtMemIntrinsic()
3541 case Intrinsic::nvvm_wmma_m8n32k16_load_a_bf16_col: in getTgtMemIntrinsic()
3542 case Intrinsic::nvvm_wmma_m8n32k16_load_a_bf16_col_stride: in getTgtMemIntrinsic()
3543 case Intrinsic::nvvm_wmma_m8n32k16_load_a_bf16_row: in getTgtMemIntrinsic()
3544 case Intrinsic::nvvm_wmma_m8n32k16_load_a_bf16_row_stride: in getTgtMemIntrinsic()
3545 case Intrinsic::nvvm_wmma_m16n16k16_load_b_s8_col: in getTgtMemIntrinsic()
3546 case Intrinsic::nvvm_wmma_m16n16k16_load_b_s8_col_stride: in getTgtMemIntrinsic()
3547 case Intrinsic::nvvm_wmma_m16n16k16_load_b_u8_col_stride: in getTgtMemIntrinsic()
3548 case Intrinsic::nvvm_wmma_m16n16k16_load_b_u8_col: in getTgtMemIntrinsic()
3549 case Intrinsic::nvvm_wmma_m16n16k16_load_b_s8_row: in getTgtMemIntrinsic()
3550 case Intrinsic::nvvm_wmma_m16n16k16_load_b_s8_row_stride: in getTgtMemIntrinsic()
3551 case Intrinsic::nvvm_wmma_m16n16k16_load_b_u8_row_stride: in getTgtMemIntrinsic()
3552 case Intrinsic::nvvm_wmma_m16n16k16_load_b_u8_row: in getTgtMemIntrinsic()
3553 case Intrinsic::nvvm_wmma_m32n8k16_load_b_bf16_col: in getTgtMemIntrinsic()
3554 case Intrinsic::nvvm_wmma_m32n8k16_load_b_bf16_col_stride: in getTgtMemIntrinsic()
3555 case Intrinsic::nvvm_wmma_m32n8k16_load_b_bf16_row: in getTgtMemIntrinsic()
3556 case Intrinsic::nvvm_wmma_m32n8k16_load_b_bf16_row_stride: { in getTgtMemIntrinsic()
3566 case Intrinsic::nvvm_wmma_m32n8k16_load_a_s8_col: in getTgtMemIntrinsic()
3567 case Intrinsic::nvvm_wmma_m32n8k16_load_a_s8_col_stride: in getTgtMemIntrinsic()
3568 case Intrinsic::nvvm_wmma_m32n8k16_load_a_u8_col_stride: in getTgtMemIntrinsic()
3569 case Intrinsic::nvvm_wmma_m32n8k16_load_a_u8_col: in getTgtMemIntrinsic()
3570 case Intrinsic::nvvm_wmma_m32n8k16_load_a_s8_row: in getTgtMemIntrinsic()
3571 case Intrinsic::nvvm_wmma_m32n8k16_load_a_s8_row_stride: in getTgtMemIntrinsic()
3572 case Intrinsic::nvvm_wmma_m32n8k16_load_a_u8_row_stride: in getTgtMemIntrinsic()
3573 case Intrinsic::nvvm_wmma_m32n8k16_load_a_u8_row: in getTgtMemIntrinsic()
3574 case Intrinsic::nvvm_wmma_m16n16k16_load_a_bf16_col: in getTgtMemIntrinsic()
3575 case Intrinsic::nvvm_wmma_m16n16k16_load_a_bf16_col_stride: in getTgtMemIntrinsic()
3576 case Intrinsic::nvvm_wmma_m16n16k16_load_a_bf16_row: in getTgtMemIntrinsic()
3577 case Intrinsic::nvvm_wmma_m16n16k16_load_a_bf16_row_stride: in getTgtMemIntrinsic()
3578 case Intrinsic::nvvm_wmma_m16n16k8_load_a_tf32_col: in getTgtMemIntrinsic()
3579 case Intrinsic::nvvm_wmma_m16n16k8_load_a_tf32_col_stride: in getTgtMemIntrinsic()
3580 case Intrinsic::nvvm_wmma_m16n16k8_load_a_tf32_row: in getTgtMemIntrinsic()
3581 case Intrinsic::nvvm_wmma_m16n16k8_load_a_tf32_row_stride: in getTgtMemIntrinsic()
3583 case Intrinsic::nvvm_wmma_m8n32k16_load_b_s8_col: in getTgtMemIntrinsic()
3584 case Intrinsic::nvvm_wmma_m8n32k16_load_b_s8_col_stride: in getTgtMemIntrinsic()
3585 case Intrinsic::nvvm_wmma_m8n32k16_load_b_u8_col_stride: in getTgtMemIntrinsic()
3586 case Intrinsic::nvvm_wmma_m8n32k16_load_b_u8_col: in getTgtMemIntrinsic()
3587 case Intrinsic::nvvm_wmma_m8n32k16_load_b_s8_row: in getTgtMemIntrinsic()
3588 case Intrinsic::nvvm_wmma_m8n32k16_load_b_s8_row_stride: in getTgtMemIntrinsic()
3589 case Intrinsic::nvvm_wmma_m8n32k16_load_b_u8_row_stride: in getTgtMemIntrinsic()
3590 case Intrinsic::nvvm_wmma_m8n32k16_load_b_u8_row: in getTgtMemIntrinsic()
3591 case Intrinsic::nvvm_wmma_m16n16k16_load_b_bf16_col: in getTgtMemIntrinsic()
3592 case Intrinsic::nvvm_wmma_m16n16k16_load_b_bf16_col_stride: in getTgtMemIntrinsic()
3593 case Intrinsic::nvvm_wmma_m16n16k16_load_b_bf16_row: in getTgtMemIntrinsic()
3594 case Intrinsic::nvvm_wmma_m16n16k16_load_b_bf16_row_stride: in getTgtMemIntrinsic()
3595 case Intrinsic::nvvm_wmma_m16n16k8_load_b_tf32_col: in getTgtMemIntrinsic()
3596 case Intrinsic::nvvm_wmma_m16n16k8_load_b_tf32_col_stride: in getTgtMemIntrinsic()
3597 case Intrinsic::nvvm_wmma_m16n16k8_load_b_tf32_row: in getTgtMemIntrinsic()
3598 case Intrinsic::nvvm_wmma_m16n16k8_load_b_tf32_row_stride: in getTgtMemIntrinsic()
3599 case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x4_b16: in getTgtMemIntrinsic()
3600 case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x4_trans_b16: { in getTgtMemIntrinsic()
3610 case Intrinsic::nvvm_wmma_m32n8k16_load_b_s8_col: in getTgtMemIntrinsic()
3611 case Intrinsic::nvvm_wmma_m32n8k16_load_b_s8_col_stride: in getTgtMemIntrinsic()
3612 case Intrinsic::nvvm_wmma_m32n8k16_load_b_u8_col_stride: in getTgtMemIntrinsic()
3613 case Intrinsic::nvvm_wmma_m32n8k16_load_b_u8_col: in getTgtMemIntrinsic()
3614 case Intrinsic::nvvm_wmma_m32n8k16_load_b_s8_row: in getTgtMemIntrinsic()
3615 case Intrinsic::nvvm_wmma_m32n8k16_load_b_s8_row_stride: in getTgtMemIntrinsic()
3616 case Intrinsic::nvvm_wmma_m32n8k16_load_b_u8_row_stride: in getTgtMemIntrinsic()
3617 case Intrinsic::nvvm_wmma_m32n8k16_load_b_u8_row: in getTgtMemIntrinsic()
3619 case Intrinsic::nvvm_wmma_m8n32k16_load_a_s8_col: in getTgtMemIntrinsic()
3620 case Intrinsic::nvvm_wmma_m8n32k16_load_a_s8_col_stride: in getTgtMemIntrinsic()
3621 case Intrinsic::nvvm_wmma_m8n32k16_load_a_u8_col_stride: in getTgtMemIntrinsic()
3622 case Intrinsic::nvvm_wmma_m8n32k16_load_a_u8_col: in getTgtMemIntrinsic()
3623 case Intrinsic::nvvm_wmma_m8n32k16_load_a_s8_row: in getTgtMemIntrinsic()
3624 case Intrinsic::nvvm_wmma_m8n32k16_load_a_s8_row_stride: in getTgtMemIntrinsic()
3625 case Intrinsic::nvvm_wmma_m8n32k16_load_a_u8_row_stride: in getTgtMemIntrinsic()
3626 case Intrinsic::nvvm_wmma_m8n32k16_load_a_u8_row: in getTgtMemIntrinsic()
3627 case Intrinsic::nvvm_wmma_m8n8k128_load_a_b1_row: in getTgtMemIntrinsic()
3628 case Intrinsic::nvvm_wmma_m8n8k128_load_a_b1_row_stride: in getTgtMemIntrinsic()
3629 case Intrinsic::nvvm_wmma_m8n8k128_load_b_b1_col: in getTgtMemIntrinsic()
3630 case Intrinsic::nvvm_wmma_m8n8k128_load_b_b1_col_stride: in getTgtMemIntrinsic()
3631 case Intrinsic::nvvm_wmma_m8n8k32_load_a_s4_row: in getTgtMemIntrinsic()
3632 case Intrinsic::nvvm_wmma_m8n8k32_load_a_s4_row_stride: in getTgtMemIntrinsic()
3633 case Intrinsic::nvvm_wmma_m8n8k32_load_a_u4_row_stride: in getTgtMemIntrinsic()
3634 case Intrinsic::nvvm_wmma_m8n8k32_load_a_u4_row: in getTgtMemIntrinsic()
3635 case Intrinsic::nvvm_wmma_m8n8k32_load_b_s4_col: in getTgtMemIntrinsic()
3636 case Intrinsic::nvvm_wmma_m8n8k32_load_b_s4_col_stride: in getTgtMemIntrinsic()
3637 case Intrinsic::nvvm_wmma_m8n8k32_load_b_u4_col_stride: in getTgtMemIntrinsic()
3638 case Intrinsic::nvvm_wmma_m8n8k32_load_b_u4_col: in getTgtMemIntrinsic()
3639 case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x1_b16: in getTgtMemIntrinsic()
3640 case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x1_trans_b16: { in getTgtMemIntrinsic()
3650 case Intrinsic::nvvm_wmma_m16n16k16_load_c_f16_col: in getTgtMemIntrinsic()
3651 case Intrinsic::nvvm_wmma_m16n16k16_load_c_f16_row: in getTgtMemIntrinsic()
3652 case Intrinsic::nvvm_wmma_m16n16k16_load_c_f16_col_stride: in getTgtMemIntrinsic()
3653 case Intrinsic::nvvm_wmma_m16n16k16_load_c_f16_row_stride: in getTgtMemIntrinsic()
3654 case Intrinsic::nvvm_wmma_m32n8k16_load_c_f16_col: in getTgtMemIntrinsic()
3655 case Intrinsic::nvvm_wmma_m32n8k16_load_c_f16_row: in getTgtMemIntrinsic()
3656 case Intrinsic::nvvm_wmma_m32n8k16_load_c_f16_col_stride: in getTgtMemIntrinsic()
3657 case Intrinsic::nvvm_wmma_m32n8k16_load_c_f16_row_stride: in getTgtMemIntrinsic()
3658 case Intrinsic::nvvm_wmma_m8n32k16_load_c_f16_col: in getTgtMemIntrinsic()
3659 case Intrinsic::nvvm_wmma_m8n32k16_load_c_f16_row: in getTgtMemIntrinsic()
3660 case Intrinsic::nvvm_wmma_m8n32k16_load_c_f16_col_stride: in getTgtMemIntrinsic()
3661 case Intrinsic::nvvm_wmma_m8n32k16_load_c_f16_row_stride: { in getTgtMemIntrinsic()
3671 case Intrinsic::nvvm_wmma_m16n16k16_load_c_f32_col: in getTgtMemIntrinsic()
3672 case Intrinsic::nvvm_wmma_m16n16k16_load_c_f32_row: in getTgtMemIntrinsic()
3673 case Intrinsic::nvvm_wmma_m16n16k16_load_c_f32_col_stride: in getTgtMemIntrinsic()
3674 case Intrinsic::nvvm_wmma_m16n16k16_load_c_f32_row_stride: in getTgtMemIntrinsic()
3675 case Intrinsic::nvvm_wmma_m32n8k16_load_c_f32_col: in getTgtMemIntrinsic()
3676 case Intrinsic::nvvm_wmma_m32n8k16_load_c_f32_row: in getTgtMemIntrinsic()
3677 case Intrinsic::nvvm_wmma_m32n8k16_load_c_f32_col_stride: in getTgtMemIntrinsic()
3678 case Intrinsic::nvvm_wmma_m32n8k16_load_c_f32_row_stride: in getTgtMemIntrinsic()
3679 case Intrinsic::nvvm_wmma_m8n32k16_load_c_f32_col: in getTgtMemIntrinsic()
3680 case Intrinsic::nvvm_wmma_m8n32k16_load_c_f32_row: in getTgtMemIntrinsic()
3681 case Intrinsic::nvvm_wmma_m8n32k16_load_c_f32_col_stride: in getTgtMemIntrinsic()
3682 case Intrinsic::nvvm_wmma_m8n32k16_load_c_f32_row_stride: in getTgtMemIntrinsic()
3683 case Intrinsic::nvvm_wmma_m16n16k8_load_c_f32_col: in getTgtMemIntrinsic()
3684 case Intrinsic::nvvm_wmma_m16n16k8_load_c_f32_row: in getTgtMemIntrinsic()
3685 case Intrinsic::nvvm_wmma_m16n16k8_load_c_f32_col_stride: in getTgtMemIntrinsic()
3686 case Intrinsic::nvvm_wmma_m16n16k8_load_c_f32_row_stride: { in getTgtMemIntrinsic()
3696 case Intrinsic::nvvm_wmma_m32n8k16_load_a_bf16_col: in getTgtMemIntrinsic()
3697 case Intrinsic::nvvm_wmma_m32n8k16_load_a_bf16_col_stride: in getTgtMemIntrinsic()
3698 case Intrinsic::nvvm_wmma_m32n8k16_load_a_bf16_row: in getTgtMemIntrinsic()
3699 case Intrinsic::nvvm_wmma_m32n8k16_load_a_bf16_row_stride: in getTgtMemIntrinsic()
3701 case Intrinsic::nvvm_wmma_m8n32k16_load_b_bf16_col: in getTgtMemIntrinsic()
3702 case Intrinsic::nvvm_wmma_m8n32k16_load_b_bf16_col_stride: in getTgtMemIntrinsic()
3703 case Intrinsic::nvvm_wmma_m8n32k16_load_b_bf16_row: in getTgtMemIntrinsic()
3704 case Intrinsic::nvvm_wmma_m8n32k16_load_b_bf16_row_stride: in getTgtMemIntrinsic()
3706 case Intrinsic::nvvm_wmma_m16n16k16_load_c_s32_col: in getTgtMemIntrinsic()
3707 case Intrinsic::nvvm_wmma_m16n16k16_load_c_s32_col_stride: in getTgtMemIntrinsic()
3708 case Intrinsic::nvvm_wmma_m16n16k16_load_c_s32_row: in getTgtMemIntrinsic()
3709 case Intrinsic::nvvm_wmma_m16n16k16_load_c_s32_row_stride: in getTgtMemIntrinsic()
3710 case Intrinsic::nvvm_wmma_m32n8k16_load_c_s32_col: in getTgtMemIntrinsic()
3711 case Intrinsic::nvvm_wmma_m32n8k16_load_c_s32_col_stride: in getTgtMemIntrinsic()
3712 case Intrinsic::nvvm_wmma_m32n8k16_load_c_s32_row: in getTgtMemIntrinsic()
3713 case Intrinsic::nvvm_wmma_m32n8k16_load_c_s32_row_stride: in getTgtMemIntrinsic()
3714 case Intrinsic::nvvm_wmma_m8n32k16_load_c_s32_col: in getTgtMemIntrinsic()
3715 case Intrinsic::nvvm_wmma_m8n32k16_load_c_s32_col_stride: in getTgtMemIntrinsic()
3716 case Intrinsic::nvvm_wmma_m8n32k16_load_c_s32_row: in getTgtMemIntrinsic()
3717 case Intrinsic::nvvm_wmma_m8n32k16_load_c_s32_row_stride: { in getTgtMemIntrinsic()
3727 case Intrinsic::nvvm_wmma_m8n8k128_load_c_s32_col: in getTgtMemIntrinsic()
3728 case Intrinsic::nvvm_wmma_m8n8k128_load_c_s32_col_stride: in getTgtMemIntrinsic()
3729 case Intrinsic::nvvm_wmma_m8n8k128_load_c_s32_row: in getTgtMemIntrinsic()
3730 case Intrinsic::nvvm_wmma_m8n8k128_load_c_s32_row_stride: in getTgtMemIntrinsic()
3731 case Intrinsic::nvvm_wmma_m8n8k32_load_c_s32_col: in getTgtMemIntrinsic()
3732 case Intrinsic::nvvm_wmma_m8n8k32_load_c_s32_col_stride: in getTgtMemIntrinsic()
3733 case Intrinsic::nvvm_wmma_m8n8k32_load_c_s32_row: in getTgtMemIntrinsic()
3734 case Intrinsic::nvvm_wmma_m8n8k32_load_c_s32_row_stride: in getTgtMemIntrinsic()
3735 case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x2_b16: in getTgtMemIntrinsic()
3736 case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x2_trans_b16: { in getTgtMemIntrinsic()
3746 case Intrinsic::nvvm_wmma_m8n8k4_load_a_f64_col: in getTgtMemIntrinsic()
3747 case Intrinsic::nvvm_wmma_m8n8k4_load_a_f64_col_stride: in getTgtMemIntrinsic()
3748 case Intrinsic::nvvm_wmma_m8n8k4_load_a_f64_row: in getTgtMemIntrinsic()
3749 case Intrinsic::nvvm_wmma_m8n8k4_load_a_f64_row_stride: in getTgtMemIntrinsic()
3751 case Intrinsic::nvvm_wmma_m8n8k4_load_b_f64_col: in getTgtMemIntrinsic()
3752 case Intrinsic::nvvm_wmma_m8n8k4_load_b_f64_col_stride: in getTgtMemIntrinsic()
3753 case Intrinsic::nvvm_wmma_m8n8k4_load_b_f64_row: in getTgtMemIntrinsic()
3754 case Intrinsic::nvvm_wmma_m8n8k4_load_b_f64_row_stride: { in getTgtMemIntrinsic()
3764 case Intrinsic::nvvm_wmma_m8n8k4_load_c_f64_col: in getTgtMemIntrinsic()
3765 case Intrinsic::nvvm_wmma_m8n8k4_load_c_f64_col_stride: in getTgtMemIntrinsic()
3766 case Intrinsic::nvvm_wmma_m8n8k4_load_c_f64_row: in getTgtMemIntrinsic()
3767 case Intrinsic::nvvm_wmma_m8n8k4_load_c_f64_row_stride: { in getTgtMemIntrinsic()
3777 case Intrinsic::nvvm_wmma_m16n16k16_store_d_f16_col: in getTgtMemIntrinsic()
3778 case Intrinsic::nvvm_wmma_m16n16k16_store_d_f16_row: in getTgtMemIntrinsic()
3779 case Intrinsic::nvvm_wmma_m16n16k16_store_d_f16_col_stride: in getTgtMemIntrinsic()
3780 case Intrinsic::nvvm_wmma_m16n16k16_store_d_f16_row_stride: in getTgtMemIntrinsic()
3781 case Intrinsic::nvvm_wmma_m32n8k16_store_d_f16_col: in getTgtMemIntrinsic()
3782 case Intrinsic::nvvm_wmma_m32n8k16_store_d_f16_row: in getTgtMemIntrinsic()
3783 case Intrinsic::nvvm_wmma_m32n8k16_store_d_f16_col_stride: in getTgtMemIntrinsic()
3784 case Intrinsic::nvvm_wmma_m32n8k16_store_d_f16_row_stride: in getTgtMemIntrinsic()
3785 case Intrinsic::nvvm_wmma_m8n32k16_store_d_f16_col: in getTgtMemIntrinsic()
3786 case Intrinsic::nvvm_wmma_m8n32k16_store_d_f16_row: in getTgtMemIntrinsic()
3787 case Intrinsic::nvvm_wmma_m8n32k16_store_d_f16_col_stride: in getTgtMemIntrinsic()
3788 case Intrinsic::nvvm_wmma_m8n32k16_store_d_f16_row_stride: { in getTgtMemIntrinsic()
3798 case Intrinsic::nvvm_wmma_m16n16k16_store_d_f32_col: in getTgtMemIntrinsic()
3799 case Intrinsic::nvvm_wmma_m16n16k16_store_d_f32_row: in getTgtMemIntrinsic()
3800 case Intrinsic::nvvm_wmma_m16n16k16_store_d_f32_col_stride: in getTgtMemIntrinsic()
3801 case Intrinsic::nvvm_wmma_m16n16k16_store_d_f32_row_stride: in getTgtMemIntrinsic()
3802 case Intrinsic::nvvm_wmma_m32n8k16_store_d_f32_col: in getTgtMemIntrinsic()
3803 case Intrinsic::nvvm_wmma_m32n8k16_store_d_f32_row: in getTgtMemIntrinsic()
3804 case Intrinsic::nvvm_wmma_m32n8k16_store_d_f32_col_stride: in getTgtMemIntrinsic()
3805 case Intrinsic::nvvm_wmma_m32n8k16_store_d_f32_row_stride: in getTgtMemIntrinsic()
3806 case Intrinsic::nvvm_wmma_m8n32k16_store_d_f32_col: in getTgtMemIntrinsic()
3807 case Intrinsic::nvvm_wmma_m8n32k16_store_d_f32_row: in getTgtMemIntrinsic()
3808 case Intrinsic::nvvm_wmma_m8n32k16_store_d_f32_col_stride: in getTgtMemIntrinsic()
3809 case Intrinsic::nvvm_wmma_m8n32k16_store_d_f32_row_stride: in getTgtMemIntrinsic()
3810 case Intrinsic::nvvm_wmma_m16n16k8_store_d_f32_col: in getTgtMemIntrinsic()
3811 case Intrinsic::nvvm_wmma_m16n16k8_store_d_f32_row: in getTgtMemIntrinsic()
3812 case Intrinsic::nvvm_wmma_m16n16k8_store_d_f32_col_stride: in getTgtMemIntrinsic()
3813 case Intrinsic::nvvm_wmma_m16n16k8_store_d_f32_row_stride: { in getTgtMemIntrinsic()
3823 case Intrinsic::nvvm_wmma_m16n16k16_store_d_s32_col: in getTgtMemIntrinsic()
3824 case Intrinsic::nvvm_wmma_m16n16k16_store_d_s32_col_stride: in getTgtMemIntrinsic()
3825 case Intrinsic::nvvm_wmma_m16n16k16_store_d_s32_row: in getTgtMemIntrinsic()
3826 case Intrinsic::nvvm_wmma_m16n16k16_store_d_s32_row_stride: in getTgtMemIntrinsic()
3827 case Intrinsic::nvvm_wmma_m32n8k16_store_d_s32_col: in getTgtMemIntrinsic()
3828 case Intrinsic::nvvm_wmma_m32n8k16_store_d_s32_col_stride: in getTgtMemIntrinsic()
3829 case Intrinsic::nvvm_wmma_m32n8k16_store_d_s32_row: in getTgtMemIntrinsic()
3830 case Intrinsic::nvvm_wmma_m32n8k16_store_d_s32_row_stride: in getTgtMemIntrinsic()
3831 case Intrinsic::nvvm_wmma_m8n32k16_store_d_s32_col: in getTgtMemIntrinsic()
3832 case Intrinsic::nvvm_wmma_m8n32k16_store_d_s32_col_stride: in getTgtMemIntrinsic()
3833 case Intrinsic::nvvm_wmma_m8n32k16_store_d_s32_row: in getTgtMemIntrinsic()
3834 case Intrinsic::nvvm_wmma_m8n32k16_store_d_s32_row_stride: { in getTgtMemIntrinsic()
3844 case Intrinsic::nvvm_wmma_m8n8k128_store_d_s32_col: in getTgtMemIntrinsic()
3845 case Intrinsic::nvvm_wmma_m8n8k128_store_d_s32_col_stride: in getTgtMemIntrinsic()
3846 case Intrinsic::nvvm_wmma_m8n8k128_store_d_s32_row: in getTgtMemIntrinsic()
3847 case Intrinsic::nvvm_wmma_m8n8k128_store_d_s32_row_stride: in getTgtMemIntrinsic()
3848 case Intrinsic::nvvm_wmma_m8n8k32_store_d_s32_col: in getTgtMemIntrinsic()
3849 case Intrinsic::nvvm_wmma_m8n8k32_store_d_s32_col_stride: in getTgtMemIntrinsic()
3850 case Intrinsic::nvvm_wmma_m8n8k32_store_d_s32_row: in getTgtMemIntrinsic()
3851 case Intrinsic::nvvm_wmma_m8n8k32_store_d_s32_row_stride: { in getTgtMemIntrinsic()
3861 case Intrinsic::nvvm_wmma_m8n8k4_store_d_f64_col: in getTgtMemIntrinsic()
3862 case Intrinsic::nvvm_wmma_m8n8k4_store_d_f64_col_stride: in getTgtMemIntrinsic()
3863 case Intrinsic::nvvm_wmma_m8n8k4_store_d_f64_row: in getTgtMemIntrinsic()
3864 case Intrinsic::nvvm_wmma_m8n8k4_store_d_f64_row_stride: { in getTgtMemIntrinsic()
3874 case Intrinsic::nvvm_atomic_load_inc_32: in getTgtMemIntrinsic()
3875 case Intrinsic::nvvm_atomic_load_dec_32: in getTgtMemIntrinsic()
3877 case Intrinsic::nvvm_atomic_add_gen_f_cta: in getTgtMemIntrinsic()
3878 case Intrinsic::nvvm_atomic_add_gen_f_sys: in getTgtMemIntrinsic()
3879 case Intrinsic::nvvm_atomic_add_gen_i_cta: in getTgtMemIntrinsic()
3880 case Intrinsic::nvvm_atomic_add_gen_i_sys: in getTgtMemIntrinsic()
3881 case Intrinsic::nvvm_atomic_and_gen_i_cta: in getTgtMemIntrinsic()
3882 case Intrinsic::nvvm_atomic_and_gen_i_sys: in getTgtMemIntrinsic()
3883 case Intrinsic::nvvm_atomic_cas_gen_i_cta: in getTgtMemIntrinsic()
3884 case Intrinsic::nvvm_atomic_cas_gen_i_sys: in getTgtMemIntrinsic()
3885 case Intrinsic::nvvm_atomic_dec_gen_i_cta: in getTgtMemIntrinsic()
3886 case Intrinsic::nvvm_atomic_dec_gen_i_sys: in getTgtMemIntrinsic()
3887 case Intrinsic::nvvm_atomic_inc_gen_i_cta: in getTgtMemIntrinsic()
3888 case Intrinsic::nvvm_atomic_inc_gen_i_sys: in getTgtMemIntrinsic()
3889 case Intrinsic::nvvm_atomic_max_gen_i_cta: in getTgtMemIntrinsic()
3890 case Intrinsic::nvvm_atomic_max_gen_i_sys: in getTgtMemIntrinsic()
3891 case Intrinsic::nvvm_atomic_min_gen_i_cta: in getTgtMemIntrinsic()
3892 case Intrinsic::nvvm_atomic_min_gen_i_sys: in getTgtMemIntrinsic()
3893 case Intrinsic::nvvm_atomic_or_gen_i_cta: in getTgtMemIntrinsic()
3894 case Intrinsic::nvvm_atomic_or_gen_i_sys: in getTgtMemIntrinsic()
3895 case Intrinsic::nvvm_atomic_exch_gen_i_cta: in getTgtMemIntrinsic()
3896 case Intrinsic::nvvm_atomic_exch_gen_i_sys: in getTgtMemIntrinsic()
3897 case Intrinsic::nvvm_atomic_xor_gen_i_cta: in getTgtMemIntrinsic()
3898 case Intrinsic::nvvm_atomic_xor_gen_i_sys: { in getTgtMemIntrinsic()
3909 case Intrinsic::nvvm_ldu_global_i: in getTgtMemIntrinsic()
3910 case Intrinsic::nvvm_ldu_global_f: in getTgtMemIntrinsic()
3911 case Intrinsic::nvvm_ldu_global_p: { in getTgtMemIntrinsic()
3914 if (Intrinsic == Intrinsic::nvvm_ldu_global_i) in getTgtMemIntrinsic()
3916 else if(Intrinsic == Intrinsic::nvvm_ldu_global_p) in getTgtMemIntrinsic()
3927 case Intrinsic::nvvm_ldg_global_i: in getTgtMemIntrinsic()
3928 case Intrinsic::nvvm_ldg_global_f: in getTgtMemIntrinsic()
3929 case Intrinsic::nvvm_ldg_global_p: { in getTgtMemIntrinsic()
3933 if (Intrinsic == Intrinsic::nvvm_ldg_global_i) in getTgtMemIntrinsic()
3935 else if(Intrinsic == Intrinsic::nvvm_ldg_global_p) in getTgtMemIntrinsic()
3947 case Intrinsic::nvvm_tex_1d_v4f32_s32: in getTgtMemIntrinsic()
3948 case Intrinsic::nvvm_tex_1d_v4f32_f32: in getTgtMemIntrinsic()
3949 case Intrinsic::nvvm_tex_1d_level_v4f32_f32: in getTgtMemIntrinsic()
3950 case Intrinsic::nvvm_tex_1d_grad_v4f32_f32: in getTgtMemIntrinsic()
3951 case Intrinsic::nvvm_tex_1d_array_v4f32_s32: in getTgtMemIntrinsic()
3952 case Intrinsic::nvvm_tex_1d_array_v4f32_f32: in getTgtMemIntrinsic()
3953 case Intrinsic::nvvm_tex_1d_array_level_v4f32_f32: in getTgtMemIntrinsic()
3954 case Intrinsic::nvvm_tex_1d_array_grad_v4f32_f32: in getTgtMemIntrinsic()
3955 case Intrinsic::nvvm_tex_2d_v4f32_s32: in getTgtMemIntrinsic()
3956 case Intrinsic::nvvm_tex_2d_v4f32_f32: in getTgtMemIntrinsic()
3957 case Intrinsic::nvvm_tex_2d_level_v4f32_f32: in getTgtMemIntrinsic()
3958 case Intrinsic::nvvm_tex_2d_grad_v4f32_f32: in getTgtMemIntrinsic()
3959 case Intrinsic::nvvm_tex_2d_array_v4f32_s32: in getTgtMemIntrinsic()
3960 case Intrinsic::nvvm_tex_2d_array_v4f32_f32: in getTgtMemIntrinsic()
3961 case Intrinsic::nvvm_tex_2d_array_level_v4f32_f32: in getTgtMemIntrinsic()
3962 case Intrinsic::nvvm_tex_2d_array_grad_v4f32_f32: in getTgtMemIntrinsic()
3963 case Intrinsic::nvvm_tex_3d_v4f32_s32: in getTgtMemIntrinsic()
3964 case Intrinsic::nvvm_tex_3d_v4f32_f32: in getTgtMemIntrinsic()
3965 case Intrinsic::nvvm_tex_3d_level_v4f32_f32: in getTgtMemIntrinsic()
3966 case Intrinsic::nvvm_tex_3d_grad_v4f32_f32: in getTgtMemIntrinsic()
3967 case Intrinsic::nvvm_tex_cube_v4f32_f32: in getTgtMemIntrinsic()
3968 case Intrinsic::nvvm_tex_cube_level_v4f32_f32: in getTgtMemIntrinsic()
3969 case Intrinsic::nvvm_tex_cube_array_v4f32_f32: in getTgtMemIntrinsic()
3970 case Intrinsic::nvvm_tex_cube_array_level_v4f32_f32: in getTgtMemIntrinsic()
3971 case Intrinsic::nvvm_tld4_r_2d_v4f32_f32: in getTgtMemIntrinsic()
3972 case Intrinsic::nvvm_tld4_g_2d_v4f32_f32: in getTgtMemIntrinsic()
3973 case Intrinsic::nvvm_tld4_b_2d_v4f32_f32: in getTgtMemIntrinsic()
3974 case Intrinsic::nvvm_tld4_a_2d_v4f32_f32: in getTgtMemIntrinsic()
3975 case Intrinsic::nvvm_tex_unified_1d_v4f32_s32: in getTgtMemIntrinsic()
3976 case Intrinsic::nvvm_tex_unified_1d_v4f32_f32: in getTgtMemIntrinsic()
3977 case Intrinsic::nvvm_tex_unified_1d_level_v4f32_f32: in getTgtMemIntrinsic()
3978 case Intrinsic::nvvm_tex_unified_1d_grad_v4f32_f32: in getTgtMemIntrinsic()
3979 case Intrinsic::nvvm_tex_unified_1d_array_v4f32_s32: in getTgtMemIntrinsic()
3980 case Intrinsic::nvvm_tex_unified_1d_array_v4f32_f32: in getTgtMemIntrinsic()
3981 case Intrinsic::nvvm_tex_unified_1d_array_level_v4f32_f32: in getTgtMemIntrinsic()
3982 case Intrinsic::nvvm_tex_unified_1d_array_grad_v4f32_f32: in getTgtMemIntrinsic()
3983 case Intrinsic::nvvm_tex_unified_2d_v4f32_s32: in getTgtMemIntrinsic()
3984 case Intrinsic::nvvm_tex_unified_2d_v4f32_f32: in getTgtMemIntrinsic()
3985 case Intrinsic::nvvm_tex_unified_2d_level_v4f32_f32: in getTgtMemIntrinsic()
3986 case Intrinsic::nvvm_tex_unified_2d_grad_v4f32_f32: in getTgtMemIntrinsic()
3987 case Intrinsic::nvvm_tex_unified_2d_array_v4f32_s32: in getTgtMemIntrinsic()
3988 case Intrinsic::nvvm_tex_unified_2d_array_v4f32_f32: in getTgtMemIntrinsic()
3989 case Intrinsic::nvvm_tex_unified_2d_array_level_v4f32_f32: in getTgtMemIntrinsic()
3990 case Intrinsic::nvvm_tex_unified_2d_array_grad_v4f32_f32: in getTgtMemIntrinsic()
3991 case Intrinsic::nvvm_tex_unified_3d_v4f32_s32: in getTgtMemIntrinsic()
3992 case Intrinsic::nvvm_tex_unified_3d_v4f32_f32: in getTgtMemIntrinsic()
3993 case Intrinsic::nvvm_tex_unified_3d_level_v4f32_f32: in getTgtMemIntrinsic()
3994 case Intrinsic::nvvm_tex_unified_3d_grad_v4f32_f32: in getTgtMemIntrinsic()
3995 case Intrinsic::nvvm_tex_unified_cube_v4f32_f32: in getTgtMemIntrinsic()
3996 case Intrinsic::nvvm_tex_unified_cube_level_v4f32_f32: in getTgtMemIntrinsic()
3997 case Intrinsic::nvvm_tex_unified_cube_array_v4f32_f32: in getTgtMemIntrinsic()
3998 case Intrinsic::nvvm_tex_unified_cube_array_level_v4f32_f32: in getTgtMemIntrinsic()
3999 case Intrinsic::nvvm_tld4_unified_r_2d_v4f32_f32: in getTgtMemIntrinsic()
4000 case Intrinsic::nvvm_tld4_unified_g_2d_v4f32_f32: in getTgtMemIntrinsic()
4001 case Intrinsic::nvvm_tld4_unified_b_2d_v4f32_f32: in getTgtMemIntrinsic()
4002 case Intrinsic::nvvm_tld4_unified_a_2d_v4f32_f32: in getTgtMemIntrinsic()
4003 Info.opc = getOpcForTextureInstr(Intrinsic); in getTgtMemIntrinsic()
4011 case Intrinsic::nvvm_tex_1d_v4s32_s32: in getTgtMemIntrinsic()
4012 case Intrinsic::nvvm_tex_1d_v4s32_f32: in getTgtMemIntrinsic()
4013 case Intrinsic::nvvm_tex_1d_level_v4s32_f32: in getTgtMemIntrinsic()
4014 case Intrinsic::nvvm_tex_1d_grad_v4s32_f32: in getTgtMemIntrinsic()
4015 case Intrinsic::nvvm_tex_1d_array_v4s32_s32: in getTgtMemIntrinsic()
4016 case Intrinsic::nvvm_tex_1d_array_v4s32_f32: in getTgtMemIntrinsic()
4017 case Intrinsic::nvvm_tex_1d_array_level_v4s32_f32: in getTgtMemIntrinsic()
4018 case Intrinsic::nvvm_tex_1d_array_grad_v4s32_f32: in getTgtMemIntrinsic()
4019 case Intrinsic::nvvm_tex_2d_v4s32_s32: in getTgtMemIntrinsic()
4020 case Intrinsic::nvvm_tex_2d_v4s32_f32: in getTgtMemIntrinsic()
4021 case Intrinsic::nvvm_tex_2d_level_v4s32_f32: in getTgtMemIntrinsic()
4022 case Intrinsic::nvvm_tex_2d_grad_v4s32_f32: in getTgtMemIntrinsic()
4023 case Intrinsic::nvvm_tex_2d_array_v4s32_s32: in getTgtMemIntrinsic()
4024 case Intrinsic::nvvm_tex_2d_array_v4s32_f32: in getTgtMemIntrinsic()
4025 case Intrinsic::nvvm_tex_2d_array_level_v4s32_f32: in getTgtMemIntrinsic()
4026 case Intrinsic::nvvm_tex_2d_array_grad_v4s32_f32: in getTgtMemIntrinsic()
4027 case Intrinsic::nvvm_tex_3d_v4s32_s32: in getTgtMemIntrinsic()
4028 case Intrinsic::nvvm_tex_3d_v4s32_f32: in getTgtMemIntrinsic()
4029 case Intrinsic::nvvm_tex_3d_level_v4s32_f32: in getTgtMemIntrinsic()
4030 case Intrinsic::nvvm_tex_3d_grad_v4s32_f32: in getTgtMemIntrinsic()
4031 case Intrinsic::nvvm_tex_cube_v4s32_f32: in getTgtMemIntrinsic()
4032 case Intrinsic::nvvm_tex_cube_level_v4s32_f32: in getTgtMemIntrinsic()
4033 case Intrinsic::nvvm_tex_cube_array_v4s32_f32: in getTgtMemIntrinsic()
4034 case Intrinsic::nvvm_tex_cube_array_level_v4s32_f32: in getTgtMemIntrinsic()
4035 case Intrinsic::nvvm_tex_cube_v4u32_f32: in getTgtMemIntrinsic()
4036 case Intrinsic::nvvm_tex_cube_level_v4u32_f32: in getTgtMemIntrinsic()
4037 case Intrinsic::nvvm_tex_cube_array_v4u32_f32: in getTgtMemIntrinsic()
4038 case Intrinsic::nvvm_tex_cube_array_level_v4u32_f32: in getTgtMemIntrinsic()
4039 case Intrinsic::nvvm_tex_1d_v4u32_s32: in getTgtMemIntrinsic()
4040 case Intrinsic::nvvm_tex_1d_v4u32_f32: in getTgtMemIntrinsic()
4041 case Intrinsic::nvvm_tex_1d_level_v4u32_f32: in getTgtMemIntrinsic()
4042 case Intrinsic::nvvm_tex_1d_grad_v4u32_f32: in getTgtMemIntrinsic()
4043 case Intrinsic::nvvm_tex_1d_array_v4u32_s32: in getTgtMemIntrinsic()
4044 case Intrinsic::nvvm_tex_1d_array_v4u32_f32: in getTgtMemIntrinsic()
4045 case Intrinsic::nvvm_tex_1d_array_level_v4u32_f32: in getTgtMemIntrinsic()
4046 case Intrinsic::nvvm_tex_1d_array_grad_v4u32_f32: in getTgtMemIntrinsic()
4047 case Intrinsic::nvvm_tex_2d_v4u32_s32: in getTgtMemIntrinsic()
4048 case Intrinsic::nvvm_tex_2d_v4u32_f32: in getTgtMemIntrinsic()
4049 case Intrinsic::nvvm_tex_2d_level_v4u32_f32: in getTgtMemIntrinsic()
4050 case Intrinsic::nvvm_tex_2d_grad_v4u32_f32: in getTgtMemIntrinsic()
4051 case Intrinsic::nvvm_tex_2d_array_v4u32_s32: in getTgtMemIntrinsic()
4052 case Intrinsic::nvvm_tex_2d_array_v4u32_f32: in getTgtMemIntrinsic()
4053 case Intrinsic::nvvm_tex_2d_array_level_v4u32_f32: in getTgtMemIntrinsic()
4054 case Intrinsic::nvvm_tex_2d_array_grad_v4u32_f32: in getTgtMemIntrinsic()
4055 case Intrinsic::nvvm_tex_3d_v4u32_s32: in getTgtMemIntrinsic()
4056 case Intrinsic::nvvm_tex_3d_v4u32_f32: in getTgtMemIntrinsic()
4057 case Intrinsic::nvvm_tex_3d_level_v4u32_f32: in getTgtMemIntrinsic()
4058 case Intrinsic::nvvm_tex_3d_grad_v4u32_f32: in getTgtMemIntrinsic()
4059 case Intrinsic::nvvm_tld4_r_2d_v4s32_f32: in getTgtMemIntrinsic()
4060 case Intrinsic::nvvm_tld4_g_2d_v4s32_f32: in getTgtMemIntrinsic()
4061 case Intrinsic::nvvm_tld4_b_2d_v4s32_f32: in getTgtMemIntrinsic()
4062 case Intrinsic::nvvm_tld4_a_2d_v4s32_f32: in getTgtMemIntrinsic()
4063 case Intrinsic::nvvm_tld4_r_2d_v4u32_f32: in getTgtMemIntrinsic()
4064 case Intrinsic::nvvm_tld4_g_2d_v4u32_f32: in getTgtMemIntrinsic()
4065 case Intrinsic::nvvm_tld4_b_2d_v4u32_f32: in getTgtMemIntrinsic()
4066 case Intrinsic::nvvm_tld4_a_2d_v4u32_f32: in getTgtMemIntrinsic()
4067 case Intrinsic::nvvm_tex_unified_1d_v4s32_s32: in getTgtMemIntrinsic()
4068 case Intrinsic::nvvm_tex_unified_1d_v4s32_f32: in getTgtMemIntrinsic()
4069 case Intrinsic::nvvm_tex_unified_1d_level_v4s32_f32: in getTgtMemIntrinsic()
4070 case Intrinsic::nvvm_tex_unified_1d_grad_v4s32_f32: in getTgtMemIntrinsic()
4071 case Intrinsic::nvvm_tex_unified_1d_array_v4s32_s32: in getTgtMemIntrinsic()
4072 case Intrinsic::nvvm_tex_unified_1d_array_v4s32_f32: in getTgtMemIntrinsic()
4073 case Intrinsic::nvvm_tex_unified_1d_array_level_v4s32_f32: in getTgtMemIntrinsic()
4074 case Intrinsic::nvvm_tex_unified_1d_array_grad_v4s32_f32: in getTgtMemIntrinsic()
4075 case Intrinsic::nvvm_tex_unified_2d_v4s32_s32: in getTgtMemIntrinsic()
4076 case Intrinsic::nvvm_tex_unified_2d_v4s32_f32: in getTgtMemIntrinsic()
4077 case Intrinsic::nvvm_tex_unified_2d_level_v4s32_f32: in getTgtMemIntrinsic()
4078 case Intrinsic::nvvm_tex_unified_2d_grad_v4s32_f32: in getTgtMemIntrinsic()
4079 case Intrinsic::nvvm_tex_unified_2d_array_v4s32_s32: in getTgtMemIntrinsic()
4080 case Intrinsic::nvvm_tex_unified_2d_array_v4s32_f32: in getTgtMemIntrinsic()
4081 case Intrinsic::nvvm_tex_unified_2d_array_level_v4s32_f32: in getTgtMemIntrinsic()
4082 case Intrinsic::nvvm_tex_unified_2d_array_grad_v4s32_f32: in getTgtMemIntrinsic()
4083 case Intrinsic::nvvm_tex_unified_3d_v4s32_s32: in getTgtMemIntrinsic()
4084 case Intrinsic::nvvm_tex_unified_3d_v4s32_f32: in getTgtMemIntrinsic()
4085 case Intrinsic::nvvm_tex_unified_3d_level_v4s32_f32: in getTgtMemIntrinsic()
4086 case Intrinsic::nvvm_tex_unified_3d_grad_v4s32_f32: in getTgtMemIntrinsic()
4087 case Intrinsic::nvvm_tex_unified_1d_v4u32_s32: in getTgtMemIntrinsic()
4088 case Intrinsic::nvvm_tex_unified_1d_v4u32_f32: in getTgtMemIntrinsic()
4089 case Intrinsic::nvvm_tex_unified_1d_level_v4u32_f32: in getTgtMemIntrinsic()
4090 case Intrinsic::nvvm_tex_unified_1d_grad_v4u32_f32: in getTgtMemIntrinsic()
4091 case Intrinsic::nvvm_tex_unified_1d_array_v4u32_s32: in getTgtMemIntrinsic()
4092 case Intrinsic::nvvm_tex_unified_1d_array_v4u32_f32: in getTgtMemIntrinsic()
4093 case Intrinsic::nvvm_tex_unified_1d_array_level_v4u32_f32: in getTgtMemIntrinsic()
4094 case Intrinsic::nvvm_tex_unified_1d_array_grad_v4u32_f32: in getTgtMemIntrinsic()
4095 case Intrinsic::nvvm_tex_unified_2d_v4u32_s32: in getTgtMemIntrinsic()
4096 case Intrinsic::nvvm_tex_unified_2d_v4u32_f32: in getTgtMemIntrinsic()
4097 case Intrinsic::nvvm_tex_unified_2d_level_v4u32_f32: in getTgtMemIntrinsic()
4098 case Intrinsic::nvvm_tex_unified_2d_grad_v4u32_f32: in getTgtMemIntrinsic()
4099 case Intrinsic::nvvm_tex_unified_2d_array_v4u32_s32: in getTgtMemIntrinsic()
4100 case Intrinsic::nvvm_tex_unified_2d_array_v4u32_f32: in getTgtMemIntrinsic()
4101 case Intrinsic::nvvm_tex_unified_2d_array_level_v4u32_f32: in getTgtMemIntrinsic()
4102 case Intrinsic::nvvm_tex_unified_2d_array_grad_v4u32_f32: in getTgtMemIntrinsic()
4103 case Intrinsic::nvvm_tex_unified_3d_v4u32_s32: in getTgtMemIntrinsic()
4104 case Intrinsic::nvvm_tex_unified_3d_v4u32_f32: in getTgtMemIntrinsic()
4105 case Intrinsic::nvvm_tex_unified_3d_level_v4u32_f32: in getTgtMemIntrinsic()
4106 case Intrinsic::nvvm_tex_unified_3d_grad_v4u32_f32: in getTgtMemIntrinsic()
4107 case Intrinsic::nvvm_tex_unified_cube_v4s32_f32: in getTgtMemIntrinsic()
4108 case Intrinsic::nvvm_tex_unified_cube_level_v4s32_f32: in getTgtMemIntrinsic()
4109 case Intrinsic::nvvm_tex_unified_cube_array_v4s32_f32: in getTgtMemIntrinsic()
4110 case Intrinsic::nvvm_tex_unified_cube_array_level_v4s32_f32: in getTgtMemIntrinsic()
4111 case Intrinsic::nvvm_tex_unified_cube_v4u32_f32: in getTgtMemIntrinsic()
4112 case Intrinsic::nvvm_tex_unified_cube_level_v4u32_f32: in getTgtMemIntrinsic()
4113 case Intrinsic::nvvm_tex_unified_cube_array_v4u32_f32: in getTgtMemIntrinsic()
4114 case Intrinsic::nvvm_tex_unified_cube_array_level_v4u32_f32: in getTgtMemIntrinsic()
4115 case Intrinsic::nvvm_tld4_unified_r_2d_v4s32_f32: in getTgtMemIntrinsic()
4116 case Intrinsic::nvvm_tld4_unified_g_2d_v4s32_f32: in getTgtMemIntrinsic()
4117 case Intrinsic::nvvm_tld4_unified_b_2d_v4s32_f32: in getTgtMemIntrinsic()
4118 case Intrinsic::nvvm_tld4_unified_a_2d_v4s32_f32: in getTgtMemIntrinsic()
4119 case Intrinsic::nvvm_tld4_unified_r_2d_v4u32_f32: in getTgtMemIntrinsic()
4120 case Intrinsic::nvvm_tld4_unified_g_2d_v4u32_f32: in getTgtMemIntrinsic()
4121 case Intrinsic::nvvm_tld4_unified_b_2d_v4u32_f32: in getTgtMemIntrinsic()
4122 case Intrinsic::nvvm_tld4_unified_a_2d_v4u32_f32: in getTgtMemIntrinsic()
4123 Info.opc = getOpcForTextureInstr(Intrinsic); in getTgtMemIntrinsic()
4131 case Intrinsic::nvvm_suld_1d_i8_clamp: in getTgtMemIntrinsic()
4132 case Intrinsic::nvvm_suld_1d_v2i8_clamp: in getTgtMemIntrinsic()
4133 case Intrinsic::nvvm_suld_1d_v4i8_clamp: in getTgtMemIntrinsic()
4134 case Intrinsic::nvvm_suld_1d_array_i8_clamp: in getTgtMemIntrinsic()
4135 case Intrinsic::nvvm_suld_1d_array_v2i8_clamp: in getTgtMemIntrinsic()
4136 case Intrinsic::nvvm_suld_1d_array_v4i8_clamp: in getTgtMemIntrinsic()
4137 case Intrinsic::nvvm_suld_2d_i8_clamp: in getTgtMemIntrinsic()
4138 case Intrinsic::nvvm_suld_2d_v2i8_clamp: in getTgtMemIntrinsic()
4139 case Intrinsic::nvvm_suld_2d_v4i8_clamp: in getTgtMemIntrinsic()
4140 case Intrinsic::nvvm_suld_2d_array_i8_clamp: in getTgtMemIntrinsic()
4141 case Intrinsic::nvvm_suld_2d_array_v2i8_clamp: in getTgtMemIntrinsic()
4142 case Intrinsic::nvvm_suld_2d_array_v4i8_clamp: in getTgtMemIntrinsic()
4143 case Intrinsic::nvvm_suld_3d_i8_clamp: in getTgtMemIntrinsic()
4144 case Intrinsic::nvvm_suld_3d_v2i8_clamp: in getTgtMemIntrinsic()
4145 case Intrinsic::nvvm_suld_3d_v4i8_clamp: in getTgtMemIntrinsic()
4146 case Intrinsic::nvvm_suld_1d_i8_trap: in getTgtMemIntrinsic()
4147 case Intrinsic::nvvm_suld_1d_v2i8_trap: in getTgtMemIntrinsic()
4148 case Intrinsic::nvvm_suld_1d_v4i8_trap: in getTgtMemIntrinsic()
4149 case Intrinsic::nvvm_suld_1d_array_i8_trap: in getTgtMemIntrinsic()
4150 case Intrinsic::nvvm_suld_1d_array_v2i8_trap: in getTgtMemIntrinsic()
4151 case Intrinsic::nvvm_suld_1d_array_v4i8_trap: in getTgtMemIntrinsic()
4152 case Intrinsic::nvvm_suld_2d_i8_trap: in getTgtMemIntrinsic()
4153 case Intrinsic::nvvm_suld_2d_v2i8_trap: in getTgtMemIntrinsic()
4154 case Intrinsic::nvvm_suld_2d_v4i8_trap: in getTgtMemIntrinsic()
4155 case Intrinsic::nvvm_suld_2d_array_i8_trap: in getTgtMemIntrinsic()
4156 case Intrinsic::nvvm_suld_2d_array_v2i8_trap: in getTgtMemIntrinsic()
4157 case Intrinsic::nvvm_suld_2d_array_v4i8_trap: in getTgtMemIntrinsic()
4158 case Intrinsic::nvvm_suld_3d_i8_trap: in getTgtMemIntrinsic()
4159 case Intrinsic::nvvm_suld_3d_v2i8_trap: in getTgtMemIntrinsic()
4160 case Intrinsic::nvvm_suld_3d_v4i8_trap: in getTgtMemIntrinsic()
4161 case Intrinsic::nvvm_suld_1d_i8_zero: in getTgtMemIntrinsic()
4162 case Intrinsic::nvvm_suld_1d_v2i8_zero: in getTgtMemIntrinsic()
4163 case Intrinsic::nvvm_suld_1d_v4i8_zero: in getTgtMemIntrinsic()
4164 case Intrinsic::nvvm_suld_1d_array_i8_zero: in getTgtMemIntrinsic()
4165 case Intrinsic::nvvm_suld_1d_array_v2i8_zero: in getTgtMemIntrinsic()
4166 case Intrinsic::nvvm_suld_1d_array_v4i8_zero: in getTgtMemIntrinsic()
4167 case Intrinsic::nvvm_suld_2d_i8_zero: in getTgtMemIntrinsic()
4168 case Intrinsic::nvvm_suld_2d_v2i8_zero: in getTgtMemIntrinsic()
4169 case Intrinsic::nvvm_suld_2d_v4i8_zero: in getTgtMemIntrinsic()
4170 case Intrinsic::nvvm_suld_2d_array_i8_zero: in getTgtMemIntrinsic()
4171 case Intrinsic::nvvm_suld_2d_array_v2i8_zero: in getTgtMemIntrinsic()
4172 case Intrinsic::nvvm_suld_2d_array_v4i8_zero: in getTgtMemIntrinsic()
4173 case Intrinsic::nvvm_suld_3d_i8_zero: in getTgtMemIntrinsic()
4174 case Intrinsic::nvvm_suld_3d_v2i8_zero: in getTgtMemIntrinsic()
4175 case Intrinsic::nvvm_suld_3d_v4i8_zero: in getTgtMemIntrinsic()
4176 Info.opc = getOpcForSurfaceInstr(Intrinsic); in getTgtMemIntrinsic()
4184 case Intrinsic::nvvm_suld_1d_i16_clamp: in getTgtMemIntrinsic()
4185 case Intrinsic::nvvm_suld_1d_v2i16_clamp: in getTgtMemIntrinsic()
4186 case Intrinsic::nvvm_suld_1d_v4i16_clamp: in getTgtMemIntrinsic()
4187 case Intrinsic::nvvm_suld_1d_array_i16_clamp: in getTgtMemIntrinsic()
4188 case Intrinsic::nvvm_suld_1d_array_v2i16_clamp: in getTgtMemIntrinsic()
4189 case Intrinsic::nvvm_suld_1d_array_v4i16_clamp: in getTgtMemIntrinsic()
4190 case Intrinsic::nvvm_suld_2d_i16_clamp: in getTgtMemIntrinsic()
4191 case Intrinsic::nvvm_suld_2d_v2i16_clamp: in getTgtMemIntrinsic()
4192 case Intrinsic::nvvm_suld_2d_v4i16_clamp: in getTgtMemIntrinsic()
4193 case Intrinsic::nvvm_suld_2d_array_i16_clamp: in getTgtMemIntrinsic()
4194 case Intrinsic::nvvm_suld_2d_array_v2i16_clamp: in getTgtMemIntrinsic()
4195 case Intrinsic::nvvm_suld_2d_array_v4i16_clamp: in getTgtMemIntrinsic()
4196 case Intrinsic::nvvm_suld_3d_i16_clamp: in getTgtMemIntrinsic()
4197 case Intrinsic::nvvm_suld_3d_v2i16_clamp: in getTgtMemIntrinsic()
4198 case Intrinsic::nvvm_suld_3d_v4i16_clamp: in getTgtMemIntrinsic()
4199 case Intrinsic::nvvm_suld_1d_i16_trap: in getTgtMemIntrinsic()
4200 case Intrinsic::nvvm_suld_1d_v2i16_trap: in getTgtMemIntrinsic()
4201 case Intrinsic::nvvm_suld_1d_v4i16_trap: in getTgtMemIntrinsic()
4202 case Intrinsic::nvvm_suld_1d_array_i16_trap: in getTgtMemIntrinsic()
4203 case Intrinsic::nvvm_suld_1d_array_v2i16_trap: in getTgtMemIntrinsic()
4204 case Intrinsic::nvvm_suld_1d_array_v4i16_trap: in getTgtMemIntrinsic()
4205 case Intrinsic::nvvm_suld_2d_i16_trap: in getTgtMemIntrinsic()
4206 case Intrinsic::nvvm_suld_2d_v2i16_trap: in getTgtMemIntrinsic()
4207 case Intrinsic::nvvm_suld_2d_v4i16_trap: in getTgtMemIntrinsic()
4208 case Intrinsic::nvvm_suld_2d_array_i16_trap: in getTgtMemIntrinsic()
4209 case Intrinsic::nvvm_suld_2d_array_v2i16_trap: in getTgtMemIntrinsic()
4210 case Intrinsic::nvvm_suld_2d_array_v4i16_trap: in getTgtMemIntrinsic()
4211 case Intrinsic::nvvm_suld_3d_i16_trap: in getTgtMemIntrinsic()
4212 case Intrinsic::nvvm_suld_3d_v2i16_trap: in getTgtMemIntrinsic()
4213 case Intrinsic::nvvm_suld_3d_v4i16_trap: in getTgtMemIntrinsic()
4214 case Intrinsic::nvvm_suld_1d_i16_zero: in getTgtMemIntrinsic()
4215 case Intrinsic::nvvm_suld_1d_v2i16_zero: in getTgtMemIntrinsic()
4216 case Intrinsic::nvvm_suld_1d_v4i16_zero: in getTgtMemIntrinsic()
4217 case Intrinsic::nvvm_suld_1d_array_i16_zero: in getTgtMemIntrinsic()
4218 case Intrinsic::nvvm_suld_1d_array_v2i16_zero: in getTgtMemIntrinsic()
4219 case Intrinsic::nvvm_suld_1d_array_v4i16_zero: in getTgtMemIntrinsic()
4220 case Intrinsic::nvvm_suld_2d_i16_zero: in getTgtMemIntrinsic()
4221 case Intrinsic::nvvm_suld_2d_v2i16_zero: in getTgtMemIntrinsic()
4222 case Intrinsic::nvvm_suld_2d_v4i16_zero: in getTgtMemIntrinsic()
4223 case Intrinsic::nvvm_suld_2d_array_i16_zero: in getTgtMemIntrinsic()
4224 case Intrinsic::nvvm_suld_2d_array_v2i16_zero: in getTgtMemIntrinsic()
4225 case Intrinsic::nvvm_suld_2d_array_v4i16_zero: in getTgtMemIntrinsic()
4226 case Intrinsic::nvvm_suld_3d_i16_zero: in getTgtMemIntrinsic()
4227 case Intrinsic::nvvm_suld_3d_v2i16_zero: in getTgtMemIntrinsic()
4228 case Intrinsic::nvvm_suld_3d_v4i16_zero: in getTgtMemIntrinsic()
4229 Info.opc = getOpcForSurfaceInstr(Intrinsic); in getTgtMemIntrinsic()
4237 case Intrinsic::nvvm_suld_1d_i32_clamp: in getTgtMemIntrinsic()
4238 case Intrinsic::nvvm_suld_1d_v2i32_clamp: in getTgtMemIntrinsic()
4239 case Intrinsic::nvvm_suld_1d_v4i32_clamp: in getTgtMemIntrinsic()
4240 case Intrinsic::nvvm_suld_1d_array_i32_clamp: in getTgtMemIntrinsic()
4241 case Intrinsic::nvvm_suld_1d_array_v2i32_clamp: in getTgtMemIntrinsic()
4242 case Intrinsic::nvvm_suld_1d_array_v4i32_clamp: in getTgtMemIntrinsic()
4243 case Intrinsic::nvvm_suld_2d_i32_clamp: in getTgtMemIntrinsic()
4244 case Intrinsic::nvvm_suld_2d_v2i32_clamp: in getTgtMemIntrinsic()
4245 case Intrinsic::nvvm_suld_2d_v4i32_clamp: in getTgtMemIntrinsic()
4246 case Intrinsic::nvvm_suld_2d_array_i32_clamp: in getTgtMemIntrinsic()
4247 case Intrinsic::nvvm_suld_2d_array_v2i32_clamp: in getTgtMemIntrinsic()
4248 case Intrinsic::nvvm_suld_2d_array_v4i32_clamp: in getTgtMemIntrinsic()
4249 case Intrinsic::nvvm_suld_3d_i32_clamp: in getTgtMemIntrinsic()
4250 case Intrinsic::nvvm_suld_3d_v2i32_clamp: in getTgtMemIntrinsic()
4251 case Intrinsic::nvvm_suld_3d_v4i32_clamp: in getTgtMemIntrinsic()
4252 case Intrinsic::nvvm_suld_1d_i32_trap: in getTgtMemIntrinsic()
4253 case Intrinsic::nvvm_suld_1d_v2i32_trap: in getTgtMemIntrinsic()
4254 case Intrinsic::nvvm_suld_1d_v4i32_trap: in getTgtMemIntrinsic()
4255 case Intrinsic::nvvm_suld_1d_array_i32_trap: in getTgtMemIntrinsic()
4256 case Intrinsic::nvvm_suld_1d_array_v2i32_trap: in getTgtMemIntrinsic()
4257 case Intrinsic::nvvm_suld_1d_array_v4i32_trap: in getTgtMemIntrinsic()
4258 case Intrinsic::nvvm_suld_2d_i32_trap: in getTgtMemIntrinsic()
4259 case Intrinsic::nvvm_suld_2d_v2i32_trap: in getTgtMemIntrinsic()
4260 case Intrinsic::nvvm_suld_2d_v4i32_trap: in getTgtMemIntrinsic()
4261 case Intrinsic::nvvm_suld_2d_array_i32_trap: in getTgtMemIntrinsic()
4262 case Intrinsic::nvvm_suld_2d_array_v2i32_trap: in getTgtMemIntrinsic()
4263 case Intrinsic::nvvm_suld_2d_array_v4i32_trap: in getTgtMemIntrinsic()
4264 case Intrinsic::nvvm_suld_3d_i32_trap: in getTgtMemIntrinsic()
4265 case Intrinsic::nvvm_suld_3d_v2i32_trap: in getTgtMemIntrinsic()
4266 case Intrinsic::nvvm_suld_3d_v4i32_trap: in getTgtMemIntrinsic()
4267 case Intrinsic::nvvm_suld_1d_i32_zero: in getTgtMemIntrinsic()
4268 case Intrinsic::nvvm_suld_1d_v2i32_zero: in getTgtMemIntrinsic()
4269 case Intrinsic::nvvm_suld_1d_v4i32_zero: in getTgtMemIntrinsic()
4270 case Intrinsic::nvvm_suld_1d_array_i32_zero: in getTgtMemIntrinsic()
4271 case Intrinsic::nvvm_suld_1d_array_v2i32_zero: in getTgtMemIntrinsic()
4272 case Intrinsic::nvvm_suld_1d_array_v4i32_zero: in getTgtMemIntrinsic()
4273 case Intrinsic::nvvm_suld_2d_i32_zero: in getTgtMemIntrinsic()
4274 case Intrinsic::nvvm_suld_2d_v2i32_zero: in getTgtMemIntrinsic()
4275 case Intrinsic::nvvm_suld_2d_v4i32_zero: in getTgtMemIntrinsic()
4276 case Intrinsic::nvvm_suld_2d_array_i32_zero: in getTgtMemIntrinsic()
4277 case Intrinsic::nvvm_suld_2d_array_v2i32_zero: in getTgtMemIntrinsic()
4278 case Intrinsic::nvvm_suld_2d_array_v4i32_zero: in getTgtMemIntrinsic()
4279 case Intrinsic::nvvm_suld_3d_i32_zero: in getTgtMemIntrinsic()
4280 case Intrinsic::nvvm_suld_3d_v2i32_zero: in getTgtMemIntrinsic()
4281 case Intrinsic::nvvm_suld_3d_v4i32_zero: in getTgtMemIntrinsic()
4282 Info.opc = getOpcForSurfaceInstr(Intrinsic); in getTgtMemIntrinsic()
4290 case Intrinsic::nvvm_suld_1d_i64_clamp: in getTgtMemIntrinsic()
4291 case Intrinsic::nvvm_suld_1d_v2i64_clamp: in getTgtMemIntrinsic()
4292 case Intrinsic::nvvm_suld_1d_array_i64_clamp: in getTgtMemIntrinsic()
4293 case Intrinsic::nvvm_suld_1d_array_v2i64_clamp: in getTgtMemIntrinsic()
4294 case Intrinsic::nvvm_suld_2d_i64_clamp: in getTgtMemIntrinsic()
4295 case Intrinsic::nvvm_suld_2d_v2i64_clamp: in getTgtMemIntrinsic()
4296 case Intrinsic::nvvm_suld_2d_array_i64_clamp: in getTgtMemIntrinsic()
4297 case Intrinsic::nvvm_suld_2d_array_v2i64_clamp: in getTgtMemIntrinsic()
4298 case Intrinsic::nvvm_suld_3d_i64_clamp: in getTgtMemIntrinsic()
4299 case Intrinsic::nvvm_suld_3d_v2i64_clamp: in getTgtMemIntrinsic()
4300 case Intrinsic::nvvm_suld_1d_i64_trap: in getTgtMemIntrinsic()
4301 case Intrinsic::nvvm_suld_1d_v2i64_trap: in getTgtMemIntrinsic()
4302 case Intrinsic::nvvm_suld_1d_array_i64_trap: in getTgtMemIntrinsic()
4303 case Intrinsic::nvvm_suld_1d_array_v2i64_trap: in getTgtMemIntrinsic()
4304 case Intrinsic::nvvm_suld_2d_i64_trap: in getTgtMemIntrinsic()
4305 case Intrinsic::nvvm_suld_2d_v2i64_trap: in getTgtMemIntrinsic()
4306 case Intrinsic::nvvm_suld_2d_array_i64_trap: in getTgtMemIntrinsic()
4307 case Intrinsic::nvvm_suld_2d_array_v2i64_trap: in getTgtMemIntrinsic()
4308 case Intrinsic::nvvm_suld_3d_i64_trap: in getTgtMemIntrinsic()
4309 case Intrinsic::nvvm_suld_3d_v2i64_trap: in getTgtMemIntrinsic()
4310 case Intrinsic::nvvm_suld_1d_i64_zero: in getTgtMemIntrinsic()
4311 case Intrinsic::nvvm_suld_1d_v2i64_zero: in getTgtMemIntrinsic()
4312 case Intrinsic::nvvm_suld_1d_array_i64_zero: in getTgtMemIntrinsic()
4313 case Intrinsic::nvvm_suld_1d_array_v2i64_zero: in getTgtMemIntrinsic()
4314 case Intrinsic::nvvm_suld_2d_i64_zero: in getTgtMemIntrinsic()
4315 case Intrinsic::nvvm_suld_2d_v2i64_zero: in getTgtMemIntrinsic()
4316 case Intrinsic::nvvm_suld_2d_array_i64_zero: in getTgtMemIntrinsic()
4317 case Intrinsic::nvvm_suld_2d_array_v2i64_zero: in getTgtMemIntrinsic()
4318 case Intrinsic::nvvm_suld_3d_i64_zero: in getTgtMemIntrinsic()
4319 case Intrinsic::nvvm_suld_3d_v2i64_zero: in getTgtMemIntrinsic()
4320 Info.opc = getOpcForSurfaceInstr(Intrinsic); in getTgtMemIntrinsic()
5062 case Intrinsic::nvvm_ldg_global_i: in ReplaceINTRINSIC_W_CHAIN()
5063 case Intrinsic::nvvm_ldg_global_f: in ReplaceINTRINSIC_W_CHAIN()
5064 case Intrinsic::nvvm_ldg_global_p: in ReplaceINTRINSIC_W_CHAIN()
5065 case Intrinsic::nvvm_ldu_global_i: in ReplaceINTRINSIC_W_CHAIN()
5066 case Intrinsic::nvvm_ldu_global_f: in ReplaceINTRINSIC_W_CHAIN()
5067 case Intrinsic::nvvm_ldu_global_p: { in ReplaceINTRINSIC_W_CHAIN()
5096 case Intrinsic::nvvm_ldg_global_i: in ReplaceINTRINSIC_W_CHAIN()
5097 case Intrinsic::nvvm_ldg_global_f: in ReplaceINTRINSIC_W_CHAIN()
5098 case Intrinsic::nvvm_ldg_global_p: in ReplaceINTRINSIC_W_CHAIN()
5101 case Intrinsic::nvvm_ldu_global_i: in ReplaceINTRINSIC_W_CHAIN()
5102 case Intrinsic::nvvm_ldu_global_f: in ReplaceINTRINSIC_W_CHAIN()
5103 case Intrinsic::nvvm_ldu_global_p: in ReplaceINTRINSIC_W_CHAIN()
5113 case Intrinsic::nvvm_ldg_global_i: in ReplaceINTRINSIC_W_CHAIN()
5114 case Intrinsic::nvvm_ldg_global_f: in ReplaceINTRINSIC_W_CHAIN()
5115 case Intrinsic::nvvm_ldg_global_p: in ReplaceINTRINSIC_W_CHAIN()
5118 case Intrinsic::nvvm_ldu_global_i: in ReplaceINTRINSIC_W_CHAIN()
5119 case Intrinsic::nvvm_ldu_global_f: in ReplaceINTRINSIC_W_CHAIN()
5120 case Intrinsic::nvvm_ldu_global_p: in ReplaceINTRINSIC_W_CHAIN()