Author: kzhuravl Date: Fri Nov 18 16:31:51 2016 New Revision: 287390 URL: http://llvm.org/viewvc/llvm-project?rev=287390&view=rev Log: [AMDGPU] Change frexp.exp builtin to return i16 for f16 input
Differential Revision: https://reviews.llvm.org/D26863 Modified: cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def cfe/trunk/lib/CodeGen/CGBuiltin.cpp cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl Modified: cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def?rev=287390&r1=287389&r2=287390&view=diff ============================================================================== --- cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def (original) +++ cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def Fri Nov 18 16:31:51 2016 @@ -92,7 +92,7 @@ TARGET_BUILTIN(__builtin_amdgcn_sinh, "h 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") Modified: cfe/trunk/lib/CodeGen/CGBuiltin.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGBuiltin.cpp?rev=287390&r1=287389&r2=287390&view=diff ============================================================================== --- cfe/trunk/lib/CodeGen/CGBuiltin.cpp (original) +++ cfe/trunk/lib/CodeGen/CGBuiltin.cpp Fri Nov 18 16:31:51 2016 @@ -8250,9 +8250,18 @@ Value *CodeGenFunction::EmitAMDGPUBuilti 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: Modified: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl?rev=287390&r1=287389&r2=287390&view=diff ============================================================================== --- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl (original) +++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl Fri Nov 18 16:31:51 2016 @@ -55,7 +55,7 @@ void test_frexp_mant_f16(global half* ou } // 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); Modified: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl?rev=287390&r1=287389&r2=287390&view=diff ============================================================================== --- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl (original) +++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl Fri Nov 18 16:31:51 2016 @@ -166,14 +166,14 @@ void test_frexp_mant_f64(global double* } // 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); _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
