Sync with Microsoft ONNX Runtime - 20082026

cmake/CMakeLists.txt

@@ -1530,12 +1530,6 @@ if (onnxruntime_USE_CUDA)
     add_definitions("-DENABLE_FP4")
     message(STATUS "CUDA Toolkit version is greater or equal than 12.8, enable -DENABLE_FP4 flag")
     if(onnxruntime_USE_FP4_QMOE)
-      if(MSVC AND ("120" IN_LIST CMAKE_CUDA_ARCHITECTURES_ORIG OR
-                   "121" IN_LIST CMAKE_CUDA_ARCHITECTURES_ORIG))
-        message(FATAL_ERROR
-          "onnxruntime_USE_FP4_QMOE with SM120/SM121 is not supported on MSVC: native sm_120a "
-          "compilation fails in CUDA/CCCL tcgen05 headers. Use a non-MSVC build for native FP4 QMoE.")
-      endif()
       add_definitions("-DUSE_FP4_QMOE")
       message(STATUS "CUDA FP4 QMoE kernels enabled")
     endif()

cmake/onnxruntime_cuda_source_filters.cmake

@@ -155,18 +155,25 @@ endfunction()
 # Usage:
 #   onnxruntime_extract_llm_sources(<cu_src_list_var>
 #       LLM_SOURCES <output_var>
-#       LLM_SM90_SOURCES <output_var>)
+#       LLM_SM90_SOURCES <output_var>
+#       LLM_FP4_SOURCES <output_var>)
 function(onnxruntime_extract_llm_sources CU_SRC_LIST)
-  cmake_parse_arguments(PARSE_ARGV 1 _LLM "" "LLM_SOURCES;LLM_SM90_SOURCES" "")
+  cmake_parse_arguments(PARSE_ARGV 1 _LLM "" "LLM_SOURCES;LLM_SM90_SOURCES;LLM_FP4_SOURCES" "")
 
   set(_list "${${CU_SRC_LIST}}")
   set(_llm_srcs)
   set(_llm_sm90_srcs)
+  set(_llm_fp4_srcs)
   foreach(_src IN LISTS _list)
     if(_src MATCHES "/contrib_ops/cuda/llm/.*\\.cu$")
       # SM90-specific fpA_intB launchers (guarded by #ifndef EXCLUDE_SM_90)
       if(_src MATCHES "fpA_intB_gemm_launcher_[0-9]+\\.generated\\.cu$")
         list(APPEND _llm_sm90_srcs "${_src}")
+      elseif(onnxruntime_USE_FP4_QMOE AND
+             _src MATCHES "/moe_gemm/(moe_gemm_kernels_(bf16|fp16|fp4)_fp4|moe_kernels)\\.cu$")
+        # These units instantiate the native NVFP4 runner and activation conversion.
+        # Keep them separate so MSVC does not compile every LLM source at sm_120a.
+        list(APPEND _llm_fp4_srcs "${_src}")
       else()
         list(APPEND _llm_srcs "${_src}")
       endif()
@@ -178,10 +185,16 @@ function(onnxruntime_extract_llm_sources CU_SRC_LIST)
   if(_llm_sm90_srcs)
     list(REMOVE_ITEM _list ${_llm_sm90_srcs})
   endif()
+  if(_llm_fp4_srcs)
+    list(REMOVE_ITEM _list ${_llm_fp4_srcs})
+  endif()
 
   set("${CU_SRC_LIST}" "${_list}" PARENT_SCOPE)
   set("${_LLM_LLM_SOURCES}" "${_llm_srcs}" PARENT_SCOPE)
   set("${_LLM_LLM_SM90_SOURCES}" "${_llm_sm90_srcs}" PARENT_SCOPE)
+  if(_LLM_LLM_FP4_SOURCES)
+    set("${_LLM_LLM_FP4_SOURCES}" "${_llm_fp4_srcs}" PARENT_SCOPE)
+  endif()
 endfunction()
 
 # Filter CMAKE_CUDA_ARCHITECTURES to only those >= a minimum SM version.
@@ -193,13 +206,16 @@ endfunction()
 #       MIN_SM <number>
 #       [EXCLUDE_SM120_REAL])
 function(onnxruntime_filter_cuda_archs OUTPUT_VAR)
-  cmake_parse_arguments(PARSE_ARGV 1 _FCA "EXCLUDE_SM120_REAL" "MIN_SM" "")
+  cmake_parse_arguments(PARSE_ARGV 1 _FCA "EXCLUDE_SM120_REAL;REPLACE_SM120_REAL_WITH_VIRTUAL" "MIN_SM" "")
 
   set(_filtered)
   foreach(_arch IN LISTS CMAKE_CUDA_ARCHITECTURES)
     string(REGEX MATCH "^([0-9]+)" _arch_num "${_arch}")
     if(_arch_num GREATER_EQUAL "${_FCA_MIN_SM}")
