From f43171b27dd59a4d71de4a891cad4006d7fc1a6c Mon Sep 17 00:00:00 2001 From: ujjawalk Date: Sat, 2 Nov 2024 15:07:34 +0530 Subject: [PATCH 1/3] Implemented isZextFree and IsTruncateFree in NVPTX target lowering. --- llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 23 +++++++++++++++++++++ llvm/lib/Target/NVPTX/NVPTXISelLowering.h | 4 ++++ llvm/test/CodeGen/NVPTX/truncate_zext.ll | 17 +++++++++++++++ 3 files changed, 44 insertions(+) create mode 100644 llvm/test/CodeGen/NVPTX/truncate_zext.ll diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp index d3bf0ecfe2cc9..b5fc975d71dfa 100644 --- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp @@ -3340,6 +3340,29 @@ bool NVPTXTargetLowering::splitValueIntoRegisterParts( return false; } +bool llvm::NVPTXTargetLowering::isTruncateFree(EVT FromVT, EVT ToVT) const { + + if (!FromVT.isSimple() || !ToVT.isSimple()) { + return false; + } + + return (FromVT.getSimpleVT() == MVT::i64 && ToVT.getSimpleVT() == MVT::i32); +} + +bool llvm::NVPTXTargetLowering::isZExtFree(EVT FromVT, EVT ToVT) const { + if (!FromVT.isSimple() || !ToVT.isSimple()) { + return false; + } + return (FromVT.getSimpleVT() == MVT::i32 && ToVT.getSimpleVT() == MVT::i64); +} + +bool llvm::NVPTXTargetLowering::isZExtFree(Type *SrcTy, Type *DstTy) const { + if (!SrcTy->isIntegerTy() || !DstTy->isIntegerTy()) + return false; + return SrcTy->getPrimitiveSizeInBits() == 32 && + DstTy->getPrimitiveSizeInBits() == 64; +} + // This creates target external symbol for a function parameter. // Name of the symbol is composed from its index and the function name. // Negative index corresponds to special parameter (unsized array) used for diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.h b/llvm/lib/Target/NVPTX/NVPTXISelLowering.h index c8b589ae39413..fa73938a35a16 100644 --- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.h +++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.h @@ -616,6 +616,10 @@ class NVPTXTargetLowering : public TargetLowering { return true; } + bool isTruncateFree(EVT FromVT, EVT ToVT) const override; + bool isZExtFree(EVT FromVT, EVT ToVT) const override; + bool isZExtFree(Type *SrcTy, Type *DstTy) const override; + private: const NVPTXSubtarget &STI; // cache the subtarget here SDValue getParamSymbol(SelectionDAG &DAG, int idx, EVT) const; diff --git a/llvm/test/CodeGen/NVPTX/truncate_zext.ll b/llvm/test/CodeGen/NVPTX/truncate_zext.ll new file mode 100644 index 0000000000000..decc02c584049 --- /dev/null +++ b/llvm/test/CodeGen/NVPTX/truncate_zext.ll @@ -0,0 +1,17 @@ +; RUN: llc -march=nvptx64 < %s | FileCheck %s + +; Test for truncation from i64 to i32 +define i32 @test_trunc_i64_to_i32(i64 %val) { + ; CHECK-LABEL: test_trunc_i64_to_i32 + ; CHECK: trunc + %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: zext + %zext = zext i32 %val to i64 + ret i64 %zext +} \ No newline at end of file From 401f834f6ecfdf3b335181b783a1575c1697cf8b Mon Sep 17 00:00:00 2001 From: ujjawalk Date: Sun, 3 Nov 2024 02:11:56 +0530 Subject: [PATCH 2/3] fixed the implementations of istruncatefree and iszextfree according to nvptx arch. --- llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 16 +++++----------- 1 file changed, 5 insertions(+), 11 deletions(-) diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp index b5fc975d71dfa..d208caebbd115 100644 --- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp @@ -3341,26 +3341,20 @@ bool NVPTXTargetLowering::splitValueIntoRegisterParts( } bool llvm::NVPTXTargetLowering::isTruncateFree(EVT FromVT, EVT ToVT) const { - - if (!FromVT.isSimple() || !ToVT.isSimple()) { + if (FromVT.isVector() || ToVT.isVector() || !FromVT.isInteger() || + !ToVT.isInteger()) { return false; } - return (FromVT.getSimpleVT() == MVT::i64 && ToVT.getSimpleVT() == MVT::i32); + return FromVT.getSizeInBits() == 64 && ToVT.getSizeInBits() == 32; } bool llvm::NVPTXTargetLowering::isZExtFree(EVT FromVT, EVT ToVT) const { - if (!FromVT.isSimple() || !ToVT.isSimple()) { - return false; - } - return (FromVT.getSimpleVT() == MVT::i32 && ToVT.getSimpleVT() == MVT::i64); + return false; } bool llvm::NVPTXTargetLowering::isZExtFree(Type *SrcTy, Type *DstTy) const { - if (!SrcTy->isIntegerTy() || !DstTy->isIntegerTy()) - return false; - return SrcTy->getPrimitiveSizeInBits() == 32 && - DstTy->getPrimitiveSizeInBits() == 64; + return false; } // This creates target external symbol for a function parameter. From 65500444989f9b2accd8b0e140fe58bc6cec2c00 Mon Sep 17 00:00:00 2001 From: ujjawalk Date: Wed, 6 Nov 2024 13:47:33 +0530 Subject: [PATCH 3/3] Implemented the changes based on the feedback accordingly. --- llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 17 -------- llvm/lib/Target/NVPTX/NVPTXISelLowering.h | 15 +++++-- llvm/test/CodeGen/NVPTX/truncate_zext.ll | 47 ++++++++++++++++++--- 3 files changed, 53 insertions(+), 26 deletions(-) diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp index 1233add1d0374..18b05b23da220 100644 --- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp @@ -3340,23 +3340,6 @@ bool NVPTXTargetLowering::splitValueIntoRegisterParts( return false; } -bool llvm::NVPTXTargetLowering::isTruncateFree(EVT FromVT, EVT ToVT) const { - if (FromVT.isVector() || ToVT.isVector() || !FromVT.isInteger() || - !ToVT.isInteger()) { - return false; - } - - return FromVT.getSizeInBits() == 64 && ToVT.getSizeInBits() == 32; -} - -bool llvm::NVPTXTargetLowering::isZExtFree(EVT FromVT, EVT ToVT) const { - return false; -} - -bool llvm::NVPTXTargetLowering::isZExtFree(Type *SrcTy, Type *DstTy) const { - return false; -} - // This creates target external symbol for a function parameter. // Name of the symbol is composed from its index and the function name. // Negative index corresponds to special parameter (unsized array) used for diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.h b/llvm/lib/Target/NVPTX/NVPTXISelLowering.h index fa73938a35a16..ee7633eb9231a 100644 --- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.h +++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.h @@ -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()) @@ -616,10 +627,6 @@ class NVPTXTargetLowering : public TargetLowering { return true; } - bool isTruncateFree(EVT FromVT, EVT ToVT) const override; - bool isZExtFree(EVT FromVT, EVT ToVT) const override; - bool isZExtFree(Type *SrcTy, Type *DstTy) const override; - private: const NVPTXSubtarget &STI; // cache the subtarget here SDValue getParamSymbol(SelectionDAG &DAG, int idx, EVT) const; diff --git a/llvm/test/CodeGen/NVPTX/truncate_zext.ll b/llvm/test/CodeGen/NVPTX/truncate_zext.ll index decc02c584049..15159916b1850 100644 --- a/llvm/test/CodeGen/NVPTX/truncate_zext.ll +++ b/llvm/test/CodeGen/NVPTX/truncate_zext.ll @@ -1,17 +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 ; Test for truncation from i64 to i32 define i32 @test_trunc_i64_to_i32(i64 %val) { - ; CHECK-LABEL: test_trunc_i64_to_i32 - ; CHECK: trunc +; 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: zext +; 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 -} \ No newline at end of file +} + +; 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 +}