The GitHub Actions job "Lint" on tvm.git/main has succeeded. Run started by GitHub user tlopex (triggered by tlopex).
Head commit for run: 46c37fc523a2ae751272cc131376633cb502908b / J.W (Jun) <[email protected]> [Fix][S-TIR][DLight] Guard non-affine reduction write-back (#20057) Fixes #20048. ## Motivation A valid Relax `conv2d` can fail during `relax.build(..., target="cuda")` when the DLight GPU Reduction rule selects its inner-spatial schedule. The rule fuses all spatial loops, but `_sch_inner_spatial` chooses its `threadIdx.x` tile from only the innermost spatial extent. For the reported shape, the non-unit spatial extents are `(2, 2, 20)` and the selected tile extent is 10. After `rfactor` and `reverse_compute_at`, the write-back block has to recover the original axes from the remaining fused loop. The resulting bindings are not quasi-affine, so `bind` rejects the block's compact-dataflow precondition. The same geometry can occur without mixed spatial/reduction indexing, for example in a plain reduction generated by `relax.op.sum`. Therefore the fix is based on the fused spatial geometry, rather than on a particular access pattern. ## Changes This PR adds an applicability check inside `Reduction`. It counts the non-unit spatial components that the write-back block must recover after fusion and rejects the inner-spatial schedule when that recovery would not remain quasi-affine. Unsupported reductions return `None`, allowing `GeneralReduction` or `Fallback` to handle them; affine cases continue to use the dedicated Reduction schedule. The check derives spatial domains from the normalized dominant-read access, which is the same order consumed by `_normalize`. Reordered or transposed accesses therefore use the same loop order for both applicability and scheduling. Thread-extent selection is shared with the scheduling path, and symbolic extents are handled conservatively. Unexpected `ScheduleError` exceptions are intentionally not caught by `_apply_rules`. `None` remains the contract for a non-applicable rule, while an exception from a rule remains visible as a scheduling bug instead of being silently converted into a fallback. ## Testing The changes were tested in a Python 3.12.13 environment with CUDA 13.0: ```bash python -m pytest \ tests/python/s_tir/dlight/test_gpu_reduction.py \ tests/python/s_tir/dlight/test_gpu_fallback.py -q ``` Result: `25 passed`. The regression tests cover the reported convolution-like access, a plain reduction without mixed spatial/reduction indices, reordered dominant-read access, an affine write-back shape, and propagation of an unexpected `ScheduleError` from a broken rule. The Relax reproduction from #20048 now builds successfully with `target="cuda"`, and the previously failing plain `relax.op.sum` shape is handled by a later schedule rule. Report URL: https://github.com/apache/tvm/actions/runs/33146364652 With regards, GitHub Actions via GitBox --------------------------------------------------------------------- To unsubscribe, e-mail: [email protected] For additional commands, e-mail: [email protected]
