diff options
| author | Konstantin Zhuravlyov <kzhuravl_dev@outlook.com> | 2016-11-18 22:31:51 +0000 |
|---|---|---|
| committer | Konstantin Zhuravlyov <kzhuravl_dev@outlook.com> | 2016-11-18 22:31:51 +0000 |
| commit | 62ae8f671cc922ec15e001e40801164073c8eb84 (patch) | |
| tree | 906aa3308a7d9b18b6bf884dd9ea8c837af5fd7b /clang/lib | |
| parent | aefee42e0fd1c0b5b12a9349800b5605a1a65d0a (diff) | |
| download | bcm5719-llvm-62ae8f671cc922ec15e001e40801164073c8eb84.tar.gz bcm5719-llvm-62ae8f671cc922ec15e001e40801164073c8eb84.zip | |
[AMDGPU] Change frexp.exp builtin to return i16 for f16 input
Differential Revision: https://reviews.llvm.org/D26863
llvm-svn: 287390
Diffstat (limited to 'clang/lib')
| -rw-r--r-- | clang/lib/CodeGen/CGBuiltin.cpp | 15 |
1 files changed, 12 insertions, 3 deletions
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp index fd8f17e0efe..2520f2e75e6 100644 --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/clang/lib/CodeGen/CGBuiltin.cpp @@ -8250,9 +8250,18 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID, 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_exph: - return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_frexp_exp); + case AMDGPU::BI__builtin_amdgcn_frexp_expf: { + Value *Src0 = EmitScalarExpr(E->getArg(0)); + Value *F = CGM.getIntrinsic(Intrinsic::amdgcn_frexp_exp, + { Builder.getInt32Ty(), Src0->getType() }); + return Builder.CreateCall(F, Src0); + } + case AMDGPU::BI__builtin_amdgcn_frexp_exph: { + Value *Src0 = EmitScalarExpr(E->getArg(0)); + Value *F = CGM.getIntrinsic(Intrinsic::amdgcn_frexp_exp, + { Builder.getInt16Ty(), Src0->getType() }); + return Builder.CreateCall(F, Src0); + } case AMDGPU::BI__builtin_amdgcn_fract: case AMDGPU::BI__builtin_amdgcn_fractf: case AMDGPU::BI__builtin_amdgcn_fracth: |