-      if(_FCA_EXCLUDE_SM120_REAL AND _arch_num GREATER_EQUAL 120 AND _arch MATCHES "-real$")
+      if(_FCA_REPLACE_SM120_REAL_WITH_VIRTUAL AND _arch_num GREATER_EQUAL 120 AND _arch MATCHES "-real$")
+        list(APPEND _filtered "${_arch_num}-virtual")
+        continue()
+      elseif(_FCA_EXCLUDE_SM120_REAL AND _arch_num GREATER_EQUAL 120 AND _arch MATCHES "-real$")
         continue()
       endif()
       list(APPEND _filtered "${_arch}")

cmake/onnxruntime_providers_cuda.cmake

@@ -85,6 +85,7 @@
   onnxruntime_extract_llm_sources(onnxruntime_cuda_contrib_ops_cu_srcs
     LLM_SOURCES onnxruntime_cuda_llm_srcs
     LLM_SM90_SOURCES onnxruntime_cuda_llm_sm90_srcs
+    LLM_FP4_SOURCES onnxruntime_cuda_llm_fp4_srcs
   )
 
   # disable contrib ops conditionally
@@ -472,7 +473,8 @@
       endif()
     endif()
 
-    if(("120" IN_LIST CMAKE_CUDA_ARCHITECTURES_ORIG OR "121" IN_LIST CMAKE_CUDA_ARCHITECTURES_ORIG) AND NOT MSVC)
+    if(("120" IN_LIST CMAKE_CUDA_ARCHITECTURES_ORIG OR "121" IN_LIST CMAKE_CUDA_ARCHITECTURES_ORIG) AND
+       (NOT MSVC OR onnxruntime_USE_FP4_QMOE))
       target_compile_definitions(${target} PRIVATE COMPILE_BLACKWELL_SM120_TMA_GROUPED_GEMMS)
     endif()
 
@@ -585,9 +587,10 @@
         endif()
       endif()
 
-      # CUDA 13 generates host stubs with 128-byte aligned by-value CUTLASS parameters for these
-      # native SM120 TMA kernels. MSVC rejects those stubs with C2719, so retain the portable path.
-      if(onnxruntime_cuda_sm120_tma_srcs AND (NOT MSVC OR CMAKE_CUDA_COMPILER_VERSION VERSION_LESS 13.0))
+      # CUDA 13 gives CUtensorMap 128-byte host alignment when MSVC reports an accurate
+      # __cplusplus value. The target-specific host flag below avoids C2719 for FP4 QMoE.
+      if(onnxruntime_cuda_sm120_tma_srcs AND
+         (NOT MSVC OR CMAKE_CUDA_COMPILER_VERSION VERSION_LESS 13.0 OR onnxruntime_USE_FP4_QMOE))
         onnxruntime_filter_cuda_archs(_ort_sm120_cuda_architectures MIN_SM 120)
         if(_ort_sm120_cuda_architectures)
           onnxruntime_add_cuda_object_library(
@@ -596,6 +599,13 @@
             CUDA_ARCHITECTURES "${_ort_sm120_cuda_architectures}"
             NVCC_THREADS "${onnxruntime_NVCC_THREADS}"
             SOURCES ${onnxruntime_cuda_sm120_tma_srcs})
+          if(MSVC AND CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL 13.0)
+            # Keep CUDA's CUtensorMap payload unchanged while avoiding an
+            # alignas(128) by-value parameter that MSVC cannot represent in
+            # NVCC-generated host stubs (C2719).
+            target_compile_options(onnxruntime_providers_cuda_sm120_tma PRIVATE
+              "$<$<COMPILE_LANGUAGE:CUDA>:SHELL:-Xcompiler /Zc:__cplusplus->")
+          endif()
           target_compile_definitions(onnxruntime_providers_cuda PRIVATE ORT_ENABLE_BLOCKQUANT_SM120)
           if(TARGET onnxruntime_providers_cuda_obj)
             target_compile_definitions(onnxruntime_providers_cuda_obj PRIVATE ORT_ENABLE_BLOCKQUANT_SM120)
@@ -616,16 +626,13 @@
       # "no kernel image is available" (CUDA 209) failure when the arch list is real-only
       # (e.g. 86-real;120-real) and therefore carries no virtual compute_120 PTX fallback.
       #
