Revert "[NVPTX] designate fabs and fneg as free (#121513)"

This reverts commit 45d4698.

NVPTX fabs & fneg are incompatible with LLVM's semantics as LLVM
guarantees the payload of NaNs to stay the same while PTX mangles NaNs.

The bad patterns are still in the NVPTX backend and should probably be
removed, since this change only exposed the bad behavior.

Author: d0k

llvm/lib/Target/NVPTX/NVPTXISelLowering.h

@@ -261,9 +261,6 @@ class NVPTXTargetLowering : public TargetLowering {
     return true;
   }
 
-  bool isFAbsFree(EVT VT) const override { return true; }
-  bool isFNegFree(EVT VT) const override { return true; }
-
 private:
   const NVPTXSubtarget &STI; // cache the subtarget here
   SDValue getParamSymbol(SelectionDAG &DAG, int idx, EVT) const;

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

@@ -182,8 +182,8 @@ define <2 x bfloat> @test_fneg(<2 x bfloat> %a) #0 {
 ; CHECK-NEXT:    .reg .b32 %r<3>;
 ; CHECK-EMPTY:
 ; CHECK-NEXT:  // %bb.0:
-; CHECK-NEXT:    ld.param.b32 %r1, [test_fneg_param_0];
-; CHECK-NEXT:    neg.bf16x2 %r2, %r1;
+; CHECK-NEXT:    ld.param.u32 %r1, [test_fneg_param_0];
+; CHECK-NEXT:    xor.b32 %r2, %r1, -2147450880;
 ; CHECK-NEXT:    st.param.b32 [func_retval0], %r2;
 ; CHECK-NEXT:    ret;
   %r = fneg <2 x bfloat> %a
@@ -532,8 +532,8 @@ define <2 x bfloat> @test_fabs(<2 x bfloat> %a) #0 {
 ; CHECK-NEXT:    .reg .b32 %r<3>;
 ; CHECK-EMPTY:
 ; CHECK-NEXT:  // %bb.0:
-; CHECK-NEXT:    ld.param.b32 %r1, [test_fabs_param_0];
-; CHECK-NEXT:    abs.bf16x2 %r2, %r1;
+; CHECK-NEXT:    ld.param.u32 %r1, [test_fabs_param_0];
+; CHECK-NEXT:    and.b32 %r2, %r1, 2147450879;
 ; CHECK-NEXT:    st.param.b32 [func_retval0], %r2;
 ; CHECK-NEXT:    ret;
   %r = call <2 x bfloat> @llvm.fabs.f16(<2 x bfloat> %a)