[Fix][TOPI] Fuse GPU sort grid dimensions to avoid CUDA launch failures on large inputs

[bujjibabukatta]

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

[gemini-code-assist]

Warning

Gemini encountered an error creating the review. You can try again by commenting /gemini review.

[tlopex]

The fused indexing itself looks correct for CUDA, and I verified that both (60, 1000, 1000) and (65536, 1000) execute successfully on CUDA. However, _get_threads is shared by non-CUDA GPU backends, while the assumption that blockIdx.x supports up to 2^31-1 is CUDA-specific.

For example, with a WebGPU target and shape (10000, 1000), main currently generates grid=(8, 10000, 1), but this PR generates grid=(80000, 1, 1). WebGPU limits every dispatch dimension to 65535, so this introduces a regression for a previously valid workload. Vulkan limits are also device-dependent and cannot generally be assumed to match CUDA.

Could we apply this fused mapping only to CUDA, while retaining an appropriate multidimensional mapping for WebGPU/Vulkan? Please also add regression coverage for both the CUDA overflow case and non-CUDA grid dimensions.