RPCS3/llvm-mirror
1
0
Fork 0
mirror of https://github.com/RPCS3/llvm-mirror.git synced 2024-10-19 11:02:59 +02:00
Code Issues Projects Releases Wiki Activity

AMDGPU: Add LLVM IR Intrinsic for v_lerp_u8

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

llvm-svn: 275197
This commit is contained in:
Wei Ding 2016-07-12 18:02:14 +00:00
parent ec9745f877
commit 3dc39df64c
3 changed files with 23 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

5
include/llvm/IR/IntrinsicsAMDGPU.td
View File

@ -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
//===----------------------------------------------------------------------===//

4
lib/Target/AMDGPU/SIInstructions.td
View File

@ -1717,6 +1717,10 @@ defm V_FMA_F32 : VOP3Inst <vop3<0x14b, 0x1cb>, "v_fma_f32",
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", []>;

14
test/CodeGen/AMDGPU/llvm.amdgcn.lerp.ll Normal file
View File

@ -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 }