From ea41f356bbeb156de36ef2eebe0dcbf5b117ce53 Mon Sep 17 00:00:00 2001 From: Wei Ding Date: Fri, 15 Jul 2016 16:43:03 +0000 Subject: [PATCH] AMDGPU: Add Clang Builtin for v_lerp_u8 Differential Revision: http://reviews.llvm.org/D22380 llvm-svn: 275577 --- clang/include/clang/Basic/BuiltinsAMDGPU.def | 1 + clang/lib/CodeGen/CGBuiltin.cpp | 2 ++ clang/test/CodeGenOpenCL/builtins-amdgcn.cl | 7 +++++++ 3 files changed, 10 insertions(+) diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def index 0c0b8865f7ec..0c8bce8e68ea 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.def +++ b/clang/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/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp index 42fa8ae79f1b..7a869f203e05 100644 --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/clang/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/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl index ce463ae9d6b9..2066f0b1dad7 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl +++ b/clang/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)