[Build] FP4 QMoE in Windows

[Tianlei Wu (tianleiwu)]

Description

Enable FP4 QMoE for Blackwell SM120 in Windows CUDA builds. Previously, MSVC builds rejected this configuration because compiling the full LLM CUDA source set for native SM120 exposed tcgen05 compilation failures.

Changes

• Split the FP4 QMoE sources that require native SM120 instructions into dedicated CUDA object targets:
  • onnxruntime_providers_cuda_llm_fp4
  • onnxruntime_providers_cuda_plugin_llm_fp4
• Compile those FP4 targets for native sm_120a (120-real) so expandInputRowsKernel can use cvt.e2m1x2.
• Compile the general LLM CUDA targets as compute_120 PTX (120-virtual) on MSVC, avoiding tcgen05 failures in unrelated translation units.
• Enable the CUTLASS SM120 grouped TMA kernels needed by FP4 QMoE on Windows.
• For CUDA 13 and MSVC, apply /Zc:__cplusplus- only to the host compilation of the SM120 TMA object target. This avoids MSVC C2719 errors in generated host stubs for by-value, 128-byte-aligned CUtensorMap parameters while preserving C++20 for CUDA device compilation.
• Apply the same source split and workaround to both the in-tree CUDA provider and the CUDA plugin provider.
• Update the FP4 QMoE documentation with the Windows build behavior and constraints.

Validation

• Built the native FP4 QMoE object target for sm_120a with MSVC and CUDA 13.
• Built the native SM120 CUTLASS TMA object target with the target-local host compiler workaround.
• The CUDA plugin packaging pipeline verified this change: the Windows x64 build passed.

[Copilot]

Pull request overview

Enables FP4 QMoE builds on Windows/MSVC by isolating native SM120 CUDA sources and applying a targeted host-alignment workaround.

Changes:

• Splits FP4 QMoE sources into dedicated object libraries.
• Uses virtual SM120 PTX for general MSVC LLM sources.
• Enables SM120 grouped-TMA kernels with /Zc:__cplusplus-.

Reviewed changes

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

Show a summary per file

File	Description
cmake/CMakeLists.txt	Removes the MSVC FP4 QMoE build prohibition.
cmake/onnxruntime_cuda_source_filters.cmake	Adds FP4 source extraction and architecture replacement.
cmake/onnxruntime_providers_cuda.cmake	Configures dedicated CUDA provider object targets.
cmake/onnxruntime_providers_cuda_plugin.cmake	Mirrors the configuration for the plugin provider.
docs/contrib_ops/cuda/moe_qmoe.md	Documents Windows SM120 build behavior.

---

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.