[llvm] [DAGCombiner] Fold and/or of NaN SETCC (PR #135645)

[via llvm-commits]

llvmbot wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-llvm-selectiondag

Author: Alex MacLean (AlexMaclean)

<details>
<summary>Changes</summary>

Fold an AND or OR of two NaN SETCC nodes into a single SETCC where possible. This optimization already exists in InstCombine but adding in here as well can allow for additional folding if more logical operations are exposed.

---
Full diff: https://github.com/llvm/llvm-project/pull/135645.diff


2 Files Affected:

- (modified) llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp (+6) 
- (added) llvm/test/CodeGen/NVPTX/and-or-setcc.ll (+45) 


``````````diff
diff --git a/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp b/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp
index 8136f1794775e..8eb3f95a30989 100644
--- a/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp
@@ -6427,6 +6427,12 @@ static SDValue foldAndOrOfSETCC(SDNode *LogicOp, SelectionDAG &DAG) {
     }
   }
 
+  if (LHS0 == LHS1 && RHS0 == RHS1 && CCL == CCR &&
+      LHS0.getValueType() == RHS0.getValueType() &&
+      ((LogicOp->getOpcode() == ISD::AND && CCL == ISD::SETO) ||
+       (LogicOp->getOpcode() == ISD::OR && CCL == ISD::SETUO)))
+    return DAG.getSetCC(DL, VT, LHS0, RHS0, CCL);
+
   if (TargetPreference == AndOrSETCCFoldKind::None)
     return SDValue();
 
diff --git a/llvm/test/CodeGen/NVPTX/and-or-setcc.ll b/llvm/test/CodeGen/NVPTX/and-or-setcc.ll
new file mode 100644
index 0000000000000..21be9df94d553
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/and-or-setcc.ll
@@ -0,0 +1,45 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
+; RUN: llc < %s -mtriple=nvptx64 | FileCheck %s
+; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 | %ptxas-verify %}
+
+target triple = "nvptx64-nvidia-cuda"
+
+define i1 @and_ord(float %a, float %b) {
+; CHECK-LABEL: and_ord(
+; CHECK:       {
+; CHECK-NEXT:    .reg .pred %p<2>;
+; CHECK-NEXT:    .reg .b32 %r<2>;
+; CHECK-NEXT:    .reg .f32 %f<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.f32 %f1, [and_ord_param_0];
+; CHECK-NEXT:    ld.param.f32 %f2, [and_ord_param_1];
+; CHECK-NEXT:    setp.num.f32 %p1, %f1, %f2;
+; CHECK-NEXT:    selp.b32 %r1, 1, 0, %p1;
+; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT:    ret;
+  %c = fcmp ord float %a, 0.0
+  %d = fcmp ord float %b, 0.0
+  %e = and i1 %c, %d
+  ret i1 %e
+}
+
+define i1 @or_uno(float %a, float %b) {
+; CHECK-LABEL: or_uno(
+; CHECK:       {
+; CHECK-NEXT:    .reg .pred %p<2>;
+; CHECK-NEXT:    .reg .b32 %r<2>;
+; CHECK-NEXT:    .reg .f32 %f<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.f32 %f1, [or_uno_param_0];
+; CHECK-NEXT:    ld.param.f32 %f2, [or_uno_param_1];
+; CHECK-NEXT:    setp.nan.f32 %p1, %f1, %f2;
+; CHECK-NEXT:    selp.b32 %r1, 1, 0, %p1;
+; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT:    ret;
+  %c = fcmp uno float %a, 0.0
+  %d = fcmp uno float %b, 0.0
+  %e = or i1 %c, %d
+  ret i1 %e
+}

``````````

</details>


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