Skip to content

[NVPTX] Implement isTruncateFree and isZExtFree for i32/i64 Optimizations #114683

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Closed
wants to merge 4 commits into from
Closed

[NVPTX] Implement isTruncateFree and isZExtFree for i32/i64 Optimizations #114683

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
f43171b
Implemented isZextFree and IsTruncateFree in NVPTX target lowering.
SergeantCooper Nov 2, 2024
401f834
fixed the implementations of istruncatefree and iszextfree according …
SergeantCooper Nov 2, 2024
519453e
Merge branch 'llvm:main' into dev
SergeantCooper Nov 6, 2024
6550044
Implemented the changes based on the feedback accordingly.
SergeantCooper Nov 6, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
11 changes: 11 additions & 0 deletions llvm/lib/Target/NVPTX/NVPTXISelLowering.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -506,6 +506,17 @@ class NVPTXTargetLowering : public TargetLowering {
DstTy->getPrimitiveSizeInBits() == 32;
}

bool isTruncateFree(EVT FromVT, EVT ToVT) const override {
if (!FromVT.isScalarInteger() || !ToVT.isScalarInteger()) {
return false;
}
return FromVT.getSizeInBits() == 64 && ToVT.getSizeInBits() == 32;
}

bool isZExtFree(EVT FromVT, EVT ToVT) const override { return false; }

bool isZExtFree(Type *SrcTy, Type *DstTy) const override { return false; }

EVT getSetCCResultType(const DataLayout &DL, LLVMContext &Ctx,
EVT VT) const override {
if (VT.isVector())
Expand Down
54 changes: 54 additions & 0 deletions llvm/test/CodeGen/NVPTX/truncate_zext.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,54 @@
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
; RUN: llc -march=nvptx64 < %s | FileCheck %s
Copy link
Member

@Artem-B Artem-B Nov 4, 2024

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

You may want to use llvm/utils/update_llc_test_checks.py to generate the checks.


Copy link
Contributor

@justinfargnoli justinfargnoli Nov 4, 2024

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This test looks like it's checking that the correct PTX is generated for trunc and zext. This is of course nice to have, but not what we need for this PR.

The purpose of the test associated with this PR should be to ensure that your implementation of isTruncateFree() has an impact. Thus, the test should fail without your change and pass with it.

I like to find a place where isTruncateFree(EVT, EVT) is used in a simple peephole and check against that. For example, you could piggyback off of this peephole (Added and tested in 3332b70).

Copy link
Member

@Artem-B Artem-B Nov 4, 2024

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

+1 to that. The test should have some code which has to choose between free ops vs an alternative which would be used otherwise (e.g. i32 trunc(add(i64,i64)) vs i32 add(trunc(i64), trunc(i64)))

; Test for truncation from i64 to i32
define i32 @test_trunc_i64_to_i32(i64 %val) {
; CHECK-LABEL: test_trunc_i64_to_i32(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u32 %r1, [test_trunc_i64_to_i32_param_0];
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
; CHECK-NEXT: ret;
%trunc = trunc i64 %val to i32
ret i32 %trunc
}

; Test for zero-extension from i32 to i64
define i64 @test_zext_i32_to_i64(i32 %val) {
; CHECK-LABEL: test_zext_i32_to_i64(
; CHECK: {
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u32 %rd1, [test_zext_i32_to_i64_param_0];
; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;
; CHECK-NEXT: ret;
%zext = zext i32 %val to i64
ret i64 %zext
}

; Test for operand truncation before select
define i32 @test_select_truncate_free(i1 %cond, i64 %a, i64 %b) {
; CHECK-LABEL: test_select_truncate_free(
; CHECK: {
; CHECK-NEXT: .reg .pred %p<2>;
; CHECK-NEXT: .reg .b16 %rs<3>;
; CHECK-NEXT: .reg .b32 %r<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u8 %rs1, [test_select_truncate_free_param_0];
; CHECK-NEXT: and.b16 %rs2, %rs1, 1;
; CHECK-NEXT: setp.eq.b16 %p1, %rs2, 1;
; CHECK-NEXT: ld.param.u32 %r1, [test_select_truncate_free_param_1];
; CHECK-NEXT: ld.param.u32 %r2, [test_select_truncate_free_param_2];
; CHECK-NEXT: selp.b32 %r3, %r1, %r2, %p1;
; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
; CHECK-NEXT: ret;

%trunc_a = trunc i64 %a to i32
%trunc_b = trunc i64 %b to i32
%result = select i1 %cond, i32 %trunc_a, i32 %trunc_b
ret i32 %result
}
Loading