[CUDA] Add NVFP4 native FP4xFP4 prefill for QMoE on SM120

[Tianlei Wu (tianleiwu)]

Description

Adds a native block-scaled FP4×FP4 (W4A4) grouped-GEMM prefill path for the nvfp4 QMoE
quantization mode on Blackwell (SM120+). Previously nvfp4 always dequantized E2M1 weights to
FP16/BF16 and ran the dense A16 MoE runner; now prefill shapes route through the native CUTLASS
block-scaled tensor-op, while decode/small-M shapes stay on the fused GEMV / dequant fallback. Also
fixes the CUDA no kernel image (error 209) failure on SM120 by making the LLM object library build
native sm_120a SASS on Linux.

Summary of Changes

Native NVFP4 FP4×FP4 prefill (QMoE)

File	Change
onnxruntime/contrib_ops/cuda/moe/moe_quantization.cc	Enable native NVFP4 CUTLASS path on SM120+ when shape-supported; construct the FP4×FP4 runner plus a dense A16 fallback runner; wire NVFP4 block scales / global scales through QuantParams::FP4; route prefill (num_rows >= ORT_FP4_PREFILL_MIN_TOKENS) to native and decode to GEMV/fallback; profile the kFP4 tactic for native prefill. Replaces manual getenv parsing with ParseEnvironmentVariableWithDefault.
onnxruntime/contrib_ops/cuda/moe/moe_quantization.h	Add native-NVFP4 state (enable_nvfp4_cutlass_gemm_, fp4_prefill_min_tokens_, fp4_native_max_tokens_per_expert_, dense fallback runner).
onnxruntime/contrib_ops/cuda/llm/moe_gemm/moe_gemm_kernels_fp4_fp4.cu	New TU instantiating the FP4×FP4 grouped-GEMM template dispatch.
onnxruntime/contrib_ops/cuda/llm/moe_gemm/moe_kernels.cu	Support the FP4-weight × FP4-activation instantiation.

Build / SM120 native SASS

File	Change
cmake/onnxruntime_providers_cuda.cmake	MSVC-gate EXCLUDE_SM120_REAL on the LLM object library: Linux now builds native sm_120a SASS (fixes CUDA error 209 on real-only arch lists); MSVC keeps the virtual compute_120 PTX to dodge the CCCL tcgen05 host-compile failure; USE_FP4_QMOE keeps 120-real everywhere.
cmake/onnxruntime_cuda_source_filters.cmake	Include the new FP4×FP4 source in the CUDA build graph.

Routing controls (environment variables)

• ORT_ENABLE_NVFP4_CUTLASS_GEMM (default 1) — master switch for the native NVFP4 path.
• ORT_FP4_PREFILL_MIN_TOKENS (default 64) — prefill/decode routing threshold.
• ORT_FP4_NATIVE_MAX_TOKENS_PER_EXPERT (default 0 = no bound) — upper bound above which prefill uses the dense A16 fallback.

Tests & docs

File	Change
onnxruntime/test/python/transformers/test_qmoe_nvfp4_cuda.py	Add a native-prefill routing helper and apply a looser tolerance only to natively-routed shapes (W4A4 activation quantization), keeping strict bounds for decode/fallback shapes.
docs/contrib_ops/cuda/moe_qmoe.md	Document native prefill accuracy & routing (§9b.5), the LLM object library / SM120 native SASS arch filtering and the CUDA 209 gotcha (§14.2), plus a Limitations note and TOC entries.

Testing

• CUDA_VISIBLE_DEVICES=0 python -m pytest onnxruntime/test/python/transformers/test_qmoe_nvfp4_cuda.py → 18/18 pass on RTX 5060 Ti (SM120). Native-routed shapes (tokens 64/128) show max_diff ≈ 0.197; decode/GEMV shapes are near-exact.
• Native-arithmetic correctness proof: forcing the same shapes through the dequant fallback with ORT_FP4_PREFILL_MIN_TOKENS=100000 drops max_diff to ≈ 1e-3, confirming the ~0.2 error is legitimate W4A4 activation quantization, not a kernel bug.
• No regressions in the existing MoE/QMoE suites (test_moe_cuda.py, test_qmoe_cuda.py, test_qmoe_fp4_cuda.py).

