303   case AMDGPU::BI__builtin_amdgcn_div_scale:  in EmitAMDGPUBuiltinExpr()
304   case AMDGPU::BI__builtin_amdgcn_div_scalef: {  in EmitAMDGPUBuiltinExpr()
328   case AMDGPU::BI__builtin_amdgcn_div_fmas:  in EmitAMDGPUBuiltinExpr()
329   case AMDGPU::BI__builtin_amdgcn_div_fmasf: {  in EmitAMDGPUBuiltinExpr()
341   case AMDGPU::BI__builtin_amdgcn_ds_swizzle:  in EmitAMDGPUBuiltinExpr()
344   case AMDGPU::BI__builtin_amdgcn_mov_dpp8:  in EmitAMDGPUBuiltinExpr()
345   case AMDGPU::BI__builtin_amdgcn_mov_dpp:  in EmitAMDGPUBuiltinExpr()
346   case AMDGPU::BI__builtin_amdgcn_update_dpp: {  in EmitAMDGPUBuiltinExpr()
359         CGM.getIntrinsic(BuiltinID == AMDGPU::BI__builtin_amdgcn_mov_dpp8  in EmitAMDGPUBuiltinExpr()
365     bool InsertOld = BuiltinID == AMDGPU::BI__builtin_amdgcn_mov_dpp;  in EmitAMDGPUBuiltinExpr()
370       if (I < (BuiltinID == AMDGPU::BI__builtin_amdgcn_update_dpp ? 2u : 1u) &&  in EmitAMDGPUBuiltinExpr()
387   case AMDGPU::BI__builtin_amdgcn_permlane16:  in EmitAMDGPUBuiltinExpr()
388   case AMDGPU::BI__builtin_amdgcn_permlanex16:  in EmitAMDGPUBuiltinExpr()
391         BuiltinID == AMDGPU::BI__builtin_amdgcn_permlane16  in EmitAMDGPUBuiltinExpr()
394   case AMDGPU::BI__builtin_amdgcn_permlane64:  in EmitAMDGPUBuiltinExpr()
397   case AMDGPU::BI__builtin_amdgcn_readlane:  in EmitAMDGPUBuiltinExpr()
400   case AMDGPU::BI__builtin_amdgcn_readfirstlane:  in EmitAMDGPUBuiltinExpr()
403   case AMDGPU::BI__builtin_amdgcn_div_fixup:  in EmitAMDGPUBuiltinExpr()
404   case AMDGPU::BI__builtin_amdgcn_div_fixupf:  in EmitAMDGPUBuiltinExpr()
405   case AMDGPU::BI__builtin_amdgcn_div_fixuph:  in EmitAMDGPUBuiltinExpr()
408   case AMDGPU::BI__builtin_amdgcn_trig_preop:  in EmitAMDGPUBuiltinExpr()
409   case AMDGPU::BI__builtin_amdgcn_trig_preopf:  in EmitAMDGPUBuiltinExpr()
411   case AMDGPU::BI__builtin_amdgcn_rcp:  in EmitAMDGPUBuiltinExpr()
412   case AMDGPU::BI__builtin_amdgcn_rcpf:  in EmitAMDGPUBuiltinExpr()
413   case AMDGPU::BI__builtin_amdgcn_rcph:  in EmitAMDGPUBuiltinExpr()
415   case AMDGPU::BI__builtin_amdgcn_sqrt:  in EmitAMDGPUBuiltinExpr()
416   case AMDGPU::BI__builtin_amdgcn_sqrtf:  in EmitAMDGPUBuiltinExpr()
417   case AMDGPU::BI__builtin_amdgcn_sqrth:  in EmitAMDGPUBuiltinExpr()
420   case AMDGPU::BI__builtin_amdgcn_rsq:  in EmitAMDGPUBuiltinExpr()
421   case AMDGPU::BI__builtin_amdgcn_rsqf:  in EmitAMDGPUBuiltinExpr()
422   case AMDGPU::BI__builtin_amdgcn_rsqh:  in EmitAMDGPUBuiltinExpr()
424   case AMDGPU::BI__builtin_amdgcn_rsq_clamp:  in EmitAMDGPUBuiltinExpr()
425   case AMDGPU::BI__builtin_amdgcn_rsq_clampf:  in EmitAMDGPUBuiltinExpr()
428   case AMDGPU::BI__builtin_amdgcn_sinf:  in EmitAMDGPUBuiltinExpr()
429   case AMDGPU::BI__builtin_amdgcn_sinh:  in EmitAMDGPUBuiltinExpr()
431   case AMDGPU::BI__builtin_amdgcn_cosf:  in EmitAMDGPUBuiltinExpr()
432   case AMDGPU::BI__builtin_amdgcn_cosh:  in EmitAMDGPUBuiltinExpr()
434   case AMDGPU::BI__builtin_amdgcn_dispatch_ptr:  in EmitAMDGPUBuiltinExpr()
436   case AMDGPU::BI__builtin_amdgcn_logf:  in EmitAMDGPUBuiltinExpr()
438   case AMDGPU::BI__builtin_amdgcn_exp2f:  in EmitAMDGPUBuiltinExpr()
441   case AMDGPU::BI__builtin_amdgcn_log_clampf:  in EmitAMDGPUBuiltinExpr()
444   case AMDGPU::BI__builtin_amdgcn_ldexp:  in EmitAMDGPUBuiltinExpr()
445   case AMDGPU::BI__builtin_amdgcn_ldexpf: {  in EmitAMDGPUBuiltinExpr()
452   case AMDGPU::BI__builtin_amdgcn_ldexph: {  in EmitAMDGPUBuiltinExpr()
461   case AMDGPU::BI__builtin_amdgcn_frexp_mant:  in EmitAMDGPUBuiltinExpr()
462   case AMDGPU::BI__builtin_amdgcn_frexp_mantf:  in EmitAMDGPUBuiltinExpr()
463   case AMDGPU::BI__builtin_amdgcn_frexp_manth:  in EmitAMDGPUBuiltinExpr()
466   case AMDGPU::BI__builtin_amdgcn_frexp_exp:  in EmitAMDGPUBuiltinExpr()
467   case AMDGPU::BI__builtin_amdgcn_frexp_expf: {  in EmitAMDGPUBuiltinExpr()
473   case AMDGPU::BI__builtin_amdgcn_frexp_exph: {  in EmitAMDGPUBuiltinExpr()
479   case AMDGPU::BI__builtin_amdgcn_fract:  in EmitAMDGPUBuiltinExpr()
480   case AMDGPU::BI__builtin_amdgcn_fractf:  in EmitAMDGPUBuiltinExpr()
481   case AMDGPU::BI__builtin_amdgcn_fracth:  in EmitAMDGPUBuiltinExpr()
484   case AMDGPU::BI__builtin_amdgcn_lerp:  in EmitAMDGPUBuiltinExpr()
487   case AMDGPU::BI__builtin_amdgcn_ubfe:  in EmitAMDGPUBuiltinExpr()
490   case AMDGPU::BI__builtin_amdgcn_sbfe:  in EmitAMDGPUBuiltinExpr()
493   case AMDGPU::BI__builtin_amdgcn_ballot_w32:  in EmitAMDGPUBuiltinExpr()
494   case AMDGPU::BI__builtin_amdgcn_ballot_w64: {  in EmitAMDGPUBuiltinExpr()
500   case AMDGPU::BI__builtin_amdgcn_tanh_bf16:  in EmitAMDGPUBuiltinExpr()
503   case AMDGPU::BI__builtin_amdgcn_uicmp:  in EmitAMDGPUBuiltinExpr()
504   case AMDGPU::BI__builtin_amdgcn_uicmpl:  in EmitAMDGPUBuiltinExpr()
505   case AMDGPU::BI__builtin_amdgcn_sicmp:  in EmitAMDGPUBuiltinExpr()
506   case AMDGPU::BI__builtin_amdgcn_sicmpl: {  in EmitAMDGPUBuiltinExpr()
516   case AMDGPU::BI__builtin_amdgcn_fcmp:  in EmitAMDGPUBuiltinExpr()
517   case AMDGPU::BI__builtin_amdgcn_fcmpf: {  in EmitAMDGPUBuiltinExpr()
527   case AMDGPU::BI__builtin_amdgcn_class:  in EmitAMDGPUBuiltinExpr()
528   case AMDGPU::BI__builtin_amdgcn_classf:  in EmitAMDGPUBuiltinExpr()
529   case AMDGPU::BI__builtin_amdgcn_classh:  in EmitAMDGPUBuiltinExpr()
531   case AMDGPU::BI__builtin_amdgcn_fmed3f:  in EmitAMDGPUBuiltinExpr()
532   case AMDGPU::BI__builtin_amdgcn_fmed3h:  in EmitAMDGPUBuiltinExpr()
535   case AMDGPU::BI__builtin_amdgcn_ds_append:  in EmitAMDGPUBuiltinExpr()
536   case AMDGPU::BI__builtin_amdgcn_ds_consume: {  in EmitAMDGPUBuiltinExpr()
537     Intrinsic::ID Intrin = BuiltinID == AMDGPU::BI__builtin_amdgcn_ds_append ?  in EmitAMDGPUBuiltinExpr()
543   case AMDGPU::BI__builtin_amdgcn_global_load_tr_b64_i32:  in EmitAMDGPUBuiltinExpr()
544   case AMDGPU::BI__builtin_amdgcn_global_load_tr_b64_v2i32:  in EmitAMDGPUBuiltinExpr()
545   case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4i16:  in EmitAMDGPUBuiltinExpr()
546   case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4f16:  in EmitAMDGPUBuiltinExpr()
547   case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4bf16:  in EmitAMDGPUBuiltinExpr()
548   case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8i16:  in EmitAMDGPUBuiltinExpr()
549   case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8f16:  in EmitAMDGPUBuiltinExpr()
550   case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8bf16:  in EmitAMDGPUBuiltinExpr()
551   case AMDGPU::BI__builtin_amdgcn_global_load_tr4_b64_v2i32:  in EmitAMDGPUBuiltinExpr()
552   case AMDGPU::BI__builtin_amdgcn_global_load_tr8_b64_v2i32:  in EmitAMDGPUBuiltinExpr()
553   case AMDGPU::BI__builtin_amdgcn_global_load_tr6_b96_v3i32:  in EmitAMDGPUBuiltinExpr()
554   case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8i16:  in EmitAMDGPUBuiltinExpr()
555   case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8f16:  in EmitAMDGPUBuiltinExpr()
556   case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8bf16:  in EmitAMDGPUBuiltinExpr()
557   case AMDGPU::BI__builtin_amdgcn_ds_load_tr4_b64_v2i32:  in EmitAMDGPUBuiltinExpr()
558   case AMDGPU::BI__builtin_amdgcn_ds_load_tr8_b64_v2i32:  in EmitAMDGPUBuiltinExpr()
559   case AMDGPU::BI__builtin_amdgcn_ds_load_tr6_b96_v3i32:  in EmitAMDGPUBuiltinExpr()
560   case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8i16:  in EmitAMDGPUBuiltinExpr()
561   case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8f16:  in EmitAMDGPUBuiltinExpr()
562   case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8bf16:  in EmitAMDGPUBuiltinExpr()
563   case AMDGPU::BI__builtin_amdgcn_ds_read_tr4_b64_v2i32:  in EmitAMDGPUBuiltinExpr()
564   case AMDGPU::BI__builtin_amdgcn_ds_read_tr8_b64_v2i32:  in EmitAMDGPUBuiltinExpr()
565   case AMDGPU::BI__builtin_amdgcn_ds_read_tr6_b96_v3i32:  in EmitAMDGPUBuiltinExpr()
566   case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4f16:  in EmitAMDGPUBuiltinExpr()
567   case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4bf16:  in EmitAMDGPUBuiltinExpr()
568   case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4i16: {  in EmitAMDGPUBuiltinExpr()
571     case AMDGPU::BI__builtin_amdgcn_global_load_tr_b64_i32:  in EmitAMDGPUBuiltinExpr()
572     case AMDGPU::BI__builtin_amdgcn_global_load_tr_b64_v2i32:  in EmitAMDGPUBuiltinExpr()
573     case AMDGPU::BI__builtin_amdgcn_global_load_tr8_b64_v2i32:  in EmitAMDGPUBuiltinExpr()
576     case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4i16:  in EmitAMDGPUBuiltinExpr()
577     case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4f16:  in EmitAMDGPUBuiltinExpr()
578     case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v4bf16:  in EmitAMDGPUBuiltinExpr()
579     case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8i16:  in EmitAMDGPUBuiltinExpr()
580     case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8f16:  in EmitAMDGPUBuiltinExpr()
581     case AMDGPU::BI__builtin_amdgcn_global_load_tr_b128_v8bf16:  in EmitAMDGPUBuiltinExpr()
582     case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8i16:  in EmitAMDGPUBuiltinExpr()
583     case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8f16:  in EmitAMDGPUBuiltinExpr()
584     case AMDGPU::BI__builtin_amdgcn_global_load_tr16_b128_v8bf16:  in EmitAMDGPUBuiltinExpr()
587     case AMDGPU::BI__builtin_amdgcn_global_load_tr4_b64_v2i32:  in EmitAMDGPUBuiltinExpr()
590     case AMDGPU::BI__builtin_amdgcn_global_load_tr6_b96_v3i32:  in EmitAMDGPUBuiltinExpr()
593     case AMDGPU::BI__builtin_amdgcn_ds_load_tr4_b64_v2i32:  in EmitAMDGPUBuiltinExpr()
596     case AMDGPU::BI__builtin_amdgcn_ds_load_tr6_b96_v3i32:  in EmitAMDGPUBuiltinExpr()
599     case AMDGPU::BI__builtin_amdgcn_ds_load_tr8_b64_v2i32:  in EmitAMDGPUBuiltinExpr()
602     case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8i16:  in EmitAMDGPUBuiltinExpr()
603     case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8f16:  in EmitAMDGPUBuiltinExpr()
604     case AMDGPU::BI__builtin_amdgcn_ds_load_tr16_b128_v8bf16:  in EmitAMDGPUBuiltinExpr()
607     case AMDGPU::BI__builtin_amdgcn_ds_read_tr4_b64_v2i32:  in EmitAMDGPUBuiltinExpr()
610     case AMDGPU::BI__builtin_amdgcn_ds_read_tr8_b64_v2i32:  in EmitAMDGPUBuiltinExpr()
613     case AMDGPU::BI__builtin_amdgcn_ds_read_tr6_b96_v3i32:  in EmitAMDGPUBuiltinExpr()
616     case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4i16:  in EmitAMDGPUBuiltinExpr()
617     case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4f16:  in EmitAMDGPUBuiltinExpr()
618     case AMDGPU::BI__builtin_amdgcn_ds_read_tr16_b64_v4bf16:  in EmitAMDGPUBuiltinExpr()
627   case AMDGPU::BI__builtin_amdgcn_load_to_lds: {  in EmitAMDGPUBuiltinExpr()
632   case AMDGPU::BI__builtin_amdgcn_get_fpenv: {  in EmitAMDGPUBuiltinExpr()
637   case AMDGPU::BI__builtin_amdgcn_set_fpenv: {  in EmitAMDGPUBuiltinExpr()
643   case AMDGPU::BI__builtin_amdgcn_read_exec:  in EmitAMDGPUBuiltinExpr()
645   case AMDGPU::BI__builtin_amdgcn_read_exec_lo:  in EmitAMDGPUBuiltinExpr()
647   case AMDGPU::BI__builtin_amdgcn_read_exec_hi:  in EmitAMDGPUBuiltinExpr()
649   case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray:  in EmitAMDGPUBuiltinExpr()
650   case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray_h:  in EmitAMDGPUBuiltinExpr()
651   case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray_l:  in EmitAMDGPUBuiltinExpr()
652   case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray_lh: {  in EmitAMDGPUBuiltinExpr()
674   case AMDGPU::BI__builtin_amdgcn_image_bvh8_intersect_ray:  in EmitAMDGPUBuiltinExpr()
675   case AMDGPU::BI__builtin_amdgcn_image_bvh_dual_intersect_ray: {  in EmitAMDGPUBuiltinExpr()
678     case AMDGPU::BI__builtin_amdgcn_image_bvh8_intersect_ray:  in EmitAMDGPUBuiltinExpr()
681     case AMDGPU::BI__builtin_amdgcn_image_bvh_dual_intersect_ray:  in EmitAMDGPUBuiltinExpr()
712   case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_rtn:  in EmitAMDGPUBuiltinExpr()
713   case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_push4_pop1_rtn:  in EmitAMDGPUBuiltinExpr()
714   case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_push8_pop1_rtn:  in EmitAMDGPUBuiltinExpr()
715   case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_push8_pop2_rtn: {  in EmitAMDGPUBuiltinExpr()
718     case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_rtn:  in EmitAMDGPUBuiltinExpr()
721     case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_push4_pop1_rtn:  in EmitAMDGPUBuiltinExpr()
724     case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_push8_pop1_rtn:  in EmitAMDGPUBuiltinExpr()
727     case AMDGPU::BI__builtin_amdgcn_ds_bvh_stack_push8_pop2_rtn:  in EmitAMDGPUBuiltinExpr()
751   case AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4:  in EmitAMDGPUBuiltinExpr()
752   case AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4: {  in EmitAMDGPUBuiltinExpr()
755         BuiltinID == AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4  in EmitAMDGPUBuiltinExpr()
765   case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32:  in EmitAMDGPUBuiltinExpr()
766   case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w32:  in EmitAMDGPUBuiltinExpr()
767   case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64:  in EmitAMDGPUBuiltinExpr()
768   case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w64:  in EmitAMDGPUBuiltinExpr()
769   case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w32:  in EmitAMDGPUBuiltinExpr()
770   case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w32:  in EmitAMDGPUBuiltinExpr()
771   case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w64:  in EmitAMDGPUBuiltinExpr()
772   case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w64:  in EmitAMDGPUBuiltinExpr()
773   case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32:  in EmitAMDGPUBuiltinExpr()
774   case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w64:  in EmitAMDGPUBuiltinExpr()
775   case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w32:  in EmitAMDGPUBuiltinExpr()
776   case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w64:  in EmitAMDGPUBuiltinExpr()
777   case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w32:  in EmitAMDGPUBuiltinExpr()
778   case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w64:  in EmitAMDGPUBuiltinExpr()
779   case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32:  in EmitAMDGPUBuiltinExpr()
780   case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w64:  in EmitAMDGPUBuiltinExpr()
781   case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32_gfx12:  in EmitAMDGPUBuiltinExpr()
782   case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64_gfx12:  in EmitAMDGPUBuiltinExpr()
783   case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w32_gfx12:  in EmitAMDGPUBuiltinExpr()
784   case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w64_gfx12:  in EmitAMDGPUBuiltinExpr()
785   case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32_gfx12:  in EmitAMDGPUBuiltinExpr()
786   case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w64_gfx12:  in EmitAMDGPUBuiltinExpr()
787   case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12:  in EmitAMDGPUBuiltinExpr()
788   case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w64_gfx12:  in EmitAMDGPUBuiltinExpr()
789   case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w32_gfx12:  in EmitAMDGPUBuiltinExpr()
790   case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w64_gfx12:  in EmitAMDGPUBuiltinExpr()
791   case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12:  in EmitAMDGPUBuiltinExpr()
792   case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w64_gfx12:  in EmitAMDGPUBuiltinExpr()
793   case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w32_gfx12:  in EmitAMDGPUBuiltinExpr()
794   case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w64_gfx12:  in EmitAMDGPUBuiltinExpr()
795   case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_bf8_w32_gfx12:  in EmitAMDGPUBuiltinExpr()
796   case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_bf8_w64_gfx12:  in EmitAMDGPUBuiltinExpr()
797   case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_fp8_w32_gfx12:  in EmitAMDGPUBuiltinExpr()
798   case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_fp8_w64_gfx12:  in EmitAMDGPUBuiltinExpr()
799   case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_bf8_w32_gfx12:  in EmitAMDGPUBuiltinExpr()
800   case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_bf8_w64_gfx12:  in EmitAMDGPUBuiltinExpr()
801   case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x32_iu4_w32_gfx12:  in EmitAMDGPUBuiltinExpr()
802   case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x32_iu4_w64_gfx12:  in EmitAMDGPUBuiltinExpr()
803   case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_f16_w32:  in EmitAMDGPUBuiltinExpr()
804   case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_f16_w64:  in EmitAMDGPUBuiltinExpr()
805   case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf16_w32:  in EmitAMDGPUBuiltinExpr()
806   case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf16_w64:  in EmitAMDGPUBuiltinExpr()
807   case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x32_f16_w32:  in EmitAMDGPUBuiltinExpr()
808   case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x32_f16_w64:  in EmitAMDGPUBuiltinExpr()
809   case AMDGPU::BI__builtin_amdgcn_swmmac_bf16_16x16x32_bf16_w32:  in EmitAMDGPUBuiltinExpr()
810   case AMDGPU::BI__builtin_amdgcn_swmmac_bf16_16x16x32_bf16_w64:  in EmitAMDGPUBuiltinExpr()
811   case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu8_w32:  in EmitAMDGPUBuiltinExpr()
812   case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu8_w64:  in EmitAMDGPUBuiltinExpr()
813   case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu4_w32:  in EmitAMDGPUBuiltinExpr()
814   case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu4_w64:  in EmitAMDGPUBuiltinExpr()
815   case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x64_iu4_w32:  in EmitAMDGPUBuiltinExpr()
816   case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x64_iu4_w64:  in EmitAMDGPUBuiltinExpr()
817   case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_fp8_w32:  in EmitAMDGPUBuiltinExpr()
818   case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_fp8_w64:  in EmitAMDGPUBuiltinExpr()
819   case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_bf8_w32:  in EmitAMDGPUBuiltinExpr()
820   case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_bf8_w64:  in EmitAMDGPUBuiltinExpr()
821   case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w32:  in EmitAMDGPUBuiltinExpr()
822   case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w64:  in EmitAMDGPUBuiltinExpr()
823   case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w32:  in EmitAMDGPUBuiltinExpr()
824   case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w64: {  in EmitAMDGPUBuiltinExpr()
841     case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w32:  in EmitAMDGPUBuiltinExpr()
842     case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w64:  in EmitAMDGPUBuiltinExpr()
843     case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12:  in EmitAMDGPUBuiltinExpr()
844     case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_f16_w64_gfx12:  in EmitAMDGPUBuiltinExpr()
848     case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32:  in EmitAMDGPUBuiltinExpr()
849     case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w64:  in EmitAMDGPUBuiltinExpr()
850     case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32_gfx12:  in EmitAMDGPUBuiltinExpr()
851     case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf16_w64_gfx12:  in EmitAMDGPUBuiltinExpr()
855     case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w32_gfx12:  in EmitAMDGPUBuiltinExpr()
856     case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w64_gfx12:  in EmitAMDGPUBuiltinExpr()
859     case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w32:  in EmitAMDGPUBuiltinExpr()
860     case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_w64:  in EmitAMDGPUBuiltinExpr()
864     case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32_gfx12:  in EmitAMDGPUBuiltinExpr()
865     case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64_gfx12:  in EmitAMDGPUBuiltinExpr()
868     case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32:  in EmitAMDGPUBuiltinExpr()
869     case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64:  in EmitAMDGPUBuiltinExpr()
873     case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w32:  in EmitAMDGPUBuiltinExpr()
874     case AMDGPU::BI__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w64:  in EmitAMDGPUBuiltinExpr()
878     case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w32:  in EmitAMDGPUBuiltinExpr()
879     case AMDGPU::BI__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w64:  in EmitAMDGPUBuiltinExpr()
883     case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32:  in EmitAMDGPUBuiltinExpr()
884     case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w64:  in EmitAMDGPUBuiltinExpr()
885     case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12:  in EmitAMDGPUBuiltinExpr()
886     case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu8_w64_gfx12:  in EmitAMDGPUBuiltinExpr()
890     case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w32:  in EmitAMDGPUBuiltinExpr()
891     case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w64:  in EmitAMDGPUBuiltinExpr()
892     case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w32_gfx12:  in EmitAMDGPUBuiltinExpr()
893     case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x16_iu4_w64_gfx12:  in EmitAMDGPUBuiltinExpr()
897     case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w32_gfx12:  in EmitAMDGPUBuiltinExpr()
898     case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w64_gfx12:  in EmitAMDGPUBuiltinExpr()
902     case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_bf8_w32_gfx12:  in EmitAMDGPUBuiltinExpr()
903     case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_fp8_bf8_w64_gfx12:  in EmitAMDGPUBuiltinExpr()
907     case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_fp8_w32_gfx12:  in EmitAMDGPUBuiltinExpr()
908     case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_fp8_w64_gfx12:  in EmitAMDGPUBuiltinExpr()
912     case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_bf8_w32_gfx12:  in EmitAMDGPUBuiltinExpr()
913     case AMDGPU::BI__builtin_amdgcn_wmma_f32_16x16x16_bf8_bf8_w64_gfx12:  in EmitAMDGPUBuiltinExpr()
917     case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x32_iu4_w32_gfx12:  in EmitAMDGPUBuiltinExpr()
918     case AMDGPU::BI__builtin_amdgcn_wmma_i32_16x16x32_iu4_w64_gfx12:  in EmitAMDGPUBuiltinExpr()
922     case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_f16_w32:  in EmitAMDGPUBuiltinExpr()
923     case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_f16_w64:  in EmitAMDGPUBuiltinExpr()
927     case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf16_w32:  in EmitAMDGPUBuiltinExpr()
928     case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf16_w64:  in EmitAMDGPUBuiltinExpr()
932     case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x32_f16_w32:  in EmitAMDGPUBuiltinExpr()
933     case AMDGPU::BI__builtin_amdgcn_swmmac_f16_16x16x32_f16_w64:  in EmitAMDGPUBuiltinExpr()
937     case AMDGPU::BI__builtin_amdgcn_swmmac_bf16_16x16x32_bf16_w32:  in EmitAMDGPUBuiltinExpr()
938     case AMDGPU::BI__builtin_amdgcn_swmmac_bf16_16x16x32_bf16_w64:  in EmitAMDGPUBuiltinExpr()
942     case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu8_w32:  in EmitAMDGPUBuiltinExpr()
943     case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu8_w64:  in EmitAMDGPUBuiltinExpr()
947     case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu4_w32:  in EmitAMDGPUBuiltinExpr()
948     case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x32_iu4_w64:  in EmitAMDGPUBuiltinExpr()
952     case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x64_iu4_w32:  in EmitAMDGPUBuiltinExpr()
953     case AMDGPU::BI__builtin_amdgcn_swmmac_i32_16x16x64_iu4_w64:  in EmitAMDGPUBuiltinExpr()
957     case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_fp8_w32:  in EmitAMDGPUBuiltinExpr()
958     case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_fp8_w64:  in EmitAMDGPUBuiltinExpr()
962     case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_bf8_w32:  in EmitAMDGPUBuiltinExpr()
963     case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_fp8_bf8_w64:  in EmitAMDGPUBuiltinExpr()
967     case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w32:  in EmitAMDGPUBuiltinExpr()
968     case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w64:  in EmitAMDGPUBuiltinExpr()
972     case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w32:  in EmitAMDGPUBuiltinExpr()
973     case AMDGPU::BI__builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w64:  in EmitAMDGPUBuiltinExpr()
993   case AMDGPU::BI__builtin_amdgcn_workgroup_size_x:  in EmitAMDGPUBuiltinExpr()
995   case AMDGPU::BI__builtin_amdgcn_workgroup_size_y:  in EmitAMDGPUBuiltinExpr()
997   case AMDGPU::BI__builtin_amdgcn_workgroup_size_z:  in EmitAMDGPUBuiltinExpr()
1001   case AMDGPU::BI__builtin_amdgcn_grid_size_x:  in EmitAMDGPUBuiltinExpr()
1003   case AMDGPU::BI__builtin_amdgcn_grid_size_y:  in EmitAMDGPUBuiltinExpr()
1005   case AMDGPU::BI__builtin_amdgcn_grid_size_z:  in EmitAMDGPUBuiltinExpr()
1009   case AMDGPU::BI__builtin_r600_recipsqrt_ieee:  in EmitAMDGPUBuiltinExpr()
1010   case AMDGPU::BI__builtin_r600_recipsqrt_ieeef:  in EmitAMDGPUBuiltinExpr()
1013   case AMDGPU::BI__builtin_amdgcn_alignbit: {  in EmitAMDGPUBuiltinExpr()
1020   case AMDGPU::BI__builtin_amdgcn_fence: {  in EmitAMDGPUBuiltinExpr()
1028   case AMDGPU::BI__builtin_amdgcn_atomic_inc32:  in EmitAMDGPUBuiltinExpr()
1029   case AMDGPU::BI__builtin_amdgcn_atomic_inc64:  in EmitAMDGPUBuiltinExpr()
1030   case AMDGPU::BI__builtin_amdgcn_atomic_dec32:  in EmitAMDGPUBuiltinExpr()
1031   case AMDGPU::BI__builtin_amdgcn_atomic_dec64:  in EmitAMDGPUBuiltinExpr()
1032   case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_f64:  in EmitAMDGPUBuiltinExpr()
1033   case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_f32:  in EmitAMDGPUBuiltinExpr()
1034   case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2f16:  in EmitAMDGPUBuiltinExpr()
1035   case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2bf16:  in EmitAMDGPUBuiltinExpr()
1036   case AMDGPU::BI__builtin_amdgcn_ds_faddf:  in EmitAMDGPUBuiltinExpr()
1037   case AMDGPU::BI__builtin_amdgcn_ds_fminf:  in EmitAMDGPUBuiltinExpr()
1038   case AMDGPU::BI__builtin_amdgcn_ds_fmaxf:  in EmitAMDGPUBuiltinExpr()
1039   case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_f32:  in EmitAMDGPUBuiltinExpr()
1040   case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_f64:  in EmitAMDGPUBuiltinExpr()
1041   case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2f16:  in EmitAMDGPUBuiltinExpr()
1042   case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2f16:  in EmitAMDGPUBuiltinExpr()
1043   case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f32:  in EmitAMDGPUBuiltinExpr()
1044   case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f64:  in EmitAMDGPUBuiltinExpr()
1045   case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2bf16:  in EmitAMDGPUBuiltinExpr()
1046   case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2bf16:  in EmitAMDGPUBuiltinExpr()
1047   case AMDGPU::BI__builtin_amdgcn_global_atomic_fmin_f64:  in EmitAMDGPUBuiltinExpr()
1048   case AMDGPU::BI__builtin_amdgcn_global_atomic_fmax_f64:  in EmitAMDGPUBuiltinExpr()
1049   case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmin_f64:  in EmitAMDGPUBuiltinExpr()
1050   case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmax_f64: {  in EmitAMDGPUBuiltinExpr()
1053     case AMDGPU::BI__builtin_amdgcn_atomic_inc32:  in EmitAMDGPUBuiltinExpr()
1054     case AMDGPU::BI__builtin_amdgcn_atomic_inc64:  in EmitAMDGPUBuiltinExpr()
1057     case AMDGPU::BI__builtin_amdgcn_atomic_dec32:  in EmitAMDGPUBuiltinExpr()
1058     case AMDGPU::BI__builtin_amdgcn_atomic_dec64:  in EmitAMDGPUBuiltinExpr()
1061     case AMDGPU::BI__builtin_amdgcn_ds_faddf:  in EmitAMDGPUBuiltinExpr()
1062     case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_f64:  in EmitAMDGPUBuiltinExpr()
1063     case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_f32:  in EmitAMDGPUBuiltinExpr()
1064     case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2f16:  in EmitAMDGPUBuiltinExpr()
1065     case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2bf16:  in EmitAMDGPUBuiltinExpr()
1066     case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_f32:  in EmitAMDGPUBuiltinExpr()
1067     case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_f64:  in EmitAMDGPUBuiltinExpr()
1068     case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2f16:  in EmitAMDGPUBuiltinExpr()
1069     case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2f16:  in EmitAMDGPUBuiltinExpr()
1070     case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f32:  in EmitAMDGPUBuiltinExpr()
1071     case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_f64:  in EmitAMDGPUBuiltinExpr()
1072     case AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2bf16:  in EmitAMDGPUBuiltinExpr()
1073     case AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2bf16:  in EmitAMDGPUBuiltinExpr()
1076     case AMDGPU::BI__builtin_amdgcn_ds_fminf:  in EmitAMDGPUBuiltinExpr()
1077     case AMDGPU::BI__builtin_amdgcn_global_atomic_fmin_f64:  in EmitAMDGPUBuiltinExpr()
1078     case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmin_f64:  in EmitAMDGPUBuiltinExpr()
1081     case AMDGPU::BI__builtin_amdgcn_global_atomic_fmax_f64:  in EmitAMDGPUBuiltinExpr()
1082     case AMDGPU::BI__builtin_amdgcn_flat_atomic_fmax_f64:  in EmitAMDGPUBuiltinExpr()
1083     case AMDGPU::BI__builtin_amdgcn_ds_fmaxf:  in EmitAMDGPUBuiltinExpr()
1095     if (BuiltinID == AMDGPU::BI__builtin_amdgcn_ds_faddf ||  in EmitAMDGPUBuiltinExpr()
1096         BuiltinID == AMDGPU::BI__builtin_amdgcn_ds_fminf ||  in EmitAMDGPUBuiltinExpr()
1097         BuiltinID == AMDGPU::BI__builtin_amdgcn_ds_fmaxf) {  in EmitAMDGPUBuiltinExpr()
1122       if (BuiltinID == AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_v2bf16 ||  in EmitAMDGPUBuiltinExpr()
1123           BuiltinID == AMDGPU::BI__builtin_amdgcn_global_atomic_fadd_v2bf16 ||  in EmitAMDGPUBuiltinExpr()
1124           BuiltinID == AMDGPU::BI__builtin_amdgcn_flat_atomic_fadd_v2bf16) {  in EmitAMDGPUBuiltinExpr()
1151   case AMDGPU::BI__builtin_amdgcn_s_sendmsg_rtn:  in EmitAMDGPUBuiltinExpr()
1152   case AMDGPU::BI__builtin_amdgcn_s_sendmsg_rtnl: {  in EmitAMDGPUBuiltinExpr()
1160   case AMDGPU::BI__builtin_amdgcn_permlane16_swap:  in EmitAMDGPUBuiltinExpr()
1161   case AMDGPU::BI__builtin_amdgcn_permlane32_swap: {  in EmitAMDGPUBuiltinExpr()
1169         CGM.getIntrinsic(BuiltinID == AMDGPU::BI__builtin_amdgcn_permlane16_swap  in EmitAMDGPUBuiltinExpr()
1186   case AMDGPU::BI__builtin_amdgcn_bitop3_b32:  in EmitAMDGPUBuiltinExpr()
1187   case AMDGPU::BI__builtin_amdgcn_bitop3_b16:  in EmitAMDGPUBuiltinExpr()
1190   case AMDGPU::BI__builtin_amdgcn_make_buffer_rsrc: {  in EmitAMDGPUBuiltinExpr()
1203   case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b8:  in EmitAMDGPUBuiltinExpr()
1204   case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b16:  in EmitAMDGPUBuiltinExpr()
1205   case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b32:  in EmitAMDGPUBuiltinExpr()
1206   case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b64:  in EmitAMDGPUBuiltinExpr()
1207   case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b96:  in EmitAMDGPUBuiltinExpr()
1208   case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_b128:  in EmitAMDGPUBuiltinExpr()
1211   case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b8:  in EmitAMDGPUBuiltinExpr()
1212   case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b16:  in EmitAMDGPUBuiltinExpr()
1213   case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b32:  in EmitAMDGPUBuiltinExpr()
1214   case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b64:  in EmitAMDGPUBuiltinExpr()
1215   case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b96:  in EmitAMDGPUBuiltinExpr()
1216   case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b128: {  in EmitAMDGPUBuiltinExpr()
1219     case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b8:  in EmitAMDGPUBuiltinExpr()
1222     case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b16:  in EmitAMDGPUBuiltinExpr()
1225     case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b32:  in EmitAMDGPUBuiltinExpr()
1228     case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b64:  in EmitAMDGPUBuiltinExpr()
1231     case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b96:  in EmitAMDGPUBuiltinExpr()
1234     case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_b128:  in EmitAMDGPUBuiltinExpr()
1244   case AMDGPU::BI__builtin_amdgcn_s_prefetch_data:  in EmitAMDGPUBuiltinExpr()