771 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mtriple=nvptx -mcpu=sm_20 | FileCheck %s3; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 | FileCheck %s4; RUN: %if ptxas-ptr32 %{ llc < %s -mtriple=nvptx -mcpu=sm_20 | %ptxas-verify %}5; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 | %ptxas-verify %}6 7;; These tests should run for all targets8 9;;===-- Basic instruction selection tests ---------------------------------===;;10 11 12;;; i6413 14define i64 @icmp_eq_i64(i64 %a, i64 %b) {15; CHECK-LABEL: icmp_eq_i64(16; CHECK: {17; CHECK-NEXT: .reg .pred %p<2>;18; CHECK-NEXT: .reg .b64 %rd<4>;19; CHECK-EMPTY:20; CHECK-NEXT: // %bb.0:21; CHECK-NEXT: ld.param.b64 %rd1, [icmp_eq_i64_param_0];22; CHECK-NEXT: ld.param.b64 %rd2, [icmp_eq_i64_param_1];23; CHECK-NEXT: setp.eq.b64 %p1, %rd1, %rd2;24; CHECK-NEXT: selp.b64 %rd3, 1, 0, %p1;25; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;26; CHECK-NEXT: ret;27 %cmp = icmp eq i64 %a, %b28 %ret = zext i1 %cmp to i6429 ret i64 %ret30}31 32define i64 @icmp_ne_i64(i64 %a, i64 %b) {33; CHECK-LABEL: icmp_ne_i64(34; CHECK: {35; CHECK-NEXT: .reg .pred %p<2>;36; CHECK-NEXT: .reg .b64 %rd<4>;37; CHECK-EMPTY:38; CHECK-NEXT: // %bb.0:39; CHECK-NEXT: ld.param.b64 %rd1, [icmp_ne_i64_param_0];40; CHECK-NEXT: ld.param.b64 %rd2, [icmp_ne_i64_param_1];41; CHECK-NEXT: setp.ne.b64 %p1, %rd1, %rd2;42; CHECK-NEXT: selp.b64 %rd3, 1, 0, %p1;43; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;44; CHECK-NEXT: ret;45 %cmp = icmp ne i64 %a, %b46 %ret = zext i1 %cmp to i6447 ret i64 %ret48}49 50define i64 @icmp_ugt_i64(i64 %a, i64 %b) {51; CHECK-LABEL: icmp_ugt_i64(52; CHECK: {53; CHECK-NEXT: .reg .pred %p<2>;54; CHECK-NEXT: .reg .b64 %rd<4>;55; CHECK-EMPTY:56; CHECK-NEXT: // %bb.0:57; CHECK-NEXT: ld.param.b64 %rd1, [icmp_ugt_i64_param_0];58; CHECK-NEXT: ld.param.b64 %rd2, [icmp_ugt_i64_param_1];59; CHECK-NEXT: setp.gt.u64 %p1, %rd1, %rd2;60; CHECK-NEXT: selp.b64 %rd3, 1, 0, %p1;61; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;62; CHECK-NEXT: ret;63 %cmp = icmp ugt i64 %a, %b64 %ret = zext i1 %cmp to i6465 ret i64 %ret66}67 68define i64 @icmp_uge_i64(i64 %a, i64 %b) {69; CHECK-LABEL: icmp_uge_i64(70; CHECK: {71; CHECK-NEXT: .reg .pred %p<2>;72; CHECK-NEXT: .reg .b64 %rd<4>;73; CHECK-EMPTY:74; CHECK-NEXT: // %bb.0:75; CHECK-NEXT: ld.param.b64 %rd1, [icmp_uge_i64_param_0];76; CHECK-NEXT: ld.param.b64 %rd2, [icmp_uge_i64_param_1];77; CHECK-NEXT: setp.ge.u64 %p1, %rd1, %rd2;78; CHECK-NEXT: selp.b64 %rd3, 1, 0, %p1;79; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;80; CHECK-NEXT: ret;81 %cmp = icmp uge i64 %a, %b82 %ret = zext i1 %cmp to i6483 ret i64 %ret84}85 86define i64 @icmp_ult_i64(i64 %a, i64 %b) {87; CHECK-LABEL: icmp_ult_i64(88; CHECK: {89; CHECK-NEXT: .reg .pred %p<2>;90; CHECK-NEXT: .reg .b64 %rd<4>;91; CHECK-EMPTY:92; CHECK-NEXT: // %bb.0:93; CHECK-NEXT: ld.param.b64 %rd1, [icmp_ult_i64_param_0];94; CHECK-NEXT: ld.param.b64 %rd2, [icmp_ult_i64_param_1];95; CHECK-NEXT: setp.lt.u64 %p1, %rd1, %rd2;96; CHECK-NEXT: selp.b64 %rd3, 1, 0, %p1;97; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;98; CHECK-NEXT: ret;99 %cmp = icmp ult i64 %a, %b100 %ret = zext i1 %cmp to i64101 ret i64 %ret102}103 104define i64 @icmp_ule_i64(i64 %a, i64 %b) {105; CHECK-LABEL: icmp_ule_i64(106; CHECK: {107; CHECK-NEXT: .reg .pred %p<2>;108; CHECK-NEXT: .reg .b64 %rd<4>;109; CHECK-EMPTY:110; CHECK-NEXT: // %bb.0:111; CHECK-NEXT: ld.param.b64 %rd1, [icmp_ule_i64_param_0];112; CHECK-NEXT: ld.param.b64 %rd2, [icmp_ule_i64_param_1];113; CHECK-NEXT: setp.le.u64 %p1, %rd1, %rd2;114; CHECK-NEXT: selp.b64 %rd3, 1, 0, %p1;115; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;116; CHECK-NEXT: ret;117 %cmp = icmp ule i64 %a, %b118 %ret = zext i1 %cmp to i64119 ret i64 %ret120}121 122define i64 @icmp_sgt_i64(i64 %a, i64 %b) {123; CHECK-LABEL: icmp_sgt_i64(124; CHECK: {125; CHECK-NEXT: .reg .pred %p<2>;126; CHECK-NEXT: .reg .b64 %rd<4>;127; CHECK-EMPTY:128; CHECK-NEXT: // %bb.0:129; CHECK-NEXT: ld.param.b64 %rd1, [icmp_sgt_i64_param_0];130; CHECK-NEXT: ld.param.b64 %rd2, [icmp_sgt_i64_param_1];131; CHECK-NEXT: setp.gt.s64 %p1, %rd1, %rd2;132; CHECK-NEXT: selp.b64 %rd3, 1, 0, %p1;133; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;134; CHECK-NEXT: ret;135 %cmp = icmp sgt i64 %a, %b136 %ret = zext i1 %cmp to i64137 ret i64 %ret138}139 140define i64 @icmp_sge_i64(i64 %a, i64 %b) {141; CHECK-LABEL: icmp_sge_i64(142; CHECK: {143; CHECK-NEXT: .reg .pred %p<2>;144; CHECK-NEXT: .reg .b64 %rd<4>;145; CHECK-EMPTY:146; CHECK-NEXT: // %bb.0:147; CHECK-NEXT: ld.param.b64 %rd1, [icmp_sge_i64_param_0];148; CHECK-NEXT: ld.param.b64 %rd2, [icmp_sge_i64_param_1];149; CHECK-NEXT: setp.ge.s64 %p1, %rd1, %rd2;150; CHECK-NEXT: selp.b64 %rd3, 1, 0, %p1;151; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;152; CHECK-NEXT: ret;153 %cmp = icmp sge i64 %a, %b154 %ret = zext i1 %cmp to i64155 ret i64 %ret156}157 158define i64 @icmp_slt_i64(i64 %a, i64 %b) {159; CHECK-LABEL: icmp_slt_i64(160; CHECK: {161; CHECK-NEXT: .reg .pred %p<2>;162; CHECK-NEXT: .reg .b64 %rd<4>;163; CHECK-EMPTY:164; CHECK-NEXT: // %bb.0:165; CHECK-NEXT: ld.param.b64 %rd1, [icmp_slt_i64_param_0];166; CHECK-NEXT: ld.param.b64 %rd2, [icmp_slt_i64_param_1];167; CHECK-NEXT: setp.lt.s64 %p1, %rd1, %rd2;168; CHECK-NEXT: selp.b64 %rd3, 1, 0, %p1;169; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;170; CHECK-NEXT: ret;171 %cmp = icmp slt i64 %a, %b172 %ret = zext i1 %cmp to i64173 ret i64 %ret174}175 176define i64 @icmp_sle_i64(i64 %a, i64 %b) {177; CHECK-LABEL: icmp_sle_i64(178; CHECK: {179; CHECK-NEXT: .reg .pred %p<2>;180; CHECK-NEXT: .reg .b64 %rd<4>;181; CHECK-EMPTY:182; CHECK-NEXT: // %bb.0:183; CHECK-NEXT: ld.param.b64 %rd1, [icmp_sle_i64_param_0];184; CHECK-NEXT: ld.param.b64 %rd2, [icmp_sle_i64_param_1];185; CHECK-NEXT: setp.le.s64 %p1, %rd1, %rd2;186; CHECK-NEXT: selp.b64 %rd3, 1, 0, %p1;187; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;188; CHECK-NEXT: ret;189 %cmp = icmp sle i64 %a, %b190 %ret = zext i1 %cmp to i64191 ret i64 %ret192}193 194;;; i32195 196define i32 @icmp_eq_i32(i32 %a, i32 %b) {197; CHECK-LABEL: icmp_eq_i32(198; CHECK: {199; CHECK-NEXT: .reg .pred %p<2>;200; CHECK-NEXT: .reg .b32 %r<4>;201; CHECK-EMPTY:202; CHECK-NEXT: // %bb.0:203; CHECK-NEXT: ld.param.b32 %r1, [icmp_eq_i32_param_0];204; CHECK-NEXT: ld.param.b32 %r2, [icmp_eq_i32_param_1];205; CHECK-NEXT: setp.eq.b32 %p1, %r1, %r2;206; CHECK-NEXT: selp.b32 %r3, 1, 0, %p1;207; CHECK-NEXT: st.param.b32 [func_retval0], %r3;208; CHECK-NEXT: ret;209 %cmp = icmp eq i32 %a, %b210 %ret = zext i1 %cmp to i32211 ret i32 %ret212}213 214define i32 @icmp_ne_i32(i32 %a, i32 %b) {215; CHECK-LABEL: icmp_ne_i32(216; CHECK: {217; CHECK-NEXT: .reg .pred %p<2>;218; CHECK-NEXT: .reg .b32 %r<4>;219; CHECK-EMPTY:220; CHECK-NEXT: // %bb.0:221; CHECK-NEXT: ld.param.b32 %r1, [icmp_ne_i32_param_0];222; CHECK-NEXT: ld.param.b32 %r2, [icmp_ne_i32_param_1];223; CHECK-NEXT: setp.ne.b32 %p1, %r1, %r2;224; CHECK-NEXT: selp.b32 %r3, 1, 0, %p1;225; CHECK-NEXT: st.param.b32 [func_retval0], %r3;226; CHECK-NEXT: ret;227 %cmp = icmp ne i32 %a, %b228 %ret = zext i1 %cmp to i32229 ret i32 %ret230}231 232define i32 @icmp_ugt_i32(i32 %a, i32 %b) {233; CHECK-LABEL: icmp_ugt_i32(234; CHECK: {235; CHECK-NEXT: .reg .pred %p<2>;236; CHECK-NEXT: .reg .b32 %r<4>;237; CHECK-EMPTY:238; CHECK-NEXT: // %bb.0:239; CHECK-NEXT: ld.param.b32 %r1, [icmp_ugt_i32_param_0];240; CHECK-NEXT: ld.param.b32 %r2, [icmp_ugt_i32_param_1];241; CHECK-NEXT: setp.gt.u32 %p1, %r1, %r2;242; CHECK-NEXT: selp.b32 %r3, 1, 0, %p1;243; CHECK-NEXT: st.param.b32 [func_retval0], %r3;244; CHECK-NEXT: ret;245 %cmp = icmp ugt i32 %a, %b246 %ret = zext i1 %cmp to i32247 ret i32 %ret248}249 250define i32 @icmp_uge_i32(i32 %a, i32 %b) {251; CHECK-LABEL: icmp_uge_i32(252; CHECK: {253; CHECK-NEXT: .reg .pred %p<2>;254; CHECK-NEXT: .reg .b32 %r<4>;255; CHECK-EMPTY:256; CHECK-NEXT: // %bb.0:257; CHECK-NEXT: ld.param.b32 %r1, [icmp_uge_i32_param_0];258; CHECK-NEXT: ld.param.b32 %r2, [icmp_uge_i32_param_1];259; CHECK-NEXT: setp.ge.u32 %p1, %r1, %r2;260; CHECK-NEXT: selp.b32 %r3, 1, 0, %p1;261; CHECK-NEXT: st.param.b32 [func_retval0], %r3;262; CHECK-NEXT: ret;263 %cmp = icmp uge i32 %a, %b264 %ret = zext i1 %cmp to i32265 ret i32 %ret266}267 268define i32 @icmp_ult_i32(i32 %a, i32 %b) {269; CHECK-LABEL: icmp_ult_i32(270; CHECK: {271; CHECK-NEXT: .reg .pred %p<2>;272; CHECK-NEXT: .reg .b32 %r<4>;273; CHECK-EMPTY:274; CHECK-NEXT: // %bb.0:275; CHECK-NEXT: ld.param.b32 %r1, [icmp_ult_i32_param_0];276; CHECK-NEXT: ld.param.b32 %r2, [icmp_ult_i32_param_1];277; CHECK-NEXT: setp.lt.u32 %p1, %r1, %r2;278; CHECK-NEXT: selp.b32 %r3, 1, 0, %p1;279; CHECK-NEXT: st.param.b32 [func_retval0], %r3;280; CHECK-NEXT: ret;281 %cmp = icmp ult i32 %a, %b282 %ret = zext i1 %cmp to i32283 ret i32 %ret284}285 286define i32 @icmp_ule_i32(i32 %a, i32 %b) {287; CHECK-LABEL: icmp_ule_i32(288; CHECK: {289; CHECK-NEXT: .reg .pred %p<2>;290; CHECK-NEXT: .reg .b32 %r<4>;291; CHECK-EMPTY:292; CHECK-NEXT: // %bb.0:293; CHECK-NEXT: ld.param.b32 %r1, [icmp_ule_i32_param_0];294; CHECK-NEXT: ld.param.b32 %r2, [icmp_ule_i32_param_1];295; CHECK-NEXT: setp.le.u32 %p1, %r1, %r2;296; CHECK-NEXT: selp.b32 %r3, 1, 0, %p1;297; CHECK-NEXT: st.param.b32 [func_retval0], %r3;298; CHECK-NEXT: ret;299 %cmp = icmp ule i32 %a, %b300 %ret = zext i1 %cmp to i32301 ret i32 %ret302}303 304define i32 @icmp_sgt_i32(i32 %a, i32 %b) {305; CHECK-LABEL: icmp_sgt_i32(306; CHECK: {307; CHECK-NEXT: .reg .pred %p<2>;308; CHECK-NEXT: .reg .b32 %r<4>;309; CHECK-EMPTY:310; CHECK-NEXT: // %bb.0:311; CHECK-NEXT: ld.param.b32 %r1, [icmp_sgt_i32_param_0];312; CHECK-NEXT: ld.param.b32 %r2, [icmp_sgt_i32_param_1];313; CHECK-NEXT: setp.gt.s32 %p1, %r1, %r2;314; CHECK-NEXT: selp.b32 %r3, 1, 0, %p1;315; CHECK-NEXT: st.param.b32 [func_retval0], %r3;316; CHECK-NEXT: ret;317 %cmp = icmp sgt i32 %a, %b318 %ret = zext i1 %cmp to i32319 ret i32 %ret320}321 322define i32 @icmp_sge_i32(i32 %a, i32 %b) {323; CHECK-LABEL: icmp_sge_i32(324; CHECK: {325; CHECK-NEXT: .reg .pred %p<2>;326; CHECK-NEXT: .reg .b32 %r<4>;327; CHECK-EMPTY:328; CHECK-NEXT: // %bb.0:329; CHECK-NEXT: ld.param.b32 %r1, [icmp_sge_i32_param_0];330; CHECK-NEXT: ld.param.b32 %r2, [icmp_sge_i32_param_1];331; CHECK-NEXT: setp.ge.s32 %p1, %r1, %r2;332; CHECK-NEXT: selp.b32 %r3, 1, 0, %p1;333; CHECK-NEXT: st.param.b32 [func_retval0], %r3;334; CHECK-NEXT: ret;335 %cmp = icmp sge i32 %a, %b336 %ret = zext i1 %cmp to i32337 ret i32 %ret338}339 340define i32 @icmp_slt_i32(i32 %a, i32 %b) {341; CHECK-LABEL: icmp_slt_i32(342; CHECK: {343; CHECK-NEXT: .reg .pred %p<2>;344; CHECK-NEXT: .reg .b32 %r<4>;345; CHECK-EMPTY:346; CHECK-NEXT: // %bb.0:347; CHECK-NEXT: ld.param.b32 %r1, [icmp_slt_i32_param_0];348; CHECK-NEXT: ld.param.b32 %r2, [icmp_slt_i32_param_1];349; CHECK-NEXT: setp.lt.s32 %p1, %r1, %r2;350; CHECK-NEXT: selp.b32 %r3, 1, 0, %p1;351; CHECK-NEXT: st.param.b32 [func_retval0], %r3;352; CHECK-NEXT: ret;353 %cmp = icmp slt i32 %a, %b354 %ret = zext i1 %cmp to i32355 ret i32 %ret356}357 358define i32 @icmp_sle_i32(i32 %a, i32 %b) {359; CHECK-LABEL: icmp_sle_i32(360; CHECK: {361; CHECK-NEXT: .reg .pred %p<2>;362; CHECK-NEXT: .reg .b32 %r<4>;363; CHECK-EMPTY:364; CHECK-NEXT: // %bb.0:365; CHECK-NEXT: ld.param.b32 %r1, [icmp_sle_i32_param_0];366; CHECK-NEXT: ld.param.b32 %r2, [icmp_sle_i32_param_1];367; CHECK-NEXT: setp.le.s32 %p1, %r1, %r2;368; CHECK-NEXT: selp.b32 %r3, 1, 0, %p1;369; CHECK-NEXT: st.param.b32 [func_retval0], %r3;370; CHECK-NEXT: ret;371 %cmp = icmp sle i32 %a, %b372 %ret = zext i1 %cmp to i32373 ret i32 %ret374}375 376 377;;; i16378 379define i16 @icmp_eq_i16(i16 %a, i16 %b) {380; CHECK-LABEL: icmp_eq_i16(381; CHECK: {382; CHECK-NEXT: .reg .pred %p<2>;383; CHECK-NEXT: .reg .b16 %rs<3>;384; CHECK-NEXT: .reg .b32 %r<2>;385; CHECK-EMPTY:386; CHECK-NEXT: // %bb.0:387; CHECK-NEXT: ld.param.b16 %rs1, [icmp_eq_i16_param_0];388; CHECK-NEXT: ld.param.b16 %rs2, [icmp_eq_i16_param_1];389; CHECK-NEXT: setp.eq.b16 %p1, %rs1, %rs2;390; CHECK-NEXT: selp.b32 %r1, 1, 0, %p1;391; CHECK-NEXT: st.param.b32 [func_retval0], %r1;392; CHECK-NEXT: ret;393 %cmp = icmp eq i16 %a, %b394 %ret = zext i1 %cmp to i16395 ret i16 %ret396}397 398define i16 @icmp_ne_i16(i16 %a, i16 %b) {399; CHECK-LABEL: icmp_ne_i16(400; CHECK: {401; CHECK-NEXT: .reg .pred %p<2>;402; CHECK-NEXT: .reg .b16 %rs<3>;403; CHECK-NEXT: .reg .b32 %r<2>;404; CHECK-EMPTY:405; CHECK-NEXT: // %bb.0:406; CHECK-NEXT: ld.param.b16 %rs1, [icmp_ne_i16_param_0];407; CHECK-NEXT: ld.param.b16 %rs2, [icmp_ne_i16_param_1];408; CHECK-NEXT: setp.ne.b16 %p1, %rs1, %rs2;409; CHECK-NEXT: selp.b32 %r1, 1, 0, %p1;410; CHECK-NEXT: st.param.b32 [func_retval0], %r1;411; CHECK-NEXT: ret;412 %cmp = icmp ne i16 %a, %b413 %ret = zext i1 %cmp to i16414 ret i16 %ret415}416 417define i16 @icmp_ugt_i16(i16 %a, i16 %b) {418; CHECK-LABEL: icmp_ugt_i16(419; CHECK: {420; CHECK-NEXT: .reg .pred %p<2>;421; CHECK-NEXT: .reg .b16 %rs<3>;422; CHECK-NEXT: .reg .b32 %r<2>;423; CHECK-EMPTY:424; CHECK-NEXT: // %bb.0:425; CHECK-NEXT: ld.param.b16 %rs1, [icmp_ugt_i16_param_0];426; CHECK-NEXT: ld.param.b16 %rs2, [icmp_ugt_i16_param_1];427; CHECK-NEXT: setp.gt.u16 %p1, %rs1, %rs2;428; CHECK-NEXT: selp.b32 %r1, 1, 0, %p1;429; CHECK-NEXT: st.param.b32 [func_retval0], %r1;430; CHECK-NEXT: ret;431 %cmp = icmp ugt i16 %a, %b432 %ret = zext i1 %cmp to i16433 ret i16 %ret434}435 436define i16 @icmp_uge_i16(i16 %a, i16 %b) {437; CHECK-LABEL: icmp_uge_i16(438; CHECK: {439; CHECK-NEXT: .reg .pred %p<2>;440; CHECK-NEXT: .reg .b16 %rs<3>;441; CHECK-NEXT: .reg .b32 %r<2>;442; CHECK-EMPTY:443; CHECK-NEXT: // %bb.0:444; CHECK-NEXT: ld.param.b16 %rs1, [icmp_uge_i16_param_0];445; CHECK-NEXT: ld.param.b16 %rs2, [icmp_uge_i16_param_1];446; CHECK-NEXT: setp.ge.u16 %p1, %rs1, %rs2;447; CHECK-NEXT: selp.b32 %r1, 1, 0, %p1;448; CHECK-NEXT: st.param.b32 [func_retval0], %r1;449; CHECK-NEXT: ret;450 %cmp = icmp uge i16 %a, %b451 %ret = zext i1 %cmp to i16452 ret i16 %ret453}454 455define i16 @icmp_ult_i16(i16 %a, i16 %b) {456; CHECK-LABEL: icmp_ult_i16(457; CHECK: {458; CHECK-NEXT: .reg .pred %p<2>;459; CHECK-NEXT: .reg .b16 %rs<3>;460; CHECK-NEXT: .reg .b32 %r<2>;461; CHECK-EMPTY:462; CHECK-NEXT: // %bb.0:463; CHECK-NEXT: ld.param.b16 %rs1, [icmp_ult_i16_param_0];464; CHECK-NEXT: ld.param.b16 %rs2, [icmp_ult_i16_param_1];465; CHECK-NEXT: setp.lt.u16 %p1, %rs1, %rs2;466; CHECK-NEXT: selp.b32 %r1, 1, 0, %p1;467; CHECK-NEXT: st.param.b32 [func_retval0], %r1;468; CHECK-NEXT: ret;469 %cmp = icmp ult i16 %a, %b470 %ret = zext i1 %cmp to i16471 ret i16 %ret472}473 474define i16 @icmp_ule_i16(i16 %a, i16 %b) {475; CHECK-LABEL: icmp_ule_i16(476; CHECK: {477; CHECK-NEXT: .reg .pred %p<2>;478; CHECK-NEXT: .reg .b16 %rs<3>;479; CHECK-NEXT: .reg .b32 %r<2>;480; CHECK-EMPTY:481; CHECK-NEXT: // %bb.0:482; CHECK-NEXT: ld.param.b16 %rs1, [icmp_ule_i16_param_0];483; CHECK-NEXT: ld.param.b16 %rs2, [icmp_ule_i16_param_1];484; CHECK-NEXT: setp.le.u16 %p1, %rs1, %rs2;485; CHECK-NEXT: selp.b32 %r1, 1, 0, %p1;486; CHECK-NEXT: st.param.b32 [func_retval0], %r1;487; CHECK-NEXT: ret;488 %cmp = icmp ule i16 %a, %b489 %ret = zext i1 %cmp to i16490 ret i16 %ret491}492 493define i16 @icmp_sgt_i16(i16 %a, i16 %b) {494; CHECK-LABEL: icmp_sgt_i16(495; CHECK: {496; CHECK-NEXT: .reg .pred %p<2>;497; CHECK-NEXT: .reg .b16 %rs<3>;498; CHECK-NEXT: .reg .b32 %r<2>;499; CHECK-EMPTY:500; CHECK-NEXT: // %bb.0:501; CHECK-NEXT: ld.param.b16 %rs1, [icmp_sgt_i16_param_0];502; CHECK-NEXT: ld.param.b16 %rs2, [icmp_sgt_i16_param_1];503; CHECK-NEXT: setp.gt.s16 %p1, %rs1, %rs2;504; CHECK-NEXT: selp.b32 %r1, 1, 0, %p1;505; CHECK-NEXT: st.param.b32 [func_retval0], %r1;506; CHECK-NEXT: ret;507 %cmp = icmp sgt i16 %a, %b508 %ret = zext i1 %cmp to i16509 ret i16 %ret510}511 512define i16 @icmp_sge_i16(i16 %a, i16 %b) {513; CHECK-LABEL: icmp_sge_i16(514; CHECK: {515; CHECK-NEXT: .reg .pred %p<2>;516; CHECK-NEXT: .reg .b16 %rs<3>;517; CHECK-NEXT: .reg .b32 %r<2>;518; CHECK-EMPTY:519; CHECK-NEXT: // %bb.0:520; CHECK-NEXT: ld.param.b16 %rs1, [icmp_sge_i16_param_0];521; CHECK-NEXT: ld.param.b16 %rs2, [icmp_sge_i16_param_1];522; CHECK-NEXT: setp.ge.s16 %p1, %rs1, %rs2;523; CHECK-NEXT: selp.b32 %r1, 1, 0, %p1;524; CHECK-NEXT: st.param.b32 [func_retval0], %r1;525; CHECK-NEXT: ret;526 %cmp = icmp sge i16 %a, %b527 %ret = zext i1 %cmp to i16528 ret i16 %ret529}530 531define i16 @icmp_slt_i16(i16 %a, i16 %b) {532; CHECK-LABEL: icmp_slt_i16(533; CHECK: {534; CHECK-NEXT: .reg .pred %p<2>;535; CHECK-NEXT: .reg .b16 %rs<3>;536; CHECK-NEXT: .reg .b32 %r<2>;537; CHECK-EMPTY:538; CHECK-NEXT: // %bb.0:539; CHECK-NEXT: ld.param.b16 %rs1, [icmp_slt_i16_param_0];540; CHECK-NEXT: ld.param.b16 %rs2, [icmp_slt_i16_param_1];541; CHECK-NEXT: setp.lt.s16 %p1, %rs1, %rs2;542; CHECK-NEXT: selp.b32 %r1, 1, 0, %p1;543; CHECK-NEXT: st.param.b32 [func_retval0], %r1;544; CHECK-NEXT: ret;545 %cmp = icmp slt i16 %a, %b546 %ret = zext i1 %cmp to i16547 ret i16 %ret548}549 550define i16 @icmp_sle_i16(i16 %a, i16 %b) {551; CHECK-LABEL: icmp_sle_i16(552; CHECK: {553; CHECK-NEXT: .reg .pred %p<2>;554; CHECK-NEXT: .reg .b16 %rs<3>;555; CHECK-NEXT: .reg .b32 %r<2>;556; CHECK-EMPTY:557; CHECK-NEXT: // %bb.0:558; CHECK-NEXT: ld.param.b16 %rs1, [icmp_sle_i16_param_0];559; CHECK-NEXT: ld.param.b16 %rs2, [icmp_sle_i16_param_1];560; CHECK-NEXT: setp.le.s16 %p1, %rs1, %rs2;561; CHECK-NEXT: selp.b32 %r1, 1, 0, %p1;562; CHECK-NEXT: st.param.b32 [func_retval0], %r1;563; CHECK-NEXT: ret;564 %cmp = icmp sle i16 %a, %b565 %ret = zext i1 %cmp to i16566 ret i16 %ret567}568 569 570;;; i8571 572define i8 @icmp_eq_i8(i8 %a, i8 %b) {573; Comparison happens in 16-bit574; CHECK-LABEL: icmp_eq_i8(575; CHECK: {576; CHECK-NEXT: .reg .pred %p<2>;577; CHECK-NEXT: .reg .b16 %rs<3>;578; CHECK-NEXT: .reg .b32 %r<2>;579; CHECK-EMPTY:580; CHECK-NEXT: // %bb.0:581; CHECK-NEXT: ld.param.b8 %rs1, [icmp_eq_i8_param_0];582; CHECK-NEXT: ld.param.b8 %rs2, [icmp_eq_i8_param_1];583; CHECK-NEXT: setp.eq.b16 %p1, %rs1, %rs2;584; CHECK-NEXT: selp.b32 %r1, 1, 0, %p1;585; CHECK-NEXT: st.param.b32 [func_retval0], %r1;586; CHECK-NEXT: ret;587 %cmp = icmp eq i8 %a, %b588 %ret = zext i1 %cmp to i8589 ret i8 %ret590}591 592define i8 @icmp_ne_i8(i8 %a, i8 %b) {593; Comparison happens in 16-bit594; CHECK-LABEL: icmp_ne_i8(595; CHECK: {596; CHECK-NEXT: .reg .pred %p<2>;597; CHECK-NEXT: .reg .b16 %rs<3>;598; CHECK-NEXT: .reg .b32 %r<2>;599; CHECK-EMPTY:600; CHECK-NEXT: // %bb.0:601; CHECK-NEXT: ld.param.b8 %rs1, [icmp_ne_i8_param_0];602; CHECK-NEXT: ld.param.b8 %rs2, [icmp_ne_i8_param_1];603; CHECK-NEXT: setp.ne.b16 %p1, %rs1, %rs2;604; CHECK-NEXT: selp.b32 %r1, 1, 0, %p1;605; CHECK-NEXT: st.param.b32 [func_retval0], %r1;606; CHECK-NEXT: ret;607 %cmp = icmp ne i8 %a, %b608 %ret = zext i1 %cmp to i8609 ret i8 %ret610}611 612define i8 @icmp_ugt_i8(i8 %a, i8 %b) {613; Comparison happens in 16-bit614; CHECK-LABEL: icmp_ugt_i8(615; CHECK: {616; CHECK-NEXT: .reg .pred %p<2>;617; CHECK-NEXT: .reg .b16 %rs<3>;618; CHECK-NEXT: .reg .b32 %r<2>;619; CHECK-EMPTY:620; CHECK-NEXT: // %bb.0:621; CHECK-NEXT: ld.param.b8 %rs1, [icmp_ugt_i8_param_0];622; CHECK-NEXT: ld.param.b8 %rs2, [icmp_ugt_i8_param_1];623; CHECK-NEXT: setp.gt.u16 %p1, %rs1, %rs2;624; CHECK-NEXT: selp.b32 %r1, 1, 0, %p1;625; CHECK-NEXT: st.param.b32 [func_retval0], %r1;626; CHECK-NEXT: ret;627 %cmp = icmp ugt i8 %a, %b628 %ret = zext i1 %cmp to i8629 ret i8 %ret630}631 632define i8 @icmp_uge_i8(i8 %a, i8 %b) {633; Comparison happens in 16-bit634; CHECK-LABEL: icmp_uge_i8(635; CHECK: {636; CHECK-NEXT: .reg .pred %p<2>;637; CHECK-NEXT: .reg .b16 %rs<3>;638; CHECK-NEXT: .reg .b32 %r<2>;639; CHECK-EMPTY:640; CHECK-NEXT: // %bb.0:641; CHECK-NEXT: ld.param.b8 %rs1, [icmp_uge_i8_param_0];642; CHECK-NEXT: ld.param.b8 %rs2, [icmp_uge_i8_param_1];643; CHECK-NEXT: setp.ge.u16 %p1, %rs1, %rs2;644; CHECK-NEXT: selp.b32 %r1, 1, 0, %p1;645; CHECK-NEXT: st.param.b32 [func_retval0], %r1;646; CHECK-NEXT: ret;647 %cmp = icmp uge i8 %a, %b648 %ret = zext i1 %cmp to i8649 ret i8 %ret650}651 652define i8 @icmp_ult_i8(i8 %a, i8 %b) {653; Comparison happens in 16-bit654; CHECK-LABEL: icmp_ult_i8(655; CHECK: {656; CHECK-NEXT: .reg .pred %p<2>;657; CHECK-NEXT: .reg .b16 %rs<3>;658; CHECK-NEXT: .reg .b32 %r<2>;659; CHECK-EMPTY:660; CHECK-NEXT: // %bb.0:661; CHECK-NEXT: ld.param.b8 %rs1, [icmp_ult_i8_param_0];662; CHECK-NEXT: ld.param.b8 %rs2, [icmp_ult_i8_param_1];663; CHECK-NEXT: setp.lt.u16 %p1, %rs1, %rs2;664; CHECK-NEXT: selp.b32 %r1, 1, 0, %p1;665; CHECK-NEXT: st.param.b32 [func_retval0], %r1;666; CHECK-NEXT: ret;667 %cmp = icmp ult i8 %a, %b668 %ret = zext i1 %cmp to i8669 ret i8 %ret670}671 672define i8 @icmp_ule_i8(i8 %a, i8 %b) {673; Comparison happens in 16-bit674; CHECK-LABEL: icmp_ule_i8(675; CHECK: {676; CHECK-NEXT: .reg .pred %p<2>;677; CHECK-NEXT: .reg .b16 %rs<3>;678; CHECK-NEXT: .reg .b32 %r<2>;679; CHECK-EMPTY:680; CHECK-NEXT: // %bb.0:681; CHECK-NEXT: ld.param.b8 %rs1, [icmp_ule_i8_param_0];682; CHECK-NEXT: ld.param.b8 %rs2, [icmp_ule_i8_param_1];683; CHECK-NEXT: setp.le.u16 %p1, %rs1, %rs2;684; CHECK-NEXT: selp.b32 %r1, 1, 0, %p1;685; CHECK-NEXT: st.param.b32 [func_retval0], %r1;686; CHECK-NEXT: ret;687 %cmp = icmp ule i8 %a, %b688 %ret = zext i1 %cmp to i8689 ret i8 %ret690}691 692define i8 @icmp_sgt_i8(i8 %a, i8 %b) {693; Comparison happens in 16-bit694; CHECK-LABEL: icmp_sgt_i8(695; CHECK: {696; CHECK-NEXT: .reg .pred %p<2>;697; CHECK-NEXT: .reg .b16 %rs<3>;698; CHECK-NEXT: .reg .b32 %r<2>;699; CHECK-EMPTY:700; CHECK-NEXT: // %bb.0:701; CHECK-NEXT: ld.param.s8 %rs1, [icmp_sgt_i8_param_0];702; CHECK-NEXT: ld.param.s8 %rs2, [icmp_sgt_i8_param_1];703; CHECK-NEXT: setp.gt.s16 %p1, %rs1, %rs2;704; CHECK-NEXT: selp.b32 %r1, 1, 0, %p1;705; CHECK-NEXT: st.param.b32 [func_retval0], %r1;706; CHECK-NEXT: ret;707 %cmp = icmp sgt i8 %a, %b708 %ret = zext i1 %cmp to i8709 ret i8 %ret710}711 712define i8 @icmp_sge_i8(i8 %a, i8 %b) {713; Comparison happens in 16-bit714; CHECK-LABEL: icmp_sge_i8(715; CHECK: {716; CHECK-NEXT: .reg .pred %p<2>;717; CHECK-NEXT: .reg .b16 %rs<3>;718; CHECK-NEXT: .reg .b32 %r<2>;719; CHECK-EMPTY:720; CHECK-NEXT: // %bb.0:721; CHECK-NEXT: ld.param.s8 %rs1, [icmp_sge_i8_param_0];722; CHECK-NEXT: ld.param.s8 %rs2, [icmp_sge_i8_param_1];723; CHECK-NEXT: setp.ge.s16 %p1, %rs1, %rs2;724; CHECK-NEXT: selp.b32 %r1, 1, 0, %p1;725; CHECK-NEXT: st.param.b32 [func_retval0], %r1;726; CHECK-NEXT: ret;727 %cmp = icmp sge i8 %a, %b728 %ret = zext i1 %cmp to i8729 ret i8 %ret730}731 732define i8 @icmp_slt_i8(i8 %a, i8 %b) {733; Comparison happens in 16-bit734; CHECK-LABEL: icmp_slt_i8(735; CHECK: {736; CHECK-NEXT: .reg .pred %p<2>;737; CHECK-NEXT: .reg .b16 %rs<3>;738; CHECK-NEXT: .reg .b32 %r<2>;739; CHECK-EMPTY:740; CHECK-NEXT: // %bb.0:741; CHECK-NEXT: ld.param.s8 %rs1, [icmp_slt_i8_param_0];742; CHECK-NEXT: ld.param.s8 %rs2, [icmp_slt_i8_param_1];743; CHECK-NEXT: setp.lt.s16 %p1, %rs1, %rs2;744; CHECK-NEXT: selp.b32 %r1, 1, 0, %p1;745; CHECK-NEXT: st.param.b32 [func_retval0], %r1;746; CHECK-NEXT: ret;747 %cmp = icmp slt i8 %a, %b748 %ret = zext i1 %cmp to i8749 ret i8 %ret750}751 752define i8 @icmp_sle_i8(i8 %a, i8 %b) {753; Comparison happens in 16-bit754; CHECK-LABEL: icmp_sle_i8(755; CHECK: {756; CHECK-NEXT: .reg .pred %p<2>;757; CHECK-NEXT: .reg .b16 %rs<3>;758; CHECK-NEXT: .reg .b32 %r<2>;759; CHECK-EMPTY:760; CHECK-NEXT: // %bb.0:761; CHECK-NEXT: ld.param.s8 %rs1, [icmp_sle_i8_param_0];762; CHECK-NEXT: ld.param.s8 %rs2, [icmp_sle_i8_param_1];763; CHECK-NEXT: setp.le.s16 %p1, %rs1, %rs2;764; CHECK-NEXT: selp.b32 %r1, 1, 0, %p1;765; CHECK-NEXT: st.param.b32 [func_retval0], %r1;766; CHECK-NEXT: ret;767 %cmp = icmp sle i8 %a, %b768 %ret = zext i1 %cmp to i8769 ret i8 %ret770}771