298 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 | FileCheck %s3; 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<3>;12; CHECK-EMPTY:13; CHECK-NEXT: // %bb.0:14; CHECK-NEXT: ld.param.b32 %r1, [icmp_i1_eq_param_0];15; CHECK-NEXT: setp.gt.s32 %p1, %r1, 1;16; CHECK-NEXT: ld.param.b32 %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: // %bb121; CHECK-NEXT: st.param.b32 [func_retval0], 1;22; CHECK-NEXT: ret;23; CHECK-NEXT: $L__BB0_2: // %bb224; CHECK-NEXT: st.param.b32 [func_retval0], 127;25; CHECK-NEXT: ret;26 %p1 = icmp sgt i32 %a, 127 %p2 = icmp sgt i32 %b, 128 %c = icmp eq i1 %p1, %p229 br i1 %c, label %bb1, label %bb230bb1:31 ret i32 132bb2:33 ret i32 12734}35 36define i32 @icmp_i1_ne(i32 %a, i32 %b) {37; CHECK-LABEL: icmp_i1_ne(38; CHECK: {39; CHECK-NEXT: .reg .pred %p<5>;40; CHECK-NEXT: .reg .b32 %r<3>;41; CHECK-EMPTY:42; CHECK-NEXT: // %bb.0:43; CHECK-NEXT: ld.param.b32 %r1, [icmp_i1_ne_param_0];44; CHECK-NEXT: setp.gt.s32 %p1, %r1, 1;45; CHECK-NEXT: ld.param.b32 %r2, [icmp_i1_ne_param_1];46; CHECK-NEXT: setp.gt.s32 %p2, %r2, 1;47; CHECK-NEXT: xor.pred %p3, %p1, %p2;48; CHECK-NEXT: not.pred %p4, %p3;49; CHECK-NEXT: @%p4 bra $L__BB1_2;50; CHECK-NEXT: // %bb.1: // %bb151; CHECK-NEXT: st.param.b32 [func_retval0], 1;52; CHECK-NEXT: ret;53; CHECK-NEXT: $L__BB1_2: // %bb254; CHECK-NEXT: st.param.b32 [func_retval0], 127;55; CHECK-NEXT: ret;56 %p1 = icmp sgt i32 %a, 157 %p2 = icmp sgt i32 %b, 158 %c = icmp ne i1 %p1, %p259 br i1 %c, label %bb1, label %bb260bb1:61 ret i32 162bb2:63 ret i32 12764}65 66define i32 @icmp_i1_sgt(i32 %a, i32 %b) {67; CHECK-LABEL: icmp_i1_sgt(68; CHECK: {69; CHECK-NEXT: .reg .pred %p<4>;70; CHECK-NEXT: .reg .b32 %r<3>;71; CHECK-EMPTY:72; CHECK-NEXT: // %bb.0:73; CHECK-NEXT: ld.param.b32 %r1, [icmp_i1_sgt_param_0];74; CHECK-NEXT: setp.gt.s32 %p1, %r1, 1;75; CHECK-NEXT: ld.param.b32 %r2, [icmp_i1_sgt_param_1];76; CHECK-NEXT: setp.lt.s32 %p2, %r2, 2;77; CHECK-NEXT: or.pred %p3, %p1, %p2;78; CHECK-NEXT: @%p3 bra $L__BB2_2;79; CHECK-NEXT: // %bb.1: // %bb180; CHECK-NEXT: st.param.b32 [func_retval0], 1;81; CHECK-NEXT: ret;82; CHECK-NEXT: $L__BB2_2: // %bb283; CHECK-NEXT: st.param.b32 [func_retval0], 127;84; CHECK-NEXT: ret;85 %p1 = icmp sgt i32 %a, 186 %p2 = icmp sgt i32 %b, 187 %c = icmp sgt i1 %p1, %p288 br i1 %c, label %bb1, label %bb289bb1:90 ret i32 191bb2:92 ret i32 12793}94 95define i32 @icmp_i1_slt(i32 %a, i32 %b) {96; CHECK-LABEL: icmp_i1_slt(97; CHECK: {98; CHECK-NEXT: .reg .pred %p<4>;99; CHECK-NEXT: .reg .b32 %r<3>;100; CHECK-EMPTY:101; CHECK-NEXT: // %bb.0:102; CHECK-NEXT: ld.param.b32 %r1, [icmp_i1_slt_param_0];103; CHECK-NEXT: setp.lt.s32 %p1, %r1, 2;104; CHECK-NEXT: ld.param.b32 %r2, [icmp_i1_slt_param_1];105; CHECK-NEXT: setp.gt.s32 %p2, %r2, 1;106; CHECK-NEXT: or.pred %p3, %p2, %p1;107; CHECK-NEXT: @%p3 bra $L__BB3_2;108; CHECK-NEXT: // %bb.1: // %bb1109; CHECK-NEXT: st.param.b32 [func_retval0], 1;110; CHECK-NEXT: ret;111; CHECK-NEXT: $L__BB3_2: // %bb2112; CHECK-NEXT: st.param.b32 [func_retval0], 127;113; CHECK-NEXT: ret;114 %p1 = icmp sgt i32 %a, 1115 %p2 = icmp sgt i32 %b, 1116 %c = icmp slt i1 %p1, %p2117 br i1 %c, label %bb1, label %bb2118bb1:119 ret i32 1120bb2:121 ret i32 127122}123 124define i32 @icmp_i1_sge(i32 %a, i32 %b) {125; CHECK-LABEL: icmp_i1_sge(126; CHECK: {127; CHECK-NEXT: .reg .pred %p<4>;128; CHECK-NEXT: .reg .b32 %r<3>;129; CHECK-EMPTY:130; CHECK-NEXT: // %bb.0:131; CHECK-NEXT: ld.param.b32 %r1, [icmp_i1_sge_param_0];132; CHECK-NEXT: setp.gt.s32 %p1, %r1, 1;133; CHECK-NEXT: ld.param.b32 %r2, [icmp_i1_sge_param_1];134; CHECK-NEXT: setp.lt.s32 %p2, %r2, 2;135; CHECK-NEXT: and.pred %p3, %p1, %p2;136; CHECK-NEXT: @%p3 bra $L__BB4_2;137; CHECK-NEXT: // %bb.1: // %bb1138; CHECK-NEXT: st.param.b32 [func_retval0], 1;139; CHECK-NEXT: ret;140; CHECK-NEXT: $L__BB4_2: // %bb2141; CHECK-NEXT: st.param.b32 [func_retval0], 127;142; CHECK-NEXT: ret;143 %p1 = icmp sgt i32 %a, 1144 %p2 = icmp sgt i32 %b, 1145 %c = icmp sge i1 %p1, %p2146 br i1 %c, label %bb1, label %bb2147bb1:148 ret i32 1149bb2:150 ret i32 127151}152 153define i32 @icmp_i1_sle(i32 %a, i32 %b) {154; CHECK-LABEL: icmp_i1_sle(155; CHECK: {156; CHECK-NEXT: .reg .pred %p<4>;157; CHECK-NEXT: .reg .b32 %r<3>;158; CHECK-EMPTY:159; CHECK-NEXT: // %bb.0:160; CHECK-NEXT: ld.param.b32 %r1, [icmp_i1_sle_param_0];161; CHECK-NEXT: setp.lt.s32 %p1, %r1, 2;162; CHECK-NEXT: ld.param.b32 %r2, [icmp_i1_sle_param_1];163; CHECK-NEXT: setp.gt.s32 %p2, %r2, 1;164; CHECK-NEXT: and.pred %p3, %p2, %p1;165; CHECK-NEXT: @%p3 bra $L__BB5_2;166; CHECK-NEXT: // %bb.1: // %bb1167; CHECK-NEXT: st.param.b32 [func_retval0], 1;168; CHECK-NEXT: ret;169; CHECK-NEXT: $L__BB5_2: // %bb2170; CHECK-NEXT: st.param.b32 [func_retval0], 127;171; CHECK-NEXT: ret;172 %p1 = icmp sgt i32 %a, 1173 %p2 = icmp sgt i32 %b, 1174 %c = icmp sle i1 %p1, %p2175 br i1 %c, label %bb1, label %bb2176bb1:177 ret i32 1178bb2:179 ret i32 127180}181 182define i32 @icmp_i1_uge(i32 %a, i32 %b) {183; CHECK-LABEL: icmp_i1_uge(184; CHECK: {185; CHECK-NEXT: .reg .pred %p<4>;186; CHECK-NEXT: .reg .b32 %r<3>;187; CHECK-EMPTY:188; CHECK-NEXT: // %bb.0:189; CHECK-NEXT: ld.param.b32 %r1, [icmp_i1_uge_param_0];190; CHECK-NEXT: setp.lt.s32 %p1, %r1, 2;191; CHECK-NEXT: ld.param.b32 %r2, [icmp_i1_uge_param_1];192; CHECK-NEXT: setp.gt.s32 %p2, %r2, 1;193; CHECK-NEXT: and.pred %p3, %p2, %p1;194; CHECK-NEXT: @%p3 bra $L__BB6_2;195; CHECK-NEXT: // %bb.1: // %bb1196; CHECK-NEXT: st.param.b32 [func_retval0], 1;197; CHECK-NEXT: ret;198; CHECK-NEXT: $L__BB6_2: // %bb2199; CHECK-NEXT: st.param.b32 [func_retval0], 127;200; CHECK-NEXT: ret;201 %p1 = icmp sgt i32 %a, 1202 %p2 = icmp sgt i32 %b, 1203 %c = icmp uge i1 %p1, %p2204 br i1 %c, label %bb1, label %bb2205bb1:206 ret i32 1207bb2:208 ret i32 127209}210 211define i32 @icmp_i1_ugt(i32 %a, i32 %b) {212; CHECK-LABEL: icmp_i1_ugt(213; CHECK: {214; CHECK-NEXT: .reg .pred %p<4>;215; CHECK-NEXT: .reg .b32 %r<3>;216; CHECK-EMPTY:217; CHECK-NEXT: // %bb.0:218; CHECK-NEXT: ld.param.b32 %r1, [icmp_i1_ugt_param_0];219; CHECK-NEXT: setp.lt.s32 %p1, %r1, 2;220; CHECK-NEXT: ld.param.b32 %r2, [icmp_i1_ugt_param_1];221; CHECK-NEXT: setp.gt.s32 %p2, %r2, 1;222; CHECK-NEXT: or.pred %p3, %p2, %p1;223; CHECK-NEXT: @%p3 bra $L__BB7_2;224; CHECK-NEXT: // %bb.1: // %bb1225; CHECK-NEXT: st.param.b32 [func_retval0], 1;226; CHECK-NEXT: ret;227; CHECK-NEXT: $L__BB7_2: // %bb2228; CHECK-NEXT: st.param.b32 [func_retval0], 127;229; CHECK-NEXT: ret;230 %p1 = icmp sgt i32 %a, 1231 %p2 = icmp sgt i32 %b, 1232 %c = icmp ugt i1 %p1, %p2233 br i1 %c, label %bb1, label %bb2234bb1:235 ret i32 1236bb2:237 ret i32 127238}239 240define i32 @icmp_i1_ule(i32 %a, i32 %b) {241; CHECK-LABEL: icmp_i1_ule(242; CHECK: {243; CHECK-NEXT: .reg .pred %p<4>;244; CHECK-NEXT: .reg .b32 %r<3>;245; CHECK-EMPTY:246; CHECK-NEXT: // %bb.0:247; CHECK-NEXT: ld.param.b32 %r1, [icmp_i1_ule_param_0];248; CHECK-NEXT: setp.gt.s32 %p1, %r1, 1;249; CHECK-NEXT: ld.param.b32 %r2, [icmp_i1_ule_param_1];250; CHECK-NEXT: setp.lt.s32 %p2, %r2, 2;251; CHECK-NEXT: and.pred %p3, %p1, %p2;252; CHECK-NEXT: @%p3 bra $L__BB8_2;253; CHECK-NEXT: // %bb.1: // %bb1254; CHECK-NEXT: st.param.b32 [func_retval0], 1;255; CHECK-NEXT: ret;256; CHECK-NEXT: $L__BB8_2: // %bb2257; CHECK-NEXT: st.param.b32 [func_retval0], 127;258; CHECK-NEXT: ret;259 %p1 = icmp sgt i32 %a, 1260 %p2 = icmp sgt i32 %b, 1261 %c = icmp ule i1 %p1, %p2262 br i1 %c, label %bb1, label %bb2263bb1:264 ret i32 1265bb2:266 ret i32 127267}268 269define i32 @icmp_i1_ult(i32 %a, i32 %b) {270; CHECK-LABEL: icmp_i1_ult(271; CHECK: {272; CHECK-NEXT: .reg .pred %p<4>;273; CHECK-NEXT: .reg .b32 %r<3>;274; CHECK-EMPTY:275; CHECK-NEXT: // %bb.0:276; CHECK-NEXT: ld.param.b32 %r1, [icmp_i1_ult_param_0];277; CHECK-NEXT: setp.gt.s32 %p1, %r1, 1;278; CHECK-NEXT: ld.param.b32 %r2, [icmp_i1_ult_param_1];279; CHECK-NEXT: setp.lt.s32 %p2, %r2, 2;280; CHECK-NEXT: or.pred %p3, %p1, %p2;281; CHECK-NEXT: @%p3 bra $L__BB9_2;282; CHECK-NEXT: // %bb.1: // %bb1283; CHECK-NEXT: st.param.b32 [func_retval0], 1;284; CHECK-NEXT: ret;285; CHECK-NEXT: $L__BB9_2: // %bb2286; CHECK-NEXT: st.param.b32 [func_retval0], 127;287; CHECK-NEXT: ret;288 %p1 = icmp sgt i32 %a, 1289 %p2 = icmp sgt i32 %b, 1290 %c = icmp ult i1 %p1, %p2291 br i1 %c, label %bb1, label %bb2292bb1:293 ret i32 1294bb2:295 ret i32 127296}297 298