Index: include/llvm/IR/IntrinsicsAMDGPU.td =================================================================== --- include/llvm/IR/IntrinsicsAMDGPU.td +++ include/llvm/IR/IntrinsicsAMDGPU.td @@ -384,6 +384,11 @@ 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 //===----------------------------------------------------------------------===// Index: lib/Target/AMDGPU/SIInstructions.td =================================================================== --- lib/Target/AMDGPU/SIInstructions.td +++ lib/Target/AMDGPU/SIInstructions.td @@ -1717,6 +1717,10 @@ defm V_FMA_F64 : VOP3Inst , "v_fma_f64", VOP_F64_F64_F64_F64, fma >; + +defm V_LERP_U8 : VOP3Inst , "v_lerp_u8", + VOP_I32_I32_I32_I32, int_amdgcn_lerp +>; } // End isCommutable = 1 //def V_LERP_U8 : VOP3_U8 <0x0000014d, "v_lerp_u8", []>; Index: test/CodeGen/AMDGPU/llvm.amdgcn.lerp.ll =================================================================== --- /dev/null +++ test/CodeGen/AMDGPU/llvm.amdgcn.lerp.ll @@ -0,0 +1,15 @@ +; RUN: llc -march=amdgcn -verify-machineinstrs < %s | FileCheck -check-prefix=SI %s +; RUN: llc -march=amdgcn -mcpu=fiji -verify-machineinstrs < %s | FileCheck -check-prefix=VI %s + +declare i32 @llvm.amdgcn.lerp(i32, i32, i32) #0 + +; FUNC-LABEL: {{^}}v_lerp: +; SI: v_lerp_u8 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}} +; VI: 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 }