615 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mtriple=nvptx64-- 2>&1 | FileCheck %s3; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64-- | %ptxas-verify %}4 5define i128 @srem_i128(i128 %lhs, i128 %rhs) {6; CHECK-LABEL: srem_i128(7; CHECK: {8; CHECK-NEXT: .reg .pred %p<20>;9; CHECK-NEXT: .reg .b32 %r<12>;10; CHECK-NEXT: .reg .b64 %rd<79>;11; CHECK-EMPTY:12; CHECK-NEXT: // %bb.0: // %_udiv-special-cases13; CHECK-NEXT: ld.param.v2.b64 {%rd8, %rd9}, [srem_i128_param_0];14; CHECK-NEXT: ld.param.v2.b64 {%rd10, %rd11}, [srem_i128_param_1];15; CHECK-NEXT: shr.s64 %rd1, %rd9, 63;16; CHECK-NEXT: sub.cc.s64 %rd12, 0, %rd8;17; CHECK-NEXT: subc.cc.s64 %rd13, 0, %rd9;18; CHECK-NEXT: setp.lt.s64 %p1, %rd9, 0;19; CHECK-NEXT: selp.b64 %rd3, %rd13, %rd9, %p1;20; CHECK-NEXT: selp.b64 %rd2, %rd12, %rd8, %p1;21; CHECK-NEXT: sub.cc.s64 %rd14, 0, %rd10;22; CHECK-NEXT: subc.cc.s64 %rd15, 0, %rd11;23; CHECK-NEXT: setp.lt.s64 %p2, %rd11, 0;24; CHECK-NEXT: selp.b64 %rd5, %rd15, %rd11, %p2;25; CHECK-NEXT: selp.b64 %rd4, %rd14, %rd10, %p2;26; CHECK-NEXT: or.b64 %rd16, %rd4, %rd5;27; CHECK-NEXT: setp.eq.b64 %p3, %rd16, 0;28; CHECK-NEXT: or.b64 %rd17, %rd2, %rd3;29; CHECK-NEXT: setp.eq.b64 %p4, %rd17, 0;30; CHECK-NEXT: or.pred %p5, %p3, %p4;31; CHECK-NEXT: setp.ne.b64 %p6, %rd5, 0;32; CHECK-NEXT: clz.b64 %r1, %rd5;33; CHECK-NEXT: cvt.u64.u32 %rd18, %r1;34; CHECK-NEXT: clz.b64 %r2, %rd4;35; CHECK-NEXT: cvt.u64.u32 %rd19, %r2;36; CHECK-NEXT: add.s64 %rd20, %rd19, 64;37; CHECK-NEXT: selp.b64 %rd21, %rd18, %rd20, %p6;38; CHECK-NEXT: setp.ne.b64 %p7, %rd3, 0;39; CHECK-NEXT: clz.b64 %r3, %rd3;40; CHECK-NEXT: cvt.u64.u32 %rd22, %r3;41; CHECK-NEXT: clz.b64 %r4, %rd2;42; CHECK-NEXT: cvt.u64.u32 %rd23, %r4;43; CHECK-NEXT: add.s64 %rd24, %rd23, 64;44; CHECK-NEXT: selp.b64 %rd25, %rd22, %rd24, %p7;45; CHECK-NEXT: mov.b64 %rd70, 0;46; CHECK-NEXT: sub.cc.s64 %rd26, %rd21, %rd25;47; CHECK-NEXT: subc.cc.s64 %rd27, %rd70, 0;48; CHECK-NEXT: setp.gt.u64 %p8, %rd26, 127;49; CHECK-NEXT: setp.eq.b64 %p9, %rd27, 0;50; CHECK-NEXT: and.pred %p10, %p9, %p8;51; CHECK-NEXT: setp.ne.b64 %p11, %rd27, 0;52; CHECK-NEXT: or.pred %p12, %p10, %p11;53; CHECK-NEXT: or.pred %p13, %p5, %p12;54; CHECK-NEXT: xor.b64 %rd28, %rd26, 127;55; CHECK-NEXT: or.b64 %rd29, %rd28, %rd27;56; CHECK-NEXT: setp.eq.b64 %p14, %rd29, 0;57; CHECK-NEXT: selp.b64 %rd78, 0, %rd3, %p13;58; CHECK-NEXT: selp.b64 %rd77, 0, %rd2, %p13;59; CHECK-NEXT: or.pred %p15, %p13, %p14;60; CHECK-NEXT: @%p15 bra $L__BB0_5;61; CHECK-NEXT: // %bb.3: // %udiv-bb162; CHECK-NEXT: add.cc.s64 %rd71, %rd26, 1;63; CHECK-NEXT: addc.cc.s64 %rd72, %rd27, 0;64; CHECK-NEXT: or.b64 %rd30, %rd71, %rd72;65; CHECK-NEXT: setp.eq.b64 %p16, %rd30, 0;66; CHECK-NEXT: cvt.u32.u64 %r5, %rd26;67; CHECK-NEXT: sub.s32 %r6, 127, %r5;68; CHECK-NEXT: shl.b64 %rd31, %rd3, %r6;69; CHECK-NEXT: sub.s32 %r7, 64, %r6;70; CHECK-NEXT: shr.u64 %rd32, %rd2, %r7;71; CHECK-NEXT: or.b64 %rd33, %rd31, %rd32;72; CHECK-NEXT: sub.s32 %r8, 63, %r5;73; CHECK-NEXT: shl.b64 %rd34, %rd2, %r8;74; CHECK-NEXT: setp.gt.s32 %p17, %r6, 63;75; CHECK-NEXT: selp.b64 %rd76, %rd34, %rd33, %p17;76; CHECK-NEXT: shl.b64 %rd75, %rd2, %r6;77; CHECK-NEXT: mov.b64 %rd69, %rd70;78; CHECK-NEXT: @%p16 bra $L__BB0_4;79; CHECK-NEXT: // %bb.1: // %udiv-preheader80; CHECK-NEXT: cvt.u32.u64 %r9, %rd71;81; CHECK-NEXT: shr.u64 %rd35, %rd2, %r9;82; CHECK-NEXT: sub.s32 %r10, 64, %r9;83; CHECK-NEXT: shl.b64 %rd36, %rd3, %r10;84; CHECK-NEXT: or.b64 %rd37, %rd35, %rd36;85; CHECK-NEXT: add.s32 %r11, %r9, -64;86; CHECK-NEXT: shr.u64 %rd38, %rd3, %r11;87; CHECK-NEXT: setp.gt.s32 %p18, %r9, 63;88; CHECK-NEXT: selp.b64 %rd73, %rd38, %rd37, %p18;89; CHECK-NEXT: shr.u64 %rd74, %rd3, %r9;90; CHECK-NEXT: add.cc.s64 %rd6, %rd4, -1;91; CHECK-NEXT: addc.cc.s64 %rd7, %rd5, -1;92; CHECK-NEXT: mov.b64 %rd69, 0;93; CHECK-NEXT: mov.b64 %rd70, %rd69;94; CHECK-NEXT: $L__BB0_2: // %udiv-do-while95; CHECK-NEXT: // =>This Inner Loop Header: Depth=196; CHECK-NEXT: shr.u64 %rd39, %rd73, 63;97; CHECK-NEXT: shl.b64 %rd40, %rd74, 1;98; CHECK-NEXT: or.b64 %rd41, %rd40, %rd39;99; CHECK-NEXT: shl.b64 %rd42, %rd73, 1;100; CHECK-NEXT: shr.u64 %rd43, %rd76, 63;101; CHECK-NEXT: or.b64 %rd44, %rd42, %rd43;102; CHECK-NEXT: shr.u64 %rd45, %rd75, 63;103; CHECK-NEXT: shl.b64 %rd46, %rd76, 1;104; CHECK-NEXT: or.b64 %rd47, %rd46, %rd45;105; CHECK-NEXT: shl.b64 %rd48, %rd75, 1;106; CHECK-NEXT: or.b64 %rd75, %rd70, %rd48;107; CHECK-NEXT: or.b64 %rd76, %rd69, %rd47;108; CHECK-NEXT: sub.cc.s64 %rd49, %rd6, %rd44;109; CHECK-NEXT: subc.cc.s64 %rd50, %rd7, %rd41;110; CHECK-NEXT: shr.s64 %rd51, %rd50, 63;111; CHECK-NEXT: and.b64 %rd70, %rd51, 1;112; CHECK-NEXT: and.b64 %rd52, %rd51, %rd4;113; CHECK-NEXT: and.b64 %rd53, %rd51, %rd5;114; CHECK-NEXT: sub.cc.s64 %rd73, %rd44, %rd52;115; CHECK-NEXT: subc.cc.s64 %rd74, %rd41, %rd53;116; CHECK-NEXT: add.cc.s64 %rd71, %rd71, -1;117; CHECK-NEXT: addc.cc.s64 %rd72, %rd72, -1;118; CHECK-NEXT: or.b64 %rd54, %rd71, %rd72;119; CHECK-NEXT: setp.eq.b64 %p19, %rd54, 0;120; CHECK-NEXT: @%p19 bra $L__BB0_4;121; CHECK-NEXT: bra.uni $L__BB0_2;122; CHECK-NEXT: $L__BB0_4: // %udiv-loop-exit123; CHECK-NEXT: shr.u64 %rd55, %rd75, 63;124; CHECK-NEXT: shl.b64 %rd56, %rd76, 1;125; CHECK-NEXT: or.b64 %rd57, %rd56, %rd55;126; CHECK-NEXT: shl.b64 %rd58, %rd75, 1;127; CHECK-NEXT: or.b64 %rd77, %rd70, %rd58;128; CHECK-NEXT: or.b64 %rd78, %rd69, %rd57;129; CHECK-NEXT: $L__BB0_5: // %udiv-end130; CHECK-NEXT: mul.hi.u64 %rd59, %rd4, %rd77;131; CHECK-NEXT: mad.lo.s64 %rd60, %rd4, %rd78, %rd59;132; CHECK-NEXT: mad.lo.s64 %rd61, %rd5, %rd77, %rd60;133; CHECK-NEXT: mul.lo.s64 %rd62, %rd4, %rd77;134; CHECK-NEXT: sub.cc.s64 %rd63, %rd2, %rd62;135; CHECK-NEXT: subc.cc.s64 %rd64, %rd3, %rd61;136; CHECK-NEXT: xor.b64 %rd65, %rd63, %rd1;137; CHECK-NEXT: xor.b64 %rd66, %rd64, %rd1;138; CHECK-NEXT: sub.cc.s64 %rd67, %rd65, %rd1;139; CHECK-NEXT: subc.cc.s64 %rd68, %rd66, %rd1;140; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd67, %rd68};141; CHECK-NEXT: ret;142 %div = srem i128 %lhs, %rhs143 ret i128 %div144}145 146define i128 @urem_i128(i128 %lhs, i128 %rhs) {147; CHECK-LABEL: urem_i128(148; CHECK: {149; CHECK-NEXT: .reg .pred %p<18>;150; CHECK-NEXT: .reg .b32 %r<12>;151; CHECK-NEXT: .reg .b64 %rd<66>;152; CHECK-EMPTY:153; CHECK-NEXT: // %bb.0: // %_udiv-special-cases154; CHECK-NEXT: ld.param.v2.b64 {%rd5, %rd6}, [urem_i128_param_0];155; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [urem_i128_param_1];156; CHECK-NEXT: or.b64 %rd7, %rd1, %rd2;157; CHECK-NEXT: setp.eq.b64 %p1, %rd7, 0;158; CHECK-NEXT: or.b64 %rd8, %rd5, %rd6;159; CHECK-NEXT: setp.eq.b64 %p2, %rd8, 0;160; CHECK-NEXT: or.pred %p3, %p1, %p2;161; CHECK-NEXT: setp.ne.b64 %p4, %rd2, 0;162; CHECK-NEXT: clz.b64 %r1, %rd2;163; CHECK-NEXT: cvt.u64.u32 %rd9, %r1;164; CHECK-NEXT: clz.b64 %r2, %rd1;165; CHECK-NEXT: cvt.u64.u32 %rd10, %r2;166; CHECK-NEXT: add.s64 %rd11, %rd10, 64;167; CHECK-NEXT: selp.b64 %rd12, %rd9, %rd11, %p4;168; CHECK-NEXT: setp.ne.b64 %p5, %rd6, 0;169; CHECK-NEXT: clz.b64 %r3, %rd6;170; CHECK-NEXT: cvt.u64.u32 %rd13, %r3;171; CHECK-NEXT: clz.b64 %r4, %rd5;172; CHECK-NEXT: cvt.u64.u32 %rd14, %r4;173; CHECK-NEXT: add.s64 %rd15, %rd14, 64;174; CHECK-NEXT: selp.b64 %rd16, %rd13, %rd15, %p5;175; CHECK-NEXT: mov.b64 %rd57, 0;176; CHECK-NEXT: sub.cc.s64 %rd17, %rd12, %rd16;177; CHECK-NEXT: subc.cc.s64 %rd18, %rd57, 0;178; CHECK-NEXT: setp.gt.u64 %p6, %rd17, 127;179; CHECK-NEXT: setp.eq.b64 %p7, %rd18, 0;180; CHECK-NEXT: and.pred %p8, %p7, %p6;181; CHECK-NEXT: setp.ne.b64 %p9, %rd18, 0;182; CHECK-NEXT: or.pred %p10, %p8, %p9;183; CHECK-NEXT: or.pred %p11, %p3, %p10;184; CHECK-NEXT: xor.b64 %rd19, %rd17, 127;185; CHECK-NEXT: or.b64 %rd20, %rd19, %rd18;186; CHECK-NEXT: setp.eq.b64 %p12, %rd20, 0;187; CHECK-NEXT: selp.b64 %rd65, 0, %rd6, %p11;188; CHECK-NEXT: selp.b64 %rd64, 0, %rd5, %p11;189; CHECK-NEXT: or.pred %p13, %p11, %p12;190; CHECK-NEXT: @%p13 bra $L__BB1_5;191; CHECK-NEXT: // %bb.3: // %udiv-bb1192; CHECK-NEXT: add.cc.s64 %rd58, %rd17, 1;193; CHECK-NEXT: addc.cc.s64 %rd59, %rd18, 0;194; CHECK-NEXT: or.b64 %rd21, %rd58, %rd59;195; CHECK-NEXT: setp.eq.b64 %p14, %rd21, 0;196; CHECK-NEXT: cvt.u32.u64 %r5, %rd17;197; CHECK-NEXT: sub.s32 %r6, 127, %r5;198; CHECK-NEXT: shl.b64 %rd22, %rd6, %r6;199; CHECK-NEXT: sub.s32 %r7, 64, %r6;200; CHECK-NEXT: shr.u64 %rd23, %rd5, %r7;201; CHECK-NEXT: or.b64 %rd24, %rd22, %rd23;202; CHECK-NEXT: sub.s32 %r8, 63, %r5;203; CHECK-NEXT: shl.b64 %rd25, %rd5, %r8;204; CHECK-NEXT: setp.gt.s32 %p15, %r6, 63;205; CHECK-NEXT: selp.b64 %rd63, %rd25, %rd24, %p15;206; CHECK-NEXT: shl.b64 %rd62, %rd5, %r6;207; CHECK-NEXT: mov.b64 %rd56, %rd57;208; CHECK-NEXT: @%p14 bra $L__BB1_4;209; CHECK-NEXT: // %bb.1: // %udiv-preheader210; CHECK-NEXT: cvt.u32.u64 %r9, %rd58;211; CHECK-NEXT: shr.u64 %rd26, %rd5, %r9;212; CHECK-NEXT: sub.s32 %r10, 64, %r9;213; CHECK-NEXT: shl.b64 %rd27, %rd6, %r10;214; CHECK-NEXT: or.b64 %rd28, %rd26, %rd27;215; CHECK-NEXT: add.s32 %r11, %r9, -64;216; CHECK-NEXT: shr.u64 %rd29, %rd6, %r11;217; CHECK-NEXT: setp.gt.s32 %p16, %r9, 63;218; CHECK-NEXT: selp.b64 %rd60, %rd29, %rd28, %p16;219; CHECK-NEXT: shr.u64 %rd61, %rd6, %r9;220; CHECK-NEXT: add.cc.s64 %rd3, %rd1, -1;221; CHECK-NEXT: addc.cc.s64 %rd4, %rd2, -1;222; CHECK-NEXT: mov.b64 %rd56, 0;223; CHECK-NEXT: mov.b64 %rd57, %rd56;224; CHECK-NEXT: $L__BB1_2: // %udiv-do-while225; CHECK-NEXT: // =>This Inner Loop Header: Depth=1226; CHECK-NEXT: shr.u64 %rd30, %rd60, 63;227; CHECK-NEXT: shl.b64 %rd31, %rd61, 1;228; CHECK-NEXT: or.b64 %rd32, %rd31, %rd30;229; CHECK-NEXT: shl.b64 %rd33, %rd60, 1;230; CHECK-NEXT: shr.u64 %rd34, %rd63, 63;231; CHECK-NEXT: or.b64 %rd35, %rd33, %rd34;232; CHECK-NEXT: shr.u64 %rd36, %rd62, 63;233; CHECK-NEXT: shl.b64 %rd37, %rd63, 1;234; CHECK-NEXT: or.b64 %rd38, %rd37, %rd36;235; CHECK-NEXT: shl.b64 %rd39, %rd62, 1;236; CHECK-NEXT: or.b64 %rd62, %rd57, %rd39;237; CHECK-NEXT: or.b64 %rd63, %rd56, %rd38;238; CHECK-NEXT: sub.cc.s64 %rd40, %rd3, %rd35;239; CHECK-NEXT: subc.cc.s64 %rd41, %rd4, %rd32;240; CHECK-NEXT: shr.s64 %rd42, %rd41, 63;241; CHECK-NEXT: and.b64 %rd57, %rd42, 1;242; CHECK-NEXT: and.b64 %rd43, %rd42, %rd1;243; CHECK-NEXT: and.b64 %rd44, %rd42, %rd2;244; CHECK-NEXT: sub.cc.s64 %rd60, %rd35, %rd43;245; CHECK-NEXT: subc.cc.s64 %rd61, %rd32, %rd44;246; CHECK-NEXT: add.cc.s64 %rd58, %rd58, -1;247; CHECK-NEXT: addc.cc.s64 %rd59, %rd59, -1;248; CHECK-NEXT: or.b64 %rd45, %rd58, %rd59;249; CHECK-NEXT: setp.eq.b64 %p17, %rd45, 0;250; CHECK-NEXT: @%p17 bra $L__BB1_4;251; CHECK-NEXT: bra.uni $L__BB1_2;252; CHECK-NEXT: $L__BB1_4: // %udiv-loop-exit253; CHECK-NEXT: shr.u64 %rd46, %rd62, 63;254; CHECK-NEXT: shl.b64 %rd47, %rd63, 1;255; CHECK-NEXT: or.b64 %rd48, %rd47, %rd46;256; CHECK-NEXT: shl.b64 %rd49, %rd62, 1;257; CHECK-NEXT: or.b64 %rd64, %rd57, %rd49;258; CHECK-NEXT: or.b64 %rd65, %rd56, %rd48;259; CHECK-NEXT: $L__BB1_5: // %udiv-end260; CHECK-NEXT: mul.hi.u64 %rd50, %rd1, %rd64;261; CHECK-NEXT: mad.lo.s64 %rd51, %rd1, %rd65, %rd50;262; CHECK-NEXT: mad.lo.s64 %rd52, %rd2, %rd64, %rd51;263; CHECK-NEXT: mul.lo.s64 %rd53, %rd1, %rd64;264; CHECK-NEXT: sub.cc.s64 %rd54, %rd5, %rd53;265; CHECK-NEXT: subc.cc.s64 %rd55, %rd6, %rd52;266; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd54, %rd55};267; CHECK-NEXT: ret;268 %div = urem i128 %lhs, %rhs269 ret i128 %div270}271 272define i128 @srem_i128_pow2k(i128 %lhs) {273; CHECK-LABEL: srem_i128_pow2k(274; CHECK: {275; CHECK-NEXT: .reg .b64 %rd<10>;276; CHECK-EMPTY:277; CHECK-NEXT: // %bb.0:278; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [srem_i128_pow2k_param_0];279; CHECK-NEXT: shr.s64 %rd3, %rd2, 63;280; CHECK-NEXT: shr.u64 %rd4, %rd3, 31;281; CHECK-NEXT: add.cc.s64 %rd5, %rd1, %rd4;282; CHECK-NEXT: addc.cc.s64 %rd6, %rd2, 0;283; CHECK-NEXT: and.b64 %rd7, %rd5, -8589934592;284; CHECK-NEXT: sub.cc.s64 %rd8, %rd1, %rd7;285; CHECK-NEXT: subc.cc.s64 %rd9, %rd2, %rd6;286; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd8, %rd9};287; CHECK-NEXT: ret;288 %div = srem i128 %lhs, 8589934592289 ret i128 %div290}291 292define i128 @urem_i128_pow2k(i128 %lhs) {293; CHECK-LABEL: urem_i128_pow2k(294; CHECK: {295; CHECK-NEXT: .reg .b64 %rd<4>;296; CHECK-EMPTY:297; CHECK-NEXT: // %bb.0:298; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [urem_i128_pow2k_param_0];299; CHECK-NEXT: and.b64 %rd3, %rd1, 8589934591;300; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd3, 0};301; CHECK-NEXT: ret;302 %div = urem i128 %lhs, 8589934592303 ret i128 %div304}305 306define i128 @sdiv_i128(i128 %lhs, i128 %rhs) {307; CHECK-LABEL: sdiv_i128(308; CHECK: {309; CHECK-NEXT: .reg .pred %p<20>;310; CHECK-NEXT: .reg .b32 %r<12>;311; CHECK-NEXT: .reg .b64 %rd<74>;312; CHECK-EMPTY:313; CHECK-NEXT: // %bb.0: // %_udiv-special-cases314; CHECK-NEXT: ld.param.v2.b64 {%rd8, %rd9}, [sdiv_i128_param_0];315; CHECK-NEXT: ld.param.v2.b64 {%rd10, %rd11}, [sdiv_i128_param_1];316; CHECK-NEXT: sub.cc.s64 %rd12, 0, %rd8;317; CHECK-NEXT: subc.cc.s64 %rd13, 0, %rd9;318; CHECK-NEXT: setp.lt.s64 %p1, %rd9, 0;319; CHECK-NEXT: selp.b64 %rd2, %rd13, %rd9, %p1;320; CHECK-NEXT: selp.b64 %rd1, %rd12, %rd8, %p1;321; CHECK-NEXT: sub.cc.s64 %rd14, 0, %rd10;322; CHECK-NEXT: subc.cc.s64 %rd15, 0, %rd11;323; CHECK-NEXT: setp.lt.s64 %p2, %rd11, 0;324; CHECK-NEXT: selp.b64 %rd4, %rd15, %rd11, %p2;325; CHECK-NEXT: selp.b64 %rd3, %rd14, %rd10, %p2;326; CHECK-NEXT: xor.b64 %rd16, %rd11, %rd9;327; CHECK-NEXT: shr.s64 %rd5, %rd16, 63;328; CHECK-NEXT: or.b64 %rd17, %rd3, %rd4;329; CHECK-NEXT: setp.eq.b64 %p3, %rd17, 0;330; CHECK-NEXT: or.b64 %rd18, %rd1, %rd2;331; CHECK-NEXT: setp.eq.b64 %p4, %rd18, 0;332; CHECK-NEXT: or.pred %p5, %p3, %p4;333; CHECK-NEXT: setp.ne.b64 %p6, %rd4, 0;334; CHECK-NEXT: clz.b64 %r1, %rd4;335; CHECK-NEXT: cvt.u64.u32 %rd19, %r1;336; CHECK-NEXT: clz.b64 %r2, %rd3;337; CHECK-NEXT: cvt.u64.u32 %rd20, %r2;338; CHECK-NEXT: add.s64 %rd21, %rd20, 64;339; CHECK-NEXT: selp.b64 %rd22, %rd19, %rd21, %p6;340; CHECK-NEXT: setp.ne.b64 %p7, %rd2, 0;341; CHECK-NEXT: clz.b64 %r3, %rd2;342; CHECK-NEXT: cvt.u64.u32 %rd23, %r3;343; CHECK-NEXT: clz.b64 %r4, %rd1;344; CHECK-NEXT: cvt.u64.u32 %rd24, %r4;345; CHECK-NEXT: add.s64 %rd25, %rd24, 64;346; CHECK-NEXT: selp.b64 %rd26, %rd23, %rd25, %p7;347; CHECK-NEXT: mov.b64 %rd65, 0;348; CHECK-NEXT: sub.cc.s64 %rd27, %rd22, %rd26;349; CHECK-NEXT: subc.cc.s64 %rd28, %rd65, 0;350; CHECK-NEXT: setp.gt.u64 %p8, %rd27, 127;351; CHECK-NEXT: setp.eq.b64 %p9, %rd28, 0;352; CHECK-NEXT: and.pred %p10, %p9, %p8;353; CHECK-NEXT: setp.ne.b64 %p11, %rd28, 0;354; CHECK-NEXT: or.pred %p12, %p10, %p11;355; CHECK-NEXT: or.pred %p13, %p5, %p12;356; CHECK-NEXT: xor.b64 %rd29, %rd27, 127;357; CHECK-NEXT: or.b64 %rd30, %rd29, %rd28;358; CHECK-NEXT: setp.eq.b64 %p14, %rd30, 0;359; CHECK-NEXT: selp.b64 %rd73, 0, %rd2, %p13;360; CHECK-NEXT: selp.b64 %rd72, 0, %rd1, %p13;361; CHECK-NEXT: or.pred %p15, %p13, %p14;362; CHECK-NEXT: @%p15 bra $L__BB4_5;363; CHECK-NEXT: // %bb.3: // %udiv-bb1364; CHECK-NEXT: add.cc.s64 %rd66, %rd27, 1;365; CHECK-NEXT: addc.cc.s64 %rd67, %rd28, 0;366; CHECK-NEXT: or.b64 %rd31, %rd66, %rd67;367; CHECK-NEXT: setp.eq.b64 %p16, %rd31, 0;368; CHECK-NEXT: cvt.u32.u64 %r5, %rd27;369; CHECK-NEXT: sub.s32 %r6, 127, %r5;370; CHECK-NEXT: shl.b64 %rd32, %rd2, %r6;371; CHECK-NEXT: sub.s32 %r7, 64, %r6;372; CHECK-NEXT: shr.u64 %rd33, %rd1, %r7;373; CHECK-NEXT: or.b64 %rd34, %rd32, %rd33;374; CHECK-NEXT: sub.s32 %r8, 63, %r5;375; CHECK-NEXT: shl.b64 %rd35, %rd1, %r8;376; CHECK-NEXT: setp.gt.s32 %p17, %r6, 63;377; CHECK-NEXT: selp.b64 %rd71, %rd35, %rd34, %p17;378; CHECK-NEXT: shl.b64 %rd70, %rd1, %r6;379; CHECK-NEXT: mov.b64 %rd64, %rd65;380; CHECK-NEXT: @%p16 bra $L__BB4_4;381; CHECK-NEXT: // %bb.1: // %udiv-preheader382; CHECK-NEXT: cvt.u32.u64 %r9, %rd66;383; CHECK-NEXT: shr.u64 %rd36, %rd1, %r9;384; CHECK-NEXT: sub.s32 %r10, 64, %r9;385; CHECK-NEXT: shl.b64 %rd37, %rd2, %r10;386; CHECK-NEXT: or.b64 %rd38, %rd36, %rd37;387; CHECK-NEXT: add.s32 %r11, %r9, -64;388; CHECK-NEXT: shr.u64 %rd39, %rd2, %r11;389; CHECK-NEXT: setp.gt.s32 %p18, %r9, 63;390; CHECK-NEXT: selp.b64 %rd68, %rd39, %rd38, %p18;391; CHECK-NEXT: shr.u64 %rd69, %rd2, %r9;392; CHECK-NEXT: add.cc.s64 %rd6, %rd3, -1;393; CHECK-NEXT: addc.cc.s64 %rd7, %rd4, -1;394; CHECK-NEXT: mov.b64 %rd64, 0;395; CHECK-NEXT: mov.b64 %rd65, %rd64;396; CHECK-NEXT: $L__BB4_2: // %udiv-do-while397; CHECK-NEXT: // =>This Inner Loop Header: Depth=1398; CHECK-NEXT: shr.u64 %rd40, %rd68, 63;399; CHECK-NEXT: shl.b64 %rd41, %rd69, 1;400; CHECK-NEXT: or.b64 %rd42, %rd41, %rd40;401; CHECK-NEXT: shl.b64 %rd43, %rd68, 1;402; CHECK-NEXT: shr.u64 %rd44, %rd71, 63;403; CHECK-NEXT: or.b64 %rd45, %rd43, %rd44;404; CHECK-NEXT: shr.u64 %rd46, %rd70, 63;405; CHECK-NEXT: shl.b64 %rd47, %rd71, 1;406; CHECK-NEXT: or.b64 %rd48, %rd47, %rd46;407; CHECK-NEXT: shl.b64 %rd49, %rd70, 1;408; CHECK-NEXT: or.b64 %rd70, %rd65, %rd49;409; CHECK-NEXT: or.b64 %rd71, %rd64, %rd48;410; CHECK-NEXT: sub.cc.s64 %rd50, %rd6, %rd45;411; CHECK-NEXT: subc.cc.s64 %rd51, %rd7, %rd42;412; CHECK-NEXT: shr.s64 %rd52, %rd51, 63;413; CHECK-NEXT: and.b64 %rd65, %rd52, 1;414; CHECK-NEXT: and.b64 %rd53, %rd52, %rd3;415; CHECK-NEXT: and.b64 %rd54, %rd52, %rd4;416; CHECK-NEXT: sub.cc.s64 %rd68, %rd45, %rd53;417; CHECK-NEXT: subc.cc.s64 %rd69, %rd42, %rd54;418; CHECK-NEXT: add.cc.s64 %rd66, %rd66, -1;419; CHECK-NEXT: addc.cc.s64 %rd67, %rd67, -1;420; CHECK-NEXT: or.b64 %rd55, %rd66, %rd67;421; CHECK-NEXT: setp.eq.b64 %p19, %rd55, 0;422; CHECK-NEXT: @%p19 bra $L__BB4_4;423; CHECK-NEXT: bra.uni $L__BB4_2;424; CHECK-NEXT: $L__BB4_4: // %udiv-loop-exit425; CHECK-NEXT: shr.u64 %rd56, %rd70, 63;426; CHECK-NEXT: shl.b64 %rd57, %rd71, 1;427; CHECK-NEXT: or.b64 %rd58, %rd57, %rd56;428; CHECK-NEXT: shl.b64 %rd59, %rd70, 1;429; CHECK-NEXT: or.b64 %rd72, %rd65, %rd59;430; CHECK-NEXT: or.b64 %rd73, %rd64, %rd58;431; CHECK-NEXT: $L__BB4_5: // %udiv-end432; CHECK-NEXT: xor.b64 %rd60, %rd72, %rd5;433; CHECK-NEXT: xor.b64 %rd61, %rd73, %rd5;434; CHECK-NEXT: sub.cc.s64 %rd62, %rd60, %rd5;435; CHECK-NEXT: subc.cc.s64 %rd63, %rd61, %rd5;436; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd62, %rd63};437; CHECK-NEXT: ret;438 %div = sdiv i128 %lhs, %rhs439 ret i128 %div440}441 442define i128 @udiv_i128(i128 %lhs, i128 %rhs) {443; CHECK-LABEL: udiv_i128(444; CHECK: {445; CHECK-NEXT: .reg .pred %p<18>;446; CHECK-NEXT: .reg .b32 %r<12>;447; CHECK-NEXT: .reg .b64 %rd<60>;448; CHECK-EMPTY:449; CHECK-NEXT: // %bb.0: // %_udiv-special-cases450; CHECK-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [udiv_i128_param_0];451; CHECK-NEXT: ld.param.v2.b64 {%rd5, %rd6}, [udiv_i128_param_1];452; CHECK-NEXT: or.b64 %rd7, %rd5, %rd6;453; CHECK-NEXT: setp.eq.b64 %p1, %rd7, 0;454; CHECK-NEXT: or.b64 %rd8, %rd3, %rd4;455; CHECK-NEXT: setp.eq.b64 %p2, %rd8, 0;456; CHECK-NEXT: or.pred %p3, %p1, %p2;457; CHECK-NEXT: setp.ne.b64 %p4, %rd6, 0;458; CHECK-NEXT: clz.b64 %r1, %rd6;459; CHECK-NEXT: cvt.u64.u32 %rd9, %r1;460; CHECK-NEXT: clz.b64 %r2, %rd5;461; CHECK-NEXT: cvt.u64.u32 %rd10, %r2;462; CHECK-NEXT: add.s64 %rd11, %rd10, 64;463; CHECK-NEXT: selp.b64 %rd12, %rd9, %rd11, %p4;464; CHECK-NEXT: setp.ne.b64 %p5, %rd4, 0;465; CHECK-NEXT: clz.b64 %r3, %rd4;466; CHECK-NEXT: cvt.u64.u32 %rd13, %r3;467; CHECK-NEXT: clz.b64 %r4, %rd3;468; CHECK-NEXT: cvt.u64.u32 %rd14, %r4;469; CHECK-NEXT: add.s64 %rd15, %rd14, 64;470; CHECK-NEXT: selp.b64 %rd16, %rd13, %rd15, %p5;471; CHECK-NEXT: mov.b64 %rd51, 0;472; CHECK-NEXT: sub.cc.s64 %rd17, %rd12, %rd16;473; CHECK-NEXT: subc.cc.s64 %rd18, %rd51, 0;474; CHECK-NEXT: setp.gt.u64 %p6, %rd17, 127;475; CHECK-NEXT: setp.eq.b64 %p7, %rd18, 0;476; CHECK-NEXT: and.pred %p8, %p7, %p6;477; CHECK-NEXT: setp.ne.b64 %p9, %rd18, 0;478; CHECK-NEXT: or.pred %p10, %p8, %p9;479; CHECK-NEXT: or.pred %p11, %p3, %p10;480; CHECK-NEXT: xor.b64 %rd19, %rd17, 127;481; CHECK-NEXT: or.b64 %rd20, %rd19, %rd18;482; CHECK-NEXT: setp.eq.b64 %p12, %rd20, 0;483; CHECK-NEXT: selp.b64 %rd59, 0, %rd4, %p11;484; CHECK-NEXT: selp.b64 %rd58, 0, %rd3, %p11;485; CHECK-NEXT: or.pred %p13, %p11, %p12;486; CHECK-NEXT: @%p13 bra $L__BB5_5;487; CHECK-NEXT: // %bb.3: // %udiv-bb1488; CHECK-NEXT: add.cc.s64 %rd52, %rd17, 1;489; CHECK-NEXT: addc.cc.s64 %rd53, %rd18, 0;490; CHECK-NEXT: or.b64 %rd21, %rd52, %rd53;491; CHECK-NEXT: setp.eq.b64 %p14, %rd21, 0;492; CHECK-NEXT: cvt.u32.u64 %r5, %rd17;493; CHECK-NEXT: sub.s32 %r6, 127, %r5;494; CHECK-NEXT: shl.b64 %rd22, %rd4, %r6;495; CHECK-NEXT: sub.s32 %r7, 64, %r6;496; CHECK-NEXT: shr.u64 %rd23, %rd3, %r7;497; CHECK-NEXT: or.b64 %rd24, %rd22, %rd23;498; CHECK-NEXT: sub.s32 %r8, 63, %r5;499; CHECK-NEXT: shl.b64 %rd25, %rd3, %r8;500; CHECK-NEXT: setp.gt.s32 %p15, %r6, 63;501; CHECK-NEXT: selp.b64 %rd57, %rd25, %rd24, %p15;502; CHECK-NEXT: shl.b64 %rd56, %rd3, %r6;503; CHECK-NEXT: mov.b64 %rd50, %rd51;504; CHECK-NEXT: @%p14 bra $L__BB5_4;505; CHECK-NEXT: // %bb.1: // %udiv-preheader506; CHECK-NEXT: cvt.u32.u64 %r9, %rd52;507; CHECK-NEXT: shr.u64 %rd26, %rd3, %r9;508; CHECK-NEXT: sub.s32 %r10, 64, %r9;509; CHECK-NEXT: shl.b64 %rd27, %rd4, %r10;510; CHECK-NEXT: or.b64 %rd28, %rd26, %rd27;511; CHECK-NEXT: add.s32 %r11, %r9, -64;512; CHECK-NEXT: shr.u64 %rd29, %rd4, %r11;513; CHECK-NEXT: setp.gt.s32 %p16, %r9, 63;514; CHECK-NEXT: selp.b64 %rd54, %rd29, %rd28, %p16;515; CHECK-NEXT: shr.u64 %rd55, %rd4, %r9;516; CHECK-NEXT: add.cc.s64 %rd1, %rd5, -1;517; CHECK-NEXT: addc.cc.s64 %rd2, %rd6, -1;518; CHECK-NEXT: mov.b64 %rd50, 0;519; CHECK-NEXT: mov.b64 %rd51, %rd50;520; CHECK-NEXT: $L__BB5_2: // %udiv-do-while521; CHECK-NEXT: // =>This Inner Loop Header: Depth=1522; CHECK-NEXT: shr.u64 %rd30, %rd54, 63;523; CHECK-NEXT: shl.b64 %rd31, %rd55, 1;524; CHECK-NEXT: or.b64 %rd32, %rd31, %rd30;525; CHECK-NEXT: shl.b64 %rd33, %rd54, 1;526; CHECK-NEXT: shr.u64 %rd34, %rd57, 63;527; CHECK-NEXT: or.b64 %rd35, %rd33, %rd34;528; CHECK-NEXT: shr.u64 %rd36, %rd56, 63;529; CHECK-NEXT: shl.b64 %rd37, %rd57, 1;530; CHECK-NEXT: or.b64 %rd38, %rd37, %rd36;531; CHECK-NEXT: shl.b64 %rd39, %rd56, 1;532; CHECK-NEXT: or.b64 %rd56, %rd51, %rd39;533; CHECK-NEXT: or.b64 %rd57, %rd50, %rd38;534; CHECK-NEXT: sub.cc.s64 %rd40, %rd1, %rd35;535; CHECK-NEXT: subc.cc.s64 %rd41, %rd2, %rd32;536; CHECK-NEXT: shr.s64 %rd42, %rd41, 63;537; CHECK-NEXT: and.b64 %rd51, %rd42, 1;538; CHECK-NEXT: and.b64 %rd43, %rd42, %rd5;539; CHECK-NEXT: and.b64 %rd44, %rd42, %rd6;540; CHECK-NEXT: sub.cc.s64 %rd54, %rd35, %rd43;541; CHECK-NEXT: subc.cc.s64 %rd55, %rd32, %rd44;542; CHECK-NEXT: add.cc.s64 %rd52, %rd52, -1;543; CHECK-NEXT: addc.cc.s64 %rd53, %rd53, -1;544; CHECK-NEXT: or.b64 %rd45, %rd52, %rd53;545; CHECK-NEXT: setp.eq.b64 %p17, %rd45, 0;546; CHECK-NEXT: @%p17 bra $L__BB5_4;547; CHECK-NEXT: bra.uni $L__BB5_2;548; CHECK-NEXT: $L__BB5_4: // %udiv-loop-exit549; CHECK-NEXT: shr.u64 %rd46, %rd56, 63;550; CHECK-NEXT: shl.b64 %rd47, %rd57, 1;551; CHECK-NEXT: or.b64 %rd48, %rd47, %rd46;552; CHECK-NEXT: shl.b64 %rd49, %rd56, 1;553; CHECK-NEXT: or.b64 %rd58, %rd51, %rd49;554; CHECK-NEXT: or.b64 %rd59, %rd50, %rd48;555; CHECK-NEXT: $L__BB5_5: // %udiv-end556; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd58, %rd59};557; CHECK-NEXT: ret;558 %div = udiv i128 %lhs, %rhs559 ret i128 %div560}561 562define i128 @sdiv_i128_pow2k(i128 %lhs) {563; CHECK-LABEL: sdiv_i128_pow2k(564; CHECK: {565; CHECK-NEXT: .reg .b64 %rd<11>;566; CHECK-EMPTY:567; CHECK-NEXT: // %bb.0:568; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [sdiv_i128_pow2k_param_0];569; CHECK-NEXT: shr.s64 %rd3, %rd2, 63;570; CHECK-NEXT: shr.u64 %rd4, %rd3, 31;571; CHECK-NEXT: add.cc.s64 %rd5, %rd1, %rd4;572; CHECK-NEXT: addc.cc.s64 %rd6, %rd2, 0;573; CHECK-NEXT: shl.b64 %rd7, %rd6, 31;574; CHECK-NEXT: shr.u64 %rd8, %rd5, 33;575; CHECK-NEXT: or.b64 %rd9, %rd8, %rd7;576; CHECK-NEXT: shr.s64 %rd10, %rd6, 33;577; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd9, %rd10};578; CHECK-NEXT: ret;579 %div = sdiv i128 %lhs, 8589934592580 ret i128 %div581}582 583define i128 @udiv_i128_pow2k(i128 %lhs) {584; CHECK-LABEL: udiv_i128_pow2k(585; CHECK: {586; CHECK-NEXT: .reg .b64 %rd<7>;587; CHECK-EMPTY:588; CHECK-NEXT: // %bb.0:589; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [udiv_i128_pow2k_param_0];590; CHECK-NEXT: shl.b64 %rd3, %rd2, 31;591; CHECK-NEXT: shr.u64 %rd4, %rd1, 33;592; CHECK-NEXT: or.b64 %rd5, %rd4, %rd3;593; CHECK-NEXT: shr.u64 %rd6, %rd2, 33;594; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd5, %rd6};595; CHECK-NEXT: ret;596 %div = udiv i128 %lhs, 8589934592597 ret i128 %div598}599 600define i128 @add_i128(i128 %lhs, i128 %rhs) {601; CHECK-LABEL: add_i128(602; CHECK: {603; CHECK-NEXT: .reg .b64 %rd<7>;604; CHECK-EMPTY:605; CHECK-NEXT: // %bb.0:606; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [add_i128_param_0];607; CHECK-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [add_i128_param_1];608; CHECK-NEXT: add.cc.s64 %rd5, %rd1, %rd3;609; CHECK-NEXT: addc.cc.s64 %rd6, %rd2, %rd4;610; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd5, %rd6};611; CHECK-NEXT: ret;612 %result = add i128 %lhs, %rhs613 ret i128 %result614}615