[DAGCombiner] Fold [us]itofp of truncate

Author: AlexMaclean

llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp

@@ -18727,6 +18727,12 @@ SDValue DAGCombiner::visitSINT_TO_FP(SDNode *N) {
   if (SDValue FTrunc = foldFPToIntToFP(N, DL, DAG, TLI))
     return FTrunc;
 
+  // fold (sint_to_fp (trunc nsw x)) -> (sint_to_fp x)
+  if (N0.getOpcode() == ISD::TRUNCATE && N0->getFlags().hasNoSignedWrap() &&
+      TLI.isTypeDesirableForOp(ISD::SINT_TO_FP,
+                               N0.getOperand(0).getValueType()))
+    return DAG.getNode(ISD::SINT_TO_FP, DL, VT, N0.getOperand(0));
+
   return SDValue();
 }
 
@@ -18764,6 +18770,12 @@ SDValue DAGCombiner::visitUINT_TO_FP(SDNode *N) {
   if (SDValue FTrunc = foldFPToIntToFP(N, DL, DAG, TLI))
     return FTrunc;
 
+  // fold (uint_to_fp (trunc nuw x)) -> (uint_to_fp x)
+  if (N0.getOpcode() == ISD::TRUNCATE && N0->getFlags().hasNoUnsignedWrap() &&
+      TLI.isTypeDesirableForOp(ISD::UINT_TO_FP,
+                               N0.getOperand(0).getValueType()))
+    return DAG.getNode(ISD::UINT_TO_FP, DL, VT, N0.getOperand(0));
+
   return SDValue();
 }
 

llvm/test/CodeGen/NVPTX/i8x4-instructions.ll

@@ -2314,23 +2314,18 @@ entry:
 define <4 x float> @test_uitofp_v4i8(<4 x i8> %a) {
 ; CHECK-LABEL: test_uitofp_v4i8(
 ; CHECK:       {
-; CHECK-NEXT:    .reg .b16 %rs<5>;
 ; CHECK-NEXT:    .reg .b32 %r<10>;
 ; CHECK-EMPTY:
 ; CHECK-NEXT:  // %bb.0:
 ; CHECK-NEXT:    ld.param.b32 %r1, [test_uitofp_v4i8_param_0];
 ; CHECK-NEXT:    prmt.b32 %r2, %r1, 0, 0x7773U;
-; CHECK-NEXT:    cvt.u16.u32 %rs1, %r2;
-; CHECK-NEXT:    cvt.rn.f32.u16 %r3, %rs1;
+; CHECK-NEXT:    cvt.rn.f32.u32 %r3, %r2;
 ; CHECK-NEXT:    prmt.b32 %r4, %r1, 0, 0x7772U;
-; CHECK-NEXT:    cvt.u16.u32 %rs2, %r4;
-; CHECK-NEXT:    cvt.rn.f32.u16 %r5, %rs2;
+; CHECK-NEXT:    cvt.rn.f32.u32 %r5, %r4;
 ; CHECK-NEXT:    prmt.b32 %r6, %r1, 0, 0x7771U;
-; CHECK-NEXT:    cvt.u16.u32 %rs3, %r6;
-; CHECK-NEXT:    cvt.rn.f32.u16 %r7, %rs3;
+; CHECK-NEXT:    cvt.rn.f32.u32 %r7, %r6;
 ; CHECK-NEXT:    prmt.b32 %r8, %r1, 0, 0x7770U;
-; CHECK-NEXT:    cvt.u16.u32 %rs4, %r8;
-; CHECK-NEXT:    cvt.rn.f32.u16 %r9, %rs4;
+; CHECK-NEXT:    cvt.rn.f32.u32 %r9, %r8;
 ; CHECK-NEXT:    st.param.v4.b32 [func_retval0], {%r9, %r7, %r5, %r3};
 ; CHECK-NEXT:    ret;
   %r = uitofp <4 x i8> %a to <4 x float>

llvm/test/CodeGen/NVPTX/trunc-tofp.ll

@@ -7,15 +7,13 @@ target triple = "nvptx64-nvidia-cuda"
 define float @uitofp_trunc_nuw(i32 %x, i32 %y) {
 ; CHECK-LABEL: uitofp_trunc_nuw(
 ; CHECK:       {
-; CHECK-NEXT:    .reg .b16 %rs<2>;
 ; CHECK-NEXT:    .reg .b32 %r<5>;
 ; CHECK-EMPTY:
 ; CHECK-NEXT:  // %bb.0:
 ; CHECK-NEXT:    ld.param.b32 %r1, [uitofp_trunc_nuw_param_0];
 ; CHECK-NEXT:    ld.param.b32 %r2, [uitofp_trunc_nuw_param_1];
 ; CHECK-NEXT:    add.s32 %r3, %r1, %r2;
-; CHECK-NEXT:    cvt.u16.u32 %rs1, %r3;
-; CHECK-NEXT:    cvt.rn.f32.u16 %r4, %rs1;
+; CHECK-NEXT:    cvt.rn.f32.u32 %r4, %r3;
 ; CHECK-NEXT:    st.param.b32 [func_retval0], %r4;
 ; CHECK-NEXT:    ret;
   %v = add i32 %x, %y
@@ -27,15 +25,13 @@ define float @uitofp_trunc_nuw(i32 %x, i32 %y) {
 define float @sitofp_trunc_nsw(i32 %x, i32 %y) {
 ; CHECK-LABEL: sitofp_trunc_nsw(
 ; CHECK:       {
-; CHECK-NEXT:    .reg .b16 %rs<2>;
 ; CHECK-NEXT:    .reg .b32 %r<5>;
 ; CHECK-EMPTY:
 ; CHECK-NEXT:  // %bb.0:
 ; CHECK-NEXT:    ld.param.b32 %r1, [sitofp_trunc_nsw_param_0];
 ; CHECK-NEXT:    ld.param.b32 %r2, [sitofp_trunc_nsw_param_1];
 ; CHECK-NEXT:    add.s32 %r3, %r1, %r2;
-; CHECK-NEXT:    cvt.u16.u32 %rs1, %r3;
-; CHECK-NEXT:    cvt.rn.f32.s16 %r4, %rs1;
+; CHECK-NEXT:    cvt.rn.f32.s32 %r4, %r3;
 ; CHECK-NEXT:    st.param.b32 [func_retval0], %r4;
 ; CHECK-NEXT:    ret;
   %v = add i32 %x, %y