diff options
| -rw-r--r-- | clang/include/clang/Basic/BuiltinsAMDGPU.def | 2 | ||||
| -rw-r--r-- | clang/lib/CodeGen/CGBuiltin.cpp | 15 | ||||
| -rw-r--r-- | clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl | 2 | ||||
| -rw-r--r-- | clang/test/CodeGenOpenCL/builtins-amdgcn.cl | 4 |
4 files changed, 16 insertions, 7 deletions
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def index 99fb9373e70..f0f63fa73a7 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.def +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def @@ -92,7 +92,7 @@ TARGET_BUILTIN(__builtin_amdgcn_sinh, "hh", "nc", "16-bit-insts") TARGET_BUILTIN(__builtin_amdgcn_cosh, "hh", "nc", "16-bit-insts") TARGET_BUILTIN(__builtin_amdgcn_ldexph, "hhi", "nc", "16-bit-insts") TARGET_BUILTIN(__builtin_amdgcn_frexp_manth, "hh", "nc", "16-bit-insts") -TARGET_BUILTIN(__builtin_amdgcn_frexp_exph, "ih", "nc", "16-bit-insts") +TARGET_BUILTIN(__builtin_amdgcn_frexp_exph, "sh", "nc", "16-bit-insts") TARGET_BUILTIN(__builtin_amdgcn_fracth, "hh", "nc", "16-bit-insts") TARGET_BUILTIN(__builtin_amdgcn_classh, "bhi", "nc", "16-bit-insts") TARGET_BUILTIN(__builtin_amdgcn_s_memrealtime, "LUi", "n", "s-memrealtime") 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: diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl index 0892c046898..0c88fdfb724 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl @@ -55,7 +55,7 @@ void test_frexp_mant_f16(global half* out, half a) } // CHECK-LABEL: @test_frexp_exp_f16 -// CHECK: call i32 @llvm.amdgcn.frexp.exp.f16 +// CHECK: call i16 @llvm.amdgcn.frexp.exp.i16.f16 void test_frexp_exp_f16(global short* out, half a) { *out = __builtin_amdgcn_frexp_exph(a); diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl index eea642296b1..ce566b6e876 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl @@ -166,14 +166,14 @@ void test_frexp_mant_f64(global double* out, double a) } // CHECK-LABEL: @test_frexp_exp_f32 -// CHECK: call i32 @llvm.amdgcn.frexp.exp.f32 +// CHECK: call i32 @llvm.amdgcn.frexp.exp.i32.f32 void test_frexp_exp_f32(global int* out, float a) { *out = __builtin_amdgcn_frexp_expf(a); } // CHECK-LABEL: @test_frexp_exp_f64 -// CHECK: call i32 @llvm.amdgcn.frexp.exp.f64 +// CHECK: call i32 @llvm.amdgcn.frexp.exp.i32.f64 void test_frexp_exp_f64(global int* out, double a) { *out = __builtin_amdgcn_frexp_exp(a); |
