CUDA: use radix TOP_K when CUB DeviceTopK is unavailable

[Rhonstin]

Overview

CUDA builds with CCCL < 3.2 have no cub::DeviceTopK, so ggml_cuda_op_top_k falls back to a full argsort + copy at ~12 bytes of temp per cell. On long-context sparse-attention top-k that kills small cards: my 10 GB card (about 600 MB free) died in cuMemCreate on QSA top-k over ~200K cols x 128 rows. That's 190+ MB of temp per layer.

This PR reuses the radix selection from #27466 (merged for HIP) on CUDA when CUB_TOP_K_AVAILABLE is off. New GGML_CUDA_TOP_K_RADIX gate. Temp drops to O(nrows): per-row state plus nrows x blocks_per_row x 256 histogram ints, single-digit MBs. Same set out (ordered-bit radix; ties break either way in both paths, and downstream reads the indices as a set plus mask). Also gated next_power_of_2, unused on that path now.

Nothing changes when CUB_TOP_K_AVAILABLE is set (that branch is untouched and still wins), or on HIP (same guards, extended).

Additional information

What I checked:

• TU builds clean under nvcc 12.4 + CUB 2.5, radix branch taken, radix kernels in the object.
• With CCCL 3.3 headers the DeviceTopK branch still wins (MaxPairs present), radix compiled out.
• Full runtime of this exact branch on my box (10 GB cards, 200K ctx) still pending; kernels are unchanged from ROCm: add radix TOP_K for long rows #27466. Sibling data point: real DeviceTopK through newer CCCL headers on the same CUDA 12.4 toolchain gave +9% prefill (220.9 vs 202.0 tok/s) at 207K context against the argsort fallback.

Related: #27466.

Requirements

• I have read and agree with the contributing guidelines
• AI usage disclosure: YES - patch drafted with an AI coding assistant; all measurements come from real runs on my own hardware, and I reviewed and take responsibility for the change.

[ggml-gh-bot]

Hi @Rhonstin, thanks for your contribution!

Per our contribution guidelines, the automated PR checker found the following issue(s) that need your attention:

• PR Template not respected: Please respect the template when creating a new pull request. Make sure to fill out all required sections.

• AI-generated content: While code is allowed to be generated by AI, please write the PR description and commit messages on your own without the help of AI.

---

Please note that maintainers reserve the right to make final decisions on PRs. If you believe there is a mistake, please comment below.

[IMbackK]

sure why not, but this is silly: #if defined(GGML_CUDA_TOP_K_RADIX) || (!defined(GGML_CUDA_USE_CUB) && defined(GGML_USE_HIP)) could you refactor that into just GGML_CUDA_USE_TOP_K_RADIX that gets set if its cuda and no cub or hip or musa?

[Rhonstin]

Done - replaced with a single GGML_CUDA_USE_TOP_K_RADIX, set when CUB DeviceTopK is unavailable or on HIP without CUB (this also covers MUSA, where GGML_CUDA_USE_CUB is never defined so the first limb fires). Verified all four use sites are gone and compile-checked both configurations of top-k.cu: with CUB DeviceTopK (CCCL 3.x headers) and without (CUDA 12.4 toolchain headers). No behavior change on any config - the new macro is exactly the old condition.