diff options
| author | Wei Ding <wei.ding2@amd.com> | 2016-07-12 18:02:14 +0000 |
|---|---|---|
| committer | Wei Ding <wei.ding2@amd.com> | 2016-07-12 18:02:14 +0000 |
| commit | 5b2636a15277b2499f2bc1f4c1f009f9ed1bd3f1 (patch) | |
| tree | 317cef08d190efcbdb640035fb9436a4f17648eb | |
| parent | 98c0f482d66624d72a4186a58f535536238bee8a (diff) | |
| download | bcm5719-llvm-5b2636a15277b2499f2bc1f4c1f009f9ed1bd3f1.tar.gz bcm5719-llvm-5b2636a15277b2499f2bc1f4c1f009f9ed1bd3f1.zip | |
AMDGPU: Add LLVM IR Intrinsic for v_lerp_u8
Differential Revision: http://reviews.llvm.org/D22239
llvm-svn: 275197
| -rw-r--r-- | llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 5 | ||||
| -rw-r--r-- | llvm/lib/Target/AMDGPU/SIInstructions.td | 4 | ||||
| -rw-r--r-- | llvm/test/CodeGen/AMDGPU/llvm.amdgcn.lerp.ll | 14 |
3 files changed, 23 insertions, 0 deletions
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td index 119ce807a45..1473146cd92 100644 --- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td +++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td @@ -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 //===----------------------------------------------------------------------===// diff --git a/llvm/lib/Target/AMDGPU/SIInstructions.td b/llvm/lib/Target/AMDGPU/SIInstructions.td index 05e0e3e504c..7cf5faa216d 100644 --- a/llvm/lib/Target/AMDGPU/SIInstructions.td +++ b/llvm/lib/Target/AMDGPU/SIInstructions.td @@ -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", []>; diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.lerp.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.lerp.ll new file mode 100644 index 00000000000..014369b4501 --- /dev/null +++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.lerp.ll @@ -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 } |
