bujjibabukatta commented on code in PR #19921:
URL: https://github.com/apache/tvm/pull/19921#discussion_r3503995795
##########
python/tvm/topi/gpu/sort.py:
##########
@@ -30,9 +30,12 @@
def _get_threads(nthread_tx, nthread_bx, nthread_by):
+ # Swap blockIdx.x and blockIdx.y mappings.
+ # The batch dimension (nthread_by) can easily exceed the CUDA limit of
65535 for gridDim.y.
+ # By mapping nthread_by to blockIdx.x (limit ~2 billion), we prevent
CUDA_ERROR_INVALID_VALUE on large inputs.
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:
Addressed issues in commit f21b756 - Fuse block dimensions for CUDA topk to
prevent crashes on large sorted dimensions
--
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]