Motivation and Context

Finish remaining works of #29697.

nvfp4 QMoE previously had no native execution path, so all shapes paid the dequant-to-A16 cost.
This PR enables the Blackwell block-scaled FP4×FP4 tensor-op for prefill while keeping the
latency-sensitive decode path on the fused GEMV, and resolves the SM120 no kernel image (CUDA 209)
failure that occurred when the LLM object library shipped no loadable SM120 image.

Checklist

• Tests added/updated
• No breaking changes (native path is opt-out via ORT_ENABLE_NVFP4_CUTLASS_GEMM=0)
• Documentation updated

[Copilot]

Pull request overview

Adds a native Blackwell (SM120+) NVFP4 QMoE FP4×FP4 block-scaled grouped-GEMM prefill path (with per-call routing between native/GEMV/dequant fallback), and updates CUDA build arch filtering to ensure SM120 has a loadable image (avoiding CUDA error 209). Updates tests and docs to reflect the new routing/accuracy characteristics and adds build/CI guidance.

Changes:

• Enable NVFP4 native FP4×FP4 CUTLASS prefill on SM120+ with routing knobs (ORT_ENABLE_NVFP4_CUTLASS_GEMM, ORT_FP4_PREFILL_MIN_TOKENS, ORT_FP4_NATIVE_MAX_TOKENS_PER_EXPERT) and dense A16 fallback runner.
• Add FP4×FP4 grouped-GEMM template instantiations/TUs and wire NVFP4 block scales through QuantParams::FP4.
• Update Python parity tests and CUDA/MoE QMoE documentation; adjust CUDA arch filtering to build native sm_120a SASS on Linux while keeping MSVC constraints.

Reviewed changes

Copilot reviewed 11 out of 11 changed files in this pull request and generated 1 comment.

Show a summary per file

File	Description
onnxruntime/test/python/transformers/test_qmoe_nvfp4_cuda.py	Adjust parity tolerances for native-prefill routing; add negative-model guards and GEMV-vs-fallback parity test.
onnxruntime/contrib_ops/cuda/moe/moe_quantization.h	Add NVFP4-native state toggles and dense fallback runner members.
onnxruntime/contrib_ops/cuda/moe/moe_quantization.cc	Implement NVFP4 SM120+ native FP4×FP4 runner + routing; plumb env vars; wire NVFP4 scales into QuantParams::FP4; update prepack logic.
onnxruntime/contrib_ops/cuda/llm/moe_gemm/moe_kernels.cu	Wire NVFP4 quant_params.fp4 stride setup and explicitly instantiate FP4×FP4 runner variants.
onnxruntime/contrib_ops/cuda/llm/moe_gemm/moe_gemm_kernels_fp4_fp4.cu	New TU for FP4×FP4 grouped-GEMM template instantiations (NVFP4 native path).
docs/OperatorKernels.md	Update QMoE type constraints to include NVFP4 scale tensor type (float8e4m3fn).
docs/ContribOperators.md	Document quant_type='nvfp4' and its required scale/global-scale inputs and types.
docs/contrib_ops/cuda/moe_qmoe.md	Document NVFP4 native prefill routing/accuracy and SM120 build/arch gotchas.
cmake/onnxruntime_providers_cuda.cmake	Adjust LLM object library arch filtering to include 120-real on Linux/non-MSVC and when FP4 QMoE requires real SM120 SASS.
cmake/onnxruntime_cuda_source_filters.cmake	Add new FP4×FP4 TU to FP4-QMoE filtering rules.
.agents/skills/ort-ci/SKILL.md	Clarify CI re-run guidance and doc-artifact download workflow.