perf(moe): implement the fused decode-MoE kernel on CUDA

[inureyes]

Summary

The fused single-token decode-MoE kernel (#268) was Metal-only (mx.fast.metal_kernel), so on the CUDA backend it threw [metal_kernel] No Metal back-end and aborted for every small-expert MoE the dispatch selected (qwen2_moe, qwen3_moe, lfm2, qwen3_vl_moe). Larger experts fell through the dff > MLXCEL_FUSED_MOE_MAX_DFF break-even to gather_qmm and survived, so the failure tracked expert size, not architecture.

This PR ports the kernel to CUDA instead of disabling it, so the affected models run faster on CUDA rather than merely working.

What changed

Port both halves to mx.fast.cuda_kernel (the CUDA analogue of metal_kernel):

• Kernel A: gate/up GEMV + SwiGLU/GeGLU → act_g[K, Dff].
• Kernel B: down GEMV × score → summed over K.
• One warp owns each output row; simd_sum becomes a __shfl_down_sync warp reduction; the quantized-affine dequant (4/8-bit gate/up, 4/6/8-bit down) is unchanged.
• run_fused_moe_two_kernel selects the cuda_kernel port when metal::is_available() is false. no_cuda.cpp / no_metal.cpp stub the unused side, so both link on either backend.
• Precise expf in the SwiGLU holds greedy parity.

Only src/lib/mlxcel-core/cpp/mlx_cxx_kernels.cpp changes; fused_moe_enabled stays default-on for both backends.

Validation (GB10 / DGX Spark, CUDA 13.0)

Greedy byte-identical to the gather_qmm path on qwen3-30b-a3b.

Model	gather_qmm	fused (CUDA)	speedup
qwen3-moe-4bit	58.15	90.34	1.55x
qwen3.5-35b-a3b-4bit	45.84	64.69	1.41x
lfm2-8b-a1b-4bit	140.91	158.22	1.13x
qwen1.5-moe-a2.7b-4bit	113.03	124.73	1.10x

The previously-crashing MoE models now run at these fused rates by default (no env var).

Supersedes the earlier CUDA fused-off workaround. Relevant to the #307 fused-MoE wiring epic: the wired families now work on CUDA.

Closes #313

[inureyes]

Force-pushed: this PR now implements the CUDA fused decode-MoE kernel instead of the earlier default-off workaround. Porting the Metal kernel to mx.fast.cuda_kernel is greedy byte-identical and faster on GB10 (+10% to +55%), so fused stays default-on for both backends and the previously-crashing MoE models run faster rather than falling back to gather_qmm. See the updated description.

[inureyes]

Kept MLXCEL_FUSED_MOE_MAX_DFF at 4096 (shared default). The break-even probe shows the CUDA crossover is ~13-14k (Dff 768 +60%, 1792 +13%, 6400 +2%, 14336 -2%), so 4096 is conservative on CUDA, but the meaningful wins are all below it and the 4096-14k range gains at most ~2% (phi-3.5-moe) while mixtral slightly prefers gather_qmm. Documented the CUDA backend + per-backend break-even in the kernel design doc and environment-variables.md (cbcce6c).