forked from OSchip/llvm-project
AMDGPU: Add frexp_mant + frexp_exp builtins
llvm-svn: 264960
This commit is contained in:
parent
8d41436004
commit
3fb963389e
|
@ -34,6 +34,10 @@ BUILTIN(__builtin_amdgcn_cosf, "ff", "nc")
|
||||||
BUILTIN(__builtin_amdgcn_log_clampf, "ff", "nc")
|
BUILTIN(__builtin_amdgcn_log_clampf, "ff", "nc")
|
||||||
BUILTIN(__builtin_amdgcn_ldexp, "ddi", "nc")
|
BUILTIN(__builtin_amdgcn_ldexp, "ddi", "nc")
|
||||||
BUILTIN(__builtin_amdgcn_ldexpf, "ffi", "nc")
|
BUILTIN(__builtin_amdgcn_ldexpf, "ffi", "nc")
|
||||||
|
BUILTIN(__builtin_amdgcn_frexp_mant, "dd", "nc")
|
||||||
|
BUILTIN(__builtin_amdgcn_frexp_mantf, "ff", "nc")
|
||||||
|
BUILTIN(__builtin_amdgcn_frexp_exp, "id", "nc")
|
||||||
|
BUILTIN(__builtin_amdgcn_frexp_expf, "if", "nc")
|
||||||
BUILTIN(__builtin_amdgcn_class, "bdi", "nc")
|
BUILTIN(__builtin_amdgcn_class, "bdi", "nc")
|
||||||
BUILTIN(__builtin_amdgcn_classf, "bfi", "nc")
|
BUILTIN(__builtin_amdgcn_classf, "bfi", "nc")
|
||||||
BUILTIN(__builtin_amdgcn_cubeid, "ffff", "nc")
|
BUILTIN(__builtin_amdgcn_cubeid, "ffff", "nc")
|
||||||
|
|
|
@ -7105,6 +7105,14 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
|
||||||
case AMDGPU::BI__builtin_amdgcn_ldexp:
|
case AMDGPU::BI__builtin_amdgcn_ldexp:
|
||||||
case AMDGPU::BI__builtin_amdgcn_ldexpf:
|
case AMDGPU::BI__builtin_amdgcn_ldexpf:
|
||||||
return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_ldexp);
|
return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_ldexp);
|
||||||
|
case AMDGPU::BI__builtin_amdgcn_frexp_mant:
|
||||||
|
case AMDGPU::BI__builtin_amdgcn_frexp_mantf: {
|
||||||
|
return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_frexp_mant);
|
||||||
|
}
|
||||||
|
case AMDGPU::BI__builtin_amdgcn_frexp_exp:
|
||||||
|
case AMDGPU::BI__builtin_amdgcn_frexp_expf: {
|
||||||
|
return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_frexp_exp);
|
||||||
|
}
|
||||||
case AMDGPU::BI__builtin_amdgcn_class:
|
case AMDGPU::BI__builtin_amdgcn_class:
|
||||||
case AMDGPU::BI__builtin_amdgcn_classf:
|
case AMDGPU::BI__builtin_amdgcn_classf:
|
||||||
return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_class);
|
return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_class);
|
||||||
|
|
|
@ -150,6 +150,34 @@ void test_ldexp_f64(global double* out, double a, int b)
|
||||||
*out = __builtin_amdgcn_ldexp(a, b);
|
*out = __builtin_amdgcn_ldexp(a, b);
|
||||||
}
|
}
|
||||||
|
|
||||||
|
// CHECK-LABEL: @test_frexp_mant_f32
|
||||||
|
// CHECK: call float @llvm.amdgcn.frexp.mant.f32
|
||||||
|
void test_frexp_mant_f32(global float* out, float a)
|
||||||
|
{
|
||||||
|
*out = __builtin_amdgcn_frexp_mantf(a);
|
||||||
|
}
|
||||||
|
|
||||||
|
// CHECK-LABEL: @test_frexp_mant_f64
|
||||||
|
// CHECK: call double @llvm.amdgcn.frexp.mant.f64
|
||||||
|
void test_frexp_mant_f64(global double* out, double a)
|
||||||
|
{
|
||||||
|
*out = __builtin_amdgcn_frexp_mant(a);
|
||||||
|
}
|
||||||
|
|
||||||
|
// CHECK-LABEL: @test_frexp_exp_f32
|
||||||
|
// CHECK: call i32 @llvm.amdgcn.frexp.exp.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
|
||||||
|
void test_frexp_exp_f64(global int* out, double a)
|
||||||
|
{
|
||||||
|
*out = __builtin_amdgcn_frexp_exp(a);
|
||||||
|
}
|
||||||
|
|
||||||
// CHECK-LABEL: @test_class_f32
|
// CHECK-LABEL: @test_class_f32
|
||||||
// CHECK: call i1 @llvm.amdgcn.class.f32
|
// CHECK: call i1 @llvm.amdgcn.class.f32
|
||||||
void test_class_f32(global float* out, float a, int b)
|
void test_class_f32(global float* out, float a, int b)
|
||||||
|
|
Loading…
Reference in New Issue