| ; 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.u32 %r1, [icmp_i1_eq_param_0]; |
| ; CHECK-NEXT: setp.gt.s32 %p1, %r1, 1; |
| ; CHECK-NEXT: ld.param.u32 %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.u32 %r1, [icmp_i1_ne_param_0]; |
| ; CHECK-NEXT: setp.gt.s32 %p1, %r1, 1; |
| ; CHECK-NEXT: ld.param.u32 %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.u32 %r1, [icmp_i1_sgt_param_0]; |
| ; CHECK-NEXT: setp.gt.s32 %p1, %r1, 1; |
| ; CHECK-NEXT: ld.param.u32 %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.u32 %r1, [icmp_i1_slt_param_0]; |
| ; CHECK-NEXT: setp.lt.s32 %p1, %r1, 2; |
| ; CHECK-NEXT: ld.param.u32 %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.u32 %r1, [icmp_i1_sge_param_0]; |
| ; CHECK-NEXT: setp.gt.s32 %p1, %r1, 1; |
| ; CHECK-NEXT: ld.param.u32 %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.u32 %r1, [icmp_i1_sle_param_0]; |
| ; CHECK-NEXT: setp.lt.s32 %p1, %r1, 2; |
| ; CHECK-NEXT: ld.param.u32 %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.u32 %r1, [icmp_i1_uge_param_0]; |
| ; CHECK-NEXT: setp.lt.s32 %p1, %r1, 2; |
| ; CHECK-NEXT: ld.param.u32 %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.u32 %r1, [icmp_i1_ugt_param_0]; |
| ; CHECK-NEXT: setp.lt.s32 %p1, %r1, 2; |
| ; CHECK-NEXT: ld.param.u32 %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.u32 %r1, [icmp_i1_ule_param_0]; |
| ; CHECK-NEXT: setp.gt.s32 %p1, %r1, 1; |
| ; CHECK-NEXT: ld.param.u32 %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.u32 %r1, [icmp_i1_ult_param_0]; |
| ; CHECK-NEXT: setp.gt.s32 %p1, %r1, 1; |
| ; CHECK-NEXT: ld.param.u32 %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 |
| } |
| |