gemini-code-assist[bot] commented on code in PR #19918:
URL: https://github.com/apache/tvm/pull/19918#discussion_r3501660486
##########
python/tvm/topi/gpu/sort.py:
##########
@@ -31,8 +31,8 @@
def _get_threads(nthread_tx, nthread_bx, nthread_by):
tx = te.thread_axis("threadIdx.x")
- bx = te.thread_axis("blockIdx.x")
- by = te.thread_axis("blockIdx.y")
+ bx = te.thread_axis("blockIdx.y")
+ by = te.thread_axis("blockIdx.x")
Review Comment:
Mapping `bx` to `blockIdx.y` and `by` to `blockIdx.x` is highly
counter-intuitive and can easily lead to confusion or bugs during future
maintenance. Adding a clear inline comment explaining this swap and its purpose
(bypassing `gridDim.y` limits for large batch sizes) will greatly improve code
maintainability.
```suggestion
# Swap blockIdx.x and blockIdx.y to map the larger batch dimension
(nthread_by)
# to blockIdx.x (limit ~2B) and the smaller chunk dimension (nthread_bx)
to
# blockIdx.y (limit 65535) to avoid exceeding CUDA gridDim.y limits.
bx = te.thread_axis("blockIdx.y")
by = te.thread_axis("blockIdx.x")
```
--
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]
