Implemented the changes based on the feedback accordingly.

Author: SergeantCooper

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

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;

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
-}
+}
+
+; 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
+}