[DAGCombiner] Fold and/or of NaN SETCC

Author: AlexMaclean

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();
 

llvm/test/CodeGen/NVPTX/and-or-setcc.ll

@@ -7,17 +7,15 @@ target triple = "nvptx64-nvidia-cuda"
 define i1 @and_ord(float %a, float %b) {
 ; CHECK-LABEL: and_ord(
 ; CHECK:       {
-; CHECK-NEXT:    .reg .pred %p<4>;
+; 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:    setp.num.f32 %p1, %f1, %f1;
 ; CHECK-NEXT:    ld.param.f32 %f2, [and_ord_param_1];
-; CHECK-NEXT:    setp.num.f32 %p2, %f2, %f2;
-; CHECK-NEXT:    and.pred %p3, %p1, %p2;
-; CHECK-NEXT:    selp.b32 %r1, 1, 0, %p3;
+; 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
@@ -29,17 +27,15 @@ define i1 @and_ord(float %a, float %b) {
 define i1 @or_uno(float %a, float %b) {
 ; CHECK-LABEL: or_uno(
 ; CHECK:       {
-; CHECK-NEXT:    .reg .pred %p<4>;
+; 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:    setp.nan.f32 %p1, %f1, %f1;
 ; CHECK-NEXT:    ld.param.f32 %f2, [or_uno_param_1];
-; CHECK-NEXT:    setp.nan.f32 %p2, %f2, %f2;
-; CHECK-NEXT:    or.pred %p3, %p1, %p2;
-; CHECK-NEXT:    selp.b32 %r1, 1, 0, %p3;
+; 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