246 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 %{ llc < %s -mcpu=sm_50 | %ptxas-verify %}4 5target triple = "nvptx64-nvidia-cuda"6 7define i32 @test_simple_rotl(i32 %x) {8; CHECK-LABEL: test_simple_rotl(9; CHECK: {10; CHECK-NEXT: .reg .b32 %r<3>;11; CHECK-EMPTY:12; CHECK-NEXT: // %bb.0:13; CHECK-NEXT: ld.param.b32 %r1, [test_simple_rotl_param_0];14; CHECK-NEXT: shf.l.wrap.b32 %r2, %r1, %r1, 7;15; CHECK-NEXT: st.param.b32 [func_retval0], %r2;16; CHECK-NEXT: ret;17 %shl = shl i32 %x, 718 %shr = lshr i32 %x, 2519 %add = add i32 %shl, %shr20 ret i32 %add21}22 23define i32 @test_simple_rotr(i32 %x) {24; CHECK-LABEL: test_simple_rotr(25; CHECK: {26; CHECK-NEXT: .reg .b32 %r<3>;27; CHECK-EMPTY:28; CHECK-NEXT: // %bb.0:29; CHECK-NEXT: ld.param.b32 %r1, [test_simple_rotr_param_0];30; CHECK-NEXT: shf.l.wrap.b32 %r2, %r1, %r1, 25;31; CHECK-NEXT: st.param.b32 [func_retval0], %r2;32; CHECK-NEXT: ret;33 %shr = lshr i32 %x, 734 %shl = shl i32 %x, 2535 %add = add i32 %shr, %shl36 ret i32 %add37}38 39define i32 @test_rotl_var(i32 %x, i32 %y) {40; CHECK-LABEL: test_rotl_var(41; CHECK: {42; CHECK-NEXT: .reg .b32 %r<4>;43; CHECK-EMPTY:44; CHECK-NEXT: // %bb.0:45; CHECK-NEXT: ld.param.b32 %r1, [test_rotl_var_param_0];46; CHECK-NEXT: ld.param.b32 %r2, [test_rotl_var_param_1];47; CHECK-NEXT: shf.l.wrap.b32 %r3, %r1, %r1, %r2;48; CHECK-NEXT: st.param.b32 [func_retval0], %r3;49; CHECK-NEXT: ret;50 %shl = shl i32 %x, %y51 %sub = sub i32 32, %y52 %shr = lshr i32 %x, %sub53 %add = add i32 %shl, %shr54 ret i32 %add55}56 57define i32 @test_rotr_var(i32 %x, i32 %y) {58; CHECK-LABEL: test_rotr_var(59; CHECK: {60; CHECK-NEXT: .reg .b32 %r<4>;61; CHECK-EMPTY:62; CHECK-NEXT: // %bb.0:63; CHECK-NEXT: ld.param.b32 %r1, [test_rotr_var_param_0];64; CHECK-NEXT: ld.param.b32 %r2, [test_rotr_var_param_1];65; CHECK-NEXT: shf.r.wrap.b32 %r3, %r1, %r1, %r2;66; CHECK-NEXT: st.param.b32 [func_retval0], %r3;67; CHECK-NEXT: ret;68 %shr = lshr i32 %x, %y69 %sub = sub i32 32, %y70 %shl = shl i32 %x, %sub71 %add = add i32 %shr, %shl72 ret i32 %add73}74 75define i32 @test_invalid_rotl_var_and(i32 %x, i32 %y) {76; CHECK-LABEL: test_invalid_rotl_var_and(77; CHECK: {78; CHECK-NEXT: .reg .b32 %r<8>;79; CHECK-EMPTY:80; CHECK-NEXT: // %bb.0:81; CHECK-NEXT: ld.param.b32 %r1, [test_invalid_rotl_var_and_param_0];82; CHECK-NEXT: ld.param.b32 %r2, [test_invalid_rotl_var_and_param_1];83; CHECK-NEXT: shl.b32 %r3, %r1, %r2;84; CHECK-NEXT: neg.s32 %r4, %r2;85; CHECK-NEXT: and.b32 %r5, %r4, 31;86; CHECK-NEXT: shr.u32 %r6, %r1, %r5;87; CHECK-NEXT: add.s32 %r7, %r6, %r3;88; CHECK-NEXT: st.param.b32 [func_retval0], %r7;89; CHECK-NEXT: ret;90 %shr = shl i32 %x, %y91 %sub = sub nsw i32 0, %y92 %and = and i32 %sub, 3193 %shl = lshr i32 %x, %and94 %add = add i32 %shl, %shr95 ret i32 %add96}97 98define i32 @test_invalid_rotr_var_and(i32 %x, i32 %y) {99; CHECK-LABEL: test_invalid_rotr_var_and(100; CHECK: {101; CHECK-NEXT: .reg .b32 %r<8>;102; CHECK-EMPTY:103; CHECK-NEXT: // %bb.0:104; CHECK-NEXT: ld.param.b32 %r1, [test_invalid_rotr_var_and_param_0];105; CHECK-NEXT: ld.param.b32 %r2, [test_invalid_rotr_var_and_param_1];106; CHECK-NEXT: shr.u32 %r3, %r1, %r2;107; CHECK-NEXT: neg.s32 %r4, %r2;108; CHECK-NEXT: and.b32 %r5, %r4, 31;109; CHECK-NEXT: shl.b32 %r6, %r1, %r5;110; CHECK-NEXT: add.s32 %r7, %r3, %r6;111; CHECK-NEXT: st.param.b32 [func_retval0], %r7;112; CHECK-NEXT: ret;113 %shr = lshr i32 %x, %y114 %sub = sub nsw i32 0, %y115 %and = and i32 %sub, 31116 %shl = shl i32 %x, %and117 %add = add i32 %shr, %shl118 ret i32 %add119}120 121define i32 @test_fshl_special_case(i32 %x0, i32 %x1, i32 %y) {122; CHECK-LABEL: test_fshl_special_case(123; CHECK: {124; CHECK-NEXT: .reg .b32 %r<5>;125; CHECK-EMPTY:126; CHECK-NEXT: // %bb.0:127; CHECK-NEXT: ld.param.b32 %r1, [test_fshl_special_case_param_0];128; CHECK-NEXT: ld.param.b32 %r2, [test_fshl_special_case_param_1];129; CHECK-NEXT: ld.param.b32 %r3, [test_fshl_special_case_param_2];130; CHECK-NEXT: shf.l.wrap.b32 %r4, %r2, %r1, %r3;131; CHECK-NEXT: st.param.b32 [func_retval0], %r4;132; CHECK-NEXT: ret;133 %shl = shl i32 %x0, %y134 %srli = lshr i32 %x1, 1135 %x = xor i32 %y, 31136 %srlo = lshr i32 %srli, %x137 %o = add i32 %shl, %srlo138 ret i32 %o139}140 141define i32 @test_fshr_special_case(i32 %x0, i32 %x1, i32 %y) {142; CHECK-LABEL: test_fshr_special_case(143; CHECK: {144; CHECK-NEXT: .reg .b32 %r<5>;145; CHECK-EMPTY:146; CHECK-NEXT: // %bb.0:147; CHECK-NEXT: ld.param.b32 %r1, [test_fshr_special_case_param_0];148; CHECK-NEXT: ld.param.b32 %r2, [test_fshr_special_case_param_1];149; CHECK-NEXT: ld.param.b32 %r3, [test_fshr_special_case_param_2];150; CHECK-NEXT: shf.r.wrap.b32 %r4, %r2, %r1, %r3;151; CHECK-NEXT: st.param.b32 [func_retval0], %r4;152; CHECK-NEXT: ret;153 %shl = lshr i32 %x1, %y154 %srli = shl i32 %x0, 1155 %x = xor i32 %y, 31156 %srlo = shl i32 %srli, %x157 %o = add i32 %shl, %srlo158 ret i32 %o159}160 161define i64 @test_rotl_udiv_special_case(i64 %i) {162; CHECK-LABEL: test_rotl_udiv_special_case(163; CHECK: {164; CHECK-NEXT: .reg .b32 %r<5>;165; CHECK-NEXT: .reg .b64 %rd<5>;166; CHECK-EMPTY:167; CHECK-NEXT: // %bb.0:168; CHECK-NEXT: ld.param.b64 %rd1, [test_rotl_udiv_special_case_param_0];169; CHECK-NEXT: mul.hi.u64 %rd2, %rd1, -6148914691236517205;170; CHECK-NEXT: shr.u64 %rd3, %rd2, 1;171; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd3;172; CHECK-NEXT: shf.l.wrap.b32 %r3, %r2, %r1, 28;173; CHECK-NEXT: shf.l.wrap.b32 %r4, %r1, %r2, 28;174; CHECK-NEXT: mov.b64 %rd4, {%r4, %r3};175; CHECK-NEXT: st.param.b64 [func_retval0], %rd4;176; CHECK-NEXT: ret;177 %lhs_div = udiv i64 %i, 3178 %rhs_div = udiv i64 %i, 48179 %lhs_shift = shl i64 %lhs_div, 60180 %out = add i64 %lhs_shift, %rhs_div181 ret i64 %out182}183 184define i32 @test_rotl_mul_special_case(i32 %i) {185; CHECK-LABEL: test_rotl_mul_special_case(186; CHECK: {187; CHECK-NEXT: .reg .b32 %r<4>;188; CHECK-EMPTY:189; CHECK-NEXT: // %bb.0:190; CHECK-NEXT: ld.param.b32 %r1, [test_rotl_mul_special_case_param_0];191; CHECK-NEXT: mul.lo.s32 %r2, %r1, 9;192; CHECK-NEXT: shf.l.wrap.b32 %r3, %r2, %r2, 7;193; CHECK-NEXT: st.param.b32 [func_retval0], %r3;194; CHECK-NEXT: ret;195 %lhs_mul = mul i32 %i, 9196 %rhs_mul = mul i32 %i, 1152197 %lhs_shift = lshr i32 %lhs_mul, 25198 %out = add i32 %lhs_shift, %rhs_mul199 ret i32 %out200}201 202define i64 @test_rotl_mul_with_mask_special_case(i64 %i) {203; CHECK-LABEL: test_rotl_mul_with_mask_special_case(204; CHECK: {205; CHECK-NEXT: .reg .b32 %r<7>;206; CHECK-NEXT: .reg .b64 %rd<5>;207; CHECK-EMPTY:208; CHECK-NEXT: // %bb.0:209; CHECK-NEXT: ld.param.b64 %rd1, [test_rotl_mul_with_mask_special_case_param_0];210; CHECK-NEXT: mul.lo.s64 %rd2, %rd1, 9;211; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd1;212; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd2;213; CHECK-NEXT: shf.l.wrap.b32 %r5, %r4, %r1, 7;214; CHECK-NEXT: shf.l.wrap.b32 %r6, %r1, %r2, 7;215; CHECK-NEXT: mov.b64 %rd3, {%r5, %r6};216; CHECK-NEXT: and.b64 %rd4, %rd3, 255;217; CHECK-NEXT: st.param.b64 [func_retval0], %rd4;218; CHECK-NEXT: ret;219 %lhs_mul = mul i64 %i, 1152220 %rhs_mul = mul i64 %i, 9221 %lhs_and = and i64 %lhs_mul, 160222 %rhs_shift = lshr i64 %rhs_mul, 57223 %out = add i64 %lhs_and, %rhs_shift224 ret i64 %out225}226 227define i32 @test_fshl_with_mask_special_case(i32 %x) {228; CHECK-LABEL: test_fshl_with_mask_special_case(229; CHECK: {230; CHECK-NEXT: .reg .b32 %r<5>;231; CHECK-EMPTY:232; CHECK-NEXT: // %bb.0:233; CHECK-NEXT: ld.param.b32 %r1, [test_fshl_with_mask_special_case_param_0];234; CHECK-NEXT: or.b32 %r2, %r1, 1;235; CHECK-NEXT: shf.l.wrap.b32 %r3, %r1, %r2, 5;236; CHECK-NEXT: and.b32 %r4, %r3, -31;237; CHECK-NEXT: st.param.b32 [func_retval0], %r4;238; CHECK-NEXT: ret;239 %or1 = or i32 %x, 1240 %sh1 = shl i32 %or1, 5241 %sh2 = lshr i32 %x, 27242 %1 = and i32 %sh2, 1243 %r = add i32 %sh1, %1244 ret i32 %r245}246