[llvm] 644de6a - Revert "[NVPTX] designate fabs and fneg as free (#121513)"

[Benjamin Kramer via llvm-commits]

Author: Benjamin Kramer
Date: 2025-01-09T16:17:51+01:00
New Revision: 644de6ad1c758d3bf754d7d50b98c555df5231b1

URL: https://github.com/llvm/llvm-project/commit/644de6ad1c758d3bf754d7d50b98c555df5231b1
DIFF: https://github.com/llvm/llvm-project/commit/644de6ad1c758d3bf754d7d50b98c555df5231b1.diff

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

This reverts commit 45d46983bf7bda53bd7ee8e36a47571b3980fbd7.

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.

Added: 
    

Modified: 
    llvm/lib/Target/NVPTX/NVPTXISelLowering.h
    llvm/test/CodeGen/NVPTX/bf16x2-instructions.ll

Removed: 
    llvm/test/CodeGen/NVPTX/fabs-fneg-free.ll


################################################################################
diff  --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.h b/llvm/lib/Target/NVPTX/NVPTXISelLowering.h
index c9b7e874556990..4a98fe21b81dc6 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.h
+++ b/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;

diff  --git a/llvm/test/CodeGen/NVPTX/bf16x2-instructions.ll b/llvm/test/CodeGen/NVPTX/bf16x2-instructions.ll
index 8be3a66b7f4836..03cdeb9683abae 100644
--- a/llvm/test/CodeGen/NVPTX/bf16x2-instructions.ll
+++ b/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)

diff  --git a/llvm/test/CodeGen/NVPTX/fabs-fneg-free.ll b/llvm/test/CodeGen/NVPTX/fabs-fneg-free.ll
deleted file mode 100644
index 9031f33939f2fe..00000000000000
--- a/llvm/test/CodeGen/NVPTX/fabs-fneg-free.ll
+++ /dev/null
@@ -1,34 +0,0 @@
-; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
-; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_35 -verify-machineinstrs | FileCheck %s
-; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_35 | %ptxas-verify %}
-target triple = "nvptx64-nvidia-cuda"
-
-define float @fabs_free(i32 %in) {
-; CHECK-LABEL: fabs_free(
-; CHECK:       {
-; CHECK-NEXT:    .reg .f32 %f<3>;
-; CHECK-EMPTY:
-; CHECK-NEXT:  // %bb.0:
-; CHECK-NEXT:    ld.param.f32 %f1, [fabs_free_param_0];
-; CHECK-NEXT:    abs.f32 %f2, %f1;
-; CHECK-NEXT:    st.param.f32 [func_retval0], %f2;
-; CHECK-NEXT:    ret;
-  %b = bitcast i32 %in to float
-  %f = call float @llvm.fabs.f32(float %b)
-  ret float %f
-}
-
-define float @fneg_free(i32 %in) {
-; CHECK-LABEL: fneg_free(
-; CHECK:       {
-; CHECK-NEXT:    .reg .f32 %f<3>;
-; CHECK-EMPTY:
-; CHECK-NEXT:  // %bb.0:
-; CHECK-NEXT:    ld.param.f32 %f1, [fneg_free_param_0];
-; CHECK-NEXT:    neg.f32 %f2, %f1;
-; CHECK-NEXT:    st.param.f32 [func_retval0], %f2;
-; CHECK-NEXT:    ret;
-  %b = bitcast i32 %in to float
-  %f = fneg float %b
-  ret float %f
-}