diff options
Diffstat (limited to 'clang/lib')
| -rw-r--r-- | clang/lib/CodeGen/CGBuiltin.cpp | 26 |
1 files changed, 12 insertions, 14 deletions
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp index 7a869f203e0..dbff809ab5e 100644 --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/clang/lib/CodeGen/CGBuiltin.cpp @@ -7671,19 +7671,6 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID, CI->setConvergent(); return CI; } - // Legacy amdgpu prefix - case AMDGPU::BI__builtin_amdgpu_rsq: - case AMDGPU::BI__builtin_amdgpu_rsqf: { - if (getTarget().getTriple().getArch() == Triple::amdgcn) - return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_rsq); - return emitUnaryBuiltin(*this, E, Intrinsic::r600_rsq); - } - case AMDGPU::BI__builtin_amdgpu_ldexp: - case AMDGPU::BI__builtin_amdgpu_ldexpf: { - if (getTarget().getTriple().getArch() == Triple::amdgcn) - return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_ldexp); - return emitFPIntBuiltin(*this, E, Intrinsic::AMDGPU_ldexp); - } // amdgcn workitem case AMDGPU::BI__builtin_amdgcn_workitem_id_x: @@ -7693,13 +7680,24 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID, case AMDGPU::BI__builtin_amdgcn_workitem_id_z: return emitRangedBuiltin(*this, Intrinsic::amdgcn_workitem_id_z, 0, 1024); - // r600 workitem + // r600 intrinsics + case AMDGPU::BI__builtin_r600_recipsqrt_ieee: + case AMDGPU::BI__builtin_r600_recipsqrt_ieeef: + return emitUnaryBuiltin(*this, E, Intrinsic::r600_recipsqrt_ieee); case AMDGPU::BI__builtin_r600_read_tidig_x: return emitRangedBuiltin(*this, Intrinsic::r600_read_tidig_x, 0, 1024); case AMDGPU::BI__builtin_r600_read_tidig_y: return emitRangedBuiltin(*this, Intrinsic::r600_read_tidig_y, 0, 1024); case AMDGPU::BI__builtin_r600_read_tidig_z: return emitRangedBuiltin(*this, Intrinsic::r600_read_tidig_z, 0, 1024); + + // Legacy amdgpu prefix + case AMDGPU::BI__builtin_amdgpu_ldexp: + case AMDGPU::BI__builtin_amdgpu_ldexpf: { + if (getTarget().getTriple().getArch() == Triple::amdgcn) + return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_ldexp); + return emitFPIntBuiltin(*this, E, Intrinsic::AMDGPU_ldexp); + } default: return nullptr; } |
