islpy 2025.2 - 2025.2.3 break doctests with current `main`

[matthiasdiener]

With islpy 2025.2 -- 2025.2.3:

$ pytest doc/tutorial.rst
========================================== test session starts ==========================================
platform darwin -- Python 3.11.9, pytest-8.4.0, pluggy-1.6.0
rootdir: /Users/mdiener/Work/emirge/loopy
configfile: pyproject.toml
plugins: github-actions-annotate-failures-0.2.0, anyio-4.9.0, mock-3.14.0, xdist-3.6.1
collected 1 item

doc/tutorial.rst F                                                                                [100%]

=============================================== FAILURES ================================================
________________________________________ [doctest] tutorial.rst _________________________________________
555
556 .. doctest::
557
558     >>> knl = lp.make_kernel(
559     ...     "{ [i]: 0<=i<n }",
560     ...     "a[i] = 0", assumptions="n>=1")
561     >>> knl = lp.split_iname(knl, "i", 16)
562     >>> knl = lp.prioritize_loops(knl, "i_outer,i_inner")
563     >>> knl = lp.set_options(knl, write_code=True)
564     >>> evt, (out,) = knl(queue, a=x_vec_dev)
Differences (unified diff with -expected +actual):
    @@ -1,6 +1,9 @@
     #define lid(N) ((int) get_local_id(N))
    -...
    +#define gid(N) ((int) get_group_id(N))
    +<BLANKLINE>
    +__kernel void __attribute__ ((reqd_work_group_size(1, 1, 1))) loopy_kernel(__global float *__restrict__ a, int const n)
    +{
       for (int i_outer = 0; i_outer <= -1 + (15 + n) / 16; ++i_outer)
    -    for (int i_inner = 0; i_inner <= ((-17 + n + -16 * i_outer >= 0) ? 15 : -1 + n + -16 * i_outer); ++i_inner)
    +    for (int i_inner = 0; i_inner <= ((-16 + n + -16 * i_outer >= 0) ? 15 : -1 + n + -16 * i_outer); ++i_inner)
           a[16 * i_outer + i_inner] = (float) (0.0f);
    -...
    +}

/Users/mdiener/Work/emirge/loopy/doc/tutorial.rst:564: DocTestFailure
======================================== short test summary info ========================================
FAILED doc/tutorial.rst::tutorial.rst
=========================================== 1 failed in 0.81s ===========================================

Note that:

• The only meaningful change is the change from -17 to -16 inside the kernel. I'm not sure whether this change affects the actual kernel execution.
• There are multiple doctest errors in this file, but only the first one is shown.

Islpy 2025.1.5 runs these tests fine.

[inducer]

I'm aware. I looked a those closely, they're equivalent. Fixes for this are coming in #937.

[matthiasdiener]

Thanks! I just noticed that it came up in inducer/cgen#54 as well.

Here is what ChatGPT says regarding the equivalence of both kernels:

---

The two OpenCL kernels you provided differ only in the expression inside the conditional of the second loop:

Key difference:

// Kernel A (original):
(-16 + n + -16 * i_outer >= 0) ? 15 : -1 + n + -16 * i_outer

// Kernel B (modified):
(-17 + n + -16 * i_outer >= 0) ? 15 : -1 + n + -16 * i_outer

The second kernel subtracts 1 more in the conditional test (-17 instead of -16).

The role of this expression

This ternary controls the upper bound of the i_inner loop. The goal of both kernels is to initialize elements of the array a[0..n-1] in chunks of 16 elements per outer iteration (i_outer). The upper bound limits how many elements are processed in the final chunk, if n is not a multiple of 16.

So are they equivalent?

They are functionally equivalent for all n >= 1. Here's why:

1. When -16 + n - 16*i_outer >= 0 is true, the conditional evaluates to 15 in both versions — both kernels do 16 iterations for i_inner.

2. When the value is false, the fallback is -1 + n - 16*i_outer — the correct remaining number of iterations.

The difference lies in the cutoff point:

Kernel A starts switching from 16 to "leftover" iterations once -16 + n - 16*i_outer < 0,
while Kernel B does so once -17 + n - 16*i_outer < 0.

That means Kernel B transitions to the "remainder" case one iteration later — but this does not result in a different number of iterations.

Let’s analyze the key transition:

Suppose n = 33. Then:

• Kernel A: (-16 + 33 - 16*i_outer) >= 0
• Kernel B: (-17 + 33 - 16*i_outer) >= 0

At i_outer = 2:

• Kernel A condition: 1 >= 0 → true → 16 iterations
• Kernel B condition: 0 >= 0 → true → 16 iterations

At i_outer = 3:

• Kernel A: -15 >= 0 → false → -1 + 33 - 48 = -16 → loop does 0 iterations
• Kernel B: -16 >= 0 → false → same -16 → 0 iterations

So both kernels yield the same behavior.

The difference of -16 vs. -17 in the conditional has no effect on the actual values produced by the ternary expression, due to the logic of the loop and how rounding in (-1 + (15 + n) / 16) is used to cap i_outer.

✅ Conclusion:

Yes, the two kernels are functionally equivalent for all n >= 1.

The slight difference in the conditional threshold does not affect the number of iterations or the final values written into a[].