[llvm] [NVPTX] Fix lowering of fabs and fneg (PR #219089)

[Lewis Crawford via llvm-commits]

================
@@ -68,3 +68,42 @@ define double @sub_f64(double %a, double %b) {
 
   ret double %r4
 }
+
+; A negated operand that is known to never be a NaN reaches isel as the native
+; neg rather than a sign-bit xor. The fold has to apply to that form too.
+define float @sub_f32_never_nan(float %a, i32 %i) {
+; CHECK-LABEL: sub_f32_never_nan(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b32 %r<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b32 %r1, [sub_f32_never_nan_param_0];
+; CHECK-NEXT:    ld.param.b32 %r2, [sub_f32_never_nan_param_1];
+; CHECK-NEXT:    cvt.rn.f32.s32 %r3, %r2;
+; CHECK-NEXT:    sub.rn.f32 %r4, %r1, %r3;
+; CHECK-NEXT:    st.param.b32 [func_retval0], %r4;
+; CHECK-NEXT:    ret;
+  %b = sitofp i32 %i to float
----------------
LewisCrawford wrote:

Why is the `sitofp` necessary here? Is the `nvvm.add.rn` not marked as a canonicalizing consumer, making the `fneg` a candidate for a native PTX `neg` already, even without ensuring a non-NaN input? If not, should it be marked as canonicalizing too (along with other similar add/mul intrinsics)?

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