llvm/test/CodeGen/NVPTX/shift-logic-cse.ll

; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
; RUN: llc < %s -mtriple=nvptx64 | FileCheck %s

; Three logic ops (and/or/xor) share %x and the same shift amount, with no
; extend. Commuting the shifts through the logic ops exposes a single shifted
; %x that is CSEd, replacing three shifts with one.
define void @mixed_logic_matches(i32 %x, ptr %p0, ptr %p1, ptr %p2) {
; CHECK-LABEL: mixed_logic_matches(
; CHECK:       {
; CHECK-NEXT:    .reg .b32 %r<6>;
; CHECK-NEXT:    .reg .b64 %rd<4>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b32 %r1, [mixed_logic_matches_param_0];
; CHECK-NEXT:    ld.param.b64 %rd1, [mixed_logic_matches_param_1];
; CHECK-NEXT:    shl.b32 %r2, %r1, 2;
; CHECK-NEXT:    and.b32 %r3, %r2, 20;
; CHECK-NEXT:    ld.param.b64 %rd2, [mixed_logic_matches_param_2];
; CHECK-NEXT:    st.b32 [%rd1], %r3;
; CHECK-NEXT:    ld.param.b64 %rd3, [mixed_logic_matches_param_3];
; CHECK-NEXT:    or.b32 %r4, %r2, 28;
; CHECK-NEXT:    st.b32 [%rd2], %r4;
; CHECK-NEXT:    xor.b32 %r5, %r2, 68;
; CHECK-NEXT:    st.b32 [%rd3], %r5;
; CHECK-NEXT:    ret;
  %and = and i32 %x, 5
  %and.shl = shl i32 %and, 2
  store i32 %and.shl, ptr %p0, align 4
  %or = or i32 %x, 7
  %or.shl = shl i32 %or, 2
  store i32 %or.shl, ptr %p1, align 4
  %xor = xor i32 %x, 17
  %xor.shl = shl i32 %xor, 2
  store i32 %xor.shl, ptr %p2, align 4
  ret void
}

; Matching zext shifts: shl (zext (and X, C)), S. The shared shl (zext X), S is
; exposed and CSEd.
define void @zext_and_matches(i32 %x, ptr %p0, ptr %p1) {
; CHECK-LABEL: zext_and_matches(
; CHECK:       {
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-NEXT:    .reg .b64 %rd<6>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b32 %r1, [zext_and_matches_param_0];
; CHECK-NEXT:    ld.param.b64 %rd1, [zext_and_matches_param_1];
; CHECK-NEXT:    ld.param.b64 %rd2, [zext_and_matches_param_2];
; CHECK-NEXT:    mul.wide.u32 %rd3, %r1, 8;
; CHECK-NEXT:    and.b64 %rd4, %rd3, 40;
; CHECK-NEXT:    st.b64 [%rd1], %rd4;
; CHECK-NEXT:    and.b64 %rd5, %rd3, 48;
; CHECK-NEXT:    st.b64 [%rd2], %rd5;
; CHECK-NEXT:    ret;
  %and0 = and i32 %x, 5
  %z0 = zext i32 %and0 to i64
  %s0 = shl i64 %z0, 3
  store i64 %s0, ptr %p0, align 8
  %and1 = and i32 %x, 6
  %z1 = zext i32 %and1 to i64
  %s1 = shl i64 %z1, 3
  store i64 %s1, ptr %p1, align 8
  ret void
}

; Matching sext shifts: shl (sext (or X, C)), S.
define void @sext_or_matches(i32 %x, ptr %p0, ptr %p1) {
; CHECK-LABEL: sext_or_matches(
; CHECK:       {
; CHECK-NEXT:    .reg .b32 %r<2>;
; CHECK-NEXT:    .reg .b64 %rd<6>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b32 %r1, [sext_or_matches_param_0];
; CHECK-NEXT:    ld.param.b64 %rd1, [sext_or_matches_param_1];
; CHECK-NEXT:    ld.param.b64 %rd2, [sext_or_matches_param_2];
; CHECK-NEXT:    mul.wide.s32 %rd3, %r1, 4;
; CHECK-NEXT:    or.b64 %rd4, %rd3, 20;
; CHECK-NEXT:    st.b64 [%rd1], %rd4;
; CHECK-NEXT:    or.b64 %rd5, %rd3, 24;
; CHECK-NEXT:    st.b64 [%rd2], %rd5;
; CHECK-NEXT:    ret;
  %or0 = or i32 %x, 5
  %e0 = sext i32 %or0 to i64
  %s0 = shl i64 %e0, 2
  store i64 %s0, ptr %p0, align 8
  %or1 = or i32 %x, 6
  %e1 = sext i32 %or1 to i64
  %s1 = shl i64 %e1, 2
  store i64 %s1, ptr %p1, align 8
  ret void
}

; The shift is implicit here: the getelementptr's element-size scaling lowers
; to shl during SelectionDAG, so ordinary address arithmetic exercises the same
; pattern. NVPTX does not fold the scaled index into the load, so the combine
; still applies and the shared index computation is CSEd.
define void @gep_index_matches(i32 %base, ptr %data, ptr %p0, ptr %p1) {
; CHECK-LABEL: gep_index_matches(
; CHECK:       {
; CHECK-NEXT:    .reg .b32 %r<4>;
; CHECK-NEXT:    .reg .b64 %rd<9>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b32 %r1, [gep_index_matches_param_0];
; CHECK-NEXT:    ld.param.b64 %rd1, [gep_index_matches_param_1];
; CHECK-NEXT:    ld.param.b64 %rd2, [gep_index_matches_param_2];
; CHECK-NEXT:    mul.wide.u32 %rd3, %r1, 4;
; CHECK-NEXT:    xor.b64 %rd4, %rd3, 20;
; CHECK-NEXT:    add.s64 %rd5, %rd1, %rd4;
; CHECK-NEXT:    ld.param.b64 %rd6, [gep_index_matches_param_3];
; CHECK-NEXT:    ld.b32 %r2, [%rd5];
; CHECK-NEXT:    st.b32 [%rd2], %r2;
; CHECK-NEXT:    xor.b64 %rd7, %rd3, 24;
; CHECK-NEXT:    add.s64 %rd8, %rd1, %rd7;
; CHECK-NEXT:    ld.b32 %r3, [%rd8];
; CHECK-NEXT:    st.b32 [%rd6], %r3;
; CHECK-NEXT:    ret;
  %x0 = xor i32 %base, 5
  %i0 = zext i32 %x0 to i64
  %a0 = getelementptr i32, ptr %data, i64 %i0
  %v0 = load i32, ptr %a0, align 4
  store i32 %v0, ptr %p0, align 4
  %x1 = xor i32 %base, 6
  %i1 = zext i32 %x1 to i64
  %a1 = getelementptr i32, ptr %data, i64 %i1
  %v1 = load i32, ptr %a1, align 4
  store i32 %v1, ptr %p1, align 4
  ret void
}

; Negative test: a lone shl of a logic op has no matching shift to share with,
; so it is left alone. Preserving it lets targets that support it fold the shift
; into an addressing mode; NVPTX simply keeps the single shift.
define i32 @isolated_no_match(i32 %x) {
; CHECK-LABEL: isolated_no_match(
; CHECK:       {
; CHECK-NEXT:    .reg .b32 %r<4>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b32 %r1, [isolated_no_match_param_0];
; CHECK-NEXT:    xor.b32 %r2, %r1, 5;
; CHECK-NEXT:    shl.b32 %r3, %r2, 2;
; CHECK-NEXT:    st.param.b32 [func_retval0], %r3;
; CHECK-NEXT:    ret;
  %xor = xor i32 %x, 5
  %shl = shl i32 %xor, 2
  ret i32 %shl
}

; Negative test: with a variable shift amount the folded constant shift cannot
; be constant-folded, so it would add a shift rather than remove one. The
; combine requires a constant shift amount and leaves these alone.
define void @variable_shift_amount(i32 %x, i32 %s, ptr %p0, ptr %p1) {
; CHECK-LABEL: variable_shift_amount(
; CHECK:       {
; CHECK-NEXT:    .reg .b32 %r<7>;
; CHECK-NEXT:    .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT:  // %bb.0:
; CHECK-NEXT:    ld.param.b32 %r1, [variable_shift_amount_param_0];
; CHECK-NEXT:    and.b32 %r2, %r1, 5;
; CHECK-NEXT:    ld.param.b32 %r3, [variable_shift_amount_param_1];
; CHECK-NEXT:    shl.b32 %r4, %r2, %r3;
; CHECK-NEXT:    ld.param.b64 %rd1, [variable_shift_amount_param_2];
; CHECK-NEXT:    st.b32 [%rd1], %r4;
; CHECK-NEXT:    ld.param.b64 %rd2, [variable_shift_amount_param_3];
; CHECK-NEXT:    or.b32 %r5, %r1, 6;
; CHECK-NEXT:    shl.b32 %r6, %r5, %r3;
; CHECK-NEXT:    st.b32 [%rd2], %r6;
; CHECK-NEXT:    ret;
  %and = and i32 %x, 5
  %and.shl = shl i32 %and, %s
  store i32 %and.shl, ptr %p0, align 4
  %or = or i32 %x, 6
  %or.shl = shl i32 %or, %s
  store i32 %or.shl, ptr %p1, align 4
  ret void
}