298 Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,  in EmitAMDGPUBuiltinExpr()  argument
302   switch (BuiltinID) {  in EmitAMDGPUBuiltinExpr()
352     getContext().GetBuiltinType(BuiltinID, Error, &ICEArguments);  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()
391         BuiltinID == AMDGPU::BI__builtin_amdgcn_permlane16  in EmitAMDGPUBuiltinExpr()
537     Intrinsic::ID Intrin = BuiltinID == AMDGPU::BI__builtin_amdgcn_ds_append ?  in EmitAMDGPUBuiltinExpr()
570     switch (BuiltinID) {  in EmitAMDGPUBuiltinExpr()
677     switch (BuiltinID) {  in EmitAMDGPUBuiltinExpr()
717     switch (BuiltinID) {  in EmitAMDGPUBuiltinExpr()
755         BuiltinID == AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4  in EmitAMDGPUBuiltinExpr()
840     switch (BuiltinID) {  in EmitAMDGPUBuiltinExpr()
1052     switch (BuiltinID) {  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()
1169         CGM.getIntrinsic(BuiltinID == AMDGPU::BI__builtin_amdgcn_permlane16_swap  in EmitAMDGPUBuiltinExpr()
1218     switch (BuiltinID) {  in EmitAMDGPUBuiltinExpr()