AMDGPU: Add Clang Builtin for v_lerp_u8

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

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@275577 91177308-0d34-0410-b5e6-96231b3b80d8

diff --git a/include/clang/Basic/BuiltinsAMDGPU.def b/include/clang/Basic/BuiltinsAMDGPU.def
index 0c0b8865f7ec30c972c660f9a47547e8f69ff70d..0c8bce8e68ea6d80651d2e7d8708c0109bf8919a 100644 (file)
--- a/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/include/clang/Basic/BuiltinsAMDGPU.def
@@ -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")

diff --git a/lib/CodeGen/CGBuiltin.cpp b/lib/CodeGen/CGBuiltin.cpp
index 42fa8ae79f1b819031c45ce4317f86692a53a9b2..7a869f203e0513ea856a9550c65982c1fa4d8e0e 100644 (file)
--- a/lib/CodeGen/CGBuiltin.cpp
+++ b/lib/CodeGen/CGBuiltin.cpp
@@ -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);

diff --git a/test/CodeGenOpenCL/builtins-amdgcn.cl b/test/CodeGenOpenCL/builtins-amdgcn.cl
index ce463ae9d6b9356c3e68c2bfb6c0f2a95d2b645f..2066f0b1dad724b348052ca45f3edd4d247bab7a 100644 (file)
--- a/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ b/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -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)