|  | ; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5 | 
|  | ; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 | FileCheck %s | 
|  | ; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 | %ptxas-verify %} | 
|  |  | 
|  | target triple = "nvptx-nvidia-cuda" | 
|  |  | 
|  | define i32 @icmp_i1_eq(i32 %a, i32 %b) { | 
|  | ; CHECK-LABEL: icmp_i1_eq( | 
|  | ; CHECK:       { | 
|  | ; CHECK-NEXT:    .reg .pred %p<4>; | 
|  | ; CHECK-NEXT:    .reg .b32 %r<5>; | 
|  | ; CHECK-EMPTY: | 
|  | ; CHECK-NEXT:  // %bb.0: | 
|  | ; CHECK-NEXT:    ld.param.b32 %r1, [icmp_i1_eq_param_0]; | 
|  | ; CHECK-NEXT:    setp.gt.s32 %p1, %r1, 1; | 
|  | ; CHECK-NEXT:    ld.param.b32 %r2, [icmp_i1_eq_param_1]; | 
|  | ; CHECK-NEXT:    setp.gt.s32 %p2, %r2, 1; | 
|  | ; CHECK-NEXT:    xor.pred %p3, %p1, %p2; | 
|  | ; CHECK-NEXT:    @%p3 bra $L__BB0_2; | 
|  | ; CHECK-NEXT:  // %bb.1: // %bb1 | 
|  | ; CHECK-NEXT:    mov.b32 %r4, 1; | 
|  | ; CHECK-NEXT:    st.param.b32 [func_retval0], %r4; | 
|  | ; CHECK-NEXT:    ret; | 
|  | ; CHECK-NEXT:  $L__BB0_2: // %bb2 | 
|  | ; CHECK-NEXT:    mov.b32 %r3, 127; | 
|  | ; CHECK-NEXT:    st.param.b32 [func_retval0], %r3; | 
|  | ; CHECK-NEXT:    ret; | 
|  | %p1 = icmp sgt i32 %a, 1 | 
|  | %p2 = icmp sgt i32 %b, 1 | 
|  | %c = icmp eq i1 %p1, %p2 | 
|  | br i1 %c, label %bb1, label %bb2 | 
|  | bb1: | 
|  | ret i32 1 | 
|  | bb2: | 
|  | ret i32 127 | 
|  | } | 
|  |  | 
|  | define i32 @icmp_i1_ne(i32 %a, i32 %b) { | 
|  | ; CHECK-LABEL: icmp_i1_ne( | 
|  | ; CHECK:       { | 
|  | ; CHECK-NEXT:    .reg .pred %p<5>; | 
|  | ; CHECK-NEXT:    .reg .b32 %r<5>; | 
|  | ; CHECK-EMPTY: | 
|  | ; CHECK-NEXT:  // %bb.0: | 
|  | ; CHECK-NEXT:    ld.param.b32 %r1, [icmp_i1_ne_param_0]; | 
|  | ; CHECK-NEXT:    setp.gt.s32 %p1, %r1, 1; | 
|  | ; CHECK-NEXT:    ld.param.b32 %r2, [icmp_i1_ne_param_1]; | 
|  | ; CHECK-NEXT:    setp.gt.s32 %p2, %r2, 1; | 
|  | ; CHECK-NEXT:    xor.pred %p3, %p1, %p2; | 
|  | ; CHECK-NEXT:    not.pred %p4, %p3; | 
|  | ; CHECK-NEXT:    @%p4 bra $L__BB1_2; | 
|  | ; CHECK-NEXT:  // %bb.1: // %bb1 | 
|  | ; CHECK-NEXT:    mov.b32 %r4, 1; | 
|  | ; CHECK-NEXT:    st.param.b32 [func_retval0], %r4; | 
|  | ; CHECK-NEXT:    ret; | 
|  | ; CHECK-NEXT:  $L__BB1_2: // %bb2 | 
|  | ; CHECK-NEXT:    mov.b32 %r3, 127; | 
|  | ; CHECK-NEXT:    st.param.b32 [func_retval0], %r3; | 
|  | ; CHECK-NEXT:    ret; | 
|  | %p1 = icmp sgt i32 %a, 1 | 
|  | %p2 = icmp sgt i32 %b, 1 | 
|  | %c = icmp ne i1 %p1, %p2 | 
|  | br i1 %c, label %bb1, label %bb2 | 
|  | bb1: | 
|  | ret i32 1 | 
|  | bb2: | 
|  | ret i32 127 | 
|  | } | 
|  |  | 
|  | define i32 @icmp_i1_sgt(i32 %a, i32 %b) { | 
|  | ; CHECK-LABEL: icmp_i1_sgt( | 
|  | ; CHECK:       { | 
|  | ; CHECK-NEXT:    .reg .pred %p<4>; | 
|  | ; CHECK-NEXT:    .reg .b32 %r<5>; | 
|  | ; CHECK-EMPTY: | 
|  | ; CHECK-NEXT:  // %bb.0: | 
|  | ; CHECK-NEXT:    ld.param.b32 %r1, [icmp_i1_sgt_param_0]; | 
|  | ; CHECK-NEXT:    setp.gt.s32 %p1, %r1, 1; | 
|  | ; CHECK-NEXT:    ld.param.b32 %r2, [icmp_i1_sgt_param_1]; | 
|  | ; CHECK-NEXT:    setp.lt.s32 %p2, %r2, 2; | 
|  | ; CHECK-NEXT:    or.pred %p3, %p1, %p2; | 
|  | ; CHECK-NEXT:    @%p3 bra $L__BB2_2; | 
|  | ; CHECK-NEXT:  // %bb.1: // %bb1 | 
|  | ; CHECK-NEXT:    mov.b32 %r4, 1; | 
|  | ; CHECK-NEXT:    st.param.b32 [func_retval0], %r4; | 
|  | ; CHECK-NEXT:    ret; | 
|  | ; CHECK-NEXT:  $L__BB2_2: // %bb2 | 
|  | ; CHECK-NEXT:    mov.b32 %r3, 127; | 
|  | ; CHECK-NEXT:    st.param.b32 [func_retval0], %r3; | 
|  | ; CHECK-NEXT:    ret; | 
|  | %p1 = icmp sgt i32 %a, 1 | 
|  | %p2 = icmp sgt i32 %b, 1 | 
|  | %c = icmp sgt i1 %p1, %p2 | 
|  | br i1 %c, label %bb1, label %bb2 | 
|  | bb1: | 
|  | ret i32 1 | 
|  | bb2: | 
|  | ret i32 127 | 
|  | } | 
|  |  | 
|  | define i32 @icmp_i1_slt(i32 %a, i32 %b) { | 
|  | ; CHECK-LABEL: icmp_i1_slt( | 
|  | ; CHECK:       { | 
|  | ; CHECK-NEXT:    .reg .pred %p<4>; | 
|  | ; CHECK-NEXT:    .reg .b32 %r<5>; | 
|  | ; CHECK-EMPTY: | 
|  | ; CHECK-NEXT:  // %bb.0: | 
|  | ; CHECK-NEXT:    ld.param.b32 %r1, [icmp_i1_slt_param_0]; | 
|  | ; CHECK-NEXT:    setp.lt.s32 %p1, %r1, 2; | 
|  | ; CHECK-NEXT:    ld.param.b32 %r2, [icmp_i1_slt_param_1]; | 
|  | ; CHECK-NEXT:    setp.gt.s32 %p2, %r2, 1; | 
|  | ; CHECK-NEXT:    or.pred %p3, %p2, %p1; | 
|  | ; CHECK-NEXT:    @%p3 bra $L__BB3_2; | 
|  | ; CHECK-NEXT:  // %bb.1: // %bb1 | 
|  | ; CHECK-NEXT:    mov.b32 %r4, 1; | 
|  | ; CHECK-NEXT:    st.param.b32 [func_retval0], %r4; | 
|  | ; CHECK-NEXT:    ret; | 
|  | ; CHECK-NEXT:  $L__BB3_2: // %bb2 | 
|  | ; CHECK-NEXT:    mov.b32 %r3, 127; | 
|  | ; CHECK-NEXT:    st.param.b32 [func_retval0], %r3; | 
|  | ; CHECK-NEXT:    ret; | 
|  | %p1 = icmp sgt i32 %a, 1 | 
|  | %p2 = icmp sgt i32 %b, 1 | 
|  | %c = icmp slt i1 %p1, %p2 | 
|  | br i1 %c, label %bb1, label %bb2 | 
|  | bb1: | 
|  | ret i32 1 | 
|  | bb2: | 
|  | ret i32 127 | 
|  | } | 
|  |  | 
|  | define i32 @icmp_i1_sge(i32 %a, i32 %b) { | 
|  | ; CHECK-LABEL: icmp_i1_sge( | 
|  | ; CHECK:       { | 
|  | ; CHECK-NEXT:    .reg .pred %p<4>; | 
|  | ; CHECK-NEXT:    .reg .b32 %r<5>; | 
|  | ; CHECK-EMPTY: | 
|  | ; CHECK-NEXT:  // %bb.0: | 
|  | ; CHECK-NEXT:    ld.param.b32 %r1, [icmp_i1_sge_param_0]; | 
|  | ; CHECK-NEXT:    setp.gt.s32 %p1, %r1, 1; | 
|  | ; CHECK-NEXT:    ld.param.b32 %r2, [icmp_i1_sge_param_1]; | 
|  | ; CHECK-NEXT:    setp.lt.s32 %p2, %r2, 2; | 
|  | ; CHECK-NEXT:    and.pred %p3, %p1, %p2; | 
|  | ; CHECK-NEXT:    @%p3 bra $L__BB4_2; | 
|  | ; CHECK-NEXT:  // %bb.1: // %bb1 | 
|  | ; CHECK-NEXT:    mov.b32 %r4, 1; | 
|  | ; CHECK-NEXT:    st.param.b32 [func_retval0], %r4; | 
|  | ; CHECK-NEXT:    ret; | 
|  | ; CHECK-NEXT:  $L__BB4_2: // %bb2 | 
|  | ; CHECK-NEXT:    mov.b32 %r3, 127; | 
|  | ; CHECK-NEXT:    st.param.b32 [func_retval0], %r3; | 
|  | ; CHECK-NEXT:    ret; | 
|  | %p1 = icmp sgt i32 %a, 1 | 
|  | %p2 = icmp sgt i32 %b, 1 | 
|  | %c = icmp sge i1 %p1, %p2 | 
|  | br i1 %c, label %bb1, label %bb2 | 
|  | bb1: | 
|  | ret i32 1 | 
|  | bb2: | 
|  | ret i32 127 | 
|  | } | 
|  |  | 
|  | define i32 @icmp_i1_sle(i32 %a, i32 %b) { | 
|  | ; CHECK-LABEL: icmp_i1_sle( | 
|  | ; CHECK:       { | 
|  | ; CHECK-NEXT:    .reg .pred %p<4>; | 
|  | ; CHECK-NEXT:    .reg .b32 %r<5>; | 
|  | ; CHECK-EMPTY: | 
|  | ; CHECK-NEXT:  // %bb.0: | 
|  | ; CHECK-NEXT:    ld.param.b32 %r1, [icmp_i1_sle_param_0]; | 
|  | ; CHECK-NEXT:    setp.lt.s32 %p1, %r1, 2; | 
|  | ; CHECK-NEXT:    ld.param.b32 %r2, [icmp_i1_sle_param_1]; | 
|  | ; CHECK-NEXT:    setp.gt.s32 %p2, %r2, 1; | 
|  | ; CHECK-NEXT:    and.pred %p3, %p2, %p1; | 
|  | ; CHECK-NEXT:    @%p3 bra $L__BB5_2; | 
|  | ; CHECK-NEXT:  // %bb.1: // %bb1 | 
|  | ; CHECK-NEXT:    mov.b32 %r4, 1; | 
|  | ; CHECK-NEXT:    st.param.b32 [func_retval0], %r4; | 
|  | ; CHECK-NEXT:    ret; | 
|  | ; CHECK-NEXT:  $L__BB5_2: // %bb2 | 
|  | ; CHECK-NEXT:    mov.b32 %r3, 127; | 
|  | ; CHECK-NEXT:    st.param.b32 [func_retval0], %r3; | 
|  | ; CHECK-NEXT:    ret; | 
|  | %p1 = icmp sgt i32 %a, 1 | 
|  | %p2 = icmp sgt i32 %b, 1 | 
|  | %c = icmp sle i1 %p1, %p2 | 
|  | br i1 %c, label %bb1, label %bb2 | 
|  | bb1: | 
|  | ret i32 1 | 
|  | bb2: | 
|  | ret i32 127 | 
|  | } | 
|  |  | 
|  | define i32 @icmp_i1_uge(i32 %a, i32 %b) { | 
|  | ; CHECK-LABEL: icmp_i1_uge( | 
|  | ; CHECK:       { | 
|  | ; CHECK-NEXT:    .reg .pred %p<4>; | 
|  | ; CHECK-NEXT:    .reg .b32 %r<5>; | 
|  | ; CHECK-EMPTY: | 
|  | ; CHECK-NEXT:  // %bb.0: | 
|  | ; CHECK-NEXT:    ld.param.b32 %r1, [icmp_i1_uge_param_0]; | 
|  | ; CHECK-NEXT:    setp.lt.s32 %p1, %r1, 2; | 
|  | ; CHECK-NEXT:    ld.param.b32 %r2, [icmp_i1_uge_param_1]; | 
|  | ; CHECK-NEXT:    setp.gt.s32 %p2, %r2, 1; | 
|  | ; CHECK-NEXT:    and.pred %p3, %p2, %p1; | 
|  | ; CHECK-NEXT:    @%p3 bra $L__BB6_2; | 
|  | ; CHECK-NEXT:  // %bb.1: // %bb1 | 
|  | ; CHECK-NEXT:    mov.b32 %r4, 1; | 
|  | ; CHECK-NEXT:    st.param.b32 [func_retval0], %r4; | 
|  | ; CHECK-NEXT:    ret; | 
|  | ; CHECK-NEXT:  $L__BB6_2: // %bb2 | 
|  | ; CHECK-NEXT:    mov.b32 %r3, 127; | 
|  | ; CHECK-NEXT:    st.param.b32 [func_retval0], %r3; | 
|  | ; CHECK-NEXT:    ret; | 
|  | %p1 = icmp sgt i32 %a, 1 | 
|  | %p2 = icmp sgt i32 %b, 1 | 
|  | %c = icmp uge i1 %p1, %p2 | 
|  | br i1 %c, label %bb1, label %bb2 | 
|  | bb1: | 
|  | ret i32 1 | 
|  | bb2: | 
|  | ret i32 127 | 
|  | } | 
|  |  | 
|  | define i32 @icmp_i1_ugt(i32 %a, i32 %b) { | 
|  | ; CHECK-LABEL: icmp_i1_ugt( | 
|  | ; CHECK:       { | 
|  | ; CHECK-NEXT:    .reg .pred %p<4>; | 
|  | ; CHECK-NEXT:    .reg .b32 %r<5>; | 
|  | ; CHECK-EMPTY: | 
|  | ; CHECK-NEXT:  // %bb.0: | 
|  | ; CHECK-NEXT:    ld.param.b32 %r1, [icmp_i1_ugt_param_0]; | 
|  | ; CHECK-NEXT:    setp.lt.s32 %p1, %r1, 2; | 
|  | ; CHECK-NEXT:    ld.param.b32 %r2, [icmp_i1_ugt_param_1]; | 
|  | ; CHECK-NEXT:    setp.gt.s32 %p2, %r2, 1; | 
|  | ; CHECK-NEXT:    or.pred %p3, %p2, %p1; | 
|  | ; CHECK-NEXT:    @%p3 bra $L__BB7_2; | 
|  | ; CHECK-NEXT:  // %bb.1: // %bb1 | 
|  | ; CHECK-NEXT:    mov.b32 %r4, 1; | 
|  | ; CHECK-NEXT:    st.param.b32 [func_retval0], %r4; | 
|  | ; CHECK-NEXT:    ret; | 
|  | ; CHECK-NEXT:  $L__BB7_2: // %bb2 | 
|  | ; CHECK-NEXT:    mov.b32 %r3, 127; | 
|  | ; CHECK-NEXT:    st.param.b32 [func_retval0], %r3; | 
|  | ; CHECK-NEXT:    ret; | 
|  | %p1 = icmp sgt i32 %a, 1 | 
|  | %p2 = icmp sgt i32 %b, 1 | 
|  | %c = icmp ugt i1 %p1, %p2 | 
|  | br i1 %c, label %bb1, label %bb2 | 
|  | bb1: | 
|  | ret i32 1 | 
|  | bb2: | 
|  | ret i32 127 | 
|  | } | 
|  |  | 
|  | define i32 @icmp_i1_ule(i32 %a, i32 %b) { | 
|  | ; CHECK-LABEL: icmp_i1_ule( | 
|  | ; CHECK:       { | 
|  | ; CHECK-NEXT:    .reg .pred %p<4>; | 
|  | ; CHECK-NEXT:    .reg .b32 %r<5>; | 
|  | ; CHECK-EMPTY: | 
|  | ; CHECK-NEXT:  // %bb.0: | 
|  | ; CHECK-NEXT:    ld.param.b32 %r1, [icmp_i1_ule_param_0]; | 
|  | ; CHECK-NEXT:    setp.gt.s32 %p1, %r1, 1; | 
|  | ; CHECK-NEXT:    ld.param.b32 %r2, [icmp_i1_ule_param_1]; | 
|  | ; CHECK-NEXT:    setp.lt.s32 %p2, %r2, 2; | 
|  | ; CHECK-NEXT:    and.pred %p3, %p1, %p2; | 
|  | ; CHECK-NEXT:    @%p3 bra $L__BB8_2; | 
|  | ; CHECK-NEXT:  // %bb.1: // %bb1 | 
|  | ; CHECK-NEXT:    mov.b32 %r4, 1; | 
|  | ; CHECK-NEXT:    st.param.b32 [func_retval0], %r4; | 
|  | ; CHECK-NEXT:    ret; | 
|  | ; CHECK-NEXT:  $L__BB8_2: // %bb2 | 
|  | ; CHECK-NEXT:    mov.b32 %r3, 127; | 
|  | ; CHECK-NEXT:    st.param.b32 [func_retval0], %r3; | 
|  | ; CHECK-NEXT:    ret; | 
|  | %p1 = icmp sgt i32 %a, 1 | 
|  | %p2 = icmp sgt i32 %b, 1 | 
|  | %c = icmp ule i1 %p1, %p2 | 
|  | br i1 %c, label %bb1, label %bb2 | 
|  | bb1: | 
|  | ret i32 1 | 
|  | bb2: | 
|  | ret i32 127 | 
|  | } | 
|  |  | 
|  | define i32 @icmp_i1_ult(i32 %a, i32 %b) { | 
|  | ; CHECK-LABEL: icmp_i1_ult( | 
|  | ; CHECK:       { | 
|  | ; CHECK-NEXT:    .reg .pred %p<4>; | 
|  | ; CHECK-NEXT:    .reg .b32 %r<5>; | 
|  | ; CHECK-EMPTY: | 
|  | ; CHECK-NEXT:  // %bb.0: | 
|  | ; CHECK-NEXT:    ld.param.b32 %r1, [icmp_i1_ult_param_0]; | 
|  | ; CHECK-NEXT:    setp.gt.s32 %p1, %r1, 1; | 
|  | ; CHECK-NEXT:    ld.param.b32 %r2, [icmp_i1_ult_param_1]; | 
|  | ; CHECK-NEXT:    setp.lt.s32 %p2, %r2, 2; | 
|  | ; CHECK-NEXT:    or.pred %p3, %p1, %p2; | 
|  | ; CHECK-NEXT:    @%p3 bra $L__BB9_2; | 
|  | ; CHECK-NEXT:  // %bb.1: // %bb1 | 
|  | ; CHECK-NEXT:    mov.b32 %r4, 1; | 
|  | ; CHECK-NEXT:    st.param.b32 [func_retval0], %r4; | 
|  | ; CHECK-NEXT:    ret; | 
|  | ; CHECK-NEXT:  $L__BB9_2: // %bb2 | 
|  | ; CHECK-NEXT:    mov.b32 %r3, 127; | 
|  | ; CHECK-NEXT:    st.param.b32 [func_retval0], %r3; | 
|  | ; CHECK-NEXT:    ret; | 
|  | %p1 = icmp sgt i32 %a, 1 | 
|  | %p2 = icmp sgt i32 %b, 1 | 
|  | %c = icmp ult i1 %p1, %p2 | 
|  | br i1 %c, label %bb1, label %bb2 | 
|  | bb1: | 
|  | ret i32 1 | 
|  | bb2: | 
|  | ret i32 127 | 
|  | } | 
|  |  |