; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
; RUN: llc < %s | FileCheck %s
; RUN: %if ptxas %{ llc < %s | %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<4>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b32 %r1, [and_ord_param_0];
; CHECK-NEXT:    ld.param.b32 %r2, [and_ord_param_1];
; CHECK-NEXT:    setp.num.f32 %p1, %r1, %r2;
; CHECK-NEXT:    selp.b32 %r3, -1, 0, %p1;
; CHECK-NEXT:    st.param.b32 [func_retval0], %r3;
; 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<4>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b32 %r1, [or_uno_param_0];
; CHECK-NEXT:    ld.param.b32 %r2, [or_uno_param_1];
; CHECK-NEXT:    setp.nan.f32 %p1, %r1, %r2;
; CHECK-NEXT:    selp.b32 %r3, -1, 0, %p1;
; CHECK-NEXT:    st.param.b32 [func_retval0], %r3;
; 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
}