parent
909d32955b
commit
2d51059ebb
|
@ -38,6 +38,8 @@ 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_fract, "dd", "nc")
|
||||
BUILTIN(__builtin_amdgcn_fractf, "ff", "nc")
|
||||
BUILTIN(__builtin_amdgcn_class, "bdi", "nc")
|
||||
BUILTIN(__builtin_amdgcn_classf, "bfi", "nc")
|
||||
BUILTIN(__builtin_amdgcn_cubeid, "ffff", "nc")
|
||||
|
|
|
@ -7129,6 +7129,9 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
|
|||
case AMDGPU::BI__builtin_amdgcn_frexp_expf: {
|
||||
return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_frexp_exp);
|
||||
}
|
||||
case AMDGPU::BI__builtin_amdgcn_fract:
|
||||
case AMDGPU::BI__builtin_amdgcn_fractf:
|
||||
return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_fract);
|
||||
case AMDGPU::BI__builtin_amdgcn_class:
|
||||
case AMDGPU::BI__builtin_amdgcn_classf:
|
||||
return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_class);
|
||||
|
|
|
@ -178,6 +178,20 @@ void test_frexp_exp_f64(global int* out, double a)
|
|||
*out = __builtin_amdgcn_frexp_exp(a);
|
||||
}
|
||||
|
||||
// CHECK-LABEL: @test_fract_f32
|
||||
// CHECK: call float @llvm.amdgcn.fract.f32
|
||||
void test_fract_f32(global int* out, float a)
|
||||
{
|
||||
*out = __builtin_amdgcn_fractf(a);
|
||||
}
|
||||
|
||||
// CHECK-LABEL: @test_fract_f64
|
||||
// CHECK: call double @llvm.amdgcn.fract.f64
|
||||
void test_fract_f64(global int* out, double a)
|
||||
{
|
||||
*out = __builtin_amdgcn_fract(a);
|
||||
}
|
||||
|
||||
// CHECK-LABEL: @test_class_f32
|
||||
// CHECK: call i1 @llvm.amdgcn.class.f32
|
||||
void test_class_f32(global float* out, float a, int b)
|
||||
|
|
Loading…
Reference in New Issue