270 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mcpu=sm_50 | FileCheck %s3; RUN: %if ptxas-sm_50 %{ llc < %s -mcpu=sm_50 | %ptxas-verify -arch=sm_50 %}4 5target triple = "nvptx64-nvidia-cuda"6 7define i1 @trunc_nsw_singed_const(i32 %a) {8; CHECK-LABEL: trunc_nsw_singed_const(9; CHECK: {10; CHECK-NEXT: .reg .pred %p<2>;11; CHECK-NEXT: .reg .b32 %r<4>;12; CHECK-EMPTY:13; CHECK-NEXT: // %bb.0:14; CHECK-NEXT: ld.param.b32 %r1, [trunc_nsw_singed_const_param_0];15; CHECK-NEXT: add.s32 %r2, %r1, 1;16; CHECK-NEXT: setp.gt.s32 %p1, %r2, -1;17; CHECK-NEXT: selp.b32 %r3, -1, 0, %p1;18; CHECK-NEXT: st.param.b32 [func_retval0], %r3;19; CHECK-NEXT: ret;20 %a2 = add i32 %a, 121 %b = trunc nsw i32 %a2 to i822 %c = icmp sgt i8 %b, -123 ret i1 %c24}25 26define i1 @trunc_nuw_singed_const(i32 %a) {27; CHECK-LABEL: trunc_nuw_singed_const(28; CHECK: {29; CHECK-NEXT: .reg .pred %p<2>;30; CHECK-NEXT: .reg .b16 %rs<4>;31; CHECK-NEXT: .reg .b32 %r<2>;32; CHECK-EMPTY:33; CHECK-NEXT: // %bb.0:34; CHECK-NEXT: ld.param.b8 %rs1, [trunc_nuw_singed_const_param_0];35; CHECK-NEXT: add.s16 %rs2, %rs1, 1;36; CHECK-NEXT: cvt.s16.s8 %rs3, %rs2;37; CHECK-NEXT: setp.lt.s16 %p1, %rs3, 100;38; CHECK-NEXT: selp.b32 %r1, -1, 0, %p1;39; CHECK-NEXT: st.param.b32 [func_retval0], %r1;40; CHECK-NEXT: ret;41 %a2 = add i32 %a, 142 %b = trunc nuw i32 %a2 to i843 %c = icmp slt i8 %b, 10044 ret i1 %c45}46 47define i1 @trunc_nsw_unsinged_const(i32 %a) {48; CHECK-LABEL: trunc_nsw_unsinged_const(49; CHECK: {50; CHECK-NEXT: .reg .pred %p<2>;51; CHECK-NEXT: .reg .b16 %rs<4>;52; CHECK-NEXT: .reg .b32 %r<2>;53; CHECK-EMPTY:54; CHECK-NEXT: // %bb.0:55; CHECK-NEXT: ld.param.b8 %rs1, [trunc_nsw_unsinged_const_param_0];56; CHECK-NEXT: add.s16 %rs2, %rs1, 1;57; CHECK-NEXT: and.b16 %rs3, %rs2, 255;58; CHECK-NEXT: setp.lt.u16 %p1, %rs3, 236;59; CHECK-NEXT: selp.b32 %r1, -1, 0, %p1;60; CHECK-NEXT: st.param.b32 [func_retval0], %r1;61; CHECK-NEXT: ret;62 %a2 = add i32 %a, 163 %b = trunc nsw i32 %a2 to i864 %c = icmp ult i8 %b, -2065 ret i1 %c66}67 68define i1 @trunc_nuw_unsinged_const(i32 %a) {69; CHECK-LABEL: trunc_nuw_unsinged_const(70; CHECK: {71; CHECK-NEXT: .reg .pred %p<2>;72; CHECK-NEXT: .reg .b32 %r<4>;73; CHECK-EMPTY:74; CHECK-NEXT: // %bb.0:75; CHECK-NEXT: ld.param.b32 %r1, [trunc_nuw_unsinged_const_param_0];76; CHECK-NEXT: add.s32 %r2, %r1, 1;77; CHECK-NEXT: setp.gt.u32 %p1, %r2, 100;78; CHECK-NEXT: selp.b32 %r3, -1, 0, %p1;79; CHECK-NEXT: st.param.b32 [func_retval0], %r3;80; CHECK-NEXT: ret;81 %a2 = add i32 %a, 182 %b = trunc nuw i32 %a2 to i883 %c = icmp ugt i8 %b, 10084 ret i1 %c85}86 87 88define i1 @trunc_nsw_eq_const(i32 %a) {89; CHECK-LABEL: trunc_nsw_eq_const(90; CHECK: {91; CHECK-NEXT: .reg .pred %p<2>;92; CHECK-NEXT: .reg .b32 %r<3>;93; CHECK-EMPTY:94; CHECK-NEXT: // %bb.0:95; CHECK-NEXT: ld.param.b32 %r1, [trunc_nsw_eq_const_param_0];96; CHECK-NEXT: setp.eq.b32 %p1, %r1, 99;97; CHECK-NEXT: selp.b32 %r2, -1, 0, %p1;98; CHECK-NEXT: st.param.b32 [func_retval0], %r2;99; CHECK-NEXT: ret;100 %a2 = add i32 %a, 1101 %b = trunc nsw i32 %a2 to i8102 %c = icmp eq i8 %b, 100103 ret i1 %c104}105 106define i1 @trunc_nuw_eq_const(i32 %a) {107; CHECK-LABEL: trunc_nuw_eq_const(108; CHECK: {109; CHECK-NEXT: .reg .pred %p<2>;110; CHECK-NEXT: .reg .b32 %r<3>;111; CHECK-EMPTY:112; CHECK-NEXT: // %bb.0:113; CHECK-NEXT: ld.param.b32 %r1, [trunc_nuw_eq_const_param_0];114; CHECK-NEXT: setp.eq.b32 %p1, %r1, 99;115; CHECK-NEXT: selp.b32 %r2, -1, 0, %p1;116; CHECK-NEXT: st.param.b32 [func_retval0], %r2;117; CHECK-NEXT: ret;118 %a2 = add i32 %a, 1119 %b = trunc nuw i32 %a2 to i8120 %c = icmp eq i8 %b, 100121 ret i1 %c122}123 124;;;125 126define i1 @trunc_nsw_singed(i32 %a1, i32 %a2) {127; CHECK-LABEL: trunc_nsw_singed(128; CHECK: {129; CHECK-NEXT: .reg .pred %p<2>;130; CHECK-NEXT: .reg .b32 %r<6>;131; CHECK-EMPTY:132; CHECK-NEXT: // %bb.0:133; CHECK-NEXT: ld.param.b32 %r1, [trunc_nsw_singed_param_0];134; CHECK-NEXT: add.s32 %r2, %r1, 1;135; CHECK-NEXT: ld.param.b32 %r3, [trunc_nsw_singed_param_1];136; CHECK-NEXT: add.s32 %r4, %r3, 7;137; CHECK-NEXT: setp.gt.s32 %p1, %r2, %r4;138; CHECK-NEXT: selp.b32 %r5, -1, 0, %p1;139; CHECK-NEXT: st.param.b32 [func_retval0], %r5;140; CHECK-NEXT: ret;141 %b1 = add i32 %a1, 1142 %b2 = add i32 %a2, 7143 %c1 = trunc nsw i32 %b1 to i8144 %c2 = trunc nsw i32 %b2 to i8145 %c = icmp sgt i8 %c1, %c2146 ret i1 %c147}148 149define i1 @trunc_nuw_singed(i32 %a1, i32 %a2) {150; CHECK-LABEL: trunc_nuw_singed(151; CHECK: {152; CHECK-NEXT: .reg .pred %p<2>;153; CHECK-NEXT: .reg .b16 %rs<7>;154; CHECK-NEXT: .reg .b32 %r<2>;155; CHECK-EMPTY:156; CHECK-NEXT: // %bb.0:157; CHECK-NEXT: ld.param.b8 %rs1, [trunc_nuw_singed_param_0];158; CHECK-NEXT: ld.param.b8 %rs2, [trunc_nuw_singed_param_1];159; CHECK-NEXT: add.s16 %rs3, %rs1, 1;160; CHECK-NEXT: cvt.s16.s8 %rs4, %rs3;161; CHECK-NEXT: add.s16 %rs5, %rs2, 6;162; CHECK-NEXT: cvt.s16.s8 %rs6, %rs5;163; CHECK-NEXT: setp.lt.s16 %p1, %rs4, %rs6;164; CHECK-NEXT: selp.b32 %r1, -1, 0, %p1;165; CHECK-NEXT: st.param.b32 [func_retval0], %r1;166; CHECK-NEXT: ret;167 %b1 = add i32 %a1, 1168 %b2 = add i32 %a2, 6169 %c1 = trunc nuw i32 %b1 to i8170 %c2 = trunc nuw i32 %b2 to i8171 %c = icmp slt i8 %c1, %c2172 ret i1 %c173}174 175define i1 @trunc_nsw_unsinged(i32 %a1, i32 %a2) {176; CHECK-LABEL: trunc_nsw_unsinged(177; CHECK: {178; CHECK-NEXT: .reg .pred %p<2>;179; CHECK-NEXT: .reg .b16 %rs<7>;180; CHECK-NEXT: .reg .b32 %r<2>;181; CHECK-EMPTY:182; CHECK-NEXT: // %bb.0:183; CHECK-NEXT: ld.param.b8 %rs1, [trunc_nsw_unsinged_param_0];184; CHECK-NEXT: ld.param.b8 %rs2, [trunc_nsw_unsinged_param_1];185; CHECK-NEXT: add.s16 %rs3, %rs1, 1;186; CHECK-NEXT: and.b16 %rs4, %rs3, 255;187; CHECK-NEXT: add.s16 %rs5, %rs2, 4;188; CHECK-NEXT: and.b16 %rs6, %rs5, 255;189; CHECK-NEXT: setp.lt.u16 %p1, %rs4, %rs6;190; CHECK-NEXT: selp.b32 %r1, -1, 0, %p1;191; CHECK-NEXT: st.param.b32 [func_retval0], %r1;192; CHECK-NEXT: ret;193 %b1 = add i32 %a1, 1194 %b2 = add i32 %a2, 4195 %c1 = trunc nsw i32 %b1 to i8196 %c2 = trunc nsw i32 %b2 to i8197 %c = icmp ult i8 %c1, %c2198 ret i1 %c199}200 201define i1 @trunc_nuw_unsinged(i32 %a1, i32 %a2) {202; CHECK-LABEL: trunc_nuw_unsinged(203; CHECK: {204; CHECK-NEXT: .reg .pred %p<2>;205; CHECK-NEXT: .reg .b32 %r<6>;206; CHECK-EMPTY:207; CHECK-NEXT: // %bb.0:208; CHECK-NEXT: ld.param.b32 %r1, [trunc_nuw_unsinged_param_0];209; CHECK-NEXT: add.s32 %r2, %r1, 1;210; CHECK-NEXT: ld.param.b32 %r3, [trunc_nuw_unsinged_param_1];211; CHECK-NEXT: add.s32 %r4, %r3, 5;212; CHECK-NEXT: setp.gt.u32 %p1, %r2, %r4;213; CHECK-NEXT: selp.b32 %r5, -1, 0, %p1;214; CHECK-NEXT: st.param.b32 [func_retval0], %r5;215; CHECK-NEXT: ret;216 %b1 = add i32 %a1, 1217 %b2 = add i32 %a2, 5218 %c1 = trunc nuw i32 %b1 to i8219 %c2 = trunc nuw i32 %b2 to i8220 %c = icmp ugt i8 %c1, %c2221 ret i1 %c222}223 224 225define i1 @trunc_nsw_eq(i32 %a1, i32 %a2) {226; CHECK-LABEL: trunc_nsw_eq(227; CHECK: {228; CHECK-NEXT: .reg .pred %p<2>;229; CHECK-NEXT: .reg .b32 %r<6>;230; CHECK-EMPTY:231; CHECK-NEXT: // %bb.0:232; CHECK-NEXT: ld.param.b32 %r1, [trunc_nsw_eq_param_0];233; CHECK-NEXT: add.s32 %r2, %r1, 1;234; CHECK-NEXT: ld.param.b32 %r3, [trunc_nsw_eq_param_1];235; CHECK-NEXT: add.s32 %r4, %r3, 3;236; CHECK-NEXT: setp.eq.b32 %p1, %r2, %r4;237; CHECK-NEXT: selp.b32 %r5, -1, 0, %p1;238; CHECK-NEXT: st.param.b32 [func_retval0], %r5;239; CHECK-NEXT: ret;240 %b1 = add i32 %a1, 1241 %b2 = add i32 %a2, 3242 %c1 = trunc nsw i32 %b1 to i8243 %c2 = trunc nsw i32 %b2 to i8244 %c = icmp eq i8 %c1, %c2245 ret i1 %c246}247 248define i1 @trunc_nuw_eq(i32 %a1, i32 %a2) {249; CHECK-LABEL: trunc_nuw_eq(250; CHECK: {251; CHECK-NEXT: .reg .pred %p<2>;252; CHECK-NEXT: .reg .b32 %r<6>;253; CHECK-EMPTY:254; CHECK-NEXT: // %bb.0:255; CHECK-NEXT: ld.param.b32 %r1, [trunc_nuw_eq_param_0];256; CHECK-NEXT: add.s32 %r2, %r1, 2;257; CHECK-NEXT: ld.param.b32 %r3, [trunc_nuw_eq_param_1];258; CHECK-NEXT: add.s32 %r4, %r3, 1;259; CHECK-NEXT: setp.eq.b32 %p1, %r2, %r4;260; CHECK-NEXT: selp.b32 %r5, -1, 0, %p1;261; CHECK-NEXT: st.param.b32 [func_retval0], %r5;262; CHECK-NEXT: ret;263 %b1 = add i32 %a1, 2264 %b2 = add i32 %a2, 1265 %c1 = trunc nuw i32 %b1 to i8266 %c2 = trunc nuw i32 %b2 to i8267 %c = icmp eq i8 %c1, %c2268 ret i1 %c269}270