bujjibabukatta opened a new pull request, #19938: URL: https://github.com/apache/tvm/pull/19938
The GPU sort/topk kernels currently map nthread_bx to blockIdx.x and nthread_by to blockIdx.y. For large inputs (e.g. topk on a tensor of shape (60, 1000, 1000)), the product nthread_by * nbz can easily exceed 65535 which is the hardware limit for gridDim.y on CUDA, causing CUDA_ERROR_INVALID_VALUE at kernel launch. This patch fuses the two block dimensions into a single blockIdx.x axis (limit 2^31-1) and reconstructs virtual bx/by indices inside the kernel with modulo and integer division. All five call sites of _get_threads are updated consistently. Fixes #19549 -- 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]
