replica
/
hanchenye-llvm-project
2
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

AMDGPU: Add Clang Builtin for v_lerp_u8 

Differential Revision: http://reviews.llvm.org/D22380

llvm-svn: 275577
This commit is contained in:
Wei Ding 2016-07-15 16:43:03 +00:00
parent dcafd5ebfd
commit ea41f356bb
3 changed files with 10 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
clang/include/clang/Basic/BuiltinsAMDGPU.def
View File

@ -61,6 +61,7 @@ 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_lerp, "UiUiUiUi", "nc")
BUILTIN(__builtin_amdgcn_class, "bdi", "nc")
BUILTIN(__builtin_amdgcn_classf, "bfi", "nc")
BUILTIN(__builtin_amdgcn_cubeid, "ffff", "nc")

2
clang/lib/CodeGen/CGBuiltin.cpp
View File

@ -7659,6 +7659,8 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
case AMDGPU::BI__builtin_amdgcn_fract:
case AMDGPU::BI__builtin_amdgcn_fractf:
return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_fract);
case AMDGPU::BI__builtin_amdgcn_lerp:
return emitTernaryBuiltin(*this, E, Intrinsic::amdgcn_lerp);
case AMDGPU::BI__builtin_amdgcn_class:
case AMDGPU::BI__builtin_amdgcn_classf:
return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_class);

7
clang/test/CodeGenOpenCL/builtins-amdgcn.cl
View File

@ -192,6 +192,13 @@ void test_fract_f64(global int* out, double a)
*out = __builtin_amdgcn_fract(a);
}
// CHECK-LABEL: @test_lerp
// CHECK: call i32 @llvm.amdgcn.lerp
void test_lerp(global int* out, int a, int b, int c)
{
*out = __builtin_amdgcn_lerp(a, b, c);
}
// CHECK-LABEL: @test_class_f32
// CHECK: call i1 @llvm.amdgcn.class.f32
void test_class_f32(global float* out, float a, int b)