chenmiaoming opened a new pull request, #20108:
URL: https://github.com/apache/tvm/pull/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`
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: [email protected]
For queries about this service, please contact Infrastructure at:
[email protected]
---------------------------------------------------------------------
To unsubscribe, e-mail: [email protected]
For additional commands, e-mail: [email protected]
