[CUDA] port Metal rhs gather_qmm optimizations

[dhiltgen]

This ports the Metal rhs gather_qmm optimizations to CUDA, with one additional optimization and one bug fix. The optimization reduces tile size when expert groups are shorter than a 64-row tile (e.g. qwen3-coder 30B incremental +23% at p512, +6% @ 8k and up)

During testing I found a data corruption bug in the qmm_sm80_mainloop where the loop kept issuing cp.async past the last K tile, without waiting at the end for the tail.

Benchmarking:

mlx-lm, python -m mlx_lm benchmark, g128, main -> branch

GPU	Model (NVFP4)	prompt	main prompt tps	branch prompt tps	main gen tps	branch gen tps
RTX 5090	Qwen3-Coder-30B-A3B	512	1383	6625	217.0	217.4
RTX 5090	Qwen3-Coder-30B-A3B	2048	1390	11466	209.9	210.2
RTX 5090	Qwen3-Coder-30B-A3B	8192	1364	10622	192.9	193.2
RTX 5090	gemma-4-26b-a4b-it	512	1775	2321	207.9	207.8
RTX 5090	gemma-4-26b-a4b-it	2048	1782	2405	203.3	203.8
RTX 5090	gemma-4-26b-a4b-it	8192	1764	2374	193.1	193.5
RTX 5090	NVIDIA-Nemotron-3-Nano-30B-A3B	512	1819	2861	212.8	212.2
RTX 5090	NVIDIA-Nemotron-3-Nano-30B-A3B	2048	1895	3205	200.7	205.1
RTX 5090	NVIDIA-Nemotron-3-Nano-30B-A3B	8192	1892	3225	200.0	202.2
RTX 5090	Qwen3.6-35B-A3B	512	697	862	185.4	184.8
RTX 5090	Qwen3.6-35B-A3B	2048	623	796	184.6	183.1
RTX 6000 Ada	Qwen3-Coder-30B-A3B	512	723	5510	155.2	156.8
RTX 6000 Ada	Qwen3-Coder-30B-A3B	2048	692	7675	149.9	149.3
RTX 6000 Ada	Qwen3-Coder-30B-A3B	8192	646	6848	134.1	133.0
RTX 6000 Ada	gemma-4-26b-a4b-it	512	892	1250	133.4	133.5
RTX 6000 Ada	gemma-4-26b-a4b-it	2048	853	1233	131.2	130.6
RTX 6000 Ada	gemma-4-26b-a4b-it	8192	810	1173	124.7	124.1

Note: Qwen3-Coder-30B-A3B-Instruct, nvfp4 via mlx_lm convert

• ☑️ I understand it is strictly prohibited to use AI to write PR description
• AI usage disclosure: co-developed with coding agent