diff options
| author | Konstantin Zhuravlyov <kzhuravl_dev@outlook.com> | 2016-11-13 02:37:05 +0000 |
|---|---|---|
| committer | Konstantin Zhuravlyov <kzhuravl_dev@outlook.com> | 2016-11-13 02:37:05 +0000 |
| commit | 81a78bb8640ba352aa6e28885facdbac3e821d6c (patch) | |
| tree | 32622e9afc06acb9df19b86bee57c0ffce0084f8 /clang/lib/CodeGen/CGBuiltin.cpp | |
| parent | e0038717b9230db2b41e219735d9e352ea96ded1 (diff) | |
| download | bcm5719-llvm-81a78bb8640ba352aa6e28885facdbac3e821d6c.tar.gz bcm5719-llvm-81a78bb8640ba352aa6e28885facdbac3e821d6c.zip | |
[AMDGPU] Add f16 builtin functions (VI+)
Differential Revision: https://reviews.llvm.org/D26476
llvm-svn: 286741
Diffstat (limited to 'clang/lib/CodeGen/CGBuiltin.cpp')
| -rw-r--r-- | clang/lib/CodeGen/CGBuiltin.cpp | 16 |
1 files changed, 12 insertions, 4 deletions
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp index 039cba5e1c6..274120aae28 100644 --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/clang/lib/CodeGen/CGBuiltin.cpp @@ -8190,38 +8190,45 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID, return emitBinaryBuiltin(*this, E, Intrinsic::amdgcn_ds_swizzle); case AMDGPU::BI__builtin_amdgcn_div_fixup: case AMDGPU::BI__builtin_amdgcn_div_fixupf: + case AMDGPU::BI__builtin_amdgcn_div_fixuph: return emitTernaryBuiltin(*this, E, Intrinsic::amdgcn_div_fixup); case AMDGPU::BI__builtin_amdgcn_trig_preop: case AMDGPU::BI__builtin_amdgcn_trig_preopf: return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_trig_preop); case AMDGPU::BI__builtin_amdgcn_rcp: case AMDGPU::BI__builtin_amdgcn_rcpf: + case AMDGPU::BI__builtin_amdgcn_rcph: return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_rcp); case AMDGPU::BI__builtin_amdgcn_rsq: case AMDGPU::BI__builtin_amdgcn_rsqf: + case AMDGPU::BI__builtin_amdgcn_rsqh: return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_rsq); case AMDGPU::BI__builtin_amdgcn_rsq_clamp: case AMDGPU::BI__builtin_amdgcn_rsq_clampf: return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_rsq_clamp); case AMDGPU::BI__builtin_amdgcn_sinf: + case AMDGPU::BI__builtin_amdgcn_sinh: return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_sin); case AMDGPU::BI__builtin_amdgcn_cosf: + case AMDGPU::BI__builtin_amdgcn_cosh: return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_cos); case AMDGPU::BI__builtin_amdgcn_log_clampf: return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_log_clamp); case AMDGPU::BI__builtin_amdgcn_ldexp: case AMDGPU::BI__builtin_amdgcn_ldexpf: + case AMDGPU::BI__builtin_amdgcn_ldexph: return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_ldexp); case AMDGPU::BI__builtin_amdgcn_frexp_mant: - case AMDGPU::BI__builtin_amdgcn_frexp_mantf: { + case AMDGPU::BI__builtin_amdgcn_frexp_mantf: + case AMDGPU::BI__builtin_amdgcn_frexp_manth: return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_frexp_mant); - } case AMDGPU::BI__builtin_amdgcn_frexp_exp: - case AMDGPU::BI__builtin_amdgcn_frexp_expf: { + case AMDGPU::BI__builtin_amdgcn_frexp_expf: + case AMDGPU::BI__builtin_amdgcn_frexp_exph: return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_frexp_exp); - } case AMDGPU::BI__builtin_amdgcn_fract: case AMDGPU::BI__builtin_amdgcn_fractf: + case AMDGPU::BI__builtin_amdgcn_fracth: return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_fract); case AMDGPU::BI__builtin_amdgcn_lerp: return emitTernaryBuiltin(*this, E, Intrinsic::amdgcn_lerp); @@ -8235,6 +8242,7 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID, return emitTernaryBuiltin(*this, E, Intrinsic::amdgcn_fcmp); case AMDGPU::BI__builtin_amdgcn_class: case AMDGPU::BI__builtin_amdgcn_classf: + case AMDGPU::BI__builtin_amdgcn_classh: return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_class); case AMDGPU::BI__builtin_amdgcn_read_exec: { |
