The GitHub Actions job "Lint" on tvm.git/main has succeeded. Run started by GitHub user tlopex (triggered by tlopex).
Head commit for run: 50c9ca7ee16a6e65cd47a6e55685181663859bef / MiaoMing Chen <[email protected]> [Fix][TOPI] Fuse GPU scan blocks to avoid CUDA gridDim.y overflow (#20108) ## Motivation Fixes #20106. The fallback TOPI GPU scan maps the per-row scan block to `blockIdx.x` and the batch to `blockIdx.y`. CUDA limits `gridDim.y` to 65,535, so a valid Relax `cumprod` with shape `(65536, 1)` fails at execution time with `CUDA_ERROR_INVALID_VALUE` and `grid=(1,65536,1)`. ## Changes Fuse the virtual scan-block and batch dimensions into a single `blockIdx.x` launch dimension, then recover both indices with integer division and modulo. Apply this mapping to the initial copy, up-sweep, and down-sweep stages. Fusing the dimensions fixes large batches without moving the same 65,535 limit onto long scan axes. Add a dynamic-shape CUDA `cumprod` regression with an input of shape `(65536, 3)`. The three-element scan axis exercises the up-sweep and down-sweep stages while retaining the batch size that previously caused the invalid launch. ## Testing - `python -m pytest tests/python/relax/test_backend_dispatch_sort_scan.py -xvs` - 10 passed, 1 skipped, 1 pre-existing xpassed - Manual CUDA checks for the issue shape `(65536, 1)` and a multi-block scan shape `(8, 2000)` - `pre-commit run --files python/tvm/topi/gpu/scan.py tests/python/relax/test_backend_dispatch_sort_scan.py` Report URL: https://github.com/apache/tvm/actions/runs/32316425714 With regards, GitHub Actions via GitBox --------------------------------------------------------------------- To unsubscribe, e-mail: [email protected] For additional commands, e-mail: [email protected]
