Loop unroll criteria off by one

Bug Description

Bug Description

Warp decides whether to unroll a statically sized range() loop in Adjoint.get_unroll_range() (warp/_src/codegen.py:4104 on main):

# test if we're above max unroll count
max_iters = abs(end - start) // abs(step)

max_unroll = adj.builder_options.get("max_unroll", 16)
...
if max_iters > max_unroll:
    # generate a dynamic loop

abs(end - start) // abs(step) computes floor(span / step), but the actual trip count of range(start, end, step) is ceil(span / step), i.e. len(range(start, end, step)). The two disagree whenever step does not evenly divide the span, so the loop count used for the max_unroll decision is one lower than the real iteration count.

There are two distinct symptoms.

1. Loops that exceed max_unroll are unrolled anyway. A loop with exactly max_unroll + 1 real iterations is counted as max_unroll, so max_iters > max_unroll is false and the loop is fully unrolled despite exceeding the limit. With the default max_unroll = 16:

2. Empty ranges emit an unreachable dynamic loop. A range that is empty because the step points the wrong way is counted as non-empty, so Warp generates a dynamic loop whose body can never execute:

Reproduction

import warp as wp

@wp.kernel
def k(x: wp.array(dtype=float), y: wp.array(dtype=float)):
    tid = wp.tid()
    acc = float(0.0)
    for i in range(0, 33, 2):   # 17 iterations, default max_unroll = 16
        acc += x[tid] * float(i)
    y[tid] = acc

wp.load_module(device="cpu")
# Inspect the generated .cpp in wp.config.kernel_cache_dir:
# the loop is fully unrolled (no `start_for_` label), even though it has 17 iterations.

Suggested fix

constant_range = range(start, end, step)
max_iters = len(constant_range)
if max_iters == 0:
    return constant_range

len() on a range object is exact and O(1) for all three sign/step combinations, and the early return removes the dead dynamic loop for empty ranges.

Breaking change

Correcting the count is a user-visible behavior change, not a silent internal cleanup. Loops in the affected band (real trip count exactly max_unroll + 1, non-dividing step) stop being unrolled and become rolled loops.

Warp documents that non-unrolled loops have gradient limitations (docs/user_guide/differentiability.rst:1349), so a differentiable kernel with a boundary-sized stepped loop can go from correct gradients to incorrect gradients purely as a result of this fix. Measured with prod = prod * (1.0 + x[tid] * i * 0.01) and tape.backward():

range(0,29,2)  15 iterations, unrolled before and after (grad correct)
range(0,33,2)  17 iterations, unrolled before / rolled after
    got  [ 6.60 12.70 25.18 50.78]
    want [ 5.63  8.92 13.90 21.32]

The rolled result is not a new defect: it is bit-identical to what Warp already produces for a genuinely dynamic range(0, n, 2) loop, so the lowering is faithful. The fix simply moves these loops across the documented boundary.

Accordingly the fix should ship with:

• a changed (or fixed) changelog fragment per changelog/README.md, stating that max_unroll now counts the actual number of range() iterations, that stepped ranges which previously slipped past the limit are no longer unrolled, and that users who relied on unrolling should raise max_unroll; and
• a regression test pinning the boundary for both dividing and non-dividing steps.

System Information

No response