1034 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mcpu=sm_90 -mattr=+ptx80 -O3 | FileCheck %s --check-prefixes=CHECK,SM903; RUN: llc < %s -mcpu=sm_20 -O3 | FileCheck %s --check-prefixes=CHECK,SM204; RUN: %if ptxas-sm_90 && ptxas-isa-8.0 %{ llc < %s -mcpu=sm_90 -mattr=+ptx80 -O3 | %ptxas-verify -arch=sm_90 %}5; RUN: %if ptxas %{ llc < %s -mcpu=sm_20 -O3 | %ptxas-verify %}6 7target triple = "nvptx64-nvidia-cuda"8 9; *************************************10; * Cases with no min/max11 12define i32 @ab_eq_i32(i32 %a, i32 %b) {13; CHECK-LABEL: ab_eq_i32(14; CHECK: {15; CHECK-NEXT: .reg .b32 %r<2>;16; CHECK-EMPTY:17; CHECK-NEXT: // %bb.0:18; CHECK-NEXT: ld.param.b32 %r1, [ab_eq_i32_param_1];19; CHECK-NEXT: st.param.b32 [func_retval0], %r1;20; CHECK-NEXT: ret;21 %cmp = icmp eq i32 %a, %b22 %sel = select i1 %cmp, i32 %a, i32 %b23 ret i32 %sel24}25 26define i64 @ab_ne_i64(i64 %a, i64 %b) {27; CHECK-LABEL: ab_ne_i64(28; CHECK: {29; CHECK-NEXT: .reg .b64 %rd<2>;30; CHECK-EMPTY:31; CHECK-NEXT: // %bb.0:32; CHECK-NEXT: ld.param.b64 %rd1, [ab_ne_i64_param_1];33; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;34; CHECK-NEXT: ret;35 %cmp = icmp ne i64 %a, %b36 %sel = select i1 %cmp, i64 %b, i64 %a37 ret i64 %sel38}39 40; *************************************41; * All variations with i1642 43; *** ab, unsigned, i1644define i16 @ab_ugt_i16(i16 %a, i16 %b) {45; CHECK-LABEL: ab_ugt_i16(46; CHECK: {47; CHECK-NEXT: .reg .b16 %rs<4>;48; CHECK-NEXT: .reg .b32 %r<2>;49; CHECK-EMPTY:50; CHECK-NEXT: // %bb.0:51; CHECK-NEXT: ld.param.b16 %rs1, [ab_ugt_i16_param_0];52; CHECK-NEXT: ld.param.b16 %rs2, [ab_ugt_i16_param_1];53; CHECK-NEXT: max.u16 %rs3, %rs1, %rs2;54; CHECK-NEXT: cvt.u32.u16 %r1, %rs3;55; CHECK-NEXT: st.param.b32 [func_retval0], %r1;56; CHECK-NEXT: ret;57 %cmp = icmp ugt i16 %a, %b58 %sel = select i1 %cmp, i16 %a, i16 %b59 ret i16 %sel60}61 62define i16 @ab_uge_i16(i16 %a, i16 %b) {63; CHECK-LABEL: ab_uge_i16(64; CHECK: {65; CHECK-NEXT: .reg .b16 %rs<4>;66; CHECK-NEXT: .reg .b32 %r<2>;67; CHECK-EMPTY:68; CHECK-NEXT: // %bb.0:69; CHECK-NEXT: ld.param.b16 %rs1, [ab_uge_i16_param_0];70; CHECK-NEXT: ld.param.b16 %rs2, [ab_uge_i16_param_1];71; CHECK-NEXT: max.u16 %rs3, %rs1, %rs2;72; CHECK-NEXT: cvt.u32.u16 %r1, %rs3;73; CHECK-NEXT: st.param.b32 [func_retval0], %r1;74; CHECK-NEXT: ret;75 %cmp = icmp uge i16 %a, %b76 %sel = select i1 %cmp, i16 %a, i16 %b77 ret i16 %sel78}79 80define i16 @ab_ult_i16(i16 %a, i16 %b) {81; CHECK-LABEL: ab_ult_i16(82; CHECK: {83; CHECK-NEXT: .reg .b16 %rs<4>;84; CHECK-NEXT: .reg .b32 %r<2>;85; CHECK-EMPTY:86; CHECK-NEXT: // %bb.0:87; CHECK-NEXT: ld.param.b16 %rs1, [ab_ult_i16_param_0];88; CHECK-NEXT: ld.param.b16 %rs2, [ab_ult_i16_param_1];89; CHECK-NEXT: min.u16 %rs3, %rs1, %rs2;90; CHECK-NEXT: cvt.u32.u16 %r1, %rs3;91; CHECK-NEXT: st.param.b32 [func_retval0], %r1;92; CHECK-NEXT: ret;93 %cmp = icmp ult i16 %a, %b94 %sel = select i1 %cmp, i16 %a, i16 %b95 ret i16 %sel96}97 98define i16 @ab_ule_i16(i16 %a, i16 %b) {99; CHECK-LABEL: ab_ule_i16(100; CHECK: {101; CHECK-NEXT: .reg .b16 %rs<4>;102; CHECK-NEXT: .reg .b32 %r<2>;103; CHECK-EMPTY:104; CHECK-NEXT: // %bb.0:105; CHECK-NEXT: ld.param.b16 %rs1, [ab_ule_i16_param_0];106; CHECK-NEXT: ld.param.b16 %rs2, [ab_ule_i16_param_1];107; CHECK-NEXT: min.u16 %rs3, %rs1, %rs2;108; CHECK-NEXT: cvt.u32.u16 %r1, %rs3;109; CHECK-NEXT: st.param.b32 [func_retval0], %r1;110; CHECK-NEXT: ret;111 %cmp = icmp ule i16 %a, %b112 %sel = select i1 %cmp, i16 %a, i16 %b113 ret i16 %sel114}115 116; *** ab, signed, i16117define i16 @ab_sgt_i16(i16 %a, i16 %b) {118; CHECK-LABEL: ab_sgt_i16(119; CHECK: {120; CHECK-NEXT: .reg .b16 %rs<4>;121; CHECK-NEXT: .reg .b32 %r<2>;122; CHECK-EMPTY:123; CHECK-NEXT: // %bb.0:124; CHECK-NEXT: ld.param.b16 %rs1, [ab_sgt_i16_param_0];125; CHECK-NEXT: ld.param.b16 %rs2, [ab_sgt_i16_param_1];126; CHECK-NEXT: max.s16 %rs3, %rs1, %rs2;127; CHECK-NEXT: cvt.u32.u16 %r1, %rs3;128; CHECK-NEXT: st.param.b32 [func_retval0], %r1;129; CHECK-NEXT: ret;130 %cmp = icmp sgt i16 %a, %b131 %sel = select i1 %cmp, i16 %a, i16 %b132 ret i16 %sel133}134 135define i16 @ab_sge_i16(i16 %a, i16 %b) {136; CHECK-LABEL: ab_sge_i16(137; CHECK: {138; CHECK-NEXT: .reg .b16 %rs<4>;139; CHECK-NEXT: .reg .b32 %r<2>;140; CHECK-EMPTY:141; CHECK-NEXT: // %bb.0:142; CHECK-NEXT: ld.param.b16 %rs1, [ab_sge_i16_param_0];143; CHECK-NEXT: ld.param.b16 %rs2, [ab_sge_i16_param_1];144; CHECK-NEXT: max.s16 %rs3, %rs1, %rs2;145; CHECK-NEXT: cvt.u32.u16 %r1, %rs3;146; CHECK-NEXT: st.param.b32 [func_retval0], %r1;147; CHECK-NEXT: ret;148 %cmp = icmp sge i16 %a, %b149 %sel = select i1 %cmp, i16 %a, i16 %b150 ret i16 %sel151}152 153define i16 @ab_slt_i16(i16 %a, i16 %b) {154; CHECK-LABEL: ab_slt_i16(155; CHECK: {156; CHECK-NEXT: .reg .b16 %rs<4>;157; CHECK-NEXT: .reg .b32 %r<2>;158; CHECK-EMPTY:159; CHECK-NEXT: // %bb.0:160; CHECK-NEXT: ld.param.b16 %rs1, [ab_slt_i16_param_0];161; CHECK-NEXT: ld.param.b16 %rs2, [ab_slt_i16_param_1];162; CHECK-NEXT: min.s16 %rs3, %rs1, %rs2;163; CHECK-NEXT: cvt.u32.u16 %r1, %rs3;164; CHECK-NEXT: st.param.b32 [func_retval0], %r1;165; CHECK-NEXT: ret;166 %cmp = icmp slt i16 %a, %b167 %sel = select i1 %cmp, i16 %a, i16 %b168 ret i16 %sel169}170 171define i16 @ab_sle_i16(i16 %a, i16 %b) {172; CHECK-LABEL: ab_sle_i16(173; CHECK: {174; CHECK-NEXT: .reg .b16 %rs<4>;175; CHECK-NEXT: .reg .b32 %r<2>;176; CHECK-EMPTY:177; CHECK-NEXT: // %bb.0:178; CHECK-NEXT: ld.param.b16 %rs1, [ab_sle_i16_param_0];179; CHECK-NEXT: ld.param.b16 %rs2, [ab_sle_i16_param_1];180; CHECK-NEXT: min.s16 %rs3, %rs1, %rs2;181; CHECK-NEXT: cvt.u32.u16 %r1, %rs3;182; CHECK-NEXT: st.param.b32 [func_retval0], %r1;183; CHECK-NEXT: ret;184 %cmp = icmp sle i16 %a, %b185 %sel = select i1 %cmp, i16 %a, i16 %b186 ret i16 %sel187}188 189; *** ba, unsigned, i16190define i16 @ba_ugt_i16(i16 %a, i16 %b) {191; CHECK-LABEL: ba_ugt_i16(192; CHECK: {193; CHECK-NEXT: .reg .b16 %rs<4>;194; CHECK-NEXT: .reg .b32 %r<2>;195; CHECK-EMPTY:196; CHECK-NEXT: // %bb.0:197; CHECK-NEXT: ld.param.b16 %rs1, [ba_ugt_i16_param_0];198; CHECK-NEXT: ld.param.b16 %rs2, [ba_ugt_i16_param_1];199; CHECK-NEXT: min.u16 %rs3, %rs1, %rs2;200; CHECK-NEXT: cvt.u32.u16 %r1, %rs3;201; CHECK-NEXT: st.param.b32 [func_retval0], %r1;202; CHECK-NEXT: ret;203 %cmp = icmp ugt i16 %a, %b204 %sel = select i1 %cmp, i16 %b, i16 %a205 ret i16 %sel206}207 208define i16 @ba_uge_i16(i16 %a, i16 %b) {209; CHECK-LABEL: ba_uge_i16(210; CHECK: {211; CHECK-NEXT: .reg .b16 %rs<4>;212; CHECK-NEXT: .reg .b32 %r<2>;213; CHECK-EMPTY:214; CHECK-NEXT: // %bb.0:215; CHECK-NEXT: ld.param.b16 %rs1, [ba_uge_i16_param_0];216; CHECK-NEXT: ld.param.b16 %rs2, [ba_uge_i16_param_1];217; CHECK-NEXT: min.u16 %rs3, %rs1, %rs2;218; CHECK-NEXT: cvt.u32.u16 %r1, %rs3;219; CHECK-NEXT: st.param.b32 [func_retval0], %r1;220; CHECK-NEXT: ret;221 %cmp = icmp uge i16 %a, %b222 %sel = select i1 %cmp, i16 %b, i16 %a223 ret i16 %sel224}225 226define i16 @ba_ult_i16(i16 %a, i16 %b) {227; CHECK-LABEL: ba_ult_i16(228; CHECK: {229; CHECK-NEXT: .reg .b16 %rs<4>;230; CHECK-NEXT: .reg .b32 %r<2>;231; CHECK-EMPTY:232; CHECK-NEXT: // %bb.0:233; CHECK-NEXT: ld.param.b16 %rs1, [ba_ult_i16_param_0];234; CHECK-NEXT: ld.param.b16 %rs2, [ba_ult_i16_param_1];235; CHECK-NEXT: max.u16 %rs3, %rs1, %rs2;236; CHECK-NEXT: cvt.u32.u16 %r1, %rs3;237; CHECK-NEXT: st.param.b32 [func_retval0], %r1;238; CHECK-NEXT: ret;239 %cmp = icmp ult i16 %a, %b240 %sel = select i1 %cmp, i16 %b, i16 %a241 ret i16 %sel242}243 244define i16 @ba_ule_i16(i16 %a, i16 %b) {245; CHECK-LABEL: ba_ule_i16(246; CHECK: {247; CHECK-NEXT: .reg .b16 %rs<4>;248; CHECK-NEXT: .reg .b32 %r<2>;249; CHECK-EMPTY:250; CHECK-NEXT: // %bb.0:251; CHECK-NEXT: ld.param.b16 %rs1, [ba_ule_i16_param_0];252; CHECK-NEXT: ld.param.b16 %rs2, [ba_ule_i16_param_1];253; CHECK-NEXT: max.u16 %rs3, %rs1, %rs2;254; CHECK-NEXT: cvt.u32.u16 %r1, %rs3;255; CHECK-NEXT: st.param.b32 [func_retval0], %r1;256; CHECK-NEXT: ret;257 %cmp = icmp ule i16 %a, %b258 %sel = select i1 %cmp, i16 %b, i16 %a259 ret i16 %sel260}261 262; *** ba, signed, i16263define i16 @ba_sgt_i16(i16 %a, i16 %b) {264; CHECK-LABEL: ba_sgt_i16(265; CHECK: {266; CHECK-NEXT: .reg .b16 %rs<4>;267; CHECK-NEXT: .reg .b32 %r<2>;268; CHECK-EMPTY:269; CHECK-NEXT: // %bb.0:270; CHECK-NEXT: ld.param.b16 %rs1, [ba_sgt_i16_param_0];271; CHECK-NEXT: ld.param.b16 %rs2, [ba_sgt_i16_param_1];272; CHECK-NEXT: min.s16 %rs3, %rs1, %rs2;273; CHECK-NEXT: cvt.u32.u16 %r1, %rs3;274; CHECK-NEXT: st.param.b32 [func_retval0], %r1;275; CHECK-NEXT: ret;276 %cmp = icmp sgt i16 %a, %b277 %sel = select i1 %cmp, i16 %b, i16 %a278 ret i16 %sel279}280 281define i16 @ba_sge_i16(i16 %a, i16 %b) {282; CHECK-LABEL: ba_sge_i16(283; CHECK: {284; CHECK-NEXT: .reg .b16 %rs<4>;285; CHECK-NEXT: .reg .b32 %r<2>;286; CHECK-EMPTY:287; CHECK-NEXT: // %bb.0:288; CHECK-NEXT: ld.param.b16 %rs1, [ba_sge_i16_param_0];289; CHECK-NEXT: ld.param.b16 %rs2, [ba_sge_i16_param_1];290; CHECK-NEXT: min.s16 %rs3, %rs1, %rs2;291; CHECK-NEXT: cvt.u32.u16 %r1, %rs3;292; CHECK-NEXT: st.param.b32 [func_retval0], %r1;293; CHECK-NEXT: ret;294 %cmp = icmp sge i16 %a, %b295 %sel = select i1 %cmp, i16 %b, i16 %a296 ret i16 %sel297}298 299define i16 @ba_slt_i16(i16 %a, i16 %b) {300; CHECK-LABEL: ba_slt_i16(301; CHECK: {302; CHECK-NEXT: .reg .b16 %rs<4>;303; CHECK-NEXT: .reg .b32 %r<2>;304; CHECK-EMPTY:305; CHECK-NEXT: // %bb.0:306; CHECK-NEXT: ld.param.b16 %rs1, [ba_slt_i16_param_0];307; CHECK-NEXT: ld.param.b16 %rs2, [ba_slt_i16_param_1];308; CHECK-NEXT: max.s16 %rs3, %rs1, %rs2;309; CHECK-NEXT: cvt.u32.u16 %r1, %rs3;310; CHECK-NEXT: st.param.b32 [func_retval0], %r1;311; CHECK-NEXT: ret;312 %cmp = icmp slt i16 %a, %b313 %sel = select i1 %cmp, i16 %b, i16 %a314 ret i16 %sel315}316 317define i16 @ba_sle_i16(i16 %a, i16 %b) {318; CHECK-LABEL: ba_sle_i16(319; CHECK: {320; CHECK-NEXT: .reg .b16 %rs<4>;321; CHECK-NEXT: .reg .b32 %r<2>;322; CHECK-EMPTY:323; CHECK-NEXT: // %bb.0:324; CHECK-NEXT: ld.param.b16 %rs1, [ba_sle_i16_param_0];325; CHECK-NEXT: ld.param.b16 %rs2, [ba_sle_i16_param_1];326; CHECK-NEXT: max.s16 %rs3, %rs1, %rs2;327; CHECK-NEXT: cvt.u32.u16 %r1, %rs3;328; CHECK-NEXT: st.param.b32 [func_retval0], %r1;329; CHECK-NEXT: ret;330 %cmp = icmp sle i16 %a, %b331 %sel = select i1 %cmp, i16 %b, i16 %a332 ret i16 %sel333}334 335; *************************************336; * All variations with i32337 338; *** ab, unsigned, i32339define i32 @ab_ugt_i32(i32 %a, i32 %b) {340; CHECK-LABEL: ab_ugt_i32(341; CHECK: {342; CHECK-NEXT: .reg .b32 %r<4>;343; CHECK-EMPTY:344; CHECK-NEXT: // %bb.0:345; CHECK-NEXT: ld.param.b32 %r1, [ab_ugt_i32_param_0];346; CHECK-NEXT: ld.param.b32 %r2, [ab_ugt_i32_param_1];347; CHECK-NEXT: max.u32 %r3, %r1, %r2;348; CHECK-NEXT: st.param.b32 [func_retval0], %r3;349; CHECK-NEXT: ret;350 %cmp = icmp ugt i32 %a, %b351 %sel = select i1 %cmp, i32 %a, i32 %b352 ret i32 %sel353}354 355define i32 @ab_uge_i32(i32 %a, i32 %b) {356; CHECK-LABEL: ab_uge_i32(357; CHECK: {358; CHECK-NEXT: .reg .b32 %r<4>;359; CHECK-EMPTY:360; CHECK-NEXT: // %bb.0:361; CHECK-NEXT: ld.param.b32 %r1, [ab_uge_i32_param_0];362; CHECK-NEXT: ld.param.b32 %r2, [ab_uge_i32_param_1];363; CHECK-NEXT: max.u32 %r3, %r1, %r2;364; CHECK-NEXT: st.param.b32 [func_retval0], %r3;365; CHECK-NEXT: ret;366 %cmp = icmp uge i32 %a, %b367 %sel = select i1 %cmp, i32 %a, i32 %b368 ret i32 %sel369}370 371define i32 @ab_ult_i32(i32 %a, i32 %b) {372; CHECK-LABEL: ab_ult_i32(373; CHECK: {374; CHECK-NEXT: .reg .b32 %r<4>;375; CHECK-EMPTY:376; CHECK-NEXT: // %bb.0:377; CHECK-NEXT: ld.param.b32 %r1, [ab_ult_i32_param_0];378; CHECK-NEXT: ld.param.b32 %r2, [ab_ult_i32_param_1];379; CHECK-NEXT: min.u32 %r3, %r1, %r2;380; CHECK-NEXT: st.param.b32 [func_retval0], %r3;381; CHECK-NEXT: ret;382 %cmp = icmp ult i32 %a, %b383 %sel = select i1 %cmp, i32 %a, i32 %b384 ret i32 %sel385}386 387define i32 @ab_ule_i32(i32 %a, i32 %b) {388; CHECK-LABEL: ab_ule_i32(389; CHECK: {390; CHECK-NEXT: .reg .b32 %r<4>;391; CHECK-EMPTY:392; CHECK-NEXT: // %bb.0:393; CHECK-NEXT: ld.param.b32 %r1, [ab_ule_i32_param_0];394; CHECK-NEXT: ld.param.b32 %r2, [ab_ule_i32_param_1];395; CHECK-NEXT: min.u32 %r3, %r1, %r2;396; CHECK-NEXT: st.param.b32 [func_retval0], %r3;397; CHECK-NEXT: ret;398 %cmp = icmp ule i32 %a, %b399 %sel = select i1 %cmp, i32 %a, i32 %b400 ret i32 %sel401}402 403; *** ab, signed, i32404define i32 @ab_sgt_i32(i32 %a, i32 %b) {405; CHECK-LABEL: ab_sgt_i32(406; CHECK: {407; CHECK-NEXT: .reg .b32 %r<4>;408; CHECK-EMPTY:409; CHECK-NEXT: // %bb.0:410; CHECK-NEXT: ld.param.b32 %r1, [ab_sgt_i32_param_0];411; CHECK-NEXT: ld.param.b32 %r2, [ab_sgt_i32_param_1];412; CHECK-NEXT: max.s32 %r3, %r1, %r2;413; CHECK-NEXT: st.param.b32 [func_retval0], %r3;414; CHECK-NEXT: ret;415 %cmp = icmp sgt i32 %a, %b416 %sel = select i1 %cmp, i32 %a, i32 %b417 ret i32 %sel418}419 420define i32 @ab_sge_i32(i32 %a, i32 %b) {421; CHECK-LABEL: ab_sge_i32(422; CHECK: {423; CHECK-NEXT: .reg .b32 %r<4>;424; CHECK-EMPTY:425; CHECK-NEXT: // %bb.0:426; CHECK-NEXT: ld.param.b32 %r1, [ab_sge_i32_param_0];427; CHECK-NEXT: ld.param.b32 %r2, [ab_sge_i32_param_1];428; CHECK-NEXT: max.s32 %r3, %r1, %r2;429; CHECK-NEXT: st.param.b32 [func_retval0], %r3;430; CHECK-NEXT: ret;431 %cmp = icmp sge i32 %a, %b432 %sel = select i1 %cmp, i32 %a, i32 %b433 ret i32 %sel434}435 436define i32 @ab_slt_i32(i32 %a, i32 %b) {437; CHECK-LABEL: ab_slt_i32(438; CHECK: {439; CHECK-NEXT: .reg .b32 %r<4>;440; CHECK-EMPTY:441; CHECK-NEXT: // %bb.0:442; CHECK-NEXT: ld.param.b32 %r1, [ab_slt_i32_param_0];443; CHECK-NEXT: ld.param.b32 %r2, [ab_slt_i32_param_1];444; CHECK-NEXT: min.s32 %r3, %r1, %r2;445; CHECK-NEXT: st.param.b32 [func_retval0], %r3;446; CHECK-NEXT: ret;447 %cmp = icmp slt i32 %a, %b448 %sel = select i1 %cmp, i32 %a, i32 %b449 ret i32 %sel450}451 452define i32 @ab_sle_i32(i32 %a, i32 %b) {453; CHECK-LABEL: ab_sle_i32(454; CHECK: {455; CHECK-NEXT: .reg .b32 %r<4>;456; CHECK-EMPTY:457; CHECK-NEXT: // %bb.0:458; CHECK-NEXT: ld.param.b32 %r1, [ab_sle_i32_param_0];459; CHECK-NEXT: ld.param.b32 %r2, [ab_sle_i32_param_1];460; CHECK-NEXT: min.s32 %r3, %r1, %r2;461; CHECK-NEXT: st.param.b32 [func_retval0], %r3;462; CHECK-NEXT: ret;463 %cmp = icmp sle i32 %a, %b464 %sel = select i1 %cmp, i32 %a, i32 %b465 ret i32 %sel466}467 468; *** ba, unsigned, i32469define i32 @ba_ugt_i32(i32 %a, i32 %b) {470; CHECK-LABEL: ba_ugt_i32(471; CHECK: {472; CHECK-NEXT: .reg .b32 %r<4>;473; CHECK-EMPTY:474; CHECK-NEXT: // %bb.0:475; CHECK-NEXT: ld.param.b32 %r1, [ba_ugt_i32_param_0];476; CHECK-NEXT: ld.param.b32 %r2, [ba_ugt_i32_param_1];477; CHECK-NEXT: min.u32 %r3, %r1, %r2;478; CHECK-NEXT: st.param.b32 [func_retval0], %r3;479; CHECK-NEXT: ret;480 %cmp = icmp ugt i32 %a, %b481 %sel = select i1 %cmp, i32 %b, i32 %a482 ret i32 %sel483}484 485define i32 @ba_uge_i32(i32 %a, i32 %b) {486; CHECK-LABEL: ba_uge_i32(487; CHECK: {488; CHECK-NEXT: .reg .b32 %r<4>;489; CHECK-EMPTY:490; CHECK-NEXT: // %bb.0:491; CHECK-NEXT: ld.param.b32 %r1, [ba_uge_i32_param_0];492; CHECK-NEXT: ld.param.b32 %r2, [ba_uge_i32_param_1];493; CHECK-NEXT: min.u32 %r3, %r1, %r2;494; CHECK-NEXT: st.param.b32 [func_retval0], %r3;495; CHECK-NEXT: ret;496 %cmp = icmp uge i32 %a, %b497 %sel = select i1 %cmp, i32 %b, i32 %a498 ret i32 %sel499}500 501define i32 @ba_ult_i32(i32 %a, i32 %b) {502; CHECK-LABEL: ba_ult_i32(503; CHECK: {504; CHECK-NEXT: .reg .b32 %r<4>;505; CHECK-EMPTY:506; CHECK-NEXT: // %bb.0:507; CHECK-NEXT: ld.param.b32 %r1, [ba_ult_i32_param_0];508; CHECK-NEXT: ld.param.b32 %r2, [ba_ult_i32_param_1];509; CHECK-NEXT: max.u32 %r3, %r1, %r2;510; CHECK-NEXT: st.param.b32 [func_retval0], %r3;511; CHECK-NEXT: ret;512 %cmp = icmp ult i32 %a, %b513 %sel = select i1 %cmp, i32 %b, i32 %a514 ret i32 %sel515}516 517define i32 @ba_ule_i32(i32 %a, i32 %b) {518; CHECK-LABEL: ba_ule_i32(519; CHECK: {520; CHECK-NEXT: .reg .b32 %r<4>;521; CHECK-EMPTY:522; CHECK-NEXT: // %bb.0:523; CHECK-NEXT: ld.param.b32 %r1, [ba_ule_i32_param_0];524; CHECK-NEXT: ld.param.b32 %r2, [ba_ule_i32_param_1];525; CHECK-NEXT: max.u32 %r3, %r1, %r2;526; CHECK-NEXT: st.param.b32 [func_retval0], %r3;527; CHECK-NEXT: ret;528 %cmp = icmp ule i32 %a, %b529 %sel = select i1 %cmp, i32 %b, i32 %a530 ret i32 %sel531}532 533; *** ba, signed, i32534define i32 @ba_sgt_i32(i32 %a, i32 %b) {535; CHECK-LABEL: ba_sgt_i32(536; CHECK: {537; CHECK-NEXT: .reg .b32 %r<4>;538; CHECK-EMPTY:539; CHECK-NEXT: // %bb.0:540; CHECK-NEXT: ld.param.b32 %r1, [ba_sgt_i32_param_0];541; CHECK-NEXT: ld.param.b32 %r2, [ba_sgt_i32_param_1];542; CHECK-NEXT: min.s32 %r3, %r1, %r2;543; CHECK-NEXT: st.param.b32 [func_retval0], %r3;544; CHECK-NEXT: ret;545 %cmp = icmp sgt i32 %a, %b546 %sel = select i1 %cmp, i32 %b, i32 %a547 ret i32 %sel548}549 550define i32 @ba_sge_i32(i32 %a, i32 %b) {551; CHECK-LABEL: ba_sge_i32(552; CHECK: {553; CHECK-NEXT: .reg .b32 %r<4>;554; CHECK-EMPTY:555; CHECK-NEXT: // %bb.0:556; CHECK-NEXT: ld.param.b32 %r1, [ba_sge_i32_param_0];557; CHECK-NEXT: ld.param.b32 %r2, [ba_sge_i32_param_1];558; CHECK-NEXT: min.s32 %r3, %r1, %r2;559; CHECK-NEXT: st.param.b32 [func_retval0], %r3;560; CHECK-NEXT: ret;561 %cmp = icmp sge i32 %a, %b562 %sel = select i1 %cmp, i32 %b, i32 %a563 ret i32 %sel564}565 566define i32 @ba_slt_i32(i32 %a, i32 %b) {567; CHECK-LABEL: ba_slt_i32(568; CHECK: {569; CHECK-NEXT: .reg .b32 %r<4>;570; CHECK-EMPTY:571; CHECK-NEXT: // %bb.0:572; CHECK-NEXT: ld.param.b32 %r1, [ba_slt_i32_param_0];573; CHECK-NEXT: ld.param.b32 %r2, [ba_slt_i32_param_1];574; CHECK-NEXT: max.s32 %r3, %r1, %r2;575; CHECK-NEXT: st.param.b32 [func_retval0], %r3;576; CHECK-NEXT: ret;577 %cmp = icmp slt i32 %a, %b578 %sel = select i1 %cmp, i32 %b, i32 %a579 ret i32 %sel580}581 582define i32 @ba_sle_i32(i32 %a, i32 %b) {583; CHECK-LABEL: ba_sle_i32(584; CHECK: {585; CHECK-NEXT: .reg .b32 %r<4>;586; CHECK-EMPTY:587; CHECK-NEXT: // %bb.0:588; CHECK-NEXT: ld.param.b32 %r1, [ba_sle_i32_param_0];589; CHECK-NEXT: ld.param.b32 %r2, [ba_sle_i32_param_1];590; CHECK-NEXT: max.s32 %r3, %r1, %r2;591; CHECK-NEXT: st.param.b32 [func_retval0], %r3;592; CHECK-NEXT: ret;593 %cmp = icmp sle i32 %a, %b594 %sel = select i1 %cmp, i32 %b, i32 %a595 ret i32 %sel596}597 598; *************************************599; * All variations with i64600 601; *** ab, unsigned, i64602define i64 @ab_ugt_i64(i64 %a, i64 %b) {603; CHECK-LABEL: ab_ugt_i64(604; CHECK: {605; CHECK-NEXT: .reg .b64 %rd<4>;606; CHECK-EMPTY:607; CHECK-NEXT: // %bb.0:608; CHECK-NEXT: ld.param.b64 %rd1, [ab_ugt_i64_param_0];609; CHECK-NEXT: ld.param.b64 %rd2, [ab_ugt_i64_param_1];610; CHECK-NEXT: max.u64 %rd3, %rd1, %rd2;611; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;612; CHECK-NEXT: ret;613 %cmp = icmp ugt i64 %a, %b614 %sel = select i1 %cmp, i64 %a, i64 %b615 ret i64 %sel616}617 618define i64 @ab_uge_i64(i64 %a, i64 %b) {619; CHECK-LABEL: ab_uge_i64(620; CHECK: {621; CHECK-NEXT: .reg .b64 %rd<4>;622; CHECK-EMPTY:623; CHECK-NEXT: // %bb.0:624; CHECK-NEXT: ld.param.b64 %rd1, [ab_uge_i64_param_0];625; CHECK-NEXT: ld.param.b64 %rd2, [ab_uge_i64_param_1];626; CHECK-NEXT: max.u64 %rd3, %rd1, %rd2;627; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;628; CHECK-NEXT: ret;629 %cmp = icmp uge i64 %a, %b630 %sel = select i1 %cmp, i64 %a, i64 %b631 ret i64 %sel632}633 634define i64 @ab_ult_i64(i64 %a, i64 %b) {635; CHECK-LABEL: ab_ult_i64(636; CHECK: {637; CHECK-NEXT: .reg .b64 %rd<4>;638; CHECK-EMPTY:639; CHECK-NEXT: // %bb.0:640; CHECK-NEXT: ld.param.b64 %rd1, [ab_ult_i64_param_0];641; CHECK-NEXT: ld.param.b64 %rd2, [ab_ult_i64_param_1];642; CHECK-NEXT: min.u64 %rd3, %rd1, %rd2;643; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;644; CHECK-NEXT: ret;645 %cmp = icmp ult i64 %a, %b646 %sel = select i1 %cmp, i64 %a, i64 %b647 ret i64 %sel648}649 650define i64 @ab_ule_i64(i64 %a, i64 %b) {651; CHECK-LABEL: ab_ule_i64(652; CHECK: {653; CHECK-NEXT: .reg .b64 %rd<4>;654; CHECK-EMPTY:655; CHECK-NEXT: // %bb.0:656; CHECK-NEXT: ld.param.b64 %rd1, [ab_ule_i64_param_0];657; CHECK-NEXT: ld.param.b64 %rd2, [ab_ule_i64_param_1];658; CHECK-NEXT: min.u64 %rd3, %rd1, %rd2;659; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;660; CHECK-NEXT: ret;661 %cmp = icmp ule i64 %a, %b662 %sel = select i1 %cmp, i64 %a, i64 %b663 ret i64 %sel664}665 666; *** ab, signed, i64667define i64 @ab_sgt_i64(i64 %a, i64 %b) {668; CHECK-LABEL: ab_sgt_i64(669; CHECK: {670; CHECK-NEXT: .reg .b64 %rd<4>;671; CHECK-EMPTY:672; CHECK-NEXT: // %bb.0:673; CHECK-NEXT: ld.param.b64 %rd1, [ab_sgt_i64_param_0];674; CHECK-NEXT: ld.param.b64 %rd2, [ab_sgt_i64_param_1];675; CHECK-NEXT: max.s64 %rd3, %rd1, %rd2;676; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;677; CHECK-NEXT: ret;678 %cmp = icmp sgt i64 %a, %b679 %sel = select i1 %cmp, i64 %a, i64 %b680 ret i64 %sel681}682 683define i64 @ab_sge_i64(i64 %a, i64 %b) {684; CHECK-LABEL: ab_sge_i64(685; CHECK: {686; CHECK-NEXT: .reg .b64 %rd<4>;687; CHECK-EMPTY:688; CHECK-NEXT: // %bb.0:689; CHECK-NEXT: ld.param.b64 %rd1, [ab_sge_i64_param_0];690; CHECK-NEXT: ld.param.b64 %rd2, [ab_sge_i64_param_1];691; CHECK-NEXT: max.s64 %rd3, %rd1, %rd2;692; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;693; CHECK-NEXT: ret;694 %cmp = icmp sge i64 %a, %b695 %sel = select i1 %cmp, i64 %a, i64 %b696 ret i64 %sel697}698 699define i64 @ab_slt_i64(i64 %a, i64 %b) {700; CHECK-LABEL: ab_slt_i64(701; CHECK: {702; CHECK-NEXT: .reg .b64 %rd<4>;703; CHECK-EMPTY:704; CHECK-NEXT: // %bb.0:705; CHECK-NEXT: ld.param.b64 %rd1, [ab_slt_i64_param_0];706; CHECK-NEXT: ld.param.b64 %rd2, [ab_slt_i64_param_1];707; CHECK-NEXT: min.s64 %rd3, %rd1, %rd2;708; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;709; CHECK-NEXT: ret;710 %cmp = icmp slt i64 %a, %b711 %sel = select i1 %cmp, i64 %a, i64 %b712 ret i64 %sel713}714 715define i64 @ab_sle_i64(i64 %a, i64 %b) {716; CHECK-LABEL: ab_sle_i64(717; CHECK: {718; CHECK-NEXT: .reg .b64 %rd<4>;719; CHECK-EMPTY:720; CHECK-NEXT: // %bb.0:721; CHECK-NEXT: ld.param.b64 %rd1, [ab_sle_i64_param_0];722; CHECK-NEXT: ld.param.b64 %rd2, [ab_sle_i64_param_1];723; CHECK-NEXT: min.s64 %rd3, %rd1, %rd2;724; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;725; CHECK-NEXT: ret;726 %cmp = icmp sle i64 %a, %b727 %sel = select i1 %cmp, i64 %a, i64 %b728 ret i64 %sel729}730 731; *** ba, unsigned, i64732define i64 @ba_ugt_i64(i64 %a, i64 %b) {733; CHECK-LABEL: ba_ugt_i64(734; CHECK: {735; CHECK-NEXT: .reg .b64 %rd<4>;736; CHECK-EMPTY:737; CHECK-NEXT: // %bb.0:738; CHECK-NEXT: ld.param.b64 %rd1, [ba_ugt_i64_param_0];739; CHECK-NEXT: ld.param.b64 %rd2, [ba_ugt_i64_param_1];740; CHECK-NEXT: min.u64 %rd3, %rd1, %rd2;741; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;742; CHECK-NEXT: ret;743 %cmp = icmp ugt i64 %a, %b744 %sel = select i1 %cmp, i64 %b, i64 %a745 ret i64 %sel746}747 748define i64 @ba_uge_i64(i64 %a, i64 %b) {749; CHECK-LABEL: ba_uge_i64(750; CHECK: {751; CHECK-NEXT: .reg .b64 %rd<4>;752; CHECK-EMPTY:753; CHECK-NEXT: // %bb.0:754; CHECK-NEXT: ld.param.b64 %rd1, [ba_uge_i64_param_0];755; CHECK-NEXT: ld.param.b64 %rd2, [ba_uge_i64_param_1];756; CHECK-NEXT: min.u64 %rd3, %rd1, %rd2;757; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;758; CHECK-NEXT: ret;759 %cmp = icmp uge i64 %a, %b760 %sel = select i1 %cmp, i64 %b, i64 %a761 ret i64 %sel762}763 764define i64 @ba_ult_i64(i64 %a, i64 %b) {765; CHECK-LABEL: ba_ult_i64(766; CHECK: {767; CHECK-NEXT: .reg .b64 %rd<4>;768; CHECK-EMPTY:769; CHECK-NEXT: // %bb.0:770; CHECK-NEXT: ld.param.b64 %rd1, [ba_ult_i64_param_0];771; CHECK-NEXT: ld.param.b64 %rd2, [ba_ult_i64_param_1];772; CHECK-NEXT: max.u64 %rd3, %rd1, %rd2;773; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;774; CHECK-NEXT: ret;775 %cmp = icmp ult i64 %a, %b776 %sel = select i1 %cmp, i64 %b, i64 %a777 ret i64 %sel778}779 780define i64 @ba_ule_i64(i64 %a, i64 %b) {781; CHECK-LABEL: ba_ule_i64(782; CHECK: {783; CHECK-NEXT: .reg .b64 %rd<4>;784; CHECK-EMPTY:785; CHECK-NEXT: // %bb.0:786; CHECK-NEXT: ld.param.b64 %rd1, [ba_ule_i64_param_0];787; CHECK-NEXT: ld.param.b64 %rd2, [ba_ule_i64_param_1];788; CHECK-NEXT: max.u64 %rd3, %rd1, %rd2;789; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;790; CHECK-NEXT: ret;791 %cmp = icmp ule i64 %a, %b792 %sel = select i1 %cmp, i64 %b, i64 %a793 ret i64 %sel794}795 796; *** ba, signed, i64797define i64 @ba_sgt_i64(i64 %a, i64 %b) {798; CHECK-LABEL: ba_sgt_i64(799; CHECK: {800; CHECK-NEXT: .reg .b64 %rd<4>;801; CHECK-EMPTY:802; CHECK-NEXT: // %bb.0:803; CHECK-NEXT: ld.param.b64 %rd1, [ba_sgt_i64_param_0];804; CHECK-NEXT: ld.param.b64 %rd2, [ba_sgt_i64_param_1];805; CHECK-NEXT: min.s64 %rd3, %rd1, %rd2;806; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;807; CHECK-NEXT: ret;808 %cmp = icmp sgt i64 %a, %b809 %sel = select i1 %cmp, i64 %b, i64 %a810 ret i64 %sel811}812 813define i64 @ba_sge_i64(i64 %a, i64 %b) {814; CHECK-LABEL: ba_sge_i64(815; CHECK: {816; CHECK-NEXT: .reg .b64 %rd<4>;817; CHECK-EMPTY:818; CHECK-NEXT: // %bb.0:819; CHECK-NEXT: ld.param.b64 %rd1, [ba_sge_i64_param_0];820; CHECK-NEXT: ld.param.b64 %rd2, [ba_sge_i64_param_1];821; CHECK-NEXT: min.s64 %rd3, %rd1, %rd2;822; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;823; CHECK-NEXT: ret;824 %cmp = icmp sge i64 %a, %b825 %sel = select i1 %cmp, i64 %b, i64 %a826 ret i64 %sel827}828 829define i64 @ba_slt_i64(i64 %a, i64 %b) {830; CHECK-LABEL: ba_slt_i64(831; CHECK: {832; CHECK-NEXT: .reg .b64 %rd<4>;833; CHECK-EMPTY:834; CHECK-NEXT: // %bb.0:835; CHECK-NEXT: ld.param.b64 %rd1, [ba_slt_i64_param_0];836; CHECK-NEXT: ld.param.b64 %rd2, [ba_slt_i64_param_1];837; CHECK-NEXT: max.s64 %rd3, %rd1, %rd2;838; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;839; CHECK-NEXT: ret;840 %cmp = icmp slt i64 %a, %b841 %sel = select i1 %cmp, i64 %b, i64 %a842 ret i64 %sel843}844 845define i64 @ba_sle_i64(i64 %a, i64 %b) {846; CHECK-LABEL: ba_sle_i64(847; CHECK: {848; CHECK-NEXT: .reg .b64 %rd<4>;849; CHECK-EMPTY:850; CHECK-NEXT: // %bb.0:851; CHECK-NEXT: ld.param.b64 %rd1, [ba_sle_i64_param_0];852; CHECK-NEXT: ld.param.b64 %rd2, [ba_sle_i64_param_1];853; CHECK-NEXT: max.s64 %rd3, %rd1, %rd2;854; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;855; CHECK-NEXT: ret;856 %cmp = icmp sle i64 %a, %b857 %sel = select i1 %cmp, i64 %b, i64 %a858 ret i64 %sel859}860 861define i32 @min_relu_s32(i32 %a, i32 %b) {862; SM90-LABEL: min_relu_s32(863; SM90: {864; SM90-NEXT: .reg .b32 %r<4>;865; SM90-EMPTY:866; SM90-NEXT: // %bb.0:867; SM90-NEXT: ld.param.b32 %r1, [min_relu_s32_param_0];868; SM90-NEXT: ld.param.b32 %r2, [min_relu_s32_param_1];869; SM90-NEXT: min.relu.s32 %r3, %r1, %r2;870; SM90-NEXT: st.param.b32 [func_retval0], %r3;871; SM90-NEXT: ret;872;873; SM20-LABEL: min_relu_s32(874; SM20: {875; SM20-NEXT: .reg .b32 %r<5>;876; SM20-EMPTY:877; SM20-NEXT: // %bb.0:878; SM20-NEXT: ld.param.b32 %r1, [min_relu_s32_param_0];879; SM20-NEXT: ld.param.b32 %r2, [min_relu_s32_param_1];880; SM20-NEXT: min.s32 %r3, %r1, %r2;881; SM20-NEXT: max.s32 %r4, %r3, 0;882; SM20-NEXT: st.param.b32 [func_retval0], %r4;883; SM20-NEXT: ret;884 %min = call i32 @llvm.smin.s32(i32 %a, i32 %b)885 %max = call i32 @llvm.smax.s32(i32 %min, i32 0)886 ret i32 %max887}888 889define i32 @max_relu_s32(i32 %a, i32 %b) {890; SM90-LABEL: max_relu_s32(891; SM90: {892; SM90-NEXT: .reg .b32 %r<4>;893; SM90-EMPTY:894; SM90-NEXT: // %bb.0:895; SM90-NEXT: ld.param.b32 %r1, [max_relu_s32_param_0];896; SM90-NEXT: ld.param.b32 %r2, [max_relu_s32_param_1];897; SM90-NEXT: max.relu.s32 %r3, %r1, %r2;898; SM90-NEXT: st.param.b32 [func_retval0], %r3;899; SM90-NEXT: ret;900;901; SM20-LABEL: max_relu_s32(902; SM20: {903; SM20-NEXT: .reg .b32 %r<5>;904; SM20-EMPTY:905; SM20-NEXT: // %bb.0:906; SM20-NEXT: ld.param.b32 %r1, [max_relu_s32_param_0];907; SM20-NEXT: ld.param.b32 %r2, [max_relu_s32_param_1];908; SM20-NEXT: max.s32 %r3, %r1, %r2;909; SM20-NEXT: max.s32 %r4, %r3, 0;910; SM20-NEXT: st.param.b32 [func_retval0], %r4;911; SM20-NEXT: ret;912 %max1 = call i32 @llvm.smax.s32(i32 %a, i32 %b)913 %max2 = call i32 @llvm.smax.s32(i32 %max1, i32 0)914 ret i32 %max2915}916 917define i32 @max_relu_s32_v2(i32 %a, i32 %b) {918; SM90-LABEL: max_relu_s32_v2(919; SM90: {920; SM90-NEXT: .reg .b32 %r<4>;921; SM90-EMPTY:922; SM90-NEXT: // %bb.0:923; SM90-NEXT: ld.param.b32 %r1, [max_relu_s32_v2_param_0];924; SM90-NEXT: ld.param.b32 %r2, [max_relu_s32_v2_param_1];925; SM90-NEXT: max.relu.s32 %r3, %r1, %r2;926; SM90-NEXT: st.param.b32 [func_retval0], %r3;927; SM90-NEXT: ret;928;929; SM20-LABEL: max_relu_s32_v2(930; SM20: {931; SM20-NEXT: .reg .b32 %r<5>;932; SM20-EMPTY:933; SM20-NEXT: // %bb.0:934; SM20-NEXT: ld.param.b32 %r1, [max_relu_s32_v2_param_0];935; SM20-NEXT: ld.param.b32 %r2, [max_relu_s32_v2_param_1];936; SM20-NEXT: max.s32 %r3, %r1, %r2;937; SM20-NEXT: max.s32 %r4, %r3, 0;938; SM20-NEXT: st.param.b32 [func_retval0], %r4;939; SM20-NEXT: ret;940 %max2 = call i32 @llvm.smax.s32(i32 %a, i32 0)941 %max1 = call i32 @llvm.smax.s32(i32 %max2, i32 %b)942 ret i32 %max1943}944 945define <2 x i16> @min_relu_s16x2(<2 x i16> %a, <2 x i16> %b) {946; SM90-LABEL: min_relu_s16x2(947; SM90: {948; SM90-NEXT: .reg .b32 %r<4>;949; SM90-EMPTY:950; SM90-NEXT: // %bb.0:951; SM90-NEXT: ld.param.b32 %r1, [min_relu_s16x2_param_0];952; SM90-NEXT: ld.param.b32 %r2, [min_relu_s16x2_param_1];953; SM90-NEXT: min.relu.s16x2 %r3, %r1, %r2;954; SM90-NEXT: st.param.b32 [func_retval0], %r3;955; SM90-NEXT: ret;956;957; SM20-LABEL: min_relu_s16x2(958; SM20: {959; SM20-NEXT: .reg .b16 %rs<9>;960; SM20-EMPTY:961; SM20-NEXT: // %bb.0:962; SM20-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [min_relu_s16x2_param_0];963; SM20-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [min_relu_s16x2_param_1];964; SM20-NEXT: min.s16 %rs5, %rs1, %rs3;965; SM20-NEXT: min.s16 %rs6, %rs2, %rs4;966; SM20-NEXT: max.s16 %rs7, %rs6, 0;967; SM20-NEXT: max.s16 %rs8, %rs5, 0;968; SM20-NEXT: st.param.v2.b16 [func_retval0], {%rs8, %rs7};969; SM20-NEXT: ret;970 %min = call <2 x i16> @llvm.smin.v2i16(<2 x i16> %a, <2 x i16> %b)971 %max = call <2 x i16> @llvm.smax.v2i16(<2 x i16> %min, <2 x i16> zeroinitializer)972 ret <2 x i16> %max973}974 975define <2 x i16> @max_relu_s16x2(<2 x i16> %a, <2 x i16> %b) {976; SM90-LABEL: max_relu_s16x2(977; SM90: {978; SM90-NEXT: .reg .b32 %r<4>;979; SM90-EMPTY:980; SM90-NEXT: // %bb.0:981; SM90-NEXT: ld.param.b32 %r1, [max_relu_s16x2_param_0];982; SM90-NEXT: ld.param.b32 %r2, [max_relu_s16x2_param_1];983; SM90-NEXT: max.relu.s16x2 %r3, %r1, %r2;984; SM90-NEXT: st.param.b32 [func_retval0], %r3;985; SM90-NEXT: ret;986;987; SM20-LABEL: max_relu_s16x2(988; SM20: {989; SM20-NEXT: .reg .b16 %rs<9>;990; SM20-EMPTY:991; SM20-NEXT: // %bb.0:992; SM20-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [max_relu_s16x2_param_0];993; SM20-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [max_relu_s16x2_param_1];994; SM20-NEXT: max.s16 %rs5, %rs1, %rs3;995; SM20-NEXT: max.s16 %rs6, %rs2, %rs4;996; SM20-NEXT: max.s16 %rs7, %rs6, 0;997; SM20-NEXT: max.s16 %rs8, %rs5, 0;998; SM20-NEXT: st.param.v2.b16 [func_retval0], {%rs8, %rs7};999; SM20-NEXT: ret;1000 %max1 = call <2 x i16> @llvm.smax.v2i16(<2 x i16> %a, <2 x i16> %b)1001 %max2 = call <2 x i16> @llvm.smax.v2i16(<2 x i16> %max1, <2 x i16> zeroinitializer)1002 ret <2 x i16> %max21003}1004 1005define <2 x i16> @max_relu_s16x2_v2(<2 x i16> %a, <2 x i16> %b) {1006; SM90-LABEL: max_relu_s16x2_v2(1007; SM90: {1008; SM90-NEXT: .reg .b32 %r<4>;1009; SM90-EMPTY:1010; SM90-NEXT: // %bb.0:1011; SM90-NEXT: ld.param.b32 %r1, [max_relu_s16x2_v2_param_0];1012; SM90-NEXT: ld.param.b32 %r2, [max_relu_s16x2_v2_param_1];1013; SM90-NEXT: max.relu.s16x2 %r3, %r1, %r2;1014; SM90-NEXT: st.param.b32 [func_retval0], %r3;1015; SM90-NEXT: ret;1016;1017; SM20-LABEL: max_relu_s16x2_v2(1018; SM20: {1019; SM20-NEXT: .reg .b16 %rs<9>;1020; SM20-EMPTY:1021; SM20-NEXT: // %bb.0:1022; SM20-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [max_relu_s16x2_v2_param_0];1023; SM20-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [max_relu_s16x2_v2_param_1];1024; SM20-NEXT: max.s16 %rs5, %rs1, %rs3;1025; SM20-NEXT: max.s16 %rs6, %rs2, %rs4;1026; SM20-NEXT: max.s16 %rs7, %rs6, 0;1027; SM20-NEXT: max.s16 %rs8, %rs5, 0;1028; SM20-NEXT: st.param.v2.b16 [func_retval0], {%rs8, %rs7};1029; SM20-NEXT: ret;1030 %max2 = call <2 x i16> @llvm.smax.v2i16(<2 x i16> %a, <2 x i16> zeroinitializer)1031 %max1 = call <2 x i16> @llvm.smax.v2i16(<2 x i16> %max2, <2 x i16> %b)1032 ret <2 x i16> %max11033}1034