AMDGPU: Add Clang Builtin for v_lerp_u8

[wdng]

Summary

Add Clang Builtin for v_lerp_u8

[wdng]

Fixed an indentation issue.

[arsenm]

LGTM

cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def

 BUILTIN(__builtin_amdgcn_ldexp, "ddi", "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_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")
 BUILTIN(__builtin_amdgcn_cubesc, "ffff", "nc")
 BUILTIN(__builtin_amdgcn_cubetc, "ffff", "nc")
 BUILTIN(__builtin_amdgcn_cubema, "ffff", "nc")
 BUILTIN(__builtin_amdgcn_s_memtime, "LUi", "n")
 BUILTIN(__builtin_amdgcn_s_sleep, "vIi", "n")

cfe/trunk/lib/CodeGen/CGBuiltin.cpp

Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
   }
   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_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);
 
   case AMDGPU::BI__builtin_amdgcn_read_exec: {
     CallInst *CI = cast<CallInst>(
       EmitSpecialRegisterBuiltin(*this, E, Int64Ty, Int64Ty, true, "exec"));
     CI->setConvergent();

cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl

 
 // 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_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)
 {
   *out = __builtin_amdgcn_classf(a, b);
 }
 
 // CHECK-LABEL: @test_class_f64