[all-commits] [llvm/llvm-project] 9ce63d: [AMDGPU][SCEV] Tighten max backedge-taken count fo...

[carlobertolli via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: 9ce63d09cff0bf07da9623b753010112db87b1d7
      https://github.com/llvm/llvm-project/commit/9ce63d09cff0bf07da9623b753010112db87b1d7
  Author: carlobertolli <carlo.bertolli at amd.com>
  Date:   2026-06-11 (Thu, 11 Jun 2026)

  Changed paths:
    M llvm/lib/Analysis/ScalarEvolution.cpp
    M llvm/test/Analysis/ScalarEvolution/shift-op.ll

  Log Message:
  -----------
  [AMDGPU][SCEV] Tighten max backedge-taken count for shift recurrence loops. (#197292)

Application code on AMDGPUs often uses the following pattern:

for (int i = blockDim.x / 2; i >= 1; i >>= 1)
  if (threadIdx.x < i) {
    <STMT based on threadIdx.x>
  }
  __syncthreads();
}

but this is currently not unrolled by LLVM for two reasons:
1. We overestimate the loop count with the number of bits in the
iteration variable (32).
2. Runtime unrolling is not applicable because of convergent
instructions.

For the specific pattern above, we can compute a tighter bound by using
range information, which for blockDim.x is (0,1025) and already emitted
by clang, resulting in 10 as upper bound as opposed to 32.
The next step in unrolling the loop is to adjust the threshold used to
control full loop unrolling with convergent instructions.

Assisted-by: Cursor



To unsubscribe from these emails, change your notification settings at https://github.com/llvm/llvm-project/settings/notifications