This is an archive of the discontinued LLVM Phabricator instance.

Differential D22239

AMDGPU: Add LLVM IR Intrinsic for v_lerp_u8
ClosedPublic

Authored by wdng on Jul 11 2016, 1:35 PM.

Details

Reviewers
cfang
• tstellarAMD
arsenm
Commits
rG5b2636a15277: AMDGPU: Add LLVM IR Intrinsic for v_lerp_u8
rL275197: AMDGPU: Add LLVM IR Intrinsic for v_lerp_u8
Summary

AMDGPU: Add LLVM IR Intrinsic implementation and a LIT test for v_lerp_u8

Diff Detail

Repository
rL LLVM

Event Timeline

wdng updated this revision to Diff 63563.Jul 11 2016, 1:35 PM
wdng retitled this revision from to AMDGPU: Add LLVM IR Intrinsic for v_lerp_u8.
wdng updated this object.
wdng added reviewers: arsenm, • tstellarAMD, cfang.
wdng set the repository for this revision to rL LLVM.
wdng added a project: Restricted Project.
Herald added subscribers: kzhuravl, arsenm. · View Herald TranscriptJul 11 2016, 1:35 PM
arsenm added inline comments.Jul 11 2016, 1:40 PM
test/CodeGen/AMDGPU/llvm.amdgcn.lerp.ll
1 ↗(On Diff #63563)

You should also add an SI run line, this also doesn't need to specify hsa and -march=amdgcn is enough

5 ↗(On Diff #63563)

you don't specify this check prefix so it should still be CHECK-LABEL

wdng updated this revision to Diff 63567.Jul 11 2016, 1:49 PM
wdng marked 2 inline comments as done.

Changes based on Matt's comments.

arsenm added inline comments.Jul 11 2016, 2:23 PM
lib/Target/AMDGPU/SIInstructions.td
2359–2366 ↗(On Diff #63567)

This pattern is redundant with the one for the instruction definition and should be removed

test/CodeGen/AMDGPU/llvm.amdgcn.lerp.ll
1 ↗(On Diff #63567)

Remove -mcpu=SI

wdng updated this revision to Diff 63579.Jul 11 2016, 2:43 PM

Removed redundant codes.

wdng updated this revision to Diff 63580.Jul 11 2016, 2:45 PM

Updated code again due to git cache is not updated properly.

arsenm added inline comments.Jul 11 2016, 2:51 PM
test/CodeGen/AMDGPU/llvm.amdgcn.lerp.ll
6 ↗(On Diff #63580)

There still is no FUNC check prefix. You don't need separate SI and VI lines because they are identitcal

wdng updated this revision to Diff 63589.Jul 11 2016, 3:00 PM

Removed redundant func check lines.

arsenm added inline comments.Jul 11 2016, 3:09 PM
test/CodeGen/AMDGPU/llvm.amdgcn.lerp.ll
1–2 ↗(On Diff #63589)

You should use either the default check prefix or GCN

wdng updated this revision to Diff 63598.Jul 11 2016, 3:22 PM
wdng marked 3 inline comments as done.

Use default CHECK for LIT tests.

arsenm added inline comments.Jul 11 2016, 3:49 PM
test/CodeGen/AMDGPU/llvm.amdgcn.lerp.ll
6 ↗(On Diff #63598)

This is still using FUNC

wdng updated this revision to Diff 63604.Jul 11 2016, 3:58 PM

Use GCN LABEL.

arsenm accepted this revision.Jul 11 2016, 4:28 PM
arsenm edited edge metadata.

LGTM

This revision is now accepted and ready to land.Jul 11 2016, 4:28 PM
Closed by commit rL275197: AMDGPU: Add LLVM IR Intrinsic for v_lerp_u8 (authored by wdng). · Explain WhyJul 12 2016, 11:09 AM
This revision was automatically updated to reflect the committed changes.
wdng marked an inline comment as done.

Revision Contents

PathSize
llvm/
trunk/
include/
llvm/
IR/
5 lines
lib/
Target/
AMDGPU/
4 lines
test/
CodeGen/
AMDGPU/
14 lines

Diff 63699

llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td

Show First 20 Lines • Show All 378 Lines • ▼ Show 20 Linesdef int_amdgcn_mbcnt_hi :
GCCBuiltin<"__builtin_amdgcn_mbcnt_hi">, GCCBuiltin<"__builtin_amdgcn_mbcnt_hi">,
Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>; Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
// llvm.amdgcn.ds.swizzle src offset // llvm.amdgcn.ds.swizzle src offset
def int_amdgcn_ds_swizzle : def int_amdgcn_ds_swizzle :
GCCBuiltin<"__builtin_amdgcn_ds_swizzle">, GCCBuiltin<"__builtin_amdgcn_ds_swizzle">,
Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrConvergent]>; 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 // CI+ Intrinsics
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
def int_amdgcn_s_dcache_inv_vol : def int_amdgcn_s_dcache_inv_vol :
GCCBuiltin<"__builtin_amdgcn_s_dcache_inv_vol">, GCCBuiltin<"__builtin_amdgcn_s_dcache_inv_vol">,
Intrinsic<[], [], []>; Intrinsic<[], [], []>;
Show All 35 Lines

llvm/trunk/lib/Target/AMDGPU/SIInstructions.td

Show First 20 Lines • Show All 1,711 Lines • ▼ Show 20 Lines
let isCommutable = 1 in { let isCommutable = 1 in {
defm V_FMA_F32 : VOP3Inst <vop3<0x14b, 0x1cb>, "v_fma_f32", defm V_FMA_F32 : VOP3Inst <vop3<0x14b, 0x1cb>, "v_fma_f32",
VOP_F32_F32_F32_F32, fma VOP_F32_F32_F32_F32, fma
>; >;
defm V_FMA_F64 : VOP3Inst <vop3<0x14c, 0x1cc>, "v_fma_f64", defm V_FMA_F64 : VOP3Inst <vop3<0x14c, 0x1cc>, "v_fma_f64",
VOP_F64_F64_F64_F64, fma 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 } // End isCommutable = 1
//def V_LERP_U8 : VOP3_U8 <0x0000014d, "v_lerp_u8", []>; //def V_LERP_U8 : VOP3_U8 <0x0000014d, "v_lerp_u8", []>;
defm V_ALIGNBIT_B32 : VOP3Inst <vop3<0x14e, 0x1ce>, "v_alignbit_b32", defm V_ALIGNBIT_B32 : VOP3Inst <vop3<0x14e, 0x1ce>, "v_alignbit_b32",
VOP_I32_I32_I32_I32 VOP_I32_I32_I32_I32
>; >;
defm V_ALIGNBYTE_B32 : VOP3Inst <vop3<0x14f, 0x1cf>, "v_alignbyte_b32", defm V_ALIGNBYTE_B32 : VOP3Inst <vop3<0x14f, 0x1cf>, "v_alignbyte_b32",
VOP_I32_I32_I32_I32 VOP_I32_I32_I32_I32
▲ Show 20 Lines • Show All 1,823 Lines • Show Last 20 Lines

llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.lerp.ll

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