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
2

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

6

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

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

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

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
7

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

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
7

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
include/
llvm/
IR/
5 lines
lib/
Target/
AMDGPU/
4 lines
test/
CodeGen/
AMDGPU/
14 lines

Diff 63589

include/llvm/IR/IntrinsicsAMDGPU.td

Context not available.
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
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
Context not available.

lib/Target/AMDGPU/SIInstructions.td

Context not available.
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", []>;
Context not available.
arsenmUnsubmitted
Not Done
ReplyInline Actions

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

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

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

  • This file was added.
; RUN: llc -march=amdgcn -verify-machineinstrs < %s | FileCheck -check-prefix=SI %s
; RUN: llc -march=amdgcn -mcpu=fiji -verify-machineinstrs < %s | FileCheck -check-prefix=SI %s
arsenmUnsubmitted
Done
ReplyInline Actions

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

arsenm: You should also add an SI run line, this also doesn't need to specify hsa and -march=amdgcn is…
arsenmUnsubmitted
Done
ReplyInline Actions

Remove -mcpu=SI

arsenm: Remove -mcpu=SI
arsenmUnsubmitted
Done
ReplyInline Actions

You should use either the default check prefix or GCN

arsenm: You should use either the default check prefix or GCN
declare i32 @llvm.amdgcn.lerp(i32, i32, i32) #0
; FUNC-LABEL: {{^}}v_lerp:
arsenmUnsubmitted
Done
ReplyInline Actions

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

arsenm: you don't specify this check prefix so it should still be CHECK-LABEL
; SI: v_lerp_u8 v{{[0-9]+}}, s{{[0-9]+}}, v{{[0-9]+}}, v{{[0-9]+}}
arsenmUnsubmitted
Done
ReplyInline Actions

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

arsenm: There still is no FUNC check prefix. You don't need separate SI and VI lines because they are…
arsenmUnsubmitted
Done
ReplyInline Actions

This is still using FUNC

arsenm: This is still using FUNC
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 }