From a26d20c07a4d930d72b47004e5cd65ed56febd97 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 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@275577 91177308-0d34-0410-b5e6-96231b3b80d8 --- include/clang/Basic/BuiltinsAMDGPU.def | 1 + lib/CodeGen/CGBuiltin.cpp | 2 ++ test/CodeGenOpenCL/builtins-amdgcn.cl | 7 +++++++ 3 files changed, 10 insertions(+) diff --git a/include/clang/Basic/BuiltinsAMDGPU.def b/include/clang/Basic/BuiltinsAMDGPU.def index 0c0b8865f7..0c8bce8e68 100644 --- 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 42fa8ae79f..7a869f203e 100644 --- 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 ce463ae9d6..2066f0b1da 100644 --- 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) -- 2.40.0