[llvm] r275197 - AMDGPU: Add LLVM IR Intrinsic for v_lerp_u8

[Wei Ding via llvm-commits]

Author: wdng
Date: Tue Jul 12 13:02:14 2016
New Revision: 275197

URL: http://llvm.org/viewvc/llvm-project?rev=275197&view=rev
Log:
AMDGPU: Add LLVM IR Intrinsic for v_lerp_u8

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

Added:
    llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.lerp.ll
Modified:
    llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td
    llvm/trunk/lib/Target/AMDGPU/SIInstructions.td

Modified: llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td?rev=275197&r1=275196&r2=275197&view=diff
==============================================================================
--- llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td (original)
+++ llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td Tue Jul 12 13:02:14 2016
@@ -384,6 +384,11 @@ def int_amdgcn_ds_swizzle :
   GCCBuiltin<"__builtin_amdgcn_ds_swizzle">,
   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
 
+// llvm.amdgcn.lerp
+def int_amdgcn_lerp :
+  GCCBuiltin<"__builtin_amdgcn_lerp">,
+  Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
+
 //===----------------------------------------------------------------------===//
 // CI+ Intrinsics
 //===----------------------------------------------------------------------===//

Modified: llvm/trunk/lib/Target/AMDGPU/SIInstructions.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/SIInstructions.td?rev=275197&r1=275196&r2=275197&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/SIInstructions.td (original)
+++ llvm/trunk/lib/Target/AMDGPU/SIInstructions.td Tue Jul 12 13:02:14 2016
@@ -1717,6 +1717,10 @@ defm V_FMA_F32 : VOP3Inst <vop3<0x14b, 0
 defm V_FMA_F64 : VOP3Inst <vop3<0x14c, 0x1cc>, "v_fma_f64",
   VOP_F64_F64_F64_F64, fma
 >;
+
+defm V_LERP_U8 : VOP3Inst <vop3<0x14d, 0x1cd>, "v_lerp_u8",
+  VOP_I32_I32_I32_I32, int_amdgcn_lerp
+>;
 } // End isCommutable = 1
 
 //def V_LERP_U8 : VOP3_U8 <0x0000014d, "v_lerp_u8", []>;

Added: llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.lerp.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.lerp.ll?rev=275197&view=auto
==============================================================================
--- llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.lerp.ll (added)
+++ llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.lerp.ll Tue Jul 12 13:02:14 2016
@@ -0,0 +1,14 @@
+; RUN: llc -march=amdgcn -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
+; RUN: llc -march=amdgcn -mcpu=fiji -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
+
+declare i32 @llvm.amdgcn.lerp(i32, i32, i32) #0
+
+; GCN-LABEL: {{^}}v_lerp:
+; GCN: v_lerp_u8 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}}
+define void @v_lerp(i32 addrspace(1)* %out, i32 %src) nounwind {
+  %result= call i32 @llvm.amdgcn.lerp(i32 %src, i32 100, i32 100) #0
+  store i32 %result, i32 addrspace(1)* %out, align 4
+  ret void
+}
+
+attributes #0 = { nounwind readnone }