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

[via llvm-commits]

llvmbot wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-backend-nvptx

Author: Baodi Shan (lwshanbd)

<details>
<summary>Changes</summary>

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.



---
Full diff: https://github.com/llvm/llvm-project/pull/125580.diff


2 Files Affected:

- (modified) llvm/lib/Target/NVPTX/NVPTXISelLowering.h (+22) 
- (added) llvm/test/CodeGen/NVPTX/free-truncate-zext.ll (+20) 


``````````diff
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.h b/llvm/lib/Target/NVPTX/NVPTXISelLowering.h
index 5adf69d621552f..f4a7ce9d45c0e4 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 00000000000000..b32fe1b727eeda
--- /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
+}

``````````

</details>


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