[CUDA] Add fused decode routing, single-pass finalize, and PDL overlap for NVFP4 QMoE

[Tianlei Wu (tianleiwu)]

Stacked on #33354 (and #33094). The base branch is tlwu/20261013/nvfp4-qmoe-decode-gemv, so the diff below is this change only. Once #33094 and #33354 merge, this PR will be retargeted to main. Opened as a draft until then.

Description

Optimizes the non-GEMV kernels and inter-kernel launch serialization along the offline K-packed (weights_prepacked=1) NVFP4 QMoE decode path on SM90+:

1. Single-CTA fused decode routing (FusedSoftmaxTopKMapsKernel): Replaces LaunchSoftmaxTopK + fusedBuildExpertMapsSortFirstToken in a single CTA launch for num_rows <= 4, num_experts == 512, and k in {1, 2, 4, 6, 8, 10}. One warp per row reproduces the SoftmaxTopKMergeKernel<T, 128, 4> reduction trees operation for operation and emits the (expert, token)-sorted expert maps bit-identically (ORT_FP4_GEMV_FUSED_ROUTE=0 disables).
2. Single-pass finalizeMoeRoutingKernel: Sizes the block to one thread per 128-bit column vector (up to 1024 threads instead of a fixed 256) so hidden dimensions up to 4096 (such as Qwen3.8's hidden=2560) reduce in a single pass, and issues batched loads of up to 16 expert rows together while preserving top-k accumulation order (7.7 us -> 4.0 us per layer on H200).
3. Programmatic dependent launch (PDL) overlap of FC1 -> FC2 -> finalize: Launches FC2 and finalizeMoeRoutingKernel with cudaLaunchAttributeProgrammaticStreamSerialization (ORT_FP4_GEMV_PDL=0 disables; ORT_FP4_GEMV_FC1_PDL=1 optionally enables PDL for FC1 behind fused routing). FC1 triggers dependents at start, FC2 prefetches its first batch of weights/scales before waiting on FC1's output, and finalizeMoeRoutingKernel stages routing metadata into shared memory before waiting on FC2's output (~2.2 us saved per layer on H200, preserved under CUDA Graph capture).

All three optimizations preserve exact floating-point evaluation order and produce bit-identical outputs.

Summary of Changes

File	Change
onnxruntime/contrib_ops/cuda/moe/qmoe_fused_route.cuh	New FusedSoftmaxTopKMapsKernel single-CTA fused softmax + top-k + expert-map kernel (E=512, rows <= 4, k in {1,2,4,6,8,10})
onnxruntime/contrib_ops/cuda/moe/qmoe_kernels.h / .cu	LaunchFusedRouteMaps declarations and launcher
onnxruntime/contrib_ops/cuda/llm/moe_gemm/moe_kernels.cu / moe_util_kernels.h	Single-pass finalizeMoeRoutingKernel (blockDim.x up to 1024, FINALIZE_LOAD_BATCH=16) and early_routing_loads PDL support
onnxruntime/contrib_ops/cuda/llm/moe_gemm/moe_gemv_fp4.h / .cu	griddepcontrol.wait / griddepcontrol.launch_dependents in MoeGemvFp4DecodeKernel and programmatic_dependent_launch attribute in LaunchMoeGemvFp4RawNPacked (ORT_FP4_GEMV_PDL)
onnxruntime/contrib_ops/cuda/moe/moe_quantization.cc	Wire LaunchFusedRouteMaps and fc2_finalize_pdl into QMoE::ComputeInternal
docs/contrib_ops/cuda/moe_qmoe.md	Document ORT_FP4_GEMV_FUSED_ROUTE, ORT_FP4_GEMV_PDL, and ORT_FP4_GEMV_FC1_PDL

Testing

Validated on NVIDIA H200 (sm_90a, CUDA 13.0):

• python -m pytest onnxruntime/test/python/transformers/test_qmoe_nvfp4_cuda.py: 178 passed, 8 skipped, 0 failed (including test_nvfp4_offline_row_major_decode_gemv_bit_exact for FP16 and BF16).
• End-to-end Qwen3.8-Flash-Next NVFP4 (v12b 3-session CUDA-graph + MTP engine, 48 MoE layers + 1 MTP layer):
  • Bit-exactness: engine_gate.py --prefill steps passes 6/6 with 100% identical token sequences (256/256 tokens on text2048 and prose2048) compared against the baseline without fused routing / PDL.
  • End-to-end decode speed (bench_engine_mtp.py, 256 generated tokens, median of 3 runs):

Context	Before (#33354 only) ms/cycle	With this PR ms/cycle	Cycle speedup	Decode throughput (tok/s)
2,048	8.274 ms	7.657 ms	-0.617 ms (-7.5%)	244.98 tok/s (was 227.64)
8,192	8.164 ms	7.836 ms	-0.328 ms (-4.0%)	232.49 tok/s (was 225.81)
32,768	8.352 ms	8.027 ms	-0.325 ms (-3.9%)	237.25 tok/s (was 229.97)

• lintrunner on all touched files: no issues.