[llvm] [NVPTX] Add truncate and zero-extend free type conversions (PR #125580)

[Baodi Shan via llvm-commits]

https://github.com/lwshanbd created https://github.com/llvm/llvm-project/pull/125580

According to #114339 
Implement isTruncateFree and isZExtFree methods for NVPTX target to indicate that truncating from i64 to i32 and zero-extending from i32 to i64 are free operations. This can help the backend make more efficient code generation decisions.



>From 076245db1f571409c29738fa6a1259e8773d4737 Mon Sep 17 00:00:00 2001
From: Baodi Shan <lwshanbd at gmail.com>
Date: Mon, 3 Feb 2025 15:03:33 -0600
Subject: [PATCH] [NVPTX] Add truncate and zero-extend free type conversions

Implement isTruncateFree and isZExtFree methods for NVPTX target to indicate
that truncating from i64 to i32 and zero-extending from i32 to i64 are
free operations. This can help the backend make more efficient code generation
decisions.
---
 llvm/lib/Target/NVPTX/NVPTXISelLowering.h     | 22 +++++++++++++++++++
 llvm/test/CodeGen/NVPTX/free-truncate-zext.ll | 20 +++++++++++++++++
 2 files changed, 42 insertions(+)
 create mode 100644 llvm/test/CodeGen/NVPTX/free-truncate-zext.ll

diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.h b/llvm/lib/Target/NVPTX/NVPTXISelLowering.h
index 5adf69d621552f3..f4a7ce9d45c0e49 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.h
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.h
@@ -150,6 +150,28 @@ class NVPTXTargetLowering : public TargetLowering {
            DstTy->getPrimitiveSizeInBits() == 32;
   }
 
+  bool isTruncateFree(EVT SrcVT, EVT DstVT) const override {
+    // Truncating from i64 to i32 is free
+    if (SrcVT.isInteger() && DstVT.isInteger())
+      return SrcVT.getSizeInBits() == 64 && DstVT.getSizeInBits() == 32;
+    return false;
+  }
+
+  bool isZExtFree(EVT FromVT, EVT ToVT) const override {
+    // Zero-extending from i32 to i64 is free
+    if (FromVT.isInteger() && ToVT.isInteger())
+      return FromVT.getSizeInBits() == 32 && ToVT.getSizeInBits() == 64;
+    return false;
+  }
+
+  bool isZExtFree(Type *SrcTy, Type *DstTy) const override {
+    // Zero-extending from i32 to i64 is free
+    if (SrcTy->isIntegerTy() && DstTy->isIntegerTy())
+      return SrcTy->getPrimitiveSizeInBits() == 32 &&
+             DstTy->getPrimitiveSizeInBits() == 64;
+    return false;
+  }
+
   EVT getSetCCResultType(const DataLayout &DL, LLVMContext &Ctx,
                          EVT VT) const override {
     if (VT.isVector())
diff --git a/llvm/test/CodeGen/NVPTX/free-truncate-zext.ll b/llvm/test/CodeGen/NVPTX/free-truncate-zext.ll
new file mode 100644
index 000000000000000..b32fe1b727eeda4
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/free-truncate-zext.ll
@@ -0,0 +1,20 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_90 | FileCheck %s
+
+define i32 @test_trunc(i64 %x, i64 %y) {
+; CHECK-LABEL: test_trunc(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.u32 %r1, [test_trunc_param_0];
+; CHECK-NEXT:    ld.param.u32 %r2, [test_trunc_param_1];
+; CHECK-NEXT:    mad.lo.s32 %r3, %r1, %r2, 123456789;
+; CHECK-NEXT:    st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT:    ret;
+  %a = mul i64 %x, %y
+  %b = add i64 %a, 123456789
+  %c = and i64 %b, -1
+  %trunc = trunc i64 %c to i32
+  ret i32 %trunc
+}