-      # The one toolchain where native sm_120a does NOT compile is MSVC/Windows: targeting it
-      # pulls in CCCL tcgen05 PTX headers that fail with the MSVC host compiler. There we fall
-      # back to excluding SM120 real and rely on virtual compute_120 PTX + JIT instead. The NVFP4
-      # QMoE native FP4xFP4 path is an exception even on MSVC in principle: it emits `cvt.e2m1x2`
-      # in expandInputRowsKernel (moe_kernels.cu), which is valid only for real sm_120a and cannot
-      # be expressed in virtual PTX -- so when that feature is enabled we keep SM120 real archs
-      # regardless of compiler (NVFP4 QMoE is currently a non-MSVC configuration in practice).
+      # Native sm_120a pulls CCCL tcgen05 PTX headers that fail with the MSVC host compiler, so
+      # the broad LLM target uses virtual compute_120 PTX on Windows. FP4 QMoE is isolated below:
+      # its activation conversion requires real sm_120a for `cvt.e2m1x2` and its smaller source
+      # set does not pull in the problematic tcgen05 path.
       if(onnxruntime_cuda_llm_srcs)
-        if(MSVC AND NOT onnxruntime_USE_FP4_QMOE)
-          onnxruntime_filter_cuda_archs(_ort_llm_cuda_architectures MIN_SM 75 EXCLUDE_SM120_REAL)
+        if(MSVC)
+          onnxruntime_filter_cuda_archs(_ort_llm_cuda_architectures MIN_SM 75 REPLACE_SM120_REAL_WITH_VIRTUAL)
         else()
           onnxruntime_filter_cuda_archs(_ort_llm_cuda_architectures MIN_SM 75)
         endif()
@@ -638,6 +645,18 @@
             SOURCES ${onnxruntime_cuda_llm_srcs})
         endif()
       endif()
+
+      if(onnxruntime_cuda_llm_fp4_srcs)
+        onnxruntime_filter_cuda_archs(_ort_llm_fp4_cuda_architectures MIN_SM 75)
+        if(_ort_llm_fp4_cuda_architectures)
+          onnxruntime_add_cuda_object_library(
+            NAME onnxruntime_providers_cuda_llm_fp4
+            PARENT onnxruntime_providers_cuda
+            CUDA_ARCHITECTURES "${_ort_llm_fp4_cuda_architectures}"
+            NVCC_THREADS "${onnxruntime_NVCC_THREADS}"
+            SOURCES ${onnxruntime_cuda_llm_fp4_srcs})
+        endif()
+      endif()
     endif()
   endif()
 

cmake/onnxruntime_providers_cuda_plugin.cmake

@@ -122,6 +122,7 @@ onnxruntime_extract_flash_attention_sources(CUDA_PLUGIN_EP_CU_SRCS
 onnxruntime_extract_llm_sources(CUDA_PLUGIN_EP_CU_SRCS
   LLM_SOURCES _cuda_plugin_llm_srcs
   LLM_SM90_SOURCES _cuda_plugin_llm_sm90_srcs
+  LLM_FP4_SOURCES _cuda_plugin_llm_fp4_srcs
 )
 
 # Create shared library target using the ORT helper function for plugins
@@ -286,7 +287,8 @@ if(ORT_HAS_SM90_OR_LATER)
     target_compile_definitions(onnxruntime_providers_cuda_plugin PRIVATE COMPILE_HOPPER_TMA_GROUPED_GEMMS)
   endif()
 endif()
-if(("120" IN_LIST CMAKE_CUDA_ARCHITECTURES_ORIG OR "121" IN_LIST CMAKE_CUDA_ARCHITECTURES_ORIG) AND NOT MSVC)
+if(("120" IN_LIST CMAKE_CUDA_ARCHITECTURES_ORIG OR "121" IN_LIST CMAKE_CUDA_ARCHITECTURES_ORIG) AND
+   (NOT MSVC OR onnxruntime_USE_FP4_QMOE))
   target_compile_definitions(onnxruntime_providers_cuda_plugin PRIVATE COMPILE_BLACKWELL_SM120_TMA_GROUPED_GEMMS)
 endif()
 
@@ -342,9 +344,10 @@ if(NOT onnxruntime_DISABLE_CONTRIB_OPS)
     endif()
   endif()
 
-  # CUDA 13 generates host stubs with 128-byte aligned by-value CUTLASS parameters for these
-  # native SM120 TMA kernels. MSVC rejects those stubs with C2719, so retain the portable path.
-  if(_cuda_plugin_sm120_tma_srcs AND (NOT MSVC OR CMAKE_CUDA_COMPILER_VERSION VERSION_LESS 13.0))
+  # CUDA 13 gives CUtensorMap 128-byte host alignment when MSVC reports an accurate
+  # __cplusplus value. The target-specific host flag below avoids C2719 for FP4 QMoE.
+  if(_cuda_plugin_sm120_tma_srcs AND
+     (NOT MSVC OR CMAKE_CUDA_COMPILER_VERSION VERSION_LESS 13.0 OR onnxruntime_USE_FP4_QMOE))
     onnxruntime_filter_cuda_archs(_plugin_sm120_cuda_architectures MIN_SM 120)
     if(_plugin_sm120_cuda_architectures)
       onnxruntime_add_cuda_plugin_object_library(
@@ -354,22 +357,25 @@ if(NOT onnxruntime_DISABLE_CONTRIB_OPS)
         NVCC_THREADS "${onnxruntime_plugin_nvcc_threads}"
         COMPILE_OPTIONS ${_cuda_plugin_shared_compile_options}
         SOURCES ${_cuda_plugin_sm120_tma_srcs})
+      if(MSVC AND CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL 13.0)
+        # CUDA 13's cuda.h gives CUtensorMap alignas(128) when MSVC reports the
+        # accurate C++ language level. MSVC cannot pass that type by value in
+        # NVCC-generated host stubs (C2719). Keep NVCC at C++20, but use the
+        # legacy host __cplusplus value so CUtensorMap retains its 128-byte size
+        # with ordinary host alignment.
+        target_compile_options(onnxruntime_providers_cuda_plugin_sm120_tma PRIVATE
+          "$<$<COMPILE_LANGUAGE:CUDA>:SHELL:-Xcompiler /Zc:__cplusplus->")
+      endif()
       target_compile_definitions(onnxruntime_providers_cuda_plugin PRIVATE ORT_ENABLE_BLOCKQUANT_SM120)
     endif()
   endif()
 
   # LLM OBJECT library: SM75+ (backward compatible with fpA_intB_gemv/gemm which support SM75).
-  # FP4 QMoE requires native SM120-family SASS for activation quantization. Other MSVC builds
-  # exclude SM120+ real architectures because CCCL tcgen05 PTX headers fail with that host compiler.
+  # FP4 QMoE sources are compiled separately at native SM120 below. Keep the broad LLM target
+  # on virtual SM120 PTX under MSVC to avoid CCCL tcgen05 host-compile failures.
   if(_cuda_plugin_llm_srcs)
-    if(MSVC AND NOT onnxruntime_USE_FP4_QMOE)
-      onnxruntime_filter_cuda_archs(_plugin_llm_cuda_architectures MIN_SM 75 EXCLUDE_SM120_REAL)
-      # A native-only Windows ARM64 build has no lower architecture left after the
-      # MSVC SM120 exclusion. Emit PTX privately for this object library so its host
-      # launchers and device kernels are still linked into the plugin.
-      if(NOT _plugin_llm_cuda_architectures AND ORT_HAS_SM120_OR_LATER)
-        set(_plugin_llm_cuda_architectures "120-virtual")
-      endif()
+    if(MSVC)
+      onnxruntime_filter_cuda_archs(_plugin_llm_cuda_architectures MIN_SM 75 REPLACE_SM120_REAL_WITH_VIRTUAL)
     else()
       onnxruntime_filter_cuda_archs(_plugin_llm_cuda_architectures MIN_SM 75)
     endif()
@@ -383,6 +389,19 @@ if(NOT onnxruntime_DISABLE_CONTRIB_OPS)
         SOURCES ${_cuda_plugin_llm_srcs})
     endif()
   endif()
+
+  if(_cuda_plugin_llm_fp4_srcs)
+    onnxruntime_filter_cuda_archs(_plugin_llm_fp4_cuda_architectures MIN_SM 75)
+    if(_plugin_llm_fp4_cuda_architectures)
+      onnxruntime_add_cuda_plugin_object_library(
+        NAME onnxruntime_providers_cuda_plugin_llm_fp4
+        PARENT onnxruntime_providers_cuda_plugin
+        CUDA_ARCHITECTURES "${_plugin_llm_fp4_cuda_architectures}"
+        NVCC_THREADS "${onnxruntime_plugin_nvcc_threads}"
+        COMPILE_OPTIONS ${_cuda_plugin_shared_compile_options}
+        SOURCES ${_cuda_plugin_llm_fp4_srcs})
+    endif()
+  endif()
 endif()
 
 # --- Find cuDNN (may be at a custom path via onnxruntime_CUDNN_HOME) ---

cmake/onnxruntime_unittests.cmake

@@ -1023,6 +1023,9 @@ if (onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS AND NOT onnxruntime_BUILD_CUDA_EP_
   if(TARGET onnxruntime_providers_cuda_llm)
     target_link_libraries(onnxruntime_providers_cuda_ut PRIVATE onnxruntime_providers_cuda_llm)
   endif()
+  if(TARGET onnxruntime_providers_cuda_llm_fp4)
+    target_link_libraries(onnxruntime_providers_cuda_ut PRIVATE onnxruntime_providers_cuda_llm_fp4)
+  endif()
   if (MSVC)
     # Cutlass code has an issue with the following:
     # warning C4100: 'magic': unreferenced formal parameter