onnxruntime
f1d07f0d - [Build] FP4 QMoE in Windows (#32163)

Commit
85 days ago
[Build] FP4 QMoE in Windows (#32163) ## 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. --------- Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com>
Author
Parents
Loading