1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5 2; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 | FileCheck %s 3; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 | %ptxas-verify %} 4 5target triple = "nvptx-nvidia-cuda" 6 7define i32 @icmp_i1_eq(i32 %a, i32 %b) { 8; CHECK-LABEL: icmp_i1_eq( 9; CHECK: { 10; CHECK-NEXT: .reg .pred %p<4>; 11; CHECK-NEXT: .reg .b32 %r<5>; 12; CHECK-EMPTY: 13; CHECK-NEXT: // %bb.0: 14; CHECK-NEXT: ld.param.u32 %r1, [icmp_i1_eq_param_0]; 15; CHECK-NEXT: setp.gt.s32 %p1, %r1, 1; 16; CHECK-NEXT: ld.param.u32 %r2, [icmp_i1_eq_param_1]; 17; CHECK-NEXT: setp.gt.s32 %p2, %r2, 1; 18; CHECK-NEXT: xor.pred %p3, %p1, %p2; 19; CHECK-NEXT: @%p3 bra $L__BB0_2; 20; CHECK-NEXT: // %bb.1: // %bb1 21; CHECK-NEXT: mov.b32 %r4, 1; 22; CHECK-NEXT: st.param.b32 [func_retval0], %r4; 23; CHECK-NEXT: ret; 24; CHECK-NEXT: $L__BB0_2: // %bb2 25; CHECK-NEXT: mov.b32 %r3, 127; 26; CHECK-NEXT: st.param.b32 [func_retval0], %r3; 27; CHECK-NEXT: ret; 28 %p1 = icmp sgt i32 %a, 1 29 %p2 = icmp sgt i32 %b, 1 30 %c = icmp eq i1 %p1, %p2 31 br i1 %c, label %bb1, label %bb2 32bb1: 33 ret i32 1 34bb2: 35 ret i32 127 36} 37 38define i32 @icmp_i1_ne(i32 %a, i32 %b) { 39; CHECK-LABEL: icmp_i1_ne( 40; CHECK: { 41; CHECK-NEXT: .reg .pred %p<5>; 42; CHECK-NEXT: .reg .b32 %r<5>; 43; CHECK-EMPTY: 44; CHECK-NEXT: // %bb.0: 45; CHECK-NEXT: ld.param.u32 %r1, [icmp_i1_ne_param_0]; 46; CHECK-NEXT: setp.gt.s32 %p1, %r1, 1; 47; CHECK-NEXT: ld.param.u32 %r2, [icmp_i1_ne_param_1]; 48; CHECK-NEXT: setp.gt.s32 %p2, %r2, 1; 49; CHECK-NEXT: xor.pred %p3, %p1, %p2; 50; CHECK-NEXT: not.pred %p4, %p3; 51; CHECK-NEXT: @%p4 bra $L__BB1_2; 52; CHECK-NEXT: // %bb.1: // %bb1 53; CHECK-NEXT: mov.b32 %r4, 1; 54; CHECK-NEXT: st.param.b32 [func_retval0], %r4; 55; CHECK-NEXT: ret; 56; CHECK-NEXT: $L__BB1_2: // %bb2 57; CHECK-NEXT: mov.b32 %r3, 127; 58; CHECK-NEXT: st.param.b32 [func_retval0], %r3; 59; CHECK-NEXT: ret; 60 %p1 = icmp sgt i32 %a, 1 61 %p2 = icmp sgt i32 %b, 1 62 %c = icmp ne i1 %p1, %p2 63 br i1 %c, label %bb1, label %bb2 64bb1: 65 ret i32 1 66bb2: 67 ret i32 127 68} 69 70define i32 @icmp_i1_sgt(i32 %a, i32 %b) { 71; CHECK-LABEL: icmp_i1_sgt( 72; CHECK: { 73; CHECK-NEXT: .reg .pred %p<4>; 74; CHECK-NEXT: .reg .b32 %r<5>; 75; CHECK-EMPTY: 76; CHECK-NEXT: // %bb.0: 77; CHECK-NEXT: ld.param.u32 %r1, [icmp_i1_sgt_param_0]; 78; CHECK-NEXT: setp.gt.s32 %p1, %r1, 1; 79; CHECK-NEXT: ld.param.u32 %r2, [icmp_i1_sgt_param_1]; 80; CHECK-NEXT: setp.lt.s32 %p2, %r2, 2; 81; CHECK-NEXT: or.pred %p3, %p1, %p2; 82; CHECK-NEXT: @%p3 bra $L__BB2_2; 83; CHECK-NEXT: // %bb.1: // %bb1 84; CHECK-NEXT: mov.b32 %r4, 1; 85; CHECK-NEXT: st.param.b32 [func_retval0], %r4; 86; CHECK-NEXT: ret; 87; CHECK-NEXT: $L__BB2_2: // %bb2 88; CHECK-NEXT: mov.b32 %r3, 127; 89; CHECK-NEXT: st.param.b32 [func_retval0], %r3; 90; CHECK-NEXT: ret; 91 %p1 = icmp sgt i32 %a, 1 92 %p2 = icmp sgt i32 %b, 1 93 %c = icmp sgt i1 %p1, %p2 94 br i1 %c, label %bb1, label %bb2 95bb1: 96 ret i32 1 97bb2: 98 ret i32 127 99} 100 101define i32 @icmp_i1_slt(i32 %a, i32 %b) { 102; CHECK-LABEL: icmp_i1_slt( 103; CHECK: { 104; CHECK-NEXT: .reg .pred %p<4>; 105; CHECK-NEXT: .reg .b32 %r<5>; 106; CHECK-EMPTY: 107; CHECK-NEXT: // %bb.0: 108; CHECK-NEXT: ld.param.u32 %r1, [icmp_i1_slt_param_0]; 109; CHECK-NEXT: setp.lt.s32 %p1, %r1, 2; 110; CHECK-NEXT: ld.param.u32 %r2, [icmp_i1_slt_param_1]; 111; CHECK-NEXT: setp.gt.s32 %p2, %r2, 1; 112; CHECK-NEXT: or.pred %p3, %p2, %p1; 113; CHECK-NEXT: @%p3 bra $L__BB3_2; 114; CHECK-NEXT: // %bb.1: // %bb1 115; CHECK-NEXT: mov.b32 %r4, 1; 116; CHECK-NEXT: st.param.b32 [func_retval0], %r4; 117; CHECK-NEXT: ret; 118; CHECK-NEXT: $L__BB3_2: // %bb2 119; CHECK-NEXT: mov.b32 %r3, 127; 120; CHECK-NEXT: st.param.b32 [func_retval0], %r3; 121; CHECK-NEXT: ret; 122 %p1 = icmp sgt i32 %a, 1 123 %p2 = icmp sgt i32 %b, 1 124 %c = icmp slt i1 %p1, %p2 125 br i1 %c, label %bb1, label %bb2 126bb1: 127 ret i32 1 128bb2: 129 ret i32 127 130} 131 132define i32 @icmp_i1_sge(i32 %a, i32 %b) { 133; CHECK-LABEL: icmp_i1_sge( 134; CHECK: { 135; CHECK-NEXT: .reg .pred %p<4>; 136; CHECK-NEXT: .reg .b32 %r<5>; 137; CHECK-EMPTY: 138; CHECK-NEXT: // %bb.0: 139; CHECK-NEXT: ld.param.u32 %r1, [icmp_i1_sge_param_0]; 140; CHECK-NEXT: setp.gt.s32 %p1, %r1, 1; 141; CHECK-NEXT: ld.param.u32 %r2, [icmp_i1_sge_param_1]; 142; CHECK-NEXT: setp.lt.s32 %p2, %r2, 2; 143; CHECK-NEXT: and.pred %p3, %p1, %p2; 144; CHECK-NEXT: @%p3 bra $L__BB4_2; 145; CHECK-NEXT: // %bb.1: // %bb1 146; CHECK-NEXT: mov.b32 %r4, 1; 147; CHECK-NEXT: st.param.b32 [func_retval0], %r4; 148; CHECK-NEXT: ret; 149; CHECK-NEXT: $L__BB4_2: // %bb2 150; CHECK-NEXT: mov.b32 %r3, 127; 151; CHECK-NEXT: st.param.b32 [func_retval0], %r3; 152; CHECK-NEXT: ret; 153 %p1 = icmp sgt i32 %a, 1 154 %p2 = icmp sgt i32 %b, 1 155 %c = icmp sge i1 %p1, %p2 156 br i1 %c, label %bb1, label %bb2 157bb1: 158 ret i32 1 159bb2: 160 ret i32 127 161} 162 163define i32 @icmp_i1_sle(i32 %a, i32 %b) { 164; CHECK-LABEL: icmp_i1_sle( 165; CHECK: { 166; CHECK-NEXT: .reg .pred %p<4>; 167; CHECK-NEXT: .reg .b32 %r<5>; 168; CHECK-EMPTY: 169; CHECK-NEXT: // %bb.0: 170; CHECK-NEXT: ld.param.u32 %r1, [icmp_i1_sle_param_0]; 171; CHECK-NEXT: setp.lt.s32 %p1, %r1, 2; 172; CHECK-NEXT: ld.param.u32 %r2, [icmp_i1_sle_param_1]; 173; CHECK-NEXT: setp.gt.s32 %p2, %r2, 1; 174; CHECK-NEXT: and.pred %p3, %p2, %p1; 175; CHECK-NEXT: @%p3 bra $L__BB5_2; 176; CHECK-NEXT: // %bb.1: // %bb1 177; CHECK-NEXT: mov.b32 %r4, 1; 178; CHECK-NEXT: st.param.b32 [func_retval0], %r4; 179; CHECK-NEXT: ret; 180; CHECK-NEXT: $L__BB5_2: // %bb2 181; CHECK-NEXT: mov.b32 %r3, 127; 182; CHECK-NEXT: st.param.b32 [func_retval0], %r3; 183; CHECK-NEXT: ret; 184 %p1 = icmp sgt i32 %a, 1 185 %p2 = icmp sgt i32 %b, 1 186 %c = icmp sle i1 %p1, %p2 187 br i1 %c, label %bb1, label %bb2 188bb1: 189 ret i32 1 190bb2: 191 ret i32 127 192} 193 194define i32 @icmp_i1_uge(i32 %a, i32 %b) { 195; CHECK-LABEL: icmp_i1_uge( 196; CHECK: { 197; CHECK-NEXT: .reg .pred %p<4>; 198; CHECK-NEXT: .reg .b32 %r<5>; 199; CHECK-EMPTY: 200; CHECK-NEXT: // %bb.0: 201; CHECK-NEXT: ld.param.u32 %r1, [icmp_i1_uge_param_0]; 202; CHECK-NEXT: setp.lt.s32 %p1, %r1, 2; 203; CHECK-NEXT: ld.param.u32 %r2, [icmp_i1_uge_param_1]; 204; CHECK-NEXT: setp.gt.s32 %p2, %r2, 1; 205; CHECK-NEXT: and.pred %p3, %p2, %p1; 206; CHECK-NEXT: @%p3 bra $L__BB6_2; 207; CHECK-NEXT: // %bb.1: // %bb1 208; CHECK-NEXT: mov.b32 %r4, 1; 209; CHECK-NEXT: st.param.b32 [func_retval0], %r4; 210; CHECK-NEXT: ret; 211; CHECK-NEXT: $L__BB6_2: // %bb2 212; CHECK-NEXT: mov.b32 %r3, 127; 213; CHECK-NEXT: st.param.b32 [func_retval0], %r3; 214; CHECK-NEXT: ret; 215 %p1 = icmp sgt i32 %a, 1 216 %p2 = icmp sgt i32 %b, 1 217 %c = icmp uge i1 %p1, %p2 218 br i1 %c, label %bb1, label %bb2 219bb1: 220 ret i32 1 221bb2: 222 ret i32 127 223} 224 225define i32 @icmp_i1_ugt(i32 %a, i32 %b) { 226; CHECK-LABEL: icmp_i1_ugt( 227; CHECK: { 228; CHECK-NEXT: .reg .pred %p<4>; 229; CHECK-NEXT: .reg .b32 %r<5>; 230; CHECK-EMPTY: 231; CHECK-NEXT: // %bb.0: 232; CHECK-NEXT: ld.param.u32 %r1, [icmp_i1_ugt_param_0]; 233; CHECK-NEXT: setp.lt.s32 %p1, %r1, 2; 234; CHECK-NEXT: ld.param.u32 %r2, [icmp_i1_ugt_param_1]; 235; CHECK-NEXT: setp.gt.s32 %p2, %r2, 1; 236; CHECK-NEXT: or.pred %p3, %p2, %p1; 237; CHECK-NEXT: @%p3 bra $L__BB7_2; 238; CHECK-NEXT: // %bb.1: // %bb1 239; CHECK-NEXT: mov.b32 %r4, 1; 240; CHECK-NEXT: st.param.b32 [func_retval0], %r4; 241; CHECK-NEXT: ret; 242; CHECK-NEXT: $L__BB7_2: // %bb2 243; CHECK-NEXT: mov.b32 %r3, 127; 244; CHECK-NEXT: st.param.b32 [func_retval0], %r3; 245; CHECK-NEXT: ret; 246 %p1 = icmp sgt i32 %a, 1 247 %p2 = icmp sgt i32 %b, 1 248 %c = icmp ugt i1 %p1, %p2 249 br i1 %c, label %bb1, label %bb2 250bb1: 251 ret i32 1 252bb2: 253 ret i32 127 254} 255 256define i32 @icmp_i1_ule(i32 %a, i32 %b) { 257; CHECK-LABEL: icmp_i1_ule( 258; CHECK: { 259; CHECK-NEXT: .reg .pred %p<4>; 260; CHECK-NEXT: .reg .b32 %r<5>; 261; CHECK-EMPTY: 262; CHECK-NEXT: // %bb.0: 263; CHECK-NEXT: ld.param.u32 %r1, [icmp_i1_ule_param_0]; 264; CHECK-NEXT: setp.gt.s32 %p1, %r1, 1; 265; CHECK-NEXT: ld.param.u32 %r2, [icmp_i1_ule_param_1]; 266; CHECK-NEXT: setp.lt.s32 %p2, %r2, 2; 267; CHECK-NEXT: and.pred %p3, %p1, %p2; 268; CHECK-NEXT: @%p3 bra $L__BB8_2; 269; CHECK-NEXT: // %bb.1: // %bb1 270; CHECK-NEXT: mov.b32 %r4, 1; 271; CHECK-NEXT: st.param.b32 [func_retval0], %r4; 272; CHECK-NEXT: ret; 273; CHECK-NEXT: $L__BB8_2: // %bb2 274; CHECK-NEXT: mov.b32 %r3, 127; 275; CHECK-NEXT: st.param.b32 [func_retval0], %r3; 276; CHECK-NEXT: ret; 277 %p1 = icmp sgt i32 %a, 1 278 %p2 = icmp sgt i32 %b, 1 279 %c = icmp ule i1 %p1, %p2 280 br i1 %c, label %bb1, label %bb2 281bb1: 282 ret i32 1 283bb2: 284 ret i32 127 285} 286 287define i32 @icmp_i1_ult(i32 %a, i32 %b) { 288; CHECK-LABEL: icmp_i1_ult( 289; CHECK: { 290; CHECK-NEXT: .reg .pred %p<4>; 291; CHECK-NEXT: .reg .b32 %r<5>; 292; CHECK-EMPTY: 293; CHECK-NEXT: // %bb.0: 294; CHECK-NEXT: ld.param.u32 %r1, [icmp_i1_ult_param_0]; 295; CHECK-NEXT: setp.gt.s32 %p1, %r1, 1; 296; CHECK-NEXT: ld.param.u32 %r2, [icmp_i1_ult_param_1]; 297; CHECK-NEXT: setp.lt.s32 %p2, %r2, 2; 298; CHECK-NEXT: or.pred %p3, %p1, %p2; 299; CHECK-NEXT: @%p3 bra $L__BB9_2; 300; CHECK-NEXT: // %bb.1: // %bb1 301; CHECK-NEXT: mov.b32 %r4, 1; 302; CHECK-NEXT: st.param.b32 [func_retval0], %r4; 303; CHECK-NEXT: ret; 304; CHECK-NEXT: $L__BB9_2: // %bb2 305; CHECK-NEXT: mov.b32 %r3, 127; 306; CHECK-NEXT: st.param.b32 [func_retval0], %r3; 307; CHECK-NEXT: ret; 308 %p1 = icmp sgt i32 %a, 1 309 %p2 = icmp sgt i32 %b, 1 310 %c = icmp ult i1 %p1, %p2 311 br i1 %c, label %bb1, label %bb2 312bb1: 313 ret i32 1 314bb2: 315 ret i32 127 316} 317 318