236 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 -O1 | FileCheck %s3; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 -O1 | FileCheck %s4; RUN: %if ptxas-ptr32 %{ llc < %s -mtriple=nvptx -mcpu=sm_20 -O1 | %ptxas-verify %}5; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 -O1 | %ptxas-verify %}6 7define i32 @test1(i32 %n, i32 %m) {8;9; CHECK-LABEL: test1(10; CHECK: {11; CHECK-NEXT: .reg .b32 %r<4>;12; CHECK-EMPTY:13; CHECK-NEXT: // %bb.0:14; CHECK-NEXT: ld.param.b32 %r1, [test1_param_0];15; CHECK-NEXT: ld.param.b32 %r2, [test1_param_1];16; CHECK-NEXT: mad.lo.s32 %r3, %r2, %r1, %r2;17; CHECK-NEXT: st.param.b32 [func_retval0], %r3;18; CHECK-NEXT: ret;19 %add = add i32 %n, 120 %mul = mul i32 %add, %m21 ret i32 %mul22}23 24define i32 @test1_rev(i32 %n, i32 %m) {25;26; CHECK-LABEL: test1_rev(27; CHECK: {28; CHECK-NEXT: .reg .b32 %r<4>;29; CHECK-EMPTY:30; CHECK-NEXT: // %bb.0:31; CHECK-NEXT: ld.param.b32 %r1, [test1_rev_param_0];32; CHECK-NEXT: ld.param.b32 %r2, [test1_rev_param_1];33; CHECK-NEXT: mad.lo.s32 %r3, %r2, %r1, %r2;34; CHECK-NEXT: st.param.b32 [func_retval0], %r3;35; CHECK-NEXT: ret;36 %add = add i32 %n, 137 %mul = mul i32 %m, %add38 ret i32 %mul39}40 41; Transpose (mul (select)) if it can then be folded to mad42define i32 @test2(i32 %n, i32 %m, i32 %s) {43;44; CHECK-LABEL: test2(45; CHECK: {46; CHECK-NEXT: .reg .pred %p<2>;47; CHECK-NEXT: .reg .b32 %r<6>;48; CHECK-EMPTY:49; CHECK-NEXT: // %bb.0:50; CHECK-NEXT: ld.param.b32 %r1, [test2_param_0];51; CHECK-NEXT: ld.param.b32 %r2, [test2_param_1];52; CHECK-NEXT: ld.param.b32 %r3, [test2_param_2];53; CHECK-NEXT: setp.lt.s32 %p1, %r3, 1;54; CHECK-NEXT: mad.lo.s32 %r4, %r2, %r1, %r2;55; CHECK-NEXT: selp.b32 %r5, %r2, %r4, %p1;56; CHECK-NEXT: st.param.b32 [func_retval0], %r5;57; CHECK-NEXT: ret;58 %add = add i32 %n, 159 %cond = icmp slt i32 %s, 160 %sel = select i1 %cond, i32 1, i32 %add61 %mul = mul i32 %sel, %m62 ret i32 %mul63}64 65;; Transpose (mul (select)) if it can then be folded to mad66define i32 @test2_rev1(i32 %n, i32 %m, i32 %s) {67;68; CHECK-LABEL: test2_rev1(69; CHECK: {70; CHECK-NEXT: .reg .pred %p<2>;71; CHECK-NEXT: .reg .b32 %r<6>;72; CHECK-EMPTY:73; CHECK-NEXT: // %bb.0:74; CHECK-NEXT: ld.param.b32 %r1, [test2_rev1_param_0];75; CHECK-NEXT: ld.param.b32 %r2, [test2_rev1_param_1];76; CHECK-NEXT: ld.param.b32 %r3, [test2_rev1_param_2];77; CHECK-NEXT: setp.lt.s32 %p1, %r3, 1;78; CHECK-NEXT: mad.lo.s32 %r4, %r2, %r1, %r2;79; CHECK-NEXT: selp.b32 %r5, %r4, %r2, %p1;80; CHECK-NEXT: st.param.b32 [func_retval0], %r5;81; CHECK-NEXT: ret;82 %add = add i32 %n, 183 %cond = icmp slt i32 %s, 184 %sel = select i1 %cond, i32 %add, i32 185 %mul = mul i32 %sel, %m86 ret i32 %mul87}88 89;; Transpose (mul (select)) if it can then be folded to mad90define i32 @test2_rev2(i32 %n, i32 %m, i32 %s) {91;92; CHECK-LABEL: test2_rev2(93; CHECK: {94; CHECK-NEXT: .reg .pred %p<2>;95; CHECK-NEXT: .reg .b32 %r<6>;96; CHECK-EMPTY:97; CHECK-NEXT: // %bb.0:98; CHECK-NEXT: ld.param.b32 %r1, [test2_rev2_param_0];99; CHECK-NEXT: ld.param.b32 %r2, [test2_rev2_param_1];100; CHECK-NEXT: ld.param.b32 %r3, [test2_rev2_param_2];101; CHECK-NEXT: setp.lt.s32 %p1, %r3, 1;102; CHECK-NEXT: mad.lo.s32 %r4, %r2, %r1, %r2;103; CHECK-NEXT: selp.b32 %r5, %r4, %r2, %p1;104; CHECK-NEXT: st.param.b32 [func_retval0], %r5;105; CHECK-NEXT: ret;106 %add = add i32 %n, 1107 %cond = icmp slt i32 %s, 1108 %sel = select i1 %cond, i32 %add, i32 1109 %mul = mul i32 %m, %sel110 ret i32 %mul111}112 113;; Leave (mul (select)) intact if it transposing is not profitable114define i32 @test3(i32 %n, i32 %m, i32 %s) {115;116; CHECK-LABEL: test3(117; CHECK: {118; CHECK-NEXT: .reg .pred %p<2>;119; CHECK-NEXT: .reg .b32 %r<7>;120; CHECK-EMPTY:121; CHECK-NEXT: // %bb.0:122; CHECK-NEXT: ld.param.b32 %r1, [test3_param_0];123; CHECK-NEXT: add.s32 %r2, %r1, 3;124; CHECK-NEXT: ld.param.b32 %r3, [test3_param_1];125; CHECK-NEXT: ld.param.b32 %r4, [test3_param_2];126; CHECK-NEXT: setp.lt.s32 %p1, %r4, 1;127; CHECK-NEXT: selp.b32 %r5, 1, %r2, %p1;128; CHECK-NEXT: mul.lo.s32 %r6, %r5, %r3;129; CHECK-NEXT: st.param.b32 [func_retval0], %r6;130; CHECK-NEXT: ret;131 %add = add i32 %n, 3132 %cond = icmp slt i32 %s, 1133 %sel = select i1 %cond, i32 1, i32 %add134 %mul = mul i32 %sel, %m135 ret i32 %mul136}137 138;; (add (select 0, (mul a, b)), c) -> (select (mad a, b, c), c)139define i32 @test4(i32 %a, i32 %b, i32 %c, i1 %p) {140; CHECK-LABEL: test4(141; CHECK: {142; CHECK-NEXT: .reg .pred %p<2>;143; CHECK-NEXT: .reg .b16 %rs<3>;144; CHECK-NEXT: .reg .b32 %r<6>;145; CHECK-EMPTY:146; CHECK-NEXT: // %bb.0:147; CHECK-NEXT: ld.param.b8 %rs1, [test4_param_3];148; CHECK-NEXT: and.b16 %rs2, %rs1, 1;149; CHECK-NEXT: setp.ne.b16 %p1, %rs2, 0;150; CHECK-NEXT: ld.param.b32 %r1, [test4_param_0];151; CHECK-NEXT: ld.param.b32 %r2, [test4_param_1];152; CHECK-NEXT: ld.param.b32 %r3, [test4_param_2];153; CHECK-NEXT: mad.lo.s32 %r4, %r1, %r2, %r3;154; CHECK-NEXT: selp.b32 %r5, %r4, %r3, %p1;155; CHECK-NEXT: st.param.b32 [func_retval0], %r5;156; CHECK-NEXT: ret;157 %mul = mul i32 %a, %b158 %sel = select i1 %p, i32 %mul, i32 0159 %add = add i32 %c, %sel160 ret i32 %add161}162 163define i32 @test4_rev(i32 %a, i32 %b, i32 %c, i1 %p) {164; CHECK-LABEL: test4_rev(165; CHECK: {166; CHECK-NEXT: .reg .pred %p<2>;167; CHECK-NEXT: .reg .b16 %rs<3>;168; CHECK-NEXT: .reg .b32 %r<6>;169; CHECK-EMPTY:170; CHECK-NEXT: // %bb.0:171; CHECK-NEXT: ld.param.b8 %rs1, [test4_rev_param_3];172; CHECK-NEXT: and.b16 %rs2, %rs1, 1;173; CHECK-NEXT: setp.ne.b16 %p1, %rs2, 0;174; CHECK-NEXT: ld.param.b32 %r1, [test4_rev_param_0];175; CHECK-NEXT: ld.param.b32 %r2, [test4_rev_param_1];176; CHECK-NEXT: ld.param.b32 %r3, [test4_rev_param_2];177; CHECK-NEXT: mad.lo.s32 %r4, %r1, %r2, %r3;178; CHECK-NEXT: selp.b32 %r5, %r3, %r4, %p1;179; CHECK-NEXT: st.param.b32 [func_retval0], %r5;180; CHECK-NEXT: ret;181 %mul = mul i32 %a, %b182 %sel = select i1 %p, i32 0, i32 %mul183 %add = add i32 %c, %sel184 ret i32 %add185}186 187declare i32 @use(i32 %0, i32 %1)188 189define i32 @test_mad_multi_use(i32 %a, i32 %b, i32 %c) {190; CHECK-LABEL: test_mad_multi_use(191; CHECK: {192; CHECK-NEXT: .reg .b32 %r<7>;193; CHECK-EMPTY:194; CHECK-NEXT: // %bb.0:195; CHECK-NEXT: ld.param.b32 %r1, [test_mad_multi_use_param_0];196; CHECK-NEXT: ld.param.b32 %r2, [test_mad_multi_use_param_1];197; CHECK-NEXT: mul.lo.s32 %r3, %r1, %r2;198; CHECK-NEXT: ld.param.b32 %r4, [test_mad_multi_use_param_2];199; CHECK-NEXT: add.s32 %r5, %r3, %r4;200; CHECK-NEXT: { // callseq 0, 0201; CHECK-NEXT: .param .b32 param0;202; CHECK-NEXT: .param .b32 param1;203; CHECK-NEXT: .param .b32 retval0;204; CHECK-NEXT: st.param.b32 [param0], %r3;205; CHECK-NEXT: st.param.b32 [param1], %r5;206; CHECK-NEXT: call.uni (retval0), use, (param0, param1);207; CHECK-NEXT: ld.param.b32 %r6, [retval0];208; CHECK-NEXT: } // callseq 0209; CHECK-NEXT: st.param.b32 [func_retval0], %r6;210; CHECK-NEXT: ret;211 %mul = mul i32 %a, %b212 %add = add i32 %mul, %c213 %res = call i32 @use(i32 %mul, i32 %add)214 ret i32 %res215}216 217;; This case relies on mad x 1 y => add x y, previously we emit:218;; mad.lo.s32 %r3, %r1, 1, %r2;219define i32 @test_mad_fold(i32 %x) {220; CHECK-LABEL: test_mad_fold(221; CHECK: {222; CHECK-NEXT: .reg .b32 %r<7>;223; CHECK-EMPTY:224; CHECK-NEXT: // %bb.0:225; CHECK-NEXT: ld.param.b32 %r1, [test_mad_fold_param_0];226; CHECK-NEXT: mul.hi.s32 %r2, %r1, -2147221471;227; CHECK-NEXT: add.s32 %r3, %r2, %r1;228; CHECK-NEXT: shr.u32 %r4, %r3, 31;229; CHECK-NEXT: shr.s32 %r5, %r3, 12;230; CHECK-NEXT: add.s32 %r6, %r5, %r4;231; CHECK-NEXT: st.param.b32 [func_retval0], %r6;232; CHECK-NEXT: ret;233 %div = sdiv i32 %x, 8191234 ret i32 %div235}236