174 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mtriple=nvptx64 | FileCheck %s3 4; Fold: srl (or (x, shl(zext(y),c1)),c1) -> or(srl(x,c1), zext(y))5; c1 <= leadingzeros(zext(y))6define i64 @test_or(i64 %x, i32 %y) {7; CHECK-LABEL: test_or(8; CHECK: {9; CHECK-NEXT: .reg .b64 %rd<5>;10; CHECK-EMPTY:11; CHECK-NEXT: // %bb.0:12; CHECK-NEXT: ld.param.b64 %rd1, [test_or_param_0];13; CHECK-NEXT: ld.param.b32 %rd2, [test_or_param_1];14; CHECK-NEXT: shr.u64 %rd3, %rd1, 5;15; CHECK-NEXT: or.b64 %rd4, %rd3, %rd2;16; CHECK-NEXT: st.param.b64 [func_retval0], %rd4;17; CHECK-NEXT: ret;18 %ext = zext i32 %y to i6419 %shl = shl i64 %ext, 520 %or = or i64 %x, %shl21 %srl = lshr i64 %or, 522 ret i64 %srl23}24 25; Fold: srl (xor (x, shl(zext(y),c1)),c1) -> xor(srl(x,c1), zext(y))26; c1 <= leadingzeros(zext(y))27define i64 @test_xor(i64 %x, i32 %y) {28; CHECK-LABEL: test_xor(29; CHECK: {30; CHECK-NEXT: .reg .b64 %rd<5>;31; CHECK-EMPTY:32; CHECK-NEXT: // %bb.0:33; CHECK-NEXT: ld.param.b64 %rd1, [test_xor_param_0];34; CHECK-NEXT: ld.param.b32 %rd2, [test_xor_param_1];35; CHECK-NEXT: shr.u64 %rd3, %rd1, 5;36; CHECK-NEXT: xor.b64 %rd4, %rd3, %rd2;37; CHECK-NEXT: st.param.b64 [func_retval0], %rd4;38; CHECK-NEXT: ret;39 %ext = zext i32 %y to i6440 %shl = shl i64 %ext, 541 %or = xor i64 %x, %shl42 %srl = lshr i64 %or, 543 ret i64 %srl44}45 46; Fold: srl (and (x, shl(zext(y),c1)),c1) -> and(srl(x,c1), zext(y))47; c1 <= leadingzeros(zext(y))48define i64 @test_and(i64 %x, i32 %y) {49; CHECK-LABEL: test_and(50; CHECK: {51; CHECK-NEXT: .reg .b64 %rd<5>;52; CHECK-EMPTY:53; CHECK-NEXT: // %bb.0:54; CHECK-NEXT: ld.param.b64 %rd1, [test_and_param_0];55; CHECK-NEXT: ld.param.b32 %rd2, [test_and_param_1];56; CHECK-NEXT: shr.u64 %rd3, %rd1, 5;57; CHECK-NEXT: and.b64 %rd4, %rd3, %rd2;58; CHECK-NEXT: st.param.b64 [func_retval0], %rd4;59; CHECK-NEXT: ret;60 %ext = zext i32 %y to i6461 %shl = shl i64 %ext, 562 %or = and i64 %x, %shl63 %srl = lshr i64 %or, 564 ret i64 %srl65}66 67; Fold: srl (or (x, shl(zext(y),c1)),c1) -> or(srl(x,c1), zext(y))68; c1 <= leadingzeros(zext(y))69; x, y - vectors70define <2 x i16> @test_vec(<2 x i16> %x, <2 x i8> %y) {71; CHECK-LABEL: test_vec(72; CHECK: {73; CHECK-NEXT: .reg .b16 %rs<7>;74; CHECK-NEXT: .reg .b32 %r<4>;75; CHECK-EMPTY:76; CHECK-NEXT: // %bb.0:77; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_vec_param_0];78; CHECK-NEXT: ld.param.v2.b8 {%rs3, %rs4}, [test_vec_param_1];79; CHECK-NEXT: mov.b32 %r1, {%rs3, %rs4};80; CHECK-NEXT: shr.u16 %rs5, %rs2, 5;81; CHECK-NEXT: shr.u16 %rs6, %rs1, 5;82; CHECK-NEXT: mov.b32 %r2, {%rs6, %rs5};83; CHECK-NEXT: or.b32 %r3, %r2, %r1;84; CHECK-NEXT: st.param.b32 [func_retval0], %r3;85; CHECK-NEXT: ret;86 %ext = zext <2 x i8> %y to <2 x i16>87 %shl = shl <2 x i16> %ext, splat(i16 5)88 %or = or <2 x i16> %x, %shl89 %srl = lshr <2 x i16> %or, splat(i16 5)90 ret <2 x i16> %srl91}92 93; Do not fold: srl (or (x, shl(zext(y),c1)),c1) -> or(srl(x,c1), zext(y))94; Reason: c1 > leadingzeros(zext(y)).95define i64 @test_negative_c(i64 %x, i32 %y) {96; CHECK-LABEL: test_negative_c(97; CHECK: {98; CHECK-NEXT: .reg .b64 %rd<6>;99; CHECK-EMPTY:100; CHECK-NEXT: // %bb.0:101; CHECK-NEXT: ld.param.b64 %rd1, [test_negative_c_param_0];102; CHECK-NEXT: ld.param.b32 %rd2, [test_negative_c_param_1];103; CHECK-NEXT: shl.b64 %rd3, %rd2, 33;104; CHECK-NEXT: or.b64 %rd4, %rd1, %rd3;105; CHECK-NEXT: shr.u64 %rd5, %rd4, 33;106; CHECK-NEXT: st.param.b64 [func_retval0], %rd5;107; CHECK-NEXT: ret;108 %ext = zext i32 %y to i64109 %shl = shl i64 %ext, 33110 %or = or i64 %x, %shl111 %srl = lshr i64 %or, 33112 ret i64 %srl113}114 115declare void @use(i64)116 117; Do not fold: srl (or (x, shl(zext(y),c1)),c1) -> or(srl(x,c1), zext(y))118; Reason: multiple usage of "or"119define i64 @test_negative_use_lop(i64 %x, i32 %y) {120; CHECK-LABEL: test_negative_use_lop(121; CHECK: {122; CHECK-NEXT: .reg .b32 %r<2>;123; CHECK-NEXT: .reg .b64 %rd<5>;124; CHECK-EMPTY:125; CHECK-NEXT: // %bb.0:126; CHECK-NEXT: ld.param.b64 %rd1, [test_negative_use_lop_param_0];127; CHECK-NEXT: ld.param.b32 %r1, [test_negative_use_lop_param_1];128; CHECK-NEXT: mul.wide.u32 %rd2, %r1, 32;129; CHECK-NEXT: or.b64 %rd3, %rd1, %rd2;130; CHECK-NEXT: shr.u64 %rd4, %rd3, 5;131; CHECK-NEXT: { // callseq 0, 0132; CHECK-NEXT: .param .b64 param0;133; CHECK-NEXT: st.param.b64 [param0], %rd3;134; CHECK-NEXT: call.uni use, (param0);135; CHECK-NEXT: } // callseq 0136; CHECK-NEXT: st.param.b64 [func_retval0], %rd4;137; CHECK-NEXT: ret;138 %ext = zext i32 %y to i64139 %shl = shl i64 %ext, 5140 %or = or i64 %x, %shl141 %srl = lshr i64 %or, 5142 call void @use(i64 %or)143 ret i64 %srl144}145 146; Do not fold: srl (or (x, shl(zext(y),c1)),c1) -> or(srl(x,c1), zext(y))147; Reason: multiple usage of "shl"148define i64 @test_negative_use_shl(i64 %x, i32 %y) {149; CHECK-LABEL: test_negative_use_shl(150; CHECK: {151; CHECK-NEXT: .reg .b32 %r<2>;152; CHECK-NEXT: .reg .b64 %rd<5>;153; CHECK-EMPTY:154; CHECK-NEXT: // %bb.0:155; CHECK-NEXT: ld.param.b64 %rd1, [test_negative_use_shl_param_0];156; CHECK-NEXT: ld.param.b32 %r1, [test_negative_use_shl_param_1];157; CHECK-NEXT: mul.wide.u32 %rd2, %r1, 32;158; CHECK-NEXT: or.b64 %rd3, %rd1, %rd2;159; CHECK-NEXT: shr.u64 %rd4, %rd3, 5;160; CHECK-NEXT: { // callseq 1, 0161; CHECK-NEXT: .param .b64 param0;162; CHECK-NEXT: st.param.b64 [param0], %rd2;163; CHECK-NEXT: call.uni use, (param0);164; CHECK-NEXT: } // callseq 1165; CHECK-NEXT: st.param.b64 [func_retval0], %rd4;166; CHECK-NEXT: ret;167 %ext = zext i32 %y to i64168 %shl = shl i64 %ext, 5169 %or = or i64 %x, %shl170 %srl = lshr i64 %or, 5171 call void @use(i64 %shl)172 ret i64 %srl173}174