Lines Matching defs:Intrinsic
1338 auto MakeIntrinsicCall = [&](Intrinsic::ID IID) {
1349 return MakeIntrinsicCall(Ftz ? Intrinsic::nvvm_rsqrt_approx_ftz_f
1350 : Intrinsic::nvvm_rsqrt_approx_f);
1352 return MakeIntrinsicCall(Intrinsic::nvvm_rsqrt_approx_d);
1357 return MakeIntrinsicCall(Ftz ? Intrinsic::nvvm_sqrt_approx_ftz_f
1358 : Intrinsic::nvvm_sqrt_approx_f);
1366 DAG.getConstant(Intrinsic::nvvm_rcp_approx_ftz_d, DL, MVT::i32),
1367 MakeIntrinsicCall(Intrinsic::nvvm_rsqrt_approx_d));
3523 static unsigned getOpcForTextureInstr(unsigned Intrinsic) {
3524 switch (Intrinsic) {
3528 case Intrinsic::nvvm_tex_1d_v4f32_s32:
3530 case Intrinsic::nvvm_tex_1d_v4f32_f32:
3532 case Intrinsic::nvvm_tex_1d_level_v4f32_f32:
3534 case Intrinsic::nvvm_tex_1d_grad_v4f32_f32:
3536 case Intrinsic::nvvm_tex_1d_v4s32_s32:
3538 case Intrinsic::nvvm_tex_1d_v4s32_f32:
3540 case Intrinsic::nvvm_tex_1d_level_v4s32_f32:
3542 case Intrinsic::nvvm_tex_1d_grad_v4s32_f32:
3544 case Intrinsic::nvvm_tex_1d_v4u32_s32:
3546 case Intrinsic::nvvm_tex_1d_v4u32_f32:
3548 case Intrinsic::nvvm_tex_1d_level_v4u32_f32:
3550 case Intrinsic::nvvm_tex_1d_grad_v4u32_f32:
3553 case Intrinsic::nvvm_tex_1d_array_v4f32_s32:
3555 case Intrinsic::nvvm_tex_1d_array_v4f32_f32:
3557 case Intrinsic::nvvm_tex_1d_array_level_v4f32_f32:
3559 case Intrinsic::nvvm_tex_1d_array_grad_v4f32_f32:
3561 case Intrinsic::nvvm_tex_1d_array_v4s32_s32:
3563 case Intrinsic::nvvm_tex_1d_array_v4s32_f32:
3565 case Intrinsic::nvvm_tex_1d_array_level_v4s32_f32:
3567 case Intrinsic::nvvm_tex_1d_array_grad_v4s32_f32:
3569 case Intrinsic::nvvm_tex_1d_array_v4u32_s32:
3571 case Intrinsic::nvvm_tex_1d_array_v4u32_f32:
3573 case Intrinsic::nvvm_tex_1d_array_level_v4u32_f32:
3575 case Intrinsic::nvvm_tex_1d_array_grad_v4u32_f32:
3578 case Intrinsic::nvvm_tex_2d_v4f32_s32:
3580 case Intrinsic::nvvm_tex_2d_v4f32_f32:
3582 case Intrinsic::nvvm_tex_2d_level_v4f32_f32:
3584 case Intrinsic::nvvm_tex_2d_grad_v4f32_f32:
3586 case Intrinsic::nvvm_tex_2d_v4s32_s32:
3588 case Intrinsic::nvvm_tex_2d_v4s32_f32:
3590 case Intrinsic::nvvm_tex_2d_level_v4s32_f32:
3592 case Intrinsic::nvvm_tex_2d_grad_v4s32_f32:
3594 case Intrinsic::nvvm_tex_2d_v4u32_s32:
3596 case Intrinsic::nvvm_tex_2d_v4u32_f32:
3598 case Intrinsic::nvvm_tex_2d_level_v4u32_f32:
3600 case Intrinsic::nvvm_tex_2d_grad_v4u32_f32:
3603 case Intrinsic::nvvm_tex_2d_array_v4f32_s32:
3605 case Intrinsic::nvvm_tex_2d_array_v4f32_f32:
3607 case Intrinsic::nvvm_tex_2d_array_level_v4f32_f32:
3609 case Intrinsic::nvvm_tex_2d_array_grad_v4f32_f32:
3611 case Intrinsic::nvvm_tex_2d_array_v4s32_s32:
3613 case Intrinsic::nvvm_tex_2d_array_v4s32_f32:
3615 case Intrinsic::nvvm_tex_2d_array_level_v4s32_f32:
3617 case Intrinsic::nvvm_tex_2d_array_grad_v4s32_f32:
3619 case Intrinsic::nvvm_tex_2d_array_v4u32_s32:
3621 case Intrinsic::nvvm_tex_2d_array_v4u32_f32:
3623 case Intrinsic::nvvm_tex_2d_array_level_v4u32_f32:
3625 case Intrinsic::nvvm_tex_2d_array_grad_v4u32_f32:
3628 case Intrinsic::nvvm_tex_3d_v4f32_s32:
3630 case Intrinsic::nvvm_tex_3d_v4f32_f32:
3632 case Intrinsic::nvvm_tex_3d_level_v4f32_f32:
3634 case Intrinsic::nvvm_tex_3d_grad_v4f32_f32:
3636 case Intrinsic::nvvm_tex_3d_v4s32_s32:
3638 case Intrinsic::nvvm_tex_3d_v4s32_f32:
3640 case Intrinsic::nvvm_tex_3d_level_v4s32_f32:
3642 case Intrinsic::nvvm_tex_3d_grad_v4s32_f32:
3644 case Intrinsic::nvvm_tex_3d_v4u32_s32:
3646 case Intrinsic::nvvm_tex_3d_v4u32_f32:
3648 case Intrinsic::nvvm_tex_3d_level_v4u32_f32:
3650 case Intrinsic::nvvm_tex_3d_grad_v4u32_f32:
3653 case Intrinsic::nvvm_tex_cube_v4f32_f32:
3655 case Intrinsic::nvvm_tex_cube_level_v4f32_f32:
3657 case Intrinsic::nvvm_tex_cube_v4s32_f32:
3659 case Intrinsic::nvvm_tex_cube_level_v4s32_f32:
3661 case Intrinsic::nvvm_tex_cube_v4u32_f32:
3663 case Intrinsic::nvvm_tex_cube_level_v4u32_f32:
3666 case Intrinsic::nvvm_tex_cube_array_v4f32_f32:
3668 case Intrinsic::nvvm_tex_cube_array_level_v4f32_f32:
3670 case Intrinsic::nvvm_tex_cube_array_v4s32_f32:
3672 case Intrinsic::nvvm_tex_cube_array_level_v4s32_f32:
3674 case Intrinsic::nvvm_tex_cube_array_v4u32_f32:
3676 case Intrinsic::nvvm_tex_cube_array_level_v4u32_f32:
3679 case Intrinsic::nvvm_tld4_r_2d_v4f32_f32:
3681 case Intrinsic::nvvm_tld4_g_2d_v4f32_f32:
3683 case Intrinsic::nvvm_tld4_b_2d_v4f32_f32:
3685 case Intrinsic::nvvm_tld4_a_2d_v4f32_f32:
3687 case Intrinsic::nvvm_tld4_r_2d_v4s32_f32:
3689 case Intrinsic::nvvm_tld4_g_2d_v4s32_f32:
3691 case Intrinsic::nvvm_tld4_b_2d_v4s32_f32:
3693 case Intrinsic::nvvm_tld4_a_2d_v4s32_f32:
3695 case Intrinsic::nvvm_tld4_r_2d_v4u32_f32:
3697 case Intrinsic::nvvm_tld4_g_2d_v4u32_f32:
3699 case Intrinsic::nvvm_tld4_b_2d_v4u32_f32:
3701 case Intrinsic::nvvm_tld4_a_2d_v4u32_f32:
3704 case Intrinsic::nvvm_tex_unified_1d_v4f32_s32:
3706 case Intrinsic::nvvm_tex_unified_1d_v4f32_f32:
3708 case Intrinsic::nvvm_tex_unified_1d_level_v4f32_f32:
3710 case Intrinsic::nvvm_tex_unified_1d_grad_v4f32_f32:
3712 case Intrinsic::nvvm_tex_unified_1d_v4s32_s32:
3714 case Intrinsic::nvvm_tex_unified_1d_v4s32_f32:
3716 case Intrinsic::nvvm_tex_unified_1d_level_v4s32_f32:
3718 case Intrinsic::nvvm_tex_unified_1d_grad_v4s32_f32:
3720 case Intrinsic::nvvm_tex_unified_1d_v4u32_s32:
3722 case Intrinsic::nvvm_tex_unified_1d_v4u32_f32:
3724 case Intrinsic::nvvm_tex_unified_1d_level_v4u32_f32:
3726 case Intrinsic::nvvm_tex_unified_1d_grad_v4u32_f32:
3729 case Intrinsic::nvvm_tex_unified_1d_array_v4f32_s32:
3731 case Intrinsic::nvvm_tex_unified_1d_array_v4f32_f32:
3733 case Intrinsic::nvvm_tex_unified_1d_array_level_v4f32_f32:
3735 case Intrinsic::nvvm_tex_unified_1d_array_grad_v4f32_f32:
3737 case Intrinsic::nvvm_tex_unified_1d_array_v4s32_s32:
3739 case Intrinsic::nvvm_tex_unified_1d_array_v4s32_f32:
3741 case Intrinsic::nvvm_tex_unified_1d_array_level_v4s32_f32:
3743 case Intrinsic::nvvm_tex_unified_1d_array_grad_v4s32_f32:
3745 case Intrinsic::nvvm_tex_unified_1d_array_v4u32_s32:
3747 case Intrinsic::nvvm_tex_unified_1d_array_v4u32_f32:
3749 case Intrinsic::nvvm_tex_unified_1d_array_level_v4u32_f32:
3751 case Intrinsic::nvvm_tex_unified_1d_array_grad_v4u32_f32:
3754 case Intrinsic::nvvm_tex_unified_2d_v4f32_s32:
3756 case Intrinsic::nvvm_tex_unified_2d_v4f32_f32:
3758 case Intrinsic::nvvm_tex_unified_2d_level_v4f32_f32:
3760 case Intrinsic::nvvm_tex_unified_2d_grad_v4f32_f32:
3762 case Intrinsic::nvvm_tex_unified_2d_v4s32_s32:
3764 case Intrinsic::nvvm_tex_unified_2d_v4s32_f32:
3766 case Intrinsic::nvvm_tex_unified_2d_level_v4s32_f32:
3768 case Intrinsic::nvvm_tex_unified_2d_grad_v4s32_f32:
3770 case Intrinsic::nvvm_tex_unified_2d_v4u32_s32:
3772 case Intrinsic::nvvm_tex_unified_2d_v4u32_f32:
3774 case Intrinsic::nvvm_tex_unified_2d_level_v4u32_f32:
3776 case Intrinsic::nvvm_tex_unified_2d_grad_v4u32_f32:
3779 case Intrinsic::nvvm_tex_unified_2d_array_v4f32_s32:
3781 case Intrinsic::nvvm_tex_unified_2d_array_v4f32_f32:
3783 case Intrinsic::nvvm_tex_unified_2d_array_level_v4f32_f32:
3785 case Intrinsic::nvvm_tex_unified_2d_array_grad_v4f32_f32:
3787 case Intrinsic::nvvm_tex_unified_2d_array_v4s32_s32:
3789 case Intrinsic::nvvm_tex_unified_2d_array_v4s32_f32:
3791 case Intrinsic::nvvm_tex_unified_2d_array_level_v4s32_f32:
3793 case Intrinsic::nvvm_tex_unified_2d_array_grad_v4s32_f32:
3795 case Intrinsic::nvvm_tex_unified_2d_array_v4u32_s32:
3797 case Intrinsic::nvvm_tex_unified_2d_array_v4u32_f32:
3799 case Intrinsic::nvvm_tex_unified_2d_array_level_v4u32_f32:
3801 case Intrinsic::nvvm_tex_unified_2d_array_grad_v4u32_f32:
3804 case Intrinsic::nvvm_tex_unified_3d_v4f32_s32:
3806 case Intrinsic::nvvm_tex_unified_3d_v4f32_f32:
3808 case Intrinsic::nvvm_tex_unified_3d_level_v4f32_f32:
3810 case Intrinsic::nvvm_tex_unified_3d_grad_v4f32_f32:
3812 case Intrinsic::nvvm_tex_unified_3d_v4s32_s32:
3814 case Intrinsic::nvvm_tex_unified_3d_v4s32_f32:
3816 case Intrinsic::nvvm_tex_unified_3d_level_v4s32_f32:
3818 case Intrinsic::nvvm_tex_unified_3d_grad_v4s32_f32:
3820 case Intrinsic::nvvm_tex_unified_3d_v4u32_s32:
3822 case Intrinsic::nvvm_tex_unified_3d_v4u32_f32:
3824 case Intrinsic::nvvm_tex_unified_3d_level_v4u32_f32:
3826 case Intrinsic::nvvm_tex_unified_3d_grad_v4u32_f32:
3829 case Intrinsic::nvvm_tex_unified_cube_v4f32_f32:
3831 case Intrinsic::nvvm_tex_unified_cube_level_v4f32_f32:
3833 case Intrinsic::nvvm_tex_unified_cube_v4s32_f32:
3835 case Intrinsic::nvvm_tex_unified_cube_level_v4s32_f32:
3837 case Intrinsic::nvvm_tex_unified_cube_v4u32_f32:
3839 case Intrinsic::nvvm_tex_unified_cube_level_v4u32_f32:
3842 case Intrinsic::nvvm_tex_unified_cube_array_v4f32_f32:
3844 case Intrinsic::nvvm_tex_unified_cube_array_level_v4f32_f32:
3846 case Intrinsic::nvvm_tex_unified_cube_array_v4s32_f32:
3848 case Intrinsic::nvvm_tex_unified_cube_array_level_v4s32_f32:
3850 case Intrinsic::nvvm_tex_unified_cube_array_v4u32_f32:
3852 case Intrinsic::nvvm_tex_unified_cube_array_level_v4u32_f32:
3855 case Intrinsic::nvvm_tex_unified_cube_grad_v4f32_f32:
3857 case Intrinsic::nvvm_tex_unified_cube_grad_v4s32_f32:
3859 case Intrinsic::nvvm_tex_unified_cube_grad_v4u32_f32:
3861 case Intrinsic::nvvm_tex_unified_cube_array_grad_v4f32_f32:
3863 case Intrinsic::nvvm_tex_unified_cube_array_grad_v4s32_f32:
3865 case Intrinsic::nvvm_tex_unified_cube_array_grad_v4u32_f32:
3868 case Intrinsic::nvvm_tld4_unified_r_2d_v4f32_f32:
3870 case Intrinsic::nvvm_tld4_unified_g_2d_v4f32_f32:
3872 case Intrinsic::nvvm_tld4_unified_b_2d_v4f32_f32:
3874 case Intrinsic::nvvm_tld4_unified_a_2d_v4f32_f32:
3876 case Intrinsic::nvvm_tld4_unified_r_2d_v4s32_f32:
3878 case Intrinsic::nvvm_tld4_unified_g_2d_v4s32_f32:
3880 case Intrinsic::nvvm_tld4_unified_b_2d_v4s32_f32:
3882 case Intrinsic::nvvm_tld4_unified_a_2d_v4s32_f32:
3884 case Intrinsic::nvvm_tld4_unified_r_2d_v4u32_f32:
3886 case Intrinsic::nvvm_tld4_unified_g_2d_v4u32_f32:
3888 case Intrinsic::nvvm_tld4_unified_b_2d_v4u32_f32:
3890 case Intrinsic::nvvm_tld4_unified_a_2d_v4u32_f32:
3895 static unsigned getOpcForSurfaceInstr(unsigned Intrinsic) {
3896 switch (Intrinsic) {
3899 case Intrinsic::nvvm_suld_1d_i8_clamp:
3901 case Intrinsic::nvvm_suld_1d_i16_clamp:
3903 case Intrinsic::nvvm_suld_1d_i32_clamp:
3905 case Intrinsic::nvvm_suld_1d_i64_clamp:
3907 case Intrinsic::nvvm_suld_1d_v2i8_clamp:
3909 case Intrinsic::nvvm_suld_1d_v2i16_clamp:
3911 case Intrinsic::nvvm_suld_1d_v2i32_clamp:
3913 case Intrinsic::nvvm_suld_1d_v2i64_clamp:
3915 case Intrinsic::nvvm_suld_1d_v4i8_clamp:
3917 case Intrinsic::nvvm_suld_1d_v4i16_clamp:
3919 case Intrinsic::nvvm_suld_1d_v4i32_clamp:
3921 case Intrinsic::nvvm_suld_1d_array_i8_clamp:
3923 case Intrinsic::nvvm_suld_1d_array_i16_clamp:
3925 case Intrinsic::nvvm_suld_1d_array_i32_clamp:
3927 case Intrinsic::nvvm_suld_1d_array_i64_clamp:
3929 case Intrinsic::nvvm_suld_1d_array_v2i8_clamp:
3931 case Intrinsic::nvvm_suld_1d_array_v2i16_clamp:
3933 case Intrinsic::nvvm_suld_1d_array_v2i32_clamp:
3935 case Intrinsic::nvvm_suld_1d_array_v2i64_clamp:
3937 case Intrinsic::nvvm_suld_1d_array_v4i8_clamp:
3939 case Intrinsic::nvvm_suld_1d_array_v4i16_clamp:
3941 case Intrinsic::nvvm_suld_1d_array_v4i32_clamp:
3943 case Intrinsic::nvvm_suld_2d_i8_clamp:
3945 case Intrinsic::nvvm_suld_2d_i16_clamp:
3947 case Intrinsic::nvvm_suld_2d_i32_clamp:
3949 case Intrinsic::nvvm_suld_2d_i64_clamp:
3951 case Intrinsic::nvvm_suld_2d_v2i8_clamp:
3953 case Intrinsic::nvvm_suld_2d_v2i16_clamp:
3955 case Intrinsic::nvvm_suld_2d_v2i32_clamp:
3957 case Intrinsic::nvvm_suld_2d_v2i64_clamp:
3959 case Intrinsic::nvvm_suld_2d_v4i8_clamp:
3961 case Intrinsic::nvvm_suld_2d_v4i16_clamp:
3963 case Intrinsic::nvvm_suld_2d_v4i32_clamp:
3965 case Intrinsic::nvvm_suld_2d_array_i8_clamp:
3967 case Intrinsic::nvvm_suld_2d_array_i16_clamp:
3969 case Intrinsic::nvvm_suld_2d_array_i32_clamp:
3971 case Intrinsic::nvvm_suld_2d_array_i64_clamp:
3973 case Intrinsic::nvvm_suld_2d_array_v2i8_clamp:
3975 case Intrinsic::nvvm_suld_2d_array_v2i16_clamp:
3977 case Intrinsic::nvvm_suld_2d_array_v2i32_clamp:
3979 case Intrinsic::nvvm_suld_2d_array_v2i64_clamp:
3981 case Intrinsic::nvvm_suld_2d_array_v4i8_clamp:
3983 case Intrinsic::nvvm_suld_2d_array_v4i16_clamp:
3985 case Intrinsic::nvvm_suld_2d_array_v4i32_clamp:
3987 case Intrinsic::nvvm_suld_3d_i8_clamp:
3989 case Intrinsic::nvvm_suld_3d_i16_clamp:
3991 case Intrinsic::nvvm_suld_3d_i32_clamp:
3993 case Intrinsic::nvvm_suld_3d_i64_clamp:
3995 case Intrinsic::nvvm_suld_3d_v2i8_clamp:
3997 case Intrinsic::nvvm_suld_3d_v2i16_clamp:
3999 case Intrinsic::nvvm_suld_3d_v2i32_clamp:
4001 case Intrinsic::nvvm_suld_3d_v2i64_clamp:
4003 case Intrinsic::nvvm_suld_3d_v4i8_clamp:
4005 case Intrinsic::nvvm_suld_3d_v4i16_clamp:
4007 case Intrinsic::nvvm_suld_3d_v4i32_clamp:
4009 case Intrinsic::nvvm_suld_1d_i8_trap:
4011 case Intrinsic::nvvm_suld_1d_i16_trap:
4013 case Intrinsic::nvvm_suld_1d_i32_trap:
4015 case Intrinsic::nvvm_suld_1d_i64_trap:
4017 case Intrinsic::nvvm_suld_1d_v2i8_trap:
4019 case Intrinsic::nvvm_suld_1d_v2i16_trap:
4021 case Intrinsic::nvvm_suld_1d_v2i32_trap:
4023 case Intrinsic::nvvm_suld_1d_v2i64_trap:
4025 case Intrinsic::nvvm_suld_1d_v4i8_trap:
4027 case Intrinsic::nvvm_suld_1d_v4i16_trap:
4029 case Intrinsic::nvvm_suld_1d_v4i32_trap:
4031 case Intrinsic::nvvm_suld_1d_array_i8_trap:
4033 case Intrinsic::nvvm_suld_1d_array_i16_trap:
4035 case Intrinsic::nvvm_suld_1d_array_i32_trap:
4037 case Intrinsic::nvvm_suld_1d_array_i64_trap:
4039 case Intrinsic::nvvm_suld_1d_array_v2i8_trap:
4041 case Intrinsic::nvvm_suld_1d_array_v2i16_trap:
4043 case Intrinsic::nvvm_suld_1d_array_v2i32_trap:
4045 case Intrinsic::nvvm_suld_1d_array_v2i64_trap:
4047 case Intrinsic::nvvm_suld_1d_array_v4i8_trap:
4049 case Intrinsic::nvvm_suld_1d_array_v4i16_trap:
4051 case Intrinsic::nvvm_suld_1d_array_v4i32_trap:
4053 case Intrinsic::nvvm_suld_2d_i8_trap:
4055 case Intrinsic::nvvm_suld_2d_i16_trap:
4057 case Intrinsic::nvvm_suld_2d_i32_trap:
4059 case Intrinsic::nvvm_suld_2d_i64_trap:
4061 case Intrinsic::nvvm_suld_2d_v2i8_trap:
4063 case Intrinsic::nvvm_suld_2d_v2i16_trap:
4065 case Intrinsic::nvvm_suld_2d_v2i32_trap:
4067 case Intrinsic::nvvm_suld_2d_v2i64_trap:
4069 case Intrinsic::nvvm_suld_2d_v4i8_trap:
4071 case Intrinsic::nvvm_suld_2d_v4i16_trap:
4073 case Intrinsic::nvvm_suld_2d_v4i32_trap:
4075 case Intrinsic::nvvm_suld_2d_array_i8_trap:
4077 case Intrinsic::nvvm_suld_2d_array_i16_trap:
4079 case Intrinsic::nvvm_suld_2d_array_i32_trap:
4081 case Intrinsic::nvvm_suld_2d_array_i64_trap:
4083 case Intrinsic::nvvm_suld_2d_array_v2i8_trap:
4085 case Intrinsic::nvvm_suld_2d_array_v2i16_trap:
4087 case Intrinsic::nvvm_suld_2d_array_v2i32_trap:
4089 case Intrinsic::nvvm_suld_2d_array_v2i64_trap:
4091 case Intrinsic::nvvm_suld_2d_array_v4i8_trap:
4093 case Intrinsic::nvvm_suld_2d_array_v4i16_trap:
4095 case Intrinsic::nvvm_suld_2d_array_v4i32_trap:
4097 case Intrinsic::nvvm_suld_3d_i8_trap:
4099 case Intrinsic::nvvm_suld_3d_i16_trap:
4101 case Intrinsic::nvvm_suld_3d_i32_trap:
4103 case Intrinsic::nvvm_suld_3d_i64_trap:
4105 case Intrinsic::nvvm_suld_3d_v2i8_trap:
4107 case Intrinsic::nvvm_suld_3d_v2i16_trap:
4109 case Intrinsic::nvvm_suld_3d_v2i32_trap:
4111 case Intrinsic::nvvm_suld_3d_v2i64_trap:
4113 case Intrinsic::nvvm_suld_3d_v4i8_trap:
4115 case Intrinsic::nvvm_suld_3d_v4i16_trap:
4117 case Intrinsic::nvvm_suld_3d_v4i32_trap:
4119 case Intrinsic::nvvm_suld_1d_i8_zero:
4121 case Intrinsic::nvvm_suld_1d_i16_zero:
4123 case Intrinsic::nvvm_suld_1d_i32_zero:
4125 case Intrinsic::nvvm_suld_1d_i64_zero:
4127 case Intrinsic::nvvm_suld_1d_v2i8_zero:
4129 case Intrinsic::nvvm_suld_1d_v2i16_zero:
4131 case Intrinsic::nvvm_suld_1d_v2i32_zero:
4133 case Intrinsic::nvvm_suld_1d_v2i64_zero:
4135 case Intrinsic::nvvm_suld_1d_v4i8_zero:
4137 case Intrinsic::nvvm_suld_1d_v4i16_zero:
4139 case Intrinsic::nvvm_suld_1d_v4i32_zero:
4141 case Intrinsic::nvvm_suld_1d_array_i8_zero:
4143 case Intrinsic::nvvm_suld_1d_array_i16_zero:
4145 case Intrinsic::nvvm_suld_1d_array_i32_zero:
4147 case Intrinsic::nvvm_suld_1d_array_i64_zero:
4149 case Intrinsic::nvvm_suld_1d_array_v2i8_zero:
4151 case Intrinsic::nvvm_suld_1d_array_v2i16_zero:
4153 case Intrinsic::nvvm_suld_1d_array_v2i32_zero:
4155 case Intrinsic::nvvm_suld_1d_array_v2i64_zero:
4157 case Intrinsic::nvvm_suld_1d_array_v4i8_zero:
4159 case Intrinsic::nvvm_suld_1d_array_v4i16_zero:
4161 case Intrinsic::nvvm_suld_1d_array_v4i32_zero:
4163 case Intrinsic::nvvm_suld_2d_i8_zero:
4165 case Intrinsic::nvvm_suld_2d_i16_zero:
4167 case Intrinsic::nvvm_suld_2d_i32_zero:
4169 case Intrinsic::nvvm_suld_2d_i64_zero:
4171 case Intrinsic::nvvm_suld_2d_v2i8_zero:
4173 case Intrinsic::nvvm_suld_2d_v2i16_zero:
4175 case Intrinsic::nvvm_suld_2d_v2i32_zero:
4177 case Intrinsic::nvvm_suld_2d_v2i64_zero:
4179 case Intrinsic::nvvm_suld_2d_v4i8_zero:
4181 case Intrinsic::nvvm_suld_2d_v4i16_zero:
4183 case Intrinsic::nvvm_suld_2d_v4i32_zero:
4185 case Intrinsic::nvvm_suld_2d_array_i8_zero:
4187 case Intrinsic::nvvm_suld_2d_array_i16_zero:
4189 case Intrinsic::nvvm_suld_2d_array_i32_zero:
4191 case Intrinsic::nvvm_suld_2d_array_i64_zero:
4193 case Intrinsic::nvvm_suld_2d_array_v2i8_zero:
4195 case Intrinsic::nvvm_suld_2d_array_v2i16_zero:
4197 case Intrinsic::nvvm_suld_2d_array_v2i32_zero:
4199 case Intrinsic::nvvm_suld_2d_array_v2i64_zero:
4201 case Intrinsic::nvvm_suld_2d_array_v4i8_zero:
4203 case Intrinsic::nvvm_suld_2d_array_v4i16_zero:
4205 case Intrinsic::nvvm_suld_2d_array_v4i32_zero:
4207 case Intrinsic::nvvm_suld_3d_i8_zero:
4209 case Intrinsic::nvvm_suld_3d_i16_zero:
4211 case Intrinsic::nvvm_suld_3d_i32_zero:
4213 case Intrinsic::nvvm_suld_3d_i64_zero:
4215 case Intrinsic::nvvm_suld_3d_v2i8_zero:
4217 case Intrinsic::nvvm_suld_3d_v2i16_zero:
4219 case Intrinsic::nvvm_suld_3d_v2i32_zero:
4221 case Intrinsic::nvvm_suld_3d_v2i64_zero:
4223 case Intrinsic::nvvm_suld_3d_v4i8_zero:
4225 case Intrinsic::nvvm_suld_3d_v4i16_zero:
4227 case Intrinsic::nvvm_suld_3d_v4i32_zero:
4239 MachineFunction &MF, unsigned Intrinsic) const {
4240 switch (Intrinsic) {
4243 case Intrinsic::nvvm_match_all_sync_i32p:
4244 case Intrinsic::nvvm_match_all_sync_i64p:
4254 case Intrinsic::nvvm_wmma_m16n16k16_load_a_f16_col:
4255 case Intrinsic::nvvm_wmma_m16n16k16_load_a_f16_row:
4256 case Intrinsic::nvvm_wmma_m16n16k16_load_a_f16_col_stride:
4257 case Intrinsic::nvvm_wmma_m16n16k16_load_a_f16_row_stride:
4258 case Intrinsic::nvvm_wmma_m16n16k16_load_b_f16_col:
4259 case Intrinsic::nvvm_wmma_m16n16k16_load_b_f16_row:
4260 case Intrinsic::nvvm_wmma_m16n16k16_load_b_f16_col_stride:
4261 case Intrinsic::nvvm_wmma_m16n16k16_load_b_f16_row_stride:
4262 case Intrinsic::nvvm_wmma_m32n8k16_load_a_f16_col:
4263 case Intrinsic::nvvm_wmma_m32n8k16_load_a_f16_row:
4264 case Intrinsic::nvvm_wmma_m32n8k16_load_a_f16_col_stride:
4265 case Intrinsic::nvvm_wmma_m32n8k16_load_a_f16_row_stride:
4266 case Intrinsic::nvvm_wmma_m32n8k16_load_b_f16_col:
4267 case Intrinsic::nvvm_wmma_m32n8k16_load_b_f16_row:
4268 case Intrinsic::nvvm_wmma_m32n8k16_load_b_f16_col_stride:
4269 case Intrinsic::nvvm_wmma_m32n8k16_load_b_f16_row_stride:
4270 case Intrinsic::nvvm_wmma_m8n32k16_load_a_f16_col:
4271 case Intrinsic::nvvm_wmma_m8n32k16_load_a_f16_row:
4272 case Intrinsic::nvvm_wmma_m8n32k16_load_a_f16_col_stride:
4273 case Intrinsic::nvvm_wmma_m8n32k16_load_a_f16_row_stride:
4274 case Intrinsic::nvvm_wmma_m8n32k16_load_b_f16_col:
4275 case Intrinsic::nvvm_wmma_m8n32k16_load_b_f16_row:
4276 case Intrinsic::nvvm_wmma_m8n32k16_load_b_f16_col_stride:
4277 case Intrinsic::nvvm_wmma_m8n32k16_load_b_f16_row_stride: {
4286 case Intrinsic::nvvm_wmma_m16n16k16_load_a_s8_col:
4287 case Intrinsic::nvvm_wmma_m16n16k16_load_a_s8_col_stride:
4288 case Intrinsic::nvvm_wmma_m16n16k16_load_a_u8_col_stride:
4289 case Intrinsic::nvvm_wmma_m16n16k16_load_a_u8_col:
4290 case Intrinsic::nvvm_wmma_m16n16k16_load_a_s8_row:
4291 case Intrinsic::nvvm_wmma_m16n16k16_load_a_s8_row_stride:
4292 case Intrinsic::nvvm_wmma_m16n16k16_load_a_u8_row_stride:
4293 case Intrinsic::nvvm_wmma_m16n16k16_load_a_u8_row:
4294 case Intrinsic::nvvm_wmma_m8n32k16_load_a_bf16_col:
4295 case Intrinsic::nvvm_wmma_m8n32k16_load_a_bf16_col_stride:
4296 case Intrinsic::nvvm_wmma_m8n32k16_load_a_bf16_row:
4297 case Intrinsic::nvvm_wmma_m8n32k16_load_a_bf16_row_stride:
4298 case Intrinsic::nvvm_wmma_m16n16k16_load_b_s8_col:
4299 case Intrinsic::nvvm_wmma_m16n16k16_load_b_s8_col_stride:
4300 case Intrinsic::nvvm_wmma_m16n16k16_load_b_u8_col_stride:
4301 case Intrinsic::nvvm_wmma_m16n16k16_load_b_u8_col:
4302 case Intrinsic::nvvm_wmma_m16n16k16_load_b_s8_row:
4303 case Intrinsic::nvvm_wmma_m16n16k16_load_b_s8_row_stride:
4304 case Intrinsic::nvvm_wmma_m16n16k16_load_b_u8_row_stride:
4305 case Intrinsic::nvvm_wmma_m16n16k16_load_b_u8_row:
4306 case Intrinsic::nvvm_wmma_m32n8k16_load_b_bf16_col:
4307 case Intrinsic::nvvm_wmma_m32n8k16_load_b_bf16_col_stride:
4308 case Intrinsic::nvvm_wmma_m32n8k16_load_b_bf16_row:
4309 case Intrinsic::nvvm_wmma_m32n8k16_load_b_bf16_row_stride: {
4319 case Intrinsic::nvvm_wmma_m32n8k16_load_a_s8_col:
4320 case Intrinsic::nvvm_wmma_m32n8k16_load_a_s8_col_stride:
4321 case Intrinsic::nvvm_wmma_m32n8k16_load_a_u8_col_stride:
4322 case Intrinsic::nvvm_wmma_m32n8k16_load_a_u8_col:
4323 case Intrinsic::nvvm_wmma_m32n8k16_load_a_s8_row:
4324 case Intrinsic::nvvm_wmma_m32n8k16_load_a_s8_row_stride:
4325 case Intrinsic::nvvm_wmma_m32n8k16_load_a_u8_row_stride:
4326 case Intrinsic::nvvm_wmma_m32n8k16_load_a_u8_row:
4327 case Intrinsic::nvvm_wmma_m16n16k16_load_a_bf16_col:
4328 case Intrinsic::nvvm_wmma_m16n16k16_load_a_bf16_col_stride:
4329 case Intrinsic::nvvm_wmma_m16n16k16_load_a_bf16_row:
4330 case Intrinsic::nvvm_wmma_m16n16k16_load_a_bf16_row_stride:
4331 case Intrinsic::nvvm_wmma_m16n16k8_load_a_tf32_col:
4332 case Intrinsic::nvvm_wmma_m16n16k8_load_a_tf32_col_stride:
4333 case Intrinsic::nvvm_wmma_m16n16k8_load_a_tf32_row:
4334 case Intrinsic::nvvm_wmma_m16n16k8_load_a_tf32_row_stride:
4336 case Intrinsic::nvvm_wmma_m8n32k16_load_b_s8_col:
4337 case Intrinsic::nvvm_wmma_m8n32k16_load_b_s8_col_stride:
4338 case Intrinsic::nvvm_wmma_m8n32k16_load_b_u8_col_stride:
4339 case Intrinsic::nvvm_wmma_m8n32k16_load_b_u8_col:
4340 case Intrinsic::nvvm_wmma_m8n32k16_load_b_s8_row:
4341 case Intrinsic::nvvm_wmma_m8n32k16_load_b_s8_row_stride:
4342 case Intrinsic::nvvm_wmma_m8n32k16_load_b_u8_row_stride:
4343 case Intrinsic::nvvm_wmma_m8n32k16_load_b_u8_row:
4344 case Intrinsic::nvvm_wmma_m16n16k16_load_b_bf16_col:
4345 case Intrinsic::nvvm_wmma_m16n16k16_load_b_bf16_col_stride:
4346 case Intrinsic::nvvm_wmma_m16n16k16_load_b_bf16_row:
4347 case Intrinsic::nvvm_wmma_m16n16k16_load_b_bf16_row_stride:
4348 case Intrinsic::nvvm_wmma_m16n16k8_load_b_tf32_col:
4349 case Intrinsic::nvvm_wmma_m16n16k8_load_b_tf32_col_stride:
4350 case Intrinsic::nvvm_wmma_m16n16k8_load_b_tf32_row:
4351 case Intrinsic::nvvm_wmma_m16n16k8_load_b_tf32_row_stride:
4352 case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x4_b16:
4353 case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x4_trans_b16: {
4363 case Intrinsic::nvvm_wmma_m32n8k16_load_b_s8_col:
4364 case Intrinsic::nvvm_wmma_m32n8k16_load_b_s8_col_stride:
4365 case Intrinsic::nvvm_wmma_m32n8k16_load_b_u8_col_stride:
4366 case Intrinsic::nvvm_wmma_m32n8k16_load_b_u8_col:
4367 case Intrinsic::nvvm_wmma_m32n8k16_load_b_s8_row:
4368 case Intrinsic::nvvm_wmma_m32n8k16_load_b_s8_row_stride:
4369 case Intrinsic::nvvm_wmma_m32n8k16_load_b_u8_row_stride:
4370 case Intrinsic::nvvm_wmma_m32n8k16_load_b_u8_row:
4372 case Intrinsic::nvvm_wmma_m8n32k16_load_a_s8_col:
4373 case Intrinsic::nvvm_wmma_m8n32k16_load_a_s8_col_stride:
4374 case Intrinsic::nvvm_wmma_m8n32k16_load_a_u8_col_stride:
4375 case Intrinsic::nvvm_wmma_m8n32k16_load_a_u8_col:
4376 case Intrinsic::nvvm_wmma_m8n32k16_load_a_s8_row:
4377 case Intrinsic::nvvm_wmma_m8n32k16_load_a_s8_row_stride:
4378 case Intrinsic::nvvm_wmma_m8n32k16_load_a_u8_row_stride:
4379 case Intrinsic::nvvm_wmma_m8n32k16_load_a_u8_row:
4380 case Intrinsic::nvvm_wmma_m8n8k128_load_a_b1_row:
4381 case Intrinsic::nvvm_wmma_m8n8k128_load_a_b1_row_stride:
4382 case Intrinsic::nvvm_wmma_m8n8k128_load_b_b1_col:
4383 case Intrinsic::nvvm_wmma_m8n8k128_load_b_b1_col_stride:
4384 case Intrinsic::nvvm_wmma_m8n8k32_load_a_s4_row:
4385 case Intrinsic::nvvm_wmma_m8n8k32_load_a_s4_row_stride:
4386 case Intrinsic::nvvm_wmma_m8n8k32_load_a_u4_row_stride:
4387 case Intrinsic::nvvm_wmma_m8n8k32_load_a_u4_row:
4388 case Intrinsic::nvvm_wmma_m8n8k32_load_b_s4_col:
4389 case Intrinsic::nvvm_wmma_m8n8k32_load_b_s4_col_stride:
4390 case Intrinsic::nvvm_wmma_m8n8k32_load_b_u4_col_stride:
4391 case Intrinsic::nvvm_wmma_m8n8k32_load_b_u4_col:
4392 case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x1_b16:
4393 case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x1_trans_b16: {
4403 case Intrinsic::nvvm_wmma_m16n16k16_load_c_f16_col:
4404 case Intrinsic::nvvm_wmma_m16n16k16_load_c_f16_row:
4405 case Intrinsic::nvvm_wmma_m16n16k16_load_c_f16_col_stride:
4406 case Intrinsic::nvvm_wmma_m16n16k16_load_c_f16_row_stride:
4407 case Intrinsic::nvvm_wmma_m32n8k16_load_c_f16_col:
4408 case Intrinsic::nvvm_wmma_m32n8k16_load_c_f16_row:
4409 case Intrinsic::nvvm_wmma_m32n8k16_load_c_f16_col_stride:
4410 case Intrinsic::nvvm_wmma_m32n8k16_load_c_f16_row_stride:
4411 case Intrinsic::nvvm_wmma_m8n32k16_load_c_f16_col:
4412 case Intrinsic::nvvm_wmma_m8n32k16_load_c_f16_row:
4413 case Intrinsic::nvvm_wmma_m8n32k16_load_c_f16_col_stride:
4414 case Intrinsic::nvvm_wmma_m8n32k16_load_c_f16_row_stride: {
4424 case Intrinsic::nvvm_wmma_m16n16k16_load_c_f32_col:
4425 case Intrinsic::nvvm_wmma_m16n16k16_load_c_f32_row:
4426 case Intrinsic::nvvm_wmma_m16n16k16_load_c_f32_col_stride:
4427 case Intrinsic::nvvm_wmma_m16n16k16_load_c_f32_row_stride:
4428 case Intrinsic::nvvm_wmma_m32n8k16_load_c_f32_col:
4429 case Intrinsic::nvvm_wmma_m32n8k16_load_c_f32_row:
4430 case Intrinsic::nvvm_wmma_m32n8k16_load_c_f32_col_stride:
4431 case Intrinsic::nvvm_wmma_m32n8k16_load_c_f32_row_stride:
4432 case Intrinsic::nvvm_wmma_m8n32k16_load_c_f32_col:
4433 case Intrinsic::nvvm_wmma_m8n32k16_load_c_f32_row:
4434 case Intrinsic::nvvm_wmma_m8n32k16_load_c_f32_col_stride:
4435 case Intrinsic::nvvm_wmma_m8n32k16_load_c_f32_row_stride:
4436 case Intrinsic::nvvm_wmma_m16n16k8_load_c_f32_col:
4437 case Intrinsic::nvvm_wmma_m16n16k8_load_c_f32_row:
4438 case Intrinsic::nvvm_wmma_m16n16k8_load_c_f32_col_stride:
4439 case Intrinsic::nvvm_wmma_m16n16k8_load_c_f32_row_stride: {
4449 case Intrinsic::nvvm_wmma_m32n8k16_load_a_bf16_col:
4450 case Intrinsic::nvvm_wmma_m32n8k16_load_a_bf16_col_stride:
4451 case Intrinsic::nvvm_wmma_m32n8k16_load_a_bf16_row:
4452 case Intrinsic::nvvm_wmma_m32n8k16_load_a_bf16_row_stride:
4454 case Intrinsic::nvvm_wmma_m8n32k16_load_b_bf16_col:
4455 case Intrinsic::nvvm_wmma_m8n32k16_load_b_bf16_col_stride:
4456 case Intrinsic::nvvm_wmma_m8n32k16_load_b_bf16_row:
4457 case Intrinsic::nvvm_wmma_m8n32k16_load_b_bf16_row_stride:
4459 case Intrinsic::nvvm_wmma_m16n16k16_load_c_s32_col:
4460 case Intrinsic::nvvm_wmma_m16n16k16_load_c_s32_col_stride:
4461 case Intrinsic::nvvm_wmma_m16n16k16_load_c_s32_row:
4462 case Intrinsic::nvvm_wmma_m16n16k16_load_c_s32_row_stride:
4463 case Intrinsic::nvvm_wmma_m32n8k16_load_c_s32_col:
4464 case Intrinsic::nvvm_wmma_m32n8k16_load_c_s32_col_stride:
4465 case Intrinsic::nvvm_wmma_m32n8k16_load_c_s32_row:
4466 case Intrinsic::nvvm_wmma_m32n8k16_load_c_s32_row_stride:
4467 case Intrinsic::nvvm_wmma_m8n32k16_load_c_s32_col:
4468 case Intrinsic::nvvm_wmma_m8n32k16_load_c_s32_col_stride:
4469 case Intrinsic::nvvm_wmma_m8n32k16_load_c_s32_row:
4470 case Intrinsic::nvvm_wmma_m8n32k16_load_c_s32_row_stride: {
4480 case Intrinsic::nvvm_wmma_m8n8k128_load_c_s32_col:
4481 case Intrinsic::nvvm_wmma_m8n8k128_load_c_s32_col_stride:
4482 case Intrinsic::nvvm_wmma_m8n8k128_load_c_s32_row:
4483 case Intrinsic::nvvm_wmma_m8n8k128_load_c_s32_row_stride:
4484 case Intrinsic::nvvm_wmma_m8n8k32_load_c_s32_col:
4485 case Intrinsic::nvvm_wmma_m8n8k32_load_c_s32_col_stride:
4486 case Intrinsic::nvvm_wmma_m8n8k32_load_c_s32_row:
4487 case Intrinsic::nvvm_wmma_m8n8k32_load_c_s32_row_stride:
4488 case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x2_b16:
4489 case Intrinsic::nvvm_ldmatrix_sync_aligned_m8n8_x2_trans_b16: {
4499 case Intrinsic::nvvm_wmma_m8n8k4_load_a_f64_col:
4500 case Intrinsic::nvvm_wmma_m8n8k4_load_a_f64_col_stride:
4501 case Intrinsic::nvvm_wmma_m8n8k4_load_a_f64_row:
4502 case Intrinsic::nvvm_wmma_m8n8k4_load_a_f64_row_stride:
4504 case Intrinsic::nvvm_wmma_m8n8k4_load_b_f64_col:
4505 case Intrinsic::nvvm_wmma_m8n8k4_load_b_f64_col_stride:
4506 case Intrinsic::nvvm_wmma_m8n8k4_load_b_f64_row:
4507 case Intrinsic::nvvm_wmma_m8n8k4_load_b_f64_row_stride: {
4517 case Intrinsic::nvvm_wmma_m8n8k4_load_c_f64_col:
4518 case Intrinsic::nvvm_wmma_m8n8k4_load_c_f64_col_stride:
4519 case Intrinsic::nvvm_wmma_m8n8k4_load_c_f64_row:
4520 case Intrinsic::nvvm_wmma_m8n8k4_load_c_f64_row_stride: {
4530 case Intrinsic::nvvm_wmma_m16n16k16_store_d_f16_col:
4531 case Intrinsic::nvvm_wmma_m16n16k16_store_d_f16_row:
4532 case Intrinsic::nvvm_wmma_m16n16k16_store_d_f16_col_stride:
4533 case Intrinsic::nvvm_wmma_m16n16k16_store_d_f16_row_stride:
4534 case Intrinsic::nvvm_wmma_m32n8k16_store_d_f16_col:
4535 case Intrinsic::nvvm_wmma_m32n8k16_store_d_f16_row:
4536 case Intrinsic::nvvm_wmma_m32n8k16_store_d_f16_col_stride:
4537 case Intrinsic::nvvm_wmma_m32n8k16_store_d_f16_row_stride:
4538 case Intrinsic::nvvm_wmma_m8n32k16_store_d_f16_col:
4539 case Intrinsic::nvvm_wmma_m8n32k16_store_d_f16_row:
4540 case Intrinsic::nvvm_wmma_m8n32k16_store_d_f16_col_stride:
4541 case Intrinsic::nvvm_wmma_m8n32k16_store_d_f16_row_stride: {
4551 case Intrinsic::nvvm_wmma_m16n16k16_store_d_f32_col:
4552 case Intrinsic::nvvm_wmma_m16n16k16_store_d_f32_row:
4553 case Intrinsic::nvvm_wmma_m16n16k16_store_d_f32_col_stride:
4554 case Intrinsic::nvvm_wmma_m16n16k16_store_d_f32_row_stride:
4555 case Intrinsic::nvvm_wmma_m32n8k16_store_d_f32_col:
4556 case Intrinsic::nvvm_wmma_m32n8k16_store_d_f32_row:
4557 case Intrinsic::nvvm_wmma_m32n8k16_store_d_f32_col_stride:
4558 case Intrinsic::nvvm_wmma_m32n8k16_store_d_f32_row_stride:
4559 case Intrinsic::nvvm_wmma_m8n32k16_store_d_f32_col:
4560 case Intrinsic::nvvm_wmma_m8n32k16_store_d_f32_row:
4561 case Intrinsic::nvvm_wmma_m8n32k16_store_d_f32_col_stride:
4562 case Intrinsic::nvvm_wmma_m8n32k16_store_d_f32_row_stride:
4563 case Intrinsic::nvvm_wmma_m16n16k8_store_d_f32_col:
4564 case Intrinsic::nvvm_wmma_m16n16k8_store_d_f32_row:
4565 case Intrinsic::nvvm_wmma_m16n16k8_store_d_f32_col_stride:
4566 case Intrinsic::nvvm_wmma_m16n16k8_store_d_f32_row_stride: {
4576 case Intrinsic::nvvm_wmma_m16n16k16_store_d_s32_col:
4577 case Intrinsic::nvvm_wmma_m16n16k16_store_d_s32_col_stride:
4578 case Intrinsic::nvvm_wmma_m16n16k16_store_d_s32_row:
4579 case Intrinsic::nvvm_wmma_m16n16k16_store_d_s32_row_stride:
4580 case Intrinsic::nvvm_wmma_m32n8k16_store_d_s32_col:
4581 case Intrinsic::nvvm_wmma_m32n8k16_store_d_s32_col_stride:
4582 case Intrinsic::nvvm_wmma_m32n8k16_store_d_s32_row:
4583 case Intrinsic::nvvm_wmma_m32n8k16_store_d_s32_row_stride:
4584 case Intrinsic::nvvm_wmma_m8n32k16_store_d_s32_col:
4585 case Intrinsic::nvvm_wmma_m8n32k16_store_d_s32_col_stride:
4586 case Intrinsic::nvvm_wmma_m8n32k16_store_d_s32_row:
4587 case Intrinsic::nvvm_wmma_m8n32k16_store_d_s32_row_stride: {
4597 case Intrinsic::nvvm_wmma_m8n8k128_store_d_s32_col:
4598 case Intrinsic::nvvm_wmma_m8n8k128_store_d_s32_col_stride:
4599 case Intrinsic::nvvm_wmma_m8n8k128_store_d_s32_row:
4600 case Intrinsic::nvvm_wmma_m8n8k128_store_d_s32_row_stride:
4601 case Intrinsic::nvvm_wmma_m8n8k32_store_d_s32_col:
4602 case Intrinsic::nvvm_wmma_m8n8k32_store_d_s32_col_stride:
4603 case Intrinsic::nvvm_wmma_m8n8k32_store_d_s32_row:
4604 case Intrinsic::nvvm_wmma_m8n8k32_store_d_s32_row_stride: {
4614 case Intrinsic::nvvm_wmma_m8n8k4_store_d_f64_col:
4615 case Intrinsic::nvvm_wmma_m8n8k4_store_d_f64_col_stride:
4616 case Intrinsic::nvvm_wmma_m8n8k4_store_d_f64_row:
4617 case Intrinsic::nvvm_wmma_m8n8k4_store_d_f64_row_stride: {
4627 case Intrinsic::nvvm_atomic_load_inc_32:
4628 case Intrinsic::nvvm_atomic_load_dec_32:
4630 case Intrinsic::nvvm_atomic_add_gen_f_cta:
4631 case Intrinsic::nvvm_atomic_add_gen_f_sys:
4632 case Intrinsic::nvvm_atomic_add_gen_i_cta:
4633 case Intrinsic::nvvm_atomic_add_gen_i_sys:
4634 case Intrinsic::nvvm_atomic_and_gen_i_cta:
4635 case Intrinsic::nvvm_atomic_and_gen_i_sys:
4636 case Intrinsic::nvvm_atomic_cas_gen_i_cta:
4637 case Intrinsic::nvvm_atomic_cas_gen_i_sys:
4638 case Intrinsic::nvvm_atomic_dec_gen_i_cta:
4639 case Intrinsic::nvvm_atomic_dec_gen_i_sys:
4640 case Intrinsic::nvvm_atomic_inc_gen_i_cta:
4641 case Intrinsic::nvvm_atomic_inc_gen_i_sys:
4642 case Intrinsic::nvvm_atomic_max_gen_i_cta:
4643 case Intrinsic::nvvm_atomic_max_gen_i_sys:
4644 case Intrinsic::nvvm_atomic_min_gen_i_cta:
4645 case Intrinsic::nvvm_atomic_min_gen_i_sys:
4646 case Intrinsic::nvvm_atomic_or_gen_i_cta:
4647 case Intrinsic::nvvm_atomic_or_gen_i_sys:
4648 case Intrinsic::nvvm_atomic_exch_gen_i_cta:
4649 case Intrinsic::nvvm_atomic_exch_gen_i_sys:
4650 case Intrinsic::nvvm_atomic_xor_gen_i_cta:
4651 case Intrinsic::nvvm_atomic_xor_gen_i_sys: {
4662 case Intrinsic::nvvm_ldu_global_i:
4663 case Intrinsic::nvvm_ldu_global_f:
4664 case Intrinsic::nvvm_ldu_global_p: {
4667 if (Intrinsic == Intrinsic::nvvm_ldu_global_i)
4669 else if(Intrinsic == Intrinsic::nvvm_ldu_global_p)
4680 case Intrinsic::nvvm_ldg_global_i:
4681 case Intrinsic::nvvm_ldg_global_f:
4682 case Intrinsic::nvvm_ldg_global_p: {
4686 if (Intrinsic == Intrinsic::nvvm_ldg_global_i)
4688 else if(Intrinsic == Intrinsic::nvvm_ldg_global_p)
4700 case Intrinsic::nvvm_tex_1d_v4f32_s32:
4701 case Intrinsic::nvvm_tex_1d_v4f32_f32:
4702 case Intrinsic::nvvm_tex_1d_level_v4f32_f32:
4703 case Intrinsic::nvvm_tex_1d_grad_v4f32_f32:
4704 case Intrinsic::nvvm_tex_1d_array_v4f32_s32:
4705 case Intrinsic::nvvm_tex_1d_array_v4f32_f32:
4706 case Intrinsic::nvvm_tex_1d_array_level_v4f32_f32:
4707 case Intrinsic::nvvm_tex_1d_array_grad_v4f32_f32:
4708 case Intrinsic::nvvm_tex_2d_v4f32_s32:
4709 case Intrinsic::nvvm_tex_2d_v4f32_f32:
4710 case Intrinsic::nvvm_tex_2d_level_v4f32_f32:
4711 case Intrinsic::nvvm_tex_2d_grad_v4f32_f32:
4712 case Intrinsic::nvvm_tex_2d_array_v4f32_s32:
4713 case Intrinsic::nvvm_tex_2d_array_v4f32_f32:
4714 case Intrinsic::nvvm_tex_2d_array_level_v4f32_f32:
4715 case Intrinsic::nvvm_tex_2d_array_grad_v4f32_f32:
4716 case Intrinsic::nvvm_tex_3d_v4f32_s32:
4717 case Intrinsic::nvvm_tex_3d_v4f32_f32:
4718 case Intrinsic::nvvm_tex_3d_level_v4f32_f32:
4719 case Intrinsic::nvvm_tex_3d_grad_v4f32_f32:
4720 case Intrinsic::nvvm_tex_cube_v4f32_f32:
4721 case Intrinsic::nvvm_tex_cube_level_v4f32_f32:
4722 case Intrinsic::nvvm_tex_cube_array_v4f32_f32:
4723 case Intrinsic::nvvm_tex_cube_array_level_v4f32_f32:
4724 case Intrinsic::nvvm_tld4_r_2d_v4f32_f32:
4725 case Intrinsic::nvvm_tld4_g_2d_v4f32_f32:
4726 case Intrinsic::nvvm_tld4_b_2d_v4f32_f32:
4727 case Intrinsic::nvvm_tld4_a_2d_v4f32_f32:
4728 case Intrinsic::nvvm_tex_unified_1d_v4f32_s32:
4729 case Intrinsic::nvvm_tex_unified_1d_v4f32_f32:
4730 case Intrinsic::nvvm_tex_unified_1d_level_v4f32_f32:
4731 case Intrinsic::nvvm_tex_unified_1d_grad_v4f32_f32:
4732 case Intrinsic::nvvm_tex_unified_1d_array_v4f32_s32:
4733 case Intrinsic::nvvm_tex_unified_1d_array_v4f32_f32:
4734 case Intrinsic::nvvm_tex_unified_1d_array_level_v4f32_f32:
4735 case Intrinsic::nvvm_tex_unified_1d_array_grad_v4f32_f32:
4736 case Intrinsic::nvvm_tex_unified_2d_v4f32_s32:
4737 case Intrinsic::nvvm_tex_unified_2d_v4f32_f32:
4738 case Intrinsic::nvvm_tex_unified_2d_level_v4f32_f32:
4739 case Intrinsic::nvvm_tex_unified_2d_grad_v4f32_f32:
4740 case Intrinsic::nvvm_tex_unified_2d_array_v4f32_s32:
4741 case Intrinsic::nvvm_tex_unified_2d_array_v4f32_f32:
4742 case Intrinsic::nvvm_tex_unified_2d_array_level_v4f32_f32:
4743 case Intrinsic::nvvm_tex_unified_2d_array_grad_v4f32_f32:
4744 case Intrinsic::nvvm_tex_unified_3d_v4f32_s32:
4745 case Intrinsic::nvvm_tex_unified_3d_v4f32_f32:
4746 case Intrinsic::nvvm_tex_unified_3d_level_v4f32_f32:
4747 case Intrinsic::nvvm_tex_unified_3d_grad_v4f32_f32:
4748 case Intrinsic::nvvm_tex_unified_cube_v4f32_f32:
4749 case Intrinsic::nvvm_tex_unified_cube_level_v4f32_f32:
4750 case Intrinsic::nvvm_tex_unified_cube_array_v4f32_f32:
4751 case Intrinsic::nvvm_tex_unified_cube_array_level_v4f32_f32:
4752 case Intrinsic::nvvm_tex_unified_cube_grad_v4f32_f32:
4753 case Intrinsic::nvvm_tex_unified_cube_array_grad_v4f32_f32:
4754 case Intrinsic::nvvm_tld4_unified_r_2d_v4f32_f32:
4755 case Intrinsic::nvvm_tld4_unified_g_2d_v4f32_f32:
4756 case Intrinsic::nvvm_tld4_unified_b_2d_v4f32_f32:
4757 case Intrinsic::nvvm_tld4_unified_a_2d_v4f32_f32:
4758 Info.opc = getOpcForTextureInstr(Intrinsic);
4766 case Intrinsic::nvvm_tex_1d_v4s32_s32:
4767 case Intrinsic::nvvm_tex_1d_v4s32_f32:
4768 case Intrinsic::nvvm_tex_1d_level_v4s32_f32:
4769 case Intrinsic::nvvm_tex_1d_grad_v4s32_f32:
4770 case Intrinsic::nvvm_tex_1d_array_v4s32_s32:
4771 case Intrinsic::nvvm_tex_1d_array_v4s32_f32:
4772 case Intrinsic::nvvm_tex_1d_array_level_v4s32_f32:
4773 case Intrinsic::nvvm_tex_1d_array_grad_v4s32_f32:
4774 case Intrinsic::nvvm_tex_2d_v4s32_s32:
4775 case Intrinsic::nvvm_tex_2d_v4s32_f32:
4776 case Intrinsic::nvvm_tex_2d_level_v4s32_f32:
4777 case Intrinsic::nvvm_tex_2d_grad_v4s32_f32:
4778 case Intrinsic::nvvm_tex_2d_array_v4s32_s32:
4779 case Intrinsic::nvvm_tex_2d_array_v4s32_f32:
4780 case Intrinsic::nvvm_tex_2d_array_level_v4s32_f32:
4781 case Intrinsic::nvvm_tex_2d_array_grad_v4s32_f32:
4782 case Intrinsic::nvvm_tex_3d_v4s32_s32:
4783 case Intrinsic::nvvm_tex_3d_v4s32_f32:
4784 case Intrinsic::nvvm_tex_3d_level_v4s32_f32:
4785 case Intrinsic::nvvm_tex_3d_grad_v4s32_f32:
4786 case Intrinsic::nvvm_tex_cube_v4s32_f32:
4787 case Intrinsic::nvvm_tex_cube_level_v4s32_f32:
4788 case Intrinsic::nvvm_tex_cube_array_v4s32_f32:
4789 case Intrinsic::nvvm_tex_cube_array_level_v4s32_f32:
4790 case Intrinsic::nvvm_tex_cube_v4u32_f32:
4791 case Intrinsic::nvvm_tex_cube_level_v4u32_f32:
4792 case Intrinsic::nvvm_tex_cube_array_v4u32_f32:
4793 case Intrinsic::nvvm_tex_cube_array_level_v4u32_f32:
4794 case Intrinsic::nvvm_tex_1d_v4u32_s32:
4795 case Intrinsic::nvvm_tex_1d_v4u32_f32:
4796 case Intrinsic::nvvm_tex_1d_level_v4u32_f32:
4797 case Intrinsic::nvvm_tex_1d_grad_v4u32_f32:
4798 case Intrinsic::nvvm_tex_1d_array_v4u32_s32:
4799 case Intrinsic::nvvm_tex_1d_array_v4u32_f32:
4800 case Intrinsic::nvvm_tex_1d_array_level_v4u32_f32:
4801 case Intrinsic::nvvm_tex_1d_array_grad_v4u32_f32:
4802 case Intrinsic::nvvm_tex_2d_v4u32_s32:
4803 case Intrinsic::nvvm_tex_2d_v4u32_f32:
4804 case Intrinsic::nvvm_tex_2d_level_v4u32_f32:
4805 case Intrinsic::nvvm_tex_2d_grad_v4u32_f32:
4806 case Intrinsic::nvvm_tex_2d_array_v4u32_s32:
4807 case Intrinsic::nvvm_tex_2d_array_v4u32_f32:
4808 case Intrinsic::nvvm_tex_2d_array_level_v4u32_f32:
4809 case Intrinsic::nvvm_tex_2d_array_grad_v4u32_f32:
4810 case Intrinsic::nvvm_tex_3d_v4u32_s32:
4811 case Intrinsic::nvvm_tex_3d_v4u32_f32:
4812 case Intrinsic::nvvm_tex_3d_level_v4u32_f32:
4813 case Intrinsic::nvvm_tex_3d_grad_v4u32_f32:
4814 case Intrinsic::nvvm_tld4_r_2d_v4s32_f32:
4815 case Intrinsic::nvvm_tld4_g_2d_v4s32_f32:
4816 case Intrinsic::nvvm_tld4_b_2d_v4s32_f32:
4817 case Intrinsic::nvvm_tld4_a_2d_v4s32_f32:
4818 case Intrinsic::nvvm_tld4_r_2d_v4u32_f32:
4819 case Intrinsic::nvvm_tld4_g_2d_v4u32_f32:
4820 case Intrinsic::nvvm_tld4_b_2d_v4u32_f32:
4821 case Intrinsic::nvvm_tld4_a_2d_v4u32_f32:
4822 case Intrinsic::nvvm_tex_unified_1d_v4s32_s32:
4823 case Intrinsic::nvvm_tex_unified_1d_v4s32_f32:
4824 case Intrinsic::nvvm_tex_unified_1d_level_v4s32_f32:
4825 case Intrinsic::nvvm_tex_unified_1d_grad_v4s32_f32:
4826 case Intrinsic::nvvm_tex_unified_1d_array_v4s32_s32:
4827 case Intrinsic::nvvm_tex_unified_1d_array_v4s32_f32:
4828 case Intrinsic::nvvm_tex_unified_1d_array_level_v4s32_f32:
4829 case Intrinsic::nvvm_tex_unified_1d_array_grad_v4s32_f32:
4830 case Intrinsic::nvvm_tex_unified_2d_v4s32_s32:
4831 case Intrinsic::nvvm_tex_unified_2d_v4s32_f32:
4832 case Intrinsic::nvvm_tex_unified_2d_level_v4s32_f32:
4833 case Intrinsic::nvvm_tex_unified_2d_grad_v4s32_f32:
4834 case Intrinsic::nvvm_tex_unified_2d_array_v4s32_s32:
4835 case Intrinsic::nvvm_tex_unified_2d_array_v4s32_f32:
4836 case Intrinsic::nvvm_tex_unified_2d_array_level_v4s32_f32:
4837 case Intrinsic::nvvm_tex_unified_2d_array_grad_v4s32_f32:
4838 case Intrinsic::nvvm_tex_unified_3d_v4s32_s32:
4839 case Intrinsic::nvvm_tex_unified_3d_v4s32_f32:
4840 case Intrinsic::nvvm_tex_unified_3d_level_v4s32_f32:
4841 case Intrinsic::nvvm_tex_unified_3d_grad_v4s32_f32:
4842 case Intrinsic::nvvm_tex_unified_1d_v4u32_s32:
4843 case Intrinsic::nvvm_tex_unified_1d_v4u32_f32:
4844 case Intrinsic::nvvm_tex_unified_1d_level_v4u32_f32:
4845 case Intrinsic::nvvm_tex_unified_1d_grad_v4u32_f32:
4846 case Intrinsic::nvvm_tex_unified_1d_array_v4u32_s32:
4847 case Intrinsic::nvvm_tex_unified_1d_array_v4u32_f32:
4848 case Intrinsic::nvvm_tex_unified_1d_array_level_v4u32_f32:
4849 case Intrinsic::nvvm_tex_unified_1d_array_grad_v4u32_f32:
4850 case Intrinsic::nvvm_tex_unified_2d_v4u32_s32:
4851 case Intrinsic::nvvm_tex_unified_2d_v4u32_f32:
4852 case Intrinsic::nvvm_tex_unified_2d_level_v4u32_f32:
4853 case Intrinsic::nvvm_tex_unified_2d_grad_v4u32_f32:
4854 case Intrinsic::nvvm_tex_unified_2d_array_v4u32_s32:
4855 case Intrinsic::nvvm_tex_unified_2d_array_v4u32_f32:
4856 case Intrinsic::nvvm_tex_unified_2d_array_level_v4u32_f32:
4857 case Intrinsic::nvvm_tex_unified_2d_array_grad_v4u32_f32:
4858 case Intrinsic::nvvm_tex_unified_3d_v4u32_s32:
4859 case Intrinsic::nvvm_tex_unified_3d_v4u32_f32:
4860 case Intrinsic::nvvm_tex_unified_3d_level_v4u32_f32:
4861 case Intrinsic::nvvm_tex_unified_3d_grad_v4u32_f32:
4862 case Intrinsic::nvvm_tex_unified_cube_v4s32_f32:
4863 case Intrinsic::nvvm_tex_unified_cube_level_v4s32_f32:
4864 case Intrinsic::nvvm_tex_unified_cube_array_v4s32_f32:
4865 case Intrinsic::nvvm_tex_unified_cube_array_level_v4s32_f32:
4866 case Intrinsic::nvvm_tex_unified_cube_v4u32_f32:
4867 case Intrinsic::nvvm_tex_unified_cube_level_v4u32_f32:
4868 case Intrinsic::nvvm_tex_unified_cube_array_v4u32_f32:
4869 case Intrinsic::nvvm_tex_unified_cube_array_level_v4u32_f32:
4870 case Intrinsic::nvvm_tex_unified_cube_grad_v4s32_f32:
4871 case Intrinsic::nvvm_tex_unified_cube_grad_v4u32_f32:
4872 case Intrinsic::nvvm_tex_unified_cube_array_grad_v4s32_f32:
4873 case Intrinsic::nvvm_tex_unified_cube_array_grad_v4u32_f32:
4874 case Intrinsic::nvvm_tld4_unified_r_2d_v4s32_f32:
4875 case Intrinsic::nvvm_tld4_unified_g_2d_v4s32_f32:
4876 case Intrinsic::nvvm_tld4_unified_b_2d_v4s32_f32:
4877 case Intrinsic::nvvm_tld4_unified_a_2d_v4s32_f32:
4878 case Intrinsic::nvvm_tld4_unified_r_2d_v4u32_f32:
4879 case Intrinsic::nvvm_tld4_unified_g_2d_v4u32_f32:
4880 case Intrinsic::nvvm_tld4_unified_b_2d_v4u32_f32:
4881 case Intrinsic::nvvm_tld4_unified_a_2d_v4u32_f32:
4882 Info.opc = getOpcForTextureInstr(Intrinsic);
4890 case Intrinsic::nvvm_suld_1d_i8_clamp:
4891 case Intrinsic::nvvm_suld_1d_v2i8_clamp:
4892 case Intrinsic::nvvm_suld_1d_v4i8_clamp:
4893 case Intrinsic::nvvm_suld_1d_array_i8_clamp:
4894 case Intrinsic::nvvm_suld_1d_array_v2i8_clamp:
4895 case Intrinsic::nvvm_suld_1d_array_v4i8_clamp:
4896 case Intrinsic::nvvm_suld_2d_i8_clamp:
4897 case Intrinsic::nvvm_suld_2d_v2i8_clamp:
4898 case Intrinsic::nvvm_suld_2d_v4i8_clamp:
4899 case Intrinsic::nvvm_suld_2d_array_i8_clamp:
4900 case Intrinsic::nvvm_suld_2d_array_v2i8_clamp:
4901 case Intrinsic::nvvm_suld_2d_array_v4i8_clamp:
4902 case Intrinsic::nvvm_suld_3d_i8_clamp:
4903 case Intrinsic::nvvm_suld_3d_v2i8_clamp:
4904 case Intrinsic::nvvm_suld_3d_v4i8_clamp:
4905 case Intrinsic::nvvm_suld_1d_i8_trap:
4906 case Intrinsic::nvvm_suld_1d_v2i8_trap:
4907 case Intrinsic::nvvm_suld_1d_v4i8_trap:
4908 case Intrinsic::nvvm_suld_1d_array_i8_trap:
4909 case Intrinsic::nvvm_suld_1d_array_v2i8_trap:
4910 case Intrinsic::nvvm_suld_1d_array_v4i8_trap:
4911 case Intrinsic::nvvm_suld_2d_i8_trap:
4912 case Intrinsic::nvvm_suld_2d_v2i8_trap:
4913 case Intrinsic::nvvm_suld_2d_v4i8_trap:
4914 case Intrinsic::nvvm_suld_2d_array_i8_trap:
4915 case Intrinsic::nvvm_suld_2d_array_v2i8_trap:
4916 case Intrinsic::nvvm_suld_2d_array_v4i8_trap:
4917 case Intrinsic::nvvm_suld_3d_i8_trap:
4918 case Intrinsic::nvvm_suld_3d_v2i8_trap:
4919 case Intrinsic::nvvm_suld_3d_v4i8_trap:
4920 case Intrinsic::nvvm_suld_1d_i8_zero:
4921 case Intrinsic::nvvm_suld_1d_v2i8_zero:
4922 case Intrinsic::nvvm_suld_1d_v4i8_zero:
4923 case Intrinsic::nvvm_suld_1d_array_i8_zero:
4924 case Intrinsic::nvvm_suld_1d_array_v2i8_zero:
4925 case Intrinsic::nvvm_suld_1d_array_v4i8_zero:
4926 case Intrinsic::nvvm_suld_2d_i8_zero:
4927 case Intrinsic::nvvm_suld_2d_v2i8_zero:
4928 case Intrinsic::nvvm_suld_2d_v4i8_zero:
4929 case Intrinsic::nvvm_suld_2d_array_i8_zero:
4930 case Intrinsic::nvvm_suld_2d_array_v2i8_zero:
4931 case Intrinsic::nvvm_suld_2d_array_v4i8_zero:
4932 case Intrinsic::nvvm_suld_3d_i8_zero:
4933 case Intrinsic::nvvm_suld_3d_v2i8_zero:
4934 case Intrinsic::nvvm_suld_3d_v4i8_zero:
4935 Info.opc = getOpcForSurfaceInstr(Intrinsic);
4943 case Intrinsic::nvvm_suld_1d_i16_clamp:
4944 case Intrinsic::nvvm_suld_1d_v2i16_clamp:
4945 case Intrinsic::nvvm_suld_1d_v4i16_clamp:
4946 case Intrinsic::nvvm_suld_1d_array_i16_clamp:
4947 case Intrinsic::nvvm_suld_1d_array_v2i16_clamp:
4948 case Intrinsic::nvvm_suld_1d_array_v4i16_clamp:
4949 case Intrinsic::nvvm_suld_2d_i16_clamp:
4950 case Intrinsic::nvvm_suld_2d_v2i16_clamp:
4951 case Intrinsic::nvvm_suld_2d_v4i16_clamp:
4952 case Intrinsic::nvvm_suld_2d_array_i16_clamp:
4953 case Intrinsic::nvvm_suld_2d_array_v2i16_clamp:
4954 case Intrinsic::nvvm_suld_2d_array_v4i16_clamp:
4955 case Intrinsic::nvvm_suld_3d_i16_clamp:
4956 case Intrinsic::nvvm_suld_3d_v2i16_clamp:
4957 case Intrinsic::nvvm_suld_3d_v4i16_clamp:
4958 case Intrinsic::nvvm_suld_1d_i16_trap:
4959 case Intrinsic::nvvm_suld_1d_v2i16_trap:
4960 case Intrinsic::nvvm_suld_1d_v4i16_trap:
4961 case Intrinsic::nvvm_suld_1d_array_i16_trap:
4962 case Intrinsic::nvvm_suld_1d_array_v2i16_trap:
4963 case Intrinsic::nvvm_suld_1d_array_v4i16_trap:
4964 case Intrinsic::nvvm_suld_2d_i16_trap:
4965 case Intrinsic::nvvm_suld_2d_v2i16_trap:
4966 case Intrinsic::nvvm_suld_2d_v4i16_trap:
4967 case Intrinsic::nvvm_suld_2d_array_i16_trap:
4968 case Intrinsic::nvvm_suld_2d_array_v2i16_trap:
4969 case Intrinsic::nvvm_suld_2d_array_v4i16_trap:
4970 case Intrinsic::nvvm_suld_3d_i16_trap:
4971 case Intrinsic::nvvm_suld_3d_v2i16_trap:
4972 case Intrinsic::nvvm_suld_3d_v4i16_trap:
4973 case Intrinsic::nvvm_suld_1d_i16_zero:
4974 case Intrinsic::nvvm_suld_1d_v2i16_zero:
4975 case Intrinsic::nvvm_suld_1d_v4i16_zero:
4976 case Intrinsic::nvvm_suld_1d_array_i16_zero:
4977 case Intrinsic::nvvm_suld_1d_array_v2i16_zero:
4978 case Intrinsic::nvvm_suld_1d_array_v4i16_zero:
4979 case Intrinsic::nvvm_suld_2d_i16_zero:
4980 case Intrinsic::nvvm_suld_2d_v2i16_zero:
4981 case Intrinsic::nvvm_suld_2d_v4i16_zero:
4982 case Intrinsic::nvvm_suld_2d_array_i16_zero:
4983 case Intrinsic::nvvm_suld_2d_array_v2i16_zero:
4984 case Intrinsic::nvvm_suld_2d_array_v4i16_zero:
4985 case Intrinsic::nvvm_suld_3d_i16_zero:
4986 case Intrinsic::nvvm_suld_3d_v2i16_zero:
4987 case Intrinsic::nvvm_suld_3d_v4i16_zero:
4988 Info.opc = getOpcForSurfaceInstr(Intrinsic);
4996 case Intrinsic::nvvm_suld_1d_i32_clamp:
4997 case Intrinsic::nvvm_suld_1d_v2i32_clamp:
4998 case Intrinsic::nvvm_suld_1d_v4i32_clamp:
4999 case Intrinsic::nvvm_suld_1d_array_i32_clamp:
5000 case Intrinsic::nvvm_suld_1d_array_v2i32_clamp:
5001 case Intrinsic::nvvm_suld_1d_array_v4i32_clamp:
5002 case Intrinsic::nvvm_suld_2d_i32_clamp:
5003 case Intrinsic::nvvm_suld_2d_v2i32_clamp:
5004 case Intrinsic::nvvm_suld_2d_v4i32_clamp:
5005 case Intrinsic::nvvm_suld_2d_array_i32_clamp:
5006 case Intrinsic::nvvm_suld_2d_array_v2i32_clamp:
5007 case Intrinsic::nvvm_suld_2d_array_v4i32_clamp:
5008 case Intrinsic::nvvm_suld_3d_i32_clamp:
5009 case Intrinsic::nvvm_suld_3d_v2i32_clamp:
5010 case Intrinsic::nvvm_suld_3d_v4i32_clamp:
5011 case Intrinsic::nvvm_suld_1d_i32_trap:
5012 case Intrinsic::nvvm_suld_1d_v2i32_trap:
5013 case Intrinsic::nvvm_suld_1d_v4i32_trap:
5014 case Intrinsic::nvvm_suld_1d_array_i32_trap:
5015 case Intrinsic::nvvm_suld_1d_array_v2i32_trap:
5016 case Intrinsic::nvvm_suld_1d_array_v4i32_trap:
5017 case Intrinsic::nvvm_suld_2d_i32_trap:
5018 case Intrinsic::nvvm_suld_2d_v2i32_trap:
5019 case Intrinsic::nvvm_suld_2d_v4i32_trap:
5020 case Intrinsic::nvvm_suld_2d_array_i32_trap:
5021 case Intrinsic::nvvm_suld_2d_array_v2i32_trap:
5022 case Intrinsic::nvvm_suld_2d_array_v4i32_trap:
5023 case Intrinsic::nvvm_suld_3d_i32_trap:
5024 case Intrinsic::nvvm_suld_3d_v2i32_trap:
5025 case Intrinsic::nvvm_suld_3d_v4i32_trap:
5026 case Intrinsic::nvvm_suld_1d_i32_zero:
5027 case Intrinsic::nvvm_suld_1d_v2i32_zero:
5028 case Intrinsic::nvvm_suld_1d_v4i32_zero:
5029 case Intrinsic::nvvm_suld_1d_array_i32_zero:
5030 case Intrinsic::nvvm_suld_1d_array_v2i32_zero:
5031 case Intrinsic::nvvm_suld_1d_array_v4i32_zero:
5032 case Intrinsic::nvvm_suld_2d_i32_zero:
5033 case Intrinsic::nvvm_suld_2d_v2i32_zero:
5034 case Intrinsic::nvvm_suld_2d_v4i32_zero:
5035 case Intrinsic::nvvm_suld_2d_array_i32_zero:
5036 case Intrinsic::nvvm_suld_2d_array_v2i32_zero:
5037 case Intrinsic::nvvm_suld_2d_array_v4i32_zero:
5038 case Intrinsic::nvvm_suld_3d_i32_zero:
5039 case Intrinsic::nvvm_suld_3d_v2i32_zero:
5040 case Intrinsic::nvvm_suld_3d_v4i32_zero:
5041 Info.opc = getOpcForSurfaceInstr(Intrinsic);
5049 case Intrinsic::nvvm_suld_1d_i64_clamp:
5050 case Intrinsic::nvvm_suld_1d_v2i64_clamp:
5051 case Intrinsic::nvvm_suld_1d_array_i64_clamp:
5052 case Intrinsic::nvvm_suld_1d_array_v2i64_clamp:
5053 case Intrinsic::nvvm_suld_2d_i64_clamp:
5054 case Intrinsic::nvvm_suld_2d_v2i64_clamp:
5055 case Intrinsic::nvvm_suld_2d_array_i64_clamp:
5056 case Intrinsic::nvvm_suld_2d_array_v2i64_clamp:
5057 case Intrinsic::nvvm_suld_3d_i64_clamp:
5058 case Intrinsic::nvvm_suld_3d_v2i64_clamp:
5059 case Intrinsic::nvvm_suld_1d_i64_trap:
5060 case Intrinsic::nvvm_suld_1d_v2i64_trap:
5061 case Intrinsic::nvvm_suld_1d_array_i64_trap:
5062 case Intrinsic::nvvm_suld_1d_array_v2i64_trap:
5063 case Intrinsic::nvvm_suld_2d_i64_trap:
5064 case Intrinsic::nvvm_suld_2d_v2i64_trap:
5065 case Intrinsic::nvvm_suld_2d_array_i64_trap:
5066 case Intrinsic::nvvm_suld_2d_array_v2i64_trap:
5067 case Intrinsic::nvvm_suld_3d_i64_trap:
5068 case Intrinsic::nvvm_suld_3d_v2i64_trap:
5069 case Intrinsic::nvvm_suld_1d_i64_zero:
5070 case Intrinsic::nvvm_suld_1d_v2i64_zero:
5071 case Intrinsic::nvvm_suld_1d_array_i64_zero:
5072 case Intrinsic::nvvm_suld_1d_array_v2i64_zero:
5073 case Intrinsic::nvvm_suld_2d_i64_zero:
5074 case Intrinsic::nvvm_suld_2d_v2i64_zero:
5075 case Intrinsic::nvvm_suld_2d_array_i64_zero:
5076 case Intrinsic::nvvm_suld_2d_array_v2i64_zero:
5077 case Intrinsic::nvvm_suld_3d_i64_zero:
5078 case Intrinsic::nvvm_suld_3d_v2i64_zero:
5079 Info.opc = getOpcForSurfaceInstr(Intrinsic);
6207 case Intrinsic::nvvm_ldg_global_i:
6208 case Intrinsic::nvvm_ldg_global_f:
6209 case Intrinsic::nvvm_ldg_global_p:
6210 case Intrinsic::nvvm_ldu_global_i:
6211 case Intrinsic::nvvm_ldu_global_f:
6212 case Intrinsic::nvvm_ldu_global_p: {
6241 case Intrinsic::nvvm_ldg_global_i:
6242 case Intrinsic::nvvm_ldg_global_f:
6243 case Intrinsic::nvvm_ldg_global_p:
6246 case Intrinsic::nvvm_ldu_global_i:
6247 case Intrinsic::nvvm_ldu_global_f:
6248 case Intrinsic::nvvm_ldu_global_p:
6258 case Intrinsic::nvvm_ldg_global_i:
6259 case Intrinsic::nvvm_ldg_global_f:
6260 case Intrinsic::nvvm_ldg_global_p:
6263 case Intrinsic::nvvm_ldu_global_i:
6264 case Intrinsic::nvvm_ldu_global_f:
6265 case Intrinsic::nvvm_ldu_global_p: