llvm/test/CodeGen/NVPTX/machine-cse-predicate-inversion.ll

; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
; RUN: llc -mcpu=sm_100 < %s | FileCheck %s
; RUN: %if ptxas-sm_100 %{ llc < %s -mcpu=sm_100 | %ptxas-verify -arch=sm_100 %}

target triple = "nvptx64-nvidia-cuda"

define i32 @test_int16_eq(i16 %a, i16 %b) {
; CHECK-LABEL: test_int16_eq(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_int16_eq_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_int16_eq_param_0];
; CHECK-NEXT:    setp.eq.b16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB0_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB0_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB0_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB0_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = icmp eq i16 %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = icmp ne i16 %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_int16_ne(i16 %a, i16 %b) {
; CHECK-LABEL: test_int16_ne(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_int16_ne_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_int16_ne_param_0];
; CHECK-NEXT:    setp.ne.b16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB1_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB1_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB1_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB1_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = icmp ne i16 %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = icmp eq i16 %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_int16_slt(i16 %a, i16 %b) {
; CHECK-LABEL: test_int16_slt(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_int16_slt_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_int16_slt_param_0];
; CHECK-NEXT:    setp.lt.s16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB2_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB2_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB2_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB2_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = icmp slt i16 %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = icmp sge i16 %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_int16_sle(i16 %a, i16 %b) {
; CHECK-LABEL: test_int16_sle(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_int16_sle_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_int16_sle_param_0];
; CHECK-NEXT:    setp.le.s16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB3_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB3_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB3_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB3_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = icmp sle i16 %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = icmp sgt i16 %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_int16_sgt(i16 %a, i16 %b) {
; CHECK-LABEL: test_int16_sgt(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_int16_sgt_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_int16_sgt_param_0];
; CHECK-NEXT:    setp.gt.s16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB4_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB4_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB4_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB4_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = icmp sgt i16 %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = icmp sle i16 %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_int16_sge(i16 %a, i16 %b) {
; CHECK-LABEL: test_int16_sge(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_int16_sge_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_int16_sge_param_0];
; CHECK-NEXT:    setp.ge.s16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB5_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB5_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB5_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB5_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = icmp sge i16 %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = icmp slt i16 %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_int16_ult(i16 %a, i16 %b) {
; CHECK-LABEL: test_int16_ult(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_int16_ult_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_int16_ult_param_0];
; CHECK-NEXT:    setp.lt.u16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB6_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB6_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB6_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB6_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = icmp ult i16 %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = icmp uge i16 %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_int16_ule(i16 %a, i16 %b) {
; CHECK-LABEL: test_int16_ule(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_int16_ule_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_int16_ule_param_0];
; CHECK-NEXT:    setp.le.u16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB7_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB7_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB7_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB7_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = icmp ule i16 %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = icmp ugt i16 %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_int16_ugt(i16 %a, i16 %b) {
; CHECK-LABEL: test_int16_ugt(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_int16_ugt_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_int16_ugt_param_0];
; CHECK-NEXT:    setp.gt.u16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB8_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB8_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB8_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB8_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = icmp ugt i16 %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = icmp ule i16 %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_int16_uge(i16 %a, i16 %b) {
; CHECK-LABEL: test_int16_uge(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_int16_uge_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_int16_uge_param_0];
; CHECK-NEXT:    setp.ge.u16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB9_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB9_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB9_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB9_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = icmp uge i16 %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = icmp ult i16 %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_int32_eq(i32 %a, i32 %b) {
; CHECK-LABEL: test_int32_eq(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<4>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b32 %r2, [test_int32_eq_param_1];
; CHECK-NEXT:    ld.param.b32 %r1, [test_int32_eq_param_0];
; CHECK-NEXT:    setp.eq.b32 %p1, %r1, %r2;
; CHECK-NEXT:    mov.b32 %r3, 0;
; CHECK-NEXT:    @%p1 bra $L__BB10_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r3, 1;
; CHECK-NEXT:  $L__BB10_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB10_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r3, 0;
; CHECK-NEXT:  $L__BB10_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r3;
; CHECK-NEXT:    ret;
entry:
  %cmp = icmp eq i32 %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = icmp ne i32 %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_int32_ne(i32 %a, i32 %b) {
; CHECK-LABEL: test_int32_ne(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<4>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b32 %r2, [test_int32_ne_param_1];
; CHECK-NEXT:    ld.param.b32 %r1, [test_int32_ne_param_0];
; CHECK-NEXT:    setp.ne.b32 %p1, %r1, %r2;
; CHECK-NEXT:    mov.b32 %r3, 0;
; CHECK-NEXT:    @%p1 bra $L__BB11_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r3, 1;
; CHECK-NEXT:  $L__BB11_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB11_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r3, 0;
; CHECK-NEXT:  $L__BB11_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r3;
; CHECK-NEXT:    ret;
entry:
  %cmp = icmp ne i32 %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = icmp eq i32 %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_int32_slt(i32 %a, i32 %b) {
; CHECK-LABEL: test_int32_slt(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<4>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b32 %r2, [test_int32_slt_param_1];
; CHECK-NEXT:    ld.param.b32 %r1, [test_int32_slt_param_0];
; CHECK-NEXT:    setp.lt.s32 %p1, %r1, %r2;
; CHECK-NEXT:    mov.b32 %r3, 0;
; CHECK-NEXT:    @%p1 bra $L__BB12_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r3, 1;
; CHECK-NEXT:  $L__BB12_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB12_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r3, 0;
; CHECK-NEXT:  $L__BB12_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r3;
; CHECK-NEXT:    ret;
entry:
  %cmp = icmp slt i32 %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = icmp sge i32 %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_int32_sle(i32 %a, i32 %b) {
; CHECK-LABEL: test_int32_sle(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<4>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b32 %r2, [test_int32_sle_param_1];
; CHECK-NEXT:    ld.param.b32 %r1, [test_int32_sle_param_0];
; CHECK-NEXT:    setp.le.s32 %p1, %r1, %r2;
; CHECK-NEXT:    mov.b32 %r3, 0;
; CHECK-NEXT:    @%p1 bra $L__BB13_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r3, 1;
; CHECK-NEXT:  $L__BB13_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB13_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r3, 0;
; CHECK-NEXT:  $L__BB13_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r3;
; CHECK-NEXT:    ret;
entry:
  %cmp = icmp sle i32 %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = icmp sgt i32 %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_int32_sgt(i32 %a, i32 %b) {
; CHECK-LABEL: test_int32_sgt(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<4>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b32 %r2, [test_int32_sgt_param_1];
; CHECK-NEXT:    ld.param.b32 %r1, [test_int32_sgt_param_0];
; CHECK-NEXT:    setp.gt.s32 %p1, %r1, %r2;
; CHECK-NEXT:    mov.b32 %r3, 0;
; CHECK-NEXT:    @%p1 bra $L__BB14_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r3, 1;
; CHECK-NEXT:  $L__BB14_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB14_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r3, 0;
; CHECK-NEXT:  $L__BB14_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r3;
; CHECK-NEXT:    ret;
entry:
  %cmp = icmp sgt i32 %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = icmp sle i32 %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_int32_sge(i32 %a, i32 %b) {
; CHECK-LABEL: test_int32_sge(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<4>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b32 %r2, [test_int32_sge_param_1];
; CHECK-NEXT:    ld.param.b32 %r1, [test_int32_sge_param_0];
; CHECK-NEXT:    setp.ge.s32 %p1, %r1, %r2;
; CHECK-NEXT:    mov.b32 %r3, 0;
; CHECK-NEXT:    @%p1 bra $L__BB15_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r3, 1;
; CHECK-NEXT:  $L__BB15_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB15_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r3, 0;
; CHECK-NEXT:  $L__BB15_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r3;
; CHECK-NEXT:    ret;
entry:
  %cmp = icmp sge i32 %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = icmp slt i32 %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_int32_ult(i32 %a, i32 %b) {
; CHECK-LABEL: test_int32_ult(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<4>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b32 %r2, [test_int32_ult_param_1];
; CHECK-NEXT:    ld.param.b32 %r1, [test_int32_ult_param_0];
; CHECK-NEXT:    setp.lt.u32 %p1, %r1, %r2;
; CHECK-NEXT:    mov.b32 %r3, 0;
; CHECK-NEXT:    @%p1 bra $L__BB16_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r3, 1;
; CHECK-NEXT:  $L__BB16_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB16_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r3, 0;
; CHECK-NEXT:  $L__BB16_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r3;
; CHECK-NEXT:    ret;
entry:
  %cmp = icmp ult i32 %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = icmp uge i32 %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_int32_ule(i32 %a, i32 %b) {
; CHECK-LABEL: test_int32_ule(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<4>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b32 %r2, [test_int32_ule_param_1];
; CHECK-NEXT:    ld.param.b32 %r1, [test_int32_ule_param_0];
; CHECK-NEXT:    setp.le.u32 %p1, %r1, %r2;
; CHECK-NEXT:    mov.b32 %r3, 0;
; CHECK-NEXT:    @%p1 bra $L__BB17_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r3, 1;
; CHECK-NEXT:  $L__BB17_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB17_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r3, 0;
; CHECK-NEXT:  $L__BB17_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r3;
; CHECK-NEXT:    ret;
entry:
  %cmp = icmp ule i32 %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = icmp ugt i32 %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_int32_ugt(i32 %a, i32 %b) {
; CHECK-LABEL: test_int32_ugt(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<4>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b32 %r2, [test_int32_ugt_param_1];
; CHECK-NEXT:    ld.param.b32 %r1, [test_int32_ugt_param_0];
; CHECK-NEXT:    setp.gt.u32 %p1, %r1, %r2;
; CHECK-NEXT:    mov.b32 %r3, 0;
; CHECK-NEXT:    @%p1 bra $L__BB18_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r3, 1;
; CHECK-NEXT:  $L__BB18_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB18_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r3, 0;
; CHECK-NEXT:  $L__BB18_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r3;
; CHECK-NEXT:    ret;
entry:
  %cmp = icmp ugt i32 %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = icmp ule i32 %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_int32_uge(i32 %a, i32 %b) {
; CHECK-LABEL: test_int32_uge(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<4>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b32 %r2, [test_int32_uge_param_1];
; CHECK-NEXT:    ld.param.b32 %r1, [test_int32_uge_param_0];
; CHECK-NEXT:    setp.ge.u32 %p1, %r1, %r2;
; CHECK-NEXT:    mov.b32 %r3, 0;
; CHECK-NEXT:    @%p1 bra $L__BB19_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r3, 1;
; CHECK-NEXT:  $L__BB19_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB19_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r3, 0;
; CHECK-NEXT:  $L__BB19_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r3;
; CHECK-NEXT:    ret;
entry:
  %cmp = icmp uge i32 %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = icmp ult i32 %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_int64_eq(i64 %a, i64 %b) {
; CHECK-LABEL: test_int64_eq(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-NEXT:    .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b64 %rd2, [test_int64_eq_param_1];
; CHECK-NEXT:    ld.param.b64 %rd1, [test_int64_eq_param_0];
; CHECK-NEXT:    setp.eq.b64 %p1, %rd1, %rd2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB20_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB20_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB20_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB20_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = icmp eq i64 %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = icmp ne i64 %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_int64_ne(i64 %a, i64 %b) {
; CHECK-LABEL: test_int64_ne(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-NEXT:    .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b64 %rd2, [test_int64_ne_param_1];
; CHECK-NEXT:    ld.param.b64 %rd1, [test_int64_ne_param_0];
; CHECK-NEXT:    setp.ne.b64 %p1, %rd1, %rd2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB21_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB21_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB21_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB21_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = icmp ne i64 %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = icmp eq i64 %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_int64_slt(i64 %a, i64 %b) {
; CHECK-LABEL: test_int64_slt(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-NEXT:    .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b64 %rd2, [test_int64_slt_param_1];
; CHECK-NEXT:    ld.param.b64 %rd1, [test_int64_slt_param_0];
; CHECK-NEXT:    setp.lt.s64 %p1, %rd1, %rd2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB22_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB22_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB22_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB22_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = icmp slt i64 %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = icmp sge i64 %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_int64_sle(i64 %a, i64 %b) {
; CHECK-LABEL: test_int64_sle(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-NEXT:    .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b64 %rd2, [test_int64_sle_param_1];
; CHECK-NEXT:    ld.param.b64 %rd1, [test_int64_sle_param_0];
; CHECK-NEXT:    setp.le.s64 %p1, %rd1, %rd2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB23_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB23_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB23_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB23_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = icmp sle i64 %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = icmp sgt i64 %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_int64_sgt(i64 %a, i64 %b) {
; CHECK-LABEL: test_int64_sgt(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-NEXT:    .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b64 %rd2, [test_int64_sgt_param_1];
; CHECK-NEXT:    ld.param.b64 %rd1, [test_int64_sgt_param_0];
; CHECK-NEXT:    setp.gt.s64 %p1, %rd1, %rd2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB24_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB24_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB24_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB24_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = icmp sgt i64 %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = icmp sle i64 %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_int64_sge(i64 %a, i64 %b) {
; CHECK-LABEL: test_int64_sge(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-NEXT:    .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b64 %rd2, [test_int64_sge_param_1];
; CHECK-NEXT:    ld.param.b64 %rd1, [test_int64_sge_param_0];
; CHECK-NEXT:    setp.ge.s64 %p1, %rd1, %rd2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB25_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB25_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB25_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB25_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = icmp sge i64 %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = icmp slt i64 %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_int64_ult(i64 %a, i64 %b) {
; CHECK-LABEL: test_int64_ult(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-NEXT:    .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b64 %rd2, [test_int64_ult_param_1];
; CHECK-NEXT:    ld.param.b64 %rd1, [test_int64_ult_param_0];
; CHECK-NEXT:    setp.lt.u64 %p1, %rd1, %rd2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB26_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB26_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB26_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB26_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = icmp ult i64 %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = icmp uge i64 %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_int64_ule(i64 %a, i64 %b) {
; CHECK-LABEL: test_int64_ule(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-NEXT:    .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b64 %rd2, [test_int64_ule_param_1];
; CHECK-NEXT:    ld.param.b64 %rd1, [test_int64_ule_param_0];
; CHECK-NEXT:    setp.le.u64 %p1, %rd1, %rd2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB27_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB27_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB27_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB27_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = icmp ule i64 %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = icmp ugt i64 %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_int64_ugt(i64 %a, i64 %b) {
; CHECK-LABEL: test_int64_ugt(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-NEXT:    .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b64 %rd2, [test_int64_ugt_param_1];
; CHECK-NEXT:    ld.param.b64 %rd1, [test_int64_ugt_param_0];
; CHECK-NEXT:    setp.gt.u64 %p1, %rd1, %rd2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB28_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB28_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB28_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB28_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = icmp ugt i64 %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = icmp ule i64 %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_int64_uge(i64 %a, i64 %b) {
; CHECK-LABEL: test_int64_uge(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-NEXT:    .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b64 %rd2, [test_int64_uge_param_1];
; CHECK-NEXT:    ld.param.b64 %rd1, [test_int64_uge_param_0];
; CHECK-NEXT:    setp.ge.u64 %p1, %rd1, %rd2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB29_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB29_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB29_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB29_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = icmp uge i64 %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = icmp ult i64 %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_bfloat16_eq(bfloat %a, bfloat %b) {
; CHECK-LABEL: test_bfloat16_eq(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_bfloat16_eq_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_bfloat16_eq_param_0];
; CHECK-NEXT:    setp.eq.bf16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB30_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB30_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB30_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB30_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp oeq bfloat %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp une bfloat %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_bfloat16_ne(bfloat %a, bfloat %b) {
; CHECK-LABEL: test_bfloat16_ne(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_bfloat16_ne_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_bfloat16_ne_param_0];
; CHECK-NEXT:    setp.ne.bf16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB31_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB31_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB31_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB31_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp one bfloat %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp ueq bfloat %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_bfloat16_oeq(bfloat %a, bfloat %b) {
; CHECK-LABEL: test_bfloat16_oeq(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_bfloat16_oeq_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_bfloat16_oeq_param_0];
; CHECK-NEXT:    setp.eq.bf16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB32_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB32_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB32_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB32_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp oeq bfloat %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp une bfloat %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_bfloat16_one(bfloat %a, bfloat %b) {
; CHECK-LABEL: test_bfloat16_one(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_bfloat16_one_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_bfloat16_one_param_0];
; CHECK-NEXT:    setp.ne.bf16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB33_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB33_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB33_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB33_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp one bfloat %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp ueq bfloat %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_bfloat16_ueq(bfloat %a, bfloat %b) {
; CHECK-LABEL: test_bfloat16_ueq(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_bfloat16_ueq_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_bfloat16_ueq_param_0];
; CHECK-NEXT:    setp.equ.bf16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB34_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB34_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB34_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB34_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp ueq bfloat %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp one bfloat %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_bfloat16_une(bfloat %a, bfloat %b) {
; CHECK-LABEL: test_bfloat16_une(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_bfloat16_une_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_bfloat16_une_param_0];
; CHECK-NEXT:    setp.neu.bf16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB35_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB35_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB35_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB35_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp une bfloat %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp oeq bfloat %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_bfloat16_olt(bfloat %a, bfloat %b) {
; CHECK-LABEL: test_bfloat16_olt(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_bfloat16_olt_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_bfloat16_olt_param_0];
; CHECK-NEXT:    setp.lt.bf16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB36_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB36_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB36_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB36_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp olt bfloat %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp uge bfloat %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_bfloat16_ole(bfloat %a, bfloat %b) {
; CHECK-LABEL: test_bfloat16_ole(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_bfloat16_ole_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_bfloat16_ole_param_0];
; CHECK-NEXT:    setp.le.bf16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB37_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB37_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB37_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB37_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp ole bfloat %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp ugt bfloat %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_bfloat16_ogt(bfloat %a, bfloat %b) {
; CHECK-LABEL: test_bfloat16_ogt(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_bfloat16_ogt_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_bfloat16_ogt_param_0];
; CHECK-NEXT:    setp.gt.bf16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB38_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB38_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB38_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB38_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp ogt bfloat %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp ule bfloat %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_bfloat16_oge(bfloat %a, bfloat %b) {
; CHECK-LABEL: test_bfloat16_oge(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_bfloat16_oge_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_bfloat16_oge_param_0];
; CHECK-NEXT:    setp.ge.bf16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB39_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB39_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB39_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB39_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp oge bfloat %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp ult bfloat %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_bfloat16_ult(bfloat %a, bfloat %b) {
; CHECK-LABEL: test_bfloat16_ult(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_bfloat16_ult_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_bfloat16_ult_param_0];
; CHECK-NEXT:    setp.ltu.bf16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB40_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB40_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB40_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB40_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp ult bfloat %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp oge bfloat %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_bfloat16_ule(bfloat %a, bfloat %b) {
; CHECK-LABEL: test_bfloat16_ule(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_bfloat16_ule_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_bfloat16_ule_param_0];
; CHECK-NEXT:    setp.leu.bf16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB41_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB41_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB41_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB41_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp ule bfloat %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp ogt bfloat %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_bfloat16_ugt(bfloat %a, bfloat %b) {
; CHECK-LABEL: test_bfloat16_ugt(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_bfloat16_ugt_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_bfloat16_ugt_param_0];
; CHECK-NEXT:    setp.gtu.bf16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB42_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB42_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB42_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB42_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp ugt bfloat %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp ole bfloat %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_bfloat16_uge(bfloat %a, bfloat %b) {
; CHECK-LABEL: test_bfloat16_uge(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_bfloat16_uge_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_bfloat16_uge_param_0];
; CHECK-NEXT:    setp.geu.bf16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB43_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB43_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB43_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB43_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp uge bfloat %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp olt bfloat %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_bfloat16_ord(bfloat %a, bfloat %b) {
; CHECK-LABEL: test_bfloat16_ord(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_bfloat16_ord_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_bfloat16_ord_param_0];
; CHECK-NEXT:    setp.num.bf16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @!%p1 bra $L__BB44_1;
; CHECK-NEXT:  $L__BB44_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB44_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB44_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
; CHECK-NEXT:  $L__BB44_1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:    bra.uni $L__BB44_2;
entry:
  %cmp = fcmp ord bfloat %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp uno bfloat %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_bfloat16_uno(bfloat %a, bfloat %b) {
; CHECK-LABEL: test_bfloat16_uno(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_bfloat16_uno_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_bfloat16_uno_param_0];
; CHECK-NEXT:    setp.nan.bf16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB45_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB45_2: // %merge1
; CHECK-NEXT:    @%p1 bra $L__BB45_3;
; CHECK-NEXT:  $L__BB45_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
; CHECK-NEXT:  $L__BB45_3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    bra.uni $L__BB45_4;
entry:
  %cmp = fcmp uno bfloat %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp ord bfloat %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float16_eq(half %a, half %b) {
; CHECK-LABEL: test_float16_eq(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_float16_eq_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_float16_eq_param_0];
; CHECK-NEXT:    setp.eq.f16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB46_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB46_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB46_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB46_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp oeq half %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp une half %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float16_ne(half %a, half %b) {
; CHECK-LABEL: test_float16_ne(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_float16_ne_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_float16_ne_param_0];
; CHECK-NEXT:    setp.ne.f16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB47_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB47_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB47_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB47_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp one half %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp ueq half %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float16_oeq(half %a, half %b) {
; CHECK-LABEL: test_float16_oeq(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_float16_oeq_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_float16_oeq_param_0];
; CHECK-NEXT:    setp.eq.f16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB48_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB48_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB48_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB48_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp oeq half %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp une half %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float16_one(half %a, half %b) {
; CHECK-LABEL: test_float16_one(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_float16_one_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_float16_one_param_0];
; CHECK-NEXT:    setp.ne.f16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB49_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB49_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB49_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB49_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp one half %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp ueq half %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float16_ueq(half %a, half %b) {
; CHECK-LABEL: test_float16_ueq(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_float16_ueq_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_float16_ueq_param_0];
; CHECK-NEXT:    setp.equ.f16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB50_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB50_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB50_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB50_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp ueq half %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp one half %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float16_une(half %a, half %b) {
; CHECK-LABEL: test_float16_une(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_float16_une_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_float16_une_param_0];
; CHECK-NEXT:    setp.neu.f16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB51_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB51_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB51_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB51_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp une half %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp oeq half %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float16_olt(half %a, half %b) {
; CHECK-LABEL: test_float16_olt(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_float16_olt_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_float16_olt_param_0];
; CHECK-NEXT:    setp.lt.f16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB52_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB52_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB52_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB52_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp olt half %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp uge half %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float16_ole(half %a, half %b) {
; CHECK-LABEL: test_float16_ole(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_float16_ole_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_float16_ole_param_0];
; CHECK-NEXT:    setp.le.f16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB53_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB53_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB53_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB53_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp ole half %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp ugt half %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float16_ogt(half %a, half %b) {
; CHECK-LABEL: test_float16_ogt(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_float16_ogt_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_float16_ogt_param_0];
; CHECK-NEXT:    setp.gt.f16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB54_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB54_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB54_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB54_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp ogt half %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp ule half %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float16_oge(half %a, half %b) {
; CHECK-LABEL: test_float16_oge(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_float16_oge_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_float16_oge_param_0];
; CHECK-NEXT:    setp.ge.f16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB55_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB55_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB55_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB55_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp oge half %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp ult half %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float16_ult(half %a, half %b) {
; CHECK-LABEL: test_float16_ult(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_float16_ult_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_float16_ult_param_0];
; CHECK-NEXT:    setp.ltu.f16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB56_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB56_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB56_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB56_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp ult half %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp oge half %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float16_ule(half %a, half %b) {
; CHECK-LABEL: test_float16_ule(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_float16_ule_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_float16_ule_param_0];
; CHECK-NEXT:    setp.leu.f16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB57_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB57_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB57_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB57_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp ule half %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp ogt half %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float16_ugt(half %a, half %b) {
; CHECK-LABEL: test_float16_ugt(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_float16_ugt_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_float16_ugt_param_0];
; CHECK-NEXT:    setp.gtu.f16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB58_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB58_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB58_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB58_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp ugt half %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp ole half %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float16_uge(half %a, half %b) {
; CHECK-LABEL: test_float16_uge(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_float16_uge_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_float16_uge_param_0];
; CHECK-NEXT:    setp.geu.f16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB59_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB59_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB59_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB59_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp uge half %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp olt half %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float16_ord(half %a, half %b) {
; CHECK-LABEL: test_float16_ord(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_float16_ord_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_float16_ord_param_0];
; CHECK-NEXT:    setp.num.f16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @!%p1 bra $L__BB60_1;
; CHECK-NEXT:  $L__BB60_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB60_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB60_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
; CHECK-NEXT:  $L__BB60_1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:    bra.uni $L__BB60_2;
entry:
  %cmp = fcmp ord half %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp uno half %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float16_uno(half %a, half %b) {
; CHECK-LABEL: test_float16_uno(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b16 %rs<3>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b16 %rs2, [test_float16_uno_param_1];
; CHECK-NEXT:    ld.param.b16 %rs1, [test_float16_uno_param_0];
; CHECK-NEXT:    setp.nan.f16 %p1, %rs1, %rs2;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB61_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB61_2: // %merge1
; CHECK-NEXT:    @%p1 bra $L__BB61_3;
; CHECK-NEXT:  $L__BB61_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
; CHECK-NEXT:  $L__BB61_3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    bra.uni $L__BB61_4;
entry:
  %cmp = fcmp uno half %a, %b
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp ord half %a, %b
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float32_eq(float %arg) {
; CHECK-LABEL: test_float32_eq(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b32 %r1, [test_float32_eq_param_0];
; CHECK-NEXT:    setp.eq.f32 %p1, %r1, 0f00000000;
; CHECK-NEXT:    mov.b32 %r2, 0;
; CHECK-NEXT:    @%p1 bra $L__BB62_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r2, 1;
; CHECK-NEXT:  $L__BB62_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB62_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r2, 0;
; CHECK-NEXT:  $L__BB62_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r2;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp oeq float %arg, 0.000000e+00
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp une float %arg, 0.000000e+00
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float32_ne(float %arg) {
; CHECK-LABEL: test_float32_ne(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b32 %r1, [test_float32_ne_param_0];
; CHECK-NEXT:    setp.ne.f32 %p1, %r1, 0f00000000;
; CHECK-NEXT:    mov.b32 %r2, 0;
; CHECK-NEXT:    @%p1 bra $L__BB63_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r2, 1;
; CHECK-NEXT:  $L__BB63_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB63_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r2, 0;
; CHECK-NEXT:  $L__BB63_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r2;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp one float %arg, 0.000000e+00
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp ueq float %arg, 0.000000e+00
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float32_oeq(float %arg) {
; CHECK-LABEL: test_float32_oeq(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b32 %r1, [test_float32_oeq_param_0];
; CHECK-NEXT:    setp.eq.f32 %p1, %r1, 0f00000000;
; CHECK-NEXT:    mov.b32 %r2, 0;
; CHECK-NEXT:    @%p1 bra $L__BB64_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r2, 1;
; CHECK-NEXT:  $L__BB64_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB64_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r2, 0;
; CHECK-NEXT:  $L__BB64_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r2;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp oeq float %arg, 0.000000e+00
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp une float %arg, 0.000000e+00
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float32_one(float %arg) {
; CHECK-LABEL: test_float32_one(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b32 %r1, [test_float32_one_param_0];
; CHECK-NEXT:    setp.ne.f32 %p1, %r1, 0f00000000;
; CHECK-NEXT:    mov.b32 %r2, 0;
; CHECK-NEXT:    @%p1 bra $L__BB65_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r2, 1;
; CHECK-NEXT:  $L__BB65_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB65_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r2, 0;
; CHECK-NEXT:  $L__BB65_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r2;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp one float %arg, 0.000000e+00
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp ueq float %arg, 0.000000e+00
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float32_ueq(float %arg) {
; CHECK-LABEL: test_float32_ueq(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b32 %r1, [test_float32_ueq_param_0];
; CHECK-NEXT:    setp.equ.f32 %p1, %r1, 0f00000000;
; CHECK-NEXT:    mov.b32 %r2, 0;
; CHECK-NEXT:    @%p1 bra $L__BB66_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r2, 1;
; CHECK-NEXT:  $L__BB66_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB66_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r2, 0;
; CHECK-NEXT:  $L__BB66_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r2;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp ueq float %arg, 0.000000e+00
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp one float %arg, 0.000000e+00
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float32_une(float %arg) {
; CHECK-LABEL: test_float32_une(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b32 %r1, [test_float32_une_param_0];
; CHECK-NEXT:    setp.neu.f32 %p1, %r1, 0f00000000;
; CHECK-NEXT:    mov.b32 %r2, 0;
; CHECK-NEXT:    @%p1 bra $L__BB67_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r2, 1;
; CHECK-NEXT:  $L__BB67_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB67_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r2, 0;
; CHECK-NEXT:  $L__BB67_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r2;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp une float %arg, 0.000000e+00
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp oeq float %arg, 0.000000e+00
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float32_olt(float %arg) {
; CHECK-LABEL: test_float32_olt(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b32 %r1, [test_float32_olt_param_0];
; CHECK-NEXT:    setp.lt.f32 %p1, %r1, 0f00000000;
; CHECK-NEXT:    mov.b32 %r2, 0;
; CHECK-NEXT:    @%p1 bra $L__BB68_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r2, 1;
; CHECK-NEXT:  $L__BB68_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB68_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r2, 0;
; CHECK-NEXT:  $L__BB68_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r2;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp olt float %arg, 0.000000e+00
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp uge float %arg, 0.000000e+00
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float32_ole(float %arg) {
; CHECK-LABEL: test_float32_ole(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b32 %r1, [test_float32_ole_param_0];
; CHECK-NEXT:    setp.le.f32 %p1, %r1, 0f00000000;
; CHECK-NEXT:    mov.b32 %r2, 0;
; CHECK-NEXT:    @%p1 bra $L__BB69_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r2, 1;
; CHECK-NEXT:  $L__BB69_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB69_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r2, 0;
; CHECK-NEXT:  $L__BB69_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r2;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp ole float %arg, 0.000000e+00
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp ugt float %arg, 0.000000e+00
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float32_ogt(float %arg) {
; CHECK-LABEL: test_float32_ogt(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b32 %r1, [test_float32_ogt_param_0];
; CHECK-NEXT:    setp.gt.f32 %p1, %r1, 0f00000000;
; CHECK-NEXT:    mov.b32 %r2, 0;
; CHECK-NEXT:    @%p1 bra $L__BB70_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r2, 1;
; CHECK-NEXT:  $L__BB70_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB70_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r2, 0;
; CHECK-NEXT:  $L__BB70_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r2;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp ogt float %arg, 0.000000e+00
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp ule float %arg, 0.000000e+00
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float32_oge(float %arg) {
; CHECK-LABEL: test_float32_oge(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b32 %r1, [test_float32_oge_param_0];
; CHECK-NEXT:    setp.ge.f32 %p1, %r1, 0f00000000;
; CHECK-NEXT:    mov.b32 %r2, 0;
; CHECK-NEXT:    @%p1 bra $L__BB71_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r2, 1;
; CHECK-NEXT:  $L__BB71_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB71_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r2, 0;
; CHECK-NEXT:  $L__BB71_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r2;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp oge float %arg, 0.000000e+00
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp ult float %arg, 0.000000e+00
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float32_ult(float %arg) {
; CHECK-LABEL: test_float32_ult(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b32 %r1, [test_float32_ult_param_0];
; CHECK-NEXT:    setp.ltu.f32 %p1, %r1, 0f00000000;
; CHECK-NEXT:    mov.b32 %r2, 0;
; CHECK-NEXT:    @%p1 bra $L__BB72_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r2, 1;
; CHECK-NEXT:  $L__BB72_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB72_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r2, 0;
; CHECK-NEXT:  $L__BB72_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r2;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp ult float %arg, 0.000000e+00
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp oge float %arg, 0.000000e+00
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float32_ule(float %arg) {
; CHECK-LABEL: test_float32_ule(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b32 %r1, [test_float32_ule_param_0];
; CHECK-NEXT:    setp.leu.f32 %p1, %r1, 0f00000000;
; CHECK-NEXT:    mov.b32 %r2, 0;
; CHECK-NEXT:    @%p1 bra $L__BB73_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r2, 1;
; CHECK-NEXT:  $L__BB73_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB73_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r2, 0;
; CHECK-NEXT:  $L__BB73_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r2;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp ule float %arg, 0.000000e+00
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp ogt float %arg, 0.000000e+00
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float32_ugt(float %arg) {
; CHECK-LABEL: test_float32_ugt(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b32 %r1, [test_float32_ugt_param_0];
; CHECK-NEXT:    setp.gtu.f32 %p1, %r1, 0f00000000;
; CHECK-NEXT:    mov.b32 %r2, 0;
; CHECK-NEXT:    @%p1 bra $L__BB74_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r2, 1;
; CHECK-NEXT:  $L__BB74_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB74_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r2, 0;
; CHECK-NEXT:  $L__BB74_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r2;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp ugt float %arg, 0.000000e+00
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp ole float %arg, 0.000000e+00
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float32_uge(float %arg) {
; CHECK-LABEL: test_float32_uge(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b32 %r1, [test_float32_uge_param_0];
; CHECK-NEXT:    setp.geu.f32 %p1, %r1, 0f00000000;
; CHECK-NEXT:    mov.b32 %r2, 0;
; CHECK-NEXT:    @%p1 bra $L__BB75_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r2, 1;
; CHECK-NEXT:  $L__BB75_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB75_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r2, 0;
; CHECK-NEXT:  $L__BB75_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r2;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp uge float %arg, 0.000000e+00
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp olt float %arg, 0.000000e+00
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float32_ord(float %arg) {
; CHECK-LABEL: test_float32_ord(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b32 %r1, [test_float32_ord_param_0];
; CHECK-NEXT:    setp.num.f32 %p1, %r1, %r1;
; CHECK-NEXT:    mov.b32 %r2, 0;
; CHECK-NEXT:    @!%p1 bra $L__BB76_1;
; CHECK-NEXT:  $L__BB76_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB76_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r2, 0;
; CHECK-NEXT:  $L__BB76_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r2;
; CHECK-NEXT:    ret;
; CHECK-NEXT:  $L__BB76_1: // %then
; CHECK-NEXT:    mov.b32 %r2, 1;
; CHECK-NEXT:    bra.uni $L__BB76_2;
entry:
  %cmp = fcmp ord float %arg, 0.000000e+00
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp uno float %arg, 0.000000e+00
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float32_uno(float %arg) {
; CHECK-LABEL: test_float32_uno(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b32 %r1, [test_float32_uno_param_0];
; CHECK-NEXT:    setp.nan.f32 %p1, %r1, %r1;
; CHECK-NEXT:    mov.b32 %r2, 0;
; CHECK-NEXT:    @%p1 bra $L__BB77_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r2, 1;
; CHECK-NEXT:  $L__BB77_2: // %merge1
; CHECK-NEXT:    @%p1 bra $L__BB77_3;
; CHECK-NEXT:  $L__BB77_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r2;
; CHECK-NEXT:    ret;
; CHECK-NEXT:  $L__BB77_3: // %else
; CHECK-NEXT:    mov.b32 %r2, 0;
; CHECK-NEXT:    bra.uni $L__BB77_4;
entry:
  %cmp = fcmp uno float %arg, 0.000000e+00
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp ord float %arg, 0.000000e+00
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float64_eq(double %arg) {
; CHECK-LABEL: test_float64_eq(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-NEXT:    .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b64 %rd1, [test_float64_eq_param_0];
; CHECK-NEXT:    setp.eq.f64 %p1, %rd1, 0d0000000000000000;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB78_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB78_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB78_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB78_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp oeq double %arg, 0.000000e+00
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp une double %arg, 0.000000e+00
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float64_ne(double %arg) {
; CHECK-LABEL: test_float64_ne(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-NEXT:    .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b64 %rd1, [test_float64_ne_param_0];
; CHECK-NEXT:    setp.ne.f64 %p1, %rd1, 0d0000000000000000;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB79_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB79_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB79_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB79_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp one double %arg, 0.000000e+00
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp ueq double %arg, 0.000000e+00
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float64_oeq(double %arg) {
; CHECK-LABEL: test_float64_oeq(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-NEXT:    .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b64 %rd1, [test_float64_oeq_param_0];
; CHECK-NEXT:    setp.eq.f64 %p1, %rd1, 0d0000000000000000;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB80_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB80_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB80_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB80_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp oeq double %arg, 0.000000e+00
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp une double %arg, 0.000000e+00
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float64_one(double %arg) {
; CHECK-LABEL: test_float64_one(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-NEXT:    .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b64 %rd1, [test_float64_one_param_0];
; CHECK-NEXT:    setp.ne.f64 %p1, %rd1, 0d0000000000000000;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB81_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB81_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB81_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB81_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp one double %arg, 0.000000e+00
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp ueq double %arg, 0.000000e+00
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float64_ueq(double %arg) {
; CHECK-LABEL: test_float64_ueq(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-NEXT:    .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b64 %rd1, [test_float64_ueq_param_0];
; CHECK-NEXT:    setp.equ.f64 %p1, %rd1, 0d0000000000000000;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB82_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB82_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB82_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB82_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp ueq double %arg, 0.000000e+00
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp one double %arg, 0.000000e+00
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float64_une(double %arg) {
; CHECK-LABEL: test_float64_une(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-NEXT:    .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b64 %rd1, [test_float64_une_param_0];
; CHECK-NEXT:    setp.neu.f64 %p1, %rd1, 0d0000000000000000;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB83_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB83_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB83_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB83_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp une double %arg, 0.000000e+00
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp oeq double %arg, 0.000000e+00
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float64_olt(double %arg) {
; CHECK-LABEL: test_float64_olt(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-NEXT:    .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b64 %rd1, [test_float64_olt_param_0];
; CHECK-NEXT:    setp.lt.f64 %p1, %rd1, 0d0000000000000000;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB84_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB84_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB84_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB84_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp olt double %arg, 0.000000e+00
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp uge double %arg, 0.000000e+00
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float64_ole(double %arg) {
; CHECK-LABEL: test_float64_ole(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-NEXT:    .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b64 %rd1, [test_float64_ole_param_0];
; CHECK-NEXT:    setp.le.f64 %p1, %rd1, 0d0000000000000000;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB85_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB85_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB85_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB85_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp ole double %arg, 0.000000e+00
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp ugt double %arg, 0.000000e+00
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float64_ogt(double %arg) {
; CHECK-LABEL: test_float64_ogt(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-NEXT:    .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b64 %rd1, [test_float64_ogt_param_0];
; CHECK-NEXT:    setp.gt.f64 %p1, %rd1, 0d0000000000000000;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB86_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB86_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB86_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB86_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp ogt double %arg, 0.000000e+00
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp ule double %arg, 0.000000e+00
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float64_oge(double %arg) {
; CHECK-LABEL: test_float64_oge(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-NEXT:    .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b64 %rd1, [test_float64_oge_param_0];
; CHECK-NEXT:    setp.ge.f64 %p1, %rd1, 0d0000000000000000;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB87_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB87_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB87_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB87_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp oge double %arg, 0.000000e+00
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp ult double %arg, 0.000000e+00
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float64_ult(double %arg) {
; CHECK-LABEL: test_float64_ult(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-NEXT:    .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b64 %rd1, [test_float64_ult_param_0];
; CHECK-NEXT:    setp.ltu.f64 %p1, %rd1, 0d0000000000000000;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB88_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB88_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB88_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB88_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp ult double %arg, 0.000000e+00
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp oge double %arg, 0.000000e+00
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float64_ule(double %arg) {
; CHECK-LABEL: test_float64_ule(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-NEXT:    .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b64 %rd1, [test_float64_ule_param_0];
; CHECK-NEXT:    setp.leu.f64 %p1, %rd1, 0d0000000000000000;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB89_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB89_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB89_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB89_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp ule double %arg, 0.000000e+00
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp ogt double %arg, 0.000000e+00
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float64_ugt(double %arg) {
; CHECK-LABEL: test_float64_ugt(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-NEXT:    .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b64 %rd1, [test_float64_ugt_param_0];
; CHECK-NEXT:    setp.gtu.f64 %p1, %rd1, 0d0000000000000000;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB90_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB90_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB90_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB90_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp ugt double %arg, 0.000000e+00
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp ole double %arg, 0.000000e+00
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float64_uge(double %arg) {
; CHECK-LABEL: test_float64_uge(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-NEXT:    .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b64 %rd1, [test_float64_uge_param_0];
; CHECK-NEXT:    setp.geu.f64 %p1, %rd1, 0d0000000000000000;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB91_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB91_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB91_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB91_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
entry:
  %cmp = fcmp uge double %arg, 0.000000e+00
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp olt double %arg, 0.000000e+00
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float64_ord(double %arg) {
; CHECK-LABEL: test_float64_ord(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-NEXT:    .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b64 %rd1, [test_float64_ord_param_0];
; CHECK-NEXT:    setp.num.f64 %p1, %rd1, %rd1;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @!%p1 bra $L__BB92_1;
; CHECK-NEXT:  $L__BB92_2: // %merge1
; CHECK-NEXT:    @!%p1 bra $L__BB92_4;
; CHECK-NEXT:  // %bb.3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:  $L__BB92_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
; CHECK-NEXT:  $L__BB92_1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:    bra.uni $L__BB92_2;
entry:
  %cmp = fcmp ord double %arg, 0.000000e+00
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp uno double %arg, 0.000000e+00
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}

define i32 @test_float64_uno(double %arg) {
; CHECK-LABEL: test_float64_uno(
; CHECK:       {
; CHECK-NEXT:    .reg .pred %p<2>;
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-NEXT:    .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0: // %entry
; CHECK-NEXT:    ld.param.b64 %rd1, [test_float64_uno_param_0];
; CHECK-NEXT:    setp.nan.f64 %p1, %rd1, %rd1;
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    @%p1 bra $L__BB93_2;
; CHECK-NEXT:  // %bb.1: // %then
; CHECK-NEXT:    mov.b32 %r1, 1;
; CHECK-NEXT:  $L__BB93_2: // %merge1
; CHECK-NEXT:    @%p1 bra $L__BB93_3;
; CHECK-NEXT:  $L__BB93_4: // %merge2
; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
; CHECK-NEXT:    ret;
; CHECK-NEXT:  $L__BB93_3: // %else
; CHECK-NEXT:    mov.b32 %r1, 0;
; CHECK-NEXT:    bra.uni $L__BB93_4;
entry:
  %cmp = fcmp uno double %arg, 0.000000e+00
  br i1 %cmp, label %merge1, label %then

then:
  br label %merge1

merge1:
  %phi1 = phi i32 [ 1, %then ], [ 0, %entry ]
  %cmp2 = fcmp ord double %arg, 0.000000e+00
  br i1 %cmp2, label %merge2, label %else

else:
  br label %merge2

merge2:
  %phi2 = phi i32 [ %phi1, %merge1 ], [ 0, %else ]
  ret i32 %phi2
